Cleanup: Convert interface_region_search.cc to C++

This commit adds modifier error messages to some of the cases
where the node group is configured improperly. It also clears the
geometry set when there is an error with the node group. This is
consistent to what we do in nodes themselves, and feels more
intuitive than passing the input geometry through the node group
silently.

Fixes T87142

intern/cycles/blender/addon/properties.py

@@ -1419,10 +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 RDNA2 architecture", icon='BLANK1')
-                # TODO: provide driver version info.
-                #if sys.platform[:3] == "win":
-                #    col.label(text="and AMD driver version ??? or newer", icon='BLANK1')
+                col.label(text="Requires discrete AMD GPU with RDNA architecture", icon='BLANK1')
+                if sys.platform[:3] == "win":
+                    col.label(text="and AMD Radeon Pro 21.Q4 driver or newer", icon='BLANK1')
             return
 
         for device in devices:

intern/cycles/blender/object.cpp

@@ -62,15 +62,15 @@ bool BlenderSync::BKE_object_is_modified(BL::Object &b_ob)
   return false;
 }
 
-bool BlenderSync::object_is_geometry(BL::Object &b_ob)
+bool BlenderSync::object_is_geometry(BObjectInfo &b_ob_info)
 {
-  BL::ID b_ob_data = b_ob.data();
+  BL::ID b_ob_data = b_ob_info.object_data;
 
   if (!b_ob_data) {
     return false;
   }
 
-  BL::Object::type_enum type = b_ob.type();
+  BL::Object::type_enum type = b_ob_info.iter_object.type();
 
   if (type == BL::Object::type_VOLUME || type == BL::Object::type_HAIR) {
     /* Will be exported attached to mesh. */
@@ -87,6 +87,24 @@ bool BlenderSync::object_is_geometry(BL::Object &b_ob)
   return b_ob_data.is_a(&RNA_Mesh);
 }
 
+bool BlenderSync::object_can_have_geometry(BL::Object &b_ob)
+{
+  BL::Object::type_enum type = b_ob.type();
+  switch (type) {
+    case BL::Object::type_MESH:
+    case BL::Object::type_CURVE:
+    case BL::Object::type_SURFACE:
+    case BL::Object::type_META:
+    case BL::Object::type_FONT:
+    case BL::Object::type_HAIR:
+    case BL::Object::type_POINTCLOUD:
+    case BL::Object::type_VOLUME:
+      return true;
+    default:
+      return false;
+  }
+}
+
 bool BlenderSync::object_is_light(BL::Object &b_ob)
 {
   BL::ID b_ob_data = b_ob.data();
@@ -189,7 +207,7 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph,
   }
 
   /* only interested in object that we can create meshes from */
-  if (!object_is_geometry(b_ob)) {
+  if (!object_is_geometry(b_ob_info)) {
     return NULL;
   }
 

intern/cycles/blender/sync.cpp

@@ -162,19 +162,19 @@ void BlenderSync::sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d
     /* Object */
     else if (b_id.is_a(&RNA_Object)) {
       BL::Object b_ob(b_id);
-      const bool is_geometry = object_is_geometry(b_ob);
-      const bool is_light = !is_geometry && object_is_light(b_ob);
+      const bool can_have_geometry = object_can_have_geometry(b_ob);
+      const bool is_light = !can_have_geometry && object_is_light(b_ob);
 
       if (b_ob.is_instancer() && b_update.is_updated_shading()) {
         /* Needed for e.g. object color updates on instancer. */
         object_map.set_recalc(b_ob);
       }
 
-      if (is_geometry || is_light) {
+      if (can_have_geometry || is_light) {
         const bool updated_geometry = b_update.is_updated_geometry();
 
         /* Geometry (mesh, hair, volume). */
-        if (is_geometry) {
+        if (can_have_geometry) {
           if (b_update.is_updated_transform() || b_update.is_updated_shading()) {
             object_map.set_recalc(b_ob);
           }

intern/cycles/blender/sync.h

@@ -208,7 +208,8 @@ class BlenderSync {
   /* util */
   void find_shader(BL::ID &id, array<Node *> &used_shaders, Shader *default_shader);
   bool BKE_object_is_modified(BL::Object &b_ob);
-  bool object_is_geometry(BL::Object &b_ob);
+  bool object_is_geometry(BObjectInfo &b_ob_info);
+  bool object_can_have_geometry(BL::Object &b_ob);
   bool object_is_light(BL::Object &b_ob);
 
   /* variables */

intern/cycles/device/hip/device_impl.cpp

@@ -154,7 +154,7 @@ bool HIPDevice::support_device(const uint /*kernel_features*/)
     hipDeviceProp_t props;
     hipGetDeviceProperties(&props, hipDevId);
 
-    set_error(string_printf("HIP backend requires AMD RDNA2 graphics card or up, but found %s.",
+    set_error(string_printf("HIP backend requires AMD RDNA graphics card or up, but found %s.",
                             props.name));
     return false;
   }

intern/cycles/device/hip/util.h

@@ -64,7 +64,7 @@ static inline bool hipSupportsDevice(const int hipDevId)
   hipDeviceGetAttribute(&major, hipDeviceAttributeComputeCapabilityMajor, hipDevId);
   hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
 
-  return (major > 10) || (major == 10 && minor >= 3);
+  return (major > 10) || (major == 10 && minor >= 1);
 }
 
 CCL_NAMESPACE_END

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),
@@ -133,6 +125,11 @@ OptiXDevice::~OptiXDevice()
     }
   }
 
+  /* Make sure denoiser is destroyed before device context! */
+  if (denoiser_.optix_denoiser != nullptr) {
+    optixDenoiserDestroy(denoiser_.optix_denoiser);
+  }
+
   optixDeviceContextDestroy(context);
 }
 
@@ -884,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;
@@ -939,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;
@@ -971,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/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
@@ -723,12 +734,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 +753,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 +786,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/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

@@ -92,12 +92,29 @@
 
 /* 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) \
   extern "C" __global__ void __launch_bounds__(block_num_threads, \
                                                GPU_MULTIPRESSOR_MAX_REGISTERS / \
                                                    (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_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
+
 /* 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

@@ -35,12 +35,29 @@
 
 /* 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) \
   extern "C" __global__ void __launch_bounds__(block_num_threads, \
                                                GPU_MULTIPRESSOR_MAX_REGISTERS / \
                                                    (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_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
+
 /* sanity checks */
 
 #if GPU_KERNEL_BLOCK_NUM_THREADS > GPU_BLOCK_MAX_THREADS

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

@@ -58,6 +58,96 @@ 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(block_num_threads, thread_num_registers)
+#define ccl_gpu_kernel_threads(block_num_threads)
+
+/* 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)
+
+// clang-format on
+
 /* make_type definitions with Metal style element initializers */
 #ifdef make_float2
 #  undef make_float2
@@ -124,3 +214,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/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);
     }
   }
 }

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__ */
 

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__

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;
   }
 }
 

release/scripts/modules/bl_i18n_utils/settings.py

@@ -264,7 +264,7 @@ PYGETTEXT_KEYWORDS = (() +
           for it in ("BMO_error_raise",)) +
 
     tuple(("{}\\((?:[^\"',]+,)\\s*" + _msg_re + r"\s*(?:\)|,)").format(it)
-          for it in ("modifier_setError",)) +
+          for it in ("BKE_modifier_set_error",)) +
 
     tuple((r"{}\(\s*" + _msg_re + r"\s*,\s*(?:" +
            r"\s*,\s*)?(?:".join(_ctxt_re_gen(i) for i in range(PYGETTEXT_MAX_MULTI_CTXT)) + r")?\s*\)").format(it)

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/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

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

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

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

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

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

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

release/scripts/presets/tracking_camera/2_slash_3_inch.py

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

release/scripts/presets/tracking_camera/APS-C.py

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

release/scripts/presets/tracking_camera/APS-C_(Canon).py

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

release/scripts/presets/tracking_camera/APS-H_(Canon).py

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

release/scripts/presets/tracking_camera/Analog_16mm.py

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

release/scripts/presets/tracking_camera/Analog_35mm.py

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

release/scripts/presets/tracking_camera/Analog_65mm.py

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

release/scripts/presets/tracking_camera/Analog_IMAX.py

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

release/scripts/presets/tracking_camera/Analog_Super_16.py

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

release/scripts/presets/tracking_camera/Analog_Super_35.py

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

release/scripts/presets/tracking_camera/Arri_Alexa_65.py

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

release/scripts/presets/tracking_camera/Arri_Alexa_LF.py

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

release/scripts/presets/tracking_camera/Arri_Alexa_Mini_&_SXT.py

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

release/scripts/presets/tracking_camera/Blackmagic_Pocket_&_Studio.py

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

release/scripts/presets/tracking_camera/Blackmagic_Pocket_4K.py

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

release/scripts/presets/tracking_camera/Blackmagic_Pocket_6k.py

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

release/scripts/presets/tracking_camera/Blackmagic_URSA_4.6K.py

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

release/scripts/presets/tracking_camera/Foveon_(Sigma).py

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

release/scripts/presets/tracking_camera/Fullframe.py

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

release/scripts/presets/tracking_camera/MFT.py

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

release/scripts/presets/tracking_camera/Medium-format_(Hasselblad).py

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

release/scripts/presets/tracking_camera/RED_Dragon_5K.py

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

release/scripts/presets/tracking_camera/RED_Dragon_6K.py

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

release/scripts/presets/tracking_camera/RED_Helium_8K.py

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

release/scripts/presets/tracking_camera/RED_Monstro_8K.py

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

release/scripts/startup/bl_ui/properties_freestyle.py

@@ -526,7 +526,7 @@ class VIEWLAYER_PT_freestyle_linestyle_strokes(ViewLayerFreestyleLineStyle, Pane
         row = layout.row(align=True)
         row.alignment = 'LEFT'
         row.label(text=lineset.name, icon='LINE_DATA')
-        row.label(text="", icon='SMALL_TRI_RIGHT_VEC')
+        row.label(text="", icon='RIGHTARROW')
         row.label(text=linestyle.name)
 
         col = layout.column(align=True)
@@ -830,7 +830,7 @@ class VIEWLAYER_PT_freestyle_linestyle_color(ViewLayerFreestyleLineStyle, Panel)
         row = layout.row(align=True)
         row.alignment = 'LEFT'
         row.label(text=lineset.name, icon='LINE_DATA')
-        row.label(text="", icon='SMALL_TRI_RIGHT_VEC')
+        row.label(text="", icon='RIGHTARROW')
         row.label(text=linestyle.name)
 
         col = layout.column()
@@ -922,7 +922,7 @@ class VIEWLAYER_PT_freestyle_linestyle_alpha(ViewLayerFreestyleLineStyle, Panel)
         row = layout.row(align=True)
         row.alignment = 'LEFT'
         row.label(text=lineset.name, icon='LINE_DATA')
-        row.label(text="", icon='SMALL_TRI_RIGHT_VEC')
+        row.label(text="", icon='RIGHTARROW')
         row.label(text=linestyle.name)
 
         col = layout.column()
@@ -1036,7 +1036,7 @@ class VIEWLAYER_PT_freestyle_linestyle_thickness(ViewLayerFreestyleLineStyle, Pa
         row = layout.row(align=True)
         row.alignment = 'LEFT'
         row.label(text=lineset.name, icon='LINE_DATA')
-        row.label(text="", icon='SMALL_TRI_RIGHT_VEC')
+        row.label(text="", icon='RIGHTARROW')
         row.label(text=linestyle.name)
 
         col = layout.column()
@@ -1182,7 +1182,7 @@ class VIEWLAYER_PT_freestyle_linestyle_geometry(ViewLayerFreestyleLineStyle, Pan
         row = layout.row(align=True)
         row.alignment = 'LEFT'
         row.label(text=lineset.name, icon='LINE_DATA')
-        row.label(text="", icon='SMALL_TRI_RIGHT_VEC')
+        row.label(text="", icon='RIGHTARROW')
         row.label(text=linestyle.name)
 
         col = layout.column()
@@ -1215,7 +1215,7 @@ class VIEWLAYER_PT_freestyle_linestyle_texture(ViewLayerFreestyleLineStyle, Pane
         row = layout.row(align=True)
         row.alignment = 'LEFT'
         row.label(text=lineset.name, icon='LINE_DATA')
-        row.label(text="", icon='SMALL_TRI_RIGHT_VEC')
+        row.label(text="", icon='RIGHTARROW')
         row.label(text=linestyle.name)
 
         layout.prop(linestyle, "use_nodes")

release/scripts/startup/bl_ui/space_sequencer.py

@@ -254,7 +254,7 @@ class SEQUENCER_MT_editor_menus(Menu):
 class SEQUENCER_PT_gizmo_display(Panel):
     bl_space_type = 'SEQUENCE_EDITOR'
     bl_region_type = 'HEADER'
-    bl_label = "Gizmo"
+    bl_label = "Gizmos"
     bl_ui_units_x = 8
 
     def draw(self, context):

release/scripts/startup/bl_ui/space_toolsystem_toolbar.py

@@ -3001,7 +3001,11 @@ class VIEW3D_PT_tools_active(ToolSelectPanelHelper, Panel):
             ),
             None,
             lambda context: (
-                (_defs_view3d_generic.cursor,)
+                (
+                    _defs_view3d_generic.cursor,
+                    None,
+                    *VIEW3D_PT_tools_active._tools_transform,
+                )
                 if context is None or context.pose_object
                 else ()
             ),

release/scripts/startup/bl_ui/space_view3d.py

@@ -5990,7 +5990,7 @@ class VIEW3D_PT_shading_render_pass(Panel):
 class VIEW3D_PT_gizmo_display(Panel):
     bl_space_type = 'VIEW_3D'
     bl_region_type = 'HEADER'
-    bl_label = "Gizmo"
+    bl_label = "Gizmos"
     bl_ui_units_x = 8
 
     def draw(self, context):

source/blender/blenkernel/BKE_asset_catalog.hh

@@ -429,7 +429,9 @@ class AssetCatalog {
    * Simple, human-readable name for the asset catalog. This is stored on assets alongside the
    * catalog ID; the catalog ID is a UUID that is not human-readable,
    * so to avoid complete data-loss when the catalog definition file gets lost,
-   * we also store a human-readable simple name for the catalog. */
+   * we also store a human-readable simple name for the catalog.
+   *
+   * It should fit in sizeof(AssetMetaData::catalog_simple_name) bytes. */
   std::string simple_name;
 
   struct Flags {

source/blender/blenkernel/BKE_layer.h

@@ -328,7 +328,7 @@ void BKE_view_layer_visible_bases_iterator_end(BLI_Iterator *iter);
   { \
     Object *_instance; \
     Base *_base; \
-    for (_base = (view_layer)->object_bases.first; _base; _base = _base->next) { \
+    for (_base = (Base *)(view_layer)->object_bases.first; _base; _base = _base->next) { \
       _instance = _base->object;
 
 #define FOREACH_OBJECT_END \

source/blender/blenkernel/BKE_mesh.h

@@ -209,6 +209,7 @@ struct Mesh *BKE_mesh_create_derived_for_modifier(struct Depsgraph *depsgraph,
                                                   struct Scene *scene,
                                                   struct Object *ob_eval,
                                                   struct ModifierData *md_eval,
+                                                  const bool use_virtual_modifiers,
                                                   const bool build_shapekey_layers);
 
 /* Copies a nomain-Mesh into an existing Mesh. */

source/blender/blenkernel/BKE_mesh_runtime.h

@@ -41,7 +41,8 @@ struct Mesh;
 struct Object;
 struct Scene;
 
-void BKE_mesh_runtime_reset(struct Mesh *mesh);
+void BKE_mesh_runtime_init_data(struct Mesh *mesh);
+void BKE_mesh_runtime_free_data(struct Mesh *mesh);
 void BKE_mesh_runtime_reset_on_copy(struct Mesh *mesh, const int flag);
 int BKE_mesh_runtime_looptri_len(const struct Mesh *mesh);
 void BKE_mesh_runtime_looptri_recalc(struct Mesh *mesh);

source/blender/blenkernel/intern/asset_catalog_test.cc

@@ -24,6 +24,7 @@
 #include "BLI_fileops.h"
 #include "BLI_path_util.h"
 
+#include "DNA_asset_types.h"
 #include "DNA_userdef_types.h"
 
 #include "testing/testing.h"
@@ -930,6 +931,28 @@ TEST_F(AssetCatalogTest, update_catalog_path_simple_name)
       << "Changing the path should update the simplename of children.";
 }
 
+TEST_F(AssetCatalogTest, update_catalog_path_longer_than_simplename)
+{
+  AssetCatalogService service(asset_library_root_);
+  service.load_from_disk(asset_library_root_ + "/" +
+                         AssetCatalogService::DEFAULT_CATALOG_FILENAME);
+  const std::string new_path =
+      "this/is/a/very/long/path/that/exceeds/the/simple-name/length/of/assets";
+  ASSERT_GT(new_path.length(), sizeof(AssetMetaData::catalog_simple_name))
+      << "This test case should work with paths longer than AssetMetaData::catalog_simple_name";
+
+  service.update_catalog_path(UUID_POSES_RUZENA, new_path);
+
+  const std::string new_simple_name = service.find_catalog(UUID_POSES_RUZENA)->simple_name;
+  EXPECT_LT(new_simple_name.length(), sizeof(AssetMetaData::catalog_simple_name))
+      << "The new simple name should fit in the asset metadata.";
+  EXPECT_EQ("...very-long-path-that-exceeds-the-simple-name-length-of-assets", new_simple_name)
+      << "Changing the path should update the simplename.";
+  EXPECT_EQ("...long-path-that-exceeds-the-simple-name-length-of-assets-face",
+            service.find_catalog(UUID_POSES_RUZENA_FACE)->simple_name)
+      << "Changing the path should update the simplename of children.";
+}
+
 TEST_F(AssetCatalogTest, update_catalog_path_add_slashes)
 {
   AssetCatalogService service(asset_library_root_);

source/blender/blenkernel/intern/image.c

@@ -1323,6 +1323,9 @@ void BKE_image_print_memlist(Main *bmain)
 
 static bool imagecache_check_dirty(ImBuf *ibuf, void *UNUSED(userkey), void *UNUSED(userdata))
 {
+  if (ibuf == NULL) {
+    return false;
+  }
   return (ibuf->userflags & IB_BITMAPDIRTY) == 0;
 }
 
@@ -1366,6 +1369,9 @@ void BKE_image_free_all_textures(Main *bmain)
 
 static bool imagecache_check_free_anim(ImBuf *ibuf, void *UNUSED(userkey), void *userdata)
 {
+  if (ibuf == NULL) {
+    return true;
+  }
   int except_frame = *(int *)userdata;
   return (ibuf->userflags & IB_BITMAPDIRTY) == 0 && (ibuf->index != IMA_NO_INDEX) &&
          (except_frame != IMA_INDEX_ENTRY(ibuf->index));

source/blender/blenkernel/intern/mesh.cc

@@ -88,7 +88,7 @@ static void mesh_init_data(ID *id)
   CustomData_reset(&mesh->pdata);
   CustomData_reset(&mesh->ldata);
 
-  BKE_mesh_runtime_reset(mesh);
+  BKE_mesh_runtime_init_data(mesh);
 
   mesh->face_sets_color_seed = BLI_hash_int(PIL_check_seconds_timer_i() & UINT_MAX);
 }
@@ -168,7 +168,7 @@ static void mesh_free_data(ID *id)
     mesh->edit_mesh = nullptr;
   }
 
-  BKE_mesh_runtime_clear_cache(mesh);
+  BKE_mesh_runtime_free_data(mesh);
   mesh_clear_geometry(mesh);
   MEM_SAFE_FREE(mesh->mat);
 }
@@ -308,7 +308,9 @@ static void mesh_blend_read_data(BlendDataReader *reader, ID *id)
 
   mesh->texflag &= ~ME_AUTOSPACE_EVALUATED;
   mesh->edit_mesh = nullptr;
-  BKE_mesh_runtime_reset(mesh);
+
+  memset(&mesh->runtime, 0, sizeof(mesh->runtime));
+  BKE_mesh_runtime_init_data(mesh);
 
   /* happens with old files */
   if (mesh->mselect == nullptr) {

source/blender/blenkernel/intern/mesh_convert.cc

@@ -1277,10 +1277,16 @@ static void add_shapekey_layers(Mesh *mesh_dest, Mesh *mesh_src)
   }
 }
 
+/**
+ * \param use_virtual_modifiers: When enabled calculate virtual-modifiers before applying `md_eval`
+ * support this since virtual-modifiers are not modifiers from a user perspective,
+ * allowing shape keys to be included with the modifier being applied, see: T91923.
+ */
 Mesh *BKE_mesh_create_derived_for_modifier(struct Depsgraph *depsgraph,
                                            Scene *scene,
                                            Object *ob_eval,
                                            ModifierData *md_eval,
+                                           const bool use_virtual_modifiers,
                                            const bool build_shapekey_layers)
 {
   Mesh *me = ob_eval->runtime.data_orig ? (Mesh *)ob_eval->runtime.data_orig :
@@ -1303,22 +1309,49 @@ Mesh *BKE_mesh_create_derived_for_modifier(struct Depsgraph *depsgraph,
     BKE_keyblock_convert_to_mesh(kb, me);
   }
 
-  if (mti->type == eModifierTypeType_OnlyDeform) {
-    int numVerts;
-    float(*deformedVerts)[3] = BKE_mesh_vert_coords_alloc(me, &numVerts);
+  Mesh *mesh_temp = (Mesh *)BKE_id_copy_ex(nullptr, &me->id, nullptr, LIB_ID_COPY_LOCALIZE);
+  int numVerts = 0;
+  float(*deformedVerts)[3] = nullptr;
 
-    result = (Mesh *)BKE_id_copy_ex(nullptr, &me->id, nullptr, LIB_ID_COPY_LOCALIZE);
+  if (use_virtual_modifiers) {
+    VirtualModifierData virtualModifierData;
+    for (ModifierData *md_eval_virt =
+             BKE_modifiers_get_virtual_modifierlist(ob_eval, &virtualModifierData);
+         md_eval_virt && (md_eval_virt != ob_eval->modifiers.first);
+         md_eval_virt = md_eval_virt->next) {
+      if (!BKE_modifier_is_enabled(scene, md_eval_virt, eModifierMode_Realtime)) {
+        continue;
+      }
+      /* All virtual modifiers are deform modifiers. */
+      const ModifierTypeInfo *mti_virt = BKE_modifier_get_info((ModifierType)md_eval_virt->type);
+      BLI_assert(mti_virt->type == eModifierTypeType_OnlyDeform);
+      if (mti_virt->type != eModifierTypeType_OnlyDeform) {
+        continue;
+      }
+
+      if (deformedVerts == nullptr) {
+        deformedVerts = BKE_mesh_vert_coords_alloc(me, &numVerts);
+      }
+      mti_virt->deformVerts(md_eval_virt, &mectx, mesh_temp, deformedVerts, numVerts);
+    }
+  }
+
+  if (mti->type == eModifierTypeType_OnlyDeform) {
+    if (deformedVerts == nullptr) {
+      deformedVerts = BKE_mesh_vert_coords_alloc(me, &numVerts);
+    }
+    result = mesh_temp;
     mti->deformVerts(md_eval, &mectx, result, deformedVerts, numVerts);
     BKE_mesh_vert_coords_apply(result, deformedVerts);
 
     if (build_shapekey_layers) {
       add_shapekey_layers(result, me);
     }
-
-    MEM_freeN(deformedVerts);
   }
   else {
-    Mesh *mesh_temp = (Mesh *)BKE_id_copy_ex(nullptr, &me->id, nullptr, LIB_ID_COPY_LOCALIZE);
+    if (deformedVerts != nullptr) {
+      BKE_mesh_vert_coords_apply(mesh_temp, deformedVerts);
+    }
 
     if (build_shapekey_layers) {
       add_shapekey_layers(mesh_temp, me);
@@ -1332,6 +1365,10 @@ Mesh *BKE_mesh_create_derived_for_modifier(struct Depsgraph *depsgraph,
     }
   }
 
+  if (deformedVerts != nullptr) {
+    MEM_freeN(deformedVerts);
+  }
+
   return result;
 }
 

source/blender/blenkernel/intern/mesh_runtime.c

@@ -45,17 +45,54 @@
  * \{ */
 
 /**
- * Default values defined at read time.
+ * \brief Initialize the runtime mutexes of the given mesh.
+ *
+ * Any existing mutexes will be overridden.
  */
-void BKE_mesh_runtime_reset(Mesh *mesh)
+static void mesh_runtime_init_mutexes(Mesh *mesh)
 {
-  memset(&mesh->runtime, 0, sizeof(mesh->runtime));
   mesh->runtime.eval_mutex = MEM_mallocN(sizeof(ThreadMutex), "mesh runtime eval_mutex");
   BLI_mutex_init(mesh->runtime.eval_mutex);
   mesh->runtime.render_mutex = MEM_mallocN(sizeof(ThreadMutex), "mesh runtime render_mutex");
   BLI_mutex_init(mesh->runtime.render_mutex);
 }
 
+/**
+ * \brief free the mutexes of the given mesh runtime.
+ */
+static void mesh_runtime_free_mutexes(Mesh *mesh)
+{
+  if (mesh->runtime.eval_mutex != NULL) {
+    BLI_mutex_end(mesh->runtime.eval_mutex);
+    MEM_freeN(mesh->runtime.eval_mutex);
+    mesh->runtime.eval_mutex = NULL;
+  }
+  if (mesh->runtime.render_mutex != NULL) {
+    BLI_mutex_end(mesh->runtime.render_mutex);
+    MEM_freeN(mesh->runtime.render_mutex);
+    mesh->runtime.render_mutex = NULL;
+  }
+}
+
+/**
+ * \brief Initialize the runtime of the given mesh.
+ *
+ * Function expects that the runtime is already cleared.
+ */
+void BKE_mesh_runtime_init_data(Mesh *mesh)
+{
+  mesh_runtime_init_mutexes(mesh);
+}
+
+/**
+ * \brief Free all data (and mutexes) inside the runtime of the given mesh.
+ */
+void BKE_mesh_runtime_free_data(Mesh *mesh)
+{
+  BKE_mesh_runtime_clear_cache(mesh);
+  mesh_runtime_free_mutexes(mesh);
+}
+
 /* Clear all pointers which we don't want to be shared on copying the datablock.
  * However, keep all the flags which defines what the mesh is (for example, that
  * it's deformed only, or that its custom data layers are out of date.) */
@@ -71,25 +108,16 @@ void BKE_mesh_runtime_reset_on_copy(Mesh *mesh, const int UNUSED(flag))
   runtime->bvh_cache = NULL;
   runtime->shrinkwrap_data = NULL;
 
-  mesh->runtime.eval_mutex = MEM_mallocN(sizeof(ThreadMutex), "mesh runtime eval_mutex");
-  BLI_mutex_init(mesh->runtime.eval_mutex);
-
-  mesh->runtime.render_mutex = MEM_mallocN(sizeof(ThreadMutex), "mesh runtime render_mutex");
-  BLI_mutex_init(mesh->runtime.render_mutex);
+  mesh_runtime_init_mutexes(mesh);
 }
 
+/**
+ * \brief This function clears runtime cache of the given mesh.
+ * 
+ * Call this function to recalculate runtime data when used.
+ */
 void BKE_mesh_runtime_clear_cache(Mesh *mesh)
 {
-  if (mesh->runtime.eval_mutex != NULL) {
-    BLI_mutex_end(mesh->runtime.eval_mutex);
-    MEM_freeN(mesh->runtime.eval_mutex);
-    mesh->runtime.eval_mutex = NULL;
-  }
-  if (mesh->runtime.render_mutex != NULL) {
-    BLI_mutex_end(mesh->runtime.render_mutex);
-    MEM_freeN(mesh->runtime.render_mutex);
-    mesh->runtime.render_mutex = NULL;
-  }
   if (mesh->runtime.mesh_eval != NULL) {
     mesh->runtime.mesh_eval->edit_mesh = NULL;
     BKE_id_free(NULL, mesh->runtime.mesh_eval);

source/blender/blenkernel/intern/object.cc

@@ -1223,7 +1223,7 @@ static IDProperty *object_asset_dimensions_property(Object *ob)
     return nullptr;
   }
 
-  IDPropertyTemplate idprop = {0};
+  IDPropertyTemplate idprop{};
   idprop.array.len = ARRAY_SIZE(dimensions);
   idprop.array.type = IDP_FLOAT;
 
@@ -2635,7 +2635,7 @@ Object **BKE_object_pose_array_get_ex(ViewLayer *view_layer,
   Object *ob_pose = BKE_object_pose_armature_get(ob_active);
   Object **objects = nullptr;
   if (ob_pose == ob_active) {
-    ObjectsInModeParams ob_params = {0};
+    ObjectsInModeParams ob_params{};
     ob_params.object_mode = OB_MODE_POSE;
     ob_params.no_dup_data = unique;
 
@@ -2682,7 +2682,7 @@ Base **BKE_object_pose_base_array_get_ex(ViewLayer *view_layer,
   }
 
   if (base_active && (base_pose == base_active)) {
-    ObjectsInModeParams ob_params = {0};
+    ObjectsInModeParams ob_params{};
     ob_params.object_mode = OB_MODE_POSE;
     ob_params.no_dup_data = unique;
 
@@ -4304,7 +4304,7 @@ void BKE_object_foreach_display_point(Object *ob,
     }
   }
   else if (ob->type == OB_GPENCIL) {
-    GPencilStrokePointIterData iter_data = {nullptr};
+    GPencilStrokePointIterData iter_data{};
     iter_data.obmat = obmat;
     iter_data.point_func_cb = func_cb;
     iter_data.user_data = user_data;

source/blender/blenkernel/intern/object_dupli.cc

@@ -166,6 +166,10 @@ static void copy_dupli_context(
   r_ctx->persistent_id[r_ctx->level] = index;
   ++r_ctx->level;
 
+  if (r_ctx->level == MAX_DUPLI_RECUR - 1) {
+    std::cerr << "Warning: Maximum instance recursion level reached.\n";
+  }
+
   r_ctx->gen = get_dupli_generator(r_ctx);
 }
 

source/blender/blenloader/intern/versioning_300.c

@@ -645,6 +645,10 @@ void do_versions_after_linking_300(Main *bmain, ReportList *UNUSED(reports))
     }
   }
 
+  if (!MAIN_VERSION_ATLEAST(bmain, 300, 25)) {
+    version_node_socket_index_animdata(bmain, NTREE_SHADER, SH_NODE_BSDF_PRINCIPLED, 4, 2, 25);
+  }
+
   if (!MAIN_VERSION_ATLEAST(bmain, 300, 26)) {
     LISTBASE_FOREACH (Scene *, scene, &bmain->scenes) {
       ToolSettings *tool_settings = scene->toolsettings;
@@ -780,7 +784,6 @@ void do_versions_after_linking_300(Main *bmain, ReportList *UNUSED(reports))
   {
     /* Keep this block, even when empty. */
 
-    version_node_socket_index_animdata(bmain, NTREE_SHADER, SH_NODE_BSDF_PRINCIPLED, 4, 2, 25);
   }
 }
 

source/blender/blenloader/intern/versioning_common.cc

@@ -127,9 +127,9 @@ void version_node_input_socket_name(bNodeTree *ntree,
 }
 
 void version_node_output_socket_name(bNodeTree *ntree,
-                                    const int node_type,
-                                    const char *old_name,
-                                    const char *new_name)
+                                     const int node_type,
+                                     const char *old_name,
+                                     const char *new_name)
 {
   LISTBASE_FOREACH (bNode *, node, &ntree->nodes) {
     if (node->type == node_type) {

source/blender/bmesh/intern/bmesh_construct.c

@@ -31,6 +31,7 @@
 
 #include "BKE_customdata.h"
 
+#include "DNA_mesh_types.h"
 #include "DNA_meshdata_types.h"
 
 #include "bmesh.h"
@@ -589,6 +590,25 @@ static BMFace *bm_mesh_copy_new_face(
   return f_new;
 }
 
+void BM_mesh_copy_init_customdata_from_mesh(BMesh *bm_dst,
+                                            const Mesh *me_src,
+                                            const BMAllocTemplate *allocsize)
+{
+  if (allocsize == NULL) {
+    allocsize = &bm_mesh_allocsize_default;
+  }
+
+  CustomData_copy(&me_src->vdata, &bm_dst->vdata, CD_MASK_BMESH.vmask, CD_CALLOC, 0);
+  CustomData_copy(&me_src->edata, &bm_dst->edata, CD_MASK_BMESH.emask, CD_CALLOC, 0);
+  CustomData_copy(&me_src->ldata, &bm_dst->ldata, CD_MASK_BMESH.lmask, CD_CALLOC, 0);
+  CustomData_copy(&me_src->pdata, &bm_dst->pdata, CD_MASK_BMESH.pmask, CD_CALLOC, 0);
+
+  CustomData_bmesh_init_pool(&bm_dst->vdata, allocsize->totvert, BM_VERT);
+  CustomData_bmesh_init_pool(&bm_dst->edata, allocsize->totedge, BM_EDGE);
+  CustomData_bmesh_init_pool(&bm_dst->ldata, allocsize->totloop, BM_LOOP);
+  CustomData_bmesh_init_pool(&bm_dst->pdata, allocsize->totface, BM_FACE);
+}
+
 void BM_mesh_copy_init_customdata(BMesh *bm_dst, BMesh *bm_src, const BMAllocTemplate *allocsize)
 {
   if (allocsize == NULL) {

source/blender/bmesh/intern/bmesh_construct.h

@@ -23,6 +23,7 @@
 #include "bmesh_core.h"
 
 struct BMAllocTemplate;
+struct Mesh;
 
 bool BM_verts_from_edges(BMVert **vert_arr, BMEdge **edge_arr, const int len);
 
@@ -66,6 +67,9 @@ void BM_elem_attrs_copy_ex(BMesh *bm_src,
 void BM_elem_attrs_copy(BMesh *bm_src, BMesh *bm_dst, const void *ele_src_v, void *ele_dst_v);
 void BM_elem_select_copy(BMesh *bm_dst, void *ele_dst_v, const void *ele_src_v);
 
+void BM_mesh_copy_init_customdata_from_mesh(BMesh *bm_dst,
+                                            const struct Mesh *me_src,
+                                            const struct BMAllocTemplate *allocsize);
 void BM_mesh_copy_init_customdata(BMesh *bm_dst,
                                   BMesh *bm_src,
                                   const struct BMAllocTemplate *allocsize);

source/blender/bmesh/intern/bmesh_mesh_normals.c

@@ -404,16 +404,13 @@ void BM_normals_loops_edges_tag(BMesh *bm, const bool do_edges)
  */
 static void bm_mesh_edges_sharp_tag(BMesh *bm,
                                     const float (*fnos)[3],
-                                    const float split_angle,
+                                    float split_angle_cos,
                                     const bool do_sharp_edges_tag)
 {
   BMIter eiter;
   BMEdge *e;
   int i;
 
-  const bool check_angle = (split_angle < (float)M_PI);
-  const float split_angle_cos = check_angle ? cosf(split_angle) : -1.0f;
-
   if (fnos) {
     BM_mesh_elem_index_ensure(bm, BM_FACE);
   }
@@ -451,7 +448,7 @@ void BM_edges_sharp_from_angle_set(BMesh *bm, const float split_angle)
     return;
   }
 
-  bm_mesh_edges_sharp_tag(bm, NULL, split_angle, true);
+  bm_mesh_edges_sharp_tag(bm, NULL, cosf(split_angle), true);
 }
 
 /** \} */
@@ -1110,11 +1107,13 @@ static void bm_mesh_loops_calc_normals__single_threaded(BMesh *bm,
                                                         const short (*clnors_data)[2],
                                                         const int cd_loop_clnors_offset,
                                                         const bool do_rebuild,
-                                                        const float split_angle)
+                                                        const float split_angle_cos)
 {
   BMIter fiter;
   BMFace *f_curr;
   const bool has_clnors = clnors_data || (cd_loop_clnors_offset != -1);
+  /* When false the caller must have already tagged the edges. */
+  const bool do_edge_tag = (split_angle_cos != EDGE_TAG_FROM_SPLIT_ANGLE_BYPASS);
 
   MLoopNorSpaceArray _lnors_spacearr = {NULL};
 
@@ -1155,7 +1154,9 @@ static void bm_mesh_loops_calc_normals__single_threaded(BMesh *bm,
 
   /* Always tag edges based on winding & sharp edge flag
    * (even when the auto-smooth angle doesn't need to be calculated). */
-  bm_mesh_edges_sharp_tag(bm, fnos, has_clnors ? (float)M_PI : split_angle, false);
+  if (do_edge_tag) {
+    bm_mesh_edges_sharp_tag(bm, fnos, has_clnors ? -1.0f : split_angle_cos, false);
+  }
 
   /* We now know edges that can be smoothed (they are tagged),
    * and edges that will be hard (they aren't).
@@ -1308,12 +1309,9 @@ static void bm_mesh_loops_calc_normals__multi_threaded(BMesh *bm,
                                                        const short (*clnors_data)[2],
                                                        const int cd_loop_clnors_offset,
                                                        const bool do_rebuild,
-                                                       const float split_angle)
+                                                       const float split_angle_cos)
 {
   const bool has_clnors = clnors_data || (cd_loop_clnors_offset != -1);
-  const bool check_angle = (split_angle < (float)M_PI);
-  const float split_angle_cos = check_angle ? cosf(split_angle) : -1.0f;
-
   MLoopNorSpaceArray _lnors_spacearr = {NULL};
 
   {
@@ -1387,7 +1385,7 @@ static void bm_mesh_loops_calc_normals(BMesh *bm,
                                        const short (*clnors_data)[2],
                                        const int cd_loop_clnors_offset,
                                        const bool do_rebuild,
-                                       const float split_angle)
+                                       const float split_angle_cos)
 {
   if (bm->totloop < BM_OMP_LIMIT) {
     bm_mesh_loops_calc_normals__single_threaded(bm,
@@ -1398,7 +1396,7 @@ static void bm_mesh_loops_calc_normals(BMesh *bm,
                                                 clnors_data,
                                                 cd_loop_clnors_offset,
                                                 do_rebuild,
-                                                split_angle);
+                                                split_angle_cos);
   }
   else {
     bm_mesh_loops_calc_normals__multi_threaded(bm,
@@ -1409,7 +1407,7 @@ static void bm_mesh_loops_calc_normals(BMesh *bm,
                                                clnors_data,
                                                cd_loop_clnors_offset,
                                                do_rebuild,
-                                               split_angle);
+                                               split_angle_cos);
   }
 }
 
@@ -1620,7 +1618,7 @@ static void bm_mesh_loops_custom_normals_set(BMesh *bm,
 
   /* Tag smooth edges and set lnos from vnos when they might be completely smooth...
    * When using custom loop normals, disable the angle feature! */
-  bm_mesh_edges_sharp_tag(bm, fnos, (float)M_PI, false);
+  bm_mesh_edges_sharp_tag(bm, fnos, -1.0f, false);
 
   /* Finish computing lnos by accumulating face normals
    * in each fan of faces defined by sharp edges. */
@@ -1751,7 +1749,7 @@ void BM_loops_calc_normal_vcos(BMesh *bm,
                                clnors_data,
                                cd_loop_clnors_offset,
                                do_rebuild,
-                               has_clnors ? (float)M_PI : split_angle);
+                               has_clnors ? -1.0f : cosf(split_angle));
   }
   else {
     BLI_assert(!r_lnors_spacearr);

source/blender/draw/engines/basic/basic_engine.c

@@ -38,7 +38,6 @@
 #include "basic_engine.h"
 #include "basic_private.h"
 
-
 #define BASIC_ENGINE "BLENDER_BASIC"
 
 /* *********** LISTS *********** */
@@ -107,14 +106,14 @@ static void basic_cache_init(void *vedata)
              BASIC_shaders_pointcloud_depth_conservative_sh_get(draw_ctx->sh_cfg) :
              BASIC_shaders_pointcloud_depth_sh_get(draw_ctx->sh_cfg);
     DRW_PASS_CREATE(psl->depth_pass_pointcloud[i], state | clip_state | infront_state);
-    stl->g_data->depth_pointcloud_shgrp[i] = grp = DRW_shgroup_create(sh, psl->depth_pass_pointcloud[i]);
+    stl->g_data->depth_pointcloud_shgrp[i] = grp = DRW_shgroup_create(
+        sh, psl->depth_pass_pointcloud[i]);
     DRW_shgroup_uniform_vec2(grp, "sizeViewport", DRW_viewport_size_get(), 1);
     DRW_shgroup_uniform_vec2(grp, "sizeViewportInv", DRW_viewport_invert_size_get(), 1);
 
     stl->g_data->depth_hair_shgrp[i] = grp = DRW_shgroup_create(
         BASIC_shaders_depth_sh_get(draw_ctx->sh_cfg), psl->depth_pass[i]);
 
-
     sh = DRW_state_is_select() ? BASIC_shaders_depth_conservative_sh_get(draw_ctx->sh_cfg) :
                                  BASIC_shaders_depth_sh_get(draw_ctx->sh_cfg);
     state |= DRW_STATE_CULL_BACK;

source/blender/draw/engines/image/image_drawing_mode.hh

@@ -151,4 +151,3 @@ class DefaultDrawingMode : public AbstractDrawingMode {
 };
 
 }  // namespace blender::draw::image_engine
-

source/blender/draw/engines/image/image_engine.h

@@ -31,4 +31,3 @@ extern DrawEngineType draw_engine_image_type;
 #ifdef __cplusplus
 }
 #endif
-

source/blender/draw/engines/image/image_private.hh

@@ -194,4 +194,3 @@ void IMAGE_shader_library_ensure(void);
 void IMAGE_shader_free(void);
 
 }  // namespace blender::draw::image_engine
-

source/blender/draw/engines/image/image_space_image.hh

@@ -180,4 +180,3 @@ class SpaceImageAccessor : public AbstractSpaceAccessor {
 };
 
 }  // namespace blender::draw::image_engine
-

source/blender/draw/engines/image/image_space_node.hh

@@ -136,4 +136,3 @@ class SpaceNodeAccessor : public AbstractSpaceAccessor {
 };
 
 }  // namespace blender::draw::image_engine
-

source/blender/draw/engines/overlay/overlay_extra.c

@@ -762,10 +762,7 @@ void OVERLAY_lightprobe_cache_populate(OVERLAY_Data *vedata, Object *ob)
         instdata.mat[1][3] = prb->grid_resolution_y;
         instdata.mat[2][3] = prb->grid_resolution_z;
         /* Put theme id in matrix. */
-        if (UNLIKELY(ob->base_flag & BASE_FROM_DUPLI)) {
-          instdata.mat[3][3] = 0.0;
-        }
-        else if (theme_id == TH_ACTIVE) {
+        if (theme_id == TH_ACTIVE) {
           instdata.mat[3][3] = 1.0;
         }
         else /* TH_SELECT */ {

source/blender/draw/engines/overlay/shaders/extra_lightprobe_grid_vert.glsl

@@ -10,9 +10,6 @@ vec4 color_from_id(float color_id)
   if (isTransform) {
     return colorTransform;
   }
-  else if (color_id == 0.0) {
-    return colorDupliSelect;
-  }
   else if (color_id == 1.0) {
     return colorActive;
   }

source/blender/draw/engines/overlay/shaders/outline_detect_frag.glsl

@@ -241,9 +241,6 @@ void main()
   else if (color_id == 1u) {
     fragColor = colorSelect;
   }
-  else if (color_id == 2u) {
-    fragColor = colorDupliSelect;
-  }
   else if (color_id == 3u) {
     fragColor = colorActive;
   }

source/blender/draw/engines/overlay/shaders/outline_prepass_vert.glsl

@@ -17,18 +17,8 @@ flat out uint objectId;
 uint outline_colorid_get(void)
 {
   int flag = int(abs(ObjectInfo.w));
-  bool is_from_dupli = (flag & DRW_BASE_FROM_DUPLI) != 0;
   bool is_active = (flag & DRW_BASE_ACTIVE) != 0;
 
-  if (is_from_dupli) {
-    if (isTransform) {
-      return 0u; /* colorTransform */
-    }
-    else {
-      return 2u; /* colorDupliSelect */
-    }
-  }
-
   if (isTransform) {
     return 0u; /* colorTransform */
   }

source/blender/draw/engines/overlay/shaders/wireframe_vert.glsl

@@ -28,27 +28,12 @@ void wire_color_get(out vec3 rim_col, out vec3 wire_col)
 {
   int flag = int(abs(ObjectInfo.w));
   bool is_selected = (flag & DRW_BASE_SELECTED) != 0;
-  bool is_from_dupli = (flag & DRW_BASE_FROM_DUPLI) != 0;
   bool is_from_set = (flag & DRW_BASE_FROM_SET) != 0;
   bool is_active = (flag & DRW_BASE_ACTIVE) != 0;
 
   if (is_from_set) {
-    rim_col = colorDupli.rgb;
-    wire_col = colorDupli.rgb;
-  }
-  else if (is_from_dupli) {
-    if (is_selected) {
-      if (isTransform) {
-        rim_col = colorTransform.rgb;
-      }
-      else {
-        rim_col = colorDupliSelect.rgb;
-      }
-    }
-    else {
-      rim_col = colorDupli.rgb;
-    }
-    wire_col = colorDupli.rgb;
+    rim_col = colorWire.rgb;
+    wire_col = colorWire.rgb;
   }
   else if (is_selected && useColoring) {
     if (isTransform) {

source/blender/draw/intern/draw_common.c

@@ -101,11 +101,6 @@ void DRW_globals_update(void)
       gb->colorEditMeshMiddle,
       dot_v3v3(gb->colorEditMeshMiddle, (float[3]){0.3333f, 0.3333f, 0.3333f})); /* Desaturate */
 
-  interp_v4_v4v4(gb->colorDupliSelect, gb->colorBackground, gb->colorSelect, 0.5f);
-  /* Was 50% in 2.7x since the background was lighter making it easier to tell the color from
-   * black, with a darker background we need a more faded color. */
-  interp_v4_v4v4(gb->colorDupli, gb->colorBackground, gb->colorWire, 0.3f);
-
 #ifdef WITH_FREESTYLE
   UI_GetThemeColor4fv(TH_FREESTYLE_EDGE_MARK, gb->colorEdgeFreestyle);
   UI_GetThemeColor4fv(TH_FREESTYLE_FACE_MARK, gb->colorFaceFreestyle);
@@ -300,7 +295,11 @@ int DRW_object_wire_theme_get(Object *ob, ViewLayer *view_layer, float **r_color
 {
   const DRWContextState *draw_ctx = DRW_context_state_get();
   const bool is_edit = (draw_ctx->object_mode & OB_MODE_EDIT) && (ob->mode & OB_MODE_EDIT);
-  const bool active = (view_layer->basact && view_layer->basact->object == ob);
+  const bool active = view_layer->basact &&
+                      ((ob->base_flag & BASE_FROM_DUPLI) ?
+                           (DRW_object_get_dupli_parent(ob) == view_layer->basact->object) :
+                           (view_layer->basact->object == ob));
+
   /* confusing logic here, there are 2 methods of setting the color
    * 'colortab[colindex]' and 'theme_id', colindex overrides theme_id.
    *
@@ -345,21 +344,7 @@ int DRW_object_wire_theme_get(Object *ob, ViewLayer *view_layer, float **r_color
 
   if (r_color != NULL) {
     if (UNLIKELY(ob->base_flag & BASE_FROM_SET)) {
-      *r_color = G_draw.block.colorDupli;
-    }
-    else if (UNLIKELY(ob->base_flag & BASE_FROM_DUPLI)) {
-      switch (theme_id) {
-        case TH_ACTIVE:
-        case TH_SELECT:
-          *r_color = G_draw.block.colorDupliSelect;
-          break;
-        case TH_TRANSFORM:
-          *r_color = G_draw.block.colorTransform;
-          break;
-        default:
-          *r_color = G_draw.block.colorDupli;
-          break;
-      }
+      *r_color = G_draw.block.colorWire;
     }
     else {
       switch (theme_id) {

source/blender/draw/intern/draw_common.h

@@ -43,8 +43,6 @@ typedef struct GlobalsUboStorage {
   float colorWireEdit[4];
   float colorActive[4];
   float colorSelect[4];
-  float colorDupliSelect[4];
-  float colorDupli[4];
   float colorLibrarySelect[4];
   float colorLibrary[4];
   float colorTransform[4];

source/blender/draw/intern/draw_manager_data.c

@@ -538,7 +538,12 @@ static void drw_call_obinfos_init(DRWObjectInfos *ob_infos, Object *ob)
   ob_infos->ob_flag += (ob->base_flag & BASE_SELECTED) ? (1 << 1) : 0;
   ob_infos->ob_flag += (ob->base_flag & BASE_FROM_DUPLI) ? (1 << 2) : 0;
   ob_infos->ob_flag += (ob->base_flag & BASE_FROM_SET) ? (1 << 3) : 0;
-  ob_infos->ob_flag += (ob == DST.draw_ctx.obact) ? (1 << 4) : 0;
+  if (ob->base_flag & BASE_FROM_DUPLI) {
+    ob_infos->ob_flag += (DRW_object_get_dupli_parent(ob) == DST.draw_ctx.obact) ? (1 << 4) : 0;
+  }
+  else {
+    ob_infos->ob_flag += (ob == DST.draw_ctx.obact) ? (1 << 4) : 0;
+  }
   /* Negative scaling. */
   ob_infos->ob_flag *= (ob->transflag & OB_NEG_SCALE) ? -1.0f : 1.0f;
   /* Object Color. */

source/blender/draw/intern/shaders/common_globals_lib.glsl

@@ -7,8 +7,6 @@ layout(std140) uniform globalsBlock
   vec4 colorWireEdit;
   vec4 colorActive;
   vec4 colorSelect;
-  vec4 colorDupliSelect;
-  vec4 colorDupli;
   vec4 colorLibrarySelect;
   vec4 colorLibrary;
   vec4 colorTransform;

source/blender/draw/tests/shaders_test.cc

@@ -13,12 +13,12 @@
 
 #include "intern/draw_manager_testing.h"
 
+#include "engines/basic/basic_private.h"
 #include "engines/eevee/eevee_private.h"
 #include "engines/gpencil/gpencil_engine.h"
 #include "engines/image/image_private.hh"
 #include "engines/overlay/overlay_private.h"
 #include "engines/workbench/workbench_private.h"
-#include "engines/basic/basic_private.h"
 #include "intern/draw_shader.h"
 
 namespace blender::draw {

source/blender/editors/armature/armature_relations.c

@@ -745,6 +745,10 @@ void ARMATURE_OT_separate(wmOperatorType *ot)
 #define ARM_PAR_CONNECT 1
 #define ARM_PAR_OFFSET 2
 
+/* armature un-parenting options */
+#define ARM_PAR_CLEAR 1
+#define ARM_PAR_CLEAR_DISCONNECT 2
+
 /* check for null, before calling! */
 static void bone_connect_to_existing_parent(EditBone *bone)
 {
@@ -904,19 +908,29 @@ static int armature_parent_set_invoke(bContext *C,
                                       wmOperator *UNUSED(op),
                                       const wmEvent *UNUSED(event))
 {
-  bool all_childbones = false;
+  /* False when all selected bones are parented to the active bone. */
+  bool enable_offset = false;
+  /* False when all selected bones are connected to the active bone. */
+  bool enable_connect = false;
   {
     Object *ob = CTX_data_edit_object(C);
     bArmature *arm = ob->data;
     EditBone *actbone = arm->act_edbone;
     LISTBASE_FOREACH (EditBone *, ebone, arm->edbo) {
-      if (EBONE_EDITABLE(ebone) && (ebone->flag & BONE_SELECTED)) {
-        if (ebone != actbone) {
-          if (ebone->parent != actbone) {
-            all_childbones = true;
-            break;
-          }
-        }
+      if (!EBONE_EDITABLE(ebone) || !(ebone->flag & BONE_SELECTED)) {
+        continue;
+      }
+      if (ebone == actbone) {
+        continue;
+      }
+
+      if (ebone->parent != actbone) {
+        enable_offset = true;
+        enable_connect = true;
+        break;
+      }
+      else if (!(ebone->flag & BONE_CONNECTED)) {
+        enable_connect = true;
       }
     }
   }
@@ -924,11 +938,14 @@ static int armature_parent_set_invoke(bContext *C,
   uiPopupMenu *pup = UI_popup_menu_begin(
       C, CTX_IFACE_(BLT_I18NCONTEXT_OPERATOR_DEFAULT, "Make Parent"), ICON_NONE);
   uiLayout *layout = UI_popup_menu_layout(pup);
-  uiItemEnumO(layout, "ARMATURE_OT_parent_set", NULL, 0, "type", ARM_PAR_CONNECT);
-  if (all_childbones) {
-    /* Object becomes parent, make the associated menus. */
-    uiItemEnumO(layout, "ARMATURE_OT_parent_set", NULL, 0, "type", ARM_PAR_OFFSET);
-  }
+
+  uiLayout *row_offset = uiLayoutRow(layout, false);
+  uiLayoutSetEnabled(row_offset, enable_offset);
+  uiItemEnumO(row_offset, "ARMATURE_OT_parent_set", NULL, 0, "type", ARM_PAR_OFFSET);
+
+  uiLayout *row_connect = uiLayoutRow(layout, false);
+  uiLayoutSetEnabled(row_connect, enable_connect);
+  uiItemEnumO(row_connect, "ARMATURE_OT_parent_set", NULL, 0, "type", ARM_PAR_CONNECT);
 
   UI_popup_menu_end(C, pup);
 
@@ -955,8 +972,8 @@ void ARMATURE_OT_parent_set(wmOperatorType *ot)
 }
 
 static const EnumPropertyItem prop_editarm_clear_parent_types[] = {
-    {1, "CLEAR", 0, "Clear Parent", ""},
-    {2, "DISCONNECT", 0, "Disconnect Bone", ""},
+    {ARM_PAR_CLEAR, "CLEAR", 0, "Clear Parent", ""},
+    {ARM_PAR_CLEAR_DISCONNECT, "DISCONNECT", 0, "Disconnect Bone", ""},
     {0, NULL, 0, NULL, NULL},
 };
 
@@ -1012,6 +1029,51 @@ static int armature_parent_clear_exec(bContext *C, wmOperator *op)
   return OPERATOR_FINISHED;
 }
 
+static int armature_parent_clear_invoke(bContext *C,
+                                        wmOperator *UNUSED(op),
+                                        const wmEvent *UNUSED(event))
+{
+  /* False when no selected bones are connected to the active bone. */
+  bool enable_disconnect = false;
+  /* False when no selected bones are parented to the active bone. */
+  bool enable_clear = false;
+  {
+    Object *ob = CTX_data_edit_object(C);
+    bArmature *arm = ob->data;
+    LISTBASE_FOREACH (EditBone *, ebone, arm->edbo) {
+      if (!EBONE_EDITABLE(ebone) || !(ebone->flag & BONE_SELECTED)) {
+        continue;
+      }
+      if (ebone->parent == NULL) {
+        continue;
+      }
+      enable_clear = true;
+
+      if (ebone->flag & BONE_CONNECTED) {
+        enable_disconnect = true;
+        break;
+      }
+    }
+  }
+
+  uiPopupMenu *pup = UI_popup_menu_begin(
+      C, CTX_IFACE_(BLT_I18NCONTEXT_OPERATOR_DEFAULT, "Clear Parent"), ICON_NONE);
+  uiLayout *layout = UI_popup_menu_layout(pup);
+
+  uiLayout *row_clear = uiLayoutRow(layout, false);
+  uiLayoutSetEnabled(row_clear, enable_clear);
+  uiItemEnumO(row_clear, "ARMATURE_OT_parent_clear", NULL, 0, "type", ARM_PAR_CLEAR);
+
+  uiLayout *row_disconnect = uiLayoutRow(layout, false);
+  uiLayoutSetEnabled(row_disconnect, enable_disconnect);
+  uiItemEnumO(
+      row_disconnect, "ARMATURE_OT_parent_clear", NULL, 0, "type", ARM_PAR_CLEAR_DISCONNECT);
+
+  UI_popup_menu_end(C, pup);
+
+  return OPERATOR_INTERFACE;
+}
+
 void ARMATURE_OT_parent_clear(wmOperatorType *ot)
 {
   /* identifiers */
@@ -1021,7 +1083,7 @@ void ARMATURE_OT_parent_clear(wmOperatorType *ot)
       "Remove the parent-child relationship between selected bones and their parents";
 
   /* api callbacks */
-  ot->invoke = WM_menu_invoke;
+  ot->invoke = armature_parent_clear_invoke;
   ot->exec = armature_parent_clear_exec;
   ot->poll = ED_operator_editarmature;