sgl-project/sglang · error · ValueError

cta_n={cta_n} invalid for use_2cta={use_2cta}: bf16 K-major

Error message

cta_n={cta_n} invalid for use_2cta={use_2cta}: bf16 K-major mma requires N ∈ [{min_n}, 256] step {step_n}

What it means

The bf16 K-major tcgen05 MMA configuration constrains cta_n to [8, 256] step 8 for 1-CTA mode and [16, 256] step 16 for 2-CTA mode; these come from the Blackwell MMA atom limits. Invalid cta_n values raise at kernel config construction.

Source

Thrown at python/sglang/kernels/ops/gemm/cutedsl_bf16_gemm.py:144

        self.cta_m = cta_m
        self.cta_n = cta_n
        self.cta_k = cta_k
        self.num_ab_stage = num_ab_stage
        self.use_2cta = use_2cta
        self.use_pdl = use_pdl
        self.pdl_launch = pdl_launch if pdl_launch is not None else use_pdl
        self.pdl_count = pdl_count
        # has_bias: when True, kernel reads a (Gemm_M, Gemm_N, Gemm_L):(1,0,0)
        # bias tensor (M-broadcast over N,L), converts it to fp32 in RMEM, and
        # adds it to the accumulator before the bf16 cast. When False, all
        # bias-related code is elided via cutlass.const_expr.
        self.has_bias = has_bias

        # 1-CTA: cta_n ∈ [8, 256] step 8 (bf16 tcgen05.mma atom limit).
        # 2-CTA: cta_n ∈ [16, 256] step 16 (bf16 K-major cluster mma).
        min_n, step_n = (16, 16) if use_2cta else (8, 8)
        if cta_n < min_n or cta_n > 256 or cta_n % step_n != 0:
            raise ValueError(
                f"cta_n={cta_n} invalid for use_2cta={use_2cta}: "
                f"bf16 K-major mma requires N ∈ [{min_n}, 256] step {step_n}"
            )

        # Fixed configuration matching the C++ / DSL kernels.
        self.threads_per_cta = 256  # 8 warps (warp 3 unused)
        if use_2cta:
            # 2-CTA cluster along M; joint MMA tile = (cta_m*2, cta_n).
            self.cluster_shape = (2, 1, 1)
            self.mma_tiler_mn = (cta_m * 2, cta_n)
            self.cta_group = tcgen05.CtaGroup.TWO
            self.tma_op = cute_ext.OperationTypeEnum.SM100_TMA_LOAD_2SM
        else:
            # 1 SM mode, 1x1 cluster, no multicast.
            self.cluster_shape = (1, 1, 1)
            self.mma_tiler_mn = (cta_m, cta_n)
            self.cta_group = tcgen05.CtaGroup.ONE
            self.tma_op = cute_ext.OperationTypeEnum.SM90_TMA_LOAD

View on GitHub (pinned to 0132848349)

Solutions

1. Set cta_n to a multiple of 8 (1-CTA) or 16 (2-CTA) within [min, 256].
2. If use_2cta=True, ensure cta_n >= 16.
3. Prefer selecting a validated tactic id via _resolve_tactic instead of manual tile sizes.

Example fix

// before
gemm = Kernel(use_2cta=True, cta_n=8)
// after
gemm = Kernel(use_2cta=True, cta_n=16)

min_n, step_n = (16, 16) if use_2cta else (8, 8)
assert min_n <= cta_n <= 256 and cta_n % step_n == 0

When it happens

Trigger: Instantiating the cuteDSL bf16 GEMM class (or selecting a TGV tactic) with cta_n=4, cta_n=12 with use_2cta=True, or cta_n > 256.

Common situations: Hand-tuning tile sizes for a custom GEMM config without respecting the MMA atom alignment, or passing a tactic-derived cta_n meant for fp8 kernels.

Related errors

• Unsupported tcgen05 MMA op kind: {type(op).__name__}
• cutedsl_bf16_gemm requires an SM10x GPU
• fp8_blockwise_scaled_mm JIT kernel requires SM120 (Blackwell
• Unsupported type {type(data)}
• num_heads must be divisible by num_epi_subtiles

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