-
Notifications
You must be signed in to change notification settings - Fork 102
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
Problem with exercise 16 #326
base: main
Are you sure you want to change the base?
Conversation
….md to use items and added the corresponding code to solution.cpp
@@ -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]); |
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: ***@***.***>
|
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