Add SYCL kernels #441
Add SYCL kernels #441
Conversation
"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. |
| } | ||
| POLYBENCH_2MM_BODY3; | ||
|
|
||
| sycl::range<3> global_dim1(1, |
There was a problem hiding this comment.
Don't we only need a 2d range?
There was a problem hiding this comment.
Technically, yes, only a 2d range is needed. I discussed the pattern and consistency across kernels with @homerdin and we agreed that this is how we should do it.
There was a problem hiding this comment.
There are places where we use kernel for a 1d loop and then the base variant uses a 1d range instead of a 3d range, it makes sense to me to use what you would naturally use instead of doing what RAJA::kernel happens to use. If it ends up making a difference we could see it.
|
@MrBurmark I addressed all your comments. Anything else in here you would like me to address? |
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