diff --git a/sycl/cmake/modules/UnifiedRuntimeTag.cmake b/sycl/cmake/modules/UnifiedRuntimeTag.cmake index c48d99c81b83f..8fff99790470a 100644 --- a/sycl/cmake/modules/UnifiedRuntimeTag.cmake +++ b/sycl/cmake/modules/UnifiedRuntimeTag.cmake @@ -1,5 +1,7 @@ -# commit f2af85f35957601dd91f81e8fad39fea413ccbf2 -# Author: Yang Zhao -# Date: Wed Nov 27 00:20:29 2024 +0800 -# [DeviceSanitizer] Support "-fsanitize-ignorelist=" to disable sanitizing on some of kernels (#2055) -set(UNIFIED_RUNTIME_TAG f2af85f35957601dd91f81e8fad39fea413ccbf2) +# commit 0a90db9b2c36960c9b28ce18557ca15760724c4d +# Merge: c4d9fdb4 6e0bdeb9 +# Author: Callum Fare +# Date: Wed Nov 27 12:16:44 2024 +0000 +# Merge pull request #2369 from Bensuo/ben/kernel-binary-update-l0 +# [CMDBUF] Implement kernel binary update for L0 adapter +set(UNIFIED_RUNTIME_TAG 0a90db9b2c36960c9b28ce18557ca15760724c4d) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index 2bc3ef1d921ab..9cca0ed2aa532 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -444,7 +444,8 @@ class command_graph : public detail::modifiable_command_graph { /// Constructor. /// @param SyclQueue Queue to use for the graph device and context. /// @param PropList Optional list of properties to pass. - command_graph(const queue &SyclQueue, const property_list &PropList = {}) + explicit command_graph(const queue &SyclQueue, + const property_list &PropList = {}) : modifiable_command_graph(SyclQueue, PropList) {} private: diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 178634322f47e..fe96da8d2cef7 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -723,7 +723,8 @@ bool device_impl::has(aspect Aspect) const { UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_KERNEL_ARGUMENTS | UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_LOCAL_WORK_SIZE | UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_SIZE | - UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_OFFSET; + UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_OFFSET | + UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_KERNEL_HANDLE; return has(aspect::ext_oneapi_limited_graph) && (UpdateCapabilities & RequiredCapabilities) == RequiredCapabilities; diff --git a/sycl/test-e2e/Graph/Inputs/Kernels/dyn_cgf_accessor.spv b/sycl/test-e2e/Graph/Inputs/Kernels/dyn_cgf_accessor.spv new file mode 100644 index 0000000000000..e683e8f9c5364 Binary files /dev/null and b/sycl/test-e2e/Graph/Inputs/Kernels/dyn_cgf_accessor.spv differ diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor.cpp index a5e5a1ea78b87..803f296d9f71a 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests using dynamic command-group objects with buffer accessors #include "../graph_common.hpp" @@ -23,14 +20,14 @@ int main() { Queue.get_device(), {exp_ext::property::graph::assume_buffer_outlives_graph{}}}; - int PatternA = 42; + const int PatternA = 42; auto CGFA = [&](handler &CGH) { CGH.require(Acc); CGH.parallel_for(Size, [=](item<1> Item) { Acc[Item.get_id()] = PatternA; }); }; - int PatternB = 0xA; + const int PatternB = 0xA; auto CGFB = [&](handler &CGH) { CGH.require(Acc); CGH.parallel_for(Size, diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps.cpp index 7b477edacff98..a8018190ab741 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests adding a dynamic command-group node to a graph using buffer // accessors for the node edges. diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.cpp index a420d7deb58de..4e9ada8a3c246 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests adding a dynamic command-group node to a graph using buffer // accessors for the node edges, but where different command-groups // use different buffers that create identical edges. diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp index 0eaa714463670..08b5fa293cf80 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp @@ -1,13 +1,11 @@ // RUN: %{build} -o %t.out -// RUN: %{run} %t.out %S/../Inputs/Kernels/update_with_indices_accessor.spv +// RUN: %{run} %t.out %S/../Inputs/Kernels/dyn_cgf_accessor.spv // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG -// 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 %} +// 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 %} // 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/update_with_indices_accessor.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// 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 %} // REQUIRES: level_zero -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 // Tests updating an accessor argument to a graph node created from SPIR-V // using dynamic command-groups. @@ -23,8 +21,12 @@ int main(int, char **argv) { return bundle.ext_oneapi_get_kernel(name); }; - kernel kernel = getKernel( - KernelBundle, "_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_"); + kernel kernelA = getKernel( + KernelBundle, + "_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_4itemILi1ELb1EEEE_"); + kernel kernelB = getKernel( + KernelBundle, + "_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_EUlNS0_4itemILi1ELb1EEEE_"); exp_ext::command_graph Graph{ Queue.get_context(), @@ -36,22 +38,27 @@ int main(int, char **argv) { BufA.set_write_back(false); BufB.set_write_back(false); + int PatternA = 42; + int PatternB = 0xA; + auto AccA = BufA.get_access(); auto AccB = BufB.get_access(); auto CGFA = [&](handler &CGH) { CGH.require(AccA); CGH.set_arg(0, AccA); - CGH.single_task(kernel); + CGH.set_arg(2, PatternA); + CGH.parallel_for(sycl::range<1>(Size), kernelA); }; auto CGFB = [&](handler &CGH) { CGH.require(AccB); CGH.set_arg(0, AccB); - CGH.single_task(kernel); + CGH.set_arg(2, PatternB); + CGH.parallel_for(sycl::range<1>(Size), kernelB); }; - auto DynamicCG = exp_ext::dynamic_command_group(Queue, {CGFA, CGFB}); + auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB}); auto DynamicCGNode = Graph.add(DynamicCG); auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{}); @@ -62,8 +69,8 @@ int main(int, char **argv) { Queue.copy(BufA.get_access(), HostDataA.data()).wait(); Queue.copy(BufB.get_access(), HostDataB.data()).wait(); for (size_t i = 0; i < Size; i++) { - assert(HostDataA[i] == i); - assert(HostDataB[i] == 0); + assert(check_value(i, PatternA, HostDataA[i], "HostDataA")); + assert(check_value(i, 0, HostDataB[i], "HostDataB")); } DynamicCG.set_active_cgf(1); @@ -74,8 +81,8 @@ int main(int, char **argv) { Queue.copy(BufA.get_access(), HostDataA.data()).wait(); Queue.copy(BufB.get_access(), HostDataB.data()).wait(); for (size_t i = 0; i < Size; i++) { - assert(HostDataA[i] == i); - assert(HostDataB[i] == i); + assert(check_value(i, PatternA, HostDataA[i], "HostDataA")); + assert(check_value(i, PatternB, HostDataB[i], "HostDataB")); } return 0; } diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_different_arg_nums.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_different_arg_nums.cpp index 9b3bea4bceaff..7288fba3a73d1 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_different_arg_nums.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_different_arg_nums.cpp @@ -1,12 +1,9 @@ // RUN: %{build} -o %t.out // RUN: env SYCL_UR_TRACE=2 %{run} %t.out | FileCheck %s // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG -// 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 %} +// 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 %} // 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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} - -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 +// 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 %} // Tests updating a dynamic command-group with command-groups containing a // different number of arguments. diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_event_deps.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_event_deps.cpp index 0964f6e0c354e..9556f97de69f1 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_event_deps.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_event_deps.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests adding a dynamic command-group node to a graph using graph limited // events for dependencies. diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange.cpp index 237e9173f253e..cbe1c2c3e117a 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests updating a dynamic command-group node where the dynamic command-groups // have different ranges/nd-ranges diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange_3D.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange_3D.cpp index 261ac6ecf5c3b..3fd32ef575cf4 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange_3D.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_ndrange_3D.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests updating a dynamic command-group node where the dynamic command-groups // have different range/nd-range dimensions diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_parameters.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_parameters.cpp index 7049b5bdde305..7f00d0f8750ce 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_parameters.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_parameters.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests updating kernel code using dynamic command-groups that have different // parameters in each command-group. diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_shared_nodes.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_shared_nodes.cpp index 1f98200791b6c..eab640b45b258 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_shared_nodes.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_shared_nodes.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests using the same dynamic command-group in more than one graph node. #include "../graph_common.hpp" diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_usm.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_usm.cpp index 100701f7b62aa..97c454b6db92a 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_usm.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_usm.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests updating usm kernel code using dynamic command-groups #include "../graph_common.hpp" diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_with_all_dyn_params.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_with_all_dyn_params.cpp index 53b34d1add289..28a55ecfeceeb 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_with_all_dyn_params.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_with_all_dyn_params.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests using a dynamic command-group object with dynamic parameters inside it #include "../graph_common.hpp" diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_with_different_type_dyn_params.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_with_different_type_dyn_params.cpp index 00482185ebc27..925839729cce8 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_with_different_type_dyn_params.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_with_different_type_dyn_params.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests using a dynamic command-group object with dynamic parameters of // different types diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_with_some_dyn_params.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_with_some_dyn_params.cpp index 3213fc4eec2fe..6ee6dafaaea60 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_with_some_dyn_params.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_with_some_dyn_params.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests using a dynamic command-group object where some but not all the // command-groups use dynamic parameters. diff --git a/sycl/test-e2e/Graph/Update/whole_update_dynamic_cgf.cpp b/sycl/test-e2e/Graph/Update/whole_update_dynamic_cgf.cpp index 0e8b87c0725f2..bf40f1baf7661 100644 --- a/sycl/test-e2e/Graph/Update/whole_update_dynamic_cgf.cpp +++ b/sycl/test-e2e/Graph/Update/whole_update_dynamic_cgf.cpp @@ -5,9 +5,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 2>&1 | FileCheck %s --implicit-check-not=LEAK %} -// XFAIL: level_zero -// XFAIL-TRACKER: OFNAAO-307 - // Tests interaction of whole graph update and dynamic command-groups #include "../graph_common.hpp"