Keyboard shortcuts

Press or to navigate between chapters

Press S or / to search in the book

Press ? to show this help

Press Esc to hide this help

The Pass Catalog

All counts, names, and symbols on this page apply to neuronx_cc 2.24.5133.0+58f8de22 (cp310 wheel; cp311/cp312 carry the same registries — only Cython hash suffixes and binary VAs drift). The four stages live in four binaries: neuronxcc/starfish/bin/hlo-opt (BuildID 93dd8bd9bd4c697b, not stripped), the MLIR pipeline inside that same hlo-opt, the Cython pass modules under neuronxcc/starfish/penguin/, and neuronxcc/starfish/lib/libwalrus.so (BuildID 92b4d331…). Treat every count as version-pinned — the re-derivation recipes in §5 re-run against any wheel.

Abstract

This appendix is the single consolidated index of every compiler pass in neuronx-cc, across all four optimization stages that a graph traverses on its way from XLA HLO to a packaged NEFF. There is no one place in the binaries where all four registries meet — each stage owns its own name→pass mechanism, in its own binary, with its own selection rule — so this page is the meeting point. It exists to answer one question for a reimplementer or a debugger: given a pass name, which stage owns it, what class does it construct, where in the pipeline does it run, and which wiki page documents it.

The four stages and their headline counts, each re-derived from the binary in §5:

  1. hlo-opt B-strand — the --passes registry of HLO-level rewrites, 112 registered names → 112 classes, keyed by pass->name(). The source of truth is 4.1.
  2. hlo2penguin C-strand (MLIR) — the imperatively-built mhlo/stablehlo pass pipeline that lowers optimized HLO to penguin.py text. 32 MHLO steps; the StableHLO fork is the same spine with 5 insertions = 37 steps. The source of truth is 4.32.
  3. Penguin middle-end — a catalog of Cython pass modules (no flat registry), sequenced at run time by SharedCodeGenFlow generators. 248 total pass modules across four directories (71 + 131 + 6 + 40), of which ~8–10 in the two transforms/ dirs are helper libraries, not Pass classes. The source of truth is 5.7.
  4. walrus backend — the GeneratorRegistration name→factory registry, 150 registered names collapsing to 121 distinct C++ classes (10 classes multi-registered), scheduled into ~180 pipeline steps. The source of truth is 8.1.

Each stage uses a structurally different name→pass mechanism, and that difference is itself the most useful thing this catalog records — a reimplementer who assumes one registry for the whole compiler will get all four wrong. §0 lays the four mechanisms side by side. §1§4 index each stage's passes by family, with the documenting page per family. The full row-level tables are not reproduced here — they live in the four normative pages — because pasting 112 + 37 + ~240 + 150 rows would be a dump, not an index; this page is the map to those tables, with the cross-stage counts reconciled in one place.

This is a reference catalog, not an algorithm walkthrough; its value is completeness and correct cross-stage ordering. Confidence is per-row, and reconciled in §5: the four name counts (112 / 32+5 / 248 / 150) come off the binary; the 121-class walrus total is derived from them; the Penguin run order is [INFERRED] from the version-counter source-order trace rather than read off a table.

Stage 1 — hlo-optxla::hilo::RegisterHiloHloPasses @ 0x1e72270llvm::StringMap; 112 names (4.1)
Stage 2 — MLIRhilo::registerMHLOPasses @ 0x1ee12f0 (32) / registerStableHLOPasses @ 0x1ee2120 (37) (4.32)
Stage 3 — PenguinSharedCodeGenFlow generators → PassConstructorBuilder; 248 modules (5.7)
Stage 4 — walrusGeneratorRegistration::registerGenerator_Hashtable; 150 names → 121 classes (8.1)
Grand pass-name total112 + 37 + ~240 + 150 ≈ ~539 named pass positions across the four stages
SelectionStage 1: --passes/--skip-pass (XLA DebugOptions) · Stage 2: per-pass CompileConfig gate byte · Stage 3: OptLevel + required/optional · Stage 4: --pass (registry) + optlevel builder

0. The Four Registry Shapes

The single most important cross-stage fact is that each stage binds a pass name to a pass object by a different mechanism. They are not four instances of one design; they are four designs. A reimplementer must reproduce each one separately.

StageBinaryName→pass mechanismKey derivationSelection / ordering
1 hlo-opthlo-opt (C++)llvm::StringMap<std::function<unique_ptr<HloPass>()>> filled by 112 RegisterHloPass callskey = pass->name() read back at registration (vtable vptr+0x10) — not a literal in the registrarXLA DebugOptions whitelist/blacklist; run-list ordered by the driver outside the registrar
2 MLIRhlo-opt (C++)no registry — imperative pm.addPass(...) / addNestedPass<func::FuncOp>(...) spinenone; passes are constructed inline by create<Name>Pass() factoriesordering is the literal addPass source order; each gated on a CompileConfig byte
3 PenguinCython .so per passno registryyield PassType(Class, kwargs, required|optional) into a PassConstructorBuilder folded via +=none; the PassType carries the class object directlyOptLevel staged-skip at flow level + per-spec required/optional + clOptBool toggles
4 walruslibwalrus.so (C++)_Hashtable<string, function<unique_ptr<BackendPass>(PassOptions)>> filled by 150 register_generator_*__ static initializerskey = the literal string passed to registerGenerator("<name>", …)--pass <name> via getGenerator; optlevel pipeline builders call addPass("<name>"); granularity per call-site

Two stages have a real string-keyed registry (1, 4); two build the pipeline imperatively with no name table at all (2, 3). The two registries even differ in where the key comes from: hlo-opt reads the key back off name() at registration (so the registrar source never shows it), while walrus passes the literal name straight into registerGenerator. That asymmetry is why recovering the 112 hlo-opt keys required walking each factory to its vtable, whereas the 150 walrus names grep straight out of the symbol table.

GOTCHA — "pass count" means four different things. Stage 1's 112 is registered names = classes (1:1). Stage 4's 150 is registered names, which is not the class count (121) — 10 classes answer to several names. Stage 3's 248 is modules, which is not the pass-class count (~8–10 are helper libs) and not the run-list length (a flow runs a subset, some passes twice). Stage 2's 32/37 is pipeline steps, every one gated, so the executed count depends on the CompileConfig. A single "total passes" number for the compiler is meaningless; always say which of name/class/module/step you mean.


1. Stage 1 — hlo-opt B-strand (112 passes)

xla::hilo::RegisterHiloHloPasses registers 112 passes into an llvm::StringMap keyed by each pass's name(). An objdump of the registrar body (0x1e722700x1e744af) shows exactly 112 call instructions targeting RegisterHloPass (0x1ebc3f0). Of the 112, 87 override their own Run (vtable+0x18) and 25 are OpExpander passes sharing xla::OpExpanderPass::Run @ 0x29f0bb0 and overriding InstructionMatchesPattern (vtable+0x28). Of provenance, 3 are stock XLA classes reused unchanged (tuple-simplifier, call-inliner, dce); the rest are Neuron-authored.

The 112-row name→class→vtable→entry table is not duplicated here — it is the canonical table in 4.1 §2. Below is the index of that table by family, pointing each family at the 4.x page that documents it.

Family (registration-order span)Representative namesDocumenting pageConf
Stock XLA (reused)tuple-simplifier call-inliner dce4.1 §4CERTAIN
Legalize / intrinsics / comparelegalize-intrinsics legalize-compareintrinsics-compare-legalizeCERTAIN
Softmax / TopK / ArgMax legalizelegalize-softmax legalize-topk legalize-aws-neuron-arg-max lower-argminmax-custom-callsoftmax-legalize, topk-legalize, argmax-argmin-legalizeCERTAIN
Quantize / scaled-matmullegalize-quantize-mx legalize-scaled-matmulmx-fp8-legalization, int8-quantize-legalizationCERTAIN
Collectives → custom-callconvert-collectives-to-custom-call convert-fs-patterns-to-cccollectives-to-customcallCERTAIN
CC-op decompose / legalizedecompose-cc-ops legalize-cpu-cc-ops legalize-ccops-for-tensorizer trivial-cc-removalccops-decompose-legalizeCERTAIN
Collective combinersall-gather-combiner reduce-scatter-combiner all-reduce-combinercollective-combinersCERTAIN
Stream-ID / channel-IDcollective-stream-id-checker aws_neuron_collective_stream_id_injector neuron_unique_channel_id_enforcercollective-stream-channel-idCERTAIN
Flip-collective OpExpandersaws_neuron_flip_all_gather_* aws_neuron_flip_reduce_*flip-collective-opexpanderCERTAIN
AllReduce→ReduceScatter / DynSliceaws_neuron_rewrite_all_reduce_dynamic_slice[_multiple_groups]allreduce-dynslice-rewritesCERTAIN
CollectivePermute → AllGathercollective-permute-to-all-gather aws_neuron_rewrite_collective_permutecollectivepermute-to-allgatherCERTAIN
Looped-einsumneuron_looped_einsum_replacer neuron_looped_einsum_token_replacerlooped-einsum-collective-matmulCERTAIN
Int / scalar-reduce decomposeaws_neuron_decompose_int_all_reduce aws_neuron_decompose_scalar_reduceint-scalar-reduce-decompositionCERTAIN
While-loop unroll / trip-countwhile_loop_unroller unroll-while-loop aws_neuron_rewrite_all_gather_trip_countwhileloop-unroll-tripcountCERTAIN
While-loop collective code-motionneuron-while-loop-all-reduce-code-motion neuron_move_*_while_loopwhileloop-collective-codemotionCERTAIN
Boundary markers / layer-cutcanonicalize-boundary-marker boundary-marker-removalboundary-markers-layer-cutCERTAIN
Control-dep reificationpreserve-control-depscontrol-dep-reificationCERTAIN
Layout passesio-layout-normalization aws_neuron_ensure_descending_layout_in_rootlayout-passesCERTAIN
I/O-alias familyremove-aliases add-must-aliases add-may-aliases flip-must-aliases flip-may-aliases aws_neuron_alias_to_must_alias aws_neuron_buffer_donation_to_aliasio-alias-familyCERTAIN
Concat optimizationssimplify-concat slice-of-concat-optimizer neuron_repeated_dus_to_concatconcat-optimizationsCERTAIN
DUS/DS simplifier / moverneuron-dus-ds-index-simplifier dynamic-slice-mover dynamic-slice-transposedus-ds-simplifierCERTAIN
Dedup (param / kernel)neuron-preprocess-kernel-duplicate-remover neuron_all_gather_duplicate_removerdedup-passesCERTAIN
InstCombine peepholeneuron-hlo-inst-comb aws_neuron_common_instruction_elimination eliminate-redundant-compare aws_neuron_resolve_self_comparisoninstcombine-peepholeCERTAIN
Precision / upcastupcast-all-to-fp32 batch-norm-training-upcast rewrite-module-dtype neuron-int-matmul-downcastprecision-upcast-passesCERTAIN
Calibration / scale(scale/zero-point flow)calibration-scale-flowHIGH
NKI / native loweringlower-to-nki-kernels lower-to-custom-native-kernelhlo-to-native-kernel-loweringCERTAIN
Misc / cleanup sweepstub metadata-naming aws_neuron_default_metadata opt-barrier-removal inline-weights replace-rng inject-prints inject-numerical-errorshlo-misc-cleanup-sweepCERTAIN

NOTE — the registration order (1..112) in 4.1 §2 is the registrar's emission order and serves as the default-pipeline proxy — it is not the run order. The executed run-list is a driver-supplied ordered name subset assembled outside hlo-opt's registrar, then filtered by --passes (whitelist) / --skip-pass (blacklist). The family table above is sorted by topic, not by registration index; cross-reference 4.1 §2 for the exact #.


2. Stage 2 — hlo2penguin C-strand MLIR pipeline (32 + 5 = 37 steps)

After the 112 B-strand passes run, optimized HLO is imported into mhlo and a fixed, imperatively-built MLIR pipeline rewrites it down to penguin.py text. There is no registry: hilo::registerMHLOPasses @ 0x1ee12f0 is a flat pm.addPass(...) spine, and the linear source order is the pipeline order. The StableHLO fork (hilo::registerStableHLOPasses @ 0x1ee2120, selected by CompileConfig byte [+0x1CA]≠0) is the same spine with 5 extra insertions — 32 → 37 steps. Every step except the terminal emitter is gated on a CompileConfig byte (default-off unless a cl::opt flips it; the lone default-on is step 6 LowerComplex, set by a static initializer).

The full ordered MHLO spine is the canonical listing in 4.32. Index of the 32 steps by family:

#MHLO step (create…Pass)Family / documenting pageGateConf
1Canonicalizerupstream MLIR canonicalizecfg[0x119]CERTAIN
2NeuronControlDepTupleSimplifiercontroldep-tuple-flatten-mlircfg[0x200]CERTAIN
3StableHLOCollectiveBroadcastToAllGathercollective bcast→all_gatheruncondCERTAIN
4VerifySupportedOpsop-legality verifiercfg[0x136]CERTAIN
5CanonicalizeConvconv-canonicalizationcfg[0x12F]CERTAIN
6–7LowerComplex / LowerComplexExtracomplex→realcfg[0x11A](on) / [0x11B]CERTAIN
8RemoveOptimizationBarriersstrip optimization_barriercfg[0x11C]CERTAIN
9–10IdentifyCrossPassTensors / Rematerializecross-pass tag + recomputecfg[0x116]CERTAIN
11CoalesceCollectiveOpsMLIR collective coalescingcfg[0x117]CERTAIN
12ConvertFSPatternToCustomCallsFS subgraph → AwsFSPattern5cfg[0x11F]CERTAIN
13RemoveDeadValuesupstream MLIR DCEuncondCERTAIN
14NeuronInstCombineneuron-instcombine-mliruncondCERTAIN
15NeuronOpFusion (6 sub-passes)op-fusion-dot-elementwise, rmsnorm-fusion-cluster-codegencfg[0x120]CERTAIN
16FoldIotapow(iota)→constcfg[0x118]CERTAIN
17–18SchedulePostorder / AnalyzeScheduleschedule-fusion-compositecfg[0x130]/[0x134]CERTAIN
19ReplaceTokenTypeWithU8!mhlo.tokentensor<ui8>uncondCERTAIN
20ScatterMotionscatter code-motionuncondCERTAIN
21ScheduleFusionschedule-fusion-compositecfg[0x135]CERTAIN
22–23HoistCompute / MemcastMotionhoist / late-cast motioncfg[0xE9]/[0xE5]CERTAIN
24PruneFunctionsdead-function elimuncondCERTAIN
25SixtyFourHack (prepended before 26)i64→i32 emulationqword_9C71338==0HIGH
26PenguinizeIOstablehlo-printer-penguinizeuncondCERTAIN
27CanonicalizeForTensorizer (8 sub-rewriters)canonicalize-for-tensorizeruncondCERTAIN
28TensorizerLegalizationtensorizer-legalizationuncondCERTAIN
29–30LegalizeAlias / VerifyAliasinginsert AliasingCopy + RAW checkuncond / cfg[0x13C]CERTAIN
31MLIRInstructionHistogramop-count diagnosticsuncondCERTAIN
32MhloToPyPenguin (terminal)mhlo-to-python-printer-driver, -heavyuncondCERTAIN

The 5 StableHLO-only insertions (32 → 37)

The StableHLO spine twins every MHLO step and adds 5 of its own (4.32):

After stepStableHLO-only insertionRoleConf
1HloLegalizeToStablehlomhlostablehlo bridgeCERTAIN
1ConvertCustomCallToAllReducecustom_call("mhlo.all_reduce")all_reduceCERTAIN
14FusionToCompositemhlo.fusionstablehlo.compositeCERTAIN
26LegalizeSRAshift_right_arithmetic signednessCERTAIN
26LegalizeScatterstablehlo.scatter → F32 + canonical bodyCERTAIN

QUIRK — step 25 is prepended, not appended. SixtyFourHack's gate (jz→altblk @ 0x1ee19a0, on global qword_9C71338==0) inserts it before step 26 PenguinizeIO, so it runs ahead of IO packaging despite appearing 25th in the source. A reimplementer who addPass-es it in textual position sequences it wrong. The ordering is read from the disassembled control flow only — there is no registration table that states it.


3. Stage 3 — Penguin middle-end (248 modules, run-list ~varies)

The Penguin middle-end is not a registry and not a flat pipeline — it is a catalog of Cython pass modules sequenced at run time by generator functions. Four directories partition the catalog; counts are ls-verified (excluding __init__):

DirectoryModulesScopeDocumenting page
penguin/transforms/71target-independent IR passes (DCE, VN, LICM, simplify, delinearize, predicate)5.7
penguin/targets/transforms/131target-oriented, shared Sunda+Tonga (alloc/spill, layout, tiling, collective, DMA, type-legalize, Tritium fusion)5.7
penguin/targets/sunda/passes/6Sunda-specific ISel / tiling5.7
penguin/targets/tonga/passes/40Tonga-specific ISel / tiling / scheduler5.7
Total248(targets/cayman/passes/, targets/core_v4/passes/ ship only __init__ — reuse Sunda/Tonga)

The targets/transforms/autotune/ subpackage (6 modules: Autotuner _Compiler _PerformanceMetric _Search _TreeSearch _TreeSearchVisualizer) is counted separately, not in the 131.

GOTCHA — module count ≠ pass count. ~8–10 modules in the two transforms/ dirs export no Pass/DotTransform subclass — they are helper libraries (DataflowUtil InstTransformUtils IPUtils LoopTransformUtils TensorContractUtils TensorOpUtils InstBuilder IntrinsicBuilder DataflowUtils). The real pass-class set is the subset whose top-level class subclasses DotTransform or BaseAnalysis. The 248 is the verifiable module anchor; the pass-class subset is a few smaller — [INFERRED], since the by-class filtering was spot-checked rather than exhaustively enumerated.

Run order — the driver, not the directory

There is no static order in the directory listing. The order is set by the four SharedCodeGenFlow generators (peephole_optimizations, legalize_pe_insts, nki_codegen, codegen_optimization_post_tritium_fusion), each yielding PassType specs that a PassConstructorBuilder folds. The canonical post-fusion list (with its required/optional qualifiers and the spill-vectorize-respill double-FastSpillGeneration) is in 5.7. Index of the high-population families:

FamilyRepresentative passesDocumenting pageConf
Alloc / spillAllocateBlocks MaxLiveSpiller FastSpillGeneration StackAllocator SpillPSumscheduling-minreg, tensor-buffer-nodeCERTAIN
LayoutLayoutTilingPipeline PAGLayout* LayoutPreprocessing NKIKernelLayoutlayout-tiling-pipeline, layout-middle-endCERTAIN
Tiling / vectorizeFlattenAxesForTiling VectorizeLoop SFKVectorizer LoopSplitting PG*Analysisloop-transform-clientsCERTAIN
Loop transformsLICM LoopFusion PerfectLoopNest FlattenLoop Delinearizationloop-transform-clientsCERTAIN
ISL / affineIslSimplifier-family, ModDivDelinear, AffinePredicateResolutionisl-simplifier, isl-codegen, affine-expr-algebraHIGH
Collective (SPMD/CC)CCOpFusion CoalesceCCOp TileCCOps SPMDCodeGen SimpleAllReduceTilingdata-movement-fusionCERTAIN
DMADMALegalizer DMALocalityOpt DataStreaming DramToDramTransposedge-level-dynamic-dma, tensorcopy-dynamic-generatorsCERTAIN
Software pipeliningSoftwarePipelineCodeGensoftware-pipeliningCERTAIN
Tritium fusion (+autotune)TritiumFusion TritiumFusionBase, autotune/Autotuner5.7 autotunerCERTAIN
ISel (per-target)SundaISel (sunda), TongaISel (tonga)5.7CERTAIN
Codegen / emitNkiCodegenPass, BirCodeGenLoop, targets/codegen/*ir-mlir-bir-mappingCERTAIN

QUIRK — a pass can appear twice in the run-list. FastSpillGeneration runs before and after SFKVectorizer (vectorization creates new live ranges to spill); nki_codegen is chained into the post-fusion list. A set-based reimplementation that de-dups specs by class silently drops the second spill. The Penguin run-list length is therefore not a function of the module count — it is the generator's yield sequence, with repeats. Membership is exact; the order is [INFERRED] from the version-counter source-order trace.


4. Stage 4 — walrus backend (150 names → 121 classes)

The walrus backend is a GeneratorRegistration registry: 150 register_generator_<name>__ static initializers each install a name → factory entry, but they construct only 121 distinct C++ pass classes — 10 classes are registered under multiple names, parameterized at construction by a memory space, optimization level, or scheduling phase. The pipeline schedules these into ~180 steps (some names at several positions, plus interposed !isRealPass check/dump/sim passes).

  • 150, straight off the binary: nm -DC libwalrus.so | rg -o 'register_generator_[a-z0-9_]+' | sort -u | wc -l = 150.
  • 121, derived: the count of distinct _M_invoke ctor targets over the 150 bodies. The 10 multi-registered groups contribute 29 surplus names (9+7+3+2+2+2+1+1+1+1); 121 + 29 = 150.

The full 150-name → class → granularity → evidence table is the canonical listing in 8.1 §"The full 150-name table". The 150→121 collapse must be represented, not flattened — here are the 10 multi-registered classes that produce the gap:

ClassNamesSurplusDiscriminatorConf
MemoryAnalysis109which prior pass the snapshot follows (post-rotation/coloring/dma/sched/unroll)CERTAIN
ColoringAllocator87Type ∈ {SB,PSUM,REG,DRAM,DRAM-shared} + post_lnk/debugCERTAIN
AddressRotation43space dram/psum/sb + post_scheduleCERTAIN
DMAOptimization32psum/sb/input-coalescingCERTAIN
FullUnroll32unrollLevel ∈ {all,some} + memloc-generationCERTAIN
SeparateLoadAndCompute32base / with_memset / post_adaCERTAIN
AntiDependencyAnalyzer21base / post_shared_dramCERTAIN
DeadCodeElim21optlevel o0 / o1CERTAIN
PerfSimPass21perf_sim / perf_sim_at_endCERTAIN
PrefetchScheudling (sic — misspelled in binary)21before_sched / after_schedCERTAIN

The other 111 names are 1:1 with their class. Index of the 150 names by family, pointing each at the documenting 8.x page:

FamilyRepresentative namesDocumenting pageConf
Pass hierarchy / registry(the BackendPass base, forks, GeneratorRegistration)8.1CERTAIN
Optlevel pipeline(the optlevel builders that call addPass("<name>"))pass-pipeline-optlevelsCERTAIN
NKI-AST→BIR / unrolltranslate_nki_ast_to_bir unroll full_unroll_* heuristic_unroll flatten_small_loopstranslate-nki-unrollCERTAIN
DMA legalizationlegalize_strided_dma legalize_cce_dma lower_generic_indirect lower_dynamic_dmadma-legalizationCERTAIN
DMA materialization / queuesdynamic_dma_setup dynamic_dma_scan alloc_queues insert_dma_switch_queue_instance optimize_queue_switchdma-materialization, dma-queuesCERTAIN
DMA engine bindingassign_hwdge_engine assign_trigger_engine infer_stream_ids label_dma_qosdma-engine-bindingCERTAIN
DMA metrics / profilerdma_metrics dma_optimization_psum dma_optimization_sb input_dma_coalescingdmametrics-profiler, alt-allocators-dma-optCERTAIN
lower_select / control / branchlower_select lower_control lower_branchlower-select-control-branchCERTAIN
Engine-lowering setlower_sync lower_act lower_dve lower_ap lower_ac lower_dma lower_symbolic_inst lower_local_collectivesengine-lowering-setCERTAIN
Coloring / DRAM / PSUM alloccoloring_allocator_* (8) coloring_allocator_with_loop linear_scan_allocator psum_legalizationallocator-drivers, dram-allocator, psum-allocatorCERTAIN
Address rotationaddress_rotation_{dram,psum,sb} address_rotation_psum_post_scheduleaddress-rotationCERTAIN
Memory analysis (10 snapshots)memory_analysis_after_*allocator-driversCERTAIN
Schedulingpre_sched post_sched prefetch_scheduling_{before,after}_sched instruction_reorderpost-sched-schedulersCERTAIN
Dependence graph / dep-optbuild_fdeps dep_opt dep_reduction order_constraints anti_dependency_analyzer*dependence-graph, dep-opt-reductionCERTAIN
Peephole / constprop / rematpeephole_opts early_peephole_opts constant_propagate remat_optimization remove_redundancies seq_inst_optpeephole-constprop-rematCERTAIN
Matmul ordering / acc-groupsorder_column_tiled_mms legalize_mm_accumulation_groupsmatmul-ordering-accgroupsCERTAIN
Loop-opt (backend)loop_optimizationloopopt-licm-fusion, loopopt-transformsCERTAIN
Value numbering / mem2regvalue_numbering vn_splitter mem2reg non_ssa_legalizationlegality-dispatchHIGH
LNC splitter / barrierlnc_splitter lnc_barriercheck lnc_verifier shared_mem_cb_insertionlnc-splitter, barriercheckCERTAIN
Local collectiveslower_local_collectives coalesce_multichannel_cc_ops sync_before_global_cclocal-collectivesCERTAIN
BIR linker / kernel inlinebir_linker inline_bir_kernel inline_nki_kernel memreserve runtime_memory_reservationbir-linker, memreserve-kernel-inlineCERTAIN
Codegen / bin emissioncodegen neff_packagercodegen-driver, bin-emissionCERTAIN
Verifiers / serdesbirverifier bir_racecheck verify_bir_serdes oom_checkerbirverifier-per-opCERTAIN
Perf-sim / metrics (!isRealPass)perf_sim perf_sim_at_end perf_sim_package_pass report_stats dumper hasherperfsim-cost-model, perf-sim-wiring, metricstore, parserdumperCERTAIN

GOTCHA — serial* is not "always serial". In 8.1's full table the granularity column shows a Fork tag only where an explicit add*ParallelPass<Class> template instance exists; serial* means "added by name (default serial child)" — but the same class may be fork-wrapped at a different call-site, because granularity is a property of the call-site, not the class. Do not read serial* as a guarantee.


5. How the counts are grounded

Each cross-stage count this catalog turns on, and what pins it in the cp310 binaries:

  1. Stage 1 = 112. objdump -d --start-address=0x1e72270 --stop-address=0x1e744b0 hlo-opt | rg -c 'call.*1ebc3f0'112 calls to RegisterHloPass. One name per call, key = pass->name().
  2. Stage 2 = 32 MHLO + 5 = 37 StableHLO. registerMHLOPasses @ 0x1ee12f0 (3096 B) is a flat 32-step addPass spine; registerStableHLOPasses @ 0x1ee2120 (3142 B) is the twin plus 5 insertions. The five inserts (HloLegalizeToStablehlo, ConvertCustomCallToAllReduce, FusionToComposite, LegalizeSRA, LegalizeScatter) are enumerated in 4.32. The step count is exact; the SixtyFourHack prepend at step 25 rests on disassembled control flow alone.
  3. Stage 3 = 248 modules (71/131/6/40). ls <dir>/*.so | grep -v __init__ | wc -l in each of the four directories → 71, 131, 6, 40, summing to 248. targets/transforms/autotune/ accounts for the 6. targets/cayman/passes/ and targets/core_v4/passes/ contain only __init__. That is the module count; the pass-class subset is slightly smaller because of helper libs, flagged in §3.
  4. Stage 4 = 150 names → 121 classes. nm -DC libwalrus.so | rg -o 'register_generator_[a-z0-9_]+' | sort -u | wc -l150. The 10 multi-registered groups (MemoryAnalysis=10, ColoringAllocator=8, AddressRotation=4, DMAOptimization/FullUnroll/SeparateLoadAndCompute=3 each, four more ×2) contribute 29 surplus names, and 121 + 29 = 150. The 150 is grep-able; the 121 is a distinct-ctor-target count derived over the 8.1 table.
  5. The four name→pass mechanisms really are different. Stage 1 = llvm::StringMap keyed by name() read-back; Stage 2 = imperative addPass with no name table; Stage 3 = PassType specs folded by PassConstructorBuilder, again no name table; Stage 4 = _Hashtable keyed by the literal registerGenerator string. Each is symbol-anchored in its normative page (4.1, 4.32, 5.7, 8.1).

No pass name on this page is fabricated: every name is the verbatim string carried by its normative source page, which in turn read it from the binary — a name() literal, a create…Pass factory symbol, a __pyx_n_s_ interned token, or a register_generator_*__ symbol. Where this catalog points a family at a documenting wiki page, the slug is present on disk; families with no dedicated detail page point at their stage's normative registry page. The libBIR and libwalrus binaries and their IDA databases are present in-corpus, so the backend class addresses cited via 8.1 are read, not inferred.

NOTE — what this page is not. It is an index, not a replacement for the four normative tables. The per-row name→class→vtable→entry detail for Stage 1 lives in 4.1 §2; the per-step MLIR spine in 4.32; the Penguin driver order in 5.7; and the full 150-name → class → granularity table in 8.1. This catalog reconciles their counts and gives one map across all four.


Stage pageRelationship
4.1 — hlo-opt Pass RegistryThe 112-name --passes registry; source of §1
4.32 — hlo2penguin MLIR PipelineThe 32+5-step MLIR spine; source of §2
5.7 — Penguin Pass Roster & Pipeline DriverThe 248-module catalog + SharedCodeGenFlow order; source of §3
8.1 — BackendPass Hierarchy & 150→121 RegistryThe 150-name walrus registry; source of §4

Cross-References