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

[Marcos Maronas via llvm-commits]

================
@@ -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);
----------------
maarquitos14 wrote:

The issue was originally found in a test similar to this --which I reduced for the sake of readability. Indeed the problem comes from an unbound array.
```
#include <sycl/builtins.hpp>
#include <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/experimental/work_group_memory.hpp>
#include <sycl/group_barrier.hpp>
#include <sycl/half_type.hpp>

// #include <cstring>

namespace syclexp = sycl::ext::oneapi::experimental;


template <typename T, size_t N>
void swap_array_2d(T (&a)[N][N], T (&b)[N][N], size_t batch_size) {
  sycl::queue q;
  const auto size = sycl::range{N, N};
  const auto wgsize = sycl::range{batch_size, batch_size};
  {
    sycl::buffer<T, 2> buf_a{a[0], sycl::range{N, N}};
    sycl::buffer<T, 2> buf_b{b[0], sycl::range{N, N}};
    q.submit([&](sycl::handler &cgh) {
      sycl::accessor acc_a{buf_a, cgh};
      sycl::accessor acc_b{buf_b, cgh};
      syclexp::work_group_memory<T[][N]> temp{N, cgh};
      sycl::nd_range<2> ndr{size, wgsize};
      cgh.parallel_for(ndr, [=](sycl::nd_item<2> it) {
        const auto i = it.get_global_id()[0];
        const auto j = it.get_global_id()[1];
        temp[i][j] = acc_a[i][j];
        acc_a[i][j] = acc_b[i][j];
        syclexp::work_group_memory<T[][N]> temp2{temp};
        acc_b[i][j] = temp2[i][j];
      });
    });
  }
}

constexpr size_t N = 32;
template <typename T> void test() {
  T intarr1[N][N];
  T intarr2[N][N];
  for (int i = 0; i < N; ++i) {
    for (int j = 0; j < N; ++j) {
      intarr1[i][j] = T(i) + T(j);
      intarr2[i][j] = T(i) * T(j);
    }
  }
  swap_array_2d(intarr1, intarr2, 8);
}

int main() {
  sycl::queue q;
  if (q.get_device().has(sycl::aspect::fp16))
    test<sycl::half>();
  return 0;
}
```


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