NVIDIA / NVIDIA/cccl

[FEA]: Make cooperative groups a part of CCCL

Open
#442 2 comments 2 reactions 1 assignee Claimed by @jrhemstad View on GitHub
2024-2025 goal
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.