From f85c9f7bc18fa42dcd0232c39c94ec480ec30466 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Thu, 26 Sep 2024 18:19:57 +0100 Subject: [PATCH] [SYCL][Graph] Reenable L0 interop test on PVC 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. --- .../interop-level-zero-launch-kernel.cpp | 3 --- .../interop-level-zero-launch-kernel.cpp | 20 ++++++++++++------- .../interop-level-zero-launch-kernel.cpp | 3 --- 3 files changed, 13 insertions(+), 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 index bd29a9cc9b853..b283697720201 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 704c15724b756..01f290b8045e2 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 653d70c80b347..894c35e995152 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"