From ae8bd6081be7d8c116e8f345e7f833bd0317b22c Mon Sep 17 00:00:00 2001 From: Tom Deakin Date: Mon, 8 Feb 2021 13:43:35 +0000 Subject: [PATCH] [SYCL 2020] Use constructor initaliser list to allocate buffers - no need to use pointers --- SYCLStream.cpp | 55 ++++++++++++++++++++++++-------------------------- SYCLStream.h | 11 ++++++---- 2 files changed, 33 insertions(+), 33 deletions(-) diff --git a/SYCLStream.cpp b/SYCLStream.cpp index 7765352..d766996 100644 --- a/SYCLStream.cpp +++ b/SYCLStream.cpp @@ -16,12 +16,15 @@ void getDeviceList(void); template SYCLStream::SYCLStream(const int ARRAY_SIZE, const int device_index) +: array_size {ARRAY_SIZE}, + d_a {ARRAY_SIZE}, + d_b {ARRAY_SIZE}, + d_c {ARRAY_SIZE}, + d_sum {1} { if (!cached) getDeviceList(); - array_size = ARRAY_SIZE; - if (device_index >= devices.size()) throw std::runtime_error("Invalid device index"); @@ -55,13 +58,7 @@ SYCLStream::SYCLStream(const int ARRAY_SIZE, const int device_index) // No longer need list of devices devices.clear(); cached = true; - - // Create buffers - // Only in the constructor at runtime do we know the size, so need to use (smart) pointers - d_a = std::make_unique>(array_size); - d_b = std::make_unique>(array_size); - d_c = std::make_unique>(array_size); - d_sum = std::make_unique>(1); + } @@ -71,8 +68,8 @@ void SYCLStream::copy() { queue->submit([&](sycl::handler &cgh) { - sycl::accessor ka {*d_a, cgh, sycl::read_only}; - sycl::accessor kc {*d_c, cgh, sycl::write_only}; + sycl::accessor ka {d_a, cgh, sycl::read_only}; + sycl::accessor kc {d_c, cgh, sycl::write_only}; cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) { kc[idx] = ka[idx]; @@ -87,8 +84,8 @@ void SYCLStream::mul() const T scalar = startScalar; queue->submit([&](sycl::handler &cgh) { - sycl::accessor kb {*d_b, cgh, sycl::write_only}; - sycl::accessor kc {*d_c, cgh, sycl::read_only}; + sycl::accessor kb {d_b, cgh, sycl::write_only}; + sycl::accessor kc {d_c, cgh, sycl::read_only}; cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) { kb[idx] = scalar * kc[idx]; @@ -102,9 +99,9 @@ void SYCLStream::add() { queue->submit([&](sycl::handler &cgh) { - sycl::accessor ka {*d_a, cgh, sycl::read_only}; - sycl::accessor kb {*d_b, cgh, sycl::read_only}; - sycl::accessor kc {*d_c, cgh, sycl::write_only}; + sycl::accessor ka {d_a, cgh, sycl::read_only}; + sycl::accessor kb {d_b, cgh, sycl::read_only}; + sycl::accessor kc {d_c, cgh, sycl::write_only}; cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) { kc[idx] = ka[idx] + kb[idx]; @@ -119,9 +116,9 @@ void SYCLStream::triad() const T scalar = startScalar; queue->submit([&](sycl::handler &cgh) { - sycl::accessor ka {*d_a, cgh, sycl::write_only}; - sycl::accessor kb {*d_b, cgh, sycl::read_only}; - sycl::accessor kc {*d_c, cgh, sycl::read_only}; + sycl::accessor ka {d_a, cgh, sycl::write_only}; + sycl::accessor kb {d_b, cgh, sycl::read_only}; + sycl::accessor kc {d_c, cgh, sycl::read_only}; cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) { ka[idx] = kb[idx] + scalar * kc[idx]; @@ -136,12 +133,12 @@ T SYCLStream::dot() queue->submit([&](sycl::handler &cgh) { - sycl::accessor ka {*d_a, cgh, sycl::read_only}; - sycl::accessor kb {*d_b, cgh, sycl::read_only}; + sycl::accessor ka {d_a, cgh, sycl::read_only}; + sycl::accessor kb {d_b, cgh, sycl::read_only}; cgh.parallel_for(sycl::range<1>{array_size}, // Reduction object, to perform summation - initialises the result to zero - sycl::reduction(*d_sum, cgh, std::plus(), sycl::property::reduction::initialize_to_identity); + sycl::reduction(d_sum, cgh, std::plus(), sycl::property::reduction::initialize_to_identity); [=](sycl::id<1> idx, auto& sum) { sum += ka[idx] * kb[idx]; @@ -151,7 +148,7 @@ T SYCLStream::dot() // Get access on the host, and return a copy of the data (single number) // This will block until the result is available, so no need to wait on the queue. - sycl::host_accessor result {*d_sum, sycl::read_only}; + sycl::host_accessor result {d_sum, sycl::read_only}; return result[0]; } @@ -161,9 +158,9 @@ void SYCLStream::init_arrays(T initA, T initB, T initC) { queue->submit([&](sycl::handler &cgh) { - sycl::accessor ka {*d_a, cgh, sycl::write_only, sycl::no_init}; - sycl::accessor kb {*d_b, cgh, sycl::write_only, sycl::no_init}; - sycl::accessor kc {*d_c, cgh, sycl::write_only, sycl::no_init}; + sycl::accessor ka {d_a, cgh, sycl::write_only, sycl::no_init}; + sycl::accessor kb {d_b, cgh, sycl::write_only, sycl::no_init}; + sycl::accessor kc {d_c, cgh, sycl::write_only, sycl::no_init}; cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) { @@ -179,9 +176,9 @@ void SYCLStream::init_arrays(T initA, T initB, T initC) template void SYCLStream::read_arrays(std::vector& a, std::vector& b, std::vector& c) { - sycl::host_accessor _a {*d_a, sycl::read_only}; - sycl::host_accessor _b {*d_b, sycl::read_only}; - sycl::host_accessor _c {*d_c, sycl::read_only}; + sycl::host_accessor _a {d_a, sycl::read_only}; + sycl::host_accessor _b {d_b, sycl::read_only}; + sycl::host_accessor _c {d_c, sycl::read_only}; for (int i = 0; i < array_size; i++) { a[i] = _a[i]; diff --git a/SYCLStream.h b/SYCLStream.h index bfe9ae5..0c73594 100644 --- a/SYCLStream.h +++ b/SYCLStream.h @@ -24,11 +24,14 @@ class SYCLStream : public Stream int array_size; // SYCL objects + // Queue is a pointer because we allow device selection std::unique_ptr queue; - std::unique_ptr> d_a; - std::unique_ptr> d_b; - std::unique_ptr> d_c; - std::unique_ptr> d_sum; + + // Buffers + sycl::buffer d_a; + sycl::buffer d_b; + sycl::buffer d_c; + sycl::buffer d_sum; public: