ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-30 11:47:48 +00:00
Code Issues Packages Projects Releases Wiki Activity

docs(fmha): explain the difference between fp8, fp8bf16 and fp8fp32 precision modes

Browse Source 
Agent-Logs-Url: https://github.com/ROCm/composable_kernel/sessions/1a0f37b1-932d-4205-9cce-a9427ed9adc2

Co-authored-by: asleepzzz <4926646+asleepzzz@users.noreply.github.com>
This commit is contained in:
copilot-swe-agent[bot]
2026-03-30 14:43:36 +00:00
parent a809502aaa
commit 6004d045ea
1 changed files with 7 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
example/ck_tile/01_fmha/README.md
View File

@@ -165,7 +165,13 @@ We support sequence padding and variable-length processing in both batch and gro
Both approaches optimize memory access patterns while supporting flexible sequence length requirements commonly found in transformer inference scenarios.
## FP8 support
FP8 FMHA kernels are supported on gfx942/gfx950 machines with ROCm 6.0+. You can select fp8 precision by setting the arg `-prec=fp8` (or `fp8bf16`, `fp8fp32`) to the `tile_example_fmha_fwd`.
FP8 FMHA kernels are supported on gfx942/gfx950 machines with ROCm 6.0+. Three fp8-based precision modes are available via `-prec`:
| `-prec` value | Q/K/V input type | Output type | Description |
|---|---|---|---|
| `fp8` | fp8 | fp8 | Fully fp8: both inputs and output are in fp8 |
| `fp8bf16` | fp8 | bf16 | Mixed precision: fp8 inputs, bf16 output — useful when the consumer expects a wider-range output format |
| `fp8fp32` | fp8 | fp32 | Mixed precision: fp8 inputs, fp32 output — highest-precision output, suitable for debugging or further fp32 processing |
The following quantization scale modes are available via `-qscale`: