[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.

Author: EwanC

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"

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);

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"