perf(rfdetr-seg): Triton fusion + pipelining cuts post/pre-proc bubbles (93 → 122 FPS, +31.5%)

[aseembits93]

Summary

Metric	Baseline	Final	Δ
Avg FPS (Tesla T4, example_video.mp4, 431 frames, 5-run avg)	93.07	122.40	+31.5%
Post-process bubble ratio (nsys)	83%	24%	-59 pts
GPU kernels per run	133,821	~7,000	-94%
DtoD/PtoP memcpys per run	439	8	-98%
Postproc kernel cub/torch ops per frame	~45	~3	-93%
Parity vs baseline	—	4/431 frames ±1 detection	within 5% tolerance
Unit tests	—	11/11 pass	✓

Flags to enable (all default off):

• RFDETR_USE_TRITON_PREPROC=true — fused Triton preprocessing kernel
• RFDETR_TRITON_FULLPOSTPROC=true — fused Triton post-processing kernels
• ENABLE_AUTO_CUDA_GRAPHS_FOR_TRT_BACKEND=true — enable CUDA-graph replay for TRT forward

---

Commit-by-commit changes

Each row shows the 5-run avg FPS after the commit lands, and the delta over the prior commit.

#	Commit	Avg FPS	Δ over prev	Δ vs baseline
0	baseline (before this PR)	93.07	—	—
1	fuse ENTIRE post-process into 2 Triton kernels	106.05	+13.9%	+13.9%
2	adapter-level refinements on top of full-fusion	107.94	+1.8%	+15.9%
3	kill bubbles in post-process (+10.3% FPS on top of v3)	119.00	+10.2%	+27.9%
4	remove CPU-side sync between filter and mask Triton kernels	119.75	+0.6%	+28.7%
5	drop unnecessary .bool() casts in postproc return path	120.37	+0.5%	+29.3%
6	batch scalar DtoH into single 72B combined transfer	120.49	+0.1%	+29.5%
7	pinned host buffers + async DtoH for final transfers	121.21	+0.6%	+30.2%
8	eliminate preproc->TRT DtoD copy + pinned HtoD	122.40	+1.0%	+31.5%

---

1. Fuse ENTIRE post-process into 2 Triton kernels (+13.9%)

Replaces the whole post-TRT chain (sigmoid + argmax + class-remap + conf threshold + xywh→xyxy + denorm + pad subtract + scale divide + clip + mask bilinear upsample + threshold > 0) with two Triton kernel launches.

• _rfdetr_fullpost_filter_kernel — grid = num_queries. Per query: fused filter + box denorm + clip. Writes padded fixed-shape outputs (keep marks filtered rows).
• _rfdetr_fullpost_mask_kernel_compact — grid = n_survivors × tile_y × tile_x. Per tile: inverse-map orig→mask coords, bilinear upsample, threshold > 0, store uint8. Uses survivor_idx so only filtered masks get processed (30× less compute than the fixed-padded variant).

Key insight: "fuse into fixed-shape padded output" only wins when the kernel cost scales with inputs (filter step = 100 queries × 91 classes, constant). Per-output-pixel kernels need explicit compaction ("zero out filtered rows" is a trap when 95% of output pixels are wasted work).

2. Adapter-level refinements on top of full-fusion (+1.8%)

Three eliminations of downstream kernels via in-kernel work:

• Filter kernel rounds + casts xyxy to int32 internally → removes 431 round_kernel_cuda per run
• Mask kernel accumulates per-survivor mask_any via tl.atomic_max → removes 314 or_kernel_cuda reduces per run
• Compact (n_survivors)-shape outputs → removes per-frame torch.zeros fill (20 ms) + 4 [keep] indexing kernels per frame

Kernels dropped: 10,900 → 4,669 (-57%). GPU postproc time: 60 ms → 17 ms (-72%).

3. Kill bubbles in post-process (+10.2%)

Biggest single jump. Five eliminations combined:

1. Atomic-counter compaction in filter kernel — each surviving query does slot = tl.atomic_add(counter_ptr, 1) and writes outputs to compact[slot], skipping filtered queries entirely. Absorbs 13 downstream kernels (keep.bool, the 3-cub nonzero() pipeline, 3 index_select gathers, torch.zeros, .bool() cast).
2. Removed TRT output-buffer clone in trt.py:711. Graph replay writes to fixed output buffers; instead of 3× DtoD memcpy per frame to protect readers, replaced with a consumer_done_event on TRTCudaGraphState that the next replay waits on.
3. Removed outer stream.synchronize() in infer_from_trt_engine. Was blocking CPU after every forward. Replaced with a cross-stream cuda.Event that postproc waits on from its own stream.
4. Cached threshold tensor — was allocating torch.tensor([thr]) (4 B HtoD) per frame. Now cached per (value, device).
5. Simplified adapter — removed the GPU-side .any() + masked_select + nonempty_idx DtoH/HtoD roundtrip. Since det.mask from the Triton path is already compact and masks2poly fast-skips empty masks, CPU-side filter is adequate.

Postproc bubble ratio: 83% → 35%.

4. Remove CPU-side sync between filter and mask Triton kernels (+0.6%)

The filter→mask bubble was 146 µs per frame caused by int(counter.item()) blocking the host between the two Triton launches. Fix: launch the mask kernel with grid=(num_queries, tile_y, tile_x) and have each program read counter[0] on the GPU and early-exit if its s index exceeds n_survivors. Host reads counter only after the mask kernel is queued, so the two kernels pipeline on the stream. Tradeoff: mask kernel runtime grew 27 → 112 µs (scheduler iterates ~100 programs, most early-exit), net gain ~61 µs/frame.

Also cached mask_bin at max (num_queries, H, W) on module scope to skip per-frame torch.empty() launch.

5. Drop unnecessary .bool() casts in postproc return path (+0.5%)

Two torch.to(torch.bool) kernels per frame from dead conversions (common.py:125 and the wrapper return). Removed, no functional impact. ~1000 kernel launches deleted per run.

6. Batch scalar DtoH into single 72B combined transfer (+0.1%)

Four separate small .cpu() calls in the adapter (4B counter + 48B xyxy + 12B conf + 12B class_id), each with ~25-40 µs Python+sync overhead. Collapsed to one transfer: filter kernel now writes a single combined (num_queries, 6) int32 buffer packing [x1, y1, x2, y2, conf_as_i32_bits, class_id] per slot. Adapter does one .cpu() and bitcasts column 4 to fp32 via numpy.view(np.float32).

7. Pinned host buffers + async DtoH for final transfers (+0.6%)

Replace the adapter's two synchronous .cpu() calls with pinned-memory .copy_(non_blocking=True) + a single stream.synchronize(). Both the 72B combined buffer and the 307KB mask buffer now pipeline on the copy engine in parallel rather than chaining. Pinned buffers cached per (name, dtype) and reused across frames.

DtoH timing: 49 µs gap between scalar + mask transfers → 15 µs.

8. Eliminate preproc → TRT DtoD copy + pinned HtoD (+1.0%)

Two preprocess-side fixes:

Eliminate the per-frame DtoD from _fast_input_buffer → graph_input_buffer (1,168,128 B per frame = 508 MB per run). Threaded a _trt_reuse_as_input_buffer tensor-attribute hint through _capture_cuda_graph. When set, the graph bakes the external tensor's address as its input buffer. On replay, skip the DtoD when data_ptr matches. Result: DtoD/PtoP memcpys per run 439 → 8 (-98%).

Pinned-memory HtoD for the raw BGR frame. Was using torch.from_numpy(ascontiguousarray(img)).to(device, non_blocking=True) — but without pinned source, non_blocking=True silently falls back to sync. Now uses a cached pinned host buffer + GPU buffer, copies the numpy frame into pinned with np.copyto, then tensor.copy_(pinned, non_blocking=True) for a genuine async HtoD.

---

Remaining bubbles (future work)

From the final v15 profile:

• ~175 µs Python interpreter overhead between _rfdetr_stretch_preprocess_kernel completing and cudaGraphLaunch — pure call-chain cost across pre_process → infer → forward → infer_from_trt_engine → 4 more wrapper layers → cuda_graph.replay. Eliminable via call-chain flattening.
• ~200 µs Python overhead in the adapter between the mask kernel ending and the first .cpu() call. Same category — CPython cost, not GPU.
• CPU-GPU pipelining: currently serial per frame. GPU is idle ~50% of wall time waiting for CPU plumbing. A ping-pong double-buffer scheme could reclaim another +40-55% FPS (estimated ceiling ~170-190 FPS).

Files changed

File	Purpose
inference_models/models/rfdetr/triton_fullpostproc.py	New. Two Triton kernels + host wrapper for the fused post-process.
inference_models/models/rfdetr/triton_preprocess.py	Existing (PR #22). Used by fast-path preproc.
inference_models/models/rfdetr/common.py	Wires RFDETR_TRITON_FULLPOSTPROC into post_process_instance_segmentation_results.
inference_models/models/rfdetr/rfdetr_instance_segmentation_trt.py	Pinned HtoD, buffer reuse, _trt_reuse_as_input_buffer marker, event-based cross-stream ordering.
inference_models/models/common/trt.py	_trt_reuse_as_input_buffer support, removed output-buffer clone, cross-stream event plumbing.
inference/core/models/inference_models_adapters.py	Batched combined-buffer DtoH, pinned async DtoH, simplified mask filter.

Test plan

• pytest tests/inference/unit_tests/models/test_rfdetr.py -x -q — 11/11 pass
• Benchmark: 5-run avg 122.40 ± 0.9 FPS
• Parity: 4/431 frames differ by ±1 detection vs baseline (same 4 frames as PR perf(rfdetr-seg): Triton fusion + CUDA graphs + scratch caching (109 → 151 FPS on vehicles_312px) #22; attributable to Triton bilinear vs torchvision antialias at mask edges)
• nsys profiles saved at /tmp/nsys_profiles/rfdetr_trt_fullpost_v{3..15}.nsys-rep
• Test on higher-res video (mask kernel compute scales with orig_h × orig_w)
• Test on Jetson Orin (pinned-memory + async-transfer wins should scale better there)

[aseembits93]

Closing in favor of #22. Both PRs cover the same optimization arc on the same files; #22's branch (optimize-rfdetr-seg) is strictly ahead of #23's branch (experiment/full-postproc-triton) — every line in #23 is in #22, plus the full Triton postproc fusion, deferred counter sync (W2), and scratch caching (W7) on top. Current benchmark on vehicles_312px.mp4 (538 frames, Tesla T4 FP16): 109 FPS → 151 FPS (+38%), 0-diff parity vs v16 best across all 538 frames.