[FEA]: Make cooperative groups a part of CCCL
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
Description
### Is this a duplicate?
- [X] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Area
General CCCL
### Is your feature request related to a problem? Please describe.
I would like the `cooperative_groups.h` header and the `cooperative_groups/` subfolder of the CUDA toolkit's `include/` folder to be maintained as part of CCCL.
### Describe the solution you'd like
Provide the cooperative groups headers either as part of one of the three libraries comprising CCCL, as a fourth library, or as part of a "miscellaneous" component of CCCL.
### Describe alternatives you've considered
_No response_
### Additional context
As documented in [this exchange](https://forums.developer.nvidia.com/t/nvc-and-cooperative-groups-with-fun-little-patch/264505/5) between Mats, I've been playing around with getting cooperative groups to work with nvc++. In the process, I realized that the CG library is essentially header-only, apart from a few clearly separated functions in `cuda_device_runtime.h` starting with `cudaCG` related to multi-grid groups. I was able to come up with a patch (attached to the aforementioned [forum thread](https://forums.developer.nvidia.com/t/nvc-and-cooperative-groups-with-fun-little-patch/264505/5)) that improved nvc++/CG interoperability that I have been using for personal projects since then.
This experience gives me a strong impression that cooperative groups would benefit from being maintained as proper open-source software. If that had been the case, I would have certainly submitted a pull request instead of engaging in the prehistoric practice of distributing patches.
I'm bothering you here because CCCL feels like a clean slate and a great opportunity to squeeze CG into an open-source repository. This argument could extend to other parts of the CUDA toolkit `include` folder, but having searched through it, I think cooperative groups really stand out in this regard as they're self-contained, widely applicable, and basically completely header-only (i.e., with all implementations fully visible).
One of the difficulties I see is deciding where exactly to put CG headers in the CCCL repository. While that would certainly be the maintainers' call, here's a few options to illustrate the idea. Personally, I would argue for the inclusion of CG in libcu++, which [references it](https://nvidia.github.io/libcudacxx/extended_api/thread_groups.html) in its documentation and provides other non-stdlib extensions, such as `cuda::pipeline` (I think, despite a [strange statement](https://nvidia.github.io/libcudacxx/extended_api.html#pipelines) in the docs that the "pipeline library is not part of the open-source libcu++ distribution"). CG could also be considered a fourth CCCL library besides CUB, Thrust, and libcu++, but that might be less convincing as it's significantly smaller in scope than those libraries. Finally, CG could be tucked away into a "misc" folder (or "core", or "toolkit_support", or something like that), which would be useful but avoid the fanfare of a fourth library.
Thank you very much for your consideration.
Best,
Matjaž
Contributor guide
Assessment
This issue has not been assessed yet.