Consider optional FP8 operand formats for subgroup matrices #11051

Summary

I propose extending the subgroup-matrix proposal with optional E4M3FN FP8 matrix operands and FP32 accumulation, including implementations that decode FP8 to a wider format rather than multiply FP8 natively.

This builds on #4195 and the D3D12 mapping work in #10964. The current proposed component types are f32, f16, u32, i32, u8, and i8; FP8 is not included. Subgroup matrices themselves are still proposal/experimental work, not a finalized core-WebGPU feature.

The concrete API/WGSL changes below are a starting design I am comfortable prototyping, not a request for the WG to invent the API. Names and scope remain open to review.

Motivation and scope

FP8 can reduce operand download size, GPU storage, and bandwidth even on hardware without native FP8 multiplication. Converting an entire operand to FP16 first can lose the resident-storage benefit. A packed-load interface could instead decode only the tiles being consumed.

Whether conversion overhead makes this worthwhile needs measurement. Native instructions are one implementation strategy, not a prerequisite for a useful storage format. No measured WebGPU performance gain is claimed here.

Initially support E4M3FN on both A and B, with FP32 accumulator/result. Defer E5M2, FNUZ encodings, mixed A/B formats, FP8 results, and general-purpose FP8 scalar arithmetic. Ordinary shader code handles application scales; matrix load/multiply must not silently quantize inputs or apply hidden scales.

Proposed WebGPU API changes

1. Add an optional GPUFeatureName, provisionally "subgroup-matrix-f8", requiring "subgroup-matrix" and its "subgroups" dependency.
   Internal FP16 widening does not by itself require the application to request "shader-f16" or enable WGSL f16; use a conforming alternative lowering if the backend lacks the necessary internal FP16 support.

2. Add "f8e4m3fn" to GPUSubgroupMatrixComponentType.

3. Reuse GPUAdapterInfo.subgroupMatrixConfigs and GPUSubgroupMatrixConfig. An initial FP8 record has:

   componentType       = "f8e4m3fn"
   resultComponentType = "f32"
   M, N, K             = a supported tile
   

   Carry over the subgroup-size constraints under discussion in [subgroup-matrix] Update proposal for D3D12 #10964. Both operands use the same component type, so this initial scope does not need separate left/right-type fields.

4. Advertise the feature only if at least one complete, validated FP8 configuration is supported. Enumerate actual supported combinations, not a cross-product of independently supported types and dimensions.

5. Permit a configuration implemented by native FP8, widening to FP16 with FP32 accumulation, or a tested software matrix lowering. A configuration promises functionality/numerical semantics, not native acceleration or a speedup. Unsupported adapters need not expose the feature.

6. Reject FP8 shader use when the device feature is not enabled. Reject a multiply at pipeline creation if the operand/result types, dimensions, or subgroup constraints do not match an exposed configuration.

I do not propose a native/emulated bit in the initial API. Its usefulness and fingerprinting implications can be discussed separately; neither the feature name nor configuration enumeration should imply native hardware.

Proposed WGSL changes

Enable, types, and operations

• Add a provisional enable subgroup_matrix_f8;, requiring the base subgroup-matrix enable and the device feature above.
• Introduce f8e4m3fn as a restricted matrix-component identifier, allowed in subgroup_matrix_left and subgroup_matrix_right only. It is not an ordinary scalar/vector/array element type; packed data stays in existing u32 buffers.
• Treat its logical component size as one byte for matrix-storage calculations. Keep the existing matrix dimension ordering and function-local lifetime rules.
• Extend subgroupMatrixMultiply and subgroupMatrixMultiplyAccumulate to accept matching E4M3FN left/right matrices and an f32 result/accumulator for an exposed configuration. Existing f32 result construction and subgroupMatrixStore remain unchanged.
• FP8 operand matrices are load-only in this first version. Do not add FP8 matrix stores, scalar constructors, scalar add/multiply, or implicit float-to-FP8 conversion. Such operations would need a separate conversion/rounding contract.

Packed load contract

Extend subgroupMatrixLoad<T, Majorness> when T has component type f8e4m3fn. Initially accept pointers to array<u32> in storage or workgroup address space, with the existing readable access modes.

Keep offset and stride in array elements, as in the existing proposal. For array<u32>, they count 32-bit words, not bytes or FP8 elements:

• Byte lane j of a word is (word >> (8 * j)) & 255, with j from 0 through 3. This fixes the packing interpretation independently of host byte order.
• For row-major loads, logical component [r,c] comes from word offset + r * stride + floor(c / 4), byte lane c % 4. Swap row/column roles for column-major.
• Each major row/column starts at a word boundary. Its final word may contain padding; padding bytes do not become matrix elements.
• Require stride >= ceil(minorDimension / 4). The addressed word span is offset + (majorDimension - 1) * stride + ceil(minorDimension / 4), with checked arithmetic.
• Adapt the proposal's shader/pipeline validation, minimum-binding-size, and robust-access rules to that padded word span. A prototype must not use the smaller logical M * N byte count as the physical bound or read outside a binding to obtain padding.
• Retain the base uniform-control-flow and pointer/offset/stride uniformity requirements. Implementations introducing workgroup staging must obey those requirements and must not introduce unsafe divergent barriers.

This is a restricted extension of the existing load builtin, not a second matrix API. Other component types keep their existing load rules.

Numerical contract

Interpret every byte as the standard E4M3FN encoding: one sign bit, four exponent bits, three mantissa bits, exponent bias 7, signed zero, subnormals, maximum finite magnitude 448, and NaN encodings with exponent 15 and mantissa 7. Exponent-15 encodings with smaller mantissas are finite. There is no infinity encoding, and FNUZ bytes must not be reinterpreted as E4M3FN.

Finite E4M3FN values are exactly representable in FP16 and FP32. Require exact finite decoding, including FP8 subnormals, before multiplication. Decode must not prematurely flush an FP8 subnormal or reinterpret it as an integer. NaN payload/sign preservation is not required.

Use the base proposal's f32 matrix-accumulation accuracy model; do not permit an FP16 accumulation chain merely because operands were widened to FP16. Reduction order need not be bitwise identical across backends. The base proposal currently leaves arithmetic ULP bounds undefined, so agreeing on useful numerical/conformance bounds remains a WG question. No FP8 re-rounding occurs during multiply.

Device availability and Qualcomm follow-up

Hardware instructions, API support, compiler/driver support, and an executed shader are different evidence levels.

• NVIDIA: NVIDIA documents FP8 Tensor Core operations on Hopper. That is hardware/library evidence, not proof of Vulkan or D3D12 availability on every NVIDIA adapter.
• AMD: AMD's precision reference documents FP8 Matrix Cores on CDNA3, CDNA4, and RDNA4, including MI300X/MI350X/Radeon RX 9070 XT families. CDNA3's FNUZ encoding is distinct from this proposal. RDNA4 discrete support does not qualify older AMD integrated devices; RDNA2/RDNA3 lack FP8 Matrix Core support in that reference.
• Intel: Intel's compiler platform definitions gate hasFP8Dpas() on Xe3P and newer cores, while DPAS documentation describes BF8/HF8 operands. These do not establish a qualifying FP8 graphics path on a particular integrated Arc GPU.
• Apple: the base proposal describes non-FP8 simdgroup matrices on Apple7-family and newer devices, including M1-class Macs. A native FP8 MSL simdgroup operation is not established here. simdgroup_float8x8 means an 8-by-8 matrix of float components, not FP8.

Qualcomm: matrix support is a separate question from FP8 support

Qualcomm's public Adreno cooperative-matrix sample queries VK_KHR_cooperative_matrix configurations and documents FP32, FP16, and integer workloads on supported Adreno GPUs. The sample explicitly rejects platforms without that extension.

Qualcomm also authored VK_QCOM_cooperative_matrix_conversion and SPV_QCOM_cooperative_matrix_conversion. The Vulkan design document describes avoiding shared-memory staging and permits FP32/FP16/INT8/UINT8 cooperative-matrix operands.

These sources establish a Qualcomm cooperative-matrix API path on supported devices, but do not establish native matrix instructions on a named Adreno SKU, availability on Snapdragon X Elite, or FP8 multiplication. A datatype-name case in sample code is not evidence of that datatype executing.

I therefore cannot substantiate a blanket claim that all Qualcomm GPUs lack subgroup-matrix support, nor claim native FP8 support. Qualcomm confirmation of device generations, hardware versus emulated execution, supported formats, and Android/Windows driver versions would resolve this. Hexagon NPU capabilities are not Adreno GPU evidence.

For Intel, AMD, and Qualcomm integrated devices without native FP8, emulation is now explicitly a candidate rather than automatic exclusion. A device without any matrix primitive may need a subgroup-shuffle/vector-arithmetic implementation, whose feasibility and benefit remain untested.

Backend mappings and emulation

D3D12 / HLSL

HLSL proposal 0035 is accepted, while the D3D12 runtime design remains labeled a draft. It includes E4M3FN/E5M2 datatypes and support flags describing emulation.

Important scope distinction: its Tier 1 mandatory FP8 input emulation is in the thread-scope vector-matrix contract. EMULATED_INPUTS is documented for that scope; the required FP8 rows have FP16 results. This is not automatically a guarantee of FP8 wave-matrix operations with FP32 accumulation. Its documented higher-precision vector-input behavior must also not silently change the FP8-by-FP8 semantics proposed here.

Investigate native linalg::Matrix wave configurations where available. Otherwise prototype decoding packed bytes to FP16 tiles and using a supported FP16-input/FP32-accumulator wave configuration, or a suitable software lowering. Query exact operations and inspect the compiler/driver path; do not infer support from Tier 1 alone. This complements #10964 rather than replacing its mapping.

Vulkan / SPIR-V

VK_EXT_shader_float8 and approved SPV_EXT_float8 provide native-format mapping targets alongside VK_KHR_cooperative_matrix. Check shaderFloat8CooperativeMatrix and exact E4M3 A/B, FP32 C/result, dimensions, and scope.

VK_KHR_cooperative_matrix enumerates implementation-supported configurations; it does not promise every FP8 configuration or offer a D3D12-equivalent mandatory FP8-emulation contract.

Without float8 support, a compiler-side prototype can read packed u32 bytes, decode into FP16 tiles, stage them safely if required, and use supported FP16 cooperative matrices with FP32 accumulation. The emitted SPIR-V then uses supported wider types, not undeclared FP8 capabilities. Driver-emulated configurations, if exposed, are another possibility, but support enumeration alone must not be called native-instruction proof. Where matrix primitives are absent, separately investigate subgroup arithmetic lowering.

Metal / MSL

The base subgroup-matrix mapping documents FP16/FP32 simdgroup matrices, not native FP8.

Prototype packed-byte decoding and bounded tile staging into supported MSL matrices with FP32 accumulation, checking legal operand/accumulator combinations rather than assuming mixed precision is available. If necessary, evaluate FP32 simdgroup operands or subgroup arithmetic instead.

No public driver-level guarantee equivalent to D3D12's required FP8 vector-matrix emulation is established here. Compiler-side widening is a candidate, not a measured Metal implementation. A separate tensor/Neural Engine API does not by itself supply the WGSL compute-shader mapping.

Prototype and evaluation plan

Compare matched, already-quantized operands and identical shapes/results across:

1. Resident FP16 operands with FP32 accumulation.
2. Resident packed FP8, decoding only the consumed tiles to a wider type.
3. Packed FP8 expanded once into a persistent full FP16 buffer.
4. Native FP8 where a qualifying API/toolchain exists.

Include a non-native-FP8 integrated GPU where accessible, and investigate a no-matrix-primitive path separately. Keep native, driver-emulated, and compiler-emulated measurements labeled.

Measure serialized/download bytes, uploads, persistent GPU operand bytes, peak GPU memory including temporary tiles, and output/accumulator buffers separately. Also measure packing/decoding, compilation, execution, and end-to-end latency across large GEMMs and small matrix/vector-like cases. Include enough repetitions to report variability and avoid treating smaller storage as an automatic throughput gain.

The theoretical operand payload is one byte per FP8 value versus two for FP16, before padding. This is up to 50% fewer operand payload bytes, not a measured 50% reduction in total application memory. The expand-once path may save download size without saving resident GPU memory.

Numerical tests should exhaust all 256 byte encodings, then cover finite extrema, NaNs, signed zero, subnormals, accumulation stress, tails, row/column-major packing, and native/emulated comparisons. Validation tests should cover disabled features, unsupported configurations, binding bounds/overflow, and nonuniform control flow.

The prototype should determine whether storage savings remain substantial after padding and scratch, and whether the latency/complexity tradeoff is useful enough to expose to web developers on emulated paths. No decision threshold or benefit is asserted without evidence.

Review and implementation process

WG review is needed before this becomes a finalized specification. It need not prevent an isolated experimental Dawn/Tint prototype or a work-in-progress CL, subject to Dawn maintainers' direction and contribution requirements. Such work would not imply WG acceptance, and must be prepared to change as the API evolves.

In particular, I would welcome review of the proposed enable/type names, packed-load rules, same-format/FP32 initial scope, numerical requirements, and emulated capability policy. Which Qualcomm devices have suitable matrix hardware/API support is also an explicit open question. If issue discussion leaves concerns unresolved, bring them to a WebGPU WG meeting.