Skip to content

Commit

Permalink
Add support for SYCL 2020 accessor constructors
Browse files Browse the repository at this point in the history 
along with new enums for modes/targets and deduction from new tags.

A file for compatibility with ComputeCpp 2.5.0 was added since they do
not provide support for some of the features of SYCL 2020 we are using.

matmul, wave_sim and syncing examples were updated to the new accessor
constructor syntax.
  • Loading branch information
@facuMH
facuMH committed Jun 1, 2021
1 parent 127231d commit 001697e
Show file tree
Hide file tree
Showing 8 changed files with 376 additions and 116 deletions.
10 changes: 5 additions & 5 deletions examples/matmul/matmul.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -7,17 +7,17 @@ constexpr size_t MAT_SIZE = 1024;
template <typename T>
void set_identity(celerity::distr_queue queue, celerity::buffer<T, 2> mat) {
queue.submit([=](celerity::handler& cgh) {
auto dw = mat.template get_access<cl::sycl::access::mode::discard_write>(cgh, celerity::access::one_to_one<2>());
celerity::accessor dw{mat, cgh, celerity::access::one_to_one<2>(), cl::sycl::write_only, cl::sycl::no_init};
cgh.parallel_for<class set_identity_kernel>(mat.get_range(), [=](cl::sycl::item<2> item) { dw[item] = item[0] == item[1]; });
});
}

template <typename T>
void multiply(celerity::distr_queue queue, celerity::buffer<T, 2> mat_a, celerity::buffer<T, 2> mat_b, celerity::buffer<T, 2> mat_c) {
queue.submit([=](celerity::handler& cgh) {
auto a = mat_a.template get_access<cl::sycl::access::mode::read>(cgh, celerity::access::slice<2>(1));
auto b = mat_b.template get_access<cl::sycl::access::mode::read>(cgh, celerity::access::slice<2>(0));
auto c = mat_c.template get_access<cl::sycl::access::mode::discard_write>(cgh, celerity::access::one_to_one<2>());
celerity::accessor a{mat_a, cgh, celerity::access::slice<2>(1), cl::sycl::read_only};
celerity::accessor b{mat_b, cgh, celerity::access::slice<2>(0), cl::sycl::read_only};
celerity::accessor c{mat_c, cgh, celerity::access::one_to_one<2>(), cl::sycl::write_only, cl::sycl::no_init};

cgh.parallel_for<class mat_mul>(cl::sycl::range<2>(MAT_SIZE, MAT_SIZE), [=](cl::sycl::item<2> item) {
auto sum = 0.f;
Expand Down Expand Up @@ -58,7 +58,7 @@ int main(int argc, char* argv[]) {
multiply(queue, mat_b_buf, mat_c_buf, mat_a_buf);

queue.submit(celerity::allow_by_ref, [&](celerity::handler& cgh) {
auto result = mat_a_buf.get_access<cl::sycl::access::mode::read, cl::sycl::access::target::host_buffer>(cgh, celerity::access::one_to_one<2>());
celerity::accessor result{mat_a_buf, cgh, celerity::access::one_to_one<2>(), cl::sycl::read_only_host_task};

cgh.host_task(range, [=, &verification_passed](celerity::partition<2> part) {
celerity::experimental::bench::end("main program");
Expand Down
4 changes: 2 additions & 2 deletions examples/syncing/syncing.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,12 +13,12 @@ int main(int argc, char* argv[]) {
std::vector<int> host_buff(N);

q.submit([=](handler& cgh) {
auto b = buff.get_access<cl::sycl::access::mode::discard_write>(cgh, access::one_to_one<1>());
celerity::accessor b{buff, cgh, access::one_to_one<1>(), cl::sycl::write_only, cl::sycl::no_init};
cgh.parallel_for<class mat_mul>(cl::sycl::range<1>(N), [=](cl::sycl::item<1> item) { b[item] = item.get_linear_id(); });
});

q.submit(celerity::allow_by_ref, [=, &host_buff](handler& cgh) {
auto b = buff.get_access<cl::sycl::access::mode::read, cl::sycl::access::target::host_buffer>(cgh, access::fixed<1>({0, N}));
celerity::accessor b{buff, cgh, access::fixed<1>({0, N}), cl::sycl::read_only_host_task};
cgh.host_task(on_master_node, [=, &host_buff] {
std::this_thread::sleep_for(std::chrono::milliseconds(10)); // give the synchronization more time to fail
for(int i = 0; i < N; i++) {
Expand Down
10 changes: 5 additions & 5 deletions examples/wave_sim/wave_sim.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@

void setup_wave(celerity::distr_queue& queue, celerity::buffer<float, 2> u, cl::sycl::float2 center, float amplitude, cl::sycl::float2 sigma) {
queue.submit([=](celerity::handler& cgh) {
auto dw_u = u.get_access<cl::sycl::access::mode::discard_write>(cgh, celerity::access::one_to_one<2>());
celerity::accessor dw_u{u, cgh, celerity::access::one_to_one<2>(), cl::sycl::write_only, cl::sycl::no_init};
cgh.parallel_for<class setup_wave>(u.get_range(), [=, c = center, a = amplitude, s = sigma](cl::sycl::item<2> item) {
const float dx = item[1] - c.x();
const float dy = item[0] - c.y();
Expand All @@ -20,7 +20,7 @@ void setup_wave(celerity::distr_queue& queue, celerity::buffer<float, 2> u, cl::

void zero(celerity::distr_queue& queue, celerity::buffer<float, 2> buf) {
queue.submit([=](celerity::handler& cgh) {
auto dw_buf = buf.get_access<cl::sycl::access::mode::discard_write>(cgh, celerity::access::one_to_one<2>());
celerity::accessor dw_buf{buf, cgh, celerity::access::one_to_one<2>(), cl::sycl::write_only, cl::sycl::no_init};
cgh.parallel_for<class zero>(buf.get_range(), [=](cl::sycl::item<2> item) { dw_buf[item] = 0.f; });
});
}
Expand All @@ -40,8 +40,8 @@ struct update_config {
template <typename T, typename Config, typename KernelName>
void step(celerity::distr_queue& queue, celerity::buffer<T, 2> up, celerity::buffer<T, 2> u, float dt, cl::sycl::float2 delta) {
queue.submit([=](celerity::handler& cgh) {
auto rw_up = up.template get_access<cl::sycl::access::mode::read_write>(cgh, celerity::access::one_to_one<2>());
auto r_u = u.template get_access<cl::sycl::access::mode::read>(cgh, celerity::access::neighborhood<2>(1, 1));
celerity::accessor rw_up{up, cgh, celerity::access::one_to_one<2>(), cl::sycl::read_write};
celerity::accessor r_u{u, cgh, celerity::access::neighborhood<2>(1, 1), cl::sycl::read_only};

const auto size = up.get_range();
cgh.parallel_for<KernelName>(size, [=](cl::sycl::item<2> item) {
Expand Down Expand Up @@ -69,7 +69,7 @@ template <typename T>
void store(celerity::distr_queue& queue, celerity::buffer<T, 2> up, std::vector<std::vector<float>>& result_frames) {
const auto range = up.get_range();
queue.submit(celerity::allow_by_ref, [=, &result_frames](celerity::handler& cgh) {
auto up_r = up.template get_access<cl::sycl::access::mode::read, cl::sycl::access::target::host_buffer>(cgh, celerity::access::fixed<2>{{{}, range}});
celerity::accessor up_r{up, cgh, celerity::access::fixed<2>{{{}, range}}, cl::sycl::read_only_host_task};
cgh.host_task(celerity::on_master_node, [=, &result_frames] {
result_frames.emplace_back();
auto& frame = *result_frames.rbegin();
Expand Down
Loading

0 comments on commit 001697e

Please sign in to comment.