Skip to content

Commit f018e76

Browse files
Bensuocallumfare
Bensuo
and
callumfare
authored
[SYCL][Graph] Update UR tag for L0 kernel binary update (#16154)
- Fix potential implicit conversion from queue to command_graph when using dynamic command groups - Remove XFAIL for dyn_cgf* tests on L0 - Fix dyn_cgf_accessor_spv test and add new prebuilt spirv binary for it --------- Co-authored-by: Callum Fare <[email protected]>
1 parent 34554e6 commit f018e76

19 files changed

+36
-67
lines changed

‎sycl/cmake/modules/UnifiedRuntimeTag.cmake‎

Lines changed: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,7 @@
1-
# commit f2af85f35957601dd91f81e8fad39fea413ccbf2
2-
# Author: Yang Zhao <[email protected]>
3-
# Date: Wed Nov 27 00:20:29 2024 +0800
4-
# [DeviceSanitizer] Support "-fsanitize-ignorelist=" to disable sanitizing on some of kernels (#2055)
5-
set(UNIFIED_RUNTIME_TAG f2af85f35957601dd91f81e8fad39fea413ccbf2)
1+
# commit 0a90db9b2c36960c9b28ce18557ca15760724c4d
2+
# Merge: c4d9fdb4 6e0bdeb9
3+
# Author: Callum Fare <[email protected]>
4+
# Date: Wed Nov 27 12:16:44 2024 +0000
5+
# Merge pull request #2369 from Bensuo/ben/kernel-binary-update-l0
6+
# [CMDBUF] Implement kernel binary update for L0 adapter
7+
set(UNIFIED_RUNTIME_TAG 0a90db9b2c36960c9b28ce18557ca15760724c4d)

‎sycl/include/sycl/ext/oneapi/experimental/graph.hpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -444,7 +444,8 @@ class command_graph : public detail::modifiable_command_graph {
444444
/// Constructor.
445445
/// @param SyclQueue Queue to use for the graph device and context.
446446
/// @param PropList Optional list of properties to pass.
447-
command_graph(const queue &SyclQueue, const property_list &PropList = {})
447+
explicit command_graph(const queue &SyclQueue,
448+
const property_list &PropList = {})
448449
: modifiable_command_graph(SyclQueue, PropList) {}
449450

450451
private:

‎sycl/source/detail/device_impl.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -723,7 +723,8 @@ bool device_impl::has(aspect Aspect) const {
723723
UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_KERNEL_ARGUMENTS |
724724
UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_LOCAL_WORK_SIZE |
725725
UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_SIZE |
726-
UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_OFFSET;
726+
UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_OFFSET |
727+
UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_KERNEL_HANDLE;
727728

728729
return has(aspect::ext_oneapi_limited_graph) &&
729730
(UpdateCapabilities & RequiredCapabilities) == RequiredCapabilities;

‎sycl/test-e2e/Graph/Inputs/Kernels/dyn_cgf_accessor.spv‎

2.84 KB
Binary file not shown.

‎sycl/test-e2e/Graph/Update/dyn_cgf_accessor.cpp‎

Lines changed: 2 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -5,9 +5,6 @@
55
// Extra run to check for immediate-command-list in Level Zero
66
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
77

8-
// XFAIL: level_zero
9-
// XFAIL-TRACKER: OFNAAO-307
10-
118
// Tests using dynamic command-group objects with buffer accessors
129

1310
#include "../graph_common.hpp"
@@ -23,14 +20,14 @@ int main() {
2320
Queue.get_device(),
2421
{exp_ext::property::graph::assume_buffer_outlives_graph{}}};
2522

26-
int PatternA = 42;
23+
const int PatternA = 42;
2724
auto CGFA = [&](handler &CGH) {
2825
CGH.require(Acc);
2926
CGH.parallel_for(Size,
3027
[=](item<1> Item) { Acc[Item.get_id()] = PatternA; });
3128
};
3229

33-
int PatternB = 0xA;
30+
const int PatternB = 0xA;
3431
auto CGFB = [&](handler &CGH) {
3532
CGH.require(Acc);
3633
CGH.parallel_for(Size,

‎sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps.cpp‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5,9 +5,6 @@
55
// Extra run to check for immediate-command-list in Level Zero
66
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
77

8-
// XFAIL: level_zero
9-
// XFAIL-TRACKER: OFNAAO-307
10-
118
// Tests adding a dynamic command-group node to a graph using buffer
129
// accessors for the node edges.
1310

‎sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.cpp‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5,9 +5,6 @@
55
// Extra run to check for immediate-command-list in Level Zero
66
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
77

8-
// XFAIL: level_zero
9-
// XFAIL-TRACKER: OFNAAO-307
10-
118
// Tests adding a dynamic command-group node to a graph using buffer
129
// accessors for the node edges, but where different command-groups
1310
// use different buffers that create identical edges.

‎sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp‎

Lines changed: 21 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,11 @@
11
// RUN: %{build} -o %t.out
2-
// RUN: %{run} %t.out %S/../Inputs/Kernels/update_with_indices_accessor.spv
2+
// RUN: %{run} %t.out %S/../Inputs/Kernels/dyn_cgf_accessor.spv
33
// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
4-
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/update_with_indices_accessor.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
4+
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/dyn_cgf_accessor.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
55
// Extra run to check for immediate-command-list in Level Zero
6-
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/update_with_indices_accessor.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
6+
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/dyn_cgf_accessor.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
77

88
// REQUIRES: level_zero
9-
// XFAIL: level_zero
10-
// XFAIL-TRACKER: OFNAAO-307
119

1210
// Tests updating an accessor argument to a graph node created from SPIR-V
1311
// using dynamic command-groups.
@@ -23,8 +21,12 @@ int main(int, char **argv) {
2321
return bundle.ext_oneapi_get_kernel(name);
2422
};
2523

26-
kernel kernel = getKernel(
27-
KernelBundle, "_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_");
24+
kernel kernelA = getKernel(
25+
KernelBundle,
26+
"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_4itemILi1ELb1EEEE_");
27+
kernel kernelB = getKernel(
28+
KernelBundle,
29+
"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_EUlNS0_4itemILi1ELb1EEEE_");
2830

2931
exp_ext::command_graph Graph{
3032
Queue.get_context(),
@@ -36,22 +38,27 @@ int main(int, char **argv) {
3638
BufA.set_write_back(false);
3739
BufB.set_write_back(false);
3840

41+
int PatternA = 42;
42+
int PatternB = 0xA;
43+
3944
auto AccA = BufA.get_access();
4045
auto AccB = BufB.get_access();
4146

4247
auto CGFA = [&](handler &CGH) {
4348
CGH.require(AccA);
4449
CGH.set_arg(0, AccA);
45-
CGH.single_task(kernel);
50+
CGH.set_arg(2, PatternA);
51+
CGH.parallel_for(sycl::range<1>(Size), kernelA);
4652
};
4753

4854
auto CGFB = [&](handler &CGH) {
4955
CGH.require(AccB);
5056
CGH.set_arg(0, AccB);
51-
CGH.single_task(kernel);
57+
CGH.set_arg(2, PatternB);
58+
CGH.parallel_for(sycl::range<1>(Size), kernelB);
5259
};
5360

54-
auto DynamicCG = exp_ext::dynamic_command_group(Queue, {CGFA, CGFB});
61+
auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB});
5562
auto DynamicCGNode = Graph.add(DynamicCG);
5663
auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
5764

@@ -62,8 +69,8 @@ int main(int, char **argv) {
6269
Queue.copy(BufA.get_access(), HostDataA.data()).wait();
6370
Queue.copy(BufB.get_access(), HostDataB.data()).wait();
6471
for (size_t i = 0; i < Size; i++) {
65-
assert(HostDataA[i] == i);
66-
assert(HostDataB[i] == 0);
72+
assert(check_value(i, PatternA, HostDataA[i], "HostDataA"));
73+
assert(check_value(i, 0, HostDataB[i], "HostDataB"));
6774
}
6875

6976
DynamicCG.set_active_cgf(1);
@@ -74,8 +81,8 @@ int main(int, char **argv) {
7481
Queue.copy(BufA.get_access(), HostDataA.data()).wait();
7582
Queue.copy(BufB.get_access(), HostDataB.data()).wait();
7683
for (size_t i = 0; i < Size; i++) {
77-
assert(HostDataA[i] == i);
78-
assert(HostDataB[i] == i);
84+
assert(check_value(i, PatternA, HostDataA[i], "HostDataA"));
85+
assert(check_value(i, PatternB, HostDataB[i], "HostDataB"));
7986
}
8087
return 0;
8188
}

‎sycl/test-e2e/Graph/Update/dyn_cgf_different_arg_nums.cpp‎

Lines changed: 2 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,9 @@
11
// RUN: %{build} -o %t.out
22
// RUN: env SYCL_UR_TRACE=2 %{run} %t.out | FileCheck %s
33
// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
4-
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
4+
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 env SYCL_UR_TRACE=2 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
55
// Extra run to check for immediate-command-list in Level Zero
6-
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
7-
8-
// XFAIL: level_zero
9-
// XFAIL-TRACKER: OFNAAO-307
6+
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 env SYCL_UR_TRACE=2 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
107

118
// Tests updating a dynamic command-group with command-groups containing a
129
// different number of arguments.

‎sycl/test-e2e/Graph/Update/dyn_cgf_event_deps.cpp‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5,9 +5,6 @@
55
// Extra run to check for immediate-command-list in Level Zero
66
// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
77

8-
// XFAIL: level_zero
9-
// XFAIL-TRACKER: OFNAAO-307
10-
118
// Tests adding a dynamic command-group node to a graph using graph limited
129
// events for dependencies.
1310

0 commit comments

Comments
 (0)