Merge branch 'master' into draw-viewport-data

This is not needed anymore with the new dependency system.

CMakeLists.txt

@@ -273,11 +273,13 @@ endif()
 
 if(UNIX AND NOT APPLE)
   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_GLEW "Use GLEW OpenGL wrapper library provided by the operating system" OFF)
+  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()
 
 
@@ -536,12 +538,14 @@ 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)
@@ -567,6 +571,10 @@ if(UNIX AND NOT APPLE)
     option(WITH_LINKER_LLD "Use ld.lld linker which is usually faster than ld.gold" OFF)
     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)
@@ -702,9 +710,12 @@ if(UNIX AND NOT APPLE)
 endif()
 
 # Installation process.
-option(POSTINSTALL_SCRIPT "Run given CMake script after installation process" OFF)
+set(POSTINSTALL_SCRIPT "" CACHE FILEPATH "Run given CMake script after installation process")
 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(...)
 
 
@@ -2068,3 +2079,8 @@ 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,6 +63,7 @@ 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,8 +25,13 @@ else()
 endif()
 
 if(WIN32)
-  set(BOOST_TOOLSET toolset=msvc-14.1)
-  set(BOOST_COMPILER_STRING -vc141)
+  if(MSVC_VERSION GREATER_EQUAL 1920) # 2019
+    set(BOOST_TOOLSET toolset=msvc-14.2)
+    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

@@ -0,0 +1,38 @@
+# ***** 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,3 +94,4 @@ download_source(POTRACE)
 download_source(HARU)
 download_source(ZSTD)
 download_source(FLEX)
+download_source(BROTLI)

build_files/build_environment/cmake/freetype.cmake

@@ -23,9 +23,12 @@ 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
-  -DCMAKE_DISABLE_FIND_PACKAGE_BrotliDec=TRUE)
+  -DPC_BROTLIDEC_INCLUDEDIR=${LIBDIR}/brotli/include
+  -DPC_BROTLIDEC_LIBDIR=${LIBDIR}/brotli/lib
+  )
 
 ExternalProject_Add(external_freetype
   URL file://${PACKAGE_DIR}/${FREETYPE_FILE}
@@ -36,6 +39,11 @@ 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/harvest.cmake

@@ -79,6 +79,8 @@ 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.10.2)
+set(FREETYPE_VERSION 2.11.0)
 set(FREETYPE_URI http://prdownloads.sourceforge.net/freetype/freetype-${FREETYPE_VERSION}.tar.gz)
-set(FREETYPE_HASH b1cb620e4c875cd4d1bfa04945400945)
+set(FREETYPE_HASH cf09172322f6b50cf8f568bf8fe14bde)
 set(FREETYPE_HASH_TYPE MD5)
 set(FREETYPE_FILE freetype-${FREETYPE_VERSION}.tar.gz)
 
@@ -500,3 +500,9 @@ 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="13.0"
+LLVM_VERSION_MEX="14.0"
 LLVM_VERSION_FOUND=""
 LLVM_FORCE_BUILD=false
 LLVM_FORCE_REBUILD=false

build_files/build_environment/windows/build_deps.cmd

@@ -1,64 +1,39 @@
 @echo off
 if NOT "%1" == "" (
-	if "%1" == "2013" (
-    echo "Building for VS2013"
-    set VSVER=12.0
-    set VSVER_SHORT=12
-    set BuildDir=VS12
-    goto par2
-  )
-	if "%1" == "2015" (
-    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
     goto par2
+  )
+	if "%1" == "2019" (
+    echo "Building for VS2019"
+    set VSVER=15.0
+    set VSVER_SHORT=15
+    set BuildDir=VS15
+    goto par2
   )
   
 )
 :usage
 
-Echo Usage build_deps 2013/2015/2017 x64/x86
+Echo Usage build_deps 2017/2019 x64
 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" == "2013" (
-			set CMAKE_BUILDER=Visual Studio 12 2013 Win64
-		)
-		if "%1" == "2015" (
-			set CMAKE_BUILDER=Visual Studio 14 2015 Win64
+		if "%1" == "2019" (
+			set CMAKE_BUILDER=Visual Studio 16 2019
+			set CMAKE_BUILD_ARCH=-A x64
 		)
 		if "%1" == "2017" (
 			set CMAKE_BUILDER=Visual Studio 15 2017 Win64
+            set CMAKE_BUILD_ARCH=
 		)
-		
     goto start
   )
 )
@@ -120,7 +95,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%" -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%" %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%/
 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
@@ -133,7 +108,7 @@ if "%NODEBUG%" == "1" goto exit
 cd %BUILD_DIR%
 mkdir %STAGING%\%BuildDir%%ARCH%D
 cd %Staging%\%BuildDir%%ARCH%D
-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%
+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%
 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

@@ -0,0 +1,83 @@
+# - 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/platform/platform_apple.cmake

@@ -166,7 +166,11 @@ 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)
@@ -480,8 +484,11 @@ 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>")
-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>")
+# llvm-ranlib doesn't support this flag. Xcode's libtool does.
+if(NOT ${CMAKE_RANLIB} MATCHES ".*llvm-ranlib$")
+  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")
@@ -508,3 +515,6 @@ 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_unix.cmake

@@ -48,6 +48,9 @@ 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}")
 
@@ -100,7 +103,22 @@ 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.
@@ -536,6 +554,21 @@ 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)
@@ -648,6 +681,9 @@ endif()
 # ----------------------------------------------------------------------------
 # Compilers
 
+# Only set the linker once.
+set(_IS_LINKER_DEFAULT ON)
+
 # GNU Compiler
 if(CMAKE_COMPILER_IS_GNUCC)
   # ffp-contract=off:
@@ -666,26 +702,89 @@ if(CMAKE_COMPILER_IS_GNUCC)
   string(PREPEND CMAKE_CXX_FLAGS_RELWITHDEBINFO "${GCC_EXTRA_FLAGS_RELEASE} ")
   unset(GCC_EXTRA_FLAGS_RELEASE)
 
-  if(WITH_LINKER_GOLD)
+  # NOTE(@campbellbarton): Eventually mold will be able to use `-fuse-ld=mold`,
+  # 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_C_FLAGS " -fuse-ld=gold")
-      string(APPEND CMAKE_CXX_FLAGS " -fuse-ld=gold")
+      string(APPEND CMAKE_EXE_LINKER_FLAGS    " -fuse-ld=gold")
+      string(APPEND CMAKE_SHARED_LINKER_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)
+  if(WITH_LINKER_LLD AND _IS_LINKER_DEFAULT)
     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_C_FLAGS " -fuse-ld=lld")
-      string(APPEND CMAKE_CXX_FLAGS " -fuse-ld=lld")
+      string(APPEND CMAKE_EXE_LINKER_FLAGS    " -fuse-ld=lld")
+      string(APPEND CMAKE_SHARED_LINKER_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()
@@ -695,6 +794,28 @@ 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
@@ -718,6 +839,8 @@ 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,7 +347,11 @@ set(FREETYPE_INCLUDE_DIRS
   ${LIBDIR}/freetype/include
   ${LIBDIR}/freetype/include/freetype2
 )
-set(FREETYPE_LIBRARY ${LIBDIR}/freetype/lib/freetype2ST.lib)
+set(FREETYPE_LIBRARIES
+  ${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,7 +3,32 @@ 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%"
@@ -11,7 +36,3 @@ 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,17 +9,11 @@ exit /b 1
 :detect_done
 echo found clang-format in %CF_PATH%
 
-if EXIST %PYTHON% (
-    set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
-    goto detect_python_done
+if NOT EXIST %PYTHON% (
+    echo python not found, required for this operation
+    exit /b 1
 )
 
-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,18 +1,8 @@
-if EXIST "%PYTHON%" (
-    goto detect_python_done
+if NOT EXIST %PYTHON% (
+    echo python not found, required for this operation
+    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,18 +1,8 @@
-if EXIST %PYTHON% (
-    goto detect_python_done
+if NOT EXIST %PYTHON% (
+    echo python not found, required for this operation
+    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,10 +1,7 @@
-if EXIST %PYTHON% (
-	goto detect_python_done
+if NOT EXIST %PYTHON% (
+    echo python not found, required for this operation
+    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/python_api/examples/bpy.types.Bone.convert_local_to_pose.py

@@ -8,27 +8,42 @@ def set_pose_matrices(obj, matrix_map):
     "Assign pose space matrices of all bones at once, ignoring constraints."
 
     def rec(pbone, parent_matrix):
-        matrix = matrix_map[pbone.name]
+        if pbone.name in matrix_map:
+            matrix = matrix_map[pbone.name]
 
-        ## Instead of:
-        # pbone.matrix = matrix
-        # bpy.context.view_layer.update()
+            ## Instead of:
+            # pbone.matrix = matrix
+            # bpy.context.view_layer.update()
 
-        # Compute and assign local matrix, using the new parent matrix
-        if pbone.parent:
-            pbone.matrix_basis = pbone.bone.convert_local_to_pose(
-                matrix,
-                pbone.bone.matrix_local,
-                parent_matrix=parent_matrix,
-                parent_matrix_local=pbone.parent.bone.matrix_local,
-                invert=True
-            )
+            # Compute and assign local matrix, using the new parent matrix
+            if pbone.parent:
+                pbone.matrix_basis = pbone.bone.convert_local_to_pose(
+                    matrix,
+                    pbone.bone.matrix_local,
+                    parent_matrix=parent_matrix,
+                    parent_matrix_local=pbone.parent.bone.matrix_local,
+                    invert=True
+                )
+            else:
+                pbone.matrix_basis = pbone.bone.convert_local_to_pose(
+                    matrix,
+                    pbone.bone.matrix_local,
+                    invert=True
+                )
         else:
-            pbone.matrix_basis = pbone.bone.convert_local_to_pose(
-                matrix,
-                pbone.bone.matrix_local,
-                invert=True
-            )
+            # Compute the updated pose matrix from local and new parent matrix
+            if pbone.parent:
+                matrix = pbone.bone.convert_local_to_pose(
+                    pbone.matrix_basis,
+                    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 its self can be re-accessed, the following example will crash.
+Only the reference to the data itself can be re-accessed, the following example will crash.
 
 .. code-block:: python
 

doc/python_api/sphinx_doc_gen.py

@@ -1762,6 +1762,7 @@ 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})
+			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)
 		elseif(WIN32)
 			set(ENV{VS100COMNTOOLS} $ENV{VS120COMNTOOLS})
-			add_custom_command(OUTPUT build COMMAND ${PYTHON_EXECUTABLE} setup.py build DEPENDS ${PYTHON_SRC} ${PYTHON_HDR})
+			add_custom_command(OUTPUT build COMMAND ${PYTHON_EXECUTABLE} setup.py build DEPENDS ${PYTHON_SRC} ${PYTHON_HDR} setup.py)
 		else()
-			add_custom_command(OUTPUT build COMMAND ${PYTHON_EXECUTABLE} setup.py build DEPENDS ${PYTHON_SRC} ${PYTHON_HDR})
+			add_custom_command(OUTPUT build COMMAND ${PYTHON_EXECUTABLE} setup.py build DEPENDS ${PYTHON_SRC} ${PYTHON_HDR} setup.py)
 		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
-	from distutils.core import Distribution
-	from distutils.command.build import build
+    import subprocess
+    from distutils.core import Distribution
+    from distutils.command.build import build
 
-	dist = Distribution()
-	cmd = build(dist)
-	cmd.finalize_options()
-	#print(cmd.build_platlib)
+    dist = Distribution()
+    cmd = build(dist)
+    cmd.finalize_options()
+    #print(cmd.build_platlib)
 
-	os.environ['PYTHONPATH'] = os.path.join(os.getcwd(), cmd.build_platlib)
-	os.environ['LD_LIBRARY_PATH'] = os.getcwd()
+    os.environ['PYTHONPATH'] = os.path.join(os.getcwd(), cmd.build_platlib)
+    os.environ['LD_LIBRARY_PATH'] = os.getcwd()
 
-	ret = subprocess.call(sys.argv[2:])
-	sys.exit(ret)
+    ret = subprocess.call(sys.argv[2:])
+    sys.exit(ret)
 
 
 # the following line is not working due to https://bugs.python.org/issue9023
@@ -43,7 +43,8 @@ audaspace = Extension(
                       library_dirs = ['.', 'Release', 'Debug'],
                       language = 'c++',
                       extra_compile_args = extra_args,
-					  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 [])]
+                      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 [])]
 )
 
 setup(
@@ -56,6 +57,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,6 +95,13 @@ 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;
 
@@ -296,13 +303,78 @@ 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_wave_format_extensible({}),
+	m_default_device_changed(false),
+	m_reference_count(1)
 {
 	// initialize COM if it hasn't happened yet
 	CoInitializeEx(nullptr, COINIT_MULTITHREADED);
@@ -327,6 +399,8 @@ WASAPIDevice::WASAPIDevice(DeviceSpecs specs, int buffersize) :
 
 	create();
 
+	m_imm_device_enumerator->RegisterEndpointNotificationCallback(this);
+
 	return;
 
 	error:
@@ -340,6 +414,8 @@ 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 : public ThreadedDevice
+class AUD_PLUGIN_API WASAPIDevice : IMMNotificationClient, public ThreadedDevice
 {
 private:
 	int m_buffersize;
@@ -48,6 +48,8 @@ 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);
 
@@ -58,6 +60,17 @@ 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;

intern/cycles/blender/addon/properties.py

@@ -667,6 +667,11 @@ 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",
@@ -1447,6 +1452,19 @@ 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_sample_all_lights(context):
+def use_multi_device(context):
     cscene = context.scene.cycles
-
-    return cscene.sample_all_lights_direct or cscene.sample_all_lights_indirect
+    if cscene.device != 'GPU':
+        return False
+    return context.preferences.addons[__package__].preferences.has_multi_device()
 
 
 def show_device_active(context):
@@ -667,6 +667,10 @@ 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
 
@@ -679,21 +683,33 @@ class CYCLES_RENDER_PT_performance_acceleration_structure(CyclesButtonsPanel, Pa
 
         col = layout.column()
 
-        use_embree = False
+        use_embree = _cycles.with_embree
+
         if use_cpu(context):
-            use_embree = _cycles.with_embree
-            if not use_embree:
+            col.prop(cscene, "debug_use_spatial_splits")
+            if 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_spatial_splits")
-        sub = col.column()
-        sub.active = not use_embree
-        sub.prop(cscene, "debug_use_hair_bvh")
-        sub = col.column()
-        sub.active = not cscene.debug_use_spatial_splits and not use_embree
-        sub.prop(cscene, "debug_bvh_time_steps")
+            col.prop(cscene, "debug_use_hair_bvh")
+
+            # CPU is used in addition to a GPU
+            if use_multi_device(context) and use_embree:
+                col.prop(cscene, "debug_use_compact_bvh")
 
 
 class CYCLES_RENDER_PT_performance_final_render(CyclesButtonsPanel, Panel):
@@ -1803,18 +1819,45 @@ 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"
@@ -2183,6 +2226,7 @@ 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/mesh.cpp

@@ -1071,7 +1071,15 @@ static void create_subd_mesh(Scene *scene,
 
   for (BL::MeshEdge &e : b_mesh.edges) {
     if (e.crease() != 0.0f) {
-      mesh->add_crease(e.vertices()[0], e.vertices()[1], e.crease());
+      mesh->add_edge_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());
+      }
     }
   }
 

intern/cycles/blender/object.cpp

@@ -529,6 +529,17 @@ 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,8 +51,6 @@ 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) {
@@ -109,7 +107,7 @@ void BlenderOutputDriver::write_render_tile(const Tile &tile)
 
   BL::RenderLayer b_rlay = *b_single_rlay;
 
-  vector<float> pixels(tile.size.x * tile.size.y * 4);
+  vector<float> pixels(static_cast<size_t>(tile.size.x) * tile.size.y * 4);
 
   /* Copy each pass. */
   for (BL::RenderPass &b_pass : b_rlay.passes) {

intern/cycles/blender/shader.cpp

@@ -689,6 +689,9 @@ 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>();
   }

intern/cycles/blender/sync.cpp

@@ -787,6 +787,7 @@ 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");
 

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 overallocating here (considering leaf size into account),
+  /* This is a bit over-allocating 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,6 +61,26 @@ 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.
@@ -75,12 +95,16 @@ 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) {
@@ -160,6 +184,10 @@ 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) {
@@ -225,6 +253,11 @@ 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;
@@ -236,12 +269,15 @@ 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:
-      /* Nothing to do here. */
+      if (kernel_embree_is_self_intersection(kg, hit, cray)) {
+        *args->valid = 0;
+        return;
+      }
       break;
   }
 }
@@ -257,6 +293,14 @@ 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)
@@ -355,10 +399,12 @@ 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) |
-                                    RTC_SCENE_FLAG_COMPACT | RTC_SCENE_FLAG_ROBUST;
+                                    (compact ? RTC_SCENE_FLAG_COMPACT : RTC_SCENE_FLAG_NONE) |
+                                    RTC_SCENE_FLAG_ROBUST;
   rtcSetSceneFlags(scene, scene_flags);
   build_quality = dynamic ? RTC_BUILD_QUALITY_LOW :
                             (params.use_spatial_split ? RTC_BUILD_QUALITY_HIGH :
@@ -503,6 +549,7 @@ 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);
@@ -765,6 +812,7 @@ 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,6 +97,9 @@ 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/device/metal/bvh.h

@@ -58,6 +58,11 @@ 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,6 +19,7 @@
 #  include "scene/hair.h"
 #  include "scene/mesh.h"
 #  include "scene/object.h"
+#  include "scene/pointcloud.h"
 
 #  include "util/progress.h"
 
@@ -475,6 +476,220 @@ 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,
@@ -491,6 +706,8 @@ 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,6 +115,8 @@ 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,17 +87,14 @@ 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;
     }
@@ -432,6 +429,25 @@ 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();
@@ -561,11 +577,7 @@ void MetalDevice::generic_free(device_memory &mem)
       mmem.mtlBuffer = nil;
     }
 
-    stats.mem_free(mem.device_size);
-    mem.device_pointer = 0;
-    mem.device_size = 0;
-
-    metal_mem_map.erase(&mem);
+    erase_allocation(mem);
   }
 }
 
@@ -954,10 +966,7 @@ void MetalDevice::tex_free(device_texture &mem)
       delayed_free_list.push_back(mmem.mtlTexture);
       mmem.mtlTexture = nil;
     }
-    stats.mem_free(mem.device_size);
-    mem.device_pointer = 0;
-    mem.device_size = 0;
-    metal_mem_map.erase(&mem);
+    erase_allocation(mem);
   }
 }
 

intern/cycles/device/metal/kernel.h

@@ -36,6 +36,8 @@ 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,6 +358,8 @@ 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);
 
@@ -400,36 +402,50 @@ bool MetalDeviceKernels::load(MetalDevice *device, int kernel_type)
     NSArray *function_list = nil;
 
     if (device->use_metalrt) {
-      id<MTLFunction> box_intersect_default = nil;
-      id<MTLFunction> box_intersect_shadow = nil;
+      id<MTLFunction> curve_intersect_default = nil;
+      id<MTLFunction> curve_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. */
-          box_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL];
-          box_intersect_shadow = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL_SHADOW];
+          curve_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL];
+          curve_intersect_shadow =
+              rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL_SHADOW];
         }
         else {
-          box_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_RIBBON];
-          box_intersect_shadow =
+          curve_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_RIBBON];
+          curve_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],
-                           box_intersect_default ?
-                               box_intersect_default :
+                           curve_intersect_default ?
+                               curve_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],
-                           box_intersect_shadow ?
-                               box_intersect_shadow :
+                           curve_intersect_shadow ?
+                               curve_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 = 6;
+  pipeline_options.numPayloadValues = 8;
   pipeline_options.numAttributeValues = 2; /* u, v */
   pipeline_options.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE;
   pipeline_options.pipelineLaunchParamsVariableName = "__params"; /* See globals.h */

intern/cycles/integrator/pass_accessor.cpp

@@ -141,6 +141,7 @@ 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. */
@@ -188,8 +189,10 @@ 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
@@ -215,7 +218,7 @@ bool PassAccessor::get_render_tile_pixels(const RenderBuffers *render_buffers,
     }
   }
 
-  pad_pixels(buffer_params, destination, pass_info.num_components);
+  pad_pixels(buffer_params, destination, num_written_components);
 
   return true;
 }

intern/cycles/integrator/path_trace.cpp

@@ -820,8 +820,15 @@ 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/kernel/bvh/bvh.h

@@ -173,15 +173,16 @@ 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_NONE;
+  uint ray_flags = OPTIX_RAY_FLAG_ENFORCE_ANYHIT;
   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,
@@ -200,7 +201,9 @@ 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);
@@ -242,6 +245,7 @@ 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;
@@ -309,6 +313,7 @@ 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 &&
@@ -356,6 +361,9 @@ 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;
 
@@ -379,7 +387,9 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
              p2,
              p3,
              p4,
-             p5);
+             p5,
+             p6,
+             p7);
 
   return p5;
 #  elif defined(__METALRT__)
@@ -417,6 +427,7 @@ 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;
@@ -460,6 +471,7 @@ 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;
@@ -532,6 +544,8 @@ 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) {
@@ -555,7 +569,9 @@ 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);
@@ -588,6 +604,7 @@ 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;
@@ -634,6 +651,7 @@ 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);
@@ -685,6 +703,8 @@ 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) {
@@ -708,7 +728,9 @@ 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);
@@ -744,6 +766,7 @@ 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;
@@ -820,6 +843,7 @@ 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,6 +22,8 @@
 #include "kernel/device/cpu/compat.h"
 #include "kernel/device/cpu/globals.h"
 
+#include "kernel/bvh/util.h"
+
 #include "util/vector.h"
 
 CCL_NAMESPACE_BEGIN
@@ -38,6 +40,9 @@ struct CCLIntersectContext {
   KernelGlobals kg;
   RayType type;
 
+  /* For avoiding self intersections */
+  const Ray *ray;
+
   /* for shadow rays */
   Intersection *isect_s;
   uint max_hits;
@@ -56,6 +61,7 @@ struct CCLIntersectContext {
   {
     kg = kg_;
     type = type_;
+    ray = NULL;
     max_hits = 1;
     num_hits = 0;
     num_recorded_hits = 0;
@@ -102,7 +108,34 @@ 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.primID = RTC_INVALID_GEOMETRY_ID;
+  rayhit.hit.instID[0] = 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,12 +148,27 @@ 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,
@@ -168,13 +183,28 @@ 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,6 +15,7 @@
  */
 
 struct MetalRTIntersectionPayload {
+  RaySelfPrimitives self;
   uint visibility;
   float u, v;
   int prim;
@@ -25,6 +26,7 @@ struct MetalRTIntersectionPayload {
 };
 
 struct MetalRTIntersectionLocalPayload {
+  RaySelfPrimitives self;
   uint local_object;
   uint lcg_state;
   short max_hits;
@@ -34,6 +36,7 @@ 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 */
-          while (prim_addr < prim_addr2) {
+          for (; prim_addr < prim_addr2; prim_addr++) {
             kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) ==
                           (type & PRIMITIVE_ALL));
             bool hit;
@@ -156,16 +156,32 @@ 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, object, prim_addr);
+                    kg, &isect, P, dir, t_max_current, visibility, prim_object, prim, prim_addr);
                 break;
               }
 #if BVH_FEATURE(BVH_MOTION)
               case PRIMITIVE_MOTION_TRIANGLE: {
-                hit = motion_triangle_intersect(
-                    kg, &isect, P, dir, t_max_current, ray->time, visibility, object, prim_addr);
+                hit = motion_triangle_intersect(kg,
+                                                &isect,
+                                                P,
+                                                dir,
+                                                t_max_current,
+                                                ray->time,
+                                                visibility,
+                                                prim_object,
+                                                prim,
+                                                prim_addr);
                 break;
               }
 #endif
@@ -182,20 +198,9 @@ 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);
-                const int curve_prim = kernel_tex_fetch(__prim_index, prim_addr);
-                hit = curve_intersect(kg,
-                                      &isect,
-                                      P,
-                                      dir,
-                                      t_max_current,
-                                      curve_object,
-                                      curve_prim,
-                                      ray->time,
-                                      curve_type);
+                hit = curve_intersect(
+                    kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, curve_type);
 
                 break;
               }
@@ -211,20 +216,9 @@ 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(kg,
-                                      &isect,
-                                      P,
-                                      dir,
-                                      t_max_current,
-                                      point_object,
-                                      point_prim,
-                                      ray->time,
-                                      point_type);
+                hit = point_intersect(
+                    kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, point_type);
                 break;
               }
 #endif /* BVH_FEATURE(BVH_POINTCLOUD) */
@@ -301,8 +295,6 @@ ccl_device_inline
                 integrator_state_write_shadow_isect(state, &isect, record_index);
               }
             }
-
-            prim_addr++;
           }
         }
         else {

intern/cycles/kernel/bvh/traversal.h

@@ -133,89 +133,91 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
           --stack_ptr;
 
           /* primitive intersection */
-          switch (type & PRIMITIVE_ALL) {
-            case PRIMITIVE_TRIANGLE: {
-              for (; prim_addr < prim_addr2; prim_addr++) {
-                kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
+          for (; prim_addr < prim_addr2; prim_addr++) {
+            kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
+
+            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: {
                 if (triangle_intersect(
-                        kg, isect, P, dir, isect->t, visibility, object, prim_addr)) {
+                        kg, isect, P, dir, isect->t, visibility, prim_object, prim, 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: {
-              for (; prim_addr < prim_addr2; prim_addr++) {
-                kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
-                if (motion_triangle_intersect(
-                        kg, isect, P, dir, isect->t, ray->time, visibility, object, prim_addr)) {
+              case PRIMITIVE_MOTION_TRIANGLE: {
+                if (motion_triangle_intersect(kg,
+                                              isect,
+                                              P,
+                                              dir,
+                                              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_MOTION_CURVE_THICK:
-            case PRIMITIVE_CURVE_RIBBON:
-            case PRIMITIVE_MOTION_CURVE_RIBBON: {
-              for (; prim_addr < prim_addr2; prim_addr++) {
+              case PRIMITIVE_CURVE_THICK:
+              case PRIMITIVE_MOTION_CURVE_THICK:
+              case PRIMITIVE_CURVE_RIBBON:
+              case PRIMITIVE_MOTION_CURVE_RIBBON: {
                 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) {
-                    continue;
+                    break;
                   }
                 }
 
-                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, curve_object, curve_prim, ray->time, curve_type);
+                    kg, isect, P, dir, isect->t, prim_object, 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_MOTION_POINT: {
-              for (; prim_addr < prim_addr2; prim_addr++) {
+              case PRIMITIVE_POINT:
+              case PRIMITIVE_MOTION_POINT: {
                 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) {
-                    continue;
+                    break;
                   }
                 }
 
-                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, point_object, point_prim, ray->time, point_type);
+                    kg, isect, P, dir, isect->t, prim_object, prim, ray->time, point_type);
                 if (hit) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
                     return true;
                 }
+                break;
               }
-              break;
-            }
 #endif /* BVH_FEATURE(BVH_POINTCLOUD) */
+            }
           }
         }
         else {

intern/cycles/kernel/bvh/util.h

@@ -21,54 +21,22 @@ 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. */
-
+ * rays leaving from a surface. This is from "A Fast and Robust Method for Avoiding
+ * Self-Intersection" see https://research.nvidia.com/publication/2019-03_A-Fast-and
+ */
 ccl_device_inline float3 ray_offset(float3 P, float3 Ng)
 {
-#ifdef __INTERSECTION_REFINE__
-  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;
+  const float int_scale = 256.0f;
+  int3 of_i = make_int3((int)(int_scale * Ng.x), (int)(int_scale * Ng.y), (int)(int_scale * Ng.z));
 
-  float3 res;
-
-  /* x component */
-  if (fabsf(P.x) < epsilon_test) {
-    res.x = P.x + Ng.x * epsilon_f;
-  }
-  else {
-    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
+  float3 p_i = make_float3(__int_as_float(__float_as_int(P.x) + ((P.x < 0) ? -of_i.x : of_i.x)),
+                           __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)));
+  const float origin = 1.0f / 32.0f;
+  const float float_scale = 1.0f / 65536.0f;
+  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,
+                     fabsf(P.z) < origin ? P.z + float_scale * Ng.z : p_i.z);
 }
 
 #if defined(__KERNEL_CPU__)
@@ -227,4 +195,25 @@ 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,14 +140,20 @@ 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 */
-                uint tri_object = (object == OBJECT_NONE) ?
-                                      kernel_tex_fetch(__prim_object, prim_addr) :
-                                      object;
-                int object_flag = kernel_tex_fetch(__object_flag, tri_object);
+                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(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(kg, isect, P, dir, isect->t, visibility, object, prim_addr);
+                triangle_intersect(
+                    kg, isect, P, dir, isect->t, visibility, prim_object, prim, prim_addr);
               }
               break;
             }
@@ -157,15 +163,27 @@ 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 */
-                uint tri_object = (object == OBJECT_NONE) ?
-                                      kernel_tex_fetch(__prim_object, prim_addr) :
-                                      object;
-                int object_flag = kernel_tex_fetch(__object_flag, tri_object);
+                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(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, isect, P, dir, isect->t, ray->time, visibility, object, prim_addr);
+                motion_triangle_intersect(kg,
+                                          isect,
+                                          P,
+                                          dir,
+                                          isect->t,
+                                          ray->time,
+                                          visibility,
+                                          prim_object,
+                                          prim,
+                                          prim_addr);
               }
               break;
             }

intern/cycles/kernel/bvh/volume_all.h

@@ -143,15 +143,19 @@ 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 */
-                uint tri_object = (object == OBJECT_NONE) ?
-                                      kernel_tex_fetch(__prim_object, prim_addr) :
-                                      object;
-                int object_flag = kernel_tex_fetch(__object_flag, tri_object);
+                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(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, object, prim_addr);
+                    kg, isect_array, P, dir, isect_t, visibility, prim_object, prim, prim_addr);
                 if (hit) {
                   /* Move on to next entry in intersections array. */
                   isect_array++;
@@ -183,15 +187,27 @@ 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 */
-                uint tri_object = (object == OBJECT_NONE) ?
-                                      kernel_tex_fetch(__prim_object, prim_addr) :
-                                      object;
-                int object_flag = kernel_tex_fetch(__object_flag, tri_object);
+                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(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, isect_array, P, dir, isect_t, ray->time, visibility, object, prim_addr);
+                hit = motion_triangle_intersect(kg,
+                                                isect_array,
+                                                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,22 +199,18 @@ 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) {
-    weights_pdf = weights / sum_weights;
+    *pdf = weights / sum_weights;
   }
   else {
-    weights_pdf = make_float3(1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f);
+    *pdf = make_float3(1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f);
   }
 
-  *pdf = weights_pdf;
-
-  /* OpenCL does not support -> on float3, so don't use pdf->x. */
-  if (rand < weights_pdf.x) {
+  if (rand < pdf->x) {
     return 0;
   }
-  else if (rand < weights_pdf.x + weights_pdf.y) {
+  else if (rand < pdf->x + pdf->y) {
     return 1;
   }
   else {

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

@@ -40,6 +40,27 @@ 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,
@@ -53,8 +74,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) {
-    /* Only intersect with matching object */
+  if ((object != payload.local_object) || intersection_skip_self_local(payload.self, prim)) {
+    /* Only intersect with matching object and skip self-intersecton. */
     result.accept = false;
     result.continue_search = true;
     return result;
@@ -166,6 +187,11 @@ 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) {
@@ -322,21 +348,35 @@ inline TReturnType metalrt_visibility_test(constant KernelParamsMetal &launch_pa
   }
 #  endif
 
-#  ifdef __VISIBILITY_FLAG__
   uint visibility = payload.visibility;
+#  ifdef __VISIBILITY_FLAG__
   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) {
-    result.accept = true;
-    result.continue_search = false;
-    return result;
+    if (intersection_skip_self_shadow(payload.self, object, prim)) {
+      result.accept = false;
+      result.continue_search = true;
+      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;
@@ -576,6 +616,150 @@ __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,6 +45,11 @@ 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__
@@ -111,6 +116,12 @@ 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 */
@@ -149,8 +160,6 @@ 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;
@@ -185,6 +194,11 @@ 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()) {
@@ -314,6 +328,12 @@ 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()
@@ -330,18 +350,31 @@ 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();
   }
-
-  /* Shadow ray early termination. */
-  if (visibility & PATH_RAY_SHADOW_OPAQUE) {
-    return optixTerminateRay();
-  }
 #endif
+
+  const int prim = optixGetPrimitiveIndex();
+  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
+
+  if (visibility & PATH_RAY_SHADOW_OPAQUE) {
+    if (intersection_skip_self_shadow(ray->self, object, prim)) {
+      return optixIgnoreIntersection();
+    }
+    else {
+      /* Shadow ray early termination. */
+      return optixTerminateRay();
+    }
+  }
+  else {
+    if (intersection_skip_self(ray->self, object, prim)) {
+      return optixIgnoreIntersection();
+    }
+  }
 }
 
 extern "C" __global__ void __closesthit__kernel_optix_hit()

intern/cycles/kernel/film/read.h

@@ -214,6 +214,21 @@ 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,6 +226,18 @@ 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_intersect.h

@@ -29,46 +29,19 @@
 
 CCL_NAMESPACE_BEGIN
 
-/* Refine triangle intersection to more precise hit point. For rays that travel
- * far the precision is often not so good, this reintersects the primitive from
- * a closer distance.
+/**
+ * Use the barycentric coordinates to get the intersection location
  */
-
-ccl_device_inline float3 motion_triangle_refine(KernelGlobals kg,
-                                                ccl_private ShaderData *sd,
-                                                float3 P,
-                                                float3 D,
-                                                float t,
-                                                const int isect_object,
-                                                const int isect_prim,
-                                                float3 verts[3])
+ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg,
+                                                       ccl_private ShaderData *sd,
+                                                       const int isect_object,
+                                                       const int isect_prim,
+                                                       const float u,
+                                                       const float v,
+                                                       float3 verts[3])
 {
-#ifdef __INTERSECTION_REFINE__
-  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;
+  float w = 1.0f - u - v;
+  float3 P = u * verts[0] + v * verts[1] + w * verts[2];
 
   if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
@@ -76,71 +49,8 @@ ccl_device_inline float3 motion_triangle_refine(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.
  */
@@ -153,14 +63,12 @@ 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, fobject, prim, time, verts);
+  motion_triangle_vertices(kg, 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)) {
@@ -175,8 +83,7 @@ ccl_device_inline bool motion_triangle_intersect(KernelGlobals kg,
       isect->u = u;
       isect->v = v;
       isect->prim = prim;
-      isect->object = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) :
-                                                object;
+      isect->object = object;
       isect->type = PRIMITIVE_MOTION_TRIANGLE;
       return true;
     }
@@ -196,25 +103,15 @@ ccl_device_inline bool motion_triangle_intersect_local(KernelGlobals kg,
                                                        float3 dir,
                                                        float time,
                                                        int object,
-                                                       int local_object,
+                                                       int prim,
                                                        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, local_object, prim, time, verts);
+  motion_triangle_vertices(kg, 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)) {
@@ -266,7 +163,7 @@ ccl_device_inline bool motion_triangle_intersect_local(KernelGlobals kg,
   isect->u = u;
   isect->v = v;
   isect->prim = prim;
-  isect->object = local_object;
+  isect->object = object;
   isect->type = PRIMITIVE_MOTION_TRIANGLE;
 
   /* Record geometric normal. */

intern/cycles/kernel/geom/motion_triangle_shader.h

@@ -68,15 +68,7 @@ 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. */
-#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);
-  }
+  sd->P = motion_triangle_point_from_uv(kg, sd, isect_object, isect_prim, sd->u, sd->v, 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 float4_to_float3(kernel_tex_fetch(__attributes_float4, desc.offset + sd->prim));
+    return kernel_tex_fetch(__attributes_float3, desc.offset + sd->prim);
   }
   else {
     return make_float3(0.0f, 0.0f, 0.0f);
@@ -109,17 +109,59 @@ 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) {
-    return kernel_tex_fetch(__points, sd->prim).w;
+    /* World space radius. */
+    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 */
 

intern/cycles/kernel/geom/shader_data.h

@@ -89,7 +89,7 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
       sd->shader = kernel_tex_fetch(__tri_shader, sd->prim);
 
       /* vectors */
-      sd->P = triangle_refine(kg, sd, ray->P, ray->D, isect->t, isect->object, isect->prim);
+      sd->P = triangle_point_from_uv(kg, sd, isect->object, isect->prim, isect->u, isect->v);
       sd->Ng = Ng;
       sd->N = Ng;
 

intern/cycles/kernel/geom/triangle_intersect.h

@@ -33,9 +33,9 @@ ccl_device_inline bool triangle_intersect(KernelGlobals kg,
                                           float tmax,
                                           uint visibility,
                                           int object,
+                                          int prim,
                                           int prim_addr)
 {
-  const int prim = kernel_tex_fetch(__prim_index, prim_addr);
   const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
   const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
                tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
@@ -49,8 +49,7 @@ ccl_device_inline bool triangle_intersect(KernelGlobals kg,
     if (kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
 #endif
     {
-      isect->object = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) :
-                                                object;
+      isect->object = object;
       isect->prim = prim;
       isect->type = PRIMITIVE_TRIANGLE;
       isect->u = u;
@@ -74,21 +73,12 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg,
                                                 float3 P,
                                                 float3 dir,
                                                 int object,
-                                                int local_object,
+                                                int prim,
                                                 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;
-    }
-  }
-
-  const int prim = kernel_tex_fetch(__prim_index, prim_addr);
   const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
   const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
                tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
@@ -139,7 +129,7 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg,
   /* Record intersection. */
   ccl_private Intersection *isect = &local_isect->hits[hit];
   isect->prim = prim;
-  isect->object = local_object;
+  isect->object = object;
   isect->type = PRIMITIVE_TRIANGLE;
   isect->u = u;
   isect->v = v;
@@ -152,58 +142,23 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg,
 }
 #endif /* __BVH_LOCAL__ */
 
-/* Refine triangle intersection to more precise hit point. For rays that travel
- * far the precision is often not so good, this reintersects the primitive from
- * a closer distance. */
-
-/* Reintersections uses the paper:
- *
- * Tomas Moeller
- * Fast, minimum storage ray/triangle intersection
- * http://www.cs.virginia.edu/~gfx/Courses/2003/ImageSynthesis/papers/Acceleration/Fast%20MinimumStorage%20RayTriangle%20Intersection.pdf
+/**
+ * Use the barycentric coordinates to get the intersection location
  */
-
-ccl_device_inline float3 triangle_refine(KernelGlobals kg,
-                                         ccl_private ShaderData *sd,
-                                         float3 P,
-                                         float3 D,
-                                         float t,
-                                         const int isect_object,
-                                         const int isect_prim)
+ccl_device_inline float3 triangle_point_from_uv(KernelGlobals kg,
+                                                ccl_private ShaderData *sd,
+                                                const int isect_object,
+                                                const int isect_prim,
+                                                const float u,
+                                                const float v)
 {
-#ifdef __INTERSECTION_REFINE__
-  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;
-
   const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
   const packed_float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
                       tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
                       tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
-  float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
-  float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
-  float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
-  float3 qvec = cross(tvec, edge1);
-  float3 pvec = cross(D, edge2);
-  float det = dot(edge1, pvec);
-  if (det != 0.0f) {
-    /* If determinant is zero it means ray lies in the plane of
-     * the triangle. It is possible in theory due to watertight
-     * nature of triangle intersection. For such cases we simply
-     * don't refine intersection hoping it'll go all fine.
-     */
-    float rt = dot(edge2, qvec) / det;
-    P = P + D * rt;
-  }
+  float w = 1.0f - u - v;
+
+  float3 P = u * tri_a + v * tri_b + w * tri_c;
 
   if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
@@ -211,65 +166,6 @@ ccl_device_inline float3 triangle_refine(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.
- */
-ccl_device_inline float3 triangle_refine_local(KernelGlobals kg,
-                                               ccl_private ShaderData *sd,
-                                               float3 P,
-                                               float3 D,
-                                               float t,
-                                               const int isect_object,
-                                               const int isect_prim)
-{
-#if defined(__KERNEL_GPU_RAYTRACING__)
-  /* t is always in world space with OptiX and MetalRT. */
-  return triangle_refine(kg, sd, P, D, t, isect_object, isect_prim);
-#else
-  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;
-
-#  ifdef __INTERSECTION_REFINE__
-  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
-  const packed_float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
-                      tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
-                      tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
-  float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
-  float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
-  float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
-  float3 qvec = cross(tvec, edge1);
-  float3 pvec = cross(D, edge2);
-  float det = dot(edge1, pvec);
-  if (det != 0.0f) {
-    /* If determinant is zero it means ray lies in the plane of
-     * the triangle. It is possible in theory due to watertight
-     * nature of triangle intersection. For such cases we simply
-     * don't refine intersection hoping it'll go all fine.
-     */
-    float rt = dot(edge2, qvec) / det;
-    P = P + D * rt;
-  }
-#  endif /* __INTERSECTION_REFINE__ */
-
-  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
-    const Transform tfm = object_get_transform(kg, sd);
-    P = transform_point(&tfm, P);
-  }
-
-  return P;
-#endif
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/integrator/intersect_closest.h

@@ -328,6 +328,12 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg,
 
   /* Scene Intersection. */
   Intersection isect ccl_optional_struct_init;
+  isect.object = OBJECT_NONE;
+  isect.prim = PRIM_NONE;
+  ray.self.object = last_isect_object;
+  ray.self.prim = last_isect_prim;
+  ray.self.light_object = OBJECT_NONE;
+  ray.self.light_prim = PRIM_NONE;
   bool hit = scene_intersect(kg, &ray, visibility, &isect);
 
   /* TODO: remove this and do it in the various intersection functions instead. */

intern/cycles/kernel/integrator/intersect_shadow.h

@@ -156,7 +156,10 @@ ccl_device void integrator_intersect_shadow(KernelGlobals kg, IntegratorShadowSt
   /* Read ray from integrator state into local memory. */
   Ray ray ccl_optional_struct_init;
   integrator_state_read_shadow_ray(kg, state, &ray);
-
+  ray.self.object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, object);
+  ray.self.prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, prim);
+  ray.self.light_object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 1, object);
+  ray.self.light_prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 1, prim);
   /* Compute visibility. */
   const uint visibility = integrate_intersect_shadow_visibility(kg, state);
 

intern/cycles/kernel/integrator/intersect_volume_stack.h

@@ -38,7 +38,10 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
   Ray volume_ray ccl_optional_struct_init;
   volume_ray.P = from_P;
   volume_ray.D = normalize_len(to_P - from_P, &volume_ray.t);
-
+  volume_ray.self.object = INTEGRATOR_STATE(state, isect, object);
+  volume_ray.self.prim = INTEGRATOR_STATE(state, isect, prim);
+  volume_ray.self.light_object = OBJECT_NONE;
+  volume_ray.self.light_prim = PRIM_NONE;
   /* Store to avoid global fetches on every intersection step. */
   const uint volume_stack_size = kernel_data.volume_stack_size;
 
@@ -68,7 +71,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
     volume_stack_enter_exit(kg, state, stack_sd);
 
     /* Move ray forward. */
-    volume_ray.P = ray_offset(stack_sd->P, -stack_sd->Ng);
+    volume_ray.P = stack_sd->P;
     if (volume_ray.t != FLT_MAX) {
       volume_ray.D = normalize_len(to_P - volume_ray.P, &volume_ray.t);
     }
@@ -91,6 +94,10 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
    * fewest hits. */
   volume_ray.D = make_float3(0.0f, 0.0f, 1.0f);
   volume_ray.t = FLT_MAX;
+  volume_ray.self.object = OBJECT_NONE;
+  volume_ray.self.prim = PRIM_NONE;
+  volume_ray.self.light_object = OBJECT_NONE;
+  volume_ray.self.light_prim = PRIM_NONE;
 
   int stack_index = 0, enclosed_index = 0;
 
@@ -203,7 +210,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
     }
 
     /* Move ray forward. */
-    volume_ray.P = ray_offset(stack_sd->P, -stack_sd->Ng);
+    volume_ray.P = stack_sd->P;
     ++step;
   }
 #endif

intern/cycles/kernel/integrator/shade_light.h

@@ -37,8 +37,9 @@ ccl_device_inline void integrate_light(KernelGlobals kg,
 
   /* Advance ray beyond light. */
   /* TODO: can we make this more numerically robust to avoid reintersecting the
-   * same light in some cases? */
-  const float3 new_ray_P = ray_offset(ray_P + ray_D * isect.t, ray_D);
+   * same light in some cases? Ray should not intersect surface anymore as the
+   * object and prim ids will prevent self intersection. */
+  const float3 new_ray_P = ray_P + ray_D * isect.t;
   INTEGRATOR_STATE_WRITE(state, ray, P) = new_ray_P;
   INTEGRATOR_STATE_WRITE(state, ray, t) -= isect.t;
 
@@ -46,7 +47,7 @@ ccl_device_inline void integrate_light(KernelGlobals kg,
   const float mis_ray_t = INTEGRATOR_STATE(state, path, mis_ray_t);
   ray_P -= ray_D * mis_ray_t;
   isect.t += mis_ray_t;
-  INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = mis_ray_t + isect.t;
+  INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = isect.t;
 
   LightSample ls ccl_optional_struct_init;
   const bool use_light_sample = light_sample_from_intersection(kg, &isect, ray_P, ray_D, &ls);

intern/cycles/kernel/integrator/shade_shadow.h

@@ -83,7 +83,10 @@ ccl_device_inline void integrate_transparent_volume_shadow(KernelGlobals kg,
   /* Setup shader data. */
   Ray ray ccl_optional_struct_init;
   integrator_state_read_shadow_ray(kg, state, &ray);
-
+  ray.self.object = OBJECT_NONE;
+  ray.self.prim = PRIM_NONE;
+  ray.self.light_object = OBJECT_NONE;
+  ray.self.light_prim = PRIM_NONE;
   /* Modify ray position and length to match current segment. */
   const float start_t = (hit == 0) ? 0.0f :
                                      INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit - 1, t);
@@ -149,7 +152,7 @@ ccl_device_inline bool integrate_transparent_shadow(KernelGlobals kg,
     const float last_hit_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, num_recorded_hits - 1, t);
     const float3 ray_P = INTEGRATOR_STATE(state, shadow_ray, P);
     const float3 ray_D = INTEGRATOR_STATE(state, shadow_ray, D);
-    INTEGRATOR_STATE_WRITE(state, shadow_ray, P) = ray_offset(ray_P + last_hit_t * ray_D, ray_D);
+    INTEGRATOR_STATE_WRITE(state, shadow_ray, P) = ray_P + last_hit_t * ray_D;
     INTEGRATOR_STATE_WRITE(state, shadow_ray, t) -= last_hit_t;
   }
 

intern/cycles/kernel/integrator/shade_surface.h

@@ -182,13 +182,18 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
 
   /* Write shadow ray and associated state to global memory. */
   integrator_state_write_shadow_ray(kg, shadow_state, &ray);
+  // Save memory by storing the light and object indices in the shadow_isect
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
 
   /* Copy state from main path to shadow path. */
   const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);
   const uint16_t transparent_bounce = INTEGRATOR_STATE(state, path, transparent_bounce);
   uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
   shadow_flag |= (is_light) ? PATH_RAY_SHADOW_FOR_LIGHT : 0;
-  shadow_flag |= PATH_RAY_SURFACE_PASS;
+  shadow_flag |= (shadow_flag & PATH_RAY_ANY_PASS) ? 0 : PATH_RAY_SURFACE_PASS;
   const float3 throughput = INTEGRATOR_STATE(state, path, throughput) * bsdf_eval_sum(&bsdf_eval);
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
@@ -266,13 +271,11 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
   }
 
   /* Setup ray. Note that clipping works through transparent bounces. */
-  INTEGRATOR_STATE_WRITE(state, ray, P) = ray_offset(sd->P,
-                                                     (label & LABEL_TRANSMIT) ? -sd->Ng : sd->Ng);
+  INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
   INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(bsdf_omega_in);
   INTEGRATOR_STATE_WRITE(state, ray, t) = (label & LABEL_TRANSPARENT) ?
                                               INTEGRATOR_STATE(state, ray, t) - sd->ray_length :
                                               FLT_MAX;
-
 #ifdef __RAY_DIFFERENTIALS__
   INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
   INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(bsdf_domega_in);
@@ -316,7 +319,7 @@ ccl_device_forceinline bool integrate_surface_volume_only_bounce(IntegratorState
   }
 
   /* Setup ray position, direction stays unchanged. */
-  INTEGRATOR_STATE_WRITE(state, ray, P) = ray_offset(sd->P, -sd->Ng);
+  INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
 
   /* Clipping works through transparent. */
   INTEGRATOR_STATE_WRITE(state, ray, t) -= sd->ray_length;
@@ -360,10 +363,14 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
   }
 
   Ray ray ccl_optional_struct_init;
-  ray.P = ray_offset(sd->P, sd->Ng);
+  ray.P = sd->P;
   ray.D = ao_D;
   ray.t = kernel_data.integrator.ao_bounces_distance;
   ray.time = sd->time;
+  ray.self.object = sd->object;
+  ray.self.prim = sd->prim;
+  ray.self.light_object = OBJECT_NONE;
+  ray.self.light_prim = PRIM_NONE;
   ray.dP = differential_zero_compact();
   ray.dD = differential_zero_compact();
 
@@ -375,6 +382,10 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
 
   /* Write shadow ray and associated state to global memory. */
   integrator_state_write_shadow_ray(kg, shadow_state, &ray);
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
 
   /* Copy state from main path to shadow path. */
   const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);

intern/cycles/kernel/integrator/shade_volume.h

@@ -791,13 +791,17 @@ ccl_device_forceinline void integrate_volume_direct_light(
 
   /* Write shadow ray and associated state to global memory. */
   integrator_state_write_shadow_ray(kg, shadow_state, &ray);
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
+  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
 
   /* Copy state from main path to shadow path. */
   const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);
   const uint16_t transparent_bounce = INTEGRATOR_STATE(state, path, transparent_bounce);
   uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
   shadow_flag |= (is_light) ? PATH_RAY_SHADOW_FOR_LIGHT : 0;
-  shadow_flag |= PATH_RAY_VOLUME_PASS;
+  shadow_flag |= (shadow_flag & PATH_RAY_ANY_PASS) ? 0 : PATH_RAY_VOLUME_PASS;
   const float3 throughput_phase = throughput * bsdf_eval_sum(&phase_eval);
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
@@ -873,11 +877,13 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
   INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
   INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(phase_omega_in);
   INTEGRATOR_STATE_WRITE(state, ray, t) = FLT_MAX;
-
 #  ifdef __RAY_DIFFERENTIALS__
   INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
   INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(phase_domega_in);
 #  endif
+  // Save memory by storing last hit prim and object in isect
+  INTEGRATOR_STATE_WRITE(state, isect, prim) = sd->prim;
+  INTEGRATOR_STATE_WRITE(state, isect, object) = sd->object;
 
   /* Update throughput. */
   const float3 throughput = INTEGRATOR_STATE(state, path, throughput);

intern/cycles/kernel/integrator/shadow_state_template.h

@@ -61,6 +61,7 @@ KERNEL_STRUCT_MEMBER(shadow_ray, packed_float3, D, KERNEL_FEATURE_PATH_TRACING)
 KERNEL_STRUCT_MEMBER(shadow_ray, float, t, KERNEL_FEATURE_PATH_TRACING)
 KERNEL_STRUCT_MEMBER(shadow_ray, float, time, KERNEL_FEATURE_PATH_TRACING)
 KERNEL_STRUCT_MEMBER(shadow_ray, float, dP, KERNEL_FEATURE_PATH_TRACING)
+KERNEL_STRUCT_MEMBER(shadow_ray, int, object, KERNEL_FEATURE_PATH_TRACING)
 KERNEL_STRUCT_END(shadow_ray)
 
 /*********************** Shadow Intersection result **************************/

intern/cycles/kernel/integrator/subsurface.h

@@ -57,7 +57,6 @@ ccl_device int subsurface_bounce(KernelGlobals kg,
 
   /* Pass along object info, reusing isect to save memory. */
   INTEGRATOR_STATE_WRITE(state, subsurface, Ng) = sd->Ng;
-  INTEGRATOR_STATE_WRITE(state, isect, object) = sd->object;
 
   uint32_t path_flag = (INTEGRATOR_STATE(state, path, flag) & ~PATH_RAY_CAMERA) |
                        ((sc->type == CLOSURE_BSSRDF_BURLEY_ID) ? PATH_RAY_SUBSURFACE_DISK :
@@ -165,10 +164,8 @@ ccl_device_inline bool subsurface_scatter(KernelGlobals kg, IntegratorState stat
 
     if (object_flag & SD_OBJECT_INTERSECTS_VOLUME) {
       float3 P = INTEGRATOR_STATE(state, ray, P);
-      const float3 Ng = INTEGRATOR_STATE(state, subsurface, Ng);
-      const float3 offset_P = ray_offset(P, -Ng);
 
-      integrator_volume_stack_update_for_subsurface(kg, state, offset_P, ray.P);
+      integrator_volume_stack_update_for_subsurface(kg, state, P, ray.P);
     }
   }
 #  endif /* __VOLUME__ */

intern/cycles/kernel/integrator/subsurface_disk.h

@@ -99,6 +99,10 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg,
   ray.dP = ray_dP;
   ray.dD = differential_zero_compact();
   ray.time = time;
+  ray.self.object = OBJECT_NONE;
+  ray.self.prim = PRIM_NONE;
+  ray.self.light_object = OBJECT_NONE;
+  ray.self.light_prim = OBJECT_NONE;
 
   /* Intersect with the same object. if multiple intersections are found it
    * will use at most BSSRDF_MAX_HITS hits, a random subset of all hits. */

intern/cycles/kernel/integrator/subsurface_random_walk.h

@@ -195,6 +195,7 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
   const float time = INTEGRATOR_STATE(state, ray, time);
   const float3 Ng = INTEGRATOR_STATE(state, subsurface, Ng);
   const int object = INTEGRATOR_STATE(state, isect, object);
+  const int prim = INTEGRATOR_STATE(state, isect, prim);
 
   /* Sample diffuse surface scatter into the object. */
   float3 D;
@@ -205,12 +206,16 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
   }
 
   /* Setup ray. */
-  ray.P = ray_offset(P, -Ng);
+  ray.P = P;
   ray.D = D;
   ray.t = FLT_MAX;
   ray.time = time;
   ray.dP = ray_dP;
   ray.dD = differential_zero_compact();
+  ray.self.object = object;
+  ray.self.prim = prim;
+  ray.self.light_object = OBJECT_NONE;
+  ray.self.light_prim = PRIM_NONE;
 
 #ifndef __KERNEL_GPU_RAYTRACING__
   /* Compute or fetch object transforms. */
@@ -377,7 +382,15 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
      * If yes, we will later use backwards guided sampling in order to have a decent
      * chance of connecting to it.
      * TODO: Maybe use less than 10 times the mean free path? */
-    ray.t = (bounce == 0) ? max(t, 10.0f / (min3(sigma_t))) : t;
+    if (bounce == 0) {
+      ray.t = max(t, 10.0f / (min3(sigma_t)));
+    }
+    else {
+      ray.t = t;
+      /* After the first bounce the object can intersect the same surface again */
+      ray.self.object = OBJECT_NONE;
+      ray.self.prim = PRIM_NONE;
+    }
     scene_intersect_local(kg, &ray, &ss_isect, object, NULL, 1);
     hit = (ss_isect.num_hits > 0);
 
@@ -408,13 +421,6 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
     if (hit) {
       t = ray.t;
     }
-    else if (bounce == 0) {
-      /* Restore original position if nothing was hit after the first bounce,
-       * without the ray_offset() that was added to avoid self-intersection.
-       * Otherwise if that offset is relatively large compared to the scattering
-       * radius, we never go back up high enough to exit the surface. */
-      ray.P = P;
-    }
 
     /* Advance to new scatter location. */
     ray.P += t * ray.D;

intern/cycles/kernel/light/light.h

@@ -418,8 +418,8 @@ ccl_device bool light_sample_from_intersection(KernelGlobals kg,
   LightType type = (LightType)klight->type;
   ls->type = type;
   ls->shader = klight->shader_id;
-  ls->object = PRIM_NONE;
-  ls->prim = PRIM_NONE;
+  ls->object = isect->object;
+  ls->prim = isect->prim;
   ls->lamp = lamp;
   /* todo: missing texture coordinates */
   ls->t = isect->t;

intern/cycles/kernel/light/sample.h

@@ -198,7 +198,7 @@ ccl_device_inline float3 shadow_ray_offset(KernelGlobals kg,
   float NL = dot(sd->N, L);
   bool transmit = (NL < 0.0f);
   float3 Ng = (transmit ? -sd->Ng : sd->Ng);
-  float3 P = ray_offset(sd->P, Ng);
+  float3 P = sd->P;
 
   if ((sd->type & PRIMITIVE_TRIANGLE) && (sd->shader & SHADER_SMOOTH_NORMAL)) {
     const float offset_cutoff =
@@ -243,7 +243,7 @@ ccl_device_inline void shadow_ray_setup(ccl_private const ShaderData *ccl_restri
     }
     else {
       /* other lights, avoid self-intersection */
-      ray->D = ray_offset(ls->P, ls->Ng) - P;
+      ray->D = ls->P - P;
       ray->D = normalize_len(ray->D, &ray->t);
     }
   }
@@ -257,6 +257,12 @@ ccl_device_inline void shadow_ray_setup(ccl_private const ShaderData *ccl_restri
   ray->dP = differential_make_compact(sd->dP);
   ray->dD = differential_zero_compact();
   ray->time = sd->time;
+
+  /* Fill in intersection surface and light details. */
+  ray->self.prim = sd->prim;
+  ray->self.object = sd->object;
+  ray->self.light_prim = ls->prim;
+  ray->self.light_object = ls->object;
 }
 
 /* Create shadow ray towards light sample. */

intern/cycles/kernel/osl/services.cpp

@@ -116,6 +116,8 @@ ustring OSLRenderServices::u_curve_tangent_normal("geom:curve_tangent_normal");
 ustring OSLRenderServices::u_curve_random("geom:curve_random");
 ustring OSLRenderServices::u_is_point("geom:is_point");
 ustring OSLRenderServices::u_point_radius("geom:point_radius");
+ustring OSLRenderServices::u_point_position("geom:point_position");
+ustring OSLRenderServices::u_point_random("geom:point_random");
 ustring OSLRenderServices::u_normal_map_normal("geom:normal_map_normal");
 ustring OSLRenderServices::u_path_ray_length("path:ray_length");
 ustring OSLRenderServices::u_path_ray_depth("path:ray_depth");
@@ -999,6 +1001,10 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg
     float3 f = curve_tangent_normal(kg, sd);
     return set_attribute_float3(f, type, derivatives, val);
   }
+  else if (name == u_curve_random) {
+    float f = curve_random(kg, sd);
+    return set_attribute_float(f, type, derivatives, val);
+  }
   /* point attributes */
   else if (name == u_is_point) {
     float f = (sd->type & PRIMITIVE_POINT) != 0;
@@ -1008,6 +1014,14 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg
     float f = point_radius(kg, sd);
     return set_attribute_float(f, type, derivatives, val);
   }
+  else if (name == u_point_position) {
+    float3 f = point_position(kg, sd);
+    return set_attribute_float3(f, type, derivatives, val);
+  }
+  else if (name == u_point_random) {
+    float f = point_random(kg, sd);
+    return set_attribute_float(f, type, derivatives, val);
+  }
   else if (name == u_normal_map_normal) {
     if (sd->type & PRIMITIVE_TRIANGLE) {
       float3 f = triangle_smooth_normal_unnormalized(kg, sd, sd->Ng, sd->prim, sd->u, sd->v);

intern/cycles/kernel/osl/services.h

@@ -298,7 +298,9 @@ class OSLRenderServices : public OSL::RendererServices {
   static ustring u_curve_tangent_normal;
   static ustring u_curve_random;
   static ustring u_is_point;
+  static ustring u_point_position;
   static ustring u_point_radius;
+  static ustring u_point_random;
   static ustring u_normal_map_normal;
   static ustring u_path_ray_length;
   static ustring u_path_ray_depth;

intern/cycles/kernel/osl/shaders/CMakeLists.txt

@@ -49,6 +49,7 @@ set(SRC_OSL
   node_glossy_bsdf.osl
   node_gradient_texture.osl
   node_hair_info.osl
+  node_point_info.osl
   node_scatter_volume.osl
   node_absorption_volume.osl
   node_principled_volume.osl

intern/cycles/kernel/osl/shaders/node_point_info.osl

@@ -0,0 +1,26 @@
+/*
+ * Copyright 2011-2022 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "stdcycles.h"
+
+shader node_point_info(output point Position = point(0.0, 0.0, 0.0),
+                       output float Radius = 0.0,
+                       output float Random = 0.0)
+{
+  getattribute("geom:point_position", Position);
+  getattribute("geom:point_radius", Radius);
+  getattribute("geom:point_random", Random);
+}

intern/cycles/kernel/svm/ao.h

@@ -70,10 +70,14 @@ ccl_device float svm_ao(
 
     /* Create ray. */
     Ray ray;
-    ray.P = ray_offset(sd->P, N);
+    ray.P = sd->P;
     ray.D = D.x * T + D.y * B + D.z * N;
     ray.t = max_dist;
     ray.time = sd->time;
+    ray.self.object = sd->object;
+    ray.self.prim = sd->prim;
+    ray.self.light_object = OBJECT_NONE;
+    ray.self.light_prim = PRIM_NONE;
     ray.dP = differential_zero_compact();
     ray.dD = differential_zero_compact();
 

intern/cycles/kernel/svm/attribute.h

@@ -87,7 +87,9 @@ ccl_device_noinline void svm_node_attr(KernelGlobals kg,
   if (node.y == ATTR_STD_GENERATED && desc.element == ATTR_ELEMENT_NONE) {
     /* No generated attribute, fall back to object coordinates. */
     float3 f = sd->P;
-    object_inverse_position_transform(kg, sd, &f);
+    if (sd->object != OBJECT_NONE) {
+      object_inverse_position_transform(kg, sd, &f);
+    }
     if (type == NODE_ATTR_OUTPUT_FLOAT) {
       stack_store_float(stack, out_offset, average(f));
     }
@@ -179,7 +181,9 @@ ccl_device_noinline void svm_node_attr_bump_dx(KernelGlobals kg,
   if (node.y == ATTR_STD_GENERATED && desc.element == ATTR_ELEMENT_NONE) {
     /* No generated attribute, fall back to object coordinates. */
     float3 f = sd->P + sd->dP.dx;
-    object_inverse_position_transform(kg, sd, &f);
+    if (sd->object != OBJECT_NONE) {
+      object_inverse_position_transform(kg, sd, &f);
+    }
     if (type == NODE_ATTR_OUTPUT_FLOAT) {
       stack_store_float(stack, out_offset, average(f));
     }
@@ -275,7 +279,9 @@ ccl_device_noinline void svm_node_attr_bump_dy(KernelGlobals kg,
   if (node.y == ATTR_STD_GENERATED && desc.element == ATTR_ELEMENT_NONE) {
     /* No generated attribute, fall back to object coordinates. */
     float3 f = sd->P + sd->dP.dy;
-    object_inverse_position_transform(kg, sd, &f);
+    if (sd->object != OBJECT_NONE) {
+      object_inverse_position_transform(kg, sd, &f);
+    }
     if (type == NODE_ATTR_OUTPUT_FLOAT) {
       stack_store_float(stack, out_offset, average(f));
     }

intern/cycles/kernel/svm/bevel.h

@@ -196,6 +196,10 @@ ccl_device float3 svm_bevel(
     ray.dP = differential_zero_compact();
     ray.dD = differential_zero_compact();
     ray.time = sd->time;
+    ray.self.object = OBJECT_NONE;
+    ray.self.prim = PRIM_NONE;
+    ray.self.light_object = OBJECT_NONE;
+    ray.self.light_prim = PRIM_NONE;
 
     /* Intersect with the same object. if multiple intersections are found it
      * will use at most LOCAL_MAX_HITS hits, a random subset of all hits. */
@@ -207,15 +211,24 @@ ccl_device float3 svm_bevel(
       /* Quickly retrieve P and Ng without setting up ShaderData. */
       float3 hit_P;
       if (sd->type == PRIMITIVE_TRIANGLE) {
-        hit_P = triangle_refine_local(
-            kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim);
+        hit_P = triangle_point_from_uv(kg,
+                                       sd,
+                                       isect.hits[hit].object,
+                                       isect.hits[hit].prim,
+                                       isect.hits[hit].u,
+                                       isect.hits[hit].v);
       }
 #  ifdef __OBJECT_MOTION__
       else if (sd->type == PRIMITIVE_MOTION_TRIANGLE) {
         float3 verts[3];
         motion_triangle_vertices(kg, sd->object, isect.hits[hit].prim, sd->time, verts);
-        hit_P = motion_triangle_refine_local(
-            kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim, verts);
+        hit_P = motion_triangle_point_from_uv(kg,
+                                              sd,
+                                              isect.hits[hit].object,
+                                              isect.hits[hit].prim,
+                                              isect.hits[hit].u,
+                                              isect.hits[hit].v,
+                                              verts);
       }
 #  endif /* __OBJECT_MOTION__ */
 

intern/cycles/kernel/svm/geometry.h

@@ -242,13 +242,6 @@ ccl_device_noinline void svm_node_hair_info(KernelGlobals kg,
       stack_store_float(stack, out_offset, data);
       break;
     }
-#  if 0
-    case NODE_INFO_CURVE_FADE: {
-      data = sd->curve_transparency;
-      stack_store_float(stack, out_offset, data);
-      break;
-    }
-#  endif
     case NODE_INFO_CURVE_TANGENT_NORMAL: {
       data3 = curve_tangent_normal(kg, sd);
       stack_store_float3(stack, out_offset, data3);
@@ -258,4 +251,28 @@ ccl_device_noinline void svm_node_hair_info(KernelGlobals kg,
 }
 #endif
 
+#ifdef __POINTCLOUD__
+
+/* Point Info */
+
+ccl_device_noinline void svm_node_point_info(KernelGlobals kg,
+                                             ccl_private ShaderData *sd,
+                                             ccl_private float *stack,
+                                             uint type,
+                                             uint out_offset)
+{
+  switch (type) {
+    case NODE_INFO_POINT_POSITION:
+      stack_store_float3(stack, out_offset, point_position(kg, sd));
+      break;
+    case NODE_INFO_POINT_RADIUS:
+      stack_store_float(stack, out_offset, point_radius(kg, sd));
+      break;
+    case NODE_INFO_POINT_RANDOM:
+      break; /* handled as attribute */
+  }
+}
+
+#endif
+
 CCL_NAMESPACE_END

intern/cycles/kernel/svm/svm.h

@@ -454,13 +454,14 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
         break;
 #if defined(__HAIR__)
       case NODE_HAIR_INFO:
-        IF_KERNEL_NODES_FEATURE(HAIR)
-        {
-          svm_node_hair_info(kg, sd, stack, node.y, node.z);
-        }
+        svm_node_hair_info(kg, sd, stack, node.y, node.z);
+        break;
+#endif
+#if defined(__POINTCLOUD__)
+      case NODE_POINT_INFO:
+        svm_node_point_info(kg, sd, stack, node.y, node.z);
         break;
 #endif
-
       case NODE_TEXTURE_MAPPING:
         offset = svm_node_texture_mapping(kg, sd, stack, node.y, node.z, offset);
         break;

intern/cycles/kernel/svm/types.h

@@ -81,6 +81,7 @@ typedef enum ShaderNodeType {
   NODE_OBJECT_INFO,
   NODE_PARTICLE_INFO,
   NODE_HAIR_INFO,
+  NODE_POINT_INFO,
   NODE_TEXTURE_MAPPING,
   NODE_MAPPING,
   NODE_MIN_MAX,
@@ -176,12 +177,16 @@ typedef enum NodeHairInfo {
   NODE_INFO_CURVE_INTERCEPT,
   NODE_INFO_CURVE_LENGTH,
   NODE_INFO_CURVE_THICKNESS,
-  /* Fade for minimum hair width transiency. */
-  // NODE_INFO_CURVE_FADE,
   NODE_INFO_CURVE_TANGENT_NORMAL,
   NODE_INFO_CURVE_RANDOM,
 } NodeHairInfo;
 
+typedef enum NodePointInfo {
+  NODE_INFO_POINT_POSITION,
+  NODE_INFO_POINT_RADIUS,
+  NODE_INFO_POINT_RANDOM,
+} NodePointInfo;
+
 typedef enum NodeLightPath {
   NODE_LP_camera = 0,
   NODE_LP_shadow,

intern/cycles/kernel/types.h

@@ -512,12 +512,21 @@ typedef struct differential {
 
 /* Ray */
 
+typedef struct RaySelfPrimitives {
+  int prim;         /* Primitive the ray is starting from */
+  int object;       /* Instance prim is a part of */
+  int light_prim;   /* Light primitive */
+  int light_object; /* Light object */
+} RaySelfPrimitives;
+
 typedef struct Ray {
   float3 P;   /* origin */
   float3 D;   /* direction */
   float t;    /* length of the ray */
   float time; /* time (for motion blur) */
 
+  RaySelfPrimitives self;
+
 #ifdef __RAY_DIFFERENTIALS__
   float dP;
   float dD;
@@ -1565,21 +1574,21 @@ enum KernelFeatureFlag : uint32_t {
   KERNEL_FEATURE_NODE_BSDF = (1U << 0U),
   KERNEL_FEATURE_NODE_EMISSION = (1U << 1U),
   KERNEL_FEATURE_NODE_VOLUME = (1U << 2U),
-  KERNEL_FEATURE_NODE_HAIR = (1U << 3U),
-  KERNEL_FEATURE_NODE_BUMP = (1U << 4U),
-  KERNEL_FEATURE_NODE_BUMP_STATE = (1U << 5U),
-  KERNEL_FEATURE_NODE_VORONOI_EXTRA = (1U << 6U),
-  KERNEL_FEATURE_NODE_RAYTRACE = (1U << 7U),
-  KERNEL_FEATURE_NODE_AOV = (1U << 8U),
-  KERNEL_FEATURE_NODE_LIGHT_PATH = (1U << 9U),
+  KERNEL_FEATURE_NODE_BUMP = (1U << 3U),
+  KERNEL_FEATURE_NODE_BUMP_STATE = (1U << 4U),
+  KERNEL_FEATURE_NODE_VORONOI_EXTRA = (1U << 5U),
+  KERNEL_FEATURE_NODE_RAYTRACE = (1U << 6U),
+  KERNEL_FEATURE_NODE_AOV = (1U << 7U),
+  KERNEL_FEATURE_NODE_LIGHT_PATH = (1U << 8U),
 
   /* Use denoising kernels and output denoising passes. */
-  KERNEL_FEATURE_DENOISING = (1U << 10U),
+  KERNEL_FEATURE_DENOISING = (1U << 9U),
 
   /* Use path tracing kernels. */
-  KERNEL_FEATURE_PATH_TRACING = (1U << 11U),
+  KERNEL_FEATURE_PATH_TRACING = (1U << 10U),
 
   /* BVH/sampling kernel features. */
+  KERNEL_FEATURE_POINTCLOUD = (1U << 11U),
   KERNEL_FEATURE_HAIR = (1U << 12U),
   KERNEL_FEATURE_HAIR_THICK = (1U << 13U),
   KERNEL_FEATURE_OBJECT_MOTION = (1U << 14U),
@@ -1616,9 +1625,6 @@ enum KernelFeatureFlag : uint32_t {
   KERNEL_FEATURE_AO_PASS = (1U << 25U),
   KERNEL_FEATURE_AO_ADDITIVE = (1U << 26U),
   KERNEL_FEATURE_AO = (KERNEL_FEATURE_AO_PASS | KERNEL_FEATURE_AO_ADDITIVE),
-
-  /* Point clouds. */
-  KERNEL_FEATURE_POINTCLOUD = (1U << 27U),
 };
 
 /* Shader node feature mask, to specialize shader evaluation for kernels. */
@@ -1628,7 +1634,7 @@ enum KernelFeatureFlag : uint32_t {
    KERNEL_FEATURE_NODE_LIGHT_PATH)
 #define KERNEL_FEATURE_NODE_MASK_SURFACE_SHADOW \
   (KERNEL_FEATURE_NODE_BSDF | KERNEL_FEATURE_NODE_EMISSION | KERNEL_FEATURE_NODE_VOLUME | \
-   KERNEL_FEATURE_NODE_HAIR | KERNEL_FEATURE_NODE_BUMP | KERNEL_FEATURE_NODE_BUMP_STATE | \
+   KERNEL_FEATURE_NODE_BUMP | KERNEL_FEATURE_NODE_BUMP_STATE | \
    KERNEL_FEATURE_NODE_VORONOI_EXTRA | KERNEL_FEATURE_NODE_LIGHT_PATH)
 #define KERNEL_FEATURE_NODE_MASK_SURFACE \
   (KERNEL_FEATURE_NODE_MASK_SURFACE_SHADOW | KERNEL_FEATURE_NODE_RAYTRACE | \

intern/cycles/scene/alembic.cpp

@@ -742,6 +742,7 @@ NODE_DEFINE(AlembicProcedural)
   NodeType *type = NodeType::add("alembic", create);
 
   SOCKET_STRING(filepath, "Filename", ustring());
+  SOCKET_STRING_ARRAY(layers, "Layers", array<ustring>());
   SOCKET_FLOAT(frame, "Frame", 1.0f);
   SOCKET_FLOAT(start_frame, "Start Frame", 1.0f);
   SOCKET_FLOAT(end_frame, "End Frame", 1.0f);
@@ -839,14 +840,26 @@ void AlembicProcedural::generate(Scene *scene, Progress &progress)
     return;
   }
 
-  if (!archive.valid()) {
+  if (!archive.valid() || filepath_is_modified() || layers_is_modified()) {
     Alembic::AbcCoreFactory::IFactory factory;
     factory.setPolicy(Alembic::Abc::ErrorHandler::kQuietNoopPolicy);
-    archive = factory.getArchive(filepath.c_str());
+
+    std::vector<std::string> filenames;
+    filenames.push_back(filepath.c_str());
+
+    for (const ustring &layer : layers) {
+      filenames.push_back(layer.c_str());
+    }
+
+    /* We need to reverse the order as overriding archives should come first. */
+    std::reverse(filenames.begin(), filenames.end());
+
+    archive = factory.getArchive(filenames);
 
     if (!archive.valid()) {
       /* avoid potential infinite update loops in viewport synchronization */
       filepath.clear();
+      layers.clear();
       clear_modified();
       return;
     }
@@ -1165,6 +1178,12 @@ void AlembicProcedural::read_subd(AlembicObject *abc_object, Abc::chrono_t frame
   cached_data.subd_creases_weight.copy_to_socket(
       frame_time, mesh, mesh->get_subd_creases_weight_socket());
 
+  cached_data.subd_vertex_crease_indices.copy_to_socket(
+      frame_time, mesh, mesh->get_subd_vert_creases_socket());
+
+  cached_data.subd_vertex_crease_weights.copy_to_socket(
+      frame_time, mesh, mesh->get_subd_vert_creases_weight_socket());
+
   mesh->set_num_subd_faces(mesh->get_subd_shader().size());
 
   /* Update attributes. */

intern/cycles/scene/alembic.h

@@ -320,6 +320,8 @@ struct CachedData {
   DataStore<int> num_ngons;
   DataStore<array<int>> subd_creases_edge;
   DataStore<array<float>> subd_creases_weight;
+  DataStore<array<int>> subd_vertex_crease_indices;
+  DataStore<array<float>> subd_vertex_crease_weights;
 
   /* hair data */
   DataStore<array<float3>> curve_keys;
@@ -479,6 +481,10 @@ class AlembicProcedural : public Procedural {
   /* The file path to the Alembic archive */
   NODE_SOCKET_API(ustring, filepath)
 
+  /* Layers for the Alembic archive. Layers are in the order in which they override data, with the
+   * latter elements overriding the former ones. */
+  NODE_SOCKET_API_ARRAY(array<ustring>, layers)
+
   /* The current frame to render. */
   NODE_SOCKET_API(float, frame)
 

intern/cycles/scene/alembic_read.cpp

@@ -478,7 +478,9 @@ static void add_subd_polygons(CachedData &cached_data, const SubDSchemaData &dat
   cached_data.uv_loops.add_data(uv_loops, time);
 }
 
-static void add_subd_creases(CachedData &cached_data, const SubDSchemaData &data, chrono_t time)
+static void add_subd_edge_creases(CachedData &cached_data,
+                                  const SubDSchemaData &data,
+                                  chrono_t time)
 {
   if (!(data.crease_indices.valid() && data.crease_indices.valid() &&
         data.crease_sharpnesses.valid())) {
@@ -517,6 +519,37 @@ static void add_subd_creases(CachedData &cached_data, const SubDSchemaData &data
   }
 }
 
+static void add_subd_vertex_creases(CachedData &cached_data,
+                                    const SubDSchemaData &data,
+                                    chrono_t time)
+{
+  if (!(data.corner_indices.valid() && data.crease_sharpnesses.valid())) {
+    return;
+  }
+
+  const ISampleSelector iss = ISampleSelector(time);
+  const Int32ArraySamplePtr creases_indices = data.crease_indices.getValue(iss);
+  const FloatArraySamplePtr creases_sharpnesses = data.crease_sharpnesses.getValue(iss);
+
+  if (!(creases_indices && creases_sharpnesses) ||
+      creases_indices->size() != creases_sharpnesses->size()) {
+    return;
+  }
+
+  array<float> sharpnesses;
+  sharpnesses.reserve(creases_indices->size());
+  array<int> indices;
+  indices.reserve(creases_indices->size());
+
+  for (size_t i = 0; i < creases_indices->size(); i++) {
+    indices.push_back_reserved((*creases_indices)[i]);
+    sharpnesses.push_back_reserved((*creases_sharpnesses)[i]);
+  }
+
+  cached_data.subd_vertex_crease_indices.add_data(indices, time);
+  cached_data.subd_vertex_crease_weights.add_data(sharpnesses, time);
+}
+
 static void read_subd_geometry(CachedData &cached_data, const SubDSchemaData &data, chrono_t time)
 {
   const ISampleSelector iss = ISampleSelector(time);
@@ -525,7 +558,8 @@ static void read_subd_geometry(CachedData &cached_data, const SubDSchemaData &da
 
   if (data.topology_variance != kHomogenousTopology || cached_data.shader.size() == 0) {
     add_subd_polygons(cached_data, data, time);
-    add_subd_creases(cached_data, data, time);
+    add_subd_edge_creases(cached_data, data, time);
+    add_subd_vertex_creases(cached_data, data, time);
   }
 }
 

intern/cycles/scene/alembic_read.h

@@ -76,9 +76,10 @@ struct SubDSchemaData {
 
   vector<FaceSetShaderIndexPair> shader_face_sets;
 
-  // Those are unsupported for now.
   Alembic::AbcGeom::IInt32ArrayProperty corner_indices;
   Alembic::AbcGeom::IFloatArrayProperty corner_sharpnesses;
+
+  // Those are unsupported for now.
   Alembic::AbcGeom::IInt32Property face_varying_interpolate_boundary;
   Alembic::AbcGeom::IInt32Property face_varying_propagate_corners;
   Alembic::AbcGeom::IInt32Property interpolate_boundary;

intern/cycles/scene/colorspace.cpp

@@ -263,7 +263,9 @@ template<typename T> inline void cast_from_float4(T *data, float4 value)
 
 /* Slower versions for other all data types, which needs to convert to float and back. */
 template<typename T, bool compress_as_srgb = false>
-inline void processor_apply_pixels(const OCIO::Processor *processor, T *pixels, size_t num_pixels)
+inline void processor_apply_pixels_rgba(const OCIO::Processor *processor,
+                                        T *pixels,
+                                        size_t num_pixels)
 {
   /* TODO: implement faster version for when we know the conversion
    * is a simple matrix transform between linear spaces. In that case
@@ -310,25 +312,79 @@ inline void processor_apply_pixels(const OCIO::Processor *processor, T *pixels,
     }
   }
 }
+
+template<typename T, bool compress_as_srgb = false>
+inline void processor_apply_pixels_grayscale(const OCIO::Processor *processor,
+                                             T *pixels,
+                                             size_t num_pixels)
+{
+  OCIO::ConstCPUProcessorRcPtr device_processor = processor->getDefaultCPUProcessor();
+
+  /* Process large images in chunks to keep temporary memory requirement down. */
+  const size_t chunk_size = std::min((size_t)(16 * 1024 * 1024), num_pixels);
+  vector<float> float_pixels(chunk_size * 3);
+
+  for (size_t j = 0; j < num_pixels; j += chunk_size) {
+    size_t width = std::min(chunk_size, num_pixels - j);
+
+    /* Convert to 3 channels, since that's the minimum required by OpenColorIO. */
+    {
+      const T *pixel = pixels + j;
+      float *fpixel = float_pixels.data();
+      for (size_t i = 0; i < width; i++, pixel++, fpixel += 3) {
+        const float f = util_image_cast_to_float<T>(*pixel);
+        fpixel[0] = f;
+        fpixel[1] = f;
+        fpixel[2] = f;
+      }
+    }
+
+    OCIO::PackedImageDesc desc((float *)float_pixels.data(), width, 1, 3);
+    device_processor->apply(desc);
+
+    {
+      T *pixel = pixels + j;
+      const float *fpixel = float_pixels.data();
+      for (size_t i = 0; i < width; i++, pixel++, fpixel += 3) {
+        float f = average(make_float3(fpixel[0], fpixel[1], fpixel[2]));
+        if (compress_as_srgb) {
+          f = color_linear_to_srgb(f);
+        }
+        *pixel = util_image_cast_from_float<T>(f);
+      }
+    }
+  }
+}
+
 #endif
 
 template<typename T>
-void ColorSpaceManager::to_scene_linear(ustring colorspace,
-                                        T *pixels,
-                                        size_t num_pixels,
-                                        bool compress_as_srgb)
+void ColorSpaceManager::to_scene_linear(
+    ustring colorspace, T *pixels, size_t num_pixels, bool is_rgba, bool compress_as_srgb)
 {
 #ifdef WITH_OCIO
   const OCIO::Processor *processor = (const OCIO::Processor *)get_processor(colorspace);
 
   if (processor) {
-    if (compress_as_srgb) {
-      /* Compress output as sRGB. */
-      processor_apply_pixels<T, true>(processor, pixels, num_pixels);
+    if (is_rgba) {
+      if (compress_as_srgb) {
+        /* Compress output as sRGB. */
+        processor_apply_pixels_rgba<T, true>(processor, pixels, num_pixels);
+      }
+      else {
+        /* Write output as scene linear directly. */
+        processor_apply_pixels_rgba<T>(processor, pixels, num_pixels);
+      }
     }
     else {
-      /* Write output as scene linear directly. */
-      processor_apply_pixels<T>(processor, pixels, num_pixels);
+      if (compress_as_srgb) {
+        /* Compress output as sRGB. */
+        processor_apply_pixels_grayscale<T, true>(processor, pixels, num_pixels);
+      }
+      else {
+        /* Write output as scene linear directly. */
+        processor_apply_pixels_grayscale<T>(processor, pixels, num_pixels);
+      }
     }
   }
 #else
@@ -348,6 +404,11 @@ void ColorSpaceManager::to_scene_linear(ColorSpaceProcessor *processor_,
 
   if (processor) {
     OCIO::ConstCPUProcessorRcPtr device_processor = processor->getDefaultCPUProcessor();
+    if (channels == 1) {
+      float3 rgb = make_float3(pixel[0], pixel[0], pixel[0]);
+      device_processor->applyRGB(&rgb.x);
+      pixel[0] = average(rgb);
+    }
     if (channels == 3) {
       device_processor->applyRGB(pixel);
     }
@@ -390,9 +451,9 @@ void ColorSpaceManager::free_memory()
 }
 
 /* Template instantiations so we don't have to inline functions. */
-template void ColorSpaceManager::to_scene_linear(ustring, uchar *, size_t, bool);
-template void ColorSpaceManager::to_scene_linear(ustring, ushort *, size_t, bool);
-template void ColorSpaceManager::to_scene_linear(ustring, half *, size_t, bool);
-template void ColorSpaceManager::to_scene_linear(ustring, float *, size_t, bool);
+template void ColorSpaceManager::to_scene_linear(ustring, uchar *, size_t, bool, bool);
+template void ColorSpaceManager::to_scene_linear(ustring, ushort *, size_t, bool, bool);
+template void ColorSpaceManager::to_scene_linear(ustring, half *, size_t, bool, bool);
+template void ColorSpaceManager::to_scene_linear(ustring, float *, size_t, bool, bool);
 
 CCL_NAMESPACE_END

intern/cycles/scene/colorspace.h

@@ -43,10 +43,8 @@ class ColorSpaceManager {
   /* Convert pixels in the specified colorspace to scene linear color for
    * rendering. Must be a colorspace returned from detect_known_colorspace. */
   template<typename T>
-  static void to_scene_linear(ustring colorspace,
-                              T *pixels,
-                              size_t num_pixels,
-                              bool compress_as_srgb);
+  static void to_scene_linear(
+      ustring colorspace, T *pixels, size_t num_pixels, bool is_rgba, bool compress_as_srgb);
 
   /* Efficiently convert pixels to scene linear colorspace at render time,
    * for OSL where the image texture cache contains original pixels. The

intern/cycles/scene/constant_fold.cpp

@@ -441,9 +441,13 @@ void ConstantFolder::fold_mapping(NodeMappingType type) const
   if (is_zero(scale_in)) {
     make_zero();
   }
-  else if ((is_zero(location_in) || type == NODE_MAPPING_TYPE_VECTOR ||
-            type == NODE_MAPPING_TYPE_NORMAL) &&
-           is_zero(rotation_in) && is_one(scale_in)) {
+  else if (
+      /* Can't constant fold since we always need to normalize the output. */
+      (type != NODE_MAPPING_TYPE_NORMAL) &&
+      /* Check all use values are zero, note location is not used by vector and normal types. */
+      (is_zero(location_in) || type == NODE_MAPPING_TYPE_VECTOR ||
+       type == NODE_MAPPING_TYPE_NORMAL) &&
+      is_zero(rotation_in) && is_one(scale_in)) {
     try_bypass_or_make_constant(vector_in);
   }
 }

intern/cycles/scene/geometry.cpp

@@ -236,6 +236,7 @@ void Geometry::compute_bvh(
 
       BVHParams bparams;
       bparams.use_spatial_split = params->use_bvh_spatial_split;
+      bparams.use_compact_structure = params->use_bvh_compact_structure;
       bparams.bvh_layout = bvh_layout;
       bparams.use_unaligned_nodes = dscene->data.bvh.have_curves &&
                                     params->use_bvh_unaligned_nodes;

intern/cycles/scene/image.cpp

@@ -576,13 +576,13 @@ bool ImageManager::file_load_image(Image *img, int texture_limit)
         pixels[i * 4 + 3] = one;
       }
     }
+  }
 
-    if (img->metadata.colorspace != u_colorspace_raw &&
-        img->metadata.colorspace != u_colorspace_srgb) {
-      /* Convert to scene linear. */
-      ColorSpaceManager::to_scene_linear(
-          img->metadata.colorspace, pixels, num_pixels, img->metadata.compress_as_srgb);
-    }
+  if (img->metadata.colorspace != u_colorspace_raw &&
+      img->metadata.colorspace != u_colorspace_srgb) {
+    /* Convert to scene linear. */
+    ColorSpaceManager::to_scene_linear(
+        img->metadata.colorspace, pixels, num_pixels, is_rgba, img->metadata.compress_as_srgb);
   }
 
   /* Make sure we don't have buggy values. */
@@ -891,6 +891,10 @@ void ImageManager::device_free(Device *device)
 void ImageManager::collect_statistics(RenderStats *stats)
 {
   foreach (const Image *image, images) {
+    if (!image) {
+      /* Image may have been freed due to lack of users. */
+      continue;
+    }
     stats->image.textures.add_entry(
         NamedSizeEntry(image->loader->name(), image->mem->memory_size()));
   }

intern/cycles/scene/mesh.cpp

@@ -141,6 +141,9 @@ NODE_DEFINE(Mesh)
   subdivision_type_enum.insert("catmull_clark", SUBDIVISION_CATMULL_CLARK);
   SOCKET_ENUM(subdivision_type, "Subdivision Type", subdivision_type_enum, SUBDIVISION_NONE);
 
+  SOCKET_INT_ARRAY(subd_vert_creases, "Subdivision Vertex Crease", array<int>());
+  SOCKET_FLOAT_ARRAY(
+      subd_vert_creases_weight, "Subdivision Vertex Crease Weights", array<float>());
   SOCKET_INT_ARRAY(subd_creases_edge, "Subdivision Crease Edges", array<int>());
   SOCKET_FLOAT_ARRAY(subd_creases_weight, "Subdivision Crease Weights", array<float>());
   SOCKET_INT_ARRAY(subd_face_corners, "Subdivision Face Corners", array<int>());
@@ -408,7 +411,7 @@ Mesh::SubdFace Mesh::get_subd_face(size_t index) const
   return s;
 }
 
-void Mesh::add_crease(int v0, int v1, float weight)
+void Mesh::add_edge_crease(int v0, int v1, float weight)
 {
   subd_creases_edge.push_back_slow(v0);
   subd_creases_edge.push_back_slow(v1);
@@ -419,6 +422,17 @@ void Mesh::add_crease(int v0, int v1, float weight)
   tag_subd_creases_weight_modified();
 }
 
+void Mesh::add_vertex_crease(int v, float weight)
+{
+  assert(v < verts.size());
+
+  subd_vert_creases.push_back_slow(v);
+  subd_vert_creases_weight.push_back_slow(weight);
+
+  tag_subd_vert_creases_modified();
+  tag_subd_vert_creases_weight_modified();
+}
+
 void Mesh::copy_center_to_motion_step(const int motion_step)
 {
   Attribute *attr_mP = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);

intern/cycles/scene/mesh.h

@@ -160,6 +160,9 @@ class Mesh : public Geometry {
   NODE_SOCKET_API_ARRAY(array<int>, subd_creases_edge)
   NODE_SOCKET_API_ARRAY(array<float>, subd_creases_weight)
 
+  NODE_SOCKET_API_ARRAY(array<int>, subd_vert_creases)
+  NODE_SOCKET_API_ARRAY(array<float>, subd_vert_creases_weight)
+
   /* Subdivisions parameters */
   NODE_SOCKET_API(float, subd_dicing_rate)
   NODE_SOCKET_API(int, subd_max_level)
@@ -210,7 +213,8 @@ class Mesh : public Geometry {
   void add_vertex_slow(float3 P);
   void add_triangle(int v0, int v1, int v2, int shader, bool smooth);
   void add_subd_face(int *corners, int num_corners, int shader_, bool smooth_);
-  void add_crease(int v0, int v1, float weight);
+  void add_edge_crease(int v0, int v1, float weight);
+  void add_vertex_crease(int v, float weight);
 
   void copy_center_to_motion_step(const int motion_step);
 

intern/cycles/scene/mesh_subdivision.cpp

@@ -82,24 +82,54 @@ template<>
 bool TopologyRefinerFactory<ccl::Mesh>::assignComponentTags(TopologyRefiner &refiner,
                                                             ccl::Mesh const &mesh)
 {
+  /* Historical maximum crease weight used at Pixar, influencing the maximum in OpenSubDiv. */
+  static constexpr float CREASE_SCALE = 10.0f;
+
   size_t num_creases = mesh.get_subd_creases_weight().size();
+  size_t num_vertex_creases = mesh.get_subd_vert_creases().size();
+
+  /* The last loop is over the vertices, so early exit to avoid iterating them needlessly. */
+  if (num_creases == 0 && num_vertex_creases == 0) {
+    return true;
+  }
 
   for (int i = 0; i < num_creases; i++) {
     ccl::Mesh::SubdEdgeCrease crease = mesh.get_subd_crease(i);
     Index edge = findBaseEdge(refiner, crease.v[0], crease.v[1]);
 
     if (edge != INDEX_INVALID) {
-      setBaseEdgeSharpness(refiner, edge, crease.crease * 10.0f);
+      setBaseEdgeSharpness(refiner, edge, crease.crease * CREASE_SCALE);
     }
   }
 
+  std::map<int, float> vertex_creases;
+
+  for (size_t i = 0; i < num_vertex_creases; ++i) {
+    const int vertex_idx = mesh.get_subd_vert_creases()[i];
+    const float weight = mesh.get_subd_vert_creases_weight()[i];
+
+    vertex_creases[vertex_idx] = weight * CREASE_SCALE;
+  }
+
   for (int i = 0; i < mesh.get_verts().size(); i++) {
+    float sharpness = 0.0f;
+    std::map<int, float>::const_iterator iter = vertex_creases.find(i);
+
+    if (iter != vertex_creases.end()) {
+      sharpness = iter->second;
+    }
+
     ConstIndexArray vert_edges = getBaseVertexEdges(refiner, i);
 
     if (vert_edges.size() == 2) {
-      float sharpness = refiner.getLevel(0).getEdgeSharpness(vert_edges[0]);
-      sharpness = ccl::min(sharpness, refiner.getLevel(0).getEdgeSharpness(vert_edges[1]));
+      const float sharpness0 = refiner.getLevel(0).getEdgeSharpness(vert_edges[0]);
+      const float sharpness1 = refiner.getLevel(0).getEdgeSharpness(vert_edges[1]);
 
+      sharpness += ccl::min(sharpness0, sharpness1);
+      sharpness = ccl::min(sharpness, CREASE_SCALE);
+    }
+
+    if (sharpness != 0.0f) {
       setBaseVertexSharpness(refiner, i, sharpness);
     }
   }