CMakeLists.txt

@@ -159,6 +159,15 @@ endif()
 
 get_blender_version()
 
+if(WIN32)
+  add_definitions(
+    # This is the app ID used for file registration, given it's used from several modules
+    # there really is no nice way to get this information consistent without a global define.
+    -DBLENDER_WIN_APPID="blender.${BLENDER_VERSION_MAJOR}.${BLENDER_VERSION_MINOR}"
+    # This is the name that will be shown in the taskbar and OpenWith windows UI
+    -DBLENDER_WIN_APPID_FRIENDLY_NAME="Blender ${BLENDER_VERSION_MAJOR}.${BLENDER_VERSION_MINOR}"
+  )
+endif()
 
 # -----------------------------------------------------------------------------
 # Declare Options
@@ -445,6 +454,9 @@ endif()
 
 option(WITH_PYTHON_INSTALL       "Copy system python into the blender install folder" ON)
 
+option(WITH_INSTALL_COPYRIGHT    "Copy the official Blender Foundation's copyright.txt into the Blender install folder" OFF)
+mark_as_advanced(WITH_INSTALL_COPYRIGHT)
+
 if((WITH_AUDASPACE AND NOT WITH_SYSTEM_AUDASPACE) OR WITH_MOD_FLUID)
   option(WITH_PYTHON_NUMPY "Include NumPy in Blender (used by Audaspace and Mantaflow)"  ON)
 endif()

build_files/utils/make_update.py

@@ -374,7 +374,7 @@ def external_script_add_origin_if_needed(args: argparse.Namespace,
         #  - Rename remote "upstream" to "origin", which takes care of changing the names of
         #    remotes the local branches are tracking.
         #
-        #  - Change the URL to the "origin", which so was was still pointing to upstream.
+        #  - Change the URL to the "origin", which was still pointing to upstream.
         #
         #  - Re-introduce the "upstream" remote, with the same URL as it had prior to rename.
 
@@ -445,10 +445,17 @@ def external_scripts_update(args: argparse.Namespace,
                     # automatically and fails when the branch is available in multiple remotes.
                     if make_utils.git_local_branch_exists(args.git_command, submodule_branch):
                         call([args.git_command, "checkout", submodule_branch])
-                    elif make_utils.git_remote_exist(args.git_command, "origin"):
+                    else:
-                        call([args.git_command, "checkout", "-t", f"origin/{submodule_branch}"])
+                        if make_utils.git_remote_branch_exists(args.git_command, "origin", submodule_branch):
-                    elif make_utils.git_remote_exist(args.git_command, "upstream"):
+                            call([args.git_command, "checkout", "-t", f"origin/{submodule_branch}"])
-                        call([args.git_command, "checkout", "-t", f"upstream/{submodule_branch}"])
+                        elif make_utils.git_remote_exist(args.git_command, "upstream"):
+                            # For the Github style of upstream workflow create a local branch from
+                            # the upstream, but do not track it, so that we stick to the paradigm
+                            # that no local branches are tracking upstream, preventing possible
+                            # accidental commit to upstream.
+                            call([args.git_command, "checkout", "-b", submodule_branch,
+                                 f"upstream/{submodule_branch}", "--no-track"])
+
                 # Don't use extra fetch since all remotes of interest have been already fetched
                 # some lines above.
                 skip_msg += work_tree_update(args, use_fetch=False)

build_files/utils/make_utils.py

@@ -60,11 +60,16 @@ def git_local_branch_exists(git_command: str, branch: str) -> bool:
     )
 
 
+def git_remote_branch_exists(git_command: str, remote: str, branch: str) -> bool:
+    return call([git_command, "rev-parse", "--verify", f"remotes/{remote}/{branch}"],
+                exit_on_error=False, silent=True) == 0
+
+
 def git_branch_exists(git_command: str, branch: str) -> bool:
     return (
         git_local_branch_exists(git_command, branch) or
-        call([git_command, "rev-parse", "--verify", "remotes/upstream/" + branch], exit_on_error=False, silent=True) == 0 or
+        git_remote_branch_exists(git_command, "upstream", branch) or
-        call([git_command, "rev-parse", "--verify", "remotes/origin/" + branch], exit_on_error=False, silent=True) == 0
+        git_remote_branch_exists(git_command, "origin", branch)
     )
 
 

doc/doxygen/doxygen.source.h

@@ -104,9 +104,9 @@
  *   merged in docs.
  */
 
-/**
+/** \defgroup gui GUI
- * \defgroup gui GUI
+ *  \ingroup blender
- * \ingroup blender */
+ */
 
 /** \defgroup wm Window Manager
  *  \ingroup gui */

doc/manpage/blender.1.py

@@ -5,7 +5,7 @@
 This script generates the blender.1 man page, embedding the help text
 from the Blender executable itself. Invoke it as follows:
 
-    blender.1.py --blender <path-to-blender> --output <output-filename>
+    ./blender.bin -b --python doc/manpage/blender.1.py -- --output <output-filename>
 
 where <path-to-blender> is the path to the Blender executable,
 and <output-filename> is where to write the generated man page.
@@ -13,8 +13,8 @@ and <output-filename> is where to write the generated man page.
 
 import argparse
 import os
-import subprocess
 import time
+import sys
 
 from typing import (
     Dict,
@@ -28,58 +28,28 @@ def man_format(data: str) -> str:
     return data
 
 
-def blender_extract_info(blender_bin: str) -> Dict[str, str]:
+def blender_extract_info() -> Dict[str, str]:
+    # Only use of `bpy` in this file.
+    import bpy  # type: ignore
+    blender_help_text = bpy.app.help_text()
+    blender_version_text = bpy.app.version_string
+    blender_build_date_text = bpy.app.build_date
 
-    blender_env = {
+    if blender_build_date_text == b'Unknown':
-        "ASAN_OPTIONS": "exitcode=0:" + os.environ.get("ASAN_OPTIONS", ""),
-    }
-
-    blender_help = subprocess.run(
-        [blender_bin, "--help"],
-        env=blender_env,
-        check=True,
-        stdout=subprocess.PIPE,
-    ).stdout.decode(encoding="utf-8")
-
-    blender_version_output = subprocess.run(
-        [blender_bin, "--version"],
-        env=blender_env,
-        check=True,
-        stdout=subprocess.PIPE,
-    ).stdout.decode(encoding="utf-8")
-
-    # Extract information from the version string.
-    # Note that some internal modules may print errors (e.g. color management),
-    # check for each lines prefix to ensure these aren't included.
-    blender_version = ""
-    blender_date = ""
-    for l in blender_version_output.split("\n"):
-        if l.startswith("Blender "):
-            # Remove 'Blender' prefix.
-            blender_version = l.split(" ", 1)[1].strip()
-        elif l.lstrip().startswith("build date:"):
-            # Remove 'build date:' prefix.
-            blender_date = l.split(":", 1)[1].strip()
-        if blender_version and blender_date:
-            break
-
-    if not blender_date:
         # Happens when built without WITH_BUILD_INFO e.g.
-        date_string = time.strftime("%B %d, %Y", time.gmtime(int(os.environ.get('SOURCE_DATE_EPOCH', time.time()))))
+        blender_date = time.strftime("%B %d, %Y", time.gmtime(int(os.environ.get('SOURCE_DATE_EPOCH', time.time()))))
     else:
-        date_string = time.strftime("%B %d, %Y", time.strptime(blender_date, "%Y-%m-%d"))
+        blender_date = time.strftime("%B %d, %Y", time.strptime(blender_build_date_text, "%Y-%m-%d"))
 
     return {
-        "help": blender_help,
+        "help": blender_help_text,
-        "version": blender_version,
+        "version": blender_version_text,
-        "date": date_string,
+        "date": blender_date,
     }
 
 
-def man_page_from_blender_help(fh: TextIO, blender_bin: str, verbose: bool) -> None:
+def man_page_from_blender_help(fh: TextIO, verbose: bool) -> None:
-    if verbose:
+    blender_info = blender_extract_info()
-        print("Extracting help text:", blender_bin)
-    blender_info = blender_extract_info(blender_bin)
 
     # Header Content.
     fh.write(
@@ -172,11 +142,6 @@ def create_argparse() -> argparse.ArgumentParser:
         required=True,
         help="The man page to write to."
     )
-    parser.add_argument(
-        "--blender",
-        required=True,
-        help="Path to the blender binary."
-    )
     parser.add_argument(
         "--verbose",
         default=False,
@@ -189,15 +154,15 @@ def create_argparse() -> argparse.ArgumentParser:
 
 
 def main() -> None:
+    argv = sys.argv[sys.argv.index("--") + 1:]
     parser = create_argparse()
-    args = parser.parse_args()
+    args = parser.parse_args(argv)
 
-    blender_bin = args.blender
     output_filename = args.output
     verbose = args.verbose
 
     with open(output_filename, "w", encoding="utf-8") as fh:
-        man_page_from_blender_help(fh, blender_bin, verbose)
+        man_page_from_blender_help(fh, verbose)
         if verbose:
             print("Written:", output_filename)
 

doc/python_api/examples/bpy.context.property.py

@@ -0,0 +1,10 @@
+"""
+Get the property associated with a hovered button.
+Returns a tuple of the datablock, data path to the property, and array index.
+"""
+
+# Example inserting keyframe for the hovered property.
+active_property = bpy.context.property
+if active_property:
+    datablock, data_path, index = active_property
+    datablock.keyframe_insert(data_path=data_path, index=index, frame=1)

doc/python_api/sphinx_doc_gen.py

@@ -1202,6 +1202,7 @@ context_type_map = {
     "particle_settings": ("ParticleSettings", False),
     "particle_system": ("ParticleSystem", False),
     "particle_system_editable": ("ParticleSystem", False),
+    "property": ("(:class:`bpy.types.ID`, :class:`string`, :class:`int`)", False),
     "pointcloud": ("PointCloud", False),
     "pose_bone": ("PoseBone", False),
     "pose_object": ("Object", False),
@@ -1347,7 +1348,11 @@ def pycontext2sphinx(basepath):
                 raise SystemExit(
                     "Error: context key %r not found in context_type_map; update %s" %
                     (member, __file__)) from None
-            fw("   :type: %s :class:`bpy.types.%s`\n\n" % ("sequence of " if is_seq else "", member_type))
+
+            if member_type.isidentifier():
+                member_type = ":class:`bpy.types.%s`" % member_type
+            fw("   :type: %s %s\n\n" % ("sequence of " if is_seq else "", member_type))
+            write_example_ref("   ", fw, "bpy.context." + member)
 
     # Generate type-map:
     # for member in sorted(unique_context_strings):

extern/CMakeLists.txt

@@ -3,6 +3,7 @@
 
 # Libs that adhere to strict flags
 add_subdirectory(curve_fit_nd)
+add_subdirectory(fmtlib)
 
 # Otherwise we get warnings here that we cant fix in external projects
 remove_strict_flags()

extern/audaspace/AUTHORS

@@ -24,3 +24,4 @@ Several people provided fixes:
 - Aaron Carlisle
 - Sebastian Parborg
 - Leon Zandman
+- Richard Antalik

extern/audaspace/bindings/python/PyDevice.cpp

@@ -124,7 +124,7 @@ Device_new(PyTypeObject* type, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Device_lock_doc,
-			 ".. classmethod:: lock()\n\n"
+			 ".. method:: lock()\n\n"
 			 "   Locks the device so that it's guaranteed, that no samples are\n"
 			 "   read from the streams until :meth:`unlock` is called.\n"
 			 "   This is useful if you want to do start/stop/pause/resume some\n"
@@ -152,7 +152,7 @@ Device_lock(Device* self)
 }
 
 PyDoc_STRVAR(M_aud_Device_play_doc,
-			 ".. classmethod:: play(sound, keep=False)\n\n"
+			 ".. method:: play(sound, keep=False)\n\n"
 			 "   Plays a sound.\n\n"
 			 "   :arg sound: The sound to play.\n"
 			 "   :type sound: :class:`Sound`\n"
@@ -212,7 +212,7 @@ Device_play(Device* self, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Device_stopAll_doc,
-			 ".. classmethod:: stopAll()\n\n"
+			 ".. method:: stopAll()\n\n"
 			 "   Stops all playing and paused sounds.");
 
 static PyObject *
@@ -231,7 +231,7 @@ Device_stopAll(Device* self)
 }
 
 PyDoc_STRVAR(M_aud_Device_unlock_doc,
-			 ".. classmethod:: unlock()\n\n"
+			 ".. method:: unlock()\n\n"
 			 "   Unlocks the device after a lock call, see :meth:`lock` for\n"
 			 "   details.");
 

extern/audaspace/bindings/python/PyDynamicMusic.cpp

@@ -60,7 +60,7 @@ DynamicMusic_dealloc(DynamicMusicP* self)
 }
 
 PyDoc_STRVAR(M_aud_DynamicMusic_addScene_doc,
-	".. classmethod:: addScene(scene)\n\n"
+	".. method:: addScene(scene)\n\n"
 	"   Adds a new scene.\n\n"
 	"   :arg scene: The scene sound.\n"
 	"   :type scene: :class:`Sound`\n"
@@ -90,7 +90,7 @@ DynamicMusic_addScene(DynamicMusicP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_DynamicMusic_addTransition_doc,
-	".. classmethod:: addTransition(ini, end, transition)\n\n"
+	".. method:: addTransition(ini, end, transition)\n\n"
 	"   Adds a new scene.\n\n"
 	"   :arg ini: the initial scene foor the transition.\n"
 	"   :type ini: int\n"
@@ -125,7 +125,7 @@ DynamicMusic_addTransition(DynamicMusicP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_DynamicMusic_resume_doc,
-	".. classmethod:: resume()\n\n"
+	".. method:: resume()\n\n"
 	"   Resumes playback of the scene.\n\n"
 	"   :return: Whether the action succeeded.\n"
 	"   :rtype: bool");
@@ -145,7 +145,7 @@ DynamicMusic_resume(DynamicMusicP* self)
 }
 
 PyDoc_STRVAR(M_aud_DynamicMusic_pause_doc,
-	".. classmethod:: pause()\n\n"
+	".. method:: pause()\n\n"
 	"   Pauses playback of the scene.\n\n"
 	"   :return: Whether the action succeeded.\n"
 	"   :rtype: bool");
@@ -165,7 +165,7 @@ DynamicMusic_pause(DynamicMusicP* self)
 }
 
 PyDoc_STRVAR(M_aud_DynamicMusic_stop_doc,
-	".. classmethod:: stop()\n\n"
+	".. method:: stop()\n\n"
 	"   Stops playback of the scene.\n\n"
 	"   :return: Whether the action succeeded.\n"
 	"   :rtype: bool\n\n");

extern/audaspace/bindings/python/PyHRTF.cpp

@@ -54,7 +54,7 @@ HRTF_dealloc(HRTFP* self)
 }
 
 PyDoc_STRVAR(M_aud_HRTF_addImpulseResponse_doc,
-	".. classmethod:: addImpulseResponseFromSound(sound, azimuth, elevation)\n\n"
+	".. method:: addImpulseResponseFromSound(sound, azimuth, elevation)\n\n"
 	"   Adds a new hrtf to the HRTF object\n\n"
 	"   :arg sound: The sound that contains the hrtf.\n"
 	"   :type sound: :class:`Sound`\n"
@@ -90,7 +90,7 @@ HRTF_addImpulseResponseFromSound(HRTFP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_HRTF_loadLeftHrtfSet_doc,
-	".. classmethod:: loadLeftHrtfSet(extension, directory)\n\n"
+	".. method:: loadLeftHrtfSet(extension, directory)\n\n"
 	"   Loads all HRTFs from a directory.\n\n"
 	"   :arg extension: The file extension of the hrtfs.\n"
 	"   :type extension: string\n"
@@ -125,7 +125,7 @@ HRTF_loadLeftHrtfSet(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_HRTF_loadRightHrtfSet_doc,
-	".. classmethod:: loadLeftHrtfSet(extension, directory)\n\n"
+	".. method:: loadLeftHrtfSet(extension, directory)\n\n"
 	"   Loads all HRTFs from a directory.\n\n"
 	"   :arg extension: The file extension of the hrtfs.\n"
 	"   :type extension: string\n"

extern/audaspace/bindings/python/PyHandle.cpp

@@ -38,7 +38,7 @@ Handle_dealloc(Handle* self)
 }
 
 PyDoc_STRVAR(M_aud_Handle_pause_doc,
-			 ".. classmethod:: pause()\n\n"
+			 ".. method:: pause()\n\n"
 			 "   Pauses playback.\n\n"
 			 "   :return: Whether the action succeeded.\n"
 			 "   :rtype: bool");
@@ -58,7 +58,7 @@ Handle_pause(Handle* self)
 }
 
 PyDoc_STRVAR(M_aud_Handle_resume_doc,
-			 ".. classmethod:: resume()\n\n"
+			 ".. method:: resume()\n\n"
 			 "   Resumes playback.\n\n"
 			 "   :return: Whether the action succeeded.\n"
 			 "   :rtype: bool");
@@ -78,7 +78,7 @@ Handle_resume(Handle* self)
 }
 
 PyDoc_STRVAR(M_aud_Handle_stop_doc,
-			 ".. classmethod:: stop()\n\n"
+			 ".. method:: stop()\n\n"
 			 "   Stops playback.\n\n"
 			 "   :return: Whether the action succeeded.\n"
 			 "   :rtype: bool\n\n"

extern/audaspace/bindings/python/PyPlaybackManager.cpp

@@ -60,7 +60,7 @@ PlaybackManager_dealloc(PlaybackManagerP* self)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_play_doc,
-	".. classmethod:: play(sound, catKey)\n\n"
+	".. method:: play(sound, catKey)\n\n"
 	"   Plays a sound through the playback manager and assigns it to a category.\n\n"
 	"   :arg sound: The sound to play.\n"
 	"   :type sound: :class:`Sound`\n"
@@ -104,7 +104,7 @@ PlaybackManager_play(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_resume_doc,
-	".. classmethod:: resume(catKey)\n\n"
+	".. method:: resume(catKey)\n\n"
 	"   Resumes playback of the catgory.\n\n"
 	"   :arg catKey: the key of the category.\n"
 	"   :type catKey: int\n"
@@ -131,7 +131,7 @@ PlaybackManager_resume(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_pause_doc,
-	".. classmethod:: pause(catKey)\n\n"
+	".. method:: pause(catKey)\n\n"
 	"   Pauses playback of the category.\n\n"
 	"   :arg catKey: the key of the category.\n"
 	"   :type catKey: int\n"
@@ -158,7 +158,7 @@ PlaybackManager_pause(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_add_category_doc,
-	".. classmethod:: addCategory(volume)\n\n"
+	".. method:: addCategory(volume)\n\n"
 	"   Adds a category with a custom volume.\n\n"
 	"   :arg volume: The volume for ther new category.\n"
 	"   :type volume: float\n"
@@ -185,7 +185,7 @@ PlaybackManager_add_category(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_get_volume_doc,
-	".. classmethod:: getVolume(catKey)\n\n"
+	".. method:: getVolume(catKey)\n\n"
 	"   Retrieves the volume of a category.\n\n"
 	"   :arg catKey: the key of the category.\n"
 	"   :type catKey: int\n"
@@ -212,7 +212,7 @@ PlaybackManager_get_volume(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_set_volume_doc,
-	".. classmethod:: setVolume(volume, catKey)\n\n"
+	".. method:: setVolume(volume, catKey)\n\n"
 	"   Changes the volume of a category.\n\n"
 	"   :arg volume: the new volume value.\n"
 	"   :type volume: float\n"
@@ -242,7 +242,7 @@ PlaybackManager_set_volume(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_stop_doc,
-	".. classmethod:: stop(catKey)\n\n"
+	".. method:: stop(catKey)\n\n"
 	"   Stops playback of the category.\n\n"
 	"   :arg catKey: the key of the category.\n"
 	"   :type catKey: int\n"
@@ -269,7 +269,7 @@ PlaybackManager_stop(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_clean_doc,
-	".. classmethod:: clean()\n\n"
+	".. method:: clean()\n\n"
 	"   Cleans all the invalid and finished sound from the playback manager.\n\n");
 
 static PyObject *

extern/audaspace/bindings/python/PySequence.cpp

@@ -99,7 +99,7 @@ Sequence_new(PyTypeObject* type, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Sequence_add_doc,
-			 ".. classmethod:: add()\n\n"
+			 ".. method:: add()\n\n"
 			 "   Adds a new entry to the sequence.\n\n"
 			 "   :arg sound: The sound this entry should play.\n"
 			 "   :type sound: :class:`Sound`\n"
@@ -151,7 +151,7 @@ Sequence_add(Sequence* self, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Sequence_remove_doc,
-			 ".. classmethod:: remove()\n\n"
+			 ".. method:: remove()\n\n"
 			 "   Removes an entry from the sequence.\n\n"
 			 "   :arg entry: The entry to remove.\n"
 			 "   :type entry: :class:`SequenceEntry`\n");
@@ -183,7 +183,7 @@ Sequence_remove(Sequence* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sequence_setAnimationData_doc,
-			 ".. classmethod:: setAnimationData()\n\n"
+			 ".. method:: setAnimationData()\n\n"
 			 "   Writes animation data to a sequence.\n\n"
 			 "   :arg type: The type of animation data.\n"
 			 "   :type type: int\n"

extern/audaspace/bindings/python/PySequenceEntry.cpp

@@ -43,7 +43,7 @@ SequenceEntry_dealloc(SequenceEntry* self)
 }
 
 PyDoc_STRVAR(M_aud_SequenceEntry_move_doc,
-			 ".. classmethod:: move()\n\n"
+			 ".. method:: move()\n\n"
 			 "   Moves the entry.\n\n"
 			 "   :arg begin: The new start time.\n"
 			 "   :type begin: double\n"
@@ -73,7 +73,7 @@ SequenceEntry_move(SequenceEntry* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_SequenceEntry_setAnimationData_doc,
-			 ".. classmethod:: setAnimationData()\n\n"
+			 ".. method:: setAnimationData()\n\n"
 			 "   Writes animation data to a sequenced entry.\n\n"
 			 "   :arg type: The type of animation data.\n"
 			 "   :type type: int\n"

extern/audaspace/bindings/python/PySound.cpp

@@ -115,7 +115,7 @@ Sound_new(PyTypeObject* type, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Sound_data_doc,
-			 ".. classmethod:: data()\n\n"
+			 ".. method:: data()\n\n"
 			 "   Retrieves the data of the sound as numpy array.\n\n"
 			 "   :return: A two dimensional numpy float array.\n"
 			 "   :rtype: :class:`numpy.ndarray`\n\n"
@@ -146,7 +146,7 @@ Sound_data(Sound* self)
 }
 
 PyDoc_STRVAR(M_aud_Sound_write_doc,
-			 ".. classmethod:: write(filename, rate, channels, format, container, codec, bitrate, buffersize)\n\n"
+			 ".. method:: write(filename, rate, channels, format, container, codec, bitrate, buffersize)\n\n"
 			 "   Writes the sound to a file.\n\n"
 			 "   :arg filename: The path to write to.\n"
 			 "   :type filename: string\n"
@@ -357,7 +357,7 @@ Sound_buffer(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_cache_doc,
-			 ".. classmethod:: cache()\n\n"
+			 ".. method:: cache()\n\n"
 			 "   Caches a sound into RAM.\n\n"
 			 "   This saves CPU usage needed for decoding and file access if the\n"
 			 "   underlying sound reads from a file on the harddisk,\n"
@@ -631,7 +631,7 @@ Sound_triangle(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_accumulate_doc,
-			 ".. classmethod:: accumulate(additive=False)\n\n"
+			 ".. method:: accumulate(additive=False)\n\n"
 			 "   Accumulates a sound by summing over positive input\n"
 			 "   differences thus generating a monotonic sigal.\n"
 			 "   If additivity is set to true negative input differences get added too,\n"
@@ -683,7 +683,7 @@ Sound_accumulate(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_ADSR_doc,
-			 ".. classmethod:: ADSR(attack, decay, sustain, release)\n\n"
+			 ".. method:: ADSR(attack, decay, sustain, release)\n\n"
 			 "   Attack-Decay-Sustain-Release envelopes the volume of a sound.\n"
 			 "   Note: there is currently no way to trigger the release with this API.\n\n"
 			 "   :arg attack: The attack time in seconds.\n"
@@ -726,7 +726,7 @@ Sound_ADSR(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_delay_doc,
-			 ".. classmethod:: delay(time)\n\n"
+			 ".. method:: delay(time)\n\n"
 			 "   Delays by playing adding silence in front of the other sound's data.\n\n"
 			 "   :arg time: How many seconds of silence should be added before the sound.\n"
 			 "   :type time: float\n"
@@ -762,7 +762,7 @@ Sound_delay(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_envelope_doc,
-			 ".. classmethod:: envelope(attack, release, threshold, arthreshold)\n\n"
+			 ".. method:: envelope(attack, release, threshold, arthreshold)\n\n"
 			 "   Delays by playing adding silence in front of the other sound's data.\n\n"
 			 "   :arg attack: The attack factor.\n"
 			 "   :type attack: float\n"
@@ -804,7 +804,7 @@ Sound_envelope(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_fadein_doc,
-			 ".. classmethod:: fadein(start, length)\n\n"
+			 ".. method:: fadein(start, length)\n\n"
 			 "   Fades a sound in by raising the volume linearly in the given\n"
 			 "   time interval.\n\n"
 			 "   :arg start: Time in seconds when the fading should start.\n"
@@ -844,7 +844,7 @@ Sound_fadein(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_fadeout_doc,
-			 ".. classmethod:: fadeout(start, length)\n\n"
+			 ".. method:: fadeout(start, length)\n\n"
 			 "   Fades a sound in by lowering the volume linearly in the given\n"
 			 "   time interval.\n\n"
 			 "   :arg start: Time in seconds when the fading should start.\n"
@@ -886,7 +886,7 @@ Sound_fadeout(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_filter_doc,
-			 ".. classmethod:: filter(b, a = (1))\n\n"
+			 ".. method:: filter(b, a = (1))\n\n"
 			 "   Filters a sound with the supplied IIR filter coefficients.\n"
 			 "   Without the second parameter you'll get a FIR filter.\n\n"
 			 "   If the first value of the a sequence is 0,\n"
@@ -986,7 +986,7 @@ Sound_filter(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_highpass_doc,
-			 ".. classmethod:: highpass(frequency, Q=0.5)\n\n"
+			 ".. method:: highpass(frequency, Q=0.5)\n\n"
 			 "   Creates a second order highpass filter based on the transfer\n"
 			 "   function :math:`H(s) = s^2 / (s^2 + s/Q + 1)`\n\n"
 			 "   :arg frequency: The cut off trequency of the highpass.\n"
@@ -1026,7 +1026,7 @@ Sound_highpass(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_limit_doc,
-			 ".. classmethod:: limit(start, end)\n\n"
+			 ".. method:: limit(start, end)\n\n"
 			 "   Limits a sound within a specific start and end time.\n\n"
 			 "   :arg start: Start time in seconds.\n"
 			 "   :type start: float\n"
@@ -1064,7 +1064,7 @@ Sound_limit(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_loop_doc,
-			 ".. classmethod:: loop(count)\n\n"
+			 ".. method:: loop(count)\n\n"
 			 "   Loops a sound.\n\n"
 			 "   :arg count: How often the sound should be looped.\n"
 			 "      Negative values mean endlessly.\n"
@@ -1104,7 +1104,7 @@ Sound_loop(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_lowpass_doc,
-			 ".. classmethod:: lowpass(frequency, Q=0.5)\n\n"
+			 ".. method:: lowpass(frequency, Q=0.5)\n\n"
 			 "   Creates a second order lowpass filter based on the transfer "
 			 "   function :math:`H(s) = 1 / (s^2 + s/Q + 1)`\n\n"
 			 "   :arg frequency: The cut off trequency of the lowpass.\n"
@@ -1144,7 +1144,7 @@ Sound_lowpass(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_modulate_doc,
-			 ".. classmethod:: modulate(sound)\n\n"
+			 ".. method:: modulate(sound)\n\n"
 			 "   Modulates two factories.\n\n"
 			 "   :arg sound: The sound to modulate over the other.\n"
 			 "   :type sound: :class:`Sound`\n"
@@ -1186,7 +1186,7 @@ Sound_modulate(Sound* self, PyObject* object)
 }
 
 PyDoc_STRVAR(M_aud_Sound_pitch_doc,
-			 ".. classmethod:: pitch(factor)\n\n"
+			 ".. method:: pitch(factor)\n\n"
 			 "   Changes the pitch of a sound with a specific factor.\n\n"
 			 "   :arg factor: The factor to change the pitch with.\n"
 			 "   :type factor: float\n"
@@ -1229,7 +1229,7 @@ Sound_pitch(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_rechannel_doc,
-			 ".. classmethod:: rechannel(channels)\n\n"
+			 ".. method:: rechannel(channels)\n\n"
 			 "   Rechannels the sound.\n\n"
 			 "   :arg channels: The new channel configuration.\n"
 			 "   :type channels: int\n"
@@ -1269,7 +1269,7 @@ Sound_rechannel(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_resample_doc,
-			 ".. classmethod:: resample(rate, high_quality)\n\n"
+			 ".. method:: resample(rate, high_quality)\n\n"
 			 "   Resamples the sound.\n\n"
 			 "   :arg rate: The new sample rate.\n"
 			 "   :type rate: double\n"
@@ -1324,7 +1324,7 @@ Sound_resample(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_reverse_doc,
-			 ".. classmethod:: reverse()\n\n"
+			 ".. method:: reverse()\n\n"
 			 "   Plays a sound reversed.\n\n"
 			 "   :return: The created :class:`Sound` object.\n"
 			 "   :rtype: :class:`Sound`\n\n"
@@ -1362,7 +1362,7 @@ Sound_reverse(Sound* self)
 }
 
 PyDoc_STRVAR(M_aud_Sound_sum_doc,
-			 ".. classmethod:: sum()\n\n"
+			 ".. method:: sum()\n\n"
 			 "   Sums the samples of a sound.\n\n"
 			 "   :return: The created :class:`Sound` object.\n"
 			 "   :rtype: :class:`Sound`");
@@ -1391,7 +1391,7 @@ Sound_sum(Sound* self)
 }
 
 PyDoc_STRVAR(M_aud_Sound_threshold_doc,
-			 ".. classmethod:: threshold(threshold = 0)\n\n"
+			 ".. method:: threshold(threshold = 0)\n\n"
 			 "   Makes a threshold wave out of an audio wave by setting all samples\n"
 			 "   with a amplitude >= threshold to 1, all <= -threshold to -1 and\n"
 			 "   all between to 0.\n\n"
@@ -1430,7 +1430,7 @@ Sound_threshold(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_volume_doc,
-			 ".. classmethod:: volume(volume)\n\n"
+			 ".. method:: volume(volume)\n\n"
 			 "   Changes the volume of a sound.\n\n"
 			 "   :arg volume: The new volume..\n"
 			 "   :type volume: float\n"
@@ -1471,7 +1471,7 @@ Sound_volume(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_join_doc,
-			 ".. classmethod:: join(sound)\n\n"
+			 ".. method:: join(sound)\n\n"
 			 "   Plays two factories in sequence.\n\n"
 			 "   :arg sound: The sound to play second.\n"
 			 "   :type sound: :class:`Sound`\n"
@@ -1514,7 +1514,7 @@ Sound_join(Sound* self, PyObject* object)
 }
 
 PyDoc_STRVAR(M_aud_Sound_mix_doc,
-			 ".. classmethod:: mix(sound)\n\n"
+			 ".. method:: mix(sound)\n\n"
 			 "   Mixes two factories.\n\n"
 			 "   :arg sound: The sound to mix over the other.\n"
 			 "   :type sound: :class:`Sound`\n"
@@ -1556,7 +1556,7 @@ Sound_mix(Sound* self, PyObject* object)
 }
 
 PyDoc_STRVAR(M_aud_Sound_pingpong_doc,
-			 ".. classmethod:: pingpong()\n\n"
+			 ".. method:: pingpong()\n\n"
 			 "   Plays a sound forward and then backward.\n"
 			 "   This is like joining a sound with its reverse.\n\n"
 			 "   :return: The created :class:`Sound` object.\n"
@@ -1622,7 +1622,7 @@ Sound_list(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_mutable_doc,
-	".. classmethod:: mutable()\n\n"
+	".. method:: mutable()\n\n"
 	"   Creates a sound that will be restarted when sought backwards.\n"
 	"   If the original sound is a sound list, the playing sound can change.\n\n"
 	"   :return: The created :class:`Sound` object.\n"
@@ -1652,7 +1652,7 @@ Sound_mutable(Sound* self)
 }
 
 PyDoc_STRVAR(M_aud_Sound_list_addSound_doc,
-	".. classmethod:: addSound(sound)\n\n"
+	".. method:: addSound(sound)\n\n"
 	"   Adds a new sound to a sound list.\n\n"
 	"   :arg sound: The sound that will be added to the list.\n"
 	"   :type sound: :class:`Sound`\n\n"
@@ -1685,7 +1685,7 @@ Sound_list_addSound(Sound* self, PyObject* object)
 #ifdef WITH_CONVOLUTION
 
 PyDoc_STRVAR(M_aud_Sound_convolver_doc,
-	".. classmethod:: convolver()\n\n"
+	".. method:: convolver()\n\n"
 	"   Creates a sound that will apply convolution to another sound.\n\n"
 	"   :arg impulseResponse: The filter with which convolve the sound.\n"
 	"   :type impulseResponse: :class:`ImpulseResponse`\n"
@@ -1734,7 +1734,7 @@ Sound_convolver(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_binaural_doc,
-	".. classmethod:: convolver()\n\n"
+	".. method:: binaural()\n\n"
 	"   Creates a binaural sound using another sound as source. The original sound must be mono\n\n"
 	"   :arg hrtfs: An HRTF set.\n"
 	"   :type hrtf: :class:`HRTF`\n"

extern/fast_float/README.blender

@@ -1,7 +1,7 @@
 Project: fast_float
 URL: https://github.com/fastfloat/fast_float
 License: MIT
-Upstream version: 4.0.0 (fbd5bd7, 2023 Mar 31)
+Upstream version: 5.0.0 (f5a3e77, 2023 May 25)
 Local modifications:
 
 - Took only the fast_float.h header and the license/readme files

extern/fast_float/README.md

@@ -1,4 +1,8 @@
+
 ## fast_float number parsing library: 4x faster than strtod
+[![Fuzzing Status](https://oss-fuzz-build-logs.storage.googleapis.com/badges/fast_float.svg)](https://bugs.chromium.org/p/oss-fuzz/issues/list?sort=-opened&can=1&q=proj:fast_float)
+[![VS17-CI](https://github.com/fastfloat/fast_float/actions/workflows/vs17-ci.yml/badge.svg)](https://github.com/fastfloat/fast_float/actions/workflows/vs17-ci.yml)
+[![Ubuntu 22.04 CI (GCC 11)](https://github.com/fastfloat/fast_float/actions/workflows/ubuntu22.yml/badge.svg)](https://github.com/fastfloat/fast_float/actions/workflows/ubuntu22.yml)
 
 The fast_float library provides fast header-only implementations for the C++ from_chars
 functions for `float` and `double` types.  These functions convert ASCII strings representing
@@ -93,6 +97,24 @@ constexpr double constexptest() {
 }
 ```
 
+## Non-ASCII Inputs
+
+We also support UTF-16 and UTF-32 inputs, as well as ASCII/UTF-8, as in the following example:
+
+``` C++
+#include "fast_float/fast_float.h"
+#include <iostream>
+
+int main() {
+    const std::u16string input =  u"3.1416 xyz ";
+    double result;
+    auto answer = fast_float::from_chars(input.data(), input.data()+input.size(), result);
+    if(answer.ec != std::errc()) { std::cerr << "parsing failure\n"; return EXIT_FAILURE; }
+    std::cout << "parsed the number " << result << std::endl;
+    return EXIT_SUCCESS;
+}
+```
+
 ## Using commas as decimal separator
 
 
@@ -189,11 +211,11 @@ It can parse random floating-point numbers at a speed of 1 GB/s on some systems.
 $ ./build/benchmarks/benchmark
 # parsing random integers in the range [0,1)
 volume = 2.09808 MB
-netlib                                  :   271.18 MB/s (+/- 1.2 %)    12.93 Mfloat/s 
+netlib                                  :   271.18 MB/s (+/- 1.2 %)    12.93 Mfloat/s
-doubleconversion                        :   225.35 MB/s (+/- 1.2 %)    10.74 Mfloat/s 
+doubleconversion                        :   225.35 MB/s (+/- 1.2 %)    10.74 Mfloat/s
-strtod                                  :   190.94 MB/s (+/- 1.6 %)     9.10 Mfloat/s 
+strtod                                  :   190.94 MB/s (+/- 1.6 %)     9.10 Mfloat/s
-abseil                                  :   430.45 MB/s (+/- 2.2 %)    20.52 Mfloat/s 
+abseil                                  :   430.45 MB/s (+/- 2.2 %)    20.52 Mfloat/s
-fastfloat                               :  1042.38 MB/s (+/- 9.9 %)    49.68 Mfloat/s 
+fastfloat                               :  1042.38 MB/s (+/- 9.9 %)    49.68 Mfloat/s
 ```
 
 See https://github.com/lemire/simple_fastfloat_benchmark for our benchmarking code.
@@ -257,7 +279,7 @@ under the Apache 2.0 license.
 
 <sup>
 Licensed under either of <a href="LICENSE-APACHE">Apache License, Version
-2.0</a> or <a href="LICENSE-MIT">MIT license</a> at your option.
+2.0</a> or <a href="LICENSE-MIT">MIT license</a> or <a href="LICENSE-BOOST">BOOST license</a> .
 </sup>
 
 <br>
@@ -265,5 +287,5 @@ Licensed under either of <a href="LICENSE-APACHE">Apache License, Version
 <sub>
 Unless you explicitly state otherwise, any contribution intentionally submitted
 for inclusion in this repository by you, as defined in the Apache-2.0 license,
-shall be dual licensed as above, without any additional terms or conditions.
+shall be triple licensed as above, without any additional terms or conditions.
 </sub>

extern/fmtlib/CMakeLists.txt

@@ -0,0 +1,20 @@
+# SPDX-License-Identifier: GPL-2.0-or-later
+
+set(INC
+  include
+)
+
+set(INC_SYS
+)
+
+set(SRC
+  include/fmt/core.h
+  include/fmt/format-inl.h
+  include/fmt/format.h
+  src/format.cc
+)
+
+set(LIB
+)
+
+blender_add_lib(extern_fmtlib "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")

extern/fmtlib/LICENSE.rst

@@ -1,4 +1,4 @@
-Copyright (c) 2012 - present, Victor Zverovich
+Copyright (c) 2012 - present, Victor Zverovich and {fmt} contributors
 
 Permission is hereby granted, free of charge, to any person obtaining
 a copy of this software and associated documentation files (the

extern/fmtlib/README.blender

@@ -1,8 +1,12 @@
 Project: {fmt}
 URL: https://github.com/fmtlib/fmt
 License: MIT
-Upstream version: 8.1.1 (b6f4cea)
+Upstream version: 10.0.0 (a0b8a92, 2023 May 10)
 Local modifications:
 
-- Took only files needed for Blender: LICENSE, README and include/fmt
+- Took only files needed for Blender:
-  folder's core.h, format-inl.h, format.h
+  - LICENSE, README
+  - include/fmt: core.h, format-inl.h, format.h
+  - src/format.cc
+- CMakeLists.txt is not from fmtlib, but
+  made for Blender codebase

extern/fmtlib/README.rst

@@ -1,5 +1,7 @@
-{fmt}
+.. image:: https://user-images.githubusercontent.com/
-=====
+           576385/156254208-f5b743a9-88cf-439d-b0c0-923d53e8d551.png
+   :width: 25%
+   :alt: {fmt}
 
 .. image:: https://github.com/fmtlib/fmt/workflows/linux/badge.svg
    :target: https://github.com/fmtlib/fmt/actions?query=workflow%3Alinux
@@ -10,9 +12,6 @@
 .. image:: https://github.com/fmtlib/fmt/workflows/windows/badge.svg
    :target: https://github.com/fmtlib/fmt/actions?query=workflow%3Awindows
 
-.. image:: https://ci.appveyor.com/api/projects/status/ehjkiefde6gucy1v?svg=true
-   :target: https://ci.appveyor.com/project/vitaut/fmt
-
 .. image:: https://oss-fuzz-build-logs.storage.googleapis.com/badges/fmt.svg
    :alt: fmt is continuously fuzzed at oss-fuzz
    :target: https://bugs.chromium.org/p/oss-fuzz/issues/list?\
@@ -26,12 +25,13 @@
 **{fmt}** is an open-source formatting library providing a fast and safe
 alternative to C stdio and C++ iostreams.
 
-If you like this project, please consider donating to the BYSOL
+If you like this project, please consider donating to one of the funds that
-Foundation that helps victims of political repressions in Belarus:
+help victims of the war in Ukraine: https://www.stopputin.net/.
-https://bysol.org/en/bs/general/.
 
 `Documentation <https://fmt.dev>`__
 
+`Cheat Sheets <https://hackingcpp.com/cpp/libs/fmt.html>`__
+
 Q&A: ask questions on `StackOverflow with the tag fmt
 <https://stackoverflow.com/questions/tagged/fmt>`_.
 
@@ -47,7 +47,8 @@ Features
 * `Format string syntax <https://fmt.dev/latest/syntax.html>`_ similar to Python's
   `format <https://docs.python.org/3/library/stdtypes.html#str.format>`_
 * Fast IEEE 754 floating-point formatter with correct rounding, shortness and
-  round-trip guarantees
+  round-trip guarantees using the `Dragonbox <https://github.com/jk-jeon/dragonbox>`_
+  algorithm
 * Safe `printf implementation
   <https://fmt.dev/latest/api.html#printf-formatting>`_ including the POSIX
   extension for positional arguments
@@ -123,7 +124,7 @@ Output::
     Default format: 42s 100ms
     strftime-like format: 03:15:30
 
-**Print a container** (`run <https://godbolt.org/z/MjsY7c>`_)
+**Print a container** (`run <https://godbolt.org/z/MxM1YqjE7>`_)
 
 .. code:: c++
 
@@ -191,24 +192,24 @@ Speed tests
 ================= ============= ===========
 Library           Method        Run Time, s
 ================= ============= ===========
-libc              printf          1.04
+libc              printf          0.91
-libc++            std::ostream    3.05
+libc++            std::ostream    2.49
-{fmt} 6.1.1       fmt::print      0.75
+{fmt} 9.1         fmt::print      0.74
-Boost Format 1.67 boost::format   7.24
+Boost Format 1.80 boost::format   6.26
-Folly Format      folly::format   2.23
+Folly Format      folly::format   1.87
 ================= ============= ===========
 
-{fmt} is the fastest of the benchmarked methods, ~35% faster than ``printf``.
+{fmt} is the fastest of the benchmarked methods, ~20% faster than ``printf``.
 
 The above results were generated by building ``tinyformat_test.cpp`` on macOS
-10.14.6 with ``clang++ -O3 -DNDEBUG -DSPEED_TEST -DHAVE_FORMAT``, and taking the
+12.6.1 with ``clang++ -O3 -DNDEBUG -DSPEED_TEST -DHAVE_FORMAT``, and taking the
 best of three runs. In the test, the format string ``"%0.10f:%04d:%+g:%s:%p:%c:%%\n"``
 or equivalent is filled 2,000,000 times with output sent to ``/dev/null``; for
 further details refer to the `source
 <https://github.com/fmtlib/format-benchmark/blob/master/src/tinyformat-test.cc>`_.
 
 {fmt} is up to 20-30x faster than ``std::ostringstream`` and ``sprintf`` on
-floating-point formatting (`dtoa-benchmark <https://github.com/fmtlib/dtoa-benchmark>`_)
+IEEE754 ``float`` and ``double`` formatting (`dtoa-benchmark <https://github.com/fmtlib/dtoa-benchmark>`_)
 and faster than `double-conversion <https://github.com/google/double-conversion>`_ and
 `ryu <https://github.com/ulfjack/ryu>`_:
 
@@ -322,8 +323,10 @@ Projects using this library
 
 * `ccache <https://ccache.dev/>`_: a compiler cache
 
-* `ClickHouse <https://github.com/ClickHouse/ClickHouse>`_: analytical database
+* `ClickHouse <https://github.com/ClickHouse/ClickHouse>`_: an analytical database
   management system
+  
+* `Contour <https://github.com/contour-terminal/contour/>`_: a modern terminal emulator
 
 * `CUAUV <https://cuauv.org/>`_: Cornell University's autonomous underwater
   vehicle
@@ -341,9 +344,12 @@ Projects using this library
 
 * `Folly <https://github.com/facebook/folly>`_: Facebook open-source library
 
+* `GemRB <https://gemrb.org/>`_: a portable open-source implementation of
+  Bioware’s Infinity Engine
+
 * `Grand Mountain Adventure
   <https://store.steampowered.com/app/1247360/Grand_Mountain_Adventure/>`_:
-  A beautiful open-world ski & snowboarding game
+  a beautiful open-world ski & snowboarding game
 
 * `HarpyWar/pvpgn <https://github.com/pvpgn/pvpgn-server>`_:
   Player vs Player Gaming Network with tweaks
@@ -357,6 +363,10 @@ Projects using this library
 
 * `Knuth <https://kth.cash/>`_: high-performance Bitcoin full-node
 
+* `libunicode <https://github.com/contour-terminal/libunicode/>`_: a modern C++17 Unicode library
+
+* `MariaDB <https://mariadb.org/>`_: relational database management system
+
 * `Microsoft Verona <https://github.com/microsoft/verona>`_:
   research programming language for concurrent ownership
 
@@ -410,6 +420,9 @@ Projects using this library
 * `TrinityCore <https://github.com/TrinityCore/TrinityCore>`_: open-source
   MMORPG framework
 
+* `🐙 userver framework <https://userver.tech/>`_: open-source asynchronous
+  framework with a rich set of abstractions and database drivers
+
 * `Windows Terminal <https://github.com/microsoft/terminal>`_: the new Windows
   terminal
 
@@ -520,8 +533,7 @@ Maintainers
 -----------
 
 The {fmt} library is maintained by Victor Zverovich (`vitaut
-<https://github.com/vitaut>`_) and Jonathan Müller (`foonathan
+<https://github.com/vitaut>`_) with contributions from many other people.
-<https://github.com/foonathan>`_) with contributions from many other people.
 See `Contributors <https://github.com/fmtlib/fmt/graphs/contributors>`_ and
 `Releases <https://github.com/fmtlib/fmt/releases>`_ for some of the names.
 Let us know if your contribution is not listed or mentioned incorrectly and

extern/fmtlib/src/format.cc

@@ -0,0 +1,43 @@
+// Formatting library for C++
+//
+// Copyright (c) 2012 - 2016, Victor Zverovich
+// All rights reserved.
+//
+// For the license information refer to format.h.
+
+#include "fmt/format-inl.h"
+
+FMT_BEGIN_NAMESPACE
+namespace detail {
+
+template FMT_API auto dragonbox::to_decimal(float x) noexcept
+    -> dragonbox::decimal_fp<float>;
+template FMT_API auto dragonbox::to_decimal(double x) noexcept
+    -> dragonbox::decimal_fp<double>;
+
+#ifndef FMT_STATIC_THOUSANDS_SEPARATOR
+template FMT_API locale_ref::locale_ref(const std::locale& loc);
+template FMT_API auto locale_ref::get<std::locale>() const -> std::locale;
+#endif
+
+// Explicit instantiations for char.
+
+template FMT_API auto thousands_sep_impl(locale_ref)
+    -> thousands_sep_result<char>;
+template FMT_API auto decimal_point_impl(locale_ref) -> char;
+
+template FMT_API void buffer<char>::append(const char*, const char*);
+
+template FMT_API void vformat_to(buffer<char>&, string_view,
+                                 typename vformat_args<>::type, locale_ref);
+
+// Explicit instantiations for wchar_t.
+
+template FMT_API auto thousands_sep_impl(locale_ref)
+    -> thousands_sep_result<wchar_t>;
+template FMT_API auto decimal_point_impl(locale_ref) -> wchar_t;
+
+template FMT_API void buffer<wchar_t>::append(const wchar_t*, const wchar_t*);
+
+}  // namespace detail
+FMT_END_NAMESPACE

extern/vulkan_memory_allocator/vk_mem_alloc_impl.cc

@@ -9,4 +9,13 @@
 
 #define VMA_IMPLEMENTATION
 
+/* 
+ * Disabling internal asserts of VMA. 
+ * 
+ * Blender can destroy logical device before all the resources are freed. This is because static
+ * resources are freed as a last step during quiting. As long as Vulkan isn't feature complete
+ * we don't want to change this behavior. So for now we just disable the asserts.
+ */
+#define VMA_ASSERT(test)
+
 #include "vk_mem_alloc.h"

intern/cycles/blender/CMakeLists.txt

@@ -26,6 +26,7 @@ set(SRC
   image.cpp
   geometry.cpp
   light.cpp
+  light_linking.cpp
   mesh.cpp
   object.cpp
   object_cull.cpp
@@ -47,6 +48,7 @@ set(SRC
   display_driver.h
   id_map.h
   image.h
+  light_linking.h
   object_cull.h
   output_driver.h
   sync.h

intern/cycles/blender/addon/ui.py

@@ -7,7 +7,7 @@ from bpy.app.translations import contexts as i18n_contexts
 from bpy_extras.node_utils import find_node_input
 from bl_ui.utils import PresetPanel
 
-from bpy.types import Panel
+from bpy.types import Panel, Menu
 
 from bl_ui.properties_grease_pencil_common import GreasePencilSimplifyPanel
 from bl_ui.properties_render import draw_curves_settings
@@ -1301,6 +1301,7 @@ class CYCLES_OBJECT_PT_lightgroup(CyclesButtonsPanel, Panel):
     bl_label = "Light Group"
     bl_parent_id = "CYCLES_OBJECT_PT_shading"
     bl_context = "object"
+    bl_options = {'DEFAULT_CLOSED'}
 
     def draw(self, context):
         layout = self.layout
@@ -1321,6 +1322,88 @@ class CYCLES_OBJECT_PT_lightgroup(CyclesButtonsPanel, Panel):
         sub.operator("scene.view_layer_add_lightgroup", icon='ADD', text="").name = ob.lightgroup
 
 
+class CYCLES_OBJECT_MT_light_linking_context_menu(Menu):
+    bl_label = "Light Linking Specials"
+
+    def draw(self, _context):
+        layout = self.layout
+
+        layout.operator("object.light_linking_receivers_select")
+
+
+class CYCLES_OBJECT_MT_shadow_linking_context_menu(Menu):
+    bl_label = "Shadow Linking Specials"
+
+    def draw(self, _context):
+        layout = self.layout
+
+        layout.operator("object.light_linking_blockers_select")
+
+
+class CYCLES_OBJECT_PT_light_linking(CyclesButtonsPanel, Panel):
+    bl_label = "Light Linking"
+    bl_parent_id = "CYCLES_OBJECT_PT_shading"
+    bl_context = "object"
+    bl_options = {'DEFAULT_CLOSED'}
+
+    def draw(self, context):
+        layout = self.layout
+        layout.use_property_split = True
+
+        object = context.object
+        light_linking = object.light_linking
+
+        col = layout.column()
+
+        col.template_ID(
+            light_linking,
+            "receiver_collection",
+            new="object.light_linking_receiver_collection_new")
+
+        if not light_linking.receiver_collection:
+            return
+
+        row = layout.row()
+        col = row.column()
+        col.template_light_linking_collection(light_linking, "receiver_collection")
+
+        col = row.column()
+        sub = col.column(align=True)
+        sub.menu("CYCLES_OBJECT_MT_light_linking_context_menu", icon='DOWNARROW_HLT', text="")
+
+
+class CYCLES_OBJECT_PT_shadow_linking(CyclesButtonsPanel, Panel):
+    bl_label = "Shadow Linking"
+    bl_parent_id = "CYCLES_OBJECT_PT_shading"
+    bl_context = "object"
+    bl_options = {'DEFAULT_CLOSED'}
+
+    def draw(self, context):
+        layout = self.layout
+        layout.use_property_split = True
+
+        object = context.object
+        light_linking = object.light_linking
+
+        col = layout.column()
+
+        col.template_ID(
+            light_linking,
+            "blocker_collection",
+            new="object.light_linking_blocker_collection_new")
+
+        if not light_linking.blocker_collection:
+            return
+
+        row = layout.row()
+        col = row.column()
+        col.template_light_linking_collection(light_linking, "blocker_collection")
+
+        col = row.column()
+        sub = col.column(align=True)
+        sub.menu("CYCLES_OBJECT_MT_shadow_linking_context_menu", icon='DOWNARROW_HLT', text="")
+
+
 class CYCLES_OBJECT_PT_visibility(CyclesButtonsPanel, Panel):
     bl_label = "Visibility"
     bl_context = "object"
@@ -2363,6 +2446,14 @@ def draw_pause(self, context):
             layout.prop(cscene, "preview_pause", icon='PLAY' if cscene.preview_pause else 'PAUSE', text="")
 
 
+def draw_make_links(self, context):
+    if context.engine == "CYCLES":
+        layout = self.layout
+        layout.separator()
+        layout.operator_menu_enum("object.light_linking_receivers_link", "link_state")
+        layout.operator_menu_enum("object.light_linking_blockers_link", "link_state")
+
+
 def get_panels():
     exclude_panels = {
         'DATA_PT_camera_dof',
@@ -2451,6 +2542,10 @@ classes = (
     CYCLES_OBJECT_PT_shading_gi_approximation,
     CYCLES_OBJECT_PT_shading_caustics,
     CYCLES_OBJECT_PT_lightgroup,
+    CYCLES_OBJECT_MT_light_linking_context_menu,
+    CYCLES_OBJECT_PT_light_linking,
+    CYCLES_OBJECT_MT_shadow_linking_context_menu,
+    CYCLES_OBJECT_PT_shadow_linking,
     CYCLES_OBJECT_PT_visibility,
     CYCLES_OBJECT_PT_visibility_ray_visibility,
     CYCLES_OBJECT_PT_visibility_culling,
@@ -2497,6 +2592,7 @@ def register():
 
     bpy.types.RENDER_PT_context.append(draw_device)
     bpy.types.VIEW3D_HT_header.append(draw_pause)
+    bpy.types.VIEW3D_MT_make_links.append(draw_make_links)
 
     for panel in get_panels():
         panel.COMPAT_ENGINES.add('CYCLES')
@@ -2510,6 +2606,7 @@ def unregister():
 
     bpy.types.RENDER_PT_context.remove(draw_device)
     bpy.types.VIEW3D_HT_header.remove(draw_pause)
+    bpy.types.VIEW3D_MT_make_links.remove(draw_make_links)
 
     for panel in get_panels():
         if 'CYCLES' in panel.COMPAT_ENGINES:

intern/cycles/blender/light.cpp

@@ -3,6 +3,7 @@
 
 #include "scene/light.h"
 
+#include "blender/light_linking.h"
 #include "blender/sync.h"
 #include "blender/util.h"
 
@@ -37,6 +38,8 @@ void BlenderSync::sync_light(BL::Object &b_parent,
     }
   }
 
+  light->name = b_light.name().c_str();
+
   /* type */
   switch (b_light.type()) {
     case BL::Light::type_POINT: {
@@ -145,8 +148,12 @@ void BlenderSync::sync_light(BL::Object &b_parent,
   light->set_use_scatter((visibility & PATH_RAY_VOLUME_SCATTER) != 0);
   light->set_is_shadow_catcher(b_ob_info.real_object.is_shadow_catcher());
 
-  /* lightgroup */
+  /* Light group and linking. */
   light->set_lightgroup(ustring(b_ob_info.real_object.lightgroup()));
+  light->set_light_set_membership(
+      BlenderLightLink::get_light_set_membership(PointerRNA_NULL, b_ob_info.real_object));
+  light->set_shadow_set_membership(
+      BlenderLightLink::get_shadow_set_membership(PointerRNA_NULL, b_ob_info.real_object));
 
   /* tag */
   light->tag_update(scene);

intern/cycles/blender/light_linking.cpp

@@ -0,0 +1,63 @@
+/* SPDX-License-Identifier: Apache-2.0
+ * Copyright 2011-2023 Blender Foundation */
+
+#include "blender/light_linking.h"
+
+#include "scene/object.h"
+
+#include "DNA_object_types.h"
+
+CCL_NAMESPACE_BEGIN
+
+static const ::Object *get_blender_object(const BL::Object &object)
+{
+  return reinterpret_cast<::Object *>(object.ptr.data);
+}
+
+static const ::LightLinking *get_light_linking(const BL::Object &object)
+{
+  const ::Object *blender_object = get_blender_object(object);
+  return blender_object->light_linking;
+}
+
+uint64_t BlenderLightLink::get_light_set_membership(const BL::Object & /*parent*/,
+                                                    const BL::Object &object)
+{
+  const ::LightLinking *light_linking = get_light_linking(object);
+  return (light_linking) ? light_linking->runtime.light_set_membership : LIGHT_LINK_MASK_ALL;
+}
+
+uint BlenderLightLink::get_receiver_light_set(const BL::Object &parent, const BL::Object &object)
+{
+  if (parent) {
+    const ::LightLinking *parent_light_linking = get_light_linking(parent);
+    if (parent_light_linking && parent_light_linking->runtime.receiver_light_set) {
+      return parent_light_linking->runtime.receiver_light_set;
+    }
+  }
+
+  const ::LightLinking *light_linking = get_light_linking(object);
+  return (light_linking) ? light_linking->runtime.receiver_light_set : 0;
+}
+
+uint64_t BlenderLightLink::get_shadow_set_membership(const BL::Object & /*parent*/,
+                                                     const BL::Object &object)
+{
+  const ::LightLinking *light_linking = get_light_linking(object);
+  return (light_linking) ? light_linking->runtime.shadow_set_membership : LIGHT_LINK_MASK_ALL;
+}
+
+uint BlenderLightLink::get_blocker_shadow_set(const BL::Object &parent, const BL::Object &object)
+{
+  if (parent) {
+    const ::LightLinking *parent_light_linking = get_light_linking(parent);
+    if (parent_light_linking && parent_light_linking->runtime.blocker_shadow_set) {
+      return parent_light_linking->runtime.blocker_shadow_set;
+    }
+  }
+
+  const ::LightLinking *light_linking = get_light_linking(object);
+  return (light_linking) ? light_linking->runtime.blocker_shadow_set : 0;
+}
+
+CCL_NAMESPACE_END

intern/cycles/blender/light_linking.h

@@ -0,0 +1,22 @@
+/* SPDX-License-Identifier: Apache-2.0
+ * Copyright 2011-2023 Blender Foundation */
+
+#pragma once
+
+#include <cstdint>
+
+#include "MEM_guardedalloc.h"
+#include "RNA_blender_cpp.h"
+
+CCL_NAMESPACE_BEGIN
+
+class BlenderLightLink {
+ public:
+  static uint64_t get_light_set_membership(const BL::Object &parent, const BL::Object &object);
+  static uint get_receiver_light_set(const BL::Object &parent, const BL::Object &object);
+
+  static uint64_t get_shadow_set_membership(const BL::Object &parent, const BL::Object &object);
+  static uint get_blocker_shadow_set(const BL::Object &parent, const BL::Object &object);
+};
+
+CCL_NAMESPACE_END

intern/cycles/blender/object.cpp

@@ -1,6 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
+#include "blender/light_linking.h"
 #include "blender/object_cull.h"
 #include "blender/sync.h"
 #include "blender/util.h"
@@ -348,9 +349,14 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph,
       object->set_random_id(hash_uint2(hash_string(object->name.c_str()), 0));
     }
 
-    /* lightgroup */
+    /* Light group and linking. */
     object->set_lightgroup(ustring(b_ob.lightgroup()));
 
+    object->set_light_set_membership(BlenderLightLink::get_light_set_membership(b_parent, b_ob));
+    object->set_receiver_light_set(BlenderLightLink::get_receiver_light_set(b_parent, b_ob));
+    object->set_shadow_set_membership(BlenderLightLink::get_shadow_set_membership(b_parent, b_ob));
+    object->set_blocker_shadow_set(BlenderLightLink::get_blocker_shadow_set(b_parent, b_ob));
+
     object->tag_update(scene);
   }
 

intern/cycles/blender/util.h

@@ -118,7 +118,7 @@ static inline BL::Mesh object_to_mesh(BL::BlendData & /*data*/,
   if ((bool)mesh && subdivision_type == Mesh::SUBDIVISION_NONE) {
     if (mesh.use_auto_smooth()) {
       mesh.calc_normals_split();
-      mesh.split_faces(false);
+      mesh.split_faces();
     }
 
     mesh.calc_loop_triangles();

intern/cycles/device/cpu/kernel.cpp

@@ -24,11 +24,13 @@ CPUKernels::CPUKernels()
       REGISTER_KERNEL(integrator_intersect_shadow),
       REGISTER_KERNEL(integrator_intersect_subsurface),
       REGISTER_KERNEL(integrator_intersect_volume_stack),
+      REGISTER_KERNEL(integrator_intersect_dedicated_light),
       REGISTER_KERNEL(integrator_shade_background),
       REGISTER_KERNEL(integrator_shade_light),
       REGISTER_KERNEL(integrator_shade_shadow),
       REGISTER_KERNEL(integrator_shade_surface),
       REGISTER_KERNEL(integrator_shade_volume),
+      REGISTER_KERNEL(integrator_shade_dedicated_light),
       REGISTER_KERNEL(integrator_megakernel),
       /* Shader evaluation. */
       REGISTER_KERNEL(shader_eval_displace),

intern/cycles/device/cpu/kernel.h

@@ -33,11 +33,13 @@ class CPUKernels {
   IntegratorFunction integrator_intersect_shadow;
   IntegratorFunction integrator_intersect_subsurface;
   IntegratorFunction integrator_intersect_volume_stack;
+  IntegratorFunction integrator_intersect_dedicated_light;
   IntegratorShadeFunction integrator_shade_background;
   IntegratorShadeFunction integrator_shade_light;
   IntegratorShadeFunction integrator_shade_shadow;
   IntegratorShadeFunction integrator_shade_surface;
   IntegratorShadeFunction integrator_shade_volume;
+  IntegratorShadeFunction integrator_shade_dedicated_light;
   IntegratorShadeFunction integrator_megakernel;
 
   /* Shader evaluation. */

intern/cycles/device/kernel.cpp

@@ -18,6 +18,7 @@ bool device_kernel_has_shading(DeviceKernel kernel)
           kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE ||
           kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME ||
           kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW ||
+          kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT ||
           kernel == DEVICE_KERNEL_SHADER_EVAL_DISPLACE ||
           kernel == DEVICE_KERNEL_SHADER_EVAL_BACKGROUND ||
           kernel == DEVICE_KERNEL_SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY);
@@ -29,6 +30,7 @@ bool device_kernel_has_intersection(DeviceKernel kernel)
           kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW ||
           kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE ||
           kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK ||
+          kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT ||
           kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE ||
           kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE);
 }
@@ -49,6 +51,8 @@ const char *device_kernel_as_string(DeviceKernel kernel)
       return "integrator_intersect_subsurface";
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
       return "integrator_intersect_volume_stack";
+    case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
+      return "integrator_intersect_dedicated_light";
     case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND:
       return "integrator_shade_background";
     case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
@@ -63,6 +67,8 @@ const char *device_kernel_as_string(DeviceKernel kernel)
       return "integrator_shade_surface_mnee";
     case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME:
       return "integrator_shade_volume";
+    case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT:
+      return "integrator_shade_dedicated_light";
     case DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL:
       return "integrator_megakernel";
     case DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY:

intern/cycles/device/metal/device_impl.mm

@@ -345,7 +345,9 @@ string MetalDevice::preprocess_source(MetalPipelineType pso_type,
     case METAL_GPU_APPLE:
       global_defines += "#define __KERNEL_METAL_APPLE__\n";
 #  ifdef WITH_NANOVDB
-      global_defines += "#define WITH_NANOVDB\n";
+      if (DebugFlags().metal.use_nanovdb) {
+        global_defines += "#define WITH_NANOVDB\n";
+      }
 #  endif
       break;
   }

intern/cycles/device/metal/kernel.mm

@@ -32,20 +32,14 @@ const char *kernel_type_as_string(MetalPipelineType pso_type)
   return "";
 }
 
-bool kernel_has_intersection(DeviceKernel device_kernel)
-{
-  return (device_kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST ||
-          device_kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW ||
-          device_kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE ||
-          device_kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK ||
-          device_kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE ||
-          device_kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE);
-}
-
 struct ShaderCache {
   ShaderCache(id<MTLDevice> _mtlDevice) : mtlDevice(_mtlDevice)
   {
     /* Initialize occupancy tuning LUT. */
+
+    // TODO: Look into tuning for DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT and
+    // DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT.
+
     if (MetalInfo::get_device_vendor(mtlDevice) == METAL_GPU_APPLE) {
       switch (MetalInfo::get_apple_gpu_architecture(mtlDevice)) {
         default:
@@ -393,6 +387,11 @@ bool MetalKernelPipeline::should_use_binary_archive() const
       }
     }
 
+    if (use_metalrt && device_kernel_has_intersection(device_kernel)) {
+      /* Binary linked functions aren't supported in binary archives. */
+      return false;
+    }
+
     if (pso_type == PSO_GENERIC) {
       /* Archive the generic kernels. */
       return true;
@@ -581,7 +580,7 @@ void MetalKernelPipeline::compile()
     [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL]];
     [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL_PRIM]];
 
-    if (kernel_has_intersection(device_kernel)) {
+    if (device_kernel_has_intersection(device_kernel)) {
       linked_functions = [[NSArray arrayWithArray:[unique_functions allObjects]]
           sortedArrayUsingComparator:^NSComparisonResult(id<MTLFunction> f1, id<MTLFunction> f2) {
             return [f1.label compare:f2.label];
@@ -682,7 +681,7 @@ void MetalKernelPipeline::compile()
     __block bool compilation_finished = false;
     __block string error_str;
 
-    if (loading_existing_archive) {
+    if (loading_existing_archive || !DebugFlags().metal.use_async_pso_creation) {
       /* Use the blocking variant of newComputePipelineStateWithDescriptor if an archive exists on
        * disk. It should load almost instantaneously, and will fail gracefully when loading a
        * corrupt archive (unlike the async variant). */
@@ -695,29 +694,6 @@ void MetalKernelPipeline::compile()
       error_str = err ? err : "nil";
     }
     else {
-      /* TODO / MetalRT workaround:
-       * Workaround for a crash when addComputePipelineFunctionsWithDescriptor is called *after*
-       * newComputePipelineStateWithDescriptor with linked functions (i.e. with MetalRT enabled).
-       * Ideally we would like to call newComputePipelineStateWithDescriptor (async) first so we
-       * can bail out if needed, but we can stop the crash by flipping the order when there are
-       * linked functions. However when addComputePipelineFunctionsWithDescriptor is called first
-       * it will block while it builds the pipeline, offering no way of bailing out. */
-      auto addComputePipelineFunctionsWithDescriptor = [&]() {
-        if (creating_new_archive && ShaderCache::running) {
-          NSError *error;
-          if (![archive addComputePipelineFunctionsWithDescriptor:computePipelineStateDescriptor
-                                                            error:&error])
-          {
-            NSString *errStr = [error localizedDescription];
-            metal_printf("Failed to add PSO to archive:\n%s\n",
-                         errStr ? [errStr UTF8String] : "nil");
-          }
-        }
-      };
-      if (linked_functions) {
-        addComputePipelineFunctionsWithDescriptor();
-      }
-
       /* Use the async variant of newComputePipelineStateWithDescriptor if no archive exists on
        * disk. This allows us to respond to app shutdown. */
       [mtlDevice
@@ -745,10 +721,16 @@ void MetalKernelPipeline::compile()
       while (ShaderCache::running && !compilation_finished) {
         std::this_thread::sleep_for(std::chrono::milliseconds(5));
       }
+    }
 
-      /* Add pipeline into the new archive (unless we did it earlier). */
+    if (creating_new_archive && pipeline) {
-      if (pipeline && !linked_functions) {
+      /* Add pipeline into the new archive. */
-        addComputePipelineFunctionsWithDescriptor();
+      NSError *error;
+      if (![archive addComputePipelineFunctionsWithDescriptor:computePipelineStateDescriptor
+                                                        error:&error])
+      {
+        NSString *errStr = [error localizedDescription];
+        metal_printf("Failed to add PSO to archive:\n%s\n", errStr ? [errStr UTF8String] : "nil");
       }
     }
 

intern/cycles/device/metal/queue.mm

@@ -531,6 +531,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
         case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
         case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
         case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
+        case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
         case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE:
         case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:
           break;

intern/cycles/device/optix/device_impl.cpp

@@ -421,6 +421,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
   group_descs[PG_RGEN_INTERSECT_VOLUME_STACK].raygen.module = optix_module;
   group_descs[PG_RGEN_INTERSECT_VOLUME_STACK].raygen.entryFunctionName =
       "__raygen__kernel_optix_integrator_intersect_volume_stack";
+  group_descs[PG_RGEN_INTERSECT_DEDICATED_LIGHT].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
+  group_descs[PG_RGEN_INTERSECT_DEDICATED_LIGHT].raygen.module = optix_module;
+  group_descs[PG_RGEN_INTERSECT_DEDICATED_LIGHT].raygen.entryFunctionName =
+      "__raygen__kernel_optix_integrator_intersect_dedicated_light";
   group_descs[PG_MISS].kind = OPTIX_PROGRAM_GROUP_KIND_MISS;
   group_descs[PG_MISS].miss.module = optix_module;
   group_descs[PG_MISS].miss.entryFunctionName = "__miss__kernel_optix_miss";
@@ -547,6 +551,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
     group_descs[PG_RGEN_SHADE_SHADOW].raygen.module = optix_module;
     group_descs[PG_RGEN_SHADE_SHADOW].raygen.entryFunctionName =
         "__raygen__kernel_optix_integrator_shade_shadow";
+    group_descs[PG_RGEN_SHADE_DEDICATED_LIGHT].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
+    group_descs[PG_RGEN_SHADE_DEDICATED_LIGHT].raygen.module = optix_module;
+    group_descs[PG_RGEN_SHADE_DEDICATED_LIGHT].raygen.entryFunctionName =
+        "__raygen__kernel_optix_integrator_shade_dedicated_light";
     group_descs[PG_RGEN_EVAL_DISPLACE].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
     group_descs[PG_RGEN_EVAL_DISPLACE].raygen.module = optix_module;
     group_descs[PG_RGEN_EVAL_DISPLACE].raygen.entryFunctionName =
@@ -659,6 +667,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
     pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_SHADOW]);
     pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_SUBSURFACE]);
     pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_VOLUME_STACK]);
+    pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_DEDICATED_LIGHT]);
     pipeline_groups.push_back(groups[PG_MISS]);
     pipeline_groups.push_back(groups[PG_HITD]);
     pipeline_groups.push_back(groups[PG_HITS]);
@@ -948,6 +957,7 @@ bool OptiXDevice::load_osl_kernels()
     pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_MNEE]);
     pipeline_groups.push_back(groups[PG_RGEN_SHADE_VOLUME]);
     pipeline_groups.push_back(groups[PG_RGEN_SHADE_SHADOW]);
+    pipeline_groups.push_back(groups[PG_RGEN_SHADE_DEDICATED_LIGHT]);
     pipeline_groups.push_back(groups[PG_RGEN_EVAL_DISPLACE]);
     pipeline_groups.push_back(groups[PG_RGEN_EVAL_BACKGROUND]);
     pipeline_groups.push_back(groups[PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY]);

intern/cycles/device/optix/device_impl.h

@@ -21,6 +21,7 @@ enum {
   PG_RGEN_INTERSECT_SHADOW,
   PG_RGEN_INTERSECT_SUBSURFACE,
   PG_RGEN_INTERSECT_VOLUME_STACK,
+  PG_RGEN_INTERSECT_DEDICATED_LIGHT,
   PG_RGEN_SHADE_BACKGROUND,
   PG_RGEN_SHADE_LIGHT,
   PG_RGEN_SHADE_SURFACE,
@@ -28,6 +29,7 @@ enum {
   PG_RGEN_SHADE_SURFACE_MNEE,
   PG_RGEN_SHADE_VOLUME,
   PG_RGEN_SHADE_SHADOW,
+  PG_RGEN_SHADE_DEDICATED_LIGHT,
   PG_RGEN_EVAL_DISPLACE,
   PG_RGEN_EVAL_BACKGROUND,
   PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY,

intern/cycles/device/optix/queue.cpp

@@ -123,6 +123,10 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel,
       pipeline = optix_device->pipelines[PIP_SHADE];
       sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_SHADOW * sizeof(SbtRecord);
       break;
+    case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT:
+      pipeline = optix_device->pipelines[PIP_SHADE];
+      sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_DEDICATED_LIGHT * sizeof(SbtRecord);
+      break;
 
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST:
       pipeline = optix_device->pipelines[PIP_INTERSECT];
@@ -140,6 +144,11 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel,
       pipeline = optix_device->pipelines[PIP_INTERSECT];
       sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_INTERSECT_VOLUME_STACK * sizeof(SbtRecord);
       break;
+    case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
+      pipeline = optix_device->pipelines[PIP_INTERSECT];
+      sbt_params.raygenRecord = sbt_data_ptr +
+                                PG_RGEN_INTERSECT_DEDICATED_LIGHT * sizeof(SbtRecord);
+      break;
 
     case DEVICE_KERNEL_SHADER_EVAL_DISPLACE:
       pipeline = optix_device->pipelines[PIP_SHADE];

intern/cycles/graph/node.cpp

@@ -68,6 +68,12 @@ void Node::set(const SocketType &input, uint value)
   set_if_different(input, value);
 }
 
+void Node::set(const SocketType &input, uint64_t value)
+{
+  assert(input.type == SocketType::UINT64);
+  set_if_different(input, value);
+}
+
 void Node::set(const SocketType &input, float value)
 {
   assert(input.type == SocketType::FLOAT);
@@ -190,6 +196,12 @@ uint Node::get_uint(const SocketType &input) const
   return get_socket_value<uint>(this, input);
 }
 
+uint64_t Node::get_uint64(const SocketType &input) const
+{
+  assert(input.type == SocketType::UINT64);
+  return get_socket_value<uint64_t>(this, input);
+}
+
 float Node::get_float(const SocketType &input) const
 {
   assert(input.type == SocketType::FLOAT);
@@ -434,6 +446,9 @@ void Node::set_value(const SocketType &socket, const Node &other, const SocketTy
       case SocketType::UINT:
         set(socket, get_socket_value<uint>(&other, socket));
         break;
+      case SocketType::UINT64:
+        set(socket, get_socket_value<uint64_t>(&other, socket));
+        break;
       case SocketType::COLOR:
       case SocketType::VECTOR:
       case SocketType::POINT:
@@ -489,6 +504,8 @@ bool Node::equals_value(const Node &other, const SocketType &socket) const
       return is_value_equal<int>(this, &other, socket);
     case SocketType::UINT:
       return is_value_equal<uint>(this, &other, socket);
+    case SocketType::UINT64:
+      return is_value_equal<uint64_t>(this, &other, socket);
     case SocketType::COLOR:
       return is_value_equal<float3>(this, &other, socket);
     case SocketType::VECTOR:
@@ -534,6 +551,7 @@ bool Node::equals_value(const Node &other, const SocketType &socket) const
       return is_array_equal<void *>(this, &other, socket);
 
     case SocketType::UNDEFINED:
+    case SocketType::NUM_TYPES:
       return true;
   }
 
@@ -608,6 +626,9 @@ void Node::hash(MD5Hash &md5)
       case SocketType::UINT:
         value_hash<uint>(this, socket, md5);
         break;
+      case SocketType::UINT64:
+        value_hash<uint64_t>(this, socket, md5);
+        break;
       case SocketType::COLOR:
         float3_hash(this, socket, md5);
         break;
@@ -673,6 +694,7 @@ void Node::hash(MD5Hash &md5)
         break;
 
       case SocketType::UNDEFINED:
+      case SocketType::NUM_TYPES:
         break;
     }
   }
@@ -697,6 +719,7 @@ size_t Node::get_total_size_in_bytes() const
       case SocketType::FLOAT:
       case SocketType::INT:
       case SocketType::UINT:
+      case SocketType::UINT64:
       case SocketType::COLOR:
       case SocketType::VECTOR:
       case SocketType::POINT:
@@ -745,6 +768,7 @@ size_t Node::get_total_size_in_bytes() const
         break;
 
       case SocketType::UNDEFINED:
+      case SocketType::NUM_TYPES:
         break;
     }
   }

intern/cycles/graph/node.h

@@ -95,6 +95,7 @@ struct Node {
   void set(const SocketType &input, bool value);
   void set(const SocketType &input, int value);
   void set(const SocketType &input, uint value);
+  void set(const SocketType &input, uint64_t value);
   void set(const SocketType &input, float value);
   void set(const SocketType &input, float2 value);
   void set(const SocketType &input, float3 value);
@@ -127,6 +128,7 @@ struct Node {
   bool get_bool(const SocketType &input) const;
   int get_int(const SocketType &input) const;
   uint get_uint(const SocketType &input) const;
+  uint64_t get_uint64(const SocketType &input) const;
   float get_float(const SocketType &input) const;
   float2 get_float2(const SocketType &input) const;
   float3 get_float3(const SocketType &input) const;

intern/cycles/graph/node_type.cpp

@@ -23,6 +23,7 @@ size_t SocketType::size(Type type)
 {
   switch (type) {
     case UNDEFINED:
+    case NUM_TYPES:
       return 0;
 
     case BOOLEAN:
@@ -33,6 +34,8 @@ size_t SocketType::size(Type type)
       return sizeof(int);
     case UINT:
       return sizeof(uint);
+    case UINT64:
+      return sizeof(uint64_t);
     case COLOR:
       return sizeof(float3);
     case VECTOR:
@@ -99,11 +102,12 @@ ustring SocketType::type_name(Type type)
 
                             ustring("boolean"),       ustring("float"),
                             ustring("int"),           ustring("uint"),
-                            ustring("color"),         ustring("vector"),
+                            ustring("uint64"),        ustring("color"),
-                            ustring("point"),         ustring("normal"),
+                            ustring("vector"),        ustring("point"),
-                            ustring("point2"),        ustring("closure"),
+                            ustring("normal"),        ustring("point2"),
-                            ustring("string"),        ustring("enum"),
+                            ustring("closure"),       ustring("string"),
-                            ustring("transform"),     ustring("node"),
+                            ustring("enum"),          ustring("transform"),
+                            ustring("node"),
 
                             ustring("array_boolean"), ustring("array_float"),
                             ustring("array_int"),     ustring("array_color"),
@@ -112,6 +116,9 @@ ustring SocketType::type_name(Type type)
                             ustring("array_string"),  ustring("array_transform"),
                             ustring("array_node")};
 
+  constexpr size_t num_names = sizeof(names) / sizeof(*names);
+  static_assert(num_names == NUM_TYPES);
+
   return names[(int)type];
 }
 

intern/cycles/graph/node_type.h

@@ -21,12 +21,13 @@ typedef uint64_t SocketModifiedFlags;
 
 struct SocketType {
   enum Type {
-    UNDEFINED,
+    UNDEFINED = 0,
 
     BOOLEAN,
     FLOAT,
     INT,
     UINT,
+    UINT64,
     COLOR,
     VECTOR,
     POINT,
@@ -49,6 +50,8 @@ struct SocketType {
     STRING_ARRAY,
     TRANSFORM_ARRAY,
     NODE_ARRAY,
+
+    NUM_TYPES,
   };
 
   enum Flags {
@@ -191,6 +194,8 @@ struct NodeType {
   SOCKET_DEFINE(name, ui_name, default_value, int, SocketType::INT, 0, ##__VA_ARGS__)
 #define SOCKET_UINT(name, ui_name, default_value, ...) \
   SOCKET_DEFINE(name, ui_name, default_value, uint, SocketType::UINT, 0, ##__VA_ARGS__)
+#define SOCKET_UINT64(name, ui_name, default_value, ...) \
+  SOCKET_DEFINE(name, ui_name, default_value, uint64_t, SocketType::UINT64, 0, ##__VA_ARGS__)
 #define SOCKET_FLOAT(name, ui_name, default_value, ...) \
   SOCKET_DEFINE(name, ui_name, default_value, float, SocketType::FLOAT, 0, ##__VA_ARGS__)
 #define SOCKET_COLOR(name, ui_name, default_value, ...) \

intern/cycles/graph/node_xml.cpp

@@ -93,6 +93,10 @@ void xml_read_node(XMLReader &reader, Node *node, xml_node xml_node)
         node->set(socket, (uint)atoi(attr.value()));
         break;
       }
+      case SocketType::UINT64: {
+        node->set(socket, (uint64_t)strtoull(attr.value(), nullptr, 10));
+        break;
+      }
       case SocketType::INT_ARRAY: {
         vector<string> tokens;
         string_split(tokens, attr.value());
@@ -213,6 +217,7 @@ void xml_read_node(XMLReader &reader, Node *node, xml_node xml_node)
       }
       case SocketType::CLOSURE:
       case SocketType::UNDEFINED:
+      case SocketType::NUM_TYPES:
         break;
     }
   }
@@ -280,6 +285,10 @@ xml_node xml_write_node(Node *node, xml_node xml_root)
         attr = node->get_uint(socket);
         break;
       }
+      case SocketType::UINT64: {
+        attr = node->get_uint64(socket);
+        break;
+      }
       case SocketType::INT_ARRAY: {
         std::stringstream ss;
         const array<int> &value = node->get_int_array(socket);
@@ -409,6 +418,7 @@ xml_node xml_write_node(Node *node, xml_node xml_root)
       }
       case SocketType::CLOSURE:
       case SocketType::UNDEFINED:
+      case SocketType::NUM_TYPES:
         break;
     }
   }

intern/cycles/hydra/field.cpp

@@ -26,12 +26,20 @@ class HdCyclesVolumeLoader : public VDBImageLoader {
   HdCyclesVolumeLoader(const std::string &filePath, const std::string &gridName)
       : VDBImageLoader(gridName)
   {
-    /* Disable delay loading and file copying, this has poor performance on network drivers. */
+    /* Disable delay loading and file copying, this has poor performance on network drives. */
     const bool delay_load = false;
-    openvdb::io::File file(filePath);
+    try {
-    file.setCopyMaxBytes(0);
+      openvdb::io::File file(filePath);
-    if (file.open(delay_load)) {
+      file.setCopyMaxBytes(0);
-      grid = file.readGrid(gridName);
+      if (file.open(delay_load)) {
+        grid = file.readGrid(gridName);
+      }
+    }
+    catch (const openvdb::IoError &e) {
+      VLOG_WARNING << "Error loading OpenVDB file: " << e.what();
+    }
+    catch (...) {
+      VLOG_WARNING << "Error loading OpenVDB file: Unknown error";
     }
   }
 };

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -508,7 +508,8 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
 
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
-    case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK: {
+    case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
+    case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT: {
       /* Ray intersection kernels with integrator state. */
       DeviceKernelArguments args(&d_path_index, &work_size);
 
@@ -521,7 +522,8 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
     case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE:
     case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE:
     case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:
-    case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME: {
+    case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME:
+    case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT: {
       /* Shading kernels with integrator state and render buffer. */
       DeviceKernelArguments args(&d_path_index, &buffers_->buffer.device_pointer, &work_size);
 
@@ -1177,7 +1179,8 @@ bool PathTraceWorkGPU::kernel_creates_shadow_paths(DeviceKernel kernel)
   return (kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE ||
           kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE ||
           kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE ||
-          kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME);
+          kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME ||
+          kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT);
 }
 
 bool PathTraceWorkGPU::kernel_creates_ao_paths(DeviceKernel kernel)

intern/cycles/kernel/CMakeLists.txt

@@ -267,6 +267,7 @@ set(SRC_KERNEL_INTEGRATOR_HEADERS
   integrator/displacement_shader.h
   integrator/init_from_bake.h
   integrator/init_from_camera.h
+  integrator/intersect_dedicated_light.h
   integrator/intersect_closest.h
   integrator/intersect_shadow.h
   integrator/intersect_subsurface.h
@@ -280,7 +281,9 @@ set(SRC_KERNEL_INTEGRATOR_HEADERS
   integrator/shade_shadow.h
   integrator/shade_surface.h
   integrator/shade_volume.h
+  integrator/shade_dedicated_light.h
   integrator/shadow_catcher.h
+  integrator/shadow_linking.h
   integrator/shadow_state_template.h
   integrator/state_flow.h
   integrator/state.h

intern/cycles/kernel/bvh/shadow_all.h

@@ -144,6 +144,12 @@ ccl_device_inline
               continue;
             }
 
+#ifdef __SHADOW_LINKING__
+            if (intersection_skip_shadow_link(kg, ray, prim_object)) {
+              continue;
+            }
+#endif
+
             switch (type & PRIMITIVE_ALL) {
               case PRIMITIVE_TRIANGLE: {
                 hit = triangle_intersect(

intern/cycles/kernel/bvh/traversal.h

@@ -129,6 +129,12 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
               continue;
             }
 
+#ifdef __SHADOW_LINKING__
+            if (intersection_skip_shadow_link(kg, ray, prim_object)) {
+              continue;
+            }
+#endif
+
             switch (type & PRIMITIVE_ALL) {
               case PRIMITIVE_TRIANGLE: {
                 if (triangle_intersect(kg,

intern/cycles/kernel/bvh/util.h

@@ -253,4 +253,41 @@ ccl_device_inline bool intersection_skip_self_local(ccl_private const RaySelfPri
   return (self.prim == prim);
 }
 
+#ifdef __SHADOW_LINKING__
+ccl_device_inline uint64_t ray_get_shadow_set_membership(KernelGlobals kg,
+                                                         ccl_private const Ray *ray)
+{
+  if (ray->self.light != LAMP_NONE) {
+    return kernel_data_fetch(lights, ray->self.light).shadow_set_membership;
+  }
+
+  if (ray->self.light_object != OBJECT_NONE) {
+    return kernel_data_fetch(objects, ray->self.light_object).shadow_set_membership;
+  }
+
+  return LIGHT_LINK_MASK_ALL;
+}
+#endif
+
+ccl_device_inline bool intersection_skip_shadow_link(KernelGlobals kg,
+                                                     ccl_private const Ray *ray,
+                                                     const int isect_object)
+{
+#ifdef __SHADOW_LINKING__
+  if (!(kernel_data.kernel_features & KERNEL_FEATURE_SHADOW_LINKING)) {
+    return false;
+  }
+
+  const uint64_t set_membership = ray_get_shadow_set_membership(kg, ray);
+  if (set_membership == LIGHT_LINK_MASK_ALL) {
+    return false;
+  }
+
+  const uint blocker_set = kernel_data_fetch(objects, isect_object).blocker_shadow_set;
+  return ((uint64_t(1) << uint64_t(blocker_set)) & set_membership) == 0;
+#else
+  return false;
+#endif
+}
+
 CCL_NAMESPACE_END

intern/cycles/kernel/camera/camera.h

@@ -12,19 +12,19 @@ CCL_NAMESPACE_BEGIN
 
 /* Perspective Camera */
 
-ccl_device float2 camera_sample_aperture(ccl_constant KernelCamera *cam, float u, float v)
+ccl_device float2 camera_sample_aperture(ccl_constant KernelCamera *cam, const float2 rand)
 {
   float blades = cam->blades;
   float2 bokeh;
 
   if (blades == 0.0f) {
     /* sample disk */
-    bokeh = concentric_sample_disk(u, v);
+    bokeh = concentric_sample_disk(rand);
   }
   else {
     /* sample polygon */
     float rotation = cam->bladesrotation;
-    bokeh = regular_polygon_sample(blades, rotation, u, v);
+    bokeh = regular_polygon_sample(blades, rotation, rand);
   }
 
   /* anamorphic lens bokeh */
@@ -34,15 +34,13 @@ ccl_device float2 camera_sample_aperture(ccl_constant KernelCamera *cam, float u
 }
 
 ccl_device void camera_sample_perspective(KernelGlobals kg,
-                                          float raster_x,
+                                          const float2 raster_xy,
-                                          float raster_y,
+                                          const float2 rand_lens,
-                                          float lens_u,
-                                          float lens_v,
                                           ccl_private Ray *ray)
 {
   /* create ray form raster position */
   ProjectionTransform rastertocamera = kernel_data.cam.rastertocamera;
-  float3 raster = make_float3(raster_x, raster_y, 0.0f);
+  const float3 raster = float2_to_float3(raster_xy);
   float3 Pcamera = transform_perspective(&rastertocamera, raster);
 
   if (kernel_data.cam.have_perspective_motion) {
@@ -71,14 +69,14 @@ ccl_device void camera_sample_perspective(KernelGlobals kg,
 
   if (aperturesize > 0.0f) {
     /* sample point on aperture */
-    float2 lensuv = camera_sample_aperture(&kernel_data.cam, lens_u, lens_v) * aperturesize;
+    float2 lens_uv = camera_sample_aperture(&kernel_data.cam, rand_lens) * aperturesize;
 
     /* compute point on plane of focus */
     float ft = kernel_data.cam.focaldistance / D.z;
     float3 Pfocus = D * ft;
 
     /* update ray for effect of lens */
-    P = make_float3(lensuv.x, lensuv.y, 0.0f);
+    P = float2_to_float3(lens_uv);
     D = normalize(Pfocus - P);
   }
 
@@ -133,7 +131,7 @@ ccl_device void camera_sample_perspective(KernelGlobals kg,
 
     float3 Px = Pnostereo;
     float3 Dx = transform_perspective(&rastertocamera,
-                                      make_float3(raster_x + 1.0f, raster_y, 0.0f));
+                                      make_float3(raster.x + 1.0f, raster.y, 0.0f));
     Dx = normalize(transform_direction(&cameratoworld, Dx));
     spherical_stereo_transform(&kernel_data.cam, &Px, &Dx);
 
@@ -144,7 +142,7 @@ ccl_device void camera_sample_perspective(KernelGlobals kg,
 
     float3 Py = Pnostereo;
     float3 Dy = transform_perspective(&rastertocamera,
-                                      make_float3(raster_x, raster_y + 1.0f, 0.0f));
+                                      make_float3(raster.x, raster.y + 1.0f, 0.0f));
     Dy = normalize(transform_direction(&cameratoworld, Dy));
     spherical_stereo_transform(&kernel_data.cam, &Py, &Dy);
 
@@ -166,15 +164,13 @@ ccl_device void camera_sample_perspective(KernelGlobals kg,
 
 /* Orthographic Camera */
 ccl_device void camera_sample_orthographic(KernelGlobals kg,
-                                           float raster_x,
+                                           const float2 raster_xy,
-                                           float raster_y,
+                                           const float2 rand_lens,
-                                           float lens_u,
-                                           float lens_v,
                                            ccl_private Ray *ray)
 {
   /* create ray form raster position */
   ProjectionTransform rastertocamera = kernel_data.cam.rastertocamera;
-  float3 Pcamera = transform_perspective(&rastertocamera, make_float3(raster_x, raster_y, 0.0f));
+  float3 Pcamera = transform_perspective(&rastertocamera, float2_to_float3(raster_xy));
 
   float3 P;
   float3 D = make_float3(0.0f, 0.0f, 1.0f);
@@ -184,15 +180,15 @@ ccl_device void camera_sample_orthographic(KernelGlobals kg,
 
   if (aperturesize > 0.0f) {
     /* sample point on aperture */
-    float2 lensuv = camera_sample_aperture(&kernel_data.cam, lens_u, lens_v) * aperturesize;
+    float2 lens_uv = camera_sample_aperture(&kernel_data.cam, rand_lens) * aperturesize;
 
     /* compute point on plane of focus */
     float3 Pfocus = D * kernel_data.cam.focaldistance;
 
     /* update ray for effect of lens */
-    float3 lensuvw = make_float3(lensuv.x, lensuv.y, 0.0f);
+    float3 lens_uvw = float2_to_float3(lens_uv);
-    P = Pcamera + lensuvw;
+    P = Pcamera + lens_uvw;
-    D = normalize(Pfocus - lensuvw);
+    D = normalize(Pfocus - lens_uvw);
   }
   else {
     P = Pcamera;
@@ -229,14 +225,12 @@ ccl_device void camera_sample_orthographic(KernelGlobals kg,
 
 ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
                                               ccl_global const DecomposedTransform *cam_motion,
-                                              float raster_x,
+                                              const float2 raster,
-                                              float raster_y,
+                                              const float2 rand_lens,
-                                              float lens_u,
-                                              float lens_v,
                                               ccl_private Ray *ray)
 {
   ProjectionTransform rastertocamera = cam->rastertocamera;
-  float3 Pcamera = transform_perspective(&rastertocamera, make_float3(raster_x, raster_y, 0.0f));
+  float3 Pcamera = transform_perspective(&rastertocamera, float2_to_float3(raster));
 
   /* create ray form raster position */
   float3 P = zero_float3();
@@ -253,7 +247,7 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
 
   if (aperturesize > 0.0f) {
     /* sample point on aperture */
-    float2 lensuv = camera_sample_aperture(cam, lens_u, lens_v) * aperturesize;
+    float2 lens_uv = camera_sample_aperture(cam, rand_lens) * aperturesize;
 
     /* compute point on plane of focus */
     float3 Dfocus = normalize(D);
@@ -265,7 +259,7 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
     V = normalize(cross(Dfocus, U));
 
     /* update ray for effect of lens */
-    P = U * lensuv.x + V * lensuv.y;
+    P = U * lens_uv.x + V * lens_uv.y;
     D = normalize(Pfocus - P);
   }
 
@@ -302,7 +296,7 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
   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 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);
@@ -314,7 +308,7 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
   dP.dx = Px - Pcenter;
   dD.dx = Dx - Dcenter;
 
-  float3 Py = transform_perspective(&rastertocamera, make_float3(raster_x, raster_y + 1.0f, 0.0f));
+  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);
@@ -341,17 +335,16 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
 ccl_device_inline void camera_sample(KernelGlobals kg,
                                      int x,
                                      int y,
-                                     float filter_u,
+                                     const float2 filter_uv,
-                                     float filter_v,
+                                     const float time,
-                                     float lens_u,
+                                     const float2 lens_uv,
-                                     float lens_v,
-                                     float time,
                                      ccl_private Ray *ray)
 {
   /* pixel filter */
   int filter_table_offset = kernel_data.tables.filter_table_offset;
-  float raster_x = x + lookup_table_read(kg, filter_u, filter_table_offset, FILTER_TABLE_SIZE);
+  const float2 raster = make_float2(
-  float raster_y = y + lookup_table_read(kg, filter_v, filter_table_offset, FILTER_TABLE_SIZE);
+      x + lookup_table_read(kg, filter_uv.x, filter_table_offset, FILTER_TABLE_SIZE),
+      y + lookup_table_read(kg, filter_uv.y, filter_table_offset, FILTER_TABLE_SIZE));
 
   /* motion blur */
   if (kernel_data.cam.shuttertime == -1.0f) {
@@ -393,14 +386,14 @@ ccl_device_inline void camera_sample(KernelGlobals kg,
 
   /* sample */
   if (kernel_data.cam.type == CAMERA_PERSPECTIVE) {
-    camera_sample_perspective(kg, raster_x, raster_y, lens_u, lens_v, ray);
+    camera_sample_perspective(kg, raster, lens_uv, ray);
   }
   else if (kernel_data.cam.type == CAMERA_ORTHOGRAPHIC) {
-    camera_sample_orthographic(kg, raster_x, raster_y, lens_u, lens_v, ray);
+    camera_sample_orthographic(kg, raster, lens_uv, ray);
   }
   else {
     ccl_global const DecomposedTransform *cam_motion = kernel_data_array(camera_motion);
-    camera_sample_panorama(&kernel_data.cam, cam_motion, raster_x, raster_y, lens_u, lens_v, ray);
+    camera_sample_panorama(&kernel_data.cam, cam_motion, raster, lens_uv, ray);
   }
 }
 

intern/cycles/kernel/closure/bsdf.h

@@ -122,73 +122,53 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
   *pdf = 0.f;
   int label = LABEL_NONE;
   const float3 Ng = (sd->type & PRIMITIVE_CURVE) ? sc->N : sd->Ng;
+  const float2 rand_xy = float3_to_float2(rand);
 
   switch (sc->type) {
     case CLOSURE_BSDF_DIFFUSE_ID:
-      label = bsdf_diffuse_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
+      label = bsdf_diffuse_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
 #if defined(__SVM__) || defined(__OSL__)
     case CLOSURE_BSDF_OREN_NAYAR_ID:
-      label = bsdf_oren_nayar_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
+      label = bsdf_oren_nayar_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
 #  ifdef __OSL__
     case CLOSURE_BSDF_PHONG_RAMP_ID:
-      label = bsdf_phong_ramp_sample(
+      label = bsdf_phong_ramp_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf, sampled_roughness);
-          sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_DIFFUSE_RAMP_ID:
-      label = bsdf_diffuse_ramp_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
+      label = bsdf_diffuse_ramp_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
 #  endif
     case CLOSURE_BSDF_TRANSLUCENT_ID:
-      label = bsdf_translucent_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
+      label = bsdf_translucent_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_TRANSPARENT_ID:
-      label = bsdf_transparent_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
+      label = bsdf_transparent_sample(sc, Ng, sd->wi, eval, wo, pdf);
       *sampled_roughness = zero_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_REFLECTION_ID:
     case CLOSURE_BSDF_REFRACTION_ID:
     case CLOSURE_BSDF_SHARP_GLASS_ID:
-      label = bsdf_microfacet_sharp_sample(sc,
+      label = bsdf_microfacet_sharp_sample(
-                                           path_flag,
+          sc, path_flag, Ng, sd->wi, rand, eval, wo, pdf, sampled_roughness, eta);
-                                           Ng,
-                                           sd->wi,
-                                           rand.x,
-                                           rand.y,
-                                           rand.z,
-                                           eval,
-                                           wo,
-                                           pdf,
-                                           sampled_roughness,
-                                           eta);
       break;
     case CLOSURE_BSDF_MICROFACET_GGX_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_GLASS_ID:
-      label = bsdf_microfacet_ggx_sample(sc,
+      label = bsdf_microfacet_ggx_sample(
-                                         path_flag,
+          sc, path_flag, Ng, sd->wi, rand, eval, wo, pdf, sampled_roughness, eta);
-                                         Ng,
-                                         sd->wi,
-                                         rand.x,
-                                         rand.y,
-                                         rand.z,
-                                         eval,
-                                         wo,
-                                         pdf,
-                                         sampled_roughness,
-                                         eta);
       break;
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_ID:
       label = bsdf_microfacet_multi_ggx_sample(kg,
@@ -221,61 +201,50 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
     case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_GLASS_ID:
-      label = bsdf_microfacet_beckmann_sample(sc,
+      label = bsdf_microfacet_beckmann_sample(
-                                              path_flag,
+          sc, path_flag, Ng, sd->wi, rand, eval, wo, pdf, sampled_roughness, eta);
-                                              Ng,
-                                              sd->wi,
-                                              rand.x,
-                                              rand.y,
-                                              rand.z,
-                                              eval,
-                                              wo,
-                                              pdf,
-                                              sampled_roughness,
-                                              eta);
       break;
     case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
       label = bsdf_ashikhmin_shirley_sample(
-          sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf, sampled_roughness);
+          sc, Ng, sd->wi, rand_xy, eval, wo, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_ASHIKHMIN_VELVET_ID:
-      label = bsdf_ashikhmin_velvet_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
+      label = bsdf_ashikhmin_velvet_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_DIFFUSE_TOON_ID:
-      label = bsdf_diffuse_toon_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
+      label = bsdf_diffuse_toon_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_GLOSSY_TOON_ID:
-      label = bsdf_glossy_toon_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
+      label = bsdf_glossy_toon_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
       // double check if this is valid
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_HAIR_REFLECTION_ID:
       label = bsdf_hair_reflection_sample(
-          sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf, sampled_roughness);
+          sc, Ng, sd->wi, rand_xy, eval, wo, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_HAIR_TRANSMISSION_ID:
       label = bsdf_hair_transmission_sample(
-          sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf, sampled_roughness);
+          sc, Ng, sd->wi, rand_xy, eval, wo, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_HAIR_PRINCIPLED_ID:
-      label = bsdf_principled_hair_sample(
+      label = bsdf_principled_hair_sample(kg, sc, sd, rand, eval, wo, pdf, sampled_roughness, eta);
-          kg, sc, sd, rand.x, rand.y, rand.z, eval, wo, pdf, sampled_roughness, eta);
       break;
     case CLOSURE_BSDF_PRINCIPLED_DIFFUSE_ID:
-      label = bsdf_principled_diffuse_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
+      label = bsdf_principled_diffuse_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_PRINCIPLED_SHEEN_ID:
-      label = bsdf_principled_sheen_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
+      label = bsdf_principled_sheen_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
@@ -686,7 +655,7 @@ ccl_device_inline Spectrum bsdf_albedo(ccl_private const ShaderData *sd,
    * extra overhead though. */
 #if defined(__SVM__) || defined(__OSL__)
   if (CLOSURE_IS_BSDF_MICROFACET(sc->type)) {
-    albedo *= microfacet_fresnel((ccl_private const MicrofacetBsdf *)sc, sd->wi, sc->N, false);
+    albedo *= bsdf_microfacet_estimate_fresnel(sd, (ccl_private const MicrofacetBsdf *)sc);
   }
   else if (sc->type == CLOSURE_BSDF_PRINCIPLED_SHEEN_ID) {
     albedo *= ((ccl_private const PrincipledSheenBsdf *)sc)->avg_value;

intern/cycles/kernel/closure/bsdf_ashikhmin_shirley.h

@@ -111,24 +111,19 @@ ccl_device_forceinline Spectrum bsdf_ashikhmin_shirley_eval(ccl_private const Sh
   return make_spectrum(out);
 }
 
-ccl_device_inline void bsdf_ashikhmin_shirley_sample_first_quadrant(float n_x,
+ccl_device_inline void bsdf_ashikhmin_shirley_sample_first_quadrant(
-                                                                    float n_y,
+    float n_x, float n_y, const float2 rand, ccl_private float *phi, ccl_private float *cos_theta)
-                                                                    float randu,
-                                                                    float randv,
-                                                                    ccl_private float *phi,
-                                                                    ccl_private float *cos_theta)
 {
-  *phi = atanf(sqrtf((n_x + 1.0f) / (n_y + 1.0f)) * tanf(M_PI_2_F * randu));
+  *phi = atanf(sqrtf((n_x + 1.0f) / (n_y + 1.0f)) * tanf(M_PI_2_F * rand.x));
   float cos_phi = cosf(*phi);
   float sin_phi = sinf(*phi);
-  *cos_theta = powf(randv, 1.0f / (n_x * cos_phi * cos_phi + n_y * sin_phi * sin_phi + 1.0f));
+  *cos_theta = powf(rand.y, 1.0f / (n_x * cos_phi * cos_phi + n_y * sin_phi * sin_phi + 1.0f));
 }
 
 ccl_device int bsdf_ashikhmin_shirley_sample(ccl_private const ShaderClosure *sc,
                                              float3 Ng,
                                              float3 wi,
-                                             float randu,
+                                             float2 rand,
-                                             float randv,
                                              ccl_private Spectrum *eval,
                                              ccl_private float3 *wo,
                                              ccl_private float *pdf,
@@ -162,32 +157,28 @@ ccl_device int bsdf_ashikhmin_shirley_sample(ccl_private const ShaderClosure *sc
   float cos_theta;
   if (n_x == n_y) {
     /* isotropic sampling */
-    phi = M_2PI_F * randu;
+    phi = M_2PI_F * rand.x;
-    cos_theta = powf(randv, 1.0f / (n_x + 1.0f));
+    cos_theta = powf(rand.y, 1.0f / (n_x + 1.0f));
   }
   else {
     /* anisotropic sampling */
-    if (randu < 0.25f) { /* first quadrant */
+    if (rand.x < 0.25f) { /* first quadrant */
-      float remapped_randu = 4.0f * randu;
+      rand.x *= 4.0f;
-      bsdf_ashikhmin_shirley_sample_first_quadrant(
+      bsdf_ashikhmin_shirley_sample_first_quadrant(n_x, n_y, rand, &phi, &cos_theta);
-          n_x, n_y, remapped_randu, randv, &phi, &cos_theta);
     }
-    else if (randu < 0.5f) { /* second quadrant */
+    else if (rand.x < 0.5f) { /* second quadrant */
-      float remapped_randu = 4.0f * (.5f - randu);
+      rand.x = 4.0f * (0.5f - rand.x);
-      bsdf_ashikhmin_shirley_sample_first_quadrant(
+      bsdf_ashikhmin_shirley_sample_first_quadrant(n_x, n_y, rand, &phi, &cos_theta);
-          n_x, n_y, remapped_randu, randv, &phi, &cos_theta);
       phi = M_PI_F - phi;
     }
-    else if (randu < 0.75f) { /* third quadrant */
+    else if (rand.x < 0.75f) { /* third quadrant */
-      float remapped_randu = 4.0f * (randu - 0.5f);
+      rand.x = 4.0f * (rand.x - 0.5f);
-      bsdf_ashikhmin_shirley_sample_first_quadrant(
+      bsdf_ashikhmin_shirley_sample_first_quadrant(n_x, n_y, rand, &phi, &cos_theta);
-          n_x, n_y, remapped_randu, randv, &phi, &cos_theta);
       phi = M_PI_F + phi;
     }
     else { /* fourth quadrant */
-      float remapped_randu = 4.0f * (1.0f - randu);
+      rand.x = 4.0f * (1.0f - rand.x);
-      bsdf_ashikhmin_shirley_sample_first_quadrant(
+      bsdf_ashikhmin_shirley_sample_first_quadrant(n_x, n_y, rand, &phi, &cos_theta);
-          n_x, n_y, remapped_randu, randv, &phi, &cos_theta);
       phi = 2.0f * M_PI_F - phi;
     }
   }

intern/cycles/kernel/closure/bsdf_ashikhmin_velvet.h

@@ -78,8 +78,7 @@ ccl_device Spectrum bsdf_ashikhmin_velvet_eval(ccl_private const ShaderClosure *
 ccl_device int bsdf_ashikhmin_velvet_sample(ccl_private const ShaderClosure *sc,
                                             float3 Ng,
                                             float3 wi,
-                                            float randu,
+                                            const float2 rand,
-                                            float randv,
                                             ccl_private Spectrum *eval,
                                             ccl_private float3 *wo,
                                             ccl_private float *pdf)
@@ -90,7 +89,7 @@ ccl_device int bsdf_ashikhmin_velvet_sample(ccl_private const ShaderClosure *sc,
 
   // we are viewing the surface from above - send a ray out with uniform
   // distribution over the hemisphere
-  sample_uniform_hemisphere(N, randu, randv, wo, pdf);
+  sample_uniform_hemisphere(N, rand, wo, pdf);
 
   if (!(dot(Ng, *wo) > 0)) {
     *pdf = 0.0f;

intern/cycles/kernel/closure/bsdf_diffuse.h

@@ -42,8 +42,7 @@ ccl_device Spectrum bsdf_diffuse_eval(ccl_private const ShaderClosure *sc,
 ccl_device int bsdf_diffuse_sample(ccl_private const ShaderClosure *sc,
                                    float3 Ng,
                                    float3 wi,
-                                   float randu,
+                                   float2 rand,
-                                   float randv,
                                    ccl_private Spectrum *eval,
                                    ccl_private float3 *wo,
                                    ccl_private float *pdf)
@@ -52,7 +51,7 @@ ccl_device int bsdf_diffuse_sample(ccl_private const ShaderClosure *sc,
   float3 N = bsdf->N;
 
   // distribution over the hemisphere
-  sample_cos_hemisphere(N, randu, randv, wo, pdf);
+  sample_cos_hemisphere(N, rand, wo, pdf);
 
   if (dot(Ng, *wo) > 0.0f) {
     *eval = make_spectrum(*pdf);
@@ -88,8 +87,7 @@ ccl_device Spectrum bsdf_translucent_eval(ccl_private const ShaderClosure *sc,
 ccl_device int bsdf_translucent_sample(ccl_private const ShaderClosure *sc,
                                        float3 Ng,
                                        float3 wi,
-                                       float randu,
+                                       float2 rand,
-                                       float randv,
                                        ccl_private Spectrum *eval,
                                        ccl_private float3 *wo,
                                        ccl_private float *pdf)
@@ -99,7 +97,7 @@ ccl_device int bsdf_translucent_sample(ccl_private const ShaderClosure *sc,
 
   // we are viewing the surface from the right side - send a ray out with cosine
   // distribution over the hemisphere
-  sample_cos_hemisphere(-N, randu, randv, wo, pdf);
+  sample_cos_hemisphere(-N, rand, wo, pdf);
   if (dot(Ng, *wo) < 0) {
     *eval = make_spectrum(*pdf);
   }

intern/cycles/kernel/closure/bsdf_diffuse_ramp.h

@@ -67,8 +67,7 @@ ccl_device Spectrum bsdf_diffuse_ramp_eval(ccl_private const ShaderClosure *sc,
 ccl_device int bsdf_diffuse_ramp_sample(ccl_private const ShaderClosure *sc,
                                         float3 Ng,
                                         float3 wi,
-                                        float randu,
+                                        float2 rand,
-                                        float randv,
                                         ccl_private Spectrum *eval,
                                         ccl_private float3 *wo,
                                         ccl_private float *pdf)
@@ -77,7 +76,7 @@ ccl_device int bsdf_diffuse_ramp_sample(ccl_private const ShaderClosure *sc,
   float3 N = bsdf->N;
 
   // distribution over the hemisphere
-  sample_cos_hemisphere(N, randu, randv, wo, pdf);
+  sample_cos_hemisphere(N, rand, wo, pdf);
 
   if (dot(Ng, *wo) > 0.0f) {
     *eval = rgb_to_spectrum(bsdf_diffuse_ramp_get_color(bsdf->colors, *pdf * M_PI_F) * M_1_PI_F);

intern/cycles/kernel/closure/bsdf_hair.h

@@ -143,8 +143,7 @@ ccl_device Spectrum bsdf_hair_transmission_eval(ccl_private const ShaderClosure
 ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
                                            float3 Ng,
                                            float3 wi,
-                                           float randu,
+                                           float2 rand,
-                                           float randv,
                                            ccl_private Spectrum *eval,
                                            ccl_private float3 *wo,
                                            ccl_private float *pdf,
@@ -165,7 +164,7 @@ ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
   float a_R = fast_atan2f(((M_PI_2_F + theta_r) * 0.5f - offset) * roughness1_inv, 1.0f);
   float b_R = fast_atan2f(((-M_PI_2_F + theta_r) * 0.5f - offset) * roughness1_inv, 1.0f);
 
-  float t = roughness1 * tanf(randu * (a_R - b_R) + b_R);
+  float t = roughness1 * tanf(rand.x * (a_R - b_R) + b_R);
 
   float theta_h = t + offset;
   float theta_i = 2 * theta_h - theta_r;
@@ -173,7 +172,7 @@ ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
   float costheta_i, sintheta_i;
   fast_sincosf(theta_i, &sintheta_i, &costheta_i);
 
-  float phi = 2 * safe_asinf(1 - 2 * randv) * roughness2;
+  float phi = 2 * safe_asinf(1 - 2 * rand.y) * roughness2;
 
   float phi_pdf = fast_cosf(phi * 0.5f) * 0.25f / roughness2;
 
@@ -196,8 +195,7 @@ ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
 ccl_device int bsdf_hair_transmission_sample(ccl_private const ShaderClosure *sc,
                                              float3 Ng,
                                              float3 wi,
-                                             float randu,
+                                             float2 rand,
-                                             float randv,
                                              ccl_private Spectrum *eval,
                                              ccl_private float3 *wo,
                                              ccl_private float *pdf,
@@ -219,7 +217,7 @@ ccl_device int bsdf_hair_transmission_sample(ccl_private const ShaderClosure *sc
   float b_TT = fast_atan2f(((-M_PI_2_F + theta_r) / 2 - offset) * roughness1_inv, 1.0f);
   float c_TT = 2 * fast_atan2f(M_PI_2_F / roughness2, 1.0f);
 
-  float t = roughness1 * tanf(randu * (a_TT - b_TT) + b_TT);
+  float t = roughness1 * tanf(rand.x * (a_TT - b_TT) + b_TT);
 
   float theta_h = t + offset;
   float theta_i = 2 * theta_h - theta_r;
@@ -227,7 +225,7 @@ ccl_device int bsdf_hair_transmission_sample(ccl_private const ShaderClosure *sc
   float costheta_i, sintheta_i;
   fast_sincosf(theta_i, &sintheta_i, &costheta_i);
 
-  float p = roughness2 * tanf(c_TT * (randv - 0.5f));
+  float p = roughness2 * tanf(c_TT * (rand.y - 0.5f));
   float phi = p + M_PI_F;
   float theta_pdf = roughness1 /
                     (2 * (t * t + roughness1 * roughness1) * (a_TT - b_TT) * costheta_i);

intern/cycles/kernel/closure/bsdf_hair_principled.h

@@ -336,9 +336,7 @@ ccl_device Spectrum bsdf_principled_hair_eval(KernelGlobals kg,
 ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
                                            ccl_private const ShaderClosure *sc,
                                            ccl_private ShaderData *sd,
-                                           float randu,
+                                           float3 rand,
-                                           float randv,
-                                           float randw,
                                            ccl_private Spectrum *eval,
                                            ccl_private float3 *wo,
                                            ccl_private float *pdf,
@@ -381,12 +379,12 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
 
   int p = 0;
   for (; p < 3; p++) {
-    if (randw < Ap_energy[p]) {
+    if (rand.z < Ap_energy[p]) {
       break;
     }
-    randw -= Ap_energy[p];
+    rand.z -= Ap_energy[p];
   }
-  randw /= Ap_energy[p];
+  rand.z /= Ap_energy[p];
 
   float v = bsdf->v;
   if (p == 1) {
@@ -396,9 +394,10 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
     v *= 4.0f;
   }
 
-  randw = max(randw, 1e-5f);
+  rand.z = max(rand.z, 1e-5f);
-  const float fac = 1.0f + v * logf(randw + (1.0f - randw) * expf(-2.0f / v));
+  const float fac = 1.0f + v * logf(rand.z + (1.0f - rand.z) * expf(-2.0f / v));
-  float sin_theta_i = -fac * sin_theta_o + cos_from_sin(fac) * cosf(M_2PI_F * randv) * cos_theta_o;
+  float sin_theta_i = -fac * sin_theta_o +
+                      cos_from_sin(fac) * cosf(M_2PI_F * rand.y) * cos_theta_o;
   float cos_theta_i = cos_from_sin(sin_theta_i);
 
   float angles[6];
@@ -410,10 +409,10 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
 
   float phi;
   if (p < 3) {
-    phi = delta_phi(p, gamma_o, gamma_t) + sample_trimmed_logistic(randu, bsdf->s);
+    phi = delta_phi(p, gamma_o, gamma_t) + sample_trimmed_logistic(rand.x, bsdf->s);
   }
   else {
-    phi = M_2PI_F * randu;
+    phi = M_2PI_F * rand.x;
   }
   const float phi_i = phi_o + phi;
 

intern/cycles/kernel/closure/bsdf_microfacet.h

@@ -76,8 +76,7 @@ static_assert(sizeof(ShaderClosure) >= sizeof(MicrofacetBsdf), "MicrofacetBsdf i
 ccl_device_forceinline float3 microfacet_beckmann_sample_vndf(const float3 wi,
                                                               const float alpha_x,
                                                               const float alpha_y,
-                                                              const float randu,
+                                                              const float2 rand)
-                                                              const float randv)
 {
   /* 1. stretch wi */
   float3 wi_ = make_float3(alpha_x * wi.x, alpha_y * wi.y, wi.z);
@@ -90,8 +89,8 @@ ccl_device_forceinline float3 microfacet_beckmann_sample_vndf(const float3 wi,
 
   if (wi_.z >= 0.99999f) {
     /* Special case (normal incidence). */
-    const float r = sqrtf(-logf(randu));
+    const float r = sqrtf(-logf(rand.x));
-    const float phi = M_2PI_F * randv;
+    const float phi = M_2PI_F * rand.y;
     slope_x = r * cosf(phi);
     slope_y = r * sinf(phi);
   }
@@ -125,8 +124,8 @@ ccl_device_forceinline float3 microfacet_beckmann_sample_vndf(const float3 wi,
      *   solve y = 1 + b + K * (1 - b * b)
      */
     const float K = tan_theta_i * SQRT_PI_INV;
-    const float y_approx = randu * (1.0f + erf_a + K * (1 - erf_a * erf_a));
+    const float y_approx = rand.x * (1.0f + erf_a + K * (1 - erf_a * erf_a));
-    const float y_exact = randu * (1.0f + erf_a + K * exp_a2);
+    const float y_exact = rand.x * (1.0f + erf_a + K * exp_a2);
     float b = K > 0 ? (0.5f - sqrtf(K * (K - y_approx + 1.0f) + 0.25f)) / K : y_approx - 1.0f;
 
     float inv_erf = fast_ierff(b);
@@ -155,7 +154,7 @@ ccl_device_forceinline float3 microfacet_beckmann_sample_vndf(const float3 wi,
     }
 
     slope_x = inv_erf;
-    slope_y = fast_ierff(2.0f * randv - 1.0f);
+    slope_y = fast_ierff(2.0f * rand.y - 1.0f);
   }
 
   /* 3. rotate */
@@ -178,8 +177,7 @@ ccl_device_forceinline float3 microfacet_beckmann_sample_vndf(const float3 wi,
 ccl_device_forceinline float3 microfacet_ggx_sample_vndf(const float3 wi,
                                                          const float alpha_x,
                                                          const float alpha_y,
-                                                         const float randu,
+                                                         const float2 rand)
-                                                         const float randv)
 {
   /* Section 3.2: Transforming the view direction to the hemisphere configuration. */
   float3 wi_ = normalize(make_float3(alpha_x * wi.x, alpha_y * wi.y, wi.z));
@@ -198,7 +196,7 @@ ccl_device_forceinline float3 microfacet_ggx_sample_vndf(const float3 wi,
   }
 
   /* Section 4.2: Parameterization of the projected area. */
-  float2 t = concentric_sample_disk(randu, randv);
+  float2 t = concentric_sample_disk(rand);
   t.y = mix(safe_sqrtf(1.0f - sqr(t.x)), t.y, 0.5f * (1.0f + wi_.z));
 
   /* Section 4.3: Reprojection onto hemisphere. */
@@ -264,19 +262,39 @@ ccl_device_forceinline Spectrum microfacet_fresnel(ccl_private const MicrofacetB
     }
   }
   else if (bsdf->fresnel_type == MicrofacetFresnel::CONSTANT) {
-    kernel_assert(!refraction);
+    /* CONSTANT is only used my MultiGGX, which doesn't call this function.
-    ccl_private FresnelConstant *fresnel = (ccl_private FresnelConstant *)bsdf->fresnel;
+     * Therefore, this case only happens when determining the albedo of a MultiGGX closure.
-    return fresnel->color;
+     * In that case, return 1.0 since the constant color is already baked into the weight. */
+    return one_spectrum();
   }
   else {
     return one_spectrum();
   }
 }
 
-ccl_device_forceinline void bsdf_microfacet_adjust_weight(ccl_private const ShaderData *sd,
+/* This function estimates the albedo of the BSDF (NOT including the bsdf->weight) as caused by
-                                                          ccl_private MicrofacetBsdf *bsdf)
+ * the applied Fresnel model for the given view direction.
+ * The base microfacet model is assumed to have an albedo of 1, but e.g. a reflection-only
+ * closure with Fresnel applied can end up having a very low overall albedo.
+ * This is used to adjust the sample weight, as well as for the Diff/Gloss/Trans Color pass
+ * and the Denoising Albedo pass. */
+ccl_device Spectrum bsdf_microfacet_estimate_fresnel(ccl_private const ShaderData *sd,
+                                                     ccl_private const MicrofacetBsdf *bsdf)
 {
-  bsdf->sample_weight *= average(microfacet_fresnel(bsdf, sd->wi, bsdf->N, false));
+  const bool is_glass = CLOSURE_IS_GLASS(bsdf->type);
+  const bool is_refractive = CLOSURE_IS_REFRACTIVE(bsdf->type);
+
+  Spectrum albedo = zero_spectrum();
+  if (!is_refractive || is_glass) {
+    /* BSDF has a reflective lobe. */
+    albedo += microfacet_fresnel(bsdf, sd->wi, bsdf->N, false);
+  }
+  if (is_refractive) {
+    /* BSDF has a refractive lobe (unless there's TIR). */
+    albedo += microfacet_fresnel(bsdf, sd->wi, bsdf->N, true);
+  }
+
+  return albedo;
 }
 
 /* Generalized Trowbridge-Reitz for clearcoat. */
@@ -398,7 +416,7 @@ ccl_device Spectrum bsdf_microfacet_eval(ccl_private const ShaderClosure *sc,
    * - Purely reflective closures can't have refraction.
    * - Purely refractive closures can't have reflection.
    */
-  if ((cos_NI <= 0) || (alpha_x * alpha_y <= 1e-7f) || ((cos_NgO < 0.0f) != is_refraction) ||
+  if ((cos_NI <= 0) || (alpha_x * alpha_y <= 5e-7f) || ((cos_NgO < 0.0f) != is_refraction) ||
       (is_refraction && !m_refractive) || (!is_refraction && m_refractive && !m_glass))
   {
     *pdf = 0.0f;
@@ -467,9 +485,7 @@ ccl_device int bsdf_microfacet_sample(ccl_private const ShaderClosure *sc,
                                       const int path_flag,
                                       float3 Ng,
                                       float3 wi,
-                                      float randu,
+                                      const float3 rand,
-                                      float randv,
-                                      float randw,
                                       ccl_private Spectrum *eval,
                                       ccl_private float3 *wo,
                                       ccl_private float *pdf,
@@ -482,7 +498,7 @@ ccl_device int bsdf_microfacet_sample(ccl_private const ShaderClosure *sc,
   const bool m_refractive = CLOSURE_IS_REFRACTIVE(bsdf->type);
   const float alpha_x = bsdf->alpha_x;
   const float alpha_y = bsdf->alpha_y;
-  bool m_singular = (m_type == MicrofacetType::SHARP) || (alpha_x * alpha_y <= 1e-7f);
+  bool m_singular = (m_type == MicrofacetType::SHARP) || (alpha_x * alpha_y <= 5e-7f);
 
   const float3 N = bsdf->N;
   const float cos_NI = dot(N, wi);
@@ -513,11 +529,11 @@ ccl_device int bsdf_microfacet_sample(ccl_private const ShaderClosure *sc,
      * space before and after sampling. */
     local_I = make_float3(dot(X, wi), dot(Y, wi), cos_NI);
     if (m_type == MicrofacetType::GGX) {
-      local_H = microfacet_ggx_sample_vndf(local_I, alpha_x, alpha_y, randu, randv);
+      local_H = microfacet_ggx_sample_vndf(local_I, alpha_x, alpha_y, float3_to_float2(rand));
     }
     else {
       /* m_type == MicrofacetType::BECKMANN */
-      local_H = microfacet_beckmann_sample_vndf(local_I, alpha_x, alpha_y, randu, randv);
+      local_H = microfacet_beckmann_sample_vndf(local_I, alpha_x, alpha_y, float3_to_float2(rand));
     }
 
     H = X * local_H.x + Y * local_H.y + N * local_H.z;
@@ -545,7 +561,7 @@ ccl_device int bsdf_microfacet_sample(ccl_private const ShaderClosure *sc,
          * excessive noise for reflection highlights. */
         float reflect_pdf = (path_flag & PATH_RAY_CAMERA) ? clamp(fresnel, 0.125f, 0.875f) :
                                                             fresnel;
-        do_refract = (randw >= reflect_pdf);
+        do_refract = (rand.z >= reflect_pdf);
         lobe_pdf = do_refract ? (1.0f - reflect_pdf) : reflect_pdf;
       }
     }
@@ -651,7 +667,7 @@ ccl_device void bsdf_microfacet_setup_fresnel_principledv1(
 
   bsdf->fresnel_type = MicrofacetFresnel::PRINCIPLED_V1;
   bsdf->fresnel = fresnel;
-  bsdf_microfacet_adjust_weight(sd, bsdf);
+  bsdf->sample_weight *= average(bsdf_microfacet_estimate_fresnel(sd, bsdf));
 }
 
 ccl_device void bsdf_microfacet_setup_fresnel_conductor(ccl_private MicrofacetBsdf *bsdf,
@@ -660,7 +676,7 @@ ccl_device void bsdf_microfacet_setup_fresnel_conductor(ccl_private MicrofacetBs
 {
   bsdf->fresnel_type = MicrofacetFresnel::CONDUCTOR;
   bsdf->fresnel = fresnel;
-  bsdf_microfacet_adjust_weight(sd, bsdf);
+  bsdf->sample_weight *= average(bsdf_microfacet_estimate_fresnel(sd, bsdf));
 }
 
 ccl_device void bsdf_microfacet_setup_fresnel_dielectric_tint(
@@ -670,7 +686,7 @@ ccl_device void bsdf_microfacet_setup_fresnel_dielectric_tint(
 {
   bsdf->fresnel_type = MicrofacetFresnel::DIELECTRIC_TINT;
   bsdf->fresnel = fresnel;
-  bsdf_microfacet_adjust_weight(sd, bsdf);
+  bsdf->sample_weight *= average(bsdf_microfacet_estimate_fresnel(sd, bsdf));
 }
 
 ccl_device void bsdf_microfacet_setup_fresnel_generalized_schlick(
@@ -680,7 +696,7 @@ ccl_device void bsdf_microfacet_setup_fresnel_generalized_schlick(
 {
   bsdf->fresnel_type = MicrofacetFresnel::GENERALIZED_SCHLICK;
   bsdf->fresnel = fresnel;
-  bsdf_microfacet_adjust_weight(sd, bsdf);
+  bsdf->sample_weight *= average(bsdf_microfacet_estimate_fresnel(sd, bsdf));
 }
 
 /* GGX microfacet with Smith shadow-masking from:
@@ -715,8 +731,7 @@ ccl_device int bsdf_microfacet_ggx_clearcoat_setup(ccl_private MicrofacetBsdf *b
 
   bsdf->fresnel_type = MicrofacetFresnel::DIELECTRIC;
   bsdf->type = CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID;
-
+  bsdf->sample_weight *= average(bsdf_microfacet_estimate_fresnel(sd, bsdf));
-  bsdf_microfacet_adjust_weight(sd, bsdf);
 
   return SD_BSDF | SD_BSDF_HAS_EVAL;
 }
@@ -764,9 +779,7 @@ ccl_device int bsdf_microfacet_ggx_sample(ccl_private const ShaderClosure *sc,
                                           const int path_flag,
                                           float3 Ng,
                                           float3 wi,
-                                          float randu,
+                                          const float3 rand,
-                                          float randv,
-                                          float randw,
                                           ccl_private Spectrum *eval,
                                           ccl_private float3 *wo,
                                           ccl_private float *pdf,
@@ -774,7 +787,7 @@ ccl_device int bsdf_microfacet_ggx_sample(ccl_private const ShaderClosure *sc,
                                           ccl_private float *eta)
 {
   return bsdf_microfacet_sample<MicrofacetType::GGX>(
-      sc, path_flag, Ng, wi, randu, randv, randw, eval, wo, pdf, sampled_roughness, eta);
+      sc, path_flag, Ng, wi, rand, eval, wo, pdf, sampled_roughness, eta);
 }
 
 /* Beckmann microfacet with Smith shadow-masking from:
@@ -833,9 +846,7 @@ ccl_device int bsdf_microfacet_beckmann_sample(ccl_private const ShaderClosure *
                                                const int path_flag,
                                                float3 Ng,
                                                float3 wi,
-                                               float randu,
+                                               const float3 rand,
-                                               float randv,
-                                               float randw,
                                                ccl_private Spectrum *eval,
                                                ccl_private float3 *wo,
                                                ccl_private float *pdf,
@@ -843,7 +854,7 @@ ccl_device int bsdf_microfacet_beckmann_sample(ccl_private const ShaderClosure *
                                                ccl_private float *eta)
 {
   return bsdf_microfacet_sample<MicrofacetType::BECKMANN>(
-      sc, path_flag, Ng, wi, randu, randv, randw, eval, wo, pdf, sampled_roughness, eta);
+      sc, path_flag, Ng, wi, rand, eval, wo, pdf, sampled_roughness, eta);
 }
 
 /* Specular interface, not really a microfacet model but close enough that sharing code makes
@@ -889,9 +900,7 @@ ccl_device int bsdf_microfacet_sharp_sample(ccl_private const ShaderClosure *sc,
                                             const int path_flag,
                                             float3 Ng,
                                             float3 wi,
-                                            float randu,
+                                            const float3 rand,
-                                            float randv,
-                                            float randw,
                                             ccl_private Spectrum *eval,
                                             ccl_private float3 *wo,
                                             ccl_private float *pdf,
@@ -899,7 +908,7 @@ ccl_device int bsdf_microfacet_sharp_sample(ccl_private const ShaderClosure *sc,
                                             ccl_private float *eta)
 {
   return bsdf_microfacet_sample<MicrofacetType::SHARP>(
-      sc, path_flag, Ng, wi, randu, randv, randw, eval, wo, pdf, sampled_roughness, eta);
+      sc, path_flag, Ng, wi, rand, eval, wo, pdf, sampled_roughness, eta);
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/closure/bsdf_microfacet_multi.h

@@ -407,8 +407,7 @@ ccl_device int bsdf_microfacet_multi_ggx_fresnel_setup(ccl_private MicrofacetBsd
 
   bsdf->fresnel_type = MicrofacetFresnel::PRINCIPLED_V1;
   bsdf->type = CLOSURE_BSDF_MICROFACET_MULTI_GGX_ID;
-
+  bsdf->sample_weight *= average(bsdf_microfacet_estimate_fresnel(sd, bsdf));
-  bsdf_microfacet_adjust_weight(sd, bsdf);
 
   return bsdf_microfacet_multi_ggx_common_setup(bsdf);
 }
@@ -619,8 +618,7 @@ ccl_device int bsdf_microfacet_multi_ggx_glass_fresnel_setup(ccl_private Microfa
 
   bsdf->fresnel_type = MicrofacetFresnel::PRINCIPLED_V1;
   bsdf->type = CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_ID;
-
+  bsdf->sample_weight *= average(bsdf_microfacet_estimate_fresnel(sd, bsdf));
-  bsdf_microfacet_adjust_weight(sd, bsdf);
 
   return SD_BSDF | SD_BSDF_HAS_EVAL | SD_BSDF_NEEDS_LCG;
 }

intern/cycles/kernel/closure/bsdf_oren_nayar.h

@@ -66,14 +66,13 @@ ccl_device Spectrum bsdf_oren_nayar_eval(ccl_private const ShaderClosure *sc,
 ccl_device int bsdf_oren_nayar_sample(ccl_private const ShaderClosure *sc,
                                       float3 Ng,
                                       float3 wi,
-                                      float randu,
+                                      float2 rand,
-                                      float randv,
                                       ccl_private Spectrum *eval,
                                       ccl_private float3 *wo,
                                       ccl_private float *pdf)
 {
   ccl_private const OrenNayarBsdf *bsdf = (ccl_private const OrenNayarBsdf *)sc;
-  sample_uniform_hemisphere(bsdf->N, randu, randv, wo, pdf);
+  sample_uniform_hemisphere(bsdf->N, rand, wo, pdf);
 
   if (dot(Ng, *wo) > 0.0f) {
     *eval = bsdf_oren_nayar_get_intensity(sc, bsdf->N, wi, *wo);

intern/cycles/kernel/closure/bsdf_phong_ramp.h

@@ -78,8 +78,7 @@ ccl_device_inline float phong_ramp_exponent_to_roughness(float exponent)
 ccl_device int bsdf_phong_ramp_sample(ccl_private const ShaderClosure *sc,
                                       float3 Ng,
                                       float3 wi,
-                                      float randu,
+                                      const float2 rand,
-                                      float randv,
                                       ccl_private Spectrum *eval,
                                       ccl_private float3 *wo,
                                       ccl_private float *pdf,
@@ -96,8 +95,8 @@ ccl_device int bsdf_phong_ramp_sample(ccl_private const ShaderClosure *sc,
     float3 R = (2 * cosNI) * bsdf->N - wi;
     float3 T, B;
     make_orthonormals(R, &T, &B);
-    float phi = M_2PI_F * randu;
+    float phi = M_2PI_F * rand.x;
-    float cosTheta = powf(randv, 1 / (m_exponent + 1));
+    float cosTheta = powf(rand.y, 1 / (m_exponent + 1));
     float sinTheta2 = 1 - cosTheta * cosTheta;
     float sinTheta = sinTheta2 > 0 ? sqrtf(sinTheta2) : 0;
     *wo = (cosf(phi) * sinTheta) * T + (sinf(phi) * sinTheta) * B + (cosTheta)*R;

intern/cycles/kernel/closure/bsdf_principled_diffuse.h

@@ -132,8 +132,7 @@ ccl_device Spectrum bsdf_principled_diffuse_eval(ccl_private const ShaderClosure
 ccl_device int bsdf_principled_diffuse_sample(ccl_private const ShaderClosure *sc,
                                               float3 Ng,
                                               float3 wi,
-                                              float randu,
+                                              float2 rand,
-                                              float randv,
                                               ccl_private Spectrum *eval,
                                               ccl_private float3 *wo,
                                               ccl_private float *pdf)
@@ -142,7 +141,7 @@ ccl_device int bsdf_principled_diffuse_sample(ccl_private const ShaderClosure *s
 
   float3 N = bsdf->N;
 
-  sample_cos_hemisphere(N, randu, randv, wo, pdf);
+  sample_cos_hemisphere(N, rand, wo, pdf);
 
   if (dot(Ng, *wo) > 0) {
     *eval = bsdf_principled_diffuse_compute_brdf(bsdf, N, wi, *wo, pdf);

intern/cycles/kernel/closure/bsdf_principled_sheen.h

@@ -84,8 +84,7 @@ ccl_device Spectrum bsdf_principled_sheen_eval(ccl_private const ShaderClosure *
 ccl_device int bsdf_principled_sheen_sample(ccl_private const ShaderClosure *sc,
                                             float3 Ng,
                                             float3 wi,
-                                            float randu,
+                                            float2 rand,
-                                            float randv,
                                             ccl_private Spectrum *eval,
                                             ccl_private float3 *wo,
                                             ccl_private float *pdf)
@@ -94,7 +93,7 @@ ccl_device int bsdf_principled_sheen_sample(ccl_private const ShaderClosure *sc,
 
   float3 N = bsdf->N;
 
-  sample_cos_hemisphere(N, randu, randv, wo, pdf);
+  sample_cos_hemisphere(N, rand, wo, pdf);
 
   if (dot(Ng, *wo) > 0) {
     float3 H = normalize(wi + *wo);

intern/cycles/kernel/closure/bsdf_toon.h

@@ -79,8 +79,7 @@ ccl_device Spectrum bsdf_diffuse_toon_eval(ccl_private const ShaderClosure *sc,
 ccl_device int bsdf_diffuse_toon_sample(ccl_private const ShaderClosure *sc,
                                         float3 Ng,
                                         float3 wi,
-                                        float randu,
+                                        float2 rand,
-                                        float randv,
                                         ccl_private Spectrum *eval,
                                         ccl_private float3 *wo,
                                         ccl_private float *pdf)
@@ -89,10 +88,10 @@ ccl_device int bsdf_diffuse_toon_sample(ccl_private const ShaderClosure *sc,
   float max_angle = bsdf->size * M_PI_2_F;
   float smooth = bsdf->smooth * M_PI_2_F;
   float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
-  float angle = sample_angle * randu;
+  float angle = sample_angle * rand.x;
 
   if (sample_angle > 0.0f) {
-    sample_uniform_cone(bsdf->N, sample_angle, randu, randv, wo, pdf);
+    sample_uniform_cone(bsdf->N, sample_angle, rand, wo, pdf);
 
     if (dot(Ng, *wo) > 0.0f) {
       *eval = make_spectrum(*pdf * bsdf_toon_get_intensity(max_angle, smooth, angle));
@@ -152,8 +151,7 @@ ccl_device Spectrum bsdf_glossy_toon_eval(ccl_private const ShaderClosure *sc,
 ccl_device int bsdf_glossy_toon_sample(ccl_private const ShaderClosure *sc,
                                        float3 Ng,
                                        float3 wi,
-                                       float randu,
+                                       float2 rand,
-                                       float randv,
                                        ccl_private Spectrum *eval,
                                        ccl_private float3 *wo,
                                        ccl_private float *pdf)
@@ -168,9 +166,9 @@ 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 * randu;
+    float angle = sample_angle * rand.x;
 
-    sample_uniform_cone(R, sample_angle, randu, randv, wo, pdf);
+    sample_uniform_cone(R, sample_angle, rand, wo, pdf);
 
     if (dot(Ng, *wo) > 0.0f) {
       float cosNO = dot(bsdf->N, *wo);

intern/cycles/kernel/closure/bsdf_transparent.h

@@ -71,8 +71,6 @@ ccl_device Spectrum bsdf_transparent_eval(ccl_private const ShaderClosure *sc,
 ccl_device int bsdf_transparent_sample(ccl_private const ShaderClosure *sc,
                                        float3 Ng,
                                        float3 wi,
-                                       float randu,
-                                       float randv,
                                        ccl_private Spectrum *eval,
                                        ccl_private float3 *wo,
                                        ccl_private float *pdf)

intern/cycles/kernel/closure/volume.h

@@ -67,21 +67,20 @@ ccl_device Spectrum volume_henyey_greenstein_eval_phase(ccl_private const Shader
   return make_spectrum(*pdf);
 }
 
-ccl_device float3
+ccl_device float3 henyey_greenstrein_sample(float3 D, float g, float2 rand, ccl_private float *pdf)
-henyey_greenstrein_sample(float3 D, float g, float randu, float randv, ccl_private float *pdf)
 {
   /* match pdf for small g */
   float cos_theta;
   bool isotropic = fabsf(g) < 1e-3f;
 
   if (isotropic) {
-    cos_theta = (1.0f - 2.0f * randu);
+    cos_theta = (1.0f - 2.0f * rand.x);
     if (pdf) {
       *pdf = M_1_PI_F * 0.25f;
     }
   }
   else {
-    float k = (1.0f - g * g) / (1.0f - g + 2.0f * g * randu);
+    float k = (1.0f - g * g) / (1.0f - g + 2.0f * g * rand.x);
     cos_theta = (1.0f + g * g - k * k) / (2.0f * g);
     if (pdf) {
       *pdf = single_peaked_henyey_greenstein(cos_theta, g);
@@ -89,7 +88,7 @@ henyey_greenstrein_sample(float3 D, float g, float randu, float randv, ccl_priva
   }
 
   float sin_theta = sin_from_cos(cos_theta);
-  float phi = M_2PI_F * randv;
+  float phi = M_2PI_F * rand.y;
   float3 dir = make_float3(sin_theta * cosf(phi), sin_theta * sinf(phi), cos_theta);
 
   float3 T, B;
@@ -101,8 +100,7 @@ henyey_greenstrein_sample(float3 D, float g, float randu, float randv, ccl_priva
 
 ccl_device int volume_henyey_greenstein_sample(ccl_private const ShaderVolumeClosure *svc,
                                                float3 wi,
-                                               float randu,
+                                               float2 rand,
-                                               float randv,
                                                ccl_private Spectrum *eval,
                                                ccl_private float3 *wo,
                                                ccl_private float *pdf)
@@ -110,7 +108,7 @@ ccl_device int volume_henyey_greenstein_sample(ccl_private const ShaderVolumeClo
   float g = svc->g;
 
   /* note that wi points towards the viewer and so is used negated */
-  *wo = henyey_greenstrein_sample(-wi, g, randu, randv, pdf);
+  *wo = henyey_greenstrein_sample(-wi, g, rand, pdf);
   *eval = make_spectrum(*pdf); /* perfect importance sampling */
 
   return LABEL_VOLUME_SCATTER;
@@ -128,13 +126,12 @@ ccl_device Spectrum volume_phase_eval(ccl_private const ShaderData *sd,
 
 ccl_device int volume_phase_sample(ccl_private const ShaderData *sd,
                                    ccl_private const ShaderVolumeClosure *svc,
-                                   float randu,
+                                   float2 rand,
-                                   float randv,
                                    ccl_private Spectrum *eval,
                                    ccl_private float3 *wo,
                                    ccl_private float *pdf)
 {
-  return volume_henyey_greenstein_sample(svc, sd->wi, randu, randv, eval, wo, pdf);
+  return volume_henyey_greenstein_sample(svc, sd->wi, rand, eval, wo, pdf);
 }
 
 /* Volume sampling utilities. */

intern/cycles/kernel/data_template.h

@@ -36,7 +36,7 @@ KERNEL_STRUCT_MEMBER(background, int, map_res_x)
 KERNEL_STRUCT_MEMBER(background, int, map_res_y)
 /* Multiple importance sampling. */
 KERNEL_STRUCT_MEMBER(background, int, use_mis)
-/* Lightgroup. */
+/* Light-group. */
 KERNEL_STRUCT_MEMBER(background, int, lightgroup)
 /* Light Index. */
 KERNEL_STRUCT_MEMBER(background, int, light_index)

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

@@ -178,14 +178,19 @@ ccl_device_inline void kernel_embree_setup_rayhit(const Ray &ray,
   rayhit.hit.instID[0] = RTC_INVALID_GEOMETRY_ID;
 }
 
+ccl_device_inline int kernel_embree_get_hit_object(const RTCHit *hit)
+{
+  return (hit->instID[0] != RTC_INVALID_GEOMETRY_ID ? hit->instID[0] : hit->geomID) / 2;
+}
+
 ccl_device_inline bool kernel_embree_is_self_intersection(const KernelGlobals kg,
                                                           const RTCHit *hit,
                                                           const Ray *ray,
                                                           const intptr_t prim_offset)
 {
-  int object, prim;
+  const int object = kernel_embree_get_hit_object(hit);
-  object = (hit->instID[0] != RTC_INVALID_GEOMETRY_ID ? hit->instID[0] : hit->geomID) / 2;
 
+  int prim;
   if ((ray->self.object == object) || (ray->self.light_object == object)) {
     prim = hit->primID + prim_offset;
   }
@@ -209,7 +214,7 @@ ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg,
 {
   isect->t = ray->tfar;
   isect->prim = hit->primID + prim_offset;
-  isect->object = hit->instID[0] != RTC_INVALID_GEOMETRY_ID ? hit->instID[0] / 2 : hit->geomID / 2;
+  isect->object = kernel_embree_get_hit_object(hit);
 
   const bool is_hair = hit->geomID & 1;
   if (is_hair) {
@@ -287,7 +292,15 @@ ccl_device_forceinline void kernel_embree_filter_intersection_func_impl(
           kg, hit, cray, reinterpret_cast<intptr_t>(args->geometryUserPtr)))
   {
     *args->valid = 0;
+    return;
   }
+
+#ifdef __SHADOW_LINKING__
+  if (intersection_skip_shadow_link(kg, cray, kernel_embree_get_hit_object(hit))) {
+    *args->valid = 0;
+    return;
+  }
+#endif
 }
 
 /* This gets called by Embree at every valid ray/object intersection.
@@ -323,6 +336,13 @@ ccl_device_forceinline void kernel_embree_filter_occluded_shadow_all_func_impl(
     return;
   }
 
+#ifdef __SHADOW_LINKING__
+  if (intersection_skip_shadow_link(kg, cray, current_isect.object)) {
+    *args->valid = 0;
+    return;
+  }
+#endif
+
   /* If no transparent shadows or max number of hits exceeded, all light is blocked. */
   const int flags = intersection_get_shader_flags(kg, current_isect.prim, current_isect.type);
   if (!(flags & (SD_HAS_TRANSPARENT_SHADOW)) || ctx->num_hits >= ctx->max_hits) {

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

@@ -28,11 +28,13 @@ KERNEL_INTEGRATOR_SHADE_FUNCTION(intersect_closest);
 KERNEL_INTEGRATOR_FUNCTION(intersect_shadow);
 KERNEL_INTEGRATOR_FUNCTION(intersect_subsurface);
 KERNEL_INTEGRATOR_FUNCTION(intersect_volume_stack);
+KERNEL_INTEGRATOR_FUNCTION(intersect_dedicated_light);
 KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_background);
 KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_light);
 KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_shadow);
 KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_surface);
 KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_volume);
+KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_dedicated_light);
 KERNEL_INTEGRATOR_SHADE_FUNCTION(megakernel);
 
 #undef KERNEL_INTEGRATOR_FUNCTION

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

@@ -102,10 +102,12 @@ DEFINE_INTEGRATOR_INIT_KERNEL(init_from_bake)
 DEFINE_INTEGRATOR_SHADE_KERNEL(intersect_closest)
 DEFINE_INTEGRATOR_KERNEL(intersect_subsurface)
 DEFINE_INTEGRATOR_KERNEL(intersect_volume_stack)
+DEFINE_INTEGRATOR_KERNEL(intersect_dedicated_light)
 DEFINE_INTEGRATOR_SHADE_KERNEL(shade_background)
 DEFINE_INTEGRATOR_SHADE_KERNEL(shade_light)
 DEFINE_INTEGRATOR_SHADE_KERNEL(shade_surface)
 DEFINE_INTEGRATOR_SHADE_KERNEL(shade_volume)
+DEFINE_INTEGRATOR_SHADE_KERNEL(shade_dedicated_light)
 DEFINE_INTEGRATOR_SHADE_KERNEL(megakernel)
 DEFINE_INTEGRATOR_SHADOW_KERNEL(intersect_shadow)
 DEFINE_INTEGRATOR_SHADOW_SHADE_KERNEL(shade_shadow)

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

@@ -27,10 +27,12 @@
 #include "kernel/integrator/init_from_bake.h"
 #include "kernel/integrator/init_from_camera.h"
 #include "kernel/integrator/intersect_closest.h"
+#include "kernel/integrator/intersect_dedicated_light.h"
 #include "kernel/integrator/intersect_shadow.h"
 #include "kernel/integrator/intersect_subsurface.h"
 #include "kernel/integrator/intersect_volume_stack.h"
 #include "kernel/integrator/shade_background.h"
+#include "kernel/integrator/shade_dedicated_light.h"
 #include "kernel/integrator/shade_light.h"
 #include "kernel/integrator/shade_shadow.h"
 #include "kernel/integrator/shade_surface.h"
@@ -196,6 +198,20 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
 }
 ccl_gpu_kernel_postfix
 
+ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
+    ccl_gpu_kernel_signature(integrator_intersect_dedicated_light,
+                             ccl_global const int *path_index_array,
+                             const int work_size)
+{
+  const int global_index = ccl_gpu_global_id_x();
+
+  if (ccl_gpu_kernel_within_bounds(global_index, work_size)) {
+    const int state = (path_index_array) ? path_index_array[global_index] : global_index;
+    ccl_gpu_kernel_call(integrator_intersect_dedicated_light(NULL, state));
+  }
+}
+ccl_gpu_kernel_postfix
+
 #  ifdef __KERNEL_ONEAPI__
 #    include "kernel/device/oneapi/context_intersect_end.h"
 #  endif
@@ -334,6 +350,21 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
 }
 ccl_gpu_kernel_postfix
 
+ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
+    ccl_gpu_kernel_signature(integrator_shade_dedicated_light,
+                             ccl_global const int *path_index_array,
+                             ccl_global float *render_buffer,
+                             const int work_size)
+{
+  const int global_index = ccl_gpu_global_id_x();
+
+  if (ccl_gpu_kernel_within_bounds(global_index, work_size)) {
+    const int state = (path_index_array) ? path_index_array[global_index] : global_index;
+    ccl_gpu_kernel_call(integrator_shade_dedicated_light(NULL, state, render_buffer));
+  }
+}
+ccl_gpu_kernel_postfix
+
 ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_queued_paths_array,
                              int num_states,

intern/cycles/kernel/device/hiprt/hiprt_kernels.h

@@ -66,6 +66,22 @@ ccl_gpu_kernel_threads(GPU_HIPRT_KERNEL_BLOCK_NUM_THREADS)
     ccl_gpu_kernel_call(integrator_intersect_volume_stack(kg, state));
   }
 }
+
+ccl_gpu_kernel_threads(GPU_HIPRT_KERNEL_BLOCK_NUM_THREADS)
+    ccl_gpu_kernel_signature(integrator_intersect_dedicated_light,
+                             ccl_global const int *path_index_array,
+                             const int work_size,
+                             ccl_global int *stack_buffer)
+{
+  const int global_index = ccl_gpu_global_id_x();
+
+  if (global_index < work_size) {
+    HIPRT_INIT_KERNEL_GLOBAL()
+    const int state = (path_index_array) ? path_index_array[global_index] : global_index;
+    ccl_gpu_kernel_call(integrator_intersect_dedicated_light(kg, state));
+  }
+}
+
 ccl_gpu_kernel_postfix
 ccl_gpu_kernel_threads(GPU_HIPRT_KERNEL_BLOCK_NUM_THREADS)
     ccl_gpu_kernel_signature(integrator_shade_surface_raytrace,

intern/cycles/kernel/device/oneapi/kernel.cpp

@@ -122,6 +122,7 @@ size_t oneapi_kernel_preferred_local_size(SyclQueue *queue,
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
     case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
+    case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
     case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND:
     case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
     case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE:
@@ -129,6 +130,7 @@ size_t oneapi_kernel_preferred_local_size(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;
 
@@ -435,6 +437,15 @@ bool oneapi_enqueue_kernel(KernelContext *kernel_context,
                       oneapi_kernel_integrator_intersect_volume_stack);
           break;
         }
+        case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT: {
+          oneapi_call(kg,
+                      cgh,
+                      global_size,
+                      local_size,
+                      args,
+                      oneapi_kernel_integrator_intersect_dedicated_light);
+          break;
+        }
         case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND: {
           oneapi_call(
               kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_shade_background);
@@ -474,6 +485,15 @@ bool oneapi_enqueue_kernel(KernelContext *kernel_context,
               kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_shade_volume);
           break;
         }
+        case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT: {
+          oneapi_call(kg,
+                      cgh,
+                      global_size,
+                      local_size,
+                      args,
+                      oneapi_kernel_integrator_shade_dedicated_light);
+          break;
+        }
         case DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY: {
           oneapi_call(
               kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_queued_paths_array);

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

@@ -183,6 +183,12 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
     return optixIgnoreIntersection();
   }
 
+#  ifdef __SHADOW_LINKING__
+  if (intersection_skip_shadow_link(nullptr, ray, object)) {
+    return optixIgnoreIntersection();
+  }
+#  endif
+
 #  ifndef __TRANSPARENT_SHADOWS__
   /* No transparent shadows support compiled in, make opaque. */
   optixSetPayload_5(true);
@@ -326,6 +332,12 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
   ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
 
   if (visibility & PATH_RAY_SHADOW_OPAQUE) {
+#ifdef __SHADOW_LINKING__
+    if (intersection_skip_shadow_link(nullptr, ray, object)) {
+      return optixIgnoreIntersection();
+    }
+#endif
+
     if (intersection_skip_self_shadow(ray->self, object, prim)) {
       return optixIgnoreIntersection();
     }

intern/cycles/kernel/device/optix/kernel.cu

@@ -18,6 +18,7 @@
 #include "kernel/integrator/intersect_shadow.h"
 #include "kernel/integrator/intersect_subsurface.h"
 #include "kernel/integrator/intersect_volume_stack.h"
+#include "kernel/integrator/intersect_dedicated_light.h"
 // clang-format on
 
 extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()
@@ -56,3 +57,11 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_volume_st
   integrator_intersect_volume_stack(nullptr, path_index);
 }
 
+extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_dedicated_light()
+{
+  const int global_index = optixGetLaunchIndex().x;
+  const int path_index = (kernel_params.path_index_array) ?
+                             kernel_params.path_index_array[global_index] :
+                             global_index;
+  integrator_intersect_dedicated_light(nullptr, path_index);
+}

intern/cycles/kernel/device/optix/kernel_osl.cu

@@ -9,6 +9,7 @@
 
 #include "kernel/bake/bake.h"
 #include "kernel/integrator/shade_background.h"
+#include "kernel/integrator/shade_dedicated_light.h"
 #include "kernel/integrator/shade_light.h"
 #include "kernel/integrator/shade_shadow.h"
 #include "kernel/integrator/shade_volume.h"
@@ -58,6 +59,15 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_shade_shadow()
   integrator_shade_shadow(nullptr, path_index, kernel_params.render_buffer);
 }
 
+extern "C" __global__ void __raygen__kernel_optix_integrator_shade_dedicated_light()
+{
+  const int global_index = optixGetLaunchIndex().x;
+  const int path_index = (kernel_params.path_index_array) ?
+                             kernel_params.path_index_array[global_index] :
+                             global_index;
+  integrator_shade_dedicated_light(nullptr, path_index, kernel_params.render_buffer);
+}
+
 extern "C" __global__ void __raygen__kernel_optix_shader_eval_displace()
 {
   KernelShaderEvalInput *const input = (KernelShaderEvalInput *)kernel_params.path_index_array;

intern/cycles/kernel/geom/object.h

@@ -288,7 +288,7 @@ ccl_device_inline float object_pass_id(KernelGlobals kg, int object)
   return kernel_data_fetch(objects, object).pass_id;
 }
 
-/* Lightgroup of lamp */
+/* Light-group of lamp. */
 
 ccl_device_inline int lamp_lightgroup(KernelGlobals kg, int lamp)
 {
@@ -298,7 +298,7 @@ ccl_device_inline int lamp_lightgroup(KernelGlobals kg, int lamp)
   return kernel_data_fetch(lights, lamp).lightgroup;
 }
 
-/* Lightgroup of object */
+/* Light-group of object. */
 
 ccl_device_inline int object_lightgroup(KernelGlobals kg, int object)
 {

intern/cycles/kernel/integrator/guiding.h

@@ -399,9 +399,8 @@ ccl_device_forceinline void guiding_record_background(KernelGlobals kg,
 #endif
 }
 
-/* Records the scattered contribution of a next event estimation
+/* Records direct lighting from either next event estimation or a dedicated BSDF
- * (i.e., a direct light estimate scattered at the current path vertex
+ * sampled shadow ray. */
- * towards the previous vertex). */
 ccl_device_forceinline void guiding_record_direct_light(KernelGlobals kg,
                                                         IntegratorShadowState state)
 {
@@ -414,7 +413,22 @@ ccl_device_forceinline void guiding_record_direct_light(KernelGlobals kg,
                                           INTEGRATOR_STATE(state, shadow_path, unlit_throughput));
 
     const float3 Lo_rgb = spectrum_to_rgb(Lo);
-    openpgl::cpp::AddScatteredContribution(state->shadow_path.path_segment, guiding_vec3f(Lo_rgb));
+
+    const float mis_weight = INTEGRATOR_STATE(state, shadow_path, guiding_mis_weight);
+
+    if (mis_weight == 0.0f) {
+      /* Scattered contribution of a next event estimation (i.e., a direct light estimate
+       * scattered at the current path vertex towards the previous vertex). */
+      openpgl::cpp::AddScatteredContribution(state->shadow_path.path_segment,
+                                             guiding_vec3f(Lo_rgb));
+    }
+    else {
+      /* Dedicated shadow ray for BSDF sampled ray direction.
+       * The mis weight was already folded into the throughput, so need to divide it out. */
+      openpgl::cpp::SetDirectContribution(state->shadow_path.path_segment,
+                                          guiding_vec3f(Lo_rgb / mis_weight));
+      openpgl::cpp::SetMiWeight(state->shadow_path.path_segment, mis_weight);
+    }
   }
 #endif
 }

intern/cycles/kernel/integrator/init_from_camera.h

@@ -32,16 +32,19 @@ ccl_device_inline void integrate_camera_sample(KernelGlobals kg,
                                     path_rng_3D(kg, rng_hash, sample, PRNG_LENS_TIME) :
                                     zero_float3();
 
+  /* We use x for time and y,z for lens because in practice with Sobol
+   * sampling this seems to give better convergence when an object is
+   * both motion blurred and out of focus, without significantly harming
+   * convergence for focal blur alone.  This is a little surprising,
+   * because one would expect using x,y for lens (the 2d part) would be
+   * best, since x,y are the best stratified.  Since it's not entirely
+   * clear why this is, this is probably worth revisiting at some point
+   * to investigate further. */
+  const float rand_time = rand_time_lens.x;
+  const float2 rand_lens = make_float2(rand_time_lens.y, rand_time_lens.z);
+
   /* Generate camera ray. */
-  camera_sample(kg,
+  camera_sample(kg, x, y, rand_filter, rand_time, rand_lens, ray);
-                x,
-                y,
-                rand_filter.x,
-                rand_filter.y,
-                rand_time_lens.y,
-                rand_time_lens.z,
-                rand_time_lens.x,
-                ray);
 }
 
 /* Return false to indicate that this pixel is finished.

intern/cycles/kernel/integrator/intersect_closest.h

@@ -353,6 +353,7 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg,
   ray.self.prim = last_isect_prim;
   ray.self.light_object = OBJECT_NONE;
   ray.self.light_prim = PRIM_NONE;
+  ray.self.light = LAMP_NONE;
   bool hit = scene_intersect(kg, &ray, visibility, &isect);
 
   /* TODO: remove this and do it in the various intersection functions instead. */

intern/cycles/kernel/integrator/intersect_dedicated_light.h

@@ -0,0 +1,205 @@
+/* SPDX-License-Identifier: Apache-2.0
+ * Copyright 2011-2023 Blender Foundation */
+
+#pragma once
+
+#include "kernel/bvh/bvh.h"
+#include "kernel/integrator/path_state.h"
+#include "kernel/integrator/shade_surface.h"
+#include "kernel/integrator/shadow_linking.h"
+#include "kernel/light/light.h"
+#include "kernel/sample/lcg.h"
+
+CCL_NAMESPACE_BEGIN
+
+#ifdef __SHADOW_LINKING__
+
+#  define SHADOW_LINK_MAX_INTERSECTION_COUNT 1024
+
+/* Intersect mesh objects.
+ *
+ * Returns the total number of emissive surfaces hit, and the intersection contains a random
+ * intersected emitter to which the dedicated shadow ray is to eb traced.
+ *
+ * NOTE: Sets the ray tmax to the maximum intersection distance (past which no lights are to be
+ * considered for shadow linking). */
+ccl_device int shadow_linking_pick_mesh_intersection(KernelGlobals kg,
+                                                     IntegratorState state,
+                                                     ccl_private Ray *ccl_restrict ray,
+                                                     const int object_receiver,
+                                                     ccl_private Intersection *ccl_restrict
+                                                         linked_isect,
+                                                     ccl_private uint *lcg_state,
+                                                     int num_hits)
+{
+  /* The tmin will be offset, so store its current value and restore later on, allowing a separate
+   * light intersection loop starting from the actual ray origin. */
+  const float old_tmin = ray->tmin;
+
+  const uint visibility = path_state_ray_visibility(state);
+
+  for (int i = 0; i < SHADOW_LINK_MAX_INTERSECTION_COUNT; i++) {
+    Intersection current_isect ccl_optional_struct_init;
+    current_isect.object = OBJECT_NONE;
+    current_isect.prim = PRIM_NONE;
+
+    const bool hit = scene_intersect(kg, ray, visibility, &current_isect);
+    if (!hit) {
+      break;
+    }
+
+    /* Only record primitives that potentially have emission.
+     * TODO: optimize with a dedicated ray visibility flag, which could then also be
+     * used once lights are in the BVH as geometry? */
+    const int shader = intersection_get_shader(kg, &current_isect);
+    const int shader_flags = kernel_data_fetch(shaders, shader).flags;
+    if (light_link_object_match(kg, object_receiver, current_isect.object) &&
+        (shader_flags & SD_HAS_EMISSION))
+    {
+      const uint64_t set_membership =
+          kernel_data_fetch(objects, current_isect.object).shadow_set_membership;
+      if (set_membership != LIGHT_LINK_MASK_ALL) {
+        ++num_hits;
+
+        if ((linked_isect->prim == PRIM_NONE) || (lcg_step_float(lcg_state) < 1.0f / num_hits)) {
+          *linked_isect = current_isect;
+        }
+      }
+    }
+
+    const uint blocker_set = kernel_data_fetch(objects, current_isect.object).blocker_shadow_set;
+    if (blocker_set == 0) {
+      /* Contribution from the lights past the default blocker is accumulated using the main path.
+       */
+      ray->tmax = current_isect.t;
+      break;
+    }
+
+    /* Move the ray forward. */
+    ray->tmin = intersection_t_offset(current_isect.t);
+  }
+
+  ray->tmin = old_tmin;
+
+  return num_hits;
+}
+
+/* Pick a light for tracing a shadow ray for the shadow linking.
+ * Picks a random light which is intersected by the given ray, and stores the intersection result.
+ * If no lights were hit false is returned.
+ *
+ * NOTE: Sets the ray tmax to the maximum intersection distance (past which no lights are to be
+ * considered for shadow linking). */
+ccl_device bool shadow_linking_pick_light_intersection(KernelGlobals kg,
+                                                       IntegratorState state,
+                                                       ccl_private Ray *ccl_restrict ray,
+                                                       ccl_private Intersection *ccl_restrict
+                                                           linked_isect)
+{
+  const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
+
+  const int last_type = INTEGRATOR_STATE(state, isect, type);
+
+  const int object_receiver = light_link_receiver_forward(kg, state);
+
+  uint lcg_state = lcg_state_init(INTEGRATOR_STATE(state, path, rng_hash),
+                                  INTEGRATOR_STATE(state, path, rng_offset),
+                                  INTEGRATOR_STATE(state, path, sample),
+                                  0x68bc21eb);
+
+  /* Indicate that no intersection has been picked yet. */
+  linked_isect->prim = PRIM_NONE;
+
+  int num_hits = 0;
+
+  // TODO: Only if there are emissive meshes in the scene?
+
+  // TODO: Only if the ray hits any light? As in, check that there is a light first, before
+  // tracing potentially expensive ray.
+
+  num_hits = shadow_linking_pick_mesh_intersection(
+      kg, state, ray, object_receiver, linked_isect, &lcg_state, num_hits);
+
+  num_hits = lights_intersect_shadow_linked(kg,
+                                            ray,
+                                            linked_isect,
+                                            ray->self.prim,
+                                            ray->self.object,
+                                            last_type,
+                                            path_flag,
+                                            object_receiver,
+                                            &lcg_state,
+                                            num_hits);
+
+  if (num_hits == 0) {
+    return false;
+  }
+
+  INTEGRATOR_STATE_WRITE(state, shadow_link, dedicated_light_weight) = num_hits;
+
+  return true;
+}
+
+/* Check whether a special shadow ray is needed to calculate direct light contribution which comes
+ * from emitters which are behind objects which are blocking light for the main path, but are
+ * excluded from blocking light via shadow linking.
+ *
+ * If a special ray is needed a blocked light kernel is scheduled and true is returned, otherwise
+ * false is returned. */
+ccl_device bool shadow_linking_intersect(KernelGlobals kg, IntegratorState state)
+{
+  /* Verify that the kernel is only scheduled if it is actually needed. */
+  kernel_assert(shadow_linking_scene_need_shadow_ray(kg, state));
+
+  /* Read ray from integrator state into local memory. */
+  Ray ray ccl_optional_struct_init;
+  integrator_state_read_ray(state, &ray);
+
+  ray.self.prim = INTEGRATOR_STATE(state, isect, prim);
+  ray.self.object = INTEGRATOR_STATE(state, isect, object);
+  ray.self.light_object = OBJECT_NONE;
+  ray.self.light_prim = PRIM_NONE;
+  ray.self.light = LAMP_NONE;
+
+  Intersection isect ccl_optional_struct_init;
+  if (!shadow_linking_pick_light_intersection(kg, state, &ray, &isect)) {
+    /* No light is hit, no need in the extra shadow ray for the direct light. */
+    return false;
+  }
+
+  /* Make a copy of primitives needed by the main path self-intersection check before writing the
+   * new intersection. Those primitives will be restored before the main path is returned to the
+   * intersect_closest state. */
+  shadow_linking_store_last_primitives(state);
+
+  /* Write intersection result into global integrator state memory, so that the
+   * shade_dedicated_light kernel can use it for calculation of the light sample. */
+  integrator_state_write_isect(state, &isect);
+
+  integrator_path_next(kg,
+                       state,
+                       DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT,
+                       DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT);
+
+  return true;
+}
+
+#endif /* __SHADOW_LINKING__ */
+
+ccl_device void integrator_intersect_dedicated_light(KernelGlobals kg, IntegratorState state)
+{
+  PROFILING_INIT(kg, PROFILING_INTERSECT_DEDICATED_LIGHT);
+
+#ifdef __SHADOW_LINKING__
+  if (shadow_linking_intersect(kg, state)) {
+    return;
+  }
+#else
+  kernel_assert(!"integrator_intersect_dedicated_light is not supposed to be scheduled");
+#endif
+
+  integrator_shade_surface_next_kernel<DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT>(kg,
+                                                                                           state);
+}
+
+CCL_NAMESPACE_END

intern/cycles/kernel/integrator/intersect_shadow.h

@@ -34,6 +34,9 @@ ccl_device bool integrate_intersect_shadow_opaque(KernelGlobals kg,
   Intersection isect;
   const bool opaque_hit = scene_intersect(kg, ray, visibility & opaque_mask, &isect);
 
+  /* Only record the number of hits if nothing was hit, so that the shadow shading kernel does not
+   * consider any intersections. There is no need to write anything to the state if the hit is
+   * opaque because in this case the path is terminated. */
   if (!opaque_hit) {
     INTEGRATOR_STATE_WRITE(state, shadow_path, num_hits) = 0;
   }
@@ -144,10 +147,7 @@ ccl_device void integrator_intersect_shadow(KernelGlobals kg, IntegratorShadowSt
   /* Read ray from integrator state into local memory. */
   Ray ray ccl_optional_struct_init;
   integrator_state_read_shadow_ray(state, &ray);
-  ray.self.object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, object);
+  integrator_state_read_shadow_ray_self(kg, state, &ray);
-  ray.self.prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, prim);
-  ray.self.light_object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 1, object);
-  ray.self.light_prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 1, prim);
   /* Compute visibility. */
   const uint visibility = integrate_intersect_shadow_visibility(kg, state);
 

intern/cycles/kernel/integrator/intersect_volume_stack.h

@@ -29,6 +29,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
   volume_ray.self.prim = INTEGRATOR_STATE(state, isect, prim);
   volume_ray.self.light_object = OBJECT_NONE;
   volume_ray.self.light_prim = PRIM_NONE;
+  volume_ray.self.light = LAMP_NONE;
   /* Store to avoid global fetches on every intersection step. */
   const uint volume_stack_size = kernel_data.volume_stack_size;
 
@@ -84,6 +85,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
   volume_ray.self.prim = PRIM_NONE;
   volume_ray.self.light_object = OBJECT_NONE;
   volume_ray.self.light_prim = PRIM_NONE;
+  volume_ray.self.light = LAMP_NONE;
 
   int stack_index = 0, enclosed_index = 0;
 

intern/cycles/kernel/integrator/megakernel.h

@@ -5,10 +5,12 @@
 
 #include "kernel/integrator/init_from_camera.h"
 #include "kernel/integrator/intersect_closest.h"
+#include "kernel/integrator/intersect_dedicated_light.h"
 #include "kernel/integrator/intersect_shadow.h"
 #include "kernel/integrator/intersect_subsurface.h"
 #include "kernel/integrator/intersect_volume_stack.h"
 #include "kernel/integrator/shade_background.h"
+#include "kernel/integrator/shade_dedicated_light.h"
 #include "kernel/integrator/shade_light.h"
 #include "kernel/integrator/shade_shadow.h"
 #include "kernel/integrator/shade_surface.h"
@@ -83,12 +85,18 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
         case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
           integrator_shade_light(kg, state, render_buffer);
           break;
+        case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT:
+          integrator_shade_dedicated_light(kg, state, render_buffer);
+          break;
         case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
           integrator_intersect_subsurface(kg, state);
           break;
         case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
           integrator_intersect_volume_stack(kg, state);
           break;
+        case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
+          integrator_intersect_dedicated_light(kg, state);
+          break;
         default:
           kernel_assert(0);
           break;

intern/cycles/kernel/integrator/mnee.h

@@ -408,6 +408,7 @@ ccl_device_forceinline bool mnee_newton_solver(KernelGlobals kg,
   Ray projection_ray;
   projection_ray.self.light_object = OBJECT_NONE;
   projection_ray.self.light_prim = PRIM_NONE;
+  projection_ray.self.light = LAMP_NONE;
   projection_ray.dP = differential_make_compact(sd->dP);
   projection_ray.dD = differential_zero_compact();
   projection_ray.tmin = 0.0f;
@@ -482,11 +483,7 @@ ccl_device_forceinline bool mnee_newton_solver(KernelGlobals kg,
         if (!hit)
           break;
 
-        int hit_object = (projection_isect.object == OBJECT_NONE) ?
+        if (projection_isect.object == mv.object) {
-                             kernel_data_fetch(prim_object, projection_isect.prim) :
-                             projection_isect.object;
-
-        if (hit_object == mv.object) {
           projection_success = true;
           break;
         }
@@ -776,9 +773,8 @@ ccl_device_forceinline bool mnee_path_contribution(KernelGlobals kg,
   /* Initialize throughput and evaluate receiver bsdf * |n.wo|. */
   surface_shader_bsdf_eval(kg, state, sd, wo, throughput, ls->shader);
 
-  /* Update light sample with new position / direct.ion
+  /* Update light sample with new position / direction and keep pdf in vertex area measure. */
-   * and keep pdf in vertex area measure */
+  light_sample_update(kg, ls, vertices[vertex_count - 1].p);
-  light_sample_update_position(kg, ls, vertices[vertex_count - 1].p);
 
   /* Save state path bounce info in case a light path node is used in the refractive interface or
    * light shader graph. */
@@ -826,6 +822,7 @@ ccl_device_forceinline bool mnee_path_contribution(KernelGlobals kg,
   Ray probe_ray;
   probe_ray.self.light_object = ls->object;
   probe_ray.self.light_prim = ls->prim;
+  probe_ray.self.light = ls->lamp;
   probe_ray.tmin = 0.0f;
   probe_ray.dP = differential_make_compact(sd->dP);
   probe_ray.dD = differential_zero_compact();
@@ -926,6 +923,7 @@ ccl_device_forceinline int kernel_path_mnee_sample(KernelGlobals kg,
   probe_ray.self.prim = sd->prim;
   probe_ray.self.light_object = ls->object;
   probe_ray.self.light_prim = ls->prim;
+  probe_ray.self.light = ls->lamp;
   probe_ray.P = sd->P;
   probe_ray.tmin = 0.0f;
   if (ls->t == FLT_MAX) {

intern/cycles/kernel/integrator/path_state.h

@@ -89,6 +89,12 @@ ccl_device_inline void path_state_init_integrator(KernelGlobals kg,
     INTEGRATOR_STATE_WRITE(state, path, denoising_feature_throughput) = one_spectrum();
   }
 #endif
+
+#ifdef __LIGHT_LINKING__
+  if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING) {
+    INTEGRATOR_STATE_WRITE(state, path, mis_ray_object) = OBJECT_NONE;
+  }
+#endif
 }
 
 ccl_device_inline void path_state_next(KernelGlobals kg,

intern/cycles/kernel/integrator/shade_background.h

@@ -22,14 +22,8 @@ ccl_device Spectrum integrator_eval_background_shader(KernelGlobals kg,
   const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
 
   /* Use visibility flag to skip lights. */
-  if (shader & SHADER_EXCLUDE_ANY) {
+  if (!is_light_shader_visible_to_path(shader, path_flag)) {
-    if (((shader & SHADER_EXCLUDE_DIFFUSE) && (path_flag & PATH_RAY_DIFFUSE)) ||
+    return zero_spectrum();
-        ((shader & SHADER_EXCLUDE_GLOSSY) && ((path_flag & (PATH_RAY_GLOSSY | PATH_RAY_REFLECT)) ==
-                                              (PATH_RAY_GLOSSY | PATH_RAY_REFLECT))) ||
-        ((shader & SHADER_EXCLUDE_TRANSMIT) && (path_flag & PATH_RAY_TRANSMIT)) ||
-        ((shader & SHADER_EXCLUDE_CAMERA) && (path_flag & PATH_RAY_CAMERA)) ||
-        ((shader & SHADER_EXCLUDE_SCATTER) && (path_flag & PATH_RAY_VOLUME_SCATTER)))
-      return zero_spectrum();
   }
 
   /* Use fast constant background color if available. */
@@ -140,16 +134,21 @@ ccl_device_inline void integrate_distant_lights(KernelGlobals kg,
       /* Use visibility flag to skip lights. */
 #ifdef __PASSES__
       const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
+      if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
+        continue;
+      }
+#endif
 
-      if (ls.shader & SHADER_EXCLUDE_ANY) {
+#ifdef __LIGHT_LINKING__
-        if (((ls.shader & SHADER_EXCLUDE_DIFFUSE) && (path_flag & PATH_RAY_DIFFUSE)) ||
+      if (!light_link_light_match(kg, light_link_receiver_forward(kg, state), lamp) &&
-            ((ls.shader & SHADER_EXCLUDE_GLOSSY) &&
+          !(path_flag & PATH_RAY_CAMERA))
-             ((path_flag & (PATH_RAY_GLOSSY | PATH_RAY_REFLECT)) ==
+      {
-              (PATH_RAY_GLOSSY | PATH_RAY_REFLECT))) ||
+        continue;
-            ((ls.shader & SHADER_EXCLUDE_TRANSMIT) && (path_flag & PATH_RAY_TRANSMIT)) ||
+      }
-            ((ls.shader & SHADER_EXCLUDE_CAMERA) && (path_flag & PATH_RAY_CAMERA)) ||
+#endif
-            ((ls.shader & SHADER_EXCLUDE_SCATTER) && (path_flag & PATH_RAY_VOLUME_SCATTER)))
+#ifdef __SHADOW_LINKING__
-          continue;
+      if (kernel_data_fetch(lights, lamp).shadow_set_membership != LIGHT_LINK_MASK_ALL) {
+        continue;
       }
 #endif
 

intern/cycles/kernel/integrator/shade_dedicated_light.h

@@ -0,0 +1,274 @@
+/* SPDX-License-Identifier: Apache-2.0
+ * Copyright 2011-2023 Blender Foundation */
+
+#pragma once
+
+#include "kernel/integrator/path_state.h"
+
+#include "kernel/light/distant.h"
+#include "kernel/light/light.h"
+#include "kernel/light/sample.h"
+
+#include "kernel/integrator/shade_surface.h"
+
+CCL_NAMESPACE_BEGIN
+
+#ifdef __SHADOW_LINKING__
+
+ccl_device_inline bool shadow_linking_light_sample_from_intersection(
+    KernelGlobals kg,
+    ccl_private const Intersection &ccl_restrict isect,
+    ccl_private const Ray &ccl_restrict ray,
+    ccl_private LightSample *ccl_restrict ls)
+{
+  const int lamp = isect.prim;
+
+  const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
+  const LightType type = LightType(klight->type);
+
+  if (type == LIGHT_DISTANT) {
+    return distant_light_sample_from_intersection(kg, ray.D, lamp, ls);
+  }
+
+  return light_sample_from_intersection(kg, &isect, ray.P, ray.D, ls);
+}
+
+ccl_device_inline float shadow_linking_light_sample_mis_weight(KernelGlobals kg,
+                                                               IntegratorState state,
+                                                               const uint32_t path_flag,
+                                                               const ccl_private LightSample *ls,
+                                                               const float3 P)
+{
+  if (ls->type == LIGHT_DISTANT) {
+    return light_sample_mis_weight_forward_distant(kg, state, path_flag, ls);
+  }
+
+  return light_sample_mis_weight_forward_lamp(kg, state, path_flag, ls, P);
+}
+
+/* Setup ray for the shadow path.
+ * Expects that the current state of the ray is the one calculated by the surface bounce, and the
+ * intersection corresponds to a point on an emitter. */
+ccl_device void shadow_linking_setup_ray_from_intersection(
+    IntegratorState state,
+    ccl_private Ray *ccl_restrict ray,
+    ccl_private const Intersection *ccl_restrict isect)
+{
+  /* The ray->tmin follows the value configured at the surface bounce.
+   * it is the same for the continued main path and for this shadow ray. There is no need to push
+   * it forward here. */
+
+  ray->tmax = isect->t;
+
+  /* Use the same self intersection primitives as the main path.
+   * Those are copied to the dedicated storage from the main intersection after the surface bounce,
+   * but before the main intersection is re-used to find light to trace a ray to. */
+  ray->self.object = INTEGRATOR_STATE(state, shadow_link, last_isect_object);
+  ray->self.prim = INTEGRATOR_STATE(state, shadow_link, last_isect_prim);
+
+  if (isect->type == PRIMITIVE_LAMP) {
+    ray->self.light_object = OBJECT_NONE;
+    ray->self.light_prim = PRIM_NONE;
+    ray->self.light = isect->prim;
+  }
+  else {
+    ray->self.light_object = isect->object;
+    ray->self.light_prim = isect->prim;
+    ray->self.light = LAMP_NONE;
+  }
+}
+
+ccl_device bool shadow_linking_shade_light(KernelGlobals kg,
+                                           IntegratorState state,
+                                           ccl_private Ray &ccl_restrict ray,
+                                           ccl_private Intersection &ccl_restrict isect,
+                                           ccl_private ShaderData *emission_sd,
+                                           ccl_private Spectrum &ccl_restrict bsdf_spectrum,
+                                           ccl_private float &mis_weight,
+                                           ccl_private int &ccl_restrict light_group)
+{
+  LightSample ls ccl_optional_struct_init;
+  const bool use_light_sample = shadow_linking_light_sample_from_intersection(kg, isect, ray, &ls);
+  if (!use_light_sample) {
+    /* No light to be sampled, so no direct light contribution either. */
+    return false;
+  }
+
+  const Spectrum light_eval = light_sample_shader_eval(kg, state, emission_sd, &ls, ray.time);
+  if (is_zero(light_eval)) {
+    return false;
+  }
+
+  const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
+  if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
+    return false;
+  }
+
+  /* MIS weighting. */
+  if (!(path_flag & PATH_RAY_MIS_SKIP)) {
+    mis_weight = shadow_linking_light_sample_mis_weight(kg, state, path_flag, &ls, ray.P);
+  }
+
+  bsdf_spectrum = light_eval * mis_weight *
+                  INTEGRATOR_STATE(state, shadow_link, dedicated_light_weight);
+
+  // TODO(: De-duplicate with the shade_surface.
+  // Possibly by ensuring ls->group is always assigned properly.
+  light_group = ls.type != LIGHT_BACKGROUND ? ls.group : kernel_data.background.lightgroup;
+
+  return true;
+}
+
+ccl_device bool shadow_linking_shade_surface_emission(KernelGlobals kg,
+                                                      IntegratorState state,
+                                                      ccl_private Ray &ccl_restrict ray,
+                                                      ccl_private Intersection &ccl_restrict isect,
+                                                      ccl_private ShaderData *emission_sd,
+                                                      ccl_global float *ccl_restrict render_buffer,
+                                                      ccl_private Spectrum &ccl_restrict
+                                                          bsdf_spectrum,
+                                                      ccl_private float &mis_weight,
+                                                      ccl_private int &ccl_restrict light_group)
+{
+  const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
+
+  integrate_surface_shader_setup(kg, state, emission_sd);
+
+#  ifdef __VOLUME__
+  if (emission_sd->flag & SD_HAS_ONLY_VOLUME) {
+    return false;
+  }
+#  endif
+
+  surface_shader_eval<KERNEL_FEATURE_NODE_MASK_SURFACE_LIGHT>(
+      kg, state, emission_sd, render_buffer, path_flag);
+  surface_shader_prepare_closures(kg, state, emission_sd, path_flag);
+
+  if ((emission_sd->flag & SD_EMISSION) == 0) {
+    return false;
+  }
+
+  const Spectrum L = surface_shader_emission(emission_sd);
+
+  const bool has_mis = !(path_flag & PATH_RAY_MIS_SKIP) &&
+                       (emission_sd->flag &
+                        ((emission_sd->flag & SD_BACKFACING) ? SD_MIS_BACK : SD_MIS_FRONT));
+
+#  ifdef __HAIR__
+  if (has_mis && (emission_sd->type & PRIMITIVE_TRIANGLE))
+#  else
+  if (has_mis)
+#  endif
+  {
+    mis_weight = light_sample_mis_weight_forward_surface(kg, state, path_flag, emission_sd);
+  }
+
+  bsdf_spectrum = L * mis_weight * INTEGRATOR_STATE(state, shadow_link, dedicated_light_weight);
+  light_group = object_lightgroup(kg, emission_sd->object);
+
+  return true;
+}
+
+ccl_device void shadow_linking_shade(KernelGlobals kg,
+                                     IntegratorState state,
+                                     ccl_global float *ccl_restrict render_buffer)
+{
+  /* Read intersection from integrator state into local memory. */
+  Intersection isect ccl_optional_struct_init;
+  integrator_state_read_isect(state, &isect);
+
+  /* Read ray from integrator state into local memory. */
+  Ray ray ccl_optional_struct_init;
+  integrator_state_read_ray(state, &ray);
+
+  ShaderDataCausticsStorage emission_sd_storage;
+  ccl_private ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
+
+  Spectrum bsdf_spectrum;
+  float mis_weight = 1.0f;
+  int light_group = LIGHTGROUP_NONE;
+
+  if (isect.type == PRIMITIVE_LAMP) {
+    if (!shadow_linking_shade_light(
+            kg, state, ray, isect, emission_sd, bsdf_spectrum, mis_weight, light_group))
+    {
+      return;
+    }
+  }
+  else {
+    if (!shadow_linking_shade_surface_emission(kg,
+                                               state,
+                                               ray,
+                                               isect,
+                                               emission_sd,
+                                               render_buffer,
+                                               bsdf_spectrum,
+                                               mis_weight,
+                                               light_group))
+    {
+      return;
+    }
+  }
+
+  if (is_zero(bsdf_spectrum)) {
+    return;
+  }
+
+  shadow_linking_setup_ray_from_intersection(state, &ray, &isect);
+
+  /* Branch off shadow kernel. */
+  IntegratorShadowState shadow_state = integrate_direct_light_shadow_init_common(
+      kg, state, &ray, bsdf_spectrum, 0, light_group);
+
+  /* The light is accumulated from the shade_surface kernel, which will make the clamping decision
+   * based on the actual value of the bounce. For the dedicated shadow ray we want to follow the
+   * main path clamping rules, which subtracts one from the bounds before accumulation. */
+  INTEGRATOR_STATE_WRITE(
+      shadow_state, shadow_path, bounce) = INTEGRATOR_STATE(shadow_state, shadow_path, bounce) - 1;
+
+  /* No need to update the volume stack as the surface bounce already performed enter-exit check.
+   */
+
+  const uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
+
+  if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
+    /* The diffuse and glossy pass weights are written into the main path as part of the path
+     * configuration at a surface bounce. */
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = INTEGRATOR_STATE(
+        state, path, pass_diffuse_weight);
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = INTEGRATOR_STATE(
+        state, path, pass_glossy_weight);
+  }
+
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, flag) = shadow_flag;
+
+#  ifdef __PATH_GUIDING__
+  if (kernel_data.integrator.train_guiding) {
+    guiding_record_light_surface_segment(kg, state, &isect);
+    INTEGRATOR_STATE(shadow_state, shadow_path, guiding_mis_weight) = mis_weight;
+  }
+#  endif
+}
+
+#endif /* __SHADOW_LINKING__ */
+
+ccl_device void integrator_shade_dedicated_light(KernelGlobals kg,
+                                                 IntegratorState state,
+                                                 ccl_global float *ccl_restrict render_buffer)
+{
+  PROFILING_INIT(kg, PROFILING_SHADE_DEDICATED_LIGHT);
+
+#ifdef __SHADOW_LINKING__
+  shadow_linking_shade(kg, state, render_buffer);
+
+  /* Restore self-intersection check primitives in the main state before returning to the
+   * intersect_closest() state. */
+  shadow_linking_restore_last_primitives(state);
+#else
+  kernel_assert(!"integrator_intersect_dedicated_light is not supposed to be scheduled");
+#endif
+
+  integrator_shade_surface_next_kernel<DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT>(kg, state);
+}
+
+CCL_NAMESPACE_END

intern/cycles/kernel/integrator/shade_light.h

@@ -37,15 +37,8 @@ ccl_device_inline void integrate_light(KernelGlobals kg,
   /* Use visibility flag to skip lights. */
 #ifdef __PASSES__
   const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
-
+  if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
-  if (ls.shader & SHADER_EXCLUDE_ANY) {
+    return;
-    if (((ls.shader & SHADER_EXCLUDE_DIFFUSE) && (path_flag & PATH_RAY_DIFFUSE)) ||
-        ((ls.shader & SHADER_EXCLUDE_GLOSSY) &&
-         ((path_flag & (PATH_RAY_GLOSSY | PATH_RAY_REFLECT)) ==
-          (PATH_RAY_GLOSSY | PATH_RAY_REFLECT))) ||
-        ((ls.shader & SHADER_EXCLUDE_TRANSMIT) && (path_flag & PATH_RAY_TRANSMIT)) ||
-        ((ls.shader & SHADER_EXCLUDE_SCATTER) && (path_flag & PATH_RAY_VOLUME_SCATTER)))
-      return;
   }
 #endif
 

intern/cycles/kernel/integrator/shade_shadow.h

@@ -75,6 +75,7 @@ ccl_device_inline void integrate_transparent_volume_shadow(KernelGlobals kg,
   ray.self.prim = PRIM_NONE;
   ray.self.light_object = OBJECT_NONE;
   ray.self.light_prim = PRIM_NONE;
+  ray.self.light = LAMP_NONE;
   /* Modify ray position and length to match current segment. */
   ray.tmin = (hit == 0) ? ray.tmin : INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit - 1, t);
   ray.tmax = (hit < num_recorded_hits) ? INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit, t) :

intern/cycles/kernel/integrator/shade_surface.h

@@ -3,20 +3,22 @@
 
 #pragma once
 
+#include "kernel/integrator/path_state.h"
+#include "kernel/integrator/surface_shader.h"
+
 #include "kernel/film/data_passes.h"
 #include "kernel/film/denoising_passes.h"
 #include "kernel/film/light_passes.h"
 
+#include "kernel/light/sample.h"
+
 #include "kernel/integrator/mnee.h"
 
 #include "kernel/integrator/guiding.h"
-#include "kernel/integrator/path_state.h"
+#include "kernel/integrator/shadow_linking.h"
 #include "kernel/integrator/subsurface.h"
-#include "kernel/integrator/surface_shader.h"
 #include "kernel/integrator/volume_stack.h"
 
-#include "kernel/light/sample.h"
-
 CCL_NAMESPACE_BEGIN
 
 ccl_device_forceinline void integrate_surface_shader_setup(KernelGlobals kg,
@@ -113,6 +115,26 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg,
 {
   const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
 
+#ifdef __LIGHT_LINKING__
+  if (!light_link_object_match(kg, light_link_receiver_forward(kg, state), sd->object) &&
+      !(path_flag & PATH_RAY_CAMERA))
+  {
+    return;
+  }
+#endif
+
+#ifdef __SHADOW_LINKING__
+  /* Indirect emission of shadow-linked emissive surfaces is done via shadow rays to dedicated
+   * light sources. */
+  if (kernel_data.kernel_features & KERNEL_FEATURE_SHADOW_LINKING) {
+    if (!(path_flag & PATH_RAY_CAMERA) &&
+        kernel_data_fetch(objects, sd->object).shadow_set_membership != LIGHT_LINK_MASK_ALL)
+    {
+      return;
+    }
+  }
+#endif
+
   /* Evaluate emissive closure. */
   Spectrum L = surface_shader_emission(sd);
   float mis_weight = 1.0f;
@@ -134,6 +156,84 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg,
       kg, state, L, mis_weight, render_buffer, object_lightgroup(kg, sd->object));
 }
 
+/* Branch off a shadow path and initialize common part of it.
+ * THe common is between the surface shading and configuration of a special shadow ray for the
+ * shadow linking. */
+ccl_device_inline IntegratorShadowState
+integrate_direct_light_shadow_init_common(KernelGlobals kg,
+                                          IntegratorState state,
+                                          ccl_private const Ray *ccl_restrict ray,
+                                          const Spectrum bsdf_spectrum,
+                                          const int light_group,
+                                          const int mnee_vertex_count)
+{
+
+  /* Branch off shadow kernel. */
+  IntegratorShadowState shadow_state = integrator_shadow_path_init(
+      kg, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, false);
+
+  /* Copy volume stack and enter/exit volume. */
+  integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
+
+  /* Write shadow ray and associated state to global memory. */
+  integrator_state_write_shadow_ray(shadow_state, ray);
+  integrator_state_write_shadow_ray_self(kg, shadow_state, ray);
+
+  /* Copy state from main path to shadow path. */
+  const Spectrum unlit_throughput = INTEGRATOR_STATE(state, path, throughput);
+  const Spectrum throughput = unlit_throughput * bsdf_spectrum;
+
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
+      state, path, render_pixel_index);
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_offset) = INTEGRATOR_STATE(
+      state, path, rng_offset);
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_hash) = INTEGRATOR_STATE(
+      state, path, rng_hash);
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, sample) = INTEGRATOR_STATE(
+      state, path, sample);
+
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transparent_bounce) = INTEGRATOR_STATE(
+      state, path, transparent_bounce);
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, glossy_bounce) = INTEGRATOR_STATE(
+      state, path, glossy_bounce);
+
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, throughput) = throughput;
+
+#ifdef __MNEE__
+  if (mnee_vertex_count > 0) {
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transmission_bounce) =
+        INTEGRATOR_STATE(state, path, transmission_bounce) + mnee_vertex_count - 1;
+    INTEGRATOR_STATE_WRITE(shadow_state,
+                           shadow_path,
+                           diffuse_bounce) = INTEGRATOR_STATE(state, path, diffuse_bounce) + 1;
+    INTEGRATOR_STATE_WRITE(shadow_state,
+                           shadow_path,
+                           bounce) = INTEGRATOR_STATE(state, path, bounce) + mnee_vertex_count;
+  }
+  else
+#endif
+  {
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transmission_bounce) = INTEGRATOR_STATE(
+        state, path, transmission_bounce);
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, diffuse_bounce) = INTEGRATOR_STATE(
+        state, path, diffuse_bounce);
+    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, bounce) = INTEGRATOR_STATE(
+        state, path, bounce);
+  }
+
+  /* Write Light-group, +1 as light-group is int but we need to encode into a uint8_t. */
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, lightgroup) = light_group;
+
+#ifdef __PATH_GUIDING__
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, unlit_throughput) = unlit_throughput;
+  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, path_segment) = INTEGRATOR_STATE(
+      state, guiding, path_segment);
+  INTEGRATOR_STATE(shadow_state, shadow_path, guiding_mis_weight) = 0.0f;
+#endif
+
+  return shadow_state;
+}
+
 /* Path tracing: sample point on light and evaluate light shader, then
  * queue shadow ray to be traced. */
 template<uint node_feature_mask>
@@ -156,12 +256,11 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
 
     if (!light_sample_from_position(kg,
                                     rng_state,
-                                    rand_light.z,
+                                    rand_light,
-                                    rand_light.x,
-                                    rand_light.y,
                                     sd->time,
                                     sd->P,
                                     sd->N,
+                                    light_link_receiver_nee(kg, sd),
                                     sd->flag,
                                     bounce,
                                     path_flag,
@@ -187,8 +286,8 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
 
   const bool is_transmission = dot(ls.D, sd->N) < 0.0f;
 
-#ifdef __MNEE__
   int mnee_vertex_count = 0;
+#ifdef __MNEE__
   IF_KERNEL_FEATURE(MNEE)
   {
     if (ls.lamp != LAMP_NONE) {
@@ -241,12 +340,18 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
     light_sample_to_surface_shadow_ray(kg, sd, &ls, &ray);
   }
 
-  /* Branch off shadow kernel. */
+  if (ray.self.object != OBJECT_NONE) {
-  IntegratorShadowState shadow_state = integrator_shadow_path_init(
+    ray.P = integrate_surface_ray_offset(kg, sd, ray.P, ray.D);
-      kg, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, false);
+  }
 
-  /* Copy volume stack and enter/exit volume. */
+  /* Branch off shadow kernel. */
-  integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
+
+  // TODO(: De-duplicate with the shade_Dedicated_light.
+  // Possibly by ensuring ls->group is always assigned properly.
+  const int light_group = ls.type != LIGHT_BACKGROUND ? ls.group :
+                                                        kernel_data.background.lightgroup;
+  IntegratorShadowState shadow_state = integrate_direct_light_shadow_init_common(
+      kg, state, &ray, bsdf_eval_sum(&bsdf_eval), mnee_vertex_count, light_group);
 
   if (is_transmission) {
 #ifdef __VOLUME__
@@ -254,22 +359,7 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
 #endif
   }
 
-  if (ray.self.object != OBJECT_NONE) {
-    ray.P = integrate_surface_ray_offset(kg, sd, ray.P, ray.D);
-  }
-
-  /* Write shadow ray and associated state to global memory. */
-  integrator_state_write_shadow_ray(shadow_state, &ray);
-  // Save memory by storing the light and object indices in the shadow_isect
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
-
-  /* Copy state from main path to shadow path. */
   uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
-  const Spectrum unlit_throughput = INTEGRATOR_STATE(state, path, throughput);
-  const Spectrum throughput = unlit_throughput * bsdf_eval_sum(&bsdf_eval);
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
     PackedSpectrum pass_diffuse_weight;
@@ -291,55 +381,7 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
     INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = pass_glossy_weight;
   }
 
-  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
-      state, path, render_pixel_index);
-  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_offset) = INTEGRATOR_STATE(
-      state, path, rng_offset);
-  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_hash) = INTEGRATOR_STATE(
-      state, path, rng_hash);
-  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, sample) = INTEGRATOR_STATE(
-      state, path, sample);
   INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, flag) = shadow_flag;
-
-  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transparent_bounce) = INTEGRATOR_STATE(
-      state, path, transparent_bounce);
-  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, glossy_bounce) = INTEGRATOR_STATE(
-      state, path, glossy_bounce);
-
-#ifdef __MNEE__
-  if (mnee_vertex_count > 0) {
-    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transmission_bounce) =
-        INTEGRATOR_STATE(state, path, transmission_bounce) + mnee_vertex_count - 1;
-    INTEGRATOR_STATE_WRITE(shadow_state,
-                           shadow_path,
-                           diffuse_bounce) = INTEGRATOR_STATE(state, path, diffuse_bounce) + 1;
-    INTEGRATOR_STATE_WRITE(shadow_state,
-                           shadow_path,
-                           bounce) = INTEGRATOR_STATE(state, path, bounce) + mnee_vertex_count;
-  }
-  else
-#endif
-  {
-    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transmission_bounce) = INTEGRATOR_STATE(
-        state, path, transmission_bounce);
-    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, diffuse_bounce) = INTEGRATOR_STATE(
-        state, path, diffuse_bounce);
-    INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, bounce) = INTEGRATOR_STATE(
-        state, path, bounce);
-  }
-
-  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, throughput) = throughput;
-
-  /* Write Lightgroup, +1 as lightgroup is int but we need to encode into a uint8_t. */
-  INTEGRATOR_STATE_WRITE(
-      shadow_state, shadow_path, lightgroup) = (ls.type != LIGHT_BACKGROUND) ?
-                                                   ls.group + 1 :
-                                                   kernel_data.background.lightgroup + 1;
-#ifdef __PATH_GUIDING__
-  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, unlit_throughput) = unlit_throughput;
-  INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, path_segment) = INTEGRATOR_STATE(
-      state, guiding, path_segment);
-#endif
 }
 
 /* Path tracing: bounce off or through surface with new direction. */
@@ -453,6 +495,11 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
     INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = fminf(
         unguided_bsdf_pdf, INTEGRATOR_STATE(state, path, min_ray_pdf));
   }
+#ifdef __LIGHT_LINKING__
+  if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING) {
+    INTEGRATOR_STATE_WRITE(state, path, mis_ray_object) = sd->object;
+  }
+#endif
 
   path_state_next(kg, state, label, sd->flag);
 
@@ -529,7 +576,7 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
 
   float3 ao_D;
   float ao_pdf;
-  sample_cos_hemisphere(ao_N, rand_bsdf.x, rand_bsdf.y, &ao_D, &ao_pdf);
+  sample_cos_hemisphere(ao_N, rand_bsdf, &ao_D, &ao_pdf);
 
   bool skip_self = true;
 
@@ -546,6 +593,7 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
   ray.self.prim = (skip_self) ? sd->prim : PRIM_NONE;
   ray.self.light_object = OBJECT_NONE;
   ray.self.light_prim = PRIM_NONE;
+  ray.self.light = LAMP_NONE;
   ray.dP = differential_zero_compact();
   ray.dD = differential_zero_compact();
 
@@ -558,10 +606,7 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
 
   /* Write shadow ray and associated state to global memory. */
   integrator_state_write_shadow_ray(shadow_state, &ray);
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
+  integrator_state_write_shadow_ray_self(kg, shadow_state, &ray);
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
 
   /* Copy state from main path to shadow path. */
   const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);
@@ -590,9 +635,9 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
 #endif /* defined(__AO__) */
 
 template<uint node_feature_mask>
-ccl_device bool integrate_surface(KernelGlobals kg,
+ccl_device int integrate_surface(KernelGlobals kg,
-                                  IntegratorState state,
+                                 IntegratorState state,
-                                  ccl_global float *ccl_restrict render_buffer)
+                                 ccl_global float *ccl_restrict render_buffer)
 
 {
   PROFILING_INIT_FOR_SHADER(kg, PROFILING_SHADE_SURFACE_SETUP);
@@ -645,7 +690,7 @@ ccl_device bool integrate_surface(KernelGlobals kg,
 
       /* Evaluate holdout. */
       if (!integrate_surface_holdout(kg, state, &sd, render_buffer)) {
-        return false;
+        return LABEL_NONE;
       }
 
       /* Write emission. */
@@ -659,7 +704,7 @@ ccl_device bool integrate_surface(KernelGlobals kg,
        *
        * Also ensure we don't do it twice for SSS at both the entry and exit point. */
       if (integrate_surface_terminate(state, path_flag)) {
-        return false;
+        return LABEL_NONE;
       }
 
       /* Write render passes. */
@@ -699,7 +744,7 @@ ccl_device bool integrate_surface(KernelGlobals kg,
   }
   else {
     if (integrate_surface_terminate(state, path_flag)) {
-      return false;
+      return LABEL_NONE;
     }
 
     PROFILING_EVENT(PROFILING_SHADE_SURFACE_INDIRECT_LIGHT);
@@ -712,7 +757,20 @@ ccl_device bool integrate_surface(KernelGlobals kg,
   }
 #endif
 
-  return continue_path_label != 0;
+  return continue_path_label;
+}
+
+template<DeviceKernel current_kernel>
+ccl_device_forceinline void integrator_shade_surface_next_kernel(KernelGlobals kg,
+                                                                 IntegratorState state)
+{
+  if (INTEGRATOR_STATE(state, path, flag) & PATH_RAY_SUBSURFACE) {
+    integrator_path_next(kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE);
+  }
+  else {
+    kernel_assert(INTEGRATOR_STATE(state, ray, tmax) != 0.0f);
+    integrator_path_next(kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
+  }
 }
 
 template<uint node_feature_mask = KERNEL_FEATURE_NODE_MASK_SURFACE & ~KERNEL_FEATURE_NODE_RAYTRACE,
@@ -721,19 +779,23 @@ ccl_device_forceinline void integrator_shade_surface(KernelGlobals kg,
                                                      IntegratorState state,
                                                      ccl_global float *ccl_restrict render_buffer)
 {
-  if (integrate_surface<node_feature_mask>(kg, state, render_buffer)) {
+  const int continue_path_label = integrate_surface<node_feature_mask>(kg, state, render_buffer);
-    if (INTEGRATOR_STATE(state, path, flag) & PATH_RAY_SUBSURFACE) {
+  if (continue_path_label == LABEL_NONE) {
-      integrator_path_next(
-          kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE);
-    }
-    else {
-      kernel_assert(INTEGRATOR_STATE(state, ray, tmax) != 0.0f);
-      integrator_path_next(kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
-    }
-  }
-  else {
     integrator_path_terminate(kg, state, current_kernel);
+    return;
   }
+
+#ifdef __SHADOW_LINKING__
+  /* No need to cast shadow linking rays at a transparent bounce: the lights will be accumulated
+   * via the main path in this case. */
+  if ((continue_path_label & LABEL_TRANSPARENT) == 0) {
+    if (shadow_linking_schedule_intersection_kernel<current_kernel>(kg, state)) {
+      return;
+    }
+  }
+#endif
+
+  integrator_shade_surface_next_kernel<current_kernel>(kg, state);
 }
 
 ccl_device_forceinline void integrator_shade_surface_raytrace(

intern/cycles/kernel/integrator/shade_volume.h

@@ -10,6 +10,7 @@
 #include "kernel/integrator/guiding.h"
 #include "kernel/integrator/intersect_closest.h"
 #include "kernel/integrator/path_state.h"
+#include "kernel/integrator/shadow_linking.h"
 #include "kernel/integrator/volume_shader.h"
 #include "kernel/integrator/volume_stack.h"
 
@@ -673,8 +674,10 @@ ccl_device_forceinline void volume_integrate_heterogeneous(
 
   /* Write accumulated emission. */
   if (!is_zero(accum_emission)) {
-    film_write_volume_emission(
+    if (light_link_object_match(kg, light_link_receiver_forward(kg, state), sd->object)) {
-        kg, state, accum_emission, render_buffer, object_lightgroup(kg, sd->object));
+      film_write_volume_emission(
+          kg, state, accum_emission, render_buffer, object_lightgroup(kg, sd->object));
+    }
   }
 
 #  ifdef __DENOISING_FEATURES__
@@ -707,13 +710,12 @@ ccl_device_forceinline bool integrate_volume_equiangular_sample_light(
 
   LightSample ls ccl_optional_struct_init;
   if (!light_sample_from_volume_segment(kg,
-                                        rand_light.z,
+                                        rand_light,
-                                        rand_light.x,
-                                        rand_light.y,
                                         sd->time,
                                         sd->P,
                                         ray->D,
                                         ray->tmax - ray->tmin,
+                                        light_link_receiver_nee(kg, sd),
                                         bounce,
                                         path_flag,
                                         &ls))
@@ -772,12 +774,11 @@ ccl_device_forceinline void integrate_volume_direct_light(
 
     if (!light_sample_from_position(kg,
                                     rng_state,
-                                    rand_light.z,
+                                    rand_light,
-                                    rand_light.x,
-                                    rand_light.y,
                                     sd->time,
                                     P,
                                     zero_float3(),
+                                    light_link_receiver_nee(kg, sd),
                                     SD_BSDF_HAS_TRANSMISSION,
                                     bounce,
                                     path_flag,
@@ -831,10 +832,7 @@ ccl_device_forceinline void integrate_volume_direct_light(
 
   /* Write shadow ray and associated state to global memory. */
   integrator_state_write_shadow_ray(shadow_state, &ray);
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
+  integrator_state_write_shadow_ray_self(kg, shadow_state, &ray);
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
-  INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
 
   /* Copy state from main path to shadow path. */
   const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);
@@ -881,7 +879,7 @@ ccl_device_forceinline void integrate_volume_direct_light(
       state, path, transmission_bounce);
   INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, throughput) = throughput_phase;
 
-  /* Write Lightgroup, +1 as lightgroup is int but we need to encode into a uint8_t. */
+  /* Write Light-group, +1 as light-group is int but we need to encode into a uint8_t. */
   INTEGRATOR_STATE_WRITE(
       shadow_state, shadow_path, lightgroup) = (ls.type != LIGHT_BACKGROUND) ?
                                                    ls.group + 1 :
@@ -891,6 +889,7 @@ ccl_device_forceinline void integrate_volume_direct_light(
   INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, unlit_throughput) = unlit_throughput;
   INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, path_segment) = INTEGRATOR_STATE(
       state, guiding, path_segment);
+  INTEGRATOR_STATE(shadow_state, shadow_path, guiding_mis_weight) = 0.0f;
 #  endif
 
   integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
@@ -985,6 +984,12 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
   INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = fminf(
       unguided_phase_pdf, INTEGRATOR_STATE(state, path, min_ray_pdf));
 
+#  ifdef __LIGHT_LINKING__
+  if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING) {
+    INTEGRATOR_STATE_WRITE(state, path, mis_ray_object) = sd->object;
+  }
+#  endif
+
   path_state_next(kg, state, label, sd->flag);
   return true;
 }
@@ -1188,27 +1193,32 @@ ccl_device void integrator_shade_volume(KernelGlobals kg,
     volume_stack_clean(kg, state);
   }
 
-  VolumeIntegrateEvent event = volume_integrate(kg, state, &ray, render_buffer);
+  const VolumeIntegrateEvent event = volume_integrate(kg, state, &ray, render_buffer);
-
+  if (event == VOLUME_PATH_MISSED) {
-  if (event == VOLUME_PATH_SCATTERED) {
-    /* Queue intersect_closest kernel. */
-    integrator_path_next(kg,
-                         state,
-                         DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME,
-                         DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
-    return;
-  }
-  else if (event == VOLUME_PATH_MISSED) {
     /* End path. */
     integrator_path_terminate(kg, state, DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME);
     return;
   }
-  else {
+
+  if (event == VOLUME_PATH_ATTENUATED) {
     /* Continue to background, light or surface. */
     integrator_intersect_next_kernel_after_volume<DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME>(
         kg, state, &isect, render_buffer);
     return;
   }
+
+#  ifdef __SHADOW_LINKING__
+  if (shadow_linking_schedule_intersection_kernel<DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME>(kg,
+                                                                                         state)) {
+    return;
+  }
+#  endif /* __SHADOW_LINKING__ */
+
+  /* Queue intersect_closest kernel. */
+  integrator_path_next(kg,
+                       state,
+                       DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME,
+                       DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
 #endif /* __VOLUME__ */
 }
 

intern/cycles/kernel/integrator/shadow_linking.h

@@ -0,0 +1,71 @@
+/* SPDX-License-Identifier: Apache-2.0
+ * Copyright 2011-2023 Blender Foundation */
+
+#pragma once
+
+#include "kernel/integrator/path_state.h"
+#include "kernel/integrator/state_util.h"
+
+CCL_NAMESPACE_BEGIN
+
+#ifdef __SHADOW_LINKING__
+
+/* Check whether special shadow rays for shadow linking are needed in the current scene
+ * configuration. */
+ccl_device_forceinline bool shadow_linking_scene_need_shadow_ray(KernelGlobals kg,
+                                                                 IntegratorState state)
+{
+  if (!(kernel_data.kernel_features & KERNEL_FEATURE_SHADOW_LINKING)) {
+    /* No shadow linking in the scene, so no need to trace any extra rays. */
+    return false;
+  }
+
+  /* The distant lights might be using shadow linking, and they are not counted as
+   * kernel_data.integrator.use_light_mis.
+   * So there is a potential to avoid extra rays from being traced, but it requires more granular
+   * flags set in the integrator. */
+
+  return true;
+}
+
+/* Shadow linking re-used the main path intersection to store information about the light to which
+ * the extra ray is to be traced (this intersection communicates light between the shadow blocker
+ * intersection and shading kernels).
+ * These utilities makes a copy of the fields from the main intersection which are needed by the
+ * intersect_closest kernel after the surface bounce. */
+
+ccl_device_forceinline void shadow_linking_store_last_primitives(IntegratorState state)
+{
+  INTEGRATOR_STATE_WRITE(state, shadow_link, last_isect_prim) = INTEGRATOR_STATE(
+      state, isect, prim);
+  INTEGRATOR_STATE_WRITE(state, shadow_link, last_isect_object) = INTEGRATOR_STATE(
+      state, isect, object);
+}
+
+ccl_device_forceinline void shadow_linking_restore_last_primitives(IntegratorState state)
+{
+  INTEGRATOR_STATE_WRITE(state, isect, prim) = INTEGRATOR_STATE(
+      state, shadow_link, last_isect_prim);
+  INTEGRATOR_STATE_WRITE(state, isect, object) = INTEGRATOR_STATE(
+      state, shadow_link, last_isect_object);
+}
+
+/* Schedule shadow linking intersection kernel if it is needed.
+ * Returns true if the shadow linking specific kernel has been scheduled, false otherwise. */
+template<DeviceKernel current_kernel>
+ccl_device_inline bool shadow_linking_schedule_intersection_kernel(KernelGlobals kg,
+                                                                   IntegratorState state)
+{
+  if (!shadow_linking_scene_need_shadow_ray(kg, state)) {
+    return false;
+  }
+
+  integrator_path_next(
+      kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT);
+
+  return true;
+}
+
+#endif /* __SHADOW_LINKING__ */
+
+CCL_NAMESPACE_END