The hlo-opt Pass Registry

All addresses on this page apply to the hlo-opt standalone ELF shipped in neuronx_cc 2.24.5133.0+58f8de22 (cp310 wheel: neuronxcc/starfish/bin/hlo-opt, BuildID xxHash 93dd8bd9bd4c697b, not stripped). The cp311/cp312 builds carry the same registry; addresses are bit-identical across the three CPython variants of this version. Other versions will differ.

Abstract

hlo-opt is the Neuron compiler's standalone HLO optimizer — the XLA-side analogue of LLVM's opt. It reads an HLO module, runs a name-selected subset of the HLO-level rewrite/legalization passes, and writes the transformed module back out. Like opt, every pass it can run lives in a single global registry keyed by a short string name, and the command line selects from that registry by name. This page documents that registry: the one function that fills it (xla::hilo::RegisterHiloHloPasses), the llvm::StringMap it fills, the 112 passes it registers, and the stock-XLA DebugOptions flag mechanism that turns a --passes / --skip-pass name list into the actual run-list.

The shape will be familiar to anyone who has read LLVM's PassRegistry or XLA's HloPassPipeline: a managed-static container, a registration function that runs once at startup, and a name → factory binding so the driver can instantiate a pass by string. The Neuron specifics are (a) the registration is one flat 8.9 KB function of 112 near-identical blocks rather than per-pass static initializers, (b) the registry value is a std::function factory (not a singleton pass), and (c) the key for each pass is not a literal in the registration code — it is whatever the constructed pass's name() virtual returns, read back at registration time. That last point is the central gotcha and the reason the 112 keys had to be recovered from each pass's vtable rather than from the registrar.

This is the index page for Part 4: every other 4.x page documents one or more passes that appear as a row in the table below. The table is the canonical name→class→entry map; the rest of the page explains how a row is bound and how --passes/--skip-pass consume it.

For reimplementation, the contract is:

• The registry container and its key derivation — an llvm::StringMap<std::function<unique_ptr<HloPassInterface>()>> whose key for each entry is the C-string returned by pass->name() (vtable slot vptr+0x10), read at registration time, not a literal in the registrar.
• The registration loop — 112 identical "build a std::function factory, call RegisterHloPass" blocks, executed once from a managed-static guard.
• The selection algorithm — the stock-XLA xla_enable_hlo_passes_only / xla_disable_hlo_passes / xla_disable_all_hlo_passes DebugOptions triad, exact-name matched against the registry keys, plumbed as a flat_hash_set<string_view> into HloPassPipeline::RunPassesInternal.
• The Neuron-vs-stock split — which classes are xla::hilo::Neuron* (authored for this compiler) versus stock xla:: passes reused unchanged.

1. How a Registry Key Is Bound

Purpose

The registry maps a --passes name to a factory that constructs the pass. The subtle part is where the name comes from: the registrar never writes the name string. It builds a factory, hands it to RegisterHloPass, and RegisterHloPass constructs a throw-away instance just to ask it its name, then uses that name as the map key.

Entry Point

_GLOBAL__sub_I_… (managed-static guard)
  └─ xla::hilo::RegisterHiloHloPasses (0x1e72270, 8895 B)   ── 112 identical blocks
       └─ ×112  xla::hilo::RegisterHloPass (0x1ebc3f0, 382 B)  ── binds one key
            ├─ call (*factory)()              ── construct temp pass
            ├─ call [vptr+0x10] == pass->name()  ── StringRef {data,len}
            ├─ llvm::StringMapImpl::hash(StringRef)
            └─ llvm::StringMapImpl::LookupBucketFor(StringRef) ── insert into hloPassRegistry

Algorithm

Each of the 112 blocks in the registrar is the same shape — assemble a std::function factory on the stack, call RegisterHloPass, then destroy the temporary:

// one of 112 blocks inside RegisterHiloHloPasses (0x1e72270)
function RegisterOneBlock():
    var_28 = &Register<Pass>_M_invoke;     // the factory lambda (_M_invoke ptr)
    var_30 = &Register<Pass>_M_manager;    // std::function manager (copy/destroy)
    *(__m128*)&var_40 = 0;                 // zero the std::function inline buffer
    RegisterHloPass(/* &std::function */ &var_40);   // -> 0x1ebc3f0
    if (var_30) (*var_30)(&var_40, DESTROY);         // tear down the temporary

RegisterHloPass does the real work — it derives the key from the constructed pass and inserts the factory:

function RegisterHloPass(factory):                  // 0x1ebc3f0
    pass = (*factory.invoke)();                      // 0x1ebc416 — construct temp pass
    vptr = *(void**)pass;
    name = ((StringRef(*)(void*))vptr[2])(pass);     // 0x1ebc437 — call [vptr+0x10] = pass->name()
                                                     //   returns StringRef{ rax=data, rdx=len }
    h      = llvm::StringMapImpl::hash(name);         // 0x1ebc446
    bucket = registry.LookupBucketFor(name, h);       // 0x1ebc456 -> hloPassRegistry @ 0x9a38ee0
    bucket.value = factory;                           // store the std::function factory
    // (temp pass destroyed on return)

QUIRK — the registry key is not visible in the registrar. The block at 0x1e72270 only stores a factory _M_invoke pointer; the human-readable --passes key is produced later, when RegisterHloPass calls pass->name() on a freshly built instance. To recover the 112 keys you must follow each factory to its <Class> ctor, find the class vtable, read the name() function at vptr+0x10, and dereference the returned StringRef. The registrar symbol name (RegisterXxx) is not the key either — e.g. RegisterNeuronAllGatherCombiner registers the key all-gather-combiner.

The name() literal form

Every <Class>::name() is the same 11-byte stub returning a {data,size} StringRef by value in {edx,eax}:

b8 <len32>      mov  eax, <StringRef.size>      ; length
ba <ptr32>      mov  edx, offset <const char[]> ; data pointer
c3              ret

RegisterHloPass consumes that pair at 0x1ebc437. All 112 declared lengths equal strlen(name) (length integrity check passes for every row). Verified verbatim against the binary for, among others, dce (len 3 @ 0x913dc30), lower-to-nki-kernels (len 20 @ 0x1f3f690), aws_neuron_default_metadata (len 27 @ 0x1f964c0), and the 59-char aws_neuron_rewrite_all_reduce_dynamic_slice_multiple_groups (@ 0x2003870).

NOTE — hloPassRegistry (@ 0x9a38ee0, .bss) is a managed static initialized through a _GLOBAL__sub_I_ guard. It is read once by GetHloPassRegistry (@ 0x1ebc570), which has exactly one call site in the ELF — the driver path that resolves a requested name to its factory. RegisterHiloHloPasses likewise has exactly one call site (the startup initializer). (CONFIRMED — call-target scan of the disassembly.)

2. The 112-Row Master Pass Index

This is the canonical table every Part-4 page slots into. Columns:

• # — index in registration order (1-based). The StringMap is unordered; this order is the registrar's emission order and serves as the default-pipeline proxy (the actual run-list is a driver-supplied name subset — see §3).
• PassName — the verbatim name() C-string = the --passes/--skip-pass key. Byte-exact.
• Class — the constructed class (from the _ZTV… vtable demangle).
• Vtable — _ZTV<Class> address (vptr = this + 0x10).
• Entry — primary body: Run (vptr+0x18) for own-Run passes; for OpExpander passes (shared base Run @ 0x29f0bb0) this is the per-pass InstructionMatchesPattern (vptr+0x28), the real per-pass entry.
• Kind — HloPass (own Run) vs OpExpander (overrides InstructionMatchesPattern/ExpandInstruction, shares the base Run). 25 of 112 are OpExpander.
• Src — N = Neuron-authored (xla::hilo::Neuron* or a Neuron-only xla::hilo::* / xla::* class), S = stock XLA class reused unchanged. See §4 for the rule.

All rows below were recovered by the method in §1; name string + length are read from the binary and length-validated, and vtable + factory + registration order are disassembly-verified. Confidence is CERTAIN for every row unless flagged.

NOTE — the Entry column is the per-pass body, not the name() stub. For the 87 own-Run passes it is the Run slot (vptr+0x18). For the 25 OpExpander passes it is InstructionMatchesPattern (vptr+0x28); their Run slot all points at the shared xla::OpExpanderPass::Run @ 0x29f0bb0, so Run is useless as a per-pass discriminator. The non-virtual ExpandInstruction body for each OpExpander is reachable from its matcher but is not exposed in the vtable; those addresses are not traced here (out of registry scope).

GOTCHA — the class name and the key name diverge for the Neuron collective passes. name() keys drop both the namespace and (for several) the Neuron prefix: xla::hilo::NeuronAllGatherCombiner → all-gather-combiner (#77), xla::hilo::DeviceAssignmentLegalization → the verb-first legalize-device-assignment (#69), xla::hilo::NeuronCollectivePermuteToAllGather → collective-permute-to-all-gather (#86). A reimplementer must read the key off name(), never off the class symbol.

3. --passes / --skip-pass Selection

Purpose

The registry is a superset. The driver narrows it to a run-list. hlo-opt does this with the stock XLA DebugOptions pass-filtering triad — there is no Neuron-specific selection engine; the registry keys are matched directly against the comma-separated name lists in DebugOptions. The --passes / --skip-pass CLI flags are surface names for these DebugOptions fields.

The selection flags

QUIRK — "enable-only" is whitelist semantics, not "also enable". xla_enable_hlo_passes_only makes the listed names the entire run-set and disables everything else; an empty list with this flag set disables all passes. xla_disable_hlo_passes is the blacklist: run everything except the named passes. Setting both is legal — disable wins on conflict (a name in both lists does not run). This is exactly stock XLA HloPassPipeline filter semantics; Neuron added no new arbitration.

Algorithm

The run-time path turns the flag lists into a flat_hash_set<string_view> of disabled names and threads it through the pipeline. Names are matched by exact string equality against each pass's name() (the same StringRef used as the registry key):

// xla::HloPassPipeline::RunPassesInternal<HloModule>(module, debug_opts, disabled_set)
function RunPassesInternal(module, opts, disabled):
    if opts.xla_disable_all_hlo_passes:                 // "*All* passes disabled…"
        return module                                   // run nothing
    enable_only = split(opts.xla_enable_hlo_passes_only)  // whitelist (may be empty)
    disable     = split(opts.xla_disable_hlo_passes)      // blacklist
    for pass in pipeline_passes:                         // driver-ordered subset of the 112
        nm = pass->name()                               // exact key, == registry key
        if enable_only nonempty and nm not in enable_only:
            log "Skipping pass: " nm                     // diag string present in binary
            continue
        if nm in disable or nm in disabled:
            log "Skipping pass " nm
            continue
        Run(pass, module, disabled)                      // OpExpander passes get disabled as the run-arg set

Considerations

• Names must match exactly. Both lists are exact-match against name() — no prefix, no fuzzy, no namespace. The diagnostics Passes enabled by --xla_enable_hlo_passes_only: and Passes disabled by --xla_disable_hlo_passes: are emitted at pipeline build so a typo'd name silently does nothing (it neither enables nor disables a real pass). Validate names against the §2 table.
• Where the run-list / order comes from. The registry is unordered. The executed pipeline is a driver-supplied ordered name subset assembled outside hlo-opt's registrar (in the Neuron pipeline-construction code / Python frontend), then filtered by the flags above. The §2 registration order is the strong default-pipeline proxy but is not itself the run order. (Run-order construction not traced on this page.)
• The disabled set is also a Run argument. OpExpanderPass::Run and the per-pass Run slots take the flat_hash_set<string_view> of disabled execution-pass names as a parameter (stock XLA signature), so sub-pipelines inside a pass honor the same blacklist.

4. Stock-vs-Neuron Split

The rule

Of the 112 registered passes, 109 are Neuron-authored and 3 are stock XLA classes reused unchanged:

All three are plain xla:: (no hilo) classes whose vtables sit in the high 0xd2…/0x48… ranges shared with the rest of the bundled XLA runtime, and their Run bodies are the large stock implementations (HloDCE::Run @ 0x913fb70, CallInliner::Run @ 0x9135790).

CORRECTION (D-A03) — earlier strand notes treated all-reduce-combiner (#79, xla::AllReduceCombiner) as a fourth stock pass because it is a plain xla:: class. The binary shows it is the stock combiner driven by a Neuron combine-key: the class is stock XLA, but its factory is constructed with Neuron-specific combiner parameters and its real entry is RunWithKeyCombiner (Neuron supplies the CombineKey callback). It is therefore classified N (Neuron-configured) here, not pure stock. The two true Neuron collective combiners, #77 all-gather-combiner (xla::hilo::NeuronAllGatherCombiner) and #78 reduce-scatter-combiner (xla::hilo::NeuronReduceScatterCombiner), are Neuron subclasses with their own Run.

Namespace topology

• xla:: (no hilo) — orders 1–43, 51–56, 58–60, 79. These are the legalize/convert/fuse/lowering passes plus the three stock passes and the stock-class combiner.
• xla::hilo:: — orders 44–50, 57, 61–78, 80–112. The alias machinery, the Neuron collective transforms, the rematerializers, the while-loop motion passes, and all 25 OpExpander pattern rewriters.

NOTE — the namespace does not perfectly track "Neuron vs stock": many xla::-namespace passes (e.g. #4 xla::Stub, #5 xla::LegalizeIntrinsics, #60 xla::UpcastAllToFP32) are Neuron-authored despite lacking the hilo namespace — their source paths are all under hilo/hlo_passes/ (e.g. hilo/hlo_passes/Stub.cc, hilo/hlo_passes/LegalizeQuantizeMX.cc, recoverable as .cc path strings in the ELF). Treat the Src column in §2, not the namespace, as the authority on provenance.

Factory shapes

Most factories (105 of 112) construct via a trivial/default ctor that stores the vtable pointer inline (mov qword ptr [rax], offset off_<ZTV+0x10>). 7 use an explicit ctor call because they take constructor arguments read from the registrar's config struct ([rbx+0xDA0] string, [rbx+0xE90]/[rbx+0xF48] ints):

• #1 TupleSimplifier(bool)
• #52 InjectNumericalErrors()
• #77 NeuronAllGatherCombiner(long,long,bool,bool) — four config fields
• #78 NeuronReduceScatterCombiner(…)
• #79 AllReduceCombiner(long,long) — Neuron combine-key supplied separately
• #86 NeuronCollectivePermuteToAllGather(std::string const&,int,int) — reads string + two ints
• #105 NeuronMoveAllGatherWhileLoop(…)

The combiner-threshold / replica-count values in that registrar config struct are not reconstructed here (out of registry scope).

5. Adversarial Self-Verification

The five strongest claims on this page, each re-challenged against the binary:

1. There are exactly 112 passes. An objdump of the RegisterHiloHloPasses body (0x1e72270–0x1e744af) shows 112 call instructions whose rel32 target is RegisterHloPass (0x1ebc3f0). Count = 112. CONFIRMED.
2. RegisterHiloHloPasses is the registrar @ 0x1e72270. nm resolves _ZN3xla4hilo21RegisterHiloHloPassesEv (weak) at that exact address; it has one call site (the startup guard). CONFIRMED.
3. The registry is an llvm::StringMap keyed by pass->name(). RegisterHloPass @ 0x1ebc3f0 calls the factory, then [vptr+0x10] (=name()), then StringMapImpl::hash / LookupBucketFor, inserting into hloPassRegistry @ 0x9a38ee0. The name() stub form (mov eax,len; mov edx,offset str; ret) was read directly for multiple rows. CONFIRMED.
4. The 112 keys and classes in §2 are real. A 12-key string sample (tuple-simplifier, call-inliner, legalize-quantize-mx, lower-to-nki-kernels, neuron-hlo-inst-comb, aws_neuron_default_metadata, neuron-preprocess-kernel-duplicate-remover, aws_neuron_rewrite_all_reduce_dynamic_slice_multiple_groups, while_loop_unroller, upcast-all-to-fp32, all-gather-combiner, aws_neuron_rewrite_all_gather_trip_count) all strings-FOUND, and 8 sampled vtables (TupleSimplifier 0x486f60, HloDCE 0xd26ec8, LegalizeQuantizeMX 0x40e168, Stub 0x4105a8, RemoveAliases 0x410188, ResolveSelfComparison 0x413080, NeuronLoopedEinsumReplacer 0x411d68, NeuronReduceScatterCombiner 0x412150) match the table. The remaining rows were not each independently re-derived in this pass; the table is carried from the disassembly-verified D-A03 enumeration at CERTAIN confidence, with the 8-row spot-check as corroboration. CONFIRMED (sampled).
5. Selection is stock-XLA DebugOptions, not a Neuron engine. The strings xla_enable_hlo_passes_only ("…unspecified passes are all disabled"), xla_disable_hlo_passes, xla_disable_all_hlo_passes, the diagnostics Passes enabled by --xla_enable_hlo_passes_only: / Passes disabled by --xla_disable_hlo_passes: / *All* passes disabled by --xla_disable_all_hlo_passes. / Skipping pass: , and the HloPassPipeline::RunPassesInternal<HloModule>(…, flat_hash_set<string_view> const&) symbol are all present. CONFIRMED.

No fabricated pass names: every key in §2 is the verbatim name() literal recovered from a real vtable. The CLI surface spellings --passes/--skip-pass are mapped to their DebugOptions backing fields above; the exact argv-parsing glue (which option library binds --passes → xla_enable_hlo_passes_only) was not traced and is tagged INFERRED for the spelling, CONFIRMED for the underlying field semantics.

Related Components

Cross-References

• walrus-driver-cli — the frontend pass vocabulary and how the driver assembles the --passes list this registry resolves
• Every Part-4 (4.x) page — each documents one or more rows of the §2 table in depth
• Part-14 pass catalog — the cross-tool index that this hlo-opt registry feeds