Skip to content

[SPIRV] Don't add CPacked and Alignement decorations for Vulkan #138711

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
May 14, 2025
Merged

[SPIRV] Don't add CPacked and Alignement decorations for Vulkan #138711

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -319,7 +319,7 @@ bool SPIRVCallLowering::lowerFormalArguments(MachineIRBuilder &MIRBuilder,
buildOpDecorate(VRegs[i][0], MIRBuilder,
SPIRV::Decoration::MaxByteOffset, {DerefBytes});
}
if (Arg.hasAttribute(Attribute::Alignment)) {
if (Arg.hasAttribute(Attribute::Alignment) && !ST->isVulkanEnv()) {
auto Alignment = static_cast<unsigned>(
Arg.getAttribute(Attribute::Alignment).getValueAsInt());
buildOpDecorate(VRegs[i][0], MIRBuilder, SPIRV::Decoration::Alignment,
Expand Down
18 changes: 12 additions & 6 deletions llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -770,7 +770,7 @@ Register SPIRVGlobalRegistry::buildGlobalVariable(
if (IsConst && ST.isOpenCLEnv())
buildOpDecorate(Reg, MIRBuilder, SPIRV::Decoration::Constant, {});

if (GVar && GVar->getAlign().valueOrOne().value() != 1) {
if (GVar && GVar->getAlign().valueOrOne().value() != 1 && !ST.isVulkanEnv()) {
unsigned Alignment = (unsigned)GVar->getAlign().valueOrOne().value();
buildOpDecorate(Reg, MIRBuilder, SPIRV::Decoration::Alignment, {Alignment});
}
Expand Down Expand Up @@ -799,6 +799,9 @@ static std::string GetSpirvImageTypeName(const SPIRVType *Type,
const std::string &Prefix,
SPIRVGlobalRegistry &GR);

// Returns a name based on the Type. Notes that this does not look at
// decorations, and will return the same string for two types that are the same
// except for decorations.
static std::string buildSpirvTypeName(const SPIRVType *Type,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry &GR) {
Expand Down Expand Up @@ -885,9 +888,9 @@ Register SPIRVGlobalRegistry::getOrCreateGlobalVariableWithBinding(
Register VarReg =
MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::iIDRegClass);

// TODO: The name should come from the llvm-ir, but how that name will be
// passed from the HLSL to the backend has not been decided. Using this place
// holder for now.
// TODO(138533): The name should come from the llvm-ir, but how that name will
// be passed from the HLSL to the backend has not been decided. Using this
// place holder for now.
std::string Name =
("__resource_" + buildSpirvTypeName(VarType, MIRBuilder, *this) + "_" +
Twine(Set) + "_" + Twine(Binding))
Expand Down Expand Up @@ -955,6 +958,8 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeStruct(
const StructType *Ty, MachineIRBuilder &MIRBuilder,
SPIRV::AccessQualifier::AccessQualifier AccQual,
bool ExplicitLayoutRequired, bool EmitIR) {
const SPIRVSubtarget &ST =
cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
SmallVector<Register, 4> FieldTypes;
constexpr unsigned MaxWordCount = UINT16_MAX;
const size_t NumElements = Ty->getNumElements();
Expand All @@ -977,7 +982,7 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeStruct(
Register ResVReg = createTypeVReg(MIRBuilder);
if (Ty->hasName())
buildOpName(ResVReg, Ty->getName(), MIRBuilder);
if (Ty->isPacked())
if (Ty->isPacked() && !ST.isVulkanEnv())
buildOpDecorate(ResVReg, MIRBuilder, SPIRV::Decoration::CPacked, {});

SPIRVType *SPVType =
Expand Down Expand Up @@ -1629,7 +1634,8 @@ SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVTypeByName(
// Unable to recognize SPIRV type name
return nullptr;

auto SpirvTy = getOrCreateSPIRVType(Ty, MIRBuilder, AQ, false, true);
const SPIRVType *SpirvTy =
getOrCreateSPIRVType(Ty, MIRBuilder, AQ, false, true);

// Handle "type*" or "type* vector[N]".
if (TypeStr.starts_with("*")) {
Expand Down
76 changes: 76 additions & 0 deletions llvm/test/CodeGen/SPIRV/global-var-name-align.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,76 @@
; Check names and decoration of global variables.

; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s -check-prefixes=CHECK,OCL
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}

; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s -check-prefixes=CHECK,OCL
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}

; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-unknown %s -o - | FileCheck %s -check-prefixes=CHECK,VK
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-unknown %s -o - -filetype=obj | spirv-val %}

; CHECK-DAG: OpName %[[#id18:]] "G1"
; CHECK-DAG: OpName %[[#id22:]] "g1"
; CHECK-DAG: OpName %[[#id23:]] "g2"
; CHECK-DAG: OpName %[[#id27:]] "g4"
; CHECK-DAG: OpName %[[#id30:]] "c1"
; CHECK-DAG: OpName %[[#id31:]] "n_t"
; CHECK-DAG: OpName %[[#id32:]] "w"
; CHECK-DAG: OpName %[[#id34:]] "a.b"
; CHECK-DAG: OpName %[[#id35:]] "e"
; CHECK-DAG: OpName %[[#id36:]] "y.z"
; CHECK-DAG: OpName %[[#id38:]] "x"

; CHECK-NOT: OpDecorate %[[#id18]] LinkageAttributes
; OCL-DAG: OpDecorate %[[#id18]] Constant
; OCL-DAG: OpDecorate %[[#id22]] Alignment 4
; VK-NOT: OpDecorate {{.*}} Constant
; VK-NOT: OpDecorate {{.*}} Alignment
; CHECK-DAG: OpDecorate %[[#id22]] LinkageAttributes "g1" Export
; OCL-DAG: OpDecorate %[[#id23]] Alignment 4
; OCL-DAG: OpDecorate %[[#id27]] Alignment 4
; VK-NOT: OpDecorate {{.*}} Constant
; VK-NOT: OpDecorate {{.*}} Alignment
; CHECK-DAG: OpDecorate %[[#id27]] LinkageAttributes "g4" Export
; OCL-DAG: OpDecorate %[[#id30]] Constant
; OCL-DAG: OpDecorate %[[#id30]] Alignment 4
; VK-NOT: OpDecorate {{.*}} Constant
; VK-NOT: OpDecorate {{.*}} Alignment
; CHECK-DAG: OpDecorate %[[#id30]] LinkageAttributes "c1" Export
; OCL-DAG: OpDecorate %[[#id31]] Constant
; VK-NOT: OpDecorate {{.*}} Constant
; VK-NOT: OpDecorate {{.*}} Alignment
; CHECK-DAG: OpDecorate %[[#id31]] LinkageAttributes "n_t" Import
; OCL-DAG: OpDecorate %[[#id32]] Constant
; OCL-DAG: OpDecorate %[[#id32]] Alignment 4
; VK-NOT: OpDecorate {{.*}} Constant
; VK-NOT: OpDecorate {{.*}} Alignment
; CHECK-DAG: OpDecorate %[[#id32]] LinkageAttributes "w" Export
; OCL-DAG: OpDecorate %[[#id34]] Constant
; OCL-DAG: OpDecorate %[[#id34]] Alignment 4
; VK-NOT: OpDecorate {{.*}} Constant
; VK-NOT: OpDecorate {{.*}} Alignment
; CHECK-DAG: OpDecorate %[[#id35]] LinkageAttributes "e" Import
; OCL-DAG: OpDecorate %[[#id36]] Alignment 4
; OCL-DAG: OpDecorate %[[#id38]] Constant
; OCL-DAG: OpDecorate %[[#id38]] Alignment 4
; VK-NOT: OpDecorate {{.*}} Constant
; VK-NOT: OpDecorate {{.*}} Alignment

%"class.sycl::_V1::nd_item" = type { i8 }

@G1 = private unnamed_addr addrspace(1) constant %"class.sycl::_V1::nd_item" poison, align 1
@g1 = addrspace(1) global i32 1, align 4
@g2 = internal addrspace(1) global i32 2, align 4
@g4 = common addrspace(1) global i32 0, align 4
@c1 = addrspace(2) constant [2 x i32] [i32 0, i32 1], align 4
@n_t = external addrspace(2) constant [256 x i32]
@w = addrspace(1) constant i32 0, align 4
@a.b = internal addrspace(2) constant [2 x i32] [i32 2, i32 3], align 4
@e = external addrspace(1) global i32
@y.z = internal addrspace(1) global i32 0, align 4
@x = internal addrspace(2) constant float 1.000000e+00, align 4

define internal spir_func void @foo() {
ret void
}
37 changes: 37 additions & 0 deletions llvm/test/CodeGen/SPIRV/hlsl-resources/Packed.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,37 @@
; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv1.6-vulkan1.3-library %s -o - | FileCheck %s
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv1.6-vulkan1.3-library %s -o - -filetype=obj | spirv-val %}

target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"

; CHECK-DAG: OpName [[unpacked:%[0-9]+]] "unpacked"
; CHECK-DAG: OpName [[packed:%[0-9]+]] "packed"

; CHECK-NOT: OpDecorate {{.*}} CPacked
; CHECK-DAG: OpMemberDecorate [[unpacked]] 0 Offset 0
; CHECK-DAG: OpMemberDecorate [[unpacked]] 1 Offset 16

; CHECK-NOT: OpDecorate {{.*}} CPacked
; CHECK-DAG: OpMemberDecorate [[packed]] 0 Offset 0
; CHECK-DAG: OpMemberDecorate [[packed]] 1 Offset 4
; CHECK-NOT: OpDecorate {{.*}} CPacked


%unpacked = type {i32, <3 x i32>}
%packed = type <{i32, <3 x i32>}>


define external i32 @unpacked_vulkan_buffer_load() {
entry:
%handle = tail call target("spirv.VulkanBuffer", [0 x %unpacked], 12, 0) @llvm.spv.resource.handlefrombinding(i32 0, i32 0, i32 1, i32 0, i1 false)
%0 = tail call noundef nonnull align 4 dereferenceable(4) ptr addrspace(11) @llvm.spv.resource.getpointer(target("spirv.VulkanBuffer", [0 x %unpacked], 12, 0) %handle, i32 1)
%1 = load i32, ptr addrspace(11) %0, align 4
ret i32 %1
}

define external i32 @packed_vulkan_buffer_load() {
entry:
%handle = tail call target("spirv.VulkanBuffer", [0 x %packed], 12, 0) @llvm.spv.resource.handlefrombinding(i32 0, i32 1, i32 1, i32 0, i1 false)
%0 = tail call noundef nonnull align 4 dereferenceable(4) ptr addrspace(11) @llvm.spv.resource.getpointer(target("spirv.VulkanBuffer", [0 x %packed], 12, 0) %handle, i32 1)
%1 = load i32, ptr addrspace(11) %0, align 4
ret i32 %1
}
Loading