Codegen collectives / send-recv + the KLR↔BIR op-join table

Version pin: neuronx_cc 2.24.5133.0+58f8de22 (cp310; cp311/cp312 twins present). Subject binaries: neuronxcc/starfish/lib/libwalrus.so (the Strand-I backend object — every KlirToBirCodegen::codegen* leaf body) and neuronxcc/starfish/lib/libBIR.so (the InstBuilder::addLocalSendRecv definition; the libwalrus copy is an extern U-row resolved through the PLT). Both retain their full dynamic symbol tables, so every leaf below is a byte-exact nm -DC row. libwalrus .text (0xf00000+) and .rodata (0x1c72000+) have VA == file-offset, so jump tables read by file offset. Evidence is nm, IDA decompiled/disasm sidecars under ida/…__libwalrus.so/ and …__libBIR.so/, and strings. The two-VA-frame artifact applies: each leaf appears twice in the IDA sidecars — once at its PLT-thunk frame (_0x6000xx) and once at its real body frame (_0xf2xxxx); addresses below are always the real nm body. Every claim is tagged CONFIRMED / STRONG / INFERRED / SPECULATIVE.

Abstract

This page does two things. First, it documents the collective / send-recv codegen leaves — the corner of the KlirToBirCodegen lowering that turns KLR collective communication operators into bir::InstCollectiveCompute (IT 48) nodes. There are four such leaves, and contrary to the naïve "all collectives share one encoder" reading, they split three ways: an explicit-permute leaf that builds the instruction itself, a thin leaf that delegates to the shared group-collective encoder, a dedicated fused permute-reduce leaf, and SendRecv, which is not a CollectiveOp at all and reaches a different InstBuilder entry point. Second, it assembles the capstone of the codegen series: the complete klr→BIR op-join table — every dispatched codegen<Op> leaf, its address, and the bir::Inst* / InstructionType opcode it emits — annotated with the five 1→many fan-outs, the fan-ins (several klr ops collapsing onto one opcode), and the BIR-only opcodes that no klr leaf ever produces. The join is the architecture-independent waist of the four-layer lowering: klr::<Op> → bir::Inst* (here) → per-arch CoreV* ISA descriptor → NEFF wire bytes.

This page is the leaf companion to the dispatch core (7.21 KlirToBirCodegen dispatch, the 77-case switch and refcount-shim mechanics) and the beta3 collective counterpart 6.5.13 BirCodeGenLoop collective lowering. The CollectiveKind enum itself — the field every leaf here writes — is catalogued in 7.9 Op-family enums, which establishes that one opcode (IT 48) multiplexes all 11 CollectiveKinds.

1. The collective / send-recv codegen leaves

1.1 The three-way split (and why the "all-shared" reading is wrong)

A KLR kernel names collectives through six klr::CollectiveOp-typed operator arms plus a separate klr::SendRecv class. The dispatch switch (the codegenOperator jump table, 7.21) routes each to a codegen<Op> leaf. The temptation is to assume all collective arms are thin shims onto one shared encoder. They are not. Reading the four leaf bodies shows three distinct lowering styles:

The four group collectives — AllReduce (k52), AllGather (k53), ReduceScatter (k54), AllToAll (k59) — are the canonical thin pattern: each bumps the shared_ptr refcount, calls the shared codegenCollectiveOp @ 0xf229a0, then patches one op-specific field. So of the seven CollectiveOp arms, five are thin (four group + PermuteImplicit) and two are dedicated (Permute, PermuteImplicitReduce). [CONFIRMED — all eight addresses are nm -DC rows; codegenSendRecv@0xf1d420, codegenCollectivePermute@0xf22110, codegenCollectiveOp@0xf229a0, codegenCollectivePermuteImplicit@0xf23a40, codegenCollectivePermuteImplicitReduce@0xf21850, codegenAllReduce@0xf23510, codegenAllGather@0xf23650, codegenReduceScatter@0xf23780, codegenAllToAll@0xf23910 all verified.]

QUIRK — the shared encoder can encode kinds 7/9/10, but only 9 reaches it. The shared codegenCollectiveOp @ 0xf229a0 carries a name switch(kind) whose arms include case 7 → "CollectivePermute", case 9 → "CollectivePermuteImplicit", case 10 → "CollectivePermuteImplicitReduce". So the encoder is capable of producing all three permute variants. But the KLR dispatch only ever reaches it for kind 9 (PermuteImplicit); kinds 7 and 10 have their own dedicated bodies that build IT48 directly and merely reuse the kind-string via insertElement, not the shared encoder. The capable-but-unreached arms are dead-on-this-path — a sign these leaves were forked off a common ancestor and the explicit/reduce variants grew their own operand-binding logic. [CONFIRMED — name strings CollectivePermute, CollectivePermuteImplicit, CollectivePermuteImplicitReduce present in .rodata; the two dedicated bodies do not call codegenCollectiveOp.]

CRITICAL DISTINCTION — Operator-kind ≠ CollectiveKind. The Operator-kind (k52..k59) is a libwalrus jump-table index into the codegenOperator switch. bir::CollectiveKind (0..10) is the wire-level algorithm selector written into InstCollectiveCompute+0xF8. They are unrelated namespaces. Dispatch-kind 55 (CollectivePermute) binds CollectiveKind 7 (Permute); dispatch-kind 42 (SendRecv) binds CollectiveKind 0 (SendRecv). The leaf is what binds the CollectiveKind, and that binding is the single decisive choice — it selects the downstream on-chip-vs-remote fate (§4). [CONFIRMED]

1.2 The klr::CollectiveOp / klr::SendRecv field layouts

The shared klr::CollectiveOp carries everything the seven arms need; each arm reads a different subset. Offsets are read from the per-field flag bytes and size reads that klr::to_string(CollectiveOp&) @ 0xf95120 walks:

// klr::CollectiveOp  (to_string prints: dsts, srcs, op, replicaGroup,
//                     concatDim, sourceTargetPairs, channel_id, num_channels)
struct CollectiveOp {
    list<shared_ptr<TensorRef>> dsts;       // +0   sentinel = obj+0
    u64                         dsts_size;  // +16
    list<shared_ptr<TensorRef>> srcs;       // +24  sentinel = obj+24
    u64                         srcs_size;  // +40
    klr::AluOp                  op;         // +48  (reduce-op value)    flag @+52
    shared_ptr<ReplicaGroup>    replicaGroup;// +56  tag = **(+56); 3 = literal
    int                         concatDim;  // +72  flag @+76  (optional)
    optional<list<list<int>>>   srcTgtPairs;// +80  flag @+104, size @+96
    int                         channel_id; // +112 flag @+116 (optional)
    int                         num_channels;//+120 flag @+124 (optional)
};

[CONFIRMED — to_string key order + per-field flag-byte/size reads.]

// klr::SendRecv  (to_string @0xf8e390 prints: dst, src, sendToRank,
//                 recvFromRank, pipeId, useGpsimdDma)
struct SendRecv {
    shared_ptr<TensorRef>  dst;          // +0
    shared_ptr<TensorRef>  src;          // +16
    shared_ptr<Immediate>  sendToRank;   // +32
    shared_ptr<Immediate>  recvFromRank; // +48
    shared_ptr<Immediate>  pipeId;       // +64
    bool                   useGpsimdDma; // (printed via to_string(int)) — IGNORED at codegen (§1.6)
};

[CONFIRMED — the shared_ptr fields are read at a2[0],a2[2],a2[4],a2[6],a2[8] = byte offsets 0/16/32/48/64 inside codegenSendRecv.]

1.3 bir::InstCollectiveCompute (IT48) — the fields the leaves write

All four lowerings populate one IT48 node. These are the offsets the codegen bodies touch, mapped to the wire schema:

The reset-then-store idiom at +280/+320/+432 is the standard std::variant<value,QuasiAffineExpr> (MaybeAffine) assignment: _M_reset only when the tag byte is neither 0 nor -1, then store scalar and clear tag. [STRONG — idiom recognized across all four bodies; offsets cross 7.9 and the IT48 schema. sameInst@0x2f4a80 keying on +248 is CONFIRMED.]

1.4 The two dedicated permute leaves

// (A) codegenCollectivePermute — CollectiveKind 7 — DEDICATED  @0xf22110  [CONFIRMED]
bir::Instruction* codegenCollectivePermute(KlirToBirCodegen *this, klr::CollectiveOp *cp) {
    assert(cp->srcs_size == 1);                 // "guaranteed by frontend API"  (line 0x77A)
    assert(cp->dsts_size == 1);                 //                               (0x77B)
    assert(has_value(cp->srcTgtPairs));         // flag @cp+104                  (0x77C)
    BasicBlock *bb = this->curBB;               // this+72
    auto name = "CollectivePermute-" + returnAndIncrementCCId();
    auto cc   = bb->insertElement<InstCollectiveCompute>(name);   // IT48
    register_inst(this, cc);                                      // sub_F20EC0
    cc->addArgument(codegenTensorRef(cp->srcs[0]));   // 1 input  (*(*(cp+24)+16))
    cc->addOutput  (codegenTensorRef(cp->dsts[0]));   // 1 output (*(*cp+16))
    cc[+248] = 7;                               // CollectiveKind = Permute            (*)
    // for each pair p in cp->srcTgtPairs: assert p.size()==2 (0x78E); emit {src,tgt}
    sub_F1F790(cc + 360, builtPairs);           // → src_target_pairs (+0x168)         (*)
    cc[+320] = -1;                              // stream_id ← unset
    return cc;
    // KEY: replica_groups (+256) left EMPTY; routing is BY src_target_pairs only.
    //      NO reduce-op (+384). NO channel_id. The simulator asserts exactly this:
    //      "Permute should not contain replica_groups".
}

// (C) codegenCollectivePermuteImplicitReduce — CollectiveKind 10 — DEDICATED  @0xf21850  [CONFIRMED]
bir::Instruction* codegenCollectivePermuteImplicitReduce(KlirToBirCodegen *this, klr::CollectiveOp *cpir) {
    assert(cpir->srcs_size == 2);               // permuted-in value + local accumulator (0x7AD)
    assert(cpir->dsts_size == 1);               //                                       (0x7AE)
    assert(has_value(cpir->op));                // flag @cpir+52                          (0x7AF)
    assert((**(cpir+56)) == /*literal*/3);      // replicaGroup->tag == literal          (0x7B0)
    assert(has_value(cpir->channel_id));        // flag @cpir+116                        (0x7B1)
    auto name = "CollectivePermuteImplicitReduce-" + returnAndIncrementCCId();
    auto cc   = this->curBB->insertElement<InstCollectiveCompute>(name);
    register_inst(this, cc);
    cc->addArgument(codegenTensorRef(cpir->srcs[0]));  // peer-routed value
    cc->addArgument(codegenTensorRef(cpir->srcs[1]));  // local accumulator  → 2 in / 1 out (*)
    cc->addOutput  (codegenTensorRef(cpir->dsts[0]));
    cc[+248] = 10;                              // CollectiveKind = PermuteReduceImplicit (*)
    cc[+384] = codegenAluOp(cpir->op);          // reduce-op (AluOpType) @+0x180          (*)
    sub_F1F790(cc + 256, replicaGroups);        // → replica_groups (+0x100)              (*)
    cc[+280] = cpir->channel_id;                // channel_id (MaybeAffine assign)        (*)
    cc[+320] = -1;                              // stream_id unset
    return cc;
    // KEY: the FUSED permute-then-reduce. 2 inputs combined by op. Routing by
    //      replica_groups (implicit). Sim @0x1ebd10: "must have 2 inputs and 1 output";
    //      "Unsupported OP type for Permute Reduce Implicit".
}

[CONFIRMED — both bodies build IT48 inline; the assert line-ids are the cpp:0x7xx references. codegenAluOp@0xf140c0, codegenTensorRef@0xf18b30, returnAndIncrementCCId@0xf13d50 are nm rows.]

1.5 The thin permute leaf and the shared group encoder

// (B) codegenCollectivePermuteImplicit — CollectiveKind 9 — THIN → shared  @0xf23a40  [CONFIRMED]
bir::Instruction* codegenCollectivePermuteImplicit(KlirToBirCodegen *this, klr::CollectiveOp *cpi) {
    assert(cpi->srcs_size == 1);                              // (0x79F)
    assert(cpi->dsts_size == 1);                              // (0x7A0)
    assert(has_value(cpi->channel_id));                      // (0x7A1)
    auto cc = codegenCollectiveOp(this, cpi, /*CollectiveKind=*/9);   // SHARED ENCODER (*)
    cc[+280] = cpi->channel_id;                              // MaybeAffine assign       (*)
    return cc;
    // KEY: routing BY replica_groups (the implicit permutation = a rotate derived from
    //      replica order), NOT explicit pairs. channel_id REQUIRED. Sim:
    //      "PermuteImplicit should not contain src_target_pairs".
}

// shared codegenCollectiveOp  @0xf229a0  (444 L)  — the group/k9 encoder  [CONFIRMED]
bir::Instruction* codegenCollectiveOp(KlirToBirCodegen *this, klr::CollectiveOp *op, CollectiveKind k) {
    const char *kindStr;
    switch (k) {                                  // capable of 2,3,4,5,7,9,10
        case 2: kindStr = "AllReduce";        break;
        case 3: kindStr = "ReduceScatter";    break;
        case 4: kindStr = "AllGather";        break;
        case 5: kindStr = "AllToAll";         break;
        case 7: kindStr = "CollectivePermute";             break;  // reachable only from dedicated body's string reuse
        case 9: kindStr = "CollectivePermuteImplicit";     break;
        case 10:kindStr = "CollectivePermuteImplicitReduce";break;
        default: assert(false);
    }
    assert(!op->srcs.empty());                    // (0x719)
    assert(!op->dsts.empty());                    // (0x71A)
    assert(op->srcs_size == op->dsts_size);       // (0x71B)
    assert((**(op+56)) == /*literal*/3);          // replicaGroup->tag == literal  (0x71C)
    auto name = string(kindStr) + "-" + returnAndIncrementCCId();
    auto cc   = this->curBB->insertElement<InstCollectiveCompute>(name);
    for (auto [src,dst] : zip(op->srcs, op->dsts)) {     // N-in / N-out coalesced
        cc->append_arg(physAP_for(src));     // arg list @cc+160
        cc->append_out(physAP_for(dst));     // out list @cc+192
    }
    cc[+248] = k;                                 // CollectiveKind
    sub_F1F790(cc + 256, op->replicaGroup);       // → replica_groups (+0x100)
    cc[+320] = -1;                                // stream_id unset
    return cc;
}

The four group leaves each call this encoder then patch one field:

codegenAllReduce     @0xf23510 (32 L): assert has_value(op); cc=codegenCollectiveOp(op,2);
                                       cc[+384]=codegenAluOp(op->op);              // reduce-op
codegenAllGather     @0xf23650 (40 L): assert has_value(concatDim); cc=codegenCollectiveOp(op,4);
                                       if(concatDim>1) sub_F12900(cc); cc[+636]=concatDim;  // cc_dim
codegenReduceScatter @0xf23780 (44 L): cc=codegenCollectiveOp(op,3) + reduce-op + cc_dim
codegenAllToAll      @0xf23910 (40 L): cc=codegenCollectiveOp(op,5)

[CONFIRMED — decompiled.]

The routing-data crux — this is the field-set the simulator keys on, and it is the entire reason the permute leaves diverge:

1.6 codegenSendRecv — the non-collective point-to-point path

klr::SendRecv is a separate class, not a CollectiveOp. Its leaf does not touch codegenCollectiveOp; it validates the ranks, guards against self-send/self-recv, and calls bir::InstBuilder::addLocalSendRecv in libBIR:

// codegenSendRecv  @0xf1d420  [CONFIRMED — disasm read]
bir::Instruction* codegenSendRecv(KlirToBirCodegen *this, klr::SendRecv *sr) {
    auto out = codegenTensorRef(sr->dst);                 // **a2   (offset 0)
    auto in  = codegenTensorRef(sr->src);                 // (*a2)[2] (offset 16)
    auto sendTo   = codegenImmediate(sr->sendToRank);     // (*a2)[4]
    assert(holds_alternative<int>(sendTo)   && "The sendToRank must be a integer");    // (0x6D1)
    auto recvFrom = codegenImmediate(sr->recvFromRank);   // (*a2)[6]
    assert(holds_alternative<int>(recvFrom) && "The recvFromRank must be a integer");  // (0x6D5)
    auto pipeId   = codegenImmediate(sr->pipeId);         // (*a2)[8]
    assert(holds_alternative<int>(pipeId)   && "The pipeId must be a integer");        // (0x6D9)

    // self-loop guard against THIS core's id:
    optional<int> myCore = bir::Module::getNeuronCoreId(this->module);   // @0xf1d60e call
    assert(myCore.has_value());
    if ((float)*myCore == (float)recvFrom)
        throw runtime_error("SendRecv instruction on core <N> is receiving from itself!");
    if ((float)*myCore == (float)sendTo)
        throw runtime_error("SendRecv instruction on core <N> is sending to itself!");

    auto name = "LocalSendRecv-" + returnAndIncrementCCId();
    auto cc = bir::InstBuilder::addLocalSendRecv(            // @0xf1d768 call
                  name, /*src*/in, /*dst*/out, inAP, outAP,
                  /*sendToRank&*/sendTo, /*recvFromRank&*/recvFrom,
                  /*pipeId*/pipeId, /*pipeId&*/…,
                  /*useGpsimdDma=*/false,                     // @0xf1d750  push 0   (*)
                  srcLoc);
    cc[+320] = -1;                                           // stream_id unset
    return cc;
}

GOTCHA — useGpsimdDma is hardcoded false at codegen. The disasm shows push 0 at 0xf1d750 immediately before the call addLocalSendRecv at 0xf1d768; that 0 is the 11th argument, the bool useGpsimdDma. The klr::SendRecv.useGpsimdDma field (present in to_string) is never read here. The GPSIMD-vs-DMA engine decision is deferred to the on-chip lowering pass (lowerSendRecv isCompatible + bytes-per-partition gate, §4), not made at codegen. [CONFIRMED via disasm — 0xf1d60e call getNeuronCoreId, 0xf1d750 push 0, 0xf1d768 call addLocalSendRecv are byte-witnessed in the body-frame sidecar; the klr-field-ignore is STRONG.]

addLocalSendRecv itself (libBIR, body @ 0x2b2780, decompiled cleanly there) is what actually mints the IT48 and stamps it as a local send-recv:

// bir::InstBuilder::addLocalSendRecv  @libBIR 0x2b2780  [CONFIRMED — decompiled in libBIR sidecar]
InstCollectiveCompute* addLocalSendRecv(name, srcML, dstML, inAP, outAP,
                                        sendToRank&, recvFromRank&, pipeId, pipeId&,
                                        useGpsimdDma, srcLoc) {
    auto v = addInstruction<InstCollectiveCompute>(bb, ip, name, inAP, outAP, …);   // IT48
    v[+248] = 0;                          // CollectiveKind = SendRecv(0)          (**)
    v[+432] = 1; /*tag +464 cleared*/     // is_local = TRUE  →  "LocalSendRecv"   (*)
    v[+472] = pipeId;                     // pipe_id
    v[+536] = 1; /*tag +568*/             // initial-corebarrier-class flag = true
    v[+504] = QuasiAffineExpr(module.name, 1 - recvFromRank);  // peer/source affine
    v[+576] = useGpsimdDma;               // = false here  → use_gpsimd_dma
    return v;                             // src/recv/send ranks woven into +480/+488
}

So an explicit klr::SendRecv always lowers to a CollectiveKind-0 (SendRecv), is_local=true IT48. There is no codegenSendRecvCCE: SendRecvCCE (CollectiveKind 1) originates upstream in the HLO collectives-to-CC path, never from a KLR leaf. klr::SendRecvCompute/klr::Send/klr::Recv have ser/des/to_string but no codegen leaf in the dispatch. [CONFIRMED — no codegenSendRecvCompute/codegenSendRecvCCE symbol exists in nm -DC.]

2. The master KLR↔BIR op-join table

This is the architecture-independent waist: every dispatched codegen<Op> leaf → its bir::Inst* / InstructionType. The Operator-kind column is the codegenOperator switch case (jpt @0x1de95f8, 7.21), not the KLR-ser wire type-id — the two are independent namespaces (e.g. NcMatMul: Operator-kind 22, ser type-id 183). The IT column uses 7.10 instruction-type ordinals. Brace {…} marks a 1→many fan-out (§3).

COMPUTE — matmul / activation / scalar / tensor

SELECT / DVE / IOTA / RNG

DATA-MOVE / collective / send-recv / rank / control

Count: ~57 dispatched leaves collapse onto ~50 distinct opcodes — the fan-ins of §3 (seven collectives → IT48, three activations → IT4, four scalar ops → IT29, two reduces → IT27) account for the gap. [Leaf addresses spot-verified byte-exact against nm -DC libwalrus.so; the bir::Inst* targets carried from the per-op Strand-I bodies.]

3. The five fan-outs (and the fan-ins)

A fan-out is a single klr leaf that emits one of several distinct opcodes based on a routing predicate read off the klr node. There are exactly five; each routing condition is transcribed from the codegen body.

(F1) NcCopy → {IT23 InstTensorCopy | IT24 …DynamicSrc | IT25 …DynamicDst} — codegenNcCopy @ 0xf2d120 reads isDynamicAccess(src)/isDynamicAccess(dst):

if (!dynSrc && !dynDst) emit InstTensorCopy;           // both STATIC (addTensorCopyCustom)
else if (dynSrc)        emit InstTensorCopyDynamicSrc;  // IT24
else /* dynDst */       emit InstTensorCopyDynamicDst;  // IT25
// guard: "Only one of src and dst can have dynamic access" (cpp:836)

NcCopy never emits GenericCopy (IT1) — that is an internal AbstractCopy node.

(F2) NcDmaCopy → {IT32 InstDMACopy | IT44 InstGenericIndirectSave} — codegenNcDmaCopy @ 0xf2fa40 falls through to codegenGenericIndirectSave @ 0xf2eb10 (which, confirming the fan-out, takes a klr::NcDmaCopy):

if (plainCopy || uniqueIndices) emit InstDMACopy;                 // IT32  (static=addUnary; dynamic=insertElement+assembleDynamicInfo)
else /* scatter w/ index vector */ codegenGenericIndirectSave();  // IT44  ("IndirectSaveAccumulate-" prefix; compute_op add → RMW field +320=4)
// IT44 then expands further at the H06 LowerGenericIndirect pass into a 5-op sequence.

[CONFIRMED — nm shows codegenGenericIndirectSave(std::shared_ptr<klr::NcDmaCopy>) @ 0xf2eb10, i.e. the same klr type as codegenNcDmaCopy.]

(F3) Transpose → {IT40 InstStreamTranspose | IT8 InstMatmult (vs identity)} — codegenTranspose @ 0xf28fb0 throws if engine==5 (PE), then forks on dst dtype:

if (dstDtype == fp32) emit InstStreamTranspose;        // IT40  DVE 32×32 tile
else { getIdentityMatrix(128×128 SBUF);                // matmul-against-identity
       addMatmult("matmult_TP_", isTranspose@+440=1); } // IT8
// (a THIRD transpose form is the DMA-descriptor one: klr::DmaTranspose → IT32.)

(F4) MatchReplace8 → {IT90 InstMatchReplace | IT91 InstMaxIndexAndMatchReplace} — codegenMatchReplace8 @ 0xf29e10 branches on the index-output Option<TensorRef> engaged byte at klr+0x50:

if (byte[klr+0x50] == 0) addMatchReplace();            // IT90  (index output ABSENT)
else { read 4th TensorRef; addInstruction<InstMaxIndexAndMatchReplace>(); }  // IT91 (sentinel→+0xF0)

(F5) CmpBranch → {IT79 InstUnconditionalBranch | IT78 InstCompareAndBranch} — codegenCmpBranch @ 0xf33a70 reads the BrCmpOp at klr+0x44:

if (byte[klr+0x44] == 1) emit InstUnconditionalBranch; // IT79  "always" arm (inlineNoReorderBlock shortcut)
else emit InstCompareAndBranch;                        // IT78  BranchCompareOp@+0xF0 = BrCmpOp − 2

Fan-ins (the inverse — many klr ops → one opcode). These are mode-discriminated by flag fields, not opcode choice:

• 7 collective arms + SendRecv → IT48 InstCollectiveCompute, per-arm CollectiveKind + routing fields (§1). This is a 7-arm fan-in, not a multi-opcode fan-out: every arm produces the same opcode.
• 3 activations → IT4 InstActivation (NcActivate/Activate share the leaf; Activate2 sets is_activate2=1).
• 4 scalar ops → IT29 InstTensorScalarPtr (TensorScalar; TensorScalarReduce acc@+0x1F0=3; NcScalarTensorTensor is_scalar_tensor_tensor@+0x1A0=1; TensorTensorScan is_tensor_tensor_scan@+0x1C8=1). TensorScalarCumulative → IT30 separately.
• 2 reduces → IT27 InstTensorReduce (TensorReduce; TensorPartitionReduce hard-codes axis=C=5).

4. Where the bound CollectiveKind decides on-chip vs remote

The four collective leaves all produce an abstract IT48; the consumer is the LowerLocalCollectives pass (real worker @ 0x1624850), which groups IT48 nodes and dispatches on CollectiveKind at CC+0xF8 (= our +248):

So the codegen-time choice of CollectiveKind directly selects the on-chip-vs-remote fate. Of the four leaves on this page, only SendRecv (kind 0) is on-chip-eligible; Permute(7), PermuteImplicit(9), PermuteReduceImplicit(10) are all ≥3 and stay as IT48 for the remote collective path (fence-only on-chip). And because codegenSendRecv hardcoded use_gpsimd_dma=false (§1.6), the actual GPSIMD-vs-DMA engine for a local SendRecv is decided here, by isCompatible plus the bytes-per-partition gate — not at codegen. [STRONG — LowerLocalCollectives dispatch keys on CC+0xF8 per the Strand-H worker; the kind-0/1/2-local vs ≥3-fence cut reconciles the earlier "{9,10} get local GPSIMD" seed reading, which is inverted.]

5. BIR-only opcodes — the beta2↔beta3 coverage asymmetry

Of the 110 InstructionType opcodes (7.10), the KLR (beta2) codegen surface emits roughly 50. The remaining ~60 are never produced by a codegenOperator leaf: they originate either (a) downstream in libwalrus lowering passes (Strand-H), (b) the Penguin / BirCodeGenLoop beta3 path, or (c) the HLO frontend (Strand-B). This is the beta2/KLR ↔ beta3/Penguin coverage asymmetry — the KLR surface is the leaf ISA a NKI kernel can name; the rest are structural/lowering opcodes minted by later passes. [STRONG]

BIR-only families, grouped by origin:

The clearest "beta3 / downstream only" families are the DMADescriptor* (IT68–72), DMATrigger (IT67), Loop/DMABlock (IT105–109), and the IndirectLoad/Save split (IT42–46).

GOTCHA — klr names that are not a bir opcode. Several roster names are not independently codegen'd: the abstract parents (MatMul/Copy/DmaCopy/ScalarTensorTensor/AffineSelect/RangeSelect/LocalGather) whose concrete Nc* child is the dispatched leaf; the sub-components (LoadStationary folded into the matmul stationary operand, MatchValueLoad inside the DVE Max/MatchReplace sequence, Send/Recv folded into the SendRecv path, SendRecvCompute originating upstream); and the structural ones (ExtendedInst → InstCustomOp via the NX-126 bridge; LncKernel is the per-LNC-core kernel-body driver codegenLncKernel @ 0xf34140, not an op — it emits many insts). The single cleanest "name ≠ opcode" remap is klr::Rng → InstMemset(mode=Random) (IT10), not InstRng (IT100) — so the bir Rng/Rand opcodes are BIR-only and the klr Rng op is klr-only-named. [STRONG]

6. The four-layer join — where this table sits

The klr→BIR map here is layer 1→2 of the full lowering. Each bir::Inst* continues through two more layers to reach NEFF wire bytes:

LAYER 1  klr::<Op>          → LAYER 2  bir::Inst* (IT)         [THIS PAGE, §2]
LAYER 2  bir::Inst* (IT)    → LAYER 3  CoreV{2,3,4}GenImpl::visitInst<X>   (per-ARCH ISA descriptor)
LAYER 3  ISA descriptor     → LAYER 4  InstGen.serialize → NEFF wire bytes  (enum marshalling per the wire crosswalk)

Layer 2→3 is where the per-arch fan-out happens: one IT can map to many ISA opcodes (gen2/3/4 differ). For example visitInstTensorReduce @ 0x12383a0 forks IT27 into 6 distinct ISA reduce opcodes by axis/transpose, and visitInstMaxIndex emits two co-issued DVE bundles from one IT89 — a 1→2 split that happens at the encoder, distinct from the codegen-time MatchReplace8 split (§3 F4). So §2 is the arch-independent waist; the ISA fan-out and the int↔wire↔JSON enum marshalling are downstream legs. See 7.9 op-family enums for the L1/L2/L3 enum crosswalk that the layer boundaries marshal across.

7. Diagnostic strings (collective / send-recv path)

All present in libwalrus.so .rodata. [CONFIRMED — strings -a hits for every row above.]

8. Function map

9. Adversarial self-verification ceiling

The five strongest claims, re-checked against the binary:

1. The three-way collective split — CONFIRMED. All eight leaf addresses are nm -DC rows; the two dedicated bodies (0xf22110, 0xf21850) do not call codegenCollectiveOp, while codegenCollectivePermuteImplicit @ 0xf23a40 does. Ceiling: high (addresses byte-exact; the "does not call shared encoder" is a body read, STRONG-to-CONFIRMED).
2. The five fan-outs — CONFIRMED for the existence of each branch. F2 is byte-anchored by codegenGenericIndirectSave(std::shared_ptr<klr::NcDmaCopy>) @ 0xf2eb10 taking the same klr type. The exact branch predicates (klr+0x50, klr+0x44, dtype fork) are carried from the per-op Strand-I bodies; offsets STRONG.
3. The BIR-only ~60 opcodes — STRONG. Grounded in the 110-opcode InstructionType census minus the ~50 leaf-reachable opcodes; per-family origin (H/N/Penguin/HLO) is inference from naming + pass roster, not a per-opcode trace. Tagged STRONG, not CONFIRMED.
4. ~57 leaves → ~50 opcodes — STRONG. The "~57" is the dispatched-leaf count and "~50" the distinct-opcode count after the §3 fan-ins; both are approximate by construction (the abstract-parent / sub-component boundary is a judgment call). The fan-in arithmetic is exact for the four families named.
5. CollectiveKind multiplexing (one opcode IT48 for all kinds) — CONFIRMED. InstCollectiveCompute is IT48; the kind is the orthogonal +0xF8 sub-field; sameInst @ 0x2f4a80 keys on it (7.9, independently verified). codegenSendRecv's hardcoded false and the getNeuronCoreId/addLocalSendRecv call sequence are byte-witnessed in the body-frame disasm (0xf1d60e, 0xf1d750, 0xf1d768).

Items tagged INFERRED and not raised above: the exact JSON key-name of the +280 channel_id field and the +504 peer-affine field; the op-family-enums cross-reference resolves the kind ordering but not these two sub-field names.

Cross-references

• 7.21 KlirToBirCodegen dispatch — the 77-case switch, refcount shims, the dispatch chain feeding every leaf here.
• 6.5.13 BirCodeGenLoop collective lowering — the beta3 collective counterpart (same IT48 model, Penguin input).
• 7.9 Op-family enums — CollectiveKind/CollectiveDimension/CollectiveComputeTypeHint; the "one opcode multiplexes 11 kinds" quirk; the L1/L2/L3 crosswalk.
• 7.10 InstructionType — the 110 L1 opcodes referenced by the IT column.
• Part 13 (distribution layer — planned) — the SPMD / replica-group / ICI cross-node path that consumes the fence-only (kind ≥3) collectives.