From 33c37c1ba8bdbb41050fbc594b8b67aefc522067 Mon Sep 17 00:00:00 2001
From: Aaron Greig <aaron.greig@codeplay.com>
Date: Mon, 30 Sep 2024 13:29:44 +0100
Subject: [PATCH] Revert previous fix, specify subgroup size in test kernel.

---
 test/conformance/device_code/subgroup.cpp             | 8 ++++++--
 test/conformance/enqueue/enqueue_adapter_opencl.match | 1 +
 test/conformance/enqueue/urEnqueueKernelLaunch.cpp    | 9 +++------
 3 files changed, 10 insertions(+), 8 deletions(-)

diff --git a/test/conformance/device_code/subgroup.cpp b/test/conformance/device_code/subgroup.cpp
index 7f41cfe0b6..3243a0e737 100644
--- a/test/conformance/device_code/subgroup.cpp
+++ b/test/conformance/device_code/subgroup.cpp
@@ -11,11 +11,15 @@ struct KernelFunctor {
     KernelFunctor(sycl::accessor<size_t, 1, sycl::access_mode::write> Acc)
         : Acc(Acc) {}
 
+    auto get(sycl::ext::oneapi::experimental::properties_tag) {
+        return sycl::ext::oneapi::experimental::properties{
+            sycl::ext::oneapi::experimental::sub_group_size<8>};
+    }
+
     void operator()(sycl::nd_item<1> NdItem) const {
         auto SG = NdItem.get_sub_group();
         if (NdItem.get_global_linear_id() == 0) {
-            // get_max_local_range is always a range<1>
-            Acc[0] = SG.get_max_local_range()[0];
+            Acc[0] = SG.get_local_linear_range();
         }
     }
 };
diff --git a/test/conformance/enqueue/enqueue_adapter_opencl.match b/test/conformance/enqueue/enqueue_adapter_opencl.match
index 8fe2fa76e1..d67d35d9b0 100644
--- a/test/conformance/enqueue/enqueue_adapter_opencl.match
+++ b/test/conformance/enqueue/enqueue_adapter_opencl.match
@@ -1,4 +1,5 @@
 {{NONDETERMINISTIC}}
 {{OPT}}urEnqueueDeviceGetGlobalVariableReadTest.Success/Intel_R__OpenCL___{{.*}}_
 urEnqueueKernelLaunchKernelWgSizeTest.Success/Intel_R__OpenCL___{{.*}}_
+urEnqueueKernelLaunchKernelSubGroupTest.Success/Intel_R__OpenCL___{{.*}}_
 {{OPT}}urEnqueueKernelLaunchUSMLinkedList.Success/Intel_R__OpenCL___{{.*}}_UsePoolEnabled
diff --git a/test/conformance/enqueue/urEnqueueKernelLaunch.cpp b/test/conformance/enqueue/urEnqueueKernelLaunch.cpp
index ff7d6eac25..deccdebd34 100644
--- a/test/conformance/enqueue/urEnqueueKernelLaunch.cpp
+++ b/test/conformance/enqueue/urEnqueueKernelLaunch.cpp
@@ -174,18 +174,15 @@ TEST_P(urEnqueueKernelLaunchKernelWgSizeTest, NonMatchingLocalSize) {
 }
 
 TEST_P(urEnqueueKernelLaunchKernelSubGroupTest, Success) {
-    uint32_t max_num_sub_groups = 0;
-    ASSERT_SUCCESS(urKernelGetSubGroupInfo(
-        kernel, device, UR_KERNEL_SUB_GROUP_INFO_MAX_NUM_SUB_GROUPS,
-        sizeof(max_num_sub_groups), &max_num_sub_groups, nullptr));
     ur_mem_handle_t buffer = nullptr;
     AddBuffer1DArg(sizeof(size_t), &buffer);
     ASSERT_SUCCESS(urEnqueueKernelLaunch(
         queue, kernel, n_dimensions, global_offset.data(), global_size.data(),
         nullptr, 0, nullptr, nullptr));
     ASSERT_SUCCESS(urQueueFinish(queue));
-    ValidateBuffer<size_t>(buffer, sizeof(size_t),
-                           static_cast<size_t>(max_num_sub_groups));
+    // We specify this subgroup size in the kernel source, and then the kernel
+    // queries for its subgroup size at runtime and writes it to the buffer.
+    ValidateBuffer<size_t>(buffer, sizeof(size_t), 8);
 }
 
 struct Pair {