Skip to content

Remove old style explicit kernel names from tests #340

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
GagaLP merged 1 commit into from
Aug 11, 2026
Merged

Remove old style explicit kernel names from tests #340

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
23 changes: 11 additions & 12 deletions test/accessor_tests.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -235,27 +235,27 @@ namespace detail {

auto buf_0 = tt.mbf.create_buffer(range<1>{1});

CHECK_THROWS(test_utils::add_compute_task<class UKN(task_reduction_conflict)>(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<class UKN(task_reduction_access_conflict)>(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<access_mode::read>(cgh, fixed<1>({0, 1}));
}));

CHECK_THROWS(test_utils::add_compute_task<class UKN(task_reduction_access_conflict)>(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<access_mode::write>(cgh, fixed<1>({0, 1}));
}));

CHECK_THROWS(test_utils::add_compute_task<class UKN(task_reduction_access_conflict)>(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<access_mode::read_write>(cgh, fixed<1>({0, 1}));
}));

CHECK_THROWS(test_utils::add_compute_task<class UKN(task_reduction_access_conflict)>(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<access_mode::discard_write>(cgh, fixed<1>({0, 1}));
}));
Expand Down Expand Up @@ -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<class UKN(device)>(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->();
Expand Down Expand Up @@ -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<class UKN(device)>(range<0>(), [=](item<0>) {
cgh.parallel_for(range<0>(), [=](item<0>) {
acc_0d = 1;
*acc_0d = 1;
*acc_0d.operator->() = 1;
Expand Down Expand Up @@ -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<class UKN(device)>(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<float, 0> local_0(cgh);
cgh.parallel_for<class UKN(device)>(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;
Expand Down Expand Up @@ -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<class UKN(device_kernel_1)>(
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) {
Expand Down Expand Up @@ -603,7 +602,7 @@ namespace detail {
.get_sycl_queue()
.submit([&](sycl::handler& cgh) {
closure_hydrator::get_instance().hydrate<target::device>(cgh, [&hydrated_acc, acc]() { hydrated_acc = acc; })(/* call to hydrate */);
cgh.single_task<class UKN(nop)>([] {});
cgh.single_task([] {});
})
.wait();
CHECK(accessor_testspy::get_pointer(hydrated_acc) == allocation);
Expand Down
49 changes: 24 additions & 25 deletions test/command_graph_general_tests.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<class UKN(task_a)>(test_range).discard_write(buf0, acc::one_to_one{}).submit();
cctx.device_compute<class UKN(task_b)>(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<await_push_command_record>().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<class UKN(task_a)>(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<await_push_command_record>().on(master_node_id).assert_count(1);

Expand All @@ -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<class UKN(task_a)>(test_range).discard_write(buf0, acc::one_to_one{}).submit();
const auto tid_b = cctx.device_compute<class UKN(task_b)>(test_range).discard_write(buf1, acc::one_to_one{}).submit();
const auto tid_c = cctx.device_compute<class UKN(task_c)>(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)));
}
Expand All @@ -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<class UKN(task_a)>(one_third, id<1>{0 * one_third}).discard_write(buf, acc::one_to_one{}).submit();
const auto tid_b = cctx.device_compute<class UKN(task_b)>(one_third, id<1>{1 * one_third}).discard_write(buf, acc::one_to_one{}).submit();
const auto tid_c = cctx.device_compute<class UKN(task_c)>(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<class UKN(task_d)>(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)));
Expand All @@ -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<class UKN(task_a)>(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.
Expand All @@ -100,7 +99,7 @@ TEST_CASE("command_graph_generator generates anti-dependencies for execution com
}
return {};
};
const auto tid_b = cctx.device_compute<class UKN(task_b)>(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)));
Expand All @@ -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<class UKN(task_a)>(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<class UKN(task_b)>(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<class UKN(task_c)>(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<class UKN(task_b)>(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<class UKN(task_c)>(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)));
}
}
Expand All @@ -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<class UKN(task_a)>(test_range).discard_write(buf0, acc::one_to_one{}).submit();
const auto tid_b = cctx.device_compute<class UKN(task_b)>(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)));
}

Expand All @@ -165,10 +164,10 @@ TEST_CASE(
};

// Both nodes write parts of the buffer.
[[maybe_unused]] const auto tid_a = cctx.device_compute<class UKN(task_a)>(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<class UKN(task_b)>(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<push_command_record>().on(writing_node);
Expand Down Expand Up @@ -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<class UKN(task_a)>(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<class UKN(task_b)>(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<class UKN(task_c)>(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))));
}
Expand Down
28 changes: 14 additions & 14 deletions test/command_graph_granularity_tests.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<class UKN(simple_1d)>(range<1>{255}).submit();
const auto simple_2d = cctx.device_compute<class UKN(simple_2d)>(range<2>{255, 19}).submit();
const auto simple_3d = cctx.device_compute<class UKN(simple_3d)>(range<3>{255, 19, 31}).submit();
const auto perfect_1d = cctx.device_compute<class UKN(perfect_1d)>(celerity::nd_range<1>{{256}, {32}}).submit();
const auto perfect_2d = cctx.device_compute<class UKN(perfect_2d)>(celerity::nd_range<2>{{256, 19}, {32, 19}}).submit();
const auto perfect_3d = cctx.device_compute<class UKN(perfect_3d)>(celerity::nd_range<3>{{256, 19, 31}, {32, 19, 31}}).submit();
const auto rebalance_1d = cctx.device_compute<class UKN(rebalance_1d)>(celerity::nd_range<1>{{320}, {32}}).submit();
const auto rebalance_2d = cctx.device_compute<class UKN(rebalance_2d)>(celerity::nd_range<2>{{320, 19}, {32, 19}}).submit();
const auto rebalance_3d = cctx.device_compute<class UKN(rebalance_3d)>(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;
Expand Down Expand Up @@ -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<class UKN(task)>(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<execution_command_record>(tid_a).on(0)->execution_range.range == range<3>{64, 1, 1});
CHECK(cctx.query<execution_command_record>(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<class UKN(task)>(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<execution_command_record>(tid_b).on(0)->execution_range.range == range<3>{96, 1, 1});
CHECK(cctx.query<execution_command_record>(tid_b).on(1)->execution_range.range == range<3>{96, 1, 1});
Expand All @@ -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<class UKN(task)>(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<execution_command_record>(tid_a).on(nid)->execution_range.range == range<3>{64, 64, 1});
Expand Down Expand Up @@ -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<class UKN(write)>(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<class UKN(read)>(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));

Expand Down
Loading
Loading