Revert previous fix, specify subgroup size in test kernel.

oneapi-src · Sep 30, 2024 · 33c37c1 · 33c37c1
1 parent a44f34b
commit 33c37c1
Show file tree

Hide file tree

Showing 3 changed files with 10 additions and 8 deletions.
diff --git 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
@@ -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
@@ -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 {