[llvm] [SPIRV] Do not use OpTypeRuntimeArray in Kernel env. (PR #149522)

[via llvm-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-spir-v

Author: Marcos Maronas (maarquitos14)

<details>
<summary>Changes</summary>

Prior to this patch, when `NumElems` was 0, `OpTypeRuntimeArray` was directly generated, but it requires `Shader` capability, so it can only be generated if `Shader` env is being used. 

Additionally, the newly added test prior to this patch was generating a module with both `Shader` and `Kernel` capabilities at the same time, but they're incompatible. This patch also fixes that.

Finally, prior to this patch, the newly added test was adding `Shader` capability to the module even with the command line flag `--avoid-spirv-capabilities=Shader`. This patch also has a fix for that.

---
Full diff: https://github.com/llvm/llvm-project/pull/149522.diff


3 Files Affected:

- (modified) llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp (+5-3) 
- (modified) llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp (+7-1) 
- (added) llvm/test/CodeGen/SPIRV/array_type.ll (+22) 


``````````diff
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 83fccdc2bdba3..982d48f2a5a76 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -828,9 +828,11 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
          "Invalid array element type");
   SPIRVType *SpvTypeInt32 = getOrCreateSPIRVIntegerType(32, MIRBuilder);
   SPIRVType *ArrayType = nullptr;
-  if (NumElems != 0) {
-    Register NumElementsVReg =
-        buildConstantInt(NumElems, MIRBuilder, SpvTypeInt32, EmitIR);
+  const SPIRVSubtarget &ST =
+      cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
+  if (NumElems != 0 || !ST.isShader()) {
+    Register NumElementsVReg = buildConstantInt(
+        NumElems ? NumElems : 1, MIRBuilder, SpvTypeInt32, EmitIR);
     ArrayType = createOpType(MIRBuilder, [&](MachineIRBuilder &MIRBuilder) {
       return MIRBuilder.buildInstr(SPIRV::OpTypeArray)
           .addDef(createTypeVReg(MIRBuilder))
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index ad976e5288927..07628c6885b81 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -744,8 +744,14 @@ void SPIRV::RequirementHandler::checkSatisfiable(
     IsSatisfiable = false;
   }
 
+  AvoidCapabilitiesSet AvoidCaps;
+  if (!ST.isShader())
+    AvoidCaps.S.insert(SPIRV::Capability::Shader);
+  else
+    AvoidCaps.S.insert(SPIRV::Capability::Kernel);
+
   for (auto Cap : MinimalCaps) {
-    if (AvailableCaps.contains(Cap))
+    if (AvailableCaps.contains(Cap) && !AvoidCaps.S.contains(Cap))
       continue;
     LLVM_DEBUG(dbgs() << "Capability not supported: "
                       << getSymbolicOperandMnemonic(
diff --git a/llvm/test/CodeGen/SPIRV/array_type.ll b/llvm/test/CodeGen/SPIRV/array_type.ll
new file mode 100644
index 0000000000000..0957a7d191922
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/array_type.ll
@@ -0,0 +1,22 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-KERNEL
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; CHECK-KERNEL: OpCapability Kernel
+; CHECK-KERNEL-NOT: OpCapability Shader 
+; CHECK-KERNEL: OpTypeArray
+; CHECK-KERNEL-NOT: OpTypeRuntimeArray
+
+%"class.sycl::_V1::detail::half_impl::half" = type { half }
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo(ptr addrspace(3) noundef align 2 %_arg_temp, ptr addrspace(1) noundef align 2 %_arg_acc_a){
+entry:
+  %0 = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %_arg_acc_a, i64 15 
+  %add.ptr.i = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %0, i64 10 
+  %4 = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %add.ptr.i, i64 20 
+  %arrayidx.i5.i = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %4, i64 35
+  %arrayidx7.i = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr addrspace(3) %_arg_temp, i64 1, i64 25, i64 30
+  %5 = load i16, ptr addrspace(1) %arrayidx.i5.i, align 2
+  store i16 %5, ptr addrspace(3) %arrayidx7.i, align 2
+  ret void
+}

``````````

</details>


https://github.com/llvm/llvm-project/pull/149522