-
Notifications
You must be signed in to change notification settings - Fork 299
Move launch API from cudax to libcu++ #6667
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
Conversation
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
| const cudax::device_buffer<int> b = cudax::make_buffer(stream, resource, a.size(), 1); | ||
|
|
||
| cudax::launch(stream, cudax::make_config(cuda::grid_dims<1>, cuda::block_dims<32>), add_kernel{}, a, b); | ||
| cudax::launch(stream, cuda::make_config(cuda::grid_dims<1>, cuda::block_dims<32>), add_kernel{}, a, b); |
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.
Note: cuda::make_config is too generic in my opinion. We should find a better name in a separate PR
| __device__ void operator()(Config config) | ||
| { | ||
| decltype(auto) dynamic_smem = cudax::device::dynamic_shared_memory_view(config); | ||
| decltype(auto) dynamic_smem = cuda::dynamic_shared_memory_view(config); |
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.
| decltype(auto) dynamic_smem = cuda::dynamic_shared_memory_view(config); | |
| decltype(auto) dynamic_smem = cuda::device::dynamic_shared_memory_view(config); |
This comment has been minimized.
This comment has been minimized.
cudax/test/CMakeLists.txt
Outdated
| launch/launch_smoke.cu | ||
| ) | ||
|
|
||
| cudax_add_catch2_test(test_target launch_configuration ${cudax_target} |
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.
Heads up that this will need to be rebased or have main merged in -- we no longer pass cudax_target here.
| //! Launches a kernel functor object on the specified stream and with specified | ||
| //! configuration. Kernel functor object is a type with __device__ operator(). |
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.
Please, fix your clang-format :)
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
🥳 CI Workflow Results🟩 Finished in 3h 04m: Pass: 100%/128 | Total: 21h 47m | Max: 1h 14m | Hits: 98%/219845See results here. |
* Move launch API from cudax to libcu++ * Review feedback and test fixes * More fixes * GCC7 fix * Update libcudacxx/include/cuda/__launch/configuration.h * Fix old GCC --------- Co-authored-by: David Bayer <[email protected]> (cherry picked from commit 4ab39a7)
|
Successfully created backport PR for |
This PR moves the parts of launch API ready for release to libcu++. Graphs related and
cuda::executionparts will remain experimental.I also commented out extended lambda testing to figure out cmake changes needed to add the compilation flag to that specific test