Skip to content

Conversation

@keyradical
Copy link
Contributor

This PR adds an algorithm for doing a GPU wide barrier in CUDA backend.

Rough outline of the algorithm:

  • Every 0th thread from each workgroup performs atomic.add(1)
  • The same thread checks the atomic result with ld.acquire in a loop until it's equal to total amount of workgroups.
  • All threads call group-wide barrier.sync

One caveat to this is that there is no initialization of the atomic start value. So if we call this barrier several times in a kernel, on the second iteration, the start value will already contain the result from previous barrier. That's why we actually spin the while loop while current value % totalWgroups != 0.

@keyradical keyradical requested a review from a team as a code owner July 29, 2024 15:19
@keyradical keyradical requested a review from MartinWehking July 29, 2024 15:19
@keyradical keyradical requested a review from a team as a code owner July 30, 2024 08:27
@keyradical keyradical requested a review from maarquitos14 July 30, 2024 08:27
sycl::group_barrier(root);

if (it.get_group(0) % 2 == 0) {
X += sycl::sin(X);
Copy link
Contributor

Choose a reason for hiding this comment

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

I don't see that we ever check neither X nor Y. Should we? Otherwise, why we need this? How do we make sure it ran?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I wanted to explicitly delay some of the workgroups by adding them more work to do, because I've seen this test passing if insufficient barrier was used. For instance on CUDA backend, doing work-group wide barrier would be enough for it to pass and that is not correct. I think this test should perform some work-group divergence to actually check that we actually perform gpu-wide barrier.

How do we make sure it ran?

The X and Y are declared as volatile and my understanding was that this would prevent compiler from removing them with some optimization.

Copy link
Contributor

Choose a reason for hiding this comment

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

Right, can we have a comment explaining it? Otherwise we risk that this code will just be removed in the future thinking it's not required.

Copy link
Contributor

@maarquitos14 maarquitos14 left a comment

Choose a reason for hiding this comment

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

LGTM.

Copy link
Contributor

@MartinWehking MartinWehking left a comment

Choose a reason for hiding this comment

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

LGTM

@keyradical
Copy link
Contributor Author

@intel/llvm-gatekeepers can you merge this, please?

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants