From 9a516f0d13bc16d21fde990324ee8b1cf96e1fa8 Mon Sep 17 00:00:00 2001 From: Duncan McBain Date: Wed, 17 Nov 2021 15:35:49 +0000 Subject: [PATCH 1/3] Fix usage of constexpr constants in MSVC Seemingly MSVC cannot consume the constexpr constants in the lambda without them being static as well. --- samples/vector-addition-examples.cpp | 10 +++++----- samples/vector-addition-tiled.cpp | 12 ++++++------ 2 files changed, 11 insertions(+), 11 deletions(-) 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..32fc85f 100644 --- a/samples/vector-addition-tiled.cpp +++ b/samples/vector-addition-tiled.cpp @@ -55,10 +55,10 @@ 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; const sycl::range<1> VecSize{N}; const sycl::range<1> TileSize{T}; @@ -84,7 +84,7 @@ int main(int argc, char* argv[]) { { auto cg = [&](sycl::handler& h) { - constexpr auto local = sycl::access::target::local; + static constexpr auto local = sycl::access::target::local; auto a = bufA.get_access(h); auto b = bufB.get_access(h); @@ -116,7 +116,7 @@ int main(int argc, char* argv[]) { { auto cg = [&](sycl::handler& h) { - constexpr auto local = sycl::access::target::local; + static constexpr auto local = sycl::access::target::local; auto a = bufA.get_access(h); auto b = bufB.get_access(h); From 062090459d679909605810efbca3326fc320fd4a Mon Sep 17 00:00:00 2001 From: HJA Bird <9040797+hjabird@users.noreply.github.com> Date: Thu, 18 Nov 2021 11:53:13 +0000 Subject: [PATCH 2/3] MSVC compiler error work around (#275) Co-authored-by: Hugh Bird --- samples/vector-addition-tiled.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/samples/vector-addition-tiled.cpp b/samples/vector-addition-tiled.cpp index 32fc85f..7ef94a7 100644 --- a/samples/vector-addition-tiled.cpp +++ b/samples/vector-addition-tiled.cpp @@ -84,13 +84,13 @@ int main(int argc, char* argv[]) { { auto cg = [&](sycl::handler& h) { - static 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); + sycl::accessor tile1(TileSize, + h); + sycl::accessor tile2(TileSize, + h); h.parallel_for( sycl::nd_range<1>(VecSize, TileSize), [=](sycl::nd_item<1> i) { @@ -116,13 +116,13 @@ int main(int argc, char* argv[]) { { auto cg = [&](sycl::handler& h) { - static 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); + sycl::accessor tile1(TileSize, + h); + sycl::accessor tile2(TileSize, + h); h.parallel_for( sycl::nd_range<1>(VecSize, TileSize), [=](sycl::nd_item<1> i) { From 082a864a62d6967aef68cfe0a81456ae1b5ed6ab Mon Sep 17 00:00:00 2001 From: Duncan McBain Date: Fri, 19 Nov 2021 13:34:49 +0000 Subject: [PATCH 3/3] Try renaming local accessor class for clarity --- samples/vector-addition-tiled.cpp | 13 +++++-------- 1 file changed, 5 insertions(+), 8 deletions(-) diff --git a/samples/vector-addition-tiled.cpp b/samples/vector-addition-tiled.cpp index 7ef94a7..1fa1c22 100644 --- a/samples/vector-addition-tiled.cpp +++ b/samples/vector-addition-tiled.cpp @@ -59,6 +59,7 @@ int main(int argc, char* argv[]) { 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}; @@ -87,10 +88,8 @@ int main(int argc, char* argv[]) { 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) { @@ -119,10 +118,8 @@ int main(int argc, char* argv[]) { 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) {