From 501c61cfbde1fc30d21c5fd05f666860b4fe5485 Mon Sep 17 00:00:00 2001 From: Tom Deakin Date: Tue, 12 Jan 2021 11:14:43 +0000 Subject: [PATCH] [SYCL 2020] update namespace from cl::sycl to sycl:: Also remove the use namespace to make it clear what comes from SYCL --- SYCLStream.cpp | 101 ++++++++++++++++++++++++------------------------- SYCLStream.h | 10 ++--- 2 files changed, 55 insertions(+), 56 deletions(-) diff --git a/SYCLStream.cpp b/SYCLStream.cpp index 8f47304..984d8d2 100644 --- a/SYCLStream.cpp +++ b/SYCLStream.cpp @@ -9,11 +9,9 @@ #include -using namespace cl::sycl; - // Cache list of devices bool cached = false; -std::vector devices; +std::vector devices; void getDeviceList(void); template @@ -26,18 +24,19 @@ SYCLStream::SYCLStream(const int ARRAY_SIZE, const int device_index) if (device_index >= devices.size()) throw std::runtime_error("Invalid device index"); - device dev = devices[device_index]; + + sycl::device dev = devices[device_index]; // Determine sensible dot kernel NDRange configuration if (dev.is_cpu()) { - dot_num_groups = dev.get_info(); - dot_wgsize = dev.get_info() * 2; + dot_num_groups = dev.get_info(); + dot_wgsize = dev.get_info() * 2; } else { - dot_num_groups = dev.get_info() * 4; - dot_wgsize = dev.get_info(); + dot_num_groups = dev.get_info() * 4; + dot_wgsize = dev.get_info(); } // Print out device information @@ -45,7 +44,7 @@ SYCLStream::SYCLStream(const int ARRAY_SIZE, const int device_index) std::cout << "Driver: " << getDeviceDriver(device_index) << std::endl; std::cout << "Reduction kernel config: " << dot_num_groups << " groups of size " << dot_wgsize << std::endl; - queue = new cl::sycl::queue(dev, cl::sycl::async_handler{[&](cl::sycl::exception_list l) + queue = new sycl::queue(dev, sycl::async_handler{[&](sycl::exception_list l) { bool error = false; for(auto e: l) @@ -54,7 +53,7 @@ SYCLStream::SYCLStream(const int ARRAY_SIZE, const int device_index) { std::rethrow_exception(e); } - catch (cl::sycl::exception e) + catch (sycl::exception e) { std::cout << e.what(); error = true; @@ -67,10 +66,10 @@ SYCLStream::SYCLStream(const int ARRAY_SIZE, const int device_index) }}); // Create buffers - d_a = new buffer(array_size); - d_b = new buffer(array_size); - d_c = new buffer(array_size); - d_sum = new buffer(dot_num_groups); + d_a = new sycl::buffer(array_size); + d_b = new sycl::buffer(array_size); + d_c = new sycl::buffer(array_size); + d_sum = new sycl::buffer(dot_num_groups); } template @@ -87,11 +86,11 @@ SYCLStream::~SYCLStream() template void SYCLStream::copy() { - queue->submit([&](handler &cgh) + queue->submit([&](sycl::handler &cgh) { - auto ka = d_a->template get_access(cgh); - auto kc = d_c->template get_access(cgh); - cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) + auto ka = d_a->template get_access(cgh); + auto kc = d_c->template get_access(cgh); + cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) { kc[idx] = ka[idx]; }); @@ -103,11 +102,11 @@ template void SYCLStream::mul() { const T scalar = startScalar; - queue->submit([&](handler &cgh) + queue->submit([&](sycl::handler &cgh) { - auto kb = d_b->template get_access(cgh); - auto kc = d_c->template get_access(cgh); - cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) + auto kb = d_b->template get_access(cgh); + auto kc = d_c->template get_access(cgh); + cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) { kb[idx] = scalar * kc[idx]; }); @@ -118,12 +117,12 @@ void SYCLStream::mul() template void SYCLStream::add() { - queue->submit([&](handler &cgh) + queue->submit([&](sycl::handler &cgh) { - auto ka = d_a->template get_access(cgh); - auto kb = d_b->template get_access(cgh); - auto kc = d_c->template get_access(cgh); - cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) + auto ka = d_a->template get_access(cgh); + auto kb = d_b->template get_access(cgh); + auto kc = d_c->template get_access(cgh); + cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) { kc[idx] = ka[idx] + kb[idx]; }); @@ -135,12 +134,12 @@ template void SYCLStream::triad() { const T scalar = startScalar; - queue->submit([&](handler &cgh) + queue->submit([&](sycl::handler &cgh) { - auto ka = d_a->template get_access(cgh); - auto kb = d_b->template get_access(cgh); - auto kc = d_c->template get_access(cgh); - cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) + auto ka = d_a->template get_access(cgh); + auto kb = d_b->template get_access(cgh); + auto kc = d_c->template get_access(cgh); + cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> idx) { ka[idx] = kb[idx] + scalar * kc[idx]; }); @@ -151,16 +150,16 @@ void SYCLStream::triad() template T SYCLStream::dot() { - queue->submit([&](handler &cgh) + queue->submit([&](sycl::handler &cgh) { - auto ka = d_a->template get_access(cgh); - auto kb = d_b->template get_access(cgh); - auto ksum = d_sum->template get_access(cgh); + auto ka = d_a->template get_access(cgh); + auto kb = d_b->template get_access(cgh); + auto ksum = d_sum->template get_access(cgh); - auto wg_sum = accessor(range<1>(dot_wgsize), cgh); + auto wg_sum = sycl::accessor(sycl::range<1>(dot_wgsize), cgh); size_t N = array_size; - cgh.parallel_for(nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item) + cgh.parallel_for(sycl::nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](sycl::nd_item<1> item) { size_t i = item.get_global_id(0); size_t li = item.get_local_id(0); @@ -173,7 +172,7 @@ T SYCLStream::dot() size_t local_size = item.get_local_range()[0]; for (int offset = local_size / 2; offset > 0; offset /= 2) { - item.barrier(cl::sycl::access::fence_space::local_space); + item.barrier(sycl::access::fence_space::local_space); if (li < offset) wg_sum[li] += wg_sum[li + offset]; } @@ -184,7 +183,7 @@ T SYCLStream::dot() }); T sum = 0.0; - auto h_sum = d_sum->template get_access(); + auto h_sum = d_sum->template get_access(); for (int i = 0; i < dot_num_groups; i++) { sum += h_sum[i]; @@ -196,12 +195,12 @@ T SYCLStream::dot() template void SYCLStream::init_arrays(T initA, T initB, T initC) { - queue->submit([&](handler &cgh) + queue->submit([&](sycl::handler &cgh) { - auto ka = d_a->template get_access(cgh); - auto kb = d_b->template get_access(cgh); - auto kc = d_c->template get_access(cgh); - cgh.parallel_for(range<1>{array_size}, [=](item<1> item) + auto ka = d_a->template get_access(cgh); + auto kb = d_b->template get_access(cgh); + auto kc = d_c->template get_access(cgh); + cgh.parallel_for(sycl::range<1>{array_size}, [=](sycl::item<1> item) { auto id = item.get_id(0); ka[id] = initA; @@ -215,9 +214,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) { - auto _a = d_a->template get_access(); - auto _b = d_b->template get_access(); - auto _c = d_c->template get_access(); + auto _a = d_a->template get_access(); + auto _b = d_b->template get_access(); + auto _c = d_c->template get_access(); for (int i = 0; i < array_size; i++) { a[i] = _a[i]; @@ -229,7 +228,7 @@ void SYCLStream::read_arrays(std::vector& a, std::vector& b, std::vecto void getDeviceList(void) { // Ask SYCL runtime for all devices in system - devices = cl::sycl::device::get_devices(); + devices = sycl::device::get_devices(); cached = true; } @@ -263,7 +262,7 @@ std::string getDeviceName(const int device) if (device < devices.size()) { - name = devices[device].get_info(); + name = devices[device].get_info(); } else { @@ -282,7 +281,7 @@ std::string getDeviceDriver(const int device) if (device < devices.size()) { - driver = devices[device].get_info(); + driver = devices[device].get_info(); } else { diff --git a/SYCLStream.h b/SYCLStream.h index cb1a45a..7b79ad2 100644 --- a/SYCLStream.h +++ b/SYCLStream.h @@ -23,11 +23,11 @@ class SYCLStream : public Stream int array_size; // SYCL objects - cl::sycl::queue *queue; - cl::sycl::buffer *d_a; - cl::sycl::buffer *d_b; - cl::sycl::buffer *d_c; - cl::sycl::buffer *d_sum; + sycl::queue *queue; + sycl::buffer *d_a; + sycl::buffer *d_b; + sycl::buffer *d_c; + sycl::buffer *d_sum; // NDRange configuration for the dot kernel size_t dot_num_groups;