[SYCL][Graph] Reenable L0 interop test on PVC (#15493)

Re-enable the interop-level-zero-launch-kernel.cpp E2E test on PVC by making the following correctness changes to the native L0 code: * Both `zeKernelSetGroupSize` and `ZeGroupCount` were set to have 1024 in the X dimensions when this is the total size of the USM allocation, and therefore the number of work items we want. Updated to use `zeKernelSuggestGroupSize` to work calculate the correct group sizes and count. * Update `zeKernelSetArgumentValue` call so that the size parameter is the size of the pointer rather than size of the USM allocation.
intel · Sep 30, 2024 · b9eb520 · b9eb520
1 parent 7b34aee
commit b9eb520
Show file tree

Hide file tree

Showing 3 changed files with 13 additions and 13 deletions.
diff --git a/sycl/test-e2e/Graph/Explicit/interop-level-zero-launch-kernel.cpp b/sycl/test-e2e/Graph/Explicit/interop-level-zero-launch-kernel.cpp
@@ -9,9 +9,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/saxpy.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// https://github.com/intel/llvm/issues/14826
-// XFAIL: arch-intel_gpu_pvc
-
 #define GRAPH_E2E_EXPLICIT
 
 #include "../Inputs/interop-level-zero-launch-kernel.cpp"
diff --git a/sycl/test-e2e/Graph/Inputs/interop-level-zero-launch-kernel.cpp b/sycl/test-e2e/Graph/Inputs/interop-level-zero-launch-kernel.cpp
@@ -115,20 +115,26 @@ int main(int, char **argv) {
               ZeContext, ZeDevice, &ZeCommandQueueDesc, &ZeCommandList);
           assert(status == ZE_RESULT_SUCCESS);
 
-          status = zeKernelSetArgumentValue(ZeKernel, 0,
-                                            Size * sizeof(uint32_t), &MemZ);
+          status = zeKernelSetArgumentValue(ZeKernel, 0, sizeof(MemZ), &MemZ);
           assert(status == ZE_RESULT_SUCCESS);
-          status = zeKernelSetArgumentValue(ZeKernel, 1,
-                                            Size * sizeof(uint32_t), &MemX);
+          status = zeKernelSetArgumentValue(ZeKernel, 1, sizeof(MemX), &MemX);
           assert(status == ZE_RESULT_SUCCESS);
-          ze_group_count_t ZeGroupCount{Size, 1, 1};
 
-          zeKernelSetGroupSize(ZeKernel, 1024, 1, 1);
+          uint32_t GroupSizeX = 32;
+          uint32_t GroupSizeY = 1;
+          uint32_t GroupSizeZ = 1;
+          status = zeKernelSuggestGroupSize(ZeKernel, Size, 1, 1, &GroupSizeX,
+                                            &GroupSizeY, &GroupSizeZ);
           assert(status == ZE_RESULT_SUCCESS);
 
+          status = zeKernelSetGroupSize(ZeKernel, GroupSizeX, GroupSizeY,
+                                        GroupSizeZ);
+          assert(status == ZE_RESULT_SUCCESS);
+
+          ze_group_count_t ZeGroupCount{
+              static_cast<uint32_t>(Size) / GroupSizeX, 1, 1};
           status = zeCommandListAppendLaunchKernel(
               ZeCommandList, ZeKernel, &ZeGroupCount, nullptr, 0, nullptr);
-
           assert(status == ZE_RESULT_SUCCESS);
 
           status = zeCommandListHostSynchronize(ZeCommandList, 0);

diff --git a/sycl/test-e2e/Graph/RecordReplay/interop-level-zero-launch-kernel.cpp b/sycl/test-e2e/Graph/RecordReplay/interop-level-zero-launch-kernel.cpp
@@ -9,9 +9,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/saxpy.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// https://github.com/intel/llvm/issues/14826
-// XFAIL: arch-intel_gpu_pvc
-
 #define GRAPH_E2E_RECORD_REPLAY
 
 #include "../Inputs/interop-level-zero-launch-kernel.cpp"