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;