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.

gsplat/gsplat/cuda/csrc/bindings.cu

Lines 50 to 52 in 7822c32

	int blocks = (num_pts + N_THREADS - 1) / N_THREADS;
	compute_cov2d_bounds_kernel<<<blocks, N_THREADS>>>(

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:

gsplat/gsplat/cuda/csrc/forward.cu

Line 286 in bae8738

int2 range = tile_bins[tile_id];

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:

gsplat/gsplat/cuda/csrc/bindings.cu

Lines 313 to 319 in bae8738

	torch::Tensor get_tile_bin_edges_tensor(
	int num_intersects, const torch::Tensor &isect_ids_sorted
	) {
	CHECK_INPUT(isect_ids_sorted);
	torch::Tensor tile_bins = torch::zeros(
	{num_intersects, 2}, isect_ids_sorted.options().dtype(torch::kInt32)
	);

I would propose to fix this by changing

gsplat/gsplat/bin_and_sort_gaussians.py

Line 110 in bae8738

tile_bins = get_tile_bin_edges(num_intersects, isect_ids_sorted)

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!