@@ -26,19 +26,21 @@ struct OutputSegmentationPointType
2626 float z;
2727 std::uint8_t class_id;
2828 float probability;
29+ float entropy;
2930} __attribute__((packed));
3031
3132__global__ void createVisualizationPointcloudKernel (
3233 const float4 * input_features, const float * colors, const std::int64_t * labels,
33- float4 * output_points, std::size_t num_points)
34+ float4 * output_points, std::size_t num_classes, std:: size_t num_points)
3435{
3536 const auto idx = static_cast <std::uint32_t >(blockIdx .x * blockDim .x + threadIdx .x );
3637 if (idx >= num_points) {
3738 return ;
3839 }
3940
4041 const auto label = labels[idx];
41- const auto color = colors[label];
42+ const auto color =
43+ label >= 0 && static_cast <std::size_t >(label) < num_classes ? colors[label] : 0 .0f ;
4244
4345 output_points[idx] =
4446 make_float4 (input_features[idx].x , input_features[idx].y , input_features[idx].z , color);
@@ -55,11 +57,78 @@ __global__ void createSegmentationPointcloudKernel(
5557
5658 const auto input_point = input_features[idx];
5759 const auto label = labels[idx];
60+ const bool has_valid_label = label >= 0 && static_cast <std::size_t >(label) < num_classes;
61+ float entropy = 0 .0f ;
62+ for (std::size_t class_idx = 0 ; class_idx < num_classes; ++class_idx) {
63+ const auto probability = pred_probs[idx * num_classes + class_idx];
64+ if (probability > 0 .0f ) {
65+ entropy -= probability * logf (probability);
66+ }
67+ }
68+ if (num_classes > 1 ) {
69+ entropy /= logf (static_cast <float >(num_classes));
70+ }
5871 output_points[idx].x = input_point.x ;
5972 output_points[idx].y = input_point.y ;
6073 output_points[idx].z = input_point.z ;
61- output_points[idx].class_id = static_cast <std::uint8_t >(label);
62- output_points[idx].probability = pred_probs[idx * num_classes + label];
74+ output_points[idx].class_id = has_valid_label ? static_cast <std::uint8_t >(label) : 255U ;
75+ output_points[idx].probability = has_valid_label ? pred_probs[idx * num_classes + label] : 0 .0f ;
76+ output_points[idx].entropy = entropy;
77+ }
78+
79+ __global__ void reconstructPartialKernel (
80+ const std::int64_t * __restrict__ inverse_map, const std::int64_t * __restrict__ voxel_labels,
81+ const float * __restrict__ voxel_probs, std::int64_t * __restrict__ output_labels,
82+ float * __restrict__ output_probs, std::size_t num_classes, std::size_t num_cropped_points,
83+ std::size_t num_voxels)
84+ {
85+ const auto point_idx = blockIdx .x * blockDim .y + threadIdx .y ;
86+ const auto class_idx = blockIdx .y * blockDim .x + threadIdx .x ;
87+
88+ if (point_idx >= num_cropped_points || class_idx >= num_classes) return ;
89+
90+ const auto voxel_idx = inverse_map[point_idx];
91+ const bool has_valid_voxel = voxel_idx >= 0 && static_cast <std::size_t >(voxel_idx) < num_voxels;
92+ if (class_idx == 0 ) {
93+ output_labels[point_idx] = has_valid_voxel ? voxel_labels[voxel_idx] : 255 ;
94+ }
95+
96+ output_probs[point_idx * num_classes + class_idx] =
97+ has_valid_voxel ? voxel_probs[static_cast <std::size_t >(voxel_idx) * num_classes + class_idx]
98+ : 0 .0f ;
99+ }
100+
101+ __global__ void reconstructFullKernel (
102+ const std::uint32_t * __restrict__ crop_mask, const std::uint32_t * __restrict__ crop_indices,
103+ const std::int64_t * __restrict__ inverse_map, const std::int64_t * __restrict__ voxel_labels,
104+ const float * __restrict__ voxel_probs, std::int64_t * __restrict__ output_labels,
105+ float * __restrict__ output_probs, std::size_t num_classes, std::size_t num_points,
106+ std::size_t num_voxels)
107+ {
108+ const auto point_idx = blockIdx .x * blockDim .y + threadIdx .y ;
109+ const auto class_idx = blockIdx .y * blockDim .x + threadIdx .x ;
110+
111+ if (point_idx >= num_points || class_idx >= num_classes) return ;
112+
113+ const auto mask = crop_mask[point_idx];
114+ if (mask == 0 ) {
115+ if (class_idx == 0 ) output_labels[point_idx] = 255 ;
116+ return ;
117+ }
118+
119+ const auto cropped_idx = crop_indices[point_idx] - 1 ;
120+ const auto voxel_idx = inverse_map[cropped_idx];
121+ if (voxel_idx < 0 || static_cast <std::size_t >(voxel_idx) >= num_voxels) {
122+ if (class_idx == 0 ) output_labels[point_idx] = 255 ;
123+ return ;
124+ }
125+
126+ if (class_idx == 0 ) {
127+ output_labels[point_idx] = voxel_labels[voxel_idx];
128+ }
129+
130+ output_probs[point_idx * num_classes + class_idx] =
131+ voxel_probs[static_cast <std::size_t >(voxel_idx) * num_classes + class_idx];
63132}
64133
65134template <typename OutputPointT>
@@ -213,13 +282,13 @@ PostprocessCuda::PostprocessCuda(const PTv3Config & config, cudaStream_t stream)
213282
214283void PostprocessCuda::createVisualizationPointcloud (
215284 const float * input_features, const std::int64_t * labels, float * output_points,
216- std::size_t num_points)
285+ std::size_t num_classes, std:: size_t num_points)
217286{
218287 auto num_blocks = divup (num_points, config_.threads_per_block_ );
219288
220289 createVisualizationPointcloudKernel<<<num_blocks, config_.threads_per_block_, 0 , stream_>>> (
221290 reinterpret_cast <const float4 *>(input_features), color_map_d_.get (), labels,
222- reinterpret_cast <float4 *>(output_points), num_points);
291+ reinterpret_cast <float4 *>(output_points), num_classes, num_points);
223292
224293 CHECK_CUDA_ERROR (cudaStreamSynchronize (stream_));
225294}
@@ -237,6 +306,37 @@ void PostprocessCuda::createSegmentationPointcloud(
237306 CHECK_CUDA_ERROR (cudaStreamSynchronize (stream_));
238307}
239308
309+ void PostprocessCuda::reconstructPartial (
310+ const std::int64_t * inverse_map, const std::int64_t * voxel_labels, const float * voxel_probs,
311+ std::int64_t * output_labels, float * output_probs, std::size_t num_classes,
312+ std::size_t num_cropped_points, std::size_t num_voxels)
313+ {
314+ auto block = dim3 (32 , 8 );
315+ auto grid = dim3 (divup (num_cropped_points, block.y ), divup (num_classes, block.x ));
316+
317+ reconstructPartialKernel<<<grid, block, 0 , stream_>>> (
318+ inverse_map, voxel_labels, voxel_probs, output_labels, output_probs, num_classes,
319+ num_cropped_points, num_voxels);
320+
321+ CHECK_CUDA_ERROR (cudaStreamSynchronize (stream_));
322+ }
323+
324+ void PostprocessCuda::reconstructFull (
325+ const std::uint32_t * crop_mask, const std::uint32_t * crop_indices,
326+ const std::int64_t * inverse_map, const std::int64_t * voxel_labels, const float * voxel_probs,
327+ std::int64_t * output_labels, float * output_probs, std::size_t num_classes,
328+ std::size_t num_points, std::size_t num_voxels)
329+ {
330+ auto block = dim3 (32 , 8 );
331+ auto grid = dim3 (divup (num_points, block.y ), divup (num_classes, block.x ));
332+
333+ reconstructFullKernel<<<grid, block, 0 , stream_>>> (
334+ crop_mask, crop_indices, inverse_map, voxel_labels, voxel_probs, output_labels, output_probs,
335+ num_classes, num_points, num_voxels);
336+
337+ CHECK_CUDA_ERROR (cudaStreamSynchronize (stream_));
338+ }
339+
240340std::size_t PostprocessCuda::createFilteredPointcloud (
241341 const void * compact_input_points, CloudFormat input_format, CloudFormat output_format,
242342 const float * pred_probs, void * output_points, std::size_t num_classes, std::size_t num_points)
0 commit comments