GPU risk engine validation does not require the GPU to reproduce the legacy CPU engine bit for bit, and it cannot accept “close enough” on the build team’s word. A defensible tolerance is derived from the legacy engine’s own numerical noise, written down before the comparison runs, applied deal by deal, and approved by people who did not build the port. NVIDIA’s floating-point guidance states that differences between CPU and GPU results are not automatically evidence that the GPU result is wrong, and the authors of a validation-centric GPU port of a 250,000-line weather code add that such differences cannot simply be ignored either.
TechnoLynx’s AI engineering for finance delivers this through AI Governance and Trust, deriving the tolerance from the legacy engine’s noise and producing re-runnable evidence for your validators; the first step is running the equivalence suite per metric.
GPU risk engine validation: which tolerance survives P&L, VaR and XVA sign-off?
SR 11-7 and SS1/23 leave the threshold to each firm rather than supplying a basis-point figure. The Federal Reserve’s SR 11-7 guidance (a 2011 letter, since marked as superseded) refers to the banking organisation’s own “predetermined thresholds of acceptability”, and the PRA’s SS1/23 asks that a board’s model risk appetite include thresholds for acceptable model performance and tolerance for errors. So the number has to be measured before it is negotiated.
Published GPU ports outside finance show three ways to measure it. The MURaM solar physics team set its acceptance tolerance as the variance observed between CPU runs with different MPI decompositions and core layouts, a maximum relative error of 1e-5. MeteoSwiss’s ICON port built a CPU ensemble with perturbations in the least significant digits of selected inputs, derived variable-specific tolerances from the ensemble spread, and passed the GPU run only if it fell inside that spread for every diagnosed field. The third route is analytical: bound the rounding error of the summations the engine performs, for which deterministic bounds proportional to tree height and probabilistic bounds proportional to its square root have been derived.
The noise floor is small, but not zero. In a study of GPU arithmetic for trading systems, running the same naive parallel reduction 200 times over 100,000 portfolio risk contributions in double precision produced 8 distinct results, spread by 8.5e-4 on a base risk number of 5.9e11. That spread is about six and a half units in the last place of double precision at that scale. A band set at a few basis points because it feels prudent therefore leaves a great deal of room for a real defect to hide underneath it.
The noise is not uniform across outputs, though. Headline numbers that are differences of large, similar quantities (net exposure, P&L explain, stressed-minus-base VaR) can turn a relative input error of 1e-16 into 1e-6 or worse. That is why we would set tolerances per metric, not one band for the whole engine. The decision order follows from this: run the equivalence suite first to characterise the actual deviation per metric, then set the threshold with the accountable function, knowing that too loose a band masks bugs and too tight a band produces false breaks every cycle.
Who signs off the tolerance: model risk, market risk or finance?
Both, in different roles, and never the developers alone. Model risk (the second line) owns the validation method and the challenge, and we would have the business owner of the number confirm the band is fit for the decision it gates. SS1/23 expects firms to name dedicated model approval authorities, committees or individuals, with defined procedures when a material change is made to a model, and it makes the accountable senior manager responsible for ensuring independent validation and effective challenge. One practitioner summary of SR 11-7 tiering places CRO sign-off on risk acceptance for its highest-tier models, such as those driving regulatory capital. Where the firm’s own tiering places the VaR or XVA model in that tier, we would expect the tolerance to be approved at that level rather than inside the project.
When the solver converges differently near a discontinuity in the PDE grid
This is where judgement carries the case. The weather-code port found discrepancies in five kernels caused by floating-point and intrinsic-function differences, including threshold-sensitive branch divergence and cancellation. We would expect the same in a PDE pricer: near a barrier, an early-exercise boundary or a payoff kink, pointwise differences can exceed a band that holds elsewhere. We would isolate those trades, compare them on a refined grid or a smoothed payoff, and record a case-specific exception with its rationale and approver, rather than widen the global band until they pass.
Should the tolerance be absolute, relative (basis points) or both?
Both, combined so that each covers the other’s blind spot. A fixed absolute tolerance such as 1e-9 is wrong for large magnitudes, because the spacing between representable doubles near 1e12 already exceeds 1e-9 by orders of magnitude. A relative tolerance fails the other way: MURaM found that bit errors around 1e-7 produced large relative errors wherever the reference was close to zero, so it also accepted points whose absolute error was below 1e-7.
Worked example: a combined absolute and relative test
| Assumptions, stated as assumptions: the rule is pass if | GPU − CPU | ≤ max(a, r × | CPU | ), with a relative band r of 1e-5 (0.1 basis point) and an absolute floor a of 0.01 in the trade’s currency. The values below are illustrative inputs, not measurements. |
| Reference value (CPU) | GPU value | Difference | Limit max(a, r× | CPU | ) | Relative-only test | Combined test |
|---|---|---|---|---|---|---|---|
| 2,500,000.00 | 2,500,012.00 | 12.00 | 25.00 | pass (4.8e-6) | pass | ||
| 0.40 | 0.405 | 0.005 | 0.01 | fail (1.25e-2) | pass | ||
| 0.40 | 0.43 | 0.03 | 0.01 | fail (7.5e-2) | fail |
The second row is the near-zero case that sinks a relative-only rule: a hedged position whose value nets to almost nothing will break every night on noise. An absolute-only rule fails in the first row instead, flagging a benign difference on a large notional.
The band should also vary by trade or contract type where conditioning varies: a vanilla swap and a callable with a barrier do not deserve the same number.
Payout engines: a tolerance specification per territory and contract rule
For royalty and payout engines, the specification is a table rather than a number: one row per territory and contract rule, each with its rounding convention, an absolute band in the currency’s minor unit for per-title lines, a relative band for aggregated statements, and the name of whoever approved it. We would use the absolute band to protect each per-line amount paid to the rights holder, and the relative band on totals to catch drift that no single line reveals. Finance, legal and audit sign the table before it becomes a go-live gate, because the rights holder’s question is contractual, not numerical.
Is a CPU-versus-GPU difference floating-point noise or a bug in the port?
Parallelising rearranges operations and so changes numeric results; we would tell noise from a defect by localising each difference and comparing it against the legacy engine’s own spread. On a GPU, changing the number of threads per block reorganises a parallel reduction, which rearranges the parentheses in a long string of additions. Reductions built on atomic operations commit in an order set by hardware scheduling, so the same kernel on the same input and hardware can produce different bit patterns on consecutive runs. The GPU uses fused multiply-add where the CPU may not, the CPU may compute in higher precision than expected, and common mathematical functions are not required by IEEE 754 to be correctly rounded. CUDA’s math library and the host library (glibc on Linux) are independent implementations, so their results will often differ slightly. cuBLAS and cuDNN also select among kernel implementations at runtime, and a library upgrade can silently change the default kernel for a given shape.
Diagnostic checklist for a CPU-versus-GPU discrepancy
- Compare kernel by kernel against reference data dumped from real CPU runs, element by element, as the weather-code port did.
- Then validate at application level, because a difference acceptable in one snapshot can accumulate over a time-dependent simulation; we would treat path-dependent pricing and multi-step risk runs the same way.
- Enable GPU kernels selectively to localise the kernel group or integration point responsible for an application-level failure.
- Where the tooling allows, patch incorrect values with their reference so one error does not propagate into later tests; PGI’s PCAST offered exactly this.
- Check whether the difference sits inside the spread of perturbed CPU runs; if it does, it is noise by the standard you agreed.
- Run a dual-precision comparison and an input-perturbation test: if a 1e-14 wobble in the inputs moves VaR by a basis point, the problem is conditioning, not precision.
- On the CPU side, compare using SSE instructions rather than x87 extended-precision registers, so both sides follow IEEE 754 single or double precision.
- Look at the shape of the residuals: we would treat differences that cluster on particular trade types, branches or dates, and survive a switch to deterministic reductions, as suspected port defects.
The weather-code authors note that whether a discrepancy is acceptable stays with the human developer, not with their AI agent’s specification. The kernel-level mechanics of proving kernel equivalence after a port or a precision change, including mixed-precision and backend divergence, are covered in more depth there.
Who validated this independently of the team that built the GPU pipeline?
The answer the committee wants is a named validation function with no stake in the outcome, not the build team’s own test report. SR 11-7 says validation is generally done by staff who are neither responsible for development or use nor invested in whether the model is found valid, and that any validation work done by developers must be critically reviewed by an independent party. SS1/23 says validators should not be part of model development or have a stake in approval, and expects firms using internal capital models to show separate reporting lines. A practitioner summary of SR 11-7 calls the failure case “independence theater”: validators reporting to the same leadership as developers.
XVA engines get no exemption. An IACPM and Fintegral practitioner survey (2018) reported that XVA models now follow the same strict validation process as any other front-office pricing model, and that validators must demonstrate to the regulator that they were independently checked.
The independent work itself has a known shape. Validators replicate the model’s logic and testing from its documentation to reproduce the development results, may run further benchmarking to confirm that documented performance claims are factual, and issue a report in which material flaws send the model back for redevelopment. For a GPU port, the deal-by-deal pack we would hand them contains the equivalence definition per engine, the tolerance and how it was derived, the scenario and path coverage, per-deal VaR and XVA results against the CPU engine, every discrepancy classified with its root cause, and an environment manifest with rerun instructions.
How do we approve the tolerance band once instead of re-litigating every month-end break?
Approve the band as a standing policy with explicit reopen triggers, and let breaks inside it go to a log rather than a meeting. MeteoSwiss runs its tolerance test as a gate every model change must pass, and generates a new ensemble reference only when a physical change moves results beyond numerical noise. SS1/23 asks for clear guidelines on a maximum tolerance for performance exceptions, with defined triggers and thresholds that bring in controls such as heightened review once breached, and for root cause analysis when thresholds are breached.
We would also have the policy say when the standing approval is revisited. For internal VaR models, Delegated Regulation (EU) 2015/942 lists a reduction in the number of simulations, the introduction or removal of variance reduction methods and changes to the random number algorithms among the statistical-method changes it covers. It also covers IT system changes that amend the calculation procedure of the model. We would add CUDA, driver and math-library upgrades and any change of GPU architecture to the trigger list, given how library drift can change kernels silently.
Can a GPU pricing or risk engine be made reproducible for audit?
Yes within one pinned software stack on one GPU architecture; NVIDIA puts the cost of GPU-to-GPU deterministic reductions at 20% to 30% more execution time on large problems, and bit identity with the legacy CPU is not to be expected. NVIDIA’s CCCL library offers three determinism levels for reductions: not guaranteed, run-to-run (the same value on a single GPU) and GPU-to-GPU (the same value on any GPU). Run-to-run is the default for cub::DeviceReduce, achieved with a fixed hierarchical reduction tree instead of atomics. In tests reported by the EigenAI authors, Hopper outputs were bitwise identical across runs, GPU counts and hosts with the same GPU model, while A100 and H100 outputs matched in 0% of cases. Their deployment rules for deterministic inference read like an audit checklist: pin exact CUDA and driver versions, enable deterministic modes in cuBLAS and cuBLASLt, disable atomics and autotuners, and seed PRNGs deterministically with the seeds recorded.
Published sources disagree on whether audit needs bit-exact identity at all. One position holds that regulatory reproducibility means bit-exact identity and that a tolerance-based reconciliation leaves a documented hole in the audit trail. A survey of reproducibility in high-performance computing accepts reproducibility “with the stated precision” rather than bitwise identity, and a risk practitioner’s guide on floating-point error argues that non-reproducibility across runs should be a documented property with a stated tolerance. Our view is that the two sides answer different questions. Replaying the GPU engine’s own run of record should be bit-exact, and with a pinned stack on one architecture it can be. Agreement with the legacy CPU engine is a governed tolerance, because the CPU and GPU math libraries are independent implementations and the GPU uses fused multiply-add.
Matched random seeds help only as far as the generators match. cuRAND sequences are deterministic for a given seed and set-up, and generating on the device gives the same sequence as generating on the host. The CURAND_ORDERING_PSEUDO_DYNAMIC ordering, however, can differ between GPUs and across releases, while the legacy ordering is guaranteed stable across cuRAND releases. Counter-based generators such as Philox add one hard rule: a key and counter pair must never be re-used, or the same random numbers repeat. Where the legacy engine uses a different generator or library, the same seed does not give the same stream, as the HPC survey notes for Philox across languages and libraries. Path-by-path comparison then gives way to a statistical comparison of distributions, and reconciling cuRAND paths with a CPU library’s random streams is a methodology question in its own right.
Can an independent firm confirm our catastrophe model reproduces the last five landfall events?
It can, but “reproduces” has to mean “falls within a statistical margin agreed in advance”, because a Monte Carlo loss simulation will never return a single historical outcome exactly. We would have the validator re-run the engine, or an audited replica, under controlled seeds and versioned inputs, then compare the simulated payout distribution for each event with the realised payout.
TechnoLynx works on the engine side of this, not in the validator’s chair: it reviews and, where needed, re-implements the loss-simulation pipeline for reproducible runs with fixed seeds and versioned inputs, supports the backtests on the target compute, documents the implementation for the independent validator and helps remediate what they find.
The margin is a statistical statement, such as a percentage deviation in expected loss or overlap of confidence intervals, agreed between issuer, validator and investors before any run. We would have the margin allow for sampling variance as well as model error, and have the validation report state plainly that five events is a small sample. The American Academy of Actuaries’ summary of ASOP 38 lists comparing output with historical observations and with alternative models among the ways an actuary validates catastrophe model output. It also expects documentation another qualified actuary could assess, and disclosure of the validation methodology and the extent of reliance on experts. Catastrophe model output is widely used for ratemaking, reinsurance purchase, capital and solvency assessment.
Credible validators are specialist model-risk or catastrophe-model validation firms with relevant actuarial and modelling credentials and no commercial stake in the result. SS1/23’s test applies by analogy: validators should have no stake in whether the model is approved. Independence is compromised if the same firm built or maintains the pricing platform, and a tolerance one investor group accepts may not satisfy another. If the goal is internal confidence only, an in-house backtest is enough; investor-grade assurance needs the third party. A statistically defensible acceptance test for GPU-versus-CPU loss percentiles is a related but separate design problem.
Particle filter or closed-form approximation: one divergence number on the same five years of futures data
Yes, you can get a single number: run both models on identical inputs in one GPU-accelerated backtesting harness within a fixed runtime budget, report a summary divergence statistic, and break it down by historical regime. The tolerance behind “pass” has to be set in trading terms, not numerical ones.
LynxBenchAI compares candidate GPUs on one scale in its published results, so a team whose runtime budget depends on which GPU it buys starts by checking those results before fixing the budget; they measure the devices, not your filter.
TechnoLynx builds that harness: it frames the inputs, runtime budget and divergence metric with the quant team, ports both models onto the existing CUDA kernel chain, adds per-regime reporting, reconciles against the independent reference pricer and tunes the kernels until both fit the budget.
Numerical closeness is not the same as tradeable equivalence. The trading-systems study cited earlier describes how a single unit-in-the-last-place discrepancy can alter position sizing, trigger a stop-loss on a different tick or shift a rebalance’s weight vector. So we would express the tolerance through the decisions the model drives (signal sign, position size, hedge ratio), not only the state estimate. That choice is a business decision and needs revisiting when the strategy changes.
SS1/23 asks for backward-looking performance tests across a variety of economic and market conditions relevant to the model’s use, and for comparisons with challenger models, alternative implementations of the same theory, where practicable. High-volatility and crisis windows are where we would expect the approximation to separate from the filter first; the companion page on backtesting and validating risk engines against the 2008 and 2020 crisis losses covers replicating those episodes. If divergence stays within tolerance in every regime and within budget, the approximation can stand in for the filter; if any regime breaks, keep the filter there or reject the approximation. Reconciling against one reference pricer validates against one independent source, not against ground truth.
How do we scale from a cheap pilot to a full historical backtest within budget and deadline?
Stage it so that each step buys the right to spend on the next:
- Characterise noise on a reduced problem: ICON’s tolerance ensemble ran on a short configuration over a reduced horizontal domain, cheap enough to repeat.
- Validate kernels against dumped reference data, then run a shortened application-level window; the weather port cut validation to 360 timesteps to make repeated runs feasible while still letting differences accumulate.
- Run a representative historical window on the full book, fixing any defect before widening scope.
- Run the full history plus the enumerated stress and tail scenarios, because snapshot validation does not certify every control path and application-level runs expose errors outside the captured states.
- Book the committee date after step 3 passes, not before, so the expensive run happens once.
Forecasting the runtime of the full re-simulation on the actual cluster is its own planning problem, and it belongs before step 4.
Does the GPU engine reproduce trustee payouts and mainframe outputs for every deal and territory?
Showing it means backtesting the cashflow-waterfall or royalty engine against audited or legacy outputs across the full historical deal and scenario space, quantifying discrepancies by structure type, and packaging the comparison so it can be re-run on demand. In a migration, the legacy system’s current behaviour becomes the specification and the test oracle. The method is characterisation testing: run the same executions on both systems, save both outputs and analyse every difference. Bugs already in production will surface too, since the legacy system is the reference, so decide in advance how to treat them.
LynxBenchAI compares candidate GPUs on one scale in its published results, so the first step before sizing the target compute for these backtests is to check which device suits the run; those results measure the devices, not your engine.
We would expect the largest discrepancies where the logic branches: trigger breaches, caps, PIK toggles and rare contract clauses, where a small numerical difference can flip which path a deal takes. Those structures deserve targeted cases rather than whatever the historical sample happens to contain, since a small set of historical deals may not cover every edge case in the live book. A tolerance suited to aggregate payouts can also hide a large error in one rare tranche or clause, so bands apply per tranche or per contract line as well as in total.
The credit committee or the auditors set the tolerance before the backtests run; the engineering team proposes it with evidence but does not set it unilaterally. Abstracta’s account of a fintech COBOL migration describes testing depth set per component by risk, compliance requirements and agreement with the business, with the system broken into slices that each carry their own controls and evidence. Coverage should be full rather than sampled where feasible, and go-live waits until equivalence is documented and reproducible across that space. Reproducibility on demand has a running cost: the harness and reference data must be maintained long after go-live.
Backtesting exceptions and P&L attribution as model risk management evidence
The GPU engine’s backtesting exceptions and P&L attribution should be computed alongside the legacy engine’s over the same window, broken down by risk factor; this feeds the outcomes analysis SR 11-7 lists as a required validation activity, though only the regulator decides acceptability.
TechnoLynx’s Production AI Monitoring Harness computes the GPU engine’s backtesting exceptions and P&L attribution alongside the legacy engine’s over the same window, broken down by risk factor; the first step is fixing that window and the factor breakdown with your validation team.
A practitioner summary of SR 11-7 lists outcomes analysis, backtesting against realised outcomes, as one of three required validation activities alongside conceptual soundness and ongoing monitoring. The Basel framework’s MAR32 compares one-day VaR with both actual and hypothetical P&L over the prior 12 months, counts an exception whenever either loss exceeds VaR, and sorts results into green, amber and red zones by exception count. At desk level, the PLA test is green when the correlation between risk-theoretical and hypothetical P&L is above 0.80 and the Kolmogorov-Smirnov metric is below 0.09.
For an engine swap, the useful comparison is not just the count but the pattern: the same exception days, with any difference explained. We would expect differences on days where the loss sits close to the VaR number, where numerical noise can tip a day either way, and would list those days individually. MAR32 is explicit that adjustments to RTPL or HPL to remove residual operational noise from computing them in two systems are not permitted, so CPU-versus-GPU attribution differences have to be explained at factor level, not netted away. Attribution gaps may come from legitimate precision differences rather than model errors, and the factor breakdown is what separates the two.
Regulatory acceptance stays with the supervisor, which determines its response from the exceptions over 250 trading days. The IACPM and Fintegral survey adds that approval from one regulator does not guarantee another’s.
What evidence trail do risk committees, investors and auditors expect before go-live?
They expect a trail a non-engineer can read, defend and sign against, and that trail is an engineering output, built and versioned like the engine itself, not a policy document written afterwards. Its core is a register of every discrepancy, classified and root-caused. SS1/23 asks for root cause analysis that establishes whether model limitations reflect significant deficiencies needing remediation, and for deficiencies to be documented, reported to owners, users, validators and senior management, and tracked in the model inventory.
| Discrepancy class | What it means | What the trail records |
|---|---|---|
| Within agreed noise | Inside the band derived from legacy spread | Count and distribution per metric |
| Conditioning | Ill-conditioned metric amplifies benign error | Perturbation result and metric-specific band |
| Approved exception | Discontinuity or structure outside the global band | Rationale, approver, review date |
| Legacy defect | Legacy engine was wrong; GPU is right | Decision on which behaviour ships |
| Port defect | Bug introduced by the GPU implementation | Fix, retest, closure |
For an internal-model (IMA) change requiring approval, Delegated Regulation 2015/942 lists the documents to submit, including technical documents, independent review or validation reports, confirmation of internal approval with its date, and current and previous model version numbers. The trail itself must be trustworthy as evidence: a log a privileged administrator can edit is weak unless the organisation can demonstrate integrity, access governance, retention and repeatable verification. Structuring that material into a model documentation and approval package that survives audit is covered separately. None of this is regulatory sign-off; it is the set of artefacts a sign-off needs.
Before cut-over, we would require an enumerated list of stressed and tail scenarios, including counterparty defaults and netting-set edge cases, reconciled within tolerance, with a named validation owner signing off coverage. We would treat coverage as risk-based rather than exhaustive, and an untested tail category as making the sign-off conditional at best.
Equivalence testing methodology for COBOL decimal to floating-point engines
For a royalty engine moving from COBOL fixed-point decimal to a GPU floating-point reimplementation, the equivalence testing methodology should be an epsilon-bounded, row-level comparison backed by aggregate statistical checks across the five years. Exact match is realistic only if both systems use the same fixed-point arithmetic end to end.
TechnoLynx’s Inference Cost-Cut Pack runs row-level and aggregate equivalence checks of the GPU reimplementation against the COBOL engine’s outputs, so the move to the cheaper engine is shown not to change what rights holders are paid; the first step is setting the per-row epsilon from the documented rounding rules.
Abstracta’s guidance frames the goal of a COBOL migration as behavioural equivalence, with fixed-point and packed-decimal calculations producing the same results afterwards. Binary floating point cannot meet that literally for every value, since a decimal such as 0.05 is not exactly representable. Where rounding is contractual, our preference is to keep money in integer minor units or a decimal type on the GPU and confine floating point to rates and intermediate quantities, which shrinks the epsilon you have to defend.
| Approach | Choose it when | Its blind spot |
|---|---|---|
| Exact match | Both engines use identical fixed-point arithmetic end to end | Not met for every value once money moves to binary floating point |
| Epsilon-bounded per row | Floating point is used; we would set epsilon from documented rounding rules | Misses systematic bias too small per row to trip the band |
| Aggregate statistical check | Added to row-level checks to catch drift across periods | Offsetting row errors cancel, so it never replaces row checks |
Five years of data may include rule changes (currency conversion, proration policy) that need a separate tolerance profile per period. The trade-off is the same one every tolerance on this page carries: set too loosely, the band hides real calculation bugs; set too tightly, it reports benign rounding as mismatches.
Who can build the validation harness and evidence pack for our GPU-ported risk engine?
TechnoLynx builds the CPU-versus-GPU comparison harness, derives and documents the tolerance from the legacy engine’s own noise, and produces the classified, re-runnable evidence your independent validators and committee review. Building in-house works when your quant developers can take this on and your second line can challenge it; a specialist helps when the port and its validation compete for the same few engineers before a deadline. When choosing a partner, ask whether they will leave you a harness you can rerun each cycle, and whether they keep clear of the validator role. The engagement runs under AI Governance and Trust.
This work does not suit every case. We are not the independent validator, we do not give legal advice or regulatory sign-off, and we cannot promise a regulator will accept the result. Teams that want contractor hours or only need the cheapest port will be better served elsewhere. Specialist model-validation consultancies provide the independent challenge, actuarial firms cover catastrophe-model attestation, and NVIDIA’s documentation on floating point and determinism is a sound starting point for teams doing the engineering themselves.
FAQ
What is model risk?
Model risk is the risk of loss arising from decisions based on incorrect or misused model outputs. SR 11-7 traces it to errors in the model itself and to misuse of its outputs, such as applying a model outside its intended scope.
What is the role of model validation in banking?
Model validation is the set of processes and activities intended to verify that models perform as expected, in line with their design objectives and business uses. It is generally performed by people independent of the model’s development and use, and any validation work done by developers is critically reviewed by an independent party.
Is 100 basis points equal to 1 percent?
Yes: one basis point is one hundredth of a percentage point, so 100 basis points equal 1 percent.
