Skip to content

feat(intrinsics): add CTA barrier reductions - #1138

Closed
almightychang wants to merge 3 commits into
NVIDIA:mainfrom
almightychang:barrier-reduction
Closed

almightychang wants to merge 3 commits into
NVIDIA:mainfrom
almightychang:barrier-reduction

Conversation

@almightychang

Copy link
Copy Markdown
Contributor

Summary

Admit 13 CTA barrier intrinsics end to end: 12 reduction forms and barrier.cta.sync.all. LLVM imported these records, but the catalog did not admit them. PR #721 already admitted the four counted barrier.cta.{sync,arrive}[.aligned].count siblings, so this PTX family was only half admitted.

Users already had sync_threads(), but there was no path corresponding to CUDA's __syncthreads_and, __syncthreads_or, or __syncthreads_count, including no inline-assembly route.

Changes

  • Add AND, OR, and POPC CTA reductions in all and count shapes, with plain barrier.red and aligned bar.red spellings: 12 intrinsics total.
  • Add numbered all-thread barrier.cta.sync.all.
  • Add raw unsafe APIs, generated dialect/import/export surfaces, LLVM-NVPTX and libNVVM lowering routes, evidence, and documentation.
  • Materialize AND/OR predicate results through selp.b32 followed by truncation to i1; keep POPC results as direct u32 values.
  • Add table-driven lowering tests for AND/OR/POPC, both argument shapes, and both backends.
  • Extend the existing counted_barrier example with compile-only coverage for all 13 additions: 4 AND forms, 4 OR forms, 4 POPC forms, and barrier_cta_sync_all.

This PR keeps the libNVVM route at minimum_sm = sm_75. That is the CUDA 12.9 evidence target, not the instruction floor. It will be lowered only after evidence is produced at a lower target. The counted-barrier siblings already use the same evidence floor.

This PR exposes only raw APIs that take a barrier ID; count forms also take the participant count. CUDA-style barrier-ID-free wrappers remain out of scope. They can be added after a separate safety and barrier-ID aliasing contract is agreed.

The current branch-generated ABI entries occupy i1016..i1028, but that range is not a merge contract. Open PRs #1101 and #1104 also start at i1016. Immediately before merge, this branch will rebase onto the latest main, append its entries after the then-current ledger end, and regenerate all generated files.

Testing

  • just check passes (the local mirror of CI: cargo oxide fmt --check, clippy, tests, CUDA-linked tests, guards, generated-intrinsic checks, ABI history, rustdoc, and doctests). Exit 0 with local CUDA 12.9 and pinned LLVM 22.1.2.
  • cargo oxide run <example> passes, or scripts/smoketest.sh -o '^<example>$' — runtime was not executed because this machine has no usable GPU (nvidia-smi reports a driver/NVML communication failure).
  • New example added (not applicable: this extends the existing counted_barrier example).

Additional checks actually run:

  • bash scripts/check-book-catalog-stamp.sh: exit 0; stamp 9b5ce71c, 588 total NVVM operations, 562 generated operations, 1028 catalog entries.
  • scripts/smoketest.sh --compile-only --no-color --only '^counted_barrier$': Pass: 1 / 1, using local CUDA 12.9 and pinned LLVM 22.1.2. The compile-only kernel covers all 13 additions listed above; it is deliberately not launched.
  • cargo run -p cuda-intrinsics-gen -- probe --all --skip-terminal: all 1028 LLVM routes passed. The pinned llc selected barrier.red.and.pred, barrier.red.or.pred, and barrier.red.popc.u32, plus aligned forms and barrier.sync.
  • CUDA 12.9.86 ptxas assembled the recorded PTX evidence. Evidence profile names identify the local CUDA 12.9.86 toolkit.
  • Sphinx -W --keep-going build: succeeded.
  • git diff --check main..HEAD: exit 0.
  • cargo run -p cuda-intrinsics-gen -- check-abi-history --base-ref b7fe543b: intrinsics/abi-v1.toml preserves 1015 base entries and has 1028 total entries. Run explicitly against the PR base, since just check defaults this to HEAD^, which is not the pre-ABI base on a three-commit branch.

Not verified locally:

  • LLVM 21 end-to-end lowering. This machine has LLVM 22, not llc-21/opt-21.
  • CUDA 13.1.1 ptxas acceptance. Local terminal evidence uses CUDA 12.9.86.
  • CUDA 13.1 libNVVM behavior. Local libNVVM evidence uses CUDA 12.9.86.

runtime_validation remains unexecuted; compile and assembly evidence does not claim runtime semantics.

just check deliberately excludes these CI jobs, as documented in the Justfile: per-example clippy, examples-compile, book, and CodeQL. The focused changed-example compile and book build above cover local substitutes; CodeQL and the exact CUDA 13.1.1/LLVM 21 environment remain CI-only.

Checklist

  • All commits signed off (git commit -s)
  • SPDX headers on new source files

Questions about the review? Ping us in #contributors on Discord.

Signed-off-by: Sam, Joochul Chang <almightychang@icloud.com>
Signed-off-by: Sam, Joochul Chang <almightychang@icloud.com>
Signed-off-by: Sam, Joochul Chang <almightychang@icloud.com>
@nihalpasham nihalpasham added enhancement New feature or request intrinsics Device intrinsics and libdevice math mappings codegen Device code-generation pipeline (Rust MIR to IR to PTX) IR-lowering Lowering between dialects (dialect-mir to LLVM dialect) cuda-feature A CUDA hardware/toolkit feature (TMA, tcgen05, WGMMA) device-apis User-facing device-side and kernel-authoring APIs (cuda-device, cuda-macros) tests Adds or improves test coverage examples Example kernels and host harnesses under crates/rustc-codegen-cuda/examples documentation Improvements or additions to documentation safety Memory safety, soundness, or undefined behavior miscompile Produces incorrect PTX/IR (wrong generated code) soundness-audit Correctness case under active maintainer audit; not cleared for merge yet labels Aug 23, 2026
@nihalpasham

nihalpasham commented Aug 23, 2026 •

Copy link
Copy Markdown
Collaborator

@nihalpasham nihalpasham added needs-implementation and removed miscompile Produces incorrect PTX/IR (wrong generated code) labels Aug 23, 2026

@nihalpasham nihalpasham left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

  • Benefit required: name the concrete workload and show whether fused barrier reduction removes synchronization/instructions, improves performance, or closes a safety/API gap. The current PR adds surface area without establishing that benefit.
  • Compiler blocker: LLVM 21 leaves all 12 reduction calls unresolved, so ptxas rejects the generated PTX.
  • libNVVM blocker: sync_all emits one operand but declares r,r,~{memory}. The extra constraint makes the inline assembly contract invalid.
  • Metadata blocker: aligned bar.red forms are assigned later PTX/SM floors even though the instruction family is valid at PTX 2.0 / sm_20.
  • Safety/test blocker: the unsafe contract is incomplete, some generated docs contain unrelated atomic text, and the example never launches a kernel or checks a result.
  • A complete repair needs an LLVM fallback or inline-PTX lowering, corrected libNVVM constraints and canonical floors, regenerated docs/ABI, plus a real GPU oracle and sanitizer run. Generator and MIR-lowering suites pass locally, but they do not cover these failures.

@nihalpasham

Copy link
Copy Markdown
Collaborator

Reviewed this as part of a full pass over the open queue. Closing it. The catalog plumbing itself follows the #721 counted-barrier route correctly, so this isn't about the approach; it's about what's broken in the current head and what's missing around it.

Correctness problems we verified by running things, not just reading

  • llc-21 silently produces garbage. The llvm.nvvm.barrier.cta.red.* records only exist from LLVM 22.1.0 (llvm-project a40f4442652c). llc-21 is still in our fallback discovery list, and it exits 0 on these calls:

    llc-22:  barrier.red.and.pred  %p2, 0, %p1;          // real instruction
    llc-21:  .extern .func llvm.nvvm.barrier.cta.red.and.all
             call.uni ...                                 // ptxas: unresolved extern
    

    No version gate, no fallback, no diagnostic on our side. The failure only shows up when ptxas rejects the PTX.

  • barrier_cta_sync_all emits invalid IR on the libNVVM route. The arity check was relaxed to allow 1 operand, but the inline-asm constraint string is still hardcoded "r,r,~{memory}" (two inputs). One operand against a two-input constraint is malformed IR. The recorded libNVVM evidence was a hand-built fixture rather than real lowering output, which is why nothing caught it.

  • Wrong hardware floors in both directions. The aligned bar.red forms claim PTX 1.0 with no sm floor at the top level and PTX 3.2 on the llvm route; the PTX ISA says PTX 2.0 / sm_20.

  • Generated safety docs are for a different instruction. The barrier-reduction ABI entries tell callers that addr must point at four writable bytes in global memory. These functions have no addr argument; the text fell through to the packed-atomic renderer arm. The PR also waters down the existing counted-barrier safety text from "compatible barrier ID and expected thread count" to a generic one-liner.

And the bar itself

Features need a concrete consumer and evidenced benefit. This adds 13 raw unsafe intrinsics while explicitly deferring the CUDA-parity wrappers (__syncthreads_and/or/count equivalents), so the one consumer story the PR tells still can't be written after merge. Nothing was ever launched on a GPU; the example is compile-only and the lowering tests assert the implementation's own format strings back at themselves.

If you want to bring this back

  • ship the wrappers (or name a real consumer kernel) with a launched numeric test (AND/OR/POPC values checked on hardware)
  • gate or fall back on LLVM < 22, or we formally raise the llc floor first
  • derive the libNVVM constraint string from operand count + regression test for the 1-operand path
  • correct the floors, add a barrier_reduction safety arm to the renderer, restore the counted-barrier contracts
  • rebase onto current main (the ledger tail moved past i1016, and it's moving again in this round)

The instruction family is worth having; this needs to arrive working and with a consumer.

@nihalpasham nihalpasham added the declined PR intentionally not merged (closed without merging) label Aug 30, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

codegen Device code-generation pipeline (Rust MIR to IR to PTX) cuda-feature A CUDA hardware/toolkit feature (TMA, tcgen05, WGMMA) declined PR intentionally not merged (closed without merging) device-apis User-facing device-side and kernel-authoring APIs (cuda-device, cuda-macros) documentation Improvements or additions to documentation enhancement New feature or request examples Example kernels and host harnesses under crates/rustc-codegen-cuda/examples intrinsics Device intrinsics and libdevice math mappings IR-lowering Lowering between dialects (dialect-mir to LLVM dialect) needs-implementation safety Memory safety, soundness, or undefined behavior soundness-audit Correctness case under active maintainer audit; not cleared for merge yet tests Adds or improves test coverage

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants