`CUDA invalid configuration` if no Gaussian is in view

[Xanthorapedia]

A CUDA invalid configuration exception will be thrown from kernel launches that perform per-point processing if the number of points is 0, which can happen if no Gaussian is in the view.

This is caused by kernel launches with 0 blocks, e.g. in the example below, blocks evaluates to 0 if num_pts is 0. https://github.com/nerfstudio-project/gsplat/blob/7822c32cdc00caf1a1f0ad8543bcf9c1233deb5e/gsplat/cuda/csrc/bindings.cu#L50-L52

I would suggest checking num_pts before the kernel launch, i.e.,

if (num_pts > 0) {
    compute_cov2d_bounds_kernel<<<(num_pts + N_THREADS - 1) / N_THREADS, N_THREADS>>>(
        num_pts,
        covs2d.contiguous().data_ptr<float>(),
        conics.contiguous().data_ptr<float>(),
        radii.contiguous().data_ptr<float>()
    );
}

I'm happy to submit a PR and please let me know if there are better alternatives.

[kerrj]

That would be great, thanks! if there are no gaussians in view it seems reasonable to just return a torch.full() with dimensions HxWxC with the background color.

[Xanthorapedia]

As a related issue, access to tile_bins in the kernel below sometimes also throws OOM at the line:

https://github.com/nerfstudio-project/gsplat/blob/bae873819fe2d50b4bf097117842d15f17d3a572/gsplat/cuda/csrc/forward.cu#L286

The reason is that the size of tile_bins returned by get_tile_bin_edges_tensor is set to the number of intersections instead of the actual number of tiles, which might be very low or even 0:

https://github.com/nerfstudio-project/gsplat/blob/bae873819fe2d50b4bf097117842d15f17d3a572/gsplat/cuda/csrc/bindings.cu#L313-L319

I would propose to fix this by changing

https://github.com/nerfstudio-project/gsplat/blob/bae873819fe2d50b4bf097117842d15f17d3a572/gsplat/bin_and_sort_gaussians.py#L110

to

    num_tiles = tile_bounds[0] * tile_bounds[1]
    tile_bins = get_tile_bin_edges(num_tiles, isect_ids_sorted)

It might also be helpful to return the tile_bins with shape (tile_bounds[0], tile_bounds[1]) or just tile_bounds instead of (tile_bounds[0] * tile_bounds[1],).

Please let me know how you like it. Thanks!