[Bf-blender-cvs] [bd2e3bb7bd0] blender-v3.0-release: Fix T93045: Cycles HIP not rendering OpenVDB volumes
Brecht Van Lommel
noreply at git.blender.org
Thu Nov 18 13:39:07 CET 2021
Commit: bd2e3bb7bd06bcbb2134e4853a72ab28f5f332b9
Author: Brecht Van Lommel
Date: Thu Nov 18 00:41:04 2021 +0100
Branches: blender-v3.0-release
https://developer.blender.org/rBbd2e3bb7bd06bcbb2134e4853a72ab28f5f332b9
Fix T93045: Cycles HIP not rendering OpenVDB volumes
Build HIP kernels with NanoVDB, and patch NanoVDB to work with HIP.
This is a header only library so no rebuild is needed. The changes are being
submitted upstream to openvdb, so this patch should be temporary.
Thanks Thomas for help testing this.
===================================================================
M build_files/build_environment/cmake/nanovdb.cmake
A build_files/build_environment/patches/nanovdb.diff
M intern/cycles/kernel/CMakeLists.txt
===================================================================
diff --git a/build_files/build_environment/cmake/nanovdb.cmake b/build_files/build_environment/cmake/nanovdb.cmake
index 0baaf80c254..66bbb9f10d7 100644
--- a/build_files/build_environment/cmake/nanovdb.cmake
+++ b/build_files/build_environment/cmake/nanovdb.cmake
@@ -42,6 +42,7 @@ ExternalProject_Add(nanovdb
URL_HASH ${NANOVDB_HASH_TYPE}=${NANOVDB_HASH}
PREFIX ${BUILD_DIR}/nanovdb
SOURCE_SUBDIR nanovdb
+ PATCH_COMMAND ${PATCH_CMD} -p 1 -d ${BUILD_DIR}/nanovdb/src/nanovdb < ${PATCH_DIR}/nanovdb.diff
CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/nanovdb ${DEFAULT_CMAKE_FLAGS} ${NANOVDB_EXTRA_ARGS}
INSTALL_DIR ${LIBDIR}/nanovdb
)
diff --git a/build_files/build_environment/patches/nanovdb.diff b/build_files/build_environment/patches/nanovdb.diff
new file mode 100644
index 00000000000..fd833e61336
--- /dev/null
+++ b/build_files/build_environment/patches/nanovdb.diff
@@ -0,0 +1,374 @@
+Index: nanovdb/nanovdb/NanoVDB.h
+===================================================================
+--- a/nanovdb/nanovdb/NanoVDB.h (revision 62751)
++++ b/nanovdb/nanovdb/NanoVDB.h (working copy)
+@@ -152,8 +152,8 @@
+
+ #endif // __CUDACC_RTC__
+
+-#ifdef __CUDACC__
+-// Only define __hostdev__ when using NVIDIA CUDA compiler
++#if defined(__CUDACC__) || defined(__HIP__)
++// Only define __hostdev__ when using NVIDIA CUDA or HIP compiler
+ #define __hostdev__ __host__ __device__
+ #else
+ #define __hostdev__
+@@ -461,7 +461,7 @@
+ /// Maximum floating-point values
+ template<typename T>
+ struct Maximum;
+-#ifdef __CUDA_ARCH__
++#if defined(__CUDA_ARCH__) || defined(__HIP__)
+ template<>
+ struct Maximum<int>
+ {
+@@ -1006,10 +1006,10 @@
+ using Vec3i = Vec3<int>;
+
+ /// @brief Return a single precision floating-point vector of this coordinate
+-Vec3f Coord::asVec3s() const { return Vec3f(float(mVec[0]), float(mVec[1]), float(mVec[2])); }
++inline __hostdev__ Vec3f Coord::asVec3s() const { return Vec3f(float(mVec[0]), float(mVec[1]), float(mVec[2])); }
+
+ /// @brief Return a double precision floating-point vector of this coordinate
+-Vec3d Coord::asVec3d() const { return Vec3d(double(mVec[0]), double(mVec[1]), double(mVec[2])); }
++inline __hostdev__ Vec3d Coord::asVec3d() const { return Vec3d(double(mVec[0]), double(mVec[1]), double(mVec[2])); }
+
+ // ----------------------------> Vec4 <--------------------------------------
+
+@@ -1820,7 +1820,7 @@
+ }; // Map
+
+ template<typename Mat4T>
+-void Map::set(const Mat4T& mat, const Mat4T& invMat, double taper)
++__hostdev__ void Map::set(const Mat4T& mat, const Mat4T& invMat, double taper)
+ {
+ float * mf = mMatF, *vf = mVecF;
+ float* mif = mInvMatF;
+@@ -2170,7 +2170,7 @@
+ }; // Class Grid
+
+ template<typename TreeT>
+-int Grid<TreeT>::findBlindDataForSemantic(GridBlindDataSemantic semantic) const
++__hostdev__ int Grid<TreeT>::findBlindDataForSemantic(GridBlindDataSemantic semantic) const
+ {
+ for (uint32_t i = 0, n = blindDataCount(); i < n; ++i)
+ if (blindMetaData(i).mSemantic == semantic)
+@@ -2328,7 +2328,7 @@
+ }; // Tree class
+
+ template<typename RootT>
+-void Tree<RootT>::extrema(ValueType& min, ValueType& max) const
++__hostdev__ void Tree<RootT>::extrema(ValueType& min, ValueType& max) const
+ {
+ min = this->root().valueMin();
+ max = this->root().valueMax();
+@@ -2336,7 +2336,7 @@
+
+ template<typename RootT>
+ template<typename NodeT>
+-const NodeT* Tree<RootT>::getNode(uint32_t i) const
++__hostdev__ const NodeT* Tree<RootT>::getNode(uint32_t i) const
+ {
+ static_assert(is_same<TreeNodeT<NodeT::LEVEL>, NodeT>::value, "Tree::getNode: unvalid node type");
+ NANOVDB_ASSERT(i < DataType::mCount[NodeT::LEVEL]);
+@@ -2345,7 +2345,7 @@
+
+ template<typename RootT>
+ template<int LEVEL>
+-const typename TreeNode<Tree<RootT>, LEVEL>::type* Tree<RootT>::getNode(uint32_t i) const
++__hostdev__ const typename TreeNode<Tree<RootT>, LEVEL>::type* Tree<RootT>::getNode(uint32_t i) const
+ {
+ NANOVDB_ASSERT(i < DataType::mCount[LEVEL]);
+ return reinterpret_cast<const TreeNodeT<LEVEL>*>(reinterpret_cast<const uint8_t*>(this) + DataType::mBytes[LEVEL]) + i;
+@@ -2353,7 +2353,7 @@
+
+ template<typename RootT>
+ template<typename NodeT>
+-NodeT* Tree<RootT>::getNode(uint32_t i)
++__hostdev__ NodeT* Tree<RootT>::getNode(uint32_t i)
+ {
+ static_assert(is_same<TreeNodeT<NodeT::LEVEL>, NodeT>::value, "Tree::getNode: invalid node type");
+ NANOVDB_ASSERT(i < DataType::mCount[NodeT::LEVEL]);
+@@ -2362,7 +2362,7 @@
+
+ template<typename RootT>
+ template<int LEVEL>
+-typename TreeNode<Tree<RootT>, LEVEL>::type* Tree<RootT>::getNode(uint32_t i)
++__hostdev__ typename TreeNode<Tree<RootT>, LEVEL>::type* Tree<RootT>::getNode(uint32_t i)
+ {
+ NANOVDB_ASSERT(i < DataType::mCount[LEVEL]);
+ return reinterpret_cast<TreeNodeT<LEVEL>*>(reinterpret_cast<uint8_t*>(this) + DataType::mBytes[LEVEL]) + i;
+@@ -2370,7 +2370,7 @@
+
+ template<typename RootT>
+ template<typename NodeT>
+-uint32_t Tree<RootT>::getNodeID(const NodeT& node) const
++__hostdev__ uint32_t Tree<RootT>::getNodeID(const NodeT& node) const
+ {
+ static_assert(is_same<TreeNodeT<NodeT::LEVEL>, NodeT>::value, "Tree::getNodeID: invalid node type");
+ const NodeT* first = reinterpret_cast<const NodeT*>(reinterpret_cast<const uint8_t*>(this) + DataType::mBytes[NodeT::LEVEL]);
+@@ -2380,7 +2380,7 @@
+
+ template<typename RootT>
+ template<typename NodeT>
+-uint32_t Tree<RootT>::getLinearOffset(const NodeT& node) const
++__hostdev__ uint32_t Tree<RootT>::getLinearOffset(const NodeT& node) const
+ {
+ return this->getNodeID(node) + DataType::mPFSum[NodeT::LEVEL];
+ }
+@@ -3366,7 +3366,7 @@
+ }; // LeafNode class
+
+ template<typename ValueT, typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
+-inline void LeafNode<ValueT, CoordT, MaskT, LOG2DIM>::updateBBox()
++inline __hostdev__ void LeafNode<ValueT, CoordT, MaskT, LOG2DIM>::updateBBox()
+ {
+ static_assert(LOG2DIM == 3, "LeafNode::updateBBox: only supports LOGDIM = 3!");
+ if (!this->isActive()) return;
+Index: nanovdb/nanovdb/util/SampleFromVoxels.h
+===================================================================
+--- a/nanovdb/nanovdb/util/SampleFromVoxels.h (revision 62751)
++++ b/nanovdb/nanovdb/util/SampleFromVoxels.h (working copy)
+@@ -22,7 +22,7 @@
+ #define NANOVDB_SAMPLE_FROM_VOXELS_H_HAS_BEEN_INCLUDED
+
+ // Only define __hostdev__ when compiling as NVIDIA CUDA
+-#ifdef __CUDACC__
++#if defined(__CUDACC__) || defined(__HIP__)
+ #define __hostdev__ __host__ __device__
+ #else
+ #include <cmath> // for floor
+@@ -136,7 +136,7 @@
+
+ template<typename TreeOrAccT>
+ template<typename Vec3T>
+-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, true>::operator()(const Vec3T& xyz) const
++__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, true>::operator()(const Vec3T& xyz) const
+ {
+ const CoordT ijk = Round<CoordT>(xyz);
+ if (ijk != mPos) {
+@@ -147,7 +147,7 @@
+ }
+
+ template<typename TreeOrAccT>
+-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, true>::operator()(const CoordT& ijk) const
++__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, true>::operator()(const CoordT& ijk) const
+ {
+ if (ijk != mPos) {
+ mPos = ijk;
+@@ -158,7 +158,7 @@
+
+ template<typename TreeOrAccT>
+ template<typename Vec3T>
+-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, false>::operator()(const Vec3T& xyz) const
++__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, false>::operator()(const Vec3T& xyz) const
+ {
+ return mAcc.getValue(Round<CoordT>(xyz));
+ }
+@@ -195,7 +195,7 @@
+ }; // TrilinearSamplerBase
+
+ template<typename TreeOrAccT>
+-void TrilinearSampler<TreeOrAccT>::stencil(CoordT& ijk, ValueT (&v)[2][2][2]) const
++__hostdev__ void TrilinearSampler<TreeOrAccT>::stencil(CoordT& ijk, ValueT (&v)[2][2][2]) const
+ {
+ v[0][0][0] = mAcc.getValue(ijk); // i, j, k
+
+@@ -224,7 +224,7 @@
+
+ template<typename TreeOrAccT>
+ template<typename RealT, template<typename...> class Vec3T>
+-typename TreeOrAccT::ValueType TrilinearSampler<TreeOrAccT>::sample(const Vec3T<RealT> &uvw, const ValueT (&v)[2][2][2])
++__hostdev__ typename TreeOrAccT::ValueType TrilinearSampler<TreeOrAccT>::sample(const Vec3T<RealT> &uvw, const ValueT (&v)[2][2][2])
+ {
+ #if 0
+ auto lerp = [](ValueT a, ValueT b, ValueT w){ return fma(w, b-a, a); };// = w*(b-a) + a
+@@ -239,7 +239,7 @@
+
+ template<typename TreeOrAccT>
+ template<typename RealT, template<typename...> class Vec3T>
+-Vec3T<typename TreeOrAccT::ValueType> TrilinearSampler<TreeOrAccT>::gradient(const Vec3T<RealT> &uvw, const ValueT (&v)[2][2][2])
++__hostdev__ Vec3T<typename TreeOrAccT::ValueType> TrilinearSampler<TreeOrAccT>::gradient(const Vec3T<RealT> &uvw, const ValueT (&v)[2][2][2])
+ {
+ static_assert(std::is_floating_point<ValueT>::value, "TrilinearSampler::gradient requires a floating-point type");
+ #if 0
+@@ -270,7 +270,7 @@
+ }
+
+ template<typename TreeOrAccT>
+-bool TrilinearSampler<TreeOrAccT>::zeroCrossing(const ValueT (&v)[2][2][2])
++__hostdev__ bool TrilinearSampler<TreeOrAccT>::zeroCrossing(const ValueT (&v)[2][2][2])
+ {
+ static_assert(std::is_floating_point<ValueT>::value, "TrilinearSampler::zeroCrossing requires a floating-point type");
+ const bool less = v[0][0][0] < ValueT(0);
+@@ -363,7 +363,7 @@
+
+ template<typename TreeOrAccT>
+ template<typename RealT, template<typename...> class Vec3T>
+-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, true>::operator()(Vec3T<RealT> xyz) const
++__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, true>::operator()(Vec3T<RealT> xyz) const
+ {
+ this->cache(xyz);
+ return BaseT::sample(xyz, mVal);
+@@ -370,7 +370,7 @@
+ }
+
+ template<typename TreeOrAccT>
+-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, true>::operator()(const CoordT &ijk) const
++__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, true>::operator()(const CoordT &ijk) const
+ {
+ return ijk == mPos ? mVal[0][0][0] : BaseT::mAcc.getValue(ijk);
+ }
+@@ -377,7 +377,7 @@
+
+ temp
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list