Support resource constructs in RAJA Teams #1029
Conversation
include/RAJA/policy/cuda/teams.hpp
Outdated
| resources::Cuda &cuda_res = res.cuda; | ||
| //resources::Cuda::get_default(); | ||
| /* Use the zero stream until resource is better supported */ | ||
| cudaStream_t stream = cuda_res.get_stream(); |
There was a problem hiding this comment.
@mdavis36 is this all we need? and does this look reasonable? Idea here is to pass in a TeamResource struct that holds all supported back-ends an we choose the right one depending on what is being called.
There was a problem hiding this comment.
If you want to support events that would need a bit more. What's making the decision as to which backend to use? Is it controlled by the caller somehow? I ask because if it is, it might make sense for it to be a variant-like type where the TeamResource only contains one resource, but could contain any of a compile-time list of them, if it's out of the caller's control then maybe tuple of them makes sense so the types are known. If the types can't be known, then some more error checking is needed to make sure the resource is the right kind.
There was a problem hiding this comment.
Yes the caller chooses the backend, its base on run-time parameter in the launch API. Basically an enum for HOST or DEVICE.
There was a problem hiding this comment.
So, perhaps, pass a resource of type rather than enum to identify type could be a nice model? If it's using the camp platform enum the number is already there in the erased resource type.
There was a problem hiding this comment.
That sounds pretty good to me, we had some other ideas as well. I'll explore them and push to this branch.
|
@artv3 should this PR get merged for the release also since you mention the RAJA::expt::Resources stuff in the docs for naming kernels. |
Actually, the Resources object used Teams is different than the CAMP resource. This PR will rename RAJA::expt::Resources to RAJA::expt::Grid and introduce support for CAMP Resource. But yes, I think we have pretty good path forward (based on conversations) to add the CAMP resource support. |
Sounds good. Thanks. |
|
@trws @MrBurmark @mdavis36 , by combining idea's I've put together support for resources in RAJA Teams. The example |
artv3
changed the title
|
Thanks all!, @mdavis36 sounds good on the example. I'll work on that to make sure behavior is as expected. |
|
Something is off, the example in resource-teams.cpp isn't running the kernels asynchronously... digging into it... |
I figured it out, just need to update with develop and it will be ready to merge. |
|
|
||
| template <typename BODY_IN> | ||
| static resources::EventProxy<resources::Resource> | ||
| exec(RAJA::resources::Resource res, LaunchContext const &ctx, BODY_IN &&body_in) |
There was a problem hiding this comment.
Should this take a Cuda resource instead of a type erased resource? I'm thinking about this with the aim of having a way to use non type erased resource with teams at some point.
There was a problem hiding this comment.
I would be in favor of doing it with the next iteration of the API since we will be making some larger design changes.
| @@ -46,6 +42,7 @@ __global__ void launch_global_fcn(LaunchContext ctx, BODY body_in) | |||
|
|
|||
| template <bool async> | |||
| struct LaunchExecute<RAJA::expt::cuda_launch_t<async, 0>> { | |||
|
|
|||
| template <typename BODY_IN> | |||
| static void exec(LaunchContext const &ctx, BODY_IN &&body_in) | |||
There was a problem hiding this comment.
Can this be removed or implemented with the resource version?
There was a problem hiding this comment.
I think we can, let me see what I can do.
There was a problem hiding this comment.
I think it can, but I'll make the change after the release along with API rework.
|
@artv3 I merged the SYCL PR and pulled develop into your PR branch. It can be merged next when tests pass. |
|
@rhornung67 Travis is hanging -- are we still good to merge? |
|
Yes. Go ahead and merge. |
Summary
TODO:
API changes:
RAJA::expt::Resources -> RAJA::expt::Grid.