Bird’s-eye-view (BEV) perception projects multi-camera image features into a shared top-down grid used by detection, occupancy prediction, trajectory and planning modules. BEV pooling turns depth-aware image features into a compact BEV tensor via gather, depth-weighting and scatter-reduce; in deployment this step often becomes a latency bottleneck because of irregular memory access patterns, repeated index reads, scatter-reduce semantics and GPU cache effects.
What BEVPoolV3 does
BEVPoolV3 builds on prior work (BEVPoolV2 and CUDA-BEVFusion’s bevpool_half_pack10_kernel, here called V2+DO) and adds four implementation changes: reduced duplicate depth loads, a five-array INT32 scatter map (ranks_depth, ranks_feat, ranks_bev, interval_starts, interval_lengths), precomputed indices to avoid runtime integer division, and interval-owned output writes that remove atomics. These changes reduce memory traffic and instruction pressure and make the operator portable and performant across different GPU memory regimes.
The optimization workflow
The post presents a repeatable workflow for optimizing BEV pooling and other gather/scatter-heavy operators on NVIDIA GPUs:
- Classify the memory regime by comparing the working set size to the GPU L2 capacity.
- Eliminate redundant scatter and index traffic.
- Map the kernel implementation to the target GPU (occupancy, vector loads, dtype specialization).
- Validate the active bottleneck with NVIDIA Nsight Compute.
The same algorithm requires different kernel choices if the working set is DRAM-bound versus largely L2-resident.
Evaluation setup and key numbers
Two NVIDIA RTX GPUs illustrate different regimes:
- NVIDIA RTX A6000 (Ampere SM86) with 6 MB L2 — the canonical working set (~49 MB) exceeds L2 so the kernel is DRAM-bound.
- NVIDIA RTX PRO 6000 Blackwell Max-Q (Blackwell SM120) with 128 MB L2 and native FP8 — the same 49 MB working set fits in L2 and becomes largely L2-resident after the initial fill.
Canonical config: derived from real nuScenes samples, ~209K scatter points, 80 feature channels, ~49 MB working set.
Measured TensorRT plugin latencies (canonical case, 100-iteration median):
- RTX PRO 6000 Blackwell Max-Q: V2 FP16 = 274.0 µs; BEVPoolV3 FP16 = 17.3 µs; BEVPoolV3 FP8 = 16.4 µs.
- RTX A6000 canonical (DRAM-adapted V3 FP16): 90.0 µs versus V2 FP16 = 1,738.0 µs.
Across multiple tested configurations on RTX PRO 6000 Blackwell Max-Q, V3 FP8 shows 10.94×–42.09× speedup over V2 FP16 depending on point count and channel width.
Technical approach
-
Reducing redundant index reads: BEVPoolV2’s tile-outer loop caused repeated loads of the same scatter indices when iterating channel tiles (e.g., C=80 with 8-channel tiles results in 10 repeated loads). A depth-outer or interval-based traversal removes that repetition.
-
Five-array scatter map: using separate INT32 arrays for ranks_depth, ranks_feat and ranks_bev plus interval_starts and interval_lengths avoids awkward 12-byte packed records and enables aligned, vectorized loads (cleaner instruction stream), which is especially beneficial when the working set lives in L2.
-
Interval-owned scatter-reduce: precompute the scatter map; assign each BEV interval to a single owner thread/block; the owner iterates points in the interval, loads depth once per point, accumulates feature channels locally, then writes the output once. This removes runtime index decoding and atomic output writes.
Code sketch (conceptual):
for each interval iv in parallel: start = interval_starts[iv] length = interval_lengths[iv] bev = ranks_bev[start]
acc[channel_tile] = 0
for offset in 0 .. length - 1:
t = start + offset
d = depth[ranks_depth[t]]
feat_row = ranks_feat[t]
for c in channel_tile:
acc[c] += d * feat[feat_row, c]
out[bev, channel_tile] = acc
Production kernels specialize this invariant for the active memory regime: on small-L2 GPUs (RTX A6000) the focus is on byte reduction (increase TILE_C, use __half2 accumulation, cache-streaming stores); on large-L2 GPUs (RTX PRO 6000 Blackwell Max-Q) the focus is on occupancy, precomputed indices, vectorized index loads and FP8 specialization.
Precision and dtype considerations
When the feature and output data are L2-resident, switching to FP8 can yield significant latency decreases. The authors tested an NVFP4 approach (camera features in E2M1 with per-16-element E4M3 microblock scales, depth and output in FP8), but decoding overhead made NVFP4 slower than the direct FP8 baseline for this scatter-reduce workload. Nsight Compute showed the kernel fully resident in L2 with low DRAM utilization and instruction-issue being the limiting factor, indicating FP8 provides the best dtype tradeoff for L2-resident scatter-reduce shapes.
Validation
BEVPoolV3 is exposed as a TensorRT IPluginV3. The plugin accepts the five-array scatter map plus depth and feat and dispatches a GPU- and dtype-appropriate kernel. Validation against an FP64 reference or the V2 path showed numerical agreement: RTX A6000 DRAM-adapted kernel passed tests at atol=1e-2 (max observed error 0.0065), and on RTX PRO 6000 Blackwell Max-Q V2 and V3 produced identical outputs for the tested cases.
Applicability beyond BEV pooling and edge considerations
The same workflow applies to sparse embeddings, voxelization, histograms, segmented reductions and other irregular memory-bound kernels. On edge-class platforms (e.g., NVIDIA DRIVE AGX Thor) the core BEVPoolV3 improvements (remove redundant scatter traffic, avoid runtime index decoding, interval-owned writes) transfer well, but FP8 gains are architecture- and kernel-specific: smaller problem sizes, cache behavior and register pressure on edge targets can reduce or negate FP8 benefits.
Practical starting steps
- Profile the BEV pooling operator in isolation; measure feature, depth, scatter-index and output sizes.
- Compare the total working set to target GPU L2 capacity.
- Use NVIDIA Nsight Compute to identify whether bandwidth, instruction issue or occupancy is the active ceiling.
- Apply byte-reduction and cache-preserving stores for DRAM-bound cases; apply occupancy, precomputed indices, vector loads and dtype specialization for L2-resident cases.
Conclusion
BEVPoolV3 shows how memory-regime aware optimizations and a small set of algorithmic invariants (own the interval, remove runtime index decoding, accumulate locally, write once) can produce large latency reductions for BEV pooling. The appropriate implementation choices depend primarily on whether the working set fits in L2 or is DRAM-bound; profiling and Nsight Compute validation are essential steps before deployment via TensorRT.
Mentioned tools and references
NVIDIA Nsight Compute, TensorRT (IPluginV3), CUDA-BEVFusion, CUDA C++ Programming Guide, TensorRT documentation, Nsight Compute Profiling Guide and NVIDIA Developer Forums are cited as relevant resources.
