-
Notifications
You must be signed in to change notification settings - Fork 117
Problem with exercise 16 #326
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.
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
base: main
Are you sure you want to change the base?
Conversation
….md to use items and added the corresponding code to solution.cpp
| ``` | ||
| 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]); |
There was a problem hiding this comment.
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 rowMajorLinearId = sycl::id(globalId[0], globalId[1]); | ||
| auto columnMajorLinearId = sycl::id(globalId[1], globalId[0]); |
There was a problem hiding this comment.
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
|
Hi,
To achieve a better match with as few changes to the code as possible, I think this can be a good solution:
```cpp
cgh.parallel_for<image_convolution>(
ndRange, [=](sycl::id<2> idx) {
auto rowMajorLinearId = idx[1] * inputImgWidth + idx[0];
auto columnMajorLinearId = idx[0] * inputImgWidth + idx[1];
auto globalId(rowMajorLinearId % inputImgWidth, rowMajorLinearId / inputImgWidth);
what do you think ?
________________________________
From: Duncan McBain ***@***.***>
Sent: 31 January 2024 18:46
To: codeplaysoftware/syclacademy ***@***.***>
Cc: Pablo ***@***.***>; Author ***@***.***>
Subject: Re: [codeplaysoftware/syclacademy] Problem with exercise 16 (PR #326)
@DuncanMcBain commented on this pull request.
________________________________
In Code_Exercises/Exercise_16_Coalesced_Global_Memory/README.md<#326 (comment)>:
@@ -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]);
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?
________________________________
In Code_Exercises/Exercise_16_Coalesced_Global_Memory/solution.cpp<#326 (comment)>:
+ auto rowMajorLinearId = sycl::id(globalId[0], globalId[1]);
+ auto columnMajorLinearId = sycl::id(globalId[1], globalId[0]);
Similarly these are 2d ids, not linear ones
—
Reply to this email directly, view it on GitHub<#326 (review)>, or unsubscribe<https://github.com/notifications/unsubscribe-auth/BE6TABJLX5HC53VYXKKNMSTYRKGQDAVCNFSM6AAAAABBTCWAX2VHI2DSMVQWIX3LMV43YUDVNRWFEZLROVSXG5CSMV3GSZLXHMYTQNJUGU3DMMZUGQ>.
You are receiving this because you authored the thread.Message ID: ***@***.***>
|
|
Update on the status will be appreciated : ) |
This PR is inspired by the Lesson 4 `index.html` changes in KhronosGroup#326 and fixes up the `handler::single_task` invocations in a couple of places to make sure they are correct. * Change `parallel_for` to `single_task` in early example where we haven't defined `bufO` yet to call `bufO.get_range()` on. * Remove `(id<1> i)` parameter for `single_task` lambda invocation, as this shouldn't have any parameters.
This PR tries to resolve the issue highlighted in KhronosGroup#326 where the "Coalesced Global Memory" exercise README was asking a user to manually calculate a linear id that would be used to index memory. However, the exercise code itself uses 2d buffers/accessors which are indexed with a `sycl::id<2>`. So the current exercise request doesn't make sense, or could only be achieved with modifications beyond what is reasonably expected for a beginner. Instead, I've updated the README to state the usage of the 2d buffer/accessors and ask the user to experiment with flipping the dimension indexes, as this is the only difference between this exercise solution, and the solution from the previous lesson ```sh $ diff Code_Exercises/Coalesced_Global_Memory/solution.cpp Code_Exercises/Image_Convolution/reference.cpp 59a60 > 77a79 > globalId = sycl::id{globalId[1], globalId[0]}; ```
This PR tries to resolve the issue highlighted in KhronosGroup#326 where the "Coalesced Global Memory" exercise README was asking a user to manually calculate a linear id that would be used to index memory. However, the exercise code itself uses 2d buffers/accessors which are indexed with a `sycl::id<2>`. So the current exercise request doesn't make sense, or could only be achieved with modifications beyond what is reasonably expected for a beginner. Instead, I've updated the README to state the usage of the 2d buffer/accessors and ask the user to experiment with flipping the dimension indexes, as this is the only difference between this exercise solution, and the solution from the previous lesson ```sh $ diff Code_Exercises/Coalesced_Global_Memory/solution.cpp Code_Exercises/Image_Convolution/reference.cpp 59a60 > 77a79 > globalId = sycl::id{globalId[1], globalId[0]}; ```
I've created new PRs to try address the issues highlighted but based on up-to-date tip: |
The solution.cpp uses items in the parallel_for but the README.md uses ids. I changed the README.md and added the code to the solution.cpp