Merge branch 'master' into tmp-workbench-shader-create-infos

# Conflicts:
#	source/blender/draw/intern/shaders/common_hair_lib.glsl

CMakeLists.txt

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

@@ -19,10 +19,13 @@
 set(FREETYPE_EXTRA_ARGS
   -DCMAKE_RELEASE_POSTFIX:STRING=2ST
   -DCMAKE_DEBUG_POSTFIX:STRING=2ST_d
-  -DFT_DISABLE_BZIP2=ON
-  -DFT_DISABLE_HARFBUZZ=ON
-  -DFT_DISABLE_PNG=ON
-  -DFT_REQUIRE_BROTLI=ON
+  -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
   -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 --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:
+  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:
 )
 
 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.1)
+set(FREETYPE_VERSION 2.11.0)
 set(FREETYPE_URI http://prdownloads.sourceforge.net/freetype/freetype-${FREETYPE_VERSION}.tar.gz)
-set(FREETYPE_HASH bd4e3b007474319909a6b79d50908e85)
+set(FREETYPE_HASH cf09172322f6b50cf8f568bf8fe14bde)
 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.10.2)
-set(PYTHON_SHORT_VERSION 3.10)
-set(PYTHON_SHORT_VERSION_NO_DOTS 310)
+set(PYTHON_VERSION 3.9.7)
+set(PYTHON_SHORT_VERSION 3.9)
+set(PYTHON_SHORT_VERSION_NO_DOTS 39)
 set(PYTHON_URI https://www.python.org/ftp/python/${PYTHON_VERSION}/Python-${PYTHON_VERSION}.tar.xz)
-set(PYTHON_HASH 14e8c22458ed7779a1957b26cde01db9)
+set(PYTHON_HASH fddb060b483bc01850a3f412eea1d954)
 set(PYTHON_HASH_TYPE MD5)
 set(PYTHON_FILE Python-${PYTHON_VERSION}.tar.xz)
 
@@ -215,20 +215,18 @@ set(NANOVDB_HASH e7b9e863ec2f3b04ead171dec2322807)
 set(NANOVDB_HASH_TYPE MD5)
 set(NANOVDB_FILE nano-vdb-${NANOVDB_GIT_UID}.tar.gz)
 
-set(IDNA_VERSION 3.3)
-set(CHARSET_NORMALIZER_VERSION 2.0.10)
-set(URLLIB3_VERSION 1.26.8)
+set(IDNA_VERSION 3.2)
+set(CHARSET_NORMALIZER_VERSION 2.0.6)
+set(URLLIB3_VERSION 1.26.7)
 set(CERTIFI_VERSION 2021.10.8)
-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(REQUESTS_VERSION 2.26.0)
+set(CYTHON_VERSION 0.29.24)
+set(ZSTANDARD_VERSION 0.15.2 )
 
-set(NUMPY_VERSION 1.22.0)
-set(NUMPY_SHORT_VERSION 1.22)
+set(NUMPY_VERSION 1.21.2)
+set(NUMPY_SHORT_VERSION 1.21)
 set(NUMPY_URI https://github.com/numpy/numpy/releases/download/v${NUMPY_VERSION}/numpy-${NUMPY_VERSION}.zip)
-set(NUMPY_HASH 252de134862a27bd66705d29622edbfe)
+set(NUMPY_HASH 5638d5dae3ca387be562912312db842e)
 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.10.2"
-PYTHON_VERSION_SHORT="3.10"
-PYTHON_VERSION_MIN="3.9"
-PYTHON_VERSION_MEX="3.12"
+PYTHON_VERSION="3.9.7"
+PYTHON_VERSION_SHORT="3.9"
+PYTHON_VERSION_MIN="3.7"
+PYTHON_VERSION_MEX="3.11"
 PYTHON_VERSION_INSTALLED=$PYTHON_VERSION_SHORT
 PYTHON_FORCE_BUILD=false
 PYTHON_FORCE_REBUILD=false
 PYTHON_SKIP=false
 
 # Additional Python modules.
-PYTHON_IDNA_VERSION="3.3"
+PYTHON_IDNA_VERSION="3.2"
 PYTHON_IDNA_VERSION_MIN="2.0"
 PYTHON_IDNA_VERSION_MEX="4.0"
 PYTHON_IDNA_NAME="idna"
 
-PYTHON_CHARSET_NORMALIZER_VERSION="2.0.10"
+PYTHON_CHARSET_NORMALIZER_VERSION="2.0.6"
 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.8"
+PYTHON_URLLIB3_VERSION="1.26.7"
 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.27.1"
+PYTHON_REQUESTS_VERSION="2.23.0"
 PYTHON_REQUESTS_VERSION_MIN="2.0"
 PYTHON_REQUESTS_VERSION_MEX="3.0"
 PYTHON_REQUESTS_NAME="requests"
 
-PYTHON_ZSTANDARD_VERSION="0.16.0"
+PYTHON_ZSTANDARD_VERSION="0.15.2"
 PYTHON_ZSTANDARD_VERSION_MIN="0.15.2"
-PYTHON_ZSTANDARD_VERSION_MEX="0.20.0"
+PYTHON_ZSTANDARD_VERSION_MEX="0.16.0"
 PYTHON_ZSTANDARD_NAME="zstandard"
 
-PYTHON_NUMPY_VERSION="1.22.0"
+PYTHON_NUMPY_VERSION="1.21.2"
 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.17.0"
+OSL_VERSION="1.11.14.1"
 OSL_VERSION_SHORT="1.11"
 OSL_VERSION_MIN="1.11"
 OSL_VERSION_MEX="2.0"

build_files/cmake/Modules/FindFFmpeg.cmake

@@ -82,6 +82,4 @@ mark_as_advanced(
 
 unset(_ffmpeg_SEARCH_DIRS)
 unset(_ffmpeg_LIBRARIES)
-# 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)
+unset(_ffmpeg_INCLUDE_DIR)

build_files/cmake/Modules/FindOSL.cmake

@@ -76,7 +76,6 @@ 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
@@ -100,7 +99,6 @@ 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,14 +87,12 @@ 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,6 +33,14 @@ 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
@@ -85,24 +93,6 @@ 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)
@@ -121,57 +111,6 @@ 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)
@@ -180,25 +119,13 @@ 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>, 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()
+  # 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)
 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,7 +110,6 @@ ENDIF()
 
 MARK_AS_ADVANCED(
   OPENIMAGEDENOISE_INCLUDE_DIR
-  OPENIMAGEDENOISE_LIBRARY
 )
 
 FOREACH(COMPONENT ${_openimagedenoise_FIND_COMPONENTS})

build_files/cmake/Modules/FindOpenImageIO.cmake

@@ -48,8 +48,6 @@ FIND_LIBRARY(OPENIMAGEIO_LIBRARY
     lib64 lib
   )
 
-set(_openimageio_LIBRARIES ${OPENIMAGEIO_LIBRARY})
-
 FIND_FILE(OPENIMAGEIO_IDIFF
   NAMES
     idiff
@@ -59,47 +57,14 @@ 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_LIBRARIES OPENIMAGEIO_INCLUDE_DIR)
+    OPENIMAGEIO_LIBRARY OPENIMAGEIO_INCLUDE_DIR)
 
 IF(OPENIMAGEIO_FOUND)
-  SET(OPENIMAGEIO_LIBRARIES ${_openimageio_LIBRARIES})
+  SET(OPENIMAGEIO_LIBRARIES ${OPENIMAGEIO_LIBRARY})
   SET(OPENIMAGEIO_INCLUDE_DIRS ${OPENIMAGEIO_INCLUDE_DIR})
   IF(EXISTS ${OPENIMAGEIO_INCLUDE_DIR}/OpenImageIO/pugixml.hpp)
     SET(OPENIMAGEIO_PUGIXML_FOUND TRUE)
@@ -113,9 +78,7 @@ 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.10 CACHE STRING "Python Version (major and minor only)")
+SET(PYTHON_VERSION 3.9 CACHE STRING "Python Version (major and minor only)")
 MARK_AS_ADVANCED(PYTHON_VERSION)
 
 

build_files/cmake/platform/platform_apple.cmake

@@ -128,20 +128,25 @@ if(WITH_CODEC_SNDFILE)
 endif()
 
 if(WITH_PYTHON)
-  # Use precompiled libraries by default.
-  set(PYTHON_VERSION 3.10)
+  # we use precompiled libraries for py 3.9 and up by default
+  set(PYTHON_VERSION 3.9)
   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

@@ -109,14 +109,9 @@ if(NOT WITH_SYSTEM_FREETYPE)
   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}
-    # )
+    list(APPEND FREETYPE_LIBRARIES
+      ${BROTLI_LIBRARIES}
+    )
   endif()
 endif()
 
@@ -362,7 +357,6 @@ 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})

build_files/cmake/platform/platform_win32.cmake

@@ -55,10 +55,6 @@ 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)
@@ -269,6 +265,12 @@ 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}")
@@ -463,7 +465,7 @@ if(WITH_JACK)
 endif()
 
 if(WITH_PYTHON)
-  set(PYTHON_VERSION 3.10) # CACHE STRING)
+  set(PYTHON_VERSION 3.9) # 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: blender-v3.1-release
+        -   branch: master
             commit_id: HEAD
             path: release/scripts/addons
-        -   branch: blender-v3.1-release
+        -   branch: master
             commit_id: HEAD
             path: release/scripts/addons_contrib
-        -   branch: blender-v3.1-release
+        -   branch: master
             commit_id: HEAD
             path: release/datafiles/locale
-        -   branch: blender-v3.1-release
+        -   branch: master
             commit_id: HEAD
             path: source/tools
     svn:
         libraries:
             darwin-arm64:
-                branch: tags/blender-3.1-release
+                branch: trunk
                 commit_id: HEAD
                 path: lib/darwin_arm64
             darwin-x86_64:
-                branch: tags/blender-3.1-release
+                branch: trunk
                 commit_id: HEAD
                 path: lib/darwin
             linux-x86_64:
-                branch: tags/blender-3.1-release
+                branch: trunk
                 commit_id: HEAD
                 path: lib/linux_centos7_x86_64
             windows-amd64:
-                branch: tags/blender-3.1-release
+                branch: trunk
                 commit_id: HEAD
                 path: lib/win64_vc15
         tests:
-            branch: tags/blender-3.1-release
+            branch: trunk
             commit_id: HEAD
             path: lib/tests
         benchmarks:

build_files/windows/autodetect_msvc.cmd

@@ -3,6 +3,9 @@ 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,3 +1,4 @@
+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,6 +19,12 @@ 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,9 +37,15 @@ 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
+	)
 )
 
 if "%WITH_ASAN%"=="1" (

build_files/windows/detect_msvc2017.cmd

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

build_files/windows/parse_arguments.cmd

@@ -50,6 +50,14 @@ 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,3 +1,4 @@
+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
 

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 it's best to be aware of ID data-blocks in Blender since you will often find properties
+Before going any further its 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 than the tooltip, though some of the pages
+Being an API reference, this link often gives little more information then 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 it's an array of three floats.
+But you now know that you have to use ``.location`` and that its 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
-than access each :class:`bpy.types.ID` instance by name.
+then 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 it's excluded from the output.
+   zooming the view for example isn't so useful to repeat so its 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 script's functionality.
+for use in cases where you can alter the list order without breaking the scripts 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 this if you can avoid it, especially when writing data in a loop.
+   This is the slowest option, do **not** use 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 this to join a list of strings (the list may be temporary). In the following example, the strings are joined with
+   Use 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 be accessed with Python too which is very useful.
+Blender's operators are tools for users to access, that can access 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 it gives an error like this:
+When calling an operator 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 it's currently not used much, if you're interested to help improve the API
+   but its 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 it's not obvious why poll fails, e.g:
+   where its 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.)
-have been recalculated and are available to the script within the active view layer.
+has been recalculated and is available to the script within 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 it's done;
+While a script executes Blender waits for it to finish and is effectively locked until its 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.
 
-Tools that lock Blender in a loop redraw are highly discouraged
+When tools 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 modal operators are Transform, Painting, Fly Navigation and File Select.
+Examples of a 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 it's not the native storage format in Object-Mode.
+  since using BMesh gives some overhead because its 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 it's important that the script
+Further access to these will crash Blender so its important 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.
 
 
-It's better practice not to reference objects by names at all,
+Its 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, it's best to use a dictionary to maintain
+If you do need to use name references, its 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 it's best to avoid referencing data by name at all.
+so its best to avoid referencing data by name at all.
 
-This can be tricky at times and not even Blender handles this correctly in some cases
+This can be tricky at times and not even Blender handles this correctly in some case
 (when selecting the modifier object for e.g. you can't select between multiple objects with the same name),
-but it's still good to try avoiding these problems in this area.
+but its 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 that all strings in blend-files
+To avoid the problem for Python integration and script authors we have decided all strings in blend-files
 **must** be ``UTF-8``, ``ASCII`` compatible.
-This means assigning strings with different encodings to an object name, for instance, will raise an error.
+This means assigning strings with different encodings to an object names for instance will raise an error.
 
-Paths are an exception to this rule since the existence of non-UTF-8 paths on the user's file system cannot be ignored.
+Paths are an exception to this rule since the existence of non-UTF-8 paths on 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 it's less trouble to read it as binary data
+   when reading some input its 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 well obviously as all of their sub-data.
+instances (Object, Scene, Mesh, Light, etc.), as weel 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 it's also possible to use a script to make library data point to newly created local data,
+But its 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 script's path from a two line code block.
+This is the equivalent to running the script directly, referencing a scripts 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-file's directory to the search path
+this following example adds the current blend-files 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>`),
-it can be hard to track down the cause of crashes.
+can make it 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,8 +417,7 @@ 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_STRING = bpy.app.version_string
-BLENDER_VERSION_DOTS = "%d.%d" % (bpy.app.version[0], bpy.app.version[1])
+BLENDER_VERSION_DOTS = bpy.app.version_string
 
 if BLENDER_REVISION != "Unknown":
     # SHA1 Git hash
@@ -1725,11 +1724,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_STRING)
+    fw("project = 'Blender %s Python API'\n" % BLENDER_VERSION_DOTS)
     fw("master_doc = 'index'\n")
     fw("copyright = u'Blender Foundation'\n")
-    fw("version = '%s'\n" % BLENDER_VERSION_DOTS)
-    fw("release = '%s'\n" % BLENDER_VERSION_DOTS)
+    fw("version = '%s'\n" % BLENDER_VERSION_HASH)
+    fw("release = '%s'\n" % BLENDER_VERSION_HASH)
 
     # Quiet file not in table-of-contents warnings.
     fw("exclude_patterns = [\n")
@@ -1750,7 +1749,6 @@ 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")
@@ -1767,15 +1765,10 @@ 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")
@@ -2132,9 +2125,6 @@ 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

@@ -1,127 +0,0 @@
-/* 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

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

@@ -1,17 +0,0 @@
-<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
-	const AVCodec* aCodec = avcodec_find_decoder(m_formatCtx->streams[m_stream]->codecpar->codec_id);
+	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,7 +23,6 @@
 extern "C" {
 #include <libavcodec/avcodec.h>
 #include <libavformat/avio.h>
-#include <libavutil/channel_layout.h>
 }
 
 AUD_NAMESPACE_BEGIN
@@ -172,66 +171,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.");
 
-	const AVOutputFormat* outputFmt = m_formatCtx->oformat;
+	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.");
 	}
 
-	AVCodecID audio_codec = AV_CODEC_ID_NONE;
+	outputFmt->audio_codec = AV_CODEC_ID_NONE;
 
 	switch(codec)
 	{
 	case CODEC_AAC:
-		audio_codec = AV_CODEC_ID_AAC;
+		outputFmt->audio_codec = AV_CODEC_ID_AAC;
 		break;
 	case CODEC_AC3:
-		audio_codec = AV_CODEC_ID_AC3;
+		outputFmt->audio_codec = AV_CODEC_ID_AC3;
 		break;
 	case CODEC_FLAC:
-		audio_codec = AV_CODEC_ID_FLAC;
+		outputFmt->audio_codec = AV_CODEC_ID_FLAC;
 		break;
 	case CODEC_MP2:
-		audio_codec = AV_CODEC_ID_MP2;
+		outputFmt->audio_codec = AV_CODEC_ID_MP2;
 		break;
 	case CODEC_MP3:
-		audio_codec = AV_CODEC_ID_MP3;
+		outputFmt->audio_codec = AV_CODEC_ID_MP3;
 		break;
 	case CODEC_OPUS:
-		audio_codec = AV_CODEC_ID_OPUS;
+		outputFmt->audio_codec = AV_CODEC_ID_OPUS;
 		break;
 	case CODEC_PCM:
 		switch(specs.format)
 		{
 		case FORMAT_U8:
-			audio_codec = AV_CODEC_ID_PCM_U8;
+			outputFmt->audio_codec = AV_CODEC_ID_PCM_U8;
 			break;
 		case FORMAT_S16:
-			audio_codec = AV_CODEC_ID_PCM_S16LE;
+			outputFmt->audio_codec = AV_CODEC_ID_PCM_S16LE;
 			break;
 		case FORMAT_S24:
-			audio_codec = AV_CODEC_ID_PCM_S24LE;
+			outputFmt->audio_codec = AV_CODEC_ID_PCM_S24LE;
 			break;
 		case FORMAT_S32:
-			audio_codec = AV_CODEC_ID_PCM_S32LE;
+			outputFmt->audio_codec = AV_CODEC_ID_PCM_S32LE;
 			break;
 		case FORMAT_FLOAT32:
-			audio_codec = AV_CODEC_ID_PCM_F32LE;
+			outputFmt->audio_codec = AV_CODEC_ID_PCM_F32LE;
 			break;
 		case FORMAT_FLOAT64:
-			audio_codec = AV_CODEC_ID_PCM_F64LE;
+			outputFmt->audio_codec = AV_CODEC_ID_PCM_F64LE;
 			break;
 		default:
-			audio_codec = AV_CODEC_ID_NONE;
+			outputFmt->audio_codec = AV_CODEC_ID_NONE;
 			break;
 		}
 		break;
 	case CODEC_VORBIS:
-		audio_codec = AV_CODEC_ID_VORBIS;
+		outputFmt->audio_codec = AV_CODEC_ID_VORBIS;
 		break;
 	default:
-		audio_codec = AV_CODEC_ID_NONE;
+		outputFmt->audio_codec = AV_CODEC_ID_NONE;
 		break;
 	}
 
@@ -269,10 +268,10 @@ FFMPEGWriter::FFMPEGWriter(std::string filename, DeviceSpecs specs, Container fo
 
 	try
 	{
-		if(audio_codec == AV_CODEC_ID_NONE)
+		if(outputFmt->audio_codec == AV_CODEC_ID_NONE)
 			AUD_THROW(FileException, "File couldn't be written, audio codec not found with ffmpeg.");
 
-		const AVCodec* codec = avcodec_find_encoder(audio_codec);
+		AVCodec* codec = avcodec_find_encoder(outputFmt->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,13 +44,15 @@ else()
 endif()
 
 if(WITH_CYCLES_STANDALONE AND WITH_CYCLES_STANDALONE_GUI)
-  add_definitions(${GL_DEFINITIONS})
-  list(APPEND INC_SYS ${GLEW_INCLUDE_DIR} ${SDL2_INCLUDE_DIRS})
-  list(APPEND LIBRARIES ${CYCLES_GL_LIBRARIES} ${SDL2_LIBRARIES})
+  list(APPEND LIBRARIES ${GLUT_LIBRARIES})
 endif()
 
+list(APPEND LIBRARIES ${CYCLES_GL_LIBRARIES})
+
 # Common configuration.
 
+add_definitions(${GL_DEFINITIONS})
+
 include_directories(${INC})
 include_directories(SYSTEM ${INC_SYS})
 
@@ -64,18 +66,6 @@ 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)
 
@@ -90,10 +80,6 @@ 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,10 +40,11 @@
 #include "app/oiio_output_driver.h"
 
 #ifdef WITH_CYCLES_STANDALONE_GUI
-#  include "opengl/display_driver.h"
-#  include "opengl/window.h"
+#  include "util/view.h"
 #endif
 
+#include "app/cycles_xml.h"
+
 CCL_NAMESPACE_BEGIN
 
 struct Options {
@@ -129,13 +130,6 @@ static void session_init()
   options.output_pass = "combined";
   options.session = new Session(options.session_params, options.scene_params);
 
-#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));
@@ -145,7 +139,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(&window_redraw));
+    options.session->progress.set_update_callback(function_bind(&view_redraw));
 #endif
 
   /* load scene */
@@ -210,10 +204,10 @@ static void display_info(Progress &progress)
       sample_time,
       interactive.c_str());
 
-  window_display_info(str.c_str());
+  view_display_info(str.c_str());
 
   if (options.show_help)
-    window_display_help();
+    view_display_help();
 }
 
 static void display()
@@ -544,7 +538,7 @@ 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 */
-    window_main_loop(title.c_str(),
+    view_main_loop(title.c_str(),
                    options.width,
                    options.height,
                    session_init,

intern/cycles/app/opengl/display_driver.cpp

@@ -1,398 +0,0 @@
-/*
- * 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

@@ -1,130 +0,0 @@
-/*
- * 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

@@ -1,210 +0,0 @@
-/*
- * 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

@@ -1,58 +0,0 @@
-/*
- * 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

@@ -1,365 +0,0 @@
-/*
- * 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

@@ -1,48 +0,0 @@
-/*
- * 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 = (
@@ -901,27 +901,27 @@ class CyclesCameraSettings(bpy.types.PropertyGroup):
 
     fisheye_polynomial_k0: FloatProperty(
         name="Fisheye Polynomial K0",
-        description="Coefficient K0 of the lens polynomial",
+        description="Coefficient K0 of the lens polinomial",
         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 polynomial",
+        description="Coefficient K1 of the lens polinomial",
         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 polynomial",
+        description="Coefficient K2 of the lens polinomial",
         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 polynomial",
+        description="Coefficient K3 of the lens polinomial",
         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 polynomial",
+        description="Coefficient K4 of the lens polinomial",
         default=camera.default_fisheye_polynomial[4], precision=6, step=0.1, subtype='ANGLE',
     )
 
@@ -1374,12 +1374,6 @@ 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]:
@@ -1498,8 +1492,7 @@ 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 with macOS 12.2 or newer", icon='BLANK1')
-                col.label(text="or AMD with macOS 12.3 or newer", icon='BLANK1')
+                col.label(text="Requires Apple Silicon and macOS 12.0 or newer", icon='BLANK1')
             return
 
         for device in devices:
@@ -1526,12 +1519,6 @@ 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

@@ -669,7 +669,7 @@ class CYCLES_RENDER_PT_performance_acceleration_structure(CyclesButtonsPanel, Pa
 
     @classmethod
     def poll(cls, context):
-        return not use_optix(context) or use_multi_device(context)
+        return not use_optix(context) or has_multi_device(context)
 
     def draw(self, context):
         import _cycles

intern/cycles/blender/device.cpp

@@ -118,10 +118,6 @@ 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,12 +480,26 @@ 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
@@ -498,6 +512,9 @@ 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 {
@@ -543,30 +560,6 @@ 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)
@@ -633,12 +626,6 @@ 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();
@@ -838,8 +825,7 @@ 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 uint gl_vertex_buffer)
+                      const DrawTile &draw_tile)
 {
   if (!draw_tile.ready_to_draw()) {
     return;
@@ -848,9 +834,9 @@ static void draw_tile(const float2 &zoom,
   const GLTexture &texture = draw_tile.texture;
 
   DCHECK_NE(texture.gl_id, 0);
-  DCHECK_NE(gl_vertex_buffer, 0);
+  DCHECK_NE(draw_tile.gl_vertex_buffer, 0);
 
-  glBindBuffer(GL_ARRAY_BUFFER, gl_vertex_buffer);
+  glBindBuffer(GL_ARRAY_BUFFER, draw_tile.gl_vertex_buffer);
 
   /* Draw at the parameters for which the texture has been updated for. This allows to always draw
    * texture during bordered-rendered camera view without flickering. The validness of the display
@@ -970,14 +956,10 @@ void BlenderDisplayDriver::draw(const Params &params)
   glEnableVertexAttribArray(texcoord_attribute);
   glEnableVertexAttribArray(position_attribute);
 
-  draw_tile(zoom_,
-            texcoord_attribute,
-            position_attribute,
-            tiles_->current_tile.tile,
-            tiles_->gl_vertex_buffer);
+  draw_tile(zoom_, texcoord_attribute, position_attribute, tiles_->current_tile.tile);
 
   for (const DrawTile &tile : tiles_->finished_tiles.tiles) {
-    draw_tile(zoom_, texcoord_attribute, position_attribute, tile, tiles_->gl_vertex_buffer);
+    draw_tile(zoom_, texcoord_attribute, position_attribute, tile);
   }
 
   display_shader_->unbind();
@@ -1080,7 +1062,6 @@ 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,13 +506,8 @@ void BlenderSession::render_frame_finish()
   session->set_output_driver(nullptr);
   session->full_buffer_written_cb = function_null;
 
-  /* 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. */
+  /* The display driver holds OpenGL resources which belong to an OpenGL context held by the render
+   * engine on Blender side. Force destruction of those resources. */
   display_driver_ = nullptr;
   session->set_display_driver(nullptr);
 

intern/cycles/blender/shader.cpp

@@ -45,8 +45,7 @@ typedef map<string, ConvertNode *> ProxyMap;
 
 void BlenderSync::find_shader(BL::ID &id, array<Node *> &used_shaders, Shader *default_shader)
 {
-  Shader *synced_shader = (id) ? shader_map.find(id) : nullptr;
-  Shader *shader = (synced_shader) ? synced_shader : default_shader;
+  Shader *shader = (id) ? shader_map.find(id) : default_shader;
 
   used_shaders.push_back_slow(shader);
   shader->tag_used(scene);

intern/cycles/bvh/embree.cpp

@@ -61,26 +61,6 @@ static_assert(Object::MAX_MOTION_STEPS == Geometry::MAX_MOTION_STEPS,
 
 #  define IS_HAIR(x) (x & 1)
 
-/* This gets called by Embree at every valid ray/object intersection.
- * Things like recording subsurface or shadow hits for later evaluation
- * as well as filtering for volume objects happen here.
- * Cycles' own BVH does that directly inside the traversal calls.
- */
-static void rtc_filter_intersection_func(const RTCFilterFunctionNArguments *args)
-{
-  /* Current implementation in Cycles assumes only single-ray intersection queries. */
-  assert(args->N == 1);
-
-  RTCHit *hit = (RTCHit *)args->hit;
-  CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt;
-  const KernelGlobalsCPU *kg = ctx->kg;
-  const Ray *cray = ctx->ray;
-
-  if (kernel_embree_is_self_intersection(kg, hit, cray)) {
-    *args->valid = 0;
-  }
-}
-
 /* This gets called by Embree at every valid ray/object intersection.
  * Things like recording subsurface or shadow hits for later evaluation
  * as well as filtering for volume objects happen here.
@@ -95,16 +75,12 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
   RTCHit *hit = (RTCHit *)args->hit;
   CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt;
   const KernelGlobalsCPU *kg = ctx->kg;
-  const Ray *cray = ctx->ray;
 
   switch (ctx->type) {
     case CCLIntersectContext::RAY_SHADOW_ALL: {
       Intersection current_isect;
       kernel_embree_convert_hit(kg, ray, hit, &current_isect);
-      if (intersection_skip_self_shadow(cray->self, current_isect.object, current_isect.prim)) {
-        *args->valid = 0;
-        return;
-      }
+
       /* If no transparent shadows or max number of hits exceeded, all light is blocked. */
       const int flags = intersection_get_shader_flags(kg, current_isect.prim, current_isect.type);
       if (!(flags & (SD_HAS_TRANSPARENT_SHADOW)) || ctx->num_hits >= ctx->max_hits) {
@@ -184,10 +160,6 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
           break;
         }
       }
-      if (intersection_skip_self_local(cray->self, current_isect.prim)) {
-        *args->valid = 0;
-        return;
-      }
 
       /* No intersection information requested, just return a hit. */
       if (ctx->max_hits == 0) {
@@ -253,11 +225,6 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
       if (ctx->num_hits < ctx->max_hits) {
         Intersection current_isect;
         kernel_embree_convert_hit(kg, ray, hit, &current_isect);
-        if (intersection_skip_self(cray->self, current_isect.object, current_isect.prim)) {
-          *args->valid = 0;
-          return;
-        }
-
         Intersection *isect = &ctx->isect_s[ctx->num_hits];
         ++ctx->num_hits;
         *isect = current_isect;
@@ -269,15 +236,12 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
         }
         /* This tells Embree to continue tracing. */
         *args->valid = 0;
-      }
         break;
       }
+    }
     case CCLIntersectContext::RAY_REGULAR:
     default:
-      if (kernel_embree_is_self_intersection(kg, hit, cray)) {
-        *args->valid = 0;
-        return;
-      }
+      /* Nothing to do here. */
       break;
   }
 }
@@ -293,14 +257,6 @@ static void rtc_filter_func_backface_cull(const RTCFilterFunctionNArguments *arg
     *args->valid = 0;
     return;
   }
-
-  CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt;
-  const KernelGlobalsCPU *kg = ctx->kg;
-  const Ray *cray = ctx->ray;
-
-  if (kernel_embree_is_self_intersection(kg, hit, cray)) {
-    *args->valid = 0;
-  }
 }
 
 static void rtc_filter_occluded_func_backface_cull(const RTCFilterFunctionNArguments *args)
@@ -471,7 +427,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, (size_t)RTC_MAX_TIME_STEP_COUNT);
+  const size_t num_motion_steps = min(num_object_motion_steps, 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);
@@ -522,7 +478,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, (size_t)RTC_MAX_TIME_STEP_COUNT);
+  num_motion_steps = min(num_motion_steps, RTC_MAX_TIME_STEP_COUNT);
 
   const size_t num_triangles = mesh->num_triangles();
 
@@ -549,7 +505,6 @@ void BVHEmbree::add_triangles(const Object *ob, const Mesh *mesh, int i)
 
   rtcSetGeometryUserData(geom_id, (void *)prim_offset);
   rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func);
-  rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_intersection_func);
   rtcSetGeometryMask(geom_id, ob->visibility_for_tracing());
 
   rtcCommitGeometry(geom_id);
@@ -775,7 +730,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, (size_t)RTC_MAX_TIME_STEP_COUNT);
+  num_motion_steps = min(num_motion_steps, RTC_MAX_TIME_STEP_COUNT);
 
   const size_t num_curves = hair->num_curves();
   size_t num_segments = 0;
@@ -812,7 +767,6 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
 
   rtcSetGeometryUserData(geom_id, (void *)prim_offset);
   if (hair->curve_shape == CURVE_RIBBON) {
-    rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_intersection_func);
     rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func);
   }
   else {

intern/cycles/bvh/params.h

@@ -142,7 +142,6 @@ 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,22 +491,26 @@ else()
 endif()
 
 ###########################################################################
-# SDL
+# GLUT
 ###########################################################################
 
 if(WITH_CYCLES_STANDALONE AND WITH_CYCLES_STANDALONE_GUI)
-  # We can't use the version from the Blender precompiled libraries because
-  # it does not include the video subsystem.
-  find_package(SDL2)
+  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)
 
-  if(NOT SDL2_FOUND)
+    if(NOT GLUT_FOUND)
       set(WITH_CYCLES_STANDALONE_GUI OFF)
-    message(STATUS "SDL not found, disabling Cycles standalone GUI")
+      message(STATUS "GLUT not found, disabling Cycles standalone GUI")
+    endif()
   endif()
 
   include_directories(
     SYSTEM
-    ${SDL2_INCLUDE_DIRS}
+    ${GLUT_INCLUDE_DIR}
   )
 endif()
 
@@ -555,10 +559,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,7 +328,6 @@ 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) {
@@ -336,7 +335,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), size_t(0));
+        int cpu_threads = max(orig_cpu_threads - (subdevices.size() - 1), 0);
 
         VLOG(1) << "CPU render threads reduced from " << orig_cpu_threads << " to " << cpu_threads
                 << ", to dedicate to GPU.";
@@ -375,7 +374,6 @@ 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,7 +79,6 @@ 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;
@@ -97,7 +96,6 @@ 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,6 +25,8 @@
 
 #  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, size_t(1));
+    data_elements = max(device_type_traits<T>::num_elements, 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((size_t)1, ob->get_motion().size());
+        num_motion_transforms += max(1, ob->get_motion().size());
       }
       else {
         num_motion_transforms++;

intern/cycles/device/metal/device.mm

@@ -39,20 +39,33 @@ 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)
 {
-  auto usable_devices = MetalInfo::get_usable_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);
   /* Devices are numbered consecutively across platforms. */
   set<string> unique_ids;
   int device_index = 0;
-  for (id<MTLDevice> &device : usable_devices) {
+  for (MetalPlatformDevice &device : usable_devices) {
     /* Compute unique ID for persistent user preferences. */
-    string device_name = [device.name UTF8String];
+    const string &device_name = device.device_name;
     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", device_index);
+      id += string_printf("_ID_%d", num_devices);
     }
     unique_ids.insert(id);
 
@@ -81,13 +94,15 @@ void device_metal_info(vector<DeviceInfo> &devices)
 string device_metal_capabilities()
 {
   string result = "";
-  auto allDevices = MTLCopyAllDevices();
-  uint32_t num_devices = allDevices.count;
+  string error_msg = "";
+  uint32_t num_devices = 0;
+  assert(device_metal_get_num_devices_safe(&num_devices));
   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,10 +53,16 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
   mtlDevId = info.num;
 
   /* select chosen device */
-  auto usable_devices = MetalInfo::get_usable_devices();
+  vector<MetalPlatformDevice> usable_devices;
+  MetalInfo::get_usable_devices(&usable_devices);
+  if (usable_devices.size() == 0) {
+    set_error("Metal: no devices found.");
+    return;
+  }
   assert(mtlDevId < usable_devices.size());
-  mtlDevice = usable_devices[mtlDevId];
-  device_name = [mtlDevice.name UTF8String];
+  MetalPlatformDevice &platform_device = usable_devices[mtlDevId];
+  mtlDevice = platform_device.device_id;
+  device_name = platform_device.device_name;
   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());
@@ -94,7 +100,6 @@ 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);
   }
@@ -450,15 +455,8 @@ 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;
     }
@@ -492,7 +490,7 @@ MetalDevice::MetalMem *MetalDevice::generic_alloc(device_memory &mem)
   mmem->mtlBuffer = metal_buffer;
   mmem->offset = 0;
   mmem->size = size;
-  if (options != MTLResourceStorageModePrivate) {
+  if (mem.type != MEM_DEVICE_ONLY) {
     mmem->hostPtr = [metal_buffer contents];
   }
   else {
@@ -761,17 +759,6 @@ 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,15 +59,10 @@ bool MetalDeviceKernel::load(MetalDevice *device,
   }
 
   bool use_binary_archive = true;
-  if (device->device_vendor == METAL_GPU_APPLE) {
-    /* Workaround for T94142: Cycles Metal crash with simultaneous viewport and final render */
+  if (getenv("CYCLES_METAL_DISABLE_BINARY_ARCHIVES")) {
     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,10 +36,33 @@ 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 vector<id<MTLDevice>> const &get_usable_devices();
+  static bool device_version_check(id<MTLDevice> device);
+  static void get_usable_devices(vector<MetalPlatformDevice> *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,45 +43,83 @@ MetalGPUVendor MetalInfo::get_vendor_from_device_name(string const &device_name)
   return METAL_GPU_UNKNOWN;
 }
 
-vector<id<MTLDevice>> const &MetalInfo::get_usable_devices()
+bool MetalInfo::device_version_check(id<MTLDevice> device)
 {
-  static vector<id<MTLDevice>> usable_devices;
-  static bool already_enumerated = false;
+  /* 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]);
 
-  if (already_enumerated) {
-    return usable_devices;
+    /* Metal Cycles works on Apple Silicon GPUs at present */
+    return (vendor == METAL_GPU_APPLE);
   }
 
-  metal_printf("Usable Metal devices:\n");
-  for (id<MTLDevice> device in MTLCopyAllDevices()) {
-    const char *device_name = [device.name UTF8String];
+  return false;
+}
 
-    MetalGPUVendor vendor = get_vendor_from_device_name(device_name);
-    bool usable = false;
+void MetalInfo::get_usable_devices(vector<MetalPlatformDevice> *usable_devices)
+{
+  static bool first_time = true;
+#  define FIRST_VLOG(severity) \
+    if (first_time) \
+    VLOG(severity)
 
-    if (@available(macos 12.2, *)) {
-      usable |= (vendor == METAL_GPU_APPLE);
+  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.3, *)) {
-      usable |= (vendor == METAL_GPU_AMD);
+    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 (usable) {
-      metal_printf("- %s\n", device_name);
-      [device retain];
-      usable_devices.push_back(device);
+    if (!device_version_check(device)) {
+      FIRST_VLOG(2) << "Ignoring device " << device_name << " due to too old compiler version.";
+      continue;
     }
-    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));
   }
-  }
-  if (usable_devices.empty()) {
-    metal_printf("   No usable Metal devices found\n");
-  }
-  already_enumerated = true;
+  first_time = false;
+}
 
-  return usable_devices;
+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;
+  }
+  return num_devices;
+}
+
+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;
 }
 
 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 = 8;
+  pipeline_options.numPayloadValues = 6;
   pipeline_options.numAttributeValues = 2; /* u, v */
   pipeline_options.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE;
   pipeline_options.pipelineLaunchParamsVariableName = "__params"; /* See globals.h */
@@ -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(), (size_t)2) - 2;
+          const size_t motion_keys = max(ob->get_motion().size(), 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(), (size_t)2) - 2;
+        size_t motion_keys = max(ob->get_motion().size(), 2) - 2;
         size_t motion_transform_size = sizeof(OptixSRTMotionTransform) +
                                        motion_keys * sizeof(OptixSRTData);
 

intern/cycles/integrator/denoiser.cpp

@@ -125,41 +125,20 @@ 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,
-                                                 DenoiserType denoiser_type)
+                                                 const uint device_type_mask)
 {
   const vector<DeviceInfo> device_infos = Device::available_devices(device_type_mask);
   if (device_infos.empty()) {
     return nullptr;
   }
 
-  const DeviceInfo denoiser_device_info = find_best_denoiser_device_info(device_infos,
-                                                                         denoiser_type);
-  if (denoiser_device_info.type == DEVICE_NONE) {
-    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();
 
   unique_ptr<Device> denoiser_device(
       Device::create(denoiser_device_info, path_trace_device->stats, path_trace_device->profiler));
@@ -207,8 +186,7 @@ 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, params_.type);
+  local_denoiser_device_ = create_denoiser_device(path_trace_device_, device_type_mask);
   denoiser_device_ = local_denoiser_device_.get();
 
   return denoiser_device_;

intern/cycles/integrator/denoiser_oidn.cpp

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

intern/cycles/kernel/bvh/embree.h

@@ -22,8 +22,6 @@
 #include "kernel/device/cpu/compat.h"
 #include "kernel/device/cpu/globals.h"
 
-#include "kernel/bvh/util.h"
-
 #include "util/vector.h"
 
 CCL_NAMESPACE_BEGIN
@@ -40,9 +38,6 @@ struct CCLIntersectContext {
   KernelGlobals kg;
   RayType type;
 
-  /* For avoiding self intersections */
-  const Ray *ray;
-
   /* for shadow rays */
   Intersection *isect_s;
   uint max_hits;
@@ -61,7 +56,6 @@ struct CCLIntersectContext {
   {
     kg = kg_;
     type = type_;
-    ray = NULL;
     max_hits = 1;
     num_hits = 0;
     num_recorded_hits = 0;
@@ -108,34 +102,7 @@ ccl_device_inline void kernel_embree_setup_rayhit(const Ray &ray,
 {
   kernel_embree_setup_ray(ray, rayhit.ray, visibility);
   rayhit.hit.geomID = RTC_INVALID_GEOMETRY_ID;
-  rayhit.hit.instID[0] = RTC_INVALID_GEOMETRY_ID;
-}
-
-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;
+  rayhit.hit.primID = RTC_INVALID_GEOMETRY_ID;
 }
 
 ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg,

intern/cycles/kernel/bvh/local.h

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

intern/cycles/kernel/bvh/shadow_all.h

@@ -160,9 +160,6 @@ 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,6 +133,8 @@ 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);
 
@@ -140,22 +142,26 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
                                             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;
             }
 #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);
+
                 if (motion_triangle_intersect(kg,
                                               isect,
                                               P,
@@ -170,6 +176,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
                     return true;
                 }
+              }
               break;
             }
 #endif /* BVH_FEATURE(BVH_MOTION) */
@@ -178,13 +185,19 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
             case PRIMITIVE_MOTION_CURVE_THICK:
             case PRIMITIVE_CURVE_RIBBON:
             case PRIMITIVE_MOTION_CURVE_RIBBON: {
+              for (; prim_addr < prim_addr2; prim_addr++) {
                 if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
                   const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
                   if (ray->time < prim_time.x || ray->time > prim_time.y) {
-                    break;
+                    continue;
                   }
                 }
 
+                const int 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);
@@ -193,19 +206,26 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
                     return true;
                 }
+              }
               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++) {
                 if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
                   const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
                   if (ray->time < prim_time.x || ray->time > prim_time.y) {
-                    break;
+                    continue;
                   }
                 }
 
+                const int 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);
@@ -214,12 +234,12 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
                     return true;
                 }
+              }
               break;
             }
 #endif /* BVH_FEATURE(BVH_POINTCLOUD) */
           }
         }
-        }
         else {
           /* instance push */
           object = kernel_tex_fetch(__prim_object, -prim_addr - 1);

intern/cycles/kernel/bvh/util.h

@@ -21,22 +21,54 @@ CCL_NAMESPACE_BEGIN
 /* Ray offset to avoid self intersection.
  *
  * This function should be used to compute a modified ray start position for
- * rays leaving from a surface. This is from "A Fast and Robust Method for Avoiding
- * Self-Intersection" see https://research.nvidia.com/publication/2019-03_A-Fast-and
- */
+ * rays leaving from a surface. */
+
 ccl_device_inline float3 ray_offset(float3 P, float3 Ng)
 {
-  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));
+#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;
 
-  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);
+  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
 }
 
 #if defined(__KERNEL_CPU__)
@@ -195,25 +227,4 @@ ccl_device_inline float intersection_curve_shadow_transparency(KernelGlobals kg,
   return (1.0f - u) * f0 + u * f1;
 }
 
-ccl_device_inline bool intersection_skip_self(ccl_private const RaySelfPrimitives &self,
-                                              const int object,
-                                              const int prim)
-{
-  return (self.prim == prim) && (self.object == object);
-}
-
-ccl_device_inline bool intersection_skip_self_shadow(ccl_private const RaySelfPrimitives &self,
-                                                     const int object,
-                                                     const int prim)
-{
-  return ((self.prim == prim) && (self.object == object)) ||
-         ((self.light_prim == prim) && (self.light_object == object));
-}
-
-ccl_device_inline bool intersection_skip_self_local(ccl_private const RaySelfPrimitives &self,
-                                                    const int prim)
-{
-  return (self.prim == prim);
-}
-
 CCL_NAMESPACE_END

intern/cycles/kernel/bvh/volume.h

@@ -144,9 +144,6 @@ 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) {
@@ -167,9 +164,6 @@ 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,9 +147,6 @@ 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;
@@ -191,9 +188,6 @@ 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,11 +295,8 @@ 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)
@@ -313,11 +310,8 @@ 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)
@@ -328,11 +322,8 @@ 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)
@@ -344,11 +335,8 @@ 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)
@@ -360,11 +348,8 @@ 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)
@@ -406,11 +391,8 @@ 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)
@@ -442,11 +424,8 @@ 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,26 +31,43 @@ CCL_NAMESPACE_BEGIN
 #  define GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE 512
 #endif
 
-#ifndef __KERNEL_METAL__
-template<uint blocksize, typename IsActiveOp>
-__device__
-#endif
-    void
-    gpu_parallel_active_index_array_impl(const uint num_states,
+#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,
-#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)
-{
+                          IsActiveOp is_active_op)
+  {
+    const uint state_index = global_index;
 #else
+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)
 {
   extern ccl_gpu_shared int warp_offset[];
@@ -62,10 +79,10 @@ __device__
   const uint num_warps = blocksize / ccl_gpu_warp_size;
 
   const uint state_index = ccl_gpu_block_idx_x * blocksize + thread_index;
+#endif
 
     /* Test if state corresponding to this thread is active. */
     const uint is_active = (state_index < num_states) ? is_active_op(state_index) : 0;
-#endif
 
     /* For each thread within a warp compute how many other active states precede it. */
     const uint thread_offset = popcount(ccl_gpu_ballot(is_active) &
@@ -100,33 +117,23 @@ __device__
       const uint block_offset = warp_offset[num_warps];
       indices[block_offset + warp_offset[warp_index] + thread_offset] = state_index;
     }
-}
+  }
 
 #ifdef __KERNEL_METAL__
+}; /* end class ActiveIndexContext */
 
-#  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, \
+/* 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)
-
-#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)
-
+                       simdgroup_offset) \
+        .active_index_array
 #endif
 
 CCL_NAMESPACE_END

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

@@ -40,27 +40,6 @@ struct TriangleIntersectionResult
 
 enum { METALRT_HIT_TRIANGLE, METALRT_HIT_BOUNDING_BOX };
 
-ccl_device_inline bool intersection_skip_self(ray_data const RaySelfPrimitives& self,
-                                              const int object,
-                                              const int prim)
-{
-  return (self.prim == prim) && (self.object == object);
-}
-
-ccl_device_inline bool intersection_skip_self_shadow(ray_data const RaySelfPrimitives& self,
-                                                     const int object,
-                                                     const int prim)
-{
-  return ((self.prim == prim) && (self.object == object)) ||
-         ((self.light_prim == prim) && (self.light_object == object));
-}
-
-ccl_device_inline bool intersection_skip_self_local(ray_data const RaySelfPrimitives& self,
-                                                    const int prim)
-{
-  return (self.prim == prim);
-}
-
 template<typename TReturn, uint intersection_type>
 TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
                           ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload,
@@ -74,8 +53,8 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
 #ifdef __BVH_LOCAL__
   uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
 
-  if ((object != payload.local_object) || intersection_skip_self_local(payload.self, prim)) {
-    /* Only intersect with matching object and skip self-intersecton. */
+  if (object != payload.local_object) {
+    /* Only intersect with matching object */
     result.accept = false;
     result.continue_search = true;
     return result;
@@ -187,11 +166,6 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
   }
 #  endif
 
-  if (intersection_skip_self_shadow(payload.self, object, prim)) {
-    /* continue search */
-    return true;
-  }
-
   float u = 0.0f, v = 0.0f;
   int type = 0;
   if (intersection_type == METALRT_HIT_TRIANGLE) {
@@ -348,35 +322,21 @@ inline TReturnType metalrt_visibility_test(constant KernelParamsMetal &launch_pa
   }
 #  endif
 
-  uint visibility = payload.visibility;
 #  ifdef __VISIBILITY_FLAG__
+  uint visibility = payload.visibility;
   if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
     result.accept = false;
     result.continue_search = true;
     return result;
   }
-#  endif
 
   /* Shadow ray early termination. */
   if (visibility & PATH_RAY_SHADOW_OPAQUE) {
-    if (intersection_skip_self_shadow(payload.self, object, prim)) {
-      result.accept = 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,11 +45,6 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
   return pointer_unpack_from_uint<T>(optixGetPayload_2(), optixGetPayload_3());
 }
 
-template<typename T> ccl_device_forceinline T *get_payload_ptr_6()
-{
-  return (T *)(((uint64_t)optixGetPayload_7() << 32) | optixGetPayload_6());
-}
-
 ccl_device_forceinline int get_object_id()
 {
 #ifdef __OBJECT_MOTION__
@@ -116,12 +111,6 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
     return optixIgnoreIntersection();
   }
 
-  const int prim = optixGetPrimitiveIndex();
-  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
-  if (intersection_skip_self_local(ray->self, prim)) {
-    return optixIgnoreIntersection();
-  }
-
   const uint max_hits = optixGetPayload_5();
   if (max_hits == 0) {
     /* Special case for when no hit information is requested, just report that something was hit */
@@ -160,6 +149,8 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
     local_isect->num_hits = 1;
   }
 
+  const int prim = optixGetPrimitiveIndex();
+
   Intersection *isect = &local_isect->hits[hit];
   isect->t = optixGetRayTmax();
   isect->prim = prim;
@@ -194,11 +185,6 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
   }
 #  endif
 
-  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
-  if (intersection_skip_self_shadow(ray->self, object, prim)) {
-    return optixIgnoreIntersection();
-  }
-
   float u = 0.0f, v = 0.0f;
   int type = 0;
   if (optixIsTriangleHit()) {
@@ -328,12 +314,6 @@ extern "C" __global__ void __anyhit__kernel_optix_volume_test()
   if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
     return optixIgnoreIntersection();
   }
-
-  const int prim = optixGetPrimitiveIndex();
-  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
-  if (intersection_skip_self(ray->self, object, prim)) {
-    return optixIgnoreIntersection();
-  }
 }
 
 extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
@@ -350,31 +330,18 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
 #  endif
 #endif
 
+#ifdef __VISIBILITY_FLAG__
   const uint object = get_object_id();
   const uint visibility = optixGetPayload_4();
-#ifdef __VISIBILITY_FLAG__
   if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
     return optixIgnoreIntersection();
   }
-#endif
 
-  const int prim = optixGetPrimitiveIndex();
-  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
-
-  if (visibility & PATH_RAY_SHADOW_OPAQUE) {
-    if (intersection_skip_self_shadow(ray->self, object, prim)) {
-      return optixIgnoreIntersection();
-    }
-    else {
   /* Shadow ray early termination. */
+  if (visibility & PATH_RAY_SHADOW_OPAQUE) {
     return optixTerminateRay();
   }
-  }
-  else {
-    if (intersection_skip_self(ray->self, object, prim)) {
-      return optixIgnoreIntersection();
-    }
-  }
+#endif
 }
 
 extern "C" __global__ void __closesthit__kernel_optix_hit()

intern/cycles/kernel/geom/motion_triangle_intersect.h

@@ -29,19 +29,46 @@
 
 CCL_NAMESPACE_BEGIN
 
-/**
- * Use the barycentric coordinates to get the intersection location
+/* 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.
  */
-ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg,
+
+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,
-                                                       const float u,
-                                                       const float v,
                                                 float3 verts[3])
 {
-  float w = 1.0f - u - v;
-  float3 P = u * verts[0] + v * verts[1] + w * verts[2];
+#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;
 
   if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
@@ -49,8 +76,71 @@ ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg,
   }
 
   return P;
+#else
+  return P + D * t;
+#endif
 }
 
+/* Same as above, except that t is assumed to be in object space
+ * for instancing.
+ */
+
+#ifdef __BVH_LOCAL__
+#  if defined(__KERNEL_CUDA__) && (defined(i386) || defined(_M_IX86))
+ccl_device_noinline
+#  else
+ccl_device_inline
+#  endif
+    float3
+    motion_triangle_refine_local(KernelGlobals kg,
+                                 ccl_private ShaderData *sd,
+                                 float3 P,
+                                 float3 D,
+                                 float t,
+                                 const int isect_object,
+                                 const int isect_prim,
+                                 float3 verts[3])
+{
+#  if defined(__KERNEL_GPU_RAYTRACING__)
+  /* t is always in world space with OptiX and MetalRT. */
+  return motion_triangle_refine(kg, sd, P, D, t, isect_object, isect_prim, verts);
+#  else
+#    ifdef __INTERSECTION_REFINE__
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+    const Transform tfm = object_get_inverse_transform(kg, sd);
+
+    P = transform_point(&tfm, P);
+    D = transform_direction(&tfm, D);
+    D = normalize(D);
+  }
+
+  P = P + D * t;
+
+  /* compute refined intersection distance */
+  const float3 e1 = verts[0] - verts[2];
+  const float3 e2 = verts[1] - verts[2];
+  const float3 s1 = cross(D, e2);
+
+  const float invdivisor = 1.0f / dot(s1, e1);
+  const float3 d = P - verts[2];
+  const float3 s2 = cross(d, e1);
+  float rt = dot(e2, s2) * invdivisor;
+
+  P = P + D * rt;
+
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+    const Transform tfm = object_get_transform(kg, sd);
+    P = transform_point(&tfm, P);
+  }
+
+  return P;
+#    else  /* __INTERSECTION_REFINE__ */
+  return P + D * t;
+#    endif /* __INTERSECTION_REFINE__ */
+#  endif
+}
+#endif /* __BVH_LOCAL__ */
+
 /* Ray intersection. We simply compute the vertex positions at the given ray
  * time and do a ray intersection with the resulting triangle.
  */

intern/cycles/kernel/geom/motion_triangle_shader.h

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

intern/cycles/kernel/geom/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_point_from_uv(kg, sd, isect->object, isect->prim, isect->u, isect->v);
+      sd->P = triangle_refine(kg, sd, ray->P, ray->D, isect->t, isect->object, isect->prim);
       sd->Ng = Ng;
       sd->N = Ng;
 
@@ -190,6 +190,10 @@ 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) {
@@ -223,16 +227,6 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals kg,
 #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();
 #endif
   }
 

intern/cycles/kernel/geom/triangle_intersect.h

@@ -142,23 +142,58 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg,
 }
 #endif /* __BVH_LOCAL__ */
 
-/**
- * Use the barycentric coordinates to get the intersection location
+/* 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
  */
-ccl_device_inline float3 triangle_point_from_uv(KernelGlobals kg,
+
+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,
-                                                const float u,
-                                                const float v)
+                                         const int isect_prim)
 {
+#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);
-  float w = 1.0f - u - v;
-
-  float3 P = u * tri_a + v * tri_b + w * tri_c;
+  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;
+  }
 
   if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
@@ -166,6 +201,65 @@ ccl_device_inline float3 triangle_point_from_uv(KernelGlobals kg,
   }
 
   return P;
+#else
+  return P + D * t;
+#endif
+}
+
+/* Same as above, except that t is assumed to be in object space for
+ * instancing.
+ */
+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,12 +328,6 @@ 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,10 +156,7 @@ 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,10 +38,7 @@ 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;
 
@@ -71,7 +68,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 = stack_sd->P;
+    volume_ray.P = ray_offset(stack_sd->P, -stack_sd->Ng);
     if (volume_ray.t != FLT_MAX) {
       volume_ray.D = normalize_len(to_P - volume_ray.P, &volume_ray.t);
     }
@@ -94,10 +91,6 @@ 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;
 
@@ -210,7 +203,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
     }
 
     /* Move ray forward. */
-    volume_ray.P = stack_sd->P;
+    volume_ray.P = ray_offset(stack_sd->P, -stack_sd->Ng);
     ++step;
   }
 #endif

intern/cycles/kernel/integrator/shade_light.h

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

intern/cycles/kernel/integrator/shade_shadow.h

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

intern/cycles/kernel/integrator/shade_surface.h

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

intern/cycles/kernel/integrator/shade_volume.h

@@ -791,36 +791,22 @@ ccl_device_forceinline void integrate_volume_direct_light(
 
   /* Write shadow ray and associated state to global memory. */
   integrator_state_write_shadow_ray(kg, shadow_state, &ray);
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
 
   /* Copy state from main path to shadow path. */
   const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);
   const uint16_t transparent_bounce = INTEGRATOR_STATE(state, path, transparent_bounce);
   uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
   shadow_flag |= (is_light) ? PATH_RAY_SHADOW_FOR_LIGHT : 0;
+  shadow_flag |= (shadow_flag & PATH_RAY_ANY_PASS) ? 0 : PATH_RAY_VOLUME_PASS;
   const float3 throughput_phase = throughput * bsdf_eval_sum(&phase_eval);
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
-    packed_float3 pass_diffuse_weight;
-    packed_float3 pass_glossy_weight;
-
-    if (shadow_flag & PATH_RAY_ANY_PASS) {
-      /* Indirect bounce, use weights from earlier surface or volume bounce. */
-      pass_diffuse_weight = INTEGRATOR_STATE(state, path, pass_diffuse_weight);
-      pass_glossy_weight = INTEGRATOR_STATE(state, path, pass_glossy_weight);
-    }
-    else {
-      /* Direct light, no diffuse/glossy distinction needed for volumes. */
-      shadow_flag |= PATH_RAY_VOLUME_PASS;
-      pass_diffuse_weight = packed_float3(one_float3());
-      pass_glossy_weight = packed_float3(zero_float3());
-    }
-
+    const packed_float3 pass_diffuse_weight = (bounce == 0) ?
+                                                  packed_float3(one_float3()) :
+                                                  INTEGRATOR_STATE(
+                                                      state, path, pass_diffuse_weight);
     INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = pass_diffuse_weight;
-    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = pass_glossy_weight;
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = zero_float3();
   }
 
   INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
@@ -887,13 +873,11 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
   INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
   INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(phase_omega_in);
   INTEGRATOR_STATE_WRITE(state, ray, t) = FLT_MAX;
+
 #  ifdef __RAY_DIFFERENTIALS__
   INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
   INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(phase_domega_in);
 #  endif
-  // Save memory by storing last hit prim and object in isect
-  INTEGRATOR_STATE_WRITE(state, isect, prim) = sd->prim;
-  INTEGRATOR_STATE_WRITE(state, isect, object) = sd->object;
 
   /* Update throughput. */
   const float3 throughput = INTEGRATOR_STATE(state, path, throughput);

intern/cycles/kernel/integrator/shadow_state_template.h

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

intern/cycles/kernel/integrator/subsurface.h

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

intern/cycles/kernel/integrator/subsurface_disk.h

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

intern/cycles/kernel/integrator/subsurface_random_walk.h

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

intern/cycles/kernel/light/light.h

@@ -113,30 +113,22 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
     ls->P = make_float3(klight->co[0], klight->co[1], klight->co[2]);
 
     if (type == LIGHT_SPOT) {
-      const float3 center = make_float3(klight->co[0], klight->co[1], klight->co[2]);
-      const float radius = klight->spot.radius;
-      const float3 dir = make_float3(
-          klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
-      /* disk oriented normal */
-      const float3 lightN = normalize(P - center);
-      ls->P = center;
+      ls->Ng = make_float3(klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
+      float radius = klight->spot.radius;
 
       if (radius > 0.0f)
-        /* disk light */
-        ls->P += disk_light_sample(lightN, randu, randv) * radius;
-
-      const float invarea = klight->spot.invarea;
-      ls->pdf = invarea;
+        /* sphere light */
+        ls->P += disk_light_sample(ls->Ng, randu, randv) * radius;
 
       ls->D = normalize_len(ls->P - P, &ls->t);
-      /* we set the light normal to the outgoing direction to support texturing */
-      ls->Ng = -ls->D;
 
+      float invarea = klight->spot.invarea;
       ls->eval_fac = (0.25f * M_1_PI_F) * invarea;
+      ls->pdf = invarea;
 
       /* spot light attenuation */
       ls->eval_fac *= spot_light_attenuation(
-          dir, klight->spot.spot_angle, klight->spot.spot_smooth, -ls->D);
+          ls->Ng, klight->spot.spot_angle, klight->spot.spot_smooth, -ls->D);
       if (!in_volume_segment && ls->eval_fac == 0.0f) {
         return false;
       }
@@ -145,33 +137,32 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
       ls->u = uv.x;
       ls->v = uv.y;
 
-      ls->pdf *= lamp_light_pdf(kg, lightN, -ls->D, ls->t);
+      ls->pdf *= lamp_light_pdf(kg, ls->Ng, -ls->D, ls->t);
     }
     else if (type == LIGHT_POINT) {
       float3 center = make_float3(klight->co[0], klight->co[1], klight->co[2]);
       float radius = klight->spot.radius;
-      /* disk oriented normal */
-      const float3 lightN = normalize(P - center);
       ls->P = center;
+      float pdf = 1.0;
 
       if (radius > 0.0f) {
-        ls->P += disk_light_sample(lightN, randu, randv) * radius;
+        ls->Ng = normalize(P - center);
+        ls->P += disk_light_sample(ls->Ng, randu, randv) * radius;
+        pdf = klight->spot.invarea;
+        ls->D = normalize_len(ls->P - P, &ls->t);
+      }
+      else {
+        ls->Ng = normalize(P - center);
       }
-      ls->pdf = klight->spot.invarea;
 
       ls->D = normalize_len(ls->P - P, &ls->t);
-      /* we set the light normal to the outgoing direction to support texturing */
-      ls->Ng = -ls->D;
-
+      ls->pdf = pdf;
       ls->eval_fac = M_1_PI_F * 0.25f * klight->spot.invarea;
-      if (!in_volume_segment && ls->eval_fac == 0.0f) {
-        return false;
-      }
 
       float2 uv = map_to_sphere(ls->Ng);
       ls->u = uv.x;
       ls->v = uv.y;
-      ls->pdf *= lamp_light_pdf(kg, lightN, -ls->D, ls->t);
+      ls->pdf *= lamp_light_pdf(kg, ls->Ng, -ls->D, ls->t);
     }
     else {
       /* area light */
@@ -272,16 +263,14 @@ ccl_device bool lights_intersect(KernelGlobals kg,
 
     if (type == LIGHT_SPOT) {
       /* Spot/Disk light. */
-      const float mis_ray_t = INTEGRATOR_STATE(state, path, mis_ray_t);
-      const float3 ray_P = ray->P - ray->D * mis_ray_t;
-
       const float3 lightP = make_float3(klight->co[0], klight->co[1], klight->co[2]);
+      const float3 lightN = make_float3(
+          klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
       const float radius = klight->spot.radius;
       if (radius == 0.0f) {
         continue;
       }
-      /* disk oriented normal */
-      const float3 lightN = normalize(ray_P - lightP);
+
       /* One sided. */
       if (dot(ray->D, lightN) >= 0.0f) {
         continue;
@@ -303,10 +292,9 @@ ccl_device bool lights_intersect(KernelGlobals kg,
         continue;
       }
 
-      /* disk oriented normal */
-      const float3 lightN = normalize(ray_P - lightP);
       float3 P;
-      if (!ray_disk_intersect(ray->P, ray->D, ray->t, lightP, lightN, radius, &P, &t)) {
+      const float3 lsN = normalize(ray_P - lightP);
+      if (!ray_disk_intersect(ray->P, ray->D, ray->t, lightP, lsN, radius, &P, &t)) {
         continue;
       }
     }
@@ -430,8 +418,8 @@ ccl_device bool light_sample_from_intersection(KernelGlobals kg,
   LightType type = (LightType)klight->type;
   ls->type = type;
   ls->shader = klight->shader_id;
-  ls->object = isect->object;
-  ls->prim = isect->prim;
+  ls->object = PRIM_NONE;
+  ls->prim = PRIM_NONE;
   ls->lamp = lamp;
   /* todo: missing texture coordinates */
   ls->t = isect->t;
@@ -439,12 +427,7 @@ ccl_device bool light_sample_from_intersection(KernelGlobals kg,
   ls->D = ray_D;
 
   if (type == LIGHT_SPOT) {
-    const float3 center = make_float3(klight->co[0], klight->co[1], klight->co[2]);
-    const float3 dir = make_float3(klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
-    /* the normal of the oriented disk */
-    const float3 lightN = normalize(ray_P - center);
-    /* we set the light normal to the outgoing direction to support texturing*/
-    ls->Ng = -ls->D;
+    ls->Ng = make_float3(klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
 
     float invarea = klight->spot.invarea;
     ls->eval_fac = (0.25f * M_1_PI_F) * invarea;
@@ -452,7 +435,7 @@ ccl_device bool light_sample_from_intersection(KernelGlobals kg,
 
     /* spot light attenuation */
     ls->eval_fac *= spot_light_attenuation(
-        dir, klight->spot.spot_angle, klight->spot.spot_smooth, -ls->D);
+        ls->Ng, klight->spot.spot_angle, klight->spot.spot_smooth, -ls->D);
 
     if (ls->eval_fac == 0.0f) {
       return false;
@@ -464,32 +447,23 @@ ccl_device bool light_sample_from_intersection(KernelGlobals kg,
 
     /* compute pdf */
     if (ls->t != FLT_MAX)
-      ls->pdf *= lamp_light_pdf(kg, lightN, -ls->D, ls->t);
-    else
-      ls->pdf = 0.f;
+      ls->pdf *= lamp_light_pdf(kg, ls->Ng, -ls->D, ls->t);
   }
   else if (type == LIGHT_POINT) {
-    const float3 center = make_float3(klight->co[0], klight->co[1], klight->co[2]);
-    const float3 lighN = normalize(ray_P - center);
-
-    /* we set the light normal to the outgoing direction to support texturing*/
-    ls->Ng = -ls->D;
+    float3 center = make_float3(klight->co[0], klight->co[1], klight->co[2]);
 
+    ls->Ng = normalize(ray_P - center);
     float invarea = klight->spot.invarea;
     ls->eval_fac = (0.25f * M_1_PI_F) * invarea;
     ls->pdf = invarea;
 
-    if (ls->eval_fac == 0.0f) {
-      return false;
-    }
-
     float2 uv = map_to_sphere(ls->Ng);
     ls->u = uv.x;
     ls->v = uv.y;
 
     /* compute pdf */
     if (ls->t != FLT_MAX)
-      ls->pdf *= lamp_light_pdf(kg, lighN, -ls->D, ls->t);
+      ls->pdf *= lamp_light_pdf(kg, ls->Ng, -ls->D, ls->t);
     else
       ls->pdf = 0.f;
   }

intern/cycles/kernel/light/sample.h

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

intern/cycles/kernel/svm/ao.h

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

intern/cycles/kernel/svm/bevel.h

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