[Tile] Mark all atomics functions as __host__ __device__ only

[miscco]

Our atomics implementation is based on PTX which is not usable in a tile context

[coderabbitai]

📝 Walkthrough

Summary by CodeRabbit

• New Features

  • Atomic operations, flags, fences, and related wrappers are now available across host and device compilation contexts.
  • Expanded host/device coverage for atomic, synchronization, and pipeline functionality.

• Tests

  • Updated CUDA tile-mode compatibility checks and diagnostics.
  • Broadened test coverage across host and device execution paths.

Walkthrough

Changes

Atomic host-device enablement

Layer / File(s)	Summary
Atomic API annotations libcudacxx/include/cuda/std/atomic	Atomic, atomic flag, and fence wrappers now use _CCCL_HOST_DEVICE_API without changing dispatch logic.
Test infrastructure and CUDA tests libcudacxx/test/support/*, libcudacxx/test/libcudacxx/cuda/*	Adds TEST_HOST_DEVICE_FUNC, applies it to shared helpers and test entry points, and updates tile expectations to force-tile.
Standard and synchronization tests libcudacxx/test/libcudacxx/{libcxx,std}/*	Retags atomic, barrier, latch, and semaphore tests for host/device compilation and updates tile diagnostics.
Atomic operation coverage libcudacxx/test/libcudacxx/std/atomics/*	Updates operation, reference, type, flag, fence, and lock-free tests with host-device annotations and force-tile gating.

Possibly related PRs

• NVIDIA/cccl#10538 — Adds the same host-device test macro for another CUDA library test area.
• NVIDIA/cccl#10540 — Adds the force-tile test infrastructure used by these updated expectations.

Suggested reviewers: davebayer, bernhardmgruber, ericniebler

✨ Finishing Touches 💡 1

🛠️ Fix failing CI checks 💡

• Fix failing CI checks

---

_{Comment @coderabbitai help to get the list of available commands.}

[coderabbitai]

🧹 Nitpick comments (1)

libcudacxx/test/libcudacxx/std/atomics/atomics.lockfree/isalwayslockfree.pass.cpp (1)

33-33: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Rename the touched helpers and constant to snake_case.

checkAlwaysLockFree, checkLongLongTypes, getSizeOf, and ExpectLockFree violate the repository rule for non-CUB C++ symbols. Rename them and update their call sites in run().

As per coding guidelines, use snake_case for all other C++ symbols.

Also applies to: 55-55, 64-66, 73-79

Source: Coding guidelines

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: dd3221ad-d2ee-4c95-be7f-daa93ed59928

📥 Commits

Reviewing files that changed from the base of the PR and between 7c153c9 and 4511657.

📒 Files selected for processing (152)

• libcudacxx/include/cuda/std/atomic
• libcudacxx/test/libcudacxx/cuda/atomics/atomic.ext/atomic_fetch_half.fail.cpp
• libcudacxx/test/libcudacxx/cuda/atomics/atomic.ext/atomic_fetch_max.pass.cpp
• libcudacxx/test/libcudacxx/cuda/atomics/atomic.ext/atomic_fetch_min.pass.cpp
• libcudacxx/test/libcudacxx/cuda/atomics/atomic.ext/atomic_helpers.h
• libcudacxx/test/libcudacxx/cuda/atomics/atomic.local.pass.cpp
• libcudacxx/test/libcudacxx/cuda/atomics/atomic.uninitialized.pass.cpp
• libcudacxx/test/libcudacxx/cuda/atomics/atomic_16b_host_constructible.pass.cpp
• libcudacxx/test/libcudacxx/cuda/atomics/atomic_16b_host_no_override.fail.cpp
• libcudacxx/test/libcudacxx/cuda/atomics/atomic_16b_ld_st.pass.cpp
• libcudacxx/test/libcudacxx/cuda/atomics/atomic_ref_small.pass.cpp
• libcudacxx/test/libcudacxx/cuda/atomics/bad_atomic_alignment.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/large_type.h
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_arrive_on.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_divergent_threads.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_group_concept.h
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_group_concept_thread_scope_block.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_group_concept_thread_scope_device.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_group_concept_thread_scope_system.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_producer_consumer.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_block_16.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_block_32.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_block_64.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_block_8.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_block_large_type.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_device_16.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_device_32.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_device_64.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_device_8.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_device_large_type.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_generic.h
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_system_16.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_system_32.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_system_64.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_system_8.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_system_large_type.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_thread.h
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_thread_16.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_thread_32.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_thread_8.pass.cpp
• libcudacxx/test/libcudacxx/cuda/pipeline/pipeline_memcpy_async_thread_scope_thread_large_type.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/atomic_cuda_float.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/atomic_cuda_generic.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/atomic_cuda_signed.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/atomic_cuda_unsigned.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/atomic_std_float.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/atomic_std_generic.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/atomic_std_signed.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/atomic_std_unsigned.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/common.h
• libcudacxx/test/libcudacxx/heterogeneous/atomic/flag.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/reference_cuda.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/atomic/reference_std.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/barrier.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/barrier_parity.cuda.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/barrier_parity.std.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/helpers.h
• libcudacxx/test/libcudacxx/heterogeneous/latch.pass.cpp
• libcudacxx/test/libcudacxx/heterogeneous/semaphore.pass.cpp
• libcudacxx/test/libcudacxx/libcxx/atomics/atomics.flag/init_bool.pass.cpp
• libcudacxx/test/libcudacxx/libcxx/atomics/diagnose_invalid_memory_order.fail.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.fences/atomic_signal_fence.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.fences/atomic_thread_fence.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.flag/atomic_flag_clear.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.flag/atomic_flag_clear_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.flag/atomic_flag_test_and_set.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.flag/atomic_flag_test_and_set_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.flag/atomic_flag_wait.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.flag/clear.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.flag/default.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.flag/test_and_set.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.general/replace_failure_order.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.lockfree/isalwayslockfree.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.lockfree/lockfree.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.order/kill_dependency.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.order/memory_order.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/address.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/address_ref.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/address_ref_constness.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/atomic_copyable.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/atomic_ref_address.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/bool.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/cstdint_typedefs.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/enum_class.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/floating_point.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/floating_point_ref.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/floating_point_ref_constness.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/16b_integral_ref.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/1b_integral_cuda.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/1b_integral_std.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/2b_integral_cuda.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/2b_integral_std.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/4b_integral_cuda.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/4b_integral_std.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/8b_integral_cuda.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/8b_integral_std.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/common.h
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/integral_ref.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral/integral_ref_constness.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/integral_typedefs.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/non_arithmetic.fail.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/non_trivial.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/supported_sizes.fail.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/supported_sizes.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/trivially_copyable.fail.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/trivially_copyable.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/trivially_copyable_ref.fail.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_compare_exchange_strong.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_compare_exchange_strong_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_compare_exchange_weak.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_compare_exchange_weak_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_exchange.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_exchange_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_fetch_add.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_fetch_add_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_fetch_and.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_fetch_and_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_fetch_or.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_fetch_or_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_fetch_sub.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_fetch_sub_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_fetch_xor.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_fetch_xor_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_helpers.h
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_init.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_is_lock_free.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_load.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_load_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_store.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/atomic_store_explicit.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.req/ctor.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.wait/atomic_ref_member_wait.pass.cpp
• libcudacxx/test/libcudacxx/std/atomics/atomics.types.operations/atomics.types.operations.wait/atomic_wait.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.barrier/arrive.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.barrier/arrive_and_drop.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.barrier/arrive_and_wait.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.barrier/completion.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.barrier/try_wait_for.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.barrier/try_wait_parity_for.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.barrier/try_wait_parity_until.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.barrier/try_wait_until.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.latch/arrive_and_wait.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.latch/count_down.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.latch/try_wait.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.semaphore/acquire.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.semaphore/release.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.semaphore/timed.pass.cpp
• libcudacxx/test/libcudacxx/std/thread/thread.semaphore/try_acquire.pass.cpp
• libcudacxx/test/support/cmpxchg_loop.h
• libcudacxx/test/support/concurrent_agents.h
• libcudacxx/test/support/cuda_space_selector.h
• libcudacxx/test/support/test_macros.h

[miscco]

pre-commit.ci autofix

[github-actions]

🥳 CI Workflow Results

🟩 Finished in 5h 04m: Pass: 100%/115 | Total: 4d 19h | Max: 4h 33m | Hits: 51%/1121303

See results here.