[hipBLASLt][gfx950] FP8 e4m3 GEMM m=8192 n=304 k=28672: heuristic returns solution 36273 (MT256x320x128), which fails with HIPBLAS_STATUS_INTERNAL_ERROR #14570

Problem Description

On gfx950, hipBLASLt 1.4.1 (ROCm 10.0.0 / HIP 7.15.26333) returns HIPBLAS_STATUS_INTERNAL_ERROR from hipblasLtMatmul for an FP8 e4m3 → bf16 GEMM with m=8192, n=304, k=28672 (TN, per-tensor scale pointers, fp32 compute). The heuristic's top pick for this problem is solution index 36273 (MT256x320x128), and that solution fails whenever it is executed.

This hits vLLM in practice: amd/Llama-3.1-70B-Instruct-FP8-KV (TP1) fails at startup on the vLLM ROCm 10 nightly images. The failure comes during CUDA-graph capture at batch size 304, where down_proj produces exactly this shape.

Operating System

Ubuntu 24.04 LTS

CPU

AMD EPYC 9655 96-Core Processor

GPU

AMD Instinct MI350X

ROCm Version

ROCm 10.0

Installation Method

Python wheels

Installed ROCm Packages / Versions

rocm==10.0.0
rocm-bootstrap==0.3.0
rocm-sdk-core==10.0.0
rocm-sdk-devel==10.0.0
rocm-sdk-libraries==10.0.0
rocm-sdk-device-gfx950==10.0.0
torch==2.13.0+rocm10.0.0
amd-torch-device-gfx950==2.13.0+rocm10.0.0

ROCm Component

hipBLASLt

Steps to Reproduce

import torch
K, M, n = 28672, 8192, 304
fp8 = torch.float8_e4m3fn
B = torch.randn(M, K, device="cuda").to(fp8).t()
A = torch.randn(n, K, device="cuda").to(fp8)
s = torch.tensor(1.0, device="cuda")
torch._scaled_mm(A, B, out_dtype=torch.bfloat16, scale_a=s, scale_b=s)
# RuntimeError: CUDA error: HIPBLAS_STATUS_INTERNAL_ERROR when calling hipblasLtMatmul with
#   transpose_mat1 t transpose_mat2 n m 8192 n 304 k 28672 mat1_ld 28672 mat2_ld 28672
#   result_ld 8192 computeType 2 scaleType 0