Merge branch 'bf-blender' into mb-0013-blender-top-bar

.gitmodules

@@ -18,6 +18,11 @@
 	path = lib/windows_x64
 	url = https://projects.blender.org/blender/lib-windows_x64.git
 	branch = main
+[submodule "lib/windows_arm64"]
+    update = none
+	path = lib/windows_arm64
+	url = https://projects.blender.org/blender/lib-windows_arm64.git
+	branch = main
 [submodule "release/datafiles/assets"]
 	path = release/datafiles/assets
 	url = https://projects.blender.org/blender/blender-assets.git

build_files/build_environment/cmake/versions.cmake

@@ -529,9 +529,9 @@ set(MATERIALX_HASH fad8f4e19305fb2ee920cbff638f3560)
 set(MATERIALX_HASH_TYPE MD5)
 set(MATERIALX_FILE materialx-v${MATERIALX_VERSION}.tar.gz)
 
-set(OIDN_VERSION 2.2.1)
+set(OIDN_VERSION 2.2.2)
 set(OIDN_URI https://github.com/OpenImageDenoise/oidn/releases/download/v${OIDN_VERSION}/oidn-${OIDN_VERSION}.src.tar.gz)
-set(OIDN_HASH a1c5299b2b640a0e0569afcf405c82bf)
+set(OIDN_HASH 40c04b0371334ab863230e99a587fd59)
 set(OIDN_HASH_TYPE MD5)
 set(OIDN_FILE oidn-${OIDN_VERSION}.src.tar.gz)
 

build_files/cmake/cmake_static_check_cppcheck.py

@@ -30,7 +30,7 @@ CHECKER_EXCLUDE_SOURCE_FILES = set(os.path.join(*f.split("/")) for f in (
     "source/blender/draw/engines/eevee_next/eevee_lut.cc",
 ))
 
-CHECKER_ARGS = [
+CHECKER_ARGS = (
     # Speed up execution.
     # As Blender has many defines, the total number of configurations is large making execution unreasonably slow.
     # This could be increased but do so with care.
@@ -39,9 +39,15 @@ CHECKER_ARGS = [
     # Enable this when includes are missing.
     #  "--check-config",
 
+    # This is slower, for a comprehensive output it is needed.
+    "--check-level=exhaustive",
+
     # Shows many pedantic issues, some are quite useful.
     "--enable=all",
 
+    # Generates many warnings, CPPCHECK known about system includes without resolving them.
+    "--suppress=missingIncludeSystem",
+
     # Also shows useful messages, even if some are false-positives.
     "--inconclusive",
 
@@ -50,7 +56,15 @@ CHECKER_ARGS = [
     *(() if USE_VERBOSE else ("--quiet",))
 
     # NOTE: `--cppcheck-build-dir=<dir>` is added later as a temporary directory.
-]
+)
+
+CHECKER_ARGS_C = (
+    "--std=c11",
+)
+
+CHECKER_ARGS_CXX = (
+    "--std=c++17",
+)
 
 
 def source_info_filter(
@@ -74,22 +88,50 @@ def source_info_filter(
     return source_info_result
 
 
-def cppcheck() -> None:
+def cppcheck(temp_dir: str) -> None:
+    temp_build_dir = os.path.join(temp_dir, "build")
+    temp_source_dir = os.path.join(temp_dir, "source")
+    del temp_dir
+
+    os.mkdir(temp_build_dir)
+    os.mkdir(temp_source_dir)
+
     source_info = project_source_info.build_info(ignore_prefix_list=CHECKER_IGNORE_PREFIX)
-    source_defines = project_source_info.build_defines_as_args()
+    cppcheck_compiler_h = os.path.join(temp_source_dir, "cppcheck_compiler.h")
+    with open(cppcheck_compiler_h, "w", encoding="utf-8") as fh:
+        fh.write(project_source_info.build_defines_as_source())
+
+        # Add additional defines.
+        fh.write("\n")
+        # Python's `pyport.h` errors without this.
+        fh.write("#define UCHAR_MAX 255\n")
+        # `intern/atomic/intern/atomic_ops_utils.h` errors with `Cannot find int size` without this.
+        fh.write("#define UINT_MAX 0xFFFFFFFF\n")
 
     # Apply exclusion.
     source_info = source_info_filter(source_info)
 
     check_commands = []
     for c, inc_dirs, defs in source_info:
+        if c.endswith(".c"):
+            checker_args_extra = CHECKER_ARGS_C
+        else:
+            checker_args_extra = CHECKER_ARGS_CXX
+
         cmd = (
-            [CHECKER_BIN] +
-            CHECKER_ARGS +
-            [c] +
-            [("-I%s" % i) for i in inc_dirs] +
-            [("-D%s" % d) for d in defs] +
-            source_defines
+            CHECKER_BIN,
+            *CHECKER_ARGS,
+            *checker_args_extra,
+            "--cppcheck-build-dir=" + temp_build_dir,
+            "--include=" + cppcheck_compiler_h,
+            # NOTE: for some reason failing to include this crease a large number of syntax errors
+            # from `intern/guardedalloc/MEM_guardedalloc.h`. Include directly to resolve.
+            "--include=" + os.path.join(
+                project_source_info.SOURCE_DIR, "source", "blender", "blenlib", "BLI_compiler_attrs.h",
+            ),
+            c,
+            *[("-I%s" % i) for i in inc_dirs],
+            *[("-D%s" % d) for d in defs],
         )
 
         check_commands.append((c, cmd))
@@ -119,8 +161,7 @@ def cppcheck() -> None:
 
 def main() -> None:
     with tempfile.TemporaryDirectory() as temp_dir:
-        CHECKER_ARGS.append("--cppcheck-build-dir=" + temp_dir)
-        cppcheck()
+        cppcheck(temp_dir)
 
 
 if __name__ == "__main__":

build_files/cmake/macros.cmake

@@ -473,6 +473,8 @@ endfunction()
 # Ninja only: assign 'heavy pool' to some targets that are especially RAM-consuming to build.
 function(setup_heavy_lib_pool)
   if(WITH_NINJA_POOL_JOBS AND NINJA_MAX_NUM_PARALLEL_COMPILE_HEAVY_JOBS)
+    set(_HEAVY_LIBS)
+    set(_TARGET)
     if(WITH_CYCLES)
       list(APPEND _HEAVY_LIBS "cycles_device" "cycles_kernel")
     endif()
@@ -483,11 +485,13 @@ function(setup_heavy_lib_pool)
       list(APPEND _HEAVY_LIBS "bf_intern_openvdb")
     endif()
 
-    foreach(TARGET ${_HEAVY_LIBS})
-      if(TARGET ${TARGET})
-        set_property(TARGET ${TARGET} PROPERTY JOB_POOL_COMPILE compile_heavy_job_pool)
+    foreach(_TARGET ${_HEAVY_LIBS})
+      if(TARGET ${_TARGET})
+        set_property(TARGET ${_TARGET} PROPERTY JOB_POOL_COMPILE compile_heavy_job_pool)
       endif()
     endforeach()
+    unset(_TARGET)
+    unset(_HEAVY_LIBS)
   endif()
 endfunction()
 

build_files/cmake/project_source_info.py

@@ -24,12 +24,12 @@ from typing import (
     Any,
     Callable,
     Generator,
+    IO,
     List,
     Optional,
     Sequence,
     Tuple,
     Union,
-    cast,
 )
 
 import shlex
@@ -120,12 +120,13 @@ def makefile_log() -> List[str]:
         time.sleep(1)
 
     # We know this is always true based on the input arguments to `Popen`.
-    stdout: IO[bytes] = process.stdout  # type: ignore
+    assert process.stdout is not None
+    stdout: IO[bytes] = process.stdout
 
     out = stdout.read()
     stdout.close()
     print("done!", len(out), "bytes")
-    return cast(List[str], out.decode("utf-8", errors="ignore").split("\n"))
+    return out.decode("utf-8", errors="ignore").split("\n")
 
 
 def build_info(
@@ -211,9 +212,10 @@ def build_defines_as_source() -> str:
     )
 
     # We know this is always true based on the input arguments to `Popen`.
-    stdout: IO[bytes] = process.stdout  # type: ignore
+    assert process.stdout is not None
+    stdout: IO[bytes] = process.stdout
 
-    return cast(str, stdout.read().strip().decode('ascii'))
+    return stdout.read().strip().decode('ascii')
 
 
 def build_defines_as_args() -> List[str]:

build_files/utils/make_bpy_wheel.py

@@ -176,7 +176,6 @@ def main() -> None:
         # Support version without a minor version "3" (add zero).
         tuple((0, 0, 0))
     )
-    python_version_str = "%d.%d" % python_version_number[:2]
 
     # Get Blender version.
     blender_version_str = str(make_utils.parse_blender_version())

build_files/utils/make_source_archive.py

@@ -33,7 +33,7 @@ def main() -> None:
     blender_srcdir = Path(__file__).absolute().parent.parent.parent
 
     cli_parser = argparse.ArgumentParser(
-        description=f"Create a tarball of the Blender sources, optionally including sources of dependencies.",
+        description="Create a tarball of the Blender sources, optionally including sources of dependencies.",
         epilog="This script is intended to be run by `make source_archive_complete`.",
     )
     cli_parser.add_argument(

build_files/utils/make_update.py

@@ -393,7 +393,7 @@ def floating_checkout_add_origin_if_needed(
         upstream_url = make_utils.git_get_remote_url(args.git_command, "upstream")
 
         call((args.git_command, "remote", "rename", "upstream", "origin"))
-        make_utils.git_set_config(args.git_command, f"remote.origin.url", origin_external_url)
+        make_utils.git_set_config(args.git_command, "remote.origin.url", origin_external_url)
 
         call((args.git_command, "remote", "add", "upstream", upstream_url))
     finally:

build_files/utils/make_utils.py

@@ -115,7 +115,8 @@ def git_branch(git_command: str) -> str:
 
     try:
         branch = subprocess.check_output([git_command, "rev-parse", "--abbrev-ref", "HEAD"])
-    except subprocess.CalledProcessError as e:
+    except subprocess.CalledProcessError:
+        # No need to print the exception, error text is written to the output already.
         sys.stderr.write("Failed to get Blender git branch\n")
         sys.exit(1)
 

doc/python_api/conf.py

@@ -0,0 +1,134 @@
+# SPDX-FileCopyrightText: 2024 Blender Authors
+#
+# SPDX-License-Identifier: GPL-2.0-or-later
+
+import time
+
+
+def has_module(module_name):
+    found = False
+    try:
+        __import__(module_name)
+        found = True
+    except ModuleNotFoundError as ex:
+        if ex.name != module_name:
+            raise ex
+    return found
+
+
+# These are substituted when this file is copied to the build directory.
+BLENDER_VERSION_STRING = "${BLENDER_VERSION_STRING}"
+BLENDER_VERSION_DOTS = "${BLENDER_VERSION_DOTS}"
+BLENDER_REVISION = "${BLENDER_REVISION}"
+BLENDER_REVISION_TIMESTAMP = "${BLENDER_REVISION_TIMESTAMP}"
+BLENDER_VERSION_DATE = time.strftime(
+    "%d/%m/%Y",
+    time.localtime(int(BLENDER_REVISION_TIMESTAMP) if BLENDER_REVISION_TIMESTAMP != "0" else None),
+)
+
+if BLENDER_REVISION != "Unknown":
+    # SHA1 GIT hash.
+    BLENDER_VERSION_HASH = BLENDER_REVISION
+    BLENDER_VERSION_HASH_HTML_LINK = "<a href=https://projects.blender.org/blender/blender/commit/%s>%s</a>" % (
+        BLENDER_VERSION_HASH, BLENDER_VERSION_HASH)
+else:
+    # Fallback: Should not be used.
+    BLENDER_VERSION_HASH = "Hash Unknown"
+    BLENDER_VERSION_HASH_HTML_LINK = BLENDER_VERSION_HASH
+
+extensions = ["sphinx.ext.intersphinx"]
+intersphinx_mapping = {"blender_manual": ("https://docs.blender.org/manual/en/dev/", None)}
+
+
+# Provides copy button next to code-blocks (nice to have but not essential).
+if has_module("sphinx_copybutton"):
+    extensions.append("sphinx_copybutton")
+
+    # Exclude line numbers, prompts, and console text.
+    copybutton_exclude = ".linenos, .gp, .go"
+
+
+project = "Blender %s Python API" % BLENDER_VERSION_STRING
+root_doc = "index"
+copyright = "Blender Authors"
+version = BLENDER_VERSION_DOTS
+release = BLENDER_VERSION_DOTS
+
+# Set this as the default is a super-set of Python3.
+highlight_language = "python3"
+# No need to detect encoding.
+highlight_options = {"default": {"encoding": "utf-8"}}
+
+# Quiet file not in table-of-contents warnings.
+exclude_patterns = [
+    "include__bmesh.rst",
+]
+
+html_title = "Blender Python API"
+
+# The fallback to a built-in theme when `furo` is not found.
+html_theme = "default"
+
+if has_module("furo"):
+    html_theme = "furo"
+    html_theme_options = {
+        "light_css_variables": {
+            "color-brand-primary": "#265787",
+            "color-brand-content": "#265787",
+        },
+    }
+
+    html_sidebars = {
+        "**": [
+            "sidebar/brand.html",
+            "sidebar/search.html",
+            "sidebar/scroll-start.html",
+            "sidebar/navigation.html",
+            "sidebar/scroll-end.html",
+            "sidebar/variant-selector.html",
+        ]
+    }
+
+# Not helpful since the source is generated, adds to upload size.
+html_copy_source = False
+html_show_sphinx = False
+html_baseurl = "https://docs.blender.org/api/current/"
+html_use_opensearch = "https://docs.blender.org/api/current"
+html_show_search_summary = True
+html_split_index = True
+html_static_path = ["static"]
+templates_path = ["templates"]
+html_context = {"commit": "%s - %s" % (BLENDER_VERSION_HASH_HTML_LINK, BLENDER_VERSION_DATE)}
+html_extra_path = ["static"]
+html_favicon = "static/favicon.ico"
+html_logo = "static/blender_logo.svg"
+# Disable default `last_updated` value, since this is the date of doc generation, not the one of the source commit.
+html_last_updated_fmt = None
+if html_theme == "furo":
+    html_css_files = ["css/theme_overrides.css", "css/version_switch.css"]
+    html_js_files = ["js/version_switch.js"]
+
+# Needed for latex, PDF generation.
+latex_elements = {
+    "papersize": "a4paper",
+}
+
+latex_documents = [
+    ("contents", "contents.tex", "Blender Index", "Blender Foundation", "manual"),
+]
+
+# Workaround for useless links leading to compile errors
+# See https://github.com/sphinx-doc/sphinx/issues/3866
+from sphinx.domains.python import PythonDomain
+
+
+class PatchedPythonDomain(PythonDomain):
+    def resolve_xref(self, env, fromdocname, builder, typ, target, node, contnode):
+        if "refspecific" in node:
+            del node["refspecific"]
+        return super(PatchedPythonDomain, self).resolve_xref(
+            env, fromdocname, builder, typ, target, node, contnode)
+
+
+def setup(app):
+    app.add_domain(PatchedPythonDomain, override=True)

doc/python_api/examples/bpy.types.Depsgraph.4.py

@@ -17,10 +17,10 @@ animation or modifiers into account:
 
 When is used on evaluated object all modifiers are taken into account.
 
-.. note:: The result mesh is owned by the object. It can be freed by calling `object.to_mesh_clear()`.
+.. note:: The result mesh is owned by the object. It can be freed by calling :meth:`~Object.to_mesh_clear`.
 .. note::
    The result mesh must be treated as temporary, and cannot be referenced from objects in the main
-   database. If the mesh intended to be used in a persistent manner use bpy.data.meshes.new_from_object()
+   database. If the mesh intended to be used in a persistent manner use :meth:`~BlendDataMeshes.new_from_object`
    instead.
 .. note:: If object does not have geometry (i.e. camera) the functions returns None.
 """

doc/python_api/examples/bpy.types.Depsgraph.7.py

@@ -15,7 +15,7 @@ If the object is a text object. The text will be converted into a 3D curve and r
 never applied on text objects and apply_modifiers will be ignored. If the object is neither a curve nor
 a text object, an error will be reported.
 
-.. note:: The resulting curve is owned by the object. It can be freed by calling `object.to_curve_clear()`.
+.. note:: The resulting curve is owned by the object. It can be freed by calling :meth:`~Object.to_curve_clear`.
 .. note::
    The resulting curve must be treated as temporary, and cannot be referenced from objects in the main
    database.

doc/python_api/examples/bpy.types.FileHandler.2.py

@@ -16,10 +16,11 @@ This ``directory`` and ``files`` properties now will be used by the
 """
 
 import bpy
+from bpy_extras.io_utils import ImportHelper
 from mathutils import Vector
 
 
-class ShaderScriptImport(bpy.types.Operator):
+class ShaderScriptImport(bpy.types.Operator, ImportHelper):
     """Test importer that creates scripts nodes from .txt files"""
     bl_idname = "shader.script_import"
     bl_label = "Import a text file as a script node"
@@ -28,8 +29,11 @@ class ShaderScriptImport(bpy.types.Operator):
     This Operator can import multiple .txt files, we need following directory and files
     properties that the file handler will use to set files path data
     """
-    directory: bpy.props.StringProperty(subtype='FILE_PATH', options={'SKIP_SAVE'})
-    files: bpy.props.CollectionProperty(type=bpy.types.OperatorFileListElement, options={'SKIP_SAVE'})
+    directory: bpy.props.StringProperty(subtype='FILE_PATH', options={'SKIP_SAVE', 'HIDDEN'})
+    files: bpy.props.CollectionProperty(type=bpy.types.OperatorFileListElement, options={'SKIP_SAVE', 'HIDDEN'})
+
+    """Allow the user to select if the node's label is set or not"""
+    set_label: bpy.props.BoolProperty(name="Set Label", default=False)
 
     @classmethod
     def poll(cls, context):
@@ -57,23 +61,26 @@ class ShaderScriptImport(bpy.types.Operator):
                 filepath = os.path.join(self.directory, file.name)
                 text_node.filepath = filepath
                 text_node.location = Vector((x, y))
+
+                # Set the node's title to the file name
+                if self.set_label:
+                    text_node.label = file.name
+
                 x += 20.0
                 y -= 20.0
         return {'FINISHED'}
 
     """
-    By default the file handler invokes the operator with the directory and files properties set.
-    In this example if this properties are set the operator is executed, if not the
-    file select window is invoked.
-    This depends on setting ``options={'SKIP_SAVE'}`` to the properties options to avoid
-    to reuse filepath data between operator calls.
+    Use ImportHelper's invoke_popup() to handle the invocation so that this operator's properties
+    are shown in a popup. This allows the user to configure additional settings on the operator like
+    the `set_label` property. Consider having a draw() method on the operator in order to layout the
+    properties in the UI appropriately.
+
+    If filepath information is not provided the file select window will be invoked instead.
     """
 
     def invoke(self, context, event):
-        if self.directory:
-            return self.execute(context)
-        context.window_manager.fileselect_add(self)
-        return {'RUNNING_MODAL'}
+        return self.invoke_popup(context)
 
 
 class SHADER_FH_script_import(bpy.types.FileHandler):

doc/python_api/examples/bpy.types.UIList.2.py

@@ -136,8 +136,9 @@ class MESH_UL_vgroups_slow(bpy.types.UIList):
     def filter_items(self, context, data, propname):
         # This function gets the collection property (as the usual tuple (data, propname)), and must return two lists:
         # * The first one is for filtering, it must contain 32bit integers were self.bitflag_filter_item marks the
-        #   matching item as filtered (i.e. to be shown), and 31 other bits are free for custom needs. Here we use the
-        #   first one to mark VGROUP_EMPTY.
+        #   matching item as filtered (i.e. to be shown). The upper 16 bits (including self.bitflag_filter_item) are
+        #   reserved for internal use, the lower 16 bits are free for custom use. Here we use the first bit to mark
+        #   VGROUP_EMPTY.
         # * The second one is for reordering, it must return a list containing the new indices of the items (which
         #   gives us a mapping org_idx -> new_idx).
         # Please note that the default UI_UL_list defines helper functions for common tasks (see its doc for more info).

doc/python_api/examples/bpy.utils.register_cli_command.1.py

@@ -1,6 +1,5 @@
 """
-Using Python Argument Parsing
------------------------------
+**Using Python Argument Parsing**
 
 This example shows how the Python ``argparse`` module can be used with a custom command.
 

doc/python_api/examples/bpy.utils.register_cli_command.py

@@ -1,6 +1,5 @@
 """
-Custom Commands
----------------
+**Custom Commands**
 
 Registering commands makes it possible to conveniently expose command line
 functionality via commands passed to (``-c`` / ``--command``).

doc/python_api/requirements.txt

@@ -1,12 +1,16 @@
-sphinx==7.1.2
+sphinx==7.2.6
 
 # Sphinx dependencies that are important
-Jinja2==3.1.2
-Pygments==2.16.1
-docutils==0.18.1
+Jinja2==3.1.3
+Pygments==2.17.2
+docutils==0.20.1
 snowballstemmer==2.2.0
 requests==2.31.0
 
 # Only needed to match the theme used for the official documentation.
 # Without this theme, the default theme will be used.
-sphinx_rtd_theme==1.3.0rc1
+furo==2024.1.29
+sphinx-basic-ng==1.0.0b2
+
+# Show a copy button (convenience only).
+sphinx-copybutton==0.5.2

doc/python_api/rst_from_bmesh_opdefines.py

@@ -252,8 +252,7 @@ def main():
                     name, tp = arg
                     tp_sub = None
                 else:
-                    print(arg)
-                    assert 0
+                    assert False, "unreachable, unsupported 'arg' length found %d" % len(arg)
 
                 tp_str = ""
 
@@ -322,8 +321,7 @@ def main():
                         #     but think the idea is that that pointer is for any type?
                         tp_str = ":class:`bpy.types.bpy_struct`"
                     else:
-                        print("Can't find", vars_dict_reverse[tp_sub])
-                        assert 0
+                        assert False, "unreachable, unknown type %r" % vars_dict_reverse[tp_sub]
 
                 elif tp == BMO_OP_SLOT_ELEMENT_BUF:
                     assert tp_sub is not None
@@ -362,11 +360,9 @@ def main():
                         elif tp_sub == BMO_OP_SLOT_SUBTYPE_MAP_INTERNAL:
                             tp_str += "unknown internal data, not compatible with python"
                         else:
-                            print("Can't find", vars_dict_reverse[tp_sub])
-                            assert 0
+                            assert False, "unreachable, unknown type %r" % vars_dict_reverse[tp_sub]
                 else:
-                    print("Can't find", vars_dict_reverse[tp])
-                    assert 0
+                    assert False, "unreachable, unknown type %r" % vars_dict_reverse[tp]
 
                 args_wash.append((name, default_value, tp_str, comment))
             return args_wash

doc/python_api/sphinx_doc_gen.py

@@ -61,7 +61,6 @@ import os
 import sys
 import inspect
 import shutil
-import time
 import logging
 import warnings
 
@@ -486,19 +485,6 @@ BLENDER_REVISION_TIMESTAMP = bpy.app.build_commit_timestamp
 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
-    BLENDER_VERSION_HASH = BLENDER_REVISION
-    BLENDER_VERSION_HASH_HTML_LINK = "<a href=https://projects.blender.org/blender/blender/commit/%s>%s</a>" % (
-        BLENDER_VERSION_HASH, BLENDER_VERSION_HASH,
-    )
-    BLENDER_VERSION_DATE = time.strftime("%d/%m/%Y", time.localtime(BLENDER_REVISION_TIMESTAMP))
-else:
-    # Fallback: Should not be used
-    BLENDER_VERSION_HASH = "Hash Unknown"
-    BLENDER_VERSION_HASH_HTML_LINK = BLENDER_VERSION_HASH
-    BLENDER_VERSION_DATE = time.strftime("%Y-%m-%d")
-
 # Example: `2_83`.
 BLENDER_VERSION_PATH = "%d_%d" % (bpy.app.version[0], bpy.app.version[1])
 
@@ -1712,7 +1698,7 @@ def pyrna2sphinx(basepath):
                     lines.append("   * :class:`%s.%s`\n" % (base.identifier, identifier))
 
             if lines:
-                fw(".. rubric:: Inherited Properties\n\n")
+                fw(title_string("Inherited Properties", "-"))
 
                 fw(".. hlist::\n")
                 fw("   :columns: 2\n\n")
@@ -1738,7 +1724,7 @@ def pyrna2sphinx(basepath):
                     lines.append("   * :class:`%s.%s`\n" % (base.identifier, identifier))
 
             if lines:
-                fw(".. rubric:: Inherited Functions\n\n")
+                fw(title_string("Inherited Functions", "-"))
 
                 fw(".. hlist::\n")
                 fw("   :columns: 2\n\n")
@@ -1750,8 +1736,7 @@ def pyrna2sphinx(basepath):
             del lines[:]
 
         if struct.references:
-            # use this otherwise it gets in the index for a normal heading.
-            fw(".. rubric:: References\n\n")
+            fw(title_string("References", "-"))
 
             fw(".. hlist::\n")
             fw("   :columns: 2\n\n")
@@ -1896,103 +1881,6 @@ def pyrna2sphinx(basepath):
         write_ops()
 
 
-def write_sphinx_conf_py(basepath):
-    """
-    Write sphinx's ``conf.py``.
-    """
-    filepath = os.path.join(basepath, "conf.py")
-    file = open(filepath, "w", encoding="utf-8")
-    fw = file.write
-
-    fw("import sys, os\n\n")
-    fw("extensions = ['sphinx.ext.intersphinx']\n\n")
-    fw("intersphinx_mapping = {'blender_manual': ('https://docs.blender.org/manual/en/dev/', None)}\n\n")
-    fw("project = 'Blender %s Python API'\n" % BLENDER_VERSION_STRING)
-    fw("root_doc = 'index'\n")
-    fw("copyright = 'Blender Authors'\n")
-    fw("version = '%s'\n" % BLENDER_VERSION_DOTS)
-    fw("release = '%s'\n" % BLENDER_VERSION_DOTS)
-
-    # Set this as the default is a super-set of Python3.
-    fw("highlight_language = 'python3'\n")
-    # No need to detect encoding.
-    fw("highlight_options = {'default': {'encoding': 'utf-8'}}\n\n")
-
-    # Quiet file not in table-of-contents warnings.
-    fw("exclude_patterns = [\n")
-    fw("    'include__bmesh.rst',\n")
-    fw("]\n\n")
-
-    fw("html_title = 'Blender Python API'\n")
-
-    fw("html_theme = 'default'\n")
-    # The theme 'sphinx_rtd_theme' is no longer distributed with sphinx by default, only use when available.
-    fw(r"""
-try:
-    __import__('sphinx_rtd_theme')
-    html_theme = 'sphinx_rtd_theme'
-except ModuleNotFoundError:
-    pass
-""")
-
-    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")
-    fw("        'navigation_depth': 1,\n")
-    fw("        'includehidden': False,\n")
-    # fw("        'titles_only': False\n")
-    fw("    }\n\n")
-
-    # not helpful since the source is generated, adds to upload size.
-    fw("html_copy_source = False\n")
-    fw("html_show_sphinx = False\n")
-    fw("html_baseurl = 'https://docs.blender.org/api/current/'\n")
-    fw("html_use_opensearch = 'https://docs.blender.org/api/current'\n")
-    fw("html_show_search_summary = True\n")
-    fw("html_split_index = True\n")
-    fw("html_static_path = ['static']\n")
-    fw("templates_path = ['templates']\n")
-    fw("html_context = {'commit': '%s - %s'}\n" % (BLENDER_VERSION_HASH_HTML_LINK, BLENDER_VERSION_DATE))
-    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")
-    # Disable default `last_updated` value, since this is the date of doc generation, not the one of the source commit.
-    fw("html_last_updated_fmt = None\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")
-    fw("  'papersize': 'a4paper',\n")
-    fw("}\n\n")
-
-    fw("latex_documents = [ ('contents', 'contents.tex', 'Blender Index', 'Blender Foundation', 'manual'), ]\n")
-
-    # Workaround for useless links leading to compile errors
-    # See https://github.com/sphinx-doc/sphinx/issues/3866
-    fw(r"""
-from sphinx.domains.python import PythonDomain
-
-class PatchedPythonDomain(PythonDomain):
-    def resolve_xref(self, env, fromdocname, builder, typ, target, node, contnode):
-        if 'refspecific' in node:
-            del node['refspecific']
-        return super(PatchedPythonDomain, self).resolve_xref(
-            env, fromdocname, builder, typ, target, node, contnode)
-""")
-    # end workaround
-
-    fw("def setup(app):\n")
-    fw("    app.add_css_file('css/theme_overrides.css')\n")
-    fw("    app.add_domain(PatchedPythonDomain, override=True)\n\n")
-
-    file.close()
-
-
 def write_rst_index(basepath):
     """
     Write the RST file of the main page, needed for sphinx: ``index.html``.
@@ -2231,7 +2119,6 @@ def write_rst_enum_items(basepath, key, key_no_prefix, enum_items):
         fw(".. _%s:\n\n" % key)
 
         fw(title_string(key_no_prefix.replace("_", " ").title(), "#"))
-        # fw(".. rubric:: %s\n\n" % key_no_prefix.replace("_", " ").title())
 
         for item in enum_items:
             identifier = item.identifier
@@ -2406,7 +2293,7 @@ def copy_handwritten_extra(basepath):
         shutil.copy2(f_src, f_dst)
 
 
-def copy_theme_assets(basepath):
+def copy_sphinx_files(basepath):
     shutil.copytree(
         os.path.join(SCRIPT_DIR, "static"),
         os.path.join(basepath, "static"),
@@ -2418,17 +2305,38 @@ def copy_theme_assets(basepath):
         copy_function=shutil.copy,
     )
 
+    shutil.copy2(os.path.join(SCRIPT_DIR, "conf.py"), basepath, )
+
+
+def format_config(basepath):
+    """
+    Updates ``conf.py`` with context information from Blender.
+    """
+    from string import Template
+
+    # Ensure the string literals can contain any characters by closing the surrounding quotes
+    # and declare a separate literal via `repr()`.
+    def declare_in_quotes(string):
+        return "\" %r \"" % (string)
+
+    substitutions = {
+        "BLENDER_VERSION_STRING": declare_in_quotes(BLENDER_VERSION_STRING),
+        "BLENDER_VERSION_DOTS": declare_in_quotes(BLENDER_VERSION_DOTS),
+        "BLENDER_REVISION_TIMESTAMP": declare_in_quotes(str(BLENDER_REVISION_TIMESTAMP)),
+        "BLENDER_REVISION": declare_in_quotes(BLENDER_REVISION),
+    }
+
+    filepath = os.path.join(basepath, "conf.py")
+
+    # Read the template string from the template file.
+    with open(filepath, 'r', encoding="utf-8") as fh:
+        template_file = fh.read()
+
+    with open(filepath, 'w', encoding="utf-8") as fh:
+        fh.write(Template(template_file).substitute(substitutions))
+
 
 def rna2sphinx(basepath):
-
-    try:
-        os.mkdir(basepath)
-    except:
-        pass
-
-    # sphinx setup
-    write_sphinx_conf_py(basepath)
-
     # main page
     write_rst_index(basepath)
 
@@ -2455,9 +2363,6 @@ def rna2sphinx(basepath):
     # copy source files referenced
     copy_handwritten_extra(basepath)
 
-    # copy extra files needed for theme
-    copy_theme_assets(basepath)
-
 
 def align_sphinx_in_to_sphinx_in_tmp(dir_src, dir_dst):
     """
@@ -2578,10 +2483,22 @@ def main():
             copy_function=shutil.copy,
         )
 
-    # Dump the API in RST files.
+    # start from a clean directory everytime
     if os.path.exists(SPHINX_IN_TMP):
         shutil.rmtree(SPHINX_IN_TMP, True)
 
+    try:
+        os.mkdir(SPHINX_IN_TMP)
+    except:
+        pass
+
+    # copy extra files needed for theme
+    copy_sphinx_files(SPHINX_IN_TMP)
+
+    # write infromation needed for 'conf.py'
+    format_config(SPHINX_IN_TMP)
+
+    # Dump the API in RST files.
     rna2sphinx(SPHINX_IN_TMP)
 
     if ARGS.changelog:

doc/python_api/static/css/theme_overrides.css

@@ -1,19 +1,340 @@
-/* Hide home icon in search area */
-.wy-side-nav-search > a:hover {background: none; opacity: 0.9}
-.wy-side-nav-search > a.icon::before {content: none}
+/*
+ * This stylesheet is applied after the theme's default one,
+ * and thus any overrides or additions can be added here.
+ *
+ * More info:
+ *		https://www.sphinx-doc.org/en/master/extdev/appapi.html#sphinx.application.Sphinx.add_css_file
+ */
 
-.wy-nav-content {
-    max-width: 1000px !important;
+body {
+	--sidebar-caption-font-size: var(--font-size--normal);
+	--toc-title-font-size: var(--font-size--normal);
+	--toc-font-size: var(--sidebar-item-font-size);
+	--admonition-font-size: var(--font-size--normal);
+	--admonition-title-font-size: var(--font-size--normal);
+	--color-api-name: #e87d0d;
 }
 
-/* Fix long titles on mobile */
-h1, h2, h3, h4, h5, h6 {word-break: break-all}
+h1,
+h2,
+h3 {
+	margin-top: 1.75rem;
+	margin-bottom: 1rem;
+}
 
-/* Temp fix for https://github.com/readthedocs/sphinx_rtd_theme/pull/1109 */
+h1 {
+	font-size: 2em;
+}
+
+h2 {
+	font-size: 1.5em;
+}
+
+h3 {
+	font-size: 1.25em;
+}
+
+h4,
+h5,
+h6,
+.rubric {
+	margin-top: 1.25rem;
+	margin-bottom: 0.75rem;
+	font-size: 1.125em;
+}
+
+/* Reduce the margins on top/bottom of horizontal lines. */
+hr.docutils {
+	margin: 1rem 0;
+}
+
+
+/* Slightly decrease text to make the text fit on one line */
+.sidebar-brand-text {
+	font-size: 1.4rem;
+}
+
+.toctree-checkbox~label .icon svg {
+	transition: transform 0.25s ease-out;
+}
+
+/* Add more visual weight to definition terms */
+dl dt {
+	font-weight: bold !important
+}
+
+/* Fixes to field list, see #104636 */
+dl.field-list {
+	display: grid;
+	grid-template-columns: auto minmax(80%, 95%);
+}
+
+dl.field-list > dd > p:last-child {
+	margin-bottom: 0;
+}
+
+@media (max-width: calc(67em / 2)) {
+	dl.field-list  {
+		grid-template-columns: unset;
+    }
+}
+
+/* TABLE & FIGURE */
+
+/* Cell's vertical align. */
+/* use "valign" class for middle align */
+table.docutils:not(.valign) td {
+	vertical-align: baseline;
+}
+
+/* Decrease whitespace above figure and add it below */
+figure {
+	padding-bottom: 0.5rem;
+}
+
+figcaption {
+	margin-bottom: 0.5rem !important;
+
+	p {
+		margin-top: 0;
+	}
+}
+
+/* Allow horizontal lists to collapse on narrow screens */
 .hlist tr {
-   display: -ms-flexbox;
-   display: flex;
-   flex-flow: row wrap;
- }
+	display: flex;
+	flex-flow: row wrap;
+}
 
-.hlist td {margin-right: auto}
+/* End TABLE & FIGURE. */
+
+/* Force admonition to span the full width if close to a figure */
+.admonition {
+	clear: both;
+}
+
+/* Use secondary font color for caption text */
+figcaption,
+caption {
+	color: var(--color-foreground-secondary);
+	font-size: var(--font-size--small)
+}
+
+/* A bit hacky, revert the themes styling of kbd */
+kbd:not(.compound) {
+	all: revert;
+}
+
+/* Only style parent kbd elements instead of the individual children */
+:not(dl.option-list)> :not(kbd):not(kbd)>kbd,
+.menuselection {
+	background-color: var(--color-background-secondary);
+	border: 1px solid var(--color-foreground-border);
+	border-radius: .2rem;
+	box-shadow: 0 .0625rem 0 rgba(0, 0, 0, .2), inset 0 0 0 .125rem var(--color-background-secondary);
+	color: var(--color-foreground-primary);
+	display: inline-block;
+	margin: 0;
+	padding: 0 .2rem;
+}
+
+.highlight .nc,
+.highlight .nn,
+.highlight .gu {
+	text-decoration-line: none !important;
+}
+
+.caption .menuselection {
+	background-color: transparent;
+	border: none;
+}
+
+a {
+	text-decoration: none;
+}
+
+/* Break long code references onto a second line */
+a > code.docutils {
+	overflow-wrap: anywhere;
+}
+
+/* Quotes for Fig. "link". */
+a[href^="#fig-"]::before {
+	content: "\201c";
+}
+
+a[href^="#fig-"]::after {
+	content: "\201d";
+}
+
+/* Mark external links. */
+a.external {
+	filter: brightness(150%);
+}
+
+/* List blender.org as internal. */
+.external[href^="https://www.blender.org"],
+.external[href^="https://docs.blender.org"],
+.external[href^="https://projects.blender.org"],
+.external[href^="https://builder.blender.org"],
+.external[href^="https://code.blender.org"],
+.external[href^="https://translate.blender.org"],
+.external[href^="https://fund.blender.org"],
+.external[href^="blender_manual_html.zip"],
+.external[href^="blender_manual_epub.zip"],
+.external[href^="https://archive.blender.org"] {
+	filter: revert;
+}
+
+/* ".. container::" lead, block text float around image. */
+.lead {
+	clear: both;
+	width: 100%;
+}
+
+/* Start reference admonition. */
+.admonition.refbox {
+	border-color: rgb(50, 50, 50);
+}
+
+.admonition.refbox>.admonition-title {
+	background-color: rgba(50, 50, 50, 0.2);
+	border-color: rgb(50, 50, 50);
+}
+
+.admonition.refbox>.admonition-title::before {
+	background-color: var(--color-content-foreground);
+}
+
+/* 'refbox' field. */
+.refbox .field-list .field-name,
+.refbox .field-list .field-body {
+	padding: 0px;
+}
+
+.refbox dl dt {
+	font-weight: normal
+}
+
+/* End reference admonition. */
+
+/* Applied on main index:sections. */
+
+.global-index-toc {
+	display: none;
+}
+
+/* Start section cards. */
+.toc-cards {
+	display: grid;
+	grid-template-columns: repeat(auto-fit, minmax(350px, 1fr));
+	grid-gap: 20px;
+	list-style-type: none;
+	margin-bottom: 24px;
+}
+
+.card {
+	border-radius: .3em;
+	user-select: none;
+}
+
+.card div.figure,
+.card figure {
+	margin-bottom: 0px;
+	display: block;
+}
+
+.card img {
+	border-top-left-radius: .3em;
+	border-top-right-radius: .3em;
+}
+
+.card dl {
+	margin-bottom: 10px
+}
+
+.card dl dt>a {
+	display: block;
+	width: 100%;
+	margin-bottom: 10px;
+}
+
+.card dl dt a em,
+.card dl dt a span {
+	font-weight: bold;
+	font-style: normal;
+	font-size: 1.3em;
+}
+
+.card dl dt {
+	padding: 0px 15px 0px !important
+}
+
+.card dl dd {
+	padding: 0px 15px 5px 15px;
+	font-style: normal;
+	margin: 0px;
+	color: var(--color-foreground-secondary);
+	font-size: 90%;
+}
+
+.card {
+	box-shadow: 0 .2rem .5rem rgba(0, 0, 0, .05), 0 0 .0625rem rgba(0, 0, 0, .1);
+}
+
+#getting-started .card {
+	box-shadow: none;
+}
+
+/* End section cards. */
+
+/* Start custom toctree. */
+/* Indent all lines following the first. */
+.toctree-wrapper * a {
+	display: block;
+	padding-top: 0.25em;
+}
+
+.toctree-wrapper ul {
+	list-style: none;
+	padding-left: 0;
+}
+
+/* Underline provided by nested ul (not li). */
+.toctree-wrapper * ul {
+	margin-bottom: 1rem !important;
+	border-top: solid var(--color-background-border) 1px;
+	padding-left: 2em;
+}
+
+/* End custom toctree. */
+
+/* Start footer contribute link */
+.footer-contribute {
+	display: block;
+	font-size: var(--font-size--small);
+}
+
+.bottom-of-page {
+	padding-bottom: 0;
+}
+
+.footer-contribute ul {
+	margin: 0;
+	padding: 0;
+	padding-bottom: 1rem
+}
+
+.footer-contribute li {
+	display: inline;
+	list-style-type: none;
+	padding-right: 1.5rem;
+}
+
+@media print {
+	.footer-contribute {
+		display: none;
+	}
+}
+
+/* End footer contribute link */

doc/python_api/static/css/version_switch.css

@@ -1,95 +1,98 @@
-/* Override RTD theme */
-.rst-versions {
-	display: none;
-	border-top: 0px;
-	overflow: visible;
-}
-.version-btn.vdeact {
-	cursor: default;
-	color: dimgray;
-}
-
-.version-btn.vdeact::after {
-	content: "";
-}
 #versionwrap {
+	margin: 0;
 	display: flex;
 	padding-top: 2px;
-	font-size: 90%;
+	padding-left: 0;
+	font-size: var(--sidebar-item-font-size);
 	justify-content: center;
 	flex-wrap: wrap;
 }
+
+#versionwrap>ul {
+	list-style: none;
+}
+
+#versionwrap>li {
+	display: flex;
+	width: 50%;
+}
+
 .version-btn {
 	display: inline-block;
-	background-color: #272525;
-	width: 140px;
+	background-color: var(--color-sidebar-background);
+	width: 100%;
 	text-align: center;
 	padding: 3px 10px;
 	margin: 0px 5px 4px;
 	vertical-align: middle;
-	color: #27AE60;
-	border: solid 1px #444444;
+	color: var(--color-link);
+	border: solid 1px var(--color-sidebar-background-border);
 	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 {
+
+.version-btn:hover,
+.version-btn:focus {
 	border-color: #525252;
 }
+
 .version-btn-open {
 	color: gray;
-	border: solid 1px gray;
+	border: solid 1px var(--color-sidebar-background-border);
 }
+
 .version-btn.wait {
 	cursor: wait;
 }
+
 .version-btn.disabled {
 	cursor: not-allowed;
 	color: dimgray;
 }
+
 .version-dialog {
 	display: none;
 	position: absolute;
 	bottom: 28px;
-	width: 140px;
+	width: 50%;
 	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-x: clip;
 	overflow-y: auto;
 	cursor: default;
 }
+
 .version-title {
 	padding: 5px;
-	color: black;
+	color: var(--color-content-foreground);
 	text-align: center;
 	font-size: 102%;
-	background-color: #27ae60;
-	border-bottom: solid 1.5px #444;
+	font-weight: 700;
+	background-color: var(--color-brand-primary);
+	border-bottom: solid 1.5px var(--color-sidebar-background-border);
 }
+
 .version-list {
+	padding-left: 0;
+	margin-top: 0;
 	margin-bottom: 4px;
 	text-align: center;
-	background-color: #000C;
-	border: solid 1px gray;
+	border: solid 1px var(--color-sidebar-background-border);
 	border-radius: 0px 0px 3px 3px;
 }
-.version-list a, .version-list span, .version-list li {
+
+.version-list a,
+.version-list span,
+.version-list li {
 	position: relative;
 	display: block;
 	font-size: 98%;
@@ -97,32 +100,21 @@
 	width: 100%;
 	margin: 0;
 	padding: 4px 0px;
-	color: #404040;
+	color: var(--color-sidebar-link-text);
 }
+
 .version-list li {
-	background-color: #ede9e9;
-	color: #404040;
+	background-color: var(--color-sidebar-background);
+	color: var(--color-sidebar-link-text);
 	padding: 1px;
 }
-.version-list li:hover, .version-list li a:focus {
-	background-color: #b9cfda;
+
+.version-list li:hover,
+.version-list li a:focus {
+	background-color: var(--color-background-hover);
 }
-.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;
+
+.version-list li.selected {
+	background: var(--color-sidebar-item-background--current);
+	font-weight: 700;
 }

doc/python_api/static/js/version_switch.js

@@ -1,63 +1,60 @@
-(function() {  // switch: v1.2
+(function() {  // switch: v1.4
 "use strict";
 
 var versionsFileUrl = "https://docs.blender.org/PROD/versions.json"
 
 var all_versions;
 
-var Popover = function() {
-  function Popover(id)
+class Popover {
+  constructor(id)
   {
     this.isOpen = false;
     this.type = (id === "version-popover");
-    this.$btn = $('#' + id);
-    this.$dialog = this.$btn.next();
-    this.$list = this.$dialog.children("ul");
+    this.btn = document.querySelector('#' + id);
+    this.dialog = this.btn.nextElementSibling;
+    this.list = this.dialog.querySelector("ul");
     this.sel = null;
-    this.beforeInit();
-  }
-
-  Popover.prototype = {
-    beforeInit : function() {
-      var that = this;
-      this.$btn.on("click", function(e) {
+    const that = this;
+    this.btnClickHandler = function(e) {
+      that.init();
+      e.preventDefault();
+      e.stopPropagation();
+    };
+    this.btnKeyHandler = function(e) {
+      if (that.btnKeyFilter(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");
+      }
+    };
+    this.btn.addEventListener("click", this.btnClickHandler);
+    this.btn.addEventListener("keydown", this.btnKeyHandler);
+  }
 
+  init()
+  {
+    this.btn.removeEventListener("click", this.btnClickHandler);
+    this.btn.removeEventListener("keydown", this.btnKeyHandler);
+
+    new Promise((resolve, reject) => {
       if (all_versions === undefined) {
-        this.$btn.addClass("wait");
-        this.loadVL(this);
+        this.btn.classList.add("wait");
+        fetch(versionsFileUrl)
+            .then((response) => response.json())
+            .then((data) => {
+              all_versions = data;
+              resolve();
+            })
+            .catch(() => {
+              console.error("Version Switch Error: versions.json could not be loaded.");
+              this.btn.classList.remove("disabled");
+            });
       }
       else {
-        this.afterLoad();
+        resolve();
       }
-    },
-    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;
+    }).then(() => {
+      let release = DOCUMENTATION_OPTIONS.VERSION;
       const m = release.match(/\d\.\d+/g);
       if (m) {
         release = m[0];
@@ -65,259 +62,274 @@ var Popover = function() {
 
       this.warnOld(release, all_versions);
 
-      var version = this.getNamed(release);
-      var list = this.buildList(version);
+      const version = this.getNamed(release);
+      this.buildList(version);
 
-      this.$list.children(":first-child").remove();
-      this.$list.append(list);
-      var that = this;
-      this.$list.on("keydown", function(e) {
+      this.list.firstElementChild.remove();
+      const that = this;
+      this.list.addEventListener("keydown", function(e) {
         that.keyMove(e);
       });
 
-      this.$btn.removeClass("wait");
+      this.btn.classList.remove("wait");
       this.btnOpenHandler();
-      this.$btn.on("mousedown", function(e) {
+      this.btn.addEventListener("mousedown", function(e) {
         that.btnOpenHandler();
         e.preventDefault()
       });
-      this.$btn.on("keydown", function(e) {
+      this.btn.addEventListener("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 === "main" || 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");
+    });
+  }
+  warnOld(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.
+    let current = all_versions.current
+    if (!current) {
+      // console.log("Version Switch Error: no 'current' in version.json.");
+      return;
     }
-  };
-  return Popover
-}();
+    const m = current.match(/\d\.\d+/g);
+    if (m) {
+      current = parseFloat(m[0]);
+    }
+    if (release < current) {
+      const currentURL = window.location.pathname.replace(release, current);
+      const warning =
+          document.querySelector("template#version-warning").firstElementChild.cloneNode(true);
+      const link = warning.querySelector('a');
+      link.setAttribute('href', currentURL);
+      link.textContent = current;
 
-$(document).ready(function() {
-  var lng_popover = new Popover("version-popover");
-});
+      let body = document.querySelector("div.body");
+      if (!body.length) {
+        body = document.querySelector("div.document");
+      }
+      body.prepend(warning);
+    }
+  }
+  buildList(v)
+  {
+    const url = new URL(window.location.href);
+    let pathSplit = [ "", "api", v ];
+    if (url.pathname.startsWith("/api/")) {
+      pathSplit.push(url.pathname.split('/').slice(4).join('/'));
+    }
+    else {
+      pathSplit.push(url.pathname.substring(1));
+    }
+    let dyn, cur;
+    if (this.type) {
+      dyn = all_versions;
+      cur = v;
+    }
+    const that = this;
+    const template = document.querySelector("template#version-entry").content;
+    for (let [ix, title] of Object.entries(dyn)) {
+      let clone;
+      if (ix === cur) {
+        clone = template.querySelector("li.selected").cloneNode(true);
+        clone.querySelector("span").innerHTML = title;
+      }
+      else {
+        pathSplit[1 + that.type] = ix;
+        let href = new URL(url);
+        href.pathname = pathSplit.join('/');
+        clone = template.firstElementChild.cloneNode(true);
+        const link = clone.querySelector("a");
+        link.href = href;
+        link.innerHTML = title;
+      }
+      that.list.append(clone);
+    };
+    return this.list;
+  }
+  getNamed(v)
+  {
+    for (let [ix, title] of Object.entries(all_versions)) {
+      if (ix === "master" || ix === "main" || ix === "latest") {
+        const m = title.match(/\d\.\d[\w\d\.]*/)[0];
+        if (parseFloat(m) == v) {
+          v = ix;
+          return false;
+        }
+      }
+    };
+    return v;
+  }
+  dialogToggle(speed)
+  {
+    const wasClose = !this.isOpen;
+    const that = this;
+    if (!this.isOpen) {
+      this.btn.classList.add("version-btn-open");
+      this.btn.setAttribute("aria-pressed", true);
+      this.dialog.setAttribute("aria-hidden", false);
+      this.dialog.style.display = "block";
+      this.dialog.animate({opacity : [ 0, 1 ], easing : [ 'ease-in', 'ease-out' ]}, speed)
+          .finished.then(() => {
+            this.focusoutHandlerPrime = function(e) {
+              that.focusoutHandler();
+              e.stopImmediatePropagation();
+            };
+            this.mouseoutHandlerPrime = function(e) {
+              that.mouseoutHandler();
+              e.stopImmediatePropagation();
+            };
+            this.btn.parentNode.addEventListener("focusout", this.focusoutHandlerPrime);
+            this.btn.parentNode.addEventListener("mouseleave", this.mouseoutHandlerPrime);
+          });
+      this.isOpen = true;
+    }
+    else {
+      this.btn.classList.remove("version-btn-open");
+      this.btn.setAttribute("aria-pressed", false);
+      this.dialog.setAttribute("aria-hidden", true);
+      this.btn.parentNode.removeEventListener("focusout", this.focusoutHandlerPrime);
+      this.btn.parentNode.removeEventListener("mouseleave", this.mouseoutHandlerPrime);
+      this.dialog.animate({opacity : [ 1, 0 ], easing : [ 'ease-in', 'ease-out' ]}, speed)
+          .finished.then(() => {
+            this.dialog.style.display = "none";
+            if (this.sel) {
+              this.sel.setAttribute("tabindex", -1);
+            }
+            this.btn.setAttribute("tabindex", 0);
+            if (document.activeElement !== null && document.activeElement !== document &&
+                document.activeElement !== document.body)
+            {
+              this.btn.focus();
+            }
+          });
+      this.isOpen = false;
+    }
+
+    if (wasClose) {
+      if (this.sel) {
+        this.sel.setAttribute("tabindex", -1);
+      }
+      if (document.activeElement !== null && document.activeElement !== document &&
+          document.activeElement !== document.body)
+      {
+        const nw = this.listEnter();
+        nw.setAttribute("tabindex", 0);
+        nw.focus();
+        this.sel = nw;
+      }
+    }
+  }
+  btnOpenHandler()
+  {
+    this.dialogToggle(300);
+  }
+  focusoutHandler()
+  {
+    const list = this.list;
+    const that = this;
+    setTimeout(function() {
+      if (!list.querySelector(":focus")) {
+        that.dialogToggle(200);
+      }
+    }, 200);
+  }
+  mouseoutHandler()
+  {
+    this.dialogToggle(200);
+  }
+  btnKeyFilter(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(e)
+  {
+    if (e.ctrlKey || e.shiftKey) {
+      return true;
+    }
+    let 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:
+        return false;
+    }
+    nw.setAttribute("tabindex", 0);
+    nw.focus();
+    if (this.sel) {
+      this.sel.setAttribute("tabindex", -1);
+    }
+    this.sel = nw;
+    e.preventDefault();
+    e.stopPropagation();
+  }
+  listPrev(nw)
+  {
+    if (nw.parentNode.previousElementSibling.length !== 0) {
+      return nw.parentNode.previousElementSibling.firstElementChild;
+    }
+    else {
+      return this.listLast();
+    }
+  }
+  listNext(nw)
+  {
+    if (nw.parentNode.nextElementSibling.length !== 0) {
+      return nw.parentNode.nextElementSibling.firstElementChild;
+    }
+    else {
+      return this.listFirst();
+    }
+  }
+  listFirst()
+  {
+    return this.list.firstElementChild.firstElementChild;
+  }
+  listLast()
+  {
+    return this.list.lastElementChild.firstElementChild;
+  }
+  listExit()
+  {
+    this.mouseoutHandler();
+    return this.btn;
+  }
+  listEnter()
+  {
+    return this.list.firstElementChild.firstElementChild;
+  }
+}
+
+document.addEventListener('DOMContentLoaded', () => { new Popover("version-popover"); });
 })();

doc/python_api/templates/base.html

@@ -0,0 +1,6 @@
+{%- extends "!base.html" -%}
+
+{%- block theme_scripts -%}
+{{ super() }}
+<script defer data-domain="docs.blender.org" src="https://analytics.blender.org/js/script.js"></script>
+{%- endblock -%}

doc/python_api/templates/components/footer_contribute.html

@@ -1,6 +1,3 @@
-{# For the "Report Issue" button on the bottom of pages. #}
-{%- extends "!footer.html" %}
-{%- block extrafooter %}
 {%- if not pagename in ("search", "404", "genindex") and hasdoc(pagename) %}
 <div class="footer-contribute">
   <ul>
@@ -13,10 +10,9 @@
                #}Permanent+Link%5D%28https%3A%2F%2Fdocs.blender.org%2F{#
                #}api%2F{{ version }}%2F{{ pagename }}{{ file_suffix }}%29%0D%0A%0D%0A%2A%2A{#
                #}Short+description+of+error%2A%2A%0D%0A%5B{#
-               #}Please+fill+out+a+short+description+of+the+error+here%5D%0D%0A"
-         class="fa fa-bug"> {{ _('Report issue on this page') }}</a>
+               #}Please+fill+out+a+short+description+of+the+error+here%5D%0D%0A" class="fa fa-bug"> {{ _('Report issue
+        on this page') }}</a>
     </li>
   </ul>
 </div>
-{%- endif %}
-{% endblock %}
+{%- endif %}

doc/python_api/templates/page.html

@@ -0,0 +1,6 @@
+{%- extends "!page.html" -%}
+
+{%- block footer -%}
+{{ super() }}
+{%- include "components/footer_contribute.html" -%}
+{%- endblock footer -%}

doc/python_api/templates/sidebar/variant-selector.html

@@ -0,0 +1,29 @@
+<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-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>
+    <template id="version-entry">
+      <li tabindex="-1" role="presentation"><a tabindex="-1" role="menuitem"></a></li>
+      <li class="selected" tabindex="-1" role="presentation"><span tabindex="-1" aria-current="page"></span></li>
+    </template>
+  </ul>
+  <template id="version-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>
+  </template>
+</div>

doc/python_api/templates/versions.html

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

extern/CMakeLists.txt

@@ -32,6 +32,7 @@ endif()
 add_subdirectory(rangetree)
 add_subdirectory(nanosvg)
 add_subdirectory(wcwidth)
+add_subdirectory(xxhash)
 
 if(WITH_BULLET)
   if(NOT WITH_SYSTEM_BULLET)
@@ -104,10 +105,6 @@ if(WITH_MOD_FLUID)
   add_subdirectory(mantaflow)
 endif()
 
-if(WITH_COMPOSITOR_CPU)
-  add_subdirectory(smaa_areatex)
-endif()
-
 if(WITH_VULKAN_BACKEND)
   add_subdirectory(vulkan_memory_allocator)
 endif()

extern/README

@@ -1,4 +1,3 @@
 When updating a library remember to:
 
 * Update the README.blender with the corresponding version.
-* Update the THIRD-PARTY-LICENSE.txt document

extern/audaspace/plugins/pulseaudio/PulseAudioDevice.cpp

@@ -124,13 +124,12 @@ void PulseAudioDevice::playing(bool playing)
 
 	AUD_pa_threaded_mainloop_lock(m_mainloop);
 	AUD_pa_stream_cork(m_stream, playing ? 0 : 1, nullptr, nullptr);
-	AUD_pa_threaded_mainloop_unlock(m_mainloop);
-
 	if(!playing)
 	{
 		AUD_pa_stream_flush(m_stream, nullptr, nullptr);
 		m_clear = true;
 	}
+	AUD_pa_threaded_mainloop_unlock(m_mainloop);
 }
 
 PulseAudioDevice::PulseAudioDevice(const std::string &name, DeviceSpecs specs, int buffersize) :

extern/smaa_areatex/CMakeLists.txt

@@ -1,5 +0,0 @@
-# SPDX-FileCopyrightText: 2017 Blender Foundation
-#
-# SPDX-License-Identifier: GPL-2.0-or-later
-
-add_executable(smaa_areatex smaa_areatex.cpp)

extern/smaa_areatex/README.blender

@@ -1,5 +0,0 @@
-Project: smaa-cpp
-URL: https://github.com/iRi-E/smaa-cpp
-License: MIT
-Upstream version: 0.4.0
-Local modifications:

extern/tinygltf/README.blender

@@ -1,5 +1,5 @@
 Project: TinyGLTF
 URL: https://github.com/syoyo/tinygltf
 License: MIT
-Upstream version: 2.8.3, 84a83d39f55d
+Upstream version: 2.8.21, 4bfc1fc1807e
 Local modifications: None

extern/xxhash/CMakeLists.txt

@@ -0,0 +1,21 @@
+# SPDX-FileCopyrightText: 2024 Blender Authors
+#
+# SPDX-License-Identifier: GPL-2.0-or-later
+
+set(INC
+  PUBLIC .
+)
+
+set(INC_SYS
+)
+
+set(SRC
+  xxhash.c
+  xxhash.h
+)
+
+set(LIB
+)
+
+blender_add_lib(extern_xxhash "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")
+add_library(bf::extern::xxhash ALIAS extern_xxhash)

extern/xxhash/LICENSE

@@ -0,0 +1,26 @@
+xxHash Library
+Copyright (c) 2012-2021 Yann Collet
+All rights reserved.
+
+BSD 2-Clause License (https://www.opensource.org/licenses/bsd-license.php)
+
+Redistribution and use in source and binary forms, with or without modification,
+are permitted provided that the following conditions are met:
+
+* Redistributions of source code must retain the above copyright notice, this
+  list of conditions and the following disclaimer.
+
+* Redistributions in binary form must reproduce the above copyright notice, this
+  list of conditions and the following disclaimer in the documentation and/or
+  other materials provided with the distribution.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR
+ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON
+ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

extern/xxhash/README.blender

@@ -0,0 +1,6 @@
+Project: xxHash
+URL: https://xxhash.com/
+License: BSD 2-Clause
+Upstream version: v0.8.2 (2023-07-21)
+Local modifications:
+* None

extern/xxhash/xxhash.c

@@ -0,0 +1,43 @@
+/*
+ * xxHash - Extremely Fast Hash algorithm
+ * Copyright (C) 2012-2021 Yann Collet
+ *
+ * BSD 2-Clause License (https://www.opensource.org/licenses/bsd-license.php)
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *
+ *    * Redistributions of source code must retain the above copyright
+ *      notice, this list of conditions and the following disclaimer.
+ *    * Redistributions in binary form must reproduce the above
+ *      copyright notice, this list of conditions and the following disclaimer
+ *      in the documentation and/or other materials provided with the
+ *      distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
+ * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
+ * A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+ * OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+ * SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
+ * LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+ * DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+ * THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ * You can contact the author at:
+ *   - xxHash homepage: https://www.xxhash.com
+ *   - xxHash source repository: https://github.com/Cyan4973/xxHash
+ */
+
+
+/*
+ * xxhash.c instantiates functions defined in xxhash.h
+ */
+
+#define XXH_STATIC_LINKING_ONLY   /* access advanced declarations */
+#define XXH_IMPLEMENTATION   /* access definitions */
+
+#include "xxhash.h"

intern/clog/CLG_log.h

@@ -109,12 +109,12 @@ typedef struct CLG_LogRef {
   struct CLG_LogRef *next;
 } CLG_LogRef;
 
-void CLG_log_str(CLG_LogType *lg,
+void CLG_log_str(const CLG_LogType *lg,
                  enum CLG_Severity severity,
                  const char *file_line,
                  const char *fn,
                  const char *message) _CLOG_ATTR_NONNULL(1, 3, 4, 5);
-void CLG_logf(CLG_LogType *lg,
+void CLG_logf(const CLG_LogType *lg,
               enum CLG_Severity severity,
               const char *file_line,
               const char *fn,
@@ -156,7 +156,7 @@ int CLG_color_support_get(CLG_LogRef *clg_ref);
 
 #define CLOG_AT_SEVERITY(clg_ref, severity, verbose_level, ...) \
   { \
-    CLG_LogType *_lg_ty = CLOG_ENSURE(clg_ref); \
+    const CLG_LogType *_lg_ty = CLOG_ENSURE(clg_ref); \
     if (((_lg_ty->flag & CLG_FLAG_USE) && (_lg_ty->level >= verbose_level)) || \
         (severity >= CLG_SEVERITY_WARN)) \
     { \
@@ -167,7 +167,7 @@ int CLG_color_support_get(CLG_LogRef *clg_ref);
 
 #define CLOG_STR_AT_SEVERITY(clg_ref, severity, verbose_level, str) \
   { \
-    CLG_LogType *_lg_ty = CLOG_ENSURE(clg_ref); \
+    const CLG_LogType *_lg_ty = CLOG_ENSURE(clg_ref); \
     if (((_lg_ty->flag & CLG_FLAG_USE) && (_lg_ty->level >= verbose_level)) || \
         (severity >= CLG_SEVERITY_WARN)) \
     { \

intern/clog/clog.c

@@ -175,9 +175,8 @@ static void clg_str_vappendf(CLogStringBuf *cstr, const char *fmt, va_list args)
 {
   /* Use limit because windows may use '-1' for a formatting error. */
   const uint len_max = 65535;
+  uint len_avail = cstr->len_alloc - cstr->len;
   while (true) {
-    uint len_avail = cstr->len_alloc - cstr->len;
-
     va_list args_cpy;
     va_copy(args_cpy, args);
     int retval = vsnprintf(cstr->data + cstr->len, len_avail, fmt, args_cpy);
@@ -188,22 +187,23 @@ static void clg_str_vappendf(CLogStringBuf *cstr, const char *fmt, va_list args)
        * message. */
       break;
     }
-    else if ((uint)retval <= len_avail) {
+
+    if ((uint)retval <= len_avail) {
       /* Copy was successful. */
       cstr->len += (uint)retval;
       break;
     }
-    else {
-      /* vsnprintf was not successful, due to lack of allocated space, retval contains expected
-       * length of the formatted string, use it to allocate required amount of memory. */
-      uint len_alloc = cstr->len + (uint)retval;
-      if (len_alloc >= len_max) {
-        /* Safe upper-limit, just in case... */
-        break;
-      }
-      clg_str_reserve(cstr, len_alloc);
-      len_avail = cstr->len_alloc - cstr->len;
+
+    /* `vsnprintf` was not successful, due to lack of allocated space, `retval` contains expected
+     * length of the formatted string, use it to allocate required amount of memory. */
+    uint len_alloc = cstr->len + (uint)retval;
+    if (len_alloc >= len_max) {
+      /* Safe upper-limit, just in case... */
+      break;
     }
+
+    clg_str_reserve(cstr, len_alloc);
+    len_avail = cstr->len_alloc - cstr->len;
   }
 }
 
@@ -429,7 +429,7 @@ static void write_severity(CLogStringBuf *cstr, enum CLG_Severity severity, bool
   }
 }
 
-static void write_type(CLogStringBuf *cstr, CLG_LogType *lg)
+static void write_type(CLogStringBuf *cstr, const CLG_LogType *lg)
 {
   clg_str_append(cstr, " (");
   clg_str_append(cstr, lg->identifier);
@@ -460,7 +460,7 @@ static void write_file_line_fn(CLogStringBuf *cstr,
   clg_str_append(cstr, ": ");
 }
 
-void CLG_log_str(CLG_LogType *lg,
+void CLG_log_str(const CLG_LogType *lg,
                  enum CLG_Severity severity,
                  const char *file_line,
                  const char *fn,
@@ -498,7 +498,7 @@ void CLG_log_str(CLG_LogType *lg,
   }
 }
 
-void CLG_logf(CLG_LogType *lg,
+void CLG_logf(const CLG_LogType *lg,
               enum CLG_Severity severity,
               const char *file_line,
               const char *fn,

intern/cycles/CMakeLists.txt

@@ -219,10 +219,6 @@ endif()
 
 if(WITH_CYCLES_OSL)
   add_definitions(-DWITH_OSL)
-  # osl 1.9.x
-  add_definitions(-DOSL_STATIC_BUILD)
-  # pre 1.9
-  add_definitions(-DOSL_STATIC_LIBRARY)
   include_directories(
     SYSTEM
     ${OSL_INCLUDE_DIR}

intern/cycles/app/io_export_cycles_xml.py

@@ -72,7 +72,7 @@ class PHYSICS_PT_fluid_export(RenderButtonsPanel, bpy.types.Panel):
 
         cycles = context.scene.cycles_xml
 
-        #layout.prop(cycles, "filepath")
+        # layout.prop(cycles, "filepath")
         layout.operator("export_mesh.cycles_xml")
 
 

intern/cycles/blender/CMakeLists.txt

@@ -142,6 +142,12 @@ if(WITH_OPENIMAGEDENOISE)
   )
 endif()
 
+if(WITH_CYCLES_OSL)
+  list(APPEND LIB
+    ${OSL_LIBRARIES}
+  )
+endif()
+
 blender_add_lib(bf_intern_cycles "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")
 
 add_dependencies(bf_intern_cycles bf_rna)

intern/cycles/blender/addon/properties.py

@@ -235,8 +235,8 @@ def enum_preview_denoiser(self, context):
         items = [
             ('AUTO',
              "Automatic",
-             ("Use the fastest available denoiser for viewport rendering "
-              "(OptiX if available, OpenImageDenoise otherwise)"),
+             ("Use GPU accelerated denoising if supported, for the best performance. "
+              "Prefer OpenImageDenoise over OptiX"),
              0)]
     else:
         items = [('AUTO', "None", "Blender was compiled without a viewport denoiser", 0)]
@@ -1594,14 +1594,15 @@ class CyclesPreferences(bpy.types.AddonPreferences):
         compute_device_type = self.get_compute_device_type()
 
         # We need non-CPU devices, used for rendering and supporting OIDN GPU denoising
-        for device in _cycles.available_devices(compute_device_type):
-            device_type = device[1]
-            if device_type == 'CPU':
-                continue
+        if compute_device_type != 'NONE':
+            for device in _cycles.available_devices(compute_device_type):
+                device_type = device[1]
+                if device_type == 'CPU':
+                    continue
 
-            has_device_oidn_support = device[5]
-            if has_device_oidn_support and self.find_existing_device_entry(device).use:
-                return True
+                has_device_oidn_support = device[5]
+                if has_device_oidn_support and self.find_existing_device_entry(device).use:
+                    return True
 
         return False
 

intern/cycles/blender/display_driver.cpp

@@ -2,11 +2,11 @@
  *
  * SPDX-License-Identifier: Apache-2.0 */
 
-#include "GPU_context.h"
-#include "GPU_immediate.h"
-#include "GPU_shader.h"
-#include "GPU_state.h"
-#include "GPU_texture.h"
+#include "GPU_context.hh"
+#include "GPU_immediate.hh"
+#include "GPU_shader.hh"
+#include "GPU_state.hh"
+#include "GPU_texture.hh"
 
 #include "RE_engine.h"
 

intern/cycles/blender/python.cpp

@@ -26,7 +26,7 @@
 #include "util/tbb.h"
 #include "util/types.h"
 
-#include "GPU_state.h"
+#include "GPU_state.hh"
 
 #ifdef WITH_OSL
 #  include "scene/osl.h"

intern/cycles/bvh/optix.cpp

@@ -29,7 +29,7 @@ BVHOptiX::~BVHOptiX()
 {
   /* Acceleration structure memory is delayed freed on device, since deleting the
    * BVH may happen while still being used for rendering. */
-  device->release_optix_bvh(this);
+  device->release_bvh(this);
 }
 
 CCL_NAMESPACE_END

intern/cycles/device/CMakeLists.txt

@@ -203,6 +203,13 @@ if(WITH_OPENIMAGEDENOISE)
   )
 endif()
 
+if(WITH_CYCLES_OSL)
+  list(APPEND LIB
+    ${OSL_LIBRARIES}
+  )
+endif()
+
+
 include_directories(${INC})
 include_directories(SYSTEM ${INC_SYS})
 

intern/cycles/device/cuda/device.cpp

@@ -198,7 +198,7 @@ void device_cuda_info(vector<DeviceInfo> &devices)
     VLOG_INFO << "Added device \"" << info.description << "\" with id \"" << info.id << "\".";
 
     if (info.denoisers & DENOISER_OPENIMAGEDENOISE)
-      VLOG_INFO << "Device with id \"" << info.id << "\" is supporting "
+      VLOG_INFO << "Device with id \"" << info.id << "\" supports "
                 << denoiserTypeToHumanReadable(DENOISER_OPENIMAGEDENOISE) << ".";
   }
 

intern/cycles/device/device.h

@@ -124,7 +124,8 @@ class DeviceInfo {
     /* Multiple Devices with the same ID would be very bad. */
     assert(id != info.id ||
            (type == info.type && num == info.num && description == info.description));
-    return id == info.id;
+    return id == info.id && use_hardware_raytracing == info.use_hardware_raytracing &&
+           kernel_optimization_level == info.kernel_optimization_level;
   }
 };
 
@@ -217,11 +218,10 @@ class Device {
   /* Get OpenShadingLanguage memory buffer. */
   virtual void *get_cpu_osl_memory();
 
-  /* acceleration structure building */
+  /* Acceleration structure building. */
   virtual void build_bvh(BVH *bvh, Progress &progress, bool refit);
-
-  /* OptiX specific destructor. */
-  virtual void release_optix_bvh(BVH * /*bvh*/){};
+  /* Used by Metal and OptiX. */
+  virtual void release_bvh(BVH * /*bvh*/) {}
 
   /* multi device */
   virtual int device_number(Device * /*sub_device*/)

intern/cycles/device/hip/device.cpp

@@ -211,7 +211,7 @@ void device_hip_info(vector<DeviceInfo> &devices)
     VLOG_INFO << "Added device \"" << info.description << "\" with id \"" << info.id << "\".";
 
     if (info.denoisers & DENOISER_OPENIMAGEDENOISE)
-      VLOG_INFO << "Device with id \"" << info.id << "\" is supporting "
+      VLOG_INFO << "Device with id \"" << info.id << "\" supports "
                 << denoiserTypeToHumanReadable(DENOISER_OPENIMAGEDENOISE) << ".";
   }
 

intern/cycles/device/metal/bvh.h

@@ -28,9 +28,9 @@ class BVHMetal : public BVH {
   API_AVAILABLE(macos(11.0))
   vector<id<MTLAccelerationStructure>> unique_blas_array;
 
-  bool motion_blur = false;
+  Device *device = nullptr;
 
-  Stats &stats;
+  bool motion_blur = false;
 
   bool build(Progress &progress, id<MTLDevice> device, id<MTLCommandQueue> queue, bool refit);
 

intern/cycles/device/metal/bvh.mm

@@ -113,15 +113,18 @@ BVHMetal::BVHMetal(const BVHParams &params_,
                    const vector<Geometry *> &geometry_,
                    const vector<Object *> &objects_,
                    Device *device)
-    : BVH(params_, geometry_, objects_), stats(device->stats)
+    : BVH(params_, geometry_, objects_), device(device)
 {
 }
 
 BVHMetal::~BVHMetal()
 {
+  /* Clear point used by enqueueing. */
+  device->release_bvh(this);
+
   if (@available(macos 12.0, *)) {
     if (accel_struct) {
-      stats.mem_free(accel_struct.allocatedSize);
+      device->stats.mem_free(accel_struct.allocatedSize);
       [accel_struct release];
     }
 
@@ -132,7 +135,7 @@ BVHMetal::~BVHMetal()
 }
 
 bool BVHMetal::build_BLAS_mesh(Progress &progress,
-                               id<MTLDevice> device,
+                               id<MTLDevice> mtl_device,
                                id<MTLCommandQueue> queue,
                                Geometry *const geom,
                                bool refit)
@@ -163,7 +166,7 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
     }
 
     MTLResourceOptions storage_mode;
-    if (device.hasUnifiedMemory) {
+    if (mtl_device.hasUnifiedMemory) {
       storage_mode = MTLResourceStorageModeShared;
     }
     else {
@@ -172,18 +175,19 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
 
     /* Upload the mesh data to the GPU */
     id<MTLBuffer> posBuf = nil;
-    id<MTLBuffer> indexBuf = [device newBufferWithBytes:tris.data()
-                                                 length:num_indices * sizeof(tris.data()[0])
-                                                options:storage_mode];
+    id<MTLBuffer> indexBuf = [mtl_device newBufferWithBytes:tris.data()
+                                                     length:num_indices * sizeof(tris.data()[0])
+                                                    options:storage_mode];
 
     if (num_motion_steps == 1) {
-      posBuf = [device newBufferWithBytes:verts.data()
-                                   length:num_verts * sizeof(verts.data()[0])
-                                  options:storage_mode];
+      posBuf = [mtl_device newBufferWithBytes:verts.data()
+                                       length:num_verts * sizeof(verts.data()[0])
+                                      options:storage_mode];
     }
     else {
-      posBuf = [device newBufferWithLength:num_verts * num_motion_steps * sizeof(verts.data()[0])
-                                   options:storage_mode];
+      posBuf = [mtl_device
+          newBufferWithLength:num_verts * num_motion_steps * sizeof(verts.data()[0])
+                      options:storage_mode];
       float3 *dest_data = (float3 *)[posBuf contents];
       size_t center_step = (num_motion_steps - 1) / 2;
       for (size_t step = 0; step < num_motion_steps; ++step) {
@@ -264,13 +268,14 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
                           MTLAccelerationStructureUsagePreferFastBuild);
     }
 
-    MTLAccelerationStructureSizes accelSizes = [device
+    MTLAccelerationStructureSizes accelSizes = [mtl_device
         accelerationStructureSizesWithDescriptor:accelDesc];
-    id<MTLAccelerationStructure> accel_uncompressed = [device
+    id<MTLAccelerationStructure> accel_uncompressed = [mtl_device
         newAccelerationStructureWithSize:accelSizes.accelerationStructureSize];
-    id<MTLBuffer> scratchBuf = [device newBufferWithLength:accelSizes.buildScratchBufferSize
-                                                   options:MTLResourceStorageModePrivate];
-    id<MTLBuffer> sizeBuf = [device newBufferWithLength:8 options:MTLResourceStorageModeShared];
+    id<MTLBuffer> scratchBuf = [mtl_device newBufferWithLength:accelSizes.buildScratchBufferSize
+                                                       options:MTLResourceStorageModePrivate];
+    id<MTLBuffer> sizeBuf = [mtl_device newBufferWithLength:8
+                                                    options:MTLResourceStorageModeShared];
     id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
     id<MTLAccelerationStructureCommandEncoder> accelEnc =
         [accelCommands accelerationStructureCommandEncoder];
@@ -314,14 +319,14 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
           id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
           id<MTLAccelerationStructureCommandEncoder> accelEnc =
               [accelCommands accelerationStructureCommandEncoder];
-          id<MTLAccelerationStructure> accel = [device
+          id<MTLAccelerationStructure> accel = [mtl_device
               newAccelerationStructureWithSize:compressed_size];
           [accelEnc copyAndCompactAccelerationStructure:accel_uncompressed
                                 toAccelerationStructure:accel];
           [accelEnc endEncoding];
           [accelCommands addCompletedHandler:^(id<MTLCommandBuffer> /*command_buffer*/) {
             uint64_t allocated_size = [accel allocatedSize];
-            stats.mem_alloc(allocated_size);
+            device->stats.mem_alloc(allocated_size);
             accel_struct = accel;
             [accel_uncompressed release];
 
@@ -336,7 +341,7 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
         accel_struct = accel_uncompressed;
 
         uint64_t allocated_size = [accel_struct allocatedSize];
-        stats.mem_alloc(allocated_size);
+        device->stats.mem_alloc(allocated_size);
 
         /* Signal that we've finished doing GPU acceleration struct build. */
         g_bvh_build_throttler.release(wired_size);
@@ -355,7 +360,7 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
 }
 
 bool BVHMetal::build_BLAS_hair(Progress &progress,
-                               id<MTLDevice> device,
+                               id<MTLDevice> mtl_device,
                                id<MTLCommandQueue> queue,
                                Geometry *const geom,
                                bool refit)
@@ -382,7 +387,7 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
     }
 
     MTLResourceOptions storage_mode;
-    if (device.hasUnifiedMemory) {
+    if (mtl_device.hasUnifiedMemory) {
       storage_mode = MTLResourceStorageModeShared;
     }
     else {
@@ -445,18 +450,18 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
       }
 
       /* Allocate and populate MTLBuffers for geometry. */
-      idxBuffer = [device newBufferWithBytes:idxData.data()
-                                      length:idxData.size() * sizeof(int)
-                                     options:storage_mode];
+      idxBuffer = [mtl_device newBufferWithBytes:idxData.data()
+                                          length:idxData.size() * sizeof(int)
+                                         options:storage_mode];
 
-      cpBuffer = [device newBufferWithBytes:cpData.data()
-                                     length:cpData.size() * sizeof(float3)
-                                    options:storage_mode];
-
-      radiusBuffer = [device newBufferWithBytes:radiusData.data()
-                                         length:radiusData.size() * sizeof(float)
+      cpBuffer = [mtl_device newBufferWithBytes:cpData.data()
+                                         length:cpData.size() * sizeof(float3)
                                         options:storage_mode];
 
+      radiusBuffer = [mtl_device newBufferWithBytes:radiusData.data()
+                                             length:radiusData.size() * sizeof(float)
+                                            options:storage_mode];
+
       std::vector<MTLMotionKeyframeData *> cp_ptrs;
       std::vector<MTLMotionKeyframeData *> radius_ptrs;
       cp_ptrs.reserve(num_motion_steps);
@@ -541,18 +546,18 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
       }
 
       /* Allocate and populate MTLBuffers for geometry. */
-      idxBuffer = [device newBufferWithBytes:idxData.data()
-                                      length:idxData.size() * sizeof(int)
-                                     options:storage_mode];
+      idxBuffer = [mtl_device newBufferWithBytes:idxData.data()
+                                          length:idxData.size() * sizeof(int)
+                                         options:storage_mode];
 
-      cpBuffer = [device newBufferWithBytes:cpData.data()
-                                     length:cpData.size() * sizeof(float3)
-                                    options:storage_mode];
-
-      radiusBuffer = [device newBufferWithBytes:radiusData.data()
-                                         length:radiusData.size() * sizeof(float)
+      cpBuffer = [mtl_device newBufferWithBytes:cpData.data()
+                                         length:cpData.size() * sizeof(float3)
                                         options:storage_mode];
 
+      radiusBuffer = [mtl_device newBufferWithBytes:radiusData.data()
+                                             length:radiusData.size() * sizeof(float)
+                                            options:storage_mode];
+
       if (storage_mode == MTLResourceStorageModeManaged) {
         [cpBuffer didModifyRange:NSMakeRange(0, cpBuffer.length)];
         [idxBuffer didModifyRange:NSMakeRange(0, idxBuffer.length)];
@@ -600,13 +605,14 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
     }
     accelDesc.usage |= MTLAccelerationStructureUsageExtendedLimits;
 
-    MTLAccelerationStructureSizes accelSizes = [device
+    MTLAccelerationStructureSizes accelSizes = [mtl_device
         accelerationStructureSizesWithDescriptor:accelDesc];
-    id<MTLAccelerationStructure> accel_uncompressed = [device
+    id<MTLAccelerationStructure> accel_uncompressed = [mtl_device
         newAccelerationStructureWithSize:accelSizes.accelerationStructureSize];
-    id<MTLBuffer> scratchBuf = [device newBufferWithLength:accelSizes.buildScratchBufferSize
-                                                   options:MTLResourceStorageModePrivate];
-    id<MTLBuffer> sizeBuf = [device newBufferWithLength:8 options:MTLResourceStorageModeShared];
+    id<MTLBuffer> scratchBuf = [mtl_device newBufferWithLength:accelSizes.buildScratchBufferSize
+                                                       options:MTLResourceStorageModePrivate];
+    id<MTLBuffer> sizeBuf = [mtl_device newBufferWithLength:8
+                                                    options:MTLResourceStorageModeShared];
     id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
     id<MTLAccelerationStructureCommandEncoder> accelEnc =
         [accelCommands accelerationStructureCommandEncoder];
@@ -651,14 +657,14 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
           id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
           id<MTLAccelerationStructureCommandEncoder> accelEnc =
               [accelCommands accelerationStructureCommandEncoder];
-          id<MTLAccelerationStructure> accel = [device
+          id<MTLAccelerationStructure> accel = [mtl_device
               newAccelerationStructureWithSize:compressed_size];
           [accelEnc copyAndCompactAccelerationStructure:accel_uncompressed
                                 toAccelerationStructure:accel];
           [accelEnc endEncoding];
           [accelCommands addCompletedHandler:^(id<MTLCommandBuffer> /*command_buffer*/) {
             uint64_t allocated_size = [accel allocatedSize];
-            stats.mem_alloc(allocated_size);
+            device->stats.mem_alloc(allocated_size);
             accel_struct = accel;
             [accel_uncompressed release];
 
@@ -673,7 +679,7 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
         accel_struct = accel_uncompressed;
 
         uint64_t allocated_size = [accel_struct allocatedSize];
-        stats.mem_alloc(allocated_size);
+        device->stats.mem_alloc(allocated_size);
 
         /* Signal that we've finished doing GPU acceleration struct build. */
         g_bvh_build_throttler.release(wired_size);
@@ -690,7 +696,7 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
   }
 #  else  /* MAC_OS_VERSION_14_0 */
   (void)progress;
-  (void)device;
+  (void)mtl_device;
   (void)queue;
   (void)geom;
   (void)(refit);
@@ -699,7 +705,7 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
 }
 
 bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
-                                     id<MTLDevice> device,
+                                     id<MTLDevice> mtl_device,
                                      id<MTLCommandQueue> queue,
                                      Geometry *const geom,
                                      bool refit)
@@ -732,7 +738,7 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
     const size_t num_aabbs = num_motion_steps * num_points;
 
     MTLResourceOptions storage_mode;
-    if (device.hasUnifiedMemory) {
+    if (mtl_device.hasUnifiedMemory) {
       storage_mode = MTLResourceStorageModeShared;
     }
     else {
@@ -740,7 +746,7 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
     }
 
     /* Allocate a GPU buffer for the AABB data and populate it */
-    id<MTLBuffer> aabbBuf = [device
+    id<MTLBuffer> aabbBuf = [mtl_device
         newBufferWithLength:num_aabbs * sizeof(MTLAxisAlignedBoundingBox)
                     options:storage_mode];
     MTLAxisAlignedBoundingBox *aabb_data = (MTLAxisAlignedBoundingBox *)[aabbBuf contents];
@@ -848,13 +854,14 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
                           MTLAccelerationStructureUsagePreferFastBuild);
     }
 
-    MTLAccelerationStructureSizes accelSizes = [device
+    MTLAccelerationStructureSizes accelSizes = [mtl_device
         accelerationStructureSizesWithDescriptor:accelDesc];
-    id<MTLAccelerationStructure> accel_uncompressed = [device
+    id<MTLAccelerationStructure> accel_uncompressed = [mtl_device
         newAccelerationStructureWithSize:accelSizes.accelerationStructureSize];
-    id<MTLBuffer> scratchBuf = [device newBufferWithLength:accelSizes.buildScratchBufferSize
-                                                   options:MTLResourceStorageModePrivate];
-    id<MTLBuffer> sizeBuf = [device newBufferWithLength:8 options:MTLResourceStorageModeShared];
+    id<MTLBuffer> scratchBuf = [mtl_device newBufferWithLength:accelSizes.buildScratchBufferSize
+                                                       options:MTLResourceStorageModePrivate];
+    id<MTLBuffer> sizeBuf = [mtl_device newBufferWithLength:8
+                                                    options:MTLResourceStorageModeShared];
     id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
     id<MTLAccelerationStructureCommandEncoder> accelEnc =
         [accelCommands accelerationStructureCommandEncoder];
@@ -897,14 +904,14 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
           id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
           id<MTLAccelerationStructureCommandEncoder> accelEnc =
               [accelCommands accelerationStructureCommandEncoder];
-          id<MTLAccelerationStructure> accel = [device
+          id<MTLAccelerationStructure> accel = [mtl_device
               newAccelerationStructureWithSize:compressed_size];
           [accelEnc copyAndCompactAccelerationStructure:accel_uncompressed
                                 toAccelerationStructure:accel];
           [accelEnc endEncoding];
           [accelCommands addCompletedHandler:^(id<MTLCommandBuffer> /*command_buffer*/) {
             uint64_t allocated_size = [accel allocatedSize];
-            stats.mem_alloc(allocated_size);
+            device->stats.mem_alloc(allocated_size);
             accel_struct = accel;
             [accel_uncompressed release];
 
@@ -919,7 +926,7 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
         accel_struct = accel_uncompressed;
 
         uint64_t allocated_size = [accel_struct allocatedSize];
-        stats.mem_alloc(allocated_size);
+        device->stats.mem_alloc(allocated_size);
 
         /* Signal that we've finished doing GPU acceleration struct build. */
         g_bvh_build_throttler.release(wired_size);
@@ -937,7 +944,7 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
 }
 
 bool BVHMetal::build_BLAS(Progress &progress,
-                          id<MTLDevice> device,
+                          id<MTLDevice> mtl_device,
                           id<MTLCommandQueue> queue,
                           bool refit)
 {
@@ -948,11 +955,11 @@ bool BVHMetal::build_BLAS(Progress &progress,
   switch (geom->geometry_type) {
     case Geometry::VOLUME:
     case Geometry::MESH:
-      return build_BLAS_mesh(progress, device, queue, geom, refit);
+      return build_BLAS_mesh(progress, mtl_device, queue, geom, refit);
     case Geometry::HAIR:
-      return build_BLAS_hair(progress, device, queue, geom, refit);
+      return build_BLAS_hair(progress, mtl_device, queue, geom, refit);
     case Geometry::POINTCLOUD:
-      return build_BLAS_pointcloud(progress, device, queue, geom, refit);
+      return build_BLAS_pointcloud(progress, mtl_device, queue, geom, refit);
     default:
       return false;
   }
@@ -960,7 +967,7 @@ bool BVHMetal::build_BLAS(Progress &progress,
 }
 
 bool BVHMetal::build_TLAS(Progress &progress,
-                          id<MTLDevice> device,
+                          id<MTLDevice> mtl_device,
                           id<MTLCommandQueue> queue,
                           bool refit)
 {
@@ -969,14 +976,14 @@ bool BVHMetal::build_TLAS(Progress &progress,
 
   if (@available(macos 12.0, *)) {
     /* Defined inside available check, for return type to be available. */
-    auto make_null_BLAS = [](id<MTLDevice> device,
+    auto make_null_BLAS = [](id<MTLDevice> mtl_device,
                              id<MTLCommandQueue> queue) -> id<MTLAccelerationStructure> {
       MTLResourceOptions storage_mode = MTLResourceStorageModeManaged;
-      if (device.hasUnifiedMemory) {
+      if (mtl_device.hasUnifiedMemory) {
         storage_mode = MTLResourceStorageModeShared;
       }
 
-      id<MTLBuffer> nullBuf = [device newBufferWithLength:sizeof(float3) options:storage_mode];
+      id<MTLBuffer> nullBuf = [mtl_device newBufferWithLength:sizeof(float3) options:storage_mode];
 
       /* Create an acceleration structure. */
       MTLAccelerationStructureTriangleGeometryDescriptor *geomDesc =
@@ -997,13 +1004,14 @@ bool BVHMetal::build_TLAS(Progress &progress,
       accelDesc.geometryDescriptors = @[ geomDesc ];
       accelDesc.usage |= MTLAccelerationStructureUsageExtendedLimits;
 
-      MTLAccelerationStructureSizes accelSizes = [device
+      MTLAccelerationStructureSizes accelSizes = [mtl_device
           accelerationStructureSizesWithDescriptor:accelDesc];
-      id<MTLAccelerationStructure> accel_struct = [device
+      id<MTLAccelerationStructure> accel_struct = [mtl_device
           newAccelerationStructureWithSize:accelSizes.accelerationStructureSize];
-      id<MTLBuffer> scratchBuf = [device newBufferWithLength:accelSizes.buildScratchBufferSize
-                                                     options:MTLResourceStorageModePrivate];
-      id<MTLBuffer> sizeBuf = [device newBufferWithLength:8 options:MTLResourceStorageModeShared];
+      id<MTLBuffer> scratchBuf = [mtl_device newBufferWithLength:accelSizes.buildScratchBufferSize
+                                                         options:MTLResourceStorageModePrivate];
+      id<MTLBuffer> sizeBuf = [mtl_device newBufferWithLength:8
+                                                      options:MTLResourceStorageModeShared];
       id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
       id<MTLAccelerationStructureCommandEncoder> accelEnc =
           [accelCommands accelerationStructureCommandEncoder];
@@ -1070,7 +1078,7 @@ bool BVHMetal::build_TLAS(Progress &progress,
     };
 
     MTLResourceOptions storage_mode;
-    if (device.hasUnifiedMemory) {
+    if (mtl_device.hasUnifiedMemory) {
       storage_mode = MTLResourceStorageModeShared;
     }
     else {
@@ -1086,12 +1094,12 @@ bool BVHMetal::build_TLAS(Progress &progress,
     }
 
     /* Allocate a GPU buffer for the instance data and populate it */
-    id<MTLBuffer> instanceBuf = [device newBufferWithLength:num_instances * instance_size
-                                                    options:storage_mode];
+    id<MTLBuffer> instanceBuf = [mtl_device newBufferWithLength:num_instances * instance_size
+                                                        options:storage_mode];
     id<MTLBuffer> motion_transforms_buf = nil;
     MTLPackedFloat4x3 *motion_transforms = nullptr;
     if (motion_blur && num_motion_transforms) {
-      motion_transforms_buf = [device
+      motion_transforms_buf = [mtl_device
           newBufferWithLength:num_motion_transforms * sizeof(MTLPackedFloat4x3)
                       options:storage_mode];
       motion_transforms = (MTLPackedFloat4x3 *)motion_transforms_buf.contents;
@@ -1108,14 +1116,14 @@ bool BVHMetal::build_TLAS(Progress &progress,
       Geometry const *geom = ob->get_geometry();
       BVHMetal const *blas = static_cast<BVHMetal const *>(geom->bvh);
       if (!blas || !blas->accel_struct) {
-        /* Place a degenerate instance, to ensure [[instance_id]] equals ob->get_device_index()
+        /* Place a degenerate instance, to ensure [[instance_id]] equals ob->get_mtl_device_index()
          * in our intersection functions */
         blas = nullptr;
 
         /* Workaround for issue in macOS <= 14.1: Insert degenerate BLAS instead of zero-filling
          * the descriptor. */
         if (!null_BLAS) {
-          null_BLAS = make_null_BLAS(device, queue);
+          null_BLAS = make_null_BLAS(mtl_device, queue);
         }
         blas_array.push_back(null_BLAS);
       }
@@ -1259,12 +1267,12 @@ bool BVHMetal::build_TLAS(Progress &progress,
                           MTLAccelerationStructureUsagePreferFastBuild);
     }
 
-    MTLAccelerationStructureSizes accelSizes = [device
+    MTLAccelerationStructureSizes accelSizes = [mtl_device
         accelerationStructureSizesWithDescriptor:accelDesc];
-    id<MTLAccelerationStructure> accel = [device
+    id<MTLAccelerationStructure> accel = [mtl_device
         newAccelerationStructureWithSize:accelSizes.accelerationStructureSize];
-    id<MTLBuffer> scratchBuf = [device newBufferWithLength:accelSizes.buildScratchBufferSize
-                                                   options:MTLResourceStorageModePrivate];
+    id<MTLBuffer> scratchBuf = [mtl_device newBufferWithLength:accelSizes.buildScratchBufferSize
+                                                       options:MTLResourceStorageModePrivate];
     id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
     id<MTLAccelerationStructureCommandEncoder> accelEnc =
         [accelCommands accelerationStructureCommandEncoder];
@@ -1292,7 +1300,7 @@ bool BVHMetal::build_TLAS(Progress &progress,
     [scratchBuf release];
 
     uint64_t allocated_size = [accel allocatedSize];
-    stats.mem_alloc(allocated_size);
+    device->stats.mem_alloc(allocated_size);
 
     /* Cache top and bottom-level acceleration structs */
     accel_struct = accel;
@@ -1309,7 +1317,7 @@ bool BVHMetal::build_TLAS(Progress &progress,
 }
 
 bool BVHMetal::build(Progress &progress,
-                     id<MTLDevice> device,
+                     id<MTLDevice> mtl_device,
                      id<MTLCommandQueue> queue,
                      bool refit)
 {
@@ -1319,7 +1327,7 @@ bool BVHMetal::build(Progress &progress,
     }
     else {
       if (accel_struct) {
-        stats.mem_free(accel_struct.allocatedSize);
+        device->stats.mem_free(accel_struct.allocatedSize);
         [accel_struct release];
         accel_struct = nil;
       }
@@ -1328,10 +1336,10 @@ bool BVHMetal::build(Progress &progress,
 
   @autoreleasepool {
     if (!params.top_level) {
-      return build_BLAS(progress, device, queue, refit);
+      return build_BLAS(progress, mtl_device, queue, refit);
     }
     else {
-      return build_TLAS(progress, device, queue, refit);
+      return build_TLAS(progress, mtl_device, queue, refit);
     }
   }
 }

intern/cycles/device/metal/device.mm

@@ -89,7 +89,7 @@ void device_metal_info(vector<DeviceInfo> &devices)
     VLOG_INFO << "Added device \"" << info.description << "\" with id \"" << info.id << "\".";
 
     if (info.denoisers & DENOISER_OPENIMAGEDENOISE)
-      VLOG_INFO << "Device with id \"" << info.id << "\" is supporting "
+      VLOG_INFO << "Device with id \"" << info.id << "\" supports "
                 << denoiserTypeToHumanReadable(DENOISER_OPENIMAGEDENOISE) << ".";
   }
 }

intern/cycles/device/metal/device_impl.h

@@ -140,6 +140,8 @@ class MetalDevice : public Device {
 
   virtual void build_bvh(BVH *bvh, Progress &progress, bool refit) override;
 
+  virtual void release_bvh(BVH *bvh) override;
+
   virtual void optimize_for_scene(Scene *scene) override;
 
   static void compile_and_load(int device_id, MetalPipelineType pso_type);

intern/cycles/device/metal/device_impl.mm

@@ -341,7 +341,7 @@ string MetalDevice::preprocess_source(MetalPipelineType pso_type,
   }
 
 #  ifdef WITH_CYCLES_DEBUG
-  global_defines += "#define __KERNEL_DEBUG__\n";
+  global_defines += "#define WITH_CYCLES_DEBUG\n";
 #  endif
 
   switch (device_vendor) {
@@ -1443,6 +1443,13 @@ void MetalDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
   }
 }
 
+void MetalDevice::release_bvh(BVH *bvh)
+{
+  if (bvhMetalRT == bvh) {
+    bvhMetalRT = nullptr;
+  }
+}
+
 CCL_NAMESPACE_END
 
 #endif

intern/cycles/device/metal/kernel.mm

@@ -106,7 +106,7 @@ struct ShaderCache {
 
   friend ShaderCache *get_shader_cache(id<MTLDevice> mtlDevice);
 
-  void compile_thread_func(int thread_index);
+  void compile_thread_func();
 
   using PipelineCollection = std::vector<unique_ptr<MetalKernelPipeline>>;
 
@@ -174,7 +174,7 @@ void ShaderCache::wait_for_all()
   }
 }
 
-void ShaderCache::compile_thread_func(int /*thread_index*/)
+void ShaderCache::compile_thread_func()
 {
   while (running) {
 
@@ -309,7 +309,7 @@ void ShaderCache::load_kernel(DeviceKernel device_kernel,
 
       metal_printf("Spawning %d Cycles kernel compilation threads\n", max_mtlcompiler_threads);
       for (int i = 0; i < max_mtlcompiler_threads; i++) {
-        compile_threads.push_back(std::thread([&] { compile_thread_func(i); }));
+        compile_threads.push_back(std::thread([this] { this->compile_thread_func(); }));
       }
     }
   }

intern/cycles/device/oneapi/device.cpp

@@ -134,7 +134,7 @@ static void device_iterator_cb(
   VLOG_INFO << "Added device \"" << info.description << "\" with id \"" << info.id << "\".";
 
   if (info.denoisers & DENOISER_OPENIMAGEDENOISE)
-    VLOG_INFO << "Device with id \"" << info.id << "\" is supporting "
+    VLOG_INFO << "Device with id \"" << info.id << "\" supports "
               << denoiserTypeToHumanReadable(DENOISER_OPENIMAGEDENOISE) << ".";
 }
 #endif

intern/cycles/device/oneapi/device_impl.cpp

@@ -773,8 +773,9 @@ void OneapiDevice::get_adjusted_global_and_local_sizes(SyclQueue *queue,
                                                        size_t &kernel_local_size)
 {
   assert(queue);
-
-  const static size_t preferred_work_group_size_intersect_shading = 32;
+  const static size_t preferred_work_group_size_intersect = 128;
+  const static size_t preferred_work_group_size_shading = 256;
+  const static size_t preferred_work_group_size_shading_simd8 = 64;
   /* Shader evaluation kernels seems to use some amount of shared memory, so better
    * to avoid usage of maximum work group sizes for them. */
   const static size_t preferred_work_group_size_shader_evaluation = 256;
@@ -783,6 +784,9 @@ void OneapiDevice::get_adjusted_global_and_local_sizes(SyclQueue *queue,
   const static size_t preferred_work_group_size_cryptomatte = 512;
   const static size_t preferred_work_group_size_default = 1024;
 
+  const sycl::device &device = reinterpret_cast<sycl::queue *>(queue)->get_device();
+  const size_t max_work_group_size = device.get_info<sycl::info::device::max_work_group_size>();
+
   size_t preferred_work_group_size = 0;
   switch (kernel) {
     case DEVICE_KERNEL_INTEGRATOR_INIT_FROM_CAMERA:
@@ -792,6 +796,9 @@ void OneapiDevice::get_adjusted_global_and_local_sizes(SyclQueue *queue,
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
+      preferred_work_group_size = preferred_work_group_size_intersect;
+      break;
+
     case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND:
     case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
     case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE:
@@ -799,9 +806,13 @@ void OneapiDevice::get_adjusted_global_and_local_sizes(SyclQueue *queue,
     case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:
     case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME:
     case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW:
-    case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT:
-      preferred_work_group_size = preferred_work_group_size_intersect_shading;
-      break;
+    case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT: {
+      const bool device_is_simd8 =
+          (device.has(sycl::aspect::ext_intel_gpu_eu_simd_width) &&
+           device.get_info<sycl::ext::intel::info::device::gpu_eu_simd_width>() == 8);
+      preferred_work_group_size = (device_is_simd8) ? preferred_work_group_size_shading_simd8 :
+                                                      preferred_work_group_size_shading;
+    } break;
 
     case DEVICE_KERNEL_CRYPTOMATTE_POSTPROCESS:
       preferred_work_group_size = preferred_work_group_size_cryptomatte;
@@ -829,11 +840,7 @@ void OneapiDevice::get_adjusted_global_and_local_sizes(SyclQueue *queue,
     preferred_work_group_size = preferred_work_group_size_default;
   }
 
-  const size_t limit_work_group_size = reinterpret_cast<sycl::queue *>(queue)
-                                           ->get_device()
-                                           .get_info<sycl::info::device::max_work_group_size>();
-
-  kernel_local_size = std::min(limit_work_group_size, preferred_work_group_size);
+  kernel_local_size = std::min(max_work_group_size, preferred_work_group_size);
 
   /* NOTE(@nsirgien): As for now non-uniform work-groups don't work on most oneAPI devices,
    * we extend work size to fit uniformity requirements. */

intern/cycles/device/optix/device_impl.cpp

@@ -1680,7 +1680,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
   }
 }
 
-void OptiXDevice::release_optix_bvh(BVH *bvh)
+void OptiXDevice::release_bvh(BVH *bvh)
 {
   thread_scoped_lock lock(delayed_free_bvh_mutex);
   /* Do delayed free of BVH memory, since geometry holding BVH might be deleted

intern/cycles/device/optix/device_impl.h

@@ -106,7 +106,7 @@ class OptiXDevice : public CUDADevice {
 
   void build_bvh(BVH *bvh, Progress &progress, bool refit) override;
 
-  void release_optix_bvh(BVH *bvh) override;
+  void release_bvh(BVH *bvh) override;
   void free_bvh_memory_delayed();
 
   void const_copy_to(const char *name, void *host, size_t size) override;

intern/cycles/integrator/denoiser_oidn_gpu.cpp

@@ -46,7 +46,7 @@ static const char *oidn_device_type_to_string(const OIDNDeviceType type)
       return "HIP";
 #  endif
 
-      /* The Metal support was added in OIDN 2.2.*/
+      /* The Metal support was added in OIDN 2.2. */
 #  if (OIDN_VERSION_MAJOR > 2) || ((OIDN_VERSION_MAJOR == 2) && (OIDN_VERSION_MINOR >= 2))
     case OIDN_DEVICE_TYPE_METAL:
       return "METAL";

intern/cycles/kernel/camera/camera.h

@@ -224,18 +224,24 @@ ccl_device void camera_sample_orthographic(KernelGlobals kg,
 
 /* Panorama Camera */
 
+ccl_device_inline float3 camera_panorama_direction(ccl_constant KernelCamera *cam,
+                                                   float x,
+                                                   float y)
+{
+  const ProjectionTransform rastertocamera = cam->rastertocamera;
+  float3 Pcamera = transform_perspective(&rastertocamera, make_float3(x, y, 0.0f));
+  return panorama_to_direction(cam, Pcamera.x, Pcamera.y);
+}
+
 ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
                                               ccl_global const DecomposedTransform *cam_motion,
                                               const float2 raster,
                                               const float2 rand_lens,
                                               ccl_private Ray *ray)
 {
-  ProjectionTransform rastertocamera = cam->rastertocamera;
-  float3 Pcamera = transform_perspective(&rastertocamera, float2_to_float3(raster));
-
   /* create ray form raster position */
   float3 P = zero_float3();
-  float3 D = panorama_to_direction(cam, Pcamera.x, Pcamera.y);
+  float3 D = camera_panorama_direction(cam, raster.x, raster.y);
 
   /* indicates ray should not receive any light, outside of the lens */
   if (is_zero(D)) {
@@ -246,6 +252,11 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
   /* modify ray for depth of field */
   float aperturesize = cam->aperturesize;
 
+#ifdef __RAY_DIFFERENTIALS__
+  /* keep pre-DoF value for differentials later */
+  float3 Dcenter = D;
+#endif
+
   if (aperturesize > 0.0f) {
     /* sample point on aperture */
     float2 lens_uv = camera_sample_aperture(cam, rand_lens) * aperturesize;
@@ -289,38 +300,32 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
    * because we don't want to be affected by depth of field. We compute
    * ray origin and direction for the center and two neighboring pixels
    * and simply take their differences. */
-  float3 Pcenter = Pcamera;
-  float3 Dcenter = panorama_to_direction(cam, Pcenter.x, Pcenter.y);
+  float3 Dx = camera_panorama_direction(cam, raster.x + 1.0f, raster.y);
+  float3 Dy = camera_panorama_direction(cam, raster.x, raster.y + 1.0f);
+
   if (use_stereo) {
+    float3 Pcenter = zero_float3();
+    float3 Px = zero_float3();
+    float3 Py = zero_float3();
     spherical_stereo_transform(cam, &Pcenter, &Dcenter);
-  }
-  Pcenter = transform_point(&cameratoworld, Pcenter);
-  Dcenter = normalize(transform_direction(&cameratoworld, Dcenter));
-
-  float3 Px = transform_perspective(&rastertocamera, make_float3(raster.x + 1.0f, raster.y, 0.0f));
-  float3 Dx = panorama_to_direction(cam, Px.x, Px.y);
-  if (use_stereo) {
     spherical_stereo_transform(cam, &Px, &Dx);
-  }
-  Px = transform_point(&cameratoworld, Px);
-  Dx = normalize(transform_direction(&cameratoworld, Dx));
-
-  differential3 dP, dD;
-  dP.dx = Px - Pcenter;
-  dD.dx = Dx - Dcenter;
-
-  float3 Py = transform_perspective(&rastertocamera, make_float3(raster.x, raster.y + 1.0f, 0.0f));
-  float3 Dy = panorama_to_direction(cam, Py.x, Py.y);
-  if (use_stereo) {
     spherical_stereo_transform(cam, &Py, &Dy);
-  }
-  Py = transform_point(&cameratoworld, Py);
-  Dy = normalize(transform_direction(&cameratoworld, Dy));
 
-  dP.dy = Py - Pcenter;
-  dD.dy = Dy - Dcenter;
+    differential3 dP;
+    Pcenter = transform_point(&cameratoworld, Pcenter);
+    dP.dx = transform_point(&cameratoworld, Px) - Pcenter;
+    dP.dy = transform_point(&cameratoworld, Py) - Pcenter;
+    ray->dP = differential_make_compact(dP);
+  }
+  else {
+    ray->dP = differential_zero_compact();
+  }
+
+  differential3 dD;
+  Dcenter = normalize(transform_direction(&cameratoworld, Dcenter));
+  dD.dx = normalize(transform_direction(&cameratoworld, Dx)) - Dcenter;
+  dD.dy = normalize(transform_direction(&cameratoworld, Dy)) - Dcenter;
   ray->dD = differential_make_compact(dD);
-  ray->dP = differential_make_compact(dP);
 #endif
 
   /* clipping */

intern/cycles/kernel/closure/bsdf_principled_hair_huang.h

@@ -18,7 +18,7 @@ typedef struct HuangHairExtra {
   /* Optional modulation factors. */
   float R, TT, TRT;
 
-  /* Local coordinate system. X is stored as `bsdf->N`.*/
+  /* Local coordinate system. X is stored as `bsdf->N`. */
   float3 Y, Z;
 
   /* Incident direction in local coordinate system. */

intern/cycles/kernel/closure/bsdf_toon.h

@@ -18,15 +18,20 @@ typedef struct ToonBsdf {
 
 static_assert(sizeof(ShaderClosure) >= sizeof(ToonBsdf), "ToonBsdf is too large!");
 
+ccl_device_inline int bsdf_toon_setup_common(ccl_private ToonBsdf *bsdf)
+{
+  bsdf->size = clamp(bsdf->size, 1e-5f, 1.0f) * M_PI_2_F;
+  bsdf->smooth = saturatef(bsdf->smooth) * M_PI_2_F;
+
+  return SD_BSDF | SD_BSDF_HAS_EVAL;
+}
+
 /* DIFFUSE TOON */
 
 ccl_device int bsdf_diffuse_toon_setup(ccl_private ToonBsdf *bsdf)
 {
   bsdf->type = CLOSURE_BSDF_DIFFUSE_TOON_ID;
-  bsdf->size = saturatef(bsdf->size);
-  bsdf->smooth = saturatef(bsdf->smooth);
-
-  return SD_BSDF | SD_BSDF_HAS_EVAL;
+  return bsdf_toon_setup_common(bsdf);
 }
 
 ccl_device float bsdf_toon_get_intensity(float max_angle, float smooth, float angle)
@@ -57,19 +62,17 @@ ccl_device Spectrum bsdf_diffuse_toon_eval(ccl_private const ShaderClosure *sc,
                                            ccl_private float *pdf)
 {
   ccl_private const ToonBsdf *bsdf = (ccl_private const ToonBsdf *)sc;
+  float max_angle = bsdf->size;
+  float smooth = bsdf->smooth;
   float cosNO = dot(bsdf->N, wo);
 
   if (cosNO >= 0.0f) {
-    float max_angle = bsdf->size * M_PI_2_F;
-    float smooth = bsdf->smooth * M_PI_2_F;
     float angle = safe_acosf(fmaxf(cosNO, 0.0f));
+    float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
 
-    float eval = bsdf_toon_get_intensity(max_angle, smooth, angle);
-
-    if (eval > 0.0f) {
-      float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
-
-      *pdf = 0.5f * M_1_PI_F / (1.0f - cosf(sample_angle));
+    if (angle < sample_angle) {
+      float eval = bsdf_toon_get_intensity(max_angle, smooth, angle);
+      *pdf = M_1_2PI_F / one_minus_cos(sample_angle);
       return make_spectrum(*pdf * eval);
     }
   }
@@ -87,29 +90,22 @@ ccl_device int bsdf_diffuse_toon_sample(ccl_private const ShaderClosure *sc,
                                         ccl_private float *pdf)
 {
   ccl_private const ToonBsdf *bsdf = (ccl_private const ToonBsdf *)sc;
-  float max_angle = bsdf->size * M_PI_2_F;
-  float smooth = bsdf->smooth * M_PI_2_F;
+  float max_angle = bsdf->size;
+  float smooth = bsdf->smooth;
   float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
-  float angle = sample_angle * rand.x;
 
-  if (sample_angle > 0.0f) {
-    float unused;
-    *wo = sample_uniform_cone(bsdf->N, one_minus_cos(sample_angle), rand, &unused, pdf);
+  float cosNO;
+  *wo = sample_uniform_cone(bsdf->N, one_minus_cos(sample_angle), rand, &cosNO, pdf);
 
-    if (dot(Ng, *wo) > 0.0f) {
-      *eval = make_spectrum(*pdf * bsdf_toon_get_intensity(max_angle, smooth, angle));
-    }
-    else {
-      *eval = zero_spectrum();
-      *pdf = 0.0f;
-    }
-  }
-  else {
-    *eval = zero_spectrum();
-    *pdf = 0.0f;
+  if (dot(Ng, *wo) > 0.0f) {
+    float angle = acosf(cosNO);
+    *eval = make_spectrum(*pdf * bsdf_toon_get_intensity(max_angle, smooth, angle));
+    return LABEL_REFLECT | LABEL_DIFFUSE;
   }
 
-  return LABEL_REFLECT | LABEL_DIFFUSE;
+  *pdf = 0.0f;
+  *eval = zero_spectrum();
+  return LABEL_NONE;
 }
 
 /* GLOSSY TOON */
@@ -117,10 +113,7 @@ ccl_device int bsdf_diffuse_toon_sample(ccl_private const ShaderClosure *sc,
 ccl_device int bsdf_glossy_toon_setup(ccl_private ToonBsdf *bsdf)
 {
   bsdf->type = CLOSURE_BSDF_GLOSSY_TOON_ID;
-  bsdf->size = saturatef(bsdf->size);
-  bsdf->smooth = saturatef(bsdf->smooth);
-
-  return SD_BSDF | SD_BSDF_HAS_EVAL;
+  return bsdf_toon_setup_common(bsdf);
 }
 
 ccl_device Spectrum bsdf_glossy_toon_eval(ccl_private const ShaderClosure *sc,
@@ -129,8 +122,8 @@ ccl_device Spectrum bsdf_glossy_toon_eval(ccl_private const ShaderClosure *sc,
                                           ccl_private float *pdf)
 {
   ccl_private const ToonBsdf *bsdf = (ccl_private const ToonBsdf *)sc;
-  float max_angle = bsdf->size * M_PI_2_F;
-  float smooth = bsdf->smooth * M_PI_2_F;
+  float max_angle = bsdf->size;
+  float smooth = bsdf->smooth;
   float cosNI = dot(bsdf->N, wi);
   float cosNO = dot(bsdf->N, wo);
 
@@ -140,12 +133,13 @@ ccl_device Spectrum bsdf_glossy_toon_eval(ccl_private const ShaderClosure *sc,
     float cosRO = dot(R, wo);
 
     float angle = safe_acosf(fmaxf(cosRO, 0.0f));
-
-    float eval = bsdf_toon_get_intensity(max_angle, smooth, angle);
     float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
 
-    *pdf = 0.5f * M_1_PI_F / (1.0f - cosf(sample_angle));
-    return make_spectrum(*pdf * eval);
+    if (angle < sample_angle) {
+      float eval = bsdf_toon_get_intensity(max_angle, smooth, angle);
+      *pdf = M_1_2PI_F / one_minus_cos(sample_angle);
+      return make_spectrum(*pdf * eval);
+    }
   }
   *pdf = 0.0f;
   return zero_spectrum();
@@ -160,8 +154,8 @@ ccl_device int bsdf_glossy_toon_sample(ccl_private const ShaderClosure *sc,
                                        ccl_private float *pdf)
 {
   ccl_private const ToonBsdf *bsdf = (ccl_private const ToonBsdf *)sc;
-  float max_angle = bsdf->size * M_PI_2_F;
-  float smooth = bsdf->smooth * M_PI_2_F;
+  float max_angle = bsdf->size;
+  float smooth = bsdf->smooth;
   float cosNI = dot(bsdf->N, wi);
 
   if (cosNI > 0) {
@@ -169,30 +163,21 @@ ccl_device int bsdf_glossy_toon_sample(ccl_private const ShaderClosure *sc,
     float3 R = (2 * cosNI) * bsdf->N - wi;
 
     float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
-    float angle = sample_angle * rand.x;
 
-    float unused;
-    *wo = sample_uniform_cone(R, one_minus_cos(sample_angle), rand, &unused, pdf);
+    float cosRO;
+    *wo = sample_uniform_cone(R, one_minus_cos(sample_angle), rand, &cosRO, pdf);
 
-    if (dot(Ng, *wo) > 0.0f) {
-      float cosNO = dot(bsdf->N, *wo);
-
-      /* make sure the direction we chose is still in the right hemisphere */
-      if (cosNO > 0) {
-        *eval = make_spectrum(*pdf * bsdf_toon_get_intensity(max_angle, smooth, angle));
-      }
-      else {
-        *pdf = 0.0f;
-        *eval = zero_spectrum();
-      }
-    }
-    else {
-      *pdf = 0.0f;
-      *eval = zero_spectrum();
+    /* make sure the direction we chose is still in the right hemisphere */
+    if (dot(Ng, *wo) > 0.0f && dot(bsdf->N, *wo) > 0.0f) {
+      float angle = acosf(cosRO);
+      *eval = make_spectrum(*pdf * bsdf_toon_get_intensity(max_angle, smooth, angle));
+      return LABEL_GLOSSY | LABEL_REFLECT;
     }
   }
 
-  return LABEL_GLOSSY | LABEL_REFLECT;
+  *pdf = 0.0f;
+  *eval = zero_spectrum();
+  return LABEL_NONE;
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/device/cpu/compat.h

@@ -18,9 +18,8 @@
 #include "util/texture.h"
 #include "util/types.h"
 
-/* On x86_64, versions of glibc < 2.16 have an issue where expf is
- * much slower than the double version.  This was fixed in glibc 2.16.
- */
+/* On x86_64, versions of GLIBC < 2.16 have an issue where `expf` is
+ * much slower than the double version. This was fixed in GLIBC 2.16. */
 #if !defined(__KERNEL_GPU__) && defined(__x86_64__) && defined(__x86_64__) && \
     defined(__GNU_LIBRARY__) && defined(__GLIBC__) && defined(__GLIBC_MINOR__) && \
     (__GLIBC__ <= 2 && __GLIBC_MINOR__ < 16)

intern/cycles/kernel/device/metal/bvh.h

@@ -144,7 +144,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
     return false;
   }
 
-#if defined(__KERNEL_DEBUG__)
+#if defined(WITH_CYCLES_DEBUG)
   if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
     isect->t = ray->tmax;
     isect->type = PRIMITIVE_NONE;
@@ -281,7 +281,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
     return false;
   }
 
-#  if defined(__KERNEL_DEBUG__)
+#  if defined(WITH_CYCLES_DEBUG)
   if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
     if (local_isect) {
       local_isect->num_hits = 0;
@@ -383,7 +383,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
     return false;
   }
 
-#  if defined(__KERNEL_DEBUG__)
+#  if defined(WITH_CYCLES_DEBUG)
   if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
     kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
     return false;
@@ -446,7 +446,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
     return false;
   }
 
-#  if defined(__KERNEL_DEBUG__)
+#  if defined(WITH_CYCLES_DEBUG)
   if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
     kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
     return false;

intern/cycles/kernel/device/oneapi/compat.h

@@ -38,7 +38,7 @@
 #define ccl_global
 #define ccl_always_inline __attribute__((always_inline))
 #define ccl_device_inline inline
-#define ccl_noinline __attribute__((noinline))
+#define ccl_noinline
 #define ccl_inline_constant const constexpr
 #define ccl_device_constant static constexpr
 #define ccl_static_constexpr static constexpr

intern/cycles/kernel/integrator/shade_surface.h

@@ -280,6 +280,16 @@ ccl_device
 
   kernel_assert(ls.pdf != 0.0f);
 
+  const bool is_transmission = dot(ls.D, sd->N) < 0.0f;
+
+  if (ls.prim != PRIM_NONE && ls.prim == sd->prim && ls.object == sd->object) {
+    /* Skip self intersection if light direction lies in the same hemisphere as the geometric
+     * normal. */
+    if (dot(ls.D, is_transmission ? -sd->Ng : sd->Ng) > 0.0f) {
+      return;
+    }
+  }
+
   /* Evaluate light shader.
    *
    * TODO: can we reuse sd memory? In theory we can move this after
@@ -292,8 +302,6 @@ ccl_device
   Ray ray ccl_optional_struct_init;
   BsdfEval bsdf_eval ccl_optional_struct_init;
 
-  const bool is_transmission = dot(ls.D, sd->N) < 0.0f;
-
   int mnee_vertex_count = 0;
 #ifdef __MNEE__
   IF_KERNEL_FEATURE(MNEE)

intern/cycles/kernel/integrator/shade_volume.h

@@ -64,6 +64,11 @@ typedef struct VolumeShaderCoefficients {
   Spectrum emission;
 } VolumeShaderCoefficients;
 
+typedef struct EquiangularCoefficients {
+  float3 P;
+  float2 t_range;
+} EquiangularCoefficients;
+
 /* Evaluate shader to get extinction coefficient at P. */
 ccl_device_inline bool shadow_volume_shader_sample(KernelGlobals kg,
                                                    IntegratorShadowState state,
@@ -264,18 +269,18 @@ ccl_device void volume_shadow_heterogeneous(KernelGlobals kg,
 #  define VOLUME_SAMPLE_PDF_CUTOFF 1e-8f
 
 ccl_device float volume_equiangular_sample(ccl_private const Ray *ccl_restrict ray,
-                                           const float3 light_P,
+                                           ccl_private const EquiangularCoefficients &coeffs,
                                            const float xi,
                                            ccl_private float *pdf)
 {
-  const float tmin = ray->tmin;
-  const float tmax = ray->tmax;
-  const float delta = dot((light_P - ray->P), ray->D);
-  const float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta);
+  const float delta = dot((coeffs.P - ray->P), ray->D);
+  const float D = safe_sqrtf(len_squared(coeffs.P - ray->P) - delta * delta);
   if (UNLIKELY(D == 0.0f)) {
     *pdf = 0.0f;
     return 0.0f;
   }
+  const float tmin = coeffs.t_range.x;
+  const float tmax = coeffs.t_range.y;
   const float theta_a = atan2f(tmin - delta, D);
   const float theta_b = atan2f(tmax - delta, D);
   const float t_ = D * tanf((xi * theta_b) + (1 - xi) * theta_a);
@@ -289,17 +294,17 @@ ccl_device float volume_equiangular_sample(ccl_private const Ray *ccl_restrict r
 }
 
 ccl_device float volume_equiangular_pdf(ccl_private const Ray *ccl_restrict ray,
-                                        const float3 light_P,
+                                        ccl_private const EquiangularCoefficients &coeffs,
                                         const float sample_t)
 {
-  const float delta = dot((light_P - ray->P), ray->D);
-  const float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta);
+  const float delta = dot((coeffs.P - ray->P), ray->D);
+  const float D = safe_sqrtf(len_squared(coeffs.P - ray->P) - delta * delta);
   if (UNLIKELY(D == 0.0f)) {
     return 0.0f;
   }
 
-  const float tmin = ray->tmin;
-  const float tmax = ray->tmax;
+  const float tmin = coeffs.t_range.x;
+  const float tmax = coeffs.t_range.y;
   const float t_ = sample_t - delta;
 
   const float theta_a = atan2f(tmin - delta, D);
@@ -313,30 +318,27 @@ ccl_device float volume_equiangular_pdf(ccl_private const Ray *ccl_restrict ray,
   return pdf;
 }
 
-ccl_device float volume_equiangular_cdf(ccl_private const Ray *ccl_restrict ray,
-                                        const float3 light_P,
-                                        const float sample_t)
+ccl_device_inline bool volume_equiangular_valid_ray_segment(KernelGlobals kg,
+                                                            const float3 ray_P,
+                                                            const float3 ray_D,
+                                                            ccl_private float2 *t_range,
+                                                            const ccl_private LightSample *ls)
 {
-  float delta = dot((light_P - ray->P), ray->D);
-  float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta);
-  if (UNLIKELY(D == 0.0f)) {
-    return 0.0f;
+  if (ls->type == LIGHT_SPOT) {
+    ccl_global const KernelLight *klight = &kernel_data_fetch(lights, ls->lamp);
+    return spot_light_valid_ray_segment(klight, ray_P, ray_D, t_range);
+  }
+  if (ls->type == LIGHT_AREA) {
+    ccl_global const KernelLight *klight = &kernel_data_fetch(lights, ls->lamp);
+    return area_light_valid_ray_segment(&klight->area, ray_P - klight->co, ray_D, t_range);
+  }
+  if (ls->type == LIGHT_TRIANGLE) {
+    return triangle_light_valid_ray_segment(kg, ray_P - ls->P, ray_D, t_range, ls);
   }
 
-  const float tmin = ray->tmin;
-  const float tmax = ray->tmax;
-  const float t_ = sample_t - delta;
-
-  const float theta_a = atan2f(tmin - delta, D);
-  const float theta_b = atan2f(tmax - delta, D);
-  if (UNLIKELY(theta_b == theta_a)) {
-    return 0.0f;
-  }
-
-  const float theta_sample = atan2f(t_, D);
-  const float cdf = (theta_sample - theta_a) / (theta_b - theta_a);
-
-  return cdf;
+  /* Point light, the whole range of the ray is visible. */
+  kernel_assert(ls->type == LIGHT_POINT);
+  return true;
 }
 
 /* Distance sampling */
@@ -429,7 +431,7 @@ typedef struct VolumeIntegrateState {
 ccl_device_forceinline void volume_integrate_step_scattering(
     ccl_private const ShaderData *sd,
     ccl_private const Ray *ray,
-    const float3 equiangular_light_P,
+    ccl_private const EquiangularCoefficients &equiangular_coeffs,
     ccl_private const VolumeShaderCoefficients &ccl_restrict coeff,
     const Spectrum transmittance,
     ccl_private VolumeIntegrateState &ccl_restrict vstate,
@@ -500,7 +502,7 @@ ccl_device_forceinline void volume_integrate_step_scattering(
 
           /* Multiple importance sampling. */
           if (vstate.use_mis) {
-            const float equiangular_pdf = volume_equiangular_pdf(ray, equiangular_light_P, new_t);
+            const float equiangular_pdf = volume_equiangular_pdf(ray, equiangular_coeffs, new_t);
             const float mis_weight = power_heuristic(vstate.distance_pdf * distance_pdf,
                                                      equiangular_pdf);
             result.direct_throughput *= 2.0f * mis_weight;
@@ -535,7 +537,7 @@ ccl_device_forceinline void volume_integrate_heterogeneous(
     ccl_global float *ccl_restrict render_buffer,
     const float object_step_size,
     const VolumeSampleMethod direct_sample_method,
-    const float3 equiangular_light_P,
+    ccl_private const EquiangularCoefficients &equiangular_coeffs,
     ccl_private VolumeIntegrateResult &result)
 {
   PROFILING_INIT(kg, PROFILING_SHADE_VOLUME_INTEGRATE);
@@ -586,7 +588,7 @@ ccl_device_forceinline void volume_integrate_heterogeneous(
   /* Equiangular sampling: compute distance and PDF in advance. */
   if (vstate.direct_sample_method == VOLUME_SAMPLE_EQUIANGULAR) {
     result.direct_t = volume_equiangular_sample(
-        ray, equiangular_light_P, vstate.rscatter, &vstate.equiangular_pdf);
+        ray, equiangular_coeffs, vstate.rscatter, &vstate.equiangular_pdf);
   }
 #  ifdef __PATH_GUIDING__
   result.direct_sample_method = vstate.direct_sample_method;
@@ -640,7 +642,7 @@ ccl_device_forceinline void volume_integrate_heterogeneous(
 
           /* Scattering and absorption. */
           volume_integrate_step_scattering(
-              sd, ray, equiangular_light_P, coeff, transmittance, vstate, result);
+              sd, ray, equiangular_coeffs, coeff, transmittance, vstate, result);
         }
         else {
           /* Absorption only. */
@@ -699,7 +701,8 @@ ccl_device_forceinline bool integrate_volume_equiangular_sample_light(
     ccl_private const Ray *ccl_restrict ray,
     ccl_private const ShaderData *ccl_restrict sd,
     ccl_private const RNGState *ccl_restrict rng_state,
-    ccl_private float3 *ccl_restrict P)
+    ccl_private EquiangularCoefficients *ccl_restrict equiangular_coeffs,
+    ccl_private LightSample &ccl_restrict ls)
 {
   /* Test if there is a light or BSDF that needs direct light. */
   if (!kernel_data.integrator.use_direct_light) {
@@ -711,7 +714,6 @@ ccl_device_forceinline bool integrate_volume_equiangular_sample_light(
   const uint bounce = INTEGRATOR_STATE(state, path, bounce);
   const float3 rand_light = path_state_rng_3D(kg, rng_state, PRNG_LIGHT);
 
-  LightSample ls ccl_optional_struct_init;
   if (!light_sample_from_volume_segment(kg,
                                         rand_light,
                                         sd->time,
@@ -734,9 +736,10 @@ ccl_device_forceinline bool integrate_volume_equiangular_sample_light(
     return false;
   }
 
-  *P = ls.P;
+  equiangular_coeffs->P = ls.P;
 
-  return true;
+  return volume_equiangular_valid_ray_segment(
+      kg, ray->P, ray->D, &equiangular_coeffs->t_range, &ls);
 }
 
 /* Path tracing: sample point on light and evaluate light shader, then
@@ -751,41 +754,26 @@ ccl_device_forceinline void integrate_volume_direct_light(
 #  ifdef __PATH_GUIDING__
     ccl_private const Spectrum unlit_throughput,
 #  endif
-    ccl_private const Spectrum throughput)
+    ccl_private const Spectrum throughput,
+    ccl_private LightSample &ccl_restrict ls)
 {
   PROFILING_INIT(kg, PROFILING_SHADE_VOLUME_DIRECT_LIGHT);
 
-  if (!kernel_data.integrator.use_direct_light) {
+  if (!kernel_data.integrator.use_direct_light || ls.emitter_id == EMITTER_NONE) {
     return;
   }
 
-  /* Sample position on the same light again, now from the shading point where we scattered.
-   *
-   * Note that this means we sample the light tree twice when equiangular sampling is used.
-   * We could consider sampling the light tree just once and use the same light position again.
-   *
-   * This would make the PDFs for MIS weights more complicated due to having to account for
-   * both distance/equiangular and direct/indirect light sampling, but could be more accurate.
-   * Additionally we could end up behind the light or outside a spot light cone, which might
-   * waste a sample. Though on the other hand it would be possible to prevent that with
-   * equiangular sampling restricted to a smaller sub-segment where the light has influence. */
-  LightSample ls ccl_optional_struct_init;
+  /* Sample position on the same light again, now from the shading point where we scattered. */
   {
     const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
     const uint bounce = INTEGRATOR_STATE(state, path, bounce);
     const float3 rand_light = path_state_rng_3D(kg, rng_state, PRNG_LIGHT);
+    const float3 N = zero_float3();
+    const int object_receiver = light_link_receiver_nee(kg, sd);
+    const int shader_flags = SD_BSDF_HAS_TRANSMISSION;
 
-    if (!light_sample_from_position(kg,
-                                    rng_state,
-                                    rand_light,
-                                    sd->time,
-                                    P,
-                                    zero_float3(),
-                                    light_link_receiver_nee(kg, sd),
-                                    SD_BSDF_HAS_TRANSMISSION,
-                                    bounce,
-                                    path_flag,
-                                    &ls))
+    if (!light_sample<false>(
+            kg, rand_light, sd->time, P, N, object_receiver, shader_flags, bounce, path_flag, &ls))
     {
       return;
     }
@@ -903,6 +891,7 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
     KernelGlobals kg,
     IntegratorState state,
     ccl_private ShaderData *sd,
+    ccl_private const Ray *ray,
     ccl_private const RNGState *rng_state,
     ccl_private const ShaderVolumePhases *phases)
 {
@@ -955,6 +944,9 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
   INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
   INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(phase_wo);
   INTEGRATOR_STATE_WRITE(state, ray, tmin) = 0.0f;
+#  ifdef __LIGHT_TREE__
+  INTEGRATOR_STATE_WRITE(state, ray, previous_dt) = ray->tmax - ray->tmin;
+#  endif
   INTEGRATOR_STATE_WRITE(state, ray, tmax) = FLT_MAX;
 #  ifdef __RAY_DIFFERENTIALS__
   INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
@@ -983,7 +975,8 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
 
   /* Update path state */
   INTEGRATOR_STATE_WRITE(state, path, mis_ray_pdf) = phase_pdf;
-  INTEGRATOR_STATE_WRITE(state, path, mis_origin_n) = zero_float3();
+  const float3 previous_P = ray->P + ray->D * ray->tmin;
+  INTEGRATOR_STATE_WRITE(state, path, mis_origin_n) = sd->P - previous_P;
   INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = fminf(
       unguided_phase_pdf, INTEGRATOR_STATE(state, path, min_ray_pdf));
 
@@ -1015,11 +1008,15 @@ ccl_device VolumeIntegrateEvent volume_integrate(KernelGlobals kg,
 
   /* Sample light ahead of volume stepping, for equiangular sampling. */
   /* TODO: distant lights are ignored now, but could instead use even distribution. */
+  LightSample ls ccl_optional_struct_init;
+  ls.emitter_id = EMITTER_NONE;
   const bool need_light_sample = !(INTEGRATOR_STATE(state, path, flag) & PATH_RAY_TERMINATE);
-  float3 equiangular_P = zero_float3();
-  const bool have_equiangular_sample = need_light_sample &&
-                                       integrate_volume_equiangular_sample_light(
-                                           kg, state, ray, &sd, &rng_state, &equiangular_P);
+
+  EquiangularCoefficients equiangular_coeffs = {zero_float3(), make_float2(ray->tmin, ray->tmax)};
+
+  const bool have_equiangular_sample =
+      need_light_sample && integrate_volume_equiangular_sample_light(
+                               kg, state, ray, &sd, &rng_state, &equiangular_coeffs, ls);
 
   VolumeSampleMethod direct_sample_method = (have_equiangular_sample) ?
                                                 volume_stack_sample_method(kg, state) :
@@ -1049,7 +1046,7 @@ ccl_device VolumeIntegrateEvent volume_integrate(KernelGlobals kg,
                                  render_buffer,
                                  step_size,
                                  direct_sample_method,
-                                 equiangular_P,
+                                 equiangular_coeffs,
                                  result);
 
   /* Perform path termination. The intersect_closest will have already marked this path
@@ -1117,7 +1114,8 @@ ccl_device VolumeIntegrateEvent volume_integrate(KernelGlobals kg,
 #  ifdef __PATH_GUIDING__
                                   unlit_throughput,
 #  endif
-                                  result.direct_throughput);
+                                  result.direct_throughput,
+                                  ls);
   }
 
   /* Indirect light.
@@ -1156,7 +1154,7 @@ ccl_device VolumeIntegrateEvent volume_integrate(KernelGlobals kg,
 #    endif
 #  endif
 
-    if (integrate_volume_phase_scatter(kg, state, &sd, &rng_state, &result.indirect_phases)) {
+    if (integrate_volume_phase_scatter(kg, state, &sd, ray, &rng_state, &result.indirect_phases)) {
       return VOLUME_PATH_SCATTERED;
     }
     else {

intern/cycles/kernel/integrator/state_template.h

@@ -75,6 +75,9 @@ KERNEL_STRUCT_MEMBER(ray, float, tmax, KERNEL_FEATURE_PATH_TRACING)
 KERNEL_STRUCT_MEMBER(ray, float, time, KERNEL_FEATURE_PATH_TRACING)
 KERNEL_STRUCT_MEMBER(ray, float, dP, KERNEL_FEATURE_PATH_TRACING)
 KERNEL_STRUCT_MEMBER(ray, float, dD, KERNEL_FEATURE_PATH_TRACING)
+#ifdef __LIGHT_TREE__
+KERNEL_STRUCT_MEMBER(ray, float, previous_dt, KERNEL_FEATURE_PATH_TRACING)
+#endif
 KERNEL_STRUCT_END(ray)
 
 /*************************** Intersection result ******************************/

intern/cycles/kernel/integrator/surface_shader.h

@@ -212,7 +212,7 @@ ccl_device_inline void surface_shader_prepare_closures(KernelGlobals kg,
 /* BSDF */
 #ifdef WITH_CYCLES_DEBUG
 ccl_device_inline void surface_shader_validate_bsdf_sample(const KernelGlobals kg,
-                                                           const ShaderClosure *sc,
+                                                           ccl_private const ShaderClosure *sc,
                                                            const float3 wo,
                                                            const int org_label,
                                                            const float2 org_roughness,

intern/cycles/kernel/light/area.h

@@ -233,6 +233,11 @@ ccl_device bool area_light_spread_clamp_light(const float3 P,
   return true;
 }
 
+ccl_device_forceinline bool area_light_is_ellipse(const ccl_global KernelAreaLight *light)
+{
+  return light->invarea < 0.0f;
+}
+
 /* Common API. */
 /* Compute `eval_fac` and `pdf`. Also sample a new position on the light if `sample_coord`. */
 template<bool in_volume_segment>
@@ -311,7 +316,7 @@ ccl_device_inline bool area_light_eval(const ccl_global KernelLight *klight,
     ls->pdf *= light_pdf_area_to_solid_angle(Ng, -ls->D, ls->t);
   }
 
-  return ls->eval_fac > 0;
+  return in_volume_segment || ls->eval_fac > 0;
 }
 
 template<bool in_volume_segment>
@@ -338,7 +343,7 @@ ccl_device_inline bool area_light_sample(const ccl_global KernelLight *klight,
   const float light_v = dot(inplane, klight->area.axis_v) / klight->area.len_v;
 
   if (!in_volume_segment) {
-    const bool is_ellipse = (klight->area.invarea < 0.0f);
+    const bool is_ellipse = area_light_is_ellipse(&klight->area);
 
     /* Sampled point lies outside of the area light. */
     if (is_ellipse && (sqr(light_u) + sqr(light_v) > 0.25f)) {
@@ -380,7 +385,7 @@ ccl_device_inline bool area_light_intersect(const ccl_global KernelLight *klight
 {
   /* Area light. */
   const float invarea = fabsf(klight->area.invarea);
-  const bool is_ellipse = (klight->area.invarea < 0.0f);
+  const bool is_ellipse = area_light_is_ellipse(&klight->area);
   if (invarea == 0.0f) {
     return false;
   }
@@ -428,6 +433,55 @@ ccl_device_inline bool area_light_sample_from_intersection(
   return area_light_eval<false>(klight, ray_P, &light_P, ls, zero_float2(), false);
 }
 
+/* Returns the maximal distance between the light center and the boundary. */
+ccl_device_forceinline float area_light_max_extent(const ccl_global KernelAreaLight *light)
+{
+  return 0.5f * (area_light_is_ellipse(light) ? fmaxf(light->len_u, light->len_v) :
+                                                len(make_float2(light->len_u, light->len_v)));
+}
+
+/* Find the ray segment lit by the area light. */
+ccl_device_inline bool area_light_valid_ray_segment(const ccl_global KernelAreaLight *light,
+                                                    float3 P,
+                                                    float3 D,
+                                                    ccl_private float2 *t_range)
+{
+  bool valid;
+  const float tan_half_spread = light->tan_half_spread;
+  float3 axis = light->dir;
+
+  const bool angle_almost_zero = (tan_half_spread < 1e-5f);
+  if (angle_almost_zero) {
+    /* Map to local coordinate of the light. Do not use `itfm` in `KernelLight` as there might be
+     * additional scaling in the light size. */
+    const Transform tfm = make_transform(light->axis_u, light->axis_v, axis);
+    P = transform_point(&tfm, P);
+    D = transform_direction(&tfm, D);
+    axis = make_float3(0.0f, 0.0f, 1.0f);
+
+    const float half_len_u = 0.5f * light->len_u;
+    const float half_len_v = 0.5f * light->len_v;
+    if (area_light_is_ellipse(light)) {
+      valid = ray_infinite_cylinder_intersect(P, D, half_len_u, half_len_v, t_range);
+    }
+    else {
+      const float3 bbox_min = make_float3(-half_len_u, -half_len_v, 0.0f);
+      const float3 bbox_max = make_float3(half_len_u, half_len_v, FLT_MAX);
+      valid = ray_aabb_intersect(bbox_min, bbox_max, P, D, t_range);
+    }
+  }
+  else {
+    /* Conservative estimation with the smallest possible cone covering the whole spread. */
+    const float3 apex_to_point = P + area_light_max_extent(light) / tan_half_spread * axis;
+    const float cos_angle_sq = 1.0f / (1.0f + sqr(tan_half_spread));
+
+    valid = ray_cone_intersect(axis, apex_to_point, D, cos_angle_sq, t_range);
+  }
+
+  /* Limit the range to the positive side of the area light. */
+  return valid && ray_plane_intersect(axis, P, D, t_range);
+}
+
 template<bool in_volume_segment>
 ccl_device_forceinline bool area_light_tree_parameters(const ccl_global KernelLight *klight,
                                                        const float3 centroid,
@@ -438,13 +492,11 @@ ccl_device_forceinline bool area_light_tree_parameters(const ccl_global KernelLi
                                                        ccl_private float2 &distance,
                                                        ccl_private float3 &point_to_centroid)
 {
-  if (!in_volume_segment) {
-    /* TODO: a cheap substitute for minimal distance between point and primitive. Does it
-     * worth the overhead to compute the accurate minimal distance? */
-    float min_distance;
-    point_to_centroid = safe_normalize_len(centroid - P, &min_distance);
-    distance = make_float2(min_distance, min_distance);
-  }
+  /* TODO: a cheap substitute for minimal distance between point and primitive. Does it worth the
+   * overhead to compute the accurate minimal distance? */
+  float min_distance;
+  point_to_centroid = safe_normalize_len(centroid - P, &min_distance);
+  distance = make_float2(min_distance, min_distance);
 
   cos_theta_u = FLT_MAX;
 
@@ -464,9 +516,8 @@ ccl_device_forceinline bool area_light_tree_parameters(const ccl_global KernelLi
   const bool shape_above_surface = dot(N, centroid - P) + fabsf(dot(N, extentu)) +
                                        fabsf(dot(N, extentv)) >
                                    0;
-  const bool in_volume = is_zero(N);
 
-  return (front_facing && shape_above_surface) || in_volume;
+  return front_facing && shape_above_surface;
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/light/common.h

@@ -12,9 +12,9 @@ CCL_NAMESPACE_BEGIN
 
 typedef struct LightSample {
   float3 P;            /* position on light, or direction for distant light */
-  float3 Ng;           /* normal on light */
-  float3 D;            /* direction from shading point to light */
+  packed_float3 Ng;    /* normal on light */
   float t;             /* distance to light (FLT_MAX for distant light) */
+  float3 D;            /* direction from shading point to light */
   float u, v;          /* parametric coordinate on primitive */
   float pdf;           /* pdf for selecting light and point on light */
   float pdf_selection; /* pdf for selecting light */
@@ -25,6 +25,7 @@ typedef struct LightSample {
   int lamp;            /* lamp id */
   int group;           /* lightgroup */
   LightType type;      /* type of light */
+  int emitter_id;      /* index in the emitter array */
 } LightSample;
 
 /* Utilities */

intern/cycles/kernel/light/distribution.h

@@ -41,36 +41,14 @@ ccl_device int light_distribution_sample(KernelGlobals kg, const float rand)
   return index;
 }
 
-template<bool in_volume_segment>
 ccl_device_noinline bool light_distribution_sample(KernelGlobals kg,
-                                                   const float3 rand,
-                                                   const float time,
-                                                   const float3 P,
-                                                   const float3 N,
-                                                   const int object_receiver,
-                                                   const int shader_flags,
-                                                   const int bounce,
-                                                   const uint32_t path_flag,
+                                                   const float rand,
                                                    ccl_private LightSample *ls)
 {
   /* Sample light index from distribution. */
-  /* The first two dimensions of the Sobol sequence have better stratification. */
-  const int index = light_distribution_sample(kg, rand.z);
-  const float pdf_selection = kernel_data.integrator.distribution_pdf_lights;
-  const float2 rand_uv = float3_to_float2(rand);
-  return light_sample<in_volume_segment>(kg,
-                                         rand_uv,
-                                         time,
-                                         P,
-                                         N,
-                                         object_receiver,
-                                         shader_flags,
-                                         bounce,
-                                         path_flag,
-                                         index,
-                                         0,
-                                         pdf_selection,
-                                         ls);
+  ls->emitter_id = light_distribution_sample(kg, rand);
+  ls->pdf_selection = kernel_data.integrator.distribution_pdf_lights;
+  return true;
 }
 
 ccl_device_inline float light_distribution_pdf_lamp(KernelGlobals kg)

intern/cycles/kernel/light/light.h

@@ -177,7 +177,7 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
 
 template<bool in_volume_segment>
 ccl_device_noinline bool light_sample(KernelGlobals kg,
-                                      const float2 rand,
+                                      const float3 rand_light,
                                       const float time,
                                       const float3 P,
                                       const float3 N,
@@ -185,33 +185,31 @@ ccl_device_noinline bool light_sample(KernelGlobals kg,
                                       const int shader_flags,
                                       const int bounce,
                                       const uint32_t path_flag,
-                                      const int emitter_index,
-                                      const int object_id,
-                                      const float pdf_selection,
                                       ccl_private LightSample *ls)
 {
+  /* The first two dimensions of the Sobol sequence have better stratification, use them to sample
+   * position on the light. */
+  const float2 rand = float3_to_float2(rand_light);
+
   int prim;
   MeshLight mesh_light;
 #ifdef __LIGHT_TREE__
   if (kernel_data.integrator.use_light_tree) {
     ccl_global const KernelLightTreeEmitter *kemitter = &kernel_data_fetch(light_tree_emitters,
-                                                                           emitter_index);
+                                                                           ls->emitter_id);
     prim = kemitter->light.id;
     mesh_light.shader_flag = kemitter->mesh_light.shader_flag;
-    mesh_light.object_id = object_id;
+    mesh_light.object_id = ls->object;
   }
   else
 #endif
   {
     ccl_global const KernelLightDistribution *kdistribution = &kernel_data_fetch(
-        light_distribution, emitter_index);
+        light_distribution, ls->emitter_id);
     prim = kdistribution->prim;
     mesh_light = kdistribution->mesh_light;
   }
 
-  /* A different value would be assigned in `triangle_light_sample()` if `!use_light_tree`. */
-  ls->pdf_selection = pdf_selection;
-
   if (prim >= 0) {
     /* Mesh light. */
     const int object = mesh_light.object_id;

intern/cycles/kernel/light/point.h

@@ -196,21 +196,22 @@ ccl_device_forceinline bool point_light_tree_parameters(const ccl_global KernelL
                                                         ccl_private float2 &distance,
                                                         ccl_private float3 &point_to_centroid)
 {
+  float min_distance;
+  point_to_centroid = safe_normalize_len(centroid - P, &min_distance);
+  distance = min_distance * one_float2();
+
   if (in_volume_segment) {
     cos_theta_u = 1.0f; /* Any value in [-1, 1], irrelevant since theta = 0 */
     return true;
   }
 
-  float dist_point_to_centroid;
-  point_to_centroid = safe_normalize_len(centroid - P, &dist_point_to_centroid);
-
   const float radius = klight->spot.radius;
 
   if (klight->spot.is_sphere) {
-    if (dist_point_to_centroid > radius) {
+    if (min_distance > radius) {
       /* Equivalent to a disk light with the same angular span. */
-      cos_theta_u = cos_from_sin(radius / dist_point_to_centroid);
-      distance = dist_point_to_centroid * make_float2(1.0f / cos_theta_u, 1.0f);
+      cos_theta_u = cos_from_sin(radius / min_distance);
+      distance.x = min_distance / cos_theta_u;
     }
     else {
       /* Similar to background light. */
@@ -220,10 +221,10 @@ ccl_device_forceinline bool point_light_tree_parameters(const ccl_global KernelL
     }
   }
   else {
-    const float hypotenus = sqrtf(sqr(radius) + sqr(dist_point_to_centroid));
-    cos_theta_u = dist_point_to_centroid / hypotenus;
+    const float hypotenus = sqrtf(sqr(radius) + sqr(min_distance));
+    cos_theta_u = min_distance / hypotenus;
 
-    distance = make_float2(hypotenus, dist_point_to_centroid);
+    distance.x = hypotenus;
   }
 
   return true;

intern/cycles/kernel/light/sample.h

@@ -329,17 +329,25 @@ ccl_device_inline bool light_sample_from_volume_segment(KernelGlobals kg,
                                                         const uint32_t path_flag,
                                                         ccl_private LightSample *ls)
 {
+  const int shader_flags = SD_BSDF_HAS_TRANSMISSION;
+
 #ifdef __LIGHT_TREE__
   if (kernel_data.integrator.use_light_tree) {
-    return light_tree_sample<true>(
-        kg, rand, time, P, D, t, object_receiver, SD_BSDF_HAS_TRANSMISSION, bounce, path_flag, ls);
+    if (!light_tree_sample<true>(kg, rand.z, P, D, t, object_receiver, shader_flags, ls)) {
+      return false;
+    }
   }
   else
 #endif
   {
-    return light_distribution_sample<true>(
-        kg, rand, time, P, D, object_receiver, SD_BSDF_HAS_TRANSMISSION, bounce, path_flag, ls);
+    if (!light_distribution_sample(kg, rand.z, ls)) {
+      return false;
+    }
   }
+
+  /* Sample position on the selected light. */
+  return light_sample<true>(
+      kg, rand, time, P, D, object_receiver, shader_flags, bounce, path_flag, ls);
 }
 
 ccl_device bool light_sample_from_position(KernelGlobals kg,
@@ -354,17 +362,24 @@ ccl_device bool light_sample_from_position(KernelGlobals kg,
                                            const uint32_t path_flag,
                                            ccl_private LightSample *ls)
 {
+  /* Randomly select a light. */
 #ifdef __LIGHT_TREE__
   if (kernel_data.integrator.use_light_tree) {
-    return light_tree_sample<false>(
-        kg, rand, time, P, N, 0.0f, object_receiver, shader_flags, bounce, path_flag, ls);
+    if (!light_tree_sample<false>(kg, rand.z, P, N, 0.0f, object_receiver, shader_flags, ls)) {
+      return false;
+    }
   }
   else
 #endif
   {
-    return light_distribution_sample<false>(
-        kg, rand, time, P, N, object_receiver, shader_flags, bounce, path_flag, ls);
+    if (!light_distribution_sample(kg, rand.z, ls)) {
+      return false;
+    }
   }
+
+  /* Sample position on the selected light. */
+  return light_sample<false>(
+      kg, rand, time, P, N, object_receiver, shader_flags, bounce, path_flag, ls);
 }
 
 /* Update light sample with new shading point position for MNEE. The position on the light is fixed
@@ -415,13 +430,15 @@ ccl_device_inline float light_sample_mis_weight_forward_surface(KernelGlobals kg
 #ifdef __LIGHT_TREE__
   if (kernel_data.integrator.use_light_tree) {
     float3 ray_P = INTEGRATOR_STATE(state, ray, P);
+    const float dt = INTEGRATOR_STATE(state, ray, previous_dt);
     const float3 N = INTEGRATOR_STATE(state, path, mis_origin_n);
+
     uint lookup_offset = kernel_data_fetch(object_lookup_offset, sd->object);
     uint prim_offset = kernel_data_fetch(object_prim_offset, sd->object);
     uint triangle = kernel_data_fetch(triangle_to_tree, sd->prim - prim_offset + lookup_offset);
 
     pdf *= light_tree_pdf(
-        kg, ray_P, N, path_flag, sd->object, triangle, light_link_receiver_forward(kg, state));
+        kg, ray_P, N, dt, path_flag, sd->object, triangle, light_link_receiver_forward(kg, state));
   }
   else
 #endif
@@ -445,9 +462,11 @@ ccl_device_inline float light_sample_mis_weight_forward_lamp(KernelGlobals kg,
 #ifdef __LIGHT_TREE__
   if (kernel_data.integrator.use_light_tree) {
     const float3 N = INTEGRATOR_STATE(state, path, mis_origin_n);
+    const float dt = INTEGRATOR_STATE(state, ray, previous_dt);
     pdf *= light_tree_pdf(kg,
                           P,
                           N,
+                          dt,
                           path_flag,
                           0,
                           kernel_data_fetch(light_to_tree, ls->lamp),
@@ -485,9 +504,10 @@ ccl_device_inline float light_sample_mis_weight_forward_background(KernelGlobals
 #ifdef __LIGHT_TREE__
   if (kernel_data.integrator.use_light_tree) {
     const float3 N = INTEGRATOR_STATE(state, path, mis_origin_n);
+    const float dt = INTEGRATOR_STATE(state, ray, previous_dt);
     uint light = kernel_data_fetch(light_to_tree, kernel_data.background.light_index);
     pdf *= light_tree_pdf(
-        kg, ray_P, N, path_flag, 0, light, light_link_receiver_forward(kg, state));
+        kg, ray_P, N, dt, path_flag, 0, light, light_link_receiver_forward(kg, state));
   }
   else
 #endif

intern/cycles/kernel/light/spot.h

@@ -9,11 +9,13 @@
 CCL_NAMESPACE_BEGIN
 
 /* Transform vector to spot light's local coordinate system. */
-ccl_device float3 spot_light_to_local(const ccl_global KernelSpotLight *spot, const float3 ray)
+ccl_device float3 spot_light_to_local(const ccl_global KernelLight *klight, const float3 ray)
 {
-  return safe_normalize(make_float3(dot(ray, spot->scaled_axis_u),
-                                    dot(ray, spot->scaled_axis_v),
-                                    dot(ray, spot->dir * spot->inv_len_z)));
+  const Transform itfm = klight->itfm;
+  float3 transformed_ray = safe_normalize(transform_direction(&itfm, ray));
+  transformed_ray.z = -transformed_ray.z;
+
+  return transformed_ray;
 }
 
 /* Compute spot light attenuation of a ray given in local coordinate system. */
@@ -58,7 +60,7 @@ ccl_device_inline bool spot_light_sample(const ccl_global KernelLight *klight,
     ls->t = FLT_MAX;
     if (d_sq > r_sq) {
       /* Outside sphere. */
-      const float one_minus_cos_half_spot_spread = 1.0f - klight->spot.cos_half_spot_angle;
+      const float one_minus_cos_half_spot_spread = 1.0f - klight->spot.cos_half_larger_spread;
       const float one_minus_cos_half_angle = sin_sqr_to_one_minus_cos(r_sq / d_sq);
 
       if (in_volume_segment || one_minus_cos_half_angle < one_minus_cos_half_spot_spread) {
@@ -92,7 +94,7 @@ ccl_device_inline bool spot_light_sample(const ccl_global KernelLight *klight,
     }
 
     /* Attenuation. */
-    const float3 local_ray = spot_light_to_local(&klight->spot, -ls->D);
+    const float3 local_ray = spot_light_to_local(klight, -ls->D);
     if (d_sq > r_sq) {
       ls->eval_fac *= spot_light_attenuation(&klight->spot, local_ray);
     }
@@ -128,7 +130,7 @@ ccl_device_inline bool spot_light_sample(const ccl_global KernelLight *klight,
     ls->Ng = -ls->D;
 
     /* Attenuation. */
-    const float3 local_ray = spot_light_to_local(&klight->spot, -ls->D);
+    const float3 local_ray = spot_light_to_local(klight, -ls->D);
     ls->eval_fac *= spot_light_attenuation(&klight->spot, local_ray);
     if (!in_volume_segment && ls->eval_fac == 0.0f) {
       return false;
@@ -145,7 +147,7 @@ ccl_device_inline bool spot_light_sample(const ccl_global KernelLight *klight,
   return true;
 }
 
-ccl_device_forceinline float spot_light_pdf(const float cos_half_spread,
+ccl_device_forceinline float spot_light_pdf(const ccl_global KernelSpotLight *spot,
                                             const float d_sq,
                                             const float r_sq,
                                             const float3 N,
@@ -153,7 +155,8 @@ ccl_device_forceinline float spot_light_pdf(const float cos_half_spread,
                                             const uint32_t path_flag)
 {
   if (d_sq > r_sq) {
-    return M_1_2PI_F / min(sin_sqr_to_one_minus_cos(r_sq / d_sq), 1.0f - cos_half_spread);
+    return M_1_2PI_F /
+           min(sin_sqr_to_one_minus_cos(r_sq / d_sq), 1.0f - spot->cos_half_larger_spread);
   }
 
   const bool has_transmission = (path_flag & PATH_RAY_MIS_HAD_TRANSMISSION);
@@ -181,7 +184,7 @@ ccl_device_forceinline void spot_light_mnee_sample_update(const ccl_global Kerne
     /* NOTE : preserve pdf in area measure. */
     const float jacobian_solid_angle_to_area = 0.5f * fabsf(d_sq - r_sq - t_sq) /
                                                (radius * ls->t * t_sq);
-    ls->pdf = spot_light_pdf(klight->spot.cos_half_spot_angle, d_sq, r_sq, N, ls->D, path_flag) *
+    ls->pdf = spot_light_pdf(&klight->spot, d_sq, r_sq, N, ls->D, path_flag) *
               jacobian_solid_angle_to_area;
 
     ls->Ng = normalize(ls->P - klight->co);
@@ -196,7 +199,7 @@ ccl_device_forceinline void spot_light_mnee_sample_update(const ccl_global Kerne
   }
 
   /* Attenuation. */
-  const float3 local_ray = spot_light_to_local(&klight->spot, -ls->D);
+  const float3 local_ray = spot_light_to_local(klight, -ls->D);
   if (use_attenuation) {
     ls->eval_fac *= spot_light_attenuation(&klight->spot, local_ray);
   }
@@ -232,7 +235,7 @@ ccl_device_inline bool spot_light_sample_from_intersection(
   ls->eval_fac = klight->spot.eval_fac;
 
   if (klight->spot.is_sphere) {
-    ls->pdf = spot_light_pdf(klight->spot.cos_half_spot_angle, d_sq, r_sq, N, ray_D, path_flag);
+    ls->pdf = spot_light_pdf(&klight->spot, d_sq, r_sq, N, ray_D, path_flag);
     ls->Ng = normalize(ls->P - klight->co);
   }
   else {
@@ -248,7 +251,7 @@ ccl_device_inline bool spot_light_sample_from_intersection(
   }
 
   /* Attenuation. */
-  const float3 local_ray = spot_light_to_local(&klight->spot, -ray_D);
+  const float3 local_ray = spot_light_to_local(klight, -ray_D);
   if (!klight->spot.is_sphere || d_sq > r_sq) {
     ls->eval_fac *= spot_light_attenuation(&klight->spot, local_ray);
   }
@@ -262,6 +265,24 @@ ccl_device_inline bool spot_light_sample_from_intersection(
   return true;
 }
 
+/* Find the ray segment lit by the spot light. */
+ccl_device_inline bool spot_light_valid_ray_segment(const ccl_global KernelLight *klight,
+                                                    const float3 P,
+                                                    const float3 D,
+                                                    ccl_private float2 *t_range)
+{
+  /* Convert to local space of the spot light. */
+  const Transform itfm = klight->itfm;
+  float3 local_P = P + klight->spot.dir * klight->spot.ray_segment_dp;
+  local_P = transform_point(&itfm, local_P);
+  const float3 local_D = transform_direction(&itfm, D);
+  const float3 axis = make_float3(0.0f, 0.0f, -1.0f);
+
+  /* Intersect the ray with the smallest enclosing cone of the light spread. */
+  return ray_cone_intersect(
+      axis, local_P, local_D, sqr(klight->spot.cos_half_spot_angle), t_range);
+}
+
 template<bool in_volume_segment>
 ccl_device_forceinline bool spot_light_tree_parameters(const ccl_global KernelLight *klight,
                                                        const float3 centroid,
@@ -270,35 +291,32 @@ ccl_device_forceinline bool spot_light_tree_parameters(const ccl_global KernelLi
                                                        ccl_private float2 &distance,
                                                        ccl_private float3 &point_to_centroid)
 {
-  float dist_point_to_centroid;
-  const float3 point_to_centroid_ = safe_normalize_len(centroid - P, &dist_point_to_centroid);
+  float min_distance;
+  point_to_centroid = safe_normalize_len(centroid - P, &min_distance);
+  distance = min_distance * one_float2();
 
   const float radius = klight->spot.radius;
 
   if (klight->spot.is_sphere) {
-    cos_theta_u = (dist_point_to_centroid > radius) ?
-                      cos_from_sin(radius / dist_point_to_centroid) :
-                      -1.0f;
+    cos_theta_u = (min_distance > radius) ? cos_from_sin(radius / min_distance) : -1.0f;
 
     if (in_volume_segment) {
       return true;
     }
 
-    distance = (dist_point_to_centroid > radius) ?
-                   dist_point_to_centroid * make_float2(1.0f / cos_theta_u, 1.0f) :
-                   one_float2() * radius / M_SQRT2_F;
+    distance = (min_distance > radius) ? min_distance * make_float2(1.0f / cos_theta_u, 1.0f) :
+                                         one_float2() * radius / M_SQRT2_F;
   }
   else {
-    const float hypotenus = sqrtf(sqr(radius) + sqr(dist_point_to_centroid));
-    cos_theta_u = dist_point_to_centroid / hypotenus;
+    const float hypotenus = sqrtf(sqr(radius) + sqr(min_distance));
+    cos_theta_u = min_distance / hypotenus;
 
     if (in_volume_segment) {
       return true;
     }
 
-    distance = make_float2(hypotenus, dist_point_to_centroid);
+    distance.x = hypotenus;
   }
-  point_to_centroid = point_to_centroid_;
 
   return true;
 }

intern/cycles/kernel/light/tree.h

@@ -148,10 +148,7 @@ ccl_device void light_tree_importance(const float3 N_or_D,
   float cos_min_incidence_angle = 1.0f;
   float cos_max_incidence_angle = 1.0f;
 
-  /* When sampling the light tree for the second time in `shade_volume.h` and when query the pdf in
-   * `sample.h`. */
-  const bool in_volume = is_zero(N_or_D);
-  if (!in_volume_segment && !in_volume) {
+  if (!in_volume_segment) {
     const float3 N = N_or_D;
     const float cos_theta_i = has_transmission ? fabsf(dot(point_to_centroid, N)) :
                                                  dot(point_to_centroid, N);
@@ -197,7 +194,7 @@ ccl_device void light_tree_importance(const float3 N_or_D,
   float cos_min_outgoing_angle;
   if ((cos_theta >= cos_theta_u) || (cos_theta_minus_theta_u >= cos_theta_o)) {
     /* theta - theta_o - theta_u <= 0 */
-    kernel_assert((fast_acosf(cos_theta) - bcone.theta_o - fast_acosf(cos_theta_u)) < 5e-4f);
+    kernel_assert((fast_acosf(cos_theta) - bcone.theta_o - fast_acosf(cos_theta_u)) < 1e-3f);
     cos_min_outgoing_angle = 1.0f;
   }
   else if ((bcone.theta_o + bcone.theta_e > M_PI_F) ||
@@ -221,9 +218,9 @@ ccl_device void light_tree_importance(const float3 N_or_D,
   max_importance = fabsf(f_a * cos_min_incidence_angle * energy * cos_min_outgoing_angle /
                          (in_volume_segment ? min_distance : sqr(min_distance)));
 
-  /* TODO: also min importance for volume? */
+  /* TODO: compute proper min importance for volume. */
   if (in_volume_segment) {
-    min_importance = max_importance;
+    min_importance = 0.0f;
     return;
   }
 
@@ -270,10 +267,10 @@ ccl_device bool compute_emitter_centroid_and_dir(KernelGlobals kg,
         /* Arbitrary centroid and direction. */
         centroid = make_float3(0.0f, 0.0f, 1.0f);
         dir = make_float3(0.0f, 0.0f, -1.0f);
-        return !in_volume_segment;
+        break;
       case LIGHT_DISTANT:
         dir = centroid;
-        return !in_volume_segment;
+        break;
       default:
         return false;
     }
@@ -323,23 +320,27 @@ ccl_device void light_tree_node_importance(KernelGlobals kg,
   float cos_theta_u;
   float distance;
   if (knode->type == LIGHT_TREE_DISTANT) {
-    if (in_volume_segment) {
-      return;
-    }
     point_to_centroid = -bcone.axis;
     cos_theta_u = fast_cosf(bcone.theta_o + bcone.theta_e);
     distance = 1.0f;
+    if (t == FLT_MAX) {
+      /* In world volume, distant light has no contribution. */
+      return;
+    }
   }
   else {
     const float3 centroid = 0.5f * (bbox.min + bbox.max);
 
     if (in_volume_segment) {
       const float3 D = N_or_D;
-      const float3 closest_point = P + dot(centroid - P, D) * D;
+      const float closest_t = clamp(dot(centroid - P, D), 0.0f, t);
+      const float3 closest_point = P + D * closest_t;
       /* Minimal distance of the ray to the cluster. */
       distance = len(centroid - closest_point);
+      /* Vector that forms a minimal angle with the emitter centroid. */
       point_to_centroid = -compute_v(centroid, P, D, bcone.axis, t);
-      cos_theta_u = light_tree_cos_bounding_box_angle(bbox, closest_point, point_to_centroid);
+      cos_theta_u = light_tree_cos_bounding_box_angle(
+          bbox, closest_point, normalize(centroid - closest_point));
     }
     else {
       const float3 N = N_or_D;
@@ -404,8 +405,8 @@ ccl_device void light_tree_emitter_importance(KernelGlobals kg,
   bcone.theta_o = kemitter->theta_o;
   bcone.theta_e = kemitter->theta_e;
   float cos_theta_u;
-  float2 distance; /* distance.x = max_distance, distance.y = mix_distance */
-  float3 centroid, point_to_centroid, P_c;
+  float2 distance; /* distance.x = max_distance, distance.y = min_distance */
+  float3 centroid, point_to_centroid, P_c = P;
 
   if (!compute_emitter_centroid_and_dir<in_volume_segment>(kg, kemitter, P, centroid, bcone.axis))
   {
@@ -414,15 +415,9 @@ ccl_device void light_tree_emitter_importance(KernelGlobals kg,
 
   if (in_volume_segment) {
     const float3 D = N_or_D;
-    /* Closest point. */
-    P_c = P + dot(centroid - P, D) * D;
-    /* Minimal distance of the ray to the cluster. */
-    distance.x = len(centroid - P_c);
-    distance.y = distance.x;
-    point_to_centroid = -compute_v(centroid, P, D, bcone.axis, t);
-  }
-  else {
-    P_c = P;
+    /* Closest point from ray to the emitter centroid. */
+    const float closest_t = clamp(dot(centroid - P, D), 0.0f, t);
+    P_c += D * closest_t;
   }
 
   /* Early out if the emitter is guaranteed to be invisible. */
@@ -456,6 +451,9 @@ ccl_device void light_tree_emitter_importance(KernelGlobals kg,
       case LIGHT_DISTANT:
         is_visible = distant_light_tree_parameters(
             centroid, bcone.theta_e, cos_theta_u, distance, point_to_centroid);
+        if (in_volume_segment) {
+          centroid = P - bcone.axis;
+        }
         break;
       default:
         return;
@@ -467,6 +465,11 @@ ccl_device void light_tree_emitter_importance(KernelGlobals kg,
     return;
   }
 
+  if (in_volume_segment) {
+    /* Vector that forms a minimal angle with the emitter centroid. */
+    point_to_centroid = -compute_v(centroid, P, N_or_D, bcone.axis, t);
+  }
+
   light_tree_importance<in_volume_segment>(N_or_D,
                                            has_transmission,
                                            point_to_centroid,
@@ -697,17 +700,16 @@ ccl_device int light_tree_root_node_index(KernelGlobals kg, const int object_rec
   return 0;
 }
 
+/* Pick a random light from the light tree from a given shading point P, write to the picked light
+ * index and the probability of picking the light. */
 template<bool in_volume_segment>
 ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
-                                           const float3 rand,
-                                           const float time,
+                                           const float rand,
                                            const float3 P,
                                            float3 N_or_D,
                                            float t,
                                            const int object_receiver,
                                            const int shader_flags,
-                                           const int bounce,
-                                           const uint32_t path_flag,
                                            ccl_private LightSample *ls)
 {
   if (!kernel_data.integrator.use_direct_light) {
@@ -718,10 +720,8 @@ ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
   float pdf_leaf = 1.0f;
   float pdf_selection = 1.0f;
   int selected_emitter = -1;
-  int object_emitter = 0;
   int node_index = light_tree_root_node_index(kg, object_receiver);
-  /* The first two dimensions of the Sobol sequence have better stratification. */
-  float rand_selection = rand.z;
+  float rand_selection = rand;
 
   float3 local_P = P;
 
@@ -743,7 +743,7 @@ ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
       }
 
       /* Continue with the picked mesh light. */
-      object_emitter = kernel_data_fetch(light_tree_emitters, selected_emitter).mesh.object_id;
+      ls->object = kernel_data_fetch(light_tree_emitters, selected_emitter).mesh.object_id;
       continue;
     }
 
@@ -766,27 +766,18 @@ ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
     pdf_leaf *= (node_index == left_index) ? left_prob : (1.0f - left_prob);
   }
 
-  pdf_selection *= pdf_leaf;
+  ls->emitter_id = selected_emitter;
+  ls->pdf_selection = pdf_selection * pdf_leaf;
 
-  return light_sample<in_volume_segment>(kg,
-                                         float3_to_float2(rand),
-                                         time,
-                                         P,
-                                         N_or_D,
-                                         object_receiver,
-                                         shader_flags,
-                                         bounce,
-                                         path_flag,
-                                         selected_emitter,
-                                         object_emitter,
-                                         pdf_selection,
-                                         ls);
+  return true;
 }
 
 /* We need to be able to find the probability of selecting a given light for MIS. */
+template<bool in_volume_segment>
 ccl_device float light_tree_pdf(KernelGlobals kg,
                                 float3 P,
                                 float3 N,
+                                const float dt,
                                 const int path_flag,
                                 const int object_emitter,
                                 const uint index_emitter,
@@ -796,7 +787,7 @@ ccl_device float light_tree_pdf(KernelGlobals kg,
 
   ccl_global const KernelLightTreeEmitter *kemitter = &kernel_data_fetch(light_tree_emitters,
                                                                          index_emitter);
-  int root_index;
+  int subtree_root_index;
   uint bit_trail, target_emitter;
 
   if (is_triangle(kemitter)) {
@@ -805,16 +796,17 @@ ccl_device float light_tree_pdf(KernelGlobals kg,
     target_emitter = kernel_data_fetch(object_to_tree, object_emitter);
     ccl_global const KernelLightTreeEmitter *kmesh = &kernel_data_fetch(light_tree_emitters,
                                                                         target_emitter);
-    root_index = kmesh->mesh.node_id;
-    ccl_global const KernelLightTreeNode *kroot = &kernel_data_fetch(light_tree_nodes, root_index);
+    subtree_root_index = kmesh->mesh.node_id;
+    ccl_global const KernelLightTreeNode *kroot = &kernel_data_fetch(light_tree_nodes,
+                                                                     subtree_root_index);
     bit_trail = kroot->bit_trail;
 
     if (kroot->type == LIGHT_TREE_INSTANCE) {
-      root_index = kroot->instance.reference;
+      subtree_root_index = kroot->instance.reference;
     }
   }
   else {
-    root_index = 0;
+    subtree_root_index = -1;
     bit_trail = kemitter->bit_trail;
     target_emitter = index_emitter;
   }
@@ -836,8 +828,8 @@ ccl_device float light_tree_pdf(KernelGlobals kg,
       for (int i = 0; i < knode->num_emitters; i++) {
         const int emitter = knode->leaf.first_emitter + i;
         float max_importance, min_importance;
-        light_tree_emitter_importance<false>(
-            kg, P, N, 0, has_transmission, emitter, max_importance, min_importance);
+        light_tree_emitter_importance<in_volume_segment>(
+            kg, P, N, dt, has_transmission, emitter, max_importance, min_importance);
         num_has_importance += (max_importance > 0);
         if (emitter == target_emitter) {
           target_max_importance = max_importance;
@@ -856,13 +848,13 @@ ccl_device float light_tree_pdf(KernelGlobals kg,
         return 0.0f;
       }
 
-      if (root_index) {
+      if (subtree_root_index != -1) {
         /* Arrived at the mesh light. Continue with the subtree. */
         float unused;
-        light_tree_to_local_space<false>(kg, object_emitter, P, N, unused);
+        light_tree_to_local_space<in_volume_segment>(kg, object_emitter, P, N, unused);
 
-        node_index = root_index;
-        root_index = 0;
+        node_index = subtree_root_index;
+        subtree_root_index = -1;
         target_emitter = index_emitter;
         bit_trail = kemitter->bit_trail;
         continue;
@@ -877,8 +869,8 @@ ccl_device float light_tree_pdf(KernelGlobals kg,
     const int right_index = knode->inner.right_child;
 
     float left_prob;
-    if (!get_left_probability<false>(
-            kg, P, N, 0, has_transmission, left_index, right_index, left_prob))
+    if (!get_left_probability<in_volume_segment>(
+            kg, P, N, dt, has_transmission, left_index, right_index, left_prob))
     {
       return 0.0f;
     }
@@ -896,4 +888,27 @@ ccl_device float light_tree_pdf(KernelGlobals kg,
   }
 }
 
+/* If the function is called in volume, retrieve the previous point in volume segment, and compute
+ * pdf from there. Otherwise compute from the current shading point. */
+ccl_device_inline float light_tree_pdf(KernelGlobals kg,
+                                       float3 P,
+                                       float3 N,
+                                       const float dt,
+                                       const int path_flag,
+                                       const int emitter_object,
+                                       const uint emitter_id,
+                                       const int object_receiver)
+{
+  if (path_flag & PATH_RAY_VOLUME_SCATTER) {
+    const float3 D_times_t = N;
+    const float3 D = normalize(D_times_t);
+    P = P - D_times_t;
+    return light_tree_pdf<true>(
+        kg, P, D, dt, path_flag, emitter_object, emitter_id, object_receiver);
+  }
+
+  return light_tree_pdf<false>(
+      kg, P, N, 0.0f, path_flag, emitter_object, emitter_id, object_receiver);
+}
+
 CCL_NAMESPACE_END

intern/cycles/kernel/light/triangle.h

@@ -146,7 +146,9 @@ ccl_device_forceinline bool triangle_light_sample(KernelGlobals kg,
   ls->shader = kernel_data_fetch(tri_shader, prim);
   const float distance_to_plane = dot(N0, V[0] - P) / dot(N0, N0);
   const int ls_shader_flag = kernel_data_fetch(shaders, ls->shader & SHADER_MASK).flags;
-  if (!(ls_shader_flag & (distance_to_plane > 0 ? SD_MIS_BACK : SD_MIS_FRONT))) {
+  if (!in_volume_segment &&
+      !(ls_shader_flag & (distance_to_plane > 0 ? SD_MIS_BACK : SD_MIS_FRONT)))
+  {
     return false;
   }
 
@@ -267,6 +269,26 @@ ccl_device_forceinline bool triangle_light_sample(KernelGlobals kg,
   return (ls->pdf > 0.0f);
 }
 
+/* Find the ray segment lit by the triangle light. */
+ccl_device_inline bool triangle_light_valid_ray_segment(KernelGlobals kg,
+                                                        const float3 P,
+                                                        const float3 D,
+                                                        ccl_private float2 *t_range,
+                                                        const ccl_private LightSample *ls)
+{
+  const int shader_flag = kernel_data_fetch(shaders, ls->shader & SHADER_MASK).flags;
+  const int SD_MIS_BOTH = SD_MIS_BACK | SD_MIS_FRONT;
+  if ((shader_flag & SD_MIS_BOTH) == SD_MIS_BOTH) {
+    /* Both sides are sampled, the complete ray segment is visible. */
+    return true;
+  }
+
+  /* Only one side is sampled, intersect the ray and the triangle light plane to find the visible
+   * ray segment. Flip normal if Emission Sampling is set to back. */
+  const float3 N = ls->Ng;
+  return ray_plane_intersect((shader_flag & SD_MIS_BACK) ? -N : N, P, D, t_range);
+}
+
 template<bool in_volume_segment>
 ccl_device_forceinline bool triangle_light_tree_parameters(
     KernelGlobals kg,
@@ -279,13 +301,11 @@ ccl_device_forceinline bool triangle_light_tree_parameters(
     ccl_private float2 &distance,
     ccl_private float3 &point_to_centroid)
 {
-  if (!in_volume_segment) {
-    /* TODO: a cheap substitute for minimal distance between point and primitive. Does it
-     * worth the overhead to compute the accurate minimal distance? */
-    float min_distance;
-    point_to_centroid = safe_normalize_len(centroid - P, &min_distance);
-    distance = make_float2(min_distance, min_distance);
-  }
+  /* TODO: a cheap substitute for minimal distance between point and primitive. Does it worth the
+   * overhead to compute the accurate minimal distance? */
+  float min_distance;
+  point_to_centroid = safe_normalize_len(centroid - P, &min_distance);
+  distance = make_float2(min_distance, min_distance);
 
   cos_theta_u = FLT_MAX;
 
@@ -305,9 +325,8 @@ ccl_device_forceinline bool triangle_light_tree_parameters(
   }
 
   const bool front_facing = bcone.theta_o != 0.0f || dot(bcone.axis, point_to_centroid) < 0;
-  const bool in_volume = is_zero(N);
 
-  return (front_facing && shape_above_surface) || in_volume;
+  return front_facing && shape_above_surface;
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/osl/shaders/node_noise.h

@@ -7,76 +7,82 @@
 
 #define vector3 point
 
-float safe_noise(float p)
+float safe_noise(float co)
 {
-  float f = noise("noise", p);
-  if (isinf(f)) {
-    return 0.5;
-  }
-  return f;
+  /* Repeat Perlin noise texture every 100000.0 on each axis to prevent floating point
+   * representation issues. */
+  float p = fmod(co, 100000.0);
+
+  return noise("noise", p);
 }
 
-float safe_noise(vector2 p)
+float safe_noise(vector2 co)
 {
-  float f = noise("noise", p.x, p.y);
-  if (isinf(f)) {
-    return 0.5;
-  }
-  return f;
+  /* Repeat Perlin noise texture every 100000.0 on each axis to prevent floating point
+   * representation issues. This causes discontinuities every 100000.0, however at such scales this
+   * usually shouldn't be noticeable. */
+  vector2 p = fmod(co, 100000.0);
+
+  return noise("noise", p.x, p.y);
 }
 
-float safe_noise(vector3 p)
+float safe_noise(vector3 co)
 {
-  float f = noise("noise", p);
-  if (isinf(f)) {
-    return 0.5;
-  }
-  return f;
+  /* Repeat Perlin noise texture every 100000.0 on each axis to prevent floating point
+   * representation issues. This causes discontinuities every 100000.0, however at such scales this
+   * usually shouldn't be noticeable. */
+  vector3 p = fmod(co, 100000.0);
+
+  return noise("noise", p);
 }
 
-float safe_noise(vector4 p)
+float safe_noise(vector4 co)
 {
-  float f = noise("noise", vector3(p.x, p.y, p.z), p.w);
-  if (isinf(f)) {
-    return 0.5;
-  }
-  return f;
+  /* Repeat Perlin noise texture every 100000.0 on each axis to prevent floating point
+   * representation issues. This causes discontinuities every 100000.0, however at such scales this
+   * usually shouldn't be noticeable. */
+  vector4 p = fmod(co, 100000.0);
+
+  return noise("noise", vector3(p.x, p.y, p.z), p.w);
 }
 
-float safe_snoise(float p)
+float safe_snoise(float co)
 {
-  float f = noise("snoise", p);
-  if (isinf(f)) {
-    return 0.0;
-  }
-  return f;
+  /* Repeat Perlin noise texture every 100000.0 on each axis to prevent floating point
+   * representation issues. */
+  float p = fmod(co, 100000.0);
+
+  return noise("snoise", p);
 }
 
-float safe_snoise(vector2 p)
+float safe_snoise(vector2 co)
 {
-  float f = noise("snoise", p.x, p.y);
-  if (isinf(f)) {
-    return 0.0;
-  }
-  return f;
+  /* Repeat Perlin noise texture every 100000.0 on each axis to prevent floating point
+   * representation issues. This causes discontinuities every 100000.0, however at such scales this
+   * usually shouldn't be noticeable. */
+  vector2 p = fmod(co, 100000.0);
+
+  return noise("snoise", p.x, p.y);
 }
 
-float safe_snoise(vector3 p)
+float safe_snoise(vector3 co)
 {
-  float f = noise("snoise", p);
-  if (isinf(f)) {
-    return 0.0;
-  }
-  return f;
+  /* Repeat Perlin noise texture every 100000.0 on each axis to prevent floating point
+   * representation issues. This causes discontinuities every 100000.0, however at such scales this
+   * usually shouldn't be noticeable. */
+  vector3 p = fmod(co, 100000.0);
+
+  return noise("snoise", p);
 }
 
-float safe_snoise(vector4 p)
+float safe_snoise(vector4 co)
 {
-  float f = noise("snoise", vector3(p.x, p.y, p.z), p.w);
-  if (isinf(f)) {
-    return 0.0;
-  }
-  return f;
+  /* Repeat Perlin noise texture every 100000.0 on each axis to prevent floating point
+   * representation issues. This causes discontinuities every 100000.0, however at such scales this
+   * usually shouldn't be noticeable. */
+  vector4 p = fmod(co, 100000.0);
+
+  return noise("snoise", vector3(p.x, p.y, p.z), p.w);
 }
 
 #define NOISE_FBM(T) \

intern/cycles/kernel/svm/noise.h

@@ -684,7 +684,12 @@ ccl_device_inline float noise_scale4(float result)
 
 ccl_device_inline float snoise_1d(float p)
 {
-  return noise_scale1(ensure_finite(perlin_1d(p)));
+  /* Repeat Perlin noise texture every 100000.0 on each axis to prevent floating point
+   * representation issues. */
+  /* The 1D variant of fmod is called fmodf. */
+  p = fmodf(p, 100000.0f);
+
+  return noise_scale1(perlin_1d(p));
 }
 
 ccl_device_inline float noise_1d(float p)
@@ -694,7 +699,12 @@ ccl_device_inline float noise_1d(float p)
 
 ccl_device_inline float snoise_2d(float2 p)
 {
-  return noise_scale2(ensure_finite(perlin_2d(p.x, p.y)));
+  /* Repeat Perlin noise texture every 100000.0f on each axis to prevent floating point
+   * representation issues. This causes discontinuities every 100000.0f, however at such scales
+   * this usually shouldn't be noticeable. */
+  p = fmod(p, 100000.0f);
+
+  return noise_scale2(perlin_2d(p.x, p.y));
 }
 
 ccl_device_inline float noise_2d(float2 p)
@@ -704,7 +714,12 @@ ccl_device_inline float noise_2d(float2 p)
 
 ccl_device_inline float snoise_3d(float3 p)
 {
-  return noise_scale3(ensure_finite(perlin_3d(p.x, p.y, p.z)));
+  /* Repeat Perlin noise texture every 100000.0f on each axis to prevent floating point
+   * representation issues. This causes discontinuities every 100000.0f, however at such scales
+   * this usually shouldn't be noticeable. */
+  p = fmod(p, 100000.0f);
+
+  return noise_scale3(perlin_3d(p.x, p.y, p.z));
 }
 
 ccl_device_inline float noise_3d(float3 p)
@@ -714,7 +729,12 @@ ccl_device_inline float noise_3d(float3 p)
 
 ccl_device_inline float snoise_4d(float4 p)
 {
-  return noise_scale4(ensure_finite(perlin_4d(p.x, p.y, p.z, p.w)));
+  /* Repeat Perlin noise texture every 100000.0f on each axis to prevent floating point
+   * representation issues. This causes discontinuities every 100000.0f, however at such scales
+   * this usually shouldn't be noticeable. */
+  p = fmod(p, 100000.0f);
+
+  return noise_scale4(perlin_4d(p.x, p.y, p.z, p.w));
 }
 
 ccl_device_inline float noise_4d(float4 p)

intern/cycles/kernel/types.h

@@ -45,6 +45,7 @@ CCL_NAMESPACE_BEGIN
 #define OBJECT_NONE (~0)
 #define PRIM_NONE (~0)
 #define LAMP_NONE (~0)
+#define EMITTER_NONE (~0)
 #define ID_NONE (0.0f)
 #define PASS_UNUSED (~0)
 #define LIGHTGROUP_NONE (~0)
@@ -1367,16 +1368,17 @@ typedef struct KernelCurveSegment {
 static_assert_align(KernelCurveSegment, 8);
 
 typedef struct KernelSpotLight {
-  packed_float3 scaled_axis_u;
-  float radius;
-  packed_float3 scaled_axis_v;
-  float eval_fac;
   packed_float3 dir;
+  float radius;
+  float eval_fac;
   float cos_half_spot_angle;
   float half_cot_half_spot_angle;
-  float inv_len_z;
   float spot_smooth;
   int is_sphere;
+  /* For non-uniform object scaling, the actual spread might be different. */
+  float cos_half_larger_spread;
+  /* Distance from the apex of the smallest enclosing cone of the light spread to light center. */
+  float ray_segment_dp;
 } KernelSpotLight;
 
 /* PointLight is SpotLight with only radius and invarea being used. */

intern/cycles/scene/CMakeLists.txt

@@ -107,6 +107,7 @@ endif()
 if(WITH_CYCLES_OSL)
   list(APPEND LIB
     cycles_kernel_osl
+    ${OSL_LIBRARIES}
   )
 endif()
 

intern/cycles/scene/light.cpp

@@ -903,14 +903,14 @@ static void background_cdf(
   for (int i = start; i < end; i++) {
     float sin_theta = sinf(M_PI_F * (i + 0.5f) / res_y);
     float3 env_color = (*pixels)[i * res_x];
-    float ave_luminance = average(env_color);
+    float ave_luminance = average(fabs(env_color));
 
     cond_cdf[i * cdf_width].x = ave_luminance * sin_theta;
     cond_cdf[i * cdf_width].y = 0.0f;
 
     for (int j = 1; j < res_x; j++) {
       env_color = (*pixels)[i * res_x + j];
-      ave_luminance = average(env_color);
+      ave_luminance = average(fabs(env_color));
 
       cond_cdf[i * cdf_width + j].x = ave_luminance * sin_theta;
       cond_cdf[i * cdf_width + j].y = cond_cdf[i * cdf_width + j - 1].y +
@@ -1346,23 +1346,25 @@ void LightManager::device_update_lights(Device *device, DeviceScene *dscene, Sce
       klights[light_index].area.normalize_spread = normalize_spread;
     }
     if (light->light_type == LIGHT_SPOT) {
-      /* Scale axes to accommodate non-uniform scaling. */
-      float3 scaled_axis_u = light->get_axisu() / len_squared(light->get_axisu());
-      float3 scaled_axis_v = light->get_axisv() / len_squared(light->get_axisv());
-      float len_z;
-      /* Keep direction normalized. */
-      float3 dir = safe_normalize_len(light->get_dir(), &len_z);
+      const float cos_half_spot_angle = cosf(light->spot_angle * 0.5f);
+      const float spot_smooth = 1.0f / ((1.0f - cos_half_spot_angle) * light->spot_smooth);
+      const float tan_half_spot_angle = tanf(light->spot_angle * 0.5f);
 
-      float cos_half_spot_angle = cosf(light->spot_angle * 0.5f);
-      float spot_smooth = 1.0f / ((1.0f - cos_half_spot_angle) * light->spot_smooth);
+      const float len_w_sq = len_squared(light->get_dir());
+      const float len_u_sq = len_squared(light->get_axisu());
+      const float len_v_sq = len_squared(light->get_axisv());
+      const float tan_sq = sqr(tan_half_spot_angle);
 
-      klights[light_index].spot.scaled_axis_u = scaled_axis_u;
-      klights[light_index].spot.scaled_axis_v = scaled_axis_v;
-      klights[light_index].spot.dir = dir;
+      klights[light_index].spot.dir = safe_normalize(light->get_dir());
       klights[light_index].spot.cos_half_spot_angle = cos_half_spot_angle;
-      klights[light_index].spot.half_cot_half_spot_angle = 0.5f / tanf(light->spot_angle * 0.5f);
-      klights[light_index].spot.inv_len_z = 1.0f / len_z;
+      klights[light_index].spot.half_cot_half_spot_angle = 0.5f / tan_half_spot_angle;
       klights[light_index].spot.spot_smooth = spot_smooth;
+      /* Choose the angle which spans a larger cone. */
+      klights[light_index].spot.cos_half_larger_spread = inversesqrtf(
+          1.0f + tan_sq * fmaxf(len_u_sq, len_v_sq) / len_w_sq);
+      /* radius / sin(half_angle_small) */
+      klights[light_index].spot.ray_segment_dp =
+          light->size * sqrtf(1.0f + len_w_sq / (tan_sq * fminf(len_u_sq, len_v_sq)));
     }
 
     klights[light_index].shader_id = shader_id;

intern/cycles/session/tile.cpp

@@ -39,7 +39,7 @@ static std::atomic<uint64_t> g_instance_index = 0;
 /* Construct names of EXR channels which will ensure order of all channels to match exact offsets
  * in render buffers corresponding to the given passes.
  *
- * Returns `std` datatypes so that it can be assigned directly to the OIIO's `ImageSpec`. */
+ * Returns `std` data-types so that it can be assigned directly to the OIIO's `ImageSpec`. */
 static std::vector<std::string> exr_channel_names_for_passes(const BufferParams &buffer_params)
 {
   static const char *component_suffixes[] = {"R", "G", "B", "A"};

intern/cycles/util/color.h

@@ -275,8 +275,8 @@ ccl_device_inline float4 improve_5throot_solution_sse2(const float4 &old_result,
 /* Calculate powf(x, 2.4). Working domain: 1e-10 < x < 1e+10 */
 ccl_device_inline float4 fastpow24_sse2(const float4 &arg)
 {
-  /* max, avg and |avg| errors were calculated in gcc without FMA instructions
-   * The final precision should be better than powf in glibc */
+  /* `max`, `avg` and |avg| errors were calculated in GCC without FMA instructions.
+   * The final precision should be better than `powf` in GLIBC. */
 
   /* Calculate x^4/5, coefficient 0.994 was constructed manually to minimize avg error */
   /* 0x3F4CCCCD = 4/5 */

intern/cycles/util/math.h

@@ -1030,6 +1030,46 @@ ccl_device_inline uint32_t reverse_integer_bits(uint32_t x)
 #endif
 }
 
+/* Check if intervals (first->x, first->y) and (second.x, second.y) intersect, and replace the
+ * first interval with their intersection. */
+ccl_device_inline bool intervals_intersect(ccl_private float2 *first, const float2 second)
+{
+  first->x = fmaxf(first->x, second.x);
+  first->y = fminf(first->y, second.y);
+
+  return first->x < first->y;
+}
+
+/* Solve quadratic equation a*x^2 + b*x + c = 0, adapted from Mitsuba 3
+ * The solution is ordered so that x1 <= x2.
+ * Returns true if at least one solution is found.  */
+ccl_device_inline bool solve_quadratic(
+    const float a, const float b, const float c, ccl_private float &x1, ccl_private float &x2)
+{
+  /* If the equation is linear, the solution is -c/b, but b has to be non-zero. */
+  const bool valid_linear = (a == 0.0f) && (b != 0.0f);
+  x1 = x2 = -c / b;
+
+  const float discriminant = sqr(b) - 4.0f * a * c;
+  /* Allow slightly negative discriminant in case of numerical precision issues. */
+  const bool valid_quadratic = (a != 0.0f) && (discriminant > -1e-5f);
+
+  if (valid_quadratic) {
+    /* Numerically stable version of (-b ± sqrt(discriminant)) / (2 * a), avoiding catastrophic
+     * cancellation when `b` is very close to `sqrt(discriminant)`, by finding the solution of
+     * greater magnitude which does not suffer from loss of precision, then using the identity
+     * x1 * x2 = c / a. */
+    const float temp = -0.5f * (b + copysignf(safe_sqrtf(discriminant), b));
+    const float r1 = temp / a;
+    const float r2 = c / temp;
+
+    x1 = fminf(r1, r2);
+    x2 = fmaxf(r1, r2);
+  }
+
+  return (valid_linear || valid_quadratic);
+}
+
 CCL_NAMESPACE_END
 
 #endif /* __UTIL_MATH_H__ */

intern/cycles/util/math_float2.h

@@ -198,6 +198,11 @@ ccl_device_inline float2 clamp(const float2 a, const float2 mn, const float2 mx)
   return min(max(a, mn), mx);
 }
 
+ccl_device_inline float2 fmod(const float2 a, const float b)
+{
+  return make_float2(fmodf(a.x, b), fmodf(a.y, b));
+}
+
 ccl_device_inline float2 fabs(const float2 a)
 {
   return make_float2(fabsf(a.x), fabsf(a.y));

intern/cycles/util/math_float3.h

@@ -309,6 +309,11 @@ ccl_device_inline float3 fabs(const float3 a)
 #  endif
 }
 
+ccl_device_inline float3 fmod(const float3 a, const float b)
+{
+  return make_float3(fmodf(a.x, b), fmodf(a.y, b), fmodf(a.z, b));
+}
+
 ccl_device_inline float3 sqrt(const float3 a)
 {
 #  ifdef __KERNEL_SSE__

intern/cycles/util/math_float4.h

@@ -465,6 +465,11 @@ ccl_device_inline float4 fabs(const float4 a)
 #  endif
 }
 
+ccl_device_inline float4 fmod(const float4 a, const float b)
+{
+  return make_float4(fmodf(a.x, b), fmodf(a.y, b), fmodf(a.z, b), fmodf(a.w, b));
+}
+
 ccl_device_inline float4 floor(const float4 a)
 {
 #  ifdef __KERNEL_SSE__

intern/cycles/util/math_intersect.h

@@ -302,6 +302,140 @@ ccl_device bool ray_quad_intersect(float3 ray_P,
   return true;
 }
 
+/* Find the ray segment that lies in the same side as the normal `N` of the plane.
+ * `P` is the vector pointing from any point on the plane to the ray origin. */
+ccl_device bool ray_plane_intersect(const float3 N,
+                                    const float3 P,
+                                    const float3 ray_D,
+                                    ccl_private float2 *t_range)
+{
+  const float DN = dot(ray_D, N);
+
+  /* Distance from P to the plane. */
+  const float t = -dot(P, N) / DN;
+
+  /* Limit the range to the positive side. */
+  if (DN > 0.0f) {
+    t_range->x = fmaxf(t_range->x, t);
+  }
+  else {
+    t_range->y = fminf(t_range->y, t);
+  }
+
+  return t_range->x < t_range->y;
+}
+
+/* Find the ray segment inside an axis-aligned bounding box. */
+ccl_device bool ray_aabb_intersect(const float3 bbox_min,
+                                   const float3 bbox_max,
+                                   const float3 ray_P,
+                                   const float3 ray_D,
+                                   ccl_private float2 *t_range)
+{
+  const float3 inv_ray_D = rcp(ray_D);
+
+  /* Absolute distances to lower and upper box coordinates; */
+  const float3 t_lower = (bbox_min - ray_P) * inv_ray_D;
+  const float3 t_upper = (bbox_max - ray_P) * inv_ray_D;
+
+  /* The four t-intervals (for x-/y-/z-slabs, and ray p(t)). */
+  const float4 tmins = float3_to_float4(min(t_lower, t_upper), t_range->x);
+  const float4 tmaxes = float3_to_float4(max(t_lower, t_upper), t_range->y);
+
+  /* Max of mins and min of maxes. */
+  const float tmin = reduce_max(tmins);
+  const float tmax = reduce_min(tmaxes);
+
+  *t_range = make_float2(tmin, tmax);
+
+  return tmin < tmax;
+}
+
+/* Find the segment of a ray defined by P + D * t that lies inside a cylinder defined by
+ * (x / len_u)^2 + (y / len_v)^2 = 1. */
+ccl_device_inline bool ray_infinite_cylinder_intersect(const float3 P,
+                                                       const float3 D,
+                                                       const float len_u,
+                                                       const float len_v,
+                                                       ccl_private float2 *t_range)
+{
+  /* Convert to a 2D problem. */
+  const float2 inv_len = 1.0f / make_float2(len_u, len_v);
+  float2 P_proj = float3_to_float2(P) * inv_len;
+  const float2 D_proj = float3_to_float2(D) * inv_len;
+
+  /* Solve quadratic equation a*t^2 + 2b*t + c = 0. */
+  const float a = dot(D_proj, D_proj);
+  float b = dot(P_proj, D_proj);
+
+  /* Move ray origin closer to the cylinder to prevent precision issue when the ray is far away. */
+  const float t_mid = -b / a;
+  P_proj += D_proj * t_mid;
+
+  /* Recompute b from the shifted origin. */
+  b = dot(P_proj, D_proj);
+  const float c = dot(P_proj, P_proj) - 1.0f;
+
+  float tmin, tmax;
+  const bool valid = solve_quadratic(a, 2.0f * b, c, tmin, tmax);
+
+  return valid && intervals_intersect(t_range, make_float2(tmin, tmax) + t_mid);
+}
+
+/* *
+ * Find the ray segment inside a single-sided cone.
+ *
+ * \param axis: a unit-length direction around which the cone has a circular symmetry
+ * \param P: the vector pointing from the cone apex to the ray origin
+ * \param D: the direction of the ray, does not need to have unit-length
+ * \param cos_angle_sq: `sqr(cos(half_aperture_of_the_cone))`
+ * \param t_range: the lower and upper bounds between which the ray lies inside the cone
+ * \return whether the intersection exists and is in the provided range
+ *
+ * See https://www.geometrictools.com/Documentation/IntersectionLineCone.pdf for illustration
+ */
+ccl_device_inline bool ray_cone_intersect(const float3 axis,
+                                          const float3 P,
+                                          float3 D,
+                                          const float cos_angle_sq,
+                                          ccl_private float2 *t_range)
+{
+  if (cos_angle_sq < 1e-4f) {
+    /* The cone is nearly a plane. */
+    return ray_plane_intersect(axis, P, D, t_range);
+  }
+
+  const float inv_len = inversesqrtf(len_squared(D));
+  D *= inv_len;
+
+  const float AD = dot(axis, D);
+  const float AP = dot(axis, P);
+
+  const float a = sqr(AD) - cos_angle_sq;
+  const float b = 2.0f * (AD * AP - cos_angle_sq * dot(D, P));
+  const float c = sqr(AP) - cos_angle_sq * dot(P, P);
+
+  float tmin = 0.0f, tmax = FLT_MAX;
+  bool valid = solve_quadratic(a, b, c, tmin, tmax);
+
+  /* Check if the intersections are in the same hemisphere as the cone. */
+  const bool tmin_valid = AP + tmin * AD > 0.0f;
+  const bool tmax_valid = AP + tmax * AD > 0.0f;
+
+  valid &= (tmin_valid || tmax_valid);
+
+  if (!tmax_valid) {
+    tmax = tmin;
+    tmin = 0.0f;
+  }
+  else if (!tmin_valid) {
+    tmin = tmax;
+    tmax = FLT_MAX;
+  }
+
+  return valid && intervals_intersect(t_range, make_float2(tmin, tmax) * inv_len);
+}
+
 CCL_NAMESPACE_END
 
 #endif /* __UTIL_MATH_INTERSECT_H__ */

intern/cycles/util/thread.cpp

@@ -14,7 +14,7 @@ CCL_NAMESPACE_BEGIN
 thread::thread(function<void()> run_cb) : run_cb_(run_cb), joined_(false)
 {
 #if defined(__APPLE__) || defined(__linux__) && !defined(__GLIBC__)
-  /* Set the stack size to 2MB to match glibc. The default 512KB on macOS is
+  /* Set the stack size to 2MB to match GLIBC. The default 512KB on macOS is
    * too small for Embree, and consistent stack size also makes things more
    * predictable in general. */
   pthread_attr_t attribute;

intern/cycles/util/transform.h

@@ -161,6 +161,17 @@ ccl_device_inline Transform make_transform(float a,
   return t;
 }
 
+ccl_device_inline Transform make_transform(const float3 x, const float3 y, const float3 z)
+{
+  Transform t;
+
+  t.x = float3_to_float4(x, 0.0f);
+  t.y = float3_to_float4(y, 0.0f);
+  t.z = float3_to_float4(z, 0.0f);
+
+  return t;
+}
+
 ccl_device_inline Transform euler_to_transform(const float3 euler)
 {
   float cx = cosf(euler.x);

intern/ghost/GHOST_Path-api.hh

@@ -10,10 +10,6 @@
 
 #include "GHOST_Types.h"
 
-#ifdef __cplusplus
-extern "C" {
-#endif
-
 GHOST_DECLARE_HANDLE(GHOST_SystemPathsHandle);
 
 /**
@@ -63,7 +59,3 @@ extern const char *GHOST_getBinaryDir();
  * Add the file to the operating system most recently used files
  */
 extern void GHOST_addToSystemRecentFiles(const char *filepath);
-
-#ifdef __cplusplus
-}
-#endif

intern/ghost/intern/GHOST_SystemWayland.cc

@@ -108,6 +108,9 @@ static bool has_libdecor = true;
 #  endif
 #endif
 
+#include "IMB_imbuf.hh"
+#include "IMB_imbuf_types.hh"
+
 /* -------------------------------------------------------------------- */
 /** \name Forward Declarations
  * \{ */
@@ -424,7 +427,7 @@ struct GWL_Cursor {
 
   /**
    * The name of the theme (set by an environment variable).
-   * When disabled, leave as an empty string and the default theme will be used.
+   * When disabled, leave as an empty string and pass in nullptr to use the default theme.
    */
   std::string theme_name;
   /**
@@ -2504,7 +2507,9 @@ static const wl_cursor *gwl_seat_cursor_find_from_shape(GWL_Seat *seat,
     if (!cursor->wl.theme) {
       /* The cursor wl_surface hasn't entered an output yet. Initialize theme with scale 1. */
       cursor->wl.theme = wl_cursor_theme_load(
-          cursor->theme_name.c_str(), cursor->theme_size, seat->system->wl_shm_get());
+          (cursor->theme_name.empty() ? nullptr : cursor->theme_name.c_str()),
+          cursor->theme_size,
+          seat->system->wl_shm_get());
     }
 
     if (cursor->wl.theme) {
@@ -3518,7 +3523,9 @@ static bool update_cursor_scale(GWL_Cursor &cursor,
     }
     wl_cursor_theme_destroy(cursor.wl.theme);
     cursor.wl.theme = wl_cursor_theme_load(
-        cursor.theme_name.c_str(), scale * cursor.theme_size, shm);
+        (cursor.theme_name.empty() ? nullptr : cursor.theme_name.c_str()),
+        scale * cursor.theme_size,
+        shm);
     if (cursor.wl.theme_cursor) {
       cursor.wl.theme_cursor = wl_cursor_theme_get_cursor(cursor.wl.theme,
                                                           cursor.wl.theme_cursor_name);
@@ -7473,6 +7480,148 @@ void GHOST_SystemWayland::putClipboard(const char *buffer, bool selection) const
   }
 }
 
+static constexpr const char *ghost_wl_mime_img_png = "image/png";
+
+GHOST_TSuccess GHOST_SystemWayland::hasClipboardImage(void) const
+{
+  GWL_Seat *seat = gwl_display_seat_active_get(display_);
+  if (UNLIKELY(!seat)) {
+    return GHOST_kFailure;
+  }
+
+  GWL_DataOffer *data_offer = seat->data_offer_copy_paste;
+  if (data_offer) {
+    if (data_offer->types.count(ghost_wl_mime_img_png)) {
+      return GHOST_kSuccess;
+    }
+  }
+
+  return GHOST_kFailure;
+}
+
+uint *GHOST_SystemWayland::getClipboardImage(int *r_width, int *r_height) const
+{
+#ifdef USE_EVENT_BACKGROUND_THREAD
+  std::lock_guard lock_server_guard{*server_mutex};
+#endif
+
+  GWL_Seat *seat = gwl_display_seat_active_get(display_);
+  if (UNLIKELY(!seat)) {
+    return nullptr;
+  }
+
+  std::mutex &mutex = seat->data_offer_copy_paste_mutex;
+  mutex.lock();
+  bool mutex_locked = true;
+
+  uint *rgba = nullptr;
+
+  GWL_DataOffer *data_offer = seat->data_offer_copy_paste;
+  if (data_offer) {
+    /* Check if the source offers a supported mime type.
+     * This check could be skipped, because the paste option is not supposed to be enabled
+     * otherwise. */
+    if (data_offer->types.count(ghost_wl_mime_img_png)) {
+      /* Receive the clipboard in a thread, performing round-trips while waiting,
+       * so pasting content from own `primary->data_source` doesn't hang. */
+      struct ThreadResult {
+        char *data = nullptr;
+        size_t data_len = 0;
+        std::atomic<bool> done = false;
+      } thread_result;
+
+      auto read_clipboard_fn = [](GWL_DataOffer *data_offer,
+                                  const char *mime_receive,
+                                  std::mutex *mutex,
+                                  ThreadResult *thread_result) {
+        thread_result->data = read_buffer_from_data_offer(
+            data_offer, mime_receive, mutex, false, &thread_result->data_len);
+        thread_result->done = true;
+      };
+      std::thread read_thread(
+          read_clipboard_fn, data_offer, ghost_wl_mime_img_png, &mutex, &thread_result);
+      read_thread.detach();
+
+      while (!thread_result.done) {
+        wl_display_roundtrip(display_->wl.display);
+      }
+
+      if (thread_result.data) {
+        /* Generate the image buffer with the received data. */
+        ImBuf *ibuf = IMB_ibImageFromMemory((uint8_t *)thread_result.data,
+                                            thread_result.data_len,
+                                            IB_rect,
+                                            nullptr,
+                                            "<clipboard>");
+        if (ibuf) {
+          *r_width = ibuf->x;
+          *r_height = ibuf->y;
+          const size_t byte_count = size_t(ibuf->x) * size_t(ibuf->y) * 4;
+          rgba = (uint *)malloc(byte_count);
+          std::memcpy(rgba, ibuf->byte_buffer.data, byte_count);
+          IMB_freeImBuf(ibuf);
+        }
+      }
+
+      /* After reading the data offer, the mutex gets unlocked. */
+      mutex_locked = false;
+    }
+  }
+
+  if (mutex_locked) {
+    mutex.unlock();
+  }
+  return rgba;
+}
+
+GHOST_TSuccess GHOST_SystemWayland::putClipboardImage(uint *rgba, int width, int height) const
+{
+#ifdef USE_EVENT_BACKGROUND_THREAD
+  std::lock_guard lock_server_guard{*server_mutex};
+#endif
+
+  /* Create a #wl_data_source object. */
+  GWL_Seat *seat = gwl_display_seat_active_get(display_);
+  if (UNLIKELY(!seat)) {
+    return GHOST_kFailure;
+  }
+  std::lock_guard lock(seat->data_source_mutex);
+
+  GWL_DataSource *data_source = seat->data_source;
+
+  /* Load buffer into an #ImBuf and convert to PNG. */
+  ImBuf *ibuf = IMB_allocFromBuffer(reinterpret_cast<uint8_t *>(rgba), nullptr, width, height, 32);
+  ibuf->ftype = IMB_FTYPE_PNG;
+  ibuf->foptions.quality = 15;
+  if (!IMB_saveiff(ibuf, "<memory>", IB_rect | IB_mem)) {
+    IMB_freeImBuf(ibuf);
+    return GHOST_kFailure;
+  }
+
+  /* Copy #ImBuf encoded_buffer to data source. */
+  GWL_SimpleBuffer *imgbuffer = &data_source->buffer_out;
+  gwl_simple_buffer_free_data(imgbuffer);
+  imgbuffer->data_size = ibuf->encoded_buffer_size;
+  char *data = static_cast<char *>(malloc(imgbuffer->data_size));
+  std::memcpy(data, ibuf->encoded_buffer.data, ibuf->encoded_buffer_size);
+  imgbuffer->data = data;
+
+  data_source->wl.source = wl_data_device_manager_create_data_source(
+      display_->wl.data_device_manager);
+  wl_data_source_add_listener(data_source->wl.source, &data_source_listener, seat);
+
+  /* Advertise the mime types supported. */
+  wl_data_source_offer(data_source->wl.source, ghost_wl_mime_img_png);
+
+  if (seat->wl.data_device) {
+    wl_data_device_set_selection(
+        seat->wl.data_device, data_source->wl.source, seat->data_source_serial);
+  }
+
+  IMB_freeImBuf(ibuf);
+  return GHOST_kSuccess;
+}
+
 uint8_t GHOST_SystemWayland::getNumDisplays() const
 {
 #ifdef USE_EVENT_BACKGROUND_THREAD
@@ -8048,9 +8197,7 @@ GHOST_TCapabilityFlag GHOST_SystemWayland::getCapabilities() const
            * is negligible. */
           GHOST_kCapabilityGPUReadFrontBuffer |
           /* This WAYLAND back-end has not yet implemented desktop color sample. */
-          GHOST_kCapabilityDesktopSample |
-          /* This WAYLAND back-end has not yet implemented image copy/paste. */
-          GHOST_kCapabilityClipboardImages));
+          GHOST_kCapabilityDesktopSample));
 }
 
 bool GHOST_SystemWayland::cursor_grab_use_software_display_get(const GHOST_TGrabCursorMode mode)
@@ -8114,8 +8261,10 @@ void GHOST_SystemWayland::setMultitouchGestures(const bool use)
   }
   m_multitouchGestures = use;
 
+#ifdef USE_EVENT_BACKGROUND_THREAD
   /* Ensure this listeners aren't removed while events are generated. */
   std::lock_guard lock_server_guard{*server_mutex};
+#endif
   for (GWL_Seat *seat : display_->seats) {
     if (use == gwl_seat_capability_pointer_multitouch_check(seat, use)) {
       continue;

intern/ghost/intern/GHOST_SystemWayland.hh

@@ -159,6 +159,27 @@ class GHOST_SystemWayland : public GHOST_System {
 
   void putClipboard(const char *buffer, bool selection) const override;
 
+  /**
+   * Returns GHOST_kSuccess if the clipboard contains an image.
+   */
+  GHOST_TSuccess hasClipboardImage() const override;
+
+  /**
+   * Get image data from the Clipboard
+   * \param r_width: the returned image width in pixels.
+   * \param r_height: the returned image height in pixels.
+   * \return pointer uint array in RGBA byte order. Caller must free.
+   */
+  uint *getClipboardImage(int *r_width, int *r_height) const override;
+
+  /**
+   * Put image data to the Clipboard
+   * \param rgba: uint array in RGBA byte order.
+   * \param width: the image width in pixels.
+   * \param height: the image height in pixels.
+   */
+  GHOST_TSuccess putClipboardImage(uint *rgba, int width, int height) const override;
+
   uint8_t getNumDisplays() const override;
 
   uint64_t getMilliSeconds() const override;