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 index bd29a9cc9b85..b28369772020 100644 --- 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 index 704c15724b75..01f290b8045e 100644 --- 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(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 index 653d70c80b34..894c35e99515 100644 --- 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"