test using compiled function in math node

CMakeLists.txt

@@ -273,13 +273,11 @@ endif()
 
 if(UNIX AND NOT APPLE)
   option(WITH_SYSTEM_GLEW "Use GLEW OpenGL wrapper library provided by the operating system" OFF)
-  option(WITH_SYSTEM_GLEW "Use GLEW OpenGL wrapper library provided by the operating system" OFF)
+  option(WITH_SYSTEM_GLES "Use OpenGL ES library provided by the operating system"           ON)
-  option(WITH_SYSTEM_FREETYPE "Use the freetype library provided by the operating system" OFF)
 else()
   # not an option for other OS's
   set(WITH_SYSTEM_GLEW OFF)
   set(WITH_SYSTEM_GLES OFF)
-  set(WITH_SYSTEM_FREETYPE OFF)
 endif()
 
 
@@ -538,14 +536,12 @@ option(WITH_OPENGL              "When off limits visibility of the opengl header
 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)
-option(WITH_GPU_SHADER_BUILDER  "Shader builder is a developer option enabling linting on GLSL during compilation"                                  OFF)
 
 mark_as_advanced(
   WITH_OPENGL
   WITH_GLEW_ES
   WITH_GL_EGL
   WITH_GL_PROFILE_ES20
-  WITH_GPU_SHADER_BUILDER
 )
 
 if(WIN32)
@@ -563,18 +559,12 @@ if(WIN32)
   set(CPACK_INSTALL_PREFIX ${CMAKE_GENERIC_PROGRAM_FILES}/${})
 endif()
 
-# Compiler tool-chain.
+# Compiler toolchain
-if(UNIX AND NOT APPLE)
+if(CMAKE_COMPILER_IS_GNUCC)
-  if(CMAKE_COMPILER_IS_GNUCC)
+  option(WITH_LINKER_GOLD "Use ld.gold linker which is usually faster than ld.bfd" ON)
-    option(WITH_LINKER_GOLD "Use ld.gold linker which is usually faster than ld.bfd" ON)
+  mark_as_advanced(WITH_LINKER_GOLD)
-    mark_as_advanced(WITH_LINKER_GOLD)
+  option(WITH_LINKER_LLD "Use ld.lld linker which is usually faster than ld.gold" OFF)
-    option(WITH_LINKER_LLD "Use ld.lld linker which is usually faster than ld.gold" OFF)
+  mark_as_advanced(WITH_LINKER_LLD)
-    mark_as_advanced(WITH_LINKER_LLD)
-  endif()
-  if(CMAKE_COMPILER_IS_GNUCC OR CMAKE_C_COMPILER_ID MATCHES "Clang")
-    option(WITH_LINKER_MOLD "Use ld.mold linker which is usually faster than ld.gold & ld.lld." OFF)
-    mark_as_advanced(WITH_LINKER_MOLD)
-  endif()
 endif()
 
 option(WITH_COMPILER_ASAN "Build and link against address sanitizer (only for Debug & RelWithDebInfo targets)." OFF)
@@ -710,12 +700,9 @@ if(UNIX AND NOT APPLE)
 endif()
 
 # Installation process.
-set(POSTINSTALL_SCRIPT "" CACHE FILEPATH "Run given CMake script after installation process")
+option(POSTINSTALL_SCRIPT "Run given CMake script after installation process" OFF)
 mark_as_advanced(POSTINSTALL_SCRIPT)
 
-set(POSTCONFIGURE_SCRIPT "" CACHE FILEPATH "Run given CMake script as the last step of CMake configuration")
-mark_as_advanced(POSTCONFIGURE_SCRIPT)
-
 # end option(...)
 
 
@@ -2079,8 +2066,3 @@ endif()
 if(0)
   print_all_vars()
 endif()
-
-# Should be the last step of configuration.
-if(POSTCONFIGURE_SCRIPT)
-  include(${POSTCONFIGURE_SCRIPT})
-endif()

build_files/build_environment/CMakeLists.txt

@@ -63,7 +63,6 @@ include(cmake/jpeg.cmake)
 include(cmake/blosc.cmake)
 include(cmake/pthreads.cmake)
 include(cmake/openexr.cmake)
-include(cmake/brotli.cmake)
 include(cmake/freetype.cmake)
 include(cmake/freeglut.cmake)
 include(cmake/glew.cmake)

build_files/build_environment/cmake/boost.cmake

@@ -25,13 +25,8 @@ else()
 endif()
 
 if(WIN32)
-  if(MSVC_VERSION GREATER_EQUAL 1920) # 2019
+  set(BOOST_TOOLSET toolset=msvc-14.1)
-    set(BOOST_TOOLSET toolset=msvc-14.2)
+  set(BOOST_COMPILER_STRING -vc141)
-    set(BOOST_COMPILER_STRING -vc142)
-  else() # 2017
-    set(BOOST_TOOLSET toolset=msvc-14.1)
-    set(BOOST_COMPILER_STRING -vc141)
-  endif()
 
   set(BOOST_CONFIGURE_COMMAND bootstrap.bat)
   set(BOOST_BUILD_COMMAND b2)

build_files/build_environment/cmake/brotli.cmake

@@ -1,38 +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.
-#
-# ***** END GPL LICENSE BLOCK *****
-
-set(BROTLI_EXTRA_ARGS
-)
-
-ExternalProject_Add(external_brotli
-  URL file://${PACKAGE_DIR}/${BROTLI_FILE}
-  DOWNLOAD_DIR ${DOWNLOAD_DIR}
-  URL_HASH ${BROTLI_HASH_TYPE}=${BROTLI_HASH}
-  PREFIX ${BUILD_DIR}/brotli
-  CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/brotli ${DEFAULT_CMAKE_FLAGS} ${BROTLI_EXTRA_ARGS}
-  INSTALL_DIR ${LIBDIR}/brotli
-)
-
-if(BUILD_MODE STREQUAL Release AND WIN32)
-  ExternalProject_Add_Step(external_brotli after_install
-    COMMAND ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/brotli/include ${HARVEST_TARGET}/brotli/include
-    COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/brotli/lib/brotlidec-static${LIBEXT} ${HARVEST_TARGET}/brotli/lib/brotlidec-static${LIBEXT}
-    COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/brotli/lib/brotlicommon-static${LIBEXT} ${HARVEST_TARGET}/brotli/lib/brotlicommon-static${LIBEXT}
-    DEPENDEES install
-  )
-endif()

build_files/build_environment/cmake/download.cmake

@@ -94,4 +94,3 @@ download_source(POTRACE)
 download_source(HARU)
 download_source(ZSTD)
 download_source(FLEX)
-download_source(BROTLI)

build_files/build_environment/cmake/freetype.cmake

@@ -23,12 +23,9 @@ set(FREETYPE_EXTRA_ARGS
   -DWITH_HarfBuzz=OFF
   -DFT_WITH_HARFBUZZ=OFF
   -DFT_WITH_BZIP2=OFF
-  -DFT_WITH_BROTLI=ON
   -DCMAKE_DISABLE_FIND_PACKAGE_HarfBuzz=TRUE
   -DCMAKE_DISABLE_FIND_PACKAGE_BZip2=TRUE
-  -DPC_BROTLIDEC_INCLUDEDIR=${LIBDIR}/brotli/include
+  -DCMAKE_DISABLE_FIND_PACKAGE_BrotliDec=TRUE)
-  -DPC_BROTLIDEC_LIBDIR=${LIBDIR}/brotli/lib
-  )
 
 ExternalProject_Add(external_freetype
   URL file://${PACKAGE_DIR}/${FREETYPE_FILE}
@@ -39,11 +36,6 @@ ExternalProject_Add(external_freetype
   INSTALL_DIR ${LIBDIR}/freetype
 )
 
-add_dependencies(
-  external_freetype
-  external_brotli
-)
-
 if(BUILD_MODE STREQUAL Release AND WIN32)
   ExternalProject_Add_Step(external_freetype after_install
     COMMAND ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/freetype ${HARVEST_TARGET}/freetype

build_files/build_environment/cmake/gmp.cmake

@@ -38,6 +38,13 @@ elseif(UNIX AND NOT APPLE)
   )
 endif()
 
+if(BLENDER_PLATFORM_ARM)
+  set(GMP_OPTIONS
+    ${GMP_OPTIONS}
+    --disable-assembly
+  )
+endif()
+
 ExternalProject_Add(external_gmp
   URL file://${PACKAGE_DIR}/${GMP_FILE}
   DOWNLOAD_DIR ${DOWNLOAD_DIR}

build_files/build_environment/cmake/harvest.cmake

@@ -79,8 +79,6 @@ endfunction()
 harvest(alembic/include alembic/include "*.h")
 harvest(alembic/lib/libAlembic.a alembic/lib/libAlembic.a)
 harvest(alembic/bin alembic/bin "*")
-harvest(brotli/include brotli/include "*.h")
-harvest(brotli/lib brotli/lib "*.a")
 harvest(boost/include boost/include "*")
 harvest(boost/lib boost/lib "*.a")
 harvest(ffmpeg/include ffmpeg/include "*.h")

build_files/build_environment/cmake/versions.cmake

@@ -83,9 +83,9 @@ else()
   set(OPENEXR_VERSION_POSTFIX)
 endif()
 
-set(FREETYPE_VERSION 2.11.0)
+set(FREETYPE_VERSION 2.10.2)
 set(FREETYPE_URI http://prdownloads.sourceforge.net/freetype/freetype-${FREETYPE_VERSION}.tar.gz)
-set(FREETYPE_HASH cf09172322f6b50cf8f568bf8fe14bde)
+set(FREETYPE_HASH b1cb620e4c875cd4d1bfa04945400945)
 set(FREETYPE_HASH_TYPE MD5)
 set(FREETYPE_FILE freetype-${FREETYPE_VERSION}.tar.gz)
 
@@ -474,9 +474,9 @@ set(ISPC_HASH 2e3abedbc0ea9aaec17d6562c632454d)
 set(ISPC_HASH_TYPE MD5)
 set(ISPC_FILE ispc-${ISPC_VERSION}.tar.gz)
 
-set(GMP_VERSION 6.2.1)
+set(GMP_VERSION 6.2.0)
 set(GMP_URI https://gmplib.org/download/gmp/gmp-${GMP_VERSION}.tar.xz)
-set(GMP_HASH 0b82665c4a92fd2ade7440c13fcaa42b)
+set(GMP_HASH a325e3f09e6d91e62101e59f9bda3ec1)
 set(GMP_HASH_TYPE MD5)
 set(GMP_FILE gmp-${GMP_VERSION}.tar.xz)
 
@@ -500,9 +500,3 @@ set(ZSTD_FILE zstd-${ZSTD_VERSION}.tar.gz)
 
 set(SSE2NEON_GIT https://github.com/DLTcollab/sse2neon.git)
 set(SSE2NEON_GIT_HASH fe5ff00bb8d19b327714a3c290f3e2ce81ba3525)
-
-set(BROTLI_VERSION v1.0.9)
-set(BROTLI_URI https://github.com/google/brotli/archive/refs/tags/${BROTLI_VERSION}.tar.gz)
-set(BROTLI_HASH f9e8d81d0405ba66d181529af42a3354f838c939095ff99930da6aa9cdf6fe46)
-set(BROTLI_HASH_TYPE SHA256)
-set(BROTLI_FILE brotli-${BROTLI_VERSION}.tar.gz)

build_files/build_environment/install_deps.sh

@@ -492,7 +492,7 @@ OIIO_SKIP=false
 LLVM_VERSION="12.0.0"
 LLVM_VERSION_SHORT="12.0"
 LLVM_VERSION_MIN="11.0"
-LLVM_VERSION_MEX="14.0"
+LLVM_VERSION_MEX="13.0"
 LLVM_VERSION_FOUND=""
 LLVM_FORCE_BUILD=false
 LLVM_FORCE_REBUILD=false
@@ -3620,8 +3620,8 @@ compile_FFmpeg() {
     fi
 
     ./configure --cc="gcc -Wl,--as-needed" \
-        --extra-ldflags="-pthread" \
+        --extra-ldflags="-pthread -static-libgcc" \
-        --prefix=$_inst --enable-shared \
+        --prefix=$_inst --enable-static \
         --disable-ffplay --disable-doc \
         --enable-gray \
         --enable-avfilter --disable-vdpau \
@@ -5721,6 +5721,76 @@ install_OTHER() {
 # ----------------------------------------------------------------------------
 # Printing User Info
 
+print_info_ffmpeglink_DEB() {
+  dpkg -L $_packages | grep -e ".*\/lib[^\/]\+\.so" | gawk '{ printf(nlines ? "'"$_ffmpeg_list_sep"'%s" : "%s", gensub(/.*lib([^\/]+)\.so/, "\\1", "g", $0)); nlines++ }'
+}
+
+print_info_ffmpeglink_RPM() {
+  rpm -ql $_packages | grep -e ".*\/lib[^\/]\+\.so" | gawk '{ printf(nlines ? "'"$_ffmpeg_list_sep"'%s" : "%s", gensub(/.*lib([^\/]+)\.so/, "\\1", "g", $0)); nlines++ }'
+}
+
+print_info_ffmpeglink_ARCH() {
+  pacman -Ql $_packages | grep -e ".*\/lib[^\/]\+\.so$" | gawk '{ printf(nlines ? "'"$_ffmpeg_list_sep"'%s" : "%s", gensub(/.*lib([^\/]+)\.so/, "\\1", "g", $0)); nlines++ }'
+}
+
+print_info_ffmpeglink() {
+  # This func must only print a ';'-separated list of libs...
+  if [ -z "$DISTRO" ]; then
+    ERROR "Failed to detect distribution type"
+    exit 1
+  fi
+
+  # Create list of packages from which to get libs names...
+  _packages=""
+
+  if [ "$THEORA_USE" = true ]; then
+    _packages="$_packages $THEORA_DEV"
+  fi
+
+  if [ "$VORBIS_USE" = true ]; then
+    _packages="$_packages $VORBIS_DEV"
+  fi
+
+  if [ "$OGG_USE" = true ]; then
+    _packages="$_packages $OGG_DEV"
+  fi
+
+  if [ "$XVID_USE" = true ]; then
+    _packages="$_packages $XVID_DEV"
+  fi
+
+  if [ "$VPX_USE" = true ]; then
+    _packages="$_packages $VPX_DEV"
+  fi
+
+  if [ "$OPUS_USE" = true ]; then
+    _packages="$_packages $OPUS_DEV"
+  fi
+
+  if [ "$MP3LAME_USE" = true ]; then
+    _packages="$_packages $MP3LAME_DEV"
+  fi
+
+  if [ "$X264_USE" = true ]; then
+    _packages="$_packages $X264_DEV"
+  fi
+
+  if [ "$OPENJPEG_USE" = true ]; then
+    _packages="$_packages $OPENJPEG_DEV"
+  fi
+
+  if [ "$DISTRO" = "DEB" ]; then
+    print_info_ffmpeglink_DEB
+  elif [ "$DISTRO" = "RPM" ]; then
+    print_info_ffmpeglink_RPM
+  elif [ "$DISTRO" = "ARCH" ]; then
+    print_info_ffmpeglink_ARCH
+  # XXX TODO!
+  else
+    PRINT "<Could not determine additional link libraries needed for ffmpeg, replace this by valid list of libs...>"
+  fi
+}
+
 print_info() {
   PRINT ""
   PRINT ""
@@ -5731,7 +5801,7 @@ print_info() {
   PRINT "If you're using CMake add this to your configuration flags:"
 
   _buildargs="-U *SNDFILE* -U PYTHON* -U *BOOST* -U *Boost* -U *TBB*"
-  _buildargs="$_buildargs -U *OPENCOLORIO* -U *OPENEXR* -U *OPENIMAGEIO* -U *LLVM* -U *CLANG* -U *CYCLES*"
+  _buildargs="$_buildargs -U *OPENCOLORIO* -U *OPENEXR* -U *OPENIMAGEIO* -U *LLVM* -U *CYCLES*"
   _buildargs="$_buildargs -U *OPENSUBDIV* -U *OPENVDB*  -U *BLOSC* -U *COLLADA* -U *FFMPEG* -U *ALEMBIC* -U *USD*"
   _buildargs="$_buildargs -U *EMBREE* -U *OPENIMAGEDENOISE* -U *OPENXR*"
 
@@ -5932,10 +6002,12 @@ print_info() {
 
   if [ "$FFMPEG_SKIP" = false ]; then
     _1="-D WITH_CODEC_FFMPEG=ON"
+    _2="-D FFMPEG_LIBRARIES='avformat;avcodec;avutil;avdevice;swscale;swresample;lzma;rt;`print_info_ffmpeglink`'"
     PRINT "  $_1"
-    _buildargs="$_buildargs $_1"
+    PRINT "  $_2"
+    _buildargs="$_buildargs $_1 $_2"
     if [ -d $INST/ffmpeg ]; then
-      _1="-D FFMPEG_ROOT_DIR=$INST/ffmpeg"
+      _1="-D FFMPEG=$INST/ffmpeg"
       PRINT "  $_1"
       _buildargs="$_buildargs $_1"
     fi

build_files/build_environment/windows/build_deps.cmd

@@ -1,14 +1,21 @@
 @echo off
 if NOT "%1" == "" (
-	if "%1" == "2017" (
+	if "%1" == "2013" (
-    echo "Building for VS2017"
+    echo "Building for VS2013"
-    set VSVER=15.0
+    set VSVER=12.0
-    set VSVER_SHORT=15
+    set VSVER_SHORT=12
-    set BuildDir=VS15
+    set BuildDir=VS12
     goto par2
   )
-	if "%1" == "2019" (
+	if "%1" == "2015" (
-    echo "Building for VS2019"
+    echo "Building for VS2015"
+    set VSVER=14.0
+    set VSVER_SHORT=14
+    set BuildDir=VS14
+    goto par2
+  )
+	if "%1" == "2017" (
+    echo "Building for VS2017"
     set VSVER=15.0
     set VSVER_SHORT=15
     set BuildDir=VS15
@@ -18,22 +25,40 @@ if NOT "%1" == "" (
 )
 :usage
 
-Echo Usage build_deps 2017/2019 x64
+Echo Usage build_deps 2013/2015/2017 x64/x86
 goto exit
 :par2
 if NOT "%2" == "" (
+	if "%2" == "x86" (
+    echo "Building for x86"
+    set HARVESTROOT=Windows_vc
+    set ARCH=86
+		if "%1" == "2013" (
+			set CMAKE_BUILDER=Visual Studio 12 2013
+		)
+		if "%1" == "2015" (
+			set CMAKE_BUILDER=Visual Studio 14 2015
+		)
+		if "%1" == "2017" (
+			set CMAKE_BUILDER=Visual Studio 15 2017
+		)
+		
+    goto start
+  )
 	if "%2" == "x64" (
     echo "Building for x64"
     set HARVESTROOT=Win64_vc
     set ARCH=64
-		if "%1" == "2019" (
+		if "%1" == "2013" (
-			set CMAKE_BUILDER=Visual Studio 16 2019
+			set CMAKE_BUILDER=Visual Studio 12 2013 Win64
-			set CMAKE_BUILD_ARCH=-A x64
+		)
+		if "%1" == "2015" (
+			set CMAKE_BUILDER=Visual Studio 14 2015 Win64
 		)
 		if "%1" == "2017" (
 			set CMAKE_BUILDER=Visual Studio 15 2017 Win64
-            set CMAKE_BUILD_ARCH=
 		)
+		
     goto start
   )
 )
@@ -95,7 +120,7 @@ set path=%BUILD_DIR%\downloads\mingw\mingw64\msys\1.0\bin\;%BUILD_DIR%\downloads
 mkdir %STAGING%\%BuildDir%%ARCH%R
 cd %Staging%\%BuildDir%%ARCH%R
 echo %DATE% %TIME% : Start > %StatusFile%
-cmake -G "%CMAKE_BUILDER%" %CMAKE_BUILD_ARCH% -Thost=x64  %SOURCE_DIR% -DPACKAGE_DIR=%BUILD_DIR%/packages -DDOWNLOAD_DIR=%BUILD_DIR%/downloads -DBUILD_MODE=Release -DHARVEST_TARGET=%HARVEST_DIR%/%HARVESTROOT%%VSVER_SHORT%/
+cmake -G "%CMAKE_BUILDER%" -Thost=x64  %SOURCE_DIR% -DPACKAGE_DIR=%BUILD_DIR%/packages -DDOWNLOAD_DIR=%BUILD_DIR%/downloads -DBUILD_MODE=Release -DHARVEST_TARGET=%HARVEST_DIR%/%HARVESTROOT%%VSVER_SHORT%/
 echo %DATE% %TIME% : Release Configuration done >> %StatusFile%
 if "%dobuild%" == "1" (
 	msbuild /m "ll.vcxproj" /p:Configuration=Release /fl /flp:logfile=BlenderDeps_llvm.log;Verbosity=normal
@@ -108,7 +133,7 @@ if "%NODEBUG%" == "1" goto exit
 cd %BUILD_DIR%
 mkdir %STAGING%\%BuildDir%%ARCH%D
 cd %Staging%\%BuildDir%%ARCH%D
-cmake -G "%CMAKE_BUILDER%" %CMAKE_BUILD_ARCH% -Thost=x64 %SOURCE_DIR% -DPACKAGE_DIR=%BUILD_DIR%/packages -DDOWNLOAD_DIR=%BUILD_DIR%/downloads -DCMAKE_BUILD_TYPE=Debug -DBUILD_MODE=Debug -DHARVEST_TARGET=%HARVEST_DIR%/%HARVESTROOT%%VSVER_SHORT%/  %CMAKE_DEBUG_OPTIONS%
+cmake -G "%CMAKE_BUILDER%" -Thost=x64 %SOURCE_DIR% -DPACKAGE_DIR=%BUILD_DIR%/packages -DDOWNLOAD_DIR=%BUILD_DIR%/downloads -DCMAKE_BUILD_TYPE=Debug -DBUILD_MODE=Debug -DHARVEST_TARGET=%HARVEST_DIR%/%HARVESTROOT%%VSVER_SHORT%/  %CMAKE_DEBUG_OPTIONS%
 echo %DATE% %TIME% : Debug Configuration done >> %StatusFile%
 if "%dobuild%" == "1" (
 	msbuild /m "ll.vcxproj" /p:Configuration=Debug /fl /flp:logfile=BlenderDeps_llvm.log;;Verbosity=normal 

build_files/cmake/Modules/FindBrotli.cmake

@@ -1,83 +0,0 @@
-# - Find Brotli library (compression for freetype/woff2).
-# This module defines
-#  BROTLI_INCLUDE_DIRS, where to find Brotli headers, Set when
-#                       BROTLI_INCLUDE_DIR is found.
-#  BROTLI_LIBRARIES, libraries to link against to use Brotli.
-#  BROTLI_ROOT_DIR, The base directory to search for Brotli.
-#                   This can also be an environment variable.
-#  BROTLI_FOUND, If false, do not try to use Brotli.
-#
-
-#=============================================================================
-# Copyright 2022 Blender Foundation.
-#
-# Distributed under the OSI-approved BSD 3-Clause License,
-# see accompanying file BSD-3-Clause-license.txt for details.
-#=============================================================================
-
-# If BROTLI_ROOT_DIR was defined in the environment, use it.
-IF(NOT BROTLI_ROOT_DIR AND NOT $ENV{BROTLI_ROOT_DIR} STREQUAL "")
-  SET(BROTLI_ROOT_DIR $ENV{BROTLI_ROOT_DIR})
-ENDIF()
-
-SET(_BROTLI_SEARCH_DIRS
-  ${BROTLI_ROOT_DIR}
-)
-
-FIND_PATH(BROTLI_INCLUDE_DIR
-  NAMES
-    brotli/decode.h
-  HINTS
-    ${_BROTLI_SEARCH_DIRS}
-  PATH_SUFFIXES
-    include
-  DOC "Brotli header files"
-)
-
-FIND_LIBRARY(BROTLI_LIBRARY_COMMON
-  NAMES
-    # Some builds use a special `-static` postfix in their static libraries names.
-    brotlicommon-static
-    brotlicommon
-  HINTS
-    ${_BROTLI_SEARCH_DIRS}
-  PATH_SUFFIXES
-    lib64 lib lib/static
-  DOC "Brotli static common library"
-)
-FIND_LIBRARY(BROTLI_LIBRARY_DEC
-  NAMES
-    # Some builds use a special `-static` postfix in their static libraries names.
-    brotlidec-static
-    brotlidec
-  HINTS
-    ${_BROTLI_SEARCH_DIRS}
-  PATH_SUFFIXES
-    lib64 lib lib/static
-  DOC "Brotli static decode library"
-)
-
-
-IF(${BROTLI_LIBRARY_COMMON_NOTFOUND} or ${BROTLI_LIBRARY_DEC_NOTFOUND})
-  set(BROTLI_FOUND FALSE)
-ELSE()
-  # handle the QUIETLY and REQUIRED arguments and set BROTLI_FOUND to TRUE if
-  # all listed variables are TRUE
-  INCLUDE(FindPackageHandleStandardArgs)
-  FIND_PACKAGE_HANDLE_STANDARD_ARGS(Brotli DEFAULT_MSG BROTLI_LIBRARY_COMMON BROTLI_LIBRARY_DEC BROTLI_INCLUDE_DIR)
-
-  IF(BROTLI_FOUND)
-    get_filename_component(BROTLI_LIBRARY_DIR ${BROTLI_LIBRARY_COMMON} DIRECTORY)
-    SET(BROTLI_INCLUDE_DIRS ${BROTLI_INCLUDE_DIR})
-    SET(BROTLI_LIBRARIES ${BROTLI_LIBRARY_DEC} ${BROTLI_LIBRARY_COMMON})
-  ENDIF()
-ENDIF()
-
-MARK_AS_ADVANCED(
-  BROTLI_INCLUDE_DIR
-  BROTLI_LIBRARY_COMMON
-  BROTLI_LIBRARY_DEC
-  BROTLI_LIBRARY_DIR
-)
-
-UNSET(_BROTLI_SEARCH_DIRS)

build_files/cmake/Modules/FindFFmpeg.cmake

@@ -33,8 +33,6 @@ if(NOT FFMPEG_FIND_COMPONENTS)
     avfilter
     avformat
     avutil
-    swscale
-    swresample
   )
 endif()
 
@@ -52,9 +50,9 @@ foreach(_component ${FFMPEG_FIND_COMPONENTS})
   string(TOUPPER ${_component} _upper_COMPONENT)
   find_library(FFMPEG_${_upper_COMPONENT}_LIBRARY
     NAMES
-      ${_component}
+      ${_upper_COMPONENT}
     HINTS
-      ${_ffmpeg_SEARCH_DIRS}
+      ${LIBDIR}/ffmpeg
     PATH_SUFFIXES
       lib64 lib
   )

build_files/cmake/Modules/FindOptiX.cmake

@@ -21,7 +21,7 @@ ENDIF()
 
 SET(_optix_SEARCH_DIRS
   ${OPTIX_ROOT_DIR}
-  "$ENV{PROGRAMDATA}/NVIDIA Corporation/OptiX SDK 7.3.0"
+  "$ENV{PROGRAMDATA}/NVIDIA Corporation/OptiX SDK 7.0.0"
 )
 
 FIND_PATH(OPTIX_INCLUDE_DIR

build_files/cmake/cmake_consistency_check.py

@@ -114,7 +114,7 @@ def is_c_header(filename: str) -> bool:
 
 def is_c(filename: str) -> bool:
     ext = splitext(filename)[1]
-    return (ext in {".c", ".cpp", ".cxx", ".m", ".mm", ".rc", ".cc", ".inl", ".metal"})
+    return (ext in {".c", ".cpp", ".cxx", ".m", ".mm", ".rc", ".cc", ".inl"})
 
 
 def is_c_any(filename: str) -> bool:

build_files/cmake/macros.cmake

@@ -488,6 +488,7 @@ function(blender_add_test_executable
 
   include_directories(${includes})
   include_directories(${includes_sys})
+  setup_libdirs()
 
   BLENDER_SRC_GTEST_EX(
     NAME ${name}
@@ -524,6 +525,83 @@ function(setup_heavy_lib_pool)
   endif()
 endfunction()
 
+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.
+  if(NOT MSVC AND NOT APPLE)
+    link_directories(${JPEG_LIBPATH} ${PNG_LIBPATH} ${ZLIB_LIBPATH} ${FREETYPE_LIBPATH})
+
+    if(WITH_PYTHON)  #  AND NOT WITH_PYTHON_MODULE  # WIN32 needs
+      link_directories(${PYTHON_LIBPATH})
+    endif()
+    if(WITH_SDL AND NOT WITH_SDL_DYNLOAD)
+      link_directories(${SDL_LIBPATH})
+    endif()
+    if(WITH_CODEC_FFMPEG)
+      link_directories(${FFMPEG_LIBPATH})
+    endif()
+    if(WITH_IMAGE_OPENEXR)
+      link_directories(${OPENEXR_LIBPATH})
+    endif()
+    if(WITH_IMAGE_TIFF)
+      link_directories(${TIFF_LIBPATH})
+    endif()
+    if(WITH_BOOST)
+      link_directories(${BOOST_LIBPATH})
+    endif()
+    if(WITH_OPENIMAGEIO)
+      link_directories(${OPENIMAGEIO_LIBPATH})
+    endif()
+    if(WITH_OPENIMAGEDENOISE)
+      link_directories(${OPENIMAGEDENOISE_LIBPATH})
+    endif()
+    if(WITH_OPENCOLORIO)
+      link_directories(${OPENCOLORIO_LIBPATH})
+    endif()
+    if(WITH_OPENVDB)
+      link_directories(${OPENVDB_LIBPATH})
+    endif()
+    if(WITH_OPENAL)
+      link_directories(${OPENAL_LIBPATH})
+    endif()
+    if(WITH_JACK AND NOT WITH_JACK_DYNLOAD)
+      link_directories(${JACK_LIBPATH})
+    endif()
+    if(WITH_PULSEAUDIO AND NOT WITH_PULSEAUDIO_DYNLOAD)
+      link_directories(${LIBPULSE_LIBPATH})
+    endif()
+    if(WITH_CODEC_SNDFILE)
+      link_directories(${LIBSNDFILE_LIBPATH})
+    endif()
+    if(WITH_FFTW3)
+      link_directories(${FFTW3_LIBPATH})
+    endif()
+    if(WITH_OPENCOLLADA)
+      link_directories(${OPENCOLLADA_LIBPATH})
+      # # Never set
+      # link_directories(${PCRE_LIBPATH})
+      # link_directories(${EXPAT_LIBPATH})
+    endif()
+    if(WITH_LLVM)
+      link_directories(${LLVM_LIBPATH})
+    endif()
+
+    if(WITH_ALEMBIC)
+      link_directories(${ALEMBIC_LIBPATH})
+    endif()
+
+    if(WITH_GMP)
+      link_directories(${GMP_LIBPATH})
+    endif()
+
+    if(WIN32 AND NOT UNIX)
+      link_directories(${PTHREADS_LIBPATH})
+    endif()
+  endif()
+endfunction()
+
 # Platform specific linker flags for targets.
 function(setup_platform_linker_flags
   target)
@@ -1214,6 +1292,29 @@ macro(openmp_delayload
     endif()
 endmacro()
 
+macro(blender_precompile_headers target cpp header)
+  if(MSVC)
+    # get the name for the pch output file
+    get_filename_component(pchbase ${cpp} NAME_WE)
+    set(pchfinal "${CMAKE_CURRENT_BINARY_DIR}/${CMAKE_CFG_INTDIR}/${pchbase}.pch")
+
+    # mark the cpp as the one outputting the pch
+    set_property(SOURCE ${cpp} APPEND PROPERTY OBJECT_OUTPUTS "${pchfinal}")
+
+    # get all sources for the target
+    get_target_property(sources ${target} SOURCES)
+
+    # make all sources depend on the pch to enforce the build order
+    foreach(src ${sources})
+      set_property(SOURCE ${src} APPEND PROPERTY OBJECT_DEPENDS "${pchfinal}")
+    endforeach()
+
+    target_sources(${target} PRIVATE ${cpp} ${header})
+    set_target_properties(${target} PROPERTIES COMPILE_FLAGS "/Yu${header} /Fp${pchfinal} /FI${header}")
+    set_source_files_properties(${cpp} PROPERTIES COMPILE_FLAGS "/Yc${header} /Fp${pchfinal}")
+  endif()
+endmacro()
+
 macro(set_and_warn_dependency
   _dependency _setting _val)
   # when $_dependency is disabled, forces $_setting = $_val

build_files/cmake/platform/platform_apple.cmake

@@ -166,18 +166,13 @@ if(WITH_FFTW3)
   find_package(Fftw3)
 endif()
 
-# FreeType compiled with Brotli compression for woff2.
 find_package(Freetype REQUIRED)
-list(APPEND FREETYPE_LIBRARIES
-  ${LIBDIR}/brotli/lib/libbrotlicommon-static.a
-  ${LIBDIR}/brotli/lib/libbrotlidec-static.a)
 
 if(WITH_IMAGE_OPENEXR)
   find_package(OpenEXR)
 endif()
 
 if(WITH_CODEC_FFMPEG)
-  set(FFMPEG_ROOT_DIR ${LIBDIR}/ffmpeg)
   set(FFMPEG_FIND_COMPONENTS
     avcodec avdevice avformat avutil
     mp3lame ogg opus swresample swscale
@@ -484,11 +479,8 @@ string(APPEND PLATFORM_LINKFLAGS " -stdlib=libc++")
 # Suppress ranlib "has no symbols" warnings (workaround for T48250)
 set(CMAKE_C_ARCHIVE_CREATE   "<CMAKE_AR> Scr <TARGET> <LINK_FLAGS> <OBJECTS>")
 set(CMAKE_CXX_ARCHIVE_CREATE "<CMAKE_AR> Scr <TARGET> <LINK_FLAGS> <OBJECTS>")
-# llvm-ranlib doesn't support this flag. Xcode's libtool does.
+set(CMAKE_C_ARCHIVE_FINISH   "<CMAKE_RANLIB> -no_warning_for_no_symbols -c <TARGET>")
-if(NOT ${CMAKE_RANLIB} MATCHES ".*llvm-ranlib$")
+set(CMAKE_CXX_ARCHIVE_FINISH "<CMAKE_RANLIB> -no_warning_for_no_symbols -c <TARGET>")
-  set(CMAKE_C_ARCHIVE_FINISH   "<CMAKE_RANLIB> -no_warning_for_no_symbols -c <TARGET>")
-  set(CMAKE_CXX_ARCHIVE_FINISH "<CMAKE_RANLIB> -no_warning_for_no_symbols -c <TARGET>")
-endif()
 
 if(WITH_COMPILER_CCACHE)
   if(NOT CMAKE_GENERATOR STREQUAL "Xcode")
@@ -515,6 +507,3 @@ list(APPEND CMAKE_BUILD_RPATH "${OpenMP_LIBRARY_DIR}")
 
 set(CMAKE_SKIP_INSTALL_RPATH FALSE)
 list(APPEND CMAKE_INSTALL_RPATH "@loader_path/../Resources/${BLENDER_VERSION}/lib")
-
-# Same as `CFBundleIdentifier` in Info.plist.
-set(CMAKE_XCODE_ATTRIBUTE_PRODUCT_BUNDLE_IDENTIFIER "org.blenderfoundation.blender")

build_files/cmake/platform/platform_apple_xcode.cmake

@@ -96,7 +96,7 @@ else()
     # Detect SDK version to use.
     if(NOT DEFINED OSX_SYSTEM)
       execute_process(
-          COMMAND xcrun --sdk macosx --show-sdk-version
+          COMMAND xcrun --show-sdk-version
           OUTPUT_VARIABLE OSX_SYSTEM
           OUTPUT_STRIP_TRAILING_WHITESPACE)
     endif()

build_files/cmake/platform/platform_unix.cmake

@@ -18,7 +18,7 @@
 # All rights reserved.
 # ***** END GPL LICENSE BLOCK *****
 
-# Libraries configuration for any *nix system including Linux and Unix (excluding APPLE).
+# Libraries configuration for any *nix system including Linux and Unix.
 
 # Detect precompiled library directory
 if(NOT DEFINED LIBDIR)
@@ -48,9 +48,6 @@ if(NOT DEFINED LIBDIR)
   unset(LIBDIR_CENTOS7_ABI)
 endif()
 
-# Support restoring this value once pre-compiled libraries have been handled.
-set(WITH_STATIC_LIBS_INIT ${WITH_STATIC_LIBS})
-
 if(EXISTS ${LIBDIR})
   message(STATUS "Using pre-compiled LIBDIR: ${LIBDIR}")
 
@@ -103,22 +100,7 @@ find_package_wrapper(JPEG REQUIRED)
 find_package_wrapper(PNG REQUIRED)
 find_package_wrapper(ZLIB REQUIRED)
 find_package_wrapper(Zstd REQUIRED)
-
+find_package_wrapper(Freetype REQUIRED)
-if(NOT WITH_SYSTEM_FREETYPE)
-  # FreeType compiled with Brotli compression for woff2.
-  find_package_wrapper(Freetype REQUIRED)
-  if(EXISTS ${LIBDIR})
-    find_package_wrapper(Brotli REQUIRED)
-
-    # NOTE: This is done on WIN32 & APPLE but fails on some Linux systems.
-    # See: https://devtalk.blender.org/t/22536
-    # So `BROTLI_LIBRARIES` need to be added  `FREETYPE_LIBRARIES`.
-    #
-    # list(APPEND FREETYPE_LIBRARIES
-    #   ${BROTLI_LIBRARIES}
-    # )
-  endif()
-endif()
 
 if(WITH_PYTHON)
   # No way to set py35, remove for now.
@@ -196,30 +178,26 @@ endif()
 
 if(WITH_CODEC_FFMPEG)
   if(EXISTS ${LIBDIR})
-    set(FFMPEG_ROOT_DIR ${LIBDIR}/ffmpeg)
+    # For precompiled lib directory, all ffmpeg dependencies are in the same folder
-    # Override FFMPEG components to also include static library dependencies
+    file(GLOB ffmpeg_libs ${LIBDIR}/ffmpeg/lib/*.a ${LIBDIR}/sndfile/lib/*.a)
-    # included with precompiled libraries, and to ensure correct link order.
+    set(FFMPEG ${LIBDIR}/ffmpeg CACHE PATH "FFMPEG Directory")
-    set(FFMPEG_FIND_COMPONENTS
+    set(FFMPEG_LIBRARIES ${ffmpeg_libs} ${ffmpeg_libs} CACHE STRING "FFMPEG Libraries")
-      avformat avcodec avdevice avutil swresample swscale
+  else()
-      sndfile
+    set(FFMPEG /usr CACHE PATH "FFMPEG Directory")
-      FLAC
+    set(FFMPEG_LIBRARIES avformat avcodec avutil avdevice swscale CACHE STRING "FFMPEG Libraries")
-      mp3lame
-      opus
-      theora theoradec theoraenc
-      vorbis vorbisenc vorbisfile ogg
-      vpx
-      x264
-      xvidcore)
-  elseif(FFMPEG)
-    # Old cache variable used for root dir, convert to new standard.
-    set(FFMPEG_ROOT_DIR ${FFMPEG})
   endif()
-  find_package(FFmpeg)
 
-  if(NOT FFMPEG_FOUND)
+  mark_as_advanced(FFMPEG)
-    set(WITH_CODEC_FFMPEG OFF)
+
-    message(STATUS "FFmpeg not found, disabling it")
+  # lame, but until we have proper find module for ffmpeg
+  set(FFMPEG_INCLUDE_DIRS ${FFMPEG}/include)
+  if(EXISTS "${FFMPEG}/include/ffmpeg/")
+    list(APPEND FFMPEG_INCLUDE_DIRS "${FFMPEG}/include/ffmpeg")
   endif()
+  # end lameness
+
+  mark_as_advanced(FFMPEG_LIBRARIES)
+  set(FFMPEG_LIBPATH ${FFMPEG}/lib)
 endif()
 
 if(WITH_FFTW3)
@@ -554,21 +532,6 @@ add_definitions(-D_LARGEFILE_SOURCE -D_FILE_OFFSET_BITS=64 -D_LARGEFILE64_SOURCE
 #
 # Keep last, so indirectly linked libraries don't override our own pre-compiled libs.
 
-if(EXISTS ${LIBDIR})
-  # Clear the prefix path as it causes the `LIBDIR` to override system locations.
-  unset(CMAKE_PREFIX_PATH)
-
-  # Since the pre-compiled `LIBDIR` directories have been handled, don't prefer static libraries.
-  set(WITH_STATIC_LIBS ${WITH_STATIC_LIBS_INIT})
-endif()
-
-if(WITH_SYSTEM_FREETYPE)
-  find_package_wrapper(Freetype)
-  if(NOT FREETYPE_FOUND)
-    message(FATAL_ERROR "Failed finding system FreeType version!")
-  endif()
-endif()
-
 if(WITH_LZO AND WITH_SYSTEM_LZO)
   find_package_wrapper(LZO)
   if(NOT LZO_FOUND)
@@ -681,9 +644,6 @@ endif()
 # ----------------------------------------------------------------------------
 # Compilers
 
-# Only set the linker once.
-set(_IS_LINKER_DEFAULT ON)
-
 # GNU Compiler
 if(CMAKE_COMPILER_IS_GNUCC)
   # ffp-contract=off:
@@ -702,89 +662,26 @@ if(CMAKE_COMPILER_IS_GNUCC)
   string(PREPEND CMAKE_CXX_FLAGS_RELWITHDEBINFO "${GCC_EXTRA_FLAGS_RELEASE} ")
   unset(GCC_EXTRA_FLAGS_RELEASE)
 
-  # NOTE(@campbellbarton): Eventually mold will be able to use `-fuse-ld=mold`,
+  if(WITH_LINKER_GOLD)
-  # however at the moment this only works for GCC 12.1+ (unreleased at time of writing).
-  # So a workaround is used here "-B" which points to another path to find system commands
-  # such as `ld`.
-  if(WITH_LINKER_MOLD AND _IS_LINKER_DEFAULT)
-    find_program(MOLD_BIN "mold")
-    mark_as_advanced(MOLD_BIN)
-    if(NOT MOLD_BIN)
-      message(STATUS "The \"mold\" binary could not be found, using system linker.")
-      set(WITH_LINKER_MOLD OFF)
-    else()
-      # By default mold installs the binary to:
-      # - `{PREFIX}/bin/mold` as well as a symbolic-link in...
-      # - `{PREFIX}/lib/mold/ld`.
-      # (where `PREFIX` is typically `/usr/`).
-      #
-      # This block of code finds `{PREFIX}/lib/mold` from the `mold` binary.
-      # Other methods of searching for the path could also be made to work,
-      # we could even make our own directory and symbolic-link, however it's more
-      # convenient to use the one provided by mold.
-      #
-      # Use the binary path to "mold", to find the common prefix which contains "lib/mold".
-      # The parent directory: e.g. `/usr/bin/mold` -> `/usr/bin/`.
-      get_filename_component(MOLD_PREFIX "${MOLD_BIN}" DIRECTORY)
-      # The common prefix path: e.g. `/usr/bin/` -> `/usr/` to use as a hint.
-      get_filename_component(MOLD_PREFIX "${MOLD_PREFIX}" DIRECTORY)
-      # Find `{PREFIX}/lib/mold/ld`, store the directory component (without the `ld`).
-      # Then pass `-B {PREFIX}/lib/mold` to GCC so the `ld` located there overrides the default.
-      find_path(
-        MOLD_BIN_DIR "ld"
-        HINTS "${MOLD_PREFIX}"
-        # The default path is `libexec`, Arch Linux for e.g.
-        # replaces this with `lib` so check both.
-        PATH_SUFFIXES "libexec/mold" "lib/mold" "lib64/mold"
-        NO_DEFAULT_PATH
-        NO_CACHE
-      )
-      if(NOT MOLD_BIN_DIR)
-        message(STATUS
-          "The mold linker could not find the directory containing the linker command "
-          "(typically "
-          "\"${MOLD_PREFIX}/libexec/mold/ld\") or "
-          "\"${MOLD_PREFIX}/lib/mold/ld\") using system linker.")
-        set(WITH_LINKER_MOLD OFF)
-      endif()
-      unset(MOLD_PREFIX)
-    endif()
-
-    if(WITH_LINKER_MOLD)
-      # GCC will search for `ld` in this directory first.
-      string(APPEND CMAKE_EXE_LINKER_FLAGS    " -B \"${MOLD_BIN_DIR}\"")
-      string(APPEND CMAKE_SHARED_LINKER_FLAGS " -B \"${MOLD_BIN_DIR}\"")
-      string(APPEND CMAKE_MODULE_LINKER_FLAGS " -B \"${MOLD_BIN_DIR}\"")
-      set(_IS_LINKER_DEFAULT OFF)
-    endif()
-    unset(MOLD_BIN)
-    unset(MOLD_BIN_DIR)
-  endif()
-
-  if(WITH_LINKER_GOLD AND _IS_LINKER_DEFAULT)
     execute_process(
       COMMAND ${CMAKE_C_COMPILER} -fuse-ld=gold -Wl,--version
       ERROR_QUIET OUTPUT_VARIABLE LD_VERSION)
     if("${LD_VERSION}" MATCHES "GNU gold")
-      string(APPEND CMAKE_EXE_LINKER_FLAGS    " -fuse-ld=gold")
+      string(APPEND CMAKE_C_FLAGS " -fuse-ld=gold")
-      string(APPEND CMAKE_SHARED_LINKER_FLAGS " -fuse-ld=gold")
+      string(APPEND CMAKE_CXX_FLAGS " -fuse-ld=gold")
-      string(APPEND CMAKE_MODULE_LINKER_FLAGS " -fuse-ld=gold")
-      set(_IS_LINKER_DEFAULT OFF)
     else()
       message(STATUS "GNU gold linker isn't available, using the default system linker.")
     endif()
     unset(LD_VERSION)
   endif()
 
-  if(WITH_LINKER_LLD AND _IS_LINKER_DEFAULT)
+  if(WITH_LINKER_LLD)
     execute_process(
       COMMAND ${CMAKE_C_COMPILER} -fuse-ld=lld -Wl,--version
       ERROR_QUIET OUTPUT_VARIABLE LD_VERSION)
     if("${LD_VERSION}" MATCHES "LLD")
-      string(APPEND CMAKE_EXE_LINKER_FLAGS    " -fuse-ld=lld")
+      string(APPEND CMAKE_C_FLAGS " -fuse-ld=lld")
-      string(APPEND CMAKE_SHARED_LINKER_FLAGS " -fuse-ld=lld")
+      string(APPEND CMAKE_CXX_FLAGS " -fuse-ld=lld")
-      string(APPEND CMAKE_MODULE_LINKER_FLAGS " -fuse-ld=lld")
-      set(_IS_LINKER_DEFAULT OFF)
     else()
       message(STATUS "LLD linker isn't available, using the default system linker.")
     endif()
@@ -794,28 +691,6 @@ if(CMAKE_COMPILER_IS_GNUCC)
 # CLang is the same as GCC for now.
 elseif(CMAKE_C_COMPILER_ID MATCHES "Clang")
   set(PLATFORM_CFLAGS "-pipe -fPIC -funsigned-char -fno-strict-aliasing")
-
-  if(WITH_LINKER_MOLD AND _IS_LINKER_DEFAULT)
-    find_program(MOLD_BIN "mold")
-    mark_as_advanced(MOLD_BIN)
-    if(NOT MOLD_BIN)
-      message(STATUS "The \"mold\" binary could not be found, using system linker.")
-      set(WITH_LINKER_MOLD OFF)
-    else()
-      if(CMAKE_C_COMPILER_VERSION VERSION_GREATER_EQUAL 12.0)
-        string(APPEND CMAKE_EXE_LINKER_FLAGS    " --ld-path=\"${MOLD_BIN}\"")
-        string(APPEND CMAKE_SHARED_LINKER_FLAGS " --ld-path=\"${MOLD_BIN}\"")
-        string(APPEND CMAKE_MODULE_LINKER_FLAGS " --ld-path=\"${MOLD_BIN}\"")
-      else()
-        string(APPEND CMAKE_EXE_LINKER_FLAGS    " -fuse-ld=\"${MOLD_BIN}\"")
-        string(APPEND CMAKE_SHARED_LINKER_FLAGS " -fuse-ld=\"${MOLD_BIN}\"")
-        string(APPEND CMAKE_MODULE_LINKER_FLAGS " -fuse-ld=\"${MOLD_BIN}\"")
-      endif()
-      set(_IS_LINKER_DEFAULT OFF)
-    endif()
-    unset(MOLD_BIN)
-  endif()
-
 # Intel C++ Compiler
 elseif(CMAKE_C_COMPILER_ID MATCHES "Intel")
   # think these next two are broken
@@ -839,8 +714,6 @@ elseif(CMAKE_C_COMPILER_ID MATCHES "Intel")
   string(APPEND PLATFORM_LINKFLAGS " -static-intel")
 endif()
 
-unset(_IS_LINKER_DEFAULT)
-
 # Avoid conflicts with Mesa llvmpipe, Luxrender, and other plug-ins that may
 # use the same libraries as Blender with a different version or build options.
 set(PLATFORM_LINKFLAGS

build_files/cmake/platform/platform_win32.cmake

@@ -347,11 +347,7 @@ set(FREETYPE_INCLUDE_DIRS
   ${LIBDIR}/freetype/include
   ${LIBDIR}/freetype/include/freetype2
 )
-set(FREETYPE_LIBRARIES
+set(FREETYPE_LIBRARY ${LIBDIR}/freetype/lib/freetype2ST.lib)
-  ${LIBDIR}/freetype/lib/freetype2ST.lib
-  ${LIBDIR}/brotli/lib/brotlidec-static.lib
-  ${LIBDIR}/brotli/lib/brotlicommon-static.lib
-)
 windows_find_package(freetype REQUIRED)
 
 if(WITH_FFTW3)

build_files/windows/find_dependencies.cmd

@@ -3,32 +3,7 @@ for %%X in (svn.exe) do (set SVN=%%~$PATH:X)
 for %%X in (cmake.exe) do (set CMAKE=%%~$PATH:X)
 for %%X in (ctest.exe) do (set CTEST=%%~$PATH:X)
 for %%X in (git.exe) do (set GIT=%%~$PATH:X)
-REM For python, default on 39 but if that does not exist also check
-REM the 310,311 and 312 folders to see if those are there, it checks
-REM this far ahead to ensure good lib folder compatiblity in the future.
 set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
-if EXIST %PYTHON% (
-	goto detect_python_done
-)
-set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\310\bin\python.exe
-if EXIST %PYTHON% (
-	goto detect_python_done
-)
-set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\311\bin\python.exe
-if EXIST %PYTHON% (
-	goto detect_python_done
-)
-set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\312\bin\python.exe
-if EXIST %PYTHON% (
-	goto detect_python_done
-)
-
-if NOT EXIST %PYTHON% (
-    echo Warning: Python not found, there is likely an issue with the library folder
-    set PYTHON=""
-)
-
-:detect_python_done
 if NOT "%verbose%" == "" (
 	echo svn    : "%SVN%"
 	echo cmake  : "%CMAKE%"
@@ -36,3 +11,7 @@ if NOT "%verbose%" == "" (
 	echo git    : "%GIT%"
 	echo python : "%PYTHON%"
 )
+if "%CMAKE%" == "" (
+	echo Cmake not found in path, required for building, exiting...
+	exit /b 1
+)

build_files/windows/format.cmd

@@ -9,11 +9,17 @@ exit /b 1
 :detect_done
 echo found clang-format in %CF_PATH%
 
-if NOT EXIST %PYTHON% (
+if EXIST %PYTHON% (
-    echo python not found, required for this operation
+    set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
-    exit /b 1
+    goto detect_python_done
 )
 
+echo python not found in lib folder
+exit /b 1
+
+:detect_python_done
+echo found python (%PYTHON%)
+
 set FORMAT_PATHS=%BLENDER_DIR%\source\tools\utils_maintenance\clang_format_paths.py
 
 REM The formatting script expects clang-format to be in the current PATH.

build_files/windows/icons.cmd

@@ -1,8 +1,18 @@
-if NOT EXIST %PYTHON% (
+if EXIST "%PYTHON%" (
-    echo python not found, required for this operation
+    goto detect_python_done
-    exit /b 1
 )
 
+set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
+if EXIST %PYTHON% (
+    goto detect_python_done
+)
+
+echo python not found at %PYTHON%
+exit /b 1
+
+:detect_python_done
+echo found python (%PYTHON%)
+
 call "%~dp0\find_inkscape.cmd"
 
 if EXIST "%INKSCAPE_BIN%" (

build_files/windows/icons_geom.cmd

@@ -1,8 +1,18 @@
-if NOT EXIST %PYTHON% (
+if EXIST %PYTHON% (
-    echo python not found, required for this operation
+    goto detect_python_done
-    exit /b 1
 )
 
+set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
+if EXIST %PYTHON% (
+    goto detect_python_done
+)
+
+echo python not found at %PYTHON%
+exit /b 1
+
+:detect_python_done
+echo found python (%PYTHON%)
+
 call "%~dp0\find_blender.cmd"
 
 if EXIST "%BLENDER_BIN%" (

build_files/windows/update_sources.cmd

@@ -1,7 +1,10 @@
-if NOT EXIST %PYTHON% (
+if EXIST %PYTHON% (
-    echo python not found, required for this operation
+	goto detect_python_done
-    exit /b 1
 )
+
+echo python not found in lib folder
+exit /b 1
+
 :detect_python_done
 
 REM Use -B to avoid writing __pycache__ in lib directory and causing update conflicts.

doc/doxygen/doxygen.intern.h

@@ -51,6 +51,9 @@
 /** \defgroup intern_mikktspace MikktSpace
  *  \ingroup intern */
 
+/** \defgroup intern_numaapi NUMA (Non Uniform Memory Architecture)
+ *  \ingroup intern */
+
 /** \defgroup intern_rigidbody Rigid-Body C-API
  *  \ingroup intern */
 

doc/python_api/examples/bpy.types.Bone.convert_local_to_pose.py

@@ -8,42 +8,27 @@ def set_pose_matrices(obj, matrix_map):
     "Assign pose space matrices of all bones at once, ignoring constraints."
 
     def rec(pbone, parent_matrix):
-        if pbone.name in matrix_map:
+        matrix = matrix_map[pbone.name]
-            matrix = matrix_map[pbone.name]
 
-            ## Instead of:
+        ## Instead of:
-            # pbone.matrix = matrix
+        # pbone.matrix = matrix
-            # bpy.context.view_layer.update()
+        # bpy.context.view_layer.update()
 
-            # Compute and assign local matrix, using the new parent matrix
+        # Compute and assign local matrix, using the new parent matrix
-            if pbone.parent:
+        if pbone.parent:
-                pbone.matrix_basis = pbone.bone.convert_local_to_pose(
+            pbone.matrix_basis = pbone.bone.convert_local_to_pose(
-                    matrix,
+                matrix,
-                    pbone.bone.matrix_local,
+                pbone.bone.matrix_local,
-                    parent_matrix=parent_matrix,
+                parent_matrix=parent_matrix,
-                    parent_matrix_local=pbone.parent.bone.matrix_local,
+                parent_matrix_local=pbone.parent.bone.matrix_local,
-                    invert=True
+                invert=True
-                )
+            )
-            else:
-                pbone.matrix_basis = pbone.bone.convert_local_to_pose(
-                    matrix,
-                    pbone.bone.matrix_local,
-                    invert=True
-                )
         else:
-            # Compute the updated pose matrix from local and new parent matrix
+            pbone.matrix_basis = pbone.bone.convert_local_to_pose(
-            if pbone.parent:
+                matrix,
-                matrix = pbone.bone.convert_local_to_pose(
+                pbone.bone.matrix_local,
-                    pbone.matrix_basis,
+                invert=True
-                    pbone.bone.matrix_local,
+            )
-                    parent_matrix=parent_matrix,
-                    parent_matrix_local=pbone.parent.bone.matrix_local,
-                )
-            else:
-                matrix = pbone.bone.convert_local_to_pose(
-                    pbone.matrix_basis,
-                    pbone.bone.matrix_local,
-                )
 
         # Recursively process children, passing the new matrix through
         for child in pbone.children:

doc/python_api/rst/info_gotcha.rst

@@ -743,7 +743,7 @@ will re-allocate objects data,
 any references to a meshes vertices/polygons/UVs, armatures bones,
 curves points, etc. cannot be accessed after switching mode.
 
-Only the reference to the data itself can be re-accessed, the following example will crash.
+Only the reference to the data its self can be re-accessed, the following example will crash.
 
 .. code-block:: python
 

doc/python_api/sphinx_doc_gen.py

@@ -1762,7 +1762,6 @@ except ModuleNotFoundError:
     fw("html_show_sphinx = False\n")
     fw("html_baseurl = 'https://docs.blender.org/api/current/'\n")
     fw("html_use_opensearch = 'https://docs.blender.org/api/current'\n")
-    fw("html_show_search_summary = True\n")
     fw("html_split_index = True\n")
     fw("html_static_path = ['static']\n")
     fw("html_extra_path = ['static/favicon.ico', 'static/blender_logo.svg']\n")

extern/audaspace/CMakeLists.txt

@@ -1092,12 +1092,12 @@ if(WITH_PYTHON)
 		configure_file(${PYTHON_SOURCE_DIRECTORY}/setup.py.in ${CMAKE_CURRENT_BINARY_DIR}/setup.py ESCAPE_QUOTES @ONLY)
 
 		if(APPLE)
-			add_custom_command(OUTPUT build COMMAND MACOSX_DEPLOYMENT_TARGET=${CMAKE_OSX_DEPLOYMENT_TARGET} ${PYTHON_EXECUTABLE} setup.py build DEPENDS ${PYTHON_SRC} ${PYTHON_HDR} setup.py)
+			add_custom_command(OUTPUT build COMMAND MACOSX_DEPLOYMENT_TARGET=${CMAKE_OSX_DEPLOYMENT_TARGET} ${PYTHON_EXECUTABLE} setup.py build DEPENDS ${PYTHON_SRC} ${PYTHON_HDR})
 		elseif(WIN32)
 			set(ENV{VS100COMNTOOLS} $ENV{VS120COMNTOOLS})
-			add_custom_command(OUTPUT build COMMAND ${PYTHON_EXECUTABLE} setup.py build DEPENDS ${PYTHON_SRC} ${PYTHON_HDR} setup.py)
+			add_custom_command(OUTPUT build COMMAND ${PYTHON_EXECUTABLE} setup.py build DEPENDS ${PYTHON_SRC} ${PYTHON_HDR})
 		else()
-			add_custom_command(OUTPUT build COMMAND ${PYTHON_EXECUTABLE} setup.py build DEPENDS ${PYTHON_SRC} ${PYTHON_HDR} setup.py)
+			add_custom_command(OUTPUT build COMMAND ${PYTHON_EXECUTABLE} setup.py build DEPENDS ${PYTHON_SRC} ${PYTHON_HDR})
 		endif()
 		add_custom_target(pythonmodule ALL DEPENDS build SOURCES ${PYTHON_SOURCE_DIRECTORY}/setup.py.in ${PYTHON_SRC} ${PYTHON_HDR})
 		add_dependencies(pythonmodule audaspace)

extern/audaspace/bindings/python/setup.py.in

@@ -8,20 +8,20 @@ import numpy
 from distutils.core import setup, Extension
 
 if len(sys.argv) > 2 and sys.argv[1] == '--build-docs':
-    import subprocess
+	import subprocess
-    from distutils.core import Distribution
+	from distutils.core import Distribution
-    from distutils.command.build import build
+	from distutils.command.build import build
 
-    dist = Distribution()
+	dist = Distribution()
-    cmd = build(dist)
+	cmd = build(dist)
-    cmd.finalize_options()
+	cmd.finalize_options()
-    #print(cmd.build_platlib)
+	#print(cmd.build_platlib)
 
-    os.environ['PYTHONPATH'] = os.path.join(os.getcwd(), cmd.build_platlib)
+	os.environ['PYTHONPATH'] = os.path.join(os.getcwd(), cmd.build_platlib)
-    os.environ['LD_LIBRARY_PATH'] = os.getcwd()
+	os.environ['LD_LIBRARY_PATH'] = os.getcwd()
 
-    ret = subprocess.call(sys.argv[2:])
+	ret = subprocess.call(sys.argv[2:])
-    sys.exit(ret)
+	sys.exit(ret)
 
 
 # the following line is not working due to https://bugs.python.org/issue9023
@@ -43,8 +43,7 @@ audaspace = Extension(
                       library_dirs = ['.', 'Release', 'Debug'],
                       language = 'c++',
                       extra_compile_args = extra_args,
-                      define_macros = [('WITH_CONVOLUTION', None)] if '@WITH_FFTW@' == 'ON' else [],
+					  sources = [os.path.join(source_directory, file) for file in ['PyAPI.cpp', 'PyDevice.cpp', 'PyHandle.cpp', 'PySound.cpp', 'PySequenceEntry.cpp', 'PySequence.cpp', 'PyPlaybackManager.cpp', 'PyDynamicMusic.cpp', 'PyThreadPool.cpp', 'PySource.cpp'] + (['PyImpulseResponse.cpp', 'PyHRTF.cpp'] if '@WITH_FFTW@' == 'ON' else [])]
-                      sources = [os.path.join(source_directory, file) for file in ['PyAPI.cpp', 'PyDevice.cpp', 'PyHandle.cpp', 'PySound.cpp', 'PySequenceEntry.cpp', 'PySequence.cpp', 'PyPlaybackManager.cpp', 'PyDynamicMusic.cpp', 'PyThreadPool.cpp', 'PySource.cpp'] + (['PyImpulseResponse.cpp', 'PyHRTF.cpp'] if '@WITH_FFTW@' == 'ON' else [])]
 )
 
 setup(
@@ -57,6 +56,6 @@ setup(
       license = 'Apache License 2.0',
       long_description = codecs.open(os.path.join(source_directory, '../../README.md'), 'r', 'utf-8').read(),
       ext_modules = [audaspace],
-      headers = [os.path.join(source_directory, file) for file in ['PyAPI.h', 'PyDevice.h', 'PyHandle.h', 'PySound.h', 'PySequenceEntry.h', 'PySequence.h', 'PyPlaybackManager.h', 'PyDynamicMusic.h', 'PyThreadPool.h', 'PySource.h'] + (['PyImpulseResponse.h', 'PyHRTF.h'] if '@WITH_FFTW@' == 'ON' else [])] + ['Audaspace.h']
+	  headers = [os.path.join(source_directory, file) for file in ['PyAPI.h', 'PyDevice.h', 'PyHandle.h', 'PySound.h', 'PySequenceEntry.h', 'PySequence.h', 'PyPlaybackManager.h', 'PyDynamicMusic.h', 'PyThreadPool.h', 'PySource.h'] + (['PyImpulseResponse.h', 'PyHRTF.h'] if '@WITH_FFTW@' == 'ON' else [])] + ['Audaspace.h']
 )
 

extern/audaspace/plugins/wasapi/WASAPIDevice.cpp

@@ -95,13 +95,6 @@ void WASAPIDevice::runMixingThread()
 				sleep_duration = std::chrono::milliseconds(buffer_size * 1000 / int(m_specs.rate) / 2);
 			}
 
-			if(m_default_device_changed)
-			{
-				m_default_device_changed = false;
-				result = AUDCLNT_E_DEVICE_INVALIDATED;
-				goto stop_thread;
-			}
-
 			if(FAILED(result = m_audio_client->GetCurrentPadding(&padding)))
 				goto stop_thread;
 
@@ -303,78 +296,13 @@ bool WASAPIDevice::setupDevice(DeviceSpecs &specs)
 	return true;
 }
 
-ULONG WASAPIDevice::AddRef()
-{
-	return InterlockedIncrement(&m_reference_count);
-}
-
-ULONG WASAPIDevice::Release()
-{
-	ULONG reference_count = InterlockedDecrement(&m_reference_count);
-
-	if(0 == reference_count)
-		delete this;
-
-	return reference_count;
-}
-
-HRESULT WASAPIDevice::QueryInterface(REFIID riid, void **ppvObject)
-{
-	if(riid == __uuidof(IMMNotificationClient))
-	{
-		*ppvObject = reinterpret_cast<IMMNotificationClient*>(this);
-		AddRef();
-	}
-	else if(riid == IID_IUnknown)
-	{
-		*ppvObject = reinterpret_cast<IUnknown*>(this);
-		AddRef();
-	}
-	else
-	{
-		*ppvObject = nullptr;
-		return E_NOINTERFACE;
-	}
-
-	return S_OK;
-}
-
-HRESULT WASAPIDevice::OnDeviceStateChanged(LPCWSTR pwstrDeviceId, DWORD dwNewState)
-{
-	return S_OK;
-}
-
-HRESULT WASAPIDevice::OnDeviceAdded(LPCWSTR pwstrDeviceId)
-{
-	return S_OK;
-}
-
-HRESULT WASAPIDevice::OnDeviceRemoved(LPCWSTR pwstrDeviceId)
-{
-	return S_OK;
-}
-
-HRESULT WASAPIDevice::OnDefaultDeviceChanged(EDataFlow flow, ERole role, LPCWSTR pwstrDeviceId)
-{
-	if(flow != EDataFlow::eCapture)
-		m_default_device_changed = true;
-
-	return S_OK;
-}
-
-HRESULT WASAPIDevice::OnPropertyValueChanged(LPCWSTR pwstrDeviceId, const PROPERTYKEY key)
-{
-	return S_OK;
-}
-
 WASAPIDevice::WASAPIDevice(DeviceSpecs specs, int buffersize) :
 	m_buffersize(buffersize),
 	m_imm_device_enumerator(nullptr),
 	m_imm_device(nullptr),
 	m_audio_client(nullptr),
-	m_wave_format_extensible({}),
+
-	m_default_device_changed(false),
+	m_wave_format_extensible({})
-	m_reference_count(1)
 {
 	// initialize COM if it hasn't happened yet
 	CoInitializeEx(nullptr, COINIT_MULTITHREADED);
@@ -399,8 +327,6 @@ WASAPIDevice::WASAPIDevice(DeviceSpecs specs, int buffersize) :
 
 	create();
 
-	m_imm_device_enumerator->RegisterEndpointNotificationCallback(this);
-
 	return;
 
 	error:
@@ -414,8 +340,6 @@ WASAPIDevice::~WASAPIDevice()
 {
 	stopMixingThread();
 
-	m_imm_device_enumerator->UnregisterEndpointNotificationCallback(this);
-
 	SafeRelease(&m_audio_client);
 	SafeRelease(&m_imm_device);
 	SafeRelease(&m_imm_device_enumerator);

extern/audaspace/plugins/wasapi/WASAPIDevice.h

@@ -40,7 +40,7 @@ AUD_NAMESPACE_BEGIN
 /**
  * This device plays back through WASAPI, the Windows audio API.
  */
-class AUD_PLUGIN_API WASAPIDevice : IMMNotificationClient, public ThreadedDevice
+class AUD_PLUGIN_API WASAPIDevice : public ThreadedDevice
 {
 private:
 	int m_buffersize;
@@ -48,8 +48,6 @@ private:
 	IMMDevice* m_imm_device;
 	IAudioClient* m_audio_client;
 	WAVEFORMATEXTENSIBLE m_wave_format_extensible;
-	bool m_default_device_changed;
-	LONG m_reference_count;
 
 	AUD_LOCAL HRESULT setupRenderClient(IAudioRenderClient*& render_client, UINT32& buffer_size);
 
@@ -60,17 +58,6 @@ private:
 
 	AUD_LOCAL bool setupDevice(DeviceSpecs& specs);
 
-	// IUnknown implementation
-	ULONG STDMETHODCALLTYPE AddRef();
-	ULONG STDMETHODCALLTYPE Release();
-	HRESULT STDMETHODCALLTYPE QueryInterface(REFIID riid, void **ppvObject);
-	// IMMNotificationClient implementation
-	HRESULT STDMETHODCALLTYPE OnDeviceStateChanged(LPCWSTR pwstrDeviceId, DWORD dwNewState);
-	HRESULT STDMETHODCALLTYPE OnDeviceAdded(LPCWSTR pwstrDeviceId);
-	HRESULT STDMETHODCALLTYPE OnDeviceRemoved(LPCWSTR pwstrDeviceId);
-	HRESULT STDMETHODCALLTYPE OnDefaultDeviceChanged(EDataFlow flow, ERole role, LPCWSTR pwstrDeviceId);
-	HRESULT STDMETHODCALLTYPE OnPropertyValueChanged(LPCWSTR pwstrDeviceId, const PROPERTYKEY key);
-
 	// delete copy constructor and operator=
 	WASAPIDevice(const WASAPIDevice&) = delete;
 	WASAPIDevice& operator=(const WASAPIDevice&) = delete;

extern/hipew/src/hipew.c

@@ -257,7 +257,7 @@ static int hipewHipInit(void) {
 #endif
   static int initialized = 0;
   static int result = 0;
-  int error;
+  int error, driver_version;
 
   if (initialized) {
     return result;
@@ -565,6 +565,8 @@ int hipewCompilerVersion(void) {
   const char *path = hipewCompilerPath();
   const char *marker = "Hip compilation tools, release ";
   FILE *pipe;
+  int major, minor;
+  char *versionstr;
   char buf[128];
   char output[65536] = "\0";
   char command[65536] = "\0";

intern/CMakeLists.txt

@@ -25,6 +25,7 @@ add_subdirectory(ghost)
 add_subdirectory(guardedalloc)
 add_subdirectory(libmv)
 add_subdirectory(memutil)
+add_subdirectory(numaapi)
 add_subdirectory(opencolorio)
 add_subdirectory(opensubdiv)
 add_subdirectory(mikktspace)

intern/cycles/app/CMakeLists.txt

@@ -51,6 +51,8 @@ list(APPEND LIBRARIES ${CYCLES_GL_LIBRARIES})
 
 # Common configuration.
 
+cycles_link_directories()
+
 add_definitions(${GL_DEFINITIONS})
 
 include_directories(${INC})

intern/cycles/app/cycles_standalone.cpp

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

intern/cycles/blender/addon/engine.py

@@ -60,8 +60,9 @@ def init():
 
     path = os.path.dirname(__file__)
     user_path = os.path.dirname(os.path.abspath(bpy.utils.user_resource('CONFIG', path='')))
+    temp_path = bpy.app.tempdir
 
-    _cycles.init(path, user_path, bpy.app.background)
+    _cycles.init(path, user_path, temp_path, bpy.app.background)
     _parse_command_line()
 
 

intern/cycles/blender/addon/properties.py

@@ -667,11 +667,6 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         description="Use special type BVH optimized for hair (uses more ram but renders faster)",
         default=True,
     )
-    debug_use_compact_bvh: BoolProperty(
-        name="Use Compact BVH",
-        description="Use compact BVH structure (uses less ram but renders slower)",
-        default=True,
-    )
     debug_bvh_time_steps: IntProperty(
         name="BVH Time Steps",
         description="Split BVH primitives by this number of time steps to speed up render time in cost of memory",
@@ -807,7 +802,7 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         name="Tile Size",
         default=2048,
         description="",
-        min=8, max=8192,
+        min=8, max=16384,
     )
 
     # Various fine-tuning debug flags
@@ -1452,19 +1447,6 @@ class CyclesPreferences(bpy.types.AddonPreferences):
                     num += 1
         return num
 
-    def has_multi_device(self):
-        import _cycles
-        compute_device_type = self.get_compute_device_type()
-        device_list = _cycles.available_devices(compute_device_type)
-        for device in device_list:
-            if device[1] == compute_device_type:
-                continue
-            for dev in self.devices:
-                if dev.use and dev.id == device[2]:
-                    return True
-
-        return False
-
     def has_active_device(self):
         return self.get_num_gpu_devices() > 0
 

intern/cycles/blender/addon/ui.py

@@ -118,11 +118,11 @@ def use_optix(context):
 
     return (get_device_type(context) == 'OPTIX' and cscene.device == 'GPU')
 
-def use_multi_device(context):
+
+def use_sample_all_lights(context):
     cscene = context.scene.cycles
-    if cscene.device != 'GPU':
+
-        return False
+    return cscene.sample_all_lights_direct or cscene.sample_all_lights_indirect
-    return context.preferences.addons[__package__].preferences.has_multi_device()
 
 
 def show_device_active(context):
@@ -667,10 +667,6 @@ class CYCLES_RENDER_PT_performance_acceleration_structure(CyclesButtonsPanel, Pa
     bl_label = "Acceleration Structure"
     bl_parent_id = "CYCLES_RENDER_PT_performance"
 
-    @classmethod
-    def poll(cls, context):
-        return not use_optix(context) or has_multi_device(context)
-
     def draw(self, context):
         import _cycles
 
@@ -683,33 +679,21 @@ class CYCLES_RENDER_PT_performance_acceleration_structure(CyclesButtonsPanel, Pa
 
         col = layout.column()
 
-        use_embree = _cycles.with_embree
+        use_embree = False
-
         if use_cpu(context):
-            col.prop(cscene, "debug_use_spatial_splits")
+            use_embree = _cycles.with_embree
-            if use_embree:
+            if not use_embree:
-                col.prop(cscene, "debug_use_compact_bvh")
-            else:
-                sub = col.column()
-                sub.active = not cscene.debug_use_spatial_splits
-                sub.prop(cscene, "debug_bvh_time_steps")
-
-                col.prop(cscene, "debug_use_hair_bvh")
-
                 sub = col.column(align=True)
                 sub.label(text="Cycles built without Embree support")
                 sub.label(text="CPU raytracing performance will be poor")
-        else:
-            col.prop(cscene, "debug_use_spatial_splits")
-            sub = col.column()
-            sub.active = not cscene.debug_use_spatial_splits
-            sub.prop(cscene, "debug_bvh_time_steps")
 
-            col.prop(cscene, "debug_use_hair_bvh")
+        col.prop(cscene, "debug_use_spatial_splits")
-
+        sub = col.column()
-            # CPU is used in addition to a GPU
+        sub.active = not use_embree
-            if use_multi_device(context) and use_embree:
+        sub.prop(cscene, "debug_use_hair_bvh")
-                col.prop(cscene, "debug_use_compact_bvh")
+        sub = col.column()
+        sub.active = not cscene.debug_use_spatial_splits and not use_embree
+        sub.prop(cscene, "debug_bvh_time_steps")
 
 
 class CYCLES_RENDER_PT_performance_final_render(CyclesButtonsPanel, Panel):
@@ -1819,45 +1803,18 @@ class CYCLES_RENDER_PT_bake_output(CyclesButtonsPanel, Panel):
         rd = scene.render
 
         if rd.use_bake_multires:
+            layout.prop(rd, "bake_margin")
             layout.prop(rd, "use_bake_clear", text="Clear Image")
+
             if rd.bake_type == 'DISPLACEMENT':
                 layout.prop(rd, "use_bake_lores_mesh")
         else:
             layout.prop(cbk, "target")
+
             if cbk.target == 'IMAGE_TEXTURES':
+                layout.prop(cbk, "margin")
                 layout.prop(cbk, "use_clear", text="Clear Image")
 
-class CYCLES_RENDER_PT_bake_output_margin(CyclesButtonsPanel, Panel):
-    bl_label = "Margin"
-    bl_context = "render"
-    bl_parent_id = "CYCLES_RENDER_PT_bake_output"
-    COMPAT_ENGINES = {'CYCLES'}
-
-    @classmethod
-    def poll(cls, context):
-        scene = context.scene
-        cbk = scene.render.bake
-        return cbk.target == 'IMAGE_TEXTURES'
-
-    def draw(self, context):
-        layout = self.layout
-        layout.use_property_split = True
-        layout.use_property_decorate = False  # No animation.
-
-        scene = context.scene
-        cscene = scene.cycles
-        cbk = scene.render.bake
-        rd = scene.render
-
-        if rd.use_bake_multires:
-            layout.prop(rd, "bake_margin_type", text="Type")
-            layout.prop(rd, "bake_margin", text="Size")
-        else:
-            if cbk.target == 'IMAGE_TEXTURES':
-                layout.prop(cbk, "margin_type", text="Type")
-                layout.prop(cbk, "margin", text="Size")
-
-
 
 class CYCLES_RENDER_PT_debug(CyclesDebugButtonsPanel, Panel):
     bl_label = "Debug"
@@ -2226,7 +2183,6 @@ classes = (
     CYCLES_RENDER_PT_bake_influence,
     CYCLES_RENDER_PT_bake_selected_to_active,
     CYCLES_RENDER_PT_bake_output,
-    CYCLES_RENDER_PT_bake_output_margin,
     CYCLES_RENDER_PT_debug,
     node_panel(CYCLES_MATERIAL_PT_settings),
     node_panel(CYCLES_MATERIAL_PT_settings_surface),

intern/cycles/blender/display_driver.cpp

@@ -272,300 +272,12 @@ uint BlenderDisplaySpaceShader::get_shader_program()
   return shader_program_;
 }
 
-/* --------------------------------------------------------------------
- * DrawTile.
- */
-
-/* Higher level representation of a texture from the graphics library. */
-class GLTexture {
- public:
-  /* Global counter for all allocated OpenGL textures used by instances of this class. */
-  static inline std::atomic<int> num_used = 0;
-
-  GLTexture() = default;
-
-  ~GLTexture()
-  {
-    assert(gl_id == 0);
-  }
-
-  GLTexture(const GLTexture &other) = delete;
-  GLTexture &operator=(GLTexture &other) = delete;
-
-  GLTexture(GLTexture &&other) noexcept
-      : gl_id(other.gl_id), width(other.width), height(other.height)
-  {
-    other.reset();
-  }
-
-  GLTexture &operator=(GLTexture &&other)
-  {
-    if (this == &other) {
-      return *this;
-    }
-
-    gl_id = other.gl_id;
-    width = other.width;
-    height = other.height;
-
-    other.reset();
-
-    return *this;
-  }
-
-  bool gl_resources_ensure()
-  {
-    if (gl_id) {
-      return true;
-    }
-
-    /* Create texture. */
-    glGenTextures(1, &gl_id);
-    if (!gl_id) {
-      LOG(ERROR) << "Error creating texture.";
-      return false;
-    }
-
-    /* Configure the texture. */
-    glActiveTexture(GL_TEXTURE0);
-    glBindTexture(GL_TEXTURE_2D, gl_id);
-
-    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
-    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
-
-    /* Clamp to edge so that precision issues when zoomed out (which forces linear interpolation)
-     * does not cause unwanted repetition. */
-    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
-    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
-
-    glBindTexture(GL_TEXTURE_2D, 0);
-
-    ++num_used;
-
-    return true;
-  }
-
-  void gl_resources_destroy()
-  {
-    if (!gl_id) {
-      return;
-    }
-
-    glDeleteTextures(1, &gl_id);
-
-    reset();
-
-    --num_used;
-  }
-
-  /* OpenGL resource IDs of the texture.
-   *
-   * NOTE: Allocated on the render engine's context. */
-  uint gl_id = 0;
-
-  /* Dimensions of the texture in pixels. */
-  int width = 0;
-  int height = 0;
-
- protected:
-  void reset()
-  {
-    gl_id = 0;
-    width = 0;
-    height = 0;
-  }
-};
-
-/* Higher level representation of a Pixel Buffer Object (PBO) from the graphics library. */
-class GLPixelBufferObject {
- public:
-  /* Global counter for all allocated OpenGL PBOs used by instances of this class. */
-  static inline std::atomic<int> num_used = 0;
-
-  GLPixelBufferObject() = default;
-
-  ~GLPixelBufferObject()
-  {
-    assert(gl_id == 0);
-  }
-
-  GLPixelBufferObject(const GLPixelBufferObject &other) = delete;
-  GLPixelBufferObject &operator=(GLPixelBufferObject &other) = delete;
-
-  GLPixelBufferObject(GLPixelBufferObject &&other) noexcept
-      : gl_id(other.gl_id), width(other.width), height(other.height)
-  {
-    other.reset();
-  }
-
-  GLPixelBufferObject &operator=(GLPixelBufferObject &&other)
-  {
-    if (this == &other) {
-      return *this;
-    }
-
-    gl_id = other.gl_id;
-    width = other.width;
-    height = other.height;
-
-    other.reset();
-
-    return *this;
-  }
-
-  bool gl_resources_ensure()
-  {
-    if (gl_id) {
-      return true;
-    }
-
-    glGenBuffers(1, &gl_id);
-    if (!gl_id) {
-      LOG(ERROR) << "Error creating texture pixel buffer object.";
-      return false;
-    }
-
-    ++num_used;
-
-    return true;
-  }
-
-  void gl_resources_destroy()
-  {
-    if (!gl_id) {
-      return;
-    }
-
-    glDeleteBuffers(1, &gl_id);
-
-    reset();
-
-    --num_used;
-  }
-
-  /* OpenGL resource IDs of the PBO.
-   *
-   * NOTE: Allocated on the render engine's context. */
-  uint gl_id = 0;
-
-  /* Dimensions of the PBO. */
-  int width = 0;
-  int height = 0;
-
- protected:
-  void reset()
-  {
-    gl_id = 0;
-    width = 0;
-    height = 0;
-  }
-};
-
-class DrawTile {
- public:
-  DrawTile() = default;
-  ~DrawTile() = default;
-
-  DrawTile(const DrawTile &other) = delete;
-  DrawTile &operator=(const DrawTile &other) = delete;
-
-  DrawTile(DrawTile &&other) noexcept = default;
-
-  DrawTile &operator=(DrawTile &&other) = default;
-
-  bool gl_resources_ensure()
-  {
-    if (!texture.gl_resources_ensure()) {
-      gl_resources_destroy();
-      return false;
-    }
-
-    if (!gl_vertex_buffer) {
-      glGenBuffers(1, &gl_vertex_buffer);
-      if (!gl_vertex_buffer) {
-        LOG(ERROR) << "Error allocating tile VBO.";
-        gl_resources_destroy();
-        return false;
-      }
-    }
-
-    return true;
-  }
-
-  void gl_resources_destroy()
-  {
-    texture.gl_resources_destroy();
-
-    if (gl_vertex_buffer) {
-      glDeleteBuffers(1, &gl_vertex_buffer);
-      gl_vertex_buffer = 0;
-    }
-  }
-
-  inline bool ready_to_draw() const
-  {
-    return texture.gl_id != 0;
-  }
-
-  /* Texture which contains pixels of the tile. */
-  GLTexture texture;
-
-  /* Display parameters the texture of this tile has been updated for. */
-  BlenderDisplayDriver::Params params;
-
-  /* OpenGL resources needed for drawing. */
-  uint gl_vertex_buffer = 0;
-};
-
-class DrawTileAndPBO {
- public:
-  bool gl_resources_ensure()
-  {
-    if (!tile.gl_resources_ensure() || !buffer_object.gl_resources_ensure()) {
-      gl_resources_destroy();
-      return false;
-    }
-
-    return true;
-  }
-
-  void gl_resources_destroy()
-  {
-    tile.gl_resources_destroy();
-    buffer_object.gl_resources_destroy();
-  }
-
-  DrawTile tile;
-  GLPixelBufferObject buffer_object;
-};
-
 /* --------------------------------------------------------------------
  * BlenderDisplayDriver.
  */
 
-struct BlenderDisplayDriver::Tiles {
-  /* Resources of a tile which is being currently rendered. */
-  DrawTileAndPBO current_tile;
-
-  /* All tiles which rendering is finished and which content will not be changed. */
-  struct {
-    vector<DrawTile> tiles;
-
-    void gl_resources_destroy_and_clear()
-    {
-      for (DrawTile &tile : tiles) {
-        tile.gl_resources_destroy();
-      }
-
-      tiles.clear();
-    }
-  } finished_tiles;
-};
-
 BlenderDisplayDriver::BlenderDisplayDriver(BL::RenderEngine &b_engine, BL::Scene &b_scene)
-    : b_engine_(b_engine),
+    : b_engine_(b_engine), display_shader_(BlenderDisplayShader::create(b_engine, b_scene))
-      display_shader_(BlenderDisplayShader::create(b_engine, b_scene)),
-      tiles_(make_unique<Tiles>())
 {
   /* Create context while on the main thread. */
   gl_context_create();
@@ -580,21 +292,6 @@ BlenderDisplayDriver::~BlenderDisplayDriver()
  * Update procedure.
  */
 
-void BlenderDisplayDriver::next_tile_begin()
-{
-  if (!tiles_->current_tile.tile.ready_to_draw()) {
-    LOG(ERROR)
-        << "Unexpectedly moving to the next tile without any data provided for current tile.";
-    return;
-  }
-
-  /* Moving to the next tile without giving render data for the current tile is not an expected
-   * situation. */
-  DCHECK(!need_clear_);
-
-  tiles_->finished_tiles.tiles.emplace_back(std::move(tiles_->current_tile.tile));
-}
-
 bool BlenderDisplayDriver::update_begin(const Params &params,
                                         int texture_width,
                                         int texture_height)
@@ -615,33 +312,24 @@ bool BlenderDisplayDriver::update_begin(const Params &params,
     glWaitSync((GLsync)gl_render_sync_, 0, GL_TIMEOUT_IGNORED);
   }
 
-  DrawTile &current_tile = tiles_->current_tile.tile;
+  if (!gl_texture_resources_ensure()) {
-  GLPixelBufferObject &current_tile_buffer_object = tiles_->current_tile.buffer_object;
-
-  /* Clear storage of all finished tiles when display clear is requested.
-   * Do it when new tile data is provided to handle the display clear flag in a single place.
-   * It also makes the logic reliable from the whether drawing did happen or not point of view. */
-  if (need_clear_) {
-    tiles_->finished_tiles.gl_resources_destroy_and_clear();
-    need_clear_ = false;
-  }
-
-  if (!tiles_->current_tile.gl_resources_ensure()) {
-    tiles_->current_tile.gl_resources_destroy();
     gl_context_disable();
     return false;
   }
 
   /* Update texture dimensions if needed. */
-  if (current_tile.texture.width != texture_width ||
+  if (texture_.width != texture_width || texture_.height != texture_height) {
-      current_tile.texture.height != texture_height) {
     glActiveTexture(GL_TEXTURE0);
-    glBindTexture(GL_TEXTURE_2D, current_tile.texture.gl_id);
+    glBindTexture(GL_TEXTURE_2D, texture_.gl_id);
     glTexImage2D(
         GL_TEXTURE_2D, 0, GL_RGBA16F, texture_width, texture_height, 0, GL_RGBA, GL_HALF_FLOAT, 0);
-    current_tile.texture.width = texture_width;
+    texture_.width = texture_width;
-    current_tile.texture.height = texture_height;
+    texture_.height = texture_height;
     glBindTexture(GL_TEXTURE_2D, 0);
+
+    /* Texture did change, and no pixel storage was provided. Tag for an explicit zeroing out to
+     * avoid undefined content. */
+    texture_.need_clear = true;
   }
 
   /* Update PBO dimensions if needed.
@@ -653,58 +341,29 @@ bool BlenderDisplayDriver::update_begin(const Params &params,
    * sending too much data to GPU when resolution divider is not 1. */
   /* TODO(sergey): Investigate whether keeping the PBO exact size of the texture makes non-interop
    * mode faster. */
-  const int buffer_width = params.size.x;
+  const int buffer_width = params.full_size.x;
-  const int buffer_height = params.size.y;
+  const int buffer_height = params.full_size.y;
-  if (current_tile_buffer_object.width != buffer_width ||
+  if (texture_.buffer_width != buffer_width || texture_.buffer_height != buffer_height) {
-      current_tile_buffer_object.height != buffer_height) {
     const size_t size_in_bytes = sizeof(half4) * buffer_width * buffer_height;
-    glBindBuffer(GL_PIXEL_UNPACK_BUFFER, current_tile_buffer_object.gl_id);
+    glBindBuffer(GL_PIXEL_UNPACK_BUFFER, texture_.gl_pbo_id);
     glBufferData(GL_PIXEL_UNPACK_BUFFER, size_in_bytes, 0, GL_DYNAMIC_DRAW);
     glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
 
-    current_tile_buffer_object.width = buffer_width;
+    texture_.buffer_width = buffer_width;
-    current_tile_buffer_object.height = buffer_height;
+    texture_.buffer_height = buffer_height;
   }
 
-  /* Store an updated parameters of the current tile.
+  /* New content will be provided to the texture in one way or another, so mark this in a
-   * In theory it is only needed once per update of the tile, but doing it on every update is
+   * centralized place. */
-   * the easiest and is not expensive. */
+  texture_.need_update = true;
-  tiles_->current_tile.tile.params = params;
+
+  texture_.params = params;
 
   return true;
 }
 
-static void update_tile_texture_pixels(const DrawTileAndPBO &tile)
-{
-  const GLTexture &texture = tile.tile.texture;
-
-  DCHECK_NE(tile.buffer_object.gl_id, 0);
-
-  glActiveTexture(GL_TEXTURE0);
-  glBindTexture(GL_TEXTURE_2D, texture.gl_id);
-  glBindBuffer(GL_PIXEL_UNPACK_BUFFER, tile.buffer_object.gl_id);
-
-  glTexSubImage2D(
-      GL_TEXTURE_2D, 0, 0, 0, texture.width, texture.height, GL_RGBA, GL_HALF_FLOAT, 0);
-
-  glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
-  glBindTexture(GL_TEXTURE_2D, 0);
-}
-
 void BlenderDisplayDriver::update_end()
 {
-  /* Unpack the PBO into the texture as soon as the new content is provided.
-   *
-   * This allows to ensure that the unpacking happens while resources like graphics interop (which
-   * lifetime is outside of control of the display driver) are still valid, as well as allows to
-   * move the tile from being current to finished immediately after this call.
-   *
-   * One concern with this approach is that if the update happens more often than drawing then
-   * doing the unpack here occupies GPU transfer for no good reason. However, the render scheduler
-   * takes care of ensuring updates don't happen that often. In regular applications redraw will
-   * happen much more often than this update. */
-  update_tile_texture_pixels(tiles_->current_tile);
-
   gl_upload_sync_ = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0);
   glFlush();
 
@@ -717,11 +376,7 @@ void BlenderDisplayDriver::update_end()
 
 half4 *BlenderDisplayDriver::map_texture_buffer()
 {
-  const uint pbo_gl_id = tiles_->current_tile.buffer_object.gl_id;
+  glBindBuffer(GL_PIXEL_UNPACK_BUFFER, texture_.gl_pbo_id);
-
-  DCHECK_NE(pbo_gl_id, 0);
-
-  glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo_gl_id);
 
   half4 *mapped_rgba_pixels = reinterpret_cast<half4 *>(
       glMapBuffer(GL_PIXEL_UNPACK_BUFFER, GL_WRITE_ONLY));
@@ -729,6 +384,15 @@ half4 *BlenderDisplayDriver::map_texture_buffer()
     LOG(ERROR) << "Error mapping BlenderDisplayDriver pixel buffer object.";
   }
 
+  if (texture_.need_clear) {
+    const int64_t texture_width = texture_.width;
+    const int64_t texture_height = texture_.height;
+    memset(reinterpret_cast<void *>(mapped_rgba_pixels),
+           0,
+           texture_width * texture_height * sizeof(half4));
+    texture_.need_clear = false;
+  }
+
   return mapped_rgba_pixels;
 }
 
@@ -747,9 +411,12 @@ BlenderDisplayDriver::GraphicsInterop BlenderDisplayDriver::graphics_interop_get
 {
   GraphicsInterop interop_dst;
 
-  interop_dst.buffer_width = tiles_->current_tile.buffer_object.width;
+  interop_dst.buffer_width = texture_.buffer_width;
-  interop_dst.buffer_height = tiles_->current_tile.buffer_object.height;
+  interop_dst.buffer_height = texture_.buffer_height;
-  interop_dst.opengl_pbo_id = tiles_->current_tile.buffer_object.gl_id;
+  interop_dst.opengl_pbo_id = texture_.gl_pbo_id;
+
+  interop_dst.need_clear = texture_.need_clear;
+  texture_.need_clear = false;
 
   return interop_dst;
 }
@@ -770,7 +437,7 @@ void BlenderDisplayDriver::graphics_interop_deactivate()
 
 void BlenderDisplayDriver::clear()
 {
-  need_clear_ = true;
+  texture_.need_clear = true;
 }
 
 void BlenderDisplayDriver::set_zoom(float zoom_x, float zoom_y)
@@ -778,155 +445,26 @@ void BlenderDisplayDriver::set_zoom(float zoom_x, float zoom_y)
   zoom_ = make_float2(zoom_x, zoom_y);
 }
 
-/* Update vertex buffer with new coordinates of vertex positions and texture coordinates.
- * This buffer is used to render texture in the viewport.
- *
- * NOTE: The buffer needs to be bound. */
-static void vertex_buffer_update(const DisplayDriver::Params &params)
-{
-  const int x = params.full_offset.x;
-  const int y = params.full_offset.y;
-
-  const int width = params.size.x;
-  const int height = params.size.y;
-
-  /* Invalidate old contents - avoids stalling if the buffer is still waiting in queue to be
-   * rendered. */
-  glBufferData(GL_ARRAY_BUFFER, 16 * sizeof(float), NULL, GL_STREAM_DRAW);
-
-  float *vpointer = reinterpret_cast<float *>(glMapBuffer(GL_ARRAY_BUFFER, GL_WRITE_ONLY));
-  if (!vpointer) {
-    return;
-  }
-
-  vpointer[0] = 0.0f;
-  vpointer[1] = 0.0f;
-  vpointer[2] = x;
-  vpointer[3] = y;
-
-  vpointer[4] = 1.0f;
-  vpointer[5] = 0.0f;
-  vpointer[6] = x + width;
-  vpointer[7] = y;
-
-  vpointer[8] = 1.0f;
-  vpointer[9] = 1.0f;
-  vpointer[10] = x + width;
-  vpointer[11] = y + height;
-
-  vpointer[12] = 0.0f;
-  vpointer[13] = 1.0f;
-  vpointer[14] = x;
-  vpointer[15] = y + height;
-
-  glUnmapBuffer(GL_ARRAY_BUFFER);
-}
-
-static void draw_tile(const float2 &zoom,
-                      const int texcoord_attribute,
-                      const int position_attribute,
-                      const DrawTile &draw_tile)
-{
-  if (!draw_tile.ready_to_draw()) {
-    return;
-  }
-
-  const GLTexture &texture = draw_tile.texture;
-
-  DCHECK_NE(texture.gl_id, 0);
-  DCHECK_NE(draw_tile.gl_vertex_buffer, 0);
-
-  glBindBuffer(GL_ARRAY_BUFFER, draw_tile.gl_vertex_buffer);
-
-  /* Draw at the parameters for which the texture has been updated for. This allows to always draw
-   * texture during bordered-rendered camera view without flickering. The validness of the display
-   * parameters for a texture is guaranteed by the initial "clear" state which makes drawing to
-   * have an early output.
-   *
-   * Such approach can cause some extra "jelly" effect during panning, but it is not more jelly
-   * than overlay of selected objects. Also, it's possible to redraw texture at an intersection of
-   * the texture draw parameters and the latest updated draw parameters (although, complexity of
-   * doing it might not worth it. */
-  vertex_buffer_update(draw_tile.params);
-
-  glBindTexture(GL_TEXTURE_2D, texture.gl_id);
-
-  /* Trick to keep sharp rendering without jagged edges on all GPUs.
-   *
-   * The idea here is to enforce driver to use linear interpolation when the image is not zoomed
-   * in.
-   * For the render result with a resolution divider in effect we always use nearest interpolation.
-   *
-   * Use explicit MIN assignment to make sure the driver does not have an undefined behavior at
-   * the zoom level 1. The MAG filter is always NEAREST. */
-  const float zoomed_width = draw_tile.params.size.x * zoom.x;
-  const float zoomed_height = draw_tile.params.size.y * zoom.y;
-  if (texture.width != draw_tile.params.size.x || texture.height != draw_tile.params.size.y) {
-    /* Resolution divider is different from 1, force nearest interpolation. */
-    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
-  }
-  else if (zoomed_width - draw_tile.params.size.x > 0.5f ||
-           zoomed_height - draw_tile.params.size.y > 0.5f) {
-    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
-  }
-  else {
-    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
-  }
-
-  glVertexAttribPointer(
-      texcoord_attribute, 2, GL_FLOAT, GL_FALSE, 4 * sizeof(float), (const GLvoid *)0);
-  glVertexAttribPointer(position_attribute,
-                        2,
-                        GL_FLOAT,
-                        GL_FALSE,
-                        4 * sizeof(float),
-                        (const GLvoid *)(sizeof(float) * 2));
-
-  glDrawArrays(GL_TRIANGLE_FAN, 0, 4);
-}
-
-void BlenderDisplayDriver::flush()
-{
-  /* This is called from the render thread that also calls update_begin/end, right before ending
-   * the render loop. We wait for any queued PBO and render commands to be done, before destroying
-   * the render thread and activating the context in the main thread to destroy resources.
-   *
-   * If we don't do this, the NVIDIA driver hangs for a few seconds for when ending 3D viewport
-   * rendering, for unknown reasons. This was found with NVIDIA driver version 470.73 and a Quadro
-   * RTX 6000 on Linux. */
-  if (!gl_context_enable()) {
-    return;
-  }
-
-  if (gl_upload_sync_) {
-    glWaitSync((GLsync)gl_upload_sync_, 0, GL_TIMEOUT_IGNORED);
-  }
-
-  if (gl_render_sync_) {
-    glWaitSync((GLsync)gl_render_sync_, 0, GL_TIMEOUT_IGNORED);
-  }
-
-  gl_context_disable();
-}
-
 void BlenderDisplayDriver::draw(const Params &params)
 {
   /* See do_update_begin() for why no locking is required here. */
   const bool transparent = true;  // TODO(sergey): Derive this from Film.
 
+  if (!gl_draw_resources_ensure()) {
+    return;
+  }
+
   if (use_gl_context_) {
     gl_context_mutex_.lock();
   }
 
-  if (need_clear_) {
+  if (texture_.need_clear) {
     /* Texture is requested to be cleared and was not yet cleared.
      *
      * Do early return which should be equivalent of drawing all-zero texture.
      * Watch out for the lock though so that the clear happening during update is properly
      * synchronized here. */
-    if (use_gl_context_) {
+    gl_context_mutex_.unlock();
-      gl_context_mutex_.unlock();
-    }
     return;
   }
 
@@ -939,37 +477,66 @@ void BlenderDisplayDriver::draw(const Params &params)
     glBlendFunc(GL_ONE, GL_ONE_MINUS_SRC_ALPHA);
   }
 
-  glActiveTexture(GL_TEXTURE0);
+  display_shader_->bind(params.full_size.x, params.full_size.y);
 
-  /* NOTE: The VAO is to be allocated on the drawing context as it is not shared across contexts.
+  glActiveTexture(GL_TEXTURE0);
-   * Simplest is to allocate it on every redraw so that it is possible to destroy it from a
+  glBindTexture(GL_TEXTURE_2D, texture_.gl_id);
-   * correct context. */
+
+  /* Trick to keep sharp rendering without jagged edges on all GPUs.
+   *
+   * The idea here is to enforce driver to use linear interpolation when the image is not zoomed
+   * in.
+   * For the render result with a resolution divider in effect we always use nearest interpolation.
+   *
+   * Use explicit MIN assignment to make sure the driver does not have an undefined behavior at
+   * the zoom level 1. The MAG filter is always NEAREST. */
+  const float zoomed_width = params.size.x * zoom_.x;
+  const float zoomed_height = params.size.y * zoom_.y;
+  if (texture_.width != params.size.x || texture_.height != params.size.y) {
+    /* Resolution divider is different from 1, force nearest interpolation. */
+    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
+  }
+  else if (zoomed_width - params.size.x > 0.5f || zoomed_height - params.size.y > 0.5f) {
+    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
+  }
+  else {
+    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
+  }
+
+  glBindBuffer(GL_ARRAY_BUFFER, vertex_buffer_);
+
+  texture_update_if_needed();
+  vertex_buffer_update(params);
+
+  /* TODO(sergey): Does it make sense/possible to cache/reuse the VAO? */
   GLuint vertex_array_object;
   glGenVertexArrays(1, &vertex_array_object);
   glBindVertexArray(vertex_array_object);
 
-  display_shader_->bind(params.full_size.x, params.full_size.y);
-
   const int texcoord_attribute = display_shader_->get_tex_coord_attrib_location();
   const int position_attribute = display_shader_->get_position_attrib_location();
 
   glEnableVertexAttribArray(texcoord_attribute);
   glEnableVertexAttribArray(position_attribute);
 
-  draw_tile(zoom_, texcoord_attribute, position_attribute, tiles_->current_tile.tile);
+  glVertexAttribPointer(
+      texcoord_attribute, 2, GL_FLOAT, GL_FALSE, 4 * sizeof(float), (const GLvoid *)0);
+  glVertexAttribPointer(position_attribute,
+                        2,
+                        GL_FLOAT,
+                        GL_FALSE,
+                        4 * sizeof(float),
+                        (const GLvoid *)(sizeof(float) * 2));
 
-  for (const DrawTile &tile : tiles_->finished_tiles.tiles) {
+  glDrawArrays(GL_TRIANGLE_FAN, 0, 4);
-    draw_tile(zoom_, texcoord_attribute, position_attribute, tile);
-  }
 
-  display_shader_->unbind();
-
-  glBindTexture(GL_TEXTURE_2D, 0);
-  glBindVertexArray(0);
   glBindBuffer(GL_ARRAY_BUFFER, 0);
+  glBindTexture(GL_TEXTURE_2D, 0);
 
   glDeleteVertexArrays(1, &vertex_array_object);
 
+  display_shader_->unbind();
+
   if (transparent) {
     glDisable(GL_BLEND);
   }
@@ -977,11 +544,6 @@ void BlenderDisplayDriver::draw(const Params &params)
   gl_render_sync_ = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0);
   glFlush();
 
-  if (VLOG_IS_ON(5)) {
-    VLOG(5) << "Number of textures: " << GLTexture::num_used;
-    VLOG(5) << "Number of PBOs: " << GLPixelBufferObject::num_used;
-  }
-
   if (use_gl_context_) {
     gl_context_mutex_.unlock();
   }
@@ -1056,16 +618,154 @@ void BlenderDisplayDriver::gl_context_dispose()
   }
 }
 
+bool BlenderDisplayDriver::gl_draw_resources_ensure()
+{
+  if (!texture_.gl_id) {
+    /* If there is no texture allocated, there is nothing to draw. Inform the draw call that it can
+     * can not continue. Note that this is not an unrecoverable error, so once the texture is known
+     * we will come back here and create all the GPU resources needed for draw. */
+    return false;
+  }
+
+  if (gl_draw_resource_creation_attempted_) {
+    return gl_draw_resources_created_;
+  }
+  gl_draw_resource_creation_attempted_ = true;
+
+  if (!vertex_buffer_) {
+    glGenBuffers(1, &vertex_buffer_);
+    if (!vertex_buffer_) {
+      LOG(ERROR) << "Error creating vertex buffer.";
+      return false;
+    }
+  }
+
+  gl_draw_resources_created_ = true;
+
+  return true;
+}
+
 void BlenderDisplayDriver::gl_resources_destroy()
 {
   gl_context_enable();
 
-  tiles_->current_tile.gl_resources_destroy();
+  if (vertex_buffer_ != 0) {
-  tiles_->finished_tiles.gl_resources_destroy_and_clear();
+    glDeleteBuffers(1, &vertex_buffer_);
+  }
+
+  if (texture_.gl_pbo_id) {
+    glDeleteBuffers(1, &texture_.gl_pbo_id);
+    texture_.gl_pbo_id = 0;
+  }
+
+  if (texture_.gl_id) {
+    glDeleteTextures(1, &texture_.gl_id);
+    texture_.gl_id = 0;
+  }
 
   gl_context_disable();
 
   gl_context_dispose();
 }
 
+bool BlenderDisplayDriver::gl_texture_resources_ensure()
+{
+  if (texture_.creation_attempted) {
+    return texture_.is_created;
+  }
+  texture_.creation_attempted = true;
+
+  DCHECK(!texture_.gl_id);
+  DCHECK(!texture_.gl_pbo_id);
+
+  /* Create texture. */
+  glGenTextures(1, &texture_.gl_id);
+  if (!texture_.gl_id) {
+    LOG(ERROR) << "Error creating texture.";
+    return false;
+  }
+
+  /* Configure the texture. */
+  glActiveTexture(GL_TEXTURE0);
+  glBindTexture(GL_TEXTURE_2D, texture_.gl_id);
+  glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
+  glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
+  glBindTexture(GL_TEXTURE_2D, 0);
+
+  /* Create PBO for the texture. */
+  glGenBuffers(1, &texture_.gl_pbo_id);
+  if (!texture_.gl_pbo_id) {
+    LOG(ERROR) << "Error creating texture pixel buffer object.";
+    return false;
+  }
+
+  /* Creation finished with a success. */
+  texture_.is_created = true;
+
+  return true;
+}
+
+void BlenderDisplayDriver::texture_update_if_needed()
+{
+  if (!texture_.need_update) {
+    return;
+  }
+
+  glBindBuffer(GL_PIXEL_UNPACK_BUFFER, texture_.gl_pbo_id);
+  glTexSubImage2D(
+      GL_TEXTURE_2D, 0, 0, 0, texture_.width, texture_.height, GL_RGBA, GL_HALF_FLOAT, 0);
+  glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
+
+  texture_.need_update = false;
+}
+
+void BlenderDisplayDriver::vertex_buffer_update(const Params & /*params*/)
+{
+  /* Draw at the parameters for which the texture has been updated for. This allows to always draw
+   * texture during bordered-rendered camera view without flickering. The validness of the display
+   * parameters for a texture is guaranteed by the initial "clear" state which makes drawing to
+   * have an early output.
+   *
+   * Such approach can cause some extra "jelly" effect during panning, but it is not more jelly
+   * than overlay of selected objects. Also, it's possible to redraw texture at an intersection of
+   * the texture draw parameters and the latest updated draw parameters (although, complexity of
+   * doing it might not worth it. */
+  const int x = texture_.params.full_offset.x;
+  const int y = texture_.params.full_offset.y;
+
+  const int width = texture_.params.size.x;
+  const int height = texture_.params.size.y;
+
+  /* Invalidate old contents - avoids stalling if the buffer is still waiting in queue to be
+   * rendered. */
+  glBufferData(GL_ARRAY_BUFFER, 16 * sizeof(float), NULL, GL_STREAM_DRAW);
+
+  float *vpointer = reinterpret_cast<float *>(glMapBuffer(GL_ARRAY_BUFFER, GL_WRITE_ONLY));
+  if (!vpointer) {
+    return;
+  }
+
+  vpointer[0] = 0.0f;
+  vpointer[1] = 0.0f;
+  vpointer[2] = x;
+  vpointer[3] = y;
+
+  vpointer[4] = 1.0f;
+  vpointer[5] = 0.0f;
+  vpointer[6] = x + width;
+  vpointer[7] = y;
+
+  vpointer[8] = 1.0f;
+  vpointer[9] = 1.0f;
+  vpointer[10] = x + width;
+  vpointer[11] = y + height;
+
+  vpointer[12] = 0.0f;
+  vpointer[13] = 1.0f;
+  vpointer[14] = x;
+  vpointer[15] = y + height;
+
+  glUnmapBuffer(GL_ARRAY_BUFFER);
+}
+
 CCL_NAMESPACE_END

intern/cycles/blender/display_driver.h

@@ -26,7 +26,6 @@
 
 #include "util/thread.h"
 #include "util/unique_ptr.h"
-#include "util/vector.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -113,8 +112,6 @@ class BlenderDisplayDriver : public DisplayDriver {
   void set_zoom(float zoom_x, float zoom_y);
 
  protected:
-  virtual void next_tile_begin() override;
-
   virtual bool update_begin(const Params &params, int texture_width, int texture_height) override;
   virtual void update_end() override;
 
@@ -125,17 +122,33 @@ class BlenderDisplayDriver : public DisplayDriver {
 
   virtual void draw(const Params &params) override;
 
-  virtual void flush() override;
-
   /* Helper function which allocates new GPU context. */
   void gl_context_create();
   bool gl_context_enable();
   void gl_context_disable();
   void gl_context_dispose();
 
+  /* Make sure texture is allocated and its initial configuration is performed. */
+  bool gl_texture_resources_ensure();
+
+  /* Ensure all runtime GPU resources needed for drawing are allocated.
+   * Returns true if all resources needed for drawing are available. */
+  bool gl_draw_resources_ensure();
+
   /* Destroy all GPU resources which are being used by this object. */
   void gl_resources_destroy();
 
+  /* Update GPU texture dimensions and content if needed (new pixel data was provided).
+   *
+   * NOTE: The texture needs to be bound. */
+  void texture_update_if_needed();
+
+  /* Update vertex buffer with new coordinates of vertex positions and texture coordinates.
+   * This buffer is used to render texture in the viewport.
+   *
+   * NOTE: The buffer needs to be bound. */
+  void vertex_buffer_update(const Params &params);
+
   BL::RenderEngine b_engine_;
 
   /* OpenGL context which is used the render engine doesn't have its own. */
@@ -146,14 +159,50 @@ class BlenderDisplayDriver : public DisplayDriver {
   /* Mutex used to guard the `gl_context_`. */
   thread_mutex gl_context_mutex_;
 
-  /* Content of the display is to be filled with zeroes. */
+  /* Texture which contains pixels of the render result. */
-  std::atomic<bool> need_clear_ = true;
+  struct {
+    /* Indicates whether texture creation was attempted and succeeded.
+     * Used to avoid multiple attempts of texture creation on GPU issues or GPU context
+     * misconfiguration. */
+    bool creation_attempted = false;
+    bool is_created = false;
+
+    /* OpenGL resource IDs of the texture itself and Pixel Buffer Object (PBO) used to write
+     * pixels to it.
+     *
+     * NOTE: Allocated on the engine's context. */
+    uint gl_id = 0;
+    uint gl_pbo_id = 0;
+
+    /* Is true when new data was written to the PBO, meaning, the texture might need to be resized
+     * and new data is to be uploaded to the GPU. */
+    bool need_update = false;
+
+    /* Content of the texture is to be filled with zeroes. */
+    std::atomic<bool> need_clear = true;
+
+    /* Dimensions of the texture in pixels. */
+    int width = 0;
+    int height = 0;
+
+    /* Dimensions of the underlying PBO. */
+    int buffer_width = 0;
+    int buffer_height = 0;
+
+    /* Display parameters the texture has been updated for. */
+    Params params;
+  } texture_;
 
   unique_ptr<BlenderDisplayShader> display_shader_;
 
-  /* Opaque storage for an internal state and data for tiles. */
+  /* Special track of whether GPU resources were attempted to be created, to avoid attempts of
-  struct Tiles;
+   * their re-creation on failure on every redraw. */
-  unique_ptr<Tiles> tiles_;
+  bool gl_draw_resource_creation_attempted_ = false;
+  bool gl_draw_resources_created_ = false;
+
+  /* Vertex buffer which hold vertices of a triangle fan which is textures with the texture
+   * holding the render result. */
+  uint vertex_buffer_ = 0;
 
   void *gl_render_sync_ = nullptr;
   void *gl_upload_sync_ = nullptr;

intern/cycles/blender/mesh.cpp

@@ -1071,15 +1071,7 @@ static void create_subd_mesh(Scene *scene,
 
   for (BL::MeshEdge &e : b_mesh.edges) {
     if (e.crease() != 0.0f) {
-      mesh->add_edge_crease(e.vertices()[0], e.vertices()[1], e.crease());
+      mesh->add_crease(e.vertices()[0], e.vertices()[1], e.crease());
-    }
-  }
-
-  for (BL::MeshVertexCreaseLayer &c : b_mesh.vertex_creases) {
-    for (int i = 0; i < c.data.length(); ++i) {
-      if (c.data[i].value() != 0.0f) {
-        mesh->add_vertex_crease(i, c.data[i].value());
-      }
     }
   }
 
@@ -1094,6 +1086,40 @@ static void create_subd_mesh(Scene *scene,
 
 /* Sync */
 
+/* Check whether some of "built-in" motion-related attributes are needed to be exported (includes
+ * things like velocity from cache modifier, fluid simulation).
+ *
+ * NOTE: This code is run prior to object motion blur initialization. so can not access properties
+ * set by `sync_object_motion_init()`. */
+static bool mesh_need_motion_attribute(BObjectInfo &b_ob_info, Scene *scene)
+{
+  const Scene::MotionType need_motion = scene->need_motion();
+  if (need_motion == Scene::MOTION_NONE) {
+    /* Simple case: neither motion pass nor motion blur is needed, no need in the motion related
+     * attributes. */
+    return false;
+  }
+
+  if (need_motion == Scene::MOTION_BLUR) {
+    /* A bit tricky and implicit case:
+     * - Motion blur is enabled in the scene, which implies specific number of time steps for
+     *   objects.
+     * - If the object has motion blur disabled on it, it will have 0 time steps.
+     * - Motion attribute expects non-zero time steps.
+     *
+     * Avoid adding motion attributes if the motion blur will enforce 0 motion steps. */
+    PointerRNA cobject = RNA_pointer_get(&b_ob_info.real_object.ptr, "cycles");
+    const bool use_motion = get_boolean(cobject, "use_motion_blur");
+    if (!use_motion) {
+      return false;
+    }
+  }
+
+  /* Motion pass which implies 3 motion steps, or motion blur which is not disabled on object
+   * level. */
+  return true;
+}
+
 void BlenderSync::sync_mesh(BL::Depsgraph b_depsgraph, BObjectInfo &b_ob_info, Mesh *mesh)
 {
   /* make a copy of the shaders as the caller in the main thread still need them for syncing the
@@ -1118,7 +1144,7 @@ void BlenderSync::sync_mesh(BL::Depsgraph b_depsgraph, BObjectInfo &b_ob_info, M
 
     if (b_mesh) {
       /* Motion blur attribute is relative to seconds, we need it relative to frames. */
-      const bool need_motion = object_need_motion_attribute(b_ob_info, scene);
+      const bool need_motion = mesh_need_motion_attribute(b_ob_info, scene);
       const float motion_scale = (need_motion) ?
                                      scene->motion_shutter_time() /
                                          (b_scene.render().fps() / b_scene.render().fps_base()) :

intern/cycles/blender/object.cpp

@@ -529,17 +529,6 @@ void BlenderSync::sync_procedural(BL::Object &b_ob,
   string absolute_path = blender_absolute_path(b_data, b_ob, b_mesh_cache.cache_file().filepath());
   procedural->set_filepath(ustring(absolute_path));
 
-  array<ustring> layers;
-  for (BL::CacheFileLayer &layer : cache_file.layers) {
-    if (layer.hide_layer()) {
-      continue;
-    }
-
-    absolute_path = blender_absolute_path(b_data, b_ob, layer.filepath());
-    layers.push_back_slow(ustring(absolute_path));
-  }
-  procedural->set_layers(layers);
-
   procedural->set_scale(cache_file.scale());
 
   procedural->set_use_prefetch(cache_file.use_prefetch());

intern/cycles/blender/output_driver.cpp

@@ -51,6 +51,8 @@ bool BlenderOutputDriver::read_render_tile(const Tile &tile)
 
   BL::RenderLayer b_rlay = *b_single_rlay;
 
+  vector<float> pixels(tile.size.x * tile.size.y * 4);
+
   /* Copy each pass.
    * TODO:copy only the required ones for better performance? */
   for (BL::RenderPass &b_pass : b_rlay.passes) {
@@ -107,7 +109,7 @@ void BlenderOutputDriver::write_render_tile(const Tile &tile)
 
   BL::RenderLayer b_rlay = *b_single_rlay;
 
-  vector<float> pixels(static_cast<size_t>(tile.size.x) * tile.size.y * 4);
+  vector<float> pixels(tile.size.x * tile.size.y * 4);
 
   /* Copy each pass. */
   for (BL::RenderPass &b_pass : b_rlay.passes) {
@@ -118,7 +120,7 @@ void BlenderOutputDriver::write_render_tile(const Tile &tile)
     b_pass.rect(&pixels[0]);
   }
 
-  b_engine_.end_result(b_rr, false, false, true);
+  b_engine_.end_result(b_rr, true, false, true);
 }
 
 CCL_NAMESPACE_END

intern/cycles/blender/pointcloud.cpp

@@ -37,52 +37,12 @@ static void fill_generic_attribute(BL::PointCloud &b_pointcloud,
   }
 }
 
-static void attr_create_motion(PointCloud *pointcloud,
+static void copy_attributes(PointCloud *pointcloud, BL::PointCloud b_pointcloud)
-                               BL::Attribute &b_attribute,
-                               const float motion_scale)
-{
-  if (!(b_attribute.domain() == BL::Attribute::domain_POINT) &&
-      (b_attribute.data_type() == BL::Attribute::data_type_FLOAT_VECTOR)) {
-    return;
-  }
-
-  BL::FloatVectorAttribute b_vector_attribute(b_attribute);
-  const int num_points = pointcloud->get_points().size();
-
-  /* Find or add attribute */
-  float3 *P = &pointcloud->get_points()[0];
-  Attribute *attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
-
-  if (!attr_mP) {
-    attr_mP = pointcloud->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
-  }
-
-  /* Only export previous and next frame, we don't have any in between data. */
-  float motion_times[2] = {-1.0f, 1.0f};
-  for (int step = 0; step < 2; step++) {
-    const float relative_time = motion_times[step] * 0.5f * motion_scale;
-    float3 *mP = attr_mP->data_float3() + step * num_points;
-
-    for (int i = 0; i < num_points; i++) {
-      mP[i] = P[i] + get_float3(b_vector_attribute.data[i].vector()) * relative_time;
-    }
-  }
-}
-
-static void copy_attributes(PointCloud *pointcloud,
-                            BL::PointCloud b_pointcloud,
-                            const bool need_motion,
-                            const float motion_scale)
 {
   AttributeSet &attributes = pointcloud->attributes;
-  static const ustring u_velocity("velocity");
   for (BL::Attribute &b_attribute : b_pointcloud.attributes) {
     const ustring name{b_attribute.name().c_str()};
 
-    if (need_motion && name == u_velocity) {
-      attr_create_motion(pointcloud, b_attribute, motion_scale);
-    }
-
     if (attributes.find(name)) {
       continue;
     }
@@ -151,11 +111,7 @@ static void copy_attributes(PointCloud *pointcloud,
   }
 }
 
-static void export_pointcloud(Scene *scene,
+static void export_pointcloud(Scene *scene, PointCloud *pointcloud, BL::PointCloud b_pointcloud)
-                              PointCloud *pointcloud,
-                              BL::PointCloud b_pointcloud,
-                              const bool need_motion,
-                              const float motion_scale)
 {
   /* TODO: optimize so we can straight memcpy arrays from Blender? */
 
@@ -185,7 +141,7 @@ static void export_pointcloud(Scene *scene,
   }
 
   /* Export attributes */
-  copy_attributes(pointcloud, b_pointcloud, need_motion, motion_scale);
+  copy_attributes(pointcloud, b_pointcloud);
 }
 
 static void export_pointcloud_motion(PointCloud *pointcloud,
@@ -237,7 +193,7 @@ static void export_pointcloud_motion(PointCloud *pointcloud,
   }
 
   /* Export attributes */
-  copy_attributes(pointcloud, b_pointcloud, false, 0.0f);
+  copy_attributes(pointcloud, b_pointcloud);
 }
 
 void BlenderSync::sync_pointcloud(PointCloud *pointcloud, BObjectInfo &b_ob_info)
@@ -251,13 +207,7 @@ void BlenderSync::sync_pointcloud(PointCloud *pointcloud, BObjectInfo &b_ob_info
 
   /* TODO: add option to filter out points in the view layer. */
   BL::PointCloud b_pointcloud(b_ob_info.object_data);
-  /* Motion blur attribute is relative to seconds, we need it relative to frames. */
+  export_pointcloud(scene, &new_pointcloud, b_pointcloud);
-  const bool need_motion = object_need_motion_attribute(b_ob_info, scene);
-  const float motion_scale = (need_motion) ?
-                                 scene->motion_shutter_time() /
-                                     (b_scene.render().fps() / b_scene.render().fps_base()) :
-                                 0.0f;
-  export_pointcloud(scene, &new_pointcloud, b_pointcloud, need_motion, motion_scale);
 
   /* update original sockets */
   for (const SocketType &socket : new_pointcloud.type->inputs) {

intern/cycles/blender/python.cpp

@@ -138,18 +138,20 @@ static const char *PyC_UnicodeAsByte(PyObject *py_str, PyObject **coerce)
 
 static PyObject *init_func(PyObject * /*self*/, PyObject *args)
 {
-  PyObject *path, *user_path;
+  PyObject *path, *user_path, *temp_path;
   int headless;
 
-  if (!PyArg_ParseTuple(args, "OOi", &path, &user_path, &headless)) {
+  if (!PyArg_ParseTuple(args, "OOOi", &path, &user_path, &temp_path, &headless)) {
     return nullptr;
   }
 
-  PyObject *path_coerce = nullptr, *user_path_coerce = nullptr;
+  PyObject *path_coerce = nullptr, *user_path_coerce = nullptr, *temp_path_coerce = nullptr;
   path_init(PyC_UnicodeAsByte(path, &path_coerce),
-            PyC_UnicodeAsByte(user_path, &user_path_coerce));
+            PyC_UnicodeAsByte(user_path, &user_path_coerce),
+            PyC_UnicodeAsByte(temp_path, &temp_path_coerce));
   Py_XDECREF(path_coerce);
   Py_XDECREF(user_path_coerce);
+  Py_XDECREF(temp_path_coerce);
 
   BlenderSession::headless = headless;
 
@@ -733,20 +735,27 @@ static bool image_parse_filepaths(PyObject *pyfilepaths, vector<string> &filepat
 
 static PyObject *denoise_func(PyObject * /*self*/, PyObject *args, PyObject *keywords)
 {
+#if 1
+  (void)args;
+  (void)keywords;
+#else
   static const char *keyword_list[] = {
-      "preferences", "scene", "view_layer", "input", "output", NULL};
+      "preferences", "scene", "view_layer", "input", "output", "tile_size", "samples", NULL};
   PyObject *pypreferences, *pyscene, *pyviewlayer;
   PyObject *pyinput, *pyoutput = NULL;
+  int tile_size = 0, samples = 0;
 
   if (!PyArg_ParseTupleAndKeywords(args,
                                    keywords,
-                                   "OOOO|O",
+                                   "OOOO|Oii",
                                    (char **)keyword_list,
                                    &pypreferences,
                                    &pyscene,
                                    &pyviewlayer,
                                    &pyinput,
-                                   &pyoutput)) {
+                                   &pyoutput,
+                                   &tile_size,
+                                   &samples)) {
     return NULL;
   }
 
@@ -768,10 +777,14 @@ static PyObject *denoise_func(PyObject * /*self*/, PyObject *args, PyObject *key
                      &RNA_ViewLayer,
                      PyLong_AsVoidPtr(pyviewlayer),
                      &viewlayerptr);
-  BL::ViewLayer b_view_layer(viewlayerptr);
+  PointerRNA cviewlayer = RNA_pointer_get(&viewlayerptr, "cycles");
 
-  DenoiseParams params = BlenderSync::get_denoise_params(b_scene, b_view_layer, true);
+  DenoiseParams params;
-  params.use = true;
+  params.radius = get_int(cviewlayer, "denoising_radius");
+  params.strength = get_float(cviewlayer, "denoising_strength");
+  params.feature_strength = get_float(cviewlayer, "denoising_feature_strength");
+  params.relative_pca = get_boolean(cviewlayer, "denoising_relative_pca");
+  params.neighbor_frames = get_int(cviewlayer, "denoising_neighbor_frames");
 
   /* Parse file paths list. */
   vector<string> input, output;
@@ -799,15 +812,24 @@ static PyObject *denoise_func(PyObject * /*self*/, PyObject *args, PyObject *key
   }
 
   /* Create denoiser. */
-  DenoiserPipeline denoiser(device, params);
+  DenoiserPipeline denoiser(device);
+  denoiser.params = params;
   denoiser.input = input;
   denoiser.output = output;
 
+  if (tile_size > 0) {
+    denoiser.tile_size = make_int2(tile_size, tile_size);
+  }
+  if (samples > 0) {
+    denoiser.samples_override = samples;
+  }
+
   /* Run denoiser. */
   if (!denoiser.run()) {
     PyErr_SetString(PyExc_ValueError, denoiser.error.c_str());
     return NULL;
   }
+#endif
 
   Py_RETURN_NONE;
 }

intern/cycles/blender/session.cpp

@@ -502,15 +502,10 @@ void BlenderSession::render_frame_finish()
     path_remove(filename);
   }
 
-  /* Clear output driver. */
+  /* Clear driver. */
   session->set_output_driver(nullptr);
   session->full_buffer_written_cb = function_null;
 
-  /* The display driver holds OpenGL resources which belong to an OpenGL context held by the render
-   * engine on Blender side. Force destruction of those resources. */
-  display_driver_ = nullptr;
-  session->set_display_driver(nullptr);
-
   /* All the files are handled.
    * Clear the list so that this session can be re-used by Persistent Data. */
   full_buffer_files_.clear();

intern/cycles/blender/shader.cpp

@@ -689,9 +689,6 @@ static ShaderNode *add_node(Scene *scene,
   else if (b_node.is_a(&RNA_ShaderNodeHairInfo)) {
     node = graph->create_node<HairInfoNode>();
   }
-  else if (b_node.is_a(&RNA_ShaderNodePointInfo)) {
-    node = graph->create_node<PointInfoNode>();
-  }
   else if (b_node.is_a(&RNA_ShaderNodeVolumeInfo)) {
     node = graph->create_node<VolumeInfoNode>();
   }
@@ -779,7 +776,7 @@ static ShaderNode *add_node(Scene *scene,
       }
       else {
         ustring filename = ustring(
-            image_user_file_path(b_image_user, b_image, b_scene.frame_current()));
+            image_user_file_path(b_image_user, b_image, b_scene.frame_current(), true));
         image->set_filename(filename);
       }
     }
@@ -816,7 +813,7 @@ static ShaderNode *add_node(Scene *scene,
       }
       else {
         env->set_filename(
-            ustring(image_user_file_path(b_image_user, b_image, b_scene.frame_current())));
+            ustring(image_user_file_path(b_image_user, b_image, b_scene.frame_current(), false)));
       }
     }
     node = env;

intern/cycles/blender/sync.cpp

@@ -787,7 +787,6 @@ SceneParams BlenderSync::get_scene_params(BL::Scene &b_scene, bool background)
     params.bvh_type = BVH_TYPE_DYNAMIC;
 
   params.use_bvh_spatial_split = RNA_boolean_get(&cscene, "debug_use_spatial_splits");
-  params.use_bvh_compact_structure = RNA_boolean_get(&cscene, "debug_use_compact_bvh");
   params.use_bvh_unaligned_nodes = RNA_boolean_get(&cscene, "debug_use_hair_bvh");
   params.num_bvh_time_steps = RNA_int_get(&cscene, "debug_bvh_time_steps");
 
@@ -833,14 +832,6 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine &b_engine,
   SessionParams params;
   PointerRNA cscene = RNA_pointer_get(&b_scene.ptr, "cycles");
 
-  if (background && !b_engine.is_preview()) {
-    /* Viewport and preview renders do not require temp directory and do request session
-     * parameters more often than the background render.
-     * Optimize RNA-C++ usage and memory allocation a bit by saving string access which we know is
-     * not needed for viewport render. */
-    params.temp_dir = b_engine.temporary_directory();
-  }
-
   /* feature set */
   params.experimental = (get_enum(cscene, "feature_set") != 0);
 

intern/cycles/blender/sync.h

@@ -105,11 +105,11 @@ class BlenderSync {
   static BufferParams get_buffer_params(
       BL::SpaceView3D &b_v3d, BL::RegionView3D &b_rv3d, Camera *cam, int width, int height);
 
+ private:
   static DenoiseParams get_denoise_params(BL::Scene &b_scene,
                                           BL::ViewLayer &b_view_layer,
                                           bool background);
 
- private:
   /* sync */
   void sync_lights(BL::Depsgraph &b_depsgraph, bool update_all);
   void sync_materials(BL::Depsgraph &b_depsgraph, bool update_all);

intern/cycles/blender/util.h

@@ -18,7 +18,6 @@
 #define __BLENDER_UTIL_H__
 
 #include "scene/mesh.h"
-#include "scene/scene.h"
 
 #include "util/algorithm.h"
 #include "util/array.h"
@@ -34,7 +33,7 @@
 
 extern "C" {
 void BKE_image_user_frame_calc(void *ima, void *iuser, int cfra);
-void BKE_image_user_file_path_ex(void *iuser, void *ima, char *path, bool resolve_udim);
+void BKE_image_user_file_path(void *iuser, void *ima, char *path);
 unsigned char *BKE_image_get_pixels_for_frame(void *image, int frame, int tile);
 float *BKE_image_get_float_pixels_for_frame(void *image, int frame, int tile);
 }
@@ -291,14 +290,25 @@ static inline int render_resolution_y(BL::RenderSettings &b_render)
   return b_render.resolution_y() * b_render.resolution_percentage() / 100;
 }
 
-static inline string image_user_file_path(BL::ImageUser &iuser, BL::Image &ima, int cfra)
+static inline string image_user_file_path(BL::ImageUser &iuser,
+                                          BL::Image &ima,
+                                          int cfra,
+                                          bool load_tiled)
 {
   char filepath[1024];
   iuser.tile(0);
   BKE_image_user_frame_calc(ima.ptr.data, iuser.ptr.data, cfra);
-  BKE_image_user_file_path_ex(iuser.ptr.data, ima.ptr.data, filepath, false);
+  BKE_image_user_file_path(iuser.ptr.data, ima.ptr.data, filepath);
 
-  return string(filepath);
+  string filepath_str = string(filepath);
+  if (load_tiled && ima.source() == BL::Image::source_TILED) {
+    string udim;
+    if (!ima.tiles.empty()) {
+      udim = to_string(ima.tiles[0].number());
+    }
+    string_replace(filepath_str, udim, "<UDIM>");
+  }
+  return filepath_str;
 }
 
 static inline int image_user_frame_number(BL::ImageUser &iuser, BL::Image &ima, int cfra)
@@ -671,40 +681,6 @@ static inline uint object_ray_visibility(BL::Object &b_ob)
   return flag;
 }
 
-/* Check whether some of "built-in" motion-related attributes are needed to be exported (includes
- * things like velocity from cache modifier, fluid simulation).
- *
- * NOTE: This code is run prior to object motion blur initialization. so can not access properties
- * set by `sync_object_motion_init()`. */
-static inline bool object_need_motion_attribute(BObjectInfo &b_ob_info, Scene *scene)
-{
-  const Scene::MotionType need_motion = scene->need_motion();
-  if (need_motion == Scene::MOTION_NONE) {
-    /* Simple case: neither motion pass nor motion blur is needed, no need in the motion related
-     * attributes. */
-    return false;
-  }
-
-  if (need_motion == Scene::MOTION_BLUR) {
-    /* A bit tricky and implicit case:
-     * - Motion blur is enabled in the scene, which implies specific number of time steps for
-     *   objects.
-     * - If the object has motion blur disabled on it, it will have 0 time steps.
-     * - Motion attribute expects non-zero time steps.
-     *
-     * Avoid adding motion attributes if the motion blur will enforce 0 motion steps. */
-    PointerRNA cobject = RNA_pointer_get(&b_ob_info.real_object.ptr, "cycles");
-    const bool use_motion = get_boolean(cobject, "use_motion_blur");
-    if (!use_motion) {
-      return false;
-    }
-  }
-
-  /* Motion pass which implies 3 motion steps, or motion blur which is not disabled on object
-   * level. */
-  return true;
-}
-
 class EdgeMap {
  public:
   EdgeMap()

intern/cycles/bvh/build.cpp

@@ -935,7 +935,7 @@ BVHNode *BVHBuild::create_object_leaf_nodes(const BVHReference *ref, int start,
 
 BVHNode *BVHBuild::create_leaf_node(const BVHRange &range, const vector<BVHReference> &references)
 {
-  /* This is a bit over-allocating here (considering leaf size into account),
+  /* This is a bit overallocating here (considering leaf size into account),
    * but chunk-based re-allocation in vector makes it difficult to use small
    * size of stack storage here. Some tweaks are possible tho.
    *

intern/cycles/bvh/embree.cpp

@@ -61,26 +61,6 @@ static_assert(Object::MAX_MOTION_STEPS == Geometry::MAX_MOTION_STEPS,
 
 #  define IS_HAIR(x) (x & 1)
 
-/* This gets called by Embree at every valid ray/object intersection.
- * Things like recording subsurface or shadow hits for later evaluation
- * as well as filtering for volume objects happen here.
- * Cycles' own BVH does that directly inside the traversal calls.
- */
-static void rtc_filter_intersection_func(const RTCFilterFunctionNArguments *args)
-{
-  /* Current implementation in Cycles assumes only single-ray intersection queries. */
-  assert(args->N == 1);
-
-  RTCHit *hit = (RTCHit *)args->hit;
-  CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt;
-  const KernelGlobalsCPU *kg = ctx->kg;
-  const Ray *cray = ctx->ray;
-
-  if (kernel_embree_is_self_intersection(kg, hit, cray)) {
-    *args->valid = 0;
-  }
-}
-
 /* This gets called by Embree at every valid ray/object intersection.
  * Things like recording subsurface or shadow hits for later evaluation
  * as well as filtering for volume objects happen here.
@@ -95,16 +75,12 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
   RTCHit *hit = (RTCHit *)args->hit;
   CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt;
   const KernelGlobalsCPU *kg = ctx->kg;
-  const Ray *cray = ctx->ray;
 
   switch (ctx->type) {
     case CCLIntersectContext::RAY_SHADOW_ALL: {
       Intersection current_isect;
       kernel_embree_convert_hit(kg, ray, hit, &current_isect);
-      if (intersection_skip_self_shadow(cray->self, current_isect.object, current_isect.prim)) {
+
-        *args->valid = 0;
-        return;
-      }
       /* If no transparent shadows or max number of hits exceeded, all light is blocked. */
       const int flags = intersection_get_shader_flags(kg, current_isect.prim, current_isect.type);
       if (!(flags & (SD_HAS_TRANSPARENT_SHADOW)) || ctx->num_hits >= ctx->max_hits) {
@@ -184,10 +160,6 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
           break;
         }
       }
-      if (intersection_skip_self_local(cray->self, current_isect.prim)) {
-        *args->valid = 0;
-        return;
-      }
 
       /* No intersection information requested, just return a hit. */
       if (ctx->max_hits == 0) {
@@ -253,11 +225,6 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
       if (ctx->num_hits < ctx->max_hits) {
         Intersection current_isect;
         kernel_embree_convert_hit(kg, ray, hit, &current_isect);
-        if (intersection_skip_self(cray->self, current_isect.object, current_isect.prim)) {
-          *args->valid = 0;
-          return;
-        }
-
         Intersection *isect = &ctx->isect_s[ctx->num_hits];
         ++ctx->num_hits;
         *isect = current_isect;
@@ -269,15 +236,12 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
         }
         /* This tells Embree to continue tracing. */
         *args->valid = 0;
+        break;
       }
-      break;
     }
     case CCLIntersectContext::RAY_REGULAR:
     default:
-      if (kernel_embree_is_self_intersection(kg, hit, cray)) {
+      /* Nothing to do here. */
-        *args->valid = 0;
-        return;
-      }
       break;
   }
 }
@@ -293,14 +257,6 @@ static void rtc_filter_func_backface_cull(const RTCFilterFunctionNArguments *arg
     *args->valid = 0;
     return;
   }
-
-  CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt;
-  const KernelGlobalsCPU *kg = ctx->kg;
-  const Ray *cray = ctx->ray;
-
-  if (kernel_embree_is_self_intersection(kg, hit, cray)) {
-    *args->valid = 0;
-  }
 }
 
 static void rtc_filter_occluded_func_backface_cull(const RTCFilterFunctionNArguments *args)
@@ -399,12 +355,10 @@ void BVHEmbree::build(Progress &progress, Stats *stats, RTCDevice rtc_device_)
   }
 
   const bool dynamic = params.bvh_type == BVH_TYPE_DYNAMIC;
-  const bool compact = params.use_compact_structure;
 
   scene = rtcNewScene(rtc_device);
   const RTCSceneFlags scene_flags = (dynamic ? RTC_SCENE_FLAG_DYNAMIC : RTC_SCENE_FLAG_NONE) |
-                                    (compact ? RTC_SCENE_FLAG_COMPACT : RTC_SCENE_FLAG_NONE) |
+                                    RTC_SCENE_FLAG_COMPACT | RTC_SCENE_FLAG_ROBUST;
-                                    RTC_SCENE_FLAG_ROBUST;
   rtcSetSceneFlags(scene, scene_flags);
   build_quality = dynamic ? RTC_BUILD_QUALITY_LOW :
                             (params.use_spatial_split ? RTC_BUILD_QUALITY_HIGH :
@@ -549,7 +503,6 @@ void BVHEmbree::add_triangles(const Object *ob, const Mesh *mesh, int i)
 
   rtcSetGeometryUserData(geom_id, (void *)prim_offset);
   rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func);
-  rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_intersection_func);
   rtcSetGeometryMask(geom_id, ob->visibility_for_tracing());
 
   rtcCommitGeometry(geom_id);
@@ -812,7 +765,6 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
 
   rtcSetGeometryUserData(geom_id, (void *)prim_offset);
   if (hair->curve_shape == CURVE_RIBBON) {
-    rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_intersection_func);
     rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func);
   }
   else {

intern/cycles/bvh/params.h

@@ -97,9 +97,6 @@ class BVHParams {
    */
   bool use_unaligned_nodes;
 
-  /* Use compact acceleration structure (Embree)*/
-  bool use_compact_structure;
-
   /* Split time range to this number of steps and create leaf node for each
    * of this time steps.
    *

intern/cycles/cmake/macros.cmake

@@ -84,6 +84,39 @@ macro(cycles_add_library target library_deps)
   cycles_set_solution_folder(${target})
 endmacro()
 
+# Cycles library dependencies common to all executables
+
+function(cycles_link_directories)
+  if(APPLE)
+    # APPLE platform uses full paths for linking libraries, and avoids link_directories.
+    return()
+  endif()
+
+  if(WITH_OPENCOLORIO)
+    link_directories(${OPENCOLORIO_LIBPATH})
+  endif()
+  if(WITH_OPENVDB)
+    link_directories(${OPENVDB_LIBPATH} ${BLOSC_LIBPATH})
+  endif()
+  if(WITH_OPENSUBDIV)
+    link_directories(${OPENSUBDIV_LIBPATH})
+  endif()
+  if(WITH_OPENIMAGEDENOISE)
+    link_directories(${OPENIMAGEDENOISE_LIBPATH})
+  endif()
+
+  link_directories(
+    ${OPENIMAGEIO_LIBPATH}
+    ${BOOST_LIBPATH}
+    ${PNG_LIBPATH}
+    ${JPEG_LIBPATH}
+    ${ZLIB_LIBPATH}
+    ${TIFF_LIBPATH}
+    ${OPENEXR_LIBPATH}
+    ${OPENJPEG_LIBPATH}
+  )
+endfunction()
+
 macro(cycles_target_link_libraries target)
   if(WITH_CYCLES_LOGGING)
     target_link_libraries(${target} ${GLOG_LIBRARIES} ${GFLAGS_LIBRARIES})
@@ -135,6 +168,12 @@ macro(cycles_target_link_libraries target)
     target_link_libraries(${target} extern_hipew)
   endif()
 
+  if(CYCLES_STANDALONE_REPOSITORY)
+    target_link_libraries(${target} extern_numaapi)
+  else()
+    target_link_libraries(${target} bf_intern_numaapi)
+  endif()
+
   if(UNIX AND NOT APPLE)
     if(CYCLES_STANDALONE_REPOSITORY)
       target_link_libraries(${target} extern_libc_compat)

intern/cycles/device/cpu/device_impl.cpp

@@ -72,7 +72,7 @@ CPUDevice::CPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_
           << " CPU kernels.";
 
   if (info.cpu_threads == 0) {
-    info.cpu_threads = TaskScheduler::max_concurrency();
+    info.cpu_threads = TaskScheduler::num_threads();
   }
 
 #ifdef WITH_OSL

intern/cycles/device/cuda/graphics_interop.cpp

@@ -45,10 +45,8 @@ void CUDADeviceGraphicsInterop::set_display_interop(
 
   need_clear_ = display_interop.need_clear;
 
-  if (!display_interop.need_recreate) {
+  if (opengl_pbo_id_ == display_interop.opengl_pbo_id && buffer_area_ == new_buffer_area) {
-    if (opengl_pbo_id_ == display_interop.opengl_pbo_id && buffer_area_ == new_buffer_area) {
+    return;
-      return;
-    }
   }
 
   CUDAContextScope scope(device_);

intern/cycles/device/denoise.cpp

@@ -76,8 +76,6 @@ NODE_DEFINE(DenoiseParams)
   SOCKET_BOOLEAN(use_pass_albedo, "Use Pass Albedo", true);
   SOCKET_BOOLEAN(use_pass_normal, "Use Pass Normal", false);
 
-  SOCKET_BOOLEAN(temporally_stable, "Temporally Stable", false);
-
   SOCKET_ENUM(prefilter, "Prefilter", *prefilter_enum, DENOISER_PREFILTER_FAST);
 
   return type;

intern/cycles/device/denoise.h

@@ -72,9 +72,6 @@ class DenoiseParams : public Node {
   bool use_pass_albedo = true;
   bool use_pass_normal = true;
 
-  /* Configure the denoiser to use motion vectors, previous image and a temporally stable model. */
-  bool temporally_stable = false;
-
   DenoiserPrefilter prefilter = DENOISER_PREFILTER_FAST;
 
   static const NodeEnum *get_type_enum();
@@ -86,8 +83,7 @@ class DenoiseParams : public Node {
   {
     return !(use == other.use && type == other.type && start_sample == other.start_sample &&
              use_pass_albedo == other.use_pass_albedo &&
-             use_pass_normal == other.use_pass_normal &&
+             use_pass_normal == other.use_pass_normal && prefilter == other.prefilter);
-             temporally_stable == other.temporally_stable && prefilter == other.prefilter);
   }
 };
 

intern/cycles/device/device.cpp

@@ -37,7 +37,6 @@
 #include "util/math.h"
 #include "util/string.h"
 #include "util/system.h"
-#include "util/task.h"
 #include "util/time.h"
 #include "util/types.h"
 #include "util/vector.h"
@@ -334,7 +333,7 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
     /* Ensure CPU device does not slow down GPU. */
     if (device.type == DEVICE_CPU && subdevices.size() > 1) {
       if (background) {
-        int orig_cpu_threads = (threads) ? threads : TaskScheduler::max_concurrency();
+        int orig_cpu_threads = (threads) ? threads : system_cpu_thread_count();
         int cpu_threads = max(orig_cpu_threads - (subdevices.size() - 1), 0);
 
         VLOG(1) << "CPU render threads reduced from " << orig_cpu_threads << " to " << cpu_threads

intern/cycles/device/metal/bvh.h

@@ -58,11 +58,6 @@ class BVHMetal : public BVH {
                        id<MTLCommandQueue> queue,
                        Geometry *const geom,
                        bool refit);
-  bool build_BLAS_pointcloud(Progress &progress,
-                             id<MTLDevice> device,
-                             id<MTLCommandQueue> queue,
-                             Geometry *const geom,
-                             bool refit);
   bool build_TLAS(Progress &progress, id<MTLDevice> device, id<MTLCommandQueue> queue, bool refit);
 };
 

intern/cycles/device/metal/bvh.mm

@@ -19,7 +19,6 @@
 #  include "scene/hair.h"
 #  include "scene/mesh.h"
 #  include "scene/object.h"
-#  include "scene/pointcloud.h"
 
 #  include "util/progress.h"
 
@@ -476,220 +475,6 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
   return false;
 }
 
-bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
-                                     id<MTLDevice> device,
-                                     id<MTLCommandQueue> queue,
-                                     Geometry *const geom,
-                                     bool refit)
-{
-  if (@available(macos 12.0, *)) {
-    /* Build BLAS for point cloud */
-    PointCloud *pointcloud = static_cast<PointCloud *>(geom);
-    if (pointcloud->num_points() == 0) {
-      return false;
-    }
-
-    /*------------------------------------------------*/
-    BVH_status("Building pointcloud BLAS | %7d points | %s",
-               (int)pointcloud->num_points(),
-               geom->name.c_str());
-    /*------------------------------------------------*/
-
-    const size_t num_points = pointcloud->get_points().size();
-    const float3 *points = pointcloud->get_points().data();
-    const float *radius = pointcloud->get_radius().data();
-
-    const bool use_fast_trace_bvh = (params.bvh_type == BVH_TYPE_STATIC);
-
-    size_t num_motion_steps = 1;
-    Attribute *motion_keys = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
-    if (motion_blur && pointcloud->get_use_motion_blur() && motion_keys) {
-      num_motion_steps = pointcloud->get_motion_steps();
-    }
-
-    const size_t num_aabbs = num_motion_steps;
-
-    MTLResourceOptions storage_mode;
-    if (device.hasUnifiedMemory) {
-      storage_mode = MTLResourceStorageModeShared;
-    }
-    else {
-      storage_mode = MTLResourceStorageModeManaged;
-    }
-
-    /* Allocate a GPU buffer for the AABB data and populate it */
-    id<MTLBuffer> aabbBuf = [device
-        newBufferWithLength:num_aabbs * sizeof(MTLAxisAlignedBoundingBox)
-                    options:storage_mode];
-    MTLAxisAlignedBoundingBox *aabb_data = (MTLAxisAlignedBoundingBox *)[aabbBuf contents];
-
-    /* Get AABBs for each motion step */
-    size_t center_step = (num_motion_steps - 1) / 2;
-    for (size_t step = 0; step < num_motion_steps; ++step) {
-      /* The center step for motion vertices is not stored in the attribute */
-      if (step != center_step) {
-        size_t attr_offset = (step > center_step) ? step - 1 : step;
-        points = motion_keys->data_float3() + attr_offset * num_points;
-      }
-
-      for (size_t j = 0; j < num_points; ++j) {
-        const PointCloud::Point point = pointcloud->get_point(j);
-        BoundBox bounds = BoundBox::empty;
-        point.bounds_grow(points, radius, bounds);
-
-        const size_t index = step * num_points + j;
-        aabb_data[index].min = (MTLPackedFloat3 &)bounds.min;
-        aabb_data[index].max = (MTLPackedFloat3 &)bounds.max;
-      }
-    }
-
-    if (storage_mode == MTLResourceStorageModeManaged) {
-      [aabbBuf didModifyRange:NSMakeRange(0, aabbBuf.length)];
-    }
-
-#  if 0
-    for (size_t i=0; i<num_aabbs && i < 400; i++) {
-      MTLAxisAlignedBoundingBox& bb = aabb_data[i];
-      printf("  %d:   %.1f,%.1f,%.1f -- %.1f,%.1f,%.1f\n", int(i), bb.min.x, bb.min.y, bb.min.z, bb.max.x, bb.max.y, bb.max.z);
-    }
-#  endif
-
-    MTLAccelerationStructureGeometryDescriptor *geomDesc;
-    if (motion_blur) {
-      std::vector<MTLMotionKeyframeData *> aabb_ptrs;
-      aabb_ptrs.reserve(num_motion_steps);
-      for (size_t step = 0; step < num_motion_steps; ++step) {
-        MTLMotionKeyframeData *k = [MTLMotionKeyframeData data];
-        k.buffer = aabbBuf;
-        k.offset = step * num_points * sizeof(MTLAxisAlignedBoundingBox);
-        aabb_ptrs.push_back(k);
-      }
-
-      MTLAccelerationStructureMotionBoundingBoxGeometryDescriptor *geomDescMotion =
-          [MTLAccelerationStructureMotionBoundingBoxGeometryDescriptor descriptor];
-      geomDescMotion.boundingBoxBuffers = [NSArray arrayWithObjects:aabb_ptrs.data()
-                                                              count:aabb_ptrs.size()];
-      geomDescMotion.boundingBoxCount = num_points;
-      geomDescMotion.boundingBoxStride = sizeof(aabb_data[0]);
-      geomDescMotion.intersectionFunctionTableOffset = 2;
-
-      /* Force a single any-hit call, so shadow record-all behavior works correctly */
-      /* (Match optix behavior: unsigned int build_flags =
-       * OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;) */
-      geomDescMotion.allowDuplicateIntersectionFunctionInvocation = false;
-      geomDescMotion.opaque = true;
-      geomDesc = geomDescMotion;
-    }
-    else {
-      MTLAccelerationStructureBoundingBoxGeometryDescriptor *geomDescNoMotion =
-          [MTLAccelerationStructureBoundingBoxGeometryDescriptor descriptor];
-      geomDescNoMotion.boundingBoxBuffer = aabbBuf;
-      geomDescNoMotion.boundingBoxBufferOffset = 0;
-      geomDescNoMotion.boundingBoxCount = int(num_aabbs);
-      geomDescNoMotion.boundingBoxStride = sizeof(aabb_data[0]);
-      geomDescNoMotion.intersectionFunctionTableOffset = 2;
-
-      /* Force a single any-hit call, so shadow record-all behavior works correctly */
-      /* (Match optix behavior: unsigned int build_flags =
-       * OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;) */
-      geomDescNoMotion.allowDuplicateIntersectionFunctionInvocation = false;
-      geomDescNoMotion.opaque = true;
-      geomDesc = geomDescNoMotion;
-    }
-
-    MTLPrimitiveAccelerationStructureDescriptor *accelDesc =
-        [MTLPrimitiveAccelerationStructureDescriptor descriptor];
-    accelDesc.geometryDescriptors = @[ geomDesc ];
-
-    if (motion_blur) {
-      accelDesc.motionStartTime = 0.0f;
-      accelDesc.motionEndTime = 1.0f;
-      accelDesc.motionStartBorderMode = MTLMotionBorderModeVanish;
-      accelDesc.motionEndBorderMode = MTLMotionBorderModeVanish;
-      accelDesc.motionKeyframeCount = num_motion_steps;
-    }
-
-    if (!use_fast_trace_bvh) {
-      accelDesc.usage |= (MTLAccelerationStructureUsageRefit |
-                          MTLAccelerationStructureUsagePreferFastBuild);
-    }
-
-    MTLAccelerationStructureSizes accelSizes = [device
-        accelerationStructureSizesWithDescriptor:accelDesc];
-    id<MTLAccelerationStructure> accel_uncompressed = [device
-        newAccelerationStructureWithSize:accelSizes.accelerationStructureSize];
-    id<MTLBuffer> scratchBuf = [device newBufferWithLength:accelSizes.buildScratchBufferSize
-                                                   options:MTLResourceStorageModePrivate];
-    id<MTLBuffer> sizeBuf = [device newBufferWithLength:8 options:MTLResourceStorageModeShared];
-    id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
-    id<MTLAccelerationStructureCommandEncoder> accelEnc =
-        [accelCommands accelerationStructureCommandEncoder];
-    if (refit) {
-      [accelEnc refitAccelerationStructure:accel_struct
-                                descriptor:accelDesc
-                               destination:accel_uncompressed
-                             scratchBuffer:scratchBuf
-                       scratchBufferOffset:0];
-    }
-    else {
-      [accelEnc buildAccelerationStructure:accel_uncompressed
-                                descriptor:accelDesc
-                             scratchBuffer:scratchBuf
-                       scratchBufferOffset:0];
-    }
-    if (use_fast_trace_bvh) {
-      [accelEnc writeCompactedAccelerationStructureSize:accel_uncompressed
-                                               toBuffer:sizeBuf
-                                                 offset:0
-                                           sizeDataType:MTLDataTypeULong];
-    }
-    [accelEnc endEncoding];
-    [accelCommands addCompletedHandler:^(id<MTLCommandBuffer> command_buffer) {
-      /* free temp resources */
-      [scratchBuf release];
-      [aabbBuf release];
-
-      if (use_fast_trace_bvh) {
-        /* Compact the accel structure */
-        uint64_t compressed_size = *(uint64_t *)sizeBuf.contents;
-
-        dispatch_async(dispatch_get_global_queue(DISPATCH_QUEUE_PRIORITY_DEFAULT, 0), ^{
-          id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
-          id<MTLAccelerationStructureCommandEncoder> accelEnc =
-              [accelCommands accelerationStructureCommandEncoder];
-          id<MTLAccelerationStructure> accel = [device
-              newAccelerationStructureWithSize:compressed_size];
-          [accelEnc copyAndCompactAccelerationStructure:accel_uncompressed
-                                toAccelerationStructure:accel];
-          [accelEnc endEncoding];
-          [accelCommands addCompletedHandler:^(id<MTLCommandBuffer> command_buffer) {
-            uint64_t allocated_size = [accel allocatedSize];
-            stats.mem_alloc(allocated_size);
-            accel_struct = accel;
-            [accel_uncompressed release];
-            accel_struct_building = false;
-          }];
-          [accelCommands commit];
-        });
-      }
-      else {
-        /* set our acceleration structure to the uncompressed structure */
-        accel_struct = accel_uncompressed;
-
-        uint64_t allocated_size = [accel_struct allocatedSize];
-        stats.mem_alloc(allocated_size);
-        accel_struct_building = false;
-      }
-      [sizeBuf release];
-    }];
-
-    accel_struct_building = true;
-    [accelCommands commit];
-    return true;
-  }
-  return false;
-}
-
 bool BVHMetal::build_BLAS(Progress &progress,
                           id<MTLDevice> device,
                           id<MTLCommandQueue> queue,
@@ -706,8 +491,6 @@ bool BVHMetal::build_BLAS(Progress &progress,
         return build_BLAS_mesh(progress, device, queue, geom, refit);
       case Geometry::HAIR:
         return build_BLAS_hair(progress, device, queue, geom, refit);
-      case Geometry::POINTCLOUD:
-        return build_BLAS_pointcloud(progress, device, queue, geom, refit);
       default:
         return false;
     }

intern/cycles/device/metal/device_impl.h

@@ -115,8 +115,6 @@ class MetalDevice : public Device {
 
   void load_texture_info();
 
-  void erase_allocation(device_memory &mem);
-
   virtual bool should_use_graphics_interop() override;
 
   virtual unique_ptr<DeviceQueue> gpu_queue_create() override;

intern/cycles/device/metal/device_impl.mm

@@ -87,14 +87,17 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
     default:
       break;
     case METAL_GPU_INTEL: {
+      use_metalrt = false;
       max_threads_per_threadgroup = 64;
       break;
     }
     case METAL_GPU_AMD: {
+      use_metalrt = false;
       max_threads_per_threadgroup = 128;
       break;
     }
     case METAL_GPU_APPLE: {
+      use_metalrt = true;
       max_threads_per_threadgroup = 512;
       break;
     }
@@ -429,25 +432,6 @@ void MetalDevice::load_texture_info()
   }
 }
 
-void MetalDevice::erase_allocation(device_memory &mem)
-{
-  stats.mem_free(mem.device_size);
-  mem.device_pointer = 0;
-  mem.device_size = 0;
-
-  auto it = metal_mem_map.find(&mem);
-  if (it != metal_mem_map.end()) {
-    MetalMem *mmem = it->second.get();
-
-    /* blank out reference to MetalMem* in the launch params (fixes crash T94736) */
-    if (mmem->pointer_index >= 0) {
-      device_ptr *pointers = (device_ptr *)&launch_params;
-      pointers[mmem->pointer_index] = 0;
-    }
-    metal_mem_map.erase(it);
-  }
-}
-
 MetalDevice::MetalMem *MetalDevice::generic_alloc(device_memory &mem)
 {
   size_t size = mem.memory_size();
@@ -577,7 +561,11 @@ void MetalDevice::generic_free(device_memory &mem)
       mmem.mtlBuffer = nil;
     }
 
-    erase_allocation(mem);
+    stats.mem_free(mem.device_size);
+    mem.device_pointer = 0;
+    mem.device_size = 0;
+
+    metal_mem_map.erase(&mem);
   }
 }
 
@@ -966,7 +954,10 @@ void MetalDevice::tex_free(device_texture &mem)
       delayed_free_list.push_back(mmem.mtlTexture);
       mmem.mtlTexture = nil;
     }
-    erase_allocation(mem);
+    stats.mem_free(mem.device_size);
+    mem.device_pointer = 0;
+    mem.device_size = 0;
+    metal_mem_map.erase(&mem);
   }
 }
 

intern/cycles/device/metal/kernel.h

@@ -36,8 +36,6 @@ enum {
   METALRT_FUNC_CURVE_RIBBON_SHADOW,
   METALRT_FUNC_CURVE_ALL,
   METALRT_FUNC_CURVE_ALL_SHADOW,
-  METALRT_FUNC_POINT,
-  METALRT_FUNC_POINT_SHADOW,
   METALRT_FUNC_NUM
 };
 

intern/cycles/device/metal/kernel.mm

@@ -358,8 +358,6 @@ bool MetalDeviceKernels::load(MetalDevice *device, int kernel_type)
           "__intersection__curve_ribbon_shadow",
           "__intersection__curve_all",
           "__intersection__curve_all_shadow",
-          "__intersection__point",
-          "__intersection__point_shadow",
       };
       assert(sizeof(function_names) / sizeof(function_names[0]) == METALRT_FUNC_NUM);
 
@@ -402,50 +400,36 @@ bool MetalDeviceKernels::load(MetalDevice *device, int kernel_type)
     NSArray *function_list = nil;
 
     if (device->use_metalrt) {
-      id<MTLFunction> curve_intersect_default = nil;
+      id<MTLFunction> box_intersect_default = nil;
-      id<MTLFunction> curve_intersect_shadow = nil;
+      id<MTLFunction> box_intersect_shadow = nil;
-      id<MTLFunction> point_intersect_default = nil;
-      id<MTLFunction> point_intersect_shadow = nil;
       if (device->kernel_features & KERNEL_FEATURE_HAIR) {
         /* Add curve intersection programs. */
         if (device->kernel_features & KERNEL_FEATURE_HAIR_THICK) {
           /* Slower programs for thick hair since that also slows down ribbons.
            * Ideally this should not be needed. */
-          curve_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL];
+          box_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL];
-          curve_intersect_shadow =
+          box_intersect_shadow = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL_SHADOW];
-              rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL_SHADOW];
         }
         else {
-          curve_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_RIBBON];
+          box_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_RIBBON];
-          curve_intersect_shadow =
+          box_intersect_shadow =
               rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_RIBBON_SHADOW];
         }
       }
-      if (device->kernel_features & KERNEL_FEATURE_POINTCLOUD) {
-        point_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_POINT];
-        point_intersect_shadow = rt_intersection_funcs[kernel_type][METALRT_FUNC_POINT_SHADOW];
-      }
       table_functions[METALRT_TABLE_DEFAULT] = [NSArray
           arrayWithObjects:rt_intersection_funcs[kernel_type][METALRT_FUNC_DEFAULT_TRI],
-                           curve_intersect_default ?
+                           box_intersect_default ?
-                               curve_intersect_default :
+                               box_intersect_default :
-                               rt_intersection_funcs[kernel_type][METALRT_FUNC_DEFAULT_BOX],
-                           point_intersect_default ?
-                               point_intersect_default :
                                rt_intersection_funcs[kernel_type][METALRT_FUNC_DEFAULT_BOX],
                            nil];
       table_functions[METALRT_TABLE_SHADOW] = [NSArray
           arrayWithObjects:rt_intersection_funcs[kernel_type][METALRT_FUNC_SHADOW_TRI],
-                           curve_intersect_shadow ?
+                           box_intersect_shadow ?
-                               curve_intersect_shadow :
+                               box_intersect_shadow :
-                               rt_intersection_funcs[kernel_type][METALRT_FUNC_SHADOW_BOX],
-                           point_intersect_shadow ?
-                               point_intersect_shadow :
                                rt_intersection_funcs[kernel_type][METALRT_FUNC_SHADOW_BOX],
                            nil];
       table_functions[METALRT_TABLE_LOCAL] = [NSArray
           arrayWithObjects:rt_intersection_funcs[kernel_type][METALRT_FUNC_LOCAL_TRI],
-                           rt_intersection_funcs[kernel_type][METALRT_FUNC_LOCAL_BOX],
                            rt_intersection_funcs[kernel_type][METALRT_FUNC_LOCAL_BOX],
                            nil];
 

intern/cycles/device/optix/device_impl.cpp

@@ -226,7 +226,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
   pipeline_options.usesMotionBlur = false;
   pipeline_options.traversableGraphFlags =
       OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_LEVEL_INSTANCING;
-  pipeline_options.numPayloadValues = 8;
+  pipeline_options.numPayloadValues = 6;
   pipeline_options.numAttributeValues = 2; /* u, v */
   pipeline_options.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE;
   pipeline_options.pipelineLaunchParamsVariableName = "__params"; /* See globals.h */
@@ -566,19 +566,6 @@ class OptiXDevice::DenoiseContext {
       }
     }
 
-    if (denoise_params.temporally_stable) {
-      prev_output.device_pointer = render_buffers->buffer.device_pointer;
-
-      prev_output.offset = buffer_params.get_pass_offset(PASS_DENOISING_PREVIOUS);
-
-      prev_output.stride = buffer_params.stride;
-      prev_output.pass_stride = buffer_params.pass_stride;
-
-      num_input_passes += 1;
-      use_pass_flow = true;
-      pass_motion = buffer_params.get_pass_offset(PASS_MOTION);
-    }
-
     use_guiding_passes = (num_input_passes - 1) > 0;
 
     if (use_guiding_passes) {
@@ -587,7 +574,6 @@ class OptiXDevice::DenoiseContext {
 
         guiding_params.pass_albedo = pass_denoising_albedo;
         guiding_params.pass_normal = pass_denoising_normal;
-        guiding_params.pass_flow = pass_motion;
 
         guiding_params.stride = buffer_params.stride;
         guiding_params.pass_stride = buffer_params.pass_stride;
@@ -602,10 +588,6 @@ class OptiXDevice::DenoiseContext {
           guiding_params.pass_normal = guiding_params.pass_stride;
           guiding_params.pass_stride += 3;
         }
-        if (use_pass_flow) {
-          guiding_params.pass_flow = guiding_params.pass_stride;
-          guiding_params.pass_stride += 2;
-        }
 
         guiding_params.stride = buffer_params.width;
 
@@ -623,16 +605,6 @@ class OptiXDevice::DenoiseContext {
   RenderBuffers *render_buffers = nullptr;
   const BufferParams &buffer_params;
 
-  /* Previous output. */
-  struct {
-    device_ptr device_pointer = 0;
-
-    int offset = PASS_UNUSED;
-
-    int stride = -1;
-    int pass_stride = -1;
-  } prev_output;
-
   /* Device-side storage of the guiding passes. */
   device_only_memory<float> guiding_buffer;
 
@@ -642,7 +614,6 @@ class OptiXDevice::DenoiseContext {
     /* NOTE: Are only initialized when the corresponding guiding pass is enabled. */
     int pass_albedo = PASS_UNUSED;
     int pass_normal = PASS_UNUSED;
-    int pass_flow = PASS_UNUSED;
 
     int stride = -1;
     int pass_stride = -1;
@@ -653,7 +624,6 @@ class OptiXDevice::DenoiseContext {
   bool use_guiding_passes = false;
   bool use_pass_albedo = false;
   bool use_pass_normal = false;
-  bool use_pass_flow = false;
 
   int num_samples = 0;
 
@@ -662,7 +632,6 @@ class OptiXDevice::DenoiseContext {
   /* NOTE: Are only initialized when the corresponding guiding pass is enabled. */
   int pass_denoising_albedo = PASS_UNUSED;
   int pass_denoising_normal = PASS_UNUSED;
-  int pass_motion = PASS_UNUSED;
 
   /* For passes which don't need albedo channel for denoising we replace the actual albedo with
    * the (0.5, 0.5, 0.5). This flag indicates that the real albedo pass has been replaced with
@@ -733,7 +702,6 @@ bool OptiXDevice::denoise_filter_guiding_preprocess(DenoiseContext &context)
                              &context.guiding_params.pass_stride,
                              &context.guiding_params.pass_albedo,
                              &context.guiding_params.pass_normal,
-                             &context.guiding_params.pass_flow,
                              &context.render_buffers->buffer.device_pointer,
                              &buffer_params.offset,
                              &buffer_params.stride,
@@ -741,7 +709,6 @@ bool OptiXDevice::denoise_filter_guiding_preprocess(DenoiseContext &context)
                              &context.pass_sample_count,
                              &context.pass_denoising_albedo,
                              &context.pass_denoising_normal,
-                             &context.pass_motion,
                              &buffer_params.full_x,
                              &buffer_params.full_y,
                              &buffer_params.width,
@@ -914,8 +881,7 @@ bool OptiXDevice::denoise_create_if_needed(DenoiseContext &context)
 {
   const bool recreate_denoiser = (denoiser_.optix_denoiser == nullptr) ||
                                  (denoiser_.use_pass_albedo != context.use_pass_albedo) ||
-                                 (denoiser_.use_pass_normal != context.use_pass_normal) ||
+                                 (denoiser_.use_pass_normal != context.use_pass_normal);
-                                 (denoiser_.use_pass_flow != context.use_pass_flow);
   if (!recreate_denoiser) {
     return true;
   }
@@ -929,14 +895,8 @@ bool OptiXDevice::denoise_create_if_needed(DenoiseContext &context)
   OptixDenoiserOptions denoiser_options = {};
   denoiser_options.guideAlbedo = context.use_pass_albedo;
   denoiser_options.guideNormal = context.use_pass_normal;
-
-  OptixDenoiserModelKind model = OPTIX_DENOISER_MODEL_KIND_HDR;
-  if (context.use_pass_flow) {
-    model = OPTIX_DENOISER_MODEL_KIND_TEMPORAL;
-  }
-
   const OptixResult result = optixDenoiserCreate(
-      this->context, model, &denoiser_options, &denoiser_.optix_denoiser);
+      this->context, OPTIX_DENOISER_MODEL_KIND_HDR, &denoiser_options, &denoiser_.optix_denoiser);
 
   if (result != OPTIX_SUCCESS) {
     set_error("Failed to create OptiX denoiser");
@@ -946,7 +906,6 @@ bool OptiXDevice::denoise_create_if_needed(DenoiseContext &context)
   /* OptiX denoiser handle was created with the requested number of input passes. */
   denoiser_.use_pass_albedo = context.use_pass_albedo;
   denoiser_.use_pass_normal = context.use_pass_normal;
-  denoiser_.use_pass_flow = context.use_pass_flow;
 
   /* OptiX denoiser has been created, but it needs configuration. */
   denoiser_.is_configured = false;
@@ -1006,10 +965,8 @@ bool OptiXDevice::denoise_run(DenoiseContext &context, const DenoisePass &pass)
   OptixImage2D color_layer = {0};
   OptixImage2D albedo_layer = {0};
   OptixImage2D normal_layer = {0};
-  OptixImage2D flow_layer = {0};
 
   OptixImage2D output_layer = {0};
-  OptixImage2D prev_output_layer = {0};
 
   /* Color pass. */
   {
@@ -1025,19 +982,6 @@ bool OptiXDevice::denoise_run(DenoiseContext &context, const DenoisePass &pass)
     color_layer.format = OPTIX_PIXEL_FORMAT_FLOAT3;
   }
 
-  /* Previous output. */
-  if (context.prev_output.offset != PASS_UNUSED) {
-    const int64_t pass_stride_in_bytes = context.prev_output.pass_stride * sizeof(float);
-
-    prev_output_layer.data = context.prev_output.device_pointer +
-                             context.prev_output.offset * sizeof(float);
-    prev_output_layer.width = width;
-    prev_output_layer.height = height;
-    prev_output_layer.rowStrideInBytes = pass_stride_in_bytes * context.prev_output.stride;
-    prev_output_layer.pixelStrideInBytes = pass_stride_in_bytes;
-    prev_output_layer.format = OPTIX_PIXEL_FORMAT_FLOAT3;
-  }
-
   /* Optional albedo and color passes. */
   if (context.num_input_passes > 1) {
     const device_ptr d_guiding_buffer = context.guiding_params.device_pointer;
@@ -1061,32 +1005,21 @@ bool OptiXDevice::denoise_run(DenoiseContext &context, const DenoisePass &pass)
       normal_layer.pixelStrideInBytes = pixel_stride_in_bytes;
       normal_layer.format = OPTIX_PIXEL_FORMAT_FLOAT3;
     }
-
-    if (context.use_pass_flow) {
-      flow_layer.data = d_guiding_buffer + context.guiding_params.pass_flow * sizeof(float);
-      flow_layer.width = width;
-      flow_layer.height = height;
-      flow_layer.rowStrideInBytes = row_stride_in_bytes;
-      flow_layer.pixelStrideInBytes = pixel_stride_in_bytes;
-      flow_layer.format = OPTIX_PIXEL_FORMAT_FLOAT2;
-    }
   }
 
   /* Denoise in-place of the noisy input in the render buffers. */
   output_layer = color_layer;
 
-  OptixDenoiserGuideLayer guide_layers = {};
+  /* Finally run denoising. */
-  guide_layers.albedo = albedo_layer;
+  OptixDenoiserParams params = {}; /* All parameters are disabled/zero. */
-  guide_layers.normal = normal_layer;
-  guide_layers.flow = flow_layer;
 
   OptixDenoiserLayer image_layers = {};
   image_layers.input = color_layer;
-  image_layers.previousOutput = prev_output_layer;
   image_layers.output = output_layer;
 
-  /* Finally run denoising. */
+  OptixDenoiserGuideLayer guide_layers = {};
-  OptixDenoiserParams params = {}; /* All parameters are disabled/zero. */
+  guide_layers.albedo = albedo_layer;
+  guide_layers.normal = normal_layer;
 
   optix_assert(optixUtilDenoiserInvokeTiled(denoiser_.optix_denoiser,
                                             denoiser_.queue.stream(),

intern/cycles/device/optix/device_impl.h

@@ -104,7 +104,6 @@ class OptiXDevice : public CUDADevice {
 
     bool use_pass_albedo = false;
     bool use_pass_normal = false;
-    bool use_pass_flow = false;
   };
   Denoiser denoiser_;
 

intern/cycles/device/queue.h

@@ -19,7 +19,6 @@
 #include "device/kernel.h"
 
 #include "device/graphics_interop.h"
-#include "util/debug.h"
 #include "util/log.h"
 #include "util/map.h"
 #include "util/string.h"
@@ -43,7 +42,7 @@ struct DeviceKernelArguments {
     KERNEL_FILM_CONVERT,
   };
 
-  static const int MAX_ARGS = 18;
+  static const int MAX_ARGS = 16;
   Type types[MAX_ARGS];
   void *values[MAX_ARGS];
   size_t sizes[MAX_ARGS];
@@ -86,8 +85,6 @@ struct DeviceKernelArguments {
   }
   void add(const Type type, const void *value, size_t size)
   {
-    assert(count < MAX_ARGS);
-
     types[count] = type;
     values[count] = (void *)value;
     sizes[count] = size;

intern/cycles/integrator/pass_accessor.cpp

@@ -141,7 +141,6 @@ bool PassAccessor::get_render_tile_pixels(const RenderBuffers *render_buffers,
   const PassType type = pass_access_info_.type;
   const PassMode mode = pass_access_info_.mode;
   const PassInfo pass_info = Pass::get_info(type, pass_access_info_.include_albedo);
-  int num_written_components = pass_info.num_components;
 
   if (pass_info.num_components == 1) {
     /* Single channel passes. */
@@ -189,10 +188,8 @@ bool PassAccessor::get_render_tile_pixels(const RenderBuffers *render_buffers,
     else if ((pass_info.divide_type != PASS_NONE || pass_info.direct_type != PASS_NONE ||
               pass_info.indirect_type != PASS_NONE) &&
              mode != PassMode::DENOISED) {
-      /* RGB lighting passes that need to divide out color and/or sum direct and indirect.
+      /* RGB lighting passes that need to divide out color and/or sum direct and indirect. */
-       * These can also optionally write alpha like the combined pass. */
       get_pass_light_path(render_buffers, buffer_params, destination);
-      num_written_components = 4;
     }
     else {
       /* Passes that need no special computation, or denoised passes that already
@@ -218,7 +215,7 @@ bool PassAccessor::get_render_tile_pixels(const RenderBuffers *render_buffers,
     }
   }
 
-  pad_pixels(buffer_params, destination, num_written_components);
+  pad_pixels(buffer_params, destination, pass_info.num_components);
 
   return true;
 }

intern/cycles/integrator/path_trace.cpp

@@ -115,9 +115,7 @@ bool PathTrace::ready_to_reset()
   return false;
 }
 
-void PathTrace::reset(const BufferParams &full_params,
+void PathTrace::reset(const BufferParams &full_params, const BufferParams &big_tile_params)
-                      const BufferParams &big_tile_params,
-                      const bool reset_rendering)
 {
   if (big_tile_params_.modified(big_tile_params)) {
     big_tile_params_ = big_tile_params;
@@ -130,7 +128,7 @@ void PathTrace::reset(const BufferParams &full_params,
    * It is requires to inform about reset whenever it happens, so that the redraw state tracking is
    * properly updated. */
   if (display_) {
-    display_->reset(big_tile_params, reset_rendering);
+    display_->reset(full_params);
   }
 
   render_state_.has_denoised_result = false;
@@ -596,15 +594,6 @@ void PathTrace::draw()
   did_draw_after_reset_ |= display_->draw();
 }
 
-void PathTrace::flush_display()
-{
-  if (!display_) {
-    return;
-  }
-
-  display_->flush();
-}
-
 void PathTrace::update_display(const RenderWork &render_work)
 {
   if (!render_work.display.update) {
@@ -633,8 +622,9 @@ void PathTrace::update_display(const RenderWork &render_work)
   if (display_) {
     VLOG(3) << "Perform copy to GPUDisplay work.";
 
-    const int texture_width = render_state_.effective_big_tile_params.window_width;
+    const int resolution_divider = render_work.resolution_divider;
-    const int texture_height = render_state_.effective_big_tile_params.window_height;
+    const int texture_width = max(1, full_params_.width / resolution_divider);
+    const int texture_height = max(1, full_params_.height / resolution_divider);
     if (!display_->update_begin(texture_width, texture_height)) {
       LOG(ERROR) << "Error beginning GPUDisplay update.";
       return;
@@ -820,15 +810,8 @@ void PathTrace::tile_buffer_read()
     return;
   }
 
-  /* Read buffers back from device. */
-  tbb::parallel_for_each(path_trace_works_, [&](unique_ptr<PathTraceWork> &path_trace_work) {
-    path_trace_work->copy_render_buffers_from_device();
-  });
-
-  /* Read (subset of) passes from output driver. */
   PathTraceTile tile(*this);
   if (output_driver_->read_render_tile(tile)) {
-    /* Copy buffers to device again. */
     tbb::parallel_for_each(path_trace_works_, [](unique_ptr<PathTraceWork> &path_trace_work) {
       path_trace_work->copy_render_buffers_to_device();
     });

intern/cycles/integrator/path_trace.h

@@ -72,9 +72,7 @@ class PathTrace {
    * render result. */
   bool ready_to_reset();
 
-  void reset(const BufferParams &full_params,
+  void reset(const BufferParams &full_params, const BufferParams &big_tile_params);
-             const BufferParams &big_tile_params,
-             bool reset_rendering);
 
   void device_free();
 
@@ -114,9 +112,6 @@ class PathTrace {
   /* Perform drawing of the current state of the DisplayDriver. */
   void draw();
 
-  /* Flush outstanding display commands before ending the render loop. */
-  void flush_display();
-
   /* Cancel rendering process as soon as possible, without waiting for full tile to be sampled.
    * Used in cases like reset of render session.
    *

intern/cycles/integrator/path_trace_display.cpp

@@ -26,20 +26,15 @@ PathTraceDisplay::PathTraceDisplay(unique_ptr<DisplayDriver> driver) : driver_(m
 {
 }
 
-void PathTraceDisplay::reset(const BufferParams &buffer_params, const bool reset_rendering)
+void PathTraceDisplay::reset(const BufferParams &buffer_params)
 {
   thread_scoped_lock lock(mutex_);
 
-  params_.full_offset = make_int2(buffer_params.full_x + buffer_params.window_x,
+  params_.full_offset = make_int2(buffer_params.full_x, buffer_params.full_y);
-                                  buffer_params.full_y + buffer_params.window_y);
   params_.full_size = make_int2(buffer_params.full_width, buffer_params.full_height);
-  params_.size = make_int2(buffer_params.window_width, buffer_params.window_height);
+  params_.size = make_int2(buffer_params.width, buffer_params.height);
 
   texture_state_.is_outdated = true;
-
-  if (!reset_rendering) {
-    driver_->next_tile_begin();
-  }
 }
 
 void PathTraceDisplay::mark_texture_updated()
@@ -253,9 +248,4 @@ bool PathTraceDisplay::draw()
   return !is_outdated;
 }
 
-void PathTraceDisplay::flush()
-{
-  driver_->flush();
-}
-
 CCL_NAMESPACE_END

intern/cycles/integrator/path_trace_display.h

@@ -38,17 +38,14 @@ class BufferParams;
 
 class PathTraceDisplay {
  public:
-  explicit PathTraceDisplay(unique_ptr<DisplayDriver> driver);
+  PathTraceDisplay(unique_ptr<DisplayDriver> driver);
   virtual ~PathTraceDisplay() = default;
 
   /* Reset the display for the new state of render session. Is called whenever session is reset,
    * which happens on changes like viewport navigation or viewport dimension change.
    *
-   * This call will configure parameters for a changed buffer and reset the texture state.
+   * This call will configure parameters for a changed buffer and reset the texture state. */
-   *
+  void reset(const BufferParams &buffer_params);
-   * When the `reset_rendering` a complete display reset happens. When it is false reset happens
-   * for a new state of the buffer parameters which is assumed to correspond to the next tile. */
-  void reset(const BufferParams &buffer_params, bool reset_rendering);
 
   /* --------------------------------------------------------------------
    * Update procedure.
@@ -154,9 +151,6 @@ class PathTraceDisplay {
    * Returns true if this call did draw an updated state of the texture. */
   bool draw();
 
-  /* Flush outstanding display commands before ending the render loop. */
-  void flush();
-
  private:
   /* Display driver implemented by the host application. */
   unique_ptr<DisplayDriver> driver_;

intern/cycles/integrator/path_trace_work.cpp

@@ -194,10 +194,10 @@ PassAccessor::Destination PathTraceWork::get_display_destination_template(
   PassAccessor::Destination destination(film_->get_display_pass());
 
   const int2 display_texture_size = display->get_texture_size();
-  const int texture_x = effective_buffer_params_.full_x - effective_big_tile_params_.full_x +
+  const int texture_x = effective_buffer_params_.full_x - effective_full_params_.full_x +
-                        effective_buffer_params_.window_x - effective_big_tile_params_.window_x;
+                        effective_buffer_params_.window_x;
-  const int texture_y = effective_buffer_params_.full_y - effective_big_tile_params_.full_y +
+  const int texture_y = effective_buffer_params_.full_y - effective_full_params_.full_y +
-                        effective_buffer_params_.window_y - effective_big_tile_params_.window_y;
+                        effective_buffer_params_.window_y;
 
   destination.offset = texture_y * display_texture_size.x + texture_x;
   destination.stride = display_texture_size.x;

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -875,10 +875,8 @@ void PathTraceWorkGPU::copy_to_display_naive(PathTraceDisplay *display,
   const int final_width = buffers_->params.window_width;
   const int final_height = buffers_->params.window_height;
 
-  const int texture_x = full_x - effective_big_tile_params_.full_x +
+  const int texture_x = full_x - effective_full_params_.full_x + effective_buffer_params_.window_x;
-                        effective_buffer_params_.window_x - effective_big_tile_params_.window_x;
+  const int texture_y = full_y - effective_full_params_.full_y + effective_buffer_params_.window_y;
-  const int texture_y = full_y - effective_big_tile_params_.full_y +
-                        effective_buffer_params_.window_y - effective_big_tile_params_.window_y;
 
   /* Re-allocate display memory if needed, and make sure the device pointer is allocated.
    *

intern/cycles/integrator/render_scheduler.cpp

@@ -406,6 +406,9 @@ bool RenderScheduler::set_postprocess_render_work(RenderWork *render_work)
     any_scheduled = true;
   }
 
+  /* Force update. */
+  any_scheduled = true;
+
   if (any_scheduled) {
     render_work->display.update = true;
   }

intern/cycles/integrator/render_scheduler.h

@@ -283,7 +283,7 @@ class RenderScheduler {
   /* Check whether timing report about the given work need to reset accumulated average time. */
   bool work_report_reset_average(const RenderWork &render_work);
 
-  /* Check whether render time limit has been reached (or exceeded), and if so store related
+  /* CHeck whether render time limit has been reached (or exceeded), and if so store related
    * information in the state so that rendering is considered finished, and is possible to report
    * average render time information. */
   void check_time_limit_reached();

intern/cycles/kernel/bvh/bvh.h

@@ -173,16 +173,15 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
   uint p3 = 0;
   uint p4 = visibility;
   uint p5 = PRIMITIVE_NONE;
-  uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
-  uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
 
   uint ray_mask = visibility & 0xFF;
-  uint ray_flags = OPTIX_RAY_FLAG_ENFORCE_ANYHIT;
+  uint ray_flags = OPTIX_RAY_FLAG_NONE;
   if (0 == ray_mask && (visibility & ~0xFF) != 0) {
     ray_mask = 0xFF;
+    ray_flags = OPTIX_RAY_FLAG_ENFORCE_ANYHIT;
   }
   else if (visibility & PATH_RAY_SHADOW_OPAQUE) {
-    ray_flags |= OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT;
+    ray_flags = OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT;
   }
 
   optixTrace(scene_intersect_valid(ray) ? kernel_data.bvh.scene : 0,
@@ -201,9 +200,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
              p2,
              p3,
              p4,
-             p5,
+             p5);
-             p6,
-             p7);
 
   isect->t = __uint_as_float(p0);
   isect->u = __uint_as_float(p1);
@@ -245,7 +242,6 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
   }
 
   MetalRTIntersectionPayload payload;
-  payload.self = ray->self;
   payload.u = 0.0f;
   payload.v = 0.0f;
   payload.visibility = visibility;
@@ -313,7 +309,6 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
     CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR);
     IntersectContext rtc_ctx(&ctx);
     RTCRayHit ray_hit;
-    ctx.ray = ray;
     kernel_embree_setup_rayhit(*ray, ray_hit, visibility);
     rtcIntersect1(kernel_data.bvh.scene, &rtc_ctx.context, &ray_hit);
     if (ray_hit.hit.geomID != RTC_INVALID_GEOMETRY_ID &&
@@ -361,9 +356,6 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
   uint p2 = pointer_pack_to_uint_0(local_isect);
   uint p3 = pointer_pack_to_uint_1(local_isect);
   uint p4 = local_object;
-  uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
-  uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
-
   /* Is set to zero on miss or if ray is aborted, so can be used as return value. */
   uint p5 = max_hits;
 
@@ -387,9 +379,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
              p2,
              p3,
              p4,
-             p5,
+             p5);
-             p6,
-             p7);
 
   return p5;
 #  elif defined(__METALRT__)
@@ -427,7 +417,6 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
   }
 
   MetalRTIntersectionLocalPayload payload;
-  payload.self = ray->self;
   payload.local_object = local_object;
   payload.max_hits = max_hits;
   payload.local_isect.num_hits = 0;
@@ -471,7 +460,6 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
         kg, has_bvh ? CCLIntersectContext::RAY_SSS : CCLIntersectContext::RAY_LOCAL);
     ctx.lcg_state = lcg_state;
     ctx.max_hits = max_hits;
-    ctx.ray = ray;
     ctx.local_isect = local_isect;
     if (local_isect) {
       local_isect->num_hits = 0;
@@ -544,8 +532,6 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
   uint p3 = max_hits;
   uint p4 = visibility;
   uint p5 = false;
-  uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
-  uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
 
   uint ray_mask = visibility & 0xFF;
   if (0 == ray_mask && (visibility & ~0xFF) != 0) {
@@ -569,9 +555,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
              p2,
              p3,
              p4,
-             p5,
+             p5);
-             p6,
-             p7);
 
   *num_recorded_hits = uint16_unpack_from_uint_0(p2);
   *throughput = __uint_as_float(p1);
@@ -604,7 +588,6 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
   }
 
   MetalRTIntersectionShadowPayload payload;
-  payload.self = ray->self;
   payload.visibility = visibility;
   payload.max_hits = max_hits;
   payload.num_hits = 0;
@@ -651,7 +634,6 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
     Intersection *isect_array = (Intersection *)state->shadow_isect;
     ctx.isect_s = isect_array;
     ctx.max_hits = max_hits;
-    ctx.ray = ray;
     IntersectContext rtc_ctx(&ctx);
     RTCRay rtc_ray;
     kernel_embree_setup_ray(*ray, rtc_ray, visibility);
@@ -703,8 +685,6 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
   uint p3 = 0;
   uint p4 = visibility;
   uint p5 = PRIMITIVE_NONE;
-  uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
-  uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
 
   uint ray_mask = visibility & 0xFF;
   if (0 == ray_mask && (visibility & ~0xFF) != 0) {
@@ -728,9 +708,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
              p2,
              p3,
              p4,
-             p5,
+             p5);
-             p6,
-             p7);
 
   isect->t = __uint_as_float(p0);
   isect->u = __uint_as_float(p1);
@@ -766,7 +744,6 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
   }
 
   MetalRTIntersectionPayload payload;
-  payload.self = ray->self;
   payload.visibility = visibility;
 
   typename metalrt_intersector_type::result_type intersection;
@@ -843,7 +820,6 @@ ccl_device_intersect uint scene_intersect_volume_all(KernelGlobals kg,
     ctx.isect_s = isect;
     ctx.max_hits = max_hits;
     ctx.num_hits = 0;
-    ctx.ray = ray;
     IntersectContext rtc_ctx(&ctx);
     RTCRay rtc_ray;
     kernel_embree_setup_ray(*ray, rtc_ray, visibility);

intern/cycles/kernel/bvh/embree.h

@@ -22,8 +22,6 @@
 #include "kernel/device/cpu/compat.h"
 #include "kernel/device/cpu/globals.h"
 
-#include "kernel/bvh/util.h"
-
 #include "util/vector.h"
 
 CCL_NAMESPACE_BEGIN
@@ -40,9 +38,6 @@ struct CCLIntersectContext {
   KernelGlobals kg;
   RayType type;
 
-  /* For avoiding self intersections */
-  const Ray *ray;
-
   /* for shadow rays */
   Intersection *isect_s;
   uint max_hits;
@@ -61,7 +56,6 @@ struct CCLIntersectContext {
   {
     kg = kg_;
     type = type_;
-    ray = NULL;
     max_hits = 1;
     num_hits = 0;
     num_recorded_hits = 0;
@@ -108,34 +102,7 @@ ccl_device_inline void kernel_embree_setup_rayhit(const Ray &ray,
 {
   kernel_embree_setup_ray(ray, rayhit.ray, visibility);
   rayhit.hit.geomID = RTC_INVALID_GEOMETRY_ID;
-  rayhit.hit.instID[0] = RTC_INVALID_GEOMETRY_ID;
+  rayhit.hit.primID = RTC_INVALID_GEOMETRY_ID;
-}
-
-ccl_device_inline bool kernel_embree_is_self_intersection(const KernelGlobals kg,
-                                                          const RTCHit *hit,
-                                                          const Ray *ray)
-{
-  bool status = false;
-  if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) {
-    const int oID = hit->instID[0] / 2;
-    if ((ray->self.object == oID) || (ray->self.light_object == oID)) {
-      RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData(
-          rtcGetGeometry(kernel_data.bvh.scene, hit->instID[0]));
-      const int pID = hit->primID +
-                      (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID));
-      status = intersection_skip_self_shadow(ray->self, oID, pID);
-    }
-  }
-  else {
-    const int oID = hit->geomID / 2;
-    if ((ray->self.object == oID) || (ray->self.light_object == oID)) {
-      const int pID = hit->primID + (intptr_t)rtcGetGeometryUserData(
-                                        rtcGetGeometry(kernel_data.bvh.scene, hit->geomID));
-      status = intersection_skip_self_shadow(ray->self, oID, pID);
-    }
-  }
-
-  return status;
 }
 
 ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg,

intern/cycles/kernel/bvh/local.h

@@ -148,27 +148,12 @@ ccl_device_inline
             /* intersect ray against primitive */
             for (; prim_addr < prim_addr2; prim_addr++) {
               kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
-
-              /* Only intersect with matching object, for instanced objects we
-               * already know we are only intersecting the right object. */
-              if (object == OBJECT_NONE) {
-                if (kernel_tex_fetch(__prim_object, prim_addr) != local_object) {
-                  continue;
-                }
-              }
-
-              /* Skip self intersection. */
-              const int prim = kernel_tex_fetch(__prim_index, prim_addr);
-              if (intersection_skip_self_local(ray->self, prim)) {
-                continue;
-              }
-
               if (triangle_intersect_local(kg,
                                            local_isect,
                                            P,
                                            dir,
+                                           object,
                                            local_object,
-                                           prim,
                                            prim_addr,
                                            isect_t,
                                            lcg_state,
@@ -183,28 +168,13 @@ ccl_device_inline
             /* intersect ray against primitive */
             for (; prim_addr < prim_addr2; prim_addr++) {
               kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
-
-              /* Only intersect with matching object, for instanced objects we
-               * already know we are only intersecting the right object. */
-              if (object == OBJECT_NONE) {
-                if (kernel_tex_fetch(__prim_object, prim_addr) != local_object) {
-                  continue;
-                }
-              }
-
-              /* Skip self intersection. */
-              const int prim = kernel_tex_fetch(__prim_index, prim_addr);
-              if (intersection_skip_self_local(ray->self, prim)) {
-                continue;
-              }
-
               if (motion_triangle_intersect_local(kg,
                                                   local_isect,
                                                   P,
                                                   dir,
                                                   ray->time,
+                                                  object,
                                                   local_object,
-                                                  prim,
                                                   prim_addr,
                                                   isect_t,
                                                   lcg_state,

intern/cycles/kernel/bvh/metal.h

@@ -15,7 +15,6 @@
  */
 
 struct MetalRTIntersectionPayload {
-  RaySelfPrimitives self;
   uint visibility;
   float u, v;
   int prim;
@@ -26,7 +25,6 @@ struct MetalRTIntersectionPayload {
 };
 
 struct MetalRTIntersectionLocalPayload {
-  RaySelfPrimitives self;
   uint local_object;
   uint lcg_state;
   short max_hits;
@@ -36,7 +34,6 @@ struct MetalRTIntersectionLocalPayload {
 };
 
 struct MetalRTIntersectionShadowPayload {
-  RaySelfPrimitives self;
   uint visibility;
 #if defined(__METALRT_MOTION__)
   float time;

intern/cycles/kernel/bvh/shadow_all.h

@@ -146,7 +146,7 @@ ccl_device_inline
           --stack_ptr;
 
           /* primitive intersection */
-          for (; prim_addr < prim_addr2; prim_addr++) {
+          while (prim_addr < prim_addr2) {
             kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) ==
                           (type & PRIMITIVE_ALL));
             bool hit;
@@ -156,32 +156,16 @@ ccl_device_inline
              * might give a few % performance improvement */
             Intersection isect ccl_optional_struct_init;
 
-            const int prim_object = (object == OBJECT_NONE) ?
-                                        kernel_tex_fetch(__prim_object, prim_addr) :
-                                        object;
-            const int prim = kernel_tex_fetch(__prim_index, prim_addr);
-            if (intersection_skip_self_shadow(ray->self, prim_object, prim)) {
-              continue;
-            }
-
             switch (type & PRIMITIVE_ALL) {
               case PRIMITIVE_TRIANGLE: {
                 hit = triangle_intersect(
-                    kg, &isect, P, dir, t_max_current, visibility, prim_object, prim, prim_addr);
+                    kg, &isect, P, dir, t_max_current, visibility, object, prim_addr);
                 break;
               }
 #if BVH_FEATURE(BVH_MOTION)
               case PRIMITIVE_MOTION_TRIANGLE: {
-                hit = motion_triangle_intersect(kg,
+                hit = motion_triangle_intersect(
-                                                &isect,
+                    kg, &isect, P, dir, t_max_current, ray->time, visibility, object, prim_addr);
-                                                P,
-                                                dir,
-                                                t_max_current,
-                                                ray->time,
-                                                visibility,
-                                                prim_object,
-                                                prim,
-                                                prim_addr);
                 break;
               }
 #endif
@@ -198,9 +182,20 @@ ccl_device_inline
                   }
                 }
 
+                const int curve_object = (object == OBJECT_NONE) ?
+                                             kernel_tex_fetch(__prim_object, prim_addr) :
+                                             object;
                 const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
-                hit = curve_intersect(
+                const int curve_prim = kernel_tex_fetch(__prim_index, prim_addr);
-                    kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, curve_type);
+                hit = curve_intersect(kg,
+                                      &isect,
+                                      P,
+                                      dir,
+                                      t_max_current,
+                                      curve_object,
+                                      curve_prim,
+                                      ray->time,
+                                      curve_type);
 
                 break;
               }
@@ -216,9 +211,20 @@ ccl_device_inline
                   }
                 }
 
+                const int point_object = (object == OBJECT_NONE) ?
+                                             kernel_tex_fetch(__prim_object, prim_addr) :
+                                             object;
+                const int point_prim = kernel_tex_fetch(__prim_index, prim_addr);
                 const int point_type = kernel_tex_fetch(__prim_type, prim_addr);
-                hit = point_intersect(
+                hit = point_intersect(kg,
-                    kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, point_type);
+                                      &isect,
+                                      P,
+                                      dir,
+                                      t_max_current,
+                                      point_object,
+                                      point_prim,
+                                      ray->time,
+                                      point_type);
                 break;
               }
 #endif /* BVH_FEATURE(BVH_POINTCLOUD) */
@@ -295,6 +301,8 @@ ccl_device_inline
                 integrator_state_write_shadow_isect(state, &isect, record_index);
               }
             }
+
+            prim_addr++;
           }
         }
         else {

intern/cycles/kernel/bvh/traversal.h

@@ -133,91 +133,89 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
           --stack_ptr;
 
           /* primitive intersection */
-          for (; prim_addr < prim_addr2; prim_addr++) {
+          switch (type & PRIMITIVE_ALL) {
-            kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
+            case PRIMITIVE_TRIANGLE: {
-
+              for (; prim_addr < prim_addr2; prim_addr++) {
-            const int prim_object = (object == OBJECT_NONE) ?
+                kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
-                                        kernel_tex_fetch(__prim_object, prim_addr) :
-                                        object;
-            const int prim = kernel_tex_fetch(__prim_index, prim_addr);
-            if (intersection_skip_self_shadow(ray->self, prim_object, prim)) {
-              continue;
-            }
-
-            switch (type & PRIMITIVE_ALL) {
-              case PRIMITIVE_TRIANGLE: {
                 if (triangle_intersect(
-                        kg, isect, P, dir, isect->t, visibility, prim_object, prim, prim_addr)) {
+                        kg, isect, P, dir, isect->t, visibility, object, prim_addr)) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
                     return true;
                 }
-                break;
               }
+              break;
+            }
 #if BVH_FEATURE(BVH_MOTION)
-              case PRIMITIVE_MOTION_TRIANGLE: {
+            case PRIMITIVE_MOTION_TRIANGLE: {
-                if (motion_triangle_intersect(kg,
+              for (; prim_addr < prim_addr2; prim_addr++) {
-                                              isect,
+                kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
-                                              P,
+                if (motion_triangle_intersect(
-                                              dir,
+                        kg, isect, P, dir, isect->t, ray->time, visibility, object, prim_addr)) {
-                                              isect->t,
-                                              ray->time,
-                                              visibility,
-                                              prim_object,
-                                              prim,
-                                              prim_addr)) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
                     return true;
                 }
-                break;
               }
+              break;
+            }
 #endif /* BVH_FEATURE(BVH_MOTION) */
 #if BVH_FEATURE(BVH_HAIR)
-              case PRIMITIVE_CURVE_THICK:
+            case PRIMITIVE_CURVE_THICK:
-              case PRIMITIVE_MOTION_CURVE_THICK:
+            case PRIMITIVE_MOTION_CURVE_THICK:
-              case PRIMITIVE_CURVE_RIBBON:
+            case PRIMITIVE_CURVE_RIBBON:
-              case PRIMITIVE_MOTION_CURVE_RIBBON: {
+            case PRIMITIVE_MOTION_CURVE_RIBBON: {
+              for (; prim_addr < prim_addr2; prim_addr++) {
                 if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
                   const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
                   if (ray->time < prim_time.x || ray->time > prim_time.y) {
-                    break;
+                    continue;
                   }
                 }
 
+                const int curve_object = (object == OBJECT_NONE) ?
+                                             kernel_tex_fetch(__prim_object, prim_addr) :
+                                             object;
+                const int curve_prim = kernel_tex_fetch(__prim_index, prim_addr);
                 const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
                 const bool hit = curve_intersect(
-                    kg, isect, P, dir, isect->t, prim_object, prim, ray->time, curve_type);
+                    kg, isect, P, dir, isect->t, curve_object, curve_prim, ray->time, curve_type);
                 if (hit) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
                     return true;
                 }
-                break;
               }
+              break;
+            }
 #endif /* BVH_FEATURE(BVH_HAIR) */
 #if BVH_FEATURE(BVH_POINTCLOUD)
-              case PRIMITIVE_POINT:
+            case PRIMITIVE_POINT:
-              case PRIMITIVE_MOTION_POINT: {
+            case PRIMITIVE_MOTION_POINT: {
+              for (; prim_addr < prim_addr2; prim_addr++) {
                 if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
                   const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
                   if (ray->time < prim_time.x || ray->time > prim_time.y) {
-                    break;
+                    continue;
                   }
                 }
 
+                const int point_object = (object == OBJECT_NONE) ?
+                                             kernel_tex_fetch(__prim_object, prim_addr) :
+                                             object;
+                const int point_prim = kernel_tex_fetch(__prim_index, prim_addr);
                 const int point_type = kernel_tex_fetch(__prim_type, prim_addr);
                 const bool hit = point_intersect(
-                    kg, isect, P, dir, isect->t, prim_object, prim, ray->time, point_type);
+                    kg, isect, P, dir, isect->t, point_object, point_prim, ray->time, point_type);
                 if (hit) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
                     return true;
                 }
-                break;
               }
-#endif /* BVH_FEATURE(BVH_POINTCLOUD) */
+              break;
             }
+#endif /* BVH_FEATURE(BVH_POINTCLOUD) */
           }
         }
         else {

intern/cycles/kernel/bvh/util.h

@@ -21,22 +21,54 @@ CCL_NAMESPACE_BEGIN
 /* Ray offset to avoid self intersection.
  *
  * This function should be used to compute a modified ray start position for
- * rays leaving from a surface. This is from "A Fast and Robust Method for Avoiding
+ * rays leaving from a surface. */
- * Self-Intersection" see https://research.nvidia.com/publication/2019-03_A-Fast-and
+
- */
 ccl_device_inline float3 ray_offset(float3 P, float3 Ng)
 {
-  const float int_scale = 256.0f;
+#ifdef __INTERSECTION_REFINE__
-  int3 of_i = make_int3((int)(int_scale * Ng.x), (int)(int_scale * Ng.y), (int)(int_scale * Ng.z));
+  const float epsilon_f = 1e-5f;
+  /* ideally this should match epsilon_f, but instancing and motion blur
+   * precision makes it problematic */
+  const float epsilon_test = 1.0f;
+  const int epsilon_i = 32;
 
-  float3 p_i = make_float3(__int_as_float(__float_as_int(P.x) + ((P.x < 0) ? -of_i.x : of_i.x)),
+  float3 res;
-                           __int_as_float(__float_as_int(P.y) + ((P.y < 0) ? -of_i.y : of_i.y)),
+
-                           __int_as_float(__float_as_int(P.z) + ((P.z < 0) ? -of_i.z : of_i.z)));
+  /* x component */
-  const float origin = 1.0f / 32.0f;
+  if (fabsf(P.x) < epsilon_test) {
-  const float float_scale = 1.0f / 65536.0f;
+    res.x = P.x + Ng.x * epsilon_f;
-  return make_float3(fabsf(P.x) < origin ? P.x + float_scale * Ng.x : p_i.x,
+  }
-                     fabsf(P.y) < origin ? P.y + float_scale * Ng.y : p_i.y,
+  else {
-                     fabsf(P.z) < origin ? P.z + float_scale * Ng.z : p_i.z);
+    uint ix = __float_as_uint(P.x);
+    ix += ((ix ^ __float_as_uint(Ng.x)) >> 31) ? -epsilon_i : epsilon_i;
+    res.x = __uint_as_float(ix);
+  }
+
+  /* y component */
+  if (fabsf(P.y) < epsilon_test) {
+    res.y = P.y + Ng.y * epsilon_f;
+  }
+  else {
+    uint iy = __float_as_uint(P.y);
+    iy += ((iy ^ __float_as_uint(Ng.y)) >> 31) ? -epsilon_i : epsilon_i;
+    res.y = __uint_as_float(iy);
+  }
+
+  /* z component */
+  if (fabsf(P.z) < epsilon_test) {
+    res.z = P.z + Ng.z * epsilon_f;
+  }
+  else {
+    uint iz = __float_as_uint(P.z);
+    iz += ((iz ^ __float_as_uint(Ng.z)) >> 31) ? -epsilon_i : epsilon_i;
+    res.z = __uint_as_float(iz);
+  }
+
+  return res;
+#else
+  const float epsilon_f = 1e-4f;
+  return P + epsilon_f * Ng;
+#endif
 }
 
 #if defined(__KERNEL_CPU__)
@@ -195,25 +227,4 @@ ccl_device_inline float intersection_curve_shadow_transparency(KernelGlobals kg,
   return (1.0f - u) * f0 + u * f1;
 }
 
-ccl_device_inline bool intersection_skip_self(ccl_private const RaySelfPrimitives &self,
-                                              const int object,
-                                              const int prim)
-{
-  return (self.prim == prim) && (self.object == object);
-}
-
-ccl_device_inline bool intersection_skip_self_shadow(ccl_private const RaySelfPrimitives &self,
-                                                     const int object,
-                                                     const int prim)
-{
-  return ((self.prim == prim) && (self.object == object)) ||
-         ((self.light_prim == prim) && (self.light_object == object));
-}
-
-ccl_device_inline bool intersection_skip_self_local(ccl_private const RaySelfPrimitives &self,
-                                                    const int prim)
-{
-  return (self.prim == prim);
-}
-
 CCL_NAMESPACE_END

intern/cycles/kernel/bvh/volume.h

@@ -140,20 +140,14 @@ ccl_device_inline
               for (; prim_addr < prim_addr2; prim_addr++) {
                 kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
                 /* only primitives from volume object */
-                const int prim_object = (object == OBJECT_NONE) ?
+                uint tri_object = (object == OBJECT_NONE) ?
-                                            kernel_tex_fetch(__prim_object, prim_addr) :
+                                      kernel_tex_fetch(__prim_object, prim_addr) :
-                                            object;
+                                      object;
-                const int prim = kernel_tex_fetch(__prim_index, prim_addr);
+                int object_flag = kernel_tex_fetch(__object_flag, tri_object);
-                if (intersection_skip_self(ray->self, prim_object, prim)) {
-                  continue;
-                }
-
-                int object_flag = kernel_tex_fetch(__object_flag, prim_object);
                 if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
                   continue;
                 }
-                triangle_intersect(
+                triangle_intersect(kg, isect, P, dir, isect->t, visibility, object, prim_addr);
-                    kg, isect, P, dir, isect->t, visibility, prim_object, prim, prim_addr);
               }
               break;
             }
@@ -163,27 +157,15 @@ ccl_device_inline
               for (; prim_addr < prim_addr2; prim_addr++) {
                 kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
                 /* only primitives from volume object */
-                const int prim_object = (object == OBJECT_NONE) ?
+                uint tri_object = (object == OBJECT_NONE) ?
-                                            kernel_tex_fetch(__prim_object, prim_addr) :
+                                      kernel_tex_fetch(__prim_object, prim_addr) :
-                                            object;
+                                      object;
-                const int prim = kernel_tex_fetch(__prim_index, prim_addr);
+                int object_flag = kernel_tex_fetch(__object_flag, tri_object);
-                if (intersection_skip_self(ray->self, prim_object, prim)) {
-                  continue;
-                }
-                int object_flag = kernel_tex_fetch(__object_flag, prim_object);
                 if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
                   continue;
                 }
-                motion_triangle_intersect(kg,
+                motion_triangle_intersect(
-                                          isect,
+                    kg, isect, P, dir, isect->t, ray->time, visibility, object, prim_addr);
-                                          P,
-                                          dir,
-                                          isect->t,
-                                          ray->time,
-                                          visibility,
-                                          prim_object,
-                                          prim,
-                                          prim_addr);
               }
               break;
             }

intern/cycles/kernel/bvh/volume_all.h

@@ -143,19 +143,15 @@ ccl_device_inline
               for (; prim_addr < prim_addr2; prim_addr++) {
                 kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
                 /* only primitives from volume object */
-                const int prim_object = (object == OBJECT_NONE) ?
+                uint tri_object = (object == OBJECT_NONE) ?
-                                            kernel_tex_fetch(__prim_object, prim_addr) :
+                                      kernel_tex_fetch(__prim_object, prim_addr) :
-                                            object;
+                                      object;
-                const int prim = kernel_tex_fetch(__prim_index, prim_addr);
+                int object_flag = kernel_tex_fetch(__object_flag, tri_object);
-                if (intersection_skip_self(ray->self, prim_object, prim)) {
-                  continue;
-                }
-                int object_flag = kernel_tex_fetch(__object_flag, prim_object);
                 if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
                   continue;
                 }
                 hit = triangle_intersect(
-                    kg, isect_array, P, dir, isect_t, visibility, prim_object, prim, prim_addr);
+                    kg, isect_array, P, dir, isect_t, visibility, object, prim_addr);
                 if (hit) {
                   /* Move on to next entry in intersections array. */
                   isect_array++;
@@ -187,27 +183,15 @@ ccl_device_inline
               for (; prim_addr < prim_addr2; prim_addr++) {
                 kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
                 /* only primitives from volume object */
-                const int prim_object = (object == OBJECT_NONE) ?
+                uint tri_object = (object == OBJECT_NONE) ?
-                                            kernel_tex_fetch(__prim_object, prim_addr) :
+                                      kernel_tex_fetch(__prim_object, prim_addr) :
-                                            object;
+                                      object;
-                const int prim = kernel_tex_fetch(__prim_index, prim_addr);
+                int object_flag = kernel_tex_fetch(__object_flag, tri_object);
-                if (intersection_skip_self(ray->self, prim_object, prim)) {
-                  continue;
-                }
-                int object_flag = kernel_tex_fetch(__object_flag, prim_object);
                 if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
                   continue;
                 }
-                hit = motion_triangle_intersect(kg,
+                hit = motion_triangle_intersect(
-                                                isect_array,
+                    kg, isect_array, P, dir, isect_t, ray->time, visibility, object, prim_addr);
-                                                P,
-                                                dir,
-                                                isect_t,
-                                                ray->time,
-                                                visibility,
-                                                prim_object,
-                                                prim,
-                                                prim_addr);
                 if (hit) {
                   /* Move on to next entry in intersections array. */
                   isect_array++;

intern/cycles/kernel/closure/volume.h

@@ -199,18 +199,22 @@ ccl_device int volume_sample_channel(float3 albedo,
    *  Tracing". Matt Jen-Yuan Chiang, Peter Kutz, Brent Burley. SIGGRAPH 2016. */
   float3 weights = fabs(throughput * albedo);
   float sum_weights = weights.x + weights.y + weights.z;
+  float3 weights_pdf;
 
   if (sum_weights > 0.0f) {
-    *pdf = weights / sum_weights;
+    weights_pdf = weights / sum_weights;
   }
   else {
-    *pdf = make_float3(1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f);
+    weights_pdf = make_float3(1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f);
   }
 
-  if (rand < pdf->x) {
+  *pdf = weights_pdf;
+
+  /* OpenCL does not support -> on float3, so don't use pdf->x. */
+  if (rand < weights_pdf.x) {
     return 0;
   }
-  else if (rand < pdf->x + pdf->y) {
+  else if (rand < weights_pdf.x + weights_pdf.y) {
     return 1;
   }
   else {

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

@@ -243,10 +243,6 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
   }
 }
 
-#ifdef __KERNEL_METAL__
-constant int __dummy_constant [[function_constant(0)]];
-#endif
-
 ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
     ccl_gpu_kernel_signature(integrator_shade_surface_raytrace,
                              ccl_global const int *path_index_array,
@@ -257,16 +253,7 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
 
   if (global_index < work_size) {
     const int state = (path_index_array) ? path_index_array[global_index] : global_index;
-
-#ifdef __KERNEL_METAL__
-    KernelGlobals kg = NULL;
-    /* Workaround Ambient Occlusion and Bevel nodes not working with Metal.
-     * Dummy offset should not affect result, but somehow fixes bug! */
-    kg += __dummy_constant;
-    ccl_gpu_kernel_call(integrator_shade_surface_raytrace(kg, state, render_buffer));
-#else
     ccl_gpu_kernel_call(integrator_shade_surface_raytrace(NULL, state, render_buffer));
-#endif
   }
 }
 
@@ -769,7 +756,6 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
                              int guiding_pass_stride,
                              int guiding_pass_albedo,
                              int guiding_pass_normal,
-                             int guiding_pass_flow,
                              ccl_global const float *render_buffer,
                              int render_offset,
                              int render_stride,
@@ -777,7 +763,6 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
                              int render_pass_sample_count,
                              int render_pass_denoising_albedo,
                              int render_pass_denoising_normal,
-                             int render_pass_motion,
                              int full_x,
                              int full_y,
                              int width,
@@ -829,17 +814,6 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
     normal_out[1] = normal_in[1] * pixel_scale;
     normal_out[2] = normal_in[2] * pixel_scale;
   }
-
-  /* Flow pass. */
-  if (guiding_pass_flow != PASS_UNUSED) {
-    kernel_assert(render_pass_motion != PASS_UNUSED);
-
-    ccl_global const float *motion_in = buffer + render_pass_motion;
-    ccl_global float *flow_out = guiding_pixel + guiding_pass_flow;
-
-    flow_out[0] = -motion_in[0] * pixel_scale;
-    flow_out[1] = -motion_in[1] * pixel_scale;
-  }
 }
 
 ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
@@ -925,6 +899,7 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
   else {
     /* Assigning to zero since this is a default alpha value for 3-component passes, and it
      * is an opaque pixel for 4 component passes. */
+
     denoised_pixel[3] = 0;
   }
 }

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

@@ -98,12 +98,8 @@ using namespace metal::raytracing;
 #define FN14(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14;
 #define FN15(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15;
 #define FN16(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16;
-#define FN17(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, p17) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16; p17;
+#define GET_LAST_ARG(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, ...) p16
-#define FN18(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, p17, p18) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16; p17; p18;
+#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)
-#define FN19(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, p17, p18, p19) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16; p17; p18; p19;
-#define FN20(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, p17, p18, p19, p20) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16; p17; p18; p19; p20;
-#define GET_LAST_ARG(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, p17, p18, p19, p20, ...) p20
-#define PARAMS_MAKER(...) GET_LAST_ARG(__VA_ARGS__, FN20, FN19, FN18, FN17, 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-> */

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

@@ -40,27 +40,6 @@ struct TriangleIntersectionResult
 
 enum { METALRT_HIT_TRIANGLE, METALRT_HIT_BOUNDING_BOX };
 
-ccl_device_inline bool intersection_skip_self(ray_data const RaySelfPrimitives& self,
-                                              const int object,
-                                              const int prim)
-{
-  return (self.prim == prim) && (self.object == object);
-}
-
-ccl_device_inline bool intersection_skip_self_shadow(ray_data const RaySelfPrimitives& self,
-                                                     const int object,
-                                                     const int prim)
-{
-  return ((self.prim == prim) && (self.object == object)) ||
-         ((self.light_prim == prim) && (self.light_object == object));
-}
-
-ccl_device_inline bool intersection_skip_self_local(ray_data const RaySelfPrimitives& self,
-                                                    const int prim)
-{
-  return (self.prim == prim);
-}
-
 template<typename TReturn, uint intersection_type>
 TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
                           ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload,
@@ -74,8 +53,8 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
 #ifdef __BVH_LOCAL__
   uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
 
-  if ((object != payload.local_object) || intersection_skip_self_local(payload.self, prim)) {
+  if (object != payload.local_object) {
-    /* Only intersect with matching object and skip self-intersecton. */
+    /* Only intersect with matching object */
     result.accept = false;
     result.continue_search = true;
     return result;
@@ -187,11 +166,6 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
   }
 #  endif
 
-  if (intersection_skip_self_shadow(payload.self, object, prim)) {
-    /* continue search */
-    return true;
-  }
-
   float u = 0.0f, v = 0.0f;
   int type = 0;
   if (intersection_type == METALRT_HIT_TRIANGLE) {
@@ -348,35 +322,21 @@ inline TReturnType metalrt_visibility_test(constant KernelParamsMetal &launch_pa
   }
 #  endif
 
-  uint visibility = payload.visibility;
 #  ifdef __VISIBILITY_FLAG__
+  uint visibility = payload.visibility;
   if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
     result.accept = false;
     result.continue_search = true;
     return result;
   }
-#  endif
 
   /* Shadow ray early termination. */
   if (visibility & PATH_RAY_SHADOW_OPAQUE) {
-    if (intersection_skip_self_shadow(payload.self, object, prim)) {
+    result.accept = true;
-      result.accept = false;
+    result.continue_search = false;
-      result.continue_search = true;
+    return result;
-      return result;
-    }
-    else {
-      result.accept = true;
-      result.continue_search = false;
-      return result;
-    }
-  }
-  else {
-    if (intersection_skip_self(payload.self, object, prim)) {
-      result.accept = false;
-      result.continue_search = true;
-      return result;
-    }
   }
+#  endif
 
   result.accept = true;
   result.continue_search = true;
@@ -616,150 +576,6 @@ __intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal
 
   return result;
 }
+
 #endif /* __HAIR__ */
-
-#ifdef __POINTCLOUD__
-ccl_device_inline
-void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal,
-                                ray_data MetalKernelContext::MetalRTIntersectionPayload &payload,
-                                const uint object,
-                                const uint prim,
-                                const uint type,
-                                const float3 ray_origin,
-                                const float3 ray_direction,
-                                float time,
-                                const float ray_tmax,
-                                thread BoundingBoxIntersectionResult &result)
-{
-#  ifdef __VISIBILITY_FLAG__
-  const uint visibility = payload.visibility;
-  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
-    return;
-  }
-#  endif
-
-  float3 P = ray_origin;
-  float3 dir = ray_direction;
-
-  /* The direction is not normalized by default, but the point intersection routine expects that */
-  float len;
-  dir = normalize_len(dir, &len);
-
-  Intersection isect;
-  isect.t = ray_tmax;
-  /* Transform maximum distance into object space. */
-  if (isect.t != FLT_MAX)
-    isect.t *= len;
-
-  MetalKernelContext context(launch_params_metal);
-  if (context.point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
-    result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
-                  launch_params_metal, payload, object, prim, isect.u);
-    if (result.accept) {
-      result.distance = isect.t / len;
-      payload.u = isect.u;
-      payload.v = isect.v;
-      payload.prim = prim;
-      payload.type = type;
-    }
-  }
-}
-
-ccl_device_inline
-void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params_metal,
-                                       ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload,
-                                       const uint object,
-                                       const uint prim,
-                                       const uint type,
-                                       const float3 ray_origin,
-                                       const float3 ray_direction,
-                                       float time,
-                                       const float ray_tmax,
-                                       thread BoundingBoxIntersectionResult &result)
-{
-  const uint visibility = payload.visibility;
-
-  float3 P = ray_origin;
-  float3 dir = ray_direction;
-
-  /* The direction is not normalized by default, but the point intersection routine expects that */
-  float len;
-  dir = normalize_len(dir, &len);
-
-  Intersection isect;
-  isect.t = ray_tmax;
-  /* Transform maximum distance into object space */
-  if (isect.t != FLT_MAX)
-    isect.t *= len;
-
-  MetalKernelContext context(launch_params_metal);
-  if (context.point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
-    result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>(
-                launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax);
-    result.accept = !result.continue_search;
-
-    if (result.accept) {
-      result.distance = isect.t / len;
-    }
-  }
-}
-
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]]
-BoundingBoxIntersectionResult
-__intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
-                             ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]],
-                             const uint object [[user_instance_id]],
-                             const uint primitive_id [[primitive_id]],
-                             const float3 ray_origin [[origin]],
-                             const float3 ray_direction [[direction]],
-                             const float ray_tmax [[max_distance]])
-{
-  const uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
-  const int type = kernel_tex_fetch(__objects, object).primitive_type;
-
-  BoundingBoxIntersectionResult result;
-  result.accept = false;
-  result.continue_search = true;
-  result.distance = ray_tmax;
-
-  metalrt_intersection_point(launch_params_metal, payload, object, prim, type, ray_origin, ray_direction,
-#  if defined(__METALRT_MOTION__)
-                             payload.time,
-#  else
-                             0.0f,
-#  endif
-                             ray_tmax, result);
-
-  return result;
-}
-
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]]
-BoundingBoxIntersectionResult
-__intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
-                                    ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]],
-                                    const uint object [[user_instance_id]],
-                                    const uint primitive_id [[primitive_id]],
-                                    const float3 ray_origin [[origin]],
-                                    const float3 ray_direction [[direction]],
-                                    const float ray_tmax [[max_distance]])
-{
-  const uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
-  const int type = kernel_tex_fetch(__objects, object).primitive_type;
-
-  BoundingBoxIntersectionResult result;
-  result.accept = false;
-  result.continue_search = true;
-  result.distance = ray_tmax;
-
-  metalrt_intersection_point_shadow(launch_params_metal, payload, object, prim, type, ray_origin, ray_direction,
-#  if defined(__METALRT_MOTION__)
-                             payload.time,
-#  else
-                             0.0f,
-#  endif
-                             ray_tmax, result);
-
-  return result;
-}
-#endif /* __POINTCLOUD__ */
 #endif /* __METALRT__ */

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

@@ -45,11 +45,6 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
   return pointer_unpack_from_uint<T>(optixGetPayload_2(), optixGetPayload_3());
 }
 
-template<typename T> ccl_device_forceinline T *get_payload_ptr_6()
-{
-  return (T *)(((uint64_t)optixGetPayload_7() << 32) | optixGetPayload_6());
-}
-
 ccl_device_forceinline int get_object_id()
 {
 #ifdef __OBJECT_MOTION__
@@ -116,12 +111,6 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
     return optixIgnoreIntersection();
   }
 
-  const int prim = optixGetPrimitiveIndex();
-  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
-  if (intersection_skip_self_local(ray->self, prim)) {
-    return optixIgnoreIntersection();
-  }
-
   const uint max_hits = optixGetPayload_5();
   if (max_hits == 0) {
     /* Special case for when no hit information is requested, just report that something was hit */
@@ -160,6 +149,8 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
     local_isect->num_hits = 1;
   }
 
+  const int prim = optixGetPrimitiveIndex();
+
   Intersection *isect = &local_isect->hits[hit];
   isect->t = optixGetRayTmax();
   isect->prim = prim;
@@ -194,11 +185,6 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
   }
 #  endif
 
-  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
-  if (intersection_skip_self_shadow(ray->self, object, prim)) {
-    return optixIgnoreIntersection();
-  }
-
   float u = 0.0f, v = 0.0f;
   int type = 0;
   if (optixIsTriangleHit()) {
@@ -328,12 +314,6 @@ extern "C" __global__ void __anyhit__kernel_optix_volume_test()
   if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
     return optixIgnoreIntersection();
   }
-
-  const int prim = optixGetPrimitiveIndex();
-  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
-  if (intersection_skip_self(ray->self, object, prim)) {
-    return optixIgnoreIntersection();
-  }
 }
 
 extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
@@ -350,31 +330,18 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
 #  endif
 #endif
 
+#ifdef __VISIBILITY_FLAG__
   const uint object = get_object_id();
   const uint visibility = optixGetPayload_4();
-#ifdef __VISIBILITY_FLAG__
   if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
     return optixIgnoreIntersection();
   }
-#endif
-
-  const int prim = optixGetPrimitiveIndex();
-  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
 
+  /* Shadow ray early termination. */
   if (visibility & PATH_RAY_SHADOW_OPAQUE) {
-    if (intersection_skip_self_shadow(ray->self, object, prim)) {
+    return optixTerminateRay();
-      return optixIgnoreIntersection();
-    }
-    else {
-      /* Shadow ray early termination. */
-      return optixTerminateRay();
-    }
-  }
-  else {
-    if (intersection_skip_self(ray->self, object, prim)) {
-      return optixIgnoreIntersection();
-    }
   }
+#endif
 }
 
 extern "C" __global__ void __closesthit__kernel_optix_hit()

intern/cycles/kernel/film/passes.h

@@ -92,14 +92,6 @@ ccl_device_forceinline void kernel_write_denoising_features_surface(
     else if (sc->type == CLOSURE_BSDF_HAIR_PRINCIPLED_ID) {
       closure_albedo *= bsdf_principled_hair_albedo(sc);
     }
-    else if (sc->type == CLOSURE_BSDF_PRINCIPLED_DIFFUSE_ID) {
-      /* BSSRDF already accounts for weight, retro-reflection would double up. */
-      ccl_private const PrincipledDiffuseBsdf *bsdf = (ccl_private const PrincipledDiffuseBsdf *)
-          sc;
-      if (bsdf->components == PRINCIPLED_DIFFUSE_RETRO_REFLECTION) {
-        continue;
-      }
-    }
 
     if (bsdf_get_specular_roughness_squared(sc) > sqr(0.075f)) {
       diffuse_albedo += closure_albedo;

intern/cycles/kernel/film/read.h

@@ -214,21 +214,6 @@ ccl_device_inline void film_get_pass_pixel_light_path(
   pixel[0] = f.x;
   pixel[1] = f.y;
   pixel[2] = f.z;
-
-  /* Optional alpha channel. */
-  if (kfilm_convert->num_components >= 4) {
-    if (kfilm_convert->pass_combined != PASS_UNUSED) {
-      float scale, scale_exposure;
-      film_get_scale_and_scale_exposure(kfilm_convert, buffer, &scale, &scale_exposure);
-
-      ccl_global const float *in_combined = buffer + kfilm_convert->pass_combined;
-      const float alpha = in_combined[3] * scale;
-      pixel[3] = film_transparency_to_alpha(alpha);
-    }
-    else {
-      pixel[3] = 1.0f;
-    }
-  }
 }
 
 ccl_device_inline void film_get_pass_pixel_float3(ccl_global const KernelFilmConvert *ccl_restrict

intern/cycles/kernel/geom/curve.h

@@ -226,18 +226,6 @@ ccl_device float curve_thickness(KernelGlobals kg, ccl_private const ShaderData
   return r * 2.0f;
 }
 
-/* Curve random */
-
-ccl_device float curve_random(KernelGlobals kg, ccl_private const ShaderData *sd)
-{
-  if (sd->type & PRIMITIVE_CURVE) {
-    const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_CURVE_RANDOM);
-    return (desc.offset != ATTR_STD_NOT_FOUND) ? curve_attribute_float(kg, sd, desc, NULL, NULL) :
-                                                 0.0f;
-  }
-  return 0.0f;
-}
-
 /* Curve location for motion pass, linear interpolation between keys and
  * ignoring radius because we do the same for the motion keys */
 

intern/cycles/kernel/geom/motion_triangle.h

@@ -116,52 +116,6 @@ ccl_device_inline void motion_triangle_vertices(
   verts[2] = (1.0f - t) * verts[2] + t * next_verts[2];
 }
 
-ccl_device_inline void motion_triangle_vertices_and_normals(
-    KernelGlobals kg, int object, int prim, float time, float3 verts[3], float3 normals[3])
-{
-  /* get motion info */
-  int numsteps, numverts;
-  object_motion_info(kg, object, &numsteps, &numverts, NULL);
-
-  /* Figure out which steps we need to fetch and their interpolation factor. */
-  int maxstep = numsteps * 2;
-  int step = min((int)(time * maxstep), maxstep - 1);
-  float t = time * maxstep - step;
-
-  /* Find attribute. */
-  int offset = intersection_find_attribute(kg, object, ATTR_STD_MOTION_VERTEX_POSITION);
-  kernel_assert(offset != ATTR_STD_NOT_FOUND);
-
-  /* Fetch vertex coordinates. */
-  float3 next_verts[3];
-  uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
-
-  motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step, verts);
-  motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step + 1, next_verts);
-
-  /* Interpolate between steps. */
-  verts[0] = (1.0f - t) * verts[0] + t * next_verts[0];
-  verts[1] = (1.0f - t) * verts[1] + t * next_verts[1];
-  verts[2] = (1.0f - t) * verts[2] + t * next_verts[2];
-
-  /* Compute smooth normal. */
-
-  /* Find attribute. */
-  offset = intersection_find_attribute(kg, object, ATTR_STD_MOTION_VERTEX_NORMAL);
-  kernel_assert(offset != ATTR_STD_NOT_FOUND);
-
-  /* Fetch vertex coordinates. */
-  float3 next_normals[3];
-  motion_triangle_normals_for_step(kg, tri_vindex, offset, numverts, numsteps, step, normals);
-  motion_triangle_normals_for_step(
-      kg, tri_vindex, offset, numverts, numsteps, step + 1, next_normals);
-
-  /* Interpolate between steps. */
-  normals[0] = (1.0f - t) * normals[0] + t * next_normals[0];
-  normals[1] = (1.0f - t) * normals[1] + t * next_normals[1];
-  normals[2] = (1.0f - t) * normals[2] + t * next_normals[2];
-}
-
 ccl_device_inline float3 motion_triangle_smooth_normal(
     KernelGlobals kg, float3 Ng, int object, int prim, float u, float v, float time)
 {

intern/cycles/kernel/geom/motion_triangle_intersect.h

@@ -29,19 +29,46 @@
 
 CCL_NAMESPACE_BEGIN
 
-/**
+/* Refine triangle intersection to more precise hit point. For rays that travel
- * Use the barycentric coordinates to get the intersection location
+ * far the precision is often not so good, this reintersects the primitive from
+ * a closer distance.
  */
-ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg,
+
-                                                       ccl_private ShaderData *sd,
+ccl_device_inline float3 motion_triangle_refine(KernelGlobals kg,
-                                                       const int isect_object,
+                                                ccl_private ShaderData *sd,
-                                                       const int isect_prim,
+                                                float3 P,
-                                                       const float u,
+                                                float3 D,
-                                                       const float v,
+                                                float t,
-                                                       float3 verts[3])
+                                                const int isect_object,
+                                                const int isect_prim,
+                                                float3 verts[3])
 {
-  float w = 1.0f - u - v;
+#ifdef __INTERSECTION_REFINE__
-  float3 P = u * verts[0] + v * verts[1] + w * verts[2];
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+    if (UNLIKELY(t == 0.0f)) {
+      return P;
+    }
+    const Transform tfm = object_get_inverse_transform(kg, sd);
+
+    P = transform_point(&tfm, P);
+    D = transform_direction(&tfm, D * t);
+    D = normalize_len(D, &t);
+  }
+
+  P = P + D * t;
+
+  /* Compute refined intersection distance. */
+  const float3 e1 = verts[0] - verts[2];
+  const float3 e2 = verts[1] - verts[2];
+  const float3 s1 = cross(D, e2);
+
+  const float invdivisor = 1.0f / dot(s1, e1);
+  const float3 d = P - verts[2];
+  const float3 s2 = cross(d, e1);
+  float rt = dot(e2, s2) * invdivisor;
+
+  /* Compute refined position. */
+  P = P + D * rt;
 
   if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
@@ -49,8 +76,71 @@ ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg,
   }
 
   return P;
+#else
+  return P + D * t;
+#endif
 }
 
+/* Same as above, except that t is assumed to be in object space
+ * for instancing.
+ */
+
+#ifdef __BVH_LOCAL__
+#  if defined(__KERNEL_CUDA__) && (defined(i386) || defined(_M_IX86))
+ccl_device_noinline
+#  else
+ccl_device_inline
+#  endif
+    float3
+    motion_triangle_refine_local(KernelGlobals kg,
+                                 ccl_private ShaderData *sd,
+                                 float3 P,
+                                 float3 D,
+                                 float t,
+                                 const int isect_object,
+                                 const int isect_prim,
+                                 float3 verts[3])
+{
+#  if defined(__KERNEL_GPU_RAYTRACING__)
+  /* t is always in world space with OptiX and MetalRT. */
+  return motion_triangle_refine(kg, sd, P, D, t, isect_object, isect_prim, verts);
+#  else
+#    ifdef __INTERSECTION_REFINE__
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+    const Transform tfm = object_get_inverse_transform(kg, sd);
+
+    P = transform_point(&tfm, P);
+    D = transform_direction(&tfm, D);
+    D = normalize(D);
+  }
+
+  P = P + D * t;
+
+  /* compute refined intersection distance */
+  const float3 e1 = verts[0] - verts[2];
+  const float3 e2 = verts[1] - verts[2];
+  const float3 s1 = cross(D, e2);
+
+  const float invdivisor = 1.0f / dot(s1, e1);
+  const float3 d = P - verts[2];
+  const float3 s2 = cross(d, e1);
+  float rt = dot(e2, s2) * invdivisor;
+
+  P = P + D * rt;
+
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+    const Transform tfm = object_get_transform(kg, sd);
+    P = transform_point(&tfm, P);
+  }
+
+  return P;
+#    else  /* __INTERSECTION_REFINE__ */
+  return P + D * t;
+#    endif /* __INTERSECTION_REFINE__ */
+#  endif
+}
+#endif /* __BVH_LOCAL__ */
+
 /* Ray intersection. We simply compute the vertex positions at the given ray
  * time and do a ray intersection with the resulting triangle.
  */
@@ -63,12 +153,14 @@ ccl_device_inline bool motion_triangle_intersect(KernelGlobals kg,
                                                  float time,
                                                  uint visibility,
                                                  int object,
-                                                 int prim,
                                                  int prim_addr)
 {
+  /* Primitive index for vertex location lookup. */
+  int prim = kernel_tex_fetch(__prim_index, prim_addr);
+  int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) : object;
   /* Get vertex locations for intersection. */
   float3 verts[3];
-  motion_triangle_vertices(kg, object, prim, time, verts);
+  motion_triangle_vertices(kg, fobject, prim, time, verts);
   /* Ray-triangle intersection, unoptimized. */
   float t, u, v;
   if (ray_triangle_intersect(P, dir, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) {
@@ -83,7 +175,8 @@ ccl_device_inline bool motion_triangle_intersect(KernelGlobals kg,
       isect->u = u;
       isect->v = v;
       isect->prim = prim;
-      isect->object = object;
+      isect->object = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) :
+                                                object;
       isect->type = PRIMITIVE_MOTION_TRIANGLE;
       return true;
     }
@@ -103,15 +196,25 @@ ccl_device_inline bool motion_triangle_intersect_local(KernelGlobals kg,
                                                        float3 dir,
                                                        float time,
                                                        int object,
-                                                       int prim,
+                                                       int local_object,
                                                        int prim_addr,
                                                        float tmax,
                                                        ccl_private uint *lcg_state,
                                                        int max_hits)
 {
+  /* Only intersect with matching object, for instanced objects we
+   * already know we are only intersecting the right object. */
+  if (object == OBJECT_NONE) {
+    if (kernel_tex_fetch(__prim_object, prim_addr) != local_object) {
+      return false;
+    }
+  }
+
+  /* Primitive index for vertex location lookup. */
+  int prim = kernel_tex_fetch(__prim_index, prim_addr);
   /* Get vertex locations for intersection. */
   float3 verts[3];
-  motion_triangle_vertices(kg, object, prim, time, verts);
+  motion_triangle_vertices(kg, local_object, prim, time, verts);
   /* Ray-triangle intersection, unoptimized. */
   float t, u, v;
   if (!ray_triangle_intersect(P, dir, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) {
@@ -163,7 +266,7 @@ ccl_device_inline bool motion_triangle_intersect_local(KernelGlobals kg,
   isect->u = u;
   isect->v = v;
   isect->prim = prim;
-  isect->object = object;
+  isect->object = local_object;
   isect->type = PRIMITIVE_MOTION_TRIANGLE;
 
   /* Record geometric normal. */

intern/cycles/kernel/geom/motion_triangle_shader.h

@@ -68,7 +68,15 @@ ccl_device_noinline void motion_triangle_shader_setup(KernelGlobals kg,
   verts[1] = (1.0f - t) * verts[1] + t * next_verts[1];
   verts[2] = (1.0f - t) * verts[2] + t * next_verts[2];
   /* Compute refined position. */
-  sd->P = motion_triangle_point_from_uv(kg, sd, isect_object, isect_prim, sd->u, sd->v, verts);
+#ifdef __BVH_LOCAL__
+  if (is_local) {
+    sd->P = motion_triangle_refine_local(kg, sd, P, D, ray_t, isect_object, isect_prim, verts);
+  }
+  else
+#endif /* __BVH_LOCAL__*/
+  {
+    sd->P = motion_triangle_refine(kg, sd, P, D, ray_t, isect_object, isect_prim, verts);
+  }
   /* Compute face normal. */
   float3 Ng;
   if (sd->object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {

intern/cycles/kernel/geom/point.h

@@ -81,7 +81,7 @@ ccl_device float3 point_attribute_float3(KernelGlobals kg,
 #  endif
 
   if (desc.element == ATTR_ELEMENT_VERTEX) {
-    return kernel_tex_fetch(__attributes_float3, desc.offset + sd->prim);
+    return float4_to_float3(kernel_tex_fetch(__attributes_float4, desc.offset + sd->prim));
   }
   else {
     return make_float3(0.0f, 0.0f, 0.0f);
@@ -109,59 +109,17 @@ ccl_device float4 point_attribute_float4(KernelGlobals kg,
   }
 }
 
-/* Point position */
-
-ccl_device float3 point_position(KernelGlobals kg, ccl_private const ShaderData *sd)
-{
-  if (sd->type & PRIMITIVE_POINT) {
-    /* World space center. */
-    float3 P = (sd->type & PRIMITIVE_MOTION) ?
-                   float4_to_float3(motion_point(kg, sd->object, sd->prim, sd->time)) :
-                   float4_to_float3(kernel_tex_fetch(__points, sd->prim));
-
-    if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
-      object_position_transform(kg, sd, &P);
-    }
-
-    return P;
-  }
-
-  return zero_float3();
-}
-
 /* Point radius */
 
 ccl_device float point_radius(KernelGlobals kg, ccl_private const ShaderData *sd)
 {
   if (sd->type & PRIMITIVE_POINT) {
-    /* World space radius. */
+    return kernel_tex_fetch(__points, sd->prim).w;
-    const float r = kernel_tex_fetch(__points, sd->prim).w;
-
-    if (sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED) {
-      return r;
-    }
-    else {
-      float3 dir = make_float3(r, r, r);
-      object_dir_transform(kg, sd, &dir);
-      return average(dir);
-    }
   }
 
   return 0.0f;
 }
 
-/* Point random */
-
-ccl_device float point_random(KernelGlobals kg, ccl_private const ShaderData *sd)
-{
-  if (sd->type & PRIMITIVE_POINT) {
-    const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_POINT_RANDOM);
-    return (desc.offset != ATTR_STD_NOT_FOUND) ? point_attribute_float(kg, sd, desc, NULL, NULL) :
-                                                 0.0f;
-  }
-  return 0.0f;
-}
-
 /* Point location for motion pass, linear interpolation between keys and
  * ignoring radius because we do the same for the motion keys */