blender/build_files/build_environment/patches/openvdb.diff

diff -ur openvdb-9.0.0/cmake/FindTBB.cmake openvdb/cmake/FindTBB.cmake
--- openvdb-9.0.0/cmake/FindTBB.cmake	2021-10-30 03:55:40.000000000 +0100
+++ openvdb/cmake/FindTBB.cmake	2022-03-31 11:33:15.592329750 +0100
@@ -252,7 +252,8 @@
       set(_TBB_LIB_NAME "${_TBB_LIB_NAME}${TBB_DEBUG_SUFFIX}")
     endif()
 
-    find_library(Tbb_${COMPONENT}_LIBRARY_${BUILD_TYPE} ${_TBB_LIB_NAME}
+    find_library(Tbb_${COMPONENT}_LIBRARY_${BUILD_TYPE}
+      NAMES ${_TBB_LIB_NAME} ${_TBB_LIB_NAME}_static
       ${_FIND_TBB_ADDITIONAL_OPTIONS}
       PATHS ${_TBB_LIBRARYDIR_SEARCH_DIRS}
       PATH_SUFFIXES ${CMAKE_INSTALL_LIBDIR} lib64 lib
diff -Naur openvdb-8.0.0/openvdb/openvdb/CMakeLists.txt openvdb/openvdb/openvdb/CMakeLists.txt
--- openvdb-8.0.0/openvdb/openvdb/CMakeLists.txt	2020-12-24 10:13:14 -0700
+++ openvdb/openvdb/openvdb/CMakeLists.txt	2021-02-05 11:18:33 -0700
@@ -107,7 +107,9 @@
   #        http://boost.2283326.n4.nabble.com/CMake-config-scripts-broken-in-1-70-td4708957.html
   #        https://github.com/boostorg/boost_install/commit/160c7cb2b2c720e74463865ef0454d4c4cd9ae7c
   set(BUILD_SHARED_LIBS ON)
-  set(Boost_USE_STATIC_LIBS OFF)
+  if(NOT WIN32) # blender links boost statically on windows
+    set(Boost_USE_STATIC_LIBS OFF)
+  endif()
 endif()
 
 find_package(Boost ${MINIMUM_BOOST_VERSION} REQUIRED COMPONENTS iostreams system)
@@ -146,6 +148,7 @@
         Boost::disable_autolinking  # add -DBOOST_ALL_NO_LIB
       )
   endif()
+  add_definitions(-D__TBB_NO_IMPLICIT_LINKAGE -DOPENVDB_OPENEXR_STATICLIB)
 endif()
 
 if(USE_EXR)
@@ -379,7 +382,12 @@
 # imported targets.
 
 if(OPENVDB_CORE_SHARED)
-  add_library(openvdb_shared SHARED ${OPENVDB_LIBRARY_SOURCE_FILES})
+  if(WIN32)
+    configure_file(version.rc.in ${CMAKE_CURRENT_BINARY_DIR}/version.rc @ONLY)
+    add_library(openvdb_shared SHARED ${OPENVDB_LIBRARY_SOURCE_FILES} ${CMAKE_CURRENT_BINARY_DIR}/version.rc)
+  else()
+    add_library(openvdb_shared SHARED ${OPENVDB_LIBRARY_SOURCE_FILES})
+  endif()
 endif()
 
 if(OPENVDB_CORE_STATIC)
diff -Naur openvdb-8.0.0/openvdb/openvdb/version.rc.in openvdb/openvdb/openvdb/version.rc.in
--- openvdb-8.0.0/openvdb/openvdb/version.rc.in	1969-12-31 17:00:00 -0700
+++ openvdb/openvdb/openvdb/version.rc.in	2021-02-05 11:18:33 -0700
@@ -0,0 +1,48 @@
+#include <winver.h>
+
+#define VER_FILEVERSION             @OpenVDB_MAJOR_VERSION@,@OpenVDB_MINOR_VERSION@,@OpenVDB_PATCH_VERSION@,0
+#define VER_FILEVERSION_STR         "@OpenVDB_MAJOR_VERSION@.@OpenVDB_MINOR_VERSION@.@OpenVDB_PATCH_VERSION@.0\0"
+
+#define VER_PRODUCTVERSION          @OpenVDB_MAJOR_VERSION@,@OpenVDB_MINOR_VERSION@,@OpenVDB_PATCH_VERSION@,0
+#define VER_PRODUCTVERSION_STR      "@OpenVDB_MAJOR_VERSION@.@OpenVDB_MINOR_VERSION@\0"
+
+#ifndef DEBUG
+#define VER_DEBUG                   0
+#else
+#define VER_DEBUG                   VS_FF_DEBUG
+#endif
+
+VS_VERSION_INFO VERSIONINFO
+FILEVERSION     VER_FILEVERSION
+PRODUCTVERSION  VER_PRODUCTVERSION
+FILEFLAGSMASK   VS_FFI_FILEFLAGSMASK
+FILEFLAGS       (VER_DEBUG)
+FILEOS          VOS__WINDOWS32
+FILETYPE        VFT_DLL
+FILESUBTYPE     VFT2_UNKNOWN
+BEGIN
+    BLOCK "StringFileInfo"
+    BEGIN
+        BLOCK "040904E4"
+        BEGIN
+            VALUE "FileDescription",  "OpenVDB"
+            VALUE "FileVersion",      VER_FILEVERSION_STR
+            VALUE "InternalName",     "OpenVDB"
+            VALUE "ProductName",      "OpenVDB"
+            VALUE "ProductVersion",   VER_PRODUCTVERSION_STR
+        END
+    END
+
+    BLOCK "VarFileInfo"
+    BEGIN
+        /* The following line should only be modified for localized versions.     */
+        /* It consists of any number of WORD,WORD pairs, with each pair           */
+        /* describing a language,codepage combination supported by the file.      */
+        /*                                                                        */
+        /* For example, a file might have values "0x409,1252" indicating that it  */
+        /* supports English language (0x409) in the Windows ANSI codepage (1252). */
+
+        VALUE "Translation", 0x409, 1252
+
+    END
+END
diff --git a/nanovdb/nanovdb/NanoVDB.h b/nanovdb/nanovdb/NanoVDB.h
index cc2e54b77..703d2eabc 100644
--- a/nanovdb/nanovdb/NanoVDB.h
+++ b/nanovdb/nanovdb/NanoVDB.h
@@ -161,8 +161,8 @@ typedef unsigned long long uint64_t;
 
 #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__
@@ -611,7 +611,7 @@ struct Delta<double>
 /// Maximum floating-point values
 template<typename T>
 struct Maximum;
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HIP__)
 template<>
 struct Maximum<int>
 {
@@ -1176,10 +1176,10 @@ using Vec3f = Vec3<float>;
 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])); }
+__hostdev__ inline 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])); }
+__hostdev__ inline Vec3d Coord::asVec3d() const { return Vec3d(double(mVec[0]), double(mVec[1]), double(mVec[2])); }
 
 // ----------------------------> Vec4 <--------------------------------------
 
@@ -2042,7 +2042,7 @@ struct Map
 }; // 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;
@@ -2486,7 +2486,7 @@ class Grid : private GridData
 }; // 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 = this->blindDataCount(); i < n; ++i)
         if (this->blindMetaData(i).mSemantic == semantic)
@@ -2671,7 +2671,7 @@ class Tree : private TreeData<RootT::LEVEL>
 }; // 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().minimum();
     max = this->root().maximum();
@@ -3880,7 +3880,7 @@ class LeafNode : private LeafData<BuildT, CoordT, MaskT, Log2Dim>
 }; // LeafNode class
 
 template<typename ValueT, typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
-inline void LeafNode<ValueT, CoordT, MaskT, LOG2DIM>::updateBBox()
+__hostdev__ inline void LeafNode<ValueT, CoordT, MaskT, LOG2DIM>::updateBBox()
 {
     static_assert(LOG2DIM == 3, "LeafNode::updateBBox: only supports LOGDIM = 3!");
     if (!this->isActive()) return;
diff --git a/nanovdb/nanovdb/util/SampleFromVoxels.h b/nanovdb/nanovdb/util/SampleFromVoxels.h
index 852123dac..e779d66cf 100644
--- a/nanovdb/nanovdb/util/SampleFromVoxels.h
+++ b/nanovdb/nanovdb/util/SampleFromVoxels.h
@@ -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 @@ class SampleFromVoxels<TreeOrAccT, 0, false>
 
 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 @@ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, true>::operator()
 }
 
 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 @@ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, true>::operator()
 
 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 @@ class TrilinearSampler
 }; // 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 @@ void TrilinearSampler<TreeOrAccT>::stencil(CoordT& ijk, ValueT (&v)[2][2][2]) co
 
 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 @@ typename TreeOrAccT::ValueType TrilinearSampler<TreeOrAccT>::sample(const Vec3T<
 
 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(is_floating_point<ValueT>::value, "TrilinearSampler::gradient requires a floating-point type");
 #if 0
@@ -270,7 +270,7 @@ Vec3T<typename TreeOrAccT::ValueType> TrilinearSampler<TreeOrAccT>::gradient(con
 }
 
 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(is_floating_point<ValueT>::value, "TrilinearSampler::zeroCrossing requires a floating-point type");
     const bool less = v[0][0][0] < ValueT(0);
@@ -363,21 +363,21 @@ class SampleFromVoxels<TreeOrAccT, 1, true> : public TrilinearSampler<TreeOrAccT
 
 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);
 }
 
 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);
 }
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-Vec3T<typename TreeOrAccT::ValueType> SampleFromVoxels<TreeOrAccT, 1, true>::gradient(Vec3T<RealT> xyz) const
+__hostdev__ Vec3T<typename TreeOrAccT::ValueType> SampleFromVoxels<TreeOrAccT, 1, true>::gradient(Vec3T<RealT> xyz) const
 {
     this->cache(xyz);
     return BaseT::gradient(xyz, mVal);
@@ -393,7 +393,7 @@ __hostdev__ bool SampleFromVoxels<TreeOrAccT, 1, true>::zeroCrossing(Vec3T<RealT
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-void SampleFromVoxels<TreeOrAccT, 1, true>::cache(Vec3T<RealT>& xyz) const
+__hostdev__ void SampleFromVoxels<TreeOrAccT, 1, true>::cache(Vec3T<RealT>& xyz) const
 {
     CoordT ijk = Floor<CoordT>(xyz);
     if (ijk != mPos) {
@@ -406,7 +406,7 @@ void SampleFromVoxels<TreeOrAccT, 1, true>::cache(Vec3T<RealT>& xyz) const
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, false>::operator()(Vec3T<RealT> xyz) const
+__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, false>::operator()(Vec3T<RealT> xyz) const
 {
     ValueT val[2][2][2];
     CoordT ijk = Floor<CoordT>(xyz);
@@ -418,7 +418,7 @@ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, false>::operator(
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, false>::operator()(Vec3T<RealT> xyz) const
+__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, false>::operator()(Vec3T<RealT> xyz) const
 {
     auto lerp = [](ValueT a, ValueT b, RealT w) { return a + ValueT(w) * (b - a); };
 
@@ -463,7 +463,7 @@ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, false>::operator(
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-inline Vec3T<typename TreeOrAccT::ValueType> SampleFromVoxels<TreeOrAccT, 1, false>::gradient(Vec3T<RealT> xyz) const
+__hostdev__ inline Vec3T<typename TreeOrAccT::ValueType> SampleFromVoxels<TreeOrAccT, 1, false>::gradient(Vec3T<RealT> xyz) const
 {
     ValueT val[2][2][2];
     CoordT ijk = Floor<CoordT>(xyz);
@@ -473,7 +473,7 @@ inline Vec3T<typename TreeOrAccT::ValueType> SampleFromVoxels<TreeOrAccT, 1, fal
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-bool SampleFromVoxels<TreeOrAccT, 1, false>::zeroCrossing(Vec3T<RealT> xyz) const
+__hostdev__ bool SampleFromVoxels<TreeOrAccT, 1, false>::zeroCrossing(Vec3T<RealT> xyz) const
 {
     ValueT val[2][2][2];
     CoordT ijk = Floor<CoordT>(xyz);
@@ -510,7 +510,7 @@ class TriquadraticSampler
 }; // TriquadraticSamplerBase
 
 template<typename TreeOrAccT>
-void TriquadraticSampler<TreeOrAccT>::stencil(const CoordT &ijk, ValueT (&v)[3][3][3]) const
+__hostdev__ void TriquadraticSampler<TreeOrAccT>::stencil(const CoordT &ijk, ValueT (&v)[3][3][3]) const
 {
     CoordT p(ijk[0] - 1, 0, 0);
     for (int dx = 0; dx < 3; ++dx, ++p[0]) {
@@ -526,7 +526,7 @@ void TriquadraticSampler<TreeOrAccT>::stencil(const CoordT &ijk, ValueT (&v)[3][
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-typename TreeOrAccT::ValueType TriquadraticSampler<TreeOrAccT>::sample(const Vec3T<RealT> &uvw, const ValueT (&v)[3][3][3])
+__hostdev__ typename TreeOrAccT::ValueType TriquadraticSampler<TreeOrAccT>::sample(const Vec3T<RealT> &uvw, const ValueT (&v)[3][3][3])
 {
     auto kernel = [](const ValueT* value, double weight)->ValueT {
         return weight * (weight * (0.5f * (value[0] + value[2]) - value[1]) +
@@ -545,7 +545,7 @@ typename TreeOrAccT::ValueType TriquadraticSampler<TreeOrAccT>::sample(const Vec
 }
 
 template<typename TreeOrAccT>
-bool TriquadraticSampler<TreeOrAccT>::zeroCrossing(const ValueT (&v)[3][3][3])
+__hostdev__ bool TriquadraticSampler<TreeOrAccT>::zeroCrossing(const ValueT (&v)[3][3][3])
 {
     static_assert(is_floating_point<ValueT>::value, "TrilinearSampler::zeroCrossing requires a floating-point type");
     const bool less = v[0][0][0] < ValueT(0);
@@ -624,14 +624,14 @@ class SampleFromVoxels<TreeOrAccT, 2, true> : public TriquadraticSampler<TreeOrA
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 2, true>::operator()(Vec3T<RealT> xyz) const
+__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 2, true>::operator()(Vec3T<RealT> xyz) const
 {
     this->cache(xyz);
     return BaseT::sample(xyz, mVal);
 }
 
 template<typename TreeOrAccT>
-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 2, true>::operator()(const CoordT &ijk) const
+__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 2, true>::operator()(const CoordT &ijk) const
 {
     return  ijk == mPos ? mVal[1][1][1] : BaseT::mAcc.getValue(ijk);
 }
@@ -646,7 +646,7 @@ __hostdev__ bool SampleFromVoxels<TreeOrAccT, 2, true>::zeroCrossing(Vec3T<RealT
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-void SampleFromVoxels<TreeOrAccT, 2, true>::cache(Vec3T<RealT>& xyz) const
+__hostdev__ void SampleFromVoxels<TreeOrAccT, 2, true>::cache(Vec3T<RealT>& xyz) const
 {
     CoordT ijk = Floor<CoordT>(xyz);
     if (ijk != mPos) {
@@ -657,7 +657,7 @@ void SampleFromVoxels<TreeOrAccT, 2, true>::cache(Vec3T<RealT>& xyz) const
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 2, false>::operator()(Vec3T<RealT> xyz) const
+__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 2, false>::operator()(Vec3T<RealT> xyz) const
 {
     ValueT val[3][3][3];
     CoordT ijk = Floor<CoordT>(xyz);
@@ -667,7 +667,7 @@ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 2, false>::operator(
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-bool SampleFromVoxels<TreeOrAccT, 2, false>::zeroCrossing(Vec3T<RealT> xyz) const
+__hostdev__ bool SampleFromVoxels<TreeOrAccT, 2, false>::zeroCrossing(Vec3T<RealT> xyz) const
 {
     ValueT val[3][3][3];
     CoordT ijk = Floor<CoordT>(xyz);
@@ -710,7 +710,7 @@ class TricubicSampler
 }; // TricubicSampler
 
 template<typename TreeOrAccT>
-void TricubicSampler<TreeOrAccT>::stencil(const CoordT& ijk, ValueT (&C)[64]) const
+__hostdev__ void TricubicSampler<TreeOrAccT>::stencil(const CoordT& ijk, ValueT (&C)[64]) const
 {
     auto fetch = [&](int i, int j, int k) -> ValueT& { return C[((i + 1) << 4) + ((j + 1) << 2) + k + 1]; };
 
@@ -929,7 +929,7 @@ class SampleFromVoxels<TreeOrAccT, 3, true> : public TricubicSampler<TreeOrAccT>
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 3, true>::operator()(Vec3T<RealT> xyz) const
+__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 3, true>::operator()(Vec3T<RealT> xyz) const
 {
     this->cache(xyz);
     return BaseT::sample(xyz, mC);
@@ -937,7 +937,7 @@ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 3, true>::operator()
 
 template<typename TreeOrAccT>
 template<typename RealT, template<typename...> class Vec3T>
-void SampleFromVoxels<TreeOrAccT, 3, true>::cache(Vec3T<RealT>& xyz) const
+__hostdev__ void SampleFromVoxels<TreeOrAccT, 3, true>::cache(Vec3T<RealT>& xyz) const
 {
     CoordT ijk = Floor<CoordT>(xyz);
     if (ijk != mPos) {