Merge branch 'blender-v3.1-release' into fix-normals-custom-data

Limit the min and max of the IDProperty for the node group input
from 0 to infinity, and the soft min and max between 0 and 1.

Thanks to @PratikPB2123 for investigation.

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()
 
 
@@ -683,7 +685,7 @@ if(WIN32 OR XCODE)
   option(IDE_GROUP_PROJECTS_IN_FOLDERS "Organize the projects according to source folder structure." ON)
   mark_as_advanced(IDE_GROUP_PROJECTS_IN_FOLDERS)
 
-  if (IDE_GROUP_PROJECTS_IN_FOLDERS)
+  if(IDE_GROUP_PROJECTS_IN_FOLDERS)
     set_property(GLOBAL PROPERTY USE_FOLDERS ON)
   endif()
 endif()

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/freetype.cmake

@@ -19,13 +19,10 @@
 set(FREETYPE_EXTRA_ARGS
   -DCMAKE_RELEASE_POSTFIX:STRING=2ST
   -DCMAKE_DEBUG_POSTFIX:STRING=2ST_d
-  -DWITH_BZip2=OFF
-  -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
+  -DFT_DISABLE_BZIP2=ON
+  -DFT_DISABLE_HARFBUZZ=ON
+  -DFT_DISABLE_PNG=ON
+  -DFT_REQUIRE_BROTLI=ON
   -DPC_BROTLIDEC_INCLUDEDIR=${LIBDIR}/brotli/include
   -DPC_BROTLIDEC_LIBDIR=${LIBDIR}/brotli/lib
   )

build_files/build_environment/cmake/python_site_packages.cmake

@@ -31,7 +31,7 @@ ExternalProject_Add(external_python_site_packages
   CONFIGURE_COMMAND ${PIP_CONFIGURE_COMMAND}
   BUILD_COMMAND ""
   PREFIX ${BUILD_DIR}/site_packages
-  INSTALL_COMMAND ${PYTHON_BINARY} -m pip install ${SITE_PACKAGES_EXTRA} cython==${CYTHON_VERSION} idna==${IDNA_VERSION} charset-normalizer==${CHARSET_NORMALIZER_VERSION} urllib3==${URLLIB3_VERSION} certifi==${CERTIFI_VERSION} requests==${REQUESTS_VERSION} zstandard==${ZSTANDARD_VERSION} --no-binary :all:
+  INSTALL_COMMAND ${PYTHON_BINARY} -m pip install --no-cache-dir ${SITE_PACKAGES_EXTRA} cython==${CYTHON_VERSION} idna==${IDNA_VERSION} charset-normalizer==${CHARSET_NORMALIZER_VERSION} urllib3==${URLLIB3_VERSION} certifi==${CERTIFI_VERSION} requests==${REQUESTS_VERSION} zstandard==${ZSTANDARD_VERSION} --no-binary :all:
 )
 
 if(USE_PIP_NUMPY)

build_files/build_environment/cmake/versions.cmake

@@ -83,9 +83,9 @@ else()
   set(OPENEXR_VERSION_POSTFIX)
 endif()
 
-set(FREETYPE_VERSION 2.11.0)
+set(FREETYPE_VERSION 2.11.1)
 set(FREETYPE_URI http://prdownloads.sourceforge.net/freetype/freetype-${FREETYPE_VERSION}.tar.gz)
-set(FREETYPE_HASH cf09172322f6b50cf8f568bf8fe14bde)
+set(FREETYPE_HASH bd4e3b007474319909a6b79d50908e85)
 set(FREETYPE_HASH_TYPE MD5)
 set(FREETYPE_FILE freetype-${FREETYPE_VERSION}.tar.gz)
 
@@ -189,11 +189,11 @@ set(OSL_HASH 1abd7ce40481771a9fa937f19595d2f2)
 set(OSL_HASH_TYPE MD5)
 set(OSL_FILE OpenShadingLanguage-${OSL_VERSION}.tar.gz)
 
-set(PYTHON_VERSION 3.9.7)
-set(PYTHON_SHORT_VERSION 3.9)
-set(PYTHON_SHORT_VERSION_NO_DOTS 39)
+set(PYTHON_VERSION 3.10.2)
+set(PYTHON_SHORT_VERSION 3.10)
+set(PYTHON_SHORT_VERSION_NO_DOTS 310)
 set(PYTHON_URI https://www.python.org/ftp/python/${PYTHON_VERSION}/Python-${PYTHON_VERSION}.tar.xz)
-set(PYTHON_HASH fddb060b483bc01850a3f412eea1d954)
+set(PYTHON_HASH 14e8c22458ed7779a1957b26cde01db9)
 set(PYTHON_HASH_TYPE MD5)
 set(PYTHON_FILE Python-${PYTHON_VERSION}.tar.xz)
 
@@ -215,18 +215,20 @@ set(NANOVDB_HASH e7b9e863ec2f3b04ead171dec2322807)
 set(NANOVDB_HASH_TYPE MD5)
 set(NANOVDB_FILE nano-vdb-${NANOVDB_GIT_UID}.tar.gz)
 
-set(IDNA_VERSION 3.2)
-set(CHARSET_NORMALIZER_VERSION 2.0.6)
-set(URLLIB3_VERSION 1.26.7)
+set(IDNA_VERSION 3.3)
+set(CHARSET_NORMALIZER_VERSION 2.0.10)
+set(URLLIB3_VERSION 1.26.8)
 set(CERTIFI_VERSION 2021.10.8)
-set(REQUESTS_VERSION 2.26.0)
-set(CYTHON_VERSION 0.29.24)
-set(ZSTANDARD_VERSION 0.15.2 )
+set(REQUESTS_VERSION 2.27.1)
+set(CYTHON_VERSION 0.29.26)
+# The version of the zstd library used to build the Python package should match ZSTD_VERSION defined below.
+# At this time of writing, 0.17.0 was already released, but built against zstd 1.5.1, while we use 1.5.0.
+set(ZSTANDARD_VERSION 0.16.0)
 
-set(NUMPY_VERSION 1.21.2)
-set(NUMPY_SHORT_VERSION 1.21)
+set(NUMPY_VERSION 1.22.0)
+set(NUMPY_SHORT_VERSION 1.22)
 set(NUMPY_URI https://github.com/numpy/numpy/releases/download/v${NUMPY_VERSION}/numpy-${NUMPY_VERSION}.zip)
-set(NUMPY_HASH 5638d5dae3ca387be562912312db842e)
+set(NUMPY_HASH 252de134862a27bd66705d29622edbfe)
 set(NUMPY_HASH_TYPE MD5)
 set(NUMPY_FILE numpy-${NUMPY_VERSION}.zip)
 

build_files/build_environment/install_deps.sh

@@ -379,27 +379,27 @@ USE_CXX11=true
 CLANG_FORMAT_VERSION_MIN="6.0"
 CLANG_FORMAT_VERSION_MEX="10.0"
 
-PYTHON_VERSION="3.9.7"
-PYTHON_VERSION_SHORT="3.9"
-PYTHON_VERSION_MIN="3.7"
-PYTHON_VERSION_MEX="3.11"
+PYTHON_VERSION="3.10.2"
+PYTHON_VERSION_SHORT="3.10"
+PYTHON_VERSION_MIN="3.9"
+PYTHON_VERSION_MEX="3.12"
 PYTHON_VERSION_INSTALLED=$PYTHON_VERSION_SHORT
 PYTHON_FORCE_BUILD=false
 PYTHON_FORCE_REBUILD=false
 PYTHON_SKIP=false
 
 # Additional Python modules.
-PYTHON_IDNA_VERSION="3.2"
+PYTHON_IDNA_VERSION="3.3"
 PYTHON_IDNA_VERSION_MIN="2.0"
 PYTHON_IDNA_VERSION_MEX="4.0"
 PYTHON_IDNA_NAME="idna"
 
-PYTHON_CHARSET_NORMALIZER_VERSION="2.0.6"
+PYTHON_CHARSET_NORMALIZER_VERSION="2.0.10"
 PYTHON_CHARSET_NORMALIZER_VERSION_MIN="2.0.6"
 PYTHON_CHARSET_NORMALIZER_VERSION_MEX="2.1.0"  # requests uses `charset_normalizer~=2.0.0`
 PYTHON_CHARSET_NORMALIZER_NAME="charset-normalizer"
 
-PYTHON_URLLIB3_VERSION="1.26.7"
+PYTHON_URLLIB3_VERSION="1.26.8"
 PYTHON_URLLIB3_VERSION_MIN="1.0"
 PYTHON_URLLIB3_VERSION_MEX="2.0"
 PYTHON_URLLIB3_NAME="urllib3"
@@ -409,17 +409,17 @@ PYTHON_CERTIFI_VERSION_MIN="2021.0"
 PYTHON_CERTIFI_VERSION_MEX="2023.0"
 PYTHON_CERTIFI_NAME="certifi"
 
-PYTHON_REQUESTS_VERSION="2.23.0"
+PYTHON_REQUESTS_VERSION="2.27.1"
 PYTHON_REQUESTS_VERSION_MIN="2.0"
 PYTHON_REQUESTS_VERSION_MEX="3.0"
 PYTHON_REQUESTS_NAME="requests"
 
-PYTHON_ZSTANDARD_VERSION="0.15.2"
+PYTHON_ZSTANDARD_VERSION="0.16.0"
 PYTHON_ZSTANDARD_VERSION_MIN="0.15.2"
-PYTHON_ZSTANDARD_VERSION_MEX="0.16.0"
+PYTHON_ZSTANDARD_VERSION_MEX="0.20.0"
 PYTHON_ZSTANDARD_NAME="zstandard"
 
-PYTHON_NUMPY_VERSION="1.21.2"
+PYTHON_NUMPY_VERSION="1.22.0"
 PYTHON_NUMPY_VERSION_MIN="1.14"
 PYTHON_NUMPY_VERSION_MEX="2.0"
 PYTHON_NUMPY_NAME="numpy"
@@ -499,7 +499,7 @@ LLVM_FORCE_REBUILD=false
 LLVM_SKIP=false
 
 # OSL needs to be compiled for now!
-OSL_VERSION="1.11.14.1"
+OSL_VERSION="1.11.17.0"
 OSL_VERSION_SHORT="1.11"
 OSL_VERSION_MIN="1.11"
 OSL_VERSION_MEX="2.0"

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/Modules/FindFFmpeg.cmake

@@ -82,4 +82,6 @@ mark_as_advanced(
 
 unset(_ffmpeg_SEARCH_DIRS)
 unset(_ffmpeg_LIBRARIES)
-unset(_ffmpeg_INCLUDE_DIR)
+# In cmake version 3.21 and up, we can instead use the NO_CACHE option for
+# find_path so we don't need to clear it from the cache here.
+unset(_ffmpeg_INCLUDE_DIR CACHE)

build_files/cmake/Modules/FindOSL.cmake

@@ -76,6 +76,7 @@ FIND_PATH(OSL_SHADER_DIR
     /usr/include/OSL/
   PATH_SUFFIXES
     share/OSL/shaders
+    shaders
 )
 
 # handle the QUIETLY and REQUIRED arguments and set OSL_FOUND to TRUE if
@@ -99,6 +100,7 @@ ENDIF()
 
 MARK_AS_ADVANCED(
   OSL_INCLUDE_DIR
+  OSL_SHADER_DIR
 )
 FOREACH(COMPONENT ${_osl_FIND_COMPONENTS})
   STRING(TOUPPER ${COMPONENT} UPPERCOMPONENT)

build_files/cmake/Modules/FindOpenColorIO.cmake

@@ -87,12 +87,14 @@ ENDIF()
 MARK_AS_ADVANCED(
   OPENCOLORIO_INCLUDE_DIR
   OPENCOLORIO_LIBRARY
-  OPENCOLORIO_OPENCOLORIO_LIBRARY
-  OPENCOLORIO_TINYXML_LIBRARY
-  OPENCOLORIO_YAML-CPP_LIBRARY
   OPENCOLORIO_VERSION
 )
 
+FOREACH(COMPONENT ${_opencolorio_FIND_COMPONENTS})
+  STRING(TOUPPER ${COMPONENT} UPPERCOMPONENT)
+  MARK_AS_ADVANCED(OPENCOLORIO_${UPPERCOMPONENT}_LIBRARY)
+ENDFOREACH()
+
 UNSET(COMPONENT)
 UNSET(UPPERCOMPONENT)
 UNSET(_opencolorio_FIND_COMPONENTS)

build_files/cmake/Modules/FindOpenEXR.cmake

@@ -33,14 +33,6 @@ ENDIF()
 # Old versions (before 2.0?) do not have any version string, just assuming this should be fine though.
 SET(_openexr_libs_ver_init "2.0")
 
-SET(_openexr_FIND_COMPONENTS
-  Half
-  Iex
-  IlmImf
-  IlmThread
-  Imath
-)
-
 SET(_openexr_SEARCH_DIRS
   ${OPENEXR_ROOT_DIR}
   /opt/lib/openexr
@@ -93,6 +85,24 @@ UNSET(_openexr_libs_ver_init)
 
 STRING(REGEX REPLACE "([0-9]+)[.]([0-9]+).*" "\\1_\\2" _openexr_libs_ver ${OPENEXR_VERSION})
 
+# Different library names in 3.0, and Imath and Half moved out.
+IF(OPENEXR_VERSION VERSION_GREATER_EQUAL "3.0.0")
+  SET(_openexr_FIND_COMPONENTS
+    Iex
+    IlmThread
+    OpenEXR
+    OpenEXRCore
+  )
+ELSE()
+  SET(_openexr_FIND_COMPONENTS
+    Half
+    Iex
+    IlmImf
+    IlmThread
+    Imath
+  )
+ENDIF()
+
 SET(_openexr_LIBRARIES)
 FOREACH(COMPONENT ${_openexr_FIND_COMPONENTS})
   STRING(TOUPPER ${COMPONENT} UPPERCOMPONENT)
@@ -111,6 +121,57 @@ ENDFOREACH()
 
 UNSET(_openexr_libs_ver)
 
+IF(OPENEXR_VERSION VERSION_GREATER_EQUAL "3.0.0")
+  # For OpenEXR 3.x, we also need to find the now separate Imath library.
+  # For simplicity we add it to the OpenEXR includes and libraries, as we
+  # have no direct dependency on Imath and it's simpler to support both
+  # 2.x and 3.x this way.
+
+  # Find include directory
+  FIND_PATH(IMATH_INCLUDE_DIR
+    NAMES
+      Imath/ImathMath.h
+    HINTS
+      ${_openexr_SEARCH_DIRS}
+    PATH_SUFFIXES
+      include
+  )
+
+  # Find version
+  FIND_FILE(_imath_config
+    NAMES
+      ImathConfig.h
+    PATHS
+      ${IMATH_INCLUDE_DIR}/Imath
+    NO_DEFAULT_PATH
+  )
+
+  # Find line with version, extract string, and format for library suffix.
+  FILE(STRINGS "${_imath_config}" _imath_build_specification
+       REGEX "^[ \t]*#define[ \t]+IMATH_VERSION_STRING[ \t]+\"[.0-9]+\".*$")
+  STRING(REGEX REPLACE ".*#define[ \t]+IMATH_VERSION_STRING[ \t]+\"([.0-9]+)\".*"
+         "\\1" _imath_libs_ver ${_imath_build_specification})
+  STRING(REGEX REPLACE "([0-9]+)[.]([0-9]+).*" "\\1_\\2" _imath_libs_ver ${_imath_libs_ver})
+
+  # Find library, with or without version number.
+  FIND_LIBRARY(IMATH_LIBRARY
+    NAMES
+      Imath-${_imath_libs_ver} Imath
+    NAMES_PER_DIR
+    HINTS
+      ${_openexr_SEARCH_DIRS}
+    PATH_SUFFIXES
+      lib64 lib
+    )
+  LIST(APPEND _openexr_LIBRARIES "${IMATH_LIBRARY}")
+
+  # In cmake version 3.21 and up, we can instead use the NO_CACHE option for
+  # FIND_FILE so we don't need to clear it from the cache here.
+  UNSET(_imath_config CACHE)
+  UNSET(_imath_libs_ver)
+  UNSET(_imath_build_specification)
+ENDIF()
+
 # handle the QUIETLY and REQUIRED arguments and set OPENEXR_FOUND to TRUE if
 # all listed variables are TRUE
 INCLUDE(FindPackageHandleStandardArgs)
@@ -119,13 +180,25 @@ FIND_PACKAGE_HANDLE_STANDARD_ARGS(OpenEXR  DEFAULT_MSG
 
 IF(OPENEXR_FOUND)
   SET(OPENEXR_LIBRARIES ${_openexr_LIBRARIES})
-  # Both include paths are needed because of dummy OSL headers mixing #include <OpenEXR/foo.h> and #include <foo.h> :(
-  SET(OPENEXR_INCLUDE_DIRS ${OPENEXR_INCLUDE_DIR} ${OPENEXR_INCLUDE_DIR}/OpenEXR)
+  # Both include paths are needed because of dummy OSL headers mixing
+  # #include <OpenEXR/foo.h> and #include <foo.h>, as well as Alembic
+  # include <half.h> directly.
+  SET(OPENEXR_INCLUDE_DIRS
+    ${OPENEXR_INCLUDE_DIR}
+    ${OPENEXR_INCLUDE_DIR}/OpenEXR)
+
+  IF(OPENEXR_VERSION VERSION_GREATER_EQUAL "3.0.0")
+    LIST(APPEND OPENEXR_INCLUDE_DIRS
+      ${IMATH_INCLUDE_DIR}
+      ${IMATH_INCLUDE_DIR}/Imath)
+  ENDIF()
 ENDIF()
 
 MARK_AS_ADVANCED(
   OPENEXR_INCLUDE_DIR
   OPENEXR_VERSION
+  IMATH_INCLUDE_DIR
+  IMATH_LIBRARY
 )
 FOREACH(COMPONENT ${_openexr_FIND_COMPONENTS})
   STRING(TOUPPER ${COMPONENT} UPPERCOMPONENT)

build_files/cmake/Modules/FindOpenImageDenoise.cmake

@@ -110,6 +110,7 @@ ENDIF()
 
 MARK_AS_ADVANCED(
   OPENIMAGEDENOISE_INCLUDE_DIR
+  OPENIMAGEDENOISE_LIBRARY
 )
 
 FOREACH(COMPONENT ${_openimagedenoise_FIND_COMPONENTS})

build_files/cmake/Modules/FindOpenImageIO.cmake

@@ -48,6 +48,8 @@ FIND_LIBRARY(OPENIMAGEIO_LIBRARY
     lib64 lib
   )
 
+set(_openimageio_LIBRARIES ${OPENIMAGEIO_LIBRARY})
+
 FIND_FILE(OPENIMAGEIO_IDIFF
   NAMES
     idiff
@@ -57,14 +59,47 @@ FIND_FILE(OPENIMAGEIO_IDIFF
     bin
 )
 
+# Additionally find util library if needed. In old versions this library was
+# included in libOpenImageIO and linking to both would duplicate symbols. In
+# new versions we need to link to both.
+FIND_FILE(_openimageio_export
+  NAMES
+    export.h
+  PATHS
+    ${OPENIMAGEIO_INCLUDE_DIR}/OpenImageIO
+  NO_DEFAULT_PATH
+)
+
+# Use existence of OIIO_UTIL_API to check if it's a separate lib.
+FILE(STRINGS "${_openimageio_export}" _openimageio_util_define
+     REGEX "^[ \t]*#[ \t]*define[ \t]+OIIO_UTIL_API.*$")
+
+IF(_openimageio_util_define)
+  FIND_LIBRARY(OPENIMAGEIO_UTIL_LIBRARY
+    NAMES
+      OpenImageIO_Util
+    HINTS
+      ${_openimageio_SEARCH_DIRS}
+    PATH_SUFFIXES
+      lib64 lib
+    )
+
+  LIST(APPEND _openimageio_LIBRARIES ${OPENIMAGEIO_UTIL_LIBRARY})
+ENDIF()
+
+# In cmake version 3.21 and up, we can instead use the NO_CACHE option for
+# FIND_FILE so we don't need to clear it from the cache here.
+UNSET(_openimageio_export CACHE)
+UNSET(_openimageio_util_define)
+
 # handle the QUIETLY and REQUIRED arguments and set OPENIMAGEIO_FOUND to TRUE if
 # all listed variables are TRUE
 INCLUDE(FindPackageHandleStandardArgs)
 FIND_PACKAGE_HANDLE_STANDARD_ARGS(OpenImageIO DEFAULT_MSG
-    OPENIMAGEIO_LIBRARY OPENIMAGEIO_INCLUDE_DIR)
+    _openimageio_LIBRARIES OPENIMAGEIO_INCLUDE_DIR)
 
 IF(OPENIMAGEIO_FOUND)
-  SET(OPENIMAGEIO_LIBRARIES ${OPENIMAGEIO_LIBRARY})
+  SET(OPENIMAGEIO_LIBRARIES ${_openimageio_LIBRARIES})
   SET(OPENIMAGEIO_INCLUDE_DIRS ${OPENIMAGEIO_INCLUDE_DIR})
   IF(EXISTS ${OPENIMAGEIO_INCLUDE_DIR}/OpenImageIO/pugixml.hpp)
     SET(OPENIMAGEIO_PUGIXML_FOUND TRUE)
@@ -78,7 +113,9 @@ ENDIF()
 MARK_AS_ADVANCED(
   OPENIMAGEIO_INCLUDE_DIR
   OPENIMAGEIO_LIBRARY
+  OPENIMAGEIO_UTIL_LIBRARY
   OPENIMAGEIO_IDIFF
 )
 
 UNSET(_openimageio_SEARCH_DIRS)
+UNSET(_openimageio_LIBRARIES)

build_files/cmake/Modules/FindPythonLibsUnix.cmake

@@ -34,7 +34,7 @@ IF(NOT PYTHON_ROOT_DIR AND NOT $ENV{PYTHON_ROOT_DIR} STREQUAL "")
   SET(PYTHON_ROOT_DIR $ENV{PYTHON_ROOT_DIR})
 ENDIF()
 
-SET(PYTHON_VERSION 3.9 CACHE STRING "Python Version (major and minor only)")
+SET(PYTHON_VERSION 3.10 CACHE STRING "Python Version (major and minor only)")
 MARK_AS_ADVANCED(PYTHON_VERSION)
 
 

build_files/cmake/macros.cmake

@@ -1197,21 +1197,21 @@ endfunction()
 macro(openmp_delayload
   projectname
   )
-    if(MSVC)
-      if(WITH_OPENMP)
-        if(MSVC_CLANG)
-          set(OPENMP_DLL_NAME "libomp")
-        elseif(MSVC_VERSION EQUAL 1800)
-          set(OPENMP_DLL_NAME "vcomp120")
-        else()
-          set(OPENMP_DLL_NAME "vcomp140")
-        endif()
-        set_property(TARGET ${projectname} APPEND_STRING  PROPERTY LINK_FLAGS_RELEASE " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
-        set_property(TARGET ${projectname} APPEND_STRING  PROPERTY LINK_FLAGS_DEBUG " /DELAYLOAD:${OPENMP_DLL_NAME}d.dll delayimp.lib")
-        set_property(TARGET ${projectname} APPEND_STRING  PROPERTY LINK_FLAGS_RELWITHDEBINFO " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
-        set_property(TARGET ${projectname} APPEND_STRING  PROPERTY LINK_FLAGS_MINSIZEREL " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
+  if(MSVC)
+    if(WITH_OPENMP)
+      if(MSVC_CLANG)
+        set(OPENMP_DLL_NAME "libomp")
+      elseif(MSVC_VERSION EQUAL 1800)
+        set(OPENMP_DLL_NAME "vcomp120")
+      else()
+        set(OPENMP_DLL_NAME "vcomp140")
       endif()
+      set_property(TARGET ${projectname} APPEND_STRING  PROPERTY LINK_FLAGS_RELEASE " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
+      set_property(TARGET ${projectname} APPEND_STRING  PROPERTY LINK_FLAGS_DEBUG " /DELAYLOAD:${OPENMP_DLL_NAME}d.dll delayimp.lib")
+      set_property(TARGET ${projectname} APPEND_STRING  PROPERTY LINK_FLAGS_RELWITHDEBINFO " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
+      set_property(TARGET ${projectname} APPEND_STRING  PROPERTY LINK_FLAGS_MINSIZEREL " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
     endif()
+  endif()
 endmacro()
 
 macro(set_and_warn_dependency

build_files/cmake/platform/platform_apple.cmake

@@ -128,25 +128,20 @@ if(WITH_CODEC_SNDFILE)
 endif()
 
 if(WITH_PYTHON)
-  # we use precompiled libraries for py 3.9 and up by default
-  set(PYTHON_VERSION 3.9)
+  # Use precompiled libraries by default.
+  set(PYTHON_VERSION 3.10)
   if(NOT WITH_PYTHON_MODULE AND NOT WITH_PYTHON_FRAMEWORK)
-    # normally cached but not since we include them with blender
+    # Normally cached but not since we include them with blender.
     set(PYTHON_INCLUDE_DIR "${LIBDIR}/python/include/python${PYTHON_VERSION}")
     set(PYTHON_EXECUTABLE "${LIBDIR}/python/bin/python${PYTHON_VERSION}")
     set(PYTHON_LIBRARY ${LIBDIR}/python/lib/libpython${PYTHON_VERSION}.a)
     set(PYTHON_LIBPATH "${LIBDIR}/python/lib/python${PYTHON_VERSION}")
-    # set(PYTHON_LINKFLAGS "-u _PyMac_Error")  # won't  build with this enabled
   else()
-    # module must be compiled against Python framework
+    # Module must be compiled against Python framework.
     set(_py_framework "/Library/Frameworks/Python.framework/Versions/${PYTHON_VERSION}")
-
     set(PYTHON_INCLUDE_DIR "${_py_framework}/include/python${PYTHON_VERSION}")
     set(PYTHON_EXECUTABLE "${_py_framework}/bin/python${PYTHON_VERSION}")
     set(PYTHON_LIBPATH "${_py_framework}/lib/python${PYTHON_VERSION}")
-    # set(PYTHON_LIBRARY python${PYTHON_VERSION})
-    # set(PYTHON_LINKFLAGS "-u _PyMac_Error -framework Python")  # won't  build with this enabled
-
     unset(_py_framework)
   endif()
 

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}")
 
@@ -101,12 +104,21 @@ find_package_wrapper(PNG REQUIRED)
 find_package_wrapper(ZLIB REQUIRED)
 find_package_wrapper(Zstd REQUIRED)
 
-# FreeType compiled with Brotli compression for woff2.
-find_package_wrapper(Freetype REQUIRED)
-list(APPEND FREETYPE_LIBRARIES
-  ${LIBDIR}/brotli/lib/libbrotlidec-static.a
-  ${LIBDIR}/brotli/lib/libbrotlicommon-static.a
-)
+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 directly after `FREETYPE_LIBRARIES`.
+    #
+    # list(APPEND FREETYPE_LIBRARIES
+    #   ${BROTLI_LIBRARIES}
+    # )
+  endif()
+endif()
 
 if(WITH_PYTHON)
   # No way to set py35, remove for now.
@@ -350,6 +362,7 @@ if(WITH_BOOST)
       find_package(IcuLinux)
     endif()
     mark_as_advanced(Boost_DIR)  # why doesn't boost do this?
+    mark_as_advanced(Boost_INCLUDE_DIR)  # why doesn't boost do this?
   endif()
 
   set(BOOST_INCLUDE_DIR ${Boost_INCLUDE_DIRS})
@@ -542,6 +555,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)

build_files/cmake/platform/platform_win32.cmake

@@ -55,6 +55,10 @@ if(CMAKE_C_COMPILER_ID MATCHES "Clang")
     message(WARNING "stripped pdb not supported with clang, disabling..")
     set(WITH_WINDOWS_STRIPPED_PDB OFF)
   endif()
+else()
+  if(CMAKE_CXX_COMPILER_VERSION VERSION_LESS 19.28.29921) # MSVC 2019 16.9.16
+    message(FATAL_ERROR "Compiler is unsupported, MSVC 2019 16.9.16 or newer is required for building blender.")
+  endif()
 endif()
 
 if(NOT WITH_PYTHON_MODULE)
@@ -265,12 +269,6 @@ if(NOT DEFINED LIBDIR)
   elseif(MSVC_VERSION GREATER 1919)
     message(STATUS "Visual Studio 2019 detected.")
     set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_BASE}_vc15)
-  elseif(MSVC_VERSION GREATER 1909)
-    message(STATUS "Visual Studio 2017 detected.")
-    set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_BASE}_vc15)
-  elseif(MSVC_VERSION EQUAL 1900)
-    message(STATUS "Visual Studio 2015 detected.")
-    set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_BASE}_vc15)
   endif()
 else()
   message(STATUS "Using pre-compiled LIBDIR: ${LIBDIR}")
@@ -465,7 +463,7 @@ if(WITH_JACK)
 endif()
 
 if(WITH_PYTHON)
-  set(PYTHON_VERSION 3.9) # CACHE STRING)
+  set(PYTHON_VERSION 3.10) # CACHE STRING)
 
   string(REPLACE "." "" _PYTHON_VERSION_NO_DOTS ${PYTHON_VERSION})
   set(PYTHON_LIBRARY ${LIBDIR}/python/${_PYTHON_VERSION_NO_DOTS}/libs/python${_PYTHON_VERSION_NO_DOTS}.lib)

build_files/config/pipeline_config.yaml

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

build_files/windows/autodetect_msvc.cmd

@@ -3,9 +3,6 @@ echo No explicit msvc version requested, autodetecting version.
 call "%~dp0\detect_msvc2019.cmd"
 if %ERRORLEVEL% EQU 0 goto DetectionComplete
 
-call "%~dp0\detect_msvc2017.cmd"
-if %ERRORLEVEL% EQU 0 goto DetectionComplete
-
 call "%~dp0\detect_msvc2022.cmd"
 if %ERRORLEVEL% EQU 0 goto DetectionComplete
 

build_files/windows/check_libraries.cmd

@@ -1,4 +1,3 @@
-if "%BUILD_VS_YEAR%"=="2017" set BUILD_VS_LIBDIRPOST=vc15
 if "%BUILD_VS_YEAR%"=="2019" set BUILD_VS_LIBDIRPOST=vc15
 if "%BUILD_VS_YEAR%"=="2022" set BUILD_VS_LIBDIRPOST=vc15
 

build_files/windows/configure_msbuild.cmd

@@ -19,12 +19,6 @@ if "%WITH_PYDEBUG%"=="1" (
 	set PYDEBUG_CMAKE_ARGS=-DWINDOWS_PYTHON_DEBUG=On
 )
 
-if "%BUILD_VS_YEAR%"=="2017" (
-	set BUILD_GENERATOR_POST=%WINDOWS_ARCH%
-) else (
-	set BUILD_PLATFORM_SELECT=-A %MSBUILD_PLATFORM%
-)
-
 set BUILD_CMAKE_ARGS=%BUILD_CMAKE_ARGS% -G "Visual Studio %BUILD_VS_VER% %BUILD_VS_YEAR%%BUILD_GENERATOR_POST%" %BUILD_PLATFORM_SELECT% %TESTS_CMAKE_ARGS% %CLANG_CMAKE_ARGS% %ASAN_CMAKE_ARGS% %PYDEBUG_CMAKE_ARGS%
 
 if NOT EXIST %BUILD_DIR%\nul (

build_files/windows/configure_ninja.cmd

@@ -37,15 +37,9 @@ set LLVM_DIR=
 :DetectionComplete	
 	set CC=%LLVM_DIR%\bin\clang-cl
 	set CXX=%LLVM_DIR%\bin\clang-cl
-	if "%BUILD_VS_YEAR%" == "2019" (
-		rem build and tested against 2019 16.2
-		set CFLAGS=-m64 -fmsc-version=1922
-		set CXXFLAGS=-m64 -fmsc-version=1922
-	) else (
-		rem build and tested against 2017 15.7
-		set CFLAGS=-m64 -fmsc-version=1914
-		set CXXFLAGS=-m64 -fmsc-version=1914
-	)
+	rem build and tested against 2019 16.2
+	set CFLAGS=-m64 -fmsc-version=1922
+	set CXXFLAGS=-m64 -fmsc-version=1922
 )
 
 if "%WITH_ASAN%"=="1" (

build_files/windows/detect_msvc2017.cmd

@@ -1,3 +0,0 @@
-set BUILD_VS_VER=15
-set BUILD_VS_YEAR=2017
-call "%~dp0\detect_msvc_vswhere.cmd"

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/parse_arguments.cmd

@@ -50,14 +50,6 @@ if NOT "%1" == "" (
 		goto ERR
 	) else if "%1" == "x64" (
 		set BUILD_ARCH=x64
-	) else if "%1" == "2017" (
-		set BUILD_VS_YEAR=2017
-	) else if "%1" == "2017pre" (
-		set BUILD_VS_YEAR=2017
-		set VSWHERE_ARGS=-prerelease
-	) else if "%1" == "2017b" (
-		set BUILD_VS_YEAR=2017
-		set VSWHERE_ARGS=-products Microsoft.VisualStudio.Product.BuildTools
 	) else if "%1" == "2019" (
 		set BUILD_VS_YEAR=2019
 	) else if "%1" == "2019pre" (

build_files/windows/show_help.cmd

@@ -24,12 +24,12 @@ echo - nobuildinfo ^(disable buildinfo^)
 echo - debug ^(Build an unoptimized debuggable build^)
 echo - packagename [newname] ^(override default cpack package name^)
 echo - builddir [newdir] ^(override default build folder^)
-echo - 2017 ^(build with visual studio 2017^)
-echo - 2017pre ^(build with visual studio 2017 pre-release^)
-echo - 2017b ^(build with visual studio 2017 Build Tools^)
 echo - 2019 ^(build with visual studio 2019^)
 echo - 2019pre ^(build with visual studio 2019 pre-release^)
 echo - 2019b ^(build with visual studio 2019 Build Tools^)
+echo - 2022 ^(build with visual studio 2022^)
+echo - 2022pre ^(build with visual studio 2022 pre-release^)
+echo - 2022b ^(build with visual studio 2022 Build Tools^)
 
 echo.
 echo Documentation Targets ^(Not associated with building^)

build_files/windows/svn_fix.cmd

@@ -1,4 +1,3 @@
-if "%BUILD_VS_YEAR%"=="2017" set BUILD_VS_LIBDIRPOST=vc15
 if "%BUILD_VS_YEAR%"=="2019" set BUILD_VS_LIBDIRPOST=vc15
 if "%BUILD_VS_YEAR%"=="2022" set BUILD_VS_LIBDIRPOST=vc15
 

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/rst/info_api_reference.rst

@@ -22,7 +22,7 @@ Data Access
 ===========
 
 The most common case for using the reference API is to find out how to access data in the blend-file.
-Before going any further its best to be aware of ID data-blocks in Blender since you will often find properties
+Before going any further it's best to be aware of ID data-blocks in Blender since you will often find properties
 relative to them.
 
 
@@ -55,9 +55,9 @@ Start by collecting the information where the data is located.
 First find this setting in the interface ``Properties editor -> Object -> Transform -> Location``.
 From the button context menu select *Online Python Reference*, this will link you to:
 :class:`bpy.types.Object.location`.
-Being an API reference, this link often gives little more information then the tooltip, though some of the pages
+Being an API reference, this link often gives little more information than the tooltip, though some of the pages
 include examples (normally at the top of the page).
-But you now know that you have to use ``.location`` and that its an array of three floats.
+But you now know that you have to use ``.location`` and that it's an array of three floats.
 
 So the next step is to find out where to access objects, go down to the bottom of the page to the references section,
 for objects there are many references, but one of the most common places to access objects is via the context.
@@ -154,7 +154,7 @@ The tooltip includes :class:`bpy.types.SubsurfModifier.levels` but you want the
 
 Note that the text copied won't include the ``bpy.data.collection["name"].`` component since its assumed that
 you won't be doing collection look-ups on every access and typically you'll want to use the context rather
-then access each :class:`bpy.types.ID` instance by name.
+than access each :class:`bpy.types.ID` instance by name.
 
 Type in the ID path into a Python console :mod:`bpy.context.active_object`.
 Include the trailing dot and don't execute the code, yet.
@@ -252,6 +252,6 @@ Each entry can be selected, then copied :kbd:`Ctrl-C`, usually to paste in the t
 .. note::
 
    Not all operators get registered for display,
-   zooming the view for example isn't so useful to repeat so its excluded from the output.
+   zooming the view for example isn't so useful to repeat so it's excluded from the output.
 
    To display *every* operator that runs see :ref:`Show All Operators <info_show_all_operators>`.

doc/python_api/rst/info_best_practice.rst

@@ -229,7 +229,7 @@ removing the last items first, which is faster (as explained above):
 
 
 This example shows a fast way of removing items,
-for use in cases where you can alter the list order without breaking the scripts functionality.
+for use in cases where you can alter the list order without breaking the script's functionality.
 This works by swapping two list items, so the item you remove is always last:
 
 .. code-block:: python
@@ -278,7 +278,7 @@ Here are three ways of joining multiple strings into one string for writing.
 This also applies to any area of your code that involves a lot of string joining:
 
 String concatenation
-   This is the slowest option, do **not** use if you can avoid it, especially when writing data in a loop.
+   This is the slowest option, do **not** use this if you can avoid it, especially when writing data in a loop.
 
    >>> file.write(str1 + " " + str2 + " " + str3 + "\n")
 
@@ -288,7 +288,7 @@ String formatting
    >>> file.write("%s %s %s\n" % (str1, str2, str3))
 
 String joining
-   Use to join a list of strings (the list may be temporary). In the following example, the strings are joined with
+   Use this to join a list of strings (the list may be temporary). In the following example, the strings are joined with
    a space " " in between, other examples are "" or ", ".
 
    >>> file.write(" ".join((str1, str2, str3, "\n")))

doc/python_api/rst/info_gotcha.rst

@@ -12,7 +12,7 @@ that can be troublesome and avoid practices that are known to cause instability.
 Using Operators
 ===============
 
-Blender's operators are tools for users to access, that can access with Python too which is very useful.
+Blender's operators are tools for users to access, that can be accessed with Python too which is very useful.
 Still operators have limitations that can make them cumbersome to script.
 
 The main limits are:
@@ -20,13 +20,13 @@ The main limits are:
 - Can't pass data such as objects, meshes or materials to operate on (operators use the context instead).
 - The return value from calling an operator is the success (if it finished or was canceled),
   in some cases it would be more logical from an API perspective to return the result of the operation.
-- Operators poll function can fail where an API function would raise an exception giving details on exactly why.
+- Operators' poll function can fail where an API function would raise an exception giving details on exactly why.
 
 
 Why does an operator's poll fail?
 ---------------------------------
 
-When calling an operator gives an error like this:
+When calling an operator it gives an error like this:
 
    >>> bpy.ops.action.clean(threshold=0.001)
    RuntimeError: Operator bpy.ops.action.clean.poll() failed, context is incorrect
@@ -49,9 +49,9 @@ you should be able to find the poll function with no knowledge of C.
 .. note::
 
    Blender does have the functionality for poll functions to describe why they fail,
-   but its currently not used much, if you're interested to help improve the API
+   but it's currently not used much, if you're interested to help improve the API
    feel free to add calls to :class:`bpy.types.Operator.poll_message_set` (``CTX_wm_operator_poll_msg_set`` in C)
-   where its not obvious why poll fails, e.g:
+   where it's not obvious why poll fails, e.g:
 
       >>> bpy.ops.gpencil.draw()
       RuntimeError: Operator bpy.ops.gpencil.draw.poll() Failed to find Grease Pencil data to draw into
@@ -107,7 +107,7 @@ In this case you need to call :class:`bpy.types.ViewLayer.update` after modifyin
 
 
 Now all dependent data (child objects, modifiers, drivers, etc.)
-has been recalculated and is available to the script within active view layer.
+have been recalculated and are available to the script within the active view layer.
 
 
 Can I redraw during script execution?
@@ -116,13 +116,13 @@ Can I redraw during script execution?
 The official answer to this is no, or... *"You don't want to do that"*.
 To give some background on the topic:
 
-While a script executes Blender waits for it to finish and is effectively locked until its done,
+While a script executes, Blender waits for it to finish and is effectively locked until it's done;
 while in this state Blender won't redraw or respond to user input.
 Normally this is not such a problem because scripts distributed with Blender
 tend not to run for an extended period of time,
 nevertheless scripts *can* take a long time to complete and it would be nice to see progress in the viewport.
 
-When tools lock Blender in a loop redraw are highly discouraged
+Tools that lock Blender in a loop redraw are highly discouraged
 since they conflict with Blender's ability to run multiple operators
 at once and update different parts of the interface as the tool runs.
 
@@ -130,7 +130,7 @@ So the solution here is to write a **modal** operator, which is an operator that
 See the modal operator template in the text editor.
 Modal operators execute on user input or setup their own timers to run frequently,
 they can handle the events or pass through to be handled by the keymap or other modal operators.
-Examples of a modal operators are Transform, Painting, Fly Navigation and File Select.
+Examples of modal operators are Transform, Painting, Fly Navigation and File Select.
 
 Writing modal operators takes more effort than a simple ``for`` loop
 that contains draw calls but is more flexible and integrates better with Blender's design.
@@ -240,7 +240,7 @@ Editing
 Editing is where the three data types vary most.
 
 - Polygons are very limited for editing,
-  changing materials and options like smooth works but for anything else
+  changing materials and options like smooth works, but for anything else
   they are too inflexible and are only intended for storage.
 - Tessfaces should not be used for editing geometry because doing so will cause existing n-gons to be tessellated.
 - BMesh-faces are by far the best way to manipulate geometry.
@@ -256,7 +256,7 @@ the choice mostly depends on whether the target format supports n-gons or not.
 - Tessfaces work well for exporting to formats which don't support n-gons,
   in fact this is the only place where their use is encouraged.
 - BMesh-Faces can work for exporting too but may not be necessary if polygons can be used
-  since using BMesh gives some overhead because its not the native storage format in Object-Mode.
+  since using BMesh gives some overhead because it's not the native storage format in Object-Mode.
 
 
 Edit Bones, Pose Bones, Bone... Bones
@@ -348,7 +348,7 @@ Armature Mode Switching
 While writing scripts that deal with armatures you may find you have to switch between modes,
 when doing so take care when switching out of Edit-Mode not to keep references
 to the edit bones or their head/tail vectors.
-Further access to these will crash Blender so its important the script
+Further access to these will crash Blender so it's important that the script
 clearly separates sections of the code which operate in different modes.
 
 This is mainly an issue with Edit-Mode since pose data can be manipulated without having to be in Pose-Mode,
@@ -386,11 +386,11 @@ Or with name assignment:
 Data names may not match the assigned values if they exceed the maximum length, are already used or an empty string.
 
 
-Its better practice not to reference objects by names at all,
+It's better practice not to reference objects by names at all,
 once created you can store the data in a list, dictionary, on a class, etc;
 there is rarely a reason to have to keep searching for the same data by name.
 
-If you do need to use name references, its best to use a dictionary to maintain
+If you do need to use name references, it's best to use a dictionary to maintain
 a mapping between the names of the imported assets and the newly created data,
 this way you don't run this risk of referencing existing data from the blend-file, or worse modifying it.
 
@@ -414,11 +414,11 @@ Library Collisions
 Blender keeps data names unique (:class:`bpy.types.ID.name`) so you can't name two objects,
 meshes, scenes, etc., the same by accident.
 However, when linking in library data from another blend-file naming collisions can occur,
-so its best to avoid referencing data by name at all.
+so it's best to avoid referencing data by name at all.
 
-This can be tricky at times and not even Blender handles this correctly in some case
+This can be tricky at times and not even Blender handles this correctly in some cases
 (when selecting the modifier object for e.g. you can't select between multiple objects with the same name),
-but its still good to try avoiding these problems in this area.
+but it's still good to try avoiding these problems in this area.
 If you need to select between local and library data, there is a feature in ``bpy.data`` members to allow for this.
 
 .. code-block:: python
@@ -467,11 +467,11 @@ writing a script in ``latin1`` or ``iso-8859-15``.
 See `PEP 263 <https://www.python.org/dev/peps/pep-0263/>`__.
 
 However, this complicates matters for Blender's Python API because ``.blend`` files don't have an explicit encoding.
-To avoid the problem for Python integration and script authors we have decided all strings in blend-files
+To avoid the problem for Python integration and script authors we have decided that all strings in blend-files
 **must** be ``UTF-8``, ``ASCII`` compatible.
-This means assigning strings with different encodings to an object names for instance will raise an error.
+This means assigning strings with different encodings to an object name, for instance, will raise an error.
 
-Paths are an exception to this rule since the existence of non-UTF-8 paths on user's file system cannot be ignored.
+Paths are an exception to this rule since the existence of non-UTF-8 paths on the user's file system cannot be ignored.
 This means seemingly harmless expressions can raise errors, e.g:
 
    >>> print(bpy.data.filepath)
@@ -505,7 +505,7 @@ to keep it short about encoding problems -- here are some suggestions:
 .. note::
 
    Sometimes it's preferable to avoid string encoding issues by using bytes instead of Python strings,
-   when reading some input its less trouble to read it as binary data
+   when reading some input it's less trouble to read it as binary data
    though you will still need to decide how to treat any strings you want to use with Blender,
    some importers do this.
 
@@ -679,7 +679,7 @@ Undo/Redo
 ---------
 
 For safety, you should assume that undo and redo always invalidates all :class:`bpy.types.ID`
-instances (Object, Scene, Mesh, Light, etc.), as weel obviously as all of their sub-data.
+instances (Object, Scene, Mesh, Light, etc.), as well obviously as all of their sub-data.
 
 This example shows how you can tell undo changes the memory locations:
 
@@ -716,7 +716,7 @@ Tools in Blender are not allowed to modify library data.
 But Python does not enforce this restriction.
 
 This can be useful in some cases, using a script to adjust material values for example.
-But its also possible to use a script to make library data point to newly created local data,
+But it's also possible to use a script to make library data point to newly created local data,
 which is not supported since a call to undo will remove the local data
 but leave the library referencing it and likely crash.
 

doc/python_api/rst/info_tips_and_tricks.rst

@@ -81,7 +81,7 @@ but reference an external file rather than including it directly.
 Executing External Scripts
 --------------------------
 
-This is the equivalent to running the script directly, referencing a scripts path from a two line code block.
+This is the equivalent to running the script directly, referencing a script's path from a two line code block.
 
 .. code-block:: python
 
@@ -124,7 +124,7 @@ small script which is often useful for testing different settings quickly.
 
 The other issue with this is the script has to be in Python's module search path.
 While this is not best practice -- for testing purposes you can extend the search path,
-this following example adds the current blend-files directory to the search path
+this following example adds the current blend-file's directory to the search path
 and then loads the script as a module.
 
 .. code-block:: python
@@ -302,7 +302,7 @@ Python Safety (Build Option)
 ----------------------------
 
 Since it's possible to access data which has been removed (see :doc:`Gotchas <info_gotcha>`),
-can make it hard to track down the cause of crashes.
+it can be hard to track down the cause of crashes.
 To raise Python exceptions on accessing freed data (rather than crashing),
 enable the CMake build option ``WITH_PYTHON_SAFETY``.
 This enables data tracking which makes data access about two times slower

doc/python_api/sphinx_doc_gen.py

@@ -417,7 +417,8 @@ MODULE_GROUPING = {
 BLENDER_REVISION = str(bpy.app.build_hash, 'utf_8')
 
 # '2.83.0 Beta' or '2.83.0' or '2.83.1'
-BLENDER_VERSION_DOTS = bpy.app.version_string
+BLENDER_VERSION_STRING = bpy.app.version_string
+BLENDER_VERSION_DOTS = "%d.%d" % (bpy.app.version[0], bpy.app.version[1])
 
 if BLENDER_REVISION != "Unknown":
     # SHA1 Git hash
@@ -1724,11 +1725,11 @@ def write_sphinx_conf_py(basepath):
     fw("import sys, os\n\n")
     fw("extensions = ['sphinx.ext.intersphinx']\n\n")
     fw("intersphinx_mapping = {'blender_manual': ('https://docs.blender.org/manual/en/dev/', None)}\n\n")
-    fw("project = 'Blender %s Python API'\n" % BLENDER_VERSION_DOTS)
+    fw("project = 'Blender %s Python API'\n" % BLENDER_VERSION_STRING)
     fw("master_doc = 'index'\n")
     fw("copyright = u'Blender Foundation'\n")
-    fw("version = '%s'\n" % BLENDER_VERSION_HASH)
-    fw("release = '%s'\n" % BLENDER_VERSION_HASH)
+    fw("version = '%s'\n" % BLENDER_VERSION_DOTS)
+    fw("release = '%s'\n" % BLENDER_VERSION_DOTS)
 
     # Quiet file not in table-of-contents warnings.
     fw("exclude_patterns = [\n")
@@ -1749,6 +1750,7 @@ except ModuleNotFoundError:
 
     fw("if html_theme == 'sphinx_rtd_theme':\n")
     fw("    html_theme_options = {\n")
+    fw("        'display_version': False,\n")
     # fw("        'analytics_id': '',\n")
     # fw("        'collapse_navigation': True,\n")
     fw("        'sticky_navigation': False,\n")
@@ -1765,10 +1767,15 @@ except ModuleNotFoundError:
     fw("html_show_search_summary = True\n")
     fw("html_split_index = True\n")
     fw("html_static_path = ['static']\n")
+    fw("templates_path = ['templates']\n")
+    fw("html_context = {'commit': '%s'}\n" % BLENDER_VERSION_HASH)
     fw("html_extra_path = ['static/favicon.ico', 'static/blender_logo.svg']\n")
     fw("html_favicon = 'static/favicon.ico'\n")
     fw("html_logo = 'static/blender_logo.svg'\n")
     fw("html_last_updated_fmt = '%m/%d/%Y'\n\n")
+    fw("if html_theme == 'sphinx_rtd_theme':\n")
+    fw("    html_css_files = ['css/version_switch.css']\n")
+    fw("    html_js_files = ['js/version_switch.js']\n")
 
     # needed for latex, pdf gen
     fw("latex_elements = {\n")
@@ -2125,6 +2132,9 @@ def copy_theme_assets(basepath):
     shutil.copytree(os.path.join(SCRIPT_DIR, "static"),
                     os.path.join(basepath, "static"),
                     copy_function=shutil.copy)
+    shutil.copytree(os.path.join(SCRIPT_DIR, "templates"),
+                    os.path.join(basepath, "templates"),
+                    copy_function=shutil.copy)
 
 
 def rna2sphinx(basepath):

doc/python_api/static/css/version_switch.css

@@ -0,0 +1,127 @@
+/* Override RTD theme */
+.rst-versions {
+	border-top: 0px;
+	overflow: visible;
+}
+.version-btn.vdeact {
+	cursor: default;
+	color: dimgray;
+}
+
+.version-btn.vdeact::after {
+	content: "";
+}
+#versionwrap {
+	display: flex;
+	padding-top: 2px;
+	font-size: 90%;
+	justify-content: center;
+	flex-wrap: wrap;
+}
+.version-btn {
+	display: inline-block;
+	background-color: #272525;
+	width: 140px;
+	text-align: center;
+	padding: 3px 10px;
+	margin: 0px 5px 4px;
+	vertical-align: middle;
+	color: #27AE60;
+	border: solid 1px #444444;
+	border-radius: 3px;
+	cursor: pointer;
+	z-index: 400;
+	transition: border-color 0.4s;
+}
+.version-btn::after {
+	content:"\f0d8";
+	display: inline;
+	font: normal normal normal 16px/1 FontAwesome;
+	color: #8d8c8c;
+	vertical-align: top;
+	padding-left: 0.5em;
+}
+.version-btn-open::after {
+	color: gray;
+}
+.version-btn:hover, .version-btn:focus {
+	border-color: #525252;
+}
+.version-btn-open {
+	color: gray;
+	border: solid 1px gray;
+}
+.version-btn.wait {
+	cursor: wait;
+}
+.version-btn.disabled {
+	cursor: not-allowed;
+	color: dimgray;
+}
+.version-dialog {
+	display: none;
+	position: absolute;
+	bottom: 28px;
+	width: 140px;
+	margin: 0 5px;
+	padding-bottom: 4px;
+	background-color: #0003;
+	border-radius: 3px;
+	box-shadow: 0 0 6px #000C;
+	z-index: 999;
+	max-height: calc(100vh - 30px);
+	overflow-y: auto;
+	cursor: default;
+}
+.version-title {
+	padding: 5px;
+	color: black;
+	text-align: center;
+	font-size: 102%;
+	background-color: #27ae60;
+	border-bottom: solid 1.5px #444;
+}
+.version-list {
+	margin-bottom: 4px;
+	text-align: center;
+	background-color: #000C;
+	border: solid 1px gray;
+	border-radius: 0px 0px 3px 3px;
+}
+.version-list a, .version-list span, .version-list li {
+	position: relative;
+	display: block;
+	font-size: 98%;
+	line-height: 1.15;
+	width: 100%;
+	margin: 0;
+	padding: 4px 0px;
+	color: #404040;
+}
+.version-list li {
+	background-color: #ede9e9;
+	color: #404040;
+	padding: 1px;
+}
+.version-list li:hover, .version-list li a:focus {
+	background-color: #b9cfda;
+}
+.version-list li.selected, .version-list li.selected:hover {
+	background-color: #8d8c8c;
+}
+.version-list li.selected span {
+	cursor: default;
+	outline-color: red;
+}
+.version-arrow {
+	position: absolute;
+	width: 8px;
+	height: 8px;
+	left: 50%;
+	bottom: 4px;
+	margin-left: -4px;
+	transform: rotate(225deg);
+	background: #ede9e9;
+	border: 1px solid gray;
+	border-width: 1px 0 0 1px;
+}

doc/python_api/static/js/version_switch.js

@@ -0,0 +1,323 @@
+(function() {  // switch: v1.2
+"use strict";
+
+var versionsFileUrl = "https://docs.blender.org/PROD/versions.json"
+
+var all_versions;
+
+var Popover = function() {
+  function Popover(id)
+  {
+    this.isOpen = false;
+    this.type = (id === "version-popover");
+    this.$btn = $('#' + id);
+    this.$dialog = this.$btn.next();
+    this.$list = this.$dialog.children("ul");
+    this.sel = null;
+    this.beforeInit();
+  }
+
+  Popover.prototype = {
+    beforeInit : function() {
+      var that = this;
+      this.$btn.on("click", function(e) {
+        that.init();
+        e.preventDefault();
+        e.stopPropagation();
+      });
+      this.$btn.on("keydown", function(e) {
+        if (that.btnKeyFilter(e)) {
+          that.init();
+          e.preventDefault();
+          e.stopPropagation();
+        }
+      });
+    },
+    init : function() {
+      this.$btn.off("click");
+      this.$btn.off("keydown");
+
+      if (all_versions === undefined) {
+        this.$btn.addClass("wait");
+        this.loadVL(this);
+      }
+      else {
+        this.afterLoad();
+      }
+    },
+    loadVL : function(that) {
+      $.getJSON(versionsFileUrl, function(data) {
+         all_versions = data;
+         that.afterLoad();
+         return true;
+       }).fail(function() {
+        console.log("Version Switch Error: versions.json could not be loaded.");
+        that.$btn.addClass("disabled");
+        return false;
+      });
+    },
+    afterLoad : function() {
+      var release = DOCUMENTATION_OPTIONS.VERSION;
+      const m = release.match(/\d\.\d+/g);
+      if (m) {
+        release = m[0];
+      }
+
+      this.warnOld(release, all_versions);
+
+      var version = this.getNamed(release);
+      var list = this.buildList(version);
+
+      this.$list.children(":first-child").remove();
+      this.$list.append(list);
+      var that = this;
+      this.$list.on("keydown", function(e) {
+        that.keyMove(e);
+      });
+
+      this.$btn.removeClass("wait");
+      this.btnOpenHandler();
+      this.$btn.on("mousedown", function(e) {
+        that.btnOpenHandler();
+        e.preventDefault()
+      });
+      this.$btn.on("keydown", function(e) {
+        if (that.btnKeyFilter(e)) {
+          that.btnOpenHandler();
+        }
+      });
+    },
+    warnOld : function(release, all_versions) {
+      // Note this is effectively disabled now, two issues must fixed:
+      // * versions.js does not contain a current entry, because that leads to
+      //   duplicate version numbers in the menu. These need to be deduplicated.
+      // * It only shows the warning after opening the menu to switch version
+      //   when versions.js is loaded. This is too late to be useful.
+      var current = all_versions.current
+      if (!current)
+      {
+        // console.log("Version Switch Error: no 'current' in version.json.");
+        return;
+      }
+      const m = current.match(/\d\.\d+/g);
+      if (m) {
+        current = parseFloat(m[0]);
+      }
+      if (release < current) {
+        var currentURL = window.location.pathname.replace(release, current);
+        var warning = $('<div class="admonition warning"> ' +
+                        '<p class="first admonition-title">Note</p> ' +
+                        '<p class="last"> ' +
+                        'You are not using the most up to date version of the documentation. ' +
+                        '<a href="#"></a> is the newest version.' +
+                        '</p>' +
+                        '</div>');
+
+        warning.find('a').attr('href', currentURL).text(current);
+
+        var body = $("div.body");
+        if (!body.length) {
+          body = $("div.document");
+        }
+        body.prepend(warning);
+      }
+    },
+    buildList : function(v) {
+      var url = new URL(window.location.href);
+      let pathSplit = [ "", "api", v ];
+      if (url.pathname.startsWith("/api/")) {
+        pathSplit.push(url.pathname.split('/').slice(3).join('/'));
+      }
+      else {
+        pathSplit.push(url.pathname.substring(1));
+      }
+      if (this.type) {
+        var dyn = all_versions;
+        var cur = v;
+      }
+      var buf = [];
+      var that = this;
+      $.each(dyn, function(ix, title) {
+        buf.push("<li");
+        if (ix === cur) {
+          buf.push(
+              ' class="selected" tabindex="-1" role="presentation"><span tabindex="-1" role="menuitem" aria-current="page">' +
+              title + '</spanp></li>');
+        }
+        else {
+          pathSplit[2 + that.type] = ix;
+          var href = new URL(url);
+          href.pathname = pathSplit.join('/');
+          buf.push(' tabindex="-1" role="presentation"><a href ="' + href + '" tabindex="-1">' +
+                   title + '</a></li>');
+        }
+      });
+      return buf.join('');
+    },
+    getNamed : function(v) {
+      $.each(all_versions, function(ix, title) {
+        if (ix === "master" || ix === "latest") {
+          var m = title.match(/\d\.\d[\w\d\.]*/)[0];
+          if (parseFloat(m) == v) {
+            v = ix;
+            return false;
+          }
+        }
+      });
+      return v;
+    },
+    dialogToggle : function(speed) {
+      var wasClose = !this.isOpen;
+      var that = this;
+      if (!this.isOpen) {
+        this.$btn.addClass("version-btn-open");
+        this.$btn.attr("aria-pressed", true);
+        this.$dialog.attr("aria-hidden", false);
+        this.$dialog.fadeIn(speed, function() {
+          that.$btn.parent().on("focusout", function(e) {
+            that.focusoutHandler();
+            e.stopImmediatePropagation();
+          })
+          that.$btn.parent().on("mouseleave", function(e) {
+            that.mouseoutHandler();
+            e.stopImmediatePropagation();
+          });
+        });
+        this.isOpen = true;
+      }
+      else {
+        this.$btn.removeClass("version-btn-open");
+        this.$btn.attr("aria-pressed", false);
+        this.$dialog.attr("aria-hidden", true);
+        this.$btn.parent().off("focusout");
+        this.$btn.parent().off("mouseleave");
+        this.$dialog.fadeOut(speed, function() {
+          if (this.$sel) {
+            this.$sel.attr("tabindex", -1);
+          }
+          that.$btn.attr("tabindex", 0);
+          if (document.activeElement !== null && document.activeElement !== document &&
+              document.activeElement !== document.body) {
+            that.$btn.focus();
+          }
+        });
+        this.isOpen = false;
+      }
+
+      if (wasClose) {
+        if (this.$sel) {
+          this.$sel.attr("tabindex", -1);
+        }
+        if (document.activeElement !== null && document.activeElement !== document &&
+            document.activeElement !== document.body) {
+          var $nw = this.listEnter();
+          $nw.attr("tabindex", 0);
+          $nw.focus();
+          this.$sel = $nw;
+        }
+      }
+    },
+    btnOpenHandler : function() {
+      this.dialogToggle(300);
+    },
+    focusoutHandler : function() {
+      var list = this.$list;
+      var that = this;
+      setTimeout(function() {
+        if (list.find(":focus").length === 0) {
+          that.dialogToggle(200);
+        }
+      }, 200);
+    },
+    mouseoutHandler : function() {
+      this.dialogToggle(200);
+    },
+    btnKeyFilter : function(e) {
+      if (e.ctrlKey || e.shiftKey) {
+        return false;
+      }
+      if (e.key === " " || e.key === "Enter" || (e.key === "ArrowDown" && e.altKey) ||
+          e.key === "ArrowDown" || e.key === "ArrowUp") {
+        return true;
+      }
+      return false;
+    },
+    keyMove : function(e) {
+      if (e.ctrlKey || e.shiftKey) {
+        return true;
+      }
+      var p = true;
+      var $nw = $(e.target);
+      switch (e.key) {
+        case "ArrowUp":
+          $nw = this.listPrev($nw);
+          break;
+        case "ArrowDown":
+          $nw = this.listNext($nw);
+          break;
+        case "Home":
+          $nw = this.listFirst();
+          break;
+        case "End":
+          $nw = this.listLast();
+          break;
+        case "Escape":
+          $nw = this.listExit();
+          break;
+        case "ArrowLeft":
+          $nw = this.listExit();
+          break;
+        case "ArrowRight":
+          $nw = this.listExit();
+          break;
+        default:
+          p = false;
+      }
+      if (p) {
+        $nw.attr("tabindex", 0);
+        $nw.focus();
+        if (this.$sel) {
+          this.$sel.attr("tabindex", -1);
+        }
+        this.$sel = $nw;
+        e.preventDefault();
+        e.stopPropagation();
+      }
+    },
+    listPrev : function($nw) {
+      if ($nw.parent().prev().length !== 0) {
+        return $nw.parent().prev().children(":first-child");
+      }
+      else {
+        return this.listLast();
+      }
+    },
+    listNext : function($nw) {
+      if ($nw.parent().next().length !== 0) {
+        return $nw.parent().next().children(":first-child");
+      }
+      else {
+        return this.listFirst();
+      }
+    },
+    listFirst : function() {
+      return this.$list.children(":first-child").children(":first-child");
+    },
+    listLast : function() {
+      return this.$list.children(":last-child").children(":first-child");
+    },
+    listExit : function() {
+      this.mouseoutHandler();
+      return this.$btn;
+    },
+    listEnter : function() {
+      return this.$list.children(":first-child").children(":first-child");
+    }
+  };
+  return Popover
+}();
+
+$(document).ready(function() {
+  var lng_popover = new Popover("version-popover");
+});
+})();

doc/python_api/templates/versions.html

@@ -0,0 +1,17 @@
+<div class="rst-versions" data-toggle="rst-versions" role="note" aria-label="document versions">
+    <ul id="versionwrap" role="presentation">
+       <li role="presentation">
+            <span id="version-popover" class="version-btn" tabindex="0" role="button" aria-label="versions selector" aria-haspopup="true" aria-controls="version-vsnlist" aria-disabled="true">
+            {{ release }} 
+            </span>
+            <div class="version-dialog" aria-hidden="true">
+                <div class="version-arrow" aria-hidden="true"></div>
+                <div class="version-title">Versions</div>
+                <ul id="version-vsnlist" class="version-list" role="menu" aria-labelledby="version-popover" aria-hidden="true">
+                    <li role="presentation">Loading...</li>
+                </ul>
+            </div>
+        </li>
+    </ul>
+</div>
+    

extern/CMakeLists.txt

@@ -113,6 +113,6 @@ if(WITH_MOD_FLUID)
   add_subdirectory(mantaflow)
 endif()
 
-if (WITH_COMPOSITOR)
+if(WITH_COMPOSITOR)
   add_subdirectory(smaa_areatex)
 endif()

extern/audaspace/plugins/ffmpeg/FFMPEGReader.cpp

@@ -177,7 +177,7 @@ void FFMPEGReader::init(int stream)
 
 	// get a decoder and open it
 #ifndef FFMPEG_OLD_CODE
-	AVCodec* aCodec = avcodec_find_decoder(m_formatCtx->streams[m_stream]->codecpar->codec_id);
+	const AVCodec* aCodec = avcodec_find_decoder(m_formatCtx->streams[m_stream]->codecpar->codec_id);
 
 	if(!aCodec)
 		AUD_THROW(FileException, "File couldn't be read, no decoder found with ffmpeg.");

extern/audaspace/plugins/ffmpeg/FFMPEGWriter.cpp

@@ -23,6 +23,7 @@
 extern "C" {
 #include <libavcodec/avcodec.h>
 #include <libavformat/avio.h>
+#include <libavutil/channel_layout.h>
 }
 
 AUD_NAMESPACE_BEGIN
@@ -171,66 +172,66 @@ FFMPEGWriter::FFMPEGWriter(std::string filename, DeviceSpecs specs, Container fo
 	if(avformat_alloc_output_context2(&m_formatCtx, nullptr, formats[format], filename.c_str()) < 0)
 		AUD_THROW(FileException, "File couldn't be written, format couldn't be found with ffmpeg.");
 
-	AVOutputFormat* outputFmt = m_formatCtx->oformat;
+	const AVOutputFormat* outputFmt = m_formatCtx->oformat;
 
 	if(!outputFmt) {
 		avformat_free_context(m_formatCtx);
 		AUD_THROW(FileException, "File couldn't be written, output format couldn't be found with ffmpeg.");
 	}
 
-	outputFmt->audio_codec = AV_CODEC_ID_NONE;
+	AVCodecID audio_codec = AV_CODEC_ID_NONE;
 
 	switch(codec)
 	{
 	case CODEC_AAC:
-		outputFmt->audio_codec = AV_CODEC_ID_AAC;
+		audio_codec = AV_CODEC_ID_AAC;
 		break;
 	case CODEC_AC3:
-		outputFmt->audio_codec = AV_CODEC_ID_AC3;
+		audio_codec = AV_CODEC_ID_AC3;
 		break;
 	case CODEC_FLAC:
-		outputFmt->audio_codec = AV_CODEC_ID_FLAC;
+		audio_codec = AV_CODEC_ID_FLAC;
 		break;
 	case CODEC_MP2:
-		outputFmt->audio_codec = AV_CODEC_ID_MP2;
+		audio_codec = AV_CODEC_ID_MP2;
 		break;
 	case CODEC_MP3:
-		outputFmt->audio_codec = AV_CODEC_ID_MP3;
+		audio_codec = AV_CODEC_ID_MP3;
 		break;
 	case CODEC_OPUS:
-		outputFmt->audio_codec = AV_CODEC_ID_OPUS;
+		audio_codec = AV_CODEC_ID_OPUS;
 		break;
 	case CODEC_PCM:
 		switch(specs.format)
 		{
 		case FORMAT_U8:
-			outputFmt->audio_codec = AV_CODEC_ID_PCM_U8;
+			audio_codec = AV_CODEC_ID_PCM_U8;
 			break;
 		case FORMAT_S16:
-			outputFmt->audio_codec = AV_CODEC_ID_PCM_S16LE;
+			audio_codec = AV_CODEC_ID_PCM_S16LE;
 			break;
 		case FORMAT_S24:
-			outputFmt->audio_codec = AV_CODEC_ID_PCM_S24LE;
+			audio_codec = AV_CODEC_ID_PCM_S24LE;
 			break;
 		case FORMAT_S32:
-			outputFmt->audio_codec = AV_CODEC_ID_PCM_S32LE;
+			audio_codec = AV_CODEC_ID_PCM_S32LE;
 			break;
 		case FORMAT_FLOAT32:
-			outputFmt->audio_codec = AV_CODEC_ID_PCM_F32LE;
+			audio_codec = AV_CODEC_ID_PCM_F32LE;
 			break;
 		case FORMAT_FLOAT64:
-			outputFmt->audio_codec = AV_CODEC_ID_PCM_F64LE;
+			audio_codec = AV_CODEC_ID_PCM_F64LE;
 			break;
 		default:
-			outputFmt->audio_codec = AV_CODEC_ID_NONE;
+			audio_codec = AV_CODEC_ID_NONE;
 			break;
 		}
 		break;
 	case CODEC_VORBIS:
-		outputFmt->audio_codec = AV_CODEC_ID_VORBIS;
+		audio_codec = AV_CODEC_ID_VORBIS;
 		break;
 	default:
-		outputFmt->audio_codec = AV_CODEC_ID_NONE;
+		audio_codec = AV_CODEC_ID_NONE;
 		break;
 	}
 
@@ -268,10 +269,10 @@ FFMPEGWriter::FFMPEGWriter(std::string filename, DeviceSpecs specs, Container fo
 
 	try
 	{
-		if(outputFmt->audio_codec == AV_CODEC_ID_NONE)
+		if(audio_codec == AV_CODEC_ID_NONE)
 			AUD_THROW(FileException, "File couldn't be written, audio codec not found with ffmpeg.");
 
-		AVCodec* codec = avcodec_find_encoder(outputFmt->audio_codec);
+		const AVCodec* codec = avcodec_find_encoder(audio_codec);
 		if(!codec)
 			AUD_THROW(FileException, "File couldn't be written, audio encoder couldn't be found with ffmpeg.");
 

intern/cycles/app/CMakeLists.txt

@@ -44,15 +44,13 @@ else()
 endif()
 
 if(WITH_CYCLES_STANDALONE AND WITH_CYCLES_STANDALONE_GUI)
-  list(APPEND LIBRARIES ${GLUT_LIBRARIES})
+  add_definitions(${GL_DEFINITIONS})
+  list(APPEND INC_SYS ${GLEW_INCLUDE_DIR} ${SDL2_INCLUDE_DIRS})
+  list(APPEND LIBRARIES ${CYCLES_GL_LIBRARIES} ${SDL2_LIBRARIES})
 endif()
 
-list(APPEND LIBRARIES ${CYCLES_GL_LIBRARIES})
-
 # Common configuration.
 
-add_definitions(${GL_DEFINITIONS})
-
 include_directories(${INC})
 include_directories(SYSTEM ${INC_SYS})
 
@@ -66,6 +64,18 @@ if(WITH_CYCLES_STANDALONE)
     oiio_output_driver.cpp
     oiio_output_driver.h
   )
+
+  if(WITH_CYCLES_STANDALONE_GUI)
+    list(APPEND SRC
+      opengl/display_driver.cpp
+      opengl/display_driver.h
+      opengl/shader.cpp
+      opengl/shader.h
+      opengl/window.cpp
+      opengl/window.h
+    )
+  endif()
+
   add_executable(cycles ${SRC} ${INC} ${INC_SYS})
   unset(SRC)
 
@@ -80,6 +90,10 @@ if(WITH_CYCLES_STANDALONE)
       # OpenImageDenoise uses BNNS from the Accelerate framework.
       set_property(TARGET cycles APPEND_STRING PROPERTY LINK_FLAGS " -framework Accelerate")
     endif()
+    if(WITH_CYCLES_STANDALONE_GUI)
+      set_property(TARGET cycles APPEND_STRING PROPERTY LINK_FLAGS
+        " -framework Cocoa -framework CoreAudio -framework AudioUnit -framework AudioToolbox -framework ForceFeedback -framework CoreVideo")
+    endif()
   endif()
 
   if(UNIX AND NOT APPLE)

intern/cycles/app/cycles_standalone.cpp

@@ -40,11 +40,10 @@
 #include "app/oiio_output_driver.h"
 
 #ifdef WITH_CYCLES_STANDALONE_GUI
-#  include "util/view.h"
+#  include "opengl/display_driver.h"
+#  include "opengl/window.h"
 #endif
 
-#include "app/cycles_xml.h"
-
 CCL_NAMESPACE_BEGIN
 
 struct Options {
@@ -130,7 +129,14 @@ static void session_init()
   options.output_pass = "combined";
   options.session = new Session(options.session_params, options.scene_params);
 
-  if (!options.output_filepath.empty()) {
+#ifdef WITH_CYCLES_STANDALONE_GUI
+  if (!options.session_params.background) {
+    options.session->set_display_driver(make_unique<OpenGLDisplayDriver>(
+        window_opengl_context_enable, window_opengl_context_disable));
+  }
+  else
+#endif
+      if (!options.output_filepath.empty()) {
     options.session->set_output_driver(make_unique<OIIOOutputDriver>(
         options.output_filepath, options.output_pass, session_print));
   }
@@ -139,7 +145,7 @@ static void session_init()
     options.session->progress.set_update_callback(function_bind(&session_print_status));
 #ifdef WITH_CYCLES_STANDALONE_GUI
   else
-    options.session->progress.set_update_callback(function_bind(&view_redraw));
+    options.session->progress.set_update_callback(function_bind(&window_redraw));
 #endif
 
   /* load scene */
@@ -204,10 +210,10 @@ static void display_info(Progress &progress)
       sample_time,
       interactive.c_str());
 
-  view_display_info(str.c_str());
+  window_display_info(str.c_str());
 
   if (options.show_help)
-    view_display_help();
+    window_display_help();
 }
 
 static void display()
@@ -538,15 +544,15 @@ int main(int argc, const char **argv)
     string title = "Cycles: " + path_filename(options.filepath);
 
     /* init/exit are callback so they run while GL is initialized */
-    view_main_loop(title.c_str(),
-                   options.width,
-                   options.height,
-                   session_init,
-                   session_exit,
-                   resize,
-                   display,
-                   keyboard,
-                   motion);
+    window_main_loop(title.c_str(),
+                     options.width,
+                     options.height,
+                     session_init,
+                     session_exit,
+                     resize,
+                     display,
+                     keyboard,
+                     motion);
   }
 #endif
 

intern/cycles/app/opengl/display_driver.cpp

@@ -0,0 +1,398 @@
+/*
+ * 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 "app/opengl/display_driver.h"
+#include "app/opengl/shader.h"
+
+#include "util/log.h"
+#include "util/string.h"
+
+#include <GL/glew.h>
+#include <SDL.h>
+
+CCL_NAMESPACE_BEGIN
+
+/* --------------------------------------------------------------------
+ * OpenGLDisplayDriver.
+ */
+
+OpenGLDisplayDriver::OpenGLDisplayDriver(const function<bool()> &gl_context_enable,
+                                         const function<void()> &gl_context_disable)
+    : gl_context_enable_(gl_context_enable), gl_context_disable_(gl_context_disable)
+{
+}
+
+OpenGLDisplayDriver::~OpenGLDisplayDriver()
+{
+}
+
+/* --------------------------------------------------------------------
+ * Update procedure.
+ */
+
+void OpenGLDisplayDriver::next_tile_begin()
+{
+  /* Assuming no tiles used in interactive display. */
+}
+
+bool OpenGLDisplayDriver::update_begin(const Params &params, int texture_width, int texture_height)
+{
+  /* Note that it's the responsibility of OpenGLDisplayDriver to ensure updating and drawing
+   * the texture does not happen at the same time. This is achieved indirectly.
+   *
+   * When enabling the OpenGL context, it uses an internal mutex lock DST.gl_context_lock.
+   * This same lock is also held when do_draw() is called, which together ensure mutual
+   * exclusion.
+   *
+   * This locking is not performed on the Cycles side, because that would cause lock inversion. */
+  if (!gl_context_enable_()) {
+    return false;
+  }
+
+  if (gl_render_sync_) {
+    glWaitSync((GLsync)gl_render_sync_, 0, GL_TIMEOUT_IGNORED);
+  }
+
+  if (!gl_texture_resources_ensure()) {
+    gl_context_disable_();
+    return false;
+  }
+
+  /* Update texture dimensions if needed. */
+  if (texture_.width != texture_width || texture_.height != texture_height) {
+    glActiveTexture(GL_TEXTURE0);
+    glBindTexture(GL_TEXTURE_2D, texture_.gl_id);
+    glTexImage2D(
+        GL_TEXTURE_2D, 0, GL_RGBA16F, texture_width, texture_height, 0, GL_RGBA, GL_HALF_FLOAT, 0);
+    texture_.width = texture_width;
+    texture_.height = texture_height;
+    glBindTexture(GL_TEXTURE_2D, 0);
+
+    /* Texture did change, and no pixel storage was provided. Tag for an explicit zeroing out to
+     * avoid undefined content. */
+    texture_.need_clear = true;
+  }
+
+  /* Update PBO dimensions if needed.
+   *
+   * NOTE: Allocate the PBO for the the size which will fit the final render resolution (as in,
+   * at a resolution divider 1. This was we don't need to recreate graphics interoperability
+   * objects which are costly and which are tied to the specific underlying buffer size.
+   * The downside of this approach is that when graphics interoperability is not used we are
+   * sending too much data to GPU when resolution divider is not 1. */
+  const int buffer_width = params.full_size.x;
+  const int buffer_height = params.full_size.y;
+  if (texture_.buffer_width != buffer_width || texture_.buffer_height != buffer_height) {
+    const size_t size_in_bytes = sizeof(half4) * buffer_width * buffer_height;
+    glBindBuffer(GL_PIXEL_UNPACK_BUFFER, texture_.gl_pbo_id);
+    glBufferData(GL_PIXEL_UNPACK_BUFFER, size_in_bytes, 0, GL_DYNAMIC_DRAW);
+    glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
+
+    texture_.buffer_width = buffer_width;
+    texture_.buffer_height = buffer_height;
+  }
+
+  /* New content will be provided to the texture in one way or another, so mark this in a
+   * centralized place. */
+  texture_.need_update = true;
+
+  return true;
+}
+
+void OpenGLDisplayDriver::update_end()
+{
+  gl_upload_sync_ = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0);
+  glFlush();
+
+  gl_context_disable_();
+}
+
+/* --------------------------------------------------------------------
+ * Texture buffer mapping.
+ */
+
+half4 *OpenGLDisplayDriver::map_texture_buffer()
+{
+  glBindBuffer(GL_PIXEL_UNPACK_BUFFER, texture_.gl_pbo_id);
+
+  half4 *mapped_rgba_pixels = reinterpret_cast<half4 *>(
+      glMapBuffer(GL_PIXEL_UNPACK_BUFFER, GL_WRITE_ONLY));
+  if (!mapped_rgba_pixels) {
+    LOG(ERROR) << "Error mapping OpenGLDisplayDriver pixel buffer object.";
+  }
+
+  if (texture_.need_clear) {
+    const int64_t texture_width = texture_.width;
+    const int64_t texture_height = texture_.height;
+    memset(reinterpret_cast<void *>(mapped_rgba_pixels),
+           0,
+           texture_width * texture_height * sizeof(half4));
+    texture_.need_clear = false;
+  }
+
+  return mapped_rgba_pixels;
+}
+
+void OpenGLDisplayDriver::unmap_texture_buffer()
+{
+  glUnmapBuffer(GL_PIXEL_UNPACK_BUFFER);
+
+  glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
+}
+
+/* --------------------------------------------------------------------
+ * Graphics interoperability.
+ */
+
+OpenGLDisplayDriver::GraphicsInterop OpenGLDisplayDriver::graphics_interop_get()
+{
+  GraphicsInterop interop_dst;
+
+  interop_dst.buffer_width = texture_.buffer_width;
+  interop_dst.buffer_height = texture_.buffer_height;
+  interop_dst.opengl_pbo_id = texture_.gl_pbo_id;
+
+  interop_dst.need_clear = texture_.need_clear;
+  texture_.need_clear = false;
+
+  return interop_dst;
+}
+
+void OpenGLDisplayDriver::graphics_interop_activate()
+{
+  gl_context_enable_();
+}
+
+void OpenGLDisplayDriver::graphics_interop_deactivate()
+{
+  gl_context_disable_();
+}
+
+/* --------------------------------------------------------------------
+ * Drawing.
+ */
+
+void OpenGLDisplayDriver::clear()
+{
+  texture_.need_clear = true;
+}
+
+void OpenGLDisplayDriver::draw(const Params &params)
+{
+  /* See do_update_begin() for why no locking is required here. */
+  if (texture_.need_clear) {
+    /* Texture is requested to be cleared and was not yet cleared.
+     * Do early return which should be equivalent of drawing all-zero texture. */
+    return;
+  }
+
+  if (!gl_draw_resources_ensure()) {
+    return;
+  }
+
+  if (gl_upload_sync_) {
+    glWaitSync((GLsync)gl_upload_sync_, 0, GL_TIMEOUT_IGNORED);
+  }
+
+  glEnable(GL_BLEND);
+  glBlendFunc(GL_ONE, GL_ONE_MINUS_SRC_ALPHA);
+
+  display_shader_.bind(params.full_size.x, params.full_size.y);
+
+  glActiveTexture(GL_TEXTURE0);
+  glBindTexture(GL_TEXTURE_2D, texture_.gl_id);
+
+  if (texture_.width != params.size.x || texture_.height != params.size.y) {
+    /* Resolution divider is different from 1, force nearest interpolation. */
+    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
+  }
+  else {
+    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
+  }
+
+  glBindBuffer(GL_ARRAY_BUFFER, vertex_buffer_);
+
+  texture_update_if_needed();
+  vertex_buffer_update(params);
+
+  GLuint vertex_array_object;
+  glGenVertexArrays(1, &vertex_array_object);
+  glBindVertexArray(vertex_array_object);
+
+  const int texcoord_attribute = display_shader_.get_tex_coord_attrib_location();
+  const int position_attribute = display_shader_.get_position_attrib_location();
+
+  glEnableVertexAttribArray(texcoord_attribute);
+  glEnableVertexAttribArray(position_attribute);
+
+  glVertexAttribPointer(
+      texcoord_attribute, 2, GL_FLOAT, GL_FALSE, 4 * sizeof(float), (const GLvoid *)0);
+  glVertexAttribPointer(position_attribute,
+                        2,
+                        GL_FLOAT,
+                        GL_FALSE,
+                        4 * sizeof(float),
+                        (const GLvoid *)(sizeof(float) * 2));
+
+  glDrawArrays(GL_TRIANGLE_FAN, 0, 4);
+
+  glBindBuffer(GL_ARRAY_BUFFER, 0);
+  glBindTexture(GL_TEXTURE_2D, 0);
+
+  glDeleteVertexArrays(1, &vertex_array_object);
+
+  display_shader_.unbind();
+
+  glDisable(GL_BLEND);
+
+  gl_render_sync_ = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0);
+  glFlush();
+}
+
+bool OpenGLDisplayDriver::gl_draw_resources_ensure()
+{
+  if (!texture_.gl_id) {
+    /* If there is no texture allocated, there is nothing to draw. Inform the draw call that it can
+     * can not continue. Note that this is not an unrecoverable error, so once the texture is known
+     * we will come back here and create all the GPU resources needed for draw. */
+    return false;
+  }
+
+  if (gl_draw_resource_creation_attempted_) {
+    return gl_draw_resources_created_;
+  }
+  gl_draw_resource_creation_attempted_ = true;
+
+  if (!vertex_buffer_) {
+    glGenBuffers(1, &vertex_buffer_);
+    if (!vertex_buffer_) {
+      LOG(ERROR) << "Error creating vertex buffer.";
+      return false;
+    }
+  }
+
+  gl_draw_resources_created_ = true;
+
+  return true;
+}
+
+void OpenGLDisplayDriver::gl_resources_destroy()
+{
+  gl_context_enable_();
+
+  if (vertex_buffer_ != 0) {
+    glDeleteBuffers(1, &vertex_buffer_);
+  }
+
+  if (texture_.gl_pbo_id) {
+    glDeleteBuffers(1, &texture_.gl_pbo_id);
+    texture_.gl_pbo_id = 0;
+  }
+
+  if (texture_.gl_id) {
+    glDeleteTextures(1, &texture_.gl_id);
+    texture_.gl_id = 0;
+  }
+
+  gl_context_disable_();
+}
+
+bool OpenGLDisplayDriver::gl_texture_resources_ensure()
+{
+  if (texture_.creation_attempted) {
+    return texture_.is_created;
+  }
+  texture_.creation_attempted = true;
+
+  DCHECK(!texture_.gl_id);
+  DCHECK(!texture_.gl_pbo_id);
+
+  /* Create texture. */
+  glGenTextures(1, &texture_.gl_id);
+  if (!texture_.gl_id) {
+    LOG(ERROR) << "Error creating texture.";
+    return false;
+  }
+
+  /* Configure the texture. */
+  glActiveTexture(GL_TEXTURE0);
+  glBindTexture(GL_TEXTURE_2D, texture_.gl_id);
+  glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
+  glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
+  glBindTexture(GL_TEXTURE_2D, 0);
+
+  /* Create PBO for the texture. */
+  glGenBuffers(1, &texture_.gl_pbo_id);
+  if (!texture_.gl_pbo_id) {
+    LOG(ERROR) << "Error creating texture pixel buffer object.";
+    return false;
+  }
+
+  /* Creation finished with a success. */
+  texture_.is_created = true;
+
+  return true;
+}
+
+void OpenGLDisplayDriver::texture_update_if_needed()
+{
+  if (!texture_.need_update) {
+    return;
+  }
+
+  glBindBuffer(GL_PIXEL_UNPACK_BUFFER, texture_.gl_pbo_id);
+  glTexSubImage2D(
+      GL_TEXTURE_2D, 0, 0, 0, texture_.width, texture_.height, GL_RGBA, GL_HALF_FLOAT, 0);
+  glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
+
+  texture_.need_update = false;
+}
+
+void OpenGLDisplayDriver::vertex_buffer_update(const Params &params)
+{
+  /* Invalidate old contents - avoids stalling if the buffer is still waiting in queue to be
+   * rendered. */
+  glBufferData(GL_ARRAY_BUFFER, 16 * sizeof(float), NULL, GL_STREAM_DRAW);
+
+  float *vpointer = reinterpret_cast<float *>(glMapBuffer(GL_ARRAY_BUFFER, GL_WRITE_ONLY));
+  if (!vpointer) {
+    return;
+  }
+
+  vpointer[0] = 0.0f;
+  vpointer[1] = 0.0f;
+  vpointer[2] = params.full_offset.x;
+  vpointer[3] = params.full_offset.y;
+
+  vpointer[4] = 1.0f;
+  vpointer[5] = 0.0f;
+  vpointer[6] = (float)params.size.x + params.full_offset.x;
+  vpointer[7] = params.full_offset.y;
+
+  vpointer[8] = 1.0f;
+  vpointer[9] = 1.0f;
+  vpointer[10] = (float)params.size.x + params.full_offset.x;
+  vpointer[11] = (float)params.size.y + params.full_offset.y;
+
+  vpointer[12] = 0.0f;
+  vpointer[13] = 1.0f;
+  vpointer[14] = params.full_offset.x;
+  vpointer[15] = (float)params.size.y + params.full_offset.y;
+
+  glUnmapBuffer(GL_ARRAY_BUFFER);
+}
+
+CCL_NAMESPACE_END

intern/cycles/app/opengl/display_driver.h

@@ -0,0 +1,130 @@
+/*
+ * 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.
+ */
+
+#pragma once
+
+#include <atomic>
+
+#include "app/opengl/shader.h"
+
+#include "session/display_driver.h"
+
+#include "util/function.h"
+#include "util/unique_ptr.h"
+
+CCL_NAMESPACE_BEGIN
+
+class OpenGLDisplayDriver : public DisplayDriver {
+ public:
+  /* Callbacks for enabling and disabling the OpenGL context. Must be provided to support enabling
+   * the context on the Cycles render thread independent of the main thread. */
+  OpenGLDisplayDriver(const function<bool()> &gl_context_enable,
+                      const function<void()> &gl_context_disable);
+  ~OpenGLDisplayDriver();
+
+  virtual void graphics_interop_activate() override;
+  virtual void graphics_interop_deactivate() override;
+
+  virtual void clear() override;
+
+  void set_zoom(float zoom_x, float zoom_y);
+
+ protected:
+  virtual void next_tile_begin() override;
+
+  virtual bool update_begin(const Params &params, int texture_width, int texture_height) override;
+  virtual void update_end() override;
+
+  virtual half4 *map_texture_buffer() override;
+  virtual void unmap_texture_buffer() override;
+
+  virtual GraphicsInterop graphics_interop_get() override;
+
+  virtual void draw(const Params &params) override;
+
+  /* Make sure texture is allocated and its initial configuration is performed. */
+  bool gl_texture_resources_ensure();
+
+  /* Ensure all runtime GPU resources needed for drawing are allocated.
+   * Returns true if all resources needed for drawing are available. */
+  bool gl_draw_resources_ensure();
+
+  /* Destroy all GPU resources which are being used by this object. */
+  void gl_resources_destroy();
+
+  /* Update GPU texture dimensions and content if needed (new pixel data was provided).
+   *
+   * NOTE: The texture needs to be bound. */
+  void texture_update_if_needed();
+
+  /* Update vertex buffer with new coordinates of vertex positions and texture coordinates.
+   * This buffer is used to render texture in the viewport.
+   *
+   * NOTE: The buffer needs to be bound. */
+  void vertex_buffer_update(const Params &params);
+
+  /* Texture which contains pixels of the render result. */
+  struct {
+    /* Indicates whether texture creation was attempted and succeeded.
+     * Used to avoid multiple attempts of texture creation on GPU issues or GPU context
+     * misconfiguration. */
+    bool creation_attempted = false;
+    bool is_created = false;
+
+    /* OpenGL resource IDs of the texture itself and Pixel Buffer Object (PBO) used to write
+     * pixels to it.
+     *
+     * NOTE: Allocated on the engine's context. */
+    uint gl_id = 0;
+    uint gl_pbo_id = 0;
+
+    /* Is true when new data was written to the PBO, meaning, the texture might need to be resized
+     * and new data is to be uploaded to the GPU. */
+    bool need_update = false;
+
+    /* Content of the texture is to be filled with zeroes. */
+    std::atomic<bool> need_clear = true;
+
+    /* Dimensions of the texture in pixels. */
+    int width = 0;
+    int height = 0;
+
+    /* Dimensions of the underlying PBO. */
+    int buffer_width = 0;
+    int buffer_height = 0;
+  } texture_;
+
+  OpenGLShader display_shader_;
+
+  /* Special track of whether GPU resources were attempted to be created, to avoid attempts of
+   * their re-creation on failure on every redraw. */
+  bool gl_draw_resource_creation_attempted_ = false;
+  bool gl_draw_resources_created_ = false;
+
+  /* Vertex buffer which hold vertices of a triangle fan which is textures with the texture
+   * holding the render result. */
+  uint vertex_buffer_ = 0;
+
+  void *gl_render_sync_ = nullptr;
+  void *gl_upload_sync_ = nullptr;
+
+  float2 zoom_ = make_float2(1.0f, 1.0f);
+
+  function<bool()> gl_context_enable_ = nullptr;
+  function<void()> gl_context_disable_ = nullptr;
+};
+
+CCL_NAMESPACE_END

intern/cycles/app/opengl/shader.cpp

@@ -0,0 +1,210 @@
+/*
+ * 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 "app/opengl/shader.h"
+
+#include "util/log.h"
+#include "util/string.h"
+
+#include <GL/glew.h>
+
+CCL_NAMESPACE_BEGIN
+
+/* --------------------------------------------------------------------
+ * OpenGLShader.
+ */
+
+static const char *VERTEX_SHADER =
+    "#version 330\n"
+    "uniform vec2 fullscreen;\n"
+    "in vec2 texCoord;\n"
+    "in vec2 pos;\n"
+    "out vec2 texCoord_interp;\n"
+    "\n"
+    "vec2 normalize_coordinates()\n"
+    "{\n"
+    "   return (vec2(2.0) * (pos / fullscreen)) - vec2(1.0);\n"
+    "}\n"
+    "\n"
+    "void main()\n"
+    "{\n"
+    "   gl_Position = vec4(normalize_coordinates(), 0.0, 1.0);\n"
+    "   texCoord_interp = texCoord;\n"
+    "}\n\0";
+
+static const char *FRAGMENT_SHADER =
+    "#version 330\n"
+    "uniform sampler2D image_texture;\n"
+    "in vec2 texCoord_interp;\n"
+    "out vec4 fragColor;\n"
+    "\n"
+    "void main()\n"
+    "{\n"
+    "   vec4 rgba = texture(image_texture, texCoord_interp);\n"
+    /* Harcoded Rec.709 gamma, should use OpenColorIO eventually. */
+    "   fragColor = pow(rgba, vec4(0.45, 0.45, 0.45, 1.0));\n"
+    "}\n\0";
+
+static void shader_print_errors(const char *task, const char *log, const char *code)
+{
+  LOG(ERROR) << "Shader: " << task << " error:";
+  LOG(ERROR) << "===== shader string ====";
+
+  stringstream stream(code);
+  string partial;
+
+  int line = 1;
+  while (getline(stream, partial, '\n')) {
+    if (line < 10) {
+      LOG(ERROR) << " " << line << " " << partial;
+    }
+    else {
+      LOG(ERROR) << line << " " << partial;
+    }
+    line++;
+  }
+  LOG(ERROR) << log;
+}
+
+static int compile_shader_program(void)
+{
+  const struct Shader {
+    const char *source;
+    const GLenum type;
+  } shaders[2] = {{VERTEX_SHADER, GL_VERTEX_SHADER}, {FRAGMENT_SHADER, GL_FRAGMENT_SHADER}};
+
+  const GLuint program = glCreateProgram();
+
+  for (int i = 0; i < 2; i++) {
+    const GLuint shader = glCreateShader(shaders[i].type);
+
+    string source_str = shaders[i].source;
+    const char *c_str = source_str.c_str();
+
+    glShaderSource(shader, 1, &c_str, NULL);
+    glCompileShader(shader);
+
+    GLint compile_status;
+    glGetShaderiv(shader, GL_COMPILE_STATUS, &compile_status);
+
+    if (!compile_status) {
+      GLchar log[5000];
+      GLsizei length = 0;
+      glGetShaderInfoLog(shader, sizeof(log), &length, log);
+      shader_print_errors("compile", log, c_str);
+      return 0;
+    }
+
+    glAttachShader(program, shader);
+  }
+
+  /* Link output. */
+  glBindFragDataLocation(program, 0, "fragColor");
+
+  /* Link and error check. */
+  glLinkProgram(program);
+
+  GLint link_status;
+  glGetProgramiv(program, GL_LINK_STATUS, &link_status);
+  if (!link_status) {
+    GLchar log[5000];
+    GLsizei length = 0;
+    glGetShaderInfoLog(program, sizeof(log), &length, log);
+    shader_print_errors("linking", log, VERTEX_SHADER);
+    shader_print_errors("linking", log, FRAGMENT_SHADER);
+    return 0;
+  }
+
+  return program;
+}
+
+int OpenGLShader::get_position_attrib_location()
+{
+  if (position_attribute_location_ == -1) {
+    const uint shader_program = get_shader_program();
+    position_attribute_location_ = glGetAttribLocation(shader_program, position_attribute_name);
+  }
+  return position_attribute_location_;
+}
+
+int OpenGLShader::get_tex_coord_attrib_location()
+{
+  if (tex_coord_attribute_location_ == -1) {
+    const uint shader_program = get_shader_program();
+    tex_coord_attribute_location_ = glGetAttribLocation(shader_program, tex_coord_attribute_name);
+  }
+  return tex_coord_attribute_location_;
+}
+
+void OpenGLShader::bind(int width, int height)
+{
+  create_shader_if_needed();
+
+  if (!shader_program_) {
+    return;
+  }
+
+  glUseProgram(shader_program_);
+  glUniform1i(image_texture_location_, 0);
+  glUniform2f(fullscreen_location_, width, height);
+}
+
+void OpenGLShader::unbind()
+{
+}
+
+uint OpenGLShader::get_shader_program()
+{
+  return shader_program_;
+}
+
+void OpenGLShader::create_shader_if_needed()
+{
+  if (shader_program_ || shader_compile_attempted_) {
+    return;
+  }
+
+  shader_compile_attempted_ = true;
+
+  shader_program_ = compile_shader_program();
+  if (!shader_program_) {
+    return;
+  }
+
+  glUseProgram(shader_program_);
+
+  image_texture_location_ = glGetUniformLocation(shader_program_, "image_texture");
+  if (image_texture_location_ < 0) {
+    LOG(ERROR) << "Shader doesn't contain the 'image_texture' uniform.";
+    destroy_shader();
+    return;
+  }
+
+  fullscreen_location_ = glGetUniformLocation(shader_program_, "fullscreen");
+  if (fullscreen_location_ < 0) {
+    LOG(ERROR) << "Shader doesn't contain the 'fullscreen' uniform.";
+    destroy_shader();
+    return;
+  }
+}
+
+void OpenGLShader::destroy_shader()
+{
+  glDeleteProgram(shader_program_);
+  shader_program_ = 0;
+}
+
+CCL_NAMESPACE_END

intern/cycles/app/opengl/shader.h

@@ -0,0 +1,58 @@
+/*
+ * Copyright 2011-2022 OpenGL 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.
+ */
+
+#pragma once
+
+#include "util/types.h"
+
+CCL_NAMESPACE_BEGIN
+
+class OpenGLShader {
+ public:
+  static constexpr const char *position_attribute_name = "pos";
+  static constexpr const char *tex_coord_attribute_name = "texCoord";
+
+  OpenGLShader() = default;
+  virtual ~OpenGLShader() = default;
+
+  /* Get attribute location for position and texture coordinate respectively.
+   * NOTE: The shader needs to be bound to have access to those. */
+  int get_position_attrib_location();
+  int get_tex_coord_attrib_location();
+
+  void bind(int width, int height);
+  void unbind();
+
+ protected:
+  uint get_shader_program();
+
+  void create_shader_if_needed();
+  void destroy_shader();
+
+  /* Cached values of various OpenGL resources. */
+  int position_attribute_location_ = -1;
+  int tex_coord_attribute_location_ = -1;
+
+  uint shader_program_ = 0;
+  int image_texture_location_ = -1;
+  int fullscreen_location_ = -1;
+
+  /* Shader compilation attempted. Which means, that if the shader program is 0 then compilation or
+   * linking has failed. Do not attempt to re-compile the shader. */
+  bool shader_compile_attempted_ = false;
+};
+
+CCL_NAMESPACE_END

intern/cycles/app/opengl/window.cpp

@@ -0,0 +1,365 @@
+/*
+ * 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 <stdio.h>
+#include <stdlib.h>
+
+#include "app/opengl/window.h"
+
+#include "util/string.h"
+#include "util/thread.h"
+#include "util/time.h"
+#include "util/version.h"
+
+#include <GL/glew.h>
+#include <SDL.h>
+
+CCL_NAMESPACE_BEGIN
+
+/* structs */
+
+struct Window {
+  WindowInitFunc initf = nullptr;
+  WindowExitFunc exitf = nullptr;
+  WindowResizeFunc resize = nullptr;
+  WindowDisplayFunc display = nullptr;
+  WindowKeyboardFunc keyboard = nullptr;
+  WindowMotionFunc motion = nullptr;
+
+  bool first_display = true;
+  bool redraw = false;
+
+  int mouseX = 0, mouseY = 0;
+  int mouseBut0 = 0, mouseBut2 = 0;
+
+  int width = 0, height = 0;
+
+  SDL_Window *window = nullptr;
+  SDL_GLContext gl_context = nullptr;
+  thread_mutex gl_context_mutex;
+} V;
+
+/* public */
+
+static void window_display_text(int x, int y, const char *text)
+{
+/* Not currently supported, need to add text rendering support. */
+#if 0
+  const char *c;
+
+  glRasterPos3f(x, y, 0);
+  glPixelStorei(GL_UNPACK_ALIGNMENT, 1);
+
+  printf("display %s\n", text);
+
+  for (c = text; *c != '\0'; c++) {
+    const uint8_t *bitmap = helvetica10_character_map[*c];
+    glBitmap(bitmap[0],
+             helvetica10_height,
+             helvetica10_x_offset,
+             helvetica10_y_offset,
+             bitmap[0],
+             0.0f,
+             bitmap + 1);
+  }
+#else
+  static string last_text = "";
+
+  if (text != last_text) {
+    printf("%s\n", text);
+    last_text = text;
+  }
+#endif
+}
+
+void window_display_info(const char *info)
+{
+  const int height = 20;
+
+#if 0
+  glEnable(GL_BLEND);
+  glBlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
+  glColor4f(0.1f, 0.1f, 0.1f, 0.8f);
+  glRectf(0.0f, V.height - height, V.width, V.height);
+  glDisable(GL_BLEND);
+
+  glColor3f(0.5f, 0.5f, 0.5f);
+#endif
+
+  window_display_text(10, 7 + V.height - height, info);
+
+#if 0
+  glColor3f(1.0f, 1.0f, 1.0f);
+#endif
+}
+
+void window_display_help()
+{
+  const int w = (int)((float)V.width / 1.15f);
+  const int h = (int)((float)V.height / 1.15f);
+
+  const int x1 = (V.width - w) / 2;
+#if 0
+  const int x2 = x1 + w;
+#endif
+
+  const int y1 = (V.height - h) / 2;
+  const int y2 = y1 + h;
+
+#if 0
+  glEnable(GL_BLEND);
+  glBlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
+  glColor4f(0.5f, 0.5f, 0.5f, 0.8f);
+  glRectf(x1, y1, x2, y2);
+  glDisable(GL_BLEND);
+
+  glColor3f(0.8f, 0.8f, 0.8f);
+#endif
+
+  string info = string("Cycles Renderer ") + CYCLES_VERSION_STRING;
+
+  window_display_text(x1 + 20, y2 - 20, info.c_str());
+  window_display_text(x1 + 20, y2 - 40, "(C) 2011-2016 Blender Foundation");
+  window_display_text(x1 + 20, y2 - 80, "Controls:");
+  window_display_text(x1 + 20, y2 - 100, "h:  Info/Help");
+  window_display_text(x1 + 20, y2 - 120, "r:  Reset");
+  window_display_text(x1 + 20, y2 - 140, "p:  Pause");
+  window_display_text(x1 + 20, y2 - 160, "esc:  Cancel");
+  window_display_text(x1 + 20, y2 - 180, "q:  Quit program");
+
+  window_display_text(x1 + 20, y2 - 210, "i:  Interactive mode");
+  window_display_text(x1 + 20, y2 - 230, "Left mouse:  Move camera");
+  window_display_text(x1 + 20, y2 - 250, "Right mouse:  Rotate camera");
+  window_display_text(x1 + 20, y2 - 270, "W/A/S/D:  Move camera");
+  window_display_text(x1 + 20, y2 - 290, "0/1/2/3:  Set max bounces");
+
+#if 0
+  glColor3f(1.0f, 1.0f, 1.0f);
+#endif
+}
+
+static void window_display()
+{
+  if (V.first_display) {
+    if (V.initf) {
+      V.initf();
+    }
+    if (V.exitf) {
+      atexit(V.exitf);
+    }
+
+    V.first_display = false;
+  }
+
+  window_opengl_context_enable();
+
+  glViewport(0, 0, V.width, V.height);
+
+  glMatrixMode(GL_PROJECTION);
+  glLoadIdentity();
+
+  glMatrixMode(GL_MODELVIEW);
+  glLoadIdentity();
+
+  glClearColor(0.05f, 0.05f, 0.05f, 0.0f);
+  glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
+
+  glMatrixMode(GL_PROJECTION);
+  glLoadIdentity();
+  glOrtho(0, V.width, 0, V.height, -1, 1);
+
+  glMatrixMode(GL_MODELVIEW);
+  glLoadIdentity();
+
+  glRasterPos3f(0, 0, 0);
+
+  if (V.display)
+    V.display();
+
+  SDL_GL_SwapWindow(V.window);
+  window_opengl_context_disable();
+}
+
+static void window_reshape(int width, int height)
+{
+  if (V.width != width || V.height != height) {
+    if (V.resize) {
+      V.resize(width, height);
+    }
+  }
+
+  V.width = width;
+  V.height = height;
+}
+
+static bool window_keyboard(unsigned char key)
+{
+  if (V.keyboard)
+    V.keyboard(key);
+
+  if (key == 'q') {
+    if (V.exitf)
+      V.exitf();
+    return true;
+  }
+
+  return false;
+}
+
+static void window_mouse(int button, int state, int x, int y)
+{
+  if (button == SDL_BUTTON_LEFT) {
+    if (state == SDL_MOUSEBUTTONDOWN) {
+      V.mouseX = x;
+      V.mouseY = y;
+      V.mouseBut0 = 1;
+    }
+    else if (state == SDL_MOUSEBUTTONUP) {
+      V.mouseBut0 = 0;
+    }
+  }
+  else if (button == SDL_BUTTON_RIGHT) {
+    if (state == SDL_MOUSEBUTTONDOWN) {
+      V.mouseX = x;
+      V.mouseY = y;
+      V.mouseBut2 = 1;
+    }
+    else if (state == SDL_MOUSEBUTTONUP) {
+      V.mouseBut2 = 0;
+    }
+  }
+}
+
+static void window_motion(int x, int y)
+{
+  const int but = V.mouseBut0 ? 0 : 2;
+  const int distX = x - V.mouseX;
+  const int distY = y - V.mouseY;
+
+  if (V.motion)
+    V.motion(distX, distY, but);
+
+  V.mouseX = x;
+  V.mouseY = y;
+}
+
+bool window_opengl_context_enable()
+{
+  V.gl_context_mutex.lock();
+  SDL_GL_MakeCurrent(V.window, V.gl_context);
+  return true;
+}
+
+void window_opengl_context_disable()
+{
+  SDL_GL_MakeCurrent(V.window, nullptr);
+  V.gl_context_mutex.unlock();
+}
+
+void window_main_loop(const char *title,
+                      int width,
+                      int height,
+                      WindowInitFunc initf,
+                      WindowExitFunc exitf,
+                      WindowResizeFunc resize,
+                      WindowDisplayFunc display,
+                      WindowKeyboardFunc keyboard,
+                      WindowMotionFunc motion)
+{
+  V.width = width;
+  V.height = height;
+  V.first_display = true;
+  V.redraw = false;
+  V.initf = initf;
+  V.exitf = exitf;
+  V.resize = resize;
+  V.display = display;
+  V.keyboard = keyboard;
+  V.motion = motion;
+
+  SDL_Init(SDL_INIT_VIDEO);
+  SDL_GL_SetAttribute(SDL_GL_CONTEXT_PROFILE_MASK, SDL_GL_CONTEXT_PROFILE_CORE);
+  SDL_GL_SetAttribute(SDL_GL_SHARE_WITH_CURRENT_CONTEXT, 1);
+  V.window = SDL_CreateWindow(title,
+                              SDL_WINDOWPOS_UNDEFINED,
+                              SDL_WINDOWPOS_UNDEFINED,
+                              width,
+                              height,
+                              SDL_WINDOW_RESIZABLE | SDL_WINDOW_OPENGL | SDL_WINDOW_SHOWN);
+  if (V.window == nullptr) {
+    fprintf(stderr, "Failed to create window: %s\n", SDL_GetError());
+    return;
+  }
+
+  SDL_RaiseWindow(V.window);
+
+  V.gl_context = SDL_GL_CreateContext(V.window);
+  glewInit();
+  SDL_GL_MakeCurrent(V.window, nullptr);
+
+  window_reshape(width, height);
+  window_display();
+
+  while (true) {
+    bool quit = false;
+    SDL_Event event;
+    while (!quit && SDL_PollEvent(&event)) {
+      if (event.type == SDL_TEXTINPUT) {
+        quit = window_keyboard(event.text.text[0]);
+      }
+      else if (event.type == SDL_MOUSEMOTION) {
+        window_motion(event.motion.x, event.motion.y);
+      }
+      else if (event.type == SDL_MOUSEBUTTONDOWN || event.type == SDL_MOUSEBUTTONUP) {
+        window_mouse(event.button.button, event.button.state, event.button.x, event.button.y);
+      }
+      else if (event.type == SDL_WINDOWEVENT) {
+        if (event.window.event == SDL_WINDOWEVENT_RESIZED ||
+            event.window.event == SDL_WINDOWEVENT_SIZE_CHANGED) {
+          window_reshape(event.window.data1, event.window.data2);
+        }
+      }
+      else if (event.type == SDL_QUIT) {
+        if (V.exitf) {
+          V.exitf();
+        }
+        quit = true;
+      }
+    }
+
+    if (quit) {
+      break;
+    }
+
+    if (V.redraw) {
+      V.redraw = false;
+      window_display();
+    }
+
+    SDL_WaitEventTimeout(NULL, 100);
+  }
+
+  SDL_GL_DeleteContext(V.gl_context);
+  SDL_DestroyWindow(V.window);
+  SDL_Quit();
+}
+
+void window_redraw()
+{
+  V.redraw = true;
+}
+
+CCL_NAMESPACE_END

intern/cycles/app/opengl/window.h

@@ -0,0 +1,48 @@
+/*
+ * 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.
+ */
+
+#pragma once
+
+/* Functions to display a simple OpenGL window using SDL, simplified to the
+ * bare minimum we need to reduce boilerplate code in tests apps. */
+
+CCL_NAMESPACE_BEGIN
+
+typedef void (*WindowInitFunc)();
+typedef void (*WindowExitFunc)();
+typedef void (*WindowResizeFunc)(int width, int height);
+typedef void (*WindowDisplayFunc)();
+typedef void (*WindowKeyboardFunc)(unsigned char key);
+typedef void (*WindowMotionFunc)(int x, int y, int button);
+
+void window_main_loop(const char *title,
+                      int width,
+                      int height,
+                      WindowInitFunc initf,
+                      WindowExitFunc exitf,
+                      WindowResizeFunc resize,
+                      WindowDisplayFunc display,
+                      WindowKeyboardFunc keyboard,
+                      WindowMotionFunc motion);
+
+void window_display_info(const char *info);
+void window_display_help();
+void window_redraw();
+
+bool window_opengl_context_enable();
+void window_opengl_context_disable();
+
+CCL_NAMESPACE_END

intern/cycles/blender/addon/properties.py

@@ -74,7 +74,7 @@ enum_panorama_types = (
                           "Similar to most fisheye modern lens, takes sensor dimensions into consideration"),
     ('MIRRORBALL', "Mirror Ball", "Uses the mirror ball mapping"),
     ('FISHEYE_LENS_POLYNOMIAL', "Fisheye Lens Polynomial",
-     "Defines the lens projection as polynomial to allow real world camera lenses to be mimicked."),
+     "Defines the lens projection as polynomial to allow real world camera lenses to be mimicked"),
 )
 
 enum_curve_shape = (
@@ -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",
@@ -896,27 +901,27 @@ class CyclesCameraSettings(bpy.types.PropertyGroup):
 
     fisheye_polynomial_k0: FloatProperty(
         name="Fisheye Polynomial K0",
-        description="Coefficient K0 of the lens polinomial",
+        description="Coefficient K0 of the lens polynomial",
         default=camera.default_fisheye_polynomial[0], precision=6, step=0.1, subtype='ANGLE',
     )
     fisheye_polynomial_k1: FloatProperty(
         name="Fisheye Polynomial K1",
-        description="Coefficient K1 of the lens polinomial",
+        description="Coefficient K1 of the lens polynomial",
         default=camera.default_fisheye_polynomial[1], precision=6, step=0.1, subtype='ANGLE',
     )
     fisheye_polynomial_k2: FloatProperty(
         name="Fisheye Polynomial K2",
-        description="Coefficient K2 of the lens polinomial",
+        description="Coefficient K2 of the lens polynomial",
         default=camera.default_fisheye_polynomial[2], precision=6, step=0.1, subtype='ANGLE',
     )
     fisheye_polynomial_k3: FloatProperty(
         name="Fisheye Polynomial K3",
-        description="Coefficient K3 of the lens polinomial",
+        description="Coefficient K3 of the lens polynomial",
         default=camera.default_fisheye_polynomial[3], precision=6, step=0.1, subtype='ANGLE',
     )
     fisheye_polynomial_k4: FloatProperty(
         name="Fisheye Polynomial K4",
-        description="Coefficient K4 of the lens polinomial",
+        description="Coefficient K4 of the lens polynomial",
         default=camera.default_fisheye_polynomial[4], precision=6, step=0.1, subtype='ANGLE',
     )
 
@@ -1369,6 +1374,12 @@ class CyclesPreferences(bpy.types.AddonPreferences):
         default=False,
     )
 
+    use_metalrt: BoolProperty(
+        name="MetalRT (Experimental)",
+        description="MetalRT for ray tracing uses less memory for scenes which use curves extensively, and can give better performance in specific cases. However this support is experimental and some scenes may render incorrectly",
+        default=False,
+    )
+
     def find_existing_device_entry(self, device):
         for device_entry in self.devices:
             if device_entry.id == device[2] and device_entry.type == device[1]:
@@ -1447,6 +1458,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
 
@@ -1474,7 +1498,8 @@ class CyclesPreferences(bpy.types.AddonPreferences):
                 if sys.platform[:3] == "win":
                     col.label(text="and AMD Radeon Pro 21.Q4 driver or newer", icon='BLANK1')
             elif device_type == 'METAL':
-                col.label(text="Requires Apple Silicon and macOS 12.0 or newer", icon='BLANK1')
+                col.label(text="Requires Apple Silicon with macOS 12.2 or newer", icon='BLANK1')
+                col.label(text="or AMD with macOS 12.3 or newer", icon='BLANK1')
             return
 
         for device in devices:
@@ -1501,6 +1526,12 @@ class CyclesPreferences(bpy.types.AddonPreferences):
             row.use_property_split = True
             row.prop(self, "peer_memory")
 
+        if compute_device_type == 'METAL':
+            row = layout.row()
+            row.use_property_split = True
+            row.prop(self, "use_metalrt")
+
+
     def draw(self, context):
         self.draw_impl(self.layout, context)
 

intern/cycles/blender/addon/ui.py

@@ -118,6 +118,12 @@ def use_optix(context):
 
     return (get_device_type(context) == 'OPTIX' and cscene.device == 'GPU')
 
+def use_multi_device(context):
+    cscene = context.scene.cycles
+    if cscene.device != 'GPU':
+        return False
+    return context.preferences.addons[__package__].preferences.has_multi_device()
+
 
 def show_device_active(context):
     cscene = context.scene.cycles
@@ -661,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 use_multi_device(context)
+
     def draw(self, context):
         import _cycles
 
@@ -673,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):

intern/cycles/blender/device.cpp

@@ -118,6 +118,10 @@ DeviceInfo blender_device_info(BL::Preferences &b_preferences, BL::Scene &b_scen
     device.has_peer_memory = false;
   }
 
+  if (get_boolean(cpreferences, "use_metalrt")) {
+    device.use_metalrt = true;
+  }
+
   return device;
 }
 

intern/cycles/blender/display_driver.cpp

@@ -480,26 +480,12 @@ class DrawTile {
       return false;
     }
 
-    if (!gl_vertex_buffer) {
-      glGenBuffers(1, &gl_vertex_buffer);
-      if (!gl_vertex_buffer) {
-        LOG(ERROR) << "Error allocating tile VBO.";
-        gl_resources_destroy();
-        return false;
-      }
-    }
-
     return true;
   }
 
   void gl_resources_destroy()
   {
     texture.gl_resources_destroy();
-
-    if (gl_vertex_buffer) {
-      glDeleteBuffers(1, &gl_vertex_buffer);
-      gl_vertex_buffer = 0;
-    }
   }
 
   inline bool ready_to_draw() const
@@ -512,9 +498,6 @@ class DrawTile {
 
   /* Display parameters the texture of this tile has been updated for. */
   BlenderDisplayDriver::Params params;
-
-  /* OpenGL resources needed for drawing. */
-  uint gl_vertex_buffer = 0;
 };
 
 class DrawTileAndPBO {
@@ -560,6 +543,30 @@ struct BlenderDisplayDriver::Tiles {
       tiles.clear();
     }
   } finished_tiles;
+
+  /* OpenGL vertex buffer needed for drawing. */
+  uint gl_vertex_buffer = 0;
+
+  bool gl_resources_ensure()
+  {
+    if (!gl_vertex_buffer) {
+      glGenBuffers(1, &gl_vertex_buffer);
+      if (!gl_vertex_buffer) {
+        LOG(ERROR) << "Error allocating tile VBO.";
+        return false;
+      }
+    }
+
+    return true;
+  }
+
+  void gl_resources_destroy()
+  {
+    if (gl_vertex_buffer) {
+      glDeleteBuffers(1, &gl_vertex_buffer);
+      gl_vertex_buffer = 0;
+    }
+  }
 };
 
 BlenderDisplayDriver::BlenderDisplayDriver(BL::RenderEngine &b_engine, BL::Scene &b_scene)
@@ -626,6 +633,12 @@ bool BlenderDisplayDriver::update_begin(const Params &params,
     need_clear_ = false;
   }
 
+  if (!tiles_->gl_resources_ensure()) {
+    tiles_->gl_resources_destroy();
+    gl_context_disable();
+    return false;
+  }
+
   if (!tiles_->current_tile.gl_resources_ensure()) {
     tiles_->current_tile.gl_resources_destroy();
     gl_context_disable();
@@ -825,7 +838,8 @@ static void vertex_buffer_update(const DisplayDriver::Params &params)
 static void draw_tile(const float2 &zoom,
                       const int texcoord_attribute,
                       const int position_attribute,
-                      const DrawTile &draw_tile)
+                      const DrawTile &draw_tile,
+                      const uint gl_vertex_buffer)
 {
   if (!draw_tile.ready_to_draw()) {
     return;
@@ -834,9 +848,9 @@ static void draw_tile(const float2 &zoom,
   const GLTexture &texture = draw_tile.texture;
 
   DCHECK_NE(texture.gl_id, 0);
-  DCHECK_NE(draw_tile.gl_vertex_buffer, 0);
+  DCHECK_NE(gl_vertex_buffer, 0);
 
-  glBindBuffer(GL_ARRAY_BUFFER, draw_tile.gl_vertex_buffer);
+  glBindBuffer(GL_ARRAY_BUFFER, gl_vertex_buffer);
 
   /* Draw at the parameters for which the texture has been updated for. This allows to always draw
    * texture during bordered-rendered camera view without flickering. The validness of the display
@@ -956,10 +970,14 @@ void BlenderDisplayDriver::draw(const Params &params)
   glEnableVertexAttribArray(texcoord_attribute);
   glEnableVertexAttribArray(position_attribute);
 
-  draw_tile(zoom_, texcoord_attribute, position_attribute, tiles_->current_tile.tile);
+  draw_tile(zoom_,
+            texcoord_attribute,
+            position_attribute,
+            tiles_->current_tile.tile,
+            tiles_->gl_vertex_buffer);
 
   for (const DrawTile &tile : tiles_->finished_tiles.tiles) {
-    draw_tile(zoom_, texcoord_attribute, position_attribute, tile);
+    draw_tile(zoom_, texcoord_attribute, position_attribute, tile, tiles_->gl_vertex_buffer);
   }
 
   display_shader_->unbind();
@@ -1062,6 +1080,7 @@ void BlenderDisplayDriver::gl_resources_destroy()
 
   tiles_->current_tile.gl_resources_destroy();
   tiles_->finished_tiles.gl_resources_destroy_and_clear();
+  tiles_->gl_resources_destroy();
 
   gl_context_disable();
 

intern/cycles/blender/session.cpp

@@ -506,8 +506,13 @@ void BlenderSession::render_frame_finish()
   session->set_output_driver(nullptr);
   session->full_buffer_written_cb = function_null;
 
-  /* The display driver holds OpenGL resources which belong to an OpenGL context held by the render
-   * engine on Blender side. Force destruction of those resources. */
+  /* The display driver is the source of drawing context for both drawing and possible graphics
+   * interop objects in the path trace. Once the frame is finished the OpenGL context might be
+   * freed form Blender side. Need to ensure that all GPU resources are freed prior to that
+   * point.
+   * Ideally would only do this when OpenGL context is actually destroyed, but there is no way to
+   * know when this happens (at least in the code at the time when this comment was written).
+   * The penalty of re-creating resources on every frame is unlikely to be noticed. */
   display_driver_ = nullptr;
   session->set_display_driver(nullptr);
 

intern/cycles/blender/shader.cpp

@@ -45,7 +45,8 @@ typedef map<string, ConvertNode *> ProxyMap;
 
 void BlenderSync::find_shader(BL::ID &id, array<Node *> &used_shaders, Shader *default_shader)
 {
-  Shader *shader = (id) ? shader_map.find(id) : default_shader;
+  Shader *synced_shader = (id) ? shader_map.find(id) : nullptr;
+  Shader *shader = (synced_shader) ? synced_shader : default_shader;
 
   used_shaders.push_back_slow(shader);
   shader->tag_used(scene);
@@ -689,6 +690,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/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 :
@@ -425,7 +471,7 @@ void BVHEmbree::add_instance(Object *ob, int i)
   assert(instance_bvh != NULL);
 
   const size_t num_object_motion_steps = ob->use_motion() ? ob->get_motion().size() : 1;
-  const size_t num_motion_steps = min(num_object_motion_steps, RTC_MAX_TIME_STEP_COUNT);
+  const size_t num_motion_steps = min(num_object_motion_steps, (size_t)RTC_MAX_TIME_STEP_COUNT);
   assert(num_object_motion_steps <= RTC_MAX_TIME_STEP_COUNT);
 
   RTCGeometry geom_id = rtcNewGeometry(rtc_device, RTC_GEOMETRY_TYPE_INSTANCE);
@@ -476,7 +522,7 @@ void BVHEmbree::add_triangles(const Object *ob, const Mesh *mesh, int i)
   }
 
   assert(num_motion_steps <= RTC_MAX_TIME_STEP_COUNT);
-  num_motion_steps = min(num_motion_steps, RTC_MAX_TIME_STEP_COUNT);
+  num_motion_steps = min(num_motion_steps, (size_t)RTC_MAX_TIME_STEP_COUNT);
 
   const size_t num_triangles = mesh->num_triangles();
 
@@ -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);
@@ -728,7 +775,7 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
   }
 
   assert(num_motion_steps <= RTC_MAX_TIME_STEP_COUNT);
-  num_motion_steps = min(num_motion_steps, RTC_MAX_TIME_STEP_COUNT);
+  num_motion_steps = min(num_motion_steps, (size_t)RTC_MAX_TIME_STEP_COUNT);
 
   const size_t num_curves = hair->num_curves();
   size_t num_segments = 0;
@@ -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.
    *
@@ -139,6 +142,7 @@ class BVHParams {
 
     top_level = false;
     bvh_layout = BVH_LAYOUT_BVH2;
+    use_compact_structure = true;
     use_unaligned_nodes = false;
 
     num_motion_curve_steps = 0;

intern/cycles/cmake/external_libs.cmake

@@ -491,26 +491,22 @@ else()
 endif()
 
 ###########################################################################
-# GLUT
+# SDL
 ###########################################################################
 
 if(WITH_CYCLES_STANDALONE AND WITH_CYCLES_STANDALONE_GUI)
-  if(MSVC AND EXISTS ${_cycles_lib_dir})
-    add_definitions(-DFREEGLUT_STATIC -DFREEGLUT_LIB_PRAGMAS=0)
-    set(GLUT_LIBRARIES "${_cycles_lib_dir}/opengl/lib/freeglut_static.lib")
-    set(GLUT_INCLUDE_DIR "${_cycles_lib_dir}/opengl/include")
-  else()
-    find_package(GLUT)
+  # We can't use the version from the Blender precompiled libraries because
+  # it does not include the video subsystem.
+  find_package(SDL2)
 
-    if(NOT GLUT_FOUND)
-      set(WITH_CYCLES_STANDALONE_GUI OFF)
-      message(STATUS "GLUT not found, disabling Cycles standalone GUI")
-    endif()
+  if(NOT SDL2_FOUND)
+    set(WITH_CYCLES_STANDALONE_GUI OFF)
+    message(STATUS "SDL not found, disabling Cycles standalone GUI")
   endif()
 
   include_directories(
     SYSTEM
-    ${GLUT_INCLUDE_DIR}
+    ${SDL2_INCLUDE_DIRS}
   )
 endif()
 
@@ -559,10 +555,10 @@ if(WITH_CYCLES_DEVICE_METAL)
   find_library(METAL_LIBRARY Metal)
 
   # This file was added in the 12.0 SDK, use it as a way to detect the version.
-  if (METAL_LIBRARY AND NOT EXISTS "${METAL_LIBRARY}/Headers/MTLFunctionStitching.h")
+  if(METAL_LIBRARY AND NOT EXISTS "${METAL_LIBRARY}/Headers/MTLFunctionStitching.h")
     message(STATUS "Metal version too old, must be SDK 12.0 or newer, disabling WITH_CYCLES_DEVICE_METAL")
     set(WITH_CYCLES_DEVICE_METAL OFF)
-  elseif (NOT METAL_LIBRARY)
+  elseif(NOT METAL_LIBRARY)
     message(STATUS "Metal not found, disabling WITH_CYCLES_DEVICE_METAL")
     set(WITH_CYCLES_DEVICE_METAL OFF)
   else()

intern/cycles/device/device.cpp

@@ -328,6 +328,7 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
   info.has_osl = true;
   info.has_profiling = true;
   info.has_peer_memory = false;
+  info.use_metalrt = false;
   info.denoisers = DENOISER_ALL;
 
   foreach (const DeviceInfo &device, subdevices) {
@@ -335,7 +336,7 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
     if (device.type == DEVICE_CPU && subdevices.size() > 1) {
       if (background) {
         int orig_cpu_threads = (threads) ? threads : TaskScheduler::max_concurrency();
-        int cpu_threads = max(orig_cpu_threads - (subdevices.size() - 1), 0);
+        int cpu_threads = max(orig_cpu_threads - (subdevices.size() - 1), size_t(0));
 
         VLOG(1) << "CPU render threads reduced from " << orig_cpu_threads << " to " << cpu_threads
                 << ", to dedicate to GPU.";
@@ -374,6 +375,7 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
     info.has_osl &= device.has_osl;
     info.has_profiling &= device.has_profiling;
     info.has_peer_memory |= device.has_peer_memory;
+    info.use_metalrt |= device.use_metalrt;
     info.denoisers &= device.denoisers;
   }
 

intern/cycles/device/device.h

@@ -79,6 +79,7 @@ class DeviceInfo {
   bool has_profiling;         /* Supports runtime collection of profiling info. */
   bool has_peer_memory;       /* GPU has P2P access to memory of another GPU. */
   bool has_gpu_queue;         /* Device supports GPU queue. */
+  bool use_metalrt;           /* Use MetalRT to accelerate ray queries (Metal only). */
   DenoiserTypeMask denoisers; /* Supported denoiser types. */
   int cpu_threads;
   vector<DeviceInfo> multi_devices;
@@ -96,6 +97,7 @@ class DeviceInfo {
     has_profiling = false;
     has_peer_memory = false;
     has_gpu_queue = false;
+    use_metalrt = false;
     denoisers = DENOISER_NONE;
   }
 

intern/cycles/device/hip/device_impl.h

@@ -25,8 +25,6 @@
 
 #  ifdef WITH_HIP_DYNLOAD
 #    include "hipew.h"
-#  else
-#    include "util/opengl.h"
 #  endif
 
 CCL_NAMESPACE_BEGIN

intern/cycles/device/memory.h

@@ -311,7 +311,7 @@ template<typename T> class device_only_memory : public device_memory {
       : device_memory(device, name, allow_host_memory_fallback ? MEM_READ_WRITE : MEM_DEVICE_ONLY)
   {
     data_type = device_type_traits<T>::data_type;
-    data_elements = max(device_type_traits<T>::num_elements, 1);
+    data_elements = max(device_type_traits<T>::num_elements, size_t(1));
   }
 
   device_only_memory(device_only_memory &&other) noexcept : device_memory(std::move(other))

intern/cycles/device/metal/bvh.mm

@@ -761,7 +761,7 @@ bool BVHMetal::build_TLAS(Progress &progress,
       num_instances++;
 
       if (ob->use_motion()) {
-        num_motion_transforms += max(1, ob->get_motion().size());
+        num_motion_transforms += max((size_t)1, ob->get_motion().size());
       }
       else {
         num_motion_transforms++;

intern/cycles/device/metal/device.mm

@@ -39,33 +39,20 @@ bool device_metal_init()
   return true;
 }
 
-static int device_metal_get_num_devices_safe(uint32_t *num_devices)
-{
-  *num_devices = MTLCopyAllDevices().count;
-  return 0;
-}
-
 void device_metal_info(vector<DeviceInfo> &devices)
 {
-  uint32_t num_devices = 0;
-  device_metal_get_num_devices_safe(&num_devices);
-  if (num_devices == 0) {
-    return;
-  }
-
-  vector<MetalPlatformDevice> usable_devices;
-  MetalInfo::get_usable_devices(&usable_devices);
+  auto usable_devices = MetalInfo::get_usable_devices();
   /* Devices are numbered consecutively across platforms. */
   set<string> unique_ids;
   int device_index = 0;
-  for (MetalPlatformDevice &device : usable_devices) {
+  for (id<MTLDevice> &device : usable_devices) {
     /* Compute unique ID for persistent user preferences. */
-    const string &device_name = device.device_name;
+    string device_name = [device.name UTF8String];
     string id = string("METAL_") + device_name;
 
     /* Hardware ID might not be unique, add device number in that case. */
     if (unique_ids.find(id) != unique_ids.end()) {
-      id += string_printf("_ID_%d", num_devices);
+      id += string_printf("_ID_%d", device_index);
     }
     unique_ids.insert(id);
 
@@ -94,15 +81,13 @@ void device_metal_info(vector<DeviceInfo> &devices)
 string device_metal_capabilities()
 {
   string result = "";
-  string error_msg = "";
-  uint32_t num_devices = 0;
-  assert(device_metal_get_num_devices_safe(&num_devices));
+  auto allDevices = MTLCopyAllDevices();
+  uint32_t num_devices = allDevices.count;
   if (num_devices == 0) {
     return "No Metal devices found\n";
   }
   result += string_printf("Number of devices: %u\n", num_devices);
 
-  NSArray<id<MTLDevice>> *allDevices = MTLCopyAllDevices();
   for (id<MTLDevice> device in allDevices) {
     result += string_printf("\t\tDevice: %s\n", [device.name UTF8String]);
   }

intern/cycles/device/metal/device_impl.mm

@@ -53,16 +53,10 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
   mtlDevId = info.num;
 
   /* select chosen device */
-  vector<MetalPlatformDevice> usable_devices;
-  MetalInfo::get_usable_devices(&usable_devices);
-  if (usable_devices.size() == 0) {
-    set_error("Metal: no devices found.");
-    return;
-  }
+  auto usable_devices = MetalInfo::get_usable_devices();
   assert(mtlDevId < usable_devices.size());
-  MetalPlatformDevice &platform_device = usable_devices[mtlDevId];
-  mtlDevice = platform_device.device_id;
-  device_name = platform_device.device_name;
+  mtlDevice = usable_devices[mtlDevId];
+  device_name = [mtlDevice.name UTF8String];
   device_vendor = MetalInfo::get_vendor_from_device_name(device_name);
   assert(device_vendor != METAL_GPU_UNKNOWN);
   metal_printf("Creating new Cycles device for Metal: %s\n", device_name.c_str());
@@ -100,6 +94,7 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
     }
   }
 
+  use_metalrt = info.use_metalrt;
   if (auto metalrt = getenv("CYCLES_METALRT")) {
     use_metalrt = (atoi(metalrt) != 0);
   }
@@ -455,8 +450,15 @@ MetalDevice::MetalMem *MetalDevice::generic_alloc(device_memory &mem)
   mem.device_pointer = 0;
 
   id<MTLBuffer> metal_buffer = nil;
+  MTLResourceOptions options = default_storage_mode;
+
+  /* Workaround for "bake" unit tests which fail if RenderBuffers is allocated with
+   * MTLResourceStorageModeShared. */
+  if (strstr(mem.name, "RenderBuffers")) {
+    options = MTLResourceStorageModeManaged;
+  }
+
   if (size > 0) {
-    MTLResourceOptions options = default_storage_mode;
     if (mem.type == MEM_DEVICE_ONLY) {
       options = MTLResourceStorageModePrivate;
     }
@@ -490,7 +492,7 @@ MetalDevice::MetalMem *MetalDevice::generic_alloc(device_memory &mem)
   mmem->mtlBuffer = metal_buffer;
   mmem->offset = 0;
   mmem->size = size;
-  if (mem.type != MEM_DEVICE_ONLY) {
+  if (options != MTLResourceStorageModePrivate) {
     mmem->hostPtr = [metal_buffer contents];
   }
   else {
@@ -759,6 +761,17 @@ void MetalDevice::tex_alloc_as_buffer(device_texture &mem)
 
 void MetalDevice::tex_alloc(device_texture &mem)
 {
+  /* Check that dimensions fit within maximum allowable size.
+     See https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf
+  */
+  if (mem.data_width > 16384 || mem.data_height > 16384) {
+    set_error(string_printf(
+        "Texture exceeds maximum allowed size of 16384 x 16384 (requested: %zu x %zu)",
+        mem.data_width,
+        mem.data_height));
+    return;
+  }
+
   MTLStorageMode storage_mode = MTLStorageModeManaged;
   if (@available(macos 10.15, *)) {
     if ([mtlDevice hasUnifiedMemory] &&

intern/cycles/device/metal/kernel.mm

@@ -59,10 +59,15 @@ bool MetalDeviceKernel::load(MetalDevice *device,
   }
 
   bool use_binary_archive = true;
-  if (getenv("CYCLES_METAL_DISABLE_BINARY_ARCHIVES")) {
+  if (device->device_vendor == METAL_GPU_APPLE) {
+    /* Workaround for T94142: Cycles Metal crash with simultaneous viewport and final render */
     use_binary_archive = false;
   }
 
+  if (auto str = getenv("CYCLES_METAL_DISABLE_BINARY_ARCHIVES")) {
+    use_binary_archive = (atoi(str) == 0);
+  }
+
   id<MTLBinaryArchive> archive = nil;
   string metalbin_path;
   if (use_binary_archive) {

intern/cycles/device/metal/util.h

@@ -36,33 +36,10 @@ enum MetalGPUVendor {
   METAL_GPU_INTEL = 3,
 };
 
-/* Retains a named MTLDevice for device enumeration. */
-struct MetalPlatformDevice {
-  MetalPlatformDevice(id<MTLDevice> device, const string &device_name)
-      : device_id(device), device_name(device_name)
-  {
-    [device_id retain];
-  }
-  ~MetalPlatformDevice()
-  {
-    [device_id release];
-  }
-  id<MTLDevice> device_id;
-  string device_name;
-};
-
 /* Contains static Metal helper functions. */
 struct MetalInfo {
-  static bool device_version_check(id<MTLDevice> device);
-  static void get_usable_devices(vector<MetalPlatformDevice> *usable_devices);
+  static vector<id<MTLDevice>> const &get_usable_devices();
   static MetalGPUVendor get_vendor_from_device_name(string const &device_name);
-
-  /* Platform information. */
-  static bool get_num_devices(uint32_t *num_platforms);
-  static uint32_t get_num_devices();
-
-  static bool get_device_name(id<MTLDevice> device_id, string *device_name);
-  static string get_device_name(id<MTLDevice> device_id);
 };
 
 /* Pool of MTLBuffers whose lifetime is linked to a single MTLCommandBuffer */

intern/cycles/device/metal/util.mm

@@ -43,83 +43,45 @@ MetalGPUVendor MetalInfo::get_vendor_from_device_name(string const &device_name)
   return METAL_GPU_UNKNOWN;
 }
 
-bool MetalInfo::device_version_check(id<MTLDevice> device)
+vector<id<MTLDevice>> const &MetalInfo::get_usable_devices()
 {
-  /* Metal Cycles doesn't work correctly on macOS versions older than 12.0 */
-  if (@available(macos 12.0, *)) {
-    MetalGPUVendor vendor = get_vendor_from_device_name([[device name] UTF8String]);
+  static vector<id<MTLDevice>> usable_devices;
+  static bool already_enumerated = false;
 
-    /* Metal Cycles works on Apple Silicon GPUs at present */
-    return (vendor == METAL_GPU_APPLE);
+  if (already_enumerated) {
+    return usable_devices;
   }
 
-  return false;
-}
+  metal_printf("Usable Metal devices:\n");
+  for (id<MTLDevice> device in MTLCopyAllDevices()) {
+    const char *device_name = [device.name UTF8String];
 
-void MetalInfo::get_usable_devices(vector<MetalPlatformDevice> *usable_devices)
-{
-  static bool first_time = true;
-#  define FIRST_VLOG(severity) \
-    if (first_time) \
-    VLOG(severity)
+    MetalGPUVendor vendor = get_vendor_from_device_name(device_name);
+    bool usable = false;
 
-  usable_devices->clear();
-
-  NSArray<id<MTLDevice>> *allDevices = MTLCopyAllDevices();
-  for (id<MTLDevice> device in allDevices) {
-    string device_name;
-    if (!get_device_name(device, &device_name)) {
-      FIRST_VLOG(2) << "Failed to get device name, ignoring.";
-      continue;
+    if (@available(macos 12.2, *)) {
+      usable |= (vendor == METAL_GPU_APPLE);
     }
 
-    static const char *forceIntelStr = getenv("CYCLES_METAL_FORCE_INTEL");
-    bool forceIntel = forceIntelStr ? (atoi(forceIntelStr) != 0) : false;
-    if (forceIntel && device_name.find("Intel") == string::npos) {
-      FIRST_VLOG(2) << "CYCLES_METAL_FORCE_INTEL causing non-Intel device " << device_name
-                    << " to be ignored.";
-      continue;
+    if (@available(macos 12.3, *)) {
+      usable |= (vendor == METAL_GPU_AMD);
     }
 
-    if (!device_version_check(device)) {
-      FIRST_VLOG(2) << "Ignoring device " << device_name << " due to too old compiler version.";
-      continue;
+    if (usable) {
+      metal_printf("- %s\n", device_name);
+      [device retain];
+      usable_devices.push_back(device);
+    }
+    else {
+      metal_printf("  (skipping \"%s\")\n", device_name);
     }
-    FIRST_VLOG(2) << "Adding new device " << device_name << ".";
-    string hardware_id;
-    usable_devices->push_back(MetalPlatformDevice(device, device_name));
   }
-  first_time = false;
-}
-
-bool MetalInfo::get_num_devices(uint32_t *num_devices)
-{
-  *num_devices = MTLCopyAllDevices().count;
-  return true;
-}
-
-uint32_t MetalInfo::get_num_devices()
-{
-  uint32_t num_devices;
-  if (!get_num_devices(&num_devices)) {
-    return 0;
+  if (usable_devices.empty()) {
+    metal_printf("   No usable Metal devices found\n");
   }
-  return num_devices;
-}
+  already_enumerated = true;
 
-bool MetalInfo::get_device_name(id<MTLDevice> device, string *platform_name)
-{
-  *platform_name = [device.name UTF8String];
-  return true;
-}
-
-string MetalInfo::get_device_name(id<MTLDevice> device)
-{
-  string platform_name;
-  if (!get_device_name(device, &platform_name)) {
-    return "";
-  }
-  return platform_name;
+  return usable_devices;
 }
 
 id<MTLBuffer> MetalBufferPool::get_buffer(id<MTLDevice> device,

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 */
@@ -1586,7 +1586,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         if (ob->is_traceable() && ob->use_motion()) {
           total_motion_transform_size = align_up(total_motion_transform_size,
                                                  OPTIX_TRANSFORM_BYTE_ALIGNMENT);
-          const size_t motion_keys = max(ob->get_motion().size(), 2) - 2;
+          const size_t motion_keys = max(ob->get_motion().size(), (size_t)2) - 2;
           total_motion_transform_size = total_motion_transform_size +
                                         sizeof(OptixSRTMotionTransform) +
                                         motion_keys * sizeof(OptixSRTData);
@@ -1660,7 +1660,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
 
       /* Insert motion traversable if object has motion. */
       if (motion_blur && ob->use_motion()) {
-        size_t motion_keys = max(ob->get_motion().size(), 2) - 2;
+        size_t motion_keys = max(ob->get_motion().size(), (size_t)2) - 2;
         size_t motion_transform_size = sizeof(OptixSRTMotionTransform) +
                                        motion_keys * sizeof(OptixSRTData);
 

intern/cycles/integrator/denoiser.cpp

@@ -125,20 +125,41 @@ static Device *find_best_device(Device *device, DenoiserType type)
   return best_device;
 }
 
+static DeviceInfo find_best_denoiser_device_info(const vector<DeviceInfo> &device_infos,
+                                                 DenoiserType denoiser_type)
+{
+  for (const DeviceInfo &device_info : device_infos) {
+    if ((device_info.denoisers & denoiser_type) == 0) {
+      continue;
+    }
+
+    /* TODO(sergey): Use one of the already configured devices, so that OptiX denoising can happen
+     * on a physical CUDA device which is already used for rendering. */
+
+    /* TODO(sergey): Choose fastest device for denoising. */
+
+    return device_info;
+  }
+
+  DeviceInfo none_device;
+  none_device.type = DEVICE_NONE;
+  return none_device;
+}
+
 static unique_ptr<Device> create_denoiser_device(Device *path_trace_device,
-                                                 const uint device_type_mask)
+                                                 const uint device_type_mask,
+                                                 DenoiserType denoiser_type)
 {
   const vector<DeviceInfo> device_infos = Device::available_devices(device_type_mask);
   if (device_infos.empty()) {
     return nullptr;
   }
 
-  /* TODO(sergey): Use one of the already configured devices, so that OptiX denoising can happen on
-   * a physical CUDA device which is already used for rendering. */
-
-  /* TODO(sergey): Choose fastest device for denoising. */
-
-  const DeviceInfo denoiser_device_info = device_infos.front();
+  const DeviceInfo denoiser_device_info = find_best_denoiser_device_info(device_infos,
+                                                                         denoiser_type);
+  if (denoiser_device_info.type == DEVICE_NONE) {
+    return nullptr;
+  }
 
   unique_ptr<Device> denoiser_device(
       Device::create(denoiser_device_info, path_trace_device->stats, path_trace_device->profiler));
@@ -186,7 +207,8 @@ Device *Denoiser::ensure_denoiser_device(Progress *progress)
   device_creation_attempted_ = true;
 
   const uint device_type_mask = get_device_type_mask();
-  local_denoiser_device_ = create_denoiser_device(path_trace_device_, device_type_mask);
+  local_denoiser_device_ = create_denoiser_device(
+      path_trace_device_, device_type_mask, params_.type);
   denoiser_device_ = local_denoiser_device_.get();
 
   return denoiser_device_;

intern/cycles/integrator/denoiser_oidn.cpp

@@ -37,8 +37,6 @@ OIDNDenoiser::OIDNDenoiser(Device *path_trace_device, const DenoiseParams &param
     : Denoiser(path_trace_device, params)
 {
   DCHECK_EQ(params.type, DENOISER_OPENIMAGEDENOISE);
-
-  DCHECK(openimagedenoise_supported()) << "OpenImageDenoiser is not supported on this platform.";
 }
 
 #ifdef WITH_OPENIMAGEDENOISE
@@ -585,6 +583,9 @@ bool OIDNDenoiser::denoise_buffer(const BufferParams &buffer_params,
                                   const int num_samples,
                                   bool allow_inplace_modification)
 {
+  DCHECK(openimagedenoise_supported())
+      << "OpenImageDenoiser is not supported on this platform or build.";
+
 #ifdef WITH_OPENIMAGEDENOISE
   thread_scoped_lock lock(mutex_);
 
@@ -635,4 +636,20 @@ uint OIDNDenoiser::get_device_type_mask() const
   return DEVICE_MASK_CPU;
 }
 
+Device *OIDNDenoiser::ensure_denoiser_device(Progress *progress)
+{
+#ifndef WITH_OPENIMAGEDENOISE
+  path_trace_device_->set_error("Build without OpenImageDenoiser");
+  return nullptr;
+#else
+  if (!openimagedenoise_supported()) {
+    path_trace_device_->set_error(
+        "OpenImageDenoiser is not supported on this CPU: missing SSE 4.1 support");
+    return nullptr;
+  }
+
+  return Denoiser::ensure_denoiser_device(progress);
+#endif
+}
+
 CCL_NAMESPACE_END

intern/cycles/integrator/denoiser_oidn.h

@@ -38,6 +38,7 @@ class OIDNDenoiser : public Denoiser {
 
  protected:
   virtual uint get_device_type_mask() const override;
+  virtual Device *ensure_denoiser_device(Progress *progress) override;
 
   /* We only perform one denoising at a time, since OpenImageDenoise itself is multithreaded.
    * Use this mutex whenever images are passed to the OIDN and needs to be denoised. */

intern/cycles/integrator/path_trace.cpp

@@ -67,14 +67,7 @@ PathTrace::PathTrace(Device *device,
 
 PathTrace::~PathTrace()
 {
-  /* Destroy any GPU resource which was used for graphics interop.
-   * Need to have access to the PathTraceDisplay as it is the only source of drawing context which
-   * is used for interop. */
-  if (display_) {
-    for (auto &&path_trace_work : path_trace_works_) {
-      path_trace_work->destroy_gpu_resources(display_.get());
-    }
-  }
+  destroy_gpu_resources();
 }
 
 void PathTrace::load_kernels()
@@ -572,6 +565,11 @@ void PathTrace::set_output_driver(unique_ptr<OutputDriver> driver)
 
 void PathTrace::set_display_driver(unique_ptr<DisplayDriver> driver)
 {
+  /* The display driver is the source of the drawing context which might be used by
+   * path trace works. Make sure there is no graphics interop using resources from
+   * the old display, as it might no longer be available after this call. */
+  destroy_gpu_resources();
+
   if (driver) {
     display_ = make_unique<PathTraceDisplay>(move(driver));
   }
@@ -1088,6 +1086,18 @@ bool PathTrace::has_denoised_result() const
   return render_state_.has_denoised_result;
 }
 
+void PathTrace::destroy_gpu_resources()
+{
+  /* Destroy any GPU resource which was used for graphics interop.
+   * Need to have access to the PathTraceDisplay as it is the only source of drawing context which
+   * is used for interop. */
+  if (display_) {
+    for (auto &&path_trace_work : path_trace_works_) {
+      path_trace_work->destroy_gpu_resources(display_.get());
+    }
+  }
+}
+
 /* --------------------------------------------------------------------
  * Report generation.
  */

intern/cycles/integrator/path_trace.h

@@ -239,6 +239,9 @@ class PathTrace {
 
   void progress_set_status(const string &status, const string &substatus = "");
 
+  /* Destroy GPU resources (such as graphics interop) used by work. */
+  void destroy_gpu_resources();
+
   /* Pointer to a device which is configured to be used for path tracing. If multiple devices
    * are configured this is a `MultiDevice`. */
   Device *device_ = nullptr;

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

@@ -157,7 +157,11 @@ ccl_device_inline
                 }
               }
 
+              /* 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,
@@ -188,7 +192,11 @@ ccl_device_inline
                 }
               }
 
+              /* 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,

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

@@ -160,6 +160,9 @@ ccl_device_inline
                                         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: {

intern/cycles/kernel/bvh/traversal.h

@@ -133,35 +133,29 @@ 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);
+            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, 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);
-
-                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);
-
+              case PRIMITIVE_MOTION_TRIANGLE: {
                 if (motion_triangle_intersect(kg,
                                               isect,
                                               P,
@@ -176,28 +170,21 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
                   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 prim_object = (object == OBJECT_NONE) ?
-                                            kernel_tex_fetch(__prim_object, prim_addr) :
-                                            object;
-                const int prim = kernel_tex_fetch(__prim_index, prim_addr);
-
                 const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
                 const bool hit = curve_intersect(
                     kg, isect, P, dir, isect->t, prim_object, prim, ray->time, curve_type);
@@ -206,26 +193,19 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
                   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 prim_object = (object == OBJECT_NONE) ?
-                                            kernel_tex_fetch(__prim_object, prim_addr) :
-                                            object;
-                const int prim = kernel_tex_fetch(__prim_index, prim_addr);
-
                 const int point_type = kernel_tex_fetch(__prim_type, prim_addr);
                 const bool hit = point_intersect(
                     kg, isect, P, dir, isect->t, prim_object, prim, ray->time, point_type);
@@ -234,10 +214,10 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
                   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

@@ -144,6 +144,9 @@ ccl_device_inline
                                             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) {
@@ -164,6 +167,9 @@ ccl_device_inline
                                             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;

intern/cycles/kernel/bvh/volume_all.h

@@ -147,6 +147,9 @@ ccl_device_inline
                                             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;
@@ -188,6 +191,9 @@ ccl_device_inline
                                             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;

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

@@ -295,8 +295,11 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
                         int kernel_index);
   ccl_gpu_kernel_lambda_pass.kernel_index = kernel_index;
 
-  gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
-      num_states, indices, num_indices, ccl_gpu_kernel_lambda_pass);
+  gpu_parallel_active_index_array(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE,
+                                  num_states,
+                                  indices,
+                                  num_indices,
+                                  ccl_gpu_kernel_lambda_pass);
 }
 
 ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
@@ -310,8 +313,11 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
                         int kernel_index);
   ccl_gpu_kernel_lambda_pass.kernel_index = kernel_index;
 
-  gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
-      num_states, indices, num_indices, ccl_gpu_kernel_lambda_pass);
+  gpu_parallel_active_index_array(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE,
+                                  num_states,
+                                  indices,
+                                  num_indices,
+                                  ccl_gpu_kernel_lambda_pass);
 }
 
 ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
@@ -322,8 +328,11 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
 {
   ccl_gpu_kernel_lambda(INTEGRATOR_STATE(state, path, queued_kernel) != 0);
 
-  gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
-      num_states, indices, num_indices, ccl_gpu_kernel_lambda_pass);
+  gpu_parallel_active_index_array(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE,
+                                  num_states,
+                                  indices,
+                                  num_indices,
+                                  ccl_gpu_kernel_lambda_pass);
 }
 
 ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
@@ -335,8 +344,11 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
 {
   ccl_gpu_kernel_lambda(INTEGRATOR_STATE(state, path, queued_kernel) == 0);
 
-  gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
-      num_states, indices + indices_offset, num_indices, ccl_gpu_kernel_lambda_pass);
+  gpu_parallel_active_index_array(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE,
+                                  num_states,
+                                  indices + indices_offset,
+                                  num_indices,
+                                  ccl_gpu_kernel_lambda_pass);
 }
 
 ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
@@ -348,8 +360,11 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
 {
   ccl_gpu_kernel_lambda(INTEGRATOR_STATE(state, shadow_path, queued_kernel) == 0);
 
-  gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
-      num_states, indices + indices_offset, num_indices, ccl_gpu_kernel_lambda_pass);
+  gpu_parallel_active_index_array(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE,
+                                  num_states,
+                                  indices + indices_offset,
+                                  num_indices,
+                                  ccl_gpu_kernel_lambda_pass);
 }
 
 ccl_gpu_kernel_threads(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
@@ -391,8 +406,11 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
                         int num_active_paths);
   ccl_gpu_kernel_lambda_pass.num_active_paths = num_active_paths;
 
-  gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
-      num_states, indices, num_indices, ccl_gpu_kernel_lambda_pass);
+  gpu_parallel_active_index_array(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE,
+                                  num_states,
+                                  indices,
+                                  num_indices,
+                                  ccl_gpu_kernel_lambda_pass);
 }
 
 ccl_gpu_kernel_threads(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
@@ -424,8 +442,11 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
                         int num_active_paths);
   ccl_gpu_kernel_lambda_pass.num_active_paths = num_active_paths;
 
-  gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
-      num_states, indices, num_indices, ccl_gpu_kernel_lambda_pass);
+  gpu_parallel_active_index_array(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE,
+                                  num_states,
+                                  indices,
+                                  num_indices,
+                                  ccl_gpu_kernel_lambda_pass);
 }
 
 ccl_gpu_kernel_threads(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)

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

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

intern/cycles/kernel/device/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;

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/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.
  */

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

@@ -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;
 
@@ -190,40 +190,46 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals kg,
 #ifdef __OBJECT_MOTION__
     shader_setup_object_transforms(kg, sd, time);
 #endif
-  }
-  else if (lamp != LAMP_NONE) {
-    sd->lamp = lamp;
-  }
 
-  /* transform into world space */
-  if (object_space) {
-    object_position_transform_auto(kg, sd, &sd->P);
-    object_normal_transform_auto(kg, sd, &sd->Ng);
-    sd->N = sd->Ng;
-    object_dir_transform_auto(kg, sd, &sd->I);
-  }
+    /* transform into world space */
+    if (object_space) {
+      object_position_transform_auto(kg, sd, &sd->P);
+      object_normal_transform_auto(kg, sd, &sd->Ng);
+      sd->N = sd->Ng;
+      object_dir_transform_auto(kg, sd, &sd->I);
+    }
 
-  if (sd->type == PRIMITIVE_TRIANGLE) {
-    /* smooth normal */
-    if (sd->shader & SHADER_SMOOTH_NORMAL) {
-      sd->N = triangle_smooth_normal(kg, Ng, sd->prim, sd->u, sd->v);
+    if (sd->type == PRIMITIVE_TRIANGLE) {
+      /* smooth normal */
+      if (sd->shader & SHADER_SMOOTH_NORMAL) {
+        sd->N = triangle_smooth_normal(kg, Ng, sd->prim, sd->u, sd->v);
+
+        if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+          object_normal_transform_auto(kg, sd, &sd->N);
+        }
+      }
+
+      /* dPdu/dPdv */
+#ifdef __DPDU__
+      triangle_dPdudv(kg, sd->prim, &sd->dPdu, &sd->dPdv);
 
       if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
-        object_normal_transform_auto(kg, sd, &sd->N);
+        object_dir_transform_auto(kg, sd, &sd->dPdu);
+        object_dir_transform_auto(kg, sd, &sd->dPdv);
       }
-    }
-
-    /* dPdu/dPdv */
-#ifdef __DPDU__
-    triangle_dPdudv(kg, sd->prim, &sd->dPdu, &sd->dPdv);
-
-    if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
-      object_dir_transform_auto(kg, sd, &sd->dPdu);
-      object_dir_transform_auto(kg, sd, &sd->dPdv);
-    }
 #endif
+    }
+    else {
+#ifdef __DPDU__
+      sd->dPdu = zero_float3();
+      sd->dPdv = zero_float3();
+#endif
+    }
   }
   else {
+    if (lamp != LAMP_NONE) {
+      sd->lamp = lamp;
+    }
 #ifdef __DPDU__
     sd->dPdu = zero_float3();
     sd->dPdv = zero_float3();

intern/cycles/kernel/geom/triangle_intersect.h

@@ -142,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);
@@ -201,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