diff --git a/samples/vector-addition-examples.cpp b/samples/vector-addition-examples.cpp index 7d1adc7..8fae57b 100644 --- a/samples/vector-addition-examples.cpp +++ b/samples/vector-addition-examples.cpp @@ -55,7 +55,7 @@ class VecAddKernelMasked; class VecAddKernelPredicated; void zeroBuffer(sycl::buffer b) { - constexpr auto dwrite = sycl::access::mode::discard_write; + static constexpr auto dwrite = sycl::access::mode::discard_write; auto h = b.get_access(); for (auto i = 0u; i < b.get_range()[0]; i++) { h[i] = 0.f; @@ -63,7 +63,7 @@ void zeroBuffer(sycl::buffer b) { } void sumBuffer(sycl::buffer b) { - constexpr auto read = sycl::access::mode::read; + static constexpr auto read = sycl::access::mode::read; auto h = b.get_access(); auto sum = 0.0f; for (auto i = 0u; i < b.get_range()[0]; i++) { @@ -78,9 +78,9 @@ void sumBuffer(sycl::buffer b) { * The general flow is that the output buffer is zeroed, the calculation * scheduled, then the sum printed for each of the functions. */ int main(int argc, char* argv[]) { - constexpr auto read = sycl::access::mode::read; - constexpr auto write = sycl::access::mode::write; - constexpr auto dwrite = sycl::access::mode::discard_write; + static constexpr auto read = sycl::access::mode::read; + static constexpr auto write = sycl::access::mode::write; + static constexpr auto dwrite = sycl::access::mode::discard_write; constexpr const size_t N = 100000; const sycl::range<1> VecSize{N}; diff --git a/samples/vector-addition-tiled.cpp b/samples/vector-addition-tiled.cpp index 63efcdb..1fa1c22 100644 --- a/samples/vector-addition-tiled.cpp +++ b/samples/vector-addition-tiled.cpp @@ -55,10 +55,11 @@ class TiledVecAddDMA; int main(int argc, char* argv[]) { constexpr const size_t N = 128000; // this is the total vector size constexpr const size_t T = 32; // this is the tile size - constexpr auto read = sycl::access::mode::read; - constexpr auto write = sycl::access::mode::write; - constexpr auto rw = sycl::access::mode::read_write; - constexpr auto dwrite = sycl::access::mode::discard_write; + static constexpr auto read = sycl::access::mode::read; + static constexpr auto write = sycl::access::mode::write; + static constexpr auto rw = sycl::access::mode::read_write; + static constexpr auto dwrite = sycl::access::mode::discard_write; + using local_acc = sycl::accessor; const sycl::range<1> VecSize{N}; const sycl::range<1> TileSize{T}; @@ -84,13 +85,11 @@ int main(int argc, char* argv[]) { { auto cg = [&](sycl::handler& h) { - constexpr auto local = sycl::access::target::local; - auto a = bufA.get_access(h); auto b = bufB.get_access(h); auto c = bufC.get_access(h); - sycl::accessor tile1(TileSize, h); - sycl::accessor tile2(TileSize, h); + local_acc tile1(TileSize, h); + local_acc tile2(TileSize, h); h.parallel_for( sycl::nd_range<1>(VecSize, TileSize), [=](sycl::nd_item<1> i) { @@ -116,13 +115,11 @@ int main(int argc, char* argv[]) { { auto cg = [&](sycl::handler& h) { - constexpr auto local = sycl::access::target::local; - auto a = bufA.get_access(h); auto b = bufB.get_access(h); auto c = bufC.get_access(h); - sycl::accessor tile1(TileSize, h); - sycl::accessor tile2(TileSize, h); + local_acc tile1(TileSize, h); + local_acc tile2(TileSize, h); h.parallel_for( sycl::nd_range<1>(VecSize, TileSize), [=](sycl::nd_item<1> i) {