sgl-project/sglang · error · RuntimeError

The layout of v is not supported

Error message

The layout of v is not supported

What it means

The kernel requires V in MN-major (column-major) layout, the opposite of Q/K, per the second operand of the PV MMA. LayoutEnum.from_tensor(v).mma_major_mode() returned K-major, so V was passed row-major.

Source

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

            )
        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,
            self.q_major_mode,
            self.k_major_mode,
            self.qk_acc_dtype,
            cta_group,

View on GitHub (pinned to 0132848349)

Solutions

1. Provide v in MN-major layout as produced by the intended cache/prepare path
2. Review how the kernel's reference caller prepares v (e.g. v.transpose or cache layout) and mirror it
3. Do not blanket-contiguous() v before this kernel

Example fix

# before
out = fmha(q, k, v.contiguous())
# after
out = fmha(q, k, v)  # keep v in the kernel-expected MN-major layout

assert v.stride(-2) == 1 or layout_matches_kernel(v), 'v must be MN-major'

When it happens

Trigger: Passing a fully contiguous/row-major v tensor where the kernel expects MN-major layout for V.

Common situations: Calling .contiguous() uniformly on q/k/v 'for safety'; V coming from a cache that stores V in the same layout as K.

Related errors

• The layout of k is not supported
• The layout of q 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/ce8ac7e6fe7d5d32. Report an issue: GitHub.