[BUG]: Compile error of libcudacxx examples

[j3soon]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Compile-time Error

Component

libcu++

Describe the bug

Examples under libcudacxx/examples/ should compile successfully, but they seem outdated and result in compile error.

How to Reproduce

In a CCCL devcontainer, run the following:

cd ~/cccl/libcudacxx/examples/
cmake .
cmake --build .

and observe the compilation error.

Expected behavior

These examples should compile successfully.

Reproduction link

No response

Operating System

Devcontainer in Ubuntu 22.04 LTS

nvidia-smi output

+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 550.120 Driver Version: 550.120 CUDA Version: 12.4 |
|-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 Quadro RTX 6000 Off | 00000000:01:00.0 Off | Off |
| 35% 33C P8 11W / 260W | 6MiB / 24576MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+

+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
+-----------------------------------------------------------------------------------------+

NVCC version

nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on Thu_Sep_12_02:18:05_PDT_2024
Cuda compilation tools, release 12.6, V12.6.77
Build cuda_12.6.r12.6/compiler.34841621_0

[j3soon]

PR #3013 is my attempt to update these examples to resolve compilation issues. However, I'm currently stuck on the following error and am unsure how to proceed. Any suggestions would be greatly appreciated.

Error snippet:

[ 90%] Building CUDA object CMakeFiles/hash_map.dir/concurrent_hash_table.cu.o
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/types/base.h(147): error: no instance of overloaded function "cuda::std::__4::__atomic_compare_exchange_cuda" matches the argument list
            argument types are: (concurrent_hash_table<int, cuda::std::__4::atomic<int>, thrust::THRUST_200500_700_NS::identity<int>, thrust::THRUST_200500_700_NS::equal_to<int>, thrust::THRUST_200500_700_NS::universal_memory_resource>::state_type *, concurrent_hash_table<int, cuda::std::__4::atomic<int>, thrust::THRUST_200500_700_NS::identity<int>, thrust::THRUST_200500_700_NS::equal_to<int>, thrust::THRUST_200500_700_NS::universal_memory_resource>::state_type *, concurrent_hash_table<int, cuda::std::__4::atomic<int>, thrust::THRUST_200500_700_NS::identity<int>, thrust::THRUST_200500_700_NS::equal_to<int>, thrust::THRUST_200500_700_NS::universal_memory_resource>::state_type, __nv_bool, cuda::std::__4::__memory_order_underlying_t, cuda::std::__4::__memory_order_underlying_t, cuda::std::__4::__thread_scope_system_tag)
    { __result = __atomic_compare_exchange_cuda( __a->get(), __expected, __val, true, static_cast<__memory_order_underlying_t>(__success), static_cast<__memory_order_underlying_t>(__failure), _Sco{}); }
                 ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_derived.h(98): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile _Tp *, _Tp *, _Tp, __nv_bool, int, int, _Sco)" failed deduction
              bool __atomic_compare_exchange_cuda(
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_derived.h(59): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Tp *, _Tp, __nv_bool, int, int, _Sco)" failed deduction
              bool __atomic_compare_exchange_cuda(
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_derived.h(37): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Tp *, _Tp, __nv_bool, int, int, _Sco)" failed deduction
              bool __atomic_compare_exchange_cuda(
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(6062): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_system_tag)" failed deduction
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_system_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(6033): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_system_tag)" failed deduction
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_system_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(5999): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_system_tag)" failed deduction
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_system_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(5970): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_system_tag)" failed deduction
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_system_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(5107): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_system_tag)" failed deduction
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_system_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(5078): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_system_tag)" failed deduction
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_system_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(5044): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_system_tag)" failed deduction
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_system_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(5015): note #3327-D: candidate function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_system_tag)" failed deduction
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_system_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(3789): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_device_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_device_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(3760): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_device_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_device_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(3726): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_device_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_device_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(3697): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_device_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_device_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(2834): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_device_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_device_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(2805): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_device_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_device_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(2771): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_device_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_device_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(2742): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_device_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_device_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(1516): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_block_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_block_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(1487): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_block_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_block_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(1453): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_block_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_block_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(1424): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_block_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_block_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(561): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_block_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_block_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(532): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_block_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_block_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(498): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_block_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_block_tag) {
                   ^
/usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/functions/cuda_ptx_generated.h(469): note #3326-D: function template "cuda::std::__4::__atomic_compare_exchange_cuda(volatile void *, _Type *, _Type, __nv_bool, int, int, cuda::std::__4::__thread_scope_block_tag)" does not match because argument #7 does not match parameter
              bool __atomic_compare_exchange_cuda(volatile void *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int __failure_memorder, __thread_scope_block_tag) {
                   ^
          detected during:
            instantiation of "__nv_bool cuda::std::__4::__atomic_compare_exchange_weak_dispatch(_Sto *, _Up *, _Up, cuda::std::__4::memory_order, cuda::std::__4::memory_order, _Sco) [with _Sto=cuda::std::__4::__atomic_storage<concurrent_hash_table<int, cuda::std::__4::atomic<int>, thrust::THRUST_200500_700_NS::identity<int>, thrust::THRUST_200500_700_NS::equal_to<int>, thrust::THRUST_200500_700_NS::universal_memory_resource>::state_type>, _Up=concurrent_hash_table<int, cuda::std::__4::atomic<int>, thrust::THRUST_200500_700_NS::identity<int>, thrust::THRUST_200500_700_NS::equal_to<int>, thrust::THRUST_200500_700_NS::universal_memory_resource>::state_type, _Sco=cuda::std::__4::__thread_scope_system_tag, <unnamed>=0]" at line 49 of /usr/local/cuda/bin/../targets/x86_64-linux/include/cuda/std/__atomic/api/owned.h
            instantiation of "__nv_bool cuda::std::__4::__atomic_common<_Tp, _Sco>::compare_exchange_weak(_Tp &, _Tp, cuda::std::__4::memory_order) noexcept [with _Tp=concurrent_hash_table<int, cuda::std::__4::atomic<int>, thrust::THRUST_200500_700_NS::identity<int>, thrust::THRUST_200500_700_NS::equal_to<int>, thrust::THRUST_200500_700_NS::universal_memory_resource>::state_type, _Sco=cuda::std::__4::__thread_scope_system_tag]" at line 127 of /home/coder/cccl/libcudacxx/examples/concurrent_hash_table.cu
            instantiation of "thrust::THRUST_200500_700_NS::pair<concurrent_hash_table<Key, Value, Hash, KeyEqual, MemoryResource>::value_iterator, __nv_bool> concurrent_hash_table<Key, Value, Hash, KeyEqual, MemoryResource>::try_emplace(UKey &&, Args &&...) [with Key=int, Value=cuda::std::__4::atomic<int>, Hash=thrust::THRUST_200500_700_NS::identity<int>, KeyEqual=thrust::THRUST_200500_700_NS::equal_to<int>, MemoryResource=thrust::THRUST_200500_700_NS::universal_memory_resource, UKey=const int &, Args=<>]" at line 156 of /home/coder/cccl/libcudacxx/examples/concurrent_hash_table.cu
            instantiation of "concurrent_hash_table<Key, Value, Hash, KeyEqual, MemoryResource>::mapped_type &concurrent_hash_table<Key, Value, Hash, KeyEqual, MemoryResource>::operator[](const concurrent_hash_table<Key, Value, Hash, KeyEqual, MemoryResource>::key_type &) [with Key=int, Value=cuda::std::__4::atomic<int>, Hash=thrust::THRUST_200500_700_NS::identity<int>, KeyEqual=thrust::THRUST_200500_700_NS::equal_to<int>, MemoryResource=thrust::THRUST_200500_700_NS::universal_memory_resource]" at line 201 of /home/coder/cccl/libcudacxx/examples/concurrent_hash_table.cu

Please see the attachment for the full compiler output.

error.txt