Skip to content

Commit

Permalink
Remove broken backwards compatibility layer for sycl::item
Browse files Browse the repository at this point in the history
  • Loading branch information
@fknorr
fknorr committed Apr 12, 2023
1 parent b29b339 commit 67ccacc
Show file tree
Hide file tree
Showing 2 changed files with 12 additions and 47 deletions.
37 changes: 12 additions & 25 deletions include/handler.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -417,45 +417,32 @@ namespace detail {
};

template <typename Kernel, int Dims, typename... Reducers>
inline void invoke_kernel_with_celerity_item(const Kernel& kernel, const id<Dims>& s_id, const range<Dims>& global_range, const id<Dims>& global_offset,
inline void invoke_kernel(const Kernel& kernel, const id<Dims>& s_id, const range<Dims>& global_range, const id<Dims>& global_offset,
const id<Dims>& chunk_offset, Reducers&... reducers) {
kernel(make_item<Dims>(s_id + chunk_offset, global_offset, global_range), reducers...);
}

template <typename Kernel, int Dims, typename... Reducers>
inline void invoke_kernel_with_celerity_nd_item(const Kernel& kernel, const cl::sycl::nd_item<Dims>& s_item, const range<Dims>& global_range,
const id<Dims>& global_offset, const id<Dims>& chunk_offset, const range<Dims>& group_range, const id<Dims>& group_offset, Reducers&... reducers) {
inline void invoke_kernel(const Kernel& kernel, const cl::sycl::nd_item<Dims>& s_item, const range<Dims>& global_range, const id<Dims>& global_offset,
const id<Dims>& chunk_offset, const range<Dims>& group_range, const id<Dims>& group_offset, Reducers&... reducers) {
kernel(make_nd_item<Dims>(s_item, global_range, global_offset, chunk_offset, group_range, group_offset), reducers...);
}

template <typename Kernel, int Dims, typename... Reducers>
[[deprecated("Support for kernels receiving cl::sycl::item<Dims> will be removed in the future, change parameter type to celerity::item<Dims>")]] //
inline void
invoke_kernel_with_sycl_item(const Kernel& kernel, const cl::sycl::item<Dims>& s_item, Reducers&... reducers) {
kernel(s_item, reducers...);
}

template <typename Kernel, int Dims>
auto bind_simple_kernel(const Kernel& kernel, const range<Dims>& global_range, const id<Dims>& global_offset, const id<Dims>& chunk_offset) {
// The current mechanism for hydrating the SYCL placeholder accessors inside Celerity accessors requires that the kernel functor
// capturing those accessors is copied at least once during submission (see also live_pass_device_handler::submit_to_sycl).
// As of SYCL 2020 kernel functors are passed as const references, so we explicitly capture by value here.
return [=](auto s_item_or_id, auto&... reducers) {
if constexpr(std::is_invocable_v<Kernel, celerity::item<Dims>, decltype(reducers)...>) {
if constexpr(CELERITY_WORKAROUND(DPCPP) && std::is_same_v<id<Dims>, decltype(s_item_or_id)>) {
// CELERITY_WORKAROUND_LESS_OR_EQUAL: DPC++ passes a sycl::id instead of a sycl::item to kernels alongside reductions
invoke_kernel_with_celerity_item(kernel, s_item_or_id, global_range, global_offset, chunk_offset, reducers...);
} else {
// Explicit item constructor: ComputeCpp does not pass a sycl::item, but an implicitly convertible sycl::item_base (?) which does not have
// `sycl::id<> get_id()`
invoke_kernel_with_celerity_item(
kernel, cl::sycl::item<Dims>{s_item_or_id}.get_id(), global_range, global_offset, chunk_offset, reducers...);
}
} else if constexpr(std::is_invocable_v<Kernel, cl::sycl::item<Dims>, decltype(reducers)...>) {
invoke_kernel_with_sycl_item(kernel, cl::sycl::item<Dims>{s_item_or_id}, reducers...);
static_assert(std::is_invocable_v<Kernel, celerity::item<Dims>, decltype(reducers)...>,
"Kernel function must be invocable with celerity::item<Dims> and as many reducer objects as reductions passed to parallel_for");
if constexpr(CELERITY_WORKAROUND(DPCPP) && std::is_same_v<id<Dims>, decltype(s_item_or_id)>) {
// CELERITY_WORKAROUND_LESS_OR_EQUAL: DPC++ passes a sycl::id instead of a sycl::item to kernels alongside reductions
invoke_kernel(kernel, s_item_or_id, global_range, global_offset, chunk_offset, reducers...);
} else {
static_assert(constexpr_false<decltype(reducers)...>, "Kernel function must be invocable with celerity::item<Dims> (or cl::sycl::item<Dims>, "
"deprecated) and as many reducer objects as reductions passed to parallel_for");
// Explicit item constructor: ComputeCpp does not pass a sycl::item, but an implicitly convertible sycl::item_base (?) which does not have
// `sycl::id<> get_id()`
invoke_kernel(kernel, cl::sycl::item<Dims>{s_item_or_id}.get_id(), global_range, global_offset, chunk_offset, reducers...);
}
};
}
Expand All @@ -466,7 +453,7 @@ namespace detail {
return [=](cl::sycl::nd_item<Dims> s_item, auto&... reducers) {
static_assert(std::is_invocable_v<Kernel, celerity::nd_item<Dims>, decltype(reducers)...>,
"Kernel function must be invocable with celerity::nd_item<Dims> or and as many reducer objects as reductions passed to parallel_for");
invoke_kernel_with_celerity_nd_item(kernel, s_item, global_range, global_offset, chunk_offset, group_range, group_offset, reducers...);
invoke_kernel(kernel, s_item, global_range, global_offset, chunk_offset, group_range, group_offset, reducers...);
};
}

Expand Down
22 changes: 0 additions & 22 deletions test/runtime_deprecation_tests.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -134,27 +134,5 @@ namespace detail {
});
}

TEST_CASE_METHOD(test_utils::runtime_fixture, "Kernels receiving cl::sycl::item<Dims> (deprecated) continue to work", "[handler][deprecated]") {
distr_queue q;

buffer<int, 1> buf1d{{1}};
q.submit([=](handler& cgh) {
accessor acc{buf1d, cgh, celerity::access::one_to_one{}, celerity::read_write, celerity::no_init};
cgh.parallel_for<class UKN(kernel)>(cl::sycl::range<1>{1}, [=](cl::sycl::item<1> id) { acc[id] = 0; });
});

buffer<int, 2> buf2d{{1, 1}};
q.submit([=](handler& cgh) {
accessor acc{buf2d, cgh, celerity::access::one_to_one{}, celerity::read_write, celerity::no_init};
cgh.parallel_for<class UKN(kernel)>(cl::sycl::range<2>{1, 1}, [=](cl::sycl::item<2> id) { acc[id] = 0; });
});

buffer<int, 3> buf3d{{1, 1, 1}};
q.submit([=](handler& cgh) {
accessor acc{buf3d, cgh, celerity::access::one_to_one{}, celerity::read_write, celerity::no_init};
cgh.parallel_for<class UKN(kernel)>(cl::sycl::range<3>{1, 1, 1}, [=](cl::sycl::item<3> id) { acc[id] = 0; });
});
}

} // namespace detail
} // namespace celerity

0 comments on commit 67ccacc

Please sign in to comment.