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

Prefer numeric_limits from standard library #3360

Merged
merged 3 commits into from
Sep 26, 2019
Merged

Prefer numeric_limits from standard library #3360

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
de69551
Use std::numeric_limits instead of custom implementation of numeric_l…
Sep 23, 2019
1cac0d8
Added hint why CUDA flag --expt-relaxed-constexpr is required
Sep 24, 2019
63f5707
Added constexpr keyword to lines with std::numeric_limits if possible
Sep 24, 2019
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: 1 addition & 1 deletion apps/src/ni_linemod.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -322,7 +322,7 @@ class NILinemod
mps_.segmentAndRefine (regions, model_coefficients, inlier_indices, labels, label_indices, boundary_indices);
PCL_DEBUG ("Number of planar regions detected: %lu for a cloud of %lu points and %lu normals.\n", regions.size (), search_.getInputCloud ()->points.size (), normal_cloud->points.size ());

double max_dist = numeric_limits<double>::max ();
double max_dist = std::numeric_limits<double>::max ();
// Compute the distances from all the planar regions to the picked point, and select the closest region
int idx = -1;
for (size_t i = 0; i < regions.size (); ++i)
Expand Down
2 changes: 1 addition & 1 deletion apps/src/pcd_select_object_plane.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -341,7 +341,7 @@ class ObjectSelection
}
print_highlight ("Number of planar regions detected: %lu for a cloud of %lu points\n", regions.size (), cloud_->size ());

double max_dist = numeric_limits<double>::max ();
double max_dist = std::numeric_limits<double>::max ();
// Compute the distances from all the planar regions to the picked point, and select the closest region
int idx = -1;
for (size_t i = 0; i < regions.size (); ++i)
Expand Down
3 changes: 3 additions & 0 deletions cmake/pcl_find_cuda.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,4 +47,7 @@ if(CUDA_FOUND)

# Prevent compilation issues between recent gcc versions and old CUDA versions
list(APPEND CUDA_NVCC_FLAGS "-D_FORCE_INLINES")

# Allow calling a constexpr __host__ function from a __device__ function.
list(APPEND CUDA_NVCC_FLAGS "--expt-relaxed-constexpr")
endif()
15 changes: 8 additions & 7 deletions gpu/features/include/pcl/gpu/features/device/eigen.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,7 +90,8 @@
#ifndef PCL_GPU_FEATURES_EIGEN_HPP_
#define PCL_GPU_FEATURES_EIGEN_HPP_

#include <pcl/gpu/utils/device/limits.hpp>
#include <limits>

#include <pcl/gpu/utils/device/algorithm.hpp>
#include <pcl/gpu/utils/device/vector_math.hpp>

Expand All @@ -113,7 +114,7 @@ namespace pcl

__device__ __forceinline__ void computeRoots3(float c0, float c1, float c2, float3& roots)
{
if ( std::abs(c0) < numeric_limits<float>::epsilon())// one root is 0 -> quadratic equation
if ( std::abs(c0) < std::numeric_limits<float>::epsilon())// one root is 0 -> quadratic equation
{
computeRoots2 (c2, c1, roots);
}
Expand Down Expand Up @@ -217,7 +218,7 @@ namespace pcl
float m0123 = fmaxf( max01, max23);
float scale = fmaxf( max45, m0123);

if (scale <= numeric_limits<float>::min())
if (scale <= std::numeric_limits<float>::min())
scale = 1.f;

mat_pkg[0] /= scale;
Expand Down Expand Up @@ -245,13 +246,13 @@ namespace pcl

computeRoots3(c0, c1, c2, evals);

if(evals.z - evals.x <= numeric_limits<float>::epsilon())
if(evals.z - evals.x <= std::numeric_limits<float>::epsilon())
{
evecs[0] = make_float3(1.f, 0.f, 0.f);
evecs[1] = make_float3(0.f, 1.f, 0.f);
evecs[2] = make_float3(0.f, 0.f, 1.f);
}
else if (evals.y - evals.x <= numeric_limits<float>::epsilon() )
else if (evals.y - evals.x <= std::numeric_limits<float>::epsilon() )
{
// first and second equal
tmp[0] = row0(); tmp[1] = row1(); tmp[2] = row2();
Expand Down Expand Up @@ -281,7 +282,7 @@ namespace pcl
evecs[1] = unitOrthogonal(evecs[2]);
evecs[0] = cross(evecs[1], evecs[2]);
}
else if (evals.z - evals.y <= numeric_limits<float>::epsilon() )
else if (evals.z - evals.y <= std::numeric_limits<float>::epsilon() )
{
// second and third equal
tmp[0] = row0(); tmp[1] = row1(); tmp[2] = row2();
Expand Down Expand Up @@ -439,7 +440,7 @@ namespace pcl
__device__ __forceinline__ static bool isMuchSmallerThan (float x, float y)
{
// copied from <eigen>/include/Eigen/src/Core/NumTraits.h
const float prec_sqr = numeric_limits<float>::epsilon() * numeric_limits<float>::epsilon();
constexpr float prec_sqr = std::numeric_limits<float>::epsilon() * std::numeric_limits<float>::epsilon();
return x * x <= prec_sqr * y * y;
}

Expand Down
7 changes: 3 additions & 4 deletions gpu/features/src/normal_3d.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -82,11 +82,10 @@ namespace pcl
int size = sizes[idx];
int lane = Warp::laneId();

if (size < MIN_NEIGHBOORS)
if ((size < MIN_NEIGHBOORS) && (lane == 0))
{
const float NaN = numeric_limits<float>::quiet_NaN();
if (lane == 0)
normals.data[idx] = make_float4(NaN, NaN, NaN, NaN);
constexpr float NaN = std::numeric_limits<float>::quiet_NaN();
normals.data[idx] = make_float4(NaN, NaN, NaN, NaN);
}

const int *ibeg = indices.ptr(idx);
Expand Down
9 changes: 4 additions & 5 deletions gpu/features/src/spinimages.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,7 +37,6 @@
#include "internal.hpp"
#include "pcl/gpu/utils/device/warp.hpp"
#include "pcl/gpu/utils/device/block.hpp"
#include "pcl/gpu/utils/device/limits.hpp"
#include "pcl/gpu/utils/device/vector_math.hpp"
#include "pcl/gpu/utils/device/functional.hpp"

Expand Down Expand Up @@ -75,7 +74,7 @@ namespace pcl

struct Div12eps
{
__device__ __forceinline__ float operator()(float v1, float v2) const { return (float)(v1 / ( v2 + numeric_limits<double>::epsilon() )); }
__device__ __forceinline__ float operator()(float v1, float v2) const { return (float)(v1 / ( v2 + std::numeric_limits<double>::epsilon() )); }
};

struct DivValIfNonZero
Expand Down Expand Up @@ -146,7 +145,7 @@ namespace pcl
float3 rotation_axis = AxesStrategy::getRotationAxes(index, origin_normal);
rotation_axis = normalized_safe(rotation_axis); //normalize if non-zero

const float eps = numeric_limits<float>::epsilon ();
constexpr float eps = std::numeric_limits<float>::epsilon ();

for(int i_neighb = threadIdx.x; i_neighb < neighb_count; i_neighb += CTA_SIZE)
{
Expand Down Expand Up @@ -179,8 +178,8 @@ namespace pcl
cos_dir_axis = fmax(-1.f, fmin(1.f, cos_dir_axis));

// compute coordinates w.r.t. the reference frame
float beta = numeric_limits<float>::quiet_NaN();
float alpha = numeric_limits<float>::quiet_NaN();
float beta;
float alpha;

if (radial) // radial spin image structure
{
Expand Down
2 changes: 1 addition & 1 deletion gpu/kinfu/src/cuda/bilateral_pyrdown.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -84,7 +84,7 @@ namespace pcl
}

int res = __float2int_rn (sum1 / sum2);
dst.ptr (y)[x] = max (0, min (res, numeric_limits<short>::max ()));
dst.ptr (y)[x] = max (0, min (res, std::numeric_limits<short>::max ()));
}

//////////////////////////////////////////////////////////////////////////////////////////////////////////////////
Expand Down
2 changes: 0 additions & 2 deletions gpu/kinfu/src/cuda/device.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,8 +38,6 @@
#ifndef PCL_GPU_KINFU_DEVICE_HPP_
#define PCL_GPU_KINFU_DEVICE_HPP_

//#include <pcl/gpu/utils/device/limits.hpp>
//#include <pcl/gpu/utils/device/vector_math.hpp>
#include "utils.hpp" //temporary reimplementing to release kinfu without pcl_gpu_utils

#include "internal.h"
Expand Down
2 changes: 1 addition & 1 deletion gpu/kinfu/src/cuda/extract.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -346,7 +346,7 @@ namespace pcl

if (idx >= points.size)
return;
const float qnan = numeric_limits<float>::quiet_NaN ();
constexpr float qnan = std::numeric_limits<float>::quiet_NaN ();
float3 n = make_float3 (qnan, qnan, qnan);

float3 point = fetchPoint (idx);
Expand Down
12 changes: 6 additions & 6 deletions gpu/kinfu/src/cuda/maps.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,7 +65,7 @@ namespace pcl
vmap.ptr (v + depth.rows * 2)[u] = vz;
}
else
vmap.ptr (v)[u] = numeric_limits<float>::quiet_NaN ();
vmap.ptr (v)[u] = std::numeric_limits<float>::quiet_NaN ();

}
}
Expand All @@ -81,7 +81,7 @@ namespace pcl

if (u == cols - 1 || v == rows - 1)
{
nmap.ptr (v)[u] = numeric_limits<float>::quiet_NaN ();
nmap.ptr (v)[u] = std::numeric_limits<float>::quiet_NaN ();
return;
}

Expand All @@ -107,7 +107,7 @@ namespace pcl
nmap.ptr (v + 2 * rows)[u] = r.z;
}
else
nmap.ptr (v)[u] = numeric_limits<float>::quiet_NaN ();
nmap.ptr (v)[u] = std::numeric_limits<float>::quiet_NaN ();
}
}
}
Expand Down Expand Up @@ -159,7 +159,7 @@ namespace pcl
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;

const float qnan = pcl::device::numeric_limits<float>::quiet_NaN ();
const float qnan = std::numeric_limits<float>::quiet_NaN ();

if (x < cols && y < rows)
{
Expand Down Expand Up @@ -238,7 +238,7 @@ namespace pcl
if (x >= dcols || y >= drows)
return;

const float qnan = numeric_limits<float>::quiet_NaN ();
const float qnan = std::numeric_limits<float>::quiet_NaN ();

int xs = x * 2;
int ys = y * 2;
Expand Down Expand Up @@ -332,7 +332,7 @@ namespace pcl
if (x >= cols || y >= rows)
return;

const float qnan = numeric_limits<float>::quiet_NaN ();
const float qnan = std::numeric_limits<float>::quiet_NaN ();

T t;
t.x = map.ptr (y)[x];
Expand Down
2 changes: 1 addition & 1 deletion gpu/kinfu/src/cuda/normals_eigen.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -58,7 +58,7 @@ namespace pcl
if (u >= cols || v >= rows)
return;

nmap.ptr (v)[u] = numeric_limits<float>::quiet_NaN ();
nmap.ptr (v)[u] = std::numeric_limits<float>::quiet_NaN ();

if (isnan (vmap.ptr (v)[u]))
return;
Expand Down
10 changes: 5 additions & 5 deletions gpu/kinfu/src/cuda/ray_caster.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -124,13 +124,13 @@ namespace pcl
int3 g = getVoxel (point);

if (g.x <= 0 || g.x >= VOLUME_X - 1)
return numeric_limits<float>::quiet_NaN ();
return std::numeric_limits<float>::quiet_NaN ();

if (g.y <= 0 || g.y >= VOLUME_Y - 1)
return numeric_limits<float>::quiet_NaN ();
return std::numeric_limits<float>::quiet_NaN ();

if (g.z <= 0 || g.z >= VOLUME_Z - 1)
return numeric_limits<float>::quiet_NaN ();
return std::numeric_limits<float>::quiet_NaN ();

float vx = (g.x + 0.5f) * cell_size.x;
float vy = (g.y + 0.5f) * cell_size.y;
Expand Down Expand Up @@ -163,8 +163,8 @@ namespace pcl
if (x >= cols || y >= rows)
return;

vmap.ptr (y)[x] = numeric_limits<float>::quiet_NaN ();
nmap.ptr (y)[x] = numeric_limits<float>::quiet_NaN ();
vmap.ptr (y)[x] = std::numeric_limits<float>::quiet_NaN ();
nmap.ptr (y)[x] = std::numeric_limits<float>::quiet_NaN ();

float3 ray_start = tcurr;
float3 ray_next = Rcurr * get_ray_next (x, y) + tcurr;
Expand Down
2 changes: 1 addition & 1 deletion gpu/kinfu/src/cuda/tsdf_volume.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -451,7 +451,7 @@ namespace pcl
bool integrate = true;
if ((x > 0 && x < VOLUME_X-2) && (y > 0 && y < VOLUME_Y-2) && (z > 0 && z < VOLUME_Z-2))
{
const float qnan = numeric_limits<float>::quiet_NaN();
const float qnan = std::numeric_limits<float>::quiet_NaN();
float3 normal = make_float3(qnan, qnan, qnan);

float Fn, Fp;
Expand Down
35 changes: 8 additions & 27 deletions gpu/kinfu/src/cuda/utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,6 +39,8 @@
#ifndef PCL_GPU_KINFU_CUDA_UTILS_HPP_
#define PCL_GPU_KINFU_CUDA_UTILS_HPP_

#include <limits>

#include <cuda.h>

namespace pcl
Expand All @@ -50,27 +52,6 @@ namespace pcl
{
T c(a); a=b; b=c;
}

template<typename T> struct numeric_limits;

template<> struct numeric_limits<float>
{
__device__ __forceinline__ static float
quiet_NaN() { return __int_as_float(0x7fffffff); /*CUDART_NAN_F*/ };
__device__ __forceinline__ static float
epsilon() { return 1.192092896e-07f/*FLT_EPSILON*/; };

__device__ __forceinline__ static float
min() { return 1.175494351e-38f/*FLT_MIN*/; };
__device__ __forceinline__ static float
max() { return 3.402823466e+38f/*FLT_MAX*/; };
};

template<> struct numeric_limits<short>
{
__device__ __forceinline__ static short
max() { return SHRT_MAX; };
};

__device__ __forceinline__ float
dot(const float3& v1, const float3& v2)
Expand Down Expand Up @@ -142,7 +123,7 @@ namespace pcl
__device__ __forceinline__ void
computeRoots3(float c0, float c1, float c2, float3& roots)
{
if ( std::abs(c0) < numeric_limits<float>::epsilon())// one root is 0 -> quadratic equation
if ( std::abs(c0) < std::numeric_limits<float>::epsilon())// one root is 0 -> quadratic equation
{
computeRoots2 (c2, c1, roots);
}
Expand Down Expand Up @@ -249,7 +230,7 @@ namespace pcl
float m0123 = fmaxf( max01, max23);
float scale = fmaxf( max45, m0123);

if (scale <= numeric_limits<float>::min())
if (scale <= std::numeric_limits<float>::min())
scale = 1.f;

mat_pkg[0] /= scale;
Expand Down Expand Up @@ -277,13 +258,13 @@ namespace pcl

computeRoots3(c0, c1, c2, evals);

if(evals.z - evals.x <= numeric_limits<float>::epsilon())
if(evals.z - evals.x <= std::numeric_limits<float>::epsilon())
{
evecs[0] = make_float3(1.f, 0.f, 0.f);
evecs[1] = make_float3(0.f, 1.f, 0.f);
evecs[2] = make_float3(0.f, 0.f, 1.f);
}
else if (evals.y - evals.x <= numeric_limits<float>::epsilon() )
else if (evals.y - evals.x <= std::numeric_limits<float>::epsilon() )
{
// first and second equal
tmp[0] = row0(); tmp[1] = row1(); tmp[2] = row2();
Expand Down Expand Up @@ -313,7 +294,7 @@ namespace pcl
evecs[1] = unitOrthogonal(evecs[2]);
evecs[0] = cross(evecs[1], evecs[2]);
}
else if (evals.z - evals.y <= numeric_limits<float>::epsilon() )
else if (evals.z - evals.y <= std::numeric_limits<float>::epsilon() )
{
// second and third equal
tmp[0] = row0(); tmp[1] = row1(); tmp[2] = row2();
Expand Down Expand Up @@ -471,7 +452,7 @@ namespace pcl
__device__ __forceinline__ static bool isMuchSmallerThan (float x, float y)
{
// copied from <eigen>/include/Eigen/src/Core/NumTraits.h
const float prec_sqr = numeric_limits<float>::epsilon() * numeric_limits<float>::epsilon();
const float prec_sqr = std::numeric_limits<float>::epsilon() * std::numeric_limits<float>::epsilon();
return x * x <= prec_sqr * y * y;
}
};
Expand Down
2 changes: 1 addition & 1 deletion gpu/kinfu_large_scale/src/cuda/bilateral_pyrdown.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -86,7 +86,7 @@ namespace pcl
}

int res = __float2int_rn (sum1 / sum2);
dst.ptr (y)[x] = max (0, min (res, numeric_limits<short>::max ()));
dst.ptr (y)[x] = max (0, min (res, std::numeric_limits<short>::max ()));
}

//////////////////////////////////////////////////////////////////////////////////////////////////////////////////
Expand Down
2 changes: 0 additions & 2 deletions gpu/kinfu_large_scale/src/cuda/device.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,8 +38,6 @@
#ifndef PCL_GPU_KINFU_DEVICE_HPP_
#define PCL_GPU_KINFU_DEVICE_HPP_

//#include <pcl/gpu/utils/device/limits.hpp>
//#include <pcl/gpu/utils/device/vector_math.hpp>
#include "utils.hpp" //temporary reimplementing to release kinfu without pcl_gpu_utils

#include "internal.h"
Expand Down
2 changes: 1 addition & 1 deletion gpu/kinfu_large_scale/src/cuda/extract.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -615,7 +615,7 @@ namespace pcl

if (idx >= points.size)
return;
const float qnan = numeric_limits<float>::quiet_NaN ();
const float qnan = std::numeric_limits<float>::quiet_NaN ();
float3 n = make_float3 (qnan, qnan, qnan);

float3 point = fetchPoint (idx);
Expand Down
Loading