The BackendPass Hierarchy and the 150-Name → 121-Class Registry

All symbols and addresses on this page apply to neuronx_cc 2.24.5133.0+58f8de22 (cp310; cp310/11/12 are byte-identical). Everything here lives in neuronxcc/starfish/lib/libwalrus.so (BuildID 92b4d331…). For .text and .rodata the virtual address equals the file offset. Other wheels differ — treat every address as version-pinned.

Abstract

The Walrus backend is a pass pipeline, and this page documents its spine: the abstract neuronxcc::backend::BackendPass base, the three run entry overloads it exposes, the fork/join granularity classes layered above it, and — the headline — the pass registry that turns a pass name into a constructed pass object. The registry is the seam the walrus_driver CLI (3.7) drives: every --pass <name> resolves through it.

The single number that organises the whole backend is the collapse 150 → 121. There are exactly 150 register_generator_<name>__ free functions — one per legal pass name — but they construct only 121 distinct C++ pass classes. Ten classes are registered under more than one name, parameterized at construction time by a memory space, an optimization level, or a scheduling phase, so the same class serves several pipeline roles under several names. Above that, the pipeline quotes roughly 180 steps, because some names appear at several pipeline positions and because non-real check/dump/sim passes are interposed between named steps. This page recovers all three counts and shows exactly where each collapse arises.

The bar for reimplementation is that a reader can reconstruct: the BackendPass interface and which of its methods is the driver entry; how GeneratorRegistration populates a name → factory hashtable at library load via 150 static initializers; how a factory lambda's _M_invoke body is the ground-truth evidence for the name → class mapping; the three fork granularities and how they are assigned per call-site rather than per class; and the full 150-name table with its 10 multi-registered classes.

The BackendPass base interface

Purpose

BackendPass is the abstract base every backend pass derives from. It is a template-method base in the classic sense: it owns one non-virtual entry method that does the bookkeeping common to all passes, and dispatches the actual per-unit work to a virtual hook the derived class overrides. A reimplementer needs three things from it — the entry overload the pipeline calls, the virtual hook it dispatches to, and the two predicate methods (isRealPass, dry_run) that the verifier/measurement front uses.

The three run overloads

The base exports three methods (nm -DC libwalrus.so), all CONFIRMED by symbol:

BackendPass::run(vector<unique_ptr<bir::Module>>&)         T  0x1736d40   ── DRIVER ENTRY
BackendPass::dry_run(vector<unique_ptr<bir::Module>>&, logging::Logger&) T 0x173af40   ── DRY-RUN
BackendPass::isRealPass() const                            W  0x87fd70    ── weak inline predicate

The vector overload at 0x1736d40 is the entry point the pipeline invokes for every pass. It is the first of the three run shapes a reader meets:

1. run(vector<unique_ptr<bir::Module>>&) — the driver entry. The vector holds one bir::Module per LNC core (the N per-core modules a Logical NeuronCore split produces). This overload is non-virtual on the base; it threads the module name into a Boost.Log scoped attribute, then dispatches one virtual hook through the object's vtable.
2. run(bir::Module&) — the per-module work hook the leaf passes override. Most transform passes implement only this (e.g. Unroll::run(bir::Module&) @ 0xb42f30). It is called once per module element of the vector.
3. runSubgraph(vector<unique_ptr<bir::Module>>&) — the per-subgraph hook, overridden by the passes that fork over subgraphs (e.g. BirSim::runSubgraph @ 0x1109440). SubgraphForkPass calls this hook on each subgraph rather than the per-module one.

The dispatch is visible in the disassembly of the vector overload — the entry loads the object's vptr and calls the slot at vt+0x10:

// neuronxcc::backend::BackendPass::run(vector<unique_ptr<bir::Module>>&)  @0x1736d40  [CONFIRMED]
//   %r15 = `this` (the pass);  %rdx = &modules vector
//   1736d70:  mov    (%r15),%rax        // %rax = vptr (load the pass's vtable)
//   1736d73:  mov    (%rbx),%rdx        // %rdx = &modules[0] (first bir::Module)
//   1736d7c:  call   *0x10(%rax)        // INDIRECT call vtable slot at +0x10 == the per-unit hook
unique_ptr<BackendPass> /*ret in *%r13*/
BackendPass_run(BackendPass *this, vector<unique_ptr<bir::Module>> *modules) {
    // ... push module name into a boost.log scoped attribute ...
    (*(void(**)(BackendPass*, bir::Module*))(*(void***)this + 2))(this, modules->begin());
    // the +0x10 slot resolves, per dynamic type, to the derived
    //   run(bir::Module&)  or  runSubgraph(...)  override
}

The call *0x10(%rax) is the entire dispatch mechanism: the base's run is fixed code, the behaviour is whatever the derived class installs at vtable slot +0x10. (Slot index = (0x10 − 0x10)/8 = 0, measured from the loaded vptr — i.e. the first non-RTTI slot.) [CONFIRMED — mov (%r15),%rax ; mov (%rbx),%rdx ; call *0x10(%rax) read verbatim from the binary at 0x1736d70–0x1736d7c.]

The two predicates

• isRealPass() @ 0x87fd70 (weak, inlined) distinguishes real transform passes from the synthetic check/dump/measurement passes (Dumper, Hasher, perf_sim*, the verifiers, report_stats) that the backend's CheckPassAdder interposes between named passes. ContainerPass overrides it at 0x1743320. A false return is the marker that a pass is structural padding in the step count, not a pipeline transform. [CONFIRMED — symbol + ContainerPass::isRealPass override present.]
• dry_run(vector<Module>&, Logger&) @ 0x173af40 is the legality/measurement traversal: run the pass without committing emission, used by the verifier and perf-sim "measure" front. [CONFIRMED symbol; STRONG role — the dry-run front itself is documented in the codegen-driver strand.]

The vtable backing all this (vtable for BackendPass @ 0x3d98548) has the standard Itanium layout — RTTI header (offset-to-top, typeinfo*) at vt+0x00, the D1/D0 destructor pair, then the per-unit hook at +0x10. The remaining slot ordering is [INFERRED]: the vtable function pointers are RELR/.rela-packed and read as zero in the raw .data.rel.ro image, so only the +0x10 slot — which the disassembly of run() proves is the dispatched hook — is CONFIRMED.

Construction signature

Every leaf pass shares a uniform construction signature: a base ctor backend::<Class>::<Class>(PassOptions const&). Multi-registered classes additionally expose a parameterized overload that selects the variant. Two CONFIRMED examples:

ColoringAllocator::ColoringAllocator(PassOptions const&, Type, bool,
                                     bir::MemoryAddressSpace, bir::MemoryAddressSpace)  @0x9870c0
DapPass::DapPass(PassOptions const&, std::string const& name, bool)                    @0x87d930

PassOptions is the one argument shared by all factory lambdas — the config bundle the CLI fills from the resolved flags. The extra ColoringAllocator arguments (a Type enum, two MemoryAddressSpace values) are exactly the discriminators that let one class be registered under eight names; see the collapse.

The granularity classes (the fork/join model)

Granularity is layered above BackendPass, not inside it. A leaf pass does not know whether it runs serially or in parallel; that is decided by how it is added to a pipeline. The spine is one container class and three forks.

BackendPass                                  (abstract base, §1)
  └─ ContainerPass : BackendPass             holds an ordered list of child pass factories
       ├─ CoreForkPass     : ContainerPass   PER-CORE parallel  (fork over per-core bir::Modules)
       ├─ ModuleForkPass   : ContainerPass   PER-MODULE parallel (one fork per module in the vector)
       └─ SubgraphForkPass : ContainerPass   PER-SUBGRAPH parallel (calls children's runSubgraph hook)
  └─ DapPass : BackendPass                    adapts an external "DAP" analyzer (NOT a ContainerPass)

ContainerPass — the holder

A ContainerPass is a BackendPass (so it has the same run entry) that holds a list of child factories and runs them in turn:

ContainerPass::addPass(function<unique_ptr<BackendPass>(PassOptions)>)  0x1739cc0   ── add by factory
ContainerPass::addPass(string const&)                                  0x1739d40   ── add by NAME (→ registry)
ContainerPass::runHelper(vector<Module>&, Logger&, string const&)      0x1739da0   ── run the children
ContainerPass::getKind() const                                        0x1737480   ── LLVM-style discriminator
ContainerPass::isRealPass() const                                     0x1743320   ── override

The by-name addPass is the bridge to the registry: handed a string, it asks GeneratorRegistration for the factory and stores it. getKind() plus a per-subclass classof() form an LLVM-style RTTI discriminator — every fork subclass is identified by getKind() and tested with <Fork>::classof(ContainerPass const*). [CONFIRMED — getKind @ 0x1737480 plus all three classof symbols present.]

The three forks

Each fork is a ContainerPass whose run(vector<Module>&) overload forks the children across one axis, runs them, and joins. All addresses CONFIRMED by symbol:

The fork machinery is TBB-backed: the per-element child invocation runs under tbb::detail::r1::execute_and_wait(parallel_for …), the same TBB arena that drives the per-MemoryLocation parallel-for elsewhere in the backend. [STRONG — TBB parallel_for present in the pass bodies; the exact fork wrapping is STRONG, not byte-exact CONFIRMED.]

DapPass — the odd one out

DapPass derives straight from BackendPass, not from ContainerPass. It adapts an external "DAP" analyzer/transform and carries its own name string in its ctor (DapPass(PassOptions const&, string const& name, bool) @ 0x87d930). It exports its own run(bir::Module&) @ 0x87d7e0, run(vector&) @ 0x87db10, and getAdapterInstance() @ 0x87da40. It is nonetheless eligible for subgraph-fork wrapping (see the template instance below). [CONFIRMED — symbols present; non-ContainerPass parentage confirmed by the absence of a DapPass::classof(ContainerPass const*) and presence of BackendPass-shaped run overloads only.]

Granularity is assigned per call-site

The decisive fact: the granularity of a pass is a property of its call-site, not of its class. The BackendPassManager exposes an add* API that wraps a serially-added pass in a fork at that pipeline position:

BackendPassManager::addPass(string)                        → SERIAL child of the current pipe
BackendPassManager::addModParallelPass[WithName]<C>        → wraps C in a ModuleForkPass
BackendPassManager::addCoreParallelPass[WithName]<C>       → wraps C in a CoreForkPass
BackendPassManager::addSubgraphParallelPass[WithName]<C>   → wraps C in a SubgraphForkPass

So the same leaf class can be added serially in one place and inside a fork in another. The CONFIRMED compile-time fork bindings (from the add*ParallelPass<Class> template instantiations in the symtab) are:

MemoryAnalysis and AntiDependencyAnalyzer appear under both module- and core-forks — they run at several pipeline points at different fork granularities. [CONFIRMED — the addModParallelPass<MemoryAnalysis,…> and addCoreParallelPass<MemoryAnalysis,…> template instances both appear in nm -DC.]

A handful of class names end in Pass but are ordinary leaf passes, not granularity bases: LowerSelectPass, OrderConstraintsPass, PerfSimPass, PerfSimPackagePass, SynchronizerPass, VNSplitterPass. Each has its own register_generator_*__ factory in the table below. Do not mistake the suffix for membership in the fork spine. [CONFIRMED — RTTI typeinfos + factory functions.]

The registration mechanism: GeneratorRegistration

Purpose

GeneratorRegistration is the backend's pass-factory registry — the analogue of LLVM's PassRegistry. It is a singleton holding a hashtable from pass name to a factory closure. Every legal --pass <name> is a key in this table; the value is a std::function that, when called with the PassOptions, constructs the pass object. This is the indirection that lets the CLI name passes as strings and lets one class answer to several names.

The API

GeneratorRegistration::getInstance()                       0x17356e0   ── the singleton
GeneratorRegistration::registerGenerator(string, function) 0x1735910   ── install a name → factory
GeneratorRegistration::getGenerator(string const&)         0x1735740   ── name → factory
GeneratorRegistration::hasGenerator(string const&)         0x1735630   ── name validity test (used by the CLI)
GeneratorRegistration::getPassNames[abi:cxx11]()           0x17357e0   ── dump all names (--list-passes)

The backing store is a _Hashtable<string, pair<const string, function<unique_ptr<BackendPass>(PassOptions const&)>>>. This exact value type is visible in the symbol the driver-entry run rehashes against — confirming the registry value really is the factory closure type. [CONFIRMED — all five API symbols, and the hashtable value type function<unique_ptr<BackendPass>(PassOptions const&)> appears verbatim in the _M_rehash symbol called from BackendPass::run.]

getPassNames() is what walrus_driver --list-passes dumps; hasGenerator is the validity gate that turns an unknown --pass into the Backend pass: <X> is not registered! error (3.7).

How a name gets into the table — the static initializers

Each register_generator_<name>__ is a static guard object in .bss (nm class B), e.g. neuronxcc::backend::register_generator_unroll__ @ 0x3dff58c. A global static initializer, run at dlopen, calls:

GeneratorRegistration::getInstance().registerGenerator("<name>",
    std::function<unique_ptr<BackendPass>(PassOptions const&)>( <lambda#1> ));

The lambda is type-erased through std::_Function_handler<…>, which gives two weak functions per name — _M_invoke (call the lambda) and _M_manager (copy/destroy) — plus the lambda's typeinfo. For unroll these are:

_Function_handler<…, register_generator_unroll__::{lambda(PassOptions const&)#1}>::_M_invoke   @0xb46400   ── FACTORY BODY
_Function_handler<…, register_generator_unroll__::{lambda(PassOptions const&)#1}>::_M_manager   @0xb45540
typeinfo for register_generator_unroll__::{lambda(PassOptions const&)#1}                        @0x3d8dce8

Pitfall. The IDA decompiled *.c sidecars for these register_generator lambdas are PLT-thunk stubs, not the real bodies. To read the constructed class you must objdump -d --start-address=0xVA the _M_invoke directly.

_M_invoke is the ground-truth name → class evidence

The _M_invoke body is the canonical proof of which class a name builds: it allocates with operator new and calls the class ctor (or make_unique<Class>), then installs that class's vtable. For unroll the body at 0xb46400 reads (CONFIRMED verbatim):

// _M_invoke for register_generator_unroll__  @0xb46400  [CONFIRMED]
//   b46414:  call  5fb160 <_Znwm@plt>                                  // operator new(sizeof(Unroll))
//   b46422:  call  612430 <neuronxcc::backend::Unroll::Unroll(PassOptions const&)@plt>   // ctor
unique_ptr<BackendPass> unroll_factory(PassOptions const& opts) {
    void *p = operator new(sizeof(backend::Unroll));   // _Znwm
    backend::Unroll::Unroll((Unroll*)p, opts);         // construct → installs vtable for Unroll
    return unique_ptr<BackendPass>((BackendPass*)p);
}

Three evidence shapes recur across the 150 bodies, and the table below tags each row by which one was observed: ctor = a direct backend::Class::Class(PassOptions) PLT call; mkuq = a make_unique<backend::Class> call; vt = the vtable for backend::Class symbol the lambda installs. All three pin the same fact — which class this name constructs. [CONFIRMED — the unroll body shown is the exemplar; every row below was resolved the same way.]

Pipeline-build-time selection

At pipeline-build time, the optlevel pipeline builders call addPass("<name>"), which does getGenerator("<name>")(passOptions) → unique_ptr<BackendPass>. The same registry also serves the arch-Generator lookup: the codegen slot's Codegen pass resolves CoreV{2,3,4}Gen through GeneratorRegistration::getGenerator. The registry is thus the single name-to-object indirection for the entire backend. [CONFIRMED — getGenerator @ 0x1735740 is the shared lookup for both pass and generator slots.]

The 150 → 121 collapse

This is the headline mechanism. 150 register names construct 121 distinct classes because 10 classes are registered under multiple names. The extra names are not aliases in the trivial sense — each one constructs the same class with a different ctor argument, so the binary genuinely needs all 150 names but only 121 vtables.

Verifying the counts

$ nm -DC libwalrus.so | rg -o 'register_generator_[a-z0-9_]+' | sort -u | wc -l
150                                          # ← 150 registered pass NAMES  [CONFIRMED on the binary]

121 is the count of distinct _M_invoke ctor targets across those 150 bodies (the table below). The two are reconciled exactly by the multi-registration groups.

The 10 multi-registered classes

These are the classes that answer to more than one name; the "extra" column is (names − 1), i.e. how many surplus names each contributes over its single class. The discriminator is the ctor flag named in the rightmost column.

Sum of the surplus column: 9 + 7 + 3 + 2 + 2 + 2 + 1 + 1 + 1 + 1 = 29. Therefore 121 + 29 = 150. [arithmetic CONFIRMED — the MemoryAnalysis group's 10 names and the ColoringAllocator/AddressRotation groups were enumerated directly from nm -DC.]

A worked example: MemoryAnalysis is a single class registered under ten snapshot names (memory_analysis_after_unroll, …_after_pre_sched, …_after_post_sched, …_after_address_rotation_{dram,psum,sb}, …_after_coloring_allocator_{sb,dram_post_lnk,dram_shared}, …_after_dma_optimization_sb). Each name's _M_invoke constructs the same MemoryAnalysis with a MemoryAnalysisRequest& that says which prior pass it is profiling — confirmed by the matching addModParallelPass<MemoryAnalysis, char const(&)[N], MemoryAnalysisRequest&> template instances (lengths 29/32/33/42/44/53 — the differing [N] are the differing name-string lengths). Ten names, one vtable.

Reconciling ~180 steps

The pipeline quotes roughly 180 steps, above the 150 names, for two reasons:

1. Repeats. Some names occupy multiple pipeline positions (e.g. non_ssa_legalization early and late, coalesce_multichannel_cc_ops at two points, perf_sim twice). The same registered name is scheduled more than once.
2. Interposed non-real passes. The CheckPassAdder inserts !isRealPass check/dump/sim passes (Dumper, Hasher, perf_sim*, the verifiers, report_stats) between named steps, padding the step count above the name count.

So the three numbers are: 150 registered names (the registry size), 121 distinct classes (the vtable count), ~180 pipeline steps (names × positions + interposed checkers). [CONFIRMED.]

The full 150-name → class → granularity table

Evidence column (ev): ctor = direct backend::Class::Class(PassOptions) PLT call in the _M_invoke body; mkuq = make_unique<Class>; vt = vtable for backend::Class installed by the lambda. Granularity: serial = added by name (default serial child); a Fork tag is shown only where an explicit add*ParallelPass<Class> template instance exists, else serial* (the call-site may still fork — granularity is per-site, see above). All rows CONFIRMED by factory-body evidence unless noted.

Total: 150 register names · 121 distinct C++ pass classes. The serial* tag means added via addPass("<name>") by default; the same class may be fork-wrapped at a different call-site. Fork tags are CONFIRMED only where an explicit add*ParallelPass<Class> template instance exists in the symtab.

Adversarial self-verification

Five strongest claims, re-checked against the binary:

1. 150 → 121 collapse. nm -DC | rg -o 'register_generator_[a-z0-9_]+' | sort -u | wc -l = 150 on the binary. The 10 multi-registered groups were enumerated directly (MemoryAnalysis = 10 names, ColoringAllocator = 8 of the 9 coloring_allocator_* names — the ninth, coloring_allocator_with_loop, is the distinct ColoringAllocatorWithLoop class, AddressRotation = 4, DeadCodeElim = 2). Surplus sums to 29; 121 + 29 = 150. The 121 itself is the distinct-ctor-target count over the table (not independently grep-able), so 121 is STRONG (derived by counting distinct classes in the verified table), while 150 is CONFIRMED on the binary.
2. The three run overloads. run(vector<Module>&) @ 0x1736d40, run(bir::Module&) per-leaf (e.g. Unroll::run(bir::Module&) @ 0xb42f30), runSubgraph (e.g. BirSim::runSubgraph @ 0x1109440) — all three symbols CONFIRMED. The dispatch mov (%r15),%rax ; call *0x10(%rax) read verbatim from 0x1736d70. CONFIRMED.
3. Pass-kind taxonomy. ContainerPass::getKind @ 0x1737480 plus CoreForkPass/ModuleForkPass/SubgraphForkPass::classof @ 0x17374d0/0x1737490/0x17374b0, and DapPass's BackendPass-only run overloads (no classof). All CONFIRMED by symbol. The TBB fork wrapping is STRONG (parallel-for present, exact arena binding not byte-verified here).
4. GeneratorRegistration mechanism. All five API symbols present; the unroll factory body at 0xb46400 shows _Znwm@plt then Unroll::Unroll(PassOptions const&)@plt, and the hashtable value type function<unique_ptr<BackendPass>(PassOptions const&)> appears verbatim in the _M_rehash symbol called from run. CONFIRMED.
5. Hierarchy root. BackendPass is the abstract base; ContainerPass : BackendPass, the three forks : ContainerPass, DapPass : BackendPass — confirmed by the classof(ContainerPass const*) signatures (forks take a ContainerPass*, so they derive from it) and DapPass's lack of one. The full vtable slot ordering beyond +0x10 is INFERRED (RELR-packed pointers read 0 in the static image); only the +0x10 dispatched-hook slot is CONFIRMED.

Re-verify ceiling. Class addresses, the 150 name count, the registry API, the fork classof/getKind symbols, and the unroll factory body are all CONFIRMED against the binary verbatim. The 121 distinct-class total is STRONG (counted over the verified table, not independently grep-confirmed). The full vtable slot layout beyond the dispatched +0x10 hook is INFERRED. The TBB fork-wrapping and the dry_run measurement role are STRONG, not byte-exact. The RTTI typeinfo class-set context (178 backend typeinfos, of which 121 are registry-reachable) is reported STRONG, sourced from the rtti table rather than re-grepped here.

Cross-references

• 3.7 — walrus_driver Backend CLI & Pass Vocabulary — the --pass / --skip-pass / --start-pass vocabulary that draws from these 150 names, and hasGenerator as the validity gate.
• 8.2 — walrus/pass-pipeline-optlevels.md (planned) — the optlevel pipeline builders that call addPass("<name>") and the add*ParallelPass granularity assignment per optlevel.
• Part 14 pass catalog (appendix/, planned) — the per-pass H-strand documentation for each of the 121 classes.
• The Compile Pipeline at a Glance — where the single WalrusDriver Job sits in the process pipeline.