Hurray Format Specification

Version: 0.1.0-draft

Hurray is a language-agnostic, zero-copy runtime interchange format for multi-dimensional tensor data, optimized for the memory layout diversity, quantization schemes, and access patterns of modern AI/ML inference pipelines and scientific arrays.

Scope and Goals

• Define a binary tensor descriptor encoding that is language- and runtime-agnostic.
• Enable zero-copy buffer sharing across runtimes, processes, and devices.
• Support the full range of quantization schemes used in modern inference.
• Be streamable (streaming format): a reader MUST be able to start processing tensor data without buffering the entire input, and a writer MUST be able to emit tensor data incrementally without buffering the entire output. Tensor descriptors always precede their data buffers; the format is self-delimiting; back-references are not permitted. File format writers operate in a single forward pass and append a footer index at the end.
• Extensible and evolvable: extension points are stable across 1.x; new named values go through the spec amendment process; backward and forward-additive compatibility is guaranteed within major version 1.x. See Versioning § Evolvability Contract.
• Use the standard numeric dtype vocabulary (shared by NumPy and the Python Array API Standard) for Tier 1 element types, enabling zero-copy interoperability without dtype translation. This is an interop convenience, not an Array API conformance claim. See Python Bindings for binding-level requirements.
• Serve as the storage foundation of an array database engine: the file format, tiled/blocked layout, Morton and Hilbert curve layouts, and footer index are designed to be compatible with sub-array queries, tile-skipping, and range-based retrieval. A concrete target use case is an embeddable SQL/MDA query engine (ISO 9075-15) backed by Hurray buffers, with zero-copy handoff from query results to an ML inference pipeline. Spec decisions that would foreclose chunk-based access, spatial locality, dimension-range indexing, or SQL/MDA interoperability MUST be evaluated against this use case before being adopted.

RFC 2119 Notice

The key words "MUST", "MUST NOT", "REQUIRED", "SHALL", "SHALL NOT", "SHOULD", "SHOULD NOT", "RECOMMENDED", "MAY", and "OPTIONAL" in these documents are to be interpreted as described in RFC 2119.

Versioning

This specification follows semantic versioning. A reader MUST reject a tensor descriptor whose major version field exceeds the reader's supported major version.

Table of Contents

Data Model — Hurray Format Specification

Status: Draft

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Scope

This section defines the logical data model of a Hurray tensor: the abstract description of its shape and element structure, independent of how it is stored in memory or encoded on the wire. All other sections of the specification operate on the model defined here.

Tensor

A tensor is a multi-dimensional array of elements, all of the same element type. A tensor is characterised by:

• a rank — the number of dimensions,
• a shape — the size of each dimension,
• an element type — defined in element-types.md.

The element type and shape fully determine the logical content of the tensor. How the elements are mapped to memory is defined separately by the tensor's memory layout (memory-layout.md).

Rank

The rank of a tensor is a non-negative integer that counts the number of dimensions. It is encoded as a uint32 in the tensor descriptor (metadata.md).

• A rank-0 tensor is a scalar: it contains exactly one element and has no shape array.
• A rank-1 tensor is a vector.
• A rank-2 tensor is a matrix.
• Higher ranks are used for batched inputs, activation maps, and weight tensors in neural networks.

The maximum rank is 64. A writer MUST NOT emit a descriptor with rank > 64. A reader MUST reject a descriptor with rank > 64. A conforming implementation MUST support tensors of rank 0 through 64 inclusive (see docs/adr/ADR-008-normative-rank-cap-64.md).

Note (non-normative): The uint32 rank field can encode values above 64, but those values are reserved and will be rejected. The cap matches PyTorch's MAX_DIMS = 64 and bounds the shape-array size at 512 bytes (64 × 8), enabling stack allocation on the descriptor-parsing hot path.

Shape

The shape of a tensor is an ordered sequence of rank dimension sizes, indexed from 0. Dimension 0 is the outermost (slowest-varying) dimension in row-major order; dimension rank − 1 is the innermost (fastest-varying).

Each dimension size is encoded as a uint64. A dimension size:

• MUST be greater than or equal to 0.
• A size of 0 denotes an empty dimension: the tensor contains no elements along that axis. The total element count of a tensor with any zero-size dimension is 0.
• The value 0xFFFFFFFFFFFFFFFF (UINT64_MAX) is the dynamic dimension sentinel — see § Dynamic Dimensions below.
• Any other value in the range [1, 0xFFFFFFFFFFFFFFFE] is a static dimension size.

For a scalar tensor (rank = 0), the shape is the empty sequence; no shape bytes are present in the descriptor.

Total Element Count

For a tensor with no dynamic dimensions, the total number of logical elements is:

element_count = product(shape[i] for i in 0 .. rank-1)

This product is 1 for a scalar. For a tensor with one or more zero-size dimensions, element_count = 0.

A reader MUST NOT compute element_count for a tensor that has any dynamic dimension without first resolving all dynamic dimensions to concrete values.

Dynamic Dimensions

A dimension whose size is 0xFFFFFFFFFFFFFFFF is dynamic: its concrete value is not known at descriptor-write time and MUST be supplied by the reader or the interchange protocol before the tensor's data buffer can be safely accessed.

Rules for dynamic dimensions:

1. A reader MUST NOT compute buffer sizes, strides, or element counts for a tensor containing a dynamic dimension without first resolving it to a concrete value.
2. A writer that sets a dimension to 0xFFFFFFFFFFFFFFFF MUST ensure the interchange channel provides a mechanism to communicate the resolved value before the data buffer is transferred (see interchange.md).
3. A quantization scheme that requires a statically known dimension size along its quantization axis (axis) MUST reject a descriptor whose shape[axis] is 0xFFFFFFFFFFFFFFFF. The specific constraint is documented per scheme in quantization/.
4. Shard descriptors MUST NOT use dynamic dimensions (see metadata.md § Shard Section).

Note (non-normative): Dynamic dimensions are intended for streaming and just-in-time dispatch scenarios, where a model producer does not know the batch size or sequence length at graph-compilation time. They are not a substitute for shape polymorphism at the type level.

Element Type

The element type of a tensor is identified by a uint8 type_tag defined in element-types.md. All elements in a tensor share the same type tag.

The element type describes the storage representation of each element in the data buffer. When a tensor is quantized (HAS_QUANTIZATION flag set), the storage type is an integer or float8 type and the quantization descriptor defines the mapping to real-valued elements. The storage type is always orthogonal to the quantization scheme: type_tag never encodes quantization semantics.

Scalar Tensors

A scalar tensor has rank = 0. Its shape is empty. It contains exactly one element. The data buffer size is sizeof(element_type) bytes (or a partial byte for sub-byte types; see element-types.md § Sub-Byte Packing).

A scalar tensor MUST NOT carry a strided, tiled, or sparse layout descriptor; only row-major (0x01) is permitted for scalars.

Note (non-normative): Scalar tensors arise as outputs of reduction operations (e.g., loss values) and as single-element configuration parameters.

Empty Tensors

A tensor is empty if element_count = 0. This occurs when at least one dimension has size 0. An empty tensor is valid; its data buffer has size 0 bytes.

An empty tensor MUST still carry a complete, valid descriptor with a correct element type, layout tag, buffer table, and any applicable quantization descriptor. A reader MUST accept an empty tensor without treating it as an error. A zero-length data buffer MAY be represented with a null pointer; the 64-byte alignment requirement does not apply to a zero-length buffer.

Note (non-normative): The decision to permit empty tensors is recorded in docs/adr/ADR-007-permit-empty-tensors.md. The primary motivation is round-trip fidelity with PyTorch, NumPy, JAX, Apache Arrow, and DLPack, all of which permit zero-size dimensions.

Relationship to Other Sections

• element-types.md defines the type_tag values, sub-byte packing rules, and the Tier 1 / Tier 2 type classification.
• memory-layout.md defines how the logical index space described by the shape is mapped to a linear buffer: strides, block shapes, sparse indices.
• metadata.md defines the binary encoding of rank, shape, and all other descriptor fields on the wire.
• quantization.md defines how storage elements are mapped to real-valued elements when HAS_QUANTIZATION is set.
• interchange.md defines how dynamic dimensions are resolved during runtime tensor exchange.

Open Questions

All open questions in this section are resolved. See docs/adr/ADR-007-permit-empty-tensors.md (OQ-2) and docs/adr/ADR-008-normative-rank-cap-64.md (OQ-1).

Element Types

Status: Draft

Scope

This section defines the complete numeric element type system for the Hurray tensor format, including type identifiers, bit widths, encoding rules, and sub-byte packing conventions.

Note (non-normative): Hurray is a strictly numeric tensor format. String types, structured types, and arbitrary user-defined element types are out of scope. The type system is designed to cover the full range of numeric precisions encountered in modern AI/ML inference, from float64 down to 2-bit quantized integers.

Normative Requirements

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Byte Order

All multi-byte fields in the Hurray format — including element data, descriptor fields, and metadata — MUST be encoded in little-endian byte order (least significant byte at the lowest address). The format does not define a big-endian encoding and contains no endianness indicator field.

Implementations running on big-endian host architectures MUST convert to little-endian on write and from little-endian on read. The wire format is always little-endian regardless of host byte order.

Note (non-normative): This matches the choice made by Apache Arrow, DLPack, and SafeTensors. All hardware platforms targeted by AI/ML inference workloads (x86, ARM, RISC-V, NVIDIA/AMD GPUs, Apple Silicon) are little-endian. A fixed byte order eliminates the need for byte-swap detection and preserves zero-copy semantics between any two conforming implementations.

Type Identifier Encoding

Every element type is identified by a type tag, encoded as a uint8 value in the binary tensor descriptor. The type tag space is partitioned as follows:

A conforming reader MUST reject a tensor descriptor containing a type tag of 0x00 or 0xFF.

A conforming reader MUST reject a tensor descriptor containing a type tag it does not recognize, unless the reader is operating in an explicitly configured permissive mode. In permissive mode, the reader MAY accept the descriptor but MUST NOT attempt to interpret the tensor data buffer.

Implementations MUST NOT assign semantics to type tags in the range 0x80 -- 0xEF; these are reserved for future versions of this specification.

Implementations MAY use type tags in the range 0xF0 -- 0xFE for private extensions. Tensors using private extension type tags MUST NOT be exchanged between independent implementations unless both parties have agreed on the semantics out of band.

Tier 1 -- Core Types

All conforming implementations MUST support every Tier 1 type. A conforming implementation MUST be able to read a tensor descriptor for any Tier 1 type and correctly interpret its metadata (shape, strides, buffer layout). Whether the implementation can perform computation on every Tier 1 type is outside the scope of this specification.

Floating-Point Types

float16 (0x01)

float16 is the IEEE 754 binary16 format: 1 sign bit, 5 exponent bits, 10 significand (mantissa) bits. Total width is 16 bits (2 bytes). The two bytes MUST be stored in little-endian order (least significant byte first).

All IEEE 754 binary16 bit patterns are valid, including positive and negative zero, infinities, and NaN values. Implementations MUST preserve the bit pattern exactly during zero-copy interchange; they MUST NOT canonicalize NaN payloads or flush subnormals to zero.

bfloat16 (0x02)

bfloat16 uses 1 sign bit, 8 exponent bits, and 7 significand (mantissa) bits. Total width is 16 bits (2 bytes). The two bytes MUST be stored in little-endian order.

Note (non-normative): bfloat16 has the same exponent range as float32 but with reduced mantissa precision. It is widely used in machine learning training and inference.

All 16-bit patterns are valid bfloat16 values, including infinities and NaN values. Implementations MUST preserve the bit pattern exactly during zero-copy interchange.

float32 (0x03)

float32 is the IEEE 754 binary32 format: 1 sign bit, 8 exponent bits, 23 significand bits. Total width is 32 bits (4 bytes). The four bytes MUST be stored in little-endian order.

All IEEE 754 binary32 bit patterns are valid. Implementations MUST preserve the bit pattern exactly during zero-copy interchange.

float64 (0x04)

float64 is the IEEE 754 binary64 format: 1 sign bit, 11 exponent bits, 52 significand bits. Total width is 64 bits (8 bytes). The eight bytes MUST be stored in little-endian order.

All IEEE 754 binary64 bit patterns are valid. Implementations MUST preserve the bit pattern exactly during zero-copy interchange.

Integer Types

All signed integer types use two's complement representation.

All multi-byte integer types MUST be stored in little-endian byte order.

int8 and uint8 occupy exactly one byte; byte order is not applicable.

All bit patterns within the specified width are valid for both signed and unsigned integer types. There are no trap representations.

Boolean Type

A bool element represents a logical true or false value. Each boolean occupies a single bit. Booleans are packed 8 per byte using LSB-first (least significant bit first) order.

Packing rule: logical element at index i within a group of 8 is stored in bit (i % 8) of byte floor(i / 8). Bit 0 is the least significant bit of the byte.

• A bit value of 0x1 represents true.
• A bit value of 0x0 represents false.

When the total number of boolean elements is not a multiple of 8, the remaining high-order bits in the final byte MUST be set to 0x0.

Example: A 1-D boolean tensor with shape [5] and values [true, false, true, true, false] is stored as a single byte: 0x0D (binary 00001101). Bits 5, 6, and 7 are padding and MUST be zero.

Note (non-normative): This packing convention is identical to the one used by Apache Arrow for boolean arrays.

Tier 2 -- Extended Types

Tier 2 types are OPTIONAL. Conforming implementations MAY support any subset of Tier 2 types, including none. Implementations that do not support a given Tier 2 type MUST still reject (or, in permissive mode, skip) descriptors using that type tag according to the rules in the Type Identifier Encoding section.

Float8 Variants

float8_e4m3 and float8_e5m2 follow the OCP (Open Compute Project) 8-bit Floating Point Specification (OFP8). Each occupies exactly 1 byte; byte order is not applicable.

float8_e4m3 (0x40): 1 sign bit, 4 exponent bits, 3 mantissa bits. Exponent bias is 7. NaN is represented by the bit patterns 0x7F and 0xFF (all exponent and mantissa bits set). There are no infinity representations.

float8_e5m2 (0x41): 1 sign bit, 5 exponent bits, 2 mantissa bits. Exponent bias is 15. This format supports infinities (0x7C and 0xFC) and NaN values (exponent all ones, mantissa non-zero).

float8_e8m0 (0x42): 8 exponent bits, no sign bit, no mantissa bits. This is a power-of-two scale factor format. The value is 2^(bits - 127) where bits is the unsigned 8-bit integer in the range [0x01, 0xFE]. The bit patterns 0x00 and 0xFF are reserved (NaN per OCP MX v1.0 § 5.6) and MUST NOT be used as data values. A reader encountering 0x00 or 0xFF in a float8_e8m0 buffer MUST treat the containing descriptor as invalid.

Note (non-normative): float8_e8m0 is primarily used as a scale factor type in microscaling (MX) quantization formats, not as a general-purpose element type.

Implementations MUST preserve all float8 bit patterns exactly during zero-copy interchange.

Sub-Byte Floating-Point Types

float4_e2m1 (0x43): 1 sign bit, 2 exponent bits, 1 mantissa bit. Exponent bias is 1. This is the MXFP4 format defined in the OCP Microscaling (MX) Specification. Each element occupies 4 bits.

float4_e2m1 uses the same LSB-first 4-bit packing as int4 (see § Sub-Byte Integer Types § 4-bit packing). When the total element count is odd, the high nibble of the final byte MUST be set to 0x0.

The complete value map, per OCP MX v1.0 § 5.2, is:

Maximum representable value: 1.5 * 2^2 = 6.0. There are no infinity or NaN representations; values outside [-6.0, 6.0] MUST be clamped by hardware.

Implementations MUST preserve all float4_e2m1 bit patterns exactly during zero-copy interchange.

Note (non-normative): float4_e2m1 has native Tensor Core support on NVIDIA Blackwell (B100/B200) and is used in production quantized LLM inference. It is typically paired with float8_e8m0 block scales under the MX quantization scheme (see quantization.md).

float6_e2m3 (0x44) and float6_e3m2 (0x45): Both are OCP MX Specification 6-bit floating-point formats. Each element occupies 6 bits.

float6_e2m3: 1 sign bit, 2 exponent bits, 3 mantissa bits. Exponent bias is 1. Maximum representable value: (1 + 7/8) * 2^(3-1) = 3.75. Zero is represented by all-zero exponent and mantissa (sign-preserving). No infinity or NaN representations; out-of-range values MUST be clamped. Subnormals: exponent all-zero, mantissa non-zero → value = (mantissa/8) * 2^(1-bias). Normative bit-pattern → value mapping per OCP MX v1.0 § 5.3.

float6_e3m2: 1 sign bit, 3 exponent bits, 2 mantissa bits. Exponent bias is 3. Maximum representable value: (1 + 3/4) * 2^(7-3) = 28.0. Zero is represented by all-zero exponent and mantissa (sign-preserving). No infinity or NaN representations; out-of-range values MUST be clamped. Subnormals: exponent all-zero, mantissa non-zero → value = (mantissa/4) * 2^(1-bias). Normative bit-pattern → value mapping per OCP MX v1.0 § 5.4.

6-bit packing

Four 6-bit elements are packed into 3 bytes using LSB-first order across byte boundaries. Given four elements at logical indices 4k, 4k+1, 4k+2, 4k+3 stored in bytes B0, B1, B2:

When the total number of elements is not a multiple of 4, the unused high-order bits in the final group of 3 bytes MUST be set to 0x0.

Buffer size in bytes for N elements: ceil(N * 6 / 8) = ceil(N / 4) * 3.

Worked example. Four 6-bit elements with bit patterns 0b000001, 0b000010, 0b000100, 0b001000 (logical indices 0, 1, 2, 3) pack into bytes B0, B1, B2 as follows.

Assembling each byte (MSB on the left):

B0 = 10_000001 = 0x81
B1 = 0100_0000 = 0x40
B2 = 001000_00 = 0x20

The 3-byte packed group on the wire is therefore 0x81 0x40 0x20.

Implementations MUST preserve all float6_e2m3 and float6_e3m2 bit patterns exactly during zero-copy interchange.

Note (non-normative): float6_e2m3 and float6_e3m2 are defined in the OCP MX specification and intended for use with MX block quantization (see quantization.md). Hardware adoption is currently limited compared to MXFP4.

Extended Floating-Point Types

float128 (0x46): 1 sign bit, 15 exponent bits, 112 significand bits. Exponent bias is 16383. Total width is 128 bits (16 bytes). The sixteen bytes MUST be stored in little-endian order.

All IEEE 754 binary128 bit patterns are valid, including positive and negative zero, infinities, and NaN values. Implementations MUST preserve the bit pattern exactly during zero-copy interchange.

Note (non-normative): float128 is rarely used in ML inference. It is included for high-precision scientific computing workloads (e.g., physics simulations, climate modelling) that may share tensor data with inference pipelines via the array database use case (Core Property 10).

Sub-Byte Integer Types

Sub-byte integer types are packed into bytes using LSB-first order, analogous to the boolean packing rule.

4-bit packing

Two 4-bit elements are packed per byte. The element at even logical index 2k occupies bits [3:0] (the low nibble) of byte k. The element at odd logical index 2k+1 occupies bits [7:4] (the high nibble) of byte k.

int4 values use two's complement representation within 4 bits. The valid bit patterns are 0x0 through 0xF, representing values -8 (0x8) through 7 (0x7).

uint4 values are unsigned. The valid bit patterns are 0x0 through 0xF, representing values 0 through 15.

When the total number of 4-bit elements is odd, the high nibble of the final byte MUST be set to 0x0.

Example: A 1-D uint4 tensor with shape [3] and values [5, 12, 3] is stored as two bytes. Byte 0: element 0 in low nibble, element 1 in high nibble = 0xC5. Byte 1: element 2 in low nibble, padding in high nibble = 0x03.

2-bit packing

Four 2-bit elements are packed per byte. The element at logical index 4k + j (where 0 <= j < 4) occupies bits [2j+1 : 2j] of byte k.

int2 values use two's complement representation within 2 bits. The valid bit patterns are 0b00 (0), 0b01 (1), 0b10 (-2), 0b11 (-1).

uint2 values are unsigned. The valid bit patterns are 0b00 (0), 0b01 (1), 0b10 (2), 0b11 (3).

When the total number of 2-bit elements is not a multiple of 4, the unused high-order bits in the final byte MUST be set to 0x0.

Example: A 1-D uint2 tensor with shape [3] and values [3, 1, 2] is stored as one byte. Element 0 in bits [1:0] = 0b11, element 1 in bits [3:2] = 0b01, element 2 in bits [5:4] = 0b10, padding in bits [7:6] = 0b00. Result: 0x27 (binary 00100111).

Complex Types

A complex64 element consists of two consecutive float32 values: the real part followed by the imaginary part. Total width is 64 bits (8 bytes). Each constituent float32 MUST be stored in little-endian order.

A complex128 element consists of two consecutive float64 values: the real part followed by the imaginary part. Total width is 128 bits (16 bytes). Each constituent float64 MUST be stored in little-endian order.

Note (non-normative): Complex types are included for signal processing and scientific computing workloads that may share tensor data with ML inference pipelines. They are not commonly needed in LLM inference.

Type Properties Summary

The following table summarizes all defined types and their properties.

The Alignment column specifies the minimum alignment requirement of the natural element type. This is the minimum alignment of an individual element within a contiguous buffer; the buffer itself has a separate, stricter alignment requirement (see buffer-protocol.md).

For sub-byte types, the alignment refers to the packed byte granularity: data for sub-byte types MUST start at a byte boundary within the buffer.

Tag 0x47 is reserved for future assignment by this specification. Implementations MUST NOT use tag 0x47. Private extensions MUST NOT assign tag 0x47.

Note (non-normative): Tag 0x47 does not appear in the table because it is intentionally reserved. It is held for a future Tier 2 type whose assignment will be defined in a later revision.

Note (non-normative): The alignment column for complex types reflects the natural alignment of the constituent floating-point element: complex64 lists alignment 4 (per float32 half) and complex128 lists alignment 8 (per float64 half). This matches the natural alignment of the constituent element type. Consumers loading a full complex value as a single 128-bit SIMD register must account for the fact that only the buffer-level alignment (64 bytes minimum, see buffer-protocol.md) guarantees register-width alignment, not the element-level alignment.

Buffer Size Calculation

For whole-byte types with bit width W >= 8, the minimum buffer size in bytes for a contiguous tensor with N total elements is:

buffer_size = N * (W / 8)

For sub-byte types with bit width B < 8, the general minimum buffer size in bytes for a contiguous tensor with N total elements is:

buffer_size = ceil(N * B / 8)

where ceil denotes the ceiling function (rounding up to the next integer).

For sub-byte bit widths B ∈ {1, 2, 4} (i.e. bool, int2/uint2, int4/uint4/float4_e2m1), 8 / B is an integer packing factor P and the formula reduces to the equivalent expression ceil(N / P). For B = 6 (float6_e2m3 and float6_e3m2), the general form ceil(N * 6 / 8) = ceil(N / 4) * 3 MUST be used, because four 6-bit elements pack into three bytes rather than into a whole number of elements per byte; see § Sub-Byte Floating-Point Types § 6-bit packing.

These formulas apply to contiguous (dense) layouts. For strided layouts, the buffer size depends on the strides; see memory-layout.md.

Interaction with Other Sections

• Quantization (quantization.md): Quantized tensors use an element type from this section as their storage type and attach a quantization descriptor that defines the dequantization mapping. The storage type tag in the tensor descriptor always refers to a type defined here.
• Memory Layout (memory-layout.md): Stride semantics for sub-byte types require special treatment. Strides are expressed in logical elements; the packing rules in this section define how logical element indices map to bit positions within the buffer.
• Metadata (metadata.md): The type tag defined here is stored as a uint8 field in the binary tensor descriptor.

Relationship to the standard numeric vocabulary

Note (non-normative): The Tier 1 numeric types — bool, int8, uint8, int16, uint16, int32, uint32, int64, uint64, float16, bfloat16, float32, float64, complex64, complex128 — use the standard numeric dtype vocabulary shared by NumPy and the Python Array API Standard (data-apis.org/array-api). This alignment is intentional and is purely an interop convenience: a Hurray tensor carrying a Tier 1 element type maps to the corresponding NumPy dtype without translation when handed to the ecosystem. It does not imply that a Hurray tensor is an Array API array (see Python Bindings and ADR-029).

Note (non-normative): Tier 2 types (float8_e4m3, float8_e5m2, float8_e8m0, float4_e2m1, float6_e2m3, float6_e3m2, float128, int4, uint4, int2, uint2) and quantized tensor types (see quantization.md) have no standard NumPy dtype and are exposed as hurray-namespaced dtype objects in the Python bindings. Requirements for the Python bindings are defined in Python Bindings.

Open Questions

[OQ-1]: Should float8_e4m3 follow the OCP OFP8 convention (no infinities, two NaN values) or the IEEE 754 draft for binary8 (which may differ)? Resolved: float8_e4m3 normatively follows OCP OFP8 (no infinities, two NaN bit patterns, exponent bias 7), matching production hardware (NVIDIA H100/H200, AMD MI300). If IEEE 754 binary8 diverges when finalized, a separate type tag will be assigned rather than redefining this one.

[OQ-2]: Should the specification define a float128 (IEEE 754 binary128) type? Resolved: float128 is added as a Tier 2 type with tag 0x46. Note: the originally proposed tag 0x05 falls in the Tier 1 range (0x01–0x3F) and was corrected to 0x46 (Tier 2 range 0x40–0x7F). Rationale: high-precision scientific computing workloads sharing tensor data with inference pipelines via the array database use case (Core Property 10).

[OQ-3]: Should the private extension range (0xF0 -- 0xFE) require implementations to include a type descriptor (name, bit width) in the tensor metadata so that readers can at least compute buffer sizes for unknown types? Resolved by ADR-001: Private extension type tags MUST carry an inline descriptor encoding at minimum the bit width, packing, and floating-point parameters (sign/exponent/mantissa widths, exponent bias, NaN/Inf flags). See docs/adr/ADR-001-private-extension-type-descriptors.md. The descriptor binary encoding will be defined in metadata.md.

[OQ-4]: Should float4_e2m1 (MXFP4) be added as a Tier 2 type? Resolved: float4_e2m1 is added as Tier 2 with tag 0x43. Packing follows the LSB-first int4 convention (two elements per byte). Rationale: native Tensor Core support on NVIDIA Blackwell and production use in quantized LLM inference.

[OQ-5]: Should float6_e2m3 and float6_e3m2 be added as Tier 2 types? Resolved: Both are added as Tier 2 with tags 0x44 and 0x45 respectively. Packing: 4 elements per 3 bytes, LSB-first across byte boundaries. Buffer size: ceil(N / 4) * 3 bytes.

Quantization — Hurray Format Specification

Status: Draft

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Scope

This section defines the binary encoding of the quantization descriptor: the opaque payload that appears inside the Quantization Section of a tensor descriptor (see metadata.md § Quantization Section). It specifies the set of supported quantization schemes, the wire format for each scheme's parameters, and the normative dequantization formula that a reader MUST apply to recover real-valued elements from the quantized storage buffer.

A Hurray quantized tensor has two parts:

1. A storage type — an integer or float8 type tag defined in element-types.md that describes how each element is encoded in the tensor's data buffer.
2. A quantization descriptor — the subject of this file, which describes how to map those storage values to real-valued elements.

Note (non-normative): The tensor descriptor's type_tag field always refers to the storage type. Hurray does not allocate separate type tags for quantized formats; a tensor is quantized if and only if the HAS_QUANTIZATION flag (bit 0 of flags) is set in the tensor descriptor. This separation keeps the type system orthogonal to the quantization scheme space.

Relationship to Other Sections

• metadata.md defines the Quantization Section framing. The section is present if and only if the HAS_QUANTIZATION flag (bit 0) is set in the tensor descriptor. It consists of a uint32 quantization_length prefix followed by quantization_length bytes of payload. This file defines the contents of those payload bytes.
• element-types.md defines the storage type tags (type_tag values) that may appear in a quantized tensor descriptor. Each scheme below lists the set of storage types it accepts.
• memory-layout.md defines the buffer table: an ordered list of buffer handles indexed from 0. Several schemes store their per-block scale and zero-point arrays in buffers other than buffer 0; those schemes reference those buffers by index.

Descriptor Header

Every quantization descriptor begins with a fixed 4-byte header.

All multi-byte fields in the quantization descriptor MUST be encoded in little-endian byte order.

Immediately following the 4-byte header, scheme-specific fields are encoded. The total length of the descriptor (header plus scheme-specific fields, plus any trailing padding chosen by the writer) MUST equal the quantization_length prefix defined in metadata.md.

A reader MUST dispatch on scheme_tag after reading the first byte. A reader that does not recognise scheme_tag MUST reject the tensor descriptor, unless operating in permissive mode. In permissive mode, the reader MAY skip past the descriptor using the quantization_length prefix but MUST NOT dereference the tensor data buffer.

A reader MUST reject a descriptor whose scheme_version exceeds the highest version defined in this specification for the given scheme_tag.

A reader MUST reject a descriptor with any reserved flags bit set.

A reader MUST NOT read beyond quantization_length bytes when parsing the descriptor.

Scheme Tag Space

A reader MUST reject a descriptor whose scheme_tag is 0x00 or 0xFF.

Implementations MUST NOT assign semantics to scheme tags in the range 0x60–0x7F; these are reserved for future nested/composite quantization schemes defined by this specification.

Tags in the range 0x80 – 0xEF MUST NOT be used by any implementation; they are reserved for future specification versions.

Tags in the range 0xF0 – 0xFE MAY be used by implementations for private schemes. Tensors using private scheme tags MUST NOT be exchanged between independent implementations unless both parties have agreed on the semantics out of band.

Note (non-normative): The three block-quantization schemes defined in this specification have different block_size lower bounds, reflecting their distinct design constraints:

• Per-block-affine (Tier 1, 0x03): block_size ≥ 2. Any positive block size with even packing makes algebraic sense; the lower bound exists only to forbid degenerate single-element blocks.
• NF4 (Tier 2, 0x04): block_size ≥ 8. Below 8 elements per block the 16-point NF4 information content provides no statistical benefit over a plain low-bit linear quantization, so the format is not meaningful at smaller block sizes.
• MXFP (Tier 2, 0x05): block_size ≥ 16. Tier 2 microscaling formats require hardware Tensor Core support; 16 is the minimum block size for which any production silicon (NVIDIA Blackwell) currently implements MX compute.

Assigned Scheme Tags

Partial Block Policy

The block-quantization schemes defined in this specification differ in whether they permit a partial trailing block — a final block along axis containing fewer than block_size valid logical elements. The active scheme's policy applies; readers implementing multiple schemes MUST apply the partial-block rule of the active scheme.

• Per-block-affine (0x03): a partial trailing block is permitted. The storage buffer MUST still allocate space for a full block of block_size elements; the unused trailing storage elements MUST be zero bytes and a reader MUST ignore them when dequantizing. The scale (and zero-point) arrays MUST contain one entry per block, including the partial final block. See quantization/per-block-affine.md § Padding for the normative encoding.
• NF4 (0x04): a partial trailing block is permitted under the same rules as per-block-affine, since NF4 derives its block index computation from per-block-affine. See quantization/per-block-affine.md § Padding (NF4 uses identical padding rules).
• MXFP (0x05): a partial trailing block is NOT permitted. shape[axis] MUST be a positive multiple of block_size. This restriction is normative under the OCP MX v1.0 specification: microscaling hardware assumes exact block alignment along the quantized axis. A reader MUST reject an MXFP descriptor whose shape[axis] is not a positive multiple of block_size. See quantization/mxfp.md § Validity Constraints.

A reader that supports multiple schemes MUST dispatch on scheme_tag and apply the corresponding partial-block rule. A reader MUST NOT cross-apply the permissive per-block-affine padding rule to MXFP, nor reject a partial final block of a per-block-affine or NF4 tensor on the basis of the MXFP exact- multiple constraint.

Zero-Point Convention

The affine schemes (0x01, 0x02, 0x03) dequantize by subtracting zero_point from the storage value. The stored zero point is the value subtracted: Hurray defines no implicit bias, and the descriptor carries no field selecting an alternative convention.

A writer converting a tensor from an external quantization toolchain MUST normalize that toolchain's zero-point values to this convention before placing them in a Hurray buffer.

Note (non-normative): Some external toolchains store a biased zero point and remove the bias in the loader rather than recording it in the file. GPTQ is the widely deployed case: its packed zero-point plane has conventionally held zero_point - 1. A descriptor built from un-normalized values is structurally valid — it satisfies every validity constraint of its scheme, and a reader has no way to detect the discrepancy — but dequantizes every affected element off by one scale step. Hurray records the zero point, not the convention that produced it, so this normalization is the writer's responsibility and cannot be recovered downstream.

Buffer Table Placement Rules

Every scheme that references a buffer (per-channel, per-block, NF4, MXFP) adds entries to the tensor descriptor's buffer table beyond the buffers used by the layout itself. The following rules apply across all quantization schemes:

1. The tensor data buffer always occupies a buffer table index determined by the layout (typically 0 for dense layouts). Quantization-parameter buffers MUST occupy distinct indices.
2. A quantization-parameter buffer MUST NOT be shared with the tensor data buffer. A reader MUST reject a descriptor that violates this rule.
3. Two quantization-parameter buffers (e.g., scale and zero-point) MAY reside in the same buffer table entry if and only if the entry's byte_size accommodates both arrays laid out end-to-end, and both arrays start at byte offset 0. Since each parameter descriptor field specifies its own *_buffer_index, this case is expressed by writing the same index into both fields; readers MUST then interpret the buffer as the concatenation [scales | zero_points]. Writers SHOULD prefer distinct buffers for clarity; sharing is permitted only as an optimization.
4. A quantization-parameter buffer's device_tag (see metadata.md buffer handle format) MUST match the tensor data buffer's device_tag. A reader MUST reject a descriptor that violates this rule.

Note (non-normative): The device-colocation rule ensures that quantized tensor kernels can dereference both the data and the quantization parameters without triggering cross-device transfers. A writer that needs to materialize quantization parameters on a different device must emit a separate tensor.

Extension Schemes (0xF0 – 0xFE)

Implementations MAY define private quantization schemes using scheme tags in 0xF0 – 0xFE. The binary encoding of an extension scheme descriptor is unconstrained beyond the fixed 4-byte header; the writer and reader MUST agree on the payload format out of band.

Note (non-normative): Extension schemes are the mechanism for implementation-specific or experimental quantization formats (e.g., GPTQ, AWQ group quantization with per-group permutations, or hardware-vendor- specific packings). Schemes that prove broadly useful are candidates for assignment in the Tier 1 or Tier 2 range through a specification revision.

What sends these formats here is the per-group permutation — GPTQ activation-order grouping reorders the quantized axis, which per-block affine cannot express. The permutation-free subset is a different case: a group-quantized tensor whose groups are contiguous along one axis is representable under per-block affine (0x03) by choosing the appropriate storage type, block size, and flags, and does not need an extension scheme. See § Zero-Point Convention for the normalization such a conversion requires.

Version Compatibility

Adding a new scheme tag MUST be accompanied by a minor version increment of the overall format (version_minor in the fixed header; see metadata.md). A reader at an earlier minor version will encounter an unrecognised scheme_tag and reject the descriptor per the rules above, which is the intended behaviour.

Adding a new field to an existing scheme (or repurposing a reserved byte) MUST be accompanied by a scheme_version increment for that scheme. A reader MUST compare scheme_version against the highest version it supports for the given scheme_tag and reject the descriptor if the version is newer than supported.

Removing a scheme or changing the dequantization formula for an existing (scheme_tag, scheme_version) pair is a backward-incompatible change and MUST be accompanied by a major version increment.

Open Questions

Double / nested quantization (deferred). This version does not define a normative encoding for double quantization (quantizing the scale buffer itself, as in bitsandbytes QLoRA "nested" quantization). The scheme-tag range 0x60–0x7F is reserved for a future nested/composite scheme (see § Scheme Tag Space). A conforming implementation MAY express double quantization today by representing the scale tensor as a separate quantized tensor using existing scheme tags, with the relationship conveyed by application-layer convention.

Note (non-normative): A recursive or two-level scale descriptor is deferred until the base quantization encoding has been validated through implementation; double quantization is primarily a weight-storage optimization rather than a runtime interchange primitive. The intended future occupant of 0x60–0x7F is a nested-scale scheme compatible with bitsandbytes-style double quantization (NF4 data with quantized float8 scales and a float32 super-scale), to be specified in a future revision once implementation experience is available.

[OQ-2]: Should MXFP block_size be fixed at 32 or parameterised? Resolved: block_size is parameterised. The field already exists in the binary encoding; the constraint is: MUST be a power of two in [16, 2048]. The lower bound of 16 excludes values with no hardware Tensor Core support. The OCP MX v1.0 canonical value of 32 is documented as the default. Rationale: avoids scheme tag proliferation for what is effectively one scheme with a size variation; future OCP revisions and hardware variants can use different block sizes under the same scheme tag.

[OQ-3]: Should per-channel affine support lower-precision scale types? Resolved: Per-channel scales remain locked to float32 (scale_type_tag = 0x03). Rationale: storage saving is negligible (~8–16 KB per layer) while accuracy cost is real for precision-critical per-channel weight quantization. A scale_type_tag field (+ 3 reserved bytes) has been added to the per-channel affine binary encoding at offset 16 to allow future relaxation without a wire-format break.

[OQ-4]: Should the descriptor include an explicit num_blocks field? Resolved: num_blocks remains derived: num_blocks = shape[axis] / block_size. Rationale: the value is fully determined by fields already present in the descriptor; an explicit field would be redundant and introduce a new mismatch failure mode.

[OQ-5]: Should the NF4 lookup table be duplicated into a normative reference appendix? Resolved: The table remains inline in quantization/nf4.md. Rationale: duplication would risk the two copies diverging; a single source of truth is safer. Spec audits (spec-checker) are the guard against accidental mutation.

Per-Tensor Affine Quantization — Hurray Format Specification

Scheme tag: 0x01 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

A single scale and zero_point pair applies to every element of the tensor. This scheme covers both asymmetric quantization (arbitrary zero_point) and symmetric quantization (zero_point = 0).

The scale and zero-point are stored inline in the quantization descriptor; no additional buffer table entries are required beyond the tensor data buffer.

Binary Encoding

Total descriptor length: 16 bytes.

All multi-byte fields MUST be encoded in little-endian byte order.

Dequantization Formula

For each storage element q:

x_real = scale * (q - zero_point)

The subtraction q - zero_point is performed in signed 32-bit integer arithmetic. The result is then cast to float32 and multiplied by scale in float32 arithmetic. The real-valued element type produced by dequantization is float32. A consumer MAY further convert to float64 or to a lower-precision float type; such conversion is out of scope for this specification.

zero_point is subtracted exactly as stored. See quantization.md § Zero-Point Convention for the normalization a writer MUST apply when converting from a toolchain that stores a biased zero point.

Validity Constraints

• scale MUST NOT be zero, NaN, or infinity. A reader MUST reject a descriptor that violates this constraint.
• zero_point MUST lie within the representable range of the storage type. For example, for a uint8 storage type, zero_point MUST be in [0, 255]. A reader MUST reject a descriptor that violates this constraint.
• The _reserved bytes MUST be 0x00. A reader MUST reject a descriptor that violates this constraint.

Valid Storage Types

The storage type (type_tag in the tensor descriptor) MUST be one of:

• int8 (0x10), uint8 (0x11)
• int16 (0x12), uint16 (0x13)
• int32 (0x14), uint32 (0x15)
• int4 (0x48), uint4 (0x49)
• int2 (0x4A), uint2 (0x4B)

A reader MUST reject a descriptor whose storage type is not in this list.

Per-Channel Affine Quantization — Hurray Format Specification

Scheme tag: 0x02 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

One scale and zero_point pair per slice along a specified axis. The scale and zero-point arrays are stored in separate buffers listed in the tensor descriptor's buffer table.

Note (non-normative): Because dense-layout descriptors require buffer_count = 0x01, a per-channel-affine quantized dense tensor requires buffer_count to be at least 2 (symmetric case) or 3 (asymmetric case). This is the mechanism by which quantization schemes extend the buffer table beyond the dense-layout minimum.

Binary Encoding

Total descriptor length: 20 bytes.

All multi-byte fields MUST be encoded in little-endian byte order.

Flags bits:

Referenced Buffers

The scale buffer MUST contain exactly shape[axis] consecutive float32 values in little-endian byte order, starting at byte offset 0 within the referenced buffer. Its byte size MUST be exactly shape[axis] * 4.

The zero_point buffer — present only if the SYMMETRIC flag is not set — MUST contain exactly shape[axis] consecutive int32 values in little-endian byte order, starting at byte offset 0 within the referenced buffer. Its byte size MUST be exactly shape[axis] * 4.

Note (non-normative): The zero_point buffer uses int32 regardless of the storage type's bit width (e.g. int4 or int2). This simplifies alignment and avoids sub-byte zero-point packing. The Validity Constraints section enforces that zero-point values lie within the representable range of the storage type, so the wider container does not introduce additional degrees of freedom on the wire.

A reader MUST reject a descriptor whose scale_buffer_index or (when the SYMMETRIC flag is not set) zero_point_buffer_index is greater than or equal to buffer_count in the buffer table.

A reader MUST reject a descriptor whose scale_buffer_index or zero_point_buffer_index equals the buffer index used by the layout for tensor data (typically 0 for dense layouts).

Dequantization Formula

For a storage element q at logical index [i_0, i_1, ..., i_{rank-1}]:

c = i_axis
x_real = scale[c] * (q - zero_point[c])

If the SYMMETRIC flag is set, zero_point[c] is treated as 0 for all c.

zero_point[c] is subtracted exactly as stored. See quantization.md § Zero-Point Convention for the normalization a writer MUST apply when converting from a toolchain that stores a biased zero point.

Validity Constraints

• axis MUST satisfy axis < rank.
• shape[axis] MUST NOT equal 0xFFFFFFFFFFFFFFFF (the dynamic dimension sentinel): per-channel quantization requires a statically known channel count.
• scale_type_tag MUST be 0x03 when scheme_version = 0x01. Future scheme versions MAY define additional values. A reader MUST reject a scheme_version = 0x01 descriptor with any other value.
• The _reserved bytes MUST be 0x00 when scheme_version = 0x01. A reader MUST reject a scheme_version = 0x01 descriptor with any non-zero reserved byte.
• Every element of the scale array MUST be a finite, non-zero float32 value.
• Every element of the zero_point array (when present) MUST lie within the representable range of the storage type.

A reader MAY defer the per-element validity check on the scale and zero-point arrays to the first dequantization attempt, but MUST perform the axis and shape checks before accepting the descriptor.

Valid Storage Types

The storage type (type_tag in the tensor descriptor) MUST be one of:

• int8 (0x10), uint8 (0x11)
• int16 (0x12), uint16 (0x13)
• int32 (0x14), uint32 (0x15)
• int4 (0x48), uint4 (0x49)
• int2 (0x4A), uint2 (0x4B)

A reader MUST reject a descriptor whose storage type is not in this list.

Worked Example

A rank-2 int8-stored weight tensor with shape [768, 1024], per-channel affine quantization along axis 0 (asymmetric), no statistics or shard sections. The tensor descriptor's buffer table carries three buffers:

• Buffer 0 — tensor data, 768 * 1024 = 786432 bytes, int8 storage.
• Buffer 1 — scale array, 768 * 4 = 3072 bytes, float32.
• Buffer 2 — zero-point array, 768 * 4 = 3072 bytes, int32.

Quantization descriptor bytes (20 total):

Offset  Value (hex)                   Field
------  ----------------------------  -----
0       02                            scheme_tag = 0x02 (per-channel affine)
1       01                            scheme_version = 1
2       00 00                         flags = 0x0000 (asymmetric)
4       00 00 00 00                   axis = 0
8       01 00 00 00                   scale_buffer_index = 1
12      02 00 00 00                   zero_point_buffer_index = 2
16      03                            scale_type_tag = 0x03 (float32)
17      00 00 00                      _reserved = 0x00

The quantization_length prefix in the tensor descriptor's Quantization Section would be 0x00000014 (20).

Dequantization of element q at logical position [c, k]:

x_real = scale[c] * (q - zero_point[c])

Per-Block Affine Quantization — Hurray Format Specification

Scheme tag: 0x03 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

The tensor is divided into fixed-size, contiguous blocks along a specified axis. Each block carries its own scale (and optionally zero_point). This scheme covers the single-level GGUF block-quantized formats — those carrying one scale, and optionally one zero point, per block (e.g., Q8_0, Q4_0, Q4_1) — at the descriptor level; those layouts are representable by choosing the appropriate storage type, block size, and flags.

The GGUF K-quant family (Q2_K–Q6_K) is not expressible under this scheme. Those formats apply a second, super-block scaling level: the per-block scales are themselves quantized and scaled by a shared factor, which this scheme's single scale array cannot represent. See quantization.md § Open Questions, where nested and two-level scale descriptors are deferred to the reserved 0x60–0x7F tag range.

Binary Encoding

Total descriptor length: 24 bytes.

All multi-byte fields MUST be encoded in little-endian byte order.

Flags bits:

Block Layout

Let S = shape[axis] (resolved; MUST NOT be the dynamic dimension sentinel) and K = block_size. The number of blocks along axis is:

num_blocks_per_axis = ceil(S / K)

The total number of blocks across the whole tensor is:

num_blocks = num_blocks_per_axis * product(shape[j] for j != axis)

Block index b at a tensor position [i_0, ..., i_{rank-1}] is computed as follows. Let outer be the linear index formed from all dimensions except axis using row-major order over those dimensions. Then:

b = outer * num_blocks_per_axis + floor(i_axis / K)

Note (non-normative): This mapping preserves the tensor's row-major traversal order along non-quantized dimensions, which matches GGUF's linear layout for 2-D weight matrices. Writers targeting column-major traversal SHOULD use a column-major tensor layout rather than altering the block mapping.

Padding

If S is not a multiple of K, the final block along axis for each outer position contains only S mod K valid elements. The storage buffer MUST still allocate space for a full block of K elements; the unused trailing elements within the final block MUST be set to 0x00 bytes by the writer. A reader MUST ignore these padding elements when dequantizing.

The scale (and zero-point) arrays MUST contain one entry per block, including the partial final block.

Referenced Buffers

The scale buffer MUST contain exactly num_blocks consecutive values of scale_type_tag in little-endian byte order, starting at byte offset 0. Its byte size MUST be exactly num_blocks * sizeof(scale_type_tag).

The zero_point buffer — present only if the SYMMETRIC flag is not set — MUST contain exactly num_blocks consecutive int32 values in little-endian byte order, starting at byte offset 0. Its byte size MUST be exactly num_blocks * 4.

A reader MUST reject a descriptor whose scale_buffer_index or (when applicable) zero_point_buffer_index is greater than or equal to buffer_count.

Dequantization Formula

Let b be the block index for a storage element q at logical position [i_0, ..., i_{rank-1}], computed as above. Let s = scale[b] and, if the SYMMETRIC flag is set, z = 0, else z = zero_point[b].

x_real = s * (q - z)

The multiplication is performed in float32 arithmetic. If scale_type_tag is float16 or bfloat16, s MUST be widened to float32 losslessly before the multiplication.

z is subtracted exactly as stored. See quantization.md § Zero-Point Convention for the normalization a writer MUST apply when converting from a toolchain that stores a biased zero point — the case that arises when a group-quantized GPTQ or AWQ tensor without activation-order permutation is mapped onto this scheme.

Validity Constraints

• axis MUST satisfy axis < rank.
• block_size MUST be a power of two and MUST be greater than or equal to 2.
• When shape[axis] is greater than 0, block_size MUST be less than or equal to shape[axis]. A reader MUST reject a descriptor whose block_size exceeds a non-zero shape[axis].
• When shape[axis] equals 0 (an empty quantization axis, per ADR-007), the upper-bound check is waived. The block_size field MUST still be a power of two greater than or equal to 2, but its value has no effect on buffer sizing: num_blocks evaluates to 0, and the scale and zero-point buffers MUST have byte size 0 (their pointers MAY be null per buffer-protocol.md).
• shape[axis] MUST NOT equal 0xFFFFFFFFFFFFFFFF.
• Every scale value MUST be a finite, non-zero value.
• scale_type_tag MUST be one of 0x01, 0x02, 0x03.

Note (non-normative): Permitting block_size to exceed shape[axis] when the axis is empty preserves the producer's declared quantization granularity across shape changes (for example, a filter that selects zero rows from an otherwise per-block-quantized weight tensor). The descriptor remains structurally valid and round-trippable; no blocks are materialized.

Note (non-normative): The case block_size = shape[axis] — a single block covering the entire quantized axis — is intentionally permitted. It is a degenerate but valid configuration that produces semantics equivalent to per-tensor affine quantization along that axis (one shared scale/zero_point pair per outer position). Readers MUST handle it identically to any other block_size value; there is no separate code path.

Valid Storage Types

The storage type (type_tag in the tensor descriptor) MUST be one of:

• int8 (0x10), uint8 (0x11)
• int4 (0x48), uint4 (0x49)
• int2 (0x4A), uint2 (0x4B)

A reader MUST reject a descriptor whose storage type is not in this list.

Note (non-normative): Wider integer storage types (int16 and above) are not permitted for per-block affine because block quantization is only beneficial at low bit widths. A writer wishing to apply per-block scaling to wider storage types should use per-channel affine (0x02) instead.

NF4 (NormalFloat4) Quantization — Hurray Format Specification

Scheme tag: 0x04 | Tier: 2

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

A non-linear 4-bit quantization scheme introduced by the QLoRA paper. Each storage code in [0, 15] decodes to one of 16 fixed real-valued levels chosen to be information-theoretically optimal for weights drawn from a standard normal distribution. Each block carries a single absmax scale; there is no per-element zero point.

The block index computation follows the same rule as Per-Block Affine (see quantization/per-block-affine.md § Block Layout).

Binary Encoding

Total descriptor length: 16 bytes.

All multi-byte fields MUST be encoded in little-endian byte order.

Lookup Table

The 16 NF4 levels are fixed by this specification and MUST NOT be altered by readers or writers. Indexed by the unsigned 4-bit storage code q:

The table values are given here as float32 decimal expansions of the exact levels from the QLoRA reference implementation. Implementations MUST use these float32 values verbatim; deriving the table from first principles at runtime is NOT RECOMMENDED and MUST match these values bit-for-bit if attempted.

Referenced Buffer

The scale buffer MUST contain exactly num_blocks consecutive float32 absmax values in little-endian byte order, starting at byte offset 0 within the referenced buffer. num_blocks is computed identically to the Per-Block Affine scheme. Its byte size MUST be exactly num_blocks * 4.

Dequantization Formula

Let b be the block index for a storage element q at logical position [i_0, ..., i_{rank-1}] (computed as in Per-Block Affine). Let s = scale[b].

x_real = s * nf4[q]

The multiplication is performed in float32 arithmetic.

Validity Constraints

• axis MUST satisfy axis < rank.
• block_size MUST be a power of two and MUST be greater than or equal to 8.
• When shape[axis] is greater than 0, block_size MUST be less than or equal to shape[axis]. A reader MUST reject a descriptor whose block_size exceeds a non-zero shape[axis].
• When shape[axis] equals 0 (an empty quantization axis, per ADR-007), the upper-bound check is waived. The block_size field MUST still be a power of two greater than or equal to 8, but its value has no effect on buffer sizing: num_blocks evaluates to 0, and the absmax scale buffer MUST have byte size 0.
• shape[axis] MUST NOT equal 0xFFFFFFFFFFFFFFFF.
• Every scale value MUST be a finite, non-negative float32.
• scale_buffer_index MUST be a valid index into the buffer table and MUST NOT refer to the tensor data buffer.

Note (non-normative): The minimum block size of 8 reflects the information-theoretic minimum for NF4: below 8 elements per block, the 16-point quantization grid provides no statistical benefit over a lower-bit format.

Valid Storage Types

The storage type (type_tag in the tensor descriptor) MUST be uint4 (0x49).

A reader MUST reject an NF4 descriptor whose storage type is not uint4.

Note (non-normative): NF4 storage codes are conceptually unsigned (they index into a signed-valued lookup table); uint4 is the correct storage type. The packing order follows the standard uint4 rule from element-types.md: element 2k in the low nibble of byte k, element 2k+1 in the high nibble.

MXFP (OCP Microscaling) Quantization — Hurray Format Specification

Scheme tag: 0x05 | Tier: 2

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

A block quantization format standardized by the Open Compute Project Microscaling specification (OCP MX v1.0). A contiguous block of elements along a chosen axis shares a single float8_e8m0 exponent-only scale. Each element within the block is stored as one of several supported narrow numeric types. This is the format used by NVIDIA Blackwell Tensor Cores for MXFP8/MXFP6/MXFP4 compute.

The block size is a descriptor field (block_size). The OCP MX v1.0 canonical value is 32; other power-of-two values are permitted by this specification to accommodate future OCP revisions and hardware variants.

The block index computation follows the same rule as Per-Block Affine (see quantization/per-block-affine.md § Block Layout), with the descriptor-specified block_size.

Binary Encoding

Total descriptor length: 16 bytes.

All multi-byte fields MUST be encoded in little-endian byte order.

Referenced Buffer

The scale buffer MUST contain exactly num_blocks consecutive float8_e8m0 values, one byte each, starting at byte offset 0 within the referenced buffer. Its byte size MUST be exactly num_blocks bytes. num_blocks is computed identically to the Per-Block Affine scheme using the descriptor-specified block_size, with the additional MXFP constraint that shape[axis] is a positive multiple of block_size (no partial trailing block — see Validity Constraints below). The full count across the whole tensor is:

num_blocks = (shape[axis] / block_size) * product(shape[j] for j ≠ axis)

(exact division — shape[axis] MUST be a positive multiple of block_size; see Validity Constraints). See quantization/per-block-affine.md § Block Layout for the derivation; MXFP differs only in disallowing partial trailing blocks.

The bit patterns 0x00 and 0xFF in any scale byte are reserved (NaN per OCP MX v1.0 § 5.6; see element-types.md § float8_e8m0) and MUST NOT appear in the scale buffer. A reader encountering 0x00 or 0xFF in the scale buffer MUST treat the descriptor as invalid.

Dequantization Formula

Let b be the block index for a storage element at logical position [i_0, ..., i_{rank-1}]. Let e = scale[b] be the float8_e8m0 byte. The shared exponent scale is:

s = 2^(e - 127)

If the storage type is a float type (float8_e4m3, float8_e5m2, float4_e2m1, float6_e2m3, float6_e3m2):

x_real = s * float_value_of(q)

where float_value_of(q) is the real number represented by the storage element q interpreted according to element-types.md and the OCP MX v1.0 value maps for the respective type.

If the storage type is an integer type (int8, int4):

x_real = s * int_value_of(q)

where int_value_of(q) is the signed integer value of q. This MXFP integer form is equivalent to per-block symmetric affine quantization with a power-of-two scale and no zero point.

All arithmetic is performed in float32 (or higher) precision; the shared exponent s is itself exactly representable in float32 for every valid float8_e8m0 bit pattern except 0xFF (which is prohibited above).

Validity Constraints

• axis MUST satisfy axis < rank.
• block_size MUST be a power of two in the range [16, 2048]. A reader MUST reject a descriptor whose block_size is not a power of two, is less than 16, or exceeds 2048. Values below 16 have no hardware Tensor Core support and are not valid under any OCP MX revision.
• shape[axis] MUST NOT equal 0xFFFFFFFFFFFFFFFF.
• shape[axis] MUST be a positive multiple of block_size. Unlike Per-Block Affine and NF4, MXFP does NOT permit partial trailing blocks; all blocks MUST be full. A reader MUST reject a descriptor whose shape[axis] is not a positive multiple of block_size.
• scale_buffer_index MUST be a valid index into the buffer table and MUST NOT refer to the tensor data buffer.

Valid Storage Types

The storage type (type_tag in the tensor descriptor) MUST be one of:

• float8_e4m3 (0x40) — MXFP8
• float8_e5m2 (0x41) — MXFP8
• float4_e2m1 (0x43) — MXFP4
• float6_e2m3 (0x44) — MXFP6
• float6_e3m2 (0x45) — MXFP6
• int8 (0x10) — MXINT8
• int4 (0x48) — MXINT4 / MXFP4 integer-valued surrogate

A reader MUST reject an MXFP descriptor whose storage type is not in this list.

Memory Layout -- Hurray Format Specification

Status: Draft

Scope

This section defines how tensor elements are arranged in memory. It specifies the addressing model that maps a tensor's logical index space to byte positions within a data buffer. Hurray supports a range of memory layouts — from simple contiguous arrangements to tiled, space-filling-curve, sparse, and composite (virtual) layouts — to accommodate the diverse access patterns required by modern AI/ML inference pipelines.

Note (non-normative): The unifying mathematical concept behind all Hurray dense layouts is the subpaving: a finite collection of non-overlapping boxes (rectangular regions) that together tile a tensor's index space. A contiguous row-major tensor is a trivial subpaving (one box). A tiled tensor is a regular subpaving. See Subpaving (Wikipedia).

Normative Requirements

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Layout Taxonomy

Every tensor descriptor MUST include a layout tag (uint8) that identifies the memory layout of the tensor's data. The tag space is partitioned as follows:

A conforming reader MUST reject a tensor descriptor containing a layout tag of 0x00 or 0xFF.

A conforming reader MUST reject a tensor descriptor containing a layout tag it does not recognise, unless operating in permissive mode. In permissive mode, the reader MAY accept the descriptor but MUST NOT dereference or interpret the tensor data buffer.

Named Layout Tags

The Type column classifies each layout's addressing model:

• Dense — every logical element exists and maps to a physical position by an affine stride formula.
• Sparse — only stored (non-zero) elements are materialised; unstored coordinates are implicitly zero.
• Indirect — every logical element exists (no implicit zeros), but the mapping from a logical index to a physical position is non-affine and resolved through an index structure (e.g. a block table) rather than an affine stride formula.
• Virtual — the descriptor owns no data buffers; it presents a logical view assembled from a set of member tensors (the head of a composite; see layouts/composite.md).

Note (non-normative): Rows marked "(reserved — planned)" name a tag that is earmarked for a layout whose spec section does not yet exist. A reader treats such a tag exactly as it treats any unrecognised tag: it MUST reject a descriptor bearing that tag unless operating in permissive mode, in which case it MUST NOT dereference the tensor data buffer (see the unrecognised-tag rule above). The reservation only records intent so the tag is not reassigned before the layout is specified.

Writers choose the layout. Hurray imposes no requirement on which layout a writer selects; any layout from the table above (or from the extension range, by prior agreement) is valid.

Common Fields

All layouts share the following fields in the tensor descriptor. These fields are defined once here; individual layout files specify additional layout-specific fields.

Rank and Shape

• rank (uint32): the number of dimensions. 0 denotes a scalar tensor.
• shape (uint64[rank]): the size of each dimension. Each value MUST be greater than or equal to 0. A size of 0 indicates an empty tensor.

Note (non-normative): Zero-size dimensions are valid for placeholder tensors or empty batches. A tensor with shape [3, 0, 5] has zero total elements.

The value 0xFFFFFFFFFFFFFFFF (UINT64_MAX) is the dynamic dimension sentinel: the dimension's size is not statically known and MUST be resolved before use.

byte_offset

• byte_offset (uint64): offset in bytes from the start of buffer 0 to the element at logical index [0, 0, ..., 0]. MUST be ≤ the buffer's byte size.

For sub-byte types (bool, int4, uint4, int2, uint2), byte_offset MUST point to a byte boundary.

For sparse layouts (COO, CSR, CSC, CSF, and future sparse tags) and indirect layouts (block-paged, and future indirect tags), the concept of a "first element at a fixed offset" does not apply: the first logical element is located through an index structure, not at a fixed offset. For these tensors, byte_offset MUST be set to 0x0000000000000000.

For the virtual layout (composite head, tag 0x0B), there is no data buffer at all; byte_offset MUST be set to 0x0000000000000000. See layouts/composite.md.

Element Address Computation

Whole-Byte Types

For element types with bit width ≥ 8, the byte address of the element at linear offset offset (as computed by the layout-specific addressing formula) is:

byte_address = byte_offset + offset * (bit_width / 8)

Sub-Byte Types

For sub-byte types (bool, int4, uint4, int2, uint2), strides are expressed in logical elements. Given a linear element offset offset:

• Packing factor P: elements per byte (8 for bool, 2 for 4-bit, 4 for 2-bit).
• Bit width B: bits per element (1, 4, or 2).
• Byte index: byte_offset + floor(offset / P).
• Bit position within byte: (offset mod P) * B (counting from LSB).

For strided layouts with sub-byte types, strides are in logical elements; the implementation MUST compute the linear element offset using the strides, then apply the packing formula above.

Note (non-normative): Sub-byte types are almost always used with contiguous layouts. Strided sub-byte tensors are supported for completeness but require bit-level manipulation per element. Writers SHOULD prefer contiguous layouts for sub-byte data.

Note (non-normative): The 6-bit types (float6_e2m3, float6_e3m2) pack 4 elements per 3 bytes (see element-types.md § Buffer Size Calculation). However, because their bit width (6) does not divide evenly into 8 bits, the standard sub-byte bit-addressing model (where bit position within a byte is well-defined) does not apply cleanly. Elements are always addressed at the group level (3 bytes per group of 4 elements); individual element extraction within a group is defined by the bit layout in element-types.md § Encoding but is not generalised here.

Accordingly, 6-bit types (float6_e2m3, float6_e3m2) MUST be accessed at the group level — one group is 4 elements packed into 3 bytes — and an implementation MUST NOT apply the single-element byte/bit-addressing model of § Sub-Byte Types to them. Single-element extraction within a group is defined normatively by element-types.md § Encoding and is intentionally not generalised into a layout addressing formula here; that encoding is sufficient for implementors.

Alignment

• The data buffer MUST be aligned to at least 64 bytes regardless of layout or element type.
• For tiled layouts, each tile's data SHOULD start at a naturally aligned boundary. Writers SHOULD insert inter-tile padding when needed; tile stride values MUST account for any padding.
• For sub-byte types, byte_offset MUST be byte-aligned.
• Page-aligned buffers (typically 4096 bytes) SHOULD be used when the tensor is shared across processes or with GPU devices. See buffer-protocol.md.

Buffer Table

Every tensor descriptor contains a buffer table: an ordered list of buffer handles. The buffer table is encoded as a uint8 count followed by that many buffer handle entries, as defined in metadata.md.

For dense layouts (tags 0x01–0x05, 0x40), the buffer table MUST contain at least one entry. Non-quantized dense tensors MUST have exactly buffer_count = 0x01. Quantized dense tensors MUST have buffer_count = 0x01 plus the number of quantization-parameter buffers required by the active scheme (see quantization.md § Buffer Table Placement Rules).

For sparse layouts (tags 0x06, 0x07, 0x08, 0x09, and future sparse tags), the buffer table MUST contain the number of entries specified by that layout's individual spec file. Each buffer holds a distinct component array (values, indices, pointers). Most sparse layouts have a fixed buffer count; CSF (0x09) is the exception, with a rank-dependent count of 2 × rank + 1 (see layouts/csf.md).

For indirect layouts (tag 0x0A, and future indirect tags), the buffer table MUST contain at least three entries (buffer_count >= 3): a values buffer plus the index/pointer buffers that resolve the logical-to-physical mapping. For block-paged (0x0A) these are buffer 0 = page_pool, buffer 1 = block_table, and buffer 2 = seq_ptr. When the tensor is quantized, the quantization-parameter buffers follow at indices 3 and up, per quantization.md § Buffer Table Placement Rules. See that layout's individual spec file for the exact buffer table.

For the virtual layout (composite head, tag 0x0B), the buffer table MUST be empty (buffer_count = 0x00): the head owns no data and supplies its logical view through its member tensors. See layouts/composite.md.

Splittability and Sharding

A tensor MAY be described as a shard: a rectangular sub-region of a larger logical tensor. A shard carries a shard descriptor in the tensor descriptor (see metadata.md § Shard Section) indicating its position within the parent index space.

Shard descriptor fields:

• parent_shape (uint64[rank]): shape of the logical parent tensor.
• shard_offset (uint64[rank]): starting index of this shard within the parent along each dimension.

The constraint shard_offset[k] + shape[k] <= parent_shape[k] MUST hold for every dimension k.

Note (non-normative): Sharding always produces rectangular sub-regions (hyperrectangles), which covers all practical partitioning patterns in ML inference (batch splitting, tensor parallelism, pipeline stages). Protocol-level splitting of tensor data across stream frames is a separate concern covered in interchange.md.

Extension Layouts

Layout tags in the range 0xF0–0xFE are reserved for implementation-private extension layouts. Tensors using these tags MUST NOT be exchanged between independent implementations unless both parties have agreed on the layout semantics out of band.

An extension layout descriptor MUST include:

• extension_layout_id (uint64): implementation-defined unique identifier.
• extension_data (byte sequence): opaque layout-specific metadata, preceded by a uint32 byte-length field.

Note (non-normative): Extension layouts are the mechanism for hardware-specific panel/pack formats used between BLAS pipeline stages. A client advertises an extension layout tag with opaque hardware metadata during capability negotiation (see interchange.md); the server transcodes and packs accordingly.

Custom Layouts

Any layout expressible as a composition of the named primitives — strides, tiling, space-filling curves — using a composite partition (layouts/composite.md) or recursive tiling is representable without the extension mechanism. Truly opaque custom layouts MUST use extension tags (0xF0–0xFE) with an out-of-band semantic agreement.

Interaction with Other Sections

• Element Types (element-types.md): defines bit widths, packing rules, and natural alignment for each element type.
• Quantization (quantization.md): block-quantized layouts interact with the memory layout tiling structure.
• Buffer Protocol (buffer-protocol.md): defines buffer ownership, alignment requirements, and device memory semantics.
• Metadata (metadata.md): defines the binary encoding of all layout-specific fields in the tensor descriptor.
• Interchange (interchange.md): defines how tensor descriptors and data buffers are transmitted. Protocol-level data splitting is distinct from logical sharding.

Row-Major (C Order) Layout — Hurray Format Specification

Layout tag: 0x01 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

In row-major layout, elements are stored with the last dimension varying fastest. Strides are implicit and MUST NOT be present in the descriptor for this layout tag.

Implicit Strides

strides[rank - 1] = 1
strides[i] = shape[i + 1] * strides[i + 1]    for i = rank - 2, ..., 0

All strides are in logical elements.

Element Address

The linear element offset of element [i_0, i_1, ..., i_{r-1}] in a row-major tensor of rank r is:

offset = sum(i_k * strides[k] for k = 0, ..., r - 1)

The byte address is computed from the element offset using the rules in memory-layout.md § Element Address Computation.

Buffer Size

For a contiguous row-major tensor, the minimum buffer size is num_elements * element_byte_width for whole-byte types, or ceil(num_elements / packing_factor) for sub-byte types, where num_elements is the product of all dimension sizes.

Additional Descriptor Fields

None. This layout has no layout-specific fields in the tensor descriptor.

Example

A rank-2 tensor with shape [3, 4] has implicit strides [4, 1]. Element [1, 2] is at linear offset 1 * 4 + 2 * 1 = 6.

Column-Major (Fortran Order) Layout — Hurray Format Specification

Layout tag: 0x02 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

In column-major layout, elements are stored with the first dimension varying fastest. Strides are implicit and MUST NOT be present in the descriptor for this layout tag.

Implicit Strides

strides[0] = 1
strides[i] = shape[i - 1] * strides[i - 1]    for i = 1, ..., rank - 1

All strides are in logical elements.

Element Address

The linear element offset of element [i_0, i_1, ..., i_{r-1}] is computed identically to row-major using the column-major strides above.

The byte address is computed from the element offset using the rules in memory-layout.md § Element Address Computation.

Buffer Size

Identical to row-major: num_elements * element_byte_width for whole-byte types, or ceil(num_elements / packing_factor) for sub-byte types.

Additional Descriptor Fields

None. This layout has no layout-specific fields in the tensor descriptor.

Example

A rank-2 tensor with shape [3, 4] has implicit strides [1, 3]. Element [1, 2] is at linear offset 1 * 1 + 2 * 3 = 7.

Strided Layout — Hurray Format Specification

Layout tag: 0x03 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

The strided layout generalises both row-major and column-major by allowing an arbitrary stride value per dimension. It is the most general of the simple (non-tiled, non-curve) layouts.

Additional Descriptor Fields

Stride Semantics

• A positive stride advances forward through the buffer.
• A negative stride advances backward. A negative stride on dimension k reverses that dimension: logical index 0 maps to the highest physical offset along that axis.
• A zero stride on dimension k means all indices along k map to the same physical element — a broadcast (virtual) dimension. Data is not physically replicated.

Negative and zero strides are valid. A conforming implementation MUST support them.

Element Address

The linear element offset of element [i_0, i_1, ..., i_{r-1}] is:

offset = sum(i_k * strides[k] for k = 0, ..., r - 1)

When negative strides are present, the offset may be negative relative to the base address at byte_offset. The byte_offset field MUST be set such that element [0, 0, ..., 0] is addressable within the buffer. The physical address of every valid element MUST lie within the buffer's bounds.

Buffer Size

The minimum buffer size must cover every addressable element:

max_offset = sum(max(0, strides[k] * (shape[k] - 1)) for all k)
min_offset = sum(min(0, strides[k] * (shape[k] - 1)) for all k)
range_elements = max_offset - min_offset + 1

Buffer size in bytes depends on the element type (see memory-layout.md § Element Address Computation).

Validity Constraints

This layout MUST NOT be used for rank-0 (scalar) tensors. See data-model.md § Scalar Tensors.

Contiguity

A strided tensor is dense (contiguous with no gaps) if and only if the absolute values of its strides form a permutation of the row-major strides for the same shape. Implementations SHOULD NOT assume density without verifying this condition.

Example

A rank-2 tensor with shape [3, 4] and strides [-4, 1] represents a row-major matrix with the row order reversed. byte_offset points to what would be element [2, 0] in a non-reversed matrix. Element [2, 3] is at offset 2 * (-4) + 3 * 1 = -5. The buffer MUST be large enough to cover the full range from offset -8 to 0 (9 elements).

Tiled / Blocked Layout — Hurray Format Specification

Layout tag: 0x04 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

A tiled layout partitions the tensor's index space into uniform rectangular tiles (blocks). This is a regular subpaving: all tiles have the same shape and tile the index space without overlap or gap (with possible padding at the boundaries).

Additional Descriptor Fields

If outer_layout is 0x03 (strided):

If inner_layout is 0x03 (strided):

If inner_layout is 0x04 (recursive tiling), the tiled layout-specific fields are encoded recursively beginning with tile_shape. A reader MUST enforce a maximum recursion depth (RECOMMENDED: 8 levels) and MUST reject descriptors that exceed it.

Element Address

To locate element [i_0, i_1, ..., i_{r-1}]:

1. Compute the tile index: t_k = floor(i_k / tile_shape[k]) for each dimension.
2. Compute the intra-tile offset: e_k = i_k mod tile_shape[k] for each dimension.
3. Compute the linear tile number using the outer layout applied to [t_0, ..., t_{r-1}]. Multiply by total_tile_elements to get the byte offset to the tile start.
4. Compute the intra-tile linear offset using the inner layout applied to [e_0, ..., e_{r-1}].
5. Final linear element offset = (3) + (4).

Boundary Padding

When a dimension size is not evenly divisible by the tile size, partial tiles exist at the boundary. The buffer MUST contain storage for full tiles, including padding. Padding element values are undefined; readers MUST NOT access elements whose logical index exceeds shape.

Number of tiles per dimension:

num_tiles[k] = ceil(shape[k] / tile_shape[k])

Total buffer elements:

total_tile_elements = product(tile_shape[k] for all k)
total_tiles        = product(num_tiles[k] for all k)
buffer_elements    = total_tiles * total_tile_elements

Validity Constraints

This layout MUST NOT be used for rank-0 (scalar) tensors. See data-model.md § Scalar Tensors.

Recursive Tiling

Note (non-normative): Recursive tiling is useful for hierarchical blocking in GEMM kernels (e.g., 128×128 L2 tiles subdivided into 32×32 L1 tiles). It is also the natural expression of cache-oblivious recursive blocking.

Example

Rank-2 tensor with shape [6, 8], tile_shape = [2, 4], outer_layout = 0x01 (row-major), inner_layout = 0x01 (row-major).

• Tile grid shape: [3, 2]. Outer strides (implicit): [2, 1] in tiles.
• Each tile: 8 elements. Inner strides (implicit): [4, 1] in elements.
• Element [3, 5]: tile [1, 1], intra-tile [1, 1]. Tile linear index = 1 * 2 + 1 = 3. Tile start = 3 * 8 = 24. Intra-tile offset = 1 * 4 + 1 = 5. Final offset = 29.

Morton (Z-Order Curve) Layout — Hurray Format Specification

Layout tag: 0x05 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

The Morton layout stores elements by interleaving the bits of their dimension indices, producing a linear order with good spatial locality for multi-dimensional access patterns.

Additional Descriptor Fields

Dimension Size Constraints

For each dimension k, shape[k] MUST satisfy shape[k] <= 2^morton_bits[k].

Morton Index Computation

For element [i_0, i_1, ..., i_{r-1}], the Morton code is computed by interleaving bits in round-robin order, starting from the least significant bit of dimension 0:

morton_code = 0
for bit_position b = 0, 1, 2, ...:
    for dimension d = 0, 1, ..., rank - 1:
        if b < morton_bits[d]:
            morton_code |= ((i_d >> b) & 1) << (b * rank + d)

The element at Morton code m is stored at linear offset m in the buffer.

Buffer Size

The buffer MUST hold 2^(sum(morton_bits[k] for all k)) elements.

Note (non-normative): For non-power-of-two dimension sizes, morton_bits[k] SHOULD be set to ceil(log2(shape[k])) to minimise padding. The worst-case padding factor per dimension is strictly less than 2×. Writers for whom padding waste is unacceptable SHOULD prefer a tiled or row-major layout.

Example

Rank-2 tensor with shape [4, 4], morton_bits = [2, 2].

Element [2, 3]: i_0 = 2 = 0b10, i_1 = 3 = 0b11. Interleaved (LSB first, dim 0 then dim 1): bit 0 of i_0=0, bit 0 of i_1=1, bit 1 of i_0=1, bit 1 of i_1=1. Morton code = 0b1110 = 14. Element [2, 3] is stored at linear offset 14.

Hilbert Curve Layout — Hurray Format Specification

Layout tag: 0x40 | Tier: 2

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

The Hilbert curve layout stores elements according to a Hilbert space-filling curve, which provides better locality than the Morton curve: consecutive Hilbert indices always correspond to physically adjacent elements (L∞ distance = 1 in index space), with no large jumps. This benefits access patterns that traverse multi-dimensional regions with spatial coherence (e.g., image convolution, volumetric sampling, point-cloud processing), at the cost of a more expensive index computation than Morton.

Note (non-normative): The Hilbert curve is particularly effective for 2D and 3D spatial tensors where Morton's large jumps at quadrant boundaries cause cache misses. For 1D access patterns or cases where index-computation cost dominates, row-major or Morton layouts are preferable.

Additional Descriptor Fields

Validity Constraints

A conforming reader MUST reject a Hilbert-curve descriptor that violates any of the following constraints:

1. shape[k] = 2^hilbert_order for every k in [0, hilbert_rank). All tensor dimensions MUST be a power of two equal to 2^hilbert_order.
2. hilbert_rank MUST equal the tensor's rank.
3. hilbert_rank MUST be greater than or equal to 2.
4. hilbert_order MUST be greater than 0.

This layout MUST NOT be used for rank-0 (scalar) tensors (constraints 2 and 3 together exclude rank 0). See data-model.md § Scalar Tensors.

Buffer Size

The buffer MUST hold exactly 2^(hilbert_rank * hilbert_order) elements.

Normative Index Mapping

The normative index mapping is the algorithm defined by Skilling (2004) (see references.md). Conforming implementations MUST use this algorithm. All arithmetic is integer arithmetic; ^ denotes bitwise XOR; array indexing is zero-based.

Let r = hilbert_rank, p = hilbert_order. Coordinates X[0..r-1] are each in [0, 2^p). Hilbert index h is in [0, 2^(r*p)).

Bit packing: bit (b * r + (r - 1 - d)) of h holds bit b of X[d], for b = 0, ..., p-1 and d = 0, ..., r-1.

CoordsToHilbert(X[0..r-1], r, p) → h:

M = 1 << (p - 1)
Q = M
while Q > 1:
    P = Q - 1
    for i = 0 to r-1:
        if X[i] & Q:
            X[0] ^= P
        else:
            t = (X[0] ^ X[i]) & P
            X[0] ^= t; X[i] ^= t
    Q >>= 1
for i = 1 to r-1:
    X[i] ^= X[i-1]
t = 0; Q = M
while Q > 1:
    if X[r-1] & Q: t ^= Q - 1
    Q >>= 1
for i = 0 to r-1:
    X[i] ^= t
h = 0
for b = 0 to p-1:
    for d = 0 to r-1:
        h |= ((X[d] >> b) & 1) << (b * r + (r - 1 - d))
return h

HilbertToCoords(h, r, p) → X[0..r-1]:

X = [0] * r
for b = 0 to p-1:
    for d = 0 to r-1:
        X[d] |= ((h >> (b * r + (r - 1 - d))) & 1) << b
t = X[r-1] >> 1
for i = r-1 downto 1:
    X[i] ^= X[i-1]
X[0] ^= t
Q = 2
while Q != (1 << p):
    P = Q - 1
    for i = r-1 downto 0:
        if X[i] & Q:
            X[0] ^= P
        else:
            t = (X[0] ^ X[i]) & P
            X[0] ^= t; X[i] ^= t
    Q <<= 1
return X

Conformance Check

Selected index mappings for r = 2, p = 2 (shape [4, 4]):

Consecutive entries differ by exactly 1 in exactly one coordinate. Implementations SHOULD validate against this table as a conformance check.

Note (non-normative): The MUST-level normative reference for the index mapping is the CoordsToHilbert / HilbertToCoords algorithm defined above. This table is provided as a SHOULD-level conformance aid; if a discrepancy is ever observed between table and algorithm, the algorithm output prevails.

COO (Coordinate) Sparse Layout — Hurray Format Specification

Layout tag: 0x06 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

The COO (Coordinate) format stores a sparse tensor as a list of explicitly enumerated non-zero elements. Each non-zero is described by its logical index tuple and its value. COO is the most general sparse format: it imposes no ordering requirement on the non-zeros and supports arbitrary rank.

Note (non-normative): COO is simple to construct (append-only) and easy to convert to other formats. It is less efficient for row-wise random access than CSR, but is well-suited for assembly (e.g., accumulating contributions from finite-element computations) and as an interchange format between sparse kernels.

Buffer Table

A COO tensor descriptor MUST have buffer_count = 2 in the buffer table.

Additional Descriptor Fields

Storage Order

Non-zeros are stored at positions 0 through nnz - 1 in values and indices. The i-th non-zero has:

• Value: values[i]
• Logical index: [indices[i * rank + 0], indices[i * rank + 1], ..., indices[i * rank + (rank-1)]]

If is_sorted = 0x01, the non-zeros MUST appear in lexicographic index order (i.e., sorted by dimension 0 first, then dimension 1, etc.). A conforming reader MAY use binary search for element lookup when is_sorted = 0x01.

If is_sorted = 0x00, no ordering guarantee is made. Readers MUST perform a linear scan to locate a specific element.

Validity Constraints

This layout MUST NOT be used for rank-0 (scalar) tensors. See data-model.md § Scalar Tensors.

A conforming writer MUST ensure:

1. Every stored index tuple is within bounds: 0 <= indices[i * rank + d] < shape[d] for all i in [0, nnz) and all d in [0, rank).
2. No two stored entries share the same index tuple (no duplicate coordinates).
3. If is_sorted = 0x01, entries are in strictly increasing lexicographic order (no ties, since duplicate coordinates are forbidden).

A conforming reader SHOULD validate constraints (1) and (2) and MUST reject descriptors that violate them, unless operating in permissive mode.

Buffer Size

• values buffer: nnz * element_byte_width bytes (or ceil(nnz / packing_factor) for sub-byte types).
• indices buffer: nnz * rank * 8 bytes (uint64 elements).

Both buffers MUST satisfy the alignment requirements in buffer-protocol.md.

Note (non-normative): The byte_offset field in the common descriptor header is not meaningful for sparse layouts — there is no single "first element" at a fixed offset. For COO tensors, byte_offset MUST be set to 0x0000000000000000.

Element Lookup

To retrieve the value at logical index idx[0..rank-1]:

1. If is_sorted = 0x01, perform a binary search on the indices buffer (comparing full rank-tuples in lexicographic order) to find a matching entry.
2. If is_sorted = 0x00, perform a linear scan over all nnz entries.
3. If a match is found at position i, return values[i].
4. If no match is found, the element is implicitly zero (or the zero value of the element type).

Interaction with Statistics Section

When the HAS_STATISTICS flag is set, the nnz field in the statistics section MUST match the nnz field in the COO descriptor, which is authoritative. The sparsity_ratio SHOULD be computed as 1.0 - (nnz / total_elements) where total_elements is the product of all shape values.

Example

Rank-2 sparse tensor with shape [4, 4], element type float32, 3 non-zeros:

Non-zeros: (0, 1) = 1.5,  (2, 0) = -0.5,  (3, 3) = 2.0
nnz = 3, is_sorted = 0x01

values  (buffer 0, 12 bytes):  [1.5, -0.5, 2.0]  as float32 LE

indices (buffer 1, 3×2×8 = 48 bytes, uint64 LE):
  entry 0: [0, 1]
  entry 1: [2, 0]
  entry 2: [3, 3]

CSR (Compressed Sparse Row) Layout — Hurray Format Specification

Layout tag: 0x07 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

The CSR (Compressed Sparse Row) format stores a sparse matrix (rank-2 tensor) by compressing the row structure: each row is represented by a contiguous slice of a flat non-zero array, with a separate pointer array indicating where each row begins. CSR is the most widely-used sparse matrix format in scientific computing and machine learning (e.g., sparse attention, graph neural networks, sparse weight matrices).

Note (non-normative): CSR is defined only for rank-2 tensors in this version of the specification. Generalisation to arbitrary rank (Compressed Sparse Fiber, CSF) is specified in CSF (Compressed Sparse Fiber).

A conforming implementation MUST reject a CSR descriptor whose rank is not 2.

Buffer Table

A CSR tensor descriptor MUST have buffer_count = 3 in the buffer table.

where nrows = shape[0].

Additional Descriptor Fields

Validity Constraints

This layout MUST NOT be used for rank-0 (scalar) tensors. See data-model.md § Scalar Tensors. CSR is further restricted to rank 2 (see § Description).

Storage Invariants

A conforming writer MUST ensure:

1. row_ptr[0] = 0 and row_ptr[nrows] = nnz.
2. row_ptr is non-decreasing: row_ptr[i] <= row_ptr[i+1] for all i.
3. Within each row i, the non-zeros in col_indices[row_ptr[i]..row_ptr[i+1]) MUST be sorted in strictly increasing order (no duplicate column indices per row).
4. All column indices are within bounds: 0 <= col_indices[j] < shape[1] for all j.

A conforming reader SHOULD validate these invariants and MUST reject descriptors that violate them, unless operating in permissive mode.

Buffer Size

• values buffer: nnz * element_byte_width bytes (or ceil(nnz / packing_factor) for sub-byte types).
• col_indices buffer: nnz * 8 bytes (uint64 elements).
• row_ptr buffer: (nrows + 1) * 8 bytes (uint64 elements), where nrows = shape[0].

All buffers MUST satisfy the alignment requirements in buffer-protocol.md.

For CSR tensors, byte_offset MUST be set to 0x0000000000000000.

Note (non-normative): The byte_offset field in the common descriptor header is not meaningful for CSR — there is no single "first element" at a fixed offset.

Element Lookup

To retrieve the value at row r, column c:

1. The non-zeros of row r occupy positions row_ptr[r] through row_ptr[r+1] - 1 in values and col_indices.
2. Perform a binary search on col_indices[row_ptr[r]..row_ptr[r+1]) for value c.
3. If found at position j, return values[j].
4. If not found, the element is implicitly zero.

Binary search is valid because column indices within each row are sorted (invariant 3).

Row Iteration

To iterate over all non-zeros in row r:

for j = row_ptr[r] to row_ptr[r+1] - 1:
    col = col_indices[j]
    val = values[j]

Interaction with Statistics Section

When the HAS_STATISTICS flag is set, the nnz field in the statistics section MUST match the nnz field in the CSR descriptor, which is authoritative. The sparsity_ratio SHOULD be 1.0 - (nnz / (shape[0] * shape[1])).

Example

Rank-2 sparse matrix with shape [4, 5], element type float32:

Dense representation:
  row 0: [1.0,  0,   0,  2.0,  0 ]
  row 1: [ 0,   0,  3.0,  0,   0 ]
  row 2: [ 0,  4.0,  0,   0,  5.0]
  row 3: [ 0,   0,   0,   0,   0 ]

nnz = 5

values      (buffer 0): [1.0, 2.0, 3.0, 4.0, 5.0]
col_indices (buffer 1): [0,   3,   2,   1,   4  ]
row_ptr     (buffer 2): [0,   2,   3,   5,   5  ]

row_ptr[3] = row_ptr[4] = 5 because row 3 has no non-zeros.

CSC (Compressed Sparse Column) Layout — Hurray Format Specification

Layout tag: 0x08 | Tier: 1

Also known as: CCS (Compressed Column Storage)

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

The CSC (Compressed Sparse Column) format is the column analog of CSR. It stores a sparse matrix (rank-2 tensor) by compressing the column structure: each column is represented by a contiguous slice of a flat non-zero array, with a separate pointer array indicating where each column begins.

CSC is preferred when the workload requires efficient column access or column iteration — e.g., sparse matrix-vector products with column-major dense vectors, column pivoting in sparse direct solvers, and graph algorithms that traverse out-edges (adjacency stored column-wise).

Note (non-normative): CSC is defined only for rank-2 tensors in this version of the specification, consistent with CSR; generalisation to arbitrary rank is provided by the Compressed Sparse Fiber (CSF) layout (see CSF (Compressed Sparse Fiber)). The format is also known as Compressed Column Storage (CCS) in some communities (e.g., MATLAB, some LAPACK interfaces). The two names refer to the same format.

A conforming implementation MUST reject a CSC descriptor whose rank is not 2.

Buffer Table

A CSC tensor descriptor MUST have buffer_count = 3 in the buffer table.

where ncols = shape[1].

Additional Descriptor Fields

Validity Constraints

This layout MUST NOT be used for rank-0 (scalar) tensors. See data-model.md § Scalar Tensors. CSC is further restricted to rank 2 (see § Description).

Storage Invariants

A conforming writer MUST ensure:

1. col_ptr[0] = 0 and col_ptr[ncols] = nnz.
2. col_ptr is non-decreasing: col_ptr[j] <= col_ptr[j+1] for all j.
3. Within each column j, the non-zeros in row_indices[col_ptr[j]..col_ptr[j+1]) MUST be sorted in strictly increasing order (no duplicate row indices per column).
4. All row indices are within bounds: 0 <= row_indices[i] < shape[0] for all i.

A conforming reader SHOULD validate these invariants and MUST reject descriptors that violate them, unless operating in permissive mode.

Buffer Size

• values buffer: nnz * element_byte_width bytes (or ceil(nnz / packing_factor) for sub-byte types).
• row_indices buffer: nnz * 8 bytes (uint64 elements).
• col_ptr buffer: (ncols + 1) * 8 bytes (uint64 elements), where ncols = shape[1].

All buffers MUST satisfy the alignment requirements in buffer-protocol.md.

Note (non-normative): As with CSR and COO, the byte_offset field in the common descriptor header is not meaningful for CSC. For CSC tensors, byte_offset MUST be set to 0x0000000000000000.

Element Lookup

To retrieve the value at row r, column c:

1. The non-zeros of column c occupy positions col_ptr[c] through col_ptr[c+1] - 1 in values and row_indices.
2. Perform a binary search on row_indices[col_ptr[c]..col_ptr[c+1]) for value r.
3. If found at position i, return values[i].
4. If not found, the element is implicitly zero.

Binary search is valid because row indices within each column are sorted (invariant 3).

Column Iteration

To iterate over all non-zeros in column c:

for i = col_ptr[c] to col_ptr[c+1] - 1:
    row = row_indices[i]
    val = values[i]

Relationship to CSR

CSC is the transpose of CSR: the CSC representation of matrix A is equivalent to the CSR representation of A^T with shape swapped. Implementations that support both formats MAY convert between them by transposing the pointer and index arrays.

Interaction with Statistics Section

When the HAS_STATISTICS flag is set, the nnz field in the statistics section MUST match the nnz field in the CSC descriptor, which is authoritative. The sparsity_ratio SHOULD be 1.0 - (nnz / (shape[0] * shape[1])).

Example

The same matrix as in the CSR example (shape [4, 5], element type float32):

Dense representation:
  row 0: [1.0,  0,   0,  2.0,  0 ]
  row 1: [ 0,   0,  3.0,  0,   0 ]
  row 2: [ 0,  4.0,  0,   0,  5.0]
  row 3: [ 0,   0,   0,   0,   0 ]

nnz = 5

values      (buffer 0): [1.0, 4.0, 3.0, 2.0, 5.0]
row_indices (buffer 1): [0,   2,   1,   0,   2  ]
col_ptr     (buffer 2): [0,   1,   2,   3,   4,   5]

col_ptr[j+1] - col_ptr[j] = number of non-zeros in column j: col 0: 1 (value 1.0 at row 0), col 1: 1 (value 4.0 at row 2), col 2: 1 (value 3.0 at row 1), col 3: 1 (value 2.0 at row 0), col 4: 1 (value 5.0 at row 2).

CSF (Compressed Sparse Fiber) Layout — Hurray Format Specification

Layout tag: 0x09 | Tier: 1

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

The CSF (Compressed Sparse Fiber) format is the rank-N generalisation of CSR/CSC. It stores a sparse tensor as a tree of rank levels, where each level compresses one mode (logical dimension) with a (pos, crd) pair: a pos pointer array delimiting each parent's children, and a crd coordinate array naming those children. A single values array holds the non-zero values at the leaves. CSF is the natural higher-rank complement to COO for structured sparsity — sparse attention masks, sparse activations, and higher-order tensor factorisations.

Every level of a CSF tree is compressed: there is no per-mode dense/compressed distinction in this version of the specification. Dense-outer rank-2 cases are covered by CSR/CSC.

Note (non-normative): CSF is defined only for rank ≥ 3 tensors in this version. Rank-2 sparse matrices are served by CSR (CSR (Compressed Sparse Row)) and CSC (CSC (Compressed Sparse Column)), whose dense outer pointer array is more compact and more interop-canonical. COO (COO (Coordinate)) (which supports any rank) and CSR/CSC are preferable for rank ≤ 2. A future revision could admit dense levels via a mode_format field; a reader of this version rejecting such a descriptor is intended versioning behaviour.

A conforming implementation MUST reject a CSF descriptor whose rank is less than 3.

Mode Ordering

A CSF descriptor carries a mode_order: uint32[rank] field, a permutation of the integers 0..rank-1. mode_order[L] is the logical dimension stored at tree level L, where level 0 is the outermost level and level rank-1 is the leaf level directly above values. The bounding size for level L — the upper bound on every coordinate stored at that level — is shape[mode_order[L]].

mode_order affects storage traversal only; it does not change the tensor's logical shape.

A conforming reader MUST honour any valid mode_order permutation for lookup and iteration. A conforming reader MUST NOT reject a CSF descriptor merely because mode_order is non-identity. A conforming reader MUST reject a descriptor whose mode_order is not a valid permutation of 0..rank-1 (e.g. contains a duplicate or an out-of-range value), unless operating in permissive mode.

When a writer has no access-pattern preference, the identity ordering [0, 1, ..., rank-1] (row-major, outer-to-inner) SHOULD be used as the default.

Note (non-normative): mode_order is a performance knob: it lets a writer match the tree's nesting to its access pattern. The identity ordering is preferred as the default because it is reproducible across writers and keeps the outermost (slowest- varying) logical dimension at the top of the tree, which is cache-friendly for row-major access.

Buffer Table

A CSF tensor descriptor MUST have buffer_count = 2 * rank + 1 in the buffer table (plus any quantization-parameter buffers, which follow at indices 2 * rank + 1 and up).

where n_L is the number of nodes stored at level L, n_{-1} = 1 (a single virtual root), and n_{rank-1} = nnz. The top-level pointer array pos_0 always has length 2 and equals [0, n_0].

Additional Descriptor Fields

Validity Constraints

This layout MUST NOT be used for rank-0 (scalar), rank-1, or rank-2 tensors. CSF requires rank >= 3 (see § Description); rank-1 and rank-2 cases are served by COO, CSR, and CSC. Rank remains capped at 64 by data-model.md.

Storage Invariants

A conforming writer MUST ensure, for every level L in 0..rank-1:

1. pos_L[0] = 0.
2. pos_L is non-decreasing: pos_L[k] <= pos_L[k+1] for all k.
3. The terminal pos entry equals the level's stored count: pos_0[1] = n_0, pos_L[n_{L-1}] = n_L, and n_{rank-1} = nnz.
4. Within each parent slice crd_L[pos_L[k]..pos_L[k+1]), the coordinates MUST be sorted in strictly increasing order (no duplicate siblings).
5. All coordinates are within bounds: 0 <= crd_L[i] < shape[mode_order[L]] for all i.

In addition, mode_order MUST be a valid permutation of 0..rank-1.

For an empty tensor (nnz = 0), pos_0 MUST be [0, 0], values and every crd_L MUST have length 0, and for each level L >= 1 pos_L MUST be [0] (length 1, since n_{L-1} = 0).

A conforming reader SHOULD validate these invariants and MUST reject descriptors that violate them, unless operating in permissive mode.

Buffer Size

• values buffer: nnz * element_byte_width bytes (or ceil(nnz / packing_factor) for sub-byte types).
• pos_L buffer: (n_{L-1} + 1) * 8 bytes (uint64 elements); pos_0 is always 2 * 8 = 16 bytes.
• crd_L buffer: n_L * 8 bytes (uint64 elements); the leaf crd_{rank-1} is nnz * 8 bytes.

All buffers MUST satisfy the alignment requirements in buffer-protocol.md.

For CSF tensors, byte_offset MUST be set to 0x0000000000000000.

Note (non-normative): As with COO, CSR, and CSC, the byte_offset field in the common descriptor header is not meaningful for CSF — there is no single "first element" at a fixed offset; the first stored element is located by descending the tree.

Element Lookup

To retrieve the value at logical index idx = [idx[0], ..., idx[rank-1]]:

1. Permute the query into storage order: the coordinate sought at level L is q_L = idx[mode_order[L]].
2. Initialise the parent position p = 0 (the virtual root).
3. For each level L from 0 to rank-1:
   • The children of the current parent occupy the slice crd_L[pos_L[p] .. pos_L[p+1]).
   • Perform a binary search on that slice for q_L.
   • If the binary search finds q_L at relative offset k within the slice (so that crd_L[pos_L[p] + k] == q_L), set p = pos_L[p] + k — the absolute index into crd_L — and continue to the next level.
   • If not found, the element is implicitly zero; stop.
4. After the leaf level (L = rank-1), p is the leaf position; return values[p].

Binary search is valid at each level because sibling coordinates within a parent slice are sorted (invariant 4).

Iteration

To iterate over all non-zeros, descend the tree in depth-first, leaf order. The following sketch accumulates the storage-order coordinate tuple c[0..rank-1]; the logical index is recovered by placing c[L] at dimension mode_order[L]:

def visit(L, p):
    for i = pos_L[p] to pos_L[p+1] - 1:
        c[L] = crd_L[i]
        if L == rank - 1:
            emit(c, values[i])      # i is the leaf position
        else:
            visit(L + 1, i)

visit(0, 0)

Relationship to CSR / CSC / COO

CSF generalises CSR/CSC to arbitrary rank but does not replace them. For rank-2 tensors, writers SHOULD prefer CSR (CSR (Compressed Sparse Row)) or CSC (CSC (Compressed Sparse Column)), whose dense outer pointer array is more compact and more interop-canonical; for rank-1 and rank-2 coordinate lists, writers SHOULD prefer COO (COO (Coordinate)). CSF MUST NOT be used below rank 3.

A rank-2 CSF tree with mode_order = [0, 1] is structurally analogous to CSR (pos_0 the trivial root pointer, pos_1 the row pointer, crd_1 the column indices), and with mode_order = [1, 0] to CSC — but CSF stores an explicit top-level pos_0/crd_0 pair that CSR/CSC fold into their dense outer array, so the layouts are not byte-identical.

Quantization Compatibility

The values buffer (buffer index 0) holds the quantized leaf values, exactly as for COO and CSR. Quantization schemes decorate the leaves: per-tensor, per-channel, per-block, NF4, and MXFP compose as for COO/CSR, following the placement rules in quantization.md. Quantization-parameter buffers occupy buffer indices 2 * rank + 1 and up. All quantization-parameter buffers MUST carry the same device_tag as the values buffer.

Sharding

A CSF tensor MUST NOT carry a shard descriptor in this version of the specification. Sharding a CSF tree (the interaction between a shard's shard_offset and the per-level pos/crd structure) is deferred to a future revision.

Interaction with Statistics Section

When the HAS_STATISTICS flag is set, the nnz field in the statistics section MUST match the nnz field in the CSF descriptor, which is authoritative. The sparsity_ratio SHOULD be 1.0 - (nnz / total_elements), where total_elements is the product of all shape entries.

Example

Rank-3 sparse tensor with shape [2, 3, 4], element type float32, mode_order = [0, 1, 2] (identity), 4 non-zeros:

Non-zeros (logical index -> value):
  (0, 0, 1) -> 1.0
  (0, 2, 3) -> 2.0
  (1, 1, 0) -> 3.0
  (1, 1, 2) -> 4.0

nnz = 4
buffer_count = 2 * 3 + 1 = 7

Level 0 (dimension 0, bound 2):
  pos_0   (buffer 1): [0, 2]
  crd_0   (buffer 2): [0, 1]              # two non-empty slices: i=0, i=1

Level 1 (dimension 1, bound 3):
  pos_1   (buffer 3): [0, 2, 3]           # parent 0 has 2 children, parent 1 has 1
  crd_1   (buffer 4): [0, 2, 1]           # i=0: j=0,2 ; i=1: j=1

Level 2 (dimension 2, bound 4):
  pos_2   (buffer 5): [0, 1, 2, 4]        # leaf delimiters
  crd_2   (buffer 6): [1, 3, 0, 2]        # the k coordinates

values    (buffer 0): [1.0, 2.0, 3.0, 4.0]

Lookup of (1, 1, 2):

q = [1, 1, 2]                 # mode_order is identity

Level 0: parent p = 0, slice crd_0[pos_0[0]..pos_0[1]) = crd_0[0..2) = [0, 1]
         search 1 -> relative offset k = 1; p = pos_0[0] + 1 = 0 + 1 = 1
Level 1: slice crd_1[pos_1[1]..pos_1[2]) = crd_1[2..3) = [1]
         search 1 -> relative offset k = 0; p = pos_1[1] + 0 = 2 + 0 = 2
Level 2: slice crd_2[pos_2[2]..pos_2[3]) = crd_2[2..4) = [0, 2]
         search 2 -> relative offset k = 1; p = pos_2[2] + 1 = 2 + 1 = 3

leaf position p = 3 -> values[3] = 4.0

Block-Paged Layout — Hurray Format Specification

Status: Draft

Layout tag: 0x0A | Tier: 1 | Type: Indirect

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

The block-paged format stores a tensor whose paged axis is divided into fixed-size pages drawn from a shared page pool, with a block table mapping each logical page position to a physical page ID. It is the interchange form of a PagedAttention KV cache: one page pool plus a ragged set of per-sequence block tables, addressed CSR-style by a seq_ptr offset array.

Block-paged is an indirect-dense layout: every logical element along the paged axis exists (there are no implicit zeros, unlike sparse layouts), but the mapping from a logical index to a physical buffer position is resolved through the block table rather than by an affine stride formula. A block-paged descriptor describes a static snapshot of one whole batch for one {kv_role, layer_index} pair. It carries no live allocator state.

Prefix sharing across sequences is expressed as static structure: two block-table entries MAY name the same physical page ID. The aliasing is internal to the block table; it creates no shared buffer ownership (see § Prefix Sharing).

Note (non-normative): Block-paged is defined only for rank-3 tensors in this version of the specification. Generalisation to other ranks is left to a future revision. The whole transformer KV cache ([layers, 2, ...]) is transmitted as a stream of descriptors, one per key/value role per layer (see ADR-024).

A conforming implementation MUST reject a block-paged descriptor whose rank is not 3.

Note (non-normative): "Page" and "block" refer to the same fixed-size physical unit. "Page" names the storage slot in the page_pool; "block" names the entry in the index structure (block_table) that maps a logical page position to a physical page ID. This dual naming matches the vLLM convention (page pool, block table).

Logical Shape

The logical shape is [total_tokens, num_heads, head_dim], where total_tokens is the sum of the per-sequence token counts across the batch. The ragged per-sequence structure is carried by seq_ptr (see § Buffer Table), not by the shape, which remains a hyperrectangle. The paged axis is the axis subdivided into pages; in this version it MUST be axis 0 (paged_axis = 0).

Buffer Table

A block-paged tensor descriptor MUST have buffer_count >= 3 in the buffer table. Buffers 0–2 are always present; buffers 3 and up are present only when the HAS_QUANTIZATION flag is set, per the placement rules in quantization.md.

where total_logical_pages = seq_ptr[num_seqs].

Note (non-normative): Buffers 1 (block_table) and 2 (seq_ptr) share the same element type, selected by block_table_index_type.

Note (non-normative): Empty batch (num_seqs = 0): block_table (buffer 1) is empty (byte_size = 0, and MAY be a null pointer per buffer-protocol.md), and seq_ptr (buffer 2) contains exactly one element, seq_ptr[0] = 0.

Additional Descriptor Fields

All multi-byte fields in the block-paged descriptor MUST be little-endian.

Note (non-normative): A uint32 block table addresses up to 0xFFFFFFFF (4294967295) pages. Larger pools MUST use uint64 (block_table_index_type = 0x01).

byte_offset

For block-paged tensors, byte_offset MUST be set to 0x0000000000000000.

Note (non-normative): The byte_offset field in the common descriptor header is not meaningful for block-paged — logical position 0 is located via the block table, not at a fixed offset, exactly as for sparse layouts.

Storage Invariants

A conforming writer MUST ensure:

1. seq_ptr[0] = 0 and seq_ptr[num_seqs] = total_logical_pages.
2. seq_ptr is non-decreasing: seq_ptr[s] <= seq_ptr[s+1] for all s.
3. All physical page IDs are within bounds: 0 <= block_table[k] < num_pages for all k.
4. The page_pool buffer size is num_pages * page_size * num_heads * head_dim * element_byte_width bytes (or the sub-byte equivalent, ceil(num_pages * page_size * num_heads * head_dim / packing_factor), where packing_factor is as defined in memory-layout.md § Sub-Byte Types).
5. When HAS_QUANTIZATION is set, scales (and zero_points, when present) are indexed by the same block_table, with one entry per page slot (num_pages * page_size entries).
6. paged_axis = 0.

A conforming reader SHOULD validate invariants (1)–(3) and (6) and MUST reject descriptors that violate them, unless operating in permissive mode. A reader is NOT required to validate that aliased page IDs carry identical content: aliasing is expressible, not validated (see § Prefix Sharing).

Partial trailing page. A sequence's token count need not be a multiple of page_size; the slots in its final page beyond the valid token count are undefined. A reader MUST NOT read past a sequence's valid token count (bounded by seq_ptr and the sequence's token count). A writer SHOULD zero those unused slots when transferring across a trust or tenant boundary.

Element Lookup

To retrieve the value at token t of sequence s, head h, dimension d:

page_in_seq    = t / page_size            (integer division)
offset_in_page = t mod page_size
phys_page      = block_table[seq_ptr[s] + page_in_seq]
flat           = ((phys_page * page_size + offset_in_page) * num_heads + h) * head_dim + d
value          = page_pool[flat]

When the tensor is quantized, dequantize value using scales[phys_page * page_size + offset_in_page] (and the corresponding zero_points entry for an asymmetric scheme), per the active scheme in quantization.md.

Prefix Sharing

Two sequences share a prefix when their block-table slices name the same physical page IDs for the shared leading positions. For example, with

seq 0 block-table slice: [12, 5, 7, 9]
seq 1 block-table slice: [12, 5, 7, 3]

sequences 0 and 1 share physical pages 12, 5, and 7. The aliasing is internal to the block_table buffer: it creates no shared buffer ownership and no wire-level reference count (see ADR-024 and ADR-009). The engine-internal copy-on-write reference counts that govern live page lifetime are out of scope; a block-paged descriptor is a static snapshot.

Quantization Compatibility

Block-paged tensors MAY be quantized using the schemes defined in quantization.md. The quantization-parameter buffers occupy buffer indices 3 and up, per quantization.md § Buffer Table Placement Rules. The following rules govern how those schemes compose with the paged structure.

Per-page-slot parameters

For a quantized block-paged tensor, scales (and zero-points, when present) are stored per page slot: the scale array has exactly num_pages * page_size entries, and it is paged through the same block_table as the values. The scale for the element at physical page p, slot i lives at scales[p * page_size + i]. This is the indexing already given in § Element Lookup: a reader dequantizes using scales[phys_page * page_size + offset_in_page] (and the corresponding zero_points entry for an asymmetric scheme). Zero-point parameters, when carried in a separate buffer, follow the same per-page-slot layout.

Note (non-normative): Per-page co-location is required so that a shared (aliased) page carries its own scales. Because two sequences may name the same physical page ID in their block_table slices (see § Prefix Sharing), storing scales per logical token would force a shared page's leading tokens to be dequantized with different scales depending on the aliasing sequence, breaking numerical coherence of the prefix. Paging the scales through the same block_table as the values keeps every aliased reference to a page numerically identical.

Scale-buffer size

The standard per-block-affine scale-buffer-size formula in quantization.md (which yields shape[axis] / block_size scale entries) does NOT apply to block-paged. A reader MUST compute the scale-buffer size as num_pages * page_size entries from the page structure, because scales are stored per physical page slot rather than per logical token. This is what allows an aliased/shared page to carry its own scales and keeps prefix sharing numerically coherent.

Per-scheme composition

• Per-tensor schemes (scheme_tag = 0x01, see quantization.md) compose normally: one scale (and zero-point) applies to the whole tensor.
• Per-channel schemes (scheme_tag = 0x02) compose normally when the quantization axis is num_heads (axis 1) or head_dim (axis 2). The per-channel parameters are indexed by channel along that axis and are independent of the paged structure.
• Per-block-affine (scheme_tag = 0x03) composes only under the constraint that the quantization descriptor's axis field MUST be 0 (the paged / token axis) and block_size MUST equal page_size. Under this constraint the per-block parameters coincide exactly with the per-page-slot parameters described above. A reader MUST reject a block-paged descriptor using scheme_tag = 0x03 whose axis field is not 0, or whose block_size is not equal to page_size.

A writer MUST NOT quantize the paged / token axis (axis 0) with the per-channel scheme (scheme_tag = 0x02); paged-axis quantization MUST use per-block-affine under the block_size == page_size constraint above, so that scales remain per-page-slot and aliasing stays coherent.

Sharding

A shard descriptor (see memory-layout.md § Splittability and Sharding) MUST NOT be applied to a block-paged tensor in this version of the specification. A reader MUST reject a block-paged descriptor that also carries a shard descriptor.

Note (non-normative): Tensor-parallel / multi-GPU sharding of a block-paged KV cache (for example, splitting along the num_heads axis) is deferred to a future revision. Before sharding can be permitted, the interaction between a shard descriptor's shard_offset and the absolute seq_ptr offsets into the block_table must be resolved, since seq_ptr indexes the whole-batch block table rather than a shard-local slice. The normative rule for this version is stated above: a reader MUST reject a block-paged descriptor that also carries a shard descriptor.

Alignment

All buffers MUST satisfy the alignment requirements in buffer-protocol.md (at least 64 bytes; page-aligned for GPU, IPC, or RDMA). All buffers in a block-paged descriptor MUST share the same device_tag and memory_class.

Framework Compliance

Note (non-normative): There is no formal PagedAttention specification. This layout is designed to be faithful to the dominant inference frameworks, principally vLLM.

In vLLM's FlashAttention / FlashInfer backend, the per-layer KV cache is stored as one tensor per key and per value with shape [num_blocks, block_size, num_kv_heads, head_dim], and the default block_size is 16. The block table is a per-sequence list of physical block numbers, and sequences may share physical blocks (copy-on-write prefix sharing). These map onto Hurray as follows:

vLLM concept	Hurray block-paged
num_blocks	num_pages
block_size (default 16)	page_size
num_kv_heads	num_heads
head_dim	head_dim
per-K / per-V tensor	one descriptor per kv_role per layer_index
per-sequence block table	block_table slice delimited by seq_ptr
shared physical block	aliased page ID in block_table

Hurray encodes the portable interchange layout ([num_blocks, block_size, num_kv_heads, head_dim] per role), not a kernel-internal reshaped layout. In particular, the older paged_attention_v1 key-cache reshape [num_blocks, num_kv_heads, head_dim/x, block_size, x] is an engine implementation detail and is out of scope for this layout. A producer that stores its cache in a kernel-internal layout MUST transpose to the interchange layout before emitting a block-paged descriptor.

Example

A KV cache for layer 3, key role, with 2 sequences, page_size = 4, num_pages = 5, num_heads = 2, head_dim = 8, element type float16. Sequence 0 has 6 tokens, sequence 1 has 3 tokens, and sequence 1 reuses page 0 (a shared prefix):

num_seqs    = 2
seq_ptr     (buffer 2): [0, 2, 3]
block_table (buffer 1): [0, 1, 0]
page_pool   (buffer 0): 5 * 4 * 2 * 8 = 320 float16 values

kv_role                = 0x00   (key)
layer_index            = 3
page_size              = 4
num_pages              = 5
paged_axis             = 0
block_table_index_type = 0x00   (uint32)

Sequence 0 occupies block_table[0..2) = [0, 1] (6 tokens span 2 pages: page 0 holds tokens 0–3, page 1 holds tokens 4–5, with slots 6–7 of page 1 undefined). Sequence 1 occupies block_table[2..3) = [0] (3 tokens in page 0, slot 3 undefined), aliasing page 0 — the shared prefix.

Composite / Virtual Tensor — Hurray Format Specification

Layout tag: 0x0B | Tier: 1 | Type: Virtual

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Description

A composite tensor is a head descriptor plus an ordered set of member tensors, combined under a declared composition rule. The head presents a single logical view — one shape and one type_tag — while the members supply the actual data, each as a complete, ordinary tensor descriptor with its own layout, buffers, quantization, statistics, and device placement.

The head is a virtual (data-less) descriptor: it owns no buffers. This is a new addressing category, Virtual, alongside Dense, Sparse, and Indirect (see memory-layout.md § Layout Taxonomy). A composite unifies three previously distinct capabilities under one primitive:

• Partition — a logical tensor whose index space is an exact, non-overlapping tiling of heterogeneous regions, each a full member tensor (see § Composition Semantics › Partition for the coverage and non-overlap constraints).
• Overlay — a base tensor spanning the whole index space plus scattered corrections at shared indices (SpQR / KVQuant outlier quantization).
• Group — several independent tensors delivered together under one logical identity (multi-output inference; weight collections).

Note (non-normative): A member is an ordinary TensorDescriptor. Everything a region needs — per-member layout, buffers, quantization, statistics, device tags — comes from the existing descriptor machinery, so the composite adds no per-member wire format of its own beyond the small Composite Member section (§ Composite Member Section). The crux of the unification is that a region is a shard is a member.

Note (non-normative): This version specifies only the v1.0-accepted scope: partition, group, and sealed overlay. Versioned / open overlay — an appendable, time-travel-capable overlay identified by an open-composite member_count sentinel and a per-member version field — is deferred to a future ADR (see ADR-027 § Status). Wire space is reserved for it (the member_count sentinel 0xFFFFFFFF in § Head Layout-Specific Fields, and the reserved padding in § Composite Member Section) so the future feature is an additive change, not a re-layout. Those reserved encodings are not usable in v1.0; the normative rejection rules for them appear in § Head Layout-Specific Fields and § Composite Member Section below.

Head Descriptor

The head is an ordinary tensor descriptor (see metadata.md) with layout_tag = 0x0B. It carries the composite's logical shape (shape) and logical element type (type_tag) — the view the composite presents to a consumer.

The head owns no data:

• buffer_count MUST be 0x00.
• byte_offset MUST be 0x0000000000000000.

A reader MUST reject a head descriptor whose buffer_count is not 0x00, or whose byte_offset is not 0x0000000000000000.

The head MUST NOT set the HAS_COMPOSITE_MEMBER flag (bit 4); that flag applies to members only (§ Composite Member Section). Because the head owns no data buffer, it MUST NOT set the HAS_QUANTIZATION flag (there is no stored data to dequantize). When a head is itself a member of an enclosing partition or overlay composite (a nested composite; see § Binding), it MUST carry a shard section exactly as any other member does.

A strict reader MUST reject an unrecognised layout_tag = 0x0B. A permissive reader MAY read the head's shape and type_tag but MUST NOT dereference tensor data through the head (there is none).

Head Layout-Specific Fields

Immediately following the head's byte_offset (per metadata.md § Layout-Specific Fields), the composite head encodes the following fields. All multi-byte fields are little-endian.

A reader MUST reject a head whose composition_rule is 0x00, 0xFF, or in the reserved range 0x04–0xEF. A reader that does not recognise an implementation-private composition_rule in 0xF0–0xFE MUST reject the head unless operating in permissive mode with an out-of-band agreement.

A reader MUST reject a head whose combine_op is not 0x00 when composition_rule is 0x01 (partition) or 0x03 (group). When composition_rule is 0x02 (overlay), a reader MUST reject a head whose combine_op is not 0x01 or 0x02.

A reader MUST reject a head whose _reserved bytes are not all 0x00.

member_count sentinel — RESERVED. The value 0xFFFFFFFF denotes an open composite and is RESERVED for a future ADR (versioned / open overlay). A strict v1.0 reader MUST reject a head whose member_count is 0xFFFFFFFF. All v1.0 composites — partition, group, and overlay — use a definite member_count.

Note (non-normative): member_count = 0x00000000 (a head with no members) is a definite count of zero. It is syntactically valid but carries no data view; writers are not expected to emit it. Validation for each rule at "the Nth member" is trivially satisfied at N = 0 except where a rule requires a base member (overlay), which such a head cannot supply — see § Validation.

Members

A member is a complete, ordinary TensorDescriptor with its own layout_tag (dense, sparse, indirect, or a nested composite), its own buffer table, and its own optional quantization, statistics, and device tags.

Shard section

For partition (0x01) and overlay (0x02) composites, every member MUST carry a shard section (HAS_SHARD flag set; see metadata.md § Shard Section, and ADR-004). Its parent_shape MUST equal the head's logical shape, and its shard_offset together with the member's own shape define the member's box in the head's index space:

member covers, along dimension k, the half-open range
    [ shard_offset[k], shard_offset[k] + shape[k] )

For group (0x03) composites, members MAY omit the shard section (group members have no spatial relationship to the head; see § Composition Semantics).

Element type across members

For partition and overlay, each member MAY declare its own stored element type and its own quantization scheme, but each member's decoded value type MUST equal the head's type_tag. Dequantization already yields a canonical real-valued view, so no new machinery is required.

Note (non-normative): SpQR example — head type_tag = float16; the base member is int4 with per-block-affine quantization decoding to float16; an outlier correction member is a float16 COO sparse tensor. The overlay combine is evaluated in float16 (the head's type). Composite versioning would change values over a fixed index space; shape evolution is out of scope.

For group composites, the head's shape and type_tag are advisory and members MAY differ arbitrarily (see § Composition Semantics).

Composite Member Section

Overlay members carry a Composite Member section, a new optional descriptor section gated by descriptor flag bit 4, HAS_COMPOSITE_MEMBER, defined in metadata.md § Composite Member Section and appended after the Extension Type section. Its single v1.0 field is member_role:

Partition and group members MUST NOT carry a Composite Member section (they MUST NOT set HAS_COMPOSITE_MEMBER). An overlay member MUST carry one.

Note (non-normative): v1.0 carries member_role only. A sealed overlay's precedence is plain stream / emission order (§ Composition Semantics), so no explicit version field is needed. The section's reserved padding leaves room for a future ADR to add a member_version field additively without reallocating the section.

Binding

A head with a definite member_count = N binds the next N self-delimiting tensors in stream / file write order as its members. This is a forward promise — the head precedes its members, which precede their data — not a back-reference. It introduces no name namespace and is streamable for both readers and writers.

• In-process: the head handle plus an array of N member handles (no wire concern).
• IPC / network streaming: the framing and "close" rules are defined in interchange.md § Composite Tensor Streaming.
• File: the head + members occupy consecutive index entries in the tensor region; the recovery rule is defined in file-format.md § Composite Tensors.

Nested composites are permitted: a member MAY itself be a head with layout_tag = 0x0B, parsed pre-order. A reader MUST enforce a maximum composite nesting depth of 8 levels and MUST reject a descriptor that exceeds it (the same recursion-depth discipline used for nested Tiled layouts, metadata.md § Layout-Specific Fields › Tiled / Blocked).

Note (non-normative): Plain sharding (members without a head; see interchange.md § Parallel Transfers) is the status quo. The head upgrades an ephemeral shard set into a persistent, composition-typed collection. Explicit member identifiers for out-of-order random access are not defined in this version (see § Deferred).

Composition Semantics

Partition (0x01)

Members' shard boxes MUST exactly cover the head's index space with no overlap. Each logical index belongs to exactly one member, so the composite view is zero-copy: value lookup at logical index idx is (1) select the member whose box contains idx, (2) compute the local index idx - shard_offset, (3) apply that member's own addressing.

Coverage constraint. The union of all member boxes MUST exactly cover every element in the head's index space: for every valid index [i_0, i_1, ..., i_{r-1}] (where 0 <= i_k < shape[k] for all k, shape being the head's shape), there MUST be exactly one member whose box — [shard_offset[k], shard_offset[k] + shape[k]) along each dimension k (§ Members › Shard section) — contains that index.

Non-overlap constraint. Two members' boxes A and B overlap if, for every dimension k:

A.shard_offset[k] < B.shard_offset[k] + B.shape[k]
AND
B.shard_offset[k] < A.shard_offset[k] + A.shape[k]

Member boxes MUST NOT overlap. A conforming writer MUST produce a partition composite whose members satisfy both constraints. A conforming reader SHOULD validate them and MUST reject a violating composite unless in permissive mode (§ Validation).

Overlay (0x02, v1.0: sealed only)

One member is the base (member_role = 0x01): it MUST be the first member and its box MUST span the whole index space (its shard shard_offset is all-zero and its shape equals the head's shape). The remaining members are corrections (member_role = 0x00) whose boxes MAY overlap one another and the base.

Precedence is emission / stream order — later wins. The writer emits the base first, then corrections in the order they take effect. A single-pass reader applies members as they arrive; no version field or reordering is required.

Reads apply the base, then all corrections in emission order, under combine_op, evaluated in the head's type_tag:

• combine_op = 0x01 (replace): within a correction's box, the topmost (latest-emitted) member covering an index wins; outside every correction's box, the base shows through.
• combine_op = 0x02 (add): the value at an index is the base value plus the sum of all covering corrections' values at that index.

A correction's box replaces or adds within its box only; outside it, lower-precedence members (down to the base) show through. Per-member storage is zero-copy, but the merged logical view is computed by the consumer — a sealed overlay is not zero-copy at the composite level.

Note (non-normative): Replace serves region overwrites; add serves residual / outlier overlays (SpQR / KVQuant). A v1.0 overlay is a complete snapshot: definite member_count, not appendable without rewriting the head. Logical delete (reverting a region to base, or masking it) is out of scope for v1.0; a region is reverted by writing a new correction that carries the desired values. A data-less tombstone member kind is reserved (see § Deferred).

Group (0x03)

Members are independent tensors under one head identity, with no spatial semantics and no ordering semantics. The head's shape and type_tag are advisory; members MAY differ arbitrarily in rank, shape, element type, layout, and device. This occupies the grouping gap (see ADR-010) using forward adjacency, not naming.

Validation

Validation is cross-member and stateful but bounded: because every v1.0 composition rule uses a definite member_count = N, a reader accumulates state only up to N, reaches one verdict at the Nth member, and is done.

Per-member checks (immediate)

On each member, a reader MUST verify:

1. For partition and overlay: the member carries a shard section whose parent_shape equals the head's shape, and whose box is in bounds (shard_offset[k] + shape[k] <= the head's shape[k] for every dimension k).
2. The member's decoded value type equals the head's type_tag (§ Members).
3. combine_op is legal for the composition rule (§ Head Layout-Specific Fields).
4. For overlay: the member carries a Composite Member section with a valid member_role (0x00 or 0x01). The first overlay member MUST be the base (member_role = 0x01) and MUST span the index space (all-zero shard_offset, shape equal to the head's shape). Every subsequent overlay member MUST have member_role = 0x00.
5. For partition and group: the member MUST NOT set HAS_COMPOSITE_MEMBER.

A per-member violation MUST cause rejection of the whole composite.

Close-time checks (at the Nth member)

• Partition: on receiving the Nth member, run the exact-cover and non-overlap checks (§ Composition Semantics › Partition) over the N boxes. A gap or an overlap MUST cause rejection.
• Sealed overlay: close at the Nth member. The base-span is already checked at the first member (per-member check 4); overlap between corrections is legal; there is no exact-cover requirement. A reader MUST reject an overlay head (composition_rule = 0x02) with member_count = 0x00000000: overlay requires a base member (§ Composition Semantics), which a zero-member composite cannot supply.
• Group: close at the Nth member. There is no coverage or overlap check (members MAY differ arbitrarily).

Torn composite

A torn composite — the stream or file ends before all N members have arrived, for any composition rule — is incomplete. A strict reader MUST reject it. A permissive reader MAY expose the arrived members as independent shard tensors but MUST NOT present the composite as complete.

Deferred

Note (non-normative): The following are reserved for a future ADR (versioned / open overlay and related work) and are not part of v1.0: the 0xFFFFFFFF open-composite member_count sentinel and its append-oriented membership rules; the member_version field (the Composite Member section's reserved padding holds room for it); time-travel reads; file append with footer regeneration; a tombstone member kind (member_role = 0x02, data-less, revert-to-base-within-box); explicit member identifiers for out-of-order random access; wall-clock timestamp versioning; an optional inline single-frame compaction of a partition composite (a future revival of this idea would need a fresh layout tag allocated from the reserved range 0x0C–0x3F; the old subpaving tag 0x06 is permanently reassigned to COO and is not available for it); and heterogeneous per-member device placement. See ADR-027 § Status and § Consequences.

Example

A sealed overlay (SpQR-style) with a float16 logical view of shape [4096, 4096], combine_op = 0x01 (replace), one base and one correction:

Head (layout_tag = 0x0B):
  shape            = [4096, 4096]
  type_tag         = float16
  buffer_count     = 0x00
  byte_offset      = 0x0000000000000000
  composition_rule = 0x02   (overlay)
  combine_op       = 0x01   (replace)
  member_count     = 2

Member 0 (base):
  shape            = [4096, 4096]
  type_tag         = int4    (stored), decodes to float16
  layout_tag       = 0x01    (row-major), per-block-affine quantization
  HAS_SHARD:  parent_shape = [4096, 4096], shard_offset = [0, 0]
  HAS_COMPOSITE_MEMBER:  member_role = 0x01  (base)

Member 1 (correction):
  shape            = [4096, 4096]
  type_tag         = float16
  layout_tag       = 0x06    (COO sparse: scattered outliers)
  HAS_SHARD:  parent_shape = [4096, 4096], shard_offset = [0, 0]
  HAS_COMPOSITE_MEMBER:  member_role = 0x00  (correction)

The merged logical view is: for each index, the correction's outlier value if present, otherwise the dequantized base value. The reader computes this merge; it is not zero-copy at the composite level.

A partition composite with a float32 logical view of shape [8, 8], split into two [8, 4] members:

Head: shape = [8, 8], type_tag = float32, buffer_count = 0,
      composition_rule = 0x01 (partition), combine_op = 0x00, member_count = 2

Member 0: shape = [8, 4], HAS_SHARD parent_shape = [8, 8], shard_offset = [0, 0]
Member 1: shape = [8, 4], HAS_SHARD parent_shape = [8, 8], shard_offset = [0, 4]

The two boxes exactly cover [8, 8] with no overlap; element [3, 6] resolves to member 1 at local index [3, 2]. The view is zero-copy.

Buffer Protocol — Hurray Format Specification

Status: Draft

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Scope

This section defines the buffer protocol: the rules governing how tensor data buffers and quantization-parameter buffers are represented, aligned, located on a device, owned, and released. It is the normative reference for all other sections that reference buffer handles, device tags, or alignment requirements.

The buffer protocol is independent of the interchange transport (in-process, IPC, or cross-machine); interchange.md defines the transport-level framing.

Buffer Handle

A buffer handle is the unit by which a tensor descriptor references a contiguous region of memory. Each handle appears as a 16-byte entry in the buffer table of the tensor descriptor (see metadata.md § Buffer Table).

All multi-byte fields MUST be encoded in little-endian byte order.

Note (non-normative): The buffer handle in the descriptor is a declaration of a buffer's properties, not a pointer. The actual pointer (or shared-memory handle, or RDMA registration) is communicated out-of-band via the interchange protocol or, for in-process use, via the C ABI defined in docs/impl/c-ffi.md.

Alignment

Minimum Alignment

The base address of every non-empty buffer MUST be aligned to at least 64 bytes. This ensures compatibility with all current SIMD instruction sets (AVX-512, NEON, SVE) without requiring per-operation alignment negotiation.

A writer MUST set alignment to the actual alignment it guarantees, which MUST be a power of two and MUST be at least 64. A writer MAY set alignment to a larger value (e.g., a page boundary of 4096 or 65536) to communicate a stronger guarantee.

A reader MAY rely on the declared alignment for SIMD loads. A reader MUST NOT rely on alignment stronger than what is declared in the alignment field.

Page Alignment for GPU and IPC

Buffers shared across process boundaries (IPC) or placed in device memory (GPU) SHOULD be aligned to the host page size, which is typically 4096 bytes. Writers targeting GPU or IPC transport MUST set alignment to at least 4096.

Buffers intended for RDMA transfer SHOULD be aligned to the RDMA provider's minimum pinnable unit, which is typically 4096 bytes. Writers targeting RDMA MUST set alignment to at least 4096.

Empty Buffers

A buffer with byte_size = 0 is an empty buffer. The 64-byte minimum alignment requirement does not apply to an empty buffer: there are no addressable bytes to align. A writer MAY set alignment to any power-of-two value (including 0x00000001) for an empty buffer. A reader MUST NOT dereference the pointer of an empty buffer.

In C ABI contexts, an empty buffer MAY be represented by a null pointer. A non-null pointer for an empty buffer is also valid; readers MUST handle both.

Device Tags

The device_tag field identifies the memory space in which the buffer resides.

A reader MUST reject a buffer handle whose device_tag is 0xFF.

Tags in the range 0x09–0xEF MUST NOT be used by any implementation; they are reserved for future specification versions.

Tags in the range 0xF0–0xFE MAY be used by implementations for private or experimental device types. Descriptors carrying private device tags MUST NOT be exchanged between independent implementations unless both parties have agreed on the semantics out of band.

Per-Device Memory Model

This subsection specifies, for each named device tag, the allocation context the buffer's base address refers to and any alignment requirement that applies in addition to the global 64-byte SIMD minimum defined in § Alignment. Where no device-specific alignment is required, only the 64-byte minimum (and the page-alignment SHOULD for GPU/IPC) applies.

0x00 CPU host memory

Buffers reside in the producer's process-addressable host memory, allocated by any standard host allocator (e.g., malloc, mmap, jemalloc, the system page allocator). No alignment requirement above the 64-byte SIMD minimum applies. Buffers shared via IPC SHOULD be page-aligned per § Page Alignment for GPU and IPC.

0x01 CUDA device memory

Buffers reside in CUDA device memory, typically allocated by cudaMalloc, cuMemAlloc, or cuMemCreate / cuMemMap on the device identified by the producer's CUDA context. The base address is a CUDA device pointer and MUST NOT be dereferenced from the host. Buffers intended for cross-process sharing or GPUDirect RDMA SHOULD be page-aligned (typically 4096 bytes); the CUDA driver returns allocations aligned to at least 256 bytes, which already satisfies the 64-byte SIMD minimum.

0x02 ROCm device memory

Buffers reside in AMD GPU device memory allocated through the HIP / ROCr runtime (e.g., hipMalloc, hsa_amd_memory_pool_allocate). The base address is a device pointer and MUST NOT be dereferenced from the host. The 64-byte minimum applies; page alignment SHOULD be used for IPC and RDMA per the global rules.

0x03 Metal device memory

Buffers reside in a Metal MTLBuffer whose storage mode is shared (Apple Silicon unified memory) or private. On Apple Silicon, unified memory permits host access to shared-mode buffers; private-mode buffers MUST NOT be dereferenced from the host. The base address conveyed via this tag is the buffer's contents pointer (for shared/managed storage) or the GPU resource handle (for private storage); the consumer MUST agree out of band on which storage mode the producer used. The 64-byte minimum applies.

0x04 Vulkan device memory

Buffers reside in a VkDeviceMemory allocation on the producer's Vulkan logical device, typically backing a VkBuffer. The base address is the result of vkMapMemory for host-visible memory, or an opaque device handle that MUST NOT be dereferenced from the host for device-local memory. Cross-process sharing requires an external memory handle (VkExternalMemoryHandleTypeFlagBits) exchanged out of band. The 64-byte minimum applies; producers SHOULD honour the device's nonCoherentAtomSize and minMemoryMapAlignment properties when applicable.

0x05 WebGPU device memory

Buffers reside in a WebGPU GPUBuffer allocated against the producer's GPUDevice. The base address is a host pointer only when the buffer was mapped via mapAsync with MAP_READ or MAP_WRITE; otherwise it is an opaque GPU resource handle that MUST NOT be dereferenced from the host. WebGPU imposes a 4-byte minimum for buffer offsets and copy sizes; the 64-byte SIMD minimum subsumes this. Cross-process sharing of WebGPU buffers is not defined by the W3C specification; implementations MUST exchange GPU resource handles out of band.

0x06 Qualcomm Hexagon (HVX/HMX) memory

Buffers reside in memory allocated via the Qualcomm AI Engine Direct (QNN) SDK or the Hexagon SDK's rpcmem / FastRPC allocator, accessible to the Hexagon DSP / NPU via the QNN HTP backend. The base address is a host pointer into the rpcmem region that the DSP can also address through its IOMMU. HVX operations are most efficient on 128-byte-aligned addresses; producers SHOULD align buffers to at least 128 bytes when targeting Hexagon, which exceeds the 64-byte SIMD minimum.

0x07 Intel Level Zero / oneAPI device memory

Buffers reside in memory allocated via the Intel Level Zero API (zeMemAllocDevice, zeMemAllocHost, zeMemAllocShared) on a Level Zero device, or equivalently via SYCL Unified Shared Memory (sycl::malloc_device, sycl::malloc_shared). The base address may or may not be host-dereferenceable depending on the allocation kind; consumers MUST agree on the kind out of band or query it via zeMemGetAllocProperties. The 64-byte minimum applies; the Level Zero spec returns allocations aligned to at least 64 bytes by default.

0x08 OpenCL device memory

Buffers reside in an OpenCL cl_mem allocation on the producer's OpenCL context and device. The base address is a host pointer only when the buffer was created with CL_MEM_USE_HOST_PTR / CL_MEM_ALLOC_HOST_PTR and is currently mapped via clEnqueueMapBuffer; otherwise it is an opaque device handle and MUST NOT be dereferenced from the host. The 64-byte minimum applies; producers SHOULD honour the device's CL_DEVICE_MEM_BASE_ADDR_ALIGN property when allocating sub-buffers.

0xF0–0xFE Implementation-private device types

Buffers carrying a private device tag MUST be exchanged only between peers that have agreed on the allocation context, alignment requirements, host-versus-device addressability, and synchronisation rules out of band. The format specification makes no statement about the memory model of private tags beyond the global buffer protocol invariants in this section.

Note (non-normative): The mapping between Hurray device tags and DLPack DLDeviceType constants is normative for Python bindings only and lives in docs/impl/python-bindings.md § Device Tag Mapping (Hurray ↔ DLPack). It is intentionally not duplicated here, since translation is the binding layer's responsibility, not a property of the buffer protocol.

Device Colocation

All buffers referenced by a single tensor descriptor (data buffer + all quantization-parameter buffers) MUST share the same device_tag AND the same memory_class. A reader MUST reject a descriptor whose buffers carry different device_tag or memory_class values.

For TENSOR_PUT transfers (see interchange.md), the client unilaterally declares the destination device_tag in the descriptor; the server MAY reject the transfer with DEVICE_UNAVAILABLE but MUST NOT silently place buffers on a different device.

Note (non-normative): Device colocation ensures that quantized tensor kernels can dereference both the data and the quantization parameters without triggering cross-device transfers. A writer that needs quantization parameters on a different device must emit a separate tensor descriptor.

When buffer handles are exchanged across machines, the device selection rules in interchange.md § Device Negotiation govern which device tag is valid for a given transfer.

Memory Class

The memory_class field identifies how a buffer is accessible — specifically, whether it can be read without copying by more than one compute unit simultaneously. memory_class is orthogonal to device_tag: the device tag names the allocator or hardware domain; the memory class names the access semantics within that domain.

Memory Class Values

A reader MUST reject a buffer handle whose memory_class value is in the range 0x04–0xEF or equals 0xFF.

Per-Device Validity

Not every (device_tag, memory_class) combination is meaningful. The following table defines the valid combinations. A reader MUST reject a buffer handle whose (device_tag, memory_class) pair is not listed as valid for the declared device. Private device tags (0xF0–0xFE) MAY be paired with any private memory class (0xF0–0xFE) or STANDARD (0x00); semantics are out of band.

Note (non-normative): Metal HOST_PINNED (MTLStorageModeManaged) is deprecated and unavailable on Apple Silicon. Producers targeting Apple Silicon MUST use UNIFIED (MTLStorageModeShared) instead. The HOST_PINNED value remains defined for discrete Metal GPU configurations.

Note (non-normative): ROCm UNIFIED requires hardware support for Heterogeneous Memory Management (HMM). Producers MUST verify hardware support before tagging a buffer UNIFIED; consumers MAY fall back to a copy-based path if UNIFIED is declared but the consumer's runtime does not support HMM on the current device.

Backward Compatibility

Existing descriptors that encode 0x00 in what was previously the first byte of _reserved[2] are implicitly STANDARD (memory_class = 0x00) — the most conservative and correct semantics for any pre-ADR-020 allocation. No existing producer or consumer is broken by this reassignment.

Readers compiled before this amendment will encounter a non-zero memory_class byte and reject it at the _reserved byte check. This is the intended fail-safe: a consumer that does not understand the memory class MUST NOT silently treat a UNIFIED buffer as STANDARD, as doing so would yield incorrect synchronisation.

Buffer Ownership and Lifetime

Ownership Model

At any instant, exactly one entity — the owner — is responsible for the buffer's memory. Ownership may be transferred between a producer and a consumer as part of the interchange protocol, but it is never shared: concurrent read/write access to the same buffer by multiple owners is a protocol error.

In-Process

In in-process exchange, the producer creates the buffer and holds ownership until the consumer signals that it has retained a reference (via the release callback mechanism described below). The consumer then owns the buffer for the duration of its use and MUST release it exactly once when done.

IPC

In IPC exchange via shared memory, the producer creates and owns the shared memory segment. The consumer maps the segment into its own address space. The producer MUST NOT unmap or destroy the segment until all consumers have unmapped it. The IPC channel MUST convey a release signal so that the producer knows when it may reclaim the segment.

Cross-Machine

In cross-machine exchange, the sender owns the source buffer and the receiver owns the destination buffer. There is no shared buffer; data is copied (or RDMA-written) from sender to receiver. See interchange.md for framing details.

Release Callback

For in-process and IPC exchange, the buffer handle is augmented at the ABI level with a release callback: a function pointer that the consumer calls exactly once when it has finished using the buffer. The release callback is not encoded in the binary descriptor; it is supplied by the producer at handoff time via the C ABI (see docs/impl/c-ffi.md).

A consumer MUST call the release callback exactly once. A consumer MUST NOT access the buffer after calling the release callback. A producer's release callback MUST be safe to call from any thread.

Reference Counting

Reference counting is an implementation detail — not a normative contract (see docs/adr/ADR-009-release-callback-not-normative-refcount.md). A producer that wishes to support multiple simultaneous consumers of the same buffer MUST implement reference counting internally. Each consumer receives a separate buffer handle whose release callback decrements the internal count; the actual deallocation occurs only when the count reaches zero. Consumers are unaware of this; they call their release callback exactly once as the normative contract requires.

Note (non-normative): This is the same model used by DLPack's DLManagedTensor.deleter. It keeps the ABI surface minimal and allows each language binding to use its own lifetime management idiom (Python GC, Rust Arc, etc.) without bridging to a C reference count.

Stream and Event Synchronisation

The sync_mode field at offset 13 of the buffer handle declares how the producer has ordered its device-side writes with respect to the moment of handoff. The release callback (see § Release Callback) governs the end of consumer access; sync_mode governs the start. The two contracts are independent and apply symmetrically.

sync_mode Values

A reader MUST reject a buffer handle whose sync_mode value is not one of the values defined for the format version it implements.

Producer Requirement

A producer of a non-CPU buffer (device_tag != 0x00) MUST ensure that, at the instant ownership of the buffer is transferred to the consumer, all device-side writes enqueued by the producer that affect the buffer's bytes have reached a point at which a properly-synchronised consumer access on the same device will observe them. The producer MUST satisfy this requirement by exactly one of the three mechanisms enumerated above, and the chosen mechanism MUST be declared in the buffer handle's sync_mode field.

For a CPU buffer (device_tag == 0x00), sync_mode MUST be SYNC_PRODUCER_SYNCED (0x00). When concurrent host-side writes exist, the producer MUST additionally issue a host memory fence (a release-store or equivalent) before handoff so that all preceding host writes are visible to the consumer's subsequent loads.

Consumer Requirement

A consumer that has received a buffer handle MUST inspect the sync_mode field and apply the matching rule before accessing the buffer's bytes:

• If sync_mode == SYNC_PRODUCER_SYNCED, the consumer MAY access the buffer immediately on any stream.
• If sync_mode == SYNC_EVENT, the consumer MUST retrieve the producer's event handle from the C ABI handoff structure and MUST issue a device-stream-wait on it on every stream that will access the buffer before enqueuing any work that touches the buffer. The consumer MUST release the event handle exactly once via the event-release callback.
• If sync_mode == SYNC_CONSUMER_STREAM, the consumer MAY access the buffer on the stream(s) it declared at handoff time, but MUST NOT access the buffer on any other stream until it has issued an inter-stream wait.

A consumer that does not recognise the declared sync_mode value MUST reject the descriptor.

Per-Transport Constraints

The set of sync_mode values that are valid depends on the interchange transport (see interchange.md):

• In-process. All three sync_mode values are valid. Event and stream handles are exchanged out of band via the C ABI; both parties share the same driver context.
• IPC (same machine, different processes). SYNC_PRODUCER_SYNCED is always valid. SYNC_EVENT is valid only if the device supports IPC-exportable events (e.g., CUDA cudaIpcEventHandle_t, ROCm hipIpcEventHandle_t). SYNC_CONSUMER_STREAM is valid only if the device supports IPC-exportable streams. When neither SYNC_EVENT nor SYNC_CONSUMER_STREAM is available on the underlying device, the producer MUST use SYNC_PRODUCER_SYNCED or fall back to a host-staged copy.
• Cross-machine (network transport). sync_mode MUST be SYNC_PRODUCER_SYNCED for every buffer handle transmitted over a network transport. SYNC_EVENT and SYNC_CONSUMER_STREAM are FORBIDDEN across machines because device event and stream handles are not valid in a different driver context on a different host. A receiver MUST reject a cross-machine TENSOR_DESCRIPTOR whose buffer handle declares any other mode. See interchange.md § RDMA Data Plane for how TENSOR_DATA_END serves as the cross-machine equivalent of SYNC_PRODUCER_SYNCED.

Relationship to the Release Callback

The sync_mode contract governs the start of consumer access; the release callback (see § Release Callback) governs the end of consumer access. The two are independent normative contracts that occupy symmetric positions at the bookends of the consumer's hold.

The event-release callback used in SYNC_EVENT mode is separate from the buffer-release callback. A consumer in SYNC_EVENT mode therefore makes two release calls per buffer: one for the event handle, called after the consumer has issued its stream-wait (typically immediately after handoff), and one for the buffer, called after all device work on the buffer is complete. Conflating the two would force the producer to keep the event alive for the buffer's entire lifetime, defeating the purpose of using events instead of full stream synchronisation.

A consumer that has issued device work using the buffer MUST NOT call the buffer-release callback until that device work has completed on the device. The consumer is free to satisfy this by host-side waiting, by recording its own completion event and waiting on it, or by deferring the release callback to a completion callback registered on its stream.

C ABI Note

The opaque event handle (for SYNC_EVENT) and the consumer stream handle (for SYNC_CONSUMER_STREAM) are NOT carried in the binary descriptor. They are exchanged out of band via the C ABI. See docs/impl/c-ffi.md for the per-mode handoff payload definitions and the ABI-side cross-check that the payload provided at handoff time matches the sync_mode declared in the descriptor.

Note (non-normative): sync_mode is a declaration of a buffer's synchronisation properties; the synchronisation handle is a transport detail, not a buffer property. Putting the discriminant in the binary descriptor lets a static inspector (hurray-inspect) surface the synchronisation contract for any buffer without reaching into the C ABI layer.

Zero-Copy Invariants

The buffer protocol is designed to preserve zero-copy access across language and runtime boundaries. The following invariants MUST hold at all times:

1. No implicit copies. Neither the producer nor the consumer MAY copy the buffer contents as part of the handoff. Copies are only permitted when explicitly requested by the interchange protocol (e.g., layout transcoding in response to a TENSOR_REQUEST that specifies a different layout tag).
2. No in-place mutation after handoff. Once a producer has handed off a buffer to a consumer, the producer MUST NOT modify the buffer's contents. Mutating a buffer that is held by a consumer is a protocol error.
3. Quantization parameter buffers are immutable. Scale and zero-point buffers MUST NOT be modified after the tensor descriptor is emitted. They are part of the tensor's logical value and MUST be treated as read-only by all consumers.
4. Pointer stability. The base address of a buffer MUST NOT change for the duration of the consumer's hold. Buffer defragmentation or garbage collection that moves the buffer is the producer's responsibility to prevent while any consumer holds a reference.

Relationship to Other Sections

• metadata.md defines the binary encoding of buffer handles in the buffer table. The alignment and device_tag fields are declared there; this file defines the normative rules they must satisfy.
• data-model.md defines empty tensors (zero-size dimensions). This file specifies the corresponding empty-buffer rules (null pointer allowed, alignment waived, no dereference).
• quantization.md defines quantization-parameter buffers. This file's device-colocation and immutability rules apply to those buffers.
• interchange.md defines in-process, IPC, and cross-machine transport. This file's alignment and ownership rules are prerequisites for all three transport modes.
• docs/impl/c-ffi.md defines the C ABI for buffer handle handoff, including the release callback signature.

Open Questions

All open questions in this section are resolved. See docs/adr/ADR-009-release-callback-not-normative-refcount.md (OQ-1) and docs/adr/ADR-020-memory-class-field.md (memory class field, device colocation extension, and supported_memory_classes interchange advertisement).

Metadata — Hurray Format Specification

Status: Draft

Scope

This section defines the binary encoding of the tensor descriptor: the self-describing header that precedes every tensor data buffer in the Hurray format. The tensor descriptor encodes all information required to interpret a tensor: its element type, rank, shape, memory layout, buffer table, and optional quantization and shard annotations.

Note (non-normative): The tensor descriptor is designed to be self-delimiting: a receiver can determine its total byte length from the first 10 bytes (the descriptor_length field occupies bytes 6–9) and MAY skip the descriptor entirely without parsing any layout-specific fields. Skipping past the whole descriptor is what only requires the first 10 bytes; locating or skipping a particular section within the descriptor (e.g. the buffer table or the quantization section) requires reading through the preceding fields. This property is essential for streaming readers.

Normative Requirements

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Overall Descriptor Structure

A tensor descriptor consists of a fixed header, followed by variable-length core fields, followed by layout-specific fields, followed by a buffer table, followed by zero or more optional sections selected by the flags field.

[Fixed Header]         20 bytes
[shape]                8 × rank bytes
[byte_offset]          8 bytes
[Layout-specific]      variable (layout-dependent)
[Buffer Table]         variable
[Quantization]         variable, present if HAS_QUANTIZATION flag is set
[Shard]                variable, present if HAS_SHARD flag is set
[Statistics]           72 bytes, present if HAS_STATISTICS flag is set
[Extension Type]       20 bytes, present if HAS_EXTENSION_TYPE flag is set
[Composite Member]     16 bytes, present if HAS_COMPOSITE_MEMBER flag is set

All multi-byte fields MUST be encoded in little-endian byte order (least significant byte at the lowest address).

Fixed Header

The fixed header occupies the first 20 bytes of every tensor descriptor.

A reader MUST reject a descriptor whose magic field is not 0x48 0x52 0x52 0x59.

A reader MUST reject a descriptor whose version_major exceeds the reader's supported major version.

A reader MUST reject a descriptor whose descriptor_length is less than 20 (the minimum valid descriptor size).

A reader MUST NOT read beyond descriptor_length bytes when parsing a descriptor.

Flags

Core Variable Fields

Immediately following the fixed header:

shape

uint64[rank] — the size of each dimension, in ascending dimension order (dimension 0 first). Each value MUST be greater than or equal to 0. A dimension size of 0 indicates an empty tensor.

The value 0xFFFFFFFFFFFFFFFF (UINT64_MAX) is the dynamic dimension sentinel: it indicates that the dimension's size is not statically known. A reader MUST NOT compute buffer sizes, strides, or element counts for a dimension carrying this sentinel without first resolving it to a concrete value.

For a scalar tensor (rank = 0), this field is absent (zero bytes).

byte_offset

uint64 — the byte offset from the start of buffer 0 in the buffer table to the element at logical index [0, 0, ..., 0]. MUST be less than or equal to the byte size of buffer 0.

For sub-byte types (bool, int4, uint4, int2, uint2), byte_offset MUST point to a byte boundary. The first element begins at bit 0 of that byte.

Layout-Specific Fields

Immediately following byte_offset, the layout-specific fields for the layout identified by layout_tag are encoded. If the layout has no additional fields, this section is absent.

Row-Major (0x01) and Column-Major (0x02)

No additional fields. Strides are implicit and computed as defined in memory-layout.md.

Strided (0x03)

Tiled / Blocked (0x04)

If outer_layout is 0x03 (strided):

If inner_layout is 0x03 (strided):

If inner_layout is 0x04 (recursive tiling), the tiled layout-specific fields are encoded recursively at this point, beginning with tile_shape.

A reader MUST enforce a maximum recursion depth for nested tiled descriptors. The RECOMMENDED limit is 8 levels. A reader MUST reject a descriptor that exceeds its configured recursion limit.

Morton (Z-Order Curve) (0x05)

COO (0x06)

See layouts/coo.md for buffer table composition, storage order, and validity constraints.

CSR (0x07)

See layouts/csr.md for buffer table composition, storage invariants, and rank-2 restriction.

CSC (0x08)

See layouts/csc.md for buffer table composition, storage invariants, and rank-2 restriction.

Hilbert Curve (0x40)

Composite / Virtual (0x0B)

The composite head is a virtual (data-less) descriptor. Its layout-specific fields encode the composition rule. All multi-byte fields are little-endian.

A composite head MUST have buffer_count = 0x00 and byte_offset = 0x0000000000000000. See layouts/composite.md for members, binding, composition semantics, and validation.

Extension Layouts (0xF0–0xFE)

A reader that does not recognise extension_layout_id MUST reject the descriptor, unless operating in permissive mode.

Buffer Table

Immediately following the layout-specific fields, the buffer table is encoded.

The maximum value is 255, imposed by the uint8 wire type. This limit applies to the sum of data and quantization-parameter buffers. Implementations that require more than 255 buffers MUST use multiple tensor descriptors.

Followed by buffer_count buffer handles, each encoded as 16 bytes:

The _reserved bytes MUST be 0x00. A conforming reader in strict mode MUST reject a descriptor containing any buffer handle whose _reserved bytes are not all 0x00.

Note (non-normative): For sparse layout tags (COO 0x06, CSR 0x07, CSC 0x08, CSF 0x09), buffer_count exceeds 1 — each entry holds a distinct component array (values, indices, pointers). For CSF the count is rank-dependent, 2·rank + 1 (one values buffer plus a pos/crd pair per level). For quantized dense tensors, quantization-parameter buffers (scales, zero-points) extend the buffer table beyond the layout baseline. The layout-defined minimum is always 0x01 for dense layouts; quantization schemes append their parameter buffers on top.

Quantization Section

Present if and only if the HAS_QUANTIZATION flag (bit 0) is set.

A reader that encounters HAS_QUANTIZATION but does not support quantized types MUST reject the descriptor unless operating in permissive mode.

Shard Section

Present if and only if the HAS_SHARD flag (bit 1) is set.

The constraint shard_offset[k] + shape[k] <= parent_shape[k] MUST hold for every dimension k. A reader MUST reject a shard descriptor that violates this constraint.

Statistics Section

Present if and only if the HAS_STATISTICS flag (bit 3) is set.

The statistics section is a fixed-size 72-byte block. All statistics are advisory: they reflect the tensor data at write time. A reader MUST NOT rely on any statistic for correctness; statistics MAY be used only as optimization hints (algorithm selection, memory pre-allocation, routing decisions).

Note (non-normative): A streaming writer that has not processed the entire tensor buffer before emitting the descriptor (e.g., a pipeline stage forwarding data on the fly) MUST omit the statistics section (HAS_STATISTICS not set) rather than emitting invalid statistics. A writer that knows only a subset of statistics (e.g., nnz is known from a sparse format but value_mean was not computed) MUST mark the unknown fields as not valid in computed_mask.

computed_mask

The computed_mask field (first 4 bytes of the section) is a bitmask indicating which statistics fields contain valid data. A reader MUST check the relevant bit before using any field. Fields whose bit is not set MUST be treated as unknown, regardless of their encoded value.

A conforming reader MUST reject a descriptor whose computed_mask has any reserved bit set (any bit greater than or equal to 6, given the six defined statistics fields above).

Field Encoding

The 72-byte statistics block is encoded as follows. All multi-byte fields are little-endian.

Note (non-normative): value_min, value_max, value_abs_max, value_mean, and value_stddev are always expressed in float64 regardless of the tensor's element type. For quantized tensors, these values reflect dequantized (real-valued) statistics, not the raw quantized storage values. For bool types, the statistics are defined over the integer domain {0, 1}.

Note (non-normative): N:M structured sparsity is particularly relevant for NVIDIA Ampere/Ada/Hopper Tensor Cores, which provide hardware-accelerated 2:4 sparsity (2 non-zeros in every group of 4 consecutive elements). Declaring the N:M pattern in the descriptor lets a receiver select the sparse kernel path without scanning the buffer.

Extension Type Section

Present if and only if the HAS_EXTENSION_TYPE flag (bit 2) is set. This flag MUST be set whenever type_tag is in the range 0xF0–0xFE, and MUST NOT be set otherwise.

Per ADR-001, extension type tags MUST carry an inline descriptor providing at minimum the bit width and packing parameters required to compute buffer sizes.

The extension type descriptor is 20 bytes:

Note (non-normative): The sign of a float extension type is carried by sign_bits; is_signed applies to integer types only. Requiring is_signed to be 0x00 for float types does not mean float extension types are unsigned — it means a reader has exactly one field to consult. A signed float sets is_float = 0x01 and sign_bits = 0x01; an unsigned float, such as a private analogue of the exponent-only float8_e8m0 (0x42), sets is_float = 0x01 and sign_bits = 0x00.

Sub-byte element widths that are not a power of two (notably 6-bit) are reserved to the built-in type tag space. Implementors requiring an interchange-portable non-power-of-two sub-byte type MUST request a built-in tag assignment through the specification governance process rather than encoding the type in the private extension range. The extension descriptor's whole-byte and power-of-two sub-byte width restriction ensures that buffer-size computation remains a single integer formula (ceil(N / packing_factor) for sub-byte, N * (bit_width / 8) for whole-byte) without rational arithmetic.

Note (non-normative): The 6-bit float6_e2m3 (0x44) and float6_e3m2 (0x45) types are built-in (Tier 2) and use a dedicated 4-elements-per-3-bytes packing defined in element-types.md. Their packing rule is not expressible as 8 / bit_width and is therefore not delegable to the generic extension descriptor. The extension descriptor is designed for private, implementation-defined types whose layout fits the simple "elements per byte" model; richer packings remain the prerogative of the standardized type system.

A reader MUST use bit_width and packing_factor to compute buffer sizes for tensors with extension type tags, even if it does not interpret the numeric semantics of the type.

Composite Member Section

Present if and only if the HAS_COMPOSITE_MEMBER flag (bit 4) is set. It appears after the Extension Type section. The section is a fixed-size 16-byte block. All multi-byte fields are little-endian.

A reader MUST reject a Composite Member section whose member_role is not 0x00 or 0x01, or whose _reserved bytes are not all 0x00.

This section is carried by the members of an overlay composite (composition_rule = 0x02). Partition and group members MUST NOT set HAS_COMPOSITE_MEMBER. The full semantics — base-vs-correction roles, precedence, and the combine operation — are defined in layouts/composite.md.

Note (non-normative): v1.0 uses member_role only. The 15 reserved bytes hold room for a future member_version field (versioned overlay), added additively without reallocating the section. See layouts/composite.md § Deferred.

Version Compatibility

A reader MUST reject a descriptor whose version_major exceeds the reader's supported major version.

A reader encountering a version_minor greater than its supported minor version SHOULD accept the descriptor but MUST NOT interpret fields beyond what its supported minor version defines. The descriptor_length field allows the reader to skip the descriptor entirely if desired.

Note (non-normative): Minor version increments add optional fields or new flag bits. A reader built against version 1.0 will correctly skip a 1.1 descriptor by consuming exactly descriptor_length bytes, because all reserved flag bits were required to be 0 in 1.0. Major version increments signal backward-incompatible changes; a reader MUST NOT attempt to parse a descriptor with an unsupported major version.

Worked Example

A rank-2 float32 tensor with shape [3, 4] in row-major layout, one buffer of 192 bytes aligned to 64 bytes on CPU, no optional sections:

Offset  Value (hex)                   Field
------  ----------------------------  -----
0       48 52 52 59                   magic = "HRRY"
4       01                            version_major = 1
5       00                            version_minor = 0
6       3D 00 00 00                   descriptor_length = 61
10      00 00 00 00                   flags = 0x00000000 (no optional sections)
14      03                            type_tag = 0x03 (float32)
15      01                            layout_tag = 0x01 (row-major)
16      02 00 00 00                   rank = 2
20      03 00 00 00 00 00 00 00       shape[0] = 3
28      04 00 00 00 00 00 00 00       shape[1] = 4
36      00 00 00 00 00 00 00 00       byte_offset = 0
                                      (no layout-specific fields for row-major)
44      01                            buffer_count = 1
45      C0 00 00 00 00 00 00 00       buffer[0].byte_size = 192
53      40 00 00 00                   buffer[0].alignment = 64
57      00                            buffer[0].device_tag = 0x00 (CPU)
58      00 00 00                      buffer[0]._reserved

Total: 61 bytes. Fixed header (20) + shape (16) + byte_offset (8) + buffer table (1 + 16) = 61.

Interaction with Other Sections

• Element Types (element-types.md): defines type_tag values and the bit-width, packing, and alignment properties used during buffer size computation.
• Memory Layout (memory-layout.md): defines layout_tag values and the layout-specific fields encoded in this descriptor.
• Quantization (quantization.md): defines the binary format of the quantization_descriptor payload in the quantization section.
• Buffer Protocol (buffer-protocol.md): defines buffer ownership, device memory semantics, and release callback conventions referenced by the buffer table entries.
• Interchange (interchange.md): the tensor descriptor defined here is transmitted verbatim in TENSOR_DESCRIPTOR and TENSOR_PUT message payloads, followed by transport-specific fields (total_data_bytes, shard_index, total_shards).

Open Questions

[OQ-1]: Should the descriptor include a CRC-32 checksum field? Resolved: No checksum in the descriptor. Integrity is delegated to the transport/storage layer (TCP, TLS, ECC, ZFS). Adding 4 bytes and a full-pass CRC on every descriptor would penalise in-process and IPC interchange where corruption is not a realistic threat. If file-level integrity is needed, it belongs in the file format footer (see file-format.md OQ-3).

[OQ-2]: The binary encoding of the quantization descriptor is deferred to quantization.md. Resolved: The encoding is fully defined in quantization.md: a fixed 4-byte header (scheme_tag, scheme_version, flags) followed by a per-scheme payload. Complete byte-level layouts are specified in quantization/per-tensor-affine.md, quantization/per-channel-affine.md, quantization/per-block-affine.md, quantization/nf4.md, and quantization/mxfp.md. The quantization_length prefix in metadata.md allows readers to skip unrecognised schemes safely.

Interchange — Hurray Format Specification

Status: Draft

This section uses RFC 2119 key words: MUST, MUST NOT, REQUIRED, SHALL, SHALL NOT, SHOULD, SHOULD NOT, RECOMMENDED, MAY, and OPTIONAL.

Scope

This section defines how Hurray tensors are exchanged between producers and consumers. Three interchange modes are in scope:

This section focuses primarily on the network transport mode, as it is the most complex and the most relevant to distributed inference pipelines. In-process and IPC modes are covered in In-Process and IPC.

A Hurray stream MAY contain zero or more tensors. Back-to-back concatenation of self-delimiting descriptor+data pairs is the canonical multi-tensor encoding: a reader processes one tensor at a time, advancing to the next descriptor after the current tensor's data is consumed. No container header, index, or tensor names are defined at this level (see docs/adr/ADR-010-multi-tensor-collections-deferred.md).

Note (non-normative): The design of the Hurray network transport protocol draws inspiration from Apache Arrow Flight, but differs in several key respects: it is tensor-focused (not columnar), it supports layout and device negotiation, it defines on-the-fly transcoding, and its data plane is designed for large buffer transfers where gRPC framing overhead is impractical.

In-Process and IPC

In-Process

Within a single address space, tensor interchange is accomplished by passing a tensor descriptor (see metadata.md) and a buffer handle (see buffer-protocol.md) by value. No serialization is required. The buffer handle carries a release callback (see buffer-protocol.md § Buffer Ownership and Lifetime and ADR-009); the receiver MUST retain the handle for the duration of its use and call the release callback exactly once when done.

IPC

For cross-process interchange on a single machine, two mechanisms are supported:

1. Shared memory: the producer maps a shared memory segment and places the data buffer there. The tensor descriptor is transmitted over any IPC channel (pipe, Unix domain socket, etc.). The byte_offset field in the descriptor identifies the buffer's position within the shared segment.
2. Unix domain socket streaming: the producer serializes the tensor using the network transport framing defined below, transmitted over a Unix domain socket. This is slower than shared memory but requires no shared-memory setup.

In both cases, buffer alignment requirements from buffer-protocol.md apply.

Network Transport Protocol

Overview

The Hurray network transport protocol is a client-server streaming protocol for transferring tensor data over a reliable, ordered byte stream (e.g., TCP). It consists of:

• A control plane: session establishment, capability negotiation, tensor requests, and error signalling. Uses message framing defined in this section.
• A data plane: tensor descriptor and buffer transmission. Uses the same message framing, but data frame payloads MAY be transferred over a separate high-throughput channel (e.g., RDMA) by prior agreement during session establishment.

The protocol is half-duplex per stream: within a single stream, the client sends a request and the server responds with a sequence of messages. Multiple streams MAY be multiplexed over a single connection.

Note (non-normative): The stream_id field is an opaque per-stream identifier. Implementations MAY use one TCP connection per stream, or MAY multiplex multiple streams over a single connection by demultiplexing on stream_id. Messages are self-framing (fixed-size header + payload_length bytes), so a receiver can always advance to the next message regardless of stream_id. Normative multiplexing rules are not defined in this version of the spec.

Message Framing

Every message on the wire consists of a message header followed by a payload. All multi-byte fields are little-endian.

Message header (12 bytes):

The payload follows immediately after the header. A receiver MUST read exactly payload_length bytes as the payload. A receiver MUST reject messages whose payload_length exceeds the receiver's configured maximum message size, and MUST send an ERROR message in response.

Message Types

A receiver that encounters an unrecognised message_type MUST send an ERROR message and close the stream.

Message type values in the range 0x000000F0–0x000000FE are reserved for implementation-private extensions. Implementations MAY use these values for private extension message types; messages using private extension types MUST NOT be sent to a peer that has not agreed to the extension out of band. The value 0x000000FF is reserved and MUST NOT be used.

Session Establishment

CLIENT_HELLO Payload

SERVER_HELLO Payload

Note (non-normative): supported_memory_classes advertises which memory access classes a peer can produce or consume across all devices. A server that supports CUDA UNIFIED (cudaMallocManaged) and CPU STANDARD would list [0x00, 0x02]. A client that cannot handle UNIFIED buffers lists only [0x00]; the server MUST NOT send a UNIFIED buffer to such a client. Memory class negotiation follows the same fallback logic as device negotiation: the server SHOULD prefer the client's advertised classes and MUST report the actual memory_class in the TENSOR_DESCRIPTOR buffer handle.

A client MUST send a CLIENT_HELLO as the first message on every new connection. The server MUST respond with a SERVER_HELLO before any other message. If the server cannot satisfy the minimum requirements (e.g., protocol version mismatch), it MUST respond with an ERROR message instead and close the connection.

Capability Flags

Device Tags

The canonical definition of device tags lives in buffer-protocol.md § Device Tags. The table below is reproduced for transport-protocol convenience and MUST stay consistent with buffer-protocol.md.