improve handling of incompatible enums

CMakeLists.txt

@@ -411,7 +411,6 @@ option(WITH_CYCLES                   "Enable Cycles Render Engine" ON)
 option(WITH_CYCLES_OSL               "Build Cycles with OpenShadingLanguage support" ON)
 option(WITH_CYCLES_EMBREE            "Build Cycles with Embree support" ON)
 option(WITH_CYCLES_LOGGING           "Build Cycles with logging support" ON)
-option(WITH_CYCLES_DEBUG             "Build Cycles with options useful for debugging (e.g., MIS)" OFF)
 
 option(WITH_CYCLES_STANDALONE        "Build Cycles standalone application" OFF)
 option(WITH_CYCLES_STANDALONE_GUI    "Build Cycles standalone with GUI" OFF)
@@ -441,11 +440,7 @@ mark_as_advanced(WITH_CYCLES_CUDA_BUILD_SERIAL)
 mark_as_advanced(WITH_CUDA_DYNLOAD)
 
 # AMD HIP
-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_DEVICE_HIP        "Enable Cycles AMD HIP support" OFF)
 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)
@@ -514,13 +509,9 @@ if(UNIX AND NOT APPLE)
 endif()
 
 
-# Vulkan
-option(WITH_VULKAN                    "Enable Vulkan backend (Experimental)" OFF)
-option(WITH_VULKAN_SHADER_COMPILATION "Temporary flag to enable vulkan shader compilation needed to continue development during the migration of GLSL to Vulkan." OFF)
-
 # OpenGL
 
-option(WITH_OPENGL              "When off limits visibility of the opengl headers to just bf_gpu (temporary option for development purposes)" ON)
+option(WITH_OPENGL              "When off limits visibility of the opengl headers to just bf_gpu and gawain (temporary option for development purposes)" ON)
 option(WITH_GLEW_ES             "Switches to experimental copy of GLEW that has support for OpenGL ES. (temporary option for development purposes)" OFF)
 option(WITH_GL_EGL              "Use the EGL OpenGL system library instead of the platform specific OpenGL system library (CGL, glX, or WGL)"       OFF)
 option(WITH_GL_PROFILE_ES20     "Support using OpenGL ES 2.0. (through either EGL or the AGL/WGL/XGL 'es20' profile)"                               OFF)
@@ -530,7 +521,6 @@ mark_as_advanced(
   WITH_GLEW_ES
   WITH_GL_EGL
   WITH_GL_PROFILE_ES20
-  WITH_VULKAN_SHADER_COMPILATION
 )
 
 if(WIN32)
@@ -1075,7 +1065,7 @@ if(MSVC)
   add_definitions(-D__LITTLE_ENDIAN__)
 
   # OSX-Note: as we do cross-compiling with specific set architecture,
-  # endianness-detection and auto-setting is counterproductive
+  # endianess-detection and auto-setting is counterproductive
   # so we just set endianness according CMAKE_OSX_ARCHITECTURES
 
 elseif(CMAKE_OSX_ARCHITECTURES MATCHES i386 OR CMAKE_OSX_ARCHITECTURES MATCHES x86_64 OR CMAKE_OSX_ARCHITECTURES MATCHES arm64)
@@ -1131,18 +1121,6 @@ if(WITH_OPENVDB)
   list(APPEND OPENVDB_LIBRARIES ${BOOST_LIBRARIES} ${TBB_LIBRARIES})
 endif()
 
-#-----------------------------------------------------------------------------
-# Configure Vulkan.
-
-if(WITH_VULKAN)
-  list(APPEND BLENDER_GL_LIBRARIES ${Vulkan_LIBRARY})
-
-  add_definitions(-DWITH_VULKAN)
-  if(WITH_VULKAN_SHADER_COMPILATION)
-    add_definitions(-DWITH_VULKAN_SHADER_COMPILATION)
-  endif()
-endif()
-
 #-----------------------------------------------------------------------------
 # Configure OpenGL.
 
@@ -1777,7 +1755,7 @@ endif()
 set(CMAKE_CXX_STANDARD 17)
 # If C++17 is not available, downgrading to an earlier standard is NOT OK.
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
-# Do not enable compiler specific language extensions.
+# Do not enable compiler specific language extentions.
 set(CMAKE_CXX_EXTENSIONS OFF)
 
 # Make MSVC properly report the value of the __cplusplus preprocessor macro

GNUmakefile

@@ -51,7 +51,7 @@ Other Convenience Targets
    * config:        Run cmake configuration tool to set build options.
    * deps:          Build library dependencies (intended only for platform maintainers).
 
-                    The existance of locally build dependencies overrides the pre-built dependencies from subversion.
+                    The existance of locally build dependancies overrides the pre-built dependencies from subversion.
                     These must be manually removed from '../lib/' to go back to using the pre-compiled libraries.
 
 Project Files

build_files/build_environment/cmake/harvest.cmake

@@ -17,7 +17,7 @@
 # ***** END GPL LICENSE BLOCK *****
 
 ########################################################################
-# Copy all generated files to the proper structure as blender prefers
+# Copy all generated files to the proper strucure as blender prefers
 ########################################################################
 
 if(NOT DEFINED HARVEST_TARGET)

build_files/build_environment/cmake/nanovdb.cmake

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

build_files/build_environment/cmake/options.cmake

@@ -39,7 +39,7 @@ endif()
 set(DOWNLOAD_DIR "${CMAKE_CURRENT_BINARY_DIR}/downloads" CACHE STRING "Path for downloaded files")
 # This path must be hard-coded like this, so that the GNUmakefile knows where it is and can pass it to make_source_archive.py:
 set(PACKAGE_DIR "${CMAKE_CURRENT_BINARY_DIR}/packages")
-option(PACKAGE_USE_UPSTREAM_SOURCES "Use sources upstream to download the package sources, when OFF the blender mirror will be used" ON)
+option(PACKAGE_USE_UPSTREAM_SOURCES "Use soures upstream to download the package sources, when OFF the blender mirror will be used" ON)
 
 file(TO_CMAKE_PATH ${DOWNLOAD_DIR} DOWNLOAD_DIR)
 file(TO_CMAKE_PATH ${PACKAGE_DIR} PACKAGE_DIR)

build_files/build_environment/cmake/package_python.cmake

@@ -24,7 +24,7 @@ if(MSVC)
     add_custom_command(
       OUTPUT ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
       COMMAND echo packaging python
-      COMMAND echo this should output at ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
+      COMMAND echo this should ouput at ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
       COMMAND ${CMAKE_COMMAND} -E make_directory ${PYTARGET}/libs
       COMMAND ${CMAKE_COMMAND} -E copy ${PYSRC}/libs/python${PYTHON_SHORT_VERSION_NO_DOTS}.lib ${PYTARGET}/libs/python${PYTHON_SHORT_VERSION_NO_DOTS}.lib
       COMMAND ${CMAKE_COMMAND} -E copy ${PYSRC}/python.exe ${PYTARGET}/bin/python.exe
@@ -43,7 +43,7 @@ if(MSVC)
     add_custom_command(
       OUTPUT ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
       COMMAND echo packaging python
-      COMMAND echo this should output at ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
+      COMMAND echo this should ouput at ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
       COMMAND ${CMAKE_COMMAND} -E make_directory ${PYTARGET}/libs
       COMMAND ${CMAKE_COMMAND} -E copy ${PYSRC}/libs/python${PYTHON_SHORT_VERSION_NO_DOTS}${PYTHON_POSTFIX}.lib ${PYTARGET}/libs/python${PYTHON_SHORT_VERSION_NO_DOTS}${PYTHON_POSTFIX}.lib
       COMMAND ${CMAKE_COMMAND} -E copy ${PYSRC}/python${PYTHON_POSTFIX}.exe ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe

build_files/build_environment/install_deps.sh

@@ -1826,7 +1826,7 @@ compile_OCIO() {
     # Force linking against static libs
     #rm -f $_inst/lib/*.so*
 
-    # Additional dependencies
+    # Additional depencencies
     #cp ext/dist/lib/libtinyxml.a $_inst/lib
     #cp ext/dist/lib/libyaml-cpp.a $_inst/lib
 

build_files/build_environment/patches/nanovdb.diff

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

build_files/cmake/Modules/FindShaderC.cmake

@@ -1,66 +0,0 @@
-# - Find SHADERC library
-# Find the native Haru includes and library
-# This module defines
-#  SHADERC_INCLUDE_DIRS, where to find hpdf.h, set when
-#                        SHADERC_INCLUDE_DIR is found.
-#  SHADERC_LIBRARIES, libraries to link against to use Haru.
-#  SHADERC_ROOT_DIR, The base directory to search for Haru.
-#                    This can also be an environment variable.
-#  SHADERC_FOUND, If false, do not try to use Haru.
-#
-# also defined, but not for general use are
-#  SHADERC_LIBRARY, where to find the Haru library.
-
-#=============================================================================
-# Copyright 2021 Blender Foundation.
-#
-# Distributed under the OSI-approved BSD 3-Clause License,
-# see accompanying file BSD-3-Clause-license.txt for details.
-#=============================================================================
-
-# If SHADERC_ROOT_DIR was defined in the environment, use it.
-if(NOT SHADERC_ROOT_DIR AND NOT $ENV{SHADERC_ROOT_DIR} STREQUAL "")
-  set(SHADERC_ROOT_DIR $ENV{SHADERC_ROOT_DIR})
-endif()
-
-set(_shaderc_SEARCH_DIRS
-  ${SHADERC_ROOT_DIR}
-  /opt/lib/haru
-)
-
-find_path(SHADERC_INCLUDE_DIR
-  NAMES
-    shaderc.hpp
-  HINTS
-    ${_shaderc_SEARCH_DIRS}
-  PATH_SUFFIXES
-    include/shaderc
-    include
-)
-
-find_library(SHADERC_LIBRARY
-  NAMES
-    shaderc_combined
-    shaderc
-  HINTS
-    ${_shaderc_SEARCH_DIRS}
-  PATH_SUFFIXES
-    lib64 lib
-)
-
-# Handle the QUIETLY and REQUIRED arguments and set SHADERC_FOUND to TRUE if
-# all listed variables are TRUE.
-include(FindPackageHandleStandardArgs)
-find_package_handle_standard_args(ShaderC DEFAULT_MSG SHADERC_LIBRARY SHADERC_INCLUDE_DIR)
-
-if(SHADERC_FOUND)
-  set(SHADERC_LIBRARIES ${SHADERC_LIBRARY})
-  set(SHADERC_INCLUDE_DIRS ${SHADERC_INCLUDE_DIR})
-endif()
-
-mark_as_advanced(
-  SHADERC_INCLUDE_DIR
-  SHADERC_LIBRARY
-)
-
-unset(_shaderc_SEARCH_DIRS)

build_files/cmake/cmake_netbeans_project.py

@@ -180,7 +180,7 @@ def create_nb_project_main():
         f.write('    </logicalFolder>\n')
 
         f.write('  </logicalFolder>\n')
-        # default, but this dir is in fact not in blender dir so we can ignore it
+        # default, but this dir is infact not in blender dir so we can ignore it
         # f.write('  <sourceFolderFilter>^(nbproject)$</sourceFolderFilter>\n')
         f.write(r'  <sourceFolderFilter>^(nbproject|__pycache__|.*\.py|.*\.html|.*\.blend)$</sourceFolderFilter>\n')
 

build_files/cmake/config/blender_release.cmake

@@ -81,5 +81,4 @@ 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()

build_files/cmake/macros.cmake

@@ -529,7 +529,7 @@ function(SETUP_LIBDIRS)
 
   # NOTE: For all new libraries, use absolute library paths.
   # This should eventually be phased out.
-  # APPLE platform uses full paths for linking libraries, and avoids link_directories.
+  # APPLE plaform uses full paths for linking libraries, and avoids link_directories.
   if(NOT MSVC AND NOT APPLE)
     link_directories(${JPEG_LIBPATH} ${PNG_LIBPATH} ${ZLIB_LIBPATH} ${FREETYPE_LIBPATH})
 

build_files/cmake/platform/platform_unix.cmake

@@ -102,11 +102,6 @@ find_package_wrapper(ZLIB REQUIRED)
 find_package_wrapper(Zstd REQUIRED)
 find_package_wrapper(Freetype REQUIRED)
 
-if(WITH_VULKAN)
-  find_package_wrapper(Vulkan REQUIRED)
-  find_package(ShaderC REQUIRED)
-endif()
-
 if(WITH_PYTHON)
   # No way to set py35, remove for now.
   # find_package(PythonLibs)

build_files/cmake/platform/platform_win32.cmake

@@ -874,32 +874,5 @@ if(WITH_HARU)
   endif()
 endif()
 
-if(WITH_VULKAN)
-  if(EXISTS ${LIBDIR}/vulkan)
-    set(Vulkan_FOUND On)
-    set(Vulkan_ROOT_DIR ${LIBDIR}/vulkan)
-    set(Vulkan_INCLUDE_DIR ${Vulkan_ROOT_DIR}/include)
-    set(Vulkan_INCLUDE_DIRS ${Vulkan_INCLUDE_DIR})
-    set(Vulkan_LIBRARY ${Vulkan_ROOT_DIR}/lib/vulkan-1.lib)
-    set(Vulkan_LIBRARIES ${Vulkan_LIBRARY})
-  else()
-    message(WARNING "vulkan was not found, disabling WITH_VULKAN")
-    set(WITH_VULKAN OFF)
-  endif()
-endif()
-
-if(WITH_VULKAN)
-  if(EXISTS ${LIBDIR}/shaderc)
-    set(SHADERC_ROOT_DIR ${LIBDIR}/shaderc)
-    set(SHADERC_INCLUDE_DIR ${SHADERC_ROOT_DIR}/include)
-    set(SHADERC_INCLUDE_DIRS ${SHADERC_INCLUDE_DIR})
-    set(SHADERC_LIBRARY optimized ${SHADERC_ROOT_DIR}/lib/shaderc_shared.lib debug ${SHADERC_ROOT_DIR}/lib/shaderc_shared_d.lib)
-    set(SHADERC_LIBRARIES ${SHADERC_LIBRARY})
-  else()
-    message(WARNING "shaderc was not found, disabling WITH_VULKAN")
-    set(WITH_VULKAN OFF)
-  endif()
-endif()
-
 set(ZSTD_INCLUDE_DIRS ${LIBDIR}/zstd/include)
 set(ZSTD_LIBRARIES ${LIBDIR}/zstd/lib/zstd_static.lib)

build_files/cmake/platform/platform_win32_bundle_crt.cmake

@@ -27,7 +27,7 @@ if(WITH_WINDOWS_BUNDLE_CRT)
   # Install the CRT to the blender.crt Sub folder.
   install(FILES ${CMAKE_INSTALL_SYSTEM_RUNTIME_LIBS} DESTINATION ./blender.crt COMPONENT Libraries)
 
-  # Generating the manifest is a relatively expensive operation since
+  # Generating the manifest is a relativly expensive operation since
   # it is collecting an sha1 hash for every file required. so only do
   # this work when the libs have either changed or the manifest does
   # not exist yet.

doc/python_api/examples/bpy.app.timers.5.py

@@ -11,7 +11,7 @@ import queue
 
 execution_queue = queue.Queue()
 
-# This function can safely be called in another thread.
+# This function can savely be called in another thread.
 # The function will be executed when the timer runs the next time.
 def run_in_main_thread(function):
     execution_queue.put(function)

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

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

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

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

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

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

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

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

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

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

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

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

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

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

doc/python_api/rst/info_gotcha.rst

@@ -728,7 +728,7 @@ Abusing RNA property callbacks
 ------------------------------
 
 Python-defined RNA properties can have custom callbacks. Trying to perform complex operations
-from there, like calling an operator, may work, but is not officially recommended nor supported.
+from there, like calling an operator, may work, but is not officialy recommended nor supported.
 
 Main reason is that those callback should be very fast, but additionally, it may for example
 create issues with undo/redo system (most operators store an history step, and editing an RNA

extern/CMakeLists.txt

@@ -116,7 +116,3 @@ endif()
 if (WITH_COMPOSITOR)
   add_subdirectory(smaa_areatex)
 endif()
-
-if(WITH_VULKAN)
-  add_subdirectory(vulkan_memory_allocator)
-endif()

extern/hipew/include/hipew.h

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

extern/hipew/src/hipew.c

@@ -71,7 +71,6 @@ thipDriverGetVersion *hipDriverGetVersion;
 thipGetDevice *hipGetDevice;
 thipGetDeviceCount *hipGetDeviceCount;
 thipGetDeviceProperties *hipGetDeviceProperties;
-thipDeviceGet* hipDeviceGet;
 thipDeviceGetName *hipDeviceGetName;
 thipDeviceGetAttribute *hipDeviceGetAttribute;
 thipDeviceComputeCapability *hipDeviceComputeCapability;
@@ -256,7 +255,6 @@ static int hipewHipInit(void) {
   HIP_LIBRARY_FIND_CHECKED(hipGetDevice);
   HIP_LIBRARY_FIND_CHECKED(hipGetDeviceCount);
   HIP_LIBRARY_FIND_CHECKED(hipGetDeviceProperties);
-  HIP_LIBRARY_FIND_CHECKED(hipDeviceGet);
   HIP_LIBRARY_FIND_CHECKED(hipDeviceGetName);
   HIP_LIBRARY_FIND_CHECKED(hipDeviceGetAttribute);
   HIP_LIBRARY_FIND_CHECKED(hipDeviceComputeCapability);

extern/vulkan_memory_allocator/CMakeLists.txt

@@ -1,42 +0,0 @@
-# ***** BEGIN GPL LICENSE BLOCK *****
-#
-# This program is free software; you can redistribute it and/or
-# modify it under the terms of the GNU General Public License
-# as published by the Free Software Foundation; either version 2
-# of the License, or (at your option) any later version.
-#
-# This program is distributed in the hope that it will be useful,
-# but WITHOUT ANY WARRANTY; without even the implied warranty of
-# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-# GNU General Public License for more details.
-#
-# You should have received a copy of the GNU General Public License
-# along with this program; if not, write to the Free Software Foundation,
-# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-#
-# The Original Code is Copyright (C) 2012, Blender Foundation
-# All rights reserved.
-# ***** END GPL LICENSE BLOCK *****
-
-
-set(INC
-  .
-)
-
-set(INC_SYS
-  ${Vulkan_INCLUDE_DIRS}
-)
-
-set(SRC
-  vk_mem_alloc_impl.cc
-
-  vk_mem_alloc.h
-)
-
-blender_add_lib(extern_vulkan_memory_allocator "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")
-
-if(CMAKE_COMPILER_IS_GNUCC OR CMAKE_C_COMPILER_ID MATCHES "Clang")
-  target_compile_options(extern_vulkan_memory_allocator
-    PRIVATE "-Wno-nullability-completeness"
-  )
-endif()

extern/vulkan_memory_allocator/LICENSE.txt

@@ -1,19 +0,0 @@
-Copyright (c) 2017-2021 Advanced Micro Devices, Inc. All rights reserved.
-
-Permission is hereby granted, free of charge, to any person obtaining a copy
-of this software and associated documentation files (the "Software"), to deal
-in the Software without restriction, including without limitation the rights
-to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-copies of the Software, and to permit persons to whom the Software is
-furnished to do so, subject to the following conditions:
-
-The above copyright notice and this permission notice shall be included in
-all copies or substantial portions of the Software.
-
-THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
-AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
-THE SOFTWARE.

extern/vulkan_memory_allocator/README.blender

@@ -1,5 +0,0 @@
-Project: VulkanMemoryAllocator
-URL: https://github.com/GPUOpen-LibrariesAndSDKs/VulkanMemoryAllocator
-License: MIT
-Upstream version: 4b047fd
-Local modifications: None

extern/vulkan_memory_allocator/README.md

@@ -1,134 +0,0 @@
-# Vulkan Memory Allocator
-
-Easy to integrate Vulkan memory allocation library.
-
-**Documentation:** See [Vulkan Memory Allocator](https://gpuopen-librariesandsdks.github.io/VulkanMemoryAllocator/html/) (generated from Doxygen-style comments in [src/vk_mem_alloc.h](src/vk_mem_alloc.h))
-
-**License:** MIT. See [LICENSE.txt](LICENSE.txt)
-
-**Changelog:** See [CHANGELOG.md](CHANGELOG.md)
-
-**Product page:** [Vulkan Memory Allocator on GPUOpen](https://gpuopen.com/gaming-product/vulkan-memory-allocator/)
-
-**Build status:**
-
-- Windows: [![Build status](https://ci.appveyor.com/api/projects/status/4vlcrb0emkaio2pn/branch/master?svg=true)](https://ci.appveyor.com/project/adam-sawicki-amd/vulkanmemoryallocator/branch/master)  
-- Linux: [![Build Status](https://travis-ci.com/GPUOpen-LibrariesAndSDKs/VulkanMemoryAllocator.svg?branch=master)](https://travis-ci.com/GPUOpen-LibrariesAndSDKs/VulkanMemoryAllocator)
-
-# Problem
-
-Memory allocation and resource (buffer and image) creation in Vulkan is difficult (comparing to older graphics API-s, like D3D11 or OpenGL) for several reasons:
-
-- It requires a lot of boilerplate code, just like everything else in Vulkan, because it is a low-level and high-performance API.
-- There is additional level of indirection: `VkDeviceMemory` is allocated separately from creating `VkBuffer`/`VkImage` and they must be bound together.
-- Driver must be queried for supported memory heaps and memory types. Different IHVs provide different types of it.
-- It is recommended practice to allocate bigger chunks of memory and assign parts of them to particular resources.
-
-# Features
-
-This library can help game developers to manage memory allocations and resource creation by offering some higher-level functions:
-
-1. Functions that help to choose correct and optimal memory type based on intended usage of the memory.
-   - Required or preferred traits of the memory are expressed using higher-level description comparing to Vulkan flags.
-2. Functions that allocate memory blocks, reserve and return parts of them (`VkDeviceMemory` + offset + size) to the user.
-   - Library keeps track of allocated memory blocks, used and unused ranges inside them, finds best matching unused ranges for new allocations, respects all the rules of alignment and buffer/image granularity.
-3. Functions that can create an image/buffer, allocate memory for it and bind them together - all in one call.
-
-Additional features:
-
-- Well-documented - description of all functions and structures provided, along with chapters that contain general description and example code.
-- Thread-safety: Library is designed to be used in multithreaded code. Access to a single device memory block referred by different buffers and textures (binding, mapping) is synchronized internally.
-- Configuration: Fill optional members of CreateInfo structure to provide custom CPU memory allocator, pointers to Vulkan functions and other parameters.
-- Customization: Predefine appropriate macros to provide your own implementation of all external facilities used by the library, from assert, mutex, and atomic, to vector and linked list. 
-- Support for memory mapping, reference-counted internally. Support for persistently mapped memory: Just allocate with appropriate flag and you get access to mapped pointer.
-- Support for non-coherent memory. Functions that flush/invalidate memory. `nonCoherentAtomSize` is respected automatically.
-- Support for resource aliasing (overlap).
-- Support for sparse binding and sparse residency: Convenience functions that allocate or free multiple memory pages at once.
-- Custom memory pools: Create a pool with desired parameters (e.g. fixed or limited maximum size) and allocate memory out of it.
-- Linear allocator: Create a pool with linear algorithm and use it for much faster allocations and deallocations in free-at-once, stack, double stack, or ring buffer fashion.
-- Support for Vulkan 1.0, 1.1, 1.2.
-- Support for extensions (and equivalent functionality included in new Vulkan versions):
-   - VK_EXT_memory_budget: Used internally if available to query for current usage and budget. If not available, it falls back to an estimation based on memory heap sizes.
-   - VK_KHR_dedicated_allocation: Just enable it and it will be used automatically by the library.
-   - VK_AMD_device_coherent_memory
-   - VK_KHR_buffer_device_address
-- Defragmentation of GPU and CPU memory: Let the library move data around to free some memory blocks and make your allocations better compacted.
-- Lost allocations: Allocate memory with appropriate flags and let the library remove allocations that are not used for many frames to make room for new ones.
-- Statistics: Obtain detailed statistics about the amount of memory used, unused, number of allocated blocks, number of allocations etc. - globally, per memory heap, and per memory type.
-- Debug annotations: Associate string with name or opaque pointer to your own data with every allocation.
-- JSON dump: Obtain a string in JSON format with detailed map of internal state, including list of allocations and gaps between them.
-- Convert this JSON dump into a picture to visualize your memory. See [tools/VmaDumpVis](tools/VmaDumpVis/README.md).
-- Debugging incorrect memory usage: Enable initialization of all allocated memory with a bit pattern to detect usage of uninitialized or freed memory. Enable validation of a magic number before and after every allocation to detect out-of-bounds memory corruption.
-- Record and replay sequence of calls to library functions to a file to check correctness, measure performance, and gather statistics.
-
-# Prequisites
-
-- Self-contained C++ library in single header file. No external dependencies other than standard C and C++ library and of course Vulkan. STL containers are not used by default.
-- Public interface in C, in same convention as Vulkan API. Implementation in C++.
-- Error handling implemented by returning `VkResult` error codes - same way as in Vulkan.
-- Interface documented using Doxygen-style comments.
-- Platform-independent, but developed and tested on Windows using Visual Studio. Continuous integration setup for Windows and Linux. Used also on Android, MacOS, and other platforms.
-
-# Example
-
-Basic usage of this library is very simple. Advanced features are optional. After you created global `VmaAllocator` object, a complete code needed to create a buffer may look like this:
-
-```cpp
-VkBufferCreateInfo bufferInfo = { VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO };
-bufferInfo.size = 65536;
-bufferInfo.usage = VK_BUFFER_USAGE_VERTEX_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT;
-
-VmaAllocationCreateInfo allocInfo = {};
-allocInfo.usage = VMA_MEMORY_USAGE_GPU_ONLY;
-
-VkBuffer buffer;
-VmaAllocation allocation;
-vmaCreateBuffer(allocator, &bufferInfo, &allocInfo, &buffer, &allocation, nullptr);
-```
-
-With this one function call:
-
-1. `VkBuffer` is created.
-2. `VkDeviceMemory` block is allocated if needed.
-3. An unused region of the memory block is bound to this buffer.
-
-`VmaAllocation` is an object that represents memory assigned to this buffer. It can be queried for parameters like Vulkan memory handle and offset.
-
-# Binaries
-
-The release comes with precompiled binary executables for "VulkanSample" application which contains test suite and "VmaReplay" tool. They are compiled using Visual Studio 2019, so they require appropriate libraries to work, including "MSVCP140.dll", "VCRUNTIME140.dll", "VCRUNTIME140_1.dll". If their launch fails with error message telling about those files missing, please download and install [Microsoft Visual C++ Redistributable for Visual Studio 2015, 2017 and 2019](https://support.microsoft.com/en-us/help/2977003/the-latest-supported-visual-c-downloads), "x64" version.
-
-# Read more
-
-See **[Documentation](https://gpuopen-librariesandsdks.github.io/VulkanMemoryAllocator/html/)**.
-
-# Software using this library
-
-- **[Detroit: Become Human](https://gpuopen.com/learn/porting-detroit-3/)**
-- **[Vulkan Samples](https://github.com/LunarG/VulkanSamples)** - official Khronos Vulkan samples. License: Apache-style.
-- **[Anvil](https://github.com/GPUOpen-LibrariesAndSDKs/Anvil)** - cross-platform framework for Vulkan. License: MIT.
-- **[Filament](https://github.com/google/filament)** - physically based rendering engine for Android, Windows, Linux and macOS, from Google. Apache License 2.0.
-- **[Atypical Games - proprietary game engine](https://developer.samsung.com/galaxy-gamedev/gamedev-blog/infinitejet.html)**
-- **[Flax Engine](https://flaxengine.com/)**
-- **[Lightweight Java Game Library (LWJGL)](https://www.lwjgl.org/)** - includes binding of the library for Java. License: BSD.
-- **[PowerVR SDK](https://github.com/powervr-graphics/Native_SDK)** - C++ cross-platform 3D graphics SDK, from Imagination. License: MIT.
-- **[Skia](https://github.com/google/skia)** - complete 2D graphic library for drawing Text, Geometries, and Images, from Google.
-- **[The Forge](https://github.com/ConfettiFX/The-Forge)** - cross-platform rendering framework. Apache License 2.0.
-- **[VK9](https://github.com/disks86/VK9)** - Direct3D 9 compatibility layer using Vulkan. Zlib lincese.
-- **[vkDOOM3](https://github.com/DustinHLand/vkDOOM3)** - Vulkan port of GPL DOOM 3 BFG Edition. License: GNU GPL.
-- **[vkQuake2](https://github.com/kondrak/vkQuake2)** - vanilla Quake 2 with Vulkan support. License: GNU GPL.
-- **[Vulkan Best Practice for Mobile Developers](https://github.com/ARM-software/vulkan_best_practice_for_mobile_developers)** from ARM. License: MIT.
-- **[RPCS3](https://github.com/RPCS3/rpcs3)** - PlayStation 3 emulator/debugger. License: GNU GPLv2.
-
-[Many other projects on GitHub](https://github.com/search?q=AMD_VULKAN_MEMORY_ALLOCATOR_H&type=Code) and some game development studios that use Vulkan in their games.
-
-# See also
-
-- **[D3D12 Memory Allocator](https://github.com/GPUOpen-LibrariesAndSDKs/D3D12MemoryAllocator)** - equivalent library for Direct3D 12. License: MIT.
-- **[Awesome Vulkan](https://github.com/vinjn/awesome-vulkan)** - a curated list of awesome Vulkan libraries, debuggers and resources.
-- **[VulkanMemoryAllocator-Hpp](https://github.com/malte-v/VulkanMemoryAllocator-Hpp)** - C++ binding for this library. License: CC0-1.0.
-- **[PyVMA](https://github.com/realitix/pyvma)** - Python wrapper for this library. Author: Jean-Sébastien B. (@realitix). License: Apache 2.0.
-- **[vk-mem](https://github.com/gwihlidal/vk-mem-rs)** - Rust binding for this library. Author: Graham Wihlidal. License: Apache 2.0 or MIT.
-- **[Haskell bindings](https://hackage.haskell.org/package/VulkanMemoryAllocator)**, **[github](https://github.com/expipiplus1/vulkan/tree/master/VulkanMemoryAllocator)** - Haskell bindings for this library. Author: Joe Hermaszewski (@expipiplus1). License BSD-3-Clause.
-- **[vma_sample_sdl](https://github.com/rextimmy/vma_sample_sdl)** - SDL port of the sample app of this library (with the goal of running it on multiple platforms, including MacOS). Author: @rextimmy. License: MIT.
-- **[vulkan-malloc](https://github.com/dylanede/vulkan-malloc)** - Vulkan memory allocation library for Rust. Based on version 1 of this library. Author: Dylan Ede (@dylanede). License: MIT / Apache 2.0.

extern/vulkan_memory_allocator/vk_mem_alloc_impl.cc

@@ -1,3 +0,0 @@
-#define VMA_IMPLEMENTATION
-
-#include "vk_mem_alloc.h"

intern/CMakeLists.txt

@@ -85,7 +85,3 @@ endif()
 if(UNIX AND NOT APPLE)
   add_subdirectory(libc_compat)
 endif()
-
-if(WITH_VULKAN)
-  add_subdirectory(shader_compiler)
-endif()

intern/cycles/CMakeLists.txt

@@ -226,9 +226,6 @@ add_definitions(
   -DCCL_NAMESPACE_END=}
 )
 
-if(WITH_CYCLES_DEBUG)
-  add_definitions(-DWITH_CYCLES_DEBUG)
-endif()
 if(WITH_CYCLES_STANDALONE_GUI)
   add_definitions(-DWITH_CYCLES_STANDALONE_GUI)
 endif()
@@ -337,7 +334,7 @@ else()
 endif()
 
 # Warnings
-if(CMAKE_COMPILER_IS_GNUCXX OR CMAKE_C_COMPILER_ID MATCHES "Clang")
+if(CMAKE_COMPILER_IS_GNUCXX)
   ADD_CHECK_CXX_COMPILER_FLAG(CMAKE_CXX_FLAGS _has_cxxflag_float_conversion "-Werror=float-conversion")
   ADD_CHECK_CXX_COMPILER_FLAG(CMAKE_CXX_FLAGS _has_cxxflag_double_promotion "-Werror=double-promotion")
   ADD_CHECK_CXX_COMPILER_FLAG(CMAKE_CXX_FLAGS _has_no_error_unused_macros "-Wno-error=unused-macros")

intern/cycles/blender/addon/properties.py

@@ -218,12 +218,6 @@ enum_denoising_prefilter = (
     ('ACCURATE', "Accurate", "Prefilter noisy guiding passes before denoising color. Improves quality when guiding passes are noisy using extra processing time", 3),
 )
 
-enum_direct_light_sampling_type = (
-    ('MULTIPLE_IMPORTANCE_SAMPLING', "Multiple Importance Sampling", "Multiple importance sampling is used to combine direct light contributions from next-event estimation and forward path tracing", 0),
-    ('FORWARD_PATH_TRACING', "Forward Path Tracing", "Direct light contributions are only sampled using forward path tracing", 1),
-    ('NEXT_EVENT_ESTIMATION', "Next-Event Estimation", "Direct light contributions are only sampled using next-event estimation", 2),
-)
-
 def update_render_passes(self, context):
     scene = context.scene
     view_layer = context.view_layer
@@ -331,13 +325,6 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         default=1024,
     )
 
-    sample_offset: IntProperty(
-        name="Sample Offset",
-        description="Number of samples to skip when starting render",
-        min=0, max=(1 << 24),
-        default=0,
-    )
-
     time_limit: FloatProperty(
         name="Time Limit",
         description="Limit the render time (excluding synchronization time)."
@@ -359,7 +346,7 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         name="Scrambling Distance",
         default=1.0,
         min=0.0, max=1.0,
-        description="Reduce randomization between pixels to improve GPU rendering performance, at the cost of possible rendering artifacts if set too low. Only works when not using adaptive sampling",
+        description="Lower values give faster rendering with GPU rendering and less noise with all devices at the cost of possible artifacts if set too low. Only works when not using adaptive sampling",
     )
     preview_scrambling_distance: BoolProperty(
         name="Scrambling Distance viewport",
@@ -367,10 +354,10 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         description="Uses the Scrambling Distance value for the viewport. Faster but may flicker",
     )
 
-    auto_scrambling_distance: BoolProperty(
-        name="Automatic Scrambling Distance",
+    adaptive_scrambling_distance: BoolProperty(
+        name="Adaptive Scrambling Distance",
         default=False,
-        description="Automatically reduce the randomization between pixels to improve GPU rendering performance, at the cost of possible rendering artifacts. Only works when not using adaptive sampling",
+        description="Uses a formula to adapt the scrambling distance strength based on the sample count",
     )
 
     use_layer_samples: EnumProperty(
@@ -428,13 +415,6 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         default=0,
     )
 
-    direct_light_sampling_type: EnumProperty(
-        name="Direct Light Sampling Type",
-        description="The type of strategy used for sampling direct light contributions",
-        items=enum_direct_light_sampling_type,
-        default='MULTIPLE_IMPORTANCE_SAMPLING',
-    )
-
     min_light_bounces: IntProperty(
         name="Min Light Bounces",
         description="Minimum number of light bounces. Setting this higher reduces noise in the first bounces, "
@@ -790,8 +770,8 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
     )
 
     use_auto_tile: BoolProperty(
-        name="Using Tiling",
-        description="Render high resolution images in tiles to reduce memory usage, using the specified tile size. Tiles are cached to disk while rendering to save memory",
+        name="Auto Tiles",
+        description="Automatically render high resolution images in tiles to reduce memory usage, using the specified tile size. Tiles are cached to disk while rendering to save memory",
         default=True,
     )
     tile_size: IntProperty(

intern/cycles/blender/addon/ui.py

@@ -290,18 +290,15 @@ class CYCLES_RENDER_PT_sampling_advanced(CyclesButtonsPanel, Panel):
         col.active = not (cscene.use_adaptive_sampling and cscene.use_preview_adaptive_sampling)
         col.prop(cscene, "sampling_pattern", text="Pattern")
 
-        col = layout.column(align=True)
-        col.prop(cscene, "sample_offset")
-
         layout.separator()
 
-        heading = layout.column(align=True, heading="Scrambling Distance")
-        heading.active = not (cscene.use_adaptive_sampling and cscene.use_preview_adaptive_sampling)
-        heading.prop(cscene, "auto_scrambling_distance", text="Automatic")
-        sub = heading.row()
+        col = layout.column(align=True)
+        col.active = not (cscene.use_adaptive_sampling and cscene.use_preview_adaptive_sampling)
+        col.prop(cscene, "scrambling_distance", text="Scrambling Distance")
+        col.prop(cscene, "adaptive_scrambling_distance", text="Adaptive")
+        sub = col.row(align=True)
         sub.active = not cscene.use_preview_adaptive_sampling
         sub.prop(cscene, "preview_scrambling_distance", text="Viewport")
-        heading.prop(cscene, "scrambling_distance", text="Multiplier")
 
         layout.separator()
 
@@ -1054,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', 'VOLUME', 'POINTCLOUD', 'HAIR'}) or
+    return ob and ((ob.type in {'MESH', 'CURVE', 'SURFACE', 'FONT', 'META', 'LIGHT'}) or
                    (ob.instance_type == 'COLLECTION' and ob.instance_collection))
 
 

intern/cycles/blender/curves.cpp

@@ -199,7 +199,7 @@ static bool ObtainCacheParticleUV(Hair *hair,
           b_mesh->uv_layers.begin(l);
 
           float2 uv = zero_float2();
-          if (!b_mesh->uv_layers.empty())
+          if (b_mesh->uv_layers.length())
             b_psys.uv_on_emitter(psmd, *b_pa, pa_no, uv_num, &uv.x);
           CData->curve_uv.push_back_slow(uv);
 
@@ -261,7 +261,7 @@ static bool ObtainCacheParticleVcol(Hair *hair,
           b_mesh->vertex_colors.begin(l);
 
           float4 vcol = make_float4(0.0f, 0.0f, 0.0f, 1.0f);
-          if (!b_mesh->vertex_colors.empty())
+          if (b_mesh->vertex_colors.length())
             b_psys.mcol_on_emitter(psmd, *b_pa, pa_no, vcol_num, &vcol.x);
           CData->curve_vcol.push_back_slow(vcol);
 

intern/cycles/blender/display_driver.cpp

@@ -334,7 +334,7 @@ bool BlenderDisplayDriver::update_begin(const Params &params,
 
   /* Update PBO dimensions if needed.
    *
-   * NOTE: Allocate the PBO for the size which will fit the final render resolution (as in,
+   * NOTE: Allocate the PBO for the the size which will fit the final render resolution (as in,
    * at a resolution divider 1. This was we don't need to recreate graphics interoperability
    * objects which are costly and which are tied to the specific underlying buffer size.
    * The downside of this approach is that when graphics interoperability is not used we are

intern/cycles/blender/mesh.cpp

@@ -555,7 +555,7 @@ static void attr_create_vertex_color(Scene *scene, Mesh *mesh, BL::Mesh &b_mesh,
 /* Create uv map attributes. */
 static void attr_create_uv_map(Scene *scene, Mesh *mesh, BL::Mesh &b_mesh)
 {
-  if (!b_mesh.uv_layers.empty()) {
+  if (b_mesh.uv_layers.length() != 0) {
     for (BL::MeshUVLoopLayer &l : b_mesh.uv_layers) {
       const bool active_render = l.active_render();
       AttributeStandard uv_std = (active_render) ? ATTR_STD_UV : ATTR_STD_NONE;
@@ -619,7 +619,7 @@ static void attr_create_uv_map(Scene *scene, Mesh *mesh, BL::Mesh &b_mesh)
 
 static void attr_create_subd_uv_map(Scene *scene, Mesh *mesh, BL::Mesh &b_mesh, bool subdivide_uvs)
 {
-  if (!b_mesh.uv_layers.empty()) {
+  if (b_mesh.uv_layers.length() != 0) {
     BL::Mesh::uv_layers_iterator l;
     int i = 0;
 
@@ -951,7 +951,7 @@ static void create_mesh(Scene *scene,
   N = attr_N->data_float3();
 
   /* create generated coordinates from undeformed coordinates */
-  const bool need_default_tangent = (subdivision == false) && (b_mesh.uv_layers.empty()) &&
+  const bool need_default_tangent = (subdivision == false) && (b_mesh.uv_layers.length() == 0) &&
                                     (mesh->need_attribute(scene, ATTR_STD_UV_TANGENT));
   if (mesh->need_attribute(scene, ATTR_STD_GENERATED) || need_default_tangent) {
     Attribute *attr = attributes.add(ATTR_STD_GENERATED);

intern/cycles/blender/object.cpp

@@ -62,15 +62,15 @@ bool BlenderSync::BKE_object_is_modified(BL::Object &b_ob)
   return false;
 }
 
-bool BlenderSync::object_is_geometry(BObjectInfo &b_ob_info)
+bool BlenderSync::object_is_geometry(BL::Object &b_ob)
 {
-  BL::ID b_ob_data = b_ob_info.object_data;
+  BL::ID b_ob_data = b_ob.data();
 
   if (!b_ob_data) {
     return false;
   }
 
-  BL::Object::type_enum type = b_ob_info.iter_object.type();
+  BL::Object::type_enum type = b_ob.type();
 
   if (type == BL::Object::type_VOLUME || type == BL::Object::type_HAIR) {
     /* Will be exported attached to mesh. */
@@ -87,24 +87,6 @@ bool BlenderSync::object_is_geometry(BObjectInfo &b_ob_info)
   return b_ob_data.is_a(&RNA_Mesh);
 }
 
-bool BlenderSync::object_can_have_geometry(BL::Object &b_ob)
-{
-  BL::Object::type_enum type = b_ob.type();
-  switch (type) {
-    case BL::Object::type_MESH:
-    case BL::Object::type_CURVE:
-    case BL::Object::type_SURFACE:
-    case BL::Object::type_META:
-    case BL::Object::type_FONT:
-    case BL::Object::type_HAIR:
-    case BL::Object::type_POINTCLOUD:
-    case BL::Object::type_VOLUME:
-      return true;
-    default:
-      return false;
-  }
-}
-
 bool BlenderSync::object_is_light(BL::Object &b_ob)
 {
   BL::ID b_ob_data = b_ob.data();
@@ -207,7 +189,7 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph,
   }
 
   /* only interested in object that we can create meshes from */
-  if (!object_is_geometry(b_ob_info)) {
+  if (!object_is_geometry(b_ob)) {
     return NULL;
   }
 
@@ -294,7 +276,7 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph,
 
   object->set_visibility(visibility);
 
-  object->set_is_shadow_catcher(b_ob.is_shadow_catcher() || b_parent.is_shadow_catcher());
+  object->set_is_shadow_catcher(b_ob.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/session.cpp

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

intern/cycles/blender/session.h

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

intern/cycles/blender/sync.cpp

@@ -162,19 +162,19 @@ void BlenderSync::sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d
     /* Object */
     else if (b_id.is_a(&RNA_Object)) {
       BL::Object b_ob(b_id);
-      const bool can_have_geometry = object_can_have_geometry(b_ob);
-      const bool is_light = !can_have_geometry && object_is_light(b_ob);
+      const bool is_geometry = object_is_geometry(b_ob);
+      const bool is_light = !is_geometry && object_is_light(b_ob);
 
       if (b_ob.is_instancer() && b_update.is_updated_shading()) {
         /* Needed for e.g. object color updates on instancer. */
         object_map.set_recalc(b_ob);
       }
 
-      if (can_have_geometry || is_light) {
+      if (is_geometry || is_light) {
         const bool updated_geometry = b_update.is_updated_geometry();
 
         /* Geometry (mesh, hair, volume). */
-        if (can_have_geometry) {
+        if (is_geometry) {
           if (b_update.is_updated_transform() || b_update.is_updated_shading()) {
             object_map.set_recalc(b_ob);
           }
@@ -365,8 +365,8 @@ void BlenderSync::sync_integrator(BL::ViewLayer &b_view_layer, bool background)
 
   int samples = get_int(cscene, "samples");
   float scrambling_distance = get_float(cscene, "scrambling_distance");
-  bool auto_scrambling_distance = get_boolean(cscene, "auto_scrambling_distance");
-  if (auto_scrambling_distance) {
+  bool adaptive_scrambling_distance = get_boolean(cscene, "adaptive_scrambling_distance");
+  if (adaptive_scrambling_distance) {
     scrambling_distance *= 4.0f / sqrtf(samples);
   }
 
@@ -392,12 +392,6 @@ void BlenderSync::sync_integrator(BL::ViewLayer &b_view_layer, bool background)
     integrator->set_ao_bounces(0);
   }
 
-#ifdef WITH_CYCLES_DEBUG
-  DirectLightSamplingType direct_light_sampling_type = (DirectLightSamplingType)get_enum(
-      cscene, "direct_light_sampling_type", DIRECT_LIGHT_SAMPLING_NUM, DIRECT_LIGHT_SAMPLING_MIS);
-  integrator->set_direct_light_sampling_type(direct_light_sampling_type);
-#endif
-
   const DenoiseParams denoise_params = get_denoise_params(b_scene, b_view_layer, background);
   integrator->set_use_denoise(denoise_params.use);
 
@@ -841,25 +835,18 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine &b_engine,
   /* samples */
   int samples = get_int(cscene, "samples");
   int preview_samples = get_int(cscene, "preview_samples");
-  int sample_offset = get_int(cscene, "sample_offset");
 
   if (background) {
     params.samples = samples;
-    params.sample_offset = sample_offset;
   }
   else {
     params.samples = preview_samples;
-    if (params.samples == 0) {
+    if (params.samples == 0)
       params.samples = INT_MAX;
-    }
-    params.sample_offset = 0;
   }
 
-  /* Clamp sample offset. */
-  params.sample_offset = clamp(params.sample_offset, 0, Integrator::MAX_SAMPLES);
-
   /* Clamp samples. */
-  params.samples = clamp(params.samples, 0, Integrator::MAX_SAMPLES - params.sample_offset);
+  params.samples = min(params.samples, Integrator::MAX_SAMPLES);
 
   /* Viewport Performance */
   params.pixel_size = b_engine.get_preview_pixel_size(b_scene);
@@ -878,7 +865,7 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine &b_engine,
 
   /* Time limit. */
   if (background) {
-    params.time_limit = (double)get_float(cscene, "time_limit");
+    params.time_limit = get_float(cscene, "time_limit");
   }
   else {
     /* For the viewport it kind of makes more sense to think in terms of the noise floor, which is

intern/cycles/blender/sync.h

@@ -208,8 +208,7 @@ class BlenderSync {
   /* util */
   void find_shader(BL::ID &id, array<Node *> &used_shaders, Shader *default_shader);
   bool BKE_object_is_modified(BL::Object &b_ob);
-  bool object_is_geometry(BObjectInfo &b_ob_info);
-  bool object_can_have_geometry(BL::Object &b_ob);
+  bool object_is_geometry(BL::Object &b_ob);
   bool object_is_light(BL::Object &b_ob);
 
   /* variables */

intern/cycles/blender/util.h

@@ -303,7 +303,7 @@ static inline string image_user_file_path(BL::ImageUser &iuser,
   string filepath_str = string(filepath);
   if (load_tiled && ima.source() == BL::Image::source_TILED) {
     string udim;
-    if (!ima.tiles.empty()) {
+    if (ima.tiles.length() > 0) {
       udim = to_string(ima.tiles[0].number());
     }
     string_replace(filepath_str, udim, "<UDIM>");
@@ -647,7 +647,7 @@ static inline Mesh::SubdivisionType object_subdivision_type(BL::Object &b_ob,
 {
   PointerRNA cobj = RNA_pointer_get(&b_ob.ptr, "cycles");
 
-  if (cobj.data && !b_ob.modifiers.empty() && experimental) {
+  if (cobj.data && b_ob.modifiers.length() > 0 && experimental) {
     BL::Modifier mod = b_ob.modifiers[b_ob.modifiers.length() - 1];
     bool enabled = preview ? mod.show_viewport() : mod.show_render();
 

intern/cycles/bvh/embree.cpp

@@ -303,7 +303,7 @@ static void rtc_error_func(void *, enum RTCError, const char *str)
   VLOG(1) << str;
 }
 
-static double progress_start_time = 0.0;
+static double progress_start_time = 0.0f;
 
 static bool rtc_progress_func(void *user_ptr, const double n)
 {

intern/cycles/bvh/node.cpp

@@ -153,7 +153,7 @@ void BVHNode::update_time()
 namespace {
 
 struct DumpTraversalContext {
-  /* Descriptor of while where writing is happening. */
+  /* Descriptor of wile where writing is happening. */
   FILE *stream;
   /* Unique identifier of the node current. */
   int id;

intern/cycles/bvh/node.h

@@ -178,7 +178,7 @@ class InnerNode : public BVHNode {
     reset_unused_children();
   }
 
-  /* NOTE: This function is only used during binary BVH builder, and it's
+  /* NOTE: This function is only used during binary BVH builder, and it
    * supposed to be configured to have 2 children which will be filled-in in a
    * bit. But this is important to have children reset to NULL. */
   explicit InnerNode(const BoundBox &bounds) : BVHNode(bounds), num_children_(0)

intern/cycles/bvh/optix.cpp

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

intern/cycles/bvh/optix.h

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

intern/cycles/cmake/macros.cmake

@@ -88,7 +88,7 @@ endmacro()
 
 function(cycles_link_directories)
   if(APPLE)
-    # APPLE platform uses full paths for linking libraries, and avoids link_directories.
+    # APPLE plaform uses full paths for linking libraries, and avoids link_directories.
     return()
   endif()
 

intern/cycles/device/cpu/device_impl.cpp

@@ -93,6 +93,11 @@ CPUDevice::~CPUDevice()
   texture_info.free();
 }
 
+bool CPUDevice::show_samples() const
+{
+  return (info.cpu_threads == 1);
+}
+
 BVHLayoutMask CPUDevice::get_bvh_layout_mask() const
 {
   BVHLayoutMask bvh_layout_mask = BVH_LAYOUT_BVH2;

intern/cycles/device/cpu/device_impl.h

@@ -60,6 +60,8 @@ class CPUDevice : public Device {
   CPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_);
   ~CPUDevice();
 
+  virtual bool show_samples() const override;
+
   virtual BVHLayoutMask get_bvh_layout_mask() const override;
 
   /* Returns true if the texture info was copied to the device (meaning, some more

intern/cycles/device/cpu/kernel.h

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

intern/cycles/device/cuda/device_impl.cpp

@@ -46,6 +46,12 @@ bool CUDADevice::have_precompiled_kernels()
   return path_exists(cubins_path);
 }
 
+bool CUDADevice::show_samples() const
+{
+  /* The CUDADevice only processes one tile at a time, so showing samples is fine. */
+  return true;
+}
+
 BVHLayoutMask CUDADevice::get_bvh_layout_mask() const
 {
   return BVH_LAYOUT_BVH2;
@@ -236,10 +242,6 @@ string CUDADevice::compile_kernel_get_common_cflags(const uint kernel_features)
   cflags += " -DWITH_NANOVDB";
 #  endif
 
-#  ifdef WITH_CYCLES_DEBUG
-  cflags += " -DWITH_CYCLES_DEBUG";
-#  endif
-
   return cflags;
 }
 
@@ -775,7 +777,6 @@ void CUDADevice::generic_free(device_memory &mem)
   if (mem.device_pointer) {
     CUDAContextScope scope(this);
     thread_scoped_lock lock(cuda_mem_map_mutex);
-    DCHECK(cuda_mem_map.find(&mem) != cuda_mem_map.end());
     const CUDAMem &cmem = cuda_mem_map[&mem];
 
     /* If cmem.use_mapped_host is true, reference counting is used
@@ -930,6 +931,7 @@ void CUDADevice::tex_alloc(device_texture &mem)
 {
   CUDAContextScope scope(this);
 
+  /* General variables for both architectures */
   string bind_name = mem.name;
   size_t dsize = datatype_size(mem.data_type);
   size_t size = mem.memory_size();
@@ -1092,6 +1094,7 @@ void CUDADevice::tex_alloc(device_texture &mem)
 
   if (mem.info.data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT &&
       mem.info.data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT3) {
+    /* Kepler+, bindless textures. */
     CUDA_RESOURCE_DESC resDesc;
     memset(&resDesc, 0, sizeof(resDesc));
 
@@ -1142,7 +1145,6 @@ void CUDADevice::tex_free(device_texture &mem)
   if (mem.device_pointer) {
     CUDAContextScope scope(this);
     thread_scoped_lock lock(cuda_mem_map_mutex);
-    DCHECK(cuda_mem_map.find(&mem) != cuda_mem_map.end());
     const CUDAMem &cmem = cuda_mem_map[&mem];
 
     if (cmem.texobject) {

intern/cycles/device/cuda/device_impl.h

@@ -76,6 +76,8 @@ class CUDADevice : public Device {
 
   static bool have_precompiled_kernels();
 
+  virtual bool show_samples() const override;
+
   virtual BVHLayoutMask get_bvh_layout_mask() const override;
 
   void set_error(const string &error) override;

intern/cycles/device/device.h

@@ -149,6 +149,10 @@ class Device {
     fprintf(stderr, "%s\n", error.c_str());
     fflush(stderr);
   }
+  virtual bool show_samples() const
+  {
+    return false;
+  }
   virtual BVHLayoutMask get_bvh_layout_mask() const = 0;
 
   /* statistics */

intern/cycles/device/hip/device_impl.cpp

@@ -47,6 +47,12 @@ bool HIPDevice::have_precompiled_kernels()
   return path_exists(fatbins_path);
 }
 
+bool HIPDevice::show_samples() const
+{
+  /* The HIPDevice only processes one tile at a time, so showing samples is fine. */
+  return true;
+}
+
 BVHLayoutMask HIPDevice::get_bvh_layout_mask() const
 {
   return BVH_LAYOUT_BVH2;
@@ -93,7 +99,7 @@ HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
   }
 
   /* Setup device and context. */
-  result = hipDeviceGet(&hipDevice, hipDevId);
+  result = hipGetDevice(&hipDevice, hipDevId);
   if (result != hipSuccess) {
     set_error(string_printf("Failed to get HIP device handle from ordinal (%s)",
                             hipewErrorString(result)));
@@ -216,6 +222,7 @@ 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\"",
@@ -227,7 +234,10 @@ 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)
+string HIPDevice::compile_kernel(const uint kernel_features,
+                                 const char *name,
+                                 const char *base,
+                                 bool force_ptx)
 {
   /* Compute kernel name. */
   int major, minor;
@@ -237,7 +247,7 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
   hipGetDeviceProperties(&props, hipDevId);
 
   /* gcnArchName can contain tokens after the arch name with features, ie.
-   * `gfx1010:sramecc-:xnack-` so we tokenize it to get the first part. */
+    "gfx1010:sramecc-:xnack-" so we tokenize it to get the first part. */
   char *arch = strtok(props.gcnArchName, ":");
   if (arch == NULL) {
     arch = props.gcnArchName;
@@ -245,11 +255,13 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
 
   /* Attempt to use kernel provided with Blender. */
   if (!use_adaptive_compilation()) {
-    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;
+    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;
+      }
     }
   }
 
@@ -286,9 +298,9 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
 
 #  ifdef _WIN32
   if (!use_adaptive_compilation() && have_precompiled_kernels()) {
-    if (!hipSupportsDevice(hipDevId)) {
+    if (major < 3) {
       set_error(
-          string_printf("HIP backend requires compute capability 10.1 or up, but found %d.%d. "
+          string_printf("HIP backend requires compute capability 3.0 or up, but found %d.%d. "
                         "Your GPU is not supported.",
                         major,
                         minor));
@@ -368,9 +380,10 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
 
 bool HIPDevice::load_kernels(const uint kernel_features)
 {
-  /* TODO(sergey): Support kernels re-load for HIP devices adaptive compile.
+  /* TODO(sergey): Support kernels re-load for CUDA devices adaptive compile.
    *
-   * Currently re-loading kernels will invalidate memory pointers.
+   * Currently re-loading kernel will invalidate memory pointers,
+   * causing problems in cuCtxSynchronize.
    */
   if (hipModule) {
     if (use_adaptive_compilation()) {
@@ -738,7 +751,6 @@ void HIPDevice::generic_free(device_memory &mem)
   if (mem.device_pointer) {
     HIPContextScope scope(this);
     thread_scoped_lock lock(hip_mem_map_mutex);
-    DCHECK(hip_mem_map.find(&mem) != hip_mem_map.end());
     const HIPMem &cmem = hip_mem_map[&mem];
 
     /* If cmem.use_mapped_host is true, reference counting is used
@@ -892,6 +904,7 @@ void HIPDevice::tex_alloc(device_texture &mem)
 {
   HIPContextScope scope(this);
 
+  /* General variables for both architectures */
   string bind_name = mem.name;
   size_t dsize = datatype_size(mem.data_type);
   size_t size = mem.memory_size();
@@ -981,16 +994,16 @@ void HIPDevice::tex_alloc(device_texture &mem)
             << string_human_readable_number(mem.memory_size()) << " bytes. ("
             << string_human_readable_size(mem.memory_size()) << ")";
 
-    hip_assert(hipArray3DCreate((hArray *)&array_3d, &desc));
+    hip_assert(hipArray3DCreate(&array_3d, &desc));
 
     if (!array_3d) {
       return;
     }
 
     HIP_MEMCPY3D param;
-    memset(&param, 0, sizeof(HIP_MEMCPY3D));
+    memset(&param, 0, sizeof(param));
     param.dstMemoryType = hipMemoryTypeArray;
-    param.dstArray = array_3d;
+    param.dstArray = &array_3d;
     param.srcMemoryType = hipMemoryTypeHost;
     param.srcHost = mem.host_pointer;
     param.srcPitch = src_pitch;
@@ -1056,13 +1069,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) {
-    /* Bindless textures. */
+    /* Kepler+, 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) {
@@ -1107,7 +1120,6 @@ void HIPDevice::tex_free(device_texture &mem)
   if (mem.device_pointer) {
     HIPContextScope scope(this);
     thread_scoped_lock lock(hip_mem_map_mutex);
-    DCHECK(hip_mem_map.find(&mem) != hip_mem_map.end());
     const HIPMem &cmem = hip_mem_map[&mem];
 
     if (cmem.texobject) {
@@ -1148,8 +1160,6 @@ 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;
@@ -1168,7 +1178,6 @@ bool HIPDevice::should_use_graphics_interop()
       return true;
     }
   }
-#  endif
 
   return false;
 }

intern/cycles/device/hip/device_impl.h

@@ -75,6 +75,8 @@ class HIPDevice : public Device {
 
   static bool have_precompiled_kernels();
 
+  virtual bool show_samples() const override;
+
   virtual BVHLayoutMask get_bvh_layout_mask() const override;
 
   void set_error(const string &error) override;
@@ -91,7 +93,10 @@ class HIPDevice : public Device {
 
   virtual string compile_kernel_get_common_cflags(const uint kernel_features);
 
-  string compile_kernel(const uint kernel_features, const char *name, const char *base = "hip");
+  string compile_kernel(const uint kernel_features,
+                        const char *name,
+                        const char *base = "hip",
+                        bool force_ptx = false);
 
   virtual bool load_kernels(const uint kernel_features) override;
   void reserve_local_memory(const uint kernel_features);

intern/cycles/device/hip/graphics_interop.h

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

intern/cycles/device/memory.cpp

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

intern/cycles/device/memory.h

@@ -81,140 +81,155 @@ static constexpr size_t datatype_size(DataType datatype)
 
 template<typename T> struct device_type_traits {
   static const DataType data_type = TYPE_UNKNOWN;
-  static const size_t num_elements = sizeof(T);
+  static const size_t num_elements_cpu = sizeof(T);
+  static const size_t num_elements_gpu = sizeof(T);
 };
 
 template<> struct device_type_traits<uchar> {
   static const DataType data_type = TYPE_UCHAR;
-  static const size_t num_elements = 1;
-  static_assert(sizeof(uchar) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 1;
+  static const size_t num_elements_gpu = 1;
+  static_assert(sizeof(uchar) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<uchar2> {
   static const DataType data_type = TYPE_UCHAR;
-  static const size_t num_elements = 2;
-  static_assert(sizeof(uchar2) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 2;
+  static const size_t num_elements_gpu = 2;
+  static_assert(sizeof(uchar2) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<uchar3> {
   static const DataType data_type = TYPE_UCHAR;
-  static const size_t num_elements = 3;
-  static_assert(sizeof(uchar3) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 3;
+  static const size_t num_elements_gpu = 3;
+  static_assert(sizeof(uchar3) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<uchar4> {
   static const DataType data_type = TYPE_UCHAR;
-  static const size_t num_elements = 4;
-  static_assert(sizeof(uchar4) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 4;
+  static const size_t num_elements_gpu = 4;
+  static_assert(sizeof(uchar4) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<uint> {
   static const DataType data_type = TYPE_UINT;
-  static const size_t num_elements = 1;
-  static_assert(sizeof(uint) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 1;
+  static const size_t num_elements_gpu = 1;
+  static_assert(sizeof(uint) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<uint2> {
   static const DataType data_type = TYPE_UINT;
-  static const size_t num_elements = 2;
-  static_assert(sizeof(uint2) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 2;
+  static const size_t num_elements_gpu = 2;
+  static_assert(sizeof(uint2) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<uint3> {
   static const DataType data_type = TYPE_UINT;
-  static const size_t num_elements = 3;
-  static_assert(sizeof(uint3) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 3;
+  static const size_t num_elements_gpu = 3;
+  static_assert(sizeof(uint3) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<uint4> {
   static const DataType data_type = TYPE_UINT;
-  static const size_t num_elements = 4;
-  static_assert(sizeof(uint4) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 4;
+  static const size_t num_elements_gpu = 4;
+  static_assert(sizeof(uint4) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<int> {
   static const DataType data_type = TYPE_INT;
-  static const size_t num_elements = 1;
-  static_assert(sizeof(int) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 1;
+  static const size_t num_elements_gpu = 1;
+  static_assert(sizeof(int) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<int2> {
   static const DataType data_type = TYPE_INT;
-  static const size_t num_elements = 2;
-  static_assert(sizeof(int2) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 2;
+  static const size_t num_elements_gpu = 2;
+  static_assert(sizeof(int2) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<int3> {
   static const DataType data_type = TYPE_INT;
-  static const size_t num_elements = 4;
-  static_assert(sizeof(int3) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 4;
+  static const size_t num_elements_gpu = 3;
+  static_assert(sizeof(int3) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<int4> {
   static const DataType data_type = TYPE_INT;
-  static const size_t num_elements = 4;
-  static_assert(sizeof(int4) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 4;
+  static const size_t num_elements_gpu = 4;
+  static_assert(sizeof(int4) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<float> {
   static const DataType data_type = TYPE_FLOAT;
-  static const size_t num_elements = 1;
-  static_assert(sizeof(float) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 1;
+  static const size_t num_elements_gpu = 1;
+  static_assert(sizeof(float) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<float2> {
   static const DataType data_type = TYPE_FLOAT;
-  static const size_t num_elements = 2;
-  static_assert(sizeof(float2) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 2;
+  static const size_t num_elements_gpu = 2;
+  static_assert(sizeof(float2) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<float3> {
-  /* float3 has different size depending on the device, can't use it for interchanging
-   * memory between CPU and GPU.
-   *
-   * Leave body empty to trigger a compile error if used. */
-};
-
-template<> struct device_type_traits<packed_float3> {
   static const DataType data_type = TYPE_FLOAT;
-  static const size_t num_elements = 3;
-  static_assert(sizeof(packed_float3) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 4;
+  static const size_t num_elements_gpu = 3;
+  static_assert(sizeof(float3) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<float4> {
   static const DataType data_type = TYPE_FLOAT;
-  static const size_t num_elements = 4;
-  static_assert(sizeof(float4) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 4;
+  static const size_t num_elements_gpu = 4;
+  static_assert(sizeof(float4) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<half> {
   static const DataType data_type = TYPE_HALF;
-  static const size_t num_elements = 1;
-  static_assert(sizeof(half) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 1;
+  static const size_t num_elements_gpu = 1;
+  static_assert(sizeof(half) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<ushort4> {
   static const DataType data_type = TYPE_UINT16;
-  static const size_t num_elements = 4;
-  static_assert(sizeof(ushort4) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 4;
+  static const size_t num_elements_gpu = 4;
+  static_assert(sizeof(ushort4) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<uint16_t> {
   static const DataType data_type = TYPE_UINT16;
-  static const size_t num_elements = 1;
-  static_assert(sizeof(uint16_t) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 1;
+  static const size_t num_elements_gpu = 1;
+  static_assert(sizeof(uint16_t) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<half4> {
   static const DataType data_type = TYPE_HALF;
-  static const size_t num_elements = 4;
-  static_assert(sizeof(half4) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 4;
+  static const size_t num_elements_gpu = 4;
+  static_assert(sizeof(half4) == num_elements_cpu * datatype_size(data_type));
 };
 
 template<> struct device_type_traits<uint64_t> {
   static const DataType data_type = TYPE_UINT64;
-  static const size_t num_elements = 1;
-  static_assert(sizeof(uint64_t) == num_elements * datatype_size(data_type));
+  static const size_t num_elements_cpu = 1;
+  static const size_t num_elements_gpu = 1;
+  static_assert(sizeof(uint64_t) == num_elements_cpu * datatype_size(data_type));
 };
 
 /* Device Memory
@@ -266,16 +281,11 @@ class device_memory {
 
   /* Only create through subclasses. */
   device_memory(Device *device, const char *name, MemoryType type);
+  device_memory(device_memory &&other) noexcept;
 
-  /* No copying and allowed.
-   *
-   * This is because device implementation might need to register device memory in an allocation
-   * map of some sort and use pointer as a key to identify blocks. Moving data from one place to
-   * another bypassing device allocation routines will make those maps hard to maintain. */
+  /* No copying allowed. */
   device_memory(const device_memory &) = delete;
-  device_memory(device_memory &&other) noexcept = delete;
   device_memory &operator=(const device_memory &) = delete;
-  device_memory &operator=(device_memory &&) = delete;
 
   /* Host allocation on the device. All host_pointer memory should be
    * allocated with these functions, for devices that support using
@@ -310,7 +320,9 @@ template<typename T> class device_only_memory : public device_memory {
       : device_memory(device, name, allow_host_memory_fallback ? MEM_READ_WRITE : MEM_DEVICE_ONLY)
   {
     data_type = device_type_traits<T>::data_type;
-    data_elements = max(device_type_traits<T>::num_elements, 1);
+    data_elements = max(device_is_cpu() ? device_type_traits<T>::num_elements_cpu :
+                                          device_type_traits<T>::num_elements_gpu,
+                        1);
   }
 
   device_only_memory(device_only_memory &&other) noexcept : device_memory(std::move(other))
@@ -366,11 +378,15 @@ template<typename T> class device_only_memory : public device_memory {
 
 template<typename T> class device_vector : public device_memory {
  public:
+  /* Can only use this for types that have the same size on CPU and GPU. */
+  static_assert(device_type_traits<T>::num_elements_cpu ==
+                device_type_traits<T>::num_elements_gpu);
+
   device_vector(Device *device, const char *name, MemoryType type)
       : device_memory(device, name, type)
   {
     data_type = device_type_traits<T>::data_type;
-    data_elements = device_type_traits<T>::num_elements;
+    data_elements = device_type_traits<T>::num_elements_cpu;
     modified = true;
     need_realloc_ = true;
 

intern/cycles/device/multi/device.cpp

@@ -109,6 +109,14 @@ class MultiDevice : public Device {
     return error_msg;
   }
 
+  virtual bool show_samples() const override
+  {
+    if (devices.size() > 1) {
+      return false;
+    }
+    return devices.front().device->show_samples();
+  }
+
   virtual BVHLayoutMask get_bvh_layout_mask() const override
   {
     BVHLayoutMask bvh_layout_mask = BVH_LAYOUT_ALL;

intern/cycles/device/optix/device_impl.cpp

@@ -886,7 +886,8 @@ 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);
+  denoiser_.state.alloc_to_device(denoiser_.scratch_offset + denoiser_.scratch_size +
+                                  sizeof(float));
 
   /* Initialize denoiser state for the current tile size. */
   const OptixResult result = optixDenoiserSetup(
@@ -970,6 +971,16 @@ 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;
@@ -1032,7 +1043,7 @@ bool OptiXDevice::build_optix_bvh(BVHOptiX *bvh,
     return false;
   }
 
-  device_only_memory<char> &out_data = *bvh->as_data;
+  device_only_memory<char> &out_data = bvh->as_data;
   if (operation == OPTIX_BUILD_OPERATION_BUILD) {
     assert(out_data.device == this);
     out_data.alloc_to_device(sizes.outputSizeInBytes);
@@ -1123,7 +1134,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;
     }
 
@@ -1344,9 +1355,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,
@@ -1379,8 +1390,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) {
@@ -1441,7 +1452,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
 
         motion_transform_offset = align_up(motion_transform_offset,
                                            OPTIX_TRANSFORM_BYTE_ALIGNMENT);
-        CUdeviceptr motion_transform_gpu = bvh_optix->motion_transform_data->device_pointer +
+        CUdeviceptr motion_transform_gpu = bvh_optix->motion_transform_data.device_pointer +
                                            motion_transform_offset;
         motion_transform_offset += motion_transform_size;
 

intern/cycles/device/optix/device_impl.h

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

intern/cycles/device/optix/queue.cpp

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

intern/cycles/doc/license/readme.txt

@@ -3,7 +3,7 @@ This program uses code from various sources, the default license is Apache 2.0
 for all code, with the following exceptions.
 
 Modified BSD License
-* Code adapted from Open Shading Language
+* Code adapated from Open Shading Language
 * Sobol direction vectors
 * Matrix inversion code from OpenEXR
 * MD5 Hash code

intern/cycles/integrator/denoiser.cpp

@@ -33,10 +33,7 @@ unique_ptr<Denoiser> Denoiser::create(Device *path_trace_device, const DenoisePa
     return make_unique<OptiXDenoiser>(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);
+  return make_unique<OIDNDenoiser>(path_trace_device, params);
 }
 
 Denoiser::Denoiser(Device *path_trace_device, const DenoiseParams &params)

intern/cycles/integrator/path_trace.cpp

@@ -296,13 +296,13 @@ static BufferParams scale_buffer_params(const BufferParams &params, int resoluti
 
   scaled_params.window_x = params.window_x / resolution_divider;
   scaled_params.window_y = params.window_y / resolution_divider;
-  scaled_params.window_width = max(1, params.window_width / resolution_divider);
-  scaled_params.window_height = max(1, params.window_height / resolution_divider);
+  scaled_params.window_width = params.window_width / resolution_divider;
+  scaled_params.window_height = params.window_height / resolution_divider;
 
   scaled_params.full_x = params.full_x / resolution_divider;
   scaled_params.full_y = params.full_y / resolution_divider;
-  scaled_params.full_width = max(1, params.full_width / resolution_divider);
-  scaled_params.full_height = max(1, params.full_height / resolution_divider);
+  scaled_params.full_width = params.full_width / resolution_divider;
+  scaled_params.full_height = params.full_height / resolution_divider;
 
   scaled_params.update_offset_stride();
 
@@ -380,10 +380,7 @@ void PathTrace::path_trace(RenderWork &render_work)
     PathTraceWork *path_trace_work = path_trace_works_[i].get();
 
     PathTraceWork::RenderStatistics statistics;
-    path_trace_work->render_samples(statistics,
-                                    render_work.path_trace.start_sample,
-                                    num_samples,
-                                    render_work.path_trace.sample_offset);
+    path_trace_work->render_samples(statistics, render_work.path_trace.start_sample, num_samples);
 
     const double work_time = time_dt() - work_start_time;
     work_balance_infos_[i].time_spent += work_time;
@@ -850,11 +847,9 @@ void PathTrace::progress_update_if_needed(const RenderWork &render_work)
 {
   if (progress_ != nullptr) {
     const int2 tile_size = get_render_tile_size();
-    const uint64_t num_samples_added = uint64_t(tile_size.x) * tile_size.y *
-                                       render_work.path_trace.num_samples;
+    const int num_samples_added = tile_size.x * tile_size.y * render_work.path_trace.num_samples;
     const int current_sample = render_work.path_trace.start_sample +
-                               render_work.path_trace.num_samples -
-                               render_work.path_trace.sample_offset;
+                               render_work.path_trace.num_samples;
     progress_->add_samples(num_samples_added, current_sample);
   }
 

intern/cycles/integrator/path_trace_display.h

@@ -76,7 +76,7 @@ class PathTraceDisplay {
 
   /* Copy buffer of rendered pixels of a given size into a given position of the texture.
    *
-   * This function does not acquire a lock. The reason for this is to allow use of this function
+   * This function does not acquire a lock. The reason for this is is to allow use of this function
    * for partial updates from different devices. In this case the caller will acquire the lock
    * once, update all the slices and release
    * the lock once. This will ensure that draw() will never use partially updated texture. */

intern/cycles/integrator/path_trace_work.h

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

intern/cycles/integrator/path_trace_work_cpu.cpp

@@ -71,17 +71,14 @@ void PathTraceWorkCPU::init_execution()
 
 void PathTraceWorkCPU::render_samples(RenderStatistics &statistics,
                                       int start_sample,
-                                      int samples_num,
-                                      int sample_offset)
+                                      int samples_num)
 {
   const int64_t image_width = effective_buffer_params_.width;
   const int64_t image_height = effective_buffer_params_.height;
   const int64_t total_pixels_num = image_width * image_height;
 
-  if (device_->profiler.active()) {
-    for (CPUKernelThreadGlobals &kernel_globals : kernel_thread_globals_) {
-      kernel_globals.start_profiling();
-    }
+  for (CPUKernelThreadGlobals &kernel_globals : kernel_thread_globals_) {
+    kernel_globals.start_profiling();
   }
 
   tbb::task_arena local_arena = local_tbb_arena_create(device_);
@@ -100,7 +97,6 @@ void PathTraceWorkCPU::render_samples(RenderStatistics &statistics,
       work_tile.w = 1;
       work_tile.h = 1;
       work_tile.start_sample = start_sample;
-      work_tile.sample_offset = sample_offset;
       work_tile.num_samples = 1;
       work_tile.offset = effective_buffer_params_.offset;
       work_tile.stride = effective_buffer_params_.stride;
@@ -110,10 +106,9 @@ void PathTraceWorkCPU::render_samples(RenderStatistics &statistics,
       render_samples_full_pipeline(kernel_globals, work_tile, samples_num);
     });
   });
-  if (device_->profiler.active()) {
-    for (CPUKernelThreadGlobals &kernel_globals : kernel_thread_globals_) {
-      kernel_globals.stop_profiling();
-    }
+
+  for (CPUKernelThreadGlobals &kernel_globals : kernel_thread_globals_) {
+    kernel_globals.stop_profiling();
   }
 
   statistics.occupancy = 1.0f;

intern/cycles/integrator/path_trace_work_cpu.h

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

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -250,8 +250,7 @@ void PathTraceWorkGPU::init_execution()
 
 void PathTraceWorkGPU::render_samples(RenderStatistics &statistics,
                                       int start_sample,
-                                      int samples_num,
-                                      int sample_offset)
+                                      int samples_num)
 {
   /* Limit number of states for the tile and rely on a greedy scheduling of tiles. This allows to
    * add more work (because tiles are smaller, so there is higher chance that more paths will
@@ -262,7 +261,6 @@ void PathTraceWorkGPU::render_samples(RenderStatistics &statistics,
   work_tile_scheduler_.reset(effective_buffer_params_,
                              start_sample,
                              samples_num,
-                             sample_offset,
                              device_scene_->data.integrator.scrambling_distance);
 
   enqueue_reset();
@@ -439,15 +437,7 @@ 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: {
-      /* 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_CLOSEST:
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK: {

intern/cycles/integrator/path_trace_work_gpu.h

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

intern/cycles/integrator/render_scheduler.cpp

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

intern/cycles/integrator/render_scheduler.h

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

intern/cycles/integrator/work_tile_scheduler.cpp

@@ -36,7 +36,6 @@ void WorkTileScheduler::set_max_num_path_states(int max_num_path_states)
 void WorkTileScheduler::reset(const BufferParams &buffer_params,
                               int sample_start,
                               int samples_num,
-                              int sample_offset,
                               float scrambling_distance)
 {
   /* Image buffer parameters. */
@@ -52,7 +51,6 @@ void WorkTileScheduler::reset(const BufferParams &buffer_params,
   /* Samples parameters. */
   sample_start_ = sample_start;
   samples_num_ = samples_num;
-  sample_offset_ = sample_offset;
 
   /* Initialize new scheduling. */
   reset_scheduler_state();
@@ -113,7 +111,6 @@ bool WorkTileScheduler::get_work(KernelWorkTile *work_tile_, const int max_work_
   work_tile.h = tile_size_.height;
   work_tile.start_sample = sample_start_ + start_sample;
   work_tile.num_samples = min(tile_size_.num_samples, samples_num_ - start_sample);
-  work_tile.sample_offset = sample_offset_;
   work_tile.offset = offset_;
   work_tile.stride = stride_;
 

intern/cycles/integrator/work_tile_scheduler.h

@@ -41,7 +41,6 @@ class WorkTileScheduler {
   void reset(const BufferParams &buffer_params,
              int sample_start,
              int samples_num,
-             int sample_offset,
              float scrambling_distance);
 
   /* Get work for a device.
@@ -80,7 +79,6 @@ class WorkTileScheduler {
    * (splitting into a smaller work tiles). */
   int sample_start_ = 0;
   int samples_num_ = 0;
-  int sample_offset_ = 0;
 
   /* Tile size which be scheduled for rendering. */
   TileSize tile_size_;

intern/cycles/kernel/CMakeLists.txt

@@ -273,7 +273,6 @@ set(SRC_KERNEL_UTIL_HEADERS
 )
 
 set(SRC_KERNEL_TYPES_HEADERS
-  tables.h
   textures.h
   types.h
 )
@@ -380,6 +379,7 @@ if(WITH_CYCLES_CUDA_BINARIES)
     ${SRC_KERNEL_HEADERS}
     ${SRC_KERNEL_DEVICE_GPU_HEADERS}
     ${SRC_KERNEL_DEVICE_CUDA_HEADERS}
+    ${SRC_KERNEL_DEVICE_METAL_HEADERS}
     ${SRC_UTIL_HEADERS}
   )
   set(cuda_cubins)
@@ -411,8 +411,12 @@ if(WITH_CYCLES_CUDA_BINARIES)
       -I ${CMAKE_CURRENT_SOURCE_DIR}/..
       -I ${CMAKE_CURRENT_SOURCE_DIR}/device/cuda
       --use_fast_math
-      -o ${CMAKE_CURRENT_BINARY_DIR}/${cuda_file}
-      -Wno-deprecated-gpu-targets)
+      -o ${CMAKE_CURRENT_BINARY_DIR}/${cuda_file})
+
+    if(${experimental})
+      set(cuda_flags ${cuda_flags} -D __KERNEL_EXPERIMENTAL__)
+      set(name ${name}_experimental)
+    endif()
 
     if(WITH_NANOVDB)
       set(cuda_flags ${cuda_flags}
@@ -420,10 +424,6 @@ if(WITH_CYCLES_CUDA_BINARIES)
         -I "${NANOVDB_INCLUDE_DIR}")
     endif()
 
-    if(WITH_CYCLES_DEBUG)
-      set(cuda_flags ${cuda_flags} -D WITH_CYCLES_DEBUG)
-    endif()
-
     if(WITH_CYCLES_CUBIN_COMPILER)
       string(SUBSTRING ${arch} 3 -1 CUDA_ARCH)
 
@@ -572,14 +572,13 @@ if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
       -ffast-math
       -o ${CMAKE_CURRENT_BINARY_DIR}/${hip_file})
 
-    if(WITH_NANOVDB)
-      set(hip_flags ${hip_flags}
-        -D WITH_NANOVDB
-        -I "${NANOVDB_INCLUDE_DIR}")
+    if(${experimental})
+      set(hip_flags ${hip_flags} -D __KERNEL_EXPERIMENTAL__)
+      set(name ${name}_experimental)
     endif()
 
     if(WITH_CYCLES_DEBUG)
-      set(hip_flags ${hip_flags} -D WITH_CYCLES_DEBUG)
+      set(hip_flags ${hip_flags} -D __KERNEL_DEBUG__)
     endif()
 
     add_custom_command(
@@ -620,10 +619,6 @@ if(WITH_CYCLES_DEVICE_OPTIX AND WITH_CYCLES_CUDA_BINARIES)
         -I "${NANOVDB_INCLUDE_DIR}")
     endif()
 
-    if(WITH_CYCLES_DEBUG)
-      set(cuda_flags ${cuda_flags} -D WITH_CYCLES_DEBUG)
-    endif()
-
     if(WITH_CYCLES_CUBIN_COMPILER)
       # Needed to find libnvrtc-builtins.so. Can't do it from inside
       # cycles_cubin_cc since the env variable is read before main()
@@ -712,7 +707,7 @@ if(WITH_COMPILER_ASAN)
     string(APPEND CMAKE_CXX_FLAGS_RELWITHDEBINFO " -fno-sanitize=all")
     string(APPEND CMAKE_CXX_FLAGS_DEBUG " -fno-sanitize=vptr")
   elseif(CMAKE_C_COMPILER_ID MATCHES "Clang")
-    # With OSL, Cycles disables rtti in some modules, which then breaks at linking
+    # With OSL, Cycles disables rtti in some modules, wich then breaks at linking
     # when trying to use vptr sanitizer (included into 'undefined' general option).
     string(APPEND CMAKE_CXX_FLAGS_RELWITHDEBINFO " -fno-sanitize=vptr")
     string(APPEND CMAKE_CXX_FLAGS_DEBUG " -fno-sanitize=vptr")

intern/cycles/kernel/bvh/util.h

@@ -97,7 +97,7 @@ ccl_device_inline void sort_intersections_and_normals(ccl_private Intersection *
     swapped = false;
     for (int j = 0; j < num_hits - 1; ++j) {
       if (hits[j].t > hits[j + 1].t) {
-        Intersection tmp_hit = hits[j];
+        struct Intersection tmp_hit = hits[j];
         float3 tmp_Ng = Ng[j];
         hits[j] = hits[j + 1];
         Ng[j] = Ng[j + 1];

intern/cycles/kernel/closure/bsdf.h

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

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

@@ -18,7 +18,6 @@
 
 #pragma once
 
-#include "kernel/tables.h"
 #include "kernel/types.h"
 #include "kernel/util/profiling.h"
 

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_SHADE_FUNCTION(intersect_closest);
+KERNEL_INTEGRATOR_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_SHADE_KERNEL(intersect_closest)
+DEFINE_INTEGRATOR_KERNEL(intersect_closest)
 DEFINE_INTEGRATOR_KERNEL(intersect_subsurface)
 DEFINE_INTEGRATOR_KERNEL(intersect_volume_stack)
 DEFINE_INTEGRATOR_SHADE_KERNEL(shade_background)

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

@@ -52,9 +52,8 @@ typedef unsigned long long uint64_t;
 #endif
 #define ccl_device_noinline __device__ __noinline__
 #define ccl_device_noinline_cpu ccl_device
-#define ccl_device_inline_method ccl_device
 #define ccl_global
-#define ccl_inline_constant __constant__
+#define ccl_static_constant __constant__
 #define ccl_device_constant __constant__ __device__
 #define ccl_constant const
 #define ccl_gpu_shared __shared__
@@ -86,6 +85,7 @@ typedef unsigned long long uint64_t;
 #define ccl_gpu_syncthreads() __syncthreads()
 #define ccl_gpu_ballot(predicate) __ballot_sync(0xFFFFFFFF, predicate)
 #define ccl_gpu_shfl_down_sync(mask, var, detla) __shfl_down_sync(mask, var, detla)
+#define ccl_gpu_popc(x) __popc(x)
 
 /* GPU texture objects */
 

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

@@ -92,19 +92,25 @@
 
 /* Compute number of threads per block and minimum blocks per multiprocessor
  * given the maximum number of registers per thread. */
-#define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
+
+#define ccl_gpu_kernel_threads(block_num_threads) \
+  extern "C" __global__ void __launch_bounds__(block_num_threads)
+
+#define ccl_gpu_kernel_threads_registers(block_num_threads, thread_num_registers) \
   extern "C" __global__ void __launch_bounds__(block_num_threads, \
                                                GPU_MULTIPRESSOR_MAX_REGISTERS / \
                                                    (block_num_threads * thread_num_registers))
 
-#define ccl_gpu_kernel_threads(block_num_threads) \
-  extern "C" __global__ void __launch_bounds__(block_num_threads)
+/* allow ccl_gpu_kernel to accept 1 or 2 parameters */
+#define SELECT_MACRO(_1, _2, NAME, ...) NAME
+#define ccl_gpu_kernel(...) \
+  SELECT_MACRO(__VA_ARGS__, ccl_gpu_kernel_threads_registers, ccl_gpu_kernel_threads)(__VA_ARGS__)
 
 #define ccl_gpu_kernel_signature(name, ...) kernel_gpu_##name(__VA_ARGS__)
 
 #define ccl_gpu_kernel_call(x) x
 
-/* Define a function object where "func" is the lambda body, and additional parameters are used to
+/* define a function object where "func" is the lambda body, and additional parameters are used to
  * specify captured state  */
 #define ccl_gpu_kernel_lambda(func, ...) \
   struct KernelLambda { \
@@ -113,7 +119,8 @@
     { \
       return (func); \
     } \
-  } ccl_gpu_kernel_lambda_pass
+  } ccl_gpu_kernel_lambda_pass; \
+  ccl_gpu_kernel_lambda_pass
 
 /* sanity checks */
 

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

@@ -21,9 +21,6 @@
 #include "kernel/device/gpu/parallel_sorted_index.h"
 #include "kernel/device/gpu/work_stealing.h"
 
-/* Include constant tables before entering Metal's context class scope (context_begin.h) */
-#include "kernel/tables.h"
-
 #ifdef __KERNEL_METAL__
 #  include "kernel/device/metal/context_begin.h"
 #endif
@@ -59,7 +56,8 @@
  */
 
 ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
-    ccl_gpu_kernel_signature(integrator_reset, int num_states)
+ccl_gpu_kernel_signature(integrator_reset,
+                          int num_states)
 {
   const int state = ccl_gpu_global_id_x();
 
@@ -134,14 +132,13 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
 ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
     ccl_gpu_kernel_signature(integrator_intersect_closest,
                              ccl_global 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;
-    ccl_gpu_kernel_call(integrator_intersect_closest(NULL, state, render_buffer));
+    ccl_gpu_kernel_call(integrator_intersect_closest(NULL, state));
   }
 }
 
@@ -268,7 +265,7 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
   }
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
+ccl_gpu_kernel(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_queued_paths_array,
                              int num_states,
                              ccl_global int *indices,
@@ -276,14 +273,14 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
                              int kernel_index)
 {
   ccl_gpu_kernel_lambda(INTEGRATOR_STATE(state, path, queued_kernel) == kernel_index,
-                        int kernel_index);
-  ccl_gpu_kernel_lambda_pass.kernel_index = kernel_index;
+                        int kernel_index)
+      .kernel_index = kernel_index;
 
   gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
       num_states, indices, num_indices, ccl_gpu_kernel_lambda_pass);
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
+ccl_gpu_kernel(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_queued_shadow_paths_array,
                              int num_states,
                              ccl_global int *indices,
@@ -291,26 +288,25 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
                              int kernel_index)
 {
   ccl_gpu_kernel_lambda(INTEGRATOR_STATE(state, shadow_path, queued_kernel) == kernel_index,
-                        int kernel_index);
-  ccl_gpu_kernel_lambda_pass.kernel_index = kernel_index;
+                        int kernel_index)
+      .kernel_index = kernel_index;
 
   gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
       num_states, indices, num_indices, ccl_gpu_kernel_lambda_pass);
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
+ccl_gpu_kernel(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_active_paths_array,
                              int num_states,
                              ccl_global int *indices,
                              ccl_global int *num_indices)
 {
   ccl_gpu_kernel_lambda(INTEGRATOR_STATE(state, path, queued_kernel) != 0);
-
   gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
       num_states, indices, num_indices, ccl_gpu_kernel_lambda_pass);
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
+ccl_gpu_kernel(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_terminated_paths_array,
                              int num_states,
                              ccl_global int *indices,
@@ -318,12 +314,11 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
                              int indices_offset)
 {
   ccl_gpu_kernel_lambda(INTEGRATOR_STATE(state, path, queued_kernel) == 0);
-
   gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
       num_states, indices + indices_offset, num_indices, ccl_gpu_kernel_lambda_pass);
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
+ccl_gpu_kernel(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_terminated_shadow_paths_array,
                              int num_states,
                              ccl_global int *indices,
@@ -331,12 +326,11 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
                              int indices_offset)
 {
   ccl_gpu_kernel_lambda(INTEGRATOR_STATE(state, shadow_path, queued_kernel) == 0);
-
   gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
       num_states, indices + indices_offset, num_indices, ccl_gpu_kernel_lambda_pass);
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
+ccl_gpu_kernel(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_sorted_paths_array,
                              int num_states,
                              int num_states_limit,
@@ -349,37 +343,37 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
   ccl_gpu_kernel_lambda((INTEGRATOR_STATE(state, path, queued_kernel) == kernel_index) ?
                             INTEGRATOR_STATE(state, path, shader_sort_key) :
                             GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY,
-                        int kernel_index);
-  ccl_gpu_kernel_lambda_pass.kernel_index = kernel_index;
-
+                        int kernel_index)
+    .kernel_index = kernel_index;
+  
   const uint state_index = ccl_gpu_global_id_x();
-  gpu_parallel_sorted_index_array(state_index,
-                                  num_states,
-                                  num_states_limit,
-                                  indices,
+  gpu_parallel_sorted_index_array(
+      state_index,
+      num_states,
+      num_states_limit,
+      indices,
                                   num_indices,
                                   key_counter,
                                   key_prefix_sum,
                                   ccl_gpu_kernel_lambda_pass);
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
+ccl_gpu_kernel(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_compact_paths_array,
                              int num_states,
                              ccl_global int *indices,
-                             ccl_global int *num_indices,
-                             int num_active_paths)
+                         ccl_global int *num_indices,
+                         int num_active_paths)
 {
-  ccl_gpu_kernel_lambda((state >= num_active_paths) &&
-                            (INTEGRATOR_STATE(state, path, queued_kernel) != 0),
-                        int num_active_paths);
-  ccl_gpu_kernel_lambda_pass.num_active_paths = num_active_paths;
-
+  ccl_gpu_kernel_lambda((state >= num_active_paths) && (INTEGRATOR_STATE(state, path, queued_kernel) != 0),
+                        int num_active_paths)
+    .num_active_paths = num_active_paths;
+  
   gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
       num_states, indices, num_indices, ccl_gpu_kernel_lambda_pass);
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
+ccl_gpu_kernel(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_compact_states,
                              ccl_global const int *active_terminated_states,
                              const int active_states_offset,
@@ -396,23 +390,22 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
   }
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
+ccl_gpu_kernel(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_compact_shadow_paths_array,
                              int num_states,
                              ccl_global int *indices,
-                             ccl_global int *num_indices,
-                             int num_active_paths)
+                         ccl_global int *num_indices,
+                         int num_active_paths)
 {
-  ccl_gpu_kernel_lambda((state >= num_active_paths) &&
-                            (INTEGRATOR_STATE(state, shadow_path, queued_kernel) != 0),
-                        int num_active_paths);
-  ccl_gpu_kernel_lambda_pass.num_active_paths = num_active_paths;
+  ccl_gpu_kernel_lambda((state >= num_active_paths) && (INTEGRATOR_STATE(state, shadow_path, queued_kernel) != 0),
+                        int num_active_paths)
+    .num_active_paths = num_active_paths;
 
   gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
       num_states, indices, num_indices, ccl_gpu_kernel_lambda_pass);
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
+ccl_gpu_kernel(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_compact_shadow_states,
                              ccl_global const int *active_terminated_states,
                              const int active_states_offset,
@@ -429,7 +422,7 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
   }
 }
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE) ccl_gpu_kernel_signature(
+ccl_gpu_kernel(GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE) ccl_gpu_kernel_signature(
     prefix_sum, ccl_global int *counter, ccl_global int *prefix_sum, int num_values)
 {
   gpu_parallel_prefix_sum(ccl_gpu_global_id_x(), counter, prefix_sum, num_values);
@@ -467,7 +460,7 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
   const auto num_active_pixels_mask = ccl_gpu_ballot(!converged);
   const int lane_id = ccl_gpu_thread_idx_x % ccl_gpu_warp_size;
   if (lane_id == 0) {
-    atomic_fetch_and_add_uint32(num_active_pixels, popcount(num_active_pixels_mask));
+    atomic_fetch_and_add_uint32(num_active_pixels, ccl_gpu_popc(num_active_pixels_mask));
   }
 }
 
@@ -527,31 +520,11 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
  * Film.
  */
 
-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
-}
-
 #define KERNEL_FILM_CONVERT_VARIANT(variant, input_channel_count) \
   ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) \
       ccl_gpu_kernel_signature(film_convert_##variant, \
                                const KernelFilmConvert kfilm_convert, \
-                               ccl_global float *pixels, \
+                               ccl_global uchar4 *rgba, \
                                ccl_global float *render_buffer, \
                                int num_pixels, \
                                int width, \
@@ -571,10 +544,20 @@ ccl_device_inline void kernel_gpu_film_convert_half_write(ccl_global uchar4 *rgb
     ccl_global const float *buffer = render_buffer + offset + x * kfilm_convert.pass_stride + \
                                      y * stride * kfilm_convert.pass_stride; \
 \
-    ccl_global float *pixel = pixels + \
-                              (render_pixel_index + rgba_offset) * kfilm_convert.pixel_stride; \
-\
+    float pixel[4]; \
     film_get_pass_pixel_##variant(&kfilm_convert, buffer, pixel); \
+\
+    film_apply_pass_pixel_overlays_rgba(&kfilm_convert, buffer, pixel); \
+\
+    if (input_channel_count == 1) { \
+      pixel[1] = pixel[2] = pixel[0]; \
+    } \
+    if (input_channel_count <= 3) { \
+      pixel[3] = 1.0f; \
+    } \
+\
+    ccl_global float *out = ((ccl_global float *)rgba) + rgba_offset + y * rgba_stride + x; \
+    *(ccl_global float4 *)out = make_float4(pixel[0], pixel[1], pixel[2], pixel[3]); \
   } \
 \
   ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) \
@@ -602,6 +585,8 @@ ccl_device_inline void kernel_gpu_film_convert_half_write(ccl_global uchar4 *rgb
 \
     float pixel[4]; \
     film_get_pass_pixel_##variant(&kfilm_convert, buffer, pixel); \
+\
+    film_apply_pass_pixel_overlays_rgba(&kfilm_convert, buffer, pixel); \
 \
     if (input_channel_count == 1) { \
       pixel[1] = pixel[2] = pixel[0]; \
@@ -610,11 +595,8 @@ ccl_device_inline void kernel_gpu_film_convert_half_write(ccl_global uchar4 *rgb
       pixel[3] = 1.0f; \
     } \
 \
-    film_apply_pass_pixel_overlays_rgba(&kfilm_convert, buffer, pixel); \
-\
-    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); \
+    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])); \
   }
 
 /* 1 channel inputs */
@@ -635,8 +617,6 @@ KERNEL_FILM_CONVERT_VARIANT(shadow_catcher_matte_with_shadow, 4)
 KERNEL_FILM_CONVERT_VARIANT(combined, 4)
 KERNEL_FILM_CONVERT_VARIANT(float4, 4)
 
-#undef KERNEL_FILM_CONVERT_VARIANT
-
 /* --------------------------------------------------------------------
  * Shader evaluation.
  */
@@ -895,6 +875,6 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
   const auto can_split_mask = ccl_gpu_ballot(can_split);
   const int lane_id = ccl_gpu_thread_idx_x % ccl_gpu_warp_size;
   if (lane_id == 0) {
-    atomic_fetch_and_add_uint32(num_possible_splits, popcount(can_split_mask));
+    atomic_fetch_and_add_uint32(num_possible_splits, ccl_gpu_popc(can_split_mask));
   }
 }

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

@@ -85,8 +85,8 @@ __device__ void gpu_parallel_active_index_array(const uint num_states,
     const uint is_active = (state_index < num_states) ? is_active_op(state_index) : 0;
 
     /* For each thread within a warp compute how many other active states precede it. */
-    const uint thread_offset = popcount(ccl_gpu_ballot(is_active) &
-                                        ccl_gpu_thread_mask(thread_warp));
+    const uint thread_offset = ccl_gpu_popc(ccl_gpu_ballot(is_active) &
+                                            ccl_gpu_thread_mask(thread_warp));
 
     /* Last thread in warp stores number of active states for each warp. */
     if (thread_warp == ccl_gpu_warp_size - 1) {

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

@@ -45,9 +45,8 @@ typedef unsigned long long uint64_t;
 #define ccl_device_forceinline __device__ __forceinline__
 #define ccl_device_noinline __device__ __noinline__
 #define ccl_device_noinline_cpu ccl_device
-#define ccl_device_inline_method ccl_device
 #define ccl_global
-#define ccl_inline_constant __constant__
+#define ccl_static_constant __constant__
 #define ccl_device_constant __constant__ __device__
 #define ccl_constant const
 #define ccl_gpu_shared __shared__
@@ -85,6 +84,7 @@ typedef unsigned long long uint64_t;
 #define ccl_gpu_syncthreads() __syncthreads()
 #define ccl_gpu_ballot(predicate) __ballot(predicate)
 #define ccl_gpu_shfl_down_sync(mask, var, detla) __shfl_down(var, detla)
+#define ccl_gpu_popc(x) __popc(x)
 
 /* GPU texture objects */
 typedef hipTextureObject_t ccl_gpu_tex_object;

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

@@ -35,19 +35,25 @@
 
 /* Compute number of threads per block and minimum blocks per multiprocessor
  * given the maximum number of registers per thread. */
-#define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
+
+#define ccl_gpu_kernel_threads(block_num_threads) \
+  extern "C" __global__ void __launch_bounds__(block_num_threads)
+
+#define ccl_gpu_kernel_threads_registers(block_num_threads, thread_num_registers) \
   extern "C" __global__ void __launch_bounds__(block_num_threads, \
                                                GPU_MULTIPRESSOR_MAX_REGISTERS / \
                                                    (block_num_threads * thread_num_registers))
 
-#define ccl_gpu_kernel_threads(block_num_threads) \
-  extern "C" __global__ void __launch_bounds__(block_num_threads)
+/* allow ccl_gpu_kernel to accept 1 or 2 parameters */
+#define SELECT_MACRO(_1, _2, NAME, ...) NAME
+#define ccl_gpu_kernel(...) \
+  SELECT_MACRO(__VA_ARGS__, ccl_gpu_kernel_threads_registers, ccl_gpu_kernel_threads)(__VA_ARGS__)
 
 #define ccl_gpu_kernel_signature(name, ...) kernel_gpu_##name(__VA_ARGS__)
 
 #define ccl_gpu_kernel_call(x) x
 
-/* Define a function object where "func" is the lambda body, and additional parameters are used to
+/* define a function object where "func" is the lambda body, and additional parameters are used to
  * specify captured state  */
 #define ccl_gpu_kernel_lambda(func, ...) \
   struct KernelLambda { \
@@ -56,7 +62,8 @@
     { \
       return (func); \
     } \
-  } ccl_gpu_kernel_lambda_pass
+  } ccl_gpu_kernel_lambda_pass; \
+  ccl_gpu_kernel_lambda_pass
 
 /* sanity checks */
 

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

@@ -34,7 +34,6 @@ using namespace metal;
 
 #pragma clang diagnostic ignored "-Wunused-variable"
 #pragma clang diagnostic ignored "-Wsign-compare"
-#pragma clang diagnostic ignored "-Wuninitialized"
 
 /* Qualifiers */
 
@@ -43,9 +42,8 @@ using namespace metal;
 #define ccl_device_forceinline ccl_device
 #define ccl_device_noinline ccl_device __attribute__((noinline))
 #define ccl_device_noinline_cpu ccl_device
-#define ccl_device_inline_method ccl_device
 #define ccl_global device
-#define ccl_inline_constant static constant constexpr
+#define ccl_static_constant static constant constexpr
 #define ccl_device_constant constant
 #define ccl_constant const device
 #define ccl_gpu_shared threadgroup
@@ -66,17 +64,15 @@ using namespace metal;
 #define ccl_gpu_thread_mask(thread_warp) uint64_t((1ull << thread_warp) - 1)
 
 #define ccl_gpu_ballot(predicate) ((uint64_t)((simd_vote::vote_t)simd_ballot(predicate)))
-#define ccl_gpu_syncthreads() threadgroup_barrier(mem_flags::mem_threadgroup);
+#define ccl_gpu_popc(x) popcount(x)
 
 // clang-format off
 
 /* kernel.h adapters */
 
-#define ccl_gpu_kernel(block_num_threads, thread_num_registers)
-#define ccl_gpu_kernel_threads(block_num_threads)
+#define ccl_gpu_kernel(...)
 
-/* Convert a comma-separated list into a semicolon-separated list
- * (so that we can generate a struct based on kernel entry-point parameters). */
+/* convert a comma-separated list into a semicolon-separated list (so that we can generate a struct based on kernel entrypoint parameters) */
 #define FN0()
 #define FN1(p1) p1;
 #define FN2(p1, p2) p1; p2;
@@ -97,8 +93,7 @@ using namespace metal;
 #define GET_LAST_ARG(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, ...) p16
 #define PARAMS_MAKER(...) GET_LAST_ARG(__VA_ARGS__, FN16, FN15, FN14, FN13, FN12, FN11, FN10, FN9, FN8, FN7, FN6, FN5, FN4, FN3, FN2, FN1, FN0)
 
-/* Generate a struct containing the entry-point parameters and a "run"
- * method which can access them implicitly via this-> */
+/* generate a struct containing the entrypoint parameters and a "run" method which can access them implicitly via this-> */
 #define ccl_gpu_kernel_signature(name, ...) \
 struct kernel_gpu_##name \
 { \
@@ -125,6 +120,7 @@ kernel void kernel_metal_##name(device const kernel_gpu_##name *params_struct, \
                                 uint simd_group_index [[simdgroup_index_in_threadgroup]], \
                                 uint num_simd_groups [[simdgroups_per_threadgroup]]) { \
   MetalKernelContext context(_launch_params_metal, _metal_ancillaries); \
+  INIT_DEBUG_BUFFER \
   params_struct->run(context, simdgroup_offset, metal_global_id, metal_local_id, metal_local_size, simdgroup_size, simd_lane_index, simd_group_index, num_simd_groups); \
 } \
 void kernel_gpu_##name::run(thread MetalKernelContext& context, \
@@ -147,35 +143,10 @@ void kernel_gpu_##name::run(thread MetalKernelContext& context, \
     ccl_private MetalKernelContext &context; \
     __VA_ARGS__; \
     int operator()(const int state) const { return (func); } \
-  } ccl_gpu_kernel_lambda_pass(context)
+  }ccl_gpu_kernel_lambda_pass(context); ccl_gpu_kernel_lambda_pass
 
 // clang-format on
 
-/* volumetric lambda functions - use function objects for lambda-like functionality */
-#define VOLUME_READ_LAMBDA(function_call) \
-  struct FnObjectRead { \
-    KernelGlobals kg; \
-    ccl_private MetalKernelContext *context; \
-    int state; \
-\
-    VolumeStack operator()(const int i) const \
-    { \
-      return context->function_call; \
-    } \
-  } volume_read_lambda_pass{kg, this, state};
-
-#define VOLUME_WRITE_LAMBDA(function_call) \
-  struct FnObjectWrite { \
-    KernelGlobals kg; \
-    ccl_private MetalKernelContext *context; \
-    int state; \
-\
-    void operator()(const int i, VolumeStack entry) const \
-    { \
-      context->function_call; \
-    } \
-  } volume_write_lambda_pass{kg, this, state};
-
 /* make_type definitions with Metal style element initializers */
 #ifdef make_float2
 #  undef make_float2
@@ -230,7 +201,6 @@ void kernel_gpu_##name::run(thread MetalKernelContext& context, \
 #define sinhf(x) sinh(float(x))
 #define coshf(x) cosh(float(x))
 #define tanhf(x) tanh(float(x))
-#define saturatef(x) saturate(float(x))
 
 /* Use native functions with possibly lower precision for performance,
  * no issues found so far. */
@@ -244,8 +214,6 @@ void kernel_gpu_##name::run(thread MetalKernelContext& context, \
 
 #define NULL 0
 
-#define __device__
-
 /* texture bindings and sampler setup */
 
 struct Texture2DParamsMetal {
@@ -260,9 +228,6 @@ struct MetalAncillaries {
   device Texture3DParamsMetal *textures_3d;
 };
 
-#include "util/half.h"
-#include "util/types.h"
-
 enum SamplerType {
   SamplerFilterNearest_AddressRepeat,
   SamplerFilterNearest_AddressClampEdge,
@@ -282,4 +247,4 @@ constant constexpr array<sampler, SamplerCount> metal_samplers = {
     sampler(address::repeat, filter::linear),
     sampler(address::clamp_to_edge, filter::linear),
     sampler(address::clamp_to_zero, filter::linear),
-};
+};

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

@@ -76,4 +76,4 @@ class MetalKernelContext {
     }
 #    include "kernel/device/gpu/image.h"
 
-  // clang-format on
+  // clang-format on

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

@@ -17,7 +17,7 @@
 ; /* end of MetalKernelContext class definition */
 
 /* Silently redirect into the MetalKernelContext instance */
-/* NOTE: These macros will need maintaining as entry-points change. */
+/* NOTE: These macros will need maintaining as entrypoints change */
 
 #undef kernel_integrator_state
 #define kernel_integrator_state context.launch_params_metal.__integrator_state

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

@@ -25,7 +25,7 @@ CCL_NAMESPACE_BEGIN
 
 typedef struct KernelParamsMetal {
 
-#define KERNEL_TEX(type, name) ccl_global const type *name;
+#define KERNEL_TEX(type, name) ccl_constant type *name;
 #include "kernel/textures.h"
 #undef KERNEL_TEX
 

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

@@ -49,11 +49,10 @@ typedef unsigned long long uint64_t;
   __device__ __forceinline__  // Function calls are bad for OptiX performance, so inline everything
 #define ccl_device_inline ccl_device
 #define ccl_device_forceinline ccl_device
-#define ccl_device_inline_method ccl_device
 #define ccl_device_noinline __device__ __noinline__
 #define ccl_device_noinline_cpu ccl_device
 #define ccl_global
-#define ccl_inline_constant __constant__
+#define ccl_static_constant __constant__
 #define ccl_device_constant __constant__ __device__
 #define ccl_constant const
 #define ccl_gpu_shared __shared__
@@ -87,6 +86,7 @@ typedef unsigned long long uint64_t;
 #define ccl_gpu_syncthreads() __syncthreads()
 #define ccl_gpu_ballot(predicate) __ballot_sync(0xFFFFFFFF, predicate)
 #define ccl_gpu_shfl_down_sync(mask, var, detla) __shfl_down_sync(mask, var, detla)
+#define ccl_gpu_popc(x) __popc(x)
 
 /* GPU texture objects */
 

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

@@ -21,8 +21,6 @@
 
 #include "kernel/device/gpu/image.h"  /* Texture lookup uses normal CUDA intrinsics. */
 
-#include "kernel/tables.h"
-
 #include "kernel/integrator/state.h"
 #include "kernel/integrator/state_flow.h"
 #include "kernel/integrator/state_util.h"
@@ -46,7 +44,7 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
 ccl_device_forceinline int get_object_id()
 {
 #ifdef __OBJECT_MOTION__
-  /* Always get the instance ID from the TLAS
+  /* Always get the the instance ID from the TLAS
    * There might be a motion transform node between TLAS and BLAS which does not have one. */
   return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
 #else
@@ -59,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, __params.render_buffer);
+  integrator_intersect_closest(nullptr, path_index);
 }
 
 extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_shadow()
@@ -161,9 +159,9 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
 
   /* Record geometric normal. */
   const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
-  const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0);
-  const float3 tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1);
-  const float3 tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
+  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0));
+  const float3 tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1));
+  const float3 tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
   local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
 
   /* Continue tracing (without this the trace call would return after the first hit). */

intern/cycles/kernel/film/accumulate.h

@@ -33,72 +33,62 @@ CCL_NAMESPACE_BEGIN
  * them separately. */
 
 ccl_device_inline void bsdf_eval_init(ccl_private BsdfEval *eval,
-                                      const ClosureType closure_type,
+                                      const bool is_diffuse,
                                       float3 value)
 {
   eval->diffuse = zero_float3();
   eval->glossy = zero_float3();
 
-  if (CLOSURE_IS_BSDF_DIFFUSE(closure_type)) {
+  if (is_diffuse) {
     eval->diffuse = value;
   }
-  else if (CLOSURE_IS_BSDF_GLOSSY(closure_type)) {
+  else {
     eval->glossy = value;
   }
-
-  eval->sum = value;
 }
 
 ccl_device_inline void bsdf_eval_accum(ccl_private BsdfEval *eval,
-                                       const ClosureType closure_type,
-                                       float3 value)
+                                       const bool is_diffuse,
+                                       float3 value,
+                                       float mis_weight)
 {
-  if (CLOSURE_IS_BSDF_DIFFUSE(closure_type)) {
+  value *= mis_weight;
+
+  if (is_diffuse) {
     eval->diffuse += value;
   }
-  else if (CLOSURE_IS_BSDF_GLOSSY(closure_type)) {
+  else {
     eval->glossy += value;
   }
-
-  eval->sum += value;
 }
 
 ccl_device_inline bool bsdf_eval_is_zero(ccl_private BsdfEval *eval)
 {
-  return is_zero(eval->sum);
+  return is_zero(eval->diffuse) && is_zero(eval->glossy);
 }
 
 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->sum;
+  return eval->diffuse + eval->glossy;
 }
 
-ccl_device_inline float3 bsdf_eval_pass_diffuse_weight(ccl_private const BsdfEval *eval)
+ccl_device_inline float3 bsdf_eval_diffuse_glossy_ratio(ccl_private const BsdfEval *eval)
 {
-  /* Ratio of diffuse weight to recover proportions for writing to render pass.
+  /* Ratio of diffuse and glossy 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->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);
+  return safe_divide_float3_float3(eval->diffuse, eval->diffuse + eval->glossy);
 }
 
 /* --------------------------------------------------------------------
@@ -151,8 +141,7 @@ ccl_device_forceinline ccl_global float *kernel_accum_pixel_render_buffer(
 ccl_device_inline int kernel_accum_sample(KernelGlobals kg,
                                           ConstIntegratorState state,
                                           ccl_global float *ccl_restrict render_buffer,
-                                          int sample,
-                                          int sample_offset)
+                                          int sample)
 {
   if (kernel_data.film.pass_sample_count == PASS_UNUSED) {
     return sample;
@@ -160,9 +149,7 @@ ccl_device_inline int kernel_accum_sample(KernelGlobals kg,
 
   ccl_global float *buffer = kernel_accum_pixel_render_buffer(kg, state, render_buffer);
 
-  return atomic_fetch_and_add_uint32(
-             (ccl_global uint *)(buffer) + kernel_data.film.pass_sample_count, 1) +
-         sample_offset;
+  return atomic_fetch_and_add_uint32((uint *)(buffer) + kernel_data.film.pass_sample_count, 1);
 }
 
 ccl_device void kernel_accum_adaptive_buffer(KernelGlobals kg,
@@ -364,48 +351,38 @@ 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 (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);
+  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);
 
-      /* 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);
-      }
-
-      /* 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;
-      }
+    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);
     }
-    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;
+
+    /* 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);
     }
   }
+  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. */
   if (pass_offset != PASS_UNUSED) {
@@ -449,60 +426,49 @@ 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 (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
-      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 (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;
+      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);
       }
 
-      /* Single write call for GPU coherence. */
+      /* 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) {
-        kernel_write_pass_float3(buffer + pass_offset, contribution);
+        contribution *= INTEGRATOR_STATE(state, shadow_path, diffuse_glossy_ratio);
       }
     }
+    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_TRANSPARENT_BACKGROUND)) {
+        (path_flag & PATH_RAY_CAMERA)) {
       const float3 unshadowed_throughput = INTEGRATOR_STATE(
           state, shadow_path, unshadowed_throughput);
       const float3 shadowed_throughput = INTEGRATOR_STATE(state, shadow_path, throughput);
@@ -553,7 +519,7 @@ ccl_device_inline void kernel_accum_background(KernelGlobals kg,
                                                const bool is_transparent_background_ray,
                                                ccl_global float *ccl_restrict render_buffer)
 {
-  float3 contribution = float3(INTEGRATOR_STATE(state, path, throughput)) * L;
+  float3 contribution = INTEGRATOR_STATE(state, path, throughput) * L;
   kernel_accum_clamp(kg, &contribution, INTEGRATOR_STATE(state, path, bounce) - 1);
 
   ccl_global float *buffer = kernel_accum_pixel_render_buffer(kg, state, render_buffer);

intern/cycles/kernel/film/passes.h

@@ -160,6 +160,40 @@ 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,
@@ -177,7 +211,7 @@ ccl_device_inline void kernel_write_data_passes(KernelGlobals kg,
 #ifdef __PASSES__
   const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
 
-  if (!(path_flag & PATH_RAY_TRANSPARENT_BACKGROUND)) {
+  if (!(path_flag & PATH_RAY_CAMERA)) {
     return;
   }