Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Problem with exercise 16 #326

Open
wants to merge 2 commits into
base: main
Choose a base branch
from
Open

Problem with exercise 16 #326

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
9062772
Changed single_task to parallel_for
lrpablo Dec 27, 2023
96ee457
The code uses items, but the README.md uses ids. I changed the README…
lrpablo Jan 9, 2024
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
4 changes: 2 additions & 2 deletions Code_Exercises/Exercise_16_Coalesced_Global_Memory/README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,8 +17,8 @@ global memory access patterns in your kernel are coalesced.
Consider two alternative ways to linearize the global id:

```
auto rowMajorLinearId = (idx[1] * width) + idx[0]; // row-major
auto columnMajorLinearId = (idx[0] * height) + idx[1]; // column-major
auto rowMajorLinearId = sycl::id(globalId[0], globalId[1]);
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't think this is the original intent, here, since it talks about the linear id presumably the intent is to linearise it manually (to demonstrate the difference in striding). Can we change the other side instead to better match this?

auto columnMajorLinearId = sycl::id(globalId[1], globalId[0]);
```

Try using both of these and compare the execution time of each.
Expand Down
6 changes: 6 additions & 0 deletions Code_Exercises/Exercise_16_Coalesced_Global_Memory/solution.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -79,6 +79,12 @@ TEST_CASE("image_convolution_coalesced", "coalesced_global_memory_solution") {
ndRange, [=](sycl::nd_item<2> item) {
auto globalId = item.get_global_id();

auto rowMajorLinearId = sycl::id(globalId[0], globalId[1]);
auto columnMajorLinearId = sycl::id(globalId[1], globalId[0]);
Comment on lines +82 to +83
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Similarly these are 2d ids, not linear ones


// Set row major or column major
globalId = rowMajorLinearId;

auto channelsStride = sycl::range(1, channels);
auto haloOffset = sycl::id(halo, halo);
auto src = (globalId + haloOffset) * channelsStride;
Expand Down
Binary file modified Lesson_Materials/Lecture_04_Handling_Errors/Lecture_04_Handling_Errors.pdf
View file
Open in desktop
Binary file not shown.
14 changes: 6 additions & 8 deletions Lesson_Materials/Lecture_04_Handling_Errors/index.html
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,7 +91,7 @@

/* Synchronous code */

cgh.single_task&lt;add&gt;(bufO.get_range(), [=](id&lt;1&gt; i) {
cgh.parallel_for&lt;add&gt;(bufO.get_range(), [=](id&lt;1&gt; i) {

/* Asynchronous code */

Expand Down Expand Up @@ -166,7 +166,7 @@
auto inB = accessor{bufB, cgh, read_only};
auto out = accessor{bufO, cgh, write_only};

cgh.single_task&lt;add&gt;(bufO.get_range(), [=](id&lt;1&gt; i) {
cgh.parallel_for&lt;add&gt;(bufO.get_range(), [=](id&lt;1&gt; i) {
out[i] = inA[i] + inB[i];
});
}).wait();
Expand All @@ -190,7 +190,6 @@
std::vector&lt;float&gt; dA{ 7, 5, 16, 8 }, dB{ 8, 16, 5, 7 }, dO{ 0, 0, 0, 0 };
try{
queue gpuQueue(gpu_selector{}, <mark>async_handler{}</mark>);

buffer bufA{dA};
buffer bufB{dB};
buffer bufO{dO};
Expand All @@ -200,7 +199,7 @@
auto inB = accessor{bufB, cgh, read_only};
auto out = accessor{bufO, cgh, write_only};

cgh.single_task&lt;add&gt;(bufO.get_range(), [=](id&lt;1&gt; i) {
cgh.parallel_for&lt;add&gt;(bufO.get_range(), [=](id&lt;1&gt; i) {
out[i] = inA[i] + inB[i];
});
}).wait();
Expand All @@ -213,8 +212,7 @@
<div class="bottom-bullets" data-markdown>
* Asynchronous errors errors that may have occurred will be thrown after a command group has been submitted to a `queue`.
* To handle these errors you must provide an async handler when constructing the queue object.
* Then you must also call the `throw_asynchronous` or `wait_and_throw` member functions of the `queue`
class.
* Then you must also call the `throw_asynchronous` or `wait_and_throw` member functions of the `queue` class.
* This will pass the exceptions to the async handler in the user thread so they can be thrown.
</div>
</section>
Expand All @@ -240,7 +238,7 @@
auto inB = accessor{bufB, cgh, read_only};
auto out = accessor{bufO, cgh, write_only};

cgh.single_task&lt;add&gt;(bufO.get_range(), [=](id&lt;1&gt; i) {
cgh.parallel_for&lt;add&gt;(bufO.get_range(), [=](id&lt;1&gt; i) {
out[i] = inA[i] + inB[i];
});
}).wait();
Expand Down Expand Up @@ -404,7 +402,7 @@
auto inB = accessor{bufB, cgh, read_only};
auto out = accessor{bufO, cgh, write_only};

cgh.single_task&lt;add&gt;(bufO.get_range(), [=](id&lt;1&gt; i) {
cgh.parallel_for&lt;add&gt;(bufO.get_range(), [=](id&lt;1&gt; i) {
out[i] = inA[i] + inB[i];
});
});
Expand Down