Add clang build fixes

* Add clang build fixes

* Add/fix comments

Signed-off-by: Matthew Cong <[email protected]>

Author: matthewdcong

nanovdb/nanovdb/cuda/DeviceBuffer.h

@@ -41,7 +41,7 @@ class DeviceBuffer
 
     /// @brief Initialize buffer
     /// @param size byte size of buffer to be initialized
-    /// @param host If true buffer is initialized only on the host/CPU, else on the device/GPU
+    /// @param device id of the device on which to initialize the buffer
     /// @note All existing buffers are first cleared
     /// @warning size is expected to be non-zero. Use clear() clear buffer!
     void init(uint64_t size, int device, cudaStream_t stream);
@@ -152,6 +152,11 @@ class DeviceBuffer
     /// @return An instance of this class using move semantics
     static DeviceBuffer create(uint64_t size, const DeviceBuffer* dummy, bool host, void* stream){return DeviceBuffer(size, host, stream);}
 
+    /// @brief Static factory method that returns an instance of this buffer
+    /// @param size byte size of buffer to be initialized
+    /// @param dummy this argument is currently ignored but required to match the API of the HostBuffer
+    /// @param device id of the device on which to initialize the buffer
+    /// @param stream cuda stream
     static DeviceBuffer create(uint64_t size, const DeviceBuffer* dummy = nullptr, int device = cudaCpuDeviceId, cudaStream_t stream = 0){return DeviceBuffer(size, device, stream);}
 
     /// @brief Static factory method that returns an instance of this buffer that wraps externally managed memory
@@ -167,7 +172,7 @@ class DeviceBuffer
     /// @param list list of device IDs and device memory pointers
     static DeviceBuffer create(uint64_t size, void* cpuData, std::initializer_list<std::pair<int,void*>> list) {return DeviceBuffer(size, cpuData, list);}
 
-    /// @brief  Static factory method that returns an instance of this buffer constructed from a HostBuffer
+    /// @brief Static factory method that returns an instance of this buffer constructed from a HostBuffer
     /// @param buffer host buffer from which to copy data
     /// @param device id of the device on which to initialize the buffer
     /// @param stream cuda stream

nanovdb/nanovdb/cuda/GridHandle.cuh

@@ -34,7 +34,8 @@ __global__ void cpyGridHandleMeta(const GridData *d_data, GridHandleMetaData *d_
 __global__ void updateGridCount(GridData *d_data, uint32_t gridIndex, uint32_t gridCount, bool *d_dirty)
 {
     NANOVDB_ASSERT(gridIndex < gridCount);
-    if (*d_dirty = d_data->mGridIndex != gridIndex || d_data->mGridCount != gridCount) {
+    *d_dirty = (d_data->mGridIndex != gridIndex) || (d_data->mGridCount != gridCount);
+    if (*d_dirty) {
         d_data->mGridIndex = gridIndex;
         d_data->mGridCount = gridCount;
         if (d_data->mChecksum.isEmpty()) *d_dirty = false;// no need to update checksum if it didn't already exist

nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh

@@ -24,7 +24,6 @@
 #include <nanovdb/NanoVDB.h>
 #include <nanovdb/cuda/DeviceBuffer.h>
 #include <nanovdb/cuda/TempPool.h>
-#include <nanovdb/cuda/UnifiedBuffer.h>
 #include <nanovdb/GridHandle.h>
 #include <nanovdb/tools/cuda/GridChecksum.cuh>
 #include <nanovdb/util/cuda/Timer.h>
@@ -789,9 +788,9 @@ struct BuildGridTreeRootFunctor
         auto &root = d_data->getRoot();
         root.mBBox = CoordBBox(); // init to empty
         root.mTableSize = d_data->nodeCount[2];
-        root.mBackground = NanoRoot<BuildT>::ValueType(0);// background_value
-        root.mMinimum = root.mMaximum = NanoRoot<BuildT>::ValueType(0);
-        root.mAverage = root.mStdDevi = NanoRoot<BuildT>::FloatType(0);
+        root.mBackground = typename NanoRoot<BuildT>::ValueType(0);// background_value
+        root.mMinimum = root.mMaximum = typename NanoRoot<BuildT>::ValueType(0);
+        root.mAverage = root.mStdDevi = typename NanoRoot<BuildT>::FloatType(0);
 
         // process Tree
         auto &tree = d_data->getTree();
@@ -940,15 +939,15 @@ struct BuildUpperNodesFunctor
         };
         const Coord ijk = keyToCoord(d_data->d_tile_keys[tid]);
 #else
-        const Coord ijk = NanoRoot<uint32_t>::KeyToCoord(d_data->d_tile_keys[tid]);
+        const Coord ijk = typename NanoRoot<uint32_t>::KeyToCoord(d_data->d_tile_keys[tid]);
 #endif
         root.tile(tid)->setChild(ijk, &upper, &root);
         upper.mBBox[0] = ijk;
         upper.mFlags = 0;
         upper.mValueMask.setOff();
         upper.mChildMask.setOff();
-        upper.mMinimum = upper.mMaximum = NanoLower<BuildT>::ValueType(0);
-        upper.mAverage = upper.mStdDevi = NanoLower<BuildT>::FloatType(0);
+        upper.mMinimum = upper.mMaximum = typename NanoLower<BuildT>::ValueType(0);
+        upper.mAverage = upper.mStdDevi = typename NanoLower<BuildT>::FloatType(0);
     }
 };
 
@@ -958,7 +957,7 @@ struct SetUpperBackgroundValuesFunctor
     __device__
     void operator()(size_t tid, PointsToGridData<BuildT> *d_data) {
         auto &upper = d_data->getUpper(tid >> 15);
-        upper.mTable[tid & 32767u].value = NanoUpper<BuildT>::ValueType(0);// background
+        upper.mTable[tid & 32767u].value = typename NanoUpper<BuildT>::ValueType(0);// background
     }
 };
 
@@ -993,8 +992,8 @@ struct BuildLowerNodesFunctor
         lower.mFlags = 0;
         lower.mValueMask.setOff();
         lower.mChildMask.setOff();
-        lower.mMinimum = lower.mMaximum = NanoLower<BuildT>::ValueType(0);// background;
-        lower.mAverage = lower.mStdDevi = NanoLower<BuildT>::FloatType(0);
+        lower.mMinimum = lower.mMaximum = typename NanoLower<BuildT>::ValueType(0);// background;
+        lower.mAverage = lower.mStdDevi = typename NanoLower<BuildT>::FloatType(0);
     }
 };
 
@@ -1004,7 +1003,7 @@ struct SetLowerBackgroundValuesFunctor
     __device__
     void operator()(size_t tid, PointsToGridData<BuildT> *d_data) {
         auto &lower = d_data->getLower(tid >> 12);
-        lower.mTable[tid & 4095u].value = NanoLower<BuildT>::ValueType(0);// background
+        lower.mTable[tid & 4095u].value = typename NanoLower<BuildT>::ValueType(0);// background
     }
 };
 
@@ -1045,8 +1044,8 @@ struct ProcessLeafMetaDataFunctor
             leaf.mOffset = tid*512u + 1u;// background is index 0
             leaf.mPrefixSum = 0u;
         } else if constexpr(!BuildTraits<BuildT>::is_special) {
-            leaf.mAverage = leaf.mStdDevi = NanoLeaf<BuildT>::FloatType(0);
-            leaf.mMinimum = leaf.mMaximum = NanoLeaf<BuildT>::ValueType(0);
+            leaf.mAverage = leaf.mStdDevi = typename NanoLeaf<BuildT>::FloatType(0);
+            leaf.mMinimum = leaf.mMaximum = typename NanoLeaf<BuildT>::ValueType(0);
         }
     }
 };
@@ -1066,7 +1065,7 @@ struct SetLeafActiveVoxelStateAndValuesFunctor
         if constexpr(util::is_same<Point, BuildT>::value) {
             leaf.mValues[n] = uint16_t(pointID + d_data->pointsPerVoxel[tid] - leaf.offset());
         } else if constexpr(!BuildTraits<BuildT>::is_special) {
-            leaf.mValues[n] = NanoLeaf<BuildT>::ValueType(1);// set value of active voxels that are not points (or index)
+            leaf.mValues[n] = typename NanoLeaf<BuildT>::ValueType(1);// set value of active voxels that are not points (or index)
         }
     }
 };
@@ -1083,7 +1082,7 @@ struct SetLeafInactiveVoxelValuesFunctor
             const uint32_t m = leaf.mValueMask.findPrev<true>(n - 1);
             leaf.mValues[n] = m < 512u ? leaf.mValues[m] : 0u;
         } else if constexpr(!BuildTraits<BuildT>::is_special) {
-            leaf.mValues[n] = NanoLeaf<BuildT>::ValueType(0);// value of inactive voxels
+            leaf.mValues[n] = typename NanoLeaf<BuildT>::ValueType(0);// value of inactive voxels
         }
     }
 };