A ported CUDA kernel is correct when its output stays inside a tolerance budget you agreed before testing. That budget is measured against a reference and tied to what the application can absorb. A match to the last bit is not the test. Bit-exact parity is the wrong default. IEEE 754-compliant code can legitimately give different results on a CPU and a GPU, because the GPU uses fused multiply-add, parallel algorithms reorder operations, the CPU may compute at a higher precision than expected, and common math functions are not required to be correctly rounded. NVIDIA’s floating-point guide says directly that a difference between CPU and GPU results does not by itself show that the GPU result is wrong.

The work has three parts. First, set the budget from the application’s own quality target. Second, isolate divergence by stage, operator, library and precision. Third, lock down the reproducibility controls, so that a pass means the same thing on every run.

Is bit-for-bit identical output realistic after porting a CUDA kernel?

cuBLAS and TensorRT document bit-for-bit reproducibility under narrow conditions. For almost every port, the tolerance band has to be defined before the first comparison runs.

TechnoLynx starts this kind of validation by cataloguing every expected source of divergence between the research prototype and the port: precision mode, library and algorithm selection, reduction order and parallel decomposition. It then turns that list into a tolerance budget for each metric before any comparison is scored.

The narrow conditions are worth knowing exactly: cuBLAS routines from one toolkit version give bit-wise identical results on every run on GPUs with the same architecture and the same number of SMs. That guarantee does not hold across toolkit versions, with multiple CUDA streams active, or when fixed-point emulation is in use. The multi-stream non-determinism comes from how the library picks internal workspace. You can remove it with a separate workspace per stream, one cuBLAS handle per stream, cublasLtMatmul() with a workspace you own, or CUBLAS_WORKSPACE_CONFIG set to :16:8 or :4096:8. Problem shape matters too. By default, splitting a matmul along M or N can give results that are not bitwise identical to the unsplit problem, because the library may choose different kernels or reduction schemes for different shapes. A TensorRT engine, once built for a specific GPU, should give bit-accurate outputs across repeated runs. So run-to-run identity on one machine is achievable, but identity across machines is a separate property that has to be engineered.

What tolerance should you accept between the reference and the ported kernel?

No universal tolerance exists. A default such as 1e-3 says nothing about whether your downstream result survives. The right figure comes from how sensitive the final output is to small perturbations. Examples are a ranking that flips, a classification near a decision boundary, or a price that moves a basis point. Tasks that sit close to a decision boundary need tight bands, comparisons in matched FP32, and repeated runs. Tasks that are robust to numeric noise can carry bands wide enough to absorb TF32, FP16 and reduction-order variation.

A large reference set changes what an exact-match test can show. Research on calibration checks for safety-critical systems (2026) points out that, given enough data, an exact test will eventually reject any approximate model, however small and harmless its deviation. Their fix is to replace the exact null hypothesis with a tolerance region set by an engineering choice of ε_tol, following the logic of equivalence testing. In that work, typical ε_tol values run from 0.01 to 0.05 and encode a practical judgement of how much deviation is acceptable. Those values belong to calibration coverage, not kernel outputs. What carries over is the structure: an explicit, pre-declared slack that does not shrink to zero as the test set grows.

Keep the data split honest. The small lab set you used to tune the port is not the set that accepts it. Run acceptance on a separate, disclosed reference set, and have the harness track both sets over time. Keep in mind that round-off in recursive summation grows as the sum gets large relative to its summands. In NVIDIA’s Monte Carlo sample, the relative difference between single-precision GPU and double-precision CPU results improved from 7.6e-7 to 6e-8 after enlarging the shared-memory reduction array from 128 to 1024 elements. A budget derived on small test cases may therefore not hold at production input sizes.

Expected sources of divergence to budget for

Source	How it shows up	Expected or suspect
Summation and reduction order	Libraries and GPUs use different summation trees; FPRev (2024) found a 5-way tree on V100, 9-way on A100 and 17-way on H100	Expected; budget for it
Fused multiply-add	GPUs with hardware FMA use it by default, and both FMA and non-FMA results are correct under IEEE 754	Expected
Reduced-precision tensor-core paths	cuBLAS may use Tensor Cores automatically, including TF32 compute for FP32 inputs, unless pedantic modes are selected	Expected if the budget allows it
Library algorithm or tactic choice	TensorRT picks each layer’s precision for performance, which can cost accuracy	Expected, but check per layer
Math-function implementations	CUDA math library and host libraries such as glibc are independent and not guaranteed to be correctly rounded, so results often differ slightly	Expected
Aggressive compiler flags	In one NVIDIA-versus-AMD study, FP32 tests built with O3 and -ffast-math showed 13,877 discrepancies versus 45 at O0	Suspect until justified

Summation order is generally unspecified and unknown for library summation functions, which makes reproducible migration hard. FPRev treats a summation-based function as a black box. It builds inputs that exploit swamping, where a much smaller number added to a much larger one contributes nothing, and infers the order of summation from the output. The same study shows that NumPy, PyTorch and JAX each reduce in different orders. Order-independent summation algorithms give the same sum regardless of order, but they are inefficient and therefore not widely used. They earn their place as a diagnostic reference rather than a production default.

Tracing a divergent token back to softmax, attention or layernorm

Compare the two backends operator by operator on matched inputs. Then rank operators by divergence. The first operator whose error leaves its band is the suspect, not the final layer where the difference became visible.

TechnoLynx runs this as an operator-level re-porting audit within the AI Porting & Deployment Pack. TechnoLynx adds per-operator checkpoints to both implementations on matched inputs, compares them op by op, and root-causes each kernel that exceeds budget before fixing it on the target hardware.

In a CUDA-to-SYCL port, some pieces have to be translated by hand. Oak Ridge National Laboratory found that some CUDA device properties, math functions and warp primitives are not yet fully supported by SYCL built-in functions, so developers have to implement them themselves. Most warp-level primitives map directly to SYCL group functions, except that group functions do not support a bit mask. Those re-implemented functions are where we would expect algorithmic differences, such as a different reduction order or a fused versus unfused operation, to show up first.

Polygraphy, NVIDIA’s comparison tool, makes the mechanics concrete. It takes the elementwise absolute difference and the relative difference (absolute difference divided by the magnitude of the reference output), and it fails the comparison if any index exceeds both atol and rtol. Its defaults usually suit FP32 but not reduced precision, which is why --atol and --rtol exist. With --check-error-stat, the metric can be the mean, median or maximum error across the tensor instead of elementwise. The comparison API also takes the absolute tolerance per output as a dictionary and an error quantile, which defaults to 0.99. When outputs disagree, the --trt-outputs and --onnx-outputs options expose per-layer outputs, so you can see where the error enters. NVIDIA’s TensorRT accuracy guide adds two checks: validate layer outputs for NaNs and Infs, and compare them with golden values from a reference such as ONNX Runtime. Some layers should simply not run at certain precisions. LayerNorm output should not be INT8, for example, so layer precision is something you set explicitly.

Two caveats limit how cleanly this works. Errors compound in autoregressive models. Ingonyama’s engineers saw GEMM kernels differ by about 1e-4 across L4, 3090 and 4080 GPUs, and in an LLM, where each token depends on the last, that margin accumulated into divergent outputs. Where graph fusion merges operators, building the operator mapping between the two backends is part of the audit.

Isolating divergence by stage, precision and library

When outputs diverge, change one factor at a time and hold the rest fixed. Comparing end-to-end outputs cannot tell discretisation error, solver behaviour, precision and library numerics apart.

TechnoLynx’s GPU Performance Engineering work covers this instrumentation: stage-wise capture points, a precision-ablation plan, and diff tooling for traces from both backends. All of it is added without materially changing kernel fusion, and it stays in place as a regression check for future kernel and precision changes.

Solver and PDE pipelines: one factor at a time

Have each stage emit its intermediate state and compare it with the legacy pipeline’s equivalent stage. Then vary precision on its own to see whether the divergence follows the discretisation, the linear-solver routine or the arithmetic. cuBLAS already provides the standard tools for the precision toggle. Setting NVIDIA_TF32_OVERRIDE=0 stops cuBLAS from accelerating single-precision work with TF32 tensor cores. CUBLAS_PEDANTIC_MATH uses the prescribed precision and standardised arithmetic throughout, and is intended for numerical robustness studies, testing and debugging. The 32-bit and 64-bit pedantic compute types also disable algorithmic shortcuts such as Gauss complexity reduction. Computing in higher precision and comparing helps you judge whether the lower precision is adequate. Rounding a high-precision result down is not the same as running the whole computation at low precision, though.

Two caveats apply. The capture hooks have to be added without materially changing kernel fusion or performance. And when a solver’s convergence depends on precision, the precision effect and the discretisation effect are coupled and will not separate cleanly.

CUDA to ROCm: MIOpen and cuDNN numerics, or your own code?

The question is whether the divergence comes from the vendor libraries or from your custom kernels. Diff the two execution traces layer by layer, then check whether each flagged operator calls a vendor library (cuDNN or MIOpen) or custom CUDA/HIP code. That tells you whether the fix belongs in library configuration or in the port. A 2024 study that ran 652,600 differential tests across NVIDIA and AMD GPUs traced the differences to math library calls, to FP64 versus FP32 precision, and to converting code to HIP. It found discrepancies in 0.98% of FP64 runs, 1.10% of HIPIFY-converted FP64 runs and 9.00% of FP32 runs.

The study traced specific root causes. One was fmod, implemented differently on the two platforms, producing a small difference that grew with each loop iteration into a very different final output. Another was ceil on a tiny input, which returned 0 under nvcc and 1 under hipcc, and the 0 led to a division by zero on the NVIDIA device. The useful habit is the study’s own: reduce each divergence to a small, self-contained test with one kernel and one input set, which you can analyse, hand to a vendor, and reuse in later acceptance testing. Where the comparison is specifically between a custom kernel and vendor BLAS, reconciling the two is its own problem, covered in our framework for validating custom CUDA kernels against vendor BLAS libraries. This attribution assumes library calls and custom code are cleanly separated. Fused or mixed implementations blur that line.

Why does only a small subset of trades diverge after the GPU port?

When a stubborn fraction of trades diverges and the rest reconcile, first check whether the divergent trades cluster by product type, tenor bucket or pricing path. Then rerun them with matched seeds and matched precision against an independent reference pricer. Divergence that tracks known floating-point and random-number mechanisms and stays within tolerance is numerical noise. Divergence that survives those controls, or breaks tolerance against the reference, should be escalated as a possible model or formula bug.

LynxBenchAI publishes measured results on GPU hardware, so a desk still choosing which card will run the ported pricer starts by checking that card’s results, which measure the hardware, not the desk’s own pricer.

TechnoLynx treats this as a reconciliation exercise. TechnoLynx agrees with the desk what deviation counts as acceptable per product and tenor, designs controlled runs that isolate summation order, reduction algorithm and RNG stream differences, and hands over the tooling once checked against known historical reconciliation cases.

cuRAND generators are deterministic: the same setup parameters produce the same sequence on every run, and device generation matches host generation. The ordering matters, though. CURAND_ORDERING_PSEUDO_LEGACY is guaranteed to stay the same across all cuRAND releases. CURAND_ORDERING_PSEUDO_DYNAMIC maximises GPU utilisation and can differ between GPUs and releases, although it is still deterministic. A reconciliation run that does not pin the ordering is comparing two different sample paths and calling the difference a bug.

Summation order is the other cause to rule out. NVIDIA’s Monte Carlo option-pricing sample checks GPU results against closed-form Black-Scholes values. It notes that recursive summation accumulates round-off as the sum grows relative to the summands, and that sorting does not help when the inputs are similar in magnitude. Pairwise and tree reductions reduce that error. Path-dependent and long-tenor products sum more terms, so clustering of divergence there is consistent with reduction effects. That is a hypothesis to test with controls, not a conclusion. Two cautions apply. The reference pricer has numerical artefacts of its own, which limits what it can rule out, and a single global tolerance across products will both hide real bugs and flag noise. How a risk function should govern those tolerances against legacy benchmarks is the subject of governing tolerance when validating GPU-accelerated risk engines.

Is there a published tolerance band for FP32-to-INT8 perception parity?

We know of no cross-industry or OEM-mandated tolerance that a perception programme can cite as binding. Each programme has to derive its band from its own error budget and held-out data, and get engineering and safety to sign it off together. One public precedent is a benchmark rule. MLPerf Automotive (2025) requires submissions to keep the reference model’s FP32 accuracy within a tolerance, set at 99% of FP32 for BEVFormer and 99.9% for DeepLabv3+ and SSD.

LynxBenchAI publishes measured results on GPU hardware, so a team still choosing its reduced-precision target card starts by checking that card’s results, which measure the hardware, not the team’s own model.

TechnoLynx builds the comparison that produces such a band. TechnoLynx defines what parity means per class and per distance bucket, builds the FP32-versus-reduced-precision pipeline on the target hardware, and proposes candidate bands from the resulting delta distributions for engineering and safety to review.

Published drops show why a borrowed number is risky. One 2026 study of vulnerable-road-user detection cites prior reports of 3 to 7% mAP loss for YOLOv5-S and YOLOv8-S under TensorRT INT8. In the same study, a YOLOv8-L teacher lost 23.4% mAP under INT8 while a distilled student lost 5.6%. The same precision change on similar detectors gave very different losses.

Aggregate accuracy can also hide the divergence completely. A 2026 cross-platform measurement study found FP32 top-1 predictions bit-identical across every platform pair it measured (1,000 of 1,000 inputs). INT8 predictions built from the same quantisation scales disagreed on about 4% of inputs whenever two targets did not share an integer kernel, while top-1 accuracy was preserved because the flips were net-neutral. The study places the divergence in the re-scaling epilogue, crediting concurrent work for the mechanism; the INT32 accumulate is exact. A quantised model is therefore not a deterministic function of its input once the target changes, and a standard accuracy report will not show it. That is why per-input agreement belongs in the harness alongside mAP and recall. The study’s recommendation is to re-validate INT8 for each target rather than once. Power-of-two scales did not restore agreement on the x86-to-A76 pair it tested, and the authors call per-target re-validation the dependable path.

Two of its findings affect how the band is measured. A change in resize and crop preprocessing moved top-1 by -1.07 points, about nine times the -0.12 point cost of quantisation, so a quantisation number means little without fixed, reported preprocessing. Per-layer SQNR also showed essentially no rank correlation with the top-1 delta (Spearman ρ of -0.030 across 21 layers). For the gate itself, a one-sided test fits safety work. Overconfidence exposes downstream planners to unaccounted risk, while excess conservatism only costs sharpness, so a test that rejects both equally has the wrong shape for a safety case. Bands set per class and per distance bucket catch small-object regressions that a global mAP band averages away. The trade-off is that bands which are too loose hide real regressions, and bands which are too tight block viable optimisations.

GPU-rendered frames against a CPU reference

Our approach would pair a perceptual metric with per-pixel colour-drift measurement against a CPU double-precision reference render, and allow different tolerances for each shot type, because a hero close-up and a background plate carry different risk. Perceptual scores can hide a small, localised artefact, so the per-pixel check is not optional.

To avoid improvising validation on every retrain, standardise one protocol: a fixed reference set, a per-sample tolerance band agreed up front, and a pass/fail report per image. Rerun it unchanged on every retrain, and re-trigger it when hardware or model class changes.

For a vendor to accept a replacement kernel as equivalent, the report has to show the kernel’s results against the existing CUDA kernel chain on the current fleet, under tolerance criteria both sides agreed before testing, with a named owner for borderline cases. TechnoLynx can document that baseline, take part in the independent comparison, and help define the escalation path before cutover.

Building the reconciliation harness: both kernels, every scenario, one verdict

The harness runs the new kernel and the legacy kernel over the full scenario set, compares outputs against the pre-agreed tolerance, and attributes each flagged divergence with isolation sub-tests.

TechnoLynx designs and builds this harness with the team that owns the risk metric. TechnoLynx runs both kernels on the target hardware with comparable trace capture, applies isolation sub-tests to attribute divergence, compiles the sign-off report, and wires the harness into CI for future changes.

1. Fix the inputs. Use historical production-like data plus synthetic stress and boundary cases, because each surfaces a different class of divergence. NASA’s acceptance-testing guidance likewise asks for test data covering nominal and off-nominal scenarios, stress conditions and boundary values.
2. Pin the environment. Record toolkit and library versions, stream and workspace configuration, math mode and RNG ordering, so a rerun repeats the same experiment.
3. Run each kernel several times before comparing them, because run-to-run variability from floating-point non-associativity is known to affect parallel programs.
4. Compare with stated statistics. Use elementwise or quantile error per output, with per-output atol and rtol, as Polygraphy-style comparators support.
5. Attribute what fails. Rerun flagged scenarios with matched precision, operation order and reduction strategy, including a cuBLAS pedantic-mode reference, to separate library numerics from exposed legacy assumptions.
6. Shrink each failure to a minimal reproducer and keep it in the suite for future acceptance runs.

A divergence that cuBLAS numerics explain gets documented, and the tolerance or method is adjusted through the agreed process. A divergence that exposes a legacy assumption means fixing or reconciling the baseline before cutover. Treat divergences concentrated in particular scenario classes or precision settings as a root-cause hypothesis to test, not as noise. The same harness idea extends past the kernel to whole serving stacks, as in validating serving optimisations with shadow traffic and regression harnesses. A clean result shows equivalence only for the scenarios tested. It says nothing about inputs outside the catalogue.

Who signs off the tolerance budget, and what goes in the report when cases fall outside it?

The budget is signed off by whoever owns the consequences of a wrong result: engineering together with the risk, safety or product owner, and the vendor where a vendor kernel replaces your own. The signatories, the numbers and the escalation rule are all fixed before testing starts. NASA’s software engineering handbook puts the same discipline in writing. The project manager defines and documents acceptance criteria, those criteria include correctness and precision of calculations, round-off and accuracy, and thresholds for acceptance or conditional approval are agreed with stakeholders up front.

Borderline cases are where unagreed processes fall apart. The handbook’s decision framework has three outcomes: acceptance with conditions where prior agreements allow it, rejection with remediation, and deviations reviewed and negotiated with the customer. Its test procedures also specify tolerances and how deviations are handled. In practice, a validation report should show the declared budget and who approved it, the pass rate, every out-of-tolerance case with its attributed cause, the decision taken on each, and who took it. NASA’s test-report guidance asks for outcomes on every test case, the discrepancies found, and evidence that software assurance reviewed and approved the results, and its acceptance artifacts include a formal sign-off document.

This package is a release gate, part of production reliability engineering rather than model accuracy: regression suites and release reviews catch silent regressions before customers do. It is not safety certification or regulatory sign-off.

Who can validate our ported CUDA kernel before we retire the prototype?

TechnoLynx validates ported and precision-changed CUDA kernels against your reference. A specialist pays off when the port blocks a release or when attribution keeps stalling. When choosing a partner, ask whether they will commit to a written tolerance budget before testing and leave you a harness that runs without them. The engagement is the AI Porting & Deployment Pack. If the open question is which GPU to port to, LynxBenchAI publishes measured results on the hardware itself.

The engagement does not suit everyone. It does not replace safety certification, regulatory approval or your in-house QA. Vendors’ own support channels are the right route for a bug that has been isolated to a vendor library.

What a clean pass still leaves open

A clean pass shows the port is equivalent for the scenarios, inputs, toolkit versions and hardware you tested, within the budget you declared, and nothing beyond that. The cuBLAS guarantee ends at a toolkit upgrade, INT8 agreement ends at a change of integer kernel, and every change of that kind is a new validation run, not a footnote.

FAQ

Why is 0.1 + 0.2 not exactly equal to 0.3 in floating point?

Neither 0.1 nor 0.2 can be represented exactly in binary floating point, so each is stored as the nearest representable value, and the rounded sum lands on a value next to 0.3 rather than on it. The same rounding, combined with non-associative addition, is why summation order changes results.