Dense rasterizers cannot render a crop into crop-sized buffers: window origin and tile-grid contract are inconsistent

[swahtz]

Summary

The dense Gaussian rasterizers accept a render window (image_width, image_height, image_origin_w, image_origin_h) but the window contract is incoherent, so there is no way to rasterize a crop of an image into crop-sized buffers. Callers that want crops must render the full image and slice, which keeps full-resolution forward and backward buffers alive. fvdb-reality-capture's crops_per_image > 1 training mode exists to reduce memory and currently cannot (openvdb/fvdb-reality-capture#338).

What the kernel does today

In src/fvdb/detail/utils/gsplat/GaussianRasterize.cuh, RasterizeCommonArgs:

• denseCoordinates() adds the window's tile origin to the block index (tileRow = blockIdx.y + mTileOriginH) and looks that absolute tile up in tile_offsets. This is correct only if tile_offsets covers the full image.
• checkInputShapes() requires tile_offsets to have the window's tile extent (tileOffsets.size(2) == renderWindow.tileExtentW(tileSize)), i.e. a grid sized to the crop. With a nonzero origin the absolute tile index then runs past the grid.
• The pixel coordinates it derives (row = tileRow * tileSize + threadIdx.y) are absolute, while pixelIndex() treats them as window-relative (row * mRenderWindow.width + col) for output, mask and background indexing. The comment above them says they are meant to be local to the crop.

The same applies to the screen-space and world-space forward and backward kernels. No gtest exercises a nonzero origin.

Reproduction

On the gsplat/test_garden_cropped.npz scene, 128x96 crop at origin (96, 64), tile size 16, via fvdb.functional:

Call	Result
Crop-sized tile grid (intersect_gaussian_tiles with the crop's tile counts) + rasterize_screen_space_gaussians_fwd with the origin	Runs, but max abs error vs. the slice of a full render is 0.80 for a mean pixel value of 0.17
Full-image tile grid + rasterize_screen_space_gaussians_fwd with the origin	ValueError: tileOffsets width must match the number of tiles in image size
Full render, then slice	Correct; the slice retains 22x its logical storage (816480 vs 36864 floats)

The first row is what fvdb-reality-capture did before openvdb/fvdb-reality-capture#338; it produced wrong pixels for any crop not at the origin. The third row is what it does now.

Proposed contract

A window renders the region [origin, origin + size) of an image whose tile intersections were computed once over the full image:

1. tile_offsets is the full-image grid [C, ceil(H / tile_size), ceil(W / tile_size)]. checkInputShapes validates that the window's tile origin plus its tile extent fits inside that grid (and, in packed mode, the analogous bound).
2. Gaussian evaluation uses absolute pixel coordinates; output, masks and backgrounds indexing uses window-relative coordinates, so outputs are allocated at [C, window_h, window_w, ...].
3. The backward kernels mirror this, including last_ids and the saved alphas being window-sized.
4. Add gtests for both dense rasterizers with a nonzero origin that compare against the slice of a full render, and a Python contract test in tests/unit/test_gaussian_splat_functional.py.

With that in place, fvdb-reality-capture can intersect once per image and rasterize each crop as a true window; its crop argument on the stage functions becomes that window instead of a slice.

Out of scope

The sparse rasterizers take explicit pixels and do not use the window origin; nothing changes there. Tile intersection is unaffected.

Related

• Publish Gaussian splatting kernels as a flat public surface in fvdb.functional #797, Publish Gaussian splatting kernels as a flat public surface in fvdb.functional #799 published the kernels as fvdb.functional; the window contract is unchanged by them.
• Composable differentiable Gaussian splatting API over fvdb.functional fvdb-reality-capture#337, Composable differentiable Gaussian splatting API over fvdb.functional fvdb-reality-capture#338 compose the pipeline downstream and document the full-size allocation on crops.