From b9ed112ba0b6144505f4379a41224fc70d94953d Mon Sep 17 00:00:00 2001 From: sarthak arora Date: Wed, 15 Apr 2026 10:31:48 -0500 Subject: [PATCH] feat: enable hybrid FP8 dtypes on Triton grouped GEMM backends The Triton JIT kernels are dtype-agnostic -- tl.dot handles mixed FP8 natively on gfx950 via v_mfma_f32_32x32x64_f8f6f4 with per-operand cbsz/blgp modifiers. This aligns the Triton dispatcher dtype gates with hipBLASLt (which already supports hybrid pairs) by including _HYBRID_SUPPORTED_DTYPES in GroupedGEMMFP8TritonBackend and GroupedGEMMFP8VariableKTritonBackend. Made-with: Cursor --- .../pytorch/kernels/grouped_gemm/grouped_gemm_fp8_impl.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/primus_turbo/pytorch/kernels/grouped_gemm/grouped_gemm_fp8_impl.py b/primus_turbo/pytorch/kernels/grouped_gemm/grouped_gemm_fp8_impl.py index e34414f07..9d2483cde 100644 --- a/primus_turbo/pytorch/kernels/grouped_gemm/grouped_gemm_fp8_impl.py +++ b/primus_turbo/pytorch/kernels/grouped_gemm/grouped_gemm_fp8_impl.py @@ -322,7 +322,7 @@ class GroupedGEMMFP8TritonBackend(KernelBackend): ScalingGranularity.BLOCKWISE, } - SUPPORTED_DTYPES = set(_COMMON_SUPPORTED_DTYPES) + SUPPORTED_DTYPES = set(_COMMON_SUPPORTED_DTYPES + _HYBRID_SUPPORTED_DTYPES) @staticmethod def can_handle( @@ -439,7 +439,7 @@ class GroupedGEMMFP8VariableKTritonBackend(KernelBackend): ScalingGranularity.BLOCKWISE, } - SUPPORTED_DTYPES = set(_COMMON_SUPPORTED_DTYPES) + SUPPORTED_DTYPES = set(_COMMON_SUPPORTED_DTYPES + _HYBRID_SUPPORTED_DTYPES) @staticmethod def can_handle(