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 43 commits into
base: main
Choose a base branch
from
Open

backport unreachable #2852

wants to merge 43 commits into from

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 and others added 29 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
@miscco @davebayer
Move implementation of _LIBCUDACXX_TEMPLATE to CCCL (NVIDIA#2832)
0a0773f 
* 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`
@ericniebler @davebayer
Replace inconsistent Doxygen macros with _CCCL_DOXYGEN_INVOKED (NVI…
3911b25 
…DIA#2921)

fixes NVIDIA#2362
@davebayer @miscco
@bernhardmgruber @ericniebler
move msvc compiler macros to new version (NVIDIA#2885)
b760c7f 
Co-authored-by: Michael Schellenberger Costa <[email protected]>
Co-authored-by: Bernhard Manfred Gruber <[email protected]>
Co-authored-by: Eric Niebler <[email protected]>
@davebayer @miscco
implement C++26 std::span's constructor from std::initializer_list (
8ee0b18 
NVIDIA#2923)

Co-authored-by: Michael Schellenberger Costa <[email protected]>
@bernhardmgruber @davebayer
Add missing qualifier for cuda namespace (NVIDIA#2940)
ef9ae72 
Fixes: NVIDIA#2939
@bernhardmgruber @davebayer
Try to fix a clang warning: (NVIDIA#2941)
5e4d6e0 
agent_histogram.cuh:827:37: warning: comparison of different enumeration types
@ericniebler @davebayer
minor consistency improvements in concepts macros (NVIDIA#2928)
6dbeb46
@miscco @davebayer
Drop some of the mdspan fold implementation (NVIDIA#2949)
2a47168 
* Drop unused macros

* Do not return a custom struct

* Replace `__MDSPAN_FOLD_AND` with `__fold_and_v` when possible
@caugonnet @davebayer
[STF] Implement CUDASTF_DOT_TIMING for the ctx.cuda_kernel construct (N…
a4477a4 
…VIDIA#2950)

* Implement CUDASTF_DOT_TIMING facility for ctx.cuda_kernel

* clang-format
@miscco @davebayer
Avoid potential null dereference in annotated_ptr (NVIDIA#2951)
09db35f 
Fixes [BUG]: UB in annotated_ptr NVIDIA#2942
@davebayer
make compiler version comparison utility generic (NVIDIA#2952)
8ca5380
@miscco @davebayer
Add SM100 descriptor to target (NVIDIA#2954)
3c6fcd2 
This is adding the missing sm_100 identifier to nv/target

Fixes NVIDIA#2890
@bernhardmgruber @davebayer
Regenerate cuda::ptx headers/docs and run format (NVIDIA#2937)
aededb9 
Overwrites all generated PTX header and documentation files and runs `pre-commit run --all-files`. Also exclude generated PTX headers from header check.
@bernhardmgruber @davebayer
Regenerate PTX test (NVIDIA#2953)
4d53204 
Overwrites all generated PTX tests and runs `pre-commit run --all-files`
@miscco @davebayer
Do not include extended floating point headers if they are not needed (…
bdff820 
…NVIDIA#2956)

Fixes NVIDIA#2933
@pciolkosz @davebayer
[CUDAX] Add copy_bytes and fill_bytes overloads for mdspan (NVIDIA#2932)
651bbcf 
* Implement copy_bytes for mdspan

* Add final conversion to mdspan and more tests

* mdspan fill_bytes

* Add docs

* Fix issues after rebase

* Help old GCC figure out the types

* Move runtime extents check to a function

* Fix clang and more old GCC fixes
@davebayer davebayer requested a review from a team as a code owner November 25, 2024 21:03
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.
6 participants
@davebayer @miscco @bernhardmgruber @ericniebler @caugonnet @pciolkosz