Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Replace uint64_type by std::uint64_t #3435

Merged
merged 1 commit into from
Oct 23, 2019
Merged

Replace uint64_type by std::uint64_t #3435

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 0 additions & 2 deletions gpu/surface/src/convex_hull.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -74,8 +74,6 @@ struct pcl::gpu::PseudoConvexHull3D::Impl

pcl::gpu::PseudoConvexHull3D::PseudoConvexHull3D(std::size_t bsize)
{
pcl::gpu::Static<sizeof(pcl::device::std::uint64_type) == 8>::check();

impl_.reset( new Impl(bsize) );
}
pcl::gpu::PseudoConvexHull3D::~PseudoConvexHull3D() {}
Expand Down
26 changes: 12 additions & 14 deletions gpu/surface/src/cuda/convex_hull.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,8 +66,6 @@ namespace pcl
{
namespace device
{
__global__ void size_check() { Static<sizeof(std::uint64_type) == 8>::check(); };

template<bool use_max>
struct IndOp
{
Expand Down Expand Up @@ -262,7 +260,7 @@ namespace pcl
}

__device__ __forceinline__
std::uint64_type
std::uint64_t
operator()(const PointType& p) const
{
float4 x = p;
Expand Down Expand Up @@ -314,13 +312,13 @@ namespace pcl
//if (neg_count == 0)
// then internal point ==>> idx = INT_MAX

std::uint64_type res = idx;
std::uint64_t res = idx;
res <<= 32;
return res + *reinterpret_cast<unsigned int*>(&dist);
}
};

__global__ void initalClassifyKernel(const InitalClassify ic, const PointType* points, int cloud_size, std::uint64_type* output)
__global__ void initalClassifyKernel(const InitalClassify ic, const PointType* points, int cloud_size, std::uint64_t* output)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;

Expand All @@ -334,7 +332,7 @@ void pcl::device::PointStream::initalClassify()
{
//thrust::device_ptr<const PointType> beg(cloud.ptr());
//thrust::device_ptr<const PointType> end = beg + cloud_size;
thrust::device_ptr<std::uint64_type> out(facets_dists.ptr());
thrust::device_ptr<std::uint64_t> out(facets_dists.ptr());

InitalClassify ic(simplex.p1, simplex.p2, simplex.p3, simplex.p4, cloud_diag);
//thrust::transform(beg, end, out, ic);
Expand All @@ -359,7 +357,7 @@ namespace pcl
__device__ int new_cloud_size;
struct SearchFacetHeads
{
std::uint64_type *facets_dists;
std::uint64_t *facets_dists;
int cloud_size;
int facet_count;
int *perm;
Expand All @@ -371,8 +369,8 @@ namespace pcl
__device__ __forceinline__
void operator()(int facet) const
{
const std::uint64_type* b = facets_dists;
const std::uint64_type* e = b + cloud_size;
const std::uint64_t* b = facets_dists;
const std::uint64_t* e = b + cloud_size;

bool last_thread = facet == facet_count;

Expand Down Expand Up @@ -588,7 +586,7 @@ namespace pcl
{
struct Classify
{
std::uint64_type* facets_dists;
std::uint64_t* facets_dists;
int* scan_buffer;

int* head_points;
Expand All @@ -613,7 +611,7 @@ namespace pcl

if (hi == perm_index)
{
std::uint64_type res = std::numeric_limits<int>::max();
std::uint64_t res = std::numeric_limits<int>::max();
res <<= 32;
facets_dists[point_idx] = res;
}
Expand Down Expand Up @@ -689,7 +687,7 @@ namespace pcl
// if (neg_count == 0)
// new_idx = INT_MAX ==>> internal point

std::uint64_type res = new_idx;
std::uint64_t res = new_idx;
res <<= 32;
res += *reinterpret_cast<unsigned int*>(&dist);

Expand Down Expand Up @@ -731,8 +729,8 @@ void pcl::device::PointStream::classify(FacetStream& fs)
cudaSafeCall( cudaGetLastError() );
cudaSafeCall( cudaDeviceSynchronize() );

thrust::device_ptr<std::uint64_type> beg(facets_dists.ptr());
thrust::device_ptr<std::uint64_type> end = beg + cloud_size;
thrust::device_ptr<std::uint64_t> beg(facets_dists.ptr());
thrust::device_ptr<std::uint64_t> end = beg + cloud_size;

thrust::device_ptr<int> pbeg(perm.ptr());
thrust::sort_by_key(beg, end, pbeg);
Expand Down
2 changes: 1 addition & 1 deletion gpu/surface/src/cuda/device.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,7 +71,7 @@ namespace pcl
struct LessThanByFacet
{
__device__ __forceinline__
bool operator()(const std::uint64_type& e1, const int& e2) const
bool operator()(const std::uint64_t& e1, const int& e2) const
{
int i1 = (int)(e1 >> 32);
return i1 < e2;
Expand Down
8 changes: 4 additions & 4 deletions gpu/surface/src/internal.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,19 +37,19 @@

#pragma once

#include <pcl/gpu/containers/device_array.h>
#include <cstdint>
#include <cuda_runtime.h>

#include <pcl/gpu/containers/device_array.h>

namespace pcl
{
namespace device
{
using std::uint64_type = unsigned long long;
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@PointCloudLibrary/maintainers allow API breakage here since this is an unstable module?

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Maybe deprecate it, and not use it. Like with pcl::int_t types

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Name of file is internal.h - this doesn't sound as it is intended to be used outside.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Maybe deprecate it, and not use it. Like with pcl::int_t types

It is relevant to mention, that I have not been able to compile PCLs CUDA code with C++14 flags. So the usual deprecated attribute might be off-limits and we got rid of our custom deprecation pragmas.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Putting a line in CHANGELOG (and/or RELEASE_NOTES) should be sufficient since this is an internal API and that is allowed to change without warnings.


using PointType = float4;
using Cloud = pcl::gpu::DeviceArray<PointType>;

using FacetsDists = DeviceArray<std::uint64_type>;
using FacetsDists = DeviceArray<std::uint64_t>;
using Perm = DeviceArray<int>;

struct InitalSimplex
Expand Down