Fix illegal memory acces in CUDA Octree builder

gpu/containers/src/error.cpp

@@ -42,5 +42,5 @@
 void pcl::gpu::error(const char *error_string, const char *file, const int line, const char *func)
 {	
     std::cout << "Error: " << error_string << "\t" << file << ":" << line << std::endl;
-    exit(0);
+    exit(EXIT_FAILURE);
 }

gpu/features/test/test_normals.cpp

@@ -371,6 +371,28 @@ TEST(PCL_FeaturesGPU, normals_highlevel_4)
     }
 }
 
+// Test from issue:
+// - https://github.com/PointCloudLibrary/pcl/issues/2371#issuecomment-577727912
+TEST(PCL_FeaturesGPU, issue_2371)
+{
+    // This number is magic, do not set to lower value.
+    // It may affect error reproducibility.
+    const std::size_t N = 1000;
+    std::vector<pcl::PointXYZ> cloud_cpu(N, {0.0, 0.0, 0.0});
+
+    pcl::gpu::NormalEstimation::PointCloud cloud_gpu;
+    cloud_gpu.upload(cloud_cpu);
+
+    pcl::gpu::NormalEstimation ne_gpu;
+    ne_gpu.setInputCloud(cloud_gpu);
+
+    const float radius_search = 2.0F;
+    const int max_results = 500;
+    ne_gpu.setRadiusSearch(radius_search, max_results);
+
+    pcl::gpu::NormalEstimation::Normals normals_gpu;
+    ne_gpu.compute(normals_gpu);
+}
 
 int main (int argc, char** argv)
 {

gpu/octree/src/cuda/octree_builder.cu

@@ -212,7 +212,7 @@ namespace pcl
 
                 __syncthreads();
 
-                while (tasks_beg < tasks_end)
+                while (tasks_beg < tasks_end && level < Morton::levels)
                 {                  
                     int task_count = tasks_end - tasks_beg;                    
                     int iters = divUp(task_count, CTA_SIZE);
@@ -245,7 +245,7 @@ namespace pcl
                                     octree.begs [offset + i] = cell_begs[i];
                                     octree.ends [offset + i] = cell_begs[i + 1];
                                     octree.codes[offset + i] = parent_code_shifted + cell_code[i];
-
+                                    octree.nodes[offset + i] = 0;
                                     octree.parent[offset + i] = task;
                                     mask |= (1 << cell_code[i]);
                                 }