CUDA Playground: learn how GPU kernels run, online

Q: Can I learn CUDA without an NVIDIA GPU?

Yes. The ideas that trip people up — the thread index, launch sizes, bounds checks, threads overwriting each other, warp divergence — are about how threads are organised, not about speed. This playground runs your kernel thread by thread in the browser and shows each of them, which is often clearer than running it on a real GPU.

Q: How do I run CUDA code online?

Paste or write a __global__ kernel in the editor, set the number of blocks and threads per block, give the arrays a length and starting values, and press Run kernel. The page shows what every thread did and the arrays before and after.

Q: How do I calculate the thread index in CUDA?

For a one-dimensional launch: int i = blockIdx.x * blockDim.x + threadIdx.x; Click any thread in the playground to see the numbers filled in, for example 2 × 8 + 5 = 21.

Q: How many blocks should I launch?

Enough to cover every element: blocks = (n + threadsPerBlock − 1) / threadsPerBlock, which rounds up. “Fit blocks to N” does this. Then add if (i < n) for the spare threads in the last block.

Q: What is warp divergence and how do I avoid it?

The 32 threads of a warp run in lockstep. If some take the if and others the else, the warp runs both, one after the other. Make branches depend on whole warps (threadIdx.x / 32) or on data that is the same for neighbouring threads. Compare the “Warp divergence” and “No divergence” examples.

Q: Why do my results change between runs on a real GPU?

Usually a race: two or more threads write the same element. The playground marks such elements in orange and lists which threads wrote them. Use atomicAdd, or give each thread its own element.

Q: What part of CUDA does the playground understand?

One __global__ kernel with int and float scalars and arrays, local variables, if/else, for, while, break and continue, arithmetic, comparisons, bit operations, casts, min, max, sqrtf and similar, and threadIdx, blockIdx, blockDim and gridDim. Shared memory, __syncthreads, atomics and 2-D launches are planned.

Q: Does it show how fast my kernel would be on a GPU?

No. It runs one thread at a time to show you what each thread does; it does not use your graphics card or estimate real timings. The warp numbers show relative cost from divergence only.

See how a CUDA kernel really runs, without an NVIDIA GPU. Write or pick a kernel, choose how many blocks and threads to launch, and watch every thread: which element it computes, which threads sit idle, which ones write past the end of an array, where threads overwrite each other, and which warps diverge.

Free, nothing to install, and it works on any laptop or phone.

Try an example:

1Kernel

Example

2Launch

kernel<<<,>>>(…)

Loading…

3Threads

wrotedid nothingerrorwrote where another thread also wrotedivergent warpOne row = one warp of 32 threads. Click a thread.

Run the kernel, then click a thread to see what it did.

4Memory

Index	Before	After	Writers	Last written by

The kernel runs in a simulator in your browser, one thread at a time, so you can see each thread. It does not use your graphics card or measure GPU speed. Nothing is uploaded.

About the CUDA Playground

CUDA is NVIDIA’s way to run code on a graphics card, where thousands of small cores work at the same time. The hard part for beginners is not the syntax but the picture: which thread handles which element, what happens to the threads left over, and why some kernels give wrong or slow results. This playground draws that picture for any kernel you write.

What you can do with it

See the thread index formula worked out for every thread of a launch.
Find out-of-bounds reads and writes, with the thread, the index and the line.
Spot threads overwriting each other’s results, and elements no thread wrote.
Watch warps diverge on an if, and see how to avoid it.
Try grid-stride loops, launch sizes and block sizes before running on a real GPU.
Teach or learn GPU programming in a class, with nothing to install.

Your first CUDA kernel: vector addition

On a CPU, a loop adds the elements one after another. In CUDA, the loop disappears: each thread adds one pair.

// CPU
for (int i = 0; i < n; i++) c[i] = a[i] + b[i];

// GPU: one thread per element
__global__ void vecAdd(const float* a, const float* b, float* c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) c[i] = a[i] + b[i];
}

// launch enough blocks of 256 threads to cover n elements
vecAdd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);

The six steps of a CUDA program

Prepare data on the hostThe CPU (the host) fills arrays in its own memory.
Allocate GPU memorycudaMalloc reserves space in the GPU’s memory (the device).
Copy the inputs overcudaMemcpy(…, cudaMemcpyHostToDevice) moves the data to the GPU.
Launch the kernelkernel<<<blocks, threadsPerBlock>>>(…) starts thousands of threads, each running the same function on its own element.
Copy the results backcudaMemcpy(…, cudaMemcpyDeviceToHost) brings the answer to the CPU.
Free GPU memorycudaFree releases what was allocated.

The playground covers step 4, the kernel itself, where most mistakes happen; it fills the arrays for you in place of steps 1 to 3.

CUDA concepts in one page

Kernel: A function marked __global__ that runs on the GPU. One launch runs it once per thread.
Thread, block, grid: Threads are grouped into blocks (up to 1,024 threads each); all the blocks of one launch form the grid. The launch <<<3, 8>>> makes 3 blocks of 8 threads.
threadIdx, blockIdx, blockDim, gridDim: Each thread’s position: its number inside the block, its block’s number, the block size and the number of blocks.
Global index: blockIdx.x * blockDim.x + threadIdx.x gives every thread a different number across the whole grid, used to pick its element.
Bounds check: Blocks come in whole sizes, so a launch usually has a few more threads than elements. if (i < n) keeps the spare threads from touching memory past the end.
Warp: The GPU runs threads in groups of 32 that execute the same instruction together. In the picture above, each row is a warp.
Warp divergence: When threads of one warp take different sides of an if, the warp runs each side in turn with part of its threads switched off, so it takes longer.
Race condition: Several threads writing the same element at the same time: updates are lost and the result changes between runs. atomicAdd or a reduction avoids it.
Streaming multiprocessor (SM): The hardware unit that runs blocks. A GPU has dozens to over a hundred; blocks are spread across them.
Shared memory: Fast memory shared by the threads of one block, used to reuse data (as in tiled matrix multiplication). Coming to this playground later.

When does a GPU help?

Work	Why	GPU?
Large arrays, the same work on every element	Images, vectors, matrices, neural-network layers: thousands of independent pieces keep thousands of threads busy.	Yes
Small jobs	Copying data to the GPU and back costs more than the work saves.	No
Each step needs the previous result	A long chain of dependent steps cannot be split across threads.	No
Lots of different branches per element	Warps diverge and most threads wait.	Poorly

Most of the work in AI models is matrix multiplication: millions of independent multiply-adds, exactly what GPUs are built for. Libraries such as cuBLAS and cuDNN, called by PyTorch and TensorFlow, contain hand-tuned kernels so most people never write CUDA themselves.

Questions

Can I learn CUDA without an NVIDIA GPU?

Yes. The ideas that trip people up — the thread index, launch sizes, bounds checks, threads overwriting each other, warp divergence — are about how threads are organised, not about speed. This playground runs your kernel thread by thread in the browser and shows each of them, which is often clearer than running it on a real GPU.

How do I run CUDA code online?