38 #ifndef PCL_GPU_KINFU_DEVICE_HPP_ 39 #define PCL_GPU_KINFU_DEVICE_HPP_ 47 #include "pointer_shift.cu" 55 #define INV_DIV 3.051850947599719e-5f 57 __device__ __forceinline__
void 62 value = make_short2 (fixedp, weight);
65 __device__ __forceinline__
void 69 tsdf = __int2float_rn (value.x) /
DIVISOR;
72 __device__ __forceinline__
float 75 return static_cast<float>(value.x) /
DIVISOR;
79 __device__ __forceinline__ float3
91 template<ScanKind Kind,
class T>
92 __device__ __forceinline__ T
93 scan_warp (
volatile T *ptr,
const unsigned int idx = threadIdx.x )
95 const unsigned int lane = idx & 31;
97 if (lane >= 1) ptr[idx] = ptr[idx - 1] + ptr[idx];
98 if (lane >= 2) ptr[idx] = ptr[idx - 2] + ptr[idx];
99 if (lane >= 4) ptr[idx] = ptr[idx - 4] + ptr[idx];
100 if (lane >= 8) ptr[idx] = ptr[idx - 8] + ptr[idx];
101 if (lane >= 16) ptr[idx] = ptr[idx - 16] + ptr[idx];
106 return (lane > 0) ? ptr[idx - 1] : 0;
This file defines compatibility wrappers for low level I/O functions.
__device__ __forceinline__ float3 operator*(const Mat33 &m, const float3 &vec)
3x3 Matrix for device code
__device__ __forceinline__ void pack_tsdf(float tsdf, int weight, short2 &value)
__device__ __forceinline__ float dot(const float3 &v1, const float3 &v2)
__device__ __forceinline__ T scan_warp(volatile T *ptr, const unsigned int idx=threadIdx.x)
__device__ __forceinline__ void unpack_tsdf(short2 value, float &tsdf, int &weight)