38#ifndef PCL_GPU_KINFU_DEVICE_HPP_
39#define PCL_GPU_KINFU_DEVICE_HPP_
45#include "pointer_shift.cu"
53 #define INV_DIV 3.051850947599719e-5f
55 __device__ __forceinline__
void
73 return static_cast<float>(value.x) /
DIVISOR;
89 template<ScanKind Kind,
class T>
93 const unsigned int lane = idx & 31;
95 if (
lane >= 1) ptr[idx] = ptr[idx - 1] + ptr[idx];
96 if (
lane >= 2) ptr[idx] = ptr[idx - 2] + ptr[idx];
97 if (
lane >= 4) ptr[idx] = ptr[idx - 4] + ptr[idx];
98 if (
lane >= 8) ptr[idx] = ptr[idx - 8] + ptr[idx];
99 if (
lane >= 16) ptr[idx] = ptr[idx - 16] + ptr[idx];
104 return (
lane > 0) ? ptr[idx - 1] : 0;
Iterator class for point clouds with or without given indices.
__device__ __forceinline__ float3 operator*(const Mat33 &m, const float3 &vec)
__device__ __forceinline__ T scan_warp(volatile T *ptr, const unsigned int idx=threadIdx.x)
__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