CMakeLists.txt

@@ -521,7 +521,8 @@ endif()
 if(NOT APPLE)
   option(WITH_CYCLES_DEVICE_HIP        "Enable Cycles AMD HIP support" ON)
   option(WITH_CYCLES_HIP_BINARIES      "Build Cycles AMD HIP binaries" OFF)
-  set(CYCLES_HIP_BINARIES_ARCH gfx900 gfx906 gfx90c gfx902 gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1034 gfx1035 gfx1100 gfx1101 gfx1102 CACHE STRING "AMD HIP architectures to build binaries for")
+  # Radeon VII (gfx906) not currently working with HIP SDK, so left out of the list.
+  set(CYCLES_HIP_BINARIES_ARCH gfx900 gfx90c gfx902 gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1034 gfx1035 gfx1100 gfx1101 gfx1102 CACHE STRING "AMD HIP architectures to build binaries for")
   mark_as_advanced(WITH_CYCLES_DEVICE_HIP)
   mark_as_advanced(CYCLES_HIP_BINARIES_ARCH)
 endif()

build_files/build_environment/cmake/dpcpp.cmake

@@ -113,8 +113,14 @@ add_dependencies(
 
 if(BUILD_MODE STREQUAL Release AND WIN32)
   ExternalProject_Add_Step(external_dpcpp after_install
-      COMMAND ${CMAKE_COMMAND} -E rm -f ${LIBDIR}/dpcpp/bin/clang-cl.exe
-      COMMAND ${CMAKE_COMMAND} -E rm -f ${LIBDIR}/dpcpp/bin/clang-cpp.exe
       COMMAND ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/dpcpp ${HARVEST_TARGET}/dpcpp
+      COMMAND ${CMAKE_COMMAND} -E rm -f ${HARVEST_TARGET}/dpcpp/bin/clang-cl.exe
+      COMMAND ${CMAKE_COMMAND} -E rm -f ${HARVEST_TARGET}/dpcpp/bin/clang-cpp.exe
+      COMMAND ${CMAKE_COMMAND} -E rm -f ${HARVEST_TARGET}/dpcpp/bin/clang.exe
+      COMMAND ${CMAKE_COMMAND} -E rm -f ${HARVEST_TARGET}/dpcpp/bin/ld.lld.exe
+      COMMAND ${CMAKE_COMMAND} -E rm -f ${HARVEST_TARGET}/dpcpp/bin/ld64.lld.exe
+      COMMAND ${CMAKE_COMMAND} -E rm -f ${HARVEST_TARGET}/dpcpp/bin/lld.exe
+      COMMAND ${CMAKE_COMMAND} -E rm -f ${HARVEST_TARGET}/dpcpp/bin/lld-link.exe
+      COMMAND ${CMAKE_COMMAND} -E rm -f ${HARVEST_TARGET}/dpcpp/bin/wasm-ld.exe
   )
 endif()

build_files/cmake/Modules/FindHIP.cmake

@@ -37,18 +37,24 @@ elseif(HIP_HIPCC_EXECUTABLE)
   set(HIP_VERSION_MINOR 0)
   set(HIP_VERSION_PATCH 0)
 
+  if(WIN32)
+    set(_hipcc_executable ${HIP_HIPCC_EXECUTABLE}.bat)
+  else()
+    set(_hipcc_executable ${HIP_HIPCC_EXECUTABLE})
+  endif()
+
   # Get version from the output.
-  execute_process(COMMAND ${HIP_HIPCC_EXECUTABLE} --version
+  execute_process(COMMAND ${_hipcc_executable} --version
-                  OUTPUT_VARIABLE HIP_VERSION_RAW
+                  OUTPUT_VARIABLE _hip_version_raw
                   ERROR_QUIET
                   OUTPUT_STRIP_TRAILING_WHITESPACE)
 
   # Parse parts.
-  if(HIP_VERSION_RAW MATCHES "HIP version: .*")
+  if(_hip_version_raw MATCHES "HIP version: .*")
     # Strip the HIP prefix and get list of individual version components.
     string(REGEX REPLACE
            ".*HIP version: ([.0-9]+).*" "\\1"
-           HIP_SEMANTIC_VERSION "${HIP_VERSION_RAW}")
+           HIP_SEMANTIC_VERSION "${_hip_version_raw}")
     string(REPLACE "." ";" HIP_VERSION_PARTS "${HIP_SEMANTIC_VERSION}")
     list(LENGTH HIP_VERSION_PARTS NUM_HIP_VERSION_PARTS)
 
@@ -71,7 +77,13 @@ elseif(HIP_HIPCC_EXECUTABLE)
 
   # Construct full semantic version.
   set(HIP_VERSION "${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_PATCH}")
-  unset(HIP_VERSION_RAW)
+  unset(_hip_version_raw)
+  unset(_hipcc_executable)
 else()
   set(HIP_FOUND FALSE)
 endif()
+
+include(FindPackageHandleStandardArgs)
+find_package_handle_standard_args(HIP
+    REQUIRED_VARS HIP_HIPCC_EXECUTABLE
+    VERSION_VAR HIP_VERSION)

build_files/cmake/config/blender_release.cmake

@@ -82,7 +82,7 @@ if(NOT APPLE)
   set(WITH_CYCLES_DEVICE_OPTIX    ON  CACHE BOOL "" FORCE)
   set(WITH_CYCLES_CUDA_BINARIES   ON  CACHE BOOL "" FORCE)
   set(WITH_CYCLES_CUBIN_COMPILER  OFF CACHE BOOL "" FORCE)
-  set(WITH_CYCLES_HIP_BINARIES    OFF CACHE BOOL "" FORCE)
+  set(WITH_CYCLES_HIP_BINARIES    ON  CACHE BOOL "" FORCE)
   set(WITH_CYCLES_DEVICE_ONEAPI   ON  CACHE BOOL "" FORCE)
   set(WITH_CYCLES_ONEAPI_BINARIES ON  CACHE BOOL "" FORCE)
 endif()

build_files/cmake/platform/platform_apple.cmake

@@ -174,7 +174,7 @@ if(SYSTEMSTUBS_LIBRARY)
   list(APPEND PLATFORM_LINKLIBS SystemStubs)
 endif()
 
-string(APPEND PLATFORM_CFLAGS " -pipe -funsigned-char -fno-strict-aliasing")
+string(APPEND PLATFORM_CFLAGS " -pipe -funsigned-char -fno-strict-aliasing -ffp-contract=off")
 set(PLATFORM_LINKFLAGS
   "-fexceptions -framework CoreServices -framework Foundation -framework IOKit -framework AppKit -framework Cocoa -framework Carbon -framework AudioUnit -framework AudioToolbox -framework CoreAudio -framework Metal -framework QuartzCore"
 )

build_files/cmake/platform/platform_unix.cmake

@@ -803,8 +803,7 @@ if(CMAKE_COMPILER_IS_GNUCC)
   # Automatically turned on when building with "-march=native". This is
   # explicitly turned off here as it will make floating point math give a bit
   # different results. This will lead to automated test failures. So disable
-  # this until we support it. Seems to default to off in clang and the intel
+  # this until we support it.
-  # compiler.
   set(PLATFORM_CFLAGS "-pipe -fPIC -funsigned-char -fno-strict-aliasing -ffp-contract=off")
 
   # `maybe-uninitialized` is unreliable in release builds, but fine in debug builds.
@@ -892,7 +891,7 @@ if(CMAKE_COMPILER_IS_GNUCC)
 
 # CLang is the same as GCC for now.
 elseif(CMAKE_C_COMPILER_ID MATCHES "Clang")
-  set(PLATFORM_CFLAGS "-pipe -fPIC -funsigned-char -fno-strict-aliasing")
+  set(PLATFORM_CFLAGS "-pipe -fPIC -funsigned-char -fno-strict-aliasing -ffp-contract=off")
 
   if(WITH_LINKER_MOLD AND _IS_LINKER_DEFAULT)
     find_program(MOLD_BIN "mold")

build_files/config/pipeline_config.yaml

@@ -9,7 +9,7 @@ buildbot:
     cuda11:
         version: '11.4.1'
     hip:
-        version: '5.3.22480'
+        version: '5.5.30571'
     optix:
         version: '7.3.0'
     ocloc:

intern/cycles/blender/addon/properties.py

@@ -1683,9 +1683,6 @@ class CyclesPreferences(bpy.types.AddonPreferences):
                 col.label(text=iface_("and NVIDIA driver version %s or newer") % driver_version,
                           icon='BLANK1', translate=False)
             elif device_type == 'HIP':
-                if True:
-                    col.label(text="HIP temporarily disabled due to compiler bugs", icon='BLANK1')
-                else:
                 import sys
                 if sys.platform[:3] == "win":
                     driver_version = "21.Q4"
@@ -1693,6 +1690,9 @@ class CyclesPreferences(bpy.types.AddonPreferences):
                     col.label(text=iface_("and AMD Radeon Pro %s driver or newer") % driver_version,
                               icon='BLANK1', translate=False)
                 elif sys.platform.startswith("linux"):
+                    if True:
+                        col.label(text="HIP temporarily disabled due to compiler bugs", icon='BLANK1')
+                    else:
                         driver_version = "22.10"
                         col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
                         col.label(text=iface_("and AMD driver version %s or newer") % driver_version, icon='BLANK1',

intern/cycles/cmake/external_libs.cmake

@@ -42,15 +42,19 @@ endif()
 ###########################################################################
 
 if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
+  if(UNIX)
+    # Disabled until there is a HIP 5.5 release for Linux.
     set(WITH_CYCLES_HIP_BINARIES OFF)
     message(STATUS "HIP temporarily disabled due to compiler bugs")
+  else()
+    # Need at least HIP 5.5 to solve compiler bug affecting the kernel.
+    find_package(HIP 5.5.0)
+    set_and_warn_library_found("HIP compiler" HIP_FOUND WITH_CYCLES_HIP_BINARIES)
 
-  # find_package(HIP)
+    if(HIP_FOUND)
-  # set_and_warn_library_found("HIP compiler" HIP_FOUND WITH_CYCLES_HIP_BINARIES)
+      message(STATUS "Found HIP ${HIP_HIPCC_EXECUTABLE} (${HIP_VERSION})")
-
+    endif()
-  # if(HIP_FOUND)
+  endif()
-  #   message(STATUS "Found HIP ${HIP_HIPCC_EXECUTABLE} (${HIP_VERSION})")
-  # endif()
 endif()
 
 if(NOT WITH_HIP_DYNLOAD)

intern/cycles/device/hip/device.cpp

@@ -137,7 +137,7 @@ void device_hip_info(vector<DeviceInfo> &devices)
     info.num = num;
 
     info.has_nanovdb = true;
-    info.has_light_tree = false;
+    info.has_light_tree = true;
     info.denoisers = 0;
 
     info.has_gpu_queue = true;

intern/cycles/device/oneapi/device_impl.cpp

@@ -866,9 +866,9 @@ char *OneapiDevice::device_capabilities()
 
     sycl::id<3> max_work_item_sizes =
         device.get_info<sycl::info::device::max_work_item_sizes<3>>();
-    WRITE_ATTR("max_work_item_sizes_dim0", ((size_t)max_work_item_sizes.get(0)))
+    WRITE_ATTR(max_work_item_sizes_dim0, ((size_t)max_work_item_sizes.get(0)))
-    WRITE_ATTR("max_work_item_sizes_dim1", ((size_t)max_work_item_sizes.get(1)))
+    WRITE_ATTR(max_work_item_sizes_dim1, ((size_t)max_work_item_sizes.get(1)))
-    WRITE_ATTR("max_work_item_sizes_dim2", ((size_t)max_work_item_sizes.get(2)))
+    WRITE_ATTR(max_work_item_sizes_dim2, ((size_t)max_work_item_sizes.get(2)))
 
     GET_NUM_ATTR(max_work_group_size)
     GET_NUM_ATTR(max_num_sub_groups)
@@ -891,7 +891,7 @@ char *OneapiDevice::device_capabilities()
     GET_NUM_ATTR(native_vector_width_half)
 
     size_t max_clock_frequency = device.get_info<sycl::info::device::max_clock_frequency>();
-    WRITE_ATTR("max_clock_frequency", max_clock_frequency)
+    WRITE_ATTR(max_clock_frequency, max_clock_frequency)
 
     GET_NUM_ATTR(address_bits)
     GET_NUM_ATTR(max_mem_alloc_size)
@@ -900,7 +900,7 @@ char *OneapiDevice::device_capabilities()
      * supported so we always return false, even if device supports HW texture usage acceleration.
      */
     bool image_support = false;
-    WRITE_ATTR("image_support", (size_t)image_support)
+    WRITE_ATTR(image_support, (size_t)image_support)
 
     GET_NUM_ATTR(max_parameter_size)
     GET_NUM_ATTR(mem_base_addr_align)

intern/cycles/kernel/types.h

@@ -79,9 +79,8 @@ CCL_NAMESPACE_BEGIN
 #define __VISIBILITY_FLAG__
 #define __VOLUME__
 
-/* TODO: solve internal compiler errors and enable light tree on HIP. */
 /* TODO: solve internal compiler perf issue and enable light tree on Metal/AMD. */
-#if defined(__KERNEL_HIP__) || defined(__KERNEL_METAL_AMD__)
+#if defined(__KERNEL_METAL_AMD__)
 #  undef __LIGHT_TREE__
 #endif
 

intern/cycles/scene/CMakeLists.txt

@@ -15,8 +15,12 @@ set(SRC
   camera.cpp
   colorspace.cpp
   constant_fold.cpp
+  devicescene.cpp
   film.cpp
   geometry.cpp
+  geometry_attributes.cpp
+  geometry_bvh.cpp
+  geometry_mesh.cpp
   hair.cpp
   image.cpp
   image_oiio.cpp
@@ -55,6 +59,7 @@ set(SRC_HEADERS
   camera.h
   colorspace.h
   constant_fold.h
+  devicescene.h
   film.h
   geometry.h
   hair.h

intern/cycles/scene/devicescene.cpp

@@ -0,0 +1,64 @@
+/* SPDX-License-Identifier: Apache-2.0
+ * Copyright 2011-2022 Blender Foundation */
+
+#include "scene/devicescene.h"
+#include "device/device.h"
+#include "device/memory.h"
+
+CCL_NAMESPACE_BEGIN
+
+DeviceScene::DeviceScene(Device *device)
+    : bvh_nodes(device, "bvh_nodes", MEM_GLOBAL),
+      bvh_leaf_nodes(device, "bvh_leaf_nodes", MEM_GLOBAL),
+      object_node(device, "object_node", MEM_GLOBAL),
+      prim_type(device, "prim_type", MEM_GLOBAL),
+      prim_visibility(device, "prim_visibility", MEM_GLOBAL),
+      prim_index(device, "prim_index", MEM_GLOBAL),
+      prim_object(device, "prim_object", MEM_GLOBAL),
+      prim_time(device, "prim_time", MEM_GLOBAL),
+      tri_verts(device, "tri_verts", MEM_GLOBAL),
+      tri_shader(device, "tri_shader", MEM_GLOBAL),
+      tri_vnormal(device, "tri_vnormal", MEM_GLOBAL),
+      tri_vindex(device, "tri_vindex", MEM_GLOBAL),
+      tri_patch(device, "tri_patch", MEM_GLOBAL),
+      tri_patch_uv(device, "tri_patch_uv", MEM_GLOBAL),
+      curves(device, "curves", MEM_GLOBAL),
+      curve_keys(device, "curve_keys", MEM_GLOBAL),
+      curve_segments(device, "curve_segments", MEM_GLOBAL),
+      patches(device, "patches", MEM_GLOBAL),
+      points(device, "points", MEM_GLOBAL),
+      points_shader(device, "points_shader", MEM_GLOBAL),
+      objects(device, "objects", MEM_GLOBAL),
+      object_motion_pass(device, "object_motion_pass", MEM_GLOBAL),
+      object_motion(device, "object_motion", MEM_GLOBAL),
+      object_flag(device, "object_flag", MEM_GLOBAL),
+      object_volume_step(device, "object_volume_step", MEM_GLOBAL),
+      object_prim_offset(device, "object_prim_offset", MEM_GLOBAL),
+      camera_motion(device, "camera_motion", MEM_GLOBAL),
+      attributes_map(device, "attributes_map", MEM_GLOBAL),
+      attributes_float(device, "attributes_float", MEM_GLOBAL),
+      attributes_float2(device, "attributes_float2", MEM_GLOBAL),
+      attributes_float3(device, "attributes_float3", MEM_GLOBAL),
+      attributes_float4(device, "attributes_float4", MEM_GLOBAL),
+      attributes_uchar4(device, "attributes_uchar4", MEM_GLOBAL),
+      light_distribution(device, "light_distribution", MEM_GLOBAL),
+      lights(device, "lights", MEM_GLOBAL),
+      light_background_marginal_cdf(device, "light_background_marginal_cdf", MEM_GLOBAL),
+      light_background_conditional_cdf(device, "light_background_conditional_cdf", MEM_GLOBAL),
+      light_tree_nodes(device, "light_tree_nodes", MEM_GLOBAL),
+      light_tree_emitters(device, "light_tree_emitters", MEM_GLOBAL),
+      light_to_tree(device, "light_to_tree", MEM_GLOBAL),
+      object_to_tree(device, "object_to_tree", MEM_GLOBAL),
+      object_lookup_offset(device, "object_lookup_offset", MEM_GLOBAL),
+      triangle_to_tree(device, "triangle_to_tree", MEM_GLOBAL),
+      particles(device, "particles", MEM_GLOBAL),
+      svm_nodes(device, "svm_nodes", MEM_GLOBAL),
+      shaders(device, "shaders", MEM_GLOBAL),
+      lookup_table(device, "lookup_table", MEM_GLOBAL),
+      sample_pattern_lut(device, "sample_pattern_lut", MEM_GLOBAL),
+      ies_lights(device, "ies", MEM_GLOBAL)
+{
+  memset((void *)&data, 0, sizeof(data));
+}
+
+CCL_NAMESPACE_END

intern/cycles/scene/devicescene.h

@@ -0,0 +1,101 @@
+/* SPDX-License-Identifier: Apache-2.0
+ * Copyright 2011-2022 Blender Foundation */
+
+#ifndef __DEVICESCENE_H__
+#define __DEVICESCENE_H__
+
+#include "device/device.h"
+#include "device/memory.h"
+
+#include "util/types.h"
+#include "util/vector.h"
+
+CCL_NAMESPACE_BEGIN
+
+class DeviceScene {
+ public:
+  /* BVH */
+  device_vector<int4> bvh_nodes;
+  device_vector<int4> bvh_leaf_nodes;
+  device_vector<int> object_node;
+  device_vector<int> prim_type;
+  device_vector<uint> prim_visibility;
+  device_vector<int> prim_index;
+  device_vector<int> prim_object;
+  device_vector<float2> prim_time;
+
+  /* mesh */
+  device_vector<packed_float3> tri_verts;
+  device_vector<uint> tri_shader;
+  device_vector<packed_float3> tri_vnormal;
+  device_vector<packed_uint3> tri_vindex;
+  device_vector<uint> tri_patch;
+  device_vector<float2> tri_patch_uv;
+
+  device_vector<KernelCurve> curves;
+  device_vector<float4> curve_keys;
+  device_vector<KernelCurveSegment> curve_segments;
+
+  device_vector<uint> patches;
+
+  /* point-cloud */
+  device_vector<float4> points;
+  device_vector<uint> points_shader;
+
+  /* objects */
+  device_vector<KernelObject> objects;
+  device_vector<Transform> object_motion_pass;
+  device_vector<DecomposedTransform> object_motion;
+  device_vector<uint> object_flag;
+  device_vector<float> object_volume_step;
+  device_vector<uint> object_prim_offset;
+
+  /* cameras */
+  device_vector<DecomposedTransform> camera_motion;
+
+  /* attributes */
+  device_vector<AttributeMap> attributes_map;
+  device_vector<float> attributes_float;
+  device_vector<float2> attributes_float2;
+  device_vector<packed_float3> attributes_float3;
+  device_vector<float4> attributes_float4;
+  device_vector<uchar4> attributes_uchar4;
+
+  /* lights */
+  device_vector<KernelLightDistribution> light_distribution;
+  device_vector<KernelLight> lights;
+  device_vector<float2> light_background_marginal_cdf;
+  device_vector<float2> light_background_conditional_cdf;
+
+  /* light tree */
+  device_vector<KernelLightTreeNode> light_tree_nodes;
+  device_vector<KernelLightTreeEmitter> light_tree_emitters;
+  device_vector<uint> light_to_tree;
+  device_vector<uint> object_to_tree;
+  device_vector<uint> object_lookup_offset;
+  device_vector<uint> triangle_to_tree;
+
+  /* particles */
+  device_vector<KernelParticle> particles;
+
+  /* shaders */
+  device_vector<int4> svm_nodes;
+  device_vector<KernelShader> shaders;
+
+  /* lookup tables */
+  device_vector<float> lookup_table;
+
+  /* integrator */
+  device_vector<float> sample_pattern_lut;
+
+  /* IES lights */
+  device_vector<float> ies_lights;
+
+  KernelData data;
+
+  DeviceScene(Device *device);
+};
+
+CCL_NAMESPACE_END
+
+#endif /*  __DEVICESCENE_H__ */

intern/cycles/scene/geometry.h

@@ -30,6 +30,38 @@ class Shader;
 class Volume;
 struct PackedBVH;
 
+/* Set of flags used to help determining what data has been modified or needs reallocation, so we
+ * can decide which device data to free or update. */
+enum {
+  DEVICE_CURVE_DATA_MODIFIED = (1 << 0),
+  DEVICE_MESH_DATA_MODIFIED = (1 << 1),
+  DEVICE_POINT_DATA_MODIFIED = (1 << 2),
+
+  ATTR_FLOAT_MODIFIED = (1 << 3),
+  ATTR_FLOAT2_MODIFIED = (1 << 4),
+  ATTR_FLOAT3_MODIFIED = (1 << 5),
+  ATTR_FLOAT4_MODIFIED = (1 << 6),
+  ATTR_UCHAR4_MODIFIED = (1 << 7),
+
+  CURVE_DATA_NEED_REALLOC = (1 << 8),
+  MESH_DATA_NEED_REALLOC = (1 << 9),
+  POINT_DATA_NEED_REALLOC = (1 << 10),
+
+  ATTR_FLOAT_NEEDS_REALLOC = (1 << 11),
+  ATTR_FLOAT2_NEEDS_REALLOC = (1 << 12),
+  ATTR_FLOAT3_NEEDS_REALLOC = (1 << 13),
+  ATTR_FLOAT4_NEEDS_REALLOC = (1 << 14),
+
+  ATTR_UCHAR4_NEEDS_REALLOC = (1 << 15),
+
+  ATTRS_NEED_REALLOC = (ATTR_FLOAT_NEEDS_REALLOC | ATTR_FLOAT2_NEEDS_REALLOC |
+                        ATTR_FLOAT3_NEEDS_REALLOC | ATTR_FLOAT4_NEEDS_REALLOC |
+                        ATTR_UCHAR4_NEEDS_REALLOC),
+  DEVICE_MESH_DATA_NEEDS_REALLOC = (MESH_DATA_NEED_REALLOC | ATTRS_NEED_REALLOC),
+  DEVICE_POINT_DATA_NEEDS_REALLOC = (POINT_DATA_NEED_REALLOC | ATTRS_NEED_REALLOC),
+  DEVICE_CURVE_DATA_NEEDS_REALLOC = (CURVE_DATA_NEED_REALLOC | ATTRS_NEED_REALLOC),
+};
+
 /* Geometry
  *
  * Base class for geometric types like Mesh and Hair. */

intern/cycles/scene/geometry_attributes.cpp

@@ -0,0 +1,722 @@
+/* SPDX-License-Identifier: Apache-2.0
+ * Copyright 2011-2022 Blender Foundation */
+
+#include "bvh/bvh.h"
+#include "bvh/bvh2.h"
+
+#include "device/device.h"
+
+#include "scene/attribute.h"
+#include "scene/camera.h"
+#include "scene/geometry.h"
+#include "scene/hair.h"
+#include "scene/light.h"
+#include "scene/mesh.h"
+#include "scene/object.h"
+#include "scene/pointcloud.h"
+#include "scene/scene.h"
+#include "scene/shader.h"
+#include "scene/shader_nodes.h"
+#include "scene/stats.h"
+#include "scene/volume.h"
+
+#include "subd/patch_table.h"
+#include "subd/split.h"
+
+#include "kernel/osl/globals.h"
+
+#include "util/foreach.h"
+#include "util/log.h"
+#include "util/progress.h"
+#include "util/task.h"
+
+CCL_NAMESPACE_BEGIN
+
+bool Geometry::need_attribute(Scene *scene, AttributeStandard std)
+{
+  if (std == ATTR_STD_NONE)
+    return false;
+
+  if (scene->need_global_attribute(std))
+    return true;
+
+  foreach (Node *node, used_shaders) {
+    Shader *shader = static_cast<Shader *>(node);
+    if (shader->attributes.find(std))
+      return true;
+  }
+
+  return false;
+}
+
+bool Geometry::need_attribute(Scene * /*scene*/, ustring name)
+{
+  if (name == ustring())
+    return false;
+
+  foreach (Node *node, used_shaders) {
+    Shader *shader = static_cast<Shader *>(node);
+    if (shader->attributes.find(name))
+      return true;
+  }
+
+  return false;
+}
+
+AttributeRequestSet Geometry::needed_attributes()
+{
+  AttributeRequestSet result;
+
+  foreach (Node *node, used_shaders) {
+    Shader *shader = static_cast<Shader *>(node);
+    result.add(shader->attributes);
+  }
+
+  return result;
+}
+
+bool Geometry::has_voxel_attributes() const
+{
+  foreach (const Attribute &attr, attributes.attributes) {
+    if (attr.element == ATTR_ELEMENT_VOXEL) {
+      return true;
+    }
+  }
+
+  return false;
+}
+
+/* Generate a normal attribute map entry from an attribute descriptor. */
+static void emit_attribute_map_entry(AttributeMap *attr_map,
+                                     size_t index,
+                                     uint64_t id,
+                                     TypeDesc type,
+                                     const AttributeDescriptor &desc)
+{
+  attr_map[index].id = id;
+  attr_map[index].element = desc.element;
+  attr_map[index].offset = as_uint(desc.offset);
+
+  if (type == TypeDesc::TypeFloat)
+    attr_map[index].type = NODE_ATTR_FLOAT;
+  else if (type == TypeDesc::TypeMatrix)
+    attr_map[index].type = NODE_ATTR_MATRIX;
+  else if (type == TypeFloat2)
+    attr_map[index].type = NODE_ATTR_FLOAT2;
+  else if (type == TypeFloat4)
+    attr_map[index].type = NODE_ATTR_FLOAT4;
+  else if (type == TypeRGBA)
+    attr_map[index].type = NODE_ATTR_RGBA;
+  else
+    attr_map[index].type = NODE_ATTR_FLOAT3;
+
+  attr_map[index].flags = desc.flags;
+}
+
+/* Generate an attribute map end marker, optionally including a link to another map.
+ * Links are used to connect object attribute maps to mesh attribute maps. */
+static void emit_attribute_map_terminator(AttributeMap *attr_map,
+                                          size_t index,
+                                          bool chain,
+                                          uint chain_link)
+{
+  for (int j = 0; j < ATTR_PRIM_TYPES; j++) {
+    attr_map[index + j].id = ATTR_STD_NONE;
+    attr_map[index + j].element = chain;                     /* link is valid flag */
+    attr_map[index + j].offset = chain ? chain_link + j : 0; /* link to the correct sub-entry */
+    attr_map[index + j].type = 0;
+    attr_map[index + j].flags = 0;
+  }
+}
+
+/* Generate all necessary attribute map entries from the attribute request. */
+static void emit_attribute_mapping(
+    AttributeMap *attr_map, size_t index, uint64_t id, AttributeRequest &req, Geometry *geom)
+{
+  emit_attribute_map_entry(attr_map, index, id, req.type, req.desc);
+
+  if (geom->is_mesh()) {
+    Mesh *mesh = static_cast<Mesh *>(geom);
+    if (mesh->get_num_subd_faces()) {
+      emit_attribute_map_entry(attr_map, index + 1, id, req.subd_type, req.subd_desc);
+    }
+  }
+}
+
+void GeometryManager::update_svm_attributes(Device *,
+                                            DeviceScene *dscene,
+                                            Scene *scene,
+                                            vector<AttributeRequestSet> &geom_attributes,
+                                            vector<AttributeRequestSet> &object_attributes)
+{
+  /* for SVM, the attributes_map table is used to lookup the offset of an
+   * attribute, based on a unique shader attribute id. */
+
+  /* compute array stride */
+  size_t attr_map_size = 0;
+
+  for (size_t i = 0; i < scene->geometry.size(); i++) {
+    Geometry *geom = scene->geometry[i];
+    geom->attr_map_offset = attr_map_size;
+
+#ifdef WITH_OSL
+    size_t attr_count = 0;
+    foreach (AttributeRequest &req, geom_attributes[i].requests) {
+      if (req.std != ATTR_STD_NONE &&
+          scene->shader_manager->get_attribute_id(req.std) != (uint64_t)req.std)
+        attr_count += 2;
+      else
+        attr_count += 1;
+    }
+#else
+    const size_t attr_count = geom_attributes[i].size();
+#endif
+
+    attr_map_size += (attr_count + 1) * ATTR_PRIM_TYPES;
+  }
+
+  for (size_t i = 0; i < scene->objects.size(); i++) {
+    Object *object = scene->objects[i];
+
+    /* only allocate a table for the object if it actually has attributes */
+    if (object_attributes[i].size() == 0) {
+      object->attr_map_offset = 0;
+    }
+    else {
+      object->attr_map_offset = attr_map_size;
+      attr_map_size += (object_attributes[i].size() + 1) * ATTR_PRIM_TYPES;
+    }
+  }
+
+  if (attr_map_size == 0)
+    return;
+
+  if (!dscene->attributes_map.need_realloc()) {
+    return;
+  }
+
+  /* create attribute map */
+  AttributeMap *attr_map = dscene->attributes_map.alloc(attr_map_size);
+  memset(attr_map, 0, dscene->attributes_map.size() * sizeof(*attr_map));
+
+  for (size_t i = 0; i < scene->geometry.size(); i++) {
+    Geometry *geom = scene->geometry[i];
+    AttributeRequestSet &attributes = geom_attributes[i];
+
+    /* set geometry attributes */
+    size_t index = geom->attr_map_offset;
+
+    foreach (AttributeRequest &req, attributes.requests) {
+      uint64_t id;
+      if (req.std == ATTR_STD_NONE)
+        id = scene->shader_manager->get_attribute_id(req.name);
+      else
+        id = scene->shader_manager->get_attribute_id(req.std);
+
+      emit_attribute_mapping(attr_map, index, id, req, geom);
+      index += ATTR_PRIM_TYPES;
+
+#ifdef WITH_OSL
+      /* Some standard attributes are explicitly referenced via their standard ID, so add those
+       * again in case they were added under a different attribute ID. */
+      if (req.std != ATTR_STD_NONE && id != (uint64_t)req.std) {
+        emit_attribute_mapping(attr_map, index, (uint64_t)req.std, req, geom);
+        index += ATTR_PRIM_TYPES;
+      }
+#endif
+    }
+
+    emit_attribute_map_terminator(attr_map, index, false, 0);
+  }
+
+  for (size_t i = 0; i < scene->objects.size(); i++) {
+    Object *object = scene->objects[i];
+    AttributeRequestSet &attributes = object_attributes[i];
+
+    /* set object attributes */
+    if (attributes.size() > 0) {
+      size_t index = object->attr_map_offset;
+
+      foreach (AttributeRequest &req, attributes.requests) {
+        uint64_t id;
+        if (req.std == ATTR_STD_NONE)
+          id = scene->shader_manager->get_attribute_id(req.name);
+        else
+          id = scene->shader_manager->get_attribute_id(req.std);
+
+        emit_attribute_mapping(attr_map, index, id, req, object->geometry);
+        index += ATTR_PRIM_TYPES;
+      }
+
+      emit_attribute_map_terminator(attr_map, index, true, object->geometry->attr_map_offset);
+    }
+  }
+
+  /* copy to device */
+  dscene->attributes_map.copy_to_device();
+}
+
+void GeometryManager::update_attribute_element_offset(Geometry *geom,
+                                                      device_vector<float> &attr_float,
+                                                      size_t &attr_float_offset,
+                                                      device_vector<float2> &attr_float2,
+                                                      size_t &attr_float2_offset,
+                                                      device_vector<packed_float3> &attr_float3,
+                                                      size_t &attr_float3_offset,
+                                                      device_vector<float4> &attr_float4,
+                                                      size_t &attr_float4_offset,
+                                                      device_vector<uchar4> &attr_uchar4,
+                                                      size_t &attr_uchar4_offset,
+                                                      Attribute *mattr,
+                                                      AttributePrimitive prim,
+                                                      TypeDesc &type,
+                                                      AttributeDescriptor &desc)
+{
+  if (mattr) {
+    /* store element and type */
+    desc.element = mattr->element;
+    desc.flags = mattr->flags;
+    type = mattr->type;
+
+    /* store attribute data in arrays */
+    size_t size = mattr->element_size(geom, prim);
+
+    AttributeElement &element = desc.element;
+    int &offset = desc.offset;
+
+    if (mattr->element == ATTR_ELEMENT_VOXEL) {
+      /* store slot in offset value */
+      ImageHandle &handle = mattr->data_voxel();
+      offset = handle.svm_slot();
+    }
+    else if (mattr->element == ATTR_ELEMENT_CORNER_BYTE) {
+      uchar4 *data = mattr->data_uchar4();
+      offset = attr_uchar4_offset;
+
+      assert(attr_uchar4.size() >= offset + size);
+      if (mattr->modified) {
+        for (size_t k = 0; k < size; k++) {
+          attr_uchar4[offset + k] = data[k];
+        }
+        attr_uchar4.tag_modified();
+      }
+      attr_uchar4_offset += size;
+    }
+    else if (mattr->type == TypeDesc::TypeFloat) {
+      float *data = mattr->data_float();
+      offset = attr_float_offset;
+
+      assert(attr_float.size() >= offset + size);
+      if (mattr->modified) {
+        for (size_t k = 0; k < size; k++) {
+          attr_float[offset + k] = data[k];
+        }
+        attr_float.tag_modified();
+      }
+      attr_float_offset += size;
+    }
+    else if (mattr->type == TypeFloat2) {
+      float2 *data = mattr->data_float2();
+      offset = attr_float2_offset;
+
+      assert(attr_float2.size() >= offset + size);
+      if (mattr->modified) {
+        for (size_t k = 0; k < size; k++) {
+          attr_float2[offset + k] = data[k];
+        }
+        attr_float2.tag_modified();
+      }
+      attr_float2_offset += size;
+    }
+    else if (mattr->type == TypeDesc::TypeMatrix) {
+      Transform *tfm = mattr->data_transform();
+      offset = attr_float4_offset;
+
+      assert(attr_float4.size() >= offset + size * 3);
+      if (mattr->modified) {
+        for (size_t k = 0; k < size * 3; k++) {
+          attr_float4[offset + k] = (&tfm->x)[k];
+        }
+        attr_float4.tag_modified();
+      }
+      attr_float4_offset += size * 3;
+    }
+    else if (mattr->type == TypeFloat4 || mattr->type == TypeRGBA) {
+      float4 *data = mattr->data_float4();
+      offset = attr_float4_offset;
+
+      assert(attr_float4.size() >= offset + size);
+      if (mattr->modified) {
+        for (size_t k = 0; k < size; k++) {
+          attr_float4[offset + k] = data[k];
+        }
+        attr_float4.tag_modified();
+      }
+      attr_float4_offset += size;
+    }
+    else {
+      float3 *data = mattr->data_float3();
+      offset = attr_float3_offset;
+
+      assert(attr_float3.size() >= offset + size);
+      if (mattr->modified) {
+        for (size_t k = 0; k < size; k++) {
+          attr_float3[offset + k] = data[k];
+        }
+        attr_float3.tag_modified();
+      }
+      attr_float3_offset += size;
+    }
+
+    /* mesh vertex/curve index is global, not per object, so we sneak
+     * a correction for that in here */
+    if (geom->is_mesh()) {
+      Mesh *mesh = static_cast<Mesh *>(geom);
+      if (mesh->subdivision_type == Mesh::SUBDIVISION_CATMULL_CLARK &&
+          desc.flags & ATTR_SUBDIVIDED) {
+        /* Indices for subdivided attributes are retrieved
+         * from patch table so no need for correction here. */
+      }
+      else if (element == ATTR_ELEMENT_VERTEX)
+        offset -= mesh->vert_offset;
+      else if (element == ATTR_ELEMENT_VERTEX_MOTION)
+        offset -= mesh->vert_offset;
+      else if (element == ATTR_ELEMENT_FACE) {
+        if (prim == ATTR_PRIM_GEOMETRY)
+          offset -= mesh->prim_offset;
+        else
+          offset -= mesh->face_offset;
+      }
+      else if (element == ATTR_ELEMENT_CORNER || element == ATTR_ELEMENT_CORNER_BYTE) {
+        if (prim == ATTR_PRIM_GEOMETRY)
+          offset -= 3 * mesh->prim_offset;
+        else
+          offset -= mesh->corner_offset;
+      }
+    }
+    else if (geom->is_hair()) {
+      Hair *hair = static_cast<Hair *>(geom);
+      if (element == ATTR_ELEMENT_CURVE)
+        offset -= hair->prim_offset;
+      else if (element == ATTR_ELEMENT_CURVE_KEY)
+        offset -= hair->curve_key_offset;
+      else if (element == ATTR_ELEMENT_CURVE_KEY_MOTION)
+        offset -= hair->curve_key_offset;
+    }
+    else if (geom->is_pointcloud()) {
+      if (element == ATTR_ELEMENT_VERTEX)
+        offset -= geom->prim_offset;
+      else if (element == ATTR_ELEMENT_VERTEX_MOTION)
+        offset -= geom->prim_offset;
+    }
+  }
+  else {
+    /* attribute not found */
+    desc.element = ATTR_ELEMENT_NONE;
+    desc.offset = 0;
+  }
+}
+
+static void update_attribute_element_size(Geometry *geom,
+                                          Attribute *mattr,
+                                          AttributePrimitive prim,
+                                          size_t *attr_float_size,
+                                          size_t *attr_float2_size,
+                                          size_t *attr_float3_size,
+                                          size_t *attr_float4_size,
+                                          size_t *attr_uchar4_size)
+{
+  if (mattr) {
+    size_t size = mattr->element_size(geom, prim);
+
+    if (mattr->element == ATTR_ELEMENT_VOXEL) {
+      /* pass */
+    }
+    else if (mattr->element == ATTR_ELEMENT_CORNER_BYTE) {
+      *attr_uchar4_size += size;
+    }
+    else if (mattr->type == TypeDesc::TypeFloat) {
+      *attr_float_size += size;
+    }
+    else if (mattr->type == TypeFloat2) {
+      *attr_float2_size += size;
+    }
+    else if (mattr->type == TypeDesc::TypeMatrix) {
+      *attr_float4_size += size * 4;
+    }
+    else if (mattr->type == TypeFloat4 || mattr->type == TypeRGBA) {
+      *attr_float4_size += size;
+    }
+    else {
+      *attr_float3_size += size;
+    }
+  }
+}
+
+void GeometryManager::device_update_attributes(Device *device,
+                                               DeviceScene *dscene,
+                                               Scene *scene,
+                                               Progress &progress)
+{
+  progress.set_status("Updating Mesh", "Computing attributes");
+
+  /* gather per mesh requested attributes. as meshes may have multiple
+   * shaders assigned, this merges the requested attributes that have
+   * been set per shader by the shader manager */
+  vector<AttributeRequestSet> geom_attributes(scene->geometry.size());
+
+  for (size_t i = 0; i < scene->geometry.size(); i++) {
+    Geometry *geom = scene->geometry[i];
+
+    geom->index = i;
+    scene->need_global_attributes(geom_attributes[i]);
+
+    foreach (Node *node, geom->get_used_shaders()) {
+      Shader *shader = static_cast<Shader *>(node);
+      geom_attributes[i].add(shader->attributes);
+    }
+
+    if (geom->is_hair() && static_cast<Hair *>(geom)->need_shadow_transparency()) {
+      geom_attributes[i].add(ATTR_STD_SHADOW_TRANSPARENCY);
+    }
+  }
+
+  /* convert object attributes to use the same data structures as geometry ones */
+  vector<AttributeRequestSet> object_attributes(scene->objects.size());
+  vector<AttributeSet> object_attribute_values;
+
+  object_attribute_values.reserve(scene->objects.size());
+
+  for (size_t i = 0; i < scene->objects.size(); i++) {
+    Object *object = scene->objects[i];
+    Geometry *geom = object->geometry;
+    size_t geom_idx = geom->index;
+
+    assert(geom_idx < scene->geometry.size() && scene->geometry[geom_idx] == geom);
+
+    object_attribute_values.push_back(AttributeSet(geom, ATTR_PRIM_GEOMETRY));
+
+    AttributeRequestSet &geom_requests = geom_attributes[geom_idx];
+    AttributeRequestSet &attributes = object_attributes[i];
+    AttributeSet &values = object_attribute_values[i];
+
+    for (size_t j = 0; j < object->attributes.size(); j++) {
+      ParamValue &param = object->attributes[j];
+
+      /* add attributes that are requested and not already handled by the mesh */
+      if (geom_requests.find(param.name()) && !geom->attributes.find(param.name())) {
+        attributes.add(param.name());
+
+        Attribute *attr = values.add(param.name(), param.type(), ATTR_ELEMENT_OBJECT);
+        assert(param.datasize() == attr->buffer.size());
+        memcpy(attr->buffer.data(), param.data(), param.datasize());
+      }
+    }
+  }
+
+  /* mesh attribute are stored in a single array per data type. here we fill
+   * those arrays, and set the offset and element type to create attribute
+   * maps next */
+
+  /* Pre-allocate attributes to avoid arrays re-allocation which would
+   * take 2x of overall attribute memory usage.
+   */
+  size_t attr_float_size = 0;
+  size_t attr_float2_size = 0;
+  size_t attr_float3_size = 0;
+  size_t attr_float4_size = 0;
+  size_t attr_uchar4_size = 0;
+
+  for (size_t i = 0; i < scene->geometry.size(); i++) {
+    Geometry *geom = scene->geometry[i];
+    AttributeRequestSet &attributes = geom_attributes[i];
+    foreach (AttributeRequest &req, attributes.requests) {
+      Attribute *attr = geom->attributes.find(req);
+
+      update_attribute_element_size(geom,
+                                    attr,
+                                    ATTR_PRIM_GEOMETRY,
+                                    &attr_float_size,
+                                    &attr_float2_size,
+                                    &attr_float3_size,
+                                    &attr_float4_size,
+                                    &attr_uchar4_size);
+
+      if (geom->is_mesh()) {
+        Mesh *mesh = static_cast<Mesh *>(geom);
+        Attribute *subd_attr = mesh->subd_attributes.find(req);
+
+        update_attribute_element_size(mesh,
+                                      subd_attr,
+                                      ATTR_PRIM_SUBD,
+                                      &attr_float_size,
+                                      &attr_float2_size,
+                                      &attr_float3_size,
+                                      &attr_float4_size,
+                                      &attr_uchar4_size);
+      }
+    }
+  }
+
+  for (size_t i = 0; i < scene->objects.size(); i++) {
+    Object *object = scene->objects[i];
+
+    foreach (Attribute &attr, object_attribute_values[i].attributes) {
+      update_attribute_element_size(object->geometry,
+                                    &attr,
+                                    ATTR_PRIM_GEOMETRY,
+                                    &attr_float_size,
+                                    &attr_float2_size,
+                                    &attr_float3_size,
+                                    &attr_float4_size,
+                                    &attr_uchar4_size);
+    }
+  }
+
+  dscene->attributes_float.alloc(attr_float_size);
+  dscene->attributes_float2.alloc(attr_float2_size);
+  dscene->attributes_float3.alloc(attr_float3_size);
+  dscene->attributes_float4.alloc(attr_float4_size);
+  dscene->attributes_uchar4.alloc(attr_uchar4_size);
+
+  /* The order of those flags needs to match that of AttrKernelDataType. */
+  const bool attributes_need_realloc[AttrKernelDataType::NUM] = {
+      dscene->attributes_float.need_realloc(),
+      dscene->attributes_float2.need_realloc(),
+      dscene->attributes_float3.need_realloc(),
+      dscene->attributes_float4.need_realloc(),
+      dscene->attributes_uchar4.need_realloc(),
+  };
+
+  size_t attr_float_offset = 0;
+  size_t attr_float2_offset = 0;
+  size_t attr_float3_offset = 0;
+  size_t attr_float4_offset = 0;
+  size_t attr_uchar4_offset = 0;
+
+  /* Fill in attributes. */
+  for (size_t i = 0; i < scene->geometry.size(); i++) {
+    Geometry *geom = scene->geometry[i];
+    AttributeRequestSet &attributes = geom_attributes[i];
+
+    /* todo: we now store std and name attributes from requests even if
+     * they actually refer to the same mesh attributes, optimize */
+    foreach (AttributeRequest &req, attributes.requests) {
+      Attribute *attr = geom->attributes.find(req);
+
+      if (attr) {
+        /* force a copy if we need to reallocate all the data */
+        attr->modified |= attributes_need_realloc[Attribute::kernel_type(*attr)];
+      }
+
+      update_attribute_element_offset(geom,
+                                      dscene->attributes_float,
+                                      attr_float_offset,
+                                      dscene->attributes_float2,
+                                      attr_float2_offset,
+                                      dscene->attributes_float3,
+                                      attr_float3_offset,
+                                      dscene->attributes_float4,
+                                      attr_float4_offset,
+                                      dscene->attributes_uchar4,
+                                      attr_uchar4_offset,
+                                      attr,
+                                      ATTR_PRIM_GEOMETRY,
+                                      req.type,
+                                      req.desc);
+
+      if (geom->is_mesh()) {
+        Mesh *mesh = static_cast<Mesh *>(geom);
+        Attribute *subd_attr = mesh->subd_attributes.find(req);
+
+        if (subd_attr) {
+          /* force a copy if we need to reallocate all the data */
+          subd_attr->modified |= attributes_need_realloc[Attribute::kernel_type(*subd_attr)];
+        }
+
+        update_attribute_element_offset(mesh,
+                                        dscene->attributes_float,
+                                        attr_float_offset,
+                                        dscene->attributes_float2,
+                                        attr_float2_offset,
+                                        dscene->attributes_float3,
+                                        attr_float3_offset,
+                                        dscene->attributes_float4,
+                                        attr_float4_offset,
+                                        dscene->attributes_uchar4,
+                                        attr_uchar4_offset,
+                                        subd_attr,
+                                        ATTR_PRIM_SUBD,
+                                        req.subd_type,
+                                        req.subd_desc);
+      }
+
+      if (progress.get_cancel())
+        return;
+    }
+  }
+
+  for (size_t i = 0; i < scene->objects.size(); i++) {
+    Object *object = scene->objects[i];
+    AttributeRequestSet &attributes = object_attributes[i];
+    AttributeSet &values = object_attribute_values[i];
+
+    foreach (AttributeRequest &req, attributes.requests) {
+      Attribute *attr = values.find(req);
+
+      if (attr) {
+        attr->modified |= attributes_need_realloc[Attribute::kernel_type(*attr)];
+      }
+
+      update_attribute_element_offset(object->geometry,
+                                      dscene->attributes_float,
+                                      attr_float_offset,
+                                      dscene->attributes_float2,
+                                      attr_float2_offset,
+                                      dscene->attributes_float3,
+                                      attr_float3_offset,
+                                      dscene->attributes_float4,
+                                      attr_float4_offset,
+                                      dscene->attributes_uchar4,
+                                      attr_uchar4_offset,
+                                      attr,
+                                      ATTR_PRIM_GEOMETRY,
+                                      req.type,
+                                      req.desc);
+
+      /* object attributes don't care about subdivision */
+      req.subd_type = req.type;
+      req.subd_desc = req.desc;
+
+      if (progress.get_cancel())
+        return;
+    }
+  }
+
+  /* create attribute lookup maps */
+  if (scene->shader_manager->use_osl())
+    update_osl_globals(device, scene);
+
+  update_svm_attributes(device, dscene, scene, geom_attributes, object_attributes);
+
+  if (progress.get_cancel())
+    return;
+
+  /* copy to device */
+  progress.set_status("Updating Mesh", "Copying Attributes to device");
+
+  dscene->attributes_float.copy_to_device_if_modified();
+  dscene->attributes_float2.copy_to_device_if_modified();
+  dscene->attributes_float3.copy_to_device_if_modified();
+  dscene->attributes_float4.copy_to_device_if_modified();
+  dscene->attributes_uchar4.copy_to_device_if_modified();
+
+  if (progress.get_cancel())
+    return;
+
+  /* After mesh attributes and patch tables have been copied to device memory,
+   * we need to update offsets in the objects. */
+  scene->object_manager->device_update_geom_offsets(device, dscene, scene);
+}
+
+CCL_NAMESPACE_END

intern/cycles/scene/geometry_bvh.cpp

@@ -0,0 +1,196 @@
+/* SPDX-License-Identifier: Apache-2.0
+ * Copyright 2011-2022 Blender Foundation */
+
+#include "bvh/bvh.h"
+#include "bvh/bvh2.h"
+
+#include "device/device.h"
+
+#include "scene/attribute.h"
+#include "scene/camera.h"
+#include "scene/geometry.h"
+#include "scene/hair.h"
+#include "scene/light.h"
+#include "scene/mesh.h"
+#include "scene/object.h"
+#include "scene/pointcloud.h"
+#include "scene/scene.h"
+#include "scene/shader.h"
+#include "scene/shader_nodes.h"
+#include "scene/stats.h"
+#include "scene/volume.h"
+
+#include "subd/patch_table.h"
+#include "subd/split.h"
+
+#include "kernel/osl/globals.h"
+
+#include "util/foreach.h"
+#include "util/log.h"
+#include "util/progress.h"
+#include "util/task.h"
+
+CCL_NAMESPACE_BEGIN
+
+void Geometry::compute_bvh(Device *device,
+                           DeviceScene *dscene,
+                           SceneParams *params,
+                           Progress *progress,
+                           size_t n,
+                           size_t total)
+{
+  if (progress->get_cancel())
+    return;
+
+  compute_bounds();
+
+  const BVHLayout bvh_layout = BVHParams::best_bvh_layout(
+      params->bvh_layout, device->get_bvh_layout_mask(dscene->data.kernel_features));
+  if (need_build_bvh(bvh_layout)) {
+    string msg = "Updating Geometry BVH ";
+    if (name.empty())
+      msg += string_printf("%u/%u", (uint)(n + 1), (uint)total);
+    else
+      msg += string_printf("%s %u/%u", name.c_str(), (uint)(n + 1), (uint)total);
+
+    Object object;
+
+    /* Ensure all visibility bits are set at the geometry level BVH. In
+     * the object level BVH is where actual visibility is tested. */
+    object.set_is_shadow_catcher(true);
+    object.set_visibility(~0);
+
+    object.set_geometry(this);
+
+    vector<Geometry *> geometry;
+    geometry.push_back(this);
+    vector<Object *> objects;
+    objects.push_back(&object);
+
+    if (bvh && !need_update_rebuild) {
+      progress->set_status(msg, "Refitting BVH");
+
+      bvh->replace_geometry(geometry, objects);
+
+      device->build_bvh(bvh, *progress, true);
+    }
+    else {
+      progress->set_status(msg, "Building BVH");
+
+      BVHParams bparams;
+      bparams.use_spatial_split = params->use_bvh_spatial_split;
+      bparams.use_compact_structure = params->use_bvh_compact_structure;
+      bparams.bvh_layout = bvh_layout;
+      bparams.use_unaligned_nodes = dscene->data.bvh.have_curves &&
+                                    params->use_bvh_unaligned_nodes;
+      bparams.num_motion_triangle_steps = params->num_bvh_time_steps;
+      bparams.num_motion_curve_steps = params->num_bvh_time_steps;
+      bparams.num_motion_point_steps = params->num_bvh_time_steps;
+      bparams.bvh_type = params->bvh_type;
+      bparams.curve_subdivisions = params->curve_subdivisions();
+
+      delete bvh;
+      bvh = BVH::create(bparams, geometry, objects, device);
+      MEM_GUARDED_CALL(progress, device->build_bvh, bvh, *progress, false);
+    }
+  }
+
+  need_update_rebuild = false;
+  need_update_bvh_for_offset = false;
+}
+
+void GeometryManager::device_update_bvh(Device *device,
+                                        DeviceScene *dscene,
+                                        Scene *scene,
+                                        Progress &progress)
+{
+  /* bvh build */
+  progress.set_status("Updating Scene BVH", "Building");
+
+  BVHParams bparams;
+  bparams.top_level = true;
+  bparams.bvh_layout = BVHParams::best_bvh_layout(
+      scene->params.bvh_layout, device->get_bvh_layout_mask(dscene->data.kernel_features));
+  bparams.use_spatial_split = scene->params.use_bvh_spatial_split;
+  bparams.use_unaligned_nodes = dscene->data.bvh.have_curves &&
+                                scene->params.use_bvh_unaligned_nodes;
+  bparams.num_motion_triangle_steps = scene->params.num_bvh_time_steps;
+  bparams.num_motion_curve_steps = scene->params.num_bvh_time_steps;
+  bparams.num_motion_point_steps = scene->params.num_bvh_time_steps;
+  bparams.bvh_type = scene->params.bvh_type;
+  bparams.curve_subdivisions = scene->params.curve_subdivisions();
+
+  VLOG_INFO << "Using " << bvh_layout_name(bparams.bvh_layout) << " layout.";
+
+  const bool can_refit = scene->bvh != nullptr &&
+                         (bparams.bvh_layout == BVHLayout::BVH_LAYOUT_OPTIX ||
+                          bparams.bvh_layout == BVHLayout::BVH_LAYOUT_METAL);
+
+  BVH *bvh = scene->bvh;
+  if (!scene->bvh) {
+    bvh = scene->bvh = BVH::create(bparams, scene->geometry, scene->objects, device);
+  }
+
+  device->build_bvh(bvh, progress, can_refit);
+
+  if (progress.get_cancel()) {
+    return;
+  }
+
+  const bool has_bvh2_layout = (bparams.bvh_layout == BVH_LAYOUT_BVH2);
+
+  PackedBVH pack;
+  if (has_bvh2_layout) {
+    pack = std::move(static_cast<BVH2 *>(bvh)->pack);
+  }
+  else {
+    pack.root_index = -1;
+  }
+
+  /* copy to device */
+  progress.set_status("Updating Scene BVH", "Copying BVH to device");
+
+  /* When using BVH2, we always have to copy/update the data as its layout is dependent on the
+   * BVH's leaf nodes which may be different when the objects or vertices move. */
+
+  if (pack.nodes.size()) {
+    dscene->bvh_nodes.steal_data(pack.nodes);
+    dscene->bvh_nodes.copy_to_device();
+  }
+  if (pack.leaf_nodes.size()) {
+    dscene->bvh_leaf_nodes.steal_data(pack.leaf_nodes);
+    dscene->bvh_leaf_nodes.copy_to_device();
+  }
+  if (pack.object_node.size()) {
+    dscene->object_node.steal_data(pack.object_node);
+    dscene->object_node.copy_to_device();
+  }
+  if (pack.prim_type.size()) {
+    dscene->prim_type.steal_data(pack.prim_type);
+    dscene->prim_type.copy_to_device();
+  }
+  if (pack.prim_visibility.size()) {
+    dscene->prim_visibility.steal_data(pack.prim_visibility);
+    dscene->prim_visibility.copy_to_device();
+  }
+  if (pack.prim_index.size()) {
+    dscene->prim_index.steal_data(pack.prim_index);
+    dscene->prim_index.copy_to_device();
+  }
+  if (pack.prim_object.size()) {
+    dscene->prim_object.steal_data(pack.prim_object);
+    dscene->prim_object.copy_to_device();
+  }
+  if (pack.prim_time.size()) {
+    dscene->prim_time.steal_data(pack.prim_time);
+    dscene->prim_time.copy_to_device();
+  }
+
+  dscene->data.bvh.root = pack.root_index;
+  dscene->data.bvh.use_bvh_steps = (scene->params.num_bvh_time_steps != 0);
+  dscene->data.bvh.curve_subdivisions = scene->params.curve_subdivisions();
+  /* The scene handle is set in 'CPUDevice::const_copy_to' and 'OptiXDevice::const_copy_to' */
+  dscene->data.device_bvh = 0;
+}
+
+CCL_NAMESPACE_END

intern/cycles/scene/geometry_mesh.cpp

@@ -0,0 +1,223 @@
+/* SPDX-License-Identifier: Apache-2.0
+ * Copyright 2011-2022 Blender Foundation */
+
+#include "bvh/bvh.h"
+#include "bvh/bvh2.h"
+
+#include "device/device.h"
+
+#include "scene/attribute.h"
+#include "scene/camera.h"
+#include "scene/geometry.h"
+#include "scene/hair.h"
+#include "scene/light.h"
+#include "scene/mesh.h"
+#include "scene/object.h"
+#include "scene/osl.h"
+#include "scene/pointcloud.h"
+#include "scene/scene.h"
+#include "scene/shader.h"
+#include "scene/shader_nodes.h"
+#include "scene/stats.h"
+#include "scene/volume.h"
+
+#include "subd/patch_table.h"
+#include "subd/split.h"
+
+#ifdef WITH_OSL
+#  include "kernel/osl/globals.h"
+#endif
+
+#include "util/foreach.h"
+#include "util/log.h"
+#include "util/progress.h"
+#include "util/task.h"
+
+CCL_NAMESPACE_BEGIN
+
+void GeometryManager::device_update_mesh(Device *,
+                                         DeviceScene *dscene,
+                                         Scene *scene,
+                                         Progress &progress)
+{
+  /* Count. */
+  size_t vert_size = 0;
+  size_t tri_size = 0;
+
+  size_t curve_key_size = 0;
+  size_t curve_size = 0;
+  size_t curve_segment_size = 0;
+
+  size_t point_size = 0;
+
+  size_t patch_size = 0;
+
+  foreach (Geometry *geom, scene->geometry) {
+    if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
+      Mesh *mesh = static_cast<Mesh *>(geom);
+
+      vert_size += mesh->verts.size();
+      tri_size += mesh->num_triangles();
+
+      if (mesh->get_num_subd_faces()) {
+        Mesh::SubdFace last = mesh->get_subd_face(mesh->get_num_subd_faces() - 1);
+        patch_size += (last.ptex_offset + last.num_ptex_faces()) * 8;
+
+        /* patch tables are stored in same array so include them in patch_size */
+        if (mesh->patch_table) {
+          mesh->patch_table_offset = patch_size;
+          patch_size += mesh->patch_table->total_size();
+        }
+      }
+    }
+    else if (geom->is_hair()) {
+      Hair *hair = static_cast<Hair *>(geom);
+
+      curve_key_size += hair->get_curve_keys().size();
+      curve_size += hair->num_curves();
+      curve_segment_size += hair->num_segments();
+    }
+    else if (geom->is_pointcloud()) {
+      PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+      point_size += pointcloud->num_points();
+    }
+  }
+
+  /* Fill in all the arrays. */
+  if (tri_size != 0) {
+    /* normals */
+    progress.set_status("Updating Mesh", "Computing normals");
+
+    packed_float3 *tri_verts = dscene->tri_verts.alloc(vert_size);
+    uint *tri_shader = dscene->tri_shader.alloc(tri_size);
+    packed_float3 *vnormal = dscene->tri_vnormal.alloc(vert_size);
+    packed_uint3 *tri_vindex = dscene->tri_vindex.alloc(tri_size);
+    uint *tri_patch = dscene->tri_patch.alloc(tri_size);
+    float2 *tri_patch_uv = dscene->tri_patch_uv.alloc(vert_size);
+
+    const bool copy_all_data = dscene->tri_shader.need_realloc() ||
+                               dscene->tri_vindex.need_realloc() ||
+                               dscene->tri_vnormal.need_realloc() ||
+                               dscene->tri_patch.need_realloc() ||
+                               dscene->tri_patch_uv.need_realloc();
+
+    foreach (Geometry *geom, scene->geometry) {
+      if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
+        Mesh *mesh = static_cast<Mesh *>(geom);
+
+        if (mesh->shader_is_modified() || mesh->smooth_is_modified() ||
+            mesh->triangles_is_modified() || copy_all_data) {
+          mesh->pack_shaders(scene, &tri_shader[mesh->prim_offset]);
+        }
+
+        if (mesh->verts_is_modified() || copy_all_data) {
+          mesh->pack_normals(&vnormal[mesh->vert_offset]);
+        }
+
+        if (mesh->verts_is_modified() || mesh->triangles_is_modified() ||
+            mesh->vert_patch_uv_is_modified() || copy_all_data) {
+          mesh->pack_verts(&tri_verts[mesh->vert_offset],
+                           &tri_vindex[mesh->prim_offset],
+                           &tri_patch[mesh->prim_offset],
+                           &tri_patch_uv[mesh->vert_offset]);
+        }
+
+        if (progress.get_cancel())
+          return;
+      }
+    }
+
+    /* vertex coordinates */
+    progress.set_status("Updating Mesh", "Copying Mesh to device");
+
+    dscene->tri_verts.copy_to_device_if_modified();
+    dscene->tri_shader.copy_to_device_if_modified();
+    dscene->tri_vnormal.copy_to_device_if_modified();
+    dscene->tri_vindex.copy_to_device_if_modified();
+    dscene->tri_patch.copy_to_device_if_modified();
+    dscene->tri_patch_uv.copy_to_device_if_modified();
+  }
+
+  if (curve_segment_size != 0) {
+    progress.set_status("Updating Mesh", "Copying Curves to device");
+
+    float4 *curve_keys = dscene->curve_keys.alloc(curve_key_size);
+    KernelCurve *curves = dscene->curves.alloc(curve_size);
+    KernelCurveSegment *curve_segments = dscene->curve_segments.alloc(curve_segment_size);
+
+    const bool copy_all_data = dscene->curve_keys.need_realloc() ||
+                               dscene->curves.need_realloc() ||
+                               dscene->curve_segments.need_realloc();
+
+    foreach (Geometry *geom, scene->geometry) {
+      if (geom->is_hair()) {
+        Hair *hair = static_cast<Hair *>(geom);
+
+        bool curve_keys_co_modified = hair->curve_radius_is_modified() ||
+                                      hair->curve_keys_is_modified();
+        bool curve_data_modified = hair->curve_shader_is_modified() ||
+                                   hair->curve_first_key_is_modified();
+
+        if (!curve_keys_co_modified && !curve_data_modified && !copy_all_data) {
+          continue;
+        }
+
+        hair->pack_curves(scene,
+                          &curve_keys[hair->curve_key_offset],
+                          &curves[hair->prim_offset],
+                          &curve_segments[hair->curve_segment_offset]);
+        if (progress.get_cancel())
+          return;
+      }
+    }
+
+    dscene->curve_keys.copy_to_device_if_modified();
+    dscene->curves.copy_to_device_if_modified();
+    dscene->curve_segments.copy_to_device_if_modified();
+  }
+
+  if (point_size != 0) {
+    progress.set_status("Updating Mesh", "Copying Point clouds to device");
+
+    float4 *points = dscene->points.alloc(point_size);
+    uint *points_shader = dscene->points_shader.alloc(point_size);
+
+    foreach (Geometry *geom, scene->geometry) {
+      if (geom->is_pointcloud()) {
+        PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+        pointcloud->pack(
+            scene, &points[pointcloud->prim_offset], &points_shader[pointcloud->prim_offset]);
+        if (progress.get_cancel())
+          return;
+      }
+    }
+
+    dscene->points.copy_to_device();
+    dscene->points_shader.copy_to_device();
+  }
+
+  if (patch_size != 0 && dscene->patches.need_realloc()) {
+    progress.set_status("Updating Mesh", "Copying Patches to device");
+
+    uint *patch_data = dscene->patches.alloc(patch_size);
+
+    foreach (Geometry *geom, scene->geometry) {
+      if (geom->is_mesh()) {
+        Mesh *mesh = static_cast<Mesh *>(geom);
+        mesh->pack_patches(&patch_data[mesh->patch_offset]);
+
+        if (mesh->patch_table) {
+          mesh->patch_table->copy_adjusting_offsets(&patch_data[mesh->patch_table_offset],
+                                                    mesh->patch_table_offset);
+        }
+
+        if (progress.get_cancel())
+          return;
+      }
+    }
+
+    dscene->patches.copy_to_device();
+  }
+}
+
+CCL_NAMESPACE_END

intern/cycles/scene/scene.cpp

@@ -24,6 +24,7 @@
 #include "scene/svm.h"
 #include "scene/tables.h"
 #include "scene/volume.h"
+#include "scene/devicescene.h"
 #include "session/session.h"
 
 #include "util/foreach.h"
@@ -33,59 +34,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-DeviceScene::DeviceScene(Device *device)
+
-    : bvh_nodes(device, "bvh_nodes", MEM_GLOBAL),
-      bvh_leaf_nodes(device, "bvh_leaf_nodes", MEM_GLOBAL),
-      object_node(device, "object_node", MEM_GLOBAL),
-      prim_type(device, "prim_type", MEM_GLOBAL),
-      prim_visibility(device, "prim_visibility", MEM_GLOBAL),
-      prim_index(device, "prim_index", MEM_GLOBAL),
-      prim_object(device, "prim_object", MEM_GLOBAL),
-      prim_time(device, "prim_time", MEM_GLOBAL),
-      tri_verts(device, "tri_verts", MEM_GLOBAL),
-      tri_shader(device, "tri_shader", MEM_GLOBAL),
-      tri_vnormal(device, "tri_vnormal", MEM_GLOBAL),
-      tri_vindex(device, "tri_vindex", MEM_GLOBAL),
-      tri_patch(device, "tri_patch", MEM_GLOBAL),
-      tri_patch_uv(device, "tri_patch_uv", MEM_GLOBAL),
-      curves(device, "curves", MEM_GLOBAL),
-      curve_keys(device, "curve_keys", MEM_GLOBAL),
-      curve_segments(device, "curve_segments", MEM_GLOBAL),
-      patches(device, "patches", MEM_GLOBAL),
-      points(device, "points", MEM_GLOBAL),
-      points_shader(device, "points_shader", MEM_GLOBAL),
-      objects(device, "objects", MEM_GLOBAL),
-      object_motion_pass(device, "object_motion_pass", MEM_GLOBAL),
-      object_motion(device, "object_motion", MEM_GLOBAL),
-      object_flag(device, "object_flag", MEM_GLOBAL),
-      object_volume_step(device, "object_volume_step", MEM_GLOBAL),
-      object_prim_offset(device, "object_prim_offset", MEM_GLOBAL),
-      camera_motion(device, "camera_motion", MEM_GLOBAL),
-      attributes_map(device, "attributes_map", MEM_GLOBAL),
-      attributes_float(device, "attributes_float", MEM_GLOBAL),
-      attributes_float2(device, "attributes_float2", MEM_GLOBAL),
-      attributes_float3(device, "attributes_float3", MEM_GLOBAL),
-      attributes_float4(device, "attributes_float4", MEM_GLOBAL),
-      attributes_uchar4(device, "attributes_uchar4", MEM_GLOBAL),
-      light_distribution(device, "light_distribution", MEM_GLOBAL),
-      lights(device, "lights", MEM_GLOBAL),
-      light_background_marginal_cdf(device, "light_background_marginal_cdf", MEM_GLOBAL),
-      light_background_conditional_cdf(device, "light_background_conditional_cdf", MEM_GLOBAL),
-      light_tree_nodes(device, "light_tree_nodes", MEM_GLOBAL),
-      light_tree_emitters(device, "light_tree_emitters", MEM_GLOBAL),
-      light_to_tree(device, "light_to_tree", MEM_GLOBAL),
-      object_to_tree(device, "object_to_tree", MEM_GLOBAL),
-      object_lookup_offset(device, "object_lookup_offset", MEM_GLOBAL),
-      triangle_to_tree(device, "triangle_to_tree", MEM_GLOBAL),
-      particles(device, "particles", MEM_GLOBAL),
-      svm_nodes(device, "svm_nodes", MEM_GLOBAL),
-      shaders(device, "shaders", MEM_GLOBAL),
-      lookup_table(device, "lookup_table", MEM_GLOBAL),
-      sample_pattern_lut(device, "sample_pattern_lut", MEM_GLOBAL),
-      ies_lights(device, "ies", MEM_GLOBAL)
-{
-  memset((void *)&data, 0, sizeof(data));
-}
 
 Scene::Scene(const SceneParams &params_, Device *device)
     : name("Scene"),

intern/cycles/scene/scene.h

@@ -6,20 +6,16 @@
 
 #include "bvh/params.h"
 
+#include "scene/devicescene.h"
 #include "scene/film.h"
 #include "scene/image.h"
 #include "scene/shader.h"
 
-#include "device/device.h"
-#include "device/memory.h"
-
 #include "util/param.h"
 #include "util/string.h"
 #include "util/system.h"
 #include "util/texture.h"
 #include "util/thread.h"
-#include "util/types.h"
-#include "util/vector.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -54,92 +50,6 @@ class RenderStats;
 class SceneUpdateStats;
 class Volume;
 
-/* Scene Device Data */
-
-class DeviceScene {
- public:
-  /* BVH */
-  device_vector<int4> bvh_nodes;
-  device_vector<int4> bvh_leaf_nodes;
-  device_vector<int> object_node;
-  device_vector<int> prim_type;
-  device_vector<uint> prim_visibility;
-  device_vector<int> prim_index;
-  device_vector<int> prim_object;
-  device_vector<float2> prim_time;
-
-  /* mesh */
-  device_vector<packed_float3> tri_verts;
-  device_vector<uint> tri_shader;
-  device_vector<packed_float3> tri_vnormal;
-  device_vector<packed_uint3> tri_vindex;
-  device_vector<uint> tri_patch;
-  device_vector<float2> tri_patch_uv;
-
-  device_vector<KernelCurve> curves;
-  device_vector<float4> curve_keys;
-  device_vector<KernelCurveSegment> curve_segments;
-
-  device_vector<uint> patches;
-
-  /* point-cloud */
-  device_vector<float4> points;
-  device_vector<uint> points_shader;
-
-  /* objects */
-  device_vector<KernelObject> objects;
-  device_vector<Transform> object_motion_pass;
-  device_vector<DecomposedTransform> object_motion;
-  device_vector<uint> object_flag;
-  device_vector<float> object_volume_step;
-  device_vector<uint> object_prim_offset;
-
-  /* cameras */
-  device_vector<DecomposedTransform> camera_motion;
-
-  /* attributes */
-  device_vector<AttributeMap> attributes_map;
-  device_vector<float> attributes_float;
-  device_vector<float2> attributes_float2;
-  device_vector<packed_float3> attributes_float3;
-  device_vector<float4> attributes_float4;
-  device_vector<uchar4> attributes_uchar4;
-
-  /* lights */
-  device_vector<KernelLightDistribution> light_distribution;
-  device_vector<KernelLight> lights;
-  device_vector<float2> light_background_marginal_cdf;
-  device_vector<float2> light_background_conditional_cdf;
-
-  /* light tree */
-  device_vector<KernelLightTreeNode> light_tree_nodes;
-  device_vector<KernelLightTreeEmitter> light_tree_emitters;
-  device_vector<uint> light_to_tree;
-  device_vector<uint> object_to_tree;
-  device_vector<uint> object_lookup_offset;
-  device_vector<uint> triangle_to_tree;
-
-  /* particles */
-  device_vector<KernelParticle> particles;
-
-  /* shaders */
-  device_vector<int4> svm_nodes;
-  device_vector<KernelShader> shaders;
-
-  /* lookup tables */
-  device_vector<float> lookup_table;
-
-  /* integrator */
-  device_vector<float> sample_pattern_lut;
-
-  /* IES lights */
-  device_vector<float> ies_lights;
-
-  KernelData data;
-
-  DeviceScene(Device *device);
-};
-
 /* Scene Parameters */
 
 class SceneParams {

intern/ghost/GHOST_Path-api.hh

@@ -61,7 +61,7 @@ extern const char *GHOST_getBinaryDir(void);
 /**
  * Add the file to the operating system most recently used files
  */
-extern void GHOST_addToSystemRecentFiles(const char *filename);
+extern void GHOST_addToSystemRecentFiles(const char *filepath);
 
 #ifdef __cplusplus
 }

intern/ghost/intern/GHOST_IXrGraphicsBinding.hh

@@ -57,4 +57,4 @@ class GHOST_IXrGraphicsBinding {
 };
 
 std::unique_ptr<GHOST_IXrGraphicsBinding> GHOST_XrGraphicsBindingCreateFromType(
-    GHOST_TXrGraphicsBinding type, GHOST_Context &ghost_ctx);
+    GHOST_TXrGraphicsBinding type, GHOST_Context &context);

intern/ghost/intern/GHOST_SystemWayland.cc

@@ -5586,8 +5586,6 @@ GHOST_SystemWayland::GHOST_SystemWayland(bool background)
 #  endif
       display_destroy_and_free_all();
       throw std::runtime_error("Wayland: unable to find libdecor!");
-
-      use_libdecor = true;
     }
   }
   else {
@@ -5610,7 +5608,7 @@ GHOST_SystemWayland::GHOST_SystemWayland(bool background)
   (void)background;
 #endif
   {
-    GWL_XDG_Decor_System &decor = *display_->xdg_decor;
+    const GWL_XDG_Decor_System &decor = *display_->xdg_decor;
     if (!decor.shell) {
       display_destroy_and_free_all();
       throw std::runtime_error("Wayland: unable to access xdg_shell!");
@@ -6071,11 +6069,9 @@ static GHOST_TSuccess getCursorPositionClientRelative_impl(
     /* As the cursor is restored at the warped location,
      * apply warping when requesting the cursor location. */
     GHOST_Rect wrap_bounds{};
-    if (win->getCursorGrabModeIsWarp()) {
     if (win->getCursorGrabBounds(wrap_bounds) == GHOST_kFailure) {
       win->getClientBounds(wrap_bounds);
     }
-    }
     int xy_wrap[2] = {
         seat_state_pointer->xy[0],
         seat_state_pointer->xy[1],
@@ -6680,10 +6676,9 @@ GHOST_TSuccess GHOST_SystemWayland::cursor_shape_custom_set(uint8_t *bitmap,
   static constexpr uint32_t transparent = 0x00000000;
 
   uint8_t datab = 0, maskb = 0;
-  uint32_t *pixel;
 
   for (int y = 0; y < sizey; ++y) {
-    pixel = &static_cast<uint32_t *>(cursor->custom_data)[y * sizex];
+    uint32_t *pixel = &static_cast<uint32_t *>(cursor->custom_data)[y * sizex];
     for (int x = 0; x < sizex; ++x) {
       if ((x % 8) == 0) {
         datab = *bitmap++;

intern/ghost/intern/GHOST_SystemWayland.hh

@@ -37,15 +37,15 @@ bool ghost_wl_output_own(const struct wl_output *wl_output);
 void ghost_wl_output_tag(struct wl_output *wl_output);
 struct GWL_Output *ghost_wl_output_user_data(struct wl_output *wl_output);
 
-bool ghost_wl_surface_own(const struct wl_surface *surface);
+bool ghost_wl_surface_own(const struct wl_surface *wl_surface);
-void ghost_wl_surface_tag(struct wl_surface *surface);
+void ghost_wl_surface_tag(struct wl_surface *wl_surface);
-GHOST_WindowWayland *ghost_wl_surface_user_data(struct wl_surface *surface);
+GHOST_WindowWayland *ghost_wl_surface_user_data(struct wl_surface *wl_surface);
 
-bool ghost_wl_surface_own_cursor_pointer(const struct wl_surface *surface);
+bool ghost_wl_surface_own_cursor_pointer(const struct wl_surface *wl_surface);
-void ghost_wl_surface_tag_cursor_pointer(struct wl_surface *surface);
+void ghost_wl_surface_tag_cursor_pointer(struct wl_surface *wl_surface);
 
-bool ghost_wl_surface_own_cursor_tablet(const struct wl_surface *surface);
+bool ghost_wl_surface_own_cursor_tablet(const struct wl_surface *wl_surface);
-void ghost_wl_surface_tag_cursor_tablet(struct wl_surface *surface);
+void ghost_wl_surface_tag_cursor_tablet(struct wl_surface *wl_surface);
 
 /* Scaling to: translates from WAYLAND into GHOST (viewport local) coordinates.
  * Scaling from: performs the reverse translation.

intern/ghost/intern/GHOST_SystemX11.cc

@@ -85,7 +85,7 @@ static uchar bit_is_on(const uchar *ptr, int bit)
 
 static GHOST_TKey ghost_key_from_keysym(const KeySym key);
 static GHOST_TKey ghost_key_from_keycode(const XkbDescPtr xkb_descr, const KeyCode keycode);
-static GHOST_TKey ghost_key_from_keysym_or_keycode(const KeySym key,
+static GHOST_TKey ghost_key_from_keysym_or_keycode(const KeySym key_sym,
                                                    const XkbDescPtr xkb_descr,
                                                    const KeyCode keycode);
 
@@ -1776,11 +1776,11 @@ bool GHOST_SystemX11::generateWindowExposeEvents()
   return anyProcessed;
 }
 
-static GHOST_TKey ghost_key_from_keysym_or_keycode(const KeySym keysym,
+static GHOST_TKey ghost_key_from_keysym_or_keycode(const KeySym key_sym,
                                                    XkbDescPtr xkb_descr,
                                                    const KeyCode keycode)
 {
-  GHOST_TKey type = ghost_key_from_keysym(keysym);
+  GHOST_TKey type = ghost_key_from_keysym(key_sym);
   if (type == GHOST_kKeyUnknown) {
     if (xkb_descr) {
       type = ghost_key_from_keycode(xkb_descr, keycode);

intern/ghost/intern/GHOST_SystemX11.hh

@@ -31,7 +31,7 @@
 #endif
 
 /* generic error handlers */
-int GHOST_X11_ApplicationErrorHandler(Display *display, XErrorEvent *theEvent);
+int GHOST_X11_ApplicationErrorHandler(Display *display, XErrorEvent *event);
 int GHOST_X11_ApplicationIOErrorHandler(Display *display);
 
 #define GHOST_X11_ERROR_HANDLERS_OVERRIDE(var) \

intern/ghost/intern/GHOST_Window.hh

@@ -155,7 +155,7 @@ class GHOST_Window : public GHOST_IWindow {
   GHOST_TSuccess getCursorGrabBounds(GHOST_Rect &bounds) const override;
 
   void getCursorGrabState(GHOST_TGrabCursorMode &mode,
-                          GHOST_TAxisFlag &axis_flag,
+                          GHOST_TAxisFlag &wrap_axis,
                           GHOST_Rect &bounds,
                           bool &use_software_cursor) override;
   /**

intern/ghost/intern/GHOST_XrAction.hh

@@ -89,11 +89,13 @@ class GHOST_XrAction {
                    const XrTime &predicted_display_time);
   void applyHapticFeedback(XrSession session,
                            const char *action_name,
-                           const char *subaction_path,
+                           const char *subaction_path_str,
                            const int64_t &duration,
                            const float &frequency,
                            const float &amplitude);
-  void stopHapticFeedback(XrSession session, const char *action_name, const char *subaction_path);
+  void stopHapticFeedback(XrSession session,
+                          const char *action_name,
+                          const char *subaction_path_str);
 
   void *getCustomdata();
   void getBindings(std::map<XrPath, std::vector<XrActionSuggestedBinding>> &r_bindings) const;

scripts/modules/bl_keymap_utils/keymap_hierarchy.py

@@ -204,6 +204,9 @@ _km_hierarchy = [
         ('Grease Pencil Stroke Sculpt (Clone)', 'EMPTY', 'WINDOW', []),
         ('Grease Pencil Stroke Weight Mode', 'EMPTY', 'WINDOW', []),
         ('Grease Pencil Stroke Weight (Draw)', 'EMPTY', 'WINDOW', []),
+        ('Grease Pencil Stroke Weight (Blur)', 'EMPTY', 'WINDOW', []),
+        ('Grease Pencil Stroke Weight (Average)', 'EMPTY', 'WINDOW', []),
+        ('Grease Pencil Stroke Weight (Smear)', 'EMPTY', 'WINDOW', []),
         ('Grease Pencil Stroke Vertex Mode', 'EMPTY', 'WINDOW', []),
         ('Grease Pencil Stroke Vertex (Draw)', 'EMPTY', 'WINDOW', []),
         ('Grease Pencil Stroke Vertex (Blur)', 'EMPTY', 'WINDOW', []),

scripts/presets/keyconfig/keymap_data/blender_default.py

@@ -4178,6 +4178,9 @@ def km_grease_pencil_stroke_weight_mode(params):
         # Brush size
         ("wm.radial_control", {"type": 'F', "value": 'PRESS'},
          {"properties": [("data_path_primary", 'tool_settings.gpencil_weight_paint.brush.size')]}),
+        # Brush weight
+        ("wm.radial_control", {"type": 'F', "value": 'PRESS', "ctrl": True},
+         {"properties": [("data_path_primary", 'tool_settings.gpencil_weight_paint.brush.weight')]}),
         # Increase/Decrease brush size
         ("brush.scale_size", {"type": 'LEFT_BRACKET', "value": 'PRESS', "repeat": True},
          {"properties": [("scalar", 0.9)]}),
@@ -4197,6 +4200,10 @@ def km_grease_pencil_stroke_weight_mode(params):
         op_menu("VIEW3D_MT_gpencil_animation", {"type": 'I', "value": 'PRESS'}),
         # Context menu
         *_template_items_context_panel("VIEW3D_PT_gpencil_weight_context_menu", params.context_menu_event),
+        # Toggle Add/Subtract for weight draw tool
+        ("gpencil.weight_toggle_direction", {"type": 'D', "value": 'PRESS'}, None),
+        # Weight sample
+        ("gpencil.weight_sample", {"type": params.action_mouse, "value": 'PRESS', "ctrl": True}, None),
     ])
 
     if params.select_mouse == 'LEFTMOUSE':
@@ -4220,6 +4227,59 @@ def km_grease_pencil_stroke_weight_draw(_params):
         # Draw
         ("gpencil.weight_paint", {"type": 'LEFTMOUSE', "value": 'PRESS'},
          {"properties": [("wait_for_input", False)]}),
+        ("gpencil.weight_paint", {"type": 'LEFTMOUSE', "value": 'PRESS', "ctrl": True},
+         {"properties": [("wait_for_input", False)]}),
+    ])
+
+    return keymap
+
+
+def km_grease_pencil_stroke_weight_blur(_params):
+    items = []
+    keymap = (
+        "Grease Pencil Stroke Weight (Blur)",
+        {"space_type": 'EMPTY', "region_type": 'WINDOW'},
+        {"items": items},
+    )
+
+    items.extend([
+        # Blur
+        ("gpencil.weight_paint", {"type": 'LEFTMOUSE', "value": 'PRESS'},
+         {"properties": [("wait_for_input", False)]}),
+    ])
+
+    return keymap
+
+
+def km_grease_pencil_stroke_weight_average(_params):
+    items = []
+    keymap = (
+        "Grease Pencil Stroke Weight (Average)",
+        {"space_type": 'EMPTY', "region_type": 'WINDOW'},
+        {"items": items},
+    )
+
+    items.extend([
+        # Average
+        ("gpencil.weight_paint", {"type": 'LEFTMOUSE', "value": 'PRESS'},
+         {"properties": [("wait_for_input", False)]}),
+    ])
+
+    return keymap
+
+
+def km_grease_pencil_stroke_weight_smear(_params):
+    items = []
+    keymap = (
+        "Grease Pencil Stroke Weight (Smear)",
+        {"space_type": 'EMPTY', "region_type": 'WINDOW'},
+        {"items": items},
+    )
+
+    items.extend([
+        # Smear
+        ("gpencil.weight_paint", {"type": 'LEFTMOUSE', "value": 'PRESS'},
+         {"properties": [("wait_for_input", False)]}),
     ])
 
     return keymap
@@ -8115,6 +8175,9 @@ def generate_keymaps(params=None):
         km_grease_pencil_stroke_sculpt_clone(params),
         km_grease_pencil_stroke_weight_mode(params),
         km_grease_pencil_stroke_weight_draw(params),
+        km_grease_pencil_stroke_weight_blur(params),
+        km_grease_pencil_stroke_weight_average(params),
+        km_grease_pencil_stroke_weight_smear(params),
         km_grease_pencil_stroke_vertex_mode(params),
         km_grease_pencil_stroke_vertex_draw(params),
         km_grease_pencil_stroke_vertex_blur(params),

scripts/presets/keyconfig/keymap_data/industry_compatible_data.py

@@ -2777,6 +2777,56 @@ def km_grease_pencil_stroke_weight_draw(_params):
         {"items": items},
     )
 
+    items.extend([
+        ("gpencil.weight_paint", {"type": 'LEFTMOUSE', "value": 'PRESS'},
+         {"properties": [("wait_for_input", False)]}),
+        ("gpencil.weight_paint", {"type": 'LEFTMOUSE', "value": 'PRESS', "ctrl": True},
+         {"properties": [("wait_for_input", False)]}),
+    ])
+
+    return keymap
+
+
+def km_grease_pencil_stroke_weight_blur(_params):
+    items = []
+    keymap = (
+        "Grease Pencil Stroke Weight (Blur)",
+        {"space_type": 'EMPTY', "region_type": 'WINDOW'},
+        {"items": items},
+    )
+
+    items.extend([
+        ("gpencil.weight_paint", {"type": 'LEFTMOUSE', "value": 'PRESS'},
+         {"properties": [("wait_for_input", False)]}),
+    ])
+
+    return keymap
+
+
+def km_grease_pencil_stroke_weight_average(_params):
+    items = []
+    keymap = (
+        "Grease Pencil Stroke Weight (Average)",
+        {"space_type": 'EMPTY', "region_type": 'WINDOW'},
+        {"items": items},
+    )
+
+    items.extend([
+        ("gpencil.weight_paint", {"type": 'LEFTMOUSE', "value": 'PRESS'},
+         {"properties": [("wait_for_input", False)]}),
+    ])
+
+    return keymap
+
+
+def km_grease_pencil_stroke_weight_smear(_params):
+    items = []
+    keymap = (
+        "Grease Pencil Stroke Weight (Smear)",
+        {"space_type": 'EMPTY', "region_type": 'WINDOW'},
+        {"items": items},
+    )
+
     items.extend([
         ("gpencil.weight_paint", {"type": 'LEFTMOUSE', "value": 'PRESS'},
          {"properties": [("wait_for_input", False)]}),
@@ -4195,6 +4245,9 @@ def generate_keymaps_impl(params=None):
         km_grease_pencil_stroke_sculpt_clone(params),
         km_grease_pencil_stroke_weight_mode(params),
         km_grease_pencil_stroke_weight_draw(params),
+        km_grease_pencil_stroke_weight_blur(params),
+        km_grease_pencil_stroke_weight_average(params),
+        km_grease_pencil_stroke_weight_smear(params),
         km_grease_pencil_stroke_vertex_mode(params),
         km_grease_pencil_stroke_vertex_draw(params),
         km_grease_pencil_stroke_vertex_blur(params),

scripts/startup/bl_ui/node_add_menu_geometry.py

@@ -220,6 +220,7 @@ class NODE_MT_geometry_node_GEO_GEOMETRY_SAMPLE(Menu):
     def draw(self, _context):
         layout = self.layout
         node_add_menu.add_node_type(layout, "GeometryNodeProximity")
+        node_add_menu.add_node_type(layout, "GeometryNodeIndexOfNearest")
         node_add_menu.add_node_type(layout, "GeometryNodeRaycast")
         node_add_menu.add_node_type(layout, "GeometryNodeSampleIndex")
         node_add_menu.add_node_type(layout, "GeometryNodeSampleNearest")
@@ -528,7 +529,6 @@ class NODE_MT_category_GEO_UTILITIES_FIELD(Menu):
         node_add_menu.add_node_type(layout, "GeometryNodeAccumulateField")
         node_add_menu.add_node_type(layout, "GeometryNodeFieldAtIndex")
         node_add_menu.add_node_type(layout, "GeometryNodeFieldOnDomain")
-        node_add_menu.add_node_type(layout, "GeometryNodeIndexOfNearest")
         node_add_menu.draw_assets_for_catalog(layout, self.bl_label)
 
 

scripts/startup/bl_ui/properties_data_curve.py

@@ -371,10 +371,18 @@ class DATA_PT_font(CurveButtonsPanelText, Panel):
                 row.prop(char, "use_small_caps", toggle=True)
             else:
                 row = layout.row(align=True)
-                row.operator("font.style_toggle", text="Bold", icon='BOLD' , depress = text.is_select_bold).style = 'BOLD'
+                row.operator(
-                row.operator("font.style_toggle", text="Italic", icon='ITALIC' , depress = text.is_select_italic).style = 'ITALIC'
+                    "font.style_toggle", text="Bold", icon='BOLD', depress=text.is_select_bold,
-                row.operator("font.style_toggle", text="Underline", icon='UNDERLINE' , depress = text.is_select_underline).style = 'UNDERLINE'
+                ).style = 'BOLD'
-                row.operator("font.style_toggle", text="Small Caps", icon='SMALL_CAPS' , depress = text.is_select_smallcaps).style = 'SMALL_CAPS'
+                row.operator(
+                    "font.style_toggle", text="Italic", icon='ITALIC', depress=text.is_select_italic,
+                ).style = 'ITALIC'
+                row.operator(
+                    "font.style_toggle", text="Underline", icon='UNDERLINE', depress=text.is_select_underline,
+                ).style = 'UNDERLINE'
+                row.operator(
+                    "font.style_toggle", text="Small Caps", icon='SMALL_CAPS', depress=text.is_select_smallcaps,
+                ).style = 'SMALL_CAPS'
 
 
 class DATA_PT_font_transform(CurveButtonsPanelText, Panel):

scripts/startup/bl_ui/properties_paint_common.py

@@ -1415,8 +1415,12 @@ def brush_basic_gpencil_weight_settings(layout, _context, brush, *, compact=Fals
     row.prop(brush, "strength", slider=True)
     row.prop(brush, "use_pressure_strength", text="")
 
+    if brush.gpencil_weight_tool in {'WEIGHT'}:
         layout.prop(brush, "weight", slider=True)
 
+        gp_settings = brush.gpencil_settings
+        layout.prop(gp_settings, "direction", expand=True, text="" if compact else "Direction")
+
 
 def brush_basic_gpencil_vertex_settings(layout, _context, brush, *, compact=False):
     gp_settings = brush.gpencil_settings

scripts/startup/bl_ui/space_filebrowser.py

@@ -706,6 +706,15 @@ class ASSETBROWSER_PT_metadata(asset_utils.AssetBrowserPanel, Panel):
     bl_label = "Asset Metadata"
     bl_options = {'HIDE_HEADER'}
 
+    @staticmethod
+    def metadata_prop(layout, asset_data, propname):
+        """
+        Only display properties that are either set or can be modified (i.e. the
+        asset is in the current file). Empty, non-editable fields are not really useful.
+        """
+        if getattr(asset_data, propname) or not asset_data.is_property_readonly(propname):
+            layout.prop(asset_data, propname)
+
     def draw(self, context):
         layout = self.layout
         wm = context.window_manager
@@ -745,10 +754,11 @@ class ASSETBROWSER_PT_metadata(asset_utils.AssetBrowserPanel, Panel):
         row.prop(wm, "asset_path_dummy", text="Source", icon='CURRENT_FILE' if is_local_asset else 'NONE')
         row.operator("asset.open_containing_blend_file", text="", icon='TOOL_SETTINGS')
 
-        layout.prop(asset_file_handle.asset_data, "description")
+        asset_data = asset_file_handle.asset_data
-        layout.prop(asset_file_handle.asset_data, "license")
+        self.metadata_prop(layout, asset_data, "description")
-        layout.prop(asset_file_handle.asset_data, "copyright")
+        self.metadata_prop(layout, asset_data, "license")
-        layout.prop(asset_file_handle.asset_data, "author")
+        self.metadata_prop(layout, asset_data, "copyright")
+        self.metadata_prop(layout, asset_data, "author")
 
 
 class ASSETBROWSER_PT_metadata_preview(asset_utils.AssetMetaDataPanel, Panel):

scripts/startup/bl_ui/space_node.py

@@ -903,7 +903,10 @@ class NodeTreeInterfacePanel(Panel):
                     props = property_row.operator_menu_enum(
                         "node.tree_socket_change_subtype",
                         "socket_subtype",
-                        text=active_socket.bl_subtype_label if active_socket.bl_subtype_label else active_socket.bl_idname
+                        text=(
+                            active_socket.bl_subtype_label if active_socket.bl_subtype_label else
+                            active_socket.bl_idname
+                        ),
                     )
 
             layout.use_property_split = True

scripts/startup/bl_ui/space_toolsystem_toolbar.py

@@ -68,6 +68,7 @@ def generate_from_enum_ex(
                 dict(
                     idname=idname_prefix + name,
                     label=name,
+                    description=enum.description,
                     icon=icon,
                     cursor=cursor,
                     data_block=idname,

scripts/startup/bl_ui/space_view3d.py

@@ -8,6 +8,7 @@ from bpy.types import (
 from bl_ui.properties_paint_common import (
     UnifiedPaintPanel,
     brush_basic_texpaint_settings,
+    brush_basic_gpencil_weight_settings,
 )
 from bl_ui.properties_grease_pencil_common import (
     AnnotationDataPanel,
@@ -414,11 +415,13 @@ class _draw_tool_settings_context_mode:
         paint = context.tool_settings.gpencil_weight_paint
         brush = paint.brush
 
-        from bl_ui.properties_paint_common import (
+        layout.template_ID_preview(paint, "brush", rows=3, cols=8, hide_buttons=True)
-            brush_basic_gpencil_weight_settings,
+
-        )
         brush_basic_gpencil_weight_settings(layout, context, brush, compact=True)
 
+        layout.popover("VIEW3D_PT_tools_grease_pencil_weight_options", text="Options")
+        layout.popover("VIEW3D_PT_tools_grease_pencil_brush_weight_falloff", text="Falloff")
+
         return True
 
     @staticmethod
@@ -7662,12 +7665,10 @@ class VIEW3D_PT_gpencil_weight_context_menu(Panel):
         tool_settings = context.tool_settings
         settings = tool_settings.gpencil_weight_paint
         brush = settings.brush
-
         layout = self.layout
 
-        layout.prop(brush, "size", slider=True)
+        # Weight settings
-        layout.prop(brush, "strength")
+        brush_basic_gpencil_weight_settings(layout, context, brush)
-        layout.prop(brush, "weight")
 
         # Layers
         draw_gpencil_layer_active(context, layout)

scripts/startup/bl_ui/space_view3d_toolbar.py

@@ -2021,6 +2021,9 @@ class VIEW3D_PT_tools_grease_pencil_weight_paint_settings(Panel, View3DPanel, Gr
     bl_label = "Brush Settings"
 
     def draw(self, context):
+        if self.is_popover:
+            return
+
         layout = self.layout
         layout.use_property_split = True
         layout.use_property_decorate = False
@@ -2029,7 +2032,6 @@ class VIEW3D_PT_tools_grease_pencil_weight_paint_settings(Panel, View3DPanel, Gr
         settings = tool_settings.gpencil_weight_paint
         brush = settings.brush
 
-        if not self.is_popover:
         from bl_ui.properties_paint_common import (
             brush_basic_gpencil_weight_settings,
         )
@@ -2038,6 +2040,7 @@ class VIEW3D_PT_tools_grease_pencil_weight_paint_settings(Panel, View3DPanel, Gr
 
 class VIEW3D_PT_tools_grease_pencil_brush_weight_falloff(GreasePencilBrushFalloff, Panel, View3DPaintPanel):
     bl_context = ".greasepencil_weight"
+    bl_parent_id = 'VIEW3D_PT_tools_grease_pencil_weight_paint_settings'
     bl_label = "Falloff"
     bl_options = {'DEFAULT_CLOSED'}
 
@@ -2049,6 +2052,20 @@ class VIEW3D_PT_tools_grease_pencil_brush_weight_falloff(GreasePencilBrushFallof
         return (brush and brush.curve)
 
 
+class VIEW3D_PT_tools_grease_pencil_weight_options(Panel, View3DPanel, GreasePencilWeightPanel):
+    bl_label = "Options"
+    bl_options = {'DEFAULT_CLOSED'}
+
+    def draw(self, context):
+        layout = self.layout
+        layout.use_property_split = True
+        layout.use_property_decorate = False
+        tool_settings = context.scene.tool_settings
+
+        col = layout.column()
+        col.prop(tool_settings, "use_auto_normalize", text="Auto Normalize")
+
+
 # Grease Pencil vertex painting tools
 class GreasePencilVertexPanel:
     bl_context = ".greasepencil_vertex"
@@ -2425,6 +2442,7 @@ classes = (
     VIEW3D_PT_tools_grease_pencil_sculpt_appearance,
     VIEW3D_PT_tools_grease_pencil_weight_paint_select,
     VIEW3D_PT_tools_grease_pencil_weight_paint_settings,
+    VIEW3D_PT_tools_grease_pencil_weight_options,
     VIEW3D_PT_tools_grease_pencil_weight_appearance,
     VIEW3D_PT_tools_grease_pencil_vertex_paint_select,
     VIEW3D_PT_tools_grease_pencil_vertex_paint_settings,

source/blender/asset_system/intern/asset_catalog.cc

@@ -260,6 +260,7 @@ void AssetCatalogService::update_catalog_path(const CatalogID catalog_id,
     }
     cat->path = new_path;
     cat->simple_name_refresh();
+    this->tag_has_unsaved_changes(cat);
 
     /* TODO(Sybren): go over all assets that are assigned to this catalog, defined in the current
      * blend file, and update the catalog simple name stored there. */

source/blender/blenkernel/BKE_curves.h

@@ -28,7 +28,7 @@ bool BKE_curves_attribute_required(const struct Curves *curves, const char *name
 
 /* Depsgraph */
 
-struct Curves *BKE_curves_copy_for_eval(struct Curves *curves_src, bool reference);
+struct Curves *BKE_curves_copy_for_eval(struct Curves *curves_src);
 
 void BKE_curves_data_update(struct Depsgraph *depsgraph,
                             struct Scene *scene,

source/blender/blenkernel/BKE_lib_id.h

@@ -152,8 +152,6 @@ enum {
   LIB_ID_COPY_CACHES = 1 << 18,
   /** Don't copy `id->adt`, used by ID data-block localization routines. */
   LIB_ID_COPY_NO_ANIMDATA = 1 << 19,
-  /** Mesh: Reference CD data layers instead of doing real copy - USE WITH CAUTION! */
-  LIB_ID_COPY_CD_REFERENCE = 1 << 20,
   /** Do not copy id->override_library, used by ID data-block override routines. */
   LIB_ID_COPY_NO_LIB_OVERRIDE = 1 << 21,
   /** When copying local sub-data (like constraints or modifiers), do not set their "library

source/blender/blenkernel/BKE_mesh.h

@@ -158,15 +158,15 @@ void BKE_mesh_ensure_skin_customdata(struct Mesh *me);
 /** Add poly offsets to describe faces to a new mesh. */
 void BKE_mesh_poly_offsets_ensure_alloc(struct Mesh *mesh);
 
-struct Mesh *BKE_mesh_new_nomain(int verts_len, int edges_len, int loops_len, int polys_len);
+struct Mesh *BKE_mesh_new_nomain(int verts_num, int edges_num, int polys_num, int loops_num);
 struct Mesh *BKE_mesh_new_nomain_from_template(
-    const struct Mesh *me_src, int verts_len, int edges_len, int loops_len, int polys_len);
+    const struct Mesh *me_src, int verts_num, int edges_num, int polys_num, int loops_num);
 struct Mesh *BKE_mesh_new_nomain_from_template_ex(const struct Mesh *me_src,
-                                                  int verts_len,
+                                                  int verts_num,
-                                                  int edges_len,
+                                                  int edges_num,
-                                                  int tessface_len,
+                                                  int tessface_num,
-                                                  int loops_len,
+                                                  int polys_num,
-                                                  int polys_len,
+                                                  int loops_num,
                                                   struct CustomData_MeshMasks mask);
 
 void BKE_mesh_eval_delete(struct Mesh *mesh_eval);
@@ -175,7 +175,7 @@ void BKE_mesh_eval_delete(struct Mesh *mesh_eval);
  * Performs copy for use during evaluation,
  * optional referencing original arrays to reduce memory.
  */
-struct Mesh *BKE_mesh_copy_for_eval(const struct Mesh *source, bool reference);
+struct Mesh *BKE_mesh_copy_for_eval(const struct Mesh *source);
 
 /**
  * These functions construct a new Mesh,

source/blender/blenkernel/BKE_node.h

@@ -298,6 +298,12 @@ typedef struct bNodeType {
   void (*freefunc_api)(struct PointerRNA *ptr);
   void (*copyfunc_api)(struct PointerRNA *ptr, const struct bNode *src_node);
 
+  /**
+   * An additional poll test for deciding whether nodes should be an option in search menus.
+   * Potentially more strict poll than #poll(), but doesn't have to check the same things.
+   */
+  bool (*add_ui_poll)(const struct bContext *C);
+
   /**
    * Can this node type be added to a node tree?
    * \param r_disabled_hint: Hint to display in the UI when the poll fails.

source/blender/blenkernel/BKE_pointcloud.h

@@ -70,8 +70,7 @@ void *BKE_pointcloud_add(struct Main *bmain, const char *name);
 void *BKE_pointcloud_add_default(struct Main *bmain, const char *name);
 struct PointCloud *BKE_pointcloud_new_nomain(int totpoint);
 void BKE_pointcloud_nomain_to_pointcloud(struct PointCloud *pointcloud_src,
-                                         struct PointCloud *pointcloud_dst,
+                                         struct PointCloud *pointcloud_dst);
-                                         bool take_ownership);
 
 struct BoundBox *BKE_pointcloud_boundbox_get(struct Object *ob);
 
@@ -79,7 +78,7 @@ bool BKE_pointcloud_attribute_required(const struct PointCloud *pointcloud, cons
 
 /* Dependency Graph */
 
-struct PointCloud *BKE_pointcloud_copy_for_eval(struct PointCloud *pointcloud_src, bool reference);
+struct PointCloud *BKE_pointcloud_copy_for_eval(struct PointCloud *pointcloud_src);
 
 void BKE_pointcloud_data_update(struct Depsgraph *depsgraph,
                                 struct Scene *scene,

source/blender/blenkernel/BKE_vfont.h

@@ -53,8 +53,9 @@ typedef struct EditFont {
   int selstart, selend;
 
   /**
-   * Combined styles (#CharInfo.flag) for selected string. A flag will be
+   * Combined styles from #CharInfo.flag for the selected range selected
-   * set only if ALL characters in the selected string have it.
+   * (only including values from #CU_CHINFO_STYLE_ALL).
+   * A flag will be set only if ALL characters in the selected string have it.
    */
   int select_char_info_flag;
 

source/blender/blenkernel/BKE_volume.h

@@ -130,7 +130,7 @@ void BKE_volume_grid_transform_matrix_set(const struct Volume *volume,
  * file path. Grids are shared with the source data-block, not copied. */
 
 struct Volume *BKE_volume_new_for_eval(const struct Volume *volume_src);
-struct Volume *BKE_volume_copy_for_eval(struct Volume *volume_src, bool reference);
+struct Volume *BKE_volume_copy_for_eval(struct Volume *volume_src);
 
 struct VolumeGrid *BKE_volume_grid_add(struct Volume *volume,
                                        const char *name,

source/blender/blenkernel/intern/DerivedMesh.cc

@@ -380,7 +380,7 @@ static Mesh *create_orco_mesh(Object *ob, Mesh *me, BMEditMesh *em, int layer)
     BKE_mesh_ensure_default_orig_index_customdata(mesh);
   }
   else {
-    mesh = BKE_mesh_copy_for_eval(me, true);
+    mesh = BKE_mesh_copy_for_eval(me);
   }
 
   orco = get_orco_coords(ob, em, layer, &free);
@@ -654,7 +654,7 @@ static void mesh_calc_modifiers(struct Depsgraph *depsgraph,
 
   if (ob->modifier_flag & OB_MODIFIER_FLAG_ADD_REST_POSITION) {
     if (mesh_final == nullptr) {
-      mesh_final = BKE_mesh_copy_for_eval(mesh_input, true);
+      mesh_final = BKE_mesh_copy_for_eval(mesh_input);
       ASSERT_IS_VALID_MESH(mesh_final);
     }
     MutableAttributeAccessor attributes = mesh_final->attributes_for_write();
@@ -685,7 +685,7 @@ static void mesh_calc_modifiers(struct Depsgraph *depsgraph,
       if (mti->type == eModifierTypeType_OnlyDeform && !sculpt_dyntopo) {
         blender::bke::ScopedModifierTimer modifier_timer{*md};
         if (!mesh_final) {
-          mesh_final = BKE_mesh_copy_for_eval(mesh_input, true);
+          mesh_final = BKE_mesh_copy_for_eval(mesh_input);
           ASSERT_IS_VALID_MESH(mesh_final);
         }
         BKE_modifier_deform_verts(md,
@@ -703,12 +703,7 @@ static void mesh_calc_modifiers(struct Depsgraph *depsgraph,
      * places that wish to use the original mesh but with deformed
      * coordinates (like vertex paint). */
     if (r_deform) {
-      if (mesh_final) {
+      mesh_deform = BKE_mesh_copy_for_eval(mesh_final ? mesh_final : mesh_input);
-        mesh_deform = BKE_mesh_copy_for_eval(mesh_final, false);
-      }
-      else {
-        mesh_deform = BKE_mesh_copy_for_eval(mesh_input, false);
-      }
     }
   }
 
@@ -779,7 +774,7 @@ static void mesh_calc_modifiers(struct Depsgraph *depsgraph,
 
     if (mti->type == eModifierTypeType_OnlyDeform) {
       if (!mesh_final) {
-        mesh_final = BKE_mesh_copy_for_eval(mesh_input, true);
+        mesh_final = BKE_mesh_copy_for_eval(mesh_input);
         ASSERT_IS_VALID_MESH(mesh_final);
       }
       BKE_modifier_deform_verts(md,
@@ -798,7 +793,7 @@ static void mesh_calc_modifiers(struct Depsgraph *depsgraph,
         }
       }
       else {
-        mesh_final = BKE_mesh_copy_for_eval(mesh_input, true);
+        mesh_final = BKE_mesh_copy_for_eval(mesh_input);
         ASSERT_IS_VALID_MESH(mesh_final);
         check_for_needs_mapping = true;
       }
@@ -966,7 +961,7 @@ static void mesh_calc_modifiers(struct Depsgraph *depsgraph,
       mesh_final = mesh_input;
     }
     else {
-      mesh_final = BKE_mesh_copy_for_eval(mesh_input, true);
+      mesh_final = BKE_mesh_copy_for_eval(mesh_input);
     }
   }
 
@@ -1011,7 +1006,7 @@ static void mesh_calc_modifiers(struct Depsgraph *depsgraph,
         /* Not yet finalized by any instance, do it now
          * Isolate since computing normals is multithreaded and we are holding a lock. */
         blender::threading::isolate_task([&] {
-          mesh_final = BKE_mesh_copy_for_eval(mesh_input, true);
+          mesh_final = BKE_mesh_copy_for_eval(mesh_input);
           mesh_calc_modifier_final_normals(
               mesh_input, &final_datamask, sculpt_dyntopo, mesh_final);
           mesh_calc_finalize(mesh_input, mesh_final);
@@ -1026,7 +1021,7 @@ static void mesh_calc_modifiers(struct Depsgraph *depsgraph,
     else if (!mesh_has_modifier_final_normals(mesh_input, &final_datamask, runtime->mesh_eval)) {
       /* Modifier stack was (re-)evaluated with a request for additional normals
        * different than the instanced mesh, can't instance anymore now. */
-      mesh_final = BKE_mesh_copy_for_eval(mesh_input, true);
+      mesh_final = BKE_mesh_copy_for_eval(mesh_input);
       mesh_calc_modifier_final_normals(mesh_input, &final_datamask, sculpt_dyntopo, mesh_final);
       mesh_calc_finalize(mesh_input, mesh_final);
     }
@@ -1255,7 +1250,7 @@ static void editbmesh_calc_modifiers(struct Depsgraph *depsgraph,
       /* apply vertex coordinates or build a DerivedMesh as necessary */
       if (mesh_final) {
         if (deformed_verts) {
-          Mesh *mesh_tmp = BKE_mesh_copy_for_eval(mesh_final, false);
+          Mesh *mesh_tmp = BKE_mesh_copy_for_eval(mesh_final);
           if (mesh_final != mesh_cage) {
             BKE_id_free(nullptr, mesh_final);
           }
@@ -1264,7 +1259,7 @@ static void editbmesh_calc_modifiers(struct Depsgraph *depsgraph,
         }
         else if (mesh_final == mesh_cage) {
           /* 'me' may be changed by this modifier, so we need to copy it. */
-          mesh_final = BKE_mesh_copy_for_eval(mesh_final, false);
+          mesh_final = BKE_mesh_copy_for_eval(mesh_final);
         }
       }
       else {
@@ -1337,7 +1332,7 @@ static void editbmesh_calc_modifiers(struct Depsgraph *depsgraph,
 
     if (r_cage && i == cageIndex) {
       if (mesh_final && deformed_verts) {
-        mesh_cage = BKE_mesh_copy_for_eval(mesh_final, false);
+        mesh_cage = BKE_mesh_copy_for_eval(mesh_final);
         BKE_mesh_vert_coords_apply(mesh_cage, deformed_verts);
       }
       else if (mesh_final) {
@@ -1373,7 +1368,7 @@ static void editbmesh_calc_modifiers(struct Depsgraph *depsgraph,
   if (mesh_final) {
     if (deformed_verts) {
       if (mesh_final == mesh_cage) {
-        mesh_final = BKE_mesh_copy_for_eval(mesh_final, false);
+        mesh_final = BKE_mesh_copy_for_eval(mesh_final);
       }
       BKE_mesh_vert_coords_apply(mesh_final, deformed_verts);
     }

source/blender/blenkernel/intern/brush.cc

@@ -1276,14 +1276,57 @@ void BKE_gpencil_brush_preset_set(Main *bmain, Brush *brush, const short type)
 
       break;
     }
-    case GP_BRUSH_PRESET_DRAW_WEIGHT: {
+    case GP_BRUSH_PRESET_WEIGHT_DRAW: {
       brush->gpencil_settings->icon_id = GP_BRUSH_ICON_GPBRUSH_WEIGHT;
       brush->gpencil_weight_tool = GPWEIGHT_TOOL_DRAW;
 
       brush->size = 25.0f;
       brush->gpencil_settings->flag |= GP_BRUSH_USE_PRESSURE;
 
-      brush->gpencil_settings->draw_strength = 0.8f;
+      brush->alpha = 0.3f;
+      brush->gpencil_settings->draw_strength = 0.3f;
+      brush->gpencil_settings->flag |= GP_BRUSH_USE_STRENGTH_PRESSURE;
+      brush->gpencil_settings->sculpt_mode_flag |= GP_SCULPT_FLAGMODE_APPLY_POSITION;
+
+      break;
+    }
+    case GP_BRUSH_PRESET_WEIGHT_BLUR: {
+      brush->gpencil_settings->icon_id = GP_BRUSH_ICON_VERTEX_BLUR;
+      brush->gpencil_weight_tool = GPWEIGHT_TOOL_BLUR;
+
+      brush->size = 50.0f;
+      brush->gpencil_settings->flag |= GP_BRUSH_USE_PRESSURE;
+
+      brush->alpha = 0.3f;
+      brush->gpencil_settings->draw_strength = 0.3f;
+      brush->gpencil_settings->flag |= GP_BRUSH_USE_STRENGTH_PRESSURE;
+      brush->gpencil_settings->sculpt_mode_flag |= GP_SCULPT_FLAGMODE_APPLY_POSITION;
+
+      break;
+    }
+    case GP_BRUSH_PRESET_WEIGHT_AVERAGE: {
+      brush->gpencil_settings->icon_id = GP_BRUSH_ICON_VERTEX_BLUR;
+      brush->gpencil_weight_tool = GPWEIGHT_TOOL_AVERAGE;
+
+      brush->size = 50.0f;
+      brush->gpencil_settings->flag |= GP_BRUSH_USE_PRESSURE;
+
+      brush->alpha = 0.3f;
+      brush->gpencil_settings->draw_strength = 0.3f;
+      brush->gpencil_settings->flag |= GP_BRUSH_USE_STRENGTH_PRESSURE;
+      brush->gpencil_settings->sculpt_mode_flag |= GP_SCULPT_FLAGMODE_APPLY_POSITION;
+
+      break;
+    }
+    case GP_BRUSH_PRESET_WEIGHT_SMEAR: {
+      brush->gpencil_settings->icon_id = GP_BRUSH_ICON_VERTEX_BLUR;
+      brush->gpencil_weight_tool = GPWEIGHT_TOOL_SMEAR;
+
+      brush->size = 50.0f;
+      brush->gpencil_settings->flag |= GP_BRUSH_USE_PRESSURE;
+
+      brush->alpha = 0.3f;
+      brush->gpencil_settings->draw_strength = 0.3f;
       brush->gpencil_settings->flag |= GP_BRUSH_USE_STRENGTH_PRESSURE;
       brush->gpencil_settings->sculpt_mode_flag |= GP_SCULPT_FLAGMODE_APPLY_POSITION;
 
@@ -1569,13 +1612,32 @@ void BKE_brush_gpencil_weight_presets(Main *bmain, ToolSettings *ts, const bool
   Paint *weightpaint = &ts->gp_weightpaint->paint;
   Brush *brush_prev = weightpaint->brush;
   Brush *brush, *deft_weight;
-  /* Vertex Draw brush. */
+
-  brush = gpencil_brush_ensure(bmain, ts, "Draw Weight", OB_MODE_WEIGHT_GPENCIL, &r_new);
+  /* Weight Draw brush. */
+  brush = gpencil_brush_ensure(bmain, ts, "Weight Draw", OB_MODE_WEIGHT_GPENCIL, &r_new);
   if ((reset) || (r_new)) {
-    BKE_gpencil_brush_preset_set(bmain, brush, GP_BRUSH_PRESET_DRAW_WEIGHT);
+    BKE_gpencil_brush_preset_set(bmain, brush, GP_BRUSH_PRESET_WEIGHT_DRAW);
   }
   deft_weight = brush; /* save default brush. */
 
+  /* Weight Blur brush. */
+  brush = gpencil_brush_ensure(bmain, ts, "Weight Blur", OB_MODE_WEIGHT_GPENCIL, &r_new);
+  if ((reset) || (r_new)) {
+    BKE_gpencil_brush_preset_set(bmain, brush, GP_BRUSH_PRESET_WEIGHT_BLUR);
+  }
+
+  /* Weight Average brush. */
+  brush = gpencil_brush_ensure(bmain, ts, "Weight Average", OB_MODE_WEIGHT_GPENCIL, &r_new);
+  if ((reset) || (r_new)) {
+    BKE_gpencil_brush_preset_set(bmain, brush, GP_BRUSH_PRESET_WEIGHT_AVERAGE);
+  }
+
+  /* Weight Smear brush. */
+  brush = gpencil_brush_ensure(bmain, ts, "Weight Smear", OB_MODE_WEIGHT_GPENCIL, &r_new);
+  if ((reset) || (r_new)) {
+    BKE_gpencil_brush_preset_set(bmain, brush, GP_BRUSH_PRESET_WEIGHT_SMEAR);
+  }
+
   /* Set default brush. */
   if (reset || brush_prev == nullptr) {
     BKE_paint_brush_set(weightpaint, deft_weight);

source/blender/blenkernel/intern/cachefile.c

@@ -399,7 +399,7 @@ bool BKE_cachefile_filepath_get(const Main *bmain,
     const int frame = (int)BKE_cachefile_time_offset(cache_file, (double)ctime, fps);
 
     char ext[32];
-    BLI_path_frame_strip(r_filepath, ext);
+    BLI_path_frame_strip(r_filepath, ext, sizeof(ext));
     BLI_path_frame(r_filepath, frame, frame_len);
     BLI_path_extension_ensure(r_filepath, FILE_MAX, ext);
 

source/blender/blenkernel/intern/cloth.cc

@@ -1176,7 +1176,7 @@ static void cloth_update_verts(Object *ob, ClothModifierData *clmd, Mesh *mesh)
 static Mesh *cloth_make_rest_mesh(ClothModifierData *clmd, Mesh *mesh)
 {
   using namespace blender;
-  Mesh *new_mesh = BKE_mesh_copy_for_eval(mesh, false);
+  Mesh *new_mesh = BKE_mesh_copy_for_eval(mesh);
   ClothVertex *verts = clmd->clothObject->verts;
   MutableSpan<float3> positions = mesh->vert_positions_for_write();
 

source/blender/blenkernel/intern/crazyspace.cc

@@ -389,7 +389,7 @@ int BKE_sculpt_get_first_deform_matrices(struct Depsgraph *depsgraph,
       if (defmats == nullptr) {
         /* NOTE: Evaluated object is re-set to its original un-deformed state. */
         Mesh *me = static_cast<Mesh *>(object_eval.data);
-        me_eval = BKE_mesh_copy_for_eval(me, true);
+        me_eval = BKE_mesh_copy_for_eval(me);
         crazyspace_init_verts_and_matrices(me_eval, &defmats, &deformedVerts);
       }
 
@@ -470,7 +470,7 @@ void BKE_crazyspace_build_sculpt(struct Depsgraph *depsgraph,
         }
 
         if (mesh_eval == nullptr) {
-          mesh_eval = BKE_mesh_copy_for_eval(mesh, true);
+          mesh_eval = BKE_mesh_copy_for_eval(mesh);
         }
 
         mti->deformVerts(md, &mectx, mesh_eval, deformedVerts, mesh_eval->totvert);

source/blender/blenkernel/intern/curve_to_mesh_convert.cc

@@ -691,7 +691,7 @@ Mesh *curve_to_mesh_sweep(const CurvesGeometry &main,
   }
 
   Mesh *mesh = BKE_mesh_new_nomain(
-      offsets.vert.last(), offsets.edge.last(), offsets.loop.last(), offsets.poly.last());
+      offsets.vert.last(), offsets.edge.last(), offsets.poly.last(), offsets.loop.last());
   mesh->flag |= ME_AUTOSMOOTH;
   mesh->smoothresh = DEG2RADF(180.0f);
   MutableSpan<float3> positions = mesh->vert_positions_for_write();

source/blender/blenkernel/intern/curves.cc

@@ -220,16 +220,10 @@ bool BKE_curves_attribute_required(const Curves * /*curves*/, const char *name)
   return STREQ(name, ATTR_POSITION);
 }
 
-Curves *BKE_curves_copy_for_eval(Curves *curves_src, bool reference)
+Curves *BKE_curves_copy_for_eval(Curves *curves_src)
 {
-  int flags = LIB_ID_COPY_LOCALIZE;
+  return reinterpret_cast<Curves *>(
-
+      BKE_id_copy_ex(nullptr, &curves_src->id, nullptr, LIB_ID_COPY_LOCALIZE));
-  if (reference) {
-    flags |= LIB_ID_COPY_CD_REFERENCE;
-  }
-
-  Curves *result = (Curves *)BKE_id_copy_ex(nullptr, &curves_src->id, nullptr, flags);
-  return result;
 }
 
 static void curves_evaluate_modifiers(struct Depsgraph *depsgraph,

source/blender/blenkernel/intern/dynamicpaint.cc

@@ -1904,7 +1904,7 @@ static void dynamic_paint_apply_surface_wave_cb(void *__restrict userdata,
  */
 static Mesh *dynamicPaint_Modifier_apply(DynamicPaintModifierData *pmd, Object *ob, Mesh *mesh)
 {
-  Mesh *result = BKE_mesh_copy_for_eval(mesh, false);
+  Mesh *result = BKE_mesh_copy_for_eval(mesh);
 
   if (pmd->canvas && !(pmd->canvas->flags & MOD_DPAINT_BAKING) &&
       pmd->type == MOD_DYNAMICPAINT_TYPE_CANVAS) {
@@ -2049,7 +2049,7 @@ static Mesh *dynamicPaint_Modifier_apply(DynamicPaintModifierData *pmd, Object *
     if (runtime_data->brush_mesh != nullptr) {
       BKE_id_free(nullptr, runtime_data->brush_mesh);
     }
-    runtime_data->brush_mesh = BKE_mesh_copy_for_eval(result, false);
+    runtime_data->brush_mesh = BKE_mesh_copy_for_eval(result);
   }
 
   return result;
@@ -2070,7 +2070,7 @@ static void canvas_copyMesh(DynamicPaintCanvasSettings *canvas, Mesh *mesh)
     BKE_id_free(nullptr, runtime->canvas_mesh);
   }
 
-  runtime->canvas_mesh = BKE_mesh_copy_for_eval(mesh, false);
+  runtime->canvas_mesh = BKE_mesh_copy_for_eval(mesh);
 }
 
 /*
@@ -3796,7 +3796,7 @@ static void dynamicPaint_brushMeshCalculateVelocity(Depsgraph *depsgraph,
                                       SUBFRAME_RECURSION,
                                       BKE_scene_ctime_get(scene),
                                       eModifierType_DynamicPaint);
-  mesh_p = BKE_mesh_copy_for_eval(dynamicPaint_brush_mesh_get(brush), false);
+  mesh_p = BKE_mesh_copy_for_eval(dynamicPaint_brush_mesh_get(brush));
   numOfVerts_p = mesh_p->totvert;
 
   float(*positions_p)[3] = BKE_mesh_vert_positions_for_write(mesh_p);
@@ -4282,7 +4282,7 @@ static bool dynamicPaint_paintMesh(Depsgraph *depsgraph,
     Bounds3D mesh_bb = {{0}};
     VolumeGrid *grid = bData->grid;
 
-    mesh = BKE_mesh_copy_for_eval(brush_mesh, false);
+    mesh = BKE_mesh_copy_for_eval(brush_mesh);
     float(*positions)[3] = BKE_mesh_vert_positions_for_write(mesh);
     const blender::Span<blender::float3> vert_normals = mesh->vert_normals();
     const blender::Span<int> corner_verts = mesh->corner_verts();

source/blender/blenkernel/intern/fluid.cc

@@ -1003,7 +1003,7 @@ static void obstacles_from_mesh(Object *coll_ob,
     float *vert_vel = nullptr;
     bool has_velocity = false;
 
-    Mesh *me = BKE_mesh_copy_for_eval(fes->mesh, false);
+    Mesh *me = BKE_mesh_copy_for_eval(fes->mesh);
     float(*positions)[3] = BKE_mesh_vert_positions_for_write(me);
 
     int min[3], max[3], res[3];
@@ -2062,7 +2062,7 @@ static void emit_from_mesh(
 
     /* Copy mesh for thread safety as we modify it.
      * Main issue is its VertArray being modified, then replaced and freed. */
-    Mesh *me = BKE_mesh_copy_for_eval(ffs->mesh, false);
+    Mesh *me = BKE_mesh_copy_for_eval(ffs->mesh);
     float(*positions)[3] = BKE_mesh_vert_positions_for_write(me);
 
     const blender::Span<int> corner_verts = me->corner_verts();
@@ -3229,7 +3229,7 @@ static Mesh *create_liquid_geometry(FluidDomainSettings *fds,
     return nullptr;
   }
 
-  me = BKE_mesh_new_nomain(num_verts, 0, num_faces * 3, num_faces);
+  me = BKE_mesh_new_nomain(num_verts, 0, num_faces, num_faces * 3);
   if (!me) {
     return nullptr;
   }
@@ -3364,10 +3364,10 @@ static Mesh *create_smoke_geometry(FluidDomainSettings *fds, Mesh *orgmesh, Obje
 
   /* Just copy existing mesh if there is no content or if the adaptive domain is not being used. */
   if (fds->total_cells <= 1 || (fds->flags & FLUID_DOMAIN_USE_ADAPTIVE_DOMAIN) == 0) {
-    return BKE_mesh_copy_for_eval(orgmesh, false);
+    return BKE_mesh_copy_for_eval(orgmesh);
   }
 
-  result = BKE_mesh_new_nomain(num_verts, 0, num_faces * 4, num_faces);
+  result = BKE_mesh_new_nomain(num_verts, 0, num_faces, num_faces * 4);
   float(*positions)[3] = BKE_mesh_vert_positions_for_write(result);
   blender::MutableSpan<int> poly_offsets = result->poly_offsets_for_write();
   blender::MutableSpan<int> corner_verts = result->corner_verts_for_write();
@@ -3590,7 +3590,7 @@ static void fluid_modifier_processFlow(FluidModifierData *fmd,
     if (fmd->flow->mesh) {
       BKE_id_free(nullptr, fmd->flow->mesh);
     }
-    fmd->flow->mesh = BKE_mesh_copy_for_eval(me, false);
+    fmd->flow->mesh = BKE_mesh_copy_for_eval(me);
   }
 
   if (scene_framenr > fmd->time) {
@@ -3617,7 +3617,7 @@ static void fluid_modifier_processEffector(FluidModifierData *fmd,
     if (fmd->effector->mesh) {
       BKE_id_free(nullptr, fmd->effector->mesh);
     }
-    fmd->effector->mesh = BKE_mesh_copy_for_eval(me, false);
+    fmd->effector->mesh = BKE_mesh_copy_for_eval(me);
   }
 
   if (scene_framenr > fmd->time) {
@@ -4125,7 +4125,7 @@ Mesh *BKE_fluid_modifier_do(
   }
 
   if (!result) {
-    result = BKE_mesh_copy_for_eval(me, false);
+    result = BKE_mesh_copy_for_eval(me);
   }
   else {
     BKE_mesh_copy_parameters_for_eval(result, me);

source/blender/blenkernel/intern/geometry_component_curves.cc

@@ -33,7 +33,7 @@ GeometryComponent *CurveComponent::copy() const
 {
   CurveComponent *new_component = new CurveComponent();
   if (curves_ != nullptr) {
-    new_component->curves_ = BKE_curves_copy_for_eval(curves_, false);
+    new_component->curves_ = BKE_curves_copy_for_eval(curves_);
     new_component->ownership_ = GeometryOwnershipType::Owned;
   }
   return new_component;
@@ -87,7 +87,7 @@ Curves *CurveComponent::get_for_write()
 {
   BLI_assert(this->is_mutable());
   if (ownership_ == GeometryOwnershipType::ReadOnly) {
-    curves_ = BKE_curves_copy_for_eval(curves_, false);
+    curves_ = BKE_curves_copy_for_eval(curves_);
     ownership_ = GeometryOwnershipType::Owned;
   }
   return curves_;
@@ -107,7 +107,7 @@ void CurveComponent::ensure_owns_direct_data()
 {
   BLI_assert(this->is_mutable());
   if (ownership_ != GeometryOwnershipType::Owned) {
-    curves_ = BKE_curves_copy_for_eval(curves_, false);
+    curves_ = BKE_curves_copy_for_eval(curves_);
     ownership_ = GeometryOwnershipType::Owned;
   }
 }

source/blender/blenkernel/intern/geometry_component_mesh.cc

@@ -36,7 +36,7 @@ GeometryComponent *MeshComponent::copy() const
 {
   MeshComponent *new_component = new MeshComponent();
   if (mesh_ != nullptr) {
-    new_component->mesh_ = BKE_mesh_copy_for_eval(mesh_, false);
+    new_component->mesh_ = BKE_mesh_copy_for_eval(mesh_);
     new_component->ownership_ = GeometryOwnershipType::Owned;
   }
   return new_component;
@@ -83,7 +83,7 @@ Mesh *MeshComponent::get_for_write()
 {
   BLI_assert(this->is_mutable());
   if (ownership_ == GeometryOwnershipType::ReadOnly) {
-    mesh_ = BKE_mesh_copy_for_eval(mesh_, false);
+    mesh_ = BKE_mesh_copy_for_eval(mesh_);
     ownership_ = GeometryOwnershipType::Owned;
   }
   return mesh_;
@@ -103,7 +103,7 @@ void MeshComponent::ensure_owns_direct_data()
 {
   BLI_assert(this->is_mutable());
   if (ownership_ != GeometryOwnershipType::Owned) {
-    mesh_ = BKE_mesh_copy_for_eval(mesh_, false);
+    mesh_ = BKE_mesh_copy_for_eval(mesh_);
     ownership_ = GeometryOwnershipType::Owned;
   }
 }

source/blender/blenkernel/intern/geometry_component_pointcloud.cc

@@ -23,7 +23,7 @@ GeometryComponent *PointCloudComponent::copy() const
 {
   PointCloudComponent *new_component = new PointCloudComponent();
   if (pointcloud_ != nullptr) {
-    new_component->pointcloud_ = BKE_pointcloud_copy_for_eval(pointcloud_, false);
+    new_component->pointcloud_ = BKE_pointcloud_copy_for_eval(pointcloud_);
     new_component->ownership_ = GeometryOwnershipType::Owned;
   }
   return new_component;
@@ -70,7 +70,7 @@ PointCloud *PointCloudComponent::get_for_write()
 {
   BLI_assert(this->is_mutable());
   if (ownership_ == GeometryOwnershipType::ReadOnly) {
-    pointcloud_ = BKE_pointcloud_copy_for_eval(pointcloud_, false);
+    pointcloud_ = BKE_pointcloud_copy_for_eval(pointcloud_);
     ownership_ = GeometryOwnershipType::Owned;
   }
   return pointcloud_;
@@ -90,7 +90,7 @@ void PointCloudComponent::ensure_owns_direct_data()
 {
   BLI_assert(this->is_mutable());
   if (ownership_ != GeometryOwnershipType::Owned) {
-    pointcloud_ = BKE_pointcloud_copy_for_eval(pointcloud_, false);
+    pointcloud_ = BKE_pointcloud_copy_for_eval(pointcloud_);
     ownership_ = GeometryOwnershipType::Owned;
   }
 }

source/blender/blenkernel/intern/geometry_component_volume.cc

@@ -21,7 +21,7 @@ GeometryComponent *VolumeComponent::copy() const
 {
   VolumeComponent *new_component = new VolumeComponent();
   if (volume_ != nullptr) {
-    new_component->volume_ = BKE_volume_copy_for_eval(volume_, false);
+    new_component->volume_ = BKE_volume_copy_for_eval(volume_);
     new_component->ownership_ = GeometryOwnershipType::Owned;
   }
   return new_component;
@@ -68,7 +68,7 @@ Volume *VolumeComponent::get_for_write()
 {
   BLI_assert(this->is_mutable());
   if (ownership_ == GeometryOwnershipType::ReadOnly) {
-    volume_ = BKE_volume_copy_for_eval(volume_, false);
+    volume_ = BKE_volume_copy_for_eval(volume_);
     ownership_ = GeometryOwnershipType::Owned;
   }
   return volume_;
@@ -83,7 +83,7 @@ void VolumeComponent::ensure_owns_direct_data()
 {
   BLI_assert(this->is_mutable());
   if (ownership_ != GeometryOwnershipType::Owned) {
-    volume_ = BKE_volume_copy_for_eval(volume_, false);
+    volume_ = BKE_volume_copy_for_eval(volume_);
     ownership_ = GeometryOwnershipType::Owned;
   }
 }

source/blender/blenkernel/intern/mball_tessellate.cc

@@ -1451,18 +1451,17 @@ Mesh *BKE_mball_polygonize(Depsgraph *depsgraph, Scene *scene, Object *ob)
 
   freepolygonize(&process);
 
-  Mesh *mesh = (Mesh *)BKE_id_new_nomain(ID_ME, ((ID *)ob->data)->name + 2);
+  int corners_num = 0;
+  for (uint i = 0; i < process.curindex; i++) {
+    const int *indices = process.indices[i];
+    const int count = indices[2] != indices[3] ? 4 : 3;
+    corners_num += count;
+  }
 
-  mesh->totvert = int(process.co.size());
+  Mesh *mesh = BKE_mesh_new_nomain(int(process.co.size()), 0, int(process.curindex), corners_num);
-  CustomData_add_layer_named(
-      &mesh->vdata, CD_PROP_FLOAT3, CD_CONSTRUCT, mesh->totvert, "position");
   mesh->vert_positions_for_write().copy_from(process.co);
-
-  mesh->totpoly = int(process.curindex);
-  BKE_mesh_poly_offsets_ensure_alloc(mesh);
   blender::MutableSpan<int> poly_offsets = mesh->poly_offsets_for_write();
-  int *corner_verts = static_cast<int *>(CustomData_add_layer_named(
+  blender::MutableSpan<int> corner_verts = mesh->corner_verts_for_write();
-      &mesh->ldata, CD_PROP_INT32, CD_CONSTRUCT, mesh->totpoly * 4, ".corner_vert"));
 
   int loop_offset = 0;
   for (int i = 0; i < mesh->totpoly; i++) {
@@ -1490,9 +1489,6 @@ Mesh *BKE_mball_polygonize(Depsgraph *depsgraph, Scene *scene, Object *ob)
          sizeof(float[3]) * size_t(mesh->totvert));
   BKE_mesh_vert_normals_clear_dirty(mesh);
 
-  mesh->totloop = loop_offset;
-  poly_offsets.last() = loop_offset;
-
   BKE_mesh_calc_edges(mesh, false, false);
 
   return mesh;

source/blender/blenkernel/intern/mesh.cc

@@ -1032,23 +1032,19 @@ static void mesh_ensure_cdlayers_primary(Mesh &mesh)
   }
 }
 
-Mesh *BKE_mesh_new_nomain(int verts_len, int edges_len, int loops_len, int polys_len)
+Mesh *BKE_mesh_new_nomain(const int verts_num,
+                          const int edges_num,
+                          const int polys_num,
+                          const int loops_num)
 {
-  Mesh *mesh = (Mesh *)BKE_libblock_alloc(
+  Mesh *mesh = static_cast<Mesh *>(BKE_libblock_alloc(
-      nullptr, ID_ME, BKE_idtype_idcode_to_name(ID_ME), LIB_ID_CREATE_LOCALIZE);
+      nullptr, ID_ME, BKE_idtype_idcode_to_name(ID_ME), LIB_ID_CREATE_LOCALIZE));
   BKE_libblock_init_empty(&mesh->id);
 
-  /* Don't use #CustomData_reset because we don't want to touch custom-data. */
+  mesh->totvert = verts_num;
-  copy_vn_i(mesh->vdata.typemap, CD_NUMTYPES, -1);
+  mesh->totedge = edges_num;
-  copy_vn_i(mesh->edata.typemap, CD_NUMTYPES, -1);
+  mesh->totpoly = polys_num;
-  copy_vn_i(mesh->fdata.typemap, CD_NUMTYPES, -1);
+  mesh->totloop = loops_num;
-  copy_vn_i(mesh->ldata.typemap, CD_NUMTYPES, -1);
-  copy_vn_i(mesh->pdata.typemap, CD_NUMTYPES, -1);
-
-  mesh->totvert = verts_len;
-  mesh->totedge = edges_len;
-  mesh->totloop = loops_len;
-  mesh->totpoly = polys_len;
 
   mesh_ensure_cdlayers_primary(*mesh);
   BKE_mesh_poly_offsets_ensure_alloc(mesh);
@@ -1112,35 +1108,35 @@ void BKE_mesh_copy_parameters_for_eval(Mesh *me_dst, const Mesh *me_src)
 }
 
 Mesh *BKE_mesh_new_nomain_from_template_ex(const Mesh *me_src,
-                                           int verts_len,
+                                           const int verts_num,
-                                           int edges_len,
+                                           const int edges_num,
-                                           int tessface_len,
+                                           const int tessface_num,
-                                           int loops_len,
+                                           const int polys_num,
-                                           int polys_len,
+                                           const int loops_num,
-                                           CustomData_MeshMasks mask)
+                                           const CustomData_MeshMasks mask)
 {
   /* Only do tessface if we are creating tessfaces or copying from mesh with only tessfaces. */
-  const bool do_tessface = (tessface_len || ((me_src->totface != 0) && (me_src->totpoly == 0)));
+  const bool do_tessface = (tessface_num || ((me_src->totface != 0) && (me_src->totpoly == 0)));
 
   Mesh *me_dst = (Mesh *)BKE_id_new_nomain(ID_ME, nullptr);
 
   me_dst->mselect = (MSelect *)MEM_dupallocN(me_src->mselect);
 
-  me_dst->totvert = verts_len;
+  me_dst->totvert = verts_num;
-  me_dst->totedge = edges_len;
+  me_dst->totedge = edges_num;
-  me_dst->totface = tessface_len;
+  me_dst->totpoly = polys_num;
-  me_dst->totloop = loops_len;
+  me_dst->totloop = loops_num;
-  me_dst->totpoly = polys_len;
+  me_dst->totface = tessface_num;
 
   BKE_mesh_copy_parameters_for_eval(me_dst, me_src);
 
-  CustomData_copy_layout(&me_src->vdata, &me_dst->vdata, mask.vmask, CD_SET_DEFAULT, verts_len);
+  CustomData_copy_layout(&me_src->vdata, &me_dst->vdata, mask.vmask, CD_SET_DEFAULT, verts_num);
-  CustomData_copy_layout(&me_src->edata, &me_dst->edata, mask.emask, CD_SET_DEFAULT, edges_len);
+  CustomData_copy_layout(&me_src->edata, &me_dst->edata, mask.emask, CD_SET_DEFAULT, edges_num);
-  CustomData_copy_layout(&me_src->ldata, &me_dst->ldata, mask.lmask, CD_SET_DEFAULT, loops_len);
+  CustomData_copy_layout(&me_src->pdata, &me_dst->pdata, mask.pmask, CD_SET_DEFAULT, polys_num);
-  CustomData_copy_layout(&me_src->pdata, &me_dst->pdata, mask.pmask, CD_SET_DEFAULT, polys_len);
+  CustomData_copy_layout(&me_src->ldata, &me_dst->ldata, mask.lmask, CD_SET_DEFAULT, loops_num);
   if (do_tessface) {
     CustomData_copy_layout(
-        &me_src->fdata, &me_dst->fdata, mask.fmask, CD_SET_DEFAULT, tessface_len);
+        &me_src->fdata, &me_dst->fdata, mask.fmask, CD_SET_DEFAULT, tessface_num);
   }
   else {
     mesh_tessface_clear_intern(me_dst, false);
@@ -1160,11 +1156,14 @@ Mesh *BKE_mesh_new_nomain_from_template_ex(const Mesh *me_src,
   return me_dst;
 }
 
-Mesh *BKE_mesh_new_nomain_from_template(
+Mesh *BKE_mesh_new_nomain_from_template(const Mesh *me_src,
-    const Mesh *me_src, int verts_len, int edges_len, int loops_len, int polys_len)
+                                        const int verts_num,
+                                        const int edges_num,
+                                        const int polys_num,
+                                        const int loops_num)
 {
   return BKE_mesh_new_nomain_from_template_ex(
-      me_src, verts_len, edges_len, 0, loops_len, polys_len, CD_MASK_EVERYTHING);
+      me_src, verts_num, edges_num, 0, polys_num, loops_num, CD_MASK_EVERYTHING);
 }
 
 void BKE_mesh_eval_delete(struct Mesh *mesh_eval)
@@ -1176,16 +1175,10 @@ void BKE_mesh_eval_delete(struct Mesh *mesh_eval)
   MEM_freeN(mesh_eval);
 }
 
-Mesh *BKE_mesh_copy_for_eval(const Mesh *source, bool reference)
+Mesh *BKE_mesh_copy_for_eval(const Mesh *source)
 {
-  int flags = LIB_ID_COPY_LOCALIZE;
+  return reinterpret_cast<Mesh *>(
-
+      BKE_id_copy_ex(nullptr, &source->id, nullptr, LIB_ID_COPY_LOCALIZE));
-  if (reference) {
-    flags |= LIB_ID_COPY_CD_REFERENCE;
-  }
-
-  Mesh *result = (Mesh *)BKE_id_copy_ex(nullptr, &source->id, nullptr, flags);
-  return result;
 }
 
 BMesh *BKE_mesh_to_bmesh_ex(const Mesh *me,

source/blender/blenkernel/intern/mesh_boolean_convert.cc

@@ -700,7 +700,7 @@ static Mesh *imesh_to_mesh(IMesh *im, MeshesToIMeshInfo &mim)
   }
   /* Will calculate edges later. */
   Mesh *result = BKE_mesh_new_nomain_from_template(
-      mim.meshes[0], out_totvert, 0, out_totloop, out_totpoly);
+      mim.meshes[0], out_totvert, 0, out_totpoly, out_totloop);
 
   merge_vertex_loop_poly_customdata_layers(result, mim);
   /* Set the vertex coordinate values and other data. */

source/blender/blenkernel/intern/mesh_calc_edges.cc

@@ -151,14 +151,13 @@ static void serialize_and_initialize_deduplicated_edges(MutableSpan<EdgeMap> edg
   });
 }
 
-static void update_edge_indices_in_poly_loops(Mesh *mesh,
+static void update_edge_indices_in_poly_loops(const OffsetIndices<int> polys,
-                                              Span<EdgeMap> edge_maps,
+                                              const Span<int> corner_verts,
-                                              uint32_t parallel_mask)
+                                              const Span<EdgeMap> edge_maps,
+                                              const uint32_t parallel_mask,
+                                              MutableSpan<int> corner_edges)
 {
-  const OffsetIndices polys = mesh->polys();
+  threading::parallel_for(polys.index_range(), 100, [&](IndexRange range) {
-  const Span<int> corner_verts = mesh->corner_verts();
-  MutableSpan<int> corner_edges = mesh->corner_edges_for_write();
-  threading::parallel_for(IndexRange(mesh->totpoly), 100, [&](IndexRange range) {
     for (const int poly_index : range) {
       const IndexRange poly = polys[poly_index];
       int prev_corner = poly.last();
@@ -239,7 +238,11 @@ void BKE_mesh_calc_edges(Mesh *mesh, bool keep_existing_edges, const bool select
   MutableSpan<int2> new_edges{
       static_cast<int2 *>(MEM_calloc_arrayN(new_totedge, sizeof(int2), __func__)), new_totedge};
   calc_edges::serialize_and_initialize_deduplicated_edges(edge_maps, new_edges);
-  calc_edges::update_edge_indices_in_poly_loops(mesh, edge_maps, parallel_mask);
+  calc_edges::update_edge_indices_in_poly_loops(mesh->polys(),
+                                                mesh->corner_verts(),
+                                                edge_maps,
+                                                parallel_mask,
+                                                mesh->corner_edges_for_write());
 
   /* Free old CustomData and assign new one. */
   CustomData_free(&mesh->edata, mesh->totedge);

source/blender/blenkernel/intern/mesh_convert.cc

@@ -62,98 +62,8 @@ using blender::MutableSpan;
 using blender::Span;
 using blender::StringRefNull;
 
-/* Define for cases when you want extra validation of mesh
- * after certain modifications.
- */
-// #undef VALIDATE_MESH
-
-#ifdef VALIDATE_MESH
-#  define ASSERT_IS_VALID_MESH(mesh) \
-    (BLI_assert((mesh == nullptr) || (BKE_mesh_is_valid(mesh) == true)))
-#else
-#  define ASSERT_IS_VALID_MESH(mesh)
-#endif
-
 static CLG_LogRef LOG = {"bke.mesh_convert"};
 
-static void poly_edgehash_insert(EdgeHash *ehash, const Span<int> poly_verts)
-{
-  int i = poly_verts.size();
-
-  int next = 0;              /* first loop */
-  int poly_corner = (i - 1); /* last loop */
-
-  while (i-- != 0) {
-    BLI_edgehash_reinsert(ehash, poly_verts[poly_corner], poly_verts[next], nullptr);
-
-    poly_corner = next;
-    next++;
-  }
-}
-
-/**
- * Specialized function to use when we _know_ existing edges don't overlap with poly edges.
- */
-static void make_edges_mdata_extend(Mesh &mesh)
-{
-  int totedge = mesh.totedge;
-
-  const blender::OffsetIndices polys = mesh.polys();
-  const Span<int> corner_verts = mesh.corner_verts();
-  MutableSpan<int> corner_edges = mesh.corner_edges_for_write();
-
-  const int eh_reserve = max_ii(totedge, BLI_EDGEHASH_SIZE_GUESS_FROM_POLYS(mesh.totpoly));
-  EdgeHash *eh = BLI_edgehash_new_ex(__func__, eh_reserve);
-
-  for (const int i : polys.index_range()) {
-    poly_edgehash_insert(eh, corner_verts.slice(polys[i]));
-  }
-
-  const int totedge_new = BLI_edgehash_len(eh);
-
-#ifdef DEBUG
-  /* ensure that there's no overlap! */
-  if (totedge_new) {
-    for (const blender::int2 &edge : mesh.edges()) {
-      BLI_assert(BLI_edgehash_haskey(eh, edge[0], edge[1]) == false);
-    }
-  }
-#endif
-
-  if (totedge_new) {
-    /* The only layer should be edges, so no other layers need to be initialized. */
-    BLI_assert(mesh.edata.totlayer == 1);
-    CustomData_realloc(&mesh.edata, totedge, totedge + totedge_new);
-    mesh.totedge += totedge_new;
-    MutableSpan<blender::int2> edges = mesh.edges_for_write();
-    blender::int2 *edge = &edges[totedge];
-
-    EdgeHashIterator *ehi;
-    uint e_index = totedge;
-    for (ehi = BLI_edgehashIterator_new(eh); BLI_edgehashIterator_isDone(ehi) == false;
-         BLI_edgehashIterator_step(ehi), ++edge, e_index++) {
-      BLI_edgehashIterator_getKey(ehi, &(*edge)[0], &(*edge)[1]);
-      BLI_edgehashIterator_setValue(ehi, POINTER_FROM_UINT(e_index));
-    }
-    BLI_edgehashIterator_free(ehi);
-
-    for (const int i : polys.index_range()) {
-      const IndexRange poly = polys[i];
-      int corner = poly.start();
-      int corner_prev = poly.start() + (poly.size() - 1);
-      int j;
-      for (j = 0; j < poly.size(); j++, corner++) {
-        /* lookup hashed edge index */
-        corner_edges[corner_prev] = POINTER_AS_UINT(
-            BLI_edgehash_lookup(eh, corner_verts[corner_prev], corner_verts[corner]));
-        corner_prev = corner;
-      }
-    }
-  }
-
-  BLI_edgehash_free(eh, nullptr);
-}
-
 static Mesh *mesh_nurbs_displist_to_mesh(const Curve *cu, const ListBase *dispbase)
 {
   using namespace blender::bke;
@@ -203,7 +113,7 @@ static Mesh *mesh_nurbs_displist_to_mesh(const Curve *cu, const ListBase *dispba
     return BKE_mesh_new_nomain(0, 0, 0, 0);
   }
 
-  Mesh *mesh = BKE_mesh_new_nomain(totvert, totedge, totloop, totpoly);
+  Mesh *mesh = BKE_mesh_new_nomain(totvert, totedge, totpoly, totloop);
   MutableSpan<float3> positions = mesh->vert_positions_for_write();
   MutableSpan<blender::int2> edges = mesh->edges_for_write();
   MutableSpan<int> poly_offsets = mesh->poly_offsets_for_write();
@@ -393,7 +303,7 @@ static Mesh *mesh_nurbs_displist_to_mesh(const Curve *cu, const ListBase *dispba
   }
 
   if (totpoly) {
-    make_edges_mdata_extend(*mesh);
+    BKE_mesh_calc_edges(mesh, true, false);
   }
 
   material_indices.finish();
@@ -783,7 +693,7 @@ static const Curves *get_evaluated_curves_from_object(const Object *object)
 static Mesh *mesh_new_from_evaluated_curve_type_object(const Object *evaluated_object)
 {
   if (const Mesh *mesh = BKE_object_get_evaluated_mesh(evaluated_object)) {
-    return BKE_mesh_copy_for_eval(mesh, false);
+    return BKE_mesh_copy_for_eval(mesh);
   }
   if (const Curves *curves = get_evaluated_curves_from_object(evaluated_object)) {
     const blender::bke::AnonymousAttributePropagationInfo propagation_info;
@@ -842,7 +752,7 @@ static Mesh *mesh_new_from_mball_object(Object *object)
     return (Mesh *)BKE_id_new_nomain(ID_ME, ((ID *)object->data)->name + 2);
   }
 
-  return BKE_mesh_copy_for_eval(mesh_eval, false);
+  return BKE_mesh_copy_for_eval(mesh_eval);
 }
 
 static Mesh *mesh_new_from_mesh(Object *object, Mesh *mesh)

source/blender/blenkernel/intern/mesh_legacy_convert.cc

@@ -2197,7 +2197,10 @@ void BKE_mesh_legacy_attribute_strings_to_flags(Mesh *mesh)
 
   CustomData_clear_layer_flag(
       vdata, CD_PROP_BYTE_COLOR, CD_FLAG_COLOR_ACTIVE | CD_FLAG_COLOR_RENDER);
+  CustomData_clear_layer_flag(
+      ldata, CD_PROP_BYTE_COLOR, CD_FLAG_COLOR_ACTIVE | CD_FLAG_COLOR_RENDER);
   CustomData_clear_layer_flag(ldata, CD_PROP_COLOR, CD_FLAG_COLOR_ACTIVE | CD_FLAG_COLOR_RENDER);
+  CustomData_clear_layer_flag(vdata, CD_PROP_COLOR, CD_FLAG_COLOR_ACTIVE | CD_FLAG_COLOR_RENDER);
 
   if (const char *name = mesh->active_color_attribute) {
     int i;

source/blender/blenkernel/intern/mesh_mirror.cc

@@ -191,7 +191,7 @@ Mesh *BKE_mesh_mirror_apply_mirror_on_axis_for_modifier(MirrorModifierData *mmd,
   const int src_loops_num = mesh->totloop;
 
   Mesh *result = BKE_mesh_new_nomain_from_template(
-      mesh, src_verts_num * 2, src_edges_num * 2, src_loops_num * 2, src_polys.size() * 2);
+      mesh, src_verts_num * 2, src_edges_num * 2, src_polys.size() * 2, src_loops_num * 2);
 
   /* Copy custom-data to original geometry. */
   CustomData_copy_data(&mesh->vdata, &result->vdata, 0, 0, src_verts_num);

source/blender/blenkernel/intern/mesh_remesh_voxel.cc

@@ -117,7 +117,7 @@ static Mesh *remesh_quadriflow(const Mesh *input_mesh,
   }
 
   /* Construct the new output mesh */
-  Mesh *mesh = BKE_mesh_new_nomain(qrd.out_totverts, 0, qrd.out_totfaces * 4, qrd.out_totfaces);
+  Mesh *mesh = BKE_mesh_new_nomain(qrd.out_totverts, 0, qrd.out_totfaces, qrd.out_totfaces * 4);
   BKE_mesh_copy_parameters(mesh, input_mesh);
   MutableSpan<int> poly_offsets = mesh->poly_offsets_for_write();
   MutableSpan<int> corner_verts = mesh->corner_verts_for_write();
@@ -222,7 +222,7 @@ static Mesh *remesh_voxel_volume_to_mesh(const openvdb::FloatGrid::Ptr level_set
       *level_set_grid, vertices, tris, quads, isovalue, adaptivity, relax_disoriented_triangles);
 
   Mesh *mesh = BKE_mesh_new_nomain(
-      vertices.size(), 0, quads.size() * 4 + tris.size() * 3, quads.size() + tris.size());
+      vertices.size(), 0, quads.size() + tris.size(), quads.size() * 4 + tris.size() * 3);
   MutableSpan<float3> vert_positions = mesh->vert_positions_for_write();
   MutableSpan<int> poly_offsets = mesh->poly_offsets_for_write();
   MutableSpan<int> mesh_corner_verts = mesh->corner_verts_for_write();
@@ -317,8 +317,7 @@ void BKE_remesh_reproject_sculpt_face_sets(Mesh *target, const Mesh *source)
   const OffsetIndices target_polys = target->polys();
   const Span<int> target_corner_verts = target->corner_verts();
 
-  const VArray src_face_sets =
+  const VArray src_face_sets = *src_attributes.lookup<int>(".sculpt_face_set", ATTR_DOMAIN_FACE);
-    *  src_attributes.lookup<int>(".sculpt_face_set", ATTR_DOMAIN_FACE);
   if (!src_face_sets) {
     return;
   }

source/blender/blenkernel/intern/multires.cc

@@ -241,7 +241,7 @@ Mesh *BKE_multires_create_mesh(struct Depsgraph *depsgraph,
   Mesh *result = mti->modifyMesh(&mmd->modifier, &modifier_ctx, deformed_mesh);
 
   if (result == deformed_mesh) {
-    result = BKE_mesh_copy_for_eval(deformed_mesh, true);
+    result = BKE_mesh_copy_for_eval(deformed_mesh);
   }
   return result;
 }

source/blender/blenkernel/intern/pointcloud.cc

@@ -249,9 +249,7 @@ PointCloud *BKE_pointcloud_new_nomain(const int totpoint)
   return pointcloud;
 }
 
-void BKE_pointcloud_nomain_to_pointcloud(PointCloud *pointcloud_src,
+void BKE_pointcloud_nomain_to_pointcloud(PointCloud *pointcloud_src, PointCloud *pointcloud_dst)
-                                         PointCloud *pointcloud_dst,
-                                         bool take_ownership)
 {
   BLI_assert(pointcloud_src->id.tag & LIB_TAG_NO_MAIN);
 
@@ -260,9 +258,7 @@ void BKE_pointcloud_nomain_to_pointcloud(PointCloud *pointcloud_src,
   const int totpoint = pointcloud_dst->totpoint = pointcloud_src->totpoint;
   CustomData_copy(&pointcloud_src->pdata, &pointcloud_dst->pdata, CD_MASK_ALL, totpoint);
 
-  if (take_ownership) {
   BKE_id_free(nullptr, pointcloud_src);
-  }
 }
 
 bool PointCloud::bounds_min_max(blender::float3 &min, blender::float3 &max) const
@@ -322,16 +318,10 @@ bool BKE_pointcloud_attribute_required(const PointCloud * /*pointcloud*/, const
 
 /* Dependency Graph */
 
-PointCloud *BKE_pointcloud_copy_for_eval(struct PointCloud *pointcloud_src, bool reference)
+PointCloud *BKE_pointcloud_copy_for_eval(struct PointCloud *pointcloud_src)
 {
-  int flags = LIB_ID_COPY_LOCALIZE;
+  return reinterpret_cast<PointCloud *>(
-
+      BKE_id_copy_ex(nullptr, &pointcloud_src->id, nullptr, LIB_ID_COPY_LOCALIZE));
-  if (reference) {
-    flags |= LIB_ID_COPY_CD_REFERENCE;
-  }
-
-  PointCloud *result = (PointCloud *)BKE_id_copy_ex(nullptr, &pointcloud_src->id, nullptr, flags);
-  return result;
 }
 
 static void pointcloud_evaluate_modifiers(struct Depsgraph *depsgraph,

source/blender/blenkernel/intern/subdiv_mesh.cc

@@ -532,7 +532,7 @@ static bool subdiv_mesh_topology_info(const SubdivForeachContext *foreach_contex
 
   SubdivMeshContext *subdiv_context = static_cast<SubdivMeshContext *>(foreach_context->user_data);
   subdiv_context->subdiv_mesh = BKE_mesh_new_nomain_from_template_ex(
-      subdiv_context->coarse_mesh, num_vertices, num_edges, 0, num_loops, num_polygons, mask);
+      subdiv_context->coarse_mesh, num_vertices, num_edges, 0, num_polygons, num_loops, mask);
   subdiv_mesh_ctx_cache_custom_data_layers(subdiv_context);
   subdiv_mesh_prepare_accumulator(subdiv_context, num_vertices);
   subdiv_context->subdiv_mesh->runtime->subsurf_face_dot_tags.clear();

source/blender/blenkernel/intern/vfont.c

@@ -1143,8 +1143,7 @@ static bool vfont_to_curve(Object *ob,
   if (ef && ef->selboxes) {
     /* Set combined style flags for the selected string. Start with all styles then
      * remove one if ANY characters do not have it. Break out if we've removed them all. */
-    ef->select_char_info_flag = CU_CHINFO_BOLD | CU_CHINFO_ITALIC | CU_CHINFO_UNDERLINE |
+    ef->select_char_info_flag = CU_CHINFO_STYLE_ALL;
-                                CU_CHINFO_SMALLCAPS;
     for (int k = selstart; k <= selend && ef->select_char_info_flag; k++) {
       info = &custrinfo[k];
       ef->select_char_info_flag &= info->flag;

source/blender/blenkernel/intern/volume.cc

@@ -768,7 +768,7 @@ static void volume_filepath_get(const Main *bmain, const Volume *volume, char r_
   int path_frame, path_digits;
   if (volume->is_sequence && BLI_path_frame_get(r_filepath, &path_frame, &path_digits)) {
     char ext[32];
-    BLI_path_frame_strip(r_filepath, ext);
+    BLI_path_frame_strip(r_filepath, ext, sizeof(ext));
     BLI_path_frame(r_filepath, volume->runtime.frame, path_digits);
     BLI_path_extension_ensure(r_filepath, FILE_MAX, ext);
   }
@@ -1523,17 +1523,10 @@ Volume *BKE_volume_new_for_eval(const Volume *volume_src)
   return volume_dst;
 }
 
-Volume *BKE_volume_copy_for_eval(Volume *volume_src, bool reference)
+Volume *BKE_volume_copy_for_eval(Volume *volume_src)
 {
-  int flags = LIB_ID_COPY_LOCALIZE;
+  return reinterpret_cast<Volume *>(
-
+      BKE_id_copy_ex(nullptr, &volume_src->id, nullptr, LIB_ID_COPY_LOCALIZE));
-  if (reference) {
-    flags |= LIB_ID_COPY_CD_REFERENCE;
-  }
-
-  Volume *result = (Volume *)BKE_id_copy_ex(nullptr, &volume_src->id, nullptr, flags);
-
-  return result;
 }
 
 #ifdef WITH_OPENVDB

source/blender/blenkernel/intern/volume_to_mesh.cc

@@ -165,7 +165,7 @@ Mesh *volume_to_mesh(const openvdb::GridBase &grid,
 
   const int tot_loops = 3 * mesh_data.tris.size() + 4 * mesh_data.quads.size();
   const int tot_polys = mesh_data.tris.size() + mesh_data.quads.size();
-  Mesh *mesh = BKE_mesh_new_nomain(mesh_data.verts.size(), 0, tot_loops, tot_polys);
+  Mesh *mesh = BKE_mesh_new_nomain(mesh_data.verts.size(), 0, tot_polys, tot_loops);
 
   fill_mesh_from_openvdb_data(mesh_data.verts,
                               mesh_data.tris,

source/blender/blenlib/BLI_fileops.h

@@ -118,7 +118,7 @@ ENUM_OPERATORS(eFileAttributes, FILE_ATTR_HARDLINK);
  * \{ */
 
 typedef enum FileExternalOperation {
-  FILE_EXTERNAL_OPERATION_OPEN = 0,
+  FILE_EXTERNAL_OPERATION_OPEN = 1,
   FILE_EXTERNAL_OPERATION_FOLDER_OPEN,
   /* Following are Windows-only: */
   FILE_EXTERNAL_OPERATION_EDIT,

source/blender/blenlib/BLI_path_util.h

@@ -60,6 +60,12 @@ void BLI_split_dir_part(const char *string, char *dir, size_t dirlen);
  * Copies the leaf filename part of string into `file`, max length `filelen`.
  */
 void BLI_split_file_part(const char *string, char *file, size_t filelen);
+/**
+ * Returns a pointer to the last extension (e.g. the position of the last period).
+ * Returns a pointer to the nil byte when no extension is found.
+ */
+const char *BLI_path_extension_or_end(const char *filepath)
+    ATTR_NONNULL() ATTR_WARN_UNUSED_RESULT ATTR_RETURNS_NONNULL;
 /**
  * Returns a pointer to the last extension (e.g. the position of the last period).
  * Returns NULL if there is no extension.
@@ -416,7 +422,7 @@ bool BLI_path_frame_get(const char *path, int *r_frame, int *r_digits_len) ATTR_
  * So:      `/some/path_123.jpeg`
  * Becomes: `/some/path_###` with `r_ext` set to `.jpeg`.
  */
-void BLI_path_frame_strip(char *path, char *r_ext) ATTR_NONNULL();
+void BLI_path_frame_strip(char *path, char *r_ext, size_t ext_maxlen) ATTR_NONNULL();
 /**
  * Check if we have '#' chars, usable for #BLI_path_frame, #BLI_path_frame_range
  */

source/blender/blenlib/intern/path_util.c

@@ -59,9 +59,9 @@ int BLI_path_sequence_decode(const char *string, char *head, char *tail, ushort
   int i;
   bool found_digit = false;
   const char *const lslash = BLI_path_slash_rfind(string);
-  const char *const extension = BLI_path_extension(lslash ? lslash : string);
+  const char *const extension = BLI_path_extension_or_end(lslash ? lslash : string);
   const uint lslash_len = lslash != NULL ? (int)(lslash - string) : 0;
-  const uint name_end = extension != NULL ? (int)((extension - string)) : strlen(string);
+  const uint name_end = (uint)(extension - string);
 
   for (i = name_end - 1; i >= (int)lslash_len; i--) {
     if (isdigit(string[i])) {
@@ -587,17 +587,30 @@ bool BLI_path_suffix(char *string, size_t maxlen, const char *suffix, const char
 #ifdef DEBUG_STRSIZE
   memset(string, 0xff, sizeof(*string) * maxlen);
 #endif
-  const size_t string_len = strlen(string);
   const size_t suffix_len = strlen(suffix);
   const size_t sep_len = strlen(sep);
-  const char *const extension = BLI_path_extension(string);
+  char *extension = (char *)BLI_path_extension_or_end(string);
-  const size_t string_end = extension != NULL ? (int)((extension - string)) : string_len;
+  const size_t extension_len = strlen(extension);
-  char extension_copy[FILE_MAX];
+  const size_t string_end = extension - string;
+  const size_t string_len = string_end + extension_len;
   if (string_len + sep_len + suffix_len >= maxlen) {
     return false;
   }
-  BLI_strncpy(extension_copy, string + string_end, sizeof(extension));
+
-  BLI_sprintf(string + string_end, "%s%s%s", sep, suffix, extension_copy);
+  if (extension_len) {
+    memmove(extension + (sep_len + suffix_len), extension, extension_len);
+  }
+  char *c = string + string_end;
+  if (sep_len) {
+    memcpy(c, sep, sep_len);
+    c += sep_len;
+  }
+  if (suffix_len) {
+    memcpy(c, suffix, suffix_len);
+    c += suffix_len;
+  }
+  c += extension_len;
+  *c = '\0';
   return true;
 }
 
@@ -750,9 +763,8 @@ bool BLI_path_frame_get(const char *path, int *r_frame, int *r_digits_len)
   *r_digits_len = 0;
 
   const char *file = BLI_path_basename(path);
-  const char *file_ext = BLI_path_extension(file);
+  const char *file_ext = BLI_path_extension_or_end(file);
-  const int file_len = strlen(file);
+  const char *c = file_ext;
-  const char *c = file_ext ? file_ext : file + file_len;
 
   /* Find start of number (if there is one). */
   int digits_len = 0;
@@ -771,7 +783,7 @@ bool BLI_path_frame_get(const char *path, int *r_frame, int *r_digits_len)
   return true;
 }
 
-void BLI_path_frame_strip(char *path, char *r_ext)
+void BLI_path_frame_strip(char *path, char *r_ext, const size_t ext_maxlen)
 {
   *r_ext = '\0';
   if (*path == '\0') {
@@ -779,10 +791,8 @@ void BLI_path_frame_strip(char *path, char *r_ext)
   }
 
   char *file = (char *)BLI_path_basename(path);
-  char *file_ext = (char *)BLI_path_extension(file);
+  char *file_ext = (char *)BLI_path_extension_or_end(file);
-  const int file_len = strlen(file);
+  char *c = file_ext;
-  char *c = file_ext ? file_ext : file + file_len;
-  char *suffix = c;
 
   /* Find start of number (if there is one). */
   int digits_len = 0;
@@ -791,10 +801,9 @@ void BLI_path_frame_strip(char *path, char *r_ext)
   }
   c++;
 
-  int suffix_length = file_len - (suffix - file);
+  BLI_strncpy(r_ext, file_ext, ext_maxlen);
-  memcpy(r_ext, suffix, suffix_length + 1);
 
-  /* replace the number with the suffix and terminate the string */
+  /* Replace the number with the suffix and terminate the string. */
   while (digits_len--) {
     *c++ = '#';
   }
@@ -1264,11 +1273,7 @@ bool BLI_path_extension_replace(char *path, size_t maxlen, const char *ext)
 #ifdef DEBUG_STRSIZE
   memset(path, 0xff, sizeof(*path) * maxlen);
 #endif
-  char *path_ext = (char *)BLI_path_extension(path);
+  char *path_ext = (char *)BLI_path_extension_or_end(path);
-  if (path_ext == NULL) {
-    path_ext = path + strlen(path);
-  }
-
   const size_t ext_len = strlen(ext);
   if ((path_ext - path) + ext_len >= maxlen) {
     return false;
@@ -1293,9 +1298,10 @@ bool BLI_path_extension_ensure(char *path, size_t maxlen, const char *ext)
 #ifdef DEBUG_STRSIZE
   memset(path, 0xff, sizeof(*path) * maxlen);
 #endif
-  /* First check the extension is already there. */
+  /* First check the extension is already there.
-  char *path_ext = (char *)BLI_path_extension(path);
+   * If `path_ext` is the end of the string this is simply checking if `ext` is also empty. */
-  if (path_ext && STREQ(path_ext, ext)) {
+  const char *path_ext = BLI_path_extension_or_end(path);
+  if (STREQ(path_ext, ext)) {
     return true;
   }
 
@@ -1369,16 +1375,17 @@ void BLI_split_file_part(const char *string, char *file, const size_t filelen)
   BLI_split_dirfile(string, NULL, file, 0, filelen);
 }
 
-const char *BLI_path_extension(const char *filepath)
+const char *BLI_path_extension_or_end(const char *filepath)
 {
   /* NOTE(@ideasman42): Skip the extension when there are no preceding non-extension characters in
    * the file name. This ignores extensions at the beginning of a string or directly after a slash.
-   * Only using trailing extension characters has the advantage that stripping the extenion
+   * Only using trailing extension characters has the advantage that stripping the extension
    * never leads to a blank string (which can't be used as a file path).
    * Matches Python's `os.path.splitext`. */
   const char *ext = NULL;
   bool has_non_ext = false;
-  for (const char *c = filepath; *c; c++) {
+  const char *c = filepath;
+  for (; *c; c++) {
     switch (*c) {
       case '.': {
         if (has_non_ext) {
@@ -1398,7 +1405,17 @@ const char *BLI_path_extension(const char *filepath)
       }
     }
   }
+  if (ext) {
     return ext;
+  }
+  BLI_assert(*c == '\0');
+  return c;
+}
+
+const char *BLI_path_extension(const char *filepath)
+{
+  const char *ext = BLI_path_extension_or_end(filepath);
+  return *ext ? ext : NULL;
 }
 
 size_t BLI_path_append(char *__restrict dst, const size_t maxlen, const char *__restrict file)

source/blender/blenlib/tests/BLI_path_util_test.cc

@@ -42,7 +42,7 @@ static char *str_replace_char_strdup(const char *str, char src, char dst)
 /** \name Tests for: #BLI_path_normalize
  * \{ */
 
-#define NORMALIZE_WITH_BASEDIR(input, input_base, output) \
+#define NORMALIZE_WITH_BASEDIR(input, input_base, output_expect) \
   { \
     char path[FILE_MAX] = input; \
     const char *input_base_test = input_base; \
@@ -57,11 +57,11 @@ static char *str_replace_char_strdup(const char *str, char src, char dst)
         free((void *)input_base_test); \
       } \
     } \
-    EXPECT_STREQ(output, path); \
+    EXPECT_STREQ(path, output_expect); \
   } \
   ((void)0)
 
-#define NORMALIZE(input, output) NORMALIZE_WITH_BASEDIR(input, nullptr, output)
+#define NORMALIZE(input, output_expect) NORMALIZE_WITH_BASEDIR(input, nullptr, output_expect)
 
 /* #BLI_path_normalize: "/./" -> "/" */
 TEST(path_util, Clean_Dot)
@@ -102,7 +102,7 @@ TEST(path_util, Clean_Parent)
 /** \name Tests for: #BLI_path_parent_dir
  * \{ */
 
-#define PARENT_DIR(input, output) \
+#define PARENT_DIR(input, output_expect) \
   { \
     char path[FILE_MAX] = input; \
     if (SEP == '\\') { \
@@ -112,7 +112,7 @@ TEST(path_util, Clean_Parent)
     if (SEP == '\\') { \
       BLI_str_replace_char(path, '\\', '/'); \
     } \
-    EXPECT_STREQ(output, path); \
+    EXPECT_STREQ(path, output_expect); \
   } \
   ((void)0)
 
@@ -177,7 +177,7 @@ TEST(path_util, ParentDir_Complex)
     } \
     else { \
       EXPECT_TRUE(ret); \
-      EXPECT_EQ(strlen(expect), len_output); \
+      EXPECT_EQ(len_output, strlen(expect)); \
       path[index_output + len_output] = '\0'; \
       EXPECT_STREQ(&path[index_output], expect); \
     } \
@@ -515,42 +515,42 @@ TEST(path_util, Frame)
     char path[FILE_MAX] = "";
     ret = BLI_path_frame(path, 123, 1);
     EXPECT_TRUE(ret);
-    EXPECT_STREQ("123", path);
+    EXPECT_STREQ(path, "123");
   }
 
   {
     char path[FILE_MAX] = "";
     ret = BLI_path_frame(path, 123, 12);
     EXPECT_TRUE(ret);
-    EXPECT_STREQ("000000000123", path);
+    EXPECT_STREQ(path, "000000000123");
   }
 
   {
     char path[FILE_MAX] = "test_";
     ret = BLI_path_frame(path, 123, 1);
     EXPECT_TRUE(ret);
-    EXPECT_STREQ("test_123", path);
+    EXPECT_STREQ(path, "test_123");
   }
 
   {
     char path[FILE_MAX] = "test_";
     ret = BLI_path_frame(path, 1, 12);
     EXPECT_TRUE(ret);
-    EXPECT_STREQ("test_000000000001", path);
+    EXPECT_STREQ(path, "test_000000000001");
   }
 
   {
     char path[FILE_MAX] = "test_############";
     ret = BLI_path_frame(path, 1, 0);
     EXPECT_TRUE(ret);
-    EXPECT_STREQ("test_000000000001", path);
+    EXPECT_STREQ(path, "test_000000000001");
   }
 
   {
     char path[FILE_MAX] = "test_#_#_middle";
     ret = BLI_path_frame(path, 123, 0);
     EXPECT_TRUE(ret);
-    EXPECT_STREQ("test_#_123_middle", path);
+    EXPECT_STREQ(path, "test_#_123_middle");
   }
 
   /* intentionally fail */
@@ -558,14 +558,14 @@ TEST(path_util, Frame)
     char path[FILE_MAX] = "";
     ret = BLI_path_frame(path, 123, 0);
     EXPECT_FALSE(ret);
-    EXPECT_STREQ("", path);
+    EXPECT_STREQ(path, "");
   }
 
   {
     char path[FILE_MAX] = "test_middle";
     ret = BLI_path_frame(path, 123, 0);
     EXPECT_FALSE(ret);
-    EXPECT_STREQ("test_middle", path);
+    EXPECT_STREQ(path, "test_middle");
   }
 
   /* negative frame numbers */
@@ -573,13 +573,13 @@ TEST(path_util, Frame)
     char path[FILE_MAX] = "test_####";
     ret = BLI_path_frame(path, -1, 4);
     EXPECT_TRUE(ret);
-    EXPECT_STREQ("test_-0001", path);
+    EXPECT_STREQ(path, "test_-0001");
   }
   {
     char path[FILE_MAX] = "test_####";
     ret = BLI_path_frame(path, -100, 4);
     EXPECT_TRUE(ret);
-    EXPECT_STREQ("test_-0100", path);
+    EXPECT_STREQ(path, "test_-0100");
   }
 }
 
@@ -595,52 +595,52 @@ TEST(path_util, SplitDirfile)
     const char *path = "";
     char dir[FILE_MAX], file[FILE_MAX];
     BLI_split_dirfile(path, dir, file, sizeof(dir), sizeof(file));
-    EXPECT_STREQ("", dir);
+    EXPECT_STREQ(dir, "");
-    EXPECT_STREQ("", file);
+    EXPECT_STREQ(file, "");
   }
 
   {
     const char *path = "/";
     char dir[FILE_MAX], file[FILE_MAX];
     BLI_split_dirfile(path, dir, file, sizeof(dir), sizeof(file));
-    EXPECT_STREQ("/", dir);
+    EXPECT_STREQ(dir, "/");
-    EXPECT_STREQ("", file);
+    EXPECT_STREQ(file, "");
   }
 
   {
     const char *path = "fileonly";
     char dir[FILE_MAX], file[FILE_MAX];
     BLI_split_dirfile(path, dir, file, sizeof(dir), sizeof(file));
-    EXPECT_STREQ("", dir);
+    EXPECT_STREQ(dir, "");
-    EXPECT_STREQ("fileonly", file);
+    EXPECT_STREQ(file, "fileonly");
   }
 
   {
     const char *path = "dironly/";
     char dir[FILE_MAX], file[FILE_MAX];
     BLI_split_dirfile(path, dir, file, sizeof(dir), sizeof(file));
-    EXPECT_STREQ("dironly/", dir);
+    EXPECT_STREQ(dir, "dironly/");
-    EXPECT_STREQ("", file);
+    EXPECT_STREQ(file, "");
   }
 
   {
     const char *path = "/a/b";
     char dir[FILE_MAX], file[FILE_MAX];
     BLI_split_dirfile(path, dir, file, sizeof(dir), sizeof(file));
-    EXPECT_STREQ("/a/", dir);
+    EXPECT_STREQ(dir, "/a/");
-    EXPECT_STREQ("b", file);
+    EXPECT_STREQ(file, "b");
   }
 
   {
     const char *path = "/dirtoobig/filetoobig";
     char dir[5], file[5];
     BLI_split_dirfile(path, dir, file, sizeof(dir), sizeof(file));
-    EXPECT_STREQ("/dir", dir);
+    EXPECT_STREQ(dir, "/dir");
-    EXPECT_STREQ("file", file);
+    EXPECT_STREQ(file, "file");
 
     BLI_split_dirfile(path, dir, file, 1, 1);
-    EXPECT_STREQ("", dir);
+    EXPECT_STREQ(dir, "");
-    EXPECT_STREQ("", file);
+    EXPECT_STREQ(file, "");
   }
 }
 
@@ -655,13 +655,13 @@ TEST(path_util, SplitDirfile)
     char path[FILE_MAX]; \
     char ext[FILE_MAX]; \
     BLI_strncpy(path, (input_path), FILE_MAX); \
-    BLI_path_frame_strip(path, ext); \
+    BLI_path_frame_strip(path, ext, sizeof(ext)); \
     EXPECT_STREQ(path, expect_path); \
     EXPECT_STREQ(ext, expect_ext); \
   } \
   ((void)0)
 
-TEST(path_util, PathFrameStrip)
+TEST(path_util, FrameStrip)
 {
   PATH_FRAME_STRIP("", "", "");
   PATH_FRAME_STRIP("nonum.abc", "nonum", ".abc");
@@ -678,29 +678,29 @@ TEST(path_util, PathFrameStrip)
 /** \name Tests for: #BLI_path_extension
  * \{ */
 
-TEST(path_util, PathExtension)
+TEST(path_util, Extension)
 {
-  EXPECT_EQ(nullptr, BLI_path_extension("some.def/file"));
+  EXPECT_EQ(BLI_path_extension("some.def/file"), nullptr);
-  EXPECT_EQ(nullptr, BLI_path_extension("Text"));
+  EXPECT_EQ(BLI_path_extension("Text"), nullptr);
-  EXPECT_EQ(nullptr, BLI_path_extension("Text…001"));
+  EXPECT_EQ(BLI_path_extension("Text…001"), nullptr);
-  EXPECT_EQ(nullptr, BLI_path_extension(".hidden"));
+  EXPECT_EQ(BLI_path_extension(".hidden"), nullptr);
-  EXPECT_EQ(nullptr, BLI_path_extension(".hidden/"));
+  EXPECT_EQ(BLI_path_extension(".hidden/"), nullptr);
-  EXPECT_EQ(nullptr, BLI_path_extension("/.hidden"));
+  EXPECT_EQ(BLI_path_extension("/.hidden"), nullptr);
-  EXPECT_EQ(nullptr, BLI_path_extension("dir/.hidden"));
+  EXPECT_EQ(BLI_path_extension("dir/.hidden"), nullptr);
-  EXPECT_EQ(nullptr, BLI_path_extension("/dir/.hidden"));
+  EXPECT_EQ(BLI_path_extension("/dir/.hidden"), nullptr);
 
-  EXPECT_EQ(nullptr, BLI_path_extension("."));
+  EXPECT_EQ(BLI_path_extension("."), nullptr);
-  EXPECT_EQ(nullptr, BLI_path_extension(".."));
+  EXPECT_EQ(BLI_path_extension(".."), nullptr);
-  EXPECT_EQ(nullptr, BLI_path_extension("..."));
+  EXPECT_EQ(BLI_path_extension("..."), nullptr);
-  EXPECT_STREQ(".", BLI_path_extension("...a."));
+  EXPECT_STREQ(BLI_path_extension("...a."), ".");
-  EXPECT_STREQ(".", BLI_path_extension("...a.."));
+  EXPECT_STREQ(BLI_path_extension("...a.."), ".");
-  EXPECT_EQ(nullptr, BLI_path_extension("...a../"));
+  EXPECT_EQ(BLI_path_extension("...a../"), nullptr);
 
-  EXPECT_STREQ(".", BLI_path_extension("some/file."));
+  EXPECT_STREQ(BLI_path_extension("some/file."), ".");
-  EXPECT_STREQ(".gz", BLI_path_extension("some/file.tar.gz"));
+  EXPECT_STREQ(BLI_path_extension("some/file.tar.gz"), ".gz");
-  EXPECT_STREQ(".abc", BLI_path_extension("some.def/file.abc"));
+  EXPECT_STREQ(BLI_path_extension("some.def/file.abc"), ".abc");
-  EXPECT_STREQ(".abc", BLI_path_extension("C:\\some.def\\file.abc"));
+  EXPECT_STREQ(BLI_path_extension("C:\\some.def\\file.abc"), ".abc");
-  EXPECT_STREQ(".001", BLI_path_extension("Text.001"));
+  EXPECT_STREQ(BLI_path_extension("Text.001"), ".001");
 }
 
 /** \} */
@@ -721,7 +721,7 @@ TEST(path_util, PathExtension)
   } \
   ((void)0)
 
-TEST(path_util, PathExtensionCheck)
+TEST(path_util, ExtensionCheck)
 {
   PATH_EXTENSION_CHECK("a/b/c.exe", ".exe", ".exe");
   PATH_EXTENSION_CHECK("correct/path/to/file.h", ".h", ".h");
@@ -752,26 +752,31 @@ TEST(path_util, PathExtensionCheck)
 /** \name Tests for: #BLI_path_extension_replace
  * \{ */
 
-#define PATH_EXTENSION_REPLACE(input_path, input_ext, expect_result, expect_path) \
+#define PATH_EXTENSION_REPLACE_WITH_MAXLEN( \
+    input_path, input_ext, expect_result, expect_path, maxlen) \
   { \
+    BLI_assert(maxlen <= FILE_MAX); \
     char path[FILE_MAX]; \
-    BLI_strncpy(path, input_path, FILE_MAX); \
+    BLI_strncpy(path, input_path, sizeof(path)); \
-    const bool ret = BLI_path_extension_replace(path, sizeof(path), input_ext); \
+    const bool ret = BLI_path_extension_replace(path, maxlen, input_ext); \
     if (expect_result) { \
       EXPECT_TRUE(ret); \
     } \
     else { \
       EXPECT_FALSE(ret); \
     } \
-    EXPECT_STREQ(expect_path, path); \
+    EXPECT_STREQ(path, expect_path); \
   } \
   ((void)0)
 
-TEST(path_util, PathExtensionReplace)
+#define PATH_EXTENSION_REPLACE(input_path, input_ext, expect_result, expect_path) \
+  PATH_EXTENSION_REPLACE_WITH_MAXLEN(input_path, input_ext, expect_result, expect_path, FILE_MAX)
+
+TEST(path_util, ExtensionReplace)
 {
   PATH_EXTENSION_REPLACE("test", ".txt", true, "test.txt");
   PATH_EXTENSION_REPLACE("test.", ".txt", true, "test.txt");
-  /* Unlike #BLI_path_extension_ensure, exceeds '.' are not stripped. */
+  /* Unlike #BLI_path_extension_ensure, excess '.' are not stripped. */
   PATH_EXTENSION_REPLACE("test..", ".txt", true, "test..txt");
 
   PATH_EXTENSION_REPLACE("test.txt", ".txt", true, "test.txt");
@@ -780,15 +785,37 @@ TEST(path_util, PathExtensionReplace)
   PATH_EXTENSION_REPLACE("test", "_txt", true, "test_txt");
   PATH_EXTENSION_REPLACE("test.ext", "_txt", true, "test_txt");
 
+  PATH_EXTENSION_REPLACE("test", "", true, "test");
+
+  /* Same as #BLI_path_extension_strip. */
+  PATH_EXTENSION_REPLACE("test.txt", "", true, "test");
+
+  /* Empty strings. */
   PATH_EXTENSION_REPLACE("test", "", true, "test");
   PATH_EXTENSION_REPLACE("", "_txt", true, "_txt");
+  PATH_EXTENSION_REPLACE("", "", true, "");
 
   /* Ensure leading '.' isn't treated as an extension. */
   PATH_EXTENSION_REPLACE(".hidden", ".hidden", true, ".hidden.hidden");
   PATH_EXTENSION_REPLACE("..hidden", ".hidden", true, "..hidden.hidden");
   PATH_EXTENSION_REPLACE("._.hidden", ".hidden", true, "._.hidden");
 }
+
+TEST(path_util, ExtensionReplace_Overflow)
+{
+  /* Small values. */
+  PATH_EXTENSION_REPLACE_WITH_MAXLEN("test", ".txt", false, "test", 0);
+  PATH_EXTENSION_REPLACE_WITH_MAXLEN("test", ".txt", false, "test", 1);
+  /* One under fails, and exactly enough space succeeds. */
+  PATH_EXTENSION_REPLACE_WITH_MAXLEN("test", ".txt", false, "test", 8);
+  PATH_EXTENSION_REPLACE_WITH_MAXLEN("test", ".txt", true, "test.txt", 9);
+
+  PATH_EXTENSION_REPLACE_WITH_MAXLEN("test.xx", ".txt", false, "test.xx", 8);
+  PATH_EXTENSION_REPLACE_WITH_MAXLEN("test.xx", ".txt", true, "test.txt", 9);
+}
+
 #undef PATH_EXTENSION_REPLACE
+#undef PATH_EXTENSION_REPLACE_WITH_MAXLEN
 
 /** \} */
 
@@ -796,22 +823,27 @@ TEST(path_util, PathExtensionReplace)
 /** \name Tests for: #BLI_path_extension_ensure
  * \{ */
 
-#define PATH_EXTENSION_ENSURE(input_path, input_ext, expect_result, expect_path) \
+#define PATH_EXTENSION_ENSURE_WITH_MAXLEN( \
+    input_path, input_ext, expect_result, expect_path, maxlen) \
   { \
+    BLI_assert(maxlen <= FILE_MAX); \
     char path[FILE_MAX]; \
-    BLI_strncpy(path, input_path, FILE_MAX); \
+    BLI_strncpy(path, input_path, sizeof(path)); \
-    const bool ret = BLI_path_extension_ensure(path, sizeof(path), input_ext); \
+    const bool ret = BLI_path_extension_ensure(path, maxlen, input_ext); \
     if (expect_result) { \
       EXPECT_TRUE(ret); \
     } \
     else { \
       EXPECT_FALSE(ret); \
     } \
-    EXPECT_STREQ(expect_path, path); \
+    EXPECT_STREQ(path, expect_path); \
   } \
   ((void)0)
 
-TEST(path_util, PathExtensionEnsure)
+#define PATH_EXTENSION_ENSURE(input_path, input_ext, expect_result, expect_path) \
+  PATH_EXTENSION_ENSURE_WITH_MAXLEN(input_path, input_ext, expect_result, expect_path, FILE_MAX)
+
+TEST(path_util, ExtensionEnsure)
 {
   PATH_EXTENSION_ENSURE("test", ".txt", true, "test.txt");
   PATH_EXTENSION_ENSURE("test.", ".txt", true, "test.txt");
@@ -823,15 +855,32 @@ TEST(path_util, PathExtensionEnsure)
   PATH_EXTENSION_ENSURE("test", "_txt", true, "test_txt");
   PATH_EXTENSION_ENSURE("test.ext", "_txt", true, "test.ext_txt");
 
+  /* An empty string does nothing (unlike replace which strips). */
+  PATH_EXTENSION_ENSURE("test.txt", "", true, "test.txt");
+
+  /* Empty strings. */
   PATH_EXTENSION_ENSURE("test", "", true, "test");
   PATH_EXTENSION_ENSURE("", "_txt", true, "_txt");
+  PATH_EXTENSION_ENSURE("", "", true, "");
 
   /* Ensure leading '.' isn't treated as an extension. */
   PATH_EXTENSION_ENSURE(".hidden", ".hidden", true, ".hidden.hidden");
   PATH_EXTENSION_ENSURE("..hidden", ".hidden", true, "..hidden.hidden");
   PATH_EXTENSION_ENSURE("._.hidden", ".hidden", true, "._.hidden");
 }
+
+TEST(path_util, ExtensionEnsure_Overflow)
+{
+  /* Small values. */
+  PATH_EXTENSION_ENSURE_WITH_MAXLEN("test", ".txt", false, "test", 0);
+  PATH_EXTENSION_ENSURE_WITH_MAXLEN("test", ".txt", false, "test", 1);
+  /* One under fails, and exactly enough space succeeds. */
+  PATH_EXTENSION_ENSURE_WITH_MAXLEN("test", ".txt", false, "test", 8);
+  PATH_EXTENSION_ENSURE_WITH_MAXLEN("test", ".txt", true, "test.txt", 9);
+}
+
 #undef PATH_EXTENSION_ENSURE
+#undef PATH_EXTENSION_ENSURE_WITH_MAXLEN
 
 /** \} */
 
@@ -851,7 +900,7 @@ TEST(path_util, PathExtensionEnsure)
   } \
   ((void)0)
 
-TEST(path_util, PathFrameCheckChars)
+TEST(path_util, FrameCheckChars)
 {
   PATH_FRAME_CHECK_CHARS("a#", true);
   PATH_FRAME_CHECK_CHARS("aaaaa#", true);
@@ -892,7 +941,7 @@ TEST(path_util, PathFrameCheckChars)
   } \
   ((void)0)
 
-TEST(path_util, PathFrameRange)
+TEST(path_util, FrameRange)
 {
   int dummy = -1;
   PATH_FRAME_RANGE("#", 1, 2, dummy, "1-2");
@@ -930,7 +979,7 @@ TEST(path_util, PathFrameRange)
   } \
   ((void)0)
 
-TEST(path_util, PathFrameGet)
+TEST(path_util, FrameGet)
 {
   PATH_FRAME_GET("001.avi", 1, 3, true);
   PATH_FRAME_GET("0000299.ext", 299, 7, true);
@@ -954,14 +1003,14 @@ TEST(path_util, PathFrameGet)
     char tail[FILE_MAX]; \
     ushort numdigits = 0; \
     const int result = BLI_path_sequence_decode(path, head, tail, &numdigits); \
-    EXPECT_EQ(expect_result, result); \
+    EXPECT_EQ(result, expect_result); \
-    EXPECT_STREQ(expect_head, head); \
+    EXPECT_STREQ(head, expect_head); \
-    EXPECT_STREQ(expect_tail, tail); \
+    EXPECT_STREQ(tail, expect_tail); \
-    EXPECT_EQ(expect_numdigits, numdigits); \
+    EXPECT_EQ(numdigits, expect_numdigits); \
   } \
   (void)0;
 
-TEST(path_util, PathSequenceDecode)
+TEST(path_util, SequenceDecode)
 {
   /* Basic use. */
   PATH_SEQ_DECODE("file_123.txt", 123, "file_", ".txt", 3);
@@ -987,12 +1036,12 @@ TEST(path_util, PathSequenceDecode)
   { \
     char path[FILE_MAX] = path_literal; \
     const bool result = BLI_path_suffix(path, path_literal_max, suffix, sep); \
-    EXPECT_EQ(expect_result, result); \
+    EXPECT_EQ(result, expect_result); \
-    EXPECT_STREQ(expect_path, path); \
+    EXPECT_STREQ(path, expect_path); \
   } \
   (void)0;
 
-TEST(path_util, PathSuffix)
+TEST(path_util, Suffix)
 {
   /* Extension. */
   PATH_SUFFIX("file.txt", FILE_MAX, "_", "123", true, "file_123.txt");
@@ -1010,6 +1059,11 @@ TEST(path_util, PathSuffix)
   PATH_SUFFIX("", FILE_MAX, "_", "123", true, "_123");
   /* Empty input/output. */
   PATH_SUFFIX("", FILE_MAX, "", "", true, "");
+
+  /* Long suffix. */
+  PATH_SUFFIX("file.txt", FILE_MAX, "_", "1234567890", true, "file_1234567890.txt");
+  /* Long extension. */
+  PATH_SUFFIX("file.txt1234567890", FILE_MAX, "_", "123", true, "file_123.txt1234567890");
 }
 
 #undef PATH_SUFFIX
@@ -1020,7 +1074,7 @@ TEST(path_util, PathSuffix)
 /** \name Tests for: #BLI_path_rel
  * \{ */
 
-#define PATH_REL(abs_path, ref_path, rel_path) \
+#define PATH_REL(abs_path, ref_path, rel_path_expect) \
   { \
     char path[FILE_MAX]; \
     const char *ref_path_test = ref_path; \
@@ -1034,7 +1088,7 @@ TEST(path_util, PathSuffix)
       BLI_str_replace_char(path, '\\', '/'); \
       free((void *)ref_path_test); \
     } \
-    EXPECT_STREQ(rel_path, path); \
+    EXPECT_STREQ(path, rel_path_expect); \
   } \
   void(0)
 
@@ -1044,17 +1098,17 @@ TEST(path_util, PathSuffix)
 #  define ABS_PREFIX ""
 #endif
 
-TEST(path_util, PathRelPath_Simple)
+TEST(path_util, RelPath_Simple)
 {
   PATH_REL(ABS_PREFIX "/foo/bar/blender.blend", ABS_PREFIX "/foo/bar/", "//blender.blend");
 }
 
-TEST(path_util, PathRelPath_SimpleSubdir)
+TEST(path_util, RelPath_SimpleSubdir)
 {
   PATH_REL(ABS_PREFIX "/foo/bar/blender.blend", ABS_PREFIX "/foo/bar", "//bar/blender.blend");
 }
 
-TEST(path_util, PathRelPath_BufferOverflowRoot)
+TEST(path_util, RelPath_BufferOverflowRoot)
 {
   char abs_path_in[FILE_MAX];
   const char *abs_prefix = ABS_PREFIX "/";
@@ -1070,7 +1124,7 @@ TEST(path_util, PathRelPath_BufferOverflowRoot)
   PATH_REL(abs_path_in, abs_prefix, abs_path_out);
 }
 
-TEST(path_util, PathRelPath_BufferOverflowSubdir)
+TEST(path_util, RelPath_BufferOverflowSubdir)
 {
   char abs_path_in[FILE_MAX];
   const char *ref_path_in = ABS_PREFIX "/foo/bar/";
@@ -1096,7 +1150,7 @@ TEST(path_util, PathRelPath_BufferOverflowSubdir)
 /** \name Tests for: #BLI_path_contains
  * \{ */
 
-TEST(path_util, PathContains)
+TEST(path_util, Contains)
 {
   EXPECT_TRUE(BLI_path_contains("/some/path", "/some/path")) << "A path contains itself";
   EXPECT_TRUE(BLI_path_contains("/some/path", "/some/path/inside"))
@@ -1119,7 +1173,7 @@ TEST(path_util, PathContains)
 }
 
 #ifdef WIN32
-TEST(path_util, PathContains_Windows_case_insensitive)
+TEST(path_util, Contains_Windows_case_insensitive)
 {
   EXPECT_TRUE(BLI_path_contains("C:\\some\\path", "c:\\SOME\\path\\inside"))
       << "On Windows path comparison should ignore case";

source/blender/blenloader/intern/versioning_300.cc

@@ -4305,6 +4305,9 @@ void blo_do_versions_300(FileData *fd, Library * /*lib*/, Main *bmain)
                                             ARRAY_SIZE(tool_replace_table));
       }
     }
+
+    /* Rename Grease Pencil weight draw brush. */
+    do_versions_rename_id(bmain, ID_BR, "Draw Weight", "Weight Draw");
   }
 
   /**

source/blender/compositor/realtime_compositor/CMakeLists.txt

@@ -118,6 +118,7 @@ set(GLSL_SRC
   shaders/compositor_glare_simple_star_vertical_pass.glsl
   shaders/compositor_glare_streaks_accumulate.glsl
   shaders/compositor_glare_streaks_filter.glsl
+  shaders/compositor_id_mask.glsl
   shaders/compositor_image_crop.glsl
   shaders/compositor_map_uv.glsl
   shaders/compositor_morphological_distance.glsl
@@ -215,6 +216,7 @@ set(SRC_SHADER_CREATE_INFOS
   shaders/infos/compositor_filter_info.hh
   shaders/infos/compositor_flip_info.hh
   shaders/infos/compositor_glare_info.hh
+  shaders/infos/compositor_id_mask_info.hh
   shaders/infos/compositor_image_crop_info.hh
   shaders/infos/compositor_map_uv_info.hh
   shaders/infos/compositor_morphological_distance_feather_info.hh

source/blender/compositor/realtime_compositor/intern/shader_node.cc

@@ -98,6 +98,7 @@ static void gpu_stack_vector_from_socket(GPUNodeStack &stack, const bNodeSocket
           return;
         default:
           BLI_assert_unreachable();
+          return;
       }
     }
     case SOCK_VECTOR: {
@@ -115,6 +116,7 @@ static void gpu_stack_vector_from_socket(GPUNodeStack &stack, const bNodeSocket
           return;
         default:
           BLI_assert_unreachable();
+          return;
       }
     }
     case SOCK_RGBA: {
@@ -131,6 +133,7 @@ static void gpu_stack_vector_from_socket(GPUNodeStack &stack, const bNodeSocket
           return;
         default:
           BLI_assert_unreachable();
+          return;
       }
     }
     default:

source/blender/compositor/realtime_compositor/shaders/compositor_id_mask.glsl

@@ -0,0 +1,11 @@
+#pragma BLENDER_REQUIRE(gpu_shader_compositor_texture_utilities.glsl)
+
+void main()
+{
+  ivec2 texel = ivec2(gl_GlobalInvocationID.xy);
+
+  float input_mask_value = texture_load(input_mask_tx, texel).x;
+  float mask = int(round(input_mask_value)) == index ? 1.0 : 0.0;
+
+  imageStore(output_mask_img, texel, vec4(mask));
+}

source/blender/compositor/realtime_compositor/shaders/infos/compositor_id_mask_info.hh

@@ -0,0 +1,11 @@
+/* SPDX-License-Identifier: GPL-2.0-or-later */
+
+#include "gpu_shader_create_info.hh"
+
+GPU_SHADER_CREATE_INFO(compositor_id_mask)
+    .local_group_size(16, 16)
+    .push_constant(Type::INT, "index")
+    .sampler(0, ImageType::FLOAT_2D, "input_mask_tx")
+    .image(0, GPU_R16F, Qualifier::WRITE, ImageType::FLOAT_2D, "output_mask_img")
+    .compute_source("compositor_id_mask.glsl")
+    .do_static_compilation(true);

source/blender/depsgraph/intern/builder/deg_builder_key.cc

@@ -84,6 +84,19 @@ RNAPathKey::RNAPathKey(ID *id, const PointerRNA &ptr, PropertyRNA *prop, RNAPoin
 {
 }
 
+RNAPathKey::RNAPathKey(const PointerRNA &target_prop,
+                       const char *rna_path_from_target_prop,
+                       const RNAPointerSource source)
+    : id(target_prop.owner_id), source(source)
+{
+  /* Try to resolve path. */
+  int index;
+  if (!RNA_path_resolve_full(&target_prop, rna_path_from_target_prop, &ptr, &prop, &index)) {
+    ptr = PointerRNA_NULL;
+    prop = nullptr;
+  }
+}
+
 string RNAPathKey::identifier() const
 {
   const char *id_name = (id) ? id->name : "<No ID>";

source/blender/depsgraph/intern/builder/deg_builder_key.h

@@ -188,6 +188,9 @@ struct PersistentOperationKey : public OperationKey {
 
 struct RNAPathKey {
   RNAPathKey(ID *id, const char *path, RNAPointerSource source);
+  RNAPathKey(const PointerRNA &target_prop,
+             const char *rna_path_from_target_prop,
+             RNAPointerSource source);
   RNAPathKey(ID *id, const PointerRNA &ptr, PropertyRNA *prop, RNAPointerSource source);
 
   string identifier() const;

source/blender/depsgraph/intern/builder/deg_builder_relations.cc

@@ -1807,7 +1807,7 @@ void DepsgraphRelationBuilder::build_driver_variables(ID *id, FCurve *fcu)
         add_relation(target_key, driver_key, "Target -> Driver");
       }
       else if (dtar->rna_path != nullptr && dtar->rna_path[0] != '\0') {
-        RNAPathKey variable_exit_key(target_id, dtar->rna_path, RNAPointerSource::EXIT);
+        RNAPathKey variable_exit_key(target_prop, dtar->rna_path, RNAPointerSource::EXIT);
         if (RNA_pointer_is_null(&variable_exit_key.ptr)) {
           continue;
         }

source/blender/draw/engines/eevee_next/shaders/eevee_depth_of_field_accumulator_lib.glsl

@@ -590,7 +590,7 @@ void dof_gather_accumulator(sampler2D color_tx,
  * The full pixel neighborhood is gathered.
  * \{ */
 
-void dof_slight_focus_gather(sampler2D depth_tx,
+void dof_slight_focus_gather(depth2D depth_tx,
                              sampler2D color_tx,
                              sampler2D bkh_lut_tx, /* Renamed because of ugly macro job. */
                              float radius,

source/blender/draw/engines/eevee_next/shaders/eevee_hiz_update_comp.glsl

@@ -62,7 +62,7 @@ void main()
   int mask_shift = 1;
 
 #define downsample_level(out_mip__, lod_) \
-  active_thread = all(lessThan(local_px, gl_WorkGroupSize.xy >> uint(mask_shift))); \
+  active_thread = all(lessThan(uvec2(local_px), gl_WorkGroupSize.xy >> uint(mask_shift))); \
   barrier(); /* Wait for previous writes to finish. */ \
   if (active_thread) { \
     max_depth = max_v4(load_local_depths(local_px)); \
@@ -89,12 +89,12 @@ void main()
   }
   finished_tile_counter = 0u;
 
-  ivec2 iter = divide_ceil(imageSize(out_mip_5), ivec2(gl_WorkGroupSize * 2u));
+  ivec2 iter = divide_ceil(imageSize(out_mip_5), ivec2(gl_WorkGroupSize.xy * 2u));
   ivec2 image_border = imageSize(out_mip_5) - 1;
   for (int y = 0; y < iter.y; y++) {
     for (int x = 0; x < iter.x; x++) {
       /* Load result of the other work groups. */
-      kernel_origin = ivec2(gl_WorkGroupSize) * ivec2(x, y);
+      kernel_origin = ivec2(gl_WorkGroupSize.xy) * ivec2(x, y);
       src_px = ivec2(kernel_origin + local_px) * 2;
       vec4 samp;
       samp.x = imageLoad(out_mip_5, min(src_px + ivec2(0, 1), image_border)).x;

source/blender/draw/engines/eevee_next/shaders/eevee_light_culling_tile_comp.glsl

@@ -168,13 +168,15 @@ void main()
         }
         /* Fallthrough to the hemispheric case. */
       case LIGHT_RECT:
-      case LIGHT_ELLIPSE:
+      case LIGHT_ELLIPSE: {
         vec3 v000 = vP - v_right * radius - v_up * radius;
         vec3 v100 = v000 + v_right * (radius * 2.0);
         vec3 v010 = v000 + v_up * (radius * 2.0);
         vec3 v001 = v000 - v_back * radius;
         Box bbox = shape_box(v000, v100, v010, v001);
         intersect_tile = intersect_tile && intersect(tile, bbox);
+        break;
+      }
       default:
         break;
     }

source/blender/draw/engines/eevee_next/shaders/eevee_motion_blur_dilate_comp.glsl

@@ -74,8 +74,10 @@ void main()
 
   vec4 max_motion = imageLoad(in_tiles_img, src_tile);
 
-  MotionPayload payload_prv = motion_blur_tile_indirection_pack_payload(max_motion.xy, src_tile);
+  MotionPayload payload_prv = motion_blur_tile_indirection_pack_payload(max_motion.xy,
-  MotionPayload payload_nxt = motion_blur_tile_indirection_pack_payload(max_motion.zw, src_tile);
+                                                                        uvec2(src_tile));
+  MotionPayload payload_nxt = motion_blur_tile_indirection_pack_payload(max_motion.zw,
+                                                                        uvec2(src_tile));
   if (true) {
     /* Rectangular area (in tiles) where the motion vector spreads. */
     MotionRect motion_rect = compute_motion_rect(src_tile, max_motion.xy);
@@ -85,17 +87,20 @@ void main()
       for (int y = 0; y < motion_rect.extent.y; y++) {
         ivec2 tile = motion_rect.bottom_left + ivec2(x, y);
         if (is_inside_motion_line(tile, motion_line)) {
-          motion_blur_tile_indirection_store(tile_indirection_buf, MOTION_PREV, tile, payload_prv);
+          motion_blur_tile_indirection_store(
+              tile_indirection_buf, MOTION_PREV, uvec2(tile), payload_prv);
           /* FIXME: This is a bit weird, but for some reason, we need the store the same vector in
            * the motion next so that weighting in gather pass is better. */
-          motion_blur_tile_indirection_store(tile_indirection_buf, MOTION_NEXT, tile, payload_nxt);
+          motion_blur_tile_indirection_store(
+              tile_indirection_buf, MOTION_NEXT, uvec2(tile), payload_nxt);
         }
       }
     }
   }
 
   if (true) {
-    MotionPayload payload = motion_blur_tile_indirection_pack_payload(max_motion.zw, src_tile);
+    MotionPayload payload = motion_blur_tile_indirection_pack_payload(max_motion.zw,
+                                                                      uvec2(src_tile));
     /* Rectangular area (in tiles) where the motion vector spreads. */
     MotionRect motion_rect = compute_motion_rect(src_tile, max_motion.zw);
     MotionLine motion_line = compute_motion_line(src_tile, max_motion.zw);
@@ -104,10 +109,12 @@ void main()
       for (int y = 0; y < motion_rect.extent.y; y++) {
         ivec2 tile = motion_rect.bottom_left + ivec2(x, y);
         if (is_inside_motion_line(tile, motion_line)) {
-          motion_blur_tile_indirection_store(tile_indirection_buf, MOTION_NEXT, tile, payload_nxt);
+          motion_blur_tile_indirection_store(
+              tile_indirection_buf, MOTION_NEXT, uvec2(tile), payload_nxt);
           /* FIXME: This is a bit weird, but for some reason, we need the store the same vector in
            * the motion next so that weighting in gather pass is better. */
-          motion_blur_tile_indirection_store(tile_indirection_buf, MOTION_PREV, tile, payload_prv);
+          motion_blur_tile_indirection_store(
+              tile_indirection_buf, MOTION_PREV, uvec2(tile), payload_prv);
         }
       }
     }

source/blender/draw/engines/eevee_next/shaders/eevee_motion_blur_gather_comp.glsl

@@ -178,10 +178,10 @@ void main()
   vec4 max_motion;
   /* Load dilation result from the indirection table. */
   ivec2 tile_prev;
-  motion_blur_tile_indirection_load(tile_indirection_buf, MOTION_PREV, tile, tile_prev);
+  motion_blur_tile_indirection_load(tile_indirection_buf, MOTION_PREV, uvec2(tile), tile_prev);
   max_motion.xy = imageLoad(in_tiles_img, tile_prev).xy;
   ivec2 tile_next;
-  motion_blur_tile_indirection_load(tile_indirection_buf, MOTION_NEXT, tile, tile_next);
+  motion_blur_tile_indirection_load(tile_indirection_buf, MOTION_NEXT, uvec2(tile), tile_next);
   max_motion.zw = imageLoad(in_tiles_img, tile_next).zw;
 
   Accumulator accum;

source/blender/draw/engines/eevee_next/shaders/eevee_nodetree_lib.glsl

@@ -242,13 +242,13 @@ void output_aov(vec4 color, float value, uint hash)
 #if defined(MAT_AOV_SUPPORT) && defined(GPU_FRAGMENT_SHADER)
   for (int i = 0; i < AOV_MAX && i < aov_buf.color_len; i++) {
     if (aov_buf.hash_color[i] == hash) {
-      imageStore(aov_color_img, ivec3(gl_FragCoord.xy, i), color);
+      imageStore(aov_color_img, ivec3(ivec2(gl_FragCoord.xy), i), color);
       return;
     }
   }
   for (int i = 0; i < AOV_MAX && i < aov_buf.value_len; i++) {
     if (aov_buf.hash_value[i] == hash) {
-      imageStore(aov_value_img, ivec3(gl_FragCoord.xy, i), vec4(value));
+      imageStore(aov_value_img, ivec3(ivec2(gl_FragCoord.xy), i), vec4(value));
       return;
     }
   }

source/blender/draw/engines/eevee_next/shaders/eevee_shadow_tag_update_comp.glsl

@@ -65,7 +65,7 @@ void main()
   }
 
   AABB aabb_tag;
-  AABB aabb_map = AABB(vec3(-0.99999), vec3(0.99999));
+  AABB aabb_map = shape_aabb(vec3(-0.99999), vec3(0.99999));
 
   /* Directionnal winmat have no correct near/far in the Z dimension at this point.
    * Do not clip in this dimension. */
@@ -87,7 +87,7 @@ void main()
     for (int y = box_min.y; y <= box_max.y; y++) {
       for (int x = box_min.x; x <= box_max.x; x++) {
         int tile_index = shadow_tile_offset(ivec2(x, y), tilemap.tiles_index, lod);
-        atomicOr(tiles_buf[tile_index], SHADOW_DO_UPDATE);
+        atomicOr(tiles_buf[tile_index], uint(SHADOW_DO_UPDATE));
       }
     }
   }

source/blender/draw/engines/eevee_next/shaders/eevee_shadow_tag_usage_lib.glsl

@@ -21,7 +21,7 @@ void shadow_tag_usage_tile(LightData light, ivec2 tile_co, int lod, int tilemap_
 
   tile_co >>= lod;
   int tile_index = shadow_tile_offset(tile_co, tilemaps_buf[tilemap_index].tiles_index, lod);
-  atomicOr(tiles_buf[tile_index], SHADOW_IS_USED);
+  atomicOr(tiles_buf[tile_index], uint(SHADOW_IS_USED));
 }
 
 void shadow_tag_usage_tilemap_directional(uint l_idx, vec3 P, vec3 V, float radius)

source/blender/draw/engines/workbench/shaders/workbench_shadow_visibility_comp.glsl

@@ -44,7 +44,7 @@ bool is_visible(IsectBox box)
 
 bool intersects_near_plane(IsectBox box)
 {
-  vec4 near_plane = drw_view_culling.planes[4];
+  vec4 near_plane = drw_view_culling.frustum_planes.planes[4];
   bool on_positive_side = false;
   bool on_negative_side = false;