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

[Marcos Maronas via llvm-commits]

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

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.

>From 2933fcf175c25242df1c9bbbcce26aa6a73a4ae3 Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Fri, 18 Jul 2025 16:33:55 +0200
Subject: [PATCH 1/2] [SPIRV] Use OpTypeArray for Kernel env.

---
 llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp |  8 ++++---
 llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp |  9 +++++++-
 llvm/test/CodeGen/SPIRV/array_type.ll         | 22 +++++++++++++++++++
 3 files changed, 35 insertions(+), 4 deletions(-)
 create mode 100644 llvm/test/CodeGen/SPIRV/array_type.ll

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..dd11b0e6bc891 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -68,6 +68,7 @@ getSymbolicOperandRequirements(SPIRV::OperandCategory::OperandCategory Category,
                                SPIRV::RequirementHandler &Reqs) {
   // A set of capabilities to avoid if there is another option.
   AvoidCapabilitiesSet AvoidCaps;
+
   if (!ST.isShader())
     AvoidCaps.S.insert(SPIRV::Capability::Shader);
   else
@@ -744,8 +745,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
+}

>From 68b8e0c323c13598b021e4ef95fba20278886e86 Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Fri, 18 Jul 2025 16:38:29 +0200
Subject: [PATCH 2/2] Remove undesired change.

---
 llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index dd11b0e6bc891..07628c6885b81 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -68,7 +68,6 @@ getSymbolicOperandRequirements(SPIRV::OperandCategory::OperandCategory Category,
                                SPIRV::RequirementHandler &Reqs) {
   // A set of capabilities to avoid if there is another option.
   AvoidCapabilitiesSet AvoidCaps;
-
   if (!ST.isShader())
     AvoidCaps.S.insert(SPIRV::Capability::Shader);
   else