{"record":{"id":"f357d2394d4d708d","repo":"sgl-project/sglang","slug":"topk-kernels-only-support-k-32-k","errorCode":null,"errorMessage":"topk kernels only support k <= 32: {k=}","messagePattern":"topk kernels only support k <= 32: (.+?)","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/moe/gate_topk.py","lineNumber":142,"sourceCode":"    Stable implementation of torch.topk(..., dim=-1) that is most efficient\n    for small values of k.\n    \"\"\"\n    assert x.is_contiguous(), f\"{x.shape=} {x.stride()=}\"\n    assert x.ndim == 2, f\"{x.shape=}\"\n    assert x.numel() <= 2**31, f\"assumes int32 indexing: {x.shape=}\"\n    n_rows, n_cols = x.shape\n    if return_values:\n        values = torch.empty((n_rows, k), dtype=x.dtype, device=x.device)\n    else:\n        values = None\n    # int32 indices: column ids fit int32 (numel <= 2**31, asserted above) and the\n    # sole caller (Inkling gate) feeds the SRT MoeRunner topk-packing, which requires\n    # int32. The kernel store casts to the buffer dtype, so this emits int32\n    # directly — no separate .to(int32) downstream.\n    indices = torch.empty((n_rows, k), dtype=torch.int32, device=x.device)\n    if k > 32:\n        # For larger topk, we need to reevaluate the kernel strategy\n        raise NotImplementedError(f\"topk kernels only support k <= 32: {k=}\")\n\n    if _impl == \"streaming\":\n        BLOCK_SIZE_N = 32\n        BLOCK_SIZE_M = 32\n        grid = (triton.cdiv(n_rows, BLOCK_SIZE_M),)\n        _streaming_topk_kernel[grid](\n            x_ptr=x,\n            stride_xm=x.stride(0),\n            values_ptr=values,\n            indices_ptr=indices,\n            M=n_rows,\n            N=n_cols,\n            N_PAD=triton.cdiv(n_cols, BLOCK_SIZE_N) * BLOCK_SIZE_N,\n            K=k,\n            K_POW2=triton.next_power_of_2(k),\n            BLOCK_SIZE_M=BLOCK_SIZE_M,\n            BLOCK_SIZE_N=BLOCK_SIZE_N,\n            RETURN_VALUES=return_values,","sourceCodeStart":124,"sourceCodeEnd":160,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/moe/gate_topk.py#L124-L160","documentation":"NotImplementedError from gate_topk: the streaming top-k Triton kernel uses power-of-two block logic sized for at most 32 selections per row, so k > 32 cannot be handled and the function refuses loudly instead of producing wrong results.","triggerScenarios":"Calling gate_topk (from the routing forward) with topk > 32, e.g. a model config with num_selected_experts=64 or a very large top-k routing setting.","commonSituations":"Enabling a new model whose router selects more than 32 experts; experimenting with high top-k for routing quality; overriding topk via CLI/config without checking kernel limits.","solutions":["Reduce k to <= 32 if the model tolerates it (check the model's routing config)","Fall back to torch.topk-based routing instead of the fused Triton kernel for k > 32","If the workload genuinely needs k > 32, extend the kernel strategy as the error message suggests (new kernel path)"],"exampleFix":"// before\nids, w = gate_topk(logits, k=64)\n// after\nids, w = gate_topk(logits, k=32)  # or torch.topk(logits, 64, dim=-1)","handlingStrategy":"fallback","validationCode":"if k > 32:\n    vals, ids = torch.topk(logits, k, dim=-1)\nelse:\n    ids = gate_topk(logits, k)","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Validate topk against 32 when parsing model routing config","Prefer torch.topk in experimental high-k branches"],"tags":["moe","topk","triton","capacity-limit"],"backgroundTag":"topk-exceeds-limit","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}