Quantized GEMV: Decode Is A Bandwidth Problem

Lesson

Single-batch LLM decode is the purest bandwidth workload in production ML: for each generated token, y = W x runs once per weight matrix, every weight participates in exactly one multiply-add, with limited arithmetic reuse; cache behavior depends on the working set. Tokens/second is, to first order, memory bandwidth divided by total weight bytes. That arithmetic is why weight-only quantization (LLM.int8, GPTQ, AWQ) became the default for local inference: store W in int8 with a per-row (or per-group) fp scale, stream a quarter of the bytes, and dequantize in registers where the multiply happens. Activations stay in floating point - hence 'weight-only' - so accuracy depends on the quantization method and data, while conversion and scaling still cost work.

This harness generates the weights FROM their int8 representation (W = q * scale exactly), so the fp32 starter and the dequantizing solution compute the same values and verify against the same double reference - both memory requests and conversion/scaling instructions may change. The kernel structure (one block per row, block-stride partial products, shared-memory tree reduction) does not change at all; only the load changes: W[idx] becomes static_cast<float>(Wq[idx]) * scale, with the row's scale read once. Note the int8 loads are still fully coalesced - 32 consecutive lanes read 32 consecutive bytes, one 32-byte sector per warp access instead of four.

The caveat the literature stresses: dequantization is not free. Clumsy dequant (per-element scale reloads, format shuffling) can eat much of the win - the SM_THROUGHPUT side channel in your profile will show it if the dequant path gets expensive. Here the dequant adds a conversion and scaling operation; measure their cost.

Metric contract: algorithmic reads/writes count source-level sweeps. GLOBAL_MEMORY_SECTORS counts SASS-derived issued/requested sectors in the selected profiling scope, including repeated launches. It is cache-blind: it is neither L2 traffic nor device-memory bytes. Trace elapsed time is a separate measurement; a global load may hit cache.

Test weights are deliberately constructed from quantized values. This does not make quantization identical to arbitrary FP32 weights. Preserve scales and use the stated tolerance; weight-byte savings differ from total traffic and elapsed-time savings.

What you will learn

• Do the decode roofline arithmetic: tokens/s ~ bandwidth / weight bytes.
• Implement weight-only dequantization in-kernel: int8 weight times per-row scale.
• Keep the reduction structure unchanged - quantization only touches the load.
• Read the 4x weight-stream drop in GLOBAL_MEMORY_SECTORS.

Starter code

    // Full-precision GEMV: one block per row, already coalesced - the only
    // lever left is the BYTES the weight stream costs.
    __shared__ float red[256];

    int row = blockIdx.x;
    if (row >= rows) return;
    int tid = threadIdx.x;
    const float* w = W + static_cast<size_t>(row) * cols;

    float partial = 0.0f;
    for (int c = tid; c < cols; c += blockDim.x) {
        partial += w[c] * x[c];
    }
    red[tid] = partial;
    __syncthreads();
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) red[tid] += red[tid + stride];
        __syncthreads();
    }
    if (tid == 0) y[row] = red[0];