Controlled by the runtime
- Buffer lifetime and address stability
- Lane ownership and queue bound
- Fixed shape and graph schedule
- GPU-stage ordering
- Warm-up separation
Inside WarpStream TRT: a bare-metal C++20, CUDA, and TensorRT 11 runtime that treats fixed memory, launch overhead, synchronization, and tail latency as one design problem.
Once a batch-one model becomes small enough, the work surrounding it starts to look uncomfortably large. A memory allocation, a CPU wake-up, three kernel launches, a layout conversion, or one accidental synchronization can consume the same order of time as the network itself.
That is the problem behind WarpStream TRT. The repository is not trying to become another universal model server. It is a constrained reference runtime for latency-critical RGB streams, built so that every address, execution context, stream, transition, and wait can be reasoned about.
The central idea is simple: if the steady-state path must be predictable, then dynamism belongs at startup. Allocate once. Resolve shapes once. Bind addresses once. Capture the GPU work once. During traffic, replay a bounded plan over memory whose ownership never changes.
Not “the lowest average kernel time.” The target is a complete producer-to-result path whose latency distribution can be explained—and whose p99 does not depend on a surprise allocator, shape update, or host round-trip.
The runtime exposes mapped pinned input memory to the producer. A fused CUDA kernel reads those RGB8 bytes, resizes and normalizes them, and writes the exact NCHW tensor expected by TensorRT. Inference produces device-resident logits. A final CUDA kernel selects Top-K and writes only the compact result into mapped host memory.
Preprocess, inference, and Top-K are captured into one lane-specific CUDA Graph. There is no CPU decision point between them. The application synchronizes only when it asks for the completed result.
| Producer format | Interleaved RGB8 |
| Network input | Batch 1 · fixed-shape FP32 NCHW |
| Network output | Fixed-size FP32 classification logits |
| Postprocess | Exact Top-K, K ≤ 32 |
| Concurrency | Fixed, preallocated execution lanes |
| Steady-state allocation | None in submit/wait path |
“No allocation in the hot loop” sounds like a micro-optimization. It is really a control over variance. General-purpose allocators carry locks, metadata, fragmentation, deferred cleanup, page behavior, and failure paths. CUDA allocation can introduce still larger synchronization and driver costs.
WarpStream moves all of that work into construction. Each lane owns a mapped input slot, a device tensor, device logits, a mapped result slot, two timing events, a non-blocking stream, one TensorRT execution context, and one instantiated CUDA Graph. The addresses bound with setTensorAddress() stay stable for the pipeline lifetime.
The implementation also makes failures explicit. A lane cannot be submitted unless it was acquired, cannot be reused while busy, and cannot silently grow a queue. Its state machine is deliberately boring:
That boundary is important. Unbounded queues are throughput-friendly until they turn burst pressure into catastrophic response time. Fixed lanes make backpressure part of the API instead of a production surprise.
WarpStream allocates input and Top-K output with cudaHostAllocMapped. The CPU and GPU receive usable views of the same pinned allocation, so a camera or decoder can write directly into the slot consumed by the preprocessing kernel. There is no staging allocation and no explicit host-to-device copy inside the runtime.
But on a discrete GPU, the bytes still travel across PCIe when the kernel reads them. Mapped memory is genuinely zero-copy in the software architecture; it is not magically device-local. A preallocated cudaMemcpyAsync into VRAM can win when the transfer engine overlaps well and the preprocessing kernel benefits from local bandwidth.
Compare mapped reads against an asynchronous copy into a fixed device input buffer. Then test GPUDirect, peer memory, or device-native decode if the producer can expose a GPU pointer. “Zero-copy” should be a measured topology choice, not a marketing adjective.
Mapped memory must also be enabled before a CUDA context is created. The runtime calls cudaSetDeviceFlags(cudaDeviceMapHost) before TensorRT initialization and fails loudly if another library already established an incompatible context.
A conventional preprocessing chain may resize into one buffer, convert integers to floats in another pass, normalize in a third, and finally transpose HWC into NCHW. That design reads and writes the same image repeatedly and launches multiple kernels before inference even begins.
WarpStream assigns one thread to each destination pixel. That thread computes the center-aligned source coordinate, fetches the four bilinear neighbors for each RGB channel, interpolates, scales from byte range to [0,1], applies ImageNet mean and standard deviation, and writes directly into the destination channel plane.
source RGB8 (HWC)
└─ bilinear sample
└─ value × 1/255
└─ (value - mean[channel]) × inverse_std[channel]
└─ destination[channel × H × W + y × W + x]No resized image exists. No normalized HWC tensor exists. The only materialized result is the tensor TensorRT consumes. This is the useful meaning of fusion: remove memory traffic and launch boundaries, not merely combine function names.
Small inference workloads can become enqueue-bound: the CPU and driver spend enough time dispatching kernels that launch overhead competes with GPU execution. CUDA Graphs let the runtime capture the sequence once and replay the instantiated graph with one host launch.
TensorRT adds two constraints that shape WarpStream’s architecture. First, a dynamic shape or profile change can trigger deferred work on the next enqueue. The runtime therefore sets the fixed input shape and performs one warm enqueue before capture. Second, graph execution retains context state and buffer addresses. That is why every lane receives its own context and fixed allocations.
// Initialization only
context.setInputShape("images", fixed_shape);
context.setTensorAddress("images", network_input);
context.setTensorAddress("logits", logits);
context.enqueueV3(stream); // flush deferred shape work
cudaStreamSynchronize(stream);
cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal);
preprocess(..., stream);
context.enqueueV3(stream);
topk(..., stream);
cudaStreamEndCapture(stream, &graph);
cudaGraphInstantiateWithFlags(&graph_exec, graph, 0);In steady state, submit() records a start event, launches that graph, and records completion. The graph does not make the model faster; it removes repeated host work around the model and makes the launch path more consistent.
Copying an entire logit vector to the host just to select five entries creates a transfer and a synchronization at the worst possible point: after the GPU is otherwise finished. WarpStream instead launches one 256-thread block and returns only K score/index pairs.
Each thread performs a strided scan over the classes. Candidates reduce inside a warp with __shfl_down_sync; warp winners land in shared memory; a final warp chooses the block winner. Selected indices remain in shared memory while the process repeats for the next rank. Ties prefer the lower class index, and NaNs are treated deterministically as negative infinity.
The algorithm is exact and intentionally optimized for small classification outputs with K ≤ 32. Its work scales with classes × K. A language-model vocabulary or very large label space should use hierarchical selection, radix methods, or a library primitive instead. A bounded specialization is better than pretending one kernel is universally optimal.
The repository’s determinism contract covers memory and execution structure: fixed shapes, fixed addresses, bounded concurrency, explicit transitions, no allocations after initialization, and no host decision between GPU stages.
It does not promise bit-identical floating-point output across GPU architectures, TensorRT versions, or tactic selections. TensorRT chooses implementations by timing candidates during engine construction. Small measurement differences can select different tactics unless the build environment and timing cache are controlled.
A production artifact pipeline should build plans for the target GPU class, persist the timing cache, record the complete software stack, and refuse to compare benchmark runs collected under different clock or power conditions.
The included benchmark records GPU events immediately before graph launch and after Top-K. That interval includes fused preprocessing, TensorRT inference, and postprocessing. A separate host wall clock measures achieved throughput across one or more lanes.
That benchmark intentionally excludes the producer’s write and the consumer’s work. It is a component benchmark, not an end-to-end service-level measurement. The application must add timestamps at capture and delivery if it wants to claim camera-to-decision latency.
A credible report should state sample count, warm-up count, precision, source and model shapes, graph mode, lane count, GPU clocks, temperature, power state, PCIe topology, model hash, TensorRT/CUDA/driver versions, and whether transfers are included. It should also compare --no-cuda-graph so the launch-overhead benefit is visible instead of assumed.
The public repository includes the runtime library, CLI, latency harness, CUDA correctness tests, host configuration tests, CMake integration, a deterministic demo ONNX generator, and a TensorRT engine-building helper.
src/pipeline.cpp | Lane ownership, state transitions, graph capture, submit/wait |
src/preprocess.cu | Fused RGB8 resize, normalization, and layout conversion |
src/topk.cu | Exact shuffle-reduced device-side Top-K |
benchmarks/ | Tail-latency distribution and throughput harness |
tools/ | Demo ONNX generation and fixed-shape TensorRT plan build |
docs/ | Architecture and reproducible benchmarking contract |
The current implementation is a foundation, not a production server. The next useful steps are precision-specialized preprocessing, detection decode and NMS, direct device-pointer registration for GPUDirect producers, NVTX instrumentation, machine-readable benchmark output, and real-GPU TensorRT CI.
Clone github.com/manishklach/warpstream-trt, generate the included demo model, build a fixed-shape plan with TensorRT 11, and compare one-lane graph replay against direct enqueue on your target GPU.
Low-latency inference is often described as a kernel problem. Kernels matter, but the system around them decides whether their speed survives contact with production. Who owns the input memory? When can an address move? Which context carries mutable state? Where does the CPU wait? What happens when all lanes are busy? Which latency interval is actually being reported?
WarpStream’s answer is to make those decisions visible and bounded. It trades generality for an execution rail that can be inspected from the producer pointer to the final Top-K entry. That is a useful baseline even if a production system later chooses copied input, dynamic batching, multiple shapes, or a different postprocessor.
The important part is not copying every choice. It is refusing to let convenience layers hide the path you are trying to optimize.