[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 #138268

Author: s-perron

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,

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("*"))

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
+}

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
+}