118 #ifndef NANOVDB_NANOVDB_H_HAS_BEEN_INCLUDED 119 #define NANOVDB_NANOVDB_H_HAS_BEEN_INCLUDED 123 #include <nanovdb/math/Math.h> 126 #define NANOVDB_DATA_ALIGNMENT 32 134 #define NANOVDB_MAGIC_NUMB 0x304244566f6e614eUL // "NanoVDB0" in hex - little endian (uint64_t) 135 #define NANOVDB_MAGIC_GRID 0x314244566f6e614eUL // "NanoVDB1" in hex - little endian (uint64_t) 136 #define NANOVDB_MAGIC_FILE 0x324244566f6e614eUL // "NanoVDB2" in hex - little endian (uint64_t) 137 #define NANOVDB_MAGIC_NODE 0x334244566f6e614eUL // "NanoVDB3" in hex - little endian (uint64_t) 138 #define NANOVDB_MAGIC_FRAG 0x344244566f6e614eUL // "NanoVDB4" in hex - little endian (uint64_t) 139 #define NANOVDB_MAGIC_MASK 0x00FFFFFFFFFFFFFFUL // use this mask to remove the number 144 #define NANOVDB_MAJOR_VERSION_NUMBER 32 // reflects changes to the ABI and hence also the file format 145 #define NANOVDB_MINOR_VERSION_NUMBER 7 // reflects changes to the API but not ABI 146 #define NANOVDB_PATCH_VERSION_NUMBER 0 // reflects changes that does not affect the ABI or API 148 #define TBB_SUPPRESS_DEPRECATED_MESSAGES 1 151 #define NANOVDB_USE_SINGLE_ROOT_KEY 160 #define NANOVDB_NEW_ACCESSOR_METHODS 162 #define NANOVDB_FPN_BRANCHLESS 164 #if !defined(NANOVDB_ALIGN) 165 #define NANOVDB_ALIGN(n) alignas(n) 166 #endif // !defined(NANOVDB_ALIGN) 208 template <
class EnumT>
557 template <
typename T>
642 switch (blindClass) {
694 : mData(major << 21 | minor << 10 | patch)
733 static const int Rank = 0;
734 static const bool IsScalar =
true;
735 static const bool IsVector =
false;
736 static const int Size = 1;
738 static T
scalar(
const T& s) {
return s; }
744 static const int Rank = 1;
745 static const bool IsScalar =
false;
746 static const bool IsVector =
true;
747 static const int Size = T::SIZE;
754 template<typename T, int = sizeof(typename TensorTraits<T>::ElementType)>
811 template<
typename BuildT>
870 template<
typename BuildT>
871 [[deprecated(
"Use toGridType<T>() instead.")]]
877 template<
typename BuildT>
892 template<
typename BuildT>
893 [[deprecated(
"Use toGridClass<T>() instead.")]]
896 return toGridClass<BuildT>();
934 BitFlags(std::initializer_list<uint8_t> list)
936 for (
auto bit : list) mFlags |=
static_cast<Type>(1 << bit);
938 template<
typename MaskT>
939 BitFlags(std::initializer_list<MaskT> list)
941 for (
auto mask : list) mFlags |=
static_cast<Type>(mask);
945 __hostdev__ void initBit(std::initializer_list<uint8_t> list)
948 for (
auto bit : list) mFlags |=
static_cast<Type>(1 << bit);
950 template<
typename MaskT>
951 __hostdev__ void initMask(std::initializer_list<MaskT> list)
954 for (
auto mask : list) mFlags |=
static_cast<Type>(mask);
963 __hostdev__ void setBitOn(uint8_t bit) { mFlags |=
static_cast<Type>(1 << bit); }
964 __hostdev__ void setBitOff(uint8_t bit) { mFlags &= ~static_cast<
Type>(1 << bit); }
966 __hostdev__ void setBitOn(std::initializer_list<uint8_t> list)
968 for (
auto bit : list) mFlags |=
static_cast<Type>(1 << bit);
970 __hostdev__ void setBitOff(std::initializer_list<uint8_t> list)
972 for (
auto bit : list) mFlags &= ~static_cast<
Type>(1 << bit);
975 template<
typename MaskT>
976 __hostdev__ void setMaskOn(MaskT mask) { mFlags |=
static_cast<Type>(mask); }
977 template<
typename MaskT>
978 __hostdev__ void setMaskOff(MaskT mask) { mFlags &= ~static_cast<
Type>(mask); }
980 template<
typename MaskT>
981 __hostdev__ void setMaskOn(std::initializer_list<MaskT> list)
983 for (
auto mask : list) mFlags |=
static_cast<Type>(mask);
985 template<
typename MaskT>
986 __hostdev__ void setMaskOff(std::initializer_list<MaskT> list)
988 for (
auto mask : list) mFlags &= ~static_cast<
Type>(mask);
991 __hostdev__ void setBit(uint8_t bit,
bool on) { on ? this->setBitOn(bit) : this->setBitOff(bit); }
992 template<
typename MaskT>
993 __hostdev__ void setMask(MaskT mask,
bool on) { on ? this->setMaskOn(mask) : this->setMaskOff(mask); }
997 __hostdev__ bool isBitOn(uint8_t bit)
const {
return 0 != (mFlags &
static_cast<Type>(1 << bit)); }
998 __hostdev__ bool isBitOff(uint8_t bit)
const {
return 0 == (mFlags &
static_cast<Type>(1 << bit)); }
999 template<
typename MaskT>
1000 __hostdev__ bool isMaskOn(MaskT mask)
const {
return 0 != (mFlags &
static_cast<Type>(mask)); }
1001 template<
typename MaskT>
1002 __hostdev__ bool isMaskOff(MaskT mask)
const {
return 0 == (mFlags &
static_cast<Type>(mask)); }
1004 template<
typename MaskT>
1005 __hostdev__ bool isMaskOn(std::initializer_list<MaskT> list)
const 1007 for (
auto mask : list) {
1008 if (0 != (mFlags & static_cast<Type>(mask)))
return true;
1013 template<
typename MaskT>
1014 __hostdev__ bool isMaskOff(std::initializer_list<MaskT> list)
const 1016 for (
auto mask : list) {
1017 if (0 == (mFlags & static_cast<Type>(mask)))
return true;
1033 template<u
int32_t LOG2DIM>
1037 static constexpr uint32_t SIZE = 1U << (3 * LOG2DIM);
1038 static constexpr uint32_t WORD_COUNT = SIZE >> 6;
1053 for (
const uint64_t *w = mWords, *q = w + WORD_COUNT; w != q; ++w)
1061 uint32_t n = i >> 6, sum =
util::countOn(mWords[n] & ((uint64_t(1) << (i & 63u)) - 1u));
1062 for (
const uint64_t* w = mWords; n--; ++w)
1087 mPos = mParent->findNext<On>(mPos + 1);
1099 const Mask* mParent;
1141 for (uint32_t i = 0; i < WORD_COUNT; ++i)
1146 const uint64_t v = on ? ~uint64_t(0) : uint64_t(0);
1147 for (uint32_t i = 0; i < WORD_COUNT; ++i)
1154 for (uint32_t i = 0; i < WORD_COUNT; ++i)
1155 mWords[i] = other.mWords[i];
1162 template<
typename WordT>
1167 return reinterpret_cast<WordT*
>(mWords)[n];
1169 template<
typename WordT>
1174 reinterpret_cast<WordT*
>(mWords)[n] = w;
1178 template<
typename MaskT = Mask>
1181 static_assert(
sizeof(
Mask) ==
sizeof(MaskT),
"Mismatching sizeof");
1182 static_assert(WORD_COUNT == MaskT::WORD_COUNT,
"Mismatching word count");
1183 static_assert(LOG2DIM == MaskT::LOG2DIM,
"Mismatching LOG2DIM");
1184 auto* src =
reinterpret_cast<const uint64_t*
>(&other);
1185 for (uint64_t *dst = mWords, *end = dst + WORD_COUNT; dst != end; ++dst)
1191 Mask& operator=(
const Mask&) =
default;
1195 for (uint32_t i = 0; i < WORD_COUNT; ++i) {
1196 if (mWords[i] != other.mWords[i])
1205 __hostdev__ bool isOn(uint32_t n)
const {
return 0 != (mWords[n >> 6] & (uint64_t(1) << (n & 63))); }
1208 __hostdev__ bool isOff(uint32_t n)
const {
return 0 == (mWords[n >> 6] & (uint64_t(1) << (n & 63))); }
1213 for (uint32_t i = 0; i < WORD_COUNT; ++i)
1214 if (mWords[i] != ~uint64_t(0))
1222 for (uint32_t i = 0; i < WORD_COUNT; ++i)
1223 if (mWords[i] != uint64_t(0))
1233 #if defined(__CUDACC__) // the following functions only run on the GPU! 1234 __device__ inline void setOnAtomic(uint32_t n)
1236 atomicOr(reinterpret_cast<unsigned long long int*>(
this) + (n >> 6), 1ull << (n & 63));
1238 __device__ inline void setOffAtomic(uint32_t n)
1240 atomicAnd(reinterpret_cast<unsigned long long int*>(
this) + (n >> 6), ~(1ull << (n & 63)));
1242 __device__ inline void setAtomic(uint32_t n,
bool on)
1244 on ? this->setOnAtomic(n) : this->setOffAtomic(n);
1269 #if 1 // switch between branchless 1270 auto& word = mWords[n >> 6];
1272 word &= ~(uint64_t(1) << n);
1273 word |= uint64_t(on) << n;
1275 on ? this->setOn(n) : this->setOff(n);
1282 for (uint32_t i = 0; i < WORD_COUNT; ++i)mWords[i] = ~uint64_t(0);
1288 for (uint32_t i = 0; i < WORD_COUNT; ++i) mWords[i] = uint64_t(0);
1294 const uint64_t v = on ? ~uint64_t(0) : uint64_t(0);
1295 for (uint32_t i = 0; i < WORD_COUNT; ++i) mWords[i] = v;
1300 uint32_t n = WORD_COUNT;
1301 for (
auto* w = mWords; n--; ++w) *w = ~*w;
1308 uint64_t* w1 = mWords;
1309 const uint64_t* w2 = other.mWords;
1310 for (uint32_t n = WORD_COUNT; n--; ++w1, ++w2) *w1 &= *w2;
1316 uint64_t* w1 = mWords;
1317 const uint64_t* w2 = other.mWords;
1318 for (uint32_t n = WORD_COUNT; n--; ++w1, ++w2) *w1 |= *w2;
1324 uint64_t* w1 = mWords;
1325 const uint64_t* w2 = other.mWords;
1326 for (uint32_t n = WORD_COUNT; n--; ++w1, ++w2) *w1 &= ~*w2;
1332 uint64_t* w1 = mWords;
1333 const uint64_t* w2 = other.mWords;
1334 for (uint32_t n = WORD_COUNT; n--; ++w1, ++w2) *w1 ^= *w2;
1343 const uint64_t* w = mWords;
1344 for (; n < WORD_COUNT && !(ON ? *w : ~*w); ++w, ++n);
1352 uint32_t n = start >> 6;
1353 if (n >= WORD_COUNT)
return SIZE;
1354 uint32_t m = start & 63u;
1355 uint64_t b = ON ? mWords[n] : ~mWords[n];
1356 if (b & (uint64_t(1u) << m))
return start;
1357 b &= ~uint64_t(0u) << m;
1358 while (!b && ++n < WORD_COUNT) b = ON ? mWords[n] : ~mWords[n];
1366 uint32_t n = start >> 6;
1367 if (n >= WORD_COUNT)
return SIZE;
1368 uint32_t m = start & 63u;
1369 uint64_t b = ON ? mWords[n] : ~mWords[n];
1370 if (b & (uint64_t(1u) << m))
return start;
1371 b &= (uint64_t(1u) << m) - 1u;
1372 while (!b && n) b = ON ? mWords[--n] : ~mWords[--n];
1377 uint64_t mWords[WORD_COUNT];
1396 : mMatF{ 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f}
1397 , mInvMatF{1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f}
1398 , mVecF{0.0f, 0.0f, 0.0f}
1400 , mMatD{ 1.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0, 1.0}
1401 , mInvMatD{1.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0, 0.0, 1.0}
1402 , mVecD{0.0, 0.0, 0.0}
1407 : mMatF{float(s), 0.0f, 0.0f, 0.0f, float(s), 0.0f, 0.0f, 0.0f, float(s)}
1408 , mInvMatF{1.0f / float(s), 0.0f, 0.0f, 0.0f, 1.0f / float(s), 0.0f, 0.0f, 0.0f, 1.0f / float(s)}
1409 , mVecF{float(t[0]), float(t[1]), float(t[2])}
1411 , mMatD{s, 0.0, 0.0, 0.0, s, 0.0, 0.0, 0.0, s}
1412 , mInvMatD{1.0 / s, 0.0, 0.0, 0.0, 1.0 / s, 0.0, 0.0, 0.0, 1.0 / s}
1413 , mVecD{t[0], t[1], t[2]}
1420 template<
typename MatT,
typename Vec3T>
1421 void set(
const MatT& mat,
const MatT& invMat,
const Vec3T& translate,
double taper = 1.0);
1426 template<
typename Mat4T>
1427 void set(
const Mat4T& mat,
const Mat4T& invMat,
double taper = 1.0) { this->
set(mat, invMat, mat[3], taper); }
1429 template<
typename Vec3T>
1430 void set(
double scale,
const Vec3T& translation,
double taper = 1.0);
1437 template<
typename Vec3T>
1445 template<
typename Vec3T>
1454 template<
typename Vec3T>
1463 template<
typename Vec3T>
1471 template<
typename Vec3T>
1474 return math::matMult(mInvMatD, Vec3T(xyz[0] - mVecD[0], xyz[1] - mVecD[1], xyz[2] - mVecD[2]));
1482 template<
typename Vec3T>
1485 return math::matMult(mInvMatF, Vec3T(xyz[0] - mVecF[0], xyz[1] - mVecF[1], xyz[2] - mVecF[2]));
1494 template<
typename Vec3T>
1503 template<
typename Vec3T>
1512 template<
typename Vec3T>
1514 template<
typename Vec3T>
1521 template<
typename MatT,
typename Vec3T>
1522 inline void Map::set(
const MatT& mat,
const MatT& invMat,
const Vec3T& translate,
double taper)
1524 float * mf = mMatF, *vf = mVecF, *mif = mInvMatF;
1525 double *md = mMatD, *vd = mVecD, *mid = mInvMatD;
1526 mTaperF =
static_cast<float>(taper);
1528 for (
int i = 0; i < 3; ++i) {
1529 *vd++ = translate[i];
1530 *vf++ =
static_cast<float>(translate[i]);
1531 for (
int j = 0; j < 3; ++j) {
1533 *mid++ = invMat[j][i];
1534 *mf++ =
static_cast<float>(mat[j][i]);
1535 *mif++ =
static_cast<float>(invMat[j][i]);
1540 template<
typename Vec3T>
1541 inline void Map::set(
double dx,
const Vec3T& trans,
double taper)
1544 const double mat[3][3] = { {dx, 0.0, 0.0},
1547 const double idx = 1.0 / dx;
1548 const double invMat[3][3] = { {idx, 0.0, 0.0},
1551 this->
set(mat, invMat, trans, taper);
1558 static const int MaxNameSize = 256;
1565 char mName[MaxNameSize];
1583 template<
typename BlindDataT>
1587 return mDataType == toGridType<BlindDataT>() ? util::PtrAdd<BlindDataT>(
this, mDataOffset) :
nullptr;
1593 auto check = [&]()->
bool{
1611 default:
return true;}
1621 return math::AlignUp<NANOVDB_DATA_ALIGNMENT>(mValueCount * mValueSize);
1629 template<
typename Gr
idOrTreeOrRootT,
int LEVEL>
1633 template<
typename Gr
idOrTreeOrRootT>
1636 static_assert(GridOrTreeOrRootT::RootNodeType::LEVEL == 3,
"Tree depth is not supported");
1637 using Type =
typename GridOrTreeOrRootT::LeafNodeType;
1638 using type =
typename GridOrTreeOrRootT::LeafNodeType;
1640 template<
typename Gr
idOrTreeOrRootT>
1643 static_assert(GridOrTreeOrRootT::RootNodeType::LEVEL == 3,
"Tree depth is not supported");
1644 using Type =
const typename GridOrTreeOrRootT::LeafNodeType;
1645 using type =
const typename GridOrTreeOrRootT::LeafNodeType;
1648 template<
typename Gr
idOrTreeOrRootT>
1651 static_assert(GridOrTreeOrRootT::RootNodeType::LEVEL == 3,
"Tree depth is not supported");
1652 using Type =
typename GridOrTreeOrRootT::RootNodeType::ChildNodeType::ChildNodeType;
1653 using type =
typename GridOrTreeOrRootT::RootNodeType::ChildNodeType::ChildNodeType;
1655 template<
typename Gr
idOrTreeOrRootT>
1658 static_assert(GridOrTreeOrRootT::RootNodeType::LEVEL == 3,
"Tree depth is not supported");
1659 using Type =
const typename GridOrTreeOrRootT::RootNodeType::ChildNodeType::ChildNodeType;
1660 using type =
const typename GridOrTreeOrRootT::RootNodeType::ChildNodeType::ChildNodeType;
1662 template<
typename Gr
idOrTreeOrRootT>
1665 static_assert(GridOrTreeOrRootT::RootNodeType::LEVEL == 3,
"Tree depth is not supported");
1666 using Type =
typename GridOrTreeOrRootT::RootNodeType::ChildNodeType;
1667 using type =
typename GridOrTreeOrRootT::RootNodeType::ChildNodeType;
1669 template<
typename Gr
idOrTreeOrRootT>
1672 static_assert(GridOrTreeOrRootT::RootNodeType::LEVEL == 3,
"Tree depth is not supported");
1673 using Type =
const typename GridOrTreeOrRootT::RootNodeType::ChildNodeType;
1674 using type =
const typename GridOrTreeOrRootT::RootNodeType::ChildNodeType;
1676 template<
typename Gr
idOrTreeOrRootT>
1679 static_assert(GridOrTreeOrRootT::RootNodeType::LEVEL == 3,
"Tree depth is not supported");
1680 using Type =
typename GridOrTreeOrRootT::RootNodeType;
1681 using type =
typename GridOrTreeOrRootT::RootNodeType;
1684 template<
typename Gr
idOrTreeOrRootT>
1687 static_assert(GridOrTreeOrRootT::RootNodeType::LEVEL == 3,
"Tree depth is not supported");
1688 using Type =
const typename GridOrTreeOrRootT::RootNodeType;
1689 using type =
const typename GridOrTreeOrRootT::RootNodeType;
1694 template<
typename BuildT>
1696 template<
typename BuildT>
1698 template<
typename BuildT>
1700 template<
typename BuildT>
1702 template<
typename BuildT>
1704 template<
typename BuildT>
1706 template<
typename BuildT>
1708 template<
typename BuildT>
1746 union { uint32_t mCRC32[2]; uint64_t
mCRC64; };
1750 static constexpr uint32_t EMPTY32 = ~uint32_t{0};
1751 static constexpr uint64_t EMPTY64 = ~uint64_t(0);
1770 [[deprecated(
"Use Checksum::data instead.")]]
1772 [[deprecated(
"Use Checksum::head and Ckecksum::tail instead.")]]
1774 [[deprecated(
"Use Checksum::head and Ckecksum::tail instead.")]]
1785 [[deprecated(
"Use Checksum::isHalf instead.")]]
1846 static const int MaxNameSize = 256;
1854 char mGridName[MaxNameSize];
1869 uint64_t gridSize = 0u,
1874 #ifdef NANOVDB_USE_NEW_MAGIC_NUMBERS 1881 mFlags.initMask(list);
1884 mGridSize = gridSize;
1885 mGridName[0] =
'\0';
1887 mWorldBBox = Vec3dBBox();
1888 mVoxelSize = map.getVoxelSize();
1889 mGridClass = gridClass;
1890 mGridType = gridType;
1891 mBlindMetadataOffset = mGridSize;
1892 mBlindMetadataCount = 0u;
1906 if (test) test = mGridCount > 0u && mGridIndex < mGridCount;
1918 const bool success = (
util::strncpy(mGridName, src, MaxNameSize)[MaxNameSize-1] ==
'\0');
1919 if (!success) mGridName[MaxNameSize-1] =
'\0';
1923 template<
typename Vec3T>
1925 template<
typename Vec3T>
1927 template<
typename Vec3T>
1929 template<
typename Vec3T>
1931 template<
typename Vec3T>
1934 template<
typename Vec3T>
1936 template<
typename Vec3T>
1938 template<
typename Vec3T>
1940 template<
typename Vec3T>
1942 template<
typename Vec3T>
1953 template <u
int32_t LEVEL>
1956 static_assert(LEVEL >= 0 && LEVEL <= 3,
"invalid LEVEL template parameter");
1957 const void *treeData =
this + 1;
1958 const uint64_t nodeOffset = *util::PtrAdd<uint64_t>(treeData, 8*LEVEL);
1959 return nodeOffset ?
util::PtrAdd(treeData, nodeOffset) :
nullptr;
1965 template <u
int32_t LEVEL>
1968 static_assert(LEVEL >= 0 && LEVEL <= 3,
"invalid LEVEL template parameter");
1969 void *treeData =
this + 1;
1970 const uint64_t nodeOffset = *util::PtrAdd<uint64_t>(treeData, 8*LEVEL);
1971 return nodeOffset ?
util::PtrAdd(treeData, nodeOffset) :
nullptr;
1976 template <u
int32_t LEVEL>
1979 static_assert(LEVEL >= 0 && LEVEL < 3,
"invalid LEVEL template parameter");
1980 return *util::PtrAdd<uint32_t>(
this + 1, 4*(8 + LEVEL));
1989 return util::PtrAdd<GridBlindMetaData>(
this, mBlindMetadataOffset) + n;
1996 for (uint32_t i = 0; i < mBlindMetadataCount; ++i) {
1997 const auto* metaData = this->blindMetaData(i);
2000 return metaData->template getBlindData<const char>();
2020 const void *root = this->nodePtr<3>();
2021 return root ? *util::PtrAdd<uint32_t>(root,
sizeof(CoordBBox)) : 0u;
2034 template<
typename BuildT,
int LEVEL0 = -1,
int LEVEL1 = -1,
int LEVEL2 = -1>
2037 template<
typename BuildT>
2044 template<
typename TreeT>
2064 Grid& operator=(
const Grid&) =
delete;
2088 template<
typename T = BuildType>
2095 template<
typename T = BuildType>
2100 __hostdev__ const TreeT&
tree()
const {
return *
reinterpret_cast<const TreeT*
>(this->treePtr()); }
2115 template<
typename Vec3T>
2119 template<
typename Vec3T>
2124 template<
typename Vec3T>
2129 template<
typename Vec3T>
2134 template<
typename Vec3T>
2138 template<
typename Vec3T>
2142 template<
typename Vec3T>
2147 template<
typename Vec3T>
2152 template<
typename Vec3T>
2157 template<
typename Vec3T>
2193 template<
typename NodeT>
2202 __hostdev__ bool isSequential()
const {
return UpperNodeType::FIXED_SIZE && LowerNodeType::FIXED_SIZE && LeafNodeType::FIXED_SIZE && this->isBreadthFirst(); }
2220 __hostdev__ int findBlindData(
const char* name)
const;
2229 [[deprecated(
"Use Grid::getBlindData<T>() instead.")]]
2232 printf(
"\nnanovdb::Grid::blindData is unsafe and hence deprecated! Please use nanovdb::Grid::getBlindData instead.\n\n");
2234 return this->blindMetaData(n).blindData();
2237 template <
typename BlindDataT>
2240 if (n >= DataType::mBlindMetadataCount)
return nullptr;
2241 return this->blindMetaData(n).template getBlindData<BlindDataT>();
2244 template <
typename BlindDataT>
2247 if (n >= DataType::mBlindMetadataCount)
return nullptr;
2248 return const_cast<BlindDataT*
>(this->blindMetaData(n).template getBlindData<BlindDataT>());
2257 template<
typename TreeT>
2260 for (uint32_t i = 0, n = this->blindDataCount(); i < n; ++i) {
2261 if (this->blindMetaData(i).mSemantic == semantic)
2267 template<
typename TreeT>
2270 auto test = [&](
int n) {
2271 const char* str = this->blindMetaData(n).mName;
2273 if (name[i] != str[i])
2275 if (name[i] ==
'\0' && str[i] ==
'\0')
2280 for (
int i = 0, n = this->blindDataCount(); i < n; ++i)
2290 int64_t mNodeOffset[4];
2291 uint32_t mNodeCount[3];
2292 uint32_t mTileCount[3];
2308 template<
typename NodeT>
2312 __hostdev__ bool isEmpty()
const {
return mNodeOffset[3] ? *util::PtrAdd<uint32_t>(
this, mNodeOffset[3] +
sizeof(CoordBBox)) == 0 :
true;}
2315 __hostdev__ CoordBBox
bbox()
const {
return mNodeOffset[3] ? *util::PtrAdd<CoordBBox>(
this, mNodeOffset[3]) : CoordBBox();}
2324 template<
typename Gr
idT>
2327 using Type =
typename GridT::TreeType;
2328 using type =
typename GridT::TreeType;
2330 template<
typename Gr
idT>
2333 using Type =
const typename GridT::TreeType;
2334 using type =
const typename GridT::TreeType;
2340 template<
typename RootT>
2343 static_assert(RootT::LEVEL == 3,
"Tree depth is not supported");
2344 static_assert(RootT::ChildNodeType::LOG2DIM == 5,
"Tree configuration is not supported");
2345 static_assert(RootT::ChildNodeType::ChildNodeType::LOG2DIM == 4,
"Tree configuration is not supported");
2346 static_assert(RootT::LeafNodeType::LOG2DIM == 3,
"Tree configuration is not supported");
2361 using Node2 =
typename RootT::ChildNodeType;
2362 using Node1 =
typename Node2::ChildNodeType;
2368 Tree& operator=(
const Tree&) =
delete;
2380 __hostdev__ const RootT&
root()
const {
return *
reinterpret_cast<const RootT*
>(DataType::getRoot());}
2417 return DataType::mTileCount[level - 1];
2420 template<
typename NodeT>
2423 static_assert(NodeT::LEVEL < 3,
"Invalid NodeT");
2424 return DataType::mNodeCount[NodeT::LEVEL];
2430 return DataType::mNodeCount[level];
2435 return DataType::mNodeCount[0] + DataType::mNodeCount[1] + DataType::mNodeCount[2];
2441 template<
typename NodeT>
2444 const int64_t nodeOffset = DataType::mNodeOffset[NodeT::LEVEL];
2445 return nodeOffset ? util::PtrAdd<NodeT>(
this, nodeOffset) :
nullptr;
2451 template<
typename NodeT>
2454 const int64_t nodeOffset = DataType::mNodeOffset[NodeT::LEVEL];
2455 return nodeOffset ? util::PtrAdd<NodeT>(
this, nodeOffset) :
nullptr;
2464 return this->
template getFirstNode<typename NodeTrait<RootT, LEVEL>::type>();
2473 return this->
template getFirstNode<typename NodeTrait<RootT, LEVEL>::type>();
2484 template<
typename OpT,
typename... ArgsT>
2487 return this->root().template get<OpT>(ijk, args...);
2490 template<
typename OpT,
typename... ArgsT>
2493 return this->root().template set<OpT>(ijk, args...);
2501 template<
typename RootT>
2504 min = this->root().minimum();
2505 max = this->root().maximum();
2513 template<
typename ChildT>
2520 static constexpr
bool FIXED_SIZE =
false;
2523 #ifdef NANOVDB_USE_SINGLE_ROOT_KEY 2525 template<
typename CoordType>
2528 static_assert(
sizeof(
CoordT) ==
sizeof(CoordType),
"Mismatching sizeof");
2529 static_assert(32 - ChildT::TOTAL <= 21,
"Cannot use 64 bit root keys");
2530 return (
KeyT(uint32_t(ijk[2]) >> ChildT::TOTAL)) |
2531 (
KeyT(uint32_t(ijk[1]) >> ChildT::TOTAL) << 21) |
2532 (
KeyT(uint32_t(ijk[0]) >> ChildT::TOTAL) << 42);
2536 static constexpr uint64_t MASK = (1u << 21) - 1;
2537 return CoordT(((key >> 42) & MASK) << ChildT::TOTAL,
2538 ((key >> 21) & MASK) << ChildT::TOTAL,
2539 (key & MASK) << ChildT::TOTAL);
2563 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
Tile 2565 template<
typename CoordType>
2568 key = CoordToKey(k);
2572 template<
typename CoordType,
typename ValueType>
2575 key = CoordToKey(k);
2596 return reinterpret_cast<const Tile*
>(
this + 1) + n;
2601 return reinterpret_cast<Tile*
>(
this + 1) + n;
2606 #if 1 // switch between linear and binary seach 2607 const auto key = CoordToKey(ijk);
2608 for (
Tile *p = reinterpret_cast<Tile*>(
this + 1), *q = p + mTableSize; p < q; ++p)
2612 #else // do not enable binary search if tiles are not guaranteed to be sorted!!!!!! 2613 int32_t low = 0, high = mTableSize;
2614 while (low != high) {
2615 int mid = low + ((high - low) >> 1);
2616 const Tile* tile = &tiles[mid];
2617 if (tile->
key == key) {
2619 }
else if (tile->
key < key) {
2631 return const_cast<RootData*
>(
this)->probeTile(ijk);
2640 return util::PtrAdd<ChildT>(
this, tile->
child);
2645 return util::PtrAdd<ChildT>(
this, tile->
child);
2668 template<
typename ChildT>
2687 static constexpr
bool FIXED_SIZE = DataType::FIXED_SIZE;
2689 static constexpr uint32_t LEVEL = 1 + ChildT::LEVEL;
2691 template<
typename RootT>
2714 return this->tile()->origin();
2719 return this->tile()->origin();
2723 template<
typename RootT>
2736 :
BaseT(parent->data(), parent->tileCount())
2739 while (*
this && !this->tile()->isChild())
2745 return *BaseT::mData->getChild(this->tile());
2750 return BaseT::mData->getChild(this->tile());
2756 while (*
this && this->tile()->isValue())
2774 template<
typename RootT>
2785 :
BaseT(parent->data(), parent->tileCount())
2788 while (*
this && this->tile()->isChild())
2794 return this->tile()->value;
2799 return this->tile()->state;
2805 while (*
this && this->tile()->isChild())
2823 template<
typename RootT>
2834 :
BaseT(parent->data(), parent->tileCount())
2837 while (*
this && !this->tile()->isActive())
2843 return this->tile()->value;
2849 while (*
this && !this->tile()->isActive())
2867 template<
typename RootT>
2879 :
BaseT(parent->data(), parent->tileCount())
2886 NodeT* child =
nullptr;
2887 auto* t = this->tile();
2889 child = BaseT::mData->getChild(t);
2898 return this->tile()->state;
2970 #ifdef NANOVDB_NEW_ACCESSOR_METHODS 2978 #else // NANOVDB_NEW_ACCESSOR_METHODS 2983 if (
const Tile* tile = DataType::probeTile(ijk)) {
2984 return tile->isChild() ? this->getChild(tile)->getValue(ijk) : tile->value;
2986 return DataType::mBackground;
2992 if (
const Tile* tile = DataType::probeTile(ijk)) {
2993 return tile->isChild() ? this->getChild(tile)->isActive(ijk) : tile->state;
3000 if (
const Tile* tile = DataType::probeTile(ijk)) {
3001 if (tile->isChild()) {
3002 const auto* child = this->getChild(tile);
3003 return child->probeValue(ijk, v);
3008 v = DataType::mBackground;
3014 const Tile* tile = DataType::probeTile(ijk);
3015 if (tile && tile->isChild()) {
3016 const auto* child = this->getChild(tile);
3017 return child->probeLeaf(ijk);
3022 #endif // NANOVDB_NEW_ACCESSOR_METHODS 3026 const Tile* tile = DataType::probeTile(ijk);
3027 return tile && tile->isChild() ? this->getChild(tile) :
nullptr;
3032 const Tile* tile = DataType::probeTile(ijk);
3033 return tile && tile->isChild() ? this->getChild(tile) :
nullptr;
3036 template<
typename OpT,
typename... ArgsT>
3039 if (
const Tile* tile = this->probeTile(ijk)) {
3040 if (tile->isChild())
3041 return this->getChild(tile)->template get<OpT>(ijk, args...);
3042 return OpT::get(*tile, args...);
3044 return OpT::get(*
this, args...);
3047 template<
typename OpT,
typename... ArgsT>
3049 __hostdev__ decltype(OpT::set(util::declval<Tile&>(), util::declval<ArgsT>()...))
3052 if (
Tile* tile = DataType::probeTile(ijk)) {
3053 if (tile->isChild())
3054 return this->getChild(tile)->template set<OpT>(ijk, args...);
3055 return OpT::set(*tile, args...);
3057 return OpT::set(*
this, args...);
3064 template<
typename,
int,
int,
int>
3069 #ifndef NANOVDB_NEW_ACCESSOR_METHODS 3071 template<
typename AccT>
3072 __hostdev__ typename AccT::NodeInfo getNodeInfoAndCache(
const CoordType& ijk,
const AccT& acc)
const 3074 using NodeInfoT =
typename AccT::NodeInfo;
3075 if (
const Tile* tile = this->probeTile(ijk)) {
3076 if (tile->isChild()) {
3077 const auto* child = this->getChild(tile);
3078 acc.insert(ijk, child);
3079 return child->getNodeInfoAndCache(ijk, acc);
3081 return NodeInfoT{LEVEL, ChildT::dim(), tile->value, tile->value, tile->value, 0, tile->origin(), tile->origin() +
CoordType(ChildT::DIM)};
3083 return NodeInfoT{LEVEL, ChildT::dim(), this->minimum(), this->maximum(), this->average(), this->stdDeviation(), this->bbox()[0], this->bbox()[1]};
3087 template<
typename AccT>
3090 if (
const Tile* tile = this->probeTile(ijk)) {
3091 if (tile->isChild()) {
3092 const auto* child = this->getChild(tile);
3093 acc.insert(ijk, child);
3094 return child->getValueAndCache(ijk, acc);
3098 return DataType::mBackground;
3101 template<
typename AccT>
3104 const Tile* tile = this->probeTile(ijk);
3105 if (tile && tile->isChild()) {
3106 const auto* child = this->getChild(tile);
3107 acc.insert(ijk, child);
3108 return child->isActiveAndCache(ijk, acc);
3113 template<
typename AccT>
3116 if (
const Tile* tile = this->probeTile(ijk)) {
3117 if (tile->isChild()) {
3118 const auto* child = this->getChild(tile);
3119 acc.insert(ijk, child);
3120 return child->probeValueAndCache(ijk, v, acc);
3125 v = DataType::mBackground;
3129 template<
typename AccT>
3132 const Tile* tile = this->probeTile(ijk);
3133 if (tile && tile->isChild()) {
3134 const auto* child = this->getChild(tile);
3135 acc.insert(ijk, child);
3136 return child->probeLeafAndCache(ijk, acc);
3140 #endif // NANOVDB_NEW_ACCESSOR_METHODS 3142 template<
typename RayT,
typename AccT>
3143 __hostdev__ uint32_t getDimAndCache(
const CoordType& ijk,
const RayT& ray,
const AccT& acc)
const 3145 if (
const Tile* tile = this->probeTile(ijk)) {
3146 if (tile->isChild()) {
3147 const auto* child = this->getChild(tile);
3148 acc.insert(ijk, child);
3149 return child->getDimAndCache(ijk, ray, acc);
3151 return 1 << ChildT::TOTAL;
3153 return ChildNodeType::dim();
3156 template<
typename OpT,
typename AccT,
typename... ArgsT>
3159 getAndCache(
const CoordType& ijk,
const AccT& acc, ArgsT&&... args)
const 3161 if (
const Tile* tile = this->probeTile(ijk)) {
3162 if (tile->isChild()) {
3163 const ChildT* child = this->getChild(tile);
3164 acc.insert(ijk, child);
3165 return child->template getAndCache<OpT>(ijk, acc, args...);
3167 return OpT::get(*tile, args...);
3169 return OpT::get(*
this, args...);
3172 template<
typename OpT,
typename AccT,
typename... ArgsT>
3174 __hostdev__ decltype(OpT::set(util::declval<Tile&>(), util::declval<ArgsT>()...))
3175 setAndCache(const
CoordType& ijk, const AccT& acc, ArgsT&&... args)
3177 if (
Tile* tile = DataType::probeTile(ijk)) {
3178 if (tile->isChild()) {
3179 ChildT* child = this->getChild(tile);
3180 acc.insert(ijk, child);
3181 return child->template setAndCache<OpT>(ijk, acc, args...);
3183 return OpT::set(*tile, args...);
3185 return OpT::set(*
this, args...);
3197 template<
typename ChildT, u
int32_t LOG2DIM>
3204 using MaskT =
typename ChildT::template MaskType<LOG2DIM>;
3205 static constexpr
bool FIXED_SIZE =
true;
3214 Tile& operator=(
const Tile&) =
delete;
3236 alignas(32) Tile mTable[1u << (3 * LOG2DIM)];
3246 template<
typename ValueT>
3250 mTable[n].value = v;
3257 return util::PtrAdd<ChildT>(
this, mTable[n].child);
3262 return util::PtrAdd<ChildT>(
this, mTable[n].child);
3268 return mTable[n].value;
3274 return mValueMask.isOn(n);
3279 template<
typename T>
3287 #if defined(__GNUC__) && !defined(__APPLE__) && !defined(__llvm__) 3288 #pragma GCC diagnostic push 3289 #pragma GCC diagnostic ignored "-Wstringop-overflow" 3295 #if defined(__GNUC__) && !defined(__APPLE__) && !defined(__llvm__) 3296 #pragma GCC diagnostic pop 3307 template<
typename ChildT, u
int32_t Log2Dim = ChildT::LOG2DIM + 1>
3318 static constexpr
bool FIXED_SIZE = DataType::FIXED_SIZE;
3319 template<u
int32_t LOG2>
3324 static constexpr uint32_t LOG2DIM = Log2Dim;
3325 static constexpr uint32_t TOTAL = LOG2DIM + ChildT::TOTAL;
3326 static constexpr uint32_t DIM = 1u << TOTAL;
3327 static constexpr uint32_t SIZE = 1u << (3 * LOG2DIM);
3328 static constexpr uint32_t MASK = (1u << TOTAL) - 1u;
3329 static constexpr uint32_t LEVEL = 1 + ChildT::LEVEL;
3330 static constexpr uint64_t NUM_VALUES = uint64_t(1) << (3 * TOTAL);
3333 template <
typename ParentT>
3348 : BaseT(parent->mChildMask.beginOn())
3356 return *mParent->getChild(BaseT::pos());
3361 return mParent->getChild(BaseT::pos());
3366 return (*this)->origin();
3390 : BaseT(parent->data()->mChildMask.beginOff())
3429 : BaseT(parent->data()->mValueMask.beginOn())
3464 , mParent(parent->data())
3471 const ChildT* child =
nullptr;
3472 if (mParent->
mChildMask.isOn(BaseT::pos())) {
3473 child = mParent->
getChild(BaseT::pos());
3475 value = mParent->
getValue(BaseT::pos());
3482 return mParent->
isActive(BaseT::pos());
3487 return mParent->offsetToGlobalCoord(BaseT::pos());
3544 return DataType::mChildMask.isOn(0) ? this->getChild(0)->getFirstValue() : DataType::getValue(0);
3551 return DataType::mChildMask.isOn(SIZE - 1) ? this->getChild(SIZE - 1)->getLastValue() : DataType::getValue(SIZE - 1);
3554 #ifdef NANOVDB_NEW_ACCESSOR_METHODS 3561 #else // NANOVDB_NEW_ACCESSOR_METHODS 3564 const uint32_t n = CoordToOffset(ijk);
3565 return DataType::mChildMask.isOn(n) ? this->getChild(n)->getValue(ijk) : DataType::getValue(n);
3569 const uint32_t n = CoordToOffset(ijk);
3570 return DataType::mChildMask.isOn(n) ? this->getChild(n)->isActive(ijk) : DataType::isActive(n);
3574 const uint32_t n = CoordToOffset(ijk);
3575 if (DataType::mChildMask.isOn(n))
3576 return this->getChild(n)->probeValue(ijk, v);
3577 v = DataType::getValue(n);
3578 return DataType::isActive(n);
3582 const uint32_t n = CoordToOffset(ijk);
3583 if (DataType::mChildMask.isOn(n))
3584 return this->getChild(n)->probeLeaf(ijk);
3588 #endif // NANOVDB_NEW_ACCESSOR_METHODS 3592 const uint32_t n = CoordToOffset(ijk);
3593 return DataType::mChildMask.isOn(n) ? this->getChild(n) :
nullptr;
3597 const uint32_t n = CoordToOffset(ijk);
3598 return DataType::mChildMask.isOn(n) ? this->getChild(n) :
nullptr;
3604 return (((ijk[0] & MASK) >> ChildT::TOTAL) << (2 * LOG2DIM)) |
3605 (((ijk[1] & MASK) >> ChildT::TOTAL) << (LOG2DIM)) |
3606 ((ijk[2] & MASK) >> ChildT::TOTAL);
3613 const uint32_t m = n & ((1 << 2 * LOG2DIM) - 1);
3614 return Coord(n >> 2 * LOG2DIM, m >> LOG2DIM, m & ((1 << LOG2DIM) - 1));
3620 ijk <<= ChildT::TOTAL;
3621 ijk += this->origin();
3627 this->localToGlobalCoord(ijk);
3634 template<
typename OpT,
typename... ArgsT>
3637 const uint32_t n = CoordToOffset(ijk);
3638 if (this->isChild(n))
3639 return this->getChild(n)->template get<OpT>(ijk, args...);
3640 return OpT::get(*
this, n, args...);
3643 template<
typename OpT,
typename... ArgsT>
3645 __hostdev__ decltype(OpT::set(util::declval<InternalNode&>(), util::declval<uint32_t>(), util::declval<ArgsT>()...))
3648 const uint32_t n = CoordToOffset(ijk);
3649 if (this->isChild(n))
3650 return this->getChild(n)->template set<OpT>(ijk, args...);
3651 return OpT::set(*
this, n, args...);
3657 template<
typename,
int,
int,
int>
3662 template<
typename, u
int32_t>
3665 #ifndef NANOVDB_NEW_ACCESSOR_METHODS 3667 template<
typename AccT>
3670 const uint32_t n = CoordToOffset(ijk);
3671 if (DataType::mChildMask.isOff(n))
3672 return DataType::getValue(n);
3673 const ChildT* child = this->getChild(n);
3674 acc.insert(ijk, child);
3675 return child->getValueAndCache(ijk, acc);
3677 template<
typename AccT>
3680 const uint32_t n = CoordToOffset(ijk);
3681 if (DataType::mChildMask.isOff(n))
3682 return DataType::isActive(n);
3683 const ChildT* child = this->getChild(n);
3684 acc.insert(ijk, child);
3685 return child->isActiveAndCache(ijk, acc);
3687 template<
typename AccT>
3690 const uint32_t n = CoordToOffset(ijk);
3691 if (DataType::mChildMask.isOff(n)) {
3692 v = DataType::getValue(n);
3693 return DataType::isActive(n);
3695 const ChildT* child = this->getChild(n);
3696 acc.insert(ijk, child);
3697 return child->probeValueAndCache(ijk, v, acc);
3699 template<
typename AccT>
3702 const uint32_t n = CoordToOffset(ijk);
3703 if (DataType::mChildMask.isOff(n))
3705 const ChildT* child = this->getChild(n);
3706 acc.insert(ijk, child);
3707 return child->probeLeafAndCache(ijk, acc);
3709 template<
typename AccT>
3710 __hostdev__ typename AccT::NodeInfo getNodeInfoAndCache(
const CoordType& ijk,
const AccT& acc)
const 3712 using NodeInfoT =
typename AccT::NodeInfo;
3713 const uint32_t n = CoordToOffset(ijk);
3714 if (DataType::mChildMask.isOff(n)) {
3715 return NodeInfoT{LEVEL, this->dim(), this->minimum(), this->maximum(), this->average(), this->stdDeviation(), this->bbox()[0], this->bbox()[1]};
3717 const ChildT* child = this->getChild(n);
3718 acc.insert(ijk, child);
3719 return child->getNodeInfoAndCache(ijk, acc);
3721 #endif // NANOVDB_NEW_ACCESSOR_METHODS 3723 template<
typename RayT,
typename AccT>
3724 __hostdev__ uint32_t getDimAndCache(
const CoordType& ijk,
const RayT& ray,
const AccT& acc)
const 3726 if (DataType::mFlags & uint32_t(1u))
3730 const uint32_t n = CoordToOffset(ijk);
3731 if (DataType::mChildMask.isOn(n)) {
3732 const ChildT* child = this->getChild(n);
3733 acc.insert(ijk, child);
3734 return child->getDimAndCache(ijk, ray, acc);
3736 return ChildNodeType::dim();
3739 template<
typename OpT,
typename AccT,
typename... ArgsT>
3742 getAndCache(
const CoordType& ijk,
const AccT& acc, ArgsT&&... args)
const 3744 const uint32_t n = CoordToOffset(ijk);
3745 if (DataType::mChildMask.isOff(n))
3746 return OpT::get(*
this, n, args...);
3747 const ChildT* child = this->getChild(n);
3748 acc.insert(ijk, child);
3749 return child->template getAndCache<OpT>(ijk, acc, args...);
3752 template<
typename OpT,
typename AccT,
typename... ArgsT>
3754 __hostdev__ decltype(OpT::set(util::declval<InternalNode&>(), util::declval<uint32_t>(), util::declval<ArgsT>()...))
3755 setAndCache(const
CoordType& ijk, const AccT& acc, ArgsT&&... args)
3757 const uint32_t n = CoordToOffset(ijk);
3758 if (DataType::mChildMask.isOff(n))
3759 return OpT::set(*
this, n, args...);
3760 ChildT* child = this->getChild(n);
3761 acc.insert(ijk, child);
3762 return child->template setAndCache<OpT>(ijk, acc, args...);
3772 template<
typename ValueT,
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
3775 static_assert(
sizeof(CoordT) ==
sizeof(Coord),
"Mismatching sizeof");
3776 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
3781 static constexpr
bool FIXED_SIZE =
true;
3784 uint8_t mBBoxDif[3];
3799 return sizeof(
LeafData) - (12 + 3 + 1 +
sizeof(MaskT<LOG2DIM>) + 2 * (
sizeof(ValueT) +
sizeof(
FloatType)) + (1u << (3 * LOG2DIM)) *
sizeof(ValueT));
3809 mValueMask.setOn(offset);
3810 mValues[offset] = value;
3824 template<
typename T>
3829 for (
auto *p = mValues, *q = p + 512; p != q; ++p)
3843 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
3846 static_assert(
sizeof(CoordT) ==
sizeof(Coord),
"Mismatching sizeof");
3847 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
3852 uint8_t mBBoxDif[3];
3869 return sizeof(
LeafFnBase) - (12 + 3 + 1 +
sizeof(MaskT<LOG2DIM>) + 2 * 4 + 4 * 2);
3874 mQuantum = (max -
min) /
float((1 << bitWidth) - 1);
3904 template<
typename T>
3913 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
3914 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<
Fp4, CoordT, MaskT, LOG2DIM>
3920 static constexpr
bool FIXED_SIZE =
true;
3921 alignas(32) uint8_t mCode[1u << (3 * LOG2DIM - 1)];
3926 static_assert(BaseT::padding() == 0,
"expected no padding in LeafFnBase");
3927 return sizeof(
LeafData) -
sizeof(
BaseT) - (1u << (3 * LOG2DIM - 1));
3934 const uint8_t c = mCode[i>>1];
3935 return ( (i&1) ? c >> 4 : c & uint8_t(15) )*BaseT::mQuantum + BaseT::mMinimum;
3937 return ((mCode[i >> 1] >> ((i & 1) << 2)) & uint8_t(15)) * BaseT::mQuantum + BaseT::mMinimum;
3950 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
3951 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<
Fp8, CoordT, MaskT, LOG2DIM>
3957 static constexpr
bool FIXED_SIZE =
true;
3958 alignas(32) uint8_t mCode[1u << 3 * LOG2DIM];
3962 static_assert(BaseT::padding() == 0,
"expected no padding in LeafFnBase");
3963 return sizeof(
LeafData) -
sizeof(
BaseT) - (1u << 3 * LOG2DIM);
3969 return mCode[i] * BaseT::mQuantum + BaseT::mMinimum;
3980 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
3981 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<
Fp16, CoordT, MaskT, LOG2DIM>
3987 static constexpr
bool FIXED_SIZE =
true;
3988 alignas(32) uint16_t mCode[1u << 3 * LOG2DIM];
3993 static_assert(BaseT::padding() == 0,
"expected no padding in LeafFnBase");
3994 return sizeof(
LeafData) -
sizeof(
BaseT) - 2 * (1u << 3 * LOG2DIM);
4000 return mCode[i] * BaseT::mQuantum + BaseT::mMinimum;
4012 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4013 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<
FpN, CoordT, MaskT, LOG2DIM>
4019 static constexpr
bool FIXED_SIZE =
false;
4022 static_assert(BaseT::padding() == 0,
"expected no padding in LeafFnBase");
4031 #ifdef NANOVDB_FPN_BRANCHLESS // faster 4032 const int b = BaseT::mFlags >> 5;
4034 uint16_t code =
reinterpret_cast<const uint16_t*
>(
this + 1)[i >> (4 - b)];
4035 const static uint8_t shift[5] = {15, 7, 3, 1, 0};
4036 const static uint16_t mask[5] = {1, 3, 15, 255, 65535};
4037 code >>= (i & shift[b]) << b;
4040 uint32_t code =
reinterpret_cast<const uint32_t*
>(
this + 1)[i >> (5 - b)];
4041 code >>= (i & ((32 >> b) - 1)) << b;
4042 code &= (1 << (1 << b)) - 1;
4044 #else // use branched version (slow) 4046 auto* values =
reinterpret_cast<const uint8_t*
>(
this + 1);
4047 switch (BaseT::mFlags >> 5) {
4049 code = float((values[i >> 3] >> (i & 7)) & uint8_t(1));
4052 code = float((values[i >> 2] >> ((i & 3) << 1)) & uint8_t(3));
4055 code = float((values[i >> 1] >> ((i & 1) << 2)) & uint8_t(15));
4058 code = float(values[i]);
4061 code = float(reinterpret_cast<const uint16_t*>(values)[i]);
4064 return float(code) * BaseT::mQuantum + BaseT::mMinimum;
4077 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4078 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<
bool, CoordT, MaskT, LOG2DIM>
4080 static_assert(
sizeof(CoordT) ==
sizeof(Coord),
"Mismatching sizeof");
4081 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
4086 static constexpr
bool FIXED_SIZE =
true;
4089 uint8_t mBBoxDif[3];
4093 uint64_t mPadding[2];
4105 mValueMask.setOn(offset);
4106 mValues.set(offset, v);
4114 template<
typename T>
4127 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4130 static_assert(
sizeof(CoordT) ==
sizeof(Coord),
"Mismatching sizeof");
4131 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
4136 static constexpr
bool FIXED_SIZE =
true;
4139 uint8_t mBBoxDif[3];
4142 uint64_t mPadding[2];
4148 return sizeof(
LeafData) - (12u + 3u + 1u +
sizeof(MaskT<LOG2DIM>) + 2 * 8u);
4163 template<
typename T>
4176 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4179 static_assert(
sizeof(CoordT) ==
sizeof(Coord),
"Mismatching sizeof");
4180 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
4184 static constexpr
bool FIXED_SIZE =
true;
4187 uint8_t mBBoxDif[3];
4193 return sizeof(
LeafIndexBase) - (12u + 3u + 1u +
sizeof(MaskT<LOG2DIM>) + 2 * 8u);
4204 template<
typename T>
4218 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4238 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4246 return util::countOn(BaseT::mValueMask.words()[7]) + (BaseT::mPrefixSum >> 54u & 511u);
4256 uint32_t n = i >> 6;
4257 const uint64_t w = BaseT::mValueMask.words()[n], mask = uint64_t(1) << (i & 63u);
4258 if (!(w & mask))
return uint64_t(0);
4259 uint64_t sum = BaseT::mOffset +
util::countOn(w & (mask - 1u));
4260 if (n--) sum += BaseT::mPrefixSum >> (9u * n) & 511u;
4267 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4269 :
public LeafData<ValueIndex, CoordT, MaskT, LOG2DIM>
4278 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4280 :
public LeafData<ValueOnIndex, CoordT, MaskT, LOG2DIM>
4291 template<
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4292 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT)
LeafData<
Point, CoordT, MaskT, LOG2DIM>
4294 static_assert(
sizeof(CoordT) ==
sizeof(Coord),
"Mismatching sizeof");
4295 static_assert(
sizeof(MaskT<LOG2DIM>) ==
sizeof(
Mask<LOG2DIM>),
"Mismatching sizeof");
4300 static constexpr
bool FIXED_SIZE =
true;
4303 uint8_t mBBoxDif[3];
4309 alignas(32) uint16_t mValues[1u << 3 * LOG2DIM];
4317 return sizeof(
LeafData) - (12u + 3u + 1u +
sizeof(MaskT<LOG2DIM>) + 2 * 8u + (1u << 3 * LOG2DIM) * 2u);
4323 __hostdev__ uint64_t
first(uint32_t i)
const {
return i ? uint64_t(mValues[i - 1u]) + mOffset : mOffset; }
4329 mValueMask.setOn(offset);
4330 mValues[offset] = value;
4344 template<
typename T>
4357 template<
typename BuildT,
4358 typename CoordT = Coord,
4359 template<u
int32_t>
class MaskT =
Mask,
4360 uint32_t Log2Dim = 3>
4366 static constexpr uint32_t TOTAL = 0;
4367 static constexpr uint32_t DIM = 1;
4376 static constexpr
bool FIXED_SIZE = DataType::FIXED_SIZE;
4377 template<u
int32_t LOG2>
4395 : BaseT(parent->data()->mValueMask.beginOn())
4403 return mParent->
getValue(BaseT::pos());
4428 : BaseT(parent->data()->mValueMask.beginOff())
4436 return mParent->
getValue(BaseT::pos());
4457 , mPos(1u << 3 * Log2Dim)
4500 static constexpr uint32_t LOG2DIM = Log2Dim;
4501 static constexpr uint32_t TOTAL = LOG2DIM;
4502 static constexpr uint32_t DIM = 1u << TOTAL;
4503 static constexpr uint32_t SIZE = 1u << 3 * LOG2DIM;
4504 static constexpr uint32_t MASK = (1u << LOG2DIM) - 1u;
4505 static constexpr uint32_t LEVEL = 0;
4506 static constexpr uint64_t NUM_VALUES = uint64_t(1) << (3 * TOTAL);
4542 const uint32_t m = n & ((1 << 2 * LOG2DIM) - 1);
4543 return CoordT(n >> 2 * LOG2DIM, m >> LOG2DIM, m & MASK);
4551 return OffsetToLocalCoord(n) + this->origin();
4560 math::BBox<CoordT> bbox(DataType::mBBoxMin, DataType::mBBoxMin);
4561 if (this->hasBBox()) {
4562 bbox.max()[0] += DataType::mBBoxDif[0];
4563 bbox.max()[1] += DataType::mBBoxDif[1];
4564 bbox.max()[2] += DataType::mBBoxDif[2];
4566 bbox = math::BBox<CoordT>();
4616 return !DataType::mValueMask.isOff();
4624 const uint32_t n = CoordToOffset(ijk);
4625 v = DataType::getValue(n);
4626 return DataType::mValueMask.isOn(n);
4634 return ((ijk[0] & MASK) << (2 * LOG2DIM)) | ((ijk[1] & MASK) << LOG2DIM) | (ijk[2] & MASK);
4646 template<
typename OpT,
typename... ArgsT>
4649 return OpT::get(*
this, CoordToOffset(ijk), args...);
4652 template<
typename OpT,
typename... ArgsT>
4655 return OpT::get(*
this, n, args...);
4658 template<
typename OpT,
typename... ArgsT>
4661 return OpT::set(*
this, CoordToOffset(ijk), args...);
4664 template<
typename OpT,
typename... ArgsT>
4667 return OpT::set(*
this, n, args...);
4673 template<
typename,
int,
int,
int>
4678 template<
typename, u
int32_t>
4681 #ifndef NANOVDB_NEW_ACCESSOR_METHODS 4683 template<
typename AccT>
4684 __hostdev__ ValueType getValueAndCache(
const CoordT& ijk,
const AccT&)
const {
return this->getValue(ijk); }
4687 template<
typename AccT>
4690 using NodeInfoT =
typename AccT::NodeInfo;
4691 return NodeInfoT{LEVEL, this->dim(), this->minimum(), this->maximum(), this->average(), this->stdDeviation(), this->bbox()[0], this->bbox()[1]};
4694 template<
typename AccT>
4695 __hostdev__ bool isActiveAndCache(
const CoordT& ijk,
const AccT&)
const {
return this->isActive(ijk); }
4697 template<
typename AccT>
4698 __hostdev__ bool probeValueAndCache(
const CoordT& ijk,
ValueType& v,
const AccT&)
const {
return this->probeValue(ijk, v); }
4700 template<
typename AccT>
4701 __hostdev__ const LeafNode* probeLeafAndCache(
const CoordT&,
const AccT&)
const {
return this; }
4704 template<
typename RayT,
typename AccT>
4705 __hostdev__ uint32_t getDimAndCache(
const CoordT&,
const RayT& ,
const AccT&)
const 4707 if (DataType::mFlags & uint8_t(1u))
4711 return ChildNodeType::dim();
4714 template<
typename OpT,
typename AccT,
typename... ArgsT>
4717 getAndCache(
const CoordType& ijk,
const AccT&, ArgsT&&... args)
const 4719 return OpT::get(*
this, CoordToOffset(ijk), args...);
4722 template<
typename OpT,
typename AccT,
typename... ArgsT>
4724 __hostdev__ decltype(OpT::set(util::declval<LeafNode&>(), util::declval<uint32_t>(), util::declval<ArgsT>()...))
4725 setAndCache(const
CoordType& ijk, const AccT&, ArgsT&&... args)
4727 return OpT::set(*
this, CoordToOffset(ijk), args...);
4734 template<
typename ValueT,
typename CoordT,
template<u
int32_t>
class MaskT, uint32_t LOG2DIM>
4737 static_assert(LOG2DIM == 3,
"LeafNode::updateBBox: only supports LOGDIM = 3!");
4738 if (DataType::mValueMask.isOff()) {
4739 DataType::mFlags &= ~uint8_t(2);
4742 auto update = [&](uint32_t
min, uint32_t
max,
int axis) {
4744 DataType::mBBoxMin[axis] = (DataType::mBBoxMin[axis] & ~MASK) +
int(min);
4745 DataType::mBBoxDif[axis] = uint8_t(max - min);
4747 uint64_t *w = DataType::mValueMask.words(), word64 = *w;
4748 uint32_t Xmin = word64 ? 0u : 8u, Xmax = Xmin;
4749 for (
int i = 1; i < 8; ++i) {
4758 update(Xmin, Xmax, 0);
4760 const uint32_t *p =
reinterpret_cast<const uint32_t*
>(&word64), word32 = p[0] | p[1];
4761 const uint16_t *q =
reinterpret_cast<const uint16_t*
>(&word32), word16 = q[0] | q[1];
4762 const uint8_t *b =
reinterpret_cast<const uint8_t*
>(&word16), byte = b[0] | b[1];
4765 DataType::mFlags |= uint8_t(2);
4773 template<
typename BuildT>
4775 template<
typename BuildT>
4777 template<
typename BuildT>
4779 template<
typename BuildT>
4781 template<
typename BuildT>
4783 template<
typename BuildT>
4787 template<
typename BuildT,
int LEVEL>
4791 template<
typename BuildT>
4797 template<
typename BuildT>
4803 template<
typename BuildT>
4809 template<
typename BuildT>
4888 template<
typename OpT,
typename GridDataT,
typename... ArgsT>
4892 switch (gridData->mGridType){
4894 return OpT::template known<float>(gridData, args...);
4896 return OpT::template known<double>(gridData, args...);
4898 return OpT::template known<int16_t>(gridData, args...);
4900 return OpT::template known<int32_t>(gridData, args...);
4902 return OpT::template known<int64_t>(gridData, args...);
4904 return OpT::template known<Vec3f>(gridData, args...);
4906 return OpT::template known<Vec3d>(gridData, args...);
4908 return OpT::template known<uint32_t>(gridData, args...);
4910 return OpT::template known<ValueMask>(gridData, args...);
4912 return OpT::template known<ValueIndex>(gridData, args...);
4914 return OpT::template known<ValueOnIndex>(gridData, args...);
4916 return OpT::template known<ValueIndexMask>(gridData, args...);
4918 return OpT::template known<ValueOnIndexMask>(gridData, args...);
4920 return OpT::template known<bool>(gridData, args...);
4922 return OpT::template known<math::Rgba8>(gridData, args...);
4924 return OpT::template known<Fp4>(gridData, args...);
4926 return OpT::template known<Fp8>(gridData, args...);
4928 return OpT::template known<Fp16>(gridData, args...);
4930 return OpT::template known<FpN>(gridData, args...);
4932 return OpT::template known<Vec4f>(gridData, args...);
4934 return OpT::template known<Vec4d>(gridData, args...);
4936 return OpT::template known<uint8_t>(gridData, args...);
4938 return OpT::unknown(gridData, args...);
4963 template<
typename BuildT>
4973 mutable const RootT* mRoot;
4979 static const int CacheLevels = 0;
4980 #ifndef NANOVDB_NEW_ACCESSOR_METHODS 5001 : ReadAccessor(grid.tree().root())
5007 : ReadAccessor(tree.root())
5021 #ifdef NANOVDB_NEW_ACCESSOR_METHODS 5024 return this->
template get<GetValue<BuildT>>(ijk);
5033 #else // NANOVDB_NEW_ACCESSOR_METHODS 5036 return mRoot->getValueAndCache(ijk, *
this);
5040 return this->getValue(
CoordType(i, j, k));
5044 return this->getValue(ijk);
5048 return this->getValue(
CoordType(i, j, k));
5053 return mRoot->getNodeInfoAndCache(ijk, *
this);
5058 return mRoot->isActiveAndCache(ijk, *
this);
5063 return mRoot->probeValueAndCache(ijk, v, *
this);
5068 return mRoot->probeLeafAndCache(ijk, *
this);
5070 #endif // NANOVDB_NEW_ACCESSOR_METHODS 5071 template<
typename RayT>
5074 return mRoot->getDimAndCache(ijk, ray, *
this);
5076 template<
typename OpT,
typename... ArgsT>
5079 return mRoot->template get<OpT>(ijk, args...);
5082 template<
typename OpT,
typename... ArgsT>
5085 return const_cast<RootT*
>(mRoot)->
template set<OpT>(ijk, args...);
5092 template<
typename, u
int32_t>
5094 template<
typename,
typename,
template<u
int32_t>
class, uint32_t>
5098 template<
typename NodeT>
5103 template<
typename BuildT,
int LEVEL0>
5106 static_assert(LEVEL0 >= 0 && LEVEL0 <= 2,
"LEVEL0 should be 0, 1, or 2");
5120 mutable CoordT mKey;
5121 mutable const RootT* mRoot;
5122 mutable const NodeT* mNode;
5129 static const int CacheLevels = 1;
5130 #ifndef NANOVDB_NEW_ACCESSOR_METHODS 5131 using NodeInfo =
typename ReadAccessor<ValueT, -1, -1, -1>::NodeInfo;
5143 : ReadAccessor(grid.tree().root())
5149 : ReadAccessor(tree.root())
5169 return (ijk[0] & int32_t(~NodeT::MASK)) == mKey[0] &&
5170 (ijk[1] & int32_t(~NodeT::MASK)) == mKey[1] &&
5171 (ijk[2] & int32_t(~NodeT::MASK)) == mKey[2];
5174 #ifdef NANOVDB_NEW_ACCESSOR_METHODS 5177 return this->
template get<GetValue<BuildT>>(ijk);
5186 #else // NANOVDB_NEW_ACCESSOR_METHODS 5189 if (this->isCached(ijk))
5190 return mNode->getValueAndCache(ijk, *
this);
5191 return mRoot->getValueAndCache(ijk, *
this);
5195 return this->getValue(
CoordType(i, j, k));
5199 return this->getValue(ijk);
5203 return this->getValue(
CoordType(i, j, k));
5208 if (this->isCached(ijk))
5209 return mNode->getNodeInfoAndCache(ijk, *
this);
5210 return mRoot->getNodeInfoAndCache(ijk, *
this);
5215 if (this->isCached(ijk))
5216 return mNode->isActiveAndCache(ijk, *
this);
5217 return mRoot->isActiveAndCache(ijk, *
this);
5222 if (this->isCached(ijk))
5223 return mNode->probeValueAndCache(ijk, v, *
this);
5224 return mRoot->probeValueAndCache(ijk, v, *
this);
5229 if (this->isCached(ijk))
5230 return mNode->probeLeafAndCache(ijk, *
this);
5231 return mRoot->probeLeafAndCache(ijk, *
this);
5233 #endif // NANOVDB_NEW_ACCESSOR_METHODS 5234 template<
typename RayT>
5237 if (this->isCached(ijk))
5238 return mNode->getDimAndCache(ijk, ray, *
this);
5239 return mRoot->getDimAndCache(ijk, ray, *
this);
5242 template<
typename OpT,
typename... ArgsT>
5245 if (this->isCached(ijk))
5246 return mNode->template getAndCache<OpT>(ijk, *
this, args...);
5247 return mRoot->template getAndCache<OpT>(ijk, *
this, args...);
5250 template<
typename OpT,
typename... ArgsT>
5253 if (this->isCached(ijk))
5254 return const_cast<NodeT*>(mNode)->template setAndCache<OpT>(ijk, *
this, args...);
5255 return const_cast<RootT*
>(mRoot)->
template setAndCache<OpT>(ijk, *
this, args...);
5262 template<
typename, u
int32_t>
5264 template<
typename,
typename,
template<u
int32_t>
class, uint32_t>
5270 mKey = ijk & ~NodeT::MASK;
5275 template<
typename OtherNodeT>
5280 template<
typename BuildT,
int LEVEL0,
int LEVEL1>
5283 static_assert(LEVEL0 >= 0 && LEVEL0 <= 2,
"LEVEL0 must be 0, 1, 2");
5284 static_assert(LEVEL1 >= 0 && LEVEL1 <= 2,
"LEVEL1 must be 0, 1, 2");
5285 static_assert(LEVEL0 < LEVEL1,
"Level 0 must be lower than level 1");
5298 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY // 44 bytes total 5299 mutable CoordT mKey;
5300 #else // 68 bytes total 5301 mutable CoordT mKeys[2];
5303 mutable const RootT* mRoot;
5304 mutable const Node1T* mNode1;
5305 mutable const Node2T* mNode2;
5312 static const int CacheLevels = 2;
5313 #ifndef NANOVDB_NEW_ACCESSOR_METHODS 5314 using NodeInfo =
typename ReadAccessor<ValueT, -1, -1, -1>::NodeInfo;
5318 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5331 : ReadAccessor(grid.tree().root())
5337 : ReadAccessor(tree.root())
5344 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5360 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5361 __hostdev__ bool isCached1(CoordValueType dirty)
const 5365 if (dirty & int32_t(~Node1T::MASK)) {
5371 __hostdev__ bool isCached2(CoordValueType dirty)
const 5375 if (dirty & int32_t(~Node2T::MASK)) {
5383 return (ijk[0] ^ mKey[0]) | (ijk[1] ^ mKey[1]) | (ijk[2] ^ mKey[2]);
5388 return (ijk[0] & int32_t(~Node1T::MASK)) == mKeys[0][0] &&
5389 (ijk[1] & int32_t(~Node1T::MASK)) == mKeys[0][1] &&
5390 (ijk[2] & int32_t(~Node1T::MASK)) == mKeys[0][2];
5394 return (ijk[0] & int32_t(~Node2T::MASK)) == mKeys[1][0] &&
5395 (ijk[1] & int32_t(~Node2T::MASK)) == mKeys[1][1] &&
5396 (ijk[2] & int32_t(~Node2T::MASK)) == mKeys[1][2];
5400 #ifdef NANOVDB_NEW_ACCESSOR_METHODS 5403 return this->
template get<GetValue<BuildT>>(ijk);
5412 #else // NANOVDB_NEW_ACCESSOR_METHODS 5416 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5417 const CoordValueType dirty = this->computeDirty(ijk);
5421 if (this->isCached1(dirty)) {
5422 return mNode1->getValueAndCache(ijk, *
this);
5423 }
else if (this->isCached2(dirty)) {
5424 return mNode2->getValueAndCache(ijk, *
this);
5426 return mRoot->getValueAndCache(ijk, *
this);
5430 return this->getValue(ijk);
5434 return this->getValue(
CoordType(i, j, k));
5438 return this->getValue(
CoordType(i, j, k));
5442 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5443 const CoordValueType dirty = this->computeDirty(ijk);
5447 if (this->isCached1(dirty)) {
5448 return mNode1->getNodeInfoAndCache(ijk, *
this);
5449 }
else if (this->isCached2(dirty)) {
5450 return mNode2->getNodeInfoAndCache(ijk, *
this);
5452 return mRoot->getNodeInfoAndCache(ijk, *
this);
5457 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5458 const CoordValueType dirty = this->computeDirty(ijk);
5462 if (this->isCached1(dirty)) {
5463 return mNode1->isActiveAndCache(ijk, *
this);
5464 }
else if (this->isCached2(dirty)) {
5465 return mNode2->isActiveAndCache(ijk, *
this);
5467 return mRoot->isActiveAndCache(ijk, *
this);
5472 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5473 const CoordValueType dirty = this->computeDirty(ijk);
5477 if (this->isCached1(dirty)) {
5478 return mNode1->probeValueAndCache(ijk, v, *
this);
5479 }
else if (this->isCached2(dirty)) {
5480 return mNode2->probeValueAndCache(ijk, v, *
this);
5482 return mRoot->probeValueAndCache(ijk, v, *
this);
5487 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5488 const CoordValueType dirty = this->computeDirty(ijk);
5492 if (this->isCached1(dirty)) {
5493 return mNode1->probeLeafAndCache(ijk, *
this);
5494 }
else if (this->isCached2(dirty)) {
5495 return mNode2->probeLeafAndCache(ijk, *
this);
5497 return mRoot->probeLeafAndCache(ijk, *
this);
5499 #endif // NANOVDB_NEW_ACCESSOR_METHODS 5501 template<
typename RayT>
5504 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5505 const CoordValueType dirty = this->computeDirty(ijk);
5509 if (this->isCached1(dirty)) {
5510 return mNode1->getDimAndCache(ijk, ray, *
this);
5511 }
else if (this->isCached2(dirty)) {
5512 return mNode2->getDimAndCache(ijk, ray, *
this);
5514 return mRoot->getDimAndCache(ijk, ray, *
this);
5517 template<
typename OpT,
typename... ArgsT>
5520 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5521 const CoordValueType dirty = this->computeDirty(ijk);
5525 if (this->isCached1(dirty)) {
5526 return mNode1->template getAndCache<OpT>(ijk, *
this, args...);
5527 }
else if (this->isCached2(dirty)) {
5528 return mNode2->template getAndCache<OpT>(ijk, *
this, args...);
5530 return mRoot->template getAndCache<OpT>(ijk, *
this, args...);
5533 template<
typename OpT,
typename... ArgsT>
5536 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5537 const CoordValueType dirty = this->computeDirty(ijk);
5541 if (this->isCached1(dirty)) {
5542 return const_cast<Node1T*
>(mNode1)->
template setAndCache<OpT>(ijk, *
this, args...);
5543 }
else if (this->isCached2(dirty)) {
5544 return const_cast<Node2T*
>(mNode2)->
template setAndCache<OpT>(ijk, *
this, args...);
5546 return const_cast<RootT*
>(mRoot)->
template setAndCache<OpT>(ijk, *
this, args...);
5553 template<
typename, u
int32_t>
5555 template<
typename,
typename,
template<u
int32_t>
class, uint32_t>
5561 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5564 mKeys[0] = ijk & ~Node1T::MASK;
5570 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5573 mKeys[1] = ijk & ~Node2T::MASK;
5577 template<
typename OtherNodeT>
5582 template<
typename BuildT>
5598 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY // 44 bytes total 5599 mutable CoordT mKey;
5600 #else // 68 bytes total 5601 mutable CoordT mKeys[3];
5603 mutable const RootT* mRoot;
5604 mutable const void* mNode[3];
5611 static const int CacheLevels = 3;
5612 #ifndef NANOVDB_NEW_ACCESSOR_METHODS 5613 using NodeInfo =
typename ReadAccessor<ValueT, -1, -1, -1>::NodeInfo;
5617 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5623 , mNode{
nullptr,
nullptr,
nullptr}
5629 : ReadAccessor(grid.tree().root())
5635 : ReadAccessor(tree.root())
5649 template<
typename NodeT>
5654 return reinterpret_cast<const T*
>(mNode[NodeT::LEVEL]);
5661 static_assert(LEVEL >= 0 && LEVEL <= 2,
"ReadAccessor::getNode: Invalid node type");
5662 return reinterpret_cast<const T*
>(mNode[LEVEL]);
5668 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5673 mNode[0] = mNode[1] = mNode[2] =
nullptr;
5676 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5677 template<
typename NodeT>
5678 __hostdev__ bool isCached(CoordValueType dirty)
const 5680 if (!mNode[NodeT::LEVEL])
5682 if (dirty & int32_t(~NodeT::MASK)) {
5683 mNode[NodeT::LEVEL] =
nullptr;
5691 return (ijk[0] ^ mKey[0]) | (ijk[1] ^ mKey[1]) | (ijk[2] ^ mKey[2]);
5694 template<
typename NodeT>
5697 return (ijk[0] & int32_t(~NodeT::MASK)) == mKeys[NodeT::LEVEL][0] &&
5698 (ijk[1] & int32_t(~NodeT::MASK)) == mKeys[NodeT::LEVEL][1] &&
5699 (ijk[2] & int32_t(~NodeT::MASK)) == mKeys[NodeT::LEVEL][2];
5703 #ifdef NANOVDB_NEW_ACCESSOR_METHODS 5706 return this->
template get<GetValue<BuildT>>(ijk);
5715 #else // NANOVDB_NEW_ACCESSOR_METHODS 5719 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5720 const CoordValueType dirty = this->computeDirty(ijk);
5724 if (this->isCached<LeafT>(dirty)) {
5725 return ((
LeafT*)mNode[0])->getValue(ijk);
5726 }
else if (this->isCached<NodeT1>(dirty)) {
5727 return ((
NodeT1*)mNode[1])->getValueAndCache(ijk, *
this);
5728 }
else if (this->isCached<NodeT2>(dirty)) {
5729 return ((
NodeT2*)mNode[2])->getValueAndCache(ijk, *
this);
5731 return mRoot->getValueAndCache(ijk, *
this);
5735 return this->getValue(ijk);
5739 return this->getValue(
CoordType(i, j, k));
5743 return this->getValue(
CoordType(i, j, k));
5748 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5749 const CoordValueType dirty = this->computeDirty(ijk);
5753 if (this->isCached<LeafT>(dirty)) {
5754 return ((
LeafT*)mNode[0])->getNodeInfoAndCache(ijk, *
this);
5755 }
else if (this->isCached<NodeT1>(dirty)) {
5756 return ((
NodeT1*)mNode[1])->getNodeInfoAndCache(ijk, *
this);
5757 }
else if (this->isCached<NodeT2>(dirty)) {
5758 return ((
NodeT2*)mNode[2])->getNodeInfoAndCache(ijk, *
this);
5760 return mRoot->getNodeInfoAndCache(ijk, *
this);
5765 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5766 const CoordValueType dirty = this->computeDirty(ijk);
5770 if (this->isCached<LeafT>(dirty)) {
5771 return ((
LeafT*)mNode[0])->isActive(ijk);
5772 }
else if (this->isCached<NodeT1>(dirty)) {
5773 return ((
NodeT1*)mNode[1])->isActiveAndCache(ijk, *
this);
5774 }
else if (this->isCached<NodeT2>(dirty)) {
5775 return ((
NodeT2*)mNode[2])->isActiveAndCache(ijk, *
this);
5777 return mRoot->isActiveAndCache(ijk, *
this);
5782 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5783 const CoordValueType dirty = this->computeDirty(ijk);
5787 if (this->isCached<LeafT>(dirty)) {
5788 return ((
LeafT*)mNode[0])->probeValue(ijk, v);
5789 }
else if (this->isCached<NodeT1>(dirty)) {
5790 return ((
NodeT1*)mNode[1])->probeValueAndCache(ijk, v, *
this);
5791 }
else if (this->isCached<NodeT2>(dirty)) {
5792 return ((
NodeT2*)mNode[2])->probeValueAndCache(ijk, v, *
this);
5794 return mRoot->probeValueAndCache(ijk, v, *
this);
5798 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5799 const CoordValueType dirty = this->computeDirty(ijk);
5803 if (this->isCached<LeafT>(dirty)) {
5804 return ((
LeafT*)mNode[0]);
5805 }
else if (this->isCached<NodeT1>(dirty)) {
5806 return ((
NodeT1*)mNode[1])->probeLeafAndCache(ijk, *
this);
5807 }
else if (this->isCached<NodeT2>(dirty)) {
5808 return ((
NodeT2*)mNode[2])->probeLeafAndCache(ijk, *
this);
5810 return mRoot->probeLeafAndCache(ijk, *
this);
5812 #endif // NANOVDB_NEW_ACCESSOR_METHODS 5814 template<
typename OpT,
typename... ArgsT>
5817 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5818 const CoordValueType dirty = this->computeDirty(ijk);
5822 if (this->isCached<LeafT>(dirty)) {
5823 return ((
const LeafT*)mNode[0])->template getAndCache<OpT>(ijk, *
this, args...);
5824 }
else if (this->isCached<NodeT1>(dirty)) {
5825 return ((
const NodeT1*)mNode[1])->template getAndCache<OpT>(ijk, *
this, args...);
5826 }
else if (this->isCached<NodeT2>(dirty)) {
5827 return ((
const NodeT2*)mNode[2])->template getAndCache<OpT>(ijk, *
this, args...);
5829 return mRoot->template getAndCache<OpT>(ijk, *
this, args...);
5832 template<
typename OpT,
typename... ArgsT>
5835 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5836 const CoordValueType dirty = this->computeDirty(ijk);
5840 if (this->isCached<LeafT>(dirty)) {
5841 return ((
LeafT*)mNode[0])->template setAndCache<OpT>(ijk, *
this, args...);
5842 }
else if (this->isCached<NodeT1>(dirty)) {
5843 return ((
NodeT1*)mNode[1])->template setAndCache<OpT>(ijk, *
this, args...);
5844 }
else if (this->isCached<NodeT2>(dirty)) {
5845 return ((
NodeT2*)mNode[2])->template setAndCache<OpT>(ijk, *
this, args...);
5847 return ((
RootT*)mRoot)->template setAndCache<OpT>(ijk, *
this, args...);
5850 template<
typename RayT>
5853 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5854 const CoordValueType dirty = this->computeDirty(ijk);
5858 if (this->isCached<LeafT>(dirty)) {
5859 return ((
LeafT*)mNode[0])->getDimAndCache(ijk, ray, *
this);
5860 }
else if (this->isCached<NodeT1>(dirty)) {
5861 return ((
NodeT1*)mNode[1])->getDimAndCache(ijk, ray, *
this);
5862 }
else if (this->isCached<NodeT2>(dirty)) {
5863 return ((
NodeT2*)mNode[2])->getDimAndCache(ijk, ray, *
this);
5865 return mRoot->getDimAndCache(ijk, ray, *
this);
5872 template<
typename, u
int32_t>
5874 template<
typename,
typename,
template<u
int32_t>
class, uint32_t>
5878 template<
typename NodeT>
5881 #ifdef NANOVDB_USE_SINGLE_ACCESSOR_KEY 5884 mKeys[NodeT::LEVEL] = ijk & ~NodeT::MASK;
5886 mNode[NodeT::LEVEL] = node;
5904 template<
int LEVEL0 = -1,
int LEVEL1 = -1,
int LEVEL2 = -1,
typename ValueT =
float>
5910 template<
int LEVEL0 = -1,
int LEVEL1 = -1,
int LEVEL2 = -1,
typename ValueT =
float>
5916 template<
int LEVEL0 = -1,
int LEVEL1 = -1,
int LEVEL2 = -1,
typename ValueT =
float>
5932 CoordBBox mIndexBBox;
5933 uint32_t mRootTableSize, mPadding{0};
5936 template<
typename T>
5939 mGridData = *grid.
data();
5940 mTreeData = *grid.
tree().data();
5942 mRootTableSize = grid.
tree().root().getTableSize();
5947 *
this = *
reinterpret_cast<const GridMetaData*
>(gridData);
5950 mGridData = *gridData;
5970 template<
typename T>
6010 template<
typename AttT,
typename BuildT = u
int32_t>
6019 : AccT(grid.tree().root())
6021 , mData(grid.template getBlindData<AttT>(0))
6039 end = begin + count;
6047 auto* leaf = this->probeLeaf(ijk);
6048 if (leaf ==
nullptr) {
6051 begin = mData + leaf->minimum();
6052 end = begin + leaf->maximum();
6053 return leaf->maximum();
6059 begin = end =
nullptr;
6060 if (
auto* leaf = this->probeLeaf(ijk)) {
6062 if (leaf->isActive(offset)) {
6063 begin = mData + leaf->minimum();
6064 end = begin + leaf->getValue(offset);
6066 begin += leaf->getValue(offset - 1);
6073 template<
typename AttT>
6082 : AccT(grid.tree().root())
6084 , mData(grid.template getBlindData<AttT>(0))
6106 end = begin + count;
6114 auto* leaf = this->probeLeaf(ijk);
6115 if (leaf ==
nullptr)
6117 begin = mData + leaf->offset();
6118 end = begin + leaf->pointCount();
6119 return leaf->pointCount();
6125 if (
auto* leaf = this->probeLeaf(ijk)) {
6127 if (leaf->isActive(n)) {
6128 begin = mData + leaf->first(n);
6129 end = mData + leaf->last(n);
6133 begin = end =
nullptr;
6141 template<
typename ChannelT,
typename IndexT = ValueIndex>
6158 :
BaseT(grid.tree().root())
6164 this->setChannel(channelID);
6169 :
BaseT(grid.tree().root())
6171 , mChannel(channelPtr)
6202 return mChannel =
const_cast<ChannelT*
>(mGrid.template getBlindData<ChannelT>(channelID));
6207 __hostdev__ uint64_t
idx(
int i,
int j,
int k)
const {
return BaseT::getValue(math::Coord(i, j, k)); }
6218 const bool isActive = BaseT::probeValue(ijk, idx);
6225 template<
typename T>
6226 __hostdev__ T&
getValue(
const math::Coord& ijk, T* channelPtr)
const {
return channelPtr[BaseT::getValue(ijk)]; }
6234 struct MiniGridHandle {
6239 BufferType(BufferType &&other) : data(other.data), size(other.size) {other.data=
nullptr; other.size=0;}
6240 ~BufferType() {std::free(data);}
6241 BufferType& operator=(
const BufferType &other) =
delete;
6242 BufferType& operator=(BufferType &&other){data=other.data; size=other.size; other.data=
nullptr; other.size=0;
return *
this;}
6243 static BufferType create(
size_t n, BufferType* dummy =
nullptr) {
return BufferType(n);}
6245 MiniGridHandle(BufferType &&buf) : buffer(std::move(buf)) {}
6246 const uint8_t* data()
const {
return buffer.data;}
6278 if (
util::streq(str,
"blosc"))
return Codec::BLOSC;
6320 uint32_t nodeCount[4];
6321 uint32_t tileCount[3];
6328 #if !defined(__CUDA_ARCH__) && !defined(__HIP__) 6351 template<
typename StreamT>
6357 #ifdef NANOVDB_USE_NEW_MAGIC_NUMBERS 6362 const char* gridName = gridData->
gridName();
6371 os.write((
const char*)&head,
sizeof(
FileHeader));
6373 os.write(gridName, nameSize);
6375 os.write((
const char*)gridData, gridData->
mGridSize);
6380 template<
typename GridHandleT,
template<
typename...>
class VecT>
6383 #ifdef NANOVDB_USE_IOSTREAMS // use this to switch between std::ofstream or FILE implementations 6384 std::ofstream os(fileName, std::ios::out | std::ios::binary | std::ios::trunc);
6388 StreamT(
const char* name) { fptr = fopen(name,
"wb"); }
6389 ~StreamT() { fclose(fptr); }
6390 void write(
const char* data,
size_t n) { fwrite(data, 1, n, fptr); }
6391 bool is_open()
const {
return fptr != NULL; }
6394 if (!os.is_open()) {
6395 fprintf(stderr,
"nanovdb::writeUncompressedGrids: Unable to open file \"%s\"for output\n", fileName);
6398 for (
auto& h : handles) {
6408 template<
typename GridHandleT,
typename StreamT,
template<
typename...>
class VecT>
6409 VecT<GridHandleT>
readUncompressedGrids(StreamT& is,
const typename GridHandleT::BufferType& pool =
typename GridHandleT::BufferType())
6411 VecT<GridHandleT> handles;
6413 is.read((
char*)&data,
sizeof(
GridData));
6415 uint64_t size = data.
mGridSize, sum = 0u;
6418 is.read((
char*)&data,
sizeof(
GridData));
6421 is.skip(-int64_t(sum +
sizeof(
GridData)));
6422 auto buffer = GridHandleT::BufferType::create(size + sum, &pool);
6423 is.read((
char*)(buffer.data()), buffer.size());
6424 handles.emplace_back(std::move(buffer));
6428 while(is.read((
char*)&head,
sizeof(
FileHeader))) {
6430 fprintf(stderr,
"nanovdb::readUncompressedGrids: invalid magic number = \"%s\"\n", (
const char*)&(head.
magic));
6434 fprintf(stderr,
"nanovdb::readUncompressedGrids: invalid major version = \"%s\"\n",
toStr(str, head.
version));
6438 fprintf(stderr,
"nanovdb::readUncompressedGrids: invalid codec = \"%s\"\n",
toStr(str, head.
codec));
6442 for (uint16_t i = 0; i < head.
gridCount; ++i) {
6445 auto buffer = GridHandleT::BufferType::create(meta.
gridSize, &pool);
6446 is.read((
char*)buffer.data(), meta.
gridSize);
6447 handles.emplace_back(std::move(buffer));
6455 template<
typename GridHandleT,
template<
typename...>
class VecT>
6456 VecT<GridHandleT>
readUncompressedGrids(
const char* fileName,
const typename GridHandleT::BufferType& buffer =
typename GridHandleT::BufferType())
6458 #ifdef NANOVDB_USE_IOSTREAMS // use this to switch between std::ifstream or FILE implementations 6459 struct StreamT :
public std::ifstream {
6460 StreamT(
const char* name) : std::ifstream(name, std::ios::in | std::ios::binary){}
6461 void skip(int64_t off) { this->seekg(off, std::ios_base::cur); }
6466 StreamT(
const char* name) { fptr = fopen(name,
"rb"); }
6467 ~StreamT() { fclose(fptr); }
6468 bool read(
char* data,
size_t n) {
6469 size_t m = fread(data, 1, n, fptr);
6472 void skip(int64_t off) { fseek(fptr, (
long int)off, SEEK_CUR); }
6473 bool is_open()
const {
return fptr != NULL; }
6476 StreamT is(fileName);
6477 if (!is.is_open()) {
6478 fprintf(stderr,
"nanovdb::readUncompressedGrids: Unable to open file \"%s\"for input\n", fileName);
6481 return readUncompressedGrids<GridHandleT, StreamT, VecT>(is, buffer);
6484 #endif // if !defined(__CUDA_ARCH__) && !defined(__HIP__) 6493 template<
typename BuildT>
6503 template<
typename BuildT>
6515 template<
typename BuildT>
6529 template<
typename BuildT>
6541 template<
typename BuildT>
6553 template<
typename BuildT>
6565 template<
typename BuildT>
6577 template<
typename BuildT>
6589 template<
typename BuildT>
6595 v = root.mBackground;
6601 return tile.state > 0u;
6605 v = node.mTable[n].value;
6606 return node.mValueMask.isOn(n);
6610 v = node.mTable[n].value;
6611 return node.mValueMask.isOn(n);
6615 v = leaf.getValue(n);
6616 return leaf.mValueMask.isOn(n);
6622 template<
typename BuildT>
6644 return NodeInfo{2u, node.
dim(), node.minimum(), node.maximum(), node.average(), node.stdDeviation(), node.bbox()};
6648 return NodeInfo{1u, node.
dim(), node.minimum(), node.maximum(), node.average(), node.stdDeviation(), node.bbox()};
6652 return NodeInfo{0u, leaf.
dim(), leaf.minimum(), leaf.maximum(), leaf.average(), leaf.stdDeviation(), leaf.bbox()};
6658 #endif // end of NANOVDB_NANOVDB_H_HAS_BEEN_INCLUDED typename FloatTraits< BuildT >::FloatType FloatType
Definition: NanoVDB.h:3779
__hostdev__ ValueType getMin() const
Definition: NanoVDB.h:3814
__hostdev__ ValueOffIterator beginValueOff() const
Definition: NanoVDB.h:4445
__hostdev__ DenseIter()
Definition: NanoVDB.h:2874
__hostdev__ const GridType & gridType() const
Definition: NanoVDB.h:2174
__hostdev__ bool probeValue(const math::Coord &ijk, typename util::remove_const< ChannelT >::type &v) const
return the state and updates the value of the specified voxel
Definition: NanoVDB.h:6215
static __hostdev__ constexpr uint32_t padding()
Definition: NanoVDB.h:3924
typename BuildT::RootType RootType
Definition: NanoVDB.h:2049
__hostdev__ const Vec3d & voxelSize() const
Return a const reference to the size of a voxel in world units.
Definition: NanoVDB.h:2109
__hostdev__ ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:5704
__hostdev__ uint32_t operator*() const
Definition: NanoVDB.h:1082
ValueT ValueType
Definition: NanoVDB.h:5126
__hostdev__ uint64_t full() const
Definition: NanoVDB.h:1777
__hostdev__ const char * shortGridName() const
Return a c-string with the name of this grid, truncated to 255 characters.
Definition: NanoVDB.h:2208
__hostdev__ util::enable_if<!util::is_same< MaskT, Mask >::value, Mask & >::type operator=(const MaskT &other)
Assignment operator that works with openvdb::util::NodeMask.
Definition: NanoVDB.h:1179
__hostdev__ const ValueType & minimum() const
Return a const reference to the minimum active value encoded in this root node and any of its child n...
Definition: NanoVDB.h:2947
bool type
Definition: NanoVDB.h:499
Visits all tile values in this node, i.e. both inactive and active tiles.
Definition: NanoVDB.h:3378
__hostdev__ math::BBox< CoordT > bbox() const
Return the bounding box in index space of active values in this leaf node.
Definition: NanoVDB.h:4558
__hostdev__ CoordT getCoord() const
Definition: NanoVDB.h:4472
uint16_t ArrayType
Definition: NanoVDB.h:4299
__hostdev__ void next()
Definition: NanoVDB.h:2709
__hostdev__ CheckMode toCheckMode(const Checksum &checksum)
Maps 64 bit checksum to CheckMode enum.
Definition: NanoVDB.h:1816
C++11 implementation of std::enable_if.
Definition: Util.h:335
FloatType mStdDevi
Definition: NanoVDB.h:3791
float type
Definition: NanoVDB.h:506
static __hostdev__ constexpr uint32_t padding()
Definition: NanoVDB.h:4020
__hostdev__ bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:5713
__hostdev__ CoordT offsetToGlobalCoord(uint32_t n) const
Definition: NanoVDB.h:4549
static __hostdev__ uint64_t memUsage()
Definition: NanoVDB.h:4144
__hostdev__ bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:5030
Type Pow2(Type x)
Return x2.
Definition: Math.h:548
__hostdev__ const MaskType< LOG2DIM > & getValueMask() const
Definition: NanoVDB.h:4514
typename util::match_const< DataType, RootT >::type DataT
Definition: NanoVDB.h:2695
void writeUncompressedGrids(const char *fileName, const VecT< GridHandleT > &handles, bool raw=false)
write multiple NanoVDB grids to a single file, without compression.
Definition: NanoVDB.h:6381
typename RootType::LeafNodeType LeafNodeType
Definition: NanoVDB.h:2354
__hostdev__ ValueType getValue(int i, int j, int k) const
Definition: NanoVDB.h:2973
__hostdev__ Vec3d getVoxelSize() const
Return a voxels size in each coordinate direction, measured at the origin.
Definition: NanoVDB.h:1518
__hostdev__ ReadAccessor(const GridT &grid)
Constructor from a grid.
Definition: NanoVDB.h:5000
StatsT mStdDevi
Definition: NanoVDB.h:3226
__hostdev__ bool hasStdDeviation() const
Definition: NanoVDB.h:2188
__hostdev__ Vec3T applyMap(const Vec3T &xyz) const
Definition: NanoVDB.h:1924
NANOVDB_HOSTDEV_DISABLE_WARNING __hostdev__ uint32_t findFirst() const
Definition: NanoVDB.h:1340
__hostdev__ TileT * tile() const
Definition: NanoVDB.h:2710
__hostdev__ bool isOff(uint32_t n) const
Return true if the given bit is NOT set.
Definition: NanoVDB.h:1208
__hostdev__ Vec3T applyMapF(const Vec3T &xyz) const
Definition: NanoVDB.h:1935
__hostdev__ const char * gridName() const
Definition: NanoVDB.h:1992
__hostdev__ ChannelT * setChannel(ChannelT *channelPtr)
Change to an external channel.
Definition: NanoVDB.h:6194
typename util::match_const< Tile, RootT >::type TileT
Definition: NanoVDB.h:2696
__hostdev__ ChildT * getChild(uint32_t n)
Returns a pointer to the child node at the specifed linear offset.
Definition: NanoVDB.h:3254
__hostdev__ ValueType operator()(const CoordType &ijk) const
Definition: NanoVDB.h:5709
__hostdev__ Vec3T applyIJTF(const Vec3T &xyz) const
Definition: NanoVDB.h:1515
VDB Tree, which is a thin wrapper around a RootNode.
Definition: NanoVDB.h:2341
__hostdev__ Vec3T applyMapF(const Vec3T &ijk) const
Apply the forward affine transformation to a vector using 32bit floating point arithmetics.
Definition: NanoVDB.h:1446
decltype(mFlags) Type
Definition: NanoVDB.h:931
__hostdev__ Vec3T indexToWorld(const Vec3T &xyz) const
index to world space transformation
Definition: NanoVDB.h:2120
math::BBox< CoordType > BBoxType
Definition: NanoVDB.h:2684
__hostdev__ Tile * tile(uint32_t n)
Definition: NanoVDB.h:2598
__hostdev__ DenseIter operator++(int)
Definition: NanoVDB.h:2906
__hostdev__ bool isActive() const
Definition: NanoVDB.h:2796
__hostdev__ ValueType operator()(int i, int j, int k) const
Definition: NanoVDB.h:5710
__hostdev__ GridClass mapToGridClass(GridClass defaultClass=GridClass::Unknown)
Definition: NanoVDB.h:894
__hostdev__ bool isChild() const
Definition: NanoVDB.h:2580
__hostdev__ ValueType getValue(int i, int j, int k) const
Definition: NanoVDB.h:2386
__hostdev__ ValueIterator()
Definition: NanoVDB.h:4455
float Type
Definition: NanoVDB.h:526
float FloatType
Definition: NanoVDB.h:3849
__hostdev__ CoordT origin() const
Return the origin in index space of this leaf node.
Definition: NanoVDB.h:4534
Highest level of the data structure. Contains a tree and a world->index transform (that currently onl...
Definition: NanoVDB.h:2045
__hostdev__ ReadAccessor(const TreeT &tree)
Constructor from a tree.
Definition: NanoVDB.h:5006
__hostdev__ ValueOnIter(RootT *parent)
Definition: NanoVDB.h:2833
Vec3dBBox mWorldBBox
Definition: NanoVDB.h:1856
__hostdev__ CoordType getOrigin() const
Definition: NanoVDB.h:3363
__hostdev__ const NodeTrait< RootT, 1 >::type * getFirstLower() const
Definition: NanoVDB.h:2480
__hostdev__ Vec3T applyIJTF(const Vec3T &xyz) const
Definition: NanoVDB.h:1943
FloatType stdDevi
Definition: NanoVDB.h:6631
__hostdev__ char * toStr(char *dst, GridType gridType)
Maps a GridType to a c-string.
Definition: NanoVDB.h:253
__hostdev__ ValueType maximum() const
Return a const reference to the maximum active value encoded in this leaf node.
Definition: NanoVDB.h:4520
__hostdev__ DenseIterator(const InternalNode *parent)
Definition: NanoVDB.h:3462
static __hostdev__ uint64_t memUsage()
Definition: NanoVDB.h:4273
__hostdev__ const DataType * data() const
Definition: NanoVDB.h:2931
__hostdev__ const MaskType< LOG2DIM > & valueMask() const
Return a const reference to the bit mask of active voxels in this internal node.
Definition: NanoVDB.h:3512
__hostdev__ const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:5411
#define NANOVDB_PATCH_VERSION_NUMBER
Definition: NanoVDB.h:146
__hostdev__ void init(std::initializer_list< GridFlags > list={GridFlags::IsBreadthFirst}, uint64_t gridSize=0u, const Map &map=Map(), GridType gridType=GridType::Unknown, GridClass gridClass=GridClass::Unknown)
Definition: NanoVDB.h:1868
__hostdev__ ValueType operator()(const CoordType &ijk) const
Definition: NanoVDB.h:5180
static __hostdev__ constexpr uint64_t memUsage()
Definition: NanoVDB.h:3923
__hostdev__ bool getValue(uint32_t i) const
Definition: NanoVDB.h:4151
__hostdev__ void setValueOnly(uint32_t offset, const ValueType &v)
Sets the value at the specified location but leaves its state unchanged.
Definition: NanoVDB.h:4604
__hostdev__ Vec3T applyInverseMap(const Vec3T &xyz) const
Apply the inverse affine mapping to a vector using 64bit floating point arithmetics.
Definition: NanoVDB.h:1472
__hostdev__ ValueOnIter()
Definition: NanoVDB.h:2829
Class to access values in channels at a specific voxel location.
Definition: NanoVDB.h:6142
__hostdev__ void setMask(uint32_t offset, bool v)
Definition: NanoVDB.h:4286
__hostdev__ void setOn(uint32_t offset)
Definition: NanoVDB.h:3812
static __hostdev__ uint32_t padding()
Definition: NanoVDB.h:4574
typename GridT::TreeType Type
Definition: NanoVDB.h:2327
__hostdev__ NodeT * operator->() const
Definition: NanoVDB.h:2747
__hostdev__ ValueType getValue(int i, int j, int k) const
Definition: NanoVDB.h:5708
char mGridName[MaxNameSize]
Definition: NanoVDB.h:1854
__hostdev__ bool operator>(const Version &rhs) const
Definition: NanoVDB.h:703
static __hostdev__ size_t memUsage(uint32_t bitWidth)
Definition: NanoVDB.h:4028
__hostdev__ void setChild(const CoordType &k, const void *ptr, const RootData *data)
Definition: NanoVDB.h:2566
__hostdev__ Version version() const
Definition: NanoVDB.h:2067
PointAccessor(const NanoGrid< Point > &grid)
Definition: NanoVDB.h:6081
__hostdev__ const ValueT & getMax() const
Definition: NanoVDB.h:2649
__hostdev__ ValueType getValue(uint32_t i) const
Definition: NanoVDB.h:3805
__hostdev__ Map(double s, const Vec3d &t=Vec3d(0.0, 0.0, 0.0))
Definition: NanoVDB.h:1406
__hostdev__ ChildNodeType * probeChild(const CoordType &ijk)
Definition: NanoVDB.h:3590
typename ChildT::CoordType CoordType
Definition: NanoVDB.h:3317
__hostdev__ void setLongGridNameOn(bool on=true)
Definition: NanoVDB.h:1913
__hostdev__ Mask(const Mask &other)
Copy constructor.
Definition: NanoVDB.h:1152
static __hostdev__ uint32_t CoordToOffset(const CoordT &ijk)
Return the linear offset corresponding to the given coordinate.
Definition: NanoVDB.h:4632
__hostdev__ uint64_t lastOffset() const
Definition: NanoVDB.h:4248
MaskT< LOG2DIM > mMask
Definition: NanoVDB.h:4272
__hostdev__ const BlindDataT * getBlindData(uint32_t n) const
Definition: NanoVDB.h:2238
#define NANOVDB_MAGIC_NUMB
Definition: NanoVDB.h:134
__hostdev__ void setWord(WordT w, uint32_t n)
Definition: NanoVDB.h:1170
GridClass
Classes (superset of OpenVDB) that are currently supported by NanoVDB.
Definition: NanoVDB.h:290
typename DataType::ValueT ValueType
Definition: NanoVDB.h:2679
__hostdev__ bool isPartial() const
return true if the 64 bit checksum is partial, i.e. of head only
Definition: NanoVDB.h:1786
static T scalar(const T &s)
Definition: NanoVDB.h:738
typename RootT::BuildType BuildType
Definition: NanoVDB.h:2356
__hostdev__ void setDev(const FloatType &)
Definition: NanoVDB.h:4161
Definition: NanoVDB.h:2775
__hostdev__ void * treePtr()
Definition: NanoVDB.h:1946
uint32_t state
Definition: NanoVDB.h:2586
BuildT BuildType
Definition: NanoVDB.h:3778
CoordT mBBoxMin
Definition: NanoVDB.h:4138
__hostdev__ void setDev(const FloatType &v)
Definition: NanoVDB.h:3822
typename UpperNodeType::ChildNodeType LowerNodeType
Definition: NanoVDB.h:2353
Return the pointer to the leaf node that contains math::Coord. Implements Tree::probeLeaf(math::Coord...
Definition: NanoVDB.h:1705
__hostdev__ bool getDev() const
Definition: NanoVDB.h:4155
#define NANOVDB_MAGIC_FRAG
Definition: NanoVDB.h:138
__hostdev__ bool isValid(GridType gridType, GridClass gridClass)
return true if the combination of GridType and GridClass is valid.
Definition: NanoVDB.h:613
static __hostdev__ bool isAligned(const void *p)
return true if the specified pointer is 32 byte aligned
Definition: NanoVDB.h:547
__hostdev__ void * getRoot()
Get a non-const void pointer to the root node (never NULL)
Definition: NanoVDB.h:2303
__hostdev__ ChildIterator beginChild()
Definition: NanoVDB.h:3374
uint8_t mFlags
Definition: NanoVDB.h:3785
__hostdev__ LeafNodeType * getFirstLeaf()
Template specializations of getFirstNode.
Definition: NanoVDB.h:2477
uint64_t mOffset
Definition: NanoVDB.h:4307
__hostdev__ void setOrigin(const T &ijk)
Definition: NanoVDB.h:3825
__hostdev__ ValueIter(RootT *parent)
Definition: NanoVDB.h:2784
static int64_t PtrDiff(const void *p, const void *q)
Compute the distance, in bytes, between two pointers, dist = p - q.
Definition: Util.h:464
__hostdev__ uint32_t gridIndex() const
Return index of this grid in the buffer.
Definition: NanoVDB.h:2080
__hostdev__ const RootT & root() const
Definition: NanoVDB.h:2380
__hostdev__ bool isEmpty() const
Return true if the root is empty, i.e. has not child nodes or constant tiles.
Definition: NanoVDB.h:2312
Definition: NanoVDB.h:2035
__hostdev__ const StatsT & stdDeviation() const
Definition: NanoVDB.h:2651
LeafNodeType Node0
Definition: NanoVDB.h:2363
__hostdev__ ValueType getValue(const CoordType &ijk) const
Return the value of the given voxel.
Definition: NanoVDB.h:2972
Checksum mChecksum
Definition: NanoVDB.h:1848
Return point to the upper internal node where math::Coord maps to one of its values, i.e. terminates.
Definition: NanoVDB.h:6578
__hostdev__ const GridClass & gridClass() const
Definition: NanoVDB.h:2175
__hostdev__ uint64_t leafPoints(const Coord &ijk, const AttT *&begin, const AttT *&end) const
Return the number of points in the leaf node containing the coordinate ijk. If this return value is l...
Definition: NanoVDB.h:6045
typename DataType::StatsT FloatType
Definition: NanoVDB.h:2680
__hostdev__ FloatType variance() const
Return the variance of all the active values encoded in this root node and any of its child nodes...
Definition: NanoVDB.h:2956
Implements Tree::getValue(math::Coord), i.e. return the value associated with a specific coordinate i...
Definition: NanoVDB.h:1695
BitFlags()
Definition: NanoVDB.h:932
__hostdev__ FloatType getAvg() const
Definition: NanoVDB.h:3816
__hostdev__ ChildIterator beginChild()
Definition: NanoVDB.h:2771
__hostdev__ bool isActive(const CoordT &ijk) const
Return true if the voxel value at the given coordinate is active.
Definition: NanoVDB.h:4608
__hostdev__ ConstDenseIterator cbeginDense() const
Definition: NanoVDB.h:2918
__hostdev__ uint64_t getValue(uint32_t i) const
Definition: NanoVDB.h:4253
ChildT ChildNodeType
Definition: NanoVDB.h:2673
#define NANOVDB_MAGIC_GRID
Definition: NanoVDB.h:135
__hostdev__ void setAvg(const FloatType &)
Definition: NanoVDB.h:4160
__hostdev__ ValueOnIterator beginValueOn() const
Definition: NanoVDB.h:3447
__hostdev__ const MaskType< LOG2DIM > & valueMask() const
Return a const reference to the bit mask of active voxels in this leaf node.
Definition: NanoVDB.h:4513
void set(const MatT &mat, const MatT &invMat, const Vec3T &translate, double taper=1.0)
Initialize the member data from 3x3 or 4x4 matrices.
Definition: NanoVDB.h:1522
static __hostdev__ KeyT CoordToKey(const CoordType &ijk)
Definition: NanoVDB.h:2526
__hostdev__ void setAvg(float avg)
Definition: NanoVDB.h:3899
MaskT< LOG2DIM > ArrayType
Definition: NanoVDB.h:4085
T Type
Definition: NanoVDB.h:463
__hostdev__ bool isActive() const
Definition: NanoVDB.h:3406
uint64_t mMagic
Definition: NanoVDB.h:1847
__hostdev__ ChannelT * setChannel(uint32_t channelID)
Change to an internal channel, assuming it exists as as blind data in the IndexGrid.
Definition: NanoVDB.h:6200
__hostdev__ void setMax(const ValueType &)
Definition: NanoVDB.h:4159
__hostdev__ bool isOff() const
Return true if none of the bits are set in this Mask.
Definition: NanoVDB.h:1220
__hostdev__ uint32_t valueCount() const
Definition: NanoVDB.h:4244
uint64_t mGridSize
Definition: NanoVDB.h:1853
__hostdev__ NodeT * probeChild(ValueType &value) const
Definition: NanoVDB.h:2883
RootT Node3
Definition: NanoVDB.h:2360
PointType
Definition: NanoVDB.h:401
__hostdev__ void toggle(uint32_t n)
Definition: NanoVDB.h:1303
Trait to map from LEVEL to node type.
Definition: NanoVDB.h:4788
__hostdev__ void setDev(const FloatType &)
Definition: NanoVDB.h:4202
__hostdev__ void setMax(const ValueT &v)
Definition: NanoVDB.h:2654
__hostdev__ void setOn(uint32_t offset)
Definition: NanoVDB.h:4203
__hostdev__ bool isFogVolume() const
Definition: NanoVDB.h:2177
__hostdev__ ValueIter()
Definition: NanoVDB.h:2780
#define NANOVDB_MINOR_VERSION_NUMBER
Definition: NanoVDB.h:145
__hostdev__ ReadAccessor(const TreeT &tree)
Constructor from a tree.
Definition: NanoVDB.h:5148
__hostdev__ WordT getWord(uint32_t n) const
Definition: NanoVDB.h:1163
__hostdev__ FloatType variance() const
Return the variance of all the active values encoded in this leaf node.
Definition: NanoVDB.h:4526
uint64_t KeyT
Return a key based on the coordinates of a voxel.
Definition: NanoVDB.h:2524
Vec3d mVoxelSize
Definition: NanoVDB.h:1857
BuildT ValueType
Definition: NanoVDB.h:3777
uint64_t mFlags
Definition: NanoVDB.h:3219
__hostdev__ const uint32_t & getTableSize() const
Definition: NanoVDB.h:2944
__hostdev__ ValueIterator()
Definition: NanoVDB.h:3384
__hostdev__ Checksum(uint32_t head, uint32_t tail)
Constructor that allows the two 32bit checksums to be initiated explicitly.
Definition: NanoVDB.h:1759
__hostdev__ uint32_t pos() const
Definition: NanoVDB.h:1111
__hostdev__ Mask()
Initialize all bits to zero.
Definition: NanoVDB.h:1139
__hostdev__ bool isCached2(const CoordType &ijk) const
Definition: NanoVDB.h:5392
__hostdev__ bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:3557
Implements Tree::getNodeInfo(math::Coord)
Definition: NanoVDB.h:1709
__hostdev__ uint64_t getMax() const
Definition: NanoVDB.h:4250
__hostdev__ void setStdDeviationOn(bool on=true)
Definition: NanoVDB.h:1915
__hostdev__ uint32_t gridCount() const
Return total number of grids in the buffer.
Definition: NanoVDB.h:2083
__hostdev__ bool isRootConnected() const
return true if RootData follows TreeData in memory without any extra padding
Definition: NanoVDB.h:2030
__hostdev__ uint64_t voxelPoints(const Coord &ijk, const AttT *&begin, const AttT *&end) const
get iterators over attributes to points at a specific voxel location
Definition: NanoVDB.h:6057
__hostdev__ ValueType operator()(int i, int j, int k) const
Definition: NanoVDB.h:5028
__hostdev__ bool isValid() const
Methods related to the classification of this grid.
Definition: NanoVDB.h:2173
__hostdev__ void setValue(uint32_t n, const ValueT &v)
Definition: NanoVDB.h:3247
ValueType minimum
Definition: NanoVDB.h:6630
__hostdev__ AccessorType getAccessor() const
Definition: NanoVDB.h:2927
ChildT UpperNodeType
Definition: NanoVDB.h:2676
uint32_t mGridCount
Definition: NanoVDB.h:1852
CoordT mBBoxMin
Definition: NanoVDB.h:3783
__hostdev__ bool isActive() const
Definition: NanoVDB.h:4477
uint64_t FloatType
Definition: NanoVDB.h:781
__hostdev__ const void * nodePtr() const
Return a non-const void pointer to the first node at LEVEL.
Definition: NanoVDB.h:1954
__hostdev__ FloatType getDev() const
Definition: NanoVDB.h:3817
__hostdev__ FloatType stdDeviation() const
Return a const reference to the standard deviation of all the active values encoded in this leaf node...
Definition: NanoVDB.h:4529
Definition: NanoVDB.h:6627
__hostdev__ bool isValid() const
return true if the magic number and the version are both valid
Definition: NanoVDB.h:1898
typename GridOrTreeOrRootT::RootNodeType::ChildNodeType::ChildNodeType type
Definition: NanoVDB.h:1653
uint64_t FloatType
Definition: NanoVDB.h:775
__hostdev__ RootT & root()
Definition: NanoVDB.h:2378
float mQuantum
Definition: NanoVDB.h:3857
char * strcpy(char *dst, const char *src)
Copy characters from src to dst.
Definition: Util.h:166
double FloatType
Definition: NanoVDB.h:805
__hostdev__ bool isIndex(GridType gridType)
Return true if the GridType maps to a special index type (not a POD integer type).
Definition: NanoVDB.h:602
Map mMap
Definition: NanoVDB.h:1855
#define NANOVDB_MAGIC_FILE
Definition: NanoVDB.h:136
__hostdev__ ValueType minimum() const
Return a const reference to the minimum active value encoded in this leaf node.
Definition: NanoVDB.h:4517
float type
Definition: NanoVDB.h:520
__hostdev__ const uint32_t & tileCount() const
Return the number of tiles encoded in this root node.
Definition: NanoVDB.h:2943
__hostdev__ Vec3T applyJacobian(const Vec3T &ijk) const
Apply the linear forward 3x3 transformation to an input 3d vector using 64bit floating point arithmet...
Definition: NanoVDB.h:1455
typename ChildT::LeafNodeType LeafNodeType
Definition: NanoVDB.h:2678
Bit-compacted representation of all three version numbers.
Definition: NanoVDB.h:678
__hostdev__ uint64_t lastOffset() const
Definition: NanoVDB.h:4227
__hostdev__ util::enable_if< BuildTraits< T >::is_index, const uint64_t & >::type valueCount() const
Return the total number of values indexed by this IndexGrid.
Definition: NanoVDB.h:2090
__hostdev__ ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:5022
__hostdev__ ChannelAccessor(const NanoGrid< IndexT > &grid, ChannelT *channelPtr)
Ctor from an IndexGrid and an external channel.
Definition: NanoVDB.h:6168
__hostdev__ bool operator>=(const Version &rhs) const
Definition: NanoVDB.h:704
typename DataType::ValueT ValueType
Definition: NanoVDB.h:3312
typename GridOrTreeOrRootT::LeafNodeType Type
Definition: NanoVDB.h:1637
static __hostdev__ uint32_t dim()
Return the dimension, in index space, of this leaf node (typically 8 as for openvdb leaf nodes!) ...
Definition: NanoVDB.h:4555
__hostdev__ ReadAccessor(const TreeT &tree)
Constructor from a tree.
Definition: NanoVDB.h:5634
Definition: NanoVDB.h:3207
static __hostdev__ uint32_t dim()
Return the dimension, in voxel units, of this internal node (typically 8*16 or 8*16*32) ...
Definition: NanoVDB.h:3506
__hostdev__ bool operator==(const Mask &other) const
Definition: NanoVDB.h:1193
__hostdev__ ChildIter & operator++()
Definition: NanoVDB.h:2752
__hostdev__ void setDev(const FloatType &)
Definition: NanoVDB.h:4342
__hostdev__ void setOn(uint32_t offset)
Definition: NanoVDB.h:4157
__hostdev__ const ChildT * probeChild(ValueType &value) const
Definition: NanoVDB.h:3468
__hostdev__ bool operator==(const Checksum &rhs) const
return true if the checksums are identical
Definition: NanoVDB.h:1806
__hostdev__ const NanoGrid< BuildT > & grid() const
Definition: NanoVDB.h:6031
char * strncpy(char *dst, const char *src, size_t max)
Copies the first num characters of src to dst. If the end of the source C string (which is signaled b...
Definition: Util.h:185
__hostdev__ DenseIterator beginAll() const
Definition: NanoVDB.h:1136
__hostdev__ ConstValueIterator cbeginValueAll() const
Definition: NanoVDB.h:2821
__hostdev__ void disable()
Definition: NanoVDB.h:1795
__hostdev__ const NanoGrid< IndexT > & grid() const
Return a const reference to the IndexGrid.
Definition: NanoVDB.h:6181
static constexpr uint32_t SIZE
Definition: NanoVDB.h:1037
uint32_t mNodeCount[3]
Definition: NanoVDB.h:2291
ValueType mMaximum
Definition: NanoVDB.h:3789
typename GridOrTreeOrRootT::LeafNodeType type
Definition: NanoVDB.h:1638
MaskT< LOG2DIM > mValueMask
Definition: NanoVDB.h:4141
static __hostdev__ CoordT KeyToCoord(const KeyT &key)
Definition: NanoVDB.h:2534
__hostdev__ const Map & map() const
Return a const reference to the Map for this grid.
Definition: NanoVDB.h:2112
__hostdev__ ValueIterator cbeginValueAll() const
Definition: NanoVDB.h:4497
__hostdev__ void setRoot(const void *root)
Definition: NanoVDB.h:2297
static __hostdev__ uint32_t wordCount()
Return the number of machine words used by this Mask.
Definition: NanoVDB.h:1047
__hostdev__ uint32_t operator*() const
Definition: NanoVDB.h:1110
__hostdev__ void setOrigin(const T &ijk)
Definition: NanoVDB.h:4115
typename DataType::BuildT BuildType
Definition: NanoVDB.h:2681
__hostdev__ void setMin(const bool &)
Definition: NanoVDB.h:4109
__hostdev__ const StatsT & average() const
Definition: NanoVDB.h:2650
__hostdev__ void setMin(const ValueType &)
Definition: NanoVDB.h:4199
__hostdev__ uint32_t tail() const
Definition: NanoVDB.h:1781
__hostdev__ bool getAvg() const
Definition: NanoVDB.h:4101
__hostdev__ bool updateBBox()
Updates the local bounding box of active voxels in this node. Return true if bbox was updated...
Definition: NanoVDB.h:4735
__hostdev__ DenseIter & operator++()
Definition: NanoVDB.h:2900
__hostdev__ bool isBreadthFirst() const
Definition: NanoVDB.h:2189
__hostdev__ bool isPointData() const
Definition: NanoVDB.h:2181
__hostdev__ uint64_t last(uint32_t i) const
Definition: NanoVDB.h:4324
bool FloatType
Definition: NanoVDB.h:769
__hostdev__ const FloatType & average() const
Return a const reference to the average of all the active values encoded in this internal node and an...
Definition: NanoVDB.h:3529
__hostdev__ Iterator & operator++()
Definition: NanoVDB.h:1085
Definition: NanoVDB.h:755
#define __hostdev__
Definition: Util.h:73
typename DataType::FloatType FloatType
Definition: NanoVDB.h:4373
#define NANOVDB_DATA_ALIGNMENT
Definition: NanoVDB.h:126
typename DataType::Tile Tile
Definition: NanoVDB.h:2686
__hostdev__ bool isValid(const GridBlindDataClass &blindClass, const GridBlindDataSemantic &blindSemantics, const GridType &blindType)
return true if the combination of GridBlindDataClass, GridBlindDataSemantic and GridType is valid...
Definition: NanoVDB.h:637
__hostdev__ DenseIterator operator++(int)
Definition: NanoVDB.h:1118
__hostdev__ uint64_t getValue(uint32_t i) const
Definition: NanoVDB.h:4233
__hostdev__ void setMax(const ValueT &v)
Definition: NanoVDB.h:3292
Definition: NanoVDB.h:925
Coord CoordType
Definition: NanoVDB.h:4375
Dummy type for a 16bit quantization of float point values.
Definition: NanoVDB.h:197
uint8_t ArrayType
Definition: NanoVDB.h:3956
typename Mask< Log2Dim >::template Iterator< On > MaskIterT
Definition: NanoVDB.h:3322
__hostdev__ bool hasLongGridName() const
Definition: NanoVDB.h:2186
__hostdev__ TreeT & tree()
Return a non-const reference to the tree.
Definition: NanoVDB.h:2103
CoordT mBBoxMin
Definition: NanoVDB.h:4302
__hostdev__ void setFirstNode(const NodeT *node)
Definition: NanoVDB.h:2309
const typename GridOrTreeOrRootT::RootNodeType::ChildNodeType Type
Definition: NanoVDB.h:1673
__hostdev__ const ValueType & maximum() const
Return a const reference to the maximum active value encoded in this internal node and any of its chi...
Definition: NanoVDB.h:3526
__hostdev__ float getValue(uint32_t i) const
Definition: NanoVDB.h:3931
static __hostdev__ uint64_t memUsage()
Definition: NanoVDB.h:3238
__hostdev__ void setMin(float min)
Definition: NanoVDB.h:3893
__hostdev__ void setValue(uint32_t offset, bool)
Definition: NanoVDB.h:4156
__hostdev__ void setMax(const ValueType &v)
Definition: NanoVDB.h:3820
__hostdev__ void setOrigin(const T &ijk)
Definition: NanoVDB.h:4164
static __hostdev__ constexpr uint32_t padding()
Return padding of this class in bytes, due to aliasing and 32B alignment.
Definition: NanoVDB.h:3797
__hostdev__ uint32_t pos() const
Definition: NanoVDB.h:2708
__hostdev__ ChildIter()
Definition: NanoVDB.h:3342
__hostdev__ void setMin(const ValueType &)
Definition: NanoVDB.h:4339
__hostdev__ BlindDataT * getBlindData(uint32_t n)
Definition: NanoVDB.h:2245
__hostdev__ void setDev(const StatsT &v)
Definition: NanoVDB.h:3294
__hostdev__ ValueType getLastValue() const
If the last entry in this node's table is a tile, return the tile's value. Otherwise, return the result of calling getLastValue() on the child.
Definition: NanoVDB.h:3549
__hostdev__ void clear()
Reset this access to its initial state, i.e. with an empty cache.
Definition: NanoVDB.h:5342
__hostdev__ bool isOn(uint32_t n) const
Return true if the given bit is set.
Definition: NanoVDB.h:1205
__hostdev__ float getValue(uint32_t i) const
Definition: NanoVDB.h:4029
__hostdev__ Vec3T applyInverseJacobian(const Vec3T &xyz) const
Apply the linear inverse 3x3 transformation to an input 3d vector using 64bit floating point arithmet...
Definition: NanoVDB.h:1495
uint64_t ValueType
Definition: NanoVDB.h:4181
uint16_t ArrayType
Definition: NanoVDB.h:3986
__hostdev__ const MaskType< LOG2DIM > & childMask() const
Return a const reference to the bit mask of child nodes in this internal node.
Definition: NanoVDB.h:3516
__hostdev__ void setMax(const ValueType &)
Definition: NanoVDB.h:4200
__hostdev__ void setAvg(const FloatType &)
Definition: NanoVDB.h:4201
ValueT value
Definition: NanoVDB.h:2587
__hostdev__ void setDev(float dev)
Definition: NanoVDB.h:3902
Node caching at all (three) tree levels.
Definition: NanoVDB.h:5583
__hostdev__ void setDev(const StatsT &v)
Definition: NanoVDB.h:2656
__hostdev__ OnIterator beginOn() const
Definition: NanoVDB.h:1132
Definition: NanoVDB.h:1697
__hostdev__ void setAvg(const FloatType &)
Definition: NanoVDB.h:4341
__hostdev__ bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:5712
__hostdev__ void setOn(uint32_t offset)
Definition: NanoVDB.h:4332
typename GridOrTreeOrRootT::RootNodeType::ChildNodeType::ChildNodeType Type
Definition: NanoVDB.h:1652
__hostdev__ bool isMaskOn(uint32_t offset) const
Definition: NanoVDB.h:4285
BuildT BuildType
Definition: NanoVDB.h:5607
Stuct with all the member data of the LeafNode (useful during serialization of an openvdb LeafNode) ...
Definition: NanoVDB.h:3773
__hostdev__ const BBoxType & bbox() const
Return a const reference to the index bounding box of all the active values in this tree...
Definition: NanoVDB.h:2934
GridBlindDataSemantic
Blind-data Semantics that are currently understood by NanoVDB.
Definition: NanoVDB.h:424
Version mVersion
Definition: NanoVDB.h:1849
__hostdev__ void setAverageOn(bool on=true)
Definition: NanoVDB.h:1914
__hostdev__ bool isSequential() const
return true if nodes at all levels can safely be accessed with simple linear offsets ...
Definition: NanoVDB.h:2202
__hostdev__ Map()
Default constructor for the identity map.
Definition: NanoVDB.h:1395
GridFlags
Grid flags which indicate what extra information is present in the grid buffer.
Definition: NanoVDB.h:327
Metafunction used to determine if the first template parameter is a specialization of the class templ...
Definition: Util.h:451
static __hostdev__ constexpr uint8_t bitWidth()
Definition: NanoVDB.h:3997
__hostdev__ uint32_t & checksum(int i)
Definition: NanoVDB.h:1773
__hostdev__ DenseIterator()
Definition: NanoVDB.h:3457
ReadAccessor< ValueT, LEVEL0, LEVEL1, LEVEL2 > createAccessor(const NanoGrid< ValueT > &grid)
Free-standing function for convenient creation of a ReadAccessor with optional and customizable node ...
Definition: NanoVDB.h:5905
RootT RootType
Definition: NanoVDB.h:2350
static __hostdev__ constexpr uint32_t padding()
Return padding of this class in bytes, due to aliasing and 32B alignment.
Definition: NanoVDB.h:3867
Definition: GridHandle.h:27
float type
Definition: NanoVDB.h:513
CoordBBox bbox
Definition: NanoVDB.h:6632
float Type
Definition: NanoVDB.h:519
__hostdev__ bool probeValue(const CoordType &ijk, ValueType &v) const
Return true if this tree is empty, i.e. contains no values or nodes.
Definition: NanoVDB.h:2395
Visits all tile values and child nodes of this node.
Definition: NanoVDB.h:3451
GridType mGridType
Definition: NanoVDB.h:1859
static __hostdev__ uint64_t memUsage()
Definition: NanoVDB.h:4284
__hostdev__ void clear()
Reset this access to its initial state, i.e. with an empty cache.
Definition: NanoVDB.h:5154
Definition: NanoVDB.h:1068
__hostdev__ bool probeValue(const CoordType &ijk, ValueType &v) const
return the state and updates the value of the specified voxel
Definition: NanoVDB.h:3559
static __hostdev__ constexpr uint32_t padding()
Definition: NanoVDB.h:4191
Define static boolean tests for template build types.
Definition: NanoVDB.h:440
__hostdev__ bool isFull() const
return true if the 64 bit checksum is fill, i.e. of both had and nodes
Definition: NanoVDB.h:1790
__hostdev__ bool hasMinMax() const
Definition: NanoVDB.h:2184
__hostdev__ ConstChildIterator cbeginChild() const
Definition: NanoVDB.h:2772
char * sprint(char *dst, T var1, Types...var2)
prints a variable number of string and/or numbers to a destination string
Definition: Util.h:286
Bit-mask to encode active states and facilitate sequential iterators and a fast codec for I/O compres...
Definition: NanoVDB.h:1034
CoordT CoordType
Definition: NanoVDB.h:5127
__hostdev__ const GridBlindMetaData * blindMetaData(uint32_t n) const
Returns a const reference to the blindMetaData at the specified linear offset.
Definition: NanoVDB.h:1986
__hostdev__ ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:5175
__hostdev__ void setOrigin(const T &ijk)
Definition: NanoVDB.h:3280
static ElementType scalar(const T &v)
Definition: NanoVDB.h:749
__hostdev__ ValueIterator beginValue() const
Definition: NanoVDB.h:3413
__hostdev__ void setMax(float max)
Definition: NanoVDB.h:3896
static __hostdev__ constexpr uint8_t bitWidth()
Definition: NanoVDB.h:3930
__hostdev__ bool getMax() const
Definition: NanoVDB.h:4100
uint64_t mData2
Definition: NanoVDB.h:1864
typename ChildT::ValueType ValueT
Definition: NanoVDB.h:2516
float mMinimum
Definition: NanoVDB.h:3856
__hostdev__ uint64_t offset() const
Definition: NanoVDB.h:4321
static __hostdev__ Coord OffsetToLocalCoord(uint32_t n)
Definition: NanoVDB.h:3610
__hostdev__ const Vec3d & voxelSize() const
Return a vector of the axial voxel sizes.
Definition: NanoVDB.h:6187
__hostdev__ constexpr uint32_t strlen()
return the number of characters (including null termination) required to convert enum type to a strin...
Definition: NanoVDB.h:209
typename NanoLeaf< BuildT >::FloatType FloatType
Definition: NanoVDB.h:6626
Definition: NanoVDB.h:4177
KeyT key
Definition: NanoVDB.h:2584
uint64_t FloatType
Definition: NanoVDB.h:787
__hostdev__ uint64_t pointCount() const
Definition: NanoVDB.h:4322
typename DataType::StatsT FloatType
Definition: NanoVDB.h:3313
__hostdev__ ValueType getValue(int i, int j, int k) const
Definition: NanoVDB.h:5405
__hostdev__ uint32_t & tail()
Definition: NanoVDB.h:1782
bool ValueType
Definition: NanoVDB.h:4082
__hostdev__ Tile * probeTile(const CoordT &ijk)
Definition: NanoVDB.h:2604
__hostdev__ uint64_t getAvg() const
Definition: NanoVDB.h:4231
__hostdev__ uint32_t getDim(const CoordType &ijk, const RayT &ray) const
Definition: NanoVDB.h:5072
__hostdev__ uint64_t checksum() const
return the 64 bit checksum of this instance
Definition: NanoVDB.h:1771
Dummy type for a voxel whose value equals an offset into an external value array of active values...
Definition: NanoVDB.h:176
__hostdev__ ValueOnIterator beginValueOn() const
Definition: NanoVDB.h:4412
Top-most node of the VDB tree structure.
Definition: NanoVDB.h:2669
int64_t child
Definition: NanoVDB.h:3210
#define NANOVDB_MAJOR_VERSION_NUMBER
Definition: NanoVDB.h:144
__hostdev__ Vec3T applyJacobianF(const Vec3T &ijk) const
Apply the linear forward 3x3 transformation to an input 3d vector using 32bit floating point arithmet...
Definition: NanoVDB.h:1464
uint8_t ArrayType
Definition: NanoVDB.h:3919
static __hostdev__ uint64_t memUsage()
Definition: NanoVDB.h:3801
Struct to derive node type from its level in a given grid, tree or root while preserving constness...
Definition: NanoVDB.h:1630
typename GridT::TreeType type
Definition: NanoVDB.h:2328
__hostdev__ Codec toCodec(const char *str)
Definition: NanoVDB.h:6274
Definition: NanoVDB.h:2724
uint32_t level
Definition: NanoVDB.h:6629
__hostdev__ float getValue(uint32_t i) const
Definition: NanoVDB.h:3967
uint32_t mTileCount[3]
Definition: NanoVDB.h:2292
typename RootT::ChildNodeType Node2
Definition: NanoVDB.h:2361
__hostdev__ CoordType getOrigin() const
Definition: NanoVDB.h:3439
__hostdev__ const ValueType & minimum() const
Return a const reference to the minimum active value encoded in this internal node and any of its chi...
Definition: NanoVDB.h:3523
ValueType mMinimum
Definition: NanoVDB.h:3788
__hostdev__ const void * blindData(uint32_t n) const
Returns a const pointer to the blindData at the specified linear offset.
Definition: NanoVDB.h:2230
__hostdev__ GridType toGridType()
Maps from a templated build type to a GridType enum.
Definition: NanoVDB.h:812
size_t strlen(const char *str)
length of a c-sting, excluding '\0'.
Definition: Util.h:153
static __hostdev__ uint32_t dim()
Definition: NanoVDB.h:4368
__hostdev__ bool isCached(const CoordType &ijk) const
Definition: NanoVDB.h:5695
uint64_t Type
Definition: NanoVDB.h:470
uint64_t type
Definition: NanoVDB.h:478
Definition: IndexIterator.h:43
__hostdev__ float getDev() const
return the quantized standard deviation of the active values in this node
Definition: NanoVDB.h:3890
__hostdev__ ValueType operator*() const
Definition: NanoVDB.h:3434
ValueT mMaximum
Definition: NanoVDB.h:3224
__hostdev__ uint64_t idx(int i, int j, int k) const
Definition: NanoVDB.h:6207
static __hostdev__ CoordT OffsetToLocalCoord(uint32_t n)
Compute the local coordinates from a linear offset.
Definition: NanoVDB.h:4539
__hostdev__ const math::BBox< CoordType > & bbox() const
Return a const reference to the bounding box in index space of active values in this internal node an...
Definition: NanoVDB.h:3538
__hostdev__ ValueType getValue(const CoordType &ijk) const
Return the value of the given voxel.
Definition: NanoVDB.h:3556
__hostdev__ const FloatType & stdDeviation() const
Return a const reference to the standard deviation of all the active values encoded in this root node...
Definition: NanoVDB.h:2959
__hostdev__ uint64_t getValue(uint32_t i) const
Definition: NanoVDB.h:4325
__hostdev__ bool operator<=(const Version &rhs) const
Definition: NanoVDB.h:702
__hostdev__ bool getValue(uint32_t i) const
Definition: NanoVDB.h:4098
T ElementType
Definition: NanoVDB.h:737
typename RootType::LeafNodeType LeafNodeType
Definition: NanoVDB.h:2053
float Type
Definition: NanoVDB.h:533
__hostdev__ uint32_t getDim(const CoordType &ijk, const RayT &ray) const
Definition: NanoVDB.h:5235
uint64_t Type
Definition: NanoVDB.h:477
__hostdev__ uint32_t getMinor() const
Definition: NanoVDB.h:707
Struct with all the member data of the RootNode (useful during serialization of an openvdb RootNode) ...
Definition: NanoVDB.h:2514
__hostdev__ CoordType getCoord() const
Definition: NanoVDB.h:3368
__hostdev__ ValueIterator operator++(int)
Definition: NanoVDB.h:4488
__hostdev__ int findBlindDataForSemantic(GridBlindDataSemantic semantic) const
Return the index of the first blind data with specified semantic if found, otherwise -1...
Definition: NanoVDB.h:2258
static __hostdev__ bool hasStats()
Definition: NanoVDB.h:3862
__hostdev__ ValueOffIterator(const LeafNode *parent)
Definition: NanoVDB.h:4427
openvdb::GridBase Grid
Definition: Utils.h:34
__hostdev__ Mask(bool on)
Definition: NanoVDB.h:1144
__hostdev__ void setOff(uint32_t n)
Set the specified bit off.
Definition: NanoVDB.h:1231
__hostdev__ const char * gridName() const
Return a c-string with the name of this grid.
Definition: NanoVDB.h:2205
__hostdev__ ValueType operator*() const
Definition: NanoVDB.h:4433
typename RootNodeType::ChildNodeType UpperNodeType
Definition: NanoVDB.h:2352
double FloatType
Definition: NanoVDB.h:763
uint64_t type
Definition: NanoVDB.h:485
__hostdev__ ChildT * getChild(const Tile *tile)
Returns a const reference to the child node in the specified tile.
Definition: NanoVDB.h:2637
__hostdev__ const Checksum & checksum() const
Return checksum of the grid buffer.
Definition: NanoVDB.h:2211
__hostdev__ ReadAccessor(const GridT &grid)
Constructor from a grid.
Definition: NanoVDB.h:5330
GridClass mGridClass
Definition: NanoVDB.h:1858
__hostdev__ Version(uint32_t data)
Constructor from a raw uint32_t data representation.
Definition: NanoVDB.h:691
Dummy type for a voxel whose value equals an offset into an external value array. ...
Definition: NanoVDB.h:173
Maps one type (e.g. the build types above) to other (actual) types.
Definition: NanoVDB.h:461
__hostdev__ const DataType * data() const
Definition: NanoVDB.h:2373
__hostdev__ ValueType getLastValue() const
Return the last value in this leaf node.
Definition: NanoVDB.h:4594
__hostdev__ void setOrigin(const T &ijk)
Definition: NanoVDB.h:3905
__hostdev__ bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:5184
__hostdev__ bool isHalf() const
Definition: NanoVDB.h:1787
__hostdev__ uint64_t getDev() const
Definition: NanoVDB.h:4232
static __hostdev__ uint64_t memUsage()
Definition: NanoVDB.h:4096
math::BBox< CoordT > mBBox
Definition: NanoVDB.h:2546
__hostdev__ ValueIterator(const LeafNode *parent)
Definition: NanoVDB.h:4460
__hostdev__ ReadAccessor(const RootT &root)
Constructor from a root node.
Definition: NanoVDB.h:4994
typename RootT::ValueType ValueType
Definition: NanoVDB.h:4976
__hostdev__ uint32_t id() const
Definition: NanoVDB.h:705
__hostdev__ size_t memUsage() const
Definition: NanoVDB.h:4027
__hostdev__ ValueType getMin() const
Definition: NanoVDB.h:4334
__hostdev__ const NodeT * getFirstNode() const
return a const pointer to the first node of the specified type
Definition: NanoVDB.h:2452
typename ChildT::CoordType CoordType
Definition: NanoVDB.h:2683
__hostdev__ uint32_t getPatch() const
Definition: NanoVDB.h:708
Definition: NanoVDB.h:2824
__hostdev__ DenseIter(RootT *parent)
Definition: NanoVDB.h:2878
__hostdev__ const FloatType & stdDeviation() const
Return a const reference to the standard deviation of all the active values encoded in this internal ...
Definition: NanoVDB.h:3535
__hostdev__ void setOn(uint32_t n)
Set the specified bit on.
Definition: NanoVDB.h:1229
__hostdev__ const uint64_t & firstOffset() const
Definition: NanoVDB.h:4198
__hostdev__ CoordT getCoord() const
Definition: NanoVDB.h:4405
__hostdev__ Vec3T applyInverseJacobianF(const Vec3T &xyz) const
Apply the linear inverse 3x3 transformation to an input 3d vector using 32bit floating point arithmet...
Definition: NanoVDB.h:1504
__hostdev__ bool isCompatible() const
Definition: NanoVDB.h:709
__hostdev__ ValueType getValue(int i, int j, int k) const
Definition: NanoVDB.h:5179
__hostdev__ ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:5401
static __hostdev__ constexpr uint8_t bitWidth()
Definition: NanoVDB.h:3966
__hostdev__ auto getNodeInfo(const CoordType &ijk) const
Definition: NanoVDB.h:5408
__hostdev__ const ValueType & background() const
Return a const reference to the background value.
Definition: NanoVDB.h:2398
const typename GridOrTreeOrRootT::LeafNodeType type
Definition: NanoVDB.h:1645
__hostdev__ int age() const
Returns the difference between major version of this instance and NANOVDB_MAJOR_VERSION_NUMBER.
Definition: NanoVDB.h:713
__hostdev__ bool isRootNext() const
return true if RootData is layout out immediately after TreeData in memory
Definition: NanoVDB.h:2318
__hostdev__ NodeT * getFirstNode()
return a pointer to the first node of the specified type
Definition: NanoVDB.h:2442
__hostdev__ const NodeTrait< RootT, 2 >::type * getFirstUpper() const
Definition: NanoVDB.h:2482
#define NANOVDB_MAGIC_NODE
Definition: NanoVDB.h:137
CheckMode
List of different modes for computing for a checksum.
Definition: NanoVDB.h:1714
__hostdev__ void setAvg(const FloatType &v)
Definition: NanoVDB.h:3821
__hostdev__ uint8_t bitWidth() const
Definition: NanoVDB.h:4026
__hostdev__ const FloatType & average() const
Return a const reference to the average of all the active values encoded in this root node and any of...
Definition: NanoVDB.h:2953
MatType scale(const Vec3< typename MatType::value_type > &s)
Return a matrix that scales by s.
Definition: Mat.h:615
__hostdev__ ValueIter operator++(int)
Definition: NanoVDB.h:2809
static __hostdev__ constexpr uint32_t padding()
Return padding of this class in bytes, due to aliasing and 32B alignment.
Definition: NanoVDB.h:3232
__hostdev__ ValueType operator*() const
Definition: NanoVDB.h:4467
__hostdev__ void extrema(ValueType &min, ValueType &max) const
Sets the extrema values of all the active values in this tree, i.e. in all nodes of the tree...
Definition: NanoVDB.h:2502
__hostdev__ uint32_t pos() const
Definition: NanoVDB.h:1083
__hostdev__ void setOn(uint32_t offset)
Definition: NanoVDB.h:3877
__hostdev__ T & getValue(const math::Coord &ijk, T *channelPtr) const
Return the value from a specified channel that maps to the specified coordinate.
Definition: NanoVDB.h:6226
typename Node2::ChildNodeType Node1
Definition: NanoVDB.h:2362
Dummy type for a 16 bit floating point values (placeholder for IEEE 754 Half)
Definition: NanoVDB.h:188
static __hostdev__ uint64_t memUsage()
return memory usage in bytes for the class
Definition: NanoVDB.h:2376
RootT RootNodeType
Definition: NanoVDB.h:2351
__hostdev__ Vec3T applyInverseMapF(const Vec3T &xyz) const
Definition: NanoVDB.h:1937
__hostdev__ uint64_t first(uint32_t i) const
Definition: NanoVDB.h:4323
__hostdev__ bool isMaskOn(uint32_t offset) const
Definition: NanoVDB.h:4274
__hostdev__ bool isGridIndex() const
Definition: NanoVDB.h:2180
__hostdev__ ReadAccessor(const TreeT &tree)
Constructor from a tree.
Definition: NanoVDB.h:5336
uint32_t countOn(uint64_t v)
Definition: Util.h:622
__hostdev__ ChannelAccessor(const NanoGrid< IndexT > &grid, uint32_t channelID=0u)
Ctor from an IndexGrid and an integer ID of an internal channel that is assumed to exist as blind dat...
Definition: NanoVDB.h:6157
__hostdev__ uint64_t gridPoints(const AttT *&begin, const AttT *&end) const
Return the total number of point in the grid and set the iterators to the complete range of points...
Definition: NanoVDB.h:6035
void ArrayType
Definition: NanoVDB.h:4183
__hostdev__ ChannelT & operator()(const math::Coord &ijk) const
Definition: NanoVDB.h:6211
__hostdev__ uint32_t countOn(uint32_t i) const
Return the number of lower set bits in mask up to but excluding the i'th bit.
Definition: NanoVDB.h:1059
__hostdev__ ChildIter()
Definition: NanoVDB.h:2731
__hostdev__ bool hasStats() const
Definition: NanoVDB.h:4196
__hostdev__ uint64_t memUsage() const
Return the actual memory footprint of this root node.
Definition: NanoVDB.h:2965
int64_t child
Definition: NanoVDB.h:2585
__hostdev__ void fill(const ValueType &v)
Definition: NanoVDB.h:3827
__hostdev__ bool getAvg() const
Definition: NanoVDB.h:4154
BuildT ArrayType
Definition: NanoVDB.h:3780
uint32_t mBlindMetadataCount
Definition: NanoVDB.h:1861
__hostdev__ OffIterator beginOff() const
Definition: NanoVDB.h:1134
__hostdev__ DenseIterator beginDense()
Definition: NanoVDB.h:2917
BuildT BuildType
Definition: NanoVDB.h:5125
__hostdev__ bool getMin() const
Definition: NanoVDB.h:4099
__hostdev__ const DataType * data() const
Definition: NanoVDB.h:4510
__hostdev__ const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:5032
__hostdev__ bool getMax() const
Definition: NanoVDB.h:4153
__hostdev__ ValueType operator*() const
Definition: NanoVDB.h:2840
bool BuildType
Definition: NanoVDB.h:4083
__hostdev__ CoordT origin() const
Definition: NanoVDB.h:2583
__hostdev__ bool operator<(const Version &rhs) const
Definition: NanoVDB.h:701
__hostdev__ bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:5409
__hostdev__ const Vec3dBBox & worldBBox() const
return AABB of active values in world space
Definition: NanoVDB.h:2012
__hostdev__ void setOn(uint32_t offset)
Definition: NanoVDB.h:4108
__hostdev__ uint8_t flags() const
Definition: NanoVDB.h:4531
__hostdev__ const ValueT & getMax() const
Definition: NanoVDB.h:3283
typename UpperNodeType::ChildNodeType LowerNodeType
Definition: NanoVDB.h:2052
__hostdev__ bool isEmpty() const
Return true if this RootNode is empty, i.e. contains no values or nodes.
Definition: NanoVDB.h:2968
VecT< GridHandleT > readUncompressedGrids(const char *fileName, const typename GridHandleT::BufferType &buffer=typename GridHandleT::BufferType())
Read a multiple un-compressed NanoVDB grids from a file and return them as a vector.
Definition: NanoVDB.h:6456
uint64_t Type
Definition: NanoVDB.h:540
CoordT mBBoxMin
Definition: NanoVDB.h:4186
__hostdev__ uint32_t checksum(int i) const
Definition: NanoVDB.h:1775
__hostdev__ bool operator!=(const Mask &other) const
Definition: NanoVDB.h:1202
CoordT CoordType
Definition: NanoVDB.h:5310
Dummy type for a variable bit quantization of floating point values.
Definition: NanoVDB.h:200
__hostdev__ Vec3T indexToWorldF(const Vec3T &xyz) const
index to world space transformation
Definition: NanoVDB.h:2143
__hostdev__ const MaskType< LOG2DIM > & getChildMask() const
Definition: NanoVDB.h:3517
StatsT mAverage
Definition: NanoVDB.h:2552
Visits all values in a leaf node, i.e. both active and inactive values.
Definition: NanoVDB.h:4449
__hostdev__ void setMin(const ValueT &v)
Definition: NanoVDB.h:3291
__hostdev__ bool hasAverage() const
Definition: NanoVDB.h:2187
__hostdev__ CoordType getCoord() const
Definition: NanoVDB.h:3405
__hostdev__ bool isActive() const
Definition: NanoVDB.h:2582
Visits active tile values of this node only.
Definition: NanoVDB.h:3417
__hostdev__ const NodeTrait< RootT, LEVEL >::type * getFirstNode() const
return a const pointer to the first node of the specified level
Definition: NanoVDB.h:2471
#define NANOVDB_HOSTDEV_DISABLE_WARNING
Definition: Util.h:94
__hostdev__ void setValueOnly(uint32_t offset, const ValueType &value)
Definition: NanoVDB.h:3806
Visits all inactive values in a leaf node.
Definition: NanoVDB.h:4416
__hostdev__ const TreeType & tree() const
Return a const reference to the tree of the IndexGrid.
Definition: NanoVDB.h:6184
Like ValueIndex but with a mutable mask.
Definition: NanoVDB.h:179
typename RootT::ValueType ValueType
Definition: NanoVDB.h:2355
static __hostdev__ uint64_t memUsage(uint32_t tableSize)
Return the expected memory footprint in bytes with the specified number of tiles. ...
Definition: NanoVDB.h:2962
Definition: NanoVDB.h:1699
__hostdev__ ValueIterator beginValue() const
Definition: NanoVDB.h:4496
static __hostdev__ bool hasStats()
Definition: NanoVDB.h:4145
float FloatType
Definition: NanoVDB.h:757
__hostdev__ CheckMode mode() const
return the mode of the 64 bit checksum
Definition: NanoVDB.h:1798
__hostdev__ bool isMask() const
Definition: NanoVDB.h:2182
__hostdev__ Vec3T applyJacobianF(const Vec3T &xyz) const
Definition: NanoVDB.h:1939
__hostdev__ auto getNodeInfo(const CoordType &ijk) const
Definition: NanoVDB.h:5182
__hostdev__ void setMax(const bool &)
Definition: NanoVDB.h:4110
typename RootNodeType::ChildNodeType UpperNodeType
Definition: NanoVDB.h:2051
typename ChildT::BuildType BuildT
Definition: NanoVDB.h:2517
typename BuildT::ValueType ValueType
Definition: NanoVDB.h:2055
__hostdev__ uint32_t nodeCount() const
Return number of nodes at LEVEL.
Definition: NanoVDB.h:1977
float ValueType
Definition: NanoVDB.h:3848
__hostdev__ Mask & operator&=(const Mask &other)
Bitwise intersection.
Definition: NanoVDB.h:1306
__hostdev__ uint32_t getDim(const CoordType &ijk, const RayT &ray) const
Definition: NanoVDB.h:5502
__hostdev__ const TreeT & tree() const
Return a const reference to the tree.
Definition: NanoVDB.h:2100
uint32_t findLowestOn(uint32_t v)
Returns the index of the lowest, i.e. least significant, on bit in the specified 32 bit word...
Definition: Util.h:502
__hostdev__ ChannelT & getValue(const math::Coord &ijk) const
Return the value from a cached channel that maps to the specified coordinate.
Definition: NanoVDB.h:6210
uint64_t mData1
Definition: NanoVDB.h:1863
BitFlags(Type mask)
Definition: NanoVDB.h:933
__hostdev__ ChildNodeType * probeChild(const CoordType &ijk)
Definition: NanoVDB.h:3030
__hostdev__ bool isValueOn() const
Definition: NanoVDB.h:3479
bool streq(const char *lhs, const char *rhs)
Test if two null-terminated byte strings are the same.
Definition: Util.h:268
__hostdev__ Vec3T worldToIndexDirF(const Vec3T &dir) const
transformation from world space direction to index space direction
Definition: NanoVDB.h:2153
__hostdev__ Iterator()
Definition: NanoVDB.h:1071
typename ChildT::template MaskType< LOG2DIM > MaskT
Definition: NanoVDB.h:3204
__hostdev__ uint64_t getDev() const
Definition: NanoVDB.h:4252
BitFlags< 32 > mFlags
Definition: NanoVDB.h:1850
__hostdev__ Vec3T applyInverseJacobianF(const Vec3T &xyz) const
Definition: NanoVDB.h:1941
__hostdev__ ValueOnIter operator++(int)
Definition: NanoVDB.h:2853
uint8_t mFlags
Definition: NanoVDB.h:4304
__hostdev__ void setAvg(const bool &)
Definition: NanoVDB.h:4111
__hostdev__ void setMin(const ValueType &v)
Definition: NanoVDB.h:3819
__hostdev__ bool getMin() const
Definition: NanoVDB.h:4152
__hostdev__ bool isStaggered() const
Definition: NanoVDB.h:2178
__hostdev__ ConstChildIterator cbeginChild() const
Definition: NanoVDB.h:3375
uint8_t mFlags
Definition: NanoVDB.h:4140
void writeUncompressedGrid(StreamT &os, const GridData *gridData, bool raw=false)
This is a standalone alternative to io::writeGrid(...,Codec::NONE) defined in util/IO.h Unlike the latter this function has no dependencies at all, not even NanoVDB.h, so it also works if client code only includes PNanoVDB.h!
Definition: NanoVDB.h:6352
__hostdev__ bool isEmpty() const
test if the grid is empty, e.i the root table has size 0
Definition: NanoVDB.h:2026
__hostdev__ uint64_t gridPoints(const AttT *&begin, const AttT *&end) const
Return the total number of point in the grid and set the iterators to the complete range of points...
Definition: NanoVDB.h:6102
__hostdev__ NodeT * operator->() const
Definition: NanoVDB.h:3358
MaskT< LOG2DIM > mValueMask
Definition: NanoVDB.h:3854
#define __device__
Definition: Util.h:79
__hostdev__ Vec3T applyInverseMap(const Vec3T &xyz) const
Definition: NanoVDB.h:1926
int64_t mBlindMetadataOffset
Definition: NanoVDB.h:1860
float mTaperF
Definition: NanoVDB.h:1388
Implements Tree::probeLeaf(math::Coord)
Definition: NanoVDB.h:1707
__hostdev__ ChildIter(RootT *parent)
Definition: NanoVDB.h:2735
__hostdev__ void setValue(uint32_t offset, const ValueType &value)
Definition: NanoVDB.h:3807
__hostdev__ CoordBBox bbox() const
Return the index bounding box of all the active values in this tree, i.e. in all nodes of the tree...
Definition: NanoVDB.h:2315
MaskT< LOG2DIM > mValueMask
Definition: NanoVDB.h:4189
typename RootT::CoordType CoordType
Definition: NanoVDB.h:4977
__hostdev__ MagicType toMagic(uint64_t magic)
maps 64 bits of magic number to enum
Definition: NanoVDB.h:368
__hostdev__ const DataType * data() const
Definition: NanoVDB.h:3503
__hostdev__ bool isLevelSet() const
Definition: NanoVDB.h:2176
RootType RootNodeType
Definition: NanoVDB.h:2050
__hostdev__ void clear()
Reset this access to its initial state, i.e. with an empty cache.
Definition: NanoVDB.h:5666
static __hostdev__ uint64_t memUsage()
Return memory usage in bytes for this class only.
Definition: NanoVDB.h:2009
__hostdev__ void setValueOnly(const CoordT &ijk, const ValueType &v)
Definition: NanoVDB.h:4605
__hostdev__ const NodeT * getNode() const
Return a const point to the cached node of the specified type.
Definition: NanoVDB.h:5650
static __hostdev__ constexpr uint32_t padding()
Definition: NanoVDB.h:3991
__hostdev__ bool isFloatingPoint(GridType gridType)
return true if the GridType maps to a floating point type
Definition: NanoVDB.h:563
CoordT mBBoxMin
Definition: NanoVDB.h:4088
__hostdev__ void setOff()
Set all bits off.
Definition: NanoVDB.h:1286
__hostdev__ void localToGlobalCoord(Coord &ijk) const
modifies local coordinates to global coordinates of a tile or child node
Definition: NanoVDB.h:3618
__hostdev__ void setAvg(const StatsT &v)
Definition: NanoVDB.h:2655
__hostdev__ CoordType getOrigin() const
Definition: NanoVDB.h:3400
__hostdev__ const LeafNodeType * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:3560
__hostdev__ bool probeValue(const CoordType &ijk, ValueType &v) const
return the state and updates the value of the specified voxel
Definition: NanoVDB.h:2976
__hostdev__ Vec3T indexToWorldGrad(const Vec3T &grad) const
transform the gradient from index space to world space.
Definition: NanoVDB.h:2135
__hostdev__ uint64_t activeVoxelCount() const
Computes a AABB of active values in world space.
Definition: NanoVDB.h:2170
__hostdev__ const LeafNode * probeLeaf(const CoordT &) const
Definition: NanoVDB.h:4629
__hostdev__ uint64_t * words()
Return a pointer to the list of words of the bit mask.
Definition: NanoVDB.h:1159
__hostdev__ void init(float min, float max, uint8_t bitWidth)
Definition: NanoVDB.h:3871
__hostdev__ bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:5031
__hostdev__ Version(uint32_t major, uint32_t minor, uint32_t patch)
Constructor from major.minor.patch version numbers.
Definition: NanoVDB.h:693
__hostdev__ Mask & operator|=(const Mask &other)
Bitwise union.
Definition: NanoVDB.h:1314
static __hostdev__ uint32_t voxelCount()
Return the total number of voxels (e.g. values) encoded in this leaf node.
Definition: NanoVDB.h:4572
typename ChildT::BuildType BuildT
Definition: NanoVDB.h:3201
__hostdev__ DataType * data()
Definition: NanoVDB.h:4508
__hostdev__ Mask & operator-=(const Mask &other)
Bitwise difference.
Definition: NanoVDB.h:1322
__hostdev__ Checksum(uint64_t checksum, CheckMode mode=CheckMode::Full)
Definition: NanoVDB.h:1764
__hostdev__ ValueIterator beginValue()
Definition: NanoVDB.h:2820
__hostdev__ const RootT & root() const
Definition: NanoVDB.h:5639
uint64_t mPointCount
Definition: NanoVDB.h:4308
__hostdev__ auto getNodeInfo(const CoordType &ijk) const
Definition: NanoVDB.h:5029
static __hostdev__ uint64_t memUsage()
Definition: NanoVDB.h:4195
__hostdev__ CoordType origin() const
Return the origin in index space of this leaf node.
Definition: NanoVDB.h:3520
__hostdev__ DenseIterator(uint32_t pos=Mask::SIZE)
Definition: NanoVDB.h:1105
__hostdev__ CoordT getCoord() const
Definition: NanoVDB.h:4438
MaskT< LOG2DIM > mMask
Definition: NanoVDB.h:4283
PointAccessor(const NanoGrid< BuildT > &grid)
Definition: NanoVDB.h:6018
__hostdev__ Vec3T applyIJT(const Vec3T &xyz) const
Apply the transposed inverse 3x3 transformation to an input 3d vector using 64bit floating point arit...
Definition: NanoVDB.h:1513
__hostdev__ void setValueOnly(uint32_t offset, uint16_t value)
Definition: NanoVDB.h:4326
__hostdev__ uint64_t getIndex(const math::Coord &ijk) const
Return the linear offset into a channel that maps to the specified coordinate.
Definition: NanoVDB.h:6206
ValueT mMinimum
Definition: NanoVDB.h:3223
__hostdev__ bool setGridName(const char *src)
Definition: NanoVDB.h:1916
__hostdev__ void setMask(uint32_t offset, bool v)
Definition: NanoVDB.h:4275
MaskT< LOG2DIM > mValueMask
Definition: NanoVDB.h:4091
__hostdev__ void localToGlobalCoord(Coord &ijk) const
Converts (in place) a local index coordinate to a global index coordinate.
Definition: NanoVDB.h:4547
__hostdev__ ReadAccessor(const RootT &root)
Constructor from a root node.
Definition: NanoVDB.h:5616
__hostdev__ void setOrigin(const T &ijk)
Definition: NanoVDB.h:4345
uint64_t ValueType
Definition: NanoVDB.h:4296
Dummy type for a 8bit quantization of float point values.
Definition: NanoVDB.h:194
__hostdev__ DataType * data()
Definition: NanoVDB.h:2371
typename NanoLeaf< BuildT >::ValueType ValueType
Definition: NanoVDB.h:6625
MagicType
Enums used to identify magic numbers recognized by NanoVDB.
Definition: NanoVDB.h:357
__hostdev__ uint32_t getDim(const CoordType &ijk, const RayT &ray) const
Definition: NanoVDB.h:5851
__hostdev__ bool isValueOn() const
Definition: NanoVDB.h:2895
Dummy type for a voxel whose value equals its binary active state.
Definition: NanoVDB.h:185
uint8_t mFlags
Definition: NanoVDB.h:4188
uint64_t mPrefixSum
Definition: NanoVDB.h:4190
__hostdev__ CoordType getCoord() const
Definition: NanoVDB.h:2716
__hostdev__ Vec3T applyJacobian(const Vec3T &xyz) const
Definition: NanoVDB.h:1928
__hostdev__ util::enable_if< util::is_same< T, Point >::value, const uint64_t & >::type pointCount() const
Return the total number of points indexed by this PointGrid.
Definition: NanoVDB.h:2097
__hostdev__ const RootT & root() const
Definition: NanoVDB.h:5353
__hostdev__ ChildIter(ParentT *parent)
Definition: NanoVDB.h:3347
uint32_t mGridIndex
Definition: NanoVDB.h:1851
__hostdev__ ValueOnIterator(const LeafNode *parent)
Definition: NanoVDB.h:4394
__hostdev__ ValueType operator()(int i, int j, int k) const
Definition: NanoVDB.h:5407
uint64_t mVoxelCount
Definition: NanoVDB.h:2293
static __hostdev__ uint32_t CoordToOffset(const CoordType &ijk)
Return the linear offset corresponding to the given coordinate.
Definition: NanoVDB.h:3602
__hostdev__ uint32_t nodeCount() const
Definition: NanoVDB.h:2421
__hostdev__ ValueType operator()(const CoordType &ijk) const
Definition: NanoVDB.h:5406
uint64_t type
Definition: NanoVDB.h:471
__hostdev__ Vec3T applyMap(const Vec3T &ijk) const
Apply the forward affine transformation to a vector using 64bit floating point arithmetics.
Definition: NanoVDB.h:1438
CoordT CoordType
Definition: NanoVDB.h:5609
__hostdev__ CoordType getCoord() const
Definition: NanoVDB.h:3444
static __hostdev__ bool hasStats()
Definition: NanoVDB.h:3803
__hostdev__ const CoordBBox & indexBBox() const
return AABB of active values in index space
Definition: NanoVDB.h:2015
__hostdev__ bool isFloatingPointVector(GridType gridType)
return true if the GridType maps to a floating point vec3.
Definition: NanoVDB.h:577
ValueT mBackground
Definition: NanoVDB.h:2549
const typename GridOrTreeOrRootT::RootNodeType::ChildNodeType type
Definition: NanoVDB.h:1674
__hostdev__ ValueOnIterator()
Definition: NanoVDB.h:3423
__hostdev__ bool isInteger(GridType gridType)
Return true if the GridType maps to a POD integer type.
Definition: NanoVDB.h:589
__hostdev__ AccessorType getAccessor() const
Definition: NanoVDB.h:2382
__hostdev__ uint64_t leafPoints(const Coord &ijk, const AttT *&begin, const AttT *&end) const
Return the number of points in the leaf node containing the coordinate ijk. If this return value is l...
Definition: NanoVDB.h:6112
NANOVDB_HOSTDEV_DISABLE_WARNING __hostdev__ uint32_t findPrev(uint32_t start) const
Definition: NanoVDB.h:1364
__hostdev__ bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:2974
__hostdev__ ValueType operator*() const
Definition: NanoVDB.h:2791
__hostdev__ const RootT & root() const
Definition: NanoVDB.h:5160
const typename GridOrTreeOrRootT::RootNodeType::ChildNodeType::ChildNodeType type
Definition: NanoVDB.h:1660
void ArrayType
Definition: NanoVDB.h:4135
Defines an affine transform and its inverse represented as a 3x3 matrix and a vec3 translation...
Definition: NanoVDB.h:1383
uint64_t FloatType
Definition: NanoVDB.h:4182
float Type
Definition: NanoVDB.h:512
__hostdev__ void clear()
Reset this access to its initial state, i.e. with an empty cache Noop since this template specializa...
Definition: NanoVDB.h:5013
__hostdev__ const Tile * probeTile(const CoordT &ijk) const
Definition: NanoVDB.h:2629
#define NANOVDB_ASSERT(x)
Definition: Util.h:50
typename GridOrTreeOrRootT::RootNodeType::ChildNodeType type
Definition: NanoVDB.h:1667
GridType
List of types that are currently supported by NanoVDB.
Definition: NanoVDB.h:219
typename BuildT::CoordType CoordType
Definition: NanoVDB.h:2057
__hostdev__ ValueOnIterator cbeginValueOn() const
Definition: NanoVDB.h:4413
__hostdev__ DenseIterator & operator++()
Definition: NanoVDB.h:1113
__hostdev__ void setOn()
Set all bits on.
Definition: NanoVDB.h:1280
__hostdev__ FloatType average() const
Return a const reference to the average of all the active values encoded in this leaf node...
Definition: NanoVDB.h:4523
Class to access points at a specific voxel location.
Definition: NanoVDB.h:6011
__hostdev__ Mask & operator^=(const Mask &other)
Bitwise XOR.
Definition: NanoVDB.h:1330
ValueT mMaximum
Definition: NanoVDB.h:2551
static __hostdev__ uint64_t alignmentPadding(const void *p)
return the smallest number of bytes that when added to the specified pointer results in a 32 byte ali...
Definition: NanoVDB.h:550
__hostdev__ ReadAccessor(const RootT &root)
Constructor from a root node.
Definition: NanoVDB.h:5134
__hostdev__ Iterator operator++(int)
Definition: NanoVDB.h:1090
ValueT mMinimum
Definition: NanoVDB.h:2550
__hostdev__ ChildIter operator++(int)
Definition: NanoVDB.h:2760
bool FloatType
Definition: NanoVDB.h:799
__hostdev__ const uint32_t & activeTileCount(uint32_t level) const
Return the total number of active tiles at the specified level of the tree.
Definition: NanoVDB.h:2414
C++11 implementation of std::is_floating_point.
Definition: Util.h:329
__hostdev__ FloatType getDev() const
Definition: NanoVDB.h:4337
__hostdev__ ValueType operator*() const
Definition: NanoVDB.h:3395
__hostdev__ const RootT & root() const
Definition: NanoVDB.h:5015
__hostdev__ DataType * data()
Definition: NanoVDB.h:3501
const typename GridOrTreeOrRootT::RootNodeType type
Definition: NanoVDB.h:1689
__hostdev__ void setAvg(const StatsT &v)
Definition: NanoVDB.h:3293
__hostdev__ ValueT getValue(uint32_t n) const
Definition: NanoVDB.h:3265
static DstT * PtrAdd(void *p, int64_t offset)
Adds a byte offset to a non-const pointer to produce another non-const pointer.
Definition: Util.h:478
__hostdev__ void setValue(const CoordT &ijk, const ValueType &v)
Sets the value at the specified location and activate its state.
Definition: NanoVDB.h:4599
__hostdev__ ValueOnIter & operator++()
Definition: NanoVDB.h:2845
__hostdev__ float getAvg() const
return the quantized average of the active values in this node
Definition: NanoVDB.h:3886
Class that encapsulates two CRC32 checksums, one for the Grid, Tree and Root node meta data and one f...
Definition: NanoVDB.h:1740
__hostdev__ const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:5714
__hostdev__ BaseIter(DataT *data=nullptr, uint32_t n=0)
Definition: NanoVDB.h:2699
__hostdev__ uint64_t activeVoxelCount() const
Return a const reference to the index bounding box of all the active values in this tree...
Definition: NanoVDB.h:2407
__hostdev__ ValueType operator()(const CoordType &ijk) const
Definition: NanoVDB.h:5027
__hostdev__ float getMax() const
return the quantized maximum of the active values in this node
Definition: NanoVDB.h:3883
typename ChildT::ValueType ValueT
Definition: NanoVDB.h:3200
__hostdev__ bool getDev() const
Definition: NanoVDB.h:4102
Implements Tree::getDim(math::Coord)
Definition: NanoVDB.h:1703
Definition: NanoVDB.h:2692
__hostdev__ const ValueType & background() const
Return the total number of active voxels in the root and all its child nodes.
Definition: NanoVDB.h:2940
__hostdev__ DenseIterator beginDense() const
Definition: NanoVDB.h:3492
Codec
Define compression codecs.
Definition: NanoVDB.h:6258
__hostdev__ Vec3T applyIJT(const Vec3T &xyz) const
Definition: NanoVDB.h:1932
__hostdev__ uint32_t countOn() const
Return the total number of set bits in this Mask.
Definition: NanoVDB.h:1050
uint8_t mFlags
Definition: NanoVDB.h:3853
__hostdev__ bool isChild(uint32_t n) const
Definition: NanoVDB.h:3277
Internal nodes of a VDB tree.
Definition: NanoVDB.h:3308
__hostdev__ ValueOnIterator()
Definition: NanoVDB.h:4389
__hostdev__ ValueType getMax() const
Definition: NanoVDB.h:4335
__hostdev__ ConstDenseIterator cbeginChildAll() const
Definition: NanoVDB.h:2919
static __hostdev__ T * alignPtr(T *p)
offset the specified pointer so it is 32 byte aligned. Works with both const and non-const pointers...
Definition: NanoVDB.h:558
__hostdev__ bool isOn() const
Return true if all the bits are set in this Mask.
Definition: NanoVDB.h:1211
ValueT ValueType
Definition: NanoVDB.h:5608
__hostdev__ ValueType getValue(int i, int j, int k) const
Definition: NanoVDB.h:5026
BuildT TreeType
Definition: NanoVDB.h:2048
Base-class for quantized float leaf nodes.
Definition: NanoVDB.h:3844
uint64_t FloatType
Definition: NanoVDB.h:793
math::BBox< CoordT > mBBox
Definition: NanoVDB.h:3218
__hostdev__ const LeafNodeType * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:2977
__hostdev__ void setMin(const ValueT &v)
Definition: NanoVDB.h:2653
__hostdev__ Vec3T worldToIndexF(const Vec3T &xyz) const
world to index space transformation
Definition: NanoVDB.h:2139
static __hostdev__ constexpr uint32_t padding()
Definition: NanoVDB.h:4095
const typename GridOrTreeOrRootT::RootNodeType::ChildNodeType::ChildNodeType Type
Definition: NanoVDB.h:1659
__hostdev__ uint64_t getMin() const
Definition: NanoVDB.h:4249
__hostdev__ CoordType getOrigin() const
Definition: NanoVDB.h:3484
__hostdev__ bool isCached(const CoordType &ijk) const
Definition: NanoVDB.h:5167
__hostdev__ Vec3T worldToIndex(const Vec3T &xyz) const
world to index space transformation
Definition: NanoVDB.h:2116
Definition: NanoVDB.h:2288
static __hostdev__ uint64_t memUsage()
Definition: NanoVDB.h:4319
C++11 implementation of std::is_same.
Definition: Util.h:314
__hostdev__ ReadAccessor(const GridT &grid)
Constructor from a grid.
Definition: NanoVDB.h:5628
static __hostdev__ uint64_t memUsage()
Definition: NanoVDB.h:3860
__hostdev__ void setValue(uint32_t offset, bool v)
Definition: NanoVDB.h:4103
__hostdev__ bool isActive() const
Return true if this node or any of its child nodes contain active values.
Definition: NanoVDB.h:3632
__hostdev__ ValueType getValue(const CoordType &ijk) const
Return the value of the given voxel (regardless of state or location in the tree.) ...
Definition: NanoVDB.h:2385
__hostdev__ uint64_t getMin() const
Definition: NanoVDB.h:4229
Struct with all the member data of the InternalNode (useful during serialization of an openvdb Intern...
Definition: NanoVDB.h:3198
static __hostdev__ constexpr int64_t memUsage()
Definition: NanoVDB.h:3959
__hostdev__ const ChildNodeType * probeChild(const CoordType &ijk) const
Definition: NanoVDB.h:3024
__hostdev__ const NanoGrid< Point > & grid() const
Definition: NanoVDB.h:6098
static __hostdev__ constexpr uint64_t memUsage()
Definition: NanoVDB.h:3990
const typename GridT::TreeType Type
Definition: NanoVDB.h:2333
Dummy type for a 4bit quantization of float point values.
Definition: NanoVDB.h:191
__hostdev__ bool operator!=(const Checksum &rhs) const
return true if the checksums are not identical
Definition: NanoVDB.h:1810
__hostdev__ uint64_t gridSize() const
Return memory usage in bytes for this class only.
Definition: NanoVDB.h:2077
typename ChildT::CoordType CoordT
Definition: NanoVDB.h:3203
uint64_t mCRC64
Definition: NanoVDB.h:1746
__hostdev__ uint64_t & full()
Definition: NanoVDB.h:1778
__hostdev__ void setMin(const ValueType &)
Definition: NanoVDB.h:4158
Return point to the lower internal node where math::Coord maps to one of its values, i.e. terminates.
Definition: NanoVDB.h:6566
uint64_t type
Definition: NanoVDB.h:492
static __hostdev__ bool hasStats()
Definition: NanoVDB.h:4097
const typename GridT::TreeType type
Definition: NanoVDB.h:2334
__hostdev__ NodeTrait< RootT, 1 >::type * getFirstLower()
Definition: NanoVDB.h:2479
__hostdev__ ValueType operator()(int i, int j, int k) const
Definition: NanoVDB.h:5181
__hostdev__ FloatType variance() const
Return the variance of all the active values encoded in this internal node and any of its child nodes...
Definition: NanoVDB.h:3532
__hostdev__ const ValueT & getMin() const
Definition: NanoVDB.h:3282
__hostdev__ uint64_t voxelPoints(const Coord &ijk, const AttT *&begin, const AttT *&end) const
get iterators over attributes to points at a specific voxel location
Definition: NanoVDB.h:6123
uint8_t mFlags
Definition: NanoVDB.h:4090
T type
Definition: Util.h:387
__hostdev__ uint64_t getAvg() const
Definition: NanoVDB.h:4251
__hostdev__ void setBBoxOn(bool on=true)
Definition: NanoVDB.h:1912
__hostdev__ bool isUnknown() const
Definition: NanoVDB.h:2183
static __hostdev__ constexpr uint32_t padding()
Return padding of this class in bytes, due to aliasing and 32B alignment.
Definition: NanoVDB.h:2558
__hostdev__ uint32_t head() const
Definition: NanoVDB.h:1779
T type
Definition: NanoVDB.h:464
__hostdev__ ValueIterator & operator++()
Definition: NanoVDB.h:4483
__hostdev__ uint32_t blindDataCount() const
Return true if this grid is empty, i.e. contains no values or nodes.
Definition: NanoVDB.h:2217
__hostdev__ void setChild(uint32_t n, const void *ptr)
Definition: NanoVDB.h:3240
__hostdev__ Vec3T applyInverseJacobian(const Vec3T &xyz) const
Definition: NanoVDB.h:1930
__hostdev__ bool operator==(const Version &rhs) const
Definition: NanoVDB.h:700
Struct with all the member data of the Grid (useful during serialization of an openvdb grid) ...
Definition: NanoVDB.h:1844
bool ValueType
Definition: NanoVDB.h:4132
typename ChildT::template MaskType< LOG2 > MaskType
Definition: NanoVDB.h:3320
auto callNanoGrid(GridDataT *gridData, ArgsT &&...args)
Below is an example of the struct used for generic programming with callNanoGrid. ...
Definition: NanoVDB.h:4889
Implements Tree::isActive(math::Coord)
Definition: NanoVDB.h:1701
__hostdev__ Vec3T applyInverseMapF(const Vec3T &xyz) const
Apply the inverse affine mapping to a vector using 32bit floating point arithmetics.
Definition: NanoVDB.h:1483
Definition: NanoVDB.h:728
bool FloatType
Definition: NanoVDB.h:4134
__hostdev__ bool probeValue(const CoordT &ijk, ValueType &v) const
Return true if the voxel value at the given coordinate is active and updates v with the value...
Definition: NanoVDB.h:4622
__hostdev__ NodeTrait< RootT, 2 >::type * getFirstUpper()
Definition: NanoVDB.h:2481
__hostdev__ void toggle()
brief Toggle the state of all bits in the mask
Definition: NanoVDB.h:1298
__hostdev__ bool isPointIndex() const
Definition: NanoVDB.h:2179
__hostdev__ void setMax(const ValueType &)
Definition: NanoVDB.h:4340
uint32_t mData0
Definition: NanoVDB.h:1862
__hostdev__ ValueType operator*() const
Definition: NanoVDB.h:4400
Dummy type for indexing points into voxels.
Definition: NanoVDB.h:203
__hostdev__ const MaskType< LOG2DIM > & getValueMask() const
Definition: NanoVDB.h:3513
__hostdev__ ValueType getValue(const CoordT &ijk) const
Return the voxel value at the given coordinate.
Definition: NanoVDB.h:4589
static __hostdev__ size_t memUsage()
Return memory usage in bytes for the class.
Definition: NanoVDB.h:3509
__hostdev__ NodeT & operator*() const
Definition: NanoVDB.h:3353
typename ChildT::FloatType StatsT
Definition: NanoVDB.h:3202
Definition: NanoVDB.h:902
__hostdev__ bool isActive(const CoordType &ijk) const
Return the active state of the given voxel (regardless of state or location in the tree...
Definition: NanoVDB.h:2389
__hostdev__ const ChildT * getChild(uint32_t n) const
Definition: NanoVDB.h:3259
uint32_t findHighestOn(uint32_t v)
Returns the index of the highest, i.e. most significant, on bit in the specified 32 bit word...
Definition: Util.h:572
DataT * mData
Definition: NanoVDB.h:2697
bool Type
Definition: NanoVDB.h:498
Definition: NanoVDB.h:1102
__hostdev__ ValueType getFirstValue() const
If the first entry in this node's table is a tile, return the tile's value. Otherwise, return the result of calling getFirstValue() on the child.
Definition: NanoVDB.h:3542
StatsT mAverage
Definition: NanoVDB.h:3225
__hostdev__ float getValue(uint32_t i) const
Definition: NanoVDB.h:3998
Definition: NanoVDB.h:2868
__hostdev__ ValueIterator cbeginValueAll() const
Definition: NanoVDB.h:3414
CoordT mBBoxMin
Definition: NanoVDB.h:3851
__hostdev__ NodeT & operator*() const
Definition: NanoVDB.h:2742
typename ChildT::CoordType CoordT
Definition: NanoVDB.h:2518
__hostdev__ uint64_t getMax() const
Definition: NanoVDB.h:4230
__hostdev__ ValueType getValue(uint32_t offset) const
Return the voxel value at the given offset.
Definition: NanoVDB.h:4586
__hostdev__ ValueIter & operator++()
Definition: NanoVDB.h:2801
MaskT mValueMask
Definition: NanoVDB.h:3220
NANOVDB_HOSTDEV_DISABLE_WARNING __hostdev__ uint32_t findNext(uint32_t start) const
Definition: NanoVDB.h:1350
__hostdev__ CoordType getOrigin() const
Definition: NanoVDB.h:2711
__hostdev__ uint32_t totalNodeCount() const
Definition: NanoVDB.h:2433
uint16_t mMin
Definition: NanoVDB.h:3858
typename ChildT::FloatType StatsT
Definition: NanoVDB.h:2519
typename GridOrTreeOrRootT::RootNodeType::ChildNodeType Type
Definition: NanoVDB.h:1666
typename FloatTraits< ValueType >::FloatType FloatType
Definition: NanoVDB.h:4298
__hostdev__ const ValueT & getMin() const
Definition: NanoVDB.h:2648
Like ValueOnIndex but with a mutable mask.
Definition: NanoVDB.h:182
const typename GridOrTreeOrRootT::LeafNodeType Type
Definition: NanoVDB.h:1644
__hostdev__ DataType * data()
Definition: NanoVDB.h:2069
MaskT< LOG2DIM > mValues
Definition: NanoVDB.h:4092
__hostdev__ int findBlindData(const char *name) const
Return the index of the first blind data with specified name if found, otherwise -1.
Definition: NanoVDB.h:2268
uint32_t mTableSize
Definition: NanoVDB.h:2547
typename BuildT::BuildType BuildType
Definition: NanoVDB.h:2056
typename T::ValueType ElementType
Definition: NanoVDB.h:748
__hostdev__ uint64_t memUsage() const
return memory usage in bytes for the leaf node
Definition: NanoVDB.h:4577
__hostdev__ bool isSequential() const
return true if the specified node type is layed out breadth-first in memory and has a fixed size...
Definition: NanoVDB.h:2194
Definition: NanoVDB.h:4364
typename RootT::CoordType CoordType
Definition: NanoVDB.h:2357
float type
Definition: NanoVDB.h:534
defines a tree type from a grid type while preserving constness
Definition: NanoVDB.h:2325
__hostdev__ bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:5410
__hostdev__ GridType mapToGridType()
Definition: NanoVDB.h:872
__hostdev__ uint32_t nodeCount(int level) const
Definition: NanoVDB.h:2427
__hostdev__ ChannelT & operator()(int i, int j, int k) const
Definition: NanoVDB.h:6212
__hostdev__ AccessorType getAccessor() const
Return a new instance of a ReadAccessor used to access values in this grid.
Definition: NanoVDB.h:2106
Visits child nodes of this node only.
Definition: NanoVDB.h:3334
__hostdev__ Coord offsetToGlobalCoord(uint32_t n) const
Definition: NanoVDB.h:3624
typename remove_const< T >::type type
Definition: Util.h:431
static __hostdev__ constexpr uint32_t padding()
Definition: NanoVDB.h:4146
__hostdev__ void setValue(uint32_t offset, uint16_t value)
Definition: NanoVDB.h:4327
__hostdev__ Checksum()
default constructor initiates checksum to EMPTY
Definition: NanoVDB.h:1754
uint64_t Type
Definition: NanoVDB.h:491
static __hostdev__ constexpr uint32_t padding()
Definition: NanoVDB.h:3960
__hostdev__ ValueIterator(const InternalNode *parent)
Definition: NanoVDB.h:3389
typename Mask< 3 >::template Iterator< ON > MaskIterT
Definition: NanoVDB.h:4380
Leaf nodes of the VDB tree. (defaults to 8x8x8 = 512 voxels)
Definition: NanoVDB.h:4361
__hostdev__ bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:5183
__hostdev__ DataType * data()
Definition: NanoVDB.h:2929
__hostdev__ const uint64_t & valueCount() const
Return total number of values indexed by the IndexGrid.
Definition: NanoVDB.h:6190
__hostdev__ NodeTrait< RootT, LEVEL >::type * getFirstNode()
return a pointer to the first node at the specified level
Definition: NanoVDB.h:2462
__hostdev__ bool isValue() const
Definition: NanoVDB.h:2581
__hostdev__ Vec3T worldToIndexDir(const Vec3T &dir) const
transformation from world space direction to index space direction
Definition: NanoVDB.h:2130
__hostdev__ DenseIterator cbeginChildAll() const
Definition: NanoVDB.h:3493
BuildT BuildType
Definition: NanoVDB.h:5308
__hostdev__ uint32_t rootTableSize() const
return the root table has size
Definition: NanoVDB.h:2018
bool FloatType
Definition: NanoVDB.h:4084
__hostdev__ bool hasBBox() const
Definition: NanoVDB.h:4619
double mTaperD
Definition: NanoVDB.h:1392
__hostdev__ CoordType getCoord() const
Definition: NanoVDB.h:3489
uint32_t dim
Definition: NanoVDB.h:6629
__hostdev__ const ValueType & maximum() const
Return a const reference to the maximum active value encoded in this root node and any of its child n...
Definition: NanoVDB.h:2950
MaskT mChildMask
Definition: NanoVDB.h:3221
__hostdev__ bool isActive(uint32_t n) const
Definition: NanoVDB.h:4609
__hostdev__ Version()
Default constructor.
Definition: NanoVDB.h:684
__hostdev__ void setMinMaxOn(bool on=true)
Definition: NanoVDB.h:1911
static __hostdev__ uint32_t valueCount()
Definition: NanoVDB.h:4225
MaskT< LOG2DIM > mValueMask
Definition: NanoVDB.h:3786
__hostdev__ const Tile * tile(uint32_t n) const
Returns a non-const reference to the tile at the specified linear offset.
Definition: NanoVDB.h:2593
__hostdev__ const StatsT & average() const
Definition: NanoVDB.h:3284
__hostdev__ ValueType getFirstValue() const
Return the first value in this leaf node.
Definition: NanoVDB.h:4592
__hostdev__ ValueOnIterator cbeginValueOn() const
Definition: NanoVDB.h:3448
typename GridOrTreeOrRootT::RootNodeType Type
Definition: NanoVDB.h:1680
__hostdev__ ValueOnIterator(const InternalNode *parent)
Definition: NanoVDB.h:3428
__hostdev__ ConstValueOnIterator cbeginValueOn() const
Definition: NanoVDB.h:2865
Definition: NanoVDB.h:2563
__hostdev__ ReadAccessor(const GridT &grid)
Constructor from a grid.
Definition: NanoVDB.h:5142
FloatType mAverage
Definition: NanoVDB.h:3790
BuildT BuildType
Definition: NanoVDB.h:4975
ValueT ValueType
Definition: NanoVDB.h:5309
__hostdev__ const ChildNodeType * probeChild(const CoordType &ijk) const
Definition: NanoVDB.h:3595
float Type
Definition: NanoVDB.h:505
typename UpperNodeType::ChildNodeType LowerNodeType
Definition: NanoVDB.h:2677
StatsT mStdDevi
Definition: NanoVDB.h:2553
__hostdev__ void setOrigin(const T &ijk)
Definition: NanoVDB.h:4205
__hostdev__ const DataType * data() const
Definition: NanoVDB.h:2071
__hostdev__ uint32_t & head()
Definition: NanoVDB.h:1780
ValueT value
Definition: NanoVDB.h:3209
static __hostdev__ constexpr uint32_t padding()
Return padding of this class in bytes, due to aliasing and 32B alignment.
Definition: NanoVDB.h:4315
__hostdev__ bool isCached1(const CoordType &ijk) const
Definition: NanoVDB.h:5386
__hostdev__ bool isActive(uint32_t n) const
Definition: NanoVDB.h:3271
__hostdev__ ValueOnIterator beginValueOn()
Definition: NanoVDB.h:2864
__hostdev__ const ChildT * getChild(const Tile *tile) const
Definition: NanoVDB.h:2642
__hostdev__ bool isEmpty() const
return true if the 64 bit checksum is disables (unset)
Definition: NanoVDB.h:1793
__hostdev__ Iterator(uint32_t pos, const Mask *parent)
Definition: NanoVDB.h:1076
__hostdev__ ReadAccessor(const RootT &root)
Constructor from a root node.
Definition: NanoVDB.h:5317
__hostdev__ ValueType getMax() const
Definition: NanoVDB.h:3815
typename GridOrTreeOrRootT::RootNodeType type
Definition: NanoVDB.h:1681
__hostdev__ void * nodePtr()
Return a non-const void pointer to the first node at LEVEL.
Definition: NanoVDB.h:1966
__hostdev__ float getMin() const
return the quantized minimum of the active values in this node
Definition: NanoVDB.h:3880
__hostdev__ const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:5185
__hostdev__ ValueOffIterator()
Definition: NanoVDB.h:4422
ChildT ChildNodeType
Definition: NanoVDB.h:3316
typename DataType::BuildT BuildType
Definition: NanoVDB.h:3314
__hostdev__ ValueOffIterator cbeginValueOff() const
Definition: NanoVDB.h:4446
uint32_t mSize
Definition: NanoVDB.h:2698
__hostdev__ GridClass toGridClass(GridClass defaultClass=GridClass::Unknown)
Maps from a templated build type to a GridClass enum.
Definition: NanoVDB.h:878
typename DataType::ValueType ValueType
Definition: NanoVDB.h:4372
float type
Definition: NanoVDB.h:527
__hostdev__ uint32_t getMajor() const
Definition: NanoVDB.h:706
__hostdev__ Vec3T indexToWorldGradF(const Vec3T &grad) const
Transforms the gradient from index space to world space.
Definition: NanoVDB.h:2158
__hostdev__ const NodeTrait< TreeT, LEVEL >::type * getNode() const
Definition: NanoVDB.h:5658
__hostdev__ bool hasBBox() const
Definition: NanoVDB.h:2185
uint64_t type
Definition: NanoVDB.h:541
__hostdev__ FloatType getAvg() const
Definition: NanoVDB.h:4336
typename ChildT::LeafNodeType LeafNodeType
Definition: NanoVDB.h:3315
__hostdev__ void setValue(const CoordType &k, bool s, const ValueType &v)
Definition: NanoVDB.h:2573
__hostdev__ auto getNodeInfo(const CoordType &ijk) const
Definition: NanoVDB.h:5711
__hostdev__ const uint64_t * words() const
Definition: NanoVDB.h:1160
__hostdev__ const GridBlindMetaData & blindMetaData(uint32_t n) const
Definition: NanoVDB.h:2251
static __hostdev__ uint32_t bitCount()
Return the number of bits available in this Mask.
Definition: NanoVDB.h:1044
__hostdev__ void setDev(const bool &)
Definition: NanoVDB.h:4112
uint64_t Type
Definition: NanoVDB.h:484
__hostdev__ Vec3T indexToWorldDir(const Vec3T &dir) const
transformation from index space direction to world space direction
Definition: NanoVDB.h:2125
__hostdev__ const void * getRoot() const
Get a const void pointer to the root node (never NULL)
Definition: NanoVDB.h:2306
__hostdev__ const StatsT & stdDeviation() const
Definition: NanoVDB.h:3285
__hostdev__ bool isActive() const
Return true if any of the voxel value are active in this leaf node.
Definition: NanoVDB.h:4612
GridBlindDataClass
Blind-data Classes that are currently supported by NanoVDB.
Definition: NanoVDB.h:416
MaskT< LOG2DIM > mValueMask
Definition: NanoVDB.h:4305
__hostdev__ const void * treePtr() const
Definition: NanoVDB.h:1949
static __hostdev__ size_t memUsage()
Return the memory footprint in bytes of this Mask.
Definition: NanoVDB.h:1041
const typename GridOrTreeOrRootT::RootNodeType Type
Definition: NanoVDB.h:1688
Visits all active values in a leaf node.
Definition: NanoVDB.h:4383
__hostdev__ const LeafNodeType * getFirstLeaf() const
Definition: NanoVDB.h:2478
__hostdev__ Vec3T indexToWorldDirF(const Vec3T &dir) const
transformation from index space direction to world space direction
Definition: NanoVDB.h:2148