diff --git a/SYCLStream.cpp b/SYCLStream.cpp index 8ab642f..dfaf13c 100644 --- a/SYCLStream.cpp +++ b/SYCLStream.cpp @@ -91,7 +91,7 @@ void SYCLStream::copy() { 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) + cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) { kc[idx] = ka[idx]; }); @@ -107,7 +107,7 @@ void SYCLStream::mul() { 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) + cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) { kb[idx] = scalar * kc[idx]; }); @@ -123,7 +123,7 @@ void SYCLStream::add() 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) + cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) { kc[idx] = ka[idx] + kb[idx]; }); @@ -140,7 +140,7 @@ void SYCLStream::triad() 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) + cgh.parallel_for(range<1>{array_size}, [=](id<1> idx) { ka[idx] = kb[idx] + scalar * kc[idx]; }); @@ -160,7 +160,7 @@ T SYCLStream::dot() auto wg_sum = accessor(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(nd_range<1>(dot_num_groups*dot_wgsize, dot_wgsize), [=](nd_item<1> item) { size_t i = item.get_global_id(0); size_t li = item.get_local_id(0); @@ -201,7 +201,7 @@ void SYCLStream::init_arrays(T initA, T initB, T initC) 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) + cgh.parallel_for(range<1>{array_size}, [=](item<1> item) { auto id = item.get_id(0); ka[id] = initA; diff --git a/SYCLStream.h b/SYCLStream.h index df10946..cb1a45a 100644 --- a/SYCLStream.h +++ b/SYCLStream.h @@ -15,16 +15,6 @@ #define IMPLEMENTATION_STRING "SYCL" -namespace sycl_kernels -{ - template class init; - template class copy; - template class mul; - template class add; - template class triad; - template class dot; -} - template class SYCLStream : public Stream { @@ -39,14 +29,6 @@ class SYCLStream : public Stream cl::sycl::buffer *d_c; cl::sycl::buffer *d_sum; - // SYCL kernel names - typedef sycl_kernels::init init_kernel; - typedef sycl_kernels::copy copy_kernel; - typedef sycl_kernels::mul mul_kernel; - typedef sycl_kernels::add add_kernel; - typedef sycl_kernels::triad triad_kernel; - typedef sycl_kernels::dot dot_kernel; - // NDRange configuration for the dot kernel size_t dot_num_groups; size_t dot_wgsize;