improve handling of incompatible enums

Selection sockets are meant to come right after the geometry,
this was missed in review of rBa7672caeb255e3.
Also, the selection value was not hidden.

CMakeLists.txt

@@ -642,7 +642,7 @@ if(WIN32)
   option(WITH_WINDOWS_PDB "Generate a pdb file for client side stacktraces" ON)
   mark_as_advanced(WITH_WINDOWS_PDB)
 
-  option(WITH_WINDOWS_STRIPPED_PDB "Use a stripped PDB file" On)
+  option(WITH_WINDOWS_STRIPPED_PDB "Use a stripped PDB file" ON)
   mark_as_advanced(WITH_WINDOWS_STRIPPED_PDB)
 
 endif()

build_files/cmake/clang_array_check.py

@@ -168,7 +168,7 @@ def function_parm_wash_tokens(parm):
     # if tokens[-1].kind == To
     # remove trailing char
     if tokens[-1].kind == TokenKind.PUNCTUATION:
-        if tokens[-1].spelling in (",", ")", ";"):
+        if tokens[-1].spelling in {",", ")", ";"}:
             tokens.pop()
         # else:
         #     print(tokens[-1].spelling)
@@ -179,7 +179,7 @@ def function_parm_wash_tokens(parm):
         t_spelling = t.spelling
         ok = True
         if t_kind == TokenKind.KEYWORD:
-            if t_spelling in ("const", "restrict", "volatile"):
+            if t_spelling in {"const", "restrict", "volatile"}:
                 ok = False
             elif t_spelling.startswith("__"):
                 ok = False  # __restrict

build_files/cmake/platform/platform_win32.cmake

@@ -27,7 +27,7 @@ if(NOT MSVC)
 endif()
 
 if(CMAKE_C_COMPILER_ID MATCHES "Clang")
-  set(MSVC_CLANG On)
+  set(MSVC_CLANG ON)
   set(VC_TOOLS_DIR $ENV{VCToolsRedistDir} CACHE STRING "Location of the msvc redistributables")
   set(MSVC_REDIST_DIR ${VC_TOOLS_DIR})
   if(DEFINED MSVC_REDIST_DIR)
@@ -53,7 +53,7 @@ if(CMAKE_C_COMPILER_ID MATCHES "Clang")
   endif()
   if(WITH_WINDOWS_STRIPPED_PDB)
     message(WARNING "stripped pdb not supported with clang, disabling..")
-    set(WITH_WINDOWS_STRIPPED_PDB Off)
+    set(WITH_WINDOWS_STRIPPED_PDB OFF)
   endif()
 endif()
 
@@ -159,7 +159,7 @@ endif()
 if(WITH_COMPILER_ASAN AND MSVC AND NOT MSVC_CLANG)
   if(CMAKE_CXX_COMPILER_VERSION VERSION_GREATER_EQUAL 19.28.29828)
     #set a flag so we don't have to do this comparison all the time
-    SET(MSVC_ASAN On)
+    SET(MSVC_ASAN ON)
     set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /fsanitize=address")
     set(CMAKE_C_FLAGS     "${CMAKE_C_FLAGS} /fsanitize=address")
     string(APPEND CMAKE_EXE_LINKER_FLAGS_DEBUG " /INCREMENTAL:NO")
@@ -179,22 +179,22 @@ endif()
 
 if(WITH_WINDOWS_SCCACHE AND CMAKE_VS_MSBUILD_COMMAND)
     message(WARNING "Disabling sccache, sccache is not supported with msbuild")
-    set(WITH_WINDOWS_SCCACHE Off)
+    set(WITH_WINDOWS_SCCACHE OFF)
 endif()
 
 # Debug Symbol format
 # sccache # MSVC_ASAN # format # why
-# On      # On        # Z7     # sccache will only play nice with Z7
-# On      # Off       # Z7     # sccache will only play nice with Z7
-# Off     # On        # Zi     # Asan will not play nice with Edit and Continue
-# Off     # Off       # ZI     # Neither asan nor sscache is enabled Edit and Continue is available
+# ON      # ON        # Z7     # sccache will only play nice with Z7
+# ON      # OFF       # Z7     # sccache will only play nice with Z7
+# OFF     # ON        # Zi     # Asan will not play nice with Edit and Continue
+# OFF     # OFF       # ZI     # Neither asan nor sscache is enabled Edit and Continue is available
 
 # Release Symbol format
 # sccache # MSVC_ASAN # format # why
-# On      # On        # Z7     # sccache will only play nice with Z7
-# On      # Off       # Z7     # sccache will only play nice with Z7
-# Off     # On        # Zi     # Asan will not play nice with Edit and Continue
-# Off     # Off       # Zi     # Edit and Continue disables some optimizations
+# ON      # ON        # Z7     # sccache will only play nice with Z7
+# ON      # OFF       # Z7     # sccache will only play nice with Z7
+# OFF     # ON        # Zi     # Asan will not play nice with Edit and Continue
+# OFF     # OFF       # Zi     # Edit and Continue disables some optimizations
 
 
 if(WITH_WINDOWS_SCCACHE)
@@ -288,7 +288,7 @@ if(CMAKE_GENERATOR MATCHES "^Visual Studio.+" AND # Only supported in the VS IDE
     "EnableMicrosoftCodeAnalysis=false"
     "EnableClangTidyCodeAnalysis=true"
   )
-  set(VS_CLANG_TIDY On)
+  set(VS_CLANG_TIDY ON)
 endif()
 
 # Mark libdir as system headers with a lower warn level, to resolve some warnings
@@ -469,7 +469,7 @@ if(WITH_PYTHON)
 
   set(PYTHON_INCLUDE_DIR ${LIBDIR}/python/${_PYTHON_VERSION_NO_DOTS}/include)
   set(PYTHON_NUMPY_INCLUDE_DIRS ${LIBDIR}/python/${_PYTHON_VERSION_NO_DOTS}/lib/site-packages/numpy/core/include)
-  set(NUMPY_FOUND On)
+  set(NUMPY_FOUND ON)
   unset(_PYTHON_VERSION_NO_DOTS)
   # uncached vars
   set(PYTHON_INCLUDE_DIRS "${PYTHON_INCLUDE_DIR}")
@@ -853,18 +853,18 @@ if(WITH_GMP)
   set(GMP_INCLUDE_DIRS ${LIBDIR}/gmp/include)
   set(GMP_LIBRARIES ${LIBDIR}/gmp/lib/libgmp-10.lib optimized ${LIBDIR}/gmp/lib/libgmpxx.lib debug ${LIBDIR}/gmp/lib/libgmpxx_d.lib)
   set(GMP_ROOT_DIR ${LIBDIR}/gmp)
-  set(GMP_FOUND On)
+  set(GMP_FOUND ON)
 endif()
 
 if(WITH_POTRACE)
   set(POTRACE_INCLUDE_DIRS ${LIBDIR}/potrace/include)
   set(POTRACE_LIBRARIES ${LIBDIR}/potrace/lib/potrace.lib)
-  set(POTRACE_FOUND On)
+  set(POTRACE_FOUND ON)
 endif()
 
 if(WITH_HARU)
   if(EXISTS ${LIBDIR}/haru)
-    set(HARU_FOUND On)
+    set(HARU_FOUND ON)
     set(HARU_ROOT_DIR ${LIBDIR}/haru)
     set(HARU_INCLUDE_DIRS ${HARU_ROOT_DIR}/include)
     set(HARU_LIBRARIES ${HARU_ROOT_DIR}/lib/libhpdfs.lib)

doc/python_api/sphinx_doc_gen.py

@@ -2254,7 +2254,7 @@ def main():
     # First monkey patch to load in fake members.
     setup_monkey_patch()
 
-    # Perform changes to Blender it's self.
+    # Perform changes to Blender itself.
     setup_data = setup_blender()
 
     # eventually, create the dirs

intern/cycles/blender/CMakeLists.txt

@@ -138,11 +138,6 @@ endif()
 
 blender_add_lib(bf_intern_cycles "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")
 
-# avoid link failure with clang 3.4 debug
-if(CMAKE_C_COMPILER_ID MATCHES "Clang" AND NOT ${CMAKE_C_COMPILER_VERSION} VERSION_LESS '3.4')
-  string(APPEND CMAKE_CXX_FLAGS_DEBUG " -gline-tables-only")
-endif()
-
 add_dependencies(bf_intern_cycles bf_rna)
 
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${ADDON_FILES}" ${CYCLES_INSTALL_PATH})

intern/cycles/blender/addon/engine.py

@@ -233,6 +233,7 @@ def list_render_passes(scene, srl):
     if crl.denoising_store_passes:
         yield ("Denoising Normal",          "XYZ", 'VECTOR')
         yield ("Denoising Albedo",          "RGB", 'COLOR')
+        yield ("Denoising Depth",           "Z", 'VALUE')
 
     # Custom AOV passes.
     for aov in srl.aovs:

intern/cycles/blender/addon/presets.py

@@ -40,10 +40,10 @@ class AddPresetIntegrator(AddPresetBase, Operator):
         "cycles.transparent_max_bounces",
         "cycles.caustics_reflective",
         "cycles.caustics_refractive",
-        "cycles.blur_glossy"
-        "cycles.use_fast_gi"
-        "cycles.ao_bounces"
-        "cycles.ao_bounces_render"
+        "cycles.blur_glossy",
+        "cycles.use_fast_gi",
+        "cycles.ao_bounces",
+        "cycles.ao_bounces_render",
     ]
 
     preset_subdir = "cycles/integrator"

intern/cycles/blender/addon/properties.py

@@ -87,7 +87,7 @@ enum_use_layer_samples = (
 
 enum_sampling_pattern = (
     ('SOBOL', "Sobol", "Use Sobol random sampling pattern", 0),
-    ('PROGRESSIVE_MUTI_JITTER', "Progressive Multi-Jitter", "Use Progressive Multi-Jitter random sampling pattern", 1),
+    ('PROGRESSIVE_MULTI_JITTER', "Progressive Multi-Jitter", "Use Progressive Multi-Jitter random sampling pattern", 1),
 )
 
 enum_volume_sampling = (
@@ -339,7 +339,7 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         name="Sampling Pattern",
         description="Random sampling pattern used by the integrator. When adaptive sampling is enabled, Progressive Multi-Jitter is always used instead of Sobol",
         items=enum_sampling_pattern,
-        default='PROGRESSIVE_MUTI_JITTER',
+        default='PROGRESSIVE_MULTI_JITTER',
     )
 
     scrambling_distance: FloatProperty(
@@ -1360,7 +1360,7 @@ class CyclesPreferences(bpy.types.AddonPreferences):
             elif entry.type == 'CPU':
                 cpu_devices.append(entry)
         # Extend all GPU devices with CPU.
-        if compute_device_type != 'CPU':
+        if len(devices) and compute_device_type != 'CPU':
             devices.extend(cpu_devices)
         return devices
 
@@ -1378,12 +1378,18 @@ class CyclesPreferences(bpy.types.AddonPreferences):
         self.refresh_devices()
         return None
 
+    def get_compute_device_type(self):
+        if self.compute_device_type == '':
+            return 'NONE'
+        return self.compute_device_type
+
     def get_num_gpu_devices(self):
         import _cycles
-        device_list = _cycles.available_devices(self.compute_device_type)
+        compute_device_type = self.get_compute_device_type()
+        device_list = _cycles.available_devices(compute_device_type)
         num = 0
         for device in device_list:
-            if device[1] != self.compute_device_type:
+            if device[1] != compute_device_type:
                 continue
             for dev in self.devices:
                 if dev.use and dev.id == device[2]:
@@ -1413,9 +1419,9 @@ class CyclesPreferences(bpy.types.AddonPreferences):
                 col.label(text="and NVIDIA driver version 470 or newer", icon='BLANK1')
             elif device_type == 'HIP':
                 import sys
-                col.label(text="Requires discrete AMD GPU with ??? architecture", icon='BLANK1')
+                col.label(text="Requires discrete AMD GPU with RDNA architecture", icon='BLANK1')
                 if sys.platform[:3] == "win":
-                    col.label(text="and AMD driver version ??? or newer", icon='BLANK1')
+                    col.label(text="and AMD Radeon Pro 21.Q4 driver or newer", icon='BLANK1')
             return
 
         for device in devices:
@@ -1425,15 +1431,16 @@ class CyclesPreferences(bpy.types.AddonPreferences):
         row = layout.row()
         row.prop(self, "compute_device_type", expand=True)
 
-        if self.compute_device_type == 'NONE':
+        compute_device_type = self.get_compute_device_type()
+        if compute_device_type == 'NONE':
             return
         row = layout.row()
-        devices = self.get_devices_for_type(self.compute_device_type)
-        self._draw_devices(row, self.compute_device_type, devices)
+        devices = self.get_devices_for_type(compute_device_type)
+        self._draw_devices(row, compute_device_type, devices)
 
         import _cycles
         has_peer_memory = 0
-        for device in _cycles.available_devices(self.compute_device_type):
+        for device in _cycles.available_devices(compute_device_type):
             if device[3] and self.find_existing_device_entry(device).use:
                 has_peer_memory += 1
         if has_peer_memory > 1:

intern/cycles/blender/addon/version_update.py

@@ -86,7 +86,7 @@ def do_versions(self):
             # Device might not currently be available so this can fail
             try:
                 if system.legacy_compute_device_type == 1:
-                    prop.compute_device_type = 'OPENCL'
+                    prop.compute_device_type = 'NONE' # Was OpenCL
                 elif system.legacy_compute_device_type == 2:
                     prop.compute_device_type = 'CUDA'
                 else:
@@ -97,6 +97,12 @@ def do_versions(self):
             # Init device list for UI
             prop.get_devices(prop.compute_device_type)
 
+    if bpy.context.preferences.version <= (3, 0, 40):
+        # Disable OpenCL device
+        prop = bpy.context.preferences.addons[__package__].preferences
+        if prop.is_property_set("compute_device_type") and prop['compute_device_type'] == 4:
+            prop.compute_device_type = 'NONE'
+
     # We don't modify startup file because it assumes to
     # have all the default values only.
     if not bpy.data.is_saved:
@@ -237,7 +243,7 @@ def do_versions(self):
                     cscene.use_preview_denoising = False
                 if not cscene.is_property_set("sampling_pattern") or \
                    cscene.get('sampling_pattern') >= 2:
-                    cscene.sampling_pattern = 'PROGRESSIVE_MUTI_JITTER'
+                    cscene.sampling_pattern = 'PROGRESSIVE_MULTI_JITTER'
 
                 # Removal of square samples.
                 cscene = scene.cycles

intern/cycles/blender/camera.cpp

@@ -639,7 +639,7 @@ void BlenderSync::sync_camera_motion(
     /* TODO(sergey): De-duplicate calculation with camera sync. */
     float fov = 2.0f * atanf((0.5f * sensor_size) / bcam.lens / aspectratio);
     if (fov != cam->get_fov()) {
-      VLOG(1) << "Camera " << b_ob.name() << " FOV change detected.";
+      VLOG(3) << "Camera " << b_ob.name() << " FOV change detected.";
       if (motion_time == 0.0f) {
         cam->set_fov(fov);
       }

intern/cycles/blender/curves.cpp

@@ -304,10 +304,6 @@ static void ExportCurveSegments(Scene *scene, Hair *hair, ParticleCurveData *CDa
     }
   }
 
-  if (num_curves > 0) {
-    VLOG(1) << "Exporting curve segments for mesh " << hair->name;
-  }
-
   hair->reserve_curves(hair->num_curves() + num_curves, hair->get_curve_keys().size() + num_keys);
 
   num_keys = 0;
@@ -356,7 +352,7 @@ static void ExportCurveSegments(Scene *scene, Hair *hair, ParticleCurveData *CDa
 
   /* check allocation */
   if ((hair->get_curve_keys().size() != num_keys) || (hair->num_curves() != num_curves)) {
-    VLOG(1) << "Allocation failed, clearing data";
+    VLOG(1) << "Hair memory allocation failed, clearing data.";
     hair->clear(true);
   }
 }
@@ -412,16 +408,11 @@ static void export_hair_motion_validate_attribute(Hair *hair,
   if (num_motion_keys != num_keys || !have_motion) {
     /* No motion or hair "topology" changed, remove attributes again. */
     if (num_motion_keys != num_keys) {
-      VLOG(1) << "Hair topology changed, removing attribute.";
-    }
-    else {
-      VLOG(1) << "No motion, removing attribute.";
+      VLOG(1) << "Hair topology changed, removing motion attribute.";
     }
     hair->attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION);
   }
   else if (motion_step > 0) {
-    VLOG(1) << "Filling in new motion vertex position for motion_step " << motion_step;
-
     /* Motion, fill up previous steps that we might have skipped because
      * they had no motion, but we need them anyway now. */
     for (int step = 0; step < motion_step; step++) {
@@ -437,16 +428,12 @@ static void export_hair_motion_validate_attribute(Hair *hair,
 
 static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, int motion_step)
 {
-  VLOG(1) << "Exporting curve motion segments for hair " << hair->name << ", motion step "
-          << motion_step;
-
   /* find attribute */
   Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
   bool new_attribute = false;
 
   /* add new attribute if it doesn't exist already */
   if (!attr_mP) {
-    VLOG(1) << "Creating new motion vertex position attribute";
     attr_mP = hair->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
     new_attribute = true;
   }
@@ -682,10 +669,6 @@ static void export_hair_curves(Scene *scene, Hair *hair, BL::Hair b_hair)
   const int num_keys = b_hair.points.length();
   const int num_curves = b_hair.curves.length();
 
-  if (num_curves > 0) {
-    VLOG(1) << "Exporting curve segments for hair " << hair->name;
-  }
-
   hair->reserve_curves(num_curves, num_keys);
 
   /* Export curves and points. */
@@ -743,15 +726,11 @@ static void export_hair_curves(Scene *scene, Hair *hair, BL::Hair b_hair)
 
 static void export_hair_curves_motion(Hair *hair, BL::Hair b_hair, int motion_step)
 {
-  VLOG(1) << "Exporting curve motion segments for hair " << hair->name << ", motion step "
-          << motion_step;
-
   /* Find or add attribute. */
   Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
   bool new_attribute = false;
 
   if (!attr_mP) {
-    VLOG(1) << "Creating new motion vertex position attribute";
     attr_mP = hair->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
     new_attribute = true;
   }

intern/cycles/blender/object.cpp

@@ -76,18 +76,15 @@ bool BlenderSync::object_is_geometry(BL::Object &b_ob)
     /* Will be exported attached to mesh. */
     return true;
   }
-  else if (type == BL::Object::type_CURVE) {
-    /* Skip exporting curves without faces, overhead can be
-     * significant if there are many for path animation. */
-    BL::Curve b_curve(b_ob_data);
 
-    return (b_curve.bevel_object() || b_curve.extrude() != 0.0f || b_curve.bevel_depth() != 0.0f ||
-            b_curve.dimensions() == BL::Curve::dimensions_2D || b_ob.modifiers.length());
-  }
-  else {
-    return (b_ob_data.is_a(&RNA_Mesh) || b_ob_data.is_a(&RNA_Curve) ||
-            b_ob_data.is_a(&RNA_MetaBall));
+  /* Other object types that are not meshes but evaluate to meshes are presented to render engines
+   * as separate instance objects. Metaballs and surface objects have not been affected by that
+   * change yet. */
+  if (type == BL::Object::type_SURFACE || type == BL::Object::type_META) {
+    return true;
   }
+
+  return b_ob_data.is_a(&RNA_Mesh);
 }
 
 bool BlenderSync::object_is_light(BL::Object &b_ob)
@@ -161,6 +158,11 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph,
   if (is_instance) {
     persistent_id_array = b_instance.persistent_id();
     persistent_id = persistent_id_array.data;
+    if (!b_ob_info.is_real_object_data()) {
+      /* Remember which object data the geometry is coming from, so that we can sync it when the
+       * object has changed. */
+      instance_geometries_by_object[b_ob_info.real_object.ptr.data].insert(b_ob_info.object_data);
+    }
   }
 
   /* light is handled separately */
@@ -560,6 +562,7 @@ void BlenderSync::sync_objects(BL::Depsgraph &b_depsgraph,
   else {
     geometry_motion_synced.clear();
   }
+  instance_geometries_by_object.clear();
 
   /* initialize culling */
   BlenderObjectCulling culling(scene, b_scene);

intern/cycles/blender/python.cpp

@@ -157,8 +157,6 @@ static PyObject *init_func(PyObject * /*self*/, PyObject *args)
 
   DebugFlags().running_inside_blender = true;
 
-  VLOG(2) << "Debug flags initialized to:\n" << DebugFlags();
-
   Py_RETURN_NONE;
 }
 
@@ -885,8 +883,6 @@ static PyObject *debug_flags_update_func(PyObject * /*self*/, PyObject *args)
 
   debug_flags_sync_from_scene(b_scene);
 
-  VLOG(2) << "Debug flags set to:\n" << DebugFlags();
-
   debug_flags_set = true;
 
   Py_RETURN_NONE;
@@ -896,7 +892,6 @@ static PyObject *debug_flags_reset_func(PyObject * /*self*/, PyObject * /*args*/
 {
   debug_flags_reset();
   if (debug_flags_set) {
-    VLOG(2) << "Debug flags reset to:\n" << DebugFlags();
     debug_flags_set = false;
   }
   Py_RETURN_NONE;

intern/cycles/blender/sync.cpp

@@ -183,6 +183,15 @@ void BlenderSync::sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d
               (object_subdivision_type(b_ob, preview, experimental) != Mesh::SUBDIVISION_NONE)) {
             BL::ID key = BKE_object_is_modified(b_ob) ? b_ob : b_ob.data();
             geometry_map.set_recalc(key);
+
+            /* Sync all contained geometry instances as well when the object changed.. */
+            map<void *, set<BL::ID>>::const_iterator instance_geometries =
+                instance_geometries_by_object.find(b_ob.ptr.data);
+            if (instance_geometries != instance_geometries_by_object.end()) {
+              for (BL::ID geometry : instance_geometries->second) {
+                geometry_map.set_recalc(geometry);
+              }
+            }
           }
 
           if (updated_geometry) {
@@ -366,7 +375,9 @@ void BlenderSync::sync_integrator(BL::ViewLayer &b_view_layer, bool background)
   if ((preview && !preview_scrambling_distance) || use_adaptive_sampling)
     scrambling_distance = 1.0f;
 
-  VLOG(1) << "Used Scrambling Distance: " << scrambling_distance;
+  if (scrambling_distance != 1.0f) {
+    VLOG(3) << "Using scrambling distance: " << scrambling_distance;
+  }
   integrator->set_scrambling_distance(scrambling_distance);
 
   if (get_boolean(cscene, "use_fast_gi")) {

intern/cycles/blender/sync.h

@@ -225,6 +225,8 @@ class BlenderSync {
   set<Geometry *> geometry_synced;
   set<Geometry *> geometry_motion_synced;
   set<Geometry *> geometry_motion_attribute_synced;
+  /** Remember which geometries come from which objects to be able to sync them after changes. */
+  map<void *, set<BL::ID>> instance_geometries_by_object;
   set<float> motion_times;
   void *world_map;
   bool world_recalc;

intern/cycles/device/cpu/device.cpp

@@ -38,7 +38,6 @@ void device_cpu_info(vector<DeviceInfo> &devices)
   info.id = "CPU";
   info.num = 0;
   info.has_osl = true;
-  info.has_half_images = true;
   info.has_nanovdb = true;
   info.has_profiling = true;
   if (openimagedenoise_supported()) {

intern/cycles/device/cpu/device_impl.cpp

@@ -68,8 +68,8 @@ CPUDevice::CPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_
 {
   /* Pick any kernel, all of them are supposed to have same level of microarchitecture
    * optimization. */
-  VLOG(1) << "Will be using " << kernels.integrator_init_from_camera.get_uarch_name()
-          << " kernels.";
+  VLOG(1) << "Using " << get_cpu_kernels().integrator_init_from_camera.get_uarch_name()
+          << " CPU kernels.";
 
   if (info.cpu_threads == 0) {
     info.cpu_threads = TaskScheduler::num_threads();
@@ -297,11 +297,6 @@ void CPUDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
     Device::build_bvh(bvh, progress, refit);
 }
 
-const CPUKernels *CPUDevice::get_cpu_kernels() const
-{
-  return &kernels;
-}
-
 void CPUDevice::get_cpu_kernel_thread_globals(
     vector<CPUKernelThreadGlobals> &kernel_thread_globals)
 {

intern/cycles/device/cpu/device_impl.h

@@ -57,8 +57,6 @@ class CPUDevice : public Device {
   RTCDevice embree_device;
 #endif
 
-  CPUKernels kernels;
-
   CPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_);
   ~CPUDevice();
 
@@ -90,7 +88,6 @@ class CPUDevice : public Device {
 
   void build_bvh(BVH *bvh, Progress &progress, bool refit) override;
 
-  virtual const CPUKernels *get_cpu_kernels() const override;
   virtual void get_cpu_kernel_thread_globals(
       vector<CPUKernelThreadGlobals> &kernel_thread_globals) override;
   virtual void *get_cpu_osl_memory() override;

intern/cycles/device/cpu/kernel.cpp

@@ -26,6 +26,9 @@ CCL_NAMESPACE_BEGIN
       KERNEL_NAME_EVAL(cpu_avx, name), KERNEL_NAME_EVAL(cpu_avx2, name)
 
 #define REGISTER_KERNEL(name) name(KERNEL_FUNCTIONS(name))
+#define REGISTER_KERNEL_FILM_CONVERT(name) \
+  film_convert_##name(KERNEL_FUNCTIONS(film_convert_##name)), \
+      film_convert_half_rgba_##name(KERNEL_FUNCTIONS(film_convert_half_rgba_##name))
 
 CPUKernels::CPUKernels()
     : /* Integrator. */
@@ -50,11 +53,25 @@ CPUKernels::CPUKernels()
       REGISTER_KERNEL(adaptive_sampling_filter_x),
       REGISTER_KERNEL(adaptive_sampling_filter_y),
       /* Cryptomatte. */
-      REGISTER_KERNEL(cryptomatte_postprocess)
+      REGISTER_KERNEL(cryptomatte_postprocess),
+      /* Film Convert. */
+      REGISTER_KERNEL_FILM_CONVERT(depth),
+      REGISTER_KERNEL_FILM_CONVERT(mist),
+      REGISTER_KERNEL_FILM_CONVERT(sample_count),
+      REGISTER_KERNEL_FILM_CONVERT(float),
+      REGISTER_KERNEL_FILM_CONVERT(light_path),
+      REGISTER_KERNEL_FILM_CONVERT(float3),
+      REGISTER_KERNEL_FILM_CONVERT(motion),
+      REGISTER_KERNEL_FILM_CONVERT(cryptomatte),
+      REGISTER_KERNEL_FILM_CONVERT(shadow_catcher),
+      REGISTER_KERNEL_FILM_CONVERT(shadow_catcher_matte_with_shadow),
+      REGISTER_KERNEL_FILM_CONVERT(combined),
+      REGISTER_KERNEL_FILM_CONVERT(float4)
 {
 }
 
 #undef REGISTER_KERNEL
+#undef REGISTER_KERNEL_FILM_CONVERT
 #undef KERNEL_FUNCTIONS
 
 CCL_NAMESPACE_END

intern/cycles/device/cpu/kernel.h

@@ -17,11 +17,13 @@
 #pragma once
 
 #include "device/cpu/kernel_function.h"
+#include "util/half.h"
 #include "util/types.h"
 
 CCL_NAMESPACE_BEGIN
 
 struct KernelGlobalsCPU;
+struct KernelFilmConvert;
 struct IntegratorStateCPU;
 struct TileInfo;
 
@@ -102,6 +104,41 @@ class CPUKernels {
 
   CryptomattePostprocessFunction cryptomatte_postprocess;
 
+  /* Film Convert. */
+  using FilmConvertFunction = CPUKernelFunction<void (*)(const KernelFilmConvert *kfilm_convert,
+                                                         const float *buffer,
+                                                         float *pixel,
+                                                         const int width,
+                                                         const int buffer_stride,
+                                                         const int pixel_stride)>;
+  using FilmConvertHalfRGBAFunction =
+      CPUKernelFunction<void (*)(const KernelFilmConvert *kfilm_convert,
+                                 const float *buffer,
+                                 half4 *pixel,
+                                 const int width,
+                                 const int buffer_stride)>;
+
+#define KERNEL_FILM_CONVERT_FUNCTION(name) \
+  FilmConvertFunction film_convert_##name; \
+  FilmConvertHalfRGBAFunction film_convert_half_rgba_##name;
+
+  KERNEL_FILM_CONVERT_FUNCTION(depth)
+  KERNEL_FILM_CONVERT_FUNCTION(mist)
+  KERNEL_FILM_CONVERT_FUNCTION(sample_count)
+  KERNEL_FILM_CONVERT_FUNCTION(float)
+
+  KERNEL_FILM_CONVERT_FUNCTION(light_path)
+  KERNEL_FILM_CONVERT_FUNCTION(float3)
+
+  KERNEL_FILM_CONVERT_FUNCTION(motion)
+  KERNEL_FILM_CONVERT_FUNCTION(cryptomatte)
+  KERNEL_FILM_CONVERT_FUNCTION(shadow_catcher)
+  KERNEL_FILM_CONVERT_FUNCTION(shadow_catcher_matte_with_shadow)
+  KERNEL_FILM_CONVERT_FUNCTION(combined)
+  KERNEL_FILM_CONVERT_FUNCTION(float4)
+
+#undef KERNEL_FILM_CONVERT_FUNCTION
+
   CPUKernels();
 };
 

intern/cycles/device/cuda/device.cpp

@@ -144,7 +144,6 @@ void device_cuda_info(vector<DeviceInfo> &devices)
     info.description = string(name);
     info.num = num;
 
-    info.has_half_images = (major >= 3);
     info.has_nanovdb = true;
     info.denoisers = 0;
 

intern/cycles/device/cuda/device_impl.cpp

@@ -378,7 +378,9 @@ string CUDADevice::compile_kernel(const uint kernel_features,
       cubin.c_str(),
       common_cflags.c_str());
 
-  printf("Compiling CUDA kernel ...\n%s\n", command.c_str());
+  printf("Compiling %sCUDA kernel ...\n%s\n",
+         (use_adaptive_compilation()) ? "adaptive " : "",
+         command.c_str());
 
 #  ifdef _WIN32
   command = "call " + command;
@@ -405,13 +407,15 @@ string CUDADevice::compile_kernel(const uint kernel_features,
 
 bool CUDADevice::load_kernels(const uint kernel_features)
 {
-  /* TODO(sergey): Support kernels re-load for CUDA devices.
+  /* TODO(sergey): Support kernels re-load for CUDA devices adaptive compile.
    *
    * Currently re-loading kernel will invalidate memory pointers,
    * causing problems in cuCtxSynchronize.
    */
   if (cuModule) {
-    VLOG(1) << "Skipping kernel reload, not currently supported.";
+    if (use_adaptive_compilation()) {
+      VLOG(1) << "Skipping CUDA kernel reload for adaptive compilation, not currently supported.";
+    }
     return true;
   }
 

intern/cycles/device/device.cpp

@@ -23,6 +23,7 @@
 #include "device/queue.h"
 
 #include "device/cpu/device.h"
+#include "device/cpu/kernel.h"
 #include "device/cuda/device.h"
 #include "device/dummy/device.h"
 #include "device/hip/device.h"
@@ -285,7 +286,6 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
   info.description = "Multi Device";
   info.num = 0;
 
-  info.has_half_images = true;
   info.has_nanovdb = true;
   info.has_osl = true;
   info.has_profiling = true;
@@ -332,7 +332,6 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
     }
 
     /* Accumulate device info. */
-    info.has_half_images &= device.has_half_images;
     info.has_nanovdb &= device.has_nanovdb;
     info.has_osl &= device.has_osl;
     info.has_profiling &= device.has_profiling;
@@ -363,10 +362,11 @@ unique_ptr<DeviceQueue> Device::gpu_queue_create()
   return nullptr;
 }
 
-const CPUKernels *Device::get_cpu_kernels() const
+const CPUKernels &Device::get_cpu_kernels()
 {
-  LOG(FATAL) << "Device does not support CPU kernels.";
-  return nullptr;
+  /* Initialize CPU kernels once and reuse. */
+  static CPUKernels kernels;
+  return kernels;
 }
 
 void Device::get_cpu_kernel_thread_globals(

intern/cycles/device/device.h

@@ -73,7 +73,6 @@ class DeviceInfo {
   int num;
   bool display_device;        /* GPU is used as a display device. */
   bool has_nanovdb;           /* Support NanoVDB volumes. */
-  bool has_half_images;       /* Support half-float textures. */
   bool has_osl;               /* Support Open Shading Language. */
   bool has_profiling;         /* Supports runtime collection of profiling info. */
   bool has_peer_memory;       /* GPU has P2P access to memory of another GPU. */
@@ -90,7 +89,6 @@ class DeviceInfo {
     num = 0;
     cpu_threads = 0;
     display_device = false;
-    has_half_images = false;
     has_nanovdb = false;
     has_osl = false;
     has_profiling = false;
@@ -180,7 +178,7 @@ class Device {
    * These may not be used on GPU or multi-devices. */
 
   /* Get CPU kernel functions for native instruction set. */
-  virtual const CPUKernels *get_cpu_kernels() const;
+  static const CPUKernels &get_cpu_kernels();
   /* Get kernel globals to pass to kernels. */
   virtual void get_cpu_kernel_thread_globals(
       vector<CPUKernelThreadGlobals> & /*kernel_thread_globals*/);

intern/cycles/device/hip/device.cpp

@@ -131,9 +131,9 @@ void device_hip_info(vector<DeviceInfo> &devices)
       continue;
     }
 
-    int major;
-    hipDeviceGetAttribute(&major, hipDeviceAttributeComputeCapabilityMajor, num);
-    // TODO : (Arya) What is the last major version we are supporting?
+    if (!hipSupportsDevice(num)) {
+      continue;
+    }
 
     DeviceInfo info;
 
@@ -141,7 +141,6 @@ void device_hip_info(vector<DeviceInfo> &devices)
     info.description = string(name);
     info.num = num;
 
-    info.has_half_images = (major >= 3);
     info.has_nanovdb = true;
     info.denoisers = 0;
 

intern/cycles/device/hip/device_impl.cpp

@@ -146,12 +146,18 @@ HIPDevice::~HIPDevice()
 
 bool HIPDevice::support_device(const uint /*kernel_features*/)
 {
-  int major, minor;
-  hipDeviceGetAttribute(&major, hipDeviceAttributeComputeCapabilityMajor, hipDevId);
-  hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
+  if (hipSupportsDevice(hipDevId)) {
+    return true;
+  }
+  else {
+    /* We only support Navi and above. */
+    hipDeviceProp_t props;
+    hipGetDeviceProperties(&props, hipDevId);
 
-  // TODO : (Arya) What versions do we plan to support?
-  return true;
+    set_error(string_printf("HIP backend requires AMD RDNA graphics card or up, but found %s.",
+                            props.name));
+    return false;
+  }
 }
 
 bool HIPDevice::check_peer_access(Device *peer_device)
@@ -240,36 +246,23 @@ string HIPDevice::compile_kernel(const uint kernel_features,
   hipDeviceProp_t props;
   hipGetDeviceProperties(&props, hipDevId);
 
+  /* gcnArchName can contain tokens after the arch name with features, ie.
+    "gfx1010:sramecc-:xnack-" so we tokenize it to get the first part. */
+  char *arch = strtok(props.gcnArchName, ":");
+  if (arch == NULL) {
+    arch = props.gcnArchName;
+  }
+
   /* Attempt to use kernel provided with Blender. */
   if (!use_adaptive_compilation()) {
     if (!force_ptx) {
-      const string fatbin = path_get(string_printf("lib/%s_%s.fatbin", name, props.gcnArchName));
+      const string fatbin = path_get(string_printf("lib/%s_%s.fatbin", name, arch));
       VLOG(1) << "Testing for pre-compiled kernel " << fatbin << ".";
       if (path_exists(fatbin)) {
         VLOG(1) << "Using precompiled kernel.";
         return fatbin;
       }
     }
-
-    /* The driver can JIT-compile PTX generated for older generations, so find the closest one. */
-    int ptx_major = major, ptx_minor = minor;
-    while (ptx_major >= 3) {
-      const string ptx = path_get(
-          string_printf("lib/%s_compute_%d%d.ptx", name, ptx_major, ptx_minor));
-      VLOG(1) << "Testing for pre-compiled kernel " << ptx << ".";
-      if (path_exists(ptx)) {
-        VLOG(1) << "Using precompiled kernel.";
-        return ptx;
-      }
-
-      if (ptx_minor > 0) {
-        ptx_minor--;
-      }
-      else {
-        ptx_major--;
-        ptx_minor = 9;
-      }
-    }
   }
 
   /* Try to use locally compiled kernel. */
@@ -292,12 +285,10 @@ string HIPDevice::compile_kernel(const uint kernel_features,
 #  ifdef _DEBUG
   options.append(" -save-temps");
 #  endif
-  options.append(" --amdgpu-target=").append(props.gcnArchName);
+  options.append(" --amdgpu-target=").append(arch);
 
   const string include_path = source_path;
-  const char *const kernel_arch = props.gcnArchName;
-  const string fatbin_file = string_printf(
-      "cycles_%s_%s_%s", name, kernel_arch, kernel_md5.c_str());
+  const string fatbin_file = string_printf("cycles_%s_%s_%s", name, arch, kernel_md5.c_str());
   const string fatbin = path_cache_get(path_join("kernels", fatbin_file));
   VLOG(1) << "Testing for locally compiled kernel " << fatbin << ".";
   if (path_exists(fatbin)) {
@@ -360,7 +351,9 @@ string HIPDevice::compile_kernel(const uint kernel_features,
                                  source_path.c_str(),
                                  fatbin.c_str());
 
-  printf("Compiling HIP kernel ...\n%s\n", command.c_str());
+  printf("Compiling %sHIP kernel ...\n%s\n",
+         (use_adaptive_compilation()) ? "adaptive " : "",
+         command.c_str());
 
 #  ifdef _WIN32
   command = "call " + command;
@@ -387,13 +380,15 @@ string HIPDevice::compile_kernel(const uint kernel_features,
 
 bool HIPDevice::load_kernels(const uint kernel_features)
 {
-  /* TODO(sergey): Support kernels re-load for HIP devices.
+  /* TODO(sergey): Support kernels re-load for CUDA devices adaptive compile.
    *
    * Currently re-loading kernel will invalidate memory pointers,
-   * causing problems in hipCtxSynchronize.
+   * causing problems in cuCtxSynchronize.
    */
   if (hipModule) {
-    VLOG(1) << "Skipping kernel reload, not currently supported.";
+    if (use_adaptive_compilation()) {
+      VLOG(1) << "Skipping HIP kernel reload for adaptive compilation, not currently supported.";
+    }
     return true;
   }
 
@@ -402,8 +397,9 @@ bool HIPDevice::load_kernels(const uint kernel_features)
     return false;
 
   /* check if GPU is supported */
-  if (!support_device(kernel_features))
+  if (!support_device(kernel_features)) {
     return false;
+  }
 
   /* get kernel */
   const char *kernel_name = "kernel";

intern/cycles/device/hip/util.h

@@ -58,6 +58,15 @@ const char *hipewCompilerPath();
 int hipewCompilerVersion();
 #  endif /* WITH_HIP_DYNLOAD */
 
+static inline bool hipSupportsDevice(const int hipDevId)
+{
+  int major, minor;
+  hipDeviceGetAttribute(&major, hipDeviceAttributeComputeCapabilityMajor, hipDevId);
+  hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
+
+  return (major > 10) || (major == 10 && minor >= 1);
+}
+
 CCL_NAMESPACE_END
 
 #endif /* WITH_HIP */

intern/cycles/device/optix/device_impl.cpp

@@ -48,14 +48,6 @@ OptiXDevice::Denoiser::Denoiser(OptiXDevice *device)
 {
 }
 
-OptiXDevice::Denoiser::~Denoiser()
-{
-  const CUDAContextScope scope(device);
-  if (optix_denoiser != nullptr) {
-    optixDenoiserDestroy(optix_denoiser);
-  }
-}
-
 OptiXDevice::OptiXDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
     : CUDADevice(info, stats, profiler),
       sbt_data(this, "__sbt", MEM_READ_ONLY),
@@ -91,6 +83,7 @@ OptiXDevice::OptiXDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
   };
 #  endif
   if (DebugFlags().optix.use_debug) {
+    VLOG(1) << "Using OptiX debug mode.";
     options.validationMode = OPTIX_DEVICE_CONTEXT_VALIDATION_MODE_ALL;
   }
   optix_assert(optixDeviceContextCreate(cuContext, &options, &context));
@@ -132,6 +125,11 @@ OptiXDevice::~OptiXDevice()
     }
   }
 
+  /* Make sure denoiser is destroyed before device context! */
+  if (denoiser_.optix_denoiser != nullptr) {
+    optixDenoiserDestroy(denoiser_.optix_denoiser);
+  }
+
   optixDeviceContextDestroy(context);
 }
 
@@ -883,27 +881,32 @@ bool OptiXDevice::denoise_configure_if_needed(DenoiseContext &context)
   optix_assert(optixDenoiserComputeMemoryResources(
       denoiser_.optix_denoiser, buffer_params.width, buffer_params.height, &sizes));
 
-  denoiser_.scratch_size = sizes.withOverlapScratchSizeInBytes;
+  /* Denoiser is invoked on whole images only, so no overlap needed (would be used for tiling). */
+  denoiser_.scratch_size = sizes.withoutOverlapScratchSizeInBytes;
   denoiser_.scratch_offset = sizes.stateSizeInBytes;
 
   /* Allocate denoiser state if tile size has changed since last setup. */
-  denoiser_.state.alloc_to_device(denoiser_.scratch_offset + denoiser_.scratch_size);
+  denoiser_.state.alloc_to_device(denoiser_.scratch_offset + denoiser_.scratch_size +
+                                  sizeof(float));
 
   /* Initialize denoiser state for the current tile size. */
-  const OptixResult result = optixDenoiserSetup(denoiser_.optix_denoiser,
-                                                denoiser_.queue.stream(),
-                                                buffer_params.width,
-                                                buffer_params.height,
-                                                denoiser_.state.device_pointer,
-                                                denoiser_.scratch_offset,
-                                                denoiser_.state.device_pointer +
-                                                    denoiser_.scratch_offset,
-                                                denoiser_.scratch_size);
+  const OptixResult result = optixDenoiserSetup(
+      denoiser_.optix_denoiser,
+      0, /* Work around bug in r495 drivers that causes artifacts when denoiser setup is called
+            on a stream that is not the default stream */
+      buffer_params.width,
+      buffer_params.height,
+      denoiser_.state.device_pointer,
+      denoiser_.scratch_offset,
+      denoiser_.state.device_pointer + denoiser_.scratch_offset,
+      denoiser_.scratch_size);
   if (result != OPTIX_SUCCESS) {
     set_error("Failed to set up OptiX denoiser");
     return false;
   }
 
+  cuda_assert(cuCtxSynchronize());
+
   denoiser_.is_configured = true;
   denoiser_.configured_size.x = buffer_params.width;
   denoiser_.configured_size.y = buffer_params.height;
@@ -938,8 +941,6 @@ bool OptiXDevice::denoise_run(DenoiseContext &context, const DenoisePass &pass)
     color_layer.format = OPTIX_PIXEL_FORMAT_FLOAT3;
   }
 
-  device_vector<float> fake_albedo(this, "fake_albedo", MEM_READ_WRITE);
-
   /* Optional albedo and color passes. */
   if (context.num_input_passes > 1) {
     const device_ptr d_guiding_buffer = context.guiding_params.device_pointer;
@@ -970,6 +971,17 @@ bool OptiXDevice::denoise_run(DenoiseContext &context, const DenoisePass &pass)
 
   /* Finally run denoising. */
   OptixDenoiserParams params = {}; /* All parameters are disabled/zero. */
+  params.hdrIntensity = denoiser_.state.device_pointer + denoiser_.scratch_offset +
+                        denoiser_.scratch_size;
+
+  optix_assert(
+      optixDenoiserComputeIntensity(denoiser_.optix_denoiser,
+                                    denoiser_.queue.stream(),
+                                    &color_layer,
+                                    params.hdrIntensity,
+                                    denoiser_.state.device_pointer + denoiser_.scratch_offset,
+                                    denoiser_.scratch_size));
+
   OptixDenoiserLayer image_layers = {};
   image_layers.input = color_layer;
   image_layers.output = output_layer;

intern/cycles/device/optix/device_impl.h

@@ -82,7 +82,6 @@ class OptiXDevice : public CUDADevice {
   class Denoiser {
    public:
     explicit Denoiser(OptiXDevice *device);
-    ~Denoiser();
 
     OptiXDevice *device;
     OptiXDeviceQueue queue;

intern/cycles/integrator/denoiser.cpp

@@ -29,23 +29,11 @@ unique_ptr<Denoiser> Denoiser::create(Device *path_trace_device, const DenoisePa
 {
   DCHECK(params.use);
 
-  switch (params.type) {
-    case DENOISER_OPTIX:
-      return make_unique<OptiXDenoiser>(path_trace_device, params);
-
-    case DENOISER_OPENIMAGEDENOISE:
-      return make_unique<OIDNDenoiser>(path_trace_device, params);
-
-    case DENOISER_NUM:
-    case DENOISER_NONE:
-    case DENOISER_ALL:
-      /* pass */
-      break;
+  if (params.type == DENOISER_OPTIX && Device::available_devices(DEVICE_MASK_OPTIX).size()) {
+    return make_unique<OptiXDenoiser>(path_trace_device, params);
   }
 
-  LOG(FATAL) << "Unhandled denoiser type " << params.type << ", should never happen.";
-
-  return nullptr;
+  return make_unique<OIDNDenoiser>(path_trace_device, params);
 }
 
 Denoiser::Denoiser(Device *path_trace_device, const DenoiseParams &params)

intern/cycles/integrator/pass_accessor.cpp

@@ -138,10 +138,6 @@ bool PassAccessor::get_render_tile_pixels(const RenderBuffers *render_buffers,
     return false;
   }
 
-  if (pass_access_info_.offset == PASS_UNUSED) {
-    return false;
-  }
-
   const PassType type = pass_access_info_.type;
   const PassMode mode = pass_access_info_.mode;
   const PassInfo pass_info = Pass::get_info(type, pass_access_info_.include_albedo);

intern/cycles/integrator/pass_accessor_cpu.cpp

@@ -14,9 +14,12 @@
  * limitations under the License.
  */
 
+#include "device/device.h"
+
 #include "integrator/pass_accessor_cpu.h"
 
 #include "session/buffers.h"
+
 #include "util/log.h"
 #include "util/tbb.h"
 
@@ -33,70 +36,16 @@ CCL_NAMESPACE_BEGIN
  * Kernel processing.
  */
 
-template<typename Processor>
-inline void PassAccessorCPU::run_get_pass_kernel_processor(const RenderBuffers *render_buffers,
-                                                           const BufferParams &buffer_params,
-                                                           const Destination &destination,
-                                                           const Processor &processor) const
-{
-  KernelFilmConvert kfilm_convert;
-  init_kernel_film_convert(&kfilm_convert, buffer_params, destination);
-
-  if (destination.pixels) {
-    /* NOTE: No overlays are applied since they are not used for final renders.
-     * Can be supported via some sort of specialization to avoid code duplication. */
-
-    run_get_pass_kernel_processor_float(
-        &kfilm_convert, render_buffers, buffer_params, destination, processor);
-  }
-
-  if (destination.pixels_half_rgba) {
-    /* TODO(sergey): Consider adding specialization to avoid per-pixel overlay check. */
-
-    if (destination.num_components == 1) {
-      run_get_pass_kernel_processor_half_rgba(&kfilm_convert,
-                                              render_buffers,
-                                              buffer_params,
-                                              destination,
-                                              [&processor](const KernelFilmConvert *kfilm_convert,
-                                                           ccl_global const float *buffer,
-                                                           float *pixel_rgba) {
-                                                float pixel;
-                                                processor(kfilm_convert, buffer, &pixel);
-
-                                                pixel_rgba[0] = pixel;
-                                                pixel_rgba[1] = pixel;
-                                                pixel_rgba[2] = pixel;
-                                                pixel_rgba[3] = 1.0f;
-                                              });
-    }
-    else if (destination.num_components == 3) {
-      run_get_pass_kernel_processor_half_rgba(&kfilm_convert,
-                                              render_buffers,
-                                              buffer_params,
-                                              destination,
-                                              [&processor](const KernelFilmConvert *kfilm_convert,
-                                                           ccl_global const float *buffer,
-                                                           float *pixel_rgba) {
-                                                processor(kfilm_convert, buffer, pixel_rgba);
-                                                pixel_rgba[3] = 1.0f;
-                                              });
-    }
-    else if (destination.num_components == 4) {
-      run_get_pass_kernel_processor_half_rgba(
-          &kfilm_convert, render_buffers, buffer_params, destination, processor);
-    }
-  }
-}
-
-template<typename Processor>
 inline void PassAccessorCPU::run_get_pass_kernel_processor_float(
     const KernelFilmConvert *kfilm_convert,
     const RenderBuffers *render_buffers,
     const BufferParams &buffer_params,
     const Destination &destination,
-    const Processor &processor) const
+    const CPUKernels::FilmConvertFunction func) const
 {
+  /* NOTE: No overlays are applied since they are not used for final renders.
+   * Can be supported via some sort of specialization to avoid code duplication. */
+
   DCHECK_EQ(destination.stride, 0) << "Custom stride for float destination is not implemented.";
 
   const int64_t pass_stride = buffer_params.pass_stride;
@@ -112,21 +61,16 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_float(
     const float *buffer = window_data + y * buffer_row_stride;
     float *pixel = destination.pixels +
                    (y * buffer_params.width + destination.offset) * pixel_stride;
-
-    for (int64_t x = 0; x < buffer_params.window_width;
-         ++x, buffer += pass_stride, pixel += pixel_stride) {
-      processor(kfilm_convert, buffer, pixel);
-    }
+    func(kfilm_convert, buffer, pixel, buffer_params.window_width, pass_stride, pixel_stride);
   });
 }
 
-template<typename Processor>
 inline void PassAccessorCPU::run_get_pass_kernel_processor_half_rgba(
     const KernelFilmConvert *kfilm_convert,
     const RenderBuffers *render_buffers,
     const BufferParams &buffer_params,
     const Destination &destination,
-    const Processor &processor) const
+    const CPUKernels::FilmConvertHalfRGBAFunction func) const
 {
   const int64_t pass_stride = buffer_params.pass_stride;
   const int64_t buffer_row_stride = buffer_params.stride * buffer_params.pass_stride;
@@ -141,16 +85,7 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_half_rgba(
   tbb::parallel_for(0, buffer_params.window_height, [&](int64_t y) {
     const float *buffer = window_data + y * buffer_row_stride;
     half4 *pixel = dst_start + y * destination_stride;
-    for (int64_t x = 0; x < buffer_params.window_width; ++x, buffer += pass_stride, ++pixel) {
-
-      float pixel_rgba[4];
-      processor(kfilm_convert, buffer, pixel_rgba);
-
-      film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel_rgba);
-
-      *pixel = float4_to_half4_display(
-          make_float4(pixel_rgba[0], pixel_rgba[1], pixel_rgba[2], pixel_rgba[3]));
-    }
+    func(kfilm_convert, buffer, pixel, buffer_params.window_width, pass_stride);
   });
 }
 
@@ -163,8 +98,25 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_half_rgba(
                                         const BufferParams &buffer_params, \
                                         const Destination &destination) const \
   { \
-    run_get_pass_kernel_processor( \
-        render_buffers, buffer_params, destination, film_get_pass_pixel_##pass); \
+    const CPUKernels &kernels = Device::get_cpu_kernels(); \
+    KernelFilmConvert kfilm_convert; \
+    init_kernel_film_convert(&kfilm_convert, buffer_params, destination); \
+\
+    if (destination.pixels) { \
+      run_get_pass_kernel_processor_float(&kfilm_convert, \
+                                          render_buffers, \
+                                          buffer_params, \
+                                          destination, \
+                                          kernels.film_convert_##pass); \
+    } \
+\
+    if (destination.pixels_half_rgba) { \
+      run_get_pass_kernel_processor_half_rgba(&kfilm_convert, \
+                                              render_buffers, \
+                                              buffer_params, \
+                                              destination, \
+                                              kernels.film_convert_half_rgba_##pass); \
+    } \
   }
 
 /* Float (scalar) passes. */

intern/cycles/integrator/pass_accessor_cpu.h

@@ -16,6 +16,8 @@
 
 #pragma once
 
+#include "device/cpu/kernel.h"
+
 #include "integrator/pass_accessor.h"
 
 CCL_NAMESPACE_BEGIN
@@ -28,25 +30,19 @@ class PassAccessorCPU : public PassAccessor {
   using PassAccessor::PassAccessor;
 
  protected:
-  template<typename Processor>
-  inline void run_get_pass_kernel_processor(const RenderBuffers *render_buffers,
-                                            const BufferParams &buffer_params,
-                                            const Destination &destination,
-                                            const Processor &processor) const;
+  inline void run_get_pass_kernel_processor_float(
+      const KernelFilmConvert *kfilm_convert,
+      const RenderBuffers *render_buffers,
+      const BufferParams &buffer_params,
+      const Destination &destination,
+      const CPUKernels::FilmConvertFunction func) const;
 
-  template<typename Processor>
-  inline void run_get_pass_kernel_processor_float(const KernelFilmConvert *kfilm_convert,
-                                                  const RenderBuffers *render_buffers,
-                                                  const BufferParams &buffer_params,
-                                                  const Destination &destination,
-                                                  const Processor &processor) const;
-
-  template<typename Processor>
-  inline void run_get_pass_kernel_processor_half_rgba(const KernelFilmConvert *kfilm_convert,
-                                                      const RenderBuffers *render_buffers,
-                                                      const BufferParams &buffer_params,
-                                                      const Destination &destination,
-                                                      const Processor &processor) const;
+  inline void run_get_pass_kernel_processor_half_rgba(
+      const KernelFilmConvert *kfilm_convert,
+      const RenderBuffers *render_buffers,
+      const BufferParams &buffer_params,
+      const Destination &destination,
+      const CPUKernels::FilmConvertHalfRGBAFunction func) const;
 
 #define DECLARE_PASS_ACCESSOR(pass) \
   virtual void get_pass_##pass(const RenderBuffers *render_buffers, \

intern/cycles/integrator/path_trace_work_cpu.cpp

@@ -58,7 +58,7 @@ PathTraceWorkCPU::PathTraceWorkCPU(Device *device,
                                    DeviceScene *device_scene,
                                    bool *cancel_requested_flag)
     : PathTraceWork(device, film, device_scene, cancel_requested_flag),
-      kernels_(*(device->get_cpu_kernels()))
+      kernels_(Device::get_cpu_kernels())
 {
   DCHECK_EQ(device->info.type, DEVICE_CPU);
 }

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -807,10 +807,10 @@ bool PathTraceWorkGPU::should_use_graphics_interop()
     interop_use_ = device->should_use_graphics_interop();
 
     if (interop_use_) {
-      VLOG(2) << "Will be using graphics interop GPU display update.";
+      VLOG(2) << "Using graphics interop GPU display update.";
     }
     else {
-      VLOG(2) << "Will be using naive GPU display update.";
+      VLOG(2) << "Using naive GPU display update.";
     }
 
     interop_use_checked_ = true;

intern/cycles/integrator/shader_eval.cpp

@@ -96,7 +96,7 @@ bool ShaderEval::eval_cpu(Device *device,
   device->get_cpu_kernel_thread_globals(kernel_thread_globals);
 
   /* Find required kernel function. */
-  const CPUKernels &kernels = *(device->get_cpu_kernels());
+  const CPUKernels &kernels = Device::get_cpu_kernels();
 
   /* Simple parallel_for over all work items. */
   KernelShaderEvalInput *input_data = input.data();

intern/cycles/kernel/CMakeLists.txt

@@ -39,6 +39,10 @@ set(SRC_KERNEL_DEVICE_HIP
   device/hip/kernel.cpp
 )
 
+set(SRC_KERNEL_DEVICE_METAL
+  device/metal/kernel.metal
+)
+
 set(SRC_KERNEL_DEVICE_OPTIX
   device/optix/kernel.cu
   device/optix/kernel_shader_raytrace.cu
@@ -79,6 +83,13 @@ set(SRC_KERNEL_DEVICE_OPTIX_HEADERS
   device/optix/globals.h
 )
 
+set(SRC_KERNEL_DEVICE_METAL_HEADERS
+  device/metal/compat.h
+  device/metal/context_begin.h
+  device/metal/context_end.h
+  device/metal/globals.h
+)
+
 set(SRC_KERNEL_CLOSURE_HEADERS
   closure/alloc.h
   closure/bsdf.h
@@ -368,6 +379,7 @@ if(WITH_CYCLES_CUDA_BINARIES)
     ${SRC_KERNEL_HEADERS}
     ${SRC_KERNEL_DEVICE_GPU_HEADERS}
     ${SRC_KERNEL_DEVICE_CUDA_HEADERS}
+    ${SRC_KERNEL_DEVICE_METAL_HEADERS}
     ${SRC_UTIL_HEADERS}
   )
   set(cuda_cubins)
@@ -723,12 +735,14 @@ cycles_add_library(cycles_kernel "${LIB}"
   ${SRC_KERNEL_DEVICE_CUDA}
   ${SRC_KERNEL_DEVICE_HIP}
   ${SRC_KERNEL_DEVICE_OPTIX}
+  ${SRC_KERNEL_DEVICE_METAL}
   ${SRC_KERNEL_HEADERS}
   ${SRC_KERNEL_DEVICE_CPU_HEADERS}
   ${SRC_KERNEL_DEVICE_GPU_HEADERS}
   ${SRC_KERNEL_DEVICE_CUDA_HEADERS}
   ${SRC_KERNEL_DEVICE_HIP_HEADERS}
   ${SRC_KERNEL_DEVICE_OPTIX_HEADERS}
+  ${SRC_KERNEL_DEVICE_METAL_HEADERS}
 )
 
 source_group("bake" FILES ${SRC_KERNEL_BAKE_HEADERS})
@@ -740,6 +754,7 @@ source_group("device\\cuda" FILES ${SRC_KERNEL_DEVICE_CUDA} ${SRC_KERNEL_DEVICE_
 source_group("device\\gpu" FILES ${SRC_KERNEL_DEVICE_GPU_HEADERS})
 source_group("device\\hip" FILES ${SRC_KERNEL_DEVICE_HIP} ${SRC_KERNEL_DEVICE_HIP_HEADERS})
 source_group("device\\optix" FILES ${SRC_KERNEL_DEVICE_OPTIX} ${SRC_KERNEL_DEVICE_OPTIX_HEADERS})
+source_group("device\\metal" FILES ${SRC_KERNEL_DEVICE_METAL} ${SRC_KERNEL_DEVICE_METAL_HEADERS})
 source_group("film" FILES ${SRC_KERNEL_FILM_HEADERS})
 source_group("geom" FILES ${SRC_KERNEL_GEOM_HEADERS})
 source_group("integrator" FILES ${SRC_KERNEL_INTEGRATOR_HEADERS})
@@ -772,6 +787,8 @@ delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_HIP}" ${CYCLES_
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_HIP_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/hip)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_OPTIX}" ${CYCLES_INSTALL_PATH}/source/kernel/device/optix)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_OPTIX_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/optix)
+delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_METAL}" ${CYCLES_INSTALL_PATH}/source/kernel/device/metal)
+delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_METAL_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/metal)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_FILM_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/film)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_GEOM_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/geom)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_INTEGRATOR_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/integrator)

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

@@ -18,6 +18,7 @@
 
 /* CPU Kernel Interface */
 
+#include "util/half.h"
 #include "util/types.h"
 
 #include "kernel/types.h"

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

@@ -52,6 +52,37 @@ KERNEL_INTEGRATOR_SHADE_FUNCTION(megakernel);
 #undef KERNEL_INTEGRATOR_INIT_FUNCTION
 #undef KERNEL_INTEGRATOR_SHADE_FUNCTION
 
+#define KERNEL_FILM_CONVERT_FUNCTION(name) \
+  void KERNEL_FUNCTION_FULL_NAME(film_convert_##name)(const KernelFilmConvert *kfilm_convert, \
+                                                      const float *buffer, \
+                                                      float *pixel, \
+                                                      const int width, \
+                                                      const int buffer_stride, \
+                                                      const int pixel_stride); \
+  void KERNEL_FUNCTION_FULL_NAME(film_convert_half_rgba_##name)( \
+      const KernelFilmConvert *kfilm_convert, \
+      const float *buffer, \
+      half4 *pixel, \
+      const int width, \
+      const int buffer_stride);
+
+KERNEL_FILM_CONVERT_FUNCTION(depth)
+KERNEL_FILM_CONVERT_FUNCTION(mist)
+KERNEL_FILM_CONVERT_FUNCTION(sample_count)
+KERNEL_FILM_CONVERT_FUNCTION(float)
+
+KERNEL_FILM_CONVERT_FUNCTION(light_path)
+KERNEL_FILM_CONVERT_FUNCTION(float3)
+
+KERNEL_FILM_CONVERT_FUNCTION(motion)
+KERNEL_FILM_CONVERT_FUNCTION(cryptomatte)
+KERNEL_FILM_CONVERT_FUNCTION(shadow_catcher)
+KERNEL_FILM_CONVERT_FUNCTION(shadow_catcher_matte_with_shadow)
+KERNEL_FILM_CONVERT_FUNCTION(combined)
+KERNEL_FILM_CONVERT_FUNCTION(float4)
+
+#undef KERNEL_FILM_CONVERT_FUNCTION
+
 /* --------------------------------------------------------------------
  * Shader evaluation.
  */

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

@@ -47,8 +47,8 @@
 #    include "kernel/integrator/megakernel.h"
 
 #    include "kernel/film/adaptive_sampling.h"
-#    include "kernel/film/read.h"
 #    include "kernel/film/id_passes.h"
+#    include "kernel/film/read.h"
 
 #    include "kernel/bake/bake.h"
 
@@ -232,6 +232,85 @@ void KERNEL_FUNCTION_FULL_NAME(cryptomatte_postprocess)(const KernelGlobalsCPU *
 #endif
 }
 
+/* --------------------------------------------------------------------
+ * Film Convert.
+ */
+
+#ifdef KERNEL_STUB
+
+#  define KERNEL_FILM_CONVERT_FUNCTION(name, is_float) \
+    void KERNEL_FUNCTION_FULL_NAME(film_convert_##name)(const KernelFilmConvert *kfilm_convert, \
+                                                        const float *buffer, \
+                                                        float *pixel, \
+                                                        const int width, \
+                                                        const int buffer_stride, \
+                                                        const int pixel_stride) \
+    { \
+      STUB_ASSERT(KERNEL_ARCH, film_convert_##name); \
+    } \
+    void KERNEL_FUNCTION_FULL_NAME(film_convert_half_rgba_##name)( \
+        const KernelFilmConvert *kfilm_convert, \
+        const float *buffer, \
+        half4 *pixel, \
+        const int width, \
+        const int buffer_stride) \
+    { \
+      STUB_ASSERT(KERNEL_ARCH, film_convert_##name); \
+    }
+
+#else
+
+#  define KERNEL_FILM_CONVERT_FUNCTION(name, is_float) \
+    void KERNEL_FUNCTION_FULL_NAME(film_convert_##name)(const KernelFilmConvert *kfilm_convert, \
+                                                        const float *buffer, \
+                                                        float *pixel, \
+                                                        const int width, \
+                                                        const int buffer_stride, \
+                                                        const int pixel_stride) \
+    { \
+      for (int i = 0; i < width; i++, buffer += buffer_stride, pixel += pixel_stride) { \
+        film_get_pass_pixel_##name(kfilm_convert, buffer, pixel); \
+      } \
+    } \
+    void KERNEL_FUNCTION_FULL_NAME(film_convert_half_rgba_##name)( \
+        const KernelFilmConvert *kfilm_convert, \
+        const float *buffer, \
+        half4 *pixel, \
+        const int width, \
+        const int buffer_stride) \
+    { \
+      for (int i = 0; i < width; i++, buffer += buffer_stride, pixel++) { \
+        float pixel_rgba[4] = {0.0f, 0.0f, 0.0f, 1.0f}; \
+        film_get_pass_pixel_##name(kfilm_convert, buffer, pixel_rgba); \
+        if (is_float) { \
+          pixel_rgba[1] = pixel_rgba[0]; \
+          pixel_rgba[2] = pixel_rgba[0]; \
+        } \
+        film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel_rgba); \
+        *pixel = float4_to_half4_display( \
+            make_float4(pixel_rgba[0], pixel_rgba[1], pixel_rgba[2], pixel_rgba[3])); \
+      } \
+    }
+
+#endif
+
+KERNEL_FILM_CONVERT_FUNCTION(depth, true)
+KERNEL_FILM_CONVERT_FUNCTION(mist, true)
+KERNEL_FILM_CONVERT_FUNCTION(sample_count, true)
+KERNEL_FILM_CONVERT_FUNCTION(float, true)
+
+KERNEL_FILM_CONVERT_FUNCTION(light_path, false)
+KERNEL_FILM_CONVERT_FUNCTION(float3, false)
+
+KERNEL_FILM_CONVERT_FUNCTION(motion, false)
+KERNEL_FILM_CONVERT_FUNCTION(cryptomatte, false)
+KERNEL_FILM_CONVERT_FUNCTION(shadow_catcher, false)
+KERNEL_FILM_CONVERT_FUNCTION(shadow_catcher_matte_with_shadow, false)
+KERNEL_FILM_CONVERT_FUNCTION(combined, false)
+KERNEL_FILM_CONVERT_FUNCTION(float4, false)
+
+#undef KERNEL_FILM_CONVERT_FUNCTION
+
 #undef KERNEL_INVOKE
 #undef DEFINE_INTEGRATOR_KERNEL
 #undef DEFINE_INTEGRATOR_SHADE_KERNEL

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

@@ -75,6 +75,7 @@ typedef unsigned long long uint64_t;
 #define ccl_gpu_block_idx_x (blockIdx.x)
 #define ccl_gpu_grid_dim_x (gridDim.x)
 #define ccl_gpu_warp_size (warpSize)
+#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
 
 #define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
 #define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)

intern/cycles/kernel/device/cuda/config.h

@@ -93,11 +93,35 @@
 /* Compute number of threads per block and minimum blocks per multiprocessor
  * given the maximum number of registers per thread. */
 
-#define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
+#define ccl_gpu_kernel_threads(block_num_threads) \
+  extern "C" __global__ void __launch_bounds__(block_num_threads)
+
+#define ccl_gpu_kernel_threads_registers(block_num_threads, thread_num_registers) \
   extern "C" __global__ void __launch_bounds__(block_num_threads, \
                                                GPU_MULTIPRESSOR_MAX_REGISTERS / \
                                                    (block_num_threads * thread_num_registers))
 
+/* allow ccl_gpu_kernel to accept 1 or 2 parameters */
+#define SELECT_MACRO(_1, _2, NAME, ...) NAME
+#define ccl_gpu_kernel(...) \
+  SELECT_MACRO(__VA_ARGS__, ccl_gpu_kernel_threads_registers, ccl_gpu_kernel_threads)(__VA_ARGS__)
+
+#define ccl_gpu_kernel_signature(name, ...) kernel_gpu_##name(__VA_ARGS__)
+
+#define ccl_gpu_kernel_call(x) x
+
+/* define a function object where "func" is the lambda body, and additional parameters are used to
+ * specify captured state  */
+#define ccl_gpu_kernel_lambda(func, ...) \
+  struct KernelLambda { \
+    __VA_ARGS__; \
+    __device__ int operator()(const int state) \
+    { \
+      return (func); \
+    } \
+  } ccl_gpu_kernel_lambda_pass; \
+  ccl_gpu_kernel_lambda_pass
+
 /* sanity checks */
 
 #if GPU_KERNEL_BLOCK_NUM_THREADS > GPU_BLOCK_MAX_THREADS

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

@@ -65,7 +65,9 @@ ccl_device float cubic_h1(float a)
 
 /* Fast bicubic texture lookup using 4 bilinear lookups, adapted from CUDA samples. */
 template<typename T>
-ccl_device_noinline T kernel_tex_image_interp_bicubic(const TextureInfo &info, float x, float y)
+ccl_device_noinline T kernel_tex_image_interp_bicubic(ccl_global const TextureInfo &info,
+                                                      float x,
+                                                      float y)
 {
   ccl_gpu_tex_object tex = (ccl_gpu_tex_object)info.data;
 
@@ -94,7 +96,7 @@ ccl_device_noinline T kernel_tex_image_interp_bicubic(const TextureInfo &info, f
 /* Fast tricubic texture lookup using 8 trilinear lookups. */
 template<typename T>
 ccl_device_noinline T
-kernel_tex_image_interp_tricubic(const TextureInfo &info, float x, float y, float z)
+kernel_tex_image_interp_tricubic(ccl_global const TextureInfo &info, float x, float y, float z)
 {
   ccl_gpu_tex_object tex = (ccl_gpu_tex_object)info.data;
 
@@ -169,7 +171,7 @@ ccl_device T kernel_tex_image_interp_tricubic_nanovdb(S &s, float x, float y, fl
 
 template<typename T>
 ccl_device_noinline T kernel_tex_image_interp_nanovdb(
-    const TextureInfo &info, float x, float y, float z, uint interpolation)
+    ccl_global const TextureInfo &info, float x, float y, float z, uint interpolation)
 {
   using namespace nanovdb;
 
@@ -191,7 +193,7 @@ ccl_device_noinline T kernel_tex_image_interp_nanovdb(
 
 ccl_device float4 kernel_tex_image_interp(KernelGlobals kg, int id, float x, float y)
 {
-  const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
+  ccl_global const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
 
   /* float4, byte4, ushort4 and half4 */
   const int texture_type = info.data_type;
@@ -226,7 +228,7 @@ ccl_device float4 kernel_tex_image_interp_3d(KernelGlobals kg,
                                              float3 P,
                                              InterpolationType interp)
 {
-  const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
+  ccl_global const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
 
   if (info.use_transform_3d) {
     P = transform_point(&info.transform_3d, P);

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

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

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

@@ -33,10 +33,12 @@ CCL_NAMESPACE_BEGIN
 #  define GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE 512
 #endif
 
-template<uint blocksize>
-__device__ void gpu_parallel_prefix_sum(int *counter, int *prefix_sum, const int num_values)
+__device__ void gpu_parallel_prefix_sum(const int global_id,
+                                        ccl_global int *counter,
+                                        ccl_global int *prefix_sum,
+                                        const int num_values)
 {
-  if (!(ccl_gpu_block_idx_x == 0 && ccl_gpu_thread_idx_x == 0)) {
+  if (global_id != 0) {
     return;
   }
 

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

@@ -33,16 +33,16 @@ CCL_NAMESPACE_BEGIN
 #endif
 #define GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY (~0)
 
-template<uint blocksize, typename GetKeyOp>
-__device__ void gpu_parallel_sorted_index_array(const uint num_states,
+template<typename GetKeyOp>
+__device__ void gpu_parallel_sorted_index_array(const uint state_index,
+                                                const uint num_states,
                                                 const int num_states_limit,
-                                                int *indices,
-                                                int *num_indices,
-                                                int *key_counter,
-                                                int *key_prefix_sum,
+                                                ccl_global int *indices,
+                                                ccl_global int *num_indices,
+                                                ccl_global int *key_counter,
+                                                ccl_global int *key_prefix_sum,
                                                 GetKeyOp get_key_op)
 {
-  const uint state_index = ccl_gpu_block_idx_x * blocksize + ccl_gpu_thread_idx_x;
   const int key = (state_index < num_states) ? get_key_op(state_index) :
                                                GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY;
 

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

@@ -74,6 +74,7 @@ typedef unsigned long long uint64_t;
 #define ccl_gpu_block_idx_x (blockIdx.x)
 #define ccl_gpu_grid_dim_x (gridDim.x)
 #define ccl_gpu_warp_size (warpSize)
+#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
 
 #define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
 #define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)

intern/cycles/kernel/device/hip/config.h

@@ -36,11 +36,35 @@
 /* Compute number of threads per block and minimum blocks per multiprocessor
  * given the maximum number of registers per thread. */
 
-#define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
+#define ccl_gpu_kernel_threads(block_num_threads) \
+  extern "C" __global__ void __launch_bounds__(block_num_threads)
+
+#define ccl_gpu_kernel_threads_registers(block_num_threads, thread_num_registers) \
   extern "C" __global__ void __launch_bounds__(block_num_threads, \
                                                GPU_MULTIPRESSOR_MAX_REGISTERS / \
                                                    (block_num_threads * thread_num_registers))
 
+/* allow ccl_gpu_kernel to accept 1 or 2 parameters */
+#define SELECT_MACRO(_1, _2, NAME, ...) NAME
+#define ccl_gpu_kernel(...) \
+  SELECT_MACRO(__VA_ARGS__, ccl_gpu_kernel_threads_registers, ccl_gpu_kernel_threads)(__VA_ARGS__)
+
+#define ccl_gpu_kernel_signature(name, ...) kernel_gpu_##name(__VA_ARGS__)
+
+#define ccl_gpu_kernel_call(x) x
+
+/* define a function object where "func" is the lambda body, and additional parameters are used to
+ * specify captured state  */
+#define ccl_gpu_kernel_lambda(func, ...) \
+  struct KernelLambda { \
+    __VA_ARGS__; \
+    __device__ int operator()(const int state) \
+    { \
+      return (func); \
+    } \
+  } ccl_gpu_kernel_lambda_pass; \
+  ccl_gpu_kernel_lambda_pass
+
 /* sanity checks */
 
 #if GPU_KERNEL_BLOCK_NUM_THREADS > GPU_BLOCK_MAX_THREADS

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

@@ -58,6 +58,95 @@ using namespace metal;
 
 #define kernel_assert(cond)
 
+#define ccl_gpu_global_id_x() metal_global_id
+#define ccl_gpu_warp_size simdgroup_size
+#define ccl_gpu_thread_idx_x simd_group_index
+#define ccl_gpu_thread_mask(thread_warp) uint64_t((1ull << thread_warp) - 1)
+
+#define ccl_gpu_ballot(predicate) ((uint64_t)((simd_vote::vote_t)simd_ballot(predicate)))
+#define ccl_gpu_popc(x) popcount(x)
+
+// clang-format off
+
+/* kernel.h adapters */
+
+#define ccl_gpu_kernel(...)
+
+/* convert a comma-separated list into a semicolon-separated list (so that we can generate a struct based on kernel entrypoint parameters) */
+#define FN0()
+#define FN1(p1) p1;
+#define FN2(p1, p2) p1; p2;
+#define FN3(p1, p2, p3) p1; p2; p3;
+#define FN4(p1, p2, p3, p4) p1; p2; p3; p4;
+#define FN5(p1, p2, p3, p4, p5) p1; p2; p3; p4; p5;
+#define FN6(p1, p2, p3, p4, p5, p6) p1; p2; p3; p4; p5; p6;
+#define FN7(p1, p2, p3, p4, p5, p6, p7) p1; p2; p3; p4; p5; p6; p7;
+#define FN8(p1, p2, p3, p4, p5, p6, p7, p8) p1; p2; p3; p4; p5; p6; p7; p8;
+#define FN9(p1, p2, p3, p4, p5, p6, p7, p8, p9) p1; p2; p3; p4; p5; p6; p7; p8; p9;
+#define FN10(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10;
+#define FN11(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11;
+#define FN12(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12;
+#define FN13(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13;
+#define FN14(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14;
+#define FN15(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15;
+#define FN16(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16;
+#define GET_LAST_ARG(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, ...) p16
+#define PARAMS_MAKER(...) GET_LAST_ARG(__VA_ARGS__, FN16, FN15, FN14, FN13, FN12, FN11, FN10, FN9, FN8, FN7, FN6, FN5, FN4, FN3, FN2, FN1, FN0)
+
+/* generate a struct containing the entrypoint parameters and a "run" method which can access them implicitly via this-> */
+#define ccl_gpu_kernel_signature(name, ...) \
+struct kernel_gpu_##name \
+{ \
+  PARAMS_MAKER(__VA_ARGS__)(__VA_ARGS__) \
+  void run(thread MetalKernelContext& context, \
+           threadgroup int *simdgroup_offset, \
+           const uint metal_global_id, \
+           const ushort metal_local_id, \
+           const ushort metal_local_size, \
+           uint simdgroup_size, \
+           uint simd_lane_index, \
+           uint simd_group_index, \
+           uint num_simd_groups) ccl_global const; \
+}; \
+kernel void kernel_metal_##name(device const kernel_gpu_##name *params_struct, \
+                                constant KernelParamsMetal &ccl_restrict   _launch_params_metal, \
+                                constant MetalAncillaries *_metal_ancillaries, \
+                                threadgroup int *simdgroup_offset[[ threadgroup(0) ]], \
+                                const uint metal_global_id [[thread_position_in_grid]], \
+                                const ushort metal_local_id   [[thread_position_in_threadgroup]], \
+                                const ushort metal_local_size [[threads_per_threadgroup]], \
+                                uint simdgroup_size [[threads_per_simdgroup]], \
+                                uint simd_lane_index [[thread_index_in_simdgroup]], \
+                                uint simd_group_index [[simdgroup_index_in_threadgroup]], \
+                                uint num_simd_groups [[simdgroups_per_threadgroup]]) { \
+  MetalKernelContext context(_launch_params_metal, _metal_ancillaries); \
+  INIT_DEBUG_BUFFER \
+  params_struct->run(context, simdgroup_offset, metal_global_id, metal_local_id, metal_local_size, simdgroup_size, simd_lane_index, simd_group_index, num_simd_groups); \
+} \
+void kernel_gpu_##name::run(thread MetalKernelContext& context, \
+                  threadgroup int *simdgroup_offset, \
+                  const uint metal_global_id, \
+                  const ushort metal_local_id, \
+                  const ushort metal_local_size, \
+                  uint simdgroup_size, \
+                  uint simd_lane_index, \
+                  uint simd_group_index, \
+                  uint num_simd_groups) ccl_global const
+
+#define ccl_gpu_kernel_call(x) context.x
+
+/* define a function object where "func" is the lambda body, and additional parameters are used to specify captured state  */
+#define ccl_gpu_kernel_lambda(func, ...) \
+  struct KernelLambda \
+  { \
+    KernelLambda(ccl_private MetalKernelContext &_context) : context(_context) {} \
+    ccl_private MetalKernelContext &context; \
+    __VA_ARGS__; \
+    int operator()(const int state) const { return (func); } \
+  }ccl_gpu_kernel_lambda_pass(context); ccl_gpu_kernel_lambda_pass
+
+// clang-format on
+
 /* make_type definitions with Metal style element initializers */
 #ifdef make_float2
 #  undef make_float2
@@ -124,3 +213,38 @@ using namespace metal;
 #define logf(x) trigmode::log(float(x))
 
 #define NULL 0
+
+/* texture bindings and sampler setup */
+
+struct Texture2DParamsMetal {
+  texture2d<float, access::sample> tex;
+};
+struct Texture3DParamsMetal {
+  texture3d<float, access::sample> tex;
+};
+
+struct MetalAncillaries {
+  device Texture2DParamsMetal *textures_2d;
+  device Texture3DParamsMetal *textures_3d;
+};
+
+enum SamplerType {
+  SamplerFilterNearest_AddressRepeat,
+  SamplerFilterNearest_AddressClampEdge,
+  SamplerFilterNearest_AddressClampZero,
+
+  SamplerFilterLinear_AddressRepeat,
+  SamplerFilterLinear_AddressClampEdge,
+  SamplerFilterLinear_AddressClampZero,
+
+  SamplerCount
+};
+
+constant constexpr array<sampler, SamplerCount> metal_samplers = {
+    sampler(address::repeat, filter::nearest),
+    sampler(address::clamp_to_edge, filter::nearest),
+    sampler(address::clamp_to_zero, filter::nearest),
+    sampler(address::repeat, filter::linear),
+    sampler(address::clamp_to_edge, filter::linear),
+    sampler(address::clamp_to_zero, filter::linear),
+};

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

@@ -0,0 +1,79 @@
+/*
+ * Copyright 2021 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+// clang-format off
+
+/* Open the Metal kernel context class
+ * Necessary to access resource bindings */
+class MetalKernelContext {
+  public:
+    constant KernelParamsMetal &launch_params_metal;
+    constant MetalAncillaries *metal_ancillaries;
+
+    MetalKernelContext(constant KernelParamsMetal &_launch_params_metal, constant MetalAncillaries * _metal_ancillaries)
+    : launch_params_metal(_launch_params_metal), metal_ancillaries(_metal_ancillaries)
+    {}
+
+    /* texture fetch adapter functions */
+    typedef uint64_t ccl_gpu_tex_object;
+
+    template<typename T>
+    inline __attribute__((__always_inline__))
+    T ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
+      kernel_assert(0);
+      return 0;
+    }
+    template<typename T>
+    inline __attribute__((__always_inline__))
+    T ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
+      kernel_assert(0);
+      return 0;
+    }
+
+    // texture2d
+    template<>
+    inline __attribute__((__always_inline__))
+    float4 ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
+      const uint tid(tex);
+      const uint sid(tex >> 32);
+      return metal_ancillaries->textures_2d[tid].tex.sample(metal_samplers[sid], float2(x, y));
+    }
+    template<>
+    inline __attribute__((__always_inline__))
+    float ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
+      const uint tid(tex);
+      const uint sid(tex >> 32);
+      return metal_ancillaries->textures_2d[tid].tex.sample(metal_samplers[sid], float2(x, y)).x;
+    }
+
+    // texture3d
+    template<>
+    inline __attribute__((__always_inline__))
+    float4 ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
+      const uint tid(tex);
+      const uint sid(tex >> 32);
+      return metal_ancillaries->textures_3d[tid].tex.sample(metal_samplers[sid], float3(x, y, z));
+    }
+    template<>
+    inline __attribute__((__always_inline__))
+    float ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
+      const uint tid(tex);
+      const uint sid(tex >> 32);
+      return metal_ancillaries->textures_3d[tid].tex.sample(metal_samplers[sid], float3(x, y, z)).x;
+    }
+#    include "kernel/device/gpu/image.h"
+
+  // clang-format on

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

@@ -0,0 +1,23 @@
+/*
+ * Copyright 2021 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+}
+; /* end of MetalKernelContext class definition */
+
+/* Silently redirect into the MetalKernelContext instance */
+/* NOTE: These macros will need maintaining as entrypoints change */
+
+#undef kernel_integrator_state
+#define kernel_integrator_state context.launch_params_metal.__integrator_state

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

@@ -0,0 +1,51 @@
+/*
+ * Copyright 2021 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+/* Constant Globals */
+
+#include "kernel/types.h"
+#include "kernel/util/profiling.h"
+
+#include "kernel/integrator/state.h"
+
+CCL_NAMESPACE_BEGIN
+
+typedef struct KernelParamsMetal {
+
+#define KERNEL_TEX(type, name) ccl_constant type *name;
+#include "kernel/textures.h"
+#undef KERNEL_TEX
+
+  const IntegratorStateGPU __integrator_state;
+  const KernelData data;
+
+} KernelParamsMetal;
+
+typedef struct KernelGlobalsGPU {
+  int unused[1];
+} KernelGlobalsGPU;
+
+typedef ccl_global const KernelGlobalsGPU *ccl_restrict KernelGlobals;
+
+#define kernel_data launch_params_metal.data
+#define kernel_integrator_state launch_params_metal.__integrator_state
+
+/* data lookup defines */
+
+#define kernel_tex_fetch(tex, index) launch_params_metal.tex[index]
+#define kernel_tex_array(tex) launch_params_metal.tex
+
+CCL_NAMESPACE_END

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

@@ -0,0 +1,25 @@
+/*
+ * Copyright 2021 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+/* Metal kernel entry points */
+
+// clang-format off
+
+#include "kernel/device/metal/compat.h"
+#include "kernel/device/metal/globals.h"
+#include "kernel/device/gpu/kernel.h"
+
+// clang-format on

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

@@ -76,6 +76,7 @@ typedef unsigned long long uint64_t;
 #define ccl_gpu_block_idx_x (blockIdx.x)
 #define ccl_gpu_grid_dim_x (gridDim.x)
 #define ccl_gpu_warp_size (warpSize)
+#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
 
 #define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
 #define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)

intern/cycles/kernel/film/accumulate.h

@@ -540,11 +540,10 @@ ccl_device_inline void kernel_accum_background(KernelGlobals kg,
 /* Write emission to render buffer. */
 ccl_device_inline void kernel_accum_emission(KernelGlobals kg,
                                              ConstIntegratorState state,
-                                             const float3 throughput,
                                              const float3 L,
                                              ccl_global float *ccl_restrict render_buffer)
 {
-  float3 contribution = throughput * L;
+  float3 contribution = L;
   kernel_accum_clamp(kg, &contribution, INTEGRATOR_STATE(state, path, bounce) - 1);
 
   ccl_global float *buffer = kernel_accum_pixel_render_buffer(kg, state, render_buffer);

intern/cycles/kernel/integrator/intersect_closest.h

@@ -31,7 +31,6 @@
 
 CCL_NAMESPACE_BEGIN
 
-template<uint32_t current_kernel>
 ccl_device_forceinline bool integrator_intersect_terminate(KernelGlobals kg,
                                                            IntegratorState state,
                                                            const int shader_flags)
@@ -63,6 +62,7 @@ ccl_device_forceinline bool integrator_intersect_terminate(KernelGlobals kg,
    * perform MIS as part of indirect rays. */
   const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
   const float probability = path_state_continuation_probability(kg, state, path_flag);
+  INTEGRATOR_STATE_WRITE(state, path, continuation_probability) = probability;
 
   if (probability != 1.0f) {
     const float terminate = path_state_rng_1D(kg, &rng_state, PRNG_TERMINATE);
@@ -85,36 +85,75 @@ ccl_device_forceinline bool integrator_intersect_terminate(KernelGlobals kg,
   return false;
 }
 
-/* Note that current_kernel is a template value since making this a variable
- * leads to poor performance with CUDA atomics. */
-template<uint32_t current_kernel>
-ccl_device_forceinline void integrator_intersect_shader_next_kernel(
-    KernelGlobals kg,
-    IntegratorState state,
-    ccl_private const Intersection *ccl_restrict isect,
-    const int shader,
-    const int shader_flags)
+#ifdef __SHADOW_CATCHER__
+/* Split path if a shadow catcher was hit. */
+ccl_device_forceinline void integrator_split_shadow_catcher(
+    KernelGlobals kg, IntegratorState state, ccl_private const Intersection *ccl_restrict isect)
 {
-  /* Note on scheduling.
-   *
-   * When there is no shadow catcher split the scheduling is simple: schedule surface shading with
-   * or without raytrace support, depending on the shader used.
-   *
-   * When there is a shadow catcher split the general idea is to have the following configuration:
-   *
-   *  - Schedule surface shading kernel (with corresponding raytrace support) for the ray which
-   *    will trace shadow catcher object.
-   *
-   *  - When no alpha-over of approximate shadow catcher is needed, schedule surface shading for
-   *    the matte ray.
-   *
-   *  - Otherwise schedule background shading kernel, so that we have a background to alpha-over
-   *    on. The background kernel will then schedule surface shading for the matte ray.
+  /* Test if we hit a shadow catcher object, and potentially split the path to continue tracing two
+   * paths from here. */
+  const int object_flags = intersection_get_object_flags(kg, isect);
+  if (!kernel_shadow_catcher_is_path_split_bounce(kg, state, object_flags)) {
+    return;
+  }
+
+  /* Mark state as having done a shadow catcher split so that it stops contributing to
+   * the shadow catcher matte pass, but keeps contributing to the combined pass. */
+  INTEGRATOR_STATE_WRITE(state, path, flag) |= PATH_RAY_SHADOW_CATCHER_HIT;
+
+  /* Copy current state to new state. */
+  state = integrator_state_shadow_catcher_split(kg, state);
+
+  /* Initialize new state.
    *
    * Note that the splitting leaves kernel and sorting counters as-is, so use INIT semantic for
    * the matte path. */
 
-  const bool use_raytrace_kernel = (shader_flags & SD_HAS_RAYTRACE);
+  /* Mark current state so that it will only track contribution of shadow catcher objects ignoring
+   * non-catcher objects. */
+  INTEGRATOR_STATE_WRITE(state, path, flag) |= PATH_RAY_SHADOW_CATCHER_PASS;
+
+  if (kernel_data.film.pass_background != PASS_UNUSED && !kernel_data.background.transparent) {
+    /* If using background pass, schedule background shading kernel so that we have a background
+     * to alpha-over on. The background kernel will then continue the path afterwards. */
+    INTEGRATOR_STATE_WRITE(state, path, flag) |= PATH_RAY_SHADOW_CATCHER_BACKGROUND;
+    INTEGRATOR_PATH_INIT(DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND);
+    return;
+  }
+
+  if (!integrator_state_volume_stack_is_empty(kg, state)) {
+    /* Volume stack is not empty. Re-init the volume stack to exclude any non-shadow catcher
+     * objects from it, and then continue shading volume and shadow catcher surface after. */
+    INTEGRATOR_PATH_INIT(DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK);
+    return;
+  }
+
+  /* Continue with shading shadow catcher surface. */
+  const int shader = intersection_get_shader(kg, isect);
+  const int flags = kernel_tex_fetch(__shaders, shader).flags;
+  const bool use_raytrace_kernel = (flags & SD_HAS_RAYTRACE);
+
+  if (use_raytrace_kernel) {
+    INTEGRATOR_PATH_INIT_SORTED(DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE, shader);
+  }
+  else {
+    INTEGRATOR_PATH_INIT_SORTED(DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE, shader);
+  }
+}
+
+/* Schedule next kernel to be executed after updating volume stack for shadow catcher. */
+template<uint32_t current_kernel>
+ccl_device_forceinline void integrator_intersect_next_kernel_after_shadow_catcher_volume(
+    KernelGlobals kg, IntegratorState state)
+{
+  /* Continue with shading shadow catcher surface. Same as integrator_split_shadow_catcher, but
+   * using NEXT instead of INIT. */
+  Intersection isect ccl_optional_struct_init;
+  integrator_state_read_isect(kg, state, &isect);
+
+  const int shader = intersection_get_shader(kg, &isect);
+  const int flags = kernel_tex_fetch(__shaders, shader).flags;
+  const bool use_raytrace_kernel = (flags & SD_HAS_RAYTRACE);
 
   if (use_raytrace_kernel) {
     INTEGRATOR_PATH_NEXT_SORTED(
@@ -123,23 +162,132 @@ ccl_device_forceinline void integrator_intersect_shader_next_kernel(
   else {
     INTEGRATOR_PATH_NEXT_SORTED(current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE, shader);
   }
+}
 
-#ifdef __SHADOW_CATCHER__
-  const int object_flags = intersection_get_object_flags(kg, isect);
-  if (kernel_shadow_catcher_split(kg, state, object_flags)) {
-    if (kernel_data.film.pass_background != PASS_UNUSED && !kernel_data.background.transparent) {
-      INTEGRATOR_STATE_WRITE(state, path, flag) |= PATH_RAY_SHADOW_CATCHER_BACKGROUND;
+/* Schedule next kernel to be executed after executing background shader for shadow catcher. */
+template<uint32_t current_kernel>
+ccl_device_forceinline void integrator_intersect_next_kernel_after_shadow_catcher_background(
+    KernelGlobals kg, IntegratorState state)
+{
+  /* Same logic as integrator_split_shadow_catcher, but using NEXT instead of INIT. */
+  if (!integrator_state_volume_stack_is_empty(kg, state)) {
+    /* Volume stack is not empty. Re-init the volume stack to exclude any non-shadow catcher
+     * objects from it, and then continue shading volume and shadow catcher surface after. */
+    INTEGRATOR_PATH_NEXT(current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK);
+    return;
+  }
 
-      INTEGRATOR_PATH_INIT(DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND);
-    }
-    else if (use_raytrace_kernel) {
-      INTEGRATOR_PATH_INIT_SORTED(DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE, shader);
+  /* Continue with shading shadow catcher surface. */
+  integrator_intersect_next_kernel_after_shadow_catcher_volume<current_kernel>(kg, state);
+}
+#endif
+
+/* Schedule next kernel to be executed after intersect closest.
+ *
+ * Note that current_kernel is a template value since making this a variable
+ * leads to poor performance with CUDA atomics. */
+template<uint32_t current_kernel>
+ccl_device_forceinline void integrator_intersect_next_kernel(
+    KernelGlobals kg,
+    IntegratorState state,
+    ccl_private const Intersection *ccl_restrict isect,
+    const bool hit)
+{
+  /* Continue with volume kernel if we are inside a volume, regardless if we hit anything. */
+#ifdef __VOLUME__
+  if (!integrator_state_volume_stack_is_empty(kg, state)) {
+    const bool hit_surface = hit && !(isect->type & PRIMITIVE_LAMP);
+    const int shader = (hit_surface) ? intersection_get_shader(kg, isect) : SHADER_NONE;
+    const int flags = (hit_surface) ? kernel_tex_fetch(__shaders, shader).flags : 0;
+
+    if (!integrator_intersect_terminate(kg, state, flags)) {
+      INTEGRATOR_PATH_NEXT(current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME);
     }
     else {
-      INTEGRATOR_PATH_INIT_SORTED(DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE, shader);
+      INTEGRATOR_PATH_TERMINATE(current_kernel);
     }
+    return;
   }
 #endif
+
+  if (hit) {
+    /* Hit a surface, continue with light or surface kernel. */
+    if (isect->type & PRIMITIVE_LAMP) {
+      INTEGRATOR_PATH_NEXT(current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT);
+    }
+    else {
+      /* Hit a surface, continue with surface kernel unless terminated. */
+      const int shader = intersection_get_shader(kg, isect);
+      const int flags = kernel_tex_fetch(__shaders, shader).flags;
+
+      if (!integrator_intersect_terminate(kg, state, flags)) {
+        const bool use_raytrace_kernel = (flags & SD_HAS_RAYTRACE);
+        if (use_raytrace_kernel) {
+          INTEGRATOR_PATH_NEXT_SORTED(
+              current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE, shader);
+        }
+        else {
+          INTEGRATOR_PATH_NEXT_SORTED(
+              current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE, shader);
+        }
+
+#ifdef __SHADOW_CATCHER__
+        /* Handle shadow catcher. */
+        integrator_split_shadow_catcher(kg, state, isect);
+#endif
+      }
+      else {
+        INTEGRATOR_PATH_TERMINATE(current_kernel);
+      }
+    }
+  }
+  else {
+    /* Nothing hit, continue with background kernel. */
+    INTEGRATOR_PATH_NEXT(current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND);
+  }
+}
+
+/* Schedule next kernel to be executed after shade volume.
+ *
+ * The logic here matches integrator_intersect_next_kernel, except that
+ * volume shading and termination testing have already been done. */
+template<uint32_t current_kernel>
+ccl_device_forceinline void integrator_intersect_next_kernel_after_volume(
+    KernelGlobals kg, IntegratorState state, ccl_private const Intersection *ccl_restrict isect)
+{
+  if (isect->prim != PRIM_NONE) {
+    /* Hit a surface, continue with light or surface kernel. */
+    if (isect->type & PRIMITIVE_LAMP) {
+      INTEGRATOR_PATH_NEXT(current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT);
+      return;
+    }
+    else {
+      /* Hit a surface, continue with surface kernel unless terminated. */
+      const int shader = intersection_get_shader(kg, isect);
+      const int flags = kernel_tex_fetch(__shaders, shader).flags;
+      const bool use_raytrace_kernel = (flags & SD_HAS_RAYTRACE);
+
+      if (use_raytrace_kernel) {
+        INTEGRATOR_PATH_NEXT_SORTED(
+            current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE, shader);
+      }
+      else {
+        INTEGRATOR_PATH_NEXT_SORTED(
+            current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE, shader);
+      }
+
+#ifdef __SHADOW_CATCHER__
+      /* Handle shadow catcher. */
+      integrator_split_shadow_catcher(kg, state, isect);
+#endif
+      return;
+    }
+  }
+  else {
+    /* Nothing hit, continue with background kernel. */
+    INTEGRATOR_PATH_NEXT(current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND);
+    return;
+  }
 }
 
 ccl_device void integrator_intersect_closest(KernelGlobals kg, IntegratorState state)
@@ -191,56 +339,9 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg, IntegratorState s
   /* Write intersection result into global integrator state memory. */
   integrator_state_write_isect(kg, state, &isect);
 
-#ifdef __VOLUME__
-  if (!integrator_state_volume_stack_is_empty(kg, state)) {
-    const bool hit_surface = hit && !(isect.type & PRIMITIVE_LAMP);
-    const int shader = (hit_surface) ? intersection_get_shader(kg, &isect) : SHADER_NONE;
-    const int flags = (hit_surface) ? kernel_tex_fetch(__shaders, shader).flags : 0;
-
-    if (!integrator_intersect_terminate<DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST>(
-            kg, state, flags)) {
-      /* Continue with volume kernel if we are inside a volume, regardless
-       * if we hit anything. */
-      INTEGRATOR_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST,
-                           DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME);
-    }
-    else {
-      INTEGRATOR_PATH_TERMINATE(DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
-    }
-    return;
-  }
-#endif
-
-  if (hit) {
-    /* Hit a surface, continue with light or surface kernel. */
-    if (isect.type & PRIMITIVE_LAMP) {
-      INTEGRATOR_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST,
-                           DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT);
-      return;
-    }
-    else {
-      /* Hit a surface, continue with surface kernel unless terminated. */
-      const int shader = intersection_get_shader(kg, &isect);
-      const int flags = kernel_tex_fetch(__shaders, shader).flags;
-
-      if (!integrator_intersect_terminate<DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST>(
-              kg, state, flags)) {
-        integrator_intersect_shader_next_kernel<DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST>(
-            kg, state, &isect, shader, flags);
-        return;
-      }
-      else {
-        INTEGRATOR_PATH_TERMINATE(DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
-        return;
-      }
-    }
-  }
-  else {
-    /* Nothing hit, continue with background kernel. */
-    INTEGRATOR_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST,
-                         DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND);
-    return;
-  }
+  /* Setup up next kernel to be executed. */
+  integrator_intersect_next_kernel<DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST>(
+      kg, state, &isect, hit);
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/integrator/intersect_volume_stack.h

@@ -42,10 +42,13 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
   /* Store to avoid global fetches on every intersection step. */
   const uint volume_stack_size = kernel_data.volume_stack_size;
 
+  const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
+  const uint32_t visibility = SHADOW_CATCHER_PATH_VISIBILITY(path_flag, PATH_RAY_ALL_VISIBILITY);
+
 #ifdef __VOLUME_RECORD_ALL__
   Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1];
   uint num_hits = scene_intersect_volume_all(
-      kg, &volume_ray, hits, 2 * volume_stack_size, PATH_RAY_ALL_VISIBILITY);
+      kg, &volume_ray, hits, 2 * volume_stack_size, visibility);
   if (num_hits > 0) {
     Intersection *isect = hits;
 
@@ -60,7 +63,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
   Intersection isect;
   int step = 0;
   while (step < 2 * volume_stack_size &&
-         scene_intersect_volume(kg, &volume_ray, &isect, PATH_RAY_ALL_VISIBILITY)) {
+         scene_intersect_volume(kg, &volume_ray, &isect, visibility)) {
     shader_setup_from_ray(kg, stack_sd, &volume_ray, &isect);
     volume_stack_enter_exit(kg, state, stack_sd);
 
@@ -74,7 +77,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
 #endif
 }
 
-ccl_device void integrator_intersect_volume_stack(KernelGlobals kg, IntegratorState state)
+ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState state)
 {
   PROFILING_INIT(kg, PROFILING_INTERSECT_VOLUME_STACK);
 
@@ -83,16 +86,26 @@ ccl_device void integrator_intersect_volume_stack(KernelGlobals kg, IntegratorSt
 
   Ray volume_ray ccl_optional_struct_init;
   integrator_state_read_ray(kg, state, &volume_ray);
+
+  /* Trace ray in random direction. Any direction works, Z up is a guess to get the
+   * fewest hits. */
+  volume_ray.D = make_float3(0.0f, 0.0f, 1.0f);
   volume_ray.t = FLT_MAX;
 
-  const uint visibility = (INTEGRATOR_STATE(state, path, flag) & PATH_RAY_ALL_VISIBILITY);
   int stack_index = 0, enclosed_index = 0;
 
-  /* Write background shader. */
+  const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
+  const uint32_t visibility = SHADOW_CATCHER_PATH_VISIBILITY(path_flag, PATH_RAY_CAMERA);
+
+  /* Initialize volume stack with background volume For shadow catcher the
+   * background volume is always assumed to be CG. */
   if (kernel_data.background.volume_shader != SHADER_NONE) {
-    const VolumeStack new_entry = {OBJECT_NONE, kernel_data.background.volume_shader};
-    integrator_state_write_volume_stack(state, stack_index, new_entry);
-    stack_index++;
+    if (!(path_flag & PATH_RAY_SHADOW_CATCHER_PASS)) {
+      INTEGRATOR_STATE_ARRAY_WRITE(state, volume_stack, stack_index, object) = OBJECT_NONE;
+      INTEGRATOR_STATE_ARRAY_WRITE(
+          state, volume_stack, stack_index, shader) = kernel_data.background.volume_shader;
+      stack_index++;
+    }
   }
 
   /* Store to avoid global fetches on every intersection step. */
@@ -147,7 +160,7 @@ ccl_device void integrator_intersect_volume_stack(KernelGlobals kg, IntegratorSt
   int enclosed_volumes[MAX_VOLUME_STACK_SIZE];
   int step = 0;
 
-  while (stack_index < volume_stack_size - 1 && enclosed_index < volume_stack_size - 1 &&
+  while (stack_index < volume_stack_size - 1 && enclosed_index < MAX_VOLUME_STACK_SIZE - 1 &&
          step < 2 * volume_stack_size) {
     Intersection isect;
     if (!scene_intersect_volume(kg, &volume_ray, &isect, visibility)) {
@@ -198,9 +211,22 @@ ccl_device void integrator_intersect_volume_stack(KernelGlobals kg, IntegratorSt
   /* Write terminator. */
   const VolumeStack new_entry = {OBJECT_NONE, SHADER_NONE};
   integrator_state_write_volume_stack(state, stack_index, new_entry);
+}
 
-  INTEGRATOR_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK,
-                       DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
+ccl_device void integrator_intersect_volume_stack(KernelGlobals kg, IntegratorState state)
+{
+  integrator_volume_stack_init(kg, state);
+
+  if (INTEGRATOR_STATE(state, path, flag) & PATH_RAY_SHADOW_CATCHER_PASS) {
+    /* Volume stack re-init for shadow catcher, continue with shading of hit. */
+    integrator_intersect_next_kernel_after_shadow_catcher_volume<
+        DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK>(kg, state);
+  }
+  else {
+    /* Volume stack init for camera rays, continue with intersection of camera ray. */
+    INTEGRATOR_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK,
+                         DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
+  }
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/integrator/path_state.h

@@ -67,6 +67,7 @@ ccl_device_inline void path_state_init_integrator(KernelGlobals kg,
   INTEGRATOR_STATE_WRITE(state, path, mis_ray_pdf) = 0.0f;
   INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = 0.0f;
   INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = FLT_MAX;
+  INTEGRATOR_STATE_WRITE(state, path, continuation_probability) = 1.0f;
   INTEGRATOR_STATE_WRITE(state, path, throughput) = make_float3(1.0f, 1.0f, 1.0f);
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_VOLUME) {

intern/cycles/kernel/integrator/shade_background.h

@@ -175,7 +175,7 @@ ccl_device_inline void integrate_distant_lights(KernelGlobals kg,
 
       /* Write to render buffer. */
       const float3 throughput = INTEGRATOR_STATE(state, path, throughput);
-      kernel_accum_emission(kg, state, throughput, light_eval, render_buffer);
+      kernel_accum_emission(kg, state, throughput * light_eval, render_buffer);
     }
   }
 }
@@ -192,23 +192,11 @@ ccl_device void integrator_shade_background(KernelGlobals kg,
 
 #ifdef __SHADOW_CATCHER__
   if (INTEGRATOR_STATE(state, path, flag) & PATH_RAY_SHADOW_CATCHER_BACKGROUND) {
+    /* Special case for shadow catcher where we want to fill the background pass
+     * behind the shadow catcher but also continue tracing the path. */
     INTEGRATOR_STATE_WRITE(state, path, flag) &= ~PATH_RAY_SHADOW_CATCHER_BACKGROUND;
-
-    const int isect_prim = INTEGRATOR_STATE(state, isect, prim);
-    const int isect_type = INTEGRATOR_STATE(state, isect, type);
-    const int shader = intersection_get_shader_from_isect_prim(kg, isect_prim, isect_type);
-    const int shader_flags = kernel_tex_fetch(__shaders, shader).flags;
-
-    if (shader_flags & SD_HAS_RAYTRACE) {
-      INTEGRATOR_PATH_NEXT_SORTED(DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND,
-                                  DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE,
-                                  shader);
-    }
-    else {
-      INTEGRATOR_PATH_NEXT_SORTED(DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND,
-                                  DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE,
-                                  shader);
-    }
+    integrator_intersect_next_kernel_after_shadow_catcher_background<
+        DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND>(kg, state);
     return;
   }
 #endif

intern/cycles/kernel/integrator/shade_light.h

@@ -90,7 +90,7 @@ ccl_device_inline void integrate_light(KernelGlobals kg,
 
   /* Write to render buffer. */
   const float3 throughput = INTEGRATOR_STATE(state, path, throughput);
-  kernel_accum_emission(kg, state, throughput, light_eval, render_buffer);
+  kernel_accum_emission(kg, state, throughput * light_eval, render_buffer);
 }
 
 ccl_device void integrator_shade_light(KernelGlobals kg,

intern/cycles/kernel/integrator/shade_surface.h

@@ -101,7 +101,7 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg,
   }
 
   const float3 throughput = INTEGRATOR_STATE(state, path, throughput);
-  kernel_accum_emission(kg, state, throughput, L, render_buffer);
+  kernel_accum_emission(kg, state, throughput * L, render_buffer);
 }
 #endif /* __EMISSION__ */
 
@@ -479,7 +479,7 @@ ccl_device bool integrate_surface(KernelGlobals kg,
     if (!(path_flag & PATH_RAY_SUBSURFACE)) {
       const float probability = (path_flag & PATH_RAY_TERMINATE_ON_NEXT_SURFACE) ?
                                     0.0f :
-                                    path_state_continuation_probability(kg, state, path_flag);
+                                    INTEGRATOR_STATE(state, path, continuation_probability);
       if (probability == 0.0f) {
         return false;
       }

intern/cycles/kernel/integrator/shade_volume.h

@@ -608,7 +608,7 @@ ccl_device_forceinline void volume_integrate_heterogeneous(
         if (!result.indirect_scatter) {
           const float3 emission = volume_emission_integrate(
               &coeff, closure_flag, transmittance, dt);
-          accum_emission += emission;
+          accum_emission += result.indirect_throughput * emission;
         }
       }
 
@@ -661,7 +661,7 @@ ccl_device_forceinline void volume_integrate_heterogeneous(
 
   /* Write accumulated emission. */
   if (!is_zero(accum_emission)) {
-    kernel_accum_emission(kg, state, result.indirect_throughput, accum_emission, render_buffer);
+    kernel_accum_emission(kg, state, accum_emission, render_buffer);
   }
 
 #  ifdef __DENOISING_FEATURES__
@@ -941,7 +941,7 @@ ccl_device VolumeIntegrateEvent volume_integrate(KernelGlobals kg,
   const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
   const float probability = (path_flag & PATH_RAY_TERMINATE_IN_NEXT_VOLUME) ?
                                 0.0f :
-                                path_state_continuation_probability(kg, state, path_flag);
+                                INTEGRATOR_STATE(state, path, continuation_probability);
   if (probability == 0.0f) {
     return VOLUME_PATH_MISSED;
   }
@@ -1023,25 +1023,9 @@ ccl_device void integrator_shade_volume(KernelGlobals kg,
   }
   else {
     /* Continue to background, light or surface. */
-    if (isect.prim == PRIM_NONE) {
-      INTEGRATOR_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME,
-                           DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND);
-      return;
-    }
-    else if (isect.type & PRIMITIVE_LAMP) {
-      INTEGRATOR_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME,
-                           DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT);
-      return;
-    }
-    else {
-      /* Hit a surface, continue with surface kernel unless terminated. */
-      const int shader = intersection_get_shader(kg, &isect);
-      const int flags = kernel_tex_fetch(__shaders, shader).flags;
-
-      integrator_intersect_shader_next_kernel<DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME>(
-          kg, state, &isect, shader, flags);
-      return;
-    }
+    integrator_intersect_next_kernel_after_volume<DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME>(
+        kg, state, &isect);
+    return;
   }
 #endif /* __VOLUME__ */
 }

intern/cycles/kernel/integrator/shadow_catcher.h

@@ -76,33 +76,6 @@ ccl_device_inline bool kernel_shadow_catcher_path_can_split(KernelGlobals kg,
   return (path_flag & PATH_RAY_TRANSPARENT_BACKGROUND) != 0;
 }
 
-/* NOTE: Leaves kernel scheduling information untouched. Use INIT semantic for one of the paths
- * after this function. */
-ccl_device_inline bool kernel_shadow_catcher_split(KernelGlobals kg,
-                                                   IntegratorState state,
-                                                   const int object_flags)
-{
-#ifdef __SHADOW_CATCHER__
-
-  if (!kernel_shadow_catcher_is_path_split_bounce(kg, state, object_flags)) {
-    return false;
-  }
-
-  /* The split is to be done. Mark the current state as such, so that it stops contributing to the
-   * shadow catcher matte pass, but keeps contributing to the combined pass. */
-  INTEGRATOR_STATE_WRITE(state, path, flag) |= PATH_RAY_SHADOW_CATCHER_HIT;
-
-  /* Split new state from the current one. This new state will only track contribution of shadow
-   * catcher objects ignoring non-catcher objects. */
-  integrator_state_shadow_catcher_split(kg, state);
-
-  return true;
-#else
-  (void)object_flags;
-  return false;
-#endif
-}
-
 #ifdef __SHADOW_CATCHER__
 
 ccl_device_forceinline bool kernel_shadow_catcher_is_matte_path(const uint32_t path_flag)

intern/cycles/kernel/integrator/shadow_state_template.h

@@ -20,7 +20,7 @@ KERNEL_STRUCT_BEGIN(shadow_path)
 /* Index of a pixel within the device render buffer. */
 KERNEL_STRUCT_MEMBER(shadow_path, uint32_t, render_pixel_index, KERNEL_FEATURE_PATH_TRACING)
 /* Current sample number. */
-KERNEL_STRUCT_MEMBER(shadow_path, uint16_t, sample, KERNEL_FEATURE_PATH_TRACING)
+KERNEL_STRUCT_MEMBER(shadow_path, uint32_t, sample, KERNEL_FEATURE_PATH_TRACING)
 /* Random number generator seed. */
 KERNEL_STRUCT_MEMBER(shadow_path, uint32_t, rng_hash, KERNEL_FEATURE_PATH_TRACING)
 /* Random number dimension offset. */

intern/cycles/kernel/integrator/state.h

@@ -173,10 +173,10 @@ typedef const IntegratorShadowStateCPU *ccl_restrict ConstIntegratorShadowState;
 
 /* Array access on GPU with Structure-of-Arrays. */
 
-typedef const int IntegratorState;
-typedef const int ConstIntegratorState;
-typedef const int IntegratorShadowState;
-typedef const int ConstIntegratorShadowState;
+typedef int IntegratorState;
+typedef int ConstIntegratorState;
+typedef int IntegratorShadowState;
+typedef int ConstIntegratorShadowState;
 
 #  define INTEGRATOR_STATE_NULL -1
 

intern/cycles/kernel/integrator/state_template.h

@@ -25,7 +25,7 @@ KERNEL_STRUCT_BEGIN(path)
  * The multiplication is delayed for later, so that state can use 32bit integer. */
 KERNEL_STRUCT_MEMBER(path, uint32_t, render_pixel_index, KERNEL_FEATURE_PATH_TRACING)
 /* Current sample number. */
-KERNEL_STRUCT_MEMBER(path, uint16_t, sample, KERNEL_FEATURE_PATH_TRACING)
+KERNEL_STRUCT_MEMBER(path, uint32_t, sample, KERNEL_FEATURE_PATH_TRACING)
 /* Current ray bounce depth. */
 KERNEL_STRUCT_MEMBER(path, uint16_t, bounce, KERNEL_FEATURE_PATH_TRACING)
 /* Current transparent ray bounce depth. */
@@ -56,6 +56,8 @@ KERNEL_STRUCT_MEMBER(path, float, mis_ray_pdf, KERNEL_FEATURE_PATH_TRACING)
 KERNEL_STRUCT_MEMBER(path, float, mis_ray_t, KERNEL_FEATURE_PATH_TRACING)
 /* Filter glossy. */
 KERNEL_STRUCT_MEMBER(path, float, min_ray_pdf, KERNEL_FEATURE_PATH_TRACING)
+/* Continuation probability for path termination. */
+KERNEL_STRUCT_MEMBER(path, float, continuation_probability, KERNEL_FEATURE_PATH_TRACING)
 /* Throughput. */
 KERNEL_STRUCT_MEMBER(path, float3, throughput, KERNEL_FEATURE_PATH_TRACING)
 /* Ratio of throughput to distinguish diffuse and glossy render passes. */

intern/cycles/kernel/integrator/state_util.h

@@ -326,8 +326,8 @@ ccl_device_inline void integrator_shadow_state_move(KernelGlobals kg,
 
 /* NOTE: Leaves kernel scheduling information untouched. Use INIT semantic for one of the paths
  * after this function. */
-ccl_device_inline void integrator_state_shadow_catcher_split(KernelGlobals kg,
-                                                             IntegratorState state)
+ccl_device_inline IntegratorState integrator_state_shadow_catcher_split(KernelGlobals kg,
+                                                                        IntegratorState state)
 {
 #if defined(__KERNEL_GPU__)
   ConstIntegratorState to_state = atomic_fetch_and_add_uint32(
@@ -337,14 +337,14 @@ ccl_device_inline void integrator_state_shadow_catcher_split(KernelGlobals kg,
 #else
   IntegratorStateCPU *ccl_restrict to_state = state + 1;
 
-  /* Only copy the required subset, since shadow intersections are big and irrelevant here. */
+  /* Only copy the required subset for performance. */
   to_state->path = state->path;
   to_state->ray = state->ray;
   to_state->isect = state->isect;
   integrator_state_copy_volume_stack(kg, to_state, state);
 #endif
 
-  INTEGRATOR_STATE_WRITE(to_state, path, flag) |= PATH_RAY_SHADOW_CATCHER_PASS;
+  return to_state;
 }
 
 #ifdef __KERNEL_CPU__

intern/cycles/kernel/light/light.h

@@ -353,8 +353,8 @@ ccl_device bool light_sample_from_distant_ray(KernelGlobals kg,
   /* compute pdf */
   float invarea = klight->distant.invarea;
   ls->pdf = invarea / (costheta * costheta * costheta);
-  ls->pdf *= kernel_data.integrator.pdf_lights;
   ls->eval_fac = ls->pdf;
+  ls->pdf *= kernel_data.integrator.pdf_lights;
 
   return true;
 }

intern/cycles/kernel/osl/services.cpp

@@ -832,16 +832,21 @@ static bool get_object_attribute(const OSLGlobals::Attribute &attr,
 {
   if (attr.type == TypeDesc::TypePoint || attr.type == TypeDesc::TypeVector ||
       attr.type == TypeDesc::TypeNormal || attr.type == TypeDesc::TypeColor) {
-    return set_attribute_float3(*(float3 *)attr.value.data(), type, derivatives, val);
+    const float *data = (const float *)attr.value.data();
+    return set_attribute_float3(make_float3(data[0], data[1], data[2]), type, derivatives, val);
   }
   else if (attr.type == TypeFloat2) {
-    return set_attribute_float2(*(float2 *)attr.value.data(), type, derivatives, val);
+    const float *data = (const float *)attr.value.data();
+    return set_attribute_float2(make_float2(data[0], data[1]), type, derivatives, val);
   }
   else if (attr.type == TypeDesc::TypeFloat) {
-    return set_attribute_float(*(float *)attr.value.data(), type, derivatives, val);
+    const float *data = (const float *)attr.value.data();
+    return set_attribute_float(data[0], type, derivatives, val);
   }
   else if (attr.type == TypeRGBA || attr.type == TypeDesc::TypeFloat4) {
-    return set_attribute_float4(*(float4 *)attr.value.data(), type, derivatives, val);
+    const float *data = (const float *)attr.value.data();
+    return set_attribute_float4(
+        make_float4(data[0], data[1], data[2], data[3]), type, derivatives, val);
   }
   else if (attr.type == type) {
     size_t datasize = attr.value.datasize();

intern/cycles/kernel/osl/shader.cpp

@@ -132,10 +132,12 @@ static void shaderdata_to_shaderglobals(const KernelGlobalsCPU *kg,
   /* Used by render-services. */
   sd->osl_globals = kg;
   if (path_flag & PATH_RAY_SHADOW) {
+    sd->osl_path_state = nullptr;
     sd->osl_shadow_path_state = (const IntegratorShadowStateCPU *)state;
   }
   else {
     sd->osl_path_state = (const IntegratorStateCPU *)state;
+    sd->osl_shadow_path_state = nullptr;
   }
 }
 

intern/cycles/kernel/osl/shaders/node_musgrave_texture.osl

@@ -697,7 +697,7 @@ shader node_musgrave_texture(
     output float Fac = 0.0)
 {
   float dimension = max(Dimension, 1e-5);
-  float octaves = clamp(Detail, 0.0, 16.0);
+  float octaves = clamp(Detail, 0.0, 15.0);
   float lacunarity = max(Lacunarity, 1e-5);
 
   vector3 s = Vector;

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

@@ -90,7 +90,7 @@ float fractal_noise(float p, float details, float roughness)
   float amp = 1.0;
   float maxamp = 0.0;
   float sum = 0.0;
-  float octaves = clamp(details, 0.0, 16.0);
+  float octaves = clamp(details, 0.0, 15.0);
   int n = (int)octaves;
   for (int i = 0; i <= n; i++) {
     float t = safe_noise(fscale * p);
@@ -119,7 +119,7 @@ float fractal_noise(vector2 p, float details, float roughness)
   float amp = 1.0;
   float maxamp = 0.0;
   float sum = 0.0;
-  float octaves = clamp(details, 0.0, 16.0);
+  float octaves = clamp(details, 0.0, 15.0);
   int n = (int)octaves;
   for (int i = 0; i <= n; i++) {
     float t = safe_noise(fscale * p);
@@ -148,7 +148,7 @@ float fractal_noise(vector3 p, float details, float roughness)
   float amp = 1.0;
   float maxamp = 0.0;
   float sum = 0.0;
-  float octaves = clamp(details, 0.0, 16.0);
+  float octaves = clamp(details, 0.0, 15.0);
   int n = (int)octaves;
   for (int i = 0; i <= n; i++) {
     float t = safe_noise(fscale * p);
@@ -177,7 +177,7 @@ float fractal_noise(vector4 p, float details, float roughness)
   float amp = 1.0;
   float maxamp = 0.0;
   float sum = 0.0;
-  float octaves = clamp(details, 0.0, 16.0);
+  float octaves = clamp(details, 0.0, 15.0);
   int n = (int)octaves;
   for (int i = 0; i <= n; i++) {
     float t = safe_noise(fscale * p);

intern/cycles/kernel/svm/fractal_noise.h

@@ -27,7 +27,7 @@ ccl_device_noinline float fractal_noise_1d(float p, float octaves, float roughne
   float amp = 1.0f;
   float maxamp = 0.0f;
   float sum = 0.0f;
-  octaves = clamp(octaves, 0.0f, 16.0f);
+  octaves = clamp(octaves, 0.0f, 15.0f);
   int n = float_to_int(octaves);
   for (int i = 0; i <= n; i++) {
     float t = noise_1d(fscale * p);
@@ -56,7 +56,7 @@ ccl_device_noinline float fractal_noise_2d(float2 p, float octaves, float roughn
   float amp = 1.0f;
   float maxamp = 0.0f;
   float sum = 0.0f;
-  octaves = clamp(octaves, 0.0f, 16.0f);
+  octaves = clamp(octaves, 0.0f, 15.0f);
   int n = float_to_int(octaves);
   for (int i = 0; i <= n; i++) {
     float t = noise_2d(fscale * p);
@@ -85,7 +85,7 @@ ccl_device_noinline float fractal_noise_3d(float3 p, float octaves, float roughn
   float amp = 1.0f;
   float maxamp = 0.0f;
   float sum = 0.0f;
-  octaves = clamp(octaves, 0.0f, 16.0f);
+  octaves = clamp(octaves, 0.0f, 15.0f);
   int n = float_to_int(octaves);
   for (int i = 0; i <= n; i++) {
     float t = noise_3d(fscale * p);
@@ -114,7 +114,7 @@ ccl_device_noinline float fractal_noise_4d(float4 p, float octaves, float roughn
   float amp = 1.0f;
   float maxamp = 0.0f;
   float sum = 0.0f;
-  octaves = clamp(octaves, 0.0f, 16.0f);
+  octaves = clamp(octaves, 0.0f, 15.0f);
   int n = float_to_int(octaves);
   for (int i = 0; i <= n; i++) {
     float t = noise_4d(fscale * p);

intern/cycles/kernel/svm/musgrave.h

@@ -737,7 +737,7 @@ ccl_device_noinline int svm_node_tex_musgrave(KernelGlobals kg,
   float gain = stack_load_float_default(stack, gain_stack_offset, defaults2.z);
 
   dimension = fmaxf(dimension, 1e-5f);
-  detail = clamp(detail, 0.0f, 16.0f);
+  detail = clamp(detail, 0.0f, 15.0f);
   lacunarity = fmaxf(lacunarity, 1e-5f);
 
   float fac;

intern/cycles/scene/constant_fold.cpp

@@ -43,7 +43,7 @@ bool ConstantFolder::all_inputs_constant() const
 
 void ConstantFolder::make_constant(float value) const
 {
-  VLOG(1) << "Folding " << node->name << "::" << output->name() << " to constant (" << value
+  VLOG(3) << "Folding " << node->name << "::" << output->name() << " to constant (" << value
           << ").";
 
   foreach (ShaderInput *sock, output->links) {
@@ -56,7 +56,7 @@ void ConstantFolder::make_constant(float value) const
 
 void ConstantFolder::make_constant(float3 value) const
 {
-  VLOG(1) << "Folding " << node->name << "::" << output->name() << " to constant " << value << ".";
+  VLOG(3) << "Folding " << node->name << "::" << output->name() << " to constant " << value << ".";
 
   foreach (ShaderInput *sock, output->links) {
     sock->set(value);
@@ -112,7 +112,7 @@ void ConstantFolder::bypass(ShaderOutput *new_output) const
 {
   assert(new_output);
 
-  VLOG(1) << "Folding " << node->name << "::" << output->name() << " to socket "
+  VLOG(3) << "Folding " << node->name << "::" << output->name() << " to socket "
           << new_output->parent->name << "::" << new_output->name() << ".";
 
   /* Remove all outgoing links from socket and connect them to new_output instead.
@@ -131,7 +131,7 @@ void ConstantFolder::discard() const
 {
   assert(output->type() == SocketType::CLOSURE);
 
-  VLOG(1) << "Discarding closure " << node->name << ".";
+  VLOG(3) << "Discarding closure " << node->name << ".";
 
   graph->disconnect(output);
 }

intern/cycles/scene/geometry.cpp

@@ -1588,9 +1588,20 @@ void GeometryManager::device_update_displacement_images(Device *device,
   set<int> bump_images;
   foreach (Geometry *geom, scene->geometry) {
     if (geom->is_modified()) {
+      /* Geometry-level check for hair shadow transparency.
+       * This matches the logic in the `Hair::update_shadow_transparency()`, avoiding access to
+       * possible non-loaded images. */
+      bool need_shadow_transparency = false;
+      if (geom->geometry_type == Geometry::HAIR) {
+        Hair *hair = static_cast<Hair *>(geom);
+        need_shadow_transparency = hair->need_shadow_transparency();
+      }
+
       foreach (Node *node, geom->get_used_shaders()) {
         Shader *shader = static_cast<Shader *>(node);
-        if (!shader->has_displacement || shader->get_displacement_method() == DISPLACE_BUMP) {
+        const bool is_true_displacement = (shader->has_displacement &&
+                                           shader->get_displacement_method() != DISPLACE_BUMP);
+        if (!is_true_displacement && !need_shadow_transparency) {
           continue;
         }
         foreach (ShaderNode *node, shader->graph->nodes) {

intern/cycles/scene/image.cpp

@@ -303,7 +303,6 @@ ImageManager::ImageManager(const DeviceInfo &info)
   animation_frame = 0;
 
   /* Set image limits */
-  features.has_half_float = info.has_half_images;
   features.has_nanovdb = info.has_nanovdb;
 }
 
@@ -357,8 +356,6 @@ void ImageManager::load_image_metadata(Image *img)
 
   metadata.detect_colorspace();
 
-  assert(features.has_half_float ||
-         (metadata.type != IMAGE_DATA_TYPE_HALF4 && metadata.type != IMAGE_DATA_TYPE_HALF));
   assert(features.has_nanovdb || (metadata.type != IMAGE_DATA_TYPE_NANOVDB_FLOAT ||
                                   metadata.type != IMAGE_DATA_TYPE_NANOVDB_FLOAT3));
 

intern/cycles/scene/image.h

@@ -100,7 +100,6 @@ class ImageMetaData {
 /* Information about supported features that Image loaders can use. */
 class ImageDeviceFeatures {
  public:
-  bool has_half_float;
   bool has_nanovdb;
 };
 

intern/cycles/scene/image_oiio.cpp

@@ -30,7 +30,8 @@ OIIOImageLoader::~OIIOImageLoader()
 {
 }
 
-bool OIIOImageLoader::load_metadata(const ImageDeviceFeatures &features, ImageMetaData &metadata)
+bool OIIOImageLoader::load_metadata(const ImageDeviceFeatures & /*features*/,
+                                    ImageMetaData &metadata)
 {
   /* Perform preliminary checks, with meaningful logging. */
   if (!path_exists(filepath.string())) {
@@ -76,7 +77,7 @@ bool OIIOImageLoader::load_metadata(const ImageDeviceFeatures &features, ImageMe
   }
 
   /* check if it's half float */
-  if (spec.format == TypeDesc::HALF && features.has_half_float) {
+  if (spec.format == TypeDesc::HALF) {
     is_half = true;
   }
 

intern/cycles/scene/shader_nodes.cpp

@@ -2397,7 +2397,7 @@ void GlossyBsdfNode::simplify_settings(Scene *scene)
      * Note: Keep the epsilon in sync with kernel!
      */
     if (!roughness_input->link && roughness <= 1e-4f) {
-      VLOG(1) << "Using sharp glossy BSDF.";
+      VLOG(3) << "Using sharp glossy BSDF.";
       distribution = CLOSURE_BSDF_REFLECTION_ID;
     }
   }
@@ -2406,7 +2406,7 @@ void GlossyBsdfNode::simplify_settings(Scene *scene)
      * benefit from closure blur to remove unwanted noise.
      */
     if (roughness_input->link == NULL && distribution == CLOSURE_BSDF_REFLECTION_ID) {
-      VLOG(1) << "Using GGX glossy with filter glossy.";
+      VLOG(3) << "Using GGX glossy with filter glossy.";
       distribution = CLOSURE_BSDF_MICROFACET_GGX_ID;
       roughness = 0.0f;
     }
@@ -2490,7 +2490,7 @@ void GlassBsdfNode::simplify_settings(Scene *scene)
      * Note: Keep the epsilon in sync with kernel!
      */
     if (!roughness_input->link && roughness <= 1e-4f) {
-      VLOG(1) << "Using sharp glass BSDF.";
+      VLOG(3) << "Using sharp glass BSDF.";
       distribution = CLOSURE_BSDF_SHARP_GLASS_ID;
     }
   }
@@ -2499,7 +2499,7 @@ void GlassBsdfNode::simplify_settings(Scene *scene)
      * benefit from closure blur to remove unwanted noise.
      */
     if (roughness_input->link == NULL && distribution == CLOSURE_BSDF_SHARP_GLASS_ID) {
-      VLOG(1) << "Using GGX glass with filter glossy.";
+      VLOG(3) << "Using GGX glass with filter glossy.";
       distribution = CLOSURE_BSDF_MICROFACET_GGX_GLASS_ID;
       roughness = 0.0f;
     }
@@ -2583,7 +2583,7 @@ void RefractionBsdfNode::simplify_settings(Scene *scene)
      * Note: Keep the epsilon in sync with kernel!
      */
     if (!roughness_input->link && roughness <= 1e-4f) {
-      VLOG(1) << "Using sharp refraction BSDF.";
+      VLOG(3) << "Using sharp refraction BSDF.";
       distribution = CLOSURE_BSDF_REFRACTION_ID;
     }
   }
@@ -2592,7 +2592,7 @@ void RefractionBsdfNode::simplify_settings(Scene *scene)
      * benefit from closure blur to remove unwanted noise.
      */
     if (roughness_input->link == NULL && distribution == CLOSURE_BSDF_REFRACTION_ID) {
-      VLOG(1) << "Using GGX refraction with filter glossy.";
+      VLOG(3) << "Using GGX refraction with filter glossy.";
       distribution = CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID;
       roughness = 0.0f;
     }

intern/cycles/scene/svm.cpp

@@ -64,7 +64,7 @@ void SVMShaderManager::device_update_shader(Scene *scene,
   compiler.background = (shader == scene->background->get_shader(scene));
   compiler.compile(shader, *svm_nodes, 0, &summary);
 
-  VLOG(2) << "Compilation summary:\n"
+  VLOG(3) << "Compilation summary:\n"
           << "Shader name: " << shader->name << "\n"
           << summary.full_report();
 }

intern/cycles/session/merge.cpp

@@ -163,8 +163,26 @@ static bool parse_channels(const ImageSpec &in_spec,
     file_layers[layername].passes.push_back(pass);
   }
 
-  /* Loop over all detected render-layers, check whether they contain a full set of input channels.
-   * Any channels that won't be processed internally are also passed through. */
+  /* If file contains a single unnamed layer, name it after the first layer metadata we find. */
+  if (file_layers.size() == 1 && file_layers.find("") != file_layers.end()) {
+    for (const ParamValue &attrib : in_spec.extra_attribs) {
+      const string attrib_name = attrib.name().string();
+      if (string_startswith(attrib_name, "cycles.") && string_endswith(attrib_name, ".samples")) {
+        /* Extract layer name. */
+        const size_t start = strlen("cycles.");
+        const size_t end = attrib_name.size() - strlen(".samples");
+        const string layername = attrib_name.substr(start, end - start);
+
+        /* Reinsert as named instead of unnamed layer. */
+        const MergeImageLayer layer = file_layers[""];
+        file_layers.clear();
+        file_layers[layername] = layer;
+      }
+    }
+  }
+
+  /* Loop over all detected render-layers, check whether they contain a full set of input
+   * channels. Any channels that won't be processed internally are also passed through. */
   for (auto &i : file_layers) {
     const string &name = i.first;
     MergeImageLayer &layer = i.second;

intern/cycles/test/render_graph_finalize_test.cpp

@@ -179,7 +179,7 @@ class RenderGraph : public testing::Test {
   virtual void SetUp()
   {
     util_logging_start();
-    util_logging_verbosity_set(1);
+    util_logging_verbosity_set(3);
 
     device_cpu = Device::create(device_info, stats, profiler);
     scene = new Scene(scene_params, device_cpu);

intern/cycles/util/debug.cpp

@@ -99,26 +99,4 @@ void DebugFlags::reset()
   optix.reset();
 }
 
-std::ostream &operator<<(std::ostream &os, DebugFlagsConstRef debug_flags)
-{
-  os << "CPU flags:\n"
-     << "  AVX2       : " << string_from_bool(debug_flags.cpu.avx2) << "\n"
-     << "  AVX        : " << string_from_bool(debug_flags.cpu.avx) << "\n"
-     << "  SSE4.1     : " << string_from_bool(debug_flags.cpu.sse41) << "\n"
-     << "  SSE3       : " << string_from_bool(debug_flags.cpu.sse3) << "\n"
-     << "  SSE2       : " << string_from_bool(debug_flags.cpu.sse2) << "\n"
-     << "  BVH layout : " << bvh_layout_name(debug_flags.cpu.bvh_layout) << "\n";
-
-  os << "CUDA flags:\n"
-     << "  Adaptive Compile : " << string_from_bool(debug_flags.cuda.adaptive_compile) << "\n";
-
-  os << "OptiX flags:\n"
-     << "  Debug : " << string_from_bool(debug_flags.optix.use_debug) << "\n";
-
-  os << "HIP flags:\n"
-     << "  HIP streams : " << string_from_bool(debug_flags.hip.adaptive_compile) << "\n";
-
-  return os;
-}
-
 CCL_NAMESPACE_END

intern/cycles/util/debug.h

@@ -160,8 +160,6 @@ inline DebugFlags &DebugFlags()
   return DebugFlags::get();
 }
 
-std::ostream &operator<<(std::ostream &os, DebugFlagsConstRef debug_flags);
-
 CCL_NAMESPACE_END
 
 #endif /* __UTIL_DEBUG_H__ */

intern/ghost/CMakeLists.txt

@@ -489,6 +489,9 @@ if(WITH_XR_OPENXR)
     intern/GHOST_XrSwapchain.h
     intern/GHOST_Xr_intern.h
     intern/GHOST_Xr_openxr_includes.h
+
+    # Header only library.
+    ../../extern/tinygltf/tiny_gltf.h
   )
   list(APPEND INC
     ../../extern/json/include

intern/locale/boost_locale_wrapper.cpp

@@ -26,8 +26,8 @@ static std::string messages_path;
 static std::string default_domain;
 static std::string locale_str;
 
-/* Note: We cannot use short stuff like boost::locale::gettext, because those return
- * std::basic_string objects, which c_ptr()-returned char* is no more valid
+/* NOTE: We cannot use short stuff like `boost::locale::gettext`, because those return
+ * `std::basic_string` objects, which c_ptr()-returned char* is no more valid
  * once deleted (which happens as soons they are out of scope of this func). */
 typedef boost::locale::message_format<char> char_message_facet;
 static std::locale locale_global;
@@ -63,7 +63,7 @@ static void bl_locale_global_cache()
 
 void bl_locale_init(const char *_messages_path, const char *_default_domain)
 {
-  // Avoid using ICU backend, we do not need its power and it's rather heavy!
+  /* Avoid using ICU backend, we do not need its power and it's rather heavy! */
   boost::locale::localization_backend_manager lman =
       boost::locale::localization_backend_manager::global();
 #if defined(_WIN32)
@@ -81,7 +81,7 @@ void bl_locale_set(const char *locale)
 {
   boost::locale::generator gen;
   std::locale _locale;
-  // Specify location of dictionaries.
+  /* Specify location of dictionaries. */
   gen.add_messages_path(messages_path);
   gen.add_messages_domain(default_domain);
   // gen.set_default_messages_domain(default_domain);
@@ -99,12 +99,12 @@ void bl_locale_set(const char *locale)
 #endif
     }
     std::locale::global(_locale);
-    // Note: boost always uses "C" LC_NUMERIC by default!
+    /* NOTE: boost always uses "C" LC_NUMERIC by default! */
 
     bl_locale_global_cache();
 
-    // Generate the locale string
-    // (useful to know which locale we are actually using in case of "default" one).
+    /* Generate the locale string
+     * (useful to know which locale we are actually using in case of "default" one). */
 #define LOCALE_INFO std::use_facet<boost::locale::info>(_locale)
 
     locale_str = LOCALE_INFO.language();
@@ -117,10 +117,9 @@ void bl_locale_set(const char *locale)
 
 #undef LOCALE_INFO
   }
-  // Extra catch on `std::runtime_error` is needed for macOS/Clang as it seems that exceptions
-  // like `boost::locale::conv::conversion_error` (which inherit from `std::runtime_error`) are
-  // not caught by their ancestor `std::exception`. See
-  // https://developer.blender.org/T88877#1177108 .
+  /* Extra catch on `std::runtime_error` is needed for macOS/Clang as it seems that exceptions
+   * like `boost::locale::conv::conversion_error` (which inherit from `std::runtime_error`) are
+   * not caught by their ancestor `std::exception`. See T88877#1177108 */
   catch (std::runtime_error const &e) {
     std::cout << "bl_locale_set(" << locale << "): " << e.what() << " \n";
   }

release/datafiles/userdef/userdef_default_theme.c

@@ -25,8 +25,8 @@ const bTheme U_theme_default = {
       .outline = RGBA(0x3d3d3dff),
       .inner = RGBA(0x545454ff),
       .inner_sel = RGBA(0x4772b3ff),
-      .item = RGBA(0xffffff80),
-      .text = RGBA(0xd9d9d9ff),
+      .item = RGBA(0x1d1d1d80),
+      .text = RGBA(0xe6e6e6ff),
       .text_sel = RGBA(0xffffffff),
       .roundness = 0.2f,
     },
@@ -43,7 +43,7 @@ const bTheme U_theme_default = {
       .outline = RGBA(0x3d3d3dff),
       .inner = RGBA(0x282828ff),
       .inner_sel = RGBA(0x4772b3ff),
-      .item = RGBA(0xffffff80),
+      .item = RGBA(0xffffffb3),
       .text = RGBA(0xe6e6e6ff),
       .text_sel = RGBA(0xffffffff),
       .roundness = 0.2f,
@@ -59,7 +59,7 @@ const bTheme U_theme_default = {
     },
     .wcol_radio = {
       .outline = RGBA(0x3d3d3dff),
-      .inner = RGBA(0x282828ff),
+      .inner = RGBA(0x545454ff),
       .inner_sel = RGBA(0x4772b3ff),
       .item = RGBA(0x252525ff),
       .text = RGBA(0xe6e6e6ff),
@@ -68,16 +68,16 @@ const bTheme U_theme_default = {
     },
     .wcol_option = {
       .outline = RGBA(0x3d3d3dff),
-      .inner = RGBA(0x282828ff),
-      .inner_sel = RGBA(0x71aaffff),
-      .item = RGBA(0x111111ff),
+      .inner = RGBA(0x545454ff),
+      .inner_sel = RGBA(0x4772b3ff),
+      .item = RGBA(0xffffffff),
       .text = RGBA(0xe6e6e6ff),
       .text_sel = RGBA(0xffffffff),
       .roundness = 0.2f,
     },
     .wcol_toggle = {
       .outline = RGBA(0x3d3d3dff),
-      .inner = RGBA(0x282828ff),
+      .inner = RGBA(0x545454ff),
       .inner_sel = RGBA(0x4772b3ff),
       .item = RGBA(0x252525ff),
       .text = RGBA(0xe6e6e6ff),
@@ -148,11 +148,11 @@ const bTheme U_theme_default = {
       .roundness = 0.2f,
     },
     .wcol_tooltip = {
-      .outline = RGBA(0x19191aff),
-      .inner = RGBA(0x181818ff),
-      .inner_sel = RGBA(0x181818ff),
-      .item = RGBA(0x181818ff),
-      .text = RGBA(0xccccccff),
+      .outline = RGBA(0x242424ff),
+      .inner = RGBA(0x1d1d1dff),
+      .inner_sel = RGBA(0x4772b3ff),
+      .item = RGBA(0xd9d9d9ff),
+      .text = RGBA(0xd9d9d9ff),
       .text_sel = RGBA(0xffffffff),
       .roundness = 0.2f,
     },
@@ -187,16 +187,16 @@ const bTheme U_theme_default = {
       .outline = RGBA(0x2d2d2dff),
       .inner = RGBA(0x2d2d2d00),
       .inner_sel = RGBA(0x484a4fff),
-      .item = RGBA(0xb3b3b3ff),
+      .item = RGBA(0x4772b3ff),
       .text = RGBA(0xccccccff),
       .text_sel = RGBA(0xffffffff),
       .roundness = 0.2f,
     },
     .wcol_pie_menu = {
-      .outline = RGBA(0x3d3d3dff),
+      .outline = RGBA(0x242424ff),
       .inner = RGBA(0x181818ff),
-      .inner_sel = RGBA(0x5680c2ff),
-      .item = RGBA(0x4772b3ff),
+      .inner_sel = RGBA(0x4772b3ff),
+      .item = RGBA(0x545454ff),
       .text = RGBA(0xd9d9d9ff),
       .text_sel = RGBA(0xffffffff),
       .roundness = 0.2f,
@@ -414,7 +414,7 @@ const bTheme U_theme_default = {
     .list = RGBA(0x1d1d1dff),
     .list_title = RGBA(0xffffffff),
     .list_text = RGBA(0xb8b8b8ff),
-    .list_text_hi = RGBA(0xffffffff),
+    .list_text_hi = RGBA(0xffaf23ff),
     .navigation_bar = RGBA(0x1d1d1dff),
     .execution_buts = RGBA(0x303030ff),
     .panelcolors = {
@@ -497,7 +497,7 @@ const bTheme U_theme_default = {
     .back = RGBA(0x30303000),
     .title = RGBA(0xeeeeeeff),
     .text = RGBA(0xa6a6a6ff),
-    .text_hi = RGBA(0x143e66ff),
+    .text_hi = RGBA(0xffffffff),
     .header = RGBA(0x303030b3),
     .header_text = RGBA(0xeeeeeeff),
     .header_text_hi = RGBA(0xffffffff),
@@ -512,7 +512,7 @@ const bTheme U_theme_default = {
     .list = RGBA(0x1d1d1dff),
     .list_title = RGBA(0xffffffff),
     .list_text = RGBA(0xb8b8b8ff),
-    .list_text_hi = RGBA(0xffffffff),
+    .list_text_hi = RGBA(0xffaf23ff),
     .navigation_bar = RGBA(0x1d1d1dff),
     .execution_buts = RGBA(0x303030ff),
     .panelcolors = {
@@ -573,8 +573,8 @@ const bTheme U_theme_default = {
     .button_text_hi = RGBA(0xffffffff),
     .list = RGBA(0x1d1d1dff),
     .list_title = RGBA(0xffffffff),
-    .list_text = RGBA(0xe5e5e5ff),
-    .list_text_hi = RGBA(0xffffffff),
+    .list_text = RGBA(0xb8b8b8ff),
+    .list_text_hi = RGBA(0xffaf23ff),
     .navigation_bar = RGBA(0x1d1d1dff),
     .execution_buts = RGBA(0x303030ff),
     .panelcolors = {
@@ -612,7 +612,7 @@ const bTheme U_theme_default = {
     .nla_sound_sel = RGBA(0x1f7a7aff),
   },
   .space_sequencer = {
-    .back = RGBA(0x30303000),
+    .back = RGBA(0x18181800),
     .title = RGBA(0xeeeeeeff),
     .text = RGBA(0xa6a6a6ff),
     .text_hi = RGBA(0xffffffff),
@@ -635,7 +635,7 @@ const bTheme U_theme_default = {
       .sub_back = RGBA(0x0000001f),
     },
     .shade1 = RGBA(0xa0a0a000),
-    .grid = RGBA(0x181818ff),
+    .grid = RGBA(0x303030ff),
     .vertex_select = RGBA(0xff8500ff),
     .bone_pose = RGBA(0x50c8ff50),
     .cframe = RGBA(0x4772b3ff),
@@ -811,8 +811,8 @@ const bTheme U_theme_default = {
     .button_text_hi = RGBA(0xffffffff),
     .list = RGBA(0x303030ff),
     .list_title = RGBA(0xffffffff),
-    .list_text = RGBA(0xccccccff),
-    .list_text_hi = RGBA(0xffffffff),
+    .list_text = RGBA(0xb8b8b8ff),
+    .list_text_hi = RGBA(0xffaf23ff),
     .navigation_bar = RGBA(0x1d1d1dff),
     .execution_buts = RGBA(0x303030ff),
     .panelcolors = {
@@ -935,7 +935,7 @@ const bTheme U_theme_default = {
     .list = RGBA(0x303030ff),
     .list_title = RGBA(0xffffff00),
     .list_text = RGBA(0xb8b8b8ff),
-    .list_text_hi = RGBA(0xffffffff),
+    .list_text_hi = RGBA(0xffaf23ff),
     .navigation_bar = RGBA(0x1d1d1dff),
     .execution_buts = RGBA(0x303030ff),
     .panelcolors = {
@@ -1046,8 +1046,8 @@ const bTheme U_theme_default = {
     .button_text_hi = RGBA(0xffffffff),
     .list = RGBA(0x303030ff),
     .list_title = RGBA(0xc3c3c3ff),
-    .list_text = RGBA(0xc3c3c3ff),
-    .list_text_hi = RGBA(0xffffffff),
+    .list_text = RGBA(0xb8b8b8ff),
+    .list_text_hi = RGBA(0xffaf23ff),
     .navigation_bar = RGBA(0x1d1d1dff),
     .execution_buts = RGBA(0x303030ff),
     .panelcolors = {

release/scripts/modules/bl_i18n_utils/bl_extract_messages.py

@@ -252,27 +252,50 @@ def dump_rna_messages(msgs, reports, settings, verbose=False):
 
     # Function definitions
     def walk_properties(cls):
+        # This handles properties whose name is the same as their identifier.
+        # Usually, it means that those are internal properties not exposed in the UI, however there are some cases
+        # where the UI label is actually defined and same as the identifier (color spaces e.g., `RGB` etc.).
+        # So we only exclude those properties in case they belong to an operator for now.
+        def prop_name_validate(cls, prop_name, prop_identifier):
+            if prop_name != prop_identifier:
+                return True
+            # Heuristic: A lot of operator's HIDDEN properties have no UI label/description.
+            # While this is not ideal (for API doc purposes, description should always be provided),
+            # for now skip those properties.
+            # NOTE: keep in sync with C code in ui_searchbox_region_draw_cb__operator().
+            if issubclass(cls, bpy.types.OperatorProperties) and "_OT_" in cls.__name__:
+                return False
+            # Heuristic: If UI label is not capitalized, it is likely a private (undocumented) property,
+            # that can be skipped.
+            if prop_name and not prop_name[0].isupper():
+                return False
+            return True
+
         bl_rna = cls.bl_rna
         # Get our parents' properties, to not export them multiple times.
         bl_rna_base = bl_rna.base
+        bl_rna_base_props = set()
         if bl_rna_base:
-            bl_rna_base_props = set(bl_rna_base.properties.values())
-        else:
-            bl_rna_base_props = set()
+            bl_rna_base_props |= set(bl_rna_base.properties.values())
+        for cls_base in cls.__bases__:
+            bl_rna_base = getattr(cls_base, "bl_rna", None)
+            if not bl_rna_base:
+                continue
+            bl_rna_base_props |= set(bl_rna_base.properties.values())
 
         props = sorted(bl_rna.properties, key=lambda p: p.identifier)
         for prop in props:
             # Only write this property if our parent hasn't got it.
             if prop in bl_rna_base_props:
                 continue
-            if prop.identifier == "rna_type":
+            if prop.identifier in {"rna_type", "bl_icon", "icon"}:
                 continue
             reports["rna_props"].append((cls, prop))
 
             msgsrc = "bpy.types.{}.{}".format(bl_rna.identifier, prop.identifier)
             msgctxt = prop.translation_context or default_context
 
-            if prop.name and (prop.name != prop.identifier or msgctxt != default_context):
+            if prop.name and prop_name_validate(cls, prop.name, prop.identifier):
                 process_msg(msgs, msgctxt, prop.name, msgsrc, reports, check_ctxt_rna, settings)
             if prop.description:
                 process_msg(msgs, default_context, prop.description, msgsrc, reports, check_ctxt_rna_tip, settings)
@@ -282,7 +305,7 @@ def dump_rna_messages(msgs, reports, settings, verbose=False):
                 for item in prop.enum_items:
                     msgsrc = "bpy.types.{}.{}:'{}'".format(bl_rna.identifier, prop.identifier, item.identifier)
                     done_items.add(item.identifier)
-                    if item.name and item.name != item.identifier:
+                    if item.name and prop_name_validate(cls, item.name, item.identifier):
                         process_msg(msgs, msgctxt, item.name, msgsrc, reports, check_ctxt_rna, settings)
                     if item.description:
                         process_msg(msgs, default_context, item.description, msgsrc, reports, check_ctxt_rna_tip,
@@ -292,7 +315,7 @@ def dump_rna_messages(msgs, reports, settings, verbose=False):
                         continue
                     msgsrc = "bpy.types.{}.{}:'{}'".format(bl_rna.identifier, prop.identifier, item.identifier)
                     done_items.add(item.identifier)
-                    if item.name and item.name != item.identifier:
+                    if item.name and prop_name_validate(cls, item.name, item.identifier):
                         process_msg(msgs, msgctxt, item.name, msgsrc, reports, check_ctxt_rna, settings)
                     if item.description:
                         process_msg(msgs, default_context, item.description, msgsrc, reports, check_ctxt_rna_tip,

release/scripts/modules/bl_i18n_utils/settings.py

@@ -195,7 +195,7 @@ DOMAIN = "blender"
 
 # Our own "gettext" stuff.
 # File type (ext) to parse.
-PYGETTEXT_ALLOWED_EXTS = {".c", ".cpp", ".cxx", ".hpp", ".hxx", ".h"}
+PYGETTEXT_ALLOWED_EXTS = {".c", ".cc", ".cpp", ".cxx", ".hh", ".hpp", ".hxx", ".h"}
 
 # Max number of contexts into a BLT_I18N_MSGID_MULTI_CTXT macro...
 PYGETTEXT_MAX_MULTI_CTXT = 16
@@ -400,6 +400,8 @@ WARN_MSGID_NOT_CAPITALIZED_ALLOWED = {
     "verts only",
     "view",
     "virtual parents",
+    "and NVIDIA driver version 470 or newer",
+    "and AMD driver version ??? or newer",
 }
 WARN_MSGID_NOT_CAPITALIZED_ALLOWED |= set(lng[2] for lng in LANGUAGES)
 

release/scripts/modules/bl_i18n_utils/utils_spell_check.py

@@ -66,6 +66,7 @@ class SpellChecker:
         "ons",  # add-ons
         "pong",  # ping pong
         "resumable",
+        "runtimes",
         "scalable",
         "shadeless",
         "shouldn",  # shouldn't
@@ -76,6 +77,12 @@ class SpellChecker:
         "vertices",
         "wasn",  # wasn't
 
+        # Brands etc.
+        "htc",
+        "huawei",
+        "vive",
+        "xbox",
+
         # Merged words
         "antialiasing", "antialias",
         "arcsine", "arccosine", "arctangent",
@@ -131,6 +138,7 @@ class SpellChecker:
         "forcefield", "forcefields",
         "fulldome", "fulldomes",
         "fullscreen",
+        "gamepad",
         "gridline", "gridlines",
         "hardlight",
         "hemi",
@@ -681,11 +689,13 @@ class SpellChecker:
         "ctrl",
         "cw", "ccw",
         "dev",
+        "dls",
         "djv",
         "dpi",
         "dvar",
         "dx",
         "eo",
+        "ewa",
         "fh",
         "fk",
         "fov",
@@ -725,6 +735,7 @@ class SpellChecker:
         "rhs",
         "rv",
         "sdl",
+        "sdls",
         "sl",
         "smpte",
         "ssao",
@@ -759,6 +770,7 @@ class SpellChecker:
         "svbvh",
 
         # Files types/formats
+        "aac",
         "avi",
         "attrac",
         "autocad",
@@ -781,6 +793,7 @@ class SpellChecker:
         "ico",
         "jpg", "jpeg", "jpegs",
         "json",
+        "lzw",
         "matroska",
         "mdd",
         "mkv",
@@ -790,6 +803,7 @@ class SpellChecker:
         "openjpeg",
         "osl",
         "oso",
+        "pcm",
         "piz",
         "png", "pngs",
         "po",

release/scripts/modules/bpy/utils/__init__.py

@@ -26,6 +26,7 @@ not associated with blenders internal data.
 __all__ = (
     "blend_paths",
     "escape_identifier",
+    "flip_name",
     "unescape_identifier",
     "keyconfig_init",
     "keyconfig_set",
@@ -61,6 +62,7 @@ from _bpy import (
     _utils_units as units,
     blend_paths,
     escape_identifier,
+    flip_name,
     unescape_identifier,
     register_class,
     resource_path,

release/scripts/modules/rna_manual_reference.py

@@ -60,10 +60,12 @@ url_manual_mapping = (
 	("bpy.types.fluiddomainsettings.sndparticle_sampling_wavecrest*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-sampling-wavecrest"),
 	("bpy.types.rigidbodyconstraint.use_override_solver_iterations*", "physics/rigid_body/constraints/introduction.html#bpy-types-rigidbodyconstraint-use-override-solver-iterations"),
 	("bpy.types.toolsettings.use_transform_correct_face_attributes*", "modeling/meshes/tools/tool_settings.html#bpy-types-toolsettings-use-transform-correct-face-attributes"),
+	("bpy.types.cyclesrendersettings.adaptive_scrambling_distance*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-adaptive-scrambling-distance"),
 	("bpy.types.cyclesrendersettings.preview_adaptive_min_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-preview-adaptive-min-samples"),
 	("bpy.types.rendersettings.use_sequencer_override_scene_strip*", "video_editing/preview/sidebar.html#bpy-types-rendersettings-use-sequencer-override-scene-strip"),
 	("bpy.types.toolsettings.use_transform_correct_keep_connected*", "modeling/meshes/tools/tool_settings.html#bpy-types-toolsettings-use-transform-correct-keep-connected"),
 	("bpy.types.cyclesrendersettings.preview_denoising_prefilter*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-preview-denoising-prefilter"),
+	("bpy.types.cyclesrendersettings.preview_scrambling_distance*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-preview-scrambling-distance"),
 	("bpy.types.fluiddomainsettings.sndparticle_potential_radius*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-potential-radius"),
 	("bpy.types.cyclesrendersettings.film_transparent_roughness*", "render/cycles/render_settings/film.html#bpy-types-cyclesrendersettings-film-transparent-roughness"),
 	("bpy.types.cyclesrendersettings.preview_adaptive_threshold*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-preview-adaptive-threshold"),
@@ -156,9 +158,11 @@ url_manual_mapping = (
 	("bpy.types.brushgpencilsettings.fill_simplify_level*", "grease_pencil/modes/draw/tools/fill.html#bpy-types-brushgpencilsettings-fill-simplify-level"),
 	("bpy.types.brushgpencilsettings.use_jitter_pressure*", "grease_pencil/modes/draw/tools/draw.html#bpy-types-brushgpencilsettings-use-jitter-pressure"),
 	("bpy.types.brushgpencilsettings.use_settings_random*", "grease_pencil/modes/draw/tools/draw.html#bpy-types-brushgpencilsettings-use-settings-random"),
+	("bpy.types.colormanagedinputcolorspacesettings.name*", "editors/image/image_settings.html#bpy-types-colormanagedinputcolorspacesettings-name"),
 	("bpy.types.cyclesrendersettings.denoising_prefilter*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-denoising-prefilter"),
 	("bpy.types.cyclesrendersettings.preview_dicing_rate*", "render/cycles/render_settings/subdivision.html#bpy-types-cyclesrendersettings-preview-dicing-rate"),
 	("bpy.types.cyclesrendersettings.sample_clamp_direct*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-sample-clamp-direct"),
+	("bpy.types.cyclesrendersettings.scrambling_distance*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-scrambling-distance"),
 	("bpy.types.cyclesworldsettings.volume_interpolation*", "render/cycles/world_settings.html#bpy-types-cyclesworldsettings-volume-interpolation"),
 	("bpy.types.fluiddomainsettings.mesh_particle_radius*", "physics/fluid/type/domain/liquid/mesh.html#bpy-types-fluiddomainsettings-mesh-particle-radius"),
 	("bpy.types.fluiddomainsettings.sndparticle_boundary*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-boundary"),
@@ -173,6 +177,7 @@ url_manual_mapping = (
 	("bpy.types.spacespreadsheet.geometry_component_type*", "editors/spreadsheet.html#bpy-types-spacespreadsheet-geometry-component-type"),
 	("bpy.types.toolsettings.use_gpencil_weight_data_add*", "grease_pencil/modes/draw/introduction.html#bpy-types-toolsettings-use-gpencil-weight-data-add"),
 	("bpy.types.view3doverlay.texture_paint_mode_opacity*", "editors/3dview/display/overlays.html#bpy-types-view3doverlay-texture-paint-mode-opacity"),
+	("bpy.ops.mesh.customdata_custom_splitnormals_clear*", "modeling/meshes/properties/custom_data.html#bpy-ops-mesh-customdata-custom-splitnormals-clear"),
 	("bpy.types.bakesettings.use_pass_ambient_occlusion*", "render/cycles/baking.html#bpy-types-bakesettings-use-pass-ambient-occlusion"),
 	("bpy.types.brush.surface_smooth_shape_preservation*", "sculpt_paint/sculpting/tools/smooth.html#bpy-types-brush-surface-smooth-shape-preservation"),
 	("bpy.types.brush.use_cloth_pin_simulation_boundary*", "sculpt_paint/sculpting/tools/cloth.html#bpy-types-brush-use-cloth-pin-simulation-boundary"),
@@ -255,6 +260,7 @@ url_manual_mapping = (
 	("bpy.types.toolsettings.use_snap_backface_culling*", "editors/3dview/controls/snapping.html#bpy-types-toolsettings-use-snap-backface-culling"),
 	("bpy.types.toolsettings.use_transform_data_origin*", "scene_layout/object/tools/tool_settings.html#bpy-types-toolsettings-use-transform-data-origin"),
 	("bpy.types.view3doverlay.sculpt_mode_mask_opacity*", "sculpt_paint/sculpting/editing/mask.html#bpy-types-view3doverlay-sculpt-mode-mask-opacity"),
+	("bpy.ops.mesh.customdata_custom_splitnormals_add*", "modeling/meshes/properties/custom_data.html#bpy-ops-mesh-customdata-custom-splitnormals-add"),
 	("bpy.ops.outliner.collection_indirect_only_clear*", "render/layers/introduction.html#bpy-ops-outliner-collection-indirect-only-clear"),
 	("bpy.types.cyclesrendersettings.max_subdivisions*", "render/cycles/render_settings/subdivision.html#bpy-types-cyclesrendersettings-max-subdivisions"),
 	("bpy.types.cyclesrendersettings.preview_denoiser*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-preview-denoiser"),
@@ -278,6 +284,7 @@ url_manual_mapping = (
 	("bpy.types.freestylelineset.select_by_collection*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset-select-by-collection"),
 	("bpy.types.freestylelineset.select_by_edge_types*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset-select-by-edge-types"),
 	("bpy.types.freestylelineset.select_by_face_marks*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset-select-by-face-marks"),
+	("bpy.types.geometrynodeinputcurvehandlepositions*", "modeling/geometry_nodes/curve/curve_handle_position.html#bpy-types-geometrynodeinputcurvehandlepositions"),
 	("bpy.types.linestyle*modifier_distancefromcamera*", "render/freestyle/view_layer/line_style/modifiers/color/distance_from_camera.html#bpy-types-linestyle-modifier-distancefromcamera"),
 	("bpy.types.linestyle*modifier_distancefromobject*", "render/freestyle/view_layer/line_style/modifiers/color/distance_from_object.html#bpy-types-linestyle-modifier-distancefromobject"),
 	("bpy.types.linestylegeometrymodifier_2dtransform*", "render/freestyle/view_layer/line_style/modifiers/geometry/2d_transform.html#bpy-types-linestylegeometrymodifier-2dtransform"),
@@ -315,6 +322,7 @@ url_manual_mapping = (
 	("bpy.types.freestylelinestyle.material_boundary*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-material-boundary"),
 	("bpy.types.freestylelinestyle.use_split_pattern*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-use-split-pattern"),
 	("bpy.types.freestylesettings.use_view_map_cache*", "render/freestyle/view_layer/freestyle.html#bpy-types-freestylesettings-use-view-map-cache"),
+	("bpy.types.geometrynodecurvehandletypeselection*", "modeling/geometry_nodes/curve/handle_type_selection.html#bpy-types-geometrynodecurvehandletypeselection"),
 	("bpy.types.greasepencil.curve_edit_corner_angle*", "grease_pencil/modes/edit/curve_editing.html#bpy-types-greasepencil-curve-edit-corner-angle"),
 	("bpy.types.linestylegeometrymodifier_tipremover*", "render/freestyle/view_layer/line_style/modifiers/geometry/tip_remover.html#bpy-types-linestylegeometrymodifier-tipremover"),
 	("bpy.types.movieclipuser.use_render_undistorted*", "editors/clip/display/clip_display.html#bpy-types-movieclipuser-use-render-undistorted"),
@@ -360,6 +368,8 @@ url_manual_mapping = (
 	("bpy.types.freestylelineset.face_mark_negation*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset-face-mark-negation"),
 	("bpy.types.freestylelinestyle.integration_type*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-integration-type"),
 	("bpy.types.freestylelinestyle.use_split_length*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-use-split-length"),
+	("bpy.types.geometrynodedistributepointsonfaces*", "modeling/geometry_nodes/point/distribute_points_on_faces.html#bpy-types-geometrynodedistributepointsonfaces"),
+	("bpy.types.geometrynodesetcurvehandlepositions*", "modeling/geometry_nodes/curve/set_handle_positions.html#bpy-types-geometrynodesetcurvehandlepositions"),
 	("bpy.types.linestylegeometrymodifier_blueprint*", "render/freestyle/view_layer/line_style/modifiers/geometry/blueprint.html#bpy-types-linestylegeometrymodifier-blueprint"),
 	("bpy.types.materialgpencilstyle.alignment_mode*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-alignment-mode"),
 	("bpy.types.particlesettings.use_modifier_stack*", "physics/particles/emitter/emission.html#bpy-types-particlesettings-use-modifier-stack"),
@@ -374,6 +384,7 @@ url_manual_mapping = (
 	("bpy.types.spacespreadsheetrowfilter.threshold*", "editors/spreadsheet.html#bpy-types-spacespreadsheetrowfilter-threshold"),
 	("bpy.types.toolsettings.use_snap_grid_absolute*", "editors/3dview/controls/snapping.html#bpy-types-toolsettings-use-snap-grid-absolute"),
 	("bpy.types.view3doverlay.show_face_orientation*", "editors/3dview/display/overlays.html#bpy-types-view3doverlay-show-face-orientation"),
+	("bpy.ops.gpencil.bake_grease_pencil_animation*", "grease_pencil/animation/tools.html#bpy-ops-gpencil-bake-grease-pencil-animation"),
 	("bpy.ops.object.blenderkit_material_thumbnail*", "addons/3d_view/blenderkit.html#bpy-ops-object-blenderkit-material-thumbnail"),
 	("bpy.ops.object.multires_higher_levels_delete*", "modeling/modifiers/generate/multiresolution.html#bpy-ops-object-multires-higher-levels-delete"),
 	("bpy.ops.object.vertex_group_copy_to_selected*", "modeling/meshes/properties/vertex_groups/vertex_groups.html#bpy-ops-object-vertex-group-copy-to-selected"),
@@ -408,7 +419,7 @@ url_manual_mapping = (
 	("bpy.types.freestylelinestyle.use_chain_count*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-use-chain-count"),
 	("bpy.types.freestylelinestyle.use_dashed_line*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-use-dashed-line"),
 	("bpy.types.freestylelinestyle.use_same_object*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-use-same-object"),
-	("bpy.types.geometrynodeattributesampletexture*", "modeling/geometry_nodes/attribute/attribute_sample_texture.html#bpy-types-geometrynodeattributesampletexture"),
+	("bpy.types.functionnodeinputspecialcharacters*", "modeling/geometry_nodes/text/special_characters.html#bpy-types-functionnodeinputspecialcharacters"),
 	("bpy.types.gpencilsculptguide.reference_point*", "grease_pencil/modes/draw/guides.html#bpy-types-gpencilsculptguide-reference-point"),
 	("bpy.types.greasepencil.edit_curve_resolution*", "grease_pencil/modes/edit/curve_editing.html#bpy-types-greasepencil-edit-curve-resolution"),
 	("bpy.types.linestylegeometrymodifier_2doffset*", "render/freestyle/view_layer/line_style/modifiers/geometry/2d_offset.html#bpy-types-linestylegeometrymodifier-2doffset"),
@@ -417,6 +428,7 @@ url_manual_mapping = (
 	("bpy.types.rendersettings.line_thickness_mode*", "render/freestyle/render.html#bpy-types-rendersettings-line-thickness-mode"),
 	("bpy.types.rendersettings.motion_blur_shutter*", "render/cycles/render_settings/motion_blur.html#bpy-types-rendersettings-motion-blur-shutter"),
 	("bpy.types.rendersettings.use_persistent_data*", "render/cycles/render_settings/performance.html#bpy-types-rendersettings-use-persistent-data"),
+	("bpy.types.sequencertimelineoverlay.show_grid*", "editors/video_sequencer/sequencer/display.html#bpy-types-sequencertimelineoverlay-show-grid"),
 	("bpy.types.sequencertoolsettings.overlap_mode*", "video_editing/sequencer/editing.html#bpy-types-sequencertoolsettings-overlap-mode"),
 	("bpy.types.spaceclipeditor.show_green_channel*", "editors/clip/display/clip_display.html#bpy-types-spaceclipeditor-show-green-channel"),
 	("bpy.types.spaceoutliner.show_restrict_column*", "editors/outliner/interface.html#bpy-types-spaceoutliner-show-restrict-column"),
@@ -443,8 +455,7 @@ url_manual_mapping = (
 	("bpy.types.freestylelineset.select_edge_mark*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset-select-edge-mark"),
 	("bpy.types.freestylelinestyle.use_length_max*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-use-length-max"),
 	("bpy.types.freestylelinestyle.use_length_min*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-use-length-min"),
-	("bpy.types.geometrynodealignrotationtovector*", "modeling/geometry_nodes/point/align_rotation_to_vector.html#bpy-types-geometrynodealignrotationtovector"),
-	("bpy.types.geometrynodeattributevectorrotate*", "modeling/geometry_nodes/attribute/attribute_vector_rotate.html#bpy-types-geometrynodeattributevectorrotate"),
+	("bpy.types.geometrynodeinputsplineresolution*", "modeling/geometry_nodes/curve/spline_resolution.html#bpy-types-geometrynodeinputsplineresolution"),
 	("bpy.types.greasepencil.curve_edit_threshold*", "grease_pencil/modes/edit/curve_editing.html#bpy-types-greasepencil-curve-edit-threshold"),
 	("bpy.types.materialgpencilstyle.stroke_style*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-stroke-style"),
 	("bpy.types.objectlineart.use_crease_override*", "scene_layout/object/properties/line_art.html#bpy-types-objectlineart-use-crease-override"),
@@ -501,6 +512,7 @@ url_manual_mapping = (
 	("bpy.types.freestylesettings.use_smoothness*", "render/freestyle/view_layer/freestyle.html#bpy-types-freestylesettings-use-smoothness"),
 	("bpy.types.geometrynodecurvequadraticbezier*", "modeling/geometry_nodes/curve_primitives/quadratic_bezier.html#bpy-types-geometrynodecurvequadraticbezier"),
 	("bpy.types.materialgpencilstyle.show_stroke*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-show-stroke"),
+	("bpy.types.mesh.use_customdata_vertex_bevel*", "modeling/meshes/properties/custom_data.html#bpy-types-mesh-use-customdata-vertex-bevel"),
 	("bpy.types.movietrackingcamera.focal_length*", "movie_clip/tracking/clip/sidebar/track/camera.html#bpy-types-movietrackingcamera-focal-length"),
 	("bpy.types.movietrackingcamera.pixel_aspect*", "movie_clip/tracking/clip/sidebar/track/camera.html#bpy-types-movietrackingcamera-pixel-aspect"),
 	("bpy.types.movietrackingcamera.sensor_width*", "movie_clip/tracking/clip/sidebar/track/camera.html#bpy-types-movietrackingcamera-sensor-width"),
@@ -546,7 +558,7 @@ url_manual_mapping = (
 	("bpy.types.freestylelinestyle.split_length*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-split-length"),
 	("bpy.types.freestylelinestyle.use_chaining*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-use-chaining"),
 	("bpy.types.freestylesettings.sphere_radius*", "render/freestyle/view_layer/freestyle.html#bpy-types-freestylesettings-sphere-radius"),
-	("bpy.types.geometrynodeattributevectormath*", "modeling/geometry_nodes/attribute/attribute_vector_math.html#bpy-types-geometrynodeattributevectormath"),
+	("bpy.types.geometrynodesetsplineresolution*", "modeling/geometry_nodes/curve/set_spline_resolution.html#bpy-types-geometrynodesetsplineresolution"),
 	("bpy.types.gpencillayer.annotation_opacity*", "interface/annotate_tool.html#bpy-types-gpencillayer-annotation-opacity"),
 	("bpy.types.gpencillayer.use_onion_skinning*", "grease_pencil/properties/layers.html#bpy-types-gpencillayer-use-onion-skinning"),
 	("bpy.types.gpencilsculptguide.use_snapping*", "grease_pencil/modes/draw/guides.html#bpy-types-gpencilsculptguide-use-snapping"),
@@ -558,6 +570,7 @@ url_manual_mapping = (
 	("bpy.types.materialgpencilstyle.fill_style*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-fill-style"),
 	("bpy.types.materialgpencilstyle.mix_factor*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-mix-factor"),
 	("bpy.types.materialgpencilstyle.pass_index*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-pass-index"),
+	("bpy.types.mesh.use_customdata_edge_crease*", "modeling/meshes/properties/custom_data.html#bpy-types-mesh-use-customdata-edge-crease"),
 	("bpy.types.nodesocketinterface.description*", "interface/controls/nodes/groups.html#bpy-types-nodesocketinterface-description"),
 	("bpy.types.rendersettings.dither_intensity*", "render/output/properties/post_processing.html#bpy-types-rendersettings-dither-intensity"),
 	("bpy.types.rendersettings.film_transparent*", "render/cycles/render_settings/film.html#bpy-types-rendersettings-film-transparent"),
@@ -613,18 +626,20 @@ url_manual_mapping = (
 	("bpy.types.freestylelinestyle.chain_count*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-chain-count"),
 	("bpy.types.freestylelinestyle.use_sorting*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-use-sorting"),
 	("bpy.types.freestylesettings.crease_angle*", "render/freestyle/view_layer/freestyle.html#bpy-types-freestylesettings-crease-angle"),
-	("bpy.types.geometrynodeattributecolorramp*", "modeling/geometry_nodes/attribute/attribute_color_ramp.html#bpy-types-geometrynodeattributecolorramp"),
-	("bpy.types.geometrynodeattributeproximity*", "modeling/geometry_nodes/attribute/attribute_proximity.html#bpy-types-geometrynodeattributeproximity"),
-	("bpy.types.geometrynodeattributerandomize*", "modeling/geometry_nodes/attribute/attribute_randomize.html#bpy-types-geometrynodeattributerandomize"),
+	("bpy.types.geometrynodealigneulertovector*", "modeling/geometry_nodes/utilities/align_euler_to_vector.html#bpy-types-geometrynodealigneulertovector"),
+	("bpy.types.geometrynodeattributestatistic*", "modeling/geometry_nodes/attribute/attribute_statistic.html#bpy-types-geometrynodeattributestatistic"),
 	("bpy.types.geometrynodecurvequadrilateral*", "modeling/geometry_nodes/curve_primitives/quadrilateral.html#bpy-types-geometrynodecurvequadrilateral"),
+	("bpy.types.geometrynodeinputmaterialindex*", "modeling/geometry_nodes/material/material_index.html#bpy-types-geometrynodeinputmaterialindex"),
 	("bpy.types.geometrynodeseparatecomponents*", "modeling/geometry_nodes/geometry/separate_components.html#bpy-types-geometrynodeseparatecomponents"),
 	("bpy.types.geometrynodesubdivisionsurface*", "modeling/geometry_nodes/mesh/subdivision_surface.html#bpy-types-geometrynodesubdivisionsurface"),
+	("bpy.types.geometrynodetranslateinstances*", "modeling/geometry_nodes/instances/translate_instances.html#bpy-types-geometrynodetranslateinstances"),
 	("bpy.types.imageformatsettings.color_mode*", "render/output/properties/output.html#bpy-types-imageformatsettings-color-mode"),
 	("bpy.types.linestyle*modifier_alongstroke*", "render/freestyle/view_layer/line_style/modifiers/color/along_stroke.html#bpy-types-linestyle-modifier-alongstroke"),
 	("bpy.types.linestyle*modifier_creaseangle*", "render/freestyle/view_layer/line_style/modifiers/color/crease_angle.html#bpy-types-linestyle-modifier-creaseangle"),
 	("bpy.types.linestylecolormodifier_tangent*", "render/freestyle/view_layer/line_style/modifiers/color/tangent.html#bpy-types-linestylecolormodifier-tangent"),
 	("bpy.types.materialgpencilstyle.mix_color*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-mix-color"),
 	("bpy.types.materialgpencilstyle.show_fill*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-show-fill"),
+	("bpy.types.mesh.use_customdata_edge_bevel*", "modeling/meshes/properties/custom_data.html#bpy-types-mesh-use-customdata-edge-bevel"),
 	("bpy.types.mesh.use_mirror_vertex_group_x*", "sculpt_paint/weight_paint/tool_settings/symmetry.html#bpy-types-mesh-use-mirror-vertex-group-x"),
 	("bpy.types.movietrackingcamera.division_k*", "movie_clip/tracking/clip/sidebar/track/camera.html#bpy-types-movietrackingcamera-division-k"),
 	("bpy.types.movietrackingobject.keyframe_a*", "movie_clip/tracking/clip/toolbar/solve.html#bpy-types-movietrackingobject-keyframe-a"),
@@ -685,9 +700,11 @@ url_manual_mapping = (
 	("bpy.types.freestylelinestyle.sort_order*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-sort-order"),
 	("bpy.types.freestylelinestyle.split_dash*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-split-dash"),
 	("bpy.types.freestylesettings.use_culling*", "render/freestyle/view_layer/freestyle.html#bpy-types-freestylesettings-use-culling"),
-	("bpy.types.geometrynodeattributecurvemap*", "modeling/geometry_nodes/attribute/attribute_curve_map.html#bpy-types-geometrynodeattributecurvemap"),
-	("bpy.types.geometrynodeattributemaprange*", "modeling/geometry_nodes/attribute/attribute_map_range.html#bpy-types-geometrynodeattributemaprange"),
-	("bpy.types.geometrynodeattributetransfer*", "modeling/geometry_nodes/attribute/attribute_transfer.html#bpy-types-geometrynodeattributetransfer"),
+	("bpy.types.geometrynodeendpointselection*", "modeling/geometry_nodes/curve/endpoint_selection.html#bpy-types-geometrynodeendpointselection"),
+	("bpy.types.geometrynodegeometryproximity*", "modeling/geometry_nodes/geometry/geometry_proximity.html#bpy-types-geometrynodegeometryproximity"),
+	("bpy.types.geometrynodeinputsplinecyclic*", "modeling/geometry_nodes/curve/is_spline_cyclic.html#bpy-types-geometrynodeinputsplinecyclic"),
+	("bpy.types.geometrynodeinstancestopoints*", "modeling/geometry_nodes/instances/instances_to_points.html#bpy-types-geometrynodeinstancestopoints"),
+	("bpy.types.geometrynodetransferattribute*", "modeling/geometry_nodes/attribute/transfer_attribute.html#bpy-types-geometrynodetransferattribute"),
 	("bpy.types.layercollection.hide_viewport*", "editors/outliner/interface.html#bpy-types-layercollection-hide-viewport"),
 	("bpy.types.layercollection.indirect_only*", "editors/outliner/interface.html#bpy-types-layercollection-indirect-only"),
 	("bpy.types.material.use_sss_translucency*", "render/eevee/materials/settings.html#bpy-types-material-use-sss-translucency"),
@@ -740,9 +757,12 @@ url_manual_mapping = (
 	("bpy.types.freestylelinestyle.angle_min*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-angle-min"),
 	("bpy.types.freestylelinestyle.split_gap*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-split-gap"),
 	("bpy.types.freestylelinestyle.use_nodes*", "render/freestyle/view_layer/line_style/texture.html#bpy-types-freestylelinestyle-use-nodes"),
-	("bpy.types.geometrynodeattributecompare*", "modeling/geometry_nodes/attribute/attribute_compare.html#bpy-types-geometrynodeattributecompare"),
-	("bpy.types.geometrynodeattributeconvert*", "modeling/geometry_nodes/attribute/attribute_convert.html#bpy-types-geometrynodeattributeconvert"),
-	("bpy.types.geometrynodeselectbymaterial*", "modeling/geometry_nodes/material/select_by_material.html#bpy-types-geometrynodeselectbymaterial"),
+	("bpy.types.geometrynodecaptureattribute*", "modeling/geometry_nodes/attribute/capture_attribute.html#bpy-types-geometrynodecaptureattribute"),
+	("bpy.types.geometrynodeinstanceonpoints*", "modeling/geometry_nodes/instances/instance_on_points.html#bpy-types-geometrynodeinstanceonpoints"),
+	("bpy.types.geometrynodepointstovertices*", "modeling/geometry_nodes/point/points_to_vertices.html#bpy-types-geometrynodepointstovertices"),
+	("bpy.types.geometrynoderealizeinstances*", "modeling/geometry_nodes/instances/realize_instances.html#bpy-types-geometrynoderealizeinstances"),
+	("bpy.types.geometrynodeseparategeometry*", "modeling/geometry_nodes/geometry/separate_geometry.html#bpy-types-geometrynodeseparategeometry"),
+	("bpy.types.geometrynodesetmaterialindex*", "modeling/geometry_nodes/material/set_material_index.html#bpy-types-geometrynodesetmaterialindex"),
 	("bpy.types.material.preview_render_type*", "render/materials/preview.html#bpy-types-material-preview-render-type"),
 	("bpy.types.materialgpencilstyle.pattern*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-pattern"),
 	("bpy.types.materialgpencilstyle.texture*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-texture"),
@@ -804,9 +824,11 @@ url_manual_mapping = (
 	("bpy.types.freestylelineset.visibility*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset-visibility"),
 	("bpy.types.freestylelinestyle.chaining*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-chaining"),
 	("bpy.types.freestylelinestyle.sort_key*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-sort-key"),
-	("bpy.types.geometrynodeattributeremove*", "modeling/geometry_nodes/attribute/attribute_remove.html#bpy-types-geometrynodeattributeremove"),
-	("bpy.types.geometrynodematerialreplace*", "modeling/geometry_nodes/material/replace.html#bpy-types-geometrynodematerialreplace"),
-	("bpy.types.geometrynodepointdistribute*", "modeling/geometry_nodes/point/point_distribute.html#bpy-types-geometrynodepointdistribute"),
+	("bpy.types.geometrynodecurvesethandles*", "modeling/geometry_nodes/curve/set_handle_type.html#bpy-types-geometrynodecurvesethandles"),
+	("bpy.types.geometrynodecurvesplinetype*", "modeling/geometry_nodes/curve/set_spline_type.html#bpy-types-geometrynodecurvesplinetype"),
+	("bpy.types.geometrynodereplacematerial*", "modeling/geometry_nodes/material/replace_material.html#bpy-types-geometrynodereplacematerial"),
+	("bpy.types.geometrynoderotateinstances*", "modeling/geometry_nodes/instances/rotate_instances.html#bpy-types-geometrynoderotateinstances"),
+	("bpy.types.geometrynodesetsplinecyclic*", "modeling/geometry_nodes/curve/set_spline_cyclic.html#bpy-types-geometrynodesetsplinecyclic"),
 	("bpy.types.gpencillayer.use_mask_layer*", "grease_pencil/properties/layers.html#bpy-types-gpencillayer-use-mask-layer"),
 	("bpy.types.greasepencil.use_curve_edit*", "grease_pencil/modes/edit/curve_editing.html#bpy-types-greasepencil-use-curve-edit"),
 	("bpy.types.imagepaint.screen_grab_size*", "sculpt_paint/texture_paint/tool_settings/options.html#bpy-types-imagepaint-screen-grab-size"),
@@ -871,14 +893,17 @@ url_manual_mapping = (
 	("bpy.types.fileselectparams.directory*", "editors/file_browser.html#bpy-types-fileselectparams-directory"),
 	("bpy.types.fluidflowsettings.use_flow*", "physics/fluid/type/flow.html#bpy-types-fluidflowsettings-use-flow"),
 	("bpy.types.fmodifierfunctiongenerator*", "editors/graph_editor/fcurves/modifiers.html#bpy-types-fmodifierfunctiongenerator"),
-	("bpy.types.geometrynodeattributeclamp*", "modeling/geometry_nodes/attribute/attribute_clamp.html#bpy-types-geometrynodeattributeclamp"),
 	("bpy.types.geometrynodecollectioninfo*", "modeling/geometry_nodes/input/collection_info.html#bpy-types-geometrynodecollectioninfo"),
-	("bpy.types.geometrynodecurveendpoints*", "modeling/geometry_nodes/curve/curve_endpoints.html#bpy-types-geometrynodecurveendpoints"),
-	("bpy.types.geometrynodecurvesubdivide*", "modeling/geometry_nodes/curve/curve_subdivide.html#bpy-types-geometrynodecurvesubdivide"),
+	("bpy.types.geometrynodecurveparameter*", "modeling/geometry_nodes/curve/curve_parameter.html#bpy-types-geometrynodecurveparameter"),
 	("bpy.types.geometrynodedeletegeometry*", "modeling/geometry_nodes/geometry/delete_geometry.html#bpy-types-geometrynodedeletegeometry"),
-	("bpy.types.geometrynodematerialassign*", "modeling/geometry_nodes/material/assign.html#bpy-types-geometrynodematerialassign"),
-	("bpy.types.geometrynodepointstovolume*", "modeling/geometry_nodes/volume/points_to_volume.html#bpy-types-geometrynodepointstovolume"),
-	("bpy.types.geometrynodepointtranslate*", "modeling/geometry_nodes/point/point_translate.html#bpy-types-geometrynodepointtranslate"),
+	("bpy.types.geometrynodeinputcurvetilt*", "modeling/geometry_nodes/curve/curve_tilt.html#bpy-types-geometrynodeinputcurvetilt"),
+	("bpy.types.geometrynodepointstovolume*", "modeling/geometry_nodes/point/points_to_volume.html#bpy-types-geometrynodepointstovolume"),
+	("bpy.types.geometrynodescaleinstances*", "modeling/geometry_nodes/instances/scale_instances.html#bpy-types-geometrynodescaleinstances"),
+	("bpy.types.geometrynodesetcurveradius*", "modeling/geometry_nodes/curve/set_curve_radius.html#bpy-types-geometrynodesetcurveradius"),
+	("bpy.types.geometrynodesetpointradius*", "modeling/geometry_nodes/point/set_point_radius.html#bpy-types-geometrynodesetpointradius"),
+	("bpy.types.geometrynodesetshadesmooth*", "modeling/geometry_nodes/mesh/set_shade_smooth.html#bpy-types-geometrynodesetshadesmooth"),
+	("bpy.types.geometrynodestringtocurves*", "modeling/geometry_nodes/text/string_to_curves.html#bpy-types-geometrynodestringtocurves"),
+	("bpy.types.geometrynodesubdividecurve*", "modeling/geometry_nodes/curve/subdivide_curve.html#bpy-types-geometrynodesubdividecurve"),
 	("bpy.types.greasepencil.use_multiedit*", "grease_pencil/multiframe.html#bpy-types-greasepencil-use-multiedit"),
 	("bpy.types.keyframe.handle_right_type*", "editors/graph_editor/fcurves/properties.html#bpy-types-keyframe-handle-right-type"),
 	("bpy.types.materialgpencilstyle.color*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-color"),
@@ -947,15 +972,16 @@ url_manual_mapping = (
 	("bpy.types.fluidflowsettings.density*", "physics/fluid/type/flow.html#bpy-types-fluidflowsettings-density"),
 	("bpy.types.freestylelineset.qi_start*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset-qi-start"),
 	("bpy.types.freestylelinestyle.rounds*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-rounds"),
-	("bpy.types.geometrynodeattributefill*", "modeling/geometry_nodes/attribute/attribute_fill.html#bpy-types-geometrynodeattributefill"),
-	("bpy.types.geometrynodeattributemath*", "modeling/geometry_nodes/attribute/attribute_math.html#bpy-types-geometrynodeattributemath"),
+	("bpy.types.functionnodecomparefloats*", "modeling/geometry_nodes/utilities/compare_floats.html#bpy-types-functionnodecomparefloats"),
 	("bpy.types.geometrynodecurvetopoints*", "modeling/geometry_nodes/curve/curve_to_points.html#bpy-types-geometrynodecurvetopoints"),
 	("bpy.types.geometrynodeinputmaterial*", "modeling/geometry_nodes/input/material.html#bpy-types-geometrynodeinputmaterial"),
+	("bpy.types.geometrynodeinputposition*", "modeling/geometry_nodes/input/position.html#bpy-types-geometrynodeinputposition"),
+	("bpy.types.geometrynodeisshadesmooth*", "modeling/geometry_nodes/mesh/is_shade_smooth.html#bpy-types-geometrynodeisshadesmooth"),
 	("bpy.types.geometrynodemeshicosphere*", "modeling/geometry_nodes/mesh_primitives/icosphere.html#bpy-types-geometrynodemeshicosphere"),
-	("bpy.types.geometrynodemeshsubdivide*", "modeling/geometry_nodes/mesh/subdivide.html#bpy-types-geometrynodemeshsubdivide"),
-	("bpy.types.geometrynodepointinstance*", "modeling/geometry_nodes/point/point_instance.html#bpy-types-geometrynodepointinstance"),
-	("bpy.types.geometrynodepointseparate*", "modeling/geometry_nodes/point/point_separate.html#bpy-types-geometrynodepointseparate"),
+	("bpy.types.geometrynodereplacestring*", "modeling/geometry_nodes/text/replace_string.html#bpy-types-geometrynodereplacestring"),
 	("bpy.types.geometrynoderesamplecurve*", "modeling/geometry_nodes/curve/resample_curve.html#bpy-types-geometrynoderesamplecurve"),
+	("bpy.types.geometrynodesubdividemesh*", "modeling/geometry_nodes/mesh/subdivide_mesh.html#bpy-types-geometrynodesubdividemesh"),
+	("bpy.types.geometrynodevaluetostring*", "modeling/geometry_nodes/text/value_to_string.html#bpy-types-geometrynodevaluetostring"),
 	("bpy.types.keyframe.handle_left_type*", "editors/graph_editor/fcurves/properties.html#bpy-types-keyframe-handle-left-type"),
 	("bpy.types.light.use_custom_distance*", "render/eevee/lighting.html#bpy-types-light-use-custom-distance"),
 	("bpy.types.materialgpencilstyle.flip*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-flip"),
@@ -988,7 +1014,8 @@ url_manual_mapping = (
 	("bpy.types.volumedisplay.slice_depth*", "modeling/volumes/properties.html#bpy-types-volumedisplay-slice-depth"),
 	("bpy.types.worldmistsettings.falloff*", "render/cycles/world_settings.html#bpy-types-worldmistsettings-falloff"),
 	("bpy.ops.clip.lock_selection_toggle*", "editors/clip/introduction.html#bpy-ops-clip-lock-selection-toggle"),
-	("bpy.ops.mesh.customdata_mask_clear*", "sculpt_paint/sculpting/editing/mask.html#bpy-ops-mesh-customdata-mask-clear"),
+	("bpy.ops.mesh.customdata_mask_clear*", "modeling/meshes/properties/custom_data.html#bpy-ops-mesh-customdata-mask-clear"),
+	("bpy.ops.mesh.customdata_skin_clear*", "modeling/meshes/properties/custom_data.html#bpy-ops-mesh-customdata-skin-clear"),
 	("bpy.ops.mesh.extrude_vertices_move*", "modeling/meshes/editing/vertex/extrude_vertices.html#bpy-ops-mesh-extrude-vertices-move"),
 	("bpy.ops.mesh.mod_weighted_strength*", "modeling/meshes/editing/mesh/normals.html#bpy-ops-mesh-mod-weighted-strength"),
 	("bpy.ops.mesh.quads_convert_to_tris*", "modeling/meshes/editing/face/triangulate_faces.html#bpy-ops-mesh-quads-convert-to-tris"),
@@ -1035,14 +1062,18 @@ url_manual_mapping = (
 	("bpy.types.freestylelineset.exclude*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset-exclude"),
 	("bpy.types.freestylelinestyle.alpha*", "render/freestyle/view_layer/line_style/alpha.html#bpy-types-freestylelinestyle-alpha"),
 	("bpy.types.freestylelinestyle.color*", "render/freestyle/view_layer/line_style/color.html#bpy-types-freestylelinestyle-color"),
-	("bpy.types.functionnodefloatcompare*", "modeling/geometry_nodes/utilities/float_compare.html#bpy-types-functionnodefloatcompare"),
-	("bpy.types.geometrynodeattributemix*", "modeling/geometry_nodes/attribute/attribute_mix.html#bpy-types-geometrynodeattributemix"),
-	("bpy.types.geometrynodecurvereverse*", "modeling/geometry_nodes/curve/curve_reverse.html#bpy-types-geometrynodecurvereverse"),
+	("bpy.types.geometrynodeinputtangent*", "modeling/geometry_nodes/curve/curve_tangent.html#bpy-types-geometrynodeinputtangent"),
 	("bpy.types.geometrynodejoingeometry*", "modeling/geometry_nodes/geometry/join_geometry.html#bpy-types-geometrynodejoingeometry"),
 	("bpy.types.geometrynodemeshcylinder*", "modeling/geometry_nodes/mesh_primitives/cylinder.html#bpy-types-geometrynodemeshcylinder"),
+	("bpy.types.geometrynodemeshtopoints*", "modeling/geometry_nodes/mesh/mesh_to_points.html#bpy-types-geometrynodemeshtopoints"),
 	("bpy.types.geometrynodemeshuvsphere*", "modeling/geometry_nodes/mesh_primitives/uv_sphere.html#bpy-types-geometrynodemeshuvsphere"),
+	("bpy.types.geometrynodereversecurve*", "modeling/geometry_nodes/curve/reverse_curve.html#bpy-types-geometrynodereversecurve"),
+	("bpy.types.geometrynodesetcurvetilt*", "modeling/geometry_nodes/curve/set_curve_tilt.html#bpy-types-geometrynodesetcurvetilt"),
+	("bpy.types.geometrynodesplinelength*", "modeling/geometry_nodes/curve/spline_length.html#bpy-types-geometrynodesplinelength"),
+	("bpy.types.geometrynodestringlength*", "modeling/geometry_nodes/text/string_length.html#bpy-types-geometrynodestringlength"),
 	("bpy.types.geometrynodevolumetomesh*", "modeling/geometry_nodes/volume/volume_to_mesh.html#bpy-types-geometrynodevolumetomesh"),
 	("bpy.types.image.use_half_precision*", "editors/image/image_settings.html#bpy-types-image-use-half-precision"),
+	("bpy.types.image.use_view_as_render*", "editors/image/image_settings.html#bpy-types-image-use-view-as-render"),
 	("bpy.types.imagepaint.interpolation*", "sculpt_paint/texture_paint/tool_settings/texture_slots.html#bpy-types-imagepaint-interpolation"),
 	("bpy.types.linestyle*modifier_noise*", "render/freestyle/view_layer/line_style/modifiers/color/noise.html#bpy-types-linestyle-modifier-noise"),
 	("bpy.types.maintainvolumeconstraint*", "animation/constraints/transform/maintain_volume.html#bpy-types-maintainvolumeconstraint"),
@@ -1125,13 +1156,22 @@ url_manual_mapping = (
 	("bpy.types.functionnodebooleanmath*", "modeling/geometry_nodes/utilities/boolean_math.html#bpy-types-functionnodebooleanmath"),
 	("bpy.types.functionnodeinputstring*", "modeling/geometry_nodes/input/string.html#bpy-types-functionnodeinputstring"),
 	("bpy.types.functionnodeinputvector*", "modeling/geometry_nodes/input/vector.html#bpy-types-functionnodeinputvector"),
-	("bpy.types.functionnoderandomfloat*", "modeling/geometry_nodes/input/random_float.html#bpy-types-functionnoderandomfloat"),
-	("bpy.types.geometrynodecurvecircle*", "modeling/geometry_nodes/curve_primitives/circle.html#bpy-types-geometrynodecurvecircle"),
+	("bpy.types.geometrynodecurvecircle*", "modeling/geometry_nodes/curve_primitives/curve_circle.html#bpy-types-geometrynodecurvecircle"),
 	("bpy.types.geometrynodecurvelength*", "modeling/geometry_nodes/curve/curve_length.html#bpy-types-geometrynodecurvelength"),
-	("bpy.types.geometrynodecurvespiral*", "modeling/geometry_nodes/curve_primitives/spiral.html#bpy-types-geometrynodecurvespiral"),
+	("bpy.types.geometrynodecurvespiral*", "modeling/geometry_nodes/curve_primitives/curve_spiral.html#bpy-types-geometrynodecurvespiral"),
 	("bpy.types.geometrynodecurvetomesh*", "modeling/geometry_nodes/curve/curve_to_mesh.html#bpy-types-geometrynodecurvetomesh"),
-	("bpy.types.geometrynodemeshtocurve*", "modeling/geometry_nodes/curve/mesh_to_curve.html#bpy-types-geometrynodemeshtocurve"),
-	("bpy.types.geometrynodepointrotate*", "modeling/geometry_nodes/point/point_rotate.html#bpy-types-geometrynodepointrotate"),
+	("bpy.types.geometrynodefilletcurve*", "modeling/geometry_nodes/curve/fillet_curve.html#bpy-types-geometrynodefilletcurve"),
+	("bpy.types.geometrynodeinputnormal*", "modeling/geometry_nodes/input/normal.html#bpy-types-geometrynodeinputnormal"),
+	("bpy.types.geometrynodeinputradius*", "modeling/geometry_nodes/input/radius.html#bpy-types-geometrynodeinputradius"),
+	("bpy.types.geometrynodejoinstrings*", "modeling/geometry_nodes/text/join_strings.html#bpy-types-geometrynodejoinstrings"),
+	("bpy.types.geometrynodemeshboolean*", "modeling/geometry_nodes/mesh/mesh_boolean.html#bpy-types-geometrynodemeshboolean"),
+	("bpy.types.geometrynodemeshtocurve*", "modeling/geometry_nodes/mesh/mesh_to_curve.html#bpy-types-geometrynodemeshtocurve"),
+	("bpy.types.geometrynoderandomvalue*", "modeling/geometry_nodes/utilities/random_value.html#bpy-types-geometrynoderandomvalue"),
+	("bpy.types.geometrynoderotateeuler*", "modeling/geometry_nodes/utilities/rotate_euler.html#bpy-types-geometrynoderotateeuler"),
+	("bpy.types.geometrynodesamplecurve*", "modeling/geometry_nodes/curve/sample_curve.html#bpy-types-geometrynodesamplecurve"),
+	("bpy.types.geometrynodesetmaterial*", "modeling/geometry_nodes/material/set_material.html#bpy-types-geometrynodesetmaterial"),
+	("bpy.types.geometrynodesetposition*", "modeling/geometry_nodes/geometry/set_position.html#bpy-types-geometrynodesetposition"),
+	("bpy.types.geometrynodeslicestring*", "modeling/geometry_nodes/text/slice_string.html#bpy-types-geometrynodeslicestring"),
 	("bpy.types.geometrynodetriangulate*", "modeling/geometry_nodes/mesh/triangulate.html#bpy-types-geometrynodetriangulate"),
 	("bpy.types.gpencillayer.blend_mode*", "grease_pencil/properties/layers.html#bpy-types-gpencillayer-blend-mode"),
 	("bpy.types.gpencillayer.use_lights*", "grease_pencil/properties/layers.html#bpy-types-gpencillayer-use-lights"),
@@ -1172,6 +1212,7 @@ url_manual_mapping = (
 	("bpy.ops.gpencil.vertex_color_set*", "grease_pencil/modes/vertex_paint/editing.html#bpy-ops-gpencil-vertex-color-set"),
 	("bpy.ops.graph.extrapolation_type*", "editors/graph_editor/channels.html#bpy-ops-graph-extrapolation-type"),
 	("bpy.ops.graph.interpolation_type*", "editors/graph_editor/fcurves/editing.html#bpy-ops-graph-interpolation-type"),
+	("bpy.ops.mesh.customdata_skin_add*", "modeling/meshes/properties/custom_data.html#bpy-ops-mesh-customdata-skin-add"),
 	("bpy.ops.mesh.dissolve_degenerate*", "modeling/meshes/editing/mesh/cleanup.html#bpy-ops-mesh-dissolve-degenerate"),
 	("bpy.ops.mesh.face_split_by_edges*", "modeling/meshes/editing/face/weld_edges_faces.html#bpy-ops-mesh-face-split-by-edges"),
 	("bpy.ops.mesh.mark_freestyle_face*", "modeling/meshes/editing/face/face_data.html#bpy-ops-mesh-mark-freestyle-face"),
@@ -1234,11 +1275,12 @@ url_manual_mapping = (
 	("bpy.types.freestylelinestyle.gap*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-gap"),
 	("bpy.types.freestylesettings.mode*", "render/freestyle/view_layer/freestyle.html#bpy-types-freestylesettings-mode"),
 	("bpy.types.geometrynodeconvexhull*", "modeling/geometry_nodes/geometry/convex_hull.html#bpy-types-geometrynodeconvexhull"),
-	("bpy.types.geometrynodefloattoint*", "modeling/geometry_nodes/utilities/float_to_int.html#bpy-types-geometrynodefloattoint"),
+	("bpy.types.geometrynodefloattoint*", "modeling/geometry_nodes/utilities/float_to_integer.html#bpy-types-geometrynodefloattoint"),
+	("bpy.types.geometrynodeinputcolor*", "modeling/geometry_nodes/input/color.html#bpy-types-geometrynodeinputcolor"),
+	("bpy.types.geometrynodeinputindex*", "modeling/geometry_nodes/input/input_index.html#bpy-types-geometrynodeinputindex"),
 	("bpy.types.geometrynodeisviewport*", "modeling/geometry_nodes/input/is_viewport.html#bpy-types-geometrynodeisviewport"),
-	("bpy.types.geometrynodemeshcircle*", "modeling/geometry_nodes/mesh_primitives/circle.html#bpy-types-geometrynodemeshcircle"),
+	("bpy.types.geometrynodemeshcircle*", "modeling/geometry_nodes/mesh_primitives/mesh_circle.html#bpy-types-geometrynodemeshcircle"),
 	("bpy.types.geometrynodeobjectinfo*", "modeling/geometry_nodes/input/object_info.html#bpy-types-geometrynodeobjectinfo"),
-	("bpy.types.geometrynodepointscale*", "modeling/geometry_nodes/point/point_scale.html#bpy-types-geometrynodepointscale"),
 	("bpy.types.imagepaint.use_occlude*", "sculpt_paint/texture_paint/tool_settings/options.html#bpy-types-imagepaint-use-occlude"),
 	("bpy.types.imagesequence.use_flip*", "video_editing/sequencer/sidebar/strip.html#bpy-types-imagesequence-use-flip"),
 	("bpy.types.keyframe.interpolation*", "editors/graph_editor/fcurves/properties.html#bpy-types-keyframe-interpolation"),
@@ -1347,11 +1389,12 @@ url_manual_mapping = (
 	("bpy.types.fcurve.auto_smoothing*", "editors/graph_editor/fcurves/properties.html#bpy-types-fcurve-auto-smoothing"),
 	("bpy.types.fluideffectorsettings*", "physics/fluid/type/effector.html#bpy-types-fluideffectorsettings"),
 	("bpy.types.followtrackconstraint*", "animation/constraints/motion_tracking/follow_track.html#bpy-types-followtrackconstraint"),
-	("bpy.types.geometrynodecurveline*", "modeling/geometry_nodes/curve_primitives/line.html#bpy-types-geometrynodecurveline"),
+	("bpy.types.geometrynodecurveline*", "modeling/geometry_nodes/curve_primitives/curve_line.html#bpy-types-geometrynodecurveline"),
 	("bpy.types.geometrynodecurvestar*", "modeling/geometry_nodes/curve_primitives/star.html#bpy-types-geometrynodecurvestar"),
-	("bpy.types.geometrynodecurvetrim*", "modeling/geometry_nodes/curve/curve_trim.html#bpy-types-geometrynodecurvetrim"),
-	("bpy.types.geometrynodeedgesplit*", "modeling/geometry_nodes/mesh/edge_split.html#bpy-types-geometrynodeedgesplit"),
+	("bpy.types.geometrynodeedgesplit*", "modeling/geometry_nodes/mesh/split_edges.html#bpy-types-geometrynodeedgesplit"),
+	("bpy.types.geometrynodefillcurve*", "modeling/geometry_nodes/curve/fill_curve.html#bpy-types-geometrynodefillcurve"),
 	("bpy.types.geometrynodetransform*", "modeling/geometry_nodes/geometry/transform.html#bpy-types-geometrynodetransform"),
+	("bpy.types.geometrynodetrimcurve*", "modeling/geometry_nodes/curve/trim_curve.html#bpy-types-geometrynodetrimcurve"),
 	("bpy.types.gpencilsculptsettings*", "grease_pencil/properties/index.html#bpy-types-gpencilsculptsettings"),
 	("bpy.types.keyframe.handle_right*", "editors/graph_editor/fcurves/properties.html#bpy-types-keyframe-handle-right"),
 	("bpy.types.light.cutoff_distance*", "render/eevee/lighting.html#bpy-types-light-cutoff-distance"),
@@ -1405,6 +1448,7 @@ url_manual_mapping = (
 	("bpy.ops.outliner.orphans_purge*", "editors/outliner/interface.html#bpy-ops-outliner-orphans-purge"),
 	("bpy.ops.paint.mask_box_gesture*", "sculpt_paint/sculpting/editing/mask.html#bpy-ops-paint-mask-box-gesture"),
 	("bpy.ops.screen.region_quadview*", "editors/3dview/navigate/views.html#bpy-ops-screen-region-quadview"),
+	("bpy.ops.screen.screenshot_area*", "interface/window_system/topbar.html#bpy-ops-screen-screenshot-area"),
 	("bpy.ops.sequencer.offset_clear*", "video_editing/sequencer/editing.html#bpy-ops-sequencer-offset-clear"),
 	("bpy.ops.spreadsheet.toggle_pin*", "editors/spreadsheet.html#bpy-ops-spreadsheet-toggle-pin"),
 	("bpy.ops.uv.follow_active_quads*", "modeling/meshes/editing/uv.html#bpy-ops-uv-follow-active-quads"),
@@ -1412,6 +1456,7 @@ url_manual_mapping = (
 	("bpy.types.bone.envelope_weight*", "animation/armatures/bones/properties/deform.html#bpy-types-bone-envelope-weight"),
 	("bpy.types.brush.use_persistent*", "sculpt_paint/sculpting/tools/layer.html#bpy-types-brush-use-persistent"),
 	("bpy.types.buildgpencilmodifier*", "grease_pencil/modifiers/generate/build.html#bpy-types-buildgpencilmodifier"),
+	("bpy.types.camera.sensor_height*", "render/cameras.html#bpy-types-camera-sensor-height"),
 	("bpy.types.colorbalancemodifier*", "video_editing/sequencer/sidebar/modifiers.html#bpy-types-colorbalancemodifier"),
 	("bpy.types.colorgpencilmodifier*", "grease_pencil/modifiers/color/hue_saturation.html#bpy-types-colorgpencilmodifier"),
 	("bpy.types.compositornodefilter*", "compositing/types/filter/filter_node.html#bpy-types-compositornodefilter"),
@@ -1433,10 +1478,11 @@ url_manual_mapping = (
 	("bpy.types.followpathconstraint*", "animation/constraints/relationship/follow_path.html#bpy-types-followpathconstraint"),
 	("bpy.types.gaussianblursequence*", "video_editing/sequencer/strips/effects/blur.html#bpy-types-gaussianblursequence"),
 	("bpy.types.geometrynodeboundbox*", "modeling/geometry_nodes/geometry/bounding_box.html#bpy-types-geometrynodeboundbox"),
+	("bpy.types.geometrynodematerial*", "-1"),
 	("bpy.types.geometrynodemeshcone*", "modeling/geometry_nodes/mesh_primitives/cone.html#bpy-types-geometrynodemeshcone"),
 	("bpy.types.geometrynodemeshcube*", "modeling/geometry_nodes/mesh_primitives/cube.html#bpy-types-geometrynodemeshcube"),
 	("bpy.types.geometrynodemeshgrid*", "modeling/geometry_nodes/mesh_primitives/grid.html#bpy-types-geometrynodemeshgrid"),
-	("bpy.types.geometrynodemeshline*", "modeling/geometry_nodes/mesh_primitives/line.html#bpy-types-geometrynodemeshline"),
+	("bpy.types.geometrynodemeshline*", "modeling/geometry_nodes/mesh_primitives/mesh_line.html#bpy-types-geometrynodemeshline"),
 	("bpy.types.gpencillayer.opacity*", "grease_pencil/properties/layers.html#bpy-types-gpencillayer-opacity"),
 	("bpy.types.image.display_aspect*", "editors/image/sidebar.html#bpy-types-image-display-aspect"),
 	("bpy.types.keyframe.handle_left*", "editors/graph_editor/fcurves/properties.html#bpy-types-keyframe-handle-left"),
@@ -1457,6 +1503,7 @@ url_manual_mapping = (
 	("bpy.types.shadernodebsdfglossy*", "render/shader_nodes/shader/glossy.html#bpy-types-shadernodebsdfglossy"),
 	("bpy.types.shadernodebsdfvelvet*", "render/shader_nodes/shader/velvet.html#bpy-types-shadernodebsdfvelvet"),
 	("bpy.types.shadernodecameradata*", "render/shader_nodes/input/camera_data.html#bpy-types-shadernodecameradata"),
+	("bpy.types.shadernodefloatcurve*", "render/shader_nodes/converter/float_curve.html#bpy-types-shadernodefloatcurve"),
 	("bpy.types.shadernodeobjectinfo*", "render/shader_nodes/input/object_info.html#bpy-types-shadernodeobjectinfo"),
 	("bpy.types.shadernodetexchecker*", "render/shader_nodes/textures/checker.html#bpy-types-shadernodetexchecker"),
 	("bpy.types.shadernodetexvoronoi*", "render/shader_nodes/textures/voronoi.html#bpy-types-shadernodetexvoronoi"),
@@ -1532,6 +1579,7 @@ url_manual_mapping = (
 	("bpy.types.brush.smooth_stroke*", "grease_pencil/modes/draw/tools/draw.html#bpy-types-brush-smooth-stroke"),
 	("bpy.types.brush.tip_roundness*", "sculpt_paint/sculpting/tools/clay_strips.html#bpy-types-brush-tip-roundness"),
 	("bpy.types.camera.display_size*", "render/cameras.html#bpy-types-camera-display-size"),
+	("bpy.types.camera.sensor_width*", "render/cameras.html#bpy-types-camera-sensor-width"),
 	("bpy.types.compositornodedblur*", "compositing/types/filter/directional_blur.html#bpy-types-compositornodedblur"),
 	("bpy.types.compositornodegamma*", "compositing/types/color/gamma.html#bpy-types-compositornodegamma"),
 	("bpy.types.compositornodeglare*", "compositing/types/filter/glare.html#bpy-types-compositornodeglare"),
@@ -1546,8 +1594,11 @@ url_manual_mapping = (
 	("bpy.types.curve.use_map_taper*", "modeling/curves/properties/geometry.html#bpy-types-curve-use-map-taper"),
 	("bpy.types.cyclesworldsettings*", "render/cycles/world_settings.html#bpy-types-cyclesworldsettings"),
 	("bpy.types.fluiddomainsettings*", "physics/fluid/type/domain/index.html#bpy-types-fluiddomainsettings"),
-	("bpy.types.geometrynodeboolean*", "modeling/geometry_nodes/mesh/boolean.html#bpy-types-geometrynodeboolean"),
+	("bpy.types.geometrynodeboolean*", "modeling/geometry_nodes/input/boolean.html#bpy-types-geometrynodeboolean"),
+	("bpy.types.geometrynodeinputid*", "modeling/geometry_nodes/input/id.html#bpy-types-geometrynodeinputid"),
+	("bpy.types.geometrynodeinteger*", "modeling/geometry_nodes/input/integer.html#bpy-types-geometrynodeinteger"),
 	("bpy.types.geometrynoderaycast*", "modeling/geometry_nodes/geometry/raycast.html#bpy-types-geometrynoderaycast"),
+	("bpy.types.geonodeimagetexture*", "modeling/geometry_nodes/texture/image.html#bpy-types-geonodeimagetexture"),
 	("bpy.types.hookgpencilmodifier*", "grease_pencil/modifiers/deform/hook.html#bpy-types-hookgpencilmodifier"),
 	("bpy.types.imageformatsettings*", "files/media/image_formats.html#bpy-types-imageformatsettings"),
 	("bpy.types.kinematicconstraint*", "animation/constraints/tracking/ik_solver.html#bpy-types-kinematicconstraint"),
@@ -1618,7 +1669,7 @@ url_manual_mapping = (
 	("bpy.ops.object.select_linked*", "scene_layout/object/selecting.html#bpy-ops-object-select-linked"),
 	("bpy.ops.object.select_mirror*", "scene_layout/object/selecting.html#bpy-ops-object-select-mirror"),
 	("bpy.ops.object.select_random*", "scene_layout/object/selecting.html#bpy-ops-object-select-random"),
-	("bpy.ops.object.transfer_mode*", "sculpt_paint/sculpting/editing/sculpt.html#bpy-ops-object-transfer-mode"),
+	("bpy.ops.object.transfer_mode*", "editors/3dview/modes.html#bpy-ops-object-transfer-mode"),
 	("bpy.ops.outliner.show_active*", "editors/outliner/editing.html#bpy-ops-outliner-show-active"),
 	("bpy.ops.paint.add_simple_uvs*", "sculpt_paint/texture_paint/tool_settings/texture_slots.html#bpy-ops-paint-add-simple-uvs"),
 	("bpy.ops.pose.rigify_generate*", "addons/rigging/rigify/basics.html#bpy-ops-pose-rigify-generate"),
@@ -1664,6 +1715,7 @@ url_manual_mapping = (
 	("bpy.types.geometrynodeviewer*", "modeling/geometry_nodes/output/viewer.html#bpy-types-geometrynodeviewer"),
 	("bpy.types.gpencilsculptguide*", "grease_pencil/modes/draw/guides.html#bpy-types-gpencilsculptguide"),
 	("bpy.types.huecorrectmodifier*", "video_editing/sequencer/sidebar/modifiers.html#bpy-types-huecorrectmodifier"),
+	("bpy.types.image.is_multiview*", "editors/image/image_settings.html#bpy-types-image-is-multiview"),
 	("bpy.types.imagepaint.stencil*", "sculpt_paint/texture_paint/tool_settings/mask.html#bpy-types-imagepaint-stencil"),
 	("bpy.types.meshdeformmodifier*", "modeling/modifiers/deform/mesh_deform.html#bpy-types-meshdeformmodifier"),
 	("bpy.types.movietrackingtrack*", "movie_clip/tracking/clip/sidebar/track/index.html#bpy-types-movietrackingtrack"),
@@ -1763,6 +1815,7 @@ url_manual_mapping = (
 	("bpy.types.armatureeditbones*", "animation/armatures/bones/editing/index.html#bpy-types-armatureeditbones"),
 	("bpy.types.brush.pose_offset*", "sculpt_paint/sculpting/tools/pose.html#bpy-types-brush-pose-offset"),
 	("bpy.types.brush.rake_factor*", "sculpt_paint/sculpting/tools/snake_hook.html#bpy-types-brush-rake-factor"),
+	("bpy.types.camera.sensor_fit*", "render/cameras.html#bpy-types-camera-sensor-fit"),
 	("bpy.types.cameradofsettings*", "render/cameras.html#bpy-types-cameradofsettings"),
 	("bpy.types.childofconstraint*", "animation/constraints/relationship/child_of.html#bpy-types-childofconstraint"),
 	("bpy.types.clamptoconstraint*", "animation/constraints/tracking/clamp_to.html#bpy-types-clamptoconstraint"),
@@ -1778,6 +1831,7 @@ url_manual_mapping = (
 	("bpy.types.fmodifierenvelope*", "editors/graph_editor/fcurves/modifiers.html#bpy-types-fmodifierenvelope"),
 	("bpy.types.freestylesettings*", "render/freestyle/view_layer/freestyle.html#bpy-types-freestylesettings"),
 	("bpy.types.geometrynodegroup*", "modeling/geometry_nodes/group.html#bpy-types-geometrynodegroup"),
+	("bpy.types.geometrynodesetid*", "modeling/geometry_nodes/geometry/set_id.html#bpy-types-geometrynodesetid"),
 	("bpy.types.gpencillayer.hide*", "grease_pencil/properties/layers.html#bpy-types-gpencillayer-hide"),
 	("bpy.types.gpencillayer.lock*", "grease_pencil/properties/layers.html#bpy-types-gpencillayer-lock"),
 	("bpy.types.imagepaint.dither*", "sculpt_paint/texture_paint/tool_settings/options.html#bpy-types-imagepaint-dither"),
@@ -1860,6 +1914,7 @@ url_manual_mapping = (
 	("bpy.ops.transform.tosphere*", "modeling/meshes/editing/mesh/transform/to_sphere.html#bpy-ops-transform-tosphere"),
 	("bpy.ops.view3d.clip_border*", "editors/3dview/navigate/regions.html#bpy-ops-view3d-clip-border"),
 	("bpy.ops.wm.previews_ensure*", "files/blend/previews.html#bpy-ops-wm-previews-ensure"),
+	("bpy.ops.wm.properties_edit*", "files/data_blocks.html#bpy-ops-wm-properties-edit"),
 	("bpy.ops.wm.search_operator*", "interface/controls/templates/operator_search.html#bpy-ops-wm-search-operator"),
 	("bpy.types.actionconstraint*", "animation/constraints/relationship/action.html#bpy-types-actionconstraint"),
 	("bpy.types.addonpreferences*", "editors/preferences/addons.html#bpy-types-addonpreferences"),
@@ -1883,8 +1938,8 @@ url_manual_mapping = (
 	("bpy.types.fileselectparams*", "editors/file_browser.html#bpy-types-fileselectparams"),
 	("bpy.types.fmodifierstepped*", "editors/graph_editor/fcurves/modifiers.html#bpy-types-fmodifierstepped"),
 	("bpy.types.freestylelineset*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset"),
+	("bpy.types.image.alpha_mode*", "editors/image/image_settings.html#bpy-types-image-alpha-mode"),
 	("bpy.types.mask.frame_start*", "movie_clip/masking/sidebar.html#bpy-types-mask-frame-start"),
-	("bpy.types.mesh.*customdata*", "modeling/meshes/properties/custom_data.html#bpy-types-mesh-customdata"),
 	("bpy.types.multicamsequence*", "video_editing/sequencer/strips/effects/multicam.html#bpy-types-multicamsequence"),
 	("bpy.types.multiplysequence*", "video_editing/sequencer/strips/effects/multiply.html#bpy-types-multiplysequence"),
 	("bpy.types.multiresmodifier*", "modeling/modifiers/generate/multiresolution.html#bpy-types-multiresmodifier"),
@@ -1903,7 +1958,7 @@ url_manual_mapping = (
 	("bpy.types.shaderfxcolorize*", "grease_pencil/visual_effects/colorize.html#bpy-types-shaderfxcolorize"),
 	("bpy.types.shaderfxpixelate*", "grease_pencil/visual_effects/pixelate.html#bpy-types-shaderfxpixelate"),
 	("bpy.types.shadernodeinvert*", "render/shader_nodes/color/invert.html#bpy-types-shadernodeinvert"),
-	("bpy.types.shadernodemixrgb*", "render/shader_nodes/color/mix.html#bpy-types-shadernodemixrgb"),
+	("bpy.types.shadernodemixrgb*", "modeling/geometry_nodes/color/mix_rgb.html#bpy-types-shadernodemixrgb"),
 	("bpy.types.shadernodenormal*", "render/shader_nodes/vector/normal.html#bpy-types-shadernodenormal"),
 	("bpy.types.shadernodescript*", "render/shader_nodes/osl.html#bpy-types-shadernodescript"),
 	("bpy.types.shadernodetexies*", "render/shader_nodes/textures/ies.html#bpy-types-shadernodetexies"),
@@ -1962,6 +2017,7 @@ url_manual_mapping = (
 	("bpy.ops.preferences.addon*", "editors/preferences/addons.html#bpy-ops-preferences-addon"),
 	("bpy.ops.scene.light_cache*", "render/eevee/render_settings/indirect_lighting.html#bpy-ops-scene-light-cache"),
 	("bpy.ops.screen.area_dupli*", "interface/window_system/areas.html#bpy-ops-screen-area-dupli"),
+	("bpy.ops.screen.screenshot*", "interface/window_system/topbar.html#bpy-ops-screen-screenshot"),
 	("bpy.ops.sculpt.dirty_mask*", "sculpt_paint/sculpting/editing/mask.html#bpy-ops-sculpt-dirty-mask"),
 	("bpy.ops.sculpt.symmetrize*", "sculpt_paint/sculpting/tool_settings/symmetry.html#bpy-ops-sculpt-symmetrize"),
 	("bpy.ops.uv.remove_doubles*", "modeling/meshes/uv/editing.html#bpy-ops-uv-remove-doubles"),
@@ -2055,6 +2111,7 @@ url_manual_mapping = (
 	("bpy.types.curvesmodifier*", "video_editing/sequencer/sidebar/modifiers.html#bpy-types-curvesmodifier"),
 	("bpy.types.ffmpegsettings*", "render/output/properties/output.html#bpy-types-ffmpegsettings"),
 	("bpy.types.fmodifiernoise*", "editors/graph_editor/fcurves/modifiers.html#bpy-types-fmodifiernoise"),
+	("bpy.types.image.filepath*", "editors/image/image_settings.html#bpy-types-image-filepath"),
 	("bpy.types.keyframe.co_ui*", "editors/graph_editor/fcurves/properties.html#bpy-types-keyframe-co-ui"),
 	("bpy.types.mask.frame_end*", "movie_clip/masking/sidebar.html#bpy-types-mask-frame-end"),
 	("bpy.types.material.paint*", "sculpt_paint/texture_paint/index.html#bpy-types-material-paint"),
@@ -2100,7 +2157,6 @@ url_manual_mapping = (
 	("bpy.ops.mask.parent_set*", "movie_clip/masking/editing.html#bpy-ops-mask-parent-set"),
 	("bpy.ops.mask.select_all*", "movie_clip/masking/selecting.html#bpy-ops-mask-select-all"),
 	("bpy.ops.mask.select_box*", "movie_clip/masking/selecting.html#bpy-ops-mask-select-box"),
-	("bpy.ops.mesh.customdata*", "modeling/meshes/properties/custom_data.html#bpy-ops-mesh-customdata"),
 	("bpy.ops.mesh.edge_split*", "modeling/meshes/editing/mesh/split.html#bpy-ops-mesh-edge-split"),
 	("bpy.ops.mesh.fill_holes*", "modeling/meshes/editing/mesh/cleanup.html#bpy-ops-mesh-fill-holes"),
 	("bpy.ops.mesh.mark_sharp*", "modeling/meshes/editing/edge/edge_data.html#bpy-ops-mesh-mark-sharp"),
@@ -2136,7 +2192,6 @@ url_manual_mapping = (
 	("bpy.types.arraymodifier*", "modeling/modifiers/generate/array.html#bpy-types-arraymodifier"),
 	("bpy.types.bevelmodifier*", "modeling/modifiers/generate/bevel.html#bpy-types-bevelmodifier"),
 	("bpy.types.buildmodifier*", "modeling/modifiers/generate/build.html#bpy-types-buildmodifier"),
-	("bpy.types.camera.sensor*", "render/cameras.html#bpy-types-camera-sensor"),
 	("bpy.types.clothmodifier*", "physics/cloth/index.html#bpy-types-clothmodifier"),
 	("bpy.types.clothsettings*", "physics/cloth/settings/index.html#bpy-types-clothsettings"),
 	("bpy.types.cloudstexture*", "render/materials/legacy_textures/types/clouds.html#bpy-types-cloudstexture"),
@@ -2228,6 +2283,7 @@ url_manual_mapping = (
 	("bpy.types.glowsequence*", "video_editing/sequencer/strips/effects/glow.html#bpy-types-glowsequence"),
 	("bpy.types.gpencillayer*", "grease_pencil/properties/layers.html#bpy-types-gpencillayer"),
 	("bpy.types.hookmodifier*", "modeling/modifiers/deform/hooks.html#bpy-types-hookmodifier"),
+	("bpy.types.image.source*", "editors/image/image_settings.html#bpy-types-image-source"),
 	("bpy.types.imagetexture*", "render/materials/legacy_textures/types/image_movie.html#bpy-types-imagetexture"),
 	("bpy.types.latticepoint*", "animation/lattice.html#bpy-types-latticepoint"),
 	("bpy.types.magictexture*", "render/materials/legacy_textures/types/magic.html#bpy-types-magictexture"),

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

@@ -141,10 +141,11 @@ class Params:
             # Use the "cursor" functionality for RMB select.
             if use_alt_tool_or_cursor:
                 self.cursor_set_event = {"type": 'LEFTMOUSE', "value": 'PRESS', "alt": True}
+                self.cursor_tweak_event = {"type": 'EVT_TWEAK_L', "value": 'ANY', "alt": True}
             else:
                 self.cursor_set_event = {"type": 'LEFTMOUSE', "value": 'CLICK'}
+                self.cursor_tweak_event = None
 
-            self.cursor_tweak_event = None
             self.use_fallback_tool = use_fallback_tool
             self.tool_modifier = {}
         else:
@@ -7623,7 +7624,7 @@ def km_sequencer_editor_tool_move(params):
         "Sequencer Tool: Move",
         {"space_type": 'SEQUENCE_EDITOR', "region_type": 'WINDOW'},
         {"items": [
-            ("transform.translate", params.tool_maybe_tweak_event,
+            ("transform.translate", {**params.tool_maybe_tweak_event, **params.tool_modifier},
              {"properties": [("release_confirm", True)]}),
         ]},
     )
@@ -7634,7 +7635,7 @@ def km_sequencer_editor_tool_rotate(params):
         "Sequencer Tool: Rotate",
         {"space_type": 'SEQUENCE_EDITOR', "region_type": 'WINDOW'},
         {"items": [
-            ("transform.rotate", params.tool_maybe_tweak_event,
+            ("transform.rotate", {**params.tool_maybe_tweak_event, **params.tool_modifier},
              {"properties": [("release_confirm", True)]}),
         ]},
     )
@@ -7645,7 +7646,7 @@ def km_sequencer_editor_tool_scale(params):
         "Sequencer Tool: Scale",
         {"space_type": 'SEQUENCE_EDITOR', "region_type": 'WINDOW'},
         {"items": [
-            ("transform.resize", params.tool_maybe_tweak_event,
+            ("transform.resize", {**params.tool_maybe_tweak_event, **params.tool_modifier},
              {"properties": [("release_confirm", True)]}),
         ]},
     )

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

@@ -4203,20 +4203,29 @@ def keymap_transform_tool_mmb(keymap):
             km_items_new = []
             for kmi in km_items:
                 ty = kmi[1]["type"]
-                if ty == 'LEFTMOUSE':
-                    kmi = (kmi[0], kmi[1].copy(), kmi[2])
-                    kmi[1]["type"] = 'MIDDLEMOUSE'
-                    km_items_new.append(kmi)
-                elif ty == 'EVT_TWEAK_L':
-                    kmi = (kmi[0], kmi[1].copy(), kmi[2])
-                    if kmi[1]["value"] == 'ANY':
+                if km_name.endswith(" (fallback)"):
+                    if ty == 'RIGHTMOUSE':
+                        kmi = (kmi[0], kmi[1].copy(), kmi[2])
+                        kmi[1]["type"] = 'LEFTMOUSE'
+                        km_items_new.append(kmi)
+                    elif ty == 'EVT_TWEAK_R':
+                        kmi = (kmi[0], kmi[1].copy(), kmi[2])
+                        kmi[1]["type"] = 'EVT_TWEAK_L'
+                        km_items_new.append(kmi)
+                else:
+                    if ty == 'LEFTMOUSE':
+                        kmi = (kmi[0], kmi[1].copy(), kmi[2])
                         kmi[1]["type"] = 'MIDDLEMOUSE'
-                        kmi[1]["value"] = 'PRESS'
-                    else:
-                        # Directional tweaking can't be replaced by middle-mouse.
-                        kmi[1]["type"] = 'EVT_TWEAK_M'
-
-                    km_items_new.append(kmi)
+                        km_items_new.append(kmi)
+                    elif ty == 'EVT_TWEAK_L':
+                        kmi = (kmi[0], kmi[1].copy(), kmi[2])
+                        if kmi[1]["value"] == 'ANY':
+                            kmi[1]["type"] = 'MIDDLEMOUSE'
+                            kmi[1]["value"] = 'PRESS'
+                        else:
+                            # Directional tweaking can't be replaced by middle-mouse.
+                            kmi[1]["type"] = 'EVT_TWEAK_M'
+                        km_items_new.append(kmi)
             km_items.extend(km_items_new)
 
 
@@ -4227,9 +4236,17 @@ def generate_keymaps(params=None):
 
     # Combine the key-map to support manipulating it, so we don't need to manually
     # define key-map here just to manipulate them.
-    blender_default = execfile(
+    blender_default_mod = execfile(
         os.path.join(os.path.dirname(__file__), "blender_default.py"),
-    ).generate_keymaps()
+    )
+
+    blender_default = blender_default_mod.generate_keymaps(
+        # Use the default key-map with only minor changes to default arguments.
+        blender_default_mod.Params(
+            # Needed so the fallback key-map items are populated.
+            use_fallback_tool=True,
+        ),
+    )
 
     keymap_existing_names = {km[0] for km in keymap}
     keymap.extend([km for km in blender_default if km[0] not in keymap_existing_names])

release/scripts/presets/tracking_camera/1_inch.py

@@ -1,4 +1,6 @@
 import bpy
-bpy.context.camera.sensor_width = 13.2
-bpy.context.camera.sensor_height = 8.80
-bpy.context.camera.sensor_fit = 'HORIZONTAL'
+camera = bpy.context.edit_movieclip.tracking.camera
+
+camera.sensor_width = 13.2
+camera.units = 'MILLIMETERS'
+camera.pixel_aspect = 1

release/scripts/presets/tracking_camera/1_slash_1.8_inch.py

@@ -1,4 +1,6 @@
 import bpy
-bpy.context.camera.sensor_width = 7.18
-bpy.context.camera.sensor_height = 5.32
-bpy.context.camera.sensor_fit = 'HORIZONTAL'
+camera = bpy.context.edit_movieclip.tracking.camera
+
+camera.sensor_width = 7.18
+camera.units = 'MILLIMETERS'
+camera.pixel_aspect = 1

release/scripts/presets/tracking_camera/1_slash_2.3_inch.py

@@ -1,4 +1,6 @@
 import bpy
-bpy.context.camera.sensor_width = 6.17
-bpy.context.camera.sensor_height = 4.55
-bpy.context.camera.sensor_fit = 'HORIZONTAL'
+camera = bpy.context.edit_movieclip.tracking.camera
+
+camera.sensor_width = 6.17
+camera.units = 'MILLIMETERS'
+camera.pixel_aspect = 1