118 #ifndef NANOVDB_NANOVDB_H_HAS_BEEN_INCLUDED 119 #define NANOVDB_NANOVDB_H_HAS_BEEN_INCLUDED 121 #define NANOVDB_MAGIC_NUMBER 0x304244566f6e614eUL // "NanoVDB0" in hex - little endian (uint64_t) 123 #define NANOVDB_MAJOR_VERSION_NUMBER 32 // reflects changes to the ABI and hence also the file format 124 #define NANOVDB_MINOR_VERSION_NUMBER 4 // reflects changes to the API but not ABI 125 #define NANOVDB_PATCH_VERSION_NUMBER 2 // reflects changes that does not affect the ABI or API 128 #define USE_SINGLE_ROOT_KEY 135 #define NANOVDB_FPN_BRANCHLESS 137 #define NANOVDB_DATA_ALIGNMENT 32 139 #if !defined(NANOVDB_ALIGN) 140 #define NANOVDB_ALIGN(n) alignas(n) 141 #endif // !defined(NANOVDB_ALIGN) 143 #ifdef __CUDACC_RTC__ 145 typedef signed char int8_t;
146 typedef short int16_t;
148 typedef long long int64_t;
149 typedef unsigned char uint8_t;
150 typedef unsigned int uint32_t;
151 typedef unsigned short uint16_t;
152 typedef unsigned long long uint64_t;
154 #define NANOVDB_ASSERT(x) 156 #define UINT64_C(x) (x ## ULL) 158 #else // !__CUDACC_RTC__ 168 #ifdef NANOVDB_USE_IOSTREAMS 173 #define NANOVDB_ASSERT(x) assert(x) 175 #define NANOVDB_ASSERT(x) 178 #if defined(NANOVDB_USE_INTRINSICS) && defined(_MSC_VER) 180 #pragma intrinsic(_BitScanReverse) 181 #pragma intrinsic(_BitScanForward) 182 #pragma intrinsic(_BitScanReverse64) 183 #pragma intrinsic(_BitScanForward64) 186 #endif // __CUDACC_RTC__ 188 #if defined(__CUDACC__) || defined(__HIP__) 190 #define __hostdev__ __host__ __device__ 197 #if defined(_MSC_VER) && defined(__CUDACC__) 198 #define NANOVDB_HOSTDEV_DISABLE_WARNING __pragma("hd_warning_disable") 199 #elif defined(__GNUC__) && defined(__CUDACC__) 200 #define NANOVDB_HOSTDEV_DISABLE_WARNING _Pragma("hd_warning_disable") 202 #define NANOVDB_HOSTDEV_DISABLE_WARNING 206 #define NANOVDB_OFFSETOF(CLASS, MEMBER) ((int)(size_t)((char*)&((CLASS*)0)->MEMBER - (char*)0)) 265 #ifndef __CUDACC_RTC__ 269 static const char * LUT[] = {
"?",
"float",
"double" ,
"int16",
"int32",
270 "int64",
"Vec3f",
"Vec3d",
"Mask",
"Half",
271 "uint32",
"bool",
"RGBA8",
"Float4",
"Float8",
272 "Float16",
"FloatN",
"Vec4f",
"Vec4d",
"Index",
"End" };
273 static_assert(
sizeof(LUT)/
sizeof(
char*) - 1 ==
int(
GridType::End),
"Unexpected size of LUT" );
274 return LUT[
static_cast<int>(gridType)];
292 #ifndef __CUDACC_RTC__ 296 static const char * LUT[] = {
"?",
"SDF",
"FOG" ,
"MAC",
"PNTIDX",
297 "PNTDAT",
"TOPO",
"VOX",
"INDEX",
"END" };
298 static_assert(
sizeof(LUT)/
sizeof(
char*) - 1 ==
int(
GridClass::End),
"Unexpected size of LUT" );
299 return LUT[
static_cast<int>(gridClass)];
316 #ifndef __CUDACC_RTC__ 320 static const char * LUT[] = {
"has long grid name",
324 "has standard deviation",
327 static_assert( 1 << (
sizeof(LUT)/
sizeof(
char*) - 1) ==
int(
GridFlags::End),
"Unexpected size of LUT" );
328 return LUT[
static_cast<int>(gridFlags)];
355 template<
typename T1,
typename T2>
358 static constexpr
bool value =
false;
370 template <
bool,
typename T =
void>
375 template <
typename T>
386 static constexpr
bool value =
false;
425 template<
typename AnyType,
template<
typename...>
class TemplateType>
430 template<
typename... Args,
template<
typename...>
class TemplateType>
439 template <
typename T>
517 template <
typename T>
525 template <
typename T>
529 return reinterpret_cast<const T*
>( (
const uint8_t*)p +
alignmentPadding(p) );
534 template <
typename T1,
typename T2>
538 return reinterpret_cast<const char*
>(p) - reinterpret_cast<const char*>(q);
541 template <
typename DstT,
typename SrcT>
545 return reinterpret_cast<DstT*
>(
reinterpret_cast<char*
>(p) + offset);
548 template <
typename DstT,
typename SrcT>
552 return reinterpret_cast<const DstT*
>(
reinterpret_cast<const char*
>(p) + offset);
565 static const int SIZE = 4;
573 __hostdev__ Rgba8(uint8_t r, uint8_t g, uint8_t b, uint8_t a = 255u) : mData{r, g, b, a} {}
576 : mData{(uint8_t(0.5f + r * 255.0f)),
577 (uint8_t(0.5f + g * 255.0f)),
578 (uint8_t(0.5f + b * 255.0f)),
579 (uint8_t(0.5f + a * 255.0f))}
586 return 0.0000153787005f*(float(mData.c[0])*mData.c[0] +
587 float(mData.c[1])*mData.c[1] +
588 float(mData.c[2])*mData.c[2]);
657 : mData( major << 21 | minor << 10 | patch )
673 #ifndef __CUDACC_RTC__ 676 char *buffer = (
char*)malloc(4 + 1 + 4 + 1 + 4 + 1);
677 snprintf(buffer, 4 + 1 + 4 + 1 + 4 + 1,
"%d.%d.%d", this->getMajor(), this->getMinor(), this->getPatch());
721 #if defined(__CUDA_ARCH__) || defined(__HIP__) 751 template<
typename Type>
757 template<
typename Type>
760 return (a < b) ? a : b;
764 return int32_t(fminf(
float(a),
float(b)));
768 return uint32_t(fminf(
float(a),
float(b)));
778 template<
typename Type>
781 return (a > b) ? a : b;
786 return int32_t(fmaxf(
float(a),
float(b)));
790 return uint32_t(fmaxf(
float(a),
float(b)));
811 return x - floorf(x);
820 return int32_t(floorf(x));
824 return int32_t(floor(x));
829 return int32_t(ceilf(x));
833 return int32_t(ceil(x));
856 return x < 0 ? -x : x;
877 template<
typename CoordT,
typename RealT,
template<
typename>
class Vec3T>
880 template<
typename CoordT,
template<
typename>
class Vec3T>
883 return CoordT(int32_t(rintf(xyz[0])), int32_t(rintf(xyz[1])), int32_t(rintf(xyz[2])));
888 template<
typename CoordT,
template<
typename>
class Vec3T>
891 return CoordT(int32_t(floor(xyz[0] + 0.5)), int32_t(floor(xyz[1] + 0.5)), int32_t(floor(xyz[2] + 0.5)));
894 template<
typename CoordT,
typename RealT,
template<
typename>
class Vec3T>
913 template <
typename T>
914 __hostdev__ inline T
Sign(
const T &x) {
return ((T(0) < x)?T(1):T(0)) - ((x < T(0))?T(1):T(0)); }
916 template<
typename Vec3T>
920 static const int hashTable[8] = {2, 1, 9, 1, 2, 9, 0, 0};
921 const int hashKey = ((v[0] < v[1]) << 2) + ((v[0] < v[2]) << 1) + (v[1] < v[2]);
922 return hashTable[hashKey];
924 if (v[0] < v[1] && v[0] < v[2])
933 template<
typename Vec3T>
937 static const int hashTable[8] = {2, 1, 9, 1, 2, 9, 0, 0};
938 const int hashKey = ((v[0] > v[1]) << 2) + ((v[0] > v[2]) << 1) + (v[1] > v[2]);
939 return hashTable[hashKey];
941 if (v[0] > v[1] && v[0] > v[2])
953 template<u
int64_t wordSize>
956 const uint64_t r = byteCount % wordSize;
957 return r ? byteCount - r + wordSize : byteCount;
992 : mVec{ptr[0], ptr[1], ptr[2]}
1019 template <
typename CoordT>
1022 static_assert(
sizeof(
Coord) ==
sizeof(CoordT),
"Mis-matched sizeof");
1041 return mVec[0] < rhs[0] ?
true : mVec[0] > rhs[0] ?
false : mVec[1] < rhs[1] ?
true : mVec[1] > rhs[1] ?
false : mVec[2] < rhs[2] ?
true :
false;
1095 if (other[0] < mVec[0])
1097 if (other[1] < mVec[1])
1099 if (other[2] < mVec[2])
1107 if (other[0] > mVec[0])
1109 if (other[1] > mVec[1])
1111 if (other[2] > mVec[2])
1118 return Coord(mVec[0] + dx, mVec[1] + dy, mVec[2] + dz);
1127 return (a[0] < b[0] || a[1] < b[1] || a[2] < b[2]);
1132 template<
typename Vec3T>
1137 template<
int Log2N = 3 + 4 + 5>
1138 __hostdev__ uint32_t
hash()
const {
return ((1 << Log2N) - 1) & (mVec[0] * 73856093 ^ mVec[1] * 19349663 ^ mVec[2] * 83492791); }
1143 (uint8_t(
bool(mVec[1] & (1u << 31))) << 1) |
1144 (uint8_t(
bool(mVec[2] & (1u << 31))) << 2)); }
1156 template<
typename T>
1162 static const int SIZE = 3;
1173 template<
typename T2>
1175 : mVec{T(v[0]), T(v[1]), T(v[2])}
1179 : mVec{T(ijk[0]), T(ijk[1]), T(ijk[2])}
1184 template<
typename Vec3T>
1194 template<
typename Vec3T>
1195 __hostdev__ T
dot(
const Vec3T& v)
const {
return mVec[0] * v[0] + mVec[1] * v[1] + mVec[2] * v[2]; }
1196 template<
typename Vec3T>
1199 return Vec3(mVec[1] * v[2] - mVec[2] * v[1],
1200 mVec[2] * v[0] - mVec[0] * v[2],
1201 mVec[0] * v[1] - mVec[1] * v[0]);
1205 return mVec[0] * mVec[0] + mVec[1] * mVec[1] + mVec[2] * mVec[2];
1241 if (other[0] < mVec[0])
1243 if (other[1] < mVec[1])
1245 if (other[2] < mVec[2])
1253 if (other[0] > mVec[0])
1255 if (other[1] > mVec[1])
1257 if (other[2] > mVec[2])
1264 return mVec[0] < mVec[1] ? (mVec[0] < mVec[2] ? mVec[0] : mVec[2]) : (mVec[1] < mVec[2] ? mVec[1] : mVec[2]);
1269 return mVec[0] > mVec[1] ? (mVec[0] > mVec[2] ? mVec[0] : mVec[2]) : (mVec[1] > mVec[2] ? mVec[1] : mVec[2]);
1276 template<
typename T1,
typename T2>
1279 return Vec3<T2>(scalar * vec[0], scalar * vec[1], scalar * vec[2]);
1281 template<
typename T1,
typename T2>
1284 return Vec3<T2>(scalar / vec[0], scalar / vec[1], scalar / vec[2]);
1301 template<
typename T>
1307 static const int SIZE = 4;
1318 template<
typename T2>
1320 : mVec{T(v[0]), T(v[1]), T(v[2]), T(v[3])}
1323 __hostdev__ bool operator==(
const Vec4& rhs)
const {
return mVec[0] == rhs[0] && mVec[1] == rhs[1] && mVec[2] == rhs[2] && mVec[3] == rhs[3]; }
1324 __hostdev__ bool operator!=(
const Vec4& rhs)
const {
return mVec[0] != rhs[0] || mVec[1] != rhs[1] || mVec[2] != rhs[2] || mVec[3] != rhs[3]; }
1325 template<
typename Vec4T>
1336 template<
typename Vec4T>
1337 __hostdev__ T
dot(
const Vec4T& v)
const {
return mVec[0] * v[0] + mVec[1] * v[1] + mVec[2] * v[2] + mVec[3] * v[3]; }
1340 return mVec[0] * mVec[0] + mVec[1] * mVec[1] + mVec[2] * mVec[2] + mVec[3] * mVec[3];
1379 if (other[0] < mVec[0])
1381 if (other[1] < mVec[1])
1383 if (other[2] < mVec[2])
1385 if (other[3] < mVec[3])
1393 if (other[0] > mVec[0])
1395 if (other[1] > mVec[1])
1397 if (other[2] > mVec[2])
1399 if (other[3] > mVec[3])
1405 template<
typename T1,
typename T2>
1408 return Vec4<T2>(scalar * vec[0], scalar * vec[1], scalar * vec[2], scalar * vec[3]);
1410 template<
typename T1,
typename T2>
1413 return Vec4<T2>(scalar / vec[0], scalar / vec[1], scalar / vec[2], scalar / vec[3]);
1428 template<
typename T>
1431 static const int Rank = 0;
1432 static const bool IsScalar =
true;
1433 static const bool IsVector =
false;
1434 static const int Size = 1;
1439 template<
typename T>
1442 static const int Rank = 1;
1443 static const bool IsScalar =
false;
1444 static const bool IsVector =
true;
1445 static const int Size = T::SIZE;
1452 template<typename T, int = sizeof(typename TensorTraits<T>::ElementType)>
1458 template<
typename T>
1485 template<
typename BuildT>
1530 template<
typename Vec3T>
1533 return Vec3T(fmaf(xyz[0], mat[0], fmaf(xyz[1], mat[1], xyz[2] * mat[2])),
1534 fmaf(xyz[0], mat[3], fmaf(xyz[1], mat[4], xyz[2] * mat[5])),
1535 fmaf(xyz[0], mat[6], fmaf(xyz[1], mat[7], xyz[2] * mat[8])));
1538 template<
typename Vec3T>
1541 return Vec3T(fma(static_cast<double>(xyz[0]), mat[0], fma(static_cast<double>(xyz[1]), mat[1], static_cast<double>(xyz[2]) * mat[2])),
1542 fma(static_cast<double>(xyz[0]), mat[3], fma(static_cast<double>(xyz[1]), mat[4], static_cast<double>(xyz[2]) * mat[5])),
1543 fma(static_cast<double>(xyz[0]), mat[6], fma(static_cast<double>(xyz[1]), mat[7], static_cast<double>(xyz[2]) * mat[8])));
1546 template<
typename Vec3T>
1549 return Vec3T(fmaf(xyz[0], mat[0], fmaf(xyz[1], mat[1], fmaf(xyz[2], mat[2], vec[0]))),
1550 fmaf(xyz[0], mat[3], fmaf(xyz[1], mat[4], fmaf(xyz[2], mat[5], vec[1]))),
1551 fmaf(xyz[0], mat[6], fmaf(xyz[1], mat[7], fmaf(xyz[2], mat[8], vec[2]))));
1554 template<
typename Vec3T>
1557 return Vec3T(fma(static_cast<double>(xyz[0]), mat[0], fma(static_cast<double>(xyz[1]), mat[1], fma(static_cast<double>(xyz[2]), mat[2], vec[0]))),
1558 fma(static_cast<double>(xyz[0]), mat[3], fma(static_cast<double>(xyz[1]), mat[4], fma(static_cast<double>(xyz[2]), mat[5], vec[1]))),
1559 fma(static_cast<double>(xyz[0]), mat[6], fma(static_cast<double>(xyz[1]), mat[7], fma(static_cast<double>(xyz[2]), mat[8], vec[2]))));
1564 template<
typename Vec3T>
1567 return Vec3T(fmaf(xyz[0], mat[0], fmaf(xyz[1], mat[3], xyz[2] * mat[6])),
1568 fmaf(xyz[0], mat[1], fmaf(xyz[1], mat[4], xyz[2] * mat[7])),
1569 fmaf(xyz[0], mat[2], fmaf(xyz[1], mat[5], xyz[2] * mat[8])));
1572 template<
typename Vec3T>
1575 return Vec3T(fma(static_cast<double>(xyz[0]), mat[0], fma(static_cast<double>(xyz[1]), mat[3], static_cast<double>(xyz[2]) * mat[6])),
1576 fma(static_cast<double>(xyz[0]), mat[1], fma(static_cast<double>(xyz[1]), mat[4], static_cast<double>(xyz[2]) * mat[7])),
1577 fma(static_cast<double>(xyz[0]), mat[2], fma(static_cast<double>(xyz[1]), mat[5], static_cast<double>(xyz[2]) * mat[8])));
1580 template<
typename Vec3T>
1583 return Vec3T(fmaf(xyz[0], mat[0], fmaf(xyz[1], mat[3], fmaf(xyz[2], mat[6], vec[0]))),
1584 fmaf(xyz[0], mat[1], fmaf(xyz[1], mat[4], fmaf(xyz[2], mat[7], vec[1]))),
1585 fmaf(xyz[0], mat[2], fmaf(xyz[1], mat[5], fmaf(xyz[2], mat[8], vec[2]))));
1588 template<
typename Vec3T>
1591 return Vec3T(fma(static_cast<double>(xyz[0]), mat[0], fma(static_cast<double>(xyz[1]), mat[3], fma(static_cast<double>(xyz[2]), mat[6], vec[0]))),
1592 fma(static_cast<double>(xyz[0]), mat[1], fma(static_cast<double>(xyz[1]), mat[4], fma(static_cast<double>(xyz[2]), mat[7], vec[1]))),
1593 fma(static_cast<double>(xyz[0]), mat[2], fma(static_cast<double>(xyz[1]), mat[5], fma(static_cast<double>(xyz[2]), mat[8], vec[2]))));
1599 template<
typename Vec3T>
1620 mCoord[0].minComponent(xyz);
1621 mCoord[1].maxComponent(xyz);
1628 mCoord[0].maxComponent(bbox.
min());
1629 mCoord[1].minComponent(bbox.
max());
1639 if (xyz[0] < mCoord[0][0] || xyz[1] < mCoord[0][1] || xyz[2] < mCoord[0][2])
1641 if (xyz[0] > mCoord[1][0] || xyz[1] > mCoord[1][1] || xyz[2] > mCoord[1][2])
1661 template<
typename Vec3T>
1668 using BaseT::mCoord;
1690 mCoord[0][1] >= mCoord[1][1] ||
1691 mCoord[0][2] >= mCoord[1][2]; }
1695 return p[0] > mCoord[0][0] && p[1] > mCoord[0][1] && p[2] > mCoord[0][2] &&
1696 p[0] < mCoord[1][0] && p[1] < mCoord[1][1] && p[2] < mCoord[1][2];
1705 template<
typename CoordT>
1710 using BaseT::mCoord;
1725 if (mPos[2] < mBBox[1][2]) {
1727 }
else if (mPos[1] < mBBox[1][1]) {
1728 mPos[2] = mBBox[0][2];
1730 }
else if (mPos[0] <= mBBox[1][0]) {
1731 mPos[2] = mBBox[0][2];
1732 mPos[1] = mBBox[0][1];
1757 template<
typename SplitT>
1759 :
BaseT(other.mCoord[0], other.mCoord[1])
1762 const int n =
MaxIndex(this->dim());
1763 mCoord[1][n] = (mCoord[0][n] + mCoord[1][n]) >> 1;
1764 other.mCoord[0][n] = mCoord[1][n] + 1;
1769 return BBox(min, min.offsetBy(dim - 1));
1773 mCoord[0][1] < mCoord[1][1] &&
1774 mCoord[0][2] < mCoord[1][2]; }
1777 mCoord[0][1] > mCoord[1][1] ||
1778 mCoord[0][2] > mCoord[1][2]; }
1780 __hostdev__ uint64_t
volume()
const {
auto d = this->dim();
return uint64_t(d[0])*uint64_t(d[1])*uint64_t(d[2]); }
1785 return !(CoordT::lessThan(b.min(), this->
min()) || CoordT::lessThan(this->
max(), b.max()));
1791 return !(CoordT::lessThan(this->
max(), b.min()) || CoordT::lessThan(b.max(), this->
min()));
1795 template<
typename RealT>
1800 Vec3<RealT>(RealT(mCoord[1][0] + 1), RealT(mCoord[1][1] + 1), RealT(mCoord[1][2] + 1)));
1805 return BBox(mCoord[0].offsetBy(-padding), mCoord[1].offsetBy(padding));
1821 #if (defined(__CUDA_ARCH__) || defined(__HIP__)) && defined(NANOVDB_USE_INTRINSICS) 1823 #elif defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS) 1824 unsigned long index;
1825 _BitScanForward(&index, v);
1826 return static_cast<uint32_t
>(index);
1827 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS) 1828 return static_cast<uint32_t
>(__builtin_ctzl(v));
1831 static const unsigned char DeBruijn[32] = {
1832 0, 1, 28, 2, 29, 14, 24, 3, 30, 22, 20, 15, 25, 17, 4, 8, 31, 27, 13, 23, 21, 19, 16, 7, 26, 12, 18, 6, 11, 5, 10, 9};
1834 #if defined(_MSC_VER) && !defined(__NVCC__) 1835 #pragma warning(push) 1836 #pragma warning(disable : 4146) 1838 return DeBruijn[uint32_t((v & -v) * 0x077CB531U) >> 27];
1839 #if defined(_MSC_VER) && !defined(__NVCC__) 1840 #pragma warning(pop) 1853 #if defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS) 1854 unsigned long index;
1855 _BitScanReverse(&index, v);
1856 return static_cast<uint32_t
>(index);
1857 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS) 1858 return sizeof(
unsigned long) * 8 - 1 - __builtin_clzl(v);
1861 static const unsigned char DeBruijn[32] = {
1862 0, 9, 1, 10, 13, 21, 2, 29, 11, 14, 16, 18, 22, 25, 3, 30, 8, 12, 20, 28, 15, 17, 24, 7, 19, 27, 23, 6, 26, 5, 4, 31};
1868 return DeBruijn[uint32_t(v * 0x07C4ACDDU) >> 27];
1879 #if (defined(__CUDA_ARCH__) || defined(__HIP__)) && defined(NANOVDB_USE_INTRINSICS) 1881 #elif defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS) 1882 unsigned long index;
1883 _BitScanForward64(&index, v);
1884 return static_cast<uint32_t
>(index);
1885 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS) 1886 return static_cast<uint32_t
>(__builtin_ctzll(v));
1889 static const unsigned char DeBruijn[64] = {
1890 0, 1, 2, 53, 3, 7, 54, 27, 4, 38, 41, 8, 34, 55, 48, 28,
1891 62, 5, 39, 46, 44, 42, 22, 9, 24, 35, 59, 56, 49, 18, 29, 11,
1892 63, 52, 6, 26, 37, 40, 33, 47, 61, 45, 43, 21, 23, 58, 17, 10,
1893 51, 25, 36, 32, 60, 20, 57, 16, 50, 31, 19, 15, 30, 14, 13, 12,
1896 #if defined(_MSC_VER) && !defined(__NVCC__) 1897 #pragma warning(push) 1898 #pragma warning(disable : 4146) 1900 return DeBruijn[uint64_t((v & -v) * UINT64_C(0x022FDD63CC95386D)) >> 58];
1901 #if defined(_MSC_VER) && !defined(__NVCC__) 1902 #pragma warning(pop) 1915 #if defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS) 1916 unsigned long index;
1917 _BitScanReverse64(&index, v);
1918 return static_cast<uint32_t
>(index);
1919 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS) 1920 return sizeof(
unsigned long) * 8 - 1 - __builtin_clzll(v);
1922 const uint32_t* p =
reinterpret_cast<const uint32_t*
>(&v);
1933 #if (defined(__CUDA_ARCH__) || defined(__HIP__)) && defined(NANOVDB_USE_INTRINSICS) 1937 #elif defined(_MSC_VER) && defined(_M_X64) && (_MSC_VER >= 1928) && defined(NANOVDB_USE_INTRINSICS) 1939 return __popcnt64(v);
1940 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS) 1942 return __builtin_popcountll(v);
1943 #else// use software implementation 1945 v = v - ((v >> 1) & uint64_t(0x5555555555555555));
1946 v = (v & uint64_t(0x3333333333333333)) + ((v >> 2) & uint64_t(0x3333333333333333));
1947 return (((v + (v >> 4)) & uint64_t(0xF0F0F0F0F0F0F0F)) * uint64_t(0x101010101010101)) >> 56;
1955 template<u
int32_t LOG2DIM>
1958 static constexpr uint32_t SIZE = 1U << (3 * LOG2DIM);
1959 static constexpr uint32_t WORD_COUNT = SIZE >> 6;
1960 uint64_t mWords[WORD_COUNT];
1975 uint32_t sum = 0, n = WORD_COUNT;
1976 for (
const uint64_t* w = mWords; n--; ++w)
1984 uint32_t n = i >> 6, sum =
CountOn( mWords[n] & ((uint64_t(1) << (i & 63u))-1u) );
1985 for (
const uint64_t* w = mWords; n--; ++w) sum +=
CountOn(*w);
2001 mPos = mParent->findNext<On>(mPos + 1);
2013 const Mask* mParent;
2026 for (uint32_t i = 0; i < WORD_COUNT; ++i)
2031 const uint64_t v = on ? ~uint64_t(0) : uint64_t(0);
2032 for (uint32_t i = 0; i < WORD_COUNT; ++i)
2039 for (uint32_t i = 0; i < WORD_COUNT; ++i)
2040 mWords[i] = other.mWords[i];
2044 template<
typename WordT>
2048 return reinterpret_cast<const WordT*
>(mWords)[n];
2052 template<
typename WordT>
2056 return reinterpret_cast<WordT*
>(mWords)[n];
2060 template<
typename MaskT>
2063 static_assert(
sizeof(
Mask) ==
sizeof(MaskT),
"Mismatching sizeof");
2064 static_assert(WORD_COUNT == MaskT::WORD_COUNT,
"Mismatching word count");
2065 static_assert(LOG2DIM == MaskT::LOG2DIM,
"Mismatching LOG2DIM");
2066 auto *src =
reinterpret_cast<const uint64_t*
>(&other);
2067 uint64_t *dst = mWords;
2068 for (uint32_t i = 0; i < WORD_COUNT; ++i) {
2076 for (uint32_t i = 0; i < WORD_COUNT; ++i) {
2077 if (mWords[i] != other.mWords[i])
return false;
2085 __hostdev__ bool isOn(uint32_t n)
const {
return 0 != (mWords[n >> 6] & (uint64_t(1) << (n & 63))); }
2088 __hostdev__ bool isOff(uint32_t n)
const {
return 0 == (mWords[n >> 6] & (uint64_t(1) << (n & 63))); }
2093 for (uint32_t i = 0; i < WORD_COUNT; ++i)
2094 if (mWords[i] != ~uint64_t(0))
2102 for (uint32_t i = 0; i < WORD_COUNT; ++i)
2103 if (mWords[i] != uint64_t(0))
2117 #if 1 // switch between branchless 2118 auto &word = mWords[n >> 6];
2120 word &= ~(uint64_t(1) << n);
2121 word |= uint64_t(On) << n;
2123 On ? this->setOn(n) : this->setOff(n);
2130 for (uint32_t i = 0; i < WORD_COUNT; ++i)
2131 mWords[i] = ~uint64_t(0);
2137 for (uint32_t i = 0; i < WORD_COUNT; ++i)
2138 mWords[i] = uint64_t(0);
2144 const uint64_t v = on ? ~uint64_t(0) : uint64_t(0);
2145 for (uint32_t i = 0; i < WORD_COUNT; ++i)
2151 uint32_t n = WORD_COUNT;
2152 for (
auto* w = mWords; n--; ++w)
2160 uint64_t *w1 = mWords;
2161 const uint64_t *w2 = other.mWords;
2162 for (uint32_t n = WORD_COUNT; n--; ++w1, ++w2) *w1 &= *w2;
2168 uint64_t *w1 = mWords;
2169 const uint64_t *w2 = other.mWords;
2170 for (uint32_t n = WORD_COUNT; n--; ++w1, ++w2) *w1 |= *w2;
2176 uint64_t *w1 = mWords;
2177 const uint64_t *w2 = other.mWords;
2178 for (uint32_t n = WORD_COUNT; n--; ++w1, ++w2) *w1 &= ~*w2;
2184 uint64_t *w1 = mWords;
2185 const uint64_t *w2 = other.mWords;
2186 for (uint32_t n = WORD_COUNT; n--; ++w1, ++w2) *w1 ^= *w2;
2197 const uint64_t* w = mWords;
2198 for (; n<WORD_COUNT && !(On ? *w : ~*w); ++w, ++n);
2199 return n==WORD_COUNT ? SIZE : (n << 6) +
FindLowestOn(On ? *w : ~*w);
2204 __hostdev__ uint32_t findNext(uint32_t start)
const 2206 uint32_t n = start >> 6;
2207 if (n >= WORD_COUNT)
2209 uint32_t m = start & 63;
2210 uint64_t b = On ? mWords[n] : ~mWords[n];
2211 if (b & (uint64_t(1) << m))
2213 b &= ~uint64_t(0) << m;
2214 while (!b && ++n < WORD_COUNT)
2215 b = On ? mWords[n] : ~mWords[n];
2235 template<
typename Mat3T,
typename Vec3T>
2236 __hostdev__ void set(
const Mat3T& mat,
const Mat3T& invMat,
const Vec3T& translate,
double taper);
2240 template<
typename Mat4T>
2241 __hostdev__ void set(
const Mat4T& mat,
const Mat4T& invMat,
double taper) {this->
set(mat, invMat, mat[3], taper);}
2243 template<
typename Vec3T>
2244 __hostdev__ void set(
double scale,
const Vec3T &translation,
double taper);
2246 template<
typename Vec3T>
2248 template<
typename Vec3T>
2251 template<
typename Vec3T>
2253 template<
typename Vec3T>
2256 template<
typename Vec3T>
2259 return matMult(mInvMatD, Vec3T(xyz[0] - mVecD[0], xyz[1] - mVecD[1], xyz[2] - mVecD[2]));
2261 template<
typename Vec3T>
2264 return matMult(mInvMatF, Vec3T(xyz[0] - mVecF[0], xyz[1] - mVecF[1], xyz[2] - mVecF[2]));
2267 template<
typename Vec3T>
2269 template<
typename Vec3T>
2272 template<
typename Vec3T>
2274 template<
typename Vec3T>
2278 template<
typename Mat3T,
typename Vec3T>
2281 float *mf = mMatF, *vf = mVecF, *mif = mInvMatF;
2282 double *md = mMatD, *vd = mVecD, *mid = mInvMatD;
2283 mTaperF =
static_cast<float>(taper);
2285 for (
int i = 0; i < 3; ++i) {
2286 *vd++ = translate[i];
2287 *vf++ =
static_cast<float>(translate[i]);
2288 for (
int j = 0; j < 3; ++j) {
2290 *mid++ = invMat[j][i];
2291 *mf++ =
static_cast<float>(mat[j][i]);
2292 *mif++ =
static_cast<float>(invMat[j][i]);
2297 template<
typename Vec3T>
2300 const double mat[3][3] = {
2304 }, idx = 1.0/dx, invMat[3][3] = {
2309 this->
set(mat, invMat, trans, taper);
2316 static const int MaxNameSize = 256;
2323 char mName[MaxNameSize];
2333 template <
typename T>
2342 template<
typename Gr
idOrTreeOrRootT,
int LEVEL>
2346 template<
typename Gr
idOrTreeOrRootT>
2349 static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3,
"Tree depth is not supported");
2350 using Type =
typename GridOrTreeOrRootT::LeafNodeType;
2351 using type =
typename GridOrTreeOrRootT::LeafNodeType;
2353 template<
typename Gr
idOrTreeOrRootT>
2356 static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3,
"Tree depth is not supported");
2357 using Type =
const typename GridOrTreeOrRootT::LeafNodeType;
2358 using type =
const typename GridOrTreeOrRootT::LeafNodeType;
2361 template<
typename Gr
idOrTreeOrRootT>
2364 static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3,
"Tree depth is not supported");
2365 using Type =
typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType;
2366 using type =
typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType;
2368 template<
typename Gr
idOrTreeOrRootT>
2371 static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3,
"Tree depth is not supported");
2372 using Type =
const typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType;
2373 using type =
const typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType;
2375 template<
typename Gr
idOrTreeOrRootT>
2378 static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3,
"Tree depth is not supported");
2379 using Type =
typename GridOrTreeOrRootT::RootType::ChildNodeType;
2380 using type =
typename GridOrTreeOrRootT::RootType::ChildNodeType;
2382 template<
typename Gr
idOrTreeOrRootT>
2385 static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3,
"Tree depth is not supported");
2386 using Type =
const typename GridOrTreeOrRootT::RootType::ChildNodeType;
2387 using type =
const typename GridOrTreeOrRootT::RootType::ChildNodeType;
2389 template<
typename Gr
idOrTreeOrRootT>
2392 static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3,
"Tree depth is not supported");
2393 using Type =
typename GridOrTreeOrRootT::RootType;
2394 using type =
typename GridOrTreeOrRootT::RootType;
2397 template<
typename Gr
idOrTreeOrRootT>
2400 static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3,
"Tree depth is not supported");
2401 using Type =
const typename GridOrTreeOrRootT::RootType;
2402 using type =
const typename GridOrTreeOrRootT::RootType;
2433 static const int MaxNameSize = 256;
2441 char mGridName[MaxNameSize];
2504 template<
typename Vec3T>
2506 template<
typename Vec3T>
2508 template<
typename Vec3T>
2510 template<
typename Vec3T>
2512 template<
typename Vec3T>
2515 template<
typename Vec3T>
2517 template<
typename Vec3T>
2519 template<
typename Vec3T>
2521 template<
typename Vec3T>
2523 template<
typename Vec3T>
2538 return PtrAdd<GridBlindMetaData>(
this, mBlindMetadataOffset) + n;
2544 template <
typename BuildT,
int LEVEL0 = -1,
int LEVEL1 = -1,
int LEVEL2 = -1>
2547 template <
typename BuildT>
2554 template<
typename TreeT>
2570 Grid& operator=(
const Grid&) =
delete;
2594 template <
typename T = BuildType>
2598 __hostdev__ const TreeT&
tree()
const {
return *
reinterpret_cast<const TreeT*
>(this->treePtr()); }
2613 template<
typename Vec3T>
2617 template<
typename Vec3T>
2622 template<
typename Vec3T>
2627 template<
typename Vec3T>
2632 template<
typename Vec3T>
2636 template<
typename Vec3T>
2640 template<
typename Vec3T>
2645 template<
typename Vec3T>
2650 template<
typename Vec3T>
2655 template<
typename Vec3T>
2691 template <
typename NodeT>
2696 template <
int LEVEL>
2702 if (this->hasLongGridName()) {
2704 const auto &metaData = this->blindMetaData(DataType::mBlindMetadataCount-1);
2706 return metaData.template getBlindData<const char>();
2708 return DataType::mGridName;
2731 if (DataType::mBlindMetadataCount == 0u) {
2735 return this->blindMetaData(n).template getBlindData<void>();
2744 template<
typename TreeT>
2747 for (uint32_t i = 0, n = this->blindDataCount(); i < n; ++i)
2748 if (this->blindMetaData(i).mSemantic == semantic)
2755 template<
int ROOT_LEVEL = 3>
2758 static_assert(
ROOT_LEVEL == 3,
"Root level is assumed to be three");
2759 uint64_t mNodeOffset[4];
2760 uint32_t mNodeCount[3];
2761 uint32_t mTileCount[3];
2764 template <
typename RootT>
2766 template <
typename RootT>
2768 template <
typename RootT>
2771 template <
typename NodeT>
2774 mNodeOffset[NodeT::LEVEL] = node ?
PtrDiff(node,
this) : 0;
2781 template<
typename Gr
idT>
2784 using Type =
typename GridT::TreeType;
2785 using type =
typename GridT::TreeType;
2787 template<
typename Gr
idT>
2790 using Type =
const typename GridT::TreeType;
2791 using type =
const typename GridT::TreeType;
2797 template<
typename RootT>
2800 static_assert(RootT::LEVEL == 3,
"Tree depth is not supported");
2801 static_assert(RootT::ChildNodeType::LOG2DIM == 5,
"Tree configuration is not supported");
2802 static_assert(RootT::ChildNodeType::ChildNodeType::LOG2DIM == 4,
"Tree configuration is not supported");
2803 static_assert(RootT::LeafNodeType::LOG2DIM == 3,
"Tree configuration is not supported");
2815 using Node2 =
typename RootT::ChildNodeType;
2816 using Node1 =
typename Node2::ChildNodeType;
2822 Tree& operator=(
const Tree&) =
delete;
2870 return DataType::mTileCount[level - 1];
2873 template<
typename NodeT>
2876 static_assert(NodeT::LEVEL < 3,
"Invalid NodeT");
2877 return DataType::mNodeCount[NodeT::LEVEL];
2883 return DataType::mNodeCount[level];
2889 template <
typename NodeT>
2892 const uint64_t offset = DataType::mNodeOffset[NodeT::LEVEL];
2893 return offset>0 ? PtrAdd<NodeT>(
this, offset) :
nullptr;
2899 template <
typename NodeT>
2902 const uint64_t offset = DataType::mNodeOffset[NodeT::LEVEL];
2903 return offset>0 ? PtrAdd<NodeT>(
this, offset) :
nullptr;
2909 template <
int LEVEL>
2913 return this->
template getFirstNode<typename NodeTrait<RootT,LEVEL>::type>();
2919 template <
int LEVEL>
2923 return this->
template getFirstNode<typename NodeTrait<RootT,LEVEL>::type>();
2939 template<
typename RootT>
2942 min = this->root().minimum();
2943 max = this->root().maximum();
2951 template<
typename ChildT>
2958 static constexpr
bool FIXED_SIZE =
false;
2961 #ifdef USE_SINGLE_ROOT_KEY 2963 template <
typename CoordType>
2966 static_assert(
sizeof(
CoordT) ==
sizeof(CoordType),
"Mismatching sizeof");
2967 static_assert(32 - ChildT::TOTAL <= 21,
"Cannot use 64 bit root keys");
2968 return (
KeyT(uint32_t(ijk[2]) >> ChildT::TOTAL)) |
2969 (
KeyT(uint32_t(ijk[1]) >> ChildT::TOTAL) << 21) |
2970 (
KeyT(uint32_t(ijk[0]) >> ChildT::TOTAL) << 42);
2974 static constexpr uint64_t MASK = (1u << 21) - 1;
2975 return CoordT(((key >> 42) & MASK) << ChildT::TOTAL,
2976 ((key >> 21) & MASK) << ChildT::TOTAL,
2977 (key & MASK) << ChildT::TOTAL);
3000 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
Tile 3002 template <
typename CoordType>
3005 key = CoordToKey(k);
3008 template <
typename CoordType,
typename ValueType>
3011 key = CoordToKey(k);
3032 return reinterpret_cast<const Tile*
>(
this + 1) + n;
3037 return reinterpret_cast<Tile*
>(
this + 1) + n;
3046 return PtrAdd<ChildT>(
this, tile->
child);
3051 return PtrAdd<ChildT>(
this, tile->
child);
3072 template<
typename ChildT>
3089 static constexpr
bool FIXED_SIZE = DataType::FIXED_SIZE;
3091 static constexpr uint32_t LEVEL = 1 + ChildT::LEVEL;
3096 uint32_t mPos, mSize;
3101 while (mPos<mSize && !mParent->tile(mPos)->isChild()) ++mPos;
3112 while (mPos < mSize && mParent->tile(mPos)->isValue()) ++mPos;
3127 uint32_t mPos, mSize;
3132 while (mPos < mSize && mParent->tile(mPos)->isChild()) ++mPos;
3143 while (mPos < mSize && mParent->tile(mPos)->isChild()) ++mPos;
3158 uint32_t mPos, mSize;
3163 while (mPos < mSize && !mParent->tile(mPos)->isActive()) ++mPos;
3173 while (mPos < mSize && !mParent->tile(mPos)->isActive()) ++mPos;
3233 if (
const Tile* tile = this->probeTile(ijk)) {
3234 return tile->isChild() ? this->getChild(tile)->getValue(ijk) : tile->value;
3236 return DataType::mBackground;
3241 if (
const Tile* tile = this->probeTile(ijk)) {
3242 return tile->isChild() ? this->getChild(tile)->isActive(ijk) : tile->state;
3252 if (
const Tile* tile = this->probeTile(ijk)) {
3253 if (tile->isChild()) {
3254 const auto *
child = this->getChild(tile);
3255 return child->probeValue(ijk, v);
3260 v = DataType::mBackground;
3266 const Tile* tile = this->probeTile(ijk);
3267 if (tile && tile->isChild()) {
3268 const auto *
child = this->getChild(tile);
3269 return child->probeLeaf(ijk);
3276 const Tile* tile = this->probeTile(ijk);
3277 if (tile && tile->isChild()) {
3278 return this->getChild(tile);
3286 const Tile* tiles =
reinterpret_cast<const Tile*
>(
this + 1);
3287 const auto key = DataType::CoordToKey(ijk);
3288 #if 1 // switch between linear and binary seach 3289 for (uint32_t i = 0; i < DataType::mTableSize; ++i) {
3290 if (tiles[i].key == key)
return &tiles[i];
3292 #else// do not enable binary search if tiles are not guaranteed to be sorted!!!!!! 3294 int32_t low = 0, high = DataType::mTableSize;
3295 while (low != high) {
3296 int mid = low + ((high - low) >> 1);
3297 const Tile* tile = &tiles[mid];
3298 if (tile->key == key) {
3300 }
else if (tile->key < key) {
3314 template<
typename,
int,
int,
int>
3321 template<
typename AccT>
3322 __hostdev__ typename AccT::NodeInfo getNodeInfoAndCache(
const CoordType& ijk,
const AccT& acc)
const 3324 using NodeInfoT =
typename AccT::NodeInfo;
3325 if (
const Tile* tile = this->probeTile(ijk)) {
3326 if (tile->isChild()) {
3327 const auto *
child = this->getChild(tile);
3328 acc.insert(ijk,
child);
3329 return child->getNodeInfoAndCache(ijk, acc);
3331 return NodeInfoT{LEVEL, ChildT::dim(), tile->value, tile->value, tile->value,
3332 0, tile->origin(), tile->origin() +
CoordType(ChildT::DIM)};
3334 return NodeInfoT{LEVEL, ChildT::dim(), this->minimum(), this->maximum(),
3335 this->average(), this->stdDeviation(), this->bbox()[0], this->bbox()[1]};
3339 template<
typename AccT>
3342 if (
const Tile* tile = this->probeTile(ijk)) {
3343 if (tile->isChild()) {
3344 const auto *
child = this->getChild(tile);
3345 acc.insert(ijk,
child);
3346 return child->getValueAndCache(ijk, acc);
3350 return DataType::mBackground;
3353 template<
typename AccT>
3356 const Tile* tile = this->probeTile(ijk);
3357 if (tile && tile->isChild()) {
3358 const auto *
child = this->getChild(tile);
3359 acc.insert(ijk,
child);
3360 return child->isActiveAndCache(ijk, acc);
3365 template<
typename AccT>
3368 if (
const Tile* tile = this->probeTile(ijk)) {
3369 if (tile->isChild()) {
3370 const auto *
child = this->getChild(tile);
3371 acc.insert(ijk,
child);
3372 return child->probeValueAndCache(ijk, v, acc);
3377 v = DataType::mBackground;
3381 template<
typename AccT>
3384 const Tile* tile = this->probeTile(ijk);
3385 if (tile && tile->isChild()) {
3386 const auto *
child = this->getChild(tile);
3387 acc.insert(ijk,
child);
3388 return child->probeLeafAndCache(ijk, acc);
3393 template<
typename RayT,
typename AccT>
3394 __hostdev__ uint32_t getDimAndCache(
const CoordType& ijk,
const RayT& ray,
const AccT& acc)
const 3396 if (
const Tile* tile = this->probeTile(ijk)) {
3397 if (tile->isChild()) {
3398 const auto *
child = this->getChild(tile);
3399 acc.insert(ijk,
child);
3400 return child->getDimAndCache(ijk, ray, acc);
3402 return 1 << ChildT::TOTAL;
3404 return ChildNodeType::dim();
3416 template<
typename ChildT, u
int32_t LOG2DIM>
3423 using MaskT =
typename ChildT::template MaskType<LOG2DIM>;
3424 static constexpr
bool FIXED_SIZE =
true;
3433 Tile& operator=(
const Tile&) =
delete;
3453 + (1u << (3 * LOG2DIM))*(sizeof(ValueT) > 8u ?
sizeof(
ValueT) : 8u));
3455 alignas(32) Tile mTable[1u << (3 * LOG2DIM)];
3462 mTable[n].child =
PtrDiff(ptr,
this);
3465 template <
typename ValueT>
3469 mTable[n].value = v;
3476 return PtrAdd<ChildT>(
this, mTable[n].child);
3481 return PtrAdd<ChildT>(
this, mTable[n].child);
3487 return mTable[n].value;
3493 return mValueMask.isOn(n);
3498 template <
typename T>
3519 template<
typename ChildT, u
int32_t Log2Dim = ChildT::LOG2DIM + 1>
3530 static constexpr
bool FIXED_SIZE = DataType::FIXED_SIZE;
3531 template<u
int32_t LOG2>
3536 static constexpr uint32_t LOG2DIM = Log2Dim;
3537 static constexpr uint32_t TOTAL = LOG2DIM + ChildT::TOTAL;
3538 static constexpr uint32_t DIM = 1u << TOTAL;
3539 static constexpr uint32_t SIZE = 1u << (3 * LOG2DIM);
3540 static constexpr uint32_t MASK = (1u << TOTAL) - 1u;
3541 static constexpr uint32_t LEVEL = 1 + ChildT::LEVEL;
3542 static constexpr uint64_t NUM_VALUES = uint64_t(1) << (3 * TOTAL);
3637 const uint32_t n = CoordToOffset(ijk);
3638 return DataType::mChildMask.isOn(n) ? this->getChild(n)->getValue(ijk) : DataType::getValue(n);
3643 const uint32_t n = CoordToOffset(ijk);
3644 return DataType::mChildMask.isOn(n) ? this->getChild(n)->isActive(ijk) : DataType::isActive(n);
3650 const uint32_t n = CoordToOffset(ijk);
3651 if (DataType::mChildMask.isOn(n))
3652 return this->getChild(n)->probeValue(ijk, v);
3653 v = DataType::getValue(n);
3654 return DataType::isActive(n);
3659 const uint32_t n = CoordToOffset(ijk);
3660 if (DataType::mChildMask.isOn(n))
3661 return this->getChild(n)->probeLeaf(ijk);
3667 const uint32_t n = CoordToOffset(ijk);
3668 return DataType::mChildMask.isOn(n) ? this->getChild(n) :
nullptr;
3675 return (((ijk[0] & MASK) >> ChildT::TOTAL) << (2 * LOG2DIM)) +
3676 (((ijk[1] & MASK) >> ChildT::TOTAL) << (LOG2DIM)) +
3677 ((ijk[2] & MASK) >> ChildT::TOTAL);
3679 return (((ijk[0] & MASK) >> ChildT::TOTAL) << (2 * LOG2DIM)) |
3680 (((ijk[1] & MASK) >> ChildT::TOTAL) << (LOG2DIM)) |
3681 ((ijk[2] & MASK) >> ChildT::TOTAL);
3689 const uint32_t m = n & ((1 << 2 * LOG2DIM) - 1);
3690 return Coord(n >> 2 * LOG2DIM, m >> LOG2DIM, m & ((1 << LOG2DIM) - 1));
3696 ijk <<= ChildT::TOTAL;
3697 ijk += this->origin();
3703 this->localToGlobalCoord(ijk);
3710 return DataType::mFlags & uint32_t(2);
3717 template<
typename,
int,
int,
int>
3722 template<
typename, u
int32_t>
3726 template<
typename AccT>
3729 const uint32_t n = CoordToOffset(ijk);
3730 if (!DataType::mChildMask.isOn(n))
3731 return DataType::getValue(n);
3732 const ChildT*
child = this->getChild(n);
3733 acc.insert(ijk, child);
3734 return child->getValueAndCache(ijk, acc);
3737 template<
typename AccT>
3738 __hostdev__ typename AccT::NodeInfo getNodeInfoAndCache(
const CoordType& ijk,
const AccT& acc)
const 3740 using NodeInfoT =
typename AccT::NodeInfo;
3741 const uint32_t n = CoordToOffset(ijk);
3742 if (!DataType::mChildMask.isOn(n)) {
3743 return NodeInfoT{LEVEL, this->dim(), this->minimum(), this->maximum(), this->average(),
3744 this->stdDeviation(), this->bbox()[0], this->bbox()[1]};
3746 const ChildT* child = this->getChild(n);
3747 acc.insert(ijk, child);
3748 return child->getNodeInfoAndCache(ijk, acc);
3751 template<
typename AccT>
3754 const uint32_t n = CoordToOffset(ijk);
3755 if (!DataType::mChildMask.isOn(n))
3756 return DataType::isActive(n);
3757 const ChildT* child = this->getChild(n);
3758 acc.insert(ijk, child);
3759 return child->isActiveAndCache(ijk, acc);
3762 template<
typename AccT>
3765 const uint32_t n = CoordToOffset(ijk);
3766 if (!DataType::mChildMask.isOn(n)) {
3767 v = DataType::getValue(n);
3768 return DataType::isActive(n);
3770 const ChildT* child = this->getChild(n);
3771 acc.insert(ijk, child);
3772 return child->probeValueAndCache(ijk, v, acc);
3775 template<
typename AccT>
3778 const uint32_t n = CoordToOffset(ijk);
3779 if (!DataType::mChildMask.isOn(n))
3781 const ChildT* child = this->getChild(n);
3782 acc.insert(ijk, child);
3783 return child->probeLeafAndCache(ijk, acc);
3786 template<
typename RayT,
typename AccT>
3787 __hostdev__ uint32_t getDimAndCache(
const CoordType& ijk,
const RayT& ray,
const AccT& acc)
const 3789 if (DataType::mFlags & uint32_t(1u))
return this->dim();
3792 const uint32_t n = CoordToOffset(ijk);
3793 if (DataType::mChildMask.isOn(n)) {
3794 const ChildT* child = this->getChild(n);
3795 acc.insert(ijk, child);
3796 return child->getDimAndCache(ijk, ray, acc);
3798 return ChildNodeType::dim();
3808 template<
typename ValueT,
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
3811 static_assert(
sizeof(CoordT) ==
sizeof(
Coord),
"Mismatching sizeof");
3812 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
3817 static constexpr
bool FIXED_SIZE =
true;
3820 uint8_t mBBoxDif[3];
3834 return sizeof(
LeafData) - (12 + 3 + 1 +
sizeof(MaskT<LOG2DIM>)
3835 + 2*(
sizeof(ValueT) +
sizeof(
FloatType))
3836 + (1u << (3 * LOG2DIM))*
sizeof(ValueT));
3845 mValueMask.setOn(offset);
3846 mValues[offset] =
value;
3859 template <
typename T>
3870 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
3873 static_assert(
sizeof(CoordT) ==
sizeof(
Coord),
"Mismatching sizeof");
3874 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
3879 uint8_t mBBoxDif[3];
3893 return sizeof(
LeafFnBase) - (12 + 3 + 1 +
sizeof(MaskT<LOG2DIM>) + 2*4 + 4*2);
3898 mQuantum = (max -
min)/
float((1 << bitWidth)-1);
3926 template <
typename T>
3933 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
3934 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<
Fp4, CoordT, MaskT, LOG2DIM>
3940 static constexpr
bool FIXED_SIZE =
true;
3941 alignas(32) uint8_t mCode[1u << (3 * LOG2DIM - 1)];
3945 static_assert(BaseT::padding()==0,
"expected no padding in LeafFnBase");
3946 return sizeof(
LeafData) -
sizeof(
BaseT) - (1u << (3 * LOG2DIM - 1));
3953 const uint8_t c = mCode[i>>1];
3954 return ( (i&1) ? c >> 4 : c & uint8_t(15) )*BaseT::mQuantum + BaseT::mMinimum;
3956 return ((mCode[i>>1] >> ((i&1)<<2)) & uint8_t(15))*BaseT::mQuantum + BaseT::mMinimum;
3967 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
3968 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<
Fp8, CoordT, MaskT, LOG2DIM>
3974 static constexpr
bool FIXED_SIZE =
true;
3975 alignas(32) uint8_t mCode[1u << 3 * LOG2DIM];
3978 static_assert(BaseT::padding()==0,
"expected no padding in LeafFnBase");
3979 return sizeof(
LeafData) -
sizeof(
BaseT) - (1u << 3 * LOG2DIM);
3985 return mCode[i]*BaseT::mQuantum + BaseT::mMinimum;
3994 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
3995 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<
Fp16, CoordT, MaskT, LOG2DIM>
4001 static constexpr
bool FIXED_SIZE =
true;
4002 alignas(32) uint16_t mCode[1u << 3 * LOG2DIM];
4006 static_assert(BaseT::padding()==0,
"expected no padding in LeafFnBase");
4007 return sizeof(
LeafData) -
sizeof(
BaseT) - 2*(1u << 3 * LOG2DIM);
4013 return mCode[i]*BaseT::mQuantum + BaseT::mMinimum;
4023 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4024 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<
FpN, CoordT, MaskT, LOG2DIM>
4030 static constexpr
bool FIXED_SIZE =
false;
4032 static_assert(BaseT::padding()==0,
"expected no padding in LeafFnBase");
4041 #ifdef NANOVDB_FPN_BRANCHLESS// faster 4042 const int b = BaseT::mFlags >> 5;
4044 uint16_t code =
reinterpret_cast<const uint16_t*
>(
this + 1)[i >> (4 - b)];
4045 const static uint8_t shift[5] = {15, 7, 3, 1, 0};
4046 const static uint16_t mask[5] = {1, 3, 15, 255, 65535};
4047 code >>= (i & shift[b]) << b;
4050 uint32_t code =
reinterpret_cast<const uint32_t*
>(
this + 1)[i >> (5 - b)];
4052 code >>= (i & ((32 >> b) - 1)) << b;
4053 code &= (1 << (1 << b)) - 1;
4055 #else// use branched version (slow) 4057 auto *values =
reinterpret_cast<const uint8_t*
>(
this+1);
4058 switch (BaseT::mFlags >> 5) {
4060 code = float((values[i>>3] >> (i&7) ) & uint8_t(1));
4063 code = float((values[i>>2] >> ((i&3)<<1)) & uint8_t(3));
4066 code = float((values[i>>1] >> ((i&1)<<2)) & uint8_t(15));
4069 code = float(values[i]);
4072 code = float(reinterpret_cast<const uint16_t*>(values)[i]);
4075 return float(code) * BaseT::mQuantum + BaseT::mMinimum;
4086 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4087 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<bool, CoordT, MaskT, LOG2DIM>
4089 static_assert(
sizeof(CoordT) ==
sizeof(
Coord),
"Mismatching sizeof");
4090 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
4095 static constexpr
bool FIXED_SIZE =
true;
4098 uint8_t mBBoxDif[3];
4102 uint64_t mPadding[2];
4115 mValueMask.setOn(offset);
4116 mValues.set(offset, v);
4124 template <
typename T>
4135 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4138 static_assert(
sizeof(CoordT) ==
sizeof(
Coord),
"Mismatching sizeof");
4139 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
4144 static constexpr
bool FIXED_SIZE =
true;
4147 uint8_t mBBoxDif[3];
4150 uint64_t mPadding[2];
4155 return sizeof(
LeafData) - (12u + 3u + 1u +
sizeof(MaskT<LOG2DIM>) + 2*8u);
4166 mValueMask.setOn(offset);
4174 template <
typename T>
4185 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4188 static_assert(
sizeof(CoordT) ==
sizeof(
Coord),
"Mismatching sizeof");
4189 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
4194 static constexpr
bool FIXED_SIZE =
true;
4197 uint8_t mBBoxDif[3];
4206 return sizeof(
LeafData) - (12u + 3u + 1u +
sizeof(MaskT<LOG2DIM>) + 2*8u);
4217 mValueMask.setOn(offset);
4222 if (mFlags & uint8_t(16u)) {
4223 return mValueMask.isOn(i) ? mValueOff + mValueMask.countOn(i) : 0;
4225 return mValueOff + i;
4228 template <
typename T>
4230 template <
typename T>
4232 template <
typename T>
4234 template <
typename T>
4236 template <
typename T>
4247 template<
typename BuildT,
4248 typename CoordT =
Coord,
4249 template<u
int32_t>
class MaskT =
Mask,
4250 uint32_t Log2Dim = 3>
4256 static constexpr uint32_t TOTAL = 0;
4257 static constexpr uint32_t DIM = 1;
4266 static constexpr
bool FIXED_SIZE = DataType::FIXED_SIZE;
4267 template<u
int32_t LOG2>
4314 __hostdev__ operator bool()
const {
return mPos < (1u << 3 * Log2Dim);}
4326 static constexpr uint32_t LOG2DIM = Log2Dim;
4327 static constexpr uint32_t TOTAL = LOG2DIM;
4328 static constexpr uint32_t DIM = 1u << TOTAL;
4329 static constexpr uint32_t SIZE = 1u << 3 * LOG2DIM;
4330 static constexpr uint32_t MASK = (1u << LOG2DIM) - 1u;
4331 static constexpr uint32_t LEVEL = 0;
4332 static constexpr uint64_t NUM_VALUES = uint64_t(1) << (3 * TOTAL);