[SPIRV] Don't add CPacked and Alignement decorations for Vulkan #138711
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Packed structs get a different layout than a struct that is not packed. This is handled by assigning different offset decoration when appropriate. The `CPacked` decoration is not required, and is in fact not valid when creating a shader. Similaraly the alignment decoration is not allows when creating a shader. We must avoid generating that decoration. Fixes llvm#138268
|
@llvm/pr-subscribers-backend-spir-v Author: Steven Perron (s-perron) ChangesPacked structs get a different layout than a struct that is not packed. Similaraly the alignment decoration is not allows when creating a Fixes #138268 Full diff: https://github.com/llvm/llvm-project/pull/138711.diff 4 Files Affected:
diff --git a/llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp b/llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
index b824b9aeda660..5991a9af6364d 100644
--- a/llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
@@ -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,
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 35ddb906c366a..536424b63ec4f 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -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});
}
@@ -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) {
@@ -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))
@@ -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();
@@ -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 =
@@ -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("*")) {
diff --git a/llvm/test/CodeGen/SPIRV/global-var-name-align.ll b/llvm/test/CodeGen/SPIRV/global-var-name-align.ll
new file mode 100644
index 0000000000000..d73c98e55b872
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/global-var-name-align.ll
@@ -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
+}
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-resources/Packed.ll b/llvm/test/CodeGen/SPIRV/hlsl-resources/Packed.ll
new file mode 100644
index 0000000000000..d5f6545180147
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/hlsl-resources/Packed.ll
@@ -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
+}
|
cassiebeckley
approved these changes
May 6, 2025
Keenuts
approved these changes
May 7, 2025
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Packed structs get a different layout than a struct that is not packed.
This is handled by assigning different offset decoration when
appropriate. The
CPackeddecoration is not required, and is in factnot valid when creating a shader.
Similaraly the alignment decoration is not allows when creating a
shader. We must avoid generating that decoration.
Fixes #138268