{"record":{"id":"490c6a9e852f5e2f","repo":"deepseek-ai/DeepSeek-V3","slug":"scaling-factor-tensors-must-be-contiguous","errorCode":null,"errorMessage":"Scaling factor tensors must be contiguous","messagePattern":"Scaling factor tensors must be contiguous","errorType":"validation","errorClass":"AssertionError","httpStatus":null,"severity":"error","filePath":"inference/kernel.py","lineNumber":189,"sourceCode":"    mask = (offs_m[:, None] < M) & (offs_n[None, :] < N)\n    tl.store(c_ptrs, c, mask=mask)\n\n\ndef fp8_gemm(a: torch.Tensor, a_s: torch.Tensor, b: torch.Tensor, b_s: torch.Tensor):\n    \"\"\"\n    Perform a matrix multiplication using FP8 precision.\n\n    Args:\n        a (torch.Tensor): The first input matrix, must be contiguous.\n        a_s (torch.Tensor): The scaling factor for the first input matrix, must be contiguous.\n        b (torch.Tensor): The second input matrix, must be contiguous.\n        b_s (torch.Tensor): The scaling factor for the second input matrix, must be contiguous.\n\n    Returns:\n        torch.Tensor: The result of the matrix multiplication.\n    \"\"\"\n    assert a.is_contiguous() and b.is_contiguous(), 'Input tensors must be contiguous'\n    assert a_s.is_contiguous() and b_s.is_contiguous(), 'Scaling factor tensors must be contiguous'\n    K = a.size(-1)\n    M = a.numel() // K\n    N = b.size(0)\n    c = a.new_empty(*a.size()[:-1], N, dtype=torch.get_default_dtype())\n    grid = lambda META: (triton.cdiv(M, META['BLOCK_SIZE_M']), triton.cdiv(N, META['BLOCK_SIZE_N']))\n    fp8_gemm_kernel[grid](a, b, c, a_s, b_s, M, N, K)\n    return c\n","sourceCodeStart":171,"sourceCodeEnd":197,"githubUrl":"https://github.com/deepseek-ai/DeepSeek-V3/blob/9b4e9788e4a3a731f7567338ed15d3ec549ce03b/inference/kernel.py#L171-L197","documentation":"Thrown by fp8_gemm (inference/kernel.py:189): the per-block scaling factors a_s and b_s are also passed to the Triton kernel as raw pointers and must be contiguous. They are produced by act_quant (contiguous by construction), so this assert typically fires only when scales are sliced, transposed, or otherwise re-viewed by custom code.","triggerScenarios":"Calling fp8_gemm with a_s/b_s obtained via slicing (e.g. s[..., 1:]), transpose, or narrow of an act_quant output, or scales loaded from a checkpoint sharded with views. The preceding assert (error 11) checks a/b; this one covers the scale tensors.","commonSituations":"Custom kernels/fusers that manipulate scale tensors for batching; multi-rank setups passing per-rank scale shards as views instead of copies.","solutions":["Pass scales directly from act_quant output without modification","If you must transform them, end with .contiguous(): a_s = a_s.transpose(-1,-2).contiguous()","Check that sharding code materializes per-rank scale slices with .contiguous() (as convert.py does)"],"exampleFix":"# before\nout = fp8_gemm(a, a_s[:, 1:], b, b_s)  # sliced scale — view\n\n# after\na_s = a_s[:, 1:].contiguous()\nout = fp8_gemm(a, a_s, b, b_s)","handlingStrategy":"type-guard","validationCode":"if not (a_s.is_contiguous() and b_s.is_contiguous()):\n    a_s, b_s = a_s.contiguous(), b_s.contiguous()\nc = fp8_gemm(a, a_s, b, b_s)","typeGuard":"def scales_ready(a_s: torch.Tensor, b_s: torch.Tensor) -> bool:\n    return a_s.is_contiguous() and b_s.is_contiguous()","tryCatchPattern":null,"preventionTips":["Pass scale tensors straight from act_quant without edits","End any scale transformation with .contiguous()","Materialize per-rank scale shards rather than passing views"],"tags":["triton","fp8","gemm","tensor-layout","deepseek"],"backgroundTag":null,"analyzedSha":"9b4e9788e4a3a731f7567338ed15d3ec549ce03b","analyzedAt":"2026-08-14T19:02:32.748Z","schemaVersion":2},"datasetVersion":"2026-08-15T17:31:12.345Z"}