From 6b89b23f2978b71aa67ed889ad09426f5a2fac1f Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Wed, 29 Mar 2023 12:04:50 +0200 Subject: [PATCH 1/2] Add missing waits in graph tests --- sycl/test/graph/graph-explicit-repeated-exec.cpp | 8 ++++++-- sycl/test/graph/graph-explicit-single-node.cpp | 4 +++- sycl/test/graph/graph-record-dotp-buffer.cpp | 2 +- sycl/test/graph/graph-record-dotp.cpp | 2 +- sycl/test/graph/graph-record-simple.cpp | 2 +- sycl/test/graph/graph-record-temp-scope.cpp | 2 +- 6 files changed, 13 insertions(+), 7 deletions(-) diff --git a/sycl/test/graph/graph-explicit-repeated-exec.cpp b/sycl/test/graph/graph-explicit-repeated-exec.cpp index e4c29204d1bd5..23b9f8fbead7b 100644 --- a/sycl/test/graph/graph-explicit-repeated-exec.cpp +++ b/sycl/test/graph/graph-explicit-repeated-exec.cpp @@ -36,13 +36,17 @@ int main() { assert(arr[i] == 0); } - q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(executable_graph); }); + q.submit([&](sycl::handler &h) { + h.ext_oneapi_graph(executable_graph); + }).wait(); for (int i = 0; i < n; i++) { assert(arr[i] == 1); } - q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(executable_graph); }); + q.submit([&](sycl::handler &h) { + h.ext_oneapi_graph(executable_graph); + }).wait(); for (int i = 0; i < n; i++) assert(arr[i] == 2); diff --git a/sycl/test/graph/graph-explicit-single-node.cpp b/sycl/test/graph/graph-explicit-single-node.cpp index 69e8d30942acc..b1461f308edc0 100644 --- a/sycl/test/graph/graph-explicit-single-node.cpp +++ b/sycl/test/graph/graph-explicit-single-node.cpp @@ -37,7 +37,9 @@ int main() { assert(arr[i] == 0); } - q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(executable_graph); }); + q.submit([&](sycl::handler &h) { + h.ext_oneapi_graph(executable_graph); + }).wait(); for (int i = 0; i < n; i++) assert(arr[i] == 1); diff --git a/sycl/test/graph/graph-record-dotp-buffer.cpp b/sycl/test/graph/graph-record-dotp-buffer.cpp index 69a060682c4f5..0ee55942f031b 100644 --- a/sycl/test/graph/graph-record-dotp-buffer.cpp +++ b/sycl/test/graph/graph-record-dotp-buffer.cpp @@ -102,7 +102,7 @@ int main() { auto exec_graph = g.finalize(q.get_context()); - q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(exec_graph); }); + q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(exec_graph); }).wait(); } assert(dotpData == host_gold_result()); diff --git a/sycl/test/graph/graph-record-dotp.cpp b/sycl/test/graph/graph-record-dotp.cpp index 78fe512e5c299..cfe9d59f39735 100644 --- a/sycl/test/graph/graph-record-dotp.cpp +++ b/sycl/test/graph/graph-record-dotp.cpp @@ -86,7 +86,7 @@ int main() { auto exec_graph = g.finalize(q.get_context()); - q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(exec_graph); }); + q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(exec_graph); }).wait(); assert(dotp[0] == host_gold_result()); diff --git a/sycl/test/graph/graph-record-simple.cpp b/sycl/test/graph/graph-record-simple.cpp index 297661caf5431..d1956d8704d70 100644 --- a/sycl/test/graph/graph-record-simple.cpp +++ b/sycl/test/graph/graph-record-simple.cpp @@ -33,7 +33,7 @@ int main() { auto exec_graph = g.finalize(q.get_context()); - q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(exec_graph); }); + q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(exec_graph); }).wait(); // Verify results for (size_t i = 0; i < n; i++) { diff --git a/sycl/test/graph/graph-record-temp-scope.cpp b/sycl/test/graph/graph-record-temp-scope.cpp index b4d660ccaec0f..a4ac296dd6d29 100644 --- a/sycl/test/graph/graph-record-temp-scope.cpp +++ b/sycl/test/graph/graph-record-temp-scope.cpp @@ -36,7 +36,7 @@ int main() { auto exec_graph = g.finalize(q.get_context()); - q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(exec_graph); }); + q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(exec_graph); }).wait(); // Verify results for (size_t i = 0; i < n; i++) { From 7d4e3155a6414b64a4d06081542d56f755b60d6e Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Wed, 29 Mar 2023 12:05:55 +0200 Subject: [PATCH 2/2] Add USM device graph test --- .../graph/graph-explicit-dotp-device-mem.cpp | 96 +++++++++++++++++++ 1 file changed, 96 insertions(+) create mode 100644 sycl/test/graph/graph-explicit-dotp-device-mem.cpp diff --git a/sycl/test/graph/graph-explicit-dotp-device-mem.cpp b/sycl/test/graph/graph-explicit-dotp-device-mem.cpp new file mode 100644 index 0000000000000..3163a6fb2ea64 --- /dev/null +++ b/sycl/test/graph/graph-explicit-dotp-device-mem.cpp @@ -0,0 +1,96 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +#include + +#include + +const size_t n = 10; + +float host_gold_result() { + float alpha = 1.0f; + float beta = 2.0f; + float gamma = 3.0f; + + float sum = 0.0f; + + for (size_t i = 0; i < n; ++i) { + sum += (alpha * 1.0f + beta * 2.0f) * (gamma * 3.0f + beta * 2.0f); + } + + return sum; +} + +int main() { + float alpha = 1.0f; + float beta = 2.0f; + float gamma = 3.0f; + + sycl::property_list properties{ + sycl::property::queue::in_order{}, + sycl::ext::oneapi::property::queue::lazy_execution{}}; + + sycl::queue q{sycl::gpu_selector_v, properties}; + + sycl::ext::oneapi::experimental::command_graph g; + + float *dotp = sycl::malloc_device(1, q); + + float *x = sycl::malloc_device(n, q); + float *y = sycl::malloc_device(n, q); + float *z = sycl::malloc_device(n, q); + + /* init data on the device */ + auto n_i = g.add([&](sycl::handler &h) { + h.parallel_for(n, [=](sycl::id<1> it) { + const size_t i = it[0]; + x[i] = 1.0f; + y[i] = 2.0f; + z[i] = 3.0f; + }); + }); + + auto node_a = g.add( + [&](sycl::handler &h) { + h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> it) { + const size_t i = it[0]; + x[i] = alpha * x[i] + beta * y[i]; + }); + }, + {n_i}); + + auto node_b = g.add( + [&](sycl::handler &h) { + h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> it) { + const size_t i = it[0]; + z[i] = gamma * z[i] + beta * y[i]; + }); + }, + {n_i}); + + auto node_c = g.add( + [&](sycl::handler &h) { + h.parallel_for(sycl::range<1>{n}, + sycl::reduction(dotp, 0.0f, std::plus()), + [=](sycl::id<1> it, auto &sum) { + const size_t i = it[0]; + sum += x[i] * z[i]; + }); + }, + {node_a, node_b}); + + auto executable_graph = g.finalize(q.get_context()); + + // Using shortcut for executing a graph of commands + q.ext_oneapi_graph(executable_graph).wait(); + + std::vector dotp_host(1); + q.copy(dotp, dotp_host.data(), 1).wait(); + + assert(dotp_host[0] == host_gold_result()); + + sycl::free(dotp, q); + sycl::free(x, q); + sycl::free(y, q); + sycl::free(z, q); + + return 0; +}