VibeOps

The moment I decided to learn CUDA, I did the wrong thing first: I opened three tabs comparing RunPod, Lambda Labs, and Vast.ai. I calculated hourly rates. I read a Reddit thread arguing about whether a spot A100 is worth the eviction risk for a beginner.

Then I closed all of it. I didn’t need any of them for the first three days.

Here is what I actually needed: to understand what a kernel is, what threadIdx means, why syncthreads() exists, and why every CUDA tutorial insists on the phrase grid of blocks of threads like it’s a mantra. None of that requires silicon. It requires reading, writing, and running small pieces of code that behave like a GPU without being one.

That is a solved problem. Most of the internet has forgotten it is solved.

Every CUDA Emulator You’ll Google Is Dead — Except One

If you search “CUDA emulator” you get a cemetery.

gpuocelot — an academic project from Georgia Tech that translated PTX to x86. Last real commit in 2012. Mirrors of mirrors on GitHub, none of them buildable against a modern toolchain.

CUDA’s own device emulation mode — a compile flag NVIDIA shipped with early CUDA and killed in version 3.1 (that was 2010). Any Stack Overflow answer suggesting -deviceemu is fifteen years stale.

MCUDA, Ocelot forks, various HPC papers — the same story. Someone published, defended a thesis, moved on. No one maintained anything.

The reason for the graveyard is boring: once cloud GPUs got cheap and Colab started giving away a T4 for free, nobody had the incentive to keep a CPU-side interpreter alive. The market spoke and it said just rent one.

Except for one project, which kept the feature alive because it was useful for a completely different reason: debugging.

Numba is a Python JIT for numeric code. It compiles @numba.jit functions to fast native x86. It also has a @cuda.jit decorator that compiles Python to PTX and runs it on an NVIDIA GPU. And — critically — it has an environment variable, NUMBA_ENABLE_CUDASIM=1, that reroutes those kernels to a CPU interpreter with the full CUDA thread model. threadIdx, blockIdx, blockDim, cuda.shared.array, cuda.syncthreads, atomics. All of it, on your laptop, with no driver, no toolkit, no cloud.

It is not a mock. It is not a wrapper around NumPy. It is a CPU interpreter that thinks like a GPU. And it exists because Numba’s authors needed a way to print from inside a kernel and step through it in pdb — something that is famously painful on real hardware.

Their debug tool is our teaching platform.

Why Slow Is Actually the Feature

The simulator has honest, obvious limitations. It runs on one CPU thread. It is slower than plain NumPy for anything you’d actually want to compute. It doesn’t support warp-level intrinsics (__shfl_*, __ballot_*). It can’t touch cuBLAS or cuDNN. Any benchmark you run in simulator mode is a lie.

This is the feature.

The wrong questions to ask when you are learning CUDA are all performance questions. Is my memory coalesced? Is my occupancy above 50%? Am I hitting the L2? Those questions matter enormously in production. They are noise when you are learning.

The right questions are correctness questions. Did every thread see the right data? Is my index arithmetic actually covering all N elements? When I call syncthreads, do the writes from the other threads actually become visible?

On a real GPU, correctness bugs hide behind wall-clock speed. A race condition that would corrupt one output in a hundred still ships a mostly-right array back to the host, and you spend an afternoon convinced your reduction is off by 0.02%. On the simulator, threads execute in a predictable order, print statements from inside a kernel land in your terminal, and the same race is visible as a wrong number in the very first run.

The simulator strips away the resource anxiety and leaves you alone with the mental model. That is exactly what you want during the first three days.

Dockerize the Whole Thing

I could have installed Numba into a virtualenv on my laptop and been done. I didn’t, because I have burned an evening too many times reconciling NumPy versions with something-that-depends-on-NumPy. A container is fifty lines of config that costs you nothing and buys you never having to think about pip pinning again.

The base image matters. The instinct is to reach for nvidia/cuda:12.4.0-devel-ubuntu22.04 because it says CUDA on the tin. Don’t. That image is 4 GB, ships an nvcc compiler you can’t run without a driver, and gives you exactly zero emulation. It is a base image for a scenario you are not in.

The right base is python:3.11-slim, 150 MB, and one pip install.

FROM python:3.11-slim
RUN useradd --create-home --shell /bin/bash cuda
WORKDIR /app
RUN pip install --no-cache-dir 'numba==0.60.0' 'numpy==2.0.2'
USER cuda

The environment flag goes in Compose, not in the Dockerfile, because I want the same image to work later on a real GPU host — I just override the env there.

services:
  cuda-lab:
    build: .
    image: cuda-lab:sim
    environment:
      NUMBA_ENABLE_CUDASIM: "1"
    volumes:
      - ./src:/app/src
    working_dir: /app

That’s it. One docker compose build, one docker compose run --rm cuda-lab python src/01_hello.py, and you’re inside the CUDA thread model.

Hello, Thread

The Python hello world is one line. The CUDA hello world is a lesson in parallel semantics, and if you write it right you learn something in the first thirty seconds that took me years to internalize.

from numba import cuda
@cuda.jit
def hello_kernel():
    tx, ty = cuda.threadIdx.x, cuda.threadIdx.y
    bx, by = cuda.blockIdx.x, cuda.blockIdx.y
    print("hello from thread (", tx, ",", ty, ") of block (", bx, ",", by, ")")
hello_kernel[(2, 2), (2, 2)]()
cuda.synchronize()

Two lines you have to read carefully. hello_kernel[(2, 2), (2, 2)]() is the launch syntax: a 2×2 grid of blocks, each block being a 2×2 grid of threads. Sixteen threads total, each running the same function.

Run it, and you get sixteen lines of output. Run it again, and you get the same sixteen lines in a different order. That is the whole first lesson right there: the order is not yours to control. Threads run in parallel. If your code depends on any particular thread finishing before any other one, your code is wrong, and it will look right on your laptop and blow up in production.

This is the kind of thing you can read fifty times in a textbook and not feel until you see it print out of order twice in a row.

From there the tutorial ladder is short and mechanical. Next kernel: vector addition, one-dimensional grid, bounds check for the tail element, host-to-device copy, device-to-host copy back, numpy.allclose for verification. The kernel after that: a per-block sum reduction using cuda.shared.array and cuda.syncthreads, which teaches you the two most important CUDA primitives after the thread index itself. Three files, maybe eighty lines each, and you have touched every concept that matters for reading a real CUDA codebase.

I put mine at src/02_vector_add.py and src/03_shared_memory.py. They run in the same container. They pass.

The Escape Hatch to Real Silicon

Here is the part that makes the whole approach worth defending: the same Python file runs on real hardware, unchanged.

Unset NUMBA_ENABLE_CUDASIM, put yourself on a machine with an NVIDIA GPU, and Numba routes the exact same @cuda.jit kernel to a real device. No conditional imports. No if simulator: ... else: .... No rewrite. The file you debugged on your laptop is the file that runs on the A100.

The migration path has two rungs.

Colab, when you want to see it run on a real GPU for the first time. Free tier. Runtime → Change runtime type → T4. !pip install numba, upload the src/ folder, !python src/02_vector_add.py. The PASS prints faster. That’s the entire difference at this stage.

RunPod (or Lambda, or Vast) when you want to do actual work. Change one line in the Dockerfile — FROM nvidia/cuda:12.4.0-runtime-ubuntu22.04 — reinstall Python and Numba on top of it, remove the NUMBA_ENABLE_CUDASIM env from Compose, launch with --gpus all. Same three scripts, same three PASS lines, orders of magnitude faster.

You go from zero silicon to Colab to a rented A100 without rewriting a single kernel. That portability is not a lucky accident; it is the reason Numba’s simulator is worth using instead of some homegrown NumPy stand-in that would only look like CUDA.

What This Is Really About

I am writing this because the default learning path for CUDA — the one the ecosystem pushes — starts with “get a GPU.” Rent one. Beg for one. Colab if you must. And that path is not wrong, exactly. It is just the wrong step one.

Step one is the mental model. Step two is the hardware.

The generalized principle is simple and applies well beyond CUDA: when a domain requires expensive or rare hardware to run, look for a correctness-focused emulator before you pay for the real thing. eBPF has bpftrace. WebAssembly has wasmtime. Embedded firmware has QEMU. CUDA has Numba’s simulator. These tools are not second-best consolation prizes. They are the first-best option for the phase of learning where correctness and mental model matter more than throughput.

The GPU is a tool. The mental model is the product. Rent the tool when you’re ready to build. Not before.

I planned this tutorial before I built it — one paragraph of intent, one page of scope, a short list of files to produce, and a verification checklist. I wrote the plan, ran it end-to-end in about ninety seconds of container build time, and the whole thing worked on the first try. That is only possible because the ground I was walking on — CPU-emulated CUDA in a small Python container — has no dependencies on the outside world. No cloud account. No driver. No GPU. Just Docker and one env var.

That is the kind of foundation you want under a learning curve.