GPU kernel engineering
CUDA Kernel Optimization: Profiling, Latency and Business Value
A real client engagement. The engineering and the results are described below.
A document-analysis company we worked with is about to expand its paid API when its latency budget runs out. The team has budget for more GPUs, but first asks whether one expensive operator is wasting the capacity it already owns.
Request a time through the inquiry form. A meeting is confirmed separately by email.
The business problem behind the technology
Can a kernel change rescue the response budget?
The client’s 100 ms compute path is too close to the product’s internal budget. One elementwise operator occupies 35% of that path; the rest is outside this optimization.
Read the client engagement ↓Client engagement / Delivered results
A larger fleet was not the first decision
A document-analysis provider we worked with plans 50 million comparable API requests a month. Its leadership wants the new workload accepted without hiding slow responses behind a higher average throughput.
The constraint
The client’s 100 ms compute path is too close to the product’s internal budget. One elementwise operator occupies 35% of that path; the rest is outside this optimization.
The engineering decision
The team profiles the application, defines the operator’s numerical contract and tests a candidate that halves only that operator’s time. It keeps the original implementation as a rollback and refuses to price fleet savings until capacity can actually retire.
The delivered outcome
In this engagement, 65 ms of unchanged work plus 17.5 ms of optimized work gives an 82.5 ms path: a 17.5% latency reduction across the request path.
| Measure / unit | Before | After | Difference |
|---|---|---|---|
| Compute-path time seconds/request | 0.1 | 0.0825 | 0.0175 |
Compute-path time. The 0.0175-second difference follows from the 35% operator share and the 2× operator speed delivered in the engagement.
The conditions behind the results
- The 35% operator share matches the separate introductory teaching assumption.
- The candidate halves operator time with identical supported inputs and numerical acceptance.
- Other computation and transfers are unchanged; queueing, retries and network time are excluded.
What this does not prove. The result was scoped to the request path; fleet savings were priced only when capacity could actually retire.
Evidence to collect for your own decision
- Collect an application profile and the operator’s share of completed request time.
- Compare numerical errors, awkward shapes and synchronized end-to-end timing.
- Check whether billable capacity can retire before claiming cash savings.
Key decisions
A useful GPU kernel is correct on awkward inputs, faster on representative workloads, and worth maintaining. Follow a small operation from thread ownership through memory traffic and profiling to an honest production comparison.
- Collect an application profile and the operator’s share of completed request time.
- Compare numerical errors, awkward shapes and synchronized end-to-end timing.
- Response-time headroom is not an invoice reduction, additional sales or a promised hardware speedup.
Follow the decision
Select a step to follow its reasoning, then continue into the technical chapters.
Problem → boundary → decision → evidence
Follow the complete request before selecting a kernel to replace.
Read every component and connection
- Find the delay · Problem
- Follow the complete request before selecting a kernel to replace.
- Prove the operation · Boundary
- Establish independent correctness and the supported input contract.
- Change the constraint · Decision
- Test memory reuse, launch overhead or arithmetic with an explicit hypothesis.
- Measure useful progress · Evidence
- Keep the application-level result and maintenance cost in the decision.
- Find the delay → Prove the operation: identify the constraint
- Prove the operation → Change the constraint: choose a bounded change
- Change the constraint → Measure useful progress: check the outcome
A conceptual decision map for this article, not a measured timeline, physical topology or a depiction of a specific client system.
Hover, focus or tap a component to inspect it. Motion adapts automatically to connection, device and accessibility signals; the component key remains readable without JavaScript.
Find the bottleneck before writing a kernel
The client’s API team starts with the 35% operator share, because that share sets the ceiling on the business case.
Begin with a trace of the complete request: preprocessing, host-to-device copies, kernels, synchronization and response handling. Identify the exact operation, shape distribution, precision and concurrency that dominate a relevant workload. A profiler hotspot in an isolated loop may disappear behind network waits in production.
Try the supported library or framework compiler first. cuBLAS, cuDNN and fused framework operations already encode substantial hardware knowledge. In this scenario, making a 35% component infinitely fast caps total speedup at 1 / 0.65, about 1.54×. That Amdahl bound assumes the rest is unchanged; it is not a promised improvement.
Reference
Choose an independently correct implementation and numerical tolerances.
Candidate
Change one bottleneck: access pattern, reuse, launch or synchronization.
Compare
Retest correctness and the original end-to-end workload.
The moving marker shows a workflow, not execution time or a speedup.
Give each thread explicit ownership
The investigation narrows to an elementwise operation. Before choosing a clever launch, you give every thread an unambiguous piece of work.
For independent elementwise addition, a thread owns an output index. A grid-stride loop lets a bounded launch cover a larger array. The unsigned size calculation avoids narrowing a large product to a 32-bit index. The example assumes valid device pointers, compatible buffer lengths and a launch configuration chosen by the host.
Do not launch a zero-block grid for an empty input; handle that in the caller. This kernel is not a complete benchmark harness. It intentionally leaves allocation, transfers, stream selection, launch-error checks and synchronization to the surrounding application.
Software ownership, from cluster to silicon
Kubernetes places the workload. The container stack exposes a device. CUDA and the driver submit work. GPU scheduling and execution happen below that boundary.
Cluster lifecycle and placement
Application and user-space libraries
Kernel and device boundary
Observe the real service
The scheduler chooses a suitable node using declared resources and placement policy. It does not place individual CUDA blocks, schedule warps or make an unready model useful.
Read every component and connection
- Kubernetes scheduler · Node / resource placement
- The scheduler chooses a suitable node using declared resources and placement policy. It does not place individual CUDA blocks, schedule warps or make an unready model useful.
- Device plugin / kubelet · GPU allocation
- A device plugin advertises supported GPU resources and participates in allocation. Kubelet starts the pod with the chosen devices. GPU sharing changes the resource contract.
- GPU Operator · Lifecycle controller
- The operator can manage drivers, Container Toolkit, device plugin, feature discovery and monitoring. This is a control-plane responsibility, not a mandatory hop for each tensor.
- OCI runtime / Toolkit · Container device access
- The configured runtime and NVIDIA Container Toolkit expose the permitted devices and libraries. Containers share their host kernel; a VM or bare-metal host sets another boundary.
- Serving framework · Batching / model / NCCL
- A serving engine schedules requests and calls framework/library kernels. NCCL coordinates supported multi-GPU collectives. Application queueing is distinct from GPU instruction scheduling.
- CUDA runtime / driver · libcudart / libcuda
- CUDA APIs manage contexts, streams, memory and launches. The user-mode driver interfaces with kernel support. Compatible machine code or supported PTX compilation is required.
- Linux NVIDIA modules · Devices / memory / I/O
- Kernel modules cooperate with the OS for device access, memory mapping and control. Linux permissions and isolation still matter; a container image does not replace the host kernel driver.
- Queues / GPU firmware · Device work submission
- Supported driver and firmware paths submit and manage device work. Firmware-mediated responsibilities vary with platform and driver; this is not a proprietary command-protocol schematic.
- Kubernetes scheduler → Device plugin / kubelet: placement / allocation
- Device plugin / kubelet → OCI runtime / Toolkit: start with devices
- GPU Operator → Device plugin / kubelet: manage plugin
- GPU Operator → OCI runtime / Toolkit: manage toolkit
- GPU Operator → Linux NVIDIA modules: manage driver
- OCI runtime / Toolkit → Serving framework: run workload
- Serving framework → CUDA runtime / driver: library calls
- CUDA runtime / driver → Linux NVIDIA modules: driver interface
- Linux NVIDIA modules → Queues / GPU firmware: submit / manage
- Queues / GPU firmware → SMs / memory / ALUs: device work
- SMs / memory / ALUs → DCGM + application telemetry: device signals
- Serving framework → DCGM + application telemetry: service signals
A typical Linux/device-plugin deployment, not a universal platform configuration. GPU Operator manages components; it is not on every inference request’s data path. Driver, CUDA, framework and hardware compatibility must be validated together.
Hover, focus or tap a component to inspect it. Motion adapts automatically to connection, device and accessibility signals; the component key remains readable without JavaScript.
__global__ void add(const float* a, const float* b,
float* out, size_t n) {
const size_t stride = size_t(blockDim.x) * gridDim.x;
for (size_t i = size_t(blockIdx.x) * blockDim.x + threadIdx.x;
i < n; i += stride) {
out[i] = a[i] + b[i];
}
}Make correctness a gate, not an afterthought
A candidate can look fast while mishandling an awkward input. The next chapter therefore starts with the cases a convenient benchmark would omit.
Compare against an independent reference across empty, tiny, large and non-tile-aligned inputs. Sizes such as 1, 31, 32, 33 and 257 expose boundary assumptions that powers of two can hide. Include supported strides, layouts and aliasing rules. If the contract excludes overlapping buffers or noncontiguous tensors, reject them rather than silently miscomputing.
Use memory and race checking where applicable. For reductions and matrix operations, choose absolute and relative error tolerances from the data range, accumulation precision and downstream application. Reordering floating-point additions can change results. Include cancellation, large magnitudes and exceptional values when the application supports them; do not loosen a failed tolerance merely to preserve a performance result.
Estimate the physical roofline
You now ask what the hardware could achieve in an ideal case. Roofline analysis turns that question into a bound, not a promised runtime.
Two ceilings bound an ideal kernel: arithmetic throughput and memory bandwidth. Count floating-point operations and bytes crossing the memory level under study. With decimal units, 0.4 TFLOP at 80 TFLOP/s takes at least 5 ms; 8 GB over 2 TB/s takes at least 4 ms. Perfect overlap gives a lower bound of max(5, 4) = 5 ms, not 9 ms.
Change the traffic to 20 GB below and the memory bound becomes 10 ms. This does not predict achieved time: dependencies, imperfect occupancy, cache behavior, synchronization and instruction mix can move execution further from either ceiling. Peak tensor throughput is not the correct ceiling for an arbitrary scalar FP32 kernel.
Count work
Use the exact operation and precision; distinguish dense from sparse throughput.
Count bytes
Include reads and writes at HBM, not only the size of the input tensor.
Compare bounds
The slower ideal resource establishes the roofline time lower bound.
Measure reality
Compare an unprofiled timing run with the same workload and correctness gate.
All rates in this model are editable ceilings, not specifications for a named GPU.
Work (TFLOP): 0.4; HBM traffic (decimal GB): 8; Compute ceiling (TFLOP/s): 80; Bandwidth ceiling (decimal TB/s): 2
max(work / compute, bytes / bandwidth), converted to ms. Requires perfect overlap and achievable ceilings for the exact precision and operation; excludes launch, transfers, synchronization and inefficiency.
Change the example assumptions
The worked result is readable without JavaScript. Inputs become available when the local WASM model loads; constrained connections and devices retain the static example.
Reduce traffic without creating a different bottleneck
The launch decision now depends on saved memory traffic without sacrificing correctness or moving the bottleneck.
Coalescing means neighboring active lanes access addresses that memory hardware can service efficiently. Tiling reuses a block of data from shared memory or registers instead of repeatedly fetching it from HBM. Fusion can avoid writing an intermediate tensor and launching another kernel. These are mechanisms to test, not a checklist of automatic wins.
Larger tiles consume registers and shared memory, potentially reducing resident work or causing register spills. Barriers must be reached by the required participating threads. A masked edge load still needs correct values for the following computation. Fusion may also increase live state and change rounding, so retest both resource usage and numerical behavior.
Conceptual Blackwell execution and memory map
Off-chip DRAM → memory controllers → caches → registers / execution → stores. Instruction issue and asynchronous copies coordinate different paths.
Large off-chip storage
On-chip reuse and staging
Inside one representative streaming multiprocessor
Device-wide work and peers
Stacked DRAM stores large device-resident arrays. DRAM cells need refresh; capacity and sustained bandwidth are different limits. HBM is not the register file.
Read every component and connection
- HBM3e DRAM · Weights / KV / arrays
- Stacked DRAM stores large device-resident arrays. DRAM cells need refresh; capacity and sustained bandwidth are different limits. HBM is not the register file.
- Memory controllers · Channels and requests
- Controllers organize reads and writes to memory channels. Access patterns, contention and the memory technology influence service time; bandwidth is not zero latency.
- L2 cache · Device-wide reuse
- L2 can satisfy repeated requests without another HBM access. Its capacity and residency behavior affect traffic; a cache hit is not a new DRAM transfer.
- L1 / shared memory · Caching / explicit tiles
- B200 combines L1, texture and shared-memory resources. Shared memory is software-managed block storage with synchronization rules; it is not an automatic replacement for registers.
- Register file · Thread operands
- The B200 tuning guide specifies 64K 32-bit registers per SM. Threads use registers for live values; spills can create device-memory traffic. Registers are not off-chip DRAM.
- ALU pipelines · Integer / floating point
- Execution pipelines perform supported arithmetic and logic on operands. CMOS gates underlie those circuits. Floating-point operations include more work than the integer full adder shown.
- Tensor cores · Matrix operations
- Specialized matrix instructions use supported operand formats and accumulation paths. Tensor throughput is not scalar ALU throughput, and not every kernel can use tensor cores.
- Load / store units · Addresses and movement
- Load/store machinery forms and services memory operations. Coalescing groups useful lane accesses; dependencies prevent a consumer from using a value before it is ready.
- Warp schedulers · Ready instruction issue
- Schedulers issue eligible warp instructions subject to dependencies and resource availability. Other ready warps can hide a wait; occupancy alone does not prove throughput.
- Instruction path · Fetch / decode / issue
- Compiled machine instructions reach the SM instruction machinery. PTX is a virtual ISA; a compatible cubin or driver compilation supplies hardware-executable code.
- Async copy / TMA · Tile movement
- Supported asynchronous transfer paths can stage tiles while computation proceeds. Barriers and producer/consumer ordering still apply; overlap is not permission to read unfinished data.
- GPU front end · Submitted work
- Device work submission and scheduling machinery distribute kernel work. Block resource requirements influence residency. Kubernetes does not choose a warp or allocate an SM register.
- NVLink interface · Peer devices
- Peer access and collectives move data between compatible GPUs. The application/runtime manages distributed work; aggregate device memory is not one automatically shared allocation.
- GPU front end → Instruction path: kernel work
- Instruction path → Warp schedulers: decoded instructions
- Warp schedulers → Load / store units: memory instruction
- HBM3e DRAM → Memory controllers: DRAM service
- Memory controllers → L2 cache: cache-line traffic
- L2 cache → L1 / shared memory: cache / tile path
- L1 / shared memory → Load / store units: load path
- Load / store units → Register file: operand load
- Register file → ALU pipelines: ALU operands
- ALU pipelines → Register file: result
- Register file → Load / store units: store
- L2 cache → Async copy / TMA: async tile copy
- Async copy / TMA → L1 / shared memory: staging
- L1 / shared memory → Tensor cores: supported matrix operands
- L2 cache → NVLink interface: peer traffic
This is a functional map, not a floorplan or cycle-accurate simulator. One representative SM is expanded; it is not the GPU’s SM count. Cache bypass, asynchronous copies, distributed shared memory and specialized tensor accumulator paths mean not every operation follows every arrow.
Hover, focus or tap a component to inspect it. Motion adapts automatically to connection, device and accessibility signals; the component key remains readable without JavaScript.
Time the device and the application separately
The timer finally enters the story, but there are two clocks to respect: completed GPU work and the request a user actually experiences.
Use device events on the appropriate stream to time GPU work, with explicit completion before reading results. Measure the full request with a host clock separately. An asynchronous launch returning quickly does not mean the GPU finished quickly. Report whether allocation, transfers, compilation and warmup are included in each number.
Run enough repetitions to observe a distribution rather than a single best sample. Compare equivalent inputs, accuracy, batch sizes, clocks, power settings and thermal conditions. Alternate baseline and candidate runs when drift matters. Profile representative kernels for explanation, then use unprofiled runs for final timing: counter collection and replay can change execution.
| Measurement | What it answers | Common mistake |
|---|---|---|
| Device elapsed time | Did this operation get faster on this stream? | Timing only the host launch |
| Request p50 / p95 / p99 | Did users receive faster complete responses? | Quoting only the fastest sample |
| Throughput at the latency target | Does more useful work fit the same service budget? | Changing concurrency or precision without disclosure |
| Accuracy and failure rate | Does the candidate still satisfy the contract? | Discarding difficult inputs or failed runs |
Use Nsight to explain the limit
Nsight Compute provides the evidence for the next hypothesis. You use CUDA profiling to explain a limit rather than collect a flattering utilization chart.
Start with a small metric set and a targeted kernel range. Inspect memory throughput, achieved compute throughput, register use, launch geometry and scheduler behavior. A low occupancy number alone is not proof of a slow kernel; sufficient eligible work and actual pipeline use matter. A cache-friendly kernel can exceed an HBM-only bandwidth estimate without violating physics.
Nsight Compute may replay kernels or the whole application to collect incompatible counters. Read its replay and overhead guidance before profiling a workload with host interactions, changing inputs or concurrency dependencies. Keep the exact profiler version, GPU, driver, binary and command with the result.
Ship only a maintainable application win
The team returns to the full request: the 17.5% application reduction must earn its maintenance cost.
Benchmark the full shape distribution, not only the shape used to design the tile. Keep a tested library path for unsupported hardware or inputs when the application contract requires it. Bound compilation and specialization costs; a fast warm kernel can still make cold starts worse.
The decision record should contain the before/after distributions, accuracy gate, memory footprint, hardware and software versions, supported inputs and rollback artifact. Only then decide whether the measured gain justifies owning custom device code. For a deeper implementation example, continue to the tiled matrix-multiplication and fused-softmax guides.
Questions behind the decision
When is a custom CUDA kernel worth the engineering cost?
Start with a repeatable application bottleneck and the fraction of end-to-end time it consumes. Compare against tuned libraries and include numerical testing, hardware support and future maintenance in the decision.
Why can a 2× faster kernel barely change inference latency?
Only the time spent in that kernel is reduced. Queueing, data movement and other operators still contribute, so report the complete request boundary as well as device time.