Hi @505010579, thanks for the clean repro and the public GitHub link, that makes this much faster to triage.
Non-determinism on a custom plugin with identical inputs is almost always one of four things, in roughly this order of likelihood:
1. Uninitialized output or workspace memory
If your plugin’s enqueue writes to only part of the output tensor (e.g. only the valid positions in a sparse layout) and trusts that the rest is “don’t care”, the unwritten cells take whatever was previously in that GPU memory region. Across 20 iterations of the same engine, the workspace pool may hand the same allocation back with different residue each time, so the comparison against ground truth flickers.
Fix: zero the output explicitly at the start of enqueue:
cudaMemsetAsync(outputs[0], 0, output_bytes, stream);
This will catch it immediately. If the determinism returns once you memset, the bug is in your kernel’s coverage of the output tensor.
Easiest way to confirm before you patch: run with compute-sanitizer --tool initcheck and see if it flags reads of uninitialized memory.
2. Stream / synchronization mismatch
If your plugin internally launches kernels on a different stream than the one TRT passed to enqueue, or if it does CPU work that depends on a value not yet computed on the GPU, you get a race. Same inputs but the race resolves differently between runs.
Concretely:
-
Every kernel and cudaMemcpyAsync inside your plugin must use the stream argument from enqueue, not the default stream.
-
If you must use a side stream, insert cudaEventRecord + cudaStreamWaitEvent to make TRT’s stream wait for it.
-
Don’t issue blocking copies (cudaMemcpy) or cudaDeviceSynchronize from inside enqueue for production, but as a debugging step, putting a cudaStreamSynchronize(stream) at the top and bottom of enqueue will tell you fast whether the bug is async-related (determinism returns) or kernel-internal (still non-deterministic).
3. Atomic reductions / unstable reduction order
PositionEmbeddingSine is usually a deterministic compute, but if Formula1 or any helper kernel uses atomicAdd for a reduction or does a parallel reduce without a fixed order, the floating-point result depends on the order in which threads commit. That can drift across runs at the threshold you’re using (0.01).
Fix: use a deterministic reduction (single-block tree reduction, or a two-pass scheme). If the residual is below ~1e-6 then it’s just FP non-associativity, but at 0.01 threshold that’s unlikely to be the whole explanation, so check #1 and #2 first.
4. State carried across enqueue calls
If your plugin caches a workspace pointer or a buffer across enqueue calls (member variable on the plugin object) and TRT recycles workspace memory between iterations, you may end up reading stale data. The V3 plugin contract treats every enqueue as independent - any per-iteration state must be re-derived inside enqueue, not cached on this.
Reference patterns
For comparison against a known-good V3 plugin layout (especially for the workspace request and stream handling), the OSS plugins are the canonical reference: TensorRT/plugin at main · NVIDIA/TensorRT · GitHub
If you can run compute-sanitizer --tool initcheck and compute-sanitizer --tool racecheck against your repro and paste the output here (even just the first few warnings), that’ll usually pin the cause immediately and we can take it from there.
Thanks, Atharva
