GRIT — Grouped Reduced-precision Interchange Type

Normative Specification, version 1

Status: normative. Descriptor version: 1. Endianness: little, everywhere. Companion artifacts: spec/reference.py (executable reference implementation), spec/conformance/vectors.json (conformance vector suite, generated by the reference).

This document is self-contained. An implementer needs nothing else — no design document, no vendor manual, no external standard — to produce a conformant implementation. Where GRIT reuses a concept from OCP MX, IEEE P3109, PTX or GGUF, the concept is restated here in full.

Key words MUST, MUST NOT, SHOULD, MAY are to be interpreted as in RFC 2119.

Table of contents

• 0. What GRIT is
• 1. Numeric preliminaries
• 2. grit_fmt_t — the 16-bit format word
• 3. grit_desc — the 64-byte descriptor
• 4. The Grade
• 5. The placement word
• 6. Fingerprinting: grade_id, gid, digest
• 7. grit_check and the boundary result
• 8. Normative decode()
• 9. Normative encode() / quantize()
• 10. The four reference grades
• 11. Supersedes edges and the frame/note split
• 12. assume_placement and the audit ledger
• 13. grit scan — the diagnostic wedge
• 14. Degenerate byte patterns
• 15. Conformance
• Appendix A. Constant tables
• Appendix B. Worked examples
• References (non-normative)

0. What GRIT is

A GRIT array is a quadruple (Grade, Placement, Planes, Shape) denoting a logical array of binary32 values (extended reals ∪ {NaN}).

• Grade — the complete numeric contract: which element format, how many elements share a scale, what the scale formats are, whether there is a zero point and what convention it uses, whether the array is N:M sparse, how bits are packed, and what rounding/saturation/accumulation the producer and the consumer promise. A Grade has a canonical string form and a 64-bit id derived from it. §4.
• Placement — the physical layout, carried as a 64-bit value, never as a type parameter, so that a new layout is a new constant rather than a new type. That much is long-standing practice in vendor runtimes and compilers (§0.2, C3); what GRIT adds is that the value is hashed, registry-free and attached to bytes at rest next to the numeric contract. §5.
• Planes — up to four separate byte buffers: data, scale0, scale1, aux (aux = sparsity metadata or zero points). §3.6.
• Shape — rank and extents, carried in the view, not in the descriptor.

decode() (§8) is the normative semantics. Two GRIT arrays are semantically equal iff their decodes are equal. Everything else in this document exists to make that statement checkable at an O(1) boundary crossing.

0.1 The three claims GRIT makes and the one it refuses to make

1. Equal gid and equal plane bytes ⇒ bit-identical decode() on every conformant implementation. This is testable, and §15 is the test.
2. grit_check failure is always a diagnostic, never undefined behaviour, for any 64-byte input including adversarial ones (§14).
3. gid certifies decode agreement and nothing else. It does not certify calibration quality, provenance, authorship, or trustworthiness. GRIT makes no social claims: a forged descriptor gains an attacker nothing, because a descriptor that disagrees with its bytes fails grit_check.

0.2 Claims and prior art

(Non-normative. Nothing in this subsection changes encoding, decoding, checking, or hashing behaviour.)

GRIT is assembled almost entirely from ingredients that already exist. This subsection states, component by component, exactly what GRIT claims and what it does not claim, and names the closest prior art for everything in the second category. It is written deliberately narrower than the marketing version of each idea, following an adversarial prior-art review (2026-08-01; roughly forty candidates run to primary sources across standards, vendor runtimes, compiler IRs, open-source frameworks, and arXiv). Full citations are in References.

The components are labelled C1–C6 for cross-reference.

The honest residual risk. The exposure here is not anticipation but obviousness-by-union: a reader can assemble the Khronos Data Format Specification (parametric binary format descriptor) + Avro (canonical form and fingerprint) + XLA/CuTe (layout as value with an algebra) + graded modal types (a contract lattice) and say "apply that to quantized tensors"; and C1's content union can be read as an aggregation of long-established block-scaling and sparse-format practice. No amount of searching settles that argument, and it is stated here rather than hidden. What the review could offer against it is that well-resourced parties recognised the gaps without producing the synthesis (ONNX issue #7691, vLLM RFC #8913, the MLIR quant-dialect RFCs, the TOSA block_scaled RFC), and that IEEE P3109 §VIII explicitly disclaims memory and wire representation — stopping exactly where C3 and C6 begin.

What GRIT claims about none of the above. GRIT claims no priority over any cited work, asserts no defect in any of them, and does not claim to be the first parameterised description of reduced-precision numeric formats. Where this document reuses vocabulary from OCP MX, IEEE P3109, PTX or GGUF, it reuses it openly (§0, opening note).

0.3 Scope and non-goals

(Non-normative. This subsection describes the boundary of version 1; it imposes no requirement beyond those already stated in §1–§15.)

Version 1 covers the block-scaled affine family: one or two levels of shared scale over contiguous groups, an optional zero point, optional N:M structured sparsity, over parameterised float/exponent-only/integer element formats. It does not cover everything deployed today. Any impression that GRIT v1 describes every shipped quantized format should be read against the following four exclusions, each of which is structural in version 1 rather than an oversight or a to-do.

(a) FNUZ FP8 is inexpressible in version 1. bias_ovr MUST be 0 (§2.1), so every format's exponent bias is exactly 2^(e−1) − 1; and e4m3fnuz MUST NOT be spelled e4m3 (§2.6). e4m3fnuz and e5m2fnuz use bias 2^(e−1) and encode their single NaN as the "negative zero" code, which no nan_enc value in §2.3 describes. There is therefore no grit_fmt_t word denoting the AMD MI300 / CDNA3 FNUZ FP8 formats, and a conformant v1 implementation cannot describe such a checkpoint. This is a real coverage hole, not a claim that the formats are unimportant. Version 1.1 intends to lift bias_ovr (adding an explicit bias field) together with the additional nan_enc value the FNUZ NaN code requires; both changes are additive and are not to be anticipated by a v1 implementation.

(b) There is no codebook / lookup-table element class. The precision code p (§2.2) admits exactly three kinds of element: exponent-only, floating point, and integer. A codebook format decodes a stored code through a table that travels with the tensor, which is a different decode equation with a different plane inventory and different laws (table identity, table sharing, table mutation). Consequently NF4 and the llama.cpp IQ* family — and vector-/codebook-quantized formats generally — are out of scope for version 1. Adding them means a new element class plus a codebook plane plus their own law set, not a new flag.

(c) GPTQ act-order (g_idx) grouping is unsupported. §3.7 defines grouping as contiguous runs, g0 = lin(i) / k0. A GPTQ checkpoint quantized with activation reordering ships a per-column g_idx vector assigning each column to an arbitrary group, which the §3.7 equation cannot express.

It cannot be moved into Placement either, and the reason is a law rather than an implementation limit. L1 (§5.4) states decode ∘ relayout(p₁→p₂) = decode: placement is decode-invariant, and that theorem is exactly what licenses erasing the placement witness at GRIT_CHECK=0 (L9). A g_idx permutation changes which shared scale multiplies which element, so it changes decode(). Encoding it as a placement would therefore make L1 false and take L9's erasure soundness down with it — a silently wrong dequantisation at the exact check level where the witness is erased. Contiguous grouping is structural in version 1.

An act-order checkpoint can still be handled outside GRIT by materialising the permutation — permuting columns so groups become contiguous, and applying the matching permutation to the activations — after which the result is an ordinary contiguous-group GRIT array. Version 1 provides no way to record that permutation, so it is not self-describing. Expressing act-order natively would require a new grade field, its own plane, and a revised law set.

(d) Dense-nibble zero-point planes are inexpressible, so GPTQ and AWQ checkpoints are not byte-representable in version 1. §3.4 stores sub-byte scale and zero-point plane values as CT_PAD8 — one value per byte, low bits, zero-extended — regardless of container, which removes an otherwise unresolvable ambiguity in the aux plane. GPTQ and AWQ both pack qzeros as dense nibbles, two per byte, so for a 4-bit group-128 projection GRIT's plane size equation predicts exactly twice the aux bytes those files actually hold (measured: 6,272 predicted against 3,136 present, on unmodified plane bytes from two real Hub checkpoints — see audit/scan-experiment.md, Finding C). The data and scale planes match the size equations exactly; only the zero-point plane diverges. Placement does not rescue this: placement_meta exists, but version 1 defines no dense-nibble metadata placement in its registry. This exclusion is the least exotic of the four — it is the two most widely deployed 4-bit weight formats on the Hub — and closing it in a later version means permitting a dense sub-byte container for zero-point planes, together with whatever disambiguation §3.4's rule was introduced to supply.

Beyond these four, and restating §0.1(3): GRIT is not a security mechanism, not a calibration-quality signal, and not a provenance system.

1. Numeric preliminaries

1.1 Working precision

All GRIT arithmetic specified in this document is performed in IEEE 754 binary32 with round-nearest-ties-to-even (RNE) unless a Grade field says otherwise. Every intermediate named in §8 and §9 is a binary32 value; an implementation MUST round to binary32 after each named step. Fused multiply-add MUST NOT be substituted for a specified multiply-then-add: FMA changes the result and is therefore a different Grade.

An implementation MAY evaluate in wider precision internally only if the observable result is bit-identical to the step-by-step binary32 evaluation. (Evaluating a*b in binary64 and rounding once to binary32 is bit-identical to binary32 multiplication when a and b are binary32; evaluating a*b*c that way is not, and is therefore forbidden.)

1.2 NaN

Every NaN produced by decode() is the canonical quiet NaN with binary32 bit pattern 0x7FC00000. Implementations MUST NOT propagate payloads or signalling bits. This makes conformance vectors comparable bit-for-bit.

1.3 Signed zero

Signed zero is preserved end to end. decode() of an element whose stored code is a negative zero yields -0.0 (bits 0x80000000) when the scale chain is a positive finite value. encode() of -0.0 yields the negative-zero code of the element format when that format has one.

1.4 Rounding modes

For a projection into a format with a finite, monotone magnitude ladder (§2.5), "ties to even significand" is exactly "ties to even stored code", because consecutive codes in the ladder differ by one in the significand.

1.5 Saturation modes

For signed integer element formats, satfinite clamps to [-(2^(w−1)−1), +(2^(w−1)−1)] — i.e. the asymmetric minimum -2^(w−1) is never produced by encode(). It remains decodable if present in a buffer.

2. grit_fmt_t — the 16-bit format word

A grit_fmt_t is a uint16_t, little-endian, describing one stored numeric format. It is parameterised, not enumerated: E2M1, E4M3, E5M2, E8M0, binary16, binary32, u2/u4/u6/u8, s4/s8 and the format nobody has invented yet are all instances. Parameterising a numeric format rather than enumerating it is not a GRIT idea — the Khronos Data Format Specification and IEEE P3109 both do it, and GRIT reuses their vocabulary openly (§0.2, C6a; References). What is GRIT's is that this word sits inside the hashed descriptor, so a format distinction is a gid distinction (§2.6).

2.1 Bit layout

 bit  15   14         13      12 11      10       9 .. 5      4 .. 0
     +----+----------+-------+--------+--------+-----------+-----------+
     |bias| unsigned | szero | nan_enc| has_inf|     p     |    wm1    |
     |_ovr|          |       |  (2b)  |        |   (5b)    |   (5b)    |
     +----+----------+-------+--------+--------+-----------+-----------+

The all-zero word 0x0000 is reserved and means ABSENT. No 1-bit format may be encoded in version 1.

2.2 The precision code p

Implementer note. p = 31 as an integer sentinel is a deliberate spec decision. Without it, s8 and a hypothetical E1M6 would share a bit pattern. Do not "optimise" it away.

2.3 nan_enc

2.4 Decoding a stored code

fmt_decode(word, code) -> binary32
  f = unpack(word); mask code to w bits
  if p == 31:                       # integer
      if unsigned: return (float)code
      if code >= 2^(w-1): code -= 2^w
      return (float)code
  if p == 0:                        # exponent-only
      if nan_enc != 0 and code == 2^w - 1: return NaN
      return 2^(code - (2^(w-1) - 1))
  # floating point
  m_bits = p - 1;  e_bits = w - (unsigned?0:1) - m_bits;  bias = 2^(e_bits-1) - 1
  m = code & (2^m_bits - 1);  e = (code >> m_bits) & (2^e_bits - 1)
  s = unsigned ? 0 : (code >> (m_bits + e_bits)) & 1
  if e == 2^e_bits - 1:
      if has_inf:            v = (m == 0) ? Inf : NaN
      elif nan_enc == 1 and m == 2^m_bits - 1: v = NaN
      elif nan_enc == 2 and m != 0:            v = NaN
      else:                  v = (1 + m/2^m_bits) * 2^(e - bias)   # E2M1 lands here
  elif e == 0:               v = (m / 2^m_bits) * 2^(1 - bias)      # subnormal
  else:                      v = (1 + m/2^m_bits) * 2^(e - bias)
  return s ? -v : v

Every registered format's every code is exactly representable in binary32; the conversion is therefore exact and rounding-free.

2.5 The magnitude ladder and fmt_encode_rne

For any format of width ≤ 8, the ladder is the list of (magnitude, code) pairs for all sign-cleared codes that decode to a finite value, sorted by increasing magnitude. fmt_encode_rne(word, x, sat):

1. x is NaN ⇒ the format's canonical NaN code (error if nan_enc = 0).
2. Let a = |x|, neg = signbit(x).
3. If a ≥ ladder.last.magnitude ⇒ ladder.last.code (this is satfinite; ±∞ lands here too).
4. Otherwise find i with ladder[i].mag ≤ a < ladder[i+1].mag. Choose the nearer; on an exact tie choose the even code.
5. OR in the sign bit if neg and the format is signed.

For integer formats, fmt_encode_rne rounds x to the nearest integer, ties to even, then clamps per §1.5. For binary16 and binary32 it is the ordinary IEEE conversion.

2.6 Registered formats

e4m3fnuz (bias 8, no signed zero) is a different format and MUST NOT be spelled e4m3. §4.4 makes this a parse error rather than a silent alias.

3. grit_desc — the 64-byte descriptor

sizeof(grit_desc_t) == 64, alignof == 8, POD, little-endian, no padding holes, no bitfields. This is the ABI. Every language binds to this.

3.1 Byte-by-byte layout

Bytes [0,28) ∪ [32,64) — 60 bytes — are the frame. Bytes [28,32) are the note. The split is normative and load-bearing; see §11.

/* grit.h -- normative C binding */
typedef uint16_t grit_fmt_t;
typedef struct grit_desc {          /* sizeof == 64, alignof == 8 */
  uint32_t   magic;                 /* 0x31545247 == 'G','R','T','1' LE     */
  uint8_t    hdr_bytes;             /* 64                                    */
  uint8_t    version;               /* 1                                     */
  uint8_t    n_levels;              /* 1 or 2                                */
  uint8_t    flags;                 /* GRIT_F_*                              */
  grit_fmt_t elem, scale0, scale1, zp_fmt;
  int16_t    group_axis;
  uint16_t   k0;                    /* elements per level-0 group            */
  uint16_t   k1;                    /* level-0 groups per level-1 group;
                                       0 == per-tensor (n_levels == 2 only)  */
  uint8_t    sp_n, sp_m;            /* N of every M survive                  */
  uint8_t    sp_meta_fmt;           /* GRIT_SPMETA_*                         */
  uint8_t    sp_order;              /* GRIT_SPORDER_*                        */
  uint8_t    container;             /* GRIT_CT_DENSE | GRIT_CT_PAD8          */
  uint8_t    frame_rsvd;            /* 0                                     */
  uint32_t   note;                  /* informational, NOT hashed (§11)       */
  uint64_t   placement_data;
  uint64_t   placement_scale;       /* covers scale0 AND scale1              */
  uint64_t   placement_meta;        /* covers the aux plane (meta or zp)     */
  uint64_t   grade_id;              /* FNV-1a-64 of the canonical grade      */
} grit_desc_t;
typedef struct grit_view {
  grit_desc_t d;
  int32_t     rank;
  int64_t     shape[GRIT_MAXRANK];
  void*       plane[4];             /* data, scale0, scale1, aux             */
  int64_t     plane_bytes[4];
  uint8_t     gid[16];              /* 0 => compute on demand                */
} grit_view_t;

3.2 magic, hdr_bytes, version

magic MUST be 0x31545247. As bytes on the wire that is 47 52 54 31, i.e. the ASCII GRT1, so a hex dump is human-recognisable. hdr_bytes MUST be 64. version MUST be 1.

A future version MAY grow the descriptor; a reader MUST use hdr_bytes to skip a longer descriptor it does not understand, and MUST NOT read past it.

3.3 flags

Version-1 constraints, all enforced by grit_check:

• GRIT_F_ZP_L1SCALED requires GRIT_F_HAS_ZP and n_levels == 2.
• GRIT_F_SPARSE and GRIT_F_HAS_ZP MUST NOT both be set (there is only one aux plane). Reserved for a future version with five planes.
• n_levels == 1 requires k1 == 0 and scale1 == 0x0000.

3.4 container

Both exist because PTX requires the padded form for mma.kind::mxf8f6f4 ("resides in the lower 6 bits of the 8-bit container") and the dense form for mma.kind::mxf4 ("no explicit padding is necessary"). One field replaces a per-backend reimplementation of that rule.

container describes the element plane only. Sub-byte scale and zero-point plane values are ALWAYS stored CT_PAD8 (one per byte, low bits, zero-extended), regardless of container. This is normative and removes an otherwise unresolvable ambiguity.

3.5 Bit packing

Within a container byte, element 2j occupies the low nibble and element 2j+1 the high nibble (LSB-first), unless GRIT_F_MSB_FIRST. Packing runs in increasing physical index order — that is, after the placement permutation. AWQ's [0,2,4,6,1,3,5,7] interleave is a Perm placement, not a packing variant.

Concurrency rule. Two threads writing elements that share a container byte is a data race, exactly as for C bitfields. Atomic element stores are provided only for widths ≥ 8. Stated, not discovered.

No &element. GRIT never claims a reference to a sub-byte element exists. Element access is through a proxy reference. This is why GRIT ships as a library today instead of dying in the RFC that asked how &u4 would work.

3.6 Planes

A plane that is not present MUST have plane_bytes == 0 and a NULL pointer.

placement_scale covers both scale0 and scale1; in version 1 the two scale planes MUST share a placement. placement_meta covers the aux plane whichever role it plays.

Serialised payload order. When the four planes are concatenated into a single byte string (as in the conformance vectors, safetensors, or a GGUF tensor blob) the order is data ‖ scale0 ‖ scale1 ‖ aux. Every plane length is derivable from the descriptor and the shape (§7.3), so the concatenation is unambiguously splittable.

3.7 Grouping

Groups are formed over runs of k0 consecutive elements along group_axis. k1 counts level-0 groups per level-1 group, not elements: k1 = 0 with n_levels == 2 means one level-1 group for the whole tensor (per-tensor scale). For element i with linear index lin(i):

g0 = lin(i) / k0
g1 = (k1 == 0) ? 0 : g0 / k1

shape[group_axis] MUST be a multiple of k0, and when k1 != 0, n_groups0 MUST be a multiple of k1.

3.8 Sparsity fields

sp_n : sp_m means sp_n of every sp_m consecutive logical elements along group_axis survive. 2:4 means two survivors per four. nnz = numel · N / M.

Spec decision, do not misread. Some prose in the wild writes the survivor count as (M−N)/M. GRIT rejects that reading. N is the number of survivors, matching PTX mma.sp and every 2:4 kernel in existence.

SPORDER_ASC_LSB is what mma.sp::ordered_metadata requires; PTX says the alternative is undefined behaviour. GRIT makes it a declared, checkable precondition (§7.4) instead of an unwritten one.

Metadata packing. For quad (group of M) index q and survivor slot j ∈ [0, N), the index occupies bits [(q·N + j)·bits, (q·N + j)·bits + bits) of the metadata plane, LSB-first within each byte. For 2:4 with SPMETA_U2: quad q occupies one nibble; survivor 0 in bits 0..1 of that nibble, survivor 1 in bits 2..3.

Data plane. Only surviving elements are stored, densely, in ascending logical order.

4. The Grade

A Grade is the complete numeric contract. It is what grade_id hashes and what a kernel signature names. It has exactly one canonical string form.

4.1 Grammar

GRADE  ::= "grit1(" [ FIELD { "," FIELD } ] ")"
FIELD  ::= KEY "=" VALUE
KEY    ::= one of the 23 keys in §4.2

Canonical form: lowercase, no whitespace anywhere, fields in the fixed order of §4.2 (not alphabetical), fields equal to their default omitted, integers in decimal with no leading zeros and no +.

4.2 Fields, in canonical order

Keys 1–11 describe the representation. Keys 12–15 are the encode projection π_enc. Keys 16–18 are the decode projection π_dec. Keys 19–23 are the accumulation contract α.

esr (the shared-scale derivation rule) has no default and is always emitted. That is deliberate: a defaulted scale rule is exactly the three-way disagreement (RNE in one framework, RTZ in the format spec, RTZ-or-RCEIL in a third) that produces unattributable eval regressions. The fork is documented, not hypothetical — see pytorch/pytorch#146414 (References), where a maintainer records that there is no consensus and that the OCP floor rule differs from a vendor's modified-ceiling rule. GRIT refuses to let it be implied.

Scale rules (esr):

Accumulation order (ao). unspec is a legal, first-class value — it is what every current tensor core honestly provides. PTX states, normatively, for every mma kind: "The accumulation order, rounding, and handling of subnormal inputs are unspecified." GRIT's contribution here is not to make hardware deterministic; it is to make nondeterminism a fact a program can see and refuse.

Honest caveat. Because unspec is what essentially every deployed tensor-core kernel reports today, the ao axis carries little discriminating information on the fast path at present. It earns its place by being writable at all — a kernel that does promise seq can say so, and a request for seq can be refused rather than silently served — not by partitioning today's kernel population.

4.3 Refinement order (⊑)

A Grade denotes a set of permitted numeric behaviours. g₁ ⊑ g₂ ⟺ behaviors(g₁) ⊆ behaviors(g₂). On the ao axis: seq ⊑ pairwise(k) ⊑ tree ⊑ unspec. On the subnormal axis, preserve and ftz are incomparable singletons whose join is any.

The order theory here is not new and is not claimed as GRIT's: a pre-ordered set of grades with subsumption and join-on-composition is standard graded modal typing, published under the same word Grade (Granule, ICFP 2019) and already applied to floating-point rounding (Numerical Fuzz, PLDI 2024). Contracts as structured values in an operation signature are likewise deployed art (StableHLO / jax.lax.DotAlgorithm). What §4.2–§4.3 contribute is the numeric vocabulary being ordered, and the fact that the ordered value is hashed and travels with data at rest. See §0.2, C2.

• Selection rule. A kernel offering grade k may serve a request r iff k ⊑ r.
• Propagation rule. grade(op(x, y, kernel)) = grade(x) ⊔ grade(y) ⊔ grade(kernel). Guarantees only ever weaken, and you can see where.

⊑, not hash equality, is the interoperability test for computation. Hash equality (plus §11's equivalence edges) is the test for interchange.

4.4 Canonicalisation

A conformant parser MUST:

1. Strip all whitespace, including newlines, before parsing.
2. Lowercase keys and values (grit1( may be spelled GRIT1().
3. Accept any input field order; emit the §4.2 order.
4. Map aliases to canonical tokens (table below). An alias that maps to a different format MUST be rejected, not silently accepted.
5. Drop fields whose value equals their default.
6. Reject: unknown keys, duplicate keys, missing required keys (e, k0, s0, esr), k1 without s1 or vice versa, zpc/zps without zp, and zps=l1 without a level-1 scale.

Rule 6's "reject unknown keys" is normative and important: an implementation that ignores a key it does not understand silently drops a contract. swizzled is not a Grade field; layout is never in the Grade.

Alias table (informative sample; the full table is in reference.py):

4.5 The decode projection

Only these fields influence decode():

{ e, k0, s0, k1, s1, zp, zpc, zps, sp, ct, po, dr, ds, dt }

The decode projection of a canonical grade string is written gritdec1(...) and contains exactly those fields, all of them explicit (defaults included), in §4.2 order. Two grades are decode-equivalent iff their decode projections are byte-identical. This is what a supersedes edge of scope decode asserts (§11.2), and it is machine-checkable rather than editorial.

5. The placement word

Layout is never a structural type parameter. There is no Grit<…, Tcgen05_128x4B>. A placement is a 64-bit value.

Carrying layout as a runtime value is well-trodden ground — cuBLASLt's cublasLtOrder_t, XLA's LayoutProto, NVIDIA's opaque CUtensorMap, Legion's layout-constraint registrar and Chapel's domain maps all do it — and §5 claims none of that. What is GRIT's is the combination in §5.1–§5.5: the value is hashed, registry-free at class 0x8, portable across processes, languages and files, sits next to the numeric contract in the same 64 bytes, and is obtainable in exactly three ways with an audited escape hatch. See §0.2, C3.

5.1 Encoding

 63 .. 60   59 .. 0
 +--------+-------------------------------+
 | class  |            payload            |
 +--------+-------------------------------+

The all-zero word 0x0000000000000000 is PLACEMENT_UNSPECIFIED: a legal, well-formed value meaning "placement not asserted." It is not a parse error. A boundary that requires a specific placement reports PlacementUnknown when it meets it.

5.2 Registered placements (version 1)

New hardware = a new 64-bit constant. Not a new type, not an ABI change, not a recompile. That is the point of layout-as-value — a property GRIT shares with, rather than takes from, cuBLASLt orders, XLA layouts and CUtensorMap (§0.2, C3). The part that is GRIT's is that the constant is hashed and travels with the bytes, so two tools can disagree about layout in a file and be told so.

5.3 The layout expression language

Class 0x8 payloads are 60-bit hashes of expressions in this grammar, which is canonically serialised and carried out of band (in a registry file, a safetensors header key, or a GGUF KV entry):

P ::= RowMajor | ColMajor
    | Tile(P, dims) | Interleave(P, n) | Swizzle(bits, base, shift)
    | Perm(π) | Pad(align) | Compose(P, P)
    | Opaque(vendor, blob)

Collisions. 2⁻⁶⁰ is an estimate of the probability of an accidental collision, under an explicitly non-adversarial assumption. FNV-1a offers no adversarial collision resistance: an attacker who wants two distinct layout expressions to hash to the same class-0x8 word can construct them.

A collision here is a spurious match, not a spurious mismatch. Two different layouts compare equal, the diagnostic diff path therefore never fires, and wrong-layout arithmetic proceeds silently — the opposite of a confusing message. Consumers MUST therefore not treat equality of two class-0x8 words as equality of placement. Before relayouting or eliding a relayout on the strength of two equal class-0x8 words, a consumer MUST fetch and compare the full canonical layout expressions, and MUST refuse the relayout when either expression is unresolvable.

5.4 Laws

• L1 (placement invariance). decode ∘ relayout(p₁→p₂) = decode. Layout is decode-invariant. This is the theorem that licenses erasing the witness at GRIT_CHECK=0.
• L2 (composition). relayout(p₂←p₁) ∘ relayout(p₁←p₀) = relayout(p₂←p₀); relayout(p←p) = id. Placements and relayouts form a groupoid, so any pair of non-opaque layouts is connected through ROW_MAJOR. GPTQ→Marlin and linear→swizzled are derivable coercions, not hand-written kernels. (Derivable layout-to-layout conversion is prior art and is not claimed here: CuTe's layout algebra and Triton's Linear Layouts derive conversions from a strictly more general representation than this groupoid. §0.2, C3.)
• L3 (opacity is one-way). OPAQUE can be produced and consumed but never relayouted out of. It is the honest hatch, and it is visibly worse in the type.
• L4 (transpose is not a view on the group axis). For the group axis, decode(transposeᵃ(x)) ≠ transposeᵃ(decode(x)) in general, because the k0-neighbourhoods change. Group-axis transpose therefore does not exist as a placement change. The only spelling is requantize(transposeᵃ(decode(x)), grade'), which returns a #[must_use] Requantized<…>. Transpose on any non-group axis is a placement change and is free.
• L5 (slice honesty). A group-axis slice whose start is not a multiple of k0 MUST return an error. Never a silent view.
• L6 (sparse metadata order). A kernel declaring requires SPORDER_ASC_LSB MUST reject an SPORDER_UNORDERED operand.
• L9 (erasure soundness). For any program well-typed at GRIT_CHECK=0, raising the check level changes observable behaviour only by turning a would-be-undefined execution into a diagnostic.

5.5 Obtaining a placement witness

A witness that bytes are in placement p is obtainable in exactly three ways:

1. Canonical construction — the constructor produced them in p.
2. relayout(x, p) — performs the permutation.
3. assume_placement(x, p, reason) — the audited escape hatch, §12.

There is no fourth way, and the third is greppable by construction.

6. Fingerprinting: grade_id, gid, digest

GRIT uses FNV-1a exclusively. It is implementable in nine lines of C with no dependencies, freestanding-safe, and requires no crypto library in a bootloader, a WebGPU shim, or a ggml build. GRIT deliberately does not use a cryptographic hash, because GRIT deliberately makes no unforgeability claim (§0.1). A forged descriptor fails grit_check; that is the whole defence.

The recipe of §6 — canonicalise a description, then fingerprint the canonical form at several widths so that independent implementations agree without a registry operator — is Apache Avro's Parsing Canonical Form and schema fingerprints (2011), reused here step for step and claimed by GRIT in no part. What §6.2–§6.3 contribute is the object being canonicalised and hashed: an arithmetic contract, plus the binary descriptor frame, plus the shape. See §0.2, C5.

6.1 FNV-1a

/* 64-bit */
uint64_t h = 0xcbf29ce484222325ULL;
for (size_t i = 0; i < n; i++) { h ^= p[i]; h *= 0x100000001b3ULL; }
/* 128-bit: offset basis 0x6c62272e07bb014262b821756295c58d
            prime        0x0000000001000000000000000000013b
   arithmetic mod 2^128; result serialised BIG-endian (16 bytes)          */

Self-check: FNV-1a-64 of "a" is 0xaf63dc4c8601ec8c; FNV-1a-128 of "a" is d228cb696f1a8caf78912b704e4a8964.

6.2 grade_id (64-bit)

grade_id = FNV1a64( canonical_grade_string, UTF-8, no NUL terminator )

It is derived, not assigned. Anyone can compute it from the string; there is no registry operator and no allocation authority.

6.3 gid (128-bit)

gid = FNV1a128(
        canonical_grade_string (UTF-8)
     || 0x00
     || desc[0 .. 28)              /* frame, part 1 */
     || desc[32 .. 64)             /* frame, part 2 -- note is skipped */
     || u32le(rank)
     || i64le(shape[0]) .. i64le(shape[rank-1])
      )

The note field at [28,32) is excluded. §11.1 explains why.

gid identifies the contract, not the bytes: two different MXFP4 tensors of the same shape and placement have the same gid and different digests. That is intended. "Do these two tools agree about what these bytes mean?" is a gid comparison; "are these the same bytes?" is a digest comparison.

6.4 Printable form: Crockford base32

gid_base32 = 26 Crockford base32 symbols, most-significant first
             alphabet "0123456789ABCDEFGHJKMNPQRSTVWXYZ"
             the 128-bit value is zero-extended on the LEFT to 130 bits
             uppercase, no hyphens, no check symbol

No hyphens on purpose: a hyphen breaks double-click selection, and the entire justification for a printable form is that somebody has to select it out of a log line at 3 a.m. The first symbol is always in 0–7.

On parse, I and L MUST be accepted as 1, O as 0, case-insensitively; U MUST be rejected. Implementations MAY display a 7-symbol prefix in compact contexts, and MUST NOT treat a prefix as an identity.

Example: gid = 42d145a24c0f96ea7e336d82a2ac740d renders as 22T52T4K0FJVN7WCVDGAHARX0D.

6.5 digest (128-bit, optional)

digest = FNV1a128( for each present plane, in order data,scale0,scale1,aux:
                     u64le(plane_bytes) || plane_bytes_content )

Length-prefixing prevents a concatenation ambiguity between planes. digest is what check level 2 verifies (§7.4). Where cryptographic binding is wanted (a signed model release), an existing signing scheme signs (gid, digest); GRIT depends on no PKI and no trusted third party.

7. grit_check and the boundary result

7.1 The must-use result type

typedef enum grit_status {
  GRIT_OK                = 0,
  GRIT_SIZE_MISMATCH     = 1,
  GRIT_GRADE_UNKNOWN     = 2,
  GRIT_PLACEMENT_UNKNOWN = 3,
  GRIT_DIGEST_MISMATCH   = 4
} grit_status_t;
#if defined(__cplusplus) && __cplusplus >= 201703L
#  define GRIT_NODISCARD [[nodiscard]]
#elif defined(__GNUC__) || defined(__clang__)
#  define GRIT_NODISCARD __attribute__((warn_unused_result))
#else
#  define GRIT_NODISCARD
#endif
GRIT_NODISCARD grit_status_t grit_check(const grit_view_t*, int level,
                                        grit_detail_t* out_detail);

The result type is must-use in every binding: GRIT_NODISCARD in C, [[nodiscard]] in C++, #[must_use] enum GritStatus (and #[must_use] Result<T, GritError> on every fallible API) in Rust, a raised exception in Python, a rejected Promise/thrown GritError in TypeScript. A failed boundary check that can be silently dropped is not a check.

The five variants are exhaustive by design: a fixed, small, closed set is what makes cross-language bindings and cross-tool diagnostics comparable. Two variants therefore carry a slightly wider meaning than their name suggests:

• GradeUnknown — "this descriptor does not denote a grade I recognise," including because its frame is unrecognisable at all (bad magic, bad version, reserved bits set, internally inconsistent flags). A descriptor whose header cannot be believed cannot have a knowable grade.
• DigestMismatch — "the plane contents disagree with what the descriptor asserts about them." That covers the plane digest and any declared content invariant, notably SPORDER_ASC_LSB metadata ordering (§7.4). Size assertions are SizeMismatch; content assertions are DigestMismatch.

grit_detail_t carries a human-readable message plus the expected/actual values and the offending plane index. Diagnostics are the product; see §7.5.

7.2 Check levels

7.3 Evaluation order (normative)

grit_check MUST evaluate in exactly this order and return the first failure. The order is normative so that two implementations report the same variant for the same broken input.

Step 1 — frame. Fail ⇒ GradeUnknown. magic == 0x31545247; hdr_bytes == 64; version == 1; n_levels ∈ {1,2}; flags & 0xC0 == 0; frame_rsvd == 0; container ∈ {0,1}; n_levels == 1 ⇒ k1 == 0 ∧ scale1 == 0; ¬(SPARSE ∧ HAS_ZP); ZP_L1SCALED ⇒ HAS_ZP ∧ n_levels == 2; elem.bias_ovr == 0; every nan_enc field ≠ 3.

Step 2 — grade. Fail ⇒ GradeUnknown. grade_id MUST resolve to a grade the implementation knows, and the descriptor's frame fields MUST agree with that grade's canonical string (n_levels, flags, elem, scale0, scale1, zp_fmt, k0, k1, sp_n, sp_m, sp_meta_fmt, sp_order, container). This catches the entire "the JSON says group_size=128 but the tensor was quantised at 32" family before a byte is read.

Step 3 — placement. Fail ⇒ PlacementUnknown. Each of placement_data, placement_scale, placement_meta MUST be PLACEMENT_UNSPECIFIED, or class 0x1 and present in the registry, or class 0x8, or class 0xF. Additionally: an aux plane is present iff placement_meta != PLACEMENT_UNSPECIFIED.

Step 4 — the size equations. Fail ⇒ SizeMismatch.

numel        = Π shape[d]
require        shape[group_axis] % k0 == 0
nnz          = SPARSE ? numel * sp_n / sp_m : numel
               (SPARSE additionally requires shape[group_axis] % sp_m == 0)
n_groups0    = numel / k0
n_groups1    = (n_levels == 2) ? (k1 == 0 ? 1 : n_groups0 / k1) : 0
               (k1 != 0 requires n_groups0 % k1 == 0)
cw(f, ct)    = (ct == CT_PAD8 && width(f) <= 8) ? 8 : width(f)
bytes(data)   = ceil( nnz       * cw(elem, container) / 8 )
bytes(scale0) = ceil( n_groups0 * cw(scale0, CT_PAD8) / 8 )
bytes(scale1) = ceil( n_groups1 * (ZP_L1SCALED ? 2 : 1) * width(scale1) / 8 )
bytes(aux)    = SPARSE ? ceil( numel / sp_m * sp_n * metabits / 8 )
              : HAS_ZP ? ceil( n_groups0 * cw(zp_fmt, CT_PAD8) / 8 )
              : 0

Every one of the four plane_bytes MUST equal its computed value exactly. Not "at least" — exactly. A too-large plane is as much a mismatch as a too-small one, because it means somebody's mental model of the layout is wrong.

These five equations catch, at O(1): wrong block size, wrong scale count, forgotten padding, dense/sparse confusion, truncated shards, and mis-served HTTP ranges.

Step 5 — content (level ≥ 2). Fail ⇒ DigestMismatch. If an expected digest was supplied, recompute per §6.5 and compare. If SPARSE ∧ sp_order == SPORDER_ASC_LSB, verify that every group's metadata indices are strictly increasing and < sp_m.

7.4 Sparsity metadata verification

PTX makes disordered mma.sp::ordered_metadata indices undefined behaviour. GRIT converts an unchecked, unnamed precondition into a DigestMismatch at the boundary. This check is O(numel/M) — cheap enough that it belongs in every model-load path, and expensive enough that it is level 2 and not level 1.

7.5 Diagnostics

Diagnostics are the product, not a courtesy. A conformant implementation SHOULD produce messages of this shape:

GritLayoutError: matmul operand B is placed SWIZZLE_128x4B (0x1000…0003) but
  kernel 'cutlass_mxfp4_sm100a' requires ROW_MAJOR (0x1000…0001).
  Fix:  B = grit.relayout(B, grit.layouts.ROW_MAJOR)
  Or:   grit.matmul(x, B, kernel='cutlass_mxfp4_sm100a_sw')
GritGradeError: requested ao=seq (deterministic); the only kernel for
  (MXFP4, 2:4, sm100a) offers ao=unspec.
  PTX 9.3, mma.sync .e2m1: "The accumulation order, rounding, and handling of
  subnormal inputs are unspecified."
  Options: grade=GRADE_MXFP4_REFERENCE (software, exact, ~180x slower)
           grade.relax(ao=unspec)      # accept nondeterminism, explicitly
GritSizeError: scale0 plane is 128 bytes; descriptor requires 256.
  numel=8192, k0=32 -> n_groups0=256, cw(e8m0)=8 -> 256 bytes.
  Likeliest cause: the checkpoint was quantised at group_size=64, not 32.

grit_explain(a, b) MUST print a field-by-field descriptor diff. That single tool resolves a "MXFP4 is type code 4 here and 39 there" incident in ten seconds.

7.6 Level 3: decode-sanity probes (informative)

Level 3 samples min(64, n_groups0) groups, decodes them, and checks distributional invariants that a structurally valid but semantically lying descriptor would violate — e.g. a claimed E8M0 scale plane whose decoded scales have no dynamic range, or an int8 plane whose decoded values never approach ±127 (suggesting the wrong scale rule). Level 3 is advisory: it reports warnings, never a grit_status_t failure, because a false positive that blocks a load is worse than the bug it hunts.

8. Normative decode()

decode() is the definition of what the bytes mean. Everything else is commentary.

8.1 The equation

For logical index i with group indices g0, g1 (§3.7), let L be the number of scale levels:

decode(i) =
   1. if SPARSE and i is structurally masked:  return +0.0     [sparsity law]
   2. S := sval(s_{L-1}, scale_{L-1}[g_{L-1}])                 [outermost first]
      for l = L-2 downto 0:  S := f32( S * sval(s_l, scale_l[g_l]) )
   3. if isnan(S): return NaN                                  [poison law]
   4. e := fmt_decode(elem, q[i])            (exact, no rounding)
      if isnan(e): return NaN
   5. if HAS_ZP:
        if ZP_L1SCALED:   Z := f32( sval(s_{L-1}, zscale[g_{L-1}])
                                    * sval(zp_fmt, zp[g0]) )
                          if isnan(Z): return NaN
                          return f32( f32(S * e) - Z )
        else:             z := fmt_decode(zp_fmt, zp[g0]) + (ZP_MINUS1 ? 1 : 0)
                          return f32( S * f32(e - z) )
      else:               return f32( S * e )
   6. project the binary32 result into dt under (dr, ds).
      For dt = binary32 this is the identity.

sval(f, code) is fmt_decode(f, code) (§2.4).

The evaluation order is fixed because reproducibility is the product. Left-to-right, outermost scale first, one binary32 rounding per named step. An implementation that multiplies the two scale levels in the other order, or that folds S * e - Z into an FMA, is non-conformant even when it is more accurate.

8.2 The poison law

If any scale in i's scale chain is NaN — including E8M0 code 0xFF, E4M3 code 0x7F/0xFF, or a binary16 NaN — then all elements of that group decode to NaN. NaN-ness propagates down the group tree, never sideways: a NaN level-1 scale poisons every element under it; a NaN level-0 scale poisons only its own k0 elements.

If the element format has a NaN encoding and q[i] is that encoding, only element i is NaN.

E2M1 has no NaN encoding: nan_enc = 0. Every 4-bit code is a finite value. A conformant decoder must not invent one.

8.3 The sparsity law

Structurally-masked positions decode to +0.0 exactly — bit pattern 0x00000000. Not "unspecified", not -0.0, not "whatever the kernel skipped".

The sparsity law wins over the poison law. A masked position under a NaN scale decodes to +0.0, not NaN, because step 1 returns before the scale chain is read. This is deliberate: a structurally-absent element has no scale.

8.4 The zero-point laws

Two conventions exist and they are not interchangeable:

• zps=shared (GPTQ shape). The zero point rides the same scale as the data: x = S · (e − z). One multiply, one rounding. Note this is not S·e − S·z, which rounds three times and gives different bits.
• zps=l1 (K-quant shape). The zero point has its own level-1 scale stored alongside the data's, and the offset is subtracted after scaling: x = (S · e) − Z. This is required for affine schemes where the min is quantised independently of the step, and it cannot be expressed by S·(e−z) without loss.

ZP_L1SCALED also changes the scale1 plane's shape: it stores two values per level-1 group, [S₁, Z₁], in that order.

ZP_MINUS1 (the GPTQ one-off). GPTQ stores zp − 1 and adds 1 back at load. When set, z = stored + 1. This field exists because the convention lived in code rather than in the file for years, silently degrading asymmetric quantisation quality across an entire ecosystem. One flag bit, one load-time check.

8.5 Semantic equality

Two GRIT arrays are semantically equal iff their decodes are equal. Any operation whose result depends on the representation beyond decode is illegal without an explicit Grade argument. That single rule is what makes group-axis transpose and misaligned slicing impossible to write by accident (L4, L5).

9. Normative encode() / quantize()

encode(vals, grade) -> planes is the projection of a binary32 array into a Grade. It is lossy; decode(encode(v)) = π_enc(v) blockwise.

encode specifies the projection, not the calibration. Choosing which values to feed it — AWQ, SmoothQuant, GPTQ's Hessian-weighted search, GGUF's iterative make_qkx refinement — is out of scope and deliberately so. Two tensors with the same gid can come from very different calibration and differ enormously in downstream accuracy. gid certifies decode agreement, nothing more. Do not over-read it.

9.1 Universal rules

1. All intermediate arithmetic is binary32 RNE, per §1.1.
2. Poison on encode. If any input in a group is NaN or ±∞, the group's scale is set to its format's NaN code and every element code in that group is set to 0. decode then yields NaN for the whole group, by §8.2. This is the only way to represent NaN under an element format with nan_enc = 0.
3. All-zero group. If amax == 0, the scale is set to the format's encoding of 1.0 and all element codes to the (possibly signed) zero code. Never a zero scale, never a NaN scale.
4. floor(log2(amax)) MUST be computed exactly from the binary32 exponent — a read of the biased exponent bit-field, adjusted for subnormals (e.g. frexp(amax).exp − 1) — never via a floating-point log2 that can be off by one at powers of two.
5. Element projection uses fmt_encode_rne (§2.5) with es saturation.
6. amax is used unmodified. amax MUST be computed from the input values exactly as the scale rule states and fed to the esr rule as-is: no pre-scaling, no epsilon nudging, no mantissa inspection, no threshold adjustment. An implementation that modifies amax before applying the esr rule is non-conformant.

9.2 esr = exp2floormaxabs (the MX rule)

amax   = max |v| over the group (all finite)
e_floor = floor(log2(amax))
X_exp   = clamp(e_floor - emax_elem, -127, +127)
scale   = 2^X_exp                       stored as E8M0 code (X_exp + 127)
q[i]    = fmt_encode_rne(elem, f32(v[i] / scale), es)

emax_elem is the maximum exponent of the element format: 2 for E2M1 (max normal 6.0 = 1.5 · 2²). The clamp is normative: an extremely small amax saturates the shared exponent at -127 and the elements underflow to zero rather than producing an unrepresentable scale.

exp2ceilmaxabs and exp2rnemaxabs differ only in step 2's rounding. They are different grades (§11). Both roundings MUST be computed exactly, from e_floor and the significand f = amax / 2^e_floor (f in [1, 2), exact in binary64 for any binary32 amax), never via a floating-point log2:

• ceil(log2(amax)) is e_floor when f = 1 (amax an exact power of two), else e_floor + 1.
• rne(log2(amax)) is e_floor + 1 when f² > 2 (that is, f > √2), else e_floor. f² is exact in binary64 for binary32 inputs, so the comparison is never off by a rounding error. A true tie (f = √2 exactly) cannot occur, because a binary32 significand is rational; an implementation that handles the case anyway MUST round the exponent half to even.

GRADE_MXFP4_RNE (§10.0) instantiates exp2rnemaxabs and is covered by conformance vectors. exp2ceilmaxabs has no registered grade in version 1; the rule above is normative for any grade that names it, and all five reference implementations (C, C++, Rust, Python, TypeScript) cover it with direct unit tests rather than vectors.

9.3 esr = affineminmax (the two-sided rule)

Per level-0 group: lo = min(0, min v), hi = max(0, max v), step = (hi − lo)/(2^w − 1), off = −lo. Zero is forced into the range so that the all-zero input is exactly representable. §10.2 gives the full two-level instantiation.

9.4 esr = ratiomaxabs

scale = amax / maxnorm_elem, then projected into the scale format under (er, es). maxnorm_elem is 6.0 for E2M1, 448.0 for E4M3, 127.0 for s8 (§1.5: the asymmetric minimum is never produced).

10. The four reference grades

These four are normative. Every conformant implementation MUST implement all four exactly, and MUST reproduce conformance/vectors.json bit-for-bit.

10.0 Summary

Two sibling grades exist to exercise §11:

10.1 GRADE_MXFP4

Representation. E2M1 elements, blocks of 32 along the group axis, one E8M0 shared scale per block. One level. No zero point, no sparsity. Dense nibble packing, LSB-first.

Descriptor.

Planes. data = numel/2 bytes; scale0 = numel/32 bytes; scale1 = aux = empty.

decode.

g0 = i / 32
S  = (scale0[g0] == 0xFF) ? NaN : 2^(scale0[g0] - 127)
if isnan(S): return NaN                            # poison law
e  = E2M1[ q[i] ]                                  # exact; no NaN in E2M1
return f32(S * e)

encode. Per block of 32:

1. Any input NaN or ±∞ ⇒ scale0 = 0xFF, all 32 codes 0. Done.
2. amax = max|v|. If amax == 0 ⇒ scale0 = 0x7F (2⁰) and each code is fmt_encode_rne(e2m1, v) — which preserves −0.0 as code 8. Done.
3. X_exp = clamp(floor(log2(amax)) − 2, −127, +127); scale0 = X_exp + 127.
4. q[i] = fmt_encode_rne(e2m1, f32(v[i] / 2^X_exp), satfinite).

Worked tie check (vector mxfp4/tie_to_even, S = 1.0):

Every one of the first seven is an exact tie broken to the even code; 6.5 saturates.

10.2 GRADE_Q4KS

Representation. GGUF Q4_K-shaped two-level affine quantisation. Unsigned 4-bit elements, sub-blocks of 32, super-blocks of 8 sub-blocks (256 elements). Per super-block: two binary16 values d (scale-of-scales) and dmin (scale-of-mins). Per sub-block: a u6 scale code sc and a u6 min code m.

Spec decision, do not misread. The elements are u4 (unsigned 0..15), not s4. Q4_K nibbles are unsigned and the affine offset carries the sign. Spelling this int4 — as informal prose often does — produces a tensor that decodes to garbage. grit_check step 2 catches it.

Decode form. x = (d · sc) · q − (dmin · m). This is affine, not symmetric; it is why zps=l1 exists (§8.4).

Descriptor.

Planes, per 256-element super-block (148 bytes total):

The canonical GRIT scale planes use CT_PAD8 (16 bytes per super-block for sc+m), not GGUF's 12-byte get_scale_min_k4 bit-packing. That packing is a placement, registered as PACK_K4_12B = 0x1000000000000020, reachable by relayout. GRIT is a descriptor language that can name ggml's ~30 block structs; it is not a 31st block struct. The 6-bit unpack that has been hand-ported into C, CUDA, Metal, GLSL, SYCL, Rust and Python becomes one placement constant.

decode.

g0 = i / 32;  g1 = g0 / 8
S = binary16_value( scale1[2*g1]     )              # d,    outermost first
S = f32( S * u6_value( scale0[g0] ) )               # * sc
if isnan(S): return NaN
Z = binary16_value( scale1[2*g1 + 1] )              # dmin
Z = f32( Z * u6_value( aux[g0] ) )                  # * m
if isnan(Z): return NaN
e = (float) q[i]                                    # u4, 0..15
return f32( f32(S * e) - Z )

encode. Per 256-element super-block:

1. Any input NaN or ±∞ ⇒ d = 0x7E00 (binary16 NaN), dmin = 0x0000, all sc, m, q = 0. Done.
2. For each of the 8 sub-blocks s: lo_s = min(0, min v), hi_s = max(0, max v), step_s = f32(f32(hi_s − lo_s) / 15), off_s = f32(−lo_s).
3. d = binary16_rne( f32(max_s step_s / 63) ), dmin = binary16_rne( f32(max_s off_s / 63) ).
4. sc_s = clamp(rne_int(f32(step_s / d_val)), 0, 63) (0 if d_val == 0); m_s = clamp(rne_int(f32(off_s / dmin_val)), 0, 63) (0 if dmin_val == 0).
5. S_s = f32(d_val · sc_s), Z_s = f32(dmin_val · m_s); q[i] = clamp(rne_int(f32(f32(v[i] + Z_s) / S_s)), 0, 15), or 0 if S_s == 0.

Note step 3 and step 5 both re-read d and dmin after the binary16 round-trip. Using the pre-rounding binary32 values gives different bits and is non-conformant.

10.3 GRADE_NVF4

Representation. NVFP4-shaped two-level scaling. E2M1 elements, blocks of 16, one E4M3 block scale per block, one binary32 per-tensor scale (k1 = 0). No zero point, no sparsity.

Descriptor.

Planes. data = numel/2; scale0 = numel/16 bytes; scale1 = 4 bytes (exactly one binary32); aux empty.

decode.

g0 = i / 16
S = binary32_value( scale1[0] )                     # per-tensor, outermost
S = f32( S * e4m3_value( scale0[g0] ) )
if isnan(S): return NaN                             # E4M3 0x7F/0xFF poison
e = E2M1[ q[i] ]
return f32(S * e)

encode.

1. amax = max |v| over all finite elements of the whole tensor. If amax == 0 ⇒ sT = 1.0. Otherwise sT = f32( f32(amax / 6.0) / 448.0 ); if that underflows to 0, use the smallest positive binary32 subnormal 2^-149. Store sT as binary32, then re-read it — subsequent steps use the stored value.
2. Per block of 16:
   • Any NaN or ±∞ ⇒ scale0 = 0x7F (E4M3 NaN), all 16 codes 0. Done.
   • bamax = max|v|. If bamax == 0 ⇒ scale0 = fmt_encode_rne(e4m3, 1.0). Otherwise scale0 = fmt_encode_rne(e4m3, f32(f32(bamax / 6.0) / sT)).
   • S = f32(sT · e4m3_value(scale0)).
   • q[i] = fmt_encode_rne(e2m1, f32(v[i] / S), satfinite), or the signed zero code if S == 0.

S == 0 is reachable and is not an error: when a block's magnitude is more than ~2¹⁸ below the tensor amax, its E4M3 block scale underflows to zero and the block decodes to ±0. That is NVFP4's honest dynamic-range limit, and the conformance vector nvf4/wide_dynamic_range pins it so nobody "fixes" it differently in each framework.

10.4 GRADE_I8S24

Representation. Signed 8-bit elements with 2:4 structured sparsity, one binary16 scale per 32 logical elements. One scale level. No zero point.

Descriptor.

Planes. data = numel/2 bytes (only survivors, s8); scale0 = numel/16 bytes (one binary16 per 32 logical elements); scale1 empty; aux = numel/8 bytes (one nibble per quad).

k0 = 32 counts logical positions: one scale covers 8 quads = 16 stored values.

decode.

for each quad q (4 logical positions):
  keep = { meta_index(q,0), meta_index(q,1) }
  for j in 0..3:
     i = 4*q + j
     if j not in keep:  out[i] = +0.0;  continue      # sparsity law, wins
     S = binary16_value( scale0[i / 32] )
     v = (int8) data[stored++]                        # ascending logical order
     if isnan(S): out[i] = NaN; continue              # poison law
     out[i] = f32(S * (float)v)

The stored cursor advances only for surviving positions and only in ascending logical order. This is the whole reason SPORDER_ASC_LSB is checkable: the data cursor and the metadata must agree, and §7.4 verifies it.

encode.

1. Selection, per quad of 4: keep the two largest by magnitude. NaN sorts as maximal magnitude (so a NaN can never be silently dropped by sparsification); ties in magnitude break to the lower logical index. Emit the two kept indices in ascending order.
2. Scale, per group of 32 logical / 16 kept values:
   • Any kept value NaN or ±∞ ⇒ scale0 = 0x7E00 (binary16 NaN), all 16 codes 0. Done.
   • amax = max |kept|. If amax == 0 ⇒ scale0 = binary16(1.0). Otherwise scale0 = binary16_rne(f32(amax / 127.0)); if that rounds to zero, use 0x0001 (the smallest positive binary16 subnormal). A zero scale is never emitted.
3. q = clamp(rne_int(f32(v / S)), −127, +127) for each kept value, in ascending logical order, where S is the decoded binary16 scale.

11. Supersedes edges and the frame/note split

GRIT's own worst risk is that grade canonicalisation is subtly wrong for years, trading silent numeric corruption for loud spurious errors. Two mechanisms bound that risk.

11.1 The frame/note split

Descriptor bytes [28,32) — note — are informational and MUST NOT be hashed into gid. Everything else is the frame and is hashed.

The rule for deciding which side a field belongs on:

A field belongs in the frame iff changing it changes what the bytes mean. A field belongs in the note iff changing it changes only what somebody said about the bytes.

note layout:

This is why the escape-hatch audit flag lives in note and not in flags: an assumption about the layout is provenance, not semantics. Two artifacts, one of which was hand-verified and one of which was machine-derived, describe the same bytes and MUST have the same gid — otherwise every escape-hatch use would fork the fingerprint namespace and destroy interchange. Conformance vector fp/frame_note_split pins this.

Corollary: the note is not authenticated by gid. It is a hint for humans and for grit scan, never an input to a correctness decision.

11.2 Supersedes equivalence edges

A supersedes edge is a registry record:

{ from: GradeName, to: GradeName, scope: "decode" | "full", reason: string }

• scope = "decode" — the two grades' decode projections (§4.5) are byte-identical. Stored bytes are interchangeable; re-quantisation is not.
• scope = "full" — encode and decode both agree. Fully interchangeable.

Equivalence is the reflexive-symmetric-transitive closure of the edge set at or above the requested scope, plus canonical-string identity. Interchange of stored bytes requires decode-scope equivalence. Re-quantisation pipelines require full scope.

An edge MUST NOT be registered without a reason string, and MUST NOT be registered unless the two grades' decode projections actually match (for decode scope) or the whole conformance suite agrees under both (for full). The projection check is mechanical, which is what stops an edge set from becoming a folklore repository.

Registered edges in version 1:

Both have decode projection equality: gritdec1(e=e2m1,k0=32,s0=e8m0,ct=dense,po=lsb,dr=rne,ds=satfinite,dt=binary32) for the MXFP4 pair.

What this buys. A descriptor-version bump, a renamed alias, or a newly explicit field no longer orphans every fingerprint minted under the old spelling. It also gives GRIT the escape valve that n_levels otherwise lacks: when a three-level scheme forces a descriptor version bump, the v1 and v2 spellings of an unchanged grade are joined by a full-scope edge rather than being two unrelated hashes.

What it does not buy. An edge is a claim, and a wrong edge is silent corruption — exactly the failure mode GRIT exists to prevent. Hence: mechanical precondition, mandatory reason string, and every edge carries its own conformance vector (vectors.json → supersedes).

12. assume_placement and the audit ledger

Every checked-witness system lives or dies on how often the escape hatch is used. GRIT's named death mode is "somebody wrote assume_placement." The hatch must exist — nobody may ever be blocked — so instead it is made measurable.

/// SAFETY: the caller asserts that `x`'s bytes are already in placement B.
/// `why` MUST be a compile-time constant so it is greppable.
pub unsafe fn assume_placement<S, G, const A: u64, const B: u64>(
    x: Grit<S, G, A>, why: &'static str
) -> Grit<S, G, B>;

Normative requirements:

1. A non-empty reason string is mandatory. An API that accepts an empty reason, or that defaults it, is non-conformant. In languages with static strings the parameter MUST be &'static str / constexpr const char* / Literal[str], so grep -rn assume_placement finds every site and its justification.

2. The NOTE_ASSUMED bit becomes sticky. It is set on the result and propagates through relayout, slice, concat and serialisation. It is cleared only by verify_placement(x, p), which actually reads the bytes and proves the claim.

3. A ledger record is appended, append-only, per process:

   { timestamp, gid_before, gid_after, placement_from, placement_to,
     reason, source_location }
   

   Exposed as grit.audit.ledger(); serialisable to JSONL so CI can diff it.

4. A rate is exposed. grit.audit.rate() returns assumed_crossings / total_boundary_crossings for the process. A build MAY set a ceiling (GRIT_AUDIT_MAX_RATE) above which assume_placement itself fails. Escape-hatch abuse becomes a number on a dashboard rather than an anecdote.

5. grit scan reports it. Any on-disk artifact carrying NOTE_ASSUMED is flagged, with its producer tag, so an assumption made once in a conversion script stays visible to everyone downstream forever.

Because note is not hashed (§11.1), none of this perturbs gid. Audit and identity are orthogonal, on purpose.

13. grit scan — the diagnostic wedge

GRIT delivers value before anyone adopts the type. grit scan requires zero code changes to any framework. Tools that print the fields of a GGUF or safetensors artifact are common; as of the 2026-08 prior-art review (§0.2, C6b) no tool was found that checks a declared numeric contract against the bytes of an existing artifact, which is what this section specifies.

grit scan [--deep] [--json] [--exit-policy graded|simple] PATH...

PATH may be a safetensors file, a GGUF file, a directory tree, or a manifest. scan recovers or infers a descriptor for every tensor it can — from a __grit__ safetensors header key, from GGUF KV entries, or by pattern-matching the well-known block structs — and then reports:

1. Descriptor-vs-bytes violations. The §7.3 size equations, run against the actual on-disk plane lengths. This alone catches wrong group_size, truncated shards, and dense/sparse confusion.
2. Cross-file gid disagreement. Tensors that ought to be the same operand across two tools, files, or revisions, with the field-by-field grit_explain diff for each disagreement.
3. Grade-drift clusters. Tensors within one checkpoint carrying more than one grade_id for the same logical role — the shape of a half-converted checkpoint.
4. Escape-hatch exposure. Every artifact with NOTE_ASSUMED set, grouped by producer tag.
5. --deep: level-2 digest and metadata-order verification; --deep additionally reports level-3 probe warnings (§7.6).

Exit codes, under the default --exit-policy=graded: 0 clean; 1 structural violation (any non-Ok grit_status_t); 2 disagreement without a structural violation (e.g., cross-file gid mismatch, role-scope grade drift, declared-vs-actual quantisation, undeclared zero-point convention); 3 warnings only (e.g., NOTE_ASSUMED, level-3 probes, inferred descriptors, unregistered grades, partial quantisation, artifact-scope grade drift); 4 could not parse the input at all. The parenthesised examples are illustrative, not exhaustive.

--exit-policy=simple collapses the same report into the three values a shell caller usually wants: 0 clean, 1 diagnostics found (any severity), 2 errors (no input could be parsed at all). The policy affects the process exit status only — the report is identical, and --json always carries both codes. Neither policy is normative; grit scan is diagnostic tooling, and nothing in this section constrains decode, encode, grit_check or any fingerprint.

grit explain A B prints the field-by-field descriptor diff for any two descriptors, canonical grade strings included. It is the ten-second resolution for "this model loads in tool X and produces garbage in tool Y."

14. Degenerate byte patterns

A 64-byte buffer arriving from a network, a mmap, an uninitialised allocation, or an attacker is still a grit_desc-shaped thing. Version 1 requires:

Both 00 × 64 and FF × 64 MUST be well-formed inputs and MUST be safe.

Concretely, for any 64-byte input whatsoever:

• grit_check MUST NOT read outside the 64 bytes plus the caller-supplied shape and plane_bytes.
• grit_check MUST NOT dereference any plane pointer at level 1.
• grit_check MUST NOT trap, panic, abort, divide by zero, or overflow. Note the size equations divide by k0, sp_m and k1; each divisor MUST be tested for zero before the division, and the frame checks in step 1 run first regardless.
• grit_check MUST terminate and MUST return exactly one grit_status_t.

Specified outcomes:

Both are pinned as conformance vectors (degenerate/all_zeros_descriptor, degenerate/all_ones_descriptor).

Additionally, degenerate field values are legal rather than exceptional and MUST NOT be special-cased into errors:

• placement_* == 0 is PLACEMENT_UNSPECIFIED — a legal value (§5.1).
• placement_* == 0xFFFFFFFFFFFFFFFF is class 0xF = OPAQUE — a legal, storable, non-relayoutable value.
• k1 == 0 with n_levels == 2 is per-tensor scaling, not "missing".
• note == 0 is the normal, unremarkable case.

15. Conformance

An implementation is conformant iff, running only the algorithms in this document, it reproduces spec/conformance/vectors.json exactly. That file is generated by spec/reference.py and re-verified by it:

python3 spec/reference.py gen      # regenerate
python3 spec/reference.py verify   # re-derive everything and check

15.1 What the suite contains

15.2 Vector schema

Positive vectors (vectors[]):

Negative vectors (negative[]) carry desc_hex, shape, payload_hex, plane_bytes, check_level, optional expected_digest, and expected_error ∈ {SizeMismatch, GradeUnknown, PlacementUnknown, DigestMismatch}.

15.3 The conformance claim

L8. Equal gid + equal bytes ⇒ identical decode() on every conformant implementation. This is a test, not a promise. Anyone can claim conformance by running the suite; anyone can refute a claim the same way.

Appendix A. Constant tables

A.1 Descriptor constants

#define GRIT_MAGIC              0x31545247u   /* 'G','R','T','1' LE */
#define GRIT_HDR_BYTES          64
#define GRIT_VERSION            1
#define GRIT_F_HAS_ZP           0x01
#define GRIT_F_ZP_MINUS1        0x02
#define GRIT_F_SPARSE           0x04
#define GRIT_F_MSB_FIRST        0x08
#define GRIT_F_INTERLEAVED      0x10
#define GRIT_F_ZP_L1SCALED      0x20
/*      reserved                0xC0  -- MUST be 0 */
#define GRIT_CT_DENSE           0
#define GRIT_CT_PAD8            1
#define GRIT_SPMETA_NONE        0
#define GRIT_SPMETA_U2          1
#define GRIT_SPMETA_U4          2
#define GRIT_SPORDER_UNORDERED  0
#define GRIT_SPORDER_ASC_LSB    1
#define GRIT_NOTE_ASSUMED       0x00000001u   /* NOT hashed */
#define GRIT_PLACEMENT_UNSPECIFIED  0x0000000000000000ull
#define GRIT_P_ROW_MAJOR            0x1000000000000001ull
#define GRIT_P_COL_MAJOR            0x1000000000000002ull
#define GRIT_P_SWIZZLE_128X4B       0x1000000000000003ull
#define GRIT_P_TILE_16X16           0x1000000000000004ull
#define GRIT_P_BLOCK_AOS            0x1000000000000005ull
#define GRIT_P_LINEAR_META          0x1000000000000011ull
#define GRIT_P_META_M16N8K32        0x1000000000000012ull
#define GRIT_P_PACK_K4_12B          0x1000000000000020ull
#define GRIT_FMT_E2M1           0x2043u
#define GRIT_FMT_E4M3           0x2887u
#define GRIT_FMT_E8M0           0x4807u
#define GRIT_FMT_BINARY16       0x356Fu
#define GRIT_FMT_BINARY32       0x371Fu
#define GRIT_FMT_U2             0x43E1u
#define GRIT_FMT_U4             0x43E3u
#define GRIT_FMT_U6             0x43E5u
#define GRIT_FMT_U8             0x43E7u
#define GRIT_FMT_S4             0x03E3u
#define GRIT_FMT_S8             0x03E7u
#define GRIT_GRADE_MXFP4        0x2370EE9EE454A183ull
#define GRIT_GRADE_Q4KS         0x1C600DE02594FB55ull
#define GRIT_GRADE_NVF4         0x00FF6F4C2870831Eull
#define GRIT_GRADE_I8S24        0xBBBB283A2539EBE8ull
#define GRIT_GRADE_MXFP4_RNE    0x254285E10D6068C8ull
#define GRIT_GRADE_NVF4_RTZ     0x006C3850A3F42DEAull

A.2 E2M1 value table (complete)

emax = 2, maxnorm = 6.0, no Inf, no NaN, bias 1, one subnormal magnitude (0.5).

A.3 E8M0 and E4M3

E8M0: code c ∈ [0,254] ⇒ 2^(c − 127); code 255 ⇒ NaN. No sign, no zero. 0x7F is 1.0. 0x00 is 2^-127 (a binary32 subnormal, exactly representable).

E4M3: bias 7, no Inf. 0x00 = +0.0, 0x80 = −0.0; 0x01 = 2^-9 = 0.001953125 (smallest subnormal); 0x38 = 1.0; 0x7E = 448.0 (max finite); 0x7F and 0xFF = NaN.

A.4 Reference descriptors (hex, 64 bytes each)

Each block below is 128 hex characters, wrapped at 32 per line (= 16 bytes per line, so line n covers descriptor offsets 16n .. 16n+15).

GRADE_MXFP4                          offsets
  47525431400101004320074800000000    0..15
  00002000000000000000000000000000   16..31
  01000000000000100100000000000010   32..47
  000000000000000083a154e49eee7023   48..63
GRADE_Q4KS
  4752543140010221e343e5436f35e543    0..15
  00002000080000000000000000000000   16..31
  01000000000000100100000000000010   32..47
  010000000000001055fb9425e00d601c   48..63
GRADE_NVF4
  4752543140010200432087281f370000    0..15
  00001000000000000000000000000000   16..31
  01000000000000100100000000000010   32..47
  00000000000000001e8370284c6fff00   48..63
GRADE_I8S24
  4752543140010104e7036f3500000000    0..15
  00002000000002040101000000000000   16..31
  01000000000000100100000000000010   32..47
  1100000000000010e8eb39253a28bbbb   48..63

Reading GRADE_MXFP4 line 1: 47525431 = magic 'GRT1'; 40 = hdr_bytes 64; 01 = version; 01 = n_levels; 00 = flags; 4320 = elem 0x2043 LE; 0748 = scale0 0x4807 LE; 0000 0000 = scale1, zp_fmt absent.

(Line breaks are cosmetic. The authoritative copies are the desc_hex fields in conformance/vectors.json.)

Appendix B. Worked examples

B.1 mxfp4/tie_to_even

shape        [32]
inputs       0.25 0.75 1.25 1.75 2.5 3.5 5.0 6.5
             -0.25 -0.75 -1.25 -1.75 -2.5 -3.5 -5.0 -6.5
             0.0 x16
amax         6.5,  floor(log2 6.5) = 2,  X_exp = 2 - 2 = 0,  scale0 = 0x7F
payload      20 42 64 76 a8 ca ec fe 00 00 00 00 00 00 00 00 | 7f
             ^ data (16 bytes)                                 ^ scale0
decode       0 1 1 2 2 4 4 6  -0 -1 -1 -2 -2 -4 -4 -6  0 x16
plane_bytes  [16, 1, 0, 0]
digest       9452a02851ba1a1d793b97df51266001
gid          42d145a24c0f96ea7e336d82a2ac740d
gid_base32   22T52T4K0FJVN7WCVDGAHARX0D

Reading the first data byte 0x20: low nibble 0 = element 0 = +0.0 (the tie at 0.25 broke to the even code); high nibble 2 = element 1 = 1.0 (the tie at 0.75 broke to the even code 2, i.e. away from 0.5).

B.2 i8s24/quad_pattern

shape        [32]
inputs       (3.0, 0.0, 0.0, 1.0) x 8
selection    per quad, keep indices {0, 3}  -> ascending, ASC_LSB satisfied
meta nibble  idx0=0 (bits 0..1), idx1=3 (bits 2..3)  -> 0b1100 = 0xC
meta plane   cc cc cc cc            (4 bytes, two quads per byte)
amax         3.0  ->  scale0 = binary16_rne(3.0/127) = 0x260c -> 0.023621...
data plane   7f 2a  x 8             (127 and 42; 3.0/S ~ 127.0, 1.0/S ~ 42.3)
decode       2.99981689453125, +0.0, +0.0, 0.9920654296875,  (repeated)
plane_bytes  [16, 2, 0, 4]
gid_base32   2X7NNMYRJ9K6508JQTR1J37VR9

Note the masked positions are +0.0 exactly (0x00000000), by the sparsity law, and that the descriptor's placement_meta is LINEAR_META — a real 2:4 kernel would carry META_M16N8K32 instead and the two are a PlacementUnknown apart, not a silent reinterpretation.

B.3 supersedes in bytes

supersedes/demo/mxfp4_rne_bytes_decode_identically carries one payload and two descriptors — GRADE_MXFP4_RNE (0x254285E10D6068C8) and GRADE_MXFP4 (0x2370EE9EE454A183) — with bit-identical decoded_f32_bits. Two gids, one meaning. The decode-scope edge says interchange is legal; the absence of a full-scope edge says re-quantising one as the other is not.

End of normative specification.

References

(Non-normative. These works are cited by §0.2 and by the softening notes in §2, §4, §5, §6 and §13. Nothing here imposes a requirement; an implementer needs only §1–§15 and the appendices. URLs were last checked 2026-08-01.)

Standards and specifications

1. Apache Avro Specification 1.11 — "Parsing Canonical Form for Schemas" and "Schema Fingerprints" (64/128/256-bit). Apache Software Foundation, 2011–. https://avro.apache.org/docs/1.11.0/spec.html
2. Khronos Data Format Specification, v1.0–v1.4. The Khronos Group. Parametric descriptors for numeric formats (including E4M3FN, E4M3FNUZ, E5M2FNUZ), canonical field ordering, totalSize skip rule. https://registry.khronos.org/DataFormat/specs/1.4/dataformat.1.4.html
3. DLPack — DLDataType, bits/lanes, and the LSB-first sub-byte packing rule. https://github.com/dmlc/dlpack/blob/main/include/dlpack/dlpack.h
4. Apache Arrow C Data Interface. Apache Software Foundation. https://arrow.apache.org/docs/format/CDataInterface.html
5. OCP Microscaling Formats (MX) Specification v1.0. Open Compute Project, 2023. https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf
6. IEEE SA P3109 Working Group, Arithmetic Formats for Machine Learning (draft; Interim Report v3.2.0, 2025–2026). §VIII explicitly disclaims memory and wire representation. https://sagroups.ieee.org/p3109wgpublic/ — see also A. Fitzgibbon et al., "Novel Aspects of IEEE SA P3109 Arithmetic Formats for Machine Learning", arXiv:2606.04028.
7. StableHLO specification — quantized tensor element types and DotAlgorithm (the dot_general algorithm attribute; also jax.lax.DotAlgorithm), OpenXLA, 2024–. https://openxla.org/stablehlo/spec
8. XLA LayoutProto — minor_to_major, tiles, element_size_in_bits, sparsity. OpenXLA, ~2017–. https://openxla.org/xla/shapes
9. MLIR quant dialect — UniformQuantizedType, UniformQuantizedPerAxisType, UniformQuantizedSubChannelType (Dec 2024). https://mlir.llvm.org/docs/Dialects/QuantDialect/ — RFCs: "Supporting sub-channel quantization in MLIR" (https://discourse.llvm.org/t/rfc-supporting-sub-channel-quantization-in-mlir/82694) and "Extending UniformQuantizedType with interface-based support for new storage types" (RFC 87803, https://discourse.llvm.org/t/rfc-extending-uniformquantizedtype-with-interface-based-support-for-new-storage-types-in-quant-dialect/87803).
10. TOSA RFC: "Expressing block-scaled operations in TOSA" (!tosa.block_scaled), Jun 2026. https://discourse.llvm.org/t/rfc-expressing-block-scaled-operations-in-tosa/91056
11. ONNX — DequantizeLinear / QuantizeLinear blocked quantization (https://onnx.ai/onnx/operators/onnx__DequantizeLinear.html) and issue #7691, "nested K-quant blocks", Feb 2026 (onnx/onnx#7691).
12. FPCore 2.0 — rounding contexts. FPBench. https://fptalks.org/spec/fpcore-2.0.html
13. LLVM Language Reference — constrained floating-point intrinsics. https://llvm.org/docs/LangRef.html#constrained-floating-point-intrinsics

Implementations, runtimes and libraries

14. vLLM / Neural Magic compressed-tensors — safetensors extension for sparse-quantized checkpoints (QuantizationArgs, QuantizationStrategy incl. TENSOR_GROUP, CompressionFormat). https://github.com/vllm-project/compressed-tensors
15. torchao — AffineQuantizedTensor, Layout dataclasses, MXTensor, NVFP4Tensor, ScaleCalculationMode. https://github.com/pytorch/ao and https://docs.pytorch.org/ao/stable/contributing/quantization_overview.html
16. ROCm TensorCast (tcast) — string-grammar DataType (element spec × scale spec). https://github.com/ROCm/tensorcast
17. Qualcomm QNN / QAIRT — Qnn_QuantizeParams_t with QNN_QUANTIZATION_ENCODING_BLOCKWISE_EXPANSION (LPBQ): per-channel float × per-block integer scales, offsets, axis and block size in a POD C struct. https://docs.qualcomm.com/nav/home/structQnn__QuantizeParams__t.html
18. NVIDIA cuBLASLt — cublasLtOrder_t (layout as a runtime constant on a descriptor, ~2018–) and the Blackwell matrix-scale modes (CUBLASLT_MATMUL_MATRIX_SCALE_VEC32_UE8M0, VEC16_UE4M3, …). https://docs.nvidia.com/cuda/cublas/
19. NVIDIA CUtensorMap / TMA — cuTensorMapEncodeTiled; a 128-byte opaque POD runtime descriptor carrying dtype, tile shape, swizzle and interleave. https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TENSOR__MEMORY.html
20. CuTe layout algebra (NVIDIA CUTLASS). https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/02_layout_algebra.html
21. Triton Linear Layouts — tensor layouts as binary matrices over 𝔽₂, with generic layout-to-layout conversion; ASPLOS 2026. arXiv:2505.23819. https://arxiv.org/abs/2505.23819
22. QONNX — IntQuant / FloatQuant custom ops, fastmachinelearning. https://github.com/fastmachinelearning/qonnx
23. Unison — content-addressed definition hashes. https://www.unison-lang.org/docs/language-reference/hashes/
24. Chapel user-defined domain maps and Legion LayoutConstraintRegistrar (layout constraint sets registered under a global id). https://chapel-lang.org/publications/hotpar10-final.pdf

Papers

25. D. Orchard, V.-B. Liepelt, H. Eades III. "Quantitative Program Reasoning with Graded Modal Types." Proc. ACM Program. Lang. 3(ICFP), 2019 (Granule). https://www.cs.kent.ac.uk/people/staff/dao7/publ/granule-icfp19.pdf
26. A. E. Kellison, J. Hsu. "Numerical Fuzz: A Type System for Rounding Error Analysis." PLDI 2024.
27. Z. Chen, A. Lafont, L. O'Connor, G. Keller, C. McLaughlin, V. Jackson, C. Rizkallah. "Dargent: A Silver Bullet for Verified Data Layout Refinement." POPL 2023. https://dl.acm.org/doi/10.1145/3571240
28. M. Noonan. "Ghosts of Departed Proofs (Functional Pearl)." Haskell Symposium 2018 — phantom witnesses. https://kataskeue.com/gdp.pdf — and R. Atkey, "Syntax and Semantics of Quantitative Type Theory", LICS 2018, for erasure.
29. "Kernel Contracts: A Specification Language for ML Kernel Correctness Across Heterogeneous Silicon." arXiv:2604.22032, Apr 2026. https://arxiv.org/abs/2604.22032
30. D. Vasilev. "An 83/84-Format Numeric Catalog with Bit-Exact Conformance Vectors: A Vendor-Neutral Reference for FP8, BF16, MXFP4, and Microscaling Formats." arXiv:2606.09686, Jun 2026 — SHA-256-fingerprinted JSON conformance packs cross-compiled to several languages, with an IEEE P3109 v3.2.0 cross-walk. Cited here as prior art for the fingerprinted conformance-vector genre; the citation is made regardless of authorship, and if it shares an author with this document it is prior art against this document all the same (cf. arXiv:2606.01138).
31. L. Wang et al. "Ladder: Enabling Efficient Low-Precision Deep Learning Computing through Hardware-aware Tensor Transformation." OSDI 2024 (tType). https://www.usenix.org/conference/osdi24/presentation/wang-lei

Ecosystem evidence cited in §4.2 and §13

32. pytorch/pytorch#146414 — no consensus on the MX shared-scale derivation rule; the OCP floor rule differs from a vendor's modified-ceiling rule (the esr fork). pytorch/pytorch#146414
33. vllm-project/vllm#8913 — RFC requesting "a more sophisticated way of describing packed layouts"; open since 2024. vllm-project/vllm#8913