-
Notifications
You must be signed in to change notification settings - Fork 5
Open
Description
CHECK_CUDA(cub::DeviceScan::InclusiveSum(geomState.scanning_space, geomState.scan_size, geomState.tiles_touched, geomState.point_offsets, P), debug)
// Retrieve total number of Gaussian instances to launch and resize aux buffers
CHECK_CUDA(cudaMemcpy(&num_rendered_gs, geomState.point_offsets + P - 1, sizeof(int), cudaMemcpyDeviceToHost), debug);
// printf("num rendered: %d\n", num_rendered);
size_t gs_binning_chunk_size = required<GS_BinningState>(num_rendered_gs);
char* gs_binning_chunkptr = gs_binningBuffer(gs_binning_chunk_size);
gs_binningState = GS_BinningState::fromChunk(gs_binning_chunkptr, num_rendered_gs);
// For each instance to be rendered, produce adequate [ tile | depth ] key
// and corresponding dublicated Gaussian indices to be sorted
duplicateWithKeys << <(P + 255) / 256, 256 >> > (
P,
geomState.means2D,
geomState.depths,
geomState.point_offsets,
gs_binningState.point_list_keys_unsorted,
gs_binningState.point_list_unsorted,
radii,
tile_grid)
CHECK_CUDA(, debug)
int bit = getHigherMsb(tile_grid.x * tile_grid.y);
// Sort complete list of (duplicated) Gaussian indices by keys
CHECK_CUDA(cub::DeviceRadixSort::SortPairs(
gs_binningState.list_sorting_space,
gs_binningState.sorting_size,
gs_binningState.point_list_keys_unsorted, gs_binningState.point_list_keys,
gs_binningState.point_list_unsorted, gs_binningState.point_list,
num_rendered_gs, 0, bit), debug)
CHECK_CUDA(cudaMemset(imgState.ranges, 0, tile_grid.x * tile_grid.y * sizeof(uint2)), debug);
// Identify start and end of per-tile workloads in sorted list
if (num_rendered_gs > 0)
identifyTileRanges << <(num_rendered_gs + 255) / 256, 256 >> > (
num_rendered_gs,
gs_binningState.point_list_keys,
imgState.ranges);
CHECK_CUDA(, debug)
// Let each tile blend its range of Gaussians independently in parallel
const float* feature_ptr = colors_precomp != nullptr ? colors_precomp : geomState.rgb;
CHECK_CUDA(FORWARD::render(
weight_background,
tile_grid, block,
imgState.ranges,
gs_binningState.point_list,
width, height,
geomState.means2D,
feature_ptr,
geomState.conic_opacity,
geomState.alpha_weight,
imgState.weight_sum,
background,
out_color
), debug)
return num_rendered_gs;
}
Reactions are currently unavailable
Metadata
Metadata
Assignees
Labels
No labels