GPU computing means arranging work so that thousands of lightweight threads apply the same operation across large data at once. A GPU is built to maximise throughput, while a CPU core is built to keep latency low for a few threads at a time. A GPU coding program worth your time teaches that execution model first. It then teaches you to tell whether a workload is limited by memory bandwidth, compute or overhead, then how to profile, and only at the end how to write a custom kernel. You do not need to write CUDA from scratch to use GPUs for machine learning. NVIDIA’s own CUDA documentation does say that library and framework users still gain from knowing what happens beneath the abstractions.
What is GPU programming, and what should a GPU coding program teach?
Good GPU code keeps the device’s multiprocessors as busy as possible. NVIDIA’s CUDA C++ Best Practices Guide says this generally takes thousands or tens of thousands of concurrent threads, doing arithmetic on large data sets such as matrices. How much an application gains depends entirely on how far it can be parallelised. Code that cannot be parallelised enough should stay on the host, unless keeping it there would force too many transfers between host and device.
For an engineer working on inference, the curriculum has a definite order, and in our view the order matters more than any single module:
- Execution model. Threads run in warps of 32. Warps are grouped into blocks (cooperative thread arrays), and every warp of a block runs on the same streaming multiprocessor. Global memory costs hundreds of clock cycles of latency, while shared memory sits on-chip with much higher bandwidth.
- Bottleneck regimes. Time on a GPU goes to compute, to moving tensors through memory, or to overhead, and each needs a different fix.
- Data movement between host and device. Memory optimisation starts by cutting these transfers, because they run at far lower bandwidth than transfers inside the device.
- Profiling. Reading Nsight Systems timelines and collecting Nsight Compute metrics for chosen kernels.
- Runtimes and compilers. TensorRT optimises inference through quantisation, layer and tensor fusion, and kernel tuning.
- Kernel authoring. CUDA C++ once a profiled hotspot shows custom code is the best approach, or the Triton language and compiler.
NVIDIA presents optimisation as a loop: find an opportunity, apply it, check the speedup, repeat. The guide also says you need not memorise every strategy before you see good speedups. A course that starts with hand-tuned matrix multiplication before you can read a profile has the order backwards.
GPU vs CPU: when GPU computing is worth the move
On a workload that suits it, the gap is very large. A study of single-precision 4096×4096 matrix multiplication found a discrete NVIDIA GPU 93.5 times faster than the eight-core CPU of the same system. The GPU used about 2% of the CPU’s energy, which is roughly a 50-fold gain in energy efficiency. It also had the highest average power draw: 46.22 W against the CPU’s 24.73 W. It still used the least energy because it finished in 0.613 seconds against the CPU’s 57.34 and dropped back to idle almost at once, which the authors call “race to idle”.
Keep in mind what that study measured. The authors chose matrix multiplication because its high arithmetic intensity makes it an ideal candidate for parallel hardware. It is one favourable workload, and it predicts little about small jobs or memory-bound ones. The architecture behind the gap is covered in our piece on how thousands of GPU cores solve problems differently.
Data movement between host and device is the other limit to watch. NVIDIA’s Best Practices Guide gives a peak theoretical bandwidth of 898 GB/s between device memory and the GPU on a Tesla V100, against 16 GB/s over PCIe x16 Gen3. The guide advises keeping host-device transfers to a minimum, even if that means running kernels on the GPU that show no speedup over the CPU.
Do you need to learn CUDA from scratch to use GPUs for machine learning?
No. The sensible place to start is the libraries. NVIDIA’s guide says parallelising a hotspot can be as simple as calling cuBLAS, cuFFT or Thrust. From Python, CuPy works as a drop-in replacement for NumPy and SciPy code on NVIDIA CUDA or AMD ROCm. Its routines run on cuBLAS, cuFFT, cuSPARSE and related libraries. When you outgrow it, CuPy also lets you write raw CUDA kernels, generate kernels from Python, or fuse several operations into one kernel.
A good library makes a large difference with no hand-written code. In frameworks such as PyTorch, the linear operator usually runs on a GEMM backend. One survey of LLM inference hardware reported Llama2-7B on a single A100 at 44.60 tokens/s with the naive HuggingFace implementation. Frameworks reached 88.45 to 93.63 tokens/s.
You can skip writing CUDA, but you cannot skip the mental model. NVIDIA’s CUDA Programming Guide makes the same point: knowing how GPUs execute code helps even developers who mainly use libraries, frameworks or DSLs. Is CUDA hard to learn? In our judgement the syntax is the easy part. The harder part is memory behaviour, which NVIDIA calls the most important area for performance. For how this plays out in training and inference pipelines, see how to use GPU programming in machine learning.
What hardware should you use to learn GPU programming?
Learning CUDA needs an NVIDIA GPU, because CUDA is NVIDIA’s proprietary platform. Two practical details matter for profiling. First, NVIDIA drivers need elevated permissions by default to read GPU performance counters, so you need root access or non-admin profiling turned on. Second, NVIDIA recommends profiling on a GPU without an active display. Display and other engines share L2, DRAM and PCIe, and on small kernels they can blur the DRAM figures.
If your hardware is AMD, HIP is the native language extension and CuPy also runs on ROCm. An integrated GPU still teaches the core lesson: in the same matrix multiplication study, an integrated GPU running OpenCL was 27.3 times faster than the CPU.
NVIDIA’s Nsight Compute guide warns that cache sizes, SM counts, register files, peak FLOP rates and memory bandwidth all change which part of a workload limits performance on a new architecture. A bottleneck you find on a laptop GPU is a hypothesis to re-check on the deployment hardware, not a conclusion.
Which languages and frameworks are used for GPU programming in 2026?
The choice comes down to how much control you need and how many hardware vendors you have to support. The main options:
| Option | What it is | Where it fits |
|---|---|---|
| CUDA C++ | NVIDIA’s proprietary language extension to C++ | NVIDIA-only fleets; best or within 3% of the best model in eleven of fifteen cases in one portability study |
| HIP | Open-standard language extension | AMD GPUs |
| SYCL | Open-standard C++ extension with explicit or implicit memory management | Mixed-vendor code; the most consistent performance portability in the same study |
| Kokkos, RAJA | C++ abstraction libraries that compile to CUDA, HIP or OpenMP backends | Portable C++ codebases; close behind SYCL on portability |
| OpenCL | Long-standing standard with large amounts of legacy scientific code | Existing OpenCL code; kernels can compile to SPIR-V and run on Vulkan |
| Triton | Open-source language and compiler for neural network kernels, written in Python | Custom deep learning kernels with less effort than CUDA |
| CuPy | NumPy/SciPy-compatible GPU array library | Array code that should run on a GPU with minimal changes |
The portability figures come from Davis et al. They compared CUDA, HIP, Kokkos, RAJA, OpenMP, OpenACC and SYCL on NVIDIA and AMD supercomputers. Their numbers are deliberately out of the box: they did not tune grid size, block size or shared memory, so the results reflect minimal porting effort. Within that limit, the portable models matched or beat vendor models on XSBench, a complex kernel, with similar development effort. A compute-bound kernel with high register pressure, however, was highly sensitive to which portable model was used.
Vendor-specific APIs are not faster in themselves: the same code performs essentially the same. Their advantage is access to lower-level features, and using that access well takes real commitment, which is not always worth it. For a single-vendor NVIDIA inference stack, we would stay in the CUDA ecosystem and reach for Triton when a kernel is needed. For a mixed fleet, SYCL or Kokkos earns its place. The trade-offs of CUDA framework and ecosystem lock-in, what makes a GPU CUDA-capable, and the kernel-authoring choice between Triton and CUDA each have their own page.
What limits GPU computing performance in practice?
Memory, in the HPC proxy kernels Davis et al. studied, though only measurement shows which limit you are hitting. “The GPU is slow” breaks down into three regimes with three different fixes.
Compute, bandwidth and overhead need different fixes
If all your time goes on memory transfers, a GPU with more FLOPS will not help. If all your time goes on large matrix multiplications, cutting framework overhead will not help. Overhead covers time in the Python interpreter, in the PyTorch framework and in launching CUDA kernels. When operators are too small, the CPU cannot queue work ahead of the GPU, and the GPU spends its time idle. Overhead does not scale with problem size, which gives a quick test. Double the batch size: if runtime rises only about 10%, you are probably overhead-bound.
Memory bandwidth is the usual ceiling
NVIDIA’s Best Practices Guide calls bandwidth one of the most important limits on performance and says almost every code change should be judged by its effect on bandwidth. Davis et al. found every kernel they examined to be memory-bound, apart from one strongly compute-bound case. The roofline chart makes this concrete. It plots a kernel’s arithmetic intensity (work per byte of memory traffic) against the GPU’s peak compute and bandwidth. Kernels that land under the sloped bandwidth boundary are memory-bound, and the distance to the boundary shows how much improvement is available.
LLM inference shows why this matters. Profiling Llama2-7B on a single A100 showed the decode stage taking more than 80% of time for typical prompt and output lengths. At batch size 1, decode turns GEMM into a matrix-vector problem with reduced dimensions. The same survey notes that quantisation eases the memory-access bottleneck in LLM inference and frees up compute that was going unused.
Host-device transfer is the slowest link
NVIDIA calls PCIe transfers between host and device the slowest link in GPU computing. Pinned host memory, allocated with cudaMallocHost() or cudaHostAlloc(), removes an intermediate copy and more than doubles transfer rates on typical systems. Batching many small transfers into one larger transfer removes the cost each transfer carries. Pinned memory is scarce, though, and using too much of it can slow the whole system. Streams let transfers overlap with computation. Transfer cost is not always the limit: in the Davis et al. runs, data movement took less than 2% of kernel time.
Occupancy is a diagnostic, not a target
Occupancy is the ratio of active warps per multiprocessor to the maximum possible. Low occupancy always hurts latency hiding, but higher occupancy does not always mean higher performance. NVIDIA’s own libraries show this. On G80, CUFFT 2.3 ran at 17% occupancy and reached 93 Gflop/s, while CUFFT 2.2 ran at 33% and reached 45 Gflop/s. Instruction-level parallelism explains it: independent operations within one thread also hide latency, so fewer warps are needed.
| Regime | How it shows up | What helps |
|---|---|---|
| Overhead-bound | Runtime barely grows with batch size; gaps on the GPU timeline | Tracing, jax.jit, CUDA Graphs, at some cost to flexibility |
| Memory-bandwidth-bound | Kernel sits under the roofline’s bandwidth slope | Coalescing, shared memory reuse, fusion, quantisation |
| Memory-latency-bound | L1TEX scoreboard stalls dominate while bandwidth is not saturated | ILP, unrolling, pipelining rather than bandwidth work |
| Compute-bound | Achieved FLOPS close to peak | Lower precision formats and specialised units such as Tensor Cores |
| Transfer-bound | Host-device transfers, the slowest link, set the pace | Fewer, batched, pinned and overlapped transfers |
Why profile a GPU workload before optimising it?
The bottleneck decides which fix works, so a fix chosen before measuring is a guess. Horace He makes this point and describes people trying speedups blindly without knowing which regime they are in. NVIDIA’s TensorRT documentation says the same: measure latency and throughput on your own model and hardware first. Its rewritten Best Practices guide is structured as a measure-then-optimise loop, including reading Nsight Systems timelines. The workload also has to be realistic. Unrealistic problem sizes send developers to optimise the wrong functions.
We hold that best practices such as “avoid branches” only make sense in context, and that the context comes from a profile, not from blog posts. NVIDIA’s guide ranks its own recommendations by priority for the same reason: apply the high-impact ones first and avoid premature optimisation.
A profiling pass that produces numbers you can trust runs in this order:
- Start with the timeline. Nsight Systems collects data in a single pass and does not lock clocks or flush caches. Gaps where the GPU waits for the CPU point to overhead.
- Narrow to the dominant kernels in Nsight Compute. Limit collection to the kernels that matter, because every extra metric and kernel adds overhead and replay passes.
- Control clocks and caches before comparing. Nsight Compute locks SM clocks by default; NVIDIA recommends locking clocks externally with
nvidia-smiand passing--clock-control none. Nsight Compute also flushes caches between replay passes. For cache-sensitive work, use--replay-mode application --cache-control none. - Read the result in roofline and issue-slot terms. The goal is higher issue-slot utilisation, not every stall counter at zero.
- Change one thing and judge it by duration. Use an optimised build. Utilisation percentages can move either way when total work changes, so absolute duration is the ground truth.
- Re-profile on new hardware. The dominant limiter can shift between GPU architectures.
The same discipline at fleet level, comparing devices on your own workload rather than on spec sheets, is a separate task covered in benchmarking GPUs against your own workload.
When does an inference team actually need custom GPU kernels?
Only when a profiled hotspot shows that the optimised libraries and the runtime fall short, and the gap is large enough to pay for the code you will then maintain. NVIDIA’s guide sets out the profile-first part of that test. Once a profile has located a hotspot and custom code has been judged the best approach, CUDA C++ exposes that portion as a kernel without major rewrites elsewhere. The approach is most straightforward when running time is concentrated in a few isolated parts of the code.
Before that point there is a ladder to climb. TensorRT compiles trained PyTorch and ONNX models into optimised engines. It supports mixed precision from FP32 down to FP8, INT8, FP4 and INT4, with optimisations for transformers and LLMs. Horace He calls operator fusion the most important optimisation in deep learning compilers, because it avoids writing data to global memory only to read it back. Above the runtime sit generated or DSL kernels. OpenAI reports Triton kernels up to twice as efficient as equivalent Torch implementations.
Hand-written kernels earn their place where the general paths are poorly matched to the shape of the work. LLM decode is one such case. TensorRT-LLM added a dedicated GEMV kernel for batch size 1 and reached 98.19 tokens/s on Llama2-7B on one A100. FlashDecoding++ targets inefficiencies in cuBLAS and CUTLASS at small decode batch sizes. Weight-only quantisation is another case. It shrinks the memory footprint, but dequantising from integer to float can cancel the compute gain. One design fuses dequantisation with GEMV, reaching 45.2 tokens/s for Llama2-7B on an RTX 3090.
TechnoLynx approaches GPU performance engineering from this same profile-first position. The question to settle before writing one is whether the profile shows a stage large enough for even a perfect kernel to matter.
FAQ
Is it better to be GPU bound or CPU bound?
For an accelerated workload, being GPU-bound is the healthier state. When operators are too small for the CPU to queue work ahead, the GPU sits idle and becomes, in Horace He’s words, “an expensive paperweight”.
How do GPU kernels work?
A kernel launches as a grid of independent thread blocks, and each block’s warps of 32 threads run on one streaming multiprocessor. Because blocks cannot wait on one another, a grid too large to fit on the hardware at once still runs correctly on any GPU.
Can a CPU work without a GPU?
Yes, even for LLM inference. The 2024 hardware survey reports 12.5 to 50 tokens/s for 6B to 20B parameter models with INT4 kernels on a single socket of a 4th Generation Intel Xeon processor.
