Bump version to 3.0.1 release

unity launches blender in background mode to do some
file conversions, ever since the launcher got introduced
this process broke.

The root cause here is: Unity looks up the default program
to launch .blend files with, which is now the launcher, then
launches it in background mode with a script to export the data.

The launcher however was designed to exit as quickly as
possible so there would not be an extra background process
lingering. It does not wait for blender to exit and does not
pass back any error codes.

This broke unity's workflow since it assumed if the process
exits and succeeds the data *must* be ready for reading which
no longer holds true.

This change keeps the launcher design as was previously,
*except* when launching in background mode, then it
waits and passes back any error codes, thus restoring
unity's workflow.

Differential Revision: https://developer.blender.org/D13894
Reviewed by: LazyDodo, Brecht

CMakeLists.txt

@@ -440,7 +440,11 @@ mark_as_advanced(WITH_CYCLES_CUDA_BUILD_SERIAL)
 mark_as_advanced(WITH_CUDA_DYNLOAD)
 
 # AMD HIP
-option(WITH_CYCLES_DEVICE_HIP        "Enable Cycles AMD HIP support" OFF)
+if(WIN32)
+  option(WITH_CYCLES_DEVICE_HIP        "Enable Cycles AMD HIP support" ON)
+else()
+  option(WITH_CYCLES_DEVICE_HIP        "Enable Cycles AMD HIP support" OFF)
+endif()
 option(WITH_CYCLES_HIP_BINARIES      "Build Cycles AMD HIP binaries" OFF)
 set(CYCLES_HIP_BINARIES_ARCH gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1034 CACHE STRING "AMD HIP architectures to build binaries for")
 mark_as_advanced(WITH_CYCLES_DEVICE_HIP)
@@ -490,7 +494,8 @@ endif()
 
 # This should be turned off when Blender enter beta/rc/release
 if("${BLENDER_VERSION_CYCLE}" STREQUAL "release" OR
-   "${BLENDER_VERSION_CYCLE}" STREQUAL "rc")
+   "${BLENDER_VERSION_CYCLE}" STREQUAL "rc" OR
+   "${BLENDER_VERSION_CYCLE}" STREQUAL "beta")
   set(WITH_EXPERIMENTAL_FEATURES OFF)
 else()
   set(WITH_EXPERIMENTAL_FEATURES ON)

build_files/build_environment/cmake/nanovdb.cmake

@@ -42,6 +42,7 @@ ExternalProject_Add(nanovdb
   URL_HASH ${NANOVDB_HASH_TYPE}=${NANOVDB_HASH}
   PREFIX ${BUILD_DIR}/nanovdb
   SOURCE_SUBDIR nanovdb
+  PATCH_COMMAND ${PATCH_CMD} -p 1 -d ${BUILD_DIR}/nanovdb/src/nanovdb < ${PATCH_DIR}/nanovdb.diff
   CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/nanovdb ${DEFAULT_CMAKE_FLAGS} ${NANOVDB_EXTRA_ARGS}
   INSTALL_DIR ${LIBDIR}/nanovdb
 )

build_files/build_environment/patches/nanovdb.diff

@@ -0,0 +1,374 @@
+Index: nanovdb/nanovdb/NanoVDB.h
+===================================================================
+--- a/nanovdb/nanovdb/NanoVDB.h	(revision 62751)
++++ b/nanovdb/nanovdb/NanoVDB.h	(working copy)
+@@ -152,8 +152,8 @@
+ 
+ #endif // __CUDACC_RTC__
+ 
+-#ifdef __CUDACC__
+-// Only define __hostdev__ when using NVIDIA CUDA compiler
++#if defined(__CUDACC__) || defined(__HIP__)
++// Only define __hostdev__ when using NVIDIA CUDA or HIP compiler
+ #define __hostdev__ __host__ __device__
+ #else
+ #define __hostdev__
+@@ -461,7 +461,7 @@
+ /// Maximum floating-point values
+ template<typename T>
+ struct Maximum;
+-#ifdef __CUDA_ARCH__
++#if defined(__CUDA_ARCH__) || defined(__HIP__)
+ template<>
+ struct Maximum<int>
+ {
+@@ -1006,10 +1006,10 @@
+ using Vec3i = Vec3<int>;
+ 
+ /// @brief Return a single precision floating-point vector of this coordinate
+-Vec3f Coord::asVec3s() const { return Vec3f(float(mVec[0]), float(mVec[1]), float(mVec[2])); }
++inline __hostdev__ Vec3f Coord::asVec3s() const { return Vec3f(float(mVec[0]), float(mVec[1]), float(mVec[2])); }
+ 
+ /// @brief Return a double precision floating-point vector of this coordinate
+-Vec3d Coord::asVec3d() const { return Vec3d(double(mVec[0]), double(mVec[1]), double(mVec[2])); }
++inline __hostdev__ Vec3d Coord::asVec3d() const { return Vec3d(double(mVec[0]), double(mVec[1]), double(mVec[2])); }
+ 
+ // ----------------------------> Vec4 <--------------------------------------
+ 
+@@ -1820,7 +1820,7 @@
+ }; // Map
+ 
+ template<typename Mat4T>
+-void Map::set(const Mat4T& mat, const Mat4T& invMat, double taper)
++__hostdev__ void Map::set(const Mat4T& mat, const Mat4T& invMat, double taper)
+ {
+     float * mf = mMatF, *vf = mVecF;
+     float*  mif = mInvMatF;
+@@ -2170,7 +2170,7 @@
+ }; // Class Grid
+ 
+ template<typename TreeT>
+-int Grid<TreeT>::findBlindDataForSemantic(GridBlindDataSemantic semantic) const
++__hostdev__ int Grid<TreeT>::findBlindDataForSemantic(GridBlindDataSemantic semantic) const
+ {
+     for (uint32_t i = 0, n = blindDataCount(); i < n; ++i)
+         if (blindMetaData(i).mSemantic == semantic)
+@@ -2328,7 +2328,7 @@
+ }; // Tree class
+ 
+ template<typename RootT>
+-void Tree<RootT>::extrema(ValueType& min, ValueType& max) const
++__hostdev__ void Tree<RootT>::extrema(ValueType& min, ValueType& max) const
+ {
+     min = this->root().valueMin();
+     max = this->root().valueMax();
+@@ -2336,7 +2336,7 @@
+ 
+ template<typename RootT>
+ template<typename NodeT>
+-const NodeT* Tree<RootT>::getNode(uint32_t i) const
++__hostdev__ const NodeT* Tree<RootT>::getNode(uint32_t i) const
+ {
+     static_assert(is_same<TreeNodeT<NodeT::LEVEL>, NodeT>::value, "Tree::getNode: unvalid node type");
+     NANOVDB_ASSERT(i < DataType::mCount[NodeT::LEVEL]);
+@@ -2345,7 +2345,7 @@
+ 
+ template<typename RootT>
+ template<int LEVEL>
+-const typename TreeNode<Tree<RootT>, LEVEL>::type* Tree<RootT>::getNode(uint32_t i) const
++__hostdev__ const typename TreeNode<Tree<RootT>, LEVEL>::type* Tree<RootT>::getNode(uint32_t i) const
+ {
+     NANOVDB_ASSERT(i < DataType::mCount[LEVEL]);
+     return reinterpret_cast<const TreeNodeT<LEVEL>*>(reinterpret_cast<const uint8_t*>(this) + DataType::mBytes[LEVEL]) + i;
+@@ -2353,7 +2353,7 @@
+ 
+ template<typename RootT>
+ template<typename NodeT>
+-NodeT* Tree<RootT>::getNode(uint32_t i)
++__hostdev__ NodeT* Tree<RootT>::getNode(uint32_t i)
+ {
+     static_assert(is_same<TreeNodeT<NodeT::LEVEL>, NodeT>::value, "Tree::getNode: invalid node type");
+     NANOVDB_ASSERT(i < DataType::mCount[NodeT::LEVEL]);
+@@ -2362,7 +2362,7 @@
+ 
+ template<typename RootT>
+ template<int LEVEL>
+-typename TreeNode<Tree<RootT>, LEVEL>::type* Tree<RootT>::getNode(uint32_t i)
++__hostdev__ typename TreeNode<Tree<RootT>, LEVEL>::type* Tree<RootT>::getNode(uint32_t i)
+ {
+     NANOVDB_ASSERT(i < DataType::mCount[LEVEL]);
+     return reinterpret_cast<TreeNodeT<LEVEL>*>(reinterpret_cast<uint8_t*>(this) + DataType::mBytes[LEVEL]) + i;
+@@ -2370,7 +2370,7 @@
+ 
+ template<typename RootT>
+ template<typename NodeT>
+-uint32_t Tree<RootT>::getNodeID(const NodeT& node) const
++__hostdev__ uint32_t Tree<RootT>::getNodeID(const NodeT& node) const
+ {
+     static_assert(is_same<TreeNodeT<NodeT::LEVEL>, NodeT>::value, "Tree::getNodeID: invalid node type");
+     const NodeT* first = reinterpret_cast<const NodeT*>(reinterpret_cast<const uint8_t*>(this) + DataType::mBytes[NodeT::LEVEL]);
+@@ -2380,7 +2380,7 @@
+ 
+ template<typename RootT>
+ template<typename NodeT>
+-uint32_t Tree<RootT>::getLinearOffset(const NodeT& node) const
++__hostdev__ uint32_t Tree<RootT>::getLinearOffset(const NodeT& node) const
+ {
+     return this->getNodeID(node) + DataType::mPFSum[NodeT::LEVEL];
+ }
+@@ -3366,7 +3366,7 @@
+ }; // LeafNode class
+ 
+ template<typename ValueT, typename CoordT, template<uint32_t> class MaskT, uint32_t LOG2DIM>
+-inline void LeafNode<ValueT, CoordT, MaskT, LOG2DIM>::updateBBox()
++inline __hostdev__ void LeafNode<ValueT, CoordT, MaskT, LOG2DIM>::updateBBox()
+ {
+     static_assert(LOG2DIM == 3, "LeafNode::updateBBox: only supports LOGDIM = 3!");
+     if (!this->isActive()) return;
+Index: nanovdb/nanovdb/util/SampleFromVoxels.h
+===================================================================
+--- a/nanovdb/nanovdb/util/SampleFromVoxels.h	(revision 62751)
++++ b/nanovdb/nanovdb/util/SampleFromVoxels.h	(working copy)
+@@ -22,7 +22,7 @@
+ #define NANOVDB_SAMPLE_FROM_VOXELS_H_HAS_BEEN_INCLUDED
+ 
+ // Only define __hostdev__ when compiling as NVIDIA CUDA
+-#ifdef __CUDACC__
++#if defined(__CUDACC__) || defined(__HIP__)
+ #define __hostdev__ __host__ __device__
+ #else
+ #include <cmath> // for floor
+@@ -136,7 +136,7 @@
+ 
+ template<typename TreeOrAccT>
+ template<typename Vec3T>
+-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, true>::operator()(const Vec3T& xyz) const
++__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, true>::operator()(const Vec3T& xyz) const
+ {
+     const CoordT ijk = Round<CoordT>(xyz);
+     if (ijk != mPos) {
+@@ -147,7 +147,7 @@
+ }
+ 
+ template<typename TreeOrAccT>
+-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, true>::operator()(const CoordT& ijk) const
++__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, true>::operator()(const CoordT& ijk) const
+ {
+     if (ijk != mPos) {
+         mPos = ijk;
+@@ -158,7 +158,7 @@
+ 
+ template<typename TreeOrAccT>
+ template<typename Vec3T>
+-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, false>::operator()(const Vec3T& xyz) const
++__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 0, false>::operator()(const Vec3T& xyz) const
+ {
+     return mAcc.getValue(Round<CoordT>(xyz));
+ }
+@@ -195,7 +195,7 @@
+ }; // TrilinearSamplerBase
+ 
+ template<typename TreeOrAccT>
+-void TrilinearSampler<TreeOrAccT>::stencil(CoordT& ijk, ValueT (&v)[2][2][2]) const
++__hostdev__ void TrilinearSampler<TreeOrAccT>::stencil(CoordT& ijk, ValueT (&v)[2][2][2]) const
+ {
+     v[0][0][0] = mAcc.getValue(ijk); // i, j, k
+ 
+@@ -224,7 +224,7 @@
+ 
+ template<typename TreeOrAccT>
+ template<typename RealT, template<typename...> class Vec3T>
+-typename TreeOrAccT::ValueType TrilinearSampler<TreeOrAccT>::sample(const Vec3T<RealT> &uvw, const ValueT (&v)[2][2][2])
++__hostdev__ typename TreeOrAccT::ValueType TrilinearSampler<TreeOrAccT>::sample(const Vec3T<RealT> &uvw, const ValueT (&v)[2][2][2])
+ {
+ #if 0
+   auto lerp = [](ValueT a, ValueT b, ValueT w){ return fma(w, b-a, a); };// = w*(b-a) + a
+@@ -239,7 +239,7 @@
+ 
+ template<typename TreeOrAccT>
+ template<typename RealT, template<typename...> class Vec3T>
+-Vec3T<typename TreeOrAccT::ValueType> TrilinearSampler<TreeOrAccT>::gradient(const Vec3T<RealT> &uvw, const ValueT (&v)[2][2][2])
++__hostdev__ Vec3T<typename TreeOrAccT::ValueType> TrilinearSampler<TreeOrAccT>::gradient(const Vec3T<RealT> &uvw, const ValueT (&v)[2][2][2])
+ {
+     static_assert(std::is_floating_point<ValueT>::value, "TrilinearSampler::gradient requires a floating-point type");
+ #if 0
+@@ -270,7 +270,7 @@
+ }
+ 
+ template<typename TreeOrAccT>
+-bool TrilinearSampler<TreeOrAccT>::zeroCrossing(const ValueT (&v)[2][2][2])
++__hostdev__ bool TrilinearSampler<TreeOrAccT>::zeroCrossing(const ValueT (&v)[2][2][2])
+ {
+     static_assert(std::is_floating_point<ValueT>::value, "TrilinearSampler::zeroCrossing requires a floating-point type");
+     const bool less = v[0][0][0] < ValueT(0);
+@@ -363,7 +363,7 @@
+ 
+ template<typename TreeOrAccT>
+ template<typename RealT, template<typename...> class Vec3T>
+-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, true>::operator()(Vec3T<RealT> xyz) const
++__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, true>::operator()(Vec3T<RealT> xyz) const
+ {
+     this->cache(xyz);
+     return BaseT::sample(xyz, mVal);
+@@ -370,7 +370,7 @@
+ }
+ 
+ template<typename TreeOrAccT>
+-typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, true>::operator()(const CoordT &ijk) const
++__hostdev__ typename TreeOrAccT::ValueType SampleFromVoxels<TreeOrAccT, 1, true>::operator()(const CoordT &ijk) const
+ {
+     return  ijk == mPos ? mVal[0][0][0] : BaseT::mAcc.getValue(ijk);
+ }
+@@ -377,7 +377,7 @@
+ 
+ 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/cmake/config/blender_release.cmake

@@ -81,4 +81,5 @@ if(NOT APPLE)
   set(WITH_CYCLES_DEVICE_OPTIX    ON  CACHE BOOL "" FORCE)
   set(WITH_CYCLES_CUDA_BINARIES   ON  CACHE BOOL "" FORCE)
   set(WITH_CYCLES_CUBIN_COMPILER  OFF CACHE BOOL "" FORCE)
+  set(WITH_CYCLES_HIP_BINARIES    ON  CACHE BOOL "" FORCE)
 endif()

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

@@ -42,8 +42,13 @@ 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
+# Register and add to the file selector (required to also use F3 search "Text Export Operator" for quick access)
 bpy.utils.register_class(ExportSomeData)
 bpy.types.TOPBAR_MT_file_export.append(menu_func)
 

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

@@ -27,8 +27,14 @@ 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,8 +41,13 @@ 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,8 +55,13 @@ 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,8 +31,13 @@ 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,8 +22,13 @@ 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,24 +106,6 @@ 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": ("MaterialTextureSlot", False),
+    "texture_slot": ("TextureSlot", False),
     "texture_user": ("ID", False),
     "texture_user_property": ("Property", False),
     "ui_list": ("UIList", False),

extern/hipew/README

@@ -0,0 +1,12 @@
+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

@@ -0,0 +1,5 @@
+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,31 +804,29 @@ typedef enum hipDeviceP2PAttr {
 } hipDeviceP2PAttr;
 
 typedef struct HIP_MEMCPY3D {
-  size_t srcXInBytes;
-  size_t srcY;
-  size_t srcZ;
-  size_t srcLOD;
+  unsigned int srcXInBytes;
+  unsigned int srcY;
+  unsigned int srcZ;
+  unsigned int srcLOD;
   hipMemoryType srcMemoryType;
   const void* srcHost;
   hipDeviceptr_t srcDevice;
-  hArray * srcArray;
-  void* reserved0;
-  size_t srcPitch;
-  size_t srcHeight;
-  size_t dstXInBytes;
-  size_t dstY;
-  size_t dstZ;
-  size_t dstLOD;
+  hArray srcArray;
+  unsigned int srcPitch;
+  unsigned int srcHeight;
+  unsigned int dstXInBytes;
+  unsigned int dstY;
+  unsigned int dstZ;
+  unsigned int dstLOD;
   hipMemoryType dstMemoryType;
   void* dstHost;
   hipDeviceptr_t dstDevice;
-  hArray * dstArray;
-  void* reserved1;
-  size_t dstPitch;
-  size_t dstHeight;
-  size_t WidthInBytes;
-  size_t Height;
-  size_t Depth;
+  hArray dstArray;
+  unsigned int dstPitch;
+  unsigned int dstHeight;
+  unsigned int WidthInBytes;
+  unsigned int Height;
+  unsigned int Depth;
 } HIP_MEMCPY3D;
 
 typedef struct HIP_MEMCPY3D_PEER_st {
@@ -879,7 +877,7 @@ typedef struct HIP_RESOURCE_DESC_st {
   hipResourceType resType;
   union {
     struct {
-      hArray * h_Array;
+      hArray h_Array;
     } array;
     struct {
       hipMipmappedArray_t hMipmappedArray;
@@ -1074,9 +1072,10 @@ 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(hipDevice_t* device, int ordinal);
+typedef hipError_t HIPAPI thipGetDevice(int* device);
 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);
@@ -1209,6 +1208,7 @@ 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,6 +1333,7 @@ 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,6 +71,7 @@ thipDriverGetVersion *hipDriverGetVersion;
 thipGetDevice *hipGetDevice;
 thipGetDeviceCount *hipGetDeviceCount;
 thipGetDeviceProperties *hipGetDeviceProperties;
+thipDeviceGet* hipDeviceGet;
 thipDeviceGetName *hipDeviceGetName;
 thipDeviceGetAttribute *hipDeviceGetAttribute;
 thipDeviceComputeCapability *hipDeviceComputeCapability;
@@ -213,6 +214,36 @@ 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
@@ -240,6 +271,14 @@ 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);
 
@@ -255,6 +294,7 @@ 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: 
+Upstream version: 3cdd4a9d7886
 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 */
-  float progress = options.session->progress.get_progress();
+  double 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);
-  float progress_val = progress.get_progress();
+  double progress_val = progress.get_progress();
 
   if (substatus != "")
     status += ": " + substatus;

intern/cycles/blender/addon/properties.py

@@ -346,7 +346,7 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         name="Scrambling Distance",
         default=1.0,
         min=0.0, max=1.0,
-        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",
+        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",
     )
     preview_scrambling_distance: BoolProperty(
         name="Scrambling Distance viewport",
@@ -354,10 +354,10 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         description="Uses the Scrambling Distance value for the viewport. Faster but may flicker",
     )
 
-    adaptive_scrambling_distance: BoolProperty(
-        name="Adaptive Scrambling Distance",
+    auto_scrambling_distance: BoolProperty(
+        name="Automatic Scrambling Distance",
         default=False,
-        description="Uses a formula to adapt the scrambling distance strength based on the sample count",
+        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",
     )
 
     use_layer_samples: EnumProperty(
@@ -770,8 +770,8 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
     )
 
     use_auto_tile: BoolProperty(
-        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",
+        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",
         default=True,
     )
     tile_size: IntProperty(

intern/cycles/blender/addon/ui.py

@@ -292,13 +292,13 @@ class CYCLES_RENDER_PT_sampling_advanced(CyclesButtonsPanel, Panel):
 
         layout.separator()
 
-        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)
+        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()
         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()
 
@@ -1051,7 +1051,7 @@ class CYCLES_OBJECT_PT_motion_blur(CyclesButtonsPanel, Panel):
 
 
 def has_geometry_visibility(ob):
-    return ob and ((ob.type in {'MESH', 'CURVE', 'SURFACE', 'FONT', 'META', 'LIGHT'}) or
+    return ob and ((ob.type in {'MESH', 'CURVE', 'SURFACE', 'FONT', 'META', 'LIGHT', 'VOLUME', 'POINTCLOUD', 'HAIR'}) or
                    (ob.instance_type == 'COLLECTION' and ob.instance_collection))
 
 

intern/cycles/blender/curves.cpp

@@ -819,11 +819,14 @@ 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 {
+    else
+#endif
+    {
       /* Particle hair. */
       bool need_undeformed = new_hair.need_attribute(scene, ATTR_STD_GENERATED);
       BL::Mesh b_mesh = object_to_mesh(
@@ -870,12 +873,15 @@ 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 {
+    else
+#endif
+    {
       /* 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,7 +31,11 @@ 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;
   }
 
@@ -215,7 +219,11 @@ 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,8 +24,14 @@ CCL_NAMESPACE_BEGIN
 
 /* Packed Images */
 
-BlenderImageLoader::BlenderImageLoader(BL::Image b_image, int frame)
-    : b_image(b_image), frame(frame), free_cache(!b_image.has_data())
+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())
 {
 }
 

intern/cycles/blender/image.h

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

intern/cycles/blender/object.cpp

@@ -294,7 +294,7 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph,
 
   object->set_visibility(visibility);
 
-  object->set_is_shadow_catcher(b_ob.is_shadow_catcher());
+  object->set_is_shadow_catcher(b_ob.is_shadow_catcher() || b_parent.is_shadow_catcher());
 
   float shadow_terminator_shading_offset = get_float(cobject, "shadow_terminator_offset");
   object->set_shadow_terminator_shading_offset(shadow_terminator_shading_offset);

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, true, false, true);
+  b_engine_.end_result(b_rr, false, 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.0f;
+  last_progress = -1.0;
   start_resize_time = 0.0;
 
   /* create session */
@@ -615,6 +615,24 @@ 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();
   }
 
@@ -841,7 +859,7 @@ void BlenderSession::get_status(string &status, string &substatus)
   session->progress.get_status(status, substatus);
 }
 
-void BlenderSession::get_progress(float &progress, double &total_time, double &render_time)
+void BlenderSession::get_progress(double &progress, double &total_time, double &render_time)
 {
   session->progress.get_time(total_time, render_time);
   progress = session->progress.get_progress();
@@ -849,10 +867,10 @@ void BlenderSession::get_progress(float &progress, double &total_time, double &r
 
 void BlenderSession::update_bake_progress()
 {
-  float progress = session->progress.get_progress();
+  double progress = session->progress.get_progress();
 
   if (progress != last_progress) {
-    b_engine.update_progress(progress);
+    b_engine.update_progress((float)progress);
     last_progress = progress;
   }
 }
@@ -861,7 +879,7 @@ void BlenderSession::update_status_progress()
 {
   string timestatus, status, substatus;
   string scene_status = "";
-  float progress;
+  double 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;
@@ -905,7 +923,7 @@ void BlenderSession::update_status_progress()
     last_status_time = current_time;
   }
   if (progress != last_progress) {
-    b_engine.update_progress(progress);
+    b_engine.update_progress((float)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(float &progress, double &total_time, double &render_time);
+  void get_progress(double &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;
-  float last_progress;
+  double last_progress;
   double last_status_time;
 
   int width, height;

intern/cycles/blender/shader.cpp

@@ -762,7 +762,8 @@ 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), image->image_params());
+            new BlenderImageLoader(b_image, image_frame, b_engine.is_preview()),
+            image->image_params());
       }
       else {
         ustring filename = ustring(
@@ -797,8 +798,9 @@ 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),
-                                                      env->image_params());
+        env->handle = scene->image_manager->add_image(
+            new BlenderImageLoader(b_image, image_frame, b_engine.is_preview()),
+            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 adaptive_scrambling_distance = get_boolean(cscene, "adaptive_scrambling_distance");
-  if (adaptive_scrambling_distance) {
+  bool auto_scrambling_distance = get_boolean(cscene, "auto_scrambling_distance");
+  if (auto_scrambling_distance) {
     scrambling_distance *= 4.0f / sqrtf(samples);
   }
 

intern/cycles/bvh/optix.cpp

@@ -30,15 +30,17 @@ BVHOptiX::BVHOptiX(const BVHParams &params_,
     : BVH(params_, geometry_, objects_),
       device(device),
       traversable_handle(0),
-      as_data(device, params_.top_level ? "optix tlas" : "optix blas", false),
-      motion_transform_data(device, "optix motion transform", false)
+      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))
 {
 }
 
 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,14 +25,16 @@
 
 #  include "device/memory.h"
 
+#  include "util/unique_ptr.h"
+
 CCL_NAMESPACE_BEGIN
 
 class BVHOptiX : public BVH {
  public:
   Device *device;
   uint64_t traversable_handle;
-  device_only_memory<char> as_data;
-  device_only_memory<char> motion_transform_data;
+  unique_ptr<device_only_memory<char>> as_data;
+  unique_ptr<device_only_memory<char>> motion_transform_data;
 
  protected:
   friend class BVH;

intern/cycles/device/cpu/device_impl.cpp

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

intern/cycles/device/cpu/kernel.h

@@ -42,7 +42,7 @@ class CPUKernels {
 
   IntegratorInitFunction integrator_init_from_camera;
   IntegratorInitFunction integrator_init_from_bake;
-  IntegratorFunction integrator_intersect_closest;
+  IntegratorShadeFunction integrator_intersect_closest;
   IntegratorFunction integrator_intersect_shadow;
   IntegratorFunction integrator_intersect_subsurface;
   IntegratorFunction integrator_intersect_volume_stack;

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) {
+  if (mem_alloc_result != CUDA_SUCCESS && can_map_host && mem.type != MEM_DEVICE_ONLY) {
     if (mem.shared_pointer) {
       /* Another device already allocated host memory. */
       mem_alloc_result = CUDA_SUCCESS;
@@ -703,8 +703,14 @@ CUDADevice::CUDAMem *CUDADevice::generic_alloc(device_memory &mem, size_t pitch_
   }
 
   if (mem_alloc_result != CUDA_SUCCESS) {
-    status = " failed, out of device and host memory";
-    set_error("System is out of GPU and shared host memory");
+    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");
+    }
   }
 
   if (mem.name) {
@@ -777,6 +783,7 @@ 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
@@ -1145,6 +1152,7 @@ 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/hip/device.cpp

@@ -57,9 +57,16 @@ bool device_hip_init()
     }
   }
   else {
-    VLOG(1) << "HIPEW initialization failed: "
-            << ((hipew_result == HIPEW_ERROR_ATEXIT_FAILED) ? "Error setting up atexit() handler" :
-                                                              "Error opening the library");
+    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";
+    }
   }
 
   return result;

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 = hipGetDevice(&hipDevice, hipDevId);
+  result = hipDeviceGet(&hipDevice, hipDevId);
   if (result != hipSuccess) {
     set_error(string_printf("Failed to get HIP device handle from ordinal (%s)",
                             hipewErrorString(result)));
@@ -222,7 +222,6 @@ string HIPDevice::compile_kernel_get_common_cflags(const uint kernel_features)
   const string include_path = source_path;
   string cflags = string_printf(
       "-m%d "
-      "--ptxas-options=\"-v\" "
       "--use_fast_math "
       "-DHIPCC "
       "-I\"%s\"",
@@ -234,10 +233,7 @@ string HIPDevice::compile_kernel_get_common_cflags(const uint kernel_features)
   return cflags;
 }
 
-string HIPDevice::compile_kernel(const uint kernel_features,
-                                 const char *name,
-                                 const char *base,
-                                 bool force_ptx)
+string HIPDevice::compile_kernel(const uint kernel_features, const char *name, const char *base)
 {
   /* Compute kernel name. */
   int major, minor;
@@ -255,13 +251,11 @@ string HIPDevice::compile_kernel(const uint kernel_features,
 
   /* Attempt to use kernel provided with Blender. */
   if (!use_adaptive_compilation()) {
-    if (!force_ptx) {
-      const string fatbin = path_get(string_printf("lib/%s_%s.fatbin", name, arch));
-      VLOG(1) << "Testing for pre-compiled kernel " << fatbin << ".";
-      if (path_exists(fatbin)) {
-        VLOG(1) << "Using precompiled kernel.";
-        return fatbin;
-      }
+    const string fatbin = path_get(string_printf("lib/%s_%s.fatbin", name, arch));
+    VLOG(1) << "Testing for pre-compiled kernel " << fatbin << ".";
+    if (path_exists(fatbin)) {
+      VLOG(1) << "Using precompiled kernel.";
+      return fatbin;
     }
   }
 
@@ -298,9 +292,9 @@ string HIPDevice::compile_kernel(const uint kernel_features,
 
 #  ifdef _WIN32
   if (!use_adaptive_compilation() && have_precompiled_kernels()) {
-    if (major < 3) {
+    if (!hipSupportsDevice(hipDevId)) {
       set_error(
-          string_printf("HIP backend requires compute capability 3.0 or up, but found %d.%d. "
+          string_printf("HIP backend requires compute capability 10.1 or up, but found %d.%d. "
                         "Your GPU is not supported.",
                         major,
                         minor));
@@ -751,6 +745,7 @@ 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
@@ -994,16 +989,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(&array_3d, &desc));
+    hip_assert(hipArray3DCreate((hArray *)&array_3d, &desc));
 
     if (!array_3d) {
       return;
     }
 
     HIP_MEMCPY3D param;
-    memset(&param, 0, sizeof(param));
+    memset(&param, 0, sizeof(HIP_MEMCPY3D));
     param.dstMemoryType = hipMemoryTypeArray;
-    param.dstArray = &array_3d;
+    param.dstArray = array_3d;
     param.srcMemoryType = hipMemoryTypeHost;
     param.srcHost = mem.host_pointer;
     param.srcPitch = src_pitch;
@@ -1069,13 +1064,13 @@ 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) {
-    /* Kepler+, bindless textures. */
+    /* 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) {
@@ -1120,6 +1115,7 @@ 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) {
@@ -1160,6 +1156,8 @@ bool HIPDevice::should_use_graphics_interop()
    * possible, but from the empiric measurements it can be considerably slower than using naive
    * pixels copy. */
 
+  /* Disable graphics interop for now, because of driver bug in 21.40. See T92972 */
+#  if 0
   HIPContextScope scope(this);
 
   int num_all_devices = 0;
@@ -1178,6 +1176,7 @@ bool HIPDevice::should_use_graphics_interop()
       return true;
     }
   }
+#  endif
 
   return false;
 }

intern/cycles/device/hip/device_impl.h

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

intern/cycles/device/memory.cpp

@@ -44,45 +44,6 @@ 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,11 +281,16 @@ class device_memory {
 
   /* Only create through subclasses. */
   device_memory(Device *device, const char *name, MemoryType type);
-  device_memory(device_memory &&other) noexcept;
 
-  /* No copying allowed. */
+  /* 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. */
   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")
+    : device(device), queue(device), state(device, "__denoiser_state", true)
 {
 }
 
 OptiXDevice::OptiXDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
     : CUDADevice(info, stats, profiler),
       sbt_data(this, "__sbt", MEM_READ_ONLY),
-      launch_params(this, "__params"),
+      launch_params(this, "__params", false),
       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"),
+        guiding_buffer(device, "denoiser guiding passes buffer", true),
         num_samples(task.num_samples)
   {
     num_input_passes = 1;
@@ -522,9 +522,9 @@ class OptiXDevice::DenoiseContext {
       }
     }
 
-    const int num_guiding_passes = num_input_passes - 1;
+    use_guiding_passes = (num_input_passes - 1) > 0;
 
-    if (num_guiding_passes) {
+    if (use_guiding_passes) {
       if (task.allow_inplace_modification) {
         guiding_params.device_pointer = render_buffers->buffer.device_pointer;
 
@@ -577,6 +577,7 @@ 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;
 
@@ -708,7 +709,7 @@ void OptiXDevice::denoise_pass(DenoiseContext &context, PassType pass_type)
       return;
     }
   }
-  else if (!context.albedo_replaced_with_fake) {
+  else if (context.use_guiding_passes && !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.";
@@ -886,8 +887,7 @@ bool OptiXDevice::denoise_configure_if_needed(DenoiseContext &context)
   denoiser_.scratch_offset = sizes.stateSizeInBytes;
 
   /* Allocate denoiser state if tile size has changed since last setup. */
-  denoiser_.state.alloc_to_device(denoiser_.scratch_offset + denoiser_.scratch_size +
-                                  sizeof(float));
+  denoiser_.state.alloc_to_device(denoiser_.scratch_offset + denoiser_.scratch_size);
 
   /* Initialize denoiser state for the current tile size. */
   const OptixResult result = optixDenoiserSetup(
@@ -971,16 +971,6 @@ bool OptiXDevice::denoise_run(DenoiseContext &context, const DenoisePass &pass)
 
   /* Finally run denoising. */
   OptixDenoiserParams params = {}; /* All parameters are disabled/zero. */
-  params.hdrIntensity = denoiser_.state.device_pointer + denoiser_.scratch_offset +
-                        denoiser_.scratch_size;
-
-  optix_assert(
-      optixDenoiserComputeIntensity(denoiser_.optix_denoiser,
-                                    denoiser_.queue.stream(),
-                                    &color_layer,
-                                    params.hdrIntensity,
-                                    denoiser_.state.device_pointer + denoiser_.scratch_offset,
-                                    denoiser_.scratch_size));
 
   OptixDenoiserLayer image_layers = {};
   image_layers.input = color_layer;
@@ -1011,6 +1001,13 @@ 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);
@@ -1036,14 +1033,15 @@ 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");
+  device_only_memory<char> temp_mem(this, "optix temp as build mem", true);
   temp_mem.alloc_to_device(align_up(sizes.tempSizeInBytes, 8) + 8);
   if (!temp_mem.device_pointer) {
     /* Make sure temporary memory allocation succeeded. */
     return false;
   }
 
-  device_only_memory<char> &out_data = bvh->as_data;
+  /* Acceleration structure memory has to be allocated on the device (not allowed on the host). */
+  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);
@@ -1091,12 +1089,13 @@ 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");
+      device_only_memory<char> compacted_data(this, "optix compacted as", false);
       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));
@@ -1107,6 +1106,8 @@ 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.
+       */
     }
   }
 
@@ -1134,7 +1135,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;
     }
 
@@ -1195,7 +1196,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 Bezier spline. */
+              /* Convert Catmull-Rom data to B-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;
@@ -1355,9 +1356,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,
@@ -1390,8 +1391,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) {
@@ -1452,7 +1453,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,6 +23,7 @@
 #  include "device/optix/queue.h"
 #  include "device/optix/util.h"
 #  include "kernel/types.h"
+#  include "util/unique_ptr.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -76,7 +77,7 @@ class OptiXDevice : public CUDADevice {
   device_only_memory<KernelParamsOptiX> launch_params;
   OptixTraversableHandle tlas_handle = 0;
 
-  vector<device_only_memory<char>> delayed_free_bvh_memory;
+  vector<unique_ptr<device_only_memory<char>>> delayed_free_bvh_memory;
   thread_mutex delayed_free_bvh_mutex;
 
   class Denoiser {

intern/cycles/device/optix/queue.cpp

@@ -73,7 +73,8 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel, const int work_size, void *a
                         sizeof(device_ptr),
                         cuda_stream_));
 
-  if (kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE) {
+  if (kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST ||
+      kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE) {
     cuda_device_assert(
         cuda_device_,
         cuMemcpyHtoDAsync(launch_params_ptr + offsetof(KernelParamsOptiX, render_buffer),

intern/cycles/integrator/denoiser.cpp

@@ -33,7 +33,10 @@ unique_ptr<Denoiser> Denoiser::create(Device *path_trace_device, const DenoisePa
     return make_unique<OptiXDenoiser>(path_trace_device, params);
   }
 
-  return make_unique<OIDNDenoiser>(path_trace_device, params);
+  /* Always fallback to OIDN. */
+  DenoiseParams oidn_params = params;
+  oidn_params.type = DENOISER_OPENIMAGEDENOISE;
+  return make_unique<OIDNDenoiser>(path_trace_device, oidn_params);
 }
 
 Denoiser::Denoiser(Device *path_trace_device, const DenoiseParams &params)

intern/cycles/integrator/denoiser_oidn.cpp

@@ -47,9 +47,6 @@ 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:
@@ -547,7 +544,6 @@ 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)
 {
@@ -582,18 +578,20 @@ 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);
 
@@ -620,6 +618,11 @@ 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 = params.window_width / resolution_divider;
-  scaled_params.window_height = params.window_height / 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.full_x = params.full_x / resolution_divider;
   scaled_params.full_y = params.full_y / resolution_divider;
-  scaled_params.full_width = params.full_width / resolution_divider;
-  scaled_params.full_height = params.full_height / 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.update_offset_stride();
 
@@ -479,7 +479,11 @@ void PathTrace::set_denoiser_params(const DenoiseParams &params)
   }
 
   denoiser_ = Denoiser::create(device_, params);
-  denoiser_->is_cancelled_cb = [this]() { return is_cancel_requested(); };
+
+  /* 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; };
 }
 
 void PathTrace::set_adaptive_sampling(const AdaptiveSampling &adaptive_sampling)
@@ -847,7 +851,8 @@ void PathTrace::progress_update_if_needed(const RenderWork &render_work)
 {
   if (progress_ != nullptr) {
     const int2 tile_size = get_render_tile_size();
-    const int num_samples_added = tile_size.x * tile_size.y * render_work.path_trace.num_samples;
+    const uint64_t num_samples_added = uint64_t(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;
     progress_->add_samples(num_samples_added, current_sample);

intern/cycles/integrator/path_trace_work_cpu.cpp

@@ -77,8 +77,10 @@ void PathTraceWorkCPU::render_samples(RenderStatistics &statistics,
   const int64_t image_height = effective_buffer_params_.height;
   const int64_t total_pixels_num = image_width * image_height;
 
-  for (CPUKernelThreadGlobals &kernel_globals : kernel_thread_globals_) {
-    kernel_globals.start_profiling();
+  if (device_->profiler.active()) {
+    for (CPUKernelThreadGlobals &kernel_globals : kernel_thread_globals_) {
+      kernel_globals.start_profiling();
+    }
   }
 
   tbb::task_arena local_arena = local_tbb_arena_create(device_);
@@ -106,9 +108,10 @@ void PathTraceWorkCPU::render_samples(RenderStatistics &statistics,
       render_samples_full_pipeline(kernel_globals, work_tile, samples_num);
     });
   });
-
-  for (CPUKernelThreadGlobals &kernel_globals : kernel_thread_globals_) {
-    kernel_globals.stop_profiling();
+  if (device_->profiler.active()) {
+    for (CPUKernelThreadGlobals &kernel_globals : kernel_thread_globals_) {
+      kernel_globals.stop_profiling();
+    }
   }
 
   statistics.occupancy = 1.0f;

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -257,7 +257,8 @@ void PathTraceWorkGPU::render_samples(RenderStatistics &statistics,
    * 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,
@@ -437,7 +438,15 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
   DCHECK_LE(work_size, max_num_paths_);
 
   switch (kernel) {
-    case DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST:
+    case DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST: {
+      /* Closest ray intersection kernels with integrator state and render buffer. */
+      void *d_render_buffer = (void *)buffers_->buffer.device_pointer;
+      void *args[] = {&d_path_index, &d_render_buffer, const_cast<int *>(&work_size)};
+
+      queue_->enqueue(kernel, work_size, args);
+      break;
+    }
+
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK: {

intern/cycles/integrator/render_scheduler.cpp

@@ -827,6 +827,26 @@ 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/tile.cpp

@@ -46,7 +46,8 @@ ccl_device_inline uint round_up_to_power_of_two(uint x)
   return next_power_of_two(x);
 }
 
-TileSize tile_calculate_best_size(const int2 &image_size,
+TileSize tile_calculate_best_size(const bool accel_rt,
+                                  const int2 &image_size,
                                   const int num_samples,
                                   const int max_num_path_states,
                                   const float scrambling_distance)
@@ -73,7 +74,7 @@ TileSize tile_calculate_best_size(const int2 &image_size,
 
   TileSize tile_size;
   const int num_path_states_per_sample = max_num_path_states / num_samples;
-  if (scrambling_distance < 0.9f) {
+  if (scrambling_distance < 0.9f && accel_rt) {
     /* 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,7 +49,8 @@ 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 int2 &image_size,
+TileSize tile_calculate_best_size(const bool accel_rt,
+                                  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,6 +28,11 @@ 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;
@@ -59,7 +64,7 @@ void WorkTileScheduler::reset(const BufferParams &buffer_params,
 void WorkTileScheduler::reset_scheduler_state()
 {
   tile_size_ = tile_calculate_best_size(
-      image_size_px_, samples_num_, max_num_path_states_, scrambling_distance_);
+      accelerated_rt_, image_size_px_, samples_num_, max_num_path_states_, scrambling_distance_);
 
   VLOG(3) << "Will schedule tiles of size " << tile_size_;
 

intern/cycles/integrator/work_tile_scheduler.h

@@ -31,6 +31,9 @@ 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
@@ -54,6 +57,9 @@ 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

intern/cycles/kernel/CMakeLists.txt

@@ -565,6 +565,12 @@ 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()

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) {
+    if (threshold_squared >= 0.0f && !(label & LABEL_DIFFUSE)) {
       if (bsdf_get_specular_roughness_squared(sc) <= threshold_squared) {
         label |= LABEL_TRANSMIT_TRANSPARENT;
       }

intern/cycles/kernel/device/cpu/kernel_arch.h

@@ -37,7 +37,7 @@
 
 KERNEL_INTEGRATOR_INIT_FUNCTION(init_from_camera);
 KERNEL_INTEGRATOR_INIT_FUNCTION(init_from_bake);
-KERNEL_INTEGRATOR_FUNCTION(intersect_closest);
+KERNEL_INTEGRATOR_SHADE_FUNCTION(intersect_closest);
 KERNEL_INTEGRATOR_FUNCTION(intersect_shadow);
 KERNEL_INTEGRATOR_FUNCTION(intersect_subsurface);
 KERNEL_INTEGRATOR_FUNCTION(intersect_volume_stack);

intern/cycles/kernel/device/cpu/kernel_arch_impl.h

@@ -112,7 +112,7 @@ CCL_NAMESPACE_BEGIN
 
 DEFINE_INTEGRATOR_INIT_KERNEL(init_from_camera)
 DEFINE_INTEGRATOR_INIT_KERNEL(init_from_bake)
-DEFINE_INTEGRATOR_KERNEL(intersect_closest)
+DEFINE_INTEGRATOR_SHADE_KERNEL(intersect_closest)
 DEFINE_INTEGRATOR_KERNEL(intersect_subsurface)
 DEFINE_INTEGRATOR_KERNEL(intersect_volume_stack)
 DEFINE_INTEGRATOR_SHADE_KERNEL(shade_background)

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

@@ -116,13 +116,15 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
 }
 
 ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
-    kernel_gpu_integrator_intersect_closest(const int *path_index_array, const int work_size)
+    kernel_gpu_integrator_intersect_closest(const int *path_index_array,
+                                            ccl_global float *render_buffer,
+                                            const int work_size)
 {
   const int global_index = ccl_gpu_global_id_x();
 
   if (global_index < work_size) {
     const int state = (path_index_array) ? path_index_array[global_index] : global_index;
-    integrator_intersect_closest(NULL, state);
+    integrator_intersect_closest(NULL, state, render_buffer);
   }
 }
 
@@ -486,6 +488,26 @@ ccl_device_inline void kernel_gpu_film_convert_common(const KernelFilmConvert *k
   processor(kfilm_convert, buffer, pixel);
 }
 
+ccl_device_inline void kernel_gpu_film_convert_half_write(ccl_global uchar4 *rgba,
+                                                          const int rgba_offset,
+                                                          const int rgba_stride,
+                                                          const int x,
+                                                          const int y,
+                                                          const half4 half_pixel)
+{
+  /* Work around HIP issue with half float display, see T92972. */
+#ifdef __KERNEL_HIP__
+  ccl_global half *out = ((ccl_global half *)rgba) + (rgba_offset + y * rgba_stride + x) * 4;
+  out[0] = half_pixel.x;
+  out[1] = half_pixel.y;
+  out[2] = half_pixel.z;
+  out[3] = half_pixel.w;
+#else
+  ccl_global half4 *out = ((ccl_global half4 *)rgba) + rgba_offset + y * rgba_stride + x;
+  *out = half_pixel;
+#endif
+}
+
 /* Common implementation for half4 destination and 4-channel input pass. */
 template<typename Processor>
 ccl_device_inline void kernel_gpu_film_convert_half_rgba_common_rgba(
@@ -516,8 +538,9 @@ ccl_device_inline void kernel_gpu_film_convert_half_rgba_common_rgba(
 
   film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel);
 
-  ccl_global half4 *out = ((ccl_global half4 *)rgba) + rgba_offset + y * rgba_stride + x;
-  *out = float4_to_half4_display(make_float4(pixel[0], pixel[1], pixel[2], pixel[3]));
+  const half4 half_pixel = float4_to_half4_display(
+      make_float4(pixel[0], pixel[1], pixel[2], pixel[3]));
+  kernel_gpu_film_convert_half_write(rgba, rgba_offset, rgba_stride, x, y, half_pixel);
 }
 
 /* Common implementation for half4 destination and 3-channel input pass. */

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

@@ -29,17 +29,20 @@ ccl_device_inline void get_work_pixel(ccl_global const KernelWorkTile *tile,
                                       ccl_private uint *y,
                                       ccl_private uint *sample)
 {
-#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 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;
+  }
 
   uint y_offset = pixel_offset / tile->w;
   uint x_offset = pixel_offset - y_offset * tile->w;

intern/cycles/kernel/device/optix/kernel.cu

@@ -57,7 +57,7 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()
   const int global_index = optixGetLaunchIndex().x;
   const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
                                                        global_index;
-  integrator_intersect_closest(nullptr, path_index);
+  integrator_intersect_closest(nullptr, path_index, __params.render_buffer);
 }
 
 extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_shadow()

intern/cycles/kernel/film/accumulate.h

@@ -33,62 +33,72 @@ CCL_NAMESPACE_BEGIN
  * them separately. */
 
 ccl_device_inline void bsdf_eval_init(ccl_private BsdfEval *eval,
-                                      const bool is_diffuse,
+                                      const ClosureType closure_type,
                                       float3 value)
 {
   eval->diffuse = zero_float3();
   eval->glossy = zero_float3();
 
-  if (is_diffuse) {
+  if (CLOSURE_IS_BSDF_DIFFUSE(closure_type)) {
     eval->diffuse = value;
   }
-  else {
+  else if (CLOSURE_IS_BSDF_GLOSSY(closure_type)) {
     eval->glossy = value;
   }
+
+  eval->sum = value;
 }
 
 ccl_device_inline void bsdf_eval_accum(ccl_private BsdfEval *eval,
-                                       const bool is_diffuse,
-                                       float3 value,
-                                       float mis_weight)
+                                       const ClosureType closure_type,
+                                       float3 value)
 {
-  value *= mis_weight;
-
-  if (is_diffuse) {
+  if (CLOSURE_IS_BSDF_DIFFUSE(closure_type)) {
     eval->diffuse += value;
   }
-  else {
+  else if (CLOSURE_IS_BSDF_GLOSSY(closure_type)) {
     eval->glossy += value;
   }
+
+  eval->sum += value;
 }
 
 ccl_device_inline bool bsdf_eval_is_zero(ccl_private BsdfEval *eval)
 {
-  return is_zero(eval->diffuse) && is_zero(eval->glossy);
+  return is_zero(eval->sum);
 }
 
 ccl_device_inline void bsdf_eval_mul(ccl_private BsdfEval *eval, float value)
 {
   eval->diffuse *= value;
   eval->glossy *= value;
+  eval->sum *= value;
 }
 
 ccl_device_inline void bsdf_eval_mul3(ccl_private BsdfEval *eval, float3 value)
 {
   eval->diffuse *= value;
   eval->glossy *= value;
+  eval->sum *= value;
 }
 
 ccl_device_inline float3 bsdf_eval_sum(ccl_private const BsdfEval *eval)
 {
-  return eval->diffuse + eval->glossy;
+  return eval->sum;
 }
 
-ccl_device_inline float3 bsdf_eval_diffuse_glossy_ratio(ccl_private const BsdfEval *eval)
+ccl_device_inline float3 bsdf_eval_pass_diffuse_weight(ccl_private const BsdfEval *eval)
 {
-  /* Ratio of diffuse and glossy to recover proportions for writing to render pass.
+  /* Ratio of diffuse weight to recover proportions for writing to render pass.
    * We assume reflection, transmission and volume scatter to be exclusive. */
-  return safe_divide_float3_float3(eval->diffuse, eval->diffuse + eval->glossy);
+  return safe_divide_float3_float3(eval->diffuse, eval->sum);
+}
+
+ccl_device_inline float3 bsdf_eval_pass_glossy_weight(ccl_private const BsdfEval *eval)
+{
+  /* Ratio of glossy weight to recover proportions for writing to render pass.
+   * We assume reflection, transmission and volume scatter to be exclusive. */
+  return safe_divide_float3_float3(eval->glossy, eval->sum);
 }
 
 /* --------------------------------------------------------------------
@@ -351,37 +361,47 @@ ccl_device_inline void kernel_accum_emission_or_background_pass(KernelGlobals kg
     /* Directly visible, write to emission or background pass. */
     pass_offset = pass;
   }
-  else if (path_flag & (PATH_RAY_REFLECT_PASS | PATH_RAY_TRANSMISSION_PASS)) {
-    /* Indirectly visible through reflection. */
-    const int glossy_pass_offset = (path_flag & PATH_RAY_REFLECT_PASS) ?
-                                       ((INTEGRATOR_STATE(state, path, bounce) == 1) ?
-                                            kernel_data.film.pass_glossy_direct :
-                                            kernel_data.film.pass_glossy_indirect) :
-                                       ((INTEGRATOR_STATE(state, path, bounce) == 1) ?
-                                            kernel_data.film.pass_transmission_direct :
-                                            kernel_data.film.pass_transmission_indirect);
+  else if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
+    if (path_flag & PATH_RAY_SURFACE_PASS) {
+      /* Indirectly visible through reflection. */
+      const float3 diffuse_weight = INTEGRATOR_STATE(state, path, pass_diffuse_weight);
+      const float3 glossy_weight = INTEGRATOR_STATE(state, path, pass_glossy_weight);
 
-    if (glossy_pass_offset != PASS_UNUSED) {
-      /* Glossy is a subset of the throughput, reconstruct it here using the
-       * diffuse-glossy ratio. */
-      const float3 ratio = INTEGRATOR_STATE(state, path, diffuse_glossy_ratio);
-      const float3 glossy_contribution = (one_float3() - ratio) * contribution;
-      kernel_write_pass_float3(buffer + glossy_pass_offset, glossy_contribution);
-    }
+      /* Glossy */
+      const int glossy_pass_offset = ((INTEGRATOR_STATE(state, path, bounce) == 1) ?
+                                          kernel_data.film.pass_glossy_direct :
+                                          kernel_data.film.pass_glossy_indirect);
+      if (glossy_pass_offset != PASS_UNUSED) {
+        kernel_write_pass_float3(buffer + glossy_pass_offset, glossy_weight * contribution);
+      }
 
-    /* Reconstruct diffuse subset of throughput. */
-    pass_offset = (INTEGRATOR_STATE(state, path, bounce) == 1) ?
-                      kernel_data.film.pass_diffuse_direct :
-                      kernel_data.film.pass_diffuse_indirect;
-    if (pass_offset != PASS_UNUSED) {
-      contribution *= INTEGRATOR_STATE(state, path, diffuse_glossy_ratio);
+      /* Transmission */
+      const int transmission_pass_offset = ((INTEGRATOR_STATE(state, path, bounce) == 1) ?
+                                                kernel_data.film.pass_transmission_direct :
+                                                kernel_data.film.pass_transmission_indirect);
+
+      if (transmission_pass_offset != PASS_UNUSED) {
+        /* Transmission is what remains if not diffuse and glossy, not stored explicitly to save
+         * GPU memory. */
+        const float3 transmission_weight = one_float3() - diffuse_weight - glossy_weight;
+        kernel_write_pass_float3(buffer + transmission_pass_offset,
+                                 transmission_weight * contribution);
+      }
+
+      /* Reconstruct diffuse subset of throughput. */
+      pass_offset = (INTEGRATOR_STATE(state, path, bounce) == 1) ?
+                        kernel_data.film.pass_diffuse_direct :
+                        kernel_data.film.pass_diffuse_indirect;
+      if (pass_offset != PASS_UNUSED) {
+        contribution *= diffuse_weight;
+      }
+    }
+    else if (path_flag & PATH_RAY_VOLUME_PASS) {
+      /* Indirectly visible through volume. */
+      pass_offset = (INTEGRATOR_STATE(state, path, bounce) == 1) ?
+                        kernel_data.film.pass_volume_direct :
+                        kernel_data.film.pass_volume_indirect;
     }
-  }
-  else if (path_flag & PATH_RAY_VOLUME_PASS) {
-    /* Indirectly visible through volume. */
-    pass_offset = (INTEGRATOR_STATE(state, path, bounce) == 1) ?
-                      kernel_data.film.pass_volume_direct :
-                      kernel_data.film.pass_volume_indirect;
   }
 
   /* Single write call for GPU coherence. */
@@ -426,49 +446,60 @@ ccl_device_inline void kernel_accum_light(KernelGlobals kg,
 #ifdef __PASSES__
   if (kernel_data.film.light_pass_flag & PASS_ANY) {
     const uint32_t path_flag = INTEGRATOR_STATE(state, shadow_path, flag);
-    int pass_offset = PASS_UNUSED;
 
-    if (path_flag & (PATH_RAY_REFLECT_PASS | PATH_RAY_TRANSMISSION_PASS)) {
-      /* Indirectly visible through reflection. */
-      const int glossy_pass_offset = (path_flag & PATH_RAY_REFLECT_PASS) ?
-                                         ((INTEGRATOR_STATE(state, shadow_path, bounce) == 0) ?
-                                              kernel_data.film.pass_glossy_direct :
-                                              kernel_data.film.pass_glossy_indirect) :
-                                         ((INTEGRATOR_STATE(state, shadow_path, bounce) == 0) ?
-                                              kernel_data.film.pass_transmission_direct :
-                                              kernel_data.film.pass_transmission_indirect);
+    if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
+      int pass_offset = PASS_UNUSED;
 
-      if (glossy_pass_offset != PASS_UNUSED) {
-        /* Glossy is a subset of the throughput, reconstruct it here using the
-         * diffuse-glossy ratio. */
-        const float3 ratio = INTEGRATOR_STATE(state, shadow_path, diffuse_glossy_ratio);
-        const float3 glossy_contribution = (one_float3() - ratio) * contribution;
-        kernel_write_pass_float3(buffer + glossy_pass_offset, glossy_contribution);
+      if (path_flag & PATH_RAY_SURFACE_PASS) {
+        /* Indirectly visible through reflection. */
+        const float3 diffuse_weight = INTEGRATOR_STATE(state, shadow_path, pass_diffuse_weight);
+        const float3 glossy_weight = INTEGRATOR_STATE(state, shadow_path, pass_glossy_weight);
+
+        /* Glossy */
+        const int glossy_pass_offset = ((INTEGRATOR_STATE(state, shadow_path, bounce) == 0) ?
+                                            kernel_data.film.pass_glossy_direct :
+                                            kernel_data.film.pass_glossy_indirect);
+        if (glossy_pass_offset != PASS_UNUSED) {
+          kernel_write_pass_float3(buffer + glossy_pass_offset, glossy_weight * contribution);
+        }
+
+        /* Transmission */
+        const int transmission_pass_offset = ((INTEGRATOR_STATE(state, shadow_path, bounce) == 0) ?
+                                                  kernel_data.film.pass_transmission_direct :
+                                                  kernel_data.film.pass_transmission_indirect);
+
+        if (transmission_pass_offset != PASS_UNUSED) {
+          /* Transmission is what remains if not diffuse and glossy, not stored explicitly to save
+           * GPU memory. */
+          const float3 transmission_weight = one_float3() - diffuse_weight - glossy_weight;
+          kernel_write_pass_float3(buffer + transmission_pass_offset,
+                                   transmission_weight * contribution);
+        }
+
+        /* Reconstruct diffuse subset of throughput. */
+        pass_offset = (INTEGRATOR_STATE(state, shadow_path, bounce) == 0) ?
+                          kernel_data.film.pass_diffuse_direct :
+                          kernel_data.film.pass_diffuse_indirect;
+        if (pass_offset != PASS_UNUSED) {
+          contribution *= diffuse_weight;
+        }
+      }
+      else if (path_flag & PATH_RAY_VOLUME_PASS) {
+        /* Indirectly visible through volume. */
+        pass_offset = (INTEGRATOR_STATE(state, shadow_path, bounce) == 0) ?
+                          kernel_data.film.pass_volume_direct :
+                          kernel_data.film.pass_volume_indirect;
       }
 
-      /* Reconstruct diffuse subset of throughput. */
-      pass_offset = (INTEGRATOR_STATE(state, shadow_path, bounce) == 0) ?
-                        kernel_data.film.pass_diffuse_direct :
-                        kernel_data.film.pass_diffuse_indirect;
+      /* Single write call for GPU coherence. */
       if (pass_offset != PASS_UNUSED) {
-        contribution *= INTEGRATOR_STATE(state, shadow_path, diffuse_glossy_ratio);
+        kernel_write_pass_float3(buffer + pass_offset, contribution);
       }
     }
-    else if (path_flag & PATH_RAY_VOLUME_PASS) {
-      /* Indirectly visible through volume. */
-      pass_offset = (INTEGRATOR_STATE(state, shadow_path, bounce) == 0) ?
-                        kernel_data.film.pass_volume_direct :
-                        kernel_data.film.pass_volume_indirect;
-    }
-
-    /* Single write call for GPU coherence. */
-    if (pass_offset != PASS_UNUSED) {
-      kernel_write_pass_float3(buffer + pass_offset, contribution);
-    }
 
     /* Write shadow pass. */
     if (kernel_data.film.pass_shadow != PASS_UNUSED && (path_flag & PATH_RAY_SHADOW_FOR_LIGHT) &&
-        (path_flag & PATH_RAY_CAMERA)) {
+        (path_flag & PATH_RAY_TRANSPARENT_BACKGROUND)) {
       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

@@ -160,40 +160,6 @@ ccl_device_forceinline void kernel_write_denoising_features_volume(KernelGlobals
 }
 #endif /* __DENOISING_FEATURES__ */
 
-#ifdef __SHADOW_CATCHER__
-
-/* Write shadow catcher passes on a bounce from the shadow catcher object. */
-ccl_device_forceinline void kernel_write_shadow_catcher_bounce_data(
-    KernelGlobals kg,
-    IntegratorState state,
-    ccl_private const ShaderData *sd,
-    ccl_global float *ccl_restrict render_buffer)
-{
-  if (!kernel_data.integrator.has_shadow_catcher) {
-    return;
-  }
-
-  kernel_assert(kernel_data.film.pass_shadow_catcher_sample_count != PASS_UNUSED);
-  kernel_assert(kernel_data.film.pass_shadow_catcher_matte != PASS_UNUSED);
-
-  if (!kernel_shadow_catcher_is_path_split_bounce(kg, state, sd->object_flag)) {
-    return;
-  }
-
-  ccl_global float *buffer = kernel_pass_pixel_render_buffer(kg, state, render_buffer);
-
-  /* Count sample for the shadow catcher object. */
-  kernel_write_pass_float(buffer + kernel_data.film.pass_shadow_catcher_sample_count, 1.0f);
-
-  /* Since the split is done, the sample does not contribute to the matte, so accumulate it as
-   * transparency to the matte. */
-  const float3 throughput = INTEGRATOR_STATE(state, path, throughput);
-  kernel_write_pass_float(buffer + kernel_data.film.pass_shadow_catcher_matte + 3,
-                          average(throughput));
-}
-
-#endif /* __SHADOW_CATCHER__ */
-
 ccl_device_inline size_t kernel_write_id_pass(ccl_global float *ccl_restrict buffer,
                                               size_t depth,
                                               float id,
@@ -211,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_CAMERA)) {
+  if (!(path_flag & PATH_RAY_TRANSPARENT_BACKGROUND)) {
     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 - average(shadow_catcher)) + alpha;
+  const float alpha_matte = (1.0f - alpha) * (1.0f - saturatef(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

@@ -70,14 +70,16 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg,
   /* Setup render buffers. */
   const int index = INTEGRATOR_STATE(state, path, render_pixel_index);
   const int pass_stride = kernel_data.film.pass_stride;
-  render_buffer += index * pass_stride;
+  ccl_global float *buffer = render_buffer + index * pass_stride;
 
-  ccl_global float *primitive = render_buffer + kernel_data.film.pass_bake_primitive;
-  ccl_global float *differential = render_buffer + kernel_data.film.pass_bake_differential;
+  ccl_global float *primitive = buffer + kernel_data.film.pass_bake_primitive;
+  ccl_global float *differential = 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/intersect_closest.h

@@ -88,7 +88,10 @@ ccl_device_forceinline bool integrator_intersect_terminate(KernelGlobals kg,
 #ifdef __SHADOW_CATCHER__
 /* Split path if a shadow catcher was hit. */
 ccl_device_forceinline void integrator_split_shadow_catcher(
-    KernelGlobals kg, IntegratorState state, ccl_private const Intersection *ccl_restrict isect)
+    KernelGlobals kg,
+    IntegratorState state,
+    ccl_private const Intersection *ccl_restrict isect,
+    ccl_global float *ccl_restrict render_buffer)
 {
   /* Test if we hit a shadow catcher object, and potentially split the path to continue tracing two
    * paths from here. */
@@ -97,6 +100,8 @@ ccl_device_forceinline void integrator_split_shadow_catcher(
     return;
   }
 
+  kernel_write_shadow_catcher_bounce_data(kg, state, render_buffer);
+
   /* Mark state as having done a shadow catcher split so that it stops contributing to
    * the shadow catcher matte pass, but keeps contributing to the combined pass. */
   INTEGRATOR_STATE_WRITE(state, path, flag) |= PATH_RAY_SHADOW_CATCHER_HIT;
@@ -191,6 +196,7 @@ ccl_device_forceinline void integrator_intersect_next_kernel(
     KernelGlobals kg,
     IntegratorState state,
     ccl_private const Intersection *ccl_restrict isect,
+    ccl_global float *ccl_restrict render_buffer,
     const bool hit)
 {
   /* Continue with volume kernel if we are inside a volume, regardless if we hit anything. */
@@ -233,7 +239,7 @@ ccl_device_forceinline void integrator_intersect_next_kernel(
 
 #ifdef __SHADOW_CATCHER__
         /* Handle shadow catcher. */
-        integrator_split_shadow_catcher(kg, state, isect);
+        integrator_split_shadow_catcher(kg, state, isect, render_buffer);
 #endif
       }
       else {
@@ -253,7 +259,10 @@ ccl_device_forceinline void integrator_intersect_next_kernel(
  * volume shading and termination testing have already been done. */
 template<uint32_t current_kernel>
 ccl_device_forceinline void integrator_intersect_next_kernel_after_volume(
-    KernelGlobals kg, IntegratorState state, ccl_private const Intersection *ccl_restrict isect)
+    KernelGlobals kg,
+    IntegratorState state,
+    ccl_private const Intersection *ccl_restrict isect,
+    ccl_global float *ccl_restrict render_buffer)
 {
   if (isect->prim != PRIM_NONE) {
     /* Hit a surface, continue with light or surface kernel. */
@@ -278,7 +287,7 @@ ccl_device_forceinline void integrator_intersect_next_kernel_after_volume(
 
 #ifdef __SHADOW_CATCHER__
       /* Handle shadow catcher. */
-      integrator_split_shadow_catcher(kg, state, isect);
+      integrator_split_shadow_catcher(kg, state, isect, render_buffer);
 #endif
       return;
     }
@@ -290,7 +299,9 @@ ccl_device_forceinline void integrator_intersect_next_kernel_after_volume(
   }
 }
 
-ccl_device void integrator_intersect_closest(KernelGlobals kg, IntegratorState state)
+ccl_device void integrator_intersect_closest(KernelGlobals kg,
+                                             IntegratorState state,
+                                             ccl_global float *ccl_restrict render_buffer)
 {
   PROFILING_INIT(kg, PROFILING_INTERSECT_CLOSEST);
 
@@ -341,7 +352,7 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg, IntegratorState s
 
   /* Setup up next kernel to be executed. */
   integrator_intersect_next_kernel<DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST>(
-      kg, state, &isect, hit);
+      kg, state, &isect, render_buffer, hit);
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/integrator/megakernel.h

@@ -76,7 +76,7 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
     if (queued_kernel) {
       switch (queued_kernel) {
         case DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST:
-          integrator_intersect_closest(kg, state);
+          integrator_intersect_closest(kg, state, render_buffer);
           break;
         case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND:
           integrator_shade_background(kg, state, render_buffer);

intern/cycles/kernel/integrator/path_state.h

@@ -70,6 +70,9 @@ 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(
@@ -122,7 +125,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 (bounce == 1) {
+    if (!(flag & PATH_RAY_ANY_PASS)) {
       flag |= PATH_RAY_VOLUME_PASS;
     }
 
@@ -184,8 +187,8 @@ ccl_device_inline void path_state_next(KernelGlobals kg, IntegratorState state,
     }
 
     /* Render pass categories. */
-    if (bounce == 1) {
-      flag |= (label & LABEL_TRANSMIT) ? PATH_RAY_TRANSMISSION_PASS : PATH_RAY_REFLECT_PASS;
+    if (!(flag & PATH_RAY_ANY_PASS) && !(flag & PATH_RAY_TRANSPARENT_BACKGROUND)) {
+      flag |= PATH_RAY_SURFACE_PASS;
     }
   }
 
@@ -208,9 +211,7 @@ ccl_device_inline bool path_state_volume_next(IntegratorState state)
   }
 
   /* Random number generator next bounce. */
-  if (volume_bounds_bounce > 1) {
-    INTEGRATOR_STATE_WRITE(state, path, rng_offset) += PRNG_BOUNCE_NUM;
-  }
+  INTEGRATOR_STATE_WRITE(state, path, rng_offset) += PRNG_BOUNCE_NUM;
 
   return true;
 }

intern/cycles/kernel/integrator/shade_surface.h

@@ -191,14 +191,18 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
   const uint16_t transparent_bounce = INTEGRATOR_STATE(state, path, transparent_bounce);
   uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
   shadow_flag |= (is_light) ? PATH_RAY_SHADOW_FOR_LIGHT : 0;
-  shadow_flag |= (is_transmission) ? PATH_RAY_TRANSMISSION_PASS : PATH_RAY_REFLECT_PASS;
+  shadow_flag |= PATH_RAY_SURFACE_PASS;
   const float3 throughput = INTEGRATOR_STATE(state, path, throughput) * bsdf_eval_sum(&bsdf_eval);
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
-    const float3 diffuse_glossy_ratio = (bounce == 0) ?
-                                            bsdf_eval_diffuse_glossy_ratio(&bsdf_eval) :
-                                            INTEGRATOR_STATE(state, path, diffuse_glossy_ratio);
-    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, diffuse_glossy_ratio) = diffuse_glossy_ratio;
+    const float3 pass_diffuse_weight = (bounce == 0) ?
+                                           bsdf_eval_pass_diffuse_weight(&bsdf_eval) :
+                                           INTEGRATOR_STATE(state, path, pass_diffuse_weight);
+    const float3 pass_glossy_weight = (bounce == 0) ?
+                                          bsdf_eval_pass_glossy_weight(&bsdf_eval) :
+                                          INTEGRATOR_STATE(state, path, pass_glossy_weight);
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = pass_diffuse_weight;
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = pass_glossy_weight;
   }
 
   INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
@@ -283,7 +287,9 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
     if (INTEGRATOR_STATE(state, path, bounce) == 0) {
-      INTEGRATOR_STATE_WRITE(state, path, diffuse_glossy_ratio) = bsdf_eval_diffuse_glossy_ratio(
+      INTEGRATOR_STATE_WRITE(state, path, pass_diffuse_weight) = bsdf_eval_pass_diffuse_weight(
+          &bsdf_eval);
+      INTEGRATOR_STATE_WRITE(state, path, pass_glossy_weight) = bsdf_eval_pass_glossy_weight(
           &bsdf_eval);
     }
   }
@@ -445,7 +451,7 @@ ccl_device bool integrate_surface(KernelGlobals kg,
     }
 #endif
 
-    shader_prepare_surface_closures(kg, state, &sd);
+    shader_prepare_surface_closures(kg, state, &sd, path_flag);
 
 #ifdef __HOLDOUT__
     /* Evaluate holdout. */
@@ -492,10 +498,6 @@ ccl_device bool integrate_surface(KernelGlobals kg,
     kernel_write_denoising_features_surface(kg, state, &sd, render_buffer);
 #endif
 
-#ifdef __SHADOW_CATCHER__
-    kernel_write_shadow_catcher_bounce_data(kg, state, &sd, render_buffer);
-#endif
-
     /* Direct light. */
     PROFILING_EVENT(PROFILING_SHADE_SURFACE_DIRECT_LIGHT);
     integrate_surface_direct_light(kg, state, &sd, &rng_state);

intern/cycles/kernel/integrator/shade_volume.h

@@ -263,6 +263,12 @@ 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,
@@ -437,7 +443,8 @@ 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) {
+    if (result.direct_t >= vstate.start_t && result.direct_t <= vstate.end_t &&
+        vstate.equiangular_pdf > VOLUME_SAMPLE_PDF_CUTOFF) {
       const float new_dt = result.direct_t - vstate.start_t;
       const float3 new_transmittance = volume_color_transmittance(coeff.sigma_t, new_dt);
 
@@ -474,26 +481,28 @@ 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);
 
-      /* 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.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);
 
-      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;
+          }
         }
       }
     }
@@ -694,8 +703,10 @@ 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);
 
-  light_distribution_sample_from_volume_segment(
-      kg, light_u, light_v, sd->time, sd->P, bounce, path_flag, ls);
+  if (!light_distribution_sample_from_volume_segment(
+          kg, light_u, light_v, sd->time, sd->P, bounce, path_flag, ls)) {
+    return false;
+  }
 
   if (ls->shader & SHADER_EXCLUDE_SCATTER) {
     return false;
@@ -794,10 +805,11 @@ ccl_device_forceinline void integrate_volume_direct_light(
   const float3 throughput_phase = throughput * bsdf_eval_sum(&phase_eval);
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
-    const float3 diffuse_glossy_ratio = (bounce == 0) ?
-                                            one_float3() :
-                                            INTEGRATOR_STATE(state, path, diffuse_glossy_ratio);
-    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, diffuse_glossy_ratio) = diffuse_glossy_ratio;
+    const float3 pass_diffuse_weight = (bounce == 0) ?
+                                           one_float3() :
+                                           INTEGRATOR_STATE(state, path, pass_diffuse_weight);
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = pass_diffuse_weight;
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = zero_float3();
   }
 
   INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
@@ -876,7 +888,8 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
   INTEGRATOR_STATE_WRITE(state, path, throughput) = throughput_phase;
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
-    INTEGRATOR_STATE_WRITE(state, path, diffuse_glossy_ratio) = one_float3();
+    INTEGRATOR_STATE_WRITE(state, path, pass_diffuse_weight) = one_float3();
+    INTEGRATOR_STATE_WRITE(state, path, pass_glossy_weight) = zero_float3();
   }
 
   /* Update path state */
@@ -1024,7 +1037,7 @@ ccl_device void integrator_shade_volume(KernelGlobals kg,
   else {
     /* Continue to background, light or surface. */
     integrator_intersect_next_kernel_after_volume<DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME>(
-        kg, state, &isect);
+        kg, state, &isect, render_buffer);
     return;
   }
 #endif /* __VOLUME__ */

intern/cycles/kernel/integrator/shader_eval.h

@@ -105,8 +105,42 @@ ccl_device_inline void shader_copy_volume_phases(ccl_private ShaderVolumePhases
 
 ccl_device_inline void shader_prepare_surface_closures(KernelGlobals kg,
                                                        ConstIntegratorState state,
-                                                       ccl_private ShaderData *sd)
+                                                       ccl_private ShaderData *sd,
+                                                       const uint32_t path_flag)
 {
+  /* Filter out closures. */
+  if (kernel_data.integrator.filter_closures) {
+    if (kernel_data.integrator.filter_closures & FILTER_CLOSURE_EMISSION) {
+      sd->closure_emission_background = zero_float3();
+    }
+
+    if (kernel_data.integrator.filter_closures & FILTER_CLOSURE_DIRECT_LIGHT) {
+      sd->flag &= ~SD_BSDF_HAS_EVAL;
+    }
+
+    if (path_flag & PATH_RAY_CAMERA) {
+      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;
+        }
+        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;
+        }
+      }
+    }
+  }
+
   /* Defensive sampling.
    *
    * We can likely also do defensive sampling at deeper bounces, particularly
@@ -209,8 +243,7 @@ ccl_device_inline float _shader_bsdf_multi_eval(KernelGlobals kg,
         float3 eval = bsdf_eval(kg, sd, sc, omega_in, is_transmission, &bsdf_pdf);
 
         if (bsdf_pdf != 0.0f) {
-          const bool is_diffuse = CLOSURE_IS_BSDF_DIFFUSE(sc->type);
-          bsdf_eval_accum(result_eval, is_diffuse, eval * sc->weight, 1.0f);
+          bsdf_eval_accum(result_eval, sc->type, eval * sc->weight);
           sum_pdf += bsdf_pdf * sc->sample_weight;
         }
       }
@@ -235,7 +268,7 @@ ccl_device_inline
                      ccl_private BsdfEval *bsdf_eval,
                      const uint light_shader_flags)
 {
-  bsdf_eval_init(bsdf_eval, false, zero_float3());
+  bsdf_eval_init(bsdf_eval, CLOSURE_NONE_ID, zero_float3());
 
   return _shader_bsdf_multi_eval(
       kg, sd, omega_in, is_transmission, NULL, bsdf_eval, 0.0f, 0.0f, light_shader_flags);
@@ -328,8 +361,7 @@ ccl_device int shader_bsdf_sample_closure(KernelGlobals kg,
   label = bsdf_sample(kg, sd, sc, randu, randv, &eval, omega_in, domega_in, pdf);
 
   if (*pdf != 0.0f) {
-    const bool is_diffuse = CLOSURE_IS_BSDF_DIFFUSE(sc->type);
-    bsdf_eval_init(bsdf_eval, is_diffuse, eval * sc->weight);
+    bsdf_eval_init(bsdf_eval, sc->type, eval * sc->weight);
 
     if (sd->num_closure > 1) {
       const bool is_transmission = shader_bsdf_is_transmission(sd, *omega_in);
@@ -655,7 +687,7 @@ ccl_device_inline float _shader_volume_phase_multi_eval(
     float3 eval = volume_phase_eval(sd, svc, omega_in, &phase_pdf);
 
     if (phase_pdf != 0.0f) {
-      bsdf_eval_accum(result_eval, false, eval, 1.0f);
+      bsdf_eval_accum(result_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, eval);
       sum_pdf += phase_pdf * svc->sample_weight;
     }
 
@@ -671,7 +703,7 @@ ccl_device float shader_volume_phase_eval(KernelGlobals kg,
                                           const float3 omega_in,
                                           ccl_private BsdfEval *phase_eval)
 {
-  bsdf_eval_init(phase_eval, false, zero_float3());
+  bsdf_eval_init(phase_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, zero_float3());
 
   return _shader_volume_phase_multi_eval(sd, phases, omega_in, -1, phase_eval, 0.0f, 0.0f);
 }
@@ -729,7 +761,7 @@ ccl_device int shader_volume_phase_sample(KernelGlobals kg,
   label = volume_phase_sample(sd, svc, randu, randv, &eval, omega_in, domega_in, pdf);
 
   if (*pdf != 0.0f) {
-    bsdf_eval_init(phase_eval, false, eval);
+    bsdf_eval_init(phase_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, eval);
   }
 
   return label;
@@ -752,7 +784,7 @@ ccl_device int shader_phase_sample_closure(KernelGlobals kg,
   label = volume_phase_sample(sd, sc, randu, randv, &eval, omega_in, domega_in, pdf);
 
   if (*pdf != 0.0f)
-    bsdf_eval_init(phase_eval, false, eval);
+    bsdf_eval_init(phase_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, eval);
 
   return label;
 }

intern/cycles/kernel/integrator/shadow_catcher.h

@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include "kernel/film/write_passes.h"
 #include "kernel/integrator/path_state.h"
 #include "kernel/integrator/state_util.h"
 
@@ -47,7 +48,7 @@ ccl_device_inline bool kernel_shadow_catcher_is_path_split_bounce(KernelGlobals
     return false;
   }
 
-  if (path_flag & PATH_RAY_SHADOW_CATCHER_PASS) {
+  if (path_flag & PATH_RAY_SHADOW_CATCHER_HIT) {
     return false;
   }
 
@@ -88,6 +89,28 @@ ccl_device_forceinline bool kernel_shadow_catcher_is_object_pass(const uint32_t
   return path_flag & PATH_RAY_SHADOW_CATCHER_PASS;
 }
 
+/* Write shadow catcher passes on a bounce from the shadow catcher object. */
+ccl_device_forceinline void kernel_write_shadow_catcher_bounce_data(
+    KernelGlobals kg, IntegratorState state, ccl_global float *ccl_restrict render_buffer)
+{
+  kernel_assert(kernel_data.film.pass_shadow_catcher_sample_count != PASS_UNUSED);
+  kernel_assert(kernel_data.film.pass_shadow_catcher_matte != PASS_UNUSED);
+
+  const uint32_t render_pixel_index = INTEGRATOR_STATE(state, path, render_pixel_index);
+  const uint64_t render_buffer_offset = (uint64_t)render_pixel_index *
+                                        kernel_data.film.pass_stride;
+  ccl_global float *buffer = render_buffer + render_buffer_offset;
+
+  /* Count sample for the shadow catcher object. */
+  kernel_write_pass_float(buffer + kernel_data.film.pass_shadow_catcher_sample_count, 1.0f);
+
+  /* Since the split is done, the sample does not contribute to the matte, so accumulate it as
+   * transparency to the matte. */
+  const float3 throughput = INTEGRATOR_STATE(state, path, throughput);
+  kernel_write_pass_float(buffer + kernel_data.film.pass_shadow_catcher_matte + 3,
+                          average(throughput));
+}
+
 #endif /* __SHADOW_CATCHER__ */
 
 CCL_NAMESPACE_END

intern/cycles/kernel/integrator/shadow_state_template.h

@@ -46,8 +46,9 @@ KERNEL_STRUCT_MEMBER(shadow_path,
                      float3,
                      unshadowed_throughput,
                      KERNEL_FEATURE_SHADOW_PASS | KERNEL_FEATURE_AO_ADDITIVE)
-/* Ratio of throughput to distinguish diffuse and glossy render passes. */
-KERNEL_STRUCT_MEMBER(shadow_path, float3, diffuse_glossy_ratio, KERNEL_FEATURE_LIGHT_PASSES)
+/* Ratio of throughput to distinguish diffuse / glossy / transmission render passes. */
+KERNEL_STRUCT_MEMBER(shadow_path, float3, pass_diffuse_weight, KERNEL_FEATURE_LIGHT_PASSES)
+KERNEL_STRUCT_MEMBER(shadow_path, float3, pass_glossy_weight, KERNEL_FEATURE_LIGHT_PASSES)
 /* Number of intersections found by ray-tracing. */
 KERNEL_STRUCT_MEMBER(shadow_path, uint16_t, num_hits, KERNEL_FEATURE_PATH_TRACING)
 KERNEL_STRUCT_END(shadow_path)

intern/cycles/kernel/integrator/state_template.h

@@ -60,8 +60,9 @@ KERNEL_STRUCT_MEMBER(path, float, min_ray_pdf, KERNEL_FEATURE_PATH_TRACING)
 KERNEL_STRUCT_MEMBER(path, float, continuation_probability, KERNEL_FEATURE_PATH_TRACING)
 /* Throughput. */
 KERNEL_STRUCT_MEMBER(path, float3, throughput, KERNEL_FEATURE_PATH_TRACING)
-/* Ratio of throughput to distinguish diffuse and glossy render passes. */
-KERNEL_STRUCT_MEMBER(path, float3, diffuse_glossy_ratio, KERNEL_FEATURE_LIGHT_PASSES)
+/* Ratio of throughput to distinguish diffuse / glossy / transmission render passes. */
+KERNEL_STRUCT_MEMBER(path, float3, pass_diffuse_weight, KERNEL_FEATURE_LIGHT_PASSES)
+KERNEL_STRUCT_MEMBER(path, float3, pass_glossy_weight, KERNEL_FEATURE_LIGHT_PASSES)
 /* Denoising. */
 KERNEL_STRUCT_MEMBER(path, float3, denoising_feature_throughput, KERNEL_FEATURE_DENOISING)
 /* Shader sorting. */

intern/cycles/kernel/integrator/subsurface.h

@@ -71,6 +71,10 @@ 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;
 
@@ -79,7 +83,8 @@ ccl_device int subsurface_bounce(KernelGlobals kg,
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
     if (INTEGRATOR_STATE(state, path, bounce) == 0) {
-      INTEGRATOR_STATE_WRITE(state, path, diffuse_glossy_ratio) = one_float3();
+      INTEGRATOR_STATE_WRITE(state, path, pass_diffuse_weight) = one_float3();
+      INTEGRATOR_STATE_WRITE(state, path, pass_glossy_weight) = zero_float3();
     }
   }
 

intern/cycles/kernel/integrator/subsurface_disk.h

@@ -47,6 +47,7 @@ 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);
@@ -123,6 +124,9 @@ 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 = true;
+    ls->pdf = 1.0f;
     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 (ls->eval_fac == 0.0f) {
+        if (!in_volume_segment && 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 (klight->area.tan_spread > 0.0f) {
+        if (!in_volume_segment && 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 (ls->pdf > 0.0f);
+  return in_volume_segment || (ls->pdf > 0.0f);
 }
 
 ccl_device bool lights_intersect(KernelGlobals kg,

intern/cycles/kernel/light/sample.h

@@ -199,6 +199,9 @@ 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,7 +23,8 @@ 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_CAMERA) && (!(path_flag & PATH_RAY_SINGLE_PASS_DONE));
+  bool is_primary = (path_flag & PATH_RAY_TRANSPARENT_BACKGROUND) &&
+                    (!(path_flag & PATH_RAY_SINGLE_PASS_DONE));
 
   return ((render_buffer != NULL) && is_primary);
 }

intern/cycles/kernel/types.h

@@ -279,17 +279,17 @@ 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_USE_FRESNEL | PATH_RAY_SUBSURFACE_BACKFACING),
 
   /* Contribute to denoising features. */
-  PATH_RAY_DENOISING_FEATURES = (1U << 23U),
+  PATH_RAY_DENOISING_FEATURES = (1U << 24U),
 
   /* Render pass categories. */
-  PATH_RAY_REFLECT_PASS = (1U << 24U),
-  PATH_RAY_TRANSMISSION_PASS = (1U << 25U),
+  PATH_RAY_SURFACE_PASS = (1U << 25U),
   PATH_RAY_VOLUME_PASS = (1U << 26U),
-  PATH_RAY_ANY_PASS = (PATH_RAY_REFLECT_PASS | PATH_RAY_TRANSMISSION_PASS | PATH_RAY_VOLUME_PASS),
+  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),
@@ -428,8 +428,20 @@ typedef enum CryptomatteType {
 typedef struct BsdfEval {
   float3 diffuse;
   float3 glossy;
+  float3 sum;
 } BsdfEval;
 
+/* Closure Filter */
+
+typedef enum FilterClosures {
+  FILTER_CLOSURE_EMISSION = (1 << 0),
+  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),
+} FilterClosures;
+
 /* Shader Flag */
 
 typedef enum ShaderFlag {
@@ -1186,7 +1198,11 @@ typedef struct KernelIntegrator {
   int has_shadow_catcher;
   float scrambling_distance;
 
+  /* Closure filter. */
+  int filter_closures;
+
   /* padding */
+  int pad1, pad2, pad3;
 } KernelIntegrator;
 static_assert_align(KernelIntegrator, 16);
 

intern/cycles/scene/film.cpp

@@ -187,8 +187,6 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
   kfilm->pass_transmission_indirect = PASS_UNUSED;
   kfilm->pass_volume_direct = PASS_UNUSED;
   kfilm->pass_volume_indirect = PASS_UNUSED;
-  kfilm->pass_volume_direct = PASS_UNUSED;
-  kfilm->pass_volume_indirect = PASS_UNUSED;
   kfilm->pass_shadow = PASS_UNUSED;
 
   /* Mark passes as unused so that the kernel knows the pass is inaccessible. */
@@ -673,13 +671,12 @@ uint Film::get_kernel_features(const Scene *scene) const
       kernel_features |= KERNEL_FEATURE_DENOISING;
     }
 
-    if (pass_type != PASS_NONE && pass_type != PASS_COMBINED &&
-        pass_type <= PASS_CATEGORY_LIGHT_END) {
+    if (pass_type >= PASS_DIFFUSE && pass_type <= PASS_VOLUME_INDIRECT) {
       kernel_features |= KERNEL_FEATURE_LIGHT_PASSES;
+    }
 
-      if (pass_type == PASS_SHADOW) {
-        kernel_features |= KERNEL_FEATURE_SHADOW_PASS;
-      }
+    if (pass_type == PASS_SHADOW) {
+      kernel_features |= KERNEL_FEATURE_SHADOW_PASS;
     }
 
     if (pass_type == PASS_AO) {

intern/cycles/scene/integrator.cpp

@@ -14,11 +14,13 @@
  * 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"
@@ -63,6 +65,14 @@ NODE_DEFINE(Integrator)
   SOCKET_BOOLEAN(caustics_reflective, "Reflective Caustics", true);
   SOCKET_BOOLEAN(caustics_refractive, "Refractive Caustics", true);
   SOCKET_FLOAT(filter_glossy, "Filter Glossy", 0.0f);
+
+  SOCKET_BOOLEAN(use_direct_light, "Use Direct Light", true);
+  SOCKET_BOOLEAN(use_indirect_light, "Use Indirect Light", true);
+  SOCKET_BOOLEAN(use_diffuse, "Use Diffuse", true);
+  SOCKET_BOOLEAN(use_glossy, "Use Glossy", true);
+  SOCKET_BOOLEAN(use_transmission, "Use Transmission", true);
+  SOCKET_BOOLEAN(use_emission, "Use Emission", true);
+
   SOCKET_INT(seed, "Seed", 0);
   SOCKET_FLOAT(sample_clamp_direct, "Sample Clamp Direct", 0.0f);
   SOCKET_FLOAT(sample_clamp_indirect, "Sample Clamp Indirect", 0.0f);
@@ -158,7 +168,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_bounces;
+  kintegrator->ao_bounces = (ao_factor != 0.0f) ? ao_bounces : 0;
   kintegrator->ao_bounces_distance = ao_distance;
   kintegrator->ao_bounces_factor = ao_factor;
   kintegrator->ao_additive_factor = ao_additive_factor;
@@ -184,6 +194,32 @@ void Integrator::device_update(Device *device, DeviceScene *dscene, Scene *scene
   kintegrator->caustics_refractive = caustics_refractive;
   kintegrator->filter_glossy = (filter_glossy == 0.0f) ? FLT_MAX : 1.0f / filter_glossy;
 
+  kintegrator->filter_closures = 0;
+  if (!use_direct_light) {
+    kintegrator->filter_closures |= FILTER_CLOSURE_DIRECT_LIGHT;
+  }
+  if (!use_indirect_light) {
+    kintegrator->min_bounce = 1;
+    kintegrator->max_bounce = 1;
+  }
+  if (!use_diffuse) {
+    kintegrator->filter_closures |= FILTER_CLOSURE_DIFFUSE;
+  }
+  if (!use_glossy) {
+    kintegrator->filter_closures |= FILTER_CLOSURE_GLOSSY;
+  }
+  if (!use_transmission) {
+    kintegrator->filter_closures |= FILTER_CLOSURE_TRANSMISSION;
+  }
+  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;
 
   kintegrator->sample_clamp_direct = (sample_clamp_direct == 0.0f) ? FLT_MAX :

intern/cycles/scene/integrator.h

@@ -56,6 +56,13 @@ class Integrator : public Node {
   NODE_SOCKET_API(bool, caustics_refractive)
   NODE_SOCKET_API(float, filter_glossy)
 
+  NODE_SOCKET_API(bool, use_direct_light);
+  NODE_SOCKET_API(bool, use_indirect_light);
+  NODE_SOCKET_API(bool, use_diffuse);
+  NODE_SOCKET_API(bool, use_glossy);
+  NODE_SOCKET_API(bool, use_transmission);
+  NODE_SOCKET_API(bool, use_emission);
+
   NODE_SOCKET_API(int, seed)
 
   NODE_SOCKET_API(float, sample_clamp_direct)

intern/cycles/scene/osl.cpp

@@ -274,19 +274,26 @@ void OSLShaderManager::shading_system_init()
 
         "diffuse_ancestor", /* PATH_RAY_DIFFUSE_ANCESTOR */
 
-        "__unused__", /* PATH_RAY_SINGLE_PASS_DONE */
-        "__unused__", /* PATH_RAY_TRANSPARENT_BACKGROUND */
-        "__unused__", /* PATH_RAY_TERMINATE_IMMEDIATE */
-        "__unused__", /* PATH_RAY_TERMINATE_AFTER_TRANSPARENT */
-        "__unused__", /* PATH_RAY_EMISSION */
-        "__unused__", /* PATH_RAY_SUBSURFACE */
-        "__unused__", /* PATH_RAY_DENOISING_FEATURES */
-        "__unused__", /* PATH_RAY_REFLECT_PASS */
-        "__unused__", /* PATH_RAY_TRANSMISSION_PASS */
-        "__unused__", /* PATH_RAY_VOLUME_PASS */
-        "__unused__", /* PATH_RAY_SHADOW_FOR_LIGHT */
-        "__unused__", /* PATH_RAY_SHADOW_CATCHER_HIT */
-        "__unused__", /* PATH_RAY_SHADOW_CATCHER_PASS */
+        /* Remaining irrelevant bits up to 32. */
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
+        "__unused__",
     };
 
     const int nraytypes = sizeof(raytypes) / sizeof(raytypes[0]);

intern/cycles/scene/shader_nodes.cpp

@@ -1503,18 +1503,21 @@ 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,
-                                           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.encode_uchar4(vector_offset, scale_ofs, distortion_ofs),
+                    compiler.encode_uchar4(detail_ofs, dscale_ofs, droughness_ofs, phase_ofs));
 
-  compiler.add_node(compiler.encode_uchar4(compiler.stack_assign_if_linked(color_out),
-                                           compiler.stack_assign_if_linked(fac_out)),
+  compiler.add_node(compiler.encode_uchar4(color_ofs, fac_ofs),
                     __float_as_int(scale),
                     __float_as_int(distortion),
                     __float_as_int(detail));
@@ -3572,45 +3575,54 @@ 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. */
-  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),
+  /* 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),
                     __float_as_uint(offset),
                     __float_as_uint(ior));
-
-  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),
+  /* data node 2 */
+  compiler.add_node(compiler.encode_uchar4(
+                        coat_ofs, melanin_ofs, melanin_redness_ofs, absorption_coefficient_ofs),
                     __float_as_uint(coat),
                     __float_as_uint(melanin),
                     __float_as_uint(melanin_redness));
 
-  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 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));
 
+  /* 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

@@ -367,14 +367,26 @@ 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(buffer_params_.width, buffer_params_.height);
+    return make_int2(image_width, image_height);
   }
 
-  /* TODO(sergey): Take available memory into account, and if there is enough memory do not tile
-   * and prefer optimal performance. */
+  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. */
+
   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);
 }
 
@@ -504,8 +516,8 @@ void Session::set_display_driver(unique_ptr<DisplayDriver> driver)
 
 double Session::get_estimated_remaining_time() const
 {
-  const float completed = progress.get_progress();
-  if (completed == 0.0f) {
+  const double completed = progress.get_progress();
+  if (completed == 0.0) {
     return 0.0;
   }
 
@@ -573,7 +585,7 @@ void Session::update_status_time(bool show_pause, bool show_done)
   }
 
   /* Sample. */
-  if (num_samples == Integrator::MAX_SAMPLES) {
+  if (!params.background && num_samples == Integrator::MAX_SAMPLES) {
     substatus = status_append(substatus, string_printf("Sample %d", current_sample));
   }
   else {

intern/cycles/session/tile.cpp

@@ -29,6 +29,7 @@
 #include "util/path.h"
 #include "util/string.h"
 #include "util/system.h"
+#include "util/time.h"
 #include "util/types.h"
 
 CCL_NAMESPACE_BEGIN
@@ -503,9 +504,9 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
     }
   }
 
-  DCHECK_EQ(tile_buffers.params.pass_stride, buffer_params_.pass_stride);
+  const double time_start = time_dt();
 
-  vector<float> pixel_storage;
+  DCHECK_EQ(tile_buffers.params.pass_stride, buffer_params_.pass_stride);
 
   const BufferParams &tile_params = tile_buffers.params;
 
@@ -515,13 +516,32 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
   const int64_t pass_stride = tile_params.pass_stride;
   const int64_t tile_row_stride = tile_params.width * pass_stride;
 
-  const int64_t xstride = pass_stride * sizeof(float);
-  const int64_t ystride = xstride * tile_params.width;
-  const int64_t zstride = ystride * tile_params.height;
-
+  vector<float> pixel_storage;
   const float *pixels = tile_buffers.buffer.data() + tile_params.window_x * pass_stride +
                         tile_params.window_y * tile_row_stride;
 
+  /* If there is an overscan used for the tile copy pixels into single continuous block of memory
+   * without any "gaps".
+   * This is a workaround for bug in OIIO (https://github.com/OpenImageIO/oiio/pull/3176).
+   * Our task reference: T93008. */
+  if (tile_params.window_x || tile_params.window_y ||
+      tile_params.window_width != tile_params.width ||
+      tile_params.window_height != tile_params.height) {
+    pixel_storage.resize(pass_stride * tile_params.window_width * tile_params.window_height);
+    float *pixels_continuous = pixel_storage.data();
+
+    const int64_t pixels_row_stride = pass_stride * tile_params.width;
+    const int64_t pixels_continuous_row_stride = pass_stride * tile_params.window_width;
+
+    for (int i = 0; i < tile_params.window_height; ++i) {
+      memcpy(pixels_continuous, pixels, sizeof(float) * pixels_continuous_row_stride);
+      pixels += pixels_row_stride;
+      pixels_continuous += pixels_continuous_row_stride;
+    }
+
+    pixels = pixel_storage.data();
+  }
+
   VLOG(3) << "Write tile at " << tile_x << ", " << tile_y;
 
   /* The image tile sizes in the OpenEXR file are different from the size of our big tiles. The
@@ -531,6 +551,11 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
    *
    * The only thing we have to ensure is that the tile_x and tile_y are a multiple of the
    * image tile size, which happens in compute_render_tile_size. */
+
+  const int64_t xstride = pass_stride * sizeof(float);
+  const int64_t ystride = xstride * tile_params.window_width;
+  const int64_t zstride = ystride * tile_params.window_height;
+
   if (!write_state_.tile_out->write_tiles(tile_x,
                                           tile_x + tile_params.window_width,
                                           tile_y,
@@ -548,6 +573,8 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
 
   ++write_state_.num_tiles_written;
 
+  VLOG(3) << "Tile written in " << time_dt() - time_start << " seconds.";
+
   return true;
 }
 
@@ -589,6 +616,9 @@ void TileManager::finish_write_tiles()
     full_buffer_written_cb(write_state_.filename);
   }
 
+  VLOG(3) << "Tile file size is "
+          << string_human_readable_number(path_file_size(write_state_.filename)) << " bytes.";
+
   /* Advance the counter upon explicit finish of the file.
    * Makes it possible to re-use tile manager for another scene, and avoids unnecessary increments
    * of the tile-file-within-session index. */

intern/cycles/test/CMakeLists.txt

@@ -54,17 +54,21 @@ set(SRC
   util_transform_test.cpp
 )
 
-if(CXX_HAS_AVX)
-  list(APPEND SRC
-    util_avxf_avx_test.cpp
-  )
-  set_source_files_properties(util_avxf_avx_test.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_AVX_KERNEL_FLAGS}")
-endif()
-if(CXX_HAS_AVX2)
-  list(APPEND SRC
-    util_avxf_avx2_test.cpp
-  )
-  set_source_files_properties(util_avxf_avx2_test.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_AVX2_KERNEL_FLAGS}")
+# Disable AVX tests on macOS. Rosetta has problems running them, and other
+# platforms should be enough to verify AVX operations are implemented correctly.
+if(NOT APPLE)
+  if(CXX_HAS_AVX)
+    list(APPEND SRC
+      util_avxf_avx_test.cpp
+    )
+    set_source_files_properties(util_avxf_avx_test.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_AVX_KERNEL_FLAGS}")
+  endif()
+  if(CXX_HAS_AVX2)
+    list(APPEND SRC
+      util_avxf_avx2_test.cpp
+    )
+    set_source_files_properties(util_avxf_avx2_test.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_AVX2_KERNEL_FLAGS}")
+  endif()
 endif()
 
 if(WITH_GTESTS)

intern/cycles/test/integrator_tile_test.cpp

@@ -24,26 +24,26 @@ CCL_NAMESPACE_BEGIN
 TEST(tile_calculate_best_size, Basic)
 {
   /* Make sure CPU-like case is handled properly. */
-  EXPECT_EQ(tile_calculate_best_size(make_int2(1920, 1080), 1, 1, 1.0f), TileSize(1, 1, 1));
-  EXPECT_EQ(tile_calculate_best_size(make_int2(1920, 1080), 100, 1, 1.0f), TileSize(1, 1, 1));
+  EXPECT_EQ(tile_calculate_best_size(false, make_int2(1920, 1080), 1, 1, 1.0f), TileSize(1, 1, 1));
+  EXPECT_EQ(tile_calculate_best_size(false, make_int2(1920, 1080), 100, 1, 1.0f), TileSize(1, 1, 1));
 
   /* Enough path states to fit an entire image with all samples. */
-  EXPECT_EQ(tile_calculate_best_size(make_int2(1920, 1080), 1, 1920 * 1080, 1.0f),
+  EXPECT_EQ(tile_calculate_best_size(false, make_int2(1920, 1080), 1, 1920 * 1080, 1.0f),
             TileSize(1920, 1080, 1));
-  EXPECT_EQ(tile_calculate_best_size(make_int2(1920, 1080), 100, 1920 * 1080 * 100, 1.0f),
+  EXPECT_EQ(tile_calculate_best_size(false, make_int2(1920, 1080), 100, 1920 * 1080 * 100, 1.0f),
             TileSize(1920, 1080, 100));
 }
 
 TEST(tile_calculate_best_size, Extreme)
 {
-  EXPECT_EQ(tile_calculate_best_size(make_int2(32, 32), 262144, 131072, 1.0f),
+  EXPECT_EQ(tile_calculate_best_size(false, make_int2(32, 32), 262144, 131072, 1.0f),
             TileSize(1, 1, 512));
-  EXPECT_EQ(tile_calculate_best_size(make_int2(32, 32), 1048576, 131072, 1.0f),
+  EXPECT_EQ(tile_calculate_best_size(false, make_int2(32, 32), 1048576, 131072, 1.0f),
             TileSize(1, 1, 1024));
-  EXPECT_EQ(tile_calculate_best_size(make_int2(32, 32), 10485760, 131072, 1.0f),
+  EXPECT_EQ(tile_calculate_best_size(false, make_int2(32, 32), 10485760, 131072, 1.0f),
             TileSize(1, 1, 4096));
 
-  EXPECT_EQ(tile_calculate_best_size(make_int2(32, 32), 8192 * 8192 * 2, 1024, 1.0f),
+  EXPECT_EQ(tile_calculate_best_size(false, make_int2(32, 32), 8192 * 8192 * 2, 1024, 1.0f),
             TileSize(1, 1, 1024));
 }
 

intern/cycles/test/util_avxf_test.h

@@ -32,9 +32,13 @@ static bool validate_cpu_capabilities()
 #endif
 }
 
-#define VALIDATECPU \
+#define INIT_AVX_TEST \
   if (!validate_cpu_capabilities()) \
-    return;
+    return; \
+\
+  const avxf avxf_a(0.1f, 0.2f, 0.3f, 0.4f, 0.5f, 0.6f, 0.7f, 0.8f); \
+  const avxf avxf_b(1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f); \
+  const avxf avxf_c(1.1f, 2.2f, 3.3f, 4.4f, 5.5f, 6.6f, 7.7f, 8.8f);
 
 #define compare_vector_scalar(a, b) \
   for (size_t index = 0; index < a.size; index++) \
@@ -49,21 +53,18 @@ static bool validate_cpu_capabilities()
     EXPECT_NEAR(a[index], b[index], abserror);
 
 #define basic_test_vv(a, b, op) \
-  VALIDATECPU \
+  INIT_AVX_TEST \
   avxf c = a op b; \
   for (size_t i = 0; i < a.size; i++) \
     EXPECT_FLOAT_EQ(c[i], a[i] op b[i]);
 
 /* vector op float tests */
 #define basic_test_vf(a, b, op) \
-  VALIDATECPU \
+  INIT_AVX_TEST \
   avxf c = a op b; \
   for (size_t i = 0; i < a.size; i++) \
     EXPECT_FLOAT_EQ(c[i], a[i] op b);
 
-static const avxf avxf_a(0.1f, 0.2f, 0.3f, 0.4f, 0.5f, 0.6f, 0.7f, 0.8f);
-static const avxf avxf_b(1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f);
-static const avxf avxf_c(1.1f, 2.2f, 3.3f, 4.4f, 5.5f, 6.6f, 7.7f, 8.8f);
 static const float float_b = 1.5f;
 
 TEST(TEST_CATEGORY_NAME, avxf_add_vv){basic_test_vv(avxf_a, avxf_b, +)} TEST(TEST_CATEGORY_NAME,
@@ -78,7 +79,7 @@ TEST(TEST_CATEGORY_NAME, avxf_add_vv){basic_test_vv(avxf_a, avxf_b, +)} TEST(TES
 
 TEST(TEST_CATEGORY_NAME, avxf_ctor)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   compare_vector_scalar(avxf(7.0f, 6.0f, 5.0f, 4.0f, 3.0f, 2.0f, 1.0f, 0.0f),
                         static_cast<float>(index));
   compare_vector_scalar(avxf(1.0f), 1.0f);
@@ -91,28 +92,28 @@ TEST(TEST_CATEGORY_NAME, avxf_ctor)
 
 TEST(TEST_CATEGORY_NAME, avxf_sqrt)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   compare_vector_vector(mm256_sqrt(avxf(1.0f, 4.0f, 9.0f, 16.0f, 25.0f, 36.0f, 49.0f, 64.0f)),
                         avxf(1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f));
 }
 
 TEST(TEST_CATEGORY_NAME, avxf_min_max)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   compare_vector_vector(min(avxf_a, avxf_b), avxf_a);
   compare_vector_vector(max(avxf_a, avxf_b), avxf_b);
 }
 
 TEST(TEST_CATEGORY_NAME, avxf_set_sign)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   avxf res = set_sign_bit<1, 0, 0, 0, 0, 0, 0, 0>(avxf_a);
   compare_vector_vector(res, avxf(0.1f, 0.2f, 0.3f, 0.4f, 0.5f, 0.6f, 0.7f, -0.8f));
 }
 
 TEST(TEST_CATEGORY_NAME, avxf_msub)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   avxf res = msub(avxf_a, avxf_b, avxf_c);
   avxf exp = avxf((avxf_a[7] * avxf_b[7]) - avxf_c[7],
                   (avxf_a[6] * avxf_b[6]) - avxf_c[6],
@@ -127,7 +128,7 @@ TEST(TEST_CATEGORY_NAME, avxf_msub)
 
 TEST(TEST_CATEGORY_NAME, avxf_madd)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   avxf res = madd(avxf_a, avxf_b, avxf_c);
   avxf exp = avxf((avxf_a[7] * avxf_b[7]) + avxf_c[7],
                   (avxf_a[6] * avxf_b[6]) + avxf_c[6],
@@ -142,7 +143,7 @@ TEST(TEST_CATEGORY_NAME, avxf_madd)
 
 TEST(TEST_CATEGORY_NAME, avxf_nmadd)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   avxf res = nmadd(avxf_a, avxf_b, avxf_c);
   avxf exp = avxf(avxf_c[7] - (avxf_a[7] * avxf_b[7]),
                   avxf_c[6] - (avxf_a[6] * avxf_b[6]),
@@ -157,7 +158,7 @@ TEST(TEST_CATEGORY_NAME, avxf_nmadd)
 
 TEST(TEST_CATEGORY_NAME, avxf_compare)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   avxf a(0.0f, 1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f);
   avxf b(7.0f, 6.0f, 5.0f, 4.0f, 3.0f, 2.0f, 1.0f, 0.0f);
   avxb res = a <= b;
@@ -176,28 +177,28 @@ TEST(TEST_CATEGORY_NAME, avxf_compare)
 
 TEST(TEST_CATEGORY_NAME, avxf_permute)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   avxf res = permute<3, 0, 1, 7, 6, 5, 2, 4>(avxf_b);
   compare_vector_vector(res, avxf(4.0f, 6.0f, 3.0f, 2.0f, 1.0f, 7.0f, 8.0f, 5.0f));
 }
 
 TEST(TEST_CATEGORY_NAME, avxf_blend)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   avxf res = blend<0, 0, 1, 0, 1, 0, 1, 0>(avxf_a, avxf_b);
   compare_vector_vector(res, avxf(0.1f, 0.2f, 3.0f, 0.4f, 5.0f, 0.6f, 7.0f, 0.8f));
 }
 
 TEST(TEST_CATEGORY_NAME, avxf_shuffle)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   avxf res = shuffle<0, 1, 2, 3, 1, 3, 2, 0>(avxf_a);
   compare_vector_vector(res, avxf(0.4f, 0.2f, 0.1f, 0.3f, 0.5f, 0.6f, 0.7f, 0.8f));
 }
 
 TEST(TEST_CATEGORY_NAME, avxf_cross)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   avxf res = cross(avxf_b, avxf_c);
   compare_vector_vector_near(res,
                              avxf(0.0f,
@@ -213,7 +214,7 @@ TEST(TEST_CATEGORY_NAME, avxf_cross)
 
 TEST(TEST_CATEGORY_NAME, avxf_dot3)
 {
-  VALIDATECPU
+  INIT_AVX_TEST
   float den, den2;
   dot3(avxf_a, avxf_b, den, den2);
   EXPECT_FLOAT_EQ(den, 14.9f);

intern/cycles/util/profiling.cpp

@@ -171,4 +171,9 @@ bool Profiler::get_object(int object, uint64_t &samples, uint64_t &hits)
   return true;
 }
 
+bool Profiler::active() const
+{
+  return (worker != nullptr);
+}
+
 CCL_NAMESPACE_END

intern/cycles/util/profiling.h

@@ -96,6 +96,8 @@ class Profiler {
   bool get_shader(int shader, uint64_t &samples, uint64_t &hits);
   bool get_object(int object, uint64_t &samples, uint64_t &hits);
 
+  bool active() const;
+
  protected:
   void run();
 

intern/cycles/util/progress.h

@@ -200,12 +200,12 @@ class Progress {
     total_pixel_samples = total_pixel_samples_;
   }
 
-  float get_progress() const
+  double get_progress() const
   {
     thread_scoped_lock lock(progress_mutex);
 
     if (total_pixel_samples > 0) {
-      return ((float)pixel_samples) / total_pixel_samples;
+      return ((double)pixel_samples) / (double)total_pixel_samples;
     }
     return 0.0f;
   }

intern/ghost/intern/GHOST_SystemWin32.cpp

@@ -1002,10 +1002,10 @@ void GHOST_SystemWin32::processWintabEvent(GHOST_WindowWin32 *window)
     DWORD pos = GetMessagePos();
     int x = GET_X_LPARAM(pos);
     int y = GET_Y_LPARAM(pos);
+    GHOST_TabletData td = wt->getLastTabletData();
 
-    /* TODO supply tablet data */
     system->pushEvent(new GHOST_EventCursor(
-        system->getMilliSeconds(), GHOST_kEventCursorMove, window, x, y, GHOST_TABLET_DATA_NONE));
+        system->getMilliSeconds(), GHOST_kEventCursorMove, window, x, y, td));
   }
 }
 
@@ -1472,6 +1472,7 @@ LRESULT WINAPI GHOST_SystemWin32::s_wndProc(HWND hwnd, UINT msg, WPARAM wParam,
         case WM_IME_SETCONTEXT: {
           GHOST_ImeWin32 *ime = window->getImeInput();
           ime->UpdateInputLanguage();
+          ime->UpdateConversionStatus(hwnd);
           ime->CreateImeWindow(hwnd);
           ime->CleanupComposition(hwnd);
           ime->CheckFirst(hwnd);

intern/ghost/intern/GHOST_WindowViewCocoa.h

@@ -510,6 +510,14 @@
 - (void)checkKeyCodeIsControlChar:(NSEvent *)event
 {
   ime.state_flag &= ~GHOST_IME_KEY_CONTROL_CHAR;
+
+  /* Don't use IME for command and ctrl key combinations, these are shortcuts. */
+  if ([event modifierFlags] & (NSEventModifierFlagCommand | NSEventModifierFlagControl)) {
+    ime.state_flag |= GHOST_IME_KEY_CONTROL_CHAR;
+    return;
+  }
+
+  /* Don't use IME for these control keys. */
   switch ([event keyCode]) {
     case kVK_ANSI_KeypadEnter:
     case kVK_ANSI_KeypadClear:

release/datafiles/blender_icons.svg

@@ -8520,26 +8520,6 @@
          inkscape:connector-curvature="0"
          sodipodi:nodetypes="sssssssssssssss" />
     </g>
-    <g
-       transform="matrix(0.92857149,0,0,0.92857137,106.93015,-501.7093)"
-       style="display:inline;opacity:0.98999999;fill:#ffffff;stroke-width:1.07692313;enable-background:new"
-       id="g8599"
-       inkscape:export-filename="C:\Users\Andrzej Ambroż\Desktop\mtrx.png"
-       inkscape:export-xdpi="96"
-       inkscape:export-ydpi="96">
-      <path
-         sodipodi:nodetypes="cccccccccccssccsssss"
-         inkscape:connector-curvature="0"
-         id="path8595"
-         transform="matrix(1.076923,0,0,1.0769231,-739.76878,506.92345)"
-         d="m 992.67578,105.26367 c -0.19597,-0.18575 -0.45872,-0.28437 -0.72851,-0.27344 -0.89539,0.0396 -1.29072,1.14623 -0.62305,1.74415 L 992.69727,108 H 986 c -1.36099,-0.0279 -1.36099,2.02794 0,2 h 3.0918 c -0.003,0.002 -0.005,0.004 -0.008,0.006 l -4.75,4.24805 c -0.99479,0.88933 0.3392,2.38151 1.33399,1.49218 l 2.33984,-2.09375 C 988.08993,116.60772 990.52575,119 993.5,119 c 3.02565,0 5.49805,-2.47428 5.49805,-5.5 0,-1.56564 -0.66395,-2.97983 -1.72241,-3.98356 z M 993.5,110 c 1.94471,0 3.5,1.5551 3.5,3.5 0,1.94489 -1.55529,3.5 -3.5,3.5 -1.94472,0 -3.5,-1.55511 -3.5,-3.5 0,-1.9449 1.55528,-3.5 3.5,-3.5 z"
-         style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:2;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:accumulate" />
-      <path
-         id="path8597"
-         d="m 331.99999,629.15424 a 1.86519,1.8457757 0 0 1 -1.86519,1.84577 1.86519,1.8457757 0 0 1 -1.86519,-1.84577 1.86519,1.8457757 0 0 1 1.86519,-1.84578 1.86519,1.8457757 0 0 1 1.86519,1.84578 z"
-         style="opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;stroke:none;stroke-width:2.15384626;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke"
-         inkscape:connector-curvature="0" />
-    </g>
     <g
        id="g12575"
        transform="matrix(0.8666665,0,0,0.8666665,-253.07368,16.407198)"
@@ -13729,11 +13709,6 @@
          id="g16343-6"
          style="display:inline;opacity:0.98999999;fill:#ffffff;enable-background:new" />
     </g>
-    <path
-       inkscape:connector-curvature="0"
-       id="path18041"
-       d="m 517.0019,53.000264 c -0.47496,-0.0438 -0.94252,0.0803 -1.30468,0.35742 -0.28815,0.22047 -0.55389,0.56737 -0.6211,1.0332 -0.0672,0.46584 0.11485,1.01133 0.57227,1.46876 l 0.14648,0.14648 h -4.29297 c -0.53506,4e-5 -1.03128,0.28661 -1.29883,0.75 -0.26752,0.46339 -0.26752,1.03661 0,1.5 0.26755,0.46339 0.76375,0.75 1.29883,0.75 h 0.79297 l -2.64648,2.64648 c -0.79091,0.79079 -0.84722,1.93194 -0.29297,2.65625 0.22044,0.28814 0.56748,0.55402 1.0332,0.6211 0.46572,0.0671 1.0091,-0.11342 1.4668,-0.57031 l 1.4043,-1.40235 c 0.55515,1.96585 2.10614,3.53021 4.16406,3.94141 2.37442,0.47444 4.78539,-0.66388 5.92773,-2.79883 1.14235,-2.13495 0.75134,-4.77035 -0.96094,-6.48242 a 0.50005,0.50005 0 0 0 -0.002,-0.004 l -4.0332,-3.96094 c -0.39539,-0.39541 -0.87856,-0.60853 -1.35352,-0.65234 z m -0.0957,1.00195 c 0.24304,0.0198 0.50958,0.1248 0.74219,0.35743 a 0.50005,0.50005 0 0 0 0.004,0.002 l 4.03125,3.96289 c 1.40369,1.40351 1.72164,3.55449 0.78516,5.30469 -0.93648,1.7502 -2.90114,2.678 -4.84766,2.28906 -1.9465,-0.38894 -3.40679,-2.00407 -3.60937,-3.97266 a 0.50005,0.50005 0 0 0 -0.85156,-0.30273 l -2.01172,2.00976 c -0.29296,0.29245 -0.47153,0.30809 -0.61719,0.28711 -0.14566,-0.021 -0.29945,-0.12932 -0.38281,-0.23828 -0.21096,-0.27568 -0.25826,-0.87658 0.20703,-1.34179 l 3.5,-3.5 a 0.50005,0.50005 0 0 0 -0.35352,-0.85352 h -2 c -0.17943,0 -0.34387,-0.0946 -0.43359,-0.25 -0.0897,-0.15539 -0.0897,-0.34461 0,-0.5 0.0897,-0.15539 0.25413,-0.24999 0.43359,-0.25 h 5.5 a 0.50005,0.50005 0 0 0 0.35352,-0.85352 l -1,-1 c -0.29258,-0.29257 -0.31008,-0.47147 -0.28906,-0.61718 0.021,-0.14571 0.13127,-0.29945 0.24023,-0.38282 0.13784,-0.10546 0.35657,-0.17021 0.59961,-0.15039 z m 1.5957,5.00391 c -1.37479,0 -2.5,1.12521 -2.5,2.5 0,1.37479 1.12521,2.5 2.5,2.5 1.37479,0 2.5,-1.12521 2.5,-2.5 0,-1.37479 -1.12521,-2.5 -2.5,-2.5 z m 0,1 c 0.83435,0 1.5,0.66565 1.5,1.5 0,0.83435 -0.66565,1.5 -1.5,1.5 -0.83435,0 -1.5,-0.66565 -1.5,-1.5 0,-0.83435 0.66565,-1.5 1.5,-1.5 z"
-       style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:0.98999999;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:accumulate" />
     <path
        id="path14479-6"
        style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:0.99999994;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
@@ -17373,6 +17348,136 @@
        d="m 417.92349,304.73964 c -0.7818,-0.0644 -0.86293,1.09626 -0.0796,1.1383 l 0.41758,0.0202 c 0.78182,0.0644 0.86296,-1.09626 0.0796,-1.13831 z m -7.87437,1.29265 c -0.65325,0.42724 0.0163,1.38626 0.65667,0.94062 l 0.34001,-0.23929 c 0.65327,-0.42727 -0.0163,-1.38631 -0.65662,-0.94061 z m 5.26412,-0.10772 c 0.785,-0.0185 0.73895,-1.18175 -0.0451,-1.14009 -0.6811,-0.0652 -1.43225,-0.0213 -2.22341,0.0851 -0.785,0.0185 -0.73896,1.18176 0.0451,1.14011 0.8585,-0.10954 1.60282,-0.14009 2.22342,-0.0852 z m -5.74172,5.34858 c -0.17789,-0.75187 -1.32618,-0.47161 -1.12597,0.27482 -0.008,0.72815 0.18352,1.43475 0.53595,2.12392 0.17789,0.75187 1.32617,0.47159 1.12598,-0.27483 -0.40688,-0.70818 -0.47775,-1.41605 -0.53596,-2.12391 z m 1.14987,4.81425 c 0.55238,0.5479 1.3799,-0.2833 0.81165,-0.81524 l -0.30437,-0.28193 c -0.55238,-0.54789 -1.37991,0.2833 -0.81163,0.81524 z m 2.55883,0.11471 c -0.78112,0.0716 -0.65484,1.22767 0.12391,1.13446 0.79706,0.0708 1.5429,0.0136 2.2124,-0.23372 0.7811,-0.0716 0.65482,-1.22768 -0.12391,-1.13445 -0.66955,0.35373 -1.42049,0.37687 -2.2124,0.23371 z m 4.35036,-1.24066 c 0.39775,-0.66505 -0.63058,-1.23994 -1.00859,-0.56384 l -0.19953,0.36135 c -0.39776,0.66506 0.63057,1.23995 1.00857,0.56383 z m -1.53457,-4.82813 c -0.44444,-0.63566 -1.409,0.0364 -0.94666,0.65956 0.53116,0.53126 0.99257,1.10609 1.28624,1.78569 0.44445,0.63565 1.40902,-0.0364 0.94667,-0.65956 -0.24301,-0.74231 -0.69323,-1.32054 -1.28625,-1.78569 z m -2.73483,-1.49223 c -0.72218,-0.30138 -1.16808,0.7761 -0.43732,1.05681 l 0.39025,0.14758 c 0.7222,0.30141 1.1681,-0.7761 0.43732,-1.0568 z m -7.60223,1.91562 c -0.52109,0.57678 0.37464,1.33651 0.87855,0.74515 l 0.26685,-0.31654 c 0.52111,-0.57679 -0.37465,-1.33654 -0.87854,-0.74516 z m 1.15912,7.09355 c -0.1906,-0.74845 -1.33363,-0.44917 -1.12109,0.29354 l 0.11543,0.39523 c 0.19062,0.74845 1.33365,0.44917 1.12109,-0.29354 z m -0.68592,-4.36328 c -0.0858,-0.76698 -1.25912,-0.62352 -1.15127,0.14077 -0.065,0.75431 -0.008,1.50847 0.28594,2.26232 0.0859,0.76696 1.25912,0.62352 1.15129,-0.14076 -0.28468,-0.81162 -0.29126,-1.53878 -0.28596,-2.26233 z m 1.97398,-4.7241 c -0.77314,0.13162 -0.55483,1.27463 0.21417,1.12135 0.7762,-0.30633 1.5005,-0.42412 2.18687,-0.40397 0.77313,-0.13163 0.55482,-1.27462 -0.21418,-1.12137 -0.74152,0.0229 -1.4733,0.13255 -2.18686,0.40399 z"
        style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1.15052;stroke-linecap:round;stroke-linejoin:miter;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:2.2;stroke-opacity:1;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
        id="path4101-2-6-9-1_GP_dotdash" />
+    <g
+       id="g7580"
+       style="display:inline;enable-background:new">
+      <g
+         id="g905">
+        <g
+           transform="matrix(0.6740384,0,0,0.6740384,192.80592,-339.68227)"
+         style="display:inline;opacity:0.99;fill:#ffffff;stroke-width:1.07692;enable-background:new"
+           id="g8599-6-7"
+         inkscape:export-filename="C:\Users\Andrzej Ambroż\Desktop\mtrx.png"
+         inkscape:export-xdpi="96"
+         inkscape:export-ydpi="96">
+        <path
+             id="path8597-7-0"
+             d="m 331.87142,629.24052 c 0,0.81936 -0.66423,1.48359 -1.48359,1.48359 -0.81937,0 -1.4836,-0.66422 -1.4836,-1.48359 0,-0.81937 0.66423,-1.4836 1.4836,-1.4836 0.81937,0 1.48359,0.66423 1.48359,1.4836 z"
+             style="opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;stroke:none;stroke-width:1.72218;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke"
+             inkscape:connector-curvature="0"
+             sodipodi:nodetypes="sssss" />
+      </g>
+      <path
+           sodipodi:nodetypes="cccccccccczzcczzzzz"
+         inkscape:connector-curvature="0"
+           id="path8595-5-9"
+           d="m 414.93725,78.999996 c -0.65344,-0.653443 -1.58833,0.255453 -0.88297,0.960812 L 415.18167,81 h -4.6057 c -0.88913,-0.01822 -0.88913,1.018254 0,1 l 2.19337,-0.004 -0.006,0.004 -3.27651,2.873468 c -0.64989,0.580999 0.2216,1.555837 0.87149,0.97484 L 412.00004,84.45 c 0,1.651946 1.15621,3.581251 3.47506,3.550001 C 417.79395,87.968751 419,86.250706 419,84.45 c 0,-1.800705 -1.00462,-2.558141 -1.45954,-3.013074 z M 415.5,82.2 c 1.22478,0 2.25,0.945047 2.25,2.25 0,1.304953 -1.05311,2.25 -2.25,2.25 -1.19689,0 -2.25,-0.960585 -2.25,-2.25 0,-1.289415 1.02522,-2.25 2.25,-2.25 z"
+           style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1.3066;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:accumulate" />
+      </g>
+      <g
+         style="display:inline;fill:#ffffff;enable-background:new"
+         id="g15543-3-3"
+         transform="translate(21.029034,-20.999943)">
+        <g
+           id="g15520-6-5"
+           transform="translate(231.97182,-397.99995)"
+           style="display:inline;opacity:0.6;fill:#ffffff;enable-background:new"
+           inkscape:export-filename="C:\Users\Andrzej Ambroż\Desktop\mtrx.png"
+           inkscape:export-xdpi="96"
+           inkscape:export-ydpi="96">
+          <g
+             id="g4103-6">
+            <path
+               style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
+               d="m 430.48047,53 c -0.15153,0.004 -0.29304,0.07659 -0.38477,0.197266 l -3.94922,3.949218 C 425.83169,57.461484 426.05468,57.99983 426.5,58 h 4 c 0.27613,-3e-5 0.49997,-0.22387 0.5,-0.5 V 54 h 6 v 2.999953 c -0.01,0.676161 1.00956,0.676161 1,0 V 53.5 c -3e-5,-0.276131 -0.22387,-0.499972 -0.5,-0.5 h -7 c -0.005,-6.2e-5 -0.009,-6.2e-5 -0.0137,0 -6.7e-4,1.8e-5 -0.001,-2.1e-5 -0.002,0 -0.001,4.1e-5 -0.003,-5e-5 -0.004,0 z M 426,59.000004 V 66.5 c 3e-5,0.276131 0.22387,0.499972 0.5,0.5 h 4.49914 c 0.67616,0.0096 0.67616,-1.009563 0,-1 H 427 v -6.999996 z"
+               transform="translate(-273.99999,439.99994)"
+               id="path15514-7-2"
+               inkscape:connector-curvature="0"
+               sodipodi:nodetypes="ccccccccccccccccccccccccc" />
+          </g>
+        </g>
+      </g>
+    </g>
+    <g
+       transform="translate(499.00086,52.015151)"
+       id="g824"
+       style="display:inline;enable-background:new">
+      <g
+         id="g4207"
+         style="display:inline;enable-background:new"
+         transform="translate(-499.00079,-52.007172)">
+        <g
+           transform="translate(105,-21.000376)"
+           style="display:inline;enable-background:new"
+           id="g7580-6">
+          <g
+             transform="matrix(0.6740384,0,0,0.6740384,192.80592,-339.68227)"
+             style="display:inline;opacity:0.99;fill:#ffffff;stroke-width:1.07692;enable-background:new"
+             id="g8599-6-7-2"
+             inkscape:export-filename="C:\Users\Andrzej Ambroż\Desktop\mtrx.png"
+             inkscape:export-xdpi="96"
+             inkscape:export-ydpi="96">
+            <path
+               id="path8597-7-0-9"
+               d="m 331.87142,629.2294 c 0,0.81936 -0.66423,1.48359 -1.48359,1.48359 -0.81937,0 -1.4836,-0.66422 -1.4836,-1.48359 0,-0.81937 0.66423,-1.4836 1.4836,-1.4836 0.81937,0 1.48359,0.66423 1.48359,1.4836 z"
+               style="opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;stroke:none;stroke-width:1.72218;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke"
+               inkscape:connector-curvature="0"
+               sodipodi:nodetypes="sssss" />
+          </g>
+          <path
+             sodipodi:nodetypes="cccccccccczzcczzzzz"
+             inkscape:connector-curvature="0"
+             id="path8595-5-9-1"
+             d="m 414.93725,78.992499 c -0.65344,-0.653443 -1.58833,0.255453 -0.88297,0.960812 l 1.12739,1.039192 h -4.6057 c -0.88913,-0.01822 -0.88913,1.018254 0,1 l 2.19337,-0.004 -0.006,0.004 -3.27651,2.873468 c -0.64989,0.580999 0.2216,1.555837 0.87149,0.97484 l 1.64161,-1.398308 c 0,1.651946 1.15632,3.581251 3.47517,3.550001 2.31885,-0.03125 3.5249,-1.749295 3.5249,-3.550001 0,-1.800705 -1.00462,-2.558141 -1.45954,-3.013074 z M 415.5,82.192503 c 1.22478,0 2.25,0.945047 2.25,2.25 0,1.304953 -1.05311,2.25 -2.25,2.25 -1.19689,0 -2.25,-0.960585 -2.25,-2.25 0,-1.289415 1.02522,-2.25 2.25,-2.25 z"
+             style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1.3066;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:accumulate" />
+        </g>
+      <g
+           transform="translate(19.000003)"
+           id="g28228-3-3"
+           style="display:inline;opacity:0.6;stroke:#ffffff;enable-background:new"
+         inkscape:export-filename="C:\Users\Andrzej Ambroż\Desktop\mtrx.png"
+         inkscape:export-xdpi="96"
+         inkscape:export-ydpi="96">
+          <g
+             id="g28217-6-6"
+             transform="translate(338.99999,-439.99995)"
+             style="display:inline;opacity:0.99;stroke:#ffffff;enable-background:new">
+        <path
+           inkscape:connector-curvature="0"
+               id="path28215-0"
+               transform="translate(-337.99999,439.99995)"
+               d="m 501.49219,52.992188 c -0.27615,0.0043 -0.49651,0.223792 -0.49219,0.499938 v 1 c -0.01,0.676161 1.00956,0.676161 1,0 v -1 c 0.004,-0.282265 -0.22554,-0.504353 -0.50781,-0.499938 z m -7.00781,0.0067 c -0.12718,0.004 -0.248,0.0564 -0.3379,0.146484 l -4,4.001116 c -0.10126,0.101337 -0.1304,0.223491 -0.13086,0.345704 L 490,57.507812 v 1.984314 c -0.01,0.676161 1.00956,0.676161 1,0 v -1.5 h 3.5 c 0.27613,-3e-5 0.49997,-0.22387 0.5,-0.5 v -3.5 h 1.5 c 0.67616,0.0096 0.67616,-1.002805 0,-0.993242 h -2 v 0.002 c -0.005,-5e-6 -0.0101,-0.0021 -0.0156,-0.002 z m 4.01562,-0.0068 c -0.67616,-0.0096 -0.67616,1.009563 0,1 h 1 c 0.67616,0.0096 0.67616,-1.009563 0,-1 z m 2.99219,3.000062 c -0.27615,0.0043 -0.49651,0.223792 -0.49219,0.499938 v 1 c -0.01,0.676161 1.00956,0.676161 1,0 v -1 c 0.004,-0.282265 -0.22554,-0.504353 -0.50781,-0.499938 z m -11,5 c -0.27615,0.0043 -0.49651,0.223792 -0.49219,0.499938 v 1 c -0.01,0.676161 1.00956,0.676161 1,0 v -1 c 0.004,-0.282265 -0.22554,-0.504353 -0.50781,-0.499938 z m 0,3 c -0.27615,0.0043 -0.49651,0.223792 -0.49219,0.499938 v 1 c -0.01,0.676161 1.00956,0.676161 1,0 v -1 c 0.004,-0.282265 -0.22554,-0.504353 -0.50781,-0.499938 z M 492.5,65.992126 c -0.67616,-0.0096 -0.67616,1.009563 0,1 h 1 c 0.67616,0.0096 0.67616,-1.009563 0,-1 z m 3,0 c -0.67616,-0.0096 -0.67616,1.009563 0,1 h 1 c 0.67616,0.0096 0.67616,-1.009563 0,-1 z"
+               style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:accumulate"
+               sodipodi:nodetypes="cccccccccccccccccccccccccccccccccccccccccccccccccccccccc" />
+      </g>
+        </g>
+      </g>
+    </g>
+    <g
+       transform="translate(-0.999993)"
+       id="g4126"
+       style="display:inline;enable-background:new">
+      <g
+         style="display:inline;fill:#ffffff;enable-background:new"
+         inkscape:export-ydpi="96"
+         inkscape:export-xdpi="96"
+         inkscape:export-filename="C:\Users\Andrzej Ambroż\Desktop\mtrx.png"
+         transform="matrix(0.65914281,0,0,0.65914281,248.47102,-214.89549)"
+         id="g12839-3-5">
+        <path
+           style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
+           d="m 352.97527,473.19666 c -0.49613,-0.0451 -1.03862,0.15972 -1.49219,0.61328 l -0.41128,0.47153 c -0.19519,0.19527 -0.19519,0.51177 0,0.70704 l 3,3 c 0.19527,0.19519 0.51177,0.19519 0.70704,0 l 0.41128,-0.47153 c 0.45356,-0.45357 0.65838,-0.99606 0.61328,-1.49219 -0.0451,-0.49613 -0.30436,-0.90592 -0.61328,-1.21485 l -1,-1 c -0.30893,-0.30892 -0.71872,-0.56817 -1.21485,-0.61328 z m -3.39644,2.7168 c -0.12989,0.002 -0.25387,0.0546 -0.34571,0.14648 l -6.22807,6.22726 c -0.0639,0.0642 -0.10909,0.14453 -0.13086,0.23243 l -1,4 c -0.0904,0.36537 0.2401,0.69582 0.60547,0.60547 l 4,-1 c 0.0879,-0.0218 0.16823,-0.067 0.23243,-0.13086 0,0 6.17898,-6.1737 6.22807,-6.22726 0.19519,-0.19527 0.19519,-0.51177 0,-0.70704 l -3,-3 c -0.0957,-0.0957 -0.22605,-0.14856 -0.36137,-0.14648 z"
+           id="path12837-6-3"
+           inkscape:connector-curvature="0"
+           sodipodi:nodetypes="cccccccscccccccccccccccc" />
+      </g>
+        <path
+         sodipodi:nodetypes="csscccccccccccssscccc"
+         style="opacity:0.6;fill:#ffffff"
+           inkscape:connector-curvature="0"
+         id="path2-6"
+         d="m 469,101 v 7.5 c 0,0.276 0.224,0.5 0.5,0.5 h 11 c 0.30423,0 0.5,-0.22782 0.5,-0.5 v -4 c 0,-0.65459 -1,-0.65682 -1,0 v 3.5 h -10 v -7 z m 4.48081,-6 c -0.151,0.004 -0.293,0.077 -0.384,0.197 l -3.95,3.949 c -0.314,0.315 -0.091,0.854 0.354,0.854 h 4 c 0.276,0 0.5,-0.224 0.5,-0.5 V 96 H 480.5 c 0.68512,0 0.64092,-1 0,-1 z" />
+    </g>
   </g>
   <g
      inkscape:groupmode="layer"