# GPU-Accelerated INT8 Quantization for KV Cache Compression in Large Language Models
**Authors**: Maanas Taneja, Purab Shingvi
(December 2025)
## Abstract
The key-value (KV) cache in large language models presents a significant memory bottleneck during inference, growing linearly with sequence length and often exceeding the memory footprint of model weights themselves. We implement and evaluate GPU-accelerated INT8 quantization for KV cache compression, achieving 4 $×$ memory reduction with minimal accuracy degradation. We develop four CUDA kernel variants—naive, tiled, coarsened, and vectorized—and benchmark them across realistic workload sizes up to 1 billion elements. Our vectorized kernel achieves up to 1,694 $×$ speedup over CPU baselines while maintaining reconstruction error below 0.004 and attention score error below 0.1 even for 8K-dimensional heads. These results demonstrate that INT8 quantization provides a practical approach for reducing memory pressure in LLM inference with negligible computational overhead (6–58ms) and minimal impact on downstream model behavior. Code available at: https://github.com/MaanasTaneja/cuda-kv-cache-compression
## 1 Introduction
Modern large language models (LLMs) based on the Transformer architecture have achieved impressive results across a wide range of tasks, from text generation to code completion. However, deploying these models efficiently presents a significant challenge: memory consumption during inference.
A key contributor to this memory bottleneck is the key-value (KV) cache. During autoregressive text generation, the model produces one token at a time, and at each step it must attend to all previously generated tokens. To avoid redundantly recomputing the key and value projections for past tokens, these values are stored in a cache. While this speeds up inference, the cache grows linearly with sequence length—for long contexts (32k, 128k tokens or more), the KV cache can consume tens of gigabytes of memory, sometimes exceeding the memory required by the model weights themselves.
This memory pressure has practical consequences: it limits the maximum context length, forces smaller batch sizes (reducing throughput), and increases serving costs on expensive GPU hardware. Quantization offers a promising solution to this problem. The idea is simple: instead of storing cached keys and values in 32-bit floating point (FP32), we compress them to 8-bit integers (INT8). This yields a 4 $×$ reduction in memory footprint. When the values are needed for attention computation, we dequantize them back to FP32. The tradeoff is a small loss in numerical precision due to rounding.
In this project, we implement and evaluate INT8 quantization and dequantization for KV cache compression. We use per-channel quantization, where each dimension of the key vectors gets its own scale factor, allowing us to preserve precision across dimensions with different value ranges. We implement a CPU baseline and four GPU kernel variants—naive, tiled, coarsened, and vectorized—to explore performance optimization strategies on parallel hardware.
To evaluate our implementation, we measure:
- Performance: How fast are the quantization and dequantization operations? What speedup do the GPU kernels achieve over the CPU?
- Reconstruction error: How closely does the dequantized matrix match the original? We report L2 error and maximum absolute error.
- Attention score error: Does quantization affect downstream computation? We measure the mean absolute difference between attention dot products computed with original and reconstructed keys.
## 2 Related Work
### 2.1 KV Cache Optimization
Recent work has explored various approaches to reducing KV cache memory consumption. FlashAttention [2] and FlashAttention-2 [3] optimize attention computation through kernel fusion and memory hierarchy exploitation but do not directly compress the cache. PagedAttention [4], implemented in the vLLM serving system, applies paging techniques to reduce memory fragmentation but maintains full precision storage.
### 2.2 Quantization Techniques
Quantization has emerged as a promising direction for LLM compression. LLM.int8() [5] demonstrates that 8-bit quantization can be applied to model weights and activations with minimal accuracy loss. SmoothQuant [6] addresses activation quantization challenges through per-channel scaling. More recent work explores extreme quantization: KIVI [7] proposes 2-bit asymmetric quantization specifically for KV caches, while KVQuant [8] achieves sub-4-bit precision through learned quantization parameters.
### 2.3 LLM Inference Systems
Production inference frameworks have begun incorporating memory optimization techniques. FlexGen [9] explores offloading strategies for memory-constrained scenarios. TensorRT-LLM and similar systems [10] integrate quantization into optimized inference pipelines. However, these systems often treat quantization as a black box, providing limited insight into the performance characteristics of different kernel implementations.
### 2.4 Our Contribution
Our work focuses on the implementation and systematic evaluation of per-channel INT8 quantization for KV cache compression. Unlike prior work that primarily evaluates end-to-end model accuracy, we provide detailed analysis of kernel-level performance characteristics across multiple GPU optimization strategies. We demonstrate that vectorized memory operations provide consistent speedups for this memory-bound workload, while traditional optimizations like shared memory tiling offer limited benefit. Our analysis provides practical guidance for implementing efficient quantization kernels in production LLM serving systems.
## 3 Background
### 3.1 Transformer Attention & KV Caching
The Transformer architecture relies on a mechanism called self-attention, which allows each token in a sequence to attend to all other tokens. For a given input, the model computes three projections: queries ( $Q$ ), keys ( $K$ ), and values ( $V$ ). The attention output is computed as:
$$
Attention(Q,K,V)=softmax≤ft(\frac{QK^T}{√{d_k}}\right)V \tag{1}
$$
where $d_k$ is the dimension of the key vectors.
While producing one token at a time, the model must compute attention between the current token’s query and the keys of all previous tokens. Naively, this would require recomputing the key and value projections for every past token at each generation step, a very expensive operation.
The KV cache solves this by storing the key and value projections from previous tokens. At each new generation step, only the current token’s $K$ and $V$ are computed and appended to the cache. This transforms the complexity from quadratic to linear in sequence length per step, but introduces a memory cost: the cache must store all past keys and values.
### 3.2 Memory Bottleneck in LLM Inference
For a model with $L$ layers, $H$ attention heads, head dimension $d$ , and sequence length $T$ , the KV cache stores:
$$
KV cache size=2× L× H× d× T×(bytes per element) \tag{2}
$$
Table 1: Example KV cache memory calculation
| Layers ( $L$ ) Heads ( $H$ ) Head dimension ( $d$ ) | 32 32 128 |
| --- | --- |
| Sequence length ( $T$ ) | 131,072 |
| Precision | FP32 (4 bytes) |
| Total memory | $≈$ 137 GB |
Even with FP16, this is nearly 70 GB, exceeding the memory of most GPUs. This memory pressure limits batch sizes, restricts context lengths, and increases deployment costs.
### 3.3 Quantization Fundamentals
Quantization reduces memory by representing values in fewer bits. The core idea is to map floating point values to a smaller set of discrete integer levels.
For linear quantization to INT8, a floating point value $x$ is mapped to an 8-bit integer $x_q$ using:
$$
x_q=round≤ft(\frac{x}{s}\right) \tag{3}
$$
where $s$ is a scale factor that maps the floating point range to the integer range $[-127,127]$ .
To recover the approximate original value (dequantization):
$$
\hat{x}=x_q× s \tag{4}
$$
The reconstruction $\hat{x}$ differs from the original $x$ due to rounding error. This error is bounded by $\frac{s}{2}$ - half the quantization step size.
Per-channel quantization uses a separate scale for each dimension (column) of the matrix:
$$
s_d=\frac{\max_t|K[t,d]|}{127} \tag{5}
$$
This allows dimensions with different value ranges to each utilize the full INT8 range, improving precision compared to a single global scale.
## 4 Method
### 4.1 Per-Channel INT8 Quantization
We apply per-channel quantization along the head dimension of the KV cache. For a key matrix $K$ with shape $(T,D)$ where $T$ is the number of tokens and $D$ is the head dimension, we compute one scale factor per column.
### 4.2 Scale Computation
For each dimension $d∈[0,D)$ , the scale is computed as:
$$
s_d=\frac{\max_t∈[0,T)|K[t,d]|}{127} \tag{6}
$$
This requires a single pass over each column to find the maximum absolute value. The scales are stored as FP32 values, adding negligible memory overhead ( $D$ floats compared to $T× D$ elements in the cache).
The scale computation algorithm is presented in Algorithm 1.
Algorithm 1 Compute Scales
1: $K$ : FP32 matrix of shape $T× D$
2: scales: FP32 array of length $D$
3: for $d=0$ to $D-1$ do
4: $max\_abs← 0$
5: for $t=0$ to $T-1$ do
6: if $|K[t,d]|>max\_abs$ then
7: $max\_abs←|K[t,d]|$
8: end if
9: end for
10: $scales[d]←max\_abs/127$
11: end for
### 4.3 Quantization and Dequantization Operations
Quantization converts each FP32 element to INT8:
$$
K_int8[t,d]=round≤ft(\frac{K[t,d]}{s_d}\right) \tag{7}
$$
The result is clamped to $[-127,127]$ to handle any numerical edge cases.
Dequantization recovers approximate FP32 values:
$$
\hat{K}[t,d]=K_int8[t,d]× s_d \tag{8}
$$
The round-trip introduces quantization error. For a value $x$ with scale $s$ :
$$
error=|x-\hat{x}|≤\frac{s}{2} \tag{9}
$$
This error propagates to attention computation, but as we measure empirically, the impact on attention scores is typically small.
## 5 Implementation
### 5.1 Data Structures
We define two matrix structures to hold original and quantized data:
⬇
1 typedef struct {
2 int rows; // T (number of tokens)
3 int columns; // D (head dimension)
4 float * data; // row-major storage
5} FP32Matrix;
6
7 typedef struct {
8 int rows;
9 int columns;
10 int8_t * data; // row-major storage
11} INT8Matrix;
Listing 1: Matrix data structures
The quantized matrix uses 4 $×$ less memory than the original.
### 5.2 CPU Baseline
The CPU implementation provides a correctness reference and performance baseline.
#### 5.2.1 Scale Computation
⬇
1 void compute_scales (FP32Matrix * K, float * scales) {
2 for (int d = 0; d < K -> columns; d ++) {
3 float max_abs = 0.0 f;
4 for (int t = 0; t < K -> rows; t ++) {
5 float val = fabsf (K -> data [t * K -> columns + d]);
6 if (val > max_abs) max_abs = val;
7 }
8 scales [d] = max_abs / 127.0 f;
9 }
10}
Listing 2: CPU scale computation
#### 5.2.2 Quantization
⬇
1 void quantize_matrix (FP32Matrix * K, INT8Matrix * K_int8, float * scales) {
2 for (int t = 0; t < K -> rows; t ++) {
3 for (int d = 0; d < K -> columns; d ++) {
4 float val = K -> data [t * K -> columns + d];
5 int q = (int) roundf (val / scales [d]);
6 // Clamp to INT8 range
7 if (q > 127) q = 127;
8 if (q < -127) q = -127;
9 K_int8 -> data [t * K -> columns + d] = (int8_t) q;
10 }
11 }
12}
Listing 3: CPU quantization
#### 5.2.3 Dequantization
⬇
1 void dequantize_matrix (INT8Matrix * K_int8, FP32Matrix * K_recon,
2 float * scales) {
3 for (int t = 0; t < K_int8 -> rows; t ++) {
4 for (int d = 0; d < K_int8 -> columns; d ++) {
5 int8_t q = K_int8 -> data [t * K_int8 -> columns + d];
6 K_recon -> data [t * K_int8 -> columns + d] = (float) q * scales [d];
7 }
8 }
9}
Listing 4: CPU dequantization
### 5.3 GPU Kernel Variants
We implement four CUDA kernel variants with increasing levels of optimization. All kernels perform the same computation but differ in how they map work to threads and utilize GPU memory hierarchy.
#### 5.3.1 Naive Kernel
The naive approach assigns one thread per matrix element. Each thread independently loads its scale, performs the computation, and writes the result.
⬇
1 __global__ void quantize_naive (float * K, int8_t * K_int8,
2 float * scales, int T, int D) {
3 int t = blockIdx. y * blockDim. y + threadIdx. y; // row
4 int d = blockIdx. x * blockDim. x + threadIdx. x; // column
5
6 if (t < T && d < D) {
7 float val = K [t * D + d];
8 float s = scales [d];
9 int q = __float2int_rn (val / s);
10 q = max (-127, min (127, q));
11 K_int8 [t * D + d] = (int8_t) q;
12 }
13}
Listing 5: Naive CUDA kernel for quantization
Characteristics:
- Simple & easy to understand
- Memory coalescing access.
- Redundant scale loads -multiple threads loading the same scale
#### 5.3.2 Tiled Kernel
The tiled kernel uses shared memory to cache scales, reducing redundant global memory accesses.
⬇
1 __global__ void quantize_tiled (float * K, int8_t * K_int8,
2 float * scales, int T, int D) {
3 __shared__ float s_scales [TILE_DIM];
4
5 int t = blockIdx. y * blockDim. y + threadIdx. y;
6 int d = blockIdx. x * blockDim. x + threadIdx. x;
7
8 if (threadIdx. y == 0 && d < D) {
9 s_scales [threadIdx. x] = scales [d];
10 }
11 __syncthreads ();
12
13 if (t < T && d < D) {
14 float val = K [t * D + d];
15 float s = s_scales [threadIdx. x];
16 int q = __float2int_rn (val / s);
17 q = max (-127, min (127, q));
18 K_int8 [t * D + d] = (int8_t) q;
19 }
20}
Listing 6: Tiled CUDA kernel for quantization
Characteristics:
- Scales loaded once per block, then reused from fast shared memory
- Requires thread synchronization (__syncthreads())
- Better memory efficiency for scale access
#### 5.3.3 Coarsened Kernel
Thread coarsening assigns multiple elements to each thread, amortizing thread management overhead and improving instruction-level parallelism.
⬇
1 __global__ void quantize_coarsened (float * K, int8_t * K_int8,
2 float * scales, int T, int D) {
3 int d = blockIdx. x * blockDim. x + threadIdx. x;
4
5 if (d < D) {
6 float s = scales [d];
7
8 for (int t = blockIdx. y; t < T; t += gridDim. y) {
9 float val = K [t * D + d];
10 int q = __float2int_rn (val / s);
11 q = max (-127, min (127, q));
12 K_int8 [t * D + d] = (int8_t) q;
13 }
14 }
15}
Listing 7: Coarsened CUDA kernel for quantization
Characteristics:
- Each thread loads its scale once, processes many elements
- Reduces total number of threads needed
- Better register utilization
#### 5.3.4 Vectorized Kernel
The vectorized kernel uses vector load/store instructions (float4, char4) to process multiple elements per memory transaction, improving memory bandwidth utilization.
⬇
1 __global__ void quantize_vectorized (float * K, int8_t * K_int8,
2 float * scales, int T, int D) {
3 int t = blockIdx. y * blockDim. y + threadIdx. y;
4 int d4 = (blockIdx. x * blockDim. x + threadIdx. x) * 4;
5
6 if (t < T && d4 + 3 < D) {
7 // Vector load: 4 floats at once
8 float4 vals = * reinterpret_cast < float4 *>(& K [t * D + d4]);
9 float4 s = * reinterpret_cast < float4 *>(& scales [d4]);
10
11 char4 q;
12 q. x = (int8_t) max (-127, min (127, __float2int_rn (vals. x / s. x)));
13 q. y = (int8_t) max (-127, min (127, __float2int_rn (vals. y / s. y)));
14 q. z = (int8_t) max (-127, min (127, __float2int_rn (vals. z / s. z)));
15 q. w = (int8_t) max (-127, min (127, __float2int_rn (vals. w / s. w)));
16
17 * reinterpret_cast < char4 *>(& K_int8 [t * D + d4]) = q;
18 }
19}
Listing 8: Vectorized CUDA kernel for quantization
Characteristics:
- 4 $×$ fewer memory transactions
- Better memory coalescing
- Requires $D$ to be divisible by 4 (or handle edge cases)
- Highest bandwidth utilization
## 6 Tests
### 6.1 Hardware Configurations
Table 2: Hardware configuration
| Server | Dell PowerEdge R740 |
| --- | --- |
| GPU | NVIDIA Tesla T4 (16 GB GDDR6) |
| CPU | 2× Intel Xeon Gold 6148 @ 2.40 GHz (40 cores) |
| RAM | 96 GB |
| OS | Ubuntu |
| CUDA Version | 12.0 |
### 6.2 Test Configurations
We benchmark across eight test cases representing different workload sizes, as shown in Table 3.
Table 3: Test configurations for benchmarking
| Small | 2,048 | 128 | Minimal test |
| --- | --- | --- | --- |
| Medium | 16,384 | 256 | Development testing |
| Large | 65,536 | 256 | Extended context |
| Very Large | 131,072 | 256 | Long context |
| Realistic Small | 131,072 | 1,024 | Small LLM workload |
| Realistic Medium | 131,072 | 2,048 | Medium LLM workload |
| Realistic Large | 131,072 | 4,096 | Large LLM workload |
| Realistic V. Large | 131,072 | 8,192 | Very large LLM workload |
## 7 Evaluation
### 7.1 Performance Results
Figure 1 shows GPU speedup across all test configurations. The vectorized kernel consistently outperforms other variants, achieving 1,694× speedup on the largest workload (131K × 8K). Speedup improves with problem size as larger workloads better amortize GPU launch overhead.
<details>
<summary>speedup_comparison.png Details</summary>
### Visual Description
## Bar Chart: GPU Kernel Speedup vs CPU Baseline
### Overview
The chart compares GPU kernel speedup relative to CPU baseline performance across eight test configurations (Small to Real-VL). Four kernel optimization strategies are visualized: Naive (blue), Tiled (red), Coarsened (green), and Vectorized (purple). Speedup values are represented as multiples of CPU performance (x-axis).
### Components/Axes
- **X-axis**: Test Configuration categories (Small, Medium, Large, V.Large, Real-S, Real-M, Real-L, Real-VL)
- **Y-axis**: Speedup over CPU (x) with logarithmic scale (0–2000)
- **Legend**: Located in top-left corner with color-coded labels:
- Naive: Blue
- Tiled: Red
- Coarsened: Green
- Vectorized: Purple
- **Bar Groups**: Four bars per configuration, grouped by optimization strategy
### Detailed Analysis
1. **Small Configuration**
- Naive: ~200x
- Tiled: ~150x
- Coarsened: ~200x
- Vectorized: ~250x
2. **Medium Configuration**
- Naive: ~450x
- Tiled: ~400x
- Coarsened: ~500x
- Vectorized: ~700x
3. **Large Configuration**
- Naive: ~700x
- Tiled: ~600x
- Coarsened: ~900x
- Vectorized: ~1200x
4. **V.Large Configuration**
- Naive: ~700x
- Tiled: ~700x
- Coarsened: ~950x
- Vectorized: ~1150x
5. **Real-S Configuration**
- Naive: ~1100x
- Tiled: ~1250x
- Coarsened: ~1200x
- Vectorized: ~1500x
6. **Real-M Configuration**
- Naive: ~1400x
- Tiled: ~1300x
- Coarsened: ~1300x
- Vectorized: ~1600x
7. **Real-L Configuration**
- Naive: ~1400x
- Tiled: ~1250x
- Coarsened: ~1350x
- Vectorized: ~1650x
8. **Real-VL Configuration**
- Naive: ~1550x
- Tiled: ~1350x
- Coarsened: ~1400x
- Vectorized: ~1700x
### Key Observations
1. **Vectorized Dominance**: Purple bars consistently show highest speedup across all configurations, with Real-VL reaching ~1700x
2. **Scaling Pattern**: Speedup increases with configuration size, particularly pronounced in Vectorized kernels
3. **Tiled vs Naive**: Tiled kernels outperform Naive in all configurations except Small
4. **Coarsened Performance**: Green bars show mid-range performance, often between Tiled and Vectorized
5. **Logarithmic Scale**: Y-axis compression emphasizes relative differences in speedup
### Interpretation
The data demonstrates that vectorized GPU kernels achieve significantly higher performance gains compared to other optimization strategies, with the advantage becoming more pronounced in larger configurations. This suggests that vectorization becomes increasingly effective as problem complexity grows. The Real-VL configuration shows the most dramatic speedup (1700x), indicating that vectorized implementations scale exceptionally well for very large workloads. While Tiled and Coarsened approaches show moderate improvements over Naive implementations, their performance plateaus at lower multiples compared to Vectorized kernels. The consistent pattern across all configurations reinforces the conclusion that vectorization is the most effective optimization strategy for GPU kernel performance relative to CPU baselines.
</details>
Figure 1: GPU kernel speedup comparison across all test configurations.
Interestingly, the tiled kernel does not outperform the naive kernel despite using shared memory. This suggests the scales array is small enough to remain in L2 cache, making explicit caching unnecessary.
Figure 2 shows execution time on a log-log scale. The gap between CPU and GPU spans three orders of magnitude. Even for the largest configuration (1 billion elements), GPU kernels complete in under 50ms compared to 79 seconds on CPU.
<details>
<summary>execution_time.png Details</summary>
### Visual Description
## Line Chart: Execution Time: CPU vs GPU Kernels
### Overview
The chart compares execution times (in seconds) for CPU and various GPU kernel implementations (Naive, Tiled, Coarsened, Vectorized) across matrix sizes ranging from 1 million to 1,000 million elements. Execution time is plotted on a logarithmic y-axis, while matrix size uses a linear x-axis.
### Components/Axes
- **X-axis**: Matrix Size (millions of elements)
- Range: 10⁰ (1M) to 10³ (1,000M)
- Labels: 10⁰, 10¹, 10², 10³
- **Y-axis**: Execution Time (seconds)
- Logarithmic scale: 10⁻³ to 10²
- Labels: 10⁻³, 10⁻², 10⁻¹, 10⁰, 10¹, 10²
- **Legend**:
- CPU: Dark blue circles
- GPU Naive: Blue squares
- GPU Tiled: Red triangles
- GPU Coarsened: Green triangles
- GPU Vectorized: Purple diamonds
### Detailed Analysis
1. **CPU Performance**
- Execution time increases exponentially with matrix size.
- At 1M elements: ~0.05 seconds (5×10⁻²)
- At 100M elements: ~10 seconds (1×10¹)
- At 1,000M elements: ~100 seconds (1×10²)
2. **GPU Performance**
- All GPU methods outperform CPU by orders of magnitude.
- **GPU Naive**:
- 1M: ~1×10⁻⁴ seconds
- 100M: ~1×10⁻¹ seconds
- 1,000M: ~5×10⁻¹ seconds
- **GPU Tiled**:
- 1M: ~5×10⁻⁵ seconds
- 100M: ~2×10⁻¹ seconds
- 1,000M: ~4×10⁻¹ seconds
- **GPU Coarsened**:
- 1M: ~3×10⁻⁵ seconds
- 100M: ~1.5×10⁻¹ seconds
- 1,000M: ~4.5×10⁻¹ seconds
- **GPU Vectorized**:
- 1M: ~2×10⁻⁵ seconds
- 100M: ~1×10⁻¹ seconds
- 1,000M: ~4×10⁻¹ seconds
### Key Observations
- **Exponential vs. Polynomial Growth**:
- CPU execution time grows exponentially (straight line on log scale).
- GPU execution times grow polynomially (curved upward on log scale).
- **GPU Efficiency**:
- At 1,000M elements, GPU methods are **200–500x faster** than CPU.
- GPU Vectorized is the fastest, followed by Tiled, Coarsened, and Naive.
- **Convergence at Large Sizes**:
- GPU methods plateau near 0.5 seconds for 1,000M elements, while CPU reaches 100 seconds.
### Interpretation
The data demonstrates that GPU acceleration drastically reduces execution time for large matrix operations, with GPU Vectorized achieving the best performance. The logarithmic scale emphasizes the exponential advantage of GPUs over CPUs as matrix size increases. While GPU optimization techniques (Tiling, Coarsening) improve performance, their differences diminish at scale, suggesting diminishing returns for advanced optimizations in this context. This highlights the critical role of parallel computing in handling large-scale numerical tasks.
</details>
Figure 2: Execution Time: CPU vs GPU
For realistic LLM workloads (Figure 3), all GPU kernels complete in 6–58 milliseconds, demonstrating that quantization overhead is negligible compared to actual attention computation.
<details>
<summary>realistic_workloads.png Details</summary>
### Visual Description
## Bar Chart: GPU Kernel Performance on Realistic LLM Workloads
### Overview
The chart compares GPU kernel execution times (in milliseconds) across four optimization techniques (Naive, Tiled, Coarsened, Vectorized) for four LLM workload configurations: Real Small (131K×1K), Real Medium (131K×2K), Real Large (131K×4K), and Real V.Large (131K×8K). Execution time increases as workload complexity increases, with optimized kernels outperforming the baseline.
### Components/Axes
- **X-axis (Test Configuration)**:
- Categories: Real Small (131K×1K), Real Medium (131K×2K), Real Large (131K×4K), Real V.Large (131K×8K)
- **Y-axis (Execution Time)**:
- Scale: 0 to 60 milliseconds (milliseconds)
- **Legend**:
- Position: Top-left
- Colors:
- Naive: Blue
- Tiled: Red
- Coarsened: Green
- Vectorized: Purple
### Detailed Analysis
1. **Real Small (131K×1K)**:
- Naive: ~7.5 ms (blue)
- Tiled: ~6.8 ms (red)
- Coarsened: ~7.0 ms (green)
- Vectorized: ~5.5 ms (purple)
2. **Real Medium (131K×2K)**:
- Naive: ~13.0 ms (blue)
- Tiled: ~14.0 ms (red)
- Coarsened: ~14.0 ms (green)
- Vectorized: ~11.5 ms (purple)
3. **Real Large (131K×4K)**:
- Naive: ~27.0 ms (blue)
- Tiled: ~30.0 ms (red)
- Coarsened: ~28.0 ms (green)
- Vectorized: ~23.0 ms (purple)
4. **Real V.Large (131K×8K)**:
- Naive: ~51.0 ms (blue)
- Tiled: ~58.0 ms (red)
- Coarsened: ~57.0 ms (green)
- Vectorized: ~47.0 ms (purple)
### Key Observations
- **Vectorized kernels** consistently achieve the lowest execution times across all workloads, outperforming other methods by 15–25% in larger configurations.
- **Tiled kernels** show mixed performance: slightly better than Naive in small workloads but worse in larger ones.
- **Coarsened kernels** match Tiled in medium workloads but lag behind Vectorized in all cases.
- Execution time scales non-linearly with workload size, with Real V.Large being ~7× slower than Real Small for Naive kernels.
### Interpretation
The data demonstrates that **kernel optimization significantly impacts performance**, particularly for large-scale LLM workloads. Vectorized implementations provide the most substantial speedup, suggesting they are the most efficient for high-complexity tasks. Tiled and Coarsened methods offer moderate improvements but fail to scale effectively. The widening performance gap between Naive and optimized kernels as workload size increases underscores the importance of algorithmic optimization in GPU-accelerated LLM inference.
</details>
Figure 3: GPU Performance on realistic LLM loads.
### 7.2 Reconstruction Error
All GPU kernels produce identical results to the CPU implementation, confirming correctness.
The maximum absolute error is constant at 0.00394 across all configurations (Figure 4, left). This matches the theoretical bound: with input values in [-1, 1] and 255 quantization levels, max error = 1/(2×127) equaling roughly 0.00394.
<details>
<summary>error_metrics.png Details</summary>
### Visual Description
## Line Charts: L2 Reconstruction Error vs Matrix Size and Attention Score Error vs Matrix Size
### Overview
The image contains two line charts comparing error metrics (L2 Reconstruction Error and Attention Score Error) against matrix size (measured in millions of elements). Both charts use logarithmic scales for the x-axis (matrix size) and linear scales for the y-axis (error metrics). The left chart shows a red line representing L2 Reconstruction Error, while the right chart uses a purple line for Attention Score Error.
### Components/Axes
#### Left Chart (L2 Reconstruction Error):
- **X-axis**: Matrix Size (millions of elements)
- Logarithmic scale: 10⁰ (1 million), 10¹ (10 million), 10² (100 million), 10³ (1,000 million).
- **Y-axis**: L2 Error
- Linear scale: 0 to 22.
- **Line**: Red, with markers at each data point.
#### Right Chart (Attention Score Error):
- **X-axis**: Matrix Size (millions of elements)
- Logarithmic scale: 10⁰ (1 million), 10¹ (10 million), 10² (100 million), 10³ (1,000 million).
- **Y-axis**: Mean |Q·K - Q·K|
- Linear scale: 0 to 0.085.
- **Line**: Purple, with markers at each data point.
### Detailed Analysis
#### Left Chart (L2 Reconstruction Error):
- **Trend**: The error increases monotonically with matrix size.
- At 1 million elements: ~1.5.
- At 10 million elements: ~5.
- At 100 million elements: ~9.
- At 1,000 million elements: ~22.
#### Right Chart (Attention Score Error):
- **Trend**: The error remains relatively stable for smaller matrices but spikes sharply at the largest size.
- At 1 million elements: ~0.01.
- At 10 million elements: ~0.015.
- At 100 million elements: ~0.017.
- At 1,000 million elements: ~0.085.
### Key Observations
1. **L2 Reconstruction Error**:
- Errors grow steadily as matrix size increases, with a ~14.5x increase from 100 million to 1,000 million elements.
2. **Attention Score Error**:
- Errors remain stable until the largest matrix size, where they increase by ~400% (from 0.017 to 0.085).
3. **Scale Sensitivity**:
- Both charts use logarithmic x-axes, emphasizing exponential growth in matrix size.
### Interpretation
The data suggests that larger matrix sizes correlate with increased errors in both L2 reconstruction and attention scores. However, the attention score error exhibits a critical threshold effect at the largest matrix size (1,000 million elements), where errors surge dramatically. This could indicate:
- **Scalability Limits**: The model may struggle with computational or memory constraints at extreme scales.
- **Attention Mechanism Instability**: The sharp rise in attention score error might reflect degradation in the model's ability to maintain coherent relationships between elements in very large matrices.
- **Practical Implications**: While the model performs reasonably well up to 100 million elements, its reliability diminishes significantly beyond this point, highlighting the need for optimization or architectural adjustments for large-scale applications.
### Spatial Grounding
- **Legend**: No explicit legend is present, but colors are visually distinct (red for L2, purple for attention score).
- **Positioning**:
- Left chart occupies the left half of the image; right chart occupies the right half.
- Both charts share identical x-axis labels and logarithmic scaling.
### Content Details
- **Data Points**:
- Left chart: (1e0, 1.5), (1e1, 5), (1e2, 9), (1e3, 22).
- Right chart: (1e0, 0.01), (1e1, 0.015), (1e2, 0.017), (1e3, 0.085).
- **Trend Verification**:
- Left chart slopes upward consistently.
- Right chart remains flat until the final data point, where it rises sharply.
### Final Notes
The absence of a legend or explicit error bars limits precision in interpreting variability. However, the clear visual trends suggest a strong relationship between matrix size and error magnitude, with critical implications for model design at scale.
</details>
Figure 4: L2 Reconstruction & Attention Score Error vs Matrix Size.
L2 error grows with matrix size because it sums errors across all elements, this is expected and does not indicate degraded per-element precision.
### 7.3 Attention Score Error
The attention score error measures whether quantization affects downstream computation. Figure 4 (right) show that error scales approximately as $√{D}$ with head dimension.
Crucially, even for the largest configuration (D=8192), mean attention error is only 0.095. Given that attention scores typically span several orders of magnitude before softmax, this error is unlikely to meaningfully alter attention distributions.
### 7.4 Discussion: Kernel Behavior and Architectural Implications
A key observation from the performance results is that the naive kernel performs competitively with, and in some cases slightly better than, the shared-memory tiled kernel. This outcome appears counterintuitive at first, since tiling is often expected to improve performance by reducing global memory traffic. However, this behavior is explained by the memory-access characteristics of the quantization workload.
KV cache quantization is a strictly element-wise operation with no data reuse across threads. Each element is read once, scaled, quantized, and written back. As a result, the computation is overwhelmingly memory-bound, with extremely low arithmetic intensity. In such cases, performance is dominated by global memory bandwidth rather than arithmetic throughput. Because the naive kernel already issues fully coalesced global memory loads and stores along the column dimension, it effectively saturates memory bandwidth without additional optimization.
The tiled kernel introduces shared memory but does not benefit from reuse, since each element is consumed by exactly one thread. Moreover, explicit shared-memory staging introduces additional instructions and synchronization overhead. Although this overhead is small, it offsets any potential benefit, resulting in performance that closely matches, but does not surpass, the naive implementation. The results therefore reinforce the principle that shared memory is beneficial primarily when data reuse exists, not simply as a default optimization.
Thread coarsening improves performance modestly by reducing instruction overhead and improving instruction-level parallelism within each thread. By allowing each thread to process multiple contiguous columns, coarsening reduces loop and address-computation overhead while maintaining coalesced memory access. However, because memory bandwidth remains the dominant bottleneck, the gains from coarsening alone are limited and plateau quickly as the workload scales.
In contrast, the vectorized kernel delivers consistent and substantial performance improvements across all problem sizes. Vectorization reduces the number of memory transactions by loading and storing multiple elements per instruction, effectively increasing memory throughput without changing the algorithmic structure. This optimization directly targets the primary bottleneck of the workload. Unlike shared memory or coarsening, vectorization improves effective bandwidth utilization rather than reducing arithmetic overhead, making it uniquely well-suited for this access pattern.
Taken together, these results confirm that KV cache quantization is fundamentally a memory-bound operation. Once memory accesses are fully coalesced, further optimization requires reducing the number of memory instructions rather than restructuring computation. Vectorization achieves this goal, while tiling and aggressive coarsening offer diminishing returns. Beyond vectorization, additional kernel-level optimizations are unlikely to yield meaningful gains without changes to the data layout or precision format.
<details>
<summary>speedup_scaling.png Details</summary>
### Visual Description
## Line Chart: Speedup Scaling with Problem Size
### Overview
The chart illustrates the performance scaling of four computational methods (Naive, Tiled, Coarsened, Vectorized) across varying matrix sizes (10⁰ to 10³ million elements). Speedup is measured relative to a baseline CPU performance (x-axis), with logarithmic scaling on the x-axis to emphasize exponential growth in problem size.
### Components/Axes
- **X-axis**: Matrix Size (millions of elements), logarithmic scale (10⁰, 10¹, 10², 10³).
- **Y-axis**: Speedup over CPU (x), linear scale (0–1800).
- **Legend**: Located in the bottom-right corner, associating colors and markers:
- **Naive**: Blue circles.
- **Tiled**: Red squares.
- **Coarsened**: Green triangles.
- **Vectorized**: Purple diamonds.
### Detailed Analysis
1. **Vectorized (Purple Diamonds)**:
- Starts at ~200 speedup at 10⁰ elements.
- Increases sharply, reaching ~1,200 at 10¹, ~1,500 at 10², and ~1,700 at 10³.
- Maintains the highest speedup across all matrix sizes.
2. **Coarsened (Green Triangles)**:
- Begins at ~200 at 10⁰, rising to ~900 at 10¹, ~1,250 at 10², and ~1,400 at 10³.
- Shows steady growth but lags behind Vectorized.
3. **Tiled (Red Squares)**:
- Starts at ~200 at 10⁰, peaking at ~700 at 10¹, ~1,250 at 10², and ~1,400 at 10³.
- Converges with Coarsened at 10³ but underperforms at smaller sizes.
4. **Naive (Blue Circles)**:
- Begins at ~200 at 10⁰, rising to ~700 at 10¹, ~1,150 at 10², and ~1,600 at 10³.
- Outperforms Tiled and Coarsened at 10³ but remains the slowest at smaller sizes.
### Key Observations
- **Vectorized dominance**: Consistently achieves the highest speedup, suggesting superior parallelization or algorithmic efficiency.
- **Convergence at large sizes**: Tiled and Coarsened methods perform similarly at 10³ elements (~1,400 speedup), indicating diminishing returns for smaller optimizations.
- **Naive improvement**: Despite being the baseline, Naive methods show significant gains at 10³ elements (~1,600 speedup), though still trailing Vectorized.
### Interpretation
The data highlights that **Vectorized approaches scale most effectively** with problem size, likely due to optimized parallel execution or hardware utilization. Coarsened and Tiled methods exhibit comparable performance at large matrix sizes, suggesting that their optimizations (e.g., memory tiling, data compression) become less impactful as problem size grows. Naive methods, while improving, remain the least efficient, emphasizing the importance of algorithmic refinement for scalability. These trends are critical for selecting computational strategies in high-performance computing, where problem size and resource constraints dictate method choice.
</details>
Figure 5: Speed Scale Ups vs Problem Size.
### 7.5 Test Case Design and Rationale
To ensure that our evaluation reflects both theoretical scalability and real-world deployment scenarios, we carefully selected a diverse set of test configurations that span from synthetic micro-benchmarks to realistic large language model (LLM) workloads.
Small and Medium Matrices.
The smallest configurations (e.g., $T=2048$ , $D=128$ and $T=16384$ , $D=256$ ) serve two purposes. First, they validate correctness and numerical stability under light workloads, where kernel launch overhead and instruction latency dominate. Second, they allow us to observe how different kernel designs behave before GPU parallelism is fully saturated. In these regimes, performance differences between naive, tiled, and coarsened kernels are intentionally modest, helping confirm that optimizations do not introduce regressions at small scale.
Large and Very Large Matrices.
Configurations such as $T=65536$ and $T=131072$ with moderate head dimensions ( $D=256$ ) are representative of intermediate transformer layers operating on long sequences. These tests are large enough to amortize kernel launch overhead and stress global memory bandwidth, making them ideal for comparing memory access patterns across kernel variants. This regime is particularly useful for validating whether shared memory tiling or coarsening provides any benefit over a well-coalesced naive implementation.
Realistic LLM Workloads.
The largest test cases ( $T=131072$ with $D∈\{1024,2048,4096,8192\}$ ) are chosen to closely mirror KV-cache sizes encountered in modern large language models during long-context inference. These configurations approach billions of elements and represent the regime where quantization overhead must be negligible relative to attention computation to be practical. Evaluating performance here allows us to draw conclusions about whether GPU-based quantization is viable in production inference pipelines.
Motivation for Dimensional Scaling.
We vary the head dimension $D$ independently of sequence length $T$ to study how performance and numerical error scale with model width. This is especially important for analyzing attention surrogate error, which depends directly on $D$ . By sweeping $D$ across powers of two up to 8192, we ensure compatibility with vectorized kernels (e.g., float4) while also capturing realistic transformer design choices.
Benchmarking Philosophy.
Finally, these test cases support two complementary evaluation perspectives. First, kernel-only timing isolates architectural effects such as memory coalescing, instruction-level parallelism, and vectorized loads. Second, end-to-end comparisons against a CPU baseline contextualize these gains in terms of practical speedups. Together, this test suite enables us to rigorously validate prior assumptions about tiling, coarsening, and vectorization under both controlled and realistic conditions.
Unit Testing and Validation
To ensure correctness and robustness of both the CPU and GPU implementations, we developed a comprehensive unit testing framework covering functionality, numerical accuracy, and edge cases. The test suite consists of 25 individual tests and is executed independently of the benchmarking harness.
Basic structural tests verify correct matrix allocation, dimension handling, and memory initialization for both FP32 and INT8 matrix types. Randomized fill routines are validated to ensure values remain within specified bounds, preventing silent data corruption at later stages.
Numerical correctness is tested using identity checks for L2 error, maximum absolute error, and attention score error, all of which must evaluate to zero when comparing a matrix against itself. Additional deterministic tests validate scale computation, quantization rounding behavior, and dequantization accuracy on hand-constructed inputs where expected outputs are analytically known.
GPU kernels are validated by direct comparison against the CPU reference implementation. For each kernel variant (naive, tiled, coarsened, and vectorized), quantized outputs are compared element-wise to the CPU result, allowing a tolerance of $± 1$ to account for minor rounding differences between CPU and GPU floating-point behavior. Cross-kernel consistency tests further verify that all GPU variants produce identical outputs for the same inputs.
Edge cases such as $1× 1$ matrices, minimal vectorizable dimensions, and structured input patterns (all zeros, all ones, alternating signs) are explicitly tested to ensure correctness under degenerate conditions. Finally, stress tests on larger matrices confirm stability and correctness under non-trivial workloads.
This validation framework provides strong confidence that all reported performance results correspond to correct and numerically consistent behavior, and that observed differences between kernel variants arise from architectural and memory-access characteristics rather than implementation errors.
## 8 Conclusion
This work implemented and evaluated GPU-accelerated INT8 quantization for KV cache compression in large language models. We developed a CPU baseline and four CUDA kernel variants—naive, tiled, coarsened, and vectorized—and evaluated them across realistic LLM workload sizes.
Performance: The vectorized kernel achieves up to 1,694 $×$ speedup over the CPU baseline, completing quantization of 1 billion elements in under 50ms.
Memory reduction: INT8 quantization provides 4 $×$ memory savings compared to FP32 storage, enabling longer context windows and larger batch sizes.
Accuracy: Maximum per-element error is bounded at 0.004, and attention score error remains below 0.1 even for 8K-dimensional heads—demonstrating minimal impact on model behavior.
Practical viability: For realistic LLM workloads, quantization overhead is just 6–58ms, making it negligible compared to actual inference time.
Our analysis reveals that KV cache quantization is fundamentally memory-bound, and vectorization provides the most effective optimization by reducing memory transaction overhead. These findings provide practical guidance for implementing efficient quantization in production LLM inference systems.
### 8.1 Limitations
This work focuses on per-channel INT8 quantization and does not explore several promising directions:
- Lower bit-widths: INT4 or INT2 quantization could provide 8 $×$ or 16 $×$ compression at the cost of increased error.
- Mixed-precision strategies: Selective quantization of less critical layers or tokens could balance memory savings with accuracy.
- Dynamic quantization: Adaptive scale computation during inference could improve precision for out-of-distribution inputs.
- End-to-end evaluation: We measure reconstruction and attention score error but do not evaluate impact on downstream task performance (e.g., perplexity on language modeling benchmarks).
- Multi-GPU scaling: Our implementation targets single-GPU scenarios and does not address distributed inference workloads.
### 8.2 Future Work
Several promising directions extend this work toward more comprehensive optimization and real-world deployment evaluation. First, exploring additional GPU optimization strategies could further improve quantization performance. Warp-level primitives such as __shfl_down_sync could enable more efficient reduction operations during scale computation, while persistent kernels that maintain thread blocks across multiple quantization operations could amortize launch overhead for streaming workloads. Additionally, investigating GPU-specific numeric formats such as FP8, which recent architectures natively support, could provide hardware-accelerated alternatives to INT8 with potentially better accuracy-performance tradeoffs.
Beyond kernel-level optimizations, the most critical next step is evaluating these quantization techniques within complete transformer blocks during actual inference. This requires integrating our kernels into existing inference engines such as vLLM, TensorRT-LLM, or HuggingFace Transformers, then measuring end-to-end latency, throughput, and memory consumption on real workloads. Such integration would reveal how quantization overhead interacts with attention computation, feedforward layers, and memory bandwidth contention in production scenarios. Furthermore, testing on diverse models ranging from smaller variants like GPT-2 to larger models approaching 70B parameters would establish whether our performance characteristics hold across different architectural scales and whether quantization enables previously infeasible batch sizes or context lengths.
Finally, comprehensive accuracy evaluation on standard benchmarks remains essential. While our attention score error analysis suggests minimal impact, measuring perplexity degradation on language modeling tasks, accuracy on question-answering benchmarks, and quality on generation tasks would quantify the practical tradeoff between memory savings and model capability. This end-to-end evaluation across GPU optimizations, inference system integration, and task-specific metrics would provide the complete picture necessary for deploying quantized KV caches in production LLM serving systems.
## Acknowledgments
We would like to thank Mohnish Sonsare for his contributions to this work, and Professor James Mooney for his guidance and support throughout this project.
## References
- [1] Ashish Vaswani, Noam Shazeer, Niki Parmar, Jakob Uszkoreit, Llion Jones, Aidan N Gomez, Łukasz Kaiser, and Illia Polosukhin. Attention is all you need. In Advances in Neural Information Processing Systems, pages 5998–6008, 2017.
- [2] Tri Dao, Daniel Y Fu, Stefano Ermon, Atri Rudra, and Christopher Ré. FlashAttention: Fast and memory-efficient exact attention with IO-awareness. In Advances in Neural Information Processing Systems, 2022.
- [3] Tri Dao. FlashAttention-2: Faster attention with better parallelism and work partitioning. In International Conference on Learning Representations, 2023.
- [4] Woosuk Kwon, Zhuohan Li, Siyuan Zhuang, Ying Sheng, Lianmin Zheng, Cody Hao Yu, Joseph E. Gonzalez, Hao Zhang, and Ion Stoica. Efficient memory management for large language model serving with PagedAttention. In Symposium on Operating Systems Principles (SOSP), 2023.
- [5] Tim Dettmers, Mike Lewis, Younes Belkada, and Luke Zettlemoyer. LLM.int8(): 8-bit matrix multiplication for transformers at scale. In Advances in Neural Information Processing Systems, 2022.
- [6] Guangxuan Xiao, Ji Lin, Mickael Seznec, Hao Wu, Julien Demouth, and Song Han. SmoothQuant: Accurate and efficient post-training quantization for large language models. In International Conference on Machine Learning, 2023.
- [7] Zirui Liu, Jiayi Yuan, Hongye Jin, Shaochen Zhong, Zhaozhuo Xu, Vladimir Braverman, Beidi Chen, and Xia Hu. KIVI: A tuning-free asymmetric 2bit quantization for KV cache. arXiv preprint arXiv:2402.02750, 2024.
- [8] Coleman Hooper, Sehoon Kim, Hiva Mohammadzadeh, Michael W. Mahoney, Yakun Sophia Shao, Kurt Keutzer, and Amir Gholami. KVQuant: Towards 10 million context length LLM inference with KV cache quantization. arXiv preprint arXiv:2401.18079, 2024.
- [9] Ying Sheng, Lianmin Zheng, Binhang Yuan, Zhuohan Li, Max Ryabinin, Beidi Chen, Percy Liang, Christopher Ré, Ion Stoica, and Ce Zhang. FlexGen: High-throughput generative inference of large language models with a single GPU. In International Conference on Machine Learning, 2023.
- [10] Reiner Pope, Sholto Douglas, Aakanksha Chowdhery, Jacob Devlin, James Bradbury, Jonathan Heek, Kefan Xiao, Shivani Agrawal, and Jeff Dean. Efficiently scaling transformer inference. In MLSys, 2022.