{"record":{"id":"dbca6d79092499b7","repo":"sgl-project/sglang","slug":"kv-canary-scatter-req-token-ids-bs-1-bs-1-exc","errorCode":null,"errorMessage":"kv-canary: scatter_req_token_ids bs+1={bs + 1} exceeds BATCH_BLOCK={_SCATTER_BATCH_BLOCK}; bump _SCATTER_BATCH_BLOCK if real workloads need this","messagePattern":"kv-canary: scatter_req_token_ids bs\\+1=(.+?) exceeds BATCH_BLOCK=(.+?); bump _SCATTER_BATCH_BLOCK if real workloads need this","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/kv_canary/scatter_req_token_ids.py","lineNumber":93,"sourceCode":"    if req_pool_indices.dtype != torch.int64:\n        raise TypeError(\n            f\"kv-canary: scatter_req_token_ids req_pool_indices must be int64, got \"\n            f\"{req_pool_indices.dtype}\"\n        )\n    if pool_out.dtype != torch.int32:\n        raise TypeError(\n            f\"kv-canary: scatter_req_token_ids pool_out must be int32, got \"\n            f\"{pool_out.dtype}\"\n        )\n\n    bs = int(req_pool_indices.shape[0])\n    if int(offsets.shape[0]) != bs + 1:\n        raise ValueError(\n            f\"kv-canary: scatter_req_token_ids offsets length {offsets.shape[0]} != \"\n            f\"bs+1 ({bs + 1})\"\n        )\n    if bs + 1 > _SCATTER_BATCH_BLOCK:\n        raise ValueError(\n            f\"kv-canary: scatter_req_token_ids bs+1={bs + 1} exceeds BATCH_BLOCK=\"\n            f\"{_SCATTER_BATCH_BLOCK}; bump _SCATTER_BATCH_BLOCK if real workloads need this\"\n        )\n\n    num_tokens = int(flat_in.shape[0])\n    if num_tokens == 0:\n        return\n\n    pool_stride0 = int(pool_out.stride(0))\n    pool_max_context_len = int(pool_out.shape[1])\n\n    grid = (triton.cdiv(num_tokens, _SCATTER_TOKEN_BLOCK),)\n    _scatter_req_token_ids_kernel[grid](\n        flat_in,\n        offsets,\n        req_pool_indices,\n        pool_out,\n        num_tokens=num_tokens,","sourceCodeStart":75,"sourceCodeEnd":111,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/kv_canary/scatter_req_token_ids.py#L75-L111","documentation":"The Triton scatter kernel keeps the whole offsets vector in registers via a single BLOCK of size _SCATTER_BATCH_BLOCK; if bs+1 exceeds that compile-time constant the kernel would silently drop requests, so the launcher raises ValueError instead. The message says to bump the constant if real workloads need larger batches.","triggerScenarios":"Calling launch_scatter_req_token_ids_kernel with a batch size larger than _SCATTER_BATCH_BLOCK - 1 (the kernel was sized for 'bs <= a few thousand').","commonSituations":"Scaling up server concurrency so the running batch exceeds the kernel's design limit; running large batch microbenchmarks.","solutions":["Split the batch into chunks of at most _SCATTER_BATCH_BLOCK - 1 requests and launch per chunk with sliced offsets/indices","Bump _SCATTER_BATCH_BLOCK in scatter_req_token_ids.py and recompile if your workload genuinely needs bigger single-launch batches","Check the current constant first to know your effective max bs"],"exampleFix":"# before\nlaunch_scatter(..., offsets=offsets, req_pool_indices=rp)  # bs too big\n# after\nCHUNK = _SCATTER_BATCH_BLOCK - 1\nfor i in range(0, bs, CHUNK):\n    launch_scatter(..., offsets=offsets[i:i+CHUNK+1], req_pool_indices=rp[i:i+CHUNK])","handlingStrategy":"validation","validationCode":"bs = req_pool_indices.shape[0]\nassert bs + 1 <= _SCATTER_BATCH_BLOCK, f\"bs={bs} exceeds kernel limit\"","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Chunk large batches at the caller; monitor running batch size vs the kernel's constant","Track _SCATTER_BATCH_BLOCK value across version upgrades"],"tags":["kv-cache","batch-size-limit","triton","capacity"],"backgroundTag":"kernel-batch-size-limit","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}