38 #ifndef PCL_GPU_KINFU_DEVICE_HPP_ 39 #define PCL_GPU_KINFU_DEVICE_HPP_ 51 #define INV_DIV 3.051850947599719e-5f 53 __device__ __forceinline__
void 58 value = make_short2 (fixedp, weight);
61 __device__ __forceinline__
void 65 tsdf = __int2float_rn (value.x) /
DIVISOR;
68 __device__ __forceinline__
float 71 return static_cast<float>(value.x) /
DIVISOR;
75 __device__ __forceinline__ float3
87 template<ScanKind Kind,
class T>
88 __device__ __forceinline__ T
89 scan_warp (
volatile T *ptr,
const unsigned int idx = threadIdx.x )
91 const unsigned int lane = idx & 31;
93 if (lane >= 1) ptr[idx] = ptr[idx - 1] + ptr[idx];
94 if (lane >= 2) ptr[idx] = ptr[idx - 2] + ptr[idx];
95 if (lane >= 4) ptr[idx] = ptr[idx - 4] + ptr[idx];
96 if (lane >= 8) ptr[idx] = ptr[idx - 8] + ptr[idx];
97 if (lane >= 16) ptr[idx] = ptr[idx - 16] + ptr[idx];
102 return (lane > 0) ? ptr[idx - 1] : 0;
__device__ __forceinline__ float3 operator*(const Mat33 &m, const float3 &vec)
__device__ __forceinline__ void pack_tsdf(float tsdf, int weight, short2 &value)
__device__ __forceinline__ float dot(const float3 &v1, const float3 &v2)
__device__ __forceinline__ void unpack_tsdf(short2 value, float &tsdf, int &weight)
3x3 Matrix for device code
__device__ __forceinline__ T scan_warp(volatile T *ptr, const unsigned int idx=threadIdx.x)