diff --git a/sycl/test-e2e/Graph/RecordReplay/in_order_queue_event_dependency.cpp b/sycl/test-e2e/Graph/RecordReplay/in_order_queue_event_dependency.cpp new file mode 100644 index 0000000000000..a2f2a0bd1fb19 --- /dev/null +++ b/sycl/test-e2e/Graph/RecordReplay/in_order_queue_event_dependency.cpp @@ -0,0 +1,55 @@ +// 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 %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{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 %} + +#include "../graph_common.hpp" + +#include + +int main() { + queue Queue1{sycl::property::queue::in_order{}}; + queue Queue2(Queue1.get_context(), Queue1.get_device(), + sycl::property::queue::in_order()); + exp_ext::command_graph Graph{Queue1}; + + std::vector Data(Size, 0.0f); + + float *DevicePtr = sycl::malloc_device(Size, Queue1); + + Graph.begin_recording(Queue1); + + Queue1.submit([&](handler &CGH) { + CGH.memcpy(DevicePtr, Data.data(), Size * sizeof(float)); + }); + + Queue1.submit([&](handler &CGH) { + CGH.parallel_for(sycl::range<1>(Size), + [=](sycl::id<1> Id) { DevicePtr[Id] += 1.0f; }); + }); + + Graph.end_recording(Queue1); + + auto GraphExec = Graph.finalize(); + + auto Event = Queue1.ext_oneapi_graph(GraphExec); + + Queue2.submit([&](sycl::handler &CGH) { +#if 1 // Setting to zero hides the fail + CGH.depends_on({Event}); +#endif + CGH.host_task([=]() { volatile float b = 3.0; }); + }); + + std::vector HostData(Size, 0.0f); + Queue1.memcpy(HostData.data(), DevicePtr, Size * sizeof(float)).wait(); + for (size_t i = 0; i < Size; ++i) { + assert(HostData[i] == 1.0f); + } + + sycl::free(DevicePtr, Queue1); + + return 0; +}