Skip to content

Commit a1b98c0

Browse files
joeatoddjoeatodd
authored
Merge pull request #157 from codeplaysoftware/accessor-ctor
Replace buffer.get_access with sycl::accessor ctor
2 parents 7d5824b + a2dfaf8 commit a1b98c0

File tree

15 files changed

+65
-79
lines changed

15 files changed

+65
-79
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,

‎Lesson_Materials/Lecture_10_Data_and_Dependencies/index.html

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -157,7 +157,7 @@
157157
});
158158

159159
gpuQueue.submit([&](sycl::handler &cgh) {
160-
auto acc = buf.get_access(cgh);
160+
sycl::accessor acc{buf, cgh};
161161

162162
cgh.parallel_for&lt;kernel_b&gt;(sycl::range{1024},
163163
[=](sycl::id&lt;1&gt; idx) {
@@ -195,7 +195,7 @@
195195
});
196196

197197
gpuQueue.submit([&](sycl::handler &cgh) {
198-
auto acc = buf.get_access(cgh);
198+
sycl::accessor acc {buf, cgh};
199199

200200
cgh.parallel_for&lt;kernel_b&gt;(sycl::range{1024},
201201
[=](sycl::id&lt;1&gt; idx) {
@@ -225,7 +225,7 @@
225225
sycl::buffer buf {data, sycl::range{1024}};
226226

227227
gpuQueue.submit([&](sycl::handler &cgh) {
228-
<mark>sycl::accessor acc {buf, cgh};</mark>
228+
<mark>sycl::accessor acc{buf, cgh};</mark>
229229

230230
cgh.parallel_for&lt;my_kernel&gt;(sycl::range{1024},
231231
[=](sycl::id&lt;1&gt; idx) {
@@ -234,7 +234,7 @@
234234
});
235235

236236
gpuQueue.submit([&](sycl::handler &cgh) {
237-
<mark>auto acc = buf.get_access(cgh);</mark>
237+
<mark>sycl::accessor acc{buf, cgh};</mark>
238238

239239
cgh.parallel_for&lt;my_kernel&gt;(sycl::range{1024},
240240
[=](sycl::id&lt;1&gt; idx) {
@@ -272,7 +272,7 @@
272272
});
273273

274274
gpuQueue.submit([&](sycl::handler &cgh) {
275-
auto acc = buf.get_access(<mark>cgh</mark>);
275+
sycl::accessor acc {buf, <mark>cgh</mark>};
276276

277277
cgh.parallel_for&lt;my_kernel&gt;(sycl::range{1024},
278278
[=](sycl::id&lt;1&gt; idx) {
@@ -424,7 +424,7 @@
424424
sycl::buffer bufB {dataB, sycl::range{1024}};
425425

426426
gpuQueue.submit([&](sycl::handler &cgh) {
427-
<mark>auto accA = bufA.get_access(cgh);</mark>
427+
<mark>sycl::accessor accA {bufA, cgh};</mark>
428428

429429
cgh.parallel_for&lt;kernel_a&gt;(sycl::range{1024},
430430
[=](sycl::id&lt;1&gt; idx) {
@@ -433,7 +433,7 @@
433433
});
434434

435435
gpuQueue.submit([&](sycl::handler &cgh) {
436-
<mark>auto accB = bufB.get_access(cgh);</mark>
436+
<mark>sycl::accessor accB {bufB, cgh};</mark>
437437

438438
cgh.parallel_for&lt;kernel_b&gt;(sycl::range{1024},
439439
[=](sycl::id&lt;1&gt; idx) {

‎Lesson_Materials/Lecture_11_In_Order_Queue/index.html

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -121,7 +121,7 @@
121121
buf = sycl::buffer(data, sycl::range{1024});
122122

123123
inOrderQueue.submit([&](sycl::handler &cgh){
124-
auto acc = buf.get_access(cgh);
124+
sycl::accessor acc{buf, cgh};
125125

126126
cgh.parallel_for&lt;kernel_a&gt;(sycl::range{1024},
127127
[=](sycl::id&lt;1&gt; idx){
@@ -130,7 +130,7 @@
130130
});
131131

132132
inOrderQueue.submit([&](sycl::handler &cgh){
133-
auto acc = buf.get_access(cgh);
133+
sycl::accessor acc{buf, cgh};
134134

135135
cgh.parallel_for&lt;kernel_b&gt;(sycl::range{1024},
136136
[=](sycl::id&lt;1&gt; idx){

‎Lesson_Materials/Lecture_14_ND_Range_Kernel/index.html

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -346,9 +346,9 @@
346346
buffer&ltfloat, 1&gt bufO(dO.data(), range&lt1&gt(dO.size()));
347347

348348
gpuQueue.submit([&](handler &cgh){
349-
auto inA = bufA.get_access&ltaccess::mode::read&gt(cgh);
350-
auto inB = bufB.get_access&ltaccess::mode::read&gt(cgh);
351-
auto out = bufO.get_access&ltaccess::mode::write&gt(cgh);
349+
sycl::accessor inA{bufA, cgh, sycl::read_only};
350+
sycl::accessor inB{bufB, cgh, sycl::read_only};
351+
sycl::accessor out{bufO, cgh, sycl::write_only};
352352
cgh.parallel_for&ltadd&gt(range&lt1&gt(dA.size()),
353353
[=](id&lt1&gt i){
354354
<mark>out[i] = inA[i] + inB[i];</mark>
@@ -376,9 +376,9 @@
376376
buffer&ltfloat, 1&gt bufO(dO.data(), range&lt1&gt(dO.size()));
377377

378378
gpuQueue.submit([&](handler &cgh){
379-
auto inA = bufA.get_access&ltaccess::mode::read&gt(cgh);
380-
auto inB = bufB.get_access&ltaccess::mode::read&gt(cgh);
381-
auto out = bufO.get_access&ltaccess::mode::write&gt(cgh);
379+
sycl::accessor inA{bufA, cgh, sycl::read_only};
380+
sycl::accessor inB{bufB, cgh, sycl::read_only};
381+
sycl::accessor out{bufO, cgh, sycl::write_only};
382382
cgh.parallel_for&ltadd&gt(rng, [=](id&lt3&gt i){
383383
<mark>auto ptrA = inA.get_pointer();</mark>
384384
<mark>auto ptrB = inB.get_pointer();</mark>

0 commit comments

Comments
 (0)