Multiple out-of-bounds accesses in QMatrix quant-metadata handling via a crafted GPTQ/EXL2 model

[professor-moody]

QMatrix (exllamav2_ext/cuda/q_matrix.cu) does not validate any of the quantization metadata it reads from a model, so a crafted GPTQ/EXL2 checkpoint triggers several out-of-bounds accesses at model-load and inference time. All are reachable from an untrusted model file (the metadata is passed through from ext.py/linear.py unvalidated) and confirmed with AddressSanitizer / compute-sanitizer on HEAD 7dc12af.

1. Heap OOB write via unbounded GPTQ g_idx (load time).
QMatrix::make_sequential builds a groups-sized histogram indexed by the raw g_idx:

uint32_t* cpu_g_idx_map = (uint32_t*) calloc(groups, sizeof(uint32_t));
for (int i = 0; i < height; i++) cpu_g_idx_map[cpu_g_idx[i]]++;   // OOB write when g_idx[i] >= groups

cpu_g_idx[i] is the GPTQ act-order permutation from the model, nothing constrains it to [0, groups). ASan: heap-buffer-overflow WRITE of size 4. A large value SEGVs on a wild write. The second loop (x-map) does more OOB read/write.

2. Device OOB read of the scale buffers via unbounded EXL2 q_group_map (inference).
The dequant/matmul kernels take the group index straight from b_q_group_map and index the scale matrices with no clamp: q_gemm_kernel.cuh int group = b_q_group_map[offset_k*2] then b_q_scale_.item4(qscales, group+g, n) / b_q_scale_max[group+g]; reconstruct_kernel (q_matrix.cu) does the same with b_q_scale_.item(group,n). MatrixView::item is a raw data[row*width+col]. compute-sanitizer with a crafted q_group_map: Invalid __global__ read, ~63 KB out of bounds. q_group_map is loaded directly from the model (linear.py) or computed from q_groups.

3. Divide-by-zero at load via EXL2 bits.
int bits = cpu_q_groups[i*2]; ... rows = qrows * 32 / bits; with no bits != 0 check. A group declaring bits = 0 -> SIGFPE at load (confirmed, exit 136). (1 << (bits-1) is also UB for bits==0/large.)

4. Device OOB via inflated row counts from EXL2 q_groups.
rows_8..rows_2 are computed from q_groups (rows = qrows*32/bits, qrows an attacker delta) with no validation; shuffle_kernel then walks the weight pointer by those counts. A crafted q_groups inflates the counts so shuffle_4bit_8 reads past the weight allocation. compute-sanitizer: Invalid __global__ read, 1 byte past a 512-byte allocation.

Impact: loading (or running) a crafted GPTQ/EXL2 model gives a heap out-of-bounds write with an attacker-controlled offset, device out-of-bounds reads, and a load-time crash, the untrusted-checkpoint threat model.

Suggested fix: validate all model-supplied quant metadata in the QMatrix constructor / make_sequential: every g_idx in [0, groups), every q_group_map group index in [0, groups), bits in a valid set, groups > 0, and the q_groups row counts consistent with the weight buffer size.