rocWMMA precision support#
This topic lists the data type support for the rocWMMA library on AMD GPUs.
This page lists the data types supported by the library itself and does not indicate hardware support. A type listed here is only usable if the GPU architecture also supports it; otherwise it is unsupported. For data type support across the other ROCm libraries and by GPU architecture, see the Data types and precision support page.
Supported data types overview#
The following table summarizes the input and output data types supported by rocWMMA. For the supported mixed-precision multiply-accumulate combinations and their per-architecture availability, see the section that follows.
Icon |
Definition |
|---|---|
✅ |
Fully supported as both an input and output type. |
⚠️ |
Partially supported as an input or output type. |
Data types not listed in the table below are not supported.
Data type |
Support |
|---|---|
int8 |
✅ |
int32 |
Output only |
float8 (E4M3) |
Input only |
float8 (E5M2) |
Input only |
float16 |
✅ |
bfloat16 |
✅ |
tensorfloat32 |
✅ |
float32 |
✅ |
float64 |
✅ |
Supported data type combinations#
rocWMMA mixed precision multiply-accumulate operations support the following data type combinations.
Data Types <Ti / To / Tc> = <Input type / Output Type / Compute Type>, where:
Input Type = Matrix A / B
Output Type = Matrix C / D
Compute Type = Math / Accumulation type
Supported data types:
i8: 8-bit precision integer
f8: 8-bit precision floating point
bf8: 8-bit precision brain floating point
f16: half-precision floating point
bf16: half-precision brain floating point
f32: single-precision floating point
i32: 32-bit precision integer
xf32: single-precision tensor floating point
f64: double-precision floating point
Note
f16 includes support for both _Float16 and __half types.
f8 NANOO (optimized) format is only supported on gfx942, otherwise f8 OCP is assumed on targets that support f8 datatypes.
f8(bf8) represents any one of mixed Matrix A / B input type: f8 / f8, f8 / bf8, bf8 / f8, bf8 / bf8.
Ti / To / Tc |
BlockM |
BlockN |
BlockK Range* (Powers of 2) |
CDNA Support |
RDNA Support |
|---|---|---|---|---|---|
bf8 / f32 / f32 |
16 |
16 |
32+ |
gfx942, gfx950, gfx1250 |
gfx1200, gfx1201 |
32 |
32 |
16+ |
- |
||
f8 / f32 / f32 |
16 |
16 |
32+ |
gfx942, gfx950, gfx1250 |
gfx1200, gfx1201 |
32 |
32 |
16+ |
- |
||
i8 / i32 / i32 |
16 |
16 |
16 |
gfx908, gfx90a, gfx1250 |
gfx11, gfx1200, gfx1201 |
32 |
gfx942, gfx950, gfx1250 |
- |
|||
64+ |
gfx950, gfx1250 |
- |
|||
32 |
32 |
8 |
gfx908, gfx90a |
- |
|
16 |
gfx942, gfx950 |
- |
|||
32+ |
gfx950 |
- |
|||
i8 / i8 / i32 |
16 |
16 |
16 |
gfx908, gfx90a, gfx1250 |
gfx11, gfx1200, gfx1201 |
32 |
gfx942, gfx950, gfx1250 |
- |
|||
64+ |
gfx950, gfx1250 |
- |
|||
32 |
32 |
8 |
gfx908, gfx90a |
- |
|
16 |
gfx942, gfx950 |
- |
|||
32+ |
gfx950 |
- |
|||
f8(bf8) / f16 / f16 |
16 |
16 |
16 |
gfx1250 |
- |
32 |
gfx1250 |
- |
|||
64 |
gfx1250 |
- |
|||
f8(bf8) / f32 / f32 |
16 |
16 |
16 |
gfx1250 |
gfx1200, gfx1201 |
32 |
gfx1250 |
- |
|||
64 |
gfx1250 |
- |
|||
f16 / f32 / f32 |
16 |
16 |
16 |
gfx9 |
gfx11, gfx1200, gfx1201 |
32+ |
gfx950 |
- |
|||
32 |
32 |
8 |
gfx9 |
- |
|
16+ |
gfx950 |
- |
|||
f16 / f16 / f32 |
16 |
16 |
16 |
gfx9, gfx1250 |
gfx11, gfx1200, gfx1201 |
32+ |
gfx950, gfx1250 |
- |
|||
32 |
32 |
8 |
gfx9 |
- |
|
16+ |
gfx950 |
- |
|||
f16 / f16 / f16** |
16 |
16 |
16 |
gfx9, gfx1250 |
gfx11, gfx1200, gfx1201 |
32+ |
gfx950, gfx1250 |
- |
|||
32 |
32 |
8 |
gfx9 |
- |
|
16+ |
gfx950 |
- |
|||
bf16 / f32 / f32 |
16 |
16 |
8 |
gfx908 |
- |
16 |
gfx90a, gfx942, gfx950 |
gfx11, gfx1200, gfx1201 |
|||
32+ |
gfx950 |
- |
|||
32 |
32 |
4+ |
gfx908 |
- |
|
8 |
gfx90a, gfx942, gfx950 |
- |
|||
16+ |
gfx950 |
- |
|||
bf16 / bf16 / f32 |
16 |
16 |
8 |
gfx908 |
- |
16 |
gfx90a, gfx942, gfx950, gfx1250 |
gfx11, gfx1200, gfx1201 |
|||
32+ |
gfx950, gfx1250 |
- |
|||
32 |
32 |
4+ |
gfx908 |
- |
|
8 |
gfx90a, gfx942, gfx950 |
- |
|||
16+ |
gfx950 |
- |
|||
bf16 / bf16 / bf16** |
16 |
16 |
8 |
gfx908 |
- |
16 |
gfx90a, gfx942, gfx950, gfx1250 |
gfx11, gfx1200, gfx1201 |
|||
32+ |
gfx950, gfx1250 |
- |
|||
32 |
32 |
4+ |
gfx908 |
- |
|
8 |
gfx90a, gfx942, gfx950 |
- |
|||
16+ |
gfx950 |
- |
|||
f32 / f32 / f32 |
16 |
16 |
4+ |
gfx9, gfx1250 |
- |
32 |
32 |
2+ |
gfx9 |
- |
|
xf32 / xf32 / xf32 |
16 |
16 |
8+ |
gfx942 |
- |
32 |
32 |
4+ |
|||
f64 / f64 / f64 |
16 |
16 |
4+ |
gfx90a, gfx942, gfx950 |
- |
Note
BlockM/N values are minimum recommended values. Below these values padding is used which may impact performance. Above this value powers of 2 are acceptable.
* BlockK range specifies the minimum recommended value. Below this value padding is used which may impact performance. Above this value powers of 2 are acceptable. In practice, BlockK values are typically 32 or less.
** On CDNA architectures, matrix unit accumulation is performed in natively 32-bit precision and then converted to the target data type.
Note
rocWMMA supports partial fragment sizes where FragMNK may be smaller than the BlockMNK sizes listed in the table above. These fragments are internally padded to nearest supported BlockMNK sizes.