NVIDIA / NVIDIA/cccl

Port `triple_chevron`

Open
#568 1 comment 0 reactions 1 assignee Assigned to @elstehle View on GitHub
good first issue
Dominant language
C++
Stars
2.5k
Forks
487
Avg merge
2d 7h
Merged PRs (30d)
296

Description

CUB uses `THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron` for launching all of its kernels. Since all of Thrust kernels will migrate to cub (https://github.com/NVIDIA/cccl/issues/26), there's no reason to keep the kernel launcher in Thrust. Unlike remaining tasks in https://github.com/NVIDIA/cccl/issues/26, this one is not blocked on VSMem work (https://github.com/NVIDIA/cccl/issues/548).

Regarding breaking API, we consider `thrust::cuda_cub` to be an implementation detail, so strictly speaking there should be no need in maintaining the API in the thrust namespace. Nevertheless, I can see people using `thrust::cuda_cub` directly, probably because there's no `::detail::` in the name. To be on the safe side, we can add `using triple_chevron = cub::detail::triple_chevron`.

On the CUB side, we'll need to make sure that `triple_chevron` header is placed in `cccl/cub/cub/detail` and the struct itself is in `cub::detail::` namespace. Besides that, I'd like to see the test coverage for it. The test should use `cdp_launch` (example is in /cccl/cub/test/catch2_test_cdp_wrapper.cu) and check that `triple_chevron` is able to launch kernels from both host and device code. Launching a wrong configuration to validate the return code (example in /cccl/cub/test/catch2_test_debug.cu) would also be helpful.

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.