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

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

Thanks to @PratikPB2123 for investigation.

.clang-format

@@ -264,9 +264,6 @@ ForEachMacros:
   - SET_SLOT_PROBING_BEGIN
   - MAP_SLOT_PROBING_BEGIN
   - VECTOR_SET_SLOT_PROBING_BEGIN
-  - LIGHT_FOREACH_BEGIN_DIRECTIONAL
-  - LIGHT_FOREACH_BEGIN_LOCAL
-  - LIGHT_FOREACH_BEGIN_LOCAL_NO_CULL
 
 StatementMacros:
   - PyObject_HEAD

CMakeLists.txt

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

build_files/build_environment/cmake/freetype.cmake

@@ -19,13 +19,10 @@
 set(FREETYPE_EXTRA_ARGS
   -DCMAKE_RELEASE_POSTFIX:STRING=2ST
   -DCMAKE_DEBUG_POSTFIX:STRING=2ST_d
-  -DWITH_BZip2=OFF
-  -DWITH_HarfBuzz=OFF
-  -DFT_WITH_HARFBUZZ=OFF
-  -DFT_WITH_BZIP2=OFF
-  -DFT_WITH_BROTLI=ON
-  -DCMAKE_DISABLE_FIND_PACKAGE_HarfBuzz=TRUE
-  -DCMAKE_DISABLE_FIND_PACKAGE_BZip2=TRUE
+  -DFT_DISABLE_BZIP2=ON
+  -DFT_DISABLE_HARFBUZZ=ON
+  -DFT_DISABLE_PNG=ON
+  -DFT_REQUIRE_BROTLI=ON
   -DPC_BROTLIDEC_INCLUDEDIR=${LIBDIR}/brotli/include
   -DPC_BROTLIDEC_LIBDIR=${LIBDIR}/brotli/lib
   )

build_files/build_environment/cmake/versions.cmake

@@ -83,9 +83,9 @@ else()
   set(OPENEXR_VERSION_POSTFIX)
 endif()
 
-set(FREETYPE_VERSION 2.11.0)
+set(FREETYPE_VERSION 2.11.1)
 set(FREETYPE_URI http://prdownloads.sourceforge.net/freetype/freetype-${FREETYPE_VERSION}.tar.gz)
-set(FREETYPE_HASH cf09172322f6b50cf8f568bf8fe14bde)
+set(FREETYPE_HASH bd4e3b007474319909a6b79d50908e85)
 set(FREETYPE_HASH_TYPE MD5)
 set(FREETYPE_FILE freetype-${FREETYPE_VERSION}.tar.gz)
 

build_files/cmake/Modules/FindFFmpeg.cmake

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

build_files/cmake/Modules/FindOSL.cmake

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

build_files/cmake/Modules/FindOpenColorIO.cmake

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

build_files/cmake/Modules/FindOpenEXR.cmake

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

build_files/cmake/Modules/FindOpenImageDenoise.cmake

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

build_files/cmake/Modules/FindOpenImageIO.cmake

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

build_files/cmake/Modules/FindPythonLibsUnix.cmake

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

build_files/cmake/macros.cmake

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

build_files/cmake/platform/platform_unix.cmake

@@ -112,7 +112,7 @@ if(NOT WITH_SYSTEM_FREETYPE)
 
     # NOTE: This is done on WIN32 & APPLE but fails on some Linux systems.
     # See: https://devtalk.blender.org/t/22536
-    # So `BROTLI_LIBRARIES` need to be added  `FREETYPE_LIBRARIES`.
+    # So `BROTLI_LIBRARIES` need to be added directly after `FREETYPE_LIBRARIES`.
     #
     # list(APPEND FREETYPE_LIBRARIES
     #   ${BROTLI_LIBRARIES}
@@ -362,6 +362,7 @@ if(WITH_BOOST)
       find_package(IcuLinux)
     endif()
     mark_as_advanced(Boost_DIR)  # why doesn't boost do this?
+    mark_as_advanced(Boost_INCLUDE_DIR)  # why doesn't boost do this?
   endif()
 
   set(BOOST_INCLUDE_DIR ${Boost_INCLUDE_DIRS})

build_files/config/pipeline_config.yaml

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

doc/python_api/rst/info_api_reference.rst

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

doc/python_api/rst/info_best_practice.rst

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

doc/python_api/rst/info_gotcha.rst

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

doc/python_api/rst/info_tips_and_tricks.rst

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

doc/python_api/sphinx_doc_gen.py

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

doc/python_api/static/css/version_switch.css

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

doc/python_api/static/js/version_switch.js

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

doc/python_api/templates/versions.html

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

extern/CMakeLists.txt

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

extern/audaspace/plugins/ffmpeg/FFMPEGReader.cpp

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

extern/audaspace/plugins/ffmpeg/FFMPEGWriter.cpp

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

intern/cycles/app/CMakeLists.txt

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

intern/cycles/app/cycles_standalone.cpp

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

intern/cycles/app/opengl/display_driver.cpp

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

intern/cycles/app/opengl/display_driver.h

@@ -0,0 +1,130 @@
+/*
+ * Copyright 2011-2022 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <atomic>
+
+#include "app/opengl/shader.h"
+
+#include "session/display_driver.h"
+
+#include "util/function.h"
+#include "util/unique_ptr.h"
+
+CCL_NAMESPACE_BEGIN
+
+class OpenGLDisplayDriver : public DisplayDriver {
+ public:
+  /* Callbacks for enabling and disabling the OpenGL context. Must be provided to support enabling
+   * the context on the Cycles render thread independent of the main thread. */
+  OpenGLDisplayDriver(const function<bool()> &gl_context_enable,
+                      const function<void()> &gl_context_disable);
+  ~OpenGLDisplayDriver();
+
+  virtual void graphics_interop_activate() override;
+  virtual void graphics_interop_deactivate() override;
+
+  virtual void clear() override;
+
+  void set_zoom(float zoom_x, float zoom_y);
+
+ protected:
+  virtual void next_tile_begin() override;
+
+  virtual bool update_begin(const Params &params, int texture_width, int texture_height) override;
+  virtual void update_end() override;
+
+  virtual half4 *map_texture_buffer() override;
+  virtual void unmap_texture_buffer() override;
+
+  virtual GraphicsInterop graphics_interop_get() override;
+
+  virtual void draw(const Params &params) override;
+
+  /* Make sure texture is allocated and its initial configuration is performed. */
+  bool gl_texture_resources_ensure();
+
+  /* Ensure all runtime GPU resources needed for drawing are allocated.
+   * Returns true if all resources needed for drawing are available. */
+  bool gl_draw_resources_ensure();
+
+  /* Destroy all GPU resources which are being used by this object. */
+  void gl_resources_destroy();
+
+  /* Update GPU texture dimensions and content if needed (new pixel data was provided).
+   *
+   * NOTE: The texture needs to be bound. */
+  void texture_update_if_needed();
+
+  /* Update vertex buffer with new coordinates of vertex positions and texture coordinates.
+   * This buffer is used to render texture in the viewport.
+   *
+   * NOTE: The buffer needs to be bound. */
+  void vertex_buffer_update(const Params &params);
+
+  /* Texture which contains pixels of the render result. */
+  struct {
+    /* Indicates whether texture creation was attempted and succeeded.
+     * Used to avoid multiple attempts of texture creation on GPU issues or GPU context
+     * misconfiguration. */
+    bool creation_attempted = false;
+    bool is_created = false;
+
+    /* OpenGL resource IDs of the texture itself and Pixel Buffer Object (PBO) used to write
+     * pixels to it.
+     *
+     * NOTE: Allocated on the engine's context. */
+    uint gl_id = 0;
+    uint gl_pbo_id = 0;
+
+    /* Is true when new data was written to the PBO, meaning, the texture might need to be resized
+     * and new data is to be uploaded to the GPU. */
+    bool need_update = false;
+
+    /* Content of the texture is to be filled with zeroes. */
+    std::atomic<bool> need_clear = true;
+
+    /* Dimensions of the texture in pixels. */
+    int width = 0;
+    int height = 0;
+
+    /* Dimensions of the underlying PBO. */
+    int buffer_width = 0;
+    int buffer_height = 0;
+  } texture_;
+
+  OpenGLShader display_shader_;
+
+  /* Special track of whether GPU resources were attempted to be created, to avoid attempts of
+   * their re-creation on failure on every redraw. */
+  bool gl_draw_resource_creation_attempted_ = false;
+  bool gl_draw_resources_created_ = false;
+
+  /* Vertex buffer which hold vertices of a triangle fan which is textures with the texture
+   * holding the render result. */
+  uint vertex_buffer_ = 0;
+
+  void *gl_render_sync_ = nullptr;
+  void *gl_upload_sync_ = nullptr;
+
+  float2 zoom_ = make_float2(1.0f, 1.0f);
+
+  function<bool()> gl_context_enable_ = nullptr;
+  function<void()> gl_context_disable_ = nullptr;
+};
+
+CCL_NAMESPACE_END

intern/cycles/app/opengl/shader.cpp

@@ -0,0 +1,210 @@
+/*
+ * Copyright 2011-2022 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "app/opengl/shader.h"
+
+#include "util/log.h"
+#include "util/string.h"
+
+#include <GL/glew.h>
+
+CCL_NAMESPACE_BEGIN
+
+/* --------------------------------------------------------------------
+ * OpenGLShader.
+ */
+
+static const char *VERTEX_SHADER =
+    "#version 330\n"
+    "uniform vec2 fullscreen;\n"
+    "in vec2 texCoord;\n"
+    "in vec2 pos;\n"
+    "out vec2 texCoord_interp;\n"
+    "\n"
+    "vec2 normalize_coordinates()\n"
+    "{\n"
+    "   return (vec2(2.0) * (pos / fullscreen)) - vec2(1.0);\n"
+    "}\n"
+    "\n"
+    "void main()\n"
+    "{\n"
+    "   gl_Position = vec4(normalize_coordinates(), 0.0, 1.0);\n"
+    "   texCoord_interp = texCoord;\n"
+    "}\n\0";
+
+static const char *FRAGMENT_SHADER =
+    "#version 330\n"
+    "uniform sampler2D image_texture;\n"
+    "in vec2 texCoord_interp;\n"
+    "out vec4 fragColor;\n"
+    "\n"
+    "void main()\n"
+    "{\n"
+    "   vec4 rgba = texture(image_texture, texCoord_interp);\n"
+    /* Harcoded Rec.709 gamma, should use OpenColorIO eventually. */
+    "   fragColor = pow(rgba, vec4(0.45, 0.45, 0.45, 1.0));\n"
+    "}\n\0";
+
+static void shader_print_errors(const char *task, const char *log, const char *code)
+{
+  LOG(ERROR) << "Shader: " << task << " error:";
+  LOG(ERROR) << "===== shader string ====";
+
+  stringstream stream(code);
+  string partial;
+
+  int line = 1;
+  while (getline(stream, partial, '\n')) {
+    if (line < 10) {
+      LOG(ERROR) << " " << line << " " << partial;
+    }
+    else {
+      LOG(ERROR) << line << " " << partial;
+    }
+    line++;
+  }
+  LOG(ERROR) << log;
+}
+
+static int compile_shader_program(void)
+{
+  const struct Shader {
+    const char *source;
+    const GLenum type;
+  } shaders[2] = {{VERTEX_SHADER, GL_VERTEX_SHADER}, {FRAGMENT_SHADER, GL_FRAGMENT_SHADER}};
+
+  const GLuint program = glCreateProgram();
+
+  for (int i = 0; i < 2; i++) {
+    const GLuint shader = glCreateShader(shaders[i].type);
+
+    string source_str = shaders[i].source;
+    const char *c_str = source_str.c_str();
+
+    glShaderSource(shader, 1, &c_str, NULL);
+    glCompileShader(shader);
+
+    GLint compile_status;
+    glGetShaderiv(shader, GL_COMPILE_STATUS, &compile_status);
+
+    if (!compile_status) {
+      GLchar log[5000];
+      GLsizei length = 0;
+      glGetShaderInfoLog(shader, sizeof(log), &length, log);
+      shader_print_errors("compile", log, c_str);
+      return 0;
+    }
+
+    glAttachShader(program, shader);
+  }
+
+  /* Link output. */
+  glBindFragDataLocation(program, 0, "fragColor");
+
+  /* Link and error check. */
+  glLinkProgram(program);
+
+  GLint link_status;
+  glGetProgramiv(program, GL_LINK_STATUS, &link_status);
+  if (!link_status) {
+    GLchar log[5000];
+    GLsizei length = 0;
+    glGetShaderInfoLog(program, sizeof(log), &length, log);
+    shader_print_errors("linking", log, VERTEX_SHADER);
+    shader_print_errors("linking", log, FRAGMENT_SHADER);
+    return 0;
+  }
+
+  return program;
+}
+
+int OpenGLShader::get_position_attrib_location()
+{
+  if (position_attribute_location_ == -1) {
+    const uint shader_program = get_shader_program();
+    position_attribute_location_ = glGetAttribLocation(shader_program, position_attribute_name);
+  }
+  return position_attribute_location_;
+}
+
+int OpenGLShader::get_tex_coord_attrib_location()
+{
+  if (tex_coord_attribute_location_ == -1) {
+    const uint shader_program = get_shader_program();
+    tex_coord_attribute_location_ = glGetAttribLocation(shader_program, tex_coord_attribute_name);
+  }
+  return tex_coord_attribute_location_;
+}
+
+void OpenGLShader::bind(int width, int height)
+{
+  create_shader_if_needed();
+
+  if (!shader_program_) {
+    return;
+  }
+
+  glUseProgram(shader_program_);
+  glUniform1i(image_texture_location_, 0);
+  glUniform2f(fullscreen_location_, width, height);
+}
+
+void OpenGLShader::unbind()
+{
+}
+
+uint OpenGLShader::get_shader_program()
+{
+  return shader_program_;
+}
+
+void OpenGLShader::create_shader_if_needed()
+{
+  if (shader_program_ || shader_compile_attempted_) {
+    return;
+  }
+
+  shader_compile_attempted_ = true;
+
+  shader_program_ = compile_shader_program();
+  if (!shader_program_) {
+    return;
+  }
+
+  glUseProgram(shader_program_);
+
+  image_texture_location_ = glGetUniformLocation(shader_program_, "image_texture");
+  if (image_texture_location_ < 0) {
+    LOG(ERROR) << "Shader doesn't contain the 'image_texture' uniform.";
+    destroy_shader();
+    return;
+  }
+
+  fullscreen_location_ = glGetUniformLocation(shader_program_, "fullscreen");
+  if (fullscreen_location_ < 0) {
+    LOG(ERROR) << "Shader doesn't contain the 'fullscreen' uniform.";
+    destroy_shader();
+    return;
+  }
+}
+
+void OpenGLShader::destroy_shader()
+{
+  glDeleteProgram(shader_program_);
+  shader_program_ = 0;
+}
+
+CCL_NAMESPACE_END

intern/cycles/app/opengl/shader.h

@@ -0,0 +1,58 @@
+/*
+ * Copyright 2011-2022 OpenGL Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include "util/types.h"
+
+CCL_NAMESPACE_BEGIN
+
+class OpenGLShader {
+ public:
+  static constexpr const char *position_attribute_name = "pos";
+  static constexpr const char *tex_coord_attribute_name = "texCoord";
+
+  OpenGLShader() = default;
+  virtual ~OpenGLShader() = default;
+
+  /* Get attribute location for position and texture coordinate respectively.
+   * NOTE: The shader needs to be bound to have access to those. */
+  int get_position_attrib_location();
+  int get_tex_coord_attrib_location();
+
+  void bind(int width, int height);
+  void unbind();
+
+ protected:
+  uint get_shader_program();
+
+  void create_shader_if_needed();
+  void destroy_shader();
+
+  /* Cached values of various OpenGL resources. */
+  int position_attribute_location_ = -1;
+  int tex_coord_attribute_location_ = -1;
+
+  uint shader_program_ = 0;
+  int image_texture_location_ = -1;
+  int fullscreen_location_ = -1;
+
+  /* Shader compilation attempted. Which means, that if the shader program is 0 then compilation or
+   * linking has failed. Do not attempt to re-compile the shader. */
+  bool shader_compile_attempted_ = false;
+};
+
+CCL_NAMESPACE_END

intern/cycles/app/opengl/window.cpp

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

intern/cycles/app/opengl/window.h

@@ -0,0 +1,48 @@
+/*
+ * Copyright 2011-2022 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+/* Functions to display a simple OpenGL window using SDL, simplified to the
+ * bare minimum we need to reduce boilerplate code in tests apps. */
+
+CCL_NAMESPACE_BEGIN
+
+typedef void (*WindowInitFunc)();
+typedef void (*WindowExitFunc)();
+typedef void (*WindowResizeFunc)(int width, int height);
+typedef void (*WindowDisplayFunc)();
+typedef void (*WindowKeyboardFunc)(unsigned char key);
+typedef void (*WindowMotionFunc)(int x, int y, int button);
+
+void window_main_loop(const char *title,
+                      int width,
+                      int height,
+                      WindowInitFunc initf,
+                      WindowExitFunc exitf,
+                      WindowResizeFunc resize,
+                      WindowDisplayFunc display,
+                      WindowKeyboardFunc keyboard,
+                      WindowMotionFunc motion);
+
+void window_display_info(const char *info);
+void window_display_help();
+void window_redraw();
+
+bool window_opengl_context_enable();
+void window_opengl_context_disable();
+
+CCL_NAMESPACE_END

intern/cycles/blender/addon/properties.py

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

intern/cycles/blender/addon/ui.py

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

intern/cycles/blender/device.cpp

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

intern/cycles/blender/display_driver.cpp

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

intern/cycles/blender/session.cpp

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

intern/cycles/blender/shader.cpp

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

intern/cycles/bvh/embree.cpp

@@ -471,7 +471,7 @@ void BVHEmbree::add_instance(Object *ob, int i)
   assert(instance_bvh != NULL);
 
   const size_t num_object_motion_steps = ob->use_motion() ? ob->get_motion().size() : 1;
-  const size_t num_motion_steps = min(num_object_motion_steps, RTC_MAX_TIME_STEP_COUNT);
+  const size_t num_motion_steps = min(num_object_motion_steps, (size_t)RTC_MAX_TIME_STEP_COUNT);
   assert(num_object_motion_steps <= RTC_MAX_TIME_STEP_COUNT);
 
   RTCGeometry geom_id = rtcNewGeometry(rtc_device, RTC_GEOMETRY_TYPE_INSTANCE);
@@ -522,7 +522,7 @@ void BVHEmbree::add_triangles(const Object *ob, const Mesh *mesh, int i)
   }
 
   assert(num_motion_steps <= RTC_MAX_TIME_STEP_COUNT);
-  num_motion_steps = min(num_motion_steps, RTC_MAX_TIME_STEP_COUNT);
+  num_motion_steps = min(num_motion_steps, (size_t)RTC_MAX_TIME_STEP_COUNT);
 
   const size_t num_triangles = mesh->num_triangles();
 
@@ -775,7 +775,7 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
   }
 
   assert(num_motion_steps <= RTC_MAX_TIME_STEP_COUNT);
-  num_motion_steps = min(num_motion_steps, RTC_MAX_TIME_STEP_COUNT);
+  num_motion_steps = min(num_motion_steps, (size_t)RTC_MAX_TIME_STEP_COUNT);
 
   const size_t num_curves = hair->num_curves();
   size_t num_segments = 0;

intern/cycles/bvh/params.h

@@ -142,6 +142,7 @@ class BVHParams {
 
     top_level = false;
     bvh_layout = BVH_LAYOUT_BVH2;
+    use_compact_structure = true;
     use_unaligned_nodes = false;
 
     num_motion_curve_steps = 0;

intern/cycles/cmake/external_libs.cmake

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

intern/cycles/device/device.cpp

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

intern/cycles/device/device.h

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

intern/cycles/device/hip/device_impl.h

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

intern/cycles/device/memory.h

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

intern/cycles/device/metal/bvh.mm

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

intern/cycles/device/metal/device.mm

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

intern/cycles/device/metal/device_impl.mm

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

intern/cycles/device/metal/kernel.mm

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

intern/cycles/device/metal/util.h

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

intern/cycles/device/metal/util.mm

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

intern/cycles/device/optix/device_impl.cpp

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

intern/cycles/integrator/denoiser.cpp

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

intern/cycles/integrator/denoiser_oidn.cpp

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

intern/cycles/integrator/denoiser_oidn.h

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

intern/cycles/integrator/path_trace.cpp

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

intern/cycles/integrator/path_trace.h

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

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

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

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

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

intern/cycles/kernel/geom/shader_data.h

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

intern/cycles/kernel/integrator/shade_surface.h

@@ -193,17 +193,24 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
   const uint16_t transparent_bounce = INTEGRATOR_STATE(state, path, transparent_bounce);
   uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
   shadow_flag |= (is_light) ? PATH_RAY_SHADOW_FOR_LIGHT : 0;
-  shadow_flag |= (shadow_flag & PATH_RAY_ANY_PASS) ? 0 : PATH_RAY_SURFACE_PASS;
   const float3 throughput = INTEGRATOR_STATE(state, path, throughput) * bsdf_eval_sum(&bsdf_eval);
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
-    const packed_float3 pass_diffuse_weight =
-        (bounce == 0) ? packed_float3(bsdf_eval_pass_diffuse_weight(&bsdf_eval)) :
-                        INTEGRATOR_STATE(state, path, pass_diffuse_weight);
-    const packed_float3 pass_glossy_weight = (bounce == 0) ?
-                                                 packed_float3(
-                                                     bsdf_eval_pass_glossy_weight(&bsdf_eval)) :
-                                                 INTEGRATOR_STATE(state, path, pass_glossy_weight);
+    packed_float3 pass_diffuse_weight;
+    packed_float3 pass_glossy_weight;
+
+    if (shadow_flag & PATH_RAY_ANY_PASS) {
+      /* Indirect bounce, use weights from earlier surface or volume bounce. */
+      pass_diffuse_weight = INTEGRATOR_STATE(state, path, pass_diffuse_weight);
+      pass_glossy_weight = INTEGRATOR_STATE(state, path, pass_glossy_weight);
+    }
+    else {
+      /* Direct light, use BSDFs at this bounce. */
+      shadow_flag |= PATH_RAY_SURFACE_PASS;
+      pass_diffuse_weight = packed_float3(bsdf_eval_pass_diffuse_weight(&bsdf_eval));
+      pass_glossy_weight = packed_float3(bsdf_eval_pass_glossy_weight(&bsdf_eval));
+    }
+
     INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = pass_diffuse_weight;
     INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = pass_glossy_weight;
   }

intern/cycles/kernel/integrator/shade_volume.h

@@ -801,16 +801,26 @@ ccl_device_forceinline void integrate_volume_direct_light(
   const uint16_t transparent_bounce = INTEGRATOR_STATE(state, path, transparent_bounce);
   uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
   shadow_flag |= (is_light) ? PATH_RAY_SHADOW_FOR_LIGHT : 0;
-  shadow_flag |= (shadow_flag & PATH_RAY_ANY_PASS) ? 0 : PATH_RAY_VOLUME_PASS;
   const float3 throughput_phase = throughput * bsdf_eval_sum(&phase_eval);
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
-    const packed_float3 pass_diffuse_weight = (bounce == 0) ?
-                                                  packed_float3(one_float3()) :
-                                                  INTEGRATOR_STATE(
-                                                      state, path, pass_diffuse_weight);
+    packed_float3 pass_diffuse_weight;
+    packed_float3 pass_glossy_weight;
+
+    if (shadow_flag & PATH_RAY_ANY_PASS) {
+      /* Indirect bounce, use weights from earlier surface or volume bounce. */
+      pass_diffuse_weight = INTEGRATOR_STATE(state, path, pass_diffuse_weight);
+      pass_glossy_weight = INTEGRATOR_STATE(state, path, pass_glossy_weight);
+    }
+    else {
+      /* Direct light, no diffuse/glossy distinction needed for volumes. */
+      shadow_flag |= PATH_RAY_VOLUME_PASS;
+      pass_diffuse_weight = packed_float3(one_float3());
+      pass_glossy_weight = packed_float3(zero_float3());
+    }
+
     INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = pass_diffuse_weight;
-    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = zero_float3();
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = pass_glossy_weight;
   }
 
   INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(

intern/cycles/kernel/light/light.h

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

intern/cycles/scene/hair.cpp

@@ -119,7 +119,7 @@ void Hair::Curve::motion_keys(const float3 *curve_keys,
 {
   /* Figure out which steps we need to fetch and their interpolation factor. */
   const size_t max_step = num_steps - 1;
-  const size_t step = min((int)(time * max_step), max_step - 1);
+  const size_t step = std::min((size_t)(time * max_step), max_step - 1);
   const float t = time * max_step - step;
   /* Fetch vertex coordinates. */
   float4 curr_keys[2];
@@ -147,7 +147,7 @@ void Hair::Curve::cardinal_motion_keys(const float3 *curve_keys,
 {
   /* Figure out which steps we need to fetch and their interpolation factor. */
   const size_t max_step = num_steps - 1;
-  const size_t step = min((int)(time * max_step), max_step - 1);
+  const size_t step = min((size_t)(time * max_step), max_step - 1);
   const float t = time * max_step - step;
   /* Fetch vertex coordinates. */
   float4 curr_keys[4];
@@ -191,8 +191,8 @@ void Hair::Curve::keys_for_step(const float3 *curve_keys,
                                 size_t k1,
                                 float4 r_keys[2]) const
 {
-  k0 = max(k0, 0);
-  k1 = min(k1, num_keys - 1);
+  k0 = max(k0, (size_t)0);
+  k1 = min(k1, (size_t)(num_keys - 1));
   const size_t center_step = ((num_steps - 1) / 2);
   if (step == center_step) {
     /* Center step: regular key location. */
@@ -237,8 +237,8 @@ void Hair::Curve::cardinal_keys_for_step(const float3 *curve_keys,
                                          size_t k3,
                                          float4 r_keys[4]) const
 {
-  k0 = max(k0, 0);
-  k3 = min(k3, num_keys - 1);
+  k0 = max(k0, (size_t)0);
+  k3 = min(k3, (size_t)(num_keys - 1));
   const size_t center_step = ((num_steps - 1) / 2);
   if (step == center_step) {
     /* Center step: regular key location. */

intern/cycles/scene/light.cpp

@@ -606,8 +606,8 @@ void LightManager::device_update_background(Device *device,
       ImageMetaData metadata;
       if (!env->handle.empty()) {
         ImageMetaData metadata = env->handle.metadata();
-        environment_res.x = max(environment_res.x, metadata.width);
-        environment_res.y = max(environment_res.y, metadata.height);
+        environment_res.x = max(environment_res.x, (int)metadata.width);
+        environment_res.y = max(environment_res.y, (int)metadata.height);
       }
     }
     if (node->type == SkyTextureNode::get_node_type()) {

intern/cycles/scene/mesh.cpp

@@ -53,7 +53,7 @@ void Mesh::Triangle::motion_verts(const float3 *verts,
 {
   /* Figure out which steps we need to fetch and their interpolation factor. */
   const size_t max_step = num_steps - 1;
-  const size_t step = min((int)(time * max_step), max_step - 1);
+  const size_t step = min((size_t)(time * max_step), max_step - 1);
   const float t = time * max_step - step;
   /* Fetch vertex coordinates. */
   float3 curr_verts[3];

intern/cycles/scene/pointcloud.cpp

@@ -55,7 +55,7 @@ float4 PointCloud::Point::motion_key(const float3 *points,
   /* Figure out which steps we need to fetch and their
    * interpolation factor. */
   const size_t max_step = num_steps - 1;
-  const size_t step = min((int)(time * max_step), max_step - 1);
+  const size_t step = min((size_t)(time * max_step), max_step - 1);
   const float t = time * max_step - step;
   /* Fetch vertex coordinates. */
   const float4 curr_key = point_for_step(

intern/cycles/scene/shader.cpp

@@ -830,28 +830,28 @@ void ShaderManager::init_xyz_transforms()
   Transform xyz_to_rgb;
 
   if (config->hasRole("aces_interchange")) {
-    /* Standard OpenColorIO role, defined as ACES2065-1. */
-    const Transform xyz_E_to_aces = make_transform(1.0498110175f,
-                                                   0.0f,
-                                                   -0.0000974845f,
-                                                   0.0f,
-                                                   -0.4959030231f,
-                                                   1.3733130458f,
-                                                   0.0982400361f,
-                                                   0.0f,
-                                                   0.0f,
-                                                   0.0f,
-                                                   0.9912520182f,
-                                                   0.0f);
-    const Transform xyz_D65_to_E = make_transform(
-        1.0521111f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.9184170f, 0.0f);
-
+    /* Standard OpenColorIO role, defined as ACES AP0 (ACES2065-1). */
     Transform aces_to_rgb;
     if (!to_scene_linear_transform(config, "aces_interchange", aces_to_rgb)) {
       return;
     }
 
-    xyz_to_rgb = aces_to_rgb * xyz_E_to_aces * xyz_D65_to_E;
+    /* This is the OpenColorIO builtin transform:
+     * UTILITY - ACES-AP0_to_CIE-XYZ-D65_BFD. */
+    const Transform ACES_AP0_to_xyz_D65 = make_transform(0.938280f,
+                                                         -0.004451f,
+                                                         0.016628f,
+                                                         0.000000f,
+                                                         0.337369f,
+                                                         0.729522f,
+                                                         -0.066890f,
+                                                         0.000000f,
+                                                         0.001174f,
+                                                         -0.003711f,
+                                                         1.091595f,
+                                                         0.000000f);
+    const Transform xyz_to_aces = transform_inverse(ACES_AP0_to_xyz_D65);
+    xyz_to_rgb = aces_to_rgb * xyz_to_aces;
   }
   else if (config->hasRole("XYZ")) {
     /* Custom role used before the standard existed. */

intern/cycles/session/session.cpp

@@ -49,12 +49,9 @@ Session::Session(const SessionParams &params_, const SceneParams &scene_params)
 {
   TaskScheduler::init(params.threads);
 
-  session_thread_ = nullptr;
-
   delayed_reset_.do_reset = false;
 
   pause_ = false;
-  cancel_ = false;
   new_work_added_ = false;
 
   device = Device::create(params.device, stats, profiler);
@@ -73,48 +70,82 @@ Session::Session(const SessionParams &params_, const SceneParams &scene_params)
     }
     full_buffer_written_cb(filename);
   };
+
+  /* Create session thread. */
+  session_thread_ = new thread(function_bind(&Session::thread_run, this));
 }
 
 Session::~Session()
 {
+  /* Cancel any ongoing render operation. */
   cancel();
 
-  /* Make sure path tracer is destroyed before the device. This is needed because destruction might
-   * need to access device for device memory free. */
-  /* TODO(sergey): Convert device to be unique_ptr, and rely on C++ to destruct objects in the
+  /* Signal session thread to end. */
+  {
+    thread_scoped_lock session_thread_lock(session_thread_mutex_);
+    session_thread_state_ = SESSION_THREAD_END;
+  }
+  session_thread_cond_.notify_all();
+
+  /* Destroy session thread. */
+  session_thread_->join();
+  delete session_thread_;
+
+  /* Destroy path tracer, before the device. This is needed because destruction might need to
+   * access device for device memory free.
+   * TODO(sergey): Convert device to be unique_ptr, and rely on C++ to destruct objects in the
    * pre-defined order. */
   path_trace_.reset();
 
+  /* Destroy scene and device. */
   delete scene;
   delete device;
 
+  /* Stop task scheduler. */
   TaskScheduler::exit();
 }
 
 void Session::start()
 {
-  if (!session_thread_) {
-    session_thread_ = new thread(function_bind(&Session::run, this));
+  {
+    /* Signal session thread to start rendering. */
+    thread_scoped_lock session_thread_lock(session_thread_mutex_);
+    if (session_thread_state_ == SESSION_THREAD_RENDER) {
+      /* Already rendering, nothing to do. */
+      return;
+    }
+    session_thread_state_ = SESSION_THREAD_RENDER;
   }
+
+  session_thread_cond_.notify_all();
 }
 
 void Session::cancel(bool quick)
 {
-  if (quick && path_trace_) {
-    path_trace_->cancel();
+  /* Check if session thread is rendering. */
+  bool rendering;
+  {
+    thread_scoped_lock session_thread_lock(session_thread_mutex_);
+    rendering = (session_thread_state_ == SESSION_THREAD_RENDER);
   }
 
-  if (session_thread_) {
-    /* wait for session thread to end */
+  if (rendering) {
+    /* Cancel path trace operations. */
+    if (quick && path_trace_) {
+      path_trace_->cancel();
+    }
+
+    /* Cancel other operations. */
     progress.set_cancel("Exiting");
 
+    /* Signal unpause in case the render was paused. */
     {
       thread_scoped_lock pause_lock(pause_mutex_);
       pause_ = false;
-      cancel_ = true;
     }
     pause_cond_.notify_all();
 
+    /* Wait for render thread to be cancelled or finished. */
     wait();
   }
 }
@@ -192,11 +223,46 @@ void Session::run_main_render_loop()
       break;
     }
   }
-
-  path_trace_->flush_display();
 }
 
-void Session::run()
+void Session::thread_run()
+{
+  while (true) {
+    {
+      thread_scoped_lock session_thread_lock(session_thread_mutex_);
+
+      if (session_thread_state_ == SESSION_THREAD_WAIT) {
+        /* Continue waiting for any signal from the main thread. */
+        session_thread_cond_.wait(session_thread_lock);
+        continue;
+      }
+      else if (session_thread_state_ == SESSION_THREAD_END) {
+        /* End thread immediately. */
+        break;
+      }
+    }
+
+    /* Execute a render. */
+    thread_render();
+
+    /* Go back from rendering to waiting. */
+    {
+      thread_scoped_lock session_thread_lock(session_thread_mutex_);
+      if (session_thread_state_ == SESSION_THREAD_RENDER) {
+        session_thread_state_ = SESSION_THREAD_WAIT;
+      }
+    }
+    session_thread_cond_.notify_all();
+  }
+
+  /* Flush any remaining operations and destroy display driver here. This ensure
+   * graphics API resources are created and destroyed all in the session thread,
+   * which can avoid problems contexts and multiple threads. */
+  path_trace_->flush_display();
+  path_trace_->set_display_driver(nullptr);
+}
+
+void Session::thread_render()
 {
   if (params.use_profiling && (params.device.type == DEVICE_CPU)) {
     profiler.start();
@@ -338,9 +404,9 @@ bool Session::run_wait_for_work(const RenderWork &render_work)
   const bool no_work = !render_work;
   update_status_time(pause_, no_work);
 
-  /* Only leave the loop when rendering is not paused. But even if the current render is un-paused
-   * but there is nothing to render keep waiting until new work is added. */
-  while (!cancel_) {
+  /* Only leave the loop when rendering is not paused. But even if the current render is
+   * un-paused but there is nothing to render keep waiting until new work is added. */
+  while (!progress.get_cancel()) {
     scoped_timer pause_timer;
 
     if (!pause_ && (render_work || new_work_added_ || delayed_reset_.do_reset)) {
@@ -427,7 +493,8 @@ void Session::do_delayed_reset()
   tile_manager_.update(buffer_params_, scene);
 
   /* Update temp directory on reset.
-   * This potentially allows to finish the existing rendering with a previously configure temporary
+   * This potentially allows to finish the existing rendering with a previously configure
+   * temporary
    * directory in the host software and switch to a new temp directory when new render starts. */
   tile_manager_.set_temp_dir(params.temp_dir);
 
@@ -544,12 +611,14 @@ double Session::get_estimated_remaining_time() const
 
 void Session::wait()
 {
-  if (session_thread_) {
-    session_thread_->join();
-    delete session_thread_;
+  /* Wait until session thread either is waiting or ending. */
+  while (true) {
+    thread_scoped_lock session_thread_lock(session_thread_mutex_);
+    if (session_thread_state_ != SESSION_THREAD_RENDER) {
+      break;
+    }
+    session_thread_cond_.wait(session_thread_lock);
   }
-
-  session_thread_ = nullptr;
 }
 
 bool Session::update_scene(int width, int height)

intern/cycles/session/session.h

@@ -172,7 +172,8 @@ class Session {
     BufferParams buffer_params;
   } delayed_reset_;
 
-  void run();
+  void thread_run();
+  void thread_render();
 
   /* Update for the new iteration of the main loop in run implementation (run_cpu and run_gpu).
    *
@@ -205,10 +206,19 @@ class Session {
 
   int2 get_effective_tile_size() const;
 
-  thread *session_thread_;
+  /* Session thread that performs rendering tasks decoupled from the thread
+   * controlling the sessions. The thread is created and destroyed along with
+   * the session. */
+  thread *session_thread_ = nullptr;
+  thread_condition_variable session_thread_cond_;
+  thread_mutex session_thread_mutex_;
+  enum {
+    SESSION_THREAD_WAIT,
+    SESSION_THREAD_RENDER,
+    SESSION_THREAD_END,
+  } session_thread_state_ = SESSION_THREAD_WAIT;
 
   bool pause_ = false;
-  bool cancel_ = false;
   bool new_work_added_ = false;
 
   thread_condition_variable pause_cond_;

intern/cycles/util/CMakeLists.txt

@@ -18,7 +18,6 @@ set(INC
 )
 
 set(INC_SYS
-  ${GLEW_INCLUDE_DIR}
 )
 
 set(SRC
@@ -45,14 +44,6 @@ set(LIB
   ${TBB_LIBRARIES}
 )
 
-if(WITH_CYCLES_STANDALONE)
-  if(WITH_CYCLES_STANDALONE_GUI)
-    list(APPEND SRC
-      view.cpp
-    )
-  endif()
-endif()
-
 set(SRC_HEADERS
   algorithm.h
   aligned_malloc.h
@@ -153,7 +144,6 @@ set(SRC_HEADERS
   unique_ptr.h
   vector.h
   version.h
-  view.h
   windows.h
   xml.h
 )

intern/cycles/util/algorithm.h

@@ -21,8 +21,6 @@
 
 CCL_NAMESPACE_BEGIN
 
-using std::max;
-using std::min;
 using std::remove;
 using std::sort;
 using std::stable_sort;

intern/cycles/util/math.h

@@ -124,7 +124,41 @@ ccl_device_inline int min(int a, int b)
   return (a < b) ? a : b;
 }
 
-ccl_device_inline uint min(uint a, uint b)
+ccl_device_inline uint32_t max(uint32_t a, uint32_t b)
+{
+  return (a > b) ? a : b;
+}
+
+ccl_device_inline uint32_t min(uint32_t a, uint32_t b)
+{
+  return (a < b) ? a : b;
+}
+
+ccl_device_inline uint64_t max(uint64_t a, uint64_t b)
+{
+  return (a > b) ? a : b;
+}
+
+ccl_device_inline uint64_t min(uint64_t a, uint64_t b)
+{
+  return (a < b) ? a : b;
+}
+
+/* NOTE: On 64bit Darwin the `size_t` is defined as `unsigned long int` and `uint64_t` is defined
+ * as `unsigned long long`. Both of the definitions are 64 bit unsigned integer, but the automatic
+ * substitution does not allow to automatically pick function defined for `uint64_t` as it is not
+ * exactly the same type definition.
+ * Work this around by adding a templated function enabled for `size_t` type which will be used
+ * when there is no explicit specialization of `min()`/`max()` above. */
+
+template<class T>
+ccl_device_inline typename std::enable_if_t<std::is_same_v<T, size_t>, T> max(T a, T b)
+{
+  return (a > b) ? a : b;
+}
+
+template<class T>
+ccl_device_inline typename std::enable_if_t<std::is_same_v<T, size_t>, T> min(T a, T b)
 {
   return (a < b) ? a : b;
 }
@@ -901,9 +935,15 @@ ccl_device_inline uint prev_power_of_two(uint x)
 ccl_device_inline uint32_t reverse_integer_bits(uint32_t x)
 {
   /* Use a native instruction if it exists. */
-#if defined(__arm__) || defined(__aarch64__)
+#if defined(__aarch64__) || defined(_M_ARM64)
+  /* Assume the rbit is always available on 64bit ARM architecture. */
   __asm__("rbit %w0, %w1" : "=r"(x) : "r"(x));
   return x;
+#elif defined(__arm__) && ((__ARM_ARCH > 7) || __ARM_ARCH == 6 && __ARM_ARCH_ISA_THUMB >= 2)
+  /* This ARM instruction is available in ARMv6T2 and above.
+   * This 32-bit Thumb instruction is available in ARMv6T2 and above. */
+  __asm__("rbit %0, %1" : "=r"(x) : "r"(x));
+  return x;
 #elif defined(__KERNEL_CUDA__)
   return __brev(x);
 #elif defined(__KERNEL_METAL__)

intern/cycles/util/murmurhash.cpp

@@ -23,7 +23,7 @@
 #include <stdlib.h>
 #include <string.h>
 
-#include "util/algorithm.h"
+#include "util/math.h"
 #include "util/murmurhash.h"
 
 #if defined(_MSC_VER)

intern/cycles/util/view.cpp

@@ -1,282 +0,0 @@
-/*
- * Copyright 2011-2013 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#include <stdio.h>
-#include <stdlib.h>
-
-#include "util/opengl.h"
-#include "util/string.h"
-#include "util/time.h"
-#include "util/version.h"
-#include "util/view.h"
-
-#ifdef __APPLE__
-#  include <GLUT/glut.h>
-#else
-#  include <GL/glut.h>
-#endif
-
-CCL_NAMESPACE_BEGIN
-
-/* structs */
-
-struct View {
-  ViewInitFunc initf;
-  ViewExitFunc exitf;
-  ViewResizeFunc resize;
-  ViewDisplayFunc display;
-  ViewKeyboardFunc keyboard;
-  ViewMotionFunc motion;
-
-  bool first_display;
-  bool redraw;
-
-  int mouseX, mouseY;
-  int mouseBut0, mouseBut2;
-
-  int width, height;
-} V;
-
-/* public */
-
-static void view_display_text(int x, int y, const char *text)
-{
-  const char *c;
-
-  glRasterPos3f(x, y, 0);
-
-  for (c = text; *c != '\0'; c++)
-    glutBitmapCharacter(GLUT_BITMAP_HELVETICA_10, *c);
-}
-
-void view_display_info(const char *info)
-{
-  const int height = 20;
-
-  glEnable(GL_BLEND);
-  glBlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
-  glColor4f(0.1f, 0.1f, 0.1f, 0.8f);
-  glRectf(0.0f, V.height - height, V.width, V.height);
-  glDisable(GL_BLEND);
-
-  glColor3f(0.5f, 0.5f, 0.5f);
-
-  view_display_text(10, 7 + V.height - height, info);
-
-  glColor3f(1.0f, 1.0f, 1.0f);
-}
-
-void view_display_help()
-{
-  const int w = (int)((float)V.width / 1.15f);
-  const int h = (int)((float)V.height / 1.15f);
-
-  const int x1 = (V.width - w) / 2;
-  const int x2 = x1 + w;
-
-  const int y1 = (V.height - h) / 2;
-  const int y2 = y1 + h;
-
-  glEnable(GL_BLEND);
-  glBlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
-  glColor4f(0.5f, 0.5f, 0.5f, 0.8f);
-  glRectf(x1, y1, x2, y2);
-  glDisable(GL_BLEND);
-
-  glColor3f(0.8f, 0.8f, 0.8f);
-
-  string info = string("Cycles Renderer ") + CYCLES_VERSION_STRING;
-
-  view_display_text(x1 + 20, y2 - 20, info.c_str());
-  view_display_text(x1 + 20, y2 - 40, "(C) 2011-2016 Blender Foundation");
-  view_display_text(x1 + 20, y2 - 80, "Controls:");
-  view_display_text(x1 + 20, y2 - 100, "h:  Info/Help");
-  view_display_text(x1 + 20, y2 - 120, "r:  Reset");
-  view_display_text(x1 + 20, y2 - 140, "p:  Pause");
-  view_display_text(x1 + 20, y2 - 160, "esc:  Cancel");
-  view_display_text(x1 + 20, y2 - 180, "q:  Quit program");
-
-  view_display_text(x1 + 20, y2 - 210, "i:  Interactive mode");
-  view_display_text(x1 + 20, y2 - 230, "Left mouse:  Move camera");
-  view_display_text(x1 + 20, y2 - 250, "Right mouse:  Rotate camera");
-  view_display_text(x1 + 20, y2 - 270, "W/A/S/D:  Move camera");
-  view_display_text(x1 + 20, y2 - 290, "0/1/2/3:  Set max bounces");
-
-  glColor3f(1.0f, 1.0f, 1.0f);
-}
-
-static void view_display()
-{
-  if (V.first_display) {
-    if (V.initf)
-      V.initf();
-    if (V.exitf)
-      atexit(V.exitf);
-
-    V.first_display = false;
-  }
-
-  glClearColor(0.05f, 0.05f, 0.05f, 0.0f);
-  glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
-
-  glMatrixMode(GL_PROJECTION);
-  glLoadIdentity();
-  glOrtho(0, V.width, 0, V.height, -1, 1);
-
-  glMatrixMode(GL_MODELVIEW);
-  glLoadIdentity();
-
-  glRasterPos3f(0, 0, 0);
-
-  if (V.display)
-    V.display();
-
-  glutSwapBuffers();
-}
-
-static void view_reshape(int width, int height)
-{
-  if (width <= 0 || height <= 0)
-    return;
-
-  V.width = width;
-  V.height = height;
-
-  glViewport(0, 0, width, height);
-
-  glMatrixMode(GL_PROJECTION);
-  glLoadIdentity();
-
-  glMatrixMode(GL_MODELVIEW);
-  glLoadIdentity();
-
-  if (V.resize)
-    V.resize(width, height);
-}
-
-static void view_keyboard(unsigned char key, int x, int y)
-{
-  if (V.keyboard)
-    V.keyboard(key);
-
-  if (key == 'm')
-    printf("mouse %d %d\n", x, y);
-  if (key == 'q') {
-    if (V.exitf)
-      V.exitf();
-    exit(0);
-  }
-}
-
-static void view_mouse(int button, int state, int x, int y)
-{
-  if (button == 0) {
-    if (state == GLUT_DOWN) {
-      V.mouseX = x;
-      V.mouseY = y;
-      V.mouseBut0 = 1;
-    }
-    else if (state == GLUT_UP) {
-      V.mouseBut0 = 0;
-    }
-  }
-  else if (button == 2) {
-    if (state == GLUT_DOWN) {
-      V.mouseX = x;
-      V.mouseY = y;
-      V.mouseBut2 = 1;
-    }
-    else if (state == GLUT_UP) {
-      V.mouseBut2 = 0;
-    }
-  }
-}
-
-static void view_motion(int x, int y)
-{
-  const int but = V.mouseBut0 ? 0 : 2;
-  const int distX = x - V.mouseX;
-  const int distY = y - V.mouseY;
-
-  if (V.motion)
-    V.motion(distX, distY, but);
-
-  V.mouseX = x;
-  V.mouseY = y;
-}
-
-static void view_idle()
-{
-  if (V.redraw) {
-    V.redraw = false;
-    glutPostRedisplay();
-  }
-
-  time_sleep(0.1);
-}
-
-void view_main_loop(const char *title,
-                    int width,
-                    int height,
-                    ViewInitFunc initf,
-                    ViewExitFunc exitf,
-                    ViewResizeFunc resize,
-                    ViewDisplayFunc display,
-                    ViewKeyboardFunc keyboard,
-                    ViewMotionFunc motion)
-{
-  const char *name = "app";
-  char *argv = (char *)name;
-  int argc = 1;
-
-  memset(&V, 0, sizeof(V));
-  V.width = width;
-  V.height = height;
-  V.first_display = true;
-  V.redraw = false;
-  V.initf = initf;
-  V.exitf = exitf;
-  V.resize = resize;
-  V.display = display;
-  V.keyboard = keyboard;
-  V.motion = motion;
-
-  glutInit(&argc, &argv);
-  glutInitWindowSize(width, height);
-  glutInitWindowPosition(0, 0);
-  glutInitDisplayMode(GLUT_RGB | GLUT_DOUBLE | GLUT_DEPTH);
-  glutCreateWindow(title);
-
-  glewInit();
-
-  view_reshape(width, height);
-
-  glutDisplayFunc(view_display);
-  glutIdleFunc(view_idle);
-  glutReshapeFunc(view_reshape);
-  glutKeyboardFunc(view_keyboard);
-  glutMouseFunc(view_mouse);
-  glutMotionFunc(view_motion);
-
-  glutMainLoop();
-}
-
-void view_redraw()
-{
-  V.redraw = true;
-}
-
-CCL_NAMESPACE_END

intern/cycles/util/view.h

@@ -1,48 +0,0 @@
-/*
- * Copyright 2011-2013 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#ifndef __UTIL_VIEW_H__
-#define __UTIL_VIEW_H__
-
-/* Functions to display a simple OpenGL window using GLUT, simplified to the
- * bare minimum we need to reduce boilerplate code in tests apps. */
-
-CCL_NAMESPACE_BEGIN
-
-typedef void (*ViewInitFunc)();
-typedef void (*ViewExitFunc)();
-typedef void (*ViewResizeFunc)(int width, int height);
-typedef void (*ViewDisplayFunc)();
-typedef void (*ViewKeyboardFunc)(unsigned char key);
-typedef void (*ViewMotionFunc)(int x, int y, int button);
-
-void view_main_loop(const char *title,
-                    int width,
-                    int height,
-                    ViewInitFunc initf,
-                    ViewExitFunc exitf,
-                    ViewResizeFunc resize,
-                    ViewDisplayFunc display,
-                    ViewKeyboardFunc keyboard,
-                    ViewMotionFunc motion);
-
-void view_display_info(const char *info);
-void view_display_help();
-void view_redraw();
-
-CCL_NAMESPACE_END
-
-#endif /*__UTIL_VIEW_H__*/

intern/ghost/CMakeLists.txt

@@ -291,7 +291,7 @@ elseif(WITH_GHOST_X11 OR WITH_GHOST_WAYLAND)
     include(CheckSymbolExists)
     set(CMAKE_REQUIRED_DEFINITIONS "-D_GNU_SOURCE")
     check_symbol_exists(memfd_create "sys/mman.h" HAVE_MEMFD_CREATE)
-    if (HAVE_MEMFD_CREATE)
+    if(HAVE_MEMFD_CREATE)
       add_definitions(-DHAVE_MEMFD_CREATE)
     endif()
 
@@ -307,7 +307,7 @@ elseif(WITH_GHOST_X11 OR WITH_GHOST_WAYLAND)
     pkg_get_variable(WAYLAND_SCANNER wayland-scanner wayland_scanner)
 
     pkg_check_modules(wayland-protocols wayland-protocols>=1.15)
-    if (${wayland-protocols_FOUND})
+    if(${wayland-protocols_FOUND})
       pkg_get_variable(WAYLAND_PROTOCOLS_DIR wayland-protocols pkgdatadir)
     else()
       find_path(WAYLAND_PROTOCOLS_DIR
@@ -316,7 +316,7 @@ elseif(WITH_GHOST_X11 OR WITH_GHOST_WAYLAND)
       )
     endif()
 
-    if (NOT EXISTS ${WAYLAND_PROTOCOLS_DIR})
+    if(NOT EXISTS ${WAYLAND_PROTOCOLS_DIR})
       message(FATAL_ERROR "path to wayland-protocols not found")
     endif()
 
@@ -518,11 +518,11 @@ if(WITH_XR_OPENXR)
     )
   elseif(UNIX AND NOT APPLE)
     list(APPEND XR_PLATFORM_DEFINES -DXR_OS_LINUX)
-    if (WITH_GHOST_WAYLAND)
+    if(WITH_GHOST_WAYLAND)
       list(APPEND XR_PLATFORM_DEFINES -DXR_USE_PLATFORM_WAYLAND)
     endif()
-    if (WITH_GHOST_X11)
-      if (WITH_GL_EGL)
+    if(WITH_GHOST_X11)
+      if(WITH_GL_EGL)
         list(APPEND XR_PLATFORM_DEFINES -DXR_USE_PLATFORM_EGL)
       else()
         list(APPEND XR_PLATFORM_DEFINES -DXR_USE_PLATFORM_XLIB)

intern/ghost/intern/GHOST_ContextD3D.cpp

@@ -124,9 +124,11 @@ class GHOST_SharedOpenGLResource {
   struct SharedData {
     HANDLE device;
     GLuint fbo;
-    HANDLE render_buf{nullptr};
+    HANDLE render_target{nullptr};
   } m_shared;
 
+  enum RenderTarget { TARGET_RENDERBUF, TARGET_TEX2D };
+
  public:
   GHOST_SharedOpenGLResource(ID3D11Device *device,
                              ID3D11DeviceContext *device_ctx,
@@ -193,37 +195,64 @@ class GHOST_SharedOpenGLResource {
     }
 
     if (m_is_initialized) {
-      if (m_shared.render_buf) {
-        wglDXUnregisterObjectNV(m_shared.device, m_shared.render_buf);
+      if (m_shared.render_target
+#if 1
+          /* TODO: #wglDXUnregisterObjectNV() causes an access violation on AMD when the shared
+           * resource is a GL texture. Since there is currently no good alternative, just skip
+           * unregistering the shared resource. */
+          && !m_use_gl_texture2d
+#endif
+      ) {
+        wglDXUnregisterObjectNV(m_shared.device, m_shared.render_target);
       }
       if (m_shared.device) {
         wglDXCloseDeviceNV(m_shared.device);
       }
       glDeleteFramebuffers(1, &m_shared.fbo);
-      glDeleteRenderbuffers(1, &m_gl_render_buf);
+      if (m_use_gl_texture2d) {
+        glDeleteTextures(1, &m_gl_render_target);
+      }
+      else {
+        glDeleteRenderbuffers(1, &m_gl_render_target);
+      }
     }
   }
 
-  void reregisterSharedObject()
+  /* Returns true if the shared object was successfully registered, false otherwise. */
+  bool reregisterSharedObject(RenderTarget target)
   {
-    if (m_shared.render_buf) {
-      wglDXUnregisterObjectNV(m_shared.device, m_shared.render_buf);
+    if (m_shared.render_target) {
+      wglDXUnregisterObjectNV(m_shared.device, m_shared.render_target);
     }
 
     if (!m_render_target_tex) {
-      return;
+      return false;
     }
 
-    m_shared.render_buf = wglDXRegisterObjectNV(m_shared.device,
-                                                m_render_target_tex,
-                                                m_gl_render_buf,
-                                                GL_RENDERBUFFER,
-                                                WGL_ACCESS_READ_WRITE_NV);
+    if (target == TARGET_TEX2D) {
+      glTexImage2D(GL_TEXTURE_2D,
+                   0,
+                   GL_RGBA8,
+                   m_cur_width,
+                   m_cur_height,
+                   0,
+                   GL_RGBA,
+                   GL_UNSIGNED_BYTE,
+                   nullptr);
+    }
 
-    if (!m_shared.render_buf) {
+    m_shared.render_target = wglDXRegisterObjectNV(m_shared.device,
+                                                   m_render_target_tex,
+                                                   m_gl_render_target,
+                                                   (target == TARGET_TEX2D) ? GL_TEXTURE_2D :
+                                                                              GL_RENDERBUFFER,
+                                                   WGL_ACCESS_READ_WRITE_NV);
+    if (!m_shared.render_target) {
       fprintf(stderr, "Error registering shared object using wglDXRegisterObjectNV()\n");
-      return;
+      return false;
     }
+
+    return true;
   }
 
   GHOST_TSuccess initialize()
@@ -235,16 +264,33 @@ class GHOST_SharedOpenGLResource {
     }
 
     /* Build the renderbuffer. */
-    glGenRenderbuffers(1, &m_gl_render_buf);
-    glBindRenderbuffer(GL_RENDERBUFFER, m_gl_render_buf);
+    glGenRenderbuffers(1, &m_gl_render_target);
+    glBindRenderbuffer(GL_RENDERBUFFER, m_gl_render_target);
 
-    reregisterSharedObject();
+    if (!reregisterSharedObject(TARGET_RENDERBUF)) {
+      glBindRenderbuffer(GL_RENDERBUFFER, 0);
+      if (m_gl_render_target) {
+        glDeleteRenderbuffers(1, &m_gl_render_target);
+      }
+      /* Fall back to texture 2d. */
+      m_use_gl_texture2d = true;
+      glGenTextures(1, &m_gl_render_target);
+      glBindTexture(GL_TEXTURE_2D, m_gl_render_target);
+
+      reregisterSharedObject(TARGET_TEX2D);
+    }
 
     /* Build the framebuffer */
     glGenFramebuffers(1, &m_shared.fbo);
     glBindFramebuffer(GL_FRAMEBUFFER, m_shared.fbo);
-    glFramebufferRenderbuffer(
-        GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_RENDERBUFFER, m_gl_render_buf);
+    if (m_use_gl_texture2d) {
+      glFramebufferTexture2D(
+          GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, m_gl_render_target, 0);
+    }
+    else {
+      glFramebufferRenderbuffer(
+          GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_RENDERBUFFER, m_gl_render_target);
+    }
     m_is_initialized = true;
 
     return GHOST_kSuccess;
@@ -259,7 +305,7 @@ class GHOST_SharedOpenGLResource {
     if ((m_cur_width != width) || (m_cur_height != height)) {
       m_cur_width = width;
       m_cur_height = height;
-      reregisterSharedObject();
+      reregisterSharedObject(m_use_gl_texture2d ? TARGET_TEX2D : TARGET_RENDERBUF);
     }
   }
 
@@ -307,18 +353,19 @@ class GHOST_SharedOpenGLResource {
  private:
   void beginGLOnly()
   {
-    wglDXLockObjectsNV(m_shared.device, 1, &m_shared.render_buf);
+    wglDXLockObjectsNV(m_shared.device, 1, &m_shared.render_target);
   }
   void endGLOnly()
   {
-    wglDXUnlockObjectsNV(m_shared.device, 1, &m_shared.render_buf);
+    wglDXUnlockObjectsNV(m_shared.device, 1, &m_shared.render_target);
   }
 
   ID3D11Device *m_device;
   ID3D11DeviceContext *m_device_ctx;
-  GLuint m_gl_render_buf;
+  GLuint m_gl_render_target;
   unsigned int m_cur_width, m_cur_height;
   bool m_is_initialized{false};
+  bool m_use_gl_texture2d{false};
 };
 
 GHOST_SharedOpenGLResource *GHOST_ContextD3D::createSharedOpenGLResource(

intern/ghost/intern/GHOST_XrGraphicsBinding.cpp

@@ -84,16 +84,26 @@ class GHOST_XrGraphicsBindingOpenGL : public GHOST_IXrGraphicsBinding {
 #endif
     static PFN_xrGetOpenGLGraphicsRequirementsKHR s_xrGetOpenGLGraphicsRequirementsKHR_fn =
         nullptr;
+    // static XrInstance s_instance = XR_NULL_HANDLE;
     XrGraphicsRequirementsOpenGLKHR gpu_requirements = {XR_TYPE_GRAPHICS_REQUIREMENTS_OPENGL_KHR};
     const XrVersion gl_version = XR_MAKE_VERSION(
         ctx_gl.m_contextMajorVersion, ctx_gl.m_contextMinorVersion, 0);
 
+    /* Although it would seem reasonable that the proc address would not change if the instance was
+     * the same, in testing, repeated calls to #xrGetInstanceProcAddress() with the same instance
+     * can still result in changes so the workaround is to simply set the function pointer every
+     * time (trivializing its 'static' designation). */
+    // if (instance != s_instance) {
+    // s_instance = instance;
+    s_xrGetOpenGLGraphicsRequirementsKHR_fn = nullptr;
+    //}
     if (!s_xrGetOpenGLGraphicsRequirementsKHR_fn &&
         XR_FAILED(xrGetInstanceProcAddr(
             instance,
             "xrGetOpenGLGraphicsRequirementsKHR",
             (PFN_xrVoidFunction *)&s_xrGetOpenGLGraphicsRequirementsKHR_fn))) {
       s_xrGetOpenGLGraphicsRequirementsKHR_fn = nullptr;
+      return false;
     }
 
     s_xrGetOpenGLGraphicsRequirementsKHR_fn(instance, system_id, &gpu_requirements);
@@ -165,11 +175,18 @@ class GHOST_XrGraphicsBindingOpenGL : public GHOST_IXrGraphicsBinding {
                                                bool &r_is_srgb_format) const override
   {
     std::vector<int64_t> gpu_binding_formats = {
+#if 0 /* RGB10A2, RGBA16 don't seem to work with Oculus head-sets, \
+       * so move them after RGBA16F for the time being. */
         GL_RGB10_A2,
         GL_RGBA16,
-        GL_RGBA16F,
-        GL_RGBA8,
-        GL_SRGB8_ALPHA8,
+#endif
+      GL_RGBA16F,
+#if 1
+      GL_RGB10_A2,
+      GL_RGBA16,
+#endif
+      GL_RGBA8,
+      GL_SRGB8_ALPHA8,
     };
 
     std::optional result = choose_swapchain_format_from_candidates(gpu_binding_formats,
@@ -305,14 +322,24 @@ class GHOST_XrGraphicsBindingD3D : public GHOST_IXrGraphicsBinding {
       std::string *r_requirement_info) const override
   {
     static PFN_xrGetD3D11GraphicsRequirementsKHR s_xrGetD3D11GraphicsRequirementsKHR_fn = nullptr;
+    // static XrInstance s_instance = XR_NULL_HANDLE;
     XrGraphicsRequirementsD3D11KHR gpu_requirements = {XR_TYPE_GRAPHICS_REQUIREMENTS_D3D11_KHR};
 
+    /* Although it would seem reasonable that the proc address would not change if the instance was
+     * the same, in testing, repeated calls to #xrGetInstanceProcAddress() with the same instance
+     * can still result in changes so the workaround is to simply set the function pointer every
+     * time (trivializing its 'static' designation). */
+    // if (instance != s_instance) {
+    // s_instance = instance;
+    s_xrGetD3D11GraphicsRequirementsKHR_fn = nullptr;
+    //}
     if (!s_xrGetD3D11GraphicsRequirementsKHR_fn &&
         XR_FAILED(xrGetInstanceProcAddr(
             instance,
             "xrGetD3D11GraphicsRequirementsKHR",
             (PFN_xrVoidFunction *)&s_xrGetD3D11GraphicsRequirementsKHR_fn))) {
       s_xrGetD3D11GraphicsRequirementsKHR_fn = nullptr;
+      return false;
     }
 
     s_xrGetD3D11GraphicsRequirementsKHR_fn(instance, system_id, &gpu_requirements);
@@ -341,14 +368,15 @@ class GHOST_XrGraphicsBindingD3D : public GHOST_IXrGraphicsBinding {
                                                bool &r_is_srgb_format) const override
   {
     std::vector<int64_t> gpu_binding_formats = {
-#  if 0 /* RGB10A2 doesn't seem to work with Oculus head-sets, \
-         * so move it after RGB16AF for the time being. */
+#  if 0 /* RGB10A2, RGBA16 don't seem to work with Oculus head-sets, \
+         * so move them after RGBA16F for the time being. */
         DXGI_FORMAT_R10G10B10A2_UNORM,
+        DXGI_FORMAT_R16G16B16A16_UNORM,
 #  endif
-      DXGI_FORMAT_R16G16B16A16_UNORM,
       DXGI_FORMAT_R16G16B16A16_FLOAT,
 #  if 1
       DXGI_FORMAT_R10G10B10A2_UNORM,
+      DXGI_FORMAT_R16G16B16A16_UNORM,
 #  endif
       DXGI_FORMAT_R8G8B8A8_UNORM,
       DXGI_FORMAT_R8G8B8A8_UNORM_SRGB,

intern/opencolorio/ocio_impl.cc

@@ -336,16 +336,18 @@ void OCIOImpl::configGetXYZtoRGB(OCIO_ConstConfigRcPtr *config_, float xyz_to_rg
   }
 
   if (config->hasRole("aces_interchange")) {
-    /* Standard OpenColorIO role, defined as ACES2065-1. */
-    const float xyz_E_to_aces[3][3] = {{1.0498110175f, -0.4959030231f, 0.0f},
-                                       {0.0f, 1.3733130458f, 0.0f},
-                                       {-0.0000974845f, 0.0982400361f, 0.9912520182f}};
-    const float xyz_D65_to_E[3][3] = {
-        {1.0521111f, 0.0f, 0.0f}, {0.0f, 1.0f, 0.0f}, {0.0f, 0.0f, 0.9184170f}};
-
+    /* Standard OpenColorIO role, defined as ACES AP0 (ACES2065-1). */
     float aces_to_rgb[3][3];
     if (to_scene_linear_matrix(config, "aces_interchange", aces_to_rgb)) {
-      mul_m3_series(xyz_to_rgb, aces_to_rgb, xyz_E_to_aces, xyz_D65_to_E);
+      /* This is the OpenColorIO builtin transform:
+       * UTILITY - ACES-AP0_to_CIE-XYZ-D65_BFD. */
+      const float ACES_AP0_to_xyz_D65[3][3] = {{0.938280f, 0.337369f, 0.001174f},
+                                               {-0.004451f, 0.729522f, -0.003711f},
+                                               {0.016628f, -0.066890f, 1.091595f}};
+      float xyz_to_aces[3][3];
+      invert_m3_m3(xyz_to_aces, ACES_AP0_to_xyz_D65);
+
+      mul_m3_m3m3(xyz_to_rgb, aces_to_rgb, xyz_to_aces);
     }
   }
   else if (config->hasRole("XYZ")) {

intern/opensubdiv/CMakeLists.txt

@@ -83,6 +83,8 @@ if(WITH_OPENSUBDIV)
     internal/evaluator/evaluator_capi.cc
     internal/evaluator/evaluator_impl.cc
     internal/evaluator/evaluator_impl.h
+    internal/evaluator/gl_compute_evaluator.cc
+    internal/evaluator/gl_compute_evaluator.h
     internal/evaluator/patch_map.cc
     internal/evaluator/patch_map.h
 
@@ -121,6 +123,8 @@ if(WITH_OPENSUBDIV)
     add_definitions(-DNOMINMAX)
     add_definitions(-D_USE_MATH_DEFINES)
   endif()
+
+  data_to_c_simple(internal/evaluator/shaders/glsl_compute_kernel.glsl SRC)
 else()
   list(APPEND SRC
     stub/opensubdiv_stub.cc

intern/opensubdiv/internal/evaluator/eval_output_gpu.h

@@ -20,13 +20,11 @@
 #define OPENSUBDIV_EVAL_OUTPUT_GPU_H_
 
 #include "internal/evaluator/eval_output.h"
+#include "internal/evaluator/gl_compute_evaluator.h"
 
-#include <opensubdiv/osd/glComputeEvaluator.h>
 #include <opensubdiv/osd/glPatchTable.h>
 #include <opensubdiv/osd/glVertexBuffer.h>
 
-using OpenSubdiv::Osd::GLComputeEvaluator;
-using OpenSubdiv::Osd::GLStencilTableSSBO;
 using OpenSubdiv::Osd::GLVertexBuffer;
 
 namespace blender {

intern/opensubdiv/internal/evaluator/gl_compute_evaluator.cc

@@ -0,0 +1,647 @@
+//
+//   Copyright 2015 Pixar
+//
+//   Licensed under the Apache License, Version 2.0 (the "Apache License")
+//   with the following modification; you may not use this file except in
+//   compliance with the Apache License and the following modification to it:
+//   Section 6. Trademarks. is deleted and replaced with:
+//
+//   6. Trademarks. This License does not grant permission to use the trade
+//      names, trademarks, service marks, or product names of the Licensor
+//      and its affiliates, except as required to comply with Section 4(c) of
+//      the License and to reproduce the content of the NOTICE file.
+//
+//   You may obtain a copy of the Apache License at
+//
+//       http://www.apache.org/licenses/LICENSE-2.0
+//
+//   Unless required by applicable law or agreed to in writing, software
+//   distributed under the Apache License with the above modification is
+//   distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+//   KIND, either express or implied. See the Apache License for the specific
+//   language governing permissions and limitations under the Apache License.
+//
+
+#include "gl_compute_evaluator.h"
+
+#include <GL/glew.h>
+
+#include <opensubdiv/far/error.h>
+#include <opensubdiv/far/patchDescriptor.h>
+#include <opensubdiv/far/stencilTable.h>
+#include <opensubdiv/osd/glslPatchShaderSource.h>
+
+#include <cassert>
+#include <cmath>
+#include <sstream>
+#include <string>
+#include <vector>
+
+using OpenSubdiv::Far::LimitStencilTable;
+using OpenSubdiv::Far::StencilTable;
+using OpenSubdiv::Osd::BufferDescriptor;
+using OpenSubdiv::Osd::PatchArray;
+using OpenSubdiv::Osd::PatchArrayVector;
+
+extern "C" char datatoc_glsl_compute_kernel_glsl[];
+
+namespace blender {
+namespace opensubdiv {
+
+template<class T> GLuint createSSBO(std::vector<T> const &src)
+{
+  if (src.empty()) {
+    return 0;
+  }
+
+  GLuint devicePtr = 0;
+
+#if defined(GL_ARB_direct_state_access)
+  if (GLEW_ARB_direct_state_access) {
+    glCreateBuffers(1, &devicePtr);
+    glNamedBufferData(devicePtr, src.size() * sizeof(T), &src.at(0), GL_STATIC_DRAW);
+  }
+  else
+#endif
+  {
+    GLint prev = 0;
+    glGetIntegerv(GL_SHADER_STORAGE_BUFFER_BINDING, &prev);
+    glGenBuffers(1, &devicePtr);
+    glBindBuffer(GL_SHADER_STORAGE_BUFFER, devicePtr);
+    glBufferData(GL_SHADER_STORAGE_BUFFER, src.size() * sizeof(T), &src.at(0), GL_STATIC_DRAW);
+    glBindBuffer(GL_SHADER_STORAGE_BUFFER, prev);
+  }
+
+  return devicePtr;
+}
+
+GLStencilTableSSBO::GLStencilTableSSBO(StencilTable const *stencilTable)
+{
+  _numStencils = stencilTable->GetNumStencils();
+  if (_numStencils > 0) {
+    _sizes = createSSBO(stencilTable->GetSizes());
+    _offsets = createSSBO(stencilTable->GetOffsets());
+    _indices = createSSBO(stencilTable->GetControlIndices());
+    _weights = createSSBO(stencilTable->GetWeights());
+    _duWeights = _dvWeights = 0;
+    _duuWeights = _duvWeights = _dvvWeights = 0;
+  }
+  else {
+    _sizes = _offsets = _indices = _weights = 0;
+    _duWeights = _dvWeights = 0;
+    _duuWeights = _duvWeights = _dvvWeights = 0;
+  }
+}
+
+GLStencilTableSSBO::GLStencilTableSSBO(LimitStencilTable const *limitStencilTable)
+{
+  _numStencils = limitStencilTable->GetNumStencils();
+  if (_numStencils > 0) {
+    _sizes = createSSBO(limitStencilTable->GetSizes());
+    _offsets = createSSBO(limitStencilTable->GetOffsets());
+    _indices = createSSBO(limitStencilTable->GetControlIndices());
+    _weights = createSSBO(limitStencilTable->GetWeights());
+    _duWeights = createSSBO(limitStencilTable->GetDuWeights());
+    _dvWeights = createSSBO(limitStencilTable->GetDvWeights());
+    _duuWeights = createSSBO(limitStencilTable->GetDuuWeights());
+    _duvWeights = createSSBO(limitStencilTable->GetDuvWeights());
+    _dvvWeights = createSSBO(limitStencilTable->GetDvvWeights());
+  }
+  else {
+    _sizes = _offsets = _indices = _weights = 0;
+    _duWeights = _dvWeights = 0;
+    _duuWeights = _duvWeights = _dvvWeights = 0;
+  }
+}
+
+GLStencilTableSSBO::~GLStencilTableSSBO()
+{
+  if (_sizes)
+    glDeleteBuffers(1, &_sizes);
+  if (_offsets)
+    glDeleteBuffers(1, &_offsets);
+  if (_indices)
+    glDeleteBuffers(1, &_indices);
+  if (_weights)
+    glDeleteBuffers(1, &_weights);
+  if (_duWeights)
+    glDeleteBuffers(1, &_duWeights);
+  if (_dvWeights)
+    glDeleteBuffers(1, &_dvWeights);
+  if (_duuWeights)
+    glDeleteBuffers(1, &_duuWeights);
+  if (_duvWeights)
+    glDeleteBuffers(1, &_duvWeights);
+  if (_dvvWeights)
+    glDeleteBuffers(1, &_dvvWeights);
+}
+
+// ---------------------------------------------------------------------------
+
+GLComputeEvaluator::GLComputeEvaluator() : _workGroupSize(64), _patchArraysSSBO(0)
+{
+  memset((void *)&_stencilKernel, 0, sizeof(_stencilKernel));
+  memset((void *)&_patchKernel, 0, sizeof(_patchKernel));
+}
+
+GLComputeEvaluator::~GLComputeEvaluator()
+{
+  if (_patchArraysSSBO) {
+    glDeleteBuffers(1, &_patchArraysSSBO);
+  }
+}
+
+static GLuint compileKernel(BufferDescriptor const &srcDesc,
+                            BufferDescriptor const &dstDesc,
+                            BufferDescriptor const &duDesc,
+                            BufferDescriptor const &dvDesc,
+                            BufferDescriptor const &duuDesc,
+                            BufferDescriptor const &duvDesc,
+                            BufferDescriptor const &dvvDesc,
+                            const char *kernelDefine,
+                            int workGroupSize)
+{
+  GLuint program = glCreateProgram();
+
+  GLuint shader = glCreateShader(GL_COMPUTE_SHADER);
+
+  std::string patchBasisShaderSource =
+      OpenSubdiv::Osd::GLSLPatchShaderSource::GetPatchBasisShaderSource();
+  const char *patchBasisShaderSourceDefine = "#define OSD_PATCH_BASIS_GLSL\n";
+
+  std::ostringstream defines;
+  defines << "#define LENGTH " << srcDesc.length << "\n"
+          << "#define SRC_STRIDE " << srcDesc.stride << "\n"
+          << "#define DST_STRIDE " << dstDesc.stride << "\n"
+          << "#define WORK_GROUP_SIZE " << workGroupSize << "\n"
+          << kernelDefine << "\n"
+          << patchBasisShaderSourceDefine << "\n";
+
+  bool deriv1 = (duDesc.length > 0 || dvDesc.length > 0);
+  bool deriv2 = (duuDesc.length > 0 || duvDesc.length > 0 || dvvDesc.length > 0);
+  if (deriv1) {
+    defines << "#define OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES\n";
+  }
+  if (deriv2) {
+    defines << "#define OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES\n";
+  }
+
+  std::string defineStr = defines.str();
+
+  const char *shaderSources[4] = {"#version 430\n", 0, 0, 0};
+
+  shaderSources[1] = defineStr.c_str();
+  shaderSources[2] = patchBasisShaderSource.c_str();
+  shaderSources[3] = datatoc_glsl_compute_kernel_glsl;
+  glShaderSource(shader, 4, shaderSources, NULL);
+  glCompileShader(shader);
+  glAttachShader(program, shader);
+
+  GLint linked = 0;
+  glLinkProgram(program);
+  glGetProgramiv(program, GL_LINK_STATUS, &linked);
+
+  if (linked == GL_FALSE) {
+    char buffer[1024];
+    glGetShaderInfoLog(shader, 1024, NULL, buffer);
+    OpenSubdiv::Far::Error(OpenSubdiv::Far::FAR_RUNTIME_ERROR, buffer);
+
+    glGetProgramInfoLog(program, 1024, NULL, buffer);
+    OpenSubdiv::Far::Error(OpenSubdiv::Far::FAR_RUNTIME_ERROR, buffer);
+
+    glDeleteProgram(program);
+    return 0;
+  }
+
+  glDeleteShader(shader);
+
+  return program;
+}
+
+bool GLComputeEvaluator::Compile(BufferDescriptor const &srcDesc,
+                                 BufferDescriptor const &dstDesc,
+                                 BufferDescriptor const &duDesc,
+                                 BufferDescriptor const &dvDesc,
+                                 BufferDescriptor const &duuDesc,
+                                 BufferDescriptor const &duvDesc,
+                                 BufferDescriptor const &dvvDesc)
+{
+
+  // create a stencil kernel
+  if (!_stencilKernel.Compile(
+          srcDesc, dstDesc, duDesc, dvDesc, duuDesc, duvDesc, dvvDesc, _workGroupSize)) {
+    return false;
+  }
+
+  // create a patch kernel
+  if (!_patchKernel.Compile(
+          srcDesc, dstDesc, duDesc, dvDesc, duuDesc, duvDesc, dvvDesc, _workGroupSize)) {
+    return false;
+  }
+
+  // create a patch arrays buffer
+  if (!_patchArraysSSBO) {
+    glGenBuffers(1, &_patchArraysSSBO);
+  }
+
+  return true;
+}
+
+/* static */
+void GLComputeEvaluator::Synchronize(void * /*kernel*/)
+{
+  // XXX: this is currently just for the performance measuring purpose.
+  // need to be reimplemented by fence and sync.
+  glFinish();
+}
+
+int GLComputeEvaluator::GetDispatchSize(int count) const
+{
+  return (count + _workGroupSize - 1) / _workGroupSize;
+}
+
+void GLComputeEvaluator::DispatchCompute(int totalDispatchSize) const
+{
+  int maxWorkGroupCount[2] = {0, 0};
+
+  glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_COUNT, 0, &maxWorkGroupCount[0]);
+  glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_COUNT, 1, &maxWorkGroupCount[1]);
+
+  const GLuint maxResX = static_cast<GLuint>(maxWorkGroupCount[0]);
+
+  const int dispatchSize = GetDispatchSize(totalDispatchSize);
+  GLuint dispatchRX = static_cast<GLuint>(dispatchSize);
+  GLuint dispatchRY = 1u;
+  if (dispatchRX > maxResX) {
+    /* Since there are some limitations with regards to the maximum work group size (could be as
+     * low as 64k elements per call), we split the number elements into a "2d" number, with the
+     * final index being computed as `res_x + res_y * max_work_group_size`. Even with a maximum
+     * work group size of 64k, that still leaves us with roughly `64k * 64k = 4` billion elements
+     * total, which should be enough. If not, we could also use the 3rd dimension. */
+    /* TODO(fclem): We could dispatch fewer groups if we compute the prime factorization and
+     * get the smallest rect fitting the requirements. */
+    dispatchRX = dispatchRY = std::ceil(std::sqrt(dispatchSize));
+    /* Avoid a completely empty dispatch line caused by rounding. */
+    if ((dispatchRX * (dispatchRY - 1)) >= dispatchSize) {
+      dispatchRY -= 1;
+    }
+  }
+
+  /* X and Y dimensions may have different limits so the above computation may not be right, but
+   * even with the standard 64k minimum on all dimensions we still have a lot of room. Therefore,
+   * we presume it all fits. */
+  assert(dispatchRY < static_cast<GLuint>(maxWorkGroupCount[1]));
+
+  glDispatchCompute(dispatchRX, dispatchRY, 1);
+}
+
+bool GLComputeEvaluator::EvalStencils(GLuint srcBuffer,
+                                      BufferDescriptor const &srcDesc,
+                                      GLuint dstBuffer,
+                                      BufferDescriptor const &dstDesc,
+                                      GLuint duBuffer,
+                                      BufferDescriptor const &duDesc,
+                                      GLuint dvBuffer,
+                                      BufferDescriptor const &dvDesc,
+                                      GLuint sizesBuffer,
+                                      GLuint offsetsBuffer,
+                                      GLuint indicesBuffer,
+                                      GLuint weightsBuffer,
+                                      GLuint duWeightsBuffer,
+                                      GLuint dvWeightsBuffer,
+                                      int start,
+                                      int end) const
+{
+
+  return EvalStencils(srcBuffer,
+                      srcDesc,
+                      dstBuffer,
+                      dstDesc,
+                      duBuffer,
+                      duDesc,
+                      dvBuffer,
+                      dvDesc,
+                      0,
+                      BufferDescriptor(),
+                      0,
+                      BufferDescriptor(),
+                      0,
+                      BufferDescriptor(),
+                      sizesBuffer,
+                      offsetsBuffer,
+                      indicesBuffer,
+                      weightsBuffer,
+                      duWeightsBuffer,
+                      dvWeightsBuffer,
+                      0,
+                      0,
+                      0,
+                      start,
+                      end);
+}
+
+bool GLComputeEvaluator::EvalStencils(GLuint srcBuffer,
+                                      BufferDescriptor const &srcDesc,
+                                      GLuint dstBuffer,
+                                      BufferDescriptor const &dstDesc,
+                                      GLuint duBuffer,
+                                      BufferDescriptor const &duDesc,
+                                      GLuint dvBuffer,
+                                      BufferDescriptor const &dvDesc,
+                                      GLuint duuBuffer,
+                                      BufferDescriptor const &duuDesc,
+                                      GLuint duvBuffer,
+                                      BufferDescriptor const &duvDesc,
+                                      GLuint dvvBuffer,
+                                      BufferDescriptor const &dvvDesc,
+                                      GLuint sizesBuffer,
+                                      GLuint offsetsBuffer,
+                                      GLuint indicesBuffer,
+                                      GLuint weightsBuffer,
+                                      GLuint duWeightsBuffer,
+                                      GLuint dvWeightsBuffer,
+                                      GLuint duuWeightsBuffer,
+                                      GLuint duvWeightsBuffer,
+                                      GLuint dvvWeightsBuffer,
+                                      int start,
+                                      int end) const
+{
+
+  if (!_stencilKernel.program)
+    return false;
+  int count = end - start;
+  if (count <= 0) {
+    return true;
+  }
+
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, srcBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, dstBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, duBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, dvBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 10, duuBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 11, duvBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 12, dvvBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 4, sizesBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 5, offsetsBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 6, indicesBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 7, weightsBuffer);
+  if (duWeightsBuffer)
+    glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 8, duWeightsBuffer);
+  if (dvWeightsBuffer)
+    glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 9, dvWeightsBuffer);
+  if (duuWeightsBuffer)
+    glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 13, duuWeightsBuffer);
+  if (duvWeightsBuffer)
+    glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 14, duvWeightsBuffer);
+  if (dvvWeightsBuffer)
+    glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 15, dvvWeightsBuffer);
+
+  glUseProgram(_stencilKernel.program);
+
+  glUniform1i(_stencilKernel.uniformStart, start);
+  glUniform1i(_stencilKernel.uniformEnd, end);
+  glUniform1i(_stencilKernel.uniformSrcOffset, srcDesc.offset);
+  glUniform1i(_stencilKernel.uniformDstOffset, dstDesc.offset);
+  if (_stencilKernel.uniformDuDesc > 0) {
+    glUniform3i(_stencilKernel.uniformDuDesc, duDesc.offset, duDesc.length, duDesc.stride);
+  }
+  if (_stencilKernel.uniformDvDesc > 0) {
+    glUniform3i(_stencilKernel.uniformDvDesc, dvDesc.offset, dvDesc.length, dvDesc.stride);
+  }
+  if (_stencilKernel.uniformDuuDesc > 0) {
+    glUniform3i(_stencilKernel.uniformDuuDesc, duuDesc.offset, duuDesc.length, duuDesc.stride);
+  }
+  if (_stencilKernel.uniformDuvDesc > 0) {
+    glUniform3i(_stencilKernel.uniformDuvDesc, duvDesc.offset, duvDesc.length, duvDesc.stride);
+  }
+  if (_stencilKernel.uniformDvvDesc > 0) {
+    glUniform3i(_stencilKernel.uniformDvvDesc, dvvDesc.offset, dvvDesc.length, dvvDesc.stride);
+  }
+
+  DispatchCompute(count);
+
+  glUseProgram(0);
+
+  glMemoryBarrier(GL_TEXTURE_FETCH_BARRIER_BIT);
+  for (int i = 0; i < 16; ++i) {
+    glBindBufferBase(GL_SHADER_STORAGE_BUFFER, i, 0);
+  }
+
+  return true;
+}
+
+bool GLComputeEvaluator::EvalPatches(GLuint srcBuffer,
+                                     BufferDescriptor const &srcDesc,
+                                     GLuint dstBuffer,
+                                     BufferDescriptor const &dstDesc,
+                                     GLuint duBuffer,
+                                     BufferDescriptor const &duDesc,
+                                     GLuint dvBuffer,
+                                     BufferDescriptor const &dvDesc,
+                                     int numPatchCoords,
+                                     GLuint patchCoordsBuffer,
+                                     const PatchArrayVector &patchArrays,
+                                     GLuint patchIndexBuffer,
+                                     GLuint patchParamsBuffer) const
+{
+
+  return EvalPatches(srcBuffer,
+                     srcDesc,
+                     dstBuffer,
+                     dstDesc,
+                     duBuffer,
+                     duDesc,
+                     dvBuffer,
+                     dvDesc,
+                     0,
+                     BufferDescriptor(),
+                     0,
+                     BufferDescriptor(),
+                     0,
+                     BufferDescriptor(),
+                     numPatchCoords,
+                     patchCoordsBuffer,
+                     patchArrays,
+                     patchIndexBuffer,
+                     patchParamsBuffer);
+}
+
+bool GLComputeEvaluator::EvalPatches(GLuint srcBuffer,
+                                     BufferDescriptor const &srcDesc,
+                                     GLuint dstBuffer,
+                                     BufferDescriptor const &dstDesc,
+                                     GLuint duBuffer,
+                                     BufferDescriptor const &duDesc,
+                                     GLuint dvBuffer,
+                                     BufferDescriptor const &dvDesc,
+                                     GLuint duuBuffer,
+                                     BufferDescriptor const &duuDesc,
+                                     GLuint duvBuffer,
+                                     BufferDescriptor const &duvDesc,
+                                     GLuint dvvBuffer,
+                                     BufferDescriptor const &dvvDesc,
+                                     int numPatchCoords,
+                                     GLuint patchCoordsBuffer,
+                                     const PatchArrayVector &patchArrays,
+                                     GLuint patchIndexBuffer,
+                                     GLuint patchParamsBuffer) const
+{
+
+  if (!_patchKernel.program)
+    return false;
+
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, srcBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, dstBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, duBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, dvBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 10, duuBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 11, duvBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 12, dvvBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 5, patchCoordsBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 6, patchIndexBuffer);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 7, patchParamsBuffer);
+
+  glUseProgram(_patchKernel.program);
+
+  glUniform1i(_patchKernel.uniformSrcOffset, srcDesc.offset);
+  glUniform1i(_patchKernel.uniformDstOffset, dstDesc.offset);
+
+  int patchArraySize = sizeof(PatchArray);
+  glBindBuffer(GL_SHADER_STORAGE_BUFFER, _patchArraysSSBO);
+  glBufferData(
+      GL_SHADER_STORAGE_BUFFER, patchArrays.size() * patchArraySize, NULL, GL_STATIC_DRAW);
+  for (int i = 0; i < (int)patchArrays.size(); ++i) {
+    glBufferSubData(
+        GL_SHADER_STORAGE_BUFFER, i * patchArraySize, sizeof(PatchArray), &patchArrays[i]);
+  }
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 4, _patchArraysSSBO);
+
+  if (_patchKernel.uniformDuDesc > 0) {
+    glUniform3i(_patchKernel.uniformDuDesc, duDesc.offset, duDesc.length, duDesc.stride);
+  }
+  if (_patchKernel.uniformDvDesc > 0) {
+    glUniform3i(_patchKernel.uniformDvDesc, dvDesc.offset, dvDesc.length, dvDesc.stride);
+  }
+  if (_patchKernel.uniformDuuDesc > 0) {
+    glUniform3i(_patchKernel.uniformDuuDesc, duuDesc.offset, duuDesc.length, duuDesc.stride);
+  }
+  if (_patchKernel.uniformDuvDesc > 0) {
+    glUniform3i(_patchKernel.uniformDuvDesc, duvDesc.offset, duvDesc.length, duvDesc.stride);
+  }
+  if (_patchKernel.uniformDvvDesc > 0) {
+    glUniform3i(_patchKernel.uniformDvvDesc, dvvDesc.offset, dvvDesc.length, dvvDesc.stride);
+  }
+
+  DispatchCompute(numPatchCoords);
+
+  glUseProgram(0);
+
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, 0);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, 0);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, 0);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, 0);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 4, 0);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 5, 0);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 6, 0);
+
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 10, 0);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 11, 0);
+  glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 12, 0);
+
+  return true;
+}
+// ---------------------------------------------------------------------------
+
+GLComputeEvaluator::_StencilKernel::_StencilKernel() : program(0)
+{
+}
+GLComputeEvaluator::_StencilKernel::~_StencilKernel()
+{
+  if (program) {
+    glDeleteProgram(program);
+  }
+}
+
+bool GLComputeEvaluator::_StencilKernel::Compile(BufferDescriptor const &srcDesc,
+                                                 BufferDescriptor const &dstDesc,
+                                                 BufferDescriptor const &duDesc,
+                                                 BufferDescriptor const &dvDesc,
+                                                 BufferDescriptor const &duuDesc,
+                                                 BufferDescriptor const &duvDesc,
+                                                 BufferDescriptor const &dvvDesc,
+                                                 int workGroupSize)
+{
+  // create stencil kernel
+  if (program) {
+    glDeleteProgram(program);
+  }
+
+  const char *kernelDefine = "#define OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_STENCILS\n";
+
+  program = compileKernel(
+      srcDesc, dstDesc, duDesc, dvDesc, duuDesc, duvDesc, dvvDesc, kernelDefine, workGroupSize);
+  if (program == 0)
+    return false;
+
+  // cache uniform locations (TODO: use uniform block)
+  uniformStart = glGetUniformLocation(program, "batchStart");
+  uniformEnd = glGetUniformLocation(program, "batchEnd");
+  uniformSrcOffset = glGetUniformLocation(program, "srcOffset");
+  uniformDstOffset = glGetUniformLocation(program, "dstOffset");
+  uniformDuDesc = glGetUniformLocation(program, "duDesc");
+  uniformDvDesc = glGetUniformLocation(program, "dvDesc");
+  uniformDuuDesc = glGetUniformLocation(program, "duuDesc");
+  uniformDuvDesc = glGetUniformLocation(program, "duvDesc");
+  uniformDvvDesc = glGetUniformLocation(program, "dvvDesc");
+
+  return true;
+}
+
+// ---------------------------------------------------------------------------
+
+GLComputeEvaluator::_PatchKernel::_PatchKernel() : program(0)
+{
+}
+GLComputeEvaluator::_PatchKernel::~_PatchKernel()
+{
+  if (program) {
+    glDeleteProgram(program);
+  }
+}
+
+bool GLComputeEvaluator::_PatchKernel::Compile(BufferDescriptor const &srcDesc,
+                                               BufferDescriptor const &dstDesc,
+                                               BufferDescriptor const &duDesc,
+                                               BufferDescriptor const &dvDesc,
+                                               BufferDescriptor const &duuDesc,
+                                               BufferDescriptor const &duvDesc,
+                                               BufferDescriptor const &dvvDesc,
+                                               int workGroupSize)
+{
+  // create stencil kernel
+  if (program) {
+    glDeleteProgram(program);
+  }
+
+  const char *kernelDefine = "#define OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_PATCHES\n";
+
+  program = compileKernel(
+      srcDesc, dstDesc, duDesc, dvDesc, duuDesc, duvDesc, dvvDesc, kernelDefine, workGroupSize);
+  if (program == 0)
+    return false;
+
+  // cache uniform locations
+  uniformSrcOffset = glGetUniformLocation(program, "srcOffset");
+  uniformDstOffset = glGetUniformLocation(program, "dstOffset");
+  uniformPatchArray = glGetUniformLocation(program, "patchArray");
+  uniformDuDesc = glGetUniformLocation(program, "duDesc");
+  uniformDvDesc = glGetUniformLocation(program, "dvDesc");
+  uniformDuuDesc = glGetUniformLocation(program, "duuDesc");
+  uniformDuvDesc = glGetUniformLocation(program, "duvDesc");
+  uniformDvvDesc = glGetUniformLocation(program, "dvvDesc");
+
+  return true;
+}
+
+}  // namespace opensubdiv
+}  // namespace blender

intern/opensubdiv/internal/evaluator/shaders/glsl_compute_kernel.glsl

@@ -0,0 +1,316 @@
+//
+//   Copyright 2013 Pixar
+//
+//   Licensed under the Apache License, Version 2.0 (the "Apache License")
+//   with the following modification; you may not use this file except in
+//   compliance with the Apache License and the following modification to it:
+//   Section 6. Trademarks. is deleted and replaced with:
+//
+//   6. Trademarks. This License does not grant permission to use the trade
+//      names, trademarks, service marks, or product names of the Licensor
+//      and its affiliates, except as required to comply with Section 4(c) of
+//      the License and to reproduce the content of the NOTICE file.
+//
+//   You may obtain a copy of the Apache License at
+//
+//       http://www.apache.org/licenses/LICENSE-2.0
+//
+//   Unless required by applicable law or agreed to in writing, software
+//   distributed under the Apache License with the above modification is
+//   distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+//   KIND, either express or implied. See the Apache License for the specific
+//   language governing permissions and limitations under the Apache License.
+//
+
+//------------------------------------------------------------------------------
+
+
+layout(local_size_x=WORK_GROUP_SIZE, local_size_y=1, local_size_z=1) in;
+layout(std430) buffer;
+
+// source and destination buffers
+
+uniform int srcOffset = 0;
+uniform int dstOffset = 0;
+layout(binding=0) buffer src_buffer      { float    srcVertexBuffer[]; };
+layout(binding=1) buffer dst_buffer      { float    dstVertexBuffer[]; };
+
+// derivative buffers (if needed)
+
+#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES)
+uniform ivec3 duDesc;
+uniform ivec3 dvDesc;
+layout(binding=2) buffer du_buffer   { float duBuffer[]; };
+layout(binding=3) buffer dv_buffer   { float dvBuffer[]; };
+#endif
+
+#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES)
+uniform ivec3 duuDesc;
+uniform ivec3 duvDesc;
+uniform ivec3 dvvDesc;
+layout(binding=10) buffer duu_buffer   { float duuBuffer[]; };
+layout(binding=11) buffer duv_buffer   { float duvBuffer[]; };
+layout(binding=12) buffer dvv_buffer   { float dvvBuffer[]; };
+#endif
+
+// stencil buffers
+
+#if defined(OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_STENCILS)
+
+uniform int batchStart = 0;
+uniform int batchEnd = 0;
+layout(binding=4) buffer stencilSizes    { int      _sizes[];   };
+layout(binding=5) buffer stencilOffsets  { int      _offsets[]; };
+layout(binding=6) buffer stencilIndices  { int      _indices[]; };
+layout(binding=7) buffer stencilWeights  { float    _weights[]; };
+
+#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES)
+layout(binding=8) buffer stencilDuWeights { float  _duWeights[]; };
+layout(binding=9) buffer stencilDvWeights { float  _dvWeights[]; };
+#endif
+
+#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES)
+layout(binding=13) buffer stencilDuuWeights { float  _duuWeights[]; };
+layout(binding=14) buffer stencilDuvWeights { float  _duvWeights[]; };
+layout(binding=15) buffer stencilDvvWeights { float  _dvvWeights[]; };
+#endif
+
+uint getGlobalInvocationIndex()
+{
+  uint invocations_per_row = gl_WorkGroupSize.x * gl_NumWorkGroups.x;
+  return gl_GlobalInvocationID.x + gl_GlobalInvocationID.y * invocations_per_row;
+}
+
+#endif
+
+// patch buffers
+
+#if defined(OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_PATCHES)
+
+layout(binding=4) buffer patchArray_buffer { OsdPatchArray patchArrayBuffer[]; };
+layout(binding=5) buffer patchCoord_buffer { OsdPatchCoord patchCoords[]; };
+layout(binding=6) buffer patchIndex_buffer { int patchIndexBuffer[]; };
+layout(binding=7) buffer patchParam_buffer { OsdPatchParam patchParamBuffer[]; };
+
+OsdPatchCoord GetPatchCoord(int coordIndex)
+{
+    return patchCoords[coordIndex];
+}
+
+OsdPatchArray GetPatchArray(int arrayIndex)
+{
+    return patchArrayBuffer[arrayIndex];
+}
+
+OsdPatchParam GetPatchParam(int patchIndex)
+{
+    return patchParamBuffer[patchIndex];
+}
+
+#endif
+
+//------------------------------------------------------------------------------
+
+struct Vertex {
+    float vertexData[LENGTH];
+};
+
+void clear(out Vertex v) {
+    for (int i = 0; i < LENGTH; ++i) {
+        v.vertexData[i] = 0;
+    }
+}
+
+Vertex readVertex(int index) {
+    Vertex v;
+    int vertexIndex = srcOffset + index * SRC_STRIDE;
+    for (int i = 0; i < LENGTH; ++i) {
+        v.vertexData[i] = srcVertexBuffer[vertexIndex + i];
+    }
+    return v;
+}
+
+void writeVertex(int index, Vertex v) {
+    int vertexIndex = dstOffset + index * DST_STRIDE;
+    for (int i = 0; i < LENGTH; ++i) {
+        dstVertexBuffer[vertexIndex + i] = v.vertexData[i];
+    }
+}
+
+void addWithWeight(inout Vertex v, const Vertex src, float weight) {
+    for (int i = 0; i < LENGTH; ++i) {
+        v.vertexData[i] += weight * src.vertexData[i];
+    }
+}
+
+#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES)
+void writeDu(int index, Vertex du) {
+    int duIndex = duDesc.x + index * duDesc.z;
+    for (int i = 0; i < LENGTH; ++i) {
+        duBuffer[duIndex + i] = du.vertexData[i];
+    }
+}
+
+void writeDv(int index, Vertex dv) {
+    int dvIndex = dvDesc.x + index * dvDesc.z;
+    for (int i = 0; i < LENGTH; ++i) {
+        dvBuffer[dvIndex + i] = dv.vertexData[i];
+    }
+}
+#endif
+
+#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES)
+void writeDuu(int index, Vertex duu) {
+    int duuIndex = duuDesc.x + index * duuDesc.z;
+    for (int i = 0; i < LENGTH; ++i) {
+        duuBuffer[duuIndex + i] = duu.vertexData[i];
+    }
+}
+
+void writeDuv(int index, Vertex duv) {
+    int duvIndex = duvDesc.x + index * duvDesc.z;
+    for (int i = 0; i < LENGTH; ++i) {
+        duvBuffer[duvIndex + i] = duv.vertexData[i];
+    }
+}
+
+void writeDvv(int index, Vertex dvv) {
+    int dvvIndex = dvvDesc.x + index * dvvDesc.z;
+    for (int i = 0; i < LENGTH; ++i) {
+        dvvBuffer[dvvIndex + i] = dvv.vertexData[i];
+    }
+}
+#endif
+
+//------------------------------------------------------------------------------
+#if defined(OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_STENCILS)
+
+void main() {
+    int current = int(getGlobalInvocationIndex()) + batchStart;
+
+    if (current>=batchEnd) {
+        return;
+    }
+
+    Vertex dst;
+    clear(dst);
+
+    int offset = _offsets[current],
+        size   = _sizes[current];
+
+    for (int stencil = 0; stencil < size; ++stencil) {
+        int vindex = offset + stencil;
+        addWithWeight(
+            dst, readVertex(_indices[vindex]), _weights[vindex]);
+    }
+
+    writeVertex(current, dst);
+
+#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES)
+    Vertex du, dv;
+    clear(du);
+    clear(dv);
+    for (int i=0; i<size; ++i) {
+        // expects the compiler optimizes readVertex out here.
+        Vertex src = readVertex(_indices[offset+i]);
+        addWithWeight(du, src, _duWeights[offset+i]);
+        addWithWeight(dv, src, _dvWeights[offset+i]);
+    }
+
+    if (duDesc.y > 0) { // length
+        writeDu(current, du);
+    }
+    if (dvDesc.y > 0) {
+        writeDv(current, dv);
+    }
+#endif
+#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES)
+    Vertex duu, duv, dvv;
+    clear(duu);
+    clear(duv);
+    clear(dvv);
+    for (int i=0; i<size; ++i) {
+        // expects the compiler optimizes readVertex out here.
+        Vertex src = readVertex(_indices[offset+i]);
+        addWithWeight(duu, src, _duuWeights[offset+i]);
+        addWithWeight(duv, src, _duvWeights[offset+i]);
+        addWithWeight(dvv, src, _dvvWeights[offset+i]);
+    }
+
+    if (duuDesc.y > 0) { // length
+        writeDuu(current, duu);
+    }
+    if (duvDesc.y > 0) {
+        writeDuv(current, duv);
+    }
+    if (dvvDesc.y > 0) {
+        writeDvv(current, dvv);
+    }
+#endif
+}
+
+#endif
+
+//------------------------------------------------------------------------------
+#if defined(OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_PATCHES)
+
+// PERFORMANCE: stride could be constant, but not as significant as length
+
+void main() {
+
+    int current = int(gl_GlobalInvocationID.x);
+
+    OsdPatchCoord coord = GetPatchCoord(current);
+    OsdPatchArray array = GetPatchArray(coord.arrayIndex);
+    OsdPatchParam param = GetPatchParam(coord.patchIndex);
+
+    int patchType = OsdPatchParamIsRegular(param) ? array.regDesc : array.desc;
+
+    float wP[20], wDu[20], wDv[20], wDuu[20], wDuv[20], wDvv[20];
+    int nPoints = OsdEvaluatePatchBasis(patchType, param,
+        coord.s, coord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
+
+    Vertex dst, du, dv, duu, duv, dvv;
+    clear(dst);
+    clear(du);
+    clear(dv);
+    clear(duu);
+    clear(duv);
+    clear(dvv);
+
+    int indexBase = array.indexBase + array.stride *
+                (coord.patchIndex - array.primitiveIdBase);
+
+    for (int cv = 0; cv < nPoints; ++cv) {
+        int index = patchIndexBuffer[indexBase + cv];
+        addWithWeight(dst, readVertex(index), wP[cv]);
+        addWithWeight(du, readVertex(index), wDu[cv]);
+        addWithWeight(dv, readVertex(index), wDv[cv]);
+        addWithWeight(duu, readVertex(index), wDuu[cv]);
+        addWithWeight(duv, readVertex(index), wDuv[cv]);
+        addWithWeight(dvv, readVertex(index), wDvv[cv]);
+    }
+    writeVertex(current, dst);
+
+#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES)
+    if (duDesc.y > 0) { // length
+        writeDu(current, du);
+    }
+    if (dvDesc.y > 0) {
+        writeDv(current, dv);
+    }
+#endif
+#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES)
+    if (duuDesc.y > 0) { // length
+        writeDuu(current, duu);
+    }
+    if (duvDesc.y > 0) { // length
+        writeDuv(current, duv);
+    }
+    if (dvvDesc.y > 0) {
+        writeDvv(current, dvv);
+    }
+#endif
+}
+
+#endif

release/datafiles/colormanagement/config.ocio

@@ -9,7 +9,7 @@
 #
 # See ocio-license.txt for details.
 
-ocio_profile_version: 1
+ocio_profile_version: 2
 
 search_path: "luts:filmic"
 strictparsing: true
@@ -100,8 +100,7 @@ colorspaces:
     from_reference: !<GroupTransform>
       children:
         - !<FileTransform> {src: srgb_to_xyz.spimtx, interpolation: linear}
-        - !<FileTransform> {src: xyz_D65_to_E.spimtx, interpolation: linear}
-        - !<FileTransform> {src: xyz_to_aces.spimtx, interpolation: linear}
+        - !<BuiltinTransform> {style: "UTILITY - ACES-AP0_to_CIE-XYZ-D65_BFD", direction: inverse}
 
   - !<ColorSpace>
     name: nuke_rec709

release/scripts/modules/bl_i18n_utils/settings.py

@@ -356,6 +356,7 @@ WARN_MSGID_NOT_CAPITALIZED_ALLOWED = {
     "y",
     "y = (Ax + B)",
     # Sub-strings.
+    "and AMD Radeon Pro 21.Q4 driver or newer",
     "available with",
     "brown fox",
     "can't save image while rendering",
@@ -378,6 +379,7 @@ WARN_MSGID_NOT_CAPITALIZED_ALLOWED = {
     "image path can't be written to",
     "in memory to enable editing!",
     "insufficient content",
+    "into",
     "jumps over",
     "left",
     "local",
@@ -387,6 +389,7 @@ WARN_MSGID_NOT_CAPITALIZED_ALLOWED = {
     "performance impact!",
     "right",
     "the lazy dog",
+    "to the top level of the tree",
     "unable to load movie clip",
     "unable to load text",
     "unable to open the file",

release/scripts/modules/bl_i18n_utils/utils_spell_check.py

@@ -76,10 +76,12 @@ class SpellChecker:
         "tangency",
         "vertices",
         "wasn",  # wasn't
+        "zig", "zag",
 
         # Brands etc.
         "htc",
         "huawei",
+        "radeon",
         "vive",
         "xbox",
 
@@ -136,6 +138,7 @@ class SpellChecker:
         "filename", "filenames",
         "filepath", "filepaths",
         "forcefield", "forcefields",
+        "framerange",
         "fulldome", "fulldomes",
         "fullscreen",
         "gamepad",
@@ -498,6 +501,7 @@ class SpellChecker:
         "framerate",
         "gimbal",
         "grayscale",
+        "icosahedron",
         "icosphere",
         "inpaint",
         "kerning",
@@ -556,6 +560,7 @@ class SpellChecker:
         "bspline",
         "bweight",
         "colorband",
+        "crazyspace",
         "datablock", "datablocks",
         "despeckle",
         "depsgraph",
@@ -730,6 +735,7 @@ class SpellChecker:
         "precisa",
         "px",
         "qmc",
+        "rdna",
         "rdp",
         "rgb", "rgba",
         "rhs",

release/scripts/presets/keyconfig/Blender.py

@@ -318,7 +318,8 @@ def load():
             use_v3d_tab_menu=kc_prefs.use_v3d_tab_menu,
             use_v3d_shade_ex_pie=kc_prefs.use_v3d_shade_ex_pie,
             use_gizmo_drag=(is_select_left and kc_prefs.gizmo_action == 'DRAG'),
-            use_fallback_tool=(True if is_select_left else (kc_prefs.rmb_action == 'FALLBACK_TOOL')),
+            use_fallback_tool=True,
+            use_fallback_tool_rmb=(False if is_select_left else kc_prefs.rmb_action == 'FALLBACK_TOOL'),
             use_alt_tool_or_cursor=(
                 (not use_mouse_emulate_3_button) and
                 (kc_prefs.use_alt_tool if is_select_left else kc_prefs.use_alt_cursor)

release/scripts/presets/keyconfig/keymap_data/blender_default.py

@@ -16,6 +16,16 @@
 #
 # ##### END GPL LICENSE BLOCK #####
 
+# ------------------------------------------------------------------------------
+# Developer Notes
+#
+# - This script should run without Blender (no references to the `bpy` module for example).
+# - All configuration must be passed into the `generate_keymaps` function (via `Params`).
+# - Supporting some combinations of options is becoming increasingly complex,
+#   especially `Params.select_mouse` & `Params.use_fallback_tool_rmb`.
+#   To ensure changes don't unintentionally break other configurations, see:
+#   `source/tools/utils/blender_keyconfig_export_permutations.py --help`
+#
 
 # ------------------------------------------------------------------------------
 # Configurable Parameters
@@ -48,6 +58,8 @@ class Params:
         "use_gizmo_drag",
         # Use the fallback tool instead of tweak for RMB select.
         "use_fallback_tool",
+        # Only set for RMB select.
+        "use_fallback_tool_rmb",
         # Use pie menu for tab by default (swap 'Tab/Ctrl-Tab').
         "use_v3d_tab_menu",
         # Use extended pie menu for shading.
@@ -65,15 +77,16 @@ class Params:
         "v3d_tilde_action",
         # Alt-MMB axis switching 'RELATIVE' or 'ABSOLUTE' axis switching.
         "v3d_alt_mmb_drag_action",
-
+        # File selector actions on single click.
         "use_file_single_click",
+
         # Convenience variables:
         # (derived from other settings).
         #
-        # This case needs to be checked often,
-        # Shorthand for: `(params.use_fallback_tool if params.select_mouse == 'RIGHTMOUSE' else False)`.
-        "use_fallback_tool_rmb",
-        # Shorthand for: `('CLICK' if params.use_fallback_tool_rmb else params.select_mouse_value)`.
+        # The fallback tool is activated on the same button as selection.
+        # Shorthand for: `(True if (select_mouse == 'LEFT') else self.use_fallback_tool_rmb)`
+        "use_fallback_tool_select_mouse",
+        # Shorthand for: `('CLICK' if self.use_fallback_tool_rmb else self.select_mouse_value)`.
         "select_mouse_value_fallback",
         # Shorthand for: `{"type": params.select_tweak, "value": 'ANY'}`.
         "select_tweak_event",
@@ -103,6 +116,7 @@ class Params:
             use_select_all_toggle=False,
             use_gizmo_drag=True,
             use_fallback_tool=False,
+            use_fallback_tool_rmb=False,
             use_v3d_tab_menu=False,
             use_v3d_shade_ex_pie=False,
             use_v3d_mmb_pan=False,
@@ -146,7 +160,6 @@ class Params:
                 self.cursor_set_event = {"type": 'LEFTMOUSE', "value": 'CLICK'}
                 self.cursor_tweak_event = None
 
-            self.use_fallback_tool = use_fallback_tool
             self.tool_modifier = {}
         else:
             # Left mouse select uses Click event for selection. This is a little
@@ -169,7 +182,6 @@ class Params:
 
             self.cursor_set_event = {"type": 'RIGHTMOUSE', "value": 'PRESS', "shift": True}
             self.cursor_tweak_event = {"type": 'EVT_TWEAK_R', "value": 'ANY', "shift": True}
-            self.use_fallback_tool = True
 
             # Use the "tool" functionality for LMB select.
             if use_alt_tool_or_cursor:
@@ -197,8 +209,11 @@ class Params:
 
         self.use_file_single_click = use_file_single_click
 
+        self.use_fallback_tool = use_fallback_tool
+        self.use_fallback_tool_rmb = use_fallback_tool_rmb
+
         # Convenience variables:
-        self.use_fallback_tool_rmb = self.use_fallback_tool if select_mouse == 'RIGHT' else False
+        self.use_fallback_tool_select_mouse = True if (select_mouse == 'LEFT') else self.use_fallback_tool_rmb
         self.select_mouse_value_fallback = 'CLICK' if self.use_fallback_tool_rmb else self.select_mouse_value
         self.select_tweak_event = {"type": self.select_tweak, "value": 'ANY'}
         self.pie_value = 'CLICK_DRAG' if use_pie_click_drag else 'PRESS'
@@ -1149,11 +1164,7 @@ def km_uv_editor(params):
     items.extend([
         # Selection modes.
         *_template_items_uv_select_mode(params),
-        *_template_uv_select(
-            type=params.select_mouse,
-            value=('CLICK' if params.use_fallback_tool_rmb else params.select_mouse_value),
-            legacy=params.legacy,
-        ),
+        *_template_uv_select(type=params.select_mouse, value=params.select_mouse_value_fallback, legacy=params.legacy),
         ("uv.mark_seam", {"type": 'E', "value": 'PRESS', "ctrl": True}, None),
         ("uv.select_loop",
          {"type": params.select_mouse, "value": params.select_mouse_value, "alt": True}, None),
@@ -2018,8 +2029,10 @@ def km_node_editor(params):
         ("node.link_make", {"type": 'F', "value": 'PRESS', "shift": True},
          {"properties": [("replace", True)]}),
         op_menu("NODE_MT_add", {"type": 'A', "value": 'PRESS', "shift": True}),
-        ("node.duplicate_move", {"type": 'D', "value": 'PRESS', "shift": True}, None),
-        ("node.duplicate_move_keep_inputs", {"type": 'D', "value": 'PRESS', "shift": True, "ctrl": True}, None),
+        ("node.duplicate_move", {"type": 'D', "value": 'PRESS', "shift": True},
+         {"properties": [("NODE_OT_translate_attach", [("TRANSFORM_OT_translate", [("view2d_edge_pan", True)])])]}),
+        ("node.duplicate_move_keep_inputs", {"type": 'D', "value": 'PRESS', "shift": True, "ctrl": True},
+         {"properties": [("NODE_OT_translate_attach", [("TRANSFORM_OT_translate", [("view2d_edge_pan", True)])])]}),
         ("node.parent_set", {"type": 'P', "value": 'PRESS', "ctrl": True}, None),
         ("node.detach", {"type": 'P', "value": 'PRESS', "alt": True}, None),
         ("node.join", {"type": 'J', "value": 'PRESS', "ctrl": True}, None),
@@ -4655,7 +4668,9 @@ def _template_paint_radial_control(paint, rotation=False, secondary_rotation=Fal
     return items
 
 
-def _template_view3d_select(*, type, value, legacy):
+def _template_view3d_select(*, type, value, legacy, exclude_mod=None):
+    # NOTE: `exclude_mod` is needed since we don't want this tool to exclude Control-RMB actions when this is used
+    # as a tool key-map with RMB-select and `use_fallback_tool_rmb` is enabled. See T92467.
     return [(
         "view3d.select",
         {"type": type, "value": value, **{m: True for m in mods}},
@@ -4669,7 +4684,7 @@ def _template_view3d_select(*, type, value, legacy):
         (("center", "enumerate"), ("ctrl", "alt")),
         (("toggle", "enumerate"), ("shift", "alt")),
         (("toggle", "center", "enumerate"), ("shift", "ctrl", "alt")),
-    )]
+    ) if exclude_mod is None or exclude_mod not in mods]
 
 
 def _template_view3d_gpencil_select(*, type, value, legacy, use_select_mouse=True):
@@ -6283,7 +6298,8 @@ def km_image_editor_tool_uv_select_box(params, *, fallback):
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
                 "uv.select_box",
                 # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if (fallback and params.use_fallback_tool_select_mouse) else
+                   params.tool_tweak_event))),
         ]},
     )
 
@@ -6295,7 +6311,8 @@ def km_image_editor_tool_uv_select_circle(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
                 "uv.select_circle",
-                **(params.select_tweak_event if fallback else {"type": params.tool_mouse, "value": 'PRESS'}),
+                **(params.select_tweak_event if (fallback and params.use_fallback_tool_select_mouse) else
+                   {"type": params.tool_mouse, "value": 'PRESS'}),
                 properties=[("wait_for_input", False)])),
             # No selection fallback since this operates on press.
         ]},
@@ -6310,7 +6327,8 @@ def km_image_editor_tool_uv_select_lasso(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
                 "uv.select_lasso",
-                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if (fallback and params.use_fallback_tool_select_mouse) else
+                   params.tool_tweak_event))),
         ]},
     )
 
@@ -6402,7 +6420,8 @@ def km_node_editor_tool_select_box(params, *, fallback):
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
                 "node.select_box",
                 # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-                **(params.select_tweak_event if fallback else params.tool_tweak_event),
+                **(params.select_tweak_event if (fallback and params.use_fallback_tool_select_mouse) else
+                   params.tool_tweak_event),
                 properties=[("tweak", True)],
             )),
         ]},
@@ -6415,7 +6434,9 @@ def km_node_editor_tool_select_lasso(params, *, fallback):
         {"space_type": 'NODE_EDITOR', "region_type": 'WINDOW'},
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
-                "node.select_lasso", **(params.select_tweak_event if fallback else params.tool_tweak_event),
+                "node.select_lasso",
+                **(params.select_tweak_event if (fallback and params.use_fallback_tool_select_mouse) else
+                   params.tool_tweak_event),
                 properties=[("tweak", True)]))
         ]},
     )
@@ -6430,7 +6451,7 @@ def km_node_editor_tool_select_circle(params, *, fallback):
                 "node.select_circle",
                 # Why circle select should be used on tweak?
                 # So that RMB or Shift-RMB is still able to set an element as active.
-                type=params.select_tweak if fallback else params.tool_mouse,
+                type=params.select_tweak if (fallback and params.use_fallback_tool_select_mouse) else params.tool_mouse,
                 value='ANY' if fallback else 'PRESS',
                 properties=[("wait_for_input", False)])),
         ]},
@@ -6471,7 +6492,7 @@ def km_3d_view_tool_select(params, *, fallback):
             *([] if (fallback and (params.select_mouse == 'RIGHTMOUSE')) else _template_items_tool_select(
                 params, "view3d.select", "view3d.cursor3d", extend="toggle")),
             *([] if (not params.use_fallback_tool_rmb) else _template_view3d_select(
-                type=params.select_mouse, value=params.select_mouse_value, legacy=params.legacy)),
+                type=params.select_mouse, value=params.select_mouse_value, legacy=params.legacy, exclude_mod="ctrl")),
         ]},
     )
 
@@ -6484,7 +6505,8 @@ def km_3d_view_tool_select_box(params, *, fallback):
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "view3d.select_box",
                 # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if (fallback and params.use_fallback_tool_select_mouse) else
+                   params.tool_tweak_event))),
         ]},
     )
 
@@ -6498,7 +6520,7 @@ def km_3d_view_tool_select_circle(params, *, fallback):
                 "view3d.select_circle",
                 # Why circle select should be used on tweak?
                 # So that RMB or Shift-RMB is still able to set an element as active.
-                type=params.select_tweak if fallback else params.tool_mouse,
+                type=params.select_tweak if (fallback and params.use_fallback_tool_select_mouse) else params.tool_mouse,
                 value='ANY' if fallback else 'PRESS',
                 properties=[("wait_for_input", False)])),
         ]},
@@ -6512,7 +6534,8 @@ def km_3d_view_tool_select_lasso(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "view3d.select_lasso",
-                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if (fallback and params.use_fallback_tool_select_mouse) else
+                   params.tool_tweak_event))),
         ]}
     )
 
@@ -7394,7 +7417,8 @@ def km_3d_view_tool_edit_gpencil_select_box(params, *, fallback):
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "gpencil.select_box",
                 # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if (fallback and params.use_fallback_tool_select_mouse) else
+                   params.tool_tweak_event))),
         ]},
     )
 
@@ -7408,7 +7432,7 @@ def km_3d_view_tool_edit_gpencil_select_circle(params, *, fallback):
                 "gpencil.select_circle",
                 # Why circle select should be used on tweak?
                 # So that RMB or Shift-RMB is still able to set an element as active.
-                type=params.select_tweak if fallback else params.tool_mouse,
+                type=params.select_tweak if (fallback and params.use_fallback_tool_select_mouse) else params.tool_mouse,
                 value='ANY' if fallback else 'PRESS',
                 properties=[("wait_for_input", False)])),
         ]},
@@ -7422,7 +7446,8 @@ def km_3d_view_tool_edit_gpencil_select_lasso(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "gpencil.select_lasso",
-                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if (fallback and params.use_fallback_tool_select_mouse) else
+                   params.tool_tweak_event))),
         ]}
     )
 
@@ -7573,7 +7598,8 @@ def km_sequencer_editor_tool_generic_select_box(params, *, fallback):
             # Don't use `tool_maybe_tweak_event`, see comment for this slot.
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
                 "sequencer.select_box",
-                **(params.select_tweak_event if fallback else params.tool_tweak_event),
+                **(params.select_tweak_event if (fallback and params.use_fallback_tool_select_mouse) else
+                   params.tool_tweak_event),
                 properties=[("tweak", params.select_mouse == 'LEFTMOUSE')])),
 
             # RMB select can already set the frame, match the tweak tool.

release/scripts/presets/keyconfig/keymap_data/industry_compatible_data.py

@@ -1126,7 +1126,8 @@ def km_node_editor(params):
          {"properties": [("replace", False)]}),
         ("node.link_make", {"type": 'L', "value": 'PRESS', "shift": True},
          {"properties": [("replace", True)]}),
-        ("node.duplicate_move", {"type": 'D', "value": 'PRESS', "ctrl": True}, None),
+        ("node.duplicate_move", {"type": 'D', "value": 'PRESS', "ctrl": True},
+         {"properties": [("NODE_OT_translate_attach", [("TRANSFORM_OT_translate", [("view2d_edge_pan", True)])])]}),
         ("node.parent_set", {"type": 'P', "value": 'PRESS'}, None),
         ("node.join", {"type": 'J', "value": 'PRESS', "ctrl": True}, None),
         ("node.hide_toggle", {"type": 'H', "value": 'PRESS', "ctrl": True}, None),

release/scripts/startup/bl_operators/screen_play_rendered_anim.py

@@ -74,6 +74,12 @@ class PlayRenderedAnim(Operator):
         # file_path = bpy.path.abspath(rd.filepath)  # UNUSED
         is_movie = rd.is_movie_format
 
+        views_format = rd.image_settings.views_format
+        if rd.use_multiview and views_format == 'INDIVIDUAL':
+            view_suffix = rd.views.active.file_suffix
+        else:
+            view_suffix = ""
+
         # try and guess a command line if it doesn't exist
         if preset == 'CUSTOM':
             player_path = prefs.filepaths.animation_player
@@ -82,16 +88,16 @@ class PlayRenderedAnim(Operator):
 
         if is_movie is False and preset in {'FRAMECYCLER', 'RV', 'MPLAYER'}:
             # replace the number with '#'
-            file_a = rd.frame_path(frame=0)
+            file_a = rd.frame_path(frame=0, view=view_suffix)
 
             # TODO, make an api call for this
             frame_tmp = 9
-            file_b = rd.frame_path(frame=frame_tmp)
+            file_b = rd.frame_path(frame=frame_tmp, view=view_suffix)
 
             while len(file_a) == len(file_b):
                 frame_tmp = (frame_tmp * 10) + 9
-                file_b = rd.frame_path(frame=frame_tmp)
-            file_b = rd.frame_path(frame=int(frame_tmp / 10))
+                file_b = rd.frame_path(frame=frame_tmp, view=view_suffix)
+            file_b = rd.frame_path(frame=int(frame_tmp / 10), view=view_suffix)
 
             file = ("".join((c if file_b[i] == c else "#")
                             for i, c in enumerate(file_a)))
@@ -100,7 +106,7 @@ class PlayRenderedAnim(Operator):
         else:
             path_valid = True
             # works for movies and images
-            file = rd.frame_path(frame=scene.frame_start, preview=scene.use_preview_range)
+            file = rd.frame_path(frame=scene.frame_start, preview=scene.use_preview_range, view=view_suffix)
             file = bpy.path.abspath(file)  # expand '//'
             if not os.path.exists(file):
                 err_msg = tip_("File %r not found") % file
@@ -109,7 +115,7 @@ class PlayRenderedAnim(Operator):
 
             # one last try for full range if we used preview range
             if scene.use_preview_range and not path_valid:
-                file = rd.frame_path(frame=scene.frame_start, preview=False)
+                file = rd.frame_path(frame=scene.frame_start, preview=False, view=view_suffix)
                 file = bpy.path.abspath(file)  # expand '//'
                 err_msg = tip_("File %r not found") % file
                 if not os.path.exists(file):

release/scripts/startup/bl_ui/properties_data_camera.py

@@ -117,21 +117,8 @@ class DATA_PT_lens(CameraButtonsPanel, Panel):
                     col.prop(ccam, "fisheye_polynomial_k2", text="K2")
                     col.prop(ccam, "fisheye_polynomial_k3", text="K3")
                     col.prop(ccam, "fisheye_polynomial_k4", text="K4")
-            elif engine == 'BLENDER_EEVEE':
-                col.prop(cam, "panorama_type")
-                if cam.panorama_type == 'FISHEYE_EQUIDISTANT':
-                    col.prop(cam, "fisheye_fov")
-                elif cam.panorama_type == 'FISHEYE_EQUISOLID':
-                    col.prop(cam, "fisheye_lens", text="Lens")
-                    col.prop(cam, "fisheye_fov")
-                elif cam.panorama_type == 'EQUIRECTANGULAR':
-                    sub = col.column(align=True)
-                    sub.prop(cam, "latitude_min", text="Latitude Min")
-                    sub.prop(cam, "latitude_max", text="Max")
-                    sub = col.column(align=True)
-                    sub.prop(cam, "longitude_min", text="Longitude Min")
-                    sub.prop(cam, "longitude_max", text="Max")
-            elif engine in {'BLENDER_RENDER', 'BLENDER_WORKBENCH'}:
+
+            elif engine in {'BLENDER_RENDER', 'BLENDER_EEVEE', 'BLENDER_WORKBENCH'}:
                 if cam.lens_unit == 'MILLIMETERS':
                     col.prop(cam, "lens")
                 elif cam.lens_unit == 'FOV':

release/scripts/startup/bl_ui/properties_material.py

@@ -227,6 +227,8 @@ def draw_material_settings(self, context):
         layout.prop(mat, "show_transparent_back")
 
     layout.prop(mat, "use_screen_refraction")
+    layout.prop(mat, "refraction_depth")
+    layout.prop(mat, "use_sss_translucency")
     layout.prop(mat, "pass_index")
 
 

release/scripts/startup/bl_ui/properties_render.py

@@ -180,37 +180,6 @@ class RENDER_PT_eevee_motion_blur(RenderButtonsPanel, Panel):
         col.prop(props, "motion_blur_max")
         col.prop(props, "motion_blur_steps", text="Steps")
 
-class RENDER_PT_eevee_motion_blur_curve(RenderButtonsPanel, Panel):
-    bl_label = "Shutter Curve"
-    bl_parent_id = "RENDER_PT_eevee_motion_blur"
-    bl_options = {'DEFAULT_CLOSED'}
-    COMPAT_ENGINES = {'BLENDER_EEVEE'}
-
-    @classmethod
-    def poll(cls, context):
-        return (context.engine in cls.COMPAT_ENGINES)
-
-    def draw(self, context):
-        layout = self.layout
-        layout.use_property_split = True
-        scene = context.scene
-        rd = scene.render
-        props = scene.eevee
-
-        layout.active = props.use_motion_blur
-        col = layout.column()
-
-        col.template_curve_mapping(rd, "motion_blur_shutter_curve")
-
-        col = layout.column(align=True)
-        row = col.row(align=True)
-        row.operator("render.shutter_curve_preset", icon='SMOOTHCURVE', text="").shape = 'SMOOTH'
-        row.operator("render.shutter_curve_preset", icon='SPHERECURVE', text="").shape = 'ROUND'
-        row.operator("render.shutter_curve_preset", icon='ROOTCURVE', text="").shape = 'ROOT'
-        row.operator("render.shutter_curve_preset", icon='SHARPCURVE', text="").shape = 'SHARP'
-        row.operator("render.shutter_curve_preset", icon='LINCURVE', text="").shape = 'LINE'
-        row.operator("render.shutter_curve_preset", icon='NOCURVE', text="").shape = 'MAX'
-
 
 class RENDER_PT_eevee_depth_of_field(RenderButtonsPanel, Panel):
     bl_label = "Depth of Field"
@@ -360,8 +329,8 @@ class RENDER_PT_eevee_subsurface_scattering(RenderButtonsPanel, Panel):
         col.prop(props, "sss_jitter_threshold")
 
 
-class RENDER_PT_eevee_raytracing(RenderButtonsPanel, Panel):
-    bl_label = "Raytracing"
+class RENDER_PT_eevee_screen_space_reflections(RenderButtonsPanel, Panel):
+    bl_label = "Screen Space Reflections"
     bl_options = {'DEFAULT_CLOSED'}
     COMPAT_ENGINES = {'BLENDER_EEVEE'}
 
@@ -383,9 +352,12 @@ class RENDER_PT_eevee_raytracing(RenderButtonsPanel, Panel):
 
         col = layout.column()
         col.active = props.use_ssr
+        col.prop(props, "use_ssr_refraction", text="Refraction")
+        col.prop(props, "use_ssr_halfres")
         col.prop(props, "ssr_quality")
         col.prop(props, "ssr_max_roughness")
         col.prop(props, "ssr_thickness")
+        col.prop(props, "ssr_border_fade")
         col.prop(props, "ssr_firefly_fac")
 
 
@@ -519,7 +491,6 @@ class RENDER_PT_eevee_film(RenderButtonsPanel, Panel):
         col = layout.column()
         col.prop(rd, "filter_size")
         col.prop(rd, "film_transparent", text="Transparent")
-        col.prop(props, "use_log_space")
 
         col = layout.column(align=False, heading="Overscan")
         row = col.row(align=True)
@@ -732,9 +703,8 @@ classes = (
     RENDER_PT_eevee_bloom,
     RENDER_PT_eevee_depth_of_field,
     RENDER_PT_eevee_subsurface_scattering,
-    RENDER_PT_eevee_raytracing,
+    RENDER_PT_eevee_screen_space_reflections,
     RENDER_PT_eevee_motion_blur,
-    RENDER_PT_eevee_motion_blur_curve,
     RENDER_PT_eevee_volumetric,
     RENDER_PT_eevee_volumetric_lighting,
     RENDER_PT_eevee_volumetric_shadows,