[uurl]

Closes #774.

Summary

Implements mbarrier_arrive_no_complete(bar, count) -> u64 through the
generated intrinsic catalog/ABI path and removes the stale handwritten stub.

The intrinsic maps to:

llvm.nvvm.mbarrier.arrive.noComplete.shared

and lowers to:

mbarrier.arrive.noComplete.shared.b64

with a dynamic u32 arrival count and returned u64 barrier state.

Changes

• adds ABI entry i1016

• adds mbarrier_arrive_no_complete to the mbarrier_basic catalog family

• exposes:

  pub unsafe fn mbarrier_arrive_no_complete(
      bar: *const Barrier,
      count: u32,
  ) -> u64

• removes the stale handwritten no-count/no-return implementation

• generates the public CUDA API and compatibility path

• adds MbarrierArriveNoCompleteSharedOp

• adds MIR importer and lowering support

• adds LLVM NVPTX lowering through:

  llvm.nvvm.mbarrier.arrive.noComplete.shared
  

• adds libNVVM lowering through inline PTX:

  mbarrier.arrive.noComplete.shared.b64
  

• adds verifier coverage for:

  • two operands
  • dynamic i32 count
  • i64 result

• adds the generated LLVM probe

• adds LLVM and CUDA 13.3/libNVVM evidence

• updates the barrier example with a regression that verifies the no-complete phase behavior using the returned barrier state

Validation

• cargo test -p cuda-intrinsics-gen

  • 293 passed

• cargo run -p cuda-intrinsics-gen -- check

  • passed

• cargo test -p dialect-nvvm

  • 51 passed

• cargo test -p mir-lower

  • 263 passed

• MIR lowering tests

  • 117 passed

• cargo fmt --all -- --check

  • passed

• git diff --check

  • passed

• compile-only smoketest coverage

  • full suite: 215 / 216 passed
  • the remaining cutile_inter_kernel failure was traced to missing CUDA toolkit host headers; after completing the isolated CUDA 13.3 toolkit, its targeted compile-only rerun passed 1 / 1
  • all 216 examples were therefore observed compiling successfully on the same code revision

LLVM NVPTX

The typed NVVM intrinsic was compiled for sm_80 and emitted:

mbarrier.arrive.noComplete.shared.b64

with the dynamic count operand and returned barrier state preserved.

CUDA 13.3 ptxas accepted the generated PTX for sm_80.

libNVVM

CUDA 13.3 libNVVM successfully compiled the production barrier no-complete regression to LTOIR.

nvJitLink subsequently produced an sm_80 cubin containing the
barrier_no_complete_test kernel.

Runtime validation

Physical runtime validation on sm_80 hardware remains unexecuted.

The evidence profiles record this explicitly as:

runtime_validation: unexecuted

Compile, PTX assembly, libNVVM, and device-link validation for sm_80 have all completed successfully.

[nihalpasham]

Putting this PR on hold until the pointer-contract refactor is stable and the branch can be reviewed on that architecture.

• importer does not recognize mbarrier_arrive_no_complete; kernels fail verification with symbol-not-found #774 already approved the exact (bar, count) -> u64 API, so that compatibility decision is settled.
• ABI ID i1016 is already assigned on current main; this entry must use a new append-only ID and all generated outputs must be refreshed.
• The shared-pointer conversion and lowering path must be rebased and revalidated against Preserve Rust reference kind and exclusivity through MIR import and lowering #413.
• The final branch needs a real sm80 fail-before/pass-after run, not only generated-code checks.

Review can resume after the refactor rebase, ABI regeneration, and runtime evidence required by the Discord acceptance criteria.

[nihalpasham]

Took this off hold and finished it ourselves; every blocker on it was churn from our side (the #1190 generator split, the LLVM 23 bump, #1104 taking your ABI id), and the design was already agreed in the #774 thread. Your commit is the base of the new head with authorship and sign-off intact.

What the port changed (mechanics only, zero design changes)

• ABI id i1016 -> i1017 (main spent i1016 on ex2_approx_f16); ledger check passes against main: "preserves 1016 base entries, 1017 total"
• generator/render/importer/lowering hunks re-homed into the post-split module layout (resolve.rs, render.rs and friends no longer exist as files)
• evidence re-recorded rather than ported: your rust-llvm-22.1.2 profile is obsolete, so the record was re-attested end to end under rust-llvm-23.1.0 with the local CUDA 13.3 ptxas, and the libNVVM record's LTOIR/cubin sizes+hashes are from local runs (nvJitLink 13.3, SASS shows the counted arrival)

Runtime evidence, the part the hold was really about

Before (old stub, on main):  kernel calling mbarrier_arrive_no_complete
  -> [rustc_codegen_cuda] Device codegen failed: ... Symbol
     cuda_device__barrier__mbarrier_arrive_no_complete not found  (exit 101)

After (this branch, RTX 5090):  cargo oxide run barrier
  init(N+1); state = arrive_no_complete(&BAR, N)
  test_wait(state) == false     // phase held open   -> out[0] = 0
  arrive(&BAR)                  // one normal arrival
  test_wait(state) == true      // phase completes   -> out[1] = 1
  all three barrier tests pass

Gates on the final head: generate/check round-trip clean, probe gate (2366 target configs across 1017 routes), full workspace tests, fmt/clippy, per-example gates, lockfile check.

One ordering note: #1139 also wants the ledger tail. This branch assumes it lands first; whichever goes second regenerates, which is one command on our side. Dropping the on-hold/Depends labels; CI approval + landing next.