Fix subpass iteration

The default shader could have been requested as deferred and then as non
defered. This could make the shader compilation assert.

.clang-format

@@ -257,10 +257,6 @@ ForEachMacros:
   - SURFACE_QUAD_ITER_BEGIN
   - foreach
   - ED_screen_areas_iter
-  - SLOT_PROBING_BEGIN
-  - SET_SLOT_PROBING_BEGIN
-  - MAP_SLOT_PROBING_BEGIN
-  - VECTOR_SET_SLOT_PROBING_BEGIN
 
 # Use once we bump the minimum version to version 8.
 # # Without this string literals that in-line 'STRINGIFY' behave strangely (a bug?).

CMakeLists.txt

@@ -281,7 +281,7 @@ option(WITH_ALEMBIC             "Enable Alembic Support" ON)
 option(WITH_ALEMBIC_HDF5        "Enable Legacy Alembic Support (not officially supported)" OFF)
 
 # Universal Scene Description support
-option(WITH_USD                 "Enable Universal Scene Description (USD) Support" ON)
+option(WITH_USD                 "Enable Universal Scene Description (USD) Support" OFF)
 
 # 3D format support
 # Disable opencollada when we don't have precompiled libs
@@ -628,6 +628,12 @@ if(WITH_PYTHON_MODULE AND WITH_PYTHON_INSTALL)
   message(FATAL_ERROR "WITH_PYTHON_MODULE requires WITH_PYTHON_INSTALL to be OFF")
 endif()
 
+
+# may as well build python module without a UI
+if(WITH_PYTHON_MODULE)
+  set(WITH_HEADLESS ON)
+endif()
+
 if(NOT WITH_PYTHON)
   set(WITH_CYCLES OFF)
   set(WITH_DRACO OFF)

GNUmakefile

@@ -32,7 +32,6 @@ Convenience Targets
    * debug:         Build a debug binary.
    * full:          Enable all supported dependencies & options.
    * lite:          Disable non essential features for a smaller binary and faster build.
-   * release        Complete build with all options enabled including CUDA and Optix, matching the releases on blender.org
    * headless:      Build without an interface (renderfarm or server automation).
    * cycles:        Build Cycles standalone only, without Blender.
    * bpy:           Build as a python module which can be loaded from python directly.
@@ -208,10 +207,6 @@ ifneq "$(findstring lite, $(MAKECMDGOALS))" ""
 	BUILD_DIR:=$(BUILD_DIR)_lite
 	CMAKE_CONFIG_ARGS:=-C"$(BLENDER_DIR)/build_files/cmake/config/blender_lite.cmake" $(CMAKE_CONFIG_ARGS)
 endif
-ifneq "$(findstring release, $(MAKECMDGOALS))" ""
-	BUILD_DIR:=$(BUILD_DIR)_release
-	CMAKE_CONFIG_ARGS:=-C"$(BLENDER_DIR)/build_files/cmake/config/blender_release.cmake" $(CMAKE_CONFIG_ARGS)
-endif
 ifneq "$(findstring cycles, $(MAKECMDGOALS))" ""
 	BUILD_DIR:=$(BUILD_DIR)_cycles
 	CMAKE_CONFIG_ARGS:=-C"$(BLENDER_DIR)/build_files/cmake/config/cycles_standalone.cmake" $(CMAKE_CONFIG_ARGS)
@@ -322,7 +317,6 @@ all: .FORCE
 debug: all
 full: all
 lite: all
-release: all
 cycles: all
 headless: all
 bpy: all

build_files/build_environment/CMakeLists.txt

@@ -130,6 +130,7 @@ if(NOT WIN32 OR ENABLE_MINGW64)
     include(cmake/vpx.cmake)
     include(cmake/x264.cmake)
     include(cmake/xvidcore.cmake)
+    include(cmake/faad.cmake)
     include(cmake/ffmpeg.cmake)
     include(cmake/fftw.cmake)
     include(cmake/sndfile.cmake)

build_files/build_environment/cmake/alembic.cmake

@@ -38,7 +38,6 @@ set(ALEMBIC_EXTRA_ARGS
   -DBoost_DEBUG=ON
   -DBOOST_ROOT=${LIBDIR}/boost
   -DBoost_NO_SYSTEM_PATHS=ON
-  -DBoost_NO_BOOST_CMAKE=ON
   -DILMBASE_ROOT=${LIBDIR}/openexr
   -DALEMBIC_ILMBASE_INCLUDE_DIRECTORY=${LIBDIR}/openexr/include/OpenEXR
   -DALEMBIC_ILMBASE_HALF_LIB=${LIBDIR}/openexr/lib/${LIBPREFIX}Half${OPENEXR_VERSION_POSTFIX}${LIBEXT}

build_files/build_environment/cmake/embree.cmake

@@ -26,20 +26,10 @@ set(EMBREE_EXTRA_ARGS
   -DEMBREE_RAY_MASK=ON
   -DEMBREE_FILTER_FUNCTION=ON
   -DEMBREE_BACKFACE_CULLING=OFF
+  -DEMBREE_TASKING_SYSTEM=INTERNAL
   -DEMBREE_MAX_ISA=AVX2
-  -DEMBREE_TASKING_SYSTEM=TBB
-  -DEMBREE_TBB_ROOT=${LIBDIR}/tbb
-  -DTBB_STATIC_LIB=${TBB_STATIC_LIBRARY}
 )
 
-if(TBB_STATIC_LIBRARY)
-  set(EMBREE_EXTRA_ARGS
-    ${EMBREE_EXTRA_ARGS}
-    -DEMBREE_TBB_LIBRARY_NAME=tbb_static
-    -DEMBREE_TBBMALLOC_LIBRARY_NAME=tbbmalloc_static
-  )
-endif()
-
 if(WIN32)
   set(EMBREE_BUILD_DIR ${BUILD_MODE}/)
 else()
@@ -54,12 +44,6 @@ ExternalProject_Add(external_embree
   CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/embree ${DEFAULT_CMAKE_FLAGS} ${EMBREE_EXTRA_ARGS}
   INSTALL_DIR ${LIBDIR}/embree
 )
-
-add_dependencies(
-  external_embree
-  external_tbb
-)
-
 if(WIN32)
 
   if(BUILD_MODE STREQUAL Release)

build_files/build_environment/cmake/faad.cmake

@@ -0,0 +1,40 @@
+# ***** BEGIN GPL LICENSE BLOCK *****
+#
+# This program is free software; you can redistribute it and/or
+# modify it under the terms of the GNU General Public License
+# as published by the Free Software Foundation; either version 2
+# of the License, or (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program; if not, write to the Free Software Foundation,
+# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+#
+# ***** END GPL LICENSE BLOCK *****
+
+set(FAAD_EXTRA_ARGS)
+
+if(WIN32)
+  set(FAAD_EXTRA_CONFIGURE "utils\\win32\\ac2ver.exe" "faad2" "configure.ac" > libfaad\\win32_ver.h)
+else()
+  set(FAAD_EXTRA_CONFIGURE echo .)
+endif()
+
+ExternalProject_Add(external_faad
+  URL ${FAAD_URI}
+  DOWNLOAD_DIR ${DOWNLOAD_DIR}
+  URL_HASH MD5=${FAAD_HASH}
+  PREFIX ${BUILD_DIR}/faad
+  CONFIGURE_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/faad/src/external_faad/ && ${FAAD_EXTRA_CONFIGURE} && ${CONFIGURE_COMMAND} --disable-shared --enable-static --prefix=${LIBDIR}/faad
+  BUILD_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/faad/src/external_faad/ && make -j${MAKE_THREADS}
+  INSTALL_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/faad/src/external_faad/ && make install
+  INSTALL_DIR ${LIBDIR}/faad
+)
+
+if(MSVC)
+  set_target_properties(external_faad PROPERTIES FOLDER Mingw)
+endif()

build_files/build_environment/cmake/ffmpeg.cmake

@@ -127,6 +127,7 @@ endif()
 add_dependencies(
   external_ffmpeg
   external_zlib
+  external_faad
   external_openjpeg
   external_xvidcore
   external_x264

build_files/build_environment/cmake/fftw.cmake

@@ -19,12 +19,8 @@
 set(FFTW_EXTRA_ARGS)
 
 if(WIN32)
+  set(FFTW3_ENV set CFLAGS=-fno-stack-check -fno-stack-protector -mno-stack-arg-probe -fno-lto &&)
   set(FFTW3_PATCH_COMMAND ${PATCH_CMD} --verbose -p 0 -N -d ${BUILD_DIR}/fftw3/src/external_fftw3 < ${PATCH_DIR}/fftw3.diff)
-  set(FFTW_EXTRA_ARGS --disable-static --enable-shared)
-  set(FFTW_INSTALL install-strip)
-else()
-  set(FFTW_EXTRA_ARGS --enable-static)
-  set(FFTW_INSTALL install)
 endif()
 
 ExternalProject_Add(external_fftw3
@@ -32,10 +28,10 @@ ExternalProject_Add(external_fftw3
   DOWNLOAD_DIR ${DOWNLOAD_DIR}
   URL_HASH MD5=${FFTW_HASH}
   PREFIX ${BUILD_DIR}/fftw3
-  CONFIGURE_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/fftw3/src/external_fftw3/ && ${CONFIGURE_COMMAND} ${FFTW_EXTRA_ARGS} --prefix=${mingw_LIBDIR}/fftw3
+  CONFIGURE_COMMAND ${CONFIGURE_ENV} && ${FFTW3_ENV} cd ${BUILD_DIR}/fftw3/src/external_fftw3/ && ${CONFIGURE_COMMAND} --enable-static --prefix=${mingw_LIBDIR}/fftw3
   PATCH_COMMAND ${FFTW3_PATCH_COMMAND}
   BUILD_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/fftw3/src/external_fftw3/ && make -j${MAKE_THREADS}
-  INSTALL_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/fftw3/src/external_fftw3/ && make ${FFTW_INSTALL}
+  INSTALL_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/fftw3/src/external_fftw3/ && make install
   INSTALL_DIR ${LIBDIR}/fftw3
 )
 
@@ -43,8 +39,7 @@ if(MSVC)
   set_target_properties(external_fftw3 PROPERTIES FOLDER Mingw)
   if(BUILD_MODE STREQUAL Release)
     ExternalProject_Add_Step(external_fftw3 after_install
-      COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/fftw3/lib/libfftw3.dll.a ${HARVEST_TARGET}/fftw3/lib/libfftw.lib
+      COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/fftw3/lib/libfftw3.a ${HARVEST_TARGET}/fftw3/lib/libfftw.lib
-      COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/fftw3/bin/libfftw3-3.dll ${HARVEST_TARGET}/fftw3/lib/libfftw3-3.dll
       COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/fftw3/include/fftw3.h ${HARVEST_TARGET}/fftw3/include/fftw3.h
       DEPENDEES install
     )

build_files/build_environment/cmake/harvest.cmake

@@ -31,6 +31,10 @@ if(BUILD_MODE STREQUAL Release)
     COMMAND # jpeg rename libfile + copy include
         ${CMAKE_COMMAND} -E copy ${LIBDIR}/jpg/lib/jpeg-static.lib ${HARVEST_TARGET}/jpeg/lib/libjpeg.lib &&
         ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/jpg/include/ ${HARVEST_TARGET}/jpeg/include/ &&
+        # OpenImageIO
+        ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/OpenImageIO/include ${HARVEST_TARGET}/OpenImageIO/include &&
+        ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/OpenImageIO/lib ${HARVEST_TARGET}/OpenImageIO/lib &&
+        ${CMAKE_COMMAND} -E copy ${LIBDIR}/OpenImageIO/bin/idiff.exe ${HARVEST_TARGET}/OpenImageIO/bin/idiff.exe &&
         # png
         ${CMAKE_COMMAND} -E copy ${LIBDIR}/png/lib/libpng16_static.lib ${HARVEST_TARGET}/png/lib/libpng.lib &&
         ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/png/include/ ${HARVEST_TARGET}/png/include/ &&
@@ -40,6 +44,10 @@ if(BUILD_MODE STREQUAL Release)
         # glew-> opengl
         ${CMAKE_COMMAND} -E copy ${LIBDIR}/glew/lib/libglew32.lib ${HARVEST_TARGET}/opengl/lib/glew.lib &&
         ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/glew/include/ ${HARVEST_TARGET}/opengl/include/ &&
+        # sndfile
+        ${CMAKE_COMMAND} -E copy ${LIBDIR}/sndfile/lib/libsndfile.dll.a ${HARVEST_TARGET}/sndfile/lib/libsndfile-1.lib &&
+        ${CMAKE_COMMAND} -E copy ${LIBDIR}/sndfile/bin/libsndfile-1.dll ${HARVEST_TARGET}/sndfile/lib/libsndfile-1.dll &&
+        ${CMAKE_COMMAND} -E copy ${LIBDIR}/sndfile/include/sndfile.h ${HARVEST_TARGET}/sndfile/include/sndfile.h &&
         # tiff
         ${CMAKE_COMMAND} -E copy ${LIBDIR}/tiff/lib/tiff.lib ${HARVEST_TARGET}/tiff/lib/libtiff.lib &&
         ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/tiff/include/ ${HARVEST_TARGET}/tiff/include/ &&
@@ -51,7 +59,9 @@ endif()
 
 if(BUILD_MODE STREQUAL Debug)
   add_custom_target(Harvest_Debug_Results
-    COMMAND
+        # OpenImageIO
+    COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/openimageio/lib/OpenImageIO.lib ${HARVEST_TARGET}/openimageio/lib/OpenImageIO_d.lib &&
+        ${CMAKE_COMMAND} -E copy ${LIBDIR}/openimageio/lib/OpenImageIO_Util.lib ${HARVEST_TARGET}/openimageio/lib/OpenImageIO_Util_d.lib &&
         # hdf5
         ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/hdf5/lib ${HARVEST_TARGET}/hdf5/lib &&
     DEPENDS Package_Python

build_files/build_environment/cmake/jpeg.cmake

@@ -42,17 +42,20 @@ if(WIN32)
     set(JPEG_LIBRARY jpeg-staticd${LIBEXT})
   endif()
 else(WIN32)
-  # cmake for unix
+  # autoconf for unix
-  set(JPEG_EXTRA_ARGS
+  if(APPLE)
-    -DWITH_JPEG8=ON
+    set(JPEG_EXTRA_ARGS --host x86_64-apple-darwin --with-jpeg8)
-    -DENABLE_STATIC=ON
+  else()
-    -DENABLE_SHARED=OFF
+    set(JPEG_EXTRA_ARGS --with-jpeg8)
-    -DCMAKE_INSTALL_LIBDIR=${LIBDIR}/jpg/lib)
+  endif()
 
   ExternalProject_Add(external_jpeg
     URL ${JPEG_URI}
     DOWNLOAD_DIR ${DOWNLOAD_DIR}
     URL_HASH MD5=${JPEG_HASH}
+    CONFIGURE_COMMAND ${CONFIGURE_ENV} && autoreconf -fiv && ${CONFIGURE_COMMAND} --prefix=${LIBDIR}/jpg NASM=yasm ${JPEG_EXTRA_ARGS}
+    BUILD_IN_SOURCE 1
+    BUILD_COMMAND ${CONFIGURE_ENV} && make install
     PREFIX ${BUILD_DIR}/jpg
     CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/jpg ${DEFAULT_CMAKE_FLAGS} ${JPEG_EXTRA_ARGS}
     INSTALL_DIR ${LIBDIR}/jpg

build_files/build_environment/cmake/openal.cmake

@@ -52,6 +52,7 @@ if(BUILD_MODE STREQUAL Release)
     PREFIX ${BUILD_DIR}/openal
     CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/openal ${DEFAULT_CMAKE_FLAGS} ${OPENAL_EXTRA_ARGS}
     INSTALL_DIR ${LIBDIR}/openal
+    PATCH_COMMAND ${PATCH_CMD} -p 1 -d ${BUILD_DIR}/openal/src/external_openal < ${PATCH_DIR}/openal.diff
   )
 
   if(WIN32)

build_files/build_environment/cmake/openimageio.cmake

@@ -32,7 +32,7 @@ endif()
 
 if(WIN32)
   set(PNG_LIBNAME libpng16_static${LIBEXT})
-  set(OIIO_SIMD_FLAGS -DUSE_SIMD=sse2)
+  set(OIIO_SIMD_FLAGS -DUSE_SIMD=sse2 -DOPJ_STATIC=1)
   set(OPENJPEG_POSTFIX _msvc)
 else()
   set(PNG_LIBNAME libpng${LIBEXT})
@@ -49,17 +49,23 @@ endif()
 
 if(MSVC)
   set(OPENJPEG_FLAGS
-    -DOpenJpeg_ROOT=${LIBDIR}/openjpeg_msvc
+    -DOPENJPEG_HOME=${LIBDIR}/openjpeg_msvc
+    -DOPENJPEG_INCLUDE_DIR=${LIBDIR}/openjpeg_msvc/include/openjpeg-${OPENJPEG_SHORT_VERSION}
+    -DOPENJPEG_LIBRARY=${LIBDIR}/openjpeg_msvc/lib/openjp2${LIBEXT}
+    -DOPENJPEG_LIBRARY_DEBUG=${LIBDIR}/openjpeg_msvc/lib/openjp2${LIBEXT}
   )
 else()
   set(OPENJPEG_FLAGS
-    -DOpenJpeg_ROOT=${LIBDIR}/openjpeg
+    -DOPENJPEG_INCLUDE_DIR=${LIBDIR}/openjpeg/include/openjpeg-${OPENJPEG_SHORT_VERSION}
+    -DOPENJPEG_LIBRARY=${LIBDIR}/openjpeg/lib/${OPENJPEG_LIBRARY}
   )
 endif()
 
 set(OPENIMAGEIO_EXTRA_ARGS
-  -DBUILD_SHARED_LIBS=OFF
+  -DBUILDSTATIC=ON
   ${OPENIMAGEIO_LINKSTATIC}
+  -DOPENEXR_INCLUDE_DIR=${LIBDIR}/openexr/include/openexr/
+  -DOPENEXR_ILMIMF_LIBRARIES=${LIBDIR}/openexr/lib/IlmImf${OPENEXR_VERSION_POSTFIX}${LIBEXT}
   -DBoost_COMPILER:STRING=${BOOST_COMPILER_STRING}
   -DBoost_USE_MULTITHREADED=ON
   -DBoost_USE_STATIC_LIBS=ON
@@ -67,16 +73,7 @@ set(OPENIMAGEIO_EXTRA_ARGS
   -DBOOST_ROOT=${LIBDIR}/boost
   -DBOOST_LIBRARYDIR=${LIBDIR}/boost/lib/
   -DBoost_NO_SYSTEM_PATHS=ON
-  -DBoost_NO_BOOST_CMAKE=ON
   -OIIO_BUILD_CPP11=ON
-  -DUSE_LIBSQUISH=OFF
-  -DUSE_QT5=OFF
-  -DUSE_NUKE=OFF
-  -DUSE_OPENVDB=OFF
-  -DUSE_BZIP2=OFF
-  -DUSE_FREETYPE=OFF
-  -DUSE_DCMTK=OFF
-  -DUSE_LIBHEIF=OFF
   -DUSE_OPENGL=OFF
   -DUSE_TBB=OFF
   -DUSE_FIELD3D=OFF
@@ -84,12 +81,15 @@ set(OPENIMAGEIO_EXTRA_ARGS
   -DUSE_PYTHON=OFF
   -DUSE_GIF=OFF
   -DUSE_OPENCV=OFF
+  -DUSE_OPENSSL=OFF
   -DUSE_OPENJPEG=ON
   -DUSE_FFMPEG=OFF
   -DUSE_PTEX=OFF
   -DUSE_FREETYPE=OFF
   -DUSE_LIBRAW=OFF
-  -DUSE_OPENCOLORIO=OFF
+  -DUSE_PYTHON=OFF
+  -DUSE_PYTHON3=OFF
+  -DUSE_OCIO=OFF
   -DUSE_WEBP=${WITH_WEBP}
   -DOIIO_BUILD_TOOLS=${OIIO_TOOLS}
   -DOIIO_BUILD_TESTS=OFF
@@ -103,13 +103,17 @@ set(OPENIMAGEIO_EXTRA_ARGS
   -DJPEG_LIBRARY=${LIBDIR}/jpg/lib/${JPEG_LIBRARY}
   -DJPEG_INCLUDE_DIR=${LIBDIR}/jpg/include
   ${OPENJPEG_FLAGS}
+  -DOCIO_PATH=${LIBDIR}/opencolorio/
   -DOpenEXR_USE_STATIC_LIBS=On
+  -DOPENEXR_HOME=${LIBDIR}/openexr/
+  -DILMBASE_INCLUDE_PATH=${LIBDIR}/openexr/
+  -DILMBASE_PACKAGE_PREFIX=${LIBDIR}/openexr/
   -DILMBASE_INCLUDE_DIR=${LIBDIR}/openexr/include/
-  -DOPENEXR_INCLUDE_DIR=${LIBDIR}/openexr/include/
   -DOPENEXR_HALF_LIBRARY=${LIBDIR}/openexr/lib/${LIBPREFIX}Half${OPENEXR_VERSION_POSTFIX}${LIBEXT}
   -DOPENEXR_IMATH_LIBRARY=${LIBDIR}/openexr/lib/${LIBPREFIX}Imath${OPENEXR_VERSION_POSTFIX}${LIBEXT}
   -DOPENEXR_ILMTHREAD_LIBRARY=${LIBDIR}/openexr/lib/${LIBPREFIX}IlmThread${OPENEXR_VERSION_POSTFIX}${LIBEXT}
   -DOPENEXR_IEX_LIBRARY=${LIBDIR}/openexr/lib/${LIBPREFIX}Iex${OPENEXR_VERSION_POSTFIX}${LIBEXT}
+  -DOPENEXR_INCLUDE_DIR=${LIBDIR}/openexr/include/
   -DOPENEXR_ILMIMF_LIBRARY=${LIBDIR}/openexr/lib/${LIBPREFIX}IlmImf${OPENEXR_VERSION_POSTFIX}${LIBEXT}
   -DSTOP_ON_WARNING=OFF
   ${WEBP_FLAGS}
@@ -121,38 +125,27 @@ ExternalProject_Add(external_openimageio
   DOWNLOAD_DIR ${DOWNLOAD_DIR}
   URL_HASH MD5=${OPENIMAGEIO_HASH}
   PREFIX ${BUILD_DIR}/openimageio
-  PATCH_COMMAND ${PATCH_CMD} -p 1 -N -d ${BUILD_DIR}/openimageio/src/external_openimageio/ < ${PATCH_DIR}/openimageio.diff
+  PATCH_COMMAND
+    ${PATCH_CMD} -p 0 -N -d ${BUILD_DIR}/openimageio/src/external_openimageio/src/include < ${PATCH_DIR}/openimageio_gdi.diff &&
+    ${PATCH_CMD} -p 1 -N -d ${BUILD_DIR}/openimageio/src/external_openimageio/ < ${PATCH_DIR}/openimageio_static_libs.diff
   CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/openimageio ${DEFAULT_CMAKE_FLAGS} ${OPENIMAGEIO_EXTRA_ARGS}
   INSTALL_DIR ${LIBDIR}/openimageio
 )
 
 add_dependencies(
   external_openimageio
-  external_png
+  external_png external_zlib
-  external_zlib
   external_openexr
   external_jpeg
   external_boost
   external_tiff
+  external_opencolorio
   external_openjpeg${OPENJPEG_POSTFIX}
   ${WEBP_DEP}
 )
-
+if(NOT WIN32)
-if(WIN32)
+  add_dependencies(
-  if(BUILD_MODE STREQUAL Release)
+    external_openimageio
-    ExternalProject_Add_Step(external_openimageio after_install
+    external_opencolorio_extra
-      COMMAND ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/OpenImageIO/include ${HARVEST_TARGET}/OpenImageIO/include
+  )
-      COMMAND ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/OpenImageIO/lib ${HARVEST_TARGET}/OpenImageIO/lib
-      COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/OpenImageIO/bin/idiff.exe ${HARVEST_TARGET}/OpenImageIO/bin/idiff.exe
-
-      DEPENDEES install
-    )
-  endif()
-  if(BUILD_MODE STREQUAL Debug)
-    ExternalProject_Add_Step(external_openimageio after_install
-      COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/openimageio/lib/OpenImageIO.lib ${HARVEST_TARGET}/openimageio/lib/OpenImageIO_d.lib
-      COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/openimageio/lib/OpenImageIO_Util.lib ${HARVEST_TARGET}/openimageio/lib/OpenImageIO_Util_d.lib
-      DEPENDEES install
-    )
-  endif()
 endif()

build_files/build_environment/cmake/opensubdiv.cmake

@@ -67,7 +67,7 @@ endif()
 ExternalProject_Add(external_opensubdiv
   URL ${OPENSUBDIV_URI}
   DOWNLOAD_DIR ${DOWNLOAD_DIR}
-  URL_HASH MD5=${OPENSUBDIV_HASH}
+  URL_HASH MD5=${OPENSUBDIV_Hash}
   PREFIX ${BUILD_DIR}/opensubdiv
   CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/opensubdiv -Wno-dev ${DEFAULT_CMAKE_FLAGS} ${OPENSUBDIV_EXTRA_ARGS}
   INSTALL_DIR ${LIBDIR}/opensubdiv

build_files/build_environment/cmake/openvdb.cmake

@@ -27,7 +27,6 @@ set(OPENVDB_EXTRA_ARGS
   -DBoost_USE_STATIC_RUNTIME=OFF
   -DBOOST_ROOT=${LIBDIR}/boost
   -DBoost_NO_SYSTEM_PATHS=ON
-  -DBoost_NO_BOOST_CMAKE=ON
   -DZLIB_LIBRARY=${LIBDIR}/zlib/lib/${ZLIB_LIBRARY}
   -DZLIB_INCLUDE_DIR=${LIBDIR}/zlib/include/
   -DBlosc_INCLUDE_DIR=${LIBDIR}/blosc/include/

build_files/build_environment/cmake/osl.cmake

@@ -44,7 +44,6 @@ set(OSL_EXTRA_ARGS
   -DBOOST_ROOT=${LIBDIR}/boost
   -DBOOST_LIBRARYDIR=${LIBDIR}/boost/lib/
   -DBoost_NO_SYSTEM_PATHS=ON
-  -DBoost_NO_BOOST_CMAKE=ON
   -DLLVM_DIRECTORY=${LIBDIR}/llvm
   -DLLVM_INCLUDES=${LIBDIR}/llvm/include
   -DLLVM_LIB_DIR=${LIBDIR}/llvm/lib

build_files/build_environment/cmake/png.cmake

@@ -25,7 +25,7 @@ set(PNG_EXTRA_ARGS
 ExternalProject_Add(external_png
   URL ${PNG_URI}
   DOWNLOAD_DIR ${DOWNLOAD_DIR}
-  URL_HASH SHA256=${PNG_HASH}
+  URL_HASH MD5=${PNG_HASH}
   PREFIX ${BUILD_DIR}/png
   CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/png ${DEFAULT_CMAKE_FLAGS} ${PNG_EXTRA_ARGS}
   INSTALL_DIR ${LIBDIR}/png

build_files/build_environment/cmake/sndfile.cmake

@@ -60,14 +60,3 @@ if(UNIX)
     external_flac
   )
 endif()
-
-if(BUILD_MODE STREQUAL Release AND WIN32)
-  ExternalProject_Add_Step(external_sndfile after_install
-      COMMAND  lib /def:${BUILD_DIR}/sndfile/src/external_sndfile/src/libsndfile-1.def /machine:x64 /out:${BUILD_DIR}/sndfile/src/external_sndfile/src/libsndfile-1.lib
-      COMMAND  ${CMAKE_COMMAND} -E copy ${LIBDIR}/sndfile/bin/libsndfile-1.dll ${HARVEST_TARGET}/sndfile/lib/libsndfile-1.dll
-      COMMAND  ${CMAKE_COMMAND} -E copy ${BUILD_DIR}/sndfile/src/external_sndfile/src/libsndfile-1.lib ${HARVEST_TARGET}/sndfile/lib/libsndfile-1.lib
-      COMMAND  ${CMAKE_COMMAND} -E copy ${LIBDIR}/sndfile/include/sndfile.h ${HARVEST_TARGET}/sndfile/include/sndfile.h
-
-    DEPENDEES install
-  )
-endif()

build_files/build_environment/cmake/usd.cmake

@@ -22,8 +22,6 @@ set(USD_EXTRA_ARGS
   -DBoost_USE_STATIC_LIBS=ON
   -DBoost_USE_STATIC_RUNTIME=OFF
   -DBOOST_ROOT=${LIBDIR}/boost
-  -DBoost_NO_SYSTEM_PATHS=ON
-  -DBoost_NO_BOOST_CMAKE=ON
   -DTBB_INCLUDE_DIRS=${LIBDIR}/tbb/include
   -DTBB_LIBRARIES=${LIBDIR}/tbb/lib/${LIBPREFIX}${TBB_LIBRARY}${LIBEXT}
   -DTbb_TBB_LIBRARY=${LIBDIR}/tbb/lib/${LIBPREFIX}${TBB_LIBRARY}${LIBEXT}

build_files/build_environment/cmake/versions.cmake

@@ -20,17 +20,17 @@ set(ZLIB_VERSION 1.2.11)
 set(ZLIB_URI https://zlib.net/zlib-${ZLIB_VERSION}.tar.gz)
 set(ZLIB_HASH 1c9f62f0778697a09d36121ead88e08e)
 
-set(OPENAL_VERSION 1.20.1)
+set(OPENAL_VERSION 1.18.2)
 set(OPENAL_URI http://openal-soft.org/openal-releases/openal-soft-${OPENAL_VERSION}.tar.bz2)
-set(OPENAL_HASH 556695068ce8375b89986083d810fd35)
+set(OPENAL_HASH d4eeb0889812e2fdeaa1843523d76190)
 
-set(PNG_VERSION 1.6.37)
+set(PNG_VERSION 1.6.35)
 set(PNG_URI http://prdownloads.sourceforge.net/libpng/libpng-${PNG_VERSION}.tar.xz)
-set(PNG_HASH 505e70834d35383537b6491e7ae8641f1a4bed1876dbfe361201fc80868d88ca)
+set(PNG_HASH 678b7e696a62a193ed3503b04bf449d6)
 
-set(JPEG_VERSION 2.0.4)
+set(JPEG_VERSION 1.5.3)
 set(JPEG_URI https://github.com/libjpeg-turbo/libjpeg-turbo/archive/${JPEG_VERSION}.tar.gz)
-set(JPEG_HASH 44c43e4a9fb352f47090804529317c88)
+set(JPEG_HASH 5b7549d440b86c98a517355c102d155e)
 
 set(BOOST_VERSION 1.70.0)
 set(BOOST_VERSION_NODOTS 1_70_0)
@@ -66,9 +66,9 @@ else()
   set(OPENEXR_VERSION_POSTFIX)
 endif()
 
-set(FREETYPE_VERSION 2.10.2)
+set(FREETYPE_VERSION 2.10.1)
 set(FREETYPE_URI http://prdownloads.sourceforge.net/freetype/freetype-${FREETYPE_VERSION}.tar.gz)
-set(FREETYPE_HASH b1cb620e4c875cd4d1bfa04945400945)
+set(FREETYPE_HASH c50a3c9e5e62bdc938a6e1598a782947)
 
 set(GLEW_VERSION 1.13.0)
 set(GLEW_URI http://prdownloads.sourceforge.net/glew/glew/${GLEW_VERSION}/glew-${GLEW_VERSION}.tgz)
@@ -101,21 +101,21 @@ set(CUEW_GIT_UID 1744972026de9cf27c8a7dc39cf39cd83d5f922f)
 set(CUEW_URI https://github.com/CudaWrangler/cuew/archive/${CUEW_GIT_UID}.zip)
 set(CUEW_HASH 86760d62978ebfd96cd93f5aa1abaf4a)
 
-set(OPENSUBDIV_VERSION v3_4_3)
+set(OPENSUBDIV_VERSION v3_4_0_RC2)
+set(OPENSUBDIV_Hash f6a10ba9efaa82fde86fe65aad346319)
 set(OPENSUBDIV_URI https://github.com/PixarAnimationStudios/OpenSubdiv/archive/${OPENSUBDIV_VERSION}.tar.gz)
-set(OPENSUBDIV_HASH 7bbfa275d021fb829e521df749160edb)
 
-set(SDL_VERSION 2.0.12)
+set(SDL_VERSION 2.0.8)
 set(SDL_URI https://www.libsdl.org/release/SDL2-${SDL_VERSION}.tar.gz)
-set(SDL_HASH 783b6f2df8ff02b19bb5ce492b99c8ff)
+set(SDL_HASH 3800d705cef742c6a634f202c37f263f)
 
 set(OPENCOLLADA_VERSION v1.6.68)
 set(OPENCOLLADA_URI https://github.com/KhronosGroup/OpenCOLLADA/archive/${OPENCOLLADA_VERSION}.tar.gz)
 set(OPENCOLLADA_HASH ee7dae874019fea7be11613d07567493)
 
-set(OPENCOLORIO_VERSION 1.1.1)
+set(OPENCOLORIO_VERSION 1.1.0)
-set(OPENCOLORIO_URI https://github.com/AcademySoftwareFoundation/OpenColorIO/archive/v${OPENCOLORIO_VERSION}.tar.gz)
+set(OPENCOLORIO_URI https://github.com/imageworks/OpenColorIO/archive/v${OPENCOLORIO_VERSION}.tar.gz)
-set(OPENCOLORIO_HASH 23d8b9ac81599305539a5a8674b94a3d)
+set(OPENCOLORIO_HASH 802d8f5b1d1fe316ec5f76511aa611b8)
 
 set(LLVM_VERSION 9.0.1)
 set(LLVM_URI https://github.com/llvm/llvm-project/releases/download/llvmorg-${LLVM_VERSION}/llvm-${LLVM_VERSION}.src.tar.xz)
@@ -127,17 +127,17 @@ set(CLANG_HASH 13468e4a44940efef1b75e8641752f90)
 set(OPENMP_URI https://github.com/llvm/llvm-project/releases/download/llvmorg-${LLVM_VERSION}/openmp-${LLVM_VERSION}.src.tar.xz)
 set(OPENMP_HASH 6eade16057edbdecb3c4eef9daa2bfcf)
 
-set(OPENIMAGEIO_VERSION 2.1.15.0)
+set(OPENIMAGEIO_VERSION 1.8.13)
 set(OPENIMAGEIO_URI https://github.com/OpenImageIO/oiio/archive/Release-${OPENIMAGEIO_VERSION}.tar.gz)
-set(OPENIMAGEIO_HASH f03aa5e3ac4795af04771ee4146e9832)
+set(OPENIMAGEIO_HASH f5526c3c9878029ee900d84856683f93)
 
-set(TIFF_VERSION 4.1.0)
+set(TIFF_VERSION 4.0.9)
 set(TIFF_URI http://download.osgeo.org/libtiff/tiff-${TIFF_VERSION}.tar.gz)
-set(TIFF_HASH 2165e7aba557463acc0664e71a3ed424)
+set(TIFF_HASH 54bad211279cc93eb4fca31ba9bfdc79)
 
-set(OSL_VERSION 1.10.10)
+set(OSL_VERSION 1.10.9)
 set(OSL_URI https://github.com/imageworks/OpenShadingLanguage/archive/Release-${OSL_VERSION}.tar.gz)
-set(OSL_HASH 00dec08a93c8084e53848b9ad047889f)
+set(OSL_HASH a94f1e8506f7e8f5e993653de5c5fa00)
 
 set(PYTHON_VERSION 3.7.4)
 set(PYTHON_SHORT_VERSION 3.7)
@@ -168,9 +168,9 @@ set(LAME_VERSION 3.100)
 set(LAME_URI http://downloads.sourceforge.net/project/lame/lame/3.100/lame-${LAME_VERSION}.tar.gz)
 set(LAME_HASH 83e260acbe4389b54fe08e0bdbf7cddb)
 
-set(OGG_VERSION 1.3.4)
+set(OGG_VERSION 1.3.3)
 set(OGG_URI http://downloads.xiph.org/releases/ogg/libogg-${OGG_VERSION}.tar.gz)
-set(OGG_HASH fe5670640bd49e828d64d2879c31cb4dde9758681bb664f9bdbf159a01b0c76e)
+set(OGG_HASH c2e8a485110b97550f453226ec644ebac6cb29d1caef2902c007edab4308d985)
 
 set(VORBIS_VERSION 1.3.6)
 set(VORBIS_URI http://downloads.xiph.org/releases/vorbis/libvorbis-${VORBIS_VERSION}.tar.gz)
@@ -180,41 +180,47 @@ set(THEORA_VERSION 1.1.1)
 set(THEORA_URI http://downloads.xiph.org/releases/theora/libtheora-${THEORA_VERSION}.tar.bz2)
 set(THEORA_HASH b6ae1ee2fa3d42ac489287d3ec34c5885730b1296f0801ae577a35193d3affbc)
 
-set(FLAC_VERSION 1.3.3)
+set(FLAC_VERSION 1.3.2)
 set(FLAC_URI http://downloads.xiph.org/releases/flac/flac-${FLAC_VERSION}.tar.xz)
-set(FLAC_HASH 213e82bd716c9de6db2f98bcadbc4c24c7e2efe8c75939a1a84e28539c4e1748)
+set(FLAC_HASH 91cfc3ed61dc40f47f050a109b08610667d73477af6ef36dcad31c31a4a8d53f)
 
-set(VPX_VERSION 1.8.2)
+set(VPX_VERSION 1.7.0)
 set(VPX_URI https://github.com/webmproject/libvpx/archive/v${VPX_VERSION}/libvpx-v${VPX_VERSION}.tar.gz)
-set(VPX_HASH 8735d9fcd1a781ae6917f28f239a8aa358ce4864ba113ea18af4bb2dc8b474ac)
+set(VPX_HASH 1fec931eb5c94279ad219a5b6e0202358e94a93a90cfb1603578c326abfc1238)
 
 set(OPUS_VERSION 1.3.1)
 set(OPUS_URI https://archive.mozilla.org/pub/opus/opus-${OPUS_VERSION}.tar.gz)
 set(OPUS_HASH 65b58e1e25b2a114157014736a3d9dfeaad8d41be1c8179866f144a2fb44ff9d)
 
-set(X264_URI https://code.videolan.org/videolan/x264/-/archive/master/x264-33f9e1474613f59392be5ab6a7e7abf60fa63622.tar.gz)
+set(X264_URI http://download.videolan.org/pub/videolan/x264/snapshots/x264-snapshot-20180811-2245-stable.tar.bz2)
-set(X264_HASH 300dfb5b6c35722516f168868ce9419252a9e9eb77a05d82c9cede925b691bd6)
+set(X264_HASH ae8a868a0e236a348b35d79f3ee80294b169d1195408b689f9851383661ed7aa)
 
-set(XVIDCORE_VERSION 1.3.7)
+set(XVIDCORE_VERSION 1.3.5)
-set(XVIDCORE_URI https://downloads.xvid.com/downloads/xvidcore-${XVIDCORE_VERSION}.tar.gz)
+set(XVIDCORE_URI http://downloads.xvid.org/downloads/xvidcore-${XVIDCORE_VERSION}.tar.gz)
-set(XVIDCORE_HASH abbdcbd39555691dd1c9b4d08f0a031376a3b211652c0d8b3b8aa9be1303ce2d)
+set(XVIDCORE_HASH 165ba6a2a447a8375f7b06db5a3c91810181f2898166e7c8137401d7fc894cf0)
 
-set(OPENJPEG_VERSION 2.3.1)
+# This has to be in sync with the version in blenders /extern folder.
+set(OPENJPEG_VERSION 2.3.0)
 set(OPENJPEG_SHORT_VERSION 2.3)
-set(OPENJPEG_URI https://github.com/uclouvain/openjpeg/archive/v${OPENJPEG_VERSION}.tar.gz)
+# Use slightly newer commit after release which includes a cmake fix
-set(OPENJPEG_HASH 63f5a4713ecafc86de51bfad89cc07bb788e9bba24ebbf0c4ca637621aadb6a9)
+set(OPENJPEG_URI https://github.com/uclouvain/openjpeg/archive/66297f07a43.zip)
+set(OPENJPEG_HASH 8242b18d908c7c42174e4231a741cfa7ce7c26b6ed5c9644feb9df7b3054310b)
 
-set(FFMPEG_VERSION 4.2.3)
+set(FAAD_VERSION 2-2.8.8)
+set(FAAD_URI http://downloads.sourceforge.net/faac/faad${FAAD_VERSION}.tar.gz)
+set(FAAD_HASH 28f6116efdbe9378269f8a6221767d1f)
+
+set(FFMPEG_VERSION 4.0.2)
 set(FFMPEG_URI http://ffmpeg.org/releases/ffmpeg-${FFMPEG_VERSION}.tar.bz2)
-set(FFMPEG_HASH 695fad11f3baf27784e24cb0e977b65a)
+set(FFMPEG_HASH 5576e8a22f80b6a336db39808f427cfb)
 
 set(FFTW_VERSION 3.3.8)
 set(FFTW_URI http://www.fftw.org/fftw-${FFTW_VERSION}.tar.gz)
 set(FFTW_HASH 8aac833c943d8e90d51b697b27d4384d)
 
-set(ICONV_VERSION 1.16)
+set(ICONV_VERSION 1.15)
 set(ICONV_URI http://ftp.gnu.org/pub/gnu/libiconv/libiconv-${ICONV_VERSION}.tar.gz)
-set(ICONV_HASH 7d2a800b952942bb2880efb00cfd524c)
+set(ICONV_HASH ace8b5f2db42f7b3b3057585e80d9808)
 
 set(LAPACK_VERSION 3.6.0)
 set(LAPACK_URI http://www.netlib.org/lapack/lapack-${LAPACK_VERSION}.tgz)
@@ -261,9 +267,9 @@ set(LCMS_VERSION 2.9)
 set(LCMS_URI https://nchc.dl.sourceforge.net/project/lcms/lcms/${LCMS_VERSION}/lcms2-${LCMS_VERSION}.tar.gz)
 set(LCMS_HASH 8de1b7724f578d2995c8fdfa35c3ad0e)
 
-set(PUGIXML_VERSION 1.10)
+set(PUGIXML_VERSION 1.9)
-set(PUGIXML_URI https://github.com/zeux/pugixml/archive/v${PUGIXML_VERSION}.tar.gz)
+set(PUGIXML_URI https://github.com/zeux/pugixml/archive/v1.9.tar.gz)
-set(PUGIXML_HASH 0c208b0664c7fb822bf1b49ad035e8fd)
+set(PUGIXML_HASH 9346ca1dce2c48f1748c12fdac41a714)
 
 set(FLEXBISON_VERSION 2.5.5)
 set(FLEXBISON_URI http://prdownloads.sourceforge.net/winflexbison//win_flex_bison-2.5.5.zip)
@@ -293,9 +299,9 @@ set(SQLITE_VERSION 3.24.0)
 set(SQLITE_URI https://www.sqlite.org/2018/sqlite-src-3240000.zip)
 set(SQLITE_HASH fb558c49ee21a837713c4f1e7e413309aabdd9c7)
 
-set(EMBREE_VERSION 3.10.0)
+set(EMBREE_VERSION 3.8.0)
 set(EMBREE_URI https://github.com/embree/embree/archive/v${EMBREE_VERSION}.zip)
-set(EMBREE_HASH 4bbe29e7eaa46417efc75fc5f1e8eb87)
+set(EMBREE_HASH ac504d5426945fe25dec1267e0c39d52)
 
 set(USD_VERSION 19.11)
 set(USD_URI https://github.com/PixarAnimationStudios/USD/archive/v${USD_VERSION}.tar.gz)
@@ -313,6 +319,6 @@ set(MESA_VERSION 18.3.1)
 set(MESA_URI ftp://ftp.freedesktop.org/pub/mesa//mesa-${MESA_VERSION}.tar.xz)
 set(MESA_HASH d60828056d77bfdbae0970f9b15fb1be)
 
-set(XR_OPENXR_SDK_VERSION 1.0.8)
+set(XR_OPENXR_SDK_VERSION 1.0.6)
 set(XR_OPENXR_SDK_URI https://github.com/KhronosGroup/OpenXR-SDK/archive/release-${XR_OPENXR_SDK_VERSION}.tar.gz)
-set(XR_OPENXR_SDK_HASH c6de63d2e0f9029aa58dfa97cad8ce07)
+set(XR_OPENXR_SDK_HASH 21daea7c3bfec365298d779a0e19caa1)

build_files/build_environment/cmake/x264.cmake

@@ -18,6 +18,9 @@
 
 if(WIN32)
   set(X264_EXTRA_ARGS --enable-win32thread --cross-prefix=${MINGW_HOST}- --host=${MINGW_HOST})
+  set(X264_PATCH_CMD ${PATCH_CMD} --verbose -p 1 -N -d ${BUILD_DIR}/x264/src/external_x264 < ${PATCH_DIR}/x264.diff)
+else()
+  set(X264_PATCH_CMD echo .)
 endif()
 
 
@@ -26,6 +29,7 @@ ExternalProject_Add(external_x264
   DOWNLOAD_DIR ${DOWNLOAD_DIR}
   URL_HASH SHA256=${X264_HASH}
   PREFIX ${BUILD_DIR}/x264
+  PATCH_COMMAND ${X264_PATCH_CMD}
   CONFIGURE_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/x264/src/external_x264/ && ${CONFIGURE_COMMAND} --prefix=${LIBDIR}/x264
     --enable-static
     --enable-pic

build_files/build_environment/install_deps.sh

@@ -378,7 +378,6 @@ CLANG_FORMAT_VERSION_MIN="6.0"
 
 PYTHON_VERSION="3.7.4"
 PYTHON_VERSION_MIN="3.7"
-PYTHON_VERSION_INSTALLED=$PYTHON_VERSION_MIN
 PYTHON_FORCE_BUILD=false
 PYTHON_FORCE_REBUILD=false
 PYTHON_SKIP=false
@@ -437,7 +436,7 @@ OSL_FORCE_REBUILD=false
 OSL_SKIP=false
 
 # OpenSubdiv needs to be compiled for now
-OSD_VERSION="3.4.3"
+OSD_VERSION="3.4.0_RC2"
 OSD_VERSION_MIN=$OSD_VERSION
 OSD_FORCE_BUILD=false
 OSD_FORCE_REBUILD=false
@@ -469,7 +468,7 @@ OPENCOLLADA_FORCE_BUILD=false
 OPENCOLLADA_FORCE_REBUILD=false
 OPENCOLLADA_SKIP=false
 
-EMBREE_VERSION="3.10.0"
+EMBREE_VERSION="3.8.0"
 EMBREE_FORCE_BUILD=false
 EMBREE_FORCE_REBUILD=false
 EMBREE_SKIP=false
@@ -479,14 +478,14 @@ OIDN_FORCE_BUILD=false
 OIDN_FORCE_REBUILD=false
 OIDN_SKIP=false
 
-FFMPEG_VERSION="4.2.3"
+FFMPEG_VERSION="4.0.2"
 FFMPEG_VERSION_MIN="2.8.4"
 FFMPEG_FORCE_BUILD=false
 FFMPEG_FORCE_REBUILD=false
 FFMPEG_SKIP=false
 _ffmpeg_list_sep=";"
 
-XR_OPENXR_VERSION="1.0.8"
+XR_OPENXR_VERSION="1.0.6"
 XR_OPENXR_FORCE_BUILD=false
 XR_OPENXR_FORCE_REBUILD=false
 XR_OPENXR_SKIP=false
@@ -1030,7 +1029,7 @@ Those libraries should be available as packages in all recent distributions (opt
     * libjpeg, libpng, libtiff, [openjpeg2], [libopenal].
     * libx11, libxcursor, libxi, libxrandr, libxinerama (and other libx... as needed).
     * libsqlite3, libbz2, libssl, libfftw3, libxml2, libtinyxml, yasm, libyaml-cpp.
-    * libsdl2, libglew, [libglewmx].\""
+    * libsdl1.2, libglew, [libglewmx].\""
 
 DEPS_SPECIFIC_INFO="\"BUILDABLE DEPENDENCIES:
 
@@ -1536,13 +1535,11 @@ _update_deps_tbb() {
   OSD_FORCE_REBUILD=true
   OPENVDB_FORCE_REBUILD=true
   USD_FORCE_REBUILD=true
-  EMBREE_FORCE_REBUILD=true
   OIDN_FORCE_REBUILD=true
   if [ "$_is_building" = true ]; then
     OSD_FORCE_BUILD=true
     OPENVDB_FORCE_BUILD=true
     USD_FORCE_BUILD=true
-    EMBREE_FORCE_BUILD=true
     OIDN_FORCE_BUILD=true
   fi
 }
@@ -2988,7 +2985,7 @@ compile_Embree() {
   fi
 
   # To be changed each time we make edits that would modify the compiled results!
-  embree_magic=10
+  embree_magic=9
   _init_embree
 
   # Clean install if needed!
@@ -3042,13 +3039,9 @@ compile_Embree() {
     cmake_d="$cmake_d -D EMBREE_RAY_MASK=ON"
     cmake_d="$cmake_d -D EMBREE_FILTER_FUNCTION=ON"
     cmake_d="$cmake_d -D EMBREE_BACKFACE_CULLING=OFF"
+    cmake_d="$cmake_d -D EMBREE_TASKING_SYSTEM=INTERNAL"
     cmake_d="$cmake_d -D EMBREE_MAX_ISA=AVX2"
 
-    cmake_d="$cmake_d -D EMBREE_TASKING_SYSTEM=TBB"
-    if [ -d $INST/tbb ]; then
-      make_d="$make_d EMBREE_TBB_ROOT=$INST/tbb"
-    fi
-
     cmake $cmake_d ../
 
     make -j$THREADS && make install
@@ -3524,7 +3517,7 @@ install_DEB() {
              libxcursor-dev libxi-dev wget libsqlite3-dev libxrandr-dev libxinerama-dev \
              libbz2-dev libncurses5-dev libssl-dev liblzma-dev libreadline-dev \
              libopenal-dev libglew-dev yasm $THEORA_DEV $VORBIS_DEV $OGG_DEV \
-             libsdl2-dev libfftw3-dev patch bzip2 libxml2-dev libtinyxml-dev libjemalloc-dev"
+             libsdl1.2-dev libfftw3-dev patch bzip2 libxml2-dev libtinyxml-dev libjemalloc-dev"
              # libglewmx-dev  (broken in deb testing currently...)
 
   VORBIS_USE=true
@@ -3675,11 +3668,9 @@ install_DEB() {
     INFO "Forced Python/NumPy building, as requested..."
     _do_compile_python=true
   else
-    check_package_version_ge_DEB python3-dev $PYTHON_VERSION_MIN
+    check_package_DEB python$PYTHON_VERSION_MIN-dev
     if [ $? -eq 0 ]; then
-      PYTHON_VERSION_INSTALLED=$(echo `get_package_version_DEB python3-dev` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
+      install_packages_DEB python$PYTHON_VERSION_MIN-dev
-      
-      install_packages_DEB python3-dev
       clean_Python
       PRINT ""
       if [ "$NUMPY_SKIP" = true ]; then
@@ -4172,7 +4163,7 @@ install_RPM() {
   THEORA_DEV="libtheora-devel"
 
   _packages="gcc gcc-c++ git make cmake tar bzip2 xz findutils flex bison \
-             libtiff-devel libjpeg-devel libpng-devel sqlite-devel fftw-devel SDL2-devel \
+             libtiff-devel libjpeg-devel libpng-devel sqlite-devel fftw-devel SDL-devel \
              libX11-devel libXi-devel libXcursor-devel libXrandr-devel libXinerama-devel \
              wget ncurses-devel readline-devel $OPENJPEG_DEV openal-soft-devel \
              glew-devel yasm $THEORA_DEV $VORBIS_DEV $OGG_DEV patch \
@@ -4299,10 +4290,8 @@ install_RPM() {
     INFO "Forced Python/NumPy building, as requested..."
     _do_compile_python=true
   else
-    check_package_version_ge_RPM python3-devel $PYTHON_VERSION_MIN
+    check_package_version_match_RPM python3-devel $PYTHON_VERSION_MIN
     if [ $? -eq 0 ]; then
-      PYTHON_VERSION_INSTALLED=$(echo `get_package_version_RPM python3-devel` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
-
       install_packages_RPM python3-devel
       clean_Python
       PRINT ""
@@ -4735,7 +4724,7 @@ install_ARCH() {
 
   _packages="$BASE_DEVEL git cmake \
              libxi libxcursor libxrandr libxinerama glew libpng libtiff wget openal \
-             $OPENJPEG_DEV $VORBIS_DEV $OGG_DEV $THEORA_DEV yasm sdl2 fftw \
+             $OPENJPEG_DEV $VORBIS_DEV $OGG_DEV $THEORA_DEV yasm sdl fftw \
              libxml2 yaml-cpp tinyxml python-requests jemalloc"
 
   OPENJPEG_USE=true
@@ -4826,8 +4815,6 @@ install_ARCH() {
   else
     check_package_version_ge_ARCH python $PYTHON_VERSION_MIN
     if [ $? -eq 0 ]; then
-      PYTHON_VERSION_INSTALLED=$(echo `get_package_version_ARCH python` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
-
       install_packages_ARCH python
       clean_Python
       PRINT ""
@@ -5418,11 +5405,11 @@ print_info() {
   PRINT "  $_1"
   _buildargs="$_buildargs $_1"
 
-  _1="-D PYTHON_VERSION=$PYTHON_VERSION_INSTALLED"
+  _1="-D PYTHON_VERSION=$PYTHON_VERSION_MIN"
   PRINT "  $_1"
   _buildargs="$_buildargs $_1"
-  if [ -d "$INST/python-$PYTHON_VERSION_INSTALLED" ]; then
+  if [ -d "$INST/python-$PYTHON_VERSION_MIN" ]; then
-    _1="-D PYTHON_ROOT_DIR=$INST/python-$PYTHON_VERSION_INSTALLED"
+    _1="-D PYTHON_ROOT_DIR=$INST/python-$PYTHON_VERSION_MIN"
     PRINT "  $_1"
     _buildargs="$_buildargs $_1"
   fi

build_files/build_environment/patches/openal.diff

@@ -0,0 +1,13 @@
+diff -Naur external_openal_original/CMakeLists.txt external_openal/CMakeLists.txt
+--- external_openal_original/CMakeLists.txt	2016-01-24 20:12:39 -0700
++++ external_openal/CMakeLists.txt	2018-06-02 12:16:52 -0600
+@@ -885,7 +885,8 @@
+ OPTION(ALSOFT_REQUIRE_MMDEVAPI "Require MMDevApi backend" OFF)
+ IF(HAVE_WINDOWS_H)
+     # Check MMSystem backend
+-    CHECK_INCLUDE_FILES("windows.h;mmsystem.h" HAVE_MMSYSTEM_H -D_WIN32_WINNT=0x0502)
++    set(CMAKE_REQUIRED_FLAGS "-D_WIN32_WINNT=0x0502")
++    CHECK_INCLUDE_FILES("windows.h;mmsystem.h" HAVE_MMSYSTEM_H)
+     IF(HAVE_MMSYSTEM_H)
+         CHECK_SHARED_FUNCTION_EXISTS(waveOutOpen "windows.h;mmsystem.h" winmm "" HAVE_LIBWINMM)
+         IF(HAVE_LIBWINMM)

build_files/build_environment/patches/openimageio.diff

@@ -1,36 +0,0 @@
-diff -Naur orig/CMakeLists.txt external_openimageio/CMakeLists.txt
---- orig/CMakeLists.txt	2020-05-10 21:43:52 -0600
-+++ external_openimageio/CMakeLists.txt	2020-05-13 17:03:35 -0600
-@@ -170,7 +170,7 @@
-     add_subdirectory (src/iinfo)
-     add_subdirectory (src/maketx)
-     add_subdirectory (src/oiiotool)
--    add_subdirectory (src/testtex)
-+    #add_subdirectory (src/testtex)
-     add_subdirectory (src/iv)
- endif ()
- 
-diff -Naur orig/src/cmake/compiler.cmake external_openimageio/src/cmake/compiler.cmake
---- orig/src/cmake/compiler.cmake	2020-05-10 21:43:52 -0600
-+++ external_openimageio/src/cmake/compiler.cmake	2020-05-13 17:02:54 -0600
-@@ -172,6 +172,7 @@
-     add_definitions (-D_CRT_NONSTDC_NO_WARNINGS)
-     add_definitions (-D_SCL_SECURE_NO_WARNINGS)
-     add_definitions (-DJAS_WIN_MSVC_BUILD)
-+    add_definitions (-DOPJ_STATIC)
- endif (MSVC)
- 
- if (${CMAKE_SYSTEM_NAME} STREQUAL "FreeBSD"
-diff -Naur orig/src/include/OpenImageIO/platform.h external_openimageio/src/include/OpenImageIO/platform.h
---- orig/src/include/OpenImageIO/platform.h	2020-05-10 21:43:52 -0600
-+++ external_openimageio/src/include/OpenImageIO/platform.h	2020-05-13 17:04:36 -0600
-@@ -41,6 +41,9 @@
- #    ifndef NOMINMAX
- #        define NOMINMAX
- #    endif
-+#    ifndef NOGDI
-+#        define NOGDI
-+#    endif
- #    include <windows.h>
- #endif
- 

build_files/build_environment/patches/openimageio_gdi.diff

@@ -0,0 +1,26 @@
+Index: OpenImageIO/osdep.h
+===================================================================
+--- OpenImageIO/osdep.h	(revision 61595)
++++ OpenImageIO/osdep.h	(working copy)
+@@ -34,6 +34,7 @@
+ # define WIN32_LEAN_AND_MEAN
+ # define VC_EXTRALEAN
+ # define NOMINMAX
++# define NOGDI
+ # include <windows.h>
+ #endif
+ 
+Index: OpenImageIO/platform.h
+===================================================================
+--- OpenImageIO/platform.h	(revision 61595)
++++ OpenImageIO/platform.h	(working copy)
+@@ -77,6 +77,9 @@
+ # ifndef NOMINMAX
+ #   define NOMINMAX
+ # endif
++# ifndef NOGDI
++#   define NOGDI
++# endif
+ # include <windows.h>
+ #endif
+ 

build_files/build_environment/patches/openimageio_static_libs.diff

@@ -0,0 +1,22 @@
+diff -Naur external_openimageio.orig/src/cmake/compiler.cmake external_openimageio/src/cmake/compiler.cmake
+--- external_openimageio.orig/src/cmake/compiler.cmake	2018-07-31 23:45:19 -0600
++++ external_openimageio/src/cmake/compiler.cmake	2018-08-16 12:50:12 -0600
+@@ -152,6 +152,7 @@
+     add_definitions (-D_CRT_NONSTDC_NO_WARNINGS)
+     add_definitions (-D_SCL_SECURE_NO_WARNINGS)
+     add_definitions (-DJAS_WIN_MSVC_BUILD)
++    add_definitions (-DOPJ_STATIC)
+ endif (MSVC)
+ 
+ # Use ccache if found
+--- external_openimageio.orig/CMakeLists.txt	2018-07-31 23:45:19 -0600
++++ external_openimageio/CMakeLists.txt	2018-08-17 15:22:56 -0600
+@@ -169,7 +169,7 @@
+   add_subdirectory (src/iinfo)
+   add_subdirectory (src/maketx)
+   add_subdirectory (src/oiiotool)
+-  add_subdirectory (src/testtex)
++# add_subdirectory (src/testtex)
+   add_subdirectory (src/iv)
+ endif ()
+ 

build_files/build_environment/patches/osl.diff

@@ -48,22 +48,3 @@ diff -Naur OpenShadingLanguage-Release-1.9.9/src/liboslexec/llvm_util.cpp extern
  
  size_t
  LLVM_Util::total_jit_memory_held ()
-diff -Naur OpenShadingLanguage-Release-1.9.9/CMakeLists.txt external_osl/CMakeLists.txt
- --- orig/CMakeLists.txt	2020-01-27 16:22:31 -0700
-+++ external_osl/CMakeLists.txt	2020-05-13 18:04:52 -0600
-@@ -102,10 +102,11 @@
- set (OPTIX_EXTRA_LIBS CACHE STRING "Extra lib targets needed for OptiX")
- set (CUDA_TARGET_ARCH "sm_35" CACHE STRING "CUDA GPU architecture (e.g. sm_35)")
- 
--# set (USE_OIIO_STATIC ON CACHE BOOL "If OIIO is built static")
--# if (USE_OIIO_STATIC)
--#     add_definitions ("-DOIIO_STATIC_BUILD=1")
--# endif ()
-+set (USE_OIIO_STATIC ON CACHE BOOL "If OIIO is built static")
-+if (USE_OIIO_STATIC)
-+    add_definitions ("-DOIIO_STATIC_BUILD=1")
-+    add_definitions ("-DOIIO_STATIC_DEFINE=1")
-+endif ()
- 
- set (OSL_NO_DEFAULT_TEXTURESYSTEM OFF CACHE BOOL "Do not use create a raw OIIO::TextureSystem")
- if (OSL_NO_DEFAULT_TEXTURESYSTEM)

build_files/build_environment/patches/x264.diff

@@ -0,0 +1,22 @@
+--- x264-snapshot-20180811-2245-stable\configure	2018-08-11 14:45:05 -0600
++++ external_x264\configure	2018-08-11 23:51:35 -0600
+@@ -396,7 +396,7 @@
+ # list of all preprocessor HAVE values we can define
+ CONFIG_HAVE="MALLOC_H ALTIVEC ALTIVEC_H MMX ARMV6 ARMV6T2 NEON BEOSTHREAD POSIXTHREAD WIN32THREAD THREAD LOG2F SWSCALE \
+              LAVF FFMS GPAC AVS GPL VECTOREXT INTERLACED CPU_COUNT OPENCL THP LSMASH X86_INLINE_ASM AS_FUNC INTEL_DISPATCHER \
+-             MSA MMAP WINRT VSX ARM_INLINE_ASM STRTOK_R BITDEPTH8 BITDEPTH10"
++             MSA MMAP WINRT VSX ARM_INLINE_ASM BITDEPTH8 BITDEPTH10"
+ 
+ # parse options
+ 
+@@ -1071,10 +1071,6 @@
+     define HAVE_LOG2F
+ fi
+ 
+-if cc_check 'string.h' '' 'strtok_r(0, 0, 0);' ; then
+-    define HAVE_STRTOK_R
+-fi
+-
+ if [ "$SYS" != "WINDOWS" ] && cpp_check "sys/mman.h unistd.h" "" "defined(MAP_PRIVATE)"; then
+     define HAVE_MMAP
+ fi

build_files/buildbot/buildbot_utils.py

@@ -78,26 +78,27 @@ class VersionInfo:
         blender_h = os.path.join(builder.blender_dir, "source", "blender", "blenkernel", "BKE_blender_version.h")
 
         version_number = int(self._parse_header_file(blender_h, 'BLENDER_VERSION'))
-        version_number_patch = int(self._parse_header_file(blender_h, 'BLENDER_VERSION_PATCH'))
+        self.version = "%d.%d" % (version_number // 100, version_number % 100)
-        version_numbers = (version_number // 100, version_number % 100, version_number_patch)
+        self.version_char = self._parse_header_file(blender_h, 'BLENDER_VERSION_CHAR')
-        self.short_version = "%d.%02d" % (version_numbers[0], version_numbers[1])
-        self.version = "%d.%02d.%d" % version_numbers
         self.version_cycle = self._parse_header_file(blender_h, 'BLENDER_VERSION_CYCLE')
         self.version_cycle_number = self._parse_header_file(blender_h, 'BLENDER_VERSION_CYCLE_NUMBER')
         self.hash = self._parse_header_file(buildinfo_h, 'BUILD_HASH')[1:-1]
 
         if self.version_cycle == "release":
             # Final release
-            self.full_version = self.version
+            self.full_version = self.version + self.version_char
             self.is_development_build = False
         elif self.version_cycle == "rc":
             # Release candidate
             version_cycle = self.version_cycle + self.version_cycle_number
-            self.full_version = self.version + version_cycle
+            if len(self.version_char) == 0:
+                self.full_version = self.version + version_cycle
+            else:
+                self.full_version = self.version + self.version_char + '-' + version_cycle
             self.is_development_build = False
         else:
             # Development build
-            self.full_version = self.version + '-' + self.hash
+            self.full_version = self.version + self.version_char + '-' + self.hash
             self.is_development_build = True
 
     def _parse_header_file(self, filename, define):

build_files/buildbot/slave_pack.py

@@ -167,7 +167,7 @@ def pack_linux(builder):
     buildbot_utils.call(builder.command_prefix + ['strip', '--strip-all', blender_executable])
 
     print("Stripping python...")
-    py_target = os.path.join(builder.install_dir, info.short_version)
+    py_target = os.path.join(builder.install_dir, info.version)
     buildbot_utils.call(builder.command_prefix + ['find', py_target, '-iname', '*.so', '-exec', 'strip', '-s', '{}', ';'])
 
     # Construct package name

build_files/cmake/config/blender_lite.cmake

@@ -52,7 +52,6 @@ set(WITH_OPENVDB             OFF CACHE BOOL "" FORCE)
 set(WITH_QUADRIFLOW          OFF CACHE BOOL "" FORCE)
 set(WITH_SDL                 OFF CACHE BOOL "" FORCE)
 set(WITH_TBB                 OFF CACHE BOOL "" FORCE)
-set(WITH_USD                 OFF CACHE BOOL "" FORCE)
 
 if(UNIX AND NOT APPLE)
   set(WITH_GHOST_XDND          OFF CACHE BOOL "" FORCE)

build_files/cmake/macros.cmake

@@ -756,7 +756,8 @@ function(get_blender_version)
   # - BLENDER_VERSION (major.minor)
   # - BLENDER_VERSION_MAJOR
   # - BLENDER_VERSION_MINOR
-  # - BLENDER_VERSION_PATCH
+  # - BLENDER_SUBVERSION (used for internal versioning mainly)
+  # - BLENDER_VERSION_CHAR (a, b, c, ...or empty string)
   # - BLENDER_VERSION_CYCLE (alpha, beta, rc, release)
 
   # So cmake depends on BKE_blender.h, beware of inf-loops!
@@ -766,15 +767,25 @@ function(get_blender_version)
   file(STRINGS ${CMAKE_SOURCE_DIR}/source/blender/blenkernel/BKE_blender_version.h _contents REGEX "^#define[ \t]+BLENDER_.*$")
 
   string(REGEX REPLACE ".*#define[ \t]+BLENDER_VERSION[ \t]+([0-9]+).*" "\\1" _out_version "${_contents}")
-  string(REGEX REPLACE ".*#define[ \t]+BLENDER_VERSION_PATCH[ \t]+([0-9]+).*" "\\1" _out_version_patch "${_contents}")
+  string(REGEX REPLACE ".*#define[ \t]+BLENDER_SUBVERSION[ \t]+([0-9]+).*" "\\1" _out_subversion "${_contents}")
+  string(REGEX REPLACE ".*#define[ \t]+BLENDER_VERSION_CHAR[ \t]+([a-z]+).*" "\\1" _out_version_char "${_contents}")
   string(REGEX REPLACE ".*#define[ \t]+BLENDER_VERSION_CYCLE[ \t]+([a-z]+).*" "\\1" _out_version_cycle "${_contents}")
 
   if(NOT ${_out_version} MATCHES "[0-9]+")
     message(FATAL_ERROR "Version parsing failed for BLENDER_VERSION")
   endif()
 
-  if(NOT ${_out_version_patch} MATCHES "[0-9]+")
+  if(NOT ${_out_subversion} MATCHES "[0-9]+")
-    message(FATAL_ERROR "Version parsing failed for BLENDER_VERSION_PATCH")
+    message(FATAL_ERROR "Version parsing failed for BLENDER_SUBVERSION")
+  endif()
+
+  # clumsy regex, only single char are ok but it could be unset
+
+  string(LENGTH "${_out_version_char}" _out_version_char_len)
+  if(NOT _out_version_char_len EQUAL 1)
+    set(_out_version_char "")
+  elseif(NOT ${_out_version_char} MATCHES "[a-z]+")
+    message(FATAL_ERROR "Version parsing failed for BLENDER_VERSION_CHAR")
   endif()
 
   if(NOT ${_out_version_cycle} MATCHES "[a-z]+")
@@ -784,11 +795,23 @@ function(get_blender_version)
   math(EXPR _out_version_major "${_out_version} / 100")
   math(EXPR _out_version_minor "${_out_version} % 100")
 
+  # for packaging, alpha to numbers
+  string(COMPARE EQUAL "${_out_version_char}" "" _out_version_char_empty)
+  if(${_out_version_char_empty})
+    set(_out_version_char_index "0")
+  else()
+    set(_char_ls a b c d e f g h i j k l m n o p q r s t u v w x y z)
+    list(FIND _char_ls ${_out_version_char} _out_version_char_index)
+    math(EXPR _out_version_char_index "${_out_version_char_index} + 1")
+  endif()
+
   # output vars
   set(BLENDER_VERSION "${_out_version_major}.${_out_version_minor}" PARENT_SCOPE)
   set(BLENDER_VERSION_MAJOR "${_out_version_major}" PARENT_SCOPE)
   set(BLENDER_VERSION_MINOR "${_out_version_minor}" PARENT_SCOPE)
-  set(BLENDER_VERSION_PATCH "${_out_version_patch}" PARENT_SCOPE)
+  set(BLENDER_SUBVERSION "${_out_subversion}" PARENT_SCOPE)
+  set(BLENDER_VERSION_CHAR "${_out_version_char}" PARENT_SCOPE)
+  set(BLENDER_VERSION_CHAR_INDEX "${_out_version_char_index}" PARENT_SCOPE)
   set(BLENDER_VERSION_CYCLE "${_out_version_cycle}" PARENT_SCOPE)
 
 endfunction()

build_files/cmake/packaging.cmake

@@ -7,7 +7,7 @@ set(PROJECT_VENDOR       "Blender Foundation")
 
 set(MAJOR_VERSION ${BLENDER_VERSION_MAJOR})
 set(MINOR_VERSION ${BLENDER_VERSION_MINOR})
-set(PATCH_VERSION ${BLENDER_VERSION_PATCH})
+set(PATCH_VERSION ${BLENDER_VERSION_CHAR_INDEX})
 
 set(CPACK_SYSTEM_NAME ${CMAKE_SYSTEM_NAME})
 set(CPACK_PACKAGE_DESCRIPTION ${PROJECT_DESCRIPTION})

build_files/cmake/platform/platform_win32.cmake

@@ -190,7 +190,7 @@ if(MSVC_VERSION GREATER 1914 AND NOT MSVC_CLANG)
   set(CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} /JMC")
 endif()
 
-set(PLATFORM_LINKFLAGS "${PLATFORM_LINKFLAGS} /SUBSYSTEM:CONSOLE /STACK:2097152")
+set(PLATFORM_LINKFLAGS "${PLATFORM_LINKFLAGS} /SUBSYSTEM:CONSOLE /STACK:2097152 ")
 set(PLATFORM_LINKFLAGS_RELEASE "/NODEFAULTLIB:libcmt.lib /NODEFAULTLIB:libcmtd.lib /NODEFAULTLIB:msvcrtd.lib")
 set(PLATFORM_LINKFLAGS_DEBUG "${PLATFORM_LINKFLAGS_DEBUG} /IGNORE:4099 /NODEFAULTLIB:libcmt.lib /NODEFAULTLIB:msvcrt.lib /NODEFAULTLIB:libcmtd.lib")
 
@@ -478,7 +478,7 @@ if(WITH_OPENIMAGEIO)
   set(OPENIMAGEIO_DEFINITIONS "-DUSE_TBB=0")
   set(OPENCOLORIO_DEFINITIONS "-DOCIO_STATIC_BUILD")
   set(OPENIMAGEIO_IDIFF "${OPENIMAGEIO}/bin/idiff.exe")
-  add_definitions(-DOIIO_STATIC_DEFINE)
+  add_definitions(-DOIIO_STATIC_BUILD)
   add_definitions(-DOIIO_NO_SSE=1)
 endif()
 

build_files/utils/make_source_archive.sh

@@ -7,14 +7,15 @@ BASE_DIR="$PWD"
 
 blender_srcdir=$(dirname -- $0)/../..
 blender_version=$(grep "BLENDER_VERSION\s" "$blender_srcdir/source/blender/blenkernel/BKE_blender_version.h" | awk '{print $3}')
-blender_version_patch=$(grep "BLENDER_VERSION_PATCH\s" "$blender_srcdir/source/blender/blenkernel/BKE_blender_version.h" | awk '{print $3}')
+blender_version_char=$(grep "BLENDER_VERSION_CHAR\s" "$blender_srcdir/source/blender/blenkernel/BKE_blender_version.h" | awk '{print $3}')
 blender_version_cycle=$(grep "BLENDER_VERSION_CYCLE\s" "$blender_srcdir/source/blender/blenkernel/BKE_blender_version.h" | awk '{print $3}')
+blender_subversion=$(grep "BLENDER_SUBVERSION\s" "$blender_srcdir/source/blender/blenkernel/BKE_blender_version.h" | awk '{print $3}')
 
-VERSION=$(expr $blender_version / 100).$(expr $blender_version % 100).$blender_version_patch
 if [ "$blender_version_cycle" = "release" ] ; then
+  VERSION=$(expr $blender_version / 100).$(expr $blender_version % 100)$blender_version_char
   SUBMODULE_EXCLUDE="^\(release/scripts/addons_contrib\)$"
 else
-  VERSION=$VERSION-$blender_version_cycle
+  VERSION=$(expr $blender_version / 100).$(expr $blender_version % 100)_$blender_subversion
   SUBMODULE_EXCLUDE="^$"  # dummy regex
 fi
 

doc/python_api/examples/gpu.1.py

@@ -119,24 +119,4 @@ Examples
 To try these examples, just copy them into Blenders text editor and execute them.
 To keep the examples relatively small, they just register a draw function that can't easily be removed anymore.
 Blender has to be restarted in order to delete the draw handlers.
-
-3D Lines with Single Color
---------------------------
 """
-
-import bpy
-import gpu
-from gpu_extras.batch import batch_for_shader
-
-coords = [(1, 1, 1), (-2, 0, 0), (-2, -1, 3), (0, 1, 1)]
-shader = gpu.shader.from_builtin('3D_UNIFORM_COLOR')
-batch = batch_for_shader(shader, 'LINES', {"pos": coords})
-
-
-def draw():
-    shader.bind()
-    shader.uniform_float("color", (1, 1, 0, 1))
-    batch.draw(shader)
-
-
-bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_VIEW')

doc/python_api/examples/gpu.10.py

@@ -1,65 +1,41 @@
 """
-Custom Shader for dotted 3D Line
+Rendering the 3D View into a Texture
---------------------------------
+------------------------------------
 
-In this example the arc length (distance to the first point on the line) is calculated in every vertex.
+The scene has to have a camera for this example to work.
-Between the vertex and fragment shader that value is automatically interpolated
+You could also make this independent of a specific camera,
-for all points that will be visible on the screen.
+but Blender does not expose good functions to create view and projection matrices yet.
-In the fragment shader the ``sin`` of the arc length is calculated.
-Based on the result a decision is made on whether the fragment should be drawn or not.
 """
 import bpy
+import bgl
 import gpu
-from random import random
+from gpu_extras.presets import draw_texture_2d
-from mathutils import Vector
-from gpu_extras.batch import batch_for_shader
 
-vertex_shader = '''
+WIDTH = 512
-    uniform mat4 u_ViewProjectionMatrix;
+HEIGHT = 256
 
-    in vec3 position;
+offscreen = gpu.types.GPUOffScreen(WIDTH, HEIGHT)
-    in float arcLength;
-
-    out float v_ArcLength;
-
-    void main()
-    {
-        v_ArcLength = arcLength;
-        gl_Position = u_ViewProjectionMatrix * vec4(position, 1.0f);
-    }
-'''
-
-fragment_shader = '''
-    uniform float u_Scale;
-
-    in float v_ArcLength;
-
-    void main()
-    {
-        if (step(sin(v_ArcLength * u_Scale), 0.5) == 1) discard;
-        gl_FragColor = vec4(1.0);
-    }
-'''
-
-coords = [Vector((random(), random(), random())) * 5 for _ in range(5)]
-
-arc_lengths = [0]
-for a, b in zip(coords[:-1], coords[1:]):
-    arc_lengths.append(arc_lengths[-1] + (a - b).length)
-
-shader = gpu.types.GPUShader(vertex_shader, fragment_shader)
-batch = batch_for_shader(
-    shader, 'LINE_STRIP',
-    {"position": coords, "arcLength": arc_lengths},
-)
 
 
 def draw():
-    shader.bind()
+    context = bpy.context
-    matrix = bpy.context.region_data.perspective_matrix
+    scene = context.scene
-    shader.uniform_float("u_ViewProjectionMatrix", matrix)
+
-    shader.uniform_float("u_Scale", 10)
+    view_matrix = scene.camera.matrix_world.inverted()
-    batch.draw(shader)
+
+    projection_matrix = scene.camera.calc_matrix_camera(
+        context.evaluated_depsgraph_get(), x=WIDTH, y=HEIGHT)
+
+    offscreen.draw_view3d(
+        scene,
+        context.view_layer,
+        context.space_data,
+        context.region,
+        view_matrix,
+        projection_matrix)
+
+    bgl.glDisable(bgl.GL_DEPTH_TEST)
+    draw_texture_2d(offscreen.color_texture, (10, 10), WIDTH, HEIGHT)
 
 
-bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_VIEW')
+bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_PIXEL')

doc/python_api/examples/gpu.11.py

@@ -0,0 +1,65 @@
+"""
+Custom Shader for dotted 3D Line
+--------------------------------
+
+In this example the arc length (distance to the first point on the line) is calculated in every vertex.
+Between the vertex and fragment shader that value is automatically interpolated
+for all points that will be visible on the screen.
+In the fragment shader the ``sin`` of the arc length is calculated.
+Based on the result a decision is made on whether the fragment should be drawn or not.
+"""
+import bpy
+import gpu
+from random import random
+from mathutils import Vector
+from gpu_extras.batch import batch_for_shader
+
+vertex_shader = '''
+    uniform mat4 u_ViewProjectionMatrix;
+
+    in vec3 position;
+    in float arcLength;
+
+    out float v_ArcLength;
+
+    void main()
+    {
+        v_ArcLength = arcLength;
+        gl_Position = u_ViewProjectionMatrix * vec4(position, 1.0f);
+    }
+'''
+
+fragment_shader = '''
+    uniform float u_Scale;
+
+    in float v_ArcLength;
+
+    void main()
+    {
+        if (step(sin(v_ArcLength * u_Scale), 0.5) == 1) discard;
+        gl_FragColor = vec4(1.0);
+    }
+'''
+
+coords = [Vector((random(), random(), random())) * 5 for _ in range(5)]
+
+arc_lengths = [0]
+for a, b in zip(coords[:-1], coords[1:]):
+    arc_lengths.append(arc_lengths[-1] + (a - b).length)
+
+shader = gpu.types.GPUShader(vertex_shader, fragment_shader)
+batch = batch_for_shader(
+    shader, 'LINE_STRIP',
+    {"position": coords, "arcLength": arc_lengths},
+)
+
+
+def draw():
+    shader.bind()
+    matrix = bpy.context.region_data.perspective_matrix
+    shader.uniform_float("u_ViewProjectionMatrix", matrix)
+    shader.uniform_float("u_Scale", 10)
+    batch.draw(shader)
+
+
+bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_VIEW')

doc/python_api/examples/gpu.2.py

@@ -1,45 +1,19 @@
 """
-Triangle with Custom Shader
+3D Lines with Single Color
----------------------------
+--------------------------
 """
 import bpy
 import gpu
 from gpu_extras.batch import batch_for_shader
 
-vertex_shader = '''
+coords = [(1, 1, 1), (-2, 0, 0), (-2, -1, 3), (0, 1, 1)]
-    uniform mat4 viewProjectionMatrix;
+shader = gpu.shader.from_builtin('3D_UNIFORM_COLOR')
-
+batch = batch_for_shader(shader, 'LINES', {"pos": coords})
-    in vec3 position;
-    out vec3 pos;
-
-    void main()
-    {
-        pos = position;
-        gl_Position = viewProjectionMatrix * vec4(position, 1.0f);
-    }
-'''
-
-fragment_shader = '''
-    uniform float brightness;
-
-    in vec3 pos;
-
-    void main()
-    {
-        gl_FragColor = vec4(pos * brightness, 1.0);
-    }
-'''
-
-coords = [(1, 1, 1), (2, 0, 0), (-2, -1, 3)]
-shader = gpu.types.GPUShader(vertex_shader, fragment_shader)
-batch = batch_for_shader(shader, 'TRIS', {"position": coords})
 
 
 def draw():
     shader.bind()
-    matrix = bpy.context.region_data.perspective_matrix
+    shader.uniform_float("color", (1, 1, 0, 1))
-    shader.uniform_float("viewProjectionMatrix", matrix)
-    shader.uniform_float("brightness", 0.5)
     batch.draw(shader)
 
 

doc/python_api/examples/gpu.3.py

@@ -1,29 +1,45 @@
 """
-Wireframe Cube using Index Buffer
+Triangle with Custom Shader
----------------------------------
+---------------------------
 """
 import bpy
 import gpu
 from gpu_extras.batch import batch_for_shader
 
-coords = (
+vertex_shader = '''
-    (-1, -1, -1), (+1, -1, -1),
+    uniform mat4 viewProjectionMatrix;
-    (-1, +1, -1), (+1, +1, -1),
-    (-1, -1, +1), (+1, -1, +1),
-    (-1, +1, +1), (+1, +1, +1))
 
-indices = (
+    in vec3 position;
-    (0, 1), (0, 2), (1, 3), (2, 3),
+    out vec3 pos;
-    (4, 5), (4, 6), (5, 7), (6, 7),
-    (0, 4), (1, 5), (2, 6), (3, 7))
 
-shader = gpu.shader.from_builtin('3D_UNIFORM_COLOR')
+    void main()
-batch = batch_for_shader(shader, 'LINES', {"pos": coords}, indices=indices)
+    {
+        pos = position;
+        gl_Position = viewProjectionMatrix * vec4(position, 1.0f);
+    }
+'''
+
+fragment_shader = '''
+    uniform float brightness;
+
+    in vec3 pos;
+
+    void main()
+    {
+        gl_FragColor = vec4(pos * brightness, 1.0);
+    }
+'''
+
+coords = [(1, 1, 1), (2, 0, 0), (-2, -1, 3)]
+shader = gpu.types.GPUShader(vertex_shader, fragment_shader)
+batch = batch_for_shader(shader, 'TRIS', {"position": coords})
 
 
 def draw():
     shader.bind()
-    shader.uniform_float("color", (1, 0, 0, 1))
+    matrix = bpy.context.region_data.perspective_matrix
+    shader.uniform_float("viewProjectionMatrix", matrix)
+    shader.uniform_float("brightness", 0.5)
     batch.draw(shader)
 
 

doc/python_api/examples/gpu.4.py

@@ -1,39 +1,30 @@
 """
-Mesh with Random Vertex Colors
+Wireframe Cube using Index Buffer
-------------------------------
+---------------------------------
 """
 import bpy
 import gpu
-import bgl
-import numpy as np
-from random import random
 from gpu_extras.batch import batch_for_shader
 
-mesh = bpy.context.active_object.data
+coords = (
-mesh.calc_loop_triangles()
+    (-1, -1, -1), (+1, -1, -1),
+    (-1, +1, -1), (+1, +1, -1),
+    (-1, -1, +1), (+1, -1, +1),
+    (-1, +1, +1), (+1, +1, +1))
 
-vertices = np.empty((len(mesh.vertices), 3), 'f')
+indices = (
-indices = np.empty((len(mesh.loop_triangles), 3), 'i')
+    (0, 1), (0, 2), (1, 3), (2, 3),
+    (4, 5), (4, 6), (5, 7), (6, 7),
+    (0, 4), (1, 5), (2, 6), (3, 7))
 
-mesh.vertices.foreach_get(
+shader = gpu.shader.from_builtin('3D_UNIFORM_COLOR')
-    "co", np.reshape(vertices, len(mesh.vertices) * 3))
+batch = batch_for_shader(shader, 'LINES', {"pos": coords}, indices=indices)
-mesh.loop_triangles.foreach_get(
-    "vertices", np.reshape(indices, len(mesh.loop_triangles) * 3))
-
-vertex_colors = [(random(), random(), random(), 1) for _ in range(len(mesh.vertices))]
-
-shader = gpu.shader.from_builtin('3D_SMOOTH_COLOR')
-batch = batch_for_shader(
-    shader, 'TRIS',
-    {"pos": vertices, "color": vertex_colors},
-    indices=indices,
-)
 
 
 def draw():
-    bgl.glEnable(bgl.GL_DEPTH_TEST)
+    shader.bind()
+    shader.uniform_float("color", (1, 0, 0, 1))
     batch.draw(shader)
-    bgl.glDisable(bgl.GL_DEPTH_TEST)
 
 
 bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_VIEW')

doc/python_api/examples/gpu.5.py

@@ -1,26 +1,39 @@
 """
-2D Rectangle
+Mesh with Random Vertex Colors
-------------
+------------------------------
 """
 import bpy
 import gpu
+import bgl
+import numpy as np
+from random import random
 from gpu_extras.batch import batch_for_shader
 
-vertices = (
+mesh = bpy.context.active_object.data
-    (100, 100), (300, 100),
+mesh.calc_loop_triangles()
-    (100, 200), (300, 200))
 
-indices = (
+vertices = np.empty((len(mesh.vertices), 3), 'f')
-    (0, 1, 2), (2, 1, 3))
+indices = np.empty((len(mesh.loop_triangles), 3), 'i')
 
-shader = gpu.shader.from_builtin('2D_UNIFORM_COLOR')
+mesh.vertices.foreach_get(
-batch = batch_for_shader(shader, 'TRIS', {"pos": vertices}, indices=indices)
+    "co", np.reshape(vertices, len(mesh.vertices) * 3))
+mesh.loop_triangles.foreach_get(
+    "vertices", np.reshape(indices, len(mesh.loop_triangles) * 3))
+
+vertex_colors = [(random(), random(), random(), 1) for _ in range(len(mesh.vertices))]
+
+shader = gpu.shader.from_builtin('3D_SMOOTH_COLOR')
+batch = batch_for_shader(
+    shader, 'TRIS',
+    {"pos": vertices, "color": vertex_colors},
+    indices=indices,
+)
 
 
 def draw():
-    shader.bind()
+    bgl.glEnable(bgl.GL_DEPTH_TEST)
-    shader.uniform_float("color", (0, 0.5, 0.5, 1.0))
     batch.draw(shader)
+    bgl.glDisable(bgl.GL_DEPTH_TEST)
 
 
-bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_PIXEL')
+bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_VIEW')

doc/python_api/examples/gpu.6.py

@@ -1,36 +1,25 @@
 """
-2D Image
+2D Rectangle
---------
+------------
-
-To use this example you have to provide an image that should be displayed.
 """
 import bpy
 import gpu
-import bgl
 from gpu_extras.batch import batch_for_shader
 
-IMAGE_NAME = "Untitled"
+vertices = (
-image = bpy.data.images[IMAGE_NAME]
+    (100, 100), (300, 100),
+    (100, 200), (300, 200))
 
-shader = gpu.shader.from_builtin('2D_IMAGE')
+indices = (
-batch = batch_for_shader(
+    (0, 1, 2), (2, 1, 3))
-    shader, 'TRI_FAN',
-    {
-        "pos": ((100, 100), (200, 100), (200, 200), (100, 200)),
-        "texCoord": ((0, 0), (1, 0), (1, 1), (0, 1)),
-    },
-)
 
-if image.gl_load():
+shader = gpu.shader.from_builtin('2D_UNIFORM_COLOR')
-    raise Exception()
+batch = batch_for_shader(shader, 'TRIS', {"pos": vertices}, indices=indices)
 
 
 def draw():
-    bgl.glActiveTexture(bgl.GL_TEXTURE0)
-    bgl.glBindTexture(bgl.GL_TEXTURE_2D, image.bindcode)
-
     shader.bind()
-    shader.uniform_int("image", 0)
+    shader.uniform_float("color", (0, 0.5, 0.5, 1.0))
     batch.draw(shader)
 
 

doc/python_api/examples/gpu.7.py

@@ -1,86 +1,37 @@
 """
-Generate a texture using Offscreen Rendering
+2D Image
---------------------------------------------
+--------
 
-#. Create an :class:`gpu.types.GPUOffScreen` object.
+To use this example you have to provide an image that should be displayed.
-#. Draw some circles into it.
-#. Make a new shader for drawing a planar texture in 3D.
-#. Draw the generated texture using the new shader.
 """
 import bpy
 import gpu
 import bgl
-from mathutils import Matrix
 from gpu_extras.batch import batch_for_shader
-from gpu_extras.presets import draw_circle_2d
 
-# Create and fill offscreen
+IMAGE_NAME = "Untitled"
-##########################################
+image = bpy.data.images[IMAGE_NAME]
 
-offscreen = gpu.types.GPUOffScreen(512, 512)
+shader = gpu.shader.from_builtin('2D_IMAGE')
-
-with offscreen.bind():
-    bgl.glClear(bgl.GL_COLOR_BUFFER_BIT)
-    with gpu.matrix.push_pop():
-        # reset matrices -> use normalized device coordinates [-1, 1]
-        gpu.matrix.load_matrix(Matrix.Identity(4))
-        gpu.matrix.load_projection_matrix(Matrix.Identity(4))
-
-        amount = 10
-        for i in range(-amount, amount + 1):
-            x_pos = i / amount
-            draw_circle_2d((x_pos, 0.0), (1, 1, 1, 1), 0.5, 200)
-
-
-# Drawing the generated texture in 3D space
-#############################################
-
-vertex_shader = '''
-    uniform mat4 modelMatrix;
-    uniform mat4 viewProjectionMatrix;
-
-    in vec2 position;
-    in vec2 uv;
-
-    out vec2 uvInterp;
-
-    void main()
-    {
-        uvInterp = uv;
-        gl_Position = viewProjectionMatrix * modelMatrix * vec4(position, 0.0, 1.0);
-    }
-'''
-
-fragment_shader = '''
-    uniform sampler2D image;
-
-    in vec2 uvInterp;
-
-    void main()
-    {
-        gl_FragColor = texture(image, uvInterp);
-    }
-'''
-
-shader = gpu.types.GPUShader(vertex_shader, fragment_shader)
 batch = batch_for_shader(
     shader, 'TRI_FAN',
     {
-        "position": ((-1, -1), (1, -1), (1, 1), (-1, 1)),
+        "pos": ((100, 100), (200, 100), (200, 200), (100, 200)),
-        "uv": ((0, 0), (1, 0), (1, 1), (0, 1)),
+        "texCoord": ((0, 0), (1, 0), (1, 1), (0, 1)),
     },
 )
 
+if image.gl_load():
+    raise Exception()
+
 
 def draw():
     bgl.glActiveTexture(bgl.GL_TEXTURE0)
-    bgl.glBindTexture(bgl.GL_TEXTURE_2D, offscreen.color_texture)
+    bgl.glBindTexture(bgl.GL_TEXTURE_2D, image.bindcode)
 
     shader.bind()
-    shader.uniform_float("modelMatrix", Matrix.Translation((1, 2, 3)) @ Matrix.Scale(3, 4))
+    shader.uniform_int("image", 0)
-    shader.uniform_float("viewProjectionMatrix", bpy.context.region_data.perspective_matrix)
-    shader.uniform_float("image", 0)
     batch.draw(shader)
 
 
-bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_VIEW')
+bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_PIXEL')

doc/python_api/examples/gpu.8.py

@@ -1,28 +1,23 @@
 """
-Copy Offscreen Rendering result back to RAM
+Generate a texture using Offscreen Rendering
--------------------------------------------
+--------------------------------------------
 
-This will create a new image with the given name.
+#. Create an :class:`gpu.types.GPUOffScreen` object.
-If it already exists, it will override the existing one.
+#. Draw some circles into it.
-
+#. Make a new shader for drawing a planar texture in 3D.
-Currently almost all of the execution time is spent in the last line.
+#. Draw the generated texture using the new shader.
-In the future this will hopefully be solved by implementing the Python buffer protocol
-for :class:`bgl.Buffer` and :class:`bpy.types.Image.pixels` (aka ``bpy_prop_array``).
 """
 import bpy
 import gpu
 import bgl
-import random
 from mathutils import Matrix
+from gpu_extras.batch import batch_for_shader
 from gpu_extras.presets import draw_circle_2d
 
-IMAGE_NAME = "Generated Image"
+# Create and fill offscreen
-WIDTH = 512
+##########################################
-HEIGHT = 512
-RING_AMOUNT = 10
 
-
+offscreen = gpu.types.GPUOffScreen(512, 512)
-offscreen = gpu.types.GPUOffScreen(WIDTH, HEIGHT)
 
 with offscreen.bind():
     bgl.glClear(bgl.GL_COLOR_BUFFER_BIT)
@@ -31,20 +26,61 @@ with offscreen.bind():
         gpu.matrix.load_matrix(Matrix.Identity(4))
         gpu.matrix.load_projection_matrix(Matrix.Identity(4))
 
-        for i in range(RING_AMOUNT):
+        amount = 10
-            draw_circle_2d(
+        for i in range(-amount, amount + 1):
-                (random.uniform(-1, 1), random.uniform(-1, 1)),
+            x_pos = i / amount
-                (1, 1, 1, 1), random.uniform(0.1, 1), 20)
+            draw_circle_2d((x_pos, 0.0), (1, 1, 1, 1), 0.5, 200)
-
-    buffer = bgl.Buffer(bgl.GL_BYTE, WIDTH * HEIGHT * 4)
-    bgl.glReadBuffer(bgl.GL_BACK)
-    bgl.glReadPixels(0, 0, WIDTH, HEIGHT, bgl.GL_RGBA, bgl.GL_UNSIGNED_BYTE, buffer)
-
-offscreen.free()
 
 
-if not IMAGE_NAME in bpy.data.images:
+# Drawing the generated texture in 3D space
-    bpy.data.images.new(IMAGE_NAME, WIDTH, HEIGHT)
+#############################################
-image = bpy.data.images[IMAGE_NAME]
+
-image.scale(WIDTH, HEIGHT)
+vertex_shader = '''
-image.pixels = [v / 255 for v in buffer]
+    uniform mat4 modelMatrix;
+    uniform mat4 viewProjectionMatrix;
+
+    in vec2 position;
+    in vec2 uv;
+
+    out vec2 uvInterp;
+
+    void main()
+    {
+        uvInterp = uv;
+        gl_Position = viewProjectionMatrix * modelMatrix * vec4(position, 0.0, 1.0);
+    }
+'''
+
+fragment_shader = '''
+    uniform sampler2D image;
+
+    in vec2 uvInterp;
+
+    void main()
+    {
+        gl_FragColor = texture(image, uvInterp);
+    }
+'''
+
+shader = gpu.types.GPUShader(vertex_shader, fragment_shader)
+batch = batch_for_shader(
+    shader, 'TRI_FAN',
+    {
+        "position": ((-1, -1), (1, -1), (1, 1), (-1, 1)),
+        "uv": ((0, 0), (1, 0), (1, 1), (0, 1)),
+    },
+)
+
+
+def draw():
+    bgl.glActiveTexture(bgl.GL_TEXTURE0)
+    bgl.glBindTexture(bgl.GL_TEXTURE_2D, offscreen.color_texture)
+
+    shader.bind()
+    shader.uniform_float("modelMatrix", Matrix.Translation((1, 2, 3)) @ Matrix.Scale(3, 4))
+    shader.uniform_float("viewProjectionMatrix", bpy.context.region_data.perspective_matrix)
+    shader.uniform_float("image", 0)
+    batch.draw(shader)
+
+
+bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_VIEW')

doc/python_api/examples/gpu.9.py

@@ -1,41 +1,50 @@
 """
-Rendering the 3D View into a Texture
+Copy Offscreen Rendering result back to RAM
-------------------------------------
+-------------------------------------------
 
-The scene has to have a camera for this example to work.
+This will create a new image with the given name.
-You could also make this independent of a specific camera,
+If it already exists, it will override the existing one.
-but Blender does not expose good functions to create view and projection matrices yet.
+
+Currently almost all of the execution time is spent in the last line.
+In the future this will hopefully be solved by implementing the Python buffer protocol
+for :class:`bgl.Buffer` and :class:`bpy.types.Image.pixels` (aka ``bpy_prop_array``).
 """
 import bpy
-import bgl
 import gpu
-from gpu_extras.presets import draw_texture_2d
+import bgl
+import random
+from mathutils import Matrix
+from gpu_extras.presets import draw_circle_2d
 
+IMAGE_NAME = "Generated Image"
 WIDTH = 512
-HEIGHT = 256
+HEIGHT = 512
+RING_AMOUNT = 10
+
 
 offscreen = gpu.types.GPUOffScreen(WIDTH, HEIGHT)
 
+with offscreen.bind():
+    bgl.glClear(bgl.GL_COLOR_BUFFER_BIT)
+    with gpu.matrix.push_pop():
+        # reset matrices -> use normalized device coordinates [-1, 1]
+        gpu.matrix.load_matrix(Matrix.Identity(4))
+        gpu.matrix.load_projection_matrix(Matrix.Identity(4))
 
-def draw():
+        for i in range(RING_AMOUNT):
-    context = bpy.context
+            draw_circle_2d(
-    scene = context.scene
+                (random.uniform(-1, 1), random.uniform(-1, 1)),
+                (1, 1, 1, 1), random.uniform(0.1, 1), 20)
 
-    view_matrix = scene.camera.matrix_world.inverted()
+    buffer = bgl.Buffer(bgl.GL_BYTE, WIDTH * HEIGHT * 4)
+    bgl.glReadBuffer(bgl.GL_BACK)
+    bgl.glReadPixels(0, 0, WIDTH, HEIGHT, bgl.GL_RGBA, bgl.GL_UNSIGNED_BYTE, buffer)
 
-    projection_matrix = scene.camera.calc_matrix_camera(
+offscreen.free()
-        context.evaluated_depsgraph_get(), x=WIDTH, y=HEIGHT)
-
-    offscreen.draw_view3d(
-        scene,
-        context.view_layer,
-        context.space_data,
-        context.region,
-        view_matrix,
-        projection_matrix)
-
-    bgl.glDisable(bgl.GL_DEPTH_TEST)
-    draw_texture_2d(offscreen.color_texture, (10, 10), WIDTH, HEIGHT)
 
 
-bpy.types.SpaceView3D.draw_handler_add(draw, (), 'WINDOW', 'POST_PIXEL')
+if not IMAGE_NAME in bpy.data.images:
+    bpy.data.images.new(IMAGE_NAME, WIDTH, HEIGHT)
+image = bpy.data.images[IMAGE_NAME]
+image.scale(WIDTH, HEIGHT)
+image.pixels = [v / 255 for v in buffer]

doc/python_api/requirements.txt

@@ -1,2 +1,2 @@
-Sphinx==3.0.3
+Sphinx==1.8.5
-sphinx_rtd_theme==0.5.0rc1
+sphinx_rtd_theme==0.4.3

doc/python_api/sphinx_doc_gen.py

@@ -403,21 +403,32 @@ MODULE_GROUPING = {
 
 # -------------------------------BLENDER----------------------------------------
 
+blender_version_strings = [str(v) for v in bpy.app.version]
+is_release = bpy.app.version_cycle in {"rc", "release"}
+
 # converting bytes to strings, due to T30154
 BLENDER_REVISION = str(bpy.app.build_hash, 'utf_8')
 
-# '2.83.0 Beta' or '2.83.0' or '2.83.1'
+if is_release:
-BLENDER_VERSION_DOTS = bpy.app.version_string
+    # '2.62a'
+    BLENDER_VERSION_DOTS = ".".join(blender_version_strings[:2]) + bpy.app.version_char
+else:
+    # '2.62.1'
+    BLENDER_VERSION_DOTS = ".".join(blender_version_strings)
 
 if BLENDER_REVISION != "Unknown":
-    # SHA1 Git hash
+    # '2.62a SHA1' (release) or '2.62.1 SHA1' (non-release)
     BLENDER_VERSION_HASH = BLENDER_REVISION
 else:
     # Fallback: Should not be used
     BLENDER_VERSION_HASH = "Hash Unknown"
 
-# '2_83'
+if is_release:
-BLENDER_VERSION_PATH = "%d_%d" % (bpy.app.version[0], bpy.app.version[1])
+    # '2_62a_release'
+    BLENDER_VERSION_PATH = "%s%s_release" % ("_".join(blender_version_strings[:2]), bpy.app.version_char)
+else:
+    # '2_62_1'
+    BLENDER_VERSION_PATH = "_".join(blender_version_strings)
 
 # --------------------------DOWNLOADABLE FILES----------------------------------
 
@@ -688,13 +699,11 @@ def py_descr2sphinx(ident, fw, descr, module_name, type_name, identifier):
         doc = undocumented_message(module_name, type_name, identifier)
 
     if type(descr) == GetSetDescriptorType:
-        fw(ident + ".. attribute:: %s\n" % identifier)
+        fw(ident + ".. attribute:: %s\n\n" % identifier)
-        fw(ident + "   :noindex:\n\n")
         write_indented_lines(ident + "   ", fw, doc, False)
         fw("\n")
     elif type(descr) == MemberDescriptorType:  # same as above but use 'data'
-        fw(ident + ".. data:: %s\n" % identifier)
+        fw(ident + ".. data:: %s\n\n" % identifier)
-        fw(ident + "   :noindex:\n\n")
         write_indented_lines(ident + "   ", fw, doc, False)
         fw("\n")
     elif type(descr) in {MethodDescriptorType, ClassMethodDescriptorType}:
@@ -734,14 +743,11 @@ def pyprop2sphinx(ident, fw, identifier, py_prop):
     '''
     # readonly properties use "data" directive, variables use "attribute" directive
     if py_prop.fset is None:
-        fw(ident + ".. data:: %s\n" % identifier)
+        fw(ident + ".. data:: %s\n\n" % identifier)
-        fw(ident + "   :noindex:\n\n")
     else:
-        fw(ident + ".. attribute:: %s\n" % identifier)
+        fw(ident + ".. attribute:: %s\n\n" % identifier)
-        fw(ident + "   :noindex:\n\n")
     write_indented_lines(ident + "   ", fw, py_prop.__doc__)
     if py_prop.fset is None:
-        fw("\n")
         fw(ident + "   (readonly)\n\n")
     else:
         fw("\n")
@@ -907,8 +913,7 @@ def pymodule2sphinx(basepath, module_name, module, title):
         elif issubclass(value_type, (bool, int, float, str, tuple)):
             # constant, not much fun we can do here except to list it.
             # TODO, figure out some way to document these!
-            fw(".. data:: %s\n" % attribute)
+            fw(".. data:: %s\n\n" % attribute)
-            fw("   :noindex:\n\n")
             write_indented_lines("   ", fw, "constant value %s" % repr(value), False)
             fw("\n")
         else:
@@ -1116,8 +1121,7 @@ def pycontext2sphinx(basepath):
 
             type_descr = prop.get_type_description(
                 class_fmt=":class:`bpy.types.%s`", collection_id=_BPY_PROP_COLLECTION_ID)
-            fw(".. data:: %s\n" % prop.identifier)
+            fw(".. data:: %s\n\n" % prop.identifier)
-            fw("   :noindex:\n\n")
             if prop.description:
                 fw("   %s\n\n" % prop.description)
 
@@ -1162,8 +1166,7 @@ def pycontext2sphinx(basepath):
         i = 0
         while char_array[i] is not None:
             member = ctypes.string_at(char_array[i]).decode(encoding="ascii")
-            fw(".. data:: %s\n" % member)
+            fw(".. data:: %s\n\n" % member)
-            fw("   :noindex:\n\n")
             member_type, is_seq = context_type_map[member]
             fw("   :type: %s :class:`bpy.types.%s`\n\n" % ("sequence of " if is_seq else "", member_type))
             unique.add(member)
@@ -1307,7 +1310,7 @@ def pyrna2sphinx(basepath):
 
         fw(title_string(title, "="))
 
-        fw(".. module:: %s.%s\n\n" % (struct_module_name, struct_id))
+        fw(".. module:: %s\n\n" % struct_module_name)
 
         # docs first?, ok
         write_example_ref("", fw, "%s.%s" % (struct_module_name, struct_id))
@@ -1366,11 +1369,9 @@ def pyrna2sphinx(basepath):
             type_descr = prop.get_type_description(class_fmt=":class:`%s`", collection_id=_BPY_PROP_COLLECTION_ID)
             # readonly properties use "data" directive, variables properties use "attribute" directive
             if 'readonly' in type_descr:
-                fw("   .. data:: %s\n" % prop.identifier)
+                fw("   .. data:: %s\n\n" % prop.identifier)
-                fw("      :noindex:\n\n")
             else:
-                fw("   .. attribute:: %s\n" % prop.identifier)
+                fw("   .. attribute:: %s\n\n" % prop.identifier)
-                fw("      :noindex:\n\n")
             if prop.description:
                 fw("      %s\n\n" % prop.description)
 
@@ -1540,7 +1541,7 @@ def pyrna2sphinx(basepath):
 
             fw(title_string(class_name, "="))
 
-            fw(".. module:: %s.%s\n" % (class_module_name, class_name))
+            fw(".. module:: %s\n" % class_module_name)
             fw("\n")
 
             if use_subclasses:

doc/python_api/sphinx_doc_gen.sh

@@ -36,10 +36,16 @@ fi
 blender_srcdir=$(dirname -- $0)/../..
 blender_version_header="$blender_srcdir/source/blender/blenkernel/BKE_blender_version.h"
 blender_version=$(grep "BLENDER_VERSION\s" "$blender_version_header" | awk '{print $3}')
+blender_version_char=$(grep "BLENDER_VERSION_CHAR\s" "$blender_version_header" | awk '{print $3}')
 blender_version_cycle=$(grep "BLENDER_VERSION_CYCLE\s" "$blender_version_header" | awk '{print $3}')
+blender_subversion=$(grep "BLENDER_SUBVERSION\s" "$blender_version_header" | awk '{print $3}')
 unset blender_version_header
 
-BLENDER_VERSION=$(expr $blender_version / 100)_$(expr $blender_version % 100)
+if [ "$blender_version_cycle" = "release" ] ; then
+  BLENDER_VERSION=$(expr $blender_version / 100)_$(expr $blender_version % 100)$blender_version_char"_release"
+else
+  BLENDER_VERSION=$(expr $blender_version / 100)_$(expr $blender_version % 100)_$blender_subversion
+fi
 
 SSH_UPLOAD_FULL=$SSH_UPLOAD/"blender_python_api_"$BLENDER_VERSION
 

doc/python_api/sphinx_doc_update.py

@@ -127,10 +127,11 @@ def main():
             "    f.write('%d\\n' % is_release)\n"
             "    f.write('%d\\n' % is_beta)\n"
             "    f.write('%s\\n' % branch)\n"
-            "    f.write('%d.%d\\n' % (bpy.app.version[0], bpy.app.version[1]))\n"
+            "    f.write('%d.%d%s\\n' % (bpy.app.version[0], bpy.app.version[1], bpy.app.version_char))\n"
-            "    f.write('%d.%d\\n' % (bpy.app.version[0], bpy.app.version[1])\n"
+            "    f.write('%d.%d%s\\n' % (bpy.app.version[0], bpy.app.version[1], bpy.app.version_char)\n"
             "            if (is_release or is_beta) else '%s\\n' % branch)\n"
-            "    f.write('%d_%d' % (bpy.app.version[0], bpy.app.version[1]))\n"
+            "    f.write('%d_%d%s_release' % (bpy.app.version[0], bpy.app.version[1], bpy.app.version_char)\n"
+            "            if is_release else '%d_%d_%d' % bpy.app.version)\n"
         )
         get_ver_cmd = (args.blender, "--background", "-noaudio", "--factory-startup", "--python-exit-code", "1",
                        "--python-expr", getver_script, "--", getver_file)

extern/mantaflow/preprocessed/gitinfo.h

@@ -1,3 +1,3 @@
 
 
-#define MANTA_GIT_VERSION "commit b61bf9efa7a1d8ca98635076a7e9f2c4dacb2914"
+#define MANTA_GIT_VERSION "commit b4a2742bd743e2913fba94dd35846042e2650212"

extern/mantaflow/preprocessed/particle.h

@@ -469,7 +469,6 @@ template<class S> class ParticleSystem : public ParticleBase {
                     const int integrationMode,
                     const bool deleteInObstacle = true,
                     const bool stopInObstacle = true,
-                    const bool skipNew = false,
                     const ParticleDataImpl<int> *ptype = NULL,
                     const int exclude = 0);
   static PyObject *_W_9(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
@@ -487,20 +486,13 @@ template<class S> class ParticleSystem : public ParticleBase {
         const int integrationMode = _args.get<int>("integrationMode", 2, &_lock);
         const bool deleteInObstacle = _args.getOpt<bool>("deleteInObstacle", 3, true, &_lock);
         const bool stopInObstacle = _args.getOpt<bool>("stopInObstacle", 4, true, &_lock);
-        const bool skipNew = _args.getOpt<bool>("skipNew", 5, false, &_lock);
         const ParticleDataImpl<int> *ptype = _args.getPtrOpt<ParticleDataImpl<int>>(
-            "ptype", 6, NULL, &_lock);
+            "ptype", 5, NULL, &_lock);
-        const int exclude = _args.getOpt<int>("exclude", 7, 0, &_lock);
+        const int exclude = _args.getOpt<int>("exclude", 6, 0, &_lock);
         pbo->_args.copy(_args);
         _retval = getPyNone();
-        pbo->advectInGrid(flags,
+        pbo->advectInGrid(
-                          vel,
+            flags, vel, integrationMode, deleteInObstacle, stopInObstacle, ptype, exclude);
-                          integrationMode,
-                          deleteInObstacle,
-                          stopInObstacle,
-                          skipNew,
-                          ptype,
-                          exclude);
         pbo->_args.check();
       }
       pbFinalizePlugin(pbo->getParent(), "ParticleSystem::advectInGrid", !noTiming);
@@ -1871,7 +1863,6 @@ template<class S> struct _GridAdvectKernel : public KernelBase {
                     const Real dt,
                     const bool deleteInObstacle,
                     const bool stopInObstacle,
-                    const bool skipNew,
                     const ParticleDataImpl<int> *ptype,
                     const int exclude,
                     std::vector<Vec3> &u)
@@ -1882,7 +1873,6 @@ template<class S> struct _GridAdvectKernel : public KernelBase {
         dt(dt),
         deleteInObstacle(deleteInObstacle),
         stopInObstacle(stopInObstacle),
-        skipNew(skipNew),
         ptype(ptype),
         exclude(exclude),
         u(u)
@@ -1895,13 +1885,11 @@ template<class S> struct _GridAdvectKernel : public KernelBase {
                  const Real dt,
                  const bool deleteInObstacle,
                  const bool stopInObstacle,
-                 const bool skipNew,
                  const ParticleDataImpl<int> *ptype,
                  const int exclude,
                  std::vector<Vec3> &u) const
   {
-    if ((p[idx].flag & ParticleBase::PDELETE) || (ptype && ((*ptype)[idx] & exclude)) ||
+    if ((p[idx].flag & ParticleBase::PDELETE) || (ptype && ((*ptype)[idx] & exclude))) {
-        (skipNew && (p[idx].flag & ParticleBase::PNEW))) {
       u[idx] = 0.;
       return;
     }
@@ -1922,7 +1910,7 @@ template<class S> struct _GridAdvectKernel : public KernelBase {
   void operator()(const tbb::blocked_range<IndexInt> &__r) const
   {
     for (IndexInt idx = __r.begin(); idx != (IndexInt)__r.end(); idx++)
-      op(idx, p, vel, flags, dt, deleteInObstacle, stopInObstacle, skipNew, ptype, exclude, u);
+      op(idx, p, vel, flags, dt, deleteInObstacle, stopInObstacle, ptype, exclude, u);
   }
   void run()
   {
@@ -1934,7 +1922,6 @@ template<class S> struct _GridAdvectKernel : public KernelBase {
   const Real dt;
   const bool deleteInObstacle;
   const bool stopInObstacle;
-  const bool skipNew;
   const ParticleDataImpl<int> *ptype;
   const int exclude;
   std::vector<Vec3> &u;
@@ -1946,7 +1933,6 @@ template<class S> struct GridAdvectKernel : public KernelBase {
                    const Real dt,
                    const bool deleteInObstacle,
                    const bool stopInObstacle,
-                   const bool skipNew,
                    const ParticleDataImpl<int> *ptype,
                    const int exclude)
       : KernelBase(p.size()),
@@ -1957,7 +1943,6 @@ template<class S> struct GridAdvectKernel : public KernelBase {
                dt,
                deleteInObstacle,
                stopInObstacle,
-               skipNew,
                ptype,
                exclude,
                u),
@@ -1967,7 +1952,6 @@ template<class S> struct GridAdvectKernel : public KernelBase {
         dt(dt),
         deleteInObstacle(deleteInObstacle),
         stopInObstacle(stopInObstacle),
-        skipNew(skipNew),
         ptype(ptype),
         exclude(exclude),
         u((size))
@@ -2017,21 +2001,16 @@ template<class S> struct GridAdvectKernel : public KernelBase {
     return stopInObstacle;
   }
   typedef bool type5;
-  inline const bool &getArg6()
+  inline const ParticleDataImpl<int> *getArg6()
-  {
-    return skipNew;
-  }
-  typedef bool type6;
-  inline const ParticleDataImpl<int> *getArg7()
   {
     return ptype;
   }
-  typedef ParticleDataImpl<int> type7;
+  typedef ParticleDataImpl<int> type6;
-  inline const int &getArg8()
+  inline const int &getArg7()
   {
     return exclude;
   }
-  typedef int type8;
+  typedef int type7;
   void runMessage()
   {
     debMsg("Executing kernel GridAdvectKernel ", 3);
@@ -2046,7 +2025,6 @@ template<class S> struct GridAdvectKernel : public KernelBase {
   const Real dt;
   const bool deleteInObstacle;
   const bool stopInObstacle;
-  const bool skipNew;
   const ParticleDataImpl<int> *ptype;
   const int exclude;
   std::vector<Vec3> u;
@@ -2217,7 +2195,6 @@ void ParticleSystem<S>::advectInGrid(const FlagGrid &flags,
                                      const int integrationMode,
                                      const bool deleteInObstacle,
                                      const bool stopInObstacle,
-                                     const bool skipNew,
                                      const ParticleDataImpl<int> *ptype,
                                      const int exclude)
 {
@@ -2231,15 +2208,8 @@ void ParticleSystem<S>::advectInGrid(const FlagGrid &flags,
   }
 
   // update positions
-  GridAdvectKernel<S> kernel(mData,
+  GridAdvectKernel<S> kernel(
-                             vel,
+      mData, vel, flags, getParent()->getDt(), deleteInObstacle, stopInObstacle, ptype, exclude);
-                             flags,
-                             getParent()->getDt(),
-                             deleteInObstacle,
-                             stopInObstacle,
-                             skipNew,
-                             ptype,
-                             exclude);
   integratePointSet(kernel, integrationMode);
 
   if (!deleteInObstacle) {
@@ -2468,15 +2438,15 @@ template<class S> void ParticleSystem<S>::compress()
 //! insert buffered positions as new particles, update additional particle data
 template<class S> void ParticleSystem<S>::insertBufferedParticles()
 {
-  // clear new flag everywhere
-  for (IndexInt i = 0; i < (IndexInt)mData.size(); ++i)
-    mData[i].flag &= ~PNEW;
-
   if (mNewBufferPos.size() == 0)
     return;
   IndexInt newCnt = mData.size();
   resizeAll(newCnt + mNewBufferPos.size());
 
+  // clear new flag everywhere
+  for (IndexInt i = 0; i < (IndexInt)mData.size(); ++i)
+    mData[i].flag &= ~PNEW;
+
   for (IndexInt i = 0; i < (IndexInt)mNewBufferPos.size(); ++i) {
     int flag = (mNewBufferFlag.size() > 0) ? mNewBufferFlag[i] : 0;
     // note, other fields are not initialized here...

intern/cycles/app/cycles_standalone.cpp

@@ -26,7 +26,6 @@
 #include "util/util_args.h"
 #include "util/util_foreach.h"
 #include "util/util_function.h"
-#include "util/util_image.h"
 #include "util/util_logging.h"
 #include "util/util_path.h"
 #include "util/util_progress.h"

intern/cycles/blender/addon/__init__.py

@@ -82,8 +82,8 @@ class CyclesRender(bpy.types.RenderEngine):
     def render(self, depsgraph):
         engine.render(self, depsgraph)
 
-    def bake(self, depsgraph, obj, pass_type, pass_filter, width, height):
+    def bake(self, depsgraph, obj, pass_type, pass_filter, object_id, pixel_array, num_pixels, depth, result):
-        engine.bake(self, depsgraph, obj, pass_type, pass_filter, width, height)
+        engine.bake(self, depsgraph, obj, pass_type, pass_filter, object_id, pixel_array, num_pixels, depth, result)
 
     # viewport render
     def view_update(self, context, depsgraph):

intern/cycles/blender/addon/engine.py

@@ -168,11 +168,11 @@ def render(engine, depsgraph):
         _cycles.render(engine.session, depsgraph.as_pointer())
 
 
-def bake(engine, depsgraph, obj, pass_type, pass_filter, width, height):
+def bake(engine, depsgraph, obj, pass_type, pass_filter, object_id, pixel_array, num_pixels, depth, result):
     import _cycles
     session = getattr(engine, "session", None)
     if session is not None:
-        _cycles.bake(engine.session, depsgraph.as_pointer(), obj.as_pointer(), pass_type, pass_filter, width, height)
+        _cycles.bake(engine.session, depsgraph.as_pointer(), obj.as_pointer(), pass_type, pass_filter, object_id, pixel_array.as_pointer(), num_pixels, depth, result.as_pointer())
 
 
 def reset(engine, data, depsgraph):

intern/cycles/blender/addon/osl.py

@@ -85,7 +85,6 @@ def update_script_node(node, report):
             # write text datablock contents to temporary file
             osl_file = tempfile.NamedTemporaryFile(mode='w', suffix=".osl", delete=False)
             osl_file.write(script.as_string())
-            osl_file.write("\n")
             osl_file.close()
 
             ok, oso_path = osl_compile(osl_file.name, report)

intern/cycles/blender/addon/properties.py

@@ -1205,13 +1205,6 @@ class CyclesObjectSettings(bpy.types.PropertyGroup):
         default=1.0,
     )
 
-    shadow_terminator_offset: FloatProperty(
-        name="Shadow Terminator Offset",
-        description="Push the shadow terminator towards the light to hide artifacts on low poly geometry",
-        min=0.0, max=1.0,
-        default=0.0,
-    )
-
     is_shadow_catcher: BoolProperty(
         name="Shadow Catcher",
         description="Only render shadows on this object, for compositing renders into real footage",
@@ -1535,12 +1528,6 @@ class CyclesPreferences(bpy.types.AddonPreferences):
 
     devices: bpy.props.CollectionProperty(type=CyclesDeviceSettings)
 
-    peer_memory: BoolProperty(
-        name="Distribute memory across devices",
-        description="Make more room for large scenes to fit by distributing memory across interconnected devices (e.g. via NVLink) rather than duplicating it",
-        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]:
@@ -1638,21 +1625,14 @@ class CyclesPreferences(bpy.types.AddonPreferences):
         row = layout.row()
         row.prop(self, "compute_device_type", expand=True)
 
-        if self.compute_device_type == 'NONE':
-            return
-        row = layout.row()
         devices = self.get_devices_for_type(self.compute_device_type)
-        self._draw_devices(row, self.compute_device_type, devices)
+        row = layout.row()
-
+        if self.compute_device_type == 'CUDA':
-        import _cycles
+            self._draw_devices(row, 'CUDA', devices)
-        has_peer_memory = 0
+        elif self.compute_device_type == 'OPTIX':
-        for device in _cycles.available_devices(self.compute_device_type):
+            self._draw_devices(row, 'OPTIX', devices)
-            if device[3] and self.find_existing_device_entry(device).use:
+        elif self.compute_device_type == 'OPENCL':
-                has_peer_memory += 1
+            self._draw_devices(row, 'OPENCL', devices)
-        if has_peer_memory > 1:
-            row = layout.row()
-            row.use_property_split = True
-            row.prop(self, "peer_memory")
 
     def draw(self, context):
         self.draw_impl(self.layout, context)

intern/cycles/blender/addon/ui.py

@@ -1209,27 +1209,6 @@ def has_geometry_visibility(ob):
     return ob and ((ob.type in {'MESH', 'CURVE', 'SURFACE', 'FONT', 'META', 'LIGHT'}) or
                     (ob.instance_type == 'COLLECTION' and ob.instance_collection))
 
-class CYCLES_OBJECT_PT_shading(CyclesButtonsPanel, Panel):
-    bl_label = "Shading"
-    bl_context = "object"
-    bl_options = {'DEFAULT_CLOSED'}
-
-    @classmethod
-    def poll(cls, context):
-        return  CyclesButtonsPanel.poll(context) and (context.object)
-
-    def draw(self, context):
-        layout = self.layout
-        layout.use_property_split = True
-
-        flow = layout.grid_flow(row_major=False, columns=0, even_columns=True, even_rows=False, align=False)
-        layout = self.layout
-        ob = context.object
-        cob = ob.cycles
-
-        if has_geometry_visibility(ob):
-            col = flow.column()
-            col.prop(cob, "shadow_terminator_offset")
 
 class CYCLES_OBJECT_PT_visibility(CyclesButtonsPanel, Panel):
     bl_label = "Visibility"
@@ -1388,7 +1367,7 @@ class CYCLES_LIGHT_PT_light(CyclesButtonsPanel, Panel):
         col.separator()
 
         if light.type in {'POINT', 'SPOT'}:
-            col.prop(light, "shadow_soft_size", text="Radius")
+            col.prop(light, "shadow_soft_size", text="Size")
         elif light.type == 'SUN':
             col.prop(light, "angle")
         elif light.type == 'AREA':
@@ -1928,15 +1907,10 @@ class CYCLES_RENDER_PT_bake_selected_to_active(CyclesButtonsPanel, Panel):
 
         col.prop(cbk, "use_cage", text="Cage")
         if cbk.use_cage:
-            col.prop(cbk, "cage_object")
-            col = layout.column()
-            col.prop(cbk, "cage_extrusion")
-            col.active = cbk.cage_object is None
-        else:
             col.prop(cbk, "cage_extrusion", text="Extrusion")
-
+            col.prop(cbk, "cage_object", text="Cage Object")
-        col = layout.column()
+        else:
-        col.prop(cbk, "max_ray_distance")
+            col.prop(cbk, "cage_extrusion", text="Ray Distance")
 
 
 class CYCLES_RENDER_PT_bake_output(CyclesButtonsPanel, Panel):
@@ -2289,7 +2263,6 @@ classes = (
     CYCLES_CAMERA_PT_dof_aperture,
     CYCLES_PT_context_material,
     CYCLES_OBJECT_PT_motion_blur,
-    CYCLES_OBJECT_PT_shading,
     CYCLES_OBJECT_PT_visibility,
     CYCLES_OBJECT_PT_visibility_ray_visibility,
     CYCLES_OBJECT_PT_visibility_culling,

intern/cycles/blender/blender_curves.cpp

@@ -897,7 +897,7 @@ static void ExportCurveTriangleVcol(ParticleCurveData *CData, int resol, uchar4
 
 /* Hair Curve Sync */
 
-void BlenderSync::sync_curve_settings(BL::Depsgraph &b_depsgraph)
+void BlenderSync::sync_curve_settings()
 {
   PointerRNA csscene = RNA_pointer_get(&b_scene.ptr, "cycles_curves");
 
@@ -949,9 +949,9 @@ void BlenderSync::sync_curve_settings(BL::Depsgraph &b_depsgraph)
   }
 
   if (curve_system_manager->modified_mesh(prev_curve_system_manager)) {
-    BL::Depsgraph::objects_iterator b_ob;
+    BL::BlendData::objects_iterator b_ob;
 
-    for (b_depsgraph.objects.begin(b_ob); b_ob != b_data.objects.end(); ++b_ob) {
+    for (b_data.objects.begin(b_ob); b_ob != b_data.objects.end(); ++b_ob) {
       if (object_is_mesh(*b_ob)) {
         BL::Object::particle_systems_iterator b_psys;
         for (b_ob->particle_systems.begin(b_psys); b_psys != b_ob->particle_systems.end();

intern/cycles/blender/blender_device.cpp

@@ -113,10 +113,6 @@ DeviceInfo blender_device_info(BL::Preferences &b_preferences, BL::Scene &b_scen
         device = Device::get_multi_device(used_devices, threads, background);
       }
       /* Else keep using the CPU device that was set before. */
-
-      if (!get_boolean(cpreferences, "peer_memory")) {
-        device.has_peer_memory = false;
-      }
     }
   }
 
@@ -141,25 +137,10 @@ DeviceInfo blender_device_info(BL::Preferences &b_preferences, BL::Scene &b_scen
         device.multi_devices.push_back(device);
       }
 
-      /* Try to use the same physical devices for denoising. */
+      /* Simply use the first available OptiX device. */
-      for (const DeviceInfo &cuda_device : device.multi_devices) {
+      const DeviceInfo optix_device = optix_devices.front();
-        if (cuda_device.type == DEVICE_CUDA) {
+      device.id += optix_device.id; /* Uniquely identify this special multi device. */
-          for (const DeviceInfo &optix_device : optix_devices) {
+      device.denoising_devices.push_back(optix_device);
-            if (cuda_device.num == optix_device.num) {
-              device.id += optix_device.id;
-              device.denoising_devices.push_back(optix_device);
-              break;
-            }
-          }
-        }
-      }
-
-      if (device.denoising_devices.empty()) {
-        /* Simply use the first available OptiX device. */
-        const DeviceInfo optix_device = optix_devices.front();
-        device.id += optix_device.id; /* Uniquely identify this special multi device. */
-        device.denoising_devices.push_back(optix_device);
-      }
     }
   }
 

intern/cycles/blender/blender_object.cpp

@@ -238,12 +238,6 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph,
     object_updated = true;
   }
 
-  float shadow_terminator_offset = get_float(cobject, "shadow_terminator_offset");
-  if (shadow_terminator_offset != object->shadow_terminator_offset) {
-    object->shadow_terminator_offset = shadow_terminator_offset;
-    object_updated = true;
-  }
-
   /* sync the asset name for Cryptomatte */
   BL::Object parent = b_ob.parent();
   ustring parent_name;
@@ -466,12 +460,9 @@ void BlenderSync::sync_motion(BL::RenderSettings &b_render,
     sync_objects(b_depsgraph, b_v3d, 0.0f);
   }
 
-  /* Insert motion times from camera. Motion times from other objects
+  /* always sample these times for camera motion */
-   * have already been added in a sync_objects call. */
+  motion_times.insert(-1.0f);
-  uint camera_motion_steps = object_motion_steps(b_cam, b_cam);
+  motion_times.insert(1.0f);
-  for (size_t step = 0; step < camera_motion_steps; step++) {
-    motion_times.insert(scene->camera->motion_time(step));
-  }
 
   /* note iteration over motion_times set happens in sorted order */
   foreach (float relative_time, motion_times) {
@@ -496,8 +487,10 @@ void BlenderSync::sync_motion(BL::RenderSettings &b_render,
     b_engine.frame_set(frame, subframe);
     python_thread_state_save(python_thread_state);
 
-    /* Syncs camera motion if relative_time is one of the camera's motion times. */
+    /* sync camera, only supports two times at the moment */
-    sync_camera_motion(b_render, b_cam, width, height, relative_time);
+    if (relative_time == -1.0f || relative_time == 1.0f) {
+      sync_camera_motion(b_render, b_cam, width, height, relative_time);
+    }
 
     /* sync object */
     sync_objects(b_depsgraph, b_v3d, relative_time);

intern/cycles/blender/blender_python.cpp

@@ -298,18 +298,22 @@ static PyObject *render_func(PyObject * /*self*/, PyObject *args)
 static PyObject *bake_func(PyObject * /*self*/, PyObject *args)
 {
   PyObject *pysession, *pydepsgraph, *pyobject;
+  PyObject *pypixel_array, *pyresult;
   const char *pass_type;
-  int pass_filter, width, height;
+  int num_pixels, depth, object_id, pass_filter;
 
   if (!PyArg_ParseTuple(args,
-                        "OOOsiii",
+                        "OOOsiiOiiO",
                         &pysession,
                         &pydepsgraph,
                         &pyobject,
                         &pass_type,
                         &pass_filter,
-                        &width,
+                        &object_id,
-                        &height))
+                        &pypixel_array,
+                        &num_pixels,
+                        &depth,
+                        &pyresult))
     return NULL;
 
   BlenderSession *session = (BlenderSession *)PyLong_AsVoidPtr(pysession);
@@ -322,9 +326,23 @@ static PyObject *bake_func(PyObject * /*self*/, PyObject *args)
   RNA_id_pointer_create((ID *)PyLong_AsVoidPtr(pyobject), &objectptr);
   BL::Object b_object(objectptr);
 
+  void *b_result = PyLong_AsVoidPtr(pyresult);
+
+  PointerRNA bakepixelptr;
+  RNA_pointer_create(NULL, &RNA_BakePixel, PyLong_AsVoidPtr(pypixel_array), &bakepixelptr);
+  BL::BakePixel b_bake_pixel(bakepixelptr);
+
   python_thread_state_save(&session->python_thread_state);
 
-  session->bake(b_depsgraph, b_object, pass_type, pass_filter, width, height);
+  session->bake(b_depsgraph,
+                b_object,
+                pass_type,
+                pass_filter,
+                object_id,
+                b_bake_pixel,
+                (size_t)num_pixels,
+                depth,
+                (float *)b_result);
 
   python_thread_state_restore(&session->python_thread_state);
 
@@ -416,11 +434,10 @@ static PyObject *available_devices_func(PyObject * /*self*/, PyObject *args)
   for (size_t i = 0; i < devices.size(); i++) {
     DeviceInfo &device = devices[i];
     string type_name = Device::string_from_type(device.type);
-    PyObject *device_tuple = PyTuple_New(4);
+    PyObject *device_tuple = PyTuple_New(3);
     PyTuple_SET_ITEM(device_tuple, 0, pyunicode_from_string(device.description.c_str()));
     PyTuple_SET_ITEM(device_tuple, 1, pyunicode_from_string(type_name.c_str()));
     PyTuple_SET_ITEM(device_tuple, 2, pyunicode_from_string(device.id.c_str()));
-    PyTuple_SET_ITEM(device_tuple, 3, PyBool_FromLong(device.has_peer_memory));
     PyTuple_SET_ITEM(ret, i, device_tuple);
   }
 

intern/cycles/blender/blender_session.cpp

@@ -168,13 +168,9 @@ void BlenderSession::create_session()
 
 void BlenderSession::reset_session(BL::BlendData &b_data, BL::Depsgraph &b_depsgraph)
 {
-  /* Update data, scene and depsgraph pointers. These can change after undo. */
   this->b_data = b_data;
   this->b_depsgraph = b_depsgraph;
   this->b_scene = b_depsgraph.scene_eval();
-  if (sync) {
-    sync->reset(this->b_data, this->b_scene);
-  }
 
   if (preview_osl) {
     PointerRNA cscene = RNA_pointer_get(&b_scene.ptr, "cycles");
@@ -251,7 +247,9 @@ void BlenderSession::reset_session(BL::BlendData &b_data, BL::Depsgraph &b_depsg
 
 void BlenderSession::free_session()
 {
-  delete sync;
+  if (sync)
+    delete sync;
+
   delete session;
 }
 
@@ -319,7 +317,6 @@ static void end_render_result(BL::RenderEngine &b_engine,
 
 void BlenderSession::do_write_update_render_tile(RenderTile &rtile,
                                                  bool do_update_only,
-                                                 bool do_read_only,
                                                  bool highlight)
 {
   int x = rtile.x - session->tile_manager.params.full_x;
@@ -345,23 +342,7 @@ void BlenderSession::do_write_update_render_tile(RenderTile &rtile,
 
   BL::RenderLayer b_rlay = *b_single_rlay;
 
-  if (do_read_only) {
+  if (do_update_only) {
-    /* copy each pass */
-    BL::RenderLayer::passes_iterator b_iter;
-
-    for (b_rlay.passes.begin(b_iter); b_iter != b_rlay.passes.end(); ++b_iter) {
-      BL::RenderPass b_pass(*b_iter);
-
-      /* find matching pass type */
-      PassType pass_type = BlenderSync::get_pass_type(b_pass);
-      int components = b_pass.channels();
-
-      rtile.buffers->set_pass_rect(pass_type, components, (float *)b_pass.rect());
-    }
-
-    end_render_result(b_engine, b_rr, false, false, false);
-  }
-  else if (do_update_only) {
     /* Sample would be zero at initial tile update, which is only needed
      * to tag tile form blender side as IN PROGRESS for proper highlight
      * no buffers should be sent to blender yet. For denoise we also
@@ -381,14 +362,9 @@ void BlenderSession::do_write_update_render_tile(RenderTile &rtile,
   }
 }
 
-void BlenderSession::read_render_tile(RenderTile &rtile)
-{
-  do_write_update_render_tile(rtile, false, true, false);
-}
-
 void BlenderSession::write_render_tile(RenderTile &rtile)
 {
-  do_write_update_render_tile(rtile, false, false, false);
+  do_write_update_render_tile(rtile, false, false);
 }
 
 void BlenderSession::update_render_tile(RenderTile &rtile, bool highlight)
@@ -398,9 +374,9 @@ void BlenderSession::update_render_tile(RenderTile &rtile, bool highlight)
    * would need to be investigated a bit further, but for now shall be fine
    */
   if (!b_engine.is_preview())
-    do_write_update_render_tile(rtile, true, false, highlight);
+    do_write_update_render_tile(rtile, true, highlight);
   else
-    do_write_update_render_tile(rtile, false, false, false);
+    do_write_update_render_tile(rtile, false, false);
 }
 
 static void add_cryptomatte_layer(BL::RenderResult &b_rr, string name, string manifest)
@@ -617,6 +593,25 @@ void BlenderSession::render(BL::Depsgraph &b_depsgraph_)
 #endif
 }
 
+static void populate_bake_data(BakeData *data,
+                               const int object_id,
+                               BL::BakePixel &pixel_array,
+                               const int num_pixels)
+{
+  BL::BakePixel bp = pixel_array;
+
+  int i;
+  for (i = 0; i < num_pixels; i++) {
+    if (bp.object_id() == object_id) {
+      data->set(i, bp.primitive_id(), bp.uv(), bp.du_dx(), bp.du_dy(), bp.dv_dx(), bp.dv_dy());
+    }
+    else {
+      data->set_null(i);
+    }
+    bp = bp.next();
+  }
+}
+
 static int bake_pass_filter_get(const int pass_filter)
 {
   int flag = BAKE_FILTER_NONE;
@@ -647,26 +642,43 @@ void BlenderSession::bake(BL::Depsgraph &b_depsgraph_,
                           BL::Object &b_object,
                           const string &pass_type,
                           const int pass_filter,
-                          const int bake_width,
+                          const int object_id,
-                          const int bake_height)
+                          BL::BakePixel &pixel_array,
+                          const size_t num_pixels,
+                          const int /*depth*/,
+                          float result[])
 {
   b_depsgraph = b_depsgraph_;
 
   ShaderEvalType shader_type = get_shader_type(pass_type);
+
+  /* Set baking flag in advance, so kernel loading can check if we need
+   * any baking capabilities.
+   */
+  scene->bake_manager->set_baking(true);
+
+  /* ensure kernels are loaded before we do any scene updates */
+  session->load_kernels();
+
+  if (shader_type == SHADER_EVAL_UV) {
+    /* force UV to be available */
+    Pass::add(PASS_UV, scene->film->passes);
+  }
+
   int bake_pass_filter = bake_pass_filter_get(pass_filter);
+  bake_pass_filter = BakeManager::shader_type_to_pass_filter(shader_type, bake_pass_filter);
 
-  /* Initialize bake manager, before we load the baking kernels. */
+  /* force use_light_pass to be true if we bake more than just colors */
-  scene->bake_manager->set(scene, b_object.name(), shader_type, bake_pass_filter);
+  if (bake_pass_filter & ~BAKE_FILTER_COLOR) {
+    Pass::add(PASS_LIGHT, scene->film->passes);
+  }
 
-  /* Passes are identified by name, so in order to return the combined pass we need to set the
+  /* create device and update scene */
-   * name. */
+  scene->film->tag_update(scene);
-  Pass::add(PASS_COMBINED, scene->film->passes, "Combined");
+  scene->integrator->tag_update(scene);
-
-  session->read_bake_tile_cb = function_bind(&BlenderSession::read_render_tile, this, _1);
-  session->write_render_tile_cb = function_bind(&BlenderSession::write_render_tile, this, _1);
 
   if (!session->progress.get_cancel()) {
-    /* Sync scene. */
+    /* update scene */
     BL::Object b_camera_override(b_engine.camera_override());
     sync->sync_camera(b_render, b_camera_override, width, height, "");
     sync->sync_data(
@@ -674,43 +686,75 @@ void BlenderSession::bake(BL::Depsgraph &b_depsgraph_,
     builtin_images_load();
   }
 
-  /* Object might have been disabled for rendering or excluded in some
+  BakeData *bake_data = NULL;
-   * other way, in that case Blender will report a warning afterwards. */
-  bool object_found = false;
-  foreach (Object *ob, scene->objects) {
-    if (ob->name == b_object.name()) {
-      object_found = true;
-      break;
-    }
-  }
 
-  if (object_found && !session->progress.get_cancel()) {
+  if (!session->progress.get_cancel()) {
-    /* Get session and buffer parameters. */
+    /* get buffer parameters */
     SessionParams session_params = BlenderSync::get_session_params(
         b_engine, b_userpref, b_scene, background);
-    session_params.progressive_refine = false;
+    BufferParams buffer_params = BlenderSync::get_buffer_params(
+        b_scene, b_render, b_v3d, b_rv3d, scene->camera, width, height);
 
-    BufferParams buffer_params;
+    scene->bake_manager->set_shader_limit((size_t)b_engine.tile_x(), (size_t)b_engine.tile_y());
-    buffer_params.width = bake_width;
-    buffer_params.height = bake_height;
-    buffer_params.passes = scene->film->passes;
 
-    /* Update session. */
+    /* set number of samples */
     session->tile_manager.set_samples(session_params.samples);
     session->reset(buffer_params, session_params.samples);
+    session->update_scene();
+
+    /* find object index. todo: is arbitrary - copied from mesh_displace.cpp */
+    size_t object_index = OBJECT_NONE;
+    int tri_offset = 0;
+
+    for (size_t i = 0; i < scene->objects.size(); i++) {
+      const Object *object = scene->objects[i];
+      const Geometry *geom = object->geometry;
+      if (object->name == b_object.name() && geom->type == Geometry::MESH) {
+        const Mesh *mesh = static_cast<const Mesh *>(geom);
+        object_index = i;
+        tri_offset = mesh->prim_offset;
+        break;
+      }
+    }
+
+    /* Object might have been disabled for rendering or excluded in some
+     * other way, in that case Blender will report a warning afterwards. */
+    if (object_index != OBJECT_NONE) {
+      int object = object_index;
+
+      bake_data = scene->bake_manager->init(object, tri_offset, num_pixels);
+      populate_bake_data(bake_data, object_id, pixel_array, num_pixels);
+    }
+
+    /* set number of samples */
+    session->tile_manager.set_samples(session_params.samples);
+    session->reset(buffer_params, session_params.samples);
+    session->update_scene();
 
     session->progress.set_update_callback(
         function_bind(&BlenderSession::update_bake_progress, this));
   }
 
   /* Perform bake. Check cancel to avoid crash with incomplete scene data. */
-  if (object_found && !session->progress.get_cancel()) {
+  if (!session->progress.get_cancel() && bake_data) {
-    session->start();
+    scene->bake_manager->bake(scene->device,
-    session->wait();
+                              &scene->dscene,
+                              scene,
+                              session->progress,
+                              shader_type,
+                              bake_pass_filter,
+                              bake_data,
+                              result);
   }
 
-  session->read_bake_tile_cb = function_null;
+  /* free all memory used (host and device), so we wouldn't leave render
-  session->write_render_tile_cb = function_null;
+   * engine with extra memory allocated
+   */
+
+  session->device_free();
+
+  delete sync;
+  sync = NULL;
 }
 
 void BlenderSession::do_write_update_render_result(BL::RenderLayer &b_rlay,

intern/cycles/blender/blender_session.h

@@ -66,12 +66,14 @@ class BlenderSession {
             BL::Object &b_object,
             const string &pass_type,
             const int custom_flag,
-            const int bake_width,
+            const int object_id,
-            const int bake_height);
+            BL::BakePixel &pixel_array,
+            const size_t num_pixels,
+            const int depth,
+            float pixels[]);
 
   void write_render_result(BL::RenderLayer &b_rlay, RenderTile &rtile);
   void write_render_tile(RenderTile &rtile);
-  void read_render_tile(RenderTile &rtile);
 
   /* update functions are used to update display buffer only after sample was rendered
    * only needed for better visual feedback */
@@ -153,10 +155,7 @@ class BlenderSession {
   void do_write_update_render_result(BL::RenderLayer &b_rlay,
                                      RenderTile &rtile,
                                      bool do_update_only);
-  void do_write_update_render_tile(RenderTile &rtile,
+  void do_write_update_render_tile(RenderTile &rtile, bool do_update_only, bool highlight);
-                                   bool do_update_only,
-                                   bool do_read_only,
-                                   bool highlight);
 
   void builtin_images_load();
 

intern/cycles/blender/blender_shader.cpp

@@ -384,16 +384,16 @@ static ShaderNode *add_node(Scene *scene,
 
     switch (b_aniso_node.distribution()) {
       case BL::ShaderNodeBsdfAnisotropic::distribution_BECKMANN:
-        aniso->distribution = CLOSURE_BSDF_MICROFACET_BECKMANN_ID;
+        aniso->distribution = CLOSURE_BSDF_MICROFACET_BECKMANN_ANISO_ID;
         break;
       case BL::ShaderNodeBsdfAnisotropic::distribution_GGX:
-        aniso->distribution = CLOSURE_BSDF_MICROFACET_GGX_ID;
+        aniso->distribution = CLOSURE_BSDF_MICROFACET_GGX_ANISO_ID;
         break;
       case BL::ShaderNodeBsdfAnisotropic::distribution_MULTI_GGX:
-        aniso->distribution = CLOSURE_BSDF_MICROFACET_MULTI_GGX_ID;
+        aniso->distribution = CLOSURE_BSDF_MICROFACET_MULTI_GGX_ANISO_ID;
         break;
       case BL::ShaderNodeBsdfAnisotropic::distribution_ASHIKHMIN_SHIRLEY:
-        aniso->distribution = CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID;
+        aniso->distribution = CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ANISO_ID;
         break;
     }
 
@@ -1231,11 +1231,12 @@ void BlenderSync::sync_materials(BL::Depsgraph &b_depsgraph, bool update_all)
     Shader *shader;
 
     /* test if we need to sync */
-    if (shader_map.add_or_update(&shader, b_mat) || update_all) {
+    if (shader_map.add_or_update(&shader, b_mat) || shader->need_sync_object || update_all) {
       ShaderGraph *graph = new ShaderGraph();
 
       shader->name = b_mat.name().c_str();
       shader->pass_id = b_mat.pass_index();
+      shader->need_sync_object = false;
 
       /* create nodes */
       if (b_mat.use_nodes() && b_mat.node_tree()) {

intern/cycles/blender/blender_sync.cpp

@@ -78,14 +78,6 @@ BlenderSync::~BlenderSync()
 {
 }
 
-void BlenderSync::reset(BL::BlendData &b_data, BL::Scene &b_scene)
-{
-  /* Update data and scene pointers in case they change in session reset,
-   * for example after undo. */
-  this->b_data = b_data;
-  this->b_scene = b_scene;
-}
-
 /* Sync */
 
 void BlenderSync::sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d)
@@ -93,6 +85,8 @@ void BlenderSync::sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d
   /* Sync recalc flags from blender to cycles. Actual update is done separate,
    * so we can do it later on if doing it immediate is not suitable. */
 
+  bool has_updated_objects = b_depsgraph.id_type_updated(BL::DriverTarget::id_type_OBJECT);
+
   if (experimental) {
     /* Mark all meshes as needing to be exported again if dicing changed. */
     PointerRNA cscene = RNA_pointer_get(&b_scene.ptr, "cycles");
@@ -195,6 +189,19 @@ void BlenderSync::sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d
   if (viewport_parameters.modified(new_viewport_parameters)) {
     world_recalc = true;
   }
+
+  /* Updates shader with object dependency if objects changed. */
+  if (has_updated_objects) {
+    if (scene->default_background->has_object_dependency) {
+      world_recalc = true;
+    }
+
+    foreach (Shader *shader, scene->shaders) {
+      if (shader->has_object_dependency) {
+        shader->need_sync_object = true;
+      }
+    }
+  }
 }
 
 void BlenderSync::sync_data(BL::RenderSettings &b_render,
@@ -212,7 +219,7 @@ void BlenderSync::sync_data(BL::RenderSettings &b_render,
   sync_film(b_v3d);
   sync_shaders(b_depsgraph, b_v3d);
   sync_images();
-  sync_curve_settings(b_depsgraph);
+  sync_curve_settings();
 
   geometry_synced.clear(); /* use for objects and motion sync */
 
@@ -489,9 +496,6 @@ PassType BlenderSync::get_pass_type(BL::RenderPass &b_pass)
   MAP_PASS("AO", PASS_AO);
   MAP_PASS("Shadow", PASS_SHADOW);
 
-  MAP_PASS("BakePrimitive", PASS_BAKE_PRIMITIVE);
-  MAP_PASS("BakeDifferential", PASS_BAKE_DIFFERENTIAL);
-
 #ifdef __KERNEL_DEBUG__
   MAP_PASS("Debug BVH Traversed Nodes", PASS_BVH_TRAVERSED_NODES);
   MAP_PASS("Debug BVH Traversed Instances", PASS_BVH_TRAVERSED_INSTANCES);

intern/cycles/blender/blender_sync.h

@@ -61,8 +61,6 @@ class BlenderSync {
               Progress &progress);
   ~BlenderSync();
 
-  void reset(BL::BlendData &b_data, BL::Scene &b_scene);
-
   /* sync */
   void sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d);
   void sync_data(BL::RenderSettings &b_render,
@@ -162,7 +160,7 @@ class BlenderSync {
   void sync_hair(Hair *hair, BL::Object &b_ob, bool motion, int motion_step = 0);
   void sync_particle_hair(
       Geometry *geom, BL::Mesh &b_mesh, BL::Object &b_ob, bool motion, int motion_step = 0);
-  void sync_curve_settings(BL::Depsgraph &b_depsgraph);
+  void sync_curve_settings();
   bool object_has_particle_hair(BL::Object b_ob);
 
   /* Camera */

intern/cycles/blender/blender_volume.cpp

@@ -34,13 +34,15 @@ CCL_NAMESPACE_BEGIN
 /* TODO: verify this is not loading unnecessary attributes. */
 class BlenderSmokeLoader : public ImageLoader {
  public:
-  BlenderSmokeLoader(BL::Object &b_ob, AttributeStandard attribute)
+  BlenderSmokeLoader(const BL::Object &b_ob, AttributeStandard attribute)
-      : b_domain(object_fluid_gas_domain_find(b_ob)), b_mesh(b_ob.data()), attribute(attribute)
+      : b_ob(b_ob), attribute(attribute)
   {
   }
 
   bool load_metadata(ImageMetaData &metadata) override
   {
+    BL::FluidDomainSettings b_domain = object_fluid_gas_domain_find(b_ob);
+
     if (!b_domain) {
       return false;
     }
@@ -77,6 +79,7 @@ class BlenderSmokeLoader : public ImageLoader {
     /* Create a matrix to transform from object space to mesh texture space.
      * This does not work with deformations but that can probably only be done
      * well with a volume grid mapping of coordinates. */
+    BL::Mesh b_mesh(b_ob.data());
     float3 loc, size;
     mesh_texture_space(b_mesh, loc, size);
     metadata.transform_3d = transform_translate(-loc) * transform_scale(size);
@@ -87,6 +90,9 @@ class BlenderSmokeLoader : public ImageLoader {
 
   bool load_pixels(const ImageMetaData &, void *pixels, const size_t, const bool) override
   {
+    /* smoke volume data */
+    BL::FluidDomainSettings b_domain = object_fluid_gas_domain_find(b_ob);
+
     if (!b_domain) {
       return false;
     }
@@ -173,11 +179,10 @@ class BlenderSmokeLoader : public ImageLoader {
   bool equals(const ImageLoader &other) const override
   {
     const BlenderSmokeLoader &other_loader = (const BlenderSmokeLoader &)other;
-    return b_domain == other_loader.b_domain && attribute == other_loader.attribute;
+    return b_ob == other_loader.b_ob && attribute == other_loader.attribute;
   }
 
-  BL::FluidDomainSettings b_domain;
+  BL::Object b_ob;
-  BL::Mesh b_mesh;
   AttributeStandard attribute;
 };
 

intern/cycles/bvh/bvh.cpp

@@ -229,6 +229,8 @@ void BVH::refit_primitives(int start, int end, BoundBox &bbox, uint &visibility)
 
         curve.bounds_grow(k, &hair->curve_keys[0], &hair->curve_radius[0], bbox);
 
+        visibility |= PATH_RAY_CURVE;
+
         /* Motion curves. */
         if (hair->use_motion_blur) {
           Attribute *attr = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
@@ -325,6 +327,9 @@ void BVH::pack_primitives()
         pack.prim_tri_index[i] = -1;
       }
       pack.prim_visibility[i] = ob->visibility_for_tracing();
+      if (pack.prim_type[i] & PRIMITIVE_ALL_CURVE) {
+        pack.prim_visibility[i] |= PATH_RAY_CURVE;
+      }
     }
     else {
       pack.prim_tri_index[i] = -1;

intern/cycles/bvh/bvh8.cpp

@@ -439,6 +439,8 @@ void BVH8::refit_node(int idx, bool leaf, BoundBox &bbox, uint &visibility)
 
           curve.bounds_grow(k, &hair->curve_keys[0], &hair->curve_radius[0], bbox);
 
+          visibility |= PATH_RAY_CURVE;
+
           /* Motion curves. */
           if (hair->use_motion_blur) {
             Attribute *attr = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);

intern/cycles/bvh/bvh_build.cpp

@@ -885,6 +885,9 @@ BVHNode *BVHBuild::create_leaf_node(const BVHRange &range, const vector<BVHRefer
 
       bounds[type_index].grow(ref.bounds());
       visibility[type_index] |= objects[ref.prim_object()]->visibility_for_tracing();
+      if (ref.prim_type() & PRIMITIVE_ALL_CURVE) {
+        visibility[type_index] |= PATH_RAY_CURVE;
+      }
       ++num_new_prims;
     }
     else {

intern/cycles/bvh/bvh_embree.cpp

@@ -16,6 +16,11 @@
 
 /* This class implements a ray accelerator for Cycles using Intel's Embree library.
  * It supports triangles, curves, object and deformation blur and instancing.
+ * Not supported are thick line segments, those have no native equivalent in Embree.
+ * They could be implemented using Embree's thick curves, at the expense of wasted memory.
+ * User defined intersections for Embree could also be an option, but since Embree only uses
+ * aligned BVHs for user geometry, this would come with reduced performance and/or higher memory
+ * usage.
  *
  * Since Embree allows object to be either curves or triangles but not both, Cycles object IDs are
  * mapped to Embree IDs by multiplying by two and adding one for curves.
@@ -770,21 +775,6 @@ void BVHEmbree::update_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair
       }
     }
   }
-#  if RTC_VERSION >= 30900
-  if (!use_curves) {
-    unsigned char *flags = (unsigned char *)rtcSetNewGeometryBuffer(geom_id,
-                                                                    RTC_BUFFER_TYPE_FLAGS,
-                                                                    0,
-                                                                    RTC_FORMAT_UCHAR,
-                                                                    sizeof(unsigned char),
-                                                                    num_keys_embree);
-    flags[0] = RTC_CURVE_FLAG_NEIGHBOR_RIGHT;
-    ::memset(flags + 1,
-             RTC_CURVE_FLAG_NEIGHBOR_RIGHT | RTC_CURVE_FLAG_NEIGHBOR_RIGHT,
-             num_keys_embree - 2);
-    flags[num_keys_embree - 1] = RTC_CURVE_FLAG_NEIGHBOR_LEFT;
-  }
-#  endif
 }
 
 void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
@@ -820,18 +810,10 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
   size_t prim_tri_index_size = pack.prim_index.size();
   pack.prim_tri_index.resize(prim_tri_index_size + num_segments);
 
-#  if RTC_VERSION >= 30900
-  enum RTCGeometryType type = (!use_curves) ?
-                                  (use_ribbons ? RTC_GEOMETRY_TYPE_FLAT_LINEAR_CURVE :
-                                                 RTC_GEOMETRY_TYPE_ROUND_LINEAR_CURVE) :
-                                  (use_ribbons ? RTC_GEOMETRY_TYPE_FLAT_CATMULL_ROM_CURVE :
-                                                 RTC_GEOMETRY_TYPE_ROUND_CATMULL_ROM_CURVE);
-#  else
   enum RTCGeometryType type = (!use_curves) ?
                                   RTC_GEOMETRY_TYPE_FLAT_LINEAR_CURVE :
                                   (use_ribbons ? RTC_GEOMETRY_TYPE_FLAT_CATMULL_ROM_CURVE :
                                                  RTC_GEOMETRY_TYPE_ROUND_CATMULL_ROM_CURVE);
-#  endif
 
   RTCGeometry geom_id = rtcNewGeometry(rtc_shared_device, type);
   rtcSetGeometryTessellationRate(geom_id, curve_subdivisions);

intern/cycles/device/cuda/device_cuda.h

@@ -51,7 +51,6 @@ class CUDADevice : public Device {
   size_t map_host_used;
   size_t map_host_limit;
   int can_map_host;
-  int pitch_alignment;
   int cuDevId;
   int cuDevArchitecture;
   bool first_error;
@@ -112,8 +111,6 @@ class CUDADevice : public Device {
 
   bool support_device(const DeviceRequestedFeatures & /*requested_features*/);
 
-  bool check_peer_access(Device *peer_device);
-
   bool use_adaptive_compilation();
 
   bool use_split_kernel();
@@ -226,7 +223,7 @@ class CUDADevice : public Device {
                               CUdeviceptr d_wtile,
                               CUstream stream = 0);
 
-  void render(DeviceTask &task, RenderTile &rtile, device_vector<WorkTile> &work_tiles);
+  void path_trace(DeviceTask &task, RenderTile &rtile, device_vector<WorkTile> &work_tiles);
 
   void film_convert(DeviceTask &task,
                     device_ptr buffer,

intern/cycles/device/cuda/device_cuda_impl.cpp

@@ -207,7 +207,6 @@ CUDADevice::CUDADevice(DeviceInfo &info, Stats &stats, Profiler &profiler, bool
   map_host_limit = 0;
   map_host_used = 0;
   can_map_host = 0;
-  pitch_alignment = 0;
 
   functions.loaded = false;
 
@@ -225,9 +224,6 @@ CUDADevice::CUDADevice(DeviceInfo &info, Stats &stats, Profiler &profiler, bool
   cuda_assert(
       cuDeviceGetAttribute(&can_map_host, CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, cuDevice));
 
-  cuda_assert(cuDeviceGetAttribute(
-      &pitch_alignment, CU_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT, cuDevice));
-
   unsigned int ctx_flags = CU_CTX_LMEM_RESIZE_TO_MAX;
   if (can_map_host) {
     ctx_flags |= CU_CTX_MAP_HOST;
@@ -290,49 +286,6 @@ bool CUDADevice::support_device(const DeviceRequestedFeatures & /*requested_feat
   return true;
 }
 
-bool CUDADevice::check_peer_access(Device *peer_device)
-{
-  if (peer_device == this) {
-    return false;
-  }
-  if (peer_device->info.type != DEVICE_CUDA && peer_device->info.type != DEVICE_OPTIX) {
-    return false;
-  }
-
-  CUDADevice *const peer_device_cuda = static_cast<CUDADevice *>(peer_device);
-
-  int can_access = 0;
-  cuda_assert(cuDeviceCanAccessPeer(&can_access, cuDevice, peer_device_cuda->cuDevice));
-  if (can_access == 0) {
-    return false;
-  }
-
-  // Ensure array access over the link is possible as well (for 3D textures)
-  cuda_assert(cuDeviceGetP2PAttribute(&can_access,
-                                      CU_DEVICE_P2P_ATTRIBUTE_ARRAY_ACCESS_ACCESS_SUPPORTED,
-                                      cuDevice,
-                                      peer_device_cuda->cuDevice));
-  if (can_access == 0) {
-    return false;
-  }
-
-  // Enable peer access in both directions
-  {
-    const CUDAContextScope scope(this);
-    if (cuda_error(cuCtxEnablePeerAccess(peer_device_cuda->cuContext, 0))) {
-      return false;
-    }
-  }
-  {
-    const CUDAContextScope scope(peer_device_cuda);
-    if (cuda_error(cuCtxEnablePeerAccess(cuContext, 0))) {
-      return false;
-    }
-  }
-
-  return true;
-}
-
 bool CUDADevice::use_adaptive_compilation()
 {
   return DebugFlags().cuda.adaptive_compile;
@@ -633,23 +586,20 @@ void CUDADevice::reserve_local_memory(const DeviceRequestedFeatures &requested_f
   cuMemGetInfo(&free_before, &total);
 
   /* Get kernel function. */
-  CUfunction cuRender;
+  CUfunction cuPathTrace;
 
-  if (requested_features.use_baking) {
+  if (requested_features.use_integrator_branched) {
-    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_bake"));
+    cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_branched_path_trace"));
-  }
-  else if (requested_features.use_integrator_branched) {
-    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_branched_path_trace"));
   }
   else {
-    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_path_trace"));
+    cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace"));
   }
 
-  cuda_assert(cuFuncSetCacheConfig(cuRender, CU_FUNC_CACHE_PREFER_L1));
+  cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1));
 
   int min_blocks, num_threads_per_block;
-  cuda_assert(
+  cuda_assert(cuOccupancyMaxPotentialBlockSize(
-      cuOccupancyMaxPotentialBlockSize(&min_blocks, &num_threads_per_block, cuRender, NULL, 0, 0));
+      &min_blocks, &num_threads_per_block, cuPathTrace, NULL, 0, 0));
 
   /* Launch kernel, using just 1 block appears sufficient to reserve
    * memory for all multiprocessors. It would be good to do this in
@@ -659,7 +609,7 @@ void CUDADevice::reserve_local_memory(const DeviceRequestedFeatures &requested_f
 
   void *args[] = {&d_work_tiles, &total_work_size};
 
-  cuda_assert(cuLaunchKernel(cuRender, 1, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
+  cuda_assert(cuLaunchKernel(cuPathTrace, 1, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
 
   cuda_assert(cuCtxSynchronize());
 
@@ -721,12 +671,6 @@ void CUDADevice::load_texture_info()
 
 void CUDADevice::move_textures_to_host(size_t size, bool for_texture)
 {
-  /* Break out of recursive call, which can happen when moving memory on a multi device. */
-  static bool any_device_moving_textures_to_host = false;
-  if (any_device_moving_textures_to_host) {
-    return;
-  }
-
   /* Signal to reallocate textures in host memory only. */
   move_texture_to_host = true;
 
@@ -740,12 +684,6 @@ void CUDADevice::move_textures_to_host(size_t size, bool for_texture)
       device_memory &mem = *pair.first;
       CUDAMem *cmem = &pair.second;
 
-      /* Can only move textures allocated on this device (and not those from peer devices).
-       * And need to ignore memory that is already on the host. */
-      if (!mem.is_resident(this) || cmem->use_mapped_host) {
-        continue;
-      }
-
       bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) &&
                         (&mem != &texture_info);
       bool is_image = is_texture && (mem.data_height > 1);
@@ -755,6 +693,11 @@ void CUDADevice::move_textures_to_host(size_t size, bool for_texture)
         continue;
       }
 
+      /* Already in host memory. */
+      if (cmem->use_mapped_host) {
+        continue;
+      }
+
       /* For other textures, only move image textures. */
       if (for_texture && !is_image) {
         continue;
@@ -777,30 +720,26 @@ void CUDADevice::move_textures_to_host(size_t size, bool for_texture)
       static thread_mutex move_mutex;
       thread_scoped_lock lock(move_mutex);
 
-      any_device_moving_textures_to_host = true;
+      /* Preserve the original device pointer, in case of multi device
+       * we can't change it because the pointer mapping would break. */
+      device_ptr prev_pointer = max_mem->device_pointer;
+      size_t prev_size = max_mem->device_size;
 
-      /* Potentially need to call back into multi device, so pointer mapping
+      mem_copy_to(*max_mem);
-       * and peer devices are updated. This is also necessary since the device
-       * pointer may just be a key here, so cannot be accessed and freed directly.
-       * Unfortunately it does mean that memory is reallocated on all other
-       * devices as well, which is potentially dangerous when still in use (since
-       * a thread rendering on another devices would only be caught in this mutex
-       * if it so happens to do an allocation at the same time as well. */
-      max_mem->device_copy_to();
       size = (max_size >= size) ? 0 : size - max_size;
 
-      any_device_moving_textures_to_host = false;
+      max_mem->device_pointer = prev_pointer;
+      max_mem->device_size = prev_size;
     }
     else {
       break;
     }
   }
 
-  /* Unset flag before texture info is reloaded, since it should stay in device memory. */
-  move_texture_to_host = false;
-
   /* Update texture info array with new pointers. */
   load_texture_info();
+
+  move_texture_to_host = false;
 }
 
 CUDADevice::CUDAMem *CUDADevice::generic_alloc(device_memory &mem, size_t pitch_padding)
@@ -866,6 +805,9 @@ CUDADevice::CUDAMem *CUDADevice::generic_alloc(device_memory &mem, size_t pitch_
       map_host_used += size;
       status = " in host memory";
     }
+    else {
+      status = " failed, out of host memory";
+    }
   }
 
   if (mem_alloc_result != CUDA_SUCCESS) {
@@ -961,7 +903,7 @@ void CUDADevice::generic_free(device_memory &mem)
     }
     else {
       /* Free device memory. */
-      cuda_assert(cuMemFree(mem.device_pointer));
+      cuMemFree(mem.device_pointer);
     }
 
     stats.mem_free(mem.device_size);
@@ -1087,17 +1029,18 @@ void CUDADevice::const_copy_to(const char *name, void *host, size_t size)
 
 void CUDADevice::global_alloc(device_memory &mem)
 {
-  if (mem.is_resident(this)) {
+  CUDAContextScope scope(this);
-    generic_alloc(mem);
+
-    generic_copy_to(mem);
+  generic_alloc(mem);
-  }
+  generic_copy_to(mem);
 
   const_copy_to(mem.name, &mem.device_pointer, sizeof(mem.device_pointer));
 }
 
 void CUDADevice::global_free(device_memory &mem)
 {
-  if (mem.is_resident(this) && mem.device_pointer) {
+  if (mem.device_pointer) {
+    CUDAContextScope scope(this);
     generic_free(mem);
   }
 }
@@ -1166,19 +1109,7 @@ void CUDADevice::tex_alloc(device_texture &mem)
   size_t src_pitch = mem.data_width * dsize * mem.data_elements;
   size_t dst_pitch = src_pitch;
 
-  if (!mem.is_resident(this)) {
+  if (mem.data_depth > 1) {
-    cmem = &cuda_mem_map[&mem];
-    cmem->texobject = 0;
-
-    if (mem.data_depth > 1) {
-      array_3d = (CUarray)mem.device_pointer;
-      cmem->array = array_3d;
-    }
-    else if (mem.data_height > 0) {
-      dst_pitch = align_up(src_pitch, pitch_alignment);
-    }
-  }
-  else if (mem.data_depth > 1) {
     /* 3D texture using array, there is no API for linear memory. */
     CUDA_ARRAY3D_DESCRIPTOR desc;
 
@@ -1222,7 +1153,10 @@ void CUDADevice::tex_alloc(device_texture &mem)
   }
   else if (mem.data_height > 0) {
     /* 2D texture, using pitch aligned linear memory. */
-    dst_pitch = align_up(src_pitch, pitch_alignment);
+    int alignment = 0;
+    cuda_assert(
+        cuDeviceGetAttribute(&alignment, CU_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT, cuDevice));
+    dst_pitch = align_up(src_pitch, alignment);
     size_t dst_size = dst_pitch * mem.data_height;
 
     cmem = generic_alloc(mem, dst_size - mem.memory_size());
@@ -1314,11 +1248,7 @@ void CUDADevice::tex_free(device_texture &mem)
       cuTexObjectDestroy(cmem.texobject);
     }
 
-    if (!mem.is_resident(this)) {
+    if (cmem.array) {
-      /* Do not free memory here, since it was allocated on a different device. */
-      cuda_mem_map.erase(cuda_mem_map.find(&mem));
-    }
-    else if (cmem.array) {
       /* Free array. */
       cuArrayDestroy(cmem.array);
       stats.mem_free(mem.device_size);
@@ -1850,7 +1780,9 @@ void CUDADevice::adaptive_sampling_post(RenderTile &rtile,
                              0));
 }
 
-void CUDADevice::render(DeviceTask &task, RenderTile &rtile, device_vector<WorkTile> &work_tiles)
+void CUDADevice::path_trace(DeviceTask &task,
+                            RenderTile &rtile,
+                            device_vector<WorkTile> &work_tiles)
 {
   scoped_timer timer(&rtile.buffers->render_time);
 
@@ -1858,24 +1790,21 @@ void CUDADevice::render(DeviceTask &task, RenderTile &rtile, device_vector<WorkT
     return;
 
   CUDAContextScope scope(this);
-  CUfunction cuRender;
+  CUfunction cuPathTrace;
 
   /* Get kernel function. */
-  if (rtile.task == RenderTile::BAKE) {
+  if (task.integrator_branched) {
-    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_bake"));
+    cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_branched_path_trace"));
-  }
-  else if (task.integrator_branched) {
-    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_branched_path_trace"));
   }
   else {
-    cuda_assert(cuModuleGetFunction(&cuRender, cuModule, "kernel_cuda_path_trace"));
+    cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace"));
   }
 
   if (have_error()) {
     return;
   }
 
-  cuda_assert(cuFuncSetCacheConfig(cuRender, CU_FUNC_CACHE_PREFER_L1));
+  cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1));
 
   /* Allocate work tile. */
   work_tiles.alloc(1);
@@ -1893,8 +1822,8 @@ void CUDADevice::render(DeviceTask &task, RenderTile &rtile, device_vector<WorkT
    * remain conservative for GPUs connected to a display to avoid driver
    * timeouts and display freezing. */
   int min_blocks, num_threads_per_block;
-  cuda_assert(
+  cuda_assert(cuOccupancyMaxPotentialBlockSize(
-      cuOccupancyMaxPotentialBlockSize(&min_blocks, &num_threads_per_block, cuRender, NULL, 0, 0));
+      &min_blocks, &num_threads_per_block, cuPathTrace, NULL, 0, 0));
   if (!info.display_device) {
     min_blocks *= 8;
   }
@@ -1922,7 +1851,7 @@ void CUDADevice::render(DeviceTask &task, RenderTile &rtile, device_vector<WorkT
     void *args[] = {&d_work_tiles, &total_work_size};
 
     cuda_assert(
-        cuLaunchKernel(cuRender, num_blocks, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
+        cuLaunchKernel(cuPathTrace, num_blocks, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
 
     /* Run the adaptive sampling kernels at selected samples aligned to step samples. */
     uint filter_sample = sample + wtile->num_samples - 1;
@@ -2028,7 +1957,10 @@ void CUDADevice::shader(DeviceTask &task)
   CUdeviceptr d_output = (CUdeviceptr)task.shader_output;
 
   /* get kernel function */
-  if (task.shader_eval_type == SHADER_EVAL_DISPLACE) {
+  if (task.shader_eval_type >= SHADER_EVAL_BAKE) {
+    cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_bake"));
+  }
+  else if (task.shader_eval_type == SHADER_EVAL_DISPLACE) {
     cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_displace"));
   }
   else {
@@ -2365,12 +2297,9 @@ void CUDADevice::thread_run(DeviceTask *task)
           split_kernel->path_trace(task, tile, void_buffer, void_buffer);
         }
         else {
-          render(*task, tile, work_tiles);
+          path_trace(*task, tile, work_tiles);
         }
       }
-      else if (tile.task == RenderTile::BAKE) {
-        render(*task, tile, work_tiles);
-      }
       else if (tile.task == RenderTile::DENOISE) {
         tile.sample = tile.start_sample + tile.num_samples;
 

intern/cycles/device/device.cpp

@@ -602,7 +602,6 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
   info.has_adaptive_stop_per_sample = true;
   info.has_osl = true;
   info.has_profiling = true;
-  info.has_peer_memory = false;
 
   foreach (const DeviceInfo &device, subdevices) {
     /* Ensure CPU device does not slow down GPU. */
@@ -646,7 +645,6 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
     info.has_adaptive_stop_per_sample &= device.has_adaptive_stop_per_sample;
     info.has_osl &= device.has_osl;
     info.has_profiling &= device.has_profiling;
-    info.has_peer_memory |= device.has_peer_memory;
   }
 
   return info;

intern/cycles/device/device.h

@@ -82,7 +82,6 @@ class DeviceInfo {
   bool has_osl;                      /* Support Open Shading Language. */
   bool use_split_kernel;             /* Use split or mega kernel. */
   bool has_profiling;                /* Supports runtime collection of profiling info. */
-  bool has_peer_memory;              /* GPU has P2P access to memory of another GPU. */
   int cpu_threads;
   vector<DeviceInfo> multi_devices;
   vector<DeviceInfo> denoising_devices;
@@ -100,7 +99,6 @@ class DeviceInfo {
     has_osl = false;
     use_split_kernel = false;
     has_profiling = false;
-    has_peer_memory = false;
   }
 
   bool operator==(const DeviceInfo &info)
@@ -437,17 +435,6 @@ class Device {
   {
   }
 
-  virtual bool is_resident(device_ptr /*key*/, Device *sub_device)
-  {
-    /* Memory is always resident if this is not a multi device, regardless of whether the pointer
-     * is valid or not (since it may not have been allocated yet). */
-    return sub_device == this;
-  }
-  virtual bool check_peer_access(Device * /*peer_device*/)
-  {
-    return false;
-  }
-
   /* static */
   static Device *create(DeviceInfo &info,
                         Stats &stats,

intern/cycles/device/device_cpu.cpp

@@ -188,7 +188,6 @@ class CPUDevice : public Device {
       convert_to_byte_kernel;
   KernelFunctions<void (*)(KernelGlobals *, uint4 *, float4 *, int, int, int, int, int)>
       shader_kernel;
-  KernelFunctions<void (*)(KernelGlobals *, float *, int, int, int, int, int)> bake_kernel;
 
   KernelFunctions<void (*)(
       int, TileInfo *, int, int, float *, float *, float *, float *, float *, int *, int, int)>
@@ -271,7 +270,6 @@ class CPUDevice : public Device {
         REGISTER_KERNEL(convert_to_half_float),
         REGISTER_KERNEL(convert_to_byte),
         REGISTER_KERNEL(shader),
-        REGISTER_KERNEL(bake),
         REGISTER_KERNEL(filter_divide_shadow),
         REGISTER_KERNEL(filter_get_feature),
         REGISTER_KERNEL(filter_write_feature),
@@ -897,7 +895,7 @@ class CPUDevice : public Device {
     }
   }
 
-  void render(DeviceTask &task, RenderTile &tile, KernelGlobals *kg)
+  void path_trace(DeviceTask &task, RenderTile &tile, KernelGlobals *kg)
   {
     const bool use_coverage = kernel_data.film.cryptomatte_passes & CRYPT_ACCURATE;
 
@@ -921,21 +919,12 @@ class CPUDevice : public Device {
           break;
       }
 
-      if (tile.task == RenderTile::PATH_TRACE) {
+      for (int y = tile.y; y < tile.y + tile.h; y++) {
-        for (int y = tile.y; y < tile.y + tile.h; y++) {
+        for (int x = tile.x; x < tile.x + tile.w; x++) {
-          for (int x = tile.x; x < tile.x + tile.w; x++) {
+          if (use_coverage) {
-            if (use_coverage) {
+            coverage.init_pixel(x, y);
-              coverage.init_pixel(x, y);
-            }
-            path_trace_kernel()(kg, render_buffer, sample, x, y, tile.offset, tile.stride);
-          }
-        }
-      }
-      else {
-        for (int y = tile.y; y < tile.y + tile.h; y++) {
-          for (int x = tile.x; x < tile.x + tile.w; x++) {
-            bake_kernel()(kg, render_buffer, sample, x, y, tile.offset, tile.stride);
           }
+          path_trace_kernel()(kg, render_buffer, sample, x, y, tile.offset, tile.stride);
         }
       }
       tile.sample = sample + 1;
@@ -1030,12 +1019,9 @@ class CPUDevice : public Device {
           split_kernel->path_trace(&task, tile, kgbuffer, void_buffer);
         }
         else {
-          render(task, tile, kg);
+          path_trace(task, tile, kg);
         }
       }
-      else if (tile.task == RenderTile::BAKE) {
-        render(task, tile, kg);
-      }
       else if (tile.task == RenderTile::DENOISE) {
         denoise(denoising, tile);
         task.update_progress(&tile, tile.w * tile.h);

intern/cycles/device/device_cuda.cpp

@@ -131,15 +131,6 @@ void device_cuda_info(vector<DeviceInfo> &devices)
     info.has_volume_decoupled = false;
     info.has_adaptive_stop_per_sample = false;
 
-    /* Check if the device has P2P access to any other device in the system. */
-    for (int peer_num = 0; peer_num < count && !info.has_peer_memory; peer_num++) {
-      if (num != peer_num) {
-        int can_access = 0;
-        cuDeviceCanAccessPeer(&can_access, num, peer_num);
-        info.has_peer_memory = (can_access != 0);
-      }
-    }
-
     int pci_location[3] = {0, 0, 0};
     cuDeviceGetAttribute(&pci_location[0], CU_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID, num);
     cuDeviceGetAttribute(&pci_location[1], CU_DEVICE_ATTRIBUTE_PCI_BUS_ID, num);

intern/cycles/device/device_memory.cpp

@@ -125,11 +125,6 @@ void device_memory::restore_device()
   device_pointer = original_device_ptr;
 }
 
-bool device_memory::is_resident(Device *sub_device) const
-{
-  return device->is_resident(device_pointer, sub_device);
-}
-
 /* Device Sub Ptr */
 
 device_sub_ptr::device_sub_ptr(device_memory &mem, int offset, int size) : device(mem.device)

intern/cycles/device/device_memory.h

@@ -230,8 +230,6 @@ class device_memory {
   void swap_device(Device *new_device, size_t new_device_size, device_ptr new_device_ptr);
   void restore_device();
 
-  bool is_resident(Device *sub_device) const;
-
  protected:
   friend class CUDADevice;
   friend class OptiXDevice;

intern/cycles/device/device_multi.cpp

@@ -34,88 +34,37 @@ CCL_NAMESPACE_BEGIN
 class MultiDevice : public Device {
  public:
   struct SubDevice {
-    Stats stats;
+    explicit SubDevice(Device *device_) : device(device_)
+    {
+    }
+
     Device *device;
     map<device_ptr, device_ptr> ptr_map;
-    int peer_island_index = -1;
   };
 
   list<SubDevice> devices, denoising_devices;
   device_ptr unique_key;
-  vector<vector<SubDevice *>> peer_islands;
-  bool matching_rendering_and_denoising_devices;
 
   MultiDevice(DeviceInfo &info, Stats &stats, Profiler &profiler, bool background_)
       : Device(info, stats, profiler, background_), unique_key(1)
   {
     foreach (DeviceInfo &subinfo, info.multi_devices) {
+      Device *device = Device::create(subinfo, sub_stats_, profiler, background);
+
       /* Always add CPU devices at the back since GPU devices can change
        * host memory pointers, which CPU uses as device pointer. */
-      SubDevice *sub;
       if (subinfo.type == DEVICE_CPU) {
-        devices.emplace_back();
+        devices.push_back(SubDevice(device));
-        sub = &devices.back();
       }
       else {
-        devices.emplace_front();
+        devices.push_front(SubDevice(device));
-        sub = &devices.front();
       }
-
-      /* The pointer to 'sub->stats' will stay valid even after new devices
-       * are added, since 'devices' is a linked list. */
-      sub->device = Device::create(subinfo, sub->stats, profiler, background);
     }
 
     foreach (DeviceInfo &subinfo, info.denoising_devices) {
-      denoising_devices.emplace_front();
+      Device *device = Device::create(subinfo, sub_stats_, profiler, background);
-      SubDevice *sub = &denoising_devices.front();
 
-      sub->device = Device::create(subinfo, sub->stats, profiler, background);
+      denoising_devices.push_back(SubDevice(device));
-    }
-
-    /* Build a list of peer islands for the available render devices */
-    foreach (SubDevice &sub, devices) {
-      /* First ensure that every device is in at least once peer island */
-      if (sub.peer_island_index < 0) {
-        peer_islands.emplace_back();
-        sub.peer_island_index = (int)peer_islands.size() - 1;
-        peer_islands[sub.peer_island_index].push_back(&sub);
-      }
-
-      if (!info.has_peer_memory) {
-        continue;
-      }
-
-      /* Second check peer access between devices and fill up the islands accordingly */
-      foreach (SubDevice &peer_sub, devices) {
-        if (peer_sub.peer_island_index < 0 &&
-            peer_sub.device->info.type == sub.device->info.type &&
-            peer_sub.device->check_peer_access(sub.device)) {
-          peer_sub.peer_island_index = sub.peer_island_index;
-          peer_islands[sub.peer_island_index].push_back(&peer_sub);
-        }
-      }
-    }
-
-    /* Try to re-use memory when denoising and render devices use the same physical devices
-     * (e.g. OptiX denoising and CUDA rendering device pointing to the same GPU).
-     * Ordering has to match as well, so that 'DeviceTask::split' behaves consistent. */
-    matching_rendering_and_denoising_devices = denoising_devices.empty() ||
-                                               (devices.size() == denoising_devices.size());
-    if (matching_rendering_and_denoising_devices) {
-      for (list<SubDevice>::iterator device_it = devices.begin(),
-                                     denoising_device_it = denoising_devices.begin();
-           device_it != devices.end() && denoising_device_it != denoising_devices.end();
-           ++device_it, ++denoising_device_it) {
-        const DeviceInfo &info = device_it->device->info;
-        const DeviceInfo &denoising_info = denoising_device_it->device->info;
-        if ((info.type != DEVICE_CUDA && info.type != DEVICE_OPTIX) ||
-            (denoising_info.type != DEVICE_CUDA && denoising_info.type != DEVICE_OPTIX) ||
-            info.num != denoising_info.num) {
-          matching_rendering_and_denoising_devices = false;
-          break;
-        }
-      }
     }
 
 #ifdef WITH_NETWORK
@@ -226,11 +175,11 @@ class MultiDevice : public Device {
 
   bool build_optix_bvh(BVH *bvh)
   {
-    /* Broadcast acceleration structure build to all render devices */
+    // Broadcast acceleration structure build to all render devices
-    foreach (SubDevice &sub, devices) {
+    foreach (SubDevice &sub, devices)
       if (!sub.device->build_optix_bvh(bvh))
         return false;
-    }
+
     return true;
   }
 
@@ -242,82 +191,17 @@ class MultiDevice : public Device {
     return devices.front().device->osl_memory();
   }
 
-  bool is_resident(device_ptr key, Device *sub_device)
-  {
-    foreach (SubDevice &sub, devices) {
-      if (sub.device == sub_device) {
-        return find_matching_mem_device(key, sub)->device == sub_device;
-      }
-    }
-    return false;
-  }
-
-  SubDevice *find_matching_mem_device(device_ptr key, SubDevice &sub)
-  {
-    assert(key != 0 && (sub.peer_island_index >= 0 || sub.ptr_map.find(key) != sub.ptr_map.end()));
-
-    /* Get the memory owner of this key (first try current device, then peer devices) */
-    SubDevice *owner_sub = &sub;
-    if (owner_sub->ptr_map.find(key) == owner_sub->ptr_map.end()) {
-      foreach (SubDevice *island_sub, peer_islands[sub.peer_island_index]) {
-        if (island_sub != owner_sub &&
-            island_sub->ptr_map.find(key) != island_sub->ptr_map.end()) {
-          owner_sub = island_sub;
-        }
-      }
-    }
-    return owner_sub;
-  }
-
-  SubDevice *find_suitable_mem_device(device_ptr key, const vector<SubDevice *> &island)
-  {
-    assert(!island.empty());
-
-    /* Get the memory owner of this key or the device with the lowest memory usage when new */
-    SubDevice *owner_sub = island.front();
-    foreach (SubDevice *island_sub, island) {
-      if (key ? (island_sub->ptr_map.find(key) != island_sub->ptr_map.end()) :
-                (island_sub->device->stats.mem_used < owner_sub->device->stats.mem_used)) {
-        owner_sub = island_sub;
-      }
-    }
-    return owner_sub;
-  }
-
-  inline device_ptr find_matching_mem(device_ptr key, SubDevice &sub)
-  {
-    return find_matching_mem_device(key, sub)->ptr_map[key];
-  }
-
   void mem_alloc(device_memory &mem)
   {
     device_ptr key = unique_key++;
 
-    if (mem.type == MEM_PIXELS) {
+    foreach (SubDevice &sub, devices) {
-      /* Always allocate pixels memory on all devices
+      mem.device = sub.device;
-       * This is necessary to ensure PBOs are registered everywhere, which FILM_CONVERT uses */
+      mem.device_pointer = 0;
-      foreach (SubDevice &sub, devices) {
+      mem.device_size = 0;
-        mem.device = sub.device;
-        mem.device_pointer = 0;
-        mem.device_size = 0;
 
-        sub.device->mem_alloc(mem);
+      sub.device->mem_alloc(mem);
-        sub.ptr_map[key] = mem.device_pointer;
+      sub.ptr_map[key] = mem.device_pointer;
-      }
-    }
-    else {
-      assert(mem.type == MEM_READ_ONLY || mem.type == MEM_READ_WRITE ||
-             mem.type == MEM_DEVICE_ONLY);
-      /* The remaining memory types can be distributed across devices */
-      foreach (const vector<SubDevice *> &island, peer_islands) {
-        SubDevice *owner_sub = find_suitable_mem_device(key, island);
-        mem.device = owner_sub->device;
-        mem.device_pointer = 0;
-        mem.device_size = 0;
-
-        owner_sub->device->mem_alloc(mem);
-        owner_sub->ptr_map[key] = mem.device_pointer;
-      }
     }
 
     mem.device = this;
@@ -331,36 +215,13 @@ class MultiDevice : public Device {
     device_ptr key = (existing_key) ? existing_key : unique_key++;
     size_t existing_size = mem.device_size;
 
-    /* The tile buffers are allocated on each device (see below), so copy to all of them */
+    foreach (SubDevice &sub, devices) {
-    if (strcmp(mem.name, "RenderBuffers") == 0) {
+      mem.device = sub.device;
-      foreach (SubDevice &sub, devices) {
+      mem.device_pointer = (existing_key) ? sub.ptr_map[existing_key] : 0;
-        mem.device = sub.device;
+      mem.device_size = existing_size;
-        mem.device_pointer = (existing_key) ? sub.ptr_map[existing_key] : 0;
-        mem.device_size = existing_size;
 
-        sub.device->mem_copy_to(mem);
+      sub.device->mem_copy_to(mem);
-        sub.ptr_map[key] = mem.device_pointer;
+      sub.ptr_map[key] = mem.device_pointer;
-      }
-    }
-    else {
-      foreach (const vector<SubDevice *> &island, peer_islands) {
-        SubDevice *owner_sub = find_suitable_mem_device(existing_key, island);
-        mem.device = owner_sub->device;
-        mem.device_pointer = (existing_key) ? owner_sub->ptr_map[existing_key] : 0;
-        mem.device_size = existing_size;
-
-        owner_sub->device->mem_copy_to(mem);
-        owner_sub->ptr_map[key] = mem.device_pointer;
-
-        if (mem.type == MEM_GLOBAL || mem.type == MEM_TEXTURE) {
-          /* Need to create texture objects and update pointer in kernel globals on all devices */
-          foreach (SubDevice *island_sub, island) {
-            if (island_sub != owner_sub) {
-              island_sub->device->mem_copy_to(mem);
-            }
-          }
-        }
-      }
     }
 
     mem.device = this;
@@ -377,11 +238,10 @@ class MultiDevice : public Device {
       int sy = y + i * sub_h;
       int sh = (i == (int)devices.size() - 1) ? h - sub_h * i : sub_h;
 
-      SubDevice *owner_sub = find_matching_mem_device(key, sub);
+      mem.device = sub.device;
-      mem.device = owner_sub->device;
+      mem.device_pointer = sub.ptr_map[key];
-      mem.device_pointer = owner_sub->ptr_map[key];
 
-      owner_sub->device->mem_copy_from(mem, sy, w, sh, elem);
+      sub.device->mem_copy_from(mem, sy, w, sh, elem);
       i++;
     }
 
@@ -395,47 +255,23 @@ class MultiDevice : public Device {
     device_ptr key = (existing_key) ? existing_key : unique_key++;
     size_t existing_size = mem.device_size;
 
-    /* This is a hack to only allocate the tile buffers on denoising devices
+    foreach (SubDevice &sub, devices) {
-     * Similarily the tile buffers also need to be allocated separately on all devices so any
+      mem.device = sub.device;
-     * overlap rendered for denoising does not interfer with each other */
+      mem.device_pointer = (existing_key) ? sub.ptr_map[existing_key] : 0;
-    if (strcmp(mem.name, "RenderBuffers") == 0) {
+      mem.device_size = existing_size;
-      vector<device_ptr> device_pointers;
-      device_pointers.reserve(devices.size());
 
-      foreach (SubDevice &sub, devices) {
+      sub.device->mem_zero(mem);
+      sub.ptr_map[key] = mem.device_pointer;
+    }
+
+    if (strcmp(mem.name, "RenderBuffers") == 0) {
+      foreach (SubDevice &sub, denoising_devices) {
         mem.device = sub.device;
         mem.device_pointer = (existing_key) ? sub.ptr_map[existing_key] : 0;
         mem.device_size = existing_size;
 
         sub.device->mem_zero(mem);
         sub.ptr_map[key] = mem.device_pointer;
-
-        device_pointers.push_back(mem.device_pointer);
-      }
-      foreach (SubDevice &sub, denoising_devices) {
-        if (matching_rendering_and_denoising_devices) {
-          sub.ptr_map[key] = device_pointers.front();
-          device_pointers.erase(device_pointers.begin());
-        }
-        else {
-          mem.device = sub.device;
-          mem.device_pointer = (existing_key) ? sub.ptr_map[existing_key] : 0;
-          mem.device_size = existing_size;
-
-          sub.device->mem_zero(mem);
-          sub.ptr_map[key] = mem.device_pointer;
-        }
-      }
-    }
-    else {
-      foreach (const vector<SubDevice *> &island, peer_islands) {
-        SubDevice *owner_sub = find_suitable_mem_device(existing_key, island);
-        mem.device = owner_sub->device;
-        mem.device_pointer = (existing_key) ? owner_sub->ptr_map[existing_key] : 0;
-        mem.device_size = existing_size;
-
-        owner_sub->device->mem_zero(mem);
-        owner_sub->ptr_map[key] = mem.device_pointer;
       }
     }
 
@@ -449,9 +285,17 @@ class MultiDevice : public Device {
     device_ptr key = mem.device_pointer;
     size_t existing_size = mem.device_size;
 
-    /* Free memory that was allocated for all devices (see above) on each device */
+    foreach (SubDevice &sub, devices) {
-    if (strcmp(mem.name, "RenderBuffers") == 0 || mem.type == MEM_PIXELS) {
+      mem.device = sub.device;
-      foreach (SubDevice &sub, devices) {
+      mem.device_pointer = sub.ptr_map[key];
+      mem.device_size = existing_size;
+
+      sub.device->mem_free(mem);
+      sub.ptr_map.erase(sub.ptr_map.find(key));
+    }
+
+    if (strcmp(mem.name, "RenderBuffers") == 0) {
+      foreach (SubDevice &sub, denoising_devices) {
         mem.device = sub.device;
         mem.device_pointer = sub.ptr_map[key];
         mem.device_size = existing_size;
@@ -459,39 +303,6 @@ class MultiDevice : public Device {
         sub.device->mem_free(mem);
         sub.ptr_map.erase(sub.ptr_map.find(key));
       }
-      foreach (SubDevice &sub, denoising_devices) {
-        if (matching_rendering_and_denoising_devices) {
-          sub.ptr_map.erase(key);
-        }
-        else {
-          mem.device = sub.device;
-          mem.device_pointer = sub.ptr_map[key];
-          mem.device_size = existing_size;
-
-          sub.device->mem_free(mem);
-          sub.ptr_map.erase(sub.ptr_map.find(key));
-        }
-      }
-    }
-    else {
-      foreach (const vector<SubDevice *> &island, peer_islands) {
-        SubDevice *owner_sub = find_matching_mem_device(key, *island.front());
-        mem.device = owner_sub->device;
-        mem.device_pointer = owner_sub->ptr_map[key];
-        mem.device_size = existing_size;
-
-        owner_sub->device->mem_free(mem);
-        owner_sub->ptr_map.erase(owner_sub->ptr_map.find(key));
-
-        if (mem.type == MEM_TEXTURE) {
-          /* Free texture objects on all devices */
-          foreach (SubDevice *island_sub, island) {
-            if (island_sub != owner_sub) {
-              island_sub->device->mem_free(mem);
-            }
-          }
-        }
-      }
     }
 
     mem.device = this;
@@ -519,8 +330,6 @@ class MultiDevice : public Device {
                    bool transparent,
                    const DeviceDrawParams &draw_params)
   {
-    assert(rgba.type == MEM_PIXELS);
-
     device_ptr key = rgba.device_pointer;
     int i = 0, sub_h = h / devices.size();
     int sub_height = height / devices.size();
@@ -549,7 +358,7 @@ class MultiDevice : public Device {
 
     foreach (SubDevice &sub, devices) {
       if (sub.device == sub_device) {
-        tile.buffer = find_matching_mem(tile.buffer, sub);
+        tile.buffer = sub.ptr_map[tile.buffer];
         return;
       }
     }
@@ -591,7 +400,7 @@ class MultiDevice : public Device {
       device_vector<float> &mem = tiles[i].buffers->buffer;
       tiles[i].buffer = mem.device_pointer;
 
-      if (mem.device == this && matching_rendering_and_denoising_devices) {
+      if (mem.device == this && denoising_devices.empty()) {
         /* Skip unnecessary copies in viewport mode (buffer covers the
          * whole image), but still need to fix up the tile device pointer. */
         map_tile(sub_device, tiles[i]);
@@ -635,7 +444,7 @@ class MultiDevice : public Device {
   {
     device_vector<float> &mem = tiles[9].buffers->buffer;
 
-    if (mem.device == this && matching_rendering_and_denoising_devices) {
+    if (mem.device == this && denoising_devices.empty()) {
       return;
     }
 
@@ -709,22 +518,17 @@ class MultiDevice : public Device {
         tasks.pop_front();
 
         if (task.buffer)
-          subtask.buffer = find_matching_mem(task.buffer, sub);
+          subtask.buffer = sub.ptr_map[task.buffer];
         if (task.rgba_byte)
           subtask.rgba_byte = sub.ptr_map[task.rgba_byte];
         if (task.rgba_half)
           subtask.rgba_half = sub.ptr_map[task.rgba_half];
         if (task.shader_input)
-          subtask.shader_input = find_matching_mem(task.shader_input, sub);
+          subtask.shader_input = sub.ptr_map[task.shader_input];
         if (task.shader_output)
-          subtask.shader_output = find_matching_mem(task.shader_output, sub);
+          subtask.shader_output = sub.ptr_map[task.shader_output];
 
         sub.device->task_add(subtask);
-
-        if (task.buffers && task.buffers->buffer.device == this) {
-          /* Synchronize access to RenderBuffers, since 'map_neighbor_tiles' is not thread-safe. */
-          sub.device->task_wait();
-        }
       }
     }
   }
@@ -744,6 +548,9 @@ class MultiDevice : public Device {
     foreach (SubDevice &sub, denoising_devices)
       sub.device->task_cancel();
   }
+
+ protected:
+  Stats sub_stats_;
 };
 
 Device *device_multi_create(DeviceInfo &info, Stats &stats, Profiler &profiler, bool background)

intern/cycles/device/device_optix.cpp

@@ -924,8 +924,7 @@ class OptiXDevice : public CUDADevice {
                              &rtiles[9].h,
                              &rtiles[9].offset,
                              &rtiles[9].stride,
-                             &task.pass_stride,
+                             &task.pass_stride};
-                             &rtile.sample};
       launch_filter_kernel(
           "kernel_cuda_filter_convert_from_rgb", rtiles[9].w, rtiles[9].h, output_args);
 #  endif
@@ -1537,22 +1536,34 @@ bool device_optix_init()
 
 void device_optix_info(const vector<DeviceInfo> &cuda_devices, vector<DeviceInfo> &devices)
 {
-  devices.reserve(cuda_devices.size());
-
   // Simply add all supported CUDA devices as OptiX devices again
-  for (DeviceInfo info : cuda_devices) {
+  for (const DeviceInfo &cuda_info : cuda_devices) {
+    DeviceInfo info = cuda_info;
     assert(info.type == DEVICE_CUDA);
-
-    int major;
-    cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, info.num);
-    if (major < 5) {
-      continue;  // Only Maxwell and up are supported by OptiX
-    }
-
     info.type = DEVICE_OPTIX;
     info.id += "_OptiX";
 
-    devices.push_back(info);
+    // Figure out RTX support
+    CUdevice cuda_device = 0;
+    CUcontext cuda_context = NULL;
+    unsigned int rtcore_version = 0;
+    if (cuDeviceGet(&cuda_device, info.num) == CUDA_SUCCESS &&
+        cuDevicePrimaryCtxRetain(&cuda_context, cuda_device) == CUDA_SUCCESS) {
+      OptixDeviceContext optix_context = NULL;
+      if (optixDeviceContextCreate(cuda_context, nullptr, &optix_context) == OPTIX_SUCCESS) {
+        optixDeviceContextGetProperty(optix_context,
+                                      OPTIX_DEVICE_PROPERTY_RTCORE_VERSION,
+                                      &rtcore_version,
+                                      sizeof(rtcore_version));
+        optixDeviceContextDestroy(optix_context);
+      }
+      cuDevicePrimaryCtxRelease(cuda_device);
+    }
+
+    // Only add devices with RTX support
+    if (rtcore_version != 0 || getenv("CYCLES_OPTIX_TEST")) {
+      devices.push_back(info);
+    }
   }
 }
 

intern/cycles/device/device_task.cpp

@@ -44,8 +44,7 @@ DeviceTask::DeviceTask(Type type_)
       shader_eval_type(0),
       shader_filter(0),
       shader_x(0),
-      shader_w(0),
+      shader_w(0)
-      buffers(nullptr)
 {
   last_update_time = time_dt();
 }

intern/cycles/device/opencl/device_opencl.h

@@ -451,7 +451,6 @@ class OpenCLDevice : public Device {
                     device_ptr rgba_half);
   void shader(DeviceTask &task);
   void update_adaptive(DeviceTask &task, RenderTile &tile, int sample);
-  void bake(DeviceTask &task, RenderTile &tile);
 
   void denoise(RenderTile &tile, DenoisingTask &denoising);
 

intern/cycles/device/opencl/device_opencl_impl.cpp

@@ -1367,9 +1367,6 @@ void OpenCLDevice::thread_run(DeviceTask *task)
          */
         clFinish(cqCommandQueue);
       }
-      else if (tile.task == RenderTile::BAKE) {
-        bake(*task, tile);
-      }
       else if (tile.task == RenderTile::DENOISE) {
         tile.sample = tile.start_sample + tile.num_samples;
         denoise(tile, denoising);
@@ -1861,7 +1858,10 @@ void OpenCLDevice::shader(DeviceTask &task)
   cl_int d_offset = task.offset;
 
   OpenCLDevice::OpenCLProgram *program = &background_program;
-  if (task.shader_eval_type == SHADER_EVAL_DISPLACE) {
+  if (task.shader_eval_type >= SHADER_EVAL_BAKE) {
+    program = &bake_program;
+  }
+  else if (task.shader_eval_type == SHADER_EVAL_DISPLACE) {
     program = &displace_program;
   }
   program->wait_for_availability();
@@ -1892,51 +1892,6 @@ void OpenCLDevice::shader(DeviceTask &task)
   }
 }
 
-void OpenCLDevice::bake(DeviceTask &task, RenderTile &rtile)
-{
-  scoped_timer timer(&rtile.buffers->render_time);
-
-  /* Cast arguments to cl types. */
-  cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
-  cl_mem d_buffer = CL_MEM_PTR(rtile.buffer);
-  cl_int d_x = rtile.x;
-  cl_int d_y = rtile.y;
-  cl_int d_w = rtile.w;
-  cl_int d_h = rtile.h;
-  cl_int d_offset = rtile.offset;
-  cl_int d_stride = rtile.stride;
-
-  bake_program.wait_for_availability();
-  cl_kernel kernel = bake_program();
-
-  cl_uint start_arg_index = kernel_set_args(kernel, 0, d_data, d_buffer);
-
-  set_kernel_arg_buffers(kernel, &start_arg_index);
-
-  start_arg_index += kernel_set_args(
-      kernel, start_arg_index, d_x, d_y, d_w, d_h, d_offset, d_stride);
-
-  int start_sample = rtile.start_sample;
-  int end_sample = rtile.start_sample + rtile.num_samples;
-
-  for (int sample = start_sample; sample < end_sample; sample++) {
-    if (task.get_cancel()) {
-      if (task.need_finish_queue == false)
-        break;
-    }
-
-    kernel_set_args(kernel, start_arg_index, sample);
-
-    enqueue_kernel(kernel, d_w, d_h);
-
-    rtile.sample = sample + 1;
-
-    task.update_progress(&rtile, rtile.w * rtile.h);
-  }
-
-  clFinish(cqCommandQueue);
-}
-
 string OpenCLDevice::kernel_build_options(const string *debug_src)
 {
   string build_options = "-cl-no-signed-zeros -cl-mad-enable ";

intern/cycles/kernel/bvh/bvh.h

@@ -300,9 +300,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals *kg,
   // Is set to zero on miss or if ray is aborted, so can be used as return value
   uint p5 = max_hits;
 
-  if (local_isect) {
+  local_isect->num_hits = 0;  // Initialize hit count to zero
-    local_isect->num_hits = 0;  // Initialize hit count to zero
-  }
   optixTrace(scene_intersect_valid(ray) ? kernel_data.bvh.scene : 0,
              ray->P,
              ray->D,
@@ -325,9 +323,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals *kg,
   return p5;
 #  else /* __KERNEL_OPTIX__ */
   if (!scene_intersect_valid(ray)) {
-    if (local_isect) {
+    local_isect->num_hits = 0;
-      local_isect->num_hits = 0;
-    }
     return false;
   }
 

intern/cycles/kernel/closure/bsdf.h

@@ -97,18 +97,6 @@ ccl_device_inline float bump_shadowing_term(float3 Ng, float3 N, float3 I)
   return -g2 * g + g2 + g;
 }
 
-/* Shadow terminator workaround, taken from Appleseed.
- * Original code is under the MIT License
- * Copyright (c) 2019 Francois Beaune, The appleseedhq Organization */
-ccl_device_inline float shift_cos_in(float cos_in, const float frequency_multiplier)
-{
-  cos_in = min(cos_in, 1.0f);
-
-  const float angle = fast_acosf(cos_in);
-  const float val = max(cosf(angle * frequency_multiplier), 0.0f) / cos_in;
-  return val;
-}
-
 ccl_device_inline int bsdf_sample(KernelGlobals *kg,
                                   ShaderData *sd,
                                   const ShaderClosure *sc,
@@ -241,6 +229,8 @@ ccl_device_inline int bsdf_sample(KernelGlobals *kg,
     case CLOSURE_BSDF_MICROFACET_GGX_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_FRESNEL_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
+    case CLOSURE_BSDF_MICROFACET_GGX_ANISO_ID:
+    case CLOSURE_BSDF_MICROFACET_GGX_ANISO_FRESNEL_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
       label = bsdf_microfacet_ggx_sample(kg,
                                          sc,
@@ -291,6 +281,7 @@ ccl_device_inline int bsdf_sample(KernelGlobals *kg,
                                                      &sd->lcg_state);
       break;
     case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
+    case CLOSURE_BSDF_MICROFACET_BECKMANN_ANISO_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID:
       label = bsdf_microfacet_beckmann_sample(kg,
                                               sc,
@@ -307,6 +298,7 @@ ccl_device_inline int bsdf_sample(KernelGlobals *kg,
                                               pdf);
       break;
     case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
+    case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ANISO_ID:
       label = bsdf_ashikhmin_shirley_sample(sc,
                                             sd->Ng,
                                             sd->I,
@@ -456,17 +448,9 @@ ccl_device_inline int bsdf_sample(KernelGlobals *kg,
       }
     }
   }
-  else {
+  else if (label & LABEL_DIFFUSE) {
-    /* Shadow terminator offset. */
+    if (!isequal_float3(sc->N, sd->N)) {
-    const float frequency_multiplier =
+      *eval *= bump_shadowing_term((label & LABEL_TRANSMIT) ? -sd->N : sd->N, sc->N, *omega_in);
-        kernel_tex_fetch(__objects, sd->object).shadow_terminator_offset;
-    if (frequency_multiplier > 1.0f) {
-      *eval *= shift_cos_in(dot(*omega_in, sc->N), frequency_multiplier);
-    }
-    if (label & LABEL_DIFFUSE) {
-      if (!isequal_float3(sc->N, sd->N)) {
-        *eval *= bump_shadowing_term((label & LABEL_TRANSMIT) ? -sd->N : sd->N, sc->N, *omega_in);
-      }
     }
   }
 
@@ -520,6 +504,8 @@ ccl_device_inline
       case CLOSURE_BSDF_MICROFACET_GGX_ID:
       case CLOSURE_BSDF_MICROFACET_GGX_FRESNEL_ID:
       case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
+      case CLOSURE_BSDF_MICROFACET_GGX_ANISO_ID:
+      case CLOSURE_BSDF_MICROFACET_GGX_ANISO_FRESNEL_ID:
       case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
         eval = bsdf_microfacet_ggx_eval_reflect(sc, sd->I, omega_in, pdf);
         break;
@@ -533,10 +519,12 @@ ccl_device_inline
             sc, sd->I, omega_in, pdf, &sd->lcg_state);
         break;
       case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
+      case CLOSURE_BSDF_MICROFACET_BECKMANN_ANISO_ID:
       case CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID:
         eval = bsdf_microfacet_beckmann_eval_reflect(sc, sd->I, omega_in, pdf);
         break;
       case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
+      case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ANISO_ID:
         eval = bsdf_ashikhmin_shirley_eval_reflect(sc, sd->I, omega_in, pdf);
         break;
       case CLOSURE_BSDF_ASHIKHMIN_VELVET_ID:
@@ -581,12 +569,6 @@ ccl_device_inline
         eval *= bump_shadowing_term(sd->N, sc->N, omega_in);
       }
     }
-    /* Shadow terminator offset. */
-    const float frequency_multiplier =
-        kernel_tex_fetch(__objects, sd->object).shadow_terminator_offset;
-    if (frequency_multiplier > 1.0f) {
-      eval *= shift_cos_in(dot(omega_in, sc->N), frequency_multiplier);
-    }
   }
   else {
     switch (sc->type) {
@@ -613,6 +595,8 @@ ccl_device_inline
       case CLOSURE_BSDF_MICROFACET_GGX_ID:
       case CLOSURE_BSDF_MICROFACET_GGX_FRESNEL_ID:
       case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
+      case CLOSURE_BSDF_MICROFACET_GGX_ANISO_ID:
+      case CLOSURE_BSDF_MICROFACET_GGX_ANISO_FRESNEL_ID:
       case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
         eval = bsdf_microfacet_ggx_eval_transmit(sc, sd->I, omega_in, pdf);
         break;
@@ -626,10 +610,12 @@ ccl_device_inline
             sc, sd->I, omega_in, pdf, &sd->lcg_state);
         break;
       case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
+      case CLOSURE_BSDF_MICROFACET_BECKMANN_ANISO_ID:
       case CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID:
         eval = bsdf_microfacet_beckmann_eval_transmit(sc, sd->I, omega_in, pdf);
         break;
       case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
+      case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ANISO_ID:
         eval = bsdf_ashikhmin_shirley_eval_transmit(sc, sd->I, omega_in, pdf);
         break;
       case CLOSURE_BSDF_ASHIKHMIN_VELVET_ID:
@@ -693,14 +679,18 @@ ccl_device void bsdf_blur(KernelGlobals *kg, ShaderClosure *sc, float roughness)
     case CLOSURE_BSDF_MICROFACET_GGX_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_FRESNEL_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
+    case CLOSURE_BSDF_MICROFACET_GGX_ANISO_ID:
+    case CLOSURE_BSDF_MICROFACET_GGX_ANISO_FRESNEL_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
       bsdf_microfacet_ggx_blur(sc, roughness);
       break;
     case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
+    case CLOSURE_BSDF_MICROFACET_BECKMANN_ANISO_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID:
       bsdf_microfacet_beckmann_blur(sc, roughness);
       break;
     case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
+    case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ANISO_ID:
       bsdf_ashikhmin_shirley_blur(sc, roughness);
       break;
     case CLOSURE_BSDF_HAIR_PRINCIPLED_ID:
@@ -729,14 +719,18 @@ ccl_device bool bsdf_merge(ShaderClosure *a, ShaderClosure *b)
     case CLOSURE_BSDF_MICROFACET_GGX_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_FRESNEL_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
+    case CLOSURE_BSDF_MICROFACET_GGX_ANISO_ID:
+    case CLOSURE_BSDF_MICROFACET_GGX_ANISO_FRESNEL_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_ID:
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_FRESNEL_ID:
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_ID:
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_FRESNEL_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
+    case CLOSURE_BSDF_MICROFACET_BECKMANN_ANISO_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID:
     case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
+    case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ANISO_ID:
       return bsdf_microfacet_merge(a, b);
     case CLOSURE_BSDF_ASHIKHMIN_VELVET_ID:
       return bsdf_ashikhmin_velvet_merge(a, b);

intern/cycles/kernel/closure/bsdf_ashikhmin_shirley.h

@@ -34,12 +34,21 @@ CCL_NAMESPACE_BEGIN
 ccl_device int bsdf_ashikhmin_shirley_setup(MicrofacetBsdf *bsdf)
 {
   bsdf->alpha_x = clamp(bsdf->alpha_x, 1e-4f, 1.0f);
-  bsdf->alpha_y = clamp(bsdf->alpha_y, 1e-4f, 1.0f);
+  bsdf->alpha_y = bsdf->alpha_x;
 
   bsdf->type = CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID;
   return SD_BSDF | SD_BSDF_HAS_EVAL;
 }
 
+ccl_device int bsdf_ashikhmin_shirley_aniso_setup(MicrofacetBsdf *bsdf)
+{
+  bsdf->alpha_x = clamp(bsdf->alpha_x, 1e-4f, 1.0f);
+  bsdf->alpha_y = clamp(bsdf->alpha_y, 1e-4f, 1.0f);
+
+  bsdf->type = CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ANISO_ID;
+  return SD_BSDF | SD_BSDF_HAS_EVAL;
+}
+
 ccl_device void bsdf_ashikhmin_shirley_blur(ShaderClosure *sc, float roughness)
 {
   MicrofacetBsdf *bsdf = (MicrofacetBsdf *)sc;

intern/cycles/kernel/closure/bsdf_microfacet.h

@@ -256,7 +256,9 @@ ccl_device_forceinline float3 reflection_color(const MicrofacetBsdf *bsdf, float
 {
   float3 F = make_float3(1.0f, 1.0f, 1.0f);
   bool use_fresnel = (bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_FRESNEL_ID ||
-                      bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID);
+                      bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID ||
+                      bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_ANISO_FRESNEL_ID);
+
   if (use_fresnel) {
     float F0 = fresnel_dielectric_cos(1.0f, bsdf->ior);
 
@@ -309,27 +311,19 @@ ccl_device int bsdf_microfacet_ggx_setup(MicrofacetBsdf *bsdf)
   bsdf->extra = NULL;
 
   bsdf->alpha_x = saturate(bsdf->alpha_x);
-  bsdf->alpha_y = saturate(bsdf->alpha_y);
+  bsdf->alpha_y = bsdf->alpha_x;
 
   bsdf->type = CLOSURE_BSDF_MICROFACET_GGX_ID;
 
   return SD_BSDF | SD_BSDF_HAS_EVAL;
 }
 
-/* Required to maintain OSL interface. */
-ccl_device int bsdf_microfacet_ggx_isotropic_setup(MicrofacetBsdf *bsdf)
-{
-  bsdf->alpha_y = bsdf->alpha_x;
-
-  return bsdf_microfacet_ggx_setup(bsdf);
-}
-
 ccl_device int bsdf_microfacet_ggx_fresnel_setup(MicrofacetBsdf *bsdf, const ShaderData *sd)
 {
   bsdf->extra->cspec0 = saturate3(bsdf->extra->cspec0);
 
   bsdf->alpha_x = saturate(bsdf->alpha_x);
-  bsdf->alpha_y = saturate(bsdf->alpha_y);
+  bsdf->alpha_y = bsdf->alpha_x;
 
   bsdf->type = CLOSURE_BSDF_MICROFACET_GGX_FRESNEL_ID;
 
@@ -367,6 +361,32 @@ ccl_device bool bsdf_microfacet_merge(const ShaderClosure *a, const ShaderClosur
            (bsdf_a->extra->clearcoat == bsdf_b->extra->clearcoat)));
 }
 
+ccl_device int bsdf_microfacet_ggx_aniso_setup(MicrofacetBsdf *bsdf)
+{
+  bsdf->extra = NULL;
+
+  bsdf->alpha_x = saturate(bsdf->alpha_x);
+  bsdf->alpha_y = saturate(bsdf->alpha_y);
+
+  bsdf->type = CLOSURE_BSDF_MICROFACET_GGX_ANISO_ID;
+
+  return SD_BSDF | SD_BSDF_HAS_EVAL;
+}
+
+ccl_device int bsdf_microfacet_ggx_aniso_fresnel_setup(MicrofacetBsdf *bsdf, const ShaderData *sd)
+{
+  bsdf->extra->cspec0 = saturate3(bsdf->extra->cspec0);
+
+  bsdf->alpha_x = saturate(bsdf->alpha_x);
+  bsdf->alpha_y = saturate(bsdf->alpha_y);
+
+  bsdf->type = CLOSURE_BSDF_MICROFACET_GGX_ANISO_FRESNEL_ID;
+
+  bsdf_microfacet_fresnel_color(sd, bsdf);
+
+  return SD_BSDF | SD_BSDF_HAS_EVAL;
+}
+
 ccl_device int bsdf_microfacet_ggx_refraction_setup(MicrofacetBsdf *bsdf)
 {
   bsdf->extra = NULL;
@@ -616,7 +636,8 @@ ccl_device int bsdf_microfacet_ggx_sample(KernelGlobals *kg,
             *eval = make_float3(1e6f, 1e6f, 1e6f);
 
             bool use_fresnel = (bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_FRESNEL_ID ||
-                                bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID);
+                                bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID ||
+                                bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_ANISO_FRESNEL_ID);
 
             /* if fresnel is used, calculate the color with reflection_color(...) */
             if (use_fresnel) {
@@ -790,18 +811,19 @@ ccl_device int bsdf_microfacet_ggx_sample(KernelGlobals *kg,
 ccl_device int bsdf_microfacet_beckmann_setup(MicrofacetBsdf *bsdf)
 {
   bsdf->alpha_x = saturate(bsdf->alpha_x);
-  bsdf->alpha_y = saturate(bsdf->alpha_y);
+  bsdf->alpha_y = bsdf->alpha_x;
 
   bsdf->type = CLOSURE_BSDF_MICROFACET_BECKMANN_ID;
   return SD_BSDF | SD_BSDF_HAS_EVAL;
 }
 
-/* Required to maintain OSL interface. */
+ccl_device int bsdf_microfacet_beckmann_aniso_setup(MicrofacetBsdf *bsdf)
-ccl_device int bsdf_microfacet_beckmann_isotropic_setup(MicrofacetBsdf *bsdf)
 {
-  bsdf->alpha_y = bsdf->alpha_x;
+  bsdf->alpha_x = saturate(bsdf->alpha_x);
+  bsdf->alpha_y = saturate(bsdf->alpha_y);
 
-  return bsdf_microfacet_beckmann_setup(bsdf);
+  bsdf->type = CLOSURE_BSDF_MICROFACET_BECKMANN_ANISO_ID;
+  return SD_BSDF | SD_BSDF_HAS_EVAL;
 }
 
 ccl_device int bsdf_microfacet_beckmann_refraction_setup(MicrofacetBsdf *bsdf)

intern/cycles/kernel/closure/bsdf_microfacet_multi.h

@@ -384,7 +384,7 @@ ccl_device int bsdf_microfacet_multi_ggx_common_setup(MicrofacetBsdf *bsdf)
   return SD_BSDF | SD_BSDF_HAS_EVAL | SD_BSDF_NEEDS_LCG;
 }
 
-ccl_device int bsdf_microfacet_multi_ggx_setup(MicrofacetBsdf *bsdf)
+ccl_device int bsdf_microfacet_multi_ggx_aniso_setup(MicrofacetBsdf *bsdf)
 {
   if (is_zero(bsdf->T))
     bsdf->T = make_float3(1.0f, 0.0f, 0.0f);
@@ -394,7 +394,8 @@ ccl_device int bsdf_microfacet_multi_ggx_setup(MicrofacetBsdf *bsdf)
   return bsdf_microfacet_multi_ggx_common_setup(bsdf);
 }
 
-ccl_device int bsdf_microfacet_multi_ggx_fresnel_setup(MicrofacetBsdf *bsdf, const ShaderData *sd)
+ccl_device int bsdf_microfacet_multi_ggx_aniso_fresnel_setup(MicrofacetBsdf *bsdf,
+                                                             const ShaderData *sd)
 {
   if (is_zero(bsdf->T))
     bsdf->T = make_float3(1.0f, 0.0f, 0.0f);
@@ -406,6 +407,26 @@ ccl_device int bsdf_microfacet_multi_ggx_fresnel_setup(MicrofacetBsdf *bsdf, con
   return bsdf_microfacet_multi_ggx_common_setup(bsdf);
 }
 
+ccl_device int bsdf_microfacet_multi_ggx_setup(MicrofacetBsdf *bsdf)
+{
+  bsdf->alpha_y = bsdf->alpha_x;
+
+  bsdf->type = CLOSURE_BSDF_MICROFACET_MULTI_GGX_ID;
+
+  return bsdf_microfacet_multi_ggx_common_setup(bsdf);
+}
+
+ccl_device int bsdf_microfacet_multi_ggx_fresnel_setup(MicrofacetBsdf *bsdf, const ShaderData *sd)
+{
+  bsdf->alpha_y = bsdf->alpha_x;
+
+  bsdf->type = CLOSURE_BSDF_MICROFACET_MULTI_GGX_FRESNEL_ID;
+
+  bsdf_microfacet_fresnel_color(sd, bsdf);
+
+  return bsdf_microfacet_multi_ggx_common_setup(bsdf);
+}
+
 ccl_device int bsdf_microfacet_multi_ggx_refraction_setup(MicrofacetBsdf *bsdf)
 {
   bsdf->alpha_y = bsdf->alpha_x;

intern/cycles/kernel/kernel_adaptive_sampling.h

@@ -185,19 +185,19 @@ ccl_device bool kernel_do_adaptive_filter_x(KernelGlobals *kg, int y, ccl_global
     ccl_global float *buffer = tile->buffer + index * kernel_data.film.pass_stride;
     ccl_global float4 *aux = (ccl_global float4 *)(buffer +
                                                    kernel_data.film.pass_adaptive_aux_buffer);
-    if ((*aux).w == 0.0f) {
+    if (aux->w == 0.0f) {
       any = true;
       if (x > tile->x && !prev) {
         index = index - 1;
         buffer = tile->buffer + index * kernel_data.film.pass_stride;
         aux = (ccl_global float4 *)(buffer + kernel_data.film.pass_adaptive_aux_buffer);
-        (*aux).w = 0.0f;
+        aux->w = 0.0f;
       }
       prev = true;
     }
     else {
       if (prev) {
-        (*aux).w = 0.0f;
+        aux->w = 0.0f;
       }
       prev = false;
     }
@@ -214,19 +214,19 @@ ccl_device bool kernel_do_adaptive_filter_y(KernelGlobals *kg, int x, ccl_global
     ccl_global float *buffer = tile->buffer + index * kernel_data.film.pass_stride;
     ccl_global float4 *aux = (ccl_global float4 *)(buffer +
                                                    kernel_data.film.pass_adaptive_aux_buffer);
-    if ((*aux).w == 0.0f) {
+    if (aux->w == 0.0f) {
       any = true;
       if (y > tile->y && !prev) {
         index = index - tile->stride;
         buffer = tile->buffer + index * kernel_data.film.pass_stride;
         aux = (ccl_global float4 *)(buffer + kernel_data.film.pass_adaptive_aux_buffer);
-        (*aux).w = 0.0f;
+        aux->w = 0.0f;
       }
       prev = true;
     }
     else {
       if (prev) {
-        (*aux).w = 0.0f;
+        aux->w = 0.0f;
       }
       prev = false;
     }

intern/cycles/kernel/kernel_bake.h

@@ -18,33 +18,19 @@ CCL_NAMESPACE_BEGIN
 
 #ifdef __BAKING__
 
-ccl_device_noinline void compute_light_pass(
+ccl_device_inline void compute_light_pass(
     KernelGlobals *kg, ShaderData *sd, PathRadiance *L, uint rng_hash, int pass_filter, int sample)
 {
   kernel_assert(kernel_data.film.use_light_pass);
 
+  PathRadiance L_sample;
+  PathState state;
+  Ray ray;
   float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
 
-  /* Emission and indirect shader data memory used by various functions. */
+  /* emission and indirect shader data memory used by various functions */
-  ShaderDataTinyStorage emission_sd_storage;
+  ShaderData emission_sd, indirect_sd;
-  ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
-  ShaderData indirect_sd;
 
-  /* Init radiance. */
-  path_radiance_init(kg, L);
-
-  /* Init path state. */
-  PathState state;
-  path_state_init(kg, emission_sd, &state, rng_hash, sample, NULL);
-
-  /* Evaluate surface shader. */
-  shader_eval_surface(kg, sd, &state, NULL, state.flag);
-
-  /* TODO, disable more closures we don't need besides transparent */
-  shader_bsdf_disable_transparency(kg, sd);
-
-  /* Init ray. */
-  Ray ray;
   ray.P = sd->P + sd->Ng;
   ray.D = -sd->Ng;
   ray.t = FLT_MAX;
@@ -52,6 +38,18 @@ ccl_device_noinline void compute_light_pass(
   ray.time = 0.5f;
 #  endif
 
+  /* init radiance */
+  path_radiance_init(kg, &L_sample);
+
+  /* init path state */
+  path_state_init(kg, &emission_sd, &state, rng_hash, sample, NULL);
+
+  /* evaluate surface shader */
+  shader_eval_surface(kg, sd, &state, NULL, state.flag);
+
+  /* TODO, disable more closures we don't need besides transparent */
+  shader_bsdf_disable_transparency(kg, sd);
+
 #  ifdef __BRANCHED_PATH__
   if (!kernel_data.integrator.branched) {
     /* regular path tracer */
@@ -59,13 +57,14 @@ ccl_device_noinline void compute_light_pass(
 
     /* sample ambient occlusion */
     if (pass_filter & BAKE_FILTER_AO) {
-      kernel_path_ao(kg, sd, emission_sd, L, &state, throughput, shader_bsdf_alpha(kg, sd));
+      kernel_path_ao(
+          kg, sd, &emission_sd, &L_sample, &state, throughput, shader_bsdf_alpha(kg, sd));
     }
 
     /* sample emission */
     if ((pass_filter & BAKE_FILTER_EMISSION) && (sd->flag & SD_EMISSION)) {
       float3 emission = indirect_primitive_emission(kg, sd, 0.0f, state.flag, state.ray_pdf);
-      path_radiance_accum_emission(kg, L, &state, throughput, emission);
+      path_radiance_accum_emission(kg, &L_sample, &state, throughput, emission);
     }
 
     bool is_sss_sample = false;
@@ -78,10 +77,12 @@ ccl_device_noinline void compute_light_pass(
       SubsurfaceIndirectRays ss_indirect;
       kernel_path_subsurface_init_indirect(&ss_indirect);
       if (kernel_path_subsurface_scatter(
-              kg, sd, emission_sd, L, &state, &ray, &throughput, &ss_indirect)) {
+              kg, sd, &emission_sd, &L_sample, &state, &ray, &throughput, &ss_indirect)) {
         while (ss_indirect.num_rays) {
-          kernel_path_subsurface_setup_indirect(kg, &ss_indirect, &state, &ray, L, &throughput);
+          kernel_path_subsurface_setup_indirect(
-          kernel_path_indirect(kg, &indirect_sd, emission_sd, &ray, throughput, &state, L);
+              kg, &ss_indirect, &state, &ray, &L_sample, &throughput);
+          kernel_path_indirect(
+              kg, &indirect_sd, &emission_sd, &ray, throughput, &state, &L_sample);
         }
         is_sss_sample = true;
       }
@@ -90,18 +91,18 @@ ccl_device_noinline void compute_light_pass(
 
     /* sample light and BSDF */
     if (!is_sss_sample && (pass_filter & (BAKE_FILTER_DIRECT | BAKE_FILTER_INDIRECT))) {
-      kernel_path_surface_connect_light(kg, sd, emission_sd, throughput, &state, L);
+      kernel_path_surface_connect_light(kg, sd, &emission_sd, throughput, &state, &L_sample);
 
-      if (kernel_path_surface_bounce(kg, sd, &throughput, &state, &L->state, &ray)) {
+      if (kernel_path_surface_bounce(kg, sd, &throughput, &state, &L_sample.state, &ray)) {
 #  ifdef __LAMP_MIS__
         state.ray_t = 0.0f;
 #  endif
         /* compute indirect light */
-        kernel_path_indirect(kg, &indirect_sd, emission_sd, &ray, throughput, &state, L);
+        kernel_path_indirect(kg, &indirect_sd, &emission_sd, &ray, throughput, &state, &L_sample);
 
         /* sum and reset indirect light pass variables for the next samples */
-        path_radiance_sum_indirect(L);
+        path_radiance_sum_indirect(&L_sample);
-        path_radiance_reset_indirect(L);
+        path_radiance_reset_indirect(&L_sample);
       }
     }
 #  ifdef __BRANCHED_PATH__
@@ -111,13 +112,13 @@ ccl_device_noinline void compute_light_pass(
 
     /* sample ambient occlusion */
     if (pass_filter & BAKE_FILTER_AO) {
-      kernel_branched_path_ao(kg, sd, emission_sd, L, &state, throughput);
+      kernel_branched_path_ao(kg, sd, &emission_sd, &L_sample, &state, throughput);
     }
 
     /* sample emission */
     if ((pass_filter & BAKE_FILTER_EMISSION) && (sd->flag & SD_EMISSION)) {
       float3 emission = indirect_primitive_emission(kg, sd, 0.0f, state.flag, state.ray_pdf);
-      path_radiance_accum_emission(kg, L, &state, throughput, emission);
+      path_radiance_accum_emission(kg, &L_sample, &state, throughput, emission);
     }
 
 #    ifdef __SUBSURFACE__
@@ -126,7 +127,7 @@ ccl_device_noinline void compute_light_pass(
       /* When mixing BSSRDF and BSDF closures we should skip BSDF lighting
        * if scattering was successful. */
       kernel_branched_path_subsurface_scatter(
-          kg, sd, &indirect_sd, emission_sd, L, &state, &ray, throughput);
+          kg, sd, &indirect_sd, &emission_sd, &L_sample, &state, &ray, throughput);
     }
 #    endif
 
@@ -137,16 +138,19 @@ ccl_device_noinline void compute_light_pass(
       if (kernel_data.integrator.use_direct_light) {
         int all = kernel_data.integrator.sample_all_lights_direct;
         kernel_branched_path_surface_connect_light(
-            kg, sd, emission_sd, &state, throughput, 1.0f, L, all);
+            kg, sd, &emission_sd, &state, throughput, 1.0f, &L_sample, all);
       }
 #    endif
 
       /* indirect light */
       kernel_branched_path_surface_indirect_light(
-          kg, sd, &indirect_sd, emission_sd, throughput, 1.0f, &state, L);
+          kg, sd, &indirect_sd, &emission_sd, throughput, 1.0f, &state, &L_sample);
     }
   }
 #  endif
+
+  /* accumulate into master L */
+  path_radiance_accum_sample(L, &L_sample);
 }
 
 /* this helps with AA but it's not the real solution as it does not AA the geometry
@@ -221,28 +225,41 @@ ccl_device float3 kernel_bake_evaluate_direct_indirect(KernelGlobals *kg,
   return out;
 }
 
-ccl_device void kernel_bake_evaluate(
+ccl_device void kernel_bake_evaluate(KernelGlobals *kg,
-    KernelGlobals *kg, ccl_global float *buffer, int sample, int x, int y, int offset, int stride)
+                                     ccl_global uint4 *input,
+                                     ccl_global float4 *output,
+                                     ShaderEvalType type,
+                                     int pass_filter,
+                                     int i,
+                                     int offset,
+                                     int sample)
 {
-  /* Setup render buffers. */
+  ShaderData sd;
-  const int index = offset + x + y * stride;
+  PathState state = {0};
-  const int pass_stride = kernel_data.film.pass_stride;
+  uint4 in = input[i * 2];
-  buffer += index * pass_stride;
+  uint4 diff = input[i * 2 + 1];
 
-  ccl_global float *primitive = buffer + kernel_data.film.pass_bake_primitive;
+  float3 out = make_float3(0.0f, 0.0f, 0.0f);
-  ccl_global float *differential = buffer + kernel_data.film.pass_bake_differential;
+
-  ccl_global float *output = buffer + kernel_data.film.pass_combined;
+  int object = in.x;
+  int prim = in.y;
 
-  int prim = __float_as_uint(primitive[1]);
   if (prim == -1)
     return;
 
-  prim += kernel_data.bake.tri_offset;
+  float u = __uint_as_float(in.z);
+  float v = __uint_as_float(in.w);
+
+  float dudx = __uint_as_float(diff.x);
+  float dudy = __uint_as_float(diff.y);
+  float dvdx = __uint_as_float(diff.z);
+  float dvdy = __uint_as_float(diff.w);
 
-  /* Random number generator. */
-  uint rng_hash = hash_uint2(x, y) ^ kernel_data.integrator.seed;
   int num_samples = kernel_data.integrator.aa_samples;
 
+  /* random number generator */
+  uint rng_hash = cmj_hash(offset + i, kernel_data.integrator.seed);
+
   float filter_x, filter_y;
   if (sample == 0) {
     filter_x = filter_y = 0.5f;
@@ -251,29 +268,23 @@ ccl_device void kernel_bake_evaluate(
     path_rng_2D(kg, rng_hash, sample, num_samples, PRNG_FILTER_U, &filter_x, &filter_y);
   }
 
-  /* Barycentric UV with subpixel offset. */
+  /* subpixel u/v offset */
-  float u = primitive[2];
-  float v = primitive[3];
-
-  float dudx = differential[0];
-  float dudy = differential[1];
-  float dvdx = differential[2];
-  float dvdy = differential[3];
-
   if (sample > 0) {
     u = bake_clamp_mirror_repeat(u + dudx * (filter_x - 0.5f) + dudy * (filter_y - 0.5f), 1.0f);
     v = bake_clamp_mirror_repeat(v + dvdx * (filter_x - 0.5f) + dvdy * (filter_y - 0.5f),
                                  1.0f - u);
   }
 
-  /* Shader data setup. */
+  /* triangle */
-  int object = kernel_data.bake.object_index;
   int shader;
   float3 P, Ng;
 
   triangle_point_normal(kg, object, prim, u, v, &P, &Ng, &shader);
 
-  ShaderData sd;
+  /* light passes */
+  PathRadiance L;
+  path_radiance_init(kg, &L);
+
   shader_setup_from_sample(
       kg,
       &sd,
@@ -291,7 +302,7 @@ ccl_device void kernel_bake_evaluate(
       LAMP_NONE);
   sd.I = sd.N;
 
-  /* Setup differentials. */
+  /* update differentials */
   sd.dP.dx = sd.dPdu * dudx + sd.dPdv * dvdx;
   sd.dP.dy = sd.dPdu * dudy + sd.dPdv * dvdy;
   sd.du.dx = dudx;
@@ -299,24 +310,17 @@ ccl_device void kernel_bake_evaluate(
   sd.dv.dx = dvdx;
   sd.dv.dy = dvdy;
 
-  /* Set RNG state for shaders that use sampling. */
+  /* set RNG state for shaders that use sampling */
-  PathState state = {0};
   state.rng_hash = rng_hash;
   state.rng_offset = 0;
   state.sample = sample;
   state.num_samples = num_samples;
   state.min_ray_pdf = FLT_MAX;
 
-  /* Light passes if we need more than color. */
+  /* light passes if we need more than color */
-  PathRadiance L;
+  if (pass_filter & ~BAKE_FILTER_COLOR)
-  int pass_filter = kernel_data.bake.pass_filter;
-
-  if (kernel_data.bake.pass_filter & ~BAKE_FILTER_COLOR)
     compute_light_pass(kg, &sd, &L, rng_hash, pass_filter, sample);
 
-  float3 out = make_float3(0.0f, 0.0f, 0.0f);
-
-  ShaderEvalType type = (ShaderEvalType)kernel_data.bake.type;
   switch (type) {
     /* data passes */
     case SHADER_EVAL_NORMAL:
@@ -437,8 +441,10 @@ ccl_device void kernel_bake_evaluate(
   }
 
   /* write output */
-  const float4 result = make_float4(out.x, out.y, out.z, 1.0f);
+  const float output_fac = 1.0f / num_samples;
-  kernel_write_pass_float4(output, result);
+  const float4 scaled_result = make_float4(out.x, out.y, out.z, 1.0f) * output_fac;
+
+  output[i] = (sample == 0) ? scaled_result : output[i] + scaled_result;
 }
 
 #endif /* __BAKING__ */

intern/cycles/kernel/kernel_camera.h

@@ -441,22 +441,8 @@ ccl_device_inline float camera_distance(KernelGlobals *kg, float3 P)
     float3 camD = make_float3(cameratoworld.x.z, cameratoworld.y.z, cameratoworld.z.z);
     return fabsf(dot((P - camP), camD));
   }
-  else {
+  else
     return len(P - camP);
-  }
-}
-
-ccl_device_inline float camera_z_depth(KernelGlobals *kg, float3 P)
-{
-  if (kernel_data.cam.type != CAMERA_PANORAMA) {
-    Transform worldtocamera = kernel_data.cam.worldtocamera;
-    return transform_point(&worldtocamera, P).z;
-  }
-  else {
-    Transform cameratoworld = kernel_data.cam.cameratoworld;
-    float3 camP = make_float3(cameratoworld.x.w, cameratoworld.y.w, cameratoworld.z.w);
-    return len(P - camP);
-  }
 }
 
 ccl_device_inline float3 camera_direction_from_point(KernelGlobals *kg, float3 P)

intern/cycles/kernel/kernel_film.h

@@ -28,13 +28,13 @@ ccl_device float4 film_get_pass_result(KernelGlobals *kg,
   int display_pass_components = kernel_data.film.display_pass_components;
 
   if (display_pass_components == 4) {
-    float4 in = *(ccl_global float4 *)(buffer + display_pass_stride +
+    ccl_global float4 *in = (ccl_global float4 *)(buffer + display_pass_stride +
-                                       index * kernel_data.film.pass_stride);
+                                                  index * kernel_data.film.pass_stride);
     float alpha = use_display_sample_scale ?
-                      (kernel_data.film.use_display_pass_alpha ? in.w : 1.0f / sample_scale) :
+                      (kernel_data.film.use_display_pass_alpha ? in->w : 1.0f / sample_scale) :
                       1.0f;
 
-    pass_result = make_float4(in.x, in.y, in.z, alpha);
+    pass_result = make_float4(in->x, in->y, in->z, alpha);
 
     int display_divide_pass_stride = kernel_data.film.display_divide_pass_stride;
     if (display_divide_pass_stride != -1) {

intern/cycles/kernel/kernel_light.h

@@ -249,13 +249,13 @@ ccl_device float3 background_map_sample(KernelGlobals *kg, float randu, float ra
   float u = (index_u + du) / res_x;
 
   /* compute pdf */
+  float denom = cdf_last_u.x * cdf_last_v.x;
   float sin_theta = sinf(M_PI_F * v);
-  float denom = (M_2PI_F * M_PI_F * sin_theta) * cdf_last_u.x * cdf_last_v.x;
 
   if (sin_theta == 0.0f || denom == 0.0f)
     *pdf = 0.0f;
   else
-    *pdf = (cdf_u.x * cdf_v.x) / denom;
+    *pdf = (cdf_u.x * cdf_v.x) / (M_2PI_F * M_PI_F * sin_theta * denom);
 
   /* compute direction */
   return equirectangular_to_direction(u, v);
@@ -284,7 +284,7 @@ ccl_device float background_map_pdf(KernelGlobals *kg, float3 direction)
                                        index_v * cdf_width + res_x);
   float2 cdf_last_v = kernel_tex_fetch(__light_background_marginal_cdf, res_y);
 
-  float denom = (M_2PI_F * M_PI_F * sin_theta) * cdf_last_u.x * cdf_last_v.x;
+  float denom = cdf_last_u.x * cdf_last_v.x;
 
   if (denom == 0.0f)
     return 0.0f;
@@ -294,7 +294,7 @@ ccl_device float background_map_pdf(KernelGlobals *kg, float3 direction)
                                   index_v * cdf_width + index_u);
   float2 cdf_v = kernel_tex_fetch(__light_background_marginal_cdf, index_v);
 
-  return (cdf_u.x * cdf_v.x) / denom;
+  return (cdf_u.x * cdf_v.x) / (M_2PI_F * M_PI_F * sin_theta * denom);
 }
 
 ccl_device_inline bool background_portal_data_fetch_and_check_side(
@@ -1041,19 +1041,11 @@ ccl_device_forceinline void triangle_light_sample(KernelGlobals *kg,
     }
   }
   else {
-    /* compute random point in triangle. From Eric Heitz's "A Low-Distortion Map Between Triangle
+    /* compute random point in triangle */
-     * and Square" */
+    randu = sqrtf(randu);
-    float u = randu;
-    float v = randv;
-    if (v > u) {
-      u *= 0.5f;
-      v -= u;
-    }
-    else {
-      v *= 0.5f;
-      u -= v;
-    }
 
+    const float u = 1.0f - randu;
+    const float v = randv * randu;
     const float t = 1.0f - u - v;
     ls->P = u * V[0] + v * V[1] + t * V[2];
     /* compute incoming direction, distance and pdf */

intern/cycles/kernel/kernel_montecarlo.h

@@ -85,9 +85,8 @@ ccl_device_inline void sample_uniform_hemisphere(
 ccl_device_inline void sample_uniform_cone(
     const float3 N, float angle, float randu, float randv, float3 *omega_in, float *pdf)
 {
-  float zMin = cosf(angle);
+  float z = cosf(angle * randu);
-  float z = zMin - zMin * randu + randu;
+  float r = sqrtf(max(0.0f, 1.0f - z * z));
-  float r = safe_sqrtf(1.0f - sqr(z));
   float phi = M_2PI_F * randv;
   float x = r * cosf(phi);
   float y = r * sinf(phi);
@@ -95,7 +94,7 @@ ccl_device_inline void sample_uniform_cone(
   float3 T, B;
   make_orthonormals(N, &T, &B);
   *omega_in = x * T + y * B + z * N;
-  *pdf = M_1_2PI_F / (1.0f - zMin);
+  *pdf = 0.5f * M_1_PI_F / (1.0f - cosf(angle));
 }
 
 /* sample uniform point on the surface of a sphere */

intern/cycles/kernel/kernel_passes.h

@@ -194,7 +194,7 @@ ccl_device_inline void kernel_write_data_passes(KernelGlobals *kg,
         average(shader_bsdf_alpha(kg, sd)) >= kernel_data.film.pass_alpha_threshold) {
       if (state->sample == 0) {
         if (flag & PASSMASK(DEPTH)) {
-          float depth = camera_z_depth(kg, sd->P);
+          float depth = camera_distance(kg, sd->P);
           kernel_write_pass_float(buffer + kernel_data.film.pass_depth, depth);
         }
         if (flag & PASSMASK(OBJECT_ID)) {

intern/cycles/kernel/kernel_path.h

@@ -662,7 +662,7 @@ ccl_device void kernel_path_trace(
   if (kernel_data.film.pass_adaptive_aux_buffer) {
     ccl_global float4 *aux = (ccl_global float4 *)(buffer +
                                                    kernel_data.film.pass_adaptive_aux_buffer);
-    if ((*aux).w > 0.0f) {
+    if (aux->w > 0.0f) {
       return;
     }
   }

intern/cycles/kernel/kernel_path_branched.h

@@ -526,7 +526,7 @@ ccl_device void kernel_branched_path_trace(
   if (kernel_data.film.pass_adaptive_aux_buffer) {
     ccl_global float4 *aux = (ccl_global float4 *)(buffer +
                                                    kernel_data.film.pass_adaptive_aux_buffer);
-    if ((*aux).w > 0.0f) {
+    if (aux->w > 0.0f) {
       return;
     }
   }

intern/cycles/kernel/kernel_types.h

@@ -273,7 +273,6 @@ enum SamplingPattern {
 /* these flags values correspond to raytypes in osl.cpp, so keep them in sync! */
 
 enum PathRayFlag {
-  /* Ray visibility. */
   PATH_RAY_CAMERA = (1 << 0),
   PATH_RAY_REFLECT = (1 << 1),
   PATH_RAY_TRANSMIT = (1 << 2),
@@ -282,7 +281,6 @@ enum PathRayFlag {
   PATH_RAY_SINGULAR = (1 << 5),
   PATH_RAY_TRANSPARENT = (1 << 6),
 
-  /* Shadow ray visibility. */
   PATH_RAY_SHADOW_OPAQUE_NON_CATCHER = (1 << 7),
   PATH_RAY_SHADOW_OPAQUE_CATCHER = (1 << 8),
   PATH_RAY_SHADOW_OPAQUE = (PATH_RAY_SHADOW_OPAQUE_NON_CATCHER | PATH_RAY_SHADOW_OPAQUE_CATCHER),
@@ -294,11 +292,8 @@ enum PathRayFlag {
                                  PATH_RAY_SHADOW_TRANSPARENT_NON_CATCHER),
   PATH_RAY_SHADOW = (PATH_RAY_SHADOW_OPAQUE | PATH_RAY_SHADOW_TRANSPARENT),
 
-  /* Unused, free to reuse. */
+  PATH_RAY_CURVE = (1 << 11),          /* visibility flag to define curve segments */
-  PATH_RAY_UNUSED = (1 << 11),
+  PATH_RAY_VOLUME_SCATTER = (1 << 12), /* volume scattering */
-
-  /* Ray visibility for volume scattering. */
-  PATH_RAY_VOLUME_SCATTER = (1 << 12),
 
   /* Special flag to tag unaligned BVH nodes. */
   PATH_RAY_NODE_UNALIGNED = (1 << 13),
@@ -400,10 +395,6 @@ typedef enum PassType {
   PASS_VOLUME_INDIRECT,
   /* No Scatter color since it's tricky to define what it would even mean. */
   PASS_CATEGORY_LIGHT_END = 63,
-
-  PASS_BAKE_PRIMITIVE,
-  PASS_BAKE_DIFFERENTIAL,
-  PASS_CATEGORY_BAKE_END = 95
 } PassType;
 
 #define PASS_ANY (~0)
@@ -1257,10 +1248,6 @@ typedef struct KernelFilm {
   float4 xyz_to_b;
   float4 rgb_to_y;
 
-  int pass_bake_primitive;
-  int pass_bake_differential;
-  int pad;
-
 #ifdef __KERNEL_DEBUG__
   int pass_bvh_traversed_nodes;
   int pass_bvh_traversed_instances;
@@ -1440,14 +1427,6 @@ typedef struct KernelTables {
 } KernelTables;
 static_assert_align(KernelTables, 16);
 
-typedef struct KernelBake {
-  int object_index;
-  int tri_offset;
-  int type;
-  int pass_filter;
-} KernelBake;
-static_assert_align(KernelBake, 16);
-
 typedef struct KernelData {
   KernelCamera cam;
   KernelFilm film;
@@ -1456,7 +1435,6 @@ typedef struct KernelData {
   KernelBVH bvh;
   KernelCurves curve;
   KernelTables tables;
-  KernelBake bake;
 } KernelData;
 static_assert_align(KernelData, 16);
 
@@ -1485,9 +1463,6 @@ typedef struct KernelObject {
 
   float cryptomatte_object;
   float cryptomatte_asset;
-
-  float shadow_terminator_offset;
-  float pad1, pad2, pad3;
 } KernelObject;
 static_assert_align(KernelObject, 16);