22 {
return data + x + y*dims.x + z*dims.y*dims.x; }
25 {
return data + x + dims.x * y; }
28 {
return ptr + dims.x * dims.y; }
31 {
return ptr + mult * dims.x * dims.y; }
38 coo.x = __fmaf_rn(f.x, __fdividef(
p.x,
p.z), c.x);
39 coo.y = __fmaf_rn(f.y, __fdividef(
p.y,
p.z), c.y);
48 float x = z * (u - c.x) * finv.x;
49 float y = z * (v - c.y) * finv.y;
50 return make_float3(x, y, z);
57 {
return make_ushort2 (__float2half_rn (tsdf), weight); }
62 return __half2float (value.x);
83 __kf_device__ float operator () (
float l,
float r)
const {
return l + r; }
84 __kf_device__ double operator () (
double l,
double r)
const {
return l + r; }
90 template<
typename T>
__kf_device__ static void StCs(
const T& val, T *ptr);
99 #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200
101 #if defined(_WIN64) || defined(__LP64__)
102 #define _ASM_PTR_ "l"
104 #define _ASM_PTR_ "r"
110 asm(
"ld.global.cs.v2.u16 {%0, %1}, [%2];" :
"=h"(
reinterpret_cast<ushort&
>(val.x)),
"=h"(
reinterpret_cast<ushort&
>(val.y)) : _ASM_PTR_(ptr));
116 short cx = val.x, cy = val.y;
117 asm(
"st.global.cs.v2.u16 [%0], {%1, %2};" : : _ASM_PTR_(ptr),
"h"(
reinterpret_cast<ushort&
>(cx)),
"h"(
reinterpret_cast<ushort&
>(cy)));