From e07608cba3db0136b19fa9af9dd39cbba011db61 Mon Sep 17 00:00:00 2001 From: Aaron Greig Date: Fri, 27 Sep 2024 16:29:55 +0100 Subject: [PATCH] Add missing code to enforce subgroup size in subgroup test. --- test/conformance/device_code/subgroup.cpp | 5 +++++ test/conformance/enqueue/enqueue_adapter_level_zero_v2.match | 1 - test/conformance/enqueue/enqueue_adapter_opencl.match | 1 - test/conformance/enqueue/urEnqueueKernelLaunch.cpp | 2 ++ 4 files changed, 7 insertions(+), 2 deletions(-) diff --git a/test/conformance/device_code/subgroup.cpp b/test/conformance/device_code/subgroup.cpp index fa4228f846..3243a0e737 100644 --- a/test/conformance/device_code/subgroup.cpp +++ b/test/conformance/device_code/subgroup.cpp @@ -11,6 +11,11 @@ struct KernelFunctor { KernelFunctor(sycl::accessor 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) { diff --git a/test/conformance/enqueue/enqueue_adapter_level_zero_v2.match b/test/conformance/enqueue/enqueue_adapter_level_zero_v2.match index df2be32782..86210364b5 100644 --- a/test/conformance/enqueue/enqueue_adapter_level_zero_v2.match +++ b/test/conformance/enqueue/enqueue_adapter_level_zero_v2.match @@ -2,7 +2,6 @@ urEnqueueDeviceGetGlobalVariableReadTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueKernelLaunchTest.InvalidKernelArgs/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueKernelLaunchKernelWgSizeTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueKernelLaunchKernelSubGroupTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueKernelLaunchWithVirtualMemory.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueKernelLaunchUSMLinkedList.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UsePoolEnabled urEnqueueKernelLaunchUSMLinkedList.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UsePoolDisabled diff --git a/test/conformance/enqueue/enqueue_adapter_opencl.match b/test/conformance/enqueue/enqueue_adapter_opencl.match index d67d35d9b0..8fe2fa76e1 100644 --- a/test/conformance/enqueue/enqueue_adapter_opencl.match +++ b/test/conformance/enqueue/enqueue_adapter_opencl.match @@ -1,5 +1,4 @@ {{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 c34240c057..deccdebd34 100644 --- a/test/conformance/enqueue/urEnqueueKernelLaunch.cpp +++ b/test/conformance/enqueue/urEnqueueKernelLaunch.cpp @@ -180,6 +180,8 @@ TEST_P(urEnqueueKernelLaunchKernelSubGroupTest, Success) { queue, kernel, n_dimensions, global_offset.data(), global_size.data(), nullptr, 0, nullptr, nullptr)); ASSERT_SUCCESS(urQueueFinish(queue)); + // 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(buffer, sizeof(size_t), 8); }