Skip to content

Commit ade5779

Browse files
joeatoddjoeatodd
authored
Merge branch 'main' into imageconv-range-fix
2 parents 1932ae5 + 91aa17c commit ade5779

File tree

24 files changed

+139
-114
lines changed

24 files changed

+139
-114
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: 6 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -60,8 +60,8 @@ TEST_CASE("image_convolution_naive", "image_convolution_reference") {
6060
auto filterWidth = filter.width();
6161
auto halo = filter.half_width();
6262

63-
auto localRange = sycl::range(32, 1);
64-
auto globalRange = sycl::range(inputImgHeight, inputImgWidth);
63+
auto globalRange = sycl::range(inputImgWidth, inputImgHeight);
64+
auto localRange = sycl::range(1, 32);
6565
auto ndRange = sycl::nd_range(globalRange, localRange);
6666

6767
auto inBufRange =
@@ -82,15 +82,14 @@ TEST_CASE("image_convolution_naive", "image_convolution_reference") {
8282
util::benchmark(
8383
[&]() {
8484
myQueue.submit([&](sycl::handler& cgh) {
85-
auto inputAcc = inBuf.get_access<sycl::access::mode::read>(cgh);
86-
auto outputAcc =
87-
outBuf.get_access<sycl::access::mode::write>(cgh);
88-
auto filterAcc =
89-
filterBuf.get_access<sycl::access::mode::read>(cgh);
85+
sycl::accessor inputAcc{inBuf, cgh, sycl::read_only};
86+
sycl::accessor outputAcc{outBuf, cgh, sycl::write_only};
87+
sycl::accessor filterAcc{filterBuf, cgh, sycl::read_only};
9088

9189
cgh.parallel_for<image_convolution>(
9290
ndRange, [=](sycl::nd_item<2> item) {
9391
auto globalId = item.get_global_id();
92+
globalId = sycl::id{globalId[1], globalId[0]};
9493

9594
auto channelsStride = sycl::range(1, channels);
9695
auto haloOffset = sycl::id(halo, halo);

‎Code_Exercises/Exercise_16_Coalesced_Global_Memory/solution.cpp

Lines changed: 5 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -60,8 +60,8 @@ TEST_CASE("image_convolution_coalesced", "coalesced_global_memory_solution") {
6060
auto filterWidth = filter.width();
6161
auto halo = filter.half_width();
6262

63-
auto localRange = sycl::range(32, 1);
64-
auto globalRange = sycl::range(inputImgHeight, inputImgWidth);
63+
auto globalRange = sycl::range(inputImgWidth, inputImgHeight);
64+
auto localRange = sycl::range(1, 32);
6565
auto ndRange = sycl::nd_range(globalRange, localRange);
6666

6767
auto inBufRange =
@@ -80,16 +80,13 @@ TEST_CASE("image_convolution_coalesced", "coalesced_global_memory_solution") {
8080
util::benchmark(
8181
[&]() {
8282
myQueue.submit([&](sycl::handler& cgh) {
83-
auto inputAcc = inBuf.get_access<sycl::access::mode::read>(cgh);
84-
auto outputAcc =
85-
outBuf.get_access<sycl::access::mode::write>(cgh);
86-
auto filterAcc =
87-
filterBuf.get_access<sycl::access::mode::read>(cgh);
83+
sycl::accessor inputAcc{inBuf, cgh, sycl::read_only};
84+
sycl::accessor outputAcc{outBuf, cgh, sycl::write_only};
85+
sycl::accessor filterAcc{filterBuf, cgh, sycl::read_only};
8886

8987
cgh.parallel_for<image_convolution>(
9088
ndRange, [=](sycl::nd_item<2> item) {
9189
auto globalId = item.get_global_id();
92-
globalId = sycl::id{globalId[1], globalId[0]};
9390

9491
auto channelsStride = sycl::range(1, channels);
9592
auto haloOffset = sycl::id(halo, halo);

‎Code_Exercises/Exercise_17_Vectors/solution.cpp

Lines changed: 5 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -61,8 +61,8 @@ TEST_CASE("image_convolution_vectorized", "vectors_solution") {
6161
auto filterWidth = filter.width();
6262
auto halo = filter.half_width();
6363

64-
auto localRange = sycl::range(32, 1);
65-
auto globalRange = sycl::range(inputImgHeight, inputImgWidth);
64+
auto globalRange = sycl::range(inputImgWidth, inputImgHeight);
65+
auto localRange = sycl::range(1, 32);
6666
auto ndRange = sycl::nd_range(globalRange, localRange);
6767

6868
auto inBufRange =
@@ -88,17 +88,13 @@ TEST_CASE("image_convolution_vectorized", "vectors_solution") {
8888
util::benchmark(
8989
[&]() {
9090
myQueue.submit([&](sycl::handler& cgh) {
91-
auto inputAcc =
92-
inBufVec.get_access<sycl::access::mode::read>(cgh);
93-
auto outputAcc =
94-
outBufVec.get_access<sycl::access::mode::write>(cgh);
95-
auto filterAcc =
96-
filterBufVec.get_access<sycl::access::mode::read>(cgh);
91+
sycl::accessor inputAcc{inBufVec, cgh, sycl::read_only};
92+
sycl::accessor outputAcc{outBufVec, cgh, sycl::write_only};
93+
sycl::accessor filterAcc{filterBufVec, cgh, sycl::read_only};
9794

9895
cgh.parallel_for<image_convolution>(
9996
ndRange, [=](sycl::nd_item<2> item) {
10097
auto globalId = item.get_global_id();
101-
globalId = sycl::id{globalId[1], globalId[0]};
10298

10399
auto haloOffset = sycl::id(halo, halo);
104100
auto src = (globalId + haloOffset);

‎Code_Exercises/Exercise_18_Local_Memory_Tiling/solution.cpp

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

9693
auto scratchpad = sycl::accessor<sycl::float4, 2,
9794
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
@@ -86,12 +86,9 @@ TEST_CASE("image_convolution_tiled", "local_memory_tiling_solution") {
8686
util::benchmark(
8787
[&] {
8888
myQueue.submit([&](sycl::handler &cgh) {
89-
auto inputAcc =
90-
inBufVec.get_access<sycl::access::mode::read>(cgh);
91-
auto outputAcc =
92-
outBufVec.get_access<sycl::access::mode::write>(cgh);
93-
auto filterAcc =
94-
filterBufVec.get_access<sycl::access::mode::read>(cgh);
89+
sycl::accessor inputAcc{inBufVec, cgh, sycl::read_only};
90+
sycl::accessor outputAcc{outBufVec, cgh, sycl::write_only};
91+
sycl::accessor filterAcc{filterBufVec, cgh, sycl::read_only};
9592

9693
auto scratchpad = sycl::accessor<sycl::float4, 2,
9794
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){

0 commit comments

Comments
 (0)