dequantize_4bit reads out of bounds when quant_state.shape exceeds the packed tensor #2106

Summary

dequantize_4bit uses quant_state.shape to decide how many 4-bit values to read from the packed tensor, but never checks that the packed tensor is actually large enough for that shape. A QuantState whose shape claims more elements than the packed weight holds makes kDequantizeBlockwise read past the end of the packed buffer.

This is reachable from a corrupt or crafted checkpoint: Params4bit.from_prequantized / QuantState.from_dict copy shape, absmax, etc. straight from the stored state with no consistency check against the packed data.

Per the GPU threat model in agents/security_guide.md (Tier 7), I'm filing this as a correctness/reliability issue (silent wrong results / out-of-bounds read / possible process crash), not a traditional security exploit — it's an out-of-bounds read, not a write, and the output buffer is correctly sized.

Reproduction

import torch, bitsandbytes.functional as F

honest = torch.randn(64, 1, device="cuda", dtype=torch.float16)
packed, state = F.quantize_4bit(honest, blocksize=64, quant_type="nf4")   # 32-byte packed buffer

claimed = (4096, 4096)                      # 16,777,216 elements vs 64 real
n = claimed[0] * claimed[1]
forged = F.QuantState(
    quant_type="nf4",
    absmax=torch.ones(-(n // -64), device="cuda", dtype=torch.float32),   # sized for the lie
    blocksize=64,
    code=state.code.clone(),
    dtype=torch.float16,
    shape=torch.Size(claimed),
)
out = F.dequantize_4bit(packed, forged)      # returns an 8 MB tensor from 32 packed bytes

compute-sanitizer (with PYTORCH_NO_CUDA_MEMORY_CACHING=1, since the caching allocator otherwise masks the few-byte-to-megabyte overrun):

Invalid __global__ read of size 1 bytes
    at void kDequantizeBlockwise<__half, (int)512, (int)64, (int)8, (int)2>(...)
    Access to 0xb0ac00500 is out of bounds
    and is 225 bytes after the nearest allocation at 0xb0ac00400 of size 32 bytes

The read walks progressively further past the 32-byte allocation, up to the claimed prod(shape).

Root cause

dequantize_4bit (bitsandbytes/functional.py) validates blocksize, quant_type, and dtype, but not A.numel() vs prod(quant_state.shape) (nor absmax.numel() vs the block count). from_prequantized / from_dict do not validate either.

Fix

PR adds a size check in dequantize_4bit so an inconsistent state raises a clear ValueError instead of reading out of bounds, plus a matching check on absmax length. The honest round-trip path is unchanged.

Environment: bitsandbytes 0.50.0 (also present on main), torch 2.11.0+cu128, CUDA 12.8, RTX 5060.