Build a GPU KV-Cache Memory Manager

Introduction

30 Second Summary

When your operating system runs dozens of apps at once, it does not reserve one giant chunk of memory per app. Instead, it breaks memory into small pages and hands them out on demand. Production LLM servers face the exact same challenge when juggling thousands of user requests on a single GPU.

In this project, you will build a C++/CUDA memory management engine that treats GPU VRAM like an OS virtual memory system. You will allocate memory in fixed-size blocks, map them through page tables, and simulate concurrent LLM sequences with zero external fragmentation.

What You'll Build

A framework-free GPU memory manager inspired by vLLM's PagedAttention system that pre-allocates VRAM, divides it into uniform blocks, and uses page tables to map logical token sequences to physical memory on demand.

By the end of this project, you'll have:

• A physical block allocator that manages GPU memory in uniform blocks with a free-list, eliminating external fragmentation by design.
• A page-table system that maps logical token sequences to non-contiguous physical GPU memory blocks, just like an OS maps virtual addresses to physical RAM.
• A multi-sequence simulation harness that runs 100 concurrent LLM sequence completions, demonstrating dynamic allocation and deallocation with near-zero memory waste.
• Secret Mission: implement copy-on-write block sharing for beam search, where forked sequences share physical blocks until one writes to them.

Are there any prerequisites?

This project assumes intermediate C++ knowledge (templates, pointers, manual memory management) and basic familiarity with how LLMs generate tokens one at a time. You do not need prior CUDA experience, but comfort with command-line build tools will help.

You will also need an NVIDIA GPU with Compute Capability 5.0 or higher. Most NVIDIA GPUs from the last eight years meet this requirement.

Before We Start

To build a GPU memory manager, you need a development environment that can compile CUDA code and run it on your NVIDIA GPU.

In this step, you will install the CUDA Toolkit (NVIDIA's compiler and runtime for GPU programming), CMake (the build system that will compile your project), and verify that your GPU is detected and ready.

In this step, get ready to:

• Install CUDA Toolkit 13.3 with Visual Studio 2022 integration.
• Install CMake 3.30+ and add it to your system PATH.
• Verify your NVIDIA GPU is detected and note its Compute Capability.

Install CUDA Toolkit 13.3

The CUDA Toolkit provides nvcc, the NVIDIA CUDA compiler that turns your .cu source files into GPU-executable code.

CUDA requires a C++ host compiler. On Windows, that means Visual Studio 2022 must be installed first so the CUDA installer can hook into it.

• If you do not already have Visual Studio 2022 installed, download it from visualstudio.microsoft.com.
• During the Visual Studio installer, select the Desktop development with C++ workload.
• Complete the Visual Studio installation before continuing.

Why Visual Studio 2022?

CUDA's compiler (nvcc) delegates host-side C++ compilation to your system's C++ compiler. On Windows, that is cl.exe from Visual Studio. CUDA 13.3 officially supports Visual Studio 2022.

With Visual Studio 2022 in place, you can now install the CUDA Toolkit.

• Open your browser and navigate to developer.nvidia.com/cuda-downloads.
• Select your operating system (Windows), architecture (x86_64), and version, then download the installer.
• Run the installer. When prompted for components, select at minimum: CUDA Runtime, CUDA compiler (nvcc), and Visual Studio Integration.
• Complete the installation and close the installer.

What about the GPU driver?

Starting with CUDA 13.1, the Windows display driver is no longer bundled with the CUDA Toolkit. Your NVIDIA GPU driver should already be installed since your GPU is working. If your driver is outdated, download the latest from the NVIDIA Driver Downloads page.

Now verify the installation. Open a new Command Prompt (it must be new so it picks up the updated PATH).

• Press the Windows key to open your search bar.
• Type cmd and press Enter to open Command Prompt.
• Check your CUDA compiler version by running:

nvcc -V

✔️ I see release 13.3

You should see output that includes a line like Cuda compilation tools, release 13.3 followed by a version number. CUDA is installed and ready.

ⓧ I see an older release version

If you see a release number lower than 13.3, you have an older CUDA Toolkit installed.

• Return to developer.nvidia.com/cuda-downloads and download the 13.3 installer.
• Run the installer. It will upgrade your existing installation.
• Close your Command Prompt and open a new one, then run nvcc -V again to confirm release 13.3 appears.

ⓧ Command not found

This means CUDA is not installed or is not on your PATH.

• Confirm you completed the CUDA installer steps above.
• Make sure you opened a new Command Prompt after installation. Existing terminals do not pick up PATH changes.
• If the issue persists, search your system for nvcc.exe (typically in C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.3\bin). Add that directory to your system PATH manually, then open a new Command Prompt and retry.

Install CMake

CMake is the build system that will configure and compile your CUDA project. It treats CUDA as a first-class language, so it knows how to invoke nvcc automatically.

• Open your browser and navigate to cmake.org.
• Download the latest CMake installer for Windows (version 3.30 or higher).
• Run the installer. When prompted, select Add CMake to the system PATH for all users.
• Complete the installation.

Open a new Command Prompt (or use the one already open) and verify CMake by running:

• Check your CMake version by running:

cmake --version

✔️ I see version 3.30 or higher

You should see a line starting with cmake version followed by a number like 3.30 or higher. CMake is ready.

ⓧ I see a version lower than 3.30

You have an older CMake installed that may not support the CUDA features this project needs.

• Return to cmake.org and download the latest installer.
• Run the installer. It will replace your existing version.
• Open a new Command Prompt and run cmake --version again to confirm the version is 3.30 or higher.

ⓧ Command not found

CMake is not installed or not on your PATH.

• Confirm you completed the CMake installer steps above.
• Make sure you selected Add CMake to the system PATH during installation. If you missed it, re-run the installer and select that option.
• Open a new Command Prompt (existing terminals do not pick up PATH changes) and try again.

Verify your GPU

The final check confirms that your system can see your NVIDIA GPU and reports its Compute Capability. This project requires Compute Capability 5.0 or higher.

• In your Command Prompt, run:

nvidia-smi

You should see a table showing your GPU name, driver version, and CUDA version. Note your GPU name for the next step.

How do I find my Compute Capability?

The nvidia-smi output shows your GPU name (for example, "GeForce RTX 3060" or "RTX 4090"). Look up your GPU's Compute Capability on NVIDIA's developer page at developer.nvidia.com/cuda-gpus.

Any GPU from the Maxwell generation (2014) or newer has Compute Capability 5.0+. If you have a GTX 900 series, GTX 1000 series, RTX 2000/3000/4000/5000 series, or any Tesla/Quadro from 2014 onward, you are good to go.

✔️ I see my GPU listed

Your GPU is detected and ready. Note the GPU name from the output. You will use this in the next step to set the correct architecture flag in your build configuration.

ⓧ Command not found or no GPU detected

If nvidia-smi is not found or shows no GPU, your NVIDIA driver may not be installed correctly.

• Download the latest driver for your GPU from nvidia.com/drivers.
• Install the driver, restart your computer, then open a new Command Prompt and run nvidia-smi again.
• If you still see no GPU, confirm your monitor is connected to the NVIDIA card (not integrated graphics) and that the card is properly seated in its PCIe slot.

Your development environment is ready. Next up, you will create the project structure and write a CMake configuration that compiles CUDA code targeting your specific GPU architecture.

Set Up the Project Structure and CMake Build

With your tools installed, you need a project structure before you can write any CUDA code. CMake needs to know where your source files live and how to compile them for your specific GPU.

In this step, you will create the folder layout, write a CMakeLists.txt that configures the CUDA build, and verify everything works with a minimal program that prints your GPU info.

In this step, get ready to:

• Create the project folder structure with source and header directories.
• Write a CMake configuration targeting your GPU architecture.
• Build and run a minimal CUDA program that confirms your setup works.

Create the project directory

Your project needs a clean folder structure so CMake can locate sources and headers. You will create a root folder on your Desktop with two subdirectories.

• Move to your Desktop by running this command:

cd ~/Desktop

• Create the project folder and move into it by running these commands:

mkdir paged-kv-cache
cd paged-kv-cache

• Create the src/ and include/ subdirectories by running:

mkdir src
mkdir include

• Confirm the folders exist by running:

ls

You should see src and include listed in the output.

Why two subdirectories?

The src/ folder holds implementation files (.cu CUDA source files). The include/ folder holds header files (.h and .cuh) that declare your classes and structs.

This separation keeps your build clean and makes it easy for CMake to set up include paths.

Write the CMake build configuration

CMake uses a CMakeLists.txt file to know which compiler to use, what language standard to target, and which GPU architecture to compile for.

• Create a new file called CMakeLists.txt in the root of your paged-kv-cache folder and paste the following code:

cmake_minimum_required(VERSION 3.18)
project(paged_kv_cache LANGUAGES CXX CUDA)

set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CUDA_STANDARD 17)

# Detect GPU architecture - default to sm_75 (Turing)
if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
    set(CMAKE_CUDA_ARCHITECTURES 75)
endif()

add_executable(paged_kv_cache
    src/main.cu
)

target_include_directories(paged_kv_cache PRIVATE include)

What does this configuration do?

• cmake_minimum_required(VERSION 3.18) ensures CMake supports CUDA as a first-class language with the architectures property.
• project(paged_kv_cache LANGUAGES CXX CUDA) tells CMake this project uses both C++ and CUDA compilers.
• CMAKE_CUDA_ARCHITECTURES 75 targets Turing GPUs (compute capability 7.5). The if(NOT DEFINED ...) guard lets you override this from the command line if needed.
• target_include_directories tells the compiler to look in include/ for header files.

Match your GPU architecture

The value 75 corresponds to Turing GPUs (RTX 20-series). If your GPU uses a different architecture, update this number to match. Common values: 86 for Ampere (RTX 30-series), 89 for Ada Lovelace (RTX 40-series).

You can find your Compute Capability from the nvidia-smi output you checked in the previous step.

• Save CMakeLists.txt.

Write a minimal CUDA program and verify the build

Before writing any real engine code, you need to prove that CMake can compile a CUDA file and that the resulting binary talks to your GPU.

• Create a new file called main.cu inside the src/ folder and paste the following code:

#include <cuda_runtime.h>
#include <cstdio>

int main() {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    printf("GPU: %s | VRAM: %zu MB | Compute: %d.%d\n",
           prop.name, prop.totalGlobalMem / (1024 * 1024),
           prop.major, prop.minor);
    return 0;
}

What does this code do?

• cudaGetDeviceProperties(&prop, 0) queries the first GPU (device 0) for its hardware properties.
• prop.name gives the GPU model name. prop.totalGlobalMem gives total VRAM in bytes.
• prop.major and prop.minor report the Compute Capability, confirming CUDA can communicate with your hardware.

• Save src/main.cu.
• Generate the build files by running this command from the paged-kv-cache root folder:

cmake -B build

• Compile the project by running:

cmake --build build --config Release

What do these commands do?

cmake -B build creates a build/ directory and generates Visual Studio project files that know how to invoke nvcc for your CUDA source.

cmake --build build --config Release invokes the actual compilers (nvcc for CUDA, cl.exe for C++ host code) to produce the executable.

• Run the compiled program by executing:

.\build\Release\paged_kv_cache.exe

You should see output showing your GPU name, VRAM size in MB, and Compute Capability, similar to:

GPU: NVIDIA GeForce RTX 3080 | VRAM: 10240 MB | Compute: 8.6

Build failing or no output?

• If CMake says it cannot find a CUDA compiler, make sure nvcc is on your system PATH. Open a new terminal and try nvcc -V to confirm.
• If you see architecture-related errors, double check that the CMAKE_CUDA_ARCHITECTURES value in CMakeLists.txt matches your GPU's Compute Capability.
• On Windows, ensure you are running from a terminal where Visual Studio's compiler (cl.exe) is available. The Developer Command Prompt for VS 2022 works best.

Help me debug my CMake CUDA build errors.

✔️ Awesome, I've got everything!

Great. Double check you have saved both files before moving on.

ⓧ I'd like to double check the full code

Here is the complete CMakeLists.txt:

cmake_minimum_required(VERSION 3.18)
project(paged_kv_cache LANGUAGES CXX CUDA)

set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CUDA_STANDARD 17)

# Detect GPU architecture - default to sm_75 (Turing)
if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
    set(CMAKE_CUDA_ARCHITECTURES 75)
endif()

add_executable(paged_kv_cache
    src/main.cu
)

target_include_directories(paged_kv_cache PRIVATE include)

Here is the complete src/main.cu:

#include <cuda_runtime.h>
#include <cstdio>

int main() {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    printf("GPU: %s | VRAM: %zu MB | Compute: %d.%d\n",
           prop.name, prop.totalGlobalMem / (1024 * 1024),
           prop.major, prop.minor);
    return 0;
}

Your build environment is proven and ready. Next up, you will write the physical block allocator that carves your GPU memory pool into uniform blocks.

Build the Physical Block Allocator

In the previous step, you set up your CUDA build environment and confirmed your GPU is ready. Now it's time to build the core memory system that makes paged KV-Cache work.

The block allocator is like an operating system's physical memory manager. It claims one large chunk of GPU memory upfront, then divides it into uniform blocks that can be handed out individually on demand. This eliminates external fragmentation entirely because every block is the same size.

In this step, get ready to:

• Define the PhysicalBlock struct and BlockAllocator class with GPU memory pool pre-allocation.
• Implement allocate() and free() methods with bounds checking and reference counting.
• Write CUDA test kernels that verify the allocator works correctly on the GPU.

Define the PhysicalBlock struct and BlockAllocator class

The allocator needs to do three things at construction time: claim a big slab of GPU memory, divide it into equal-sized blocks, and track which blocks are free. This mirrors how an OS kernel manages physical page frames.

• Create a new file called block_allocator.cuh inside your include folder.
• Paste the following code into the file. This defines the PhysicalBlock struct and the BlockAllocator constructor that pre-allocates the GPU memory pool:

#ifndef BLOCK_ALLOCATOR_CUH
#define BLOCK_ALLOCATOR_CUH

#include <cuda_runtime.h>
#include <vector>
#include <stdexcept>
#include <cstdint>
#include <cstdio>

struct PhysicalBlock {
    int block_id;
    int ref_count;
    float* device_ptr;
};

class BlockAllocator {
public:
    BlockAllocator(size_t total_pool_bytes, size_t block_size_bytes)
        : block_size_bytes_(block_size_bytes), pool_base_(nullptr) {

        num_blocks_ = total_pool_bytes / block_size_bytes;
        if (num_blocks_ == 0) {
            throw std::runtime_error("Pool too small for even one block");
        }

        cudaError_t err = cudaMalloc(&pool_base_, num_blocks_ * block_size_bytes_);
        if (err != cudaSuccess) {
            throw std::runtime_error(std::string("cudaMalloc failed: ") + cudaGetErrorString(err));
        }

        // Initialize free list with all block IDs
        free_list_.reserve(num_blocks_);
        for (int i = static_cast<int>(num_blocks_) - 1; i >= 0; --i) {
            free_list_.push_back(i);
        }

        // Initialize block metadata
        blocks_.resize(num_blocks_);
        for (size_t i = 0; i < num_blocks_; ++i) {
            blocks_[i].block_id = static_cast<int>(i);
            blocks_[i].ref_count = 0;
            blocks_[i].device_ptr = reinterpret_cast<float*>(
                reinterpret_cast<uint8_t*>(pool_base_) + i * block_size_bytes_
            );
        }
    }

    ~BlockAllocator() {
        if (pool_base_) {
            cudaFree(pool_base_);
        }
    }

    // Non-copyable
    BlockAllocator(const BlockAllocator&) = delete;
    BlockAllocator& operator=(const BlockAllocator&) = delete;

What does this code do?

The PhysicalBlock struct holds metadata for each block: its ID, a reference count (used later for copy-on-write), and a raw pointer to its location in GPU memory.

The constructor does the heavy lifting. It calls cudaMalloc once to claim the entire memory pool, then populates the free-list in reverse order so block 0 is allocated first. Each block's device_ptr is computed as an offset from the pool base using byte arithmetic.

Why one big cudaMalloc instead of many small ones?

Calling cudaMalloc is expensive. It requires a driver round-trip and can stall the GPU pipeline. By allocating once and sub-dividing manually, the allocator avoids this overhead entirely during runtime.

This is the same pattern production systems like vLLM use. The block engine claims a contiguous chunk of GPU DRAM upfront, then manages sub-allocation through the free-list.

Implement allocate() and free()

With the pool in place, the allocator needs methods to hand out blocks and reclaim them. The free-list acts as a stack. Allocation pops from the back, and freeing pushes back.

• In the same block_allocator.cuh file, paste the following code directly below the line BlockAllocator& operator=(const BlockAllocator&) = delete; to complete the class:

    int allocate() {
        if (free_list_.empty()) {
            return -1; // No free blocks
        }
        int block_id = free_list_.back();
        free_list_.pop_back();
        blocks_[block_id].ref_count = 1;
        return block_id;
    }

    void free(int block_id) {
        if (block_id < 0 || block_id >= static_cast<int>(num_blocks_)) {
            return;
        }
        blocks_[block_id].ref_count--;
        if (blocks_[block_id].ref_count <= 0) {
            blocks_[block_id].ref_count = 0;
            free_list_.push_back(block_id);
        }
    }

    void increment_ref(int block_id) {
        if (block_id >= 0 && block_id < static_cast<int>(num_blocks_)) {
            blocks_[block_id].ref_count++;
        }
    }

    int get_ref_count(int block_id) const {
        if (block_id < 0 || block_id >= static_cast<int>(num_blocks_)) return 0;
        return blocks_[block_id].ref_count;
    }

    float* get_block_ptr(int block_id) const {
        if (block_id < 0 || block_id >= static_cast<int>(num_blocks_)) return nullptr;
        return blocks_[block_id].device_ptr;
    }

    size_t num_free_blocks() const { return free_list_.size(); }
    size_t num_total_blocks() const { return num_blocks_; }
    size_t num_allocated_blocks() const { return num_blocks_ - free_list_.size(); }
    size_t block_size_bytes() const { return block_size_bytes_; }
    void* pool_base() const { return pool_base_; }

private:
    void* pool_base_;
    size_t num_blocks_;
    size_t block_size_bytes_;
    std::vector<int> free_list_;
    std::vector<PhysicalBlock> blocks_;
};

#endif // BLOCK_ALLOCATOR_CUH

What does this code do?

• allocate() pops the last ID from the free-list and sets the block's ref_count to 1. Returns -1 if the pool is exhausted.
• free() decrements ref_count. Only when it hits zero does the block return to the free-list. This supports reference counting for later copy-on-write.
• increment_ref() and get_ref_count() manage shared ownership. These become critical in the Secret Mission when multiple sequences share prefix blocks.
• get_block_ptr() returns the raw GPU pointer for a block ID, letting CUDA kernels read and write block memory directly.

• Save include/block_allocator.cuh.

✔️ Awesome, I've got everything!

Great. Double check you've saved the file before moving on.

ⓧ I'd like to double check the full code

#ifndef BLOCK_ALLOCATOR_CUH
#define BLOCK_ALLOCATOR_CUH

#include <cuda_runtime.h>
#include <vector>
#include <stdexcept>
#include <cstdint>
#include <cstdio>

struct PhysicalBlock {
    int block_id;
    int ref_count;
    float* device_ptr;
};

class BlockAllocator {
public:
    BlockAllocator(size_t total_pool_bytes, size_t block_size_bytes)
        : block_size_bytes_(block_size_bytes), pool_base_(nullptr) {

        num_blocks_ = total_pool_bytes / block_size_bytes;
        if (num_blocks_ == 0) {
            throw std::runtime_error("Pool too small for even one block");
        }

        cudaError_t err = cudaMalloc(&pool_base_, num_blocks_ * block_size_bytes_);
        if (err != cudaSuccess) {
            throw std::runtime_error(std::string("cudaMalloc failed: ") + cudaGetErrorString(err));
        }

        // Initialize free list with all block IDs
        free_list_.reserve(num_blocks_);
        for (int i = static_cast<int>(num_blocks_) - 1; i >= 0; --i) {
            free_list_.push_back(i);
        }

        // Initialize block metadata
        blocks_.resize(num_blocks_);
        for (size_t i = 0; i < num_blocks_; ++i) {
            blocks_[i].block_id = static_cast<int>(i);
            blocks_[i].ref_count = 0;
            blocks_[i].device_ptr = reinterpret_cast<float*>(
                reinterpret_cast<uint8_t*>(pool_base_) + i * block_size_bytes_
            );
        }
    }

    ~BlockAllocator() {
        if (pool_base_) {
            cudaFree(pool_base_);
        }
    }

    // Non-copyable
    BlockAllocator(const BlockAllocator&) = delete;
    BlockAllocator& operator=(const BlockAllocator&) = delete;

    int allocate() {
        if (free_list_.empty()) {
            return -1; // No free blocks
        }
        int block_id = free_list_.back();
        free_list_.pop_back();
        blocks_[block_id].ref_count = 1;
        return block_id;
    }

    void free(int block_id) {
        if (block_id < 0 || block_id >= static_cast<int>(num_blocks_)) {
            return;
        }
        blocks_[block_id].ref_count--;
        if (blocks_[block_id].ref_count <= 0) {
            blocks_[block_id].ref_count = 0;
            free_list_.push_back(block_id);
        }
    }

    void increment_ref(int block_id) {
        if (block_id >= 0 && block_id < static_cast<int>(num_blocks_)) {
            blocks_[block_id].ref_count++;
        }
    }

    int get_ref_count(int block_id) const {
        if (block_id < 0 || block_id >= static_cast<int>(num_blocks_)) return 0;
        return blocks_[block_id].ref_count;
    }

    float* get_block_ptr(int block_id) const {
        if (block_id < 0 || block_id >= static_cast<int>(num_blocks_)) return nullptr;
        return blocks_[block_id].device_ptr;
    }

    size_t num_free_blocks() const { return free_list_.size(); }
    size_t num_total_blocks() const { return num_blocks_; }
    size_t num_allocated_blocks() const { return num_blocks_ - free_list_.size(); }
    size_t block_size_bytes() const { return block_size_bytes_; }
    void* pool_base() const { return pool_base_; }

private:
    void* pool_base_;
    size_t num_blocks_;
    size_t block_size_bytes_;
    std::vector<int> free_list_;
    std::vector<PhysicalBlock> blocks_;
};

#endif // BLOCK_ALLOCATOR_CUH

Write a unit test with CUDA kernels

The allocator manages memory on the host side, but the actual reads and writes happen on the GPU. To prove the system works, you need CUDA kernels that write patterns into allocated blocks and verify them.

• Open src/main.cu and replace its entire contents with the following. Start with the includes, CUDA kernels, and a helper function:

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>
#include <ctime>
#include <vector>
#include "block_allocator.cuh"

// Simple kernel to write a pattern to a block's memory
__global__ void write_pattern_kernel(float* block_ptr, int num_elements, float value) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < num_elements) {
        block_ptr[idx] = value;
    }
}

// Kernel to verify a pattern
__global__ void verify_pattern_kernel(float* block_ptr, int num_elements,
                                      float expected, int* errors) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < num_elements) {
        if (block_ptr[idx] != expected) {
            atomicAdd(errors, 1);
        }
    }
}

void print_separator() {
    printf("============================================================\n");
}

What do these kernels do?

• write_pattern_kernel fills every float in a block with a known value. Each GPU thread handles one element.
• verify_pattern_kernel checks every float against the expected value and atomically increments an error counter if any mismatch is found. This proves the memory at each block's pointer is valid and writable.

• Below print_separator(), paste the test function that exercises the allocator end-to-end:

void test_block_allocator() {
    printf("\n[TEST] Block Allocator - Basic Allocation and Free\n");
    print_separator();

    size_t pool_size = 64 * 1024 * 1024;
    size_t block_size = 16 * 128 * sizeof(float);
    BlockAllocator alloc(pool_size, block_size);

    printf("  Pool size: %zu MB\n", pool_size / (1024 * 1024));
    printf("  Block size: %zu bytes\n", block_size);
    printf("  Total blocks: %zu\n", alloc.num_total_blocks());
    printf("  Free blocks: %zu\n", alloc.num_free_blocks());

    std::vector<int> allocated;
    for (int i = 0; i < 5; ++i) {
        int id = alloc.allocate();
        if (id >= 0) allocated.push_back(id);
    }
    printf("  After allocating 5: free=%zu, allocated=%zu\n",
           alloc.num_free_blocks(), alloc.num_allocated_blocks());

    int elements_per_block = static_cast<int>(block_size / sizeof(float));
    int threads = 256;
    int blocks = (elements_per_block + threads - 1) / threads;

    int* d_errors;
    cudaMalloc(&d_errors, sizeof(int));

    for (size_t i = 0; i < allocated.size(); ++i) {
        float* ptr = alloc.get_block_ptr(allocated[i]);
        float pattern = static_cast<float>(i + 1) * 1.5f;
        write_pattern_kernel<<<blocks, threads>>>(ptr, elements_per_block, pattern);
    }
    cudaDeviceSynchronize();

    int total_errors = 0;
    for (size_t i = 0; i < allocated.size(); ++i) {
        float* ptr = alloc.get_block_ptr(allocated[i]);
        float pattern = static_cast<float>(i + 1) * 1.5f;
        int zero = 0;
        cudaMemcpy(d_errors, &zero, sizeof(int), cudaMemcpyHostToDevice);
        verify_pattern_kernel<<<blocks, threads>>>(ptr, elements_per_block, pattern, d_errors);
        int err_count = 0;
        cudaMemcpy(&err_count, d_errors, sizeof(int), cudaMemcpyDeviceToHost);
        total_errors += err_count;
    }
    printf("  Pattern verification errors: %d\n", total_errors);

    for (int id : allocated) {
        alloc.free(id);
    }
    printf("  After freeing all: free=%zu\n", alloc.num_free_blocks());

    cudaFree(d_errors);
    printf("  [PASS] Block allocator test complete\n");
}

What does this test do?

• Creates a 64 MB pool divided into blocks of 16 tokens * 128 floats (8192 bytes each). This simulates a KV-Cache where each block stores 16 tokens of 128-dimensional key/value vectors.
• Allocates 5 blocks, writes a unique float pattern to each using the GPU kernel, then reads back and verifies every value matches.
• Frees all 5 blocks and confirms the free-list count is restored. This proves blocks are properly recycled.

• Finally, paste the updated main() function below test_block_allocator():

int main() {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    printf("=== Paged KV-Cache Engine ===\n");
    printf("GPU: %s | VRAM: %zu MB | Compute: %d.%d\n\n",
           prop.name, prop.totalGlobalMem / (1024 * 1024),
           prop.major, prop.minor);

    test_block_allocator();

    printf("\n=== All tests passed! ===\n");
    return 0;
}

What changed in main()?

The new main() still prints GPU info, but now calls test_block_allocator() to run the full allocation test. A banner and closing message frame the output so you can see all tests at a glance as you add more in later steps.

• Save src/main.cu.
• Rebuild and run the project by running these commands:

cmake --build build --config Release

.\build\Release\paged_kv_cache.exe

• You should see output showing your GPU info, the pool and block statistics, "Pattern verification errors: 0", and "[PASS] Block allocator test complete" at the end.

Build failing or seeing errors?

• If you see "cannot open include file 'block_allocator.cuh'", confirm your file is at include/block_allocator.cuh (not in src/) and that your CMakeLists.txt has the target_include_directories(paged_kv_cache PRIVATE include) line.
• If you see pattern verification errors greater than 0, your GPU memory may be corrupted. Try deleting the build folder, re-running cmake -B build, and rebuilding.
• If cudaMalloc fails, another process may be using your GPU memory. Close any other GPU-heavy applications and try again.

Help me debug my block allocator build error.

✔️ Awesome, I've got everything!

Great. Make sure you've saved both files before moving on.

ⓧ I'd like to double check the full code

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>
#include <ctime>
#include <vector>
#include "block_allocator.cuh"

// Simple kernel to write a pattern to a block's memory
__global__ void write_pattern_kernel(float* block_ptr, int num_elements, float value) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < num_elements) {
        block_ptr[idx] = value;
    }
}

// Kernel to verify a pattern
__global__ void verify_pattern_kernel(float* block_ptr, int num_elements,
                                      float expected, int* errors) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < num_elements) {
        if (block_ptr[idx] != expected) {
            atomicAdd(errors, 1);
        }
    }
}

void print_separator() {
    printf("============================================================\n");
}

void test_block_allocator() {
    printf("\n[TEST] Block Allocator - Basic Allocation and Free\n");
    print_separator();

    size_t pool_size = 64 * 1024 * 1024;
    size_t block_size = 16 * 128 * sizeof(float);
    BlockAllocator alloc(pool_size, block_size);

    printf("  Pool size: %zu MB\n", pool_size / (1024 * 1024));
    printf("  Block size: %zu bytes\n", block_size);
    printf("  Total blocks: %zu\n", alloc.num_total_blocks());
    printf("  Free blocks: %zu\n", alloc.num_free_blocks());

    std::vector<int> allocated;
    for (int i = 0; i < 5; ++i) {
        int id = alloc.allocate();
        if (id >= 0) allocated.push_back(id);
    }
    printf("  After allocating 5: free=%zu, allocated=%zu\n",
           alloc.num_free_blocks(), alloc.num_allocated_blocks());

    int elements_per_block = static_cast<int>(block_size / sizeof(float));
    int threads = 256;
    int blocks = (elements_per_block + threads - 1) / threads;

    int* d_errors;
    cudaMalloc(&d_errors, sizeof(int));

    for (size_t i = 0; i < allocated.size(); ++i) {
        float* ptr = alloc.get_block_ptr(allocated[i]);
        float pattern = static_cast<float>(i + 1) * 1.5f;
        write_pattern_kernel<<<blocks, threads>>>(ptr, elements_per_block, pattern);
    }
    cudaDeviceSynchronize();

    int total_errors = 0;
    for (size_t i = 0; i < allocated.size(); ++i) {
        float* ptr = alloc.get_block_ptr(allocated[i]);
        float pattern = static_cast<float>(i + 1) * 1.5f;
        int zero = 0;
        cudaMemcpy(d_errors, &zero, sizeof(int), cudaMemcpyHostToDevice);
        verify_pattern_kernel<<<blocks, threads>>>(ptr, elements_per_block, pattern, d_errors);
        int err_count = 0;
        cudaMemcpy(&err_count, d_errors, sizeof(int), cudaMemcpyDeviceToHost);
        total_errors += err_count;
    }
    printf("  Pattern verification errors: %d\n", total_errors);

    for (int id : allocated) {
        alloc.free(id);
    }
    printf("  After freeing all: free=%zu\n", alloc.num_free_blocks());

    cudaFree(d_errors);
    printf("  [PASS] Block allocator test complete\n");
}

int main() {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    printf("=== Paged KV-Cache Engine ===\n");
    printf("GPU: %s | VRAM: %zu MB | Compute: %d.%d\n\n",
           prop.name, prop.totalGlobalMem / (1024 * 1024),
           prop.major, prop.minor);

    test_block_allocator();

    printf("\n=== All tests passed! ===\n");
    return 0;
}

Your physical block allocator is working. It can claim GPU memory, divide it into blocks, hand them out, and reclaim them with zero external fragmentation. Next up, you'll build the page table that maps logical token sequences to these physical blocks, letting non-contiguous memory appear contiguous to the attention kernel.

Implement the Page Table

Your block allocator hands out physical GPU memory blocks, but there's a problem. When a sequence needs three blocks, the allocator might return blocks 7, 2, and 15. They're scattered across the pool.

A page table solves this by creating a logical-to-physical mapping. The sequence thinks it owns blocks 0, 1, 2 in order, but the page table transparently redirects those to the actual physical locations. This is the same indirection that operating systems use for virtual memory.

In this step, get ready to:

• Define a PageTable class that maps logical block indices to physical block IDs.
• Implement the lookup logic for translating logical indices into device pointers.
• Test the mapping by writing and reading patterns through non-contiguous physical blocks.

Define the PageTable class

In vLLM's architecture, a block table is a per-sequence array where each index represents a logical block and its value is the physical block ID on the GPU. Your PageTable class implements this exact pattern using a std::vector<int>.

• Create a new file called page_table.h inside your include/ folder.
• Add the following code to define the PageTable class:

#ifndef PAGE_TABLE_H
#define PAGE_TABLE_H

#include <vector>
#include <cstddef>

class PageTable {
public:
    PageTable() : num_filled_in_last_block_(0) {}

    void append_block(int physical_block_id) {
        entries_.push_back(physical_block_id);
        num_filled_in_last_block_ = 0;
    }

    int get_physical_id(size_t logical_index) const {
        if (logical_index >= entries_.size()) return -1;
        return entries_[logical_index];
    }

    size_t num_blocks() const { return entries_.size(); }

    void set_filled_in_last(int count) { num_filled_in_last_block_ = count; }
    int filled_in_last() const { return num_filled_in_last_block_; }

    const std::vector<int>& entries() const { return entries_; }

    void clear() {
        entries_.clear();
        num_filled_in_last_block_ = 0;
    }

private:
    std::vector<int> entries_;
    int num_filled_in_last_block_;
};

#endif // PAGE_TABLE_H

What does this code do?

The PageTable stores a mapping from logical block indices to physical block IDs. When a sequence grows, you call append_block() with whatever physical block the allocator gave you. Later, get_physical_id() translates a logical index back to the physical block ID.

The num_filled_in_last_block_ field tracks how many token slots are actually used in the most recent block. This matters because the last block in a sequence is usually only partially full.

The entries() method returns a reference to the underlying contiguous vector. Because std::vector<int> stores data contiguously, you can copy its .data() pointer directly to GPU memory for device-side lookups during attention kernels.

• Save include/page_table.h.

How does this map to vLLM's architecture?

In vLLM, the block_tables tensor has shape [num_seqs, max_num_blocks_per_seq] and maps logical block indices to physical block IDs. Your PageTable class is the single-sequence equivalent. The key insight from the PagedAttention paper is that "one can think of blocks as pages, tokens as bytes, and requests as processes."

Implement the block-table lookup logic

The page table translates logical indices to physical block IDs, but turning a block ID into an actual GPU pointer requires one more step. The pattern is: look up the physical block ID from the page table, then compute the device address from the pool base pointer plus the block's offset.

How does the address lookup work?

To find the GPU memory address for any token in a sequence, the lookup follows three steps:

• Compute which logical block the token falls in: logical_block = token_index / block_size.
• Translate logical to physical: physical_id = page_table.get_physical_id(logical_block).
• Compute the device pointer: pool_base + physical_id * block_size_bytes + token_offset.

Your BlockAllocator already handles step 3 via get_block_ptr(). In the next step, you will combine both classes into a KVCacheManager that encapsulates the full lookup in a single get_token_ptr() method.

Test the page-table mapping with non-contiguous blocks

The real test of a page table is whether it correctly resolves logical indices to scattered physical blocks. You will allocate blocks with gaps between them, build a page table, write distinct patterns to each physical block, then read them back using only the logical indices.

• In src/main.cu, add the include for your new header at the top of the file, below the existing #include "block_allocator.cuh" line:

#include "page_table.h"

• Add the following test_page_table_mapping() function below the existing test_block_allocator() function:

void test_page_table_mapping() {
    printf("\n[TEST] Page Table - Logical to Physical Mapping\n");
    print_separator();

    size_t pool_size = 64 * 1024 * 1024;
    size_t tokens_per_block = 16;
    size_t token_kv_bytes = 128 * sizeof(float); // 128-dim KV vector
    size_t block_bytes = tokens_per_block * token_kv_bytes;

    BlockAllocator alloc(pool_size, block_bytes);

    // Allocate 3 non-contiguous blocks
    int b0 = alloc.allocate(); // Gets block from end of free list
    // Skip some blocks to get non-contiguous allocation
    std::vector<int> skipped;
    for (int i = 0; i < 5; ++i) skipped.push_back(alloc.allocate());
    int b1 = alloc.allocate();
    for (int i = 0; i < 3; ++i) skipped.push_back(alloc.allocate());
    int b2 = alloc.allocate();

    // Free skipped blocks
    for (int id : skipped) alloc.free(id);

    printf("  Logical block 0 -> Physical block %d\n", b0);
    printf("  Logical block 1 -> Physical block %d\n", b1);
    printf("  Logical block 2 -> Physical block %d\n", b2);

    // Build page table
    PageTable pt;
    pt.append_block(b0);
    pt.append_block(b1);
    pt.append_block(b2);

    // Verify mapping
    bool pass = (pt.get_physical_id(0) == b0 &&
                 pt.get_physical_id(1) == b1 &&
                 pt.get_physical_id(2) == b2);
    printf("  Page table entries: %zu\n", pt.num_blocks());
    printf("  Mapping verification: %s\n", pass ? "PASS" : "FAIL");

    // Write distinct values to each block via page table indirection
    int elements = static_cast<int>(block_bytes / sizeof(float));
    int threads = 256;
    int grid = (elements + threads - 1) / threads;

    write_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(b0), elements, 100.0f);
    write_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(b1), elements, 200.0f);
    write_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(b2), elements, 300.0f);
    cudaDeviceSynchronize();

    // Verify via page table lookup
    int* d_errors;
    cudaMalloc(&d_errors, sizeof(int));
    int zero = 0;
    int total_errors = 0;

    for (int logical = 0; logical < 3; ++logical) {
        int phys = pt.get_physical_id(logical);
        float expected = (logical + 1) * 100.0f;
        cudaMemcpy(d_errors, &zero, sizeof(int), cudaMemcpyHostToDevice);
        verify_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(phys), elements, expected, d_errors);
        int err_count = 0;
        cudaMemcpy(&err_count, d_errors, sizeof(int), cudaMemcpyDeviceToHost);
        total_errors += err_count;
    }
    printf("  Indirection read-back errors: %d\n", total_errors);

    alloc.free(b0);
    alloc.free(b1);
    alloc.free(b2);
    cudaFree(d_errors);
    printf("  [PASS] Page table test complete\n");
}

What does this test do?

The test deliberately creates gaps in allocation. It allocates block b0, skips 5 blocks, allocates b1, skips 3 more, then allocates b2. The skipped blocks are freed immediately. This simulates a real scenario where free blocks are scattered across the pool.

After building the page table mapping (logical 0 to b0, logical 1 to b1, logical 2 to b2), the test writes unique float patterns to each physical block. It then reads them back using only the logical index and the page table. If the indirection is correct, every value matches.

• In the main() function, find this line:

    test_block_allocator();

• Add the call to your new test function directly below it:

    test_block_allocator();
    test_page_table_mapping();

• Save src/main.cu.
• Rebuild and run your project by running these commands:

cmake --build build --config Release
.\build\Release\paged_kv_cache.exe

• You should see the page table test print three logical-to-physical mappings with non-sequential physical block IDs, a PASS for mapping verification, 0 indirection read-back errors, and [PASS] Page table test complete.

Seeing compile errors or test failures?

• If you see an error about page_table.h not found, make sure the file is in the include/ directory (not src/).
• If the mapping verification prints FAIL, check that you are calling append_block() in the correct order (b0, b1, b2).
• If you see read-back errors, confirm that cudaDeviceSynchronize() is called between the write and verify kernels.

Why is my page table test failing?

✔️ Awesome, I've got everything!

Great. Double check that both files are saved before moving on.

ⓧ I'd like to double check the full code

Here is the complete include/page_table.h:

#ifndef PAGE_TABLE_H
#define PAGE_TABLE_H

#include <vector>
#include <cstddef>

class PageTable {
public:
    PageTable() : num_filled_in_last_block_(0) {}

    void append_block(int physical_block_id) {
        entries_.push_back(physical_block_id);
        num_filled_in_last_block_ = 0;
    }

    int get_physical_id(size_t logical_index) const {
        if (logical_index >= entries_.size()) return -1;
        return entries_[logical_index];
    }

    size_t num_blocks() const { return entries_.size(); }

    void set_filled_in_last(int count) { num_filled_in_last_block_ = count; }
    int filled_in_last() const { return num_filled_in_last_block_; }

    const std::vector<int>& entries() const { return entries_; }

    void clear() {
        entries_.clear();
        num_filled_in_last_block_ = 0;
    }

private:
    std::vector<int> entries_;
    int num_filled_in_last_block_;
};

#endif // PAGE_TABLE_H

Here is the complete src/main.cu:

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>
#include <ctime>
#include <vector>
#include "block_allocator.cuh"
#include "page_table.h"

__global__ void write_pattern_kernel(float* block_ptr, int num_elements, float value) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < num_elements) {
        block_ptr[idx] = value;
    }
}

__global__ void verify_pattern_kernel(float* block_ptr, int num_elements,
                                      float expected, int* errors) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < num_elements) {
        if (block_ptr[idx] != expected) {
            atomicAdd(errors, 1);
        }
    }
}

void print_separator() {
    printf("============================================================\n");
}

void test_block_allocator() {
    // ... (same as before)
}

void test_page_table_mapping() {
    printf("\n[TEST] Page Table - Logical to Physical Mapping\n");
    print_separator();

    size_t pool_size = 64 * 1024 * 1024;
    size_t tokens_per_block = 16;
    size_t token_kv_bytes = 128 * sizeof(float);
    size_t block_bytes = tokens_per_block * token_kv_bytes;

    BlockAllocator alloc(pool_size, block_bytes);

    int b0 = alloc.allocate();
    std::vector<int> skipped;
    for (int i = 0; i < 5; ++i) skipped.push_back(alloc.allocate());
    int b1 = alloc.allocate();
    for (int i = 0; i < 3; ++i) skipped.push_back(alloc.allocate());
    int b2 = alloc.allocate();

    for (int id : skipped) alloc.free(id);

    printf("  Logical block 0 -> Physical block %d\n", b0);
    printf("  Logical block 1 -> Physical block %d\n", b1);
    printf("  Logical block 2 -> Physical block %d\n", b2);

    PageTable pt;
    pt.append_block(b0);
    pt.append_block(b1);
    pt.append_block(b2);

    bool pass = (pt.get_physical_id(0) == b0 &&
                 pt.get_physical_id(1) == b1 &&
                 pt.get_physical_id(2) == b2);
    printf("  Page table entries: %zu\n", pt.num_blocks());
    printf("  Mapping verification: %s\n", pass ? "PASS" : "FAIL");

    int elements = static_cast<int>(block_bytes / sizeof(float));
    int threads = 256;
    int grid = (elements + threads - 1) / threads;

    write_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(b0), elements, 100.0f);
    write_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(b1), elements, 200.0f);
    write_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(b2), elements, 300.0f);
    cudaDeviceSynchronize();

    int* d_errors;
    cudaMalloc(&d_errors, sizeof(int));
    int zero = 0;
    int total_errors = 0;

    for (int logical = 0; logical < 3; ++logical) {
        int phys = pt.get_physical_id(logical);
        float expected = (logical + 1) * 100.0f;
        cudaMemcpy(d_errors, &zero, sizeof(int), cudaMemcpyHostToDevice);
        verify_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(phys), elements, expected, d_errors);
        int err_count = 0;
        cudaMemcpy(&err_count, d_errors, sizeof(int), cudaMemcpyDeviceToHost);
        total_errors += err_count;
    }
    printf("  Indirection read-back errors: %d\n", total_errors);

    alloc.free(b0);
    alloc.free(b1);
    alloc.free(b2);
    cudaFree(d_errors);
    printf("  [PASS] Page table test complete\n");
}

int main() {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    printf("=== Paged KV-Cache Engine ===\n");
    printf("GPU: %s | VRAM: %zu MB | Compute: %d.%d\n\n",
           prop.name, prop.totalGlobalMem / (1024 * 1024),
           prop.major, prop.minor);

    test_block_allocator();
    test_page_table_mapping();

    printf("\n=== All tests passed! ===\n");
    return 0;
}

Your page table now translates logical token positions to scattered physical blocks without the sequence ever knowing its memory is non-contiguous. Next up, you will combine the BlockAllocator and PageTable into a single KVCacheManager that automatically grows sequences as tokens are generated.

Build the KV-Cache Manager

Your block allocator handles raw GPU memory, and your page table maps logical blocks to physical ones. Now you need a top-level coordinator that ties these two together into a usable API.

In a production LLM serving engine, the KV-Cache manager is the single entry point for all sequence memory operations. It decides when to allocate new blocks, tracks which sequences are active, and reports utilization metrics. In this step, you will build that manager.

In this step, get ready to:

• Define the KVCacheManager class with a memory pool and sequence registry.
• Implement append-token logic that automatically allocates new blocks when needed.
• Implement memory utilization and fragmentation reporting.

Define the KVCacheManager class

The manager owns a BlockAllocator and a map of sequence IDs to their individual page tables. It exposes a clean interface so the simulation harness never touches raw block IDs directly.

• Create a new file called kv_cache_manager.h inside the include/ folder.
• Add the following code to define the CacheStats struct, the class constructor, and the sequence registration method:

#ifndef KV_CACHE_MANAGER_H
#define KV_CACHE_MANAGER_H

#include "block_allocator.cuh"
#include "page_table.h"
#include <unordered_map>
#include <cstdio>

struct CacheStats {
    size_t total_blocks;
    size_t free_blocks;
    size_t allocated_blocks;
    size_t active_sequences;
    float utilization_percent;
    float internal_frag_percent;
};

class KVCacheManager {
public:
    KVCacheManager(size_t total_pool_bytes, size_t block_size_tokens,
                   size_t token_kv_size_bytes)
        : block_size_tokens_(block_size_tokens),
          token_kv_size_bytes_(token_kv_size_bytes),
          allocator_(total_pool_bytes, block_size_tokens * token_kv_size_bytes) {
    }

    bool register_sequence(int seq_id) {
        if (sequences_.count(seq_id)) return false;
        sequences_[seq_id] = PageTable();
        return true;
    }

What does this code do?

The CacheStats struct is a plain data container for reporting memory metrics back to the caller.

The constructor takes three configuration values: total pool size in bytes, tokens per block, and bytes of KV data per token. It computes the block size in bytes (block_size_tokens * token_kv_size_bytes) and passes that down to the BlockAllocator it owns.

The register_sequence() method creates an empty PageTable for a new sequence and stores it in the sequences_ map. It returns false if the ID is already taken.

Why does the manager own the allocator?

In vLLM's architecture, the block space manager is the single authority over physical memory. No other component calls cudaMalloc directly. By making the KVCacheManager own the BlockAllocator, you enforce the same single-owner guarantee.

• Save kv_cache_manager.h. The file is not complete yet, so it will not compile on its own.

Implement the append token logic

Every time the LLM generates a token, the serving engine must store that token's key and value vectors. The manager checks whether the current block has room. If not, it grabs a fresh physical block from the allocator.

• In kv_cache_manager.h, directly below the closing brace of register_sequence, add the following methods:

    // Returns true if successful, false if out of memory
    bool append_token(int seq_id) {
        auto it = sequences_.find(seq_id);
        if (it == sequences_.end()) return false;

        PageTable& pt = it->second;

        // Check if we need a new block
        if (pt.num_blocks() == 0 || pt.filled_in_last() >= static_cast<int>(block_size_tokens_)) {
            int new_block = allocator_.allocate();
            if (new_block < 0) return false; // OOM
            pt.append_block(new_block);
        }

        pt.set_filled_in_last(pt.filled_in_last() + 1);
        return true;
    }

    // Append multiple tokens at once (e.g., prompt prefill)
    bool append_tokens(int seq_id, int count) {
        for (int i = 0; i < count; ++i) {
            if (!append_token(seq_id)) return false;
        }
        return true;
    }

    void free_sequence(int seq_id) {
        auto it = sequences_.find(seq_id);
        if (it == sequences_.end()) return;

        const PageTable& pt = it->second;
        for (int phys_id : pt.entries()) {
            allocator_.free(phys_id);
        }
        sequences_.erase(it);
    }

    const PageTable* get_page_table(int seq_id) const {
        auto it = sequences_.find(seq_id);
        if (it == sequences_.end()) return nullptr;
        return &it->second;
    }

    float* get_token_ptr(int seq_id, int token_index) const {
        auto it = sequences_.find(seq_id);
        if (it == sequences_.end()) return nullptr;

        const PageTable& pt = it->second;
        size_t logical_block = token_index / block_size_tokens_;
        size_t offset_in_block = token_index % block_size_tokens_;

        int phys_id = pt.get_physical_id(logical_block);
        if (phys_id < 0) return nullptr;

        float* block_ptr = allocator_.get_block_ptr(phys_id);
        if (!block_ptr) return nullptr;

        size_t float_offset = offset_in_block * (token_kv_size_bytes_ / sizeof(float));
        return block_ptr + float_offset;
    }

What does this code do?

• append_token() is the core allocation path. It checks filled_in_last() against block_size_tokens_ to decide whether the current block is full. If full, it calls allocator_.allocate() for a fresh block. This is the paging mechanism in action.
• append_tokens() is a convenience wrapper for prompt prefill, where many tokens arrive at once.
• free_sequence() iterates all physical blocks in the sequence's page table and returns each to the allocator's free list. This is how blocks get recycled for new requests.
• get_token_ptr() performs the full logical-to-physical address translation: divide the token index by block size to find the logical block, look up the physical block ID in the page table, then compute the byte offset within that block.

• Save kv_cache_manager.h. The file still needs the get_stats() method and closing braces before it will compile.

Implement get_stats() and verify the build

The stats method measures how efficiently your engine uses memory. It calculates both utilization (what fraction of blocks are in use) and internal fragmentation (how many token slots are wasted in partially-filled last blocks).

• Below get_token_ptr, add the get_stats() method, the accessor methods, and the private members to close out the class:

    CacheStats get_stats() const {
        CacheStats stats;
        stats.total_blocks = allocator_.num_total_blocks();
        stats.free_blocks = allocator_.num_free_blocks();
        stats.allocated_blocks = allocator_.num_allocated_blocks();
        stats.active_sequences = sequences_.size();

        stats.utilization_percent = (stats.total_blocks > 0)
            ? 100.0f * stats.allocated_blocks / stats.total_blocks
            : 0.0f;

        // Internal fragmentation: wasted slots in partially-filled last blocks
        size_t total_slots_allocated = stats.allocated_blocks * block_size_tokens_;
        size_t total_slots_used = 0;
        for (const auto& pair : sequences_) {
            const PageTable& pt = pair.second;
            if (pt.num_blocks() > 0) {
                total_slots_used += (pt.num_blocks() - 1) * block_size_tokens_ + pt.filled_in_last();
            }
        }
        stats.internal_frag_percent = (total_slots_allocated > 0)
            ? 100.0f * (total_slots_allocated - total_slots_used) / total_slots_allocated
            : 0.0f;

        return stats;
    }

    size_t block_size_tokens() const { return block_size_tokens_; }
    const BlockAllocator& allocator() const { return allocator_; }

private:
    size_t block_size_tokens_;
    size_t token_kv_size_bytes_;
    BlockAllocator allocator_;
    std::unordered_map<int, PageTable> sequences_;
};

#endif // KV_CACHE_MANAGER_H

What does this code do?

• The utilization percentage is simply allocated blocks divided by total blocks. Zero external fragmentation is guaranteed by design because all blocks are the same size.
• Internal fragmentation measures the wasted slots inside allocated blocks. Only the last block of each sequence can be partially filled, so the formula sums total allocated slots minus actually-used slots across all sequences.
• The private section holds block_size_tokens_ (tokens per block), token_kv_size_bytes_ (bytes per token's KV data), the BlockAllocator instance, and the sequence map.

• Save kv_cache_manager.h.

Now update src/main.cu so it includes the new manager header instead of the individual component headers.

• In src/main.cu, find these two include lines near the top of the file:

#include "block_allocator.cuh"
#include "page_table.h"

• Replace both lines with a single include for the manager:

#include "kv_cache_manager.h"

Why replace both includes?

The kv_cache_manager.h header already includes both block_allocator.cuh and page_table.h internally. Keeping the old includes would still compile (thanks to header guards), but removing them makes the dependency chain clearer.

• Save src/main.cu.
• Rebuild the project by running:

cmake --build build --config Release

• You should see a clean compilation with no errors. Run the executable to confirm all existing tests still pass.

Seeing a compile error?

• If you see an error about an incomplete type or missing method, check that you closed the class with }; and the header guard with #endif at the very end of kv_cache_manager.h.
• If you see a linker error about multiple definitions, confirm that all method implementations are inside the class body (inline by default for header-only classes).
• If you see an error about unordered_map not being found, make sure #include <unordered_map> is at the top of the file.

Help me debug my KVCacheManager compile error.

✔️ Awesome, I've got everything!

Great. Double check you have saved both include/kv_cache_manager.h and src/main.cu.

ⓧ I'd like to double check the full code

Here is the complete include/kv_cache_manager.h:

#ifndef KV_CACHE_MANAGER_H
#define KV_CACHE_MANAGER_H

#include "block_allocator.cuh"
#include "page_table.h"
#include <unordered_map>
#include <cstdio>

struct CacheStats {
    size_t total_blocks;
    size_t free_blocks;
    size_t allocated_blocks;
    size_t active_sequences;
    float utilization_percent;
    float internal_frag_percent;
};

class KVCacheManager {
public:
    KVCacheManager(size_t total_pool_bytes, size_t block_size_tokens,
                   size_t token_kv_size_bytes)
        : block_size_tokens_(block_size_tokens),
          token_kv_size_bytes_(token_kv_size_bytes),
          allocator_(total_pool_bytes, block_size_tokens * token_kv_size_bytes) {
    }

    bool register_sequence(int seq_id) {
        if (sequences_.count(seq_id)) return false;
        sequences_[seq_id] = PageTable();
        return true;
    }

    // Returns true if successful, false if out of memory
    bool append_token(int seq_id) {
        auto it = sequences_.find(seq_id);
        if (it == sequences_.end()) return false;

        PageTable& pt = it->second;

        // Check if we need a new block
        if (pt.num_blocks() == 0 || pt.filled_in_last() >= static_cast<int>(block_size_tokens_)) {
            int new_block = allocator_.allocate();
            if (new_block < 0) return false; // OOM
            pt.append_block(new_block);
        }

        pt.set_filled_in_last(pt.filled_in_last() + 1);
        return true;
    }

    // Append multiple tokens at once (e.g., prompt prefill)
    bool append_tokens(int seq_id, int count) {
        for (int i = 0; i < count; ++i) {
            if (!append_token(seq_id)) return false;
        }
        return true;
    }

    void free_sequence(int seq_id) {
        auto it = sequences_.find(seq_id);
        if (it == sequences_.end()) return;

        const PageTable& pt = it->second;
        for (int phys_id : pt.entries()) {
            allocator_.free(phys_id);
        }
        sequences_.erase(it);
    }

    const PageTable* get_page_table(int seq_id) const {
        auto it = sequences_.find(seq_id);
        if (it == sequences_.end()) return nullptr;
        return &it->second;
    }

    float* get_token_ptr(int seq_id, int token_index) const {
        auto it = sequences_.find(seq_id);
        if (it == sequences_.end()) return nullptr;

        const PageTable& pt = it->second;
        size_t logical_block = token_index / block_size_tokens_;
        size_t offset_in_block = token_index % block_size_tokens_;

        int phys_id = pt.get_physical_id(logical_block);
        if (phys_id < 0) return nullptr;

        float* block_ptr = allocator_.get_block_ptr(phys_id);
        if (!block_ptr) return nullptr;

        size_t float_offset = offset_in_block * (token_kv_size_bytes_ / sizeof(float));
        return block_ptr + float_offset;
    }

    CacheStats get_stats() const {
        CacheStats stats;
        stats.total_blocks = allocator_.num_total_blocks();
        stats.free_blocks = allocator_.num_free_blocks();
        stats.allocated_blocks = allocator_.num_allocated_blocks();
        stats.active_sequences = sequences_.size();

        stats.utilization_percent = (stats.total_blocks > 0)
            ? 100.0f * stats.allocated_blocks / stats.total_blocks
            : 0.0f;

        // Internal fragmentation: wasted slots in partially-filled last blocks
        size_t total_slots_allocated = stats.allocated_blocks * block_size_tokens_;
        size_t total_slots_used = 0;
        for (const auto& pair : sequences_) {
            const PageTable& pt = pair.second;
            if (pt.num_blocks() > 0) {
                total_slots_used += (pt.num_blocks() - 1) * block_size_tokens_ + pt.filled_in_last();
            }
        }
        stats.internal_frag_percent = (total_slots_allocated > 0)
            ? 100.0f * (total_slots_allocated - total_slots_used) / total_slots_allocated
            : 0.0f;

        return stats;
    }

    size_t block_size_tokens() const { return block_size_tokens_; }
    const BlockAllocator& allocator() const { return allocator_; }

private:
    size_t block_size_tokens_;
    size_t token_kv_size_bytes_;
    BlockAllocator allocator_;
    std::unordered_map<int, PageTable> sequences_;
};

#endif // KV_CACHE_MANAGER_H

And the updated includes at the top of src/main.cu:

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>
#include <ctime>
#include <vector>
#include "kv_cache_manager.h"

Your KV-Cache manager is ready. Next up, you will build a multi-sequence simulation that puts this manager under real load and proves it eliminates fragmentation.

Run the Multi-Sequence Simulation and Metrics

Your memory manager now handles the full lifecycle of a single sequence. In production, though, an LLM serving engine juggles dozens of requests simultaneously, each growing at different rates and finishing at unpredictable times.

In this step, you will build a simulation harness that exercises your paged memory system under realistic multi-sequence load. The simulation proves that blocks freed by completed sequences are immediately reusable by new arrivals, eliminating the external fragmentation that plagues contiguous allocation schemes.

In this step, get ready to:

• Build a simulation harness that runs 8 concurrent sequences with random lengths.
• Simulate sequence completion and immediate block reuse for 100 total completions.
• Print a metrics dashboard comparing paged allocation to naive contiguous pre-allocation.

Build the simulation harness

A real LLM serving engine handles many requests at once. Each request starts with a prompt (prefill phase), then generates tokens one at a time (decode phase) until it hits its target length or a stop token.

Your simulation mimics this pattern. Each sequence gets a random prompt length and a random target length, and the manager allocates blocks on demand as tokens arrive.

• In src/main.cu, find the line that says void test_page_table_mapping() { and scroll down past the closing brace of that function.
• Add the beginning of the run_multi_sequence_simulation() function by copying the following code below test_page_table_mapping() and above int main():

void run_multi_sequence_simulation() {
    printf("\n[SIM] Multi-Sequence KV-Cache Simulation\n");
    print_separator();

    // Configuration
    size_t pool_size = 256 * 1024 * 1024; // 256 MB pool
    size_t tokens_per_block = 16;
    size_t token_kv_bytes = 128 * sizeof(float); // Simulated: 128-dim KV per token
    int num_concurrent = 8;
    int total_completions = 100;

    KVCacheManager manager(pool_size, tokens_per_block, token_kv_bytes);

    printf("  Pool: %zu MB | Block: %zu tokens | Token KV: %zu bytes\n",
           pool_size / (1024 * 1024), tokens_per_block, token_kv_bytes);
    printf("  Simulating %d concurrent sequences, %d total completions\n\n",
           num_concurrent, total_completions);

    srand(42);

    // Per-sequence tracking struct
    struct SeqInfo {
        int seq_id;
        int current_len;
        int target_len;
        bool active;
    };

    std::vector<SeqInfo> slots(num_concurrent);
    int next_seq_id = 0;
    int completions = 0;
    size_t total_allocs = 0;
    size_t total_frees = 0;
    float peak_utilization = 0.0f;
    size_t naive_memory_needed = 0; // For comparison

    // Initialize first batch of sequences
    for (int i = 0; i < num_concurrent; ++i) {
        slots[i].seq_id = next_seq_id++;
        slots[i].current_len = 0;
        slots[i].target_len = 32 + rand() % 225; // 32-256 tokens
        slots[i].active = true;
        manager.register_sequence(slots[i].seq_id);

        // Prefill with random prompt
        int prompt_len = 4 + rand() % 29; // 4-32 tokens
        manager.append_tokens(slots[i].seq_id, prompt_len);
        slots[i].current_len = prompt_len;
        naive_memory_needed += slots[i].target_len * token_kv_bytes;
    }

What does this code do?

This sets up the simulation environment. The KVCacheManager gets a 256 MB pool divided into blocks of 16 tokens each, with 128 floats (512 bytes) of KV data per token.

The SeqInfo struct tracks each active sequence's ID, current length, target length, and whether it is still generating. The initialization loop registers 8 sequences with random prompt and target lengths, simulating the prefill phase where a batch of requests arrives simultaneously.

Simulate sequence completion and reuse

The core loop generates one token per sequence per iteration. When a sequence reaches its target length, its blocks return to the free pool and a new sequence immediately takes its slot.

• Directly below the initialization loop you just added (still inside run_multi_sequence_simulation()), add the simulation loop:

    // Simulation loop
    while (completions < total_completions) {
        for (int i = 0; i < num_concurrent; ++i) {
            if (!slots[i].active) continue;

            // Generate one token
            if (manager.append_token(slots[i].seq_id)) {
                slots[i].current_len++;
                total_allocs++;
            }

            // Check if sequence is complete
            if (slots[i].current_len >= slots[i].target_len) {
                manager.free_sequence(slots[i].seq_id);
                total_frees++;
                completions++;
                slots[i].active = false;

                // Start a new sequence in this slot
                if (completions < total_completions) {
                    slots[i].seq_id = next_seq_id++;
                    slots[i].current_len = 0;
                    slots[i].target_len = 32 + rand() % 225;
                    slots[i].active = true;
                    manager.register_sequence(slots[i].seq_id);

                    int prompt_len = 4 + rand() % 29;
                    manager.append_tokens(slots[i].seq_id, prompt_len);
                    slots[i].current_len = prompt_len;
                    naive_memory_needed += slots[i].target_len * token_kv_bytes;
                }
            }
        }

        // Track peak utilization across all iterations
        CacheStats stats = manager.get_stats();
        if (stats.utilization_percent > peak_utilization) {
            peak_utilization = stats.utilization_percent;
        }
    }

What does this code do?

Each iteration of the outer while loop simulates one "time step" across all active sequences. append_token() grows each sequence by one token, automatically allocating new physical blocks when the current block fills up.

When a sequence hits its target length, free_sequence() returns all its physical blocks to the free-list. A new sequence immediately takes the vacated slot and begins allocating from the same free pool. This proves block reuse without external fragmentation.

Why is external fragmentation zero?

Because every block is the same size, any freed block can satisfy any new allocation request. There is never a situation where free memory exists but cannot be used because it is split into pieces too small for a request. This is the core insight of PagedAttention.

Print the metrics dashboard

The final section of the function prints a comprehensive report comparing your paged system against naive pre-allocation. This quantifies exactly how much VRAM the paging approach saves.

• Add the metrics printing code directly below the simulation loop (still inside the function), then close the function with a brace:

    // Final stats
    CacheStats final_stats = manager.get_stats();
    printf("  --- Simulation Results ---\n");
    printf("  Sequences completed: %d\n", completions);
    printf("  Total token allocations: %zu\n", total_allocs);
    printf("  Total sequence frees: %zu\n", total_frees);
    printf("  Peak memory utilization: %.1f%%\n", peak_utilization);
    printf("  Final utilization: %.1f%%\n", final_stats.utilization_percent);
    printf("  Internal fragmentation: %.1f%%\n", final_stats.internal_frag_percent);
    printf("  External fragmentation: 0%% (by design - uniform blocks)\n");
    printf("\n  --- Paged vs Naive Comparison ---\n");
    printf("  Paged pool size: %zu MB\n", pool_size / (1024 * 1024));
    printf("  Naive pre-alloc would need: %zu MB\n", naive_memory_needed / (1024 * 1024));
    printf("  Memory saved: %.1fx reduction\n",
           static_cast<float>(naive_memory_needed) / pool_size);
    printf("\n  [PASS] Simulation complete\n");
}

What does this code do?

The dashboard reports two categories of metrics. The first section shows runtime stats: how many sequences completed, how many token allocations occurred, and the peak and final memory utilization percentages.

The second section compares paged vs naive allocation. The naive approach would pre-allocate max_seq_len * num_concurrent_seqs memory upfront. Your paged system only allocates blocks on demand, so the ratio shows how many times more memory the naive approach wastes.

• Now scroll down to the main() function and find this line:

    test_page_table_mapping();

    printf("\n=== All tests passed! ===\n");

• Add a call to run_multi_sequence_simulation() between test_page_table_mapping(); and the final printf, so it looks like this:

    test_page_table_mapping();
    run_multi_sequence_simulation();

    printf("\n=== All tests passed! ===\n");

• Save src/main.cu.
• Rebuild and run your project by running these commands:

cmake --build build --config Release
.\build\Release\paged_kv_cache.exe

• You should see the two earlier tests pass, followed by the simulation output showing [SIM] Multi-Sequence KV-Cache Simulation, a metrics dashboard, and [PASS] Simulation complete at the end.

What should I see?

The simulation prints a configuration summary, followed by results showing 100 completed sequences, a peak utilization percentage (typically 5-15% because 256 MB is generous for 8 short sequences), internal fragmentation from partially-filled last blocks, and 0% external fragmentation.

The "Paged vs Naive" comparison shows the memory saved. The naive approach would pre-allocate max_seq_len for all concurrent sequences upfront. Your paged engine only allocated what was actually needed, resulting in a significant memory reduction ratio.

Build failing or simulation not printing?

• If you see a compile error about SeqInfo or slots, make sure the struct definition and vector are inside the function body, not outside it.
• If the program runs but you do not see simulation output, confirm you added the run_multi_sequence_simulation(); call inside main().
• If you see a mismatched brace error, verify that your function ends with a closing } after the final printf statement.

Help me debug my multi-sequence simulation build error.

✔️ Awesome, I've got everything!

Great. Double check you have saved src/main.cu and confirmed the simulation completes with [PASS] Simulation complete in the output.

ⓧ I'd like to double check the full code

Here is the complete src/main.cu file as it should look now:

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>
#include <ctime>
#include <vector>
#include "kv_cache_manager.h"

__global__ void write_pattern_kernel(float* block_ptr, int num_elements, float value) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < num_elements) {
        block_ptr[idx] = value;
    }
}

__global__ void verify_pattern_kernel(float* block_ptr, int num_elements,
                                      float expected, int* errors) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < num_elements) {
        if (block_ptr[idx] != expected) {
            atomicAdd(errors, 1);
        }
    }
}

void print_separator() {
    printf("============================================================\n");
}

void test_block_allocator() {
    printf("\n[TEST] Block Allocator - Basic Allocation and Free\n");
    print_separator();

    size_t pool_size = 64 * 1024 * 1024;
    size_t block_size = 16 * 128 * sizeof(float);
    BlockAllocator alloc(pool_size, block_size);

    printf("  Pool size: %zu MB\n", pool_size / (1024 * 1024));
    printf("  Block size: %zu bytes\n", block_size);
    printf("  Total blocks: %zu\n", alloc.num_total_blocks());
    printf("  Free blocks: %zu\n", alloc.num_free_blocks());

    std::vector<int> allocated;
    for (int i = 0; i < 5; ++i) {
        int id = alloc.allocate();
        if (id >= 0) allocated.push_back(id);
    }
    printf("  After allocating 5: free=%zu, allocated=%zu\n",
           alloc.num_free_blocks(), alloc.num_allocated_blocks());

    int elements_per_block = static_cast<int>(block_size / sizeof(float));
    int threads = 256;
    int blocks = (elements_per_block + threads - 1) / threads;

    int* d_errors;
    cudaMalloc(&d_errors, sizeof(int));

    for (size_t i = 0; i < allocated.size(); ++i) {
        float* ptr = alloc.get_block_ptr(allocated[i]);
        float pattern = static_cast<float>(i + 1) * 1.5f;
        write_pattern_kernel<<<blocks, threads>>>(ptr, elements_per_block, pattern);
    }
    cudaDeviceSynchronize();

    int total_errors = 0;
    for (size_t i = 0; i < allocated.size(); ++i) {
        float* ptr = alloc.get_block_ptr(allocated[i]);
        float pattern = static_cast<float>(i + 1) * 1.5f;
        int zero = 0;
        cudaMemcpy(d_errors, &zero, sizeof(int), cudaMemcpyHostToDevice);
        verify_pattern_kernel<<<blocks, threads>>>(ptr, elements_per_block, pattern, d_errors);
        int err_count = 0;
        cudaMemcpy(&err_count, d_errors, sizeof(int), cudaMemcpyDeviceToHost);
        total_errors += err_count;
    }
    printf("  Pattern verification errors: %d\n", total_errors);

    for (int id : allocated) {
        alloc.free(id);
    }
    printf("  After freeing all: free=%zu\n", alloc.num_free_blocks());

    cudaFree(d_errors);
    printf("  [PASS] Block allocator test complete\n");
}

void test_page_table_mapping() {
    printf("\n[TEST] Page Table - Logical to Physical Mapping\n");
    print_separator();

    size_t pool_size = 64 * 1024 * 1024;
    size_t tokens_per_block = 16;
    size_t token_kv_bytes = 128 * sizeof(float);
    size_t block_bytes = tokens_per_block * token_kv_bytes;

    BlockAllocator alloc(pool_size, block_bytes);

    int b0 = alloc.allocate();
    std::vector<int> skipped;
    for (int i = 0; i < 5; ++i) skipped.push_back(alloc.allocate());
    int b1 = alloc.allocate();
    for (int i = 0; i < 3; ++i) skipped.push_back(alloc.allocate());
    int b2 = alloc.allocate();

    for (int id : skipped) alloc.free(id);

    printf("  Logical block 0 -> Physical block %d\n", b0);
    printf("  Logical block 1 -> Physical block %d\n", b1);
    printf("  Logical block 2 -> Physical block %d\n", b2);

    PageTable pt;
    pt.append_block(b0);
    pt.append_block(b1);
    pt.append_block(b2);

    bool pass = (pt.get_physical_id(0) == b0 &&
                 pt.get_physical_id(1) == b1 &&
                 pt.get_physical_id(2) == b2);
    printf("  Page table entries: %zu\n", pt.num_blocks());
    printf("  Mapping verification: %s\n", pass ? "PASS" : "FAIL");

    int elements = static_cast<int>(block_bytes / sizeof(float));
    int threads = 256;
    int grid = (elements + threads - 1) / threads;

    write_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(b0), elements, 100.0f);
    write_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(b1), elements, 200.0f);
    write_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(b2), elements, 300.0f);
    cudaDeviceSynchronize();

    int* d_errors;
    cudaMalloc(&d_errors, sizeof(int));
    int zero = 0;
    int total_errors = 0;

    for (int logical = 0; logical < 3; ++logical) {
        int phys = pt.get_physical_id(logical);
        float expected = (logical + 1) * 100.0f;
        cudaMemcpy(d_errors, &zero, sizeof(int), cudaMemcpyHostToDevice);
        verify_pattern_kernel<<<grid, threads>>>(alloc.get_block_ptr(phys), elements, expected, d_errors);
        int err_count = 0;
        cudaMemcpy(&err_count, d_errors, sizeof(int), cudaMemcpyDeviceToHost);
        total_errors += err_count;
    }
    printf("  Indirection read-back errors: %d\n", total_errors);

    alloc.free(b0);
    alloc.free(b1);
    alloc.free(b2);
    cudaFree(d_errors);
    printf("  [PASS] Page table test complete\n");
}

void run_multi_sequence_simulation() {
    printf("\n[SIM] Multi-Sequence KV-Cache Simulation\n");
    print_separator();

    size_t pool_size = 256 * 1024 * 1024;
    size_t tokens_per_block = 16;
    size_t token_kv_bytes = 128 * sizeof(float);
    int num_concurrent = 8;
    int total_completions = 100;

    KVCacheManager manager(pool_size, tokens_per_block, token_kv_bytes);

    printf("  Pool: %zu MB | Block: %zu tokens | Token KV: %zu bytes\n",
           pool_size / (1024 * 1024), tokens_per_block, token_kv_bytes);
    printf("  Simulating %d concurrent sequences, %d total completions\n\n",
           num_concurrent, total_completions);

    srand(42);

    struct SeqInfo {
        int seq_id;
        int current_len;
        int target_len;
        bool active;
    };

    std::vector<SeqInfo> slots(num_concurrent);
    int next_seq_id = 0;
    int completions = 0;
    size_t total_allocs = 0;
    size_t total_frees = 0;
    float peak_utilization = 0.0f;
    size_t naive_memory_needed = 0;

    for (int i = 0; i < num_concurrent; ++i) {
        slots[i].seq_id = next_seq_id++;
        slots[i].current_len = 0;
        slots[i].target_len = 32 + rand() % 225;
        slots[i].active = true;
        manager.register_sequence(slots[i].seq_id);

        int prompt_len = 4 + rand() % 29;
        manager.append_tokens(slots[i].seq_id, prompt_len);
        slots[i].current_len = prompt_len;
        naive_memory_needed += slots[i].target_len * token_kv_bytes;
    }

    while (completions < total_completions) {
        for (int i = 0; i < num_concurrent; ++i) {
            if (!slots[i].active) continue;

            if (manager.append_token(slots[i].seq_id)) {
                slots[i].current_len++;
                total_allocs++;
            }

            if (slots[i].current_len >= slots[i].target_len) {
                manager.free_sequence(slots[i].seq_id);
                total_frees++;
                completions++;
                slots[i].active = false;

                if (completions < total_completions) {
                    slots[i].seq_id = next_seq_id++;
                    slots[i].current_len = 0;
                    slots[i].target_len = 32 + rand() % 225;
                    slots[i].active = true;
                    manager.register_sequence(slots[i].seq_id);

                    int prompt_len = 4 + rand() % 29;
                    manager.append_tokens(slots[i].seq_id, prompt_len);
                    slots[i].current_len = prompt_len;
                    naive_memory_needed += slots[i].target_len * token_kv_bytes;
                }
            }
        }

        CacheStats stats = manager.get_stats();
        if (stats.utilization_percent > peak_utilization) {
            peak_utilization = stats.utilization_percent;
        }
    }

    CacheStats final_stats = manager.get_stats();
    printf("  --- Simulation Results ---\n");
    printf("  Sequences completed: %d\n", completions);
    printf("  Total token allocations: %zu\n", total_allocs);
    printf("  Total sequence frees: %zu\n", total_frees);
    printf("  Peak memory utilization: %.1f%%\n", peak_utilization);
    printf("  Final utilization: %.1f%%\n", final_stats.utilization_percent);
    printf("  Internal fragmentation: %.1f%%\n", final_stats.internal_frag_percent);
    printf("  External fragmentation: 0%% (by design - uniform blocks)\n");
    printf("\n  --- Paged vs Naive Comparison ---\n");
    printf("  Paged pool size: %zu MB\n", pool_size / (1024 * 1024));
    printf("  Naive pre-alloc would need: %zu MB\n", naive_memory_needed / (1024 * 1024));
    printf("  Memory saved: %.1fx reduction\n",
           static_cast<float>(naive_memory_needed) / pool_size);
    printf("\n  [PASS] Simulation complete\n");
}

int main() {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    printf("=== Paged KV-Cache Engine ===\n");
    printf("GPU: %s | VRAM: %zu MB | Compute: %d.%d\n\n",
           prop.name, prop.totalGlobalMem / (1024 * 1024),
           prop.major, prop.minor);

    test_block_allocator();
    test_page_table_mapping();
    run_multi_sequence_simulation();

    printf("\n=== All tests passed! ===\n");
    return 0;
}

Your paged KV-Cache engine is fully operational. You have proven that block-based allocation eliminates external fragmentation, reuses memory immediately, and requires a fraction of the VRAM that naive contiguous pre-allocation demands. This is the exact architectural pattern that powers production LLM serving engines like vLLM.

Secret mission

Copy-on-Write Block Sharing

In production LLM serving, beam search generates multiple candidate sequences from a shared prefix. Without memory sharing, each beam duplicates the entire KV-Cache of the common prefix. Your mission: implement copy-on-write semantics so beams share physical blocks until one diverges, saving massive amounts of GPU memory.

Clean Up Your Resources

Clean Up Your Resources

This project runs entirely on your local machine using GPU memory that is freed when the program exits. There are no ongoing costs or cloud resources to worry about.

Decide whether to keep your project files and tools installed, free up disk space by removing build artifacts, or delete everything entirely.

Resources you used:

• paged_kv_cache project directory on your Desktop (source files and compiled binaries).
• CUDA Toolkit 13.3 (GPU compiler and runtime libraries).
• CMake (build system tool).
• Visual Studio 2022 (C++ host compiler).

Keep everything running

No action needed. Choose this if you plan to extend the project further or experiment with different block sizes and pool configurations.

• Your project files stay on your Desktop.
• CUDA Toolkit, CMake, and Visual Studio 2022 remain installed and ready for future GPU projects.
• No GPU memory is held between runs. The program allocates and frees VRAM each time it executes.

Pause - I'll come back to this later

Remove the compiled build artifacts to reclaim disk space while keeping your source code intact.

• Delete the build directory inside your project folder by running:

macOS

cd ~/Desktop/paged_kv_cache
rm -rf build

Windows

cd $env:USERPROFILE\Desktop\paged_kv_cache
rmdir /s /q build

• Your source files (src/, include/, CMakeLists.txt) remain untouched.
• When you return, rebuild with cmake -B build and cmake --build build --config Release.

Delete - I don't want to use this again

Remove all project files and optionally uninstall the development tools.

• Delete the entire project directory by running:

macOS

cd ~/Desktop
rm -rf paged_kv_cache

Windows

cd $env:USERPROFILE\Desktop
rmdir /s /q paged_kv_cache

• Confirm the folder is gone by running ls (macOS) or dir (Windows) and verifying paged_kv_cache no longer appears.

To uninstall the development tools (optional, only if you do not plan to do other GPU or C++ projects):

• Uninstall CUDA Toolkit: open Windows Settings, go to Apps, search for NVIDIA CUDA, and click Uninstall on each CUDA component.
• Uninstall CMake: in the same Apps list, search for CMake and click Uninstall.
• Uninstall Visual Studio 2022: open the Visual Studio Installer app, select your installation, and click More then Uninstall.

Nice Work!

Nice Work!

You just built a GPU virtual memory manager from scratch. This is the exact same architectural pattern that powers vLLM and enables production LLM serving at scale.

You've learned how to:

• Built a physical block allocator that pre-allocates a contiguous GPU memory pool with cudaMalloc and manages it in uniform blocks using a free-list.
• Implemented a page-table system that maps logical token sequences to non-contiguous physical GPU memory blocks, eliminating external fragmentation by design.
• Created a KV-Cache manager that dynamically grows sequences without pre-allocation and simulated 100 concurrent sequence completions demonstrating near-zero memory waste.
• Implemented copy-on-write block sharing for beam search, proving you understand how production systems avoid duplicating shared prefixes.

Ready to quiz yourself?