fix build with CUDA (PointCloudLibrary#5976)

* fix build with CUDA * update cmake/pcl_find_cuda.cmake * fix build with CUDA 12.4
larshg · Mar 9, 2024 · ccc4eee · ccc4eee
1 parent eebcfcb
commit ccc4eee
Show file tree

Hide file tree

Showing 5 changed files with 215 additions and 211 deletions.
diff --git a/cmake/pcl_find_cuda.cmake b/cmake/pcl_find_cuda.cmake
@@ -12,6 +12,7 @@ if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.11)
 
     if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.17)
       find_package(CUDAToolkit QUIET)
+      set(CUDA_TOOLKIT_INCLUDE ${CUDAToolkit_INCLUDE_DIRS})
     else()
       set(CUDA_FIND_QUIETLY TRUE)
       find_package(CUDA 9.0)

diff --git a/gpu/features/src/centroid.cu b/gpu/features/src/centroid.cu
@@ -36,6 +36,7 @@
 
 #include "internal.hpp"
 
+#include <thrust/tuple.h>
 #include <thrust/device_ptr.h>
 #include <thrust/transform_reduce.h>
 #include <thrust/iterator/permutation_iterator.h>
@@ -56,26 +57,26 @@ namespace pcl
 
         struct PlusFloat3
         {
-            __device__ __forceinline__ float3 operator()(const float3& e1, const float3& e2) const 
-            { 
-                return make_float3(e1.x + e2.x, e1.y + e2.y, e1.z + e2.z); 
+            __device__ __forceinline__ float3 operator()(const float3& e1, const float3& e2) const
+            {
+                return make_float3(e1.x + e2.x, e1.y + e2.y, e1.z + e2.z);
             }
         };
-        
+
         struct TupleDistCvt
         {
             float3 pivot_;
             TupleDistCvt(const float3& pivot) : pivot_(pivot) {}
-            __device__ __forceinline__ thrust::tuple<float, int> operator()(const thrust::tuple<float4, int>& t) const 
-            { 
-                float4 point = t.get<0>();
+            __device__ __forceinline__ thrust::tuple<float, int> operator()(const thrust::tuple<float4, int>& t) const
+            {
+                float4 point = thrust::get<0>(t);
 
                 float dx = pivot_.x - point.x;
                 float dy = pivot_.y - point.y;
                 float dz = pivot_.z - point.z;
                 float dist = sqrt(dx*dx + dy*dy + dz*dz);
 
-                return thrust::tuple<float, int>(dist, t.get<1>());                                
+                return thrust::tuple<float, int>(dist, thrust::get<1>(t));
             }
         };
 
@@ -87,7 +88,7 @@ void pcl::device::compute3DCentroid(const DeviceArray<PointT>& cloud, float3& ce
 {
     thrust::device_ptr<PointT> src_beg((PointT*)cloud.ptr());
     thrust::device_ptr<PointT> src_end = src_beg + cloud.size();
-                
+
     centroid = transform_reduce(src_beg, src_beg, PointT2float3<PointT>(), make_float3(0.f, 0.f, 0.f), PlusFloat3());
     centroid *= 1.f/cloud.size();
 }
@@ -99,13 +100,13 @@ void pcl::device::compute3DCentroid(const DeviceArray<PointT>& cloud, const Indi
         compute3DCentroid(cloud, centroid);
     else
     {
-        thrust::device_ptr<PointT> src_beg((PointT*)cloud.ptr());    
+        thrust::device_ptr<PointT> src_beg((PointT*)cloud.ptr());
         thrust::device_ptr<int> map_beg((int*)indices.ptr());
         thrust::device_ptr<int> map_end = map_beg + indices.size();
 
 
         centroid = transform_reduce(make_permutation_iterator(src_beg, map_beg),
-            make_permutation_iterator(src_beg, map_end), 
+            make_permutation_iterator(src_beg, map_end),
             PointT2float3<PointT>(), make_float3(0.f, 0.f, 0.f), PlusFloat3());
 
         centroid *= 1.f/indices.size();
@@ -114,7 +115,7 @@ void pcl::device::compute3DCentroid(const DeviceArray<PointT>& cloud, const Indi
 
 template<typename PointT>
 float3 pcl::device::getMaxDistance(const DeviceArray<PointT>& cloud, const float3& pivot)
-{   
+{
     thrust::device_ptr<PointT> src_beg((PointT*)cloud.ptr());
     thrust::device_ptr<PointT> src_end = src_beg + cloud.size();
 
@@ -123,26 +124,26 @@ float3 pcl::device::getMaxDistance(const DeviceArray<PointT>& cloud, const float
 
     thrust::tuple<float, int> init(0.f, 0);
     thrust::maximum<thrust::tuple<float, int>> op;
-    
+
     thrust::tuple<float, int> res =
         transform_reduce(
         make_zip_iterator(make_tuple( src_beg, cf )),
         make_zip_iterator(make_tuple( src_beg, ce )),
         TupleDistCvt(pivot), init, op);
 
-    float4 point = src_beg[res.get<1>()];
+    float4 point = src_beg[thrust::get<1>(res)];
 
     return make_float3(point.x, point.y, point.z);
 }
 
 
 template<typename PointT>
 float3 pcl::device::getMaxDistance(const DeviceArray<PointT>& cloud, const Indices& indices, const float3& pivot)
-{   
+{
     if (indices.empty())
         return getMaxDistance(cloud, pivot);
 
-    thrust::device_ptr<PointT> src_beg((PointT*)cloud.ptr());    
+    thrust::device_ptr<PointT> src_beg((PointT*)cloud.ptr());
     thrust::device_ptr<int> map_beg((int*)indices.ptr());
     thrust::device_ptr<int> map_end = map_beg + indices.size();
 
@@ -151,13 +152,13 @@ float3 pcl::device::getMaxDistance(const DeviceArray<PointT>& cloud, const Indic
 
     thrust::tuple<float, int> init(0.f, 0);
     thrust::maximum<thrust::tuple<float, int>> op;
-    
+
     thrust::tuple<float, int> res = transform_reduce(
         make_zip_iterator(make_tuple( make_permutation_iterator(src_beg, map_beg), cf )),
         make_zip_iterator(make_tuple( make_permutation_iterator(src_beg, map_end), ce )),
         TupleDistCvt(pivot), init, op);
 
-    float4 point = src_beg[map_beg[res.get<1>()]];
+    float4 point = src_beg[map_beg[thrust::get<1>(res)]];
 
     return make_float3(point.x, point.y, point.z);
 }