Skip to content
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

Port triple_chevron #568

Open
gevtushenko opened this issue Oct 16, 2023 · 1 comment · May be fixed by #592
Open

Port triple_chevron #568

gevtushenko opened this issue Oct 16, 2023 · 1 comment · May be fixed by #592
Assignees
Labels
good first issue Good for newcomers.

Comments

@gevtushenko
Copy link
Collaborator

gevtushenko commented Oct 16, 2023

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 (#26), there's no reason to keep the kernel launcher in Thrust. Unlike remaining tasks in #26, this one is not blocked on VSMem work (#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.

@ZelboK
Copy link
Contributor

ZelboK commented Oct 16, 2023

Thank you for writing this out!

@gevtushenko gevtushenko linked a pull request Oct 23, 2023 that will close this issue
@gevtushenko gevtushenko linked a pull request Oct 23, 2023 that will close this issue
@jrhemstad jrhemstad assigned elstehle and unassigned ZelboK Dec 6, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
good first issue Good for newcomers.
Projects
Status: Todo
Development

Successfully merging a pull request may close this issue.

3 participants