Add test cases for zeCommandListAppendLaunchCooperativeKernel with SVM (#285)

KonstantyMisiak · web-flow · commit ae3cdb1d2c3c · 2025-09-24T08:37:25.000+02:00
* Add test cases for zeCommandListAppendLaunchCooperativeKernel with SVM

Signed-off-by: Misiak, Konstanty &lt;konstanty.misiak@intel.com&gt;
diff --git a/conformance_tests/core/test_memory/CMakeLists.txt b/conformance_tests/core/test_memory/CMakeLists.txt
@@ -1,4 +1,4 @@
-# Copyright (C) 2019 Intel Corporation
+# Copyright (C) 2019-2025 Intel Corporation
 # SPDX-License-Identifier: MIT
 
 if(UNIX)
@@ -24,6 +24,7 @@ add_lzt_test(
     level_zero_tests::random
     ${OS_SPECIFIC_LIBS}
   KERNELS
+    cooperative_reduction
     memory_add
     unified_mem_test
     write_memory_pattern
diff --git a/conformance_tests/core/test_memory/kernels/cooperative_reduction.cl b/conformance_tests/core/test_memory/kernels/cooperative_reduction.cl
@@ -0,0 +1,61 @@
+/*
+ *
+ * Copyright (C) 2025 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+__kernel void cooperative_reduction(__global const int* input, __global int* output,
+                                    __local int* local_sums) {
+  int gid = get_global_id(0);
+  int group_id = get_group_id(0);
+  int lid = get_local_id(0);
+  int local_size = get_local_size(0);
+
+  local_sums[lid] = input[gid];
+  barrier(CLK_LOCAL_MEM_FENCE);
+
+  if (lid == 0) {
+    for (uint i = 1; i < local_size; i++) {
+      local_sums[0] += local_sums[i];
+    }
+    output[group_id] = local_sums[0];
+  }
+
+  // Device-wide barrier: synchronize all workgroups
+  global_barrier();
+
+  if (group_id == 0 && lid == 0) {
+    for (uint i = 1; i < get_num_groups(0); i++) {
+      output[0] += output[i];
+    }
+  }
+}
+
+__kernel void cooperative_reduction_atomic(__global const int* input, __global int* output,
+                                           __local int* local_sums) {
+  int gid = get_global_id(0);
+  int group_id = get_group_id(0);
+  int lid = get_local_id(0);
+  int local_size = get_local_size(0);
+
+  local_sums[lid] = input[gid];
+  barrier(CLK_LOCAL_MEM_FENCE);
+
+  if (lid == 0) {
+    for (uint i = 1; i < local_size; i++) {
+      atomic_add(&local_sums[0], local_sums[i]);
+    }
+    output[group_id] = local_sums[0];
+  }
+
+  // Device-wide barrier: synchronize all workgroups
+  global_barrier();
+
+  if (group_id == 0 && lid == 0) {
+    for (uint i = 1; i < get_num_groups(0); i++) {
+      atomic_add(&output[0], output[i]);
+    }
+  }
+}
diff --git a/conformance_tests/core/test_memory/kernels/cooperative_reduction.spv b/conformance_tests/core/test_memory/kernels/cooperative_reduction.spv
diff --git a/conformance_tests/core/test_memory/src/test_svm.cpp b/conformance_tests/core/test_memory/src/test_svm.cpp
@@ -18,17 +18,21 @@ class SharedSystemMemoryTests
 protected:
   void SetUp() override {
     device = lzt::zeDevice::get_instance()->get_device();
-    module = lzt::create_module(device, "memory_add.spv");
-  }
 
-  void TearDown() override { lzt::destroy_module(module); }
+    bool is_dst_shared_system = std::get<0>(GetParam()).first;
+    bool is_src_shared_system = std::get<0>(GetParam()).second;
+    if (is_dst_shared_system || is_src_shared_system) {
+      SKIP_IF_SHARED_SYSTEM_ALLOC_UNSUPPORTED();
+    }
+  }
 
   ze_device_handle_t device;
-  ze_module_handle_t module;
 };
 
+class SharedSystemMemoryLaunchKernelTests : public SharedSystemMemoryTests {};
+
 LZT_TEST_P(
-    SharedSystemMemoryTests,
+    SharedSystemMemoryLaunchKernelTests,
     GivenSharedSystemMemoryAllocationsAsKernelArgumentsWhenKernelExecutesThenValuesAreCorrect) {
   bool is_dst_shared_system = std::get<0>(GetParam()).first;
   bool is_src_shared_system = std::get<0>(GetParam()).second;
@@ -39,12 +43,9 @@ LZT_TEST_P(
   constexpr size_t group_size = 32;
   ASSERT_EQ(buffer_size % (sizeof(int) * group_size), 0);
 
-  if (is_dst_shared_system || is_src_shared_system) {
-    SKIP_IF_SHARED_SYSTEM_ALLOC_UNSUPPORTED();
-  }
-
   constexpr int source_value = 1234;
   constexpr int add_value = 5678;
+  const size_t num_elements = buffer_size / sizeof(int);
 
   void *result = lzt::allocate_shared_memory_with_allocator_selector(
       buffer_size, 1, 0, 0, device, is_dst_shared_system);
@@ -53,10 +54,12 @@ LZT_TEST_P(
 
   memset(result, 0, buffer_size);
   int *source_as_int = reinterpret_cast<int *>(source);
-  for (size_t i = 0; i < buffer_size / sizeof(int); i++) {
+  for (size_t i = 0; i < num_elements; i++) {
     source_as_int[i] = source_value;
   }
 
+  ze_module_handle_t module = lzt::create_module(device, "memory_add.spv");
+
   const char *funcion_name =
       use_atomic_kernel ? "memory_atomic_add" : "memory_add";
   ze_kernel_handle_t function = lzt::create_function(module, funcion_name);
@@ -78,12 +81,13 @@ LZT_TEST_P(
   lzt::execute_and_sync_command_bundle(cmd_bundle, UINT64_MAX);
 
   int *result_as_int = reinterpret_cast<int *>(result);
-  for (size_t i = 0; i < buffer_size / sizeof(int); i++) {
+  for (size_t i = 0; i < num_elements; i++) {
     EXPECT_EQ(result_as_int[i], source_value + add_value) << "index = " << i;
   }
 
   lzt::destroy_command_bundle(cmd_bundle);
   lzt::destroy_function(function);
+  lzt::destroy_module(module);
 
   lzt::free_memory_with_allocator_selector(source, is_src_shared_system);
   lzt::free_memory_with_allocator_selector(result, is_dst_shared_system);
@@ -107,9 +111,13 @@ struct SharedSystemMemoryTestsNameSuffix {
         return "_4KB";
       case 0x1800u:
         return "_6KB";
-      case 0x100000u:
+      case 0x1'0000u:
+        return "_64KB";
+      case 0x1'0800u:
+        return "_66KB";
+      case 0x10'0000u:
         return "_1MB";
-      case 0x100800u:
+      case 0x10'0800u:
         return "_1MB2KB";
       case 0x4000'0000u:
         return "_1GB";
@@ -131,11 +139,118 @@ struct SharedSystemMemoryTestsNameSuffix {
 };
 
 INSTANTIATE_TEST_SUITE_P(
-    ParamSVMAllocationTests, SharedSystemMemoryTests,
+    ParamSVMAllocationLaunchKernelTests, SharedSystemMemoryLaunchKernelTests,
+    testing::Combine(testing::Values(std::make_pair(true, false),
+                                     std::make_pair(false, true),
+                                     std::make_pair(true, true)),
+                     testing::Bool(), testing::Bool(),
+                     testing::Values(0x80u, 0x1000u, 0x1800u, 0x10'0000u,
+                                     0x10'0800u, 0x4000'0000u, 0x4000'0800u)),
+    SharedSystemMemoryTestsNameSuffix());
+
+class SharedSystemMemoryLaunchCooperativeKernelTests
+    : public SharedSystemMemoryTests {};
+
+LZT_TEST_P(
+    SharedSystemMemoryLaunchCooperativeKernelTests,
+    GivenSharedSystemMemoryAllocationsAsKernelArgumentsWhenCooperativeKernelExecutesThenValueIsCorrect) {
+  int ordinal = -1;
+  auto command_queue_group_properties =
+      lzt::get_command_queue_group_properties(device);
+  for (int i = 0; i < command_queue_group_properties.size(); i++) {
+    if (command_queue_group_properties[i].flags &
+        ZE_COMMAND_QUEUE_GROUP_PROPERTY_FLAG_COOPERATIVE_KERNELS) {
+      ordinal = i;
+      break;
+    }
+  }
+  if (ordinal < 0) {
+    LOG_WARNING << "No command queues that support cooperative kernels";
+    GTEST_SKIP();
+  }
+
+  const bool is_dst_shared_system = std::get<0>(GetParam()).first;
+  const bool is_src_shared_system = std::get<0>(GetParam()).second;
+  const bool use_atomic_kernel = std::get<1>(GetParam());
+  const bool use_immediate_cmdlist = std::get<2>(GetParam());
+  const size_t buffer_size = std::get<3>(GetParam());
+  const size_t num_elements = buffer_size / sizeof(int);
+  LOG_INFO << "Num elements: " << num_elements;
+
+  auto compute_properties = lzt::get_compute_properties(device);
+
+  void *input = lzt::allocate_shared_memory_with_allocator_selector(
+      buffer_size, 1, 0, 0, device, is_src_shared_system);
+
+  int *input_as_int = reinterpret_cast<int *>(input);
+  for (size_t i = 0; i < num_elements; i++) {
+    input_as_int[i] = 1;
+  }
+
+  ze_module_handle_t module =
+      lzt::create_module(device, "cooperative_reduction.spv");
+  const char *function_name = use_atomic_kernel ? "cooperative_reduction_atomic"
+                                                : "cooperative_reduction";
+  ze_kernel_handle_t function = lzt::create_function(module, function_name);
+
+  uint32_t max_coop_group_count = 1;
+  lzt::suggest_max_cooperative_group_count(function, max_coop_group_count);
+  ASSERT_GT(max_coop_group_count, 0);
+
+  uint32_t suggested_group_count = [](uint32_t n) {
+    n |= n >> 1;
+    n |= n >> 2;
+    n |= n >> 4;
+    n |= n >> 8;
+    n |= n >> 16;
+    return n - (n >> 1);
+  }(max_coop_group_count);
+
+  uint32_t group_count = (num_elements < suggested_group_count)
+                             ? num_elements
+                             : suggested_group_count;
+  LOG_INFO << "Group count: " << group_count;
+
+  void *output = lzt::allocate_shared_memory_with_allocator_selector(
+      group_count * sizeof(int), 1, 0, 0, device, is_dst_shared_system);
+
+  uint32_t group_size = num_elements / group_count;
+  LOG_INFO << "Group size: " << group_size;
+  ASSERT_LE(group_size, compute_properties.maxGroupSizeX);
+
+  lzt::set_group_size(function, group_size, 1, 1);
+  lzt::set_argument_value(function, 0, sizeof(input), &input);
+  lzt::set_argument_value(function, 1, sizeof(output), &output);
+  lzt::set_argument_value(function, 2, group_size * sizeof(int), nullptr);
+
+  lzt::zeCommandBundle cmd_bundle = lzt::create_command_bundle(
+      lzt::get_default_context(), device, 0, ordinal, use_immediate_cmdlist);
+
+  ze_group_count_t thread_group_dimensions = {group_count, 1, 1};
+  lzt::append_launch_cooperative_function(
+      cmd_bundle.list, function, &thread_group_dimensions, nullptr, 0, nullptr);
+
+  lzt::close_command_list(cmd_bundle.list);
+  lzt::execute_and_sync_command_bundle(cmd_bundle, UINT64_MAX);
+
+  int *result = reinterpret_cast<int *>(output);
+  EXPECT_EQ(result[0], num_elements);
+
+  lzt::destroy_command_bundle(cmd_bundle);
+  lzt::destroy_function(function);
+  lzt::destroy_module(module);
+
+  lzt::free_memory_with_allocator_selector(output, is_dst_shared_system);
+  lzt::free_memory_with_allocator_selector(input, is_src_shared_system);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    ParamSVMAllocationLaunchCooperativeKernelTests,
+    SharedSystemMemoryLaunchCooperativeKernelTests,
     testing::Combine(testing::Values(std::make_pair(true, false),
                                      std::make_pair(false, true),
                                      std::make_pair(true, true)),
                      testing::Bool(), testing::Bool(),
-                     testing::Values(0x80u, 0x1000u, 0x1800u, 0x100000u,
-                                     0x100800u, 0x4000'0000u, 0x4000'0800u)),
+                     testing::Values(0x80u, 0x1000u, 0x1800u, 0x1'0000u,
+                                     0x1'0800u)),
     SharedSystemMemoryTestsNameSuffix());
diff --git a/scripts/level_zero_report_utils.py b/scripts/level_zero_report_utils.py
@@ -58,7 +58,8 @@ def assign_test_feature_tag(test_feature: str, test_name: str, test_section: str
                     (test_name.find("Thread") != -1) or \
                     (test_name.find("Affinity") != -1) or \
                     (test_name.find("Luid") != -1) or \
-                    (test_name.find("ParamSVMAllocationTests_SharedSystemMemoryTests") != -1) or \
+                    (test_name.find("SharedSystemMemoryLaunchKernelTests") != -1) or \
+                    (test_name.find("SharedSystemMemoryLaunchCooperativeKernelTests") != -1) or \
                     (re.search('concurrent', test_name, re.IGNORECASE)) or \
                     (re.search('context', test_name, re.IGNORECASE)) or \
                     (re.search('KernelOffset', test_name, re.IGNORECASE)) or \
diff --git a/utils/test_harness/include/test_harness/test_harness_cmdlist.hpp b/utils/test_harness/include/test_harness/test_harness_cmdlist.hpp
@@ -160,6 +160,12 @@ void append_launch_function(ze_command_list_handle_t hCommandList,
                             ze_event_handle_t hSignalEvent,
                             uint32_t numWaitEvents,
                             ze_event_handle_t *phWaitEvents);
+void append_launch_cooperative_function(ze_command_list_handle_t hCommandList,
+                                        ze_kernel_handle_t hFunction,
+                                        const ze_group_count_t *pLaunchFuncArgs,
+                                        ze_event_handle_t hSignalEvent,
+                                        uint32_t numWaitEvents,
+                                        ze_event_handle_t *phWaitEvents);
 void append_signal_event(ze_command_list_handle_t hCommandList,
                          ze_event_handle_t hEvent);
 void append_wait_on_events(ze_command_list_handle_t hCommandList,
diff --git a/utils/test_harness/include/test_harness/test_harness_module.hpp b/utils/test_harness/include/test_harness/test_harness_module.hpp
@@ -65,6 +65,8 @@ void suggest_group_size(ze_kernel_handle_t hFunction, uint32_t globalSizeX,
                         uint32_t globalSizeY, uint32_t globalSizeZ,
                         uint32_t &groupSizeX, uint32_t &groupSizeY,
                         uint32_t &groupSizeZ);
+void suggest_max_cooperative_group_count(ze_kernel_handle_t hFunction,
+                                         uint32_t &max_cooperative_group_count);
 void set_group_size(ze_kernel_handle_t hFunction, uint32_t groupSizeX,
                     uint32_t groupSizeY, uint32_t groupSizeZ);
 ze_kernel_handle_t create_function(ze_module_handle_t module,
diff --git a/utils/test_harness/src/test_harness_cmdlist.cpp b/utils/test_harness/src/test_harness_cmdlist.cpp
@@ -450,6 +450,31 @@ void append_launch_function(ze_command_list_handle_t hCommandList,
   }
 }
 
+void append_launch_cooperative_function(ze_command_list_handle_t hCommandList,
+                                        ze_kernel_handle_t hFunction,
+                                        const ze_group_count_t *pLaunchFuncArgs,
+                                        ze_event_handle_t hSignalEvent,
+                                        uint32_t numWaitEvents,
+                                        ze_event_handle_t *phWaitEvents) {
+  auto command_list_initial = hCommandList;
+  auto function_initial = hFunction;
+  auto signal_event_initial = hSignalEvent;
+  std::vector<ze_event_handle_t> wait_events_initial(numWaitEvents);
+  if (phWaitEvents) {
+    std::memcpy(wait_events_initial.data(), phWaitEvents,
+                sizeof(ze_event_handle_t) * numWaitEvents);
+  }
+  EXPECT_ZE_RESULT_SUCCESS(zeCommandListAppendLaunchCooperativeKernel(
+      hCommandList, hFunction, pLaunchFuncArgs, hSignalEvent, numWaitEvents,
+      phWaitEvents));
+  EXPECT_EQ(hCommandList, command_list_initial);
+  EXPECT_EQ(hFunction, function_initial);
+  EXPECT_EQ(hSignalEvent, signal_event_initial);
+  for (int i = 0; i < numWaitEvents && phWaitEvents; i++) {
+    EXPECT_EQ(phWaitEvents[i], wait_events_initial[i]);
+  }
+}
+
 void append_signal_event(ze_command_list_handle_t hCommandList,
                          ze_event_handle_t hEvent) {
   auto command_list_initial = hCommandList;
diff --git a/utils/test_harness/src/test_harness_module.cpp b/utils/test_harness/src/test_harness_module.cpp
@@ -232,6 +232,14 @@ void suggest_group_size(ze_kernel_handle_t hFunction, uint32_t globalSizeX,
   EXPECT_EQ(hFunction, function_initial);
 }
 
+void suggest_max_cooperative_group_count(
+    ze_kernel_handle_t hFunction, uint32_t &max_cooperative_group_count) {
+  auto function_initial = hFunction;
+  EXPECT_ZE_RESULT_SUCCESS(zeKernelSuggestMaxCooperativeGroupCount(
+      hFunction, &max_cooperative_group_count));
+  EXPECT_EQ(hFunction, function_initial);
+}
+
 void destroy_module(ze_module_handle_t module) {
   EXPECT_ZE_RESULT_SUCCESS(zeModuleDestroy(module));
 }

Original file line number	Diff line number	Diff line change
`@@ -232,6 +232,14 @@ void suggest_group_size(ze_kernel_handle_t hFunction, uint32_t globalSizeX,`
`232`	`232`	`EXPECT_EQ(hFunction, function_initial);`
`233`	`233`	`}`
`234`	`234`
	`235`	`+void suggest_max_cooperative_group_count(`
	`236`	`+ ze_kernel_handle_t hFunction, uint32_t &max_cooperative_group_count) {`
	`237`	`+ auto function_initial = hFunction;`
	`238`	`+ EXPECT_ZE_RESULT_SUCCESS(zeKernelSuggestMaxCooperativeGroupCount(`
	`239`	`+ hFunction, &max_cooperative_group_count));`
	`240`	`+ EXPECT_EQ(hFunction, function_initial);`
	`241`	`+}`
	`242`	`+`
`235`	`243`	`void destroy_module(ze_module_handle_t module) {`
`236`	`244`	`EXPECT_ZE_RESULT_SUCCESS(zeModuleDestroy(module));`
`237`	`245`	`}`