An more and more widespread design sample for autonomous autos (AVs), robotics, and spatial AI techniques is fowl\u2019s-eye-view (BEV) notion. BEV fashions venture multicamera picture options right into a shared top-down grid, offering downstream notion and planning modules with a standard spatial format for reasoning about lanes, autos, pedestrians, and free house.<\/p>\n
A key operation on this pipeline is BEV pooling, which gathers picture options, weights them with depth data, and scatter-reduces them into BEV grid cells. For builders, the sensible worth of BEV notion is that it converts many camera-specific views into one spatially constant illustration of the scene. As an alternative of reasoning individually over every digicam picture, downstream modules can function on a unified top-down function map aligned to the world across the automobile or robotic. BEV pooling is the step that makes this illustration usable in actual time: it turns depth-aware picture options right into a compact BEV tensor that may feed detection, occupancy, trajectory prediction, mapping, and planning workloads.\u00a0<\/p>\n
Conceptually, that is easy. In deployment, nonetheless, BEV pooling can turn into a latency bottleneck as a result of it combines irregular reminiscence entry, repeated index reads, scatter-reduce conduct, and GPU-specific cache results.<\/p>\n
This put up makes use of BEVPoolV3 as a case examine in optimizing BEV pooling and different gather- or scatter-heavy operators for NVIDIA GPUs. It walks by way of a sensible workflow you possibly can apply to your workloads: classify the reminiscence regime, take away redundant scatter visitors, map the kernel implementation to the goal GPU, and validate the lively bottleneck with NVIDIA Nsight Compute. The efficiency outcomes present why this workflow issues: the identical BEV pooling operator can require completely different optimization methods relying on whether or not the working set is DRAM-bound or largely L2-resident.<\/p>\n
How does BEVPoolV3 cut back BEV pooling latency on NVIDIA RTX GPUs?\u00a0<\/h2>\n
Prior work has already made necessary progress. BEVPoolV2, known as V2 on this put up, launched an environment friendly deployment-oriented BEV pooling formulation for BEVDet-style fashions. CUDA-BEVFusion contains bevpool_half_pack10_kernel, referred to right here as V2+DO, which makes use of depth-outer traversal to take away a lot of the V2 repeated tile-outer index loading.\u00a0<\/p>\n
BEVPoolV3 continues this optimization route with 4 further adjustments: lowered duplicate depth masses, a five-array INT32 scatter map, precomputed indices that take away runtime integer division, and interval-owned output writes.<\/p>\n
This put up makes use of BEVPoolV3 as a case examine in the right way to optimize BEV pooling and different gather- or scatter-heavy operators for NVIDIA GPUs. You’ll learn to classify a BEV pooling workload by reminiscence regime, determine redundant scatter visitors, map the kernel implementation to the goal GPU, and validate the lively bottleneck with Nsight Compute. The efficiency outcomes on two NVIDIA RTX GPUs present why this workflow issues: the identical BEV pooling algorithm will be DRAM-bound on one GPU and largely L2-resident on one other, requiring completely different optimization selections.\u00a0<\/p>\n
The analysis compares two NVIDIA RTX GPUs that signify completely different reminiscence regimes: NVIDIA RTX A6000, an NVIDIA Ampere SM86 GPU with a 6 MB L2 cache and no native FP8 ISA, and NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Version, an NVIDIA Blackwell SM120 GPU with a 128 MB L2 cache and native FP8 help. The canonical config used right here is derived from actual nuScenes samples and accommodates about 209K scatter factors, 80 function channels, and a 49 MB BEV pooling working set. That working set exceeds RTX A6000 L2 cache however matches inside RTX PRO 6000 Blackwell Max-Q L2 cache, making RTX A6000 DRAM-bound and RTX PRO 6000 Blackwell Max-Q largely L2-resident after the preliminary fill.<\/p>\n
![\"Four-panel](\"https:\/\/developer-blogs.nvidia.com\/wp-content\/uploads\/2026\/06\/bev-pooling-convert-multicamera-image-features-1.webp\")
Within the canonical config, the V2-style NVIDIA TensorRT plugin path takes 274.0 \u00b5s on RTX PRO 6000 Blackwell Max-Q. BEVPoolV3 reduces that to 17.3 \u00b5s in FP16 and 16.4 \u00b5s in FP8. On RTX A6000, the DRAM-adapted BEVPoolV3 FP16 path reaches 90.0 \u00b5s. Past the speedup, this put up exhibits a repeatable workflow for optimizing scatter-reduce kernels: classify the working set, take away redundant reminiscence visitors, match the launch form to the goal GPU, and validate the outcome with Nsight Compute.\u00a0<\/p>\n
![\"Bar](\"https:\/\/developer-blogs.nvidia.com\/wp-content\/uploads\/2026\/06\/canonical-speedup-bevpoolv3-1.webp\")
Stipulations\u00a0<\/h2>\n
This put up discusses CUDA kernel conduct, TensorRT plugin integration, and GPU profiling within the context of BEV pooling. Useful conditions embody:<\/p>\n
CUDA kernel ideas akin to warp scheduling, atomics, vectorized world masses, and DRAM\/L2\/L1 cache conduct<\/p>\n
TensorRT plugin integration, particularly the IPluginV3 interface<\/p>\n
Nsight Compute profiling for validating reminiscence conduct, occupancy, and instruction-issue bottlenecks<\/p>\n
The BEV-pooling kernel in CUDA-BEVFusion because the prior depth-outer reference implementation<\/p>\n
For associated background data, see the CUDA C++ Programming Information, TensorRT plugin documentation, TensorRT samples, and Nsight Compute Profiling Information.<\/p>\n
Classify the reminiscence regime<\/h2>\n
Step one is to categorise whether or not the BEV-pooling working set matches in L2. Within the canonical config, the principle arrays complete about 49 MB, dominated by function information and output. That single quantity determines the reminiscence regime: it’s bigger than the RTX A6000 6 MB L2 cache, however smaller than RTX PRO 6000 Blackwell Max-Q 128 MB L2 cache.<\/p>\n This match\/no-fit resolution adjustments the optimization goal. On RTX A6000, function gathers and output visitors spill past L2, so the small-L2 path prioritizes byte discount and cache-streaming output shops. On RTX PRO 6000 Blackwell Max-Q, the canonical working set matches in L2, so the large-L2 path shifts towards instruction effectivity, occupancy, precomputed indices, vectorized masses, and FP8 specialization.<\/p>\n The BEV scatter-reduce will be summarized as:<\/p>\n BEVPoolV2 iterates over channel tiles outdoors the scatter loop. For C=80 and an 8-channel tile, the identical scatter indices are loaded 10 instances. That produces roughly 25.1 MB of index visitors for indices that solely want 2.51 MB when learn as soon as. A depth-outer loop order fixes most of that drawback by iterating over every BEV interval first and accumulating all channels for that interval in a single move.<\/p>\n BEVPoolV3 extends the depth-outer optimization route utilized in CUDA-BEVFusion bevpool_half_pack10_kernel, referred to right here as V2+DO. V2+DO is a helpful baseline as a result of it already removes the repeated tile-outer index masses in BEVPoolV2 and demonstrates the worth of interval-based traversal. BEVPoolV3 retains that route and provides 4 implementation adjustments that enhance portability and efficiency throughout GPU reminiscence regimes: lowered duplicate depth masses inside every interval; a five-array INT32 scatter map \u00b5sing ranks_depth, ranks_feat, ranks_bev, interval_starts, and interval_lengths; precomputed specific indices that take away runtime integer division; and interval-owned output writes that keep away from atomics relative to the V2-style path.<\/p>\n The five-array scatter map is very necessary on large-L2 GPUs. Packing (ranks_depth, ranks_feat, ranks_bev) into an int3 array offers a 12-byte report. That format is inconvenient for aligned reminiscence transactions and doesn’t map cleanly to a 16-byte LDG.128 load. Separate INT32 arrays let adjoining threads merge aligned masses and keep away from area coupling. The whole logical bytes might look comparable, however the instruction stream is way cleaner.<\/p>\n In manufacturing, BEVPoolV3 makes use of a number of specialised kernels, however the core implementation concept is less complicated to grasp as a small logic sketch. The scatter map is ready forward of time, every BEV interval is assigned to 1 proprietor, the proprietor walks the factors in that interval, accumulates the related function channels, and writes the output as soon as.<\/p>\n This construction removes the inner-loop decoding work that seems when the scatter map is packed right into a single report. As an alternative of reconstructing indices at runtime, the kernel reads specific arrays akin to ranks_depth, ranks_feat, ranks_bev, interval_starts, and interval_lengths.<\/p>\n for every interval iv in parallel: acc[channel_tile] = 0<\/p>\n for offset in 0 .. size – 1: for c in channel_tile: out[bev, channel_tile] = acc\n<\/p><\/div>\n This code sketch captures the widespread BEVPoolV3 construction: the scatter map is specific, runtime index decoding is eliminated, depth is loaded within the interval proprietor loop, and every output cell is written as soon as after native accumulation.<\/p>\n The manufacturing kernels specialize this construction for the goal reminiscence regime. On small-L2 GPUs akin to RTX A6000, the implementation prioritizes byte discount, FP16 half2 accumulation, and cache-streaming output shops so the output tensor doesn’t evict helpful index information from L2. On large-L2 GPUs akin to RTX PRO 6000 Blackwell Max-Q, the implementation first matches a high-occupancy launch envelope, then reduces instruction overhead with precomputed indices, vectorized index masses, and FP8-specialized interior loops the place the working set is L2-resident.<\/p>\n The algorithmic invariant stays the identical: personal the interval, keep away from runtime index decoding, accumulate regionally, and write as soon as. The architecture-specific work adjustments how that invariant is carried out, not what the BEV-pooling operator computes.<\/p>\n Absolutely the latency outcomes on RTX PRO 6000 Blackwell Max-Q present how the large-L2 path behaves throughout completely different level counts and channel widths. The identical optimization sample additionally holds on the RTX A6000 DRAM-bound path when measured as speedup over the V2 FP16 baseline. On RTX A6000, the DRAM-adapted V3 FP16 path reaches speedups of 11s to 22x over V2 throughout the examined configurations. On RTX PRO 6000 Blackwell Max-Q, V3 FP8 reaches speedups of 11x to 42x over V2, with the biggest positive aspects showing at bigger level counts and wider channel configurations.<\/p>\n BEVPoolV3 is uncovered as a TensorRT IPluginV3 operator. The plugin accepts the five-array scatter map plus depth and feat, then dispatches the suitable kernel for the GPU class and dtype. The benchmark path used ONNX-to-TensorRT builds and CUDA Graph replay with trtexec.<\/p>\n For validation, evaluate in opposition to an FP64 reference or an current trusted V2 path. The RTX A6000 DRAM-adapted kernel handed all examined output parts throughout the six configurations at atol=1e-2, with most noticed error of 0.0065. On RTX PRO 6000 Blackwell Max-Q, V2 and V3 produced similar outputs for the examined configurations, indicating that the optimized scatter-map and launch adjustments preserved the numerical conduct of the reference path.<\/p>\n The 4 BEVPoolV3 algorithmic adjustments are moveable, however the manufacturing kernel should match the lively GPU bottleneck. The important thing resolution is whether or not the BEV-pooling working set matches in L2.<\/p>\n On RTX A6000, the canonical working set exceeds L2, so the kernel is restricted by random-gather DRAM visitors. The FP16 path subsequently prioritizes byte discount and cache preservation. Rising TILE_C from 8 to 16 cuts the C=80 tile passes from 10 to five, lowering loop overhead and repeated scalar work. Utilizing __half2 accumulation with __hfma2 avoids pointless FP16-to-FP32 widening and packing. Cache-streaming output shops forestall the 12.8 MB output tensor from evicting the smaller L2-resident index arrays. After these adjustments, the RTX A6000 path reaches 90.0 \u00b5s within the canonical config, in contrast with 1,738.0 \u00b5s for V2 FP16.<\/p>\n On RTX PRO 6000 Blackwell Max-Q, the canonical working set matches in L2, so the limiting components shift towards instruction subject, occupancy, and dependency latency. The manufacturing kernel first matches the high-occupancy V2+DO-style launch envelope, then removes inner-loop overhead with the five-array scatter map and precomputed indices. This avoids runtime integer division and reduces scatter-map load strain. Within the canonical config, V3 FP16 reaches 17.3 \u00b5s versus 37.8 \u00b5s for V2+DO FP16, a 2.18x speedup on the similar dtype.<\/p>\n The FP8 path additional specializes within the large-L2 case. As a result of function and output information are served from L2, lowering their dtype can translate into actual latency positive aspects. The manufacturing FP8 path makes use of per-channel-count entry factors, LDG.64 index packing for C=80, and wider function masses for C=128 and C=256. Extra aggressive combos, akin to including loop unrolling on high of the packed-index path, didn’t compose cleanly as a result of they elevated register strain and spill visitors.<\/p>\n The precision ladder has a sensible vacation spot, and our NVFP4 analysis helps make clear precisely the place every format shines: we examined an NVFP4 path that shops digicam options in E2M1 with per-16-element E4M3 microblock scales whereas conserving depth and output in FP8, and even with an aggressively optimized implementation that includes __half2 packed accumulators, fused scale\u2013depth coefficients, and a half-precision LUT, the decode overhead causes it to run notably slower than the FP8 baseline.\u00a0<\/p>\n Profiling with Nsight Compute exhibits the kernel is totally resident in L2 cache, with low DRAM bandwidth utilization and smsp__issue_active hovering properly under peak throughput, whereas the ALU pipeline carries considerably extra directions than the FMA pipeline.\u00a0<\/p>\n This means that this scatter-reduce regime has already captured the out there byte-efficiency advantages at FP8, whereas the NVFP4 further per-element nibble extraction, worth decode, and per-microblock scale fold introduce inner-loop work that the FP8 path avoids by way of a single scalar FP8 to half conversion. The result’s a crisp workload-placement story: NVFP4 stays an extremely highly effective match for compute-bound matrix multiplication shapes flowing by way of Tensor Cores by way of MMA.type::nvfp4, whereas for L2-resident scatter-reduce workloads, FP8 is good on the dtype ladder.<\/p>\n The identical evaluation applies past BEV pooling. For sparse embeddings, voxelization, histograms, segmented reductions, and different gather- or scatter-heavy operators, first classify the reminiscence regime, then use Nsight Compute to find out whether or not the lively ceiling is bandwidth, instruction subject, or occupancy.<\/p>\n Desk 1 summarizes RTX PRO 6000 Blackwell Max-Q TensorRT plugin-path latency, reported as 100-iteration median latency.<\/p>\n The identical evaluation can prolong to edge-class NVIDIA platforms, together with NVIDIA DRIVE AGX Thor. In early edge-oriented experiments, the FP16 BEVPoolV3 path carries over properly as a result of the core enhancements\u2014eradicating redundant scatter visitors, avoiding runtime index decoding, and utilizing interval-owned writes\u2014are architecture-independent.<\/p>\n FP8 speedup, nonetheless, is just not automated. On edge-class targets, smaller drawback sizes, reminiscence hierarchy conduct, register strain, and FP8 conversion overhead can restrict or offset the theoretical dtype bandwidth profit. This makes FP8 a kernel- and architecture-specific optimization fairly than a assured drop-in alternative for FP16.<\/p>\n To use the BEVPoolV3 workflow to your personal BEV notion or collect\/scatter-heavy workload, begin by profiling the operator in isolation. Measure the function, depth, scatter-index, and output tensor sizes, then evaluate the entire working set with the goal GPU L2 cache capability.<\/p>\n Use NVIDIA Nsight Compute to validate whether or not the lively bottleneck is reminiscence bandwidth, instruction subject, occupancy, or dependency latency. Then select the optimization technique that matches the reminiscence regime: byte discount and cache-preserving shops for DRAM-bound workloads, or occupancy, precomputed indices, vectorized masses, and dtype specialization for L2-resident workloads.<\/p>\n The identical strategy applies to sparse embeddings, voxelization, histograms, segmented reductions, and different irregular memory-bound kernels. Use the BEVPoolV3 outcomes as a information for profiling your personal operator, choosing the precise implementation technique for the goal GPU, and validating the outcome earlier than deploying by way of TensorRT. For associated assets, see the TensorRT documentation, CUDA C++ Programming Information, Nsight Compute documentation, and NVIDIA Developer Boards.<\/p>\n<\/div>\n![\"Two](\"https:\/\/developer-blogs.nvidia.com\/wp-content\/uploads\/2026\/06\/canonical-bev-pooling-working-set-1.webp\")
Take away redundant scatter visitors<\/h2>\n
![\"Stacked](\"https:\/\/developer-blogs.nvidia.com\/wp-content\/uploads\/2026\/06\/analytical-memory-traffic-chart.webp\")
Implement interval-owned scatter-reduce<\/h2>\n
\n\/\/ 2. Learn specific indices instantly, with no runtime index division.
\n\/\/ 3. Let one interval proprietor accumulate the output cell.
\n\/\/ 4. Load every depth worth as soon as per scatter level within the proprietor loop.<\/p>\n
\n begin = interval_starts[iv]
\n size = interval_lengths[iv]
\n bev = ranks_bev[start]<\/p>\n
\n t = begin + offset
\n d = depth[ranks_depth[t]]
\n feat_row = ranks_feat[t]<\/p>\n
\n acc[c] += d * feat[feat_row, c]<\/p>\n![\"Grouped](\"https:\/\/developer-blogs.nvidia.com\/wp-content\/uploads\/2026\/06\/latency-analysis-tensorrt.webp\")
![\"Grouped](\"https:\/\/developer-blogs.nvidia.com\/wp-content\/uploads\/2026\/06\/cross-gpu-speedup-over-v2-fp16-1.webp\")
Deploy and validate the TensorRT plugin<\/h2>\n
Map the algorithm onto the {hardware}<\/h2>\n
Concerns for edge-class platforms<\/h2>\n
Get began with BEV pooling optimization<\/h2>\n