OpenVDB  9.1.0
NanoVDB.h
Go to the documentation of this file.
1 // Copyright Contributors to the OpenVDB Project
2 // SPDX-License-Identifier: MPL-2.0
3 
4 /*!
5  \file NanoVDB.h
6 
7  \author Ken Museth
8 
9  \date January 8, 2020
10 
11  \brief Implements a light-weight self-contained VDB data-structure in a
12  single file! In other words, this is a significantly watered-down
13  version of the OpenVDB implementation, with few dependencies - so
14  a one-stop-shop for a minimalistic VDB data structure that run on
15  most platforms!
16 
17  \note It is important to note that NanoVDB (by design) is a read-only
18  sparse GPU (and CPU) friendly data structure intended for applications
19  like rendering and collision detection. As such it obviously lacks
20  a lot of the functionality and features of OpenVDB grids. NanoVDB
21  is essentially a compact linearized (or serialized) representation of
22  an OpenVDB tree with getValue methods only. For best performance use
23  the ReadAccessor::getValue method as opposed to the Tree::getValue
24  method. Note that since a ReadAccessor caches previous access patterns
25  it is by design not thread-safe, so use one instantiation per thread
26  (it is very light-weight). Also, it is not safe to copy accessors between
27  the GPU and CPU! In fact, client code should only interface
28  with the API of the Grid class (all other nodes of the NanoVDB data
29  structure can safely be ignored by most client codes)!
30 
31 
32  \warning NanoVDB grids can only be constructed via tools like openToNanoVDB
33  or the GridBuilder. This explains why none of the grid nodes defined below
34  have public constructors or destructors.
35 
36  \details Please see the following paper for more details on the data structure:
37  K. Museth, “VDB: High-Resolution Sparse Volumes with Dynamic Topology”,
38  ACM Transactions on Graphics 32(3), 2013, which can be found here:
39  http://www.museth.org/Ken/Publications_files/Museth_TOG13.pdf
40 
41 
42  Overview: This file implements the following fundamental class that when combined
43  forms the backbone of the VDB tree data structure:
44 
45  Coord- a signed integer coordinate
46  Vec3 - a 3D vector
47  Vec4 - a 4D vector
48  BBox - a bounding box
49  Mask - a bitmask essential to the non-root tree nodes
50  Map - an affine coordinate transformation
51  Grid - contains a Tree and a map for world<->index transformations. Use
52  this class as the main API with client code!
53  Tree - contains a RootNode and getValue methods that should only be used for debugging
54  RootNode - the top-level node of the VDB data structure
55  InternalNode - the internal nodes of the VDB data structure
56  LeafNode - the lowest level tree nodes that encode voxel values and state
57  ReadAccessor - implements accelerated random access operations
58 
59  Semantics: A VDB data structure encodes values and (binary) states associated with
60  signed integer coordinates. Values encoded at the leaf node level are
61  denoted voxel values, and values associated with other tree nodes are referred
62  to as tile values, which by design cover a larger coordinate index domain.
63 
64 
65  Memory layout:
66 
67  GridData is always at the very beginning of the buffer immediately followed by TreeData!
68  The remaining nodes and blind-data are allowed to be scattered thoughout the buffer,
69  though in practice they are arranged as:
70 
71  GridData: 672 bytes (e.g. magic, checksum, major, flags, index, count, size, name, map, world bbox, voxel size, class, type, offset, count)
72 
73  TreeData: 64 bytes (node counts and byte offsets)
74 
75  ... optional padding ...
76 
77  RootData: size depends on ValueType (index bbox, voxel count, tile count, min/max/avg/standard deviation)
78 
79  Array of: RootData::Tile
80 
81  ... optional padding ...
82 
83  Array of: Upper InternalNodes of size 32^3: bbox, two bit masks, 32768 tile values, and min/max/avg/standard deviation values
84 
85  ... optional padding ...
86 
87  Array of: Lower InternalNodes of size 16^3: bbox, two bit masks, 4096 tile values, and min/max/avg/standard deviation values
88 
89  ... optional padding ...
90 
91  Array of: LeafNodes of size 8^3: bbox, bit masks, 512 voxel values, and min/max/avg/standard deviation values
92 
93 
94  Example layout: ("---" implies it has a custom offset, "..." implies zero or more)
95  [GridData(672B)][TreeData(64B)]---[RootData][N x Root::Tile]---[NodeData<5>]---[ModeData<4>]---[LeafData<3>]---[BLINDMETA...]---[BLIND0]---[BLIND1]---etc.
96 
97 */
98 
99 #ifndef NANOVDB_NANOVDB_H_HAS_BEEN_INCLUDED
100 #define NANOVDB_NANOVDB_H_HAS_BEEN_INCLUDED
101 
102 #define NANOVDB_MAGIC_NUMBER 0x304244566f6e614eUL // "NanoVDB0" in hex - little endian (uint64_t)
103 
104 #define NANOVDB_MAJOR_VERSION_NUMBER 32 // reflects changes to the ABI and hence also the file format
105 #define NANOVDB_MINOR_VERSION_NUMBER 3 // reflects changes to the API but not ABI
106 #define NANOVDB_PATCH_VERSION_NUMBER 3 // reflects changes that does not affect the ABI or API
107 
108 // This replaces a Coord key at the root level with a single uint64_t
109 #define USE_SINGLE_ROOT_KEY
110 
111 // This replaces three levels of Coord keys in the ReadAccessor with one Coord
112 //#define USE_SINGLE_ACCESSOR_KEY
113 
114 #define NANOVDB_FPN_BRANCHLESS
115 
116 #define NANOVDB_DATA_ALIGNMENT 32
117 
118 #if !defined(NANOVDB_ALIGN)
119 #define NANOVDB_ALIGN(n) alignas(n)
120 #endif // !defined(NANOVDB_ALIGN)
121 
122 #ifdef __CUDACC_RTC__
123 
124 typedef signed char int8_t;
125 typedef short int16_t;
126 typedef int int32_t;
127 typedef long long int64_t;
128 typedef unsigned char uint8_t;
129 typedef unsigned int uint32_t;
130 typedef unsigned short uint16_t;
131 typedef unsigned long long uint64_t;
132 
133 #define NANOVDB_ASSERT(x)
134 
135 #define UINT64_C(x) (x ## ULL)
136 
137 #else // __CUDACC_RTC__
138 
139 #include <stdlib.h> // for abs in clang7
140 #include <stdint.h> // for types like int32_t etc
141 #include <stddef.h> // for size_t type
142 #include <cassert> // for assert
143 #include <cstdio> // for sprinf
144 #include <cmath> // for sqrt and fma
145 #include <limits> // for numeric_limits
146 
147 // All asserts can be disabled here, even for debug builds
148 #if 1
149 #define NANOVDB_ASSERT(x) assert(x)
150 #else
151 #define NANOVDB_ASSERT(x)
152 #endif
153 
154 #if defined(NANOVDB_USE_INTRINSICS) && defined(_MSC_VER)
155 #include <intrin.h>
156 #pragma intrinsic(_BitScanReverse)
157 #pragma intrinsic(_BitScanForward)
158 #pragma intrinsic(_BitScanReverse64)
159 #pragma intrinsic(_BitScanForward64)
160 #endif
161 
162 #endif // __CUDACC_RTC__
163 
164 #if defined(__CUDACC__) || defined(__HIP__)
165 // Only define __hostdev__ when using NVIDIA CUDA or HIP compiler
166 #define __hostdev__ __host__ __device__
167 #else
168 #define __hostdev__
169 #endif
170 
171 // The following macro will suppress annoying warnings when nvcc
172 // compiles functions that call (host) intrinsics (which is perfectly valid)
173 #if defined(_MSC_VER) && defined(__CUDACC__)
174 #define NANOVDB_HOSTDEV_DISABLE_WARNING __pragma("hd_warning_disable")
175 #elif defined(__GNUC__) && defined(__CUDACC__)
176 #define NANOVDB_HOSTDEV_DISABLE_WARNING _Pragma("hd_warning_disable")
177 #else
178 #define NANOVDB_HOSTDEV_DISABLE_WARNING
179 #endif
180 
181 // A portable implementation of offsetof - unfortunately it doesn't work with static_assert
182 #define NANOVDB_OFFSETOF(CLASS, MEMBER) ((int)(size_t)((char*)&((CLASS*)0)->MEMBER - (char*)0))
183 
184 namespace nanovdb {
185 
186 // --------------------------> Build types <------------------------------------
187 
188 /// @brief Dummy type for a voxel with a binary mask value, e.g. the active state
189 class ValueMask {};
190 
191 /// @brief Dummy type for a 16 bit floating point values
192 class Half {};
193 
194 /// @brief Dummy type for a 4bit quantization of float point values
195 class Fp4 {};
196 
197 /// @brief Dummy type for a 8bit quantization of float point values
198 class Fp8 {};
199 
200 /// @brief Dummy type for a 16bit quantization of float point values
201 class Fp16 {};
202 
203 /// @brief Dummy type for a variable bit quantization of floating point values
204 class FpN {};
205 
206 // --------------------------> GridType <------------------------------------
207 
208 /// @brief List of types that are currently supported by NanoVDB
209 ///
210 /// @note To expand on this list do:
211 /// 1) Add the new type between Unknown and End in the enum below
212 /// 2) Add the new type to OpenToNanoVDB::processGrid that maps OpenVDB types to GridType
213 /// 3) Verify that the ConvertTrait in NanoToOpenVDB.h works correctly with the new type
214 /// 4) Add the new type to mapToGridType (defined below) that maps NanoVDB types to GridType
215 /// 5) Add the new type to toStr (defined below)
216 enum class GridType : uint32_t { Unknown = 0,
217  Float = 1, // single precision floating point value
218  Double = 2,// double precision floating point value
219  Int16 = 3,// half precision signed integer value
220  Int32 = 4,// single precision signed integer value
221  Int64 = 5,// double precision signed integer value
222  Vec3f = 6,// single precision floating 3D vector
223  Vec3d = 7,// double precision floating 3D vector
224  Mask = 8,// no value, just the active state
225  Half = 9,// half precision floating point value
226  UInt32 = 10,// single precision unsigned integer value
227  Boolean = 11,// boolean value, encoded in bit array
228  RGBA8 = 12,// RGBA packed into 32bit word in reverse-order. R in low bits.
229  Fp4 = 13,// 4bit quantization of float point value
230  Fp8 = 14,// 8bit quantization of float point value
231  Fp16 = 15,// 16bit quantization of float point value
232  FpN = 16,// variable bit quantization of floating point value
233  Vec4f = 17,// single precision floating 4D vector
234  Vec4d = 18,// double precision floating 4D vector
235  End = 19 };
236 
237 #ifndef __CUDACC_RTC__
238 /// @brief Retuns a c-string used to describe a GridType
239 inline const char* toStr(GridType gridType)
240 {
241  static const char * LUT[] = { "?", "float", "double" , "int16", "int32",
242  "int64", "Vec3f", "Vec3d", "Mask", "Half",
243  "uint32", "bool", "RGBA8", "Float4", "Float8",
244  "Float16", "FloatN", "Vec4f", "Vec4d", "End" };
245  static_assert( sizeof(LUT)/sizeof(char*) - 1 == int(GridType::End), "Unexpected size of LUT" );
246  return LUT[static_cast<int>(gridType)];
247 }
248 #endif
249 
250 // --------------------------> GridClass <------------------------------------
251 
252 /// @brief Classes (defined in OpenVDB) that are currently supported by NanoVDB
253 enum class GridClass : uint32_t { Unknown = 0,
254  LevelSet = 1, // narrow band level set, e.g. SDF
255  FogVolume = 2, // fog volume, e.g. density
256  Staggered = 3, // staggered MAC grid, e.g. velocity
257  PointIndex = 4, // point index grid
258  PointData = 5, // point data grid
259  Topology = 6, // grid with active states only (no values)
260  VoxelVolume = 7, // volume of geometric cubes, e.g. minecraft
261  End = 8 };
262 
263 #ifndef __CUDACC_RTC__
264 /// @brief Retuns a c-string used to describe a GridClass
265 inline const char* toStr(GridClass gridClass)
266 {
267  static const char * LUT[] = { "?", "SDF", "FOG" , "MAC", "PNTIDX",
268  "PNTDAT", "TOPO", "VOX", "END" };
269  static_assert( sizeof(LUT)/sizeof(char*) - 1 == int(GridClass::End), "Unexpected size of LUT" );
270  return LUT[static_cast<int>(gridClass)];
271 }
272 #endif
273 
274 // --------------------------> GridFlags <------------------------------------
275 
276 /// @brief Grid flags which indicate what extra information is present in the grid buffer.
277 enum class GridFlags : uint32_t {
278  HasLongGridName = 1 << 0,// grid name is longer than 256 characters
279  HasBBox = 1 << 1,// nodes contain bounding-boxes of active values
280  HasMinMax = 1 << 2,// nodes contain min/max of active values
281  HasAverage = 1 << 3,// nodes contain averages of active values
282  HasStdDeviation = 1 << 4,// nodes contain standard deviations of active values
283  IsBreadthFirst = 1 << 5,// nodes are arranged breadth-first in memory
284  End = 1 << 6,
285 };
286 
287 #ifndef __CUDACC_RTC__
288 /// @brief Retuns a c-string used to describe a GridFlags
289 inline const char* toStr(GridFlags gridFlags)
290 {
291  static const char * LUT[] = { "has long grid name",
292  "has bbox",
293  "has min/max",
294  "has average",
295  "has standard deviation",
296  "is breadth-first",
297  "end" };
298  static_assert( 1 << (sizeof(LUT)/sizeof(char*) - 1) == int(GridFlags::End), "Unexpected size of LUT" );
299  return LUT[static_cast<int>(gridFlags)];
300 }
301 #endif
302 
303 // --------------------------> GridBlindData enums <------------------------------------
304 
305 /// @brief Blind-data Classes that are currently supported by NanoVDB
306 enum class GridBlindDataClass : uint32_t { Unknown = 0,
307  IndexArray = 1,
308  AttributeArray = 2,
309  GridName = 3,
310  End = 4 };
311 
312 /// @brief Blind-data Semantics that are currently understood by NanoVDB
313 enum class GridBlindDataSemantic : uint32_t { Unknown = 0,
314  PointPosition = 1,
315  PointColor = 2,
316  PointNormal = 3,
317  PointRadius = 4,
318  PointVelocity = 5,
319  PointId = 6,
320  End = 7 };
321 
322 // --------------------------> is_same <------------------------------------
323 
324 /// @brief C++11 implementation of std::is_same
325 template<typename T1, typename T2>
326 struct is_same
327 {
328  static constexpr bool value = false;
329 };
330 
331 template<typename T>
332 struct is_same<T, T>
333 {
334  static constexpr bool value = true;
335 };
336 
337 // --------------------------> enable_if <------------------------------------
338 
339 /// @brief C++11 implementation of std::enable_if
340 template <bool, typename T = void>
341 struct enable_if
342 {
343 };
344 
345 template <typename T>
346 struct enable_if<true, T>
347 {
348  using type = T;
349 };
350 
351 // --------------------------> is_floating_point <------------------------------------
352 
353 /// @brief C++11 implementation of std::is_floating_point
354 template<typename T>
356 {
358 };
359 
360 // --------------------------> is_specialization <------------------------------------
361 
362 /// @brief Metafunction used to determine if the first template
363 /// parameter is a specialization of the class template
364 /// given in the second template parameter.
365 ///
366 /// @details is_specialization<Vec3<float>, Vec3>::value == true;
367 template<typename AnyType, template<typename...> class TemplateType>
369 {
370  static const bool value = false;
371 };
372 template<typename... Args, template<typename...> class TemplateType>
373 struct is_specialization<TemplateType<Args...>, TemplateType>
374 {
375  static const bool value = true;
376 };
377 
378 // --------------------------> Value Map <------------------------------------
379 
380 /// @brief Maps one type (e.g. the build types above) to other (actual) types
381 template <typename T>
383 {
384  using Type = T;
385  using type = T;
386 };
387 
388 template<>
390 {
391  using Type = bool;
392  using type = bool;
393 };
394 
395 template<>
397 {
398  using Type = float;
399  using type = float;
400 };
401 
402 template<>
404 {
405  using Type = float;
406  using type = float;
407 };
408 
409 template<>
411 {
412  using Type = float;
413  using type = float;
414 };
415 
416 template<>
418 {
419  using Type = float;
420  using type = float;
421 };
422 
423 template<>
425 {
426  using Type = float;
427  using type = float;
428 };
429 
430 // --------------------------> PtrDiff PtrAdd <------------------------------------
431 
432 template <typename T1, typename T2>
433 __hostdev__ inline static int64_t PtrDiff(const T1* p, const T2* q)
434 {
435  NANOVDB_ASSERT(p && q);
436  return reinterpret_cast<const char*>(p) - reinterpret_cast<const char*>(q);
437 }
438 
439 template <typename DstT, typename SrcT>
440 __hostdev__ inline static DstT* PtrAdd(SrcT *p, int64_t offset)
441 {
442  NANOVDB_ASSERT(p);
443  return reinterpret_cast<DstT*>(reinterpret_cast<char*>(p) + offset);
444 }
445 
446 template <typename DstT, typename SrcT>
447 __hostdev__ inline static const DstT* PtrAdd(const SrcT *p, int64_t offset)
448 {
449  NANOVDB_ASSERT(p);
450  return reinterpret_cast<const DstT*>(reinterpret_cast<const char*>(p) + offset);
451 }
452 // --------------------------> Rgba8 <------------------------------------
453 
454 /// @brief 8-bit red, green, blue, alpha packed into 32 bit unsigned int
455 class Rgba8
456 {
457  union {
458  uint8_t c[4];// 4 color channels of red, green, blue and alpha components.
459  uint32_t packed;// 32 bit packed representation
460  } mData;
461 public:
462  static const int SIZE = 4;
463  using ValueType = uint8_t;
464 
465  Rgba8(const Rgba8&) = default;
466  Rgba8(Rgba8&&) = default;
467  Rgba8& operator=(Rgba8&&) = default;
468  Rgba8& operator=(const Rgba8&) = default;
469  __hostdev__ Rgba8() : mData{0,0,0,0} {static_assert(sizeof(uint32_t) == sizeof(Rgba8),"Unexpected sizeof");}
470  __hostdev__ Rgba8(uint8_t r, uint8_t g, uint8_t b, uint8_t a = 255u) : mData{r, g, b, a} {}
471  explicit __hostdev__ Rgba8(uint8_t v) : Rgba8(v,v,v,v) {}
472  __hostdev__ Rgba8(float r, float g, float b, float a = 1.0f)
473  : mData{(uint8_t(0.5f + r * 255.0f)),// round to nearest
474  (uint8_t(0.5f + g * 255.0f)),// round to nearest
475  (uint8_t(0.5f + b * 255.0f)),// round to nearest
476  (uint8_t(0.5f + a * 255.0f))}// round to nearest
477  {
478  }
479  __hostdev__ bool operator<(const Rgba8& rhs) const { return mData.packed < rhs.mData.packed; }
480  __hostdev__ bool operator==(const Rgba8& rhs) const { return mData.packed == rhs.mData.packed; }
481  __hostdev__ float lengthSqr() const
482  {
483  return 0.0000153787005f*(float(mData.c[0])*mData.c[0] +
484  float(mData.c[1])*mData.c[1] +
485  float(mData.c[2])*mData.c[2]);//1/255^2
486  }
487  __hostdev__ float length() const { return sqrtf(this->lengthSqr() ); }
488  __hostdev__ const uint8_t& operator[](int n) const { return mData.c[n]; }
489  __hostdev__ uint8_t& operator[](int n) { return mData.c[n]; }
490  __hostdev__ const uint32_t& packed() const { return mData.packed; }
491  __hostdev__ uint32_t& packed() { return mData.packed; }
492  __hostdev__ const uint8_t& r() const { return mData.c[0]; }
493  __hostdev__ const uint8_t& g() const { return mData.c[1]; }
494  __hostdev__ const uint8_t& b() const { return mData.c[2]; }
495  __hostdev__ const uint8_t& a() const { return mData.c[3]; }
496  __hostdev__ uint8_t& r() { return mData.c[0]; }
497  __hostdev__ uint8_t& g() { return mData.c[1]; }
498  __hostdev__ uint8_t& b() { return mData.c[2]; }
499  __hostdev__ uint8_t& a() { return mData.c[3]; }
500 };// Rgba8
501 
502 using PackedRGBA8 = Rgba8;// for backwards compatibility
503 
504 // --------------------------> isValue(GridType, GridClass) <------------------------------------
505 
506 /// @brief return true if the GridType maps to a floating point value.
507 __hostdev__ inline bool isFloatingPoint(GridType gridType)
508 {
509  return gridType == GridType::Float ||
510  gridType == GridType::Double ||
511  gridType == GridType::Fp4 ||
512  gridType == GridType::Fp8 ||
513  gridType == GridType::Fp16 ||
514  gridType == GridType::FpN;
515 }
516 
517 // --------------------------> isValue(GridType, GridClass) <------------------------------------
518 
519 /// @brief return true if the combination of GridType and GridClass is valid.
520 __hostdev__ inline bool isValid(GridType gridType, GridClass gridClass)
521 {
522  if (gridClass == GridClass::LevelSet || gridClass == GridClass::FogVolume) {
523  return isFloatingPoint(gridType);
524  } else if (gridClass == GridClass::Staggered) {
525  return gridType == GridType::Vec3f || gridType == GridType::Vec3d ||
526  gridType == GridType::Vec4f || gridType == GridType::Vec4d;
527  } else if (gridClass == GridClass::PointIndex || gridClass == GridClass::PointData) {
528  return gridType == GridType::UInt32;
529  } else if (gridClass == GridClass::VoxelVolume) {
530  return gridType == GridType::RGBA8 || gridType == GridType::Float || gridType == GridType::Double || gridType == GridType::Vec3f || gridType == GridType::Vec3d || gridType == GridType::UInt32;
531  }
532  return gridClass < GridClass::End && gridType < GridType::End;// any valid combination
533 }
534 
535 // ----------------------------> Version class <-------------------------------------
536 
537 /// @brief Bit-compacted representation of all three version numbers
538 ///
539 /// @details major is the top 11 bits, minor is the 11 middle bits and patch is the lower 10 bits
540 class Version
541 {
542  uint32_t mData;// 11 + 11 + 10 bit packing of major + minor + patch
543 public:
544  __hostdev__ Version() : mData( uint32_t(NANOVDB_MAJOR_VERSION_NUMBER) << 21 |
545  uint32_t(NANOVDB_MINOR_VERSION_NUMBER) << 10 |
546  uint32_t(NANOVDB_PATCH_VERSION_NUMBER) )
547  {
548  }
549  __hostdev__ Version(uint32_t major, uint32_t minor, uint32_t patch)
550  : mData( major << 21 | minor << 10 | patch )
551  {
552  NANOVDB_ASSERT(major < (1u << 11));// max value of major is 2047
553  NANOVDB_ASSERT(minor < (1u << 11));// max value of minor is 2047
554  NANOVDB_ASSERT(patch < (1u << 10));// max value of patch is 1023
555  }
556  __hostdev__ bool operator==(const Version &rhs) const {return mData == rhs.mData;}
557  __hostdev__ bool operator< (const Version &rhs) const {return mData < rhs.mData;}
558  __hostdev__ bool operator<=(const Version &rhs) const {return mData <= rhs.mData;}
559  __hostdev__ bool operator> (const Version &rhs) const {return mData > rhs.mData;}
560  __hostdev__ bool operator>=(const Version &rhs) const {return mData >= rhs.mData;}
561  __hostdev__ uint32_t id() const { return mData; }
562  __hostdev__ uint32_t getMajor() const { return (mData >> 21) & ((1u << 11) - 1);}
563  __hostdev__ uint32_t getMinor() const { return (mData >> 10) & ((1u << 11) - 1);}
564  __hostdev__ uint32_t getPatch() const { return mData & ((1u << 10) - 1);}
565 
566 #ifndef __CUDACC_RTC__
567  const char* c_str() const
568  {
569  char *buffer = (char*)malloc(4 + 1 + 4 + 1 + 4 + 1);// xxxx.xxxx.xxxx\n
570  sprintf(buffer, "%d.%d.%d", this->getMajor(), this->getMinor(), this->getPatch());
571  return buffer;
572  }
573 #endif
574 };// Version
575 
576 // ----------------------------> Various math functions <-------------------------------------
577 
578 //@{
579 /// Tolerance for floating-point comparison
580 template<typename T>
581 struct Tolerance;
582 template<>
583 struct Tolerance<float>
584 {
585  __hostdev__ static float value() { return 1e-8f; }
586 };
587 template<>
588 struct Tolerance<double>
589 {
590  __hostdev__ static double value() { return 1e-15; }
591 };
592 //@}
593 
594 //@{
595 /// Delta for small floating-point offsets
596 template<typename T>
597 struct Delta;
598 template<>
599 struct Delta<float>
600 {
601  __hostdev__ static float value() { return 1e-5f; }
602 };
603 template<>
604 struct Delta<double>
605 {
606  __hostdev__ static double value() { return 1e-9; }
607 };
608 //@}
609 
610 //@{
611 /// Maximum floating-point values
612 template<typename T>
613 struct Maximum;
614 #if defined(__CUDA_ARCH__) || defined(__HIP__)
615 template<>
616 struct Maximum<int>
617 {
618  __hostdev__ static int value() { return 2147483647; }
619 };
620 template<>
621 struct Maximum<uint32_t>
622 {
623  __hostdev__ static uint32_t value() { return 4294967295; }
624 };
625 template<>
626 struct Maximum<float>
627 {
628  __hostdev__ static float value() { return 1e+38f; }
629 };
630 template<>
631 struct Maximum<double>
632 {
633  __hostdev__ static double value() { return 1e+308; }
634 };
635 #else
636 template<typename T>
637 struct Maximum
638 {
639  static T value() { return std::numeric_limits<T>::max(); }
640 };
641 #endif
642 //@}
643 
644 template<typename Type>
645 __hostdev__ inline bool isApproxZero(const Type& x)
646 {
647  return !(x > Tolerance<Type>::value()) && !(x < -Tolerance<Type>::value());
648 }
649 
650 template<typename Type>
651 __hostdev__ inline Type Min(Type a, Type b)
652 {
653  return (a < b) ? a : b;
654 }
655 __hostdev__ inline int32_t Min(int32_t a, int32_t b)
656 {
657  return int32_t(fminf(float(a), float(b)));
658 }
659 __hostdev__ inline uint32_t Min(uint32_t a, uint32_t b)
660 {
661  return uint32_t(fminf(float(a), float(b)));
662 }
663 __hostdev__ inline float Min(float a, float b)
664 {
665  return fminf(a, b);
666 }
667 __hostdev__ inline double Min(double a, double b)
668 {
669  return fmin(a, b);
670 }
671 template<typename Type>
672 __hostdev__ inline Type Max(Type a, Type b)
673 {
674  return (a > b) ? a : b;
675 }
676 
677 __hostdev__ inline int32_t Max(int32_t a, int32_t b)
678 {
679  return int32_t(fmaxf(float(a), float(b)));
680 }
681 __hostdev__ inline uint32_t Max(uint32_t a, uint32_t b)
682 {
683  return uint32_t(fmaxf(float(a), float(b)));
684 }
685 __hostdev__ inline float Max(float a, float b)
686 {
687  return fmaxf(a, b);
688 }
689 __hostdev__ inline double Max(double a, double b)
690 {
691  return fmax(a, b);
692 }
693 __hostdev__ inline float Clamp(float x, float a, float b)
694 {
695  return Max(Min(x, b), a);
696 }
697 __hostdev__ inline double Clamp(double x, double a, double b)
698 {
699  return Max(Min(x, b), a);
700 }
701 
702 __hostdev__ inline float Fract(float x)
703 {
704  return x - floorf(x);
705 }
706 __hostdev__ inline double Fract(double x)
707 {
708  return x - floor(x);
709 }
710 
711 __hostdev__ inline int32_t Floor(float x)
712 {
713  return int32_t(floorf(x));
714 }
715 __hostdev__ inline int32_t Floor(double x)
716 {
717  return int32_t(floor(x));
718 }
719 
720 __hostdev__ inline int32_t Ceil(float x)
721 {
722  return int32_t(ceilf(x));
723 }
724 __hostdev__ inline int32_t Ceil(double x)
725 {
726  return int32_t(ceil(x));
727 }
728 
729 template<typename T>
730 __hostdev__ inline T Pow2(T x)
731 {
732  return x * x;
733 }
734 
735 template<typename T>
736 __hostdev__ inline T Pow3(T x)
737 {
738  return x * x * x;
739 }
740 
741 template<typename T>
742 __hostdev__ inline T Pow4(T x)
743 {
744  return Pow2(x * x);
745 }
746 template<typename T>
747 __hostdev__ inline T Abs(T x)
748 {
749  return x < 0 ? -x : x;
750 }
751 
752 template<>
753 __hostdev__ inline float Abs(float x)
754 {
755  return fabs(x);
756 }
757 
758 template<>
759 __hostdev__ inline double Abs(double x)
760 {
761  return fabs(x);
762 }
763 
764 template<>
765 __hostdev__ inline int Abs(int x)
766 {
767  return abs(x);
768 }
769 
770 template<typename CoordT, typename RealT, template<typename> class Vec3T>
771 __hostdev__ inline CoordT Round(const Vec3T<RealT>& xyz);
772 
773 template<typename CoordT, template<typename> class Vec3T>
774 __hostdev__ inline CoordT Round(const Vec3T<float>& xyz)
775 {
776  return CoordT(int32_t(rintf(xyz[0])), int32_t(rintf(xyz[1])), int32_t(rintf(xyz[2])));
777  //return CoordT(int32_t(roundf(xyz[0])), int32_t(roundf(xyz[1])), int32_t(roundf(xyz[2])) );
778  //return CoordT(int32_t(floorf(xyz[0] + 0.5f)), int32_t(floorf(xyz[1] + 0.5f)), int32_t(floorf(xyz[2] + 0.5f)));
779 }
780 
781 template<typename CoordT, template<typename> class Vec3T>
782 __hostdev__ inline CoordT Round(const Vec3T<double>& xyz)
783 {
784  return CoordT(int32_t(floor(xyz[0] + 0.5)), int32_t(floor(xyz[1] + 0.5)), int32_t(floor(xyz[2] + 0.5)));
785 }
786 
787 template<typename CoordT, typename RealT, template<typename> class Vec3T>
788 __hostdev__ inline CoordT RoundDown(const Vec3T<RealT>& xyz)
789 {
790  return CoordT(Floor(xyz[0]), Floor(xyz[1]), Floor(xyz[2]));
791 }
792 
793 //@{
794 /// Return the square root of a floating-point value.
795 __hostdev__ inline float Sqrt(float x)
796 {
797  return sqrtf(x);
798 }
799 __hostdev__ inline double Sqrt(double x)
800 {
801  return sqrt(x);
802 }
803 //@}
804 
805 /// Return the sign of the given value as an integer (either -1, 0 or 1).
806 template <typename T>
807 __hostdev__ inline T Sign(const T &x) { return ((T(0) < x)?T(1):T(0)) - ((x < T(0))?T(1):T(0)); }
808 
809 template<typename Vec3T>
810 __hostdev__ inline int MinIndex(const Vec3T& v)
811 {
812 #if 0
813  static const int hashTable[8] = {2, 1, 9, 1, 2, 9, 0, 0}; //9 are dummy values
814  const int hashKey = ((v[0] < v[1]) << 2) + ((v[0] < v[2]) << 1) + (v[1] < v[2]); // ?*4+?*2+?*1
815  return hashTable[hashKey];
816 #else
817  if (v[0] < v[1] && v[0] < v[2])
818  return 0;
819  if (v[1] < v[2])
820  return 1;
821  else
822  return 2;
823 #endif
824 }
825 
826 template<typename Vec3T>
827 __hostdev__ inline int MaxIndex(const Vec3T& v)
828 {
829 #if 0
830  static const int hashTable[8] = {2, 1, 9, 1, 2, 9, 0, 0}; //9 are dummy values
831  const int hashKey = ((v[0] > v[1]) << 2) + ((v[0] > v[2]) << 1) + (v[1] > v[2]); // ?*4+?*2+?*1
832  return hashTable[hashKey];
833 #else
834  if (v[0] > v[1] && v[0] > v[2])
835  return 0;
836  if (v[1] > v[2])
837  return 1;
838  else
839  return 2;
840 #endif
841 }
842 
843 /// @brief round up byteSize to the nearest wordSize, e.g. to align to machine word: AlignUp<sizeof(size_t)(n)
844 ///
845 /// @details both wordSize and byteSize are in byte units
846 template<uint64_t wordSize>
847 __hostdev__ inline uint64_t AlignUp(uint64_t byteCount)
848 {
849  const uint64_t r = byteCount % wordSize;
850  return r ? byteCount - r + wordSize : byteCount;
851 }
852 
853 // ------------------------------> Coord <--------------------------------------
854 
855 // forward decleration so we can define Coord::asVec3s and Coord::asVec3d
856 template<typename> class Vec3;
857 
858 /// @brief Signed (i, j, k) 32-bit integer coordinate class, similar to openvdb::math::Coord
859 class Coord
860 {
861  int32_t mVec[3]; // private member data - three signed index coordinates
862 public:
863  using ValueType = int32_t;
864  using IndexType = uint32_t;
865 
866  /// @brief Initialize all coordinates to zero.
868  : mVec{0, 0, 0}
869  {
870  }
871 
872  /// @brief Initializes all coordinates to the given signed integer.
874  : mVec{n, n, n}
875  {
876  }
877 
878  /// @brief Initializes coordinate to the given signed integers.
880  : mVec{i, j, k}
881  {
882  }
883 
885  : mVec{ptr[0], ptr[1], ptr[2]}
886  {
887  }
888 
889  __hostdev__ int32_t x() const { return mVec[0]; }
890  __hostdev__ int32_t y() const { return mVec[1]; }
891  __hostdev__ int32_t z() const { return mVec[2]; }
892 
893  __hostdev__ int32_t& x() { return mVec[0]; }
894  __hostdev__ int32_t& y() { return mVec[1]; }
895  __hostdev__ int32_t& z() { return mVec[2]; }
896 
897  __hostdev__ static Coord max() { return Coord(int32_t((1u << 31) - 1)); }
898 
899  __hostdev__ static Coord min() { return Coord(-int32_t((1u << 31) - 1) - 1); }
900 
901  __hostdev__ static size_t memUsage() { return sizeof(Coord); }
902 
903  /// @brief Return a const reference to the given Coord component.
904  /// @warning The argument is assumed to be 0, 1, or 2.
905  __hostdev__ const ValueType& operator[](IndexType i) const { return mVec[i]; }
906 
907  /// @brief Return a non-const reference to the given Coord component.
908  /// @warning The argument is assumed to be 0, 1, or 2.
909  __hostdev__ ValueType& operator[](IndexType i) { return mVec[i]; }
910 
911  /// @brief Assignment operator that works with openvdb::Coord
912  template <typename CoordT>
913  __hostdev__ Coord& operator=(const CoordT &other)
914  {
915  static_assert(sizeof(Coord) == sizeof(CoordT), "Mis-matched sizeof");
916  mVec[0] = other[0];
917  mVec[1] = other[1];
918  mVec[2] = other[2];
919  return *this;
920  }
921 
922  /// @brief Return a new instance with coordinates masked by the given unsigned integer.
923  __hostdev__ Coord operator&(IndexType n) const { return Coord(mVec[0] & n, mVec[1] & n, mVec[2] & n); }
924 
925  // @brief Return a new instance with coordinates left-shifted by the given unsigned integer.
926  __hostdev__ Coord operator<<(IndexType n) const { return Coord(mVec[0] << n, mVec[1] << n, mVec[2] << n); }
927 
928  // @brief Return a new instance with coordinates right-shifted by the given unsigned integer.
929  __hostdev__ Coord operator>>(IndexType n) const { return Coord(mVec[0] >> n, mVec[1] >> n, mVec[2] >> n); }
930 
931  /// @brief Return true if this Coord is lexicographically less than the given Coord.
932  __hostdev__ bool operator<(const Coord& rhs) const
933  {
934  return mVec[0] < rhs[0] ? true : mVec[0] > rhs[0] ? false : mVec[1] < rhs[1] ? true : mVec[1] > rhs[1] ? false : mVec[2] < rhs[2] ? true : false;
935  }
936 
937  // @brief Return true if the Coord components are identical.
938  __hostdev__ bool operator==(const Coord& rhs) const { return mVec[0] == rhs[0] && mVec[1] == rhs[1] && mVec[2] == rhs[2]; }
939  __hostdev__ bool operator!=(const Coord& rhs) const { return mVec[0] != rhs[0] || mVec[1] != rhs[1] || mVec[2] != rhs[2]; }
941  {
942  mVec[0] &= n;
943  mVec[1] &= n;
944  mVec[2] &= n;
945  return *this;
946  }
948  {
949  mVec[0] <<= n;
950  mVec[1] <<= n;
951  mVec[2] <<= n;
952  return *this;
953  }
955  {
956  mVec[0] += n;
957  mVec[1] += n;
958  mVec[2] += n;
959  return *this;
960  }
961  __hostdev__ Coord operator+(const Coord& rhs) const { return Coord(mVec[0] + rhs[0], mVec[1] + rhs[1], mVec[2] + rhs[2]); }
962  __hostdev__ Coord operator-(const Coord& rhs) const { return Coord(mVec[0] - rhs[0], mVec[1] - rhs[1], mVec[2] - rhs[2]); }
964  {
965  mVec[0] += rhs[0];
966  mVec[1] += rhs[1];
967  mVec[2] += rhs[2];
968  return *this;
969  }
971  {
972  mVec[0] -= rhs[0];
973  mVec[1] -= rhs[1];
974  mVec[2] -= rhs[2];
975  return *this;
976  }
977 
978  /// @brief Perform a component-wise minimum with the other Coord.
980  {
981  if (other[0] < mVec[0])
982  mVec[0] = other[0];
983  if (other[1] < mVec[1])
984  mVec[1] = other[1];
985  if (other[2] < mVec[2])
986  mVec[2] = other[2];
987  return *this;
988  }
989 
990  /// @brief Perform a component-wise maximum with the other Coord.
992  {
993  if (other[0] > mVec[0])
994  mVec[0] = other[0];
995  if (other[1] > mVec[1])
996  mVec[1] = other[1];
997  if (other[2] > mVec[2])
998  mVec[2] = other[2];
999  return *this;
1000  }
1001 
1003  {
1004  return Coord(mVec[0] + dx, mVec[1] + dy, mVec[2] + dz);
1005  }
1006 
1007  __hostdev__ Coord offsetBy(ValueType n) const { return this->offsetBy(n, n, n); }
1008 
1009  /// Return true if any of the components of @a a are smaller than the
1010  /// corresponding components of @a b.
1011  __hostdev__ static inline bool lessThan(const Coord& a, const Coord& b)
1012  {
1013  return (a[0] < b[0] || a[1] < b[1] || a[2] < b[2]);
1014  }
1015 
1016  /// @brief Return the largest integer coordinates that are not greater
1017  /// than @a xyz (node centered conversion).
1018  template<typename Vec3T>
1019  __hostdev__ static Coord Floor(const Vec3T& xyz) { return Coord(nanovdb::Floor(xyz[0]), nanovdb::Floor(xyz[1]), nanovdb::Floor(xyz[2])); }
1020 
1021  /// @brief Return a hash key derived from the existing coordinates.
1022  /// @details For details on this hash function please see the VDB paper.
1023  template<int Log2N = 3 + 4 + 5>
1024  __hostdev__ uint32_t hash() const { return ((1 << Log2N) - 1) & (mVec[0] * 73856093 ^ mVec[1] * 19349663 ^ mVec[2] * 83492791); }
1025 
1026  /// @brief Return the octant of this Coord
1027  //__hostdev__ size_t octant() const { return (uint32_t(mVec[0])>>31) | ((uint32_t(mVec[1])>>31)<<1) | ((uint32_t(mVec[2])>>31)<<2); }
1028  __hostdev__ uint8_t octant() const { return (uint8_t(bool(mVec[0] & (1u << 31)))) |
1029  (uint8_t(bool(mVec[1] & (1u << 31))) << 1) |
1030  (uint8_t(bool(mVec[2] & (1u << 31))) << 2); }
1031 
1032  /// @brief Return a single precision floating-point vector of this coordinate
1033  __hostdev__ inline Vec3<float> asVec3s() const;
1034 
1035  /// @brief Return a double precision floating-point vector of this coordinate
1036  __hostdev__ inline Vec3<double> asVec3d() const;
1037 }; // Coord class
1038 
1039 // ----------------------------> Vec3 <--------------------------------------
1040 
1041 /// @brief A simple vector class with three double components, similar to openvdb::math::Vec3
1042 template<typename T>
1043 class Vec3
1044 {
1045  T mVec[3];
1046 
1047 public:
1048  static const int SIZE = 3;
1049  using ValueType = T;
1050  Vec3() = default;
1051  __hostdev__ explicit Vec3(T x)
1052  : mVec{x, x, x}
1053  {
1054  }
1055  __hostdev__ Vec3(T x, T y, T z)
1056  : mVec{x, y, z}
1057  {
1058  }
1059  template<typename T2>
1060  __hostdev__ explicit Vec3(const Vec3<T2>& v)
1061  : mVec{T(v[0]), T(v[1]), T(v[2])}
1062  {
1063  }
1064  __hostdev__ explicit Vec3(const Coord& ijk)
1065  : mVec{T(ijk[0]), T(ijk[1]), T(ijk[2])}
1066  {
1067  }
1068  __hostdev__ bool operator==(const Vec3& rhs) const { return mVec[0] == rhs[0] && mVec[1] == rhs[1] && mVec[2] == rhs[2]; }
1069  __hostdev__ bool operator!=(const Vec3& rhs) const { return mVec[0] != rhs[0] || mVec[1] != rhs[1] || mVec[2] != rhs[2]; }
1070  template<typename Vec3T>
1071  __hostdev__ Vec3& operator=(const Vec3T& rhs)
1072  {
1073  mVec[0] = rhs[0];
1074  mVec[1] = rhs[1];
1075  mVec[2] = rhs[2];
1076  return *this;
1077  }
1078  __hostdev__ const T& operator[](int i) const { return mVec[i]; }
1079  __hostdev__ T& operator[](int i) { return mVec[i]; }
1080  template<typename Vec3T>
1081  __hostdev__ T dot(const Vec3T& v) const { return mVec[0] * v[0] + mVec[1] * v[1] + mVec[2] * v[2]; }
1082  template<typename Vec3T>
1083  __hostdev__ Vec3 cross(const Vec3T& v) const
1084  {
1085  return Vec3(mVec[1] * v[2] - mVec[2] * v[1],
1086  mVec[2] * v[0] - mVec[0] * v[2],
1087  mVec[0] * v[1] - mVec[1] * v[0]);
1088  }
1090  {
1091  return mVec[0] * mVec[0] + mVec[1] * mVec[1] + mVec[2] * mVec[2]; // 5 flops
1092  }
1093  __hostdev__ T length() const { return Sqrt(this->lengthSqr()); }
1094  __hostdev__ Vec3 operator-() const { return Vec3(-mVec[0], -mVec[1], -mVec[2]); }
1095  __hostdev__ Vec3 operator*(const Vec3& v) const { return Vec3(mVec[0] * v[0], mVec[1] * v[1], mVec[2] * v[2]); }
1096  __hostdev__ Vec3 operator/(const Vec3& v) const { return Vec3(mVec[0] / v[0], mVec[1] / v[1], mVec[2] / v[2]); }
1097  __hostdev__ Vec3 operator+(const Vec3& v) const { return Vec3(mVec[0] + v[0], mVec[1] + v[1], mVec[2] + v[2]); }
1098  __hostdev__ Vec3 operator-(const Vec3& v) const { return Vec3(mVec[0] - v[0], mVec[1] - v[1], mVec[2] - v[2]); }
1099  __hostdev__ Vec3 operator*(const T& s) const { return Vec3(s * mVec[0], s * mVec[1], s * mVec[2]); }
1100  __hostdev__ Vec3 operator/(const T& s) const { return (T(1) / s) * (*this); }
1102  {
1103  mVec[0] += v[0];
1104  mVec[1] += v[1];
1105  mVec[2] += v[2];
1106  return *this;
1107  }
1109  {
1110  mVec[0] -= v[0];
1111  mVec[1] -= v[1];
1112  mVec[2] -= v[2];
1113  return *this;
1114  }
1116  {
1117  mVec[0] *= s;
1118  mVec[1] *= s;
1119  mVec[2] *= s;
1120  return *this;
1121  }
1122  __hostdev__ Vec3& operator/=(const T& s) { return (*this) *= T(1) / s; }
1123  __hostdev__ Vec3& normalize() { return (*this) /= this->length(); }
1124  /// @brief Perform a component-wise minimum with the other Coord.
1126  {
1127  if (other[0] < mVec[0])
1128  mVec[0] = other[0];
1129  if (other[1] < mVec[1])
1130  mVec[1] = other[1];
1131  if (other[2] < mVec[2])
1132  mVec[2] = other[2];
1133  return *this;
1134  }
1135 
1136  /// @brief Perform a component-wise maximum with the other Coord.
1138  {
1139  if (other[0] > mVec[0])
1140  mVec[0] = other[0];
1141  if (other[1] > mVec[1])
1142  mVec[1] = other[1];
1143  if (other[2] > mVec[2])
1144  mVec[2] = other[2];
1145  return *this;
1146  }
1147  /// @brief Return the smallest vector component
1149  {
1150  return mVec[0] < mVec[1] ? (mVec[0] < mVec[2] ? mVec[0] : mVec[2]) : (mVec[1] < mVec[2] ? mVec[1] : mVec[2]);
1151  }
1152  /// @brief Return the largest vector component
1154  {
1155  return mVec[0] > mVec[1] ? (mVec[0] > mVec[2] ? mVec[0] : mVec[2]) : (mVec[1] > mVec[2] ? mVec[1] : mVec[2]);
1156  }
1157  __hostdev__ Coord floor() const { return Coord(Floor(mVec[0]), Floor(mVec[1]), Floor(mVec[2])); }
1158  __hostdev__ Coord ceil() const { return Coord(Ceil(mVec[0]), Ceil(mVec[1]), Ceil(mVec[2])); }
1159  __hostdev__ Coord round() const { return Coord(Floor(mVec[0] + 0.5), Floor(mVec[1] + 0.5), Floor(mVec[2] + 0.5)); }
1160 }; // Vec3<T>
1161 
1162 template<typename T1, typename T2>
1163 __hostdev__ inline Vec3<T2> operator*(T1 scalar, const Vec3<T2>& vec)
1164 {
1165  return Vec3<T2>(scalar * vec[0], scalar * vec[1], scalar * vec[2]);
1166 }
1167 template<typename T1, typename T2>
1168 __hostdev__ inline Vec3<T2> operator/(T1 scalar, const Vec3<T2>& vec)
1169 {
1170  return Vec3<T2>(scalar / vec[0], scalar / vec[1], scalar / vec[2]);
1171 }
1172 
1177 
1178 /// @brief Return a single precision floating-point vector of this coordinate
1179 __hostdev__ inline Vec3f Coord::asVec3s() const { return Vec3f(float(mVec[0]), float(mVec[1]), float(mVec[2])); }
1180 
1181 /// @brief Return a double precision floating-point vector of this coordinate
1182 __hostdev__ inline Vec3d Coord::asVec3d() const { return Vec3d(double(mVec[0]), double(mVec[1]), double(mVec[2])); }
1183 
1184 // ----------------------------> Vec4 <--------------------------------------
1185 
1186 /// @brief A simple vector class with three double components, similar to openvdb::math::Vec4
1187 template<typename T>
1188 class Vec4
1189 {
1190  T mVec[4];
1191 
1192 public:
1193  static const int SIZE = 4;
1194  using ValueType = T;
1195  Vec4() = default;
1196  __hostdev__ explicit Vec4(T x)
1197  : mVec{x, x, x, x}
1198  {
1199  }
1200  __hostdev__ Vec4(T x, T y, T z, T w)
1201  : mVec{x, y, z, w}
1202  {
1203  }
1204  template<typename T2>
1205  __hostdev__ explicit Vec4(const Vec4<T2>& v)
1206  : mVec{T(v[0]), T(v[1]), T(v[2]), T(v[3])}
1207  {
1208  }
1209  __hostdev__ bool operator==(const Vec4& rhs) const { return mVec[0] == rhs[0] && mVec[1] == rhs[1] && mVec[2] == rhs[2] && mVec[3] == rhs[3]; }
1210  __hostdev__ bool operator!=(const Vec4& rhs) const { return mVec[0] != rhs[0] || mVec[1] != rhs[1] || mVec[2] != rhs[2] || mVec[3] != rhs[3]; }
1211  template<typename Vec4T>
1212  __hostdev__ Vec4& operator=(const Vec4T& rhs)
1213  {
1214  mVec[0] = rhs[0];
1215  mVec[1] = rhs[1];
1216  mVec[2] = rhs[2];
1217  mVec[3] = rhs[3];
1218  return *this;
1219  }
1220  __hostdev__ const T& operator[](int i) const { return mVec[i]; }
1221  __hostdev__ T& operator[](int i) { return mVec[i]; }
1222  template<typename Vec4T>
1223  __hostdev__ T dot(const Vec4T& v) const { return mVec[0] * v[0] + mVec[1] * v[1] + mVec[2] * v[2] + mVec[3] * v[3]; }
1225  {
1226  return mVec[0] * mVec[0] + mVec[1] * mVec[1] + mVec[2] * mVec[2] + mVec[3] * mVec[3]; // 7 flops
1227  }
1228  __hostdev__ T length() const { return Sqrt(this->lengthSqr()); }
1229  __hostdev__ Vec4 operator-() const { return Vec4(-mVec[0], -mVec[1], -mVec[2], -mVec[3]); }
1230  __hostdev__ Vec4 operator*(const Vec4& v) const { return Vec4(mVec[0] * v[0], mVec[1] * v[1], mVec[2] * v[2], mVec[3] * v[3]); }
1231  __hostdev__ Vec4 operator/(const Vec4& v) const { return Vec4(mVec[0] / v[0], mVec[1] / v[1], mVec[2] / v[2], mVec[3] / v[3]); }
1232  __hostdev__ Vec4 operator+(const Vec4& v) const { return Vec4(mVec[0] + v[0], mVec[1] + v[1], mVec[2] + v[2], mVec[3] + v[3]); }
1233  __hostdev__ Vec4 operator-(const Vec4& v) const { return Vec4(mVec[0] - v[0], mVec[1] - v[1], mVec[2] - v[2], mVec[3] - v[3]); }
1234  __hostdev__ Vec4 operator*(const T& s) const { return Vec4(s * mVec[0], s * mVec[1], s * mVec[2], s * mVec[3]); }
1235  __hostdev__ Vec4 operator/(const T& s) const { return (T(1) / s) * (*this); }
1237  {
1238  mVec[0] += v[0];
1239  mVec[1] += v[1];
1240  mVec[2] += v[2];
1241  mVec[3] += v[3];
1242  return *this;
1243  }
1245  {
1246  mVec[0] -= v[0];
1247  mVec[1] -= v[1];
1248  mVec[2] -= v[2];
1249  mVec[3] -= v[3];
1250  return *this;
1251  }
1253  {
1254  mVec[0] *= s;
1255  mVec[1] *= s;
1256  mVec[2] *= s;
1257  mVec[3] *= s;
1258  return *this;
1259  }
1260  __hostdev__ Vec4& operator/=(const T& s) { return (*this) *= T(1) / s; }
1261  __hostdev__ Vec4& normalize() { return (*this) /= this->length(); }
1262  /// @brief Perform a component-wise minimum with the other Coord.
1264  {
1265  if (other[0] < mVec[0])
1266  mVec[0] = other[0];
1267  if (other[1] < mVec[1])
1268  mVec[1] = other[1];
1269  if (other[2] < mVec[2])
1270  mVec[2] = other[2];
1271  if (other[3] < mVec[3])
1272  mVec[3] = other[3];
1273  return *this;
1274  }
1275 
1276  /// @brief Perform a component-wise maximum with the other Coord.
1278  {
1279  if (other[0] > mVec[0])
1280  mVec[0] = other[0];
1281  if (other[1] > mVec[1])
1282  mVec[1] = other[1];
1283  if (other[2] > mVec[2])
1284  mVec[2] = other[2];
1285  if (other[3] > mVec[3])
1286  mVec[3] = other[3];
1287  return *this;
1288  }
1289 }; // Vec4<T>
1290 
1291 template<typename T1, typename T2>
1292 __hostdev__ inline Vec4<T2> operator*(T1 scalar, const Vec4<T2>& vec)
1293 {
1294  return Vec4<T2>(scalar * vec[0], scalar * vec[1], scalar * vec[2], scalar * vec[3]);
1295 }
1296 template<typename T1, typename T2>
1297 __hostdev__ inline Vec4<T2> operator/(T1 scalar, const Vec3<T2>& vec)
1298 {
1299  return Vec4<T2>(scalar / vec[0], scalar / vec[1], scalar / vec[2], scalar / vec[3]);
1300 }
1301 
1306 
1307 // ----------------------------> TensorTraits <--------------------------------------
1308 
1309 template<typename T, int Rank = (is_specialization<T, Vec3>::value ||
1310  is_specialization<T, Vec4>::value ||
1311  is_same<T, Rgba8>::value) ? 1 : 0>
1313 
1314 template<typename T>
1315 struct TensorTraits<T, 0>
1316 {
1317  static const int Rank = 0; // i.e. scalar
1318  static const bool IsScalar = true;
1319  static const bool IsVector = false;
1320  static const int Size = 1;
1321  using ElementType = T;
1322  static T scalar(const T& s) { return s; }
1323 };
1324 
1325 template<typename T>
1326 struct TensorTraits<T, 1>
1327 {
1328  static const int Rank = 1; // i.e. vector
1329  static const bool IsScalar = false;
1330  static const bool IsVector = true;
1331  static const int Size = T::SIZE;
1332  using ElementType = typename T::ValueType;
1333  static ElementType scalar(const T& v) { return v.length(); }
1334 };
1335 
1336 // ----------------------------> FloatTraits <--------------------------------------
1337 
1338 template<typename T, int = sizeof(typename TensorTraits<T>::ElementType)>
1340 {
1341  using FloatType = float;
1342 };
1343 
1344 template<typename T>
1345 struct FloatTraits<T, 8>
1346 {
1347  using FloatType = double;
1348 };
1349 
1350 template<>
1351 struct FloatTraits<bool, 1>
1352 {
1353  using FloatType = bool;
1354 };
1355 
1356 template<>
1358 {
1359  using FloatType = bool;
1360 };
1361 
1362 // ----------------------------> mapping ValueType -> GridType <--------------------------------------
1363 
1364 /// @brief Maps from a templated value type to a GridType enum
1365 template<typename BuildT>
1367 {
1368  if (is_same<BuildT, float>::value) { // resolved at compile-time
1369  return GridType::Float;
1370  } else if (is_same<BuildT, double>::value) {
1371  return GridType::Double;
1372  } else if (is_same<BuildT, int16_t>::value) {
1373  return GridType::Int16;
1374  } else if (is_same<BuildT, int32_t>::value) {
1375  return GridType::Int32;
1376  } else if (is_same<BuildT, int64_t>::value) {
1377  return GridType::Int64;
1378  } else if (is_same<BuildT, Vec3f>::value) {
1379  return GridType::Vec3f;
1380  } else if (is_same<BuildT, Vec3d>::value) {
1381  return GridType::Vec3d;
1382  } else if (is_same<BuildT, uint32_t>::value) {
1383  return GridType::UInt32;
1384  } else if (is_same<BuildT, ValueMask>::value) {
1385  return GridType::Mask;
1386  } else if (is_same<BuildT, bool>::value) {
1387  return GridType::Boolean;
1388  } else if (is_same<BuildT, Rgba8>::value) {
1389  return GridType::RGBA8;
1390  } else if (is_same<BuildT, Fp4>::value) {
1391  return GridType::Fp4;
1392  } else if (is_same<BuildT, Fp8>::value) {
1393  return GridType::Fp8;
1394  } else if (is_same<BuildT, Fp16>::value) {
1395  return GridType::Fp16;
1396  } else if (is_same<BuildT, FpN>::value) {
1397  return GridType::FpN;
1398  } else if (is_same<BuildT, Vec4f>::value) {
1399  return GridType::Vec4f;
1400  } else if (is_same<BuildT, Vec4d>::value) {
1401  return GridType::Vec4d;
1402  }
1403  return GridType::Unknown;
1404 }
1405 
1406 // ----------------------------> matMult <--------------------------------------
1407 
1408 template<typename Vec3T>
1409 __hostdev__ inline Vec3T matMult(const float* mat, const Vec3T& xyz)
1410 {
1411  return Vec3T(fmaf(xyz[0], mat[0], fmaf(xyz[1], mat[1], xyz[2] * mat[2])),
1412  fmaf(xyz[0], mat[3], fmaf(xyz[1], mat[4], xyz[2] * mat[5])),
1413  fmaf(xyz[0], mat[6], fmaf(xyz[1], mat[7], xyz[2] * mat[8]))); // 6 fmaf + 3 mult = 9 flops
1414 }
1415 
1416 template<typename Vec3T>
1417 __hostdev__ inline Vec3T matMult(const double* mat, const Vec3T& xyz)
1418 {
1419  return Vec3T(fma(static_cast<double>(xyz[0]), mat[0], fma(static_cast<double>(xyz[1]), mat[1], static_cast<double>(xyz[2]) * mat[2])),
1420  fma(static_cast<double>(xyz[0]), mat[3], fma(static_cast<double>(xyz[1]), mat[4], static_cast<double>(xyz[2]) * mat[5])),
1421  fma(static_cast<double>(xyz[0]), mat[6], fma(static_cast<double>(xyz[1]), mat[7], static_cast<double>(xyz[2]) * mat[8]))); // 6 fmaf + 3 mult = 9 flops
1422 }
1423 
1424 template<typename Vec3T>
1425 __hostdev__ inline Vec3T matMult(const float* mat, const float* vec, const Vec3T& xyz)
1426 {
1427  return Vec3T(fmaf(xyz[0], mat[0], fmaf(xyz[1], mat[1], fmaf(xyz[2], mat[2], vec[0]))),
1428  fmaf(xyz[0], mat[3], fmaf(xyz[1], mat[4], fmaf(xyz[2], mat[5], vec[1]))),
1429  fmaf(xyz[0], mat[6], fmaf(xyz[1], mat[7], fmaf(xyz[2], mat[8], vec[2])))); // 9 fmaf = 9 flops
1430 }
1431 
1432 template<typename Vec3T>
1433 __hostdev__ inline Vec3T matMult(const double* mat, const double* vec, const Vec3T& xyz)
1434 {
1435  return Vec3T(fma(static_cast<double>(xyz[0]), mat[0], fma(static_cast<double>(xyz[1]), mat[1], fma(static_cast<double>(xyz[2]), mat[2], vec[0]))),
1436  fma(static_cast<double>(xyz[0]), mat[3], fma(static_cast<double>(xyz[1]), mat[4], fma(static_cast<double>(xyz[2]), mat[5], vec[1]))),
1437  fma(static_cast<double>(xyz[0]), mat[6], fma(static_cast<double>(xyz[1]), mat[7], fma(static_cast<double>(xyz[2]), mat[8], vec[2])))); // 9 fma = 9 flops
1438 }
1439 
1440 // matMultT: Multiply with the transpose:
1441 
1442 template<typename Vec3T>
1443 __hostdev__ inline Vec3T matMultT(const float* mat, const Vec3T& xyz)
1444 {
1445  return Vec3T(fmaf(xyz[0], mat[0], fmaf(xyz[1], mat[3], xyz[2] * mat[6])),
1446  fmaf(xyz[0], mat[1], fmaf(xyz[1], mat[4], xyz[2] * mat[7])),
1447  fmaf(xyz[0], mat[2], fmaf(xyz[1], mat[5], xyz[2] * mat[8]))); // 6 fmaf + 3 mult = 9 flops
1448 }
1449 
1450 template<typename Vec3T>
1451 __hostdev__ inline Vec3T matMultT(const double* mat, const Vec3T& xyz)
1452 {
1453  return Vec3T(fma(static_cast<double>(xyz[0]), mat[0], fma(static_cast<double>(xyz[1]), mat[3], static_cast<double>(xyz[2]) * mat[6])),
1454  fma(static_cast<double>(xyz[0]), mat[1], fma(static_cast<double>(xyz[1]), mat[4], static_cast<double>(xyz[2]) * mat[7])),
1455  fma(static_cast<double>(xyz[0]), mat[2], fma(static_cast<double>(xyz[1]), mat[5], static_cast<double>(xyz[2]) * mat[8]))); // 6 fmaf + 3 mult = 9 flops
1456 }
1457 
1458 template<typename Vec3T>
1459 __hostdev__ inline Vec3T matMultT(const float* mat, const float* vec, const Vec3T& xyz)
1460 {
1461  return Vec3T(fmaf(xyz[0], mat[0], fmaf(xyz[1], mat[3], fmaf(xyz[2], mat[6], vec[0]))),
1462  fmaf(xyz[0], mat[1], fmaf(xyz[1], mat[4], fmaf(xyz[2], mat[7], vec[1]))),
1463  fmaf(xyz[0], mat[2], fmaf(xyz[1], mat[5], fmaf(xyz[2], mat[8], vec[2])))); // 9 fmaf = 9 flops
1464 }
1465 
1466 template<typename Vec3T>
1467 __hostdev__ inline Vec3T matMultT(const double* mat, const double* vec, const Vec3T& xyz)
1468 {
1469  return Vec3T(fma(static_cast<double>(xyz[0]), mat[0], fma(static_cast<double>(xyz[1]), mat[3], fma(static_cast<double>(xyz[2]), mat[6], vec[0]))),
1470  fma(static_cast<double>(xyz[0]), mat[1], fma(static_cast<double>(xyz[1]), mat[4], fma(static_cast<double>(xyz[2]), mat[7], vec[1]))),
1471  fma(static_cast<double>(xyz[0]), mat[2], fma(static_cast<double>(xyz[1]), mat[5], fma(static_cast<double>(xyz[2]), mat[8], vec[2])))); // 9 fma = 9 flops
1472 }
1473 
1474 // ----------------------------> BBox <-------------------------------------
1475 
1476 // Base-class for static polymorphism (cannot be constructed directly)
1477 template<typename Vec3T>
1478 struct BaseBBox
1479 {
1480  Vec3T mCoord[2];
1481  __hostdev__ bool operator==(const BaseBBox& rhs) const { return mCoord[0] == rhs.mCoord[0] && mCoord[1] == rhs.mCoord[1]; };
1482  __hostdev__ bool operator!=(const BaseBBox& rhs) const { return mCoord[0] != rhs.mCoord[0] || mCoord[1] != rhs.mCoord[1]; };
1483  __hostdev__ const Vec3T& operator[](int i) const { return mCoord[i]; }
1484  __hostdev__ Vec3T& operator[](int i) { return mCoord[i]; }
1485  __hostdev__ Vec3T& min() { return mCoord[0]; }
1486  __hostdev__ Vec3T& max() { return mCoord[1]; }
1487  __hostdev__ const Vec3T& min() const { return mCoord[0]; }
1488  __hostdev__ const Vec3T& max() const { return mCoord[1]; }
1489  __hostdev__ Coord& translate(const Vec3T& xyz)
1490  {
1491  mCoord[0] += xyz;
1492  mCoord[1] += xyz;
1493  return *this;
1494  }
1495  // @brief Expand this bounding box to enclose point (i, j, k).
1496  __hostdev__ BaseBBox& expand(const Vec3T& xyz)
1497  {
1498  mCoord[0].minComponent(xyz);
1499  mCoord[1].maxComponent(xyz);
1500  return *this;
1501  }
1502  //__hostdev__ BaseBBox expandBy(typename Vec3T::ValueType padding) const
1503  //{
1504  // return BaseBBox(mCoord[0].offsetBy(-padding),mCoord[1].offsetBy(padding));
1505  //}
1506  __hostdev__ bool isInside(const Vec3T& xyz)
1507  {
1508  if (xyz[0] < mCoord[0][0] || xyz[1] < mCoord[0][1] || xyz[2] < mCoord[0][2])
1509  return false;
1510  if (xyz[0] > mCoord[1][0] || xyz[1] > mCoord[1][1] || xyz[2] > mCoord[1][2])
1511  return false;
1512  return true;
1513  }
1514 
1515 protected:
1517  __hostdev__ BaseBBox(const Vec3T& min, const Vec3T& max)
1518  : mCoord{min, max}
1519  {
1520  }
1521 }; // BaseBBox
1522 
1524 struct BBox;
1525 
1526 /// @brief Partial template specialization for floating point coordinate types.
1527 ///
1528 /// @note Min is inclusive and max is exclusive. If min = max the dimension of
1529 /// the bounding box is zero and therefore it is also empty.
1530 template<typename Vec3T>
1531 struct BBox<Vec3T, true> : public BaseBBox<Vec3T>
1532 {
1533  using Vec3Type = Vec3T;
1534  using ValueType = typename Vec3T::ValueType;
1535  static_assert(is_floating_point<ValueType>::value, "Expected a floating point coordinate type");
1537  using BaseT::mCoord;
1539  : BaseT(Vec3T( Maximum<typename Vec3T::ValueType>::value()),
1540  Vec3T(-Maximum<typename Vec3T::ValueType>::value()))
1541  {
1542  }
1543  __hostdev__ BBox(const Vec3T& min, const Vec3T& max)
1544  : BaseT(min, max)
1545  {
1546  }
1547  __hostdev__ BBox(const Coord& min, const Coord& max)
1548  : BaseT(Vec3T(ValueType(min[0]), ValueType(min[1]), ValueType(min[2])),
1549  Vec3T(ValueType(max[0] + 1), ValueType(max[1] + 1), ValueType(max[2] + 1)))
1550  {
1551  }
1552  __hostdev__ BBox(const BaseBBox<Coord>& bbox) : BBox(bbox[0], bbox[1]) {}
1553  __hostdev__ bool empty() const { return mCoord[0][0] >= mCoord[1][0] ||
1554  mCoord[0][1] >= mCoord[1][1] ||
1555  mCoord[0][2] >= mCoord[1][2]; }
1556  __hostdev__ Vec3T dim() const { return this->empty() ? Vec3T(0) : this->max() - this->min(); }
1557  __hostdev__ bool isInside(const Vec3T& p) const
1558  {
1559  return p[0] > mCoord[0][0] && p[1] > mCoord[0][1] && p[2] > mCoord[0][2] &&
1560  p[0] < mCoord[1][0] && p[1] < mCoord[1][1] && p[2] < mCoord[1][2];
1561  }
1562 };// BBox<Vec3T, true>
1563 
1564 /// @brief Partial template specialization for integer coordinate types
1565 ///
1566 /// @note Both min and max are INCLUDED in the bbox so dim = max - min + 1. So,
1567 /// if min = max the bounding box contains exactly one point and dim = 1!
1568 template<typename CoordT>
1569 struct BBox<CoordT, false> : public BaseBBox<CoordT>
1570 {
1571  static_assert(is_same<int, typename CoordT::ValueType>::value, "Expected \"int\" coordinate type");
1573  using BaseT::mCoord;
1574  /// @brief Iterator over the domain covered by a BBox
1575  /// @details z is the fastest-moving coordinate.
1576  class Iterator
1577  {
1578  const BBox& mBBox;
1579  CoordT mPos;
1580  public:
1582  : mBBox(b)
1583  , mPos(b.min())
1584  {
1585  }
1587  {
1588  if (mPos[2] < mBBox[1][2]) {// this is the most common case
1589  ++mPos[2];
1590  } else if (mPos[1] < mBBox[1][1]) {
1591  mPos[2] = mBBox[0][2];
1592  ++mPos[1];
1593  } else if (mPos[0] <= mBBox[1][0]) {
1594  mPos[2] = mBBox[0][2];
1595  mPos[1] = mBBox[0][1];
1596  ++mPos[0];
1597  }
1598  return *this;
1599  }
1600  __hostdev__ Iterator operator++(int)
1601  {
1602  auto tmp = *this;
1603  ++(*this);
1604  return tmp;
1605  }
1606  /// @brief Return @c true if the iterator still points to a valid coordinate.
1607  __hostdev__ operator bool() const { return mPos[0] <= mBBox[1][0]; }
1608  __hostdev__ const CoordT& operator*() const { return mPos; }
1609  }; // Iterator
1610  __hostdev__ Iterator begin() const { return Iterator{*this}; }
1612  : BaseT(CoordT::max(), CoordT::min())
1613  {
1614  }
1615  __hostdev__ BBox(const CoordT& min, const CoordT& max)
1616  : BaseT(min, max)
1617  {
1618  }
1619  template<typename SplitT>
1620  __hostdev__ BBox(BBox& other, const SplitT&)
1621  : BaseT(other.mCoord[0], other.mCoord[1])
1622  {
1623  NANOVDB_ASSERT(this->is_divisible());
1624  const int n = MaxIndex(this->dim());
1625  mCoord[1][n] = (mCoord[0][n] + mCoord[1][n]) >> 1;
1626  other.mCoord[0][n] = mCoord[1][n] + 1;
1627  }
1628  __hostdev__ bool is_divisible() const { return mCoord[0][0] < mCoord[1][0] &&
1629  mCoord[0][1] < mCoord[1][1] &&
1630  mCoord[0][2] < mCoord[1][2]; }
1631  /// @brief Return true if this bounding box is empty, i.e. uninitialized
1632  __hostdev__ bool empty() const { return mCoord[0][0] > mCoord[1][0] ||
1633  mCoord[0][1] > mCoord[1][1] ||
1634  mCoord[0][2] > mCoord[1][2]; }
1635  __hostdev__ CoordT dim() const { return this->empty() ? Coord(0) : this->max() - this->min() + Coord(1); }
1636  __hostdev__ uint64_t volume() const { auto d = this->dim(); return uint64_t(d[0])*uint64_t(d[1])*uint64_t(d[2]); }
1637  __hostdev__ bool isInside(const CoordT& p) const { return !(CoordT::lessThan(p, this->min()) || CoordT::lessThan(this->max(), p)); }
1638  __hostdev__ bool isInside(const BBox& b) const
1639  {
1640  return !(CoordT::lessThan(b.min(), this->min()) || CoordT::lessThan(this->max(), b.max()));
1641  }
1642 
1643  /// @warning This converts a CoordBBox into a floating-point bounding box which implies that max += 1 !
1644  template<typename RealT>
1646  {
1647  static_assert(is_floating_point<RealT>::value, "CoordBBox::asReal: Expected a floating point coordinate");
1648  return BBox<Vec3<RealT>>(Vec3<RealT>(RealT(mCoord[0][0]), RealT(mCoord[0][1]), RealT(mCoord[0][2])),
1649  Vec3<RealT>(RealT(mCoord[1][0] + 1), RealT(mCoord[1][1] + 1), RealT(mCoord[1][2] + 1)));
1650  }
1651  /// @brief Return a new instance that is expanded by the specified padding.
1652  __hostdev__ BBox expandBy(typename CoordT::ValueType padding) const
1653  {
1654  return BBox(mCoord[0].offsetBy(-padding), mCoord[1].offsetBy(padding));
1655  }
1656 };// BBox<CoordT, false>
1657 
1660 
1661 // -------------------> Find lowest and highest bit in a word <----------------------------
1662 
1663 /// @brief Returns the index of the lowest, i.e. least significant, on bit in the specified 32 bit word
1664 ///
1665 /// @warning Assumes that at least one bit is set in the word, i.e. @a v != uint32_t(0)!
1667 __hostdev__ static inline uint32_t FindLowestOn(uint32_t v)
1668 {
1669  NANOVDB_ASSERT(v);
1670 #if defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS)
1671  unsigned long index;
1672  _BitScanForward(&index, v);
1673  return static_cast<uint32_t>(index);
1674 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS)
1675  return static_cast<uint32_t>(__builtin_ctzl(v));
1676 #else
1677  static const unsigned char DeBruijn[32] = {
1678  0, 1, 28, 2, 29, 14, 24, 3, 30, 22, 20, 15, 25, 17, 4, 8, 31, 27, 13, 23, 21, 19, 16, 7, 26, 12, 18, 6, 11, 5, 10, 9};
1679 // disable unary minus on unsigned warning
1680 #if defined(_MSC_VER) && !defined(__NVCC__)
1681 #pragma warning(push)
1682 #pragma warning(disable : 4146)
1683 #endif
1684  return DeBruijn[uint32_t((v & -v) * 0x077CB531U) >> 27];
1685 #if defined(_MSC_VER) && !defined(__NVCC__)
1686 #pragma warning(pop)
1687 #endif
1688 
1689 #endif
1690 }
1691 
1692 /// @brief Returns the index of the highest, i.e. most significant, on bit in the specified 32 bit word
1693 ///
1694 /// @warning Assumes that at least one bit is set in the word, i.e. @a v != uint32_t(0)!
1696 __hostdev__ static inline uint32_t FindHighestOn(uint32_t v)
1697 {
1698  NANOVDB_ASSERT(v);
1699 #if defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS)
1700  unsigned long index;
1701  _BitScanReverse(&index, v);
1702  return static_cast<uint32_t>(index);
1703 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS)
1704  return sizeof(unsigned long) * 8 - 1 - __builtin_clzl(v);
1705 
1706 #else
1707  static const unsigned char DeBruijn[32] = {
1708  0, 9, 1, 10, 13, 21, 2, 29, 11, 14, 16, 18, 22, 25, 3, 30, 8, 12, 20, 28, 15, 17, 24, 7, 19, 27, 23, 6, 26, 5, 4, 31};
1709  v |= v >> 1; // first round down to one less than a power of 2
1710  v |= v >> 2;
1711  v |= v >> 4;
1712  v |= v >> 8;
1713  v |= v >> 16;
1714  return DeBruijn[uint32_t(v * 0x07C4ACDDU) >> 27];
1715 #endif
1716 }
1717 
1718 /// @brief Returns the index of the lowest, i.e. least significant, on bit in the specified 64 bit word
1719 ///
1720 /// @warning Assumes that at least one bit is set in the word, i.e. @a v != uint32_t(0)!
1722 __hostdev__ static inline uint32_t FindLowestOn(uint64_t v)
1723 {
1724  NANOVDB_ASSERT(v);
1725 #if defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS)
1726  unsigned long index;
1727  _BitScanForward64(&index, v);
1728  return static_cast<uint32_t>(index);
1729 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS)
1730  return static_cast<uint32_t>(__builtin_ctzll(v));
1731 #else
1732  static const unsigned char DeBruijn[64] = {
1733  0, 1, 2, 53, 3, 7, 54, 27, 4, 38, 41, 8, 34, 55, 48, 28,
1734  62, 5, 39, 46, 44, 42, 22, 9, 24, 35, 59, 56, 49, 18, 29, 11,
1735  63, 52, 6, 26, 37, 40, 33, 47, 61, 45, 43, 21, 23, 58, 17, 10,
1736  51, 25, 36, 32, 60, 20, 57, 16, 50, 31, 19, 15, 30, 14, 13, 12,
1737  };
1738 // disable unary minus on unsigned warning
1739 #if defined(_MSC_VER) && !defined(__NVCC__)
1740 #pragma warning(push)
1741 #pragma warning(disable : 4146)
1742 #endif
1743  return DeBruijn[uint64_t((v & -v) * UINT64_C(0x022FDD63CC95386D)) >> 58];
1744 #if defined(_MSC_VER) && !defined(__NVCC__)
1745 #pragma warning(pop)
1746 #endif
1747 
1748 #endif
1749 }
1750 
1751 /// @brief Returns the index of the highest, i.e. most significant, on bit in the specified 64 bit word
1752 ///
1753 /// @warning Assumes that at least one bit is set in the word, i.e. @a v != uint32_t(0)!
1755 __hostdev__ static inline uint32_t FindHighestOn(uint64_t v)
1756 {
1757  NANOVDB_ASSERT(v);
1758 #if defined(_MSC_VER) && defined(NANOVDB_USE_INTRINSICS)
1759  unsigned long index;
1760  _BitScanReverse64(&index, v);
1761  return static_cast<uint32_t>(index);
1762 #elif (defined(__GNUC__) || defined(__clang__)) && defined(NANOVDB_USE_INTRINSICS)
1763  return sizeof(unsigned long) * 8 - 1 - __builtin_clzll(v);
1764 #else
1765  const uint32_t* p = reinterpret_cast<const uint32_t*>(&v);
1766  return p[1] ? 32u + FindHighestOn(p[1]) : FindHighestOn(p[0]);
1767 #endif
1768 }
1769 
1770 // ----------------------------> CountOn <--------------------------------------
1771 
1772 /// @return Number of bits that are on in the specified 64-bit word
1774 __hostdev__ inline uint32_t CountOn(uint64_t v)
1775 {
1776 // __popcnt* intrinsic support was added in VS 2019 16.8
1777 #if defined(_MSC_VER) && defined(_M_X64) && (_MSC_VER >= 1928)
1778  v = __popcnt64(v);
1779 #elif (defined(__GNUC__) || defined(__clang__))
1780  v = __builtin_popcountll(v);
1781 #else
1782  // Software Implementation
1783  v = v - ((v >> 1) & uint64_t(0x5555555555555555));
1784  v = (v & uint64_t(0x3333333333333333)) + ((v >> 2) & uint64_t(0x3333333333333333));
1785  v = (((v + (v >> 4)) & uint64_t(0xF0F0F0F0F0F0F0F)) * uint64_t(0x101010101010101)) >> 56;
1786 #endif
1787  return static_cast<uint32_t>(v);
1788 }
1789 
1790 // ----------------------------> Mask <--------------------------------------
1791 
1792 /// @brief Bit-mask to encode active states and facilitate sequential iterators
1793 /// and a fast codec for I/O compression.
1794 template<uint32_t LOG2DIM>
1795 class Mask
1796 {
1797  static constexpr uint32_t SIZE = 1U << (3 * LOG2DIM); // Number of bits in mask
1798  static constexpr uint32_t WORD_COUNT = SIZE >> 6; // Number of 64 bit words
1799  uint64_t mWords[WORD_COUNT];
1800 
1801 public:
1802  /// @brief Return the memory footprint in bytes of this Mask
1803  __hostdev__ static size_t memUsage() { return sizeof(Mask); }
1804 
1805  /// @brief Return the number of bits available in this Mask
1806  __hostdev__ static uint32_t bitCount() { return SIZE; }
1807 
1808  /// @brief Return the number of machine words used by this Mask
1809  __hostdev__ static uint32_t wordCount() { return WORD_COUNT; }
1810 
1811  __hostdev__ uint32_t countOn() const
1812  {
1813  uint32_t sum = 0, n = WORD_COUNT;
1814  for (const uint64_t* w = mWords; n--; ++w)
1815  sum += CountOn(*w);
1816  return sum;
1817  }
1818 
1819  class Iterator
1820  {
1821  public:
1823  : mPos(Mask::SIZE)
1824  , mParent(nullptr)
1825  {
1826  }
1827  __hostdev__ Iterator(uint32_t pos, const Mask* parent)
1828  : mPos(pos)
1829  , mParent(parent)
1830  {
1831  }
1832  Iterator& operator=(const Iterator&) = default;
1833  __hostdev__ uint32_t operator*() const { return mPos; }
1834  __hostdev__ operator bool() const { return mPos != Mask::SIZE; }
1836  {
1837  mPos = mParent->findNextOn(mPos + 1);
1838  return *this;
1839  }
1840 
1841  private:
1842  uint32_t mPos;
1843  const Mask* mParent;
1844  }; // Member class MaskIterator
1845 
1846  /// @brief Initialize all bits to zero.
1848  {
1849  for (uint32_t i = 0; i < WORD_COUNT; ++i)
1850  mWords[i] = 0;
1851  }
1853  {
1854  const uint64_t v = on ? ~uint64_t(0) : uint64_t(0);
1855  for (uint32_t i = 0; i < WORD_COUNT; ++i)
1856  mWords[i] = v;
1857  }
1858 
1859  /// @brief Copy constructor
1860  __hostdev__ Mask(const Mask& other)
1861  {
1862  for (uint32_t i = 0; i < WORD_COUNT; ++i)
1863  mWords[i] = other.mWords[i];
1864  }
1865 
1866  /// @brief Return the <i>n</i>th word of the bit mask, for a word of arbitrary size.
1867  template<typename WordT>
1868  __hostdev__ WordT getWord(int n) const
1869  {
1870  NANOVDB_ASSERT(n * 8 * sizeof(WordT) < SIZE);
1871  return reinterpret_cast<const WordT*>(mWords)[n];
1872  }
1873 
1874  /// @brief Assignment operator that works with openvdb::util::NodeMask
1875  template<typename MaskT>
1876  __hostdev__ Mask& operator=(const MaskT& other)
1877  {
1878  static_assert(sizeof(Mask) == sizeof(MaskT), "Mismatching sizeof");
1879  static_assert(WORD_COUNT == MaskT::WORD_COUNT, "Mismatching word count");
1880  static_assert(LOG2DIM == MaskT::LOG2DIM, "Mismatching LOG2DIM");
1881  auto *src = reinterpret_cast<const uint64_t*>(&other);
1882  uint64_t *dst = mWords;
1883  for (uint32_t i = 0; i < WORD_COUNT; ++i) {
1884  *dst++ = *src++;
1885  }
1886  return *this;
1887  }
1888 
1889  __hostdev__ bool operator==(const Mask& other) const
1890  {
1891  for (uint32_t i = 0; i < WORD_COUNT; ++i) {
1892  if (mWords[i] != other.mWords[i]) return false;
1893  }
1894  return true;
1895  }
1896 
1897  __hostdev__ bool operator!=(const Mask& other) const { return !((*this) == other); }
1898 
1899  __hostdev__ Iterator beginOn() const { return Iterator(this->findFirstOn(), this); }
1900 
1901  /// @brief Return true if the given bit is set.
1902  __hostdev__ bool isOn(uint32_t n) const { return 0 != (mWords[n >> 6] & (uint64_t(1) << (n & 63))); }
1903 
1904  __hostdev__ bool isOn() const
1905  {
1906  for (uint32_t i = 0; i < WORD_COUNT; ++i)
1907  if (mWords[i] != ~uint64_t(0))
1908  return false;
1909  return true;
1910  }
1911 
1912  __hostdev__ bool isOff() const
1913  {
1914  for (uint32_t i = 0; i < WORD_COUNT; ++i)
1915  if (mWords[i] != uint64_t(0))
1916  return false;
1917  return true;
1918  }
1919 
1920  /// @brief Set the given bit on.
1921  __hostdev__ void setOn(uint32_t n) { mWords[n >> 6] |= uint64_t(1) << (n & 63); }
1922  __hostdev__ void setOff(uint32_t n) { mWords[n >> 6] &= ~(uint64_t(1) << (n & 63)); }
1923 
1924  __hostdev__ void set(uint32_t n, bool On)
1925  {
1926 #if 1 // switch between branchless
1927  auto &word = mWords[n >> 6];
1928  n &= 63;
1929  word &= ~(uint64_t(1) << n);
1930  word |= uint64_t(On) << n;
1931 #else
1932  On ? this->setOn(n) : this->setOff(n);
1933 #endif
1934  }
1935 
1936  /// @brief Set all bits on
1938  {
1939  for (uint32_t i = 0; i < WORD_COUNT; ++i)
1940  mWords[i] = ~uint64_t(0);
1941  }
1942 
1943  /// @brief Set all bits off
1945  {
1946  for (uint32_t i = 0; i < WORD_COUNT; ++i)
1947  mWords[i] = uint64_t(0);
1948  }
1949 
1950  /// @brief Set all bits off
1951  __hostdev__ void set(bool on)
1952  {
1953  const uint64_t v = on ? ~uint64_t(0) : uint64_t(0);
1954  for (uint32_t i = 0; i < WORD_COUNT; ++i)
1955  mWords[i] = v;
1956  }
1957  /// brief Toggle the state of all bits in the mask
1959  {
1960  uint32_t n = WORD_COUNT;
1961  for (auto* w = mWords; n--; ++w)
1962  *w = ~*w;
1963  }
1964  __hostdev__ void toggle(uint32_t n) { mWords[n >> 6] ^= uint64_t(1) << (n & 63); }
1965 
1966 private:
1967 
1969  __hostdev__ uint32_t findFirstOn() const
1970  {
1971  uint32_t n = 0;
1972  const uint64_t* w = mWords;
1973  for (; n < WORD_COUNT && !*w; ++w, ++n)
1974  ;
1975  return n == WORD_COUNT ? SIZE : (n << 6) + FindLowestOn(*w);
1976  }
1977 
1979  __hostdev__ uint32_t findNextOn(uint32_t start) const
1980  {
1981  uint32_t n = start >> 6; // initiate
1982  if (n >= WORD_COUNT)
1983  return SIZE; // check for out of bounds
1984  uint32_t m = start & 63;
1985  uint64_t b = mWords[n];
1986  if (b & (uint64_t(1) << m))
1987  return start; // simple case: start is on
1988  b &= ~uint64_t(0) << m; // mask out lower bits
1989  while (!b && ++n < WORD_COUNT)
1990  b = mWords[n]; // find next non-zero word
1991  return (!b ? SIZE : (n << 6) + FindLowestOn(b)); // catch last word=0
1992  }
1993 }; // Mask class
1994 
1995 // ----------------------------> Map <--------------------------------------
1996 
1997 /// @brief Defines an affine transform and its inverse represented as a 3x3 matrix and a vec3 translation
1998 struct Map
1999 {
2000  float mMatF[9]; // 9*4B <- 3x3 matrix
2001  float mInvMatF[9]; // 9*4B <- 3x3 matrix
2002  float mVecF[3]; // 3*4B <- translation
2003  float mTaperF; // 4B, placeholder for taper value
2004  double mMatD[9]; // 9*8B <- 3x3 matrix
2005  double mInvMatD[9]; // 9*8B <- 3x3 matrix
2006  double mVecD[3]; // 3*8B <- translation
2007  double mTaperD; // 8B, placeholder for taper value
2008 
2009  // This method can only be called on the host to initialize the member data
2010  template<typename Mat4T>
2011  __hostdev__ void set(const Mat4T& mat, const Mat4T& invMat, double taper);
2012 
2013  template<typename Vec3T>
2014  __hostdev__ Vec3T applyMap(const Vec3T& xyz) const { return matMult(mMatD, mVecD, xyz); }
2015  template<typename Vec3T>
2016  __hostdev__ Vec3T applyMapF(const Vec3T& xyz) const { return matMult(mMatF, mVecF, xyz); }
2017 
2018  template<typename Vec3T>
2019  __hostdev__ Vec3T applyJacobian(const Vec3T& xyz) const { return matMult(mMatD, xyz); }
2020  template<typename Vec3T>
2021  __hostdev__ Vec3T applyJacobianF(const Vec3T& xyz) const { return matMult(mMatF, xyz); }
2022 
2023  template<typename Vec3T>
2024  __hostdev__ Vec3T applyInverseMap(const Vec3T& xyz) const
2025  {
2026  return matMult(mInvMatD, Vec3T(xyz[0] - mVecD[0], xyz[1] - mVecD[1], xyz[2] - mVecD[2]));
2027  }
2028  template<typename Vec3T>
2029  __hostdev__ Vec3T applyInverseMapF(const Vec3T& xyz) const
2030  {
2031  return matMult(mInvMatF, Vec3T(xyz[0] - mVecF[0], xyz[1] - mVecF[1], xyz[2] - mVecF[2]));
2032  }
2033 
2034  template<typename Vec3T>
2035  __hostdev__ Vec3T applyInverseJacobian(const Vec3T& xyz) const { return matMult(mInvMatD, xyz); }
2036  template<typename Vec3T>
2037  __hostdev__ Vec3T applyInverseJacobianF(const Vec3T& xyz) const { return matMult(mInvMatF, xyz); }
2038 
2039  template<typename Vec3T>
2040  __hostdev__ Vec3T applyIJT(const Vec3T& xyz) const { return matMultT(mInvMatD, xyz); }
2041  template<typename Vec3T>
2042  __hostdev__ Vec3T applyIJTF(const Vec3T& xyz) const { return matMultT(mInvMatF, xyz); }
2043 }; // Map
2044 
2045 template<typename Mat4T>
2046 __hostdev__ void Map::set(const Mat4T& mat, const Mat4T& invMat, double taper)
2047 {
2048  float * mf = mMatF, *vf = mVecF;
2049  float* mif = mInvMatF;
2050  double *md = mMatD, *vd = mVecD;
2051  double* mid = mInvMatD;
2052  mTaperF = static_cast<float>(taper);
2053  mTaperD = taper;
2054  for (int i = 0; i < 3; ++i) {
2055  *vd++ = mat[3][i]; //translation
2056  *vf++ = static_cast<float>(mat[3][i]);
2057  for (int j = 0; j < 3; ++j) {
2058  *md++ = mat[j][i]; //transposed
2059  *mid++ = invMat[j][i];
2060  *mf++ = static_cast<float>(mat[j][i]);
2061  *mif++ = static_cast<float>(invMat[j][i]);
2062  }
2063  }
2064 }
2065 
2066 // ----------------------------> GridBlindMetaData <--------------------------------------
2067 
2068 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) GridBlindMetaData
2069 {
2070  static const int MaxNameSize = 256;// due to NULL termination the maximum length is one less!
2071  int64_t mByteOffset; // byte offset to the blind data, relative to the GridData.
2072  uint64_t mElementCount; // number of elements, e.g. point count
2073  uint32_t mFlags; // flags
2074  GridBlindDataSemantic mSemantic; // semantic meaning of the data.
2076  GridType mDataType; // 4 bytes
2077  char mName[MaxNameSize];// note this include the NULL termination
2078 
2079  /// @brief return memory usage in bytes for the class (note this computes for all blindMetaData structures.)
2080  __hostdev__ static uint64_t memUsage(uint64_t blindDataCount = 0)
2081  {
2082  return blindDataCount * sizeof(GridBlindMetaData);
2083  }
2084 
2085  __hostdev__ void setBlindData(void *ptr) { mByteOffset = PtrDiff(ptr, this); }
2086 
2087  template <typename T>
2088  __hostdev__ const T* getBlindData() const { return PtrAdd<T>(this, mByteOffset); }
2089 
2090 }; // GridBlindMetaData
2091 
2092 // ----------------------------> NodeTrait <--------------------------------------
2093 
2094 /// @brief Struct to derive node type from its level in a given
2095 /// grid, tree or root while perserving constness
2096 template<typename GridOrTreeOrRootT, int LEVEL>
2097 struct NodeTrait;
2098 
2099 // Partial template specialization of above Node struct
2100 template<typename GridOrTreeOrRootT>
2101 struct NodeTrait<GridOrTreeOrRootT, 0>
2102 {
2103  static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3, "Tree depth is not supported");
2104  using Type = typename GridOrTreeOrRootT::LeafNodeType;
2105  using type = typename GridOrTreeOrRootT::LeafNodeType;
2106 };
2107 template<typename GridOrTreeOrRootT>
2108 struct NodeTrait<const GridOrTreeOrRootT, 0>
2109 {
2110  static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3, "Tree depth is not supported");
2111  using Type = const typename GridOrTreeOrRootT::LeafNodeType;
2112  using type = const typename GridOrTreeOrRootT::LeafNodeType;
2113 };
2114 
2115 template<typename GridOrTreeOrRootT>
2116 struct NodeTrait<GridOrTreeOrRootT, 1>
2117 {
2118  static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3, "Tree depth is not supported");
2119  using Type = typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType;
2120  using type = typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType;
2121 };
2122 template<typename GridOrTreeOrRootT>
2123 struct NodeTrait<const GridOrTreeOrRootT, 1>
2124 {
2125  static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3, "Tree depth is not supported");
2126  using Type = const typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType;
2127  using type = const typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType;
2128 };
2129 template<typename GridOrTreeOrRootT>
2130 struct NodeTrait<GridOrTreeOrRootT, 2>
2131 {
2132  static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3, "Tree depth is not supported");
2133  using Type = typename GridOrTreeOrRootT::RootType::ChildNodeType;
2134  using type = typename GridOrTreeOrRootT::RootType::ChildNodeType;
2135 };
2136 template<typename GridOrTreeOrRootT>
2137 struct NodeTrait<const GridOrTreeOrRootT, 2>
2138 {
2139  static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3, "Tree depth is not supported");
2140  using Type = const typename GridOrTreeOrRootT::RootType::ChildNodeType;
2141  using type = const typename GridOrTreeOrRootT::RootType::ChildNodeType;
2142 };
2143 template<typename GridOrTreeOrRootT>
2144 struct NodeTrait<GridOrTreeOrRootT, 3>
2145 {
2146  static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3, "Tree depth is not supported");
2147  using Type = typename GridOrTreeOrRootT::RootType;
2148  using type = typename GridOrTreeOrRootT::RootType;
2149 };
2150 
2151 template<typename GridOrTreeOrRootT>
2152 struct NodeTrait<const GridOrTreeOrRootT, 3>
2153 {
2154  static_assert(GridOrTreeOrRootT::RootType::LEVEL == 3, "Tree depth is not supported");
2155  using Type = const typename GridOrTreeOrRootT::RootType;
2156  using type = const typename GridOrTreeOrRootT::RootType;
2157 };
2158 
2159 // ----------------------------> Grid <--------------------------------------
2160 
2161 /*
2162  The following class and comment is for internal use only
2163 
2164  Memory layout:
2165 
2166  Grid -> 39 x double (world bbox and affine transformation)
2167  Tree -> Root 3 x ValueType + int32_t + N x Tiles (background,min,max,tileCount + tileCount x Tiles)
2168 
2169  N2 upper InternalNodes each with 2 bit masks, N2 tiles, and min/max values
2170 
2171  N1 lower InternalNodes each with 2 bit masks, N1 tiles, and min/max values
2172 
2173  N0 LeafNodes each with a bit mask, N0 ValueTypes and min/max
2174 
2175  Example layout: ("---" implies it has a custom offset, "..." implies zero or more)
2176  [GridData][TreeData]---[RootData][ROOT TILES...]---[NodeData<5>]---[ModeData<4>]---[LeafData<3>]---[BLINDMETA...]---[BLIND0]---[BLIND1]---etc.
2177 */
2178 
2179 /// @brief Struct with all the member data of the Grid (useful during serialization of an openvdb grid)
2180 ///
2181 /// @note The transform is assumed to be affine (so linear) and have uniform scale! So frustum transforms
2182 /// and non-uniform scaling are not supported (primarily because they complicate ray-tracing in index space)
2183 ///
2184 /// @note No client code should (or can) interface with this struct so it can safely be ignored!
2185 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) GridData
2186 {// sizeof(GridData) = 672B
2187  static const int MaxNameSize = 256;// due to NULL termination the maximum length is one less
2188  uint64_t mMagic; // 8B magic to validate it is valid grid data.
2189  uint64_t mChecksum; // 8B. Checksum of grid buffer.
2190  Version mVersion;// 4B major, minor, and patch version numbers
2191  uint32_t mFlags; // 4B. flags for grid.
2192  uint32_t mGridIndex;// 4B. Index of this grid in the buffer
2193  uint32_t mGridCount; // 4B. Total number of grids in the buffer
2194  uint64_t mGridSize; // 8B. byte count of this entire grid occupied in the buffer.
2195  char mGridName[MaxNameSize]; // 256B
2196  Map mMap; // 264B. affine transformation between index and world space in both single and double precision
2197  BBox<Vec3R> mWorldBBox; // 48B. floating-point AABB of active values in WORLD SPACE (2 x 3 doubles)
2198  Vec3R mVoxelSize; // 24B. size of a voxel in world units
2201  int64_t mBlindMetadataOffset; // 8B. offset of GridBlindMetaData structures that follow this grid.
2202  uint32_t mBlindMetadataCount; // 4B. count of GridBlindMetaData structures that follow this grid.
2203 
2204 
2205  // Set and unset various bit flags
2206  __hostdev__ void setFlagsOff() { mFlags = uint32_t(0); }
2207  __hostdev__ void setMinMaxOn(bool on = true)
2208  {
2209  if (on) {
2210  mFlags |= static_cast<uint32_t>(GridFlags::HasMinMax);
2211  } else {
2212  mFlags &= ~static_cast<uint32_t>(GridFlags::HasMinMax);
2213  }
2214  }
2215  __hostdev__ void setBBoxOn(bool on = true)
2216  {
2217  if (on) {
2218  mFlags |= static_cast<uint32_t>(GridFlags::HasBBox);
2219  } else {
2220  mFlags &= ~static_cast<uint32_t>(GridFlags::HasBBox);
2221  }
2222  }
2223  __hostdev__ void setLongGridNameOn(bool on = true)
2224  {
2225  if (on) {
2226  mFlags |= static_cast<uint32_t>(GridFlags::HasLongGridName);
2227  } else {
2228  mFlags &= ~static_cast<uint32_t>(GridFlags::HasLongGridName);
2229  }
2230  }
2231  __hostdev__ void setAverageOn(bool on = true)
2232  {
2233  if (on) {
2234  mFlags |= static_cast<uint32_t>(GridFlags::HasAverage);
2235  } else {
2236  mFlags &= ~static_cast<uint32_t>(GridFlags::HasAverage);
2237  }
2238  }
2239  __hostdev__ void setStdDeviationOn(bool on = true)
2240  {
2241  if (on) {
2242  mFlags |= static_cast<uint32_t>(GridFlags::HasStdDeviation);
2243  } else {
2244  mFlags &= ~static_cast<uint32_t>(GridFlags::HasStdDeviation);
2245  }
2246  }
2247  __hostdev__ void setBreadthFirstOn(bool on = true)
2248  {
2249  if (on) {
2250  mFlags |= static_cast<uint32_t>(GridFlags::IsBreadthFirst);
2251  } else {
2252  mFlags &= ~static_cast<uint32_t>(GridFlags::IsBreadthFirst);
2253  }
2254  }
2255 
2256  // Affine transformations based on double precision
2257  template<typename Vec3T>
2258  __hostdev__ Vec3T applyMap(const Vec3T& xyz) const { return mMap.applyMap(xyz); } // Pos: index -> world
2259  template<typename Vec3T>
2260  __hostdev__ Vec3T applyInverseMap(const Vec3T& xyz) const { return mMap.applyInverseMap(xyz); } // Pos: world -> index
2261  template<typename Vec3T>
2262  __hostdev__ Vec3T applyJacobian(const Vec3T& xyz) const { return mMap.applyJacobian(xyz); } // Dir: index -> world
2263  template<typename Vec3T>
2264  __hostdev__ Vec3T applyInverseJacobian(const Vec3T& xyz) const { return mMap.applyInverseJacobian(xyz); } // Dir: world -> index
2265  template<typename Vec3T>
2266  __hostdev__ Vec3T applyIJT(const Vec3T& xyz) const { return mMap.applyIJT(xyz); }
2267  // Affine transformations based on single precision
2268  template<typename Vec3T>
2269  __hostdev__ Vec3T applyMapF(const Vec3T& xyz) const { return mMap.applyMapF(xyz); } // Pos: index -> world
2270  template<typename Vec3T>
2271  __hostdev__ Vec3T applyInverseMapF(const Vec3T& xyz) const { return mMap.applyInverseMapF(xyz); } // Pos: world -> index
2272  template<typename Vec3T>
2273  __hostdev__ Vec3T applyJacobianF(const Vec3T& xyz) const { return mMap.applyJacobianF(xyz); } // Dir: index -> world
2274  template<typename Vec3T>
2275  __hostdev__ Vec3T applyInverseJacobianF(const Vec3T& xyz) const { return mMap.applyInverseJacobianF(xyz); } // Dir: world -> index
2276  template<typename Vec3T>
2277  __hostdev__ Vec3T applyIJTF(const Vec3T& xyz) const { return mMap.applyIJTF(xyz); }
2278 
2279  // @brief Return a non-const void pointer to the tree
2280  __hostdev__ void* treePtr() { return this + 1; }
2281 
2282  // @brief Return a const void pointer to the tree
2283  __hostdev__ const void* treePtr() const { return this + 1; }
2284 
2285  /// @brief Returns a const reference to the blindMetaData at the specified linear offset.
2286  ///
2287  /// @warning The linear offset is assumed to be in the valid range
2289  {
2290  NANOVDB_ASSERT(n < mBlindMetadataCount);
2291  return PtrAdd<GridBlindMetaData>(this, mBlindMetadataOffset) + n;
2292  }
2293 
2294 }; // GridData
2295 
2296 // Forward declaration of accelerated random access class
2297 template <typename BuildT, int LEVEL0 = -1, int LEVEL1 = -1, int LEVEL2 = -1>
2299 
2300 template <typename BuildT>
2302 
2303 /// @brief Highest level of the data structure. Contains a tree and a world->index
2304 /// transform (that currently only supports uniform scaling and translation).
2305 ///
2306 /// @note This the API of this class to interface with client code
2307 template<typename TreeT>
2308 class Grid : private GridData
2309 {
2310 public:
2311  using TreeType = TreeT;
2312  using RootType = typename TreeT::RootType;
2314  using ValueType = typename TreeT::ValueType;
2315  using BuildType = typename TreeT::BuildType;// in rare cases BuildType != ValueType, e.g. then BuildType = ValueMask and ValueType = bool
2316  using CoordType = typename TreeT::CoordType;
2318 
2319  /// @brief Disallow constructions, copy and assignment
2320  ///
2321  /// @note Only a Serializer, defined elsewhere, can instantiate this class
2322  Grid(const Grid&) = delete;
2323  Grid& operator=(const Grid&) = delete;
2324  ~Grid() = delete;
2325 
2326  __hostdev__ Version version() const { return DataType::mVersion; }
2327 
2328  __hostdev__ DataType* data() { return reinterpret_cast<DataType*>(this); }
2329 
2330  __hostdev__ const DataType* data() const { return reinterpret_cast<const DataType*>(this); }
2331 
2332  /// @brief Return memory usage in bytes for this class only.
2333  __hostdev__ static uint64_t memUsage() { return sizeof(GridData); }
2334 
2335  /// @brief Return the memory footprint of the entire grid, i.e. including all nodes and blind data
2336  __hostdev__ uint64_t gridSize() const { return DataType::mGridSize; }
2337 
2338  /// @brief Return index of this grid in the buffer
2339  __hostdev__ uint32_t gridIndex() const { return DataType::mGridIndex; }
2340 
2341  /// @brief Return total number of grids in the buffer
2342  __hostdev__ uint32_t gridCount() const { return DataType::mGridCount; }
2343 
2344  /// @brief Return a const reference to the tree
2345  __hostdev__ const TreeT& tree() const { return *reinterpret_cast<const TreeT*>(this->treePtr()); }
2346 
2347  /// @brief Return a non-const reference to the tree
2348  __hostdev__ TreeT& tree() { return *reinterpret_cast<TreeT*>(this->treePtr()); }
2349 
2350  /// @brief Return a new instance of a ReadAccessor used to access values in this grid
2351  __hostdev__ AccessorType getAccessor() const { return AccessorType(this->tree().root()); }
2352 
2353  /// @brief Return a const reference to the size of a voxel in world units
2354  __hostdev__ const Vec3R& voxelSize() const { return DataType::mVoxelSize; }
2355 
2356  /// @brief Return a const reference to the Map for this grid
2357  __hostdev__ const Map& map() const { return DataType::mMap; }
2358 
2359  /// @brief world to index space transformation
2360  template<typename Vec3T>
2361  __hostdev__ Vec3T worldToIndex(const Vec3T& xyz) const { return this->applyInverseMap(xyz); }
2362 
2363  /// @brief index to world space transformation
2364  template<typename Vec3T>
2365  __hostdev__ Vec3T indexToWorld(const Vec3T& xyz) const { return this->applyMap(xyz); }
2366 
2367  /// @brief transformation from index space direction to world space direction
2368  /// @warning assumes dir to be normalized
2369  template<typename Vec3T>
2370  __hostdev__ Vec3T indexToWorldDir(const Vec3T& dir) const { return this->applyJacobian(dir); }
2371 
2372  /// @brief transformation from world space direction to index space direction
2373  /// @warning assumes dir to be normalized
2374  template<typename Vec3T>
2375  __hostdev__ Vec3T worldToIndexDir(const Vec3T& dir) const { return this->applyInverseJacobian(dir); }
2376 
2377  /// @brief transform the gradient from index space to world space.
2378  /// @details Applies the inverse jacobian transform map.
2379  template<typename Vec3T>
2380  __hostdev__ Vec3T indexToWorldGrad(const Vec3T& grad) const { return this->applyIJT(grad); }
2381 
2382  /// @brief world to index space transformation
2383  template<typename Vec3T>
2384  __hostdev__ Vec3T worldToIndexF(const Vec3T& xyz) const { return this->applyInverseMapF(xyz); }
2385 
2386  /// @brief index to world space transformation
2387  template<typename Vec3T>
2388  __hostdev__ Vec3T indexToWorldF(const Vec3T& xyz) const { return this->applyMapF(xyz); }
2389 
2390  /// @brief transformation from index space direction to world space direction
2391  /// @warning assumes dir to be normalized
2392  template<typename Vec3T>
2393  __hostdev__ Vec3T indexToWorldDirF(const Vec3T& dir) const { return this->applyJacobianF(dir); }
2394 
2395  /// @brief transformation from world space direction to index space direction
2396  /// @warning assumes dir to be normalized
2397  template<typename Vec3T>
2398  __hostdev__ Vec3T worldToIndexDirF(const Vec3T& dir) const { return this->applyInverseJacobianF(dir); }
2399 
2400  /// @brief Transforms the gradient from index space to world space.
2401  /// @details Applies the inverse jacobian transform map.
2402  template<typename Vec3T>
2403  __hostdev__ Vec3T indexToWorldGradF(const Vec3T& grad) const { return DataType::applyIJTF(grad); }
2404 
2405  /// @brief Computes a AABB of active values in world space
2406  __hostdev__ const BBox<Vec3R>& worldBBox() const { return DataType::mWorldBBox; }
2407 
2408  /// @brief Computes a AABB of active values in index space
2409  ///
2410  /// @note This method is returning a floating point bounding box and not a CoordBBox. This makes
2411  /// it more useful for clipping rays.
2412  __hostdev__ const BBox<CoordType>& indexBBox() const { return this->tree().bbox(); }
2413 
2414  /// @brief Return the total number of active voxels in this tree.
2415  __hostdev__ uint64_t activeVoxelCount() const { return this->tree().activeVoxelCount(); }
2416 
2417  /// @brief Methods related to the classification of this grid
2418  __hostdev__ bool isValid() const { return DataType::mMagic == NANOVDB_MAGIC_NUMBER; }
2419  __hostdev__ const GridType& gridType() const { return DataType::mGridType; }
2420  __hostdev__ const GridClass& gridClass() const { return DataType::mGridClass; }
2421  __hostdev__ bool isLevelSet() const { return DataType::mGridClass == GridClass::LevelSet; }
2422  __hostdev__ bool isFogVolume() const { return DataType::mGridClass == GridClass::FogVolume; }
2423  __hostdev__ bool isStaggered() const { return DataType::mGridClass == GridClass::Staggered; }
2424  __hostdev__ bool isPointIndex() const { return DataType::mGridClass == GridClass::PointIndex; }
2425  __hostdev__ bool isPointData() const { return DataType::mGridClass == GridClass::PointData; }
2426  __hostdev__ bool isMask() const { return DataType::mGridClass == GridClass::Topology; }
2427  __hostdev__ bool isUnknown() const { return DataType::mGridClass == GridClass::Unknown; }
2428  __hostdev__ bool hasMinMax() const { return DataType::mFlags & static_cast<uint32_t>(GridFlags::HasMinMax); }
2429  __hostdev__ bool hasBBox() const { return DataType::mFlags & static_cast<uint32_t>(GridFlags::HasBBox); }
2430  __hostdev__ bool hasLongGridName() const { return DataType::mFlags & static_cast<uint32_t>(GridFlags::HasLongGridName); }
2431  __hostdev__ bool hasAverage() const { return DataType::mFlags & static_cast<uint32_t>(GridFlags::HasAverage); }
2432  __hostdev__ bool hasStdDeviation() const { return DataType::mFlags & static_cast<uint32_t>(GridFlags::HasStdDeviation); }
2433  __hostdev__ bool isBreadthFirst() const { return DataType::mFlags & static_cast<uint32_t>(GridFlags::IsBreadthFirst); }
2434 
2435  /// @brief return true if the specified node type is layed out breadth-first in memory and has a fixed size.
2436  /// This allows for sequential access to the nodes.
2437  template <typename NodeT>
2438  __hostdev__ bool isSequential() const { return NodeT::FIXED_SIZE && this->isBreadthFirst(); }
2439 
2440  /// @brief return true if the specified node level is layed out breadth-first in memory and has a fixed size.
2441  /// This allows for sequential access to the nodes.
2442  template <int LEVEL>
2443  __hostdev__ bool isSequential() const { return NodeTrait<TreeT,LEVEL>::type::FIXED_SIZE && this->isBreadthFirst(); }
2444 
2445  /// @brief Return a c-string with the name of this grid
2446  __hostdev__ const char* gridName() const
2447  {
2448  if (this->hasLongGridName()) {
2449  const auto &metaData = this->blindMetaData(DataType::mBlindMetadataCount-1);// always the last
2450  NANOVDB_ASSERT(metaData.mDataClass == GridBlindDataClass::GridName);
2451  return metaData.template getBlindData<const char>();
2452  }
2453  return DataType::mGridName;
2454  }
2455 
2456  /// @brief Return a c-string with the name of this grid, truncated to 255 characters
2457  __hostdev__ const char* shortGridName() const { return DataType::mGridName; }
2458 
2459  /// @brief Return checksum of the grid buffer.
2460  __hostdev__ uint64_t checksum() const { return DataType::mChecksum; }
2461 
2462  /// @brief Return true if this grid is empty, i.e. contains no values or nodes.
2463  __hostdev__ bool isEmpty() const { return this->tree().isEmpty(); }
2464 
2465  /// @brief Return the count of blind-data encoded in this grid
2466  __hostdev__ int blindDataCount() const { return DataType::mBlindMetadataCount; }
2467 
2468  /// @brief Return the index of the blind data with specified semantic if found, otherwise -1.
2469  __hostdev__ int findBlindDataForSemantic(GridBlindDataSemantic semantic) const;
2470 
2471  /// @brief Returns a const pointer to the blindData at the specified linear offset.
2472  ///
2473  /// @warning Point might be NULL and the linear offset is assumed to be in the valid range
2474  __hostdev__ const void* blindData(uint32_t n) const
2475  {
2476  if (DataType::mBlindMetadataCount == 0) {
2477  return nullptr;
2478  }
2479  NANOVDB_ASSERT(n < DataType::mBlindMetadataCount);
2480  return this->blindMetaData(n).template getBlindData<void>();
2481  }
2482 
2483  __hostdev__ const GridBlindMetaData& blindMetaData(int n) const { return *DataType::blindMetaData(n); }
2484 
2485 private:
2486  static_assert(sizeof(GridData) % NANOVDB_DATA_ALIGNMENT == 0, "sizeof(GridData) is misaligned");
2487 }; // Class Grid
2488 
2489 template<typename TreeT>
2491 {
2492  for (uint32_t i = 0, n = this->blindDataCount(); i < n; ++i)
2493  if (this->blindMetaData(i).mSemantic == semantic)
2494  return int(i);
2495  return -1;
2496 }
2497 
2498 // ----------------------------> Tree <--------------------------------------
2499 
2500 template<int ROOT_LEVEL = 3>
2501 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) TreeData
2502 {// sizeof(TreeData<3>) == 64B
2503  static_assert(ROOT_LEVEL == 3, "Root level is assumed to be three");
2504  uint64_t mNodeOffset[4];//32B, byte offset from this tree to first leaf, lower, upper and root node
2505  uint32_t mNodeCount[3];// 12B, total number of nodes of type: leaf, lower internal, upper internal
2506  uint32_t mTileCount[3];// 12B, total number of tiles of type: leaf, lower internal, upper internal (node, only active tiles!)
2507  uint64_t mVoxelCount;// 8B, total number of active voxels in the root and all its child nodes.
2508 
2509  template <typename RootT>
2510  __hostdev__ void setRoot(const RootT* root) { mNodeOffset[3] = PtrDiff(root, this); }
2511  template <typename RootT>
2512  __hostdev__ RootT* getRoot() { return PtrAdd<RootT>(this, mNodeOffset[3]); }
2513  template <typename RootT>
2514  __hostdev__ const RootT* getRoot() const { return PtrAdd<RootT>(this, mNodeOffset[3]); }
2515 
2516  template <typename NodeT>
2517  __hostdev__ void setFirstNode(const NodeT* node)
2518  {
2519  mNodeOffset[NodeT::LEVEL] = node ? PtrDiff(node, this) : 0;
2520  }
2521 };
2522 
2523 // ----------------------------> GridTree <--------------------------------------
2524 
2525 /// @brief defines a tree type from a grid type while perserving constness
2526 template<typename GridT>
2527 struct GridTree
2528 {
2529  using Type = typename GridT::TreeType;
2530  using type = typename GridT::TreeType;
2531 };
2532 template<typename GridT>
2533 struct GridTree<const GridT>
2534 {
2535  using Type = const typename GridT::TreeType;
2536  using type = const typename GridT::TreeType;
2537 };
2538 
2539 // ----------------------------> Tree <--------------------------------------
2540 
2541 /// @brief VDB Tree, which is a thin wrapper around a RootNode.
2542 template<typename RootT>
2543 class Tree : private TreeData<RootT::LEVEL>
2544 {
2545  static_assert(RootT::LEVEL == 3, "Tree depth is not supported");
2546  static_assert(RootT::ChildNodeType::LOG2DIM == 5, "Tree configuration is not supported");
2547  static_assert(RootT::ChildNodeType::ChildNodeType::LOG2DIM == 4, "Tree configuration is not supported");
2548  static_assert(RootT::LeafNodeType::LOG2DIM == 3, "Tree configuration is not supported");
2549 
2550 public:
2552  using RootType = RootT;
2553  using LeafNodeType = typename RootT::LeafNodeType;
2554  using ValueType = typename RootT::ValueType;
2555  using BuildType = typename RootT::BuildType;// in rare cases BuildType != ValueType, e.g. then BuildType = ValueMask and ValueType = bool
2556  using CoordType = typename RootT::CoordType;
2558 
2559  using Node3 = RootT;
2560  using Node2 = typename RootT::ChildNodeType;
2561  using Node1 = typename Node2::ChildNodeType;
2563 
2564  /// @brief This class cannot be constructed or deleted
2565  Tree() = delete;
2566  Tree(const Tree&) = delete;
2567  Tree& operator=(const Tree&) = delete;
2568  ~Tree() = delete;
2569 
2570  __hostdev__ DataType* data() { return reinterpret_cast<DataType*>(this); }
2571 
2572  __hostdev__ const DataType* data() const { return reinterpret_cast<const DataType*>(this); }
2573 
2574  /// @brief return memory usage in bytes for the class
2575  __hostdev__ static uint64_t memUsage() { return sizeof(DataType); }
2576 
2577  __hostdev__ RootT& root() { return *DataType::template getRoot<RootT>(); }
2578 
2579  __hostdev__ const RootT& root() const { return *DataType::template getRoot<RootT>(); }
2580 
2581  __hostdev__ AccessorType getAccessor() const { return AccessorType(this->root()); }
2582 
2583  /// @brief Return the value of the given voxel (regardless of state or location in the tree.)
2584  __hostdev__ ValueType getValue(const CoordType& ijk) const { return this->root().getValue(ijk); }
2585 
2586  /// @brief Return the active state of the given voxel (regardless of state or location in the tree.)
2587  __hostdev__ bool isActive(const CoordType& ijk) const { return this->root().isActive(ijk); }
2588 
2589  /// @brief Return true if this tree is empty, i.e. contains no values or nodes
2590  __hostdev__ bool isEmpty() const { return this->root().isEmpty(); }
2591 
2592  /// @brief Combines the previous two methods in a single call
2593  __hostdev__ bool probeValue(const CoordType& ijk, ValueType& v) const { return this->root().probeValue(ijk, v); }
2594 
2595  /// @brief Return a const reference to the background value.
2596  __hostdev__ const ValueType& background() const { return this->root().background(); }
2597 
2598  /// @brief Sets the extrema values of all the active values in this tree, i.e. in all nodes of the tree
2599  __hostdev__ void extrema(ValueType& min, ValueType& max) const;
2600 
2601  /// @brief Return a const reference to the index bounding box of all the active values in this tree, i.e. in all nodes of the tree
2602  __hostdev__ const BBox<CoordType>& bbox() const { return this->root().bbox(); }
2603 
2604  /// @brief Return the total number of active voxels in this tree.
2605  __hostdev__ uint64_t activeVoxelCount() const { return DataType::mVoxelCount; }
2606 
2607  /// @brief Return the total number of active tiles at the specified level of the tree.
2608  ///
2609  /// @details n = 0 corresponds to leaf level tiles.
2610  __hostdev__ const uint32_t& activeTileCount(uint32_t n) const
2611  {
2612  NANOVDB_ASSERT(n < 3);
2613  return DataType::mTileCount[n];
2614  }
2615 
2616  template<typename NodeT>
2617  __hostdev__ uint32_t nodeCount() const
2618  {
2619  static_assert(NodeT::LEVEL < 3, "Invalid NodeT");
2620  return DataType::mNodeCount[NodeT::LEVEL];
2621  }
2622 
2623  __hostdev__ uint32_t nodeCount(int level) const
2624  {
2625  NANOVDB_ASSERT(level < 3);
2626  return DataType::mNodeCount[level];
2627  }
2628 
2629  /// @brief return a pointer to the first node of the specified type
2630  ///
2631  /// @warning Note it may return NULL if no nodes exist
2632  template <typename NodeT>
2634  {
2635  const uint64_t offset = DataType::mNodeOffset[NodeT::LEVEL];
2636  return offset>0 ? PtrAdd<NodeT>(this, offset) : nullptr;
2637  }
2638 
2639  /// @brief return a const pointer to the first node of the specified type
2640  ///
2641  /// @warning Note it may return NULL if no nodes exist
2642  template <typename NodeT>
2643  __hostdev__ const NodeT* getFirstNode() const
2644  {
2645  const uint64_t offset = DataType::mNodeOffset[NodeT::LEVEL];
2646  return offset>0 ? PtrAdd<NodeT>(this, offset) : nullptr;
2647  }
2648 
2649  /// @brief return a pointer to the first node at the specified level
2650  ///
2651  /// @warning Note it may return NULL if no nodes exist
2652  template <int LEVEL>
2655  {
2656  return this->template getFirstNode<typename NodeTrait<RootT,LEVEL>::type>();
2657  }
2658 
2659  /// @brief return a const pointer to the first node of the specified level
2660  ///
2661  /// @warning Note it may return NULL if no nodes exist
2662  template <int LEVEL>
2665  {
2666  return this->template getFirstNode<typename NodeTrait<RootT,LEVEL>::type>();
2667  }
2668 
2669 private:
2670  static_assert(sizeof(DataType) % NANOVDB_DATA_ALIGNMENT == 0, "sizeof(TreeData) is misaligned");
2671 
2672 }; // Tree class
2673 
2674 template<typename RootT>
2676 {
2677  min = this->root().minimum();
2678  max = this->root().maximum();
2679 }
2680 
2681 // --------------------------> RootNode <------------------------------------
2682 
2683 /// @brief Struct with all the member data of the RootNode (useful during serialization of an openvdb RootNode)
2684 ///
2685 /// @note No client code should (or can) interface with this struct so it can safely be ignored!
2686 template<typename ChildT>
2687 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) RootData
2688 {
2689  using ValueT = typename ChildT::ValueType;
2690  using BuildT = typename ChildT::BuildType;// in rare cases BuildType != ValueType, e.g. then BuildType = ValueMask and ValueType = bool
2691  using CoordT = typename ChildT::CoordType;
2692  using StatsT = typename ChildT::FloatType;
2693  static constexpr bool FIXED_SIZE = false;
2694 
2695  /// @brief Return a key based on the coordinates of a voxel
2696 #ifdef USE_SINGLE_ROOT_KEY
2697  using KeyT = uint64_t;
2698  template <typename CoordType>
2699  __hostdev__ static KeyT CoordToKey(const CoordType& ijk)
2700  {
2701  static_assert(sizeof(CoordT) == sizeof(CoordType), "Mismatching sizeof");
2702  static_assert(32 - ChildT::TOTAL <= 21, "Cannot use 64 bit root keys");
2703  return (KeyT(uint32_t(ijk[2]) >> ChildT::TOTAL)) | // z is the lower 21 bits
2704  (KeyT(uint32_t(ijk[1]) >> ChildT::TOTAL) << 21) | // y is the middle 21 bits
2705  (KeyT(uint32_t(ijk[0]) >> ChildT::TOTAL) << 42); // x is the upper 21 bits
2706  }
2707  __hostdev__ static CoordT KeyToCoord(const KeyT& key)
2708  {
2709  static constexpr uint64_t MASK = (1u << 21) - 1;
2710  return CoordT(((key >> 42) & MASK) << ChildT::TOTAL,
2711  ((key >> 21) & MASK) << ChildT::TOTAL,
2712  (key & MASK) << ChildT::TOTAL);
2713  }
2714 #else
2715  using KeyT = CoordT;
2716  __hostdev__ static KeyT CoordToKey(const CoordT& ijk) { return ijk & ~ChildT::MASK; }
2717  __hostdev__ static CoordT KeyToCoord(const KeyT& key) { return key; }
2718 #endif
2719  BBox<CoordT> mBBox; // 24B. AABB if active values in index space.
2720  uint32_t mTableSize; // 4B. number of tiles and child pointers in the root node
2721 
2722  ValueT mBackground; // background value, i.e. value of any unset voxel
2723  ValueT mMinimum; // typically 4B, minmum of all the active values
2724  ValueT mMaximum; // typically 4B, maximum of all the active values
2725  StatsT mAverage; // typically 4B, average of all the active values in this node and its child nodes
2726  StatsT mStdDevi; // typically 4B, standard deviation of all the active values in this node and its child nodes
2727 
2728  struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) Tile
2729  {
2730  template <typename CoordType>
2731  __hostdev__ void setChild(const CoordType& k, const ChildT *ptr, const RootData *data)
2732  {
2733  key = CoordToKey(k);
2734  child = PtrDiff(ptr, data);
2735  }
2736  template <typename CoordType, typename ValueType>
2737  __hostdev__ void setValue(const CoordType& k, bool s, const ValueType &v)
2738  {
2739  key = CoordToKey(k);
2740  state = s;
2741  value = v;
2742  child = 0;
2743  }
2744  __hostdev__ bool isChild() const { return child; }
2745  __hostdev__ CoordT origin() const { return KeyToCoord(key); }
2746  KeyT key; // USE_SINGLE_ROOT_KEY ? 8B : 12B
2747  int64_t child; // 8B. signed byte offset from this node to the child node. 0 means it is a constant tile, so use value.
2748  uint32_t state; // 4B. state of tile value
2749  ValueT value; // value of tile (i.e. no child node)
2750  }; // Tile
2751 
2752  /// @brief Returns a non-const reference to the tile at the specified linear offset.
2753  ///
2754  /// @warning The linear offset is assumed to be in the valid range
2755  __hostdev__ const Tile* tile(uint32_t n) const
2756  {
2757  NANOVDB_ASSERT(n < mTableSize);
2758  return reinterpret_cast<const Tile*>(this + 1) + n;
2759  }
2760  __hostdev__ Tile* tile(uint32_t n)
2761  {
2762  NANOVDB_ASSERT(n < mTableSize);
2763  return reinterpret_cast<Tile*>(this + 1) + n;
2764  }
2765 
2766  /// @brief Returns a const reference to the child node in the specified tile.
2767  ///
2768  /// @warning A child node is assumed to exist in the specified tile
2769  __hostdev__ ChildT* getChild(const Tile* tile)
2770  {
2771  NANOVDB_ASSERT(tile->child);
2772  return PtrAdd<ChildT>(this, tile->child);
2773  }
2774  __hostdev__ const ChildT* getChild(const Tile* tile) const
2775  {
2776  NANOVDB_ASSERT(tile->child);
2777  return PtrAdd<ChildT>(this, tile->child);
2778  }
2779 
2780  __hostdev__ const ValueT& getMin() const { return mMinimum; }
2781  __hostdev__ const ValueT& getMax() const { return mMaximum; }
2782  __hostdev__ const StatsT& average() const { return mAverage; }
2783  __hostdev__ const StatsT& stdDeviation() const { return mStdDevi; }
2784 
2785  __hostdev__ void setMin(const ValueT& v) { mMinimum = v; }
2786  __hostdev__ void setMax(const ValueT& v) { mMaximum = v; }
2787  __hostdev__ void setAvg(const StatsT& v) { mAverage = v; }
2788  __hostdev__ void setDev(const StatsT& v) { mStdDevi = v; }
2789 
2790  /// @brief This class cannot be constructed or deleted
2791  RootData() = delete;
2792  RootData(const RootData&) = delete;
2793  RootData& operator=(const RootData&) = delete;
2794  ~RootData() = delete;
2795 }; // RootData
2796 
2797 /// @brief Top-most node of the VDB tree structure.
2798 template<typename ChildT>
2799 class RootNode : private RootData<ChildT>
2800 {
2801 public:
2803  using LeafNodeType = typename ChildT::LeafNodeType;
2804  using ChildNodeType = ChildT;
2805  using RootType = RootNode<ChildT>;// this allows RootNode to behave like a Tree
2806 
2807  using ValueType = typename DataType::ValueT;
2808  using FloatType = typename DataType::StatsT;
2809  using BuildType = typename DataType::BuildT;// in rare cases BuildType != ValueType, e.g. then BuildType = ValueMask and ValueType = bool
2810 
2811  using CoordType = typename ChildT::CoordType;
2813  using Tile = typename DataType::Tile;
2814  static constexpr bool FIXED_SIZE = DataType::FIXED_SIZE;
2815 
2816  static constexpr uint32_t LEVEL = 1 + ChildT::LEVEL; // level 0 = leaf
2817 
2818  /// @brief This class cannot be constructed or deleted
2819  RootNode() = delete;
2820  RootNode(const RootNode&) = delete;
2821  RootNode& operator=(const RootNode&) = delete;
2822  ~RootNode() = delete;
2823 
2825 
2826  __hostdev__ DataType* data() { return reinterpret_cast<DataType*>(this); }
2827 
2828  __hostdev__ const DataType* data() const { return reinterpret_cast<const DataType*>(this); }
2829 
2830  /// @brief Return a const reference to the index bounding box of all the active values in this tree, i.e. in all nodes of the tree
2831  __hostdev__ const BBox<CoordType>& bbox() const { return DataType::mBBox; }
2832 
2833  /// @brief Return the total number of active voxels in the root and all its child nodes.
2834 
2835  /// @brief Return a const reference to the background value, i.e. the value associated with
2836  /// any coordinate location that has not been set explicitly.
2837  __hostdev__ const ValueType& background() const { return DataType::mBackground; }
2838 
2839  /// @brief Return the number of tiles encoded in this root node
2840  __hostdev__ const uint32_t& tileCount() const { return DataType::mTableSize; }
2841 
2842  /// @brief Return a const reference to the minimum active value encoded in this root node and any of its child nodes
2843  __hostdev__ const ValueType& minimum() const { return this->getMin(); }
2844 
2845  /// @brief Return a const reference to the maximum active value encoded in this root node and any of its child nodes
2846  __hostdev__ const ValueType& maximum() const { return this->getMax(); }
2847 
2848  /// @brief Return a const reference to the average of all the active values encoded in this root node and any of its child nodes
2849  __hostdev__ const FloatType& average() const { return DataType::mAverage; }
2850 
2851  /// @brief Return the variance of all the active values encoded in this root node and any of its child nodes
2852  __hostdev__ FloatType variance() const { return DataType::mStdDevi * DataType::mStdDevi; }
2853 
2854  /// @brief Return a const reference to the standard deviation of all the active values encoded in this root node and any of its child nodes
2855  __hostdev__ const FloatType& stdDeviation() const { return DataType::mStdDevi; }
2856 
2857  /// @brief Return the expected memory footprint in bytes with the specified number of tiles
2858  __hostdev__ static uint64_t memUsage(uint32_t tableSize) { return sizeof(RootNode) + tableSize * sizeof(Tile); }
2859 
2860  /// @brief Return the actual memory footprint of this root node
2861  __hostdev__ uint64_t memUsage() const { return sizeof(RootNode) + DataType::mTableSize * sizeof(Tile); }
2862 
2863  /// @brief Return the value of the given voxel
2865  {
2866  if (const Tile* tile = this->findTile(ijk)) {
2867  return tile->isChild() ? this->getChild(tile)->getValue(ijk) : tile->value;
2868  }
2869  return DataType::mBackground;
2870  }
2871 
2872  __hostdev__ bool isActive(const CoordType& ijk) const
2873  {
2874  if (const Tile* tile = this->findTile(ijk)) {
2875  return tile->isChild() ? this->getChild(tile)->isActive(ijk) : tile->state;
2876  }
2877  return false;
2878  }
2879 
2880  /// @brief Return true if this RootNode is empty, i.e. contains no values or nodes
2881  __hostdev__ bool isEmpty() const { return DataType::mTableSize == uint32_t(0); }
2882 
2883  __hostdev__ bool probeValue(const CoordType& ijk, ValueType& v) const
2884  {
2885  if (const Tile* tile = this->findTile(ijk)) {
2886  if (tile->isChild()) {
2887  const auto *child = this->getChild(tile);
2888  return child->probeValue(ijk, v);
2889  }
2890  v = tile->value;
2891  return tile->state;
2892  }
2893  v = DataType::mBackground;
2894  return false;
2895  }
2896 
2897  __hostdev__ const LeafNodeType* probeLeaf(const CoordType& ijk) const
2898  {
2899  const Tile* tile = this->findTile(ijk);
2900  if (tile && tile->isChild()) {
2901  const auto *child = this->getChild(tile);
2902  return child->probeLeaf(ijk);
2903  }
2904  return nullptr;
2905  }
2906 
2907 private:
2908  static_assert(sizeof(DataType) % NANOVDB_DATA_ALIGNMENT == 0, "sizeof(RootData) is misaligned");
2909  static_assert(sizeof(typename DataType::Tile) % NANOVDB_DATA_ALIGNMENT == 0, "sizeof(RootData::Tile) is misaligned");
2910 
2911  template<typename, int, int, int>
2912  friend class ReadAccessor;
2913 
2914  template<typename>
2915  friend class Tree;
2916 
2917  /// @brief Private method to find a Tile of this root node by means of binary-search. This is obviously
2918  /// much slower then direct lookup into a linear array (as in the other nodes) which is exactly
2919  /// why it is important to use the ReadAccessor which amortizes this overhead by node caching and
2920  /// inverse tree traversal!
2921  __hostdev__ const Tile* findTile(const CoordType& ijk) const
2922  {
2923  const Tile* tiles = reinterpret_cast<const Tile*>(this + 1);
2924  const auto key = DataType::CoordToKey(ijk);
2925 #if 1 // switch between linear and binary seach
2926  for (uint32_t i = 0; i < DataType::mTableSize; ++i) {
2927  if (tiles[i].key == key) return &tiles[i];
2928  }
2929 #else// do not enable binary search if tiles are not guaranteed to be sorted!!!!!!
2930  // binary-search of pre-sorted elements
2931  int32_t low = 0, high = DataType::mTableSize; // low is inclusive and high is exclusive
2932  while (low != high) {
2933  int mid = low + ((high - low) >> 1);
2934  const Tile* tile = &tiles[mid];
2935  if (tile->key == key) {
2936  return tile;
2937  } else if (tile->key < key) {
2938  low = mid + 1;
2939  } else {
2940  high = mid;
2941  }
2942  }
2943 #endif
2944  return nullptr;
2945  }
2946 
2947  /// @brief Private method to return node information and update a ReadAccessor
2948  template<typename AccT>
2949  __hostdev__ typename AccT::NodeInfo getNodeInfoAndCache(const CoordType& ijk, const AccT& acc) const
2950  {
2951  using NodeInfoT = typename AccT::NodeInfo;
2952  if (const Tile* tile = this->findTile(ijk)) {
2953  if (tile->isChild()) {
2954  const auto *child = this->getChild(tile);
2955  acc.insert(ijk, child);
2956  return child->getNodeInfoAndCache(ijk, acc);
2957  }
2958  return NodeInfoT{LEVEL, ChildT::dim(), tile->value, tile->value, tile->value,
2959  0, tile->origin(), tile->origin() + CoordType(ChildT::DIM)};
2960  }
2961  return NodeInfoT{LEVEL, ChildT::dim(), this->minimum(), this->maximum(),
2962  this->average(), this->stdDeviation(), this->bbox()[0], this->bbox()[1]};
2963  }
2964 
2965  /// @brief Private method to return a voxel value and update a ReadAccessor
2966  template<typename AccT>
2967  __hostdev__ ValueType getValueAndCache(const CoordType& ijk, const AccT& acc) const
2968  {
2969  if (const Tile* tile = this->findTile(ijk)) {
2970  if (tile->isChild()) {
2971  const auto *child = this->getChild(tile);
2972  acc.insert(ijk, child);
2973  return child->getValueAndCache(ijk, acc);
2974  }
2975  return tile->value;
2976  }
2977  return DataType::mBackground;
2978  }
2979 
2980  template<typename AccT>
2981  __hostdev__ bool isActiveAndCache(const CoordType& ijk, const AccT& acc) const
2982  {
2983  const Tile* tile = this->findTile(ijk);
2984  if (tile && tile->isChild()) {
2985  const auto *child = this->getChild(tile);
2986  acc.insert(ijk, child);
2987  return child->isActiveAndCache(ijk, acc);
2988  }
2989  return false;
2990  }
2991 
2992  template<typename AccT>
2993  __hostdev__ bool probeValueAndCache(const CoordType& ijk, ValueType& v, const AccT& acc) const
2994  {
2995  if (const Tile* tile = this->findTile(ijk)) {
2996  if (tile->isChild()) {
2997  const auto *child = this->getChild(tile);
2998  acc.insert(ijk, child);
2999  return child->probeValueAndCache(ijk, v, acc);
3000  }
3001  v = tile->value;
3002  return tile->state;
3003  }
3004  v = DataType::mBackground;
3005  return false;
3006  }
3007 
3008  template<typename AccT>
3009  __hostdev__ const LeafNodeType* probeLeafAndCache(const CoordType& ijk, const AccT& acc) const
3010  {
3011  const Tile* tile = this->findTile(ijk);
3012  if (tile && tile->isChild()) {
3013  const auto *child = this->getChild(tile);
3014  acc.insert(ijk, child);
3015  return child->probeLeafAndCache(ijk, acc);
3016  }
3017  return nullptr;
3018  }
3019 
3020  template<typename RayT, typename AccT>
3021  __hostdev__ uint32_t getDimAndCache(const CoordType& ijk, const RayT& ray, const AccT& acc) const
3022  {
3023  if (const Tile* tile = this->findTile(ijk)) {
3024  if (tile->isChild()) {
3025  const auto *child = this->getChild(tile);
3026  acc.insert(ijk, child);
3027  return child->getDimAndCache(ijk, ray, acc);
3028  }
3029  return 1 << ChildT::TOTAL; //tile value
3030  }
3031  return ChildNodeType::dim(); // background
3032  }
3033 }; // RootNode class
3034 
3035 // After the RootNode the memory layout is assumed to be the sorted Tiles
3036 
3037 // --------------------------> InternalNode <------------------------------------
3038 
3039 /// @brief Struct with all the member data of the InternalNode (useful during serialization of an openvdb InternalNode)
3040 ///
3041 /// @note No client code should (or can) interface with this struct so it can safely be ignored!
3042 template<typename ChildT, uint32_t LOG2DIM>
3043 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) InternalData
3044 {
3045  using ValueT = typename ChildT::ValueType;
3046  using BuildT = typename ChildT::BuildType;// in rare cases BuildType != ValueType, e.g. then BuildType = ValueMask and ValueType = bool
3047  using StatsT = typename ChildT::FloatType;
3048  using CoordT = typename ChildT::CoordType;
3049  using MaskT = typename ChildT::template MaskType<LOG2DIM>;
3050  static constexpr bool FIXED_SIZE = true;
3051 
3052  union Tile
3053  {
3055  int64_t child;//signed 64 bit byte offset relative to the InternalData!!
3056  /// @brief This class cannot be constructed or deleted
3057  Tile() = delete;
3058  Tile(const Tile&) = delete;
3059  Tile& operator=(const Tile&) = delete;
3060  ~Tile() = delete;
3061  };
3062 
3063  BBox<CoordT> mBBox; // 24B. node bounding box. |
3064  uint64_t mFlags; // 8B. node flags. | 32B aligned
3065  MaskT mValueMask; // LOG2DIM(5): 4096B, LOG2DIM(4): 512B | 32B aligned
3066  MaskT mChildMask; // LOG2DIM(5): 4096B, LOG2DIM(4): 512B | 32B aligned
3067 
3068  ValueT mMinimum; // typically 4B
3069  ValueT mMaximum; // typically 4B
3070  StatsT mAverage; // typically 4B, average of all the active values in this node and its child nodes
3071  StatsT mStdDevi; // typically 4B, standard deviation of all the active values in this node and its child nodes
3072  alignas(32) Tile mTable[1u << (3 * LOG2DIM)]; // sizeof(ValueT) x (16*16*16 or 32*32*32)
3073 
3074  __hostdev__ void setChild(uint32_t n, const void *ptr)
3075  {
3076  NANOVDB_ASSERT(mChildMask.isOn(n));
3077  mTable[n].child = PtrDiff(ptr, this);
3078  }
3079 
3080  template <typename ValueT>
3081  __hostdev__ void setValue(uint32_t n, const ValueT &v)
3082  {
3083  NANOVDB_ASSERT(!mChildMask.isOn(n));
3084  mTable[n].value = v;
3085  }
3086 
3087  /// @brief Returns a pointer to the child node at the specifed linear offset.
3088  __hostdev__ ChildT* getChild(uint32_t n)
3089  {
3090  NANOVDB_ASSERT(mChildMask.isOn(n));
3091  return PtrAdd<ChildT>(this, mTable[n].child);
3092  }
3093  __hostdev__ const ChildT* getChild(uint32_t n) const
3094  {
3095  NANOVDB_ASSERT(mChildMask.isOn(n));
3096  return PtrAdd<ChildT>(this, mTable[n].child);
3097  }
3098 
3099  template <typename T>
3100  __hostdev__ void setOrigin(const T& ijk) { mBBox[0] = ijk; }
3101 
3102  __hostdev__ const ValueT& getMin() const { return mMinimum; }
3103  __hostdev__ const ValueT& getMax() const { return mMaximum; }
3104  __hostdev__ const StatsT& average() const { return mAverage; }
3105  __hostdev__ const StatsT& stdDeviation() const { return mStdDevi; }
3106 
3107  __hostdev__ void setMin(const ValueT& v) { mMinimum = v; }
3108  __hostdev__ void setMax(const ValueT& v) { mMaximum = v; }
3109  __hostdev__ void setAvg(const StatsT& v) { mAverage = v; }
3110  __hostdev__ void setDev(const StatsT& v) { mStdDevi = v; }
3111 
3112  /// @brief This class cannot be constructed or deleted
3113  InternalData() = delete;
3114  InternalData(const InternalData&) = delete;
3116  ~InternalData() = delete;
3117 }; // InternalData
3118 
3119 /// @brief Internal nodes of a VDB treedim(),
3120 template<typename ChildT, uint32_t Log2Dim = ChildT::LOG2DIM + 1>
3121 class InternalNode : private InternalData<ChildT, Log2Dim>
3122 {
3123 public:
3125  using ValueType = typename DataType::ValueT;
3126  using FloatType = typename DataType::StatsT;
3127  using BuildType = typename DataType::BuildT; // in rare cases BuildType != ValueType, e.g. then BuildType = ValueMask and ValueType = bool
3128  using LeafNodeType = typename ChildT::LeafNodeType;
3129  using ChildNodeType = ChildT;
3130  using CoordType = typename ChildT::CoordType;
3131  static constexpr bool FIXED_SIZE = DataType::FIXED_SIZE;
3132  template<uint32_t LOG2>
3133  using MaskType = typename ChildT::template MaskType<LOG2>;
3134 
3135  static constexpr uint32_t LOG2DIM = Log2Dim;
3136  static constexpr uint32_t TOTAL = LOG2DIM + ChildT::TOTAL; // dimension in index space
3137  static constexpr uint32_t DIM = 1u << TOTAL; // number of voxels along each axis of this node
3138  static constexpr uint32_t SIZE = 1u << (3 * LOG2DIM); // number of tile values (or child pointers)
3139  static constexpr uint32_t MASK = (1u << TOTAL) - 1u;
3140  static constexpr uint32_t LEVEL = 1 + ChildT::LEVEL; // level 0 = leaf
3141  static constexpr uint64_t NUM_VALUES = uint64_t(1) << (3 * TOTAL); // total voxel count represented by this node
3142 
3143  /// @brief This class cannot be constructed or deleted
3144  InternalNode() = delete;
3145  InternalNode(const InternalNode&) = delete;
3147  ~InternalNode() = delete;
3148 
3149  __hostdev__ DataType* data() { return reinterpret_cast<DataType*>(this); }
3150 
3151  __hostdev__ const DataType* data() const { return reinterpret_cast<const DataType*>(this); }
3152 
3153  /// @brief Return the dimension, in voxel units, of this internal node (typically 8*16 or 8*16*32)
3154  __hostdev__ static uint32_t dim() { return 1u << TOTAL; }
3155 
3156  /// @brief Return memory usage in bytes for the class
3157  __hostdev__ static size_t memUsage() { return sizeof(DataType); }
3158 
3159  /// @brief Return a const reference to the bit mask of active voxels in this internal node
3160  __hostdev__ const MaskType<LOG2DIM>& valueMask() const { return DataType::mValueMask; }
3161 
3162  /// @brief Return a const reference to the bit mask of child nodes in this internal node
3163  __hostdev__ const MaskType<LOG2DIM>& childMask() const { return DataType::mChildMask; }
3164 
3165  /// @brief Return the origin in index space of this leaf node
3166  __hostdev__ CoordType origin() const { return DataType::mBBox.min() & ~MASK; }
3167 
3168  /// @brief Return a const reference to the minimum active value encoded in this internal node and any of its child nodes
3169  __hostdev__ const ValueType& minimum() const { return this->getMin(); }
3170 
3171  /// @brief Return a const reference to the maximum active value encoded in this internal node and any of its child nodes
3172  __hostdev__ const ValueType& maximum() const { return this->getMax(); }
3173 
3174  /// @brief Return a const reference to the average of all the active values encoded in this internal node and any of its child nodes
3175  __hostdev__ const FloatType& average() const { return DataType::mAverage; }
3176 
3177  /// @brief Return the variance of all the active values encoded in this internal node and any of its child nodes
3178  __hostdev__ FloatType variance() const { return DataType::mStdDevi*DataType::mStdDevi; }
3179 
3180  /// @brief Return a const reference to the standard deviation of all the active values encoded in this internal node and any of its child nodes
3181  __hostdev__ const FloatType& stdDeviation() const { return DataType::mStdDevi; }
3182 
3183  /// @brief Return a const reference to the bounding box in index space of active values in this internal node and any of its child nodes
3184  __hostdev__ const BBox<CoordType>& bbox() const { return DataType::mBBox; }
3185 
3186  /// @brief Return the value of the given voxel
3188  {
3189  const uint32_t n = CoordToOffset(ijk);
3190  return DataType::mChildMask.isOn(n) ? this->getChild(n)->getValue(ijk) : DataType::mTable[n].value;
3191  }
3192 
3193  __hostdev__ bool isActive(const CoordType& ijk) const
3194  {
3195  const uint32_t n = CoordToOffset(ijk);
3196  return DataType::mChildMask.isOn(n) ? this->getChild(n)->isActive(ijk) : DataType::mValueMask.isOn(n);
3197  }
3198 
3199  __hostdev__ bool probeValue(const CoordType& ijk, ValueType& v) const
3200  {
3201  const uint32_t n = CoordToOffset(ijk);
3202  if (DataType::mChildMask.isOn(n))
3203  return this->getChild(n)->probeValue(ijk, v);
3204  v = DataType::mTable[n].value;
3205  return DataType::mValueMask.isOn(n);
3206  }
3207 
3208  __hostdev__ const LeafNodeType* probeLeaf(const CoordType& ijk) const
3209  {
3210  const uint32_t n = CoordToOffset(ijk);
3211  if (DataType::mChildMask.isOn(n))
3212  return this->getChild(n)->probeLeaf(ijk);
3213  return nullptr;
3214  }
3215 
3216  /// @brief Return the linear offset corresponding to the given coordinate
3217  __hostdev__ static uint32_t CoordToOffset(const CoordType& ijk)
3218  {
3219 #if 0
3220  return (((ijk[0] & MASK) >> ChildT::TOTAL) << (2 * LOG2DIM)) +
3221  (((ijk[1] & MASK) >> ChildT::TOTAL) << (LOG2DIM)) +
3222  ((ijk[2] & MASK) >> ChildT::TOTAL);
3223 #else
3224  return (((ijk[0] & MASK) >> ChildT::TOTAL) << (2 * LOG2DIM)) |
3225  (((ijk[1] & MASK) >> ChildT::TOTAL) << (LOG2DIM)) |
3226  ((ijk[2] & MASK) >> ChildT::TOTAL);
3227 #endif
3228  }
3229 
3230  /// @return the local coordinate of the n'th tile or child node
3232  {
3233  NANOVDB_ASSERT(n < SIZE);
3234  const uint32_t m = n & ((1 << 2 * LOG2DIM) - 1);
3235  return Coord(n >> 2 * LOG2DIM, m >> LOG2DIM, m & ((1 << LOG2DIM) - 1));
3236  }
3237 
3238  /// @brief modifies local coordinates to global coordinates of a tile or child node
3240  {
3241  ijk <<= ChildT::TOTAL;
3242  ijk += this->origin();
3243  }
3244 
3246  {
3248  this->localToGlobalCoord(ijk);
3249  return ijk;
3250  }
3251 
3252  /// @brief Retrun true if this node or any of its child nodes contain active values
3253  __hostdev__ bool isActive() const
3254  {
3255  return DataType::mFlags & uint32_t(2);
3256  }
3257 
3258 private:
3259  static_assert(sizeof(DataType) % NANOVDB_DATA_ALIGNMENT == 0, "sizeof(InternalData) is misaligned");
3260  //static_assert(offsetof(DataType, mTable) % 32 == 0, "InternalData::mTable is misaligned");
3261 
3262  template<typename, int, int, int>
3263  friend class ReadAccessor;
3264 
3265  template<typename>
3266  friend class RootNode;
3267  template<typename, uint32_t>
3268  friend class InternalNode;
3269 
3270  /// @brief Private read access method used by the ReadAccessor
3271  template<typename AccT>
3272  __hostdev__ ValueType getValueAndCache(const CoordType& ijk, const AccT& acc) const
3273  {
3274  const uint32_t n = CoordToOffset(ijk);
3275  if (!DataType::mChildMask.isOn(n))
3276  return DataType::mTable[n].value;
3277  const ChildT* child = this->getChild(n);
3278  acc.insert(ijk, child);
3279  return child->getValueAndCache(ijk, acc);
3280  }
3281 
3282  template<typename AccT>
3283  __hostdev__ typename AccT::NodeInfo getNodeInfoAndCache(const CoordType& ijk, const AccT& acc) const
3284  {
3285  using NodeInfoT = typename AccT::NodeInfo;
3286  const uint32_t n = CoordToOffset(ijk);
3287  if (!DataType::mChildMask.isOn(n)) {
3288  return NodeInfoT{LEVEL, this->dim(), this->minimum(), this->maximum(), this->average(),
3289  this->stdDeviation(), this->bbox()[0], this->bbox()[1]};
3290  }
3291  const ChildT* child = this->getChild(n);
3292  acc.insert(ijk, child);
3293  return child->getNodeInfoAndCache(ijk, acc);
3294  }
3295 
3296  template<typename AccT>
3297  __hostdev__ bool isActiveAndCache(const CoordType& ijk, const AccT& acc) const
3298  {
3299  const uint32_t n = CoordToOffset(ijk);
3300  if (!DataType::mChildMask.isOn(n))
3301  return DataType::mValueMask.isOn(n);
3302  const ChildT* child = this->getChild(n);
3303  acc.insert(ijk, child);
3304  return child->isActiveAndCache(ijk, acc);
3305  }
3306 
3307  template<typename AccT>
3308  __hostdev__ bool probeValueAndCache(const CoordType& ijk, ValueType& v, const AccT& acc) const
3309  {
3310  const uint32_t n = CoordToOffset(ijk);
3311  if (!DataType::mChildMask.isOn(n)) {
3312  v = DataType::mTable[n].value;
3313  return DataType::mValueMask.isOn(n);
3314  }
3315  const ChildT* child = this->getChild(n);
3316  acc.insert(ijk, child);
3317  return child->probeValueAndCache(ijk, v, acc);
3318  }
3319 
3320  template<typename AccT>
3321  __hostdev__ const LeafNodeType* probeLeafAndCache(const CoordType& ijk, const AccT& acc) const
3322  {
3323  const uint32_t n = CoordToOffset(ijk);
3324  if (!DataType::mChildMask.isOn(n))
3325  return nullptr;
3326  const ChildT* child = this->getChild(n);
3327  acc.insert(ijk, child);
3328  return child->probeLeafAndCache(ijk, acc);
3329  }
3330 
3331  template<typename RayT, typename AccT>
3332  __hostdev__ uint32_t getDimAndCache(const CoordType& ijk, const RayT& ray, const AccT& acc) const
3333  {
3334  if (DataType::mFlags & uint32_t(1))
3335  this->dim(); //ship this node if first bit is set
3336  //if (!ray.intersects( this->bbox() )) return 1<<TOTAL;
3337 
3338  const uint32_t n = CoordToOffset(ijk);
3339  if (DataType::mChildMask.isOn(n)) {
3340  const ChildT* child = this->getChild(n);
3341  acc.insert(ijk, child);
3342  return child->getDimAndCache(ijk, ray, acc);
3343  }
3344  return ChildNodeType::dim(); // tile value
3345  }
3346 
3347 }; // InternalNode class
3348 
3349 // --------------------------> LeafNode <------------------------------------
3350 
3351 /// @brief Stuct with all the member data of the LeafNode (useful during serialization of an openvdb LeafNode)
3352 ///
3353 /// @note No client code should (or can) interface with this struct so it can safely be ignored!
3354 template<typename ValueT, typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3355 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData
3356 {
3357  static_assert(sizeof(CoordT) == sizeof(Coord), "Mismatching sizeof");
3358  static_assert(sizeof(MaskT<LOG2DIM>) == sizeof(Mask<LOG2DIM>), "Mismatching sizeof");
3359  using ValueType = ValueT;
3360  using BuildType = ValueT;
3362  using ArrayType = ValueT;// type used for the internal mValue array
3363  static constexpr bool FIXED_SIZE = true;
3364 
3365  CoordT mBBoxMin; // 12B.
3366  uint8_t mBBoxDif[3]; // 3B.
3367  uint8_t mFlags; // 1B.
3368  MaskT<LOG2DIM> mValueMask; // LOG2DIM(3): 64B.
3369 
3370  ValueType mMinimum; // typically 4B
3371  ValueType mMaximum; // typically 4B
3372  FloatType mAverage; // typically 4B, average of all the active values in this node and its child nodes
3373  FloatType mStdDevi; // typically 4B, standard deviation of all the active values in this node and its child nodes
3374  alignas(32) ValueType mValues[1u << 3 * LOG2DIM];
3375 
3376  //__hostdev__ const ValueType* values() const { return mValues; }
3377  __hostdev__ ValueType getValue(uint32_t i) const { return mValues[i]; }
3378  __hostdev__ void setValueOnly(uint32_t offset, const ValueType& value) { mValues[offset] = value; }
3379  __hostdev__ void setValue(uint32_t offset, const ValueType& value)
3380  {
3381  mValueMask.setOn(offset);
3382  mValues[offset] = value;
3383  }
3384 
3385  __hostdev__ ValueType getMin() const { return mMinimum; }
3386  __hostdev__ ValueType getMax() const { return mMaximum; }
3387  __hostdev__ FloatType getAvg() const { return mAverage; }
3388  __hostdev__ FloatType getDev() const { return mStdDevi; }
3389 
3390  __hostdev__ void setMin(const ValueType& v) { mMinimum = v; }
3391  __hostdev__ void setMax(const ValueType& v) { mMaximum = v; }
3392  __hostdev__ void setAvg(const FloatType& v) { mAverage = v; }
3393  __hostdev__ void setDev(const FloatType& v) { mStdDevi = v; }
3394 
3395  template <typename T>
3396  __hostdev__ void setOrigin(const T& ijk) { mBBoxMin = ijk; }
3397 
3398  /// @brief This class cannot be constructed or deleted
3399  LeafData() = delete;
3400  LeafData(const LeafData&) = delete;
3401  LeafData& operator=(const LeafData&) = delete;
3402  ~LeafData() = delete;
3403 }; // LeafData<ValueT>
3404 
3405 /// @brief Base-class for quantized float leaf nodes
3406 template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3407 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafFnBase
3408 {
3409  static_assert(sizeof(CoordT) == sizeof(Coord), "Mismatching sizeof");
3410  static_assert(sizeof(MaskT<LOG2DIM>) == sizeof(Mask<LOG2DIM>), "Mismatching sizeof");
3411  using ValueType = float;
3412  using FloatType = float;
3413 
3414  CoordT mBBoxMin; // 12B.
3415  uint8_t mBBoxDif[3]; // 3B.
3416  uint8_t mFlags; // 1B.
3417  MaskT<LOG2DIM> mValueMask; // LOG2DIM(3): 64B.
3418 
3419  float mMinimum; // 4B - minimum of ALL values in this node
3420  float mQuantum; // = (max - min)/15 4B
3421  uint16_t mMin, mMax, mAvg, mDev;// quantized representations of statistics of active values
3422 
3423  void init(float min, float max, uint8_t bitWidth)
3424  {
3425  mMinimum = min;
3426  mQuantum = (max - min)/float((1 << bitWidth)-1);
3427  }
3428 
3429  /// @brief return the quantized minimum of the active values in this node
3430  __hostdev__ float getMin() const { return mMin*mQuantum + mMinimum; }
3431 
3432  /// @brief return the quantized maximum of the active values in this node
3433  __hostdev__ float getMax() const { return mMax*mQuantum + mMinimum; }
3434 
3435  /// @brief return the quantized average of the active values in this node
3436  __hostdev__ float getAvg() const { return mAvg*mQuantum + mMinimum; }
3437  /// @brief return the quantized standard deviation of the active values in this node
3438 
3439  /// @note 0 <= StdDev <= max-min or 0 <= StdDev/(max-min) <= 1
3440  __hostdev__ float getDev() const { return mDev*mQuantum; }
3441 
3442  /// @note min <= X <= max or 0 <= (X-min)/(min-max) <= 1
3443  __hostdev__ void setMin(float min) { mMin = uint16_t((min - mMinimum)/mQuantum + 0.5f); }
3444 
3445  /// @note min <= X <= max or 0 <= (X-min)/(min-max) <= 1
3446  __hostdev__ void setMax(float max) { mMax = uint16_t((max - mMinimum)/mQuantum + 0.5f); }
3447 
3448  /// @note min <= avg <= max or 0 <= (avg-min)/(min-max) <= 1
3449  __hostdev__ void setAvg(float avg) { mAvg = uint16_t((avg - mMinimum)/mQuantum + 0.5f); }
3450 
3451  /// @note 0 <= StdDev <= max-min or 0 <= StdDev/(max-min) <= 1
3452  __hostdev__ void setDev(float dev) { mDev = uint16_t(dev/mQuantum + 0.5f); }
3453 
3454  template <typename T>
3455  __hostdev__ void setOrigin(const T& ijk) { mBBoxMin = ijk; }
3456 };// LeafFnBase
3457 
3458 /// @brief Stuct with all the member data of the LeafNode (useful during serialization of an openvdb LeafNode)
3459 ///
3460 /// @note No client code should (or can) interface with this struct so it can safely be ignored!
3461 template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3462 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData<Fp4, CoordT, MaskT, LOG2DIM>
3463  : public LeafFnBase<CoordT, MaskT, LOG2DIM>
3464 {
3466  using BuildType = Fp4;
3467  using ArrayType = uint8_t;// type used for the internal mValue array
3468  static constexpr bool FIXED_SIZE = true;
3469  alignas(32) uint8_t mCode[1u << (3 * LOG2DIM - 1)];
3470 
3471  __hostdev__ static constexpr uint8_t bitWidth() { return 4u; }
3472  __hostdev__ float getValue(uint32_t i) const
3473  {
3474 #if 0
3475  const uint8_t c = mCode[i>>1];
3476  return ( (i&1) ? c >> 4 : c & uint8_t(15) )*BaseT::mQuantum + BaseT::mMinimum;
3477 #else
3478  return ((mCode[i>>1] >> ((i&1)<<2)) & uint8_t(15))*BaseT::mQuantum + BaseT::mMinimum;
3479 #endif
3480  }
3481 
3482  /// @brief This class cannot be constructed or deleted
3483  LeafData() = delete;
3484  LeafData(const LeafData&) = delete;
3485  LeafData& operator=(const LeafData&) = delete;
3486  ~LeafData() = delete;
3487 }; // LeafData<Fp4>
3488 
3489 template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3490 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData<Fp8, CoordT, MaskT, LOG2DIM>
3491  : public LeafFnBase<CoordT, MaskT, LOG2DIM>
3492 {
3494  using BuildType = Fp8;
3495  using ArrayType = uint8_t;// type used for the internal mValue array
3496  static constexpr bool FIXED_SIZE = true;
3497  alignas(32) uint8_t mCode[1u << 3 * LOG2DIM];
3498 
3499  __hostdev__ static constexpr uint8_t bitWidth() { return 8u; }
3500  __hostdev__ float getValue(uint32_t i) const
3501  {
3502  return mCode[i]*BaseT::mQuantum + BaseT::mMinimum;// code * (max-min)/255 + min
3503  }
3504  /// @brief This class cannot be constructed or deleted
3505  LeafData() = delete;
3506  LeafData(const LeafData&) = delete;
3507  LeafData& operator=(const LeafData&) = delete;
3508  ~LeafData() = delete;
3509 }; // LeafData<Fp8>
3510 
3511 template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3512 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData<Fp16, CoordT, MaskT, LOG2DIM>
3513  : public LeafFnBase<CoordT, MaskT, LOG2DIM>
3514 {
3516  using BuildType = Fp16;
3517  using ArrayType = uint16_t;// type used for the internal mValue array
3518  static constexpr bool FIXED_SIZE = true;
3519  alignas(32) uint16_t mCode[1u << 3 * LOG2DIM];
3520 
3521  __hostdev__ static constexpr uint8_t bitWidth() { return 16u; }
3522  __hostdev__ float getValue(uint32_t i) const
3523  {
3524  return mCode[i]*BaseT::mQuantum + BaseT::mMinimum;// code * (max-min)/65535 + min
3525  }
3526 
3527  /// @brief This class cannot be constructed or deleted
3528  LeafData() = delete;
3529  LeafData(const LeafData&) = delete;
3530  LeafData& operator=(const LeafData&) = delete;
3531  ~LeafData() = delete;
3532 }; // LeafData<Fp16>
3533 
3534 template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3535 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData<FpN, CoordT, MaskT, LOG2DIM>
3536  : public LeafFnBase<CoordT, MaskT, LOG2DIM>
3537 {
3539  using BuildType = FpN;
3540  static constexpr bool FIXED_SIZE = false;
3541 
3542  __hostdev__ uint8_t bitWidth() const { return 1 << (BaseT::mFlags >> 5); }// 4,8,16,32 = 2^(2,3,4,5)
3543  __hostdev__ size_t memUsage() const { return sizeof(*this) + this->bitWidth()*64; }
3544  __hostdev__ static size_t memUsage(uint32_t bitWidth) { return 96u + bitWidth*64; }
3545  __hostdev__ float getValue(uint32_t i) const
3546  {
3547 #ifdef NANOVDB_FPN_BRANCHLESS// faster
3548  const int b = BaseT::mFlags >> 5;// b = 0, 1, 2, 3, 4 corresponding to 1, 2, 4, 8, 16 bits
3549 #if 0// use LUT
3550  uint16_t code = reinterpret_cast<const uint16_t*>(this + 1)[i >> (4 - b)];
3551  const static uint8_t shift[5] = {15, 7, 3, 1, 0};
3552  const static uint16_t mask[5] = {1, 3, 15, 255, 65535};
3553  code >>= (i & shift[b]) << b;
3554  code &= mask[b];
3555 #else// no LUT
3556  uint32_t code = reinterpret_cast<const uint32_t*>(this + 1)[i >> (5 - b)];
3557  //code >>= (i & ((16 >> b) - 1)) << b;
3558  code >>= (i & ((32 >> b) - 1)) << b;
3559  code &= (1 << (1 << b)) - 1;
3560 #endif
3561 #else// use branched version (slow)
3562  float code;
3563  auto *values = reinterpret_cast<const uint8_t*>(this+1);
3564  switch (BaseT::mFlags >> 5) {
3565  case 0u:// 1 bit float
3566  code = float((values[i>>3] >> (i&7) ) & uint8_t(1));
3567  break;
3568  case 1u:// 2 bits float
3569  code = float((values[i>>2] >> ((i&3)<<1)) & uint8_t(3));
3570  break;
3571  case 2u:// 4 bits float
3572  code = float((values[i>>1] >> ((i&1)<<2)) & uint8_t(15));
3573  break;
3574  case 3u:// 8 bits float
3575  code = float(values[i]);
3576  break;
3577  default:// 16 bits float
3578  code = float(reinterpret_cast<const uint16_t*>(values)[i]);
3579  }
3580 #endif
3581  return float(code) * BaseT::mQuantum + BaseT::mMinimum;// code * (max-min)/UNITS + min
3582  }
3583 
3584  /// @brief This class cannot be constructed or deleted
3585  LeafData() = delete;
3586  LeafData(const LeafData&) = delete;
3587  LeafData& operator=(const LeafData&) = delete;
3588  ~LeafData() = delete;
3589 }; // LeafData<FpN>
3590 
3591 // Partial template specialization of LeafData with bool
3592 template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3593 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData<bool, CoordT, MaskT, LOG2DIM>
3594 {
3595  static_assert(sizeof(CoordT) == sizeof(Coord), "Mismatching sizeof");
3596  static_assert(sizeof(MaskT<LOG2DIM>) == sizeof(Mask<LOG2DIM>), "Mismatching sizeof");
3597  using ValueType = bool;
3598  using BuildType = bool;
3599  using FloatType = bool;// dummy value type
3600  using ArrayType = MaskT<LOG2DIM>;// type used for the internal mValue array
3601  static constexpr bool FIXED_SIZE = true;
3602 
3603  CoordT mBBoxMin; // 12B.
3604  uint8_t mBBoxDif[3]; // 3B.
3605  uint8_t mFlags; // 1B.
3606  MaskT<LOG2DIM> mValueMask; // LOG2DIM(3): 64B.
3607  MaskT<LOG2DIM> mValues; // LOG2DIM(3): 64B.
3608 
3609  //__hostdev__ const ValueType* values() const { return nullptr; }
3610  __hostdev__ bool getValue(uint32_t i) const { return mValues.isOn(i); }
3611  __hostdev__ bool getMin() const { return false; }// dummy
3612  __hostdev__ bool getMax() const { return false; }// dummy
3613  __hostdev__ bool getAvg() const { return false; }// dummy
3614  __hostdev__ bool getDev() const { return false; }// dummy
3615  __hostdev__ void setValue(uint32_t offset, bool v)
3616  {
3617  mValueMask.setOn(offset);
3618  mValues.set(offset, v);
3619  }
3620 
3621  __hostdev__ void setMin(const bool&) {}// no-op
3622  __hostdev__ void setMax(const bool&) {}// no-op
3623  __hostdev__ void setAvg(const bool&) {}// no-op
3624  __hostdev__ void setDev(const bool&) {}// no-op
3625 
3626  template <typename T>
3627  __hostdev__ void setOrigin(const T& ijk) { mBBoxMin = ijk; }
3628 
3629  /// @brief This class cannot be constructed or deleted
3630  LeafData() = delete;
3631  LeafData(const LeafData&) = delete;
3632  LeafData& operator=(const LeafData&) = delete;
3633  ~LeafData() = delete;
3634 }; // LeafData<bool>
3635 
3636 // Partial template specialization of LeafData with ValueMask
3637 template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3638 struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData<ValueMask, CoordT, MaskT, LOG2DIM>
3639 {
3640  static_assert(sizeof(CoordT) == sizeof(Coord), "Mismatching sizeof");
3641  static_assert(sizeof(MaskT<LOG2DIM>) == sizeof(Mask<LOG2DIM>), "Mismatching sizeof");
3642  using ValueType = bool;
3644  using FloatType = bool;// dummy value type
3645  using ArrayType = void;// type used for the internal mValue array - void means missing
3646  static constexpr bool FIXED_SIZE = true;
3647 
3648  CoordT mBBoxMin; // 12B.
3649  uint8_t mBBoxDif[3]; // 3B.
3650  uint8_t mFlags; // 1B.
3651  MaskT<LOG2DIM> mValueMask; // LOG2DIM(3): 64B.
3652 
3653  //__hostdev__ const ValueType* values() const { return nullptr; }
3654  __hostdev__ bool getValue(uint32_t i) const { return mValueMask.isOn(i); }
3655  __hostdev__ bool getMin() const { return false; }// dummy
3656  __hostdev__ bool getMax() const { return false; }// dummy
3657  __hostdev__ bool getAvg() const { return false; }// dummy
3658  __hostdev__ bool getDev() const { return false; }// dummy
3659  __hostdev__ void setValue(uint32_t offset, bool)
3660  {
3661  mValueMask.setOn(offset);
3662  }
3663 
3664  __hostdev__ void setMin(const ValueType&) {}// no-op
3665  __hostdev__ void setMax(const ValueType&) {}// no-op
3666  __hostdev__ void setAvg(const FloatType&) {}// no-op
3667  __hostdev__ void setDev(const FloatType&) {}// no-op
3668 
3669  template <typename T>
3670  __hostdev__ void setOrigin(const T& ijk) { mBBoxMin = ijk; }
3671 
3672  /// @brief This class cannot be constructed or deleted
3673  LeafData() = delete;
3674  LeafData(const LeafData&) = delete;
3675  LeafData& operator=(const LeafData&) = delete;
3676  ~LeafData() = delete;
3677 }; // LeafData<ValueMask>
3678 
3679 /// @brief Leaf nodes of the VDB tree. (defaults to 8x8x8 = 512 voxels)
3680 template<typename BuildT,
3681  typename CoordT = Coord,
3682  template<uint32_t> class MaskT = Mask,
3683  uint32_t Log2Dim = 3>
3684 class LeafNode : private LeafData<BuildT, CoordT, MaskT, Log2Dim>
3685 {
3686 public:
3688  {
3689  __hostdev__ static uint32_t dim() { return 1u; }
3690  }; // Voxel
3693  using ValueType = typename DataType::ValueType;
3694  using FloatType = typename DataType::FloatType;
3695  using BuildType = typename DataType::BuildType;
3696  using CoordType = CoordT;
3697  static constexpr bool FIXED_SIZE = DataType::FIXED_SIZE;
3698  template<uint32_t LOG2>
3699  using MaskType = MaskT<LOG2>;
3700 
3701  static_assert(is_same<ValueType,typename BuildToValueMap<BuildType>::Type>::value, "Mismatching BuildType");
3702  static constexpr uint32_t LOG2DIM = Log2Dim;
3703  static constexpr uint32_t TOTAL = LOG2DIM; // needed by parent nodes
3704  static constexpr uint32_t DIM = 1u << TOTAL; // number of voxels along each axis of this node
3705  static constexpr uint32_t SIZE = 1u << 3 * LOG2DIM; // total number of voxels represented by this node
3706  static constexpr uint32_t MASK = (1u << LOG2DIM) - 1u; // mask for bit operations
3707  static constexpr uint32_t LEVEL = 0; // level 0 = leaf
3708  static constexpr uint64_t NUM_VALUES = uint64_t(1) << (3 * TOTAL); // total voxel count represented by this node
3709 
3710  __hostdev__ DataType* data() { return reinterpret_cast<DataType*>(this); }
3711 
3712  __hostdev__ const DataType* data() const { return reinterpret_cast<const DataType*>(this); }
3713 
3714  /// @brief Return a const reference to the bit mask of active voxels in this leaf node
3715  __hostdev__ const MaskType<LOG2DIM>& valueMask() const { return DataType::mValueMask; }
3716 
3717  /// @brief Return a const reference to the minimum active value encoded in this leaf node
3718  __hostdev__ ValueType minimum() const { return this->getMin(); }
3719 
3720  /// @brief Return a const reference to the maximum active value encoded in this leaf node
3721  __hostdev__ ValueType maximum() const { return this->getMax(); }
3722 
3723  /// @brief Return a const reference to the average of all the active values encoded in this leaf node
3724  __hostdev__ FloatType average() const { return DataType::getAvg(); }
3725 
3726  /// @brief Return the variance of all the active values encoded in this leaf node
3727  __hostdev__ FloatType variance() const { return DataType::getDev()*DataType::getDev(); }
3728 
3729  /// @brief Return a const reference to the standard deviation of all the active values encoded in this leaf node
3730  __hostdev__ FloatType stdDeviation() const { return DataType::getDev(); }
3731 
3732  __hostdev__ uint8_t flags() const { return DataType::mFlags; }
3733 
3734  /// @brief Return the origin in index space of this leaf node
3735  __hostdev__ CoordT origin() const { return DataType::mBBoxMin & ~MASK; }
3736 
3737  __hostdev__ static CoordT OffsetToLocalCoord(uint32_t n)
3738  {
3739  NANOVDB_ASSERT(n < SIZE);
3740  const uint32_t m = n & ((1 << 2 * LOG2DIM) - 1);
3741  return CoordT(n >> 2 * LOG2DIM, m >> LOG2DIM, m & MASK);
3742  }
3743 
3744  /// @brief Converts (in place) a local index coordinate to a global index coordinate
3745  __hostdev__ void localToGlobalCoord(Coord& ijk) const { ijk += this->origin(); }
3746 
3747  __hostdev__ CoordT offsetToGlobalCoord(uint32_t n) const
3748  {
3749  return OffsetToLocalCoord(n) + this->origin();
3750  }
3751 
3752  /// @brief Return the dimension, in index space, of this leaf node (typically 8 as for openvdb leaf nodes!)
3753  __hostdev__ static uint32_t dim() { return 1u << LOG2DIM; }
3754 
3755  /// @brief Return the bounding box in index space of active values in this leaf node
3757  {
3758  BBox<CoordT> bbox(DataType::mBBoxMin, DataType::mBBoxMin);
3759  if ( this->isActive() ) {
3760  bbox.max()[0] += DataType::mBBoxDif[0];
3761  bbox.max()[1] += DataType::mBBoxDif[1];
3762  bbox.max()[2] += DataType::mBBoxDif[2];
3763  } else {// very rare case
3764  bbox = BBox<CoordT>();// invalid
3765  }
3766  return bbox;
3767  }
3768 
3769  /// @brief Return the total number of voxels (e.g. values) encoded in this leaf node
3770  __hostdev__ static uint32_t voxelCount() { return 1u << (3 * LOG2DIM); }
3771 
3772  /// @brief return memory usage in bytes for the class
3773  __hostdev__ static uint64_t memUsage() { return sizeof(LeafNodeType); }
3774 
3775  /// @brief This class cannot be constructed or deleted
3776  LeafNode() = delete;
3777  LeafNode(const LeafNode&) = delete;
3778  LeafNode& operator=(const LeafNode&) = delete;
3779  ~LeafNode() = delete;
3780 
3781  /// @brief Return the voxel value at the given offset.
3782  __hostdev__ ValueType getValue(uint32_t offset) const { return DataType::getValue(offset); }
3783 
3784  /// @brief Return the voxel value at the given coordinate.
3785  __hostdev__ ValueType getValue(const CoordT& ijk) const { return DataType::getValue(CoordToOffset(ijk)); }
3786 
3787  /// @brief Sets the value at the specified location and activate its state.
3788  ///
3789  /// @note This is safe since it does not change the topology of the tree (unlike setValue methods on the other nodes)
3790  __hostdev__ void setValue(const CoordT& ijk, const ValueType& v) { DataType::setValue(CoordToOffset(ijk), v); }
3791 
3792  /// @brief Sets the value at the specified location but leaves its state unchanged.
3793  ///
3794  /// @note This is safe since it does not change the topology of the tree (unlike setValue methods on the other nodes)
3795  __hostdev__ void setValueOnly(uint32_t offset, const ValueType& v) { DataType::setValueOnly(offset, v); }
3796  __hostdev__ void setValueOnly(const CoordT& ijk, const ValueType& v) { DataType::setValueOnly(CoordToOffset(ijk), v); }
3797 
3798  /// @brief Return @c true if the voxel value at the given coordinate is active.
3799  __hostdev__ bool isActive(const CoordT& ijk) const { return DataType::mValueMask.isOn(CoordToOffset(ijk)); }
3800  __hostdev__ bool isActive(uint32_t n) const { return DataType::mValueMask.isOn(n); }
3801 
3802  /// @brief Return @c true if any of the voxel value are active in this leaf node.
3803  __hostdev__ bool isActive() const
3804  {
3805  NANOVDB_ASSERT( bool(DataType::mFlags & uint8_t(2)) != DataType::mValueMask.isOff() );
3806  return DataType::mFlags & uint8_t(2);
3807  }
3808 
3809 
3810  /// @brief Return @c true if the voxel value at the given coordinate is active and updates @c v with the value.
3811  __hostdev__ bool probeValue(const CoordT& ijk, ValueType& v) const
3812  {
3813  const uint32_t n = CoordToOffset(ijk);
3814  v = DataType::getValue(n);
3815  return DataType::mValueMask.isOn(n);
3816  }
3817 
3818  __hostdev__ const LeafNode* probeLeaf(const CoordT&) const { return this; }
3819 
3820  /// @brief Return the linear offset corresponding to the given coordinate
3821  __hostdev__ static uint32_t CoordToOffset(const CoordT& ijk)
3822  {
3823  #if 0
3824  return ((ijk[0] & MASK) << (2 * LOG2DIM)) + ((ijk[1] & MASK) << LOG2DIM) + (ijk[2] & MASK);
3825  #else
3826  return ((ijk[0] & MASK) << (2 * LOG2DIM)) | ((ijk[1] & MASK) << LOG2DIM) | (ijk[2] & MASK);
3827  #endif
3828  }
3829 
3830  /// @brief Updates the local bounding box of active voxels in this node.
3831  ///
3832  /// @warning It assumes that the origin and value mask have already been set.
3833  ///
3834  /// @details This method is based on few (intrinsic) bit operations and hence is relatively fast.
3835  /// However, it should only only be called of either the value mask has changed or if the
3836  /// active bounding box is still undefined. e.g. during constrution of this node.
3837  __hostdev__ void updateBBox();
3838 
3839 private:
3840  static_assert(sizeof(DataType) % NANOVDB_DATA_ALIGNMENT == 0, "sizeof(LeafData) is misaligned");
3841  //static_assert(offsetof(DataType, mValues) % 32 == 0, "LeafData::mValues is misaligned");
3842 
3843  template<typename, int, int, int>
3844  friend class ReadAccessor;
3845 
3846  template<typename>
3847  friend class RootNode;
3848  template<typename, uint32_t>
3849  friend class InternalNode;
3850 
3851  /// @brief Private method to return a voxel value and update a (dummy) ReadAccessor
3852  template<typename AccT>
3853  __hostdev__ ValueType getValueAndCache(const CoordT& ijk, const AccT&) const { return this->getValue(ijk); }
3854 
3855  /// @brief Return the node information.
3856  template<typename AccT>
3857  __hostdev__ typename AccT::NodeInfo getNodeInfoAndCache(const CoordType& /*ijk*/, const AccT& /*acc*/) const {
3858  using NodeInfoT = typename AccT::NodeInfo;
3859  return NodeInfoT{LEVEL, this->dim(), this->minimum(), this->maximum(),
3860  this->average(), this->stdDeviation(), this->bbox()[0], this->bbox()[1]};
3861  }
3862 
3863  template<typename AccT>
3864  __hostdev__ bool isActiveAndCache(const CoordT& ijk, const AccT&) const { return this->isActive(ijk); }
3865 
3866  template<typename AccT>
3867  __hostdev__ bool probeValueAndCache(const CoordT& ijk, ValueType& v, const AccT&) const { return this->probeValue(ijk, v); }
3868 
3869  template<typename AccT>
3870  __hostdev__ const LeafNode* probeLeafAndCache(const CoordT&, const AccT&) const { return this; }
3871 
3872  template<typename RayT, typename AccT>
3873  __hostdev__ uint32_t getDimAndCache(const CoordT&, const RayT& /*ray*/, const AccT&) const
3874  {
3875  if (DataType::mFlags & uint8_t(1))
3876  return this->dim(); // skip this node if first bit is set
3877  //if (!ray.intersects( this->bbox() )) return 1 << LOG2DIM;
3878  return ChildNodeType::dim();
3879  }
3880 
3881 }; // LeafNode class
3882 
3883 template<typename ValueT, typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3885 {
3886  static_assert(LOG2DIM == 3, "LeafNode::updateBBox: only supports LOGDIM = 3!");
3887  if (!this->isActive()) return;
3888  auto update = [&](uint32_t min, uint32_t max, int axis) {
3889  NANOVDB_ASSERT(min <= max && max < 8);
3890  DataType::mBBoxMin[axis] = (DataType::mBBoxMin[axis] & ~MASK) + int(min);
3891  DataType::mBBoxDif[axis] = uint8_t(max - min);
3892  };
3893  uint64_t word64 = DataType::mValueMask.template getWord<uint64_t>(0);
3894  uint32_t Xmin = word64 ? 0u : 8u;
3895  uint32_t Xmax = Xmin;
3896  for (int i = 1; i < 8; ++i) { // last loop over 8 64 words
3897  if (uint64_t w = DataType::mValueMask.template getWord<uint64_t>(i)) { // skip if word has no set bits
3898  word64 |= w; // union 8 x 64 bits words into one 64 bit word
3899  if (Xmin == 8) {
3900  Xmin = i; // only set once
3901  }
3902  Xmax = i;
3903  }
3904  }
3905  NANOVDB_ASSERT(word64);
3906  update(Xmin, Xmax, 0);
3907  update(FindLowestOn(word64) >> 3, FindHighestOn(word64) >> 3, 1);
3908  const uint32_t *p = reinterpret_cast<const uint32_t*>(&word64), word32 = p[0] | p[1];
3909  const uint16_t *q = reinterpret_cast<const uint16_t*>(&word32), word16 = q[0] | q[1];
3910  const uint8_t *b = reinterpret_cast<const uint8_t* >(&word16), byte = b[0] | b[1];
3911  NANOVDB_ASSERT(byte);
3912  update(FindLowestOn(static_cast<uint32_t>(byte)), FindHighestOn(static_cast<uint32_t>(byte)), 2);
3913 } // LeafNode::updateBBox
3914 
3915 // --------------------------> Template specializations and traits <------------------------------------
3916 
3917 /// @brief Template specializations to the default configuration used in OpenVDB:
3918 /// Root -> 32^3 -> 16^3 -> 8^3
3919 template<typename BuildT>
3921 template<typename BuildT>
3923 template<typename BuildT>
3925 template<typename BuildT>
3927 template<typename BuildT>
3929 template<typename BuildT>
3931 
3932 /// @brief Trait to map from LEVEL to node type
3933 template<typename BuildT, int LEVEL>
3934 struct NanoNode;
3935 
3936 // Partial template specialization of above Node struct
3937 template<typename BuildT>
3938 struct NanoNode<BuildT, 0>
3939 {
3942 };
3943 template<typename BuildT>
3944 struct NanoNode<BuildT, 1>
3945 {
3948 };
3949 template<typename BuildT>
3950 struct NanoNode<BuildT, 2>
3951 {
3954 };
3955 template<typename BuildT>
3956 struct NanoNode<BuildT, 3>
3957 {
3960 };
3961 
3974 
3987 
3988 // --------------------------> ReadAccessor <------------------------------------
3989 
3990 /// @brief A read-only value accessor with three levels of node caching. This allows for
3991 /// inverse tree traversal during lookup, which is on average significantly faster
3992 /// than calling the equivalent method on the tree (i.e. top-down traversal).
3993 ///
3994 /// @note By virtue of the fact that a value accessor accelerates random access operations
3995 /// by re-using cached access patterns, this access should be reused for multiple access
3996 /// operations. In other words, never create an instance of this accessor for a single
3997 /// acccess only. In general avoid single access operations with this accessor, and
3998 /// if that is not possible call the corresponding method on the tree instead.
3999 ///
4000 /// @warning Since this ReadAccessor internally caches raw pointers to the nodes of the tree
4001 /// structure, it is not safe to copy between host and device, or even to share among
4002 /// multiple threads on the same host or device. However, it is light-weight so simple
4003 /// instantiate one per thread (on the host and/or device).
4004 ///
4005 /// @details Used to accelerated random access into a VDB tree. Provides on average
4006 /// O(1) random access operations by means of inverse tree traversal,
4007 /// which amortizes the non-const time complexity of the root node.
4008 
4009 template <typename BuildT>
4010 class ReadAccessor<BuildT, -1, -1, -1>
4011 {
4012  using RootT = NanoRoot<BuildT>; // root node
4013  using LeafT = NanoLeaf<BuildT>; // Leaf node
4014  using FloatType = typename RootT::FloatType;
4015  using CoordValueType = typename RootT::CoordType::ValueType;
4016 
4017  mutable const RootT* mRoot; // 8 bytes (mutable to allow for access methods to be const)
4018 public:
4019  using ValueType = typename RootT::ValueType;
4020  using CoordType = typename RootT::CoordType;
4021 
4022  static const int CacheLevels = 0;
4023 
4024  struct NodeInfo {
4025  uint32_t mLevel; // 4B
4026  uint32_t mDim; // 4B
4027  ValueType mMinimum; // typically 4B
4028  ValueType mMaximum; // typically 4B
4029  FloatType mAverage; // typically 4B
4030  FloatType mStdDevi; // typically 4B
4033  };
4034 
4035  /// @brief Constructor from a root node
4036  __hostdev__ ReadAccessor(const RootT& root) : mRoot{&root} {}
4037 
4038  __hostdev__ const RootT& root() const { return *mRoot; }
4039 
4040  /// @brief Defaults constructors
4041  ReadAccessor(const ReadAccessor&) = default;
4042  ~ReadAccessor() = default;
4044 
4046  {
4047  return mRoot->getValueAndCache(ijk, *this);
4048  }
4049 
4050  __hostdev__ NodeInfo getNodeInfo(const CoordType& ijk) const
4051  {
4052  return mRoot->getNodeInfoAndCache(ijk, *this);
4053  }
4054 
4055  __hostdev__ bool isActive(const CoordType& ijk) const
4056  {
4057  return mRoot->isActiveAndCache(ijk, *this);
4058  }
4059 
4060  __hostdev__ bool probeValue(const CoordType& ijk, ValueType& v) const
4061  {
4062  return mRoot->probeValueAndCache(ijk, v, *this);
4063  }
4064 
4065  __hostdev__ const LeafT* probeLeaf(const CoordType& ijk) const
4066  {
4067  return mRoot->probeLeafAndCache(ijk, *this);
4068  }
4069 
4070  template<typename RayT>
4071  __hostdev__ uint32_t getDim(const CoordType& ijk, const RayT& ray) const
4072  {
4073  return mRoot->getDimAndCache(ijk, ray, *this);
4074  }
4075 
4076 private:
4077  /// @brief Allow nodes to insert themselves into the cache.
4078  template<typename>
4079  friend class RootNode;
4080  template<typename, uint32_t>
4081  friend class InternalNode;
4082  template<typename, typename, template<uint32_t> class, uint32_t>
4083  friend class LeafNode;
4084 
4085  /// @brief No-op
4086  template<typename NodeT>
4087  __hostdev__ void insert(const CoordType&, const NodeT*) const {}
4088 }; // ReadAccessor<ValueT, -1, -1, -1> class
4089 
4090 /// @brief Node caching at a single tree level
4091 template <typename BuildT, int LEVEL0>
4092 class ReadAccessor<BuildT, LEVEL0, -1, -1>//e.g. 0, 1, 2
4093 {
4094  static_assert(LEVEL0 >= 0 && LEVEL0 <= 2, "LEVEL0 should be 0, 1, or 2");
4095 
4096  using TreeT = NanoTree<BuildT>;
4097  using RootT = NanoRoot<BuildT>; // root node
4098  using LeafT = NanoLeaf<BuildT>; // Leaf node
4099  using NodeT = typename NodeTrait<TreeT, LEVEL0>::type;
4100  using CoordT = typename RootT::CoordType;
4101  using ValueT = typename RootT::ValueType;
4102 
4103  using FloatType = typename RootT::FloatType;
4104  using CoordValueType = typename RootT::CoordT::ValueType;
4105 
4106  // All member data are mutable to allow for access methods to be const
4107  mutable CoordT mKey; // 3*4 = 12 bytes
4108  mutable const RootT* mRoot; // 8 bytes
4109  mutable const NodeT* mNode; // 8 bytes
4110 
4111 public:
4112  using ValueType = ValueT;
4113  using CoordType = CoordT;
4114 
4115  static const int CacheLevels = 1;
4116 
4117  using NodeInfo = typename ReadAccessor<ValueT, -1, -1, -1>::NodeInfo;
4118 
4119  /// @brief Constructor from a root node
4121  : mKey(CoordType::max())
4122  , mRoot(&root)
4123  , mNode(nullptr)
4124  {
4125  }
4126 
4127  __hostdev__ const RootT& root() const { return *mRoot; }
4128 
4129  /// @brief Defaults constructors
4130  ReadAccessor(const ReadAccessor&) = default;
4131  ~ReadAccessor() = default;
4133 
4134  __hostdev__ bool isCached(const CoordType& ijk) const
4135  {
4136  return (ijk[0] & int32_t(~NodeT::MASK)) == mKey[0] &&
4137  (ijk[1] & int32_t(~NodeT::MASK)) == mKey[1] &&
4138  (ijk[2] & int32_t(~NodeT::MASK)) == mKey[2];
4139  }
4140 
4142  {
4143  if (this->isCached(ijk)) {
4144  return mNode->getValueAndCache(ijk, *this);
4145  }
4146  return mRoot->getValueAndCache(ijk, *this);
4147  }
4148 
4150  {
4151  if (this->isCached(ijk)) {
4152  return mNode->getNodeInfoAndCache(ijk, *this);
4153  }
4154  return mRoot->getNodeInfoAndCache(ijk, *this);
4155  }
4156 
4157  __hostdev__ bool isActive(const CoordType& ijk) const
4158  {
4159  if (this->isCached(ijk)) {
4160  return mNode->isActiveAndCache(ijk, *this);
4161  }
4162  return mRoot->isActiveAndCache(ijk, *this);
4163  }
4164 
4165  __hostdev__ bool probeValue(const CoordType& ijk, ValueType& v) const
4166  {
4167  if (this->isCached(ijk)) {
4168  return mNode->probeValueAndCache(ijk, v, *this);
4169  }
4170  return mRoot->probeValueAndCache(ijk, v, *this);
4171  }
4172 
4173  __hostdev__ const LeafT* probeLeaf(const CoordType& ijk) const
4174  {
4175  if (this->isCached(ijk)) {
4176  return mNode->probeLeafAndCache(ijk, *this);
4177  }
4178  return mRoot->probeLeafAndCache(ijk, *this);
4179  }
4180 
4181  template<typename RayT>
4182  __hostdev__ uint32_t getDim(const CoordType& ijk, const RayT& ray) const
4183  {
4184  if (this->isCached(ijk)) {
4185  return mNode->getDimAndCache(ijk, ray, *this);
4186  }
4187  return mRoot->getDimAndCache(ijk, ray, *this);
4188  }
4189 
4190 private:
4191  /// @brief Allow nodes to insert themselves into the cache.
4192  template<typename>
4193  friend class RootNode;
4194  template<typename, uint32_t>
4195  friend class InternalNode;
4196  template<typename, typename, template<uint32_t> class, uint32_t>
4197  friend class LeafNode;
4198 
4199  /// @brief Inserts a leaf node and key pair into this ReadAccessor
4200  __hostdev__ void insert(const CoordType& ijk, const NodeT* node) const
4201  {
4202  mKey = ijk & ~NodeT::MASK;
4203  mNode = node;
4204  }
4205 
4206  // no-op
4207  template<typename OtherNodeT>
4208  __hostdev__ void insert(const CoordType&, const OtherNodeT*) const {}
4209 
4210 }; // ReadAccessor<ValueT, LEVEL0>
4211 
4212 template <typename BuildT, int LEVEL0, int LEVEL1>
4213 class ReadAccessor<BuildT, LEVEL0, LEVEL1, -1>//e.g. (0,1), (1,2), (0,2)
4214 {
4215  static_assert(LEVEL0 >= 0 && LEVEL0 <= 2, "LEVEL0 must be 0, 1, 2");
4216  static_assert(LEVEL1 >= 0 && LEVEL1 <= 2, "LEVEL1 must be 0, 1, 2");
4217  static_assert(LEVEL0 < LEVEL1, "Level 0 must be lower than level 1");
4218  using TreeT = NanoTree<BuildT>;
4219  using RootT = NanoRoot<BuildT>;
4220  using LeafT = NanoLeaf<BuildT>;
4221  using Node1T = typename NodeTrait<TreeT, LEVEL0>::type;
4222  using Node2T = typename NodeTrait<TreeT, LEVEL1>::type;
4223  using CoordT = typename RootT::CoordType;
4224  using ValueT = typename RootT::ValueType;
4225  using FloatType = typename RootT::FloatType;
4226  using CoordValueType = typename RootT::CoordT::ValueType;
4227 
4228  // All member data are mutable to allow for access methods to be const
4229 #ifdef USE_SINGLE_ACCESSOR_KEY // 44 bytes total
4230  mutable CoordT mKey; // 3*4 = 12 bytes
4231 #else // 68 bytes total
4232  mutable CoordT mKeys[2]; // 2*3*4 = 24 bytes
4233 #endif
4234  mutable const RootT* mRoot;
4235  mutable const Node1T* mNode1;
4236  mutable const Node2T* mNode2;
4237 
4238 public:
4239  using ValueType = ValueT;
4240  using CoordType = CoordT;
4241 
4242  static const int CacheLevels = 2;
4243 
4244  using NodeInfo = typename ReadAccessor<ValueT,-1,-1,-1>::NodeInfo;
4245 
4246  /// @brief Constructor from a root node
4248 #ifdef USE_SINGLE_ACCESSOR_KEY
4249  : mKey(CoordType::max())
4250 #else
4251  : mKeys{CoordType::max(), CoordType::max()}
4252 #endif
4253  , mRoot(&root)
4254  , mNode1(nullptr)
4255  , mNode2(nullptr)
4256  {
4257  }
4258 
4259  __hostdev__ const RootT& root() const { return *mRoot; }
4260 
4261  /// @brief Defaults constructors
4262  ReadAccessor(const ReadAccessor&) = default;
4263  ~ReadAccessor() = default;
4265 
4266 #ifdef USE_SINGLE_ACCESSOR_KEY
4267  __hostdev__ bool isCached1(CoordValueType dirty) const
4268  {
4269  if (!mNode1)
4270  return false;
4271  if (dirty & int32_t(~Node1T::MASK)) {
4272  mNode1 = nullptr;
4273  return false;
4274  }
4275  return true;
4276  }
4277  __hostdev__ bool isCached2(CoordValueType dirty) const
4278  {
4279  if (!mNode2)
4280  return false;
4281  if (dirty & int32_t(~Node2T::MASK)) {
4282  mNode2 = nullptr;
4283  return false;
4284  }
4285  return true;
4286  }
4287  __hostdev__ CoordValueType computeDirty(const CoordType& ijk) const
4288  {
4289  return (ijk[0] ^ mKey[0]) | (ijk[1] ^ mKey[1]) | (ijk[2] ^ mKey[2]);
4290  }
4291 #else
4292  __hostdev__ bool isCached1(const CoordType& ijk) const
4293  {
4294  return (ijk[0] & int32_t(~Node1T::MASK)) == mKeys[0][0] &&
4295  (ijk[1] & int32_t(~Node1T::MASK)) == mKeys[0][1] &&
4296  (ijk[2] & int32_t(~Node1T::MASK)) == mKeys[0][2];
4297  }
4298  __hostdev__ bool isCached2(const CoordType& ijk) const
4299  {
4300  return (ijk[0] & int32_t(~Node2T::MASK)) == mKeys[1][0] &&
4301  (ijk[1] & int32_t(~Node2T::MASK)) == mKeys[1][1] &&
4302  (ijk[2] & int32_t(~Node2T::MASK)) == mKeys[1][2];
4303  }
4304 #endif
4305 
4307  {
4308 #ifdef USE_SINGLE_ACCESSOR_KEY
4309  const CoordValueType dirty = this->computeDirty(ijk);
4310 #else
4311  auto&& dirty = ijk;
4312 #endif
4313  if (this->isCached1(dirty)) {
4314  return mNode1->getValueAndCache(ijk, *this);
4315  } else if (this->isCached2(dirty)) {
4316  return mNode2->getValueAndCache(ijk, *this);
4317  }
4318  return mRoot->getValueAndCache(ijk, *this);
4319  }
4320 
4322  {
4323 #ifdef USE_SINGLE_ACCESSOR_KEY
4324  const CoordValueType dirty = this->computeDirty(ijk);
4325 #else
4326  auto&& dirty = ijk;
4327 #endif
4328  if (this->isCached1(dirty)) {
4329  return mNode1->getNodeInfoAndCache(ijk, *this);
4330  } else if (this->isCached2(dirty)) {
4331  return mNode2->getNodeInfoAndCache(ijk, *this);
4332  }
4333  return mRoot->getNodeInfoAndCache(ijk, *this);
4334  }
4335 
4336  __hostdev__ bool isActive(const CoordType& ijk) const
4337  {
4338 #ifdef USE_SINGLE_ACCESSOR_KEY
4339  const CoordValueType dirty = this->computeDirty(ijk);
4340 #else
4341  auto&& dirty = ijk;
4342 #endif
4343  if (this->isCached1(dirty)) {
4344  return mNode1->isActiveAndCache(ijk, *this);
4345  } else if (this->isCached2(dirty)) {
4346  return mNode2->isActiveAndCache(ijk, *this);
4347  }
4348  return mRoot->isActiveAndCache(ijk, *this);
4349  }
4350 
4351  __hostdev__ bool probeValue(const CoordType& ijk, ValueType& v) const
4352  {
4353 #ifdef USE_SINGLE_ACCESSOR_KEY
4354  const CoordValueType dirty = this->computeDirty(ijk);
4355 #else
4356  auto&& dirty = ijk;
4357 #endif
4358  if (this->isCached1(dirty)) {
4359  return mNode1->probeValueAndCache(ijk, v, *this);
4360  } else if (this->isCached2(dirty)) {
4361  return mNode2->probeValueAndCache(ijk, v, *this);
4362  }
4363  return mRoot->probeValueAndCache(ijk, v, *this);
4364  }
4365 
4366  __hostdev__ const LeafT* probeLeaf(const CoordType& ijk) const
4367  {
4368 #ifdef USE_SINGLE_ACCESSOR_KEY
4369  const CoordValueType dirty = this->computeDirty(ijk);
4370 #else
4371  auto&& dirty = ijk;
4372 #endif
4373  if (this->isCached1(dirty)) {
4374  return mNode1->probeLeafAndCache(ijk, *this);
4375  } else if (this->isCached2(dirty)) {
4376  return mNode2->probeLeafAndCache(ijk, *this);
4377  }
4378  return mRoot->probeLeafAndCache(ijk, *this);
4379  }
4380 
4381  template<typename RayT>
4382  __hostdev__ uint32_t getDim(const CoordType& ijk, const RayT& ray) const
4383  {
4384 #ifdef USE_SINGLE_ACCESSOR_KEY
4385  const CoordValueType dirty = this->computeDirty(ijk);
4386 #else
4387  auto&& dirty = ijk;
4388 #endif
4389  if (this->isCached1(dirty)) {
4390  return mNode1->getDimAndCache(ijk, ray, *this);
4391  } else if (this->isCached2(dirty)) {
4392  return mNode2->getDimAndCache(ijk, ray, *this);
4393  }
4394  return mRoot->getDimAndCache(ijk, ray, *this);
4395  }
4396 
4397 private:
4398  /// @brief Allow nodes to insert themselves into the cache.
4399  template<typename>
4400  friend class RootNode;
4401  template<typename, uint32_t>
4402  friend class InternalNode;
4403  template<typename, typename, template<uint32_t> class, uint32_t>
4404  friend class LeafNode;
4405 
4406  /// @brief Inserts a leaf node and key pair into this ReadAccessor
4407  __hostdev__ void insert(const CoordType& ijk, const Node1T* node) const
4408  {
4409 #ifdef USE_SINGLE_ACCESSOR_KEY
4410  mKey = ijk;
4411 #else
4412  mKeys[0] = ijk & ~Node1T::MASK;
4413 #endif
4414  mNode1 = node;
4415  }
4416  __hostdev__ void insert(const CoordType& ijk, const Node2T* node) const
4417  {
4418 #ifdef USE_SINGLE_ACCESSOR_KEY
4419  mKey = ijk;
4420 #else
4421  mKeys[1] = ijk & ~Node2T::MASK;
4422 #endif
4423  mNode2 = node;
4424  }
4425  template <typename OtherNodeT>
4426  __hostdev__ void insert(const CoordType&, const OtherNodeT*) const {}
4427 }; // ReadAccessor<BuildT, LEVEL0, LEVEL1>
4428 
4429 
4430 /// @brief Node caching at all (three) tree levels
4431 template <typename BuildT>
4432 class ReadAccessor<BuildT, 0, 1, 2>
4433 {
4434  using TreeT = NanoTree<BuildT>;
4435  using RootT = NanoRoot<BuildT>; // root node
4436  using NodeT2 = NanoUpper<BuildT>; // upper internal node
4437  using NodeT1 = NanoLower<BuildT>; // lower internal node
4438  using LeafT = NanoLeaf< BuildT>; // Leaf node
4439  using CoordT = typename RootT::CoordType;
4440  using ValueT = typename RootT::ValueType;
4441 
4442  using FloatType = typename RootT::FloatType;
4443  using CoordValueType = typename RootT::CoordT::ValueType;
4444 
4445  // All member data are mutable to allow for access methods to be const
4446 #ifdef USE_SINGLE_ACCESSOR_KEY // 44 bytes total
4447  mutable CoordT mKey; // 3*4 = 12 bytes
4448 #else // 68 bytes total
4449  mutable CoordT mKeys[3]; // 3*3*4 = 36 bytes
4450 #endif
4451  mutable const RootT* mRoot;
4452  mutable const void* mNode[3]; // 4*8 = 32 bytes
4453 
4454 public:
4455  using ValueType = ValueT;
4456  using CoordType = CoordT;
4457 
4458  static const int CacheLevels = 3;
4459 
4460  using NodeInfo = typename ReadAccessor<ValueT, -1, -1, -1>::NodeInfo;
4461 
4462  /// @brief Constructor from a root node
4464 #ifdef USE_SINGLE_ACCESSOR_KEY
4465  : mKey(CoordType::max())
4466 #else
4468 #endif
4469  , mRoot(&root)
4470  , mNode{nullptr, nullptr, nullptr}
4471  {
4472  }
4473 
4474  __hostdev__ const RootT& root() const { return *mRoot; }
4475 
4476  /// @brief Defaults constructors
4477  ReadAccessor(const ReadAccessor&) = default;
4478  ~ReadAccessor() = default;
4480 
4481  /// @brief Return a const point to the cached node of the specified type
4482  ///
4483  /// @warning The return value could be NULL.
4484  template<typename NodeT>
4485  __hostdev__ const NodeT* getNode() const
4486  {
4487  using T = typename NodeTrait<TreeT, NodeT::LEVEL>::type;
4488  static_assert(is_same<T, NodeT>::value, "ReadAccessor::getNode: Invalid node type");
4489  return reinterpret_cast<const T*>(mNode[NodeT::LEVEL]);
4490  }
4491 
4492 #ifdef USE_SINGLE_ACCESSOR_KEY
4493  template<typename NodeT>
4494  __hostdev__ bool isCached(CoordValueType dirty) const
4495  {
4496  if (!mNode[NodeT::LEVEL])
4497  return false;
4498  if (dirty & int32_t(~NodeT::MASK)) {
4499  mNode[NodeT::LEVEL] = nullptr;
4500  return false;
4501  }
4502  return true;
4503  }
4504 
4505  __hostdev__ CoordValueType computeDirty(const CoordType& ijk) const
4506  {
4507  return (ijk[0] ^ mKey[0]) | (ijk[1] ^ mKey[1]) | (ijk[2] ^ mKey[2]);
4508  }
4509 #else
4510  template<typename NodeT>
4511  __hostdev__ bool isCached(const CoordType& ijk) const
4512  {
4513  return (ijk[0] & int32_t(~NodeT::MASK)) == mKeys[NodeT::LEVEL][0] && (ijk[1] & int32_t(~NodeT::MASK)) == mKeys[NodeT::LEVEL][1] && (ijk[2] & int32_t(~NodeT::MASK)) == mKeys[NodeT::LEVEL][2];
4514  }
4515 #endif
4516 
4518  {
4519 #ifdef USE_SINGLE_ACCESSOR_KEY
4520  const CoordValueType dirty = this->computeDirty(ijk);
4521 #else
4522  auto&& dirty = ijk;
4523 #endif
4524  if (this->isCached<LeafT>(dirty)) {
4525  return ((LeafT*)mNode[0])->getValue(ijk);
4526  } else if (this->isCached<NodeT1>(dirty)) {
4527  return ((NodeT1*)mNode[1])->getValueAndCache(ijk, *this);
4528  } else if (this->isCached<NodeT2>(dirty)) {
4529  return ((NodeT2*)mNode[2])->getValueAndCache(ijk, *this);
4530  }
4531  return mRoot->getValueAndCache(ijk, *this);
4532  }
4533 
4535  {
4536 #ifdef USE_SINGLE_ACCESSOR_KEY
4537  const CoordValueType dirty = this->computeDirty(ijk);
4538 #else
4539  auto&& dirty = ijk;
4540 #endif
4541  if (this->isCached<LeafT>(dirty)) {
4542  return ((LeafT*)mNode[0])->getNodeInfoAndCache(ijk, *this);
4543  } else if (this->isCached<NodeT1>(dirty)) {
4544  return ((NodeT1*)mNode[1])->getNodeInfoAndCache(ijk, *this);
4545  } else if (this->isCached<NodeT2>(dirty)) {
4546  return ((NodeT2*)mNode[2])->getNodeInfoAndCache(ijk, *this);
4547  }
4548  return mRoot->getNodeInfoAndCache(ijk, *this);
4549  }
4550 
4551  __hostdev__ bool isActive(const CoordType& ijk) const
4552  {
4553 #ifdef USE_SINGLE_ACCESSOR_KEY
4554  const CoordValueType dirty = this->computeDirty(ijk);
4555 #else
4556  auto&& dirty = ijk;
4557 #endif
4558  if (this->isCached<LeafT>(dirty)) {
4559  return ((LeafT*)mNode[0])->isActive(ijk);
4560  } else if (this->isCached<NodeT1>(dirty)) {
4561  return ((NodeT1*)mNode[1])->isActiveAndCache(ijk, *this);
4562  } else if (this->isCached<NodeT2>(dirty)) {
4563  return ((NodeT2*)mNode[2])->isActiveAndCache(ijk, *this);
4564  }
4565  return mRoot->isActiveAndCache(ijk, *this);
4566  }
4567 
4568  __hostdev__ bool probeValue(const CoordType& ijk, ValueType& v) const
4569  {
4570 #ifdef USE_SINGLE_ACCESSOR_KEY
4571  const CoordValueType dirty = this->computeDirty(ijk);
4572 #else
4573  auto&& dirty = ijk;
4574 #endif
4575  if (this->isCached<LeafT>(dirty)) {
4576  return ((LeafT*)mNode[0])->probeValue(ijk, v);
4577  } else if (this->isCached<NodeT1>(dirty)) {
4578  return ((NodeT1*)mNode[1])->probeValueAndCache(ijk, v, *this);
4579  } else if (this->isCached<NodeT2>(dirty)) {
4580  return ((NodeT2*)mNode[2])->probeValueAndCache(ijk, v, *this);
4581  }
4582  return mRoot->probeValueAndCache(ijk, v, *this);
4583  }
4584 
4585  __hostdev__ const LeafT* probeLeaf(const CoordType& ijk) const
4586  {
4587 #ifdef USE_SINGLE_ACCESSOR_KEY
4588  const CoordValueType dirty = this->computeDirty(ijk);
4589 #else
4590  auto&& dirty = ijk;
4591 #endif
4592  if (this->isCached<LeafT>(dirty)) {
4593  return ((LeafT*)mNode[0]);
4594  } else if (this->isCached<NodeT1>(dirty)) {
4595  return ((NodeT1*)mNode[1])->probeLeafAndCache(ijk, *this);
4596  } else if (this->isCached<NodeT2>(dirty)) {
4597  return ((NodeT2*)mNode[2])->probeLeafAndCache(ijk, *this);
4598  }
4599  return mRoot->probeLeafAndCache(ijk, *this);
4600  }
4601 
4602  template<typename RayT>
4603  __hostdev__ uint32_t getDim(const CoordType& ijk, const RayT& ray) const
4604  {
4605 #ifdef USE_SINGLE_ACCESSOR_KEY
4606  const CoordValueType dirty = this->computeDirty(ijk);
4607 #else
4608  auto&& dirty = ijk;
4609 #endif
4610  if (this->isCached<LeafT>(dirty)) {
4611  return ((LeafT*)mNode[0])->getDimAndCache(ijk, ray, *this);
4612  } else if (this->isCached<NodeT1>(dirty)) {
4613  return ((NodeT1*)mNode[1])->getDimAndCache(ijk, ray, *this);
4614  } else if (this->isCached<NodeT2>(dirty)) {
4615  return ((NodeT2*)mNode[2])->getDimAndCache(ijk, ray, *this);
4616  }
4617  return mRoot->getDimAndCache(ijk, ray, *this);
4618  }
4619 
4620 private:
4621  /// @brief Allow nodes to insert themselves into the cache.
4622  template<typename>
4623  friend class RootNode;
4624  template<typename, uint32_t>
4625  friend class InternalNode;
4626  template<typename, typename, template<uint32_t> class, uint32_t>
4627  friend class LeafNode;
4628 
4629  /// @brief Inserts a leaf node and key pair into this ReadAccessor
4630  template<typename NodeT>
4631  __hostdev__ void insert(const CoordType& ijk, const NodeT* node) const
4632  {
4633 #ifdef USE_SINGLE_ACCESSOR_KEY
4634  mKey = ijk;
4635 #else
4636  mKeys[NodeT::LEVEL] = ijk & ~NodeT::MASK;
4637 #endif
4638  mNode[NodeT::LEVEL] = node;
4639  }
4640 }; // ReadAccessor<BuildT, 0, 1, 2>
4641 
4642 //////////////////////////////////////////////////
4643 
4644 /// @brief Free-standing function for convenient creation of a ReadAccessor with
4645 /// optional and customizable node caching.
4646 ///
4647 /// @details createAccessor<>(grid): No caching of nodes and hence it's thread-safe but slow
4648 /// createAccessor<0>(grid): Caching of leaf nodes only
4649 /// createAccessor<1>(grid): Caching of lower internal nodes only
4650 /// createAccessor<2>(grid): Caching of upper internal nodes only
4651 /// createAccessor<0,1>(grid): Caching of leaf and lower internal nodes
4652 /// createAccessor<0,2>(grid): Caching of leaf and upper internal nodes
4653 /// createAccessor<1,2>(grid): Caching of lower and upper internal nodes
4654 /// createAccessor<0,1,2>(grid): Caching of all nodes at all tree levels
4655 
4656 template <int LEVEL0 = -1, int LEVEL1 = -1, int LEVEL2 = -1, typename ValueT = float>
4658 {
4659  return ReadAccessor<ValueT, LEVEL0, LEVEL1, LEVEL2>(grid.tree().root());
4660 }
4661 
4662 template <int LEVEL0 = -1, int LEVEL1 = -1, int LEVEL2 = -1, typename ValueT = float>
4664 {
4665  return ReadAccessor<ValueT, LEVEL0, LEVEL1, LEVEL2>(tree().root());
4666 }
4667 
4668 template <int LEVEL0 = -1, int LEVEL1 = -1, int LEVEL2 = -1, typename ValueT = float>
4670 {
4672 }
4673 
4674 //////////////////////////////////////////////////
4675 
4676 /// @brief This is a convenient class that allows for access to grid meta-data
4677 /// that are independent of the value type of a grid. That is, this class
4678 /// can be used to get information about a grid without actually knowing
4679 /// its ValueType.
4681 {
4682  // We cast to a grid templated on a dummy ValueType which is safe because we are very
4683  // careful only to call certain methods which are known to be invariant to the ValueType!
4684  // In other words, don't use this technique unless you are intimately familiar with the
4685  // memory-layout of the data structure and the reasons why certain methods are safe
4686  // to call and others are not!
4687  using GridT = NanoGrid<int>;
4688  __hostdev__ const GridT& grid() const { return *reinterpret_cast<const GridT*>(this); }
4689 
4690 public:
4691  __hostdev__ bool isValid() const { return this->grid().isValid(); }
4692  __hostdev__ uint64_t gridSize() const { return this->grid().gridSize(); }
4693  __hostdev__ uint32_t gridIndex() const { return this->grid().gridIndex(); }
4694  __hostdev__ uint32_t gridCount() const { return this->grid().gridCount(); }
4695  __hostdev__ const char* shortGridName() const { return this->grid().shortGridName(); }
4696  __hostdev__ GridType gridType() const { return this->grid().gridType(); }
4697  __hostdev__ GridClass gridClass() const { return this->grid().gridClass(); }
4698  __hostdev__ bool isLevelSet() const { return this->grid().isLevelSet(); }
4699  __hostdev__ bool isFogVolume() const { return this->grid().isFogVolume(); }
4700  __hostdev__ bool isPointIndex() const { return this->grid().isPointIndex(); }
4701  __hostdev__ bool isPointData() const { return this->grid().isPointData(); }
4702  __hostdev__ bool isMask() const { return this->grid().isMask(); }
4703  __hostdev__ bool isStaggered() const { return this->grid().isStaggered(); }
4704  __hostdev__ bool isUnknown() const { return this->grid().isUnknown(); }
4705  __hostdev__ const Map& map() const { return this->grid().map(); }
4706  __hostdev__ const BBox<Vec3R>& worldBBox() const { return this->grid().worldBBox(); }
4707  __hostdev__ const BBox<Coord>& indexBBox() const { return this->grid().indexBBox(); }
4708  __hostdev__ Vec3R voxelSize() const { return this->grid().voxelSize(); }
4709  __hostdev__ int blindDataCount() const { return this->grid().blindDataCount(); }
4710  __hostdev__ const GridBlindMetaData& blindMetaData(int n) const { return this->grid().blindMetaData(n); }
4711  __hostdev__ uint64_t activeVoxelCount() const { return this->grid().activeVoxelCount(); }
4712  __hostdev__ uint32_t activeTileCount(uint32_t n) const { return this->grid().tree().activeTileCount(n); }
4713  __hostdev__ uint32_t nodeCount(uint32_t level) const { return this->grid().tree().nodeCount(level); }
4714  __hostdev__ uint64_t checksum() const { return this->grid().checksum(); }
4715  __hostdev__ bool isEmpty() const { return this->grid().isEmpty(); }
4716  __hostdev__ Version version() const { return this->grid().version(); }
4717 }; // GridMetaData
4718 
4719 /// @brief Class to access points at a specific voxel location
4720 template<typename AttT>
4721 class PointAccessor : public DefaultReadAccessor<uint32_t>
4722 {
4724  const UInt32Grid* mGrid;
4725  const AttT* mData;
4726 
4727 public:
4729 
4731  : AccT(grid.tree().root())
4732  , mGrid(&grid)
4733  , mData(reinterpret_cast<const AttT*>(grid.blindData(0)))
4734  {
4738  NANOVDB_ASSERT(grid.blindDataCount() >= 1);
4739  }
4740  /// @brief Return the total number of point in the grid and set the
4741  /// iterators to the complete range of points.
4742  __hostdev__ uint64_t gridPoints(const AttT*& begin, const AttT*& end) const
4743  {
4744  const uint64_t count = mGrid->blindMetaData(0).mElementCount;
4745  begin = mData;
4746  end = begin + count;
4747  return count;
4748  }
4749  /// @brief Return the number of points in the leaf node containing the coordinate @a ijk.
4750  /// If this return value is larger than zero then the iterators @a begin and @a end
4751  /// will point to all the attributes contained within that leaf node.
4752  __hostdev__ uint64_t leafPoints(const Coord& ijk, const AttT*& begin, const AttT*& end) const
4753  {
4754  auto* leaf = this->probeLeaf(ijk);
4755  if (leaf == nullptr) {
4756  return 0;
4757  }
4758  begin = mData + leaf->minimum();
4759  end = begin + leaf->maximum();
4760  return leaf->maximum();
4761  }
4762 
4763  /// @brief get iterators over offsets to points at a specific voxel location
4764  __hostdev__ uint64_t voxelPoints(const Coord& ijk, const AttT*& begin, const AttT*& end) const
4765  {
4766  auto* leaf = this->probeLeaf(ijk);
4767  if (leaf == nullptr)
4768  return 0;
4769  const uint32_t offset = LeafNodeType::CoordToOffset(ijk);
4770  if (leaf->isActive(offset)) {
4771  auto* p = mData + leaf->minimum();
4772  begin = p + (offset == 0 ? 0 : leaf->getValue(offset - 1));
4773  end = p + leaf->getValue(offset);
4774  return end - begin;
4775  }
4776  return 0;
4777  }
4778 }; // PointAccessor
4779 
4780 } // namespace nanovdb
4781 
4782 #endif // end of NANOVDB_NANOVDB_H_HAS_BEEN_INCLUDED
#define ROOT_LEVEL
Definition: CNanoVDB.h:53
ValueT value
Definition: GridBuilder.h:1287
ChildT * child
Definition: GridBuilder.h:1286
#define NANOVDB_HOSTDEV_DISABLE_WARNING
Definition: NanoVDB.h:178
#define NANOVDB_MINOR_VERSION_NUMBER
Definition: NanoVDB.h:105
#define NANOVDB_DATA_ALIGNMENT
Definition: NanoVDB.h:116
#define __hostdev__
Definition: NanoVDB.h:168
#define NANOVDB_MAJOR_VERSION_NUMBER
Definition: NanoVDB.h:104
#define NANOVDB_ASSERT(x)
Definition: NanoVDB.h:149
#define NANOVDB_MAGIC_NUMBER
Definition: NanoVDB.h:102
#define NANOVDB_PATCH_VERSION_NUMBER
Definition: NanoVDB.h:106
Iterator(const BBox &b)
Definition: NanoVDB.h:1581
Iterator operator++(int)
Definition: NanoVDB.h:1600
Iterator & operator++()
Definition: NanoVDB.h:1586
const CoordT & operator*() const
Definition: NanoVDB.h:1608
Signed (i, j, k) 32-bit integer coordinate class, similar to openvdb::math::Coord.
Definition: NanoVDB.h:860
Coord operator-(const Coord &rhs) const
Definition: NanoVDB.h:962
Vec3< double > asVec3d() const
Return a double precision floating-point vector of this coordinate.
Definition: NanoVDB.h:1182
uint32_t IndexType
Definition: NanoVDB.h:864
Coord operator<<(IndexType n) const
Definition: NanoVDB.h:926
static Coord Floor(const Vec3T &xyz)
Return the largest integer coordinates that are not greater than xyz (node centered conversion).
Definition: NanoVDB.h:1019
Coord(ValueType n)
Initializes all coordinates to the given signed integer.
Definition: NanoVDB.h:873
Coord operator+(const Coord &rhs) const
Definition: NanoVDB.h:961
int32_t & y()
Definition: NanoVDB.h:894
Coord & operator+=(int n)
Definition: NanoVDB.h:954
ValueType & operator[](IndexType i)
Return a non-const reference to the given Coord component.
Definition: NanoVDB.h:909
static size_t memUsage()
Definition: NanoVDB.h:901
bool operator<(const Coord &rhs) const
Return true if this Coord is lexicographically less than the given Coord.
Definition: NanoVDB.h:932
int32_t & z()
Definition: NanoVDB.h:895
int32_t y() const
Definition: NanoVDB.h:890
Coord operator&(IndexType n) const
Return a new instance with coordinates masked by the given unsigned integer.
Definition: NanoVDB.h:923
Coord()
Initialize all coordinates to zero.
Definition: NanoVDB.h:867
uint8_t octant() const
Return the octant of this Coord.
Definition: NanoVDB.h:1028
int32_t z() const
Definition: NanoVDB.h:891
int32_t x() const
Definition: NanoVDB.h:889
Coord & operator=(const CoordT &other)
Assignment operator that works with openvdb::Coord.
Definition: NanoVDB.h:913
Coord & operator<<=(uint32_t n)
Definition: NanoVDB.h:947
Coord(ValueType i, ValueType j, ValueType k)
Initializes coordinate to the given signed integers.
Definition: NanoVDB.h:879
const ValueType & operator[](IndexType i) const
Return a const reference to the given Coord component.
Definition: NanoVDB.h:905
Coord offsetBy(ValueType n) const
Definition: NanoVDB.h:1007
Coord(ValueType *ptr)
Definition: NanoVDB.h:884
bool operator!=(const Coord &rhs) const
Definition: NanoVDB.h:939
static Coord min()
Definition: NanoVDB.h:899
Coord & operator&=(int n)
Definition: NanoVDB.h:940
Coord & operator-=(const Coord &rhs)
Definition: NanoVDB.h:970
bool operator==(const Coord &rhs) const
Definition: NanoVDB.h:938
Vec3< float > asVec3s() const
Return a single precision floating-point vector of this coordinate.
Definition: NanoVDB.h:1179
int32_t & x()
Definition: NanoVDB.h:893
Coord offsetBy(ValueType dx, ValueType dy, ValueType dz) const
Definition: NanoVDB.h:1002
Coord & maxComponent(const Coord &other)
Perform a component-wise maximum with the other Coord.
Definition: NanoVDB.h:991
static Coord max()
Definition: NanoVDB.h:897
static bool lessThan(const Coord &a, const Coord &b)
Definition: NanoVDB.h:1011
Coord operator>>(IndexType n) const
Definition: NanoVDB.h:929
uint32_t hash() const
Return a hash key derived from the existing coordinates.
Definition: NanoVDB.h:1024
Coord & operator+=(const Coord &rhs)
Definition: NanoVDB.h:963
int32_t ValueType
Definition: NanoVDB.h:863
Coord & minComponent(const Coord &other)
Perform a component-wise minimum with the other Coord.
Definition: NanoVDB.h:979
Dummy type for a 16bit quantization of float point values.
Definition: NanoVDB.h:201
Dummy type for a 4bit quantization of float point values.
Definition: NanoVDB.h:195
Dummy type for a 8bit quantization of float point values.
Definition: NanoVDB.h:198
Dummy type for a variable bit quantization of floating point values.
Definition: NanoVDB.h:204
This is a convenient class that allows for access to grid meta-data that are independent of the value...
Definition: NanoVDB.h:4681
bool isPointData() const
Definition: NanoVDB.h:4701
uint32_t gridIndex() const
Definition: NanoVDB.h:4693
uint64_t activeVoxelCount() const
Definition: NanoVDB.h:4711
const char * shortGridName() const
Definition: NanoVDB.h:4695
const Map & map() const
Definition: NanoVDB.h:4705
bool isStaggered() const
Definition: NanoVDB.h:4703
bool isValid() const
Definition: NanoVDB.h:4691
GridType gridType() const
Definition: NanoVDB.h:4696
const BBox< Vec3R > & worldBBox() const
Definition: NanoVDB.h:4706
uint64_t gridSize() const
Definition: NanoVDB.h:4692
int blindDataCount() const
Definition: NanoVDB.h:4709
uint64_t checksum() const
Definition: NanoVDB.h:4714
bool isFogVolume() const
Definition: NanoVDB.h:4699
uint32_t activeTileCount(uint32_t n) const
Definition: NanoVDB.h:4712
const GridBlindMetaData & blindMetaData(int n) const
Definition: NanoVDB.h:4710
bool isMask() const
Definition: NanoVDB.h:4702
uint32_t gridCount() const
Definition: NanoVDB.h:4694
GridClass gridClass() const
Definition: NanoVDB.h:4697
Version version() const
Definition: NanoVDB.h:4716
bool isEmpty() const
Definition: NanoVDB.h:4715
const BBox< Coord > & indexBBox() const
Definition: NanoVDB.h:4707
Vec3R voxelSize() const
Definition: NanoVDB.h:4708
bool isLevelSet() const
Definition: NanoVDB.h:4698
uint32_t nodeCount(uint32_t level) const
Definition: NanoVDB.h:4713
bool isPointIndex() const
Definition: NanoVDB.h:4700
bool isUnknown() const
Definition: NanoVDB.h:4704
Highest level of the data structure. Contains a tree and a world->index transform (that currently onl...
Definition: NanoVDB.h:2309
Vec3T indexToWorld(const Vec3T &xyz) const
index to world space transformation
Definition: NanoVDB.h:2365
Vec3T indexToWorldF(const Vec3T &xyz) const
index to world space transformation
Definition: NanoVDB.h:2388
typename TreeT::ValueType ValueType
Definition: NanoVDB.h:2314
int findBlindDataForSemantic(GridBlindDataSemantic semantic) const
Return the index of the blind data with specified semantic if found, otherwise -1.
Definition: NanoVDB.h:2490
const Vec3R & voxelSize() const
Return a const reference to the size of a voxel in world units.
Definition: NanoVDB.h:2354
typename TreeT::RootType RootType
Definition: NanoVDB.h:2312
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:2438
~Grid()=delete
const DataType * data() const
Definition: NanoVDB.h:2330
bool isPointData() const
Definition: NanoVDB.h:2425
uint32_t gridIndex() const
Return index of this grid in the buffer.
Definition: NanoVDB.h:2339
bool isBreadthFirst() const
Definition: NanoVDB.h:2433
uint64_t activeVoxelCount() const
Return the total number of active voxels in this tree.
Definition: NanoVDB.h:2415
const char * shortGridName() const
Return a c-string with the name of this grid, truncated to 255 characters.
Definition: NanoVDB.h:2457
const void * blindData(uint32_t n) const
Returns a const pointer to the blindData at the specified linear offset.
Definition: NanoVDB.h:2474
Vec3T indexToWorldDirF(const Vec3T &dir) const
transformation from index space direction to world space direction
Definition: NanoVDB.h:2393
typename TreeT::CoordType CoordType
Definition: NanoVDB.h:2316
Vec3T worldToIndexF(const Vec3T &xyz) const
world to index space transformation
Definition: NanoVDB.h:2384
const Map & map() const
Return a const reference to the Map for this grid.
Definition: NanoVDB.h:2357
Vec3T worldToIndexDirF(const Vec3T &dir) const
transformation from world space direction to index space direction
Definition: NanoVDB.h:2398
TreeT & tree()
Return a non-const reference to the tree.
Definition: NanoVDB.h:2348
bool isStaggered() const
Definition: NanoVDB.h:2423
const GridType & gridType() const
Definition: NanoVDB.h:2419
bool isValid() const
Methods related to the classification of this grid.
Definition: NanoVDB.h:2418
bool hasAverage() const
Definition: NanoVDB.h:2431
bool hasMinMax() const
Definition: NanoVDB.h:2428
AccessorType getAccessor() const
Return a new instance of a ReadAccessor used to access values in this grid.
Definition: NanoVDB.h:2351
const BBox< Vec3R > & worldBBox() const
Computes a AABB of active values in world space.
Definition: NanoVDB.h:2406
uint64_t gridSize() const
Return the memory footprint of the entire grid, i.e. including all nodes and blind data.
Definition: NanoVDB.h:2336
Vec3T indexToWorldDir(const Vec3T &dir) const
transformation from index space direction to world space direction
Definition: NanoVDB.h:2370
int blindDataCount() const
Return the count of blind-data encoded in this grid.
Definition: NanoVDB.h:2466
Grid & operator=(const Grid &)=delete
Vec3T indexToWorldGradF(const Vec3T &grad) const
Transforms the gradient from index space to world space.
Definition: NanoVDB.h:2403
uint64_t checksum() const
Return checksum of the grid buffer.
Definition: NanoVDB.h:2460
bool isFogVolume() const
Definition: NanoVDB.h:2422
const GridClass & gridClass() const
Definition: NanoVDB.h:2420
const GridBlindMetaData & blindMetaData(int n) const
Definition: NanoVDB.h:2483
bool isMask() const
Definition: NanoVDB.h:2426
bool hasLongGridName() const
Definition: NanoVDB.h:2430
static uint64_t memUsage()
Return memory usage in bytes for this class only.
Definition: NanoVDB.h:2333
Vec3T worldToIndex(const Vec3T &xyz) const
world to index space transformation
Definition: NanoVDB.h:2361
uint32_t gridCount() const
Return total number of grids in the buffer.
Definition: NanoVDB.h:2342
Vec3T worldToIndexDir(const Vec3T &dir) const
transformation from world space direction to index space direction
Definition: NanoVDB.h:2375
Version version() const
Definition: NanoVDB.h:2326
bool isEmpty() const
Return true if this grid is empty, i.e. contains no values or nodes.
Definition: NanoVDB.h:2463
DataType * data()
Definition: NanoVDB.h:2328
const BBox< CoordType > & indexBBox() const
Computes a AABB of active values in index space.
Definition: NanoVDB.h:2412
bool hasStdDeviation() const
Definition: NanoVDB.h:2432
const char * gridName() const
Return a c-string with the name of this grid.
Definition: NanoVDB.h:2446
const TreeT & tree() const
Return a const reference to the tree.
Definition: NanoVDB.h:2345
bool hasBBox() const
Definition: NanoVDB.h:2429
bool isLevelSet() const
Definition: NanoVDB.h:2421
TreeT TreeType
Definition: NanoVDB.h:2311
Grid(const Grid &)=delete
Disallow constructions, copy and assignment.
typename TreeT::BuildType BuildType
Definition: NanoVDB.h:2315
Vec3T indexToWorldGrad(const Vec3T &grad) const
transform the gradient from index space to world space.
Definition: NanoVDB.h:2380
bool isPointIndex() const
Definition: NanoVDB.h:2424
bool isUnknown() const
Definition: NanoVDB.h:2427
Dummy type for a 16 bit floating point values.
Definition: NanoVDB.h:192
Internal nodes of a VDB treedim(),.
Definition: NanoVDB.h:3122
Coord offsetToGlobalCoord(uint32_t n) const
Definition: NanoVDB.h:3245
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:3169
const LeafNodeType * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:3208
const DataType * data() const
Definition: NanoVDB.h:3151
bool isActive() const
Retrun true if this node or any of its child nodes contain active values.
Definition: NanoVDB.h:3253
InternalNode & operator=(const InternalNode &)=delete
static uint32_t dim()
Return the dimension, in voxel units, of this internal node (typically 8*16 or 8*16*32)
Definition: NanoVDB.h:3154
static size_t memUsage()
Return memory usage in bytes for the class.
Definition: NanoVDB.h:3157
const MaskType< LOG2DIM > & valueMask() const
Return a const reference to the bit mask of active voxels in this internal node.
Definition: NanoVDB.h:3160
CoordType origin() const
Return the origin in index space of this leaf node.
Definition: NanoVDB.h:3166
typename ChildT::CoordType CoordType
Definition: NanoVDB.h:3130
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:3178
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:3175
const 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:3184
void localToGlobalCoord(Coord &ijk) const
modifies local coordinates to global coordinates of a tile or child node
Definition: NanoVDB.h:3239
typename DataType::BuildT BuildType
Definition: NanoVDB.h:3127
ChildT ChildNodeType
Definition: NanoVDB.h:3129
bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:3193
typename DataType::ValueT ValueType
Definition: NanoVDB.h:3125
static uint32_t CoordToOffset(const CoordType &ijk)
Return the linear offset corresponding to the given coordinate.
Definition: NanoVDB.h:3217
typename DataType::StatsT FloatType
Definition: NanoVDB.h:3126
const FloatType & stdDeviation() const
Return a const reference to the standard deviation of all the active values encoded in this internal ...
Definition: NanoVDB.h:3181
static Coord OffsetToLocalCoord(uint32_t n)
Definition: NanoVDB.h:3231
bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:3199
typename ChildT::LeafNodeType LeafNodeType
Definition: NanoVDB.h:3128
InternalNode()=delete
This class cannot be constructed or deleted.
ValueType getValue(const CoordType &ijk) const
Return the value of the given voxel.
Definition: NanoVDB.h:3187
const MaskType< LOG2DIM > & childMask() const
Return a const reference to the bit mask of child nodes in this internal node.
Definition: NanoVDB.h:3163
DataType * data()
Definition: NanoVDB.h:3149
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:3172
typename ChildT::template MaskType< LOG2 > MaskType
Definition: NanoVDB.h:3133
InternalNode(const InternalNode &)=delete
Leaf nodes of the VDB tree. (defaults to 8x8x8 = 512 voxels)
Definition: NanoVDB.h:3685
const LeafNode * probeLeaf(const CoordT &) const
Definition: NanoVDB.h:3818
void setValue(const CoordT &ijk, const ValueType &v)
Sets the value at the specified location and activate its state.
Definition: NanoVDB.h:3790
void setValueOnly(uint32_t offset, const ValueType &v)
Sets the value at the specified location but leaves its state unchanged.
Definition: NanoVDB.h:3795
bool isActive(const CoordT &ijk) const
Return true if the voxel value at the given coordinate is active.
Definition: NanoVDB.h:3799
typename DataType::BuildType BuildType
Definition: NanoVDB.h:3695
const DataType * data() const
Definition: NanoVDB.h:3712
bool isActive() const
Return true if any of the voxel value are active in this leaf node.
Definition: NanoVDB.h:3803
CoordT CoordType
Definition: NanoVDB.h:3696
static uint32_t dim()
Return the dimension, in index space, of this leaf node (typically 8 as for openvdb leaf nodes!...
Definition: NanoVDB.h:3753
void updateBBox()
Updates the local bounding box of active voxels in this node.
Definition: NanoVDB.h:3884
const MaskType< LOG2DIM > & valueMask() const
Return a const reference to the bit mask of active voxels in this leaf node.
Definition: NanoVDB.h:3715
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:3811
FloatType variance() const
Return the variance of all the active values encoded in this leaf node.
Definition: NanoVDB.h:3727
bool isActive(uint32_t n) const
Definition: NanoVDB.h:3800
uint8_t flags() const
Definition: NanoVDB.h:3732
LeafNode & operator=(const LeafNode &)=delete
void localToGlobalCoord(Coord &ijk) const
Converts (in place) a local index coordinate to a global index coordinate.
Definition: NanoVDB.h:3745
ValueType maximum() const
Return a const reference to the maximum active value encoded in this leaf node.
Definition: NanoVDB.h:3721
typename DataType::FloatType FloatType
Definition: NanoVDB.h:3694
CoordT origin() const
Return the origin in index space of this leaf node.
Definition: NanoVDB.h:3735
BBox< CoordT > bbox() const
Return the bounding box in index space of active values in this leaf node.
Definition: NanoVDB.h:3756
static uint32_t CoordToOffset(const CoordT &ijk)
Return the linear offset corresponding to the given coordinate.
Definition: NanoVDB.h:3821
CoordT offsetToGlobalCoord(uint32_t n) const
Definition: NanoVDB.h:3747
void setValueOnly(const CoordT &ijk, const ValueType &v)
Definition: NanoVDB.h:3796
ValueType getValue(uint32_t offset) const
Return the voxel value at the given offset.
Definition: NanoVDB.h:3782
FloatType average() const
Return a const reference to the average of all the active values encoded in this leaf node.
Definition: NanoVDB.h:3724
ValueType getValue(const CoordT &ijk) const
Return the voxel value at the given coordinate.
Definition: NanoVDB.h:3785
MaskT< LOG2 > MaskType
Definition: NanoVDB.h:3699
static uint32_t voxelCount()
Return the total number of voxels (e.g. values) encoded in this leaf node.
Definition: NanoVDB.h:3770
LeafNode(const LeafNode &)=delete
static uint64_t memUsage()
return memory usage in bytes for the class
Definition: NanoVDB.h:3773
LeafNode()=delete
This class cannot be constructed or deleted.
ValueType minimum() const
Return a const reference to the minimum active value encoded in this leaf node.
Definition: NanoVDB.h:3718
DataType * data()
Definition: NanoVDB.h:3710
typename DataType::ValueType ValueType
Definition: NanoVDB.h:3693
static CoordT OffsetToLocalCoord(uint32_t n)
Definition: NanoVDB.h:3737
FloatType stdDeviation() const
Return a const reference to the standard deviation of all the active values encoded in this leaf node...
Definition: NanoVDB.h:3730
Definition: NanoVDB.h:1820
Iterator()
Definition: NanoVDB.h:1822
Iterator & operator=(const Iterator &)=default
uint32_t operator*() const
Definition: NanoVDB.h:1833
Iterator(uint32_t pos, const Mask *parent)
Definition: NanoVDB.h:1827
Iterator & operator++()
Definition: NanoVDB.h:1835
Bit-mask to encode active states and facilitate sequential iterators and a fast codec for I/O compres...
Definition: NanoVDB.h:1796
Iterator beginOn() const
Definition: NanoVDB.h:1899
void set(uint32_t n, bool On)
Definition: NanoVDB.h:1924
bool operator==(const Mask &other) const
Definition: NanoVDB.h:1889
static size_t memUsage()
Return the memory footprint in bytes of this Mask.
Definition: NanoVDB.h:1803
void setOn()
Set all bits on.
Definition: NanoVDB.h:1937
void toggle(uint32_t n)
Definition: NanoVDB.h:1964
uint32_t countOn() const
Definition: NanoVDB.h:1811
void setOff(uint32_t n)
Definition: NanoVDB.h:1922
Mask(const Mask &other)
Copy constructor.
Definition: NanoVDB.h:1860
Mask(bool on)
Definition: NanoVDB.h:1852
bool isOn(uint32_t n) const
Return true if the given bit is set.
Definition: NanoVDB.h:1902
void set(bool on)
Set all bits off.
Definition: NanoVDB.h:1951
void setOff()
Set all bits off.
Definition: NanoVDB.h:1944
static uint32_t bitCount()
Return the number of bits available in this Mask.
Definition: NanoVDB.h:1806
bool isOff() const
Definition: NanoVDB.h:1912
bool operator!=(const Mask &other) const
Definition: NanoVDB.h:1897
void toggle()
brief Toggle the state of all bits in the mask
Definition: NanoVDB.h:1958
Mask & operator=(const MaskT &other)
Assignment operator that works with openvdb::util::NodeMask.
Definition: NanoVDB.h:1876
static uint32_t wordCount()
Return the number of machine words used by this Mask.
Definition: NanoVDB.h:1809
WordT getWord(int n) const
Return the nth word of the bit mask, for a word of arbitrary size.
Definition: NanoVDB.h:1868
bool isOn() const
Definition: NanoVDB.h:1904
Mask()
Initialize all bits to zero.
Definition: NanoVDB.h:1847
void setOn(uint32_t n)
Set the given bit on.
Definition: NanoVDB.h:1921
Class to access points at a specific voxel location.
Definition: NanoVDB.h:4722
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:4742
typename NanoRoot< uint32_t >::LeafNodeType LeafNodeType
Definition: NanoVDB.h:4728
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:4752
PointAccessor(const UInt32Grid &grid)
Definition: NanoVDB.h:4730
uint64_t voxelPoints(const Coord &ijk, const AttT *&begin, const AttT *&end) const
get iterators over offsets to points at a specific voxel location
Definition: NanoVDB.h:4764
NodeInfo getNodeInfo(const CoordType &ijk) const
Definition: NanoVDB.h:4050
bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:4055
ReadAccessor & operator=(const ReadAccessor &)=default
const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:4065
bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:4060
typename RootT::CoordType CoordType
Definition: NanoVDB.h:4020
const RootT & root() const
Definition: NanoVDB.h:4038
ReadAccessor(const ReadAccessor &)=default
Defaults constructors.
ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:4045
typename RootT::ValueType ValueType
Definition: NanoVDB.h:4019
uint32_t getDim(const CoordType &ijk, const RayT &ray) const
Definition: NanoVDB.h:4071
ReadAccessor(const RootT &root)
Constructor from a root node.
Definition: NanoVDB.h:4036
Node caching at all (three) tree levels.
Definition: NanoVDB.h:4433
typename ReadAccessor< ValueT, -1, -1, -1 >::NodeInfo NodeInfo
Definition: NanoVDB.h:4460
CoordT CoordType
Definition: NanoVDB.h:4456
NodeInfo getNodeInfo(const CoordType &ijk) const
Definition: NanoVDB.h:4534
ValueT ValueType
Definition: NanoVDB.h:4455
bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:4551
ReadAccessor & operator=(const ReadAccessor &)=default
const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:4585
bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:4568
const NodeT * getNode() const
Return a const point to the cached node of the specified type.
Definition: NanoVDB.h:4485
const RootT & root() const
Definition: NanoVDB.h:4474
ReadAccessor(const ReadAccessor &)=default
Defaults constructors.
ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:4517
bool isCached(const CoordType &ijk) const
Definition: NanoVDB.h:4511
uint32_t getDim(const CoordType &ijk, const RayT &ray) const
Definition: NanoVDB.h:4603
ReadAccessor(const RootT &root)
Constructor from a root node.
Definition: NanoVDB.h:4463
typename ReadAccessor< ValueT, -1, -1, -1 >::NodeInfo NodeInfo
Definition: NanoVDB.h:4117
CoordT CoordType
Definition: NanoVDB.h:4113
bool isCached(const CoordType &ijk) const
Definition: NanoVDB.h:4134
NodeInfo getNodeInfo(const CoordType &ijk) const
Definition: NanoVDB.h:4149
ValueT ValueType
Definition: NanoVDB.h:4112
bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:4157
ReadAccessor & operator=(const ReadAccessor &)=default
const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:4173
bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:4165
const RootT & root() const
Definition: NanoVDB.h:4127
ReadAccessor(const ReadAccessor &)=default
Defaults constructors.
ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:4141
uint32_t getDim(const CoordType &ijk, const RayT &ray) const
Definition: NanoVDB.h:4182
ReadAccessor(const RootT &root)
Constructor from a root node.
Definition: NanoVDB.h:4120
CoordT CoordType
Definition: NanoVDB.h:4240
NodeInfo getNodeInfo(const CoordType &ijk) const
Definition: NanoVDB.h:4321
ValueT ValueType
Definition: NanoVDB.h:4239
bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:4336
ReadAccessor & operator=(const ReadAccessor &)=default
const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:4366
bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:4351
typename ReadAccessor< ValueT,-1,-1,-1 >::NodeInfo NodeInfo
Definition: NanoVDB.h:4244
const RootT & root() const
Definition: NanoVDB.h:4259
bool isCached1(const CoordType &ijk) const
Definition: NanoVDB.h:4292
ReadAccessor(const ReadAccessor &)=default
Defaults constructors.
ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:4306
bool isCached2(const CoordType &ijk) const
Definition: NanoVDB.h:4298
uint32_t getDim(const CoordType &ijk, const RayT &ray) const
Definition: NanoVDB.h:4382
ReadAccessor(const RootT &root)
Constructor from a root node.
Definition: NanoVDB.h:4247
Definition: NanoVDB.h:2298
8-bit red, green, blue, alpha packed into 32 bit unsigned int
Definition: NanoVDB.h:456
Rgba8(uint8_t r, uint8_t g, uint8_t b, uint8_t a=255u)
Definition: NanoVDB.h:470
uint32_t & packed()
Definition: NanoVDB.h:491
Rgba8(float r, float g, float b, float a=1.0f)
Definition: NanoVDB.h:472
uint32_t packed
Definition: NanoVDB.h:459
const uint8_t & r() const
Definition: NanoVDB.h:492
Rgba8 & operator=(const Rgba8 &)=default
bool operator<(const Rgba8 &rhs) const
Definition: NanoVDB.h:479
Rgba8(Rgba8 &&)=default
Rgba8()
Definition: NanoVDB.h:469
const uint8_t & b() const
Definition: NanoVDB.h:494
Rgba8(uint8_t v)
Definition: NanoVDB.h:471
uint8_t & r()
Definition: NanoVDB.h:496
const uint8_t & g() const
Definition: NanoVDB.h:493
float lengthSqr() const
Definition: NanoVDB.h:481
const uint8_t & a() const
Definition: NanoVDB.h:495
uint8_t & g()
Definition: NanoVDB.h:497
float length() const
Definition: NanoVDB.h:487
bool operator==(const Rgba8 &rhs) const
Definition: NanoVDB.h:480
const uint8_t & operator[](int n) const
Definition: NanoVDB.h:488
uint8_t & operator[](int n)
Definition: NanoVDB.h:489
uint8_t & b()
Definition: NanoVDB.h:498
uint8_t ValueType
Definition: NanoVDB.h:463
Rgba8(const Rgba8 &)=default
uint8_t & a()
Definition: NanoVDB.h:499
uint8_t c[4]
Definition: NanoVDB.h:458
Rgba8 & operator=(Rgba8 &&)=default
static const int SIZE
Definition: NanoVDB.h:462
const uint32_t & packed() const
Definition: NanoVDB.h:490
Top-most node of the VDB tree structure.
Definition: NanoVDB.h:2800
const ValueType & background() const
Return the total number of active voxels in the root and all its child nodes.
Definition: NanoVDB.h:2837
static uint64_t memUsage(uint32_t tableSize)
Return the expected memory footprint in bytes with the specified number of tiles.
Definition: NanoVDB.h:2858
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:2843
const LeafNodeType * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:2897
const DataType * data() const
Definition: NanoVDB.h:2828
typename ChildT::CoordType CoordType
Definition: NanoVDB.h:2811
RootNode & operator=(const RootNode &)=delete
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:2852
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:2849
const BBox< CoordType > & bbox() const
Return a const reference to the index bounding box of all the active values in this tree,...
Definition: NanoVDB.h:2831
AccessorType getAccessor() const
Definition: NanoVDB.h:2824
typename DataType::BuildT BuildType
Definition: NanoVDB.h:2809
RootNode(const RootNode &)=delete
ChildT ChildNodeType
Definition: NanoVDB.h:2804
bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:2872
typename DataType::ValueT ValueType
Definition: NanoVDB.h:2807
typename DataType::StatsT FloatType
Definition: NanoVDB.h:2808
const uint32_t & tileCount() const
Return the number of tiles encoded in this root node.
Definition: NanoVDB.h:2840
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:2855
uint64_t memUsage() const
Return the actual memory footprint of this root node.
Definition: NanoVDB.h:2861
bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:2883
typename ChildT::LeafNodeType LeafNodeType
Definition: NanoVDB.h:2803
typename DataType::Tile Tile
Definition: NanoVDB.h:2813
ValueType getValue(const CoordType &ijk) const
Return the value of the given voxel.
Definition: NanoVDB.h:2864
bool isEmpty() const
Return true if this RootNode is empty, i.e. contains no values or nodes.
Definition: NanoVDB.h:2881
DataType * data()
Definition: NanoVDB.h:2826
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:2846
RootNode()=delete
This class cannot be constructed or deleted.
VDB Tree, which is a thin wrapper around a RootNode.
Definition: NanoVDB.h:2544
typename RootT::ChildNodeType Node2
Definition: NanoVDB.h:2560
NodeTrait< RootT, LEVEL >::type * getFirstNode()
return a pointer to the first node at the specified level
Definition: NanoVDB.h:2654
const ValueType & background() const
Return a const reference to the background value.
Definition: NanoVDB.h:2596
const NodeT * getFirstNode() const
return a const pointer to the first node of the specified type
Definition: NanoVDB.h:2643
RootT Node3
Definition: NanoVDB.h:2559
const DataType * data() const
Definition: NanoVDB.h:2572
uint64_t activeVoxelCount() const
Return the total number of active voxels in this tree.
Definition: NanoVDB.h:2605
typename RootT::LeafNodeType LeafNodeType
Definition: NanoVDB.h:2553
Tree & operator=(const Tree &)=delete
const BBox< CoordType > & bbox() const
Return a const reference to the index bounding box of all the active values in this tree,...
Definition: NanoVDB.h:2602
uint32_t nodeCount() const
Definition: NanoVDB.h:2617
Tree()=delete
This class cannot be constructed or deleted.
AccessorType getAccessor() const
Definition: NanoVDB.h:2581
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:2587
NodeT * getFirstNode()
return a pointer to the first node of the specified type
Definition: NanoVDB.h:2633
~Tree()=delete
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:2675
bool probeValue(const CoordType &ijk, ValueType &v) const
Combines the previous two methods in a single call.
Definition: NanoVDB.h:2593
typename RootT::CoordType CoordType
Definition: NanoVDB.h:2556
const NodeTrait< RootT, LEVEL >::type * getFirstNode() const
return a const pointer to the first node of the specified level
Definition: NanoVDB.h:2664
const uint32_t & activeTileCount(uint32_t n) const
Return the total number of active tiles at the specified level of the tree.
Definition: NanoVDB.h:2610
Tree(const Tree &)=delete
static uint64_t memUsage()
return memory usage in bytes for the class
Definition: NanoVDB.h:2575
LeafNodeType Node0
Definition: NanoVDB.h:2562
const RootT & root() const
Definition: NanoVDB.h:2579
uint32_t nodeCount(int level) const
Definition: NanoVDB.h:2623
typename Node2::ChildNodeType Node1
Definition: NanoVDB.h:2561
ValueType getValue(const CoordType &ijk) const
Return the value of the given voxel (regardless of state or location in the tree.)
Definition: NanoVDB.h:2584
RootT RootType
Definition: NanoVDB.h:2552
bool isEmpty() const
Return true if this tree is empty, i.e. contains no values or nodes.
Definition: NanoVDB.h:2590
DataType * data()
Definition: NanoVDB.h:2570
typename RootT::ValueType ValueType
Definition: NanoVDB.h:2554
RootT & root()
Definition: NanoVDB.h:2577
typename RootT::BuildType BuildType
Definition: NanoVDB.h:2555
Dummy type for a voxel with a binary mask value, e.g. the active state.
Definition: NanoVDB.h:189
Vec3(const Vec3< T2 > &v)
Definition: NanoVDB.h:1060
Coord round() const
Definition: NanoVDB.h:1159
T & operator[](int i)
Definition: NanoVDB.h:1079
Vec3(T x)
Definition: NanoVDB.h:1051
ValueType min() const
Return the smallest vector component.
Definition: NanoVDB.h:1148
const T & operator[](int i) const
Definition: NanoVDB.h:1078
Vec3 & operator/=(const T &s)
Definition: NanoVDB.h:1122
Vec3(const Coord &ijk)
Definition: NanoVDB.h:1064
bool operator==(const Vec3 &rhs) const
Definition: NanoVDB.h:1068
Vec3 operator-() const
Definition: NanoVDB.h:1094
T length() const
Definition: NanoVDB.h:1093
Vec3 operator*(const Vec3 &v) const
Definition: NanoVDB.h:1095
Coord floor() const
Definition: NanoVDB.h:1157
Vec3 operator*(const T &s) const
Definition: NanoVDB.h:1099
Vec3(T x, T y, T z)
Definition: NanoVDB.h:1055
Vec3 operator/(const Vec3 &v) const
Definition: NanoVDB.h:1096
T dot(const Vec3T &v) const
Definition: NanoVDB.h:1081
Vec3 & normalize()
Definition: NanoVDB.h:1123
Vec3 & minComponent(const Vec3 &other)
Perform a component-wise minimum with the other Coord.
Definition: NanoVDB.h:1125
T lengthSqr() const
Definition: NanoVDB.h:1089
Vec3 & operator*=(const T &s)
Definition: NanoVDB.h:1115
Coord ceil() const
Definition: NanoVDB.h:1158
Vec3()=default
Vec3 cross(const Vec3T &v) const
Definition: NanoVDB.h:1083
Vec3 operator/(const T &s) const
Definition: NanoVDB.h:1100
Vec3 & operator-=(const Vec3 &v)
Definition: NanoVDB.h:1108
Vec3 & maxComponent(const Vec3 &other)
Perform a component-wise maximum with the other Coord.
Definition: NanoVDB.h:1137
bool operator!=(const Vec3 &rhs) const
Definition: NanoVDB.h:1069
Vec3 operator+(const Vec3 &v) const
Definition: NanoVDB.h:1097
Vec3 operator-(const Vec3 &v) const
Definition: NanoVDB.h:1098
ValueType max() const
Return the largest vector component.
Definition: NanoVDB.h:1153
T ValueType
Definition: NanoVDB.h:1049
static const int SIZE
Definition: NanoVDB.h:1048
Vec3 & operator=(const Vec3T &rhs)
Definition: NanoVDB.h:1071
Vec3 & operator+=(const Vec3 &v)
Definition: NanoVDB.h:1101
A simple vector class with three double components, similar to openvdb::math::Vec4.
Definition: NanoVDB.h:1189
T & operator[](int i)
Definition: NanoVDB.h:1221
Vec4 & operator*=(const T &s)
Definition: NanoVDB.h:1252
const T & operator[](int i) const
Definition: NanoVDB.h:1220
Vec4 & minComponent(const Vec4 &other)
Perform a component-wise minimum with the other Coord.
Definition: NanoVDB.h:1263
Vec4 operator/(const Vec4 &v) const
Definition: NanoVDB.h:1231
Vec4 operator*(const T &s) const
Definition: NanoVDB.h:1234
T length() const
Definition: NanoVDB.h:1228
Vec4 operator-(const Vec4 &v) const
Definition: NanoVDB.h:1233
Vec4 operator*(const Vec4 &v) const
Definition: NanoVDB.h:1230
Vec4(T x, T y, T z, T w)
Definition: NanoVDB.h:1200
bool operator!=(const Vec4 &rhs) const
Definition: NanoVDB.h:1210
Vec4 & operator=(const Vec4T &rhs)
Definition: NanoVDB.h:1212
Vec4(const Vec4< T2 > &v)
Definition: NanoVDB.h:1205
T lengthSqr() const
Definition: NanoVDB.h:1224
Vec4 & operator-=(const Vec4 &v)
Definition: NanoVDB.h:1244
Vec4 & operator+=(const Vec4 &v)
Definition: NanoVDB.h:1236
Vec4()=default
bool operator==(const Vec4 &rhs) const
Definition: NanoVDB.h:1209
Vec4 operator-() const
Definition: NanoVDB.h:1229
Vec4 operator/(const T &s) const
Definition: NanoVDB.h:1235
Vec4 & maxComponent(const Vec4 &other)
Perform a component-wise maximum with the other Coord.
Definition: NanoVDB.h:1277
T dot(const Vec4T &v) const
Definition: NanoVDB.h:1223
T ValueType
Definition: NanoVDB.h:1194
Vec4(T x)
Definition: NanoVDB.h:1196
static const int SIZE
Definition: NanoVDB.h:1193
Vec4 & operator/=(const T &s)
Definition: NanoVDB.h:1260
Vec4 operator+(const Vec4 &v) const
Definition: NanoVDB.h:1232
Vec4 & normalize()
Definition: NanoVDB.h:1261
Bit-compacted representation of all three version numbers.
Definition: NanoVDB.h:541
const char * c_str() const
Definition: NanoVDB.h:567
bool operator<(const Version &rhs) const
Definition: NanoVDB.h:557
uint32_t getPatch() const
Definition: NanoVDB.h:564
bool operator<=(const Version &rhs) const
Definition: NanoVDB.h:558
Version()
Definition: NanoVDB.h:544
bool operator==(const Version &rhs) const
Definition: NanoVDB.h:556
Version(uint32_t major, uint32_t minor, uint32_t patch)
Definition: NanoVDB.h:549
uint32_t getMajor() const
Definition: NanoVDB.h:562
bool operator>=(const Version &rhs) const
Definition: NanoVDB.h:560
uint32_t getMinor() const
Definition: NanoVDB.h:563
uint32_t id() const
Definition: NanoVDB.h:561
bool operator>(const Version &rhs) const
Definition: NanoVDB.h:559
Definition: NanoVDB.h:184
bool isValid(GridType gridType, GridClass gridClass)
return true if the combination of GridType and GridClass is valid.
Definition: NanoVDB.h:520
T Pow4(T x)
Definition: NanoVDB.h:742
GridType mapToGridType()
Maps from a templated value type to a GridType enum.
Definition: NanoVDB.h:1366
Vec3< T2 > operator*(T1 scalar, const Vec3< T2 > &vec)
Definition: NanoVDB.h:1163
float Clamp(float x, float a, float b)
Definition: NanoVDB.h:693
int32_t Floor(float x)
Definition: NanoVDB.h:711
int MinIndex(const Vec3T &v)
Definition: NanoVDB.h:810
Vec3< float > Vec3f
Definition: NanoVDB.h:1175
CoordT RoundDown(const Vec3T< RealT > &xyz)
Definition: NanoVDB.h:788
GridClass
Classes (defined in OpenVDB) that are currently supported by NanoVDB.
Definition: NanoVDB.h:253
GridType
List of types that are currently supported by NanoVDB.
Definition: NanoVDB.h:216
uint32_t CountOn(uint64_t v)
Definition: NanoVDB.h:1774
Vec3T matMult(const float *mat, const Vec3T &xyz)
Definition: NanoVDB.h:1409
Type Min(Type a, Type b)
Definition: NanoVDB.h:651
GridFlags
Grid flags which indicate what extra information is present in the grid buffer.
Definition: NanoVDB.h:277
static 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: NanoVDB.h:1696
float Sqrt(float x)
Return the square root of a floating-point value.
Definition: NanoVDB.h:795
T Pow3(T x)
Definition: NanoVDB.h:736
GridBlindDataClass
Blind-data Classes that are currently supported by NanoVDB.
Definition: NanoVDB.h:306
float Fract(float x)
Definition: NanoVDB.h:702
Vec3< T2 > operator/(T1 scalar, const Vec3< T2 > &vec)
Definition: NanoVDB.h:1168
Vec3T matMultT(const float *mat, const Vec3T &xyz)
Definition: NanoVDB.h:1443
Vec3< double > Vec3d
Definition: NanoVDB.h:1174
static int64_t PtrDiff(const T1 *p, const T2 *q)
Definition: NanoVDB.h:433
const char * toStr(GridType gridType)
Retuns a c-string used to describe a GridType.
Definition: NanoVDB.h:239
T Sign(const T &x)
Return the sign of the given value as an integer (either -1, 0 or 1).
Definition: NanoVDB.h:807
Type Max(Type a, Type b)
Definition: NanoVDB.h:672
T Abs(T x)
Definition: NanoVDB.h:747
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:4657
GridBlindDataSemantic
Blind-data Semantics that are currently understood by NanoVDB.
Definition: NanoVDB.h:313
CoordT Round(const Vec3T< RealT > &xyz)
bool isFloatingPoint(GridType gridType)
return true if the GridType maps to a floating point value.
Definition: NanoVDB.h:507
static DstT * PtrAdd(SrcT *p, int64_t offset)
Definition: NanoVDB.h:440
int32_t Ceil(float x)
Definition: NanoVDB.h:720
bool isApproxZero(const Type &x)
Definition: NanoVDB.h:645
int MaxIndex(const Vec3T &v)
Definition: NanoVDB.h:827
T Pow2(T x)
Definition: NanoVDB.h:730
uint64_t AlignUp(uint64_t byteCount)
round up byteSize to the nearest wordSize, e.g. to align to machine word: AlignUp<sizeof(size_t)(n)
Definition: NanoVDB.h:847
static 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: NanoVDB.h:1667
const std::enable_if<!VecTraits< T >::IsVec, T >::type & max(const T &a, const T &b)
Definition: Composite.h:107
const std::enable_if<!VecTraits< T >::IsVec, T >::type & min(const T &a, const T &b)
Definition: Composite.h:103
math::Extrema extrema(const IterT &iter, bool threaded=true)
Iterate over a scalar grid and compute extrema (min/max) of the values of the voxels that are visited...
Definition: Statistics.h:354
Iterator begin() const
Definition: NanoVDB.h:1610
BBox()
Definition: NanoVDB.h:1611
bool is_divisible() const
Definition: NanoVDB.h:1628
BBox(const CoordT &min, const CoordT &max)
Definition: NanoVDB.h:1615
CoordT dim() const
Definition: NanoVDB.h:1635
bool empty() const
Return true if this bounding box is empty, i.e. uninitialized.
Definition: NanoVDB.h:1632
BBox expandBy(typename CoordT::ValueType padding) const
Return a new instance that is expanded by the specified padding.
Definition: NanoVDB.h:1652
bool isInside(const BBox &b) const
Definition: NanoVDB.h:1638
uint64_t volume() const
Definition: NanoVDB.h:1636
BBox(BBox &other, const SplitT &)
Definition: NanoVDB.h:1620
BBox< Vec3< RealT > > asReal() const
Definition: NanoVDB.h:1645
bool isInside(const CoordT &p) const
Definition: NanoVDB.h:1637
Vec3T dim() const
Definition: NanoVDB.h:1556
bool isInside(const Vec3T &p) const
Definition: NanoVDB.h:1557
BBox()
Definition: NanoVDB.h:1538
BBox(const Coord &min, const Coord &max)
Definition: NanoVDB.h:1547
bool empty() const
Definition: NanoVDB.h:1553
BBox(const Vec3T &min, const Vec3T &max)
Definition: NanoVDB.h:1543
Vec3T Vec3Type
Definition: NanoVDB.h:1533
typename Vec3T::ValueType ValueType
Definition: NanoVDB.h:1534
BBox(const BaseBBox< Coord > &bbox)
Definition: NanoVDB.h:1552
Definition: NanoVDB.h:1524
Definition: NanoVDB.h:1479
Vec3T mCoord[2]
Definition: NanoVDB.h:1480
BaseBBox()
Definition: NanoVDB.h:1516
const Vec3T & max() const
Definition: NanoVDB.h:1488
Vec3T & operator[](int i)
Definition: NanoVDB.h:1484
bool operator!=(const BaseBBox &rhs) const
Definition: NanoVDB.h:1482
Coord & translate(const Vec3T &xyz)
Definition: NanoVDB.h:1489
Vec3T & max()
Definition: NanoVDB.h:1486
bool isInside(const Vec3T &xyz)
Definition: NanoVDB.h:1506
BaseBBox(const Vec3T &min, const Vec3T &max)
Definition: NanoVDB.h:1517
Vec3T & min()
Definition: NanoVDB.h:1485
const Vec3T & min() const
Definition: NanoVDB.h:1487
BaseBBox & expand(const Vec3T &xyz)
Definition: NanoVDB.h:1496
bool operator==(const BaseBBox &rhs) const
Definition: NanoVDB.h:1481
const Vec3T & operator[](int i) const
Definition: NanoVDB.h:1483
float type
Definition: NanoVDB.h:420
float Type
Definition: NanoVDB.h:419
float type
Definition: NanoVDB.h:406
float Type
Definition: NanoVDB.h:405
float type
Definition: NanoVDB.h:413
float Type
Definition: NanoVDB.h:412
float type
Definition: NanoVDB.h:427
float Type
Definition: NanoVDB.h:426
float type
Definition: NanoVDB.h:399
float Type
Definition: NanoVDB.h:398
bool Type
Definition: NanoVDB.h:391
bool type
Definition: NanoVDB.h:392
Maps one type (e.g. the build types above) to other (actual) types.
Definition: NanoVDB.h:383
T Type
Definition: NanoVDB.h:384
T type
Definition: NanoVDB.h:385
static double value()
Definition: NanoVDB.h:606
static float value()
Definition: NanoVDB.h:601
Delta for small floating-point offsets.
Definition: NanoVDB.h:597
double FloatType
Definition: NanoVDB.h:1347
bool FloatType
Definition: NanoVDB.h:1359
bool FloatType
Definition: NanoVDB.h:1353
Definition: NanoVDB.h:1340
float FloatType
Definition: NanoVDB.h:1341
Definition: NanoVDB.h:2069
uint32_t mFlags
Definition: NanoVDB.h:2073
static uint64_t memUsage(uint64_t blindDataCount=0)
return memory usage in bytes for the class (note this computes for all blindMetaData structures....
Definition: NanoVDB.h:2080
void setBlindData(void *ptr)
Definition: NanoVDB.h:2085
GridType mDataType
Definition: NanoVDB.h:2076
GridBlindDataSemantic mSemantic
Definition: NanoVDB.h:2074
int64_t mByteOffset
Definition: NanoVDB.h:2071
GridBlindDataClass mDataClass
Definition: NanoVDB.h:2075
const T * getBlindData() const
Definition: NanoVDB.h:2088
uint64_t mElementCount
Definition: NanoVDB.h:2072
Struct with all the member data of the Grid (useful during serialization of an openvdb grid)
Definition: NanoVDB.h:2186
uint32_t mFlags
Definition: NanoVDB.h:2191
void setMinMaxOn(bool on=true)
Definition: NanoVDB.h:2207
void setAverageOn(bool on=true)
Definition: NanoVDB.h:2231
Vec3T applyInverseJacobianF(const Vec3T &xyz) const
Definition: NanoVDB.h:2275
Vec3T applyIJTF(const Vec3T &xyz) const
Definition: NanoVDB.h:2277
uint32_t mBlindMetadataCount
Definition: NanoVDB.h:2202
Version mVersion
Definition: NanoVDB.h:2190
GridType mGridType
Definition: NanoVDB.h:2200
uint64_t mMagic
Definition: NanoVDB.h:2188
const void * treePtr() const
Definition: NanoVDB.h:2283
void setStdDeviationOn(bool on=true)
Definition: NanoVDB.h:2239
Vec3T applyIJT(const Vec3T &xyz) const
Definition: NanoVDB.h:2266
GridClass mGridClass
Definition: NanoVDB.h:2199
void setBBoxOn(bool on=true)
Definition: NanoVDB.h:2215
void setLongGridNameOn(bool on=true)
Definition: NanoVDB.h:2223
Vec3T applyJacobianF(const Vec3T &xyz) const
Definition: NanoVDB.h:2273
uint64_t mGridSize
Definition: NanoVDB.h:2194
BBox< Vec3R > mWorldBBox
Definition: NanoVDB.h:2197
Vec3T applyInverseJacobian(const Vec3T &xyz) const
Definition: NanoVDB.h:2264
void * treePtr()
Definition: NanoVDB.h:2280
void setBreadthFirstOn(bool on=true)
Definition: NanoVDB.h:2247
uint64_t mChecksum
Definition: NanoVDB.h:2189
Vec3T applyMapF(const Vec3T &xyz) const
Definition: NanoVDB.h:2269
uint32_t mGridCount
Definition: NanoVDB.h:2193
Vec3R mVoxelSize
Definition: NanoVDB.h:2198
Vec3T applyInverseMap(const Vec3T &xyz) const
Definition: NanoVDB.h:2260
Map mMap
Definition: NanoVDB.h:2196
Vec3T applyJacobian(const Vec3T &xyz) const
Definition: NanoVDB.h:2262
void setFlagsOff()
Definition: NanoVDB.h:2206
int64_t mBlindMetadataOffset
Definition: NanoVDB.h:2201
const GridBlindMetaData * blindMetaData(uint32_t n) const
Returns a const reference to the blindMetaData at the specified linear offset.
Definition: NanoVDB.h:2288
uint32_t mGridIndex
Definition: NanoVDB.h:2192
Vec3T applyMap(const Vec3T &xyz) const
Definition: NanoVDB.h:2258
Vec3T applyInverseMapF(const Vec3T &xyz) const
Definition: NanoVDB.h:2271
const typename GridT::TreeType Type
Definition: NanoVDB.h:2535
const typename GridT::TreeType type
Definition: NanoVDB.h:2536
defines a tree type from a grid type while perserving constness
Definition: NanoVDB.h:2528
typename GridT::TreeType Type
Definition: NanoVDB.h:2529
typename GridT::TreeType type
Definition: NanoVDB.h:2530
Struct with all the member data of the InternalNode (useful during serialization of an openvdb Intern...
Definition: NanoVDB.h:3044
void setOrigin(const T &ijk)
Definition: NanoVDB.h:3100
StatsT mAverage
Definition: NanoVDB.h:3070
typename ChildT::CoordType CoordT
Definition: NanoVDB.h:3048
void setChild(uint32_t n, const void *ptr)
Definition: NanoVDB.h:3074
MaskT mChildMask
Definition: NanoVDB.h:3066
void setValue(uint32_t n, const ValueT &v)
Definition: NanoVDB.h:3081
ChildT * getChild(uint32_t n)
Returns a pointer to the child node at the specifed linear offset.
Definition: NanoVDB.h:3088
const ChildT * getChild(uint32_t n) const
Definition: NanoVDB.h:3093
InternalData(const InternalData &)=delete
void setDev(const StatsT &v)
Definition: NanoVDB.h:3110
void setMin(const ValueT &v)
Definition: NanoVDB.h:3107
typename ChildT::FloatType StatsT
Definition: NanoVDB.h:3047
const ValueT & getMin() const
Definition: NanoVDB.h:3102
const StatsT & stdDeviation() const
Definition: NanoVDB.h:3105
typename ChildT::BuildType BuildT
Definition: NanoVDB.h:3046
void setMax(const ValueT &v)
Definition: NanoVDB.h:3108
StatsT mStdDevi
Definition: NanoVDB.h:3071
BBox< CoordT > mBBox
Definition: NanoVDB.h:3063
const ValueT & getMax() const
Definition: NanoVDB.h:3103
ValueT mMaximum
Definition: NanoVDB.h:3069
MaskT mValueMask
Definition: NanoVDB.h:3065
typename ChildT::template MaskType< LOG2DIM > MaskT
Definition: NanoVDB.h:3049
void setAvg(const StatsT &v)
Definition: NanoVDB.h:3109
typename ChildT::ValueType ValueT
Definition: NanoVDB.h:3045
InternalData & operator=(const InternalData &)=delete
ValueT mMinimum
Definition: NanoVDB.h:3068
const StatsT & average() const
Definition: NanoVDB.h:3104
InternalData()=delete
This class cannot be constructed or deleted.
uint64_t mFlags
Definition: NanoVDB.h:3064
static constexpr uint8_t bitWidth()
Definition: NanoVDB.h:3521
float getValue(uint32_t i) const
Definition: NanoVDB.h:3522
LeafData & operator=(const LeafData &)=delete
uint16_t ArrayType
Definition: NanoVDB.h:3517
LeafData()=delete
This class cannot be constructed or deleted.
static constexpr uint8_t bitWidth()
Definition: NanoVDB.h:3471
float getValue(uint32_t i) const
Definition: NanoVDB.h:3472
LeafData & operator=(const LeafData &)=delete
LeafData()=delete
This class cannot be constructed or deleted.
uint8_t ArrayType
Definition: NanoVDB.h:3467
static constexpr uint8_t bitWidth()
Definition: NanoVDB.h:3499
float getValue(uint32_t i) const
Definition: NanoVDB.h:3500
LeafData & operator=(const LeafData &)=delete
LeafData()=delete
This class cannot be constructed or deleted.
uint8_t ArrayType
Definition: NanoVDB.h:3495
size_t memUsage() const
Definition: NanoVDB.h:3543
float getValue(uint32_t i) const
Definition: NanoVDB.h:3545
uint8_t bitWidth() const
Definition: NanoVDB.h:3542
LeafData & operator=(const LeafData &)=delete
LeafData()=delete
This class cannot be constructed or deleted.
static size_t memUsage(uint32_t bitWidth)
Definition: NanoVDB.h:3544
void setOrigin(const T &ijk)
Definition: NanoVDB.h:3670
bool getDev() const
Definition: NanoVDB.h:3658
bool getMin() const
Definition: NanoVDB.h:3655
void setMax(const ValueType &)
Definition: NanoVDB.h:3665
bool getValue(uint32_t i) const
Definition: NanoVDB.h:3654
bool getMax() const
Definition: NanoVDB.h:3656
void setDev(const FloatType &)
Definition: NanoVDB.h:3667
bool getAvg() const
Definition: NanoVDB.h:3657
MaskT< LOG2DIM > mValueMask
Definition: NanoVDB.h:3651
void setValue(uint32_t offset, bool)
Definition: NanoVDB.h:3659
LeafData & operator=(const LeafData &)=delete
LeafData()=delete
This class cannot be constructed or deleted.
void setAvg(const FloatType &)
Definition: NanoVDB.h:3666
void setMin(const ValueType &)
Definition: NanoVDB.h:3664
void setOrigin(const T &ijk)
Definition: NanoVDB.h:3627
bool getDev() const
Definition: NanoVDB.h:3614
bool BuildType
Definition: NanoVDB.h:3598
void setMin(const bool &)
Definition: NanoVDB.h:3621
bool getMin() const
Definition: NanoVDB.h:3611
void setMax(const bool &)
Definition: NanoVDB.h:3622
void setValue(uint32_t offset, bool v)
Definition: NanoVDB.h:3615
bool getValue(uint32_t i) const
Definition: NanoVDB.h:3610
uint8_t mFlags
Definition: NanoVDB.h:3605
bool getMax() const
Definition: NanoVDB.h:3612
MaskT< LOG2DIM > ArrayType
Definition: NanoVDB.h:3600
bool getAvg() const
Definition: NanoVDB.h:3613
MaskT< LOG2DIM > mValueMask
Definition: NanoVDB.h:3606
MaskT< LOG2DIM > mValues
Definition: NanoVDB.h:3607
LeafData & operator=(const LeafData &)=delete
CoordT mBBoxMin
Definition: NanoVDB.h:3603
bool ValueType
Definition: NanoVDB.h:3597
void setDev(const bool &)
Definition: NanoVDB.h:3624
LeafData()=delete
This class cannot be constructed or deleted.
void setAvg(const bool &)
Definition: NanoVDB.h:3623
bool FloatType
Definition: NanoVDB.h:3599
Stuct with all the member data of the LeafNode (useful during serialization of an openvdb LeafNode)
Definition: NanoVDB.h:3356
void setOrigin(const T &ijk)
Definition: NanoVDB.h:3396
ValueType mMaximum
Definition: NanoVDB.h:3371
FloatType getDev() const
Definition: NanoVDB.h:3388
typename FloatTraits< ValueT >::FloatType FloatType
Definition: NanoVDB.h:3361
void setMin(const ValueType &v)
Definition: NanoVDB.h:3390
void setMax(const ValueType &v)
Definition: NanoVDB.h:3391
FloatType mAverage
Definition: NanoVDB.h:3372
void setDev(const FloatType &v)
Definition: NanoVDB.h:3393
uint8_t mFlags
Definition: NanoVDB.h:3367
LeafData(const LeafData &)=delete
ValueType getMin() const
Definition: NanoVDB.h:3385
ValueType getValue(uint32_t i) const
Definition: NanoVDB.h:3377
void setValue(uint32_t offset, const ValueType &value)
Definition: NanoVDB.h:3379
ValueT ValueType
Definition: NanoVDB.h:3359
MaskT< LOG2DIM > mValueMask
Definition: NanoVDB.h:3368
void setValueOnly(uint32_t offset, const ValueType &value)
Definition: NanoVDB.h:3378
LeafData & operator=(const LeafData &)=delete
CoordT mBBoxMin
Definition: NanoVDB.h:3365
ValueType getMax() const
Definition: NanoVDB.h:3386
LeafData()=delete
This class cannot be constructed or deleted.
ValueT ArrayType
Definition: NanoVDB.h:3362
FloatType getAvg() const
Definition: NanoVDB.h:3387
ValueType mMinimum
Definition: NanoVDB.h:3370
void setAvg(const FloatType &v)
Definition: NanoVDB.h:3392
ValueT BuildType
Definition: NanoVDB.h:3360
FloatType mStdDevi
Definition: NanoVDB.h:3373
Base-class for quantized float leaf nodes.
Definition: NanoVDB.h:3408
void setOrigin(const T &ijk)
Definition: NanoVDB.h:3455
float ValueType
Definition: NanoVDB.h:3411
float getMin() const
return the quantized minimum of the active values in this node
Definition: NanoVDB.h:3430
void setDev(float dev)
Definition: NanoVDB.h:3452
void setMin(float min)
Definition: NanoVDB.h:3443
float getAvg() const
return the quantized average of the active values in this node
Definition: NanoVDB.h:3436
uint8_t mFlags
Definition: NanoVDB.h:3416
float mQuantum
Definition: NanoVDB.h:3420
MaskT< LOG2DIM > mValueMask
Definition: NanoVDB.h:3417
void init(float min, float max, uint8_t bitWidth)
Definition: NanoVDB.h:3423
uint16_t mAvg
Definition: NanoVDB.h:3421
CoordT mBBoxMin
Definition: NanoVDB.h:3414
float mMinimum
Definition: NanoVDB.h:3419
float FloatType
Definition: NanoVDB.h:3412
void setMax(float max)
Definition: NanoVDB.h:3446
void setAvg(float avg)
Definition: NanoVDB.h:3449
float getDev() const
return the quantized standard deviation of the active values in this node
Definition: NanoVDB.h:3440
float getMax() const
return the quantized maximum of the active values in this node
Definition: NanoVDB.h:3433
Definition: NanoVDB.h:3688
static uint32_t dim()
Definition: NanoVDB.h:3689
Defines an affine transform and its inverse represented as a 3x3 matrix and a vec3 translation.
Definition: NanoVDB.h:1999
double mTaperD
Definition: NanoVDB.h:2007
Vec3T applyInverseJacobianF(const Vec3T &xyz) const
Definition: NanoVDB.h:2037
Vec3T applyIJTF(const Vec3T &xyz) const
Definition: NanoVDB.h:2042
double mVecD[3]
Definition: NanoVDB.h:2006
float mInvMatF[9]
Definition: NanoVDB.h:2001
Vec3T applyIJT(const Vec3T &xyz) const
Definition: NanoVDB.h:2040
Vec3T applyJacobianF(const Vec3T &xyz) const
Definition: NanoVDB.h:2021
Vec3T applyInverseJacobian(const Vec3T &xyz) const
Definition: NanoVDB.h:2035
void set(const Mat4T &mat, const Mat4T &invMat, double taper)
Definition: NanoVDB.h:2046
double mInvMatD[9]
Definition: NanoVDB.h:2005
Vec3T applyMapF(const Vec3T &xyz) const
Definition: NanoVDB.h:2016
float mMatF[9]
Definition: NanoVDB.h:2000
Vec3T applyInverseMap(const Vec3T &xyz) const
Definition: NanoVDB.h:2024
double mMatD[9]
Definition: NanoVDB.h:2004
Vec3T applyJacobian(const Vec3T &xyz) const
Definition: NanoVDB.h:2019
Vec3T applyMap(const Vec3T &xyz) const
Definition: NanoVDB.h:2014
Vec3T applyInverseMapF(const Vec3T &xyz) const
Definition: NanoVDB.h:2029
float mTaperF
Definition: NanoVDB.h:2003
float mVecF[3]
Definition: NanoVDB.h:2002
Maximum floating-point values.
Definition: NanoVDB.h:638
static T value()
Definition: NanoVDB.h:639
Trait to map from LEVEL to node type.
Definition: NanoVDB.h:3934
typename GridOrTreeOrRootT::LeafNodeType type
Definition: NanoVDB.h:2105
typename GridOrTreeOrRootT::LeafNodeType Type
Definition: NanoVDB.h:2104
typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType Type
Definition: NanoVDB.h:2119
typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType type
Definition: NanoVDB.h:2120
typename GridOrTreeOrRootT::RootType::ChildNodeType Type
Definition: NanoVDB.h:2133
typename GridOrTreeOrRootT::RootType::ChildNodeType type
Definition: NanoVDB.h:2134
typename GridOrTreeOrRootT::RootType type
Definition: NanoVDB.h:2148
typename GridOrTreeOrRootT::RootType Type
Definition: NanoVDB.h:2147
const typename GridOrTreeOrRootT::LeafNodeType Type
Definition: NanoVDB.h:2111
const typename GridOrTreeOrRootT::LeafNodeType type
Definition: NanoVDB.h:2112
const typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType Type
Definition: NanoVDB.h:2126
const typename GridOrTreeOrRootT::RootType::ChildNodeType::ChildNodeType type
Definition: NanoVDB.h:2127
const typename GridOrTreeOrRootT::RootType::ChildNodeType type
Definition: NanoVDB.h:2141
const typename GridOrTreeOrRootT::RootType::ChildNodeType Type
Definition: NanoVDB.h:2140
const typename GridOrTreeOrRootT::RootType type
Definition: NanoVDB.h:2156
const typename GridOrTreeOrRootT::RootType Type
Definition: NanoVDB.h:2155
Struct to derive node type from its level in a given grid, tree or root while perserving constness.
Definition: NanoVDB.h:2097
uint32_t mLevel
Definition: NanoVDB.h:4025
ValueType mMaximum
Definition: NanoVDB.h:4028
FloatType mAverage
Definition: NanoVDB.h:4029
CoordType mBBoxMax
Definition: NanoVDB.h:4032
ValueType mMinimum
Definition: NanoVDB.h:4027
CoordType mBBoxMin
Definition: NanoVDB.h:4031
FloatType mStdDevi
Definition: NanoVDB.h:4030
Definition: NanoVDB.h:2729
ValueT value
Definition: NanoVDB.h:2749
void setChild(const CoordType &k, const ChildT *ptr, const RootData *data)
Definition: NanoVDB.h:2731
uint32_t state
Definition: NanoVDB.h:2748
KeyT key
Definition: NanoVDB.h:2746
void setValue(const CoordType &k, bool s, const ValueType &v)
Definition: NanoVDB.h:2737
CoordT origin() const
Definition: NanoVDB.h:2745
bool isChild() const
Definition: NanoVDB.h:2744
int64_t child
Definition: NanoVDB.h:2747
Struct with all the member data of the RootNode (useful during serialization of an openvdb RootNode)
Definition: NanoVDB.h:2688
const ChildT * getChild(const Tile *tile) const
Definition: NanoVDB.h:2774
StatsT mAverage
Definition: NanoVDB.h:2725
typename ChildT::CoordType CoordT
Definition: NanoVDB.h:2691
static CoordT KeyToCoord(const KeyT &key)
Definition: NanoVDB.h:2707
RootData()=delete
This class cannot be constructed or deleted.
ValueT mBackground
Definition: NanoVDB.h:2722
const Tile * tile(uint32_t n) const
Returns a non-const reference to the tile at the specified linear offset.
Definition: NanoVDB.h:2755
uint32_t mTableSize
Definition: NanoVDB.h:2720
void setDev(const StatsT &v)
Definition: NanoVDB.h:2788
uint64_t KeyT
Return a key based on the coordinates of a voxel.
Definition: NanoVDB.h:2697
void setMin(const ValueT &v)
Definition: NanoVDB.h:2785
Tile * tile(uint32_t n)
Definition: NanoVDB.h:2760
typename ChildT::FloatType StatsT
Definition: NanoVDB.h:2692
const ValueT & getMin() const
Definition: NanoVDB.h:2780
const StatsT & stdDeviation() const
Definition: NanoVDB.h:2783
typename ChildT::BuildType BuildT
Definition: NanoVDB.h:2690
ChildT * getChild(const Tile *tile)
Returns a const reference to the child node in the specified tile.
Definition: NanoVDB.h:2769
void setMax(const ValueT &v)
Definition: NanoVDB.h:2786
RootData & operator=(const RootData &)=delete
StatsT mStdDevi
Definition: NanoVDB.h:2726
RootData(const RootData &)=delete
static KeyT CoordToKey(const CoordType &ijk)
Definition: NanoVDB.h:2699
BBox< CoordT > mBBox
Definition: NanoVDB.h:2719
const ValueT & getMax() const
Definition: NanoVDB.h:2781
ValueT mMaximum
Definition: NanoVDB.h:2724
void setAvg(const StatsT &v)
Definition: NanoVDB.h:2787
typename ChildT::ValueType ValueT
Definition: NanoVDB.h:2689
ValueT mMinimum
Definition: NanoVDB.h:2723
const StatsT & average() const
Definition: NanoVDB.h:2782
T ElementType
Definition: NanoVDB.h:1321
static T scalar(const T &s)
Definition: NanoVDB.h:1322
static ElementType scalar(const T &v)
Definition: NanoVDB.h:1333
typename T::ValueType ElementType
Definition: NanoVDB.h:1332
Definition: NanoVDB.h:1312
static double value()
Definition: NanoVDB.h:590
static float value()
Definition: NanoVDB.h:585
Tolerance for floating-point comparison.
Definition: NanoVDB.h:581
Definition: NanoVDB.h:2502
void setFirstNode(const NodeT *node)
Definition: NanoVDB.h:2517
RootT * getRoot()
Definition: NanoVDB.h:2512
const RootT * getRoot() const
Definition: NanoVDB.h:2514
void setRoot(const RootT *root)
Definition: NanoVDB.h:2510
uint64_t mVoxelCount
Definition: NanoVDB.h:2507
T type
Definition: NanoVDB.h:348
C++11 implementation of std::enable_if.
Definition: NanoVDB.h:342
C++11 implementation of std::is_floating_point.
Definition: NanoVDB.h:356
static const bool value
Definition: NanoVDB.h:357
C++11 implementation of std::is_same.
Definition: NanoVDB.h:327
static constexpr bool value
Definition: NanoVDB.h:328
Metafunction used to determine if the first template parameter is a specialization of the class templ...
Definition: NanoVDB.h:369
static const bool value
Definition: NanoVDB.h:370
Definition: NanoVDB.h:3053
ValueT value
Definition: NanoVDB.h:3054
Tile(const Tile &)=delete
Tile()=delete
This class cannot be constructed or deleted.
int64_t child
Definition: NanoVDB.h:3055
Tile & operator=(const Tile &)=delete