blackwell-isa

Public research artifacts for the NVIDIA Blackwell SASS instruction set.

This repository publishes machine-readable ISA databases for Blackwell GPUs: consumer SM120 (RTX 50-series, including RTX 5090 / GB202), datacenter SM103a (B300 SXM6), and SM121a (DGX Spark / GB10, Grace-Blackwell unified). It is intended for researchers and implementers building assemblers, disassemblers, compiler backends, binary analysis tools, and instruction schedulers for Blackwell-class targets.

Artifacts:

• sm120.json — canonical machine-readable ISA database
• sm103a.json — canonical SM103a ("B300", Blackwell datacenter) ISA database; same schema, validated bit-exact against an SM103a corpus (see "SM103a notes" below); refreshed from the production cubit table including the tcgen05 class set (UTCHMMA/UTCQMMA/UTCIMMA/UTCOMMA, LDTM/STTM, UTCCP, UTCSHIFT, UTCBAR, UTCATOMSWS, UVIRTCOUNT)
• sm100a.json — SM100a ("B200", Blackwell datacenter) ISA database derived from the SM103a canonical revision (encoding layer byte-identical on all 119 shared tcgen05 probe pairs) plus the tcgen05 acceptance deltas; offline-validated by decoder/encoder replay on an archived 20.0M-instruction sm_100 corpus (99.9219% EXACT-minus-sched on 1,134 cubins); scheduling/pipeline sections inherited from SM103a, unvalidated for B200 timing
• sm121a.json — canonical SM121a ("GB10", DGX Spark Grace-Blackwell) ISA database; same schema, delta-driven vs SM120 with EXACT gates per entry; 15,443 encoding keys (md5 88142da9). Counting note: the SM121a table keeps a finer per-key granularity — its instructions dict holds one key per encoding row, so raw key counts are not comparable with the canonical instruction-form counts of sm120.json/sm103a.json
• SM120_ISA_REFERENCE.html — generated, searchable HTML reference (browse online)
• SM103A_ISA_REFERENCE.html — generated, searchable HTML reference for SM103a (browse online)
• SM100A_ISA_REFERENCE.html — generated, searchable HTML reference for SM100a (browse online)
• SM121A_ISA_REFERENCE.html — generated, searchable HTML reference for SM121a (browse online)

Artifact Scope

This is a data release, not a compiler or assembler. It contains:

• 128-bit SASS encoding templates and variable bit-field mappings
• operand extraction rules for registers, predicates, immediates, guards, and modifiers
• instruction form metadata (base_op, operand signature, modifier groups)
• scheduling metadata: pipeline class, latency, throughput, stall defaults, control-word class
• pipeline/resource configuration used by an SM120 scheduler
• a generated HTML reference for manual inspection

It does not contain the reverse-engineering pipeline, probe programs, proprietary cubin corpora, driver dumps, firmware notes, or private working material. Those are intentionally excluded from this public artifact.

SM103a notes

SM103a (sm103a.json) targets the B300-class Blackwell parts (ELF e_flags SM field 0x67, distinct from SM120 0x78 and SM100 0x64). Empirical hardware finding: CS2R Rd, SRZ clears the even-aligned 64-bit register pair Rd:Rd+1, not only Rd (cuda-gdb-proven; binary patchers must not substitute 32-bit zero-idioms blindly — consumers of Rd^1 observe the side effect).

Measured tcgen05 facts encoded in this release:

• For all 119 probe pairs shared with SM100a in the tcgen05 corpus, the 128-bit instruction words are byte-identical between sm_100a and sm_103a; the encoding layer transfers 1:1 to SM100a.
• ptxas 13.3 accepts tcgen05.mma kind::i8 on sm_100a but rejects it on sm_103a.
• tcgen05.ld.red (LDTM.STAT) is accepted on sm_103a only.
• tcgen05.wait::alloc/dealloc and tcgen05.commit complete_tx::bytes are rejected by ptxas 13.3 on both targets; TMEM allocation runs through UTCATOMSWS/UTCBAR arbitration and UVIRTCOUNT.DEALLOC.SMPOOL pool release as encoded here.

Coverage note: SM103a encoding records stand on a 1.94M-sample nvdisasm-pair corpus plus the tcgen05 PTX probe corpus. Per-instruction scheduling metadata is thinner than in sm120.json (baseline defaults except where hardware-probed); consult ctrl_classes for control-word behavior, which is documented per class.

Validation Status

• 47,244 real instructions decoded across 178 cubins with 100% decode coverage
• 5,014 / 5,014 roundtrip fuzz cases passing through the companion assembler
• 936 / 936 targeted ground-truth instructions round-tripped, including the 25-type QMMA.SF ptxas corpus and LDG.E.LTC128B.128
• all 36 dense QMMA.SF...E8 type pairs covered: 25 emitted by ptxas and all 11 combinations containing undocumented E3M4 executed on RTX PRO 6000
• selected semantic bit fields validated by hardware patch tests on RTX 5090 and RTX PRO 6000
• SM120 (O2 consolidation revision): roundtrip over the 240-file sm120 corpus (1,076,075 instructions) at 99.86% match (up from 94.0% under the previous canonical revision); the 108 dropped _? placeholder rows were census-proven never hit by the decoder on 1.64M natural instructions
• SM103a: disassemble→assemble roundtrip byte-exact on 323/323 cubins of the publish corpus (tcgen05 probe kernels, FA4 kernels, instruction samples) run under the exact published table; per-cubin verdicts identical to the production cubit table
• SM100a: decoder/encoder replay over the archived b4 sm_100 corpus (1,134 cubins, 20,024,512 instructions, cuBLAS/cuBLASLt/cuSOLVER/cuSPARSE/curand + probes): 99.9219% EXACT-minus-sched, 0 fatal files; residual gap classes enumerated in the O2 validation report
• SM121a: differential legality/bit-flip scan of 260,352 probes (1,017 canonical forms × 128 bits × 2 contexts) with 0 decode/legality differences vs SM120; nvdisasm -b SM120 vs -b SM121a over 648,992 raw instructions with 0 decode differences; full-table census over the 92,098-word corpus with 0 failures; 27/27 canonical identity anchors round-trip perfectly

Data Model

At a high level:

sm120.json
├── _meta
│   ├── architecture, codename, gpu, instruction_width
│   ├── stats
│   ├── pipe_classes
│   ├── ctrl_classes / ctrl_epochs
│   └── hardware_config
│
├── instructions
│   └── INSKEY
│       ├── base_op          # opcode family, e.g. IADD3, QMMA, LDG
│       ├── operand_sig      # operand signature, e.g. R_P_P_R_R_R
│       ├── mod_groups       # encoding variants
│       ├── scheduling       # pipeline / latency / throughput metadata
│       ├── ctrl_class       # scheduler control-word class
│       ├── mercury          # recovered NVIDIA pattern metadata, when known
│       └── _discovery       # provenance for recently added or unusual entries
│
├── sched_only               # scheduling entries without full encoding records
└── pipeline_config          # resource defaults, latency tables, per-opcode records

A typical encoding entry:

{
  "and_base": "0x0000000000000000000000000000082e",
  "variable_mask": "0x1c00000000000000000000000000f000",
  "fields": [
    { "shift": 12, "bits": 4, "token_idx": 0, "extraction": "guard" },
    { "shift": 16, "bits": 8, "token_idx": 1, "extraction": "reg" }
  ]
}

and_base is the constant part of the 128-bit instruction word. variable_mask identifies operand-dependent bits. fields tells an assembler/disassembler where to read or write each operand field.

Minimal Usage Example

import json

db = json.load(open("sm120.json"))
insn = db["instructions"]["IADD3_R_P_P_R_R_R"]
variant = insn["mod_groups"][""]

print(insn["base_op"])                 # IADD3
print(variant["and_base"])             # 128-bit instruction template
print(variant["fields"][0])            # one encoded field description
print(insn["scheduling"]["pipe_name"]) # INT_ARITH

For an end-to-end assembler/disassembler using this data, see cubit.

Selected Research Results

These are included to orient readers; the primary contribution is the data itself.

• SM120 is not SM100. SM120 uses register-accumulator QMMA/HMMA/IMMA/DMMA forms, while SM100 exposes the TMEM-based tcgen05.mma model. Treating consumer Blackwell as a small B200 gives wrong code-generation assumptions.

• Block-scaled MMA is encoded in SASS. QMMA.SF, QMMA.SF.SP, and OMMA.SF forms are present. Current compute_120f PTX exposes the documented mxf8f6f4 type pairs; the database additionally records undocumented E3M4 and hidden/sparse forms outside the normal compiler path.

• There is an undocumented FP8-like type code. Type code 2 in the block-scaled MMA type field corresponds to E3M4; it appears in dense and sparse instruction forms and is accepted by hardware.

• Scheduling metadata matters. Several instruction families are not assigned to the pipeline one would infer from the mnemonic alone. Schedulers should consume the scheduling and ctrl_class fields rather than deriving hazards from opcode names.

• Control-word handling is instruction-class dependent. The upper control bits are not a uniform free field. ctrl_class and ctrl_epochs describe which bits are fixed discriminators and which bits a scheduler can safely modify.

• The public memory hierarchy numbers are easy to misread. The database records RTX 5090 L2 as 96 MB; 128 MB is the access-policy window, not the cache size.

Limitations

• sm120.json targets SM120 / consumer Blackwell; sm103a.json targets SM103a (B300); sm121a.json targets SM121A (GB10, DGX Spark). Together they are not a complete SM100/B200 ISA, though the measured tcgen05 encoding layer is byte-identical between SM100a and SM103a.
• The tables use different key-granularity conventions: canonical forms with nested variants (sm103a.json) vs enumerated variant keys (sm120.json, sm121a.json). Raw key counts are not comparable across tables; the mnemonic-family coverage is essentially the same (~150 families each).
• SM121A entries are validated on GB10 (Grace-Blackwell, unified memory) only.
• Some keys are marked partial (?): the encoding is valid, but the full operand syntax is not yet determined.
• Some entries are scheduling-only: they describe pipeline behavior without a complete encoding record.
• Some instruction names use provisional labels (INVALID*, recovered modifier names, or internal pattern names) where NVIDIA has no public terminology.
• Metadata under _discovery records provenance and confidence for unusual entries; users should inspect it before relying on those forms in production code.

Related Work

• cubit — companion SM120 / SM103a / SM121A SASS assembler/disassembler using this database (per-arch tables: sm120.json, sm103a.json, sm121a.json).

Citation

If this artifact is useful in a paper or project, please cite the repository and the commit hash of sm120.json used in your work.

License

MIT — see LICENSE.

Reverse engineering for interoperability is protected under EU Directive 2009/24/EC Article 6 and US DMCA §1201(f).