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

Fix usage of constexpr constants in MSVC #276

Merged
merged 3 commits into from
Nov 19, 2021
Merged

Fix usage of constexpr constants in MSVC #276

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
9a516f0
Fix usage of constexpr constants in MSVC
DuncanMcBain Nov 17, 2021
0620904
MSVC compiler error work around (#275)
hjabird Nov 18, 2021
082a864
Try renaming local accessor class for clarity
DuncanMcBain Nov 19, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 5 additions & 5 deletions samples/vector-addition-examples.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -55,15 +55,15 @@ class VecAddKernelMasked;
class VecAddKernelPredicated;

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

void sumBuffer(sycl::buffer<float, 1> b) {
constexpr auto read = sycl::access::mode::read;
static constexpr auto read = sycl::access::mode::read;
auto h = b.get_access<read>();
auto sum = 0.0f;
for (auto i = 0u; i < b.get_range()[0]; i++) {
Expand All @@ -78,9 +78,9 @@ void sumBuffer(sycl::buffer<float, 1> 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};

Expand Down
21 changes: 9 additions & 12 deletions samples/vector-addition-tiled.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<float, 1, rw, sycl::access::target::local>;
const sycl::range<1> VecSize{N};
const sycl::range<1> TileSize{T};

Expand All @@ -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<read>(h);
auto b = bufB.get_access<read>(h);
auto c = bufC.get_access<dwrite>(h);
sycl::accessor<float, 1, rw, local> tile1(TileSize, h);
sycl::accessor<float, 1, rw, local> tile2(TileSize, h);
local_acc tile1(TileSize, h);
local_acc tile2(TileSize, h);

h.parallel_for<TiledVecAdd>(
sycl::nd_range<1>(VecSize, TileSize), [=](sycl::nd_item<1> i) {
Expand All @@ -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<read>(h);
auto b = bufB.get_access<read>(h);
auto c = bufC.get_access<write>(h);
sycl::accessor<float, 1, rw, local> tile1(TileSize, h);
sycl::accessor<float, 1, rw, local> tile2(TileSize, h);
local_acc tile1(TileSize, h);
local_acc tile2(TileSize, h);

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