NVIDIA / NVIDIA/cccl

Specializations for `atomic[_ref]<T>` for __half and bfloat16

Open
#6,161 0 comments 0 reactions 0 assignees View on GitHub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

Current state:
- `cuda::atomic[_ref]` and `cuda::std::atomic[_ref]` compile for `__half` / `__nv_bfloat16` basic ops: load/store/exchange/CAS.
- They do not provide arithmetic ops: `fetch_add`, `fetch_sub`, `+=`, `-=`, or free `cuda::std::atomic_fetch_add[_explicit]` / `atomic_fetch_sub[_explicit]`.

Why:
- `__atomic_impl` / `__atomic_ref_impl` select arithmetic support via `cuda::std::is_floating_point_v`.
- `is_floating_point_v` only covers `float`, `double`, and `long double`.
- `__half` / `__nv_bfloat16` are only classified as extended floating-point types.
- Simply changing the trait is likely insufficient: the small-type storage path proxies `<4B` types through 32-bit integer atomics.

Work:
- Add explicit arithmetic atomic support for `__half` and `__nv_bfloat16`.
- Use native device support where available: `__half` on sm_70+, `__nv_bfloat16` on sm_80+.
- Decide fallback / compile-time behavior for older architectures.
- Cover `cuda::atomic`, `cuda::atomic_ref`, `cuda::std::atomic`, and `cuda::std::atomic_ref`.
- Add heterogeneous tests for member ops and free-function ops.

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.