Implement generated mbarrier_arrive_no_complete intrinsic #1100

Summary

Add support for the counted shared-memory mbarrier.arrive.noComplete operation through the generated intrinsic catalog/ABI pipeline.

The public API should expose:

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

The intrinsic corresponds to the NVVM intrinsic:

llvm.nvvm.mbarrier.arrive.noComplete.shared

and should lower to:

mbarrier.arrive.noComplete.shared.b64 state, [bar], count;

Requirements

• implement the intrinsic through the generated catalog/ABI path

• remove the stale handwritten mbarrier_arrive_no_complete stub

• accept a dynamic u32 arrival count

• return the opaque u64 barrier state

• add an NVVM dialect operation with two operands and one result

• support both intrinsic backends:

  • LLVM NVPTX via the typed NVVM intrinsic
  • libNVVM via inline PTX

• preserve convergent and shared-memory side-effect semantics

• require PTX 7.0+ and sm_80+

• generate the corresponding LLVM probe

• add importer, verifier, lowering, and codegen coverage

• add a barrier regression exercising the returned state and dynamic count

Validation

The implementation should verify that:

1. mbarrier_arrive_no_complete(bar, count) does not complete the current phase when the supplied count leaves outstanding arrivals.
2. The returned state can be passed to mbarrier_test_wait.
3. mbarrier_test_wait remains false before the final arrival.
4. A subsequent normal arrival completes the phase.
5. mbarrier_test_wait becomes true for the previously returned state.

Backend evidence should confirm the expected PTX for sm_80.

Physical runtime validation should be performed on sm_80 hardware when available.