importer does not recognize mbarrier_arrive_no_complete; kernels fail verification with symbol-not-found #774

Activity

1. nihalpasham

   added

   bugSomething isn't workingSomething isn't working

   intrinsicsDevice intrinsics and libdevice math mappingsDevice intrinsics and libdevice math mappings

   on Aug 10, 2026
2. vyncint commented on Aug 10, 2026

   

   vyncint

   on Aug 10, 2026

   Contributor

   More actions

   Looked into routing this through the generated path, and hit a blocker worth
   raising before anyone writes the record: the current signature cannot express
   the intrinsic, so this needs an API decision rather than just a table entry.

   intrinsics/imported.json has exactly two noComplete arrive records, and both
   take a count:

   int_nvvm_mbarrier_arrive_noComplete         args=['ptr', 'i32']         results=['i64']
   int_nvvm_mbarrier_arrive_noComplete_shared  args=['shared_ptr', 'i32']  results=['i64']
   

   There is no zero-count variant. The stub is

   pub unsafe fn mbarrier_arrive_no_complete(bar: *const Barrier)   // no count, no token

   so it is short one operand and drops the i64 state. Its doc block also shows

   mbarrier.arrive.noComplete.shared.b64 _, [addr];
   

   which is not a form the ISA accepts — the real instruction is
   mbarrier.arrive.noComplete.shared.b64 $state, [$addr], $count;. The
   already-generated sibling shows the shape the family uses:
   crates/cuda-device/src/generated/mbarrier_basic.rs has
   pub unsafe fn mbarrier_arrive(bar: *const Barrier) -> u64, i.e. generated and
   token-returning, while no_complete is hand-written in barrier.rs and
   predates that path.

   So "route it through the generated path" also means changing the public
   signature. Two ways, and the choice is yours:

   1. Match the intrinsic — mbarrier_arrive_no_complete(bar: *const Barrier, count: u32) -> u64. Consistent with mbarrier_arrive, exposes the count the
      ISA requires, lets the caller keep or drop the token.
   2. Keep the fire-and-forget shape — leave the signature alone and have the
      generated wrapper pass count = 1 and discard the state. Preserves the
      documented intent and breaks no source, but bakes in a count the caller
      cannot choose.

   Either is a real API break only on paper: the function has zero call sites in
   the tree and currently fails module verification for anyone who calls it, so
   nothing working depends on today's shape.

   Happy to implement whichever you prefer — including the ledger entry, dialect op,
   regeneration, and an sm_80 example (the A10G I test on runs mbarrier, so this one
   gets real execution evidence rather than compile-only). Just say which signature.
3. nihalpasham commented on Aug 14, 2026

   

   nihalpasham

   on Aug 14, 2026

   CollaboratorAuthor

   More actions

   I recommend we take option 1 and make this the exact low-level primitive:

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

   The count is a defining operand of the PTX/NVVM operation, and keeping the returned state preserves the real contract; callers that only need signaling can explicitly discard it. I don’t want the public intrinsic to silently bake in count = 1.

   The current function does not work through device codegen and has no in-tree callers, so correcting its signature is the right tradeoff. We should route it through the generated catalog/ABI path, remove the stale handwritten stub, and document the important safety precondition: the count must be valid and this operation must not complete the current barrier phase.

   Exact dynamic-count and returned-state checks for both backends, plus the sm_80 run you proposed, would cover it well.
4. uurl

   marked Implement generated mbarrier_arrive_no_complete intrinsic #1100 as a duplicate of this issue on Aug 20, 2026
5. nihalpasham

   added a commit that references this issue on Aug 30, 2026

   feat(intrinsics): route mbarrier_arrive_no_complete through the gener…

   ...

   9231490
6. nihalpasham

   closed this as completedin #1101on Aug 30, 2026
7. nihalpasham

   added a commit that references this issue on Aug 30, 2026

   Implement generated mbarrier_arrive_no_complete intrinsic (#1101)

   ...

   1c0b81a