mirror/vcpkg
1
0
Fork 0
mirror of https://github.com/microsoft/vcpkg.git synced 2024-11-24 03:09:00 +08:00
Code Issues Actions 1 Packages Projects Releases Wiki Activity

[pcl] Fix cuda building compatability issues with cuda 10.1 (#7388)

Browse Source
This commit is contained in:
contre 2019-07-24 10:58:19 -07:00 committed by Curtis J Bezault
parent cf32345161
commit 265921b4a3
3 changed files with 135 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
ports/pcl/CONTROL
View File

@ -1,5 +1,5 @@
Source: pcl
Version: 1.9.1-4
Version: 1.9.1-5
Homepage: https://github.com/PointCloudLibrary/pcl
Description: Point Cloud Library (PCL) is open source library for 2D/3D image and point cloud processing.
Build-Depends: eigen3, flann, qhull, vtk, libpng, boost-system, boost-filesystem, boost-thread, boost-date-time, boost-iostreams, boost-random, boost-foreach, boost-dynamic-bitset, boost-property-map, boost-graph, boost-multi-array, boost-signals2, boost-ptr-container, boost-uuid, boost-interprocess, boost-asio

133
ports/pcl/cuda_10_1.patch Normal file
View File

@ -0,0 +1,133 @@
diff --git a/cuda/common/include/pcl/cuda/thrust.h b/cuda/common/include/pcl/cuda/thrust.h
index 57586ab..af073d7 100644
--- a/cuda/common/include/pcl/cuda/thrust.h
+++ b/cuda/common/include/pcl/cuda/thrust.h
@@ -42,6 +42,7 @@
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
+#include <thrust/device_malloc.h>
#include <thrust/copy.h>
#include <thrust/device_ptr.h>
#include <thrust/sequence.h>
diff --git a/gpu/features/src/fpfh.cu b/gpu/features/src/fpfh.cu
index 8d34f76..59a5f0c 100644
--- a/gpu/features/src/fpfh.cu
+++ b/gpu/features/src/fpfh.cu
@@ -231,12 +231,12 @@ namespace pcl
int *sinds = sindices + Warp::WARP_SIZE * warp_idx;
int size = sizes[idx];
- for(int i = lane; __any(i < size); i += Warp::STRIDE)
+ for(int i = lane; __any_sync(0xFFFFFFFF, i < size); i += Warp::STRIDE)
{
if (i < size)
sinds[lane] = ginds[i];
- int inds_num = __popc(__ballot(i < size));
+ int inds_num = __popc(__ballot_sync(0xFFFFFFFF, i < size));
for(int j = 0; j < inds_num; ++j)
{
diff --git a/gpu/octree/src/cuda/approx_nsearch.cu b/gpu/octree/src/cuda/approx_nsearch.cu
index e457255..3e1adfe 100644
--- a/gpu/octree/src/cuda/approx_nsearch.cu
+++ b/gpu/octree/src/cuda/approx_nsearch.cu
@@ -141,7 +141,7 @@ namespace pcl { namespace device { namespace appnearest_search
{
__shared__ volatile int per_warp_buffer[KernelPolicy::WARPS_COUNT];
- int mask = __ballot(node_idx != -1);
+ int mask = __ballot_sync(0xFFFFFFFF, node_idx != -1);
while(mask)
{
@@ -275,7 +275,7 @@ namespace pcl { namespace device { namespace appnearest_search
bool active = query_index < batch.queries_num;
- if (__all(active == false))
+ if (__all_sync(0xFFFFFFFF, active == false))
return;
Warp_appNearestSearch search(batch, query_index);
diff --git a/gpu/octree/src/cuda/knn_search.cu b/gpu/octree/src/cuda/knn_search.cu
index a99655d..b55e3c1 100644
--- a/gpu/octree/src/cuda/knn_search.cu
+++ b/gpu/octree/src/cuda/knn_search.cu
@@ -106,7 +106,7 @@ namespace pcl { namespace device { namespace knn_search
else
query_index = -1;
- while(__any(active))
+ while(__any_sync(0xFFFFFFFF, active))
{
int leaf = -1;
@@ -163,7 +163,7 @@ namespace pcl { namespace device { namespace knn_search
__device__ __forceinline__ void processLeaf(int node_idx)
{
- int mask = __ballot(node_idx != -1);
+ int mask = __ballot_sync(0xFFFFFFFF, node_idx != -1);
unsigned int laneId = Warp::laneId();
unsigned int warpId = Warp::id();
@@ -310,7 +310,7 @@ namespace pcl { namespace device { namespace knn_search
bool active = query_index < batch.queries_num;
- if (__all(active == false))
+ if (__all_sync(0xFFFFFFFF, active == false))
return;
Warp_knnSearch search(batch, query_index);
diff --git a/gpu/octree/src/cuda/radius_search.cu b/gpu/octree/src/cuda/radius_search.cu
index f90273f..8ae84e7 100644
--- a/gpu/octree/src/cuda/radius_search.cu
+++ b/gpu/octree/src/cuda/radius_search.cu
@@ -155,7 +155,7 @@ namespace pcl
else
query_index = -1;
- while(__any(active))
+ while(__any_sync(0xFFFFFFFF, active))
{
int leaf = -1;
@@ -217,7 +217,7 @@ namespace pcl
__device__ __forceinline__ void processLeaf(int leaf)
{
- int mask = __ballot(leaf != -1);
+ int mask = __ballot_sync(0xFFFFFFFF, leaf != -1);
while(mask)
{
@@ -255,7 +255,7 @@ namespace pcl
int *out = batch.output + active_query_index * batch.max_results + active_found_count;
int length_left = batch.max_results - active_found_count;
- int test = __any(active_lane == laneId && (leaf & KernelPolicy::CHECK_FLAG));
+ int test = __any_sync(0xFFFFFFFF, active_lane == laneId && (leaf & KernelPolicy::CHECK_FLAG));
if (test)
{
@@ -329,7 +329,7 @@ namespace pcl
total_new += new_nodes;
out += new_nodes;
- if (__all(idx >= length) || __any(out_of_bounds) || total_new == length_left)
+ if (__all_sync(0xFFFFFFFF, idx >= length) || __any_sync(0xFFFFFFFF, out_of_bounds) || total_new == length_left)
break;
}
return min(total_new, length_left);
@@ -343,7 +343,7 @@ namespace pcl
bool active = query_index < batch.queries.size;
- if (__all(active == false))
+ if (__all_sync(0xFFFFFFFF, active == false))
return;
Warp_radiusSearch<BatchType> search(batch, query_index);

1
ports/pcl/portfile.cmake
View File

@ -11,6 +11,7 @@ vcpkg_from_github(
pcl_config.patch
use_flann_targets.patch
boost-1.70.patch
cuda_10_1.patch
)
file(REMOVE ${SOURCE_PATH}/cmake/Modules/FindFLANN.cmake)