Skip to content
This repository was archived by the owner on Aug 11, 2023. It is now read-only.

Commit dcad1ce

Browse files
DuncanMcBainhjabird
DuncanMcBain
and
hjabird
authored
Fix usage of constexpr constants in MSVC (#276)
Seemingly MSVC cannot consume the constexpr constants in the lambda without them being static as well. Co-authored-by: Hugh Bird <hugh.bird@codeplay.com>
1 parent a86b393 commit dcad1ce

2 files changed

Lines changed: 14 additions & 17 deletions

File tree

‎samples/vector-addition-examples.cpp‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -55,15 +55,15 @@ class VecAddKernelMasked;
5555
class VecAddKernelPredicated;
5656

5757
void zeroBuffer(sycl::buffer<float, 1> b) {
58-
constexpr auto dwrite = sycl::access::mode::discard_write;
58+
static constexpr auto dwrite = sycl::access::mode::discard_write;
5959
auto h = b.get_access<dwrite>();
6060
for (auto i = 0u; i < b.get_range()[0]; i++) {
6161
h[i] = 0.f;
6262
}
6363
}
6464

6565
void sumBuffer(sycl::buffer<float, 1> b) {
66-
constexpr auto read = sycl::access::mode::read;
66+
static constexpr auto read = sycl::access::mode::read;
6767
auto h = b.get_access<read>();
6868
auto sum = 0.0f;
6969
for (auto i = 0u; i < b.get_range()[0]; i++) {
@@ -78,9 +78,9 @@ void sumBuffer(sycl::buffer<float, 1> b) {
7878
* The general flow is that the output buffer is zeroed, the calculation
7979
* scheduled, then the sum printed for each of the functions. */
8080
int main(int argc, char* argv[]) {
81-
constexpr auto read = sycl::access::mode::read;
82-
constexpr auto write = sycl::access::mode::write;
83-
constexpr auto dwrite = sycl::access::mode::discard_write;
81+
static constexpr auto read = sycl::access::mode::read;
82+
static constexpr auto write = sycl::access::mode::write;
83+
static constexpr auto dwrite = sycl::access::mode::discard_write;
8484
constexpr const size_t N = 100000;
8585
const sycl::range<1> VecSize{N};
8686

‎samples/vector-addition-tiled.cpp‎

Lines changed: 9 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -55,10 +55,11 @@ class TiledVecAddDMA;
5555
int main(int argc, char* argv[]) {
5656
constexpr const size_t N = 128000; // this is the total vector size
5757
constexpr const size_t T = 32; // this is the tile size
58-
constexpr auto read = sycl::access::mode::read;
59-
constexpr auto write = sycl::access::mode::write;
60-
constexpr auto rw = sycl::access::mode::read_write;
61-
constexpr auto dwrite = sycl::access::mode::discard_write;
58+
static constexpr auto read = sycl::access::mode::read;
59+
static constexpr auto write = sycl::access::mode::write;
60+
static constexpr auto rw = sycl::access::mode::read_write;
61+
static constexpr auto dwrite = sycl::access::mode::discard_write;
62+
using local_acc = sycl::accessor<float, 1, rw, sycl::access::target::local>;
6263
const sycl::range<1> VecSize{N};
6364
const sycl::range<1> TileSize{T};
6465

@@ -84,13 +85,11 @@ int main(int argc, char* argv[]) {
8485

8586
{
8687
auto cg = [&](sycl::handler& h) {
87-
constexpr auto local = sycl::access::target::local;
88-
8988
auto a = bufA.get_access<read>(h);
9089
auto b = bufB.get_access<read>(h);
9190
auto c = bufC.get_access<dwrite>(h);
92-
sycl::accessor<float, 1, rw, local> tile1(TileSize, h);
93-
sycl::accessor<float, 1, rw, local> tile2(TileSize, h);
91+
local_acc tile1(TileSize, h);
92+
local_acc tile2(TileSize, h);
9493

9594
h.parallel_for<TiledVecAdd>(
9695
sycl::nd_range<1>(VecSize, TileSize), [=](sycl::nd_item<1> i) {
@@ -116,13 +115,11 @@ int main(int argc, char* argv[]) {
116115

117116
{
118117
auto cg = [&](sycl::handler& h) {
119-
constexpr auto local = sycl::access::target::local;
120-
121118
auto a = bufA.get_access<read>(h);
122119
auto b = bufB.get_access<read>(h);
123120
auto c = bufC.get_access<write>(h);
124-
sycl::accessor<float, 1, rw, local> tile1(TileSize, h);
125-
sycl::accessor<float, 1, rw, local> tile2(TileSize, h);
121+
local_acc tile1(TileSize, h);
122+
local_acc tile2(TileSize, h);
126123

127124
h.parallel_for<TiledVecAddDMA>(
128125
sycl::nd_range<1>(VecSize, TileSize), [=](sycl::nd_item<1> i) {

0 commit comments

Comments
 (0)