From 6a3d087a46f5459d7f74aabd014530bd6cc1177d Mon Sep 17 00:00:00 2001 From: Michael Aziz Date: Fri, 3 May 2024 19:14:52 -0700 Subject: [PATCH 1/3] Fix root group test case Includes the following changes: - Fix the `get_local_range`, `get_max_local_range`, `get_local_linear_range` tests. The expected values now match the extension specification. - Re-enable the test case on L0. It should pass once oneapi-src/unified-runtime#1581 is merged. Signed-off-by: Michael Aziz --- sycl/plugins/unified_runtime/CMakeLists.txt | 4 +-- sycl/test-e2e/GroupAlgorithm/root_group.cpp | 37 ++++++++++----------- 2 files changed, 20 insertions(+), 21 deletions(-) diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt index 19ae36bdc46a4..47583b867acf5 100644 --- a/sycl/plugins/unified_runtime/CMakeLists.txt +++ b/sycl/plugins/unified_runtime/CMakeLists.txt @@ -94,7 +94,7 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") + set(UNIFIED_RUNTIME_REPO "https://github.com/0x12CC/unified-runtime.git") # commit ebf873fb5996c9ddca32bbb7c9330d3ffe15473c # Merge: 633ec408 8f375039 # Author: Kenneth Benzie (Benie) @@ -102,7 +102,7 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT) # Merge pull request #1535 from przemektmalon/przemek/sampled-image-fetch # # [Bindless][Exp] Add device queries for sampled image fetch - set(UNIFIED_RUNTIME_TAG ebf873fb5996c9ddca32bbb7c9330d3ffe15473c) + set(UNIFIED_RUNTIME_TAG d3d3f6e519e804d109117e99ce6d84d10dba9131) fetch_adapter_source(level_zero ${UNIFIED_RUNTIME_REPO} diff --git a/sycl/test-e2e/GroupAlgorithm/root_group.cpp b/sycl/test-e2e/GroupAlgorithm/root_group.cpp index 94486874bc4da..955682bec2d76 100644 --- a/sycl/test-e2e/GroupAlgorithm/root_group.cpp +++ b/sycl/test-e2e/GroupAlgorithm/root_group.cpp @@ -1,5 +1,5 @@ -// Fails with opencl and level_zero on linux, enable when fixed. -// XFAIL: opencl || (linux && level_zero) +// Fails with opencl, enable when fixed. +// XFAIL: opencl // RUN: %{build} -I . -o %t.out // RUN: %{run} %t.out @@ -47,20 +47,19 @@ void testRootGroup() { const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize}; q.submit([&](sycl::handler &h) { sycl::accessor data{dataBuf, h}; - h.parallel_for( - range, props, [=](sycl::nd_item<1> it) { - auto root = it.ext_oneapi_get_root_group(); - data[root.get_local_id()] = root.get_local_id(); - sycl::group_barrier(root); + h.parallel_for< + class RootGroupKernel>(range, props, [=](sycl::nd_item<1> it) { + auto root = it.ext_oneapi_get_root_group(); + data[root.get_local_id()] = root.get_local_id(); + sycl::group_barrier(root); - root = - sycl::ext::oneapi::experimental::this_work_item::get_root_group< - 1>(); - int sum = data[root.get_local_id()] + - data[root.get_local_range() - root.get_local_id() - 1]; - sycl::group_barrier(root); - data[root.get_local_id()] = sum; - }); + root = + sycl::ext::oneapi::experimental::this_work_item::get_root_group<1>(); + int sum = data[root.get_local_id()] + + data[root.get_local_range() - root.get_local_id() - 1]; + sycl::group_barrier(root); + data[root.get_local_id()] = sum; + }); }); sycl::host_accessor data{dataBuf}; const int workItemCount = static_cast(range.get_global_range().size()); @@ -94,15 +93,15 @@ void testRootGroupFunctions() { ? root.get_local_id() == sycl::id<1>(0) : root.get_local_id() == sycl::id<1>(3); testResults[2] = root.get_group_range() == sycl::range<1>(1); - testResults[3] = - root.get_local_range() == sycl::range<1>(WorkGroupSize); + testResults[3] = root.get_local_range() == it.get_global_range(); testResults[4] = - root.get_max_local_range() == sycl::range<1>(WorkGroupSize); + root.get_max_local_range() == root.get_local_range(); testResults[5] = root.get_group_linear_id() == 0; testResults[6] = root.get_local_linear_id() == root.get_local_id().get(0); testResults[7] = root.get_group_linear_range() == 1; - testResults[8] = root.get_local_linear_range() == WorkGroupSize; + testResults[8] = + root.get_local_linear_range() == root.get_local_range().size(); } }); }); From e13d6d7927176b98e776f2459c4d67cee01b4a7e Mon Sep 17 00:00:00 2001 From: Michael Aziz Date: Mon, 6 May 2024 13:53:18 -0700 Subject: [PATCH 2/3] Update L0 adapter source Signed-off-by: Michael Aziz --- sycl/plugins/unified_runtime/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt index 97a717878abd6..9d03e9a5096a1 100644 --- a/sycl/plugins/unified_runtime/CMakeLists.txt +++ b/sycl/plugins/unified_runtime/CMakeLists.txt @@ -104,14 +104,14 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT) set(UNIFIED_RUNTIME_TAG d3d3f6e519e804d109117e99ce6d84d10dba9131) fetch_adapter_source(level_zero - "https://github.com/oneapi-src/unified-runtime.git" + "https://github.com/0x12CC/unified-runtime.git" # commit fb342f06e1ac244c5995b7be458d0c32a1ba17ab # Merge: a44e81b5 7186d6ce # Author: Kenneth Benzie (Benie) # Date: Fri May 3 11:55:26 2024 +0100 # Merge pull request #1549 from igchor/event_fix # Fix adding event to queue cache - fb342f06e1ac244c5995b7be458d0c32a1ba17ab + d3d3f6e519e804d109117e99ce6d84d10dba9131 ) fetch_adapter_source(opencl From c13fa1f4bdfee7f91fc1c302b869ec77ef1cdfb7 Mon Sep 17 00:00:00 2001 From: "Kenneth Benzie (Benie)" Date: Mon, 27 May 2024 10:44:04 +0100 Subject: [PATCH 3/3] [UR] Bump main tag to 905804c2 --- sycl/plugins/unified_runtime/CMakeLists.txt | 24 ++++++++------------- 1 file changed, 9 insertions(+), 15 deletions(-) diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt index 9d03e9a5096a1..2a3ab53851443 100644 --- a/sycl/plugins/unified_runtime/CMakeLists.txt +++ b/sycl/plugins/unified_runtime/CMakeLists.txt @@ -94,24 +94,18 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/0x12CC/unified-runtime.git") - # commit ebf873fb5996c9ddca32bbb7c9330d3ffe15473c - # Merge: 633ec408 8f375039 + set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") + # commit 905804c2e93dd046140057fd07a5d6191063bedc + # Merge: 0a11fb44 d3d3f6e5 # Author: Kenneth Benzie (Benie) - # Date: Thu May 2 10:56:55 2024 +0100 - # Merge pull request #1535 from przemektmalon/przemek/sampled-image-fetch - # [Bindless][Exp] Add device queries for sampled image fetch - set(UNIFIED_RUNTIME_TAG d3d3f6e519e804d109117e99ce6d84d10dba9131) + # Date: Mon May 27 10:34:13 2024 +0100 + # Merge pull request #1581 from 0x12CC/l0_cooperative_kernels + # Implement L0 cooperative kernel functions + set(UNIFIED_RUNTIME_TAG 905804c2e93dd046140057fd07a5d6191063bedc) fetch_adapter_source(level_zero - "https://github.com/0x12CC/unified-runtime.git" - # commit fb342f06e1ac244c5995b7be458d0c32a1ba17ab - # Merge: a44e81b5 7186d6ce - # Author: Kenneth Benzie (Benie) - # Date: Fri May 3 11:55:26 2024 +0100 - # Merge pull request #1549 from igchor/event_fix - # Fix adding event to queue cache - d3d3f6e519e804d109117e99ce6d84d10dba9131 + ${UNIFIED_RUNTIME_REPO} + ${UNIFIED_RUNTIME_TAG} ) fetch_adapter_source(opencl