fixes

Those functions were more useful when `FieldInferencingInterface`
was still declared further down in `node.cc`.

CMakeLists.txt

@@ -494,8 +494,7 @@ endif()
 
 # This should be turned off when Blender enter beta/rc/release
 if("${BLENDER_VERSION_CYCLE}" STREQUAL "release" OR
-   "${BLENDER_VERSION_CYCLE}" STREQUAL "rc" OR
-   "${BLENDER_VERSION_CYCLE}" STREQUAL "beta")
+   "${BLENDER_VERSION_CYCLE}" STREQUAL "rc")
   set(WITH_EXPERIMENTAL_FEATURES OFF)
 else()
   set(WITH_EXPERIMENTAL_FEATURES ON)

build_files/build_environment/cmake/nanovdb.cmake

@@ -42,7 +42,6 @@ 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
 )

build_files/build_environment/patches/nanovdb.diff

@@ -1,374 +0,0 @@
-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 @@
- 
- 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 @@
- 
- 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 @@
- 
- 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 @@
- 
- 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 @@
- 
- template<typename TreeOrAccT>
- template<typename RealT, template<typename...> class Vec3T>
--inline Vec3T<typename TreeOrAccT::ValueType> SampleFromVoxels<TreeOrAccT, 1, false>::gradient(Vec3T<RealT> xyz) const
-+inline __hostdev__ 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 @@
- 
- 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 @@
- }; // 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 @@
- 
- 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 @@
- }
- 
- 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(std::is_floating_point<ValueT>::value, "TrilinearSampler::zeroCrossing requires a floating-point type");
-     const bool less = v[0][0][0] < ValueT(0);
-@@ -624,7 +624,7 @@
- 
- 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);
-@@ -631,7 +631,7 @@
- }
- 
- 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 @@
- 
- 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 @@
- 
- 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 @@
- 
- 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 @@
- }; // 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 @@
- 
- 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 @@
- 
- 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) {

build_files/config/pipeline_config.yaml

@@ -5,38 +5,38 @@
 update-code:
     git:
         submodules:
-        -   branch: blender-v3.0-release
+        -   branch: master
             commit_id: HEAD
             path: release/scripts/addons
-        -   branch: blender-v3.0-release
+        -   branch: master
             commit_id: HEAD
             path: release/scripts/addons_contrib
-        -   branch: blender-v3.0-release
+        -   branch: master
             commit_id: HEAD
             path: release/datafiles/locale
-        -   branch: blender-v3.0-release
+        -   branch: master
             commit_id: HEAD
             path: source/tools
     svn:
         libraries:
             darwin-arm64:
-                branch: tags/blender-3.0-release
+                branch: trunk
                 commit_id: HEAD
                 path: lib/darwin_arm64
             darwin-x86_64:
-                branch: tags/blender-3.0-release
+                branch: trunk
                 commit_id: HEAD
                 path: lib/darwin
             linux-x86_64:
-                branch: tags/blender-3.0-release
+                branch: trunk
                 commit_id: HEAD
                 path: lib/linux_centos7_x86_64
             windows-amd64:
-                branch: tags/blender-3.0-release
+                branch: trunk
                 commit_id: HEAD
                 path: lib/win64_vc15
         tests:
-            branch: tags/blender-3.0-release
+            branch: trunk
             commit_id: HEAD
             path: lib/tests
         benchmarks:

doc/doxygen/Doxyfile

@@ -38,7 +38,7 @@ PROJECT_NAME           = Blender
 # could be handy for archiving the generated documentation or if some version
 # control system is used.
 
-PROJECT_NUMBER         = V3.0
+PROJECT_NUMBER         = V3.1
 
 # Using the PROJECT_BRIEF tag one can provide an optional one line description
 # for a project that appears at the top of each page and should give viewer a

doc/python_api/examples/bpy.types.Operator.1.py

@@ -42,13 +42,8 @@ class SimpleMouseOperator(bpy.types.Operator):
         self.y = event.mouse_y
         return self.execute(context)
 
-# Only needed if you want to add into a dynamic menu
-def menu_func(self, context):
-    self.layout.operator(SimpleMouseOperator.bl_idname, text="Simple Mouse Operator")
 
-# Register and add to the view menu (required to also use F3 search "Simple Mouse Operator" for quick access)
 bpy.utils.register_class(SimpleMouseOperator)
-bpy.types.VIEW3D_MT_view.append(menu_func)
 
 # Test call to the newly defined operator.
 # Here we call the operator and invoke it, meaning that the settings are taken

doc/python_api/examples/bpy.types.Operator.2.py

@@ -43,7 +43,7 @@ def menu_func(self, context):
     self.layout.operator(ExportSomeData.bl_idname, text="Text Export Operator")
 
 
-# Register and add to the file selector (required to also use F3 search "Text Export Operator" for quick access)
+# Register and add to the file selector
 bpy.utils.register_class(ExportSomeData)
 bpy.types.TOPBAR_MT_file_export.append(menu_func)
 

doc/python_api/examples/bpy.types.Operator.3.py

@@ -27,14 +27,8 @@ class DialogOperator(bpy.types.Operator):
         wm = context.window_manager
         return wm.invoke_props_dialog(self)
 
-# Only needed if you want to add into a dynamic menu
-def menu_func(self, context):
-    self.layout.operator(DialogOperator.bl_idname, text="Dialog Operator")
 
-
-# Register and add to the object menu (required to also use F3 search "Dialog Operator" for quick access)
 bpy.utils.register_class(DialogOperator)
-bpy.types.VIEW3D_MT_object.append(menu_func)
 
 # Test call.
 bpy.ops.object.dialog_operator('INVOKE_DEFAULT')

doc/python_api/examples/bpy.types.Operator.4.py

@@ -41,13 +41,8 @@ class CustomDrawOperator(bpy.types.Operator):
 
         col.prop(self, "my_string")
 
-# Only needed if you want to add into a dynamic menu
-def menu_func(self, context):
-    self.layout.operator(CustomDrawOperator.bl_idname, text="Custom Draw Operator")
 
-# Register and add to the object menu (required to also use F3 search "Custom Draw Operator" for quick access)
 bpy.utils.register_class(CustomDrawOperator)
-bpy.types.VIEW3D_MT_object.append(menu_func)
 
 # test call
 bpy.ops.object.custom_draw('INVOKE_DEFAULT')

doc/python_api/examples/bpy.types.Operator.5.py

@@ -55,13 +55,8 @@ class ModalOperator(bpy.types.Operator):
         context.window_manager.modal_handler_add(self)
         return {'RUNNING_MODAL'}
 
-# Only needed if you want to add into a dynamic menu
-def menu_func(self, context):
-    self.layout.operator(ModalOperator.bl_idname, text="Modal Operator")
 
-# Register and add to the object menu (required to also use F3 search "Modal Operator" for quick access)
 bpy.utils.register_class(ModalOperator)
-bpy.types.VIEW3D_MT_object.append(menu_func)
 
 # test call
 bpy.ops.object.modal_operator('INVOKE_DEFAULT')

doc/python_api/examples/bpy.types.Operator.6.py

@@ -31,13 +31,8 @@ class SearchEnumOperator(bpy.types.Operator):
         context.window_manager.invoke_search_popup(self)
         return {'RUNNING_MODAL'}
 
-# Only needed if you want to add into a dynamic menu
-def menu_func(self, context):
-    self.layout.operator(SearchEnumOperator.bl_idname, text="Search Enum Operator")
 
-# Register and add to the object menu (required to also use F3 search "Search Enum Operator" for quick access)
 bpy.utils.register_class(SearchEnumOperator)
-bpy.types.VIEW3D_MT_object.append(menu_func)
 
 # test call
 bpy.ops.object.search_enum_operator('INVOKE_DEFAULT')

doc/python_api/examples/bpy.types.Operator.py

@@ -22,13 +22,8 @@ class HelloWorldOperator(bpy.types.Operator):
         print("Hello World")
         return {'FINISHED'}
 
-# Only needed if you want to add into a dynamic menu
-def menu_func(self, context):
-    self.layout.operator(HelloWorldOperator.bl_idname, text="Hello World Operator")
 
-# Register and add to the view menu (required to also use F3 search "Hello World Operator" for quick access)
 bpy.utils.register_class(HelloWorldOperator)
-bpy.types.VIEW3D_MT_view.append(menu_func)
 
 # test call to the newly defined operator
 bpy.ops.wm.hello_world()

doc/python_api/rst/bgl.rst

@@ -106,6 +106,24 @@ including advanced features.
       floating-point values. These values are interpreted as a plane equation.
 
 
+.. function:: glColor (red, green, blue, alpha):
+
+   B{glColor3b, glColor3d, glColor3f, glColor3i, glColor3s, glColor3ub, glColor3ui, glColor3us,
+   glColor4b, glColor4d, glColor4f, glColor4i, glColor4s, glColor4ub, glColor4ui, glColor4us,
+   glColor3bv, glColor3dv, glColor3fv, glColor3iv, glColor3sv, glColor3ubv, glColor3uiv,
+   glColor3usv, glColor4bv, glColor4dv, glColor4fv, glColor4iv, glColor4sv, glColor4ubv,
+   glColor4uiv, glColor4usv}
+
+   Set a new color.
+
+   .. seealso:: `OpenGL Docs <https://khronos.org/registry/OpenGL-Refpages/gl4/html/glColor.xhtml>`__
+
+   :type red, green, blue, alpha: Depends on function prototype.
+   :arg red, green, blue: Specify new red, green, and blue values for the current color.
+   :arg alpha: Specifies a new alpha value for the current color. Included only in the
+      four-argument glColor4 commands. (With '4' colors only)
+
+
 .. function:: glColorMask(red, green, blue, alpha):
 
    Enable and disable writing of frame buffer color components

doc/python_api/sphinx_doc_gen.py

@@ -1123,7 +1123,7 @@ context_type_map = {
     "soft_body": ("SoftBodyModifier", False),
     "speaker": ("Speaker", False),
     "texture": ("Texture", False),
-    "texture_slot": ("TextureSlot", False),
+    "texture_slot": ("MaterialTextureSlot", False),
     "texture_user": ("ID", False),
     "texture_user_property": ("Property", False),
     "ui_list": ("UIList", False),
@@ -1224,7 +1224,10 @@ def pycontext2sphinx(basepath):
         while char_array[i] is not None:
             member = ctypes.string_at(char_array[i]).decode(encoding="ascii")
             fw(".. data:: %s\n\n" % member)
-            member_type, is_seq = context_type_map[member]
+            try:
+                member_type, is_seq = context_type_map[member]
+            except KeyError:
+                raise SystemExit("Error: context key %r not found in context_type_map; update %s" % (member, __file__)) from None
             fw("   :type: %s :class:`bpy.types.%s`\n\n" % ("sequence of " if is_seq else "", member_type))
             unique.add(member)
             i += 1
@@ -2251,7 +2254,7 @@ def main():
     # First monkey patch to load in fake members.
     setup_monkey_patch()
 
-    # Perform changes to Blender it's self.
+    # Perform changes to Blender itself.
     setup_data = setup_blender()
 
     # eventually, create the dirs

extern/hipew/README

@@ -1,12 +0,0 @@
-The HIP Extension Wrangler Library (HIPEW) is a cross-platform open-source
-C/C++ library to dynamically load the HIP library.
-
-HIP (Heterogeneous-Compute Interface for Portability) is an API for C++
-programming on AMD GPUs.
-
-It is maintained as part of the Blender project, but included in extern/
-for consistency with CUEW and CLEW libraries.
-
-LICENSE
-
-HIPEW is released under the Apache 2.0 license.

extern/hipew/README.blender

@@ -1,5 +0,0 @@
-Project: Blender
-URL: https://git.blender.org/blender.git
-License: Apache 2.0
-Upstream version: N/A
-Local modifications: None

extern/hipew/include/hipew.h

@@ -804,29 +804,31 @@ typedef enum hipDeviceP2PAttr {
 } hipDeviceP2PAttr;
 
 typedef struct HIP_MEMCPY3D {
-  unsigned int srcXInBytes;
-  unsigned int srcY;
-  unsigned int srcZ;
-  unsigned int srcLOD;
+  size_t srcXInBytes;
+  size_t srcY;
+  size_t srcZ;
+  size_t srcLOD;
   hipMemoryType srcMemoryType;
   const void* srcHost;
   hipDeviceptr_t srcDevice;
-  hArray srcArray;
-  unsigned int srcPitch;
-  unsigned int srcHeight;
-  unsigned int dstXInBytes;
-  unsigned int dstY;
-  unsigned int dstZ;
-  unsigned int dstLOD;
+  hArray * srcArray;
+  void* reserved0;
+  size_t srcPitch;
+  size_t srcHeight;
+  size_t dstXInBytes;
+  size_t dstY;
+  size_t dstZ;
+  size_t dstLOD;
   hipMemoryType dstMemoryType;
   void* dstHost;
   hipDeviceptr_t dstDevice;
-  hArray dstArray;
-  unsigned int dstPitch;
-  unsigned int dstHeight;
-  unsigned int WidthInBytes;
-  unsigned int Height;
-  unsigned int Depth;
+  hArray * dstArray;
+  void* reserved1;
+  size_t dstPitch;
+  size_t dstHeight;
+  size_t WidthInBytes;
+  size_t Height;
+  size_t Depth;
 } HIP_MEMCPY3D;
 
 typedef struct HIP_MEMCPY3D_PEER_st {
@@ -877,7 +879,7 @@ typedef struct HIP_RESOURCE_DESC_st {
   hipResourceType resType;
   union {
     struct {
-      hArray h_Array;
+      hArray * h_Array;
     } array;
     struct {
       hipMipmappedArray_t hMipmappedArray;
@@ -1072,10 +1074,9 @@ typedef enum hiprtcResult {
 typedef hipError_t HIPAPI thipGetErrorName(hipError_t error, const char** pStr);
 typedef hipError_t HIPAPI thipInit(unsigned int Flags);
 typedef hipError_t HIPAPI thipDriverGetVersion(int* driverVersion);
-typedef hipError_t HIPAPI thipGetDevice(int* device);
+typedef hipError_t HIPAPI thipGetDevice(hipDevice_t* device, int ordinal);
 typedef hipError_t HIPAPI thipGetDeviceCount(int* count);
 typedef hipError_t HIPAPI thipGetDeviceProperties(hipDeviceProp_t* props, int deviceId);
-typedef hipError_t HIPAPI thipDeviceGet(hipDevice_t* device, int ordinal);
 typedef hipError_t HIPAPI thipDeviceGetName(char* name, int len, hipDevice_t dev);
 typedef hipError_t HIPAPI thipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attrib, hipDevice_t dev);
 typedef hipError_t HIPAPI thipDeviceComputeCapability(int* major, int* minor, hipDevice_t dev);
@@ -1208,7 +1209,6 @@ extern thipDriverGetVersion *hipDriverGetVersion;
 extern thipGetDevice *hipGetDevice;
 extern thipGetDeviceCount *hipGetDeviceCount;
 extern thipGetDeviceProperties *hipGetDeviceProperties;
-extern thipDeviceGet* hipDeviceGet;
 extern thipDeviceGetName *hipDeviceGetName;
 extern thipDeviceGetAttribute *hipDeviceGetAttribute;
 extern thipDeviceComputeCapability *hipDeviceComputeCapability;
@@ -1333,7 +1333,6 @@ enum {
   HIPEW_SUCCESS = 0,
   HIPEW_ERROR_OPEN_FAILED = -1,
   HIPEW_ERROR_ATEXIT_FAILED = -2,
-  HIPEW_ERROR_OLD_DRIVER = -3,
 };
 
 enum {

extern/hipew/src/hipew.c

@@ -71,7 +71,6 @@ thipDriverGetVersion *hipDriverGetVersion;
 thipGetDevice *hipGetDevice;
 thipGetDeviceCount *hipGetDeviceCount;
 thipGetDeviceProperties *hipGetDeviceProperties;
-thipDeviceGet* hipDeviceGet;
 thipDeviceGetName *hipDeviceGetName;
 thipDeviceGetAttribute *hipDeviceGetAttribute;
 thipDeviceComputeCapability *hipDeviceComputeCapability;
@@ -214,36 +213,6 @@ static void hipewHipExit(void) {
   }
 }
 
-#ifdef _WIN32
-static int hipewHasOldDriver(const char *hip_path) {
-  DWORD verHandle = 0;
-  DWORD verSize = GetFileVersionInfoSize(hip_path, &verHandle);
-  int old_driver = 0;
-  if (verSize != 0) {
-    LPSTR verData = (LPSTR)malloc(verSize);
-    if (GetFileVersionInfo(hip_path, verHandle, verSize, verData)) {
-      LPBYTE lpBuffer = NULL;
-      UINT size = 0;
-      if (VerQueryValue(verData, "\\", (VOID FAR * FAR *)&lpBuffer, &size)) {
-        if (size) {
-          VS_FIXEDFILEINFO *verInfo = (VS_FIXEDFILEINFO *)lpBuffer;
-          /* Magic value from
-           * https://docs.microsoft.com/en-us/windows/win32/api/verrsrc/ns-verrsrc-vs_fixedfileinfo */
-          if (verInfo->dwSignature == 0xfeef04bd) {
-            unsigned int fileVersionLS0 = (verInfo->dwFileVersionLS >> 16) & 0xffff;
-            unsigned int fileversionLS1 = (verInfo->dwFileVersionLS >> 0) & 0xffff;
-            /* Corresponds to versions older than AMD Radeon Pro 21.Q4. */
-            old_driver = ((fileVersionLS0 < 3354) || (fileVersionLS0 == 3354 && fileversionLS1 < 13));
-          }
-        }
-      }
-    }
-    free(verData);
-  }
-  return old_driver;
-}
-#endif
-
 static int hipewHipInit(void) {
   /* Library paths. */
 #ifdef _WIN32
@@ -271,14 +240,6 @@ static int hipewHipInit(void) {
     return result;
   }
 
-#ifdef _WIN32
-  /* Test for driver version. */
-  if(hipewHasOldDriver(hip_paths[0])) {
-     result = HIPEW_ERROR_OLD_DRIVER;
-     return result;
-  }
-#endif
-
   /* Load library. */
   hip_lib = dynamic_library_open_find(hip_paths);
 
@@ -294,7 +255,6 @@ static int hipewHipInit(void) {
   HIP_LIBRARY_FIND_CHECKED(hipGetDevice);
   HIP_LIBRARY_FIND_CHECKED(hipGetDeviceCount);
   HIP_LIBRARY_FIND_CHECKED(hipGetDeviceProperties);
-  HIP_LIBRARY_FIND_CHECKED(hipDeviceGet);
   HIP_LIBRARY_FIND_CHECKED(hipDeviceGetName);
   HIP_LIBRARY_FIND_CHECKED(hipDeviceGetAttribute);
   HIP_LIBRARY_FIND_CHECKED(hipDeviceComputeCapability);

extern/nanosvg/README.blender

@@ -1,7 +1,7 @@
 Project: NanoSVG
 URL: https://github.com/memononen/nanosvg
 License: zlib
-Upstream version: 3cdd4a9d7886
+Upstream version: 
 Local modifications: Added some functionality to manage grease pencil layers
 
 Added a fix to SVG import arc and float errors (https://developer.blender.org/rB11dc674c78b49fc4e0b7c134c375b6c8b8eacbcc)

intern/cycles/app/cycles_standalone.cpp

@@ -82,7 +82,7 @@ static void session_print_status()
   string status, substatus;
 
   /* get status */
-  double progress = options.session->progress.get_progress();
+  float progress = options.session->progress.get_progress();
   options.session->progress.get_status(status, substatus);
 
   if (substatus != "")
@@ -183,7 +183,7 @@ static void display_info(Progress &progress)
 
   progress.get_time(total_time, sample_time);
   progress.get_status(status, substatus);
-  double progress_val = progress.get_progress();
+  float progress_val = progress.get_progress();
 
   if (substatus != "")
     status += ": " + substatus;

intern/cycles/blender/CMakeLists.txt

@@ -138,11 +138,6 @@ endif()
 
 blender_add_lib(bf_intern_cycles "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")
 
-# avoid link failure with clang 3.4 debug
-if(CMAKE_C_COMPILER_ID MATCHES "Clang" AND NOT ${CMAKE_C_COMPILER_VERSION} VERSION_LESS '3.4')
-  string(APPEND CMAKE_CXX_FLAGS_DEBUG " -gline-tables-only")
-endif()
-
 add_dependencies(bf_intern_cycles bf_rna)
 
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${ADDON_FILES}" ${CYCLES_INSTALL_PATH})

intern/cycles/blender/addon/properties.py

@@ -325,6 +325,13 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         default=1024,
     )
 
+    sample_offset: IntProperty(
+        name="Sample Offset",
+        description="Number of samples to skip when starting render",
+        min=0, max=(1 << 24),
+        default=0,
+    )
+
     time_limit: FloatProperty(
         name="Time Limit",
         description="Limit the render time (excluding synchronization time)."
@@ -346,7 +353,7 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         name="Scrambling Distance",
         default=1.0,
         min=0.0, max=1.0,
-        description="Reduce randomization between pixels to improve GPU rendering performance, at the cost of possible rendering artifacts if set too low. Only works when not using adaptive sampling",
+        description="Lower values give faster rendering with GPU rendering and less noise with all devices at the cost of possible artifacts if set too low. Only works when not using adaptive sampling",
     )
     preview_scrambling_distance: BoolProperty(
         name="Scrambling Distance viewport",
@@ -354,10 +361,10 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         description="Uses the Scrambling Distance value for the viewport. Faster but may flicker",
     )
 
-    auto_scrambling_distance: BoolProperty(
-        name="Automatic Scrambling Distance",
+    adaptive_scrambling_distance: BoolProperty(
+        name="Adaptive Scrambling Distance",
         default=False,
-        description="Automatically reduce the randomization between pixels to improve GPU rendering performance, at the cost of possible rendering artifacts. Only works when not using adaptive sampling",
+        description="Uses a formula to adapt the scrambling distance strength based on the sample count",
     )
 
     use_layer_samples: EnumProperty(
@@ -770,8 +777,8 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
     )
 
     use_auto_tile: BoolProperty(
-        name="Use Tiling",
-        description="Render high resolution images in tiles to reduce memory usage, using the specified tile size. Tiles are cached to disk while rendering to save memory",
+        name="Auto Tiles",
+        description="Automatically render high resolution images in tiles to reduce memory usage, using the specified tile size. Tiles are cached to disk while rendering to save memory",
         default=True,
     )
     tile_size: IntProperty(

intern/cycles/blender/addon/ui.py

@@ -290,15 +290,18 @@ class CYCLES_RENDER_PT_sampling_advanced(CyclesButtonsPanel, Panel):
         col.active = not (cscene.use_adaptive_sampling and cscene.use_preview_adaptive_sampling)
         col.prop(cscene, "sampling_pattern", text="Pattern")
 
+        col = layout.column(align=True)
+        col.prop(cscene, "sample_offset")
+
         layout.separator()
 
-        heading = layout.column(align=True, heading="Scrambling Distance")
-        heading.active = not (cscene.use_adaptive_sampling and cscene.use_preview_adaptive_sampling)
-        heading.prop(cscene, "auto_scrambling_distance", text="Automatic")
-        sub = heading.row()
+        col = layout.column(align=True)
+        col.active = not (cscene.use_adaptive_sampling and cscene.use_preview_adaptive_sampling)
+        col.prop(cscene, "scrambling_distance", text="Scrambling Distance")
+        col.prop(cscene, "adaptive_scrambling_distance", text="Adaptive")
+        sub = col.row(align=True)
         sub.active = not cscene.use_preview_adaptive_sampling
         sub.prop(cscene, "preview_scrambling_distance", text="Viewport")
-        heading.prop(cscene, "scrambling_distance", text="Multiplier")
 
         layout.separator()
 

intern/cycles/blender/curves.cpp

@@ -819,14 +819,11 @@ void BlenderSync::sync_hair(BL::Depsgraph b_depsgraph, BObjectInfo &b_ob_info, H
   new_hair.set_used_shaders(used_shaders);
 
   if (view_layer.use_hair) {
-#ifdef WITH_HAIR_NODES
     if (b_ob_info.object_data.is_a(&RNA_Hair)) {
       /* Hair object. */
       sync_hair(&new_hair, b_ob_info, false);
     }
-    else
-#endif
-    {
+    else {
       /* Particle hair. */
       bool need_undeformed = new_hair.need_attribute(scene, ATTR_STD_GENERATED);
       BL::Mesh b_mesh = object_to_mesh(
@@ -873,15 +870,12 @@ void BlenderSync::sync_hair_motion(BL::Depsgraph b_depsgraph,
 
   /* Export deformed coordinates. */
   if (ccl::BKE_object_is_deform_modified(b_ob_info, b_scene, preview)) {
-#ifdef WITH_HAIR_NODES
     if (b_ob_info.object_data.is_a(&RNA_Hair)) {
       /* Hair object. */
       sync_hair(hair, b_ob_info, true, motion_step);
       return;
     }
-    else
-#endif
-    {
+    else {
       /* Particle hair. */
       BL::Mesh b_mesh = object_to_mesh(
           b_data, b_ob_info, b_depsgraph, false, Mesh::SUBDIVISION_NONE);

intern/cycles/blender/geometry.cpp

@@ -31,11 +31,7 @@ CCL_NAMESPACE_BEGIN
 
 static Geometry::Type determine_geom_type(BObjectInfo &b_ob_info, bool use_particle_hair)
 {
-#ifdef WITH_HAIR_NODES
   if (b_ob_info.object_data.is_a(&RNA_Hair) || use_particle_hair) {
-#else
-  if (use_particle_hair) {
-#endif
     return Geometry::HAIR;
   }
 
@@ -219,11 +215,7 @@ void BlenderSync::sync_geometry_motion(BL::Depsgraph &b_depsgraph,
     if (progress.get_cancel())
       return;
 
-#ifdef WITH_HAIR_NODES
     if (b_ob_info.object_data.is_a(&RNA_Hair) || use_particle_hair) {
-#else
-    if (use_particle_hair) {
-#endif
       Hair *hair = static_cast<Hair *>(geom);
       sync_hair_motion(b_depsgraph, b_ob_info, hair, motion_step);
     }

intern/cycles/blender/image.cpp

@@ -24,14 +24,8 @@ CCL_NAMESPACE_BEGIN
 
 /* Packed Images */
 
-BlenderImageLoader::BlenderImageLoader(BL::Image b_image,
-                                       const int frame,
-                                       const bool is_preview_render)
-    : b_image(b_image),
-      frame(frame),
-      /* Don't free cache for preview render to avoid race condition from T93560, to be fixed
-         properly later as we are close to release. */
-      free_cache(!is_preview_render && !b_image.has_data())
+BlenderImageLoader::BlenderImageLoader(BL::Image b_image, int frame)
+    : b_image(b_image), frame(frame), free_cache(!b_image.has_data())
 {
 }
 

intern/cycles/blender/image.h

@@ -25,7 +25,7 @@ CCL_NAMESPACE_BEGIN
 
 class BlenderImageLoader : public ImageLoader {
  public:
-  BlenderImageLoader(BL::Image b_image, const int frame, const bool is_preview_render);
+  BlenderImageLoader(BL::Image b_image, int frame);
 
   bool load_metadata(const ImageDeviceFeatures &features, ImageMetaData &metadata) override;
   bool load_pixels(const ImageMetaData &metadata,

intern/cycles/blender/output_driver.cpp

@@ -120,7 +120,7 @@ void BlenderOutputDriver::write_render_tile(const Tile &tile)
     b_pass.rect(&pixels[0]);
   }
 
-  b_engine_.end_result(b_rr, false, false, true);
+  b_engine_.end_result(b_rr, true, false, true);
 }
 
 CCL_NAMESPACE_END

intern/cycles/blender/session.cpp

@@ -129,7 +129,7 @@ void BlenderSession::create_session()
   /* reset status/progress */
   last_status = "";
   last_error = "";
-  last_progress = -1.0;
+  last_progress = -1.0f;
   start_resize_time = 0.0;
 
   /* create session */
@@ -606,6 +606,19 @@ void BlenderSession::bake(BL::Depsgraph &b_depsgraph_,
   pass->set_type(bake_type_to_pass(bake_type, bake_filter));
   pass->set_include_albedo((bake_filter & BL::BakeSettings::pass_filter_COLOR));
 
+  if (pass->get_type() == PASS_COMBINED) {
+    /* Filtering settings for combined pass. */
+    Integrator *integrator = scene->integrator;
+    integrator->set_use_direct_light((bake_filter & BL::BakeSettings::pass_filter_DIRECT) != 0);
+    integrator->set_use_indirect_light((bake_filter & BL::BakeSettings::pass_filter_INDIRECT) !=
+                                       0);
+    integrator->set_use_diffuse((bake_filter & BL::BakeSettings::pass_filter_DIFFUSE) != 0);
+    integrator->set_use_glossy((bake_filter & BL::BakeSettings::pass_filter_GLOSSY) != 0);
+    integrator->set_use_transmission((bake_filter & BL::BakeSettings::pass_filter_TRANSMISSION) !=
+                                     0);
+    integrator->set_use_emission((bake_filter & BL::BakeSettings::pass_filter_EMIT) != 0);
+  }
+
   session->set_display_driver(nullptr);
   session->set_output_driver(make_unique<BlenderOutputDriver>(b_engine));
 
@@ -615,24 +628,6 @@ void BlenderSession::bake(BL::Depsgraph &b_depsgraph_,
     sync->sync_camera(b_render, b_camera_override, width, height, "");
     sync->sync_data(
         b_render, b_depsgraph, b_v3d, b_camera_override, width, height, &python_thread_state);
-
-    /* Filtering settings for combined pass. */
-    if (pass->get_type() == PASS_COMBINED) {
-      Integrator *integrator = scene->integrator;
-      integrator->set_use_direct_light((bake_filter & BL::BakeSettings::pass_filter_DIRECT) != 0);
-      integrator->set_use_indirect_light((bake_filter & BL::BakeSettings::pass_filter_INDIRECT) !=
-                                         0);
-      integrator->set_use_diffuse((bake_filter & BL::BakeSettings::pass_filter_DIFFUSE) != 0);
-      integrator->set_use_glossy((bake_filter & BL::BakeSettings::pass_filter_GLOSSY) != 0);
-      integrator->set_use_transmission(
-          (bake_filter & BL::BakeSettings::pass_filter_TRANSMISSION) != 0);
-      integrator->set_use_emission((bake_filter & BL::BakeSettings::pass_filter_EMIT) != 0);
-    }
-
-    /* Always use transpanent background for baking. */
-    scene->background->set_transparent(true);
-
-    /* Load built-in images from Blender. */
     builtin_images_load();
   }
 
@@ -859,7 +854,7 @@ void BlenderSession::get_status(string &status, string &substatus)
   session->progress.get_status(status, substatus);
 }
 
-void BlenderSession::get_progress(double &progress, double &total_time, double &render_time)
+void BlenderSession::get_progress(float &progress, double &total_time, double &render_time)
 {
   session->progress.get_time(total_time, render_time);
   progress = session->progress.get_progress();
@@ -867,10 +862,10 @@ void BlenderSession::get_progress(double &progress, double &total_time, double &
 
 void BlenderSession::update_bake_progress()
 {
-  double progress = session->progress.get_progress();
+  float progress = session->progress.get_progress();
 
   if (progress != last_progress) {
-    b_engine.update_progress((float)progress);
+    b_engine.update_progress(progress);
     last_progress = progress;
   }
 }
@@ -879,7 +874,7 @@ void BlenderSession::update_status_progress()
 {
   string timestatus, status, substatus;
   string scene_status = "";
-  double progress;
+  float progress;
   double total_time, remaining_time = 0, render_time;
   float mem_used = (float)session->stats.mem_used / 1024.0f / 1024.0f;
   float mem_peak = (float)session->stats.mem_peak / 1024.0f / 1024.0f;
@@ -923,7 +918,7 @@ void BlenderSession::update_status_progress()
     last_status_time = current_time;
   }
   if (progress != last_progress) {
-    b_engine.update_progress((float)progress);
+    b_engine.update_progress(progress);
     last_progress = progress;
   }
 

intern/cycles/blender/session.h

@@ -82,7 +82,7 @@ class BlenderSession {
   void tag_redraw();
   void tag_update();
   void get_status(string &status, string &substatus);
-  void get_progress(double &progress, double &total_time, double &render_time);
+  void get_progress(float &progress, double &total_time, double &render_time);
   void test_cancel();
   void update_status_progress();
   void update_bake_progress();
@@ -108,7 +108,7 @@ class BlenderSession {
 
   string last_status;
   string last_error;
-  double last_progress;
+  float last_progress;
   double last_status_time;
 
   int width, height;

intern/cycles/blender/shader.cpp

@@ -762,8 +762,7 @@ static ShaderNode *add_node(Scene *scene,
         int scene_frame = b_scene.frame_current();
         int image_frame = image_user_frame_number(b_image_user, b_image, scene_frame);
         image->handle = scene->image_manager->add_image(
-            new BlenderImageLoader(b_image, image_frame, b_engine.is_preview()),
-            image->image_params());
+            new BlenderImageLoader(b_image, image_frame), image->image_params());
       }
       else {
         ustring filename = ustring(
@@ -798,9 +797,8 @@ static ShaderNode *add_node(Scene *scene,
       if (is_builtin) {
         int scene_frame = b_scene.frame_current();
         int image_frame = image_user_frame_number(b_image_user, b_image, scene_frame);
-        env->handle = scene->image_manager->add_image(
-            new BlenderImageLoader(b_image, image_frame, b_engine.is_preview()),
-            env->image_params());
+        env->handle = scene->image_manager->add_image(new BlenderImageLoader(b_image, image_frame),
+                                                      env->image_params());
       }
       else {
         env->set_filename(

intern/cycles/blender/sync.cpp

@@ -365,8 +365,8 @@ void BlenderSync::sync_integrator(BL::ViewLayer &b_view_layer, bool background)
 
   int samples = get_int(cscene, "samples");
   float scrambling_distance = get_float(cscene, "scrambling_distance");
-  bool auto_scrambling_distance = get_boolean(cscene, "auto_scrambling_distance");
-  if (auto_scrambling_distance) {
+  bool adaptive_scrambling_distance = get_boolean(cscene, "adaptive_scrambling_distance");
+  if (adaptive_scrambling_distance) {
     scrambling_distance *= 4.0f / sqrtf(samples);
   }
 
@@ -835,18 +835,25 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine &b_engine,
   /* samples */
   int samples = get_int(cscene, "samples");
   int preview_samples = get_int(cscene, "preview_samples");
+  int sample_offset = get_int(cscene, "sample_offset");
 
   if (background) {
     params.samples = samples;
+    params.sample_offset = sample_offset;
   }
   else {
     params.samples = preview_samples;
-    if (params.samples == 0)
+    if (params.samples == 0) {
       params.samples = INT_MAX;
+    }
+    params.sample_offset = 0;
   }
 
+  /* Clamp sample offset. */
+  params.sample_offset = clamp(params.sample_offset, 0, Integrator::MAX_SAMPLES);
+
   /* Clamp samples. */
-  params.samples = min(params.samples, Integrator::MAX_SAMPLES);
+  params.samples = clamp(params.samples, 0, Integrator::MAX_SAMPLES - params.sample_offset);
 
   /* Viewport Performance */
   params.pixel_size = b_engine.get_preview_pixel_size(b_scene);

intern/cycles/bvh/optix.cpp

@@ -30,17 +30,15 @@ BVHOptiX::BVHOptiX(const BVHParams &params_,
     : BVH(params_, geometry_, objects_),
       device(device),
       traversable_handle(0),
-      as_data(make_unique<device_only_memory<char>>(
-          device, params.top_level ? "optix tlas" : "optix blas", false)),
-      motion_transform_data(
-          make_unique<device_only_memory<char>>(device, "optix motion transform", false))
+      as_data(device, params_.top_level ? "optix tlas" : "optix blas", false),
+      motion_transform_data(device, "optix motion transform", false)
 {
 }
 
 BVHOptiX::~BVHOptiX()
 {
-  /* Acceleration structure memory is delayed freed on device, since deleting the
-   * BVH may happen while still being used for rendering. */
+  // Acceleration structure memory is delayed freed on device, since deleting the
+  // BVH may happen while still being used for rendering.
   device->release_optix_bvh(this);
 }
 

intern/cycles/bvh/optix.h

@@ -25,16 +25,14 @@
 
 #  include "device/memory.h"
 
-#  include "util/unique_ptr.h"
-
 CCL_NAMESPACE_BEGIN
 
 class BVHOptiX : public BVH {
  public:
   Device *device;
   uint64_t traversable_handle;
-  unique_ptr<device_only_memory<char>> as_data;
-  unique_ptr<device_only_memory<char>> motion_transform_data;
+  device_only_memory<char> as_data;
+  device_only_memory<char> motion_transform_data;
 
  protected:
   friend class BVH;

intern/cycles/device/cpu/device.cpp

@@ -38,7 +38,6 @@ void device_cpu_info(vector<DeviceInfo> &devices)
   info.id = "CPU";
   info.num = 0;
   info.has_osl = true;
-  info.has_half_images = true;
   info.has_nanovdb = true;
   info.has_profiling = true;
   if (openimagedenoise_supported()) {

intern/cycles/device/cpu/device_impl.cpp

@@ -134,7 +134,8 @@ void CPUDevice::mem_alloc(device_memory &mem)
               << string_human_readable_size(mem.memory_size()) << ")";
     }
 
-    if (mem.type == MEM_DEVICE_ONLY || !mem.host_pointer) {
+    if (mem.type == MEM_DEVICE_ONLY) {
+      assert(!mem.host_pointer);
       size_t alignment = MIN_ALIGNMENT_CPU_DATA_TYPES;
       void *data = util_aligned_malloc(mem.memory_size(), alignment);
       mem.device_pointer = (device_ptr)data;
@@ -193,7 +194,7 @@ void CPUDevice::mem_free(device_memory &mem)
     tex_free((device_texture &)mem);
   }
   else if (mem.device_pointer) {
-    if (mem.type == MEM_DEVICE_ONLY || !mem.host_pointer) {
+    if (mem.type == MEM_DEVICE_ONLY) {
       util_aligned_free((void *)mem.device_pointer);
     }
     mem.device_pointer = 0;

intern/cycles/device/cuda/device.cpp

@@ -144,7 +144,6 @@ void device_cuda_info(vector<DeviceInfo> &devices)
     info.description = string(name);
     info.num = num;
 
-    info.has_half_images = (major >= 3);
     info.has_nanovdb = true;
     info.denoisers = 0;
 

intern/cycles/device/cuda/device_impl.cpp

@@ -680,7 +680,7 @@ CUDADevice::CUDAMem *CUDADevice::generic_alloc(device_memory &mem, size_t pitch_
 
   void *shared_pointer = 0;
 
-  if (mem_alloc_result != CUDA_SUCCESS && can_map_host && mem.type != MEM_DEVICE_ONLY) {
+  if (mem_alloc_result != CUDA_SUCCESS && can_map_host) {
     if (mem.shared_pointer) {
       /* Another device already allocated host memory. */
       mem_alloc_result = CUDA_SUCCESS;
@@ -703,14 +703,8 @@ CUDADevice::CUDAMem *CUDADevice::generic_alloc(device_memory &mem, size_t pitch_
   }
 
   if (mem_alloc_result != CUDA_SUCCESS) {
-    if (mem.type == MEM_DEVICE_ONLY) {
-      status = " failed, out of device memory";
-      set_error("System is out of GPU memory");
-    }
-    else {
-      status = " failed, out of device and host memory";
-      set_error("System is out of GPU and shared host memory");
-    }
+    status = " failed, out of device and host memory";
+    set_error("System is out of GPU and shared host memory");
   }
 
   if (mem.name) {
@@ -783,7 +777,6 @@ void CUDADevice::generic_free(device_memory &mem)
   if (mem.device_pointer) {
     CUDAContextScope scope(this);
     thread_scoped_lock lock(cuda_mem_map_mutex);
-    DCHECK(cuda_mem_map.find(&mem) != cuda_mem_map.end());
     const CUDAMem &cmem = cuda_mem_map[&mem];
 
     /* If cmem.use_mapped_host is true, reference counting is used
@@ -938,7 +931,6 @@ void CUDADevice::tex_alloc(device_texture &mem)
 {
   CUDAContextScope scope(this);
 
-  /* General variables for both architectures */
   string bind_name = mem.name;
   size_t dsize = datatype_size(mem.data_type);
   size_t size = mem.memory_size();
@@ -1101,7 +1093,6 @@ void CUDADevice::tex_alloc(device_texture &mem)
 
   if (mem.info.data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT &&
       mem.info.data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT3) {
-    /* Kepler+, bindless textures. */
     CUDA_RESOURCE_DESC resDesc;
     memset(&resDesc, 0, sizeof(resDesc));
 
@@ -1152,7 +1143,6 @@ void CUDADevice::tex_free(device_texture &mem)
   if (mem.device_pointer) {
     CUDAContextScope scope(this);
     thread_scoped_lock lock(cuda_mem_map_mutex);
-    DCHECK(cuda_mem_map.find(&mem) != cuda_mem_map.end());
     const CUDAMem &cmem = cuda_mem_map[&mem];
 
     if (cmem.texobject) {

intern/cycles/device/device.cpp

@@ -286,7 +286,6 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
   info.description = "Multi Device";
   info.num = 0;
 
-  info.has_half_images = true;
   info.has_nanovdb = true;
   info.has_osl = true;
   info.has_profiling = true;
@@ -333,7 +332,6 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
     }
 
     /* Accumulate device info. */
-    info.has_half_images &= device.has_half_images;
     info.has_nanovdb &= device.has_nanovdb;
     info.has_osl &= device.has_osl;
     info.has_profiling &= device.has_profiling;

intern/cycles/device/device.h

@@ -73,7 +73,6 @@ class DeviceInfo {
   int num;
   bool display_device;        /* GPU is used as a display device. */
   bool has_nanovdb;           /* Support NanoVDB volumes. */
-  bool has_half_images;       /* Support half-float textures. */
   bool has_osl;               /* Support Open Shading Language. */
   bool has_profiling;         /* Supports runtime collection of profiling info. */
   bool has_peer_memory;       /* GPU has P2P access to memory of another GPU. */
@@ -90,7 +89,6 @@ class DeviceInfo {
     num = 0;
     cpu_threads = 0;
     display_device = false;
-    has_half_images = false;
     has_nanovdb = false;
     has_osl = false;
     has_profiling = false;

intern/cycles/device/hip/device.cpp

@@ -57,16 +57,9 @@ bool device_hip_init()
     }
   }
   else {
-    if (hipew_result == HIPEW_ERROR_ATEXIT_FAILED) {
-      VLOG(1) << "HIPEW initialization failed: Error setting up atexit() handler";
-    }
-    else if (hipew_result == HIPEW_ERROR_OLD_DRIVER) {
-      VLOG(1) << "HIPEW initialization failed: Driver version too old, requires AMD Radeon Pro "
-                 "21.Q4 driver or newer";
-    }
-    else {
-      VLOG(1) << "HIPEW initialization failed: Error opening HIP dynamic library";
-    }
+    VLOG(1) << "HIPEW initialization failed: "
+            << ((hipew_result == HIPEW_ERROR_ATEXIT_FAILED) ? "Error setting up atexit() handler" :
+                                                              "Error opening the library");
   }
 
   return result;
@@ -148,7 +141,6 @@ void device_hip_info(vector<DeviceInfo> &devices)
     info.description = string(name);
     info.num = num;
 
-    info.has_half_images = true;
     info.has_nanovdb = true;
     info.denoisers = 0;
 

intern/cycles/device/hip/device_impl.cpp

@@ -99,7 +99,7 @@ HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
   }
 
   /* Setup device and context. */
-  result = hipDeviceGet(&hipDevice, hipDevId);
+  result = hipGetDevice(&hipDevice, hipDevId);
   if (result != hipSuccess) {
     set_error(string_printf("Failed to get HIP device handle from ordinal (%s)",
                             hipewErrorString(result)));
@@ -243,7 +243,7 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
   hipGetDeviceProperties(&props, hipDevId);
 
   /* gcnArchName can contain tokens after the arch name with features, ie.
-    "gfx1010:sramecc-:xnack-" so we tokenize it to get the first part. */
+   * `gfx1010:sramecc-:xnack-` so we tokenize it to get the first part. */
   char *arch = strtok(props.gcnArchName, ":");
   if (arch == NULL) {
     arch = props.gcnArchName;
@@ -374,10 +374,9 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
 
 bool HIPDevice::load_kernels(const uint kernel_features)
 {
-  /* TODO(sergey): Support kernels re-load for CUDA devices adaptive compile.
+  /* TODO(sergey): Support kernels re-load for HIP devices adaptive compile.
    *
-   * Currently re-loading kernel will invalidate memory pointers,
-   * causing problems in cuCtxSynchronize.
+   * Currently re-loading kernels will invalidate memory pointers.
    */
   if (hipModule) {
     if (use_adaptive_compilation()) {
@@ -745,7 +744,6 @@ void HIPDevice::generic_free(device_memory &mem)
   if (mem.device_pointer) {
     HIPContextScope scope(this);
     thread_scoped_lock lock(hip_mem_map_mutex);
-    DCHECK(hip_mem_map.find(&mem) != hip_mem_map.end());
     const HIPMem &cmem = hip_mem_map[&mem];
 
     /* If cmem.use_mapped_host is true, reference counting is used
@@ -899,7 +897,6 @@ void HIPDevice::tex_alloc(device_texture &mem)
 {
   HIPContextScope scope(this);
 
-  /* General variables for both architectures */
   string bind_name = mem.name;
   size_t dsize = datatype_size(mem.data_type);
   size_t size = mem.memory_size();
@@ -989,16 +986,16 @@ void HIPDevice::tex_alloc(device_texture &mem)
             << string_human_readable_number(mem.memory_size()) << " bytes. ("
             << string_human_readable_size(mem.memory_size()) << ")";
 
-    hip_assert(hipArray3DCreate((hArray *)&array_3d, &desc));
+    hip_assert(hipArray3DCreate(&array_3d, &desc));
 
     if (!array_3d) {
       return;
     }
 
     HIP_MEMCPY3D param;
-    memset(&param, 0, sizeof(HIP_MEMCPY3D));
+    memset(&param, 0, sizeof(param));
     param.dstMemoryType = hipMemoryTypeArray;
-    param.dstArray = array_3d;
+    param.dstArray = &array_3d;
     param.srcMemoryType = hipMemoryTypeHost;
     param.srcHost = mem.host_pointer;
     param.srcPitch = src_pitch;
@@ -1064,13 +1061,12 @@ void HIPDevice::tex_alloc(device_texture &mem)
 
   if (mem.info.data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT &&
       mem.info.data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT3) {
-    /* Bindless textures. */
     hipResourceDesc resDesc;
     memset(&resDesc, 0, sizeof(resDesc));
 
     if (array_3d) {
       resDesc.resType = hipResourceTypeArray;
-      resDesc.res.array.h_Array = array_3d;
+      resDesc.res.array.h_Array = &array_3d;
       resDesc.flags = 0;
     }
     else if (mem.data_height > 0) {
@@ -1115,7 +1111,6 @@ void HIPDevice::tex_free(device_texture &mem)
   if (mem.device_pointer) {
     HIPContextScope scope(this);
     thread_scoped_lock lock(hip_mem_map_mutex);
-    DCHECK(hip_mem_map.find(&mem) != hip_mem_map.end());
     const HIPMem &cmem = hip_mem_map[&mem];
 
     if (cmem.texobject) {

intern/cycles/device/hip/device_impl.h

@@ -93,9 +93,7 @@ class HIPDevice : public Device {
 
   virtual string compile_kernel_get_common_cflags(const uint kernel_features);
 
-  string compile_kernel(const uint kernel_features,
-                        const char *name,
-                        const char *base = "hip");
+  string compile_kernel(const uint kernel_features, const char *name, const char *base = "hip");
 
   virtual bool load_kernels(const uint kernel_features) override;
   void reserve_local_memory(const uint kernel_features);

intern/cycles/device/hip/graphics_interop.h

@@ -48,7 +48,7 @@ class HIPDeviceGraphicsInterop : public DeviceGraphicsInterop {
   HIPDeviceQueue *queue_ = nullptr;
   HIPDevice *device_ = nullptr;
 
-  /* OpenGL PBO which is currently registered as the destination for the CUDA buffer. */
+  /* OpenGL PBO which is currently registered as the destination for the HIP buffer. */
   uint opengl_pbo_id_ = 0;
   /* Buffer area in pixels of the corresponding PBO. */
   int64_t buffer_area_ = 0;

intern/cycles/device/memory.cpp

@@ -44,6 +44,45 @@ device_memory::device_memory(Device *device, const char *name, MemoryType type)
 {
 }
 
+device_memory::device_memory(device_memory &&other) noexcept
+    : data_type(other.data_type),
+      data_elements(other.data_elements),
+      data_size(other.data_size),
+      device_size(other.device_size),
+      data_width(other.data_width),
+      data_height(other.data_height),
+      data_depth(other.data_depth),
+      type(other.type),
+      name(other.name),
+      device(other.device),
+      device_pointer(other.device_pointer),
+      host_pointer(other.host_pointer),
+      shared_pointer(other.shared_pointer),
+      shared_counter(other.shared_counter),
+      original_device_ptr(other.original_device_ptr),
+      original_device_size(other.original_device_size),
+      original_device(other.original_device),
+      need_realloc_(other.need_realloc_),
+      modified(other.modified)
+{
+  other.data_elements = 0;
+  other.data_size = 0;
+  other.device_size = 0;
+  other.data_width = 0;
+  other.data_height = 0;
+  other.data_depth = 0;
+  other.device = 0;
+  other.device_pointer = 0;
+  other.host_pointer = 0;
+  other.shared_pointer = 0;
+  other.shared_counter = 0;
+  other.original_device_ptr = 0;
+  other.original_device_size = 0;
+  other.original_device = 0;
+  other.need_realloc_ = false;
+  other.modified = false;
+}
+
 device_memory::~device_memory()
 {
   assert(shared_pointer == 0);

intern/cycles/device/memory.h

@@ -281,16 +281,11 @@ class device_memory {
 
   /* Only create through subclasses. */
   device_memory(Device *device, const char *name, MemoryType type);
+  device_memory(device_memory &&other) noexcept;
 
-  /* No copying and allowed.
-   *
-   * This is because device implementation might need to register device memory in an allocation
-   * map of some sort and use pointer as a key to identify blocks. Moving data from one place to
-   * another bypassing device allocation routines will make those maps hard to maintain. */
+  /* No copying allowed. */
   device_memory(const device_memory &) = delete;
-  device_memory(device_memory &&other) noexcept = delete;
   device_memory &operator=(const device_memory &) = delete;
-  device_memory &operator=(device_memory &&) = delete;
 
   /* Host allocation on the device. All host_pointer memory should be
    * allocated with these functions, for devices that support using

intern/cycles/device/optix/device_impl.cpp

@@ -44,14 +44,14 @@
 CCL_NAMESPACE_BEGIN
 
 OptiXDevice::Denoiser::Denoiser(OptiXDevice *device)
-    : device(device), queue(device), state(device, "__denoiser_state", true)
+    : device(device), queue(device), state(device, "__denoiser_state")
 {
 }
 
 OptiXDevice::OptiXDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
     : CUDADevice(info, stats, profiler),
       sbt_data(this, "__sbt", MEM_READ_ONLY),
-      launch_params(this, "__params", false),
+      launch_params(this, "__params"),
       denoiser_(this)
 {
   /* Make the CUDA context current. */
@@ -507,7 +507,7 @@ class OptiXDevice::DenoiseContext {
       : denoise_params(task.params),
         render_buffers(task.render_buffers),
         buffer_params(task.buffer_params),
-        guiding_buffer(device, "denoiser guiding passes buffer", true),
+        guiding_buffer(device, "denoiser guiding passes buffer"),
         num_samples(task.num_samples)
   {
     num_input_passes = 1;
@@ -522,9 +522,9 @@ class OptiXDevice::DenoiseContext {
       }
     }
 
-    use_guiding_passes = (num_input_passes - 1) > 0;
+    const int num_guiding_passes = num_input_passes - 1;
 
-    if (use_guiding_passes) {
+    if (num_guiding_passes) {
       if (task.allow_inplace_modification) {
         guiding_params.device_pointer = render_buffers->buffer.device_pointer;
 
@@ -577,7 +577,6 @@ class OptiXDevice::DenoiseContext {
 
   /* Number of input passes. Including the color and extra auxiliary passes. */
   int num_input_passes = 0;
-  bool use_guiding_passes = false;
   bool use_pass_albedo = false;
   bool use_pass_normal = false;
 
@@ -709,7 +708,7 @@ void OptiXDevice::denoise_pass(DenoiseContext &context, PassType pass_type)
       return;
     }
   }
-  else if (context.use_guiding_passes && !context.albedo_replaced_with_fake) {
+  else if (!context.albedo_replaced_with_fake) {
     context.albedo_replaced_with_fake = true;
     if (!denoise_filter_guiding_set_fake_albedo(context)) {
       LOG(ERROR) << "Error replacing real albedo with the fake one.";
@@ -1001,13 +1000,6 @@ bool OptiXDevice::build_optix_bvh(BVHOptiX *bvh,
                                   const OptixBuildInput &build_input,
                                   uint16_t num_motion_steps)
 {
-  /* Allocate and build acceleration structures only one at a time, to prevent parallel builds
-   * from running out of memory (since both original and compacted acceleration structure memory
-   * may be allocated at the same time for the duration of this function). The builds would
-   * otherwise happen on the same CUDA stream anyway. */
-  static thread_mutex mutex;
-  thread_scoped_lock lock(mutex);
-
   const CUDAContextScope scope(this);
 
   const bool use_fast_trace_bvh = (bvh->params.bvh_type == BVH_TYPE_STATIC);
@@ -1033,15 +1025,14 @@ bool OptiXDevice::build_optix_bvh(BVHOptiX *bvh,
   optix_assert(optixAccelComputeMemoryUsage(context, &options, &build_input, 1, &sizes));
 
   /* Allocate required output buffers. */
-  device_only_memory<char> temp_mem(this, "optix temp as build mem", true);
+  device_only_memory<char> temp_mem(this, "optix temp as build mem");
   temp_mem.alloc_to_device(align_up(sizes.tempSizeInBytes, 8) + 8);
   if (!temp_mem.device_pointer) {
     /* Make sure temporary memory allocation succeeded. */
     return false;
   }
 
-  /* Acceleration structure memory has to be allocated on the device (not allowed on the host). */
-  device_only_memory<char> &out_data = *bvh->as_data;
+  device_only_memory<char> &out_data = bvh->as_data;
   if (operation == OPTIX_BUILD_OPERATION_BUILD) {
     assert(out_data.device == this);
     out_data.alloc_to_device(sizes.outputSizeInBytes);
@@ -1089,13 +1080,12 @@ bool OptiXDevice::build_optix_bvh(BVHOptiX *bvh,
 
     /* There is no point compacting if the size does not change. */
     if (compacted_size < sizes.outputSizeInBytes) {
-      device_only_memory<char> compacted_data(this, "optix compacted as", false);
+      device_only_memory<char> compacted_data(this, "optix compacted as");
       compacted_data.alloc_to_device(compacted_size);
-      if (!compacted_data.device_pointer) {
+      if (!compacted_data.device_pointer)
         /* Do not compact if memory allocation for compacted acceleration structure fails.
          * Can just use the uncompacted one then, so succeed here regardless. */
         return !have_error();
-      }
 
       optix_assert(optixAccelCompact(
           context, NULL, out_handle, compacted_data.device_pointer, compacted_size, &out_handle));
@@ -1106,8 +1096,6 @@ bool OptiXDevice::build_optix_bvh(BVHOptiX *bvh,
 
       std::swap(out_data.device_size, compacted_data.device_size);
       std::swap(out_data.device_pointer, compacted_data.device_pointer);
-      /* Original acceleration structure memory is freed when 'compacted_data' goes out of scope.
-       */
     }
   }
 
@@ -1135,7 +1123,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
       operation = OPTIX_BUILD_OPERATION_UPDATE;
     }
     else {
-      bvh_optix->as_data->free();
+      bvh_optix->as_data.free();
       bvh_optix->traversable_handle = 0;
     }
 
@@ -1196,7 +1184,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
               const float4 pw = make_float4(
                   curve_radius[ka], curve_radius[k0], curve_radius[k1], curve_radius[kb]);
 
-              /* Convert Catmull-Rom data to B-spline. */
+              /* Convert Catmull-Rom data to Bezier spline. */
               static const float4 cr2bsp0 = make_float4(+7, -4, +5, -2) / 6.f;
               static const float4 cr2bsp1 = make_float4(-2, 11, -4, +1) / 6.f;
               static const float4 cr2bsp2 = make_float4(+1, -4, 11, -2) / 6.f;
@@ -1356,9 +1344,9 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
     unsigned int num_instances = 0;
     unsigned int max_num_instances = 0xFFFFFFFF;
 
-    bvh_optix->as_data->free();
+    bvh_optix->as_data.free();
     bvh_optix->traversable_handle = 0;
-    bvh_optix->motion_transform_data->free();
+    bvh_optix->motion_transform_data.free();
 
     optixDeviceContextGetProperty(context,
                                   OPTIX_DEVICE_PROPERTY_LIMIT_MAX_INSTANCE_ID,
@@ -1391,8 +1379,8 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         }
       }
 
-      assert(bvh_optix->motion_transform_data->device == this);
-      bvh_optix->motion_transform_data->alloc_to_device(total_motion_transform_size);
+      assert(bvh_optix->motion_transform_data.device == this);
+      bvh_optix->motion_transform_data.alloc_to_device(total_motion_transform_size);
     }
 
     for (Object *ob : bvh->objects) {
@@ -1453,7 +1441,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
 
         motion_transform_offset = align_up(motion_transform_offset,
                                            OPTIX_TRANSFORM_BYTE_ALIGNMENT);
-        CUdeviceptr motion_transform_gpu = bvh_optix->motion_transform_data->device_pointer +
+        CUdeviceptr motion_transform_gpu = bvh_optix->motion_transform_data.device_pointer +
                                            motion_transform_offset;
         motion_transform_offset += motion_transform_size;
 

intern/cycles/device/optix/device_impl.h

@@ -23,7 +23,6 @@
 #  include "device/optix/queue.h"
 #  include "device/optix/util.h"
 #  include "kernel/types.h"
-#  include "util/unique_ptr.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -77,7 +76,7 @@ class OptiXDevice : public CUDADevice {
   device_only_memory<KernelParamsOptiX> launch_params;
   OptixTraversableHandle tlas_handle = 0;
 
-  vector<unique_ptr<device_only_memory<char>>> delayed_free_bvh_memory;
+  vector<device_only_memory<char>> delayed_free_bvh_memory;
   thread_mutex delayed_free_bvh_mutex;
 
   class Denoiser {

intern/cycles/integrator/denoiser_oidn.cpp

@@ -47,6 +47,9 @@ static bool oidn_progress_monitor_function(void *user_ptr, double /*n*/)
   OIDNDenoiser *oidn_denoiser = reinterpret_cast<OIDNDenoiser *>(user_ptr);
   return !oidn_denoiser->is_cancelled();
 }
+#endif
+
+#ifdef WITH_OPENIMAGEDENOISE
 
 class OIDNPass {
  public:
@@ -544,6 +547,7 @@ class OIDNDenoiseContext {
    * the fake values and denoising of passes which do need albedo can no longer happen. */
   bool albedo_replaced_with_fake_ = false;
 };
+#endif
 
 static unique_ptr<DeviceQueue> create_device_queue(const RenderBuffers *render_buffers)
 {
@@ -578,20 +582,18 @@ static void copy_render_buffers_to_device(unique_ptr<DeviceQueue> &queue,
   }
 }
 
-#endif
-
 bool OIDNDenoiser::denoise_buffer(const BufferParams &buffer_params,
                                   RenderBuffers *render_buffers,
                                   const int num_samples,
                                   bool allow_inplace_modification)
 {
-#ifdef WITH_OPENIMAGEDENOISE
   thread_scoped_lock lock(mutex_);
 
   /* Make sure the host-side data is available for denoising. */
   unique_ptr<DeviceQueue> queue = create_device_queue(render_buffers);
   copy_render_buffers_from_device(queue, render_buffers);
 
+#ifdef WITH_OPENIMAGEDENOISE
   OIDNDenoiseContext context(
       this, params_, buffer_params, render_buffers, num_samples, allow_inplace_modification);
 
@@ -618,11 +620,6 @@ bool OIDNDenoiser::denoise_buffer(const BufferParams &buffer_params,
      * copies data from the device it doesn't overwrite the denoiser buffers. */
     copy_render_buffers_to_device(queue, render_buffers);
   }
-#else
-  (void)buffer_params;
-  (void)render_buffers;
-  (void)num_samples;
-  (void)allow_inplace_modification;
 #endif
 
   /* This code is not supposed to run when compiled without OIDN support, so can assume if we made

intern/cycles/integrator/path_trace.cpp

@@ -296,13 +296,13 @@ static BufferParams scale_buffer_params(const BufferParams &params, int resoluti
 
   scaled_params.window_x = params.window_x / resolution_divider;
   scaled_params.window_y = params.window_y / resolution_divider;
-  scaled_params.window_width = max(1, params.window_width / resolution_divider);
-  scaled_params.window_height = max(1, params.window_height / resolution_divider);
+  scaled_params.window_width = params.window_width / resolution_divider;
+  scaled_params.window_height = params.window_height / resolution_divider;
 
   scaled_params.full_x = params.full_x / resolution_divider;
   scaled_params.full_y = params.full_y / resolution_divider;
-  scaled_params.full_width = max(1, params.full_width / resolution_divider);
-  scaled_params.full_height = max(1, params.full_height / resolution_divider);
+  scaled_params.full_width = params.full_width / resolution_divider;
+  scaled_params.full_height = params.full_height / resolution_divider;
 
   scaled_params.update_offset_stride();
 
@@ -380,7 +380,10 @@ void PathTrace::path_trace(RenderWork &render_work)
     PathTraceWork *path_trace_work = path_trace_works_[i].get();
 
     PathTraceWork::RenderStatistics statistics;
-    path_trace_work->render_samples(statistics, render_work.path_trace.start_sample, num_samples);
+    path_trace_work->render_samples(statistics,
+                                    render_work.path_trace.start_sample,
+                                    num_samples,
+                                    render_work.path_trace.sample_offset);
 
     const double work_time = time_dt() - work_start_time;
     work_balance_infos_[i].time_spent += work_time;
@@ -479,11 +482,7 @@ void PathTrace::set_denoiser_params(const DenoiseParams &params)
   }
 
   denoiser_ = Denoiser::create(device_, params);
-
-  /* Only take into account the "immediate" cancel to have interactive rendering responding to
-   * navigation as quickly as possible, but allow to run denoiser after user hit Esc button while
-   * doing offline rendering. */
-  denoiser_->is_cancelled_cb = [this]() { return render_cancel_.is_requested; };
+  denoiser_->is_cancelled_cb = [this]() { return is_cancel_requested(); };
 }
 
 void PathTrace::set_adaptive_sampling(const AdaptiveSampling &adaptive_sampling)
@@ -851,10 +850,10 @@ void PathTrace::progress_update_if_needed(const RenderWork &render_work)
 {
   if (progress_ != nullptr) {
     const int2 tile_size = get_render_tile_size();
-    const uint64_t num_samples_added = uint64_t(tile_size.x) * tile_size.y *
-                                       render_work.path_trace.num_samples;
+    const int num_samples_added = tile_size.x * tile_size.y * render_work.path_trace.num_samples;
     const int current_sample = render_work.path_trace.start_sample +
-                               render_work.path_trace.num_samples;
+                               render_work.path_trace.num_samples -
+                               render_work.path_trace.sample_offset;
     progress_->add_samples(num_samples_added, current_sample);
   }
 

intern/cycles/integrator/path_trace_work.h

@@ -75,7 +75,10 @@ class PathTraceWork {
 
   /* Render given number of samples as a synchronous blocking call.
    * The samples are added to the render buffer associated with this work. */
-  virtual void render_samples(RenderStatistics &statistics, int start_sample, int samples_num) = 0;
+  virtual void render_samples(RenderStatistics &statistics,
+                              int start_sample,
+                              int samples_num,
+                              int sample_offset) = 0;
 
   /* Copy render result from this work to the corresponding place of the GPU display.
    *

intern/cycles/integrator/path_trace_work_cpu.cpp

@@ -71,7 +71,8 @@ void PathTraceWorkCPU::init_execution()
 
 void PathTraceWorkCPU::render_samples(RenderStatistics &statistics,
                                       int start_sample,
-                                      int samples_num)
+                                      int samples_num,
+                                      int sample_offset)
 {
   const int64_t image_width = effective_buffer_params_.width;
   const int64_t image_height = effective_buffer_params_.height;
@@ -99,6 +100,7 @@ void PathTraceWorkCPU::render_samples(RenderStatistics &statistics,
       work_tile.w = 1;
       work_tile.h = 1;
       work_tile.start_sample = start_sample;
+      work_tile.sample_offset = sample_offset;
       work_tile.num_samples = 1;
       work_tile.offset = effective_buffer_params_.offset;
       work_tile.stride = effective_buffer_params_.stride;

intern/cycles/integrator/path_trace_work_cpu.h

@@ -48,7 +48,8 @@ class PathTraceWorkCPU : public PathTraceWork {
 
   virtual void render_samples(RenderStatistics &statistics,
                               int start_sample,
-                              int samples_num) override;
+                              int samples_num,
+                              int sample_offset) override;
 
   virtual void copy_to_display(PathTraceDisplay *display,
                                PassMode pass_mode,

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -250,18 +250,19 @@ void PathTraceWorkGPU::init_execution()
 
 void PathTraceWorkGPU::render_samples(RenderStatistics &statistics,
                                       int start_sample,
-                                      int samples_num)
+                                      int samples_num,
+                                      int sample_offset)
 {
   /* Limit number of states for the tile and rely on a greedy scheduling of tiles. This allows to
    * add more work (because tiles are smaller, so there is higher chance that more paths will
    * become busy after adding new tiles). This is especially important for the shadow catcher which
    * schedules work in halves of available number of paths. */
   work_tile_scheduler_.set_max_num_path_states(max_num_paths_ / 8);
-  work_tile_scheduler_.set_accelerated_rt((device_->get_bvh_layout_mask() & BVH_LAYOUT_OPTIX) !=
-                                          0);
+
   work_tile_scheduler_.reset(effective_buffer_params_,
                              start_sample,
                              samples_num,
+                             sample_offset,
                              device_scene_->data.integrator.scrambling_distance);
 
   enqueue_reset();

intern/cycles/integrator/path_trace_work_gpu.h

@@ -46,7 +46,8 @@ class PathTraceWorkGPU : public PathTraceWork {
 
   virtual void render_samples(RenderStatistics &statistics,
                               int start_sample,
-                              int samples_num) override;
+                              int samples_num,
+                              int sample_offset) override;
 
   virtual void copy_to_display(PathTraceDisplay *display,
                                PassMode pass_mode,

intern/cycles/integrator/render_scheduler.cpp

@@ -88,6 +88,16 @@ int RenderScheduler::get_num_samples() const
   return num_samples_;
 }
 
+void RenderScheduler::set_sample_offset(int sample_offset)
+{
+  sample_offset_ = sample_offset;
+}
+
+int RenderScheduler::get_sample_offset() const
+{
+  return sample_offset_;
+}
+
 void RenderScheduler::set_time_limit(double time_limit)
 {
   time_limit_ = time_limit;
@@ -110,13 +120,15 @@ int RenderScheduler::get_num_rendered_samples() const
   return state_.num_rendered_samples;
 }
 
-void RenderScheduler::reset(const BufferParams &buffer_params, int num_samples)
+void RenderScheduler::reset(const BufferParams &buffer_params, int num_samples, int sample_offset)
 {
   buffer_params_ = buffer_params;
 
   update_start_resolution_divider();
 
   set_num_samples(num_samples);
+  set_start_sample(sample_offset);
+  set_sample_offset(sample_offset);
 
   /* In background mode never do lower resolution render preview, as it is not really supported
    * by the software. */
@@ -171,7 +183,7 @@ void RenderScheduler::reset(const BufferParams &buffer_params, int num_samples)
 
 void RenderScheduler::reset_for_next_tile()
 {
-  reset(buffer_params_, num_samples_);
+  reset(buffer_params_, num_samples_, sample_offset_);
 }
 
 bool RenderScheduler::render_work_reschedule_on_converge(RenderWork &render_work)
@@ -317,6 +329,7 @@ RenderWork RenderScheduler::get_render_work()
 
   render_work.path_trace.start_sample = get_start_sample_to_path_trace();
   render_work.path_trace.num_samples = get_num_samples_to_path_trace();
+  render_work.path_trace.sample_offset = get_sample_offset();
 
   render_work.init_render_buffers = (render_work.path_trace.start_sample == get_start_sample());
 
@@ -827,26 +840,6 @@ int RenderScheduler::get_num_samples_to_path_trace() const
       num_samples_to_occupy = lround(state_.occupancy_num_samples * 0.7f / state_.occupancy);
     }
 
-    /* When time limit is used clamp the calculated number of samples to keep occupancy.
-     * This is because time limit causes the last render iteration to happen with less number of
-     * samples, which conflicts with the occupancy (lower number of samples causes lower
-     * occupancy, also the calculation is based on number of previously rendered samples).
-     *
-     * When time limit is not used the number of samples per render iteration is either increasing
-     * or stays the same, so there is no need to clamp number of samples calculated for occupancy.
-     */
-    if (time_limit_ && state_.start_render_time) {
-      const double remaining_render_time = max(
-          0.0, time_limit_ - (time_dt() - state_.start_render_time));
-      const double time_per_sample_average = path_trace_time_.get_average();
-      const double predicted_render_time = num_samples_to_occupy * time_per_sample_average;
-
-      if (predicted_render_time > remaining_render_time) {
-        num_samples_to_occupy = lround(num_samples_to_occupy *
-                                       (remaining_render_time / predicted_render_time));
-      }
-    }
-
     num_samples_to_render = max(num_samples_to_render,
                                 min(num_samples_to_occupy, max_num_samples_to_render));
   }

intern/cycles/integrator/render_scheduler.h

@@ -39,6 +39,7 @@ class RenderWork {
   struct {
     int start_sample = 0;
     int num_samples = 0;
+    int sample_offset = 0;
   } path_trace;
 
   struct {
@@ -125,6 +126,9 @@ class RenderScheduler {
   void set_num_samples(int num_samples);
   int get_num_samples() const;
 
+  void set_sample_offset(int sample_offset);
+  int get_sample_offset() const;
+
   /* Time limit for the path tracing tasks, in minutes.
    * Zero disables the limit. */
   void set_time_limit(double time_limit);
@@ -150,7 +154,7 @@ class RenderScheduler {
 
   /* Reset scheduler, indicating that rendering will happen from scratch.
    * Resets current rendered state, as well as scheduling information. */
-  void reset(const BufferParams &buffer_params, int num_samples);
+  void reset(const BufferParams &buffer_params, int num_samples, int sample_offset);
 
   /* Reset scheduler upon switching to a next tile.
    * Will keep the same number of samples and full-frame render parameters, but will reset progress
@@ -419,6 +423,8 @@ class RenderScheduler {
   int start_sample_ = 0;
   int num_samples_ = 0;
 
+  int sample_offset_ = 0;
+
   /* Limit in seconds for how long path tracing is allowed to happen.
    * Zero means no limit is applied. */
   double time_limit_ = 0.0;

intern/cycles/integrator/tile.cpp

@@ -46,8 +46,7 @@ ccl_device_inline uint round_up_to_power_of_two(uint x)
   return next_power_of_two(x);
 }
 
-TileSize tile_calculate_best_size(const bool accel_rt,
-                                  const int2 &image_size,
+TileSize tile_calculate_best_size(const int2 &image_size,
                                   const int num_samples,
                                   const int max_num_path_states,
                                   const float scrambling_distance)
@@ -74,7 +73,7 @@ TileSize tile_calculate_best_size(const bool accel_rt,
 
   TileSize tile_size;
   const int num_path_states_per_sample = max_num_path_states / num_samples;
-  if (scrambling_distance < 0.9f && accel_rt) {
+  if (scrambling_distance < 0.9f) {
     /* Prefer large tiles for scrambling distance, bounded by max num path states. */
     tile_size.width = min(image_size.x, max_num_path_states);
     tile_size.height = min(image_size.y, max(max_num_path_states / tile_size.width, 1));

intern/cycles/integrator/tile.h

@@ -49,8 +49,7 @@ std::ostream &operator<<(std::ostream &os, const TileSize &tile_size);
  * of active path states.
  * Will attempt to provide best guess to keep path tracing threads of a device as localized as
  * possible, and have as many threads active for every tile as possible. */
-TileSize tile_calculate_best_size(const bool accel_rt,
-                                  const int2 &image_size,
+TileSize tile_calculate_best_size(const int2 &image_size,
                                   const int num_samples,
                                   const int max_num_path_states,
                                   const float scrambling_distance);

intern/cycles/integrator/work_tile_scheduler.cpp

@@ -28,11 +28,6 @@ WorkTileScheduler::WorkTileScheduler()
 {
 }
 
-void WorkTileScheduler::set_accelerated_rt(bool accelerated_rt)
-{
-  accelerated_rt_ = accelerated_rt;
-}
-
 void WorkTileScheduler::set_max_num_path_states(int max_num_path_states)
 {
   max_num_path_states_ = max_num_path_states;
@@ -41,6 +36,7 @@ void WorkTileScheduler::set_max_num_path_states(int max_num_path_states)
 void WorkTileScheduler::reset(const BufferParams &buffer_params,
                               int sample_start,
                               int samples_num,
+                              int sample_offset,
                               float scrambling_distance)
 {
   /* Image buffer parameters. */
@@ -56,6 +52,7 @@ void WorkTileScheduler::reset(const BufferParams &buffer_params,
   /* Samples parameters. */
   sample_start_ = sample_start;
   samples_num_ = samples_num;
+  sample_offset_ = sample_offset;
 
   /* Initialize new scheduling. */
   reset_scheduler_state();
@@ -64,7 +61,7 @@ void WorkTileScheduler::reset(const BufferParams &buffer_params,
 void WorkTileScheduler::reset_scheduler_state()
 {
   tile_size_ = tile_calculate_best_size(
-      accelerated_rt_, image_size_px_, samples_num_, max_num_path_states_, scrambling_distance_);
+      image_size_px_, samples_num_, max_num_path_states_, scrambling_distance_);
 
   VLOG(3) << "Will schedule tiles of size " << tile_size_;
 
@@ -116,6 +113,7 @@ bool WorkTileScheduler::get_work(KernelWorkTile *work_tile_, const int max_work_
   work_tile.h = tile_size_.height;
   work_tile.start_sample = sample_start_ + start_sample;
   work_tile.num_samples = min(tile_size_.num_samples, samples_num_ - start_sample);
+  work_tile.sample_offset = sample_offset_;
   work_tile.offset = offset_;
   work_tile.stride = stride_;
 

intern/cycles/integrator/work_tile_scheduler.h

@@ -31,9 +31,6 @@ class WorkTileScheduler {
  public:
   WorkTileScheduler();
 
-  /* To indicate if there is accelerated RT support. */
-  void set_accelerated_rt(bool state);
-
   /* MAximum path states which are allowed to be used by a single scheduled work tile.
    *
    * Affects the scheduled work size: the work size will be as big as possible, but will not exceed
@@ -44,6 +41,7 @@ class WorkTileScheduler {
   void reset(const BufferParams &buffer_params,
              int sample_start,
              int samples_num,
+             int sample_offset,
              float scrambling_distance);
 
   /* Get work for a device.
@@ -57,9 +55,6 @@ class WorkTileScheduler {
  protected:
   void reset_scheduler_state();
 
-  /* Used to indicate if there is accelerated ray tracing. */
-  bool accelerated_rt_ = false;
-
   /* Maximum allowed path states to be used.
    *
    * TODO(sergey): Naming can be improved. The fact that this is a limiting factor based on the
@@ -85,6 +80,7 @@ class WorkTileScheduler {
    * (splitting into a smaller work tiles). */
   int sample_start_ = 0;
   int samples_num_ = 0;
+  int sample_offset_ = 0;
 
   /* Tile size which be scheduled for rendering. */
   TileSize tile_size_;

intern/cycles/kernel/CMakeLists.txt

@@ -39,6 +39,10 @@ set(SRC_KERNEL_DEVICE_HIP
   device/hip/kernel.cpp
 )
 
+set(SRC_KERNEL_DEVICE_METAL
+  device/metal/kernel.metal
+)
+
 set(SRC_KERNEL_DEVICE_OPTIX
   device/optix/kernel.cu
   device/optix/kernel_shader_raytrace.cu
@@ -79,6 +83,13 @@ set(SRC_KERNEL_DEVICE_OPTIX_HEADERS
   device/optix/globals.h
 )
 
+set(SRC_KERNEL_DEVICE_METAL_HEADERS
+  device/metal/compat.h
+  device/metal/context_begin.h
+  device/metal/context_end.h
+  device/metal/globals.h
+)
+
 set(SRC_KERNEL_CLOSURE_HEADERS
   closure/alloc.h
   closure/bsdf.h
@@ -565,12 +576,6 @@ if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
       set(name ${name}_experimental)
     endif()
 
-    if(WITH_NANOVDB)
-      set(hip_flags ${hip_flags}
-        -D WITH_NANOVDB
-        -I "${NANOVDB_INCLUDE_DIR}")
-    endif()
-
     if(WITH_CYCLES_DEBUG)
       set(hip_flags ${hip_flags} -D __KERNEL_DEBUG__)
     endif()
@@ -729,12 +734,14 @@ cycles_add_library(cycles_kernel "${LIB}"
   ${SRC_KERNEL_DEVICE_CUDA}
   ${SRC_KERNEL_DEVICE_HIP}
   ${SRC_KERNEL_DEVICE_OPTIX}
+  ${SRC_KERNEL_DEVICE_METAL}
   ${SRC_KERNEL_HEADERS}
   ${SRC_KERNEL_DEVICE_CPU_HEADERS}
   ${SRC_KERNEL_DEVICE_GPU_HEADERS}
   ${SRC_KERNEL_DEVICE_CUDA_HEADERS}
   ${SRC_KERNEL_DEVICE_HIP_HEADERS}
   ${SRC_KERNEL_DEVICE_OPTIX_HEADERS}
+  ${SRC_KERNEL_DEVICE_METAL_HEADERS}
 )
 
 source_group("bake" FILES ${SRC_KERNEL_BAKE_HEADERS})
@@ -746,6 +753,7 @@ source_group("device\\cuda" FILES ${SRC_KERNEL_DEVICE_CUDA} ${SRC_KERNEL_DEVICE_
 source_group("device\\gpu" FILES ${SRC_KERNEL_DEVICE_GPU_HEADERS})
 source_group("device\\hip" FILES ${SRC_KERNEL_DEVICE_HIP} ${SRC_KERNEL_DEVICE_HIP_HEADERS})
 source_group("device\\optix" FILES ${SRC_KERNEL_DEVICE_OPTIX} ${SRC_KERNEL_DEVICE_OPTIX_HEADERS})
+source_group("device\\metal" FILES ${SRC_KERNEL_DEVICE_METAL} ${SRC_KERNEL_DEVICE_METAL_HEADERS})
 source_group("film" FILES ${SRC_KERNEL_FILM_HEADERS})
 source_group("geom" FILES ${SRC_KERNEL_GEOM_HEADERS})
 source_group("integrator" FILES ${SRC_KERNEL_INTEGRATOR_HEADERS})
@@ -778,6 +786,8 @@ delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_HIP}" ${CYCLES_
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_HIP_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/hip)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_OPTIX}" ${CYCLES_INSTALL_PATH}/source/kernel/device/optix)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_OPTIX_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/optix)
+delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_METAL}" ${CYCLES_INSTALL_PATH}/source/kernel/device/metal)
+delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_METAL_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/metal)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_FILM_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/film)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_GEOM_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/geom)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_INTEGRATOR_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/integrator)

intern/cycles/kernel/closure/bsdf.h

@@ -438,7 +438,7 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
   if (label & LABEL_TRANSMIT) {
     float threshold_squared = kernel_data.background.transparent_roughness_squared_threshold;
 
-    if (threshold_squared >= 0.0f && !(label & LABEL_DIFFUSE)) {
+    if (threshold_squared >= 0.0f) {
       if (bsdf_get_specular_roughness_squared(sc) <= threshold_squared) {
         label |= LABEL_TRANSMIT_TRANSPARENT;
       }

intern/cycles/kernel/device/cuda/compat.h

@@ -75,6 +75,7 @@ typedef unsigned long long uint64_t;
 #define ccl_gpu_block_idx_x (blockIdx.x)
 #define ccl_gpu_grid_dim_x (gridDim.x)
 #define ccl_gpu_warp_size (warpSize)
+#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
 
 #define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
 #define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)

intern/cycles/kernel/device/cuda/config.h

@@ -92,12 +92,29 @@
 
 /* Compute number of threads per block and minimum blocks per multiprocessor
  * given the maximum number of registers per thread. */
-
 #define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
   extern "C" __global__ void __launch_bounds__(block_num_threads, \
                                                GPU_MULTIPRESSOR_MAX_REGISTERS / \
                                                    (block_num_threads * thread_num_registers))
 
+#define ccl_gpu_kernel_threads(block_num_threads) \
+  extern "C" __global__ void __launch_bounds__(block_num_threads)
+
+#define ccl_gpu_kernel_signature(name, ...) kernel_gpu_##name(__VA_ARGS__)
+
+#define ccl_gpu_kernel_call(x) x
+
+/* Define a function object where "func" is the lambda body, and additional parameters are used to
+ * specify captured state  */
+#define ccl_gpu_kernel_lambda(func, ...) \
+  struct KernelLambda { \
+    __VA_ARGS__; \
+    __device__ int operator()(const int state) \
+    { \
+      return (func); \
+    } \
+  } ccl_gpu_kernel_lambda_pass
+
 /* sanity checks */
 
 #if GPU_KERNEL_BLOCK_NUM_THREADS > GPU_BLOCK_MAX_THREADS

intern/cycles/kernel/device/gpu/image.h

@@ -65,7 +65,9 @@ ccl_device float cubic_h1(float a)
 
 /* Fast bicubic texture lookup using 4 bilinear lookups, adapted from CUDA samples. */
 template<typename T>
-ccl_device_noinline T kernel_tex_image_interp_bicubic(const TextureInfo &info, float x, float y)
+ccl_device_noinline T kernel_tex_image_interp_bicubic(ccl_global const TextureInfo &info,
+                                                      float x,
+                                                      float y)
 {
   ccl_gpu_tex_object tex = (ccl_gpu_tex_object)info.data;
 
@@ -94,7 +96,7 @@ ccl_device_noinline T kernel_tex_image_interp_bicubic(const TextureInfo &info, f
 /* Fast tricubic texture lookup using 8 trilinear lookups. */
 template<typename T>
 ccl_device_noinline T
-kernel_tex_image_interp_tricubic(const TextureInfo &info, float x, float y, float z)
+kernel_tex_image_interp_tricubic(ccl_global const TextureInfo &info, float x, float y, float z)
 {
   ccl_gpu_tex_object tex = (ccl_gpu_tex_object)info.data;
 
@@ -169,7 +171,7 @@ ccl_device T kernel_tex_image_interp_tricubic_nanovdb(S &s, float x, float y, fl
 
 template<typename T>
 ccl_device_noinline T kernel_tex_image_interp_nanovdb(
-    const TextureInfo &info, float x, float y, float z, uint interpolation)
+    ccl_global const TextureInfo &info, float x, float y, float z, uint interpolation)
 {
   using namespace nanovdb;
 
@@ -191,7 +193,7 @@ ccl_device_noinline T kernel_tex_image_interp_nanovdb(
 
 ccl_device float4 kernel_tex_image_interp(KernelGlobals kg, int id, float x, float y)
 {
-  const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
+  ccl_global const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
 
   /* float4, byte4, ushort4 and half4 */
   const int texture_type = info.data_type;
@@ -226,7 +228,7 @@ ccl_device float4 kernel_tex_image_interp_3d(KernelGlobals kg,
                                              float3 P,
                                              InterpolationType interp)
 {
-  const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
+  ccl_global const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
 
   if (info.use_transform_3d) {
     P = transform_point(&info.transform_3d, P);

intern/cycles/kernel/device/gpu/parallel_active_index.h

@@ -31,10 +31,43 @@ CCL_NAMESPACE_BEGIN
 #  define GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE 512
 #endif
 
+#ifdef __KERNEL_METAL__
+struct ActiveIndexContext {
+  ActiveIndexContext(int _thread_index,
+                     int _global_index,
+                     int _threadgroup_size,
+                     int _simdgroup_size,
+                     int _simd_lane_index,
+                     int _simd_group_index,
+                     int _num_simd_groups,
+                     threadgroup int *_simdgroup_offset)
+      : thread_index(_thread_index),
+        global_index(_global_index),
+        blocksize(_threadgroup_size),
+        ccl_gpu_warp_size(_simdgroup_size),
+        thread_warp(_simd_lane_index),
+        warp_index(_simd_group_index),
+        num_warps(_num_simd_groups),
+        warp_offset(_simdgroup_offset)
+  {
+  }
+
+  const int thread_index, global_index, blocksize, ccl_gpu_warp_size, thread_warp, warp_index,
+      num_warps;
+  threadgroup int *warp_offset;
+
+  template<uint blocksizeDummy, typename IsActiveOp>
+  void active_index_array(const uint num_states,
+                          ccl_global int *indices,
+                          ccl_global int *num_indices,
+                          IsActiveOp is_active_op)
+  {
+    const uint state_index = global_index;
+#else
 template<uint blocksize, typename IsActiveOp>
 __device__ void gpu_parallel_active_index_array(const uint num_states,
-                                                int *indices,
-                                                int *num_indices,
+                                                ccl_global int *indices,
+                                                ccl_global int *num_indices,
                                                 IsActiveOp is_active_op)
 {
   extern ccl_gpu_shared int warp_offset[];
@@ -45,43 +78,62 @@ __device__ void gpu_parallel_active_index_array(const uint num_states,
   const uint warp_index = thread_index / ccl_gpu_warp_size;
   const uint num_warps = blocksize / ccl_gpu_warp_size;
 
-  /* Test if state corresponding to this thread is active. */
   const uint state_index = ccl_gpu_block_idx_x * blocksize + thread_index;
-  const uint is_active = (state_index < num_states) ? is_active_op(state_index) : 0;
+#endif
 
-  /* For each thread within a warp compute how many other active states precede it. */
-  const uint thread_mask = 0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp);
-  const uint thread_offset = ccl_gpu_popc(ccl_gpu_ballot(is_active) & thread_mask);
+    /* Test if state corresponding to this thread is active. */
+    const uint is_active = (state_index < num_states) ? is_active_op(state_index) : 0;
 
-  /* Last thread in warp stores number of active states for each warp. */
-  if (thread_warp == ccl_gpu_warp_size - 1) {
-    warp_offset[warp_index] = thread_offset + is_active;
-  }
+    /* For each thread within a warp compute how many other active states precede it. */
+    const uint thread_offset = ccl_gpu_popc(ccl_gpu_ballot(is_active) &
+                                            ccl_gpu_thread_mask(thread_warp));
 
-  ccl_gpu_syncthreads();
-
-  /* Last thread in block converts per-warp sizes to offsets, increments global size of
-   * index array and gets offset to write to. */
-  if (thread_index == blocksize - 1) {
-    /* TODO: parallelize this. */
-    int offset = 0;
-    for (int i = 0; i < num_warps; i++) {
-      int num_active = warp_offset[i];
-      warp_offset[i] = offset;
-      offset += num_active;
+    /* Last thread in warp stores number of active states for each warp. */
+    if (thread_warp == ccl_gpu_warp_size - 1) {
+      warp_offset[warp_index] = thread_offset + is_active;
     }
 
-    const uint block_num_active = warp_offset[warp_index] + thread_offset + is_active;
-    warp_offset[num_warps] = atomic_fetch_and_add_uint32(num_indices, block_num_active);
+    ccl_gpu_syncthreads();
+
+    /* Last thread in block converts per-warp sizes to offsets, increments global size of
+     * index array and gets offset to write to. */
+    if (thread_index == blocksize - 1) {
+      /* TODO: parallelize this. */
+      int offset = 0;
+      for (int i = 0; i < num_warps; i++) {
+        int num_active = warp_offset[i];
+        warp_offset[i] = offset;
+        offset += num_active;
+      }
+
+      const uint block_num_active = warp_offset[warp_index] + thread_offset + is_active;
+      warp_offset[num_warps] = atomic_fetch_and_add_uint32(num_indices, block_num_active);
+    }
+
+    ccl_gpu_syncthreads();
+
+    /* Write to index array. */
+    if (is_active) {
+      const uint block_offset = warp_offset[num_warps];
+      indices[block_offset + warp_offset[warp_index] + thread_offset] = state_index;
+    }
   }
 
-  ccl_gpu_syncthreads();
+#ifdef __KERNEL_METAL__
+}; /* end class ActiveIndexContext */
 
-  /* Write to index array. */
-  if (is_active) {
-    const uint block_offset = warp_offset[num_warps];
-    indices[block_offset + warp_offset[warp_index] + thread_offset] = state_index;
-  }
-}
+/* inject the required thread params into a struct, and redirect to its templated member function
+ */
+#  define gpu_parallel_active_index_array \
+    ActiveIndexContext(metal_local_id, \
+                       metal_global_id, \
+                       metal_local_size, \
+                       simdgroup_size, \
+                       simd_lane_index, \
+                       simd_group_index, \
+                       num_simd_groups, \
+                       simdgroup_offset) \
+        .active_index_array
+#endif
 
 CCL_NAMESPACE_END

intern/cycles/kernel/device/gpu/parallel_prefix_sum.h

@@ -33,10 +33,12 @@ CCL_NAMESPACE_BEGIN
 #  define GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE 512
 #endif
 
-template<uint blocksize>
-__device__ void gpu_parallel_prefix_sum(int *counter, int *prefix_sum, const int num_values)
+__device__ void gpu_parallel_prefix_sum(const int global_id,
+                                        ccl_global int *counter,
+                                        ccl_global int *prefix_sum,
+                                        const int num_values)
 {
-  if (!(ccl_gpu_block_idx_x == 0 && ccl_gpu_thread_idx_x == 0)) {
+  if (global_id != 0) {
     return;
   }
 

intern/cycles/kernel/device/gpu/parallel_sorted_index.h

@@ -33,16 +33,16 @@ CCL_NAMESPACE_BEGIN
 #endif
 #define GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY (~0)
 
-template<uint blocksize, typename GetKeyOp>
-__device__ void gpu_parallel_sorted_index_array(const uint num_states,
+template<typename GetKeyOp>
+__device__ void gpu_parallel_sorted_index_array(const uint state_index,
+                                                const uint num_states,
                                                 const int num_states_limit,
-                                                int *indices,
-                                                int *num_indices,
-                                                int *key_counter,
-                                                int *key_prefix_sum,
+                                                ccl_global int *indices,
+                                                ccl_global int *num_indices,
+                                                ccl_global int *key_counter,
+                                                ccl_global int *key_prefix_sum,
                                                 GetKeyOp get_key_op)
 {
-  const uint state_index = ccl_gpu_block_idx_x * blocksize + ccl_gpu_thread_idx_x;
   const int key = (state_index < num_states) ? get_key_op(state_index) :
                                                GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY;
 

intern/cycles/kernel/device/gpu/work_stealing.h

@@ -29,20 +29,17 @@ ccl_device_inline void get_work_pixel(ccl_global const KernelWorkTile *tile,
                                       ccl_private uint *y,
                                       ccl_private uint *sample)
 {
-  uint sample_offset, pixel_offset;
-
-  if (kernel_data.integrator.scrambling_distance < 0.9f) {
-    /* Keep threads for the same sample together. */
-    uint tile_pixels = tile->w * tile->h;
-    sample_offset = global_work_index / tile_pixels;
-    pixel_offset = global_work_index - sample_offset * tile_pixels;
-  }
-  else {
-    /* Keeping threads for the same pixel together.
-     * Appears to improve performance by a few % on CUDA and OptiX. */
-    sample_offset = global_work_index % tile->num_samples;
-    pixel_offset = global_work_index / tile->num_samples;
-  }
+#if 0
+  /* Keep threads for the same sample together. */
+  uint tile_pixels = tile->w * tile->h;
+  uint sample_offset = global_work_index / tile_pixels;
+  uint pixel_offset = global_work_index - sample_offset * tile_pixels;
+#else
+  /* Keeping threads for the same pixel together.
+   * Appears to improve performance by a few % on CUDA and OptiX. */
+  uint sample_offset = global_work_index % tile->num_samples;
+  uint pixel_offset = global_work_index / tile->num_samples;
+#endif
 
   uint y_offset = pixel_offset / tile->w;
   uint x_offset = pixel_offset - y_offset * tile->w;

intern/cycles/kernel/device/hip/compat.h

@@ -74,6 +74,7 @@ typedef unsigned long long uint64_t;
 #define ccl_gpu_block_idx_x (blockIdx.x)
 #define ccl_gpu_grid_dim_x (gridDim.x)
 #define ccl_gpu_warp_size (warpSize)
+#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
 
 #define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
 #define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)

intern/cycles/kernel/device/hip/config.h

@@ -35,12 +35,29 @@
 
 /* Compute number of threads per block and minimum blocks per multiprocessor
  * given the maximum number of registers per thread. */
-
 #define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
   extern "C" __global__ void __launch_bounds__(block_num_threads, \
                                                GPU_MULTIPRESSOR_MAX_REGISTERS / \
                                                    (block_num_threads * thread_num_registers))
 
+#define ccl_gpu_kernel_threads(block_num_threads) \
+  extern "C" __global__ void __launch_bounds__(block_num_threads)
+
+#define ccl_gpu_kernel_signature(name, ...) kernel_gpu_##name(__VA_ARGS__)
+
+#define ccl_gpu_kernel_call(x) x
+
+/* Define a function object where "func" is the lambda body, and additional parameters are used to
+ * specify captured state  */
+#define ccl_gpu_kernel_lambda(func, ...) \
+  struct KernelLambda { \
+    __VA_ARGS__; \
+    __device__ int operator()(const int state) \
+    { \
+      return (func); \
+    } \
+  } ccl_gpu_kernel_lambda_pass
+
 /* sanity checks */
 
 #if GPU_KERNEL_BLOCK_NUM_THREADS > GPU_BLOCK_MAX_THREADS

intern/cycles/kernel/device/metal/compat.h

@@ -58,6 +58,98 @@ using namespace metal;
 
 #define kernel_assert(cond)
 
+#define ccl_gpu_global_id_x() metal_global_id
+#define ccl_gpu_warp_size simdgroup_size
+#define ccl_gpu_thread_idx_x simd_group_index
+#define ccl_gpu_thread_mask(thread_warp) uint64_t((1ull << thread_warp) - 1)
+
+#define ccl_gpu_ballot(predicate) ((uint64_t)((simd_vote::vote_t)simd_ballot(predicate)))
+#define ccl_gpu_popc(x) popcount(x)
+
+// clang-format off
+
+/* kernel.h adapters */
+
+#define ccl_gpu_kernel(block_num_threads, thread_num_registers)
+#define ccl_gpu_kernel_threads(block_num_threads)
+
+/* Convert a comma-separated list into a semicolon-separated list
+ * (so that we can generate a struct based on kernel entry-point parameters). */
+#define FN0()
+#define FN1(p1) p1;
+#define FN2(p1, p2) p1; p2;
+#define FN3(p1, p2, p3) p1; p2; p3;
+#define FN4(p1, p2, p3, p4) p1; p2; p3; p4;
+#define FN5(p1, p2, p3, p4, p5) p1; p2; p3; p4; p5;
+#define FN6(p1, p2, p3, p4, p5, p6) p1; p2; p3; p4; p5; p6;
+#define FN7(p1, p2, p3, p4, p5, p6, p7) p1; p2; p3; p4; p5; p6; p7;
+#define FN8(p1, p2, p3, p4, p5, p6, p7, p8) p1; p2; p3; p4; p5; p6; p7; p8;
+#define FN9(p1, p2, p3, p4, p5, p6, p7, p8, p9) p1; p2; p3; p4; p5; p6; p7; p8; p9;
+#define FN10(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10;
+#define FN11(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11;
+#define FN12(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12;
+#define FN13(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13;
+#define FN14(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14;
+#define FN15(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15;
+#define FN16(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16;
+#define GET_LAST_ARG(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, ...) p16
+#define PARAMS_MAKER(...) GET_LAST_ARG(__VA_ARGS__, FN16, FN15, FN14, FN13, FN12, FN11, FN10, FN9, FN8, FN7, FN6, FN5, FN4, FN3, FN2, FN1, FN0)
+
+/* Generate a struct containing the entry-point parameters and a "run"
+ * method which can access them implicitly via this-> */
+#define ccl_gpu_kernel_signature(name, ...) \
+struct kernel_gpu_##name \
+{ \
+  PARAMS_MAKER(__VA_ARGS__)(__VA_ARGS__) \
+  void run(thread MetalKernelContext& context, \
+           threadgroup int *simdgroup_offset, \
+           const uint metal_global_id, \
+           const ushort metal_local_id, \
+           const ushort metal_local_size, \
+           uint simdgroup_size, \
+           uint simd_lane_index, \
+           uint simd_group_index, \
+           uint num_simd_groups) ccl_global const; \
+}; \
+kernel void kernel_metal_##name(device const kernel_gpu_##name *params_struct, \
+                                constant KernelParamsMetal &ccl_restrict   _launch_params_metal, \
+                                constant MetalAncillaries *_metal_ancillaries, \
+                                threadgroup int *simdgroup_offset[[ threadgroup(0) ]], \
+                                const uint metal_global_id [[thread_position_in_grid]], \
+                                const ushort metal_local_id   [[thread_position_in_threadgroup]], \
+                                const ushort metal_local_size [[threads_per_threadgroup]], \
+                                uint simdgroup_size [[threads_per_simdgroup]], \
+                                uint simd_lane_index [[thread_index_in_simdgroup]], \
+                                uint simd_group_index [[simdgroup_index_in_threadgroup]], \
+                                uint num_simd_groups [[simdgroups_per_threadgroup]]) { \
+  MetalKernelContext context(_launch_params_metal, _metal_ancillaries); \
+  INIT_DEBUG_BUFFER \
+  params_struct->run(context, simdgroup_offset, metal_global_id, metal_local_id, metal_local_size, simdgroup_size, simd_lane_index, simd_group_index, num_simd_groups); \
+} \
+void kernel_gpu_##name::run(thread MetalKernelContext& context, \
+                  threadgroup int *simdgroup_offset, \
+                  const uint metal_global_id, \
+                  const ushort metal_local_id, \
+                  const ushort metal_local_size, \
+                  uint simdgroup_size, \
+                  uint simd_lane_index, \
+                  uint simd_group_index, \
+                  uint num_simd_groups) ccl_global const
+
+#define ccl_gpu_kernel_call(x) context.x
+
+/* define a function object where "func" is the lambda body, and additional parameters are used to specify captured state  */
+#define ccl_gpu_kernel_lambda(func, ...) \
+  struct KernelLambda \
+  { \
+    KernelLambda(ccl_private MetalKernelContext &_context) : context(_context) {} \
+    ccl_private MetalKernelContext &context; \
+    __VA_ARGS__; \
+    int operator()(const int state) const { return (func); } \
+  } ccl_gpu_kernel_lambda_pass(context)
+
+// clang-format on
+
 /* make_type definitions with Metal style element initializers */
 #ifdef make_float2
 #  undef make_float2
@@ -124,3 +216,38 @@ using namespace metal;
 #define logf(x) trigmode::log(float(x))
 
 #define NULL 0
+
+/* texture bindings and sampler setup */
+
+struct Texture2DParamsMetal {
+  texture2d<float, access::sample> tex;
+};
+struct Texture3DParamsMetal {
+  texture3d<float, access::sample> tex;
+};
+
+struct MetalAncillaries {
+  device Texture2DParamsMetal *textures_2d;
+  device Texture3DParamsMetal *textures_3d;
+};
+
+enum SamplerType {
+  SamplerFilterNearest_AddressRepeat,
+  SamplerFilterNearest_AddressClampEdge,
+  SamplerFilterNearest_AddressClampZero,
+
+  SamplerFilterLinear_AddressRepeat,
+  SamplerFilterLinear_AddressClampEdge,
+  SamplerFilterLinear_AddressClampZero,
+
+  SamplerCount
+};
+
+constant constexpr array<sampler, SamplerCount> metal_samplers = {
+    sampler(address::repeat, filter::nearest),
+    sampler(address::clamp_to_edge, filter::nearest),
+    sampler(address::clamp_to_zero, filter::nearest),
+    sampler(address::repeat, filter::linear),
+    sampler(address::clamp_to_edge, filter::linear),
+    sampler(address::clamp_to_zero, filter::linear),
+};

intern/cycles/kernel/device/metal/context_begin.h

@@ -0,0 +1,79 @@
+/*
+ * Copyright 2021 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+// clang-format off
+
+/* Open the Metal kernel context class
+ * Necessary to access resource bindings */
+class MetalKernelContext {
+  public:
+    constant KernelParamsMetal &launch_params_metal;
+    constant MetalAncillaries *metal_ancillaries;
+
+    MetalKernelContext(constant KernelParamsMetal &_launch_params_metal, constant MetalAncillaries * _metal_ancillaries)
+    : launch_params_metal(_launch_params_metal), metal_ancillaries(_metal_ancillaries)
+    {}
+
+    /* texture fetch adapter functions */
+    typedef uint64_t ccl_gpu_tex_object;
+
+    template<typename T>
+    inline __attribute__((__always_inline__))
+    T ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
+      kernel_assert(0);
+      return 0;
+    }
+    template<typename T>
+    inline __attribute__((__always_inline__))
+    T ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
+      kernel_assert(0);
+      return 0;
+    }
+
+    // texture2d
+    template<>
+    inline __attribute__((__always_inline__))
+    float4 ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
+      const uint tid(tex);
+      const uint sid(tex >> 32);
+      return metal_ancillaries->textures_2d[tid].tex.sample(metal_samplers[sid], float2(x, y));
+    }
+    template<>
+    inline __attribute__((__always_inline__))
+    float ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
+      const uint tid(tex);
+      const uint sid(tex >> 32);
+      return metal_ancillaries->textures_2d[tid].tex.sample(metal_samplers[sid], float2(x, y)).x;
+    }
+
+    // texture3d
+    template<>
+    inline __attribute__((__always_inline__))
+    float4 ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
+      const uint tid(tex);
+      const uint sid(tex >> 32);
+      return metal_ancillaries->textures_3d[tid].tex.sample(metal_samplers[sid], float3(x, y, z));
+    }
+    template<>
+    inline __attribute__((__always_inline__))
+    float ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
+      const uint tid(tex);
+      const uint sid(tex >> 32);
+      return metal_ancillaries->textures_3d[tid].tex.sample(metal_samplers[sid], float3(x, y, z)).x;
+    }
+#    include "kernel/device/gpu/image.h"
+
+  // clang-format on

intern/cycles/kernel/device/metal/context_end.h

@@ -0,0 +1,23 @@
+/*
+ * Copyright 2021 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+}
+; /* end of MetalKernelContext class definition */
+
+/* Silently redirect into the MetalKernelContext instance */
+/* NOTE: These macros will need maintaining as entry-points change. */
+
+#undef kernel_integrator_state
+#define kernel_integrator_state context.launch_params_metal.__integrator_state

intern/cycles/kernel/device/metal/globals.h

@@ -0,0 +1,51 @@
+/*
+ * Copyright 2021 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+/* Constant Globals */
+
+#include "kernel/types.h"
+#include "kernel/util/profiling.h"
+
+#include "kernel/integrator/state.h"
+
+CCL_NAMESPACE_BEGIN
+
+typedef struct KernelParamsMetal {
+
+#define KERNEL_TEX(type, name) ccl_constant type *name;
+#include "kernel/textures.h"
+#undef KERNEL_TEX
+
+  const IntegratorStateGPU __integrator_state;
+  const KernelData data;
+
+} KernelParamsMetal;
+
+typedef struct KernelGlobalsGPU {
+  int unused[1];
+} KernelGlobalsGPU;
+
+typedef ccl_global const KernelGlobalsGPU *ccl_restrict KernelGlobals;
+
+#define kernel_data launch_params_metal.data
+#define kernel_integrator_state launch_params_metal.__integrator_state
+
+/* data lookup defines */
+
+#define kernel_tex_fetch(tex, index) launch_params_metal.tex[index]
+#define kernel_tex_array(tex) launch_params_metal.tex
+
+CCL_NAMESPACE_END

intern/cycles/kernel/device/metal/kernel.metal

@@ -0,0 +1,25 @@
+/*
+ * Copyright 2021 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+/* Metal kernel entry points */
+
+// clang-format off
+
+#include "kernel/device/metal/compat.h"
+#include "kernel/device/metal/globals.h"
+#include "kernel/device/gpu/kernel.h"
+
+// clang-format on

intern/cycles/kernel/device/optix/compat.h

@@ -76,6 +76,7 @@ typedef unsigned long long uint64_t;
 #define ccl_gpu_block_idx_x (blockIdx.x)
 #define ccl_gpu_grid_dim_x (gridDim.x)
 #define ccl_gpu_warp_size (warpSize)
+#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
 
 #define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
 #define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)

intern/cycles/kernel/film/accumulate.h

@@ -151,7 +151,8 @@ ccl_device_forceinline ccl_global float *kernel_accum_pixel_render_buffer(
 ccl_device_inline int kernel_accum_sample(KernelGlobals kg,
                                           ConstIntegratorState state,
                                           ccl_global float *ccl_restrict render_buffer,
-                                          int sample)
+                                          int sample,
+                                          int sample_offset)
 {
   if (kernel_data.film.pass_sample_count == PASS_UNUSED) {
     return sample;
@@ -159,7 +160,8 @@ ccl_device_inline int kernel_accum_sample(KernelGlobals kg,
 
   ccl_global float *buffer = kernel_accum_pixel_render_buffer(kg, state, render_buffer);
 
-  return atomic_fetch_and_add_uint32((uint *)(buffer) + kernel_data.film.pass_sample_count, 1);
+  return atomic_fetch_and_add_uint32((uint *)(buffer) + kernel_data.film.pass_sample_count, 1) +
+         sample_offset;
 }
 
 ccl_device void kernel_accum_adaptive_buffer(KernelGlobals kg,
@@ -499,7 +501,7 @@ ccl_device_inline void kernel_accum_light(KernelGlobals kg,
 
     /* Write shadow pass. */
     if (kernel_data.film.pass_shadow != PASS_UNUSED && (path_flag & PATH_RAY_SHADOW_FOR_LIGHT) &&
-        (path_flag & PATH_RAY_TRANSPARENT_BACKGROUND)) {
+        (path_flag & PATH_RAY_CAMERA)) {
       const float3 unshadowed_throughput = INTEGRATOR_STATE(
           state, shadow_path, unshadowed_throughput);
       const float3 shadowed_throughput = INTEGRATOR_STATE(state, shadow_path, throughput);

intern/cycles/kernel/film/passes.h

@@ -177,7 +177,7 @@ ccl_device_inline void kernel_write_data_passes(KernelGlobals kg,
 #ifdef __PASSES__
   const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
 
-  if (!(path_flag & PATH_RAY_TRANSPARENT_BACKGROUND)) {
+  if (!(path_flag & PATH_RAY_CAMERA)) {
     return;
   }
 

intern/cycles/kernel/film/read.h

@@ -460,7 +460,7 @@ ccl_device_inline float4 film_calculate_shadow_catcher_matte_with_shadow(
   const float transparency = in_matte[3] * scale;
   const float alpha = saturatef(1.0f - transparency);
 
-  const float alpha_matte = (1.0f - alpha) * (1.0f - saturatef(average(shadow_catcher))) + alpha;
+  const float alpha_matte = (1.0f - alpha) * (1.0f - average(shadow_catcher)) + alpha;
 
   if (kfilm_convert->use_approximate_shadow_catcher_background) {
     kernel_assert(kfilm_convert->pass_background != PASS_UNUSED);

intern/cycles/kernel/integrator/init_from_bake.h

@@ -65,21 +65,20 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg,
   }
 
   /* Always count the sample, even if the camera sample will reject the ray. */
-  const int sample = kernel_accum_sample(kg, state, render_buffer, scheduled_sample);
+  const int sample = kernel_accum_sample(
+      kg, state, render_buffer, scheduled_sample, tile->sample_offset);
 
   /* Setup render buffers. */
   const int index = INTEGRATOR_STATE(state, path, render_pixel_index);
   const int pass_stride = kernel_data.film.pass_stride;
-  ccl_global float *buffer = render_buffer + index * pass_stride;
+  render_buffer += index * pass_stride;
 
-  ccl_global float *primitive = buffer + kernel_data.film.pass_bake_primitive;
-  ccl_global float *differential = buffer + kernel_data.film.pass_bake_differential;
+  ccl_global float *primitive = render_buffer + kernel_data.film.pass_bake_primitive;
+  ccl_global float *differential = render_buffer + kernel_data.film.pass_bake_differential;
 
   const int seed = __float_as_uint(primitive[0]);
   int prim = __float_as_uint(primitive[1]);
   if (prim == -1) {
-    /* Accumulate transparency for empty pixels. */
-    kernel_accum_transparent(kg, state, 0, 1.0f, buffer);
     return false;
   }
 

intern/cycles/kernel/integrator/init_from_camera.h

@@ -89,7 +89,8 @@ ccl_device bool integrator_init_from_camera(KernelGlobals kg,
    * This logic allows to both count actual number of samples per pixel, and to add samples to this
    * pixel after it was converged and samples were added somewhere else (in which case the
    * `scheduled_sample` will be different from actual number of samples in this pixel). */
-  const int sample = kernel_accum_sample(kg, state, render_buffer, scheduled_sample);
+  const int sample = kernel_accum_sample(
+      kg, state, render_buffer, scheduled_sample, tile->sample_offset);
 
   /* Initialize random number seed for path. */
   const uint rng_hash = path_rng_hash_init(kg, sample, x, y);

intern/cycles/kernel/integrator/path_state.h

@@ -70,9 +70,6 @@ ccl_device_inline void path_state_init_integrator(KernelGlobals kg,
   INTEGRATOR_STATE_WRITE(state, path, continuation_probability) = 1.0f;
   INTEGRATOR_STATE_WRITE(state, path, throughput) = make_float3(1.0f, 1.0f, 1.0f);
 
-  INTEGRATOR_STATE_WRITE(state, isect, object) = OBJECT_NONE;
-  INTEGRATOR_STATE_WRITE(state, isect, prim) = PRIM_NONE;
-
   if (kernel_data.kernel_features & KERNEL_FEATURE_VOLUME) {
     INTEGRATOR_STATE_ARRAY_WRITE(state, volume_stack, 0, object) = OBJECT_NONE;
     INTEGRATOR_STATE_ARRAY_WRITE(
@@ -125,7 +122,7 @@ ccl_device_inline void path_state_next(KernelGlobals kg, IntegratorState state,
     /* volume scatter */
     flag |= PATH_RAY_VOLUME_SCATTER;
     flag &= ~PATH_RAY_TRANSPARENT_BACKGROUND;
-    if (!(flag & PATH_RAY_ANY_PASS)) {
+    if (bounce == 1) {
       flag |= PATH_RAY_VOLUME_PASS;
     }
 
@@ -187,7 +184,7 @@ ccl_device_inline void path_state_next(KernelGlobals kg, IntegratorState state,
     }
 
     /* Render pass categories. */
-    if (!(flag & PATH_RAY_ANY_PASS) && !(flag & PATH_RAY_TRANSPARENT_BACKGROUND)) {
+    if (bounce == 1) {
       flag |= PATH_RAY_SURFACE_PASS;
     }
   }
@@ -211,7 +208,9 @@ ccl_device_inline bool path_state_volume_next(IntegratorState state)
   }
 
   /* Random number generator next bounce. */
-  INTEGRATOR_STATE_WRITE(state, path, rng_offset) += PRNG_BOUNCE_NUM;
+  if (volume_bounds_bounce > 1) {
+    INTEGRATOR_STATE_WRITE(state, path, rng_offset) += PRNG_BOUNCE_NUM;
+  }
 
   return true;
 }

intern/cycles/kernel/integrator/shade_volume.h

@@ -263,12 +263,6 @@ ccl_device void volume_shadow_heterogeneous(KernelGlobals kg,
 /* Equi-angular sampling as in:
  * "Importance Sampling Techniques for Path Tracing in Participating Media" */
 
-/* Below this pdf we ignore samples, as they tend to lead to very long distances.
- * This can cause performance issues with BVH traversal in OptiX, leading it to
- * traverse many nodes. Since these contribute very little to the image, just ignore
- * those samples. */
-#  define VOLUME_SAMPLE_PDF_CUTOFF 1e-8f
-
 ccl_device float volume_equiangular_sample(ccl_private const Ray *ccl_restrict ray,
                                            const float3 light_P,
                                            const float xi,
@@ -443,8 +437,7 @@ ccl_device_forceinline void volume_integrate_step_scattering(
 
   /* Equiangular sampling for direct lighting. */
   if (vstate.direct_sample_method == VOLUME_SAMPLE_EQUIANGULAR && !result.direct_scatter) {
-    if (result.direct_t >= vstate.start_t && result.direct_t <= vstate.end_t &&
-        vstate.equiangular_pdf > VOLUME_SAMPLE_PDF_CUTOFF) {
+    if (result.direct_t >= vstate.start_t && result.direct_t <= vstate.end_t) {
       const float new_dt = result.direct_t - vstate.start_t;
       const float3 new_transmittance = volume_color_transmittance(coeff.sigma_t, new_dt);
 
@@ -481,28 +474,26 @@ ccl_device_forceinline void volume_integrate_step_scattering(
       const float3 new_transmittance = volume_color_transmittance(coeff.sigma_t, new_dt);
       const float distance_pdf = dot(channel_pdf, coeff.sigma_t * new_transmittance);
 
-      if (vstate.distance_pdf * distance_pdf > VOLUME_SAMPLE_PDF_CUTOFF) {
-        /* throughput */
-        result.indirect_scatter = true;
-        result.indirect_t = new_t;
-        result.indirect_throughput *= coeff.sigma_s * new_transmittance / distance_pdf;
-        shader_copy_volume_phases(&result.indirect_phases, sd);
+      /* throughput */
+      result.indirect_scatter = true;
+      result.indirect_t = new_t;
+      result.indirect_throughput *= coeff.sigma_s * new_transmittance / distance_pdf;
+      shader_copy_volume_phases(&result.indirect_phases, sd);
 
-        if (vstate.direct_sample_method != VOLUME_SAMPLE_EQUIANGULAR) {
-          /* If using distance sampling for direct light, just copy parameters
-           * of indirect light since we scatter at the same point then. */
-          result.direct_scatter = true;
-          result.direct_t = result.indirect_t;
-          result.direct_throughput = result.indirect_throughput;
-          shader_copy_volume_phases(&result.direct_phases, sd);
+      if (vstate.direct_sample_method != VOLUME_SAMPLE_EQUIANGULAR) {
+        /* If using distance sampling for direct light, just copy parameters
+         * of indirect light since we scatter at the same point then. */
+        result.direct_scatter = true;
+        result.direct_t = result.indirect_t;
+        result.direct_throughput = result.indirect_throughput;
+        shader_copy_volume_phases(&result.direct_phases, sd);
 
-          /* Multiple importance sampling. */
-          if (vstate.use_mis) {
-            const float equiangular_pdf = volume_equiangular_pdf(ray, equiangular_light_P, new_t);
-            const float mis_weight = power_heuristic(vstate.distance_pdf * distance_pdf,
-                                                     equiangular_pdf);
-            result.direct_throughput *= 2.0f * mis_weight;
-          }
+        /* Multiple importance sampling. */
+        if (vstate.use_mis) {
+          const float equiangular_pdf = volume_equiangular_pdf(ray, equiangular_light_P, new_t);
+          const float mis_weight = power_heuristic(vstate.distance_pdf * distance_pdf,
+                                                   equiangular_pdf);
+          result.direct_throughput *= 2.0f * mis_weight;
         }
       }
     }
@@ -703,10 +694,8 @@ ccl_device_forceinline bool integrate_volume_sample_light(
   float light_u, light_v;
   path_state_rng_2D(kg, rng_state, PRNG_LIGHT_U, &light_u, &light_v);
 
-  if (!light_distribution_sample_from_volume_segment(
-          kg, light_u, light_v, sd->time, sd->P, bounce, path_flag, ls)) {
-    return false;
-  }
+  light_distribution_sample_from_volume_segment(
+      kg, light_u, light_v, sd->time, sd->P, bounce, path_flag, ls);
 
   if (ls->shader & SHADER_EXCLUDE_SCATTER) {
     return false;

intern/cycles/kernel/integrator/shader_eval.h

@@ -122,20 +122,23 @@ ccl_device_inline void shader_prepare_surface_closures(KernelGlobals kg,
       for (int i = 0; i < sd->num_closure; i++) {
         ccl_private ShaderClosure *sc = &sd->closure[i];
 
-        if ((CLOSURE_IS_BSDF_DIFFUSE(sc->type) &&
-             (kernel_data.integrator.filter_closures & FILTER_CLOSURE_DIFFUSE)) ||
-            (CLOSURE_IS_BSDF_GLOSSY(sc->type) &&
-             (kernel_data.integrator.filter_closures & FILTER_CLOSURE_GLOSSY)) ||
-            (CLOSURE_IS_BSDF_TRANSMISSION(sc->type) &&
-             (kernel_data.integrator.filter_closures & FILTER_CLOSURE_TRANSMISSION))) {
-          sc->type = CLOSURE_NONE_ID;
-          sc->sample_weight = 0.0f;
+        if (CLOSURE_IS_BSDF_DIFFUSE(sc->type)) {
+          if (kernel_data.integrator.filter_closures & FILTER_CLOSURE_DIFFUSE) {
+            sc->type = CLOSURE_NONE_ID;
+            sc->sample_weight = 0.0f;
+          }
         }
-        else if ((CLOSURE_IS_BSDF_TRANSPARENT(sc->type) &&
-                  (kernel_data.integrator.filter_closures & FILTER_CLOSURE_TRANSPARENT))) {
-          sc->type = CLOSURE_HOLDOUT_ID;
-          sc->sample_weight = 0.0f;
-          sd->flag |= SD_HOLDOUT;
+        else if (CLOSURE_IS_BSDF_GLOSSY(sc->type)) {
+          if (kernel_data.integrator.filter_closures & FILTER_CLOSURE_GLOSSY) {
+            sc->type = CLOSURE_NONE_ID;
+            sc->sample_weight = 0.0f;
+          }
+        }
+        else if (CLOSURE_IS_BSDF_TRANSMISSION(sc->type)) {
+          if (kernel_data.integrator.filter_closures & FILTER_CLOSURE_TRANSMISSION) {
+            sc->type = CLOSURE_NONE_ID;
+            sc->sample_weight = 0.0f;
+          }
         }
       }
     }

intern/cycles/kernel/integrator/subsurface.h

@@ -71,10 +71,6 @@ ccl_device int subsurface_bounce(KernelGlobals kg,
   }
 #  endif
 
-  if (sd->flag & SD_BACKFACING) {
-    path_flag |= PATH_RAY_SUBSURFACE_BACKFACING;
-  }
-
   INTEGRATOR_STATE_WRITE(state, path, throughput) *= weight;
   INTEGRATOR_STATE_WRITE(state, path, flag) = path_flag;
 

intern/cycles/kernel/integrator/subsurface_disk.h

@@ -47,7 +47,6 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg,
   const float time = INTEGRATOR_STATE(state, ray, time);
   const float3 Ng = INTEGRATOR_STATE(state, subsurface, Ng);
   const int object = INTEGRATOR_STATE(state, isect, object);
-  const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
 
   /* Read subsurface scattering parameters. */
   const float3 radius = INTEGRATOR_STATE(state, subsurface, radius);
@@ -124,9 +123,6 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg,
     const int object = ss_isect.hits[hit].object;
     const int object_flag = kernel_tex_fetch(__object_flag, object);
     float3 hit_Ng = ss_isect.Ng[hit];
-    if (path_flag & PATH_RAY_SUBSURFACE_BACKFACING) {
-      hit_Ng = -hit_Ng;
-    }
     if (object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
       hit_Ng = -hit_Ng;
     }

intern/cycles/kernel/light/light.h

@@ -73,7 +73,7 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
     ls->P = zero_float3();
     ls->Ng = zero_float3();
     ls->D = zero_float3();
-    ls->pdf = 1.0f;
+    ls->pdf = true;
     ls->t = FLT_MAX;
     return true;
   }
@@ -131,7 +131,7 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
         float3 dir = make_float3(klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
         ls->eval_fac *= spot_light_attenuation(
             dir, klight->spot.spot_angle, klight->spot.spot_smooth, ls->Ng);
-        if (!in_volume_segment && ls->eval_fac == 0.0f) {
+        if (ls->eval_fac == 0.0f) {
           return false;
         }
       }
@@ -170,7 +170,7 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
         float3 sample_axisu = axisu;
         float3 sample_axisv = axisv;
 
-        if (!in_volume_segment && klight->area.tan_spread > 0.0f) {
+        if (klight->area.tan_spread > 0.0f) {
           if (!light_spread_clamp_area_light(
                   P, Ng, &ls->P, &sample_axisu, &sample_axisv, klight->area.tan_spread)) {
             return false;
@@ -203,7 +203,7 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
 
   ls->pdf *= kernel_data.integrator.pdf_lights;
 
-  return in_volume_segment || (ls->pdf > 0.0f);
+  return (ls->pdf > 0.0f);
 }
 
 ccl_device bool lights_intersect(KernelGlobals kg,

intern/cycles/kernel/light/sample.h

@@ -199,9 +199,6 @@ ccl_device_inline float3 shadow_ray_offset(KernelGlobals kg,
     if (offset_cutoff > 0.0f) {
       float NgL = dot(Ng, L);
       float offset_amount = 0.0f;
-      if (NL < 0) {
-        NL = -NL;
-      }
       if (NL < offset_cutoff) {
         offset_amount = clamp(2.0f - (NgL + NL) / offset_cutoff, 0.0f, 1.0f);
       }

intern/cycles/kernel/svm/aov.h

@@ -23,8 +23,7 @@ CCL_NAMESPACE_BEGIN
 ccl_device_inline bool svm_node_aov_check(const uint32_t path_flag,
                                           ccl_global float *render_buffer)
 {
-  bool is_primary = (path_flag & PATH_RAY_TRANSPARENT_BACKGROUND) &&
-                    (!(path_flag & PATH_RAY_SINGLE_PASS_DONE));
+  bool is_primary = (path_flag & PATH_RAY_CAMERA) && (!(path_flag & PATH_RAY_SINGLE_PASS_DONE));
 
   return ((render_buffer != NULL) && is_primary);
 }

intern/cycles/kernel/types.h

@@ -279,34 +279,33 @@ enum PathRayFlag {
   PATH_RAY_SUBSURFACE_RANDOM_WALK = (1U << 20U),
   PATH_RAY_SUBSURFACE_DISK = (1U << 21U),
   PATH_RAY_SUBSURFACE_USE_FRESNEL = (1U << 22U),
-  PATH_RAY_SUBSURFACE_BACKFACING = (1U << 23U),
   PATH_RAY_SUBSURFACE = (PATH_RAY_SUBSURFACE_RANDOM_WALK | PATH_RAY_SUBSURFACE_DISK |
-                         PATH_RAY_SUBSURFACE_USE_FRESNEL | PATH_RAY_SUBSURFACE_BACKFACING),
+                         PATH_RAY_SUBSURFACE_USE_FRESNEL),
 
   /* Contribute to denoising features. */
-  PATH_RAY_DENOISING_FEATURES = (1U << 24U),
+  PATH_RAY_DENOISING_FEATURES = (1U << 23U),
 
   /* Render pass categories. */
-  PATH_RAY_SURFACE_PASS = (1U << 25U),
-  PATH_RAY_VOLUME_PASS = (1U << 26U),
+  PATH_RAY_SURFACE_PASS = (1U << 24U),
+  PATH_RAY_VOLUME_PASS = (1U << 25U),
   PATH_RAY_ANY_PASS = (PATH_RAY_SURFACE_PASS | PATH_RAY_VOLUME_PASS),
 
   /* Shadow ray is for a light or surface, or AO. */
-  PATH_RAY_SHADOW_FOR_LIGHT = (1U << 27U),
-  PATH_RAY_SHADOW_FOR_AO = (1U << 28U),
+  PATH_RAY_SHADOW_FOR_LIGHT = (1U << 26U),
+  PATH_RAY_SHADOW_FOR_AO = (1U << 27U),
 
   /* A shadow catcher object was hit and the path was split into two. */
-  PATH_RAY_SHADOW_CATCHER_HIT = (1U << 29U),
+  PATH_RAY_SHADOW_CATCHER_HIT = (1U << 28U),
 
   /* A shadow catcher object was hit and this path traces only shadow catchers, writing them into
    * their dedicated pass for later division.
    *
    * NOTE: Is not covered with `PATH_RAY_ANY_PASS` because shadow catcher does special handling
    * which is separate from the light passes. */
-  PATH_RAY_SHADOW_CATCHER_PASS = (1U << 30U),
+  PATH_RAY_SHADOW_CATCHER_PASS = (1U << 29U),
 
   /* Path is evaluating background for an approximate shadow catcher with non-transparent film. */
-  PATH_RAY_SHADOW_CATCHER_BACKGROUND = (1U << 31U),
+  PATH_RAY_SHADOW_CATCHER_BACKGROUND = (1U << 30U),
 };
 
 /* Configure ray visibility bits for rays and objects respectively,
@@ -438,8 +437,7 @@ typedef enum FilterClosures {
   FILTER_CLOSURE_DIFFUSE = (1 << 1),
   FILTER_CLOSURE_GLOSSY = (1 << 2),
   FILTER_CLOSURE_TRANSMISSION = (1 << 3),
-  FILTER_CLOSURE_TRANSPARENT = (1 << 4),
-  FILTER_CLOSURE_DIRECT_LIGHT = (1 << 5),
+  FILTER_CLOSURE_DIRECT_LIGHT = (1 << 4),
 } FilterClosures;
 
 /* Shader Flag */
@@ -1426,6 +1424,7 @@ typedef struct KernelWorkTile {
 
   uint start_sample;
   uint num_samples;
+  uint sample_offset;
 
   int offset;
   uint stride;

intern/cycles/scene/image.cpp

@@ -303,7 +303,6 @@ ImageManager::ImageManager(const DeviceInfo &info)
   animation_frame = 0;
 
   /* Set image limits */
-  features.has_half_float = info.has_half_images;
   features.has_nanovdb = info.has_nanovdb;
 }
 
@@ -357,8 +356,6 @@ void ImageManager::load_image_metadata(Image *img)
 
   metadata.detect_colorspace();
 
-  assert(features.has_half_float ||
-         (metadata.type != IMAGE_DATA_TYPE_HALF4 && metadata.type != IMAGE_DATA_TYPE_HALF));
   assert(features.has_nanovdb || (metadata.type != IMAGE_DATA_TYPE_NANOVDB_FLOAT ||
                                   metadata.type != IMAGE_DATA_TYPE_NANOVDB_FLOAT3));
 

intern/cycles/scene/image.h

@@ -100,7 +100,6 @@ class ImageMetaData {
 /* Information about supported features that Image loaders can use. */
 class ImageDeviceFeatures {
  public:
-  bool has_half_float;
   bool has_nanovdb;
 };
 

intern/cycles/scene/image_oiio.cpp

@@ -30,7 +30,8 @@ OIIOImageLoader::~OIIOImageLoader()
 {
 }
 
-bool OIIOImageLoader::load_metadata(const ImageDeviceFeatures &features, ImageMetaData &metadata)
+bool OIIOImageLoader::load_metadata(const ImageDeviceFeatures & /*features*/,
+                                    ImageMetaData &metadata)
 {
   /* Perform preliminary checks, with meaningful logging. */
   if (!path_exists(filepath.string())) {
@@ -76,7 +77,7 @@ bool OIIOImageLoader::load_metadata(const ImageDeviceFeatures &features, ImageMe
   }
 
   /* check if it's half float */
-  if (spec.format == TypeDesc::HALF && features.has_half_float) {
+  if (spec.format == TypeDesc::HALF) {
     is_half = true;
   }
 

intern/cycles/scene/integrator.cpp

@@ -14,13 +14,11 @@
  * limitations under the License.
  */
 
+#include "scene/integrator.h"
 #include "device/device.h"
-
 #include "scene/background.h"
-#include "scene/bake.h"
 #include "scene/camera.h"
 #include "scene/film.h"
-#include "scene/integrator.h"
 #include "scene/jitter.h"
 #include "scene/light.h"
 #include "scene/object.h"
@@ -168,7 +166,7 @@ void Integrator::device_update(Device *device, DeviceScene *dscene, Scene *scene
   kintegrator->transparent_min_bounce = transparent_min_bounce + 1;
   kintegrator->transparent_max_bounce = transparent_max_bounce + 1;
 
-  kintegrator->ao_bounces = (ao_factor != 0.0f) ? ao_bounces : 0;
+  kintegrator->ao_bounces = ao_bounces;
   kintegrator->ao_bounces_distance = ao_distance;
   kintegrator->ao_bounces_factor = ao_factor;
   kintegrator->ao_additive_factor = ao_additive_factor;
@@ -214,11 +212,6 @@ void Integrator::device_update(Device *device, DeviceScene *dscene, Scene *scene
   if (!use_emission) {
     kintegrator->filter_closures |= FILTER_CLOSURE_EMISSION;
   }
-  if (scene->bake_manager->get_baking()) {
-    /* Baking does not need to trace through transparency, we only want to bake
-     * the object itself. */
-    kintegrator->filter_closures |= FILTER_CLOSURE_TRANSPARENT;
-  }
 
   kintegrator->seed = seed;
 

intern/cycles/scene/shader_nodes.cpp

@@ -1503,21 +1503,18 @@ void WaveTextureNode::compile(SVMCompiler &compiler)
 
   int vector_offset = tex_mapping.compile_begin(compiler, vector_in);
 
-  int scale_ofs = compiler.stack_assign_if_linked(scale_in);
-  int distortion_ofs = compiler.stack_assign_if_linked(distortion_in);
-  int detail_ofs = compiler.stack_assign_if_linked(detail_in);
-  int dscale_ofs = compiler.stack_assign_if_linked(dscale_in);
-  int droughness_ofs = compiler.stack_assign_if_linked(droughness_in);
-  int phase_ofs = compiler.stack_assign_if_linked(phase_in);
-  int color_ofs = compiler.stack_assign_if_linked(color_out);
-  int fac_ofs = compiler.stack_assign_if_linked(fac_out);
-
   compiler.add_node(NODE_TEX_WAVE,
                     compiler.encode_uchar4(wave_type, bands_direction, rings_direction, profile),
-                    compiler.encode_uchar4(vector_offset, scale_ofs, distortion_ofs),
-                    compiler.encode_uchar4(detail_ofs, dscale_ofs, droughness_ofs, phase_ofs));
+                    compiler.encode_uchar4(vector_offset,
+                                           compiler.stack_assign_if_linked(scale_in),
+                                           compiler.stack_assign_if_linked(distortion_in)),
+                    compiler.encode_uchar4(compiler.stack_assign_if_linked(detail_in),
+                                           compiler.stack_assign_if_linked(dscale_in),
+                                           compiler.stack_assign_if_linked(droughness_in),
+                                           compiler.stack_assign_if_linked(phase_in)));
 
-  compiler.add_node(compiler.encode_uchar4(color_ofs, fac_ofs),
+  compiler.add_node(compiler.encode_uchar4(compiler.stack_assign_if_linked(color_out),
+                                           compiler.stack_assign_if_linked(fac_out)),
                     __float_as_int(scale),
                     __float_as_int(distortion),
                     __float_as_int(detail));
@@ -3575,54 +3572,45 @@ void PrincipledHairBsdfNode::compile(SVMCompiler &compiler)
   int tint_ofs = compiler.stack_assign(input("Tint"));
   int absorption_coefficient_ofs = compiler.stack_assign(input("Absorption Coefficient"));
 
-  int roughness_ofs = compiler.stack_assign_if_linked(roughness_in);
-  int radial_roughness_ofs = compiler.stack_assign_if_linked(radial_roughness_in);
-
-  int normal_ofs = compiler.stack_assign_if_linked(input("Normal"));
-  int offset_ofs = compiler.stack_assign_if_linked(offset_in);
-  int ior_ofs = compiler.stack_assign_if_linked(ior_in);
-
-  int coat_ofs = compiler.stack_assign_if_linked(coat_in);
-  int melanin_ofs = compiler.stack_assign_if_linked(melanin_in);
-  int melanin_redness_ofs = compiler.stack_assign_if_linked(melanin_redness_in);
-
   ShaderInput *random_in = input("Random");
   int attr_random = random_in->link ? SVM_STACK_INVALID :
                                       compiler.attribute(ATTR_STD_CURVE_RANDOM);
-  int random_in_ofs = compiler.stack_assign_if_linked(random_in);
-  int random_color_ofs = compiler.stack_assign_if_linked(random_color_in);
-  int random_roughness_ofs = compiler.stack_assign_if_linked(random_roughness_in);
 
   /* Encode all parameters into data nodes. */
-  /* node */
-  compiler.add_node(
-      NODE_CLOSURE_BSDF,
-      /* Socket IDs can be packed 4 at a time into a single data packet */
-      compiler.encode_uchar4(
-          closure, roughness_ofs, radial_roughness_ofs, compiler.closure_mix_weight_offset()),
-      /* The rest are stored as unsigned integers */
-      __float_as_uint(roughness),
-      __float_as_uint(radial_roughness));
-  /* data node */
-  compiler.add_node(normal_ofs,
-                    compiler.encode_uchar4(offset_ofs, ior_ofs, color_ofs, parametrization),
+  compiler.add_node(NODE_CLOSURE_BSDF,
+                    /* Socket IDs can be packed 4 at a time into a single data packet */
+                    compiler.encode_uchar4(closure,
+                                           compiler.stack_assign_if_linked(roughness_in),
+                                           compiler.stack_assign_if_linked(radial_roughness_in),
+                                           compiler.closure_mix_weight_offset()),
+                    /* The rest are stored as unsigned integers */
+                    __float_as_uint(roughness),
+                    __float_as_uint(radial_roughness));
+
+  compiler.add_node(compiler.stack_assign_if_linked(input("Normal")),
+                    compiler.encode_uchar4(compiler.stack_assign_if_linked(offset_in),
+                                           compiler.stack_assign_if_linked(ior_in),
+                                           color_ofs,
+                                           parametrization),
                     __float_as_uint(offset),
                     __float_as_uint(ior));
-  /* data node 2 */
-  compiler.add_node(compiler.encode_uchar4(
-                        coat_ofs, melanin_ofs, melanin_redness_ofs, absorption_coefficient_ofs),
+
+  compiler.add_node(compiler.encode_uchar4(compiler.stack_assign_if_linked(coat_in),
+                                           compiler.stack_assign_if_linked(melanin_in),
+                                           compiler.stack_assign_if_linked(melanin_redness_in),
+                                           absorption_coefficient_ofs),
                     __float_as_uint(coat),
                     __float_as_uint(melanin),
                     __float_as_uint(melanin_redness));
 
-  /* data node 3 */
-  compiler.add_node(
-      compiler.encode_uchar4(tint_ofs, random_in_ofs, random_color_ofs, random_roughness_ofs),
-      __float_as_uint(random),
-      __float_as_uint(random_color),
-      __float_as_uint(random_roughness));
+  compiler.add_node(compiler.encode_uchar4(tint_ofs,
+                                           compiler.stack_assign_if_linked(random_in),
+                                           compiler.stack_assign_if_linked(random_color_in),
+                                           compiler.stack_assign_if_linked(random_roughness_in)),
+                    __float_as_uint(random),
+                    __float_as_uint(random_color),
+                    __float_as_uint(random_roughness));
 
-  /* data node 4 */
   compiler.add_node(
       compiler.encode_uchar4(
           SVM_STACK_INVALID, SVM_STACK_INVALID, SVM_STACK_INVALID, SVM_STACK_INVALID),

intern/cycles/session/session.cpp

@@ -262,6 +262,7 @@ RenderWork Session::run_update_for_next_iteration()
   }
 
   render_scheduler_.set_num_samples(params.samples);
+  render_scheduler_.set_start_sample(params.sample_offset);
   render_scheduler_.set_time_limit(params.time_limit);
 
   while (have_tiles) {
@@ -367,26 +368,14 @@ void Session::draw()
 
 int2 Session::get_effective_tile_size() const
 {
-  const int image_width = buffer_params_.width;
-  const int image_height = buffer_params_.height;
-
   /* No support yet for baking with tiles. */
   if (!params.use_auto_tile || scene->bake_manager->get_baking()) {
-    return make_int2(image_width, image_height);
+    return make_int2(buffer_params_.width, buffer_params_.height);
   }
 
-  const int64_t image_area = static_cast<int64_t>(image_width) * image_height;
-
-  /* TODO(sergey): Take available memory into account, and if there is enough memory do not
-   * tile and prefer optimal performance. */
-
+  /* TODO(sergey): Take available memory into account, and if there is enough memory do not tile
+   * and prefer optimal performance. */
   const int tile_size = tile_manager_.compute_render_tile_size(params.tile_size);
-  const int64_t actual_tile_area = static_cast<int64_t>(tile_size) * tile_size;
-
-  if (actual_tile_area >= image_area) {
-    return make_int2(image_width, image_height);
-  }
-
   return make_int2(tile_size, tile_size);
 }
 
@@ -409,7 +398,7 @@ void Session::do_delayed_reset()
 
   /* Tile and work scheduling. */
   tile_manager_.reset_scheduling(buffer_params_, get_effective_tile_size());
-  render_scheduler_.reset(buffer_params_, params.samples);
+  render_scheduler_.reset(buffer_params_, params.samples, params.sample_offset);
 
   /* Passes. */
   /* When multiple tiles are used SAMPLE_COUNT pass is used to keep track of possible partial
@@ -516,8 +505,8 @@ void Session::set_display_driver(unique_ptr<DisplayDriver> driver)
 
 double Session::get_estimated_remaining_time() const
 {
-  const double completed = progress.get_progress();
-  if (completed == 0.0) {
+  const float completed = progress.get_progress();
+  if (completed == 0.0f) {
     return 0.0;
   }
 
@@ -585,7 +574,7 @@ void Session::update_status_time(bool show_pause, bool show_done)
   }
 
   /* Sample. */
-  if (!params.background && num_samples == Integrator::MAX_SAMPLES) {
+  if (num_samples == Integrator::MAX_SAMPLES) {
     substatus = status_append(substatus, string_printf("Sample %d", current_sample));
   }
   else {

intern/cycles/session/session.h

@@ -54,6 +54,7 @@ class SessionParams {
 
   bool experimental;
   int samples;
+  int sample_offset;
   int pixel_size;
   int threads;
 
@@ -75,6 +76,7 @@ class SessionParams {
 
     experimental = false;
     samples = 1024;
+    sample_offset = 0;
     pixel_size = 1;
     threads = 0;
     time_limit = 0.0;