Skip to content

Disable no_unique_address for NVCC and CUDA 12.9 - #435

Merged
crtrott merged 1 commit into
kokkos:stablefrom
tpadioleau:disable-no_unique_address-cuda-12-9
Oct 30, 2025
Merged

Disable no_unique_address for NVCC and CUDA 12.9#435
crtrott merged 1 commit into
kokkos:stablefrom
tpadioleau:disable-no_unique_address-cuda-12-9

Conversation

@tpadioleau

@tpadioleau tpadioleau commented Oct 27, 2025

Copy link
Copy Markdown
Member

Following #434, I need to disable [[no_unique_address]] to make the Kokkos tests pass on a CUDA 12.9 environment on a CEA computer.

With [[no_unique_address]] I have the following errors:

[ RUN      ] cuda_DeathTest.mdspan_space_aware_accessor_invalid_access_from_device
/home/TP026749/kokkos/core/unit_test/TestSpaceAwareAccessorAccessViolation.hpp:31: Failure
Death test: { Kokkos::parallel_for(Kokkos::RangePolicy<ExecutionSpace>(s, 0, 1), *this); Kokkos::fence(); }
    Result: failed to die.
 Error msg:
[  DEATH   ] 
[  FAILED  ] cuda_DeathTest.mdspan_space_aware_accessor_invalid_access_from_device (147 ms)
[ RUN      ] cuda.mdspan_minimal_functional
/home/TP026749/kokkos/core/unit_test/TestMDSpan.hpp:46: Failure
Expected equality of these values:
  errors
    Which is: 2
  0
[  FAILED  ] cuda.mdspan_minimal_functional (0 ms)
[ RUN      ] cuda.mdspan_space_aware_accessor
cudaStreamSynchronize(stream) error( cudaErrorIllegalAddress): an illegal memory access was encountered /home/TP026749/kokkos/core/src/Cuda/Kokkos_Cuda_Instance.cpp:157
Backtrace: ...

@tpadioleau

Copy link
Copy Markdown
Member Author

@dalg24 may I get your review please ?

@dalg24

dalg24 commented Oct 30, 2025

Copy link
Copy Markdown
Member

Deferring to Christian who owns that code and who should be able to reproduce without fighting with machine access.

@dalg24
dalg24 requested a review from crtrott October 30, 2025 11:45
@crtrott

crtrott commented Oct 30, 2025

Copy link
Copy Markdown
Member

Hm I don't even get there. For me the compiler crashes.

[ 22%] Building CUDA object tests/CMakeFiles/test_mdspan_at.dir/test_mdspan_at.cpp.o
nvcc warning : Support for offline compilation for architectures prior to '<compute/sm/lto>_75' will be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning).
nvcc error   : '"$CICC_PATH/cicc"' died due to signal 11 (Invalid memory reference)
nvcc error   : '"$CICC_PATH/cicc"' core dumped
make[2]: *** [tests/CMakeFiles/test_mdspan_at.dir/build.make:80: tests/CMakeFiles/test_mdspan_at.dir/test_mdspan_at.cpp.o] Error 139
make[1]: *** [CMakeFiles/Makefile2:449: tests/CMakeFiles/test_mdspan_at.dir/all] Error 2

@crtrott

crtrott commented Oct 30, 2025

Copy link
Copy Markdown
Member

This change fixes the compiler crash I see with CUDA 12.9 and GCC 11.5 as host compiler. Testing then passes. I don't see issues with CUDA 12.8.

@crtrott
crtrott merged commit 404130d into kokkos:stable Oct 30, 2025
15 checks passed
@tpadioleau
tpadioleau deleted the disable-no_unique_address-cuda-12-9 branch October 31, 2025 13:29
@tpadioleau

Copy link
Copy Markdown
Member Author

thanks!

Comment thread include/experimental/__p0009_bits/config.hpp
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants