Skip to content

Commit dfeb909

Browse files
joeatoddjoeatodd
committed
Replace get_access with CTAD accessor ctor
1 parent 7d5824b commit dfeb909

File tree

12 files changed

+50
-64
lines changed

12 files changed

+50
-64
lines changed

‎Code_Exercises/Exercise_06_Vector_Add/README.md

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,6 @@ Remember to ensure the `range` provided to the buffer is the size of the arrays.
2222
### 3. ) Create accessors
2323

2424
Create `accessor`s to each of the `buffer`s within the command group function.
25-
The simplest way to do this is to call `get_access` on the `buffer`s.
2625

2726
### 4. ) Write the kernel function
2827

‎Code_Exercises/Exercise_06_Vector_Add/solution.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -44,9 +44,9 @@ TEST_CASE("vector_add", "vector_add_solution") {
4444

4545
defaultQueue
4646
.submit([&](sycl::handler& cgh) {
47-
auto accA = bufA.get_access<sycl::access::mode::read>(cgh);
48-
auto accB = bufB.get_access<sycl::access::mode::read>(cgh);
49-
auto accR = bufR.get_access<sycl::access::mode::write>(cgh);
47+
sycl::accessor accA{bufA, cgh, sycl::read_only};
48+
sycl::accessor accB{bufB, cgh, sycl::read_only};
49+
sycl::accessor accR{bufR, cgh, sycl::write_only};
5050

5151
cgh.parallel_for<vector_add>(
5252
sycl::range{dataSize},

‎Code_Exercises/Exercise_10_Managing_Dependencies/solution.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -63,35 +63,35 @@ TEST_CASE("buffer_accessor_diamond", "managing_dependencies_solution") {
6363
auto bufOut = sycl::buffer{out, sycl::range{dataSize}};
6464

6565
defaultQueue.submit([&](sycl::handler& cgh) {
66-
auto acc = bufInA.get_access<sycl::access::mode::read_write>(cgh);
66+
sycl::accessor acc{bufInA, cgh, sycl::read_write};
6767

6868
cgh.parallel_for<kernel_a_1>(sycl::range{dataSize}, [=](sycl::id<1> idx) {
6969
acc[idx] = acc[idx] * 2.0f;
7070
});
7171
});
7272

7373
defaultQueue.submit([&](sycl::handler& cgh) {
74-
auto accIn = bufInA.get_access<sycl::access::mode::read>(cgh);
75-
auto accOut = bufInB.get_access<sycl::access::mode::write>(cgh);
74+
sycl::accessor accIn{bufInA, cgh, sycl::read_only};
75+
sycl::accessor accOut{bufInB, cgh, sycl::write_only};
7676

7777
cgh.parallel_for<kernel_b_1>(sycl::range{dataSize}, [=](sycl::id<1> idx) {
7878
accOut[idx] += accIn[idx];
7979
});
8080
});
8181

8282
defaultQueue.submit([&](sycl::handler& cgh) {
83-
auto accIn = bufInA.get_access<sycl::access::mode::read>(cgh);
84-
auto accOut = bufInC.get_access<sycl::access::mode::write>(cgh);
83+
sycl::accessor accIn{bufInA, cgh, sycl::read_only};
84+
sycl::accessor accOut{bufInC, cgh, sycl::write_only};
8585

8686
cgh.parallel_for<kernel_c_1>(sycl::range{dataSize}, [=](sycl::id<1> idx) {
8787
accOut[idx] -= accIn[idx];
8888
});
8989
});
9090

9191
defaultQueue.submit([&](sycl::handler& cgh) {
92-
auto accInA = bufInB.get_access<sycl::access::mode::read>(cgh);
93-
auto accInB = bufInC.get_access<sycl::access::mode::read>(cgh);
94-
auto accOut = bufOut.get_access<sycl::access::mode::write>(cgh);
92+
sycl::accessor accInA{bufInB, cgh, sycl::read_only};
93+
sycl::accessor accInB{bufInC, cgh, sycl::read_only};
94+
sycl::accessor accOut{bufOut, cgh, sycl::write_only};
9595

9696
cgh.parallel_for<kernel_d_1>(sycl::range{dataSize}, [=](sycl::id<1> idx) {
9797
accOut[idx] = accInA[idx] + accInB[idx];

‎Code_Exercises/Exercise_11_In_Order_Queue/solution.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -67,35 +67,35 @@ TEST_CASE("buffer_accessor_in_order_queue", "in_order_queue_solution") {
6767
auto bufOut = sycl::buffer{out, sycl::range{dataSize}};
6868

6969
inOrderQueue.submit([&](sycl::handler& cgh) {
70-
auto acc = bufInA.get_access<sycl::access::mode::read_write>(cgh);
70+
sycl::accessor acc{bufInA, cgh, sycl::read_write};
7171

7272
cgh.parallel_for<kernel_a_1>(sycl::range{dataSize}, [=](sycl::id<1> idx) {
7373
acc[idx] = acc[idx] * 2.0f;
7474
});
7575
});
7676

7777
inOrderQueue.submit([&](sycl::handler& cgh) {
78-
auto accIn = bufInA.get_access<sycl::access::mode::read>(cgh);
79-
auto accOut = bufInB.get_access<sycl::access::mode::write>(cgh);
78+
sycl::accessor accIn{bufInA, cgh, sycl::read_only};
79+
sycl::accessor accOut{bufInB, cgh, sycl::write_only};
8080

8181
cgh.parallel_for<kernel_b_1>(sycl::range{dataSize}, [=](sycl::id<1> idx) {
8282
accOut[idx] += accIn[idx];
8383
});
8484
});
8585

8686
inOrderQueue.submit([&](sycl::handler& cgh) {
87-
auto accIn = bufInA.get_access<sycl::access::mode::read>(cgh);
88-
auto accOut = bufInC.get_access<sycl::access::mode::write>(cgh);
87+
sycl::accessor accIn{bufInA, cgh, sycl::read_only};
88+
sycl::accessor accOut{bufInC, cgh, sycl::write_only};
8989

9090
cgh.parallel_for<kernel_c_1>(sycl::range{dataSize}, [=](sycl::id<1> idx) {
9191
accOut[idx] -= accIn[idx];
9292
});
9393
});
9494

9595
inOrderQueue.submit([&](sycl::handler& cgh) {
96-
auto accInA = bufInB.get_access<sycl::access::mode::read>(cgh);
97-
auto accInB = bufInC.get_access<sycl::access::mode::read>(cgh);
98-
auto accOut = bufOut.get_access<sycl::access::mode::write>(cgh);
96+
sycl::accessor accInA{bufInB, cgh, sycl::read_only};
97+
sycl::accessor accInB{bufInC, cgh, sycl::read_only};
98+
sycl::accessor accOut{bufOut, cgh, sycl::write_only};
9999

100100
cgh.parallel_for<kernel_d_1>(sycl::range{dataSize}, [=](sycl::id<1> idx) {
101101
accOut[idx] = accInA[idx] + accInB[idx];

‎Code_Exercises/Exercise_12_Temporary_Data/solution.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -61,17 +61,17 @@ TEST_CASE("buffer_accessor_temporary_data", "temporary_data_solution") {
6161
bufOut.set_final_data(out);
6262

6363
gpuQueue.submit([&](sycl::handler& cgh) {
64-
auto accIn = bufIn.get_access<sycl::access::mode::read>(cgh);
65-
auto accOut = bufInt.get_access<sycl::access::mode::write>(cgh);
64+
sycl::accessor accIn{bufIn, cgh, sycl::read_only};
65+
sycl::accessor accOut{bufInt, cgh, sycl::write_only};
6666

6767
cgh.parallel_for<kernel_a_1>(sycl::range{dataSize}, [=](sycl::id<1> idx) {
6868
accOut[idx] = accIn[idx] * 8.0f;
6969
});
7070
});
7171

7272
gpuQueue.submit([&](sycl::handler& cgh) {
73-
auto accIn = bufInt.get_access<sycl::access::mode::read>(cgh);
74-
auto accOut = bufOut.get_access<sycl::access::mode::write>(cgh);
73+
sycl::accessor accIn{bufInt, cgh, sycl::read_only};
74+
sycl::accessor accOut{bufOut, cgh, sycl::write_only};
7575

7676
cgh.parallel_for<kernel_b_1>(sycl::range{dataSize}, [=](sycl::id<1> idx) {
7777
accOut[idx] = accIn[idx] / 2.0f;

‎Code_Exercises/Exercise_13_Load_Balancing/solution.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -101,19 +101,19 @@ TEST_CASE("load_balancing", "load_balancing_solution") {
101101
sycl::buffer{r + dataSizeFirst, sycl::range{dataSizeSecond}};
102102

103103
Q1.submit([&](sycl::handler& cgh) {
104-
auto accA = bufFirstA.get_access<sycl::access::mode::read>(cgh);
105-
auto accB = bufFirstB.get_access<sycl::access::mode::read>(cgh);
106-
auto accR = bufFirstR.get_access<sycl::access::mode::write>(cgh);
104+
sycl::accessor accA{bufFirstA, cgh, sycl::read_only};
105+
sycl::accessor accB{bufFirstB, cgh, sycl::read_only};
106+
sycl::accessor accR{bufFirstR, cgh, sycl::write_only};
107107

108108
cgh.parallel_for<vector_add_first>(
109109
sycl::range{dataSizeFirst},
110110
[=](sycl::id<1> idx) { accR[idx] = accA[idx] + accB[idx]; });
111111
});
112112

113113
Q2.submit([&](sycl::handler& cgh) {
114-
auto accA = bufSecondA.get_access<sycl::access::mode::read>(cgh);
115-
auto accB = bufSecondB.get_access<sycl::access::mode::read>(cgh);
116-
auto accR = bufSecondR.get_access<sycl::access::mode::write>(cgh);
114+
sycl::accessor accA{bufSecondA, cgh, sycl::read_only};
115+
sycl::accessor accB{bufSecondB, cgh, sycl::read_only};
116+
sycl::accessor accR{bufSecondR, cgh, sycl::write_only};
117117

118118
cgh.parallel_for<vector_add_second>(
119119
sycl::range{dataSizeSecond},

‎Code_Exercises/Exercise_14_ND_Range_Kernel/solution.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -44,9 +44,9 @@ TEST_CASE("range_kernel_with_item", "nd_range_kernel_solution") {
4444
auto bufR = sycl::buffer{r, sycl::range{dataSize}};
4545

4646
gpuQueue.submit([&](sycl::handler& cgh) {
47-
auto accA = bufA.get_access<sycl::access::mode::read>(cgh);
48-
auto accB = bufB.get_access<sycl::access::mode::read>(cgh);
49-
auto accR = bufR.get_access<sycl::access::mode::write>(cgh);
47+
sycl::accessor accA{bufA, cgh, sycl::read_only};
48+
sycl::accessor accB{bufB, cgh, sycl::read_only};
49+
sycl::accessor accR{bufR, cgh, sycl::write_only};
5050

5151
cgh.parallel_for<vector_add_1>(
5252
sycl::range{dataSize}, [=](sycl::item<1> itm) {
@@ -90,9 +90,9 @@ TEST_CASE("nd_range_kernel", "nd_range_kernel_solution") {
9090
auto bufR = sycl::buffer{r, sycl::range{dataSize}};
9191

9292
gpuQueue.submit([&](sycl::handler& cgh) {
93-
auto accA = bufA.get_access<sycl::access::mode::read_write>(cgh);
94-
auto accB = bufB.get_access<sycl::access::mode::read_write>(cgh);
95-
auto accR = bufR.get_access<sycl::access::mode::read_write>(cgh);
93+
sycl::accessor accA{bufA, cgh, sycl::read_write};
94+
sycl::accessor accB{bufB, cgh, sycl::read_write};
95+
sycl::accessor accR{bufR, cgh, sycl::read_write};
9696

9797
auto ndRange =
9898
sycl::nd_range{sycl::range{dataSize}, sycl::range{workGroupSize}};

‎Code_Exercises/Exercise_15_Image_Convolution/reference.cpp

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -77,11 +77,9 @@ TEST_CASE("image_convolution_naive", "image_convolution_reference") {
7777
util::benchmark(
7878
[&]() {
7979
myQueue.submit([&](sycl::handler& cgh) {
80-
auto inputAcc = inBuf.get_access<sycl::access::mode::read>(cgh);
81-
auto outputAcc =
82-
outBuf.get_access<sycl::access::mode::write>(cgh);
83-
auto filterAcc =
84-
filterBuf.get_access<sycl::access::mode::read>(cgh);
80+
sycl::accessor inputAcc{inBuf, cgh, sycl::read_only};
81+
sycl::accessor outputAcc{outBuf, cgh, sycl::write_only};
82+
sycl::accessor filterAcc{filterBuf, cgh, sycl::read_only};
8583

8684
cgh.parallel_for<image_convolution>(
8785
ndRange, [=](sycl::nd_item<2> item) {

‎Code_Exercises/Exercise_16_Coalesced_Global_Memory/solution.cpp

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -77,11 +77,9 @@ TEST_CASE("image_convolution_coalesced", "coalesced_global_memory_solution") {
7777
util::benchmark(
7878
[&]() {
7979
myQueue.submit([&](sycl::handler& cgh) {
80-
auto inputAcc = inBuf.get_access<sycl::access::mode::read>(cgh);
81-
auto outputAcc =
82-
outBuf.get_access<sycl::access::mode::write>(cgh);
83-
auto filterAcc =
84-
filterBuf.get_access<sycl::access::mode::read>(cgh);
80+
sycl::accessor inputAcc{inBuf, cgh, sycl::read_only};
81+
sycl::accessor outputAcc{outBuf, cgh, sycl::write_only};
82+
sycl::accessor filterAcc{filterBuf, cgh, sycl::read_only};
8583

8684
cgh.parallel_for<image_convolution>(
8785
ndRange, [=](sycl::nd_item<2> item) {

‎Code_Exercises/Exercise_17_Vectors/solution.cpp

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -85,12 +85,9 @@ TEST_CASE("image_convolution_vectorized", "vectors_solution") {
8585
util::benchmark(
8686
[&]() {
8787
myQueue.submit([&](sycl::handler& cgh) {
88-
auto inputAcc =
89-
inBufVec.get_access<sycl::access::mode::read>(cgh);
90-
auto outputAcc =
91-
outBufVec.get_access<sycl::access::mode::write>(cgh);
92-
auto filterAcc =
93-
filterBufVec.get_access<sycl::access::mode::read>(cgh);
88+
sycl::accessor inputAcc{inBufVec, cgh, sycl::read_only};
89+
sycl::accessor outputAcc{outBufVec, cgh, sycl::write_only};
90+
sycl::accessor filterAcc{filterBufVec, cgh, sycl::read_only};
9491

9592
cgh.parallel_for<image_convolution>(
9693
ndRange, [=](sycl::nd_item<2> item) {

‎Code_Exercises/Exercise_18_Local_Memory_Tiling/solution.cpp

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -83,12 +83,9 @@ TEST_CASE("image_convolution_tiled", "local_memory_tiling_solution") {
8383
util::benchmark(
8484
[&] {
8585
myQueue.submit([&](sycl::handler &cgh) {
86-
auto inputAcc =
87-
inBufVec.get_access<sycl::access::mode::read>(cgh);
88-
auto outputAcc =
89-
outBufVec.get_access<sycl::access::mode::write>(cgh);
90-
auto filterAcc =
91-
filterBufVec.get_access<sycl::access::mode::read>(cgh);
86+
sycl::accessor inputAcc{inBufVec, cgh, sycl::read_only};
87+
sycl::accessor outputAcc{outBufVec, cgh, sycl::write_only};
88+
sycl::accessor filterAcc{filterBufVec, cgh, sycl::read_only};
9289

9390
auto scratchpad = sycl::accessor<sycl::float4, 2,
9491
sycl::access::mode::read_write,

‎Code_Exercises/Exercise_19_Work_Group_Sizes/solution.cpp

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -83,12 +83,9 @@ TEST_CASE("image_convolution_tiled", "local_memory_tiling_solution") {
8383
util::benchmark(
8484
[&] {
8585
myQueue.submit([&](sycl::handler &cgh) {
86-
auto inputAcc =
87-
inBufVec.get_access<sycl::access::mode::read>(cgh);
88-
auto outputAcc =
89-
outBufVec.get_access<sycl::access::mode::write>(cgh);
90-
auto filterAcc =
91-
filterBufVec.get_access<sycl::access::mode::read>(cgh);
86+
sycl::accessor inputAcc{inBufVec, cgh, sycl::read_only};
87+
sycl::accessor outputAcc{outBufVec, cgh, sycl::write_only};
88+
sycl::accessor filterAcc{filterBufVec, cgh, sycl::read_only};
9289

9390
auto scratchpad = sycl::accessor<sycl::float4, 2,
9491
sycl::access::mode::read_write,

0 commit comments

Comments
 (0)