deepseek-ai/DeepSeek-V3 · error · AssertionError

Input tensors must be contiguous

Error message

Input tensors must be contiguous

What it means

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.

Source

Thrown at inference/kernel.py:104

    tl.store(y_ptr + offs, y, mask=mask)


def weight_dequant(x: torch.Tensor, s: torch.Tensor, block_size: int = 128) -> torch.Tensor:
    """
    Dequantizes the given weight tensor using the provided scale tensor.

    Args:
        x (torch.Tensor): The quantized weight tensor of shape (M, N).
        s (torch.Tensor): The scale tensor of shape (M//block_size, N//block_size).
        block_size (int, optional): The block size to use for dequantization. Defaults to 128.

    Returns:
        torch.Tensor: The dequantized weight tensor of the same shape as `x`.

    Raises:
        AssertionError: If `x` or `s` are not contiguous or if their dimensions are not 2.
    """
    assert x.is_contiguous() and s.is_contiguous(), 'Input tensors must be contiguous'
    assert x.dim() == 2 and s.dim() == 2, 'Input tensors must have 2 dimensions'
    M, N = x.size()
    y = torch.empty_like(x, dtype=torch.get_default_dtype())
    grid = lambda meta: (triton.cdiv(M, meta['BLOCK_SIZE']), triton.cdiv(N, meta['BLOCK_SIZE']))
    weight_dequant_kernel[grid](x, s, y, M, N, BLOCK_SIZE=block_size)
    return y


fp8_gemm_configs = [
    Config({'BLOCK_SIZE_M': block_m, 'BLOCK_SIZE_N': block_n, 'BLOCK_SIZE_K': 128}, num_stages=num_stages, num_warps=8)
    for block_m in [16, 32, 64] for block_n in [32, 64, 128] for num_stages in [3, 4, 5, 6]
]

@triton.autotune(configs=fp8_gemm_configs, key=['N', 'K'])
@triton.jit
def fp8_gemm_kernel(a_ptr, b_ptr, c_ptr,
                    a_s_ptr, b_s_ptr,
                    M, N: tl.constexpr, K: tl.constexpr,

View on GitHub (pinned to 9b4e9788e4)

Solutions

1. Call .contiguous() on both tensors before weight_dequant: weight_dequant(x.contiguous(), s.contiguous())
2. Load tensors directly from safetensors without intermediate views
3. If slicing shards, materialize with .clone() instead of passing views

Example fix

# before
w = big_w.narrow(0, 0, 512)  # non-contiguous view
bf16 = weight_dequant(w, s)

# after
w = big_w.narrow(0, 0, 512).contiguous()
bf16 = weight_dequant(w, s)

if not (x.is_contiguous() and s.is_contiguous()):
    x, s = x.contiguous(), s.contiguous()
y = weight_dequant(x, s)

def dequant_ready(x: torch.Tensor, s: torch.Tensor) -> bool:
    return x.is_contiguous() and s.is_contiguous() and x.dim() == 2 and s.dim() == 2

When it happens

Trigger: 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.

Common situations: 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.

Related errors

• Input tensor must be contiguous
• Input tensors must have 2 dimensions
• Scaling factor tensors must be contiguous
• Last dimension size must be divisible by block_size (block_s
• Warning: Missing scale_inv tensor for ${weight_name}, skippi

AI-assisted analysis of deepseek-ai/DeepSeek-V3@9b4e9788e4 (2026-08-14). Data as JSON: /api/errors/65a8126a9a5f055e. Report an issue: GitHub.