The short version: GPUs give us enormous parallel throughput, but that throughput is exactly what makes observing single operations in a kernel nearly impossible. Thousands of lanes execute at once, with a device state that's hard to inspect and hardware debuggers that are extremely heavyweight. By moving PTX into a deterministic CPU interpreter where every thread, register write, memory access and synchronization event can be captured, warpsim allows us to inspect and find extremely nuanced bugs. The goal is not to run GPU code fast on the CPU, but to make GPU kernel execution understandable.
The problem: GPU execution is fast and opaque
If you have ever had to find the exact instruction that wrote a NaN into a buffer, you've felt our pain. On real hardware the kernel runs in microseconds across thousands of lanes and by the time you notice the bad value the state that produced it is gone, so you can try to add printfs, you can reach for a hardware profiler, you can bisect launch parameters, but all of it fights the same fact: the hardware is optimized for execution, not explainability.
That's the tradeoff we accept for throughput and most of the time it's the right trade, but when we're testing a small kernel for functional correctness, chasing a numerical bug or teaching a person or agent how PTX maps onto the CUDA thread model, throughput is not what we are trying to optimize for, instead we are hoping to see one instruction and one register at a time.
To this extent we built warpsim which is a Rust-based simulator that parses PTX, lowers it into an internal instruction representation and then executes it against simulated GPU state on the CPU. Off the bat: It's not a CUDA replacement and by no means a performance emulator, what it is, is a deterministic PTX interpreter which is the whole point: the same kernel produces the same execution every time, which is how we make bugs reproducible and traces meaningful.
The model is super simple: allocate device memory, write inputs, launch with grid and block dimensions, read results, but the big difference is that everything happens inside a Rust process and everything that happens is observable.
The memory model and why pointer tags matter
One of the decision we made early on is that the simulator keeps separate memory spaces that map directly onto PTX's: global (device memory), const (initialized from PTX constant declarations), shared (per-CTA), local (per-thread), param (per-thread launch and function parameters) and generic.
When the host passes a simulated device pointer as a kernel argument, the VM records not just the numeric address but metadata such as that the pointer refers to global memory. When PTX later perform address-space conversions, that tag lets our VM enforce state-space correctness and catch invalid pointer usage, so that a class of bugs that is nearly invisible in raw PTX, such as loading through the wrong address space or losing pointer provenance across generic conversion, are trivial to point at directly.
Threading and synchronization
The VM does not run GPU threads truly in parallel. It simulates many CUDA threads cooperatively inside a single CPU interpreter and it does so under an explicit scheduler. For each CTA the scheduler maintains runnable and blocked thread queues and synchronization instructions can park a thread until enough participants arrive.
The modeled synchronization covers CTA barriers (bar.sync), nonblocking arrivals (bar.arrive), warp barriers (bar.warp.sync), memory barriers and selected collectives such as shuffle, vote and MMA-style warp operations. When a collective completes, blocked threads move back into the runnable queue and because the scheduler is explicit, the trace can show exactly when a thread waited and what released it.
There are two scheduler modes and the choice matters for what you're doing. The default is RunUntilBlocked: a thread runs until it exits or blocks. The other is SequentialExecution, which executes one instruction from each runnable thread at a time. Sequential execution is the one you want for tracing, because it makes inter-thread execution order easy to follow visually instead of showing you one thread running to completion before the next one starts.
To be clear here: The simulator is deterministic, which is what makes it good for reproducing bugs, but it is not a model of real GPU scheduling (though we are working on that), timing, memory consistency or SIMT reconvergence, it's built to tell us what the threads compute, not how fast or in what hardware order.
What we can actually do with it
The most direct use-case is debugging and this is where turning a kernel launch into a regular Rust-controlled execution pays off. You can find the exact instruction that wrote a bad value, see which thread took which branch, watch how %tid, %ctaid and pointer arithmetic combine into an address, catch out-of-bounds or misaligned or wrong-state-space accesses, chase where a NaN was created or stored and understand why a barrier or warp collective blocked. For example, we build a small hook that prints diagnostics whenever a NaN value is stored to global memory or produce diagnostic whenever a floating-point op produces a NaN.
But there are many other use-cases for the tech that we're working on, for example:
- Dense reward signals instead of a binary pass/fail "Output mismatch" is a terrible learning signal, but "Thread (14, 0, 0) wrote a NaN at line 3042 because the mbarrier phase was stale and here's the register state" is a great one. An agent can get instruction-level divergence localization against a reference: run both, diff elementwise, report the first instruction where state diverges, which is basically an automatic root-cause bisection.
- Denser debugging info. The ability to go back and ask "who last wrote this shared memory byte?", trace registers or step single-warps is extremely valuable for agents, because they're actually pretty good at interactive debugging loops when the tools are scriptable and a simulator makes everything scriptable in a way that (cudad-)gdb never will be.
- Partial reward signals Because it's trivial to see where execution diverges from the reference, we can reward partial progress such as "the mainloop is now correct, the epilogue is still broken", which lets us train the model to predict trace features (which barrier fails, which register holds garbage) as auxilary objectives
- Self-play: bug-injector vs bug-fixer By making it trivial to confirm and localize the break, we can mutate correct kernels into slightly correct ones with dense labels, allowing us to generate unbounded curriculum with automatic difficulty scaling, which is what makes self-play work elsewhere
The honest limits
warpsim is about functional behavior, not GPU speed. It does not model true hardware parallelism, real warp scheduling, cache hierarchy, timing or throughput. Some PTX operations are unsupported or only partially modeled. It does not fully model SIMT divergence and reconvergence the way NVIDIA hardware does.
We specifically design warpsim because CUDA is built to make individual instruction behavior disappear into aggregate throughput, whereas warpsim is built to make it reappear in a deterministic, inspectable, replayable and diffable manner, not to run GPU code fast on the CPU, but to make kernel execution understandable.
