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

Appendix — Error-Message Catalog

Every code, count, string, and address on this page applies to neuronx_cc 2.24.5133.0+58f8de22 (cp310 wheel; the catalog is byte-identical across cp310/cp311/cp312 — md5 2c25078c8065b0381143fabf7aff60be over the sorted (cat,idx,cause,res) tuples). Other versions will differ.

Abstract

This appendix is the generated reference dump behind the diagnostics narrative pages (§3.20, §6.4.5). Those pages explain how neuronx-cc's four diagnostic systems are wired; this page is the lookup table — given an error code or message fragment, find its cause, its resolution, and which subsystem emits it. It reproduces three distinct catalogs in full:

  1. The ErrorMessages.so flat dict — NEURON_ASSERT_ERROR_MESSAGES: Dict[(category:str, index:int), (cause:str, resolution:str)], the backend BIR / walrus / codegen / scheduler / NKI / ISA message bodies. 2556 entries across 353 categories, dumped verbatim by importing the shipped .so (D-AG09).
  2. The 224 native NCC_<CAT><idx> codes baked per-call-site into the HLO/penguin frontend binaries (hlo2penguin, hlo-opt, hlo-neff-wrapper, xla_infergoldens) — a separate catalog whose code→text binding lives in those binaries' own rodata (D-A13 §2).
  3. The NKI err_* funnel — 78 err_* diagnostic leaves + 24 assert_* + 32 check_* validators in nki/.../sema.so, the front-end semantic-analysis layer (D-W05).

NOTE — the central artifact here is the ErrorMessages.so dict, and it is CONFIRMED ground truth, not reverse-engineered text. The .so imports standalone (it depends only on typing), so D-AG09 dumped the live Python object rather than reconstructing it from strings. Every (category,index)→(cause,resolution) row in §4 is the exact runtime value. This closes the gap D-A13 §7.1 flagged as "the single biggest incompleteness."

How to use this page

You have…Look in…
A [NCC_<CAT><idx>] token from a backend failure (BIR/walrus/codegen/ISA/scheduler)§4 — find category <CAT>, then index <idx>
A [NCC_<CAT><idx>] from a frontend failure (EVRF/PYP/ISPP/MOD/HPM/OPT/…)§3 — these are not in the ErrorMessages dict
A [F134]/[F137]/[F139] fatal§2
An NKI tracing error (err_*, raised from nl.*/nisa.*)§5
A free-form error:/warning: (MLIR/LLVM)not a closed catalog — see §3.20 §6
ErrorMessages.so entries2556 (cat,idx)→(cause,res) pairs
Categories353 (252 I… internal · 24 E… external · 77 bare/legacy)
Bespoke resolutions461; the other 2095 are {RESOLUTION_CONTACT_SUPPORT} only
Generated 901/902 family446 entries (217 passes × 2 wrappers + extras)
OBSOLETE codes4 — BIR199, BIR202, BIR203, BIR204
Native NCC_* codes224 (separate catalog, frontend binaries)
NKI err_* / assert_* / check_*78 / 24 / 32 (+1 warn_*)
Cross-ABImd5 2c25078c8065b0381143fabf7aff60be (cp310 = cp311 = cp312)

1. Count Reconciliation — the binary says 2556

The "≈2556" figure in the task brief is CONFIRMED against the binary, and it supersedes an earlier strings-based estimate. The reconciliation:

SourceCountMethodStatus
D-A13 §5 (strings sweep)~600 cause/resolution pairsstrings/grep -a on ErrorMessages.so; the (cat,idx)↔text pairing was not recoverable from the string sectionSuperseded estimate
D-AG09 (runtime dump)2556 entriesimport neuronxcc.logging.ErrorMessages; len(NEURON_ASSERT_ERROR_MESSAGES) → 2556CONFIRMED ground truth

CORRECTION (D-A13 §5 → D-AG09) — D-A13 estimated "~600 cause/resolution pairs" because strings sees the interned-string blob and the category tokens, but the (cat,idx)→text adjacency is encoded in a frozen Cython constant-tuple's pointer layout, not interleaved in the string section. The live import recovers the binding losslessly: the true count is 2556, not ~600. The gap between the two figures is mostly the 446-entry generated 901/902 per-pass family plus the fact that many distinct (cat,idx) keys share an identical cause string (e.g. five BIR0xx codes all say Requested Argument index {index} out of bounds), which strings deduplicates and the dict does not.

Independent re-verification (this page): I parsed D-AG09 §8 and counted 2556 CAUSE: lines and 353 ---- … (N codes) ---- headers; the per-header counts sum to 2556 with zero mismatch. First-letter histogram of category headers: I=252 E=24 X=15 L=12 B=9 S=8 M=5 N=4 D=4 C=4 T=3 P=3 A=3 O=2 J=2 V=1 R=1 G=1 — exactly D-AG09 §1's histogram. Bespoke (non-{RESOLUTION_CONTACT_SUPPORT}) resolutions: 461. (CONFIRMED — derived from the dump, cross-checked against the runtime totals.)

Container shape & render path

NEURON_ASSERT_ERROR_MESSAGES : Dict[(category:str, index:int), (cause:str, resolution:str)]
NEURON_ASSERT_GLOBAL_KEYWORDS: {'RESOLUTION_CONTACT_SUPPORT': "Please open a support ticket at
   https://github.com/aws-neuron/aws-neuron-sdk/issues/new. You may also be able to obtain more
   information using the 'XLA_IR_DEBUG' and 'XLA_HLO_DEBUG' environment variables."}

  logging/__init__.so  ──  register_namespace("neuronxcc", NEURON_ASSERT_ERROR_MESSAGES)   (one call)
                                │
  neuron_assert(category, index, condition, **kwargs)        ── Assert.so / neuronxlogger.error
                                │ if not condition:
  cause, res = namespaces["neuronxcc"][(category, index)]     ── the (cat,idx) lookup
                                │
  "[{INTERNAL_ERROR|ERROR}] [NCC_<CAT><idx:03>] {cause.format(**kw)} - {res.format(**kw, RESOLUTION_CONTACT_SUPPORT=…)}"

QUIRK — the [INTERNAL_ERROR] vs [ERROR] prefix is not chosen by parsing the category's first letter. It is chosen by which assert function the caller invokes — neuron_internal_assert emits [INTERNAL_ERROR], neuron_external_assert emits [ERROR]. The I…/E… category naming is a strong convention that co-varies with the prefix (I-cats are raised internally, E-cats externally), but no code branches on CAT[0]. A reimplementer that switches on the first letter will be usually right and occasionally wrong — drive off the assert function, not the name. (D-AG09 §4, confirmed from neuronxlogger/error.py.)

NOTE — the lone "namespace" string registered with neuronxlogger is the literal "neuronxcc"; all 353 categories are an internal sub-key of the one flat dict, not 353 separate register_namespace calls. The per-error "namespace" is the category embedded in the key tuple.

The generated 901/902 per-pass family

446 of the 2556 entries are a templated per-compiler-pass catch-all pair. For a pass class <PassName> with category abbreviation <CAT>:

(<CAT>, 901) = ("<PassName> assertion error: {baseMessage}", "{RESOLUTION_CONTACT_SUPPORT}")
(<CAT>, 902) = ("<PassName> error: {baseMessage}",           "{RESOLUTION_CONTACT_SUPPORT}")

901 = a failed C++ assert inside the pass ({baseMessage} = the assert text); 902 = a non-assert pass error ({baseMessage} = the thrown message). 217 categories carry exactly {901, 902} and nothing else (pure per-pass wrappers); others add 901/902 on top of hand-authored low-index codes. In the catalog below these appear as ordinary rows — the <PassName> is read directly from the 901/902 cause text. (D-AG09 §5.)

2. Driver Fatal Exit Codes

The only three top-level [F1xx] codes. Emitted by driver/CommandDriver.so handleError after the subcommand child process exits abnormally. The codes are the POSIX 128 + signal exit codes of the killed child. (D-A13 §1, CONFIRMED bodies; signal identity INFERRED from the names + OOM wording.)

CodeMessage bodyTriggerSeverity
[F134]neuronx-cc terminated abnormally - Please open a support ticket at https://github.com/aws-neuron/aws-neuron-sdk/issues/newchild exit 134 = 128 + 6 (SIGABRT — failed C++ assert / std::abort / uncaught native exception)FATAL
[F137]neuronx-cc was forcibly killed - This most commonly occurs due to insufficient system memory. Using a smaller data type, dimensions, batch size, or a larger instance type may help.child exit 137 = 128 + 9 (SIGKILL — overwhelmingly the Linux OOM-killer)FATAL
[F139]neuronx-cc terminated abnormally - Please open a support ticket at https://github.com/aws-neuron/aws-neuron-sdk/issues/newchild exit 139 = 128 + 11 (SIGSEGV — segfault)FATAL

grep -a '\[F[0-9]{2,3}\]' over the entire package returns exactly these three, all in CommandDriver.so — a closed set. (CONFIRMED.) A structured [NCC_…] line is typically printed by the failing subprocess before it exits, so a user often sees both an [NCC_…] line and a trailing [F1xx].

3. The 224 Native NCC_* Codes (Frontend Binaries)

These are a separate catalog from the ErrorMessages.so dict in §4. The fully-rendered token NCC_<CAT><idx> is baked per-call-site into the rodata of hlo2penguin / hlo-opt / hlo-neff-wrapper / xla_infergoldens; the cause/resolution body for these frontend categories is not in NEURON_ASSERT_ERROR_MESSAGES (the dict holds the backend bodies). The two catalogs overlap only at IIOT, which appears in the dict solely as the generated (IIOT,901)/(IIOT,902) pair. (D-AG09 §6, by category presence/absence in the dumped dict.)

GOTCHA — do not look up NCC_EVRF001, NCC_PYP001, NCC_ISPP001, NCC_MOD001, NCC_HPM001, NCC_OPT001, NCC_INL001, NCC_INFG001, NCC_IHCA001, NCC_ITUP001, or NCC_IHNW001 in §4 — those frontend categories have zero entries in the dict. Their text lives in the native HLO binaries' rodata and is out of scope for the runtime dump. The ErrorMessages dict is the backend (BIR/walrus/codegen) catalog.

3a. Code format & severity convention

[<prefix>] [NCC_<CATEGORY><III>] <cause> - <resolution>
   <prefix>   = "[INTERNAL_ERROR]"  (I-categories)  |  "[ERROR]"  (E-categories)
   <CATEGORY> = 2–5 uppercase letters; first letter convention: I…=INTERNAL, E…=EXTERNAL
   <III>      = 3-digit zero-padded index  (ErrorCode.__str__ = f"{category}{index:03}")

Index sets are sparse (gaps = retired/obsolete codes); the catalog ranges are enumerated-with-gaps, not contiguous. (D-A13 §2a.)

3b. Native category → emitting binary

BinaryCategories present
hlo-optCTR EUDT EUOC HPM IARE INL ITUP MOD OPT
hlo2penguinCTR DRV ECFT EHCA EMOD ESFH ESPP ETUP EUET EUOC EVRF HPM IARC IARE IHCA IMOD IMUT INL ISPP ITUP IVRF MOD NCO PYP SPP
hlo-neff-wrapperIHNW IIOT
xla_infergoldensCTR EUDT HPM INFG ITUP
libwalrus.sointernal NCC_I… walrus asserts (bodies in §4 via the linked logging::NeuronAssertion)

3c. Native code families (category decode)

CategoryDecodeIndex rangeComponentSeverity
OPTopt-driver / HLO optimization driver001-003hlo-optERROR(ext)
HPMHlo-Pass-Metadata / pass-manager invariant001-003all nativeINTERNAL
EVRF / IVRFHLO/penguin IR VeRiFierEVRF 001-058 · IVRF 003/015/100hlo2penguinext / int
PYPPYthon Pass (pybind HLO→penguin frontend)001-058 (no 002)hlo2penguinINTERNAL
ISPP / ESPP / SPPSetup-Penguin-Pass / penguin pipelineISPP 001-061 · ESPP 003-059hlo2penguinint / ext
MOD / IMOD / EMOD(modular) Module compilationMOD 001-016 · IMOD 020 · EMOD 018-019hlo2penguin/hlo-optmixed
IHCA / EHCAHlo-Custom-call Attr handlingIHCA 001-004 · EHCA 005hlo2penguinint / ext
ITUP / ETUPTUPle flatten/handlingITUP 001 · ETUP 002hlo-opt/hlo2penguin/xla_infergoldensint / ext
INLIN****Line pass001-002hlo-opt/hlo2penguinINTERNAL
INFGINFer-Goldens001-002xla_infergoldensINTERNAL
IHNWHlo-Neff-Wrapper001-003hlo-neff-wrapperINTERNAL
IIOTIo-transpOse Transpose001-002hlo-neff-wrapperINTERNAL
EUDT / EUOC / EUETUnsupported-Dtype / Op / …EUDT 001 · EUOC 001-002 · EUET 001hlo-opt/hlo2penguin/xla_infergoldensERROR(ext)
ECFTCustom-call / Cost FaTal001-003hlo2penguinERROR(ext)
CTR(compiler) ConTRol / contextcommon infraall nativemixed

3d. The handful with verbatim bodies (frontend rodata)

CodeMessage / triggerComponentSeverity
NCC_OPT001-003No valid Hlo passes specified, exiting... / AddPass cannot be called after Run / Error: Required pass not found! Possible causes: … (exact code↔message binding not individually pinned)hlo-opt (offsets 0x28ae5/0x30844/0x2c6f5)ERROR(ext)
NCC_HPM001-003HloPassMetadata for currently running pass not found, either because the pass did not call RecordPassStart or because a pass is creating/switching modules without using HloModuleGroup::ReplaceModule.hlo-opt + all nativeINTERNAL
NCC_IIOT001Duplicate io_transpose found for input: <name> (trigger: combineIoTransposes 0x1e50760 dup-check)hlo-neff-wrapperINTERNAL
NCC_IIOT002Forbidden io_transpose found for input: <name> (trigger: io_transpose on an IntermediateIOs cross-partition tensor)hlo-neff-wrapperINTERNAL
NCC_IHNW003Reshape size mismatch for input: <name> (trigger: processParameters 0x1e5b1d0 — declared reshape ≠ actual param shape)hlo-neff-wrapperINTERNAL

NOTE — the n001/n002/n003 "numbered" codes attributed to StaticIOTranspose (missing-netlist / dup-forbidden / reshape-mismatch) are that module's ad-hoc numbered diagnostics — not NCC_ codes, and distinct from NCC_IIOT/NCC_IHNW. See §3.20 §6b/§6d. The S2-02 claim that CC-ops legalizers emit n00x codes is uncorroborated (D-A13 §7.4).

4. ErrorMessages.so — the Full 2556-Entry Catalog

The complete verbatim runtime dump of NEURON_ASSERT_ERROR_MESSAGES (cp310; byte-identical cp311/cp312), grouped by category in ascending index. The rendered code token is NCC_<Code> (render prefix and [INTERNAL_ERROR]/[ERROR] bracket per §1). Resolution column: *(contact)* abbreviates the auto-appended {RESOLUTION_CONTACT_SUPPORT} support sentence (the 2095 entries that carry only it); a verbatim string means a bespoke resolution (the 461 hand-authored ones). {kwarg} placeholders are runtime-substituted str.format slots. (CONFIRMED — verbatim from the imported .so, D-AG09 §8.)

NOTE — this is the complete table: all 2556 entries across all 353 categories are reproduced below, none dropped. The (category, index) key in each subsection heading plus row is the exact dict key; the cause/resolution cells are the exact tuple value. A reader can look up any backend NCC_<CAT><idx> code by jumping to category <CAT> and finding index <idx>.

4a. Category index (jump table)

Severity classCategories
INTERNAL (I…, 252)IACF IACI IACT IADE IADI IADQ IADR IADV IAGO IAHE IALB IALS IANK IANS IAPR IATL IATN IBCC IBCG IBFC IBIR IBRD IBRW IBSN IBVF ICCC ICCF ICDG ICDL ICIO ICIR ICMA ICMC ICNE ICOE IDCE IDDI IDDT IDEC IDEL IDFS IDFV IDGE IDGM IDIE IDLI IDLO IDMA IDML IDMP IDNT IDSE IDSH IDST IDTP IDVR IEAD IEBN IEDV IEID IEIL IEIM IFAT IFBD IFFD IFGC IFML IFSG IFTA IGAS IGCA IGLO IHAG IHFC IICB IICR IIFG IIGL IIIC IIIN IIIT IIIV IILT IINK IINL IINT IIOT IIPS IISA IISM IIST ILBR ILCB ILCC ILCM ILCX ILDM ILFD ILFU ILIN ILIS ILKK ILLI ILLO ILLP ILLR ILLT ILMM ILNI ILNK ILNM ILOA ILOP ILPP ILPR ILPT ILPX ILRA ILSA ILSC ILSM ILSP ILSR ILSX ILTA ILTO ILTR ILTS ILTY IMCE IMDD IMDT IMGN IMPR IMS IMSA IMSH INBU INCP INDI INDR INDV INEC INIC INIS INKN INLC INLF INLI INPN INSM INSP INST INTC INVN INVR IOAC IOLV IONK IOPF IPAA IPAS IPCC IPDE IPER IPF IPFN IPLF IPLN IPLO IPLX IPMG IPMN IPNI IPNK IPSF IPST IRAC IRCP IRCX IRMT IRNS IROE IROI IRPR IRPX IRRM IRRW IRSA IRSP IRSW ISAB ISAG ISAR ISAU ISCH ISDC ISDD ISFD ISFV ISI ISIM ISIS ISMP ISMT ISMX ISNT ISOE ISPA ISPC ISPR ISPS ISSA ISSD ISSL ISSS ISST ISTA ISTL ISTN ISTP ISUP ISUS ISYC ITCC ITCO ITCT ITDM ITEN ITIN ITLA ITLX ITNM ITOE ITOS ITOT ITPO ITPR ITRF ITSH ITST IUBK IURM IUTL IVDM IVLP IVMM IVNU IWCO IXCG IXLV IZST
EXTERNAL (E…, 24)EAE EARG EBCG EBIR EBRD EBVF EDGE EDMA EDP EDRV EDVR EGCA EIL ELMM ELUR ENKI EOCP EOOM ERP ESL ESMP ETST EXSP EXTP
bare (77)ADA ALR ALS BFD BHT BIB BIR BIRSERDES BLN BMT BUG BVF CCO CFG CFP CUT DAE DMA DRV DVR GCA JIO JOB LCL LDD LDM LEG LKK LKN LLC LLR LNK LRS LSA LUR MFP MLC MLO MMP MSA NAS NKI NKS NLA OAC OQS PER PGT PRS RDP SCH SCP SDD SFP SIM SIO SQI SSA TCE TEN TST VNS XAQ XBI XCG XDR XEI XGM XLB XLS XLV XNP XRA XRO XTP XVL XXU

4b. Full catalog (all 2556 entries, by category)

ADA — bare (1 codes)

CodeCauseResolution
ADA001ADA computation should be called only once per batch. The history should be empty, but not.(contact)

ALR — bare (1 codes)

CodeCauseResolution
ALR001Live-in registers are only supported in legacy loops-on-chip flow. Register {reg_name} defined by Instruction {def_name} in BasicBlock {def_block_name} is last used by Instruction {use_name} in BasicBlock {use_block_name}(contact)

ALS — bare (12 codes)

CodeCauseResolution
ALS010Instruction must exist(contact)
ALS011No local semaphore update found in instruction updates(contact)
ALS012DMA Block has more than one local semaphore update(contact)
ALS013DMA Block was assigned a different local semaphore on reuse after queue switch(contact)
ALS014No CoreBarrier that the remote SB-to-SB DMA instruction is dependent on(contact)
ALS015No CoreBarrier which is dependent on the given remote SB-to-SB DMA instruction(contact)
ALS016Closest CoreBarrier after the given remote SB-to-SB DMA instruction has no dependency relation with it(contact)
ALS017Closest CoreBarrier before the given remote SB-to-SB DMA instruction has no dependency relation with it(contact)
ALS018Got a pair of DMA transpose {transpose} and remote SB-to-SB DMA {sb2sb} with no CoreBarrier in between(contact)
ALS019Got a pair of remote SB-to-SB DMA {sb2sb} and DMA transpose {transpose} with no CoreBarrier in between(contact)
ALS020Only Call/SwitchQueueInstance/TensorCompletion instructions are allowed in main. But Instruction '{name}' is of type '{type}'(contact)
ALS021Expect InstSwitchQueueInstance/InstCall/InstTensorCompletion in main to not have any dependencies or dependents(contact)

BFD — bare (4 codes)

CodeCauseResolution
BFD701{file} at {line} Memloc has no writer or reader(contact)
BFD702unexpected num of writers(contact)
BFD703unexpected instruction type(contact)
BFD704expected a loop instruction(contact)

BHT — bare (2 codes)

CodeCauseResolution
BHT001branch {branchName}'s target does not have first instruction(contact)
BHT002no scope when attempting to insert branch hint for {branchName}(contact)

BIB — bare (1 codes)

CodeCauseResolution
BIB001Cannot create non-FP32 PSUM tensor {tname} in CoreV2(contact)

BIR — bare (354 codes)

CodeCauseResolution
BIR001No hardware model found(contact)
BIR002No hardware model found(contact)
BIR003Invalid AccessPattern(contact)
BIR004Different MemoryLocations for sameAccessPatterns: {memloc}(contact)
BIR005Range of toIntervalSet comes from different MemoryLocations(contact)
BIR006Requested Argument index {index} out of bounds ({num_args})(contact)
BIR007Requested Argument index {index} out of bounds ({num_args})(contact)
BIR008Requested Output index {index} out of bounds ({num_outputs})(contact)
BIR009Requested Output index {index} out of bounds ({num_outputs})(contact)
BIR010Requested Argument index {index} out of bounds ({num_args})(contact)
BIR011Requested Argument index {index} out of bounds ({num_args})(contact)
BIR012Requested Output index {index} out of bounds ({num_outputs})(contact)
BIR013Requested Output index {index} out of bounds ({num_outputs})(contact)
BIR014Instruction with name {name} not found(contact)
BIR015Json field with name {name} not found(contact)
BIR016Register with name {name} not found(contact)
BIR017Register with name {name} not found(contact)
BIR018ctx is nulltpr(contact)
BIR019ctx is nulltpr(contact)
BIR020ctx is nulltpr(contact)
BIR021ctx is nulltpr(contact)
BIR022ctx is nulltpr(contact)
BIR023Unknown dtype '{dtype}'(contact)
BIR024Unknown dtype string '{dtype}'(contact)
BIR025Malformed nki_functions json field(contact)
BIR026Number of input arguments must be the same as recvFomRank size(contact)
BIR027Output dtype must match input(contact)
BIR028Invalid compression ratio: {compress_ratio}! Valid compression ratios are 2, 3, or 4.(contact)
BIR029Expected first free dimension of MatmultSparse ({ifmap_1_num}) to equal compression ratio ({compress_ratio})!(contact)
BIR030DynamicAPINFO::IndirectDimMaxIndex must be > 0(contact)
BIR031For CoreV3+, Matmult in transpose mode must have same input and output dtype (moving input {intype} != output {outtype})(contact)
BIR032For CoreV3+, Matmult in transpose mode must have same input and output dtype (stationary input {intype} != output {outtype})(contact)
BIR033neuron_core_id is not set(contact)
BIR034Mixing of 32-bit and non-32-bit Matmult inputs not supported ({intype1}, {intype2})(contact)
BIR035Transpose matmult requires matching input types ({intype1}, {intype2})(contact)
BIR036MemoryLocation with name {name} not found(contact)
BIR037Size of input tensors for instruction exceed state buffer partition size {input1} + {input2} > {partitionSize}(contact)
BIR038Output size doesn't fit PSUM {outputSize} bytes > {psumSize} bytes(contact)
BIR039Matmult moving input tile size must be <= {sbufMaxRows}x{psumMaxCols} but it was = {ifmapRows}x{ifmapCols}(contact)
BIR040Matmult stationary input tile size must be <= {sbufMaxRows}x{psumMaxRows} but it was = {weightRows}x{weightCols}(contact)
BIR041Matmult output tile size must be <= {psumMaxRows}x{psumMaxCols} But it was = {outputRows}x{outputCols}(contact)
BIR042No support for uninitialized read check in loops(contact)
BIR043Lowered precise schedule loops only support a loop count of one(contact)
BIR044split_BB is deprecated, but got the split_BB attribute on function {function_name}(contact)
BIR045Cannot find already deserialized instruction(contact)
BIR046'unroll_dependencies' must be a JSON array(contact)
BIR047'unroll_dependencies' element must be a JSON array(contact)
BIR048Cannot find unrolled dependency(contact)
BIR049Invalid number of AccessPattern dimensions; must be in [2, 5](contact)
BIR050PE Tile position and size must both be fully specified(contact)
BIR051Matmult instruction does not support PE column tile position 96 on CoreV2(contact)
BIR052PE tile position ({row_tile_pos}, {col_tile_pos}) must be >= start partition of access pattern ({ap_row_pos}, {ap_col_pos})(contact)
BIR053PE tile position ({row_tile_pos}, {col_tile_pos}) + accessed partition number ({ap_row_partitions}, {ap_col_partitions}) must be <= PE array size ({pe_row_size}, {pe_col_size})(contact)
BIR054PE tile size ({row_tile_size}, {col_tile_size}) must be >= accessed partition number ({ap_row_partitions}, {ap_col_partitions})(contact)
BIR055Symbolic Matmult second input base partition mismatch across iterations(contact)
BIR056Symbolic Matmult output base partition mismatch across iterations(contact)
BIR057Matmult output PE tile position cannot be allocated to 96 on CoreV2(contact)
BIR058Matmult instruction does not support PE tile size 32x32 on CoreV2(contact)
BIR059DMACopy does not support {cce_op} in CCE mode in instruction {name}(contact)
BIR060DMACopy with SW DGE does not support {cce_op} in CCE mode in instruction {name}(contact)
BIR061DMACopy with HW DGE does not support {cce_op} in CCE mode in instruction {name}(contact)
BIR062Unexpected DGE mode {dge_type} in DMACopy instruction {name}(contact)
BIR063DMADescriptorCCE does not support {op} in instruction {name}(contact)
BIR064GenericIndirectSave does not support {op} in instruction {name}(contact)
BIR065IndirectSaveAccumulate does not support {op} in instruction {name}(contact)
BIR066{kind} CollectiveCompute does not support {op} in instruction {name}(contact)
BIR067{kind} CollectiveCompute does not support {op} in instruction {name}(contact)
BIR068Unexpected CollectiveCompute kind {kind} in instruction {name}(contact)
BIR069Unexpected instruction type {opcode} in check CCE AluType in instruction {name}(contact)
BIR070DoWhile instruction missing loop condition(contact)
BIR071DoWhile instruction condition must be an IntRuntimeValue(contact)
BIR072CompareAndBranch instruction has: {num_args} arguments. Expected 2 arguments.(contact)
BIR073CompareAndBranch expects both args to have same dtype. Dtype for arg1: {dtype_1}; Dtype for arg2: {dtype_2}(contact)
BIR074{function_name}: Function had already been visited by BIR verifier(contact)
BIR075Collective compute instructions are not allowed in DoWhile loops(contact)
BIR076--dge-levels dst_reduce not enabled, cannot have DGE instruction with a destination reduction operation(contact)
BIR077DMACopy does not support {cce_op} with {mode} mode in instruction {name}(contact)
BIR078Empty basic blocks are not allowed: {block_name}(contact)
BIR079Basic block: {block_name} does not end with Terminator instruction. Last instruction: {inst_name} is not a terminator(contact)
BIR080Basic block: {block_name} non-terminator Instruction {inst_name} is not allowed after the first terminator Instruction(contact)
BIR081Basic block: {block_name} is a loop block that must contain only sequencer/terminator instructions. Instruction {inst_name} is not a sequencer or terminator instruction(contact)
BIR082Control flow graph of function {function_name} has multiple exit nodes: {block_a}, {block_b}(contact)
BIR083Control flow graph of function {function_name} does not have an exit node(contact)
BIR084Control flow graph of function {function_name} has multiple entry nodes: {block_a}, {block_b}(contact)
BIR085Control flow graph of function {function_name} does not have an entry node(contact)
BIR086Found duplicate targets for a back edge from {source_block} to {target_block_1} and {target_block_2} in control flow graph of function {function_name}(contact)
BIR087Symbolic Access Pattern input arguments are not allowed in DoWhile loops post unroll stage(contact)
BIR088Symbolic Access Pattern output arguments are not allowed in DoWhile loops post unroll stage(contact)
BIR089Control flow graph basic block count {cfg_block_count} mismatches the function {function_name} it is derived from: {function_block_count}(contact)
BIR090Function {function_name} basic blocks not in topological order: {block1} found before predecessor {block2}(contact)
BIR091Unknown ArgumentKind(contact)
BIR092Live-in/Live-out MemoryLocation: {memloc_name} not allocated to MemoryType: DRAM in function: {func_name}(contact)
BIR093InstDoWhile instruction incompatible with legacy loop implementations '--loops-on-chip' and '--loops-in-backend'(contact)
BIR095Different function attribute: stage_unroll_completed values for function: {func1_name} and function: {func2_name}(contact)
BIR096Invalid Module JSON filepath {fp}(contact)
BIR097Unreachable block {unreachable_block_name} found in {function_name}(contact)
BIR098Unsupported DGE Type: expect one of DGE types {valid_dge_types} but got {target_dge_type}(contact)
BIR099Gather instruction expects maximum 4 dimensions in params access pattern but found {dimensions}(contact)
BIR100Gather instruction expects maximum 4 dimensions in indices access pattern but found {dimensions}(contact)
BIR101Gather instruction expects maximum 4 dimensions in outputs access pattern but found {dimensions}(contact)
BIR102Gather instruction expects the same data type for params: {params_dtype} and outputs: {outputs_dtype}(contact)
BIR103Gather instruction expects the same number of elements per partition for indices: {indices_per_partition} and outputs: {outputs_per_partition}(contact)
BIR104Gather instruction expects the same number of partitions for params: {params_partitions} and indices: {indices_partitions}(contact)
BIR105Gather instruction expects the same number of partitions for params: {params_partitions} and outputs: {outputs_partitions}(contact)
BIR106Gather instruction expects type uint32 for indices: {indices_dtype}(contact)
BIR107Only 1 terminator is allowed on each engine, but engine {engine_name} has 2(contact)
BIR108If there is a terminator on EngineType::ALL/Unassigned, there must not be terminators on any other engine(contact)
BIR109If there is not a terminator on EngineType::ALL/Unassigned, there must be a terminator on all {num_datapath_engines} engines, but there are only terminators on {actual} engines(contact)
BIR110Matmult cannot write to multiple PSUM banks for architecture {arch}(contact)
BIR111Unsupported data type Dtype::{actual} for InstCompareAndBranch 1st argument. The supported types are Dtype::float32/int32/uint32(contact)
BIR112Unsupported data type Dtype::{actual} for InstCompareAndBranch 2nd argument. The supported types are Dtype::float32/int32/uint32(contact)
BIR113SB CollectiveKind {kind} requires non-empty replica groups(contact)
BIR114Unhandled data type {dtype} for bir::ImmediateValue::toString(contact)
BIR115Unhandled data type {dtype} for bir::ImmediateValue::createFromJson(contact)
BIR116Unhandled data type {dtype} for bir::ImmediateValue::dataToJson(contact)
BIR117Unhandled data type {dtype} for bir::ImmediateValue::clone(contact)
BIR118Unhandled data type {dtype} for bir::SymbolicImmediateValue::createFromJson(contact)
BIR119Unhandled data type {dtype} for bir::SymbolicImmediateValue::toJson(contact)
BIR120Unhandled data type {dtype} for bir::ImmediateArray::createFromJson(contact)
BIR121Unhandled data type {dtype} for bir::ImmediateArray::toJson(contact)
BIR122GetSequenceBounds instruction can only be assigned to Pooling engine, but it is assigned to {actual}(contact)
BIR123GetSequenceBounds instruction expects source AccessPattern to be three dimension, but is {actual}(contact)
BIR124GetSequenceBounds instruction expects destination AccessPattern to be three dimension, but is {actual}(contact)
BIR125GetSequenceBounds instruction expects partition dimension of source AccessPatttern to be 1, but it is {actual}(contact)
BIR126GetSequenceBounds instruction expects partition dimension of destination AccessPatttern to be 1, but it is {actual}(contact)
BIR127GetSequenceBounds instruction expects last dimension of destination AccessPattern to be 2, but it is {actual}(contact)
BIR128GetSequenceBounds instruction expects data type of input AccessPattern to be Dtype::int32 or Dtype::float32, but it is {acutal}RESOLUTION_CONTACT_SUPPORT
BIR129GetSequenceBounds instruction expects data type of output AccessPattern to be Dtype::int32 or Dtype::float32, but it is {acutal}RESOLUTION_CONTACT_SUPPORT
BIR130GetSequenceBounds instruction expects destination AccessPattern to be twice the size of source AccessPattern, but source has {source_num} and destination has {dst_num}RESOLUTION_CONTACT_SUPPORT
BIR131Custom operations not yet supported for TRN2(contact)
BIR132Unsupported expression kind ({kind}) in iterator(contact)
BIR133ScalarTensorTensor instruction requires 2D or 3D inputs, but received inputs with {pattern1_size} and {pattern2_size} dimensions respectively(contact)
BIR134ScalarTensorTensor instruction requires 2D or 3D output, but received output with {output_pattern_size} dimensions(contact)
BIR135Lower bound expression for DynamicForLoop is of unexpected kind {kind}(contact)
BIR136Lower bound expression for DynamicForLoop expected to be constant(contact)
BIR137Upper bound expression for DynamicForLoop is of unexpected kind {kind}(contact)
BIR138Upper bound expression for DynamicForLoop expected to contain only one term(contact)
BIR139Upper bound expression for DynamicForLoop expected to contain a single integer runtime value(contact)
BIR140Stride expression for DynamicForLoop is of unexpected kind {kind}(contact)
BIR141Stride expression for DynamicForLoop expected to be constant(contact)
BIR142One of InstCompareAndBranch's targets must be the next BasicBlock ({next}), but they are {true} and {false}(contact)
BIR143InstTensorLoad/InstTensorSave must have at least 1 input, but it has none(contact)
BIR144InstTensorLoad/InstTensorSave must have at least 1 output, but it has none(contact)
BIR145InstTensorLoad/InstTensorSave can only reference Const/Internal/Pointer tensors, but it is referencing a {type} tensor(contact)
BIR146Attempted to find loop header for {bb} which is not in a loop body(contact)
BIR147Attempted to find loopy body for {bb} is not a loop header(contact)
BIR148Upper bound expression for DynamicForLoop references undefined register(contact)
BIR149DynamicForLoop instruction not in a top-level basic block(contact)
BIR150Unknown argument type for transpose matrix multiplication outputs!(contact)
BIR151PSUM partition on transpose matrix multiplication outputs must be 0, got partition {partition} instead!(contact)
BIR152Array data structure expects even number of elements but got: {nibbleCount}(contact)
BIR153Array index out of of bounds: Attempted to fetch nibble at index {index} but array size is {nibbleCount}(contact)
BIR154Array index out of of bounds: Attempted to set nibble at index {index} but array size is {nibbleCount}(contact)
BIR155Output of xbar transpose must have steps must be 32 byte aligned. Output dimension {dim} has step size of {step} which is not divisible by 32 when multiplied with the dtype size {dtype_size}.(contact)
BIR156Tensor size must be < 4GB, but it is {actual} bytes(contact)
BIR157{inst_name}: Replication not supported for MatmultMX instructions(contact)
BIR160DEPRECATED, see BIR606 -- {inst_name}: MatmultMx is only supported for engine: Tensor. Got engine: {engine}(contact)
BIR161DEPRECATED, see BIR530 -- {inst_name}: MatmultMx is only supported for arch: core_v4 or newer. Got older arch: {arch}(contact)
BIR162DEPRECATED, see BIR606 -- {inst_name}: QuantizeMx is only supported for engine: Vector. Got engine: {engine}(contact)
BIR163DEPRECATED, see BIR530 -- {inst_name}: QuantizeMx is only supported for arch: core_v4 or newer. Got arch: {arch}(contact)
BIR164{inst_name}: QuantizeMx has invalid input access pattern: {iap}. Innermost dimension invalid. Step of innermost dimension must be 1 and num of innermost dimension should be a multiple of 4(contact)
BIR165{inst_name}: QuantizeMx has invalid input access pattern: {iap}. Step for all dimensions except innermost dimension should be an even number(contact)
BIR166{inst_name}: Input AP and Output AP should access same number of elements. Input AP accesses {numIap} elements and Output AP accesses {numOap} elements(contact)
BIR167{inst_name}: Starting quadrants for ifmap: {if_quad} and ifmap_scales: {iscales_quad} dont match(contact)
BIR168{inst_name}: Starting quadrants for weights: {w_quad} and weights_scales: {wscales_quad} dont match(contact)
BIR169{inst_name}: Start partition for ifmap_scales: {if_start} and weights_scales: {w_start} must be a multiple of 4(contact)
BIR170{inst_name}: Expected num_active_channels: {channels} to be either 32. 64. 96 or 128(contact)
BIR171Unsupported TransposeOps::{actual} for architecture {arch}. Only XZYW is supported.(contact)
BIR172Transpose dst must be in SB, but it is in {actual}(contact)
BIR173DGE xbar transpose src must be in SB or DRAM, but it is in {actual}(contact)
BIR174DGE xbar transpose datatype must be 2 bytes, but it is {type} ({size} bytes)(contact)
BIR175DGE xbar transpose src must have 16 partitions, but it has {actual}(contact)
BIR176DGE xbar transpose expected src tile to have {expected} columns, but it has {actual}(contact)
BIR177DGE xbar transpose src columns must be consecutive. Last dimension step should be 1 but it is {actual}(contact)
BIR178DGE xbar transpose dst columns must be consecutive. Last dimension step should be 1 but it is {actual}(contact)
BIR179DGE xbar transpose dst rows must be consecutive. Partition dimension step should be {expected} but it is {actual}(contact)
BIR180DGE xbar transpose dst tile starting address must be a multiple of 32, but it is {actual} (remainder of {remainder})(contact)
BIR181Unsupported TransposeOps::{actual}. Only XZYW and XYZW are supported.(contact)
BIR182Instruction reads un-initialized memory locations: {details}(contact)
BIR183Num elements per partition for ifmap {num} should be 4x num elements per partition for ifmap_scales {scales_num}(contact)
BIR184Num elements per partition for weights {num} should be 4x num elements per partition for weights_scales {scales_num}(contact)
BIR185Start partition of ifmap_scales: {part} should be < 16(contact)
BIR186Start partition of weights_scales: {part} should be < 16(contact)
BIR187Max number of elements per partition for ifmap: {num} should be <= {max_ifmap_elements} ({max_elements_per_psum_bank} * 4(PE quad-row mode) * 8(PSUM banks)) for output dtype {output_dtype_name}(contact)
BIR188Max number of elements per partition for weights: {num} should be <= 512 (128 * 4)(contact)
BIR189Num active rows for instruction must be 32, 64 or 128. Got: {numRows}(contact)
BIR190Number of partitions accessed by ifmap {ifmap} should equal number of partitions accessed by weights {weights}(contact)
BIR191Number of elements per partition in weights {weights} should be 4x number of output partitions {op}(contact)
BIR192Number of elements per partition in ifmap {ifmap} should equal to 4x number of elements per partition in output {op}(contact)
BIR193Partitions accessed by ifmap {ifmap} should match partitions accessed by corresponding scales {scales}(contact)
BIR194Partitions accessed by weights {weights} should match partitions accessed by corresponding scales {scales}(contact)
BIR195PE Row tile size {row_tile_size} of MatmultMx must match number of input partitions accessed {num_partitions}(contact)
BIR196MatmultMx does not support column tiling(contact)
BIR197More than one entry found in CurrentCallsites(contact)
BIR198More than one entry found in CurrentCallers(contact)
BIR199OBSOLETE ERROR CODE. DO NOT USE. - More than one entry found in CurrentNKICallsites(contact)
BIR200Expected CurrentCallsites size to be 0 before visiting InstCall {name}(contact)
BIR201Expected CurrentCallers size to be 0 before visiting InstCall {name}(contact)
BIR202OBSOLETE ERROR CODE. DO NOT USE. - Expected CurrentNKICallsites size to be 0 before visiting {name}(contact)
BIR203OBSOLETE ERROR CODE. DO NOT USE. - Expected CurrentNKICallers size to be 0 before visiting NKI Function {name}(contact)
BIR204OBSOLETE ERROR CODE. DO NOT USE. - More than one entry found in CurrentNKICallers(contact)
BIR205NKI Kernel function pointer is empty in instruction {name}(contact)
BIR206DGE xbar transpose dst start partition must be a multiple of {expected}, but it is {actual} (remainder of {remainder})(contact)
BIR207CopyPredicated only allows max as Reduce Operator. Given: {actual}(contact)
BIR208CopyPredicated only allows EngineAccumulationType of Idle, ZeroAccumulate, or AddAccumulate. Given: {actual}(contact)
BIR209CopyPredicatedReduce has AccumulationCmd set to non-Idle ({accumulation_cmd}), but does not have a corresponding ReductionCmd ({reduce_cmd}).(contact)
BIR210CopyPredicated is set to use reduction mode, requires 3 arguments. (pred_tensor, src_tensor, immediate) Received: {actual}(contact)
BIR211CopyPredicatedReduce requires src1 to be a tensor(contact)
BIR212CopyPredicatedReduce requires src2 to be a scalar(contact)
BIR213CopyPredicatedReduce allows any dtype permitted by DtypePair, except float32r, uint64, int64, int32, uint32. given: src type: {src_type}, predicate type: {pred_type}(contact)
BIR214CopyPredicated requires predicates to have same number of elements as Output. Given predicates: {pred_shape}, and output={out_shape} elements respectively(contact)
BIR215CopyPredicated requires that src data (if a tensor) must have the same shape as output and predicate tensor. Given src: {src_shape}, predicate={pred_shape}, and output={out_shape} elements respectively.(contact)
BIR216CopyPredicatedReduce immediate value must be float32. But it is {actual}.(contact)
BIR217CopyPredicatedScalar Immediate value type ({imm_type}) must match output type ({out_type}).(contact)
BIR218Non-scalar and Non-Reduce CopyPredicated cannot have reverse predication(contact)
BIR219Start partition {start_partition} of QuantizeMx scales be a multiple of 32 (start of quadrant) plus 0, 4, 8, or 12(contact)
BIR220Number of data partitions {op} should match number of scale partitions {scales}(contact)
BIR221{inst_name}: ActivationCopy cannot have an MX (microscaled) source or destination(contact)
BIR222cast_fp32_to_float4e2m1fn_x2 expects count {count} to be a multiple of 2(contact)
BIR223cast_fp32_to_float8e4m3fn_x4 expects count {count} to be a multiple of 4(contact)
BIR224cast_fp32_to_float8e5m2_x4 expects count {count} to be a multiple of 4(contact)
BIR225{inst_name}: QuantizeMx input elements per partition {inp_elts} should be 4x number of output scales per partition {scales}(contact)
BIR226InstBNStatsAggregate must output 2 elements per partition, but it outputs {actual}(contact)
BIR227Axis {axis} not found in current loop nest(contact)
BIR228State buffer allocation failed: total size of tensors that need to be allocated in state buffer to execute {inst} exceeds state buffer per-partition capacity.\nTotal accessed bytes per partition: {instruction_accessed}\nTensor details:\n{memloc_accessed}\nSB partition size: {partition_size}(contact)
BIR229State buffer allocation failed: total size of tensors that need to be allocated in state buffer to execute {inst} exceeds state buffer capacity.\nTensor details:\n{memloc_accessed}\nTotal available SB Size: {total_size}(contact)
BIR230Indirect tensor {name} of Vector DGE DMACopy must have MemoryType::DRAM, but it has MemoryType::{actual}Do not load the indirect tensor into SB first
BIR231Non-indirect tensor {name} of Vector DGE DMACopy must have MemoryType::SB, but it has MemoryType::{actual}Use an SB tensor as the non-indirect tensor
BIR232Index tensor {name} of Vector DGE DMACopy must be in SB, but it is in {actual}Load the indices into SB first, then use them in this instruction
BIR233TransposeDMA with dynamic offset cannot fallback to static DMA(contact)
BIR234CCE op is only valid on DMACopys that have mode=CCE or use vector DGE (dst_reduce)This is likely a bug in NKI. You may be able to get around it by not setting dst_rmw_op or using vector DGE on the nisa.dma_copy
BIR250DMA that is not valid for DGE must have DGEType::None, but it has DGEType::{actual}(contact)
BIR251Static DMA that is in a dynamic loop can only read from Internal tensor, but it reads from {actual} tensor(contact)
BIR252Static DMA that is in a dynamic loop can only write to Internal tensor, but it writes to {actual} tensor(contact)
BIR253Instruction Type InstTensorScalarPtr does not support EngineAccumulationType Accumulate, only supports Idle, Zero, ZeroAccumulate, and AddAccumulate(contact)
BIR254Instruction Type InstTensorTensor does not support EngineAccumulationType Accumulate, only supports Idle, Zero, ZeroAccumulate, and AddAccumulate(contact)
BIR255Instruction Type InstRangeSelect does not support EngineAccumulationType Accumulate, only supports Idle, Zero, ZeroAccumulate, and AddAccumulate(contact)
BIR256(Instruction: {inst}) ShardId can only be in one block dimension.Consider rewriting the ShardId-related access to use ShardId in only one block dimension.
BIR257(Instruction: {inst}) Block dimension expression with ShardId can only be an AffineExpr or ShardId; got a non-AffineExpr.Consider rewriting the ShardId-related access to use an affine expression instead of a more complex expression.
BIR258(Instruction: {inst}) Block dimension expression with ShardId can only be an AffineExpr with a single ShardId index; got multiple ShardId indicesConsider rewriting the ShardId-related access to use only one ShardId as an index in the affine expression.
BIR259Number of packed elements per partition for weights ({num}) should be even(contact)
BIR293Found MemoryLocationSet that has not been read, {name}(contact)
BIR294Found MemoryLocationSet that has not been read, {name}(contact)
BIR295Fail to find tensor scope parent(contact)
BIR296tensor_scope_parent can only be Function or InstLoop/InstDynamicForLoop/InstDoWhile(contact)
BIR297For specific instructions, base partition for access is expected to be equal if both inputs are in SB, but {instr} violates this constraint.(contact)
BIR298Could not find indirection argument {argId} parent={parent} location={location}{additionalInfo}(contact)
BIR299Unexpected instruction argument type: {type}(contact)
BIR300Unexpected instruction argument type string: {name}(contact)
BIR301Unexpected instruction argument type ({type}) when adding argument or output(contact)
BIR309Expected start address {addr} of output to be 2B aligned (NOT 8B/4B aligned) if outputAP step_elem_x=-1(contact)
BIR310Expected matmult instructions in same accumulation group: {instr1} with output dtype {dtype1} and {instr2} with output dtype: {dtype2} to have same output dtypes(contact)
BIR311Only Matmult and Memset instructions can write BF16 outputs to PSUM(contact)
BIR312Expected start address {addr} of output to be 4B aligned if outputAP step_elem_x=1(contact)
BIR400Couldn't open file '{file}' to write BIR json(contact)
BIR401Memset mode 'Random' is only supported on Vector Engine for NeuronCore V2(contact)
BIR402Memset mode 'Random' is only supported on Vector Engine and GpSIMD Engine for NeuronCore V3(contact)
BIR403Seed datatype must be 32 bits(contact)
BIR404TensorIndirect Symbolic AP must not directly call "Unroll" API since it needs additional dynamic AP handling(contact)
BIR406Requested Indirect Argument index {index} out of bounds ({num_indirect_args})(contact)
BIR407Requested Indirect Argument index {index} out of bounds ({num_indirect_args})(contact)
BIR410Inferring psum_zero_region for matmul but output {mloc} is not allocated(contact)
BIR411Inferring psum_zero_region for matmul but its size or start addr not consistent with psum zero region(contact)
BIR412Instruction: {opcode} has invalid memory location type: {actual_type}. Supported memory location types are: {expected_type}(contact)
BIR420SB CC kind {kind} not supported for dimension {dim}(contact)
BIR421Argument or output {arg} not contiguous for SB global CC(contact)
BIR422Argument or output {arg} not on SB for SB global CC(contact)
BIR423{kind} on SB input and output shapes mismatch, got:\n input 0 shape: {input0_shape}\n input 1 shape: {input1_shape}\n output shape: {output_shape}(contact)
BIR424AllReduce on SB input and output shapes mismatch, got:\n input shape: {input_shape}\n output shape: {output_shape}(contact)
BIR425AllGather on SB on {dim} input, output, and replica groups shapes mismatch, got:\n input shape: {input_shape}\n output shape: {output_shape}\n replica groups shape: {replica_groups_shape}(contact)
BIR426ReduceScatter on SB on {dim} input, output, and replica groups shapes mismatch, got:\n input shape: {input_shape}\n output shape: {output_shape}\n replica groups shape: {replica_groups_shape}(contact)
BIR427SB CC should only have a single output.(contact)
BIR428SB CC is only supported on trn2+.(contact)
BIR429Access must be aligned to datatype, {mem_type} start address {start_addr} for memory location {mem} must be {dtype_size} bytes aligned for data type {dtype}.(contact)
BIR430Expect no InstLoop after unroll, violation: {inst}.(contact)
BIR431Expect no InstDynamicForLoop after unroll, violation: {inst}.(contact)
BIR432Expect no InstDoWhile after unroll, violation: {inst}.(contact)
BIR433Expect no InstGenericIndirectLoad/Save after unroll, violation: {inst}.(contact)
BIR434Expect no SymbolicAccessPattern after unroll, violation: {inst}.(contact)
BIR435Expect no tensor scope parent on memory location sets after unroll, violation: {mset}.(contact)
BIR436Fine grained InstCollectiveCompute cannot enter as coalesced, this is an internal feature, violation: {inst}, {kind}.(contact)
BIR437Consecutive InstCollectiveCompute cannot have the same kind after coalesce_multichannel_cc_ops, violation: {inst1}, {inst2}.(contact)
BIR438CCE DMA instructions must have at most {arg_limit} arguments after legalization, violation: {inst} with {num_args} arguments.(contact)
BIR439Expect no InstAbstractCopy after lower_ac, violation: {inst}.(contact)
BIR440Expect all {mem_type} memory locations to be allocated after {pass_name}, violation: {mem}.(contact)
BIR441Expect {inst} write to PSUM with 4-byte aligned dataType, violation: {mem}. {exceptions_msg}(contact)
BIR444Invalid InstTensorScalarPtr {tsp} on Activation Engine, op0: {op0}, op1: {op1}, rev0: {rev0}, rev0: {rev1}.(contact)
BIR450GPSIMDSB2SB is only available for LNC=2(contact)
BIR451GPSIMDSB2SB is only available in trn2 and later(contact)
BIR452GPSIMDSB2SB cannot be used to cast(contact)
BIR453GPSIMDSB2SB can only move data between access patterns of the same size(contact)
BIR454GPSIMDSB2SB access patterns must have a valid number of dimensions(contact)
BIR455GPSIMDSB2SB can only move data between SBs(contact)
BIR456GPSIMDSB2SB does not support dynamic access patterns(contact)
BIR457GPSIMDSB2SB tensors must have partition dimension that's a multiple of 16(contact)
BIR458GPSIMDSB2SB is only available in GPSIMD engine(contact)
BIR459GPSIMDSB2SB must have a source and destination(contact)
BIR460GPSIMDSB2SB can only transfer less than {maxBPP} bytes/partition, but is trying to transfer {bpp} bytes/partition{maybeOverhead}(contact)
BIR480Expected attribute max (value: {max_val}) to be greater than min (value: {min_val}) for random number generation(contact)
BIR481Invalid seed for engine: GPSIMD. Seed must be an nx6 tensor. Got seed tensor of dimensions nx{dim}(contact)
BIR482Invalid seed for engine: Vector. Seed must be an nx24 tensor. Got seed tensor of dimensions nx{dim}(contact)
BIR500DGE xbar dynamic transpose can only use DGEType::SWDGE, but it uses DGEType::{actual}(contact)
BIR501DGE xbar dynamic transpose must be on EngineType::Pool, but it is on EngineType::{actual}(contact)
BIR503DGE xbar dynamic transpose source tensor ('{tensor_name}') address must be {expected} byte aligned, but it is {actual}(contact)
BIR504DGE xbar dynamic transpose index tensor ('{tensor_name}') must have Dtype::uint32, but it has Dtype::{actual}(contact)
BIR505DGE xbar dynamic transpose index tensor ('{tensor_name}') must have MemoryType::SB, but it has MemoryType::{actual}(contact)
BIR506DGE xbar dynamic transpose must use a multiple of 16 indices, but it uses {actual}(contact)
BIR507DGE xbar dynamic transpose index tensor ('{tensor_name}') must have 1 element per partition, but it has {actual}(contact)
BIR508DGE xbar dynamic transpose index tensor ('{tensor_name}') address must be {expected} byte aligned, but it is {actual}(contact)
BIR510DGE xbar dynamic transpose only supports TransposeOps::XZYW, but it is TransposeOps::{actual}(contact)
BIR511DMA transpose input and output should have the same type, but it has in Dtype::{in} out Dtype::{out}(contact)
BIR512DMA transpose outer dimensions are not transposed. [{in0}, {in3}] (in) -> [{out0}, {out3}] (out)(contact)
BIR514DGE xbar dynamic transpose src dimension 1 or 2 should be padding (Num=1), but neither is (dim1={dim1}, dim2={dim2})(contact)
BIR515DGE xbar dynamic transpose dst dimension 1 or 2 should be padding (Num=1), but neither is (dim1={dim1}, dim2={dim2})(contact)
BIR516Unable write json to {path}.Please ensure diretory exists and disk space is available.
BIR518DGE xbar dynamic transpose must have 1 AccessPattern output, but it has {actual}(contact)
BIR519DMA transpose with a vector of dynamic offsets is not supported on trn1Try compiling for trn2 instead
BIR520DGE xbar dynamic transpose src actual pattern num partitions ({src}) doesn't match number of indices ({idx})(contact)
BIR521DGE xbar dynamic transpose outer dimensions are not transposed. [{in0}, {in3}] (in) -> [{out0}, {out3}] (out)(contact)
BIR530{opcode} is not supported on arch {arch}, must be {min_arch} or greaterDo not use {opcode} or switch to at least {min_arch} architecture
BIR531{opcode} is not supported on arch {arch}, must be {max_arch} or lessDo not use {opcode} or switch to at most {max_arch} architecture
BIR532{opcode} instruction must have enum of type {enumType} initializedInitialize enum values for instruction
BIR533{enumVal} is not a valid enum value for field {fieldName} on arch {arch}Choose a valid value for field {fieldName}
BIR534TensorTensor only supports EngineAccumulationType::Idle, but got {actual}Use ScalarTensorTensor instead, or give a supproted AccumulationEngine.
BIR535Compilation with --target=trn1 generated local collective compute instruction. Local collectives supported only for trn2+. Please re-compile by setting compilation target to trn2+(contact)
BIR600DMA trigger has DMA blocks specifying QoS class {actual} but only {mappedCount} classes are mapped by runtime(contact)
BIR601Instruction specifies DMA QoS class {actual} but only {mappedCount} classes are mapped by runtime(contact)
BIR602Command line specifies {given} QoS classes are available on the target but the ISA encoding limits this to {max}(contact)
BIR603Failed to encode DMA QoS priority for JSON/ISA from value {value}(contact)
BIR604Failed to decode DMA QoS priority for JSON/ISA from value {value}(contact)
BIR605Bias parameter for InstActivation can only be an ImmediateValue if the activation function is Copy or Reciprocal, but got {actual}. This restriction does not apply to trn2+ architectures.Use a vector-immediate (Nx1 tensor) as the bias parameter instead, or recompile for trn2+
BIR606Engine {engine} is invalid for {opcode} instruction on arch {arch}. Must be {validEngines}.Set engine to {validEngines}
BIR700Unrecognized InstIO type {type} at index {index}(contact)
BIR701Tensor indirection is only supported on core_v4 and later(contact)
BIR702Indirection tensor for argument {index} must be in {expected} memory but is instead in {actual}(contact)
BIR703Indirection tensor for output {index} must be in {expected} memory but is instead in {actual}(contact)
BIR704Indirection tensor for output {index} must be in the same partition as the argument, partition {expected}, but is instead in partition {actual}(contact)
BIR705Tensor indirection in tensor engine can only gather data of FP16, BF16, and FP8 data types, but argument {index} is of type {actual}(contact)
BIR706Tensor indirection in scalar engine cannot scatter to PSUM, but output {index} does(contact)
BIR707Scatter tensor indirection requires that destination's indirection tensor be in the same partition as any source, but the source data is in {src} and the destination's indirection tensor is in {dst}(contact)
BIR709Tensor indirection is not supported for TensorReduce when axis type is C or XYZWC(contact)
BIR710Tensor indirection is only supported for CopyPredicated without a reduction operation(contact)
BIR711Tensor indirection is only supported for CopyPredicated involving a scalar true value(contact)
BIR712Tensor indirection is not supported for tensor-scalar-addr cases of TensorScalarPtr(contact)
BIR713Tensor indirection is not supported for scalar-tensor-tensor cases of TensorScalarPtr(contact)
BIR714Tensor indirection is not supported for tensor-tensor-scan cases of TensorScalarPtr(contact)
BIR715Tensor indirection is not supported for transpose matmul(contact)
BIR716Tensor indirection is not supported for row-tiled matmul(contact)
BIR717Tensor indirection is not supported for column-tiled matmul(contact)
BIR718{ioDirn} {ioIndex} can't be tensor-indirect because it maps to a {ioDirn} on the TensorTensor ISA that does not support it(contact)
BIR719Argument {index} doesn't support tensor indirection(contact)
BIR720Output {index} doesn't support tensor indirection(contact)
BIR721AccessPattern::getNumElementsPerPartitionConsideringIndirection(): AP must have more than 1 dim(contact)
BIR722Argument {index} doesn't support tensor indirection(contact)

BIRSERDES — bare (1 codes)

CodeCauseResolution
BIRSERDES001BIR signatures mismatch before and after serialization + deserialization(contact)

BLN — bare (1 codes)

CodeCauseResolution
BLN001

BMT — bare (1 codes)

CodeCauseResolution
BMT001Unknown BackendMetricType {unknown_type}.(contact)

BUG — bare (1 codes)

CodeCauseResolution
BUG001Walrus driver ran with flag --bugpoint-throw-after to trigger this error for testing(contact)

BVF — bare (29 codes)

CodeCauseResolution
BVF001InstTensorScalarCache must have 2 or 3 inputs, but has {actual}(contact)
BVF002InstTensorScalarCache first input must be an AccessPattern but it is a {actual}(contact)
BVF003InstTensorScalarCache must have 1 or 2 outputs, but has {actual}(contact)
BVF004InstTensorScalarCache first output must be an AccessPattern but it is a {actual}(contact)
BVF005InstTensorScalarCache first immediate value must have Dtype::float32, but it has Dtype::{actual}(contact)
BVF006InstTensorScalarCache second immediate value must have Dtype::float32, but it has Dtype::{actual}(contact)
BVF007InstTensorScalarCache with 3 arguments must have EngineAccumulationType::LoadAccumulate, but it has EngineAccumulationType::{actual}(contact)
BVF008InstTensorScalarCache with EngineAccumulationType::LoadAccumulate must have 3 arguments, but it has {actual}(contact)
BVF009InstTensorScalarCache reduce output must be an AccessPattern but it is a {actual}(contact)
BVF010InstTensorScalarCache reduce output must have 1 element per partition, but it has {actual}(contact)
BVF011InstTensorScalarCache reduce output must have same number of partitions as first argument ({expected}), but it has {actual}(contact)
BVF012InstTensorScalarCache reduce output must have a floating-point Dtype, but it has Dtype::{actual}(contact)
BVF013InstTensorScalarCache first input has invalid type Dtype::{actual}(contact)
BVF014InstTensorScalarCache first output has invalid type Dtype::{actual}(contact)
BVF015InstTensorScalarCache op0 must be arithmetic, but it is AluOpType::{actual}(contact)
BVF016InstTensorScalarCache op1 must be add/subtract/mult/min/max, but it is AluOpType::{actual}(contact)
BVF017InstTensorScalarCache with TSCMode::TensorScan must have AluOpType::bypass for op0, but it has AluOpType::{actual}(contact)
BVF018InstTensorScalarCache with TSCMode::TensorScan cannot have reverse0 = true(contact)
BVF019InstTensorScalarCache on Sunda cannot have reverse0 = true(contact)
BVF020InstTensorScalarCache on Sunda cannot have reverse1 = true(contact)
BVF021InstTensorScalarCache must be on the DVE engine, but it is on {actual}(contact)
BVF022InstTensorScalarCache EngineAccumulationType must be one of ZeroAccumulate/AddAccumulate/LoadAccumulate, but it is {actual}(contact)
BVF023InstTensorScalarCache with TSCMode::TensorScan must have an ImmediateValue as its 2nd input, but it has a {actual}(contact)
BVF024InstTensorScalarCache with TSCMode::TensorScan must have immediate value of 0 as its 2nd input, but it has {actual}(contact)
BVF025InstTensorScalarCache does not support AluOpType::rsqrt(contact)
BVF026InstTensorScalarCache SymbolicImmediateValue values must all be 0, but one of them is {actual}(contact)
BVF027Instruction can only read one of its non-scalar inputs from PSUM, but inputs {args} are read from PSUMCopy tensor(s) from PSUM to SB prior to using this instruction
BVF028Instruction can only read one of its inputs from PSUM, but inputs {args} are read from PSUMCopy tensor(s) from PSUM to SB prior to using this instruction
BVF029Invalid in-place write at {inst}: {arg} and {out} overlap in memory, and read is not guaranteed to complete before write at overlapping location(contact)

CCO — bare (10 codes)

CodeCauseResolution
CCO001Expected all consecutive CC compute ops to be of the same Op, but got {first} != {curr}(contact)
CCO002Expected all consecutive CC compute ops to be of the same Engine, but got {first} != {curr}(contact)
CCO003Expected all consecutive CC compute ops to be of the same Kind, but got {first} != {curr}(contact)
CCO004Expected all consecutive CC compute ops to be of the same ReplicaGroups(contact)
CCO005Expected all consecutive CC compute ops to be of the same HasOrder(contact)
CCO006Expected all consecutive CC compute ops to be of the same CanReadUninit(contact)
CCO007Expected all consecutive CC compute ops to be of the same IsLocal(contact)
CCO008Expected all consecutive CC compute ops to be of the same SrcTargetPairs(contact)
CCO009Expected all consecutive CC compute ops to be of channel id starting at 0 and increasing by 1, but got first: {first}, curr: {curr}(contact)
CCO010Expected CC op to not have both ReplicaGroups and SrcTargetPairs(contact)

CFG — bare (19 codes)

CodeCauseResolution
CFG001Control flow graph of function {function_name} has multiple exit nodes: {block_a}, {block_b}(contact)
CFG002Control flow graph of function {function_name} does not have an exit node(contact)
CFG003Control flow graph of function {function_name} has multiple entry nodes: {block_a}, {block_b}(contact)
CFG004Control flow graph of function {function_name} does not have an entry node(contact)
CFG005Unreachable block {unreachable_block_name} found in {function_name}(contact)
CFG006Found duplicate targets for a back edge from {source_block} to {target_block_1} and {target_block_2} in control flow graph of function {function_name}(contact)
CFG007Unable to find predecessor of non-entry block {block_name}(contact)
CFG008Block {block_name} is not the entry block, but is listed as having no immediate dominator(contact)
CFG009Unable to find successor of non-exit block {block_name}(contact)
CFG010Block {block_name} is not the exit block, but is listed as having no immediate postdominator(contact)
CFG011Function {function_name} basic blocks not in topological order: {block1} found before predecessor {block2}(contact)
CFG012Cannot find Basic Block {bb_name} in cfg in getPredecessors() call(contact)
CFG013Cannot find Basic Block {bb_name} in cfg in getSuccessors() call(contact)
CFG014Cannot find Basic Block {bb_name} in cfg in getOutDegrees() call(contact)
CFG015Loop must be reducible, detected violation back-edge: {back_edge} where {dst} does not dominate {src}.(contact)
CFG016Loop headers require exactly one external in-edge as we currently do not support IF statements.\nDetected the following loop header: {header} with entries: {entry_edges}.(contact)
CFG017Loop structures require exactly one non-exit external out-edge as we currently do not support BREAK statements.\nDetected the following loop structure:\n {basic_blocks}\nwith exits:\n{exit_edges}.(contact)
CFG018Loop structures require exactly one internal back-edge as we currently do not support CONTINUE statements or nested loops.\nDetected the following loop structure:\n {basic_blocks}\nwith internal back-edges:\n{back_edges}.(contact)
CFG019Basic blocks require exactly one forward edge with the exception of loop entries and exit blocks, as we currently do not support IF statements.\nDetected {basic_block} with forward edges: {forward_edges}.(contact)

CFP — bare (2 codes)

CodeCauseResolution
CFP001CoreForkPass called with single module {name}(contact)
CFP002Compilation failed for modules targeting the following cores:{errors}.(contact)

CUT — bare (8 codes)

CodeCauseResolution
CUT001Missing {name} vertex in NX wavegraph(contact)
CUT002Length of wavegraphs differ after transform(contact)
CUT003Missing {name} in graph.(contact)
CUT004Input must be multiple of 2(contact)
CUT005Invalid input type {dtype}(contact)
CUT006Can't convert to NC format, need (N,C,1,1) dims(contact)
CUT007Unable to compute left padding(contact)
CUT008Unable to compute right padding(contact)

DAE — bare (2 codes)

CodeCauseResolution
DAE001Lock directory does not exist(contact)
DAE002Unknown verbosity level, choose from debug|info|warning|user(contact)

DMA — bare (15 codes)

CodeCauseResolution
DMA126Unrecognized Architecture: {target}(contact)
DMA127Unrecognized Architecture: {target}(contact)
DMA128Undefined DRAM Memloc {mem}(contact)
DMA129Undefined SB Memloc {mem}(contact)
DMA130Unrecognized Architecture: {target}(contact)
DMA131Unrecognized Architecture: {arch}\n(contact)
DMA132Unrecognized Architecture: {arch}\n(contact)
DMA133DRAM Memloc with zero access {mem}(contact)
DMA134Attempted to merge load for a memloc that shouldn't be removed: {mem}(contact)
DMA135Attempted to transform Load to TensorCopy for an instruction that shouldn't be touched: {inst}(contact)
DMA136Undefined local memloc: {mem}(contact)
DMA137Try to access removed memorylocation(contact)
DMA138Could not find PSUM partition range index {end} among live track of size {live_size}; partition range indexes into ranges 0-32, 32-64, 64-96, 96-128.(contact)
DMA139Need to build valid CFG for DMA optimization pass(contact)
DMA140Found Instruction with unsupported DGE type DGEType::HWDGE on Sunda Arch(contact)

DRV — bare (78 codes)

CodeCauseResolution
DRV001Deprecated store true action does not permit any arguments(contact)
DRV002Enable disable action does not permit any arguments(contact)
DRV003Enable disable action expected a default value(contact)
DRV004Enable disable action expected one option string(contact)
DRV005{initial_option} does not start with '--enable' or '--disable'(contact)
DRV006No-able action does not permit an argument(contact)
DRV007No-able action expected a default value(contact)
DRV008No-able action expected one option string(contact)
DRV009{initial_option} does not start with '--' or '--no-'(contact)
DRV010Set argument action expected a const to use for the value of the action(contact)
DRV011While parsing {arg_string} {a}(contact)
DRV013Unable to find base classes of Job(contact)
DRV014Unknown verbosity level, choose from debug|info|warning|user(contact)
DRV015Start time not set in TimeRegion context exit(contact)
DRV016Delta could not be computed in TimeRegion context exit(contact)
DRV017XLA flow only takes 1 model file(contact)
DRV018Meta module not specified(contact)
DRV019Meta module '{abs_meta_module}' not found(contact)
DRV020No fp32-cast option in args: {codegen_args}(contact)
DRV021Expected target of 'cpu' or 'tonga', but got {target}(contact)
DRV022Insufficient target info(contact)
DRV023Expected to find 'dataflow' or 'functions' in ir_type data(contact)
DRV024Expected ir_type data to be a dict or list(contact)
DRV025HhSubgraphs expected hhRes to be a tuple(contact)
DRV026NEFF archive without accelerator or CPU compute ops is not supported. Please check that your model has non-constant computation that is dependent on the inputs listed in --io-config(contact)
DRV027Output {a} not found in network(contact)
DRV028CPU nodes not supported(contact)
DRV029Expected length of from format and shape to match(contact)
DRV030Expected length of from and to formats(contact)
DRV031Expected layout to be 'native'(contact)
DRV032--jf_data_layout=tonga not supported with --io-config(contact)
DRV033Expected shape to have a non-zero length(contact)
DRV034Unhandled data type {a}(contact)
DRV035Missing act_info for used function set {used_act_set}(contact)
DRV036Expected NEFF hash-length to be 32(contact)
DRV037Expected NEFF version string to be 'Dynamic'(contact)
DRV038Expected fewer than 64 feature bits in NEFF(contact)
DRV039Expected 16-byte UUID(contact)
DRV040Header size is {a}(contact)
DRV041Only versions 2+ supports non-zero feature bits(contact)
DRV042Experimental --autoloop-spec {autoloop_file} not found(contact)
DRV043Could not read kelpInfo for NEFF(contact)
DRV044Input {a} for sg{idoffset} not found in network inputs or previous graph(contact)
DRV045Output {a} not found in network(contact)
DRV046Expected single input BIR file(contact)
DRV047Meta module doesn't exist(contact)
DRV048Module has no reference_inputs(contact)
DRV049Module has no model object(contact)
DRV050Expected module to contain a 'vm_inputs' entry(contact)
DRV051Expected module to contain a 'mod' entry(contact)
DRV052InferGoldens failed in creating inference io context using {a}(contact)
DRV053BIRSim output {tensor_name} is not in output name map!(contact)
DRV054Multiple MLA nodes are not supported(contact)
DRV055Kelf json file does not match expected format(contact)
DRV056Compiler is not saving temps(contact)
DRV057Expected lengths of source and target layouts to match(contact)
DRV058Expected non-negative indices in transform(contact)
DRV059Expected two entries in random range match(contact)
DRV060Expected two entries in line space range match(contact)
DRV061No value available for tensor {name}(contact)
DRV062Got {a} outputs but expected {b} - {outputs}(contact)
DRV063{metadata} does not exist(contact)
DRV064{weight_file} does not exist(contact)
DRV065Model has {a} inputs while {b} given by --images(contact)
DRV066Invalid keep rate index(contact)
DRV067Tensor {name} has NaN(contact)
DRV068HLO module {input} doesn't exist(contact)
DRV069{metadata} does not exist(contact)
DRV070Discovered outputs length is {a} while only {b} golden output produced(contact)
DRV071Multiple entries added for same file key(contact)
DRV072Unknown --list argument {arg}, choose from supported|all|missing(contact)
DRV073Illegal arguments. Please use --option-preset dma_tensor=,...(contact)
DRV074Unsupported arguments in --option-preset!(contact)
DRV075Unsupported --output type {arg}, choose from text|json(contact)
DRV076Unexpected NEFF header(contact)
DRV077Compilation artifacts already exist in the output directory ({working_dir}).Please ensure to run the compiler with a clean artifacts directory to avoid unintended mixing of compilation artifacts between multiple models.
DRV078Discovered intermediates length is {a} while only {b} golden intermediates produced(contact)
DRV079Only accepts at most 1 model file(contact)

DVR — bare (24 codes)

CodeCauseResolution
DVR000Expect input file specified with -iSee neuronx-cc --help for CLI usage
DVR001Internal validation failed(contact)
DVR002Internal transformation failed(contact)
DVR003Cannot specify --optlevel and --pass at the same timeSee neuronx-cc --help for CLI usage
DVR005Cannot override storage allocator with --allocator when passing --pass optionsSee neuronx-cc --help for CLI usage
DVR006Internal pipeline construction error(contact)
DVR007--enable_partitioner may not be used with -O0See neuronx-cc --help for CLI usage
DVR008--enable_partitioner may not be used with -O6See neuronx-cc --help for CLI usage
DVR009--enable_partitioner may not be used with -O8See neuronx-cc --help for CLI usage
DVR010Unsupported --allocator '{allocator}'See neuronx-cc --help for CLI usage
DVR011Unspported --optlevel '{optlevel}'See neuronx-cc --help for CLI usage
DVR012Unspported --optlevel '{optlevel}'See neuronx-cc --help for CLI usage
DVR013--enable-call-graph may not be used with -O0See neuronx-cc --help for CLI usage
DVR014--enable-call-graph may not be used with -O6See neuronx-cc --help for CLI usage
DVR015--enable-call-graph may not be used with -O8See neuronx-cc --help for CLI usage
DVR016--enable_partitioner may not be used with -O8See neuronx-cc --help for CLI usage
DVR017--enable-call-graph may not be used with -O8See neuronx-cc --help for CLI usage
DVR018Expected subgraph name in the format 'ncXX/sgXX', but got {subgraph_name}(contact)
DVR019Expected NC prefix of 'ncXX', but got {nc_name} from subgraph {subgraph_name}(contact)
DVR020Expected NC index below {vnc_nc_count}, but got {nc_idx} from subgraph {subgraph_name}(contact)
DVR021Linear-scan allocator does not currently support VNC, -O2 or -O3 is recommended for VNC.See neuronx-cc --help for CLI usage
DVR022Expected numeric NC index in {nc_name}, but got {nc_num} in subgraph {subgraph_name}(contact)
DVR023Internal error: Duplicate attempts to setup console logging.(contact)
DVR024Internal error: Duplicate attempts to setup file logging.(contact)

EAE — EXTERNAL (43 codes)

CodeCauseResolution
EAE001EngineType::ALL is not supported in legacy loops-on-chip flow(contact)
EAE003InstCompareAndBranch must have EngineType::ALL in non-loops-on-chip flow, but it has EngineType::{actual}(contact)
EAE004InstCompareAndBranch must not have a RegisterAccess argument before expand_inst_late(contact)
EAE005InstCompareAndBranch must have 2 arguments, but it has {actual}(contact)
EAE006Unsupported Instruction type on EngineType::ALL: {actual}(contact)
EAE007InstUnconditionalBranch must have EngineType::ALL in non-loops-on-chip flow, but it has EngineType::{actual}(contact)
EAE008InstUnconditionalBranch must have 0 arguments, but it has {actual}(contact)
EAE009InstUnconditionalBranch must have 0 outputs, but it has {actual}(contact)
EAE010InstCompareAndBranch must have 0 outputs, but it has {actual}(contact)
EAE011InstTensorLoad must have EngineType::ALL in non-loops-on-chip flow, but it has EngineType::{actual}(contact)
EAE012InstTensorLoad must have 1 argument, but it has {actual}(contact)
EAE013InstTensorLoad must have 1 output, but it has {actual}(contact)
EAE014InstTensorLoad argument must be a PhysicalAccessPattern, but it is a {actual}(contact)
EAE015InstTensorLoad output must be a RegisterAccess, but it is a {actual}(contact)
EAE016InstRegisterAlu must have EngineType::ALL in non-loops-on-chip flow, but it has EngineType::{actual}(contact)
EAE017InstRegisterAlu must have 2 arguments, but it has {actual}(contact)
EAE018InstRegisterAlu must have 1 output, but it has {actual}(contact)
EAE019InstRegisterAlu arguments must be RegisterAccesses, but they are {actual0} and {actual1}(contact)
EAE020InstRegisterAlu output must be a RegisterAccess, but it is a {actual}(contact)
EAE021InstCompareAndBranch cannot have any dependencies(contact)
EAE022InstCompareAndBranch cannot have any descendents(contact)
EAE023InstUnconditionalBranch cannot have any dependencies(contact)
EAE024InstUnconditionalBranch cannot have any descendents(contact)
EAE025InstRegisterMove must have EngineType::ALL in non-loops-on-chip flow, but it has EngineType::{actual}(contact)
EAE026InstRegisterMove must have 1 argument, but it has {actual}(contact)
EAE027InstRegisterMove must have 1 output, but it has {actual}(contact)
EAE028InstRegisterMove argument must be a RegisterAccess or ImmediateValue, but it is a {actual}(contact)
EAE029InstRegisterMove output must be a RegisterAccess, but it is a {actual}(contact)
EAE030InstCompareAndBranch 1st argument must be a PhysicalAccessPattern/RegisterAccess/ImmediateValue, but it is a {actual}(contact)
EAE031InstCompareAndBranch 2nd argument must be a PhysicalAccessPattern/RegisterAccess/ImmediateValue, but it is a {actual}(contact)
EAE032Register ({reg}) not found in oldReg2Engine2NewReg!(contact)
EAE033InstExit must have EngineType::ALL, but it has EngineType::{actual}(contact)
EAE034InstExit must have 0 arguments, but it has {actual}(contact)
EAE035InstExit must have 0 outputs, but it has {actual}(contact)
EAE036InstExit cannot have any dependencies(contact)
EAE037InstExit cannot have any descendents(contact)
EAE038Late expansion instruction must have 0 arguments, but it has {actual}Remove all arguments from Instruction.
EAE039Late expansion instruction must have 0 outputs, but it has {actual}Remove all outputs from Instruction
EAE040Late expansion instruction cannot have any dependenciesRemove all dependencies from Instruction.
EAE041Late expansion instruction cannot have any descendentsRemove all descendents from Instruction.
EAE042Instruction {opcode} must have EngineType::ALL in in non-loops-on-chip flow, but it has EngineType::{actual}Reassign to ALL so that the instruction can be expanded properly.
EAE043Late expansion instruction cannot have any Synchronization info.Remove all Synchronization information from instruciton.
EAE044Instruction {opcode} must have EngineType::ALL in in non-loops-on-chip flow, but it has EngineType::{actual}Reassign to ALL so that the instruction can be expanded properly.

EARG — EXTERNAL (2 codes)

CodeCauseResolution
EARG001Illegal argument(s) - Logical Neuron Core size {lnc} is not supported on trn1. The specified architecture only supports Logical Neuron Core size of 1.Update the Logical Neuron Core size to 1 or check that the target architecture is correct.
EARG002Illegal argument(s) - the following argument(s) are unrecognized: {unrecognized_args}.Check that only recognized arguments are used: {usage}

EBCG — EXTERNAL (1 codes)

CodeCauseResolution
EBCG001Custom operations not yet supported for {hardware_target}(contact)

EBIR — EXTERNAL (21 codes)

CodeCauseResolution
EBIR006Integer TensorTensor Ops has PSUM argument.Move argument to SB to allow for execution on GPSIMD engine (which cannot access PSUM), or change operands to float so they can execute on Vector Engine
EBIR010All of InstRegisterAlu's input registers must have the same bitwidth. Got {str}Ensure all Register inputs to InstRegisterAlu have the same bitwidth
EBIR011InstTensorLoad/InstTensorSave register and tensor must have the same number of bytes. Got src={src_type} dst={dst_type}Ensure a 32-bit index is accompanied by a 32-bit register, and a 64-bit index is accompanied by a 64-bit register
EBIR012InstRegisterMove src and dst must have the same bitwidth. Got src={src_type}, dst={dst_type}Ensure InstRegisterMove's input and output have the same type (or at least the same number of bits)
EBIR013InstRegisterAlu with 64-bit operands can only perform AluOpType::add, subtract, or mult. Got AluOpType::{actual}Use 32-bit math or use one of add/subtract/mult
EBIR014NKI Kernel call references undefined KernelEnsure the NKI kernel call references a valid Kernel.
EBIR015NKI Kernel number of partitions {kernel_num_partitions} doesn't match architecture number of partitions {arch_num_partitions}. This is not supported today.Update the kernel's SBUF num partitions to match the target architecture num partitions.
EBIR016NKI Kernel has invalid SBUF shape: sb_num_partitions={sb_num_partitions}, sb_per_partition_size={sb_per_partition_size}.If SBUF is allocated (non-zero size per partition), there must be at least one partition.
EBIR017NKI Kernel PSUM bank count {psum_bank_count} exceeds architecture limit of {arch_psum_banks}.Reduce the number of PSUM banks requested by the kernel.
EBIR018NKI Kernel PSUM partition count {psum_num_partitions} exceeds architecture limit of {arch_num_partitions}.Reduce the number of PSUM partitions requested by the kernel.
EBIR019NKI Kernel has insufficient outputs: {actual_outputs} provided but {expected_min_outputs} required ({num_sb_buffers} SBUF + {num_psum_buffers} PSUM buffers).Add the required SBUF / PSUM outputs to the kernel.
EBIR020NKI Kernel has invalid address rotation scope: {scope}.Address rotation scope must be None (0), Global (1), or Kernel (2).
EBIR021Invalid architecture revision v2 on {archlevel}Revision v2 is only available on trn3
EBIR022DVE exponential is not supported on trn3preMove Exponential operation from Vector to Scalar engine, or use '--target trn3'
EBIR023MLP kernel intermediate size {intermediate_size} exceeds the maximum supported value of 4096.Consider tiling large intermediate tensors in your kernel to stay within the supported limit, or increase tensor parallelism to shard the intermediate dimension across more cores.
EBIR024Invalid dtype for SelectReduce input on_true tensor. Given: src type: {src_type}Use a supported dtype such as uint8, int8, uint16, int16, float16, bfloat16, float32, or float8 variants (float8e3, float8e4, float8_e4m3fn, float8e5)
EBIR025Invalid dtype for SelectReduce input predicates tensor. Given: predicate type: {pred_type}Use a supported dtype such as uint8, int8, uint16, int16, uint32, int32 for predicate tensors.
EBIR030InstDMACopy number of in dimensions ({in}) must equal the number of out dimensions ({out})Make sure the in/out AccessPatterns have the same number of dimensions
EBIR031Number of elements in dimension {index} of InstDMACopy's input and output AccessPatterns must match, but got in={in} and out={out}Make sure the number of elements is the same in all dimensions of the in/out AccessPatterns
EBIR032InstDMACopy number of in dimensions ({in}) must equal the number of out dimensions ({out})Make sure the in/out AccessPatterns have the same number of dimensions
EBIR033Number of elements in dimension {index} of InstDMACopy's input and output AccessPatterns must match, but got in={in} and out={out}Make sure the number of elements is the same in all dimensions of the in/out AccessPatterns

EBRD — EXTERNAL (3 codes)

CodeCauseResolution
EBRD001Mismatched number of modules {num_modules} and condensed BIR input files {num_condensed}Make sure there is 1 condensed file input for each module (only 1 module currently supported)
EBRD002Unable to open condensed format file {file_name}Ensure your condensed .bir file exists and that the full path is provided
EBRD003Unable to find instruction in condensed format fileTry regenerating your condensed format file?

EBVF — EXTERNAL (1 codes)

CodeCauseResolution
EBVF030Instructions generated by compiler {total_count} exceeds the typical limit of {max_instruction_limit}. Input computation graph is too big due to large operatorsConsider using smaller batches or sequence length, or applying tensor parellelism. For further troubleshooting visit https://awsdocs-neuron.readthedocs-hosted.com/en/latest/libraries/nxd-training/app_notes/nxd-training-tp-appnote.html

EDGE — EXTERNAL (1 codes)

CodeCauseResolution
EDGE001DMAs that use HWDGE can only be assigned to EngineType::Activation (nki.isa.engine.scalar) or EngineType::SP (nki.isa.engine.sync), but got EngineType::{actual}If you want this DMA to use HWDGE, set its engine to engine=nki.isa.engine.scalar or engine=nki.isa.engine.sync instead. If you want this DMA to be on the engine you've selected, then do not set dge_mode=dge_mode.hwdge.

EDMA — EXTERNAL (1 codes)

CodeCauseResolution
EDMA002Kernel partition size {kernel_num_partitions} does not match architecture partition size {arch_num_partitions}.Ensure the kernel's partition configuration matches the target architecture.

EDP — EXTERNAL (3 codes)

CodeCauseResolution
EDP001InstDevicePrint must have 1 argument, but it has {actual}(contact)
EDP002InstDevicePrint argument must be a PhysicalAccessPattern, but it is a {actual}(contact)
EDP003InstDevicePrint can only print a tensor in HBM or SBUF, but tensor {name} is in {type}(contact)

EDRV — EXTERNAL (1 codes)

CodeCauseResolution
EDRV080No subgraph directories (sgXX) found in NC directory: {nc_dir}(contact)

EDVR — EXTERNAL (1 codes)

CodeCauseResolution
EDVR001Transpose DMA with dynamic offset cannot exist without transpose DGE level enabledPlease add transpose to the --internal-enable-dge-levels flag

EGCA — EXTERNAL (7 codes)

CodeCauseResolution
EGCA110Memory allocation failed for {memloc}. Spilling is disabled for kernels today.Please reduce memory usage by optimizing tensor sizes or splitting operations into smaller chunks.
EGCA111Memory allocation failed for {memloc}. Spilling is disabled for kernels today.Please reduce memory usage by optimizing tensor sizes or splitting operations into smaller chunks.
EGCA112Kernel {function} requires loop unrolling for proper memory allocation in graph compiler flow.Please use new NKI flow when doing partially allocated kernel.
EGCA113Kernel {function} requires loop unrolling for proper memory allocation in graph compiler flow.Please use new NKI flow when doing partially allocated kernel.
EGCA114Conflicting partition constraints detected in the same constraint group.Please examine tensor reader/writer to ensure they do not have conflicting constraint requirements.
EGCA115If partition num > 96, tensor {tensor} base partition constraint must be 0Ensure >96 partition tensors do not have non-0 tile position.
EGCA116Conflicting base partition constraints from matmul instructions.Tensors accessed by matmuls have incompatible partition requirements based on tile positions.

EIL — EXTERNAL (11 codes)

CodeCauseResolution
EIL001InstCompareAndBranch should have exactly 2 arguments (got {actual})(contact)
EIL002InstCompareAndBranch can only handle PhysicalAccessPatterns (got a {actual})(contact)
EIL003Loop condition MemoryLocation '{name}' must be <= 4 bytes (it is {actual_size} bytes)(contact)
EIL004Loop condition MemoryLocation '{name}' must have exactly 1 element (it has {actual_size} elements)(contact)
EIL005InstCompareAndBranch cannot be on the ALL engine, it should have been expanded by expand_all_engine(contact)
EIL006InstCompareAndBranch cannot have any dependencies(contact)
EIL007InstCompareAndBranch cannot have any descendents(contact)
EIL008InstCompareAndBranch that is not expanded by expand_inst_late must not have SyncInfo(contact)
EIL009No dynamic axis register (IntRuntimeValue) found in expression with hasRuntimeValue true!(contact)
EIL010InstCompareAndBranch must only have ordered dependencies, but it has a {actual} dependency(contact)
EIL011Register within dynamic axis of instruction does not match engine of instruction!(contact)

ELMM — EXTERNAL (1 codes)

CodeCauseResolution
ELMM001Matrix Multiplication with write indirection of type bfloat16 cannot use less quadrants than previous matrix multiplicationsBefore the culprit matrix multiplication, insert a dummy matrix multiplication that covers all quadrants

ELUR — EXTERNAL (1 codes)

CodeCauseResolution
ELUR015Instructions generated by compiler {total_count} exceeds the typical limit of {max_instruction_limit}. Input computation graph is too big due to large operatorsConsider using smaller batches or sequence length, or applying tensor parellelism. For further troubleshooting visit https://awsdocs-neuron.readthedocs-hosted.com/en/latest/libraries/nxd-training/app_notes/nxd-training-tp-appnote.html

ENKI — EXTERNAL (1 codes)

CodeCauseResolution
ENKI001

EOCP — EXTERNAL (1 codes)

CodeCauseResolution
EOCP001Tensor {legacy_tensor} has legacy type {legacy_type} but tensor {ocp_tensor} has OCP-compliant type {ocp_type}. Mixed use of the two mutually-exclusive types is not supportedEnsure that only one of the two types is used.

EOOM — EXTERNAL (2 codes)

CodeCauseResolution
EOOM001Maximum peak HBM usage of {usage} (at instruction {inst}) exceeds HBM limit of {maxsize} for {arch}. This consists of {hbm_io} I/O tensors, {hbm_intermediate} intermediate tensors, and {hbm_internal} internal (scratchpad) allocations (of which {sbuf} is anticipated spills from SBUF)Check the target architecture flag to ensure you are compiling for correct target, or consider using smaller batches or applying model/tensor parallelism.
EOOM002Maximum peak HBM usage of {usage} exceeds HBM limit of {maxsize} for {arch}. This consists of {constants} model constants, {unallocated} I/O tensors, {allocated} internal (scratchpad) tensors, {dma_ring_io} DMA ring I/O, and {dma_ring_spill} DMA ring spillsCheck the target architecture flag to ensure you are compiling for correct target, or consider using smaller batches or applying model/tensor parallelism.

ERP — EXTERNAL (1 codes)

CodeCauseResolution
ERP001{inst_name}: Replication not supported for MatmultMX instructions(contact)

ESL — EXTERNAL (17 codes)

CodeCauseResolution
ESL001Expected to find a first definition for shared tensor {name} on at least one core.(contact)
ESL002Expected to find a definition before read for shared tensor {name}.(contact)
ESL003Expected shared memory definition for {name} to have a synchronization primitive before it.(contact)
ESL004Expected shared memory location for {name} on core {core_id}, but found none.(contact)
ESL005Expected core barrier ahead of shared memory definition for {name}, but none found on core {core_id}.(contact)
ESL006Expected shared memory reference for {name} to have a synchronization primitive after it.(contact)
ESL007Expected shared memory location for {name} on core {core_id}, but found none.(contact)
ESL008Expected core barrier after shared memory read for {name}, but none found on core {core_id} exists.(contact)
ESL009Expected to find {expected_count} functions in {sgId} modules for all cores, but found {actual_count} in module for core {ncId}.(contact)
ESL010Expected to find a function named {function_name} in {sgId} modules for all cores, but it does not exist in module for core {ncId}.(contact)
ESL011Expected the {function_name} function in {sgId} modules of all cores to have {expected_size} blocks, but core {ncId} has {actual_size} blocks.(contact)
ESL012Did not expect any shared internal tensors to be live on input to function {functionName} in subgraph {sgId}, but found {liveInShared} live.(contact)
ESL013Expected to find shared tensor {name}, but it is missing in function {functionName} in subgraph {sgId} on core {ncId}.(contact)
ESL014Encountered an incomplete GPSIMDSB2SB instruction(contact)
ESL015Encountered an incomplete GPSIMDSB2SB instruction(contact)
ESL016Encountered an incomplete GPSIMDSB2SB instruction(contact)
ESL017Encountered an incomplete GPSIMDSB2SB instruction(contact)

ESMP — EXTERNAL (2 codes)

CodeCauseResolution
ESMP001Constant simplification would produce negative index values. This will cause out-of-bounds access. Tensor: {tensor_name}, Op ID: {op_id}, Const values: {const_values}Review the indexing logic to ensure indices are non-negative. Check the base offset and step calculations in the index expression.
ESMP002Constant simplification would produce index values (max: {max_index}) exceeding the target tensor dimension size ({dim_size}) at dimension {dim}. This will cause out-of-bounds access. Source tensor: {tensor_name}, Target tensor: {target_tensor_name}, Op ID: {op_id}, Const values: {const_values}Review the indexing logic to ensure indices remain within the bounds of the target tensor dimension being indexed. Check that the index values are appropriate for the tensor shape at the specified dimension.

ETST — EXTERNAL (1 codes)

CodeCauseResolution
ETST001Example external error message 1.(contact)

EXSP — EXTERNAL (1 codes)

CodeCauseResolution
EXSP001Size of HBM memory required for the model exceeds HBM limit of {hardware_target}. Needed more than {usage_in_comma} ({usage_in_gb}GB) vs. available bytes {threshold_in_comma} ({threshold_in_gb}GB). This may be due to compiler allocating high scratch pad memory for this model.Check the target architecture flag to ensure you are compiling for correct target, or consider using smaller batches or applying model/tensor parallelism.

EXTP — EXTERNAL (2 codes)

CodeCauseResolution
EXTP003Instructions generated by compiler {insts} exceeds the typical limit of {threshold}. Input computation graph is too big due to large operatorsConsider using smaller batches or sequence length, or applying tensor parellelism. For further troubleshooting visit https://awsdocs-neuron.readthedocs-hosted.com/en/latest/libraries/nxd-training/app_notes/nxd-training-tp-appnote.html
EXTP004Instructions generated by compiler {insts} exceeds the typical limit of {threshold}. Input computation graph is too big due to a large number of operatorsConsider using compiler options --optlevel=1, or applying pipeline parallelism. For further troubleshooting visit https://awsdocs-neuron.readthedocs-hosted.com/en/latest/libraries/nxd-training/app_notes/nxd-training-pp-appnote.html.

GCA — bare (89 codes)

CodeCauseResolution
GCA013couldn't allocate every tensor in PSUM, likely due to illegal accumulation group !(contact)
GCA014Spill insertion at instruction {inst}, illegal accumulation group!(contact)
GCA022DRAM usage for Internal DRAM tensor exceeds {hbm_limit}GB of device space limit, cannot fit into device, model requires too much HBM memory !(contact)
GCA023Allocation of {loc} crossing runtime page boundary with address range of [{long_address},{end_address})(contact)
GCA024Illegal IR, encountered undefined use: {first_bad_use}!(contact)
GCA025Allocator bit_matrix and adjacency vectors require {required_mem} bytes, which is more than 64GB limit(contact)
GCA026Unrecognized architecture: {arch}(contact)
GCA027Expected reload to PSUM(contact)
GCA028Expected reload to PSUM(contact)
GCA029Unrecognized architecture: {arch}(contact)
GCA030{error_message}(contact)
GCA031{error_message}(contact)
GCA032Illegal tensor height(contact)
GCA034Illegal tensor height(contact)
GCA035{error_message}(contact)
GCA036Unexpected value for ml_base = {ml_base}(contact)
GCA037{node} is too big for SB, requires {bytes_per_partition} bytes(contact)
GCA038The state buffer: {stateBuffer}\nin sb_select, a bad allocation\n {node_index} {node_name} [{first_address} {last_address}) x [{first_partition} {last_partition})\n {neighbor} {neighbor_name} [{firsta} {lasta}) x [{firstp} {lastp})\n(contact)
GCA039In sb_select, a call to allocate() failed, in {memloc_name}(contact)
GCA040In sb_select, a call to allocate() failed, in {memloc_name}(contact)
GCA041Illegal tensor height(contact)
GCA042Expected reload to SB(contact)
GCA043Unexpected value for ml_base = {ml_base}(contact)
GCA044{error_message}(contact)
GCA045Couldn't color the DRAM even with 100GB of DRAM space assumption, model needs too much HBM memory !(contact)
GCA046Some infinite-cost nodes remain = {size}(contact)
GCA047Cannot spill must-be-pinned or pre-allocated tensor, spilled tensor = {tensor_name}(contact)
GCA048Instruction must have exactly 1 argument(contact)
GCA049Instruction must have exactly 1 output(contact)
GCA050Unexpected SB allocator run after linking(contact)
GCA051Unexpected address space {addr_space} for SB allocator(contact)
GCA052Unexpected PSUM allocator run after linking(contact)
GCA053Unexpected address space {addr_space} for PSUM allocator(contact)
GCA054Illegal IR, attempt to use improperly defined tensor(contact)
GCA055Illegal IR, encountered undefined use: {first_bad_use}, and unused define: {first_bad_def}!(contact)
GCA056Illegal IR, encountered unused define: {first_bad_def}!(contact)
GCA057Identity matrix of dtype {dtype} not found(contact)
GCA058Could not find memloc {loc_name} in accumulation group cache(contact)
GCA059Could not find next use(contact)
GCA060Next use instruction {instruction_name} has not been visited(contact)
GCA061Illegal first use of internal tensor(contact)
GCA062Could not find next use(contact)
GCA063Next use instruction {instruction_name} has not been visited(contact)
GCA064Illegal first use of internal tensor(contact)
GCA065Shouldn't need initial loads with a single basic block(contact)
GCA066Cannot have both transpose MM and MM write to the same memloc(contact)
GCA067Could not find first accumulation group inst of memloc(contact)
GCA068Tried to insert PSUM spill at end of block(contact)
GCA069Illegal PSUM spilling instruction type for instruction {spill_instruction_name}(contact)
GCA070Illegal PSUM spilling instruction type for instruction {spill_instruction_name}(contact)
GCA071Tried to insert PSUM spill at end of block(contact)
GCA072Tried to insert reload to PSUM at end of block(contact)
GCA073Expected identity matrix to be in DRAM but got {mem_type}(contact)
GCA074Cannot have both transpose MM and MM write to the same memloc(contact)
GCA075Expected a 128x128 identity matrix but got total size {size_in_elements}(contact)
GCA076Expected reload to PSUM(contact)
GCA077Identity matrix is too small. Need num_partitions at least {lhs_mm_num_partitions} but got {identity_sb_num_partitions}(contact)
GCA078{memloc_name}: Allocated base partition mismatch in MatMult output location(contact)
GCA079{inst_name}: Allocated base partition + access pattern must be <= PE array size(contact)
GCA080{memloc_name}: PSUM memory location accessed partitions must be > 0(contact)
GCA081{memloc_name}: PSUM memory location cannot have uninitialized DebugInfo(contact)
GCA082{memloc_name}: location must have valid def/use pair(contact)
GCA083{memloc_name}: location has no access in current BB. Need at least one access if location is only in live-in(contact)
GCA084{memloc_name}: location has no access in current BB. Need at least one access if location is only in live-out(contact)
GCA085{memloc_name}: location must have valid def/use pair(contact)
GCA086{memloc_name}: find no live interval for global location(contact)
GCA087{memloc_name}: location with only use and no def at allocation stage(contact)
GCA088{inst_name}: Requesting more PSUM memory than hardware constraint(contact)
GCA089Live-in registers are only supported in legacy loops-on-chip flow. Register {reg_name} defined by Instruction {def_name} in BasicBlock {def_block_name} is last used by Instruction {use_name} in BasicBlock {use_block_name}(contact)
GCA090{memloc_name}: no def/use pair find for memorylocation(contact)
GCA091{memloc_name}: Currently cannot spill multiple-BB access memorylocation(contact)
GCA092{memloc_name}: Currently cannot spill multiple-BB access memorylocation(contact)
GCA093{memloc_name}: location has no definition(contact)
GCA094Unable to allocate register '{reg}' on engine '{engine}' because there are too many registers live concurrently. Register spill/reload is not supported(contact)
GCA095{reg_name}: register with only use and no def at allocation stage(contact)
GCA096{reg_name}: register has no definition(contact)
GCA097{reg_name}: register has no access in current BB. Need at least one access if register is only in live-in(contact)
GCA098{reg_name}: register has no access in current BB. Need at least one access if register is only in live-out(contact)
GCA099{reg_name}: find no live interval for global register(contact)
GCA100{reg_name}: register must have valid def/use pair(contact)
GCA101Must have valid first_def and last_use. {error_str}(contact)
GCA102{bb_name}: Strongly connected graph cannot be empty(contact)
GCA103{memloc_name}: Cannot spill memorylocation defined by DMA skipping inside of loop, without initial define outside of loop. Please do a dummy define outside of loop body for this tensor(contact)
GCA104No def/use pair found for register '{reg_name}'(contact)
GCA105Register '{reg_name}' must have valid def/use pair(contact)
GCA106Each register must have at least one live interval. There are {num_regs} registers but only {num_intervals} intervals.(contact)
GCA107GCA reload instruction can only be inserted in basicblock before next user(contact)
GCA108{memloc_name}:Tensor with loop-carried depednency has to be defined outside of loop(contact)
GCA109Infinite memloc picked to spill, cannot allocate. memloc name {mam_name}. This is due to single instruction accessing memlocs needing more than total PSUM size.(contact)

IACF — INTERNAL (2 codes)

CodeCauseResolution
IACF901AutoCastFP32 assertion error: {baseMessage}(contact)
IACF902AutoCastFP32 error: {baseMessage}(contact)

IACI — INTERNAL (2 codes)

CodeCauseResolution
IACI901AutoCastInputs assertion error: {baseMessage}(contact)
IACI902AutoCastInputs error: {baseMessage}(contact)

IACT — INTERNAL (14 codes)

CodeCauseResolution
IACT001Expect that the start time of previous activation instruction {prevInst} to be less than or equal to the current activation instruction {currInst}, but it is not.(contact)
IACT002The activation table load move candidate should not be nullptr(contact)
IACT003The activation table load instruction should not be nullptr(contact)
IACT005
IACT006
IACT007
IACT008
IACT009
IACT010
IACT011
IACT012
IACT013
IACT901AutoCastTCInputs assertion error: {baseMessage}(contact)
IACT902AutoCastTCInputs error: {baseMessage}(contact)

IADE — INTERNAL (2 codes)

CodeCauseResolution
IADE901AliasDependencyElimination assertion error: {baseMessage}(contact)
IADE902AliasDependencyElimination error: {baseMessage}(contact)

IADI — INTERNAL (2 codes)

CodeCauseResolution
IADI901AliasDependencyInduction assertion error: {baseMessage}(contact)
IADI902AliasDependencyInduction error: {baseMessage}(contact)

IADQ — INTERNAL (2 codes)

CodeCauseResolution
IADQ901AssignDMAQoSLabels assertion error: {baseMessage}(contact)
IADQ902AssignDMAQoSLabels error: {baseMessage}(contact)

IADR — INTERNAL (2 codes)

CodeCauseResolution
IADR901AliasDependencyReset assertion error: {baseMessage}(contact)
IADR902AliasDependencyReset error: {baseMessage}(contact)

IADV — INTERNAL (2 codes)

CodeCauseResolution
IADV901AliasDependencyVerificationPass assertion error: {baseMessage}(contact)
IADV902AliasDependencyVerificationPass error: {baseMessage}(contact)

IAGO — INTERNAL (2 codes)

CodeCauseResolution
IAGO901AGOrderingAnalysis assertion error: {baseMessage}(contact)
IAGO902AGOrderingAnalysis error: {baseMessage}(contact)

IAHE — INTERNAL (1 codes)

CodeCauseResolution
IAHE001

IALB — INTERNAL (2 codes)

CodeCauseResolution
IALB901AllocateBlocks assertion error: {baseMessage}(contact)
IALB902AllocateBlocks error: {baseMessage}(contact)

IALS — INTERNAL (2 codes)

CodeCauseResolution
IALS001
IALS002

IANK — INTERNAL (2 codes)

CodeCauseResolution
IANK901AnalyzeKernel assertion error: {baseMessage}(contact)
IANK902AnalyzeKernel error: {baseMessage}(contact)

IANS — INTERNAL (2 codes)

CodeCauseResolution
IANS901AnnotateNoSpill assertion error: {baseMessage}(contact)
IANS902AnnotateNoSpill error: {baseMessage}(contact)

IAPR — INTERNAL (2 codes)

CodeCauseResolution
IAPR901AffinePredicateResolution assertion error: {baseMessage}(contact)
IAPR902AffinePredicateResolution error: {baseMessage}(contact)

IATL — INTERNAL (40 codes)

CodeCauseResolution
IATL001
IATL002
IATL003
IATL004
IATL005
IATL006
IATL007
IATL008
IATL009
IATL010
IATL011
IATL012
IATL013
IATL014
IATL015
IATL016
IATL017
IATL018
IATL019
IATL020
IATL021
IATL022
IATL023
IATL024
IATL025
IATL026
IATL027
IATL028
IATL029
IATL030
IATL031
IATL032
IATL033
IATL034
IATL035
IATL036
IATL037
IATL038
IATL039
IATL040

IATN — INTERNAL (2 codes)

CodeCauseResolution
IATN901Autotuner assertion error: {baseMessage}(contact)
IATN902Autotuner error: {baseMessage}(contact)

IBCC — INTERNAL (2 codes)

CodeCauseResolution
IBCC901BucketizeCCOp assertion error: {baseMessage}(contact)
IBCC902BucketizeCCOp error: {baseMessage}(contact)

IBCG — INTERNAL (2 codes)

CodeCauseResolution
IBCG901BIRCodeGenLoop assertion error: {baseMessage}(contact)
IBCG902BIRCodeGenLoop error: {baseMessage}(contact)

IBFC — INTERNAL (2 codes)

CodeCauseResolution
IBFC901BFComputeCutting assertion error: {baseMessage}(contact)
IBFC902BFComputeCutting error: {baseMessage}(contact)

IBIR — INTERNAL (55 codes)

CodeCauseResolution
IBIR094All Matmult tiling parameters must have matching sizes. Given: Size={row_size}x{col_size}, Position={row_pos}x{col_pos}Adjust Matmult inputs to have matching expressions.
IBIR158Access pattern out of bounds on instruction '{inst}'. Pattern: {pattern}(contact)
IBIR238Instruction Activate2 is supported only for trn3 architecture. Got architecture: {arch}(contact)
IBIR239Activate2 params imm0, dtype: {imm0_dtype}; imm1, dtype: {imm1_dtype}; relu_param, dtype: {relu_param_dtype} have mismatching dtypes(contact)
IBIR240Invalid operator combination ({op0}, {op1}) for Activate2. Valid operator params are (mult,add), (mult, subtract), (mult, bypass), (add, bypass), (subtract, bypass), (bypass, bypass)(contact)
IBIR241Invalid accumulation type: {acc}. Must be one of Idle, Zero, AddAccumulate or ZeroAccumulate(contact)
IBIR242Expected src_tensor (quadrant: {src_quadrant}), imm0 (quadrant: {imm0_quadrant}), imm1 (quadrant: {imm1_quadrant}) and relu_param (quadrant: {relu_param_quadrant}) to be in the same sbuf quadrant(contact)
IBIR243Access pattern out of bounds. Pattern: {pattern}(contact)
IBIR246
IBIR265
IBIR266
IBIR267
IBIR269
IBIR270
IBIR272
IBIR273
IBIR274
IBIR275
IBIR276
IBIR277
IBIR279
IBIR280Function {func} requires at least one BasicBlock.(contact)
IBIR281Register '{name}' which uses {num} physical 32-bit registers must have its lowest physical register allocated to an even ID, but it is '{actual}'(contact)
IBIR282
IBIR283DMA Transpose is not allowed for MX data types.Consider using another data type instead.
IBIR284DMA non-bypass CCE is not allowed for MX data types.Consider using another data type instead.
IBIR285{instruction_type_name} instructions don't support data type casting for MX data types (src data type: {in_dtype}, dst data dtype: {out_dtype}).Consider using another data type instead.
IBIR286MX data type is not supported for compute instructions other than QuantizeMx and MatmultMx but got {instruction_type_name}.Consider using another data type instead.
IBIR289
IBIR290bir::PhysicalAccessPattern::getAddresses forbids setting startingAddress >= 0 and perPartition = false!(contact)
IBIR291PE Tile position and size must both be fully specified(contact)
IBIR292
IBIR302Instruction has invalid tile group row position {tileGroupRowPosition}, should be one of 0, 32, 64, 96(contact)
IBIR303Instruction has invalid tile group row position {tileGroupColPosition}, should be one of 0, 32, 64, 96(contact)
IBIR304Instruction has invalid row tile position ({tileGroupRowPosition}) and tile size ({tileGroupRowSize}) specified, expect tile size to be one of {possibleTileSizes} at tile position(contact)
IBIR305Instruction has invalid col tile position ({tileGroupColPosition}) and tile size ({tileGroupColSize}) specified, expect tile size to be one of {possibleTileSizes} at tile position(contact)
IBIR308
IBIR313
IBIR314
IBIR315
IBIR316
IBIR317
IBIR318
IBIR319
IBIR320
IBIR322
IBIR351
IBIR352
IBIR353
IBIR354
IBIR355
IBIR356Unsupported expression kind ({kind}) in visitor(contact)
IBIR357
IBIR366
IBIR367

IBRD — INTERNAL (4 codes)

CodeCauseResolution
IBRD001Duplicate instruction name {inst_name} found in condensed format fileChanges in condensed format dumping or parsing are probably needed
IBRD002Instruction has no source linePossible issue in source map loading
IBRD003Unreachable debugger stop reason in printUpdate method to handle new reasons
IBRD004Core index out of boundsCheck core index is within valid range

IBRW — INTERNAL (2 codes)

CodeCauseResolution
IBRW901BroadcastWeights assertion error: {baseMessage}(contact)
IBRW902BroadcastWeights error: {baseMessage}(contact)

IBSN — INTERNAL (15 codes)

CodeCauseResolution
IBSN001PoisonTracker::mergePoison argument [const boost::icl::interval_set<uint64_t>& boundingIntervals] must fully contain argument [const boost::icl::interval_set<uint64_t>& poisonedIntervals](contact)
IBSN002PoisonTracker::checkPoisonedBytesNotNullptr assertion failed! Called from {callee}(contact)
IBSN008birsan::InstVisitor::elementToElementMap expects the same number of elements in the inputAP and the outputAP (got {inputNumElements} != {outputNumElements})(contact)
IBSN009birsan::InstVisitor::transformReduction produced and invalid outputIdx from inputIdxToOutputIdxs (inputIdxToOutputIdxs({inputIdx}) = {outputIdx} >= {outputIdxMax})(contact)
IBSN010birsan::InstVisitor::elementToElementByteCopy produced and invalid inputIdx from outputIdxToInputIdxMap (outputIdxToInputIdxMap[{outputIdx}] = {inputIdx} >= {inputIdxMax})(contact)
IBSN011birsan::InstVisitor::transformSelect expected predsAddresses.size() == dstAddresses.size(), but got {predsSize} != {dstSize}(contact)
IBSN012birsan::InstVisitor::handleDMAVectorIndirect does not support both vectorSrcIndirect and vectorDstIndirect(contact)
IBSN013birsan::InstVisitor::handleDMATranspose found InstDMACopy without bir::CopyMode::Transpose(contact)
IBSN014birsan::InstVisitor::handleDMATranspose only handles InstDMACopy with bir::CopyMode::Transpose and InstDMADescriptorTranspose(contact)
IBSN015birsan::InstVisitor::handleDMAReplicate found InstDMACopy without bir::CopyMode::Replicate(contact)
IBSN016birsan::InstVisitor::handleDMAReplicate only handles InstDMACopy with bir::CopyMode::Replicate and InstDMADescriptorReplicate(contact)
IBSN017birsan::InstVisitor::handleDMAReplicate requires input/output APs to have two dimensions(contact)
IBSN018birsan::InstVisitor::handleDMACCE found InstDMACopy without bir::CopyMode::CCE(contact)
IBSN019birsan::InstVisitor::handleDMACCE only handles InstDMACopy with bir::CopyMode::CCE and InstDMADescriptorCCE(contact)
IBSN020birsan::InstVisitor::transformBasicBlockHolder expects a bir::InstructionBasicBlockHolder(contact)

IBVF — INTERNAL (9 codes)

CodeCauseResolution
IBVF032The bias argument must be of type float32 when activation function type is of Copy or Reciprocal, but is {dtype}Check that the activation function type and bias arugment dtype
IBVF033The bias argument must be of type float32 on CoreV2+, but is {dtype}Check that the activation instruction bias arugment dtype
IBVF034
IBVF036The bias argument must be scalar or of shape at least N ({input_partitions}) x 1 but has shape {shape}Check the activation instruction bias argument shape
IBVF037The bias argument tensor may not be indirectUse a non-indirect tensor for bias
IBVF038Input and output dtypes must be valid and the same for non-bitvec operations: {inputDtype} != {outputDtype}(contact)
IBVF039
IBVF040
IBVF041

ICCC — INTERNAL (2 codes)

CodeCauseResolution
ICCC901CoalesceCCOp assertion error: {baseMessage}(contact)
ICCC902CoalesceCCOp error: {baseMessage}(contact)

ICCF — INTERNAL (2 codes)

CodeCauseResolution
ICCF901CCOpFusion assertion error: {baseMessage}(contact)
ICCF902CCOpFusion error: {baseMessage}(contact)

ICDG — INTERNAL (2 codes)

CodeCauseResolution
ICDG901CanonicalizeDAG assertion error: {baseMessage}(contact)
ICDG902CanonicalizeDAG error: {baseMessage}(contact)

ICDL — INTERNAL (2 codes)

CodeCauseResolution
ICDL901ConcatDelinearizer assertion error: {baseMessage}(contact)
ICDL902ConcatDelinearizer error: {baseMessage}(contact)

ICIO — INTERNAL (2 codes)

CodeCauseResolution
ICIO901ConvertIOBufferToMustAlias assertion error: {baseMessage}(contact)
ICIO902ConvertIOBufferToMustAlias error: {baseMessage}(contact)

ICIR — INTERNAL (2 codes)

CodeCauseResolution
ICIR901CanonicalizeIR assertion error: {baseMessage}(contact)
ICIR902CanonicalizeIR error: {baseMessage}(contact)

ICMA — INTERNAL (2 codes)

CodeCauseResolution
ICMA901ConvertMustAliasToIOBuffer assertion error: {baseMessage}(contact)
ICMA902ConvertMustAliasToIOBuffer error: {baseMessage}(contact)

ICMC — INTERNAL (2 codes)

CodeCauseResolution
ICMC901CommuteConcat assertion error: {baseMessage}(contact)
ICMC902CommuteConcat error: {baseMessage}(contact)

ICNE — INTERNAL (2 codes)

CodeCauseResolution
ICNE901ConcatOpElimination assertion error: {baseMessage}(contact)
ICNE902ConcatOpElimination error: {baseMessage}(contact)

ICOE — INTERNAL (2 codes)

CodeCauseResolution
ICOE901CastOpElimination assertion error: {baseMessage}(contact)
ICOE902CastOpElimination error: {baseMessage}(contact)

IDCE — INTERNAL (2 codes)

CodeCauseResolution
IDCE901DeadCodeElimination assertion error: {baseMessage}(contact)
IDCE902DeadCodeElimination error: {baseMessage}(contact)

IDDI — INTERNAL (2 codes)

CodeCauseResolution
IDDI901DetailDynInst assertion error: {baseMessage}(contact)
IDDI902DetailDynInst error: {baseMessage}(contact)

IDDT — INTERNAL (2 codes)

CodeCauseResolution
IDDT901DramToDramTranspose assertion error: {baseMessage}(contact)
IDDT902DramToDramTranspose error: {baseMessage}(contact)

IDEC — INTERNAL (2 codes)

CodeCauseResolution
IDEC901DeConcat assertion error: {baseMessage}(contact)
IDEC902DeConcat error: {baseMessage}(contact)

IDEL — INTERNAL (2 codes)

CodeCauseResolution
IDEL901Delinearization assertion error: {baseMessage}(contact)
IDEL902Delinearization error: {baseMessage}(contact)

IDFS — INTERNAL (2 codes)

CodeCauseResolution
IDFS901DFSStackAllocator assertion error: {baseMessage}(contact)
IDFS902DFSStackAllocator error: {baseMessage}(contact)

IDFV — INTERNAL (2 codes)

CodeCauseResolution
IDFV901DataflowViewer assertion error: {baseMessage}(contact)
IDFV902DataflowViewer error: {baseMessage}(contact)

IDGE — INTERNAL (4 codes)

CodeCauseResolution
IDGE001
IDGE002
IDGE003
IDGE004

IDGM — INTERNAL (2 codes)

CodeCauseResolution
IDGM901DumpGraphAndMetadata assertion error: {baseMessage}(contact)
IDGM902DumpGraphAndMetadata error: {baseMessage}(contact)

IDIE — INTERNAL (2 codes)

CodeCauseResolution
IDIE901DynamicInstEstimator assertion error: {baseMessage}(contact)
IDIE902DynamicInstEstimator error: {baseMessage}(contact)

IDLI — INTERNAL (2 codes)

CodeCauseResolution
IDLI901DelinearIndices assertion error: {baseMessage}(contact)
IDLI902DelinearIndices error: {baseMessage}(contact)

IDLO — INTERNAL (2 codes)

CodeCauseResolution
IDLO901DataLocalityOpt assertion error: {baseMessage}(contact)
IDLO902DataLocalityOpt error: {baseMessage}(contact)

IDMA — INTERNAL (4 codes)

CodeCauseResolution
IDMA001
IDMA002
IDMA003
IDMA004

IDML — INTERNAL (2 codes)

CodeCauseResolution
IDML901DMALocalityOpt assertion error: {baseMessage}(contact)
IDML902DMALocalityOpt error: {baseMessage}(contact)

IDMP — INTERNAL (2 codes)

CodeCauseResolution
IDMP901DMAProfiler assertion error: {baseMessage}(contact)
IDMP902DMAProfiler error: {baseMessage}(contact)

IDNT — INTERNAL (2 codes)

CodeCauseResolution
IDNT901DoNothing assertion error: {baseMessage}(contact)
IDNT902DoNothing error: {baseMessage}(contact)

IDSE — INTERNAL (2 codes)

CodeCauseResolution
IDSE901DeadStoreElimination assertion error: {baseMessage}(contact)
IDSE902DeadStoreElimination error: {baseMessage}(contact)

IDSH — INTERNAL (2 codes)

CodeCauseResolution
IDSH901DimensionSharding assertion error: {baseMessage}(contact)
IDSH902DimensionSharding error: {baseMessage}(contact)

IDST — INTERNAL (2 codes)

CodeCauseResolution
IDST901DataStreaming assertion error: {baseMessage}(contact)
IDST902DataStreaming error: {baseMessage}(contact)

IDTP — INTERNAL (2 codes)

CodeCauseResolution
IDTP901DMATilingProfiler assertion error: {baseMessage}(contact)
IDTP902DMATilingProfiler error: {baseMessage}(contact)

IDVR — INTERNAL (1 codes)

CodeCauseResolution
IDVR004

IEAD — INTERNAL (3 codes)

CodeCauseResolution
IEAD001EnforceAluDTAcc: promoting tensor {tensor_name} from {src_dtype} to {dst_dtype} would overflow SB partition ({promoted_bytes} > {limit_bytes} bytes)Remove the --accumulate-on-alu-dtype flag to disable this optimization
IEAD901EnforceAluDTAcc assertion error: {baseMessage}(contact)
IEAD902EnforceAluDTAcc error: {baseMessage}(contact)

IEBN — INTERNAL (2 codes)

CodeCauseResolution
IEBN901ExpandBatchNorm assertion error: {baseMessage}(contact)
IEBN902ExpandBatchNorm error: {baseMessage}(contact)

IEDV — INTERNAL (2 codes)

CodeCauseResolution
IEDV901EliminateDivs assertion error: {baseMessage}(contact)
IEDV902EliminateDivs error: {baseMessage}(contact)

IEID — INTERNAL (4 codes)

CodeCauseResolution
IEID001
IEID002
IEID003
IEID004

IEIL — INTERNAL (2 codes)

CodeCauseResolution
IEIL012
IEIL013

IEIM — INTERNAL (2 codes)

CodeCauseResolution
IEIM901ExpandISAMacro assertion error: {baseMessage}(contact)
IEIM902ExpandISAMacro error: {baseMessage}(contact)

IFAT — INTERNAL (2 codes)

CodeCauseResolution
IFAT901FlattenAxesForTiling assertion error: {baseMessage}(contact)
IFAT902FlattenAxesForTiling error: {baseMessage}(contact)

IFBD — INTERNAL (2 codes)

CodeCauseResolution
IFBD901FactorizeBlkDims assertion error: {baseMessage}(contact)
IFBD902FactorizeBlkDims error: {baseMessage}(contact)

IFFD — INTERNAL (2 codes)

CodeCauseResolution
IFFD901FactorizeFreeDims assertion error: {baseMessage}(contact)
IFFD902FactorizeFreeDims error: {baseMessage}(contact)

IFGC — INTERNAL (2 codes)

CodeCauseResolution
IFGC901FineGrainedCCOpFusion assertion error: {baseMessage}(contact)
IFGC902FineGrainedCCOpFusion error: {baseMessage}(contact)

IFML — INTERNAL (2 codes)

CodeCauseResolution
IFML901FlattenMacroLoop assertion error: {baseMessage}(contact)
IFML902FlattenMacroLoop error: {baseMessage}(contact)

IFSG — INTERNAL (2 codes)

CodeCauseResolution
IFSG901FastSpillGeneration assertion error: {baseMessage}(contact)
IFSG902FastSpillGeneration error: {baseMessage}(contact)

IFTA — INTERNAL (2 codes)

CodeCauseResolution
IFTA901FactorizeThreadAxesInFreeDims assertion error: {baseMessage}(contact)
IFTA902FactorizeThreadAxesInFreeDims error: {baseMessage}(contact)

IGAS — INTERNAL (2 codes)

CodeCauseResolution
IGAS901GenericAccessSimplifier assertion error: {baseMessage}(contact)
IGAS902GenericAccessSimplifier error: {baseMessage}(contact)

IGCA — INTERNAL (7 codes)

CodeCauseResolution
IGCA001Pre-allocated or already allocated tensor {tensor_name} should not be in the allocation stack during DRAM allocation.(contact)
IGCA003
IGCA004Constraint group has mixed constrained and unconstrained tensors.(contact)
IGCA117Conflicting AP-level partition constraints detected within a constraint group.(contact)
IGCA118Conflicting hard partition constraints in a constraint group.(contact)
IGCA119No valid quadrant-aligned base partition assignment exists for a constraint group.(contact)
IGCA120Cannot assign computed base partition to tensor.(contact)

IGLO — INTERNAL (2 codes)

CodeCauseResolution
IGLO901GlobalLayoutOpt assertion error: {baseMessage}(contact)
IGLO902GlobalLayoutOpt error: {baseMessage}(contact)

IHAG — INTERNAL (2 codes)

CodeCauseResolution
IHAG901HoistAllGather assertion error: {baseMessage}(contact)
IHAG902HoistAllGather error: {baseMessage}(contact)

IHFC — INTERNAL (2 codes)

CodeCauseResolution
IHFC901HoistFSDPCollectives assertion error: {baseMessage}(contact)
IHFC902HoistFSDPCollectives error: {baseMessage}(contact)

IICB — INTERNAL (2 codes)

CodeCauseResolution
IICB901InsertCoreBarrier assertion error: {baseMessage}(contact)
IICB902InsertCoreBarrier error: {baseMessage}(contact)

IICR — INTERNAL (2 codes)

CodeCauseResolution
IICR901InsertConflictResolutionOps assertion error: {baseMessage}(contact)
IICR902InsertConflictResolutionOps error: {baseMessage}(contact)

IIFG — INTERNAL (2 codes)

CodeCauseResolution
IIFG901InlineFusionGroup assertion error: {baseMessage}(contact)
IIFG902InlineFusionGroup error: {baseMessage}(contact)

IIGL — INTERNAL (2 codes)

CodeCauseResolution
IIGL901IPGlobalLayoutOpt assertion error: {baseMessage}(contact)
IIGL902IPGlobalLayoutOpt error: {baseMessage}(contact)

IIIC — INTERNAL (2 codes)

CodeCauseResolution
IIIC901InferIntrinsicOnCC assertion error: {baseMessage}(contact)
IIIC902InferIntrinsicOnCC error: {baseMessage}(contact)

IIIN — INTERNAL (2 codes)

CodeCauseResolution
IIIN901InferIntrinsic assertion error: {baseMessage}(contact)
IIIN902InferIntrinsic error: {baseMessage}(contact)

IIIT — INTERNAL (2 codes)

CodeCauseResolution
IIIT901InsertIOTransposes assertion error: {baseMessage}(contact)
IIIT902InsertIOTransposes error: {baseMessage}(contact)

IIIV — INTERNAL (2 codes)

CodeCauseResolution
IIIV901InferInitValue assertion error: {baseMessage}(contact)
IIIV902InferInitValue error: {baseMessage}(contact)

IILT — INTERNAL (2 codes)

CodeCauseResolution
IILT901InsertLocalTransposes assertion error: {baseMessage}(contact)
IILT902InsertLocalTransposes error: {baseMessage}(contact)

IINK — INTERNAL (2 codes)

CodeCauseResolution
IINK901InlineNativeKernels assertion error: {baseMessage}(contact)
IINK902InlineNativeKernels error: {baseMessage}(contact)

IINL — INTERNAL (2 codes)

CodeCauseResolution
IINL901InferNonlocalTensors assertion error: {baseMessage}(contact)
IINL902InferNonlocalTensors error: {baseMessage}(contact)

IINT — INTERNAL (2 codes)

CodeCauseResolution
IINT901InferNeuronTensor assertion error: {baseMessage}(contact)
IINT902InferNeuronTensor error: {baseMessage}(contact)

IIOT — INTERNAL (2 codes)

CodeCauseResolution
IIOT901InsertOfflaodedTransposes assertion error: {baseMessage}(contact)
IIOT902InsertOfflaodedTransposes error: {baseMessage}(contact)

IIPS — INTERNAL (2 codes)

CodeCauseResolution
IIPS901IPSimplifier assertion error: {baseMessage}(contact)
IIPS902IPSimplifier error: {baseMessage}(contact)

IISA — INTERNAL (2 codes)

CodeCauseResolution
IISA901InferShardAxis assertion error: {baseMessage}(contact)
IISA902InferShardAxis error: {baseMessage}(contact)

IISM — INTERNAL (2 codes)

CodeCauseResolution
IISM901InferSharedMemLoc assertion error: {baseMessage}(contact)
IISM902InferSharedMemLoc error: {baseMessage}(contact)

IIST — INTERNAL (2 codes)

CodeCauseResolution
IIST901IPSubgraphTensorAnalysis assertion error: {baseMessage}(contact)
IIST902IPSubgraphTensorAnalysis error: {baseMessage}(contact)

ILBR — INTERNAL (2 codes)

CodeCauseResolution
ILBR901LowerBroadcast assertion error: {baseMessage}(contact)
ILBR902LowerBroadcast error: {baseMessage}(contact)

ILCB — INTERNAL (2 codes)

CodeCauseResolution
ILCB901LowerCCOpBlockAxis assertion error: {baseMessage}(contact)
ILCB902LowerCCOpBlockAxis error: {baseMessage}(contact)

ILCC — INTERNAL (2 codes)

CodeCauseResolution
ILCC901LegalizeCCOpLayout assertion error: {baseMessage}(contact)
ILCC902LegalizeCCOpLayout error: {baseMessage}(contact)

ILCM — INTERNAL (2 codes)

CodeCauseResolution
ILCM901LICM assertion error: {baseMessage}(contact)
ILCM902LICM error: {baseMessage}(contact)

ILCX — INTERNAL (2 codes)

CodeCauseResolution
ILCX901LowerComplexBroadcast assertion error: {baseMessage}(contact)
ILCX902LowerComplexBroadcast error: {baseMessage}(contact)

ILDM — INTERNAL (1 codes)

CodeCauseResolution
ILDM002

ILFD — INTERNAL (2 codes)

CodeCauseResolution
ILFD901LinearizeFreeDim assertion error: {baseMessage}(contact)
ILFD902LinearizeFreeDim error: {baseMessage}(contact)

ILFU — INTERNAL (2 codes)

CodeCauseResolution
ILFU901LoopFusion assertion error: {baseMessage}(contact)
ILFU902LoopFusion error: {baseMessage}(contact)

ILIN — INTERNAL (2 codes)

CodeCauseResolution
ILIN901LowerIntrinsics assertion error: {baseMessage}(contact)
ILIN902LowerIntrinsics error: {baseMessage}(contact)

ILIS — INTERNAL (2 codes)

CodeCauseResolution
ILIS901LegalizeInstSize assertion error: {baseMessage}(contact)
ILIS902LegalizeInstSize error: {baseMessage}(contact)

ILKK — INTERNAL (3 codes)

CodeCauseResolution
ILKK003
ILKK004
ILKK005

ILLI — INTERNAL (2 codes)

CodeCauseResolution
ILLI901LateLegalizeInst assertion error: {baseMessage}(contact)
ILLI902LateLegalizeInst error: {baseMessage}(contact)

ILLO — INTERNAL (2 codes)

CodeCauseResolution
ILLO901LateLowerTensorOp assertion error: {baseMessage}(contact)
ILLO902LateLowerTensorOp error: {baseMessage}(contact)

ILLP — INTERNAL (2 codes)

CodeCauseResolution
ILLP901LateLegalizePostSplit assertion error: {baseMessage}(contact)
ILLP902LateLegalizePostSplit error: {baseMessage}(contact)

ILLR — INTERNAL (2 codes)

CodeCauseResolution
ILLR901LateLowerReshapeOp assertion error: {baseMessage}(contact)
ILLR902LateLowerReshapeOp error: {baseMessage}(contact)

ILLT — INTERNAL (2 codes)

CodeCauseResolution
ILLT901LocalLegalizeType assertion error: {baseMessage}(contact)
ILLT902LocalLegalizeType error: {baseMessage}(contact)

ILMM — INTERNAL (3 codes)

CodeCauseResolution
ILMM001
ILMM002
ILMM003

ILNI — INTERNAL (2 codes)

CodeCauseResolution
ILNI901LateNeuronInstComb assertion error: {baseMessage}(contact)
ILNI902LateNeuronInstComb error: {baseMessage}(contact)

ILNK — INTERNAL (4 codes)

CodeCauseResolution
ILNK002
ILNK003
ILNK004
ILNK005

ILNM — INTERNAL (2 codes)

CodeCauseResolution
ILNM901LegalizeNeuronMacro assertion error: {baseMessage}(contact)
ILNM902LegalizeNeuronMacro error: {baseMessage}(contact)

ILOA — INTERNAL (2 codes)

CodeCauseResolution
ILOA901LegalizeOpLevelAlias assertion error: {baseMessage}(contact)
ILOA902LegalizeOpLevelAlias error: {baseMessage}(contact)

ILOP — INTERNAL (2 codes)

CodeCauseResolution
ILOP901LocalLayoutOpt assertion error: {baseMessage}(contact)
ILOP902LocalLayoutOpt error: {baseMessage}(contact)

ILPP — INTERNAL (2 codes)

CodeCauseResolution
ILPP901LayoutPreprocessing assertion error: {baseMessage}(contact)
ILPP902LayoutPreprocessing error: {baseMessage}(contact)

ILPR — INTERNAL (2 codes)

CodeCauseResolution
ILPR901LegalizePartitionReduce assertion error: {baseMessage}(contact)
ILPR902LegalizePartitionReduce error: {baseMessage}(contact)

ILPT — INTERNAL (2 codes)

CodeCauseResolution
ILPT901LegalizePartitionTile assertion error: {baseMessage}(contact)
ILPT902LegalizePartitionTile error: {baseMessage}(contact)

ILPX — INTERNAL (2 codes)

CodeCauseResolution
ILPX901LowerPartitionTile assertion error: {baseMessage}(contact)
ILPX902LowerPartitionTile error: {baseMessage}(contact)

ILRA — INTERNAL (2 codes)

CodeCauseResolution
ILRA901LayoutRequirementAnalysis assertion error: {baseMessage}(contact)
ILRA902LayoutRequirementAnalysis error: {baseMessage}(contact)

ILSA — INTERNAL (2 codes)

CodeCauseResolution
ILSA901LegalizeSundaAccess assertion error: {baseMessage}(contact)
ILSA902LegalizeSundaAccess error: {baseMessage}(contact)

ILSC — INTERNAL (2 codes)

CodeCauseResolution
ILSC901LNCShardingConstraintOp assertion error: {baseMessage}(contact)
ILSC902LNCShardingConstraintOp error: {baseMessage}(contact)

ILSM — INTERNAL (2 codes)

CodeCauseResolution
ILSM901LegalizeSundaMacro assertion error: {baseMessage}(contact)
ILSM902LegalizeSundaMacro error: {baseMessage}(contact)

ILSP — INTERNAL (2 codes)

CodeCauseResolution
ILSP901LoopSplitting assertion error: {baseMessage}(contact)
ILSP902LoopSplitting error: {baseMessage}(contact)

ILSR — INTERNAL (2 codes)

CodeCauseResolution
ILSR901LowerToSendRecv assertion error: {baseMessage}(contact)
ILSR902LowerToSendRecv error: {baseMessage}(contact)

ILSX — INTERNAL (2 codes)

CodeCauseResolution
ILSX901LowerShardAxis assertion error: {baseMessage}(contact)
ILSX902LowerShardAxis error: {baseMessage}(contact)

ILTA — INTERNAL (2 codes)

CodeCauseResolution
ILTA901LegalizeTongaAccess assertion error: {baseMessage}(contact)
ILTA902LegalizeTongaAccess error: {baseMessage}(contact)

ILTO — INTERNAL (2 codes)

CodeCauseResolution
ILTO901LowerTensorOp assertion error: {baseMessage}(contact)
ILTO902LowerTensorOp error: {baseMessage}(contact)

ILTR — INTERNAL (2 codes)

CodeCauseResolution
ILTR901LowerTranspose assertion error: {baseMessage}(contact)
ILTR902LowerTranspose error: {baseMessage}(contact)

ILTS — INTERNAL (2 codes)

CodeCauseResolution
ILTS901LegalizeTongaStore assertion error: {baseMessage}(contact)
ILTS902LegalizeTongaStore error: {baseMessage}(contact)

ILTY — INTERNAL (2 codes)

CodeCauseResolution
ILTY901LegalizeType assertion error: {baseMessage}(contact)
ILTY902LegalizeType error: {baseMessage}(contact)

IMCE — INTERNAL (2 codes)

CodeCauseResolution
IMCE901MemcpyElimination assertion error: {baseMessage}(contact)
IMCE902MemcpyElimination error: {baseMessage}(contact)

IMDD — INTERNAL (2 codes)

CodeCauseResolution
IMDD901ModDivDelinear assertion error: {baseMessage}(contact)
IMDD902ModDivDelinear error: {baseMessage}(contact)

IMDT — INTERNAL (2 codes)

CodeCauseResolution
IMDT901MutateDataType assertion error: {baseMessage}(contact)
IMDT902MutateDataType error: {baseMessage}(contact)

IMGN — INTERNAL (2 codes)

CodeCauseResolution
IMGN901MacroGeneration assertion error: {baseMessage}(contact)
IMGN902MacroGeneration error: {baseMessage}(contact)

IMPR — INTERNAL (2 codes)

CodeCauseResolution
IMPR901MaskPropagation assertion error: {baseMessage}(contact)
IMPR902MaskPropagation error: {baseMessage}(contact)

IMS — INTERNAL (5 codes)

CodeCauseResolution
IMS705common loop not found(contact)
IMS706unexpected dead mem locs(contact)
IMS707missing a matmult(contact)
IMS708No loop instruction created(contact)
IMS709No loop instruction created(contact)

IMSA — INTERNAL (2 codes)

CodeCauseResolution
IMSA901MeshAxis assertion error: {baseMessage}(contact)
IMSA902MeshAxis error: {baseMessage}(contact)

IMSH — INTERNAL (2 codes)

CodeCauseResolution
IMSH901Mesh assertion error: {baseMessage}(contact)
IMSH902Mesh error: {baseMessage}(contact)

INBU — INTERNAL (2 codes)

CodeCauseResolution
INBU901NeuronBufferUsageAnalysis assertion error: {baseMessage}(contact)
INBU902NeuronBufferUsageAnalysis error: {baseMessage}(contact)

INCP — INTERNAL (2 codes)

CodeCauseResolution
INCP901NkiCodegenPass assertion error: {baseMessage}(contact)
INCP902NkiCodegenPass error: {baseMessage}(contact)

INDI — INTERNAL (2 codes)

CodeCauseResolution
INDI901NeuronAliasDependencyInduction assertion error: {baseMessage}(contact)
INDI902NeuronAliasDependencyInduction error: {baseMessage}(contact)

INDR — INTERNAL (2 codes)

CodeCauseResolution
INDR901NeuronAliasDependencyReset assertion error: {baseMessage}(contact)
INDR902NeuronAliasDependencyReset error: {baseMessage}(contact)

INDV — INTERNAL (2 codes)

CodeCauseResolution
INDV901NeuronAliasDependencyVerificationPass assertion error: {baseMessage}(contact)
INDV902NeuronAliasDependencyVerificationPass error: {baseMessage}(contact)

INEC — INTERNAL (2 codes)

CodeCauseResolution
INEC901NeuronEliminateCasts assertion error: {baseMessage}(contact)
INEC902NeuronEliminateCasts error: {baseMessage}(contact)

INIC — INTERNAL (2 codes)

CodeCauseResolution
INIC901NeuronInstComb assertion error: {baseMessage}(contact)
INIC902NeuronInstComb error: {baseMessage}(contact)

INIS — INTERNAL (2 codes)

CodeCauseResolution
INIS901NeuronISel assertion error: {baseMessage}(contact)
INIS902NeuronISel error: {baseMessage}(contact)

INKN — INTERNAL (2 codes)

CodeCauseResolution
INKN901InlineNKIKernels assertion error: {baseMessage}(contact)
INKN902InlineNKIKernels error: {baseMessage}(contact)

INLC — INTERNAL (2 codes)

CodeCauseResolution
INLC901NeuronLICM assertion error: {baseMessage}(contact)
INLC902NeuronLICM error: {baseMessage}(contact)

INLF — INTERNAL (2 codes)

CodeCauseResolution
INLF901NeuronLoopFusion assertion error: {baseMessage}(contact)
INLF902NeuronLoopFusion error: {baseMessage}(contact)

INLI — INTERNAL (2 codes)

CodeCauseResolution
INLI901NeuronLoopInterchange assertion error: {baseMessage}(contact)
INLI902NeuronLoopInterchange error: {baseMessage}(contact)

INPN — INTERNAL (2 codes)

CodeCauseResolution
INPN901NeuronPerfectLoopNest assertion error: {baseMessage}(contact)
INPN902NeuronPerfectLoopNest error: {baseMessage}(contact)

INSM — INTERNAL (2 codes)

CodeCauseResolution
INSM901NeuronSimplifier assertion error: {baseMessage}(contact)
INSM902NeuronSimplifier error: {baseMessage}(contact)

INSP — INTERNAL (2 codes)

CodeCauseResolution
INSP901NeuronSimplifyPredicates assertion error: {baseMessage}(contact)
INSP902NeuronSimplifyPredicates error: {baseMessage}(contact)

INST — INTERNAL (2 codes)

CodeCauseResolution
INST901NeuronSizeTiling assertion error: {baseMessage}(contact)
INST902NeuronSizeTiling error: {baseMessage}(contact)

INTC — INTERNAL (2 codes)

CodeCauseResolution
INTC901NKITensorCanonicalization assertion error: {baseMessage}(contact)
INTC902NKITensorCanonicalization error: {baseMessage}(contact)

INVN — INTERNAL (2 codes)

CodeCauseResolution
INVN901NeuronValueNumbering assertion error: {baseMessage}(contact)
INVN902NeuronValueNumbering error: {baseMessage}(contact)

INVR — INTERNAL (2 codes)

CodeCauseResolution
INVR901NeuronVerifier assertion error: {baseMessage}(contact)
INVR902NeuronVerifier error: {baseMessage}(contact)

IOAC — INTERNAL (2 codes)

CodeCauseResolution
IOAC901OptimizeAliasedCopyChain assertion error: {baseMessage}(contact)
IOAC902OptimizeAliasedCopyChain error: {baseMessage}(contact)

IOLV — INTERNAL (2 codes)

CodeCauseResolution
IOLV901OperatorLayoutVisualizer assertion error: {baseMessage}(contact)
IOLV902OperatorLayoutVisualizer error: {baseMessage}(contact)

IONK — INTERNAL (2 codes)

CodeCauseResolution
IONK901OptimizeNKIKernels assertion error: {baseMessage}(contact)
IONK902OptimizeNKIKernels error: {baseMessage}(contact)

IOPF — INTERNAL (2 codes)

CodeCauseResolution
IOPF901OperatorFusion assertion error: {baseMessage}(contact)
IOPF902OperatorFusion error: {baseMessage}(contact)

IPAA — INTERNAL (2 codes)

CodeCauseResolution
IPAA901ParAxesAnnotation assertion error: {baseMessage}(contact)
IPAA902ParAxesAnnotation error: {baseMessage}(contact)

IPAS — INTERNAL (2 codes)

CodeCauseResolution
IPAS901PredicateAffineSelect assertion error: {baseMessage}(contact)
IPAS902PredicateAffineSelect error: {baseMessage}(contact)

IPCC — INTERNAL (2 codes)

CodeCauseResolution
IPCC901PComputeCutting assertion error: {baseMessage}(contact)
IPCC902PComputeCutting error: {baseMessage}(contact)

IPDE — INTERNAL (2 codes)

CodeCauseResolution
IPDE901PadElimination assertion error: {baseMessage}(contact)
IPDE902PadElimination error: {baseMessage}(contact)

IPER — INTERNAL (1 codes)

CodeCauseResolution
IPER012

IPF — INTERNAL (17 codes)

CodeCauseResolution
IPF001InsertPTCOMFlat can only run on a program with 1 subgraph, but there are {actual} subgraphs.(contact)
IPF002InsertPTCOMFlat can only handle modules with a single Function, but module {name} has {actual}(contact)
IPF003No BasicBlock named '{bb_name}' found on core0(contact)
IPF004No CoreBarrier with matching id ({id}) found for '{cb_name}' in Module '{mod_name}'(contact)
IPF005Invalid value '{value}' for '--ptcom-corebarrier-delay-count'; it must be >0(contact)
IPF006Malformed BIR - exit block must not be a part of any cycles(contact)
IPF007Malformed BIR - entry block must not be a part of any cycles(contact)
IPF008findCBOnCore0 expected a CoreBarrier but got null(contact)
IPF009expected CFG of function '{function}' to have a single exit node but it has {actual}(contact)
IPF010findNewLastWrite expected a MemoryLocation but got null(contact)
IPF011expected 1 Module for LNC1 but got {actual}(contact)
IPF012unable to find a valid extendedWrite instruction(contact)
IPF013expected extendedWrite to be an Instruction but got null(contact)
IPF014findMemLocOnCore0 expected >0 Modules(contact)
IPF015findMemLocOnCore0 expected just 1 Function on core0 but got {actual}(contact)
IPF016findMemLocOnCore0 expected to find a MemoryLocation on core0, but found null(contact)
IPF017identifyPTCOMBlocks skipped over ExternalOutput MemoryLocation '{name}'(contact)

IPFN — INTERNAL (2 codes)

CodeCauseResolution
IPFN901PrintFunction assertion error: {baseMessage}(contact)
IPFN902PrintFunction error: {baseMessage}(contact)

IPLF — INTERNAL (2 codes)

CodeCauseResolution
IPLF901PartialLoopFusion assertion error: {baseMessage}(contact)
IPLF902PartialLoopFusion error: {baseMessage}(contact)

IPLN — INTERNAL (2 codes)

CodeCauseResolution
IPLN901PerfectLoopNest assertion error: {baseMessage}(contact)
IPLN902PerfectLoopNest error: {baseMessage}(contact)

IPLO — INTERNAL (2 codes)

CodeCauseResolution
IPLO901ParLayoutOpt assertion error: {baseMessage}(contact)
IPLO902ParLayoutOpt error: {baseMessage}(contact)

IPLX — INTERNAL (2 codes)

CodeCauseResolution
IPLX901PartitionLocalityOpt assertion error: {baseMessage}(contact)
IPLX902PartitionLocalityOpt error: {baseMessage}(contact)

IPMG — INTERNAL (1 codes)

CodeCauseResolution
IPMG001DRAM2DBlockTensor requires npartitions=1: {error_msg}(contact)

IPMN — INTERNAL (2 codes)

CodeCauseResolution
IPMN901PassManager assertion error: {baseMessage}(contact)
IPMN902PassManager error: {baseMessage}(contact)

IPNI — INTERNAL (2 codes)

CodeCauseResolution
IPNI901PeepholeNeuronInstComb assertion error: {baseMessage}(contact)
IPNI902PeepholeNeuronInstComb error: {baseMessage}(contact)

IPNK — INTERNAL (2 codes)

CodeCauseResolution
IPNK901PrintNKIKernelLayout assertion error: {baseMessage}(contact)
IPNK902PrintNKIKernelLayout error: {baseMessage}(contact)

IPSF — INTERNAL (2 codes)

CodeCauseResolution
IPSF901PartialSimdFusion assertion error: {baseMessage}(contact)
IPSF902PartialSimdFusion error: {baseMessage}(contact)

IPST — INTERNAL (2 codes)

CodeCauseResolution
IPST901InferPSumTensor assertion error: {baseMessage}(contact)
IPST902InferPSumTensor error: {baseMessage}(contact)

IRAC — INTERNAL (2 codes)

CodeCauseResolution
IRAC901ResolveAccessConflict assertion error: {baseMessage}(contact)
IRAC902ResolveAccessConflict error: {baseMessage}(contact)

IRCP — INTERNAL (2 codes)

CodeCauseResolution
IRCP901Recompute assertion error: {baseMessage}(contact)
IRCP902Recompute error: {baseMessage}(contact)

IRCX — INTERNAL (2 codes)

CodeCauseResolution
IRCX901ResolveComplicatePredicates assertion error: {baseMessage}(contact)
IRCX902ResolveComplicatePredicates error: {baseMessage}(contact)

IRMT — INTERNAL (2 codes)

CodeCauseResolution
IRMT901Rematerialization assertion error: {baseMessage}(contact)
IRMT902Rematerialization error: {baseMessage}(contact)

IRNS — INTERNAL (2 codes)

CodeCauseResolution
IRNS901RelocateNKIStack assertion error: {baseMessage}(contact)
IRNS902RelocateNKIStack error: {baseMessage}(contact)

IROE — INTERNAL (2 codes)

CodeCauseResolution
IROE901ReshapeOpElimination assertion error: {baseMessage}(contact)
IROE902ReshapeOpElimination error: {baseMessage}(contact)

IROI — INTERNAL (2 codes)

CodeCauseResolution
IROI901RecognizeOpIdiom assertion error: {baseMessage}(contact)
IROI902RecognizeOpIdiom error: {baseMessage}(contact)

IRPR — INTERNAL (2 codes)

CodeCauseResolution
IRPR901ResolvePredicates assertion error: {baseMessage}(contact)
IRPR902ResolvePredicates error: {baseMessage}(contact)

IRPX — INTERNAL (2 codes)

CodeCauseResolution
IRPX901RelaxPredicates assertion error: {baseMessage}(contact)
IRPX902RelaxPredicates error: {baseMessage}(contact)

IRRM — INTERNAL (2 codes)

CodeCauseResolution
IRRM901RewriteReplicationMatmul assertion error: {baseMessage}(contact)
IRRM902RewriteReplicationMatmul error: {baseMessage}(contact)

IRRW — INTERNAL (2 codes)

CodeCauseResolution
IRRW901RewriteWeights assertion error: {baseMessage}(contact)
IRRW902RewriteWeights error: {baseMessage}(contact)

IRSA — INTERNAL (2 codes)

CodeCauseResolution
IRSA901ReinsertShardAxis assertion error: {baseMessage}(contact)
IRSA902ReinsertShardAxis error: {baseMessage}(contact)

IRSP — INTERNAL (2 codes)

CodeCauseResolution
IRSP901RemoveShardedPartitionAxes assertion error: {baseMessage}(contact)
IRSP902RemoveShardedPartitionAxes error: {baseMessage}(contact)

IRSW — INTERNAL (2 codes)

CodeCauseResolution
IRSW901ReshapeWeights assertion error: {baseMessage}(contact)
IRSW902ReshapeWeights error: {baseMessage}(contact)

ISAB — INTERNAL (2 codes)

CodeCauseResolution
ISAB901InsertImplicitShardAxisBeforeISel assertion error: {baseMessage}(contact)
ISAB902InsertImplicitShardAxisBeforeISel error: {baseMessage}(contact)

ISAG — INTERNAL (2 codes)

CodeCauseResolution
ISAG901SplitAccGrp assertion error: {baseMessage}(contact)
ISAG902SplitAccGrp error: {baseMessage}(contact)

ISAR — INTERNAL (2 codes)

CodeCauseResolution
ISAR901SimpleAllReduceTiling assertion error: {baseMessage}(contact)
ISAR902SimpleAllReduceTiling error: {baseMessage}(contact)

ISAU — INTERNAL (2 codes)

CodeCauseResolution
ISAU901SplitAPUnionSets assertion error: {baseMessage}(contact)
ISAU902SplitAPUnionSets error: {baseMessage}(contact)

ISCH — INTERNAL (3 codes)

CodeCauseResolution
ISCH001
ISCH002
ISCH003

ISDC — INTERNAL (2 codes)

CodeCauseResolution
ISDC901SpmdDCE assertion error: {baseMessage}(contact)
ISDC902SpmdDCE error: {baseMessage}(contact)

ISDD — INTERNAL (11 codes)

CodeCauseResolution
ISDD013
ISDD014
ISDD015
ISDD016
ISDD017
ISDD018
ISDD019
ISDD020
ISDD021
ISDD901SoftmaxDivisionDelay assertion error: {baseMessage}(contact)
ISDD902SoftmaxDivisionDelay error: {baseMessage}(contact)

ISFD — INTERNAL (2 codes)

CodeCauseResolution
ISFD901SplitFreeDim assertion error: {baseMessage}(contact)
ISFD902SplitFreeDim error: {baseMessage}(contact)

ISFV — INTERNAL (2 codes)

CodeCauseResolution
ISFV901SFKVectorizer assertion error: {baseMessage}(contact)
ISFV902SFKVectorizer error: {baseMessage}(contact)

ISI — INTERNAL (3 codes)

CodeCauseResolution
ISI001All Collective Compute Instructions using the same src-target pairs must have the same stream ID(contact)
ISI002All Collective Compute Instructions using the same replica groups must have the same stream ID(contact)
ISI003All GetCurProcessingRankID Instructions using the same replica groups must have the same stream ID(contact)

ISIM — INTERNAL (145 codes)

CodeCauseResolution
ISIM137
ISIM138
ISIM139
ISIM140
ISIM141
ISIM142
ISIM143
ISIM144
ISIM145
ISIM146
ISIM147
ISIM148
ISIM149
ISIM150
ISIM151
ISIM152
ISIM153
ISIM154
ISIM155
ISIM156
ISIM157
ISIM158
ISIM159
ISIM160
ISIM161
ISIM162
ISIM163
ISIM164
ISIM165
ISIM166
ISIM167
ISIM168
ISIM169
ISIM170
ISIM171
ISIM172
ISIM173
ISIM174
ISIM175
ISIM176
ISIM177
ISIM178
ISIM179
ISIM180
ISIM181
ISIM182
ISIM183
ISIM184
ISIM185
ISIM186
ISIM187
ISIM188
ISIM189
ISIM190
ISIM191
ISIM192
ISIM193
ISIM194
ISIM195
ISIM196
ISIM197
ISIM198
ISIM199
ISIM200
ISIM201
ISIM202
ISIM203
ISIM204
ISIM205
ISIM206
ISIM207
ISIM208
ISIM209
ISIM210
ISIM211
ISIM212
ISIM213
ISIM214
ISIM215
ISIM216
ISIM217
ISIM218
ISIM219
ISIM220
ISIM221
ISIM222
ISIM223
ISIM224
ISIM225
ISIM226
ISIM227
ISIM228
ISIM229
ISIM230
ISIM231
ISIM232
ISIM233
ISIM234
ISIM235
ISIM236
ISIM237
ISIM238
ISIM239
ISIM240
ISIM241
ISIM242
ISIM243
ISIM244
ISIM245
ISIM246
ISIM247Functions may not have multiple shard ids, but {func_name} has {num_shard_ids}.(contact)
ISIM248Functions may not have multiple shard ids, but {func_name} has {num_shard_ids}.(contact)
ISIM249
ISIM250
ISIM251
ISIM252
ISIM253
ISIM254
ISIM255
ISIM256
ISIM257
ISIM258
ISIM259
ISIM260
ISIM261
ISIM262Load from memory with unsupported data type {data_type} into register(contact)
ISIM263The instruction cannot have an unassigned engine(contact)
ISIM264Expected all cores to reach collective op {ccOPInstIdx} but core {core_id} is at {coreCCOPInstIdx}(contact)
ISIM267
ISIM268
ISIM269
ISIM270
ISIM271
ISIM272
ISIM273
ISIM275
ISIM276{inst_type} BIRSim simulation internal error: not all engines are waiting on the same {inst_type}, inst1.name = {inst1_name}, inst2.name = {inst2_name}(contact)
ISIM284{inst_type} BIRSim simulation internal error: all engines waiting on all-engine {inst_type} but there are pending sync waiters(contact)
ISIM285
ISIM286
ISIM287
ISIM291
ISIM293Random mode not supported for int64/uint64 memset(contact)
ISIM901SimulatorPass assertion error: {baseMessage}(contact)
ISIM902SimulatorPass error: {baseMessage}(contact)

ISIS — INTERNAL (2 codes)

CodeCauseResolution
ISIS901SundaISel assertion error: {baseMessage}(contact)
ISIS902SundaISel error: {baseMessage}(contact)

ISMP — INTERNAL (2 codes)

CodeCauseResolution
ISMP901Simplifier assertion error: {baseMessage}(contact)
ISMP902Simplifier error: {baseMessage}(contact)

ISMT — INTERNAL (2 codes)

CodeCauseResolution
ISMT901SMTAllocationPass assertion error: {baseMessage}(contact)
ISMT902SMTAllocationPass error: {baseMessage}(contact)

ISMX — INTERNAL (2 codes)

CodeCauseResolution
ISMX901SimplifyMacroPredicates assertion error: {baseMessage}(contact)
ISMX902SimplifyMacroPredicates error: {baseMessage}(contact)

ISNT — INTERNAL (2 codes)

CodeCauseResolution
ISNT901SimplifyNeuronTensor assertion error: {baseMessage}(contact)
ISNT902SimplifyNeuronTensor error: {baseMessage}(contact)

ISOE — INTERNAL (2 codes)

CodeCauseResolution
ISOE901SliceOpElimination assertion error: {baseMessage}(contact)
ISOE902SliceOpElimination error: {baseMessage}(contact)

ISPA — INTERNAL (2 codes)

CodeCauseResolution
ISPA901ShardingPropagationAnalysis assertion error: {baseMessage}(contact)
ISPA902ShardingPropagationAnalysis error: {baseMessage}(contact)

ISPC — INTERNAL (2 codes)

CodeCauseResolution
ISPC901SPMDCodeGen assertion error: {baseMessage}(contact)
ISPC902SPMDCodeGen error: {baseMessage}(contact)

ISPR — INTERNAL (2 codes)

CodeCauseResolution
ISPR901SimplifyPredicates assertion error: {baseMessage}(contact)
ISPR902SimplifyPredicates error: {baseMessage}(contact)

ISPS — INTERNAL (2 codes)

CodeCauseResolution
ISPS901SpillPSum assertion error: {baseMessage}(contact)
ISPS902SpillPSum error: {baseMessage}(contact)

ISSA — INTERNAL (5 codes)

CodeCauseResolution
ISSA004
ISSA005
ISSA006
ISSA007
ISSA008

ISSD — INTERNAL (1 codes)

CodeCauseResolution
ISSD001

ISSL — INTERNAL (2 codes)

CodeCauseResolution
ISSL901SimplifySlice assertion error: {baseMessage}(contact)
ISSL902SimplifySlice error: {baseMessage}(contact)

ISSS — INTERNAL (2 codes)

CodeCauseResolution
ISSS901SymbolicStackAllocator assertion error: {baseMessage}(contact)
ISSS902SymbolicStackAllocator error: {baseMessage}(contact)

ISST — INTERNAL (2 codes)

CodeCauseResolution
ISST901SundaSizeTiling assertion error: {baseMessage}(contact)
ISST902SundaSizeTiling error: {baseMessage}(contact)

ISTA — INTERNAL (2 codes)

CodeCauseResolution
ISTA901StackAllocator assertion error: {baseMessage}(contact)
ISTA902StackAllocator error: {baseMessage}(contact)

ISTL — INTERNAL (2 codes)

CodeCauseResolution
ISTL901StaticTransposeLocalTensor assertion error: {baseMessage}(contact)
ISTL902StaticTransposeLocalTensor error: {baseMessage}(contact)

ISTN — INTERNAL (2 codes)

CodeCauseResolution
ISTN901SimplifyTensor assertion error: {baseMessage}(contact)
ISTN902SimplifyTensor error: {baseMessage}(contact)

ISTP — INTERNAL (2 codes)

CodeCauseResolution
ISTP901StaticProfiler assertion error: {baseMessage}(contact)
ISTP902StaticProfiler error: {baseMessage}(contact)

ISUP — INTERNAL (2 codes)

CodeCauseResolution
ISUP901SuperSimplifier assertion error: {baseMessage}(contact)
ISUP902SuperSimplifier error: {baseMessage}(contact)

ISUS — INTERNAL (2 codes)

CodeCauseResolution
ISUS901SplitUnionSets assertion error: {baseMessage}(contact)
ISUS902SplitUnionSets error: {baseMessage}(contact)

ISYC — INTERNAL (1 codes)

CodeCauseResolution
ISYC001

ITCC — INTERNAL (2 codes)

CodeCauseResolution
ITCC901TileCCOps assertion error: {baseMessage}(contact)
ITCC902TileCCOps error: {baseMessage}(contact)

ITCO — INTERNAL (2 codes)

CodeCauseResolution
ITCO901TransformConvOp assertion error: {baseMessage}(contact)
ITCO902TransformConvOp error: {baseMessage}(contact)

ITCT — INTERNAL (2 codes)

CodeCauseResolution
ITCT901TCTransform assertion error: {baseMessage}(contact)
ITCT902TCTransform error: {baseMessage}(contact)

ITDM — INTERNAL (2 codes)

CodeCauseResolution
ITDM901TransformerDimensionMatcher assertion error: {baseMessage}(contact)
ITDM902TransformerDimensionMatcher error: {baseMessage}(contact)

ITEN — INTERNAL (23 codes)

CodeCauseResolution
ITEN404Internal tensorizer error: {baseMessage}(contact)
ITEN405Internal tensorizer maximum recursion depth exceeded, {stack}(contact)
ITEN406Too many partition dimensions detected at {partition_set}. This is usually due to unsupported (strided) access pattern(contact)
ITEN407QuantizeMXTensorOp verification error: {error_msg}(contact)
ITEN408QuantizeMXTensorOp verification error: {error_msg}(contact)
ITEN409QuantizeMXTensorOp verification error: {error_msg}(contact)
ITEN410QuantizeMXTensorOp verification error: {error_msg}(contact)
ITEN411QuantizeMXTensorOp verification error: {error_msg}(contact)
ITEN412QuantizeMXTensorOp verification error: {error_msg}(contact)
ITEN413QuantizeMXTensorOp verification error: {error_msg}(contact)
ITEN414QuantizeMXTensorOp verification error: {error_msg}(contact)
ITEN420ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN421ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN422ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN423ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN424ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN425ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN426ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN427ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN428ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN429ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN430ScaledTensorContractTensorOp verifier error : {error_msg}(contact)
ITEN431ScaledTensorContractTensorOp verifier error : {error_msg}(contact)

ITIN — INTERNAL (2 codes)

CodeCauseResolution
ITIN901TensorInitialization assertion error: {baseMessage}(contact)
ITIN902TensorInitialization error: {baseMessage}(contact)

ITLA — INTERNAL (2 codes)

CodeCauseResolution
ITLA901TensorLifetimeAnalysis assertion error: {baseMessage}(contact)
ITLA902TensorLifetimeAnalysis error: {baseMessage}(contact)

ITLX — INTERNAL (2 codes)

CodeCauseResolution
ITLX901TransformLayout assertion error: {baseMessage}(contact)
ITLX902TransformLayout error: {baseMessage}(contact)

ITNM — INTERNAL (2 codes)

CodeCauseResolution
ITNM901TensorNamer assertion error: {baseMessage}(contact)
ITNM902TensorNamer error: {baseMessage}(contact)

ITOE — INTERNAL (2 codes)

CodeCauseResolution
ITOE901TransposeOpElimination assertion error: {baseMessage}(contact)
ITOE902TransposeOpElimination error: {baseMessage}(contact)

ITOS — INTERNAL (2 codes)

CodeCauseResolution
ITOS901TensorOpSimplifier assertion error: {baseMessage}(contact)
ITOS902TensorOpSimplifier error: {baseMessage}(contact)

ITOT — INTERNAL (2 codes)

CodeCauseResolution
ITOT901TensorOpTransform assertion error: {baseMessage}(contact)
ITOT902TensorOpTransform error: {baseMessage}(contact)

ITPO — INTERNAL (2 codes)

CodeCauseResolution
ITPO901TensorParallelOpt assertion error: {baseMessage}(contact)
ITPO902TensorParallelOpt error: {baseMessage}(contact)

ITPR — INTERNAL (2 codes)

CodeCauseResolution
ITPR901TilingProfiler assertion error: {baseMessage}(contact)
ITPR902TilingProfiler error: {baseMessage}(contact)

ITRF — INTERNAL (2 codes)

CodeCauseResolution
ITRF901TritiumFusion assertion error: {baseMessage}(contact)
ITRF902TritiumFusion error: {baseMessage}(contact)

ITSH — INTERNAL (2 codes)

CodeCauseResolution
ITSH901TensorSharding assertion error: {baseMessage}(contact)
ITSH902TensorSharding error: {baseMessage}(contact)

ITST — INTERNAL (1 codes)

CodeCauseResolution
ITST001Example internal error message 1.(contact)

IUBK — INTERNAL (2 codes)

CodeCauseResolution
IUBK901UnshardBIRKernelForSimulation assertion error: {baseMessage}(contact)
IUBK902UnshardBIRKernelForSimulation error: {baseMessage}(contact)

IURM — INTERNAL (2 codes)

CodeCauseResolution
IURM901UnrollReduceMacro assertion error: {baseMessage}(contact)
IURM902UnrollReduceMacro error: {baseMessage}(contact)

IUTL — INTERNAL (3 codes)

CodeCauseResolution
IUTL001
IUTL002
IUTL003

IVDM — INTERNAL (2 codes)

CodeCauseResolution
IVDM901VectorizeDMA assertion error: {baseMessage}(contact)
IVDM902VectorizeDMA error: {baseMessage}(contact)

IVLP — INTERNAL (2 codes)

CodeCauseResolution
IVLP901VectorizeLoop assertion error: {baseMessage}(contact)
IVLP902VectorizeLoop error: {baseMessage}(contact)

IVMM — INTERNAL (2 codes)

CodeCauseResolution
IVMM901VectorizeMatMult assertion error: {baseMessage}(contact)
IVMM902VectorizeMatMult error: {baseMessage}(contact)

IVNU — INTERNAL (2 codes)

CodeCauseResolution
IVNU901ValueNumbering assertion error: {baseMessage}(contact)
IVNU902ValueNumbering error: {baseMessage}(contact)

IWCO — INTERNAL (2 codes)

CodeCauseResolution
IWCO901WeightCoalescing assertion error: {baseMessage}(contact)
IWCO902WeightCoalescing error: {baseMessage}(contact)

IXCG — INTERNAL (32 codes)

CodeCauseResolution
IXCG013ISA check {checkName} failed: {msg}(contact)
IXCG014ISA check {checkName} failed: {msg}(contact)
IXCG015ISA check {checkName} failed: {msg}(contact)
IXCG016
IXCG017
IXCG018Instruction engine check failed ({engine})(contact)
IXCG019ISA check {checkName} failed: {msg}(contact)
IXCG020ISA check failed(contact)
IXCG021
IXCG022
IXCG023
IXCG024
IXCG025
IXCG026
IXCG027
IXCG028
IXCG029Unknown sync wait mode
IXCG030Unknown sync update mode
IXCG031Unimplemented Dtype {type}
IXCG034
IXCG035
IXCG036
IXCG037
IXCG038
IXCG039
IXCG040
IXCG041
IXCG042
IXCG043Unknown EngineType
IXCG044Instruction engine check failed ({engine})
IXCG045ISA check {checkName} failed: {msg}
IXCG046ISA check failed: {msg}

IXLV — INTERNAL (1 codes)

CodeCauseResolution
IXLV035

IZST — INTERNAL (2 codes)

CodeCauseResolution
IZST901ZeroSizeTensorElimination assertion error: {baseMessage}(contact)
IZST902ZeroSizeTensorElimination error: {baseMessage}(contact)

JIO — bare (18 codes)

CodeCauseResolution
JIO001Unable open json file {path}Please ensure file exists and disk space is available.
JIO002Unable parse empty json file {path}Please ensure disk space is available.
JIO003Encountered parsing error, {error_msg}, while attempting to read {path}(contact)
JIO004Encountered json error, {error_msg}, while attempting to parse {path}.(contact)
JIO005Encountered exception, {error_msg}, while attempting to parse {path}. IO diagnostic checks: {diagnostics}(contact)
JIO006Encountered unexpected exception while attempting to parse {path}. IO diagnostic checks: {diagnostics}(contact)
JIO007Unable write json to {path}. IO diagnostic checks: {diagnostics}Please ensure diretory exists and disk space is available.
JIO008Unable write json to {path} due to {error_msg}Please ensure json fields contain encodable unicode characters.
JIO009Encountered json error, {error_msg}, while attempting to write json to {path}(contact)
JIO010Encountered exception, {error_msg}, while attempting to write json to {path}. IO diagnostic checks: {diagnostics}(contact)
JIO011Encountered unexpected exception while attempting to write json to {path}. IO diagnostic checks: {diagnostics}(contact)
JIO012Unable write json to {path} because of error writing to file. IO diagnostic checks: {diagnostics}Please ensure diretory exists and disk space is available.
JIO013Unable write json to {path}. IO diagnostic checks: {diagnostics}Please ensure diretory exists and disk space is available.
JIO014Unable write json to {path} due to {error_msg}Please ensure json fields contain encodable unicode characters.
JIO015Encountered json error, {error_msg}, while attempting to write json to {path}(contact)
JIO016Encountered exception, {error_msg}, while attempting to write json to {path}. IO diagnostic checks: {diagnostics}(contact)
JIO017Encountered unexpected exception while attempting to write json to {path}. IO diagnostic checks: {diagnostics}(contact)
JIO018Unable write json to {path} because of error writing to file. IO diagnostic checks: {diagnostics}Please ensure diretory exists and disk space is available.

JOB — bare (1 codes)

CodeCauseResolution
JOB001Expected Job {name} to supply a 3-letter, uppercase abbreviation, but got {abbrev}(contact)

LCL — bare (26 codes)

CodeCauseResolution
LCL731Unrecognized AxisListType(contact)
LCL732User step must fit reduced location, load reduction error(contact)
LCL741Expected a matmult(contact)
LCL742Unexpected argument(contact)
LCL743Unexpected number of arguments(contact)
LCL744Expected a matmult(contact)
LCL745Unexpected instruction type(contact)
LCL746Unexpected access pattern(contact)
LCL747Expected a tensor scalar pointer(contact)
LCL748Expected a load(contact)
LCL749unexpected null pointer(contact)
LCL750Index not in the split range(contact)
LCL751Unexpected node count(contact)
LCL752Expected a load(contact)
LCL755Cluster not found(contact)
LCL756Cluster not found(contact)
LCL757Unexpected access pattern(contact)
LCL758User access pattern read range exceeds load range, load reduction error(contact)
LCL759Start address not aligned(contact)
LCL760Memory location is not dead, load reduction error(contact)
LCL761Expected a load(contact)
LCL762Load reads from DRAM, the under partition AP cannot have a stride >1(contact)
LCL763Expected a load(contact)
LCL764Expected a load(contact)
LCL765Overlapping live range among cloned loads(contact)
LCL777Expected a clean cluster(contact)

LDD — bare (1 codes)

CodeCauseResolution
LDD001SWDGE DMAs slow dimension must be a multiple of 16 to avoid performance penalties (it is {actual})(contact)

LDM — bare (1 codes)

CodeCauseResolution
LDM001Instruction {name} not found in any DMA blocks(contact)

LEG — bare (2 codes)

CodeCauseResolution
LEG194A PSUM node that cannot be in PSUM(contact)
LEG195Tensorizer generated tensor is out of bounds post PSUM legalization.(contact)

LKK — bare (2 codes)

CodeCauseResolution
LKK001Only user-allocated kernels may have dependencies, violation: {inst} -> {edge_kind} -> {dep}.(contact)
LKK002All dependencies coming from NKI should be EdgeKind::Ordered, violation: {inst} -> {edge_kind} -> {dep}.(contact)

LKN — bare (1 codes)

CodeCauseResolution
LKN001Internal error with kernel lowering.(contact)

LLC — bare (64 codes)

CodeCauseResolution
LLC001Expected local collective compute operation to have a single replication group, but found {size} groups(contact)
LLC002Unsupported local collective compute operation {kind}(contact)
LLC003Expected module name in the format 'ncXX/sgXX', but got {mod_name}(contact)
LLC004Expected NC prefix of 'ncXX', but got {nc_name} from module {mod_name}(contact)
LLC005Local CollectiveCompute input ({input_dim}) and output ({output_dim}) dimensions differ.(contact)
LLC006Expected numeric NC index in {nc_name}, but got {nc_num} in module {mod_name}(contact)
LLC007Expected to find a splittable entry in AP, but none found.(contact)
LLC008AllReduce local collective operation currently only supports addition(contact)
LLC009ReduceScatter local collective operation currently only supports addition(contact)
LLC010Expected memory location used in local collective to have tensor id.(contact)
LLC011Input and output tensors have different dimensions {to_split_size} vs {related_size}(contact)
LLC012Input and output tensors have different dimensions {pattern1_size} vs {pattern2_size}(contact)
LLC013Unable to find suitable split dimension for output tensor {tensor} AP(contact)
LLC014Unable to find suitable split dimension for input tensor {tensor} AP(contact)
LLC015Unable to find suitable split dimension for input tensor {input} AP and output tensor {output} AP(contact)
LLC016Splitting tensor {tensor} AP has a max offset of {max_byte_offset} (in bytes), but referenced memory location is only {memloc_size} bytes.(contact)
LLC017Expected {expected_mod_count} modules, but got {actual_mod_count} modules.(contact)
LLC018Expected {expected_ccop_count} collective operations, but found {actual_ccop_count} in core {core_id} subgraph(s).(contact)
LLC019Expected one input and one output for local collective, but got {arg_count} arguments and {out_count} outputs.(contact)
LLC020Expected local collective to share input dimension {expected_arg_dims}, but core {core_id} has input dimension {actual_arg_dims}.(contact)
LLC021Expected local collective to have matching replica groups {expected_cores}, but core {core_id} has replica group {actual_cores}.(contact)
LLC022Expected local collective to have matching kinds {expected_kind}, but core {core_id} has kind {actual_kind}.(contact)
LLC023Expected scatter dimension to match {expected_scatter_dim} for all cores, but core {core_id} has scatter dimension {actual_scatter_dim}.(contact)
LLC024Expected scatter dimension to match {expected_scatter_dim} for all cores, but core {core_id} has scatter dimension {actual_scatter_dim}.(contact)
LLC025Expected gather dimension to match {expected_gather_dim} for all cores, but core {core_id} has gather dimension {actual_gather_dim}.(contact)
LLC026Expected gather dimension to match {expected_gather_dim} for all cores, but core {core_id} has gather dimension {actual_gather_dim}.(contact)
LLC027Expected address space to match {expected_addr_space} for all cores, but core {core_id} has address space {actual_addr_space}.(contact)
LLC028Expected scatter dimension to match {expected_scatter_dim} for all cores, but core {core_id} has scatter dimension {actual_scatter_dim}.(contact)
LLC029Expected gather dimension to match {expected_gather_dim} for all cores, but core {core_id} has gather dimension {actual_gather_dim}.(contact)
LLC030Shared remote memory locations ({remote}) are not currently supported in local collective operations.(contact)
LLC031Expected address space to match {expected_addr_space} for all cores, but core {core_id} has address space {actual_addr_space}.(contact)
LLC032Expected local collective to share output dimension {expected_out_dims}, but core {core_id} has output dimension {actual_out_dims}.(contact)
LLC033legalizeAPs() failed: Input/output dimension expected to be {expected_dims}, but core {core_id} could not be shrunk beyond {actual_dims}.(contact)
LLC034legalizeAPs() failed: Input/output dimension {dimension} expected to be {expected_num}, but core {core_id} has {actual_num}.(contact)
LLC035Expected shared memory definition for {name} to have a synchronization primitive before it.(contact)
LLC036Expected core barrier ahead of shared memory definition for {name}, but no core barrier at index {index} on core {core_id} exists.(contact)
LLC037Expected shared memory reference for {name} to have a synchronization primitive after it.(contact)
LLC038Expected core barrier after shared memory reference for {name}, but no core barrier at index {index} on core {core_id} exists.(contact)
LLC039Expected function {name} to have a single basic block at this point, but it has {count}.(contact)
LLC040Expected function {name} to have a single basic block at this point, but it has {count}.(contact)
LLC041Expected shared memory location for {name} con core {core_id}, but found none.(contact)
LLC042Expected shared memory location for {name} con core {core_id}, but found none.(contact)
LLC043Updating dependencies failed: expect core barrier at the beginning.(contact)
LLC044Updating dependencies failed: expect core barrier at the end when remote semaphores disabled!(contact)
LLC045Expected to find a first definition for shared tensor {name} on at least one core.(contact)
LLC046lower_local_collectives failed with the following errors:{errors}(contact)
LLC047Expected to find a definition before read for shared tensor {name}.(contact)
LLC048Expected matching SendRecv to {expected} initial core barrier, but found core {ncId} {actual} initial core barrier(contact)
LLC049Expected matching SendRecvCCE to {expected} initial core barrier, but found core {ncId} {actual} initial core barrier(contact)
LLC050Expected {expected} CoreBarrier insturctions in subgraph {sgId}, but found {actual} on core {ncId}(contact)
LLC051Expected CoreBarriers {cb_index} in subgraph {sgId} to be named {expected_name} and have id {expected}, but found {actual_name} and {actual} on core {ncId}(contact)
LLC052Expected {expected_ccop_count} local collective operations, but found {actual_ccop_count} in core {core_id} subgraph(s).(contact)
LLC053Expected exactly one LNC-predicated global collective on core 0 corresponding to this instruction, but got {num_ccops}.(contact)
LLC054Expected the module of this CollectiveCompute instruction to have a valid NeuronCoreId.(contact)
LLC055Expected the module of this CollectiveCompute instruction to have a NeuronCoreId of 0.Try assigning LNC-predicated CollectiveCompute instructions to core 0 only.
LLC056CollectiveCompute {ccName} on core 0 reads or writes to the shared MemoryLocation {accessName}. Expected all other cores to have no real reads or writes to this shared MemoryLocation, but got this instruction reading or writing it.Try removing all reads and writes to that shared MemoryLocation; otherwise don't predicate the CollectiveCompute.
LLC057LNC-predicated global collective involving SB is not supported(contact)
LLC058Could not find Function named {function_name} on core {nc_id}(contact)
LLC059Could not find MemoryLocation named {memloc_name} on core {nc_id}(contact)
LLC060Can't map sendrecv to GPSIMDSB2SB because peer IDs aren't a swap(contact)
LLC061Can't map sendrecv to GPSIMDSB2SB as requested(contact)
LLC062Internal invariant violated(contact)
LLC063Could not find BasicBlock named {block_name} on core {nc_id}(contact)
LLC064Could not find corresponding CollectiveCompute instruction on core 0 to this CollectiveCompute instruction on core {ncid}Try assigning LNC-predicated CollectiveCompute instructions to core 0 only.

LLR — bare (2 codes)

CodeCauseResolution
LLR739last instruction, doesn't need spilling(contact)
LLR740expected a save instruction(contact)

LNK — bare (5 codes)

CodeCauseResolution
LNK001MemoryLocation '{name}' cannot be written by more than 1 function call(contact)
LNK020tensor_map verification failed after linking(contact)
LNK021tensor_map verification failed: output {out} was used as input earlier(contact)
LNK022Queue '{queue}' does not exist in linked module(contact)
LNK023Could not write tensor_map file '{path}' after linking(contact)

LRS — bare (2 codes)

CodeCauseResolution
LRS778Unexpected AP(contact)
LRS779Split spill instruction is not save(contact)

LSA — bare (37 codes)

CodeCauseResolution
LSA047Internal invariant violated(contact)
LSA048Internal invariant violated(contact)
LSA049Internal invariant violated(contact)
LSA050Internal invariant violated(contact)
LSA051Internal invariant violated(contact)
LSA052Internal invariant violated(contact)
LSA053Internal invariant violated(contact)
LSA054Internal invariant violated(contact)
LSA055Internal invariant violated(contact)
LSA056Internal invariant violated(contact)
LSA057Internal invariant violated(contact)
LSA058Internal invariant violated(contact)
LSA059Internal invariant violated(contact)
LSA060Internal invariant violated(contact)
LSA061Internal invariant violated(contact)
LSA062Internal invariant violated(contact)
LSA063Internal invariant violated(contact)
LSA064Internal invariant violated(contact)
LSA065Internal invariant violated(contact)
LSA066Internal invariant violated(contact)
LSA067Illegal instruction type for spilling(contact)
LSA068Internal invariant violated(contact)
LSA069Internal invariant violated(contact)
LSA070Internal invariant violated(contact)
LSA072Illegal instruction type for spilling(contact)
LSA073Internal invariant violated(contact)
LSA074Internal invariant violated(contact)
LSA075Cannot deallocate unknown base partition(contact)
LSA076Cannot deallocate unknown memory type(contact)
LSA077Internal invariant violated(contact)
LSA078Could not find {memoryLocation} in block list(contact)
LSA079Internal invariant violated(contact)
LSA080Internal invariant violated(contact)
LSA081Internal invariant violated(contact)
LSA082Internal invariant violated(contact)
LSA083Internal invariant violated(contact)
LSA084Internal invariant violated(contact)

LUR — bare (12 codes)

CodeCauseResolution
LUR001Expected neuron_core_id module attirbute in module {name} when --lnc=2(contact)
LUR002Expected number of index map entries to match number of dependency loopnest axes(contact)
LUR003MemoryLocation of {memloc_size} bytes needed to be offset by the reserved DGE scratchpad size, and got new address {new_address}, which exceeds the SB size {sb_size}. Please allocate tensors taking the reserved DGE scratchpad size into account.(contact)
LUR004Failed to allocate MemoryLocation of {memloc_size} bytes to new address {new_address}, which needed to be offset by the reserved DGE scratchpad size. Consider turning off DGE if possible, to work around this error.(contact)
LUR005Dynamic axis already exists, please ensure that nested dynamic loops do not use the same name for their axis(contact)
LUR006Dynamic axis not found! Please ensure all instructions in the dynamic for loop use the correct axis(contact)
LUR007Axis in affine expression with runtime value is not dynamic!(contact)
LUR016Incoming IR error, tensor address expression and block-dimension are not compatible(contact)
LUR017Tensor labelled as allocated but no address expression assigned in incoming IR(contact)
LUR018PSUM Tensor labelled as allocated but no bank expression/address expr assigned in incoming IR(contact)
LUR019Local tensor not allocated in smt allocation flow(contact)
LUR020Local tensor not allocated in smt allocation flow(contact)

MFP — bare (2 codes)

CodeCauseResolution
MFP001ModuleForkPass called with single module {name}(contact)
MFP002Compilation failed for the following modules:{errors}.(contact)

MLC — bare (1 codes)

CodeCauseResolution
MLC001Block AP request is invalid, partition or free size requested is larger than MemoryLocation(contact)

MLO — bare (3 codes)

CodeCauseResolution
MLO001Expected SB <-> DRAM or PSUM <-> SB, but got {src_mem_type} -> {dst_mem_type}(contact)
MLO002OptinPass must be a number(contact)
MLO003OptinPass index {actual} from bir.json is not valid. Must be in [0, {max}).(contact)

MMP — bare (72 codes)

CodeCauseResolution
MMP110Cannot mmpacking multi-block functions.(contact)
MMP111(contact)
MMP112(contact)
MMP115(contact)
MMP117Cannot mmpacking multi-block functions.(contact)
MMP121Instruction getPrivateIndex() not following tensorizer order?(contact)
MMP123(contact)
MMP125(contact)
MMP129(contact)
MMP131(contact)
MMP133(contact)
MMP135(contact)
MMP137do not compute_partition_constraint if there're no shared mm pairs.(contact)
MMP139(contact)
MMP141(contact)
MMP143(contact)
MMP151(contact)
MMP155(contact)
MMP157(contact)
MMP161Can not handle empty accumulative groups.(contact)
MMP163Can not handle empty accumulative groups.(contact)
MMP171Cannot mmpacking multi-block functions.(contact)
MMP173Contradicting base partition found for an memloc.(contact)
MMP177(contact)
MMP181Cannot compute feasible col_grps for empty AGs!(contact)
MMP185(contact)
MMP189(contact)
MMP193(contact)
MMP201(contact)
MMP207(contact)
MMP215(contact)
MMP219(contact)
MMP223(contact)
MMP227(contact)
MMP229psum_id not found!(contact)
MMP231(contact)
MMP237(contact)
MMP241(contact)
MMP245(contact)
MMP251(contact)
MMP253(contact)
MMP255Invalid topo_order or g has cycles.(contact)
MMP257max_dist < min_dist(contact)
MMP259Rhs psum_id already belongs to this agpack!(contact)
MMP261(contact)
MMP263(contact)
MMP265(contact)
MMP267(contact)
MMP269(contact)
MMP271(contact)
MMP273(contact)
MMP275Should not invoke compute_pe_grp_occupancy on empty AGPacks.(contact)
MMP277get_pack_id(.) is undefined for standalone MMs!(contact)
MMP279use_ml is not used by mm.(contact)
MMP280(contact)
MMP281(contact)
MMP283(contact)
MMP285(contact)
MMP291(contact)
MMP295(contact)
MMP297Rpo doesn't cover all f.allocs()!\n(contact)
MMP301(contact)
MMP303(contact)
MMP305(contact)
MMP307Ready queue node already in Active.\n(contact)
MMP311(contact)
MMP313(contact)
MMP315(contact)
MMP317(contact)
MMP319undefined accumulation group!(contact)
MMP320Invalid number of matmult rows {num}, must fit in 32 bits(contact)
MMP321Invalid number of matmult cols {num}, must fit in 32 bits(contact)

MSA — bare (2 codes)

CodeCauseResolution
MSA300Memloc legalization pass cannot be apply to allocated memlocs(contact)
MSA301Memloc data type not divisible by round up size(contact)

NAS — bare (1 codes)

CodeCauseResolution
NAS001No message provided for {code}.Add {code} to message file.

NKI — bare (16 codes)

CodeCauseResolution
NKI001NKI Kernel Instruction does not have a valid function(contact)
NKI002Cannot find lowering function {fname} after unroll(contact)
NKI003Insufficient number of outputs param, expect {num_sb_buf} SB buffers and {num_psum_buf} PSUM buffers, but only got {num_actual} outputs(contact)
NKI004SBUF scratchpad memset must have exactly one memory location(contact)
NKI005DRAM location of kind {mem_kind} mapping failed. Only input/output/const DRAM location is supported!(contact)
NKI006Input/Output DRAM memory location set should only has one location, but got {memloc_count} locations(contact)
NKI007Failed to map input scratchpad memory location to {target_memloc_name}(contact)
NKI008Failed to map output scratchpad memory location to {target_memloc_name}(contact)
NKI009Unrecognized DRAM input location(contact)
NKI010Unrecognized DRAM output location(contact)
NKI011Unrecognized dependency(contact)
NKI012Instruction which copied into a loop instruction is not a loop instruction(contact)
NKI013Expected copied and original instruction names to match, but got {copied_name} != {original_name}(contact)
NKI014Expected scheduling unit instruction(contact)
NKI015Expected exactly one block in scheduling unit, but got {num_blocks}(contact)
NKI016Kernel validation exception: {error_text}Please check the validation message and adjust kernel inputs accordingly

NKS — bare (4 codes)

CodeCauseResolution
NKS000Unsupported KLR version/patch. Supported (12,0,0); found: ({major},{minor},{patch})(contact)
NKS001Incorrect kernel tag(contact)
NKS002BIRSim did not finish successfully(contact)
NKS003Could not open input KLR file: {file}(contact)

NLA — bare (2 codes)

CodeCauseResolution
NLA001Unhandled exception with message: {msg}(contact)
NLA002Unhandled exception with message: {msg}(contact)

OAC — bare (2 codes)

CodeCauseResolution
OAC001Loop header basic block {bb_name} has incorrect number of predecessors ({pred})
OAC002Loop header {bb_name} cannot contain predecessor {pred_name} inside of its loop(contact)

OQS — bare (12 codes)

CodeCauseResolution
OQS001Main function should be a single basic block in modular flow(contact)
OQS002Previous function scope stacks not fully processed(contact)
OQS003Failed to generate CFG for function {fn_name}(contact)
OQS004Unexpected unclosed scope in function {fn_name}, malformed CFG found(contact)
OQS005Failed to find CFG for basic block {bb_name} in function {fn_name}(contact)
OQS006Function scope for function {fn_name} unexpectedly ended in basic block {bb_name}, malformed CFG found(contact)
OQS007Main function in modular flow cannot contain and DMAs or SQI instructions(contact)
OQS008Loop header {bb_name} cannot contain predecessor {pred_name} inside of its loop(contact)
OQS009Loop header (or Function entry block) {bb_name} is a header for multiple loops at once, check if CFG is malformed(contact)
OQS010SQI cannot reuse an instance in a queue, {prev_instr} also uses instance {instance}(contact)
OQS011SQI instruction cannot have inputs/outputs(contact)
OQS012Basic block {bb} found in loop body of processed scope with header {loop_header}, toplogical ordering does not keep loop bodies contiguous(contact)

PER — bare (11 codes)

CodeCauseResolution
PER001Could not find PerfSim event corresponding to the instruction(contact)
PER002File cost setup failed: {filename}(contact)
PER003Invalid cost file: {filename}(contact)
PER004Instruction {inst_name} was not ready at ready time {ready_time}(contact)
PER005Can't dump trace to an empty filename(contact)
PER006Trace dumping failed: {filename}(contact)
PER007Loop or scheduling unit modeling currently unsupported(contact)
PER008Loop modeling unsupported(contact)
PER009Got stream_id {stream_id} when max_stream_id is 1(contact)
PER010Got stream_id {stream_id} when max_stream_id is 2(contact)
PER011Tried to get the CC stream bandwidth of a non-CC timeline: {name}(contact)

PGT — bare (1 codes)

CodeCauseResolution
PGT002Too many instructions after unroll!Compiling under --optlevel=1 may result in smaller graphs. If you are using a transformer model, try using a smaller context_length_estimate value.

PRS — bare (15 codes)

CodeCauseResolution
PRS001INTERNAL: trying to move producer backwards will move it further from consumers, while this optimization is trying to do the opposite(contact)
PRS002INTERNAL: instructions to move are not in ascending order, which may cause 'extra jumps' when moving them forward(contact)
PRS003INTERNAL: trying to move consumer forwards will move it further from producers, while this optimization is trying to do the opposite(contact)
PRS004INTERNAL: instructions to move are not in descending order, which may cause 'extra jumps' when moving them backward(contact)
PRS710unexpected addresses(contact)
PRS711unexpected addresses(contact)
PRS712unexpected size relationships(contact)
PRS713bytes per partition should be > 0(contact)
PRS714SpillSave AP should only has two dimensions(contact)
PRS715expected a load instruciton(contact)
PRS717mloc with read out of range {loc}(contact)
PRS718non empty memloc {mem}(contact)
PRS771Unexpected number of locations(contact)
PRS772Unexpected level {level}(contact)
PRS773Expected 1 function in the module(contact)

RDP — bare (7 codes)

CodeCauseResolution
RDP716Unexpected engine type(contact)
RDP733Unexpected engine type(contact)
RDP734Unexpected DMA instruction type(contact)
RDP735Unexpected engine type(contact)
RDP736Instruction not found in map(contact)
RDP737unexpected lifetime metric calculation mode {mode}(contact)
RDP738Instruction not found in map(contact)

SCH — bare (43 codes)

CodeCauseResolution
SCH711Unexpected dependence within accumulation group(contact)
SCH712Scheduler deadocked(contact)
SCH713Violation of accumulation group interleaving(contact)
SCH714Violation of accumulation group interleaving: got 'earlier' read {er} happening after 'later' write {lw}(contact)
SCH715Violation of accumulation group interleaving(contact)
SCH716Memory Location not found in the set(contact)
SCH717Accumulation Group object not found in the set(contact)
SCH718Cycle in output dependencies {y} : {y_name} > {x} : {x_name}(contact)
SCH719topological order violations(contact)
SCH720
SCH721
SCH722Unexpected policy value {policy}(contact)
SCH723expected one top block in a function(contact)
SCH724Unexpected opcode(contact)
SCH725cycle in dependence graph(contact)
SCH726Expected a matmult(contact)
SCH727All siblings of a scheduled matmult should be ready(contact)
SCH728Unexpected engine(contact)
SCH729Unexpected engine(contact)
SCH730Instruction not found in ready list(contact)
SCH731Ready list is empty(contact)
SCH732Ready list is empty(contact)
SCH733Unexpected negative value(contact)
SCH734Instruction should have been scheduled at this point(contact)
SCH735Predecessor should already be scheduled(contact)
SCH736Predecessor should already be scheduled(contact)
SCH737Predecessor should already be scheduled(contact)
SCH738Unexpected call to this function in this context(contact)
SCH739Nothing ready(contact)
SCH768Unexpected instruction mix(contact)
SCH769Run out of instructions(contact)
SCH770Unexpected dependence(contact)
SCH771CoreBarrier is not assigned to an engine(contact)
SCH900Time-aware scheduler failed: could not find any ready inst(contact)
SCH901Time-aware scheduler only supports scheduling units with unrolled loopcount=1 and precise schedule, but got block {block_name}(contact)
SCH902Expected more than one module with LNC > 1(contact)
SCH903Expected the same number of blocks for all cores, but got {core_0_num_blocks} blocks for core 0 and {core_n_num_blocks} blocks for core {core_id}(contact)
SCH906When scheduling instruction, found {remaining_preds} remaining predecessors which were not scheduled(contact)
SCH907LNC-aware scheduler failed: could not find any ready instruction. The number of scheduled instructions so far is {num_scheduled_insts}.(contact)
SCH908When trying to schedule instruction, could not find any ready instruction at start time {start_time}(contact)
SCH909LNC-aware scheduler failed: got a null candidate(contact)
SCH910LNC-aware scheduler failed: could not find LNC sync instruction candidate {cb_candidate_name}(contact)
SCH911LNC-aware scheduler failed: could not find LNC sync instruction candidate corresponding to the current instruction, on neuron core id {ncid}(contact)

SCP — bare (1 codes)

CodeCauseResolution
SCP001Invalid predicate value for Select {name}(contact)

SDD — bare (12 codes)

CodeCauseResolution
SDD001Dynamic DMA can only have one output, got {num_outputs} outputs instead(contact)
SDD002Scalar DGE Operation should have exactly 2 arguments, got {num_args} outputs instead(contact)
SDD003Actual input ({input_dim}D) and actual output ({output_dim}D) of dynamic DMA must have same number of dimensions(contact)
SDD004Dynamic DMA shape has {num_dim} dimensions, must be either 2D or 3D(contact)
SDD005Input ({input}) and output ({output}) at dimension {i} of dynamic DMA must have same number of elements(contact)
SDD006Last dimension of access pattern is not contiguous, step size {last_step} is not equal to 1(contact)
SDD007Max Descriptor Size violated! Access pattern has {last_num} elements, maximum is {max_num}(contact)
SDD008Partition dimension element number ({part_dim}) must be divisible by 16 or equal to 1 with the free dimension ({free_dim}) divisible by 16(contact)
SDD009Generated offsets from unrolling don't match ({in_offsets} vs {out_offsets}), check to see if input and output nums match(contact)
SDD010Dynamic instruction must be either SW DGE or HW DGE(contact)
SDD011Vector DGE shape has {num_dim} dimensions, must be 2D(contact)
SDD012Transpose instruction does not have legal shape(contact)

SFP — bare (2 codes)

CodeCauseResolution
SFP001SubgraphForkPass called with single module {name}(contact)
SFP002Compilation failed for the following subgraphs:{errors}.(contact)

SIM — bare (141 codes)

CodeCauseResolution
SIM001BirSim failed with the following errors:{errors}(contact)
SIM002Expected collective or core barrier instruction type for {inst}(contact)
SIM003Core {coreId} has a different collective or core barrier instruction type than core 0 for {inst}(contact)
SIM004Expected Core Barrier Instruction type for {inst} in {coreId}(contact)
SIM005DMA CCE: Scale array must match number of inputs(contact)
SIM006Tensor {name} is not found in remote core {core}(contact)
SIM007Function {name} is not found in remote core {core}(contact)
SIM008Unable split {dims} free dimensions into {count} parts for local collective operation.(contact)
SIM009Unsupported OP type {op} for collective compute reduction.(contact)
SIM010Expected replica size ({replica_count}) to match number of NCs per VC ({nc_count}) for local collective.(contact)
SIM011Expected input size ({in_count}) to match output size ({out_count})(contact)
SIM012Expected input size ({in_count}) to match output size ({out_count})(contact)
SIM013Expected input size ({in_count}) to match output size ({out_count})(contact)
SIM014Expected input size ({in_count}) to match output size ({out_count})(contact)
SIM015Expected input size ({in_count}) to match output size ({out_count})(contact)
SIM016Expected input size ({in_count}) to match output size ({out_count})(contact)
SIM017Expected to only handle Permute or PermulteImplicit(contact)
SIM019Uninitialized read: {name}\nMemory location: {location} ({location_type})\nPartition index: {partition_index}\nAddress interval: {address_interval}Check whether the memory was clobbered or left uninitialized. {prev_info}
SIM020Expected instruction using CC get-rank to provide replica groups, but got {opcode}(contact)
SIM021Expected insruction to provide at least one replica group(contact)
SIM022Unsupported NCCL topology {actual_tp_rank}, exected one of {expected_tp_ranks}(contact)
SIM023Unsupported NCCL Channel {actual_channel}, expected either 0, 1(contact)
SIM024Expected dma copy insruction to provide at least one replica group(contact)
SIM025Unsupported NCCL Channel {actual_channel}, expected either 0, 1(contact)
SIM026Expected insruction to provide at least one replica group(contact)
SIM027Expected dma copy insruction to provide at least one replica group(contact)
SIM028Expected instruction using CC get-rank to provide replica groups, but got {opcode}(contact)
SIM029Unsupported NCCL topology {actual_tp_rank}, exected one of {expected_tp_ranks}(contact)
SIM030Unsupported NCCL topology {actual_tp_rank}, exected one of {expected_tp_ranks}(contact)
SIM031Unsupported NCCL Channel {actual_channel}, expected either 0, 1(contact)
SIM032Unsupported NCCL topology {actual_tp_rank}, exected one of {expected_tp_ranks}(contact)
SIM033Instruction missing BasicBlock parent.(contact)
SIM034BasicBlock missing Function parent.(contact)
SIM035Function missing Module parent.(contact)
SIM036NCCL Topology is not available for {arch} architecture.(contact)
SIM037NCCL Topology is not available for {arch} architecture.(contact)
SIM038NCCL Topology is not available for {arch} architecture.(contact)
SIM039Expected pelican affine expr of kind CCGetRankExpr.(contact)
SIM040InstTensorScalarAffineSelect with NaN fill value should have exactly 2 arguments, but found {args}(contact)
SIM041InstTensorScalarAffineSelect with non-NaN fill value '{value}' should have exactly 1 argument, but found {args}(contact)
SIM042Remote semaphore update '{update}' not supported (yet)(contact)
SIM043Engine is out of buffer loop(contact)
SIM044Engine is out of buffer: {type}(contact)
SIM045More than one Return instruction in BasicBlock(contact)
SIM046Engine unassigned(contact)
SIM047Unable to read PWP tables from {path}(contact)
SIM048Memory location not found(contact)
SIM049Memory object not found(contact)
SIM050Memory location not found in Memory map(contact)
SIM051Alias memory location not found(contact)
SIM052Memory location '{location}' not found(contact)
SIM053Cannot access PSUM location {location} via DMA indirection(contact)
SIM054DRAM '{location}' must be allocated in order to reach here!(contact)
SIM055doCopyIndirect accumulate: entry size ({entry}) must be multiple of dtype ({dtype})(contact)
SIM056Output of {inst} has NaN when writing to {mloc} in subgraph {sg}(contact)
SIM057Output of {inst} has NaN when writing to {mloc} in subgraph {sg}(contact)
SIM058No such HW arch '{arch}' yet!(contact)
SIM059Memory type '{type}' not supported yet!(contact)
SIM060Indirect read should not read from Psum!(contact)
SIM061Target memory type '{type}' not supported yet!(contact)
SIM062Memory type '{type}' not supported yet!(contact)
SIM063Output of {inst} has NaN when writing to {mloc} in subgraph {sg}(contact)
SIM064Indirect write should not write to Psum!(contact)
SIM065Memory type '{type}' not supported yet!(contact)
SIM066Target memory type '{type}' not supported yet!(contact)
SIM067Target output has NaN(contact)
SIM068Memory type '{type}' not supported yet!(contact)
SIM069writeAccumulate shouldn't happen for SB!(contact)
SIM070writeAccumulate shouldn't happen for DRAM!(contact)
SIM071Memory type '{type}' not supported yet!(contact)
SIM072Memory object for '{location}' is null in readLocal(contact)
SIM073Memory location not found in Memory map(contact)
SIM074Uninitialized read: {inst}\nMemory location: {loc} ({type})Check if this Memory location has empty filename or has not been written to yet or the read/write reference count doesn't match
SIM075Uninitialized access of aliased input location: {alias}\nActual location accessed: {loc}Check if this Memory location has empty filename or has not been written to yet or the read/write reference count doesn't match
SIM076InstVisitor.writeOutputs(): Attempted to write output at MemoryLocation {name} to file, but the MemoryLocation has no associated filename.(contact)
SIM077InstVisitor.writeOutputs(): Invalid output filename {file} provided.(contact)
SIM078SB Data pinned by memory location {prev_loc}, which was written by {prev_inst}, is overwritten by instruction {cur_inst}, at {type} memory Partition index: {partition} and Address interval: {addr}(contact)
SIM079Execution hanged waiting on {waiters}(contact)
SIM080All VNC cores hanged(contact)
SIM081Double Event clear to event {id}(contact)
SIM082SendRecv does not support more than 1 input/output argument(contact)
SIM083SendRecv is not supported for global collective(contact)
SIM084Multiple LNCs are not supported yet(contact)
SIM085Number of input arguments must be the same as recvFomRank size(contact)
SIM086SendRecvCCE is not supported for global collective(contact)
SIM087Multiple LNCs are not supported yet(contact)
SIM088Expected input size ({in_count}) to match output size ({out_count})(contact)
SIM089Indirect index ({index}) must be less than IndirectDimMaxIndex ({max_index})(contact)
SIM090cast_accumulate shouldn't be called for Overwrite op(contact)
SIM091BIRSIM mismatch for tensors (symbolic memory simulation).List of failing tensors and respective histograms can be found in log-neuron-cc.txt.
SIM092BIRSIM mismatch for tensors (physical memory simulation).List of failing tensors and respective histograms can be found in log-neuron-cc.txt.
SIM093Expected instruction of type (DMACopy), got ({inst_type}) instead(contact)
SIM094Kernel simulation not supported for {kernel_name} yet(contact)
SIM095InstRangeSelect expects the input and output access patterns have the same data types, but input is {inType} type and output is {outType} type(contact)
SIM096InstRangeSelect expects the input and output access patterns to have the same number of elements, but input has {inputSize} elements and output has {outputSize} elements(contact)
SIM097InstRangeSelect expects the number of partitions of {bound} ({boundNumPartitions}) to be equal to input access pattern ({numInElementsPerPartition}) and shape to be one element per partition(contact)
SIM098InstRangeSelect expects {name} access patterns to be 3D, but it is {size}D(contact)
SIM099InstRangeSelect expects the {compare_op} to be one of isEQ, IsGT, IsGE, IsLT, IsLE(contact)
SIM100InstRangeSelect only supports max reduce op, but got {op} op(contact)
SIM101Non-scalar CopyPredicated cannot have reverse pred(contact)
SIM102Invalid AluOpType '{given}' for MemoryReductionOp. Must be 'bypass' or 'add'(contact)
SIM103Duplicate indices {indices} not allowed on DGE DstReduce instruction(contact)
SIM104Invalid CCE op type '{given}' for DGE DstReduce. Must be 'add'(contact)
SIM105InstCompareAndBranch does not support arguments of type {dtype}(contact)
SIM106Cannot read scalar value from {kind} argument, must be Physical AP, Register or Immediate(contact)
SIM107Dynamic offset of scalar DGE is out of bounds(contact)
SIM108InstDoWhile does not support arguments of type {dtype}(contact)
SIM109Golden: {golden_size} and BIRSim output: {birsim_output_size} sizes do not match.(contact)
SIM110Parameter index {paramsIndex} is equal to or greater than parameters per partition {paramsPerPartition} at partition {partitionOffset} and index {indicesOffset}(contact)
SIM111Simulation with kernel inline only works after unroll pass.(contact)
SIM112Static loop axis not found in loop map!(contact)
SIM113Vector DGE instruction must have Indirect Arg ID as a top level expression!(contact)
SIM114CopyPredicatedReduce requires third Immediate Argument to be Scalar.(contact)
SIM115CopyPredicatedScalar Immediate value type ({imm_type}) must match output type ({out_type})(contact)
SIM116CopyPredicatedScalar Immediate value type ({imm_type}) must match output type ({out_type})(contact)
SIM117CopyPredicated instruction failed. Predicate and Out must have same shape but have: {pred_part}x{pred_num} and {out_part}x{out_num} respectively(contact)
SIM118CopyPredicated instruction failed. Predicate and Src must have same number of elements but have: {pred_num} and {src_num} respectively(contact)
SIM119CopyPredicated instruction failed. Predicate and OptionalImmediate must have same number of elements but have: {pred_num} and {opt_num} respectively(contact)
SIM120Reduction Outputs must be cast to float32 before usage. Given: {actual}(contact)
SIM121Must be FP32. Given: {actual}(contact)
SIM122Gold {gold_file} not found for output {output_name}(contact)
SIM123DevicePrint only supports 2D tensor(contact)
SIM124DevicePrint does not support arguments of type {dtype}(contact)
SIM125Cannot find Memory location {loc} information in tensormap(contact)
SIM126No golden found for output location {loc}(contact)
SIM127AllEngineBarrier BIRSim simulation internal error: expected AllEngineBarrier in instruction stream(contact)
SIM128Dynamic DMA transpose indices can only be OOB after first OOB index found. But index {index} is {actual}, which is not OOB(contact)
SIM129Dynamic DMA transpose number of non-OOB indices must be a multiple of 16, but it is {actual}(contact)
SIM130AllEngineBarrier instruction should be assigned to engine ALL, but it is assigned to engine {engine}(contact)
SIM131AllEngineBarrier BIRSim simulation internal error: waiters should be empty(contact)
SIM132AllEngineBarrier BIRSim simulation internal error: unexpected end of instruction stream(contact)
SIM133AllEngineBarrier BIRSim simulation internal error: not all engines are waiting on the same AllEngineBarrier, inst1.name = {inst1_name}, inst2.name = {inst2_name}(contact)
SIM134AllEngineBarrier BIRSim simulation internal error: unexpected nullptr instruction(contact)
SIM135AllEngineBarrier instruction should not have wait/update conditions(contact)
SIM200GPSIMDSB2SB instructions are only supported for LNC=2 (not {LNC}).(contact)
SIM201Only local SendRecvs can be mapped to use GPSIMD's DMA engine.(contact)
SIM202Unknown shared memory location {location}.(contact)
SIM203Invalid seed for engine: GPSIMD. Seed must be an nx6 tensor. Got seed tensor of dimensions nx{dim}(contact)
SIM204Invalid seed for engine: Vector. Seed must be an nx24 tensor. Got seed tensor of dimensions nx{dim}(contact)
SIM205Output dtype must be int32. Got dtype: {type}(contact)
SIM206Input dtype must be int32. Got dtype: {type}(contact)

SIO — bare (1 codes)

CodeCauseResolution
SIO001Register {reg_name} is used in multiple BasicBlocks - Registers must be defined and used within a single BasicBlock(contact)

SQI — bare (7 codes)

CodeCauseResolution
SQI001Expect only Spill/Reload queues are used in loop body, but found {queueName} of type {queueType}(contact)
SQI002SQI switches to a queue that does not exist(contact)
SQI003SQI instance should not be NULL(contact)
SQI004SQI instance's parent should be a DMAQueue(contact)
SQI005RQI resets to a queue that does not exist(contact)
SQI006RQI instance should not be NULL(contact)
SQI007RQI instance's parent should be a DMAQueue(contact)

SSA — bare (3 codes)

CodeCauseResolution
SSA001Got a formerly shared memory location ({mloc_name}) which is unallocated and does not have a remote local target; all such memory locations should be allocated at this point(contact)
SSA002Memory location ({mloc_name}) is unallocated, and its corresponding remote target ({remote_mloc_name}) on core {core} is also unallocated(contact)
SSA003For the current memory location, could not find its remote target {remote_name}(contact)

TCE — bare (3 codes)

CodeCauseResolution
TCE720Unexpected Access Pattern(contact)
TCE721Expected >= 2 arguments(contact)
TCE722Unexpected copy type(contact)

TEN — bare (2 codes)

CodeCauseResolution
TEN404Internal tensorizer error: {baseMessage}(contact)
TEN405Internal tensorizer maximum recursion depth exceeded, {stack}(contact)

TST — bare (7 codes)

CodeCauseResolution
TST001Testing basic assertionSkip this test
TST002Testing with string ({msg}) sourceSkip this test
TST003Testing with DebugLocation sourceSkip this test
TST004Testing with DebugLocation and file info sourceSkip this test
TST005Testing with Instruction no DebugLocation sourceSkip this test
TST006Testing with Instruction w/DebugLocation sourceSkip this test
TST007Testing with Instruction w/DebugLocation and file info sourceSkip this test

VNS — bare (25 codes)

CodeCauseResolution
VNS601
VNS604(contact)
VNS605This huge DRAM tensor has a cross-partition access! You guessed the wrong block dims maybe.(contact)
VNS606This huge DRAM tensor has a cross-partition access! You guessed the wrong block dims maybe.(contact)
VNS607(contact)
VNS608Cannot vnsplit multi-block functions.(contact)
VNS609Cannot vnsplit multi-block functions.(contact)
VNS610Cannot vnsplit multi-block functions.(contact)
VNS611Try to remove instruction not in bb.(contact)
VNS612Try to remove MemoryLocation not in function.(contact)
VNS613(contact)
VNS614Unitialized vnode.(contact)
VNS615Empty writeAPs in collectStats()!(contact)
VNS616Uninitialized read(contact)
VNS617Unsupported output ap type.(contact)
VNS618unsupported arg type!(contact)
VNS620Cannot shrink dn on multi-block functions.(contact)
VNS621Write AP into DN is never read.\n(contact)
VNS622ap's parent inst has to be either AbstractCopy or Load!\n(contact)
VNS623We've checked pred_pap is also partition contiguous earlier?!(contact)
VNS624ap's parent inst has to be either AbstractCopy or Load!\n(contact)
VNS625ap and its pred ap must have the same Pattern size.\n(contact)
VNS627Cannot vertical fuse multi-block functions.(contact)
VNS628try to remove MemoryLocation not in function.(contact)
VNS629try to remove instruction not in bb.(contact)

XAQ — bare (2 codes)

CodeCauseResolution
XAQ001Unexpected tensor kind {kind}(contact)
XAQ002Expected at most {expected} HWDGE queues (there are {num})(contact)

XBI — bare (2 codes)

CodeCauseResolution
XBI001Internal compilation error: unexpected parallelism variant(contact)
XBI002Internal compilation error: "skip-pass-before" specified but not visited(contact)

XCG — bare (217 codes)

CodeCauseResolution
XCG001DMA block must have sync info(contact)
XCG002DMA block must have one local sync update(contact)
XCG003DMA block must by synchronized using semaphore(contact)
XCG004InstGroupResetSemaphores must be on EngineType::ALL, but it is on EngineType::{actual}(contact)
XCG005InstTensorCompletion must be on a datapath engine, but it is on EngineType::{actual}(contact)
XCG006InstAllEngineBarrier must be a datapath engine. Given EngineType::{actual}.Assign to EngineType::ALL and allow it to be expanded to all engines by the expand_all_engine pass.
XCG008One of CompareAndBranch's target must be the lexicographical next block in the BasicBlockHolder ({next})!(contact)
XCG009CompareAndBranch's second argument must be a RegisterAccess or an ImmediateValue (got {actual})(contact)
XCG010InstCompareAndBranch cannot be on EngineType::ALL; it must be on a TPB engine, and should have been expanded by expand_all_engine(contact)
XCG011InstUnconditionalBranch cannot be on EngineType::ALL; it must be on a TPB engine, and should have been expanded by expand_all_engine(contact)
XCG012InstDevicePrint should not make it to codegen. It should have been expanded and removed by expand_device_print(contact)
XCG033Internal compilation failed: {msg}(contact)
XCG100Unknown sync wait mode(contact)
XCG101Unknown sync update mode(contact)
XCG102Unknown EngineType(contact)
XCG103Cannot convert Accum cmd to Reduce cmd. Unknown Accum cmd type(contact)
XCG149Couldn't encode completion info for extended instruction because it has no APs(contact)
XCG150Couldn't encode completion info for extended instruction because number of partitions/ports must be 0 < x <= 128 and a multiple of 16(contact)
XCG151Couldn't encode completion info for extended instruction because base partition must be a multiple of 16(contact)
XCG152Couldn't encode completion info for extended instruction because instruction reads and writes a different number of partitions(contact)
XCG164ISA check failed(contact)
XCG166Instruction engine check failed ({engine})(contact)
XCG400LNC {actual} is not supported for GPSIMDSB2SB, which only supports LNC {supported}(contact)
XCG800Unknown sync wait mode(contact)
XCG801Unknown sync update mode(contact)
XCG802Unknown EngineType(contact)
XCG804Function name too long: {funcName}(contact)
XCG805Unimplemented Activation function '{op}'(contact)
XCG806Unimplemented ALU opcode '{op}'(contact)
XCG807Unimplemented ALU opcode '{op}'(contact)
XCG808Unimplemented ALU opcode '{op}'(contact)
XCG809Unimplemented Axis List Type '{type}'(contact)
XCG810Unimplemented Dtype {type}(contact)
XCG811Unimplemented Compare Branch operator '{op}'(contact)
XCG812Queue instance name too long: {queName}(contact)
XCG813Expect axis type to be XYZWC(contact)
XCG814Unimplemented Engine Accumulator Command {acc}(contact)
XCG815Estimated peak HBM usage ({value}) exceeds {hbm_limit}GB. Neff might be unable to load on chip. If you believe this estimation to be inaccurate, you can disable the check using: --internal-backend-options=' --disable-hbm-usage-check '(contact)
XCG816TensorScalarPtr must have 2 or 3 inputs(contact)
XCG817TensorScalarPtr must have 2 or 3 inputs(contact)
XCG818Pool input AP must have 5 dimensions(contact)
XCG819Expect BNStatsAggr output 2 elements(contact)
XCG820Expect BNGradients mean and var type to be same(contact)
XCG821Malformed access pattern(contact)
XCG822Params should be 4 floats(contact)
XCG823StreamShuffle Mask must have 32 elements(contact)
XCG824StreamTranspose src pattern num_elem must be 1(contact)
XCG825StreamTranspose op must have 1-D mem_pattern or read a factor of 32 elements(contact)
XCG826Collective {kind} does not support an ALU op(contact)
XCG827AllToAll does not support an ALU op(contact)
XCG828float32r matmult inputs must have same dtype(contact)
XCG829two pass matmult doesn't reuse previous weight(contact)
XCG830two pass matmult doesn't reuse previous weight(contact)
XCG831float32r matmult inputs must have same dtype(contact)
XCG832TensorSave destination must be SB on trn1, due to a HW bug(contact)
XCG833Unimplemented Compare Branch operator {op}(contact)
XCG834Function name too long: {targetName}(contact)
XCG835Temporary Limitation: Tensor library only supports 1 partition now!(contact)
XCG836Queue instance name too long: {Name}(contact)
XCG837MAX8 Instruction writes exactly 8 output elements per partition(contact)
XCG838MAX8 Instruction expects at least 8 input elements per partition(contact)
XCG839MAX8 Instruction supports at most 16384 input elements per partition(contact)
XCG840MAX8 Instruction doesn't allow more than 5 dimensions in input access pattern(contact)
XCG841MAX8 Instruction doesn't allow more than 3 dimensions in output access pattern(contact)
XCG842MAX8 Instruction expects the same number of partitions in input and output(contact)
XCG843MAX8 Instruction output is restricted to 2 dimensions(contact)
XCG844MATCH_VALUE_LOAD expects exactly 8 input elements per partition(contact)
XCG845MATCH_VALUE_LOAD Instruction doesn't allow more than 3 dimensions in input access pattern(contact)
XCG846expecting physical ap at in0 of MATCH_VALUE_LOAD Instruction(contact)
XCG847FIND_INDEX8 Instruction writes exactly 8 output elements per partition(contact)
XCG848FIND_INDEX8 Instruction expects at least 8 input elements per partition(contact)
XCG849FIND_INDEX8 Instruction supports at most 16384 input elements per partition(contact)
XCG850FIND_INDEX8 Instruction doesn't allow more than 5 dimensions in input access pattern(contact)
XCG851FIND_INDEX8 Instruction doesn't allow more than 3 dimensions in output access pattern(contact)
XCG852FIND_INDEX8 Instruction expects the same number of partitions in input and output(contact)
XCG853FIND_INDEX8 Instruction output is restricted to 2 dimensions(contact)
XCG854MATCH_VALUE_LOAD expects exactly 8 input elements per partition(contact)
XCG855MATCH_VALUE_LOAD Instruction doesn't allow more than 3 dimensions in input access pattern(contact)
XCG856MATCH_REPLACE8 Instruction expects at least 8 input elements per partition(contact)
XCG857MATCH_REPLACE8 Instruction supports at most 16384 input elements per partition(contact)
XCG858MATCH_REPLACE8 Instruction expects at least 8 input elements per partition(contact)
XCG859MATCH_REPLACE8 Instruction supports at most 16384 input elements per partition(contact)
XCG860MATCH_REPLACE8 Instruction doesn't allow more than 5 dimensions in input access pattern(contact)
XCG861MATCH_REPLACE8 Instruction doesn't allow more than 5 dimensions in output access pattern(contact)
XCG862MATCH_REPLACE8 Instruction expects the same number of partitions in input and output(contact)
XCG863ISA check failed(contact)
XCG864ISA check failed(contact)
XCG865Unsupported MatmultPerfMode {mode}(contact)
XCG866Target does not support Matmul replication(contact)
XCG867Matmult's Fmap ({ifmap}) and Weight ({weight}) partitions must match(contact)
XCG868Matmult's Output ({output}) partitions must match Weight ({weight}) elements per partition(contact)
XCG869Matmult's Output ({output}) and Fmap ({ifmap}) elements per partition must match(contact)
XCG870Invalid register ({regId})(contact)
XCG871Could not open {file} - {strerr}Verify that your filesystem is writeable
XCG872ISA mem_pattern should only be SB/PSUM accesses(contact)
XCG873InstTensorScalarAffineSelect should have NaN fill value in codegen, instead has {value}(contact)
XCG874CollectiveKind {kind} not supported(contact)
XCG875CollectiveKind {kind} not supported(contact)
XCG876CollectiveKind {kind} requires 1 arguments per output, got num_arguments={num_arguments}, num_outputs={num_outputs}(contact)
XCG877CollectiveKind {kind} requires 2 arguments per output, got num_arguments={num_arguments}, num_outputs={num_outputs}(contact)
XCG878Tensor2D AP must have 2 dimensions, got {dim}(contact)
XCG879CollectiveKind {kind} does not support arguments outside of DRAM(contact)
XCG880CollectiveKind {kind} does not support outputs outside of DRAM(contact)
XCG890Matmul accumulation flag not valid, may caused by is still auto but should be specific value(contact)
XCG891Matmul accumulation flag not valid, may caused by is still auto but should be specific value(contact)
XCG892MatMulSparse accumulation flag not valid, may caused by is still auto but should be specific value(contact)
XCG893MatmulMX accumulation flag not valid, may caused by is still auto but should be specific value(contact)
XCG900Unknown sync wait mode(contact)
XCG901Unknown sync update mode(contact)
XCG902Unknown EngineType(contact)
XCG903Unimplemented Activation function {op}(contact)
XCG904Unimplemented ALU opcode {op}(contact)
XCG905Unimplemented ALU opcode {op}(contact)
XCG906Unimplemented ALU opcode {op}(contact)
XCG907Unimplemented Axis List Type {type}(contact)
XCG908Unimplemented Dtype {type}(contact)
XCG909Unimplemented Compare Branch operator {op}(contact)
XCG910Queue name too long: {queName}(contact)
XCG911Expect axis type to be XYZWC(contact)
XCG914Permute does not support an ALU op(contact)
XCG915
XCG916TensorScalarPtr must have 2 or 3 inputs(contact)
XCG917TensorScalarPtr must have 2 or 3 inputs(contact)
XCG918Unimplemented: TensorSave dest is not DRAM(contact)
XCG919bias can only take imm_fp32 when activation function has >= 128 code(contact)
XCG920Pool input AP must have 5 dimensions(contact)
XCG921out.getNumElementsPerPartition() == 2(contact)
XCG922Expect BNGradients mean and var type to be same(contact)
XCG923Malformed access pattern(contact)
XCG924Params should be 4 floats(contact)
XCG925StreamShuffle Mask must have 32 elements(contact)
XCG926StreamTranspose src pattern num_elem must be 1(contact)
XCG927StreamTranspose op must have 1-D mem_pattern or read a factor of 32 elements(contact)
XCG928AllGather does not support an ALU op(contact)
XCG929Queue instance name too long: {Name}(contact)
XCG930float32r matmult inputs must have same dtype(contact)
XCG931float32r matmult inputs must have same dtype(contact)
XCG932two pass matmult doesn't reuse previous weight(contact)
XCG933two pass matmult doesn't reuse previous weight(contact)
XCG934float32r matmult inputs must have same dtype(contact)
XCG935float32r matmult inputs must have same dtype(contact)
XCG937MAX8 Instruction writes exactly 8 output elements per partition(contact)
XCG938MAX8 Instruction expects at least 8 input elements per partition(contact)
XCG939MAX8 Instruction supports at most 16384 input elements per partition(contact)
XCG940MAX8 Instruction doesn't allow more than 5 dimensions in input access pattern(contact)
XCG941MAX8 Instruction doesn't allow more than 3 dimensions in output access pattern(contact)
XCG942MAX8 Instruction expects the same number of partitions in input and output(contact)
XCG943MAX8 Instruction output is restricted to 2 dimensions(contact)
XCG944MATCH_VALUE_LOAD expects exactly 8 input elements per partition(contact)
XCG945MATCH_VALUE_LOAD Instruction doesn't allow more than 3 dimensions in input access pattern(contact)
XCG946expecting physical ap at in0 of MATCH_VALUE_LOAD Instruction(contact)
XCG947FIND_INDEX8 Instruction writes exactly 8 output elements per partition(contact)
XCG948FIND_INDEX8 Instruction expects at least 8 input elements per partition(contact)
XCG949FIND_INDEX8 Instruction supports at most 16384 input elements per partition(contact)
XCG950FIND_INDEX8 Instruction doesn't allow more than 5 dimensions in input access pattern(contact)
XCG951FIND_INDEX8 Instruction doesn't allow more than 3 dimensions in output access pattern(contact)
XCG952FIND_INDEX8 Instruction expects the same number of partitions in input and output(contact)
XCG953FIND_INDEX8 Instruction output is restricted to 2 dimensions(contact)
XCG954MATCH_VALUE_LOAD expects exactly 8 input elements per partition(contact)
XCG955MATCH_VALUE_LOAD Instruction doesn't allow more than 3 dimensions in input access pattern(contact)
XCG956MATCH_REPLACE8 Instruction expects at least 8 input elements per partition(contact)
XCG957MATCH_REPLACE8 Instruction supports at most 16384 input elements per partition(contact)
XCG958MATCH_REPLACE8 Instruction expects at least 8 input elements per partition(contact)
XCG959MATCH_REPLACE8 Instruction supports at most 16384 input elements per partition(contact)
XCG960MATCH_REPLACE8 Instruction doesn't allow more than 5 dimensions in input access pattern(contact)
XCG961MATCH_REPLACE8 Instruction doesn't allow more than 5 dimensions in output access pattern(contact)
XCG962MATCH_REPLACE8 Instruction expects the same number of partitions in input and output(contact)
XCG963Unimplemented Compare Branch operator {op}(contact)
XCG964Function name too long: {targetName}(contact)
XCG965Instruction engine check failed ({engine})(contact)
XCG966Instruction engine check failed ({engine})(contact)
XCG967Value that is out-of-bounds for corresponding ISA field found: {details}(contact)
XCG970Innermost dimension step of transpose matmult output must be 1(contact)
XCG971Outer dimension step sizes of transpose matmult output must be even or 1(contact)
XCG972Invalid ALU {op} for RangeSelect instruction, expect one of IsEQ, isGT, IsGE, IsLE, IsLT(contact)
XCG973RangeSelect instruction expects the input and output access patterns to have the same data type (input is of {dtype0} type, but output is of {dtype1} type)(contact)
XCG974RangeSelect instruction expects the input and output access patterns to be one of float32, bfloat32, float16, float8 data types, but is {dtype} type(contact)
XCG975RangeSelect instruction expects the input and output access patterns to access the same number of elements (input has {inNum} elements, output has {outNum} elements)(contact)
XCG976RangeSelect instruction expects bound to be float32 type, but is {dtype} type(contact)
XCG977RangeSelect instruction base to be within DVE_RANGE = {DVE_RANGE}(contact)
XCG978RangeSelect instruction bound ({boundName}) + num bytes in partition of access pattern ({apName}) to be within DVE_RANGE = {DVE_RANGE}(contact)
XCG979RangeSelect instruction expects number of partitions accessed to be > 0 and <= 128, but {numPartitions} were accessed(contact)
XCG980InstRangeSelect only supports max reduce op, but received {opType} op(contact)
XCG981Matmult instruction must have PE tile position fully set or fully unset (let compiler decide)Consult nc_matmult documentation for valid tile positions
XCG983Matmult instruction must have PE tile size fully set or fully unset (let compiler decide)Consult nc_matmult documentation for valid tile sizes
XCG985CopyPredicated does not support casting in scalar mode(contact)
XCG986Invalid AluOpType '{given}' for DGE_COMPUTE_OP. Must be 'bypass' or 'add'(contact)
XCG987DVE_READ_INDICES Instruction expects total number of elements to be 8, but got {actual}(contact)
XCG988DVE_READ_INDICES Instruction expects AP type to be Dtype::Uint16 or Dtype::Uint32(contact)
XCG989DVE_READ_INDICES Instruction expects number of partitions to be between 1 - 128(contact)
XCG990Detected duplicate engine ({engine}) and label ({label}) for basic block {bb1} and {bb2}. Basic block must have unique (engine, label) pair(contact)
XCG991Expect ResetQueueInstance to have same instance as active instance on queue, but is not(contact)
XCG992Expect module attribute neff_feature_SQI_no_rearm to be set by previous SwitchQueueInstance instruction, but is not(contact)
XCG993Unimplemented ALU opcode '{op}'(contact)
XCG994CopyPredicatedReduce instruction immediate value must be a float32, but it is a {actual}(contact)
XCG995DveReadAccumulator output must be a floating-point type, but it is Dtype::{actual}(contact)
XCG996Unsupported branch hint outcome type {hint_outcome}(contact)
XCG997Expect function program counter for all datapath engines to be 1 for the preamble, but found {count} for {engine}(contact)
XCG998Instruction is not assigned an engine when updating program counter(contact)
XCG999branch label for BasicBlock {BasicBlockName} was not generated in enterFunction(contact)
XCG1000Cannot find program counter for location instruction {instName} engine {engineType}(contact)
XCG1001Expect there to be only one ISA instruction from BIR BranchHint instruction, but found multiple at PCs {oldPC} and {newPC}(contact)
XCG1002InstExit must have a datapath engine type after expansion, but it has EngineType::{actual}Please be sure that Instruction is assigned to ALL so that it can properly be expanded to all datapath engines by expand_all_engine
XCG1003CollectiveCompute does not have a streamID set. Make sure a compatible opt level which runs `InferStreamIds' is set.(contact)
XCG1004Matmult instruction has invalid PE row tile position with respect to the accessed start partition of stationary tensorConsult nc_matmult documentation for valid tile positions
XCG1005Matmult instruction has invalid PE column tile position with respect to the accessed start partition of PSUM output tensorConsult nc_matmult documentation for valid tile positions
XCG1006Matmult instruction has invalid PE row tile size with respect to the accessed start partition of stationary tensorConsult nc_matmult documentation for valid tile sizes
XCG1007Matmult instruction has invalid PE column tile size with respect to the accessed start partition of PSUM output tensorConsult nc_matmult documentation for valid tile sizes
XCG1008Compiler generates invalid PE row tile size {row_tile_size} for Matmult instruction with accessed start partition {in1_base_partition} of stationary tensor(contact)
XCG1009Compiler generates invalid PE column tile size {col_tile_size} for Matmult instruction with accessed start partition {out_base_partition} of PSUM output tensor(contact)
XCG1010Vector DGE using shape register has {dim} dimensions, shape register supports only 5 dimensionsConsult nc_matmult documentation for valid tile sizes
XCG1011Unrecognized UniqueTensorsType(contact)
XCG1012InstAllEngineBarrier must be a datapath engine. Given EngineType::{actual}.Assign to EngineType::ALL and allow it to be expanded to all engines by the expand_all_engine pass.
XCG1013CollectiveDimension {dimension} not supported(contact)
XCG1014Unrecognized EventSemaphoreClearMode {mode}(contact)
XCG1015InstGroupResetSemaphores must be a datapath engine. Given EngineType::{actual}.Assign to EngineType::ALL and allow it to be expanded to all engines by the expand_all_engine pass.
XCG1016InstGroupResetSemaphores has an empty SemaGroup vector. Should have been eliminated before this point.Ensure that instruction has EngineType::ALL so that expand_all_engines can run and expand or eliminate as needed.

XDR — bare (1 codes)

CodeCauseResolution
XDR001No state for predecessor instruction ({name}) available.Make sure graph is topologically sorted.

XEI — bare (8 codes)

CodeCauseResolution
XEI001Internal testing error, used in walrus pass test on module: {name}.Should only occur during walrus_pass_test.
XEI002Internal testing error, used in walrus pass test on modules: {names}.Should only occur during walrus_pass_test.
XEI418I'm a teapotLorem ipsum dolor sit amet, consectetur adipiscing elit, sed do eiusmod tempor incididunt ut labore et dolore magna aliqua.
XEI419I'm a teapotLorem ipsum dolor sit amet, consectetur adipiscing elit, sed do eiusmod tempor incididunt ut labore et dolore magna aliqua.
XEI420I'm a teapotLorem ipsum dolor sit amet, consectetur adipiscing elit, sed do eiusmod tempor incididunt ut labore et dolore magna aliqua.
XEI421I'm a teapotLorem ipsum dolor sit amet, consectetur adipiscing elit, sed do eiusmod tempor incididunt ut labore et dolore magna aliqua.
XEI422I'm a teapotLorem ipsum dolor sit amet, consectetur adipiscing elit, sed do eiusmod tempor incididunt ut labore et dolore magna aliqua.
XEI423A small formatted test with the number {test_arg_1} and some {test_arg_2}.(contact)

XGM — bare (6 codes)

CodeCauseResolution
XGM000Expected the subgraph {sgId} for all cores to have {expected} functions, but on core {ncId} it has {actual} functions(contact)
XGM001Expected subgraph{sgId} on core {ncId} to have a function named {name}(contact)
XGM002Expected function {name} in subgraph {sgId} to have {expected} basic blocks, but on core {ncId} it has {actual} basic blocks(contact)
XGM003Expected the subgraph {sgId} for all cores to have {expected} functions, but on core {ncId} it has {actual} functions(contact)
XGM004Expected subgraph{sgId} on core {ncId} to have a function named {name}(contact)
XGM005Expected function {name} in subgraph {sgId} to have {expected} basic blocks, but on core {ncId} it has {actual} basic blocks(contact)

XLB — bare (35 codes)

CodeCauseResolution
XLB000The number of modules should be the same as LNC. Current module size: {module}, current LNC: {lnc}(contact)
XLB001Identical function name must exist. The function name {funcName} is missing in module {modName}(contact)
XLB002The number of basicblocks should be the same across LNC modules. bbCnt: {bbNum1} in module {modName1}, {bbNum2} in module {modName2}(contact)
XLB003parallelBB size ({parallelBBSize}) should be the same as the number of LNC modules ({modNum})(contact)
XLB004The number of basicblocks should be the same across LNC modules. bbCnt: {bbNum1} in module {modName1}, {bbNum2} in module {modName2}(contact)
XLB005DummyCall should not appear in the middle of BIRs. LowerDMADummycall inst: {instName} has predecessor.(contact)
XLB006Non DMA instNode or non dummy core barrier instNode should have instruction pointer. InstNode {instNodeName} does not have instruction(contact)
XLB007Matmult instruction's engine should be always PE. Currently instruction {instName}'s engine is {engineName}(contact)
XLB008Dynamic DMA instruction should have an assigned queue. Currently instruction {instName} does not have an assigned queue.(contact)
XLB009Missing barrier check only supports LNC=2. The number of LNC modules: {modNum}(contact)
XLB010InstBarrierRangeMap should have been initialized(contact)
XLB011InstBarrierRangeMap should have been initialized(contact)
XLB012Max barrier id should be greater than Min barrier id. Current Max barrier id: {maxId} Min barrier id: {minId} for engine {engineName}.(contact)
XLB013Engine name should not be duplicated. Engine {engineName} appears twice.(contact)
XLB014Barrier range start should less than or equal to end. Current start id: {startId}, end id: {endId} of instruction {instName} for engine {engineName}(contact)
XLB015Barrier range start should less than or equal to end. Current start id: {startId}, end id: {endId} of instruction {instName} for engine {engineName}(contact)
XLB016Tensor {name} is not found in remote core {core}(contact)
XLB017Function {name} is not found in remote core {core}(contact)
XLB018Symbolic Access pattern exists on an argument of inst: {readerInst}.(contact)
XLB019Symbolic Access pattern exists on an output of inst: {writerInst}.(contact)
XLB020Unexpected Memory spec which is shared but allocated - Name: {memoryName}, type: {memoryType}, kind: {kind}, addressSpace: {addressSpace}, isAllocated: {isAllocated}.(contact)
XLB021Symbolic Access pattern exists on an argument of inst: {readerInst}.(contact)
XLB022Symbolic Access pattern exists on an output of inst: {writerInst}.(contact)
XLB023Symbolic Access pattern exists on an argument of inst: {readerInst}.(contact)
XLB024Symbolic Access pattern exists on an output of inst: {writerInst}.(contact)
XLB025Remote target for non-SB and non-DRAM - Name: {memoryName}, type: {memoryType}, kind: {kind}, addressSpace: {addressSpace}, isAllocated: {isAllocated}.(contact)
XLB026Unexpected Memory spec - Name: {memoryName}, type: {memoryType}, kind: {kind}, addressSpace: {addressSpace}, isAllocated: {isAllocated}.(contact)
XLB027CoreBarrier instruction happens other than ACT/PE/DVE/SP/Pool engines. engine name: {engineName}.(contact)
XLB028Collective Compute instruction should have an assigned queue. Currently instruction {instName} does not have an assigned queue.(contact)
XLB029CoreBarrier appears in core 0 for {engineName}, but core 1 does not have any CoreBarrier in the same engine.(contact)
XLB030Dynamic offset expression size is not 1 from dynamic AP.(contact)
XLB031Dynamic offset expression size is not 1 from dynamic AP.(contact)
XLB032Indirect arg id does not fit to the number of arguments in input of {instName}.(contact)
XLB033Indirect arg id does not fit to the number of arguments in input of {instName}.(contact)
XLB034Found {races} pair(s) of instructions missing core barriers.(contact)

XLS — bare (5 codes)

CodeCauseResolution
XLS001Unable to clone or move module artifact directory when concretizing module from source={src} to destination={dst}. Error obtained: {error}(contact)
XLS002The artifact directory path {path} associated with the input module does not exist on disk.(contact)
XLS003A directory already exists at {path}. Therefore, cannot create an artifact directory for concrete module(contact)
XLS004Unable to clone or move module artifact directory when concretizing module from source={src} to destination={dst}(contact)
XLS005Error when copying {src} to {dst}. Error obtained: {error}(contact)

XLV — bare (35 codes)

CodeCauseResolution
XLV000Expected {expected} core barriers in subgraph {sgId}, but found {actual} on core {ncId}(contact)
XLV001Expected core barriers to have id {expected}, but core barrier {expected} in subgraph {sgId} on core {ncId} has id {actual}(contact)
XLV002Expected core barriers to have name '{expected}', but core barrier {expectedIdx} in subgraph {sgId} on core {ncId} has name '{actual}'. This likely indicates mismatching names on a local SendRecv or AllReduce instruction in the input program - the corresponding instructions on different cores must have the same name.(contact)
XLV003lnc_verifier failed with the following errors:{errors}(contact)
XLV004Subgraph {sgId} expected to find function {name}, but it is missing on core {ncId}.(contact)
XLV005Function {name} in subgraph {sgId} has {actual_size} basic blocks for core {ncId}, but expected {expected_size} basic blocks.(contact)
XLV006Subgraph {sgId} was expected to have {expected_size} functions, but core {ncId} has {actual_size} functions.(contact)
XLV007Expected block name {expected_block_name}, but got {actual_block_name} on core {ncId} in function {name} of subgraph {sgId}.(contact)
XLV008Remote target memory location {remote_target_name} from subgraph {sgId} not found in core {ncId}.(contact)
XLV009Shared tensor with address {address} of subgraph {sgId} missing on core {ncId}.(contact)
XLV010The shared memory location {name} in subgraph {sgId} on core {ncId} doesn't have remoteLocalTarget set.(contact)
XLV011Expected {expected} GPSIMDSB2SBs in subgraph {sgId}, but found {actual} on core {ncId}(contact)
XLV012Modules may not have both concrete neuron core ids and symbolic shard ids, but both are in function {func_name} for core {nc_id}(contact)
XLV013Only one symbolic shard id is expected per function. Multiple found instead in function {func_name}.(contact)
XLV014The execution of the core barrier instruction {cb} is masked out on some lnc cores.(contact)
XLV015Mismatching GPSIMDSB2SB instructions across LNC cores.(contact)
XLV016Multiple modules without neuron core ids found.(contact)
XLV017Data tensor in tensor indirect AP must be in either PSUM or SBUF, tensor was in {mem_type} instead(contact)
XLV018Expected {indices} elements in tensor indirect AP, but found {cols} instead(contact)
XLV019Index tensor in tensor indirect AP must be in SBUF and have type uint16, index is in {idx_mem} and has type {dtype} instead(contact)
XLV020Index AP for tensor indirection must access number of partitions divisible by 16 and start at partition divisible by 32, Index AP accesses {idx_accessed} starting at partition {idx_base} instead(contact)
XLV021Index AP for tensor indirection must access more or equal partitions than the tensor indirect AP, index AP: {idx_accessed} vs tensor indirect AP {parts}(contact)
XLV022Indirect Tensor AP cannot be on activation engine with PSUM indirect destination(contact)
XLV023First source AP and index AP for tensor indirection must start in the same quadrant for scatter style indirection. Data AP starts at partition {base} vs index AP {idx_base}(contact)
XLV024Source data AP must be contiguous for scatter style indirection(contact)
XLV025Data AP and index AP for gather style tensor indirection must start in same quadrant, data AP starts at {base} vs index AP starts at {idx_base}(contact)
XLV026For scatter and gather tensor indirection, index tensors must be in same quadrant, src index base partition {src_idx_base} vs dst index base partition {dst_idx_base}(contact)
XLV027Data tensor in symbolic tensor indirect AP must be in either PSUM or SBUF, tensor was in {mem_type} instead(contact)
XLV028Number of indices ({indices}) cannot exceed the total number of elements per partition group ({elems_per_partition_group}) in tensor indirect index AP(contact)
XLV029Expected {indices} elements in symbolic tensor indirect AP, but found {cols} instead(contact)
XLV030Index tensor in symbolic tensor indirect AP must be in SBUF and have type uint16, index is in {idx_mem} and has type {dtype} instead(contact)
XLV031Index AP for symbolic tensor indirection must access more or equal partitions than the symbolic tensor indirect AP, index AP: {idx_accessed} vs tensor indirect AP {parts}(contact)
XLV032Symbolic indirect Tensor AP cannot be on activation engine with PSUM indirect destination(contact)
XLV033Source symbolic data AP must be contiguous for scatter style indirection(contact)
XLV034Number of indices ({indices}) cannot exceed the total number of elements per partition group ({elems_per_partition_group}) in symbolic tensor indirect index AP(contact)

XNP — bare (17 codes)

CodeCauseResolution
XNP001Unknown EngineType: {engine}(contact)
XNP002Unknown DMA queue type(contact)
XNP003Wrong arch type is used, must use CoreV2 or newer(contact)
XNP004Expected all subgraphs to have the same arch {target_arch}, but module {mod_name} targets {arch}(contact)
XNP016Expected description file '{path}' does not exist(contact)
XNP017Missing file name for DVE table entry.(contact)
XNP018Missing act_info for used func set {usedActSet}(contact)
XNP019Unable to open file '{path}'(contact)
XNP020Unable to open file '{srcPath}'(contact)
XNP021Compiler did not generate an activation entry.(contact)
XNP022Unexpected non-CC tensor {memLoc} between ranges [0 - {range_start}) or [{range_end} - {high_water_mark}), but got {addr}(contact)
XNP023Expected function {function} to have hbm related function attributes(contact)
XNP024Expected collective tensor addresses to start at {range_start}, but found one at {addr}.(contact)
XNP025Expected non-collective tensor addresses to start at {main_start}, but found one at {addr}.(contact)
XNP026DRAM tensor '{name}' has unexpected MemoryAddressSpace::{space}, only Local is supported at this point.(contact)
XNP027Expected {vnc_nc_count} modules, but got {module_size} in neff_packager(contact)
XNP028Unrecognized runtime reserve memory type(contact)

XRA — bare (2 codes)

CodeCauseResolution
XRA001Expected number of target cores {core_count} to match number of subgraphs {mod_count}(contact)
XRA002Expected to find {remote} on target core {core}, but got null for {local}(contact)

XRO — bare (2 codes)

CodeCauseResolution
XRO001Undefined DRAM Memloc {mem}(contact)
XRO002Undefined SB Memloc {mem}(contact)

XTP — bare (1 codes)

CodeCauseResolution
XTP002Number of instructions ({insts}) is over the threshold ({threshold}). Tiling could potentially do a better job.(contact)

XVL — bare (4 codes)

CodeCauseResolution
XVL001Expected number of target cores {core_count} to match number of subgraphs {mod_count}(contact)
XVL002Expected to find {remote} on target core {core}, but got null for {local}(contact)
XVL003Expected all subgraphs to share the same number of DMA blocks that modify remote tensors, but found {sizes}(contact)
XVL004Expected to find local semphore update in instruction {name}(contact)

XXU — bare (3 codes)

CodeCauseResolution
XXU001Immediate value for addRegAluOneRegArgOneImm ({value}) does not fit in 32-bits(contact)
XXU002{inst_name}: Invalid operation - tile position only available for Matmult(contact)
XXU003{inst_name}: Source and destination must be allocated before calculating tile position(contact)

5. NKI err_* Funnel (sema)

The NKI front-end semantic-analysis layer (nki/compiler/backends/neuron/sema.so, md5 755c15de65e6cc00083fab0e5838e4a0) validates every nl.*/nisa.* API call's operands during kernel tracing, before lowering to penguin.ir. On failure it raises a user-facing exception. The ~150 diagnostic symbols form a three-tier funnel: check_* validators (32) → assert_* engine helpers (24) → err_* diagnostic leaves (78). Every err_* is wrapped by sema_err_url, which appends the docs link https://awsdocs-neuron.readthedocs-hosted.com/en/latest/nki/api/nki.errors.html. (D-W05; full per-diagnostic reference on §6.4.5.)

NOTE — these are a fourth, separate surface: err_* raise Python exceptions (NKISyntaxError / ValueError / TypeError / IndexError / RuntimeError) that propagate up to the nki.jit driver and are then surfaced by the A08/A13 logger layers — i.e. the ErrorMessages/NCC_* machinery sits above this layer and reports what it raises. All 78 err_* are hard compile-time errors; the lone non-fatal diagnostic is warn_block_dimension_is_deprecated (DeprecationWarning).

5a. err_* leaves by concern (78)

Addresses are sema.so file VAs; line numbers are sema.py source lines (Cython AddTraceback). exc = the raised exception class (NKISyntaxError* = a Cython-cached class slot whose literal name was not recovered; NKISyntaxError is the documented NKI default, class at 0xf6c20).

err_*Addr · sema.py:LexcMessage (interpolated)
err_num_partition_exceed_arch_limit0x99e90 · 2420ValueErrornumber of partitions <par_dim> exceed architecture limitation of <max_p> (128)
err_num_partition_mismatch0x807f0 · 2398ValueErrornumber of partitions mismatch in parameters (<shapes>)
err_size_of_dimension_exceed_arch_limit0x77500 · 2366ValueErrorsize of dimension <dim> in '<name><shape>' of '<api>' exceed architecture limitation of <max_size>
err_exceed_max_supported_dimension0xbf3a0 · 2343ValueError'<name><shape>' … exceed max supported number of dimension <max_rank>
err_stack_overflow_sbuf0x813a0 · 1763RuntimeErrorstack overflow: required sbuf size <sbuf_size> exceeds available sbuf size <target>
err_stack_overflow_psum0x81f60 · 1756RuntimeErrorstack overflow: required psum banks <n> exceeds available psum banks <target>
err_valid_size0x63bd0 · 2867NKISyntaxError*<name> size must be in [<valid_size>]
err_tile_shape_mismatch0x4ba30 · 1330NKISyntaxError*tile shape mismatch, expected '<expected_shape>' … <name>/<tile>
err_annotation_shape_mismatch0x4cd30 · 2441TypeErrorshape of \` does not match the expected shape of <annotation_shape>`
err_param_shape_mismatch0x90430 · 360ValueErrorParameter shapes (<shapes>) … has mismatched shapes
err_param_shape_incompatible_with_numpy0x8f6d0 · 354ValueErrorParameter shapes (<shapes>) … could not be broadcast together
err_param_shape_incompatible_with_matmul0xee450 · 343ValueErrormatmul contraction-shape incompatibility (<transpose_x>, <shapes>)
err_store_dst_shape_smaller_than_other_shape0xbe290 · 1958ValueErrorIllegal assignment destination shape in '<expr>': … '<dst_name>' is smaller than other parameter shape <shapes>
err_indices_shape_mismatch0x705c0 · 1985IndexErrorshape mismatch: indices indexing tensor <tensor_name> … <shapes>
err_leading_dimension_of_tensor_must_be_partition0x7f070 · 2519NKISyntaxError*The leading dimension of SBUF/PSUM tensors must be the partition dimension. The tensor has shape <tensor_shape>
err_supported_operand_dtype0xce320 · 332TypeError'<name>' … expected one of the following dtypes: <expected_dtypes>
err_unsupported_dtype_value0xd3ec0 · 327ValueErrorUnsupported dtype '<dtype_value>' …, expected one of: <expected_dtypes>
err_unsupported_args0x628b0 · 313TypeError'<api>' with unsupported arguments on nki tensor: <e>/<sign>
err_activation_bias_invalid_type0x7fc30 · 1573TypeError'bias' param of '<instr_name>' must be a vector of type float32, float16, or bfloat16, got '<dtype>'
err_activation_scale_invalid_type0x84330 · 1558TypeError'scale' param of '<instr_name>' must be a scalar or vector of type float32, got '<dtype>'
err_activation_scale_scalar_or_vector0x4c6a0 · 1540NKISyntaxError*'scale' param of 'activation' can only be a scalar or a vector in partition dimension, scale.shape=<shape>
err_src_dst_same_dtype0x50c70 · 2861NKISyntaxError*<api> src and dst must have same data type but got <src>/<dst>
err_tile_index_add_non_integer_scalar0xa73e0 · 1752TypeErrortile_index can only add with integer scalar, got scalar with dtype <dtype>
err_unexpected_type_of_operand0x8a800 · 1339TypeErroroperand type not among <expected_types>
err_operand_cannot_be_none0x941e0 · 323ValueError'<name>' cannot be None.
err_incorrect_target_type0xd5170 · 1481NKISyntaxError*target type <target_type> not in <expected_targets> for op <name>
err_unsupported_memory0xda290 · 1810TypeErrorExpected operand '<name>' of '<api>' to be in address space '<expected_addr_space>', but got a <tile> instead.
err_shared_hbm_must_in_kernel_level0xeb420 · 2080RuntimeErrorshared_hbm buffer can only be created top level kernel scope <scope>
err_hbm_tensor_with_init_value_not_supported0x44850 · 1499NKISyntaxError*Creating HBM tensor with init value is not supported.
err_tensor_creation_on_scratchpad_with_init_value_not_supported0x44a20 · 1485NKISyntaxError*Creating SBUF/PSUM tensor with init value is not supported in allocated kernels.
err_bias_tensor_must_be_specified_in_allocation0x44680 · 1521NKISyntaxError*Bias tensor for activation op must be specified in allocated kernel!
err_transpose_on_tensor_engine_not_allowed_in_allocated_kernel0x444b0 · 1590NKISyntaxError*feature gate (nc_transpose on TensorEngine / matmul w/o transpose_x in allocated kernel)
err_instruction_unsupported_op0x5d880 · 1399NKISyntaxError*<op_name> is not supported in '<inst_name>'
err_instruction_engine_unsupported_op_comb0xba320 · 1409NKISyntaxError*unsupported (op0, op1) pair on engine
err_instruction_supported_dtypes0x98ac0 · 1420NKISyntaxError*<inst>.<param> on <target> dtype <dtype>, supported: <supported_dtypes>
err_instruction_unsupported_dtypes0x975f0 · 1430NKISyntaxError*dtype <dtype> ∈ per-target blacklist for inst/param/target
err_tensor_int32_add_multiply_supported_engine0xb2490 · 2898NKISyntaxError*int32 add/multiply only on engine(s) <multiple_engines>
err_par_reduce_unsupported_op0xeace0 · 1454NKISyntaxError*par_reduce does not support operator=<op_name>
err_reduce_unsupported_negate0x76650 · 1458NKISyntaxError*negate option can only be used with arithmetic ops, unsupported op=<op_name>
err_reduce_bitvec_op_invalid_dtype0x85c20 · 1440NKISyntaxError*bitvec reduce op requires integer dtype, got <dtype>
err_par_reduce_bitvec_op_invalid_dtype0x84ee0 · 1447NKISyntaxError*partition bitvec reduce requires integer dtype, got <dtype>
err_bitvec_operand_must_be_integer0xbc170 · 1462NKISyntaxError*bitvec op operand <operand_name> dtype <dtype> must be integer
err_mask_not_equal_not_supported0x44bf0 · 1468NKISyntaxError*'not equal' mask is not supported
err_logic_or_not_supported0x94f50 · 339ValueError'logical or' is not supported for the \``
err_exact_arch_support_for_inst0xbd240 · 2910NKISyntaxError*inst <inst_name> available only on <available_targets>
err_min_arch_support_for_inst0xe3620 · 2919NKISyntaxError*inst <inst_name> requires min target <min_target>
err_min_arch_support_for_hwdge0x69130 · 2927NKISyntaxError*HW DGE mode <dge_mode> for inst <inst_name> requires min target <min_target>
err_multi_cores_spmd_not_supported0x7d0a0 · 2064RuntimeErrorSPMD grid with multi neuron cores is not supported on <target>
err_tensor_access_out_of_bound0x5ae50 · 1840IndexErrorOut-of-bound access for tensor \<tensor_name>` on dimension … index range <oob_range> exceed dimension size of `
err_tensor_access_out_of_bound_1range_string0xa1500 · 1885NKISyntaxError*helper formatting <range_tuple> for the OOB message
err_tensor_index_not_supported0x9cda0 · 2124NKISyntaxError*tensor <tensor_name> indexed with unsupported index <index> (<type_name>)
err_cannot_assign_to_index0x46510 · 2008TypeError'index' tensor does not support item assignment
err_unsupported_mixing_basic_advanced_tensor_indexing0x43d70 · 1934NKISyntaxError*Mixing basic tensor indexing and advanced tensor indexing is not supported.
err_indirect_indices_free_dim0x439d0 · 2133NKISyntaxError*Indirect indexing on free dimension not supported, must be partition dimension or block dimension.
err_indirect_indices_sbuf0x43800 · 2158NKISyntaxError*Indirect indices tile must be on SBUF.
err_indirect_indices_are_not_supported_with_nki_api0xaafb0 · 2459NKISyntaxError*cannot use indirect indexing access with the current NKI API for <indirect_operands>
err_copy_dynamic_indirect_indices_not_natively_supported0xebba0 · 2484NKISyntaxError*cannot copy a tensor with indirect memory reference access to \`(usenisa.tensor_copy_dynamic_src`)
err_atomic_rmw_add_only0x46b10 · 2163NKISyntaxError*atomic_rmw \op` param only supports 'add' operation currently. … op=`
err_atomic_rmw_dynamic_index0x43630 · 2170NKISyntaxError*atomic_rmw \dst` only supports indirect dynamic indexing currently.`
err_1d_arange_not_supported0xa2240 · 1716NKISyntaxError*<tensor_name> with 1d arange is not supported.
err_unexpected_output_dependencies0xec2e0 · 1895NKISyntaxError*Unexpected output dependencies <missing_indices> (parallel iterations write same location)
err_dynamic_control_flow_not_supported0x442e0 · 1620NKISyntaxError*dynamic control-flow depending on tensor value is not supported.
err_control_flow_condition_depending_on_arange0x43f40 · 1688NKISyntaxError*Control-flow depending on \nl.arange` or `nl.mgrid` is not supported.`
err_unsupported_expression_in_mask0x44110 · 1635ValueErrorNKI mask expressions must be affine expressions of static loop indices. …
err_while_loop_requires_unconditional_entry0x43290 · 2584ValueErrorTraditional while loops are not supported in NKI. Use the do-while pattern instead: …
err_ambiguous_tensor_truth_value0x43460 · 2533ValueErrorcannot evaluate truth value of a multi-element tensor (use ~/&/`
err_nki_api_outside_of_nki_kernel0x43ba0 · 2027RuntimeErrorcalling NKI API outside of NKI kernels is not supported.
err_nested_kernel_with_spmd_grid0xb4270 · 2038NKISyntaxError*<func_name>[<grid>]) inside another kernel is not supported.
err_program_id_axis_out_of_bounds0x83730 · 1323IndexErroraxis in \program_id` is out-of-bound(s) of the spmd launch grid (axis=, rank )`
err_local_variable_used_out_of_scope0xb4f00 · 1257NKISyntaxError*Local variable '<name>' is referenced outside of its parent scope <parent_scope>
err_tensor_output_not_written_to0xeca20 · 2187ValueError<tensor_name> … never written to (often a dead loop)
err_cannot_update_immutable_parameter0x45f10 · 2287TypeErrorCannot update immutable parameter \<tensor_name>``
err_mutable_parameter_not_returned0xaa470 · 2242NKISyntaxError*<tensor_names> … mutable kernel parameter not returned
err_failed_to_infer_tile_from_local_tensor0x64f90 · 1770TypeErrorFailed to infer tile from tensor '<tensor_name>': the first dimension … is not the partition dimension
err_tiled_offloaded_memcpy_same_shape0x4aec0 · 2175NKISyntaxError*src and dst must have the same shape, src.shape=<src> dst.shape=<dst>
err_tiled_offloaded_fma_same_length0x95cc0 · 2181NKISyntaxError*srcs and scales must have the same length, len(srcs)=<srcs>, len(scales)=<scales>
err_expected_constant_value0x7c4a0 · 1472NKISyntaxError*expected a compile-time constant <expected>, got <value>
err_nki_param_not_hashable0x82b20 · 2071ValueError'<name>' to be hashable, but got type <ty>
err_unsupported_memory (addr cross-listed above)———
warn_block_dimension_is_deprecated0xd5ea0 · 2512DeprecationWarningBlock dimension is deprecated. The leading dimension of <tensor>… (non-fatal)

5b. check_* validators → err_* dispatch (32)

check_* validatorAddrDispatches to
check_tensor0x61bb0err_unexpected_type_of_operand
check_tile0x61010assert_is_tile
check_dtype0xc0510err_operand_cannot_be_none, err_unsupported_dtype_value
check_tensor_dtype0xcf5c0err_supported_operand_dtype
check_tensor_addr_space0xd2c60err_unsupported_memory
check_param_type0xcaae0err_operand_cannot_be_none, err_unexpected_type_of_operand
check_parameter_shapes_np_broadcast0x789a0err_param_shape_incompatible_with_numpy
check_parameter_shapes_without_broadcast0xabeb0err_param_shape_mismatch
check_par_dim_sizes0x6f060err_num_partition_exceed_arch_limit
check_par_dim_size_match0x8c0e0err_num_partition_mismatch
check_free_dim_size_sbuf0x712e0assert_free_dim_sbuf
check_matmul_high_level_shape0xf2a30err_param_shape_incompatible_with_matmul
check_matmul_f_dim_sizes0xe7f90err_size_of_dimension_exceed_arch_limit
check_transpose_shape0x73e60err_exceed_max_supported_dimension, err_size_of_dimension_exceed_arch_limit
check_store_shape0xa5ac0err_store_dst_shape_smaller_than_other_shape
check_tensor_reduce_supported_ops0x5ec40err_instruction_unsupported_op

The remaining validators (check_matmul_mx_low_level_shape, check_param_against_dtype_hints, check_modified_type, check_parameter_types, check_and_update_list_param, check_and_update_param, check_direct_type_match, check_tensor_tile_promotion, check_modified_tensor_type, check_parameter_shapes, check_parameter_shapes_no_broadcast, check_parameter_shapes_low_level, check_matmul_low_level_shape, check_local_gather_shape, check_stream_shuffle_shape, check_quantize_mx_shape) perform leaf checks or dispatch via a Cython-cached slot; several (check_quantize_mx_shape, check_matmul_mx_low_level_shape, check_local_gather_shape) raise inline messages with no fronting err_* symbol (D-W05 §5d).

5c. assert_* → err_* (the named-err subset)

The 24 assert_* funnel through one primitive nki_assert(exception, condition, msg) (0xbb9a0, sema.py:1233). Most build an inline message; six dispatch to a named err_*:

assert_*Raises
assert_num_partition / assert_par_dim_sbuf / assert_par_dim_psumerr_num_partition_exceed_arch_limit
assert_tile_has_expected_shapeerr_tile_shape_mismatch
assert_is_tileerr_failed_to_infer_tile_from_local_tensor | err_unexpected_type_of_operand
assert_constant_valueerr_expected_constant_value
assert_tensor_in_valid_addr_spaceserr_unsupported_memory

6. Adversarial Self-Verification

The catalog was extracted programmatically from D-AG09 §8 (itself a verbatim runtime dump of the live dict), not hand-typed. The verification done before publishing:

  • Total count. Parsed D-AG09 §8 and counted 2556 CAUSE: lines and 353 ---- (N codes) ---- headers; the per-header counts sum to 2556 with zero mismatch — and the markdown body below renders 2556 | \Code` |rows. This matches the runtimelen(NEURON_ASSERT_ERROR_MESSAGES) == 2556` (D-AG09 §0) and the brief's "≈2556". CONFIRMED — the binary says 2556, the page reproduces 2556.
  • Histogram. First-letter category histogram independently re-derived: I=252 E=24 X=15 L=12 B=9 S=8 M=5 N=4 D=4 C=4 T=3 P=3 A=3 O=2 J=2 V=1 R=1 G=1 — exactly D-AG09 §1. (CONFIRMED.)
  • Resolution split. 2095 rows carry only {RESOLUTION_CONTACT_SUPPORT} (rendered *(contact)*); 461 carry a bespoke resolution. Matches D-AG09 §1. (CONFIRMED.)
  • Spot-check of 5 representative rows against the string table / dump:
KeyCause (verbatim)ResolutionTag
(ADA, 1)ADA computation should be called only once per batch. The history should be empty, but not.*(contact)*CONFIRMED
(BIR, 18)ctx is nulltpr (typo preserved from binary)*(contact)*CONFIRMED
(EARG, 1)Illegal argument(s) - Logical Neuron Core size {lnc} is not supported on trn1. …Update the Logical Neuron Core size to 1 or check that the target architecture is correct.CONFIRMED
(NKI, 16)Kernel validation exception: {error_text}Please check the validation message and adjust kernel inputs accordinglyCONFIRMED
(IIOT, 901)InsertOfflaodedTransposes assertion error: {baseMessage} (generated 901 family; spelling Offlaoded preserved)*(contact)*CONFIRMED

All five match the dump byte-for-byte, including preserved typos (nulltpr, Offlaoded, diretory, parellelism) — which is itself evidence the table was dumped, not paraphrased. The (NKI, 16) body confirms the pin to nki/.../kernel_assert.py:26 (D-A13 §4). (CONFIRMED.)

NOTE — what stays opaque (honest bounds). This page gives every ErrorMessages code's text but not its native C++ call site: which libwalrus/libBIR/codegen instruction-check invokes neuron_assert("…", CAT, idx, …) for a given (CAT, idx) lives in the statically-linked C++, not in ErrorMessages.so (D-AG09 §9). The only py call sites pinned to source are DAE001/DAE002 (cli/Daemon.py) and NKI016 (kernel_assert.py). The 224 frontend NCC_* codes' bodies (§3) are a separate catalog in the HLO binaries' rodata — out of scope for the dict dump and not reproduced verbatim here. {kwarg} values are runtime-supplied; only the static templates are recovered.

Cross-References