OpenVDB 9.1.0
Loading...
Searching...
No Matches
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
124typedef signed char int8_t;
125typedef short int16_t;
126typedef int int32_t;
127typedef long long int64_t;
128typedef unsigned char uint8_t;
129typedef unsigned int uint32_t;
130typedef unsigned short uint16_t;
131typedef 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
184namespace nanovdb {
185
186// --------------------------> Build types <------------------------------------
187
188/// @brief Dummy type for a voxel with a binary mask value, e.g. the active state
189class ValueMask {};
190
191/// @brief Dummy type for a 16 bit floating point values
192class Half {};
193
194/// @brief Dummy type for a 4bit quantization of float point values
195class Fp4 {};
196
197/// @brief Dummy type for a 8bit quantization of float point values
198class Fp8 {};
199
200/// @brief Dummy type for a 16bit quantization of float point values
201class Fp16 {};
202
203/// @brief Dummy type for a variable bit quantization of floating point values
204class 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)
216enum 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
239inline 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
253enum 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
265inline 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.
277enum 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
289inline 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
306enum 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
313enum 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
325template<typename T1, typename T2>
327{
328 static constexpr bool value = false;
329};
330
331template<typename T>
332struct 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
340template <bool, typename T = void>
342{
343};
344
345template <typename T>
346struct 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
354template<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;
367template<typename AnyType, template<typename...> class TemplateType>
369{
370 static const bool value = false;
371};
372template<typename... Args, template<typename...> class TemplateType>
373struct 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
381template <typename T>
383{
384 using Type = T;
385 using type = T;
386};
387
388template<>
390{
391 using Type = bool;
392 using type = bool;
393};
394
395template<>
397{
398 using Type = float;
399 using type = float;
400};
401
402template<>
404{
405 using Type = float;
406 using type = float;
407};
408
409template<>
411{
412 using Type = float;
413 using type = float;
414};
415
416template<>
418{
419 using Type = float;
420 using type = float;
421};
422
423template<>
425{
426 using Type = float;
427 using type = float;
428};
429
430// --------------------------> PtrDiff PtrAdd <------------------------------------
431
432template <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
439template <typename DstT, typename SrcT>
440__hostdev__ inline static DstT* PtrAdd(SrcT *p, int64_t offset)
441{
443 return reinterpret_cast<DstT*>(reinterpret_cast<char*>(p) + offset);
444}
445
446template <typename DstT, typename SrcT>
447__hostdev__ inline static const DstT* PtrAdd(const SrcT *p, int64_t offset)
448{
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
455class 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;
461public:
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
502using PackedRGBA8 = Rgba8;// for backwards compatibility
503
504// --------------------------> isValue(GridType, GridClass) <------------------------------------
505
506/// @brief return true if the GridType maps to a floating point value.
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
541{
542 uint32_t mData;// 11 + 11 + 10 bit packing of major + minor + patch
543public:
545 uint32_t(NANOVDB_MINOR_VERSION_NUMBER) << 10 |
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
580template<typename T>
582template<>
583struct Tolerance<float>
584{
585 __hostdev__ static float value() { return 1e-8f; }
586};
587template<>
588struct Tolerance<double>
589{
590 __hostdev__ static double value() { return 1e-15; }
591};
592//@}
593
594//@{
595/// Delta for small floating-point offsets
596template<typename T>
597struct Delta;
598template<>
599struct Delta<float>
600{
601 __hostdev__ static float value() { return 1e-5f; }
602};
603template<>
604struct Delta<double>
605{
606 __hostdev__ static double value() { return 1e-9; }
607};
608//@}
609
610//@{
611/// Maximum floating-point values
612template<typename T>
613struct Maximum;
614#if defined(__CUDA_ARCH__) || defined(__HIP__)
615template<>
616struct Maximum<int>
617{
618 __hostdev__ static int value() { return 2147483647; }
619};
620template<>
621struct Maximum<uint32_t>
622{
623 __hostdev__ static uint32_t value() { return 4294967295; }
624};
625template<>
626struct Maximum<float>
627{
628 __hostdev__ static float value() { return 1e+38f; }
629};
630template<>
631struct Maximum<double>
632{
633 __hostdev__ static double value() { return 1e+308; }
634};
635#else
636template<typename T>
638{
639 static T value() { return std::numeric_limits<T>::max(); }
640};
641#endif
642//@}
643
644template<typename Type>
645__hostdev__ inline bool isApproxZero(const Type& x)
646{
647 return !(x > Tolerance<Type>::value()) && !(x < -Tolerance<Type>::value());
648}
649
650template<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}
671template<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
729template<typename T>
730__hostdev__ inline T Pow2(T x)
731{
732 return x * x;
733}
734
735template<typename T>
736__hostdev__ inline T Pow3(T x)
737{
738 return x * x * x;
739}
740
741template<typename T>
742__hostdev__ inline T Pow4(T x)
743{
744 return Pow2(x * x);
745}
746template<typename T>
747__hostdev__ inline T Abs(T x)
748{
749 return x < 0 ? -x : x;
750}
751
752template<>
753__hostdev__ inline float Abs(float x)
754{
755 return fabs(x);
756}
757
758template<>
759__hostdev__ inline double Abs(double x)
760{
761 return fabs(x);
762}
763
764template<>
765__hostdev__ inline int Abs(int x)
766{
767 return abs(x);
768}
769
770template<typename CoordT, typename RealT, template<typename> class Vec3T>
771__hostdev__ inline CoordT Round(const Vec3T<RealT>& xyz);
772
773template<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
781template<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
787template<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).
806template <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
809template<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
826template<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
846template<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
856template<typename> class Vec3;
857
858/// @brief Signed (i, j, k) 32-bit integer coordinate class, similar to openvdb::math::Coord
859class Coord
860{
861 int32_t mVec[3]; // private member data - three signed index coordinates
862public:
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.
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
1042template<typename T>
1043class Vec3
1044{
1045 T mVec[3];
1046
1047public:
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
1162template<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}
1167template<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
1187template<typename T>
1188class Vec4
1189{
1190 T mVec[4];
1191
1192public:
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
1291template<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}
1296template<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
1309template<typename T, int Rank = (is_specialization<T, Vec3>::value ||
1310 is_specialization<T, Vec4>::value ||
1311 is_same<T, Rgba8>::value) ? 1 : 0>
1313
1314template<typename T>
1315struct 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
1325template<typename T>
1326struct 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
1338template<typename T, int = sizeof(typename TensorTraits<T>::ElementType)>
1340{
1341 using FloatType = float;
1342};
1343
1344template<typename T>
1345struct FloatTraits<T, 8>
1346{
1347 using FloatType = double;
1348};
1349
1350template<>
1351struct FloatTraits<bool, 1>
1352{
1353 using FloatType = bool;
1354};
1355
1356template<>
1358{
1359 using FloatType = bool;
1360};
1361
1362// ----------------------------> mapping ValueType -> GridType <--------------------------------------
1363
1364/// @brief Maps from a templated value type to a GridType enum
1365template<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;
1373 return GridType::Int16;
1375 return GridType::Int32;
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;
1383 return GridType::UInt32;
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
1408template<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
1416template<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
1424template<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
1432template<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
1442template<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
1450template<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
1458template<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
1466template<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)
1477template<typename Vec3T>
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
1515protected:
1517 __hostdev__ BaseBBox(const Vec3T& min, const Vec3T& max)
1518 : mCoord{min, max}
1519 {
1520 }
1521}; // BaseBBox
1522
1523template<typename Vec3T, bool = is_floating_point<typename Vec3T::ValueType>::value>
1524struct 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.
1530template<typename Vec3T>
1531struct 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!
1568template<typename CoordT>
1569struct 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 }
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.
1794template<uint32_t LOG2DIM>
1795class 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
1801public:
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
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
1966private:
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
1998struct 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
2045template<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
2068struct 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.
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
2096template<typename GridOrTreeOrRootT, int LEVEL>
2098
2099// Partial template specialization of above Node struct
2100template<typename GridOrTreeOrRootT>
2101struct 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};
2107template<typename GridOrTreeOrRootT>
2108struct 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
2115template<typename GridOrTreeOrRootT>
2116struct 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};
2122template<typename GridOrTreeOrRootT>
2123struct 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};
2129template<typename GridOrTreeOrRootT>
2130struct 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};
2136template<typename GridOrTreeOrRootT>
2137struct 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};
2143template<typename GridOrTreeOrRootT>
2144struct 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
2151template<typename GridOrTreeOrRootT>
2152struct 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!
2185struct 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
2297template <typename BuildT, int LEVEL0 = -1, int LEVEL1 = -1, int LEVEL2 = -1>
2299
2300template <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
2307template<typename TreeT>
2308class Grid : private GridData
2309{
2310public:
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
2485private:
2486 static_assert(sizeof(GridData) % NANOVDB_DATA_ALIGNMENT == 0, "sizeof(GridData) is misaligned");
2487}; // Class Grid
2488
2489template<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
2500template<int ROOT_LEVEL = 3>
2501struct 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
2526template<typename GridT>
2528{
2529 using Type = typename GridT::TreeType;
2530 using type = typename GridT::TreeType;
2531};
2532template<typename GridT>
2533struct 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.
2542template<typename RootT>
2543class 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
2550public:
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
2669private:
2670 static_assert(sizeof(DataType) % NANOVDB_DATA_ALIGNMENT == 0, "sizeof(TreeData) is misaligned");
2671
2672}; // Tree class
2673
2674template<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!
2686template<typename ChildT>
2687struct 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 }
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.
2798template<typename ChildT>
2799class RootNode : private RootData<ChildT>
2800{
2801public:
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
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
2907private:
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!
3042template<typename ChildT, uint32_t LOG2DIM>
3043struct 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(),
3120template<typename ChildT, uint32_t Log2Dim = ChildT::LOG2DIM + 1>
3121class InternalNode : private InternalData<ChildT, Log2Dim>
3122{
3123public:
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
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 {
3247 Coord ijk = InternalNode::OffsetToLocalCoord(n);
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
3254 {
3255 return DataType::mFlags & uint32_t(2);
3256 }
3257
3258private:
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!
3354template<typename ValueT, typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3355struct 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
3406template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3407struct 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!
3461template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3462struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData<Fp4, CoordT, MaskT, LOG2DIM>
3463 : public LeafFnBase<CoordT, MaskT, LOG2DIM>
3464{
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
3489template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3490struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData<Fp8, CoordT, MaskT, LOG2DIM>
3491 : public LeafFnBase<CoordT, MaskT, LOG2DIM>
3492{
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
3511template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3512struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData<Fp16, CoordT, MaskT, LOG2DIM>
3513 : public LeafFnBase<CoordT, MaskT, LOG2DIM>
3514{
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
3534template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3535struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) LeafData<FpN, CoordT, MaskT, LOG2DIM>
3536 : public LeafFnBase<CoordT, MaskT, LOG2DIM>
3537{
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
3592template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3593struct 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
3637template<typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
3638struct 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)
3680template<typename BuildT,
3681 typename CoordT = Coord,
3682 template<uint32_t> class MaskT = Mask,
3683 uint32_t Log2Dim = 3>
3684class LeafNode : private LeafData<BuildT, CoordT, MaskT, Log2Dim>
3685{
3686public:
3688 {
3689 __hostdev__ static uint32_t dim() { return 1u; }
3690 }; // Voxel
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.
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
3839private:
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
3883template<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
3919template<typename BuildT>
3921template<typename BuildT>
3923template<typename BuildT>
3925template<typename BuildT>
3927template<typename BuildT>
3929template<typename BuildT>
3931
3932/// @brief Trait to map from LEVEL to node type
3933template<typename BuildT, int LEVEL>
3935
3936// Partial template specialization of above Node struct
3937template<typename BuildT>
3938struct NanoNode<BuildT, 0>
3939{
3942};
3943template<typename BuildT>
3944struct NanoNode<BuildT, 1>
3945{
3948};
3949template<typename BuildT>
3950struct NanoNode<BuildT, 2>
3951{
3954};
3955template<typename BuildT>
3956struct 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
4009template <typename BuildT>
4010class 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)
4018public:
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
4076private:
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
4091template <typename BuildT, int LEVEL0>
4092class 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
4111public:
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
4190private:
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
4212template <typename BuildT, int LEVEL0, int LEVEL1>
4213class 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
4238public:
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
4397private:
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
4431template <typename BuildT>
4432class 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
4454public:
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
4467 : mKeys{CoordType::max(), CoordType::max(), CoordType::max()}
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
4620private:
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
4656template <int LEVEL0 = -1, int LEVEL1 = -1, int LEVEL2 = -1, typename ValueT = float>
4658{
4660}
4661
4662template <int LEVEL0 = -1, int LEVEL1 = -1, int LEVEL2 = -1, typename ValueT = float>
4664{
4666}
4667
4668template <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
4690public:
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
4720template<typename AttT>
4721class PointAccessor : public DefaultReadAccessor<uint32_t>
4722{
4724 const UInt32Grid* mGrid;
4725 const AttT* mData;
4726
4727public:
4729
4731 : AccT(grid.tree().root())
4732 , mGrid(&grid)
4733 , mData(reinterpret_cast<const AttT*>(grid.blindData(0)))
4734 {
4735 NANOVDB_ASSERT(grid.gridType() == GridType::UInt32);
4736 NANOVDB_ASSERT((grid.gridClass() == GridClass::PointIndex && is_same<uint32_t, AttT>::value) ||
4737 (grid.gridClass() == GridClass::PointData && is_same<Vec3f, AttT>::value));
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
const CoordT & operator*() const
Definition: NanoVDB.h:1608
Iterator(const BBox &b)
Definition: NanoVDB.h:1581
Iterator operator++(int)
Definition: NanoVDB.h:1600
Iterator & operator++()
Definition: NanoVDB.h:1586
Signed (i, j, k) 32-bit integer coordinate class, similar to openvdb::math::Coord.
Definition: NanoVDB.h:860
Coord & operator&=(int n)
Definition: NanoVDB.h:940
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)
Definition: NanoVDB.h:970
Coord & minComponent(const Coord &other)
Perform a component-wise minimum with the other Coord.
Definition: NanoVDB.h:979
Coord operator+(const Coord &rhs) const
Definition: NanoVDB.h:961
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
Coord & operator+=(const Coord &rhs)
Definition: NanoVDB.h:963
Coord & operator=(const CoordT &other)
Assignment operator that works with openvdb::Coord.
Definition: NanoVDB.h:913
int32_t y() const
Definition: NanoVDB.h:890
Coord & maxComponent(const Coord &other)
Perform a component-wise maximum with the other Coord.
Definition: NanoVDB.h:991
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
const ValueType & operator[](IndexType i) const
Return a const reference to the given Coord component.
Definition: NanoVDB.h:905
int32_t z() const
Definition: NanoVDB.h:891
Coord & operator<<=(uint32_t n)
Definition: NanoVDB.h:947
int32_t x() const
Definition: NanoVDB.h:889
Coord(ValueType i, ValueType j, ValueType k)
Initializes coordinate to the given signed integers.
Definition: NanoVDB.h:879
int32_t & y()
Definition: NanoVDB.h:894
Coord offsetBy(ValueType n) const
Definition: NanoVDB.h:1007
int32_t & x()
Definition: NanoVDB.h:893
Coord(ValueType *ptr)
Definition: NanoVDB.h:884
bool operator!=(const Coord &rhs) const
Definition: NanoVDB.h:939
static Coord min()
Definition: NanoVDB.h:899
ValueType & operator[](IndexType i)
Return a non-const reference to the given Coord component.
Definition: NanoVDB.h:909
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
Coord offsetBy(ValueType dx, ValueType dy, ValueType dz) const
Definition: NanoVDB.h:1002
static Coord max()
Definition: NanoVDB.h:897
int32_t & z()
Definition: NanoVDB.h:895
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
int32_t ValueType
Definition: NanoVDB.h:863
Coord & operator+=(int n)
Definition: NanoVDB.h:954
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
const GridBlindMetaData & blindMetaData(int n) const
Definition: NanoVDB.h:4710
bool isPointData() const
Definition: NanoVDB.h:4701
uint32_t gridIndex() const
Definition: NanoVDB.h:4693
const char * shortGridName() const
Definition: NanoVDB.h:4695
uint64_t activeVoxelCount() const
Definition: NanoVDB.h:4711
bool isStaggered() const
Definition: NanoVDB.h:4703
bool isValid() const
Definition: NanoVDB.h:4691
const BBox< Vec3R > & worldBBox() const
Definition: NanoVDB.h:4706
GridType gridType() const
Definition: NanoVDB.h:4696
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
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
Vec3R voxelSize() const
Definition: NanoVDB.h:4708
const BBox< Coord > & indexBBox() const
Definition: NanoVDB.h:4707
const Map & map() const
Definition: NanoVDB.h:4705
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
const GridBlindMetaData & blindMetaData(int n) const
Definition: NanoVDB.h:2483
Vec3T indexToWorldF(const Vec3T &xyz) const
index to world space transformation
Definition: NanoVDB.h:2388
const GridType & gridType() const
Definition: NanoVDB.h:2419
typename TreeT::ValueType ValueType
Definition: NanoVDB.h:2314
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
bool isPointData() const
Definition: NanoVDB.h:2425
uint32_t gridIndex() const
Return index of this grid in the buffer.
Definition: NanoVDB.h:2339
const char * shortGridName() const
Return a c-string with the name of this grid, truncated to 255 characters.
Definition: NanoVDB.h:2457
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
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
Vec3T worldToIndexDirF(const Vec3T &dir) const
transformation from world space direction to index space direction
Definition: NanoVDB.h:2398
const Vec3R & voxelSize() const
Return a const reference to the size of a voxel in world units.
Definition: NanoVDB.h:2354
bool isStaggered() const
Definition: NanoVDB.h:2423
const GridClass & gridClass() const
Definition: NanoVDB.h:2420
bool isValid() const
Methods related to the classification of this grid.
Definition: NanoVDB.h:2418
DataType * data()
Definition: NanoVDB.h:2328
bool hasAverage() const
Definition: NanoVDB.h:2431
bool hasMinMax() const
Definition: NanoVDB.h:2428
const char * gridName() const
Return a c-string with the name of this grid.
Definition: NanoVDB.h:2446
const BBox< Vec3R > & worldBBox() const
Computes a AABB of active values in world space.
Definition: NanoVDB.h:2406
AccessorType getAccessor() const
Return a new instance of a ReadAccessor used to access values in this grid.
Definition: NanoVDB.h:2351
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
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 TreeT & tree() const
Return a const reference to the tree.
Definition: NanoVDB.h:2345
bool isMask() const
Definition: NanoVDB.h:2426
const BBox< CoordType > & indexBBox() const
Computes a AABB of active values in index space.
Definition: NanoVDB.h:2412
bool hasLongGridName() const
Definition: NanoVDB.h:2430
const void * blindData(uint32_t n) const
Returns a const pointer to the blindData at the specified linear offset.
Definition: NanoVDB.h:2474
static uint64_t memUsage()
Return memory usage in bytes for this class only.
Definition: NanoVDB.h:2333
TreeT & tree()
Return a non-const reference to the tree.
Definition: NanoVDB.h:2348
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
const Map & map() const
Return a const reference to the Map for this grid.
Definition: NanoVDB.h:2357
bool hasStdDeviation() const
Definition: NanoVDB.h:2432
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
Grid & operator=(const Grid &)=delete
Vec3T indexToWorldGrad(const Vec3T &grad) const
transform the gradient from index space to world space.
Definition: NanoVDB.h:2380
const DataType * data() const
Definition: NanoVDB.h:2330
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 MaskType< LOG2DIM > & valueMask() const
Return a const reference to the bit mask of active voxels in this internal node.
Definition: NanoVDB.h:3160
bool isActive() const
Retrun true if this node or any of its child nodes contain active values.
Definition: NanoVDB.h:3253
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
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
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
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 MaskType< LOG2DIM > & childMask() const
Return a const reference to the bit mask of child nodes in this internal node.
Definition: NanoVDB.h:3163
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
DataType * data()
Definition: NanoVDB.h:3149
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
const LeafNodeType * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:3208
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
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
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
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
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
typename ChildT::template MaskType< LOG2 > MaskType
Definition: NanoVDB.h:3133
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 DataType * data() const
Definition: NanoVDB.h:3151
InternalNode & operator=(const InternalNode &)=delete
InternalNode(const InternalNode &)=delete
Leaf nodes of the VDB tree. (defaults to 8x8x8 = 512 voxels)
Definition: NanoVDB.h:3685
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 MaskType< LOG2DIM > & valueMask() const
Return a const reference to the bit mask of active voxels in this leaf node.
Definition: NanoVDB.h:3715
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
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
LeafNode & operator=(const LeafNode &)=delete
bool isActive(uint32_t n) const
Definition: NanoVDB.h:3800
DataType * data()
Definition: NanoVDB.h:3710
uint8_t flags() const
Definition: NanoVDB.h:3732
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
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
const LeafNode * probeLeaf(const CoordT &) const
Definition: NanoVDB.h:3818
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
BBox< CoordT > bbox() const
Return the bounding box in index space of active values in this leaf node.
Definition: NanoVDB.h:3756
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
typename DataType::ValueType ValueType
Definition: NanoVDB.h:3693
static CoordT OffsetToLocalCoord(uint32_t n)
Definition: NanoVDB.h:3737
const DataType * data() const
Definition: NanoVDB.h:3712
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
Mask & operator=(const MaskT &other)
Assignment operator that works with openvdb::util::NodeMask.
Definition: NanoVDB.h:1876
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
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
ReadAccessor & operator=(const ReadAccessor &)=default
NodeInfo getNodeInfo(const CoordType &ijk) const
Definition: NanoVDB.h:4050
bool isActive(const CoordType &ijk) const
Definition: NanoVDB.h:4055
bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:4060
typename RootT::CoordType CoordType
Definition: NanoVDB.h:4020
ReadAccessor(const ReadAccessor &)=default
Defaults constructors.
const RootT & root() const
Definition: NanoVDB.h:4038
ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:4045
typename RootT::ValueType ValueType
Definition: NanoVDB.h:4019
const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:4065
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
ReadAccessor & operator=(const ReadAccessor &)=default
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
const NodeT * getNode() const
Return a const point to the cached node of the specified type.
Definition: NanoVDB.h:4485
bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:4568
ReadAccessor(const ReadAccessor &)=default
Defaults constructors.
const RootT & root() const
Definition: NanoVDB.h:4474
ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:4517
bool isCached(const CoordType &ijk) const
Definition: NanoVDB.h:4511
const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:4585
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
ReadAccessor & operator=(const ReadAccessor &)=default
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
bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:4165
ReadAccessor(const ReadAccessor &)=default
Defaults constructors.
const RootT & root() const
Definition: NanoVDB.h:4127
ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:4141
const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:4173
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
ReadAccessor & operator=(const ReadAccessor &)=default
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
bool probeValue(const CoordType &ijk, ValueType &v) const
Definition: NanoVDB.h:4351
typename ReadAccessor< ValueT,-1,-1,-1 >::NodeInfo NodeInfo
Definition: NanoVDB.h:4244
bool isCached1(const CoordType &ijk) const
Definition: NanoVDB.h:4292
ReadAccessor(const ReadAccessor &)=default
Defaults constructors.
const RootT & root() const
Definition: NanoVDB.h:4259
ValueType getValue(const CoordType &ijk) const
Definition: NanoVDB.h:4306
bool isCached2(const CoordType &ijk) const
Definition: NanoVDB.h:4298
const LeafT * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:4366
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
const uint32_t & packed() const
Definition: NanoVDB.h:490
Rgba8(uint8_t r, uint8_t g, uint8_t b, uint8_t a=255u)
Definition: NanoVDB.h:470
Rgba8 & operator=(const Rgba8 &)=default
const uint8_t & g() const
Definition: NanoVDB.h:493
uint8_t & g()
Definition: NanoVDB.h:497
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 & a() const
Definition: NanoVDB.h:495
bool operator<(const Rgba8 &rhs) const
Definition: NanoVDB.h:479
Rgba8(Rgba8 &&)=default
const uint8_t & r() const
Definition: NanoVDB.h:492
Rgba8()
Definition: NanoVDB.h:469
Rgba8(uint8_t v)
Definition: NanoVDB.h:471
float lengthSqr() const
Definition: NanoVDB.h:481
Rgba8 & operator=(Rgba8 &&)=default
const uint8_t & b() const
Definition: NanoVDB.h:494
uint32_t & packed()
Definition: NanoVDB.h:491
uint8_t & r()
Definition: NanoVDB.h:496
uint8_t & operator[](int n)
Definition: NanoVDB.h:489
float length() const
Definition: NanoVDB.h:487
const uint8_t & operator[](int n) const
Definition: NanoVDB.h:488
bool operator==(const Rgba8 &rhs) const
Definition: NanoVDB.h:480
uint8_t ValueType
Definition: NanoVDB.h:463
Rgba8(const Rgba8 &)=default
uint8_t & b()
Definition: NanoVDB.h:498
uint8_t c[4]
Definition: NanoVDB.h:458
uint8_t & a()
Definition: NanoVDB.h:499
static const int SIZE
Definition: NanoVDB.h:462
Top-most node of the VDB tree structure.
Definition: NanoVDB.h:2800
const uint32_t & tileCount() const
Return the number of tiles encoded in this root node.
Definition: NanoVDB.h:2840
static uint64_t memUsage(uint32_t tableSize)
Return the expected memory footprint in bytes with the specified number of tiles.
Definition: NanoVDB.h:2858
typename ChildT::CoordType CoordType
Definition: NanoVDB.h:2811
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
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
RootNode & operator=(const RootNode &)=delete
DataType * data()
Definition: NanoVDB.h:2826
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
const LeafNodeType * probeLeaf(const CoordType &ijk) const
Definition: NanoVDB.h:2897
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
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
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
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
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
const ValueType & background() const
Return the total number of active voxels in the root and all its child nodes.
Definition: NanoVDB.h:2837
bool isEmpty() const
Return true if this RootNode is empty, i.e. contains no values or nodes.
Definition: NanoVDB.h:2881
RootNode()=delete
This class cannot be constructed or deleted.
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 DataType * data() const
Definition: NanoVDB.h:2828
VDB Tree, which is a thin wrapper around a RootNode.
Definition: NanoVDB.h:2544
typename RootT::ChildNodeType Node2
Definition: NanoVDB.h:2560
RootT Node3
Definition: NanoVDB.h:2559
const NodeT * getFirstNode() const
return a const pointer to the first node of the specified type
Definition: NanoVDB.h:2643
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
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
DataType * data()
Definition: NanoVDB.h:2570
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
NodeTrait< RootT, LEVEL >::type * getFirstNode()
return a pointer to the first node at the specified level
Definition: NanoVDB.h:2654
~Tree()=delete
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
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
RootT & root()
Definition: NanoVDB.h:2577
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
NodeT * getFirstNode()
return a pointer to the first node of the specified type
Definition: NanoVDB.h:2633
uint32_t nodeCount(int level) const
Definition: NanoVDB.h:2623
typename Node2::ChildNodeType Node1
Definition: NanoVDB.h:2561
const NodeTrait< RootT, LEVEL >::type * getFirstNode() const
return a const pointer to the first node of the specified level
Definition: NanoVDB.h:2664
const RootT & root() const
Definition: NanoVDB.h:2579
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
const ValueType & background() const
Return a const reference to the background value.
Definition: NanoVDB.h:2596
bool isEmpty() const
Return true if this tree is empty, i.e. contains no values or nodes.
Definition: NanoVDB.h:2590
Tree & operator=(const Tree &)=delete
typename RootT::ValueType ValueType
Definition: NanoVDB.h:2554
const DataType * data() const
Definition: NanoVDB.h:2572
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
Vec3(T x)
Definition: NanoVDB.h:1051
ValueType min() const
Return the smallest vector component.
Definition: NanoVDB.h:1148
Vec3 & minComponent(const Vec3 &other)
Perform a component-wise minimum with the other Coord.
Definition: NanoVDB.h:1125
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
Vec3 & maxComponent(const Vec3 &other)
Perform a component-wise maximum with the other Coord.
Definition: NanoVDB.h:1137
Vec3 & operator-=(const Vec3 &v)
Definition: NanoVDB.h:1108
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 & operator=(const Vec3T &rhs)
Definition: NanoVDB.h:1071
Vec3 & operator*=(const T &s)
Definition: NanoVDB.h:1115
Vec3 & normalize()
Definition: NanoVDB.h:1123
T lengthSqr() const
Definition: NanoVDB.h:1089
Vec3 & operator+=(const Vec3 &v)
Definition: NanoVDB.h:1101
Coord ceil() const
Definition: NanoVDB.h:1158
Vec3()=default
const T & operator[](int i) const
Definition: NanoVDB.h:1078
Vec3 cross(const Vec3T &v) const
Definition: NanoVDB.h:1083
Vec3 operator/(const T &s) const
Definition: NanoVDB.h:1100
bool operator!=(const Vec3 &rhs) const
Definition: NanoVDB.h:1069
Vec3 operator+(const Vec3 &v) const
Definition: NanoVDB.h:1097
T & operator[](int i)
Definition: NanoVDB.h:1079
Vec3 operator-(const Vec3 &v) const
Definition: NanoVDB.h:1098
ValueType max() const
Return the largest vector component.
Definition: NanoVDB.h:1153
Vec3 & operator/=(const T &s)
Definition: NanoVDB.h:1122
T ValueType
Definition: NanoVDB.h:1049
static const int SIZE
Definition: NanoVDB.h:1048
A simple vector class with three double components, similar to openvdb::math::Vec4.
Definition: NanoVDB.h:1189
Vec4 & normalize()
Definition: NanoVDB.h:1261
Vec4 & operator/=(const T &s)
Definition: NanoVDB.h:1260
Vec4 operator/(const Vec4 &v) const
Definition: NanoVDB.h:1231
Vec4 & operator+=(const Vec4 &v)
Definition: NanoVDB.h:1236
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 T &s)
Definition: NanoVDB.h:1252
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 & maxComponent(const Vec4 &other)
Perform a component-wise maximum with the other Coord.
Definition: NanoVDB.h:1277
Vec4(const Vec4< T2 > &v)
Definition: NanoVDB.h:1205
T lengthSqr() const
Definition: NanoVDB.h:1224
const T & operator[](int i) const
Definition: NanoVDB.h:1220
Vec4 & operator=(const Vec4T &rhs)
Definition: NanoVDB.h:1212
Vec4 & operator-=(const Vec4 &v)
Definition: NanoVDB.h:1244
Vec4()=default
bool operator==(const Vec4 &rhs) const
Definition: NanoVDB.h:1209
Vec4 & minComponent(const Vec4 &other)
Perform a component-wise minimum with the other Coord.
Definition: NanoVDB.h:1263
T & operator[](int i)
Definition: NanoVDB.h:1221
Vec4 operator-() const
Definition: NanoVDB.h:1229
Vec4 operator/(const T &s) const
Definition: NanoVDB.h:1235
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 Vec4 &v) const
Definition: NanoVDB.h:1232
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
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
uint32_t CountOn(uint64_t v)
Definition: NanoVDB.h:1774
const char * toStr(GridType gridType)
Retuns a c-string used to describe a GridType.
Definition: NanoVDB.h:239
float Fract(float x)
Definition: NanoVDB.h:702
int MinIndex(const Vec3T &v)
Definition: NanoVDB.h:810
static int64_t PtrDiff(const T1 *p, const T2 *q)
Definition: NanoVDB.h:433
T Abs(T x)
Definition: NanoVDB.h:747
bool isApproxZero(const Type &x)
Definition: NanoVDB.h:645
Type Min(Type a, Type b)
Definition: NanoVDB.h:651
T Sign(const T &x)
Return the sign of the given value as an integer (either -1, 0 or 1).
Definition: NanoVDB.h:807
T Pow4(T x)
Definition: NanoVDB.h:742
Vec3< float > Vec3f
Definition: NanoVDB.h:1175
Vec3T matMult(const float *mat, const Vec3T &xyz)
Definition: NanoVDB.h:1409
static DstT * PtrAdd(SrcT *p, int64_t offset)
Definition: NanoVDB.h:440
GridClass
Classes (defined in OpenVDB) that are currently supported by NanoVDB.
Definition: NanoVDB.h:253
Vec3T matMultT(const float *mat, const Vec3T &xyz)
Definition: NanoVDB.h:1443
GridType
List of types that are currently supported by NanoVDB.
Definition: NanoVDB.h:216
float Sqrt(float x)
Return the square root of a floating-point value.
Definition: NanoVDB.h:795
int MaxIndex(const Vec3T &v)
Definition: NanoVDB.h:827
Vec3< T2 > operator/(T1 scalar, const Vec3< T2 > &vec)
Definition: NanoVDB.h:1168
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
GridFlags
Grid flags which indicate what extra information is present in the grid buffer.
Definition: NanoVDB.h:277
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
Vec3< T2 > operator*(T1 scalar, const Vec3< T2 > &vec)
Definition: NanoVDB.h:1163
int32_t Ceil(float x)
Definition: NanoVDB.h:720
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
int32_t Floor(float x)
Definition: NanoVDB.h:711
GridBlindDataClass
Blind-data Classes that are currently supported by NanoVDB.
Definition: NanoVDB.h:306
Vec3< double > Vec3d
Definition: NanoVDB.h:1174
GridType mapToGridType()
Maps from a templated value type to a GridType enum.
Definition: NanoVDB.h:1366
CoordT Round(const Vec3T< RealT > &xyz)
T Pow3(T x)
Definition: NanoVDB.h:736
CoordT RoundDown(const Vec3T< RealT > &xyz)
Definition: NanoVDB.h:788
GridBlindDataSemantic
Blind-data Semantics that are currently understood by NanoVDB.
Definition: NanoVDB.h:313
bool isFloatingPoint(GridType gridType)
return true if the GridType maps to a floating point value.
Definition: NanoVDB.h:507
T Pow2(T x)
Definition: NanoVDB.h:730
float Clamp(float x, float a, float b)
Definition: NanoVDB.h:693
Type Max(Type a, Type b)
Definition: NanoVDB.h:672
bool isValid(GridType gridType, GridClass gridClass)
return true if the combination of GridType and GridClass is valid.
Definition: NanoVDB.h:520
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
Coord & translate(const Vec3T &xyz)
Definition: NanoVDB.h:1489
BaseBBox & expand(const Vec3T &xyz)
Definition: NanoVDB.h:1496
bool operator!=(const BaseBBox &rhs) const
Definition: NanoVDB.h:1482
const Vec3T & operator[](int i) const
Definition: NanoVDB.h:1483
Vec3T & operator[](int i)
Definition: NanoVDB.h:1484
bool isInside(const Vec3T &xyz)
Definition: NanoVDB.h:1506
const Vec3T & min() const
Definition: NanoVDB.h:1487
BaseBBox(const Vec3T &min, const Vec3T &max)
Definition: NanoVDB.h:1517
Vec3T & min()
Definition: NanoVDB.h:1485
bool operator==(const BaseBBox &rhs) const
Definition: NanoVDB.h:1481
Vec3T & max()
Definition: NanoVDB.h:1486
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
uint64_t mElementCount
Definition: NanoVDB.h:2072
const T * getBlindData() const
Definition: NanoVDB.h:2088
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
const void * treePtr() const
Definition: NanoVDB.h:2283
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
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
const GridBlindMetaData * blindMetaData(uint32_t n) const
Returns a const reference to the blindMetaData at the specified linear offset.
Definition: NanoVDB.h:2288
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
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
InternalData(const InternalData &)=delete
const ValueT & getMax() const
Definition: NanoVDB.h:3103
void setDev(const StatsT &v)
Definition: NanoVDB.h:3110
const ChildT * getChild(uint32_t n) const
Definition: NanoVDB.h:3093
void setMin(const ValueT &v)
Definition: NanoVDB.h:3107
typename ChildT::FloatType StatsT
Definition: NanoVDB.h:3047
typename ChildT::BuildType BuildT
Definition: NanoVDB.h:3046
const ValueT & getMin() const
Definition: NanoVDB.h:3102
void setMax(const ValueT &v)
Definition: NanoVDB.h:3108
StatsT mStdDevi
Definition: NanoVDB.h:3071
BBox< CoordT > mBBox
Definition: NanoVDB.h:3063
ChildT * getChild(uint32_t n)
Returns a pointer to the child node at the specifed linear offset.
Definition: NanoVDB.h:3088
ValueT mMaximum
Definition: NanoVDB.h:3069
const StatsT & stdDeviation() const
Definition: NanoVDB.h:3105
const StatsT & average() const
Definition: NanoVDB.h:3104
MaskT mValueMask
Definition: NanoVDB.h:3065
typename ChildT::template MaskType< LOG2DIM > MaskT
Definition: NanoVDB.h:3049
InternalData & operator=(const InternalData &)=delete
void setAvg(const StatsT &v)
Definition: NanoVDB.h:3109
typename ChildT::ValueType ValueT
Definition: NanoVDB.h:3045
ValueT mMinimum
Definition: NanoVDB.h:3068
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
LeafData & operator=(const LeafData &)=delete
uint8_t bitWidth() const
Definition: NanoVDB.h:3542
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
LeafData & operator=(const LeafData &)=delete
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()=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
LeafData & operator=(const LeafData &)=delete
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
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 & operator=(const LeafData &)=delete
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
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
StatsT mAverage
Definition: NanoVDB.h:2725
typename ChildT::CoordType CoordT
Definition: NanoVDB.h:2691
static CoordT KeyToCoord(const KeyT &key)
Definition: NanoVDB.h:2707
const ChildT * getChild(const Tile *tile) const
Definition: NanoVDB.h:2774
RootData()=delete
This class cannot be constructed or deleted.
ValueT mBackground
Definition: NanoVDB.h:2722
Tile * tile(uint32_t n)
Definition: NanoVDB.h:2760
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
const ValueT & getMax() const
Definition: NanoVDB.h:2781
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
typename ChildT::FloatType StatsT
Definition: NanoVDB.h:2692
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
const ValueT & getMin() const
Definition: NanoVDB.h:2780
void setMax(const ValueT &v)
Definition: NanoVDB.h:2786
StatsT mStdDevi
Definition: NanoVDB.h:2726
RootData(const RootData &)=delete
static KeyT CoordToKey(const CoordType &ijk)
Definition: NanoVDB.h:2699
RootData & operator=(const RootData &)=delete
BBox< CoordT > mBBox
Definition: NanoVDB.h:2719
ValueT mMaximum
Definition: NanoVDB.h:2724
const StatsT & stdDeviation() const
Definition: NanoVDB.h:2783
const StatsT & average() const
Definition: NanoVDB.h:2782
void setAvg(const StatsT &v)
Definition: NanoVDB.h:2787
typename ChildT::ValueType ValueT
Definition: NanoVDB.h:2689
ValueT mMinimum
Definition: NanoVDB.h:2723
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
const RootT * getRoot() const
Definition: NanoVDB.h:2514
void setFirstNode(const NodeT *node)
Definition: NanoVDB.h:2517
RootT * getRoot()
Definition: NanoVDB.h:2512
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 & operator=(const Tile &)=delete
Tile()=delete
This class cannot be constructed or deleted.
int64_t child
Definition: NanoVDB.h:3055