{"record":{"id":"65a8126a9a5f055e","repo":"deepseek-ai/DeepSeek-V3","slug":"input-tensors-must-be-contiguous","errorCode":null,"errorMessage":"Input tensors must be contiguous","messagePattern":"Input tensors must be contiguous","errorType":"validation","errorClass":"AssertionError","httpStatus":null,"severity":"error","filePath":"inference/kernel.py","lineNumber":104,"sourceCode":"    tl.store(y_ptr + offs, y, mask=mask)\n\n\ndef weight_dequant(x: torch.Tensor, s: torch.Tensor, block_size: int = 128) -> torch.Tensor:\n    \"\"\"\n    Dequantizes the given weight tensor using the provided scale tensor.\n\n    Args:\n        x (torch.Tensor): The quantized weight tensor of shape (M, N).\n        s (torch.Tensor): The scale tensor of shape (M//block_size, N//block_size).\n        block_size (int, optional): The block size to use for dequantization. Defaults to 128.\n\n    Returns:\n        torch.Tensor: The dequantized weight tensor of the same shape as `x`.\n\n    Raises:\n        AssertionError: If `x` or `s` are not contiguous or if their dimensions are not 2.\n    \"\"\"\n    assert x.is_contiguous() and s.is_contiguous(), 'Input tensors must be contiguous'\n    assert x.dim() == 2 and s.dim() == 2, 'Input tensors must have 2 dimensions'\n    M, N = x.size()\n    y = torch.empty_like(x, dtype=torch.get_default_dtype())\n    grid = lambda meta: (triton.cdiv(M, meta['BLOCK_SIZE']), triton.cdiv(N, meta['BLOCK_SIZE']))\n    weight_dequant_kernel[grid](x, s, y, M, N, BLOCK_SIZE=block_size)\n    return y\n\n\nfp8_gemm_configs = [\n    Config({'BLOCK_SIZE_M': block_m, 'BLOCK_SIZE_N': block_n, 'BLOCK_SIZE_K': 128}, num_stages=num_stages, num_warps=8)\n    for block_m in [16, 32, 64] for block_n in [32, 64, 128] for num_stages in [3, 4, 5, 6]\n]\n\n@triton.autotune(configs=fp8_gemm_configs, key=['N', 'K'])\n@triton.jit\ndef fp8_gemm_kernel(a_ptr, b_ptr, c_ptr,\n                    a_s_ptr, b_s_ptr,\n                    M, N: tl.constexpr, K: tl.constexpr,","sourceCodeStart":86,"sourceCodeEnd":122,"githubUrl":"https://github.com/deepseek-ai/DeepSeek-V3/blob/9b4e9788e4a3a731f7567338ed15d3ec549ce03b/inference/kernel.py#L86-L122","documentation":"Thrown by weight_dequant (inference/kernel.py:104): the Triton weight_dequant_kernel launches on a 2-D grid of raw pointers, so both the FP8 weight tensor x and its scale tensor s must be contiguous with no stride holes. It dequantizes FP8 checkpoint weights back to bf16/fp32 during FP8-to-BF16 conversion.","triggerScenarios":"Calling weight_dequant(x, s) with tensors produced by narrow/slice/transpose — e.g. slicing a shard out of a loaded safetensors dict then passing the view directly, or transposing weights to match an expected layout first.","commonSituations":"Mostly hit when writing custom conversion or evaluation scripts around fp8_cast_bf16.py that pre-manipulate tensors; the stock script passes freshly loaded (contiguous) tensors and never fires this.","solutions":["Call .contiguous() on both tensors before weight_dequant: weight_dequant(x.contiguous(), s.contiguous())","Load tensors directly from safetensors without intermediate views","If slicing shards, materialize with .clone() instead of passing views"],"exampleFix":"# before\nw = big_w.narrow(0, 0, 512)  # non-contiguous view\nbf16 = weight_dequant(w, s)\n\n# after\nw = big_w.narrow(0, 0, 512).contiguous()\nbf16 = weight_dequant(w, s)","handlingStrategy":"type-guard","validationCode":"if not (x.is_contiguous() and s.is_contiguous()):\n    x, s = x.contiguous(), s.contiguous()\ny = weight_dequant(x, s)","typeGuard":"def dequant_ready(x: torch.Tensor, s: torch.Tensor) -> bool:\n    return x.is_contiguous() and s.is_contiguous() and x.dim() == 2 and s.dim() == 2","tryCatchPattern":null,"preventionTips":["Materialize shard slices with .contiguous()/.clone() before kernel calls","Load weights straight from safetensors without intermediate views","Centralize contiguity checks in a wrapper around Triton launches"],"tags":["triton","fp8","dequantization","tensor-layout","deepseek"],"backgroundTag":null,"analyzedSha":"9b4e9788e4a3a731f7567338ed15d3ec549ce03b","analyzedAt":"2026-08-14T19:02:32.748Z","schemaVersion":2},"datasetVersion":"2026-08-15T22:17:37.221Z"}