[llvm] 8732d0f - [SPIRV] Don't add CPacked and Alignement decorations for Vulkan (#138711)
via llvm-commits
llvm-commits at lists.llvm.org
Wed May 14 06:18:19 PDT 2025
Author: Steven Perron
Date: 2025-05-14T09:18:15-04:00
New Revision: 8732d0f38961687b8d2f588d35f9164784660862
URL: https://github.com/llvm/llvm-project/commit/8732d0f38961687b8d2f588d35f9164784660862
DIFF: https://github.com/llvm/llvm-project/commit/8732d0f38961687b8d2f588d35f9164784660862.diff
LOG: [SPIRV] Don't add CPacked and Alignement decorations for Vulkan (#138711)
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 https://github.com/llvm/llvm-project/issues/138268
Added:
llvm/test/CodeGen/SPIRV/global-var-name-align.ll
llvm/test/CodeGen/SPIRV/hlsl-resources/Packed.ll
Modified:
llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
Removed:
################################################################################
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 251828b6bc35b..ac397fc486e19 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.consume_front("*"))
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 }
+
+ at G1 = private unnamed_addr addrspace(1) constant %"class.sycl::_V1::nd_item" poison, align 1
+ at g1 = addrspace(1) global i32 1, align 4
+ at g2 = internal addrspace(1) global i32 2, align 4
+ at g4 = common addrspace(1) global i32 0, align 4
+ at c1 = addrspace(2) constant [2 x i32] [i32 0, i32 1], align 4
+ at n_t = external addrspace(2) constant [256 x i32]
+ at w = addrspace(1) constant i32 0, align 4
+ at a.b = internal addrspace(2) constant [2 x i32] [i32 2, i32 3], align 4
+ at e = external addrspace(1) global i32
+ at y.z = internal addrspace(1) global i32 0, align 4
+ at 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
+}
More information about the llvm-commits
mailing list