bug-fix to nanovdb related to new I/O (#1699)

kmuseth · web-flow · commit 863738248f8d · 2023-10-31T00:49:47.000-07:00
* bug-fix to nanovdb related to new I/O

Signed-off-by: Ken Museth &lt;ken.museth@gmail.com&gt;

* improved GridData::isValid

Signed-off-by: Ken Museth &lt;ken.museth@gmail.com&gt;

---------

Signed-off-by: Ken Museth &lt;ken.museth@gmail.com&gt;
diff --git a/nanovdb/nanovdb/NanoVDB.h b/nanovdb/nanovdb/NanoVDB.h
@@ -249,60 +249,37 @@ namespace nanovdb {
 // --------------------------> Build types <------------------------------------
 
 /// @brief Dummy type for a voxel whose value equals an offset into an external value array
-class ValueIndex
-{
-};
+class ValueIndex{};
 
 /// @brief Dummy type for a voxel whose value equals an offset into an external value array of active values
-class ValueOnIndex
-{
-};
+class ValueOnIndex{};
 
 /// @brief Like @c ValueIndex but with a mutable mask
-class ValueIndexMask
-{
-};
+class ValueIndexMask{};
 
 /// @brief Like @c ValueOnIndex but with a mutable mask
-class ValueOnIndexMask
-{
-};
+class ValueOnIndexMask{};
 
 /// @brief Dummy type for a voxel whose value equals its binary active state
-class ValueMask
-{
-};
+class ValueMask{};
 
-/// @brief Dummy type for a 16 bit floating point values
-class Half
-{
-};
+/// @brief Dummy type for a 16 bit floating point values (placeholder for IEEE 754 Half)
+class Half{};
 
 /// @brief Dummy type for a 4bit quantization of float point values
-class Fp4
-{
-};
+class Fp4{};
 
 /// @brief Dummy type for a 8bit quantization of float point values
-class Fp8
-{
-};
+class Fp8{};
 
 /// @brief Dummy type for a 16bit quantization of float point values
-class Fp16
-{
-};
+class Fp16{};
 
 /// @brief Dummy type for a variable bit quantization of floating point values
-class FpN
-{
-};
+class FpN{};
 
-/// @dummy type for indexing points into voxels
-class Point
-{
-};
-//using Points = Point;// for backwards compatibility
+/// @brief Dummy type for indexing points into voxels
+class Point{};
 
 // --------------------------> GridType <------------------------------------
 
@@ -760,7 +737,7 @@ __hostdev__ inline static T* alignPtr(T* p)
     return reinterpret_cast<T*>( (uint8_t*)p + alignmentPadding(p) );
 }
 
-/// @brief offset the specified pointer so it is aligned.
+/// @brief offset the specified const pointer so it is aligned.
 template <typename T>
 __hostdev__ inline static const T* alignPtr(const T* p)
 {
@@ -863,10 +840,10 @@ __hostdev__ inline bool isIndex(GridType gridType)
 // --------------------------> memcpy64 <------------------------------------
 
 /// @brief copy 64 bit words from @c src to @c dst
-/// @param dst pointer to destination
-/// @param src pointer to source
+/// @param dst 64 bit aligned pointer to destination
+/// @param src 64 bit aligned pointer to source
 /// @param word_count number of 64 bit words to be copied
-/// @return destination pointer
+/// @return destination pointer @c dst
 /// @warning @c src and @c dst cannot overlap and should both be 64 bit aligned
 __hostdev__ inline static void* memcpy64(void *dst, const void *src, size_t word_count)
 {
@@ -948,13 +925,16 @@ class Version
 {
     uint32_t mData; // 11 + 11 + 10 bit packing of major + minor + patch
 public:
+    /// @brief Default constructor
     __hostdev__ Version()
         : mData(uint32_t(NANOVDB_MAJOR_VERSION_NUMBER) << 21 |
                 uint32_t(NANOVDB_MINOR_VERSION_NUMBER) << 10 |
                 uint32_t(NANOVDB_PATCH_VERSION_NUMBER))
     {
     }
+    /// @brief Constructor from a raw uint32_t data representation
     __hostdev__ Version(uint32_t data) : mData(data) {}
+    /// @brief Constructor from major.minor.patch version numbers
     __hostdev__ Version(uint32_t major, uint32_t minor, uint32_t patch)
         : mData(major << 21 | minor << 10 | patch)
     {
@@ -970,14 +950,15 @@ class Version
     __hostdev__ uint32_t id() const { return mData; }
     __hostdev__ uint32_t getMajor() const { return (mData >> 21) & ((1u << 11) - 1); }
     __hostdev__ uint32_t getMinor() const { return (mData >> 10) & ((1u << 11) - 1); }
-    __hostdev__ uint32_t getPatch() const { return mData & ((1u << 10) - 1); }
-    __hostdev__ bool isCompatible() const { return this->getMajor() == uint32_t(NANOVDB_MAJOR_VERSION_NUMBER);}
-    /// @brief Check the major version of this instance relative to NANOVDB_MAJOR_VERSION_NUMBER
-    /// @return return 0 if the major version equals NANOVDB_MAJOR_VERSION_NUMBER, else a negative age if it is
-    ///         older, i.e. smaller, and a positive age if it's newer, i.e.e larger.
+    __hostdev__ uint32_t getPatch() const { return  mData        & ((1u << 10) - 1); }
+    __hostdev__ bool isCompatible() const { return this->getMajor() == uint32_t(NANOVDB_MAJOR_VERSION_NUMBER); }
+    /// @brief Returns the difference between major version of this instance and NANOVDB_MAJOR_VERSION_NUMBER
+    /// @return return 0 if the major version equals NANOVDB_MAJOR_VERSION_NUMBER, else a negative age if this
+    ///         instance has a smaller major verion (is older), and a positive age if it is newer, i.e. larger.
     __hostdev__ int age() const {return int(this->getMajor()) - int(NANOVDB_MAJOR_VERSION_NUMBER);}
 
 #ifndef __CUDACC_RTC__
+    /// @brief returns a c-string of the semantic version, i.e. major.minor.patch
     const char* c_str() const
     {
         char* buffer = (char*)malloc(4 + 1 + 4 + 1 + 4 + 1); // xxxx.xxxx.xxxx\0
@@ -990,7 +971,7 @@ class Version
 // ----------------------------> Various math functions <-------------------------------------
 
 //@{
-/// @brief  Pi constant taken from Boost to match old behaviour
+/// @brief Pi constant taken from Boost to match old behaviour
 template<typename T>
 inline __hostdev__ constexpr T pi()
 {
@@ -3560,13 +3541,18 @@ struct NANOVDB_ALIGN(NANOVDB_DATA_ALIGNMENT) GridData
         mGridType = gridType;
         mBlindMetadataOffset = mGridSize; // i.e. no blind data
         mBlindMetadataCount = 0u; // i.e. no blind data
-        mData0 = 0u;
+        mData0 = 0u; // zero padding
         mData1 = 0u; // only used for index and point grids
-        mData2 = 0u;
+        mData2 = NANOVDB_MAGIC_GRID; // since version 32.6.0 (might be removed in the future)
     }
     /// @brief return true if the magic number and the version are both valid
     __hostdev__ bool isValid() const {
-        return mMagic == NANOVDB_MAGIC_GRID || (mMagic == NANOVDB_MAGIC_NUMBER && mVersion.isCompatible());
+        if (mMagic == NANOVDB_MAGIC_GRID || mData2 == NANOVDB_MAGIC_GRID) return true;
+        bool test = mMagic == NANOVDB_MAGIC_NUMBER;// could be GridData or io::FileHeader
+        if (test) test = mVersion.isCompatible();
+        if (test) test = mGridCount > 0u && mGridIndex < mGridCount;
+        if (test) test = mGridClass < GridClass::End && mGridType < GridType::End;
+        return test;
     }
     // Set and unset various bit flags
     __hostdev__ void setMinMaxOn(bool on = true) { mFlags.setMask(GridFlags::HasMinMax, on); }
@@ -7980,20 +7966,20 @@ VecT<GridHandleT> readUncompressedGrids(StreamT& is, const typename GridHandleT:
 {
     VecT<GridHandleT> handles;
     GridData data;
-    is.read((char*)&data, 40);// we only need to load the first 40 bytes
-    if (data.mMagic == NANOVDB_MAGIC_GRID || data.isValid()) {// stream contains a raw grid buffer
+    is.read((char*)&data, sizeof(GridData));
+    if (data.isValid()) {// stream contains a raw grid buffer
         uint64_t size = data.mGridSize, sum = 0u;
         while(data.mGridIndex + 1u < data.mGridCount) {
-            is.skip(data.mGridSize - 40);// skip grid
-            is.read((char*)&data, 40);// read 40 bytes
+            is.skip(data.mGridSize - sizeof(GridData));// skip grid
+            is.read((char*)&data, sizeof(GridData));// read sizeof(GridData) bytes
             sum += data.mGridSize;
         }
-        is.skip(-int64_t(sum + 40));// rewind to start
+        is.skip(-int64_t(sum + sizeof(GridData)));// rewind to start
         auto buffer = GridHandleT::BufferType::create(size + sum, &pool);
         is.read((char*)(buffer.data()), buffer.size());
         handles.emplace_back(std::move(buffer));
     } else {// Header0, MetaData0, gridName0, Grid0...HeaderN, MetaDataN, gridNameN, GridN
-        is.skip(-40);// rewind
+        is.skip(-sizeof(GridData));// rewind
         FileHeader head;
         while(is.read((char*)&head, sizeof(FileHeader))) {
             if (!head.isValid()) {
diff --git a/nanovdb/nanovdb/examples/ex_voxels_to_grid_cuda/ex_voxels_to_grid_cuda.cu b/nanovdb/nanovdb/examples/ex_voxels_to_grid_cuda/ex_voxels_to_grid_cuda.cu
@@ -17,7 +17,7 @@ int main()
 
         // Generate a NanoVDB grid that contains the list of voxels on the device
         auto handle = cudaVoxelsToGrid<float>(d_coords, numVoxels);
-        auto *grid = handle.deviceGrid<float>();
+        auto *d_grid = handle.deviceGrid<float>();
 
         // Define a list of values and copy them to the device
         float values[numVoxels] = {1.4f, 6.7f, -5.0f}, *d_values;
@@ -29,13 +29,13 @@ int main()
         cudaLambdaKernel<<<numBlocks, numThreads>>>(numVoxels, [=] __device__(size_t tid) {
             using OpT = SetVoxel<float>;// defines type of random-access operation (set value)
             const Coord &ijk = d_coords[tid];
-            grid->tree().set<OpT>(ijk, d_values[tid]);// normally one should use a ValueAccessor
-            printf("GPU: voxel # %lu, grid(%4i,%4i,%4i) = %5.1f\n", tid, ijk[0], ijk[1], ijk[2], grid->tree().getValue(ijk));
+            d_grid->tree().set<OpT>(ijk, d_values[tid]);// normally one should use a ValueAccessor
+            printf("GPU: voxel # %lu, grid(%4i,%4i,%4i) = %5.1f\n", tid, ijk[0], ijk[1], ijk[2], d_grid->tree().getValue(ijk));
         }); cudaCheckError();
 
         // Copy grid from GPU to CPU and print the voxel values for validation
         handle.deviceDownload();// creates a copy on the CPU
-        grid = handle.grid<float>();
+        auto *grid = handle.grid<float>();
         for (size_t i=0; i<numVoxels; ++i) {
             const Coord &ijk = coords[i];
             printf("CPU: voxel # %lu, grid(%4i,%4i,%4i) = %5.1f\n", i, ijk[0], ijk[1], ijk[2], grid->tree().getValue(ijk));
diff --git a/nanovdb/nanovdb/util/GridHandle.h b/nanovdb/nanovdb/util/GridHandle.h
@@ -363,20 +363,20 @@ template<typename BufferT>
 void GridHandle<BufferT>::read(std::istream& is, const BufferT& pool)
 {
     GridData data;
-    is.read((char*)&data, 40);// only 40 bytes are required for all the data we need in GridData
+    is.read((char*)&data, sizeof(GridData));
     if (data.isValid()) {
         uint64_t size = data.mGridSize, sum = 0u;
         while(data.mGridIndex + 1u < data.mGridCount) {// loop over remaining raw grids in stream
-            is.seekg(data.mGridSize - 40, std::ios::cur);// skip grid
-            is.read((char*)&data, 40);// read 40 bytes of the next GridData
+            is.seekg(data.mGridSize - sizeof(GridData), std::ios::cur);// skip grid
+            is.read((char*)&data, sizeof(GridData));
             sum += data.mGridSize;
         }
-        is.seekg(-int64_t(sum + 40), std::ios::cur);// rewind to start
         auto buffer = BufferT::create(size + sum, &pool);
+        is.seekg(-int64_t(sum + sizeof(GridData)), std::ios::cur);// rewind to start
         is.read((char*)(buffer.data()), buffer.size());
         *this = GridHandle(std::move(buffer));
     } else {
-        is.seekg(-40, std::ios::cur);// rewind
+        is.seekg(-sizeof(GridData), std::ios::cur);// rewind
         throw std::logic_error("This stream does not contain a valid raw grid buffer");
     }
 }// void GridHandle<BufferT>::read(std::istream& is, const BufferT& pool)
@@ -385,20 +385,20 @@ template<typename BufferT>
 void GridHandle<BufferT>::read(std::istream& is, uint32_t n, const BufferT& pool)
 {
     GridData data;
-    is.read((char*)&data, 40);// only 40 bytes are required for all the data we need in GridData
+    is.read((char*)&data, sizeof(GridData));
     if (data.isValid()) {
         if (n>=data.mGridCount) throw std::runtime_error("stream does not contain a #" + std::to_string(n) + " grid");
         while(data.mGridIndex != n) {
-            is.seekg(data.mGridSize - 40, std::ios::cur);// skip grid
-            is.read((char*)&data, 40);// read 40 bytes
+            is.seekg(data.mGridSize - sizeof(GridData), std::ios::cur);// skip grid
+            is.read((char*)&data, sizeof(GridData));
         }
         auto buffer = BufferT::create(data.mGridSize, &pool);
-        is.seekg(-40, std::ios::cur);// rewind
+        is.seekg(-sizeof(GridData), std::ios::cur);// rewind
         is.read((char*)(buffer.data()), data.mGridSize);
         updateGridCount((GridData*)buffer.data(), 0u, 1u);
         *this = GridHandle(std::move(buffer));
     } else {
-        is.seekg(-40, std::ios::cur);// rewind 40 bytes to undo initial read
+        is.seekg(-sizeof(GridData), std::ios::cur);// rewind sizeof(GridData) bytes to undo initial read
         throw std::logic_error("This file does not contain a valid raw buffer");
     }
 }// void GridHandle<BufferT>::read(std::istream& is, uint32_t n, const BufferT& pool)
@@ -414,7 +414,7 @@ void GridHandle<BufferT>::read(std::istream& is, const std::string &gridName, co
         uint32_t n = 0;
         while(data.mGridName != gridName && n++ < data.mGridCount) {
             is.seekg(data.mGridSize, std::ios::cur);// skip grid
-            is.read((char*)&data, byteSize);// read 40 bytes
+            is.read((char*)&data, byteSize);// read sizeof(GridData) bytes
             is.seekg(-byteSize, std::ios::cur);// rewind
         }
         if (n>data.mGridCount) throw std::runtime_error("No raw grid named \""+gridName+"\"");
diff --git a/nanovdb/nanovdb/util/cuda/CudaPointsToGrid.cuh b/nanovdb/nanovdb/util/cuda/CudaPointsToGrid.cuh
@@ -552,28 +552,23 @@ void CudaPointsToGrid<BuildT, AllocT>::countNodes(const PtrT points, size_t poin
     if (mVerbose==2) mTimer.restart("Generate tile keys");
     cudaLambdaKernel<<<numBlocks(pointCount), mNumThreads, 0, mStream>>>(pointCount, [=] __device__(size_t tid, const Data *d_data, const PtrT points) {
         auto coordToKey = [](const Coord &ijk)->uint64_t{
-            //  int32_t has a range of -2^31 to 2^31 - 1
-            // uint32_t has a range of 0 to 2^32 - 1
+            // Note: int32_t has a range of -2^31 to 2^31 - 1 whereas uint32_t has a range of 0 to 2^32 - 1
             static constexpr int64_t offset = 1 << 31;
             return (uint64_t(uint32_t(int64_t(ijk[2]) + offset) >> 12)      ) | // z is the lower 21 bits
                    (uint64_t(uint32_t(int64_t(ijk[1]) + offset) >> 12) << 21) | // y is the middle 21 bits
                    (uint64_t(uint32_t(int64_t(ijk[0]) + offset) >> 12) << 42); //  x is the upper 21 bits
-        };
+        };// coordToKey lambda functor
         d_indx[tid] = uint32_t(tid);
         uint64_t &key = d_keys[tid];
         if constexpr(is_same<BuildT, Point>::value) {// points are in world space
             if constexpr(is_same<Vec3T, Vec3f>::value) {
                 key = coordToKey(d_data->map.applyInverseMapF(points[tid]).round());
-                //key = NanoRoot<Point>::CoordToKey(d_data->map.applyInverseMapF(points[tid]).round());
             } else {// points are Vec3d
-                //key = NanoRoot<Point>::CoordToKey(d_data->map.applyInverseMap(points[tid]).round());
                 key = coordToKey(d_data->map.applyInverseMap(points[tid]).round());
             }
         } else if constexpr(is_same<Vec3T, Coord>::value) {// points Coord are in index space
-            //key = NanoRoot<BuildT>::CoordToKey(points[tid]);
             key = coordToKey(points[tid]);
         } else {// points are Vec3f or Vec3d in index space
-            //key = NanoRoot<BuildT>::CoordToKey(points[tid].round());
             key = coordToKey(points[tid].round());
         }
     }, mDeviceData, points);
@@ -605,7 +600,7 @@ void CudaPointsToGrid<BuildT, AllocT>::countNodes(const PtrT points, size_t poin
                     uint64_t(NanoUpper<BuildT>::CoordToOffset(ijk)) << 21 | // lower offset: 32^3 = 2^15,   i.e. next 15 bits
                     uint64_t(NanoLower<BuildT>::CoordToOffset(ijk)) <<  9 | // leaf  offset: 16^3 = 2^12,   i.e. next 12 bits
                     uint64_t(NanoLeaf< BuildT>::CoordToOffset(ijk));        // voxel offset:  8^3 =  2^9,   i.e. first 9 bits
-            };
+            };// voxelKey lambda functor
             tid += offset;
             Vec3T p = points[d_indx[tid]];
             if constexpr(is_same<BuildT, Point>::value) p = is_same<Vec3T, Vec3f>::value ? d_data->map.applyInverseMapF(p) : d_data->map.applyInverseMap(p);
