libcu++: support subword extended floating-point atomic min and max.

[griwes]

Description

Resolves #10810.

cuda::atomic<__half> and cuda::atomic<__nv_bfloat16> support basic atomic operations but currently reject fetch_min and fetch_max because the small-storage dispatch doesn't consider them to be floating point types. This PR changes the check to accept extended fp, and enables those operations using a widened CAS loop that decodes the stored value, performs a floating-point comparison, and writes the selected value back through the existing 32-bit representation.

Checklist

• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[coderabbitai]

📝 Walkthrough

Summary by CodeRabbit

• New Features

  • Added atomic fetch-min and fetch-max support for two-byte extended floating-point types, including CUDA half and bfloat16 values.
  • Extended support across local, shared, and global memory scenarios.
  • Improved compatibility across supported CUDA toolchain versions.

• Bug Fixes

  • Improved extended floating-point comparisons for more consistent atomic operation results.

• Tests

  • Expanded coverage for half and bfloat16 atomic operations across supported memory types.

Walkthrough

Changes

The small atomic dispatch supports fetch_min and fetch_max for two-byte extended floating-point types through compare-exchange loops. Tests cover __half and conditional __nv_bfloat16 behavior across memory scopes.

Assessment against linked issues

Objective	Addressed	Explanation
Enable owning cuda::atomic<__half> and cuda::atomic<__nv_bfloat16> fetch_min and fetch_max operations [#10810]	✅
Use floating-point comparisons within the widened CAS loop while preserving the existing atomic storage path [#10810]	✅
Convert the half compile-failure test into positive behavioral coverage and add bfloat16 coverage [#10810]	✅

Suggested reviewers: pciolkosz, miscco

Merge Risk: 🟡 Moderate · up to 05d3f

The PR enables extended floating-point atomic min/max, but it is not merge-ready until the CAS-loop memory-order handling is corrected or explicitly accepted; supported operations may otherwise use invalid failure orders, and several tests may exercise the wrong atomic scope.

---

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

[coderabbitai]

Actionable comments posted: 2

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 0bee177c-bd68-4a92-a4ab-5a391c37bb1e

📥 Commits

Reviewing files that changed from the base of the PR and between 2deb899 and feca2cd.

📒 Files selected for processing (2)

• libcudacxx/include/cuda/std/__atomic/types/small.h
• libcudacxx/test/libcudacxx/cuda/atomics/atomic.ext/atomic_fetch_extended_floating_point.pass.cpp

[coderabbitai]

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)

libcudacxx/test/libcudacxx/cuda/atomics/atomic.ext/atomic_fetch_extended_floating_point.pass.cpp (1)

82-112: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

important: Use ThreadScope in all four fetch_max cases. cuda::atomic<T> defaults to cuda::thread_scope_system, while these TestFn instantiations pass cuda::thread_scope::thread_scope_thread. Change the aliases to cuda::atomic<T, ThreadScope>.

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: f5c551be-6e53-490d-80dd-ad498c610e5f

📥 Commits

Reviewing files that changed from the base of the PR and between ab55f2b and 05d3f3c.

📒 Files selected for processing (2)

• libcudacxx/include/cuda/std/__atomic/types/small.h
• libcudacxx/test/libcudacxx/cuda/atomics/atomic.ext/atomic_fetch_extended_floating_point.pass.cpp

Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.

[github-actions]

⏱️ CCCL compile-time benchmark comparison: Public headers compile-time bench

Result: 0 regression row(s), 3 improvement row(s) above threshold.

Run	Value
Config	public-headers-gcc13
Baseline	origin/main
Preset	all-dev
Targets	cub.headers.base, thrust.cpp.cuda.headers.base, libcudacxx.test.public_headers
GPU / launch args	rtx2080 / --cuda 13.3 --host gcc13

Artifacts: reports and traces

Direct file processing

-f file-processing exclusive --sort total

🟢 Direct file processing — Improvements

Rank	Improvement impact	Selected Δ	Baseline	Current	Event	Matched traces
1	0.612848	-0.612848	5.353479	4.740631	Processing Header File: libcudacxx/include/cuda/std/__cccl/prologue.h	552
2	0.250167	-0.250167	0.250841	0.000674	Processing Header File: cub/cub/device/device_run_length_encode.cuh	2
3	0.200135	-0.200135	1.663065	1.462930	Processing Header File: libcudacxx/include/cuda/std/__cccl/epilogue.h	552

[github-actions]

🥳 CI Workflow Results

🟩 Finished in 2h 52m: Pass: 100%/117 | Total: 4d 07h | Max: 2h 52m | Hits: 60%/638898

See results here.

[miscco]

Looks "good" for extended floating point levels at least ^^

[miscco]

There is _LIBCUDACXX_HAS_NVBF16() which does the additional CTK check