Request type
Add missing documentation
Documentation location
include/trtmc/runtime/trt_module.h — the forward() declaration
website/docs/architecture/runtime-lifecycle.md — "Runtime source map", which
points at that header as the backend abstraction
Problem or missing content
ITrtModule::forward() carries two contracts that are not written down
anywhere, and both fail by producing plausible wrong output rather than an
error.
1. Returned tensors point into a per-module staging buffer that the next call
overwrites. In src/runtime/backend/trt_module_impl.cpp:
auto& staging = host_output_staging_[name];
cudaMemcpy(staging.data(), entry.d_ptr, runtime_nbytes, cudaMemcpyDeviceToHost);
Tensor t;
t.data = staging.data();
host_output_staging_ is a member map keyed by tensor name
(src/runtime/backend/trt_module_impl.h). So a caller that holds a Tensor
from one forward() across a second forward() on the same module silently
reads the second call's data.
The declaration in include/trtmc/runtime/trt_module.h is grouped under a bare
// Forward passes comment and says nothing about this.
2. Outputs come back in the engine's own dtype (t.dtype = entry.dtype). A
bf16 engine returns half-width data, so a caller that allocates or reinterprets
as float32 gets garbage rather than a shape or type error.
Why this is worth a few lines in the header: it is a shared, model-agnostic
runtime contract, and the failure mode is silent. Concretely, in #1123
classifier-free guidance was a complete no-op for some time — the conditional
and unconditional branches ran through one module, the first pointer was still
held, and after the second pass both pointers referred to the unconditional
result. The pipeline ran, produced audio, and sounded plausible. It was found
only by measuring conditional_vs_unconditional rms == 0, not by reading the
code. Any family that runs two branches through one module and compares them is
exposed to the same thing.
Verification and search performed
grep -rn "host_output_staging" website/docs plugins # 0 results
grep -rn "output buffer" website/docs # 0 results
grep -n "forward\|lifetime\|reuse\|buffer" include/trtmc/runtime/trt_module.h
# -> only the declarations and an unrelated "Direct buffer access (KV cache binding)"
grep -rn "next forward\|reused across\|alias" src/runtime/models/*/pipeline.cpp
# -> 2 hits, both in the single family added by #1123, both written after hitting the bug
Checked and found not to cover it: website/docs/architecture/runtime-lifecycle.md,
website/docs/architecture/runtime-plugins.md, the fp16-trt-network skill
(covers dtype inside the network, not on the runtime output side), and
website/docs/wiki/TRT-Internals.md (an archived redirect stub — this is not a
request to restore it).
Suggested correction
A short comment block above the forward() declarations in
include/trtmc/runtime/trt_module.h, stating that returned Tensor::data
points into per-module storage reused by the next call on that module, that a
caller must copy anything it needs to keep across calls, and that outputs carry
the engine's dtype and may be half-width. Optionally one cross-referencing
sentence in runtime-lifecycle.md, which already names this header.
Scoping note: this describes the host forward() path. forward_device() and
externally bound buffers behave differently, so the wording should say which
path it covers — a maintainer confirming the intended contract is more useful
here than my inferring it from the implementation, which is why this is an issue
rather than a PR.
Happy to send the patch if the wording direction looks right.
Submission checks
Request type
Add missing documentation
Documentation location
include/trtmc/runtime/trt_module.h— theforward()declarationwebsite/docs/architecture/runtime-lifecycle.md— "Runtime source map", whichpoints at that header as the backend abstraction
Problem or missing content
ITrtModule::forward()carries two contracts that are not written downanywhere, and both fail by producing plausible wrong output rather than an
error.
1. Returned tensors point into a per-module staging buffer that the next call
overwrites. In
src/runtime/backend/trt_module_impl.cpp:host_output_staging_is a member map keyed by tensor name(
src/runtime/backend/trt_module_impl.h). So a caller that holds aTensorfrom one
forward()across a secondforward()on the same module silentlyreads the second call's data.
The declaration in
include/trtmc/runtime/trt_module.his grouped under a bare// Forward passescomment and says nothing about this.2. Outputs come back in the engine's own dtype (
t.dtype = entry.dtype). Abf16 engine returns half-width data, so a caller that allocates or reinterprets
as float32 gets garbage rather than a shape or type error.
Why this is worth a few lines in the header: it is a shared, model-agnostic
runtime contract, and the failure mode is silent. Concretely, in #1123
classifier-free guidance was a complete no-op for some time — the conditional
and unconditional branches ran through one module, the first pointer was still
held, and after the second pass both pointers referred to the unconditional
result. The pipeline ran, produced audio, and sounded plausible. It was found
only by measuring
conditional_vs_unconditional rms == 0, not by reading thecode. Any family that runs two branches through one module and compares them is
exposed to the same thing.
Verification and search performed
Checked and found not to cover it:
website/docs/architecture/runtime-lifecycle.md,website/docs/architecture/runtime-plugins.md, thefp16-trt-networkskill(covers dtype inside the network, not on the runtime output side), and
website/docs/wiki/TRT-Internals.md(an archived redirect stub — this is not arequest to restore it).
Suggested correction
A short comment block above the
forward()declarations ininclude/trtmc/runtime/trt_module.h, stating that returnedTensor::datapoints into per-module storage reused by the next call on that module, that a
caller must copy anything it needs to keep across calls, and that outputs carry
the engine's dtype and may be half-width. Optionally one cross-referencing
sentence in
runtime-lifecycle.md, which already names this header.Scoping note: this describes the host
forward()path.forward_device()andexternally bound buffers behave differently, so the wording should say which
path it covers — a maintainer confirming the intended contract is more useful
here than my inferring it from the implementation, which is why this is an issue
rather than a PR.
Happy to send the patch if the wording direction looks right.
Submission checks