From 1e7a766273416264bc11b0be13fb9f9a161ebb8f Mon Sep 17 00:00:00 2001 From: Gabriel Mitterrutzner Date: Mon, 10 Aug 2026 18:22:30 +0200 Subject: [PATCH] Remove old style explicit kernel names from tests --- test/accessor_tests.cc | 23 ++++---- test/command_graph_general_tests.cc | 49 ++++++++-------- test/command_graph_granularity_tests.cc | 28 +++++----- test/command_graph_reduction_tests.cc | 35 ++++++------ test/command_graph_transfer_tests.cc | 29 +++++----- test/graph_compaction_tests.cc | 12 ++-- test/runtime_tests.cc | 41 +++++++------- test/system/distr_tests.cc | 24 ++++---- test/task_graph_tests.cc | 74 ++++++++++--------------- test/test_utils.h | 6 -- 10 files changed, 146 insertions(+), 175 deletions(-) diff --git a/test/accessor_tests.cc b/test/accessor_tests.cc index eff797ed9..6eea4f692 100644 --- a/test/accessor_tests.cc +++ b/test/accessor_tests.cc @@ -235,27 +235,27 @@ namespace detail { auto buf_0 = tt.mbf.create_buffer(range<1>{1}); - CHECK_THROWS(test_utils::add_compute_task(tt.tm, [&](handler& cgh) { + CHECK_THROWS(test_utils::add_compute_task(tt.tm, [&](handler& cgh) { test_utils::add_reduction(cgh, tt.mrf, buf_0, false); test_utils::add_reduction(cgh, tt.mrf, buf_0, false); })); - CHECK_THROWS(test_utils::add_compute_task(tt.tm, [&](handler& cgh) { + CHECK_THROWS(test_utils::add_compute_task(tt.tm, [&](handler& cgh) { test_utils::add_reduction(cgh, tt.mrf, buf_0, false); buf_0.get_access(cgh, fixed<1>({0, 1})); })); - CHECK_THROWS(test_utils::add_compute_task(tt.tm, [&](handler& cgh) { + CHECK_THROWS(test_utils::add_compute_task(tt.tm, [&](handler& cgh) { test_utils::add_reduction(cgh, tt.mrf, buf_0, false); buf_0.get_access(cgh, fixed<1>({0, 1})); })); - CHECK_THROWS(test_utils::add_compute_task(tt.tm, [&](handler& cgh) { + CHECK_THROWS(test_utils::add_compute_task(tt.tm, [&](handler& cgh) { test_utils::add_reduction(cgh, tt.mrf, buf_0, false); buf_0.get_access(cgh, fixed<1>({0, 1})); })); - CHECK_THROWS(test_utils::add_compute_task(tt.tm, [&](handler& cgh) { + CHECK_THROWS(test_utils::add_compute_task(tt.tm, [&](handler& cgh) { test_utils::add_reduction(cgh, tt.mrf, buf_0, false); buf_0.get_access(cgh, fixed<1>({0, 1})); })); @@ -411,7 +411,7 @@ namespace detail { q.submit([&](handler& cgh) { accessor acc_0(buf_0, cgh, read_only); accessor acc_1(buf_1, cgh, one_to_one(), write_only, no_init); - cgh.parallel_for(buf_1.get_range(), [=](item<1> it) { + cgh.parallel_for(buf_1.get_range(), [=](item<1> it) { acc_1[it] = acc_0; acc_1[it] = *acc_0; acc_1[it] = *acc_0.operator->(); @@ -440,7 +440,7 @@ namespace detail { accessor acc_1d(buf_1d, cgh, all(), write_only, no_init); accessor acc_2d(buf_2d, cgh, all(), write_only, no_init); accessor acc_3d(buf_3d, cgh, all(), write_only, no_init); - cgh.parallel_for(range<0>(), [=](item<0>) { + cgh.parallel_for(range<0>(), [=](item<0>) { acc_0d = 1; *acc_0d = 1; *acc_0d.operator->() = 1; @@ -495,12 +495,12 @@ namespace detail { queue q; q.submit([&](handler& cgh) { accessor acc_1(buf_1, cgh, one_to_one(), write_only, no_init); - cgh.parallel_for(buf_1.get_range(), [=](item<1> it) { acc_1[it] = value_a; }); + cgh.parallel_for(buf_1.get_range(), [=](item<1> it) { acc_1[it] = value_a; }); }); q.submit([&](handler& cgh) { accessor acc_1(buf_1, cgh, one_to_one(), write_only); local_accessor local_0(cgh); - cgh.parallel_for(nd_range(buf_1.get_range(), buf_1.get_range()), [=](nd_item<1> it) { + cgh.parallel_for(nd_range(buf_1.get_range(), buf_1.get_range()), [=](nd_item<1> it) { if(it.get_local_id() == 0) { local_0 = value_b; *local_0 = value_b; @@ -540,8 +540,7 @@ namespace detail { device_acc_1 = decltype(device_acc_1)(buf_1, cgh, all()); local_acc_0 = decltype(local_acc_0)(cgh); local_acc_1 = decltype(local_acc_1)(1, cgh); - cgh.parallel_for( - nd_range(1, 1), [=](nd_item<1> /* it */) { (void)device_acc_0, (void)local_acc_0, (void)device_acc_1, (void)local_acc_1; }); + cgh.parallel_for(nd_range(1, 1), [=](nd_item<1> /* it */) { (void)device_acc_0, (void)local_acc_0, (void)device_acc_1, (void)local_acc_1; }); }); q.submit([&](handler& cgh) { @@ -603,7 +602,7 @@ namespace detail { .get_sycl_queue() .submit([&](sycl::handler& cgh) { closure_hydrator::get_instance().hydrate(cgh, [&hydrated_acc, acc]() { hydrated_acc = acc; })(/* call to hydrate */); - cgh.single_task([] {}); + cgh.single_task([] {}); }) .wait(); CHECK(accessor_testspy::get_pointer(hydrated_acc) == allocation); diff --git a/test/command_graph_general_tests.cc b/test/command_graph_general_tests.cc index b51ac7a77..e65008b48 100644 --- a/test/command_graph_general_tests.cc +++ b/test/command_graph_general_tests.cc @@ -31,14 +31,14 @@ TEST_CASE("command_graph_generator generates dependencies for execution commands auto buf1 = cctx.create_buffer(test_range); SECTION("if data is produced remotely") { - cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); - cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); const auto tid_c = cctx.master_node_host_task().read(buf0, acc::all{}).read(buf1, acc::all{}).submit(); CHECK(cctx.query().on(master_node_id).assert_count(2).successors().contains(cctx.query(tid_c).on(master_node_id))); } SECTION("if data is produced remotely but already available from an earlier task") { - cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); cctx.master_node_host_task().read(buf0, acc::all{}).submit(); const auto await_pushes = cctx.query().on(master_node_id).assert_count(1); @@ -50,9 +50,9 @@ TEST_CASE("command_graph_generator generates dependencies for execution commands } SECTION("if data is produced locally") { - const auto tid_a = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); - const auto tid_b = cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); - const auto tid_c = cctx.device_compute(test_range).read(buf0, acc::one_to_one{}).read(buf1, acc::one_to_one{}).submit(); + const auto tid_a = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + const auto tid_b = cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); + const auto tid_c = cctx.device_compute(test_range).read(buf0, acc::one_to_one{}).read(buf1, acc::one_to_one{}).submit(); CHECK(cctx.query(tid_a).successors().contains(cctx.query(tid_c))); CHECK(cctx.query(tid_b).successors().contains(cctx.query(tid_c))); } @@ -66,11 +66,11 @@ TEST_CASE( const range<1> one_third = {test_range / 3}; auto buf = cctx.create_buffer(test_range); - const auto tid_a = cctx.device_compute(one_third, id<1>{0 * one_third}).discard_write(buf, acc::one_to_one{}).submit(); - const auto tid_b = cctx.device_compute(one_third, id<1>{1 * one_third}).discard_write(buf, acc::one_to_one{}).submit(); - const auto tid_c = cctx.device_compute(one_third, id<1>{2 * one_third}).discard_write(buf, acc::one_to_one{}).submit(); + const auto tid_a = cctx.device_compute(one_third, id<1>{0 * one_third}).discard_write(buf, acc::one_to_one{}).submit(); + const auto tid_b = cctx.device_compute(one_third, id<1>{1 * one_third}).discard_write(buf, acc::one_to_one{}).submit(); + const auto tid_c = cctx.device_compute(one_third, id<1>{2 * one_third}).discard_write(buf, acc::one_to_one{}).submit(); - const auto tid_d = cctx.device_compute(test_range).read(buf, acc::one_to_one{}).submit(); + const auto tid_d = cctx.device_compute(test_range).read(buf, acc::one_to_one{}).submit(); CHECK(cctx.query(tid_a).successors().contains(cctx.query(tid_d))); CHECK(cctx.query(tid_b).successors().contains(cctx.query(tid_d))); CHECK(cctx.query(tid_c).successors().contains(cctx.query(tid_d))); @@ -86,8 +86,7 @@ TEST_CASE("command_graph_generator generates anti-dependencies for execution com auto buf1 = cctx.create_buffer(test_range); // Initialize both buffers - const auto tid_a = - cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).discard_write(buf1, acc::one_to_one{}).submit(); + const auto tid_a = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).discard_write(buf1, acc::one_to_one{}).submit(); // Read from buf0 but overwrite buf1 // Importantly, we only read on node 1, making it so that node 0 does not have a true dependency on the previous execution command. @@ -100,7 +99,7 @@ TEST_CASE("command_graph_generator generates anti-dependencies for execution com } return {}; }; - const auto tid_b = cctx.device_compute(test_range).read(buf0, node_1_writes).discard_write(buf1, acc::one_to_one{}).submit(); + const auto tid_b = cctx.device_compute(test_range).read(buf0, node_1_writes).discard_write(buf1, acc::one_to_one{}).submit(); CHECK(cctx.query(tid_a).on(0).successors().contains(cctx.query(tid_b).on(0))); CHECK(cctx.query(tid_a).on(1).successors().contains(cctx.query(tid_b).on(1))); @@ -114,22 +113,22 @@ TEST_CASE("command_graph_generator correctly handles anti-dependency edge cases" auto buf1 = cctx.create_buffer(test_range); // task_a writes both buffers - cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).discard_write(buf1, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).discard_write(buf1, acc::one_to_one{}).submit(); SECTION("correctly handles false anti-dependencies that consume a different buffer from the last writer") { // task_b reads buf0 - const auto tid_b = cctx.device_compute(test_range).read(buf0, acc::one_to_one{}).submit(); + const auto tid_b = cctx.device_compute(test_range).read(buf0, acc::one_to_one{}).submit(); // task_c writes buf1, initially making task_b a potential anti-dependency (as it is a successor of task_a). - const auto tid_c = cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); + const auto tid_c = cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); // However, since the two tasks don't actually touch the same buffers at all, nothing needs to be done. CHECK(cctx.query(tid_b).is_concurrent_with(cctx.query(tid_c))); } SECTION("does not consider anti-successors of last writer as potential anti-dependencies") { // task_b writes buf0, making task_a an anti-dependency - const auto tid_b = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + const auto tid_b = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); // task_c writes buf1. Since task_b is not a true successor of task_a, we don't consider it as a potential anti-dependency. - const auto tid_c = cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); + const auto tid_c = cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); CHECK(cctx.query(tid_b).is_concurrent_with(cctx.query(tid_c))); } } @@ -141,8 +140,8 @@ TEST_CASE("command_graph_generator generates anti-dependencies onto the original const range<1> test_range = {128}; auto buf0 = cctx.create_buffer(test_range); - const auto tid_a = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); - const auto tid_b = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + const auto tid_a = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + const auto tid_b = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); CHECK(cctx.query(tid_a).successors().contains(cctx.query(tid_b))); } @@ -165,10 +164,10 @@ TEST_CASE( }; // Both nodes write parts of the buffer. - [[maybe_unused]] const auto tid_a = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + [[maybe_unused]] const auto tid_a = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); // Both nodes read the full buffer, but writing_node also writes to it. - const auto tid_b = cctx.device_compute(test_range).read(buf0, acc::all{}).discard_write(buf0, only_one_writes).submit(); + const auto tid_b = cctx.device_compute(test_range).read(buf0, acc::all{}).discard_write(buf0, only_one_writes).submit(); // Each node pushes data to the other. const auto push_w = cctx.query().on(writing_node); @@ -197,15 +196,15 @@ TEST_CASE("command_graph_generator generates anti-dependencies for commands acce auto buf1 = cctx.create_buffer(test_range, true); // task_a reads from host-initialized buffer 0 - const auto tid_a = cctx.device_compute(test_range).read(buf0, acc::one_to_one{}).submit(); + const auto tid_a = cctx.device_compute(test_range).read(buf0, acc::one_to_one{}).submit(); // task_b writes to the same buffer 0 - const auto tid_b = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + const auto tid_b = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); // task_b should have an anti-dependency onto task_a CHECK(cctx.query(tid_a).successors().contains(cctx.query(tid_b))); // task_c writes to a different buffer 1 - const auto tid_c = cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); + const auto tid_c = cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); // task_c should not have any anti-dependencies at all CHECK(cctx.query(tid_c).is_concurrent_with(union_of(cctx.query(tid_a), cctx.query(tid_b)))); } diff --git a/test/command_graph_granularity_tests.cc b/test/command_graph_granularity_tests.cc index a03bfe036..4c8b20538 100644 --- a/test/command_graph_granularity_tests.cc +++ b/test/command_graph_granularity_tests.cc @@ -13,15 +13,15 @@ TEST_CASE("command_graph_generator respects task granularity when splitting", "[ const size_t num_nodes = 4; cdag_test_context cctx(num_nodes); - const auto simple_1d = cctx.device_compute(range<1>{255}).submit(); - const auto simple_2d = cctx.device_compute(range<2>{255, 19}).submit(); - const auto simple_3d = cctx.device_compute(range<3>{255, 19, 31}).submit(); - const auto perfect_1d = cctx.device_compute(celerity::nd_range<1>{{256}, {32}}).submit(); - const auto perfect_2d = cctx.device_compute(celerity::nd_range<2>{{256, 19}, {32, 19}}).submit(); - const auto perfect_3d = cctx.device_compute(celerity::nd_range<3>{{256, 19, 31}, {32, 19, 31}}).submit(); - const auto rebalance_1d = cctx.device_compute(celerity::nd_range<1>{{320}, {32}}).submit(); - const auto rebalance_2d = cctx.device_compute(celerity::nd_range<2>{{320, 19}, {32, 19}}).submit(); - const auto rebalance_3d = cctx.device_compute(celerity::nd_range<3>{{320, 19, 31}, {32, 19, 31}}).submit(); + const auto simple_1d = cctx.device_compute(range<1>{255}).submit(); + const auto simple_2d = cctx.device_compute(range<2>{255, 19}).submit(); + const auto simple_3d = cctx.device_compute(range<3>{255, 19, 31}).submit(); + const auto perfect_1d = cctx.device_compute(celerity::nd_range<1>{{256}, {32}}).submit(); + const auto perfect_2d = cctx.device_compute(celerity::nd_range<2>{{256, 19}, {32, 19}}).submit(); + const auto perfect_3d = cctx.device_compute(celerity::nd_range<3>{{256, 19, 31}, {32, 19, 31}}).submit(); + const auto rebalance_1d = cctx.device_compute(celerity::nd_range<1>{{320}, {32}}).submit(); + const auto rebalance_2d = cctx.device_compute(celerity::nd_range<2>{{320, 19}, {32, 19}}).submit(); + const auto rebalance_3d = cctx.device_compute(celerity::nd_range<3>{{320, 19, 31}, {32, 19, 31}}).submit(); for(auto tid : {simple_1d, simple_2d, simple_3d}) { size_t total_range_dim0 = 0; @@ -57,13 +57,13 @@ TEST_CASE("command_graph_generator respects split constraints", "[command_graph_ cdag_test_context cctx(num_nodes); // Split constraints use the same underlying mechanisms as task granularity (tested above), so we'll keep this brief - const auto tid_a = cctx.device_compute(range<1>{128}).constrain_split(range<1>{64}).submit(); + const auto tid_a = cctx.device_compute(range<1>{128}).constrain_split(range<1>{64}).submit(); REQUIRE(cctx.query(tid_a).total_count() == 2); CHECK(cctx.query(tid_a).on(0)->execution_range.range == range<3>{64, 1, 1}); CHECK(cctx.query(tid_a).on(1)->execution_range.range == range<3>{64, 1, 1}); // The more interesting aspect is that a constrained nd-range kernel uses the least common multiple of the two constraints - const auto tid_b = cctx.device_compute(nd_range<1>{{192}, {32}}).constrain_split(range<1>{3}).submit(); + const auto tid_b = cctx.device_compute(nd_range<1>{{192}, {32}}).constrain_split(range<1>{3}).submit(); REQUIRE(cctx.query(tid_b).total_count() == 2); CHECK(cctx.query(tid_b).on(0)->execution_range.range == range<3>{96, 1, 1}); CHECK(cctx.query(tid_b).on(1)->execution_range.range == range<3>{96, 1, 1}); @@ -72,7 +72,7 @@ TEST_CASE("command_graph_generator respects split constraints", "[command_graph_ TEST_CASE("command_graph_generator creates 2-dimensional chunks when providing the split_2d hint", "[command_graph_generator][split][task-hints]") { const size_t num_nodes = 4; cdag_test_context cctx(num_nodes); - const auto tid_a = cctx.device_compute(range<2>{128, 128}).hint(experimental::hints::split_2d{}).submit(); + const auto tid_a = cctx.device_compute(range<2>{128, 128}).hint(experimental::hints::split_2d{}).submit(); REQUIRE(cctx.query(tid_a).total_count() == 4); for(node_id nid = 0; nid < 4; ++nid) { CHECK(cctx.query(tid_a).on(nid)->execution_range.range == range<3>{64, 64, 1}); @@ -120,14 +120,14 @@ TEST_CASE("buffer accesses with empty ranges do not generate pushes or data-flow const range<1> buf_range{16}; auto buf = cctx.create_buffer(buf_range); - const auto write_tid = cctx.device_compute(buf_range).discard_write(buf, acc::one_to_one{}).submit(); + const auto write_tid = cctx.device_compute(buf_range).discard_write(buf, acc::one_to_one{}).submit(); const auto read_rm = [&](chunk<1> chnk) { const auto chunk_end = chnk.offset[0] + chnk.range[0]; const auto window_start = 4; const auto window_length = chunk_end > window_start ? chunk_end - window_start : 0; return subrange<1>{window_start, window_length}; }; - const auto read_tid = cctx.device_compute(buf_range / 2).read(buf, read_rm).submit(); + const auto read_tid = cctx.device_compute(buf_range / 2).read(buf, read_rm).submit(); CHECK(has_dependency(cctx.get_task_graph(), read_tid, write_tid)); diff --git a/test/command_graph_reduction_tests.cc b/test/command_graph_reduction_tests.cc index ef120c707..c233fc898 100644 --- a/test/command_graph_reduction_tests.cc +++ b/test/command_graph_reduction_tests.cc @@ -17,11 +17,10 @@ TEST_CASE("command_graph_generator generates reduction command trees", "[command auto buf0 = cctx.create_buffer(test_range); auto buf1 = cctx.create_buffer(range<1>{1}); - const auto tid_initialize = cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); - const auto tid_produce = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); - const auto tid_reduce = - cctx.device_compute(test_range).read(buf0, acc::one_to_one{}).reduce(buf1, true /* include_current_buffer_value */).submit(); - const auto tid_consume = cctx.device_compute(test_range).read(buf1, acc::all{}).submit(); + const auto tid_initialize = cctx.device_compute(test_range).discard_write(buf1, acc::one_to_one{}).submit(); + const auto tid_produce = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + const auto tid_reduce = cctx.device_compute(test_range).read(buf0, acc::one_to_one{}).reduce(buf1, true /* include_current_buffer_value */).submit(); + const auto tid_consume = cctx.device_compute(test_range).read(buf1, acc::all{}).submit(); CHECK(has_dependency(cctx.get_task_graph(), tid_reduce, tid_initialize)); CHECK(has_dependency(cctx.get_task_graph(), tid_reduce, tid_produce)); @@ -44,8 +43,8 @@ TEST_CASE("single-node configurations do not generate reduction commands", "[com const range<1> test_range = {64}; auto buf0 = cctx.create_buffer(range<1>(1)); - cctx.device_compute(test_range).reduce(buf0, false /* include_current_buffer_value */).submit(); - cctx.device_compute(test_range).read(buf0, acc::all{}).submit(); + cctx.device_compute(test_range).reduce(buf0, false /* include_current_buffer_value */).submit(); + cctx.device_compute(test_range).read(buf0, acc::all{}).submit(); CHECK(cctx.query().total_count() == 0); } @@ -56,10 +55,10 @@ TEST_CASE( const range<1> test_range = {64}; auto buf0 = cctx.create_buffer(range<1>(1)); - const auto tid_reduce = cctx.device_compute(test_range).reduce(buf0, false /* include_current_buffer_value */).submit(); - const auto tid_discard = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + const auto tid_reduce = cctx.device_compute(test_range).reduce(buf0, false /* include_current_buffer_value */).submit(); + const auto tid_discard = cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); // Now consume the result to check that the buffer was no longer in a pending reduction state (=> regression test) - cctx.device_compute(test_range).read(buf0, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).read(buf0, acc::one_to_one{}).submit(); CHECK(cctx.query().total_count() == 0); // On node 0 (where buf0 is actually being overwritten) there should be an anti-dependency between the two CHECK(cctx.query(tid_reduce).on(0).successors().contains(cctx.query(tid_discard).on(0))); @@ -85,7 +84,7 @@ TEST_CASE("command_graph_generator does not generate multiple reduction commands const range<1> test_range = {64}; auto buf0 = cctx.create_buffer(range<1>(1)); - cctx.device_compute(test_range).reduce(buf0, false /* include_current_buffer_value */).submit(); + cctx.device_compute(test_range).reduce(buf0, false /* include_current_buffer_value */).submit(); SECTION("in a single task") { cctx.master_node_host_task().read(buf0, acc::all{}).read_write(buf0, acc::all{}).write(buf0, acc::all{}).submit(); @@ -120,7 +119,7 @@ TEST_CASE("command_graph_generator forwards final reduction result if required b const range<1> test_range = {64}; auto buf0 = cctx.create_buffer(range<1>(1)); - cctx.device_compute(test_range).reduce(buf0, false /* include_current_buffer_value */).submit(); + cctx.device_compute(test_range).reduce(buf0, false /* include_current_buffer_value */).submit(); cctx.master_node_host_task().read(buf0, acc::all{}).submit(); cctx.collective_host_task().read(buf0, acc::all{}).submit(); @@ -137,8 +136,8 @@ TEST_CASE("multiple chained reductions produce appropriate data transfers", "[co auto buf0 = cctx.create_buffer(range<1>(1)); - cctx.device_compute(range<1>(num_nodes)).reduce(buf0, false /* include_current_buffer_value */).submit(); - cctx.device_compute(range<1>(num_nodes)).reduce(buf0, true /* include_current_buffer_value */).submit(); + cctx.device_compute(range<1>(num_nodes)).reduce(buf0, false /* include_current_buffer_value */).submit(); + cctx.device_compute(range<1>(num_nodes)).reduce(buf0, true /* include_current_buffer_value */).submit(); const auto reduction1 = cctx.query(); CHECK(reduction1.total_count() == 1); cctx.master_node_host_task().read(buf0, acc::all{}).submit(); @@ -169,8 +168,8 @@ TEST_CASE("reductions that overwrite the previous buffer contents do not generat return {}; }; // Node 1 initializes the buffer, then both nodes reduce into it without keeping the data from task_a. - cctx.device_compute(test_range).discard_write(buf0, only1).submit(); - cctx.device_compute(test_range).reduce(buf0, false /* include_current_buffer_value */).submit(); + cctx.device_compute(test_range).discard_write(buf0, only1).submit(); + cctx.device_compute(test_range).reduce(buf0, false /* include_current_buffer_value */).submit(); // This should not generate any data transfers. CHECK(cctx.query().total_count() == 0); CHECK(cctx.query().total_count() == 0); @@ -180,7 +179,7 @@ TEST_CASE("nodes that do not own pending reduction don't include it in final red cdag_test_context cctx(3); auto buf0 = cctx.create_buffer(range<1>(1)); - cctx.device_compute(nd_range<1>(64, 32)).reduce(buf0, false /* include_current_buffer_value */).submit(); + cctx.device_compute(nd_range<1>(64, 32)).reduce(buf0, false /* include_current_buffer_value */).submit(); CHECK(cctx.query().total_count() == 2); cctx.master_node_host_task().read(buf0, acc::all{}).submit(); @@ -197,7 +196,7 @@ TEST_CASE("reductions that do not include the current value generate anti-depend auto buf0 = cctx.create_buffer(range<1>(1)); const auto tid_write = cctx.master_node_host_task().discard_write(buf0, acc::all{}).submit(); - const auto tid_reduce = cctx.device_compute(range<1>(1)).reduce(buf0, false /* include_current_buffer_value */).submit(); + const auto tid_reduce = cctx.device_compute(range<1>(1)).reduce(buf0, false /* include_current_buffer_value */).submit(); CHECK(cctx.query(tid_write).successors().contains(cctx.query(tid_reduce))); } diff --git a/test/command_graph_transfer_tests.cc b/test/command_graph_transfer_tests.cc index 152c8800d..a078c9d9f 100644 --- a/test/command_graph_transfer_tests.cc +++ b/test/command_graph_transfer_tests.cc @@ -147,7 +147,7 @@ TEST_CASE("command_graph_generator doesn't generate data transfer commands for t auto buf0 = cctx.create_buffer(test_range); SECTION("when used in the same task") { - cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); // Both of theses are consumer modes, meaning that both have a requirement on the buffer range produced in task_a cctx.master_node_host_task().read(buf0, acc::all{}).write(buf0, acc::all{}).submit(); CHECK(cctx.query().on(1).count() == 1); @@ -165,7 +165,7 @@ TEST_CASE("command_graph_generator doesn't generate data transfer commands for t SECTION("when used in consecutive tasks") { auto buf1 = cctx.create_buffer(test_range); - cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); cctx.master_node_host_task().read(buf0, acc::all{}).discard_write(buf1, acc::all{}).submit(); CHECK(cctx.query().on(1).count() == 1); CHECK(cctx.query().on(0).count() == 1); @@ -176,7 +176,7 @@ TEST_CASE("command_graph_generator doesn't generate data transfer commands for t } SECTION("when used in parallel tasks") { - cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf0, acc::one_to_one{}).submit(); cctx.master_node_host_task().read(buf0, acc::all{}).submit(); CHECK(cctx.query().on(1).count() == 1); CHECK(cctx.query().on(0).count() == 1); @@ -258,7 +258,7 @@ TEST_CASE("command_graph_generator does not unnecessarily divide push commands d if(chnk.offset[0] + chnk.range[0] >= 64) return subrange<1>{32, 64}; return subrange<1>{0, 0}; }; - cctx.device_compute(test_range).read(buf, rm).submit(); + cctx.device_compute(test_range).read(buf, rm).submit(); const auto pushes2 = difference_of(cctx.query(), pushes1); CHECK(pushes2.total_count() == 2); } @@ -287,7 +287,7 @@ TEST_CASE("command_graph_generator generates anti-dependencies for await_push co // Node 0 starts by reading from buf (which is host-initialized) const auto tid_a = cctx.master_node_host_task().read(buf, acc::all{}).submit(); // Then both nodes write to it - cctx.device_compute(test_range).discard_write(buf, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf, acc::one_to_one{}).submit(); // Node 0 reads it again, generating a transfer cctx.master_node_host_task().read(buf, acc::all{}).submit(); // The await_push command has to wait until task_a is complete @@ -296,25 +296,25 @@ TEST_CASE("command_graph_generator generates anti-dependencies for await_push co SECTION("if writing to region used by push command") { // Both nodes write to buffer - cctx.device_compute(test_range).discard_write(buf, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf, acc::one_to_one{}).submit(); // Node 0 reads and writes the buffer, generating a push cctx.master_node_host_task().read_write(buf, acc::all{}).submit(); // Finally, both nodes read the buffer again, requiring an await_push on node 1 // Note that in this example the await_push is never at risk of actually running concurrently with the first push to node 0, as they are effectively // in a distributed dependency relationship, however more complex examples could give rise to situations where this can happen. - cctx.device_compute(test_range).read(buf, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).read(buf, acc::one_to_one{}).submit(); CHECK(cctx.query().on(1).successors().contains(cctx.query().on(1))); } SECTION("if writing to region written by another await_push command") { // Both nodes write to buffer - cctx.device_compute(test_range).discard_write(buf, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf, acc::one_to_one{}).submit(); // Node 0 reads the whole buffer const auto tid_b = cctx.master_node_host_task().read(buf, acc::all{}).submit(); const auto first_await_push = cctx.query(); CHECK(first_await_push.total_count() == 1); // Both nodes write it again - cctx.device_compute(test_range).discard_write(buf, acc::one_to_one{}).submit(); + cctx.device_compute(test_range).discard_write(buf, acc::one_to_one{}).submit(); // Node 0 reads it again cctx.master_node_host_task().read(buf, acc::all{}).submit(); const auto second_await_push = difference_of(cctx.query(), first_await_push); @@ -331,12 +331,11 @@ TEST_CASE("command_graph_generator generates anti-dependencies with subrange pre SECTION("for execution commands") { // task_a writes the first half - const auto tid_a = cctx.device_compute(range<1>(test_range[0] / 2)).discard_write(buf, acc::one_to_one{}).submit(); + const auto tid_a = cctx.device_compute(range<1>(test_range[0] / 2)).discard_write(buf, acc::one_to_one{}).submit(); // task_b reads the first half - const auto tid_b = cctx.device_compute(range<1>(test_range[0] / 2)).read(buf, acc::one_to_one{}).submit(); + const auto tid_b = cctx.device_compute(range<1>(test_range[0] / 2)).read(buf, acc::one_to_one{}).submit(); // task_c writes the second half - const auto tid_c = - cctx.device_compute(range<1>(test_range[0] / 2), id<1>(test_range[0] / 2)).discard_write(buf, acc::one_to_one{}).submit(); + const auto tid_c = cctx.device_compute(range<1>(test_range[0] / 2), id<1>(test_range[0] / 2)).discard_write(buf, acc::one_to_one{}).submit(); // task_c should not have an anti-dependency onto task_b (or task_a) CHECK(cctx.query(tid_a).is_concurrent_with(cctx.query(tid_c))); CHECK(cctx.query(tid_b).is_concurrent_with(cctx.query(tid_c))); @@ -344,11 +343,11 @@ TEST_CASE("command_graph_generator generates anti-dependencies with subrange pre SECTION("for await_push commands") { // task_a writes the full buffer - const auto tid_a = cctx.device_compute(test_range).discard_write(buf, acc::one_to_one{}).submit(); + const auto tid_a = cctx.device_compute(test_range).discard_write(buf, acc::one_to_one{}).submit(); // task_b reads the second half const auto tid_b = cctx.master_node_host_task().read(buf, acc::fixed<1>{{test_range[0] / 2, test_range[0] / 2}}).submit(); // task_c writes to the first half - cctx.device_compute(range<1>(test_range[0] / 2)).discard_write(buf, acc::one_to_one{}).submit(); + cctx.device_compute(range<1>(test_range[0] / 2)).discard_write(buf, acc::one_to_one{}).submit(); // task_d reads the first half const auto tid_d = cctx.master_node_host_task().read(buf, acc::fixed<1>{{0, test_range[0] / 2}}).submit(); // This should generate an await_push command that does NOT have an anti-dependency onto task_b, only task_a diff --git a/test/graph_compaction_tests.cc b/test/graph_compaction_tests.cc index 60a97b939..abf22bc2d 100644 --- a/test/graph_compaction_tests.cc +++ b/test/graph_compaction_tests.cc @@ -70,7 +70,7 @@ TEST_CASE("horizons prevent tracking data structures from growing indefinitely", ret.offset = id<2>(t, 0); return ret; }; - cctx.device_compute(range<1>(buffer_width)).read(buf_a, read_accessor).discard_write(buf_a, write_accessor).submit(); + cctx.device_compute(range<1>(buffer_width)).read(buf_a, read_accessor).discard_write(buf_a, write_accessor).submit(); auto& ggen = cctx.get_graph_generator(0); @@ -106,10 +106,10 @@ TEST_CASE("horizons correctly deal with antidependencies", "[horizon][command-gr auto buf_b = cctx.create_buffer<1>(full_range); // write to buf_a and buf_b - cctx.device_compute(full_range).discard_write(buf_a, acc::one_to_one{}).discard_write(buf_b, acc::one_to_one{}).submit(); + cctx.device_compute(full_range).discard_write(buf_a, acc::one_to_one{}).discard_write(buf_b, acc::one_to_one{}).submit(); // then read from buf_b to later induce anti-dependence - cctx.device_compute(full_range).read(buf_b, acc::one_to_one{}).submit(); + cctx.device_compute(full_range).read(buf_b, acc::one_to_one{}).submit(); // here, the first horizon should have been generated const auto first_horizon = cctx.query(); @@ -118,11 +118,11 @@ TEST_CASE("horizons correctly deal with antidependencies", "[horizon][command-gr // do 3 more read/writes on buf_a to generate another horizon and apply the first one task_id buf_a_rw = -1; for(int i = 0; i < 3; ++i) { - buf_a_rw = cctx.device_compute(full_range).read_write(buf_a, acc::one_to_one{}).submit(); + buf_a_rw = cctx.device_compute(full_range).read_write(buf_a, acc::one_to_one{}).submit(); } // now, do a write on buf_b which should generate an anti-dependency on the first horizon - auto write_b_after_first_horizon = cctx.device_compute(full_range) + auto write_b_after_first_horizon = cctx.device_compute(full_range) // introduce an artificial true dependency to avoid the fallback epoch dependency generated for ordering .read(buf_a, acc::one_to_one{}) .discard_write(buf_b, acc::one_to_one{}) @@ -159,7 +159,7 @@ TEST_CASE("previous horizons are used as last writers for host-initialized buffe // We need 7 tasks to generate a pseudo-critical path length of 6 (3x2 horizon step size), // and another one that triggers the actual deferred deletion. for(int i = 0; i < 8; ++i) { - cctx.device_compute(buf_range).discard_write(buf, acc::one_to_one{}).submit(); + cctx.device_compute(buf_range).discard_write(buf, acc::one_to_one{}).submit(); const auto current_horizon = task_manager_testspy::get_current_horizon(cctx.get_task_manager()); if(current_horizon != nullptr && current_horizon->get_id() > last_horizon_reached) { cctx.get_task_graph().erase_before_epoch(last_horizon_reached); diff --git a/test/runtime_tests.cc b/test/runtime_tests.cc index 59477f298..c8dd90a9d 100644 --- a/test/runtime_tests.cc +++ b/test/runtime_tests.cc @@ -178,31 +178,31 @@ namespace detail { CHECK_THROWS_WITH(q.submit([&](handler& cgh) { auto acc = buf.get_access(cgh, one_to_one{}); - cgh.parallel_for(range<1>{10}, [=](celerity::item<1>) { (void)acc; }); + cgh.parallel_for(range<1>{10}, [=](celerity::item<1>) { (void)acc; }); }), "Invalid range mapper dimensionality: 1-dimensional kernel submitted with a requirement whose range mapper is neither invocable for chunk<1> nor " "(chunk<1>, range<2>) to produce subrange<2>"); CHECK_NOTHROW(q.submit([&](handler& cgh) { auto acc = buf.get_access(cgh, one_to_one{}); - cgh.parallel_for(range<2>{10, 10}, [=](celerity::item<2>) { (void)acc; }); + cgh.parallel_for(range<2>{10, 10}, [=](celerity::item<2>) { (void)acc; }); })); CHECK_THROWS_WITH(q.submit([&](handler& cgh) { auto acc = buf.get_access(cgh, one_to_one{}); - cgh.parallel_for(range<3>{10, 10, 10}, [=](celerity::item<3>) { (void)acc; }); + cgh.parallel_for(range<3>{10, 10, 10}, [=](celerity::item<3>) { (void)acc; }); }), "Invalid range mapper dimensionality: 3-dimensional kernel submitted with a requirement whose range mapper is neither invocable for chunk<3> nor " "(chunk<3>, range<2>) to produce subrange<2>"); CHECK_NOTHROW(q.submit([&](handler& cgh) { auto acc = buf.get_access(cgh, all{}); - cgh.parallel_for(range<3>{10, 10, 10}, [=](celerity::item<3>) { (void)acc; }); + cgh.parallel_for(range<3>{10, 10, 10}, [=](celerity::item<3>) { (void)acc; }); })); CHECK_NOTHROW(q.submit([&](handler& cgh) { auto acc = buf.get_access(cgh, all{}); - cgh.parallel_for(range<3>{10, 10, 10}, [=](celerity::item<3>) { (void)acc; }); + cgh.parallel_for(range<3>{10, 10, 10}, [=](celerity::item<3>) { (void)acc; }); })); } @@ -252,37 +252,37 @@ namespace detail { buffer buf_1{range<1>{2}}; CHECK_THROWS(q.submit([&](handler& cgh) { // - cgh.parallel_for( + cgh.parallel_for( range<1>{1}, reduction(buf_1, cgh, sycl::plus{}, property::reduction::initialize_to_identity()), [=](celerity::item<1>, auto&) {}); })); buffer buf_4{range<1>{1}}; CHECK_NOTHROW(q.submit([&](handler& cgh) { // - cgh.parallel_for( + cgh.parallel_for( range<1>{1}, reduction(buf_4, cgh, sycl::plus{}, property::reduction::initialize_to_identity()), [=](celerity::item<1>, auto&) {}); })); buffer buf_2{range<2>{1, 2}}; CHECK_THROWS(q.submit([&](handler& cgh) { // - cgh.parallel_for( + cgh.parallel_for( range<2>{1, 1}, reduction(buf_2, cgh, sycl::plus{}, property::reduction::initialize_to_identity()), [=](celerity::item<2>, auto&) {}); })); buffer buf_3{range<3>{1, 2, 1}}; CHECK_THROWS(q.submit([&](handler& cgh) { // - cgh.parallel_for( + cgh.parallel_for( range<3>{1, 1, 1}, reduction(buf_3, cgh, sycl::plus{}, property::reduction::initialize_to_identity()), [=](celerity::item<3>, auto&) {}); })); buffer buf_5{range<2>{1, 1}}; CHECK_NOTHROW(q.submit([&](handler& cgh) { // - cgh.parallel_for( + cgh.parallel_for( range<2>{1, 1}, reduction(buf_5, cgh, sycl::plus{}, property::reduction::initialize_to_identity()), [=](celerity::item<2>, auto&) {}); })); buffer buf_6{range<3>{1, 1, 1}}; CHECK_NOTHROW(q.submit([&](handler& cgh) { // - cgh.parallel_for( + cgh.parallel_for( range<3>{1, 1, 1}, reduction(buf_6, cgh, sycl::plus{}, property::reduction::initialize_to_identity()), [=](celerity::item<3>, auto&) {}); })); } @@ -305,7 +305,7 @@ namespace detail { q.submit([&](handler& cgh) { local_accessor la{32, cgh}; accessor ga{out, cgh, celerity::access::one_to_one{}, write_only, no_init}; - cgh.parallel_for(celerity::nd_range<1>{64, 32}, [=](nd_item<1> item) { + cgh.parallel_for(celerity::nd_range<1>{64, 32}, [=](nd_item<1> item) { la[item.get_local_id()] = static_cast(item.get_global_linear_id()); group_barrier(item.get_group()); ga[item.get_global_id()] = la[item.get_local_range(0) - 1 - item.get_local_id(0)]; @@ -352,14 +352,13 @@ namespace detail { }); // with name - q.submit([=](handler& cgh) { cgh.parallel_for(range<1>{64}, [=](item<1> item) {}); }); - q.submit([=](handler& cgh) { cgh.parallel_for(celerity::nd_range<1>{64, 32}, [=](nd_item<1> item) {}); }); + q.submit([=](handler& cgh) { cgh.parallel_for(range<1>{64}, [=](item<1> item) {}); }); + q.submit([=](handler& cgh) { cgh.parallel_for(celerity::nd_range<1>{64, 32}, [=](nd_item<1> item) {}); }); q.submit([&](handler& cgh) { - cgh.parallel_for( - range<1>{64}, reduction(b, cgh, sycl::plus{}), [=](item<1> item, auto& r) { r += static_cast(item.get_linear_id()); }); + cgh.parallel_for(range<1>{64}, reduction(b, cgh, sycl::plus{}), [=](item<1> item, auto& r) { r += static_cast(item.get_linear_id()); }); }); q.submit([&](handler& cgh) { - cgh.parallel_for(celerity::nd_range<1>{64, 32}, reduction(b, cgh, sycl::plus{}), + cgh.parallel_for(celerity::nd_range<1>{64, 32}, reduction(b, cgh, sycl::plus{}), [=](nd_item<1> item, auto& r) { r += static_cast(item.get_global_linear_id()); }); }); } @@ -806,7 +805,7 @@ namespace detail { q.submit([&](handler& cgh) { accessor acc(buf, cgh, one_to_one(), write_only, no_init); - cgh.parallel_for(buf.get_range(), [=](celerity::item<2> item) { acc[item] = static_cast(item.get_linear_id()); }); + cgh.parallel_for(buf.get_range(), [=](celerity::item<2> item) { acc[item] = static_cast(item.get_linear_id()); }); }); const auto check_snapshot = [&](const subrange<2>& sr, const std::vector& expected_data) { @@ -831,7 +830,7 @@ namespace detail { q.submit([&](handler& cgh) { accessor acc(buf, cgh, write_only, no_init); - cgh.parallel_for(buf.get_range(), [=](celerity::item<0> item) { *acc = 42; }); + cgh.parallel_for(buf.get_range(), [=](celerity::item<0> item) { *acc = 42; }); }); const auto snapshot = q.fence(buf).get(); @@ -850,11 +849,11 @@ namespace detail { q.submit([&](handler& cgh) { accessor acc_a(buf_a, cgh, write_only, no_init); - cgh.parallel_for(range<0>(), [=](item<0> /* it */) { *acc_a = value_b; }); + cgh.parallel_for(range<0>(), [=](item<0> /* it */) { *acc_a = value_b; }); }); q.submit([&](handler& cgh) { accessor acc_b(buf_b, cgh, write_only, no_init); - cgh.parallel_for(nd_range<0>(), [=](nd_item<0> /* it */) { *acc_b = value_b; }); + cgh.parallel_for(nd_range<0>(), [=](nd_item<0> /* it */) { *acc_b = value_b; }); }); q.submit([&](handler& cgh) { accessor acc_c(buf_c, cgh, write_only_host_task, no_init); diff --git a/test/system/distr_tests.cc b/test/system/distr_tests.cc index 47d1c7a7f..db8bc9f77 100644 --- a/test/system/distr_tests.cc +++ b/test/system/distr_tests.cc @@ -31,7 +31,7 @@ namespace detail { q.submit([&](handler& cgh) { auto sum_r = reduction(sum_buf, cgh, sycl::plus{}, initialize_to_identity); auto max_r = reduction(max_buf, cgh, size_t{0}, unknown_identity_maximum{}, initialize_to_identity); - cgh.parallel_for(range{N}, id{1}, sum_r, max_r, [=](celerity::item<1> item, auto& sum, auto& max) { + cgh.parallel_for(range{N}, id{1}, sum_r, max_r, [=](celerity::item<1> item, auto& sum, auto& max) { sum += item.get_id(0); max.combine(item.get_id(0)); }); @@ -58,7 +58,7 @@ namespace detail { const int init = 42; buffer sum(&init, range{1}); q.submit([&](handler& cgh) { - cgh.parallel_for( + cgh.parallel_for( range{N}, reduction(sum, cgh, sycl::plus{} /* don't initialize to identity */), [=](celerity::item<1> item, auto& sum) { sum += 1; }); }); @@ -75,12 +75,12 @@ namespace detail { buffer sum(range(1)); q.submit([&](handler& cgh) { - cgh.parallel_for(range{N}, reduction(sum, cgh, sycl::plus{}, sycl::property::reduction::initialize_to_identity{}), + cgh.parallel_for(range{N}, reduction(sum, cgh, sycl::plus{}, sycl::property::reduction::initialize_to_identity{}), [=](celerity::item<1> item, auto& sum) { sum += 1; }); }); q.submit([&](handler& cgh) { - cgh.parallel_for( + cgh.parallel_for( range{N}, reduction(sum, cgh, sycl::plus{} /* include previous reduction result */), [=](celerity::item<1> item, auto& sum) { sum += 2; }); }); @@ -98,7 +98,7 @@ namespace detail { buffer sum(range(1)); q.submit([&](handler& cgh) { - cgh.parallel_for(range{N}, reduction(sum, cgh, sycl::plus{}, sycl::property::reduction::initialize_to_identity{}), + cgh.parallel_for(range{N}, reduction(sum, cgh, sycl::plus{}, sycl::property::reduction::initialize_to_identity{}), [=](celerity::item<1> item, auto& sum) { sum += static_cast(item.get_linear_id()); }); }); @@ -134,7 +134,7 @@ namespace detail { queue q; buffer sum(range(1)); q.submit([&](handler& cgh) { - cgh.parallel_for(range{100}, reduction(sum, cgh, sycl::plus{}, sycl::property::reduction::initialize_to_identity{}), + cgh.parallel_for(range{100}, reduction(sum, cgh, sycl::plus{}, sycl::property::reduction::initialize_to_identity{}), [](celerity::item<1> item, auto& sum) {}); }); q.submit([&](handler& cgh) { @@ -256,24 +256,24 @@ namespace detail { buffer buff_a(N); q.submit([&](handler& cgh) { accessor write_a{buff_a, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init}; - cgh.parallel_for(range<1>{N}, [=](celerity::item<1> item) { (void)write_a; }); + cgh.parallel_for(range<1>{N}, [=](celerity::item<1> item) { (void)write_a; }); }); buffer buff_b(N); q.submit([&](handler& cgh) { accessor write_b{buff_b, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init}; - cgh.parallel_for(range<1>{N}, [=](celerity::item<1> item) { (void)write_b; }); + cgh.parallel_for(range<1>{N}, [=](celerity::item<1> item) { (void)write_b; }); }); q.submit([&](handler& cgh) { accessor read_write_a{buff_a, cgh, celerity::access::one_to_one{}, celerity::read_write}; - cgh.parallel_for(range<1>{N}, [=](celerity::item<1> item) { (void)read_write_a; }); + cgh.parallel_for(range<1>{N}, [=](celerity::item<1> item) { (void)read_write_a; }); }); q.submit([&](handler& cgh) { accessor read_write_a{buff_a, cgh, celerity::access::one_to_one{}, celerity::read_write}; accessor read_write_b{buff_b, cgh, celerity::access::one_to_one{}, celerity::read_write}; - cgh.parallel_for(range<1>{N}, [=](celerity::item<1> item) { + cgh.parallel_for(range<1>{N}, [=](celerity::item<1> item) { (void)read_write_a; (void)read_write_b; }); @@ -281,7 +281,7 @@ namespace detail { q.submit([&](handler& cgh) { accessor write_a{buff_a, cgh, celerity::access::one_to_one{}, celerity::write_only, celerity::no_init}; - cgh.parallel_for(range<1>{N}, [=](celerity::item<1> item) { (void)write_a; }); + cgh.parallel_for(range<1>{N}, [=](celerity::item<1> item) { (void)write_a; }); }); q.wait(); @@ -328,7 +328,7 @@ namespace detail { accessor acc{buf, cgh, chunk_check_rm, write_only, no_init}; // The kernel has a size of 1 in dimension 0, so it will not be split into // more than one chunk (assuming current naive split behavior). - cgh.parallel_for(buf.get_range(), [=](item<2> it) { acc[it] = 0; }); + cgh.parallel_for(buf.get_range(), [=](item<2> it) { acc[it] = 0; }); }); } diff --git a/test/task_graph_tests.cc b/test/task_graph_tests.cc index 738811346..5a77c024f 100644 --- a/test/task_graph_tests.cc +++ b/test/task_graph_tests.cc @@ -160,15 +160,12 @@ namespace detail { auto tt = test_utils::task_test_context{}; auto buf = tt.mbf.create_buffer(range<1>(128), true /* mark_as_host_initialized */); - const auto tid_a = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 64}}); }); - const auto tid_b = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 128}}); }); + const auto tid_a = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 64}}); }); + const auto tid_b = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 128}}); }); CHECK(test_utils::has_dependency(tt.tdag, tid_b, tid_a)); CHECK(test_utils::has_dependency(tt.tdag, tid_b, tt.initial_epoch_task)); // for read of the host-initialized part - const auto tid_c = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{64, 128}}); }); + const auto tid_c = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{64, 128}}); }); CHECK_FALSE(test_utils::has_dependency(tt.tdag, tid_c, tid_a)); CHECK(test_utils::has_dependency(tt.tdag, tid_c, tt.initial_epoch_task)); // for read of the host-initialized part } @@ -178,20 +175,16 @@ namespace detail { auto buf = tt.mbf.create_buffer(range<1>(128)); // Write to the full buffer - const auto tid_a = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 128}}); }); + const auto tid_a = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 128}}); }); // Read the first half of the buffer - const auto tid_b = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 64}}); }); + const auto tid_b = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 64}}); }); CHECK(test_utils::has_dependency(tt.tdag, tid_b, tid_a)); // Overwrite the second half - no anti-dependency onto task_b should exist (but onto task_a) - const auto tid_c = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{64, 64}}); }); + const auto tid_c = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{64, 64}}); }); REQUIRE(test_utils::has_dependency(tt.tdag, tid_c, tid_a, dependency_kind::anti_dep)); REQUIRE_FALSE(test_utils::has_dependency(tt.tdag, tid_c, tid_b, dependency_kind::anti_dep)); // Overwrite the first half - now only an anti-dependency onto task_b should exist - const auto tid_d = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 64}}); }); + const auto tid_d = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 64}}); }); REQUIRE_FALSE(test_utils::has_dependency(tt.tdag, tid_d, tid_a, dependency_kind::anti_dep)); REQUIRE(test_utils::has_dependency(tt.tdag, tid_d, tid_b, dependency_kind::anti_dep)); } @@ -206,24 +199,24 @@ namespace detail { auto non_host_init_buf = tt.mbf.create_buffer(range<1>(128), false /* mark_as_host_initialized */); auto artificial_dependency_buf = tt.mbf.create_buffer(range<1>(1), false /* mark_as_host_initialized */); - const auto tid_a = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { + const auto tid_a = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { host_init_buf.get_access(cgh, fixed<1>{{0, 128}}); artificial_dependency_buf.get_access(cgh, all{}); }); CHECK(test_utils::has_dependency(tt.tdag, tid_a, tt.initial_epoch_task)); - const auto tid_b = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { + const auto tid_b = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { non_host_init_buf.get_access(cgh, fixed<1>{{0, 128}}); // introduce an arbitrary true-dependency to avoid the fallback epoch dependency that is generated for tasks without other true-dependencies artificial_dependency_buf.get_access(cgh, all{}); }); CHECK_FALSE(test_utils::has_dependency(tt.tdag, tid_b, tt.initial_epoch_task)); - const auto tid_c = test_utils::add_compute_task( - tt.tm, [&](handler& cgh) { host_init_buf.get_access(cgh, fixed<1>{{0, 128}}); }); + const auto tid_c = + test_utils::add_compute_task(tt.tm, [&](handler& cgh) { host_init_buf.get_access(cgh, fixed<1>{{0, 128}}); }); CHECK(test_utils::has_dependency(tt.tdag, tid_c, tid_a, dependency_kind::anti_dep)); - const auto tid_d = test_utils::add_compute_task( - tt.tm, [&](handler& cgh) { non_host_init_buf.get_access(cgh, fixed<1>{{0, 128}}); }); + const auto tid_d = + test_utils::add_compute_task(tt.tm, [&](handler& cgh) { non_host_init_buf.get_access(cgh, fixed<1>{{0, 128}}); }); // Since task b is essentially reading uninitialized garbage, it doesn't make a difference if we write into it concurrently CHECK_FALSE(test_utils::has_dependency(tt.tdag, tid_d, tid_b, dependency_kind::anti_dep)); } @@ -249,13 +242,12 @@ namespace detail { auto tt = test_utils::task_test_context{}; auto buf = tt.mbf.create_buffer(range<1>(128), true); - const auto tid_a = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { + const auto tid_a = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { for(const auto& m : mode_set) { dispatch_get_access(buf, cgh, m, fixed<1>{{0, 128}}); } }); - const auto tid_b = test_utils::add_compute_task( - tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 128}}); }); + const auto tid_b = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<1>{{0, 128}}); }); REQUIRE(test_utils::has_dependency(tt.tdag, tid_b, tid_a, dependency_kind::anti_dep)); } } @@ -272,15 +264,12 @@ namespace detail { auto tt = test_utils::task_test_context{}; auto buf = tt.mbf.create_buffer(range<1>(128), true /* mark_as_host_initialized */); - const task_id tid_a = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { dispatch_get_access(buf, cgh, producer_mode, all()); }); + const task_id tid_a = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { dispatch_get_access(buf, cgh, producer_mode, all()); }); - const task_id tid_b = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { dispatch_get_access(buf, cgh, consumer_mode, all()); }); + const task_id tid_b = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { dispatch_get_access(buf, cgh, consumer_mode, all()); }); CHECK(test_utils::has_dependency(tt.tdag, tid_b, tid_a)); - const task_id tid_c = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { dispatch_get_access(buf, cgh, producer_mode, all()); }); + const task_id tid_c = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { dispatch_get_access(buf, cgh, producer_mode, all()); }); const bool pure_consumer = consumer_mode == access_mode::read; const bool pure_producer = producer_mode == access_mode::discard_read_write || producer_mode == access_mode::discard_write; CHECK( @@ -593,10 +582,8 @@ namespace detail { CAPTURE(read_empty); CAPTURE(write_empty); - const auto write_tid = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<2>{write_sr}); }); - const auto read_tid = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<2>{read_sr}); }); + const auto write_tid = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<2>{write_sr}); }); + const auto read_tid = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<2>{read_sr}); }); CHECK(test_utils::has_any_dependency(tt.tdag, read_tid, write_tid) == (!write_empty && !read_empty)); } @@ -672,22 +659,18 @@ namespace detail { auto tt = test_utils::task_test_context{}; auto buf_a = tt.mbf.create_buffer(range<1>(1)); - const auto tid_a = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_a.get_access(cgh, all{}); }); + const auto tid_a = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_a.get_access(cgh, all{}); }); auto buf_b = tt.mbf.create_buffer(range<1>(1)); - const auto tid_b = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_b.get_access(cgh, all{}); }); + const auto tid_b = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_b.get_access(cgh, all{}); }); const auto tid_epoch = tt.tm.generate_epoch_task(epoch_action::none); - const auto tid_c = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_a.get_access(cgh, all{}); }); - const auto tid_d = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_b.get_access(cgh, all{}); }); - const auto tid_e = test_utils::add_compute_task(tt.tm, [&](handler& cgh) {}); - const auto tid_f = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_b.get_access(cgh, all{}); }); - const auto tid_g = - test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_b.get_access(cgh, all{}); }); + const auto tid_c = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_a.get_access(cgh, all{}); }); + const auto tid_d = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_b.get_access(cgh, all{}); }); + const auto tid_e = test_utils::add_compute_task(tt.tm, [&](handler& cgh) {}); + const auto tid_f = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_b.get_access(cgh, all{}); }); + const auto tid_g = test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf_b.get_access(cgh, all{}); }); CHECK(test_utils::has_dependency(tt.tdag, tid_epoch, tid_a)); CHECK(test_utils::has_dependency(tt.tdag, tid_epoch, tid_b)); @@ -771,8 +754,7 @@ namespace detail { SECTION("on a partially initialized buffer") { auto buf = tt.mbf.create_buffer<2>({64, 64}); - test_utils::add_compute_task( - tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<2>({{0, 0}, {32, 32}})); }); + test_utils::add_compute_task(tt.tm, [&](handler& cgh) { buf.get_access(cgh, fixed<2>({{0, 0}, {32, 32}})); }); CHECK_THROWS_WITH((test_utils::add_compute_task( tt.tm, [&](handler& cgh) { debug::set_task_name(cgh, "uninit_read"), buf.get_access(cgh, all{}); })), diff --git a/test/test_utils.h b/test/test_utils.h index 2ab218b2d..ba33f608b 100644 --- a/test/test_utils.h +++ b/test/test_utils.h @@ -41,12 +41,6 @@ #include "testspy/scheduler_testspy.h" #include "types.h" -// To avoid having to come up with tons of unique kernel names, we simply use the CPP counter. -// This is non-standard but widely supported. -#define _UKN_CONCAT2(x, y) x##_##y -#define _UKN_CONCAT(x, y) _UKN_CONCAT2(x, y) -#define UKN(name) _UKN_CONCAT(name, __COUNTER__) - /** * REQUIRE_LOOP is a utility macro for performing Catch2 REQUIRE assertions inside of loops. * The advantage over using a regular REQUIRE is that the number of reported assertions is much lower,