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.

Sign up for GitHub

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

Jump to bottom

backport unreachable #2852

Open
wants to merge 27 commits into
base: main
Choose a base branch
from
Open

backport unreachable #2852

wants to merge 27 commits into from
+14,257 −13,707

Conversation

davebayer
Copy link
Contributor

@davebayer davebayer commented Nov 18, 2024 •
edited
Loading

This PR backports cuda::std::unreachable to C++11, replacing old implementation (_CCCL_UNREACHABLE).

@davebayer davebayer requested review from a team as code owners November 18, 2024 12:08
@copy-pr-bot copy-pr-bot
Copy link

copy-pr-bot bot commented Nov 18, 2024

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@miscco
Copy link
Collaborator

miscco commented Nov 18, 2024 •
edited
Loading

It also fixes the implementation, because for device code the implementation was never right as __CUDA_ARCH__ is never defined outside device code.

Can you elaborate on this a bit. Outside of __CUDA_ARCH__ the macro was defined to just be __builtin_unreachable for non MSVC

@davebayer
Copy link
Contributor Author

davebayer commented Nov 18, 2024 •
edited
Loading

It also fixes the implementation, because for device code the implementation was never right as __CUDA_ARCH__ is never defined outside device code.

Can you elaborate on this a bit. Outside of __CUDA_ARCH__ the macro was defined to just be __builtin_unreachable for non MSVC

I didn't know NVHPC does not define __CUDA_ARCH__ in device code.. But I think at least for _CCCL_CUDACC_BELOW(11, 2) and _CCCL_CUDACC_BELOW(11, 3) the old implementation of unreachable should not work, because __CUDA_ARCH__ is never defined outside __device__ or __global__ scope for nvcc.

@miscco
Copy link
Collaborator

miscco commented Nov 18, 2024

/ok to test

@davebayer davebayer changed the title backport and fix unreachable backport unreachable Nov 19, 2024
@miscco
Copy link
Collaborator

miscco commented Nov 19, 2024

pre-commit.ci autofix

@miscco
Copy link
Collaborator

miscco commented Nov 20, 2024

/ok to test

@davebayer
Copy link
Contributor Author

pre-commit.ci autofix

miscco and others added 13 commits November 22, 2024 16:23
@miscco @davebayer
Move implementation of _LIBCUDACXX_TEMPLATE to CCCL (NVIDIA#2832)
897086a 
* Move implementation of `_LIBCUDACXX_TEMPLATE` to CCCL

We have emulation for concepts in LIBCUDACXX that was guarded behind C++14

But there is nothing that requires C++14 for just the template headers and we want to use them universally throughout the codebase

Consequently move them to CCCL proper and enable them unconditionally. To ensure that we do not add any hidden dependencies this also adds a barebones implementation of `enable_if_t` and a trailing `enable_if_t`
@miscco @davebayer
Try to work around issue with NVHPC in conjunction of older CTK versi…
786d442 
…ons (NVIDIA#2889)

NVHPC can consume older CTK headers for stdpar, so we need to try and avoid using those
@bernhardmgruber @davebayer
Refactoring (NVIDIA#2905)
6b5fa22
@ericniebler @davebayer
add "interface" to _CCCL_PUSH_MACROS (NVIDIA#2919)
44f0331
@ericniebler @davebayer
Replace inconsistent Doxygen macros with _CCCL_DOXYGEN_INVOKED (NVI…
c9a6e6a 
…DIA#2921)

fixes NVIDIA#2362
@davebayer @bernhardmgruber
implement C++26 std::span::at (NVIDIA#2924)
a50019d 
Co-authored-by: Bernhard Manfred Gruber <[email protected]>
@davebayer @miscco
@bernhardmgruber @ericniebler
move msvc compiler macros to new version (NVIDIA#2885)
d0f5bd2 
Co-authored-by: Michael Schellenberger Costa <[email protected]>
Co-authored-by: Bernhard Manfred Gruber <[email protected]>
Co-authored-by: Eric Niebler <[email protected]>
@bernhardmgruber @davebayer
Reorganize PTX tests to match generator (NVIDIA#2930)
5b57a4c
@bernhardmgruber @ahendriksen @davebayer
Reorganize PTX docs to match generator (NVIDIA#2929)
96e8199 
Co-authored-by: Allard Hendriksen <[email protected]>
@miscco @davebayer
Improve build instructions for libcu++ (NVIDIA#2881)
92a22f2 
* Improve build instructions for libcu++
* Add section about the options for the build script
* Delegate more to the contributor guidelines
@bernhardmgruber @davebayer
Reorganize PTX headers to match generator (NVIDIA#2925)
14484a6
@davebayer @miscco
implement C++26 std::span's constructor from std::initializer_list (
e897f1f 
NVIDIA#2923)

Co-authored-by: Michael Schellenberger Costa <[email protected]>
@davebayer
Add tuple protocol to cuda::std::complex from C++26 (NVIDIA#2882)
3b23083
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@bernhardmgruber bernhardmgruber bernhardmgruber left review comments

@wmaxey wmaxey Awaiting requested review from wmaxey wmaxey is a code owner automatically assigned from NVIDIA/cccl-libcudacxx-codeowners

@gonidelis gonidelis Awaiting requested review from gonidelis gonidelis is a code owner automatically assigned from NVIDIA/cccl-codeowners

@miscco miscco Awaiting requested review from miscco

Requested changes must be addressed to merge this pull request.

Assignees
No one assigned
Labels
None yet
Projects
CCCL
Status: In Progress
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
4 participants
@davebayer @miscco @bernhardmgruber @ericniebler