{"record":{"id":"91130ef6af216146","repo":"sgl-project/sglang","slug":"cta-n-cta-n-invalid-for-use-2cta-use-2cta-bf1","errorCode":null,"errorMessage":"cta_n={cta_n} invalid for use_2cta={use_2cta}: bf16 K-major mma requires N ∈ [{min_n}, 256] step {step_n}","messagePattern":"cta_n=(.+?) invalid for use_2cta=(.+?): bf16 K-major mma requires N ∈ \\[(.+?), 256\\] step (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/gemm/cutedsl_bf16_gemm.py","lineNumber":144,"sourceCode":"        self.cta_m = cta_m\n        self.cta_n = cta_n\n        self.cta_k = cta_k\n        self.num_ab_stage = num_ab_stage\n        self.use_2cta = use_2cta\n        self.use_pdl = use_pdl\n        self.pdl_launch = pdl_launch if pdl_launch is not None else use_pdl\n        self.pdl_count = pdl_count\n        # has_bias: when True, kernel reads a (Gemm_M, Gemm_N, Gemm_L):(1,0,0)\n        # bias tensor (M-broadcast over N,L), converts it to fp32 in RMEM, and\n        # adds it to the accumulator before the bf16 cast. When False, all\n        # bias-related code is elided via cutlass.const_expr.\n        self.has_bias = has_bias\n\n        # 1-CTA: cta_n ∈ [8, 256] step 8 (bf16 tcgen05.mma atom limit).\n        # 2-CTA: cta_n ∈ [16, 256] step 16 (bf16 K-major cluster mma).\n        min_n, step_n = (16, 16) if use_2cta else (8, 8)\n        if cta_n < min_n or cta_n > 256 or cta_n % step_n != 0:\n            raise ValueError(\n                f\"cta_n={cta_n} invalid for use_2cta={use_2cta}: \"\n                f\"bf16 K-major mma requires N ∈ [{min_n}, 256] step {step_n}\"\n            )\n\n        # Fixed configuration matching the C++ / DSL kernels.\n        self.threads_per_cta = 256  # 8 warps (warp 3 unused)\n        if use_2cta:\n            # 2-CTA cluster along M; joint MMA tile = (cta_m*2, cta_n).\n            self.cluster_shape = (2, 1, 1)\n            self.mma_tiler_mn = (cta_m * 2, cta_n)\n            self.cta_group = tcgen05.CtaGroup.TWO\n            self.tma_op = cute_ext.OperationTypeEnum.SM100_TMA_LOAD_2SM\n        else:\n            # 1 SM mode, 1x1 cluster, no multicast.\n            self.cluster_shape = (1, 1, 1)\n            self.mma_tiler_mn = (cta_m, cta_n)\n            self.cta_group = tcgen05.CtaGroup.ONE\n            self.tma_op = cute_ext.OperationTypeEnum.SM90_TMA_LOAD","sourceCodeStart":126,"sourceCodeEnd":162,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/gemm/cutedsl_bf16_gemm.py#L126-L162","documentation":"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.","triggerScenarios":"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.","commonSituations":"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.","solutions":["Set cta_n to a multiple of 8 (1-CTA) or 16 (2-CTA) within [min, 256].","If use_2cta=True, ensure cta_n >= 16.","Prefer selecting a validated tactic id via _resolve_tactic instead of manual tile sizes."],"exampleFix":"// before\ngemm = Kernel(use_2cta=True, cta_n=8)\n// after\ngemm = Kernel(use_2cta=True, cta_n=16)","handlingStrategy":"validation","validationCode":"min_n, step_n = (16, 16) if use_2cta else (8, 8)\nassert min_n <= cta_n <= 256 and cta_n % step_n == 0","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Prefer tactic ids over manual tile sizes.","Document the MMA atom alignment when adding configs."],"tags":["gemm","blackwell","cutedsl","tile-config","mma"],"backgroundTag":"invalid-gemm-tile-config","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}