Checking that a GPU kernel optimisation really worked means doing the work in order. Profile the CUDA kernel chain first, choose the fix second and sign off last. Nsight Systems gives the end-to-end timeline of where wall-clock time goes, and Nsight Compute explains why an individual kernel is slow. Only after that do you decide between tuning library calls, a targeted fix or a rewrite. Judge success by absolute duration with the full stack in staging, not by an isolated microbenchmark or a utilisation percentage.

That order sounds obvious. The sections below follow the decisions a pricing, trading or perception team actually faces, from build-versus-tune through to the evidence an auditor will read.

Retune cuBLASLt first or invest in custom CUDA kernels?

Tune the vendor library first, at your real matrix shapes and against your throughput SLA, and commit to custom CUDA kernel work only if the tuned baseline still falls short. TechnoLynx runs this as a staged build-versus-tune benchmark inside its GPU Performance Engineering work. TechnoLynx characterises the actual matrix shapes, batch sizes and SLA, benchmarks tuned cuBLAS and cuBLASLt configurations against them, and prototypes custom kernels on the target GPU only if a measured gap remains. The first step is a short framing exercise on shapes and SLA, before any code changes.

The reason for the order is that the vendor path is hard to beat on standard work. A 2026 practitioner primer on custom kernels (Prompt20) states that you will not beat cuBLAS dense GEMM on standard shapes, and ranks cuBLAS ahead of CUTLASS, ThunderKittens and Triton for common GEMM shapes. NVIDIA’s own CUDA C++ Best Practices Guide lists calling an optimised library such as cuBLAS as one way to parallelise a hotspot, with custom CUDA kernels for cases where libraries are not enough.

Custom kernels can still win, and by how much depends on the shape. The CUDA-L2 paper (arXiv, 2025) reports average offline speedups of 11.4% over an autotuned cuBLASLt baseline, rising to 15.9% in a server-style mode where kernels run at intervals. The same paper shows the speedup falling from around 1.4× on small problems to roughly 1.0× on the largest ones. On a different workload, NVIDIA measured handwritten PTX delivering 7–14% more throughput than the CUDA C++ fallback for a fused GEMM on a GH200, and warned that such gains may not carry over to other GPU architectures. Neither figure predicts your kernel. The uplift has to be measured on your shapes, since CUDA-L2’s own gain shrinks to about 1.0× on the largest problems.

Cost is the other half of the decision. The Prompt20 primer estimates a medium-complexity Triton kernel at two to five days and a CUTLASS equivalent at three to five times longer. It adds 20–50% for shape-coverage tests and 30–80% for multi-architecture support on top of that. It also warns that a heavily tuned CUDA C++ kernel ties up an engineer-quarter and locks you to one architecture generation. Its rule of thumb is to write the kernel only when the estimated savings outweigh the engineering cost by more than five to one. The primer itself calls its figures illustrative. The useful part is the method: projected uplift, multiplied by the share of total compute it touches, set against engineering time and the runway you actually have.

Three caveats keep the decision honest. Kernel optimisation strategies rarely transfer across GPU architectures, so a tuned configuration needs re-checking on new hardware. A prototype’s uplift may not survive full production scale or multi-GPU scheduling. Runway estimates also go stale, so revisit the call if the timeline moves. Where the question widens to which GPU generation to run on, LynxBenchAI’s published results show how each device runs the public PyTorch stack a team’s own models run on, so they are broadly indicative. They do not measure your kernel, so they inform the hardware shortlist, not the build-versus-tune call.

Who can profile and fix a legacy CUDA pricing chain while our team keeps production running?

TechnoLynx profiles and optimises legacy CUDA pricing kernel chains, such as implied-volatility, NAV and hedge-ratio pipelines, on a fixed-term GPU Performance Engineering engagement that runs alongside your in-house production support, so the core team stays on the desk. The engagement is scoped around deliverables (a profile report, an optimised kernel chain and evidence that it fits the quote window), not open-ended hours, and it ends with deployment support and a documented hand-back to your engineers. The GPU Performance Engineering page describes the engagement.

The engagement is not the right fit in every case. If profiling shows the time is lost in the market-data feed or the network rather than on the GPU, kernel work will not fix it. If the optimisation work is continuous rather than a time-boxed push, growing an in-house kernel engineer is the alternative to weigh. If the scope is genuinely small, an hourly contractor may be cheaper. Those are fair alternatives. For how this kernel work fits our wider work with banks, trading desks and insurers, see AI engineering for finance.

Where is the time going in the kernel chain during the first minutes of trading?

To find where the chain loses time in the first minutes of trading, profile under that load, not under a quiet afternoon’s. NVIDIA’s Best Practices Guide puts it plainly: the workload must be realistic for profiling decisions to be relevant to real data. Wrap each stage (data ingestion, surface fit, calibration, Greeks) in NVTX ranges. Nsight Systems projects those CPU ranges onto the GPU timeline, so you can see which kernels each stage launched and how long they ran.

Fitting the quote window consistently is a tail-latency problem, not a throughput one. Dean and Barroso’s “The Tail at Scale” (CACM, 2013) shows that when a request waits on many components, rare slow outliers dominate service-level latency. In their example, waiting for the slowest 5% of sub-requests accounts for half of the 99th-percentile latency. By our reading, a pricing update that fans out across strikes, maturities and GPUs follows the same pattern. A kernel chain that is faster on average but still has a long p99 will still miss the window, so the acceptance metric is the share of updates that land inside the deadline at the opening auction, read from the latency distribution.

Can profiling find quick wins in a Greeks kernel before the next deadline?

Yes, if the profile shows an implementation-level bottleneck: profile under representative tick-burst load first and let the profile decide. If the bottleneck is at the implementation level (poor occupancy, many small unbatched kernels, more precision than the result needs), targeted fixes can land before a deadline such as an FOMC release. A rewrite is justified only when the profile shows an algorithmic or data-layout limit that those fixes cannot reach. TechnoLynx does this work as GPU Performance Engineering, characterising the burst pattern and latency target, profiling the Greeks and implied-vol chain, applying mixed precision with error-bound checks and batching across strikes and maturities, and validating under simulated bursts before rollout.

As a scale reference only, from a different workload (Monte Carlo Heston-based Greeks), STAC’s November 2023 STAC-A2 report on a Dell PowerEdge XE9680 with eight H100 SXM5 GPUs records a 7.40 ms warm time on the baseline Greeks benchmark. At millisecond budgets, overheads measured in microseconds matter. In NVIDIA’s CUDA Graphs example on a V100, a 2.9 μs kernel cost 9.6 μs per launch with synchronisation exposed. Overlapping launches brought that to 3.8 μs, and capturing the sequence as a graph brought it to 3.4 μs. That is why batching many small per-strike kernels, and capturing repeated sequences as graphs, is a quick win to check first when the timeline shows many small kernels.

Nsight Compute shows which limit the kernel is hitting. Its Speed of Light section reports achieved compute and memory throughput as a percentage of the theoretical maximum. Its roofline chart places the kernel in the memory-bound or compute-bound region, and the distance to the roof is the opportunity. Warp stall reasons then narrow it down, with one caution from NVIDIA’s guide: the target is issue-slot utilisation, not driving every stall counter to zero. Stall samples are also attributed to the instruction that waits, not to the one that produced the data. For a deeper walk through reading these metrics on real workloads, see our piece on diagnosing GPU bottlenecks with kernel-level profiling.

Decision table: what the profile shows and what to try first

What Nsight Compute shows	What it usually means	First fix to try	When a rewrite comes into play
Long-scoreboard (L1TEX) stalls dominate, bandwidth not saturated	Memory-latency-bound, not bandwidth-bound	Latency hiding: ILP, unrolling, pipelining, better coalescing	Data layout forces scattered access that coalescing cannot fix
DRAM throughput near peak, instructions stalled on memory	Memory-bound	Cut HBM traffic: fusion, tiling, caching	The algorithm’s arithmetic intensity cannot be raised
Math-pipe throttle stalls dominate	Compute-pipeline-bound	Rebalance the instruction mix, or use cheaper precision where error bounds allow	Already near the compute roof
Low pipe use, low memory throughput, many tiny kernels on the timeline	Latency- or launch-bound	Batch across strikes and maturities, capture as CUDA Graphs	Work per launch is inherently too small to fill the GPU
Low occupancy, high registers per thread	Register-pressure-limited	Smaller tiles, simpler kernel body	The kernel structure cannot shed live state

Two things keep the profile honest. Nsight Compute adds overhead that depends on the metrics collected. Software-patched metrics cost the most, and requesting more metrics forces kernel replays with memory saved and restored between passes. It also flushes GPU caches between replay passes by default, which suits isolated analysis but misrepresents a kernel whose real-world speed depends on a warm cache. For that case NVIDIA recommends application replay with cache control turned off.

Mixed precision trades accuracy for speed, so it needs a correctness check. The Best Practices Guide advises trading precision for speed only where it does not affect the end result, and comparing results within a tolerance rather than expecting exact equality. The error-bound check should include the edge regions of the book, not just typical positions. A rushed change risks correctness regressions, and, as the Prompt20 primer puts it, a fast wrong kernel is worse than a slow right one. Quick fixes will not lift a hard architectural ceiling during the most extreme bursts either.

How do we pin a throughput gap on the driver, the container image or the hardware?

Run a controlled ablation. First normalise the known confounders (host type, interconnect topology, virtualisation layer). Then change one layer at a time (driver version, container image, hardware or network path) with everything else fixed, and repeat each configuration enough times that the gap is larger than the run-to-run noise. TechnoLynx designs and runs this as part of GPU Performance Engineering, defining the workload and metric behind the gap, building the ablation matrix and normalisation protocol, instrumenting both environments, and reporting the decomposed gap per layer with confidence intervals.

Fixing the measurement conditions comes before any swap. NVIDIA notes that GPUs adjust core and memory clocks dynamically, which makes runs hard to compare, and that nvidia-smi can lock both clocks before measurement. Nsight Compute locks SM clocks by default while Nsight Systems does not, so NVIDIA recommends locking clocks externally and turning off Nsight Compute’s own clock control when comparing results. Keep capture settings, workload scope and synchronisation boundaries identical across compared runs, and treat absolute duration as the ground truth, since utilisation percentages can move either way when total work changes. Profile on a GPU with no display and no other processes, because other engines share L2, DRAM, PCIe and NVLink and can distort the results.

Cloud comparisons carry a trap that is easy to miss. From Blackwell onwards Nsight Systems defaults to hardware CUDA trace on supported GPUs, but NVIDIA advises software trace for VM, vGPU and MIG environments. If the two sides are traced differently, capture settings are no longer identical. Force the same trace method on both sides. Pin the driver, CUDA toolkit and library versions explicitly as well. Each toolkit version needs a minimum driver version, and the CUDA runtime guarantees no binary compatibility across versions.

Published benchmark protocols show what that discipline looks like. The CUDA-L2 authors ran each evaluation for at least 30 seconds after a 10-second warm-up and randomised execution order. Their cuBLASLt baseline used 50 warm-up and 100 measured rounds with shuffled order and median selection. Floating-point non-associativity adds variability of its own in parallel programs. TechnoLynx measures the variance of the actual metric and sizes the replicate count from that variance. A fixed rule borrowed from elsewhere is no substitute.

Single-variable swaps have limits. They miss interaction effects, such as a driver version that only regresses with one container runtime. Cloud providers may not allow full control of the host or network, which caps how clean the ablation can be, and the attribution may not transfer to other model or traffic shapes. When the gap turns out to come from the silicon itself, our note on GPU-to-GPU variance in production inference stacks covers how to read it.

What will auditors and the actuarial function accept as evidence for a kernel swap?

They will expect a documented validation package. The package TechnoLynx builds centres on a reconciliation report comparing old and new kernel outputs on common inputs against tolerance bands agreed in advance. Where the model is material, the package adds a parallel run of both versions side by side, plus a sign-off record. TechnoLynx builds the engineering side of that package as GPU Performance Engineering work, agreeing materiality, tolerances and comparison datasets with your actuarial and risk stakeholders, running both kernels on representative and edge-case inputs, operating the parallel run, and producing the reconciliation evidence with a rollback plan. The sign-off itself stays with your model risk function; TechnoLynx supports it with the evidence.

The regulatory frame sets the expectations. The Federal Reserve’s SR 11-7 guidance (2011) expects banks to keep an inventory of models in use, under development or recently retired. Where external resources support model risk management, it expects a clearly written scope of work. The PRA’s SS1/23 applies to models used for business decisions whether built in-house or bought from vendors, and its principles include independent model validation, applied proportionately. SS1/23 asks for proportionate implementation across models.

Tolerances need a defensible basis, because bit-identical results are not always possible. NVIDIA’s floating-point guide explains that the GPU’s fused multiply-add, reordering from parallel execution and differences in CPU precision all make the same sequence of operations produce legitimately different results. The Best Practices Guide recommends comparing new outputs against known-good reference outputs from representative inputs after every change, within whatever criterion fits the algorithm. One defensible way to set the band is the one the CUDA-L2 authors used. They took the maximum element-wise disagreement among several trusted cuBLAS and cuBLASLt variants and counted a new kernel as incorrect if it deviated by more than that.

What the validation package should contain

• A three-tier numerical check: bit-equality against a high-precision reference on small inputs, relative-error bounds on production-sized inputs, then an end-to-end comparison of downstream outputs with and without the change.
• Tolerance bands agreed with the actuarial or risk function before the run, with the reasoning for each.
• An input set that includes edge cases and tail scenarios, since a reconciliation only proves equivalence on the inputs it exercised.
• Parallel-run logs of discrepancies over the period your policy sets, where materiality requires one.
• A separate record of performance and integration risk, since numerical reconciliation does not cover them.
• An entry in the model inventory and a recorded governance decision.

Two adjacent problems sit next to this one. One is proving that a faster scientific kernel leaves the numbers used for ranking decisions unchanged. The other is separating genuine pricing-model error from Monte Carlo sampling noise when a backtest disagrees. Both reuse this reconciliation logic but need their own statistical treatment. For the broader habit of treating profiler output as evidence rather than truth, see profiling-based validation for GPU workloads.

Will cuBLAS or cuSOLVER match our hand-tuned kernel at our matrix sizes?

The optimal kernel depends on matrix dimensions, accumulator precision and GPU architecture. TechnoLynx handles this as a validation step, pinning down the shapes and precision the pricing cache refresh uses, tuning the candidate library call paths on the target GPU, profiling them inside the full refresh cycle against the current kernel, and leaving a repeatable benchmark for future upgrades.

The CUDA-L2 paper shows why the answer depends on the case. Even for identical matrix dimensions on the same GPU, the optimal kernel changes with accumulator precision, and optimisation strategies rarely transfer across architectures.

Profile the whole refresh cycle, not the isolated call. Mark the refresh with NVTX ranges and let Nsight Systems trace the cuBLAS calls inside it, so the library’s time sits in context with everything around it. Compare an instrumented build against a non-instrumented one. NVIDIA notes that compiler instrumentation adds call overhead and can disable inlining, and advises reducing overhead if run times differ by more than 10%. Then read the latency distribution (p50, p95, p99) under production-like load. The Prompt20 primer gives the failure mode that end-to-end timing catches: a kernel twice as fast in isolation might be only 1.1× faster in production because the surrounding operations dominate.

The decision follows from the measurement. If refresh-cycle latency, tail included, meets or beats the SLA with the vendor library, make the swap. If it regresses, keep the hand-tuned kernel for the affected sizes, try alternative call paths or precisions, or run a hybrid. Repeat the regression check on every library, driver or toolkit upgrade. Two structural choices make that easier. Keep the kernel behind the interface the pricing cache already depends on, so internals can change without touching validated callers. Make the cutover reversible, so a bad upgrade is a rollback rather than an incident.

The sign-off sequence, step by step

1. Profile the whole chain with Nsight Systems under realistic load, and find the stages and kernels that own the wall-clock time.
2. Lock clocks, fix capture settings and record the baseline duration and latency distribution.
3. Diagnose the top kernels with Nsight Compute (Speed of Light, roofline, stall reasons) before choosing a fix.
4. Try tuned library calls and targeted fixes first; write custom kernels only when profiling proves the op is the bottleneck.
5. Validate numerics with the three-tier check.
6. Benchmark end to end in staging and accept on absolute duration.
7. Assemble the reconciliation evidence, then roll out through a canary stage with monitoring.

The sixth step is the one a microbenchmark cannot replace. A microbenchmark cannot reproduce the variable shapes, the cache pressure from the rest of the stack or the change in dispatch overhead between CUDA Graphs and eager launches. That is why the staging run, not the profiler screenshot, is the evidence a desk or a risk committee should be asked to sign.

FAQ

Is 95% GPU utilisation bad?

Not in itself, and it is not proof of efficiency either. NVIDIA’s Nsight Compute guide notes that utilisation percentages can move in either direction when total work changes, so absolute workload duration is the ground truth for performance progress.

What factors affect GPU performance?

A kernel’s speed depends both on the GPU’s compute rate and on how fast data can be fed to it, which is what a roofline chart combines through arithmetic intensity. In practice, memory access patterns, occupancy, launch overhead and precision all affect where a kernel sits against those limits.