Add SYCL kernels #441
Open
Add SYCL kernels #441
+2,434
−168
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
"non-uniform workgroup size" issue.
Base SYCL vairant is good for each; checksum issues with RAJA_SYCL variants
Also, change SYCL exec policies to sync for correctness. This needed to be investigated to resolve.
rhornung67
changed the title
|
@artv3 @MrBurmark @homerdin this should be good to go. I'll be making more modifications in subsequent PRs. |
MrBurmark
reviewed
Jun 5, 2024
| [=] (sycl::nd_item<1> item ) { | ||
|
|
||
| Index_type i = item.get_global_id(0); | ||
| if (i > 0 && i < iend) { |
There was a problem hiding this comment.
do all of the variants have both checks?
MrBurmark
reviewed
Jun 5, 2024
| } | ||
| POLYBENCH_2MM_BODY3; | ||
|
|
||
| sycl::range<3> global_dim1(1, |
There was a problem hiding this comment.
Don't we only need a 2d range?
MrBurmark
reviewed
Jun 5, 2024
| }); | ||
|
|
||
| qu->submit([&] (sycl::handler& h) { | ||
| h.parallel_for(sycl::nd_range<3>( global_dim2, wkgroup_dim), |
MrBurmark
reviewed
Jun 5, 2024
| using EXEC_POL = | ||
| RAJA::KernelPolicy< | ||
| RAJA::statement::SyclKernelAsync< | ||
| RAJA::statement::For<0, RAJA::sycl_global_0<work_group_size>, |
There was a problem hiding this comment.
In kernel do we always use a 3d grid, if so should we use sycl_global_2 here?
There was a problem hiding this comment.
Maybe we could set the sycl indexing dimension as an argument to SyclKernelAsync.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Add SYCL kernel variants
This PR adds SYCL variants of many kernels. There are a few issues still to be worked out in subsequent PRs.
Resolves #433