[SYCL][E2E] Test kernel with kernel_device_specific::work_group_size WG size and large register file

Author: againull

sycl/test-e2e/Basic/kernel_max_wg_size.cpp

@@ -0,0 +1,54 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// REQUIRES: gpu && gpu-intel-pvc
+// UNSUPPORTED: cuda || hip
+
+// Currently fails because of issue in UR Level Zero adapter.
+// XFAIL: level_zero
+
+// clang-format off
+#include <sycl/sycl.hpp>
+#include <sycl/ext/intel/experimental/grf_size_properties.hpp>
+// clang-format on
+
+using namespace sycl;
+
+// Test that kernel can be submitted with work group size returned by
+// info::kernel_device_specific::work_group_size when large register file is
+// used.
+
+class MyKernel;
+namespace syclex = sycl::ext::oneapi::experimental;
+namespace intelex = sycl::ext::intel::experimental;
+
+__attribute__((noinline)) void f(int *Result, nd_item<1> &index) {
+  Result[index.get_global_id()] = index.get_global_id();
+}
+
+int main() {
+  queue myQueue;
+  auto myContext = myQueue.get_context();
+  auto myDev = myQueue.get_device();
+
+  kernel_id kernelId = get_kernel_id<MyKernel>();
+  auto myBundle =
+      get_kernel_bundle<bundle_state::executable>(myContext, {kernelId});
+
+  kernel myKernel = myBundle.get_kernel(kernelId);
+  size_t maxWgSize =
+      myKernel.get_info<info::kernel_device_specific::work_group_size>(myDev);
+
+  nd_range myRange{range{maxWgSize}, range{maxWgSize}};
+
+  int *Result = sycl::malloc_shared<int>(maxWgSize, myQueue);
+  syclex::properties kernel_properties{intelex::grf_size<256>};
+  myQueue.submit([&](handler &cgh) {
+    cgh.use_kernel_bundle(myBundle);
+    cgh.parallel_for<MyKernel>(myRange, kernel_properties,
+                               ([=](nd_item<1> index) { f(Result, index); }));
+  });
+
+  myQueue.wait();
+  return 0;
+}