GPU kernel bottlenecks fall into three main categories: launch and host-side overhead, memory bandwidth and arithmetic throughput. The way to find out which one applies is to profile the kernel on the device you will actually deploy on. Measure its throughput as a percentage of that card’s peak, place it on a roofline, and let the result guide what to optimise. A datasheet’s peak FLOPS or a best-practice rule from a blog post should not make that decision. The same goes for an edge module and a per-frame latency budget: you settle it by measuring on that module under its worst operating conditions, not on reference hardware.

Kernel duration measured on your own device and workload is the ground truth; vendor peaks are ceilings, not predictions. NVIDIA’s Nsight Compute guide makes the same point in its own terms: utilisation percentages can move either way when total work changes, so duration is the measure of real progress.

Will the edge GPU actually hit our per-frame latency budget?

The measurement has to cover the worst conditions the device will face, and it has to happen before procurement or a fleet rollout commits. Figures from a different GPU are no substitute, since NVIDIA advises re-evaluating the dominant limiter on each new architecture.

TechnoLynx does this work as a production AI reliability engagement. The first step is to set the latency target and the worst-case operating envelope the hardware has to meet. TechnoLynx then ports and optimises the model on the actual module, profiles it under realistic and worst-case load, and confirms through a pilot that it meets the target with margin before any fleet-wide rollout.

Where the time goes once camera I/O and preprocessing are added

Inference is not always the largest stage. One analysis of DNN serving overheads for computer vision looked at inference servers. It found that non-inference work dominates request latency for most models under 5 GFLOPs, and still takes 16 to 49% of latency for models above 10 GFLOPs. In the same study, the preprocessing share grew with image size, reaching up to 56% (CPU preprocessing) and 49% (GPU preprocessing) on medium images. With large images, end-to-end throughput for ViT-base fell to 19.5% of what ViT inference alone could reach. On Jetson specifically, a 2025 arXiv study of concurrent vision workloads found that CPU-side events such as thread scheduling and context switching often became the bottleneck and held back GPU performance.

That is why you need a per-stage budget covering decode, preprocess, inference, NMS, postprocess and transfer: it shows where the headroom or the deficit sits against the target. NVIDIA’s TensorRT documentation describes a workflow that fits. Start with the built-in profiler or trtexec --dumpProfile to find the most expensive engine layers, then move to Nsight Systems to see why those layers are expensive at the kernel and stream level. The same documentation advises comparing the reported Enqueue Time with GPU Compute Time on the deployment system, because kernel-launch overhead varies with the CPU, driver, operating system and capture mode. It also suggests checking in Nsight Systems whether host-to-device or device-to-host copies take longer than the GPU execution itself, which tells you whether PCIe transfer is the limit.

Thermal and power limits that bench runs miss

NVIDIA’s Jetson Linux guide says Orin NX and Orin Nano modules have average and instantaneous power limits on the input rail. A module that exceeds its configured limit throttles the GPU and CPU clocks in hardware. For discrete cards, the TensorRT documentation says thermal throttling starts at a device-specific slowdown temperature, which nvidia-smi -q -d TEMPERATURE reports. During a run, nvidia-smi dmon shows the clock dropping as the temperature approaches that threshold. The same guide warns that on GPUs prone to power throttling, such as the T4, poor cooling can settle the clock at a lower frequency even before thermal throttling begins.

The only way to get the sustained figure is to run the module in that envelope long enough to reach thermal equilibrium. Salt exposure and weeks of unattended operation call for extended field testing, not a short bench run. Whether the sensors and enclosures themselves survive that environment is a separate validation question with its own method.

Published numbers, ports and pilots

Published results are a reasonable first filter between candidate devices. LynxBenchAI’s published results show how each device it covers runs the public PyTorch stack your own models run on, so they give a broad indication for custom PyTorch models. They do not run or measure your own workload or application.

Comparing devices from different vendors raises a tooling problem. Vendor profilers are siloed: Nsight will not profile an AMD GPU, and rocprof will not work on NVIDIA. Cross-platform tools such as HPCToolkit can trace GPU work on NVIDIA, AMD and Intel GPUs. For the deepest NVIDIA-specific metrics, such as warp stall reasons, developers still rely on Nsight Compute. A fair capacity comparison therefore uses end-to-end duration as the common measure and keeps the vendor-specific counters for diagnosing each side.

If the target hardware, measured end-to-end and per stage under worst-case conditions, clears the target with margin, pilot it on one line, site or point of presence, then scale with a rollback plan in place. If it falls short, measure the size of the gap and decide whether optimisation (quantisation, kernel re-porting, batching) can close it or whether you need a higher hardware tier. Either way, a single-unit pilot reduces the risk of a fleet-wide commitment built on bench numbers.

Is my kernel limited by launch overhead, memory bandwidth or compute?

You classify a kernel by comparing what it achieves against the device’s own ceilings. Nsight Compute’s GPU Speed Of Light section reports compute and memory throughput as a percentage of the theoretical maximum, and its roofline chart places the kernel in either the memory-bound or the compute-bound region. A GPU profiling guide for ML kernels describes three regimes. In a memory-bound kernel the compute units sit idle waiting on VRAM. In a compute-bound kernel the memory bus waits on the maths. In a latency-bound kernel too few threads are in flight to hide memory access, even though the bus is not full. AMD’s HIP documentation adds a third category, overhead-bound, which covers host-side scheduling, kernel launch overhead and small array operations. Horace He’s “Making Deep Learning Go Brrrr” splits time the same way: compute, memory transfer, and everything else counted as overhead.

The roofline gives you a visual way to tell memory-bound and compute-bound kernels apart. Williams, Waterman and Patterson’s 2008 roofline paper defines attainable performance as the lower of two limits: peak floating-point performance, or peak memory bandwidth multiplied by operational intensity (the FLOPs performed per byte of DRAM traffic). The ridge point, where the two limits meet, is the minimum intensity a kernel needs to reach the machine’s peak. Kernels to the left of the ridge are memory-bound and kernels to the right are compute-bound.

A GPU/CPU bottleneck needs a different test, because overhead does not grow with problem size while compute and memory do. Double the batch size: if runtime rises by only about 10%, the kernel is probably overhead-bound. In trtexec, add --noCudaGraph and check whether Enqueue Time is close to, or longer than, GPU Compute Time. If it is, the workload is enqueue-bound.

Regime diagnosis table

Regime	What the profiler shows	What tends to help	What will not help
Overhead / launch-bound	Runtime barely grows when the batch doubles; Enqueue Time at or above GPU Compute Time; gaps between kernels on the GPU timeline	CUDA Graphs, operator fusion, tracing or compiling away Python dispatch	Kernel-level tuning of maths or memory access
Memory-bandwidth-bound	DRAM throughput above roughly 80% in Speed Of Light; kernel left of the ridge point	Tiling and data reuse, coalescing, operator fusion	A GPU with more FLOPS
Memory-latency-bound	L1TEX scoreboard stalls dominate while bandwidth is not saturated	Latency hiding: ILP, unrolling, pipelining	Bandwidth-focused rewrites
Compute-bound	SM throughput above roughly 80%; kernel right of the ridge; math-pipe throttle stalls dominate	Tensor Cores or FP16, simpler algorithms, a better instruction mix	Rewriting host code to cut overhead

The 80% thresholds come from the Vulkan guide. Treat them as a starting point for stall analysis, not a verdict.

What occupancy and memory bandwidth is my kernel actually reaching?

Nsight Compute reports theoretical occupancy per SM. It defines occupancy as the ratio of active warps per multiprocessor to the maximum possible. Occupancy tells you how fully the hardware is loaded, not how efficiently it is used. AMD’s HIP guide points out that high-performance MFMA GEMMs often run at low occupancy, because a few warps are enough to saturate the pipelines. Low occupancy always weakens latency hiding, but past a certain point extra occupancy stops paying off. NVIDIA’s CUDA Best Practices Guide notes that raising occupancy from 66% to 100% generally does not produce a similar gain in performance. Simon Boehm’s CUDA matmul worklog shows the trap in practice: a kernel at 66% occupancy was still slow, and the profiler’s warp-state sampling eventually traced the problem to shared-memory stalls.

The CUDA Best Practices Guide computes theoretical bandwidth from the specification; for a Tesla V100 it comes to 898 GB/s. Effective bandwidth is the bytes read plus the bytes written, divided by the measured time. When effective bandwidth sits far below theoretical, the guide makes raising it the main optimisation goal. Aligned accesses reached about 790 GB/s on the V100. The profiler reports requested and actual throughput separately, and the gap between them estimates how much bandwidth poor coalescing wastes. In Boehm’s worklog, coalescing global memory accesses raised throughput from 15 GB/s to 110 GB/s and took the kernel from about 300 to 2,000 GFLOPS.

Rules such as “maximise occupancy” or “avoid branches” only make sense in context, and that context comes from profiling.

Why is my kernel nowhere near the card’s advertised TFLOPS?

AMD’s HIP documentation gives the reasons applications reach only a fraction of theoretical peak: imperfect resource utilisation, inefficient memory access, control-flow divergence, synchronisation overhead and launch costs. Boehm’s naive matmul reached about 300 GFLOPS on an A6000 that is advertised at almost 30 TFLOPS.

An investigation published on thonking.ai explains that when a GPU hits its power limit, the voltage regulator throttles its voltage and clock. It concludes that H100 matmuls are mainly power-limited, and that the H100’s real advantage over the A100 has usually been closer to 2x than to the theoretical 3x. Input data affects power draw too. Zeros or NaNs flip fewer bits and draw less power, so NVIDIA’s TensorRT documentation recommends benchmarking with inputs that match the real value distribution.

To compare against peak for your own shapes, work out each kernel’s arithmetic intensity, read the matching ceiling from the device’s roofline, and judge the kernel against that ceiling rather than the headline FLOPS. The roofline paper found that the ridge point predicted performance better than clock rate or peak performance. An NVIDIA walkthrough of Nsight Compute roofline analysis shows that a better roofline position does not guarantee a faster kernel. Unrolling moved a V100 kernel from 7.39 FLOP/byte (just below the 7.5 balance point) to about 20, which made it clearly compute-bound, yet runtime got 10% worse, from 1.74 s to 1.92 s. A later change raised throughput from 2.5 to 2.9 TFLOP/s and halved the runtime, even though the kernel dropped back into the bandwidth-bound region.

Clocks, caches and warm-up decide whether two runs are comparable

NVIDIA’s TensorRT documentation warns that two runs of the same engine on the same GPU can differ by tens of percent if the clock floats in one run and is locked in the other, or if thermal throttling starts partway through. Locked clocks give more deterministic tactic selection and steadier measurements. With floating clocks, measurements taken over short windows such as 20 ms are unstable.

The two Nsight tools handle this differently. Nsight Compute locks SM clocks by default and Nsight Systems does not. So when you compare results across the two, the Nsight Compute guide recommends locking clocks externally with nvidia-smi and passing --clock-control none to ncu. Nsight Compute also flushes all GPU caches between replay passes. For cache-sensitive workloads the guide suggests --replay-mode application --cache-control none. The Vulkan profiling guide adds three rules: never profile a Debug build, run 10 to 20 warm-up inferences before timing, and lock clocks with nvidia-smi -lgc. Other engines and processes can contaminate DRAM results, so profile on a GPU with no active display. Workloads shorter than about 20 µs may not reach a steady state, which makes multi-pass metrics unreliable. Finally, keep capture settings identical across runs unless you are deliberately changing one variable.

How should I benchmark cuBLAS against CUTLASS or a hand-tuned kernel on our shapes?

Benchmark on your production shapes and your deployment GPU, using the measurement hygiene above. A library’s reputation on square matrices tells you little about your aspect ratios. The ISAAC paper (SC17) found that vendor libraries tuned for certain inputs, such as square matrices, generally lose their optimal performance across the range of problems seen in practice. In the same paper, cuBLAS reached over 90% of peak on Maxwell for large square matrices but saturated at 85% on the P100. Boehm’s autotuned kernel came within a few percent of cuBLAS at sizes 2048 and 4096 but fell well behind on smaller matrices, because cuBLAS contains hundreds of SGEMM implementations, not one.

Vendor APIs are not magically faster than standard ones. What they offer is access to lower-level features. Making use of that access, whether through a vendor API or your own hand-tuned kernel, takes real commitment, and it is not always worth it. ISAAC is an example: it got bounds checking down to 2% overhead from 15 to 20% only after moving code generation from CUDA C to PTX, where predication is exposed. Boehm reports that his first six kernels reached 80% of peak FLOPs in two weekends, and that it took four more weekends to reach 94%.

Procedure for a fair library-versus-custom comparison

1. Take the actual (M, N, K) shapes and data types from a production trace, not a representative square size.
2. Lock clocks and use the same cache-control settings for every candidate, as described in the clocks and caches section.
3. Warm up each kernel before timing.
4. Fill the inputs with values drawn from the real distribution. The thonking.ai investigation traced an apparent CUTLASS advantage over cuBLAS to the CUTLASS profiler initialising its inputs with integers by default.
5. Time on the device with CUDA events rather than host timers. For CUTLASS, the CUTLASS Profiler runs GEMM and convolution kernels at fixed shapes, with a set profiling duration or iteration count. Its CUDA-graph launch option helps when launch time is the bottleneck.
6. Compare duration first, then open the slower candidate in Nsight Compute to see which limiter it hits.
7. Repeat on every GPU model you deploy on. In Boehm’s tests, the best parameter set on an A6000 came out 6% worse than the autotuned optimum on an A100.

Nsight Compute or Nsight Systems: which one answers a kernel-level bottleneck question?

Nsight Compute answers the kernel-level question. It profiles individual kernels with instruction-level statistics, memory transactions, occupancy and stall reasons, and it maps metrics back to source lines or SASS. Nsight Systems answers the question that comes first: where time goes across the CPU-GPU boundary, with CPU threads, kernel launches and memory copies on a single timeline. NVIDIA’s TensorRT workflow starts with the built-in profiler and then moves to Nsight Systems for the kernel and stream level.

Nsight Compute may replay a kernel several times, because the hardware can collect only a limited number of counters per pass. The Jetson study describes Nsight Systems as a medium-overhead profiler whose overhead scales with the number of metrics collected, and it used lighter tools when measuring throughput so that profiling would not disturb the inference loop. One timeline trap: the CPU can look idle inside cudaStreamSynchronize() while the GPU is busy, so read the CUDA HW rows alongside the thread rows. On AMD hardware the equivalents are rocprofv3 for roofline data and Omniperf for kernel-level counter analysis. Choosing tools is only part of the job, since staffing and build configuration matter too; both are covered in our guide to finding and fixing GPU bottlenecks.

When does batching more work per launch pay off, and when does memory cancel it?

Reducing launches per unit of work pays off in the CPU-bound region and gives little once the GPU is saturated. A 2025 arXiv study of kernel batching with CUDA Graphs explains that launch cost stays roughly constant as GPUs get faster, because it comes from the OS, the driver, the API and CPU-GPU communication. When you capture a kernel sequence in a graph, the CPU pays the launch cost only once. The same study found that graphs with one or two nodes are unlikely to help. Very large graphs carry heavy creation overhead and can use a lot of memory, so the batch size has to balance memory use against performance.

Another arXiv study, this one of LLM inference, sets the boundary. Kernel fusion to cut launches helps only in the CPU-bound region, where the GPU is underused. Once kernel queuing dominates, fusing to reduce launches gives no benefit. The study also found that each application and system pair has a batch-size sweet spot where both CPU and GPU are well used, and it recommends operating there rather than aiming for maximum GPU saturation. Too much computation per kernel leaves the CPU idle and stretches the latency to batch completion. In practice, step the batch size up, and at each step compare Enqueue Time against GPU Compute Time, runtime growth and device memory. Stop at the batch size where runtime starts to scale with the batch, because from there the launch overhead is hidden or amortised by longer-running kernels.

How do I get per-function cycle counts on a Cortex-M7 without a trace probe?

Use the DWT cycle counter, which needs no trace probe. The Data Watchpoint and Trace unit on Cortex-M3, M4, M7 and M33 includes a 32-bit free-running counter, DWT->CYCCNT, that ticks once per core clock and needs no pins and no extra hardware. The MCU on Eclipse write-up outlines the steps. Set the TRCENA bit in DEMCR (0xE000EDFC), enable CYCCNTENA in the DWT control register (0xE0001000), reset CYCCNT (0xE0001004), call the function, then read the counter. This works without an active debug session, as long as your code initialises the unit itself. To convert cycles to time, divide by the clock frequency: Memfault’s example converts 38,298,322 cycles to about 23 ms at 168 MHz.

A few details decide whether you can trust the counts. The 32-bit counter wraps after 2^32 cycles, which at 168 MHz is about 25.6 s (4,294,967,296 ÷ 168,000,000), so time longer regions in pieces. Cache misses on the M7 add to execution time, so record the minimum, maximum and mean over many runs. That statistics code adds roughly 10 to 15 cycles of overhead per interrupt. The M0 and M0+ have no DWT unit, and SysTick is the fallback on those cores. If you want a per-function distribution without instrumenting each call, ITM PC sampling over SWO produces a sampled profile, but it needs a debugger that supports it. The principle is the same as on the GPU: count cycles on the real target and trust the count over the estimate.

A profile describes one device in one state

A profile is evidence about one device under one set of conditions. The Nsight Compute guide advises re-evaluating the dominant limiter on every new GPU architecture, because cache sizes, SM counts, register files, peak FLOP rates and memory bandwidth all shift which part of a workload limits performance. So the useful output of profiling is not a one-off verdict that a kernel is “memory-bound”. It is a method you can rerun: locked clocks, real shapes, real input values, the stage budget and the roofline position, repeated whenever the hardware, the model or the environment changes. The harder question is not which profiler to use. It is whether your test conditions resemble the hottest, busiest hour the deployed system will face.

FAQ

Is 95% GPU utilisation bad?

Not on its own, but it also proves little. The “GPU-Util” figure in nvidia-smi roughly measures the share of time any kernel is running, not how well that kernel uses the hardware. The Jetson study found that throughput can still be limited by kernel operations and CPU-side scheduling even at 99% GPU utilisation and full Tensor Core utilisation.

Does NVIDIA have a benchmark tool?

Yes. trtexec reports throughput, latency and a per-layer breakdown for TensorRT engines. The CUTLASS Profiler runs and profiles the GEMM and convolution kernels in the CUTLASS library. Nsight Systems traces the whole-system timeline, and Nsight Compute examines individual kernels in depth.