Repository navigation
[Common] Optimize grouped MXFP8 requantization - #3560
Conversation
[Medium impact] The change reaches prequantized MXFP8 inputs used by grouped linear and grouped MLP workflows. Fix the shared descriptor race before merging so varying-width calls can run safely on separate CUDA streams. Findings
SummaryReplaces the earlier fused path with grouped CUDA requantization and an optional CuTeDSL backend.
Diagram%%{init: {'theme': 'neutral'}}%%
flowchart TD
A["nvte_group_requantize"] --> B["Check grouped input and output"]
B --> C{"CuTeDSL supports this call?"}
C -->|Yes| D["Run CuTeDSL kernel"]
C -->|No| E{"Member widths vary?"}
E -->|No| F["Run CUDA with call-local descriptors"]
E -->|Yes| G["Update shared g_tensor_maps"]
G --> H["Run CUDA using shared descriptors"]
D --> I["Columnwise MXFP8 and optional BF16 output"]
F --> I
H --> I
Reviews (13) · Last reviewed commit: "Merge branch 'main' into pr_requantize_m..." · Reviewed by Greptile |
|
@Oleg-Goncharov This PR seems to have some unrelated changes (like the nccl ep submodule). Could you fix that? |
de60011 to
b18d27e
Compare
2ecaf5f to
1e63df7
Compare
|
/te-ci |
|
@Oleg-Goncharov Please sign every commit in this PR. See https://github.com/NVIDIA/TransformerEngine/blob/main/CONTRIBUTING.rst#sign-your-work |
Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com>
1e63df7 to
ba75d00
Compare
Keep the replacement grouped MXFP8 requantization implementation and retain the new CUDA sources from main. Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
|
Hi @Oleg-Goncharov, what is the current state of this PR? Could we prioritize merging it soon? |
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
|
/te-ci |
|
Moved to TE2.22 |
|
/te-ci |
| update_tma_descriptors<IType, OType><<<num_tensors, THREADS_PER_WARP, 0, stream>>>( | ||
| tensor_map_input, empty_tensor_map, empty_tensor_map, tensor_map_output, | ||
| reinterpret_cast<const IType *>(input.data.dptr), nullptr, nullptr, | ||
| reinterpret_cast<OType *>(output->columnwise_data.dptr), shape_rep, num_tensors, | ||
| first_logical_dim, last_logical_dim, offsets_ptr, first_dims_ptr, last_dims_ptr, false, | ||
| true, false); |
There was a problem hiding this comment.
Concurrent calls overwrite shared descriptors
For groups with varying widths, two nvte_group_requantize calls on different CUDA streams overwrite the same g_tensor_maps entries. Each call updates those entries before launching its kernel, but stream ordering does not protect them from the other call. One kernel can therefore read or write the other call's buffers, producing wrong output or an illegal memory access.
Keep the descriptors in storage owned by each call until its kernel finishes.
Description
Add a fused grouped MXFP8 requantization path for Blackwell GPUs. The new operation converts compact rowwise MXFP8 input into columnwise E4M3 MXFP8 without materializing an intermediate high-precision tensor in global memory.
This reduces memory traffic compared with separate grouped dequantization and quantization calls. The kernel uses a tiled TMA pipeline and packed data movement/conversion operations, and supports both compact and GEMM-swizzled output scaling factors. The quantization configuration selects either an FP32 intermediate or a faster BF16 intermediate.
Type of change
Changes
Please list the changes introduced in this PR:
nvte_group_requantize_mxfp8C API for fused grouped rowwise-to-columnwise MXFP8 requantization.Checklist: