[Common/PyTorch] Guard FP8 per-tensor scaling grouped GEMM on Hopper behind cuBLAS 13.5+#3181
Conversation
cuBLAS 13.4 has no SM90 grouped GEMM algorithms for the per-tensor (PER_BATCH_SCALAR_32F) FP8 scale mode, so the heuristic query fails with CUBLAS_STATUS_NOT_SUPPORTED and users get a cryptic "Unable to find suitable cuBLAS grouped GEMM algorithm" error (see NVIDIA#3176). The support is available starting with cuBLAS 13.5. - Add CUBLAS_FP8_TENSOR_SCALING_GROUPED_GEMM_HOPPER_VERSION (130500) and a runtime check with a clear error message on the FP8 tensor-scaling grouped GEMM path. - Enable the FP8Current grouped GEMM C++ tests on Hopper with cuBLAS 13.5+ (previously skipped as Blackwell-only, so this path had no C++ coverage on SM90). - Skip FP8 current scaling GroupedLinear tests on Hopper when cuBLAS < 13.5 instead of failing in algorithm selection. Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
Greptile SummaryThis PR adds clearer gating for FP8 tensor-scaling grouped GEMM on Hopper. The main changes are:
Confidence Score: 4/5This is close, but the fallback path should be fixed before merging.
transformer_engine/pytorch/ops/basic/grouped_linear.py Important Files Changed
Reviews (2): Last reviewed commit: "Fall back to legacy GroupedLinear path w..." | Re-trigger Greptile |
![greptile-apps[bot]](https://avatars.githubusercontent.com/in/867647?s=60&v=4){kind=small}
| if (sm < 100) { | ||
| NVTE_CHECK(transformer_engine::cuda::cublas_version() >= | ||
| CUBLAS_FP8_TENSOR_SCALING_GROUPED_GEMM_HOPPER_VERSION, | ||
| "FP8 tensor scaling grouped GEMM on Hopper (SM90) requires cuBLAS 13.5+, " | ||
| "but run-time cuBLAS version is ", | ||
| transformer_engine::cuda::cublas_version()); | ||
| } |
There was a problem hiding this comment.
When FP8 current-scaling grouped MLP tests run on Hopper with cuBLAS 13.4, this new guard fails before algorithm selection. The grouped-linear tests now skip that runtime state, but the existing grouped-MLP current-scaling path still allows it, so the affected test run still ends in an NVTE_CHECK failure instead of a skip.
…vailable Route around the fused GroupedTensor path instead of erroring out when the cuBLAS version is too old: grouped GEMM needs cuBLAS 13.3+ (13.4+ on Hopper), and FP8 per-tensor current scaling on Hopper needs 13.5+. Both the module (_is_grouped_tensor_path_supported) and the ops (_is_graph_safe_path_supported) checks now fall back to the legacy multi-stream path in these cases. Signed-off-by: Pawel Gadzinski <pgadzinski@nvidia.com>
| ): | ||
| return False |
There was a problem hiding this comment.
Returning False here only disables the graph-safe grouped path. The caller can then use the split-quantize fallback, which still calls general_grouped_gemm with FP8 tensors for current-scaling grouped MLP cases. On Hopper with cuBLAS 13.4, that fallback reaches the new C++ cuBLAS 13.5 check and raises NVTE_CHECK instead of avoiding the unsupported grouped GEMM path. The Hopper/cuBLAS<13.5 case needs to skip this path or route it away from FP8 grouped GEMM before the fallback call.
|
/te-ci |
4 participants
Description
FP8 per-tensor (current/delayed scaling) grouped GEMM on Hopper requires cuBLAS 13.5+: cuBLAS 13.4 has no SM90 grouped GEMM algorithms for the
PER_BATCH_SCALAR_32Fscale mode, socublasLtMatmulAlgoGetHeuristicreturnsCUBLAS_STATUS_NOT_SUPPORTEDand the user gets a cryptic "Unable to find suitable cuBLAS grouped GEMM algorithm" error. This is exactly what happens in #3176 (NGC PyTorch 26.04 ships cuBLAS 13.4.0; reproduced on H100 — fails on cuBLAS 13.4.0, passes on 13.5.1 and 13.6.0). BF16 and FP8 block-scaling grouped GEMMs on Hopper work on 13.4 and keep their existing 13.4 gate.This PR makes GroupedLinear fall back to the legacy multi-stream path when the fused grouped GEMM is not available for the current cuBLAS version, adds a dedicated version guard with a clear error message in the common library, and closes the test-coverage gap: the
FP8Currentgrouped GEMM C++ tests were skipped on Hopper as Blackwell-only, so the only coverage of this path on SM90 was the PyTorch GroupedLinear tests, which CI runs on containers with cuBLAS 13.6.Fixes #3176
Type of change
Changes
cublaslt_grouped_gemm.cu: addCUBLAS_FP8_TENSOR_SCALING_GROUPED_GEMM_HOPPER_VERSION(130500) and a runtimeNVTE_CHECKon the FP8 tensor-scaling path for SM < 100 with an actionable error message.transformer_engine/pytorch/module/grouped_linear.pyandtransformer_engine/pytorch/ops/basic/grouped_linear.py: route to the legacy multi-stream path instead of the fused GroupedTensor path when cuBLAS is too old (13.3+ general, 13.4+ Hopper, 13.5+ FP8 per-tensor scaling on Hopper).tests/cpp/operator/test_grouped_gemm.cu: runFP8Currentcases on Hopper with cuBLAS 13.5+ (previously skipped as Blackwell-only); skip with a clear reason on older cuBLAS.tests/pytorch/test_grouped_linear.py: skip FP8 current scaling GroupedTensor tests on Hopper when cuBLAS < 13.5.Checklist:
🤖 Generated with Claude Code