GPU Lab New to kernels? Begin with Hello, GPU →
HANDS-ON PROTOTYPE · CUDA C++ · BROWSER SIMULATION
Repair a kernel
Repair a kernel, explain its failures, then build vector addition.
Begin with y = 2 * x, then transfer what you learn to y = x + b: the elementwise addition used in residual connections in neural networks.
A kernel is a function run many times; one run is an thread. An thread with a valid element index should write one output value. Array indices start at zero.
Guided challenges
Start a challenge, predict its behaviour, then run, repair and explain. Use the editor below freely between challenges.
Your task: pass all three checks
For x = [0.5, -1, 1.5, -2, …], the output must start [1, -2, 3, -4, …]. Check lengths 8, 32 and 35, including the last element of each vector.
1. Build your kernel
A block is a group of threads. threadIdx.x is a thread’s index within its block, blockIdx.x identifies the block, and blockDim.x is the number of threads per block. Combining them gives an index across blocks.
Choose CUDA C++ statements. The browser simulates these supported choices on the CPU; it does not compile CUDA or execute GPU code.
n is the vector length. factor is 2.
Browser simulation · no CUDA compilation or GPU execution.
In PyTorch: y = 2 * x. This is the operation we are implementing, not PyTorch’s internal kernel source.
Start the first guided challenge, or explore the editor below.
Enable JavaScript to build the source.Run this kernel on an NVIDIA GPU
Download a complete CUDA C++ program for 35 elements, including your kernel, allocation, launch and output read-back. Invalid accesses and duplicate writes must be repaired before download.
Run safe downloaded kernels with GPGPU-Sim without a GPU. For native execution, use an NVIDIA GPU, compatible driver, CUDA Toolkit and host C++ compiler:
nvcc -std=c++17 kernel.cu -o kernel ./kernel
These new generated programs have not yet been compiled or tested on an NVIDIA GPU.
Hint: where should the first thread write?
The first output is y[0]. Work out i when blockIdx.x * blockDim.x + threadIdx.x is zero. Then check whether that expression still reaches the right elements in a second block.
2. Compare finished outputs
| Length | Browser model | Inspect |
|---|
3. Explain the result
Blocks → threads → output elements
Select an thread to inspect its operations. These groups show membership, not execution order or physical GPU placement. “Duplicate” means multiple threads target the same element.
This is a replay of the browser model, not a captured GPU trace. It visits threads by global ID for inspection; real threads can run in a different order. Each step is a source operation, not a GPU instruction or clock cycle.
Input x · highlighted at a read
Each tile shows its element index and value. FP32 elements occupy four bytes; offsets are relative to this buffer.
Second input b
Model output y after this step
A dash means no thread has written that element yet. Finished-run checks above are independent of the replay position.
Keep investigating
- Pass all three lengths. Which element did the starting kernel miss?
- Inspect length 35. Jump to the last thread: why does its guard skip the memory access?
- Remove the bounds check and run the model. Find the first invalid access, then repair it.
- Try the local index. Why can a single-block test pass while length 35 fails?
Unsafe configurations are diagnosed by the model and cannot be downloaded until repaired.
What the simulation establishes
The browser computes indexing, guards, FP32 arithmetic, reads, writes, duplicate writes and unfilled outputs for the supported expressions. It does not compile CUDA, simulate SM scheduling, caches or warps, or predict time. Replay order is chosen for inspection, not GPU execution order.
Downloaded programs use CUDA C++ and require an NVIDIA GPU and CUDA Toolkit. Browser checks do not establish hardware correctness or performance.
Run the example on your desktop
See the complete example index and setup appendix for CPU PyTorch, the supported GPGPU-Sim programs and optional NVIDIA GPU execution.
The browser still requires no installation. Downloads are generated from the displayed source; CPU checks and simulator checks do not establish GPU performance.
References
Use these official references to check the concepts and code on this page.
Function names in the repair examples are chosen for this course. The CUDA guide documents their qualifiers, indexing variables and launch syntax.
- CUDA: built-in variables (new tab)
Check the meanings of threadIdx, blockIdx and blockDim.
- CUDA: writing and launching kernels (new tab)
Relate block counts, thread counts and bounds checks to the repair.
These browser exercises do not run an installed PyTorch or CUDA toolkit. For optional local execution, choose documentation matching your installed tools. Links open online in a new tab; the questions and answer comparisons below work offline.
Use the documentation
Use the references to explain each answer, then apply it to this page’s example. No installation is needed.