NKI Diagnostic Catalog: the err_/check_/assert_ Funnel

All addresses on this page apply to neuronxcc/nki/compiler/backends/neuron/sema.cpython-310-x86_64-linux-gnu.so from neuronx_cc 2.24.5133.0+58f8de22 (cp310 wheel, md5 755c15de65e6cc00083fab0e5838e4a0, BuildID 688cef0a3cfa701d6f5799767be408075bec1498). The module is not stripped and ships debug_info; all symbol names below are taken verbatim from its symbol table. cp311/cp312 differ only in Python-ABI tags.

Abstract

sema is the NKI (Neuron Kernel Interface) front-end semantic-analysis layer. It runs during kernel tracing — when a @nki.jit / @nki.trace body executes symbolically to build the IR — and validates every nl.* / nisa.* call's operands (shape, dtype, address space, engine, op-legality, architecture limits) before the kernel is lowered to penguin.ir. When a check fails, sema raises a user-facing exception whose message and documentation link are exactly what the kernel author sees in their traceback. This page is the complete catalog of that diagnostic surface.

The surface is a three-tier funnel, mirroring how a SelectionDAG legalizer separates a per-operand predicate from the diagnostic it emits. The bottom tier — the err_* diagnostic builders — are the leaves: each constructs one formatted message, appends the canonical docs URL via the sema_err_url decorator, and the caller raises it. Above them, the assert_* engine evaluates a single boolean predicate and, on failure, looks up an err_* global and calls it (or builds an inline message and raises through the shared nki_assert primitive). At the top, the check_* validators — bound to the NKIFunc class — validate a whole API parameter set and dispatch to assert_*/err_*. The reader who has met the compiler-driver error catalog in 3.20 frontend/diagnostic-error-catalog should read this as the layer below it: sema raises; the driver catches and surfaces.

This is a catalog / reference page, so it omits a reimplementation contract and instead front-loads the funnel mechanism (one annotated example), then gives the full err_* table grouped by concern, then the assert_* and check_* layers and the sema_err_url mechanism. Every message template is re-grounded against the shipped binary's rodata; the docstrings (__doc__) reproduced here are the verbatim source of the rendered nki.errors.html page.

CORRECTION (W05-COUNT) — the backing report headlines "78 err_*". The binary supports a more precise statement. nm over the module-definition symbols (__pyx_mdef_…_<idx>err_*) yields 70 distinct module-level err_* builders — the functions wrapped by sema_err_url. A strings | sort -u over rodata yields 77 distinct err_* names, the extra seven being class-method and Cython-collision-suffixed forms (e.g. err_param_shape_incompatible_wit_2) plus the format helper err_tensor_access_out_of_bound_1range_string. The "78" figure counts that _1range_string helper as a 78th leaf on top of the 77. None of the three counts is wrong; they count different things. The catalog below tabulates the full named surface and labels the helper explicitly, so the reimplementer can reconcile either way. (err_msg appears in rodata but is a field name, not a builder, and is excluded from all counts.)

The Three-Tier Funnel

Purpose

NKI separates what is illegal from how the illegality is reported, so that one diagnostic builder serves many predicates and one predicate can be reused across many APIs. The check_* layer knows the API contract (a matmul needs compatible contraction dims); the assert_* layer knows the primitive predicate (par_dim <= max_p); the err_* layer knows the prose. A reimplementer who collapses these into per-API inline raises will produce a working but unmaintainable surface — the NKI design is deliberately three-tier, and the nki.errors.html docs page is generated from the err_* docstrings, so the split is also the docs-generation boundary.

Pipeline

NKI API call (nl.* / nisa.*) during tracing
           │
           ▼
NKIFunc.<api>  →  check_<concern>(params)        ── tier C: VALIDATORS  (NKIFunc methods)
           │       validates a whole parameter set
           ▼
assert_<predicate>(cond, …, api_name=…)          ── tier B: ASSERT ENGINE  (24 helpers)
           │       if not cond:  f = <err_* global>;  return f(…)
           ▼
err_<diagnostic>(<interpolated fields>)          ── tier A: DIAGNOSTIC LEAVES  (70 builders)
           │       build message  →  sema_err_url.wrapper appends docs URL
           ▼
raise  NKISyntaxError / ValueError / TypeError / IndexError / RuntimeError

Algorithm — the assert→err hop

The canonical funnel is assert_num_partition (sema.py:2877, wrapper 0x96c80) dispatching to err_num_partition_exceed_arch_limit (sema.py:2420, 0x99e90). Both symbols and the docs-rendered message are confirmed in the binary.

function assert_num_partition(par_dim, shapes, max_p, api_name):     // 0x96c80
    cond = PyObject_RichCompare(par_dim, max_p, Py_LE);              // par_dim <= max_p (128)
    if (PyObject_IsTrue(cond))                                       // predicate holds
        return None;
    f = GetModuleGlobalName("err_num_partition_exceed_arch_limit");  // cached global slot, 0x99e90
    return f(shapes, max_p, api_name);                              // tp_call → builds + raises

function err_num_partition_exceed_arch_limit(shapes, max_p, api_name):   // 0x99e90
    // __doc__ (rodata): "Number of partitions exceeds architecture limitation … of 128"
    msg = "number of partitions "        // frag 0x10ac70
        + str(par_dim)
        + " exceed architecture limitation of "   // frag 0x106ae0
        + str(max_p);
    exc = ValueError(msg);                // ValueError recovered as a literal in-wrapper global
    return sema_err_url_wrapper(exc);     // appends docs URL, attaches long __doc__

The verbatim docstring example confirms the chain end-to-end (rodata): x = nl.zeros(shape=[256, 1024], dtype=np.float32, buffer=nl.sbuf) # Error: number of partitions 256 exceed architecture limitation of 128.

NOTE — the err_* global is reached through a Cython-cached module-global slot, not a direct call relocation. That is why many err_* triggers in the catalog below are tagged INFERRED for the trigger even when the builder and template are CONFIRMED: the predicate that funnels into a given err_* is sometimes only resolvable by name-and-semantics, not by a single literal call-site instruction. Three assert_* (assert_num_partition, assert_par_dim_sbuf, assert_par_dim_psum) all funnel to the same err_num_partition_exceed_arch_limit leaf — predicate fan-in is many-to-one.

The nki_assert raise primitive

Asserts that do not route to a named err_* build their own inline message string and raise through one shared primitive:

function nki_assert(exception, condition, msg):       // 0xbb9a0  ·  sema.py:1233
    if (not PyObject_IsTrue(condition)):
        raise exception(msg);                         // exception is a passed-in class object
    return None;

nki_assert is the W04 "assert engine" core: every assert_* helper either funnels to an err_* builder or calls nki_assert(<ExcClass>, cond, <inline msg>) directly. The exception class is an argument, so the same primitive raises ValueError, TypeError, or NKISyntaxError depending on caller intent.

The sema_err_url decorator

Every module-level err_* builder is wrapped by sema_err_url (0x9aed0; closure body sema_err_url.<locals>.wrapper @ 0xf9180). The wrapper takes the builder's raw message and appends the canonical documentation link:

function sema_err_url_wrapper(exc):                   // 0xf9180  (closure over wrapped err_*)
    long_doc  = wrapped.__doc__;                      // the RST docstring (→ nki.errors.html)
    short_msg = exc.args[0];                          // the runtime f-string
    url       = "https://awsdocs-neuron.readthedocs-hosted.com/en/latest/nki/api/nki.errors.html";
    exc.args  = (short_msg + "\n" + url,);            // user sees msg + link
    return exc;

Both halves are confirmed: the URL string is a single rodata literal, and the sema_err_url / sema_err_url.<locals>.wrapper symbols are present. The long RST docstrings reproduced in the catalog below are exactly the per-error sections that nki.errors.html renders; the short f-strings are what lands in the raised exception's args.

QUIRK — the decorator means there is no such thing as a "bare" NKI diagnostic. Every leaf the user sees carries the same trailing URL, and the docs site is not hand-written — it is the __doc__ corpus of these 70 builders. A reimplementer who wants byte-identical diagnostics must reproduce both the short message and the docstring, because the docstring is published.

Exception types

sema raises six exception classes plus one warning. The default is NKISyntaxError (defined in this module, __init__ @ 0xf6c20); where a concrete builtin was recovered as a literal in-wrapper global it is shown explicitly in the catalog. All are confirmed as rodata literals.

Concrete builtins were recovered as in-wrapper literals for ~23 of the leaves (ValueError ×12, TypeError ×8, RuntimeError ×4, IndexError ×3 — overlapping); the rest construct via a Cython-cached class slot and default to NKISyntaxError, tagged NKISyntaxError* below.

err_* Catalog — Group A: Architecture Limits

Partition counts, dimension ranks, dimension sizes, SBUF/PSUM capacity. The hardware invariant behind this whole group is the 128-partition limit and the fixed SBUF/PSUM budgets — these are the legalization gates that map directly onto the SBUF/PSUM geometry.

err_* Catalog — Group B: Shape / Annotation / Broadcast / Matmul

NOTE — err_param_shape_incompatible_with_matmul is reached from check_matmul_high_level_shape via the Cython-collision slot err_param_shape_incompatible_wit_2 (the _2 suffix disambiguates it from ..._with_numpy, which truncates to ..._wit). Both truncations are visible as cached-slot names; the builders are distinct.

err_* Catalog — Group C: Dtype / Operand-Type

err_* Catalog — Group D: Address Space / Memory

HBM / SBUF / PSUM placement. NKI enforces per-API memory-residency requirements; these leaves are the placement legalizer.

err_* Catalog — Group E: Engine / Op-Legality / Arch-Support

Instruction selection. These leaves gate which op runs on which engine, which dtypes a given instruction supports per target, and the minimum NeuronCore generation an instruction requires.

err_* Catalog — Group F: Indexing / Indirect / Out-of-Bound / Atomic RMW

err_* Catalog — Group G: Control-Flow / Tracing / SPMD / Kernel Structure

err_* Catalog — Group H: Activation-Shape / Tiled-Offload / Constant / Hashable

warn_* — the lone non-fatal diagnostic

NOTE — warn_block_dimension_is_deprecated is the only diagnostic that does not abort tracing. It emits through warnings.warn (DeprecationWarning, rodata 0x10b470) and the kernel continues — block dim is still accepted but points the author at the migration guide. Its sibling hard error err_leading_dimension_of_tensor_must_be_partition fires only once block dim is fully removed for a given op. The two share a docstring, so the docs page renders the warning and the error from one source.

The assert_* Engine (24)

The assert_* helpers are the W04 layer: each validates one predicate and, on failure, either funnels to a named err_* or builds an inline message and raises through nki_assert (0xbb9a0). Five funnel to named leaves; the rest raise inline.

Named-err asserts (funnel to an err_* leaf)

Inline-message asserts (raise via nki_assert directly)

These 17 build their own message string and raise through nki_assert. Recovered inline fragments (CONFIRMED rodata) anchor them:

The check_* Validators (NKIFunc methods)

The check_* validators are NKIFunc methods that validate a whole API parameter set and dispatch to assert_*/err_*. The binary has 63 check_* symbols; the named, diagnostic-relevant ones with a recovered dispatch target:

Validators with their own inline messages

Several check_* raise CONFIRMED inline messages with no err_* symbol fronting them — they belong to the same diagnostic surface:

check_quantize_mx_shape (1025):
  "quantize_mx src must have the last free dimension size to be multiple of 4."          0x1063c0
  "quantize_mx src must have 32-multiple partition dimension size."                       0x106420
  "quantize_mx must have same access shape for src and dst except the last free dimension." 0x106460
  "quantize_mx dst must access F // 4 elements per partition for src with F elements …"  0x108d20
check_matmul_mx_low_level_shape (0xe1310) / check_matmul_low_level_shape:
  "nc_matmul_mx stationary must have an even number of packed (x4) elements per partition" 0x106800
  "nc_matmul_mx stationary/moving must access 32/64/128 partitions."                       0x106860
  "nc_matmul_mx moving_scale/stationary_scale must access same elements per partition …"  0x1068c0
check_local_gather_shape (946):
  "<…> can only have input data of type uint32/uint16/uint8, unsupported dtype="          0x105fa0
  "<…> can only have input data of type int32/uint32/uint16/uint8, unsupported dtype="    0x1093c0

Raising-Op Index

The inverse map: which user-facing NKI API can trip which diagnostics. A kernel author who hits a class of errors can scan here for the full set their API surface can raise.

Adversarial Self-Verification

The five strongest load-paths of this page, re-challenged against the binary:

1. "78 err_*" — re-grounded and corrected. nm over __pyx_mdef_…err_* returns 70 module-level builders; strings | rg -o '\berr_…' | sort -u returns 77 distinct names (78 if the _1range_string helper is counted as a separate leaf). The page documents all three figures and the CORRECTION callout reconciles them. The "78" headline counts the format helper as a leaf — defensible but imprecise; the page favors the symbol-grounded 70 + named 77.
2. The funnel assert→err hop is real, not inferred. assert_num_partition (0x96c80) and err_num_partition_exceed_arch_limit (0x99e90) are both present as symbols, and the docstring example # Error: number of partitions 256 exceed architecture limitation of 128. is a verbatim rodata literal — the predicate, builder, and rendered message all confirmed in one chain.
3. sema_err_url appends one URL. sema_err_url and sema_err_url.<locals>.wrapper are both symbols; the URL https://awsdocs-neuron.readthedocs-hosted.com/en/latest/nki/api/nki.errors.html is a single rodata literal. CONFIRMED.
4. nki_assert is the shared raise primitive. __pyx_pw_…_23nki_assert at 0xbb9a0 matches the report's address exactly; NKISyntaxError.__init__ is present, confirming the default class is module-local. CONFIRMED.
5. Message templates are verbatim, not paraphrased. Spot-checks of Traditional while loops are not supported, number of partitions mismatch in parameters (, stack overflow: required sbuf size , and the tensor_tensor docstring example all return exact rodata matches. The per-row templates that could not be matched to a single literal (reassembled from interpolation fragments) carry their fragment addresses inline and inherit the report's INFERRED tagging for the trigger even when the template is CONFIRMED.

GOTCHA — the IDA sidecar for sema.so is absent in this corpus (only missing_addresses/* index files ship for it). All grounding on this page therefore comes from the shipped sema.cpython-310-…so binary itself (nm, strings, DWARF debug_info — the module is not stripped), not from a decompilation database. The addresses are file VAs from the symbol table and are directly checkable with nm; the per-fragment rodata offsets are from strings -t x. A reimplementer can reproduce every cited symbol and string from the wheel with two commands.

Related Components

Cross-References

• NKI Semantic Legality — the check_*/assert_* legality layer that raises these err_* leaves
• Frontend Diagnostic & Error Catalog — the compiler-driver error catalog that sits above this layer and surfaces what sema raises
• SBUF/PSUM Geometry — the 128-partition limit and SBUF/PSUM budgets enforced by Group A
• Range/Loop Semantics — affine_range/sequential_range, source of err_unexpected_output_dependencies
• SPMD Programming Model — program_id and grid launch, source of the SPMD-structure diagnostics