From 068ad69d2860687d31a7fa975e65e89a3f231c35 Mon Sep 17 00:00:00 2001 From: Matthew Michel Date: Fri, 23 Oct 2026 06:50:36 -0700 Subject: [PATCH 1/3] [SYCL][Graph][Test] Test cooperative kernels under native recording Co-Authored-By: Claude Opus 5.5 --- .../NativeRecording/concurrent_root_group.cpp | 138 ++++++++++++++++++ 1 file changed, 138 insertions(+) create mode 100644 sycl/test-e2e/Graph/RecordReplay/NativeRecording/concurrent_root_group.cpp diff --git a/sycl/test-e2e/Graph/RecordReplay/NativeRecording/concurrent_root_group.cpp b/sycl/test-e2e/Graph/RecordReplay/NativeRecording/concurrent_root_group.cpp new file mode 100644 index 0000000000000..c9d729708d6eb --- /dev/null +++ b/sycl/test-e2e/Graph/RecordReplay/NativeRecording/concurrent_root_group.cpp @@ -0,0 +1,138 @@ +// REQUIRES: level_zero_v2_adapter && (arch-intel_gpu_bmg_g21 || arch-intel_gpu_bmg_g31 || arch-intel_gpu_cri) +// REQUIRES-INTEL-DRIVER: lin: 40105 + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG +// RUN: %if level_zero %{%{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} + +// Tests cooperative kernels recorded with native recording. Each work-item +// writes In, synchronizes across the whole root group, then reads the value +// written by the matching work-item of the previous work-group. The root-group +// barrier comes either from use_root_sync or from a hand-written atomic barrier +// in a kernel with a concurrent work-group forward progress guarantee. + +#include "../../graph_common.hpp" + +#include +#include +#include +#include +#include +#include + +template struct RotateKernel { + int *In; + int *Out; + const int *Seed; + uint32_t *Counter; + size_t NumGroups; + + void operator()(nd_item<1> Item) const { + const size_t I = Item.get_global_linear_id(); + const size_t N = Item.get_global_range(0); + const size_t WGSize = Item.get_local_range(0); + // Delay the writes of odd work-groups so that the even work-groups reading + // them see stale data unless the barrier really synchronizes the device. + if (Item.get_group(0) % 2 == 1) { + volatile float X = 1.0f; + for (int K = 0; K < 100; ++K) + X = X + sycl::sin(X); + } + In[I] = static_cast(I) + *Seed; + if constexpr (UseCustomRootGroupBarrier) { + group_barrier(Item.get_group(), memory_scope::device); + if (Item.get_local_linear_id() == 0) { + atomic_ref + Arrived{*Counter}; + Arrived.fetch_add(1); + while (Arrived.load(memory_order::acquire) < NumGroups) { + } + } + group_barrier(Item.get_group(), memory_scope::device); + } else { + group_barrier(Item.ext_oneapi_get_root_group()); + } + Out[I] = In[(I + N - WGSize) % N]; + } + + auto get(exp_ext::properties_tag) const { + if constexpr (UseCustomRootGroupBarrier) + return exp_ext::properties{exp_ext::work_group_progress< + exp_ext::forward_progress_guarantee::concurrent, + exp_ext::execution_scope::root_group>}; + else + return exp_ext::properties{exp_ext::use_root_sync}; + } +}; + +template class RotateKernelName; + +template void run() { + using KernelName = RotateKernelName; + queue Queue{property::queue::in_order{}}; + + constexpr size_t WGSize = 64; + auto Bundle = + get_kernel_bundle(Queue.get_context()); + const size_t MaxGroups = + Bundle.get_kernel() + .template ext_oneapi_get_info< + exp_ext::info::kernel_queue_specific::max_num_work_groups>( + Queue, range<1>{WGSize}, 0); + assert(MaxGroups >= 2); + const size_t NumGroups = std::min(MaxGroups, 256); + const size_t N = NumGroups * WGSize; + + int *In = malloc_device(N, Queue); + int *Out = malloc_device(N, Queue); + int *Seed = malloc_device(1, Queue); + uint32_t *Counter = malloc_device(1, Queue); + + exp_ext::command_graph Graph{ + Queue.get_context(), + Queue.get_device(), + {exp_ext::property::graph::enable_native_recording{}}}; + + Graph.begin_recording(Queue); + Queue.memset(Counter, 0, sizeof(uint32_t)); + Queue.parallel_for(nd_range<1>{N, WGSize}, + RotateKernel{ + In, Out, Seed, Counter, NumGroups}); + Graph.end_recording(Queue); + + auto ExecGraph = Graph.finalize(); + + std::vector HostOut(N); + for (int Replay = 0; Replay < 3; ++Replay) { + const int SeedValue = Replay * 100000; + Queue.copy(&SeedValue, Seed, 1); + Queue.memset(Out, 0, N * sizeof(int)); + Queue.ext_oneapi_graph(ExecGraph); + Queue.copy(Out, HostOut.data(), N).wait(); + + for (size_t i = 0; i < N; i++) { + const int Expected = static_cast((i + N - WGSize) % N) + SeedValue; + assert(check_value(i, Expected, HostOut[i], "HostOut")); + } + } + + free(In, Queue); + free(Out, Queue); + free(Seed, Queue); + free(Counter, Queue); +} + +void test(bool UseCustomRootGroupBarrier) { + if (UseCustomRootGroupBarrier) + run(); + else + run(); +} + +int main() { + test(/*UseCustomRootGroupBarrier=*/false); + test(/*UseCustomRootGroupBarrier=*/true); + return 0; +} From f66e418c36b7c08cac9e8ac868e8284c697bbbc7 Mon Sep 17 00:00:00 2001 From: Matthew Michel Date: Tue, 27 Oct 2026 13:46:50 -0700 Subject: [PATCH 2/3] Spec update for forward_progress / root_group interaction with SYCL graph --- .../sycl_ext_oneapi_graph.asciidoc | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index e59242312c644..d44c46bdb746b 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -2496,6 +2496,15 @@ e.g. directly to a queue without returning an event. link:../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc[sycl_ext_oneapi_free_function_kernels] can be used with SYCL Graphs. +==== sycl_ext_oneapi_forward_progress + +The forward progress properties defined by +link:../experimental/sycl_ext_oneapi_forward_progress.asciidoc[sycl_ext_oneapi_forward_progress] +are supported in graph nodes only when using +<>. Finalizing a graph containing a kernel +using these properties without native recording throws a synchronous exception +with error code `errc::feature_not_supported`. + ==== sycl_ext_oneapi_kernel_compiler_spirv The kernels loaded using @@ -2552,6 +2561,15 @@ link:../experimental/sycl_ext_oneapi_reusable_events.asciidoc[ sycl_ext_oneapi_reusable_events] is described in <>. +==== sycl_ext_oneapi_root_group + +The `use_root_sync` property defined by +link:../experimental/sycl_ext_oneapi_root_group.asciidoc[sycl_ext_oneapi_root_group] +is supported in graph nodes only when using +<>. Finalizing a graph containing a kernel +using this property without native recording throws a synchronous exception +with error code `errc::feature_not_supported`. + ==== sycl_ext_oneapi_work_group_memory Using the `work_group_memory` object defined in From c67c9a4ee89ac441080fdb77c9e0355622dc13a1 Mon Sep 17 00:00:00 2001 From: Matthew Michel Date: Wed, 28 Oct 2026 11:43:54 -0700 Subject: [PATCH 3/3] Update leak check to skip on Windows --- .../RecordReplay/NativeRecording/concurrent_root_group.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/Graph/RecordReplay/NativeRecording/concurrent_root_group.cpp b/sycl/test-e2e/Graph/RecordReplay/NativeRecording/concurrent_root_group.cpp index c9d729708d6eb..eccb9dfbcdad0 100644 --- a/sycl/test-e2e/Graph/RecordReplay/NativeRecording/concurrent_root_group.cpp +++ b/sycl/test-e2e/Graph/RecordReplay/NativeRecording/concurrent_root_group.cpp @@ -4,7 +4,7 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG -// RUN: %if level_zero %{%{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// RUN: %if level_zero %{%{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=%{leak_check_not} %} // Tests cooperative kernels recorded with native recording. Each work-item // writes In, synchronizes across the whole root group, then reads the value