sgl-project/sglang · error · RuntimeError

The layout of q is not supported

Error message

The layout of q is not supported

What it means

The SM100 hd256 FMHA kernel requires Q in K-major (row-major per head) MMA layout; it inspects the tensor's stride pattern via LayoutEnum.from_tensor(q).mma_major_mode(). A MN-major (column-major) Q is not supported by this tcgen05 kernel configuration.

Source

Thrown at python/sglang/kernels/ops/attention/flash_attn/cute/sm100_hd256_2cta_fmha_forward.py:427

            self.tile_sched_params, grid = compute_grid_clc(
                (s_q, o.shape[1], o.shape[2]) if cum_seqlen_q is not None else o.shape,
                self.cta_tiler,
                (*self.cluster_shape_mn, 1),
            )
        else:
            self.tile_sched_params, grid = compute_grid(
                (s_q, o.shape[1], o.shape[2]) if cum_seqlen_q is not None else o.shape,
                self.cta_tiler,
                self.is_persistent,
            )

        self.q_major_mode = utils.LayoutEnum.from_tensor(q).mma_major_mode()
        self.k_major_mode = utils.LayoutEnum.from_tensor(k).mma_major_mode()
        self.v_major_mode = utils.LayoutEnum.from_tensor(v).mma_major_mode()
        self.o_layout = utils.LayoutEnum.from_tensor(o)

        if cutlass.const_expr(self.q_major_mode != tcgen05.OperandMajorMode.K):
            raise RuntimeError("The layout of q is not supported")
        if cutlass.const_expr(self.k_major_mode != tcgen05.OperandMajorMode.K):
            raise RuntimeError("The layout of k is not supported")
        if cutlass.const_expr(self.v_major_mode != tcgen05.OperandMajorMode.MN):
            raise RuntimeError("The layout of v is not supported")

        # check type consistency
        if cutlass.const_expr(self.q_dtype != self.k_dtype):
            raise TypeError(f"Type mismatch: {self.q_dtype} != {self.k_dtype}")
        if cutlass.const_expr(self.q_dtype != self.v_dtype):
            raise TypeError(f"Type mismatch: {self.q_dtype} != {self.v_dtype}")
        self._setup_attributes()

        cta_group = tcgen05.CtaGroup.TWO
        # the intermediate tensor p is from tmem & k-major
        p_source = tcgen05.OperandSource.TMEM
        p_major_mode = tcgen05.OperandMajorMode.K
        qk_tiled_mma = sm100_utils.make_trivial_tiled_mma(
            self.q_dtype,

View on GitHub (pinned to 0132848349)

Solutions

1. Make q contiguous in the last dim: q = q.contiguous()
2. Check that qkv projection produces row-major output
3. Verify no transpose was applied to q before the call

Example fix

# before
out = fmha(q.transpose(-1,-2).contiguous(), k, v)  # wrong layout
# after
out = fmha(q.contiguous(), k, v)

assert q.stride(-1) == 1, 'q must be K-major (row-major)'

def q_is_k_major(q: torch.Tensor) -> bool:
    return q.stride(-1) == 1

When it happens

Trigger: Passing q whose last dim stride != 1 (transposed/column-major) to the hd256 2cta fmha forward call.

Common situations: Upstream code transposes q for a different attention backend; a projection output happens to be non-contiguous; mixing kernels with differing layout requirements.

Related errors

• The layout of k is not supported
• The layout of v is not supported
• hd256 forward varlen expects k rank 3 or 5, got rank {k_rank
• hd256 forward non-varlen expects k rank 4 or 5, got rank {k_
• Type mismatch: {self.q_dtype} != {self.k_dtype}

AI-assisted analysis of sgl-project/sglang@0132848349 (2026-08-28). Data as JSON: /api/errors/96d381f266455a5d. Report an issue: GitHub.