Cleanup: Convert camelCase naming to snake_case in Compositor

To follow the style guide.

.clang-format

@@ -180,6 +180,7 @@ ForEachMacros:
   - CTX_DATA_BEGIN_WITH_ID
   - DEG_OBJECT_ITER_BEGIN
   - DEG_OBJECT_ITER_FOR_RENDER_ENGINE_BEGIN
+  - DRW_ENABLED_ENGINE_ITER
   - DRIVER_TARGETS_LOOPER_BEGIN
   - DRIVER_TARGETS_USED_LOOPER_BEGIN
   - FOREACH_BASE_IN_EDIT_MODE_BEGIN

CMakeLists.txt

@@ -406,6 +406,7 @@ mark_as_advanced(WITH_CYCLES_CUDA_BUILD_SERIAL)
 set(CYCLES_TEST_DEVICES CPU CACHE STRING "Run regression tests on the specified device types (CPU CUDA OPTIX)" )
 set(CYCLES_CUDA_BINARIES_ARCH sm_30 sm_35 sm_37 sm_50 sm_52 sm_60 sm_61 sm_70 sm_75 sm_86 compute_75 CACHE STRING "CUDA architectures to build binaries for")
 mark_as_advanced(CYCLES_CUDA_BINARIES_ARCH)
+option(WITH_CYCLES_HIP_BINARIES     "Build Cycles HIP binaries" OFF)
 unset(PLATFORM_DEFAULT)
 option(WITH_CYCLES_LOGGING      "Build Cycles with logging support" ON)
 option(WITH_CYCLES_DEBUG_NAN    "Build Cycles with additional asserts for detecting NaNs and invalid values" OFF)

build_files/cmake/Modules/FindHIP.cmake

@@ -0,0 +1,81 @@
+# - Find HIP compiler
+#
+# This module defines
+#  HIP_HIPCC_EXECUTABLE, the full path to the hipcc executable
+#  HIP_VERSION, the HIP compiler version
+#
+#  HIP_FOUND, if the HIP toolkit is found.
+
+#=============================================================================
+# Copyright 2021 Blender Foundation.
+#
+# Distributed under the OSI-approved BSD 3-Clause License,
+# see accompanying file BSD-3-Clause-license.txt for details.
+#=============================================================================
+
+# If HIP_ROOT_DIR was defined in the environment, use it.
+if(NOT HIP_ROOT_DIR AND NOT $ENV{HIP_ROOT_DIR} STREQUAL "")
+  set(HIP_ROOT_DIR $ENV{HIP_ROOT_DIR})
+endif()
+
+set(_hip_SEARCH_DIRS
+  ${HIP_ROOT_DIR}
+)
+
+find_program(HIP_HIPCC_EXECUTABLE
+  NAMES
+    hipcc
+  HINTS
+    ${_hip_SEARCH_DIRS}
+  PATH_SUFFIXES
+    bin
+)
+
+if(HIP_HIPCC_EXECUTABLE AND NOT EXISTS ${HIP_HIPCC_EXECUTABLE})
+  message(WARNING "Cached or directly specified hipcc executable does not exist.")
+  set(HIP_FOUND FALSE)
+elseif(HIP_HIPCC_EXECUTABLE)
+  set(HIP_FOUND TRUE)
+
+  set(HIP_VERSION_MAJOR 0)
+  set(HIP_VERSION_MINOR 0)
+  set(HIP_VERSION_PATCH 0)
+
+  # Get version from the output.
+  execute_process(COMMAND ${HIP_HIPCC_EXECUTABLE} --version
+                  OUTPUT_VARIABLE HIP_VERSION_RAW
+                  ERROR_QUIET
+                  OUTPUT_STRIP_TRAILING_WHITESPACE)
+
+  # Parse parts.
+  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}")
+    string(REPLACE "." ";" HIP_VERSION_PARTS "${HIP_SEMANTIC_VERSION}")
+    list(LENGTH HIP_VERSION_PARTS NUM_HIP_VERSION_PARTS)
+
+    # Extract components into corresponding variables.
+    if(NUM_HIP_VERSION_PARTS GREATER 0)
+      list(GET HIP_VERSION_PARTS 0 HIP_VERSION_MAJOR)
+    endif()
+    if(NUM_HIP_VERSION_PARTS GREATER 1)
+      list(GET HIP_VERSION_PARTS 1 HIP_VERSION_MINOR)
+    endif()
+    if(NUM_HIP_VERSION_PARTS GREATER 2)
+      list(GET HIP_VERSION_PARTS 2 HIP_VERSION_PATCH)
+    endif()
+
+    # Unset temp variables.
+    unset(NUM_HIP_VERSION_PARTS)
+    unset(HIP_SEMANTIC_VERSION)
+    unset(HIP_VERSION_PARTS)
+  endif()
+
+  # Construct full semantic version.
+  set(HIP_VERSION "${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_PATCH}")
+  unset(HIP_VERSION_RAW)
+else()
+  set(HIP_FOUND FALSE)
+endif()

extern/cuew/include/cuew.h

@@ -609,6 +609,7 @@ typedef enum cudaError_enum {
   CUDA_ERROR_INVALID_GRAPHICS_CONTEXT = 219,
   CUDA_ERROR_NVLINK_UNCORRECTABLE = 220,
   CUDA_ERROR_JIT_COMPILER_NOT_FOUND = 221,
+  CUDA_ERROR_UNSUPPORTED_PTX_VERSION = 222,
   CUDA_ERROR_INVALID_SOURCE = 300,
   CUDA_ERROR_FILE_NOT_FOUND = 301,
   CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND = 302,

extern/cuew/src/cuew.c

@@ -736,6 +736,7 @@ const char *cuewErrorString(CUresult result) {
     case CUDA_ERROR_INVALID_GRAPHICS_CONTEXT: return "Invalid graphics context";
     case CUDA_ERROR_NVLINK_UNCORRECTABLE: return "Nvlink uncorrectable";
     case CUDA_ERROR_JIT_COMPILER_NOT_FOUND: return "Jit compiler not found";
+    case CUDA_ERROR_UNSUPPORTED_PTX_VERSION: return "Unsupported PTX version";
     case CUDA_ERROR_INVALID_SOURCE: return "Invalid source";
     case CUDA_ERROR_FILE_NOT_FOUND: return "File not found";
     case CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND: return "Link to a shared object failed to resolve";

extern/smaa_areatex/smaa_areatex.cpp

@@ -574,6 +574,8 @@ Dbl2 AreaDiag::area(Dbl2 p1, Dbl2 p2, int left)
 	Dbl2 d = p2 - p1;
 	if (d.x == 0.0)
 		return Dbl2(0.0, 1.0);
+	if (d.y == 0.0)
+		return Dbl2(1.0, 0.0);
 
 	double x1 = (double)(1 + left);
 	double x2 = x1 + 1.0;

intern/cycles/blender/blender_camera.cpp

@@ -927,6 +927,9 @@ BufferParams BlenderSync::get_buffer_params(
     params.height = height;
   }
 
+  params.window_width = params.width;
+  params.window_height = params.height;
+
   return params;
 }
 

intern/cycles/bvh/bvh.h

@@ -50,10 +50,6 @@ struct PackedBVH {
   array<int4> leaf_nodes;
   /* object index to BVH node index mapping for instances */
   array<int> object_node;
-  /* Mapping from primitive index to index in triangle array. */
-  array<uint> prim_tri_index;
-  /* Continuous storage of triangle vertices. */
-  array<float4> prim_tri_verts;
   /* primitive type - triangle or strand */
   array<int> prim_type;
   /* visibility visibilitys for primitives */

intern/cycles/bvh/bvh2.cpp

@@ -439,61 +439,20 @@ void BVH2::refit_primitives(int start, int end, BoundBox &bbox, uint &visibility
 
 /* Triangles */
 
-void BVH2::pack_triangle(int idx, float4 tri_verts[3])
-{
-  int tob = pack.prim_object[idx];
-  assert(tob >= 0 && tob < objects.size());
-  const Mesh *mesh = static_cast<const Mesh *>(objects[tob]->get_geometry());
-
-  int tidx = pack.prim_index[idx];
-  Mesh::Triangle t = mesh->get_triangle(tidx);
-  const float3 *vpos = &mesh->verts[0];
-  float3 v0 = vpos[t.v[0]];
-  float3 v1 = vpos[t.v[1]];
-  float3 v2 = vpos[t.v[2]];
-
-  tri_verts[0] = float3_to_float4(v0);
-  tri_verts[1] = float3_to_float4(v1);
-  tri_verts[2] = float3_to_float4(v2);
-}
-
 void BVH2::pack_primitives()
 {
   const size_t tidx_size = pack.prim_index.size();
-  size_t num_prim_triangles = 0;
-  /* Count number of triangles primitives in BVH. */
-  for (unsigned int i = 0; i < tidx_size; i++) {
-    if ((pack.prim_index[i] != -1)) {
-      if ((pack.prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
-        ++num_prim_triangles;
-      }
-    }
-  }
   /* Reserve size for arrays. */
-  pack.prim_tri_index.clear();
-  pack.prim_tri_index.resize(tidx_size);
-  pack.prim_tri_verts.clear();
-  pack.prim_tri_verts.resize(num_prim_triangles * 3);
   pack.prim_visibility.clear();
   pack.prim_visibility.resize(tidx_size);
   /* Fill in all the arrays. */
-  size_t prim_triangle_index = 0;
   for (unsigned int i = 0; i < tidx_size; i++) {
     if (pack.prim_index[i] != -1) {
       int tob = pack.prim_object[i];
       Object *ob = objects[tob];
-      if ((pack.prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
-        pack_triangle(i, (float4 *)&pack.prim_tri_verts[3 * prim_triangle_index]);
-        pack.prim_tri_index[i] = 3 * prim_triangle_index;
-        ++prim_triangle_index;
-      }
-      else {
-        pack.prim_tri_index[i] = -1;
-      }
       pack.prim_visibility[i] = ob->visibility_for_tracing();
     }
     else {
-      pack.prim_tri_index[i] = -1;
       pack.prim_visibility[i] = 0;
     }
   }
@@ -522,10 +481,8 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
 
   /* reserve */
   size_t prim_index_size = pack.prim_index.size();
-  size_t prim_tri_verts_size = pack.prim_tri_verts.size();
 
   size_t pack_prim_index_offset = prim_index_size;
-  size_t pack_prim_tri_verts_offset = prim_tri_verts_size;
   size_t pack_nodes_offset = nodes_size;
   size_t pack_leaf_nodes_offset = leaf_nodes_size;
   size_t object_offset = 0;
@@ -535,7 +492,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
 
     if (geom->need_build_bvh(params.bvh_layout)) {
       prim_index_size += bvh->pack.prim_index.size();
-      prim_tri_verts_size += bvh->pack.prim_tri_verts.size();
       nodes_size += bvh->pack.nodes.size();
       leaf_nodes_size += bvh->pack.leaf_nodes.size();
     }
@@ -545,8 +501,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
   pack.prim_type.resize(prim_index_size);
   pack.prim_object.resize(prim_index_size);
   pack.prim_visibility.resize(prim_index_size);
-  pack.prim_tri_verts.resize(prim_tri_verts_size);
-  pack.prim_tri_index.resize(prim_index_size);
   pack.nodes.resize(nodes_size);
   pack.leaf_nodes.resize(leaf_nodes_size);
   pack.object_node.resize(objects.size());
@@ -559,8 +513,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
   int *pack_prim_type = (pack.prim_type.size()) ? &pack.prim_type[0] : NULL;
   int *pack_prim_object = (pack.prim_object.size()) ? &pack.prim_object[0] : NULL;
   uint *pack_prim_visibility = (pack.prim_visibility.size()) ? &pack.prim_visibility[0] : NULL;
-  float4 *pack_prim_tri_verts = (pack.prim_tri_verts.size()) ? &pack.prim_tri_verts[0] : NULL;
-  uint *pack_prim_tri_index = (pack.prim_tri_index.size()) ? &pack.prim_tri_index[0] : NULL;
   int4 *pack_nodes = (pack.nodes.size()) ? &pack.nodes[0] : NULL;
   int4 *pack_leaf_nodes = (pack.leaf_nodes.size()) ? &pack.leaf_nodes[0] : NULL;
   float2 *pack_prim_time = (pack.prim_time.size()) ? &pack.prim_time[0] : NULL;
@@ -609,18 +561,14 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
       int *bvh_prim_index = &bvh->pack.prim_index[0];
       int *bvh_prim_type = &bvh->pack.prim_type[0];
       uint *bvh_prim_visibility = &bvh->pack.prim_visibility[0];
-      uint *bvh_prim_tri_index = &bvh->pack.prim_tri_index[0];
       float2 *bvh_prim_time = bvh->pack.prim_time.size() ? &bvh->pack.prim_time[0] : NULL;
 
       for (size_t i = 0; i < bvh_prim_index_size; i++) {
         if (bvh->pack.prim_type[i] & PRIMITIVE_ALL_CURVE) {
           pack_prim_index[pack_prim_index_offset] = bvh_prim_index[i] + geom_prim_offset;
-          pack_prim_tri_index[pack_prim_index_offset] = -1;
         }
         else {
           pack_prim_index[pack_prim_index_offset] = bvh_prim_index[i] + geom_prim_offset;
-          pack_prim_tri_index[pack_prim_index_offset] = bvh_prim_tri_index[i] +
-                                                        pack_prim_tri_verts_offset;
         }
 
         pack_prim_type[pack_prim_index_offset] = bvh_prim_type[i];
@@ -633,15 +581,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
       }
     }
 
-    /* Merge triangle vertices data. */
-    if (bvh->pack.prim_tri_verts.size()) {
-      const size_t prim_tri_size = bvh->pack.prim_tri_verts.size();
-      memcpy(pack_prim_tri_verts + pack_prim_tri_verts_offset,
-             &bvh->pack.prim_tri_verts[0],
-             prim_tri_size * sizeof(float4));
-      pack_prim_tri_verts_offset += prim_tri_size;
-    }
-
     /* merge nodes */
     if (bvh->pack.leaf_nodes.size()) {
       int4 *leaf_nodes_offset = &bvh->pack.leaf_nodes[0];

intern/cycles/bvh/bvh_build.cpp

@@ -67,8 +67,12 @@ BVHBuild::~BVHBuild()
 
 /* Adding References */
 
-void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *mesh, int i)
+void BVHBuild::add_reference_triangles(BoundBox &root,
+                                       BoundBox &center,
+                                       Mesh *mesh,
+                                       int object_index)
 {
+  const PrimitiveType primitive_type = mesh->primitive_type();
   const Attribute *attr_mP = NULL;
   if (mesh->has_motion_blur()) {
     attr_mP = mesh->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
@@ -81,7 +85,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
       BoundBox bounds = BoundBox::empty;
       t.bounds_grow(verts, bounds);
       if (bounds.valid() && t.valid(verts)) {
-        references.push_back(BVHReference(bounds, j, i, PRIMITIVE_TRIANGLE));
+        references.push_back(BVHReference(bounds, j, object_index, primitive_type));
         root.grow(bounds);
         center.grow(bounds.center2());
       }
@@ -101,7 +105,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
         t.bounds_grow(vert_steps + step * num_verts, bounds);
       }
       if (bounds.valid()) {
-        references.push_back(BVHReference(bounds, j, i, PRIMITIVE_MOTION_TRIANGLE));
+        references.push_back(BVHReference(bounds, j, object_index, primitive_type));
         root.grow(bounds);
         center.grow(bounds.center2());
       }
@@ -140,7 +144,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
         if (bounds.valid()) {
           const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1;
           references.push_back(
-              BVHReference(bounds, j, i, PRIMITIVE_MOTION_TRIANGLE, prev_time, curr_time));
+              BVHReference(bounds, j, object_index, primitive_type, prev_time, curr_time));
           root.grow(bounds);
           center.grow(bounds.center2());
         }
@@ -153,18 +157,14 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
   }
 }
 
-void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair, int i)
+void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair, int object_index)
 {
   const Attribute *curve_attr_mP = NULL;
   if (hair->has_motion_blur()) {
     curve_attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
   }
 
-  const PrimitiveType primitive_type =
-      (curve_attr_mP != NULL) ?
-          ((hair->curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON :
-                                                 PRIMITIVE_MOTION_CURVE_THICK) :
-          ((hair->curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON : PRIMITIVE_CURVE_THICK);
+  const PrimitiveType primitive_type = hair->primitive_type();
 
   const size_t num_curves = hair->num_curves();
   for (uint j = 0; j < num_curves; j++) {
@@ -177,7 +177,7 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
         curve.bounds_grow(k, &hair->get_curve_keys()[0], curve_radius, bounds);
         if (bounds.valid()) {
           int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
-          references.push_back(BVHReference(bounds, j, i, packed_type));
+          references.push_back(BVHReference(bounds, j, object_index, packed_type));
           root.grow(bounds);
           center.grow(bounds.center2());
         }
@@ -198,7 +198,7 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
         }
         if (bounds.valid()) {
           int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
-          references.push_back(BVHReference(bounds, j, i, packed_type));
+          references.push_back(BVHReference(bounds, j, object_index, packed_type));
           root.grow(bounds);
           center.grow(bounds.center2());
         }
@@ -254,7 +254,8 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
           if (bounds.valid()) {
             const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1;
             int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
-            references.push_back(BVHReference(bounds, j, i, packed_type, prev_time, curr_time));
+            references.push_back(
+                BVHReference(bounds, j, object_index, packed_type, prev_time, curr_time));
             root.grow(bounds);
             center.grow(bounds.center2());
           }
@@ -268,15 +269,18 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
   }
 }
 
-void BVHBuild::add_reference_geometry(BoundBox &root, BoundBox &center, Geometry *geom, int i)
+void BVHBuild::add_reference_geometry(BoundBox &root,
+                                      BoundBox &center,
+                                      Geometry *geom,
+                                      int object_index)
 {
   if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
     Mesh *mesh = static_cast<Mesh *>(geom);
-    add_reference_triangles(root, center, mesh, i);
+    add_reference_triangles(root, center, mesh, object_index);
   }
   else if (geom->geometry_type == Geometry::HAIR) {
     Hair *hair = static_cast<Hair *>(geom);
-    add_reference_curves(root, center, hair, i);
+    add_reference_curves(root, center, hair, object_index);
   }
 }
 

intern/cycles/bvh/bvh_embree.cpp

@@ -89,20 +89,9 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
 
       /* Test if we need to record this transparent intersection. */
       if (ctx->num_hits < ctx->max_hits || ray->tfar < ctx->max_t) {
-        /* Skip already recorded intersections. */
-        int num_recorded_hits = min(ctx->num_hits, ctx->max_hits);
-
-        for (int i = 0; i < num_recorded_hits; ++i) {
-          if (current_isect.object == ctx->isect_s[i].object &&
-              current_isect.prim == ctx->isect_s[i].prim && current_isect.t == ctx->isect_s[i].t) {
-            /* This intersection was already recorded, skip it. */
-            *args->valid = 0;
-            return;
-          }
-        }
-
         /* If maximum number of hits was reached, replace the intersection with the
          * highest distance. We want to find the N closest intersections. */
+        const int num_recorded_hits = min(ctx->num_hits, ctx->max_hits);
         int isect_index = num_recorded_hits;
         if (num_recorded_hits + 1 >= ctx->max_hits) {
           float max_t = ctx->isect_s[0].t;
@@ -147,10 +136,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
       }
       else {
         kernel_embree_convert_hit(kg, ray, hit, &current_isect);
-        int object = (current_isect.object == OBJECT_NONE) ?
-                         kernel_tex_fetch(__prim_object, current_isect.prim) :
-                         current_isect.object;
-        if (ctx->local_object_id != object) {
+        if (ctx->local_object_id != current_isect.object) {
           /* This tells Embree to continue tracing. */
           *args->valid = 0;
           break;
@@ -169,41 +155,49 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
         break;
       }
 
-      /* See triangle_intersect_subsurface() for the native equivalent. */
-      for (int i = min(ctx->max_hits, ctx->local_isect->num_hits) - 1; i >= 0; --i) {
-        if (ctx->local_isect->hits[i].t == ray->tfar) {
-          /* This tells Embree to continue tracing. */
-          *args->valid = 0;
-          break;
-        }
-      }
-
+      LocalIntersection *local_isect = ctx->local_isect;
       int hit_idx = 0;
 
       if (ctx->lcg_state) {
+        /* See triangle_intersect_subsurface() for the native equivalent. */
+        for (int i = min(ctx->max_hits, local_isect->num_hits) - 1; i >= 0; --i) {
+          if (local_isect->hits[i].t == ray->tfar) {
+            /* This tells Embree to continue tracing. */
+            *args->valid = 0;
+            return;
+          }
+        }
 
-        ++ctx->local_isect->num_hits;
-        if (ctx->local_isect->num_hits <= ctx->max_hits) {
-          hit_idx = ctx->local_isect->num_hits - 1;
+        local_isect->num_hits++;
+
+        if (local_isect->num_hits <= ctx->max_hits) {
+          hit_idx = local_isect->num_hits - 1;
         }
         else {
           /* reservoir sampling: if we are at the maximum number of
            * hits, randomly replace element or skip it */
-          hit_idx = lcg_step_uint(ctx->lcg_state) % ctx->local_isect->num_hits;
+          hit_idx = lcg_step_uint(ctx->lcg_state) % local_isect->num_hits;
 
           if (hit_idx >= ctx->max_hits) {
             /* This tells Embree to continue tracing. */
             *args->valid = 0;
-            break;
+            return;
           }
         }
       }
       else {
-        ctx->local_isect->num_hits = 1;
+        /* Record closest intersection only. */
+        if (local_isect->num_hits && current_isect.t > local_isect->hits[0].t) {
+          *args->valid = 0;
+          return;
+        }
+
+        local_isect->num_hits = 1;
       }
+
       /* record intersection */
-      ctx->local_isect->hits[hit_idx] = current_isect;
-      ctx->local_isect->Ng[hit_idx] = normalize(make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z));
+      local_isect->hits[hit_idx] = current_isect;
+      local_isect->Ng[hit_idx] = normalize(make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z));
       /* This tells Embree to continue tracing. */
       *args->valid = 0;
       break;
@@ -213,21 +207,11 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
       if (ctx->num_hits < ctx->max_hits) {
         Intersection current_isect;
         kernel_embree_convert_hit(kg, ray, hit, &current_isect);
-        for (size_t i = 0; i < ctx->num_hits; ++i) {
-          if (current_isect.object == ctx->isect_s[i].object &&
-              current_isect.prim == ctx->isect_s[i].prim && current_isect.t == ctx->isect_s[i].t) {
-            /* This intersection was already recorded, skip it. */
-            *args->valid = 0;
-            break;
-          }
-        }
         Intersection *isect = &ctx->isect_s[ctx->num_hits];
         ++ctx->num_hits;
         *isect = current_isect;
         /* Only primitives from volume object. */
-        uint tri_object = (isect->object == OBJECT_NONE) ?
-                              kernel_tex_fetch(__prim_object, isect->prim) :
-                              isect->object;
+        uint tri_object = isect->object;
         int object_flag = kernel_tex_fetch(__object_flag, tri_object);
         if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
           --ctx->num_hits;
@@ -249,7 +233,7 @@ static void rtc_filter_func_thick_curve(const RTCFilterFunctionNArguments *args)
   const RTCRay *ray = (RTCRay *)args->ray;
   RTCHit *hit = (RTCHit *)args->hit;
 
-  /* Always ignore backfacing intersections. */
+  /* Always ignore back-facing intersections. */
   if (dot(make_float3(ray->dir_x, ray->dir_y, ray->dir_z),
           make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)) > 0.0f) {
     *args->valid = 0;
@@ -262,7 +246,7 @@ static void rtc_filter_occluded_func_thick_curve(const RTCFilterFunctionNArgumen
   const RTCRay *ray = (RTCRay *)args->ray;
   RTCHit *hit = (RTCHit *)args->hit;
 
-  /* Always ignore backfacing intersections. */
+  /* Always ignore back-facing intersections. */
   if (dot(make_float3(ray->dir_x, ray->dir_y, ray->dir_z),
           make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)) > 0.0f) {
     *args->valid = 0;
@@ -456,7 +440,7 @@ void BVHEmbree::add_instance(Object *ob, int i)
 
 void BVHEmbree::add_triangles(const Object *ob, const Mesh *mesh, int i)
 {
-  size_t prim_offset = mesh->optix_prim_offset;
+  size_t prim_offset = mesh->prim_offset;
 
   const Attribute *attr_mP = NULL;
   size_t num_motion_steps = 1;
@@ -625,7 +609,7 @@ void BVHEmbree::set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, c
 
 void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
 {
-  size_t prim_offset = hair->optix_prim_offset;
+  size_t prim_offset = hair->curve_segment_offset;
 
   const Attribute *attr_mP = NULL;
   size_t num_motion_steps = 1;
@@ -702,7 +686,7 @@ void BVHEmbree::refit(Progress &progress)
         if (mesh->num_triangles() > 0) {
           RTCGeometry geom = rtcGetGeometry(scene, geom_id);
           set_tri_vertex_buffer(geom, mesh, true);
-          rtcSetGeometryUserData(geom, (void *)mesh->optix_prim_offset);
+          rtcSetGeometryUserData(geom, (void *)mesh->prim_offset);
           rtcCommitGeometry(geom);
         }
       }
@@ -711,7 +695,7 @@ void BVHEmbree::refit(Progress &progress)
         if (hair->num_curves() > 0) {
           RTCGeometry geom = rtcGetGeometry(scene, geom_id + 1);
           set_curve_vertex_buffer(geom, hair, true);
-          rtcSetGeometryUserData(geom, (void *)hair->optix_prim_offset);
+          rtcSetGeometryUserData(geom, (void *)hair->curve_segment_offset);
           rtcCommitGeometry(geom);
         }
       }

intern/cycles/cmake/external_libs.cmake

@@ -521,7 +521,7 @@ endif()
 if(WITH_CYCLES_CUDA_BINARIES OR NOT WITH_CUDA_DYNLOAD)
   find_package(CUDA) # Try to auto locate CUDA toolkit
   if(CUDA_FOUND)
-    message(STATUS "CUDA nvcc = ${CUDA_NVCC_EXECUTABLE}")
+    message(STATUS "Found CUDA ${CUDA_NVCC_EXECUTABLE} (${CUDA_VERSION})")
   else()
     message(STATUS "CUDA compiler not found, disabling WITH_CYCLES_CUDA_BINARIES")
     set(WITH_CYCLES_CUDA_BINARIES OFF)
@@ -537,6 +537,16 @@ endif()
 # HIP
 ###########################################################################
 
+if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
+  find_package(HIP)
+  if(HIP_FOUND)
+    message(STATUS "Found HIP ${HIP_HIPCC_EXECUTABLE} (${HIP_VERSION})")
+  else()
+    message(STATUS "HIP compiler not found, disabling WITH_CYCLES_HIP_BINARIES")
+    set(WITH_CYCLES_HIP_BINARIES OFF)
+  endif()
+endif()
+
 if(NOT WITH_HIP_DYNLOAD)
   set(WITH_HIP_DYNLOAD ON)
 endif()

intern/cycles/device/optix/device_impl.cpp

@@ -768,7 +768,13 @@ void OptiXDevice::denoise_color_read(DenoiseContext &context, const DenoisePass
   destination.num_components = 3;
   destination.pixel_stride = context.buffer_params.pass_stride;
 
-  pass_accessor.get_render_tile_pixels(context.render_buffers, context.buffer_params, destination);
+  BufferParams buffer_params = context.buffer_params;
+  buffer_params.window_x = 0;
+  buffer_params.window_y = 0;
+  buffer_params.window_width = buffer_params.width;
+  buffer_params.window_height = buffer_params.height;
+
+  pass_accessor.get_render_tile_pixels(context.render_buffers, buffer_params, destination);
 }
 
 bool OptiXDevice::denoise_filter_color_preprocess(DenoiseContext &context, const DenoisePass &pass)
@@ -1246,7 +1252,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         build_input.curveArray.indexBuffer = (CUdeviceptr)index_data.device_pointer;
         build_input.curveArray.indexStrideInBytes = sizeof(int);
         build_input.curveArray.flag = build_flags;
-        build_input.curveArray.primitiveIndexOffset = hair->optix_prim_offset;
+        build_input.curveArray.primitiveIndexOffset = hair->curve_segment_offset;
       }
       else {
         /* Disable visibility test any-hit program, since it is already checked during
@@ -1259,7 +1265,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         build_input.customPrimitiveArray.strideInBytes = sizeof(OptixAabb);
         build_input.customPrimitiveArray.flags = &build_flags;
         build_input.customPrimitiveArray.numSbtRecords = 1;
-        build_input.customPrimitiveArray.primitiveIndexOffset = hair->optix_prim_offset;
+        build_input.customPrimitiveArray.primitiveIndexOffset = hair->curve_segment_offset;
       }
 
       if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
@@ -1328,7 +1334,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
        * buffers for that purpose. OptiX does not allow this to be zero though, so just pass in
        * one and rely on that having the same meaning in this case. */
       build_input.triangleArray.numSbtRecords = 1;
-      build_input.triangleArray.primitiveIndexOffset = mesh->optix_prim_offset;
+      build_input.triangleArray.primitiveIndexOffset = mesh->prim_offset;
 
       if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
         progress.set_error("Failed to build OptiX acceleration structure");
@@ -1395,8 +1401,8 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
       instance.transform[5] = 1.0f;
       instance.transform[10] = 1.0f;
 
-      /* Set user instance ID to object index (but leave low bit blank). */
-      instance.instanceId = ob->get_device_index() << 1;
+      /* Set user instance ID to object index. */
+      instance.instanceId = ob->get_device_index();
 
       /* Add some of the object visibility bits to the mask.
        * __prim_visibility contains the combined visibility bits of all instances, so is not
@@ -1508,9 +1514,6 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         else {
           /* Disable instance transform if geometry already has it applied to vertex data. */
           instance.flags |= OPTIX_INSTANCE_FLAG_DISABLE_TRANSFORM;
-          /* Non-instanced objects read ID from 'prim_object', so distinguish
-           * them from instanced objects with the low bit set. */
-          instance.instanceId |= 1;
         }
       }
     }

intern/cycles/integrator/denoiser_oidn.cpp

@@ -289,7 +289,13 @@ class OIDNDenoiseContext {
      * pixels. */
     const PassAccessorCPU pass_accessor(pass_access_info, 1.0f, num_samples_);
 
-    pass_accessor.get_render_tile_pixels(render_buffers_, buffer_params_, destination);
+    BufferParams buffer_params = buffer_params_;
+    buffer_params.window_x = 0;
+    buffer_params.window_y = 0;
+    buffer_params.window_width = buffer_params.width;
+    buffer_params.window_height = buffer_params.height;
+
+    pass_accessor.get_render_tile_pixels(render_buffers_, buffer_params, destination);
   }
 
   /* Read pass pixels using PassAccessor into a temporary buffer which is owned by the pass.. */

intern/cycles/integrator/pass_accessor_cpu.cpp

@@ -99,17 +99,22 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_float(
 {
   DCHECK_EQ(destination.stride, 0) << "Custom stride for float destination is not implemented.";
 
-  const float *buffer_data = render_buffers->buffer.data();
+  const int64_t pass_stride = buffer_params.pass_stride;
+  const int64_t buffer_row_stride = buffer_params.stride * buffer_params.pass_stride;
+
+  const float *window_data = render_buffers->buffer.data() + buffer_params.window_x * pass_stride +
+                             buffer_params.window_y * buffer_row_stride;
+
   const int pixel_stride = destination.pixel_stride ? destination.pixel_stride :
                                                       destination.num_components;
 
-  tbb::parallel_for(0, buffer_params.height, [&](int64_t y) {
-    int64_t pixel_index = y * buffer_params.width;
-    for (int64_t x = 0; x < buffer_params.width; ++x, ++pixel_index) {
-      const int64_t input_pixel_offset = pixel_index * buffer_params.pass_stride;
-      const float *buffer = buffer_data + input_pixel_offset;
-      float *pixel = destination.pixels + (pixel_index + destination.offset) * pixel_stride;
+  tbb::parallel_for(0, buffer_params.window_height, [&](int64_t y) {
+    const float *buffer = window_data + y * buffer_row_stride;
+    float *pixel = destination.pixels +
+                   (y * buffer_params.width + destination.offset) * pixel_stride;
 
+    for (int64_t x = 0; x < buffer_params.window_width;
+         ++x, buffer += pass_stride, pixel += pixel_stride) {
       processor(kfilm_convert, buffer, pixel);
     }
   });
@@ -123,26 +128,28 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_half_rgba(
     const Destination &destination,
     const Processor &processor) const
 {
-  const float *buffer_data = render_buffers->buffer.data();
+  const int64_t pass_stride = buffer_params.pass_stride;
+  const int64_t buffer_row_stride = buffer_params.stride * buffer_params.pass_stride;
+
+  const float *window_data = render_buffers->buffer.data() + buffer_params.window_x * pass_stride +
+                             buffer_params.window_y * buffer_row_stride;
 
   half4 *dst_start = destination.pixels_half_rgba + destination.offset;
   const int destination_stride = destination.stride != 0 ? destination.stride :
                                                            buffer_params.width;
 
-  tbb::parallel_for(0, buffer_params.height, [&](int64_t y) {
-    int64_t pixel_index = y * buffer_params.width;
-    half4 *dst_row_start = dst_start + y * destination_stride;
-    for (int64_t x = 0; x < buffer_params.width; ++x, ++pixel_index) {
-      const int64_t input_pixel_offset = pixel_index * buffer_params.pass_stride;
-      const float *buffer = buffer_data + input_pixel_offset;
+  tbb::parallel_for(0, buffer_params.window_height, [&](int64_t y) {
+    const float *buffer = window_data + y * buffer_row_stride;
+    half4 *pixel = dst_start + y * destination_stride;
+    for (int64_t x = 0; x < buffer_params.window_width; ++x, buffer += pass_stride, ++pixel) {
 
-      float pixel[4];
-      processor(kfilm_convert, buffer, pixel);
+      float pixel_rgba[4];
+      processor(kfilm_convert, buffer, pixel_rgba);
 
-      film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel);
+      film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel_rgba);
 
-      half4 *pixel_half_rgba = dst_row_start + x;
-      float4_store_half(&pixel_half_rgba->x, make_float4(pixel[0], pixel[1], pixel[2], pixel[3]));
+      float4_store_half(&pixel->x,
+                        make_float4(pixel_rgba[0], pixel_rgba[1], pixel_rgba[2], pixel_rgba[3]));
     }
   });
 }

intern/cycles/integrator/pass_accessor_gpu.cpp

@@ -43,10 +43,13 @@ void PassAccessorGPU::run_film_convert_kernels(DeviceKernel kernel,
   KernelFilmConvert kfilm_convert;
   init_kernel_film_convert(&kfilm_convert, buffer_params, destination);
 
-  const int work_size = buffer_params.width * buffer_params.height;
+  const int work_size = buffer_params.window_width * buffer_params.window_height;
 
   const int destination_stride = destination.stride != 0 ? destination.stride :
-                                                           buffer_params.width;
+                                                           buffer_params.window_width;
+
+  const int offset = buffer_params.window_x * buffer_params.pass_stride +
+                     buffer_params.window_y * buffer_params.stride * buffer_params.pass_stride;
 
   if (destination.d_pixels) {
     DCHECK_EQ(destination.stride, 0) << "Custom stride for float destination is not implemented.";
@@ -55,8 +58,8 @@ void PassAccessorGPU::run_film_convert_kernels(DeviceKernel kernel,
                     const_cast<device_ptr *>(&destination.d_pixels),
                     const_cast<device_ptr *>(&render_buffers->buffer.device_pointer),
                     const_cast<int *>(&work_size),
-                    const_cast<int *>(&buffer_params.width),
-                    const_cast<int *>(&buffer_params.offset),
+                    const_cast<int *>(&buffer_params.window_width),
+                    const_cast<int *>(&offset),
                     const_cast<int *>(&buffer_params.stride),
                     const_cast<int *>(&destination.offset),
                     const_cast<int *>(&destination_stride)};
@@ -70,8 +73,8 @@ void PassAccessorGPU::run_film_convert_kernels(DeviceKernel kernel,
                     const_cast<device_ptr *>(&destination.d_pixels_half_rgba),
                     const_cast<device_ptr *>(&render_buffers->buffer.device_pointer),
                     const_cast<int *>(&work_size),
-                    const_cast<int *>(&buffer_params.width),
-                    const_cast<int *>(&buffer_params.offset),
+                    const_cast<int *>(&buffer_params.window_width),
+                    const_cast<int *>(&offset),
                     const_cast<int *>(&buffer_params.stride),
                     const_cast<int *>(&destination.offset),
                     const_cast<int *>(&destination_stride)};

intern/cycles/integrator/path_trace.cpp

@@ -234,42 +234,53 @@ template<typename Callback>
 static void foreach_sliced_buffer_params(const vector<unique_ptr<PathTraceWork>> &path_trace_works,
                                          const vector<WorkBalanceInfo> &work_balance_infos,
                                          const BufferParams &buffer_params,
+                                         const int overscan,
                                          const Callback &callback)
 {
   const int num_works = path_trace_works.size();
-  const int height = buffer_params.height;
+  const int window_height = buffer_params.window_height;
 
   int current_y = 0;
   for (int i = 0; i < num_works; ++i) {
     const double weight = work_balance_infos[i].weight;
-    const int slice_height = max(lround(height * weight), 1);
+    const int slice_window_full_y = buffer_params.full_y + buffer_params.window_y + current_y;
+    const int slice_window_height = max(lround(window_height * weight), 1);
 
     /* Disallow negative values to deal with situations when there are more compute devices than
      * scan-lines. */
-    const int remaining_height = max(0, height - current_y);
+    const int remaining_window_height = max(0, window_height - current_y);
+
+    BufferParams slice_params = buffer_params;
+
+    slice_params.full_y = max(slice_window_full_y - overscan, buffer_params.full_y);
+    slice_params.window_y = slice_window_full_y - slice_params.full_y;
 
-    BufferParams slide_params = buffer_params;
-    slide_params.full_y = buffer_params.full_y + current_y;
     if (i < num_works - 1) {
-      slide_params.height = min(slice_height, remaining_height);
+      slice_params.window_height = min(slice_window_height, remaining_window_height);
     }
     else {
-      slide_params.height = remaining_height;
+      slice_params.window_height = remaining_window_height;
     }
 
-    slide_params.update_offset_stride();
+    slice_params.height = slice_params.window_y + slice_params.window_height + overscan;
+    slice_params.height = min(slice_params.height,
+                              buffer_params.height + buffer_params.full_y - slice_params.full_y);
 
-    callback(path_trace_works[i].get(), slide_params);
+    slice_params.update_offset_stride();
 
-    current_y += slide_params.height;
+    callback(path_trace_works[i].get(), slice_params);
+
+    current_y += slice_params.window_height;
   }
 }
 
 void PathTrace::update_allocated_work_buffer_params()
 {
+  const int overscan = tile_manager_.get_tile_overscan();
   foreach_sliced_buffer_params(path_trace_works_,
                                work_balance_infos_,
                                big_tile_params_,
+                               overscan,
                                [](PathTraceWork *path_trace_work, const BufferParams &params) {
                                  RenderBuffers *buffers = path_trace_work->get_render_buffers();
                                  buffers->reset(params);
@@ -282,6 +293,12 @@ static BufferParams scale_buffer_params(const BufferParams &params, int resoluti
 
   scaled_params.width = max(1, params.width / resolution_divider);
   scaled_params.height = max(1, params.height / resolution_divider);
+
+  scaled_params.window_x = params.window_x / resolution_divider;
+  scaled_params.window_y = params.window_y / resolution_divider;
+  scaled_params.window_width = params.window_width / resolution_divider;
+  scaled_params.window_height = params.window_height / resolution_divider;
+
   scaled_params.full_x = params.full_x / resolution_divider;
   scaled_params.full_y = params.full_y / resolution_divider;
   scaled_params.full_width = params.full_width / resolution_divider;
@@ -300,9 +317,12 @@ void PathTrace::update_effective_work_buffer_params(const RenderWork &render_wor
   const BufferParams scaled_big_tile_params = scale_buffer_params(big_tile_params_,
                                                                   resolution_divider);
 
+  const int overscan = tile_manager_.get_tile_overscan();
+
   foreach_sliced_buffer_params(path_trace_works_,
                                work_balance_infos_,
                                scaled_big_tile_params,
+                               overscan,
                                [&](PathTraceWork *path_trace_work, const BufferParams params) {
                                  path_trace_work->set_effective_buffer_params(
                                      scaled_full_params, scaled_big_tile_params, params);
@@ -1005,12 +1025,12 @@ bool PathTrace::set_render_tile_pixels(PassAccessor &pass_accessor,
 int2 PathTrace::get_render_tile_size() const
 {
   if (full_frame_state_.render_buffers) {
-    return make_int2(full_frame_state_.render_buffers->params.width,
-                     full_frame_state_.render_buffers->params.height);
+    return make_int2(full_frame_state_.render_buffers->params.window_width,
+                     full_frame_state_.render_buffers->params.window_height);
   }
 
   const Tile &tile = tile_manager_.get_current_tile();
-  return make_int2(tile.width, tile.height);
+  return make_int2(tile.window_width, tile.window_height);
 }
 
 int2 PathTrace::get_render_tile_offset() const
@@ -1020,7 +1040,7 @@ int2 PathTrace::get_render_tile_offset() const
   }
 
   const Tile &tile = tile_manager_.get_current_tile();
-  return make_int2(tile.x, tile.y);
+  return make_int2(tile.x + tile.window_x, tile.y + tile.window_y);
 }
 
 int2 PathTrace::get_render_size() const

intern/cycles/integrator/path_trace_work.cpp

@@ -134,7 +134,8 @@ void PathTraceWork::copy_from_denoised_render_buffers(const RenderBuffers *rende
 bool PathTraceWork::get_render_tile_pixels(const PassAccessor &pass_accessor,
                                            const PassAccessor::Destination &destination)
 {
-  const int offset_y = effective_buffer_params_.full_y - effective_big_tile_params_.full_y;
+  const int offset_y = (effective_buffer_params_.full_y + effective_buffer_params_.window_y) -
+                       (effective_big_tile_params_.full_y + effective_big_tile_params_.window_y);
   const int width = effective_buffer_params_.width;
 
   PassAccessor::Destination slice_destination = destination;
@@ -191,8 +192,10 @@ PassAccessor::Destination PathTraceWork::get_display_destination_template(
   PassAccessor::Destination destination(film_->get_display_pass());
 
   const int2 display_texture_size = display->get_texture_size();
-  const int texture_x = effective_buffer_params_.full_x - effective_full_params_.full_x;
-  const int texture_y = effective_buffer_params_.full_y - effective_full_params_.full_y;
+  const int texture_x = effective_buffer_params_.full_x - effective_full_params_.full_x +
+                        effective_buffer_params_.window_x;
+  const int texture_y = effective_buffer_params_.full_y - effective_full_params_.full_y +
+                        effective_buffer_params_.window_y;
 
   destination.offset = texture_y * display_texture_size.x + texture_x;
   destination.stride = display_texture_size.x;

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -23,6 +23,7 @@
 #include "render/buffers.h"
 #include "render/scene.h"
 #include "util/util_logging.h"
+#include "util/util_string.h"
 #include "util/util_tbb.h"
 #include "util/util_time.h"
 
@@ -30,6 +31,38 @@
 
 CCL_NAMESPACE_BEGIN
 
+static size_t estimate_single_state_size()
+{
+  size_t state_size = 0;
+
+#define KERNEL_STRUCT_BEGIN(name) for (int array_index = 0;; array_index++) {
+#define KERNEL_STRUCT_MEMBER(parent_struct, type, name, feature) state_size += sizeof(type);
+#define KERNEL_STRUCT_ARRAY_MEMBER(parent_struct, type, name, feature) state_size += sizeof(type);
+#define KERNEL_STRUCT_END(name) \
+  break; \
+  }
+#define KERNEL_STRUCT_END_ARRAY(name, cpu_array_size, gpu_array_size) \
+  if (array_index >= gpu_array_size - 1) { \
+    break; \
+  } \
+  }
+/* TODO(sergey): Look into better estimation for fields which depend on scene features. Maybe
+ * maximum state calculation should happen as `alloc_work_memory()`, so that we can react to an
+ * updated scene state here.
+ * For until then use common value. Currently this size is only used for logging, but is weak to
+ * rely on this. */
+#define KERNEL_STRUCT_VOLUME_STACK_SIZE 4
+#include "kernel/integrator/integrator_state_template.h"
+#undef KERNEL_STRUCT_BEGIN
+#undef KERNEL_STRUCT_MEMBER
+#undef KERNEL_STRUCT_ARRAY_MEMBER
+#undef KERNEL_STRUCT_END
+#undef KERNEL_STRUCT_END_ARRAY
+#undef KERNEL_STRUCT_VOLUME_STACK_SIZE
+
+  return state_size;
+}
+
 PathTraceWorkGPU::PathTraceWorkGPU(Device *device,
                                    Film *film,
                                    DeviceScene *device_scene,
@@ -47,7 +80,7 @@ PathTraceWorkGPU::PathTraceWorkGPU(Device *device,
       num_queued_paths_(device, "num_queued_paths", MEM_READ_WRITE),
       work_tiles_(device, "work_tiles", MEM_READ_WRITE),
       display_rgba_half_(device, "display buffer half", MEM_READ_WRITE),
-      max_num_paths_(queue_->num_concurrent_states(sizeof(IntegratorStateCPU))),
+      max_num_paths_(queue_->num_concurrent_states(estimate_single_state_size())),
       min_num_active_paths_(queue_->num_concurrent_busy_states()),
       max_active_path_index_(0)
 {
@@ -96,16 +129,27 @@ void PathTraceWorkGPU::alloc_integrator_soa()
   break; \
   }
 #define KERNEL_STRUCT_END_ARRAY(name, cpu_array_size, gpu_array_size) \
-  if (array_index == gpu_array_size - 1) { \
+  if (array_index >= gpu_array_size - 1) { \
     break; \
   } \
   }
+#define KERNEL_STRUCT_VOLUME_STACK_SIZE (device_scene_->data.volume_stack_size)
 #include "kernel/integrator/integrator_state_template.h"
 #undef KERNEL_STRUCT_BEGIN
 #undef KERNEL_STRUCT_MEMBER
 #undef KERNEL_STRUCT_ARRAY_MEMBER
 #undef KERNEL_STRUCT_END
 #undef KERNEL_STRUCT_END_ARRAY
+#undef KERNEL_STRUCT_VOLUME_STACK_SIZE
+
+  if (VLOG_IS_ON(3)) {
+    size_t total_soa_size = 0;
+    for (auto &&soa_memory : integrator_state_soa_) {
+      total_soa_size += soa_memory->memory_size();
+    }
+
+    VLOG(3) << "GPU SoA state size: " << string_human_readable_size(total_soa_size);
+  }
 }
 
 void PathTraceWorkGPU::alloc_integrator_queue()
@@ -712,13 +756,13 @@ void PathTraceWorkGPU::copy_to_display_naive(PathTraceDisplay *display,
 {
   const int full_x = effective_buffer_params_.full_x;
   const int full_y = effective_buffer_params_.full_y;
-  const int width = effective_buffer_params_.width;
-  const int height = effective_buffer_params_.height;
-  const int final_width = buffers_->params.width;
-  const int final_height = buffers_->params.height;
+  const int width = effective_buffer_params_.window_width;
+  const int height = effective_buffer_params_.window_height;
+  const int final_width = buffers_->params.window_width;
+  const int final_height = buffers_->params.window_height;
 
-  const int texture_x = full_x - effective_full_params_.full_x;
-  const int texture_y = full_y - effective_full_params_.full_y;
+  const int texture_x = full_x - effective_full_params_.full_x + effective_buffer_params_.window_x;
+  const int texture_y = full_y - effective_full_params_.full_y + effective_buffer_params_.window_y;
 
   /* Re-allocate display memory if needed, and make sure the device pointer is allocated.
    *

intern/cycles/kernel/CMakeLists.txt

@@ -404,16 +404,27 @@ if(WITH_CYCLES_CUDA_BINARIES)
             -cuda-toolkit-dir "${cuda_toolkit_root_dir}"
         DEPENDS ${kernel_sources} cycles_cubin_cc)
     else()
-      add_custom_command(
-        OUTPUT ${cuda_file}
-        COMMAND ${cuda_nvcc_executable}
+      set(_cuda_nvcc_args
             -arch=${arch}
             ${CUDA_NVCC_FLAGS}
             --${format}
             ${CMAKE_CURRENT_SOURCE_DIR}${cuda_kernel_src}
             --ptxas-options="-v"
-            ${cuda_flags}
-        DEPENDS ${kernel_sources})
+            ${cuda_flags})
+
+      if(WITH_COMPILER_CCACHE AND CCACHE_PROGRAM)
+        add_custom_command(
+          OUTPUT ${cuda_file}
+          COMMAND ${CCACHE_PROGRAM} ${cuda_nvcc_executable} ${_cuda_nvcc_args}
+          DEPENDS ${kernel_sources})
+      else()
+        add_custom_command(
+          OUTPUT ${cuda_file}
+          COMMAND ${cuda_nvcc_executable} ${_cuda_nvcc_args}
+          DEPENDS ${kernel_sources})
+      endif()
+
+      unset(_cuda_nvcc_args)
     endif()
     delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${cuda_file}" ${CYCLES_INSTALL_PATH}/lib)
     list(APPEND cuda_cubins ${cuda_file})
@@ -472,20 +483,10 @@ endif()
 
 # HIP module
 
-if(WITH_CYCLES_HIP_BINARIES)
+if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
   # 64 bit only
   set(HIP_BITS 64)
 
-  # HIP version
-  execute_process(COMMAND ${HIP_HIPCC_EXECUTABLE} "--version" OUTPUT_VARIABLE HIPCC_OUT)
-  string(REGEX REPLACE ".*release ([0-9]+)\\.([0-9]+).*" "\\1" HIP_VERSION_MAJOR "${HIPCC_OUT}")
-  string(REGEX REPLACE ".*release ([0-9]+)\\.([0-9]+).*" "\\2" HIP_VERSION_MINOR "${HIPCC_OUT}")
-  set(HIP_VERSION "${HIP_VERSION_MAJOR}${HIP_VERSION_MINOR}")
-
-
-  message(WARNING
-    "HIP version ${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR} detected")
-
   # build for each arch
   set(hip_sources device/hip/kernel.cpp
     ${SRC_HEADERS}
@@ -542,23 +543,24 @@ if(WITH_CYCLES_HIP_BINARIES)
         -D WITH_NANOVDB
         -I "${NANOVDB_INCLUDE_DIR}")
     endif()
+
+    add_custom_command(
+      OUTPUT ${hip_file}
+      COMMAND ${HIP_HIPCC_EXECUTABLE}
+          -arch=${arch}
+          ${HIP_HIPCC_FLAGS}
+          --${format}
+          ${CMAKE_CURRENT_SOURCE_DIR}${hip_kernel_src}
+          ${hip_flags}
+      DEPENDS ${kernel_sources})
+    delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${hip_file}" ${CYCLES_INSTALL_PATH}/lib)
+    list(APPEND hip_fatbins ${hip_file})
   endmacro()
 
   set(prev_arch "none")
   foreach(arch ${CYCLES_HIP_BINARIES_ARCH})
-      set(hip_hipcc_executable ${HIP_HIPCC_EXECUTABLE})
-      set(hip_toolkit_root_dir ${HIP_TOOLKIT_ROOT_DIR})
-    if(DEFINED hip_hipcc_executable AND DEFINED hip_toolkit_root_dir)
-      # Compile regular kernel
-      CYCLES_HIP_KERNEL_ADD(${arch} ${prev_arch} kernel "" "${hip_sources}" FALSE)
-
-      if(WITH_CYCLES_HIP_BUILD_SERIAL)
-        set(prev_arch ${arch})
-      endif()
-
-      unset(hip_hipcc_executable)
-      unset(hip_toolkit_root_dir)
-    endif()
+    # Compile regular kernel
+    CYCLES_HIP_KERNEL_ADD(${arch} ${prev_arch} kernel "" "${hip_sources}" FALSE)
   endforeach()
 
   add_custom_target(cycles_kernel_hip ALL DEPENDS ${hip_fatbins})

intern/cycles/kernel/bvh/bvh_embree.h

@@ -106,9 +106,6 @@ ccl_device_inline void kernel_embree_convert_hit(const KernelGlobals *kg,
                                                  const RTCHit *hit,
                                                  Intersection *isect)
 {
-  bool is_hair = hit->geomID & 1;
-  isect->u = is_hair ? hit->u : 1.0f - hit->v - hit->u;
-  isect->v = is_hair ? hit->v : hit->u;
   isect->t = ray->tfar;
   isect->Ng = make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z);
   if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) {
@@ -121,27 +118,37 @@ ccl_device_inline void kernel_embree_convert_hit(const KernelGlobals *kg,
   else {
     isect->prim = hit->primID + (intptr_t)rtcGetGeometryUserData(
                                     rtcGetGeometry(kernel_data.bvh.scene, hit->geomID));
-    isect->object = OBJECT_NONE;
+    isect->object = hit->geomID / 2;
+  }
+
+  const bool is_hair = hit->geomID & 1;
+  if (is_hair) {
+    const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, isect->prim);
+    isect->type = segment.type;
+    isect->prim = segment.prim;
+    isect->u = hit->u;
+    isect->v = hit->v;
+  }
+  else {
+    isect->type = kernel_tex_fetch(__objects, isect->object).primitive_type;
+    isect->u = 1.0f - hit->v - hit->u;
+    isect->v = hit->u;
   }
-  isect->type = kernel_tex_fetch(__prim_type, isect->prim);
 }
 
-ccl_device_inline void kernel_embree_convert_sss_hit(const KernelGlobals *kg,
-                                                     const RTCRay *ray,
-                                                     const RTCHit *hit,
-                                                     Intersection *isect,
-                                                     int local_object_id)
+ccl_device_inline void kernel_embree_convert_sss_hit(
+    const KernelGlobals *kg, const RTCRay *ray, const RTCHit *hit, Intersection *isect, int object)
 {
   isect->u = 1.0f - hit->v - hit->u;
   isect->v = hit->u;
   isect->t = ray->tfar;
   isect->Ng = make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z);
   RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData(
-      rtcGetGeometry(kernel_data.bvh.scene, local_object_id * 2));
+      rtcGetGeometry(kernel_data.bvh.scene, object * 2));
   isect->prim = hit->primID +
                 (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID));
-  isect->object = local_object_id;
-  isect->type = kernel_tex_fetch(__prim_type, isect->prim);
+  isect->object = object;
+  isect->type = kernel_tex_fetch(__objects, object).primitive_type;
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/bvh/bvh_shadow_all.h

@@ -130,7 +130,6 @@ ccl_device_inline
         if (prim_addr >= 0) {
           const int prim_addr2 = __float_as_int(leaf.y);
           const uint type = __float_as_int(leaf.w);
-          const uint p_type = type & PRIMITIVE_ALL;
 
           /* pop */
           node_addr = traversal_stack[stack_ptr];
@@ -138,14 +137,15 @@ ccl_device_inline
 
           /* primitive intersection */
           while (prim_addr < prim_addr2) {
-            kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) == p_type);
+            kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) ==
+                          (type & PRIMITIVE_ALL));
             bool hit;
 
             /* todo: specialized intersect functions which don't fill in
              * isect unless needed and check SD_HAS_TRANSPARENT_SHADOW?
              * might give a few % performance improvement */
 
-            switch (p_type) {
+            switch (type & PRIMITIVE_ALL) {
               case PRIMITIVE_TRIANGLE: {
                 hit = triangle_intersect(
                     kg, isect, P, dir, isect_t, visibility, object, prim_addr);
@@ -163,17 +163,20 @@ ccl_device_inline
               case PRIMITIVE_MOTION_CURVE_THICK:
               case PRIMITIVE_CURVE_RIBBON:
               case PRIMITIVE_MOTION_CURVE_RIBBON: {
-                const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
-                hit = curve_intersect(kg,
-                                      isect,
-                                      P,
-                                      dir,
-                                      isect_t,
-                                      visibility,
-                                      object,
-                                      prim_addr,
-                                      ray->time,
-                                      curve_type);
+                if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) {
+                  const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
+                  if (ray->time < prim_time.x || ray->time > prim_time.y) {
+                    hit = false;
+                    break;
+                  }
+                }
+
+                const int curve_object = kernel_tex_fetch(__prim_object, prim_addr);
+                const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
+                const int curve_prim = kernel_tex_fetch(__prim_index, prim_addr);
+                hit = curve_intersect(
+                    kg, isect, P, dir, isect_t, curve_object, curve_prim, ray->time, curve_type);
+
                 break;
               }
 #endif

intern/cycles/kernel/bvh/bvh_traversal.h

@@ -165,18 +165,18 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(const KernelGlobals *kg,
             case PRIMITIVE_CURVE_RIBBON:
             case PRIMITIVE_MOTION_CURVE_RIBBON: {
               for (; prim_addr < prim_addr2; prim_addr++) {
-                const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
-                kernel_assert((curve_type & PRIMITIVE_ALL) == (type & PRIMITIVE_ALL));
-                const bool hit = curve_intersect(kg,
-                                                 isect,
-                                                 P,
-                                                 dir,
-                                                 isect->t,
-                                                 visibility,
-                                                 object,
-                                                 prim_addr,
-                                                 ray->time,
-                                                 curve_type);
+                if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) {
+                  const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
+                  if (ray->time < prim_time.x || ray->time > prim_time.y) {
+                    continue;
+                  }
+                }
+
+                const int curve_object = kernel_tex_fetch(__prim_object, prim_addr);
+                const int curve_prim = kernel_tex_fetch(__prim_index, prim_addr);
+                const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
+                const bool hit = curve_intersect(
+                    kg, isect, P, dir, isect->t, curve_object, curve_prim, ray->time, curve_type);
                 if (hit) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)

intern/cycles/kernel/bvh/bvh_util.h

@@ -118,19 +118,18 @@ ccl_device_inline void sort_intersections(Intersection *hits, uint num_hits)
 ccl_device_forceinline int intersection_get_shader_flags(const KernelGlobals *ccl_restrict kg,
                                                          const Intersection *ccl_restrict isect)
 {
-  const int prim = kernel_tex_fetch(__prim_index, isect->prim);
+  const int prim = isect->prim;
   int shader = 0;
 
 #ifdef __HAIR__
-  if (kernel_tex_fetch(__prim_type, isect->prim) & PRIMITIVE_ALL_TRIANGLE)
+  if (isect->type & PRIMITIVE_ALL_TRIANGLE)
 #endif
   {
     shader = kernel_tex_fetch(__tri_shader, prim);
   }
 #ifdef __HAIR__
   else {
-    float4 str = kernel_tex_fetch(__curves, prim);
-    shader = __float_as_int(str.z);
+    shader = kernel_tex_fetch(__curves, prim).shader_id;
   }
 #endif
 
@@ -138,21 +137,19 @@ ccl_device_forceinline int intersection_get_shader_flags(const KernelGlobals *cc
 }
 
 ccl_device_forceinline int intersection_get_shader_from_isect_prim(
-    const KernelGlobals *ccl_restrict kg, const int isect_prim)
+    const KernelGlobals *ccl_restrict kg, const int prim, const int isect_type)
 {
-  const int prim = kernel_tex_fetch(__prim_index, isect_prim);
   int shader = 0;
 
 #ifdef __HAIR__
-  if (kernel_tex_fetch(__prim_type, isect_prim) & PRIMITIVE_ALL_TRIANGLE)
+  if (isect_type & PRIMITIVE_ALL_TRIANGLE)
 #endif
   {
     shader = kernel_tex_fetch(__tri_shader, prim);
   }
 #ifdef __HAIR__
   else {
-    float4 str = kernel_tex_fetch(__curves, prim);
-    shader = __float_as_int(str.z);
+    shader = kernel_tex_fetch(__curves, prim).shader_id;
   }
 #endif
 
@@ -162,25 +159,13 @@ ccl_device_forceinline int intersection_get_shader_from_isect_prim(
 ccl_device_forceinline int intersection_get_shader(const KernelGlobals *ccl_restrict kg,
                                                    const Intersection *ccl_restrict isect)
 {
-  return intersection_get_shader_from_isect_prim(kg, isect->prim);
-}
-
-ccl_device_forceinline int intersection_get_object(const KernelGlobals *ccl_restrict kg,
-                                                   const Intersection *ccl_restrict isect)
-{
-  if (isect->object != OBJECT_NONE) {
-    return isect->object;
-  }
-
-  return kernel_tex_fetch(__prim_object, isect->prim);
+  return intersection_get_shader_from_isect_prim(kg, isect->prim, isect->type);
 }
 
 ccl_device_forceinline int intersection_get_object_flags(const KernelGlobals *ccl_restrict kg,
                                                          const Intersection *ccl_restrict isect)
 {
-  const int object = intersection_get_object(kg, isect);
-
-  return kernel_tex_fetch(__object_flag, object);
+  return kernel_tex_fetch(__object_flag, isect->object);
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/closure/bsdf_principled_diffuse.h

@@ -35,21 +35,25 @@ static_assert(sizeof(ShaderClosure) >= sizeof(PrincipledDiffuseBsdf),
               "PrincipledDiffuseBsdf is too large!");
 
 ccl_device float3 calculate_principled_diffuse_brdf(
-    const PrincipledDiffuseBsdf *bsdf, float3 N, float3 V, float3 L, float3 H, float *pdf)
+    const PrincipledDiffuseBsdf *bsdf, float3 N, float3 V, float3 L, float *pdf)
 {
   float NdotL = dot(N, L);
-  float NdotV = dot(N, V);
 
-  if (NdotL <= 0 || NdotV <= 0) {
-    *pdf = 0.0f;
+  if (NdotL <= 0) {
     return make_float3(0.0f, 0.0f, 0.0f);
   }
 
-  float LdotH = dot(L, H);
+  float NdotV = dot(N, V);
+
+  /* H = normalize(L + V);  // Bissector of an angle between L and V
+   * LH2 = 2 * dot(L, H)^2 = 2cos(x)^2 = cos(2x) + 1 = dot(L, V) + 1,
+   * half-angle x between L and V is at most 90 deg
+   */
+  float LH2 = dot(L, V) + 1;
 
   float FL = schlick_fresnel(NdotL), FV = schlick_fresnel(NdotV);
-  const float Fd90 = 0.5f + 2.0f * LdotH * LdotH * bsdf->roughness;
-  float Fd = (1.0f * (1.0f - FL) + Fd90 * FL) * (1.0f * (1.0f - FV) + Fd90 * FV);
+  const float Fd90 = 0.5f + LH2 * bsdf->roughness;
+  float Fd = (1.0f - FL + Fd90 * FL) * (1.0f - FV + Fd90 * FV);
 
   float value = M_1_PI_F * NdotL * Fd;
 
@@ -72,11 +76,10 @@ ccl_device float3 bsdf_principled_diffuse_eval_reflect(const ShaderClosure *sc,
   float3 N = bsdf->N;
   float3 V = I;         // outgoing
   float3 L = omega_in;  // incoming
-  float3 H = normalize(L + V);
 
   if (dot(N, omega_in) > 0.0f) {
     *pdf = fmaxf(dot(N, omega_in), 0.0f) * M_1_PI_F;
-    return calculate_principled_diffuse_brdf(bsdf, N, V, L, H, pdf);
+    return calculate_principled_diffuse_brdf(bsdf, N, V, L, pdf);
   }
   else {
     *pdf = 0.0f;
@@ -112,9 +115,7 @@ ccl_device int bsdf_principled_diffuse_sample(const ShaderClosure *sc,
   sample_cos_hemisphere(N, randu, randv, omega_in, pdf);
 
   if (dot(Ng, *omega_in) > 0) {
-    float3 H = normalize(I + *omega_in);
-
-    *eval = calculate_principled_diffuse_brdf(bsdf, N, I, *omega_in, H, pdf);
+    *eval = calculate_principled_diffuse_brdf(bsdf, N, I, *omega_in, pdf);
 
 #ifdef __RAY_DIFFERENTIALS__
     // TODO: find a better approximation for the diffuse bounce

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

@@ -424,8 +424,12 @@ ccl_device_inline void kernel_gpu_film_convert_common(const KernelFilmConvert *k
     return;
   }
 
-  const uint64_t render_buffer_offset = (uint64_t)render_pixel_index * kfilm_convert->pass_stride;
-  ccl_global const float *buffer = render_buffer + render_buffer_offset;
+  const int x = render_pixel_index % width;
+  const int y = render_pixel_index / width;
+
+  ccl_global const float *buffer = render_buffer + offset + x * kfilm_convert->pass_stride +
+                                   y * stride * kfilm_convert->pass_stride;
+
   ccl_global float *pixel = pixels +
                             (render_pixel_index + dst_offset) * kfilm_convert->pixel_stride;
 
@@ -451,17 +455,17 @@ ccl_device_inline void kernel_gpu_film_convert_half_rgba_common_rgba(
     return;
   }
 
-  const uint64_t render_buffer_offset = (uint64_t)render_pixel_index * kfilm_convert->pass_stride;
-  ccl_global const float *buffer = render_buffer + render_buffer_offset;
+  const int x = render_pixel_index % width;
+  const int y = render_pixel_index / width;
+
+  ccl_global const float *buffer = render_buffer + offset + x * kfilm_convert->pass_stride +
+                                   y * stride * kfilm_convert->pass_stride;
 
   float pixel[4];
   processor(kfilm_convert, buffer, pixel);
 
   film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel);
 
-  const int x = render_pixel_index % width;
-  const int y = render_pixel_index / width;
-
   ccl_global half4 *out = ((ccl_global half4 *)rgba) + rgba_offset + y * rgba_stride + x;
   float4_store_half((ccl_global half *)out, make_float4(pixel[0], pixel[1], pixel[2], pixel[3]));
 }

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

@@ -41,22 +41,15 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
   return (T *)(((uint64_t)optixGetPayload_3() << 32) | optixGetPayload_2());
 }
 
-template<bool always = false> ccl_device_forceinline uint get_object_id()
+ccl_device_forceinline int get_object_id()
 {
 #ifdef __OBJECT_MOTION__
-  /* Always get the the instance ID from the TLAS.
+  /* Always get the the instance ID from the TLAS
    * There might be a motion transform node between TLAS and BLAS which does not have one. */
-  uint object = optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
+  return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
 #else
-  uint object = optixGetInstanceId();
+  return optixGetInstanceId();
 #endif
-  /* Choose between always returning object ID or only for instances. */
-  if (always || (object & 1) == 0)
-    /* Can just remove the low bit since instance always contains object ID. */
-    return object >> 1;
-  else
-    /* Set to OBJECT_NONE if this is not an instanced object. */
-    return OBJECT_NONE;
 }
 
 extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()
@@ -108,7 +101,7 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
 #endif
 
 #ifdef __BVH_LOCAL__
-  const uint object = get_object_id<true>();
+  const int object = get_object_id();
   if (object != optixGetPayload_4() /* local_object */) {
     /* Only intersect with matching object. */
     return optixIgnoreIntersection();
@@ -152,21 +145,23 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
     local_isect->num_hits = 1;
   }
 
+  const int prim = optixGetPrimitiveIndex();
+
   Intersection *isect = &local_isect->hits[hit];
   isect->t = optixGetRayTmax();
-  isect->prim = optixGetPrimitiveIndex();
+  isect->prim = prim;
   isect->object = get_object_id();
-  isect->type = kernel_tex_fetch(__prim_type, isect->prim);
+  isect->type = kernel_tex_fetch(__objects, isect->object).primitive_type;
 
   const float2 barycentrics = optixGetTriangleBarycentrics();
   isect->u = 1.0f - barycentrics.y - barycentrics.x;
   isect->v = barycentrics.x;
 
   /* Record geometric normal. */
-  const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim);
-  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0));
-  const float3 tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1));
-  const float3 tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));
+  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
+  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0));
+  const float3 tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1));
+  const float3 tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
   local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
 
   /* Continue tracing (without this the trace call would return after the first hit). */
@@ -179,25 +174,32 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
 #ifdef __SHADOW_RECORD_ALL__
   bool ignore_intersection = false;
 
-  const uint prim = optixGetPrimitiveIndex();
+  int prim = optixGetPrimitiveIndex();
+  const uint object = get_object_id();
 #  ifdef __VISIBILITY_FLAG__
   const uint visibility = optixGetPayload_4();
-  if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
+  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
     ignore_intersection = true;
   }
 #  endif
 
   float u = 0.0f, v = 0.0f;
+  int type = 0;
   if (optixIsTriangleHit()) {
     const float2 barycentrics = optixGetTriangleBarycentrics();
     u = 1.0f - barycentrics.y - barycentrics.x;
     v = barycentrics.x;
+    type = kernel_tex_fetch(__objects, object).primitive_type;
   }
 #  ifdef __HAIR__
   else {
     u = __uint_as_float(optixGetAttribute_0());
     v = __uint_as_float(optixGetAttribute_1());
 
+    const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
+    type = segment.type;
+    prim = segment.prim;
+
     /* Filter out curve endcaps. */
     if (u == 0.0f || u == 1.0f) {
       ignore_intersection = true;
@@ -245,8 +247,8 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
     isect->v = v;
     isect->t = optixGetRayTmax();
     isect->prim = prim;
-    isect->object = get_object_id();
-    isect->type = kernel_tex_fetch(__prim_type, prim);
+    isect->object = object;
+    isect->type = type;
 
 #  ifdef __TRANSPARENT_SHADOWS__
     /* Detect if this surface has a shader with transparent shadows. */
@@ -274,15 +276,14 @@ extern "C" __global__ void __anyhit__kernel_optix_volume_test()
   }
 #endif
 
+  const uint object = get_object_id();
 #ifdef __VISIBILITY_FLAG__
-  const uint prim = optixGetPrimitiveIndex();
   const uint visibility = optixGetPayload_4();
-  if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
+  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
     return optixIgnoreIntersection();
   }
 #endif
 
-  const uint object = get_object_id<true>();
   if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
     return optixIgnoreIntersection();
   }
@@ -301,9 +302,9 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
 #endif
 
 #ifdef __VISIBILITY_FLAG__
-  const uint prim = optixGetPrimitiveIndex();
+  const uint object = get_object_id();
   const uint visibility = optixGetPayload_4();
-  if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
+  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
     return optixIgnoreIntersection();
   }
 
@@ -316,28 +317,39 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
 
 extern "C" __global__ void __closesthit__kernel_optix_hit()
 {
+  const int object = get_object_id();
+  const int prim = optixGetPrimitiveIndex();
+
   optixSetPayload_0(__float_as_uint(optixGetRayTmax())); /* Intersection distance */
-  optixSetPayload_3(optixGetPrimitiveIndex());
-  optixSetPayload_4(get_object_id());
-  /* Can be PRIMITIVE_TRIANGLE and PRIMITIVE_MOTION_TRIANGLE or curve type and segment index. */
-  optixSetPayload_5(kernel_tex_fetch(__prim_type, optixGetPrimitiveIndex()));
+  optixSetPayload_4(object);
 
   if (optixIsTriangleHit()) {
     const float2 barycentrics = optixGetTriangleBarycentrics();
     optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x));
     optixSetPayload_2(__float_as_uint(barycentrics.x));
+    optixSetPayload_3(prim);
+    optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type);
   }
   else {
+    const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
     optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */
     optixSetPayload_2(optixGetAttribute_1());
+    optixSetPayload_3(segment.prim);
+    optixSetPayload_5(segment.type);
   }
 }
 
 #ifdef __HAIR__
-ccl_device_inline void optix_intersection_curve(const uint prim, const uint type)
+ccl_device_inline void optix_intersection_curve(const int prim, const int type)
 {
-  const uint object = get_object_id<true>();
+  const int object = get_object_id();
+
+#  ifdef __VISIBILITY_FLAG__
   const uint visibility = optixGetPayload_4();
+  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
+    return;
+  }
+#  endif
 
   float3 P = optixGetObjectRayOrigin();
   float3 dir = optixGetObjectRayDirection();
@@ -358,7 +370,7 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
   if (isect.t != FLT_MAX)
     isect.t *= len;
 
-  if (curve_intersect(NULL, &isect, P, dir, isect.t, visibility, object, prim, time, type)) {
+  if (curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
     optixReportIntersection(isect.t / len,
                             type & PRIMITIVE_ALL,
                             __float_as_int(isect.u),  /* Attribute_0 */
@@ -368,9 +380,9 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
 
 extern "C" __global__ void __intersection__curve_ribbon()
 {
-  const uint prim = optixGetPrimitiveIndex();
-  const uint type = kernel_tex_fetch(__prim_type, prim);
-
+  const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, optixGetPrimitiveIndex());
+  const int prim = segment.prim;
+  const int type = segment.type;
   if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) {
     optix_intersection_curve(prim, type);
   }

intern/cycles/kernel/geom/geom_curve.h

@@ -34,8 +34,8 @@ ccl_device float curve_attribute_float(const KernelGlobals *kg,
                                        float *dy)
 {
   if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
-    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     int k1 = k0 + 1;
 
     float f0 = kernel_tex_fetch(__attributes_float, desc.offset + k0);
@@ -76,8 +76,8 @@ ccl_device float2 curve_attribute_float2(const KernelGlobals *kg,
                                          float2 *dy)
 {
   if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
-    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     int k1 = k0 + 1;
 
     float2 f0 = kernel_tex_fetch(__attributes_float2, desc.offset + k0);
@@ -122,8 +122,8 @@ ccl_device float3 curve_attribute_float3(const KernelGlobals *kg,
                                          float3 *dy)
 {
   if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
-    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     int k1 = k0 + 1;
 
     float3 f0 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + k0));
@@ -164,8 +164,8 @@ ccl_device float4 curve_attribute_float4(const KernelGlobals *kg,
                                          float4 *dy)
 {
   if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
-    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     int k1 = k0 + 1;
 
     float4 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + k0);
@@ -206,8 +206,8 @@ ccl_device float curve_thickness(const KernelGlobals *kg, const ShaderData *sd)
   float r = 0.0f;
 
   if (sd->type & PRIMITIVE_ALL_CURVE) {
-    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     int k1 = k0 + 1;
 
     float4 P_curve[2];
@@ -231,8 +231,8 @@ ccl_device float curve_thickness(const KernelGlobals *kg, const ShaderData *sd)
 
 ccl_device float3 curve_motion_center_location(const KernelGlobals *kg, const ShaderData *sd)
 {
-  float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-  int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+  KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+  int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
   int k1 = k0 + 1;
 
   float4 P_curve[2];

intern/cycles/kernel/geom/geom_curve_intersect.h

@@ -630,33 +630,19 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
                                             const float3 P,
                                             const float3 dir,
                                             const float tmax,
-                                            uint visibility,
                                             int object,
-                                            int curveAddr,
+                                            int prim,
                                             float time,
                                             int type)
 {
   const bool is_motion = (type & PRIMITIVE_ALL_MOTION);
 
-#  ifndef __KERNEL_OPTIX__ /* See OptiX motion flag OPTIX_MOTION_FLAG_[START|END]_VANISH */
-  if (is_motion && kernel_data.bvh.use_bvh_steps) {
-    const float2 prim_time = kernel_tex_fetch(__prim_time, curveAddr);
-    if (time < prim_time.x || time > prim_time.y) {
-      return false;
-    }
-  }
-#  endif
+  KernelCurve kcurve = kernel_tex_fetch(__curves, prim);
 
-  int segment = PRIMITIVE_UNPACK_SEGMENT(type);
-  int prim = kernel_tex_fetch(__prim_index, curveAddr);
-
-  float4 v00 = kernel_tex_fetch(__curves, prim);
-
-  int k0 = __float_as_int(v00.x) + segment;
+  int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
   int k1 = k0 + 1;
-
-  int ka = max(k0 - 1, __float_as_int(v00.x));
-  int kb = min(k1 + 1, __float_as_int(v00.x) + __float_as_int(v00.y) - 1);
+  int ka = max(k0 - 1, kcurve.first_key);
+  int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
 
   float4 curve[4];
   if (!is_motion) {
@@ -666,21 +652,14 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
     curve[3] = kernel_tex_fetch(__curve_keys, kb);
   }
   else {
-    int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
-    motion_curve_keys(kg, fobject, prim, time, ka, k0, k1, kb, curve);
+    motion_curve_keys(kg, object, prim, time, ka, k0, k1, kb, curve);
   }
 
-#  ifdef __VISIBILITY_FLAG__
-  if (!(kernel_tex_fetch(__prim_visibility, curveAddr) & visibility)) {
-    return false;
-  }
-#  endif
-
   if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) {
     /* todo: adaptive number of subdivisions could help performance here. */
     const int subdivisions = kernel_data.bvh.curve_subdivisions;
     if (ribbon_intersect(P, dir, tmax, subdivisions, curve, isect)) {
-      isect->prim = curveAddr;
+      isect->prim = prim;
       isect->object = object;
       isect->type = type;
       return true;
@@ -690,7 +669,7 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
   }
   else {
     if (curve_intersect_recursive(P, dir, tmax, curve, isect)) {
-      isect->prim = curveAddr;
+      isect->prim = prim;
       isect->object = object;
       isect->type = type;
       return true;
@@ -708,7 +687,7 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
                                           const int isect_object,
                                           const int isect_prim)
 {
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_inverse_transform(kg, sd);
 
     P = transform_point(&tfm, P);
@@ -716,14 +695,12 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
     D = safe_normalize_len(D, &t);
   }
 
-  int prim = kernel_tex_fetch(__prim_index, isect_prim);
-  float4 v00 = kernel_tex_fetch(__curves, prim);
+  KernelCurve kcurve = kernel_tex_fetch(__curves, isect_prim);
 
-  int k0 = __float_as_int(v00.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+  int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
   int k1 = k0 + 1;
-
-  int ka = max(k0 - 1, __float_as_int(v00.x));
-  int kb = min(k1 + 1, __float_as_int(v00.x) + __float_as_int(v00.y) - 1);
+  int ka = max(k0 - 1, kcurve.first_key);
+  int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
 
   float4 P_curve[4];
 
@@ -780,15 +757,13 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
   sd->dPdv = cross(dPdu, sd->Ng);
 #  endif
 
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }
 
   sd->P = P;
-
-  float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-  sd->shader = __float_as_int(curvedata.z);
+  sd->shader = kernel_tex_fetch(__curves, sd->prim).shader_id;
 }
 
 #endif

intern/cycles/kernel/geom/geom_motion_triangle.h

@@ -72,9 +72,9 @@ ccl_device_inline void motion_triangle_verts_for_step(const KernelGlobals *kg,
 {
   if (step == numsteps) {
     /* center step: regular vertex location */
-    verts[0] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-    verts[1] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-    verts[2] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+    verts[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+    verts[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+    verts[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
   }
   else {
     /* center step not store in this array */

intern/cycles/kernel/geom/geom_motion_triangle_intersect.h

@@ -44,7 +44,7 @@ ccl_device_inline float3 motion_triangle_refine(const KernelGlobals *kg,
                                                 float3 verts[3])
 {
 #ifdef __INTERSECTION_REFINE__
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     if (UNLIKELY(t == 0.0f)) {
       return P;
     }
@@ -70,7 +70,7 @@ ccl_device_inline float3 motion_triangle_refine(const KernelGlobals *kg,
   /* Compute refined position. */
   P = P + D * rt;
 
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }
@@ -106,7 +106,7 @@ ccl_device_inline
   return motion_triangle_refine(kg, sd, P, D, t, isect_object, isect_prim, verts);
 #  else
 #    ifdef __INTERSECTION_REFINE__
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_inverse_transform(kg, sd);
 
     P = transform_point(&tfm, P);
@@ -128,7 +128,7 @@ ccl_device_inline
 
   P = P + D * rt;
 
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }
@@ -186,8 +186,9 @@ ccl_device_inline bool motion_triangle_intersect(const KernelGlobals *kg,
       isect->t = t;
       isect->u = u;
       isect->v = v;
-      isect->prim = prim_addr;
-      isect->object = object;
+      isect->prim = prim;
+      isect->object = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) :
+                                                object;
       isect->type = PRIMITIVE_MOTION_TRIANGLE;
       return true;
     }
@@ -288,8 +289,8 @@ ccl_device_inline bool motion_triangle_intersect_local(const KernelGlobals *kg,
   isect->t = t;
   isect->u = u;
   isect->v = v;
-  isect->prim = prim_addr;
-  isect->object = object;
+  isect->prim = prim;
+  isect->object = local_object;
   isect->type = PRIMITIVE_MOTION_TRIANGLE;
 
   /* Record geometric normal. */

intern/cycles/kernel/geom/geom_shader_data.h

@@ -37,7 +37,7 @@ ccl_device void shader_setup_object_transforms(const KernelGlobals *ccl_restrict
 #endif
 
 /* TODO: break this up if it helps reduce register pressure to load data from
- * global memory as we write it to shaderdata. */
+ * global memory as we write it to shader-data. */
 ccl_device_inline void shader_setup_from_ray(const KernelGlobals *ccl_restrict kg,
                                              ShaderData *ccl_restrict sd,
                                              const Ray *ccl_restrict ray,
@@ -52,10 +52,9 @@ ccl_device_inline void shader_setup_from_ray(const KernelGlobals *ccl_restrict k
   sd->v = isect->v;
   sd->ray_length = isect->t;
   sd->type = isect->type;
-  sd->object = (isect->object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, isect->prim) :
-                                                isect->object;
+  sd->object = isect->object;
   sd->object_flag = kernel_tex_fetch(__object_flag, sd->object);
-  sd->prim = kernel_tex_fetch(__prim_index, isect->prim);
+  sd->prim = isect->prim;
   sd->lamp = LAMP_NONE;
   sd->flag = 0;
 

intern/cycles/kernel/geom/geom_triangle.h

@@ -29,9 +29,9 @@ ccl_device_inline float3 triangle_normal(const KernelGlobals *kg, ShaderData *sd
 {
   /* load triangle vertices */
   const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
-  const float3 v0 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-  const float3 v1 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-  const float3 v2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+  const float3 v0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+  const float3 v1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+  const float3 v2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
 
   /* return normal */
   if (sd->object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
@@ -54,9 +54,9 @@ ccl_device_inline void triangle_point_normal(const KernelGlobals *kg,
 {
   /* load triangle vertices */
   const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
-  float3 v0 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-  float3 v1 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-  float3 v2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+  float3 v0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+  float3 v1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+  float3 v2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
   /* compute point */
   float t = 1.0f - u - v;
   *P = (u * v0 + v * v1 + t * v2);
@@ -78,9 +78,9 @@ ccl_device_inline void triangle_point_normal(const KernelGlobals *kg,
 ccl_device_inline void triangle_vertices(const KernelGlobals *kg, int prim, float3 P[3])
 {
   const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
-  P[0] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-  P[1] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-  P[2] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+  P[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+  P[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+  P[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
 }
 
 /* Triangle vertex locations and vertex normals */
@@ -91,9 +91,9 @@ ccl_device_inline void triangle_vertices_and_normals(const KernelGlobals *kg,
                                                      float3 N[3])
 {
   const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
-  P[0] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-  P[1] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-  P[2] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+  P[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+  P[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+  P[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
   N[0] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.x));
   N[1] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.y));
   N[2] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.z));
@@ -145,9 +145,9 @@ ccl_device_inline void triangle_dPdudv(const KernelGlobals *kg,
 {
   /* fetch triangle vertex coordinates */
   const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
-  const float3 p0 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-  const float3 p1 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-  const float3 p2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+  const float3 p0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+  const float3 p1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+  const float3 p2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
 
   /* compute derivatives of P w.r.t. uv */
   *dPdu = (p0 - p2);

intern/cycles/kernel/geom/geom_triangle_intersect.h

@@ -35,13 +35,14 @@ ccl_device_inline bool triangle_intersect(const KernelGlobals *kg,
                                           int object,
                                           int prim_addr)
 {
-  const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
+  const int prim = kernel_tex_fetch(__prim_index, prim_addr);
+  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
 #if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
-  const ssef *ssef_verts = (ssef *)&kg->__prim_tri_verts.data[tri_vindex];
+  const ssef *ssef_verts = (ssef *)&kg->__tri_verts.data[tri_vindex];
 #else
-  const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0),
-               tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1),
-               tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2);
+  const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
+               tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
+               tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
 #endif
   float t, u, v;
   if (ray_triangle_intersect(P,
@@ -64,8 +65,9 @@ ccl_device_inline bool triangle_intersect(const KernelGlobals *kg,
     if (kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
 #endif
     {
-      isect->prim = prim_addr;
-      isect->object = object;
+      isect->object = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) :
+                                                object;
+      isect->prim = prim;
       isect->type = PRIMITIVE_TRIANGLE;
       isect->u = u;
       isect->v = v;
@@ -102,13 +104,14 @@ ccl_device_inline bool triangle_intersect_local(const KernelGlobals *kg,
     }
   }
 
-  const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
+  const int prim = kernel_tex_fetch(__prim_index, prim_addr);
+  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
 #  if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
-  const ssef *ssef_verts = (ssef *)&kg->__prim_tri_verts.data[tri_vindex];
+  const ssef *ssef_verts = (ssef *)&kg->__tri_verts.data[tri_vindex];
 #  else
-  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0)),
-               tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1)),
-               tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));
+  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0)),
+               tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1)),
+               tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
 #  endif
   float t, u, v;
   if (!ray_triangle_intersect(P,
@@ -167,8 +170,8 @@ ccl_device_inline bool triangle_intersect_local(const KernelGlobals *kg,
 
   /* Record intersection. */
   Intersection *isect = &local_isect->hits[hit];
-  isect->prim = prim_addr;
-  isect->object = object;
+  isect->prim = prim;
+  isect->object = local_object;
   isect->type = PRIMITIVE_TRIANGLE;
   isect->u = u;
   isect->v = v;
@@ -176,9 +179,9 @@ ccl_device_inline bool triangle_intersect_local(const KernelGlobals *kg,
 
   /* Record geometric normal. */
 #  if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
-  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0)),
-               tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1)),
-               tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));
+  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0)),
+               tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1)),
+               tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
 #  endif
   local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
 
@@ -206,7 +209,7 @@ ccl_device_inline float3 triangle_refine(const KernelGlobals *kg,
                                          const int isect_prim)
 {
 #ifdef __INTERSECTION_REFINE__
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     if (UNLIKELY(t == 0.0f)) {
       return P;
     }
@@ -219,10 +222,10 @@ ccl_device_inline float3 triangle_refine(const KernelGlobals *kg,
 
   P = P + D * t;
 
-  const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect_prim);
-  const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0),
-               tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1),
-               tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2);
+  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
+  const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
+               tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
+               tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
   float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
   float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
   float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
@@ -239,7 +242,7 @@ ccl_device_inline float3 triangle_refine(const KernelGlobals *kg,
     P = P + D * rt;
   }
 
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }
@@ -265,7 +268,7 @@ ccl_device_inline float3 triangle_refine_local(const KernelGlobals *kg,
   /* t is always in world space with OptiX. */
   return triangle_refine(kg, sd, P, D, t, isect_object, isect_prim);
 #else
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_inverse_transform(kg, sd);
 
     P = transform_point(&tfm, P);
@@ -276,10 +279,10 @@ ccl_device_inline float3 triangle_refine_local(const KernelGlobals *kg,
   P = P + D * t;
 
 #  ifdef __INTERSECTION_REFINE__
-  const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect_prim);
-  const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0),
-               tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1),
-               tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2);
+  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
+  const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
+               tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
+               tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
   float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
   float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
   float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
@@ -297,7 +300,7 @@ ccl_device_inline float3 triangle_refine_local(const KernelGlobals *kg,
   }
 #  endif /* __INTERSECTION_REFINE__ */
 
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }

intern/cycles/kernel/integrator/integrator_intersect_closest.h

@@ -160,10 +160,7 @@ ccl_device void integrator_intersect_closest(INTEGRATOR_STATE_ARGS)
   if (path_state_ao_bounce(INTEGRATOR_STATE_PASS)) {
     ray.t = kernel_data.integrator.ao_bounces_distance;
 
-    const int last_object = last_isect_object != OBJECT_NONE ?
-                                last_isect_object :
-                                kernel_tex_fetch(__prim_object, last_isect_prim);
-    const float object_ao_distance = kernel_tex_fetch(__objects, last_object).ao_distance;
+    const float object_ao_distance = kernel_tex_fetch(__objects, last_isect_object).ao_distance;
     if (object_ao_distance != 0.0f) {
       ray.t = object_ao_distance;
     }

intern/cycles/kernel/integrator/integrator_intersect_volume_stack.h

@@ -38,10 +38,13 @@ ccl_device void integrator_volume_stack_update_for_subsurface(INTEGRATOR_STATE_A
   volume_ray.P = from_P;
   volume_ray.D = normalize_len(to_P - from_P, &volume_ray.t);
 
+  /* Store to avoid global fetches on every intersection step. */
+  const uint volume_stack_size = kernel_data.volume_stack_size;
+
 #ifdef __VOLUME_RECORD_ALL__
-  Intersection hits[2 * VOLUME_STACK_SIZE + 1];
+  Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1];
   uint num_hits = scene_intersect_volume_all(
-      kg, &volume_ray, hits, 2 * VOLUME_STACK_SIZE, PATH_RAY_ALL_VISIBILITY);
+      kg, &volume_ray, hits, 2 * volume_stack_size, PATH_RAY_ALL_VISIBILITY);
   if (num_hits > 0) {
     Intersection *isect = hits;
 
@@ -55,7 +58,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(INTEGRATOR_STATE_A
 #else
   Intersection isect;
   int step = 0;
-  while (step < 2 * VOLUME_STACK_SIZE &&
+  while (step < 2 * volume_stack_size &&
          scene_intersect_volume(kg, &volume_ray, &isect, PATH_RAY_ALL_VISIBILITY)) {
     shader_setup_from_ray(kg, stack_sd, &volume_ray, &isect);
     volume_stack_enter_exit(INTEGRATOR_STATE_PASS, stack_sd);
@@ -91,12 +94,15 @@ ccl_device void integrator_intersect_volume_stack(INTEGRATOR_STATE_ARGS)
     stack_index++;
   }
 
+  /* Store to avoid global fetches on every intersection step. */
+  const uint volume_stack_size = kernel_data.volume_stack_size;
+
 #ifdef __VOLUME_RECORD_ALL__
-  Intersection hits[2 * VOLUME_STACK_SIZE + 1];
+  Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1];
   uint num_hits = scene_intersect_volume_all(
-      kg, &volume_ray, hits, 2 * VOLUME_STACK_SIZE, visibility);
+      kg, &volume_ray, hits, 2 * volume_stack_size, visibility);
   if (num_hits > 0) {
-    int enclosed_volumes[VOLUME_STACK_SIZE];
+    int enclosed_volumes[MAX_VOLUME_STACK_SIZE];
     Intersection *isect = hits;
 
     qsort(hits, num_hits, sizeof(Intersection), intersections_compare);
@@ -121,7 +127,7 @@ ccl_device void integrator_intersect_volume_stack(INTEGRATOR_STATE_ARGS)
             break;
           }
         }
-        if (need_add && stack_index < VOLUME_STACK_SIZE - 1) {
+        if (need_add && stack_index < volume_stack_size - 1) {
           const VolumeStack new_entry = {stack_sd->object, stack_sd->shader};
           integrator_state_write_volume_stack(INTEGRATOR_STATE_PASS, stack_index, new_entry);
           ++stack_index;
@@ -136,11 +142,12 @@ ccl_device void integrator_intersect_volume_stack(INTEGRATOR_STATE_ARGS)
     }
   }
 #else
-  int enclosed_volumes[VOLUME_STACK_SIZE];
+  /* CUDA does not support defintion of a variable size arrays, so use the maximum possible. */
+  int enclosed_volumes[MAX_VOLUME_STACK_SIZE];
   int step = 0;
 
-  while (stack_index < VOLUME_STACK_SIZE - 1 && enclosed_index < VOLUME_STACK_SIZE - 1 &&
-         step < 2 * VOLUME_STACK_SIZE) {
+  while (stack_index < volume_stack_size - 1 && enclosed_index < volume_stack_size - 1 &&
+         step < 2 * volume_stack_size) {
     Intersection isect;
     if (!scene_intersect_volume(kg, &volume_ray, &isect, visibility)) {
       break;

intern/cycles/kernel/integrator/integrator_shade_background.h

@@ -192,7 +192,8 @@ ccl_device void integrator_shade_background(INTEGRATOR_STATE_ARGS,
     INTEGRATOR_STATE_WRITE(path, flag) &= ~PATH_RAY_SHADOW_CATCHER_BACKGROUND;
 
     const int isect_prim = INTEGRATOR_STATE(isect, prim);
-    const int shader = intersection_get_shader_from_isect_prim(kg, isect_prim);
+    const int isect_type = INTEGRATOR_STATE(isect, type);
+    const int shader = intersection_get_shader_from_isect_prim(kg, isect_prim, isect_type);
     const int shader_flags = kernel_tex_fetch(__shaders, shader).flags;
 
     if ((shader_flags & SD_HAS_RAYTRACE) || (kernel_data.film.pass_ao != PASS_UNUSED)) {

intern/cycles/kernel/integrator/integrator_state.h

@@ -59,8 +59,6 @@ CCL_NAMESPACE_BEGIN
  *
  * TODO: these could be made dynamic depending on the features used in the scene. */
 
-#define INTEGRATOR_VOLUME_STACK_SIZE VOLUME_STACK_SIZE
-
 #define INTEGRATOR_SHADOW_ISECT_SIZE_CPU 1024
 #define INTEGRATOR_SHADOW_ISECT_SIZE_GPU 4
 
@@ -85,12 +83,14 @@ typedef struct IntegratorStateCPU {
 #define KERNEL_STRUCT_END_ARRAY(name, cpu_size, gpu_size) \
   } \
   name[cpu_size];
+#define KERNEL_STRUCT_VOLUME_STACK_SIZE MAX_VOLUME_STACK_SIZE
 #include "kernel/integrator/integrator_state_template.h"
 #undef KERNEL_STRUCT_BEGIN
 #undef KERNEL_STRUCT_MEMBER
 #undef KERNEL_STRUCT_ARRAY_MEMBER
 #undef KERNEL_STRUCT_END
 #undef KERNEL_STRUCT_END_ARRAY
+#undef KERNEL_STRUCT_VOLUME_STACK_SIZE
 } IntegratorStateCPU;
 
 /* Path Queue
@@ -114,12 +114,14 @@ typedef struct IntegratorStateGPU {
 #define KERNEL_STRUCT_END_ARRAY(name, cpu_size, gpu_size) \
   } \
   name[gpu_size];
+#define KERNEL_STRUCT_VOLUME_STACK_SIZE MAX_VOLUME_STACK_SIZE
 #include "kernel/integrator/integrator_state_template.h"
 #undef KERNEL_STRUCT_BEGIN
 #undef KERNEL_STRUCT_MEMBER
 #undef KERNEL_STRUCT_ARRAY_MEMBER
 #undef KERNEL_STRUCT_END
 #undef KERNEL_STRUCT_END_ARRAY
+#undef KERNEL_STRUCT_VOLUME_STACK_SIZE
 
   /* Count number of queued kernels. */
   IntegratorQueueCounter *queue_counter;

intern/cycles/kernel/integrator/integrator_state_template.h

@@ -107,7 +107,9 @@ KERNEL_STRUCT_END(subsurface)
 KERNEL_STRUCT_BEGIN(volume_stack)
 KERNEL_STRUCT_ARRAY_MEMBER(volume_stack, int, object, KERNEL_FEATURE_VOLUME)
 KERNEL_STRUCT_ARRAY_MEMBER(volume_stack, int, shader, KERNEL_FEATURE_VOLUME)
-KERNEL_STRUCT_END_ARRAY(volume_stack, INTEGRATOR_VOLUME_STACK_SIZE, INTEGRATOR_VOLUME_STACK_SIZE)
+KERNEL_STRUCT_END_ARRAY(volume_stack,
+                        KERNEL_STRUCT_VOLUME_STACK_SIZE,
+                        KERNEL_STRUCT_VOLUME_STACK_SIZE)
 
 /********************************* Shadow Path State **************************/
 
@@ -163,5 +165,5 @@ KERNEL_STRUCT_BEGIN(shadow_volume_stack)
 KERNEL_STRUCT_ARRAY_MEMBER(shadow_volume_stack, int, object, KERNEL_FEATURE_VOLUME)
 KERNEL_STRUCT_ARRAY_MEMBER(shadow_volume_stack, int, shader, KERNEL_FEATURE_VOLUME)
 KERNEL_STRUCT_END_ARRAY(shadow_volume_stack,
-                        INTEGRATOR_VOLUME_STACK_SIZE,
-                        INTEGRATOR_VOLUME_STACK_SIZE)
+                        KERNEL_STRUCT_VOLUME_STACK_SIZE,
+                        KERNEL_STRUCT_VOLUME_STACK_SIZE)

intern/cycles/kernel/integrator/integrator_state_util.h

@@ -155,7 +155,7 @@ ccl_device_forceinline void integrator_state_read_shadow_isect(INTEGRATOR_STATE_
 ccl_device_forceinline void integrator_state_copy_volume_stack_to_shadow(INTEGRATOR_STATE_ARGS)
 {
   if (kernel_data.kernel_features & KERNEL_FEATURE_VOLUME) {
-    for (int i = 0; i < INTEGRATOR_VOLUME_STACK_SIZE; i++) {
+    for (int i = 0; i < kernel_data.volume_stack_size; i++) {
       INTEGRATOR_STATE_ARRAY_WRITE(shadow_volume_stack, i, object) = INTEGRATOR_STATE_ARRAY(
           volume_stack, i, object);
       INTEGRATOR_STATE_ARRAY_WRITE(shadow_volume_stack, i, shader) = INTEGRATOR_STATE_ARRAY(
@@ -223,6 +223,8 @@ ccl_device_inline void integrator_state_copy_only(const IntegratorState to_state
     while (index < gpu_array_size) \
       ;
 
+#  define KERNEL_STRUCT_VOLUME_STACK_SIZE kernel_data.volume_stack_size
+
 #  include "kernel/integrator/integrator_state_template.h"
 
 #  undef KERNEL_STRUCT_BEGIN
@@ -230,6 +232,7 @@ ccl_device_inline void integrator_state_copy_only(const IntegratorState to_state
 #  undef KERNEL_STRUCT_ARRAY_MEMBER
 #  undef KERNEL_STRUCT_END
 #  undef KERNEL_STRUCT_END_ARRAY
+#  undef KERNEL_STRUCT_VOLUME_STACK_SIZE
 }
 
 ccl_device_inline void integrator_state_move(const IntegratorState to_state,

intern/cycles/kernel/integrator/integrator_subsurface.h

@@ -577,7 +577,7 @@ ccl_device_inline bool subsurface_scatter(INTEGRATOR_STATE_ARGS)
 #  ifdef __VOLUME__
   /* Update volume stack if needed. */
   if (kernel_data.integrator.use_volumes) {
-    const int object = intersection_get_object(kg, &ss_isect.hits[0]);
+    const int object = ss_isect.hits[0].object;
     const int object_flag = kernel_tex_fetch(__object_flag, object);
 
     if (object_flag & SD_OBJECT_INTERSECTS_VOLUME) {

intern/cycles/kernel/integrator/integrator_volume_stack.h

@@ -72,7 +72,7 @@ ccl_device void volume_stack_enter_exit(INTEGRATOR_STATE_ARGS,
     }
 
     /* If we exceed the stack limit, ignore. */
-    if (i >= VOLUME_STACK_SIZE - 1) {
+    if (i >= kernel_data.volume_stack_size - 1) {
       return;
     }
 

intern/cycles/kernel/kernel_shader.h

@@ -780,8 +780,8 @@ ccl_device_inline void shader_eval_volume(INTEGRATOR_STATE_CONST_ARGS,
       break;
     }
 
-    /* setup shaderdata from stack. it's mostly setup already in
-     * shader_setup_from_volume, this switching should be quick */
+    /* Setup shader-data from stack. it's mostly setup already in
+     * shader_setup_from_volume, this switching should be quick. */
     sd->object = entry.object;
     sd->lamp = LAMP_NONE;
     sd->shader = entry.shader;

intern/cycles/kernel/kernel_textures.h

@@ -18,11 +18,9 @@
 #  define KERNEL_TEX(type, name)
 #endif
 
-/* bvh */
+/* BVH2, not used for OptiX or Embree. */
 KERNEL_TEX(float4, __bvh_nodes)
 KERNEL_TEX(float4, __bvh_leaf_nodes)
-KERNEL_TEX(float4, __prim_tri_verts)
-KERNEL_TEX(uint, __prim_tri_index)
 KERNEL_TEX(uint, __prim_type)
 KERNEL_TEX(uint, __prim_visibility)
 KERNEL_TEX(uint, __prim_index)
@@ -46,10 +44,12 @@ KERNEL_TEX(float4, __tri_vnormal)
 KERNEL_TEX(uint4, __tri_vindex)
 KERNEL_TEX(uint, __tri_patch)
 KERNEL_TEX(float2, __tri_patch_uv)
+KERNEL_TEX(float4, __tri_verts)
 
 /* curves */
-KERNEL_TEX(float4, __curves)
+KERNEL_TEX(KernelCurve, __curves)
 KERNEL_TEX(float4, __curve_keys)
+KERNEL_TEX(KernelCurveSegment, __curve_segments)
 
 /* patches */
 KERNEL_TEX(uint, __patches)

intern/cycles/kernel/kernel_types.h

@@ -61,8 +61,6 @@ CCL_NAMESPACE_BEGIN
 #define ID_NONE (0.0f)
 #define PASS_UNUSED (~0)
 
-#define VOLUME_STACK_SIZE 4
-
 /* Kernel features */
 #define __SOBOL__
 #define __DPDU__
@@ -190,7 +188,7 @@ enum SamplingPattern {
   SAMPLING_NUM_PATTERNS,
 };
 
-/* these flags values correspond to raytypes in osl.cpp, so keep them in sync! */
+/* These flags values correspond to `raytypes` in `osl.cpp`, so keep them in sync! */
 
 enum PathRayFlag {
   /* --------------------------------------------------------------------
@@ -608,6 +606,12 @@ typedef struct AttributeDescriptor {
 #  define MAX_CLOSURE __MAX_CLOSURE__
 #endif
 
+#ifndef __MAX_VOLUME_STACK_SIZE__
+#  define MAX_VOLUME_STACK_SIZE 32
+#else
+#  define MAX_VOLUME_STACK_SIZE __MAX_VOLUME_STACK_SIZE__
+#endif
+
 #define MAX_VOLUME_CLOSURE 8
 
 /* This struct is the base class for all closures. The common members are
@@ -862,9 +866,6 @@ typedef struct VolumeStack {
 
 /* Struct to gather multiple nearby intersections. */
 typedef struct LocalIntersection {
-  Ray ray;
-  float3 weight[LOCAL_MAX_HITS];
-
   int num_hits;
   struct Intersection hits[LOCAL_MAX_HITS];
   float3 Ng[LOCAL_MAX_HITS];
@@ -1223,7 +1224,7 @@ typedef struct KernelData {
   uint kernel_features;
   uint max_closures;
   uint max_shaders;
-  uint pad;
+  uint volume_stack_size;
 
   KernelCamera cam;
   KernelFilm film;
@@ -1266,10 +1267,25 @@ typedef struct KernelObject {
 
   float ao_distance;
 
-  float pad1, pad2;
+  uint visibility;
+  int primitive_type;
 } KernelObject;
 static_assert_align(KernelObject, 16);
 
+typedef struct KernelCurve {
+  int shader_id;
+  int first_key;
+  int num_keys;
+  int type;
+} KernelCurve;
+static_assert_align(KernelCurve, 16);
+
+typedef struct KernelCurveSegment {
+  int prim;
+  int type;
+} KernelCurveSegment;
+static_assert_align(KernelCurveSegment, 8);
+
 typedef struct KernelSpotLight {
   float radius;
   float invarea;

intern/cycles/kernel/svm/svm_bevel.h

@@ -182,17 +182,17 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
     float3 disk_P = (disk_r * cosf(phi)) * disk_T + (disk_r * sinf(phi)) * disk_B;
 
     /* Create ray. */
-    Ray *ray = &isect.ray;
-    ray->P = sd->P + disk_N * disk_height + disk_P;
-    ray->D = -disk_N;
-    ray->t = 2.0f * disk_height;
-    ray->dP = differential_zero_compact();
-    ray->dD = differential_zero_compact();
-    ray->time = sd->time;
+    Ray ray ccl_optional_struct_init;
+    ray.P = sd->P + disk_N * disk_height + disk_P;
+    ray.D = -disk_N;
+    ray.t = 2.0f * disk_height;
+    ray.dP = differential_zero_compact();
+    ray.dD = differential_zero_compact();
+    ray.time = sd->time;
 
     /* Intersect with the same object. if multiple intersections are found it
      * will use at most LOCAL_MAX_HITS hits, a random subset of all hits. */
-    scene_intersect_local(kg, ray, &isect, sd->object, &lcg_state, LOCAL_MAX_HITS);
+    scene_intersect_local(kg, &ray, &isect, sd->object, &lcg_state, LOCAL_MAX_HITS);
 
     int num_eval_hits = min(isect.num_hits, LOCAL_MAX_HITS);
 
@@ -201,23 +201,20 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
       float3 hit_P;
       if (sd->type & PRIMITIVE_TRIANGLE) {
         hit_P = triangle_refine_local(
-            kg, sd, ray->P, ray->D, ray->t, isect.hits[hit].object, isect.hits[hit].prim);
+            kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim);
       }
 #  ifdef __OBJECT_MOTION__
       else if (sd->type & PRIMITIVE_MOTION_TRIANGLE) {
         float3 verts[3];
-        motion_triangle_vertices(
-            kg, sd->object, kernel_tex_fetch(__prim_index, isect.hits[hit].prim), sd->time, verts);
+        motion_triangle_vertices(kg, sd->object, isect.hits[hit].prim, sd->time, verts);
         hit_P = motion_triangle_refine_local(
-            kg, sd, ray->P, ray->D, ray->t, isect.hits[hit].object, isect.hits[hit].prim, verts);
+            kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim, verts);
       }
 #  endif /* __OBJECT_MOTION__ */
 
       /* Get geometric normal. */
       float3 hit_Ng = isect.Ng[hit];
-      int object = (isect.hits[hit].object == OBJECT_NONE) ?
-                       kernel_tex_fetch(__prim_object, isect.hits[hit].prim) :
-                       isect.hits[hit].object;
+      int object = isect.hits[hit].object;
       int object_flag = kernel_tex_fetch(__object_flag, object);
       if (object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
         hit_Ng = -hit_Ng;
@@ -225,7 +222,7 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
 
       /* Compute smooth normal. */
       float3 N = hit_Ng;
-      int prim = kernel_tex_fetch(__prim_index, isect.hits[hit].prim);
+      int prim = isect.hits[hit].prim;
       int shader = kernel_tex_fetch(__tri_shader, prim);
 
       if (shader & SHADER_SMOOTH_NORMAL) {

intern/cycles/kernel/svm/svm_closure.h

@@ -85,6 +85,9 @@ ccl_device_noinline int svm_node_closure_bsdf(
   }
 
   float3 N = stack_valid(data_node.x) ? stack_load_float3(stack, data_node.x) : sd->N;
+  if (!(sd->type & PRIMITIVE_ALL_CURVE)) {
+    N = ensure_valid_reflection(sd->Ng, sd->I, N);
+  }
 
   float param1 = (stack_valid(param1_offset)) ? stack_load_float(stack, param1_offset) :
                                                 __uint_as_float(node.z);
@@ -166,6 +169,9 @@ ccl_device_noinline int svm_node_closure_bsdf(
       float3 clearcoat_normal = stack_valid(data_cn_ssr.x) ?
                                     stack_load_float3(stack, data_cn_ssr.x) :
                                     sd->N;
+      if (!(sd->type & PRIMITIVE_ALL_CURVE)) {
+        clearcoat_normal = ensure_valid_reflection(sd->Ng, sd->I, clearcoat_normal);
+      }
       float3 subsurface_radius = stack_valid(data_cn_ssr.y) ?
                                      stack_load_float3(stack, data_cn_ssr.y) :
                                      make_float3(1.0f, 1.0f, 1.0f);

intern/cycles/kernel/svm/svm_math_util.h

@@ -200,43 +200,43 @@ ccl_device float svm_math(NodeMathType type, float a, float b, float c)
   }
 }
 
-/* Calculate color in range 800..12000 using an approximation
- * a/x+bx+c for R and G and ((at + b)t + c)t + d) for B
- * Max absolute error for RGB is (0.00095, 0.00077, 0.00057),
- * which is enough to get the same 8 bit/channel color.
- */
-
-ccl_static_constant float blackbody_table_r[6][3] = {
-    {2.52432244e+03f, -1.06185848e-03f, 3.11067539e+00f},
-    {3.37763626e+03f, -4.34581697e-04f, 1.64843306e+00f},
-    {4.10671449e+03f, -8.61949938e-05f, 6.41423749e-01f},
-    {4.66849800e+03f, 2.85655028e-05f, 1.29075375e-01f},
-    {4.60124770e+03f, 2.89727618e-05f, 1.48001316e-01f},
-    {3.78765709e+03f, 9.36026367e-06f, 3.98995841e-01f},
-};
-
-ccl_static_constant float blackbody_table_g[6][3] = {
-    {-7.50343014e+02f, 3.15679613e-04f, 4.73464526e-01f},
-    {-1.00402363e+03f, 1.29189794e-04f, 9.08181524e-01f},
-    {-1.22075471e+03f, 2.56245413e-05f, 1.20753416e+00f},
-    {-1.42546105e+03f, -4.01730887e-05f, 1.44002695e+00f},
-    {-1.18134453e+03f, -2.18913373e-05f, 1.30656109e+00f},
-    {-5.00279505e+02f, -4.59745390e-06f, 1.09090465e+00f},
-};
-
-ccl_static_constant float blackbody_table_b[6][4] = {
-    {0.0f, 0.0f, 0.0f, 0.0f}, /* zeros should be optimized by compiler */
-    {0.0f, 0.0f, 0.0f, 0.0f},
-    {0.0f, 0.0f, 0.0f, 0.0f},
-    {-2.02524603e-11f, 1.79435860e-07f, -2.60561875e-04f, -1.41761141e-02f},
-    {-2.22463426e-13f, -1.55078698e-08f, 3.81675160e-04f, -7.30646033e-01f},
-    {6.72595954e-13f, -2.73059993e-08f, 4.24068546e-04f, -7.52204323e-01f},
-};
-
 ccl_device float3 svm_math_blackbody_color(float t)
 {
   /* TODO(lukas): Reimplement in XYZ. */
 
+  /* Calculate color in range 800..12000 using an approximation
+   * a/x+bx+c for R and G and ((at + b)t + c)t + d) for B
+   * Max absolute error for RGB is (0.00095, 0.00077, 0.00057),
+   * which is enough to get the same 8 bit/channel color.
+   */
+
+  const float blackbody_table_r[6][3] = {
+      {2.52432244e+03f, -1.06185848e-03f, 3.11067539e+00f},
+      {3.37763626e+03f, -4.34581697e-04f, 1.64843306e+00f},
+      {4.10671449e+03f, -8.61949938e-05f, 6.41423749e-01f},
+      {4.66849800e+03f, 2.85655028e-05f, 1.29075375e-01f},
+      {4.60124770e+03f, 2.89727618e-05f, 1.48001316e-01f},
+      {3.78765709e+03f, 9.36026367e-06f, 3.98995841e-01f},
+  };
+
+  const float blackbody_table_g[6][3] = {
+      {-7.50343014e+02f, 3.15679613e-04f, 4.73464526e-01f},
+      {-1.00402363e+03f, 1.29189794e-04f, 9.08181524e-01f},
+      {-1.22075471e+03f, 2.56245413e-05f, 1.20753416e+00f},
+      {-1.42546105e+03f, -4.01730887e-05f, 1.44002695e+00f},
+      {-1.18134453e+03f, -2.18913373e-05f, 1.30656109e+00f},
+      {-5.00279505e+02f, -4.59745390e-06f, 1.09090465e+00f},
+  };
+
+  const float blackbody_table_b[6][4] = {
+      {0.0f, 0.0f, 0.0f, 0.0f}, /* zeros should be optimized by compiler */
+      {0.0f, 0.0f, 0.0f, 0.0f},
+      {0.0f, 0.0f, 0.0f, 0.0f},
+      {-2.02524603e-11f, 1.79435860e-07f, -2.60561875e-04f, -1.41761141e-02f},
+      {-2.22463426e-13f, -1.55078698e-08f, 3.81675160e-04f, -7.30646033e-01f},
+      {6.72595954e-13f, -2.73059993e-08f, 4.24068546e-04f, -7.52204323e-01f},
+  };
+
   if (t >= 12000.0f) {
     return make_float3(0.826270103f, 0.994478524f, 1.56626022f);
   }

intern/cycles/kernel/svm/svm_tex_coord.h

@@ -347,8 +347,6 @@ ccl_device_noinline void svm_node_normal_map(const KernelGlobals *kg,
     N = safe_normalize(sd->N + (N - sd->N) * strength);
   }
 
-  N = ensure_valid_reflection(sd->Ng, sd->I, N);
-
   if (is_zero(N)) {
     N = sd->N;
   }

intern/cycles/kernel/svm/svm_wavelength.h

@@ -34,44 +34,44 @@ CCL_NAMESPACE_BEGIN
 
 /* Wavelength to RGB */
 
-// CIE colour matching functions xBar, yBar, and zBar for
-//   wavelengths from 380 through 780 nanometers, every 5
-//   nanometers.  For a wavelength lambda in this range:
-//        cie_colour_match[(lambda - 380) / 5][0] = xBar
-//        cie_colour_match[(lambda - 380) / 5][1] = yBar
-//        cie_colour_match[(lambda - 380) / 5][2] = zBar
-ccl_static_constant float cie_colour_match[81][3] = {
-    {0.0014f, 0.0000f, 0.0065f}, {0.0022f, 0.0001f, 0.0105f}, {0.0042f, 0.0001f, 0.0201f},
-    {0.0076f, 0.0002f, 0.0362f}, {0.0143f, 0.0004f, 0.0679f}, {0.0232f, 0.0006f, 0.1102f},
-    {0.0435f, 0.0012f, 0.2074f}, {0.0776f, 0.0022f, 0.3713f}, {0.1344f, 0.0040f, 0.6456f},
-    {0.2148f, 0.0073f, 1.0391f}, {0.2839f, 0.0116f, 1.3856f}, {0.3285f, 0.0168f, 1.6230f},
-    {0.3483f, 0.0230f, 1.7471f}, {0.3481f, 0.0298f, 1.7826f}, {0.3362f, 0.0380f, 1.7721f},
-    {0.3187f, 0.0480f, 1.7441f}, {0.2908f, 0.0600f, 1.6692f}, {0.2511f, 0.0739f, 1.5281f},
-    {0.1954f, 0.0910f, 1.2876f}, {0.1421f, 0.1126f, 1.0419f}, {0.0956f, 0.1390f, 0.8130f},
-    {0.0580f, 0.1693f, 0.6162f}, {0.0320f, 0.2080f, 0.4652f}, {0.0147f, 0.2586f, 0.3533f},
-    {0.0049f, 0.3230f, 0.2720f}, {0.0024f, 0.4073f, 0.2123f}, {0.0093f, 0.5030f, 0.1582f},
-    {0.0291f, 0.6082f, 0.1117f}, {0.0633f, 0.7100f, 0.0782f}, {0.1096f, 0.7932f, 0.0573f},
-    {0.1655f, 0.8620f, 0.0422f}, {0.2257f, 0.9149f, 0.0298f}, {0.2904f, 0.9540f, 0.0203f},
-    {0.3597f, 0.9803f, 0.0134f}, {0.4334f, 0.9950f, 0.0087f}, {0.5121f, 1.0000f, 0.0057f},
-    {0.5945f, 0.9950f, 0.0039f}, {0.6784f, 0.9786f, 0.0027f}, {0.7621f, 0.9520f, 0.0021f},
-    {0.8425f, 0.9154f, 0.0018f}, {0.9163f, 0.8700f, 0.0017f}, {0.9786f, 0.8163f, 0.0014f},
-    {1.0263f, 0.7570f, 0.0011f}, {1.0567f, 0.6949f, 0.0010f}, {1.0622f, 0.6310f, 0.0008f},
-    {1.0456f, 0.5668f, 0.0006f}, {1.0026f, 0.5030f, 0.0003f}, {0.9384f, 0.4412f, 0.0002f},
-    {0.8544f, 0.3810f, 0.0002f}, {0.7514f, 0.3210f, 0.0001f}, {0.6424f, 0.2650f, 0.0000f},
-    {0.5419f, 0.2170f, 0.0000f}, {0.4479f, 0.1750f, 0.0000f}, {0.3608f, 0.1382f, 0.0000f},
-    {0.2835f, 0.1070f, 0.0000f}, {0.2187f, 0.0816f, 0.0000f}, {0.1649f, 0.0610f, 0.0000f},
-    {0.1212f, 0.0446f, 0.0000f}, {0.0874f, 0.0320f, 0.0000f}, {0.0636f, 0.0232f, 0.0000f},
-    {0.0468f, 0.0170f, 0.0000f}, {0.0329f, 0.0119f, 0.0000f}, {0.0227f, 0.0082f, 0.0000f},
-    {0.0158f, 0.0057f, 0.0000f}, {0.0114f, 0.0041f, 0.0000f}, {0.0081f, 0.0029f, 0.0000f},
-    {0.0058f, 0.0021f, 0.0000f}, {0.0041f, 0.0015f, 0.0000f}, {0.0029f, 0.0010f, 0.0000f},
-    {0.0020f, 0.0007f, 0.0000f}, {0.0014f, 0.0005f, 0.0000f}, {0.0010f, 0.0004f, 0.0000f},
-    {0.0007f, 0.0002f, 0.0000f}, {0.0005f, 0.0002f, 0.0000f}, {0.0003f, 0.0001f, 0.0000f},
-    {0.0002f, 0.0001f, 0.0000f}, {0.0002f, 0.0001f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f},
-    {0.0001f, 0.0000f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f}, {0.0000f, 0.0000f, 0.0000f}};
-
 ccl_device_noinline void svm_node_wavelength(
     const KernelGlobals *kg, ShaderData *sd, float *stack, uint wavelength, uint color_out)
 {
+  // CIE colour matching functions xBar, yBar, and zBar for
+  //   wavelengths from 380 through 780 nanometers, every 5
+  //   nanometers.  For a wavelength lambda in this range:
+  //        cie_colour_match[(lambda - 380) / 5][0] = xBar
+  //        cie_colour_match[(lambda - 380) / 5][1] = yBar
+  //        cie_colour_match[(lambda - 380) / 5][2] = zBar
+  const float cie_colour_match[81][3] = {
+      {0.0014f, 0.0000f, 0.0065f}, {0.0022f, 0.0001f, 0.0105f}, {0.0042f, 0.0001f, 0.0201f},
+      {0.0076f, 0.0002f, 0.0362f}, {0.0143f, 0.0004f, 0.0679f}, {0.0232f, 0.0006f, 0.1102f},
+      {0.0435f, 0.0012f, 0.2074f}, {0.0776f, 0.0022f, 0.3713f}, {0.1344f, 0.0040f, 0.6456f},
+      {0.2148f, 0.0073f, 1.0391f}, {0.2839f, 0.0116f, 1.3856f}, {0.3285f, 0.0168f, 1.6230f},
+      {0.3483f, 0.0230f, 1.7471f}, {0.3481f, 0.0298f, 1.7826f}, {0.3362f, 0.0380f, 1.7721f},
+      {0.3187f, 0.0480f, 1.7441f}, {0.2908f, 0.0600f, 1.6692f}, {0.2511f, 0.0739f, 1.5281f},
+      {0.1954f, 0.0910f, 1.2876f}, {0.1421f, 0.1126f, 1.0419f}, {0.0956f, 0.1390f, 0.8130f},
+      {0.0580f, 0.1693f, 0.6162f}, {0.0320f, 0.2080f, 0.4652f}, {0.0147f, 0.2586f, 0.3533f},
+      {0.0049f, 0.3230f, 0.2720f}, {0.0024f, 0.4073f, 0.2123f}, {0.0093f, 0.5030f, 0.1582f},
+      {0.0291f, 0.6082f, 0.1117f}, {0.0633f, 0.7100f, 0.0782f}, {0.1096f, 0.7932f, 0.0573f},
+      {0.1655f, 0.8620f, 0.0422f}, {0.2257f, 0.9149f, 0.0298f}, {0.2904f, 0.9540f, 0.0203f},
+      {0.3597f, 0.9803f, 0.0134f}, {0.4334f, 0.9950f, 0.0087f}, {0.5121f, 1.0000f, 0.0057f},
+      {0.5945f, 0.9950f, 0.0039f}, {0.6784f, 0.9786f, 0.0027f}, {0.7621f, 0.9520f, 0.0021f},
+      {0.8425f, 0.9154f, 0.0018f}, {0.9163f, 0.8700f, 0.0017f}, {0.9786f, 0.8163f, 0.0014f},
+      {1.0263f, 0.7570f, 0.0011f}, {1.0567f, 0.6949f, 0.0010f}, {1.0622f, 0.6310f, 0.0008f},
+      {1.0456f, 0.5668f, 0.0006f}, {1.0026f, 0.5030f, 0.0003f}, {0.9384f, 0.4412f, 0.0002f},
+      {0.8544f, 0.3810f, 0.0002f}, {0.7514f, 0.3210f, 0.0001f}, {0.6424f, 0.2650f, 0.0000f},
+      {0.5419f, 0.2170f, 0.0000f}, {0.4479f, 0.1750f, 0.0000f}, {0.3608f, 0.1382f, 0.0000f},
+      {0.2835f, 0.1070f, 0.0000f}, {0.2187f, 0.0816f, 0.0000f}, {0.1649f, 0.0610f, 0.0000f},
+      {0.1212f, 0.0446f, 0.0000f}, {0.0874f, 0.0320f, 0.0000f}, {0.0636f, 0.0232f, 0.0000f},
+      {0.0468f, 0.0170f, 0.0000f}, {0.0329f, 0.0119f, 0.0000f}, {0.0227f, 0.0082f, 0.0000f},
+      {0.0158f, 0.0057f, 0.0000f}, {0.0114f, 0.0041f, 0.0000f}, {0.0081f, 0.0029f, 0.0000f},
+      {0.0058f, 0.0021f, 0.0000f}, {0.0041f, 0.0015f, 0.0000f}, {0.0029f, 0.0010f, 0.0000f},
+      {0.0020f, 0.0007f, 0.0000f}, {0.0014f, 0.0005f, 0.0000f}, {0.0010f, 0.0004f, 0.0000f},
+      {0.0007f, 0.0002f, 0.0000f}, {0.0005f, 0.0002f, 0.0000f}, {0.0003f, 0.0001f, 0.0000f},
+      {0.0002f, 0.0001f, 0.0000f}, {0.0002f, 0.0001f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f},
+      {0.0001f, 0.0000f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f}, {0.0000f, 0.0000f, 0.0000f}};
+
   float lambda_nm = stack_load_float(stack, wavelength);
   float ii = (lambda_nm - 380.0f) * (1.0f / 5.0f);  // scaled 0..80
   int i = float_to_int(ii);

intern/cycles/render/buffers.cpp

@@ -97,6 +97,11 @@ NODE_DEFINE(BufferParams)
   SOCKET_INT(width, "Width", 0);
   SOCKET_INT(height, "Height", 0);
 
+  SOCKET_INT(window_x, "Window X", 0);
+  SOCKET_INT(window_y, "Window Y", 0);
+  SOCKET_INT(window_width, "Window Width", 0);
+  SOCKET_INT(window_height, "Window Height", 0);
+
   SOCKET_INT(full_x, "Full X", 0);
   SOCKET_INT(full_y, "Full Y", 0);
   SOCKET_INT(full_width, "Full Width", 0);
@@ -233,13 +238,31 @@ void BufferParams::update_offset_stride()
 
 bool BufferParams::modified(const BufferParams &other) const
 {
-  if (!(width == other.width && height == other.height && full_x == other.full_x &&
-        full_y == other.full_y && full_width == other.full_width &&
-        full_height == other.full_height && offset == other.offset && stride == other.stride &&
-        pass_stride == other.pass_stride && layer == other.layer && view == other.view &&
-        exposure == other.exposure &&
-        use_approximate_shadow_catcher == other.use_approximate_shadow_catcher &&
-        use_transparent_background == other.use_transparent_background)) {
+  if (width != other.width || height != other.height) {
+    return true;
+  }
+
+  if (full_x != other.full_x || full_y != other.full_y || full_width != other.full_width ||
+      full_height != other.full_height) {
+    return true;
+  }
+
+  if (window_x != other.window_x || window_y != other.window_y ||
+      window_width != other.window_width || window_height != other.window_height) {
+    return true;
+  }
+
+  if (offset != other.offset || stride != other.stride || pass_stride != other.pass_stride) {
+    return true;
+  }
+
+  if (layer != other.layer || view != other.view) {
+    return false;
+  }
+
+  if (exposure != other.exposure ||
+      use_approximate_shadow_catcher != other.use_approximate_shadow_catcher ||
+      use_transparent_background != other.use_transparent_background) {
     return true;
   }
 

intern/cycles/render/buffers.h

@@ -82,6 +82,15 @@ class BufferParams : public Node {
   int width = 0;
   int height = 0;
 
+  /* Windows defines which part of the buffers is visible. The part outside of the window is
+   * considered an "overscan".
+   *
+   * Window X and Y are relative to the position of the buffer in the full buffer. */
+  int window_x = 0;
+  int window_y = 0;
+  int window_width = 0;
+  int window_height = 0;
+
   /* Offset into and width/height of the full buffer. */
   int full_x = 0;
   int full_y = 0;

intern/cycles/render/geometry.cpp

@@ -46,12 +46,6 @@ CCL_NAMESPACE_BEGIN
 
 /* Geometry */
 
-PackFlags operator|=(PackFlags &pack_flags, uint32_t value)
-{
-  pack_flags = (PackFlags)((uint32_t)pack_flags | value);
-  return pack_flags;
-}
-
 NODE_ABSTRACT_DEFINE(Geometry)
 {
   NodeType *type = NodeType::add("geometry_base", NULL);
@@ -79,7 +73,6 @@ Geometry::Geometry(const NodeType *node_type, const Type type)
 
   bvh = NULL;
   attr_map_offset = 0;
-  optix_prim_offset = 0;
   prim_offset = 0;
 }
 
@@ -707,9 +700,9 @@ void GeometryManager::update_attribute_element_offset(Geometry *geom,
       if (element == ATTR_ELEMENT_CURVE)
         offset -= hair->prim_offset;
       else if (element == ATTR_ELEMENT_CURVE_KEY)
-        offset -= hair->curvekey_offset;
+        offset -= hair->curve_key_offset;
       else if (element == ATTR_ELEMENT_CURVE_KEY_MOTION)
-        offset -= hair->curvekey_offset;
+        offset -= hair->curve_key_offset;
     }
   }
   else {
@@ -972,28 +965,22 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout)
   size_t vert_size = 0;
   size_t tri_size = 0;
 
-  size_t curve_key_size = 0;
   size_t curve_size = 0;
+  size_t curve_key_size = 0;
+  size_t curve_segment_size = 0;
 
   size_t patch_size = 0;
   size_t face_size = 0;
   size_t corner_size = 0;
 
-  size_t optix_prim_size = 0;
-
   foreach (Geometry *geom, scene->geometry) {
-    if (geom->optix_prim_offset != optix_prim_size) {
-      /* Need to rebuild BVH in OptiX, since refit only allows modified mesh data there */
-      const bool has_optix_bvh = bvh_layout == BVH_LAYOUT_OPTIX ||
-                                 bvh_layout == BVH_LAYOUT_MULTI_OPTIX ||
-                                 bvh_layout == BVH_LAYOUT_MULTI_OPTIX_EMBREE;
-      geom->need_update_rebuild |= has_optix_bvh;
-      geom->need_update_bvh_for_offset = true;
-    }
+    bool prim_offset_changed = false;
 
     if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
       Mesh *mesh = static_cast<Mesh *>(geom);
 
+      prim_offset_changed = (mesh->prim_offset != tri_size);
+
       mesh->vert_offset = vert_size;
       mesh->prim_offset = tri_size;
 
@@ -1017,27 +1004,35 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout)
 
       face_size += mesh->get_num_subd_faces();
       corner_size += mesh->subd_face_corners.size();
-
-      mesh->optix_prim_offset = optix_prim_size;
-      optix_prim_size += mesh->num_triangles();
     }
     else if (geom->is_hair()) {
       Hair *hair = static_cast<Hair *>(geom);
 
-      hair->curvekey_offset = curve_key_size;
+      prim_offset_changed = (hair->curve_segment_offset != curve_segment_size);
+      hair->curve_key_offset = curve_key_size;
+      hair->curve_segment_offset = curve_segment_size;
       hair->prim_offset = curve_size;
 
-      curve_key_size += hair->get_curve_keys().size();
       curve_size += hair->num_curves();
+      curve_key_size += hair->get_curve_keys().size();
+      curve_segment_size += hair->num_segments();
+    }
 
-      hair->optix_prim_offset = optix_prim_size;
-      optix_prim_size += hair->num_segments();
+    if (prim_offset_changed) {
+      /* Need to rebuild BVH in OptiX, since refit only allows modified mesh data there */
+      const bool has_optix_bvh = bvh_layout == BVH_LAYOUT_OPTIX ||
+                                 bvh_layout == BVH_LAYOUT_MULTI_OPTIX ||
+                                 bvh_layout == BVH_LAYOUT_MULTI_OPTIX_EMBREE;
+      geom->need_update_rebuild |= has_optix_bvh;
+      geom->need_update_bvh_for_offset = true;
     }
   }
 }
 
-void GeometryManager::device_update_mesh(
-    Device *, DeviceScene *dscene, Scene *scene, bool for_displacement, Progress &progress)
+void GeometryManager::device_update_mesh(Device *,
+                                         DeviceScene *dscene,
+                                         Scene *scene,
+                                         Progress &progress)
 {
   /* Count. */
   size_t vert_size = 0;
@@ -1045,6 +1040,7 @@ void GeometryManager::device_update_mesh(
 
   size_t curve_key_size = 0;
   size_t curve_size = 0;
+  size_t curve_segment_size = 0;
 
   size_t patch_size = 0;
 
@@ -1071,31 +1067,7 @@ void GeometryManager::device_update_mesh(
 
       curve_key_size += hair->get_curve_keys().size();
       curve_size += hair->num_curves();
-    }
-  }
-
-  /* Create mapping from triangle to primitive triangle array. */
-  vector<uint> tri_prim_index(tri_size);
-  if (for_displacement) {
-    /* For displacement kernels we do some trickery to make them believe
-     * we've got all required data ready. However, that data is different
-     * from final render kernels since we don't have BVH yet, so can't
-     * really use same semantic of arrays.
-     */
-    foreach (Geometry *geom, scene->geometry) {
-      if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
-        Mesh *mesh = static_cast<Mesh *>(geom);
-        for (size_t i = 0; i < mesh->num_triangles(); ++i) {
-          tri_prim_index[i + mesh->prim_offset] = 3 * (i + mesh->prim_offset);
-        }
-      }
-    }
-  }
-  else {
-    for (size_t i = 0; i < dscene->prim_index.size(); ++i) {
-      if ((dscene->prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
-        tri_prim_index[dscene->prim_index[i]] = dscene->prim_tri_index[i];
-      }
+      curve_segment_size += hair->num_segments();
     }
   }
 
@@ -1104,6 +1076,7 @@ void GeometryManager::device_update_mesh(
     /* normals */
     progress.set_status("Updating Mesh", "Computing normals");
 
+    float4 *tri_verts = dscene->tri_verts.alloc(tri_size * 3);
     uint *tri_shader = dscene->tri_shader.alloc(tri_size);
     float4 *vnormal = dscene->tri_vnormal.alloc(vert_size);
     uint4 *tri_vindex = dscene->tri_vindex.alloc(tri_size);
@@ -1129,13 +1102,12 @@ void GeometryManager::device_update_mesh(
           mesh->pack_normals(&vnormal[mesh->vert_offset]);
         }
 
-        if (mesh->triangles_is_modified() || mesh->vert_patch_uv_is_modified() || copy_all_data) {
-          mesh->pack_verts(tri_prim_index,
+        if (mesh->verts_is_modified() || mesh->triangles_is_modified() ||
+            mesh->vert_patch_uv_is_modified() || copy_all_data) {
+          mesh->pack_verts(&tri_verts[mesh->prim_offset * 3],
                            &tri_vindex[mesh->prim_offset],
                            &tri_patch[mesh->prim_offset],
-                           &tri_patch_uv[mesh->vert_offset],
-                           mesh->vert_offset,
-                           mesh->prim_offset);
+                           &tri_patch_uv[mesh->vert_offset]);
         }
 
         if (progress.get_cancel())
@@ -1146,6 +1118,7 @@ void GeometryManager::device_update_mesh(
     /* 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();
@@ -1153,13 +1126,16 @@ void GeometryManager::device_update_mesh(
     dscene->tri_patch_uv.copy_to_device_if_modified();
   }
 
-  if (curve_size != 0) {
-    progress.set_status("Updating Mesh", "Copying Strands to device");
+  if (curve_segment_size != 0) {
+    progress.set_status("Updating Mesh", "Copying Curves to device");
 
     float4 *curve_keys = dscene->curve_keys.alloc(curve_key_size);
-    float4 *curves = dscene->curves.alloc(curve_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();
+    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()) {
@@ -1175,9 +1151,9 @@ void GeometryManager::device_update_mesh(
         }
 
         hair->pack_curves(scene,
-                          &curve_keys[hair->curvekey_offset],
+                          &curve_keys[hair->curve_key_offset],
                           &curves[hair->prim_offset],
-                          hair->curvekey_offset);
+                          &curve_segments[hair->curve_segment_offset]);
         if (progress.get_cancel())
           return;
       }
@@ -1185,6 +1161,7 @@ void GeometryManager::device_update_mesh(
 
     dscene->curve_keys.copy_to_device_if_modified();
     dscene->curves.copy_to_device_if_modified();
+    dscene->curve_segments.copy_to_device_if_modified();
   }
 
   if (patch_size != 0 && dscene->patches.need_realloc()) {
@@ -1195,10 +1172,7 @@ void GeometryManager::device_update_mesh(
     foreach (Geometry *geom, scene->geometry) {
       if (geom->is_mesh()) {
         Mesh *mesh = static_cast<Mesh *>(geom);
-        mesh->pack_patches(&patch_data[mesh->patch_offset],
-                           mesh->vert_offset,
-                           mesh->face_offset,
-                           mesh->corner_offset);
+        mesh->pack_patches(&patch_data[mesh->patch_offset]);
 
         if (mesh->patch_table) {
           mesh->patch_table->copy_adjusting_offsets(&patch_data[mesh->patch_table_offset],
@@ -1212,23 +1186,6 @@ void GeometryManager::device_update_mesh(
 
     dscene->patches.copy_to_device();
   }
-
-  if (for_displacement) {
-    float4 *prim_tri_verts = dscene->prim_tri_verts.alloc(tri_size * 3);
-    foreach (Geometry *geom, scene->geometry) {
-      if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
-        Mesh *mesh = static_cast<Mesh *>(geom);
-        for (size_t i = 0; i < mesh->num_triangles(); ++i) {
-          Mesh::Triangle t = mesh->get_triangle(i);
-          size_t offset = 3 * (i + mesh->prim_offset);
-          prim_tri_verts[offset + 0] = float3_to_float4(mesh->verts[t.v[0]]);
-          prim_tri_verts[offset + 1] = float3_to_float4(mesh->verts[t.v[1]]);
-          prim_tri_verts[offset + 2] = float3_to_float4(mesh->verts[t.v[2]]);
-        }
-      }
-    }
-    dscene->prim_tri_verts.copy_to_device();
-  }
 }
 
 void GeometryManager::device_update_bvh(Device *device,
@@ -1256,16 +1213,6 @@ void GeometryManager::device_update_bvh(Device *device,
   const bool can_refit = scene->bvh != nullptr &&
                          (bparams.bvh_layout == BVHLayout::BVH_LAYOUT_OPTIX);
 
-  PackFlags pack_flags = PackFlags::PACK_NONE;
-
-  if (scene->bvh == nullptr) {
-    pack_flags |= PackFlags::PACK_ALL;
-  }
-
-  if (dscene->prim_visibility.is_modified()) {
-    pack_flags |= PackFlags::PACK_VISIBILITY;
-  }
-
   BVH *bvh = scene->bvh;
   if (!scene->bvh) {
     bvh = scene->bvh = BVH::create(bparams, scene->geometry, scene->objects, device);
@@ -1284,77 +1231,7 @@ void GeometryManager::device_update_bvh(Device *device,
     pack = std::move(static_cast<BVH2 *>(bvh)->pack);
   }
   else {
-    progress.set_status("Updating Scene BVH", "Packing BVH primitives");
-
-    size_t num_prims = 0;
-    size_t num_tri_verts = 0;
-    foreach (Geometry *geom, scene->geometry) {
-      if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
-        Mesh *mesh = static_cast<Mesh *>(geom);
-        num_prims += mesh->num_triangles();
-        num_tri_verts += 3 * mesh->num_triangles();
-      }
-      else if (geom->is_hair()) {
-        Hair *hair = static_cast<Hair *>(geom);
-        num_prims += hair->num_segments();
-      }
-    }
-
     pack.root_index = -1;
-
-    if (pack_flags != PackFlags::PACK_ALL) {
-      /* if we do not need to recreate the BVH, then only the vertices are updated, so we can
-       * safely retake the memory */
-      dscene->prim_tri_verts.give_data(pack.prim_tri_verts);
-
-      if ((pack_flags & PackFlags::PACK_VISIBILITY) != 0) {
-        dscene->prim_visibility.give_data(pack.prim_visibility);
-      }
-    }
-    else {
-      /* It is not strictly necessary to skip those resizes we if do not have to repack, as the OS
-       * will not allocate pages if we do not touch them, however it does help catching bugs. */
-      pack.prim_tri_index.resize(num_prims);
-      pack.prim_tri_verts.resize(num_tri_verts);
-      pack.prim_type.resize(num_prims);
-      pack.prim_index.resize(num_prims);
-      pack.prim_object.resize(num_prims);
-      pack.prim_visibility.resize(num_prims);
-    }
-
-    // Merge visibility flags of all objects and find object index for non-instanced geometry
-    unordered_map<const Geometry *, pair<int, uint>> geometry_to_object_info;
-    geometry_to_object_info.reserve(scene->geometry.size());
-    foreach (Object *ob, scene->objects) {
-      const Geometry *const geom = ob->get_geometry();
-      pair<int, uint> &info = geometry_to_object_info[geom];
-      info.second |= ob->visibility_for_tracing();
-      if (!geom->is_instanced()) {
-        info.first = ob->get_device_index();
-      }
-    }
-
-    TaskPool pool;
-    // Iterate over scene mesh list instead of objects, since 'optix_prim_offset' was calculated
-    // based on that list, which may be ordered differently from the object list.
-    foreach (Geometry *geom, scene->geometry) {
-      /* Make a copy of the pack_flags so the current geometry's flags do not pollute the others'.
-       */
-      PackFlags geom_pack_flags = pack_flags;
-
-      if (geom->is_modified()) {
-        geom_pack_flags |= PackFlags::PACK_VERTICES;
-      }
-
-      if (geom_pack_flags == PACK_NONE) {
-        continue;
-      }
-
-      const pair<int, uint> &info = geometry_to_object_info[geom];
-      pool.push(function_bind(
-          &Geometry::pack_primitives, geom, &pack, info.first, info.second, geom_pack_flags));
-    }
-    pool.wait_work();
   }
 
   /* copy to device */
@@ -1375,31 +1252,23 @@ void GeometryManager::device_update_bvh(Device *device,
     dscene->object_node.steal_data(pack.object_node);
     dscene->object_node.copy_to_device();
   }
-  if (pack.prim_tri_index.size() && (dscene->prim_tri_index.need_realloc() || has_bvh2_layout)) {
-    dscene->prim_tri_index.steal_data(pack.prim_tri_index);
-    dscene->prim_tri_index.copy_to_device();
-  }
-  if (pack.prim_tri_verts.size()) {
-    dscene->prim_tri_verts.steal_data(pack.prim_tri_verts);
-    dscene->prim_tri_verts.copy_to_device();
-  }
-  if (pack.prim_type.size() && (dscene->prim_type.need_realloc() || has_bvh2_layout)) {
+  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.is_modified() || has_bvh2_layout)) {
+  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.need_realloc() || has_bvh2_layout)) {
+  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.need_realloc() || has_bvh2_layout)) {
+  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.need_realloc() || has_bvh2_layout)) {
+  if (pack.prim_time.size()) {
     dscene->prim_time.steal_data(pack.prim_time);
     dscene->prim_time.copy_to_device();
   }
@@ -1629,8 +1498,6 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
     dscene->bvh_nodes.tag_realloc();
     dscene->bvh_leaf_nodes.tag_realloc();
     dscene->object_node.tag_realloc();
-    dscene->prim_tri_verts.tag_realloc();
-    dscene->prim_tri_index.tag_realloc();
     dscene->prim_type.tag_realloc();
     dscene->prim_visibility.tag_realloc();
     dscene->prim_index.tag_realloc();
@@ -1649,6 +1516,7 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
     if (device_update_flags & DEVICE_CURVE_DATA_NEEDS_REALLOC) {
       dscene->curves.tag_realloc();
       dscene->curve_keys.tag_realloc();
+      dscene->curve_segments.tag_realloc();
     }
   }
 
@@ -1691,6 +1559,7 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
   if (device_update_flags & DEVICE_MESH_DATA_MODIFIED) {
     /* if anything else than vertices or shaders are modified, we would need to reallocate, so
      * these are the only arrays that can be updated */
+    dscene->tri_verts.tag_modified();
     dscene->tri_vnormal.tag_modified();
     dscene->tri_shader.tag_modified();
   }
@@ -1698,6 +1567,7 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
   if (device_update_flags & DEVICE_CURVE_DATA_MODIFIED) {
     dscene->curve_keys.tag_modified();
     dscene->curves.tag_modified();
+    dscene->curve_segments.tag_modified();
   }
 
   need_flags_update = false;
@@ -1906,7 +1776,7 @@ void GeometryManager::device_update(Device *device,
             {"device_update (displacement: copy meshes to device)", time});
       }
     });
-    device_update_mesh(device, dscene, scene, true, progress);
+    device_update_mesh(device, dscene, scene, progress);
   }
   if (progress.get_cancel()) {
     return;
@@ -2058,7 +1928,7 @@ void GeometryManager::device_update(Device *device,
             {"device_update (copy meshes to device)", time});
       }
     });
-    device_update_mesh(device, dscene, scene, false, progress);
+    device_update_mesh(device, dscene, scene, progress);
     if (progress.get_cancel()) {
       return;
     }
@@ -2091,13 +1961,12 @@ void GeometryManager::device_update(Device *device,
   dscene->bvh_nodes.clear_modified();
   dscene->bvh_leaf_nodes.clear_modified();
   dscene->object_node.clear_modified();
-  dscene->prim_tri_verts.clear_modified();
-  dscene->prim_tri_index.clear_modified();
   dscene->prim_type.clear_modified();
   dscene->prim_visibility.clear_modified();
   dscene->prim_index.clear_modified();
   dscene->prim_object.clear_modified();
   dscene->prim_time.clear_modified();
+  dscene->tri_verts.clear_modified();
   dscene->tri_shader.clear_modified();
   dscene->tri_vindex.clear_modified();
   dscene->tri_patch.clear_modified();
@@ -2105,6 +1974,7 @@ void GeometryManager::device_update(Device *device,
   dscene->tri_patch_uv.clear_modified();
   dscene->curves.clear_modified();
   dscene->curve_keys.clear_modified();
+  dscene->curve_segments.clear_modified();
   dscene->patches.clear_modified();
   dscene->attributes_map.clear_modified();
   dscene->attributes_float.clear_modified();
@@ -2118,13 +1988,12 @@ void GeometryManager::device_free(Device *device, DeviceScene *dscene, bool forc
   dscene->bvh_nodes.free_if_need_realloc(force_free);
   dscene->bvh_leaf_nodes.free_if_need_realloc(force_free);
   dscene->object_node.free_if_need_realloc(force_free);
-  dscene->prim_tri_verts.free_if_need_realloc(force_free);
-  dscene->prim_tri_index.free_if_need_realloc(force_free);
   dscene->prim_type.free_if_need_realloc(force_free);
   dscene->prim_visibility.free_if_need_realloc(force_free);
   dscene->prim_index.free_if_need_realloc(force_free);
   dscene->prim_object.free_if_need_realloc(force_free);
   dscene->prim_time.free_if_need_realloc(force_free);
+  dscene->tri_verts.free_if_need_realloc(force_free);
   dscene->tri_shader.free_if_need_realloc(force_free);
   dscene->tri_vnormal.free_if_need_realloc(force_free);
   dscene->tri_vindex.free_if_need_realloc(force_free);
@@ -2132,6 +2001,7 @@ void GeometryManager::device_free(Device *device, DeviceScene *dscene, bool forc
   dscene->tri_patch_uv.free_if_need_realloc(force_free);
   dscene->curves.free_if_need_realloc(force_free);
   dscene->curve_keys.free_if_need_realloc(force_free);
+  dscene->curve_segments.free_if_need_realloc(force_free);
   dscene->patches.free_if_need_realloc(force_free);
   dscene->attributes_map.free_if_need_realloc(force_free);
   dscene->attributes_float.free_if_need_realloc(force_free);

intern/cycles/render/geometry.h

@@ -43,24 +43,6 @@ class Shader;
 class Volume;
 struct PackedBVH;
 
-/* Flags used to determine which geometry data need to be packed. */
-enum PackFlags : uint32_t {
-  PACK_NONE = 0u,
-
-  /* Pack the geometry information (e.g. triangle or curve keys indices). */
-  PACK_GEOMETRY = (1u << 0),
-
-  /* Pack the vertices, for Meshes and Volumes' bounding meshes. */
-  PACK_VERTICES = (1u << 1),
-
-  /* Pack the visibility flags for each triangle or curve. */
-  PACK_VISIBILITY = (1u << 2),
-
-  PACK_ALL = (PACK_GEOMETRY | PACK_VERTICES | PACK_VISIBILITY),
-};
-
-PackFlags operator|=(PackFlags &pack_flags, uint32_t value);
-
 /* Geometry
  *
  * Base class for geometric types like Mesh and Hair. */
@@ -100,7 +82,6 @@ class Geometry : public Node {
   BVH *bvh;
   size_t attr_map_offset;
   size_t prim_offset;
-  size_t optix_prim_offset;
 
   /* Shader Properties */
   bool has_volume;         /* Set in the device_update_flags(). */
@@ -144,10 +125,7 @@ class Geometry : public Node {
                    int n,
                    int total);
 
-  virtual void pack_primitives(PackedBVH *pack,
-                               int object,
-                               uint visibility,
-                               PackFlags pack_flags) = 0;
+  virtual PrimitiveType primitive_type() const = 0;
 
   /* Check whether the geometry should have own BVH built separately. Briefly,
    * own BVH is needed for geometry, if:
@@ -260,11 +238,7 @@ class GeometryManager {
 
   void device_update_object(Device *device, DeviceScene *dscene, Scene *scene, Progress &progress);
 
-  void device_update_mesh(Device *device,
-                          DeviceScene *dscene,
-                          Scene *scene,
-                          bool for_displacement,
-                          Progress &progress);
+  void device_update_mesh(Device *device, DeviceScene *dscene, Scene *scene, Progress &progress);
 
   void device_update_attributes(Device *device,
                                 DeviceScene *dscene,

intern/cycles/render/graph.cpp

@@ -1149,7 +1149,9 @@ int ShaderGraph::get_num_closures()
       num_closures += 8;
     }
     else if (CLOSURE_IS_VOLUME(closure_type)) {
-      num_closures += VOLUME_STACK_SIZE;
+      /* TODO(sergey): Verify this is still needed, since we have special minimized volume storage
+       * for the volume steps. */
+      num_closures += MAX_VOLUME_STACK_SIZE;
     }
     else if (closure_type == CLOSURE_BSDF_HAIR_PRINCIPLED_ID) {
       num_closures += 4;

intern/cycles/render/hair.cpp

@@ -295,7 +295,8 @@ NODE_DEFINE(Hair)
 
 Hair::Hair() : Geometry(get_node_type(), Geometry::HAIR)
 {
-  curvekey_offset = 0;
+  curve_key_offset = 0;
+  curve_segment_offset = 0;
   curve_shape = CURVE_RIBBON;
 }
 
@@ -462,8 +463,8 @@ void Hair::apply_transform(const Transform &tfm, const bool apply_to_motion)
 
 void Hair::pack_curves(Scene *scene,
                        float4 *curve_key_co,
-                       float4 *curve_data,
-                       size_t curvekey_offset)
+                       KernelCurve *curves,
+                       KernelCurveSegment *curve_segments)
 {
   size_t curve_keys_size = curve_keys.size();
 
@@ -477,7 +478,10 @@ void Hair::pack_curves(Scene *scene,
   }
 
   /* pack curve segments */
+  const PrimitiveType type = primitive_type();
+
   size_t curve_num = num_curves();
+  size_t index = 0;
 
   for (size_t i = 0; i < curve_num; i++) {
     Curve curve = get_curve(i);
@@ -487,56 +491,24 @@ void Hair::pack_curves(Scene *scene,
                          scene->default_surface;
     shader_id = scene->shader_manager->get_shader_id(shader, false);
 
-    curve_data[i] = make_float4(__int_as_float(curve.first_key + curvekey_offset),
-                                __int_as_float(curve.num_keys),
-                                __int_as_float(shader_id),
-                                0.0f);
+    curves[i].shader_id = shader_id;
+    curves[i].first_key = curve_key_offset + curve.first_key;
+    curves[i].num_keys = curve.num_keys;
+    curves[i].type = type;
+
+    for (int k = 0; k < curve.num_segments(); ++k, ++index) {
+      curve_segments[index].prim = prim_offset + i;
+      curve_segments[index].type = PRIMITIVE_PACK_SEGMENT(type, k);
+    }
   }
 }
 
-void Hair::pack_primitives(PackedBVH *pack, int object, uint visibility, PackFlags pack_flags)
+PrimitiveType Hair::primitive_type() const
 {
-  if (curve_first_key.empty())
-    return;
-
-  /* Separate loop as other arrays are not initialized if their packing is not required. */
-  if ((pack_flags & PACK_VISIBILITY) != 0) {
-    unsigned int *prim_visibility = &pack->prim_visibility[optix_prim_offset];
-
-    size_t index = 0;
-    for (size_t j = 0; j < num_curves(); ++j) {
-      Curve curve = get_curve(j);
-      for (size_t k = 0; k < curve.num_segments(); ++k, ++index) {
-        prim_visibility[index] = visibility;
-      }
-    }
-  }
-
-  if ((pack_flags & PACK_GEOMETRY) != 0) {
-    unsigned int *prim_tri_index = &pack->prim_tri_index[optix_prim_offset];
-    int *prim_type = &pack->prim_type[optix_prim_offset];
-    int *prim_index = &pack->prim_index[optix_prim_offset];
-    int *prim_object = &pack->prim_object[optix_prim_offset];
-    // 'pack->prim_time' is unused by Embree and OptiX
-
-    uint type = has_motion_blur() ?
-                    ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON :
-                                                     PRIMITIVE_MOTION_CURVE_THICK) :
-                    ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON :
-                                                     PRIMITIVE_CURVE_THICK);
-
-    size_t index = 0;
-    for (size_t j = 0; j < num_curves(); ++j) {
-      Curve curve = get_curve(j);
-      for (size_t k = 0; k < curve.num_segments(); ++k, ++index) {
-        prim_tri_index[index] = -1;
-        prim_type[index] = PRIMITIVE_PACK_SEGMENT(type, k);
-        // Each curve segment points back to its curve index
-        prim_index[index] = j + prim_offset;
-        prim_object[index] = object;
-      }
-    }
-  }
+  return has_motion_blur() ?
+             ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON :
+                                              PRIMITIVE_MOTION_CURVE_THICK) :
+             ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON : PRIMITIVE_CURVE_THICK);
 }
 
 CCL_NAMESPACE_END

intern/cycles/render/hair.h

@@ -21,6 +21,8 @@
 
 CCL_NAMESPACE_BEGIN
 
+struct KernelCurveSegment;
+
 class Hair : public Geometry {
  public:
   NODE_DECLARE
@@ -95,7 +97,8 @@ class Hair : public Geometry {
   NODE_SOCKET_API_ARRAY(array<int>, curve_shader)
 
   /* BVH */
-  size_t curvekey_offset;
+  size_t curve_key_offset;
+  size_t curve_segment_offset;
   CurveShapeType curve_shape;
 
   /* Constructor/Destructor */
@@ -144,12 +147,12 @@ class Hair : public Geometry {
   void get_uv_tiles(ustring map, unordered_set<int> &tiles) override;
 
   /* BVH */
-  void pack_curves(Scene *scene, float4 *curve_key_co, float4 *curve_data, size_t curvekey_offset);
+  void pack_curves(Scene *scene,
+                   float4 *curve_key_co,
+                   KernelCurve *curve,
+                   KernelCurveSegment *curve_segments);
 
-  void pack_primitives(PackedBVH *pack,
-                       int object,
-                       uint visibility,
-                       PackFlags pack_flags) override;
+  PrimitiveType primitive_type() const override;
 };
 
 CCL_NAMESPACE_END

intern/cycles/render/mesh.cpp

@@ -729,12 +729,7 @@ void Mesh::pack_normals(float4 *vnormal)
   }
 }
 
-void Mesh::pack_verts(const vector<uint> &tri_prim_index,
-                      uint4 *tri_vindex,
-                      uint *tri_patch,
-                      float2 *tri_patch_uv,
-                      size_t vert_offset,
-                      size_t tri_offset)
+void Mesh::pack_verts(float4 *tri_verts, uint4 *tri_vindex, uint *tri_patch, float2 *tri_patch_uv)
 {
   size_t verts_size = verts.size();
 
@@ -749,17 +744,19 @@ void Mesh::pack_verts(const vector<uint> &tri_prim_index,
   size_t triangles_size = num_triangles();
 
   for (size_t i = 0; i < triangles_size; i++) {
-    Triangle t = get_triangle(i);
-    tri_vindex[i] = make_uint4(t.v[0] + vert_offset,
-                               t.v[1] + vert_offset,
-                               t.v[2] + vert_offset,
-                               tri_prim_index[i + tri_offset]);
+    const Triangle t = get_triangle(i);
+    tri_vindex[i] = make_uint4(
+        t.v[0] + vert_offset, t.v[1] + vert_offset, t.v[2] + vert_offset, 3 * (prim_offset + i));
 
     tri_patch[i] = (!get_num_subd_faces()) ? -1 : (triangle_patch[i] * 8 + patch_offset);
+
+    tri_verts[i * 3] = float3_to_float4(verts[t.v[0]]);
+    tri_verts[i * 3 + 1] = float3_to_float4(verts[t.v[1]]);
+    tri_verts[i * 3 + 2] = float3_to_float4(verts[t.v[2]]);
   }
 }
 
-void Mesh::pack_patches(uint *patch_data, uint vert_offset, uint face_offset, uint corner_offset)
+void Mesh::pack_patches(uint *patch_data)
 {
   size_t num_faces = get_num_subd_faces();
   int ngons = 0;
@@ -805,53 +802,9 @@ void Mesh::pack_patches(uint *patch_data, uint vert_offset, uint face_offset, ui
   }
 }
 
-void Mesh::pack_primitives(ccl::PackedBVH *pack, int object, uint visibility, PackFlags pack_flags)
+PrimitiveType Mesh::primitive_type() const
 {
-  if (triangles.empty())
-    return;
-
-  const size_t num_prims = num_triangles();
-
-  /* Use prim_offset for indexing as it is computed per geometry type, and prim_tri_verts does not
-   * contain data for Hair geometries. */
-  float4 *prim_tri_verts = &pack->prim_tri_verts[prim_offset * 3];
-  // 'pack->prim_time' is unused by Embree and OptiX
-
-  uint type = has_motion_blur() ? PRIMITIVE_MOTION_TRIANGLE : PRIMITIVE_TRIANGLE;
-
-  /* Separate loop as other arrays are not initialized if their packing is not required. */
-  if ((pack_flags & PackFlags::PACK_VISIBILITY) != 0) {
-    unsigned int *prim_visibility = &pack->prim_visibility[optix_prim_offset];
-    for (size_t k = 0; k < num_prims; ++k) {
-      prim_visibility[k] = visibility;
-    }
-  }
-
-  if ((pack_flags & PackFlags::PACK_GEOMETRY) != 0) {
-    /* Use optix_prim_offset for indexing as those arrays also contain data for Hair geometries. */
-    unsigned int *prim_tri_index = &pack->prim_tri_index[optix_prim_offset];
-    int *prim_type = &pack->prim_type[optix_prim_offset];
-    int *prim_index = &pack->prim_index[optix_prim_offset];
-    int *prim_object = &pack->prim_object[optix_prim_offset];
-
-    for (size_t k = 0; k < num_prims; ++k) {
-      if ((pack_flags & PackFlags::PACK_GEOMETRY) != 0) {
-        prim_tri_index[k] = (prim_offset + k) * 3;
-        prim_type[k] = type;
-        prim_index[k] = prim_offset + k;
-        prim_object[k] = object;
-      }
-    }
-  }
-
-  if ((pack_flags & PackFlags::PACK_VERTICES) != 0) {
-    for (size_t k = 0; k < num_prims; ++k) {
-      const Mesh::Triangle t = get_triangle(k);
-      prim_tri_verts[k * 3] = float3_to_float4(verts[t.v[0]]);
-      prim_tri_verts[k * 3 + 1] = float3_to_float4(verts[t.v[1]]);
-      prim_tri_verts[k * 3 + 2] = float3_to_float4(verts[t.v[2]]);
-    }
-  }
+  return has_motion_blur() ? PRIMITIVE_MOTION_TRIANGLE : PRIMITIVE_TRIANGLE;
 }
 
 CCL_NAMESPACE_END

intern/cycles/render/mesh.h

@@ -224,18 +224,10 @@ class Mesh : public Geometry {
 
   void pack_shaders(Scene *scene, uint *shader);
   void pack_normals(float4 *vnormal);
-  void pack_verts(const vector<uint> &tri_prim_index,
-                  uint4 *tri_vindex,
-                  uint *tri_patch,
-                  float2 *tri_patch_uv,
-                  size_t vert_offset,
-                  size_t tri_offset);
-  void pack_patches(uint *patch_data, uint vert_offset, uint face_offset, uint corner_offset);
+  void pack_verts(float4 *tri_verts, uint4 *tri_vindex, uint *tri_patch, float2 *tri_patch_uv);
+  void pack_patches(uint *patch_data);
 
-  void pack_primitives(PackedBVH *pack,
-                       int object,
-                       uint visibility,
-                       PackFlags pack_flags) override;
+  PrimitiveType primitive_type() const override;
 
   void tessellate(DiagSplit *split);
 

intern/cycles/render/object.cpp

@@ -60,6 +60,7 @@ struct UpdateObjectTransformState {
 
   /* Packed object arrays. Those will be filled in. */
   uint *object_flag;
+  uint *object_visibility;
   KernelObject *objects;
   Transform *object_motion_pass;
   DecomposedTransform *object_motion;
@@ -366,6 +367,22 @@ float Object::compute_volume_step_size() const
   return step_size;
 }
 
+bool Object::check_is_volume() const
+{
+  if (geometry->geometry_type == Geometry::VOLUME) {
+    return true;
+  }
+
+  for (Node *node : get_geometry()->get_used_shaders()) {
+    const Shader *shader = static_cast<const Shader *>(node);
+    if (shader->has_volume_connected) {
+      return true;
+    }
+  }
+
+  return false;
+}
+
 int Object::get_device_index() const
 {
   return index;
@@ -512,6 +529,9 @@ void ObjectManager::device_update_object_transform(UpdateObjectTransformState *s
                                              (1.0f - 0.5f * ob->shadow_terminator_shading_offset);
   kobject.shadow_terminator_geometry_offset = ob->shadow_terminator_geometry_offset;
 
+  kobject.visibility = ob->visibility_for_tracing();
+  kobject.primitive_type = geom->primitive_type();
+
   /* Object flag. */
   if (ob->use_holdout) {
     flag |= SD_OBJECT_HOLDOUT_MASK;

intern/cycles/render/object.h

@@ -109,6 +109,13 @@ class Object : public Node {
   /* Compute step size from attributes, shaders, transforms. */
   float compute_volume_step_size() const;
 
+  /* Check whether this object requires volume sampling (and hence might require space in the
+   * volume stack).
+   *
+   * Note that this is a naive iteration over sharders, which allows to access information prior
+   * to `scene_update()`. */
+  bool check_is_volume() const;
+
  protected:
   /* Specifies the position of the object in scene->objects and
    * in the device vectors. Gets set in device_update. */

intern/cycles/render/scene.cpp

@@ -49,13 +49,12 @@ 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_tri_index(device, "__prim_tri_index", MEM_GLOBAL),
-      prim_tri_verts(device, "__prim_tri_verts", 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),
@@ -63,6 +62,7 @@ DeviceScene::DeviceScene(Device *device)
       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),
       objects(device, "__objects", MEM_GLOBAL),
       object_motion_pass(device, "__object_motion_pass", MEM_GLOBAL),
@@ -527,6 +527,8 @@ void Scene::update_kernel_features()
   const uint max_closures = (params.background) ? get_max_closure_count() : MAX_CLOSURE;
   dscene.data.max_closures = max_closures;
   dscene.data.max_shaders = shaders.size();
+
+  dscene.data.volume_stack_size = get_volume_stack_size();
 }
 
 bool Scene::update(Progress &progress)
@@ -642,6 +644,33 @@ int Scene::get_max_closure_count()
   return max_closure_global;
 }
 
+int Scene::get_volume_stack_size() const
+{
+  /* Quick non-expensive check. Can over-estimate maximum possible nested level, but does not
+   * require expensive calculation during pre-processing. */
+  int num_volume_objects = 0;
+  for (const Object *object : objects) {
+    if (object->check_is_volume()) {
+      ++num_volume_objects;
+    }
+
+    if (num_volume_objects == MAX_VOLUME_STACK_SIZE) {
+      break;
+    }
+  }
+
+  /* Count background world for the stack. */
+  const Shader *background_shader = background->get_shader(this);
+  if (background_shader && background_shader->has_volume_connected) {
+    ++num_volume_objects;
+  }
+
+  /* Space for terminator. */
+  ++num_volume_objects;
+
+  return min(num_volume_objects, MAX_VOLUME_STACK_SIZE);
+}
+
 bool Scene::has_shadow_catcher()
 {
   if (shadow_catcher_modified_) {

intern/cycles/render/scene.h

@@ -74,8 +74,6 @@ class DeviceScene {
   device_vector<int4> bvh_nodes;
   device_vector<int4> bvh_leaf_nodes;
   device_vector<int> object_node;
-  device_vector<uint> prim_tri_index;
-  device_vector<float4> prim_tri_verts;
   device_vector<int> prim_type;
   device_vector<uint> prim_visibility;
   device_vector<int> prim_index;
@@ -83,14 +81,16 @@ class DeviceScene {
   device_vector<float2> prim_time;
 
   /* mesh */
+  device_vector<float4> tri_verts;
   device_vector<uint> tri_shader;
   device_vector<float4> tri_vnormal;
   device_vector<uint4> tri_vindex;
   device_vector<uint> tri_patch;
   device_vector<float2> tri_patch_uv;
 
-  device_vector<float4> curves;
+  device_vector<KernelCurve> curves;
   device_vector<float4> curve_keys;
+  device_vector<KernelCurveSegment> curve_segments;
 
   device_vector<uint> patches;
 
@@ -344,6 +344,9 @@ class Scene : public NodeOwner {
   /* Get maximum number of closures to be used in kernel. */
   int get_max_closure_count();
 
+  /* Get size of a volume stack needed to render this scene.  */
+  int get_volume_stack_size() const;
+
   template<typename T> void delete_node_impl(T *node)
   {
     delete node;

intern/cycles/render/session.cpp

@@ -157,6 +157,13 @@ void Session::run_main_render_loop()
       continue;
     }
 
+    /* Stop rendering if error happened during scene update or other step of preparing scene
+     * for render. */
+    if (device->have_error()) {
+      progress.set_error(device->error_message());
+      break;
+    }
+
     {
       /* buffers mutex is locked entirely while rendering each
        * sample, and released/reacquired on each iteration to allow
@@ -172,10 +179,9 @@ void Session::run_main_render_loop()
       /* update status and timing */
       update_status_time();
 
+      /* Stop rendering if error happened during path tracing. */
       if (device->have_error()) {
-        const string &error_message = device->error_message();
-        progress.set_error(error_message);
-        progress.set_cancel(error_message);
+        progress.set_error(device->error_message());
         break;
       }
     }
@@ -280,12 +286,20 @@ RenderWork Session::run_update_for_next_iteration()
       BufferParams tile_params = buffer_params_;
 
       const Tile &tile = tile_manager_.get_current_tile();
+
       tile_params.width = tile.width;
       tile_params.height = tile.height;
+
+      tile_params.window_x = tile.window_x;
+      tile_params.window_y = tile.window_y;
+      tile_params.window_width = tile.window_width;
+      tile_params.window_height = tile.window_height;
+
       tile_params.full_x = tile.x + buffer_params_.full_x;
       tile_params.full_y = tile.y + buffer_params_.full_y;
       tile_params.full_width = buffer_params_.full_width;
       tile_params.full_height = buffer_params_.full_height;
+
       tile_params.update_offset_stride();
 
       path_trace_->reset(buffer_params_, tile_params);

intern/cycles/render/tile.cpp

@@ -372,8 +372,17 @@ void TileManager::update(const BufferParams &params, const Scene *scene)
   configure_image_spec_from_buffer(&write_state_.image_spec, buffer_params_, tile_size_);
 
   const DenoiseParams denoise_params = scene->integrator->get_denoise_params();
+  const AdaptiveSampling adaptive_sampling = scene->integrator->get_adaptive_sampling();
+
   node_to_image_spec_atttributes(
       &write_state_.image_spec, &denoise_params, ATTR_DENOISE_SOCKET_PREFIX);
+
+  if (adaptive_sampling.use) {
+    overscan_ = 4;
+  }
+  else {
+    overscan_ = 0;
+  }
 }
 
 bool TileManager::done()
@@ -399,18 +408,25 @@ Tile TileManager::get_tile_for_index(int index) const
   /* TODO(sergey): Consider using hilbert spiral, or. maybe, even configurable. Not sure this
    * brings a lot of value since this is only applicable to BIG tiles. */
 
-  const int tile_y = index / tile_state_.num_tiles_x;
-  const int tile_x = index - tile_y * tile_state_.num_tiles_x;
+  const int tile_index_y = index / tile_state_.num_tiles_x;
+  const int tile_index_x = index - tile_index_y * tile_state_.num_tiles_x;
+
+  const int tile_window_x = tile_index_x * tile_size_.x;
+  const int tile_window_y = tile_index_y * tile_size_.y;
 
   Tile tile;
 
-  tile.x = tile_x * tile_size_.x;
-  tile.y = tile_y * tile_size_.y;
-  tile.width = tile_size_.x;
-  tile.height = tile_size_.y;
+  tile.x = max(0, tile_window_x - overscan_);
+  tile.y = max(0, tile_window_y - overscan_);
 
-  tile.width = min(tile.width, buffer_params_.width - tile.x);
-  tile.height = min(tile.height, buffer_params_.height - tile.y);
+  tile.window_x = tile_window_x - tile.x;
+  tile.window_y = tile_window_y - tile.y;
+  tile.window_width = min(tile_size_.x, buffer_params_.width - tile_window_x);
+  tile.window_height = min(tile_size_.y, buffer_params_.height - tile_window_y);
+
+  tile.width = min(buffer_params_.width - tile.x, tile.window_x + tile.window_width + overscan_);
+  tile.height = min(buffer_params_.height - tile.y,
+                    tile.window_y + tile.window_height + overscan_);
 
   return tile;
 }
@@ -483,11 +499,22 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
 
   DCHECK_EQ(tile_buffers.params.pass_stride, buffer_params_.pass_stride);
 
+  vector<float> pixel_storage;
+
   const BufferParams &tile_params = tile_buffers.params;
 
-  const float *pixels = tile_buffers.buffer.data();
-  const int tile_x = tile_params.full_x - buffer_params_.full_x;
-  const int tile_y = tile_params.full_y - buffer_params_.full_y;
+  const int tile_x = tile_params.full_x - buffer_params_.full_x + tile_params.window_x;
+  const int tile_y = tile_params.full_y - buffer_params_.full_y + tile_params.window_y;
+
+  const int64_t pass_stride = tile_params.pass_stride;
+  const int64_t tile_row_stride = tile_params.width * pass_stride;
+
+  const int64_t xstride = pass_stride * sizeof(float);
+  const int64_t ystride = xstride * tile_params.width;
+  const int64_t zstride = ystride * tile_params.height;
+
+  const float *pixels = tile_buffers.buffer.data() + tile_params.window_x * pass_stride +
+                        tile_params.window_y * tile_row_stride;
 
   VLOG(3) << "Write tile at " << tile_x << ", " << tile_y;
 
@@ -499,13 +526,16 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
    * The only thing we have to ensure is that the tile_x and tile_y are a multiple of the
    * image tile size, which happens in compute_render_tile_size. */
   if (!write_state_.tile_out->write_tiles(tile_x,
-                                          tile_x + tile_params.width,
+                                          tile_x + tile_params.window_width,
                                           tile_y,
-                                          tile_y + tile_params.height,
+                                          tile_y + tile_params.window_height,
                                           0,
                                           1,
                                           TypeDesc::FLOAT,
-                                          pixels)) {
+                                          pixels,
+                                          xstride,
+                                          ystride,
+                                          zstride)) {
     LOG(ERROR) << "Error writing tile " << write_state_.tile_out->geterror();
     return false;
   }
@@ -531,12 +561,15 @@ void TileManager::finish_write_tiles()
          ++tile_index) {
       const Tile tile = get_tile_for_index(tile_index);
 
-      VLOG(3) << "Write dummy tile at " << tile.x << ", " << tile.y;
+      const int tile_x = tile.x + tile.window_x;
+      const int tile_y = tile.y + tile.window_y;
 
-      write_state_.tile_out->write_tiles(tile.x,
-                                         tile.x + tile.width,
-                                         tile.y,
-                                         tile.y + tile.height,
+      VLOG(3) << "Write dummy tile at " << tile_x << ", " << tile_y;
+
+      write_state_.tile_out->write_tiles(tile_x,
+                                         tile_x + tile.window_width,
+                                         tile_y,
+                                         tile_y + tile.window_height,
                                          0,
                                          1,
                                          TypeDesc::FLOAT,

intern/cycles/render/tile.h

@@ -35,6 +35,9 @@ class Tile {
   int x = 0, y = 0;
   int width = 0, height = 0;
 
+  int window_x = 0, window_y = 0;
+  int window_width = 0, window_height = 0;
+
   Tile()
   {
   }
@@ -78,6 +81,11 @@ class TileManager {
     return tile_state_.num_tiles > 1;
   }
 
+  inline int get_tile_overscan() const
+  {
+    return overscan_;
+  }
+
   bool next();
   bool done();
 
@@ -128,6 +136,9 @@ class TileManager {
 
   int2 tile_size_ = make_int2(0, 0);
 
+  /* Number of extra pixels around the actual tile to render. */
+  int overscan_ = 0;
+
   BufferParams buffer_params_;
 
   /* Tile scheduling state. */

intern/ghost/intern/GHOST_SystemPathsUnix.cpp

@@ -142,7 +142,7 @@ const char *GHOST_SystemPathsUnix::getUserSpecialDir(GHOST_TUserSpecialDirTypes
   }
 
   static string path = "";
-  /* Pipe stderr to /dev/null to avoid error prints. We will fail gracefully still. */
+  /* Pipe `stderr` to `/dev/null` to avoid error prints. We will fail gracefully still. */
   string command = string("xdg-user-dir ") + type_str + " 2> /dev/null";
 
   FILE *fstream = popen(command.c_str(), "r");
@@ -152,7 +152,7 @@ const char *GHOST_SystemPathsUnix::getUserSpecialDir(GHOST_TUserSpecialDirTypes
   std::stringstream path_stream;
   while (!feof(fstream)) {
     char c = fgetc(fstream);
-    /* xdg-user-dir ends the path with '\n'. */
+    /* `xdg-user-dir` ends the path with '\n'. */
     if (c == '\n') {
       break;
     }

intern/ghost/intern/GHOST_SystemX11.cpp

@@ -1044,7 +1044,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
 
 #ifdef USE_NON_LATIN_KB_WORKAROUND
       /* XXX: Code below is kinda awfully convoluted... Issues are:
-       * - In keyboards like latin ones, numbers need a 'Shift' to be accessed but key_sym
+       * - In keyboards like Latin ones, numbers need a 'Shift' to be accessed but key_sym
        *   is unmodified (or anyone swapping the keys with `xmodmap`).
        * - #XLookupKeysym seems to always use first defined key-map (see T47228), which generates
        *   key-codes unusable by ghost_key_from_keysym for non-Latin-compatible key-maps.
@@ -1131,7 +1131,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
         }
       }
 #else
-      /* In keyboards like latin ones,
+      /* In keyboards like Latin ones,
        * numbers needs a 'Shift' to be accessed but key_sym
        * is unmodified (or anyone swapping the keys with xmodmap).
        *

release/scripts/modules/bl_keymap_utils/keymap_hierarchy.py

@@ -168,9 +168,13 @@ _km_hierarchy = [
     ('Node Editor', 'NODE_EDITOR', 'WINDOW', [
         ('Node Generic', 'NODE_EDITOR', 'WINDOW', []),
     ]),
-    ('Sequencer', 'SEQUENCE_EDITOR', 'WINDOW', [
-        ('SequencerCommon', 'SEQUENCE_EDITOR', 'WINDOW', []),
-        ('SequencerPreview', 'SEQUENCE_EDITOR', 'WINDOW', []),
+    ('SequencerCommon', 'SEQUENCE_EDITOR', 'WINDOW', [
+        ('Sequencer', 'SEQUENCE_EDITOR', 'WINDOW', [
+            _km_expand_from_toolsystem('SEQUENCE_EDITOR', 'SEQUENCER'),
+        ]),
+        ('SequencerPreview', 'SEQUENCE_EDITOR', 'WINDOW', [
+            _km_expand_from_toolsystem('SEQUENCE_EDITOR', 'PREVIEW'),
+        ]),
     ]),
 
     ('File Browser', 'FILE_BROWSER', 'WINDOW', [

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

@@ -75,6 +75,8 @@ class Params:
         "use_fallback_tool_rmb",
         # Shorthand for: `('CLICK' if params.use_fallback_tool_rmb else params.select_mouse_value)`.
         "select_mouse_value_fallback",
+        # Shorthand for: `{"type": params.select_tweak, "value": 'ANY'}`.
+        "select_tweak_event",
         # Shorthand for: `('CLICK_DRAG' if params.use_pie_click_drag else 'PRESS')`
         "pie_value",
         # Shorthand for: `{"type": params.tool_tweak, "value": 'ANY'}`.
@@ -197,6 +199,7 @@ class Params:
         # Convenience variables:
         self.use_fallback_tool_rmb = self.use_fallback_tool if select_mouse == 'RIGHT' else False
         self.select_mouse_value_fallback = 'CLICK' if self.use_fallback_tool_rmb else self.select_mouse_value
+        self.select_tweak_event = {"type": self.select_tweak, "value": 'ANY'}
         self.pie_value = 'CLICK_DRAG' if use_pie_click_drag else 'PRESS'
         self.tool_tweak_event = {"type": self.tool_tweak, "value": 'ANY'}
         self.tool_maybe_tweak_event = {"type": self.tool_maybe_tweak, "value": self.tool_maybe_tweak_value}
@@ -2682,9 +2685,6 @@ def km_sequencercommon(params):
         ("wm.context_toggle_enum", {"type": 'TAB', "value": 'PRESS', "ctrl": True},
          {"properties": [("data_path", 'space_data.view_type'), ("value_1", 'SEQUENCER'), ("value_2", 'PREVIEW')]}),
         ("sequencer.refresh_all", {"type": 'R', "value": 'PRESS', "ctrl": True}, None),
-        ("sequencer.select", {"type": params.select_mouse, "value": 'PRESS'}, None),
-        ("sequencer.select", {"type": params.select_mouse, "value": 'PRESS', "shift": True},
-         {"properties": [("extend", True)]}),
     ])
 
     if params.select_mouse == 'LEFTMOUSE' and not params.legacy:
@@ -2767,16 +2767,11 @@ def km_sequencer(params):
              for i in range(10)
              )
         ),
-        ("sequencer.select", {"type": params.select_mouse, "value": 'PRESS', "alt": True},
-         {"properties": [("linked_handle", True)]}),
-        ("sequencer.select", {"type": params.select_mouse, "value": 'PRESS', "shift": True, "alt": True},
-         {"properties": [("extend", True), ("linked_handle", True)]}),
-        ("sequencer.select",
-         {"type": params.select_mouse, "value": 'PRESS' if params.legacy else 'CLICK', "ctrl": True},
-         {"properties": [("side_of_frame", True), ("linked_time", True)]}),
-        ("sequencer.select",
-         {"type": params.select_mouse, "value": 'PRESS' if params.legacy else 'CLICK', "ctrl": True, "shift": True},
-         {"properties": [("side_of_frame", True), ("linked_time", True), ("extend", True)]}),
+        *_template_sequencer_timeline_select(
+            type=params.select_mouse,
+            value=params.select_mouse_value_fallback,
+            legacy=params.legacy,
+        ),
         ("sequencer.select_more", {"type": 'NUMPAD_PLUS', "value": 'PRESS', "ctrl": True, "repeat": True}, None),
         ("sequencer.select_less", {"type": 'NUMPAD_MINUS', "value": 'PRESS', "ctrl": True, "repeat": True}, None),
         ("sequencer.select_linked_pick", {"type": 'L', "value": 'PRESS'}, None),
@@ -2825,6 +2820,14 @@ def km_sequencerpreview(params):
     )
 
     items.extend([
+        # Selection.
+        *_template_sequencer_preview_select(
+            type=params.select_mouse,
+            value=params.select_mouse_value_fallback,
+            legacy=params.legacy,
+        ),
+        op_menu_pie("SEQUENCER_MT_pivot_pie", {"type": 'PERIOD', "value": 'PRESS'}),
+
         ("sequencer.view_all_preview", {"type": 'HOME', "value": 'PRESS'}, None),
         ("sequencer.view_all_preview", {"type": 'NDOF_BUTTON_FIT', "value": 'PRESS'}, None),
         ("sequencer.view_ghost_border", {"type": 'O', "value": 'PRESS'}, None),
@@ -2842,7 +2845,7 @@ def km_sequencerpreview(params):
          {"properties": [("ratio", 0.25)]}),
         ("sequencer.view_zoom_ratio", {"type": 'NUMPAD_8', "value": 'PRESS'},
          {"properties": [("ratio", 0.125)]}),
-        ("sequencer.sample", {"type": params.action_mouse, "value": 'PRESS'}, None),
+        ("transform.translate", {"type": params.select_tweak, "value": 'ANY'}, None),
         op_tool_optional(
             ("transform.translate", {"type": 'G', "value": 'PRESS'}, None),
             (op_tool_cycle, "builtin.move"), params),
@@ -2858,8 +2861,21 @@ def km_sequencerpreview(params):
          {"properties": [("property", 'SCALE')]}),
         ("sequencer.strip_transform_clear", {"type": 'R', "alt": True, "value": 'PRESS'},
          {"properties": [("property", 'ROTATION')]}),
+        *_template_items_context_menu("SEQUENCER_MT_preview_context_menu", params.context_menu_event),
     ])
 
+    # 2D cursor.
+    if params.cursor_tweak_event:
+        items.extend([
+            ("sequencer.cursor_set", params.cursor_set_event, None),
+            ("transform.translate", params.cursor_tweak_event,
+             {"properties": [("release_confirm", True), ("cursor_transform", True)]}),
+        ])
+    else:
+        items.extend([
+            ("sequencer.cursor_set", params.cursor_set_event, None),
+        ])
+
     return keymap
 
 
@@ -4634,6 +4650,62 @@ def _template_uv_select_for_fallback(params, fallback):
     return []
 
 
+def _template_sequencer_generic_select(*, type, value, legacy):
+    return [(
+        "sequencer.select",
+        {"type": type, "value": value, **{m: True for m in mods}},
+        {"properties": [(c, True) for c in props]},
+    ) for props, mods in (
+        (("deselect_all",) if not legacy else (), ()),
+        (("toggle",), ("shift",)),
+    )]
+
+
+def _template_sequencer_preview_select(*, type, value, legacy):
+    return _template_sequencer_generic_select(
+        type=type, value=value, legacy=legacy,
+    ) + [(
+        "sequencer.select",
+        {"type": type, "value": value, **{m: True for m in mods}},
+        {"properties": [(c, True) for c in props]},
+    ) for props, mods in (
+        (("center",), ("ctrl",)),
+        # TODO:
+        # (("enumerate",), ("alt",)),
+            (("toggle", "center"), ("shift", "ctrl")),
+        # (("center", "enumerate"), ("ctrl", "alt")),
+        # (("toggle", "enumerate"), ("shift", "alt")),
+        # (("toggle", "center", "enumerate"), ("shift", "ctrl", "alt")),
+    )]
+
+
+def _template_sequencer_timeline_select(*, type, value, legacy):
+    return _template_sequencer_generic_select(
+        type=type, value=value, legacy=legacy,
+    ) + [(
+        "sequencer.select",
+        {"type": type, "value": value, **{m: True for m in mods}},
+        {"properties": [(c, True) for c in props]},
+    ) for props, mods in (
+        (("linked_handle",), ("alt",)),
+        (("linked_handle", "extend"), ("shift", "alt",)),
+
+        (("side_of_frame", "linked_time"), ("ctrl",)),
+        (("side_of_frame", "linked_time", "extend"), ("ctrl", "shift")),
+    )]
+
+
+def _template_sequencer_select_for_fallback(params, fallback):
+    if (not fallback) and params.use_fallback_tool_rmb:
+        # Needed so we have immediate select+tweak when the default select tool is active.
+        return _template_sequencer_generic_select(
+            type=params.select_mouse,
+            value=params.select_mouse_value,
+            legacy=params.legacy,
+        )
+    return []
+
+
 def km_image_paint(params):
     items = []
     keymap = (
@@ -6172,7 +6244,7 @@ def km_image_editor_tool_uv_select_box(params, *, fallback):
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
                 "uv.select_box",
                 # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
             *_template_uv_select_for_fallback(params, fallback),
         ]},
     )
@@ -6185,8 +6257,7 @@ def km_image_editor_tool_uv_select_circle(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
                 "uv.select_circle",
-                type=params.select_tweak if fallback else params.tool_mouse,
-                value='ANY' if fallback else 'PRESS',
+                **(params.select_tweak_event if fallback else {"type": params.tool_mouse, "value": 'PRESS'}),
                 properties=[("wait_for_input", False)])),
             # No selection fallback since this operates on press.
         ]},
@@ -6201,7 +6272,7 @@ def km_image_editor_tool_uv_select_lasso(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
                 "uv.select_lasso",
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
             *_template_uv_select_for_fallback(params, fallback),
         ]},
     )
@@ -6365,7 +6436,7 @@ def km_3d_view_tool_select_box(params, *, fallback):
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "view3d.select_box",
                 # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
             *_template_view3d_select_for_fallback(params, fallback),
         ]},
     )
@@ -6395,7 +6466,7 @@ def km_3d_view_tool_select_lasso(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "view3d.select_lasso",
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
             *_template_view3d_select_for_fallback(params, fallback),
         ]}
     )
@@ -7251,7 +7322,7 @@ def km_3d_view_tool_edit_gpencil_select_box(params, *, fallback):
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "gpencil.select_box",
                 # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
             *_template_view3d_gpencil_select_for_fallback(params, fallback),
         ]},
     )
@@ -7281,7 +7352,7 @@ def km_3d_view_tool_edit_gpencil_select_lasso(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "gpencil.select_lasso",
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **(params.select_tweak_event if fallback else params.tool_tweak_event))),
             *_template_view3d_gpencil_select_for_fallback(params, fallback),
         ]}
     )
@@ -7405,11 +7476,16 @@ def km_3d_view_tool_sculpt_gpencil_select_lasso(params):
 
 def km_sequencer_editor_tool_select(params, *, fallback):
     return (
-        # TODO, fall-back tool support.
-        _fallback_id("Sequencer Tool: Select", fallback),
+        _fallback_id("Sequencer Tool: Tweak", fallback),
         {"space_type": 'SEQUENCE_EDITOR', "region_type": 'WINDOW'},
         {"items": [
-            ("sequencer.select", {"type": params.select_mouse, "value": 'PRESS'}, None),
+            *([] if fallback else
+              _template_items_tool_select(params, "sequencer.select", "sequencer.cursor_set", extend="toggle")
+              ),
+            *([] if (not params.use_fallback_tool_rmb) else _template_sequencer_generic_select(
+                type=params.select_mouse, value=params.select_mouse_value, legacy=params.legacy)),
+
+            # Ignored for preview.
             *_template_items_change_frame(params),
         ]},
     )
@@ -7417,16 +7493,18 @@ def km_sequencer_editor_tool_select(params, *, fallback):
 
 def km_sequencer_editor_tool_select_box(params, *, fallback):
     return (
-        # TODO, fall-back tool support.
         _fallback_id("Sequencer Tool: Select Box", fallback),
         {"space_type": 'SEQUENCE_EDITOR', "region_type": 'WINDOW'},
         {"items": [
             # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-            *_template_items_tool_select_actions_simple(
-                "sequencer.select_box", **params.tool_tweak_event,
-                properties=[("tweak", params.select_mouse == 'LEFTMOUSE')],
-            ),
+            *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
+                "sequencer.select_box",
+                **(params.select_tweak_event if fallback else params.tool_tweak_event),
+                properties=[("tweak", params.select_mouse == 'LEFTMOUSE')])),
+            *_template_sequencer_select_for_fallback(params, fallback),
+
             # RMB select can already set the frame, match the tweak tool.
+            # Ignored for preview.
             *(_template_items_change_frame(params)
               if params.select_mouse == 'LEFTMOUSE' else []),
         ]},
@@ -7443,6 +7521,19 @@ def km_sequencer_editor_tool_generic_sample(params):
     )
 
 
+def km_sequencer_editor_tool_cursor(params):
+    return (
+        "Sequencer Tool: Cursor",
+        {"space_type": 'SEQUENCE_EDITOR', "region_type": 'WINDOW'},
+        {"items": [
+            ("sequencer.cursor_set", {"type": params.tool_mouse, "value": 'PRESS'}, None),
+            # Don't use `tool_maybe_tweak_event` since it conflicts with `PRESS` that places the cursor.
+            ("transform.translate", params.tool_tweak_event,
+             {"properties": [("release_confirm", True), ("cursor_transform", True)]}),
+        ]},
+    )
+
+
 def km_sequencer_editor_tool_blade(_params):
     return (
         "Sequencer Tool: Blade",
@@ -7745,6 +7836,7 @@ def generate_keymaps(params=None):
         *(km_sequencer_editor_tool_select_box(params, fallback=fallback) for fallback in (False, True)),
         km_sequencer_editor_tool_blade(params),
         km_sequencer_editor_tool_generic_sample(params),
+        km_sequencer_editor_tool_cursor(params),
         km_sequencer_editor_tool_scale(params),
         km_sequencer_editor_tool_rotate(params),
         km_sequencer_editor_tool_move(params),

release/scripts/startup/bl_operators/rigidbody.py

@@ -176,7 +176,7 @@ class BakeToKeyframes(Operator):
                         # NOTE: assume that on first frame, the starting rotation is appropriate
                         obj.rotation_euler = mat.to_euler(rot_mode, obj.rotation_euler)
 
-                bpy.ops.anim.keyframe_insert(type='BUILTIN_KSI_LocRot', confirm_success=False)
+                bpy.ops.anim.keyframe_insert(type='BUILTIN_KSI_LocRot')
 
             # remove baked objects from simulation
             bpy.ops.rigidbody.objects_remove()

release/scripts/startup/bl_ui/properties_paint_common.py

@@ -1151,7 +1151,8 @@ def brush_basic__draw_color_selector(context, layout, brush, gp_settings, props)
             if len(txt_ma) > maxw:
                 txt_ma = txt_ma[:maxw - 5] + '..' + txt_ma[-3:]
 
-    sub = row.row()
+    sub = row.row(align=True)
+    sub.enabled = not gp_settings.use_material_pin
     sub.ui_units_x = 8
     sub.popover(
         panel="TOPBAR_PT_gpencil_materials",

release/scripts/startup/bl_ui/space_filebrowser.py

@@ -36,12 +36,6 @@ class FILEBROWSER_HT_header(Header):
         space_data = context.space_data
         params = space_data.params
 
-        row = layout.row(align=True)
-        row.prop(params, "asset_library_ref", text="")
-        # External libraries don't auto-refresh, add refresh button.
-        if params.asset_library_ref != 'LOCAL':
-            row.operator("file.refresh", text="", icon='FILE_REFRESH')
-
         layout.separator_spacer()
 
         layout.prop(params, "import_type", text="")

release/scripts/startup/bl_ui/space_sequencer.py

@@ -1021,6 +1021,43 @@ class SEQUENCER_MT_context_menu(Menu):
         layout.menu("SEQUENCER_MT_strip_lock_mute")
 
 
+class SEQUENCER_MT_preview_context_menu(Menu):
+    bl_label = "Sequencer Preview Context Menu"
+
+    def draw(self, context):
+        layout = self.layout
+
+        layout.operator_context = 'INVOKE_REGION_WIN'
+
+        props = layout.operator("wm.call_panel", text="Rename...")
+        props.name = "TOPBAR_PT_name"
+        props.keep_open = False
+
+        # TODO: support in preview.
+        # layout.operator("sequencer.delete", text="Delete")
+
+        strip = context.active_sequence_strip
+
+        if strip:
+            pass
+
+
+class SEQUENCER_MT_pivot_pie(Menu):
+    bl_label = "Pivot Point"
+
+    def draw(self, context):
+        layout = self.layout
+        pie = layout.menu_pie()
+
+        tool_settings = context.tool_settings
+        sequencer_tool_settings = context.tool_settings.sequencer_tool_settings
+
+        pie.prop_enum(sequencer_tool_settings, "pivot_point", value='CENTER')
+        pie.prop_enum(sequencer_tool_settings, "pivot_point", value='CURSOR')
+        pie.prop_enum(sequencer_tool_settings, "pivot_point", value='INDIVIDUAL_ORIGINS')
+        pie.prop_enum(sequencer_tool_settings, "pivot_point", value='MEDIAN')
+
+
 class SequencerButtonsPanel:
     bl_space_type = 'SEQUENCE_EDITOR'
     bl_region_type = 'UI'
@@ -1051,7 +1088,7 @@ class SequencerButtonsPanel_Output:
 class SequencerColorTagPicker:
     bl_space_type = 'SEQUENCE_EDITOR'
     bl_region_type = 'UI'
-    
+
     @staticmethod
     def has_sequencer(context):
         return (context.space_data.view_type in {'SEQUENCER', 'SEQUENCER_PREVIEW'})
@@ -1758,7 +1795,6 @@ class SEQUENCER_PT_adjust_sound(SequencerButtonsPanel, Panel):
 
     def draw(self, context):
         layout = self.layout
-        layout.use_property_split = False
 
         st = context.space_data
         overlay_settings = st.timeline_overlay
@@ -1768,16 +1804,7 @@ class SEQUENCER_PT_adjust_sound(SequencerButtonsPanel, Panel):
         layout.active = not strip.mute
 
         if sound is not None:
-            col = layout.column()
-
-            split = col.split(factor=0.4)
-            split.label(text="")
-            split.prop(sound, "use_mono")
-            if overlay_settings.waveform_display_type == 'DEFAULT_WAVEFORMS':
-                split = col.split(factor=0.4)
-                split.label(text="")
-                split.prop(strip, "show_waveform")
-
+            layout.use_property_split = True
             col = layout.column()
 
             split = col.split(factor=0.4)
@@ -1790,15 +1817,37 @@ class SEQUENCER_PT_adjust_sound(SequencerButtonsPanel, Panel):
             split.label(text="Pitch")
             split.prop(strip, "pitch", text="")
 
+            audio_channels = context.scene.render.ffmpeg.audio_channels
+            pan_enabled = sound.use_mono and audio_channels != 'MONO'
+            pan_text = "%.2f°" % (strip.pan * 90)
+
             split = col.split(factor=0.4)
             split.alignment = 'RIGHT'
             split.label(text="Pan")
-            audio_channels = context.scene.render.ffmpeg.audio_channels
-            pan_text = ""
+            split.prop(strip, "pan", text="")
+            split.enabled = pan_enabled
+
             if audio_channels != 'MONO' and audio_channels != 'STEREO':
-                pan_text = "%.2f°" % (strip.pan * 90)
-            split.prop(strip, "pan", text=pan_text)
-            split.enabled = sound.use_mono and audio_channels != 'MONO'
+                split = col.split(factor=0.4)
+                split.alignment = 'RIGHT'
+                split.label(text="Pan Angle")
+                split.enabled = pan_enabled
+                subsplit = split.row()
+                subsplit.alignment = 'CENTER'
+                subsplit.label(text=pan_text)
+                subsplit.label(text=" ")  # Compensate for no decorate.
+                subsplit.enabled = pan_enabled
+
+            layout.use_property_split = False
+            col = layout.column()
+
+            split = col.split(factor=0.4)
+            split.label(text="")
+            split.prop(sound, "use_mono")
+            if overlay_settings.waveform_display_type == 'DEFAULT_WAVEFORMS':
+                split = col.split(factor=0.4)
+                split.label(text="")
+                split.prop(strip, "show_waveform")
 
 
 
@@ -2152,6 +2201,22 @@ class SEQUENCER_PT_view(SequencerButtonsPanel_Output, Panel):
             col.prop(st, "show_separate_color")
 
 
+class SEQUENCER_PT_view_cursor(SequencerButtonsPanel_Output, Panel):
+    bl_category = "View"
+    bl_label = "2D Cursor"
+
+    def draw(self, context):
+        layout = self.layout
+
+        st = context.space_data
+
+        layout.use_property_split = True
+        layout.use_property_decorate = False
+
+        col = layout.column()
+        col.prop(st, "cursor_location", text="Location")
+
+
 class SEQUENCER_PT_frame_overlay(SequencerButtonsPanel_Output, Panel):
     bl_label = "Frame Overlay"
     bl_category = "View"
@@ -2430,6 +2495,8 @@ classes = (
     SEQUENCER_MT_strip_lock_mute,
     SEQUENCER_MT_color_tag_picker,
     SEQUENCER_MT_context_menu,
+    SEQUENCER_MT_preview_context_menu,
+    SEQUENCER_MT_pivot_pie,
 
     SEQUENCER_PT_color_tag_picker,
 
@@ -2466,6 +2533,7 @@ classes = (
     SEQUENCER_PT_custom_props,
 
     SEQUENCER_PT_view,
+    SEQUENCER_PT_view_cursor,
     SEQUENCER_PT_frame_overlay,
     SEQUENCER_PT_view_safe_areas,
     SEQUENCER_PT_view_safe_areas_center_cut,

release/scripts/startup/bl_ui/space_toolsystem_toolbar.py

@@ -2440,6 +2440,19 @@ class _defs_node_edit:
 
 class _defs_sequencer_generic:
 
+    @ToolDef.from_fn
+    def cursor():
+        return dict(
+            idname="builtin.cursor",
+            label="Cursor",
+            description=(
+                "Set the cursor location, drag to transform"
+            ),
+            icon="ops.generic.cursor",
+            keymap="Sequencer Tool: Cursor",
+            options={'KEYMAP_FALLBACK'},
+        )
+
     @ToolDef.from_fn
     def blade():
         def draw_settings(_context, layout, tool):
@@ -2508,10 +2521,10 @@ class _defs_sequencer_select:
     def select():
         return dict(
             idname="builtin.select",
-            label="Select",
+            label="Tweak",
             icon="ops.generic.select",
             widget=None,
-            keymap="Sequencer Tool: Select",
+            keymap="Sequencer Tool: Tweak",
         )
 
     @ToolDef.from_fn
@@ -3094,6 +3107,8 @@ class SEQUENCER_PT_tools_active(ToolSelectPanelHelper, Panel):
         ],
         'PREVIEW': [
             *_tools_select,
+            _defs_sequencer_generic.cursor,
+            None,
             _defs_sequencer_generic.translate,
             _defs_sequencer_generic.rotate,
             _defs_sequencer_generic.scale,
@@ -3106,6 +3121,8 @@ class SEQUENCER_PT_tools_active(ToolSelectPanelHelper, Panel):
         ],
         'SEQUENCER_PREVIEW': [
             *_tools_select,
+            _defs_sequencer_generic.cursor,
+            None,
             _defs_sequencer_generic.translate,
             _defs_sequencer_generic.rotate,
             _defs_sequencer_generic.scale,

release/scripts/startup/bl_ui/space_topbar.py

@@ -634,6 +634,8 @@ class TOPBAR_MT_window(Menu):
 
         layout = self.layout
 
+        operator_context_default = layout.operator_context
+
         layout.operator("wm.window_new")
         layout.operator("wm.window_new_main")
 
@@ -655,7 +657,15 @@ class TOPBAR_MT_window(Menu):
         layout.separator()
 
         layout.operator("screen.screenshot")
+
+        # Showing the status in the area doesn't work well in this case.
+        # - From the top-bar, the text replaces the file-menu (not so bad but strange).
+        # - From menu-search it replaces the area that the user may want to screen-shot.
+        # Setting the context to screen causes the status to show in the global status-bar.
+        print(layout.operator_context)
+        layout.operator_context = 'INVOKE_SCREEN'
         layout.operator("screen.screenshot_area")
+        layout.operator_context = operator_context_default
 
         if sys.platform[:3] == "win":
             layout.separator()

source/blender/blenfont/intern/blf_internal_types.h

@@ -152,7 +152,7 @@ typedef struct FontBufInfoBLF {
   struct ColorManagedDisplay *display;
 
   /* and the color, the alphas is get from the glyph!
-   * color is srgb space */
+   * color is sRGB space */
   float col_init[4];
   /* cached conversion from 'col_init' */
   unsigned char col_char[4];

source/blender/blenkernel/BKE_DerivedMesh.h

@@ -157,15 +157,6 @@ struct DerivedMesh {
   int (*getNumLoops)(DerivedMesh *dm);
   int (*getNumPolys)(DerivedMesh *dm);
 
-  /** Copy a single vert/edge/tessellated face from the derived mesh into
-   * `*r_{vert/edge/face}`. note that the current implementation
-   * of this function can be quite slow, iterating over all
-   * elements (editmesh)
-   */
-  void (*getVert)(DerivedMesh *dm, int index, struct MVert *r_vert);
-  void (*getEdge)(DerivedMesh *dm, int index, struct MEdge *r_edge);
-  void (*getTessFace)(DerivedMesh *dm, int index, struct MFace *r_face);
-
   /** Return a pointer to the entire array of verts/edges/face from the
    * derived mesh. if such an array does not exist yet, it will be created,
    * and freed on the next ->release(). consider using getVert/Edge/Face if

source/blender/blenkernel/BKE_appdir.h

@@ -16,7 +16,7 @@
 #pragma once
 
 /** \file
- * \ingroup bli
+ * \ingroup bke
  */
 
 #include <stddef.h>

source/blender/blenkernel/BKE_attribute_access.hh

@@ -373,9 +373,9 @@ class CustomDataAttributes {
                          const AttributeDomain domain) const;
 };
 
-/* --------------------------------------------------------------------
- * #AttributeIDRef inline methods.
- */
+/* -------------------------------------------------------------------- */
+/** \name #AttributeIDRef Inline Methods
+ * \{ */
 
 inline AttributeIDRef::AttributeIDRef() = default;
 
@@ -438,9 +438,11 @@ inline const AnonymousAttributeID &AttributeIDRef::anonymous_id() const
   return *anonymous_id_;
 }
 
-/* --------------------------------------------------------------------
- * #OutputAttribute inline methods.
- */
+/** \} */
+
+/* -------------------------------------------------------------------- */
+/** \name #OutputAttribute Inline Methods
+ * \{ */
 
 inline OutputAttribute::OutputAttribute() = default;
 inline OutputAttribute::OutputAttribute(OutputAttribute &&other) = default;
@@ -496,11 +498,15 @@ template<typename T> inline MutableSpan<T> OutputAttribute::as_span()
   return this->as_span().typed<T>();
 }
 
+/** \} */
+
 }  // namespace blender::bke
 
-/* --------------------------------------------------------------------
- * Extern template instantiations that are defined in `intern/extern_implementations.cc`.
- */
+/* -------------------------------------------------------------------- */
+/** \name External Template Instantiations
+ *
+ * Defined in `intern/extern_implementations.cc`.
+ * \{ */
 
 namespace blender::bke {
 extern template class OutputAttribute_Typed<float>;
@@ -509,3 +515,5 @@ extern template class OutputAttribute_Typed<float3>;
 extern template class OutputAttribute_Typed<bool>;
 extern template class OutputAttribute_Typed<ColorGeometry4f>;
 }  // namespace blender::bke
+
+/** \} */

source/blender/blenkernel/BKE_curve_to_mesh.hh

@@ -20,7 +20,7 @@ struct Mesh;
 struct CurveEval;
 
 /** \file
- * \ingroup geo
+ * \ingroup bke
  */
 
 namespace blender::bke {

source/blender/blenkernel/BKE_image.h

@@ -253,13 +253,13 @@ bool BKE_image_is_stereo(struct Image *ima);
 struct RenderResult *BKE_image_acquire_renderresult(struct Scene *scene, struct Image *ima);
 void BKE_image_release_renderresult(struct Scene *scene, struct Image *ima);
 
-/* for multilayer images as well as for singlelayer */
+/* For multi-layer images as well as for single-layer. */
 bool BKE_image_is_openexr(struct Image *ima);
 
-/* for multiple slot render, call this before render */
+/* For multiple slot render, call this before render. */
 void BKE_image_backup_render(struct Scene *scene, struct Image *ima, bool free_current_slot);
 
-/* for singlelayer openexr saving */
+/* For single-layer OpenEXR saving */
 bool BKE_image_save_openexr_multiview(struct Image *ima,
                                       struct ImBuf *ibuf,
                                       const char *filepath,
@@ -285,22 +285,22 @@ void BKE_image_packfiles_from_mem(struct ReportList *reports,
                                   char *data,
                                   const size_t data_len);
 
-/* prints memory statistics for images */
+/* Prints memory statistics for images. */
 void BKE_image_print_memlist(struct Main *bmain);
 
-/* merge source into dest, and free source */
+/* Merge source into dest, and free source. */
 void BKE_image_merge(struct Main *bmain, struct Image *dest, struct Image *source);
 
-/* scale the image */
+/* Scale the image. */
 bool BKE_image_scale(struct Image *image, int width, int height);
 
-/* check if texture has alpha (depth=32) */
+/* Check if texture has alpha (depth=32). */
 bool BKE_image_has_alpha(struct Image *image);
 
-/* check if texture has gpu texture code */
+/* Check if texture has GPU texture code. */
 bool BKE_image_has_opengl_texture(struct Image *ima);
 
-/* get tile index for tiled images */
+/* Get tile index for tiled images. */
 void BKE_image_get_tile_label(struct Image *ima,
                               struct ImageTile *tile,
                               char *label,
@@ -369,10 +369,10 @@ struct ImBuf *BKE_image_get_first_ibuf(struct Image *image);
 /* Not to be use directly. */
 struct GPUTexture *BKE_image_create_gpu_texture_from_ibuf(struct Image *image, struct ImBuf *ibuf);
 
-/* Get the GPUTexture for a given `Image`.
+/* Get the #GPUTexture for a given `Image`.
  *
  * `iuser` and `ibuf` are mutual exclusive parameters. The caller can pass the `ibuf` when already
- * available. It is also required when requesting the GPUTexture for a render result. */
+ * available. It is also required when requesting the #GPUTexture for a render result. */
 struct GPUTexture *BKE_image_get_gpu_texture(struct Image *image,
                                              struct ImageUser *iuser,
                                              struct ImBuf *ibuf);

source/blender/blenkernel/BKE_main.h

@@ -201,6 +201,8 @@ typedef struct Main {
 struct Main *BKE_main_new(void);
 void BKE_main_free(struct Main *mainvar);
 
+bool BKE_main_is_empty(struct Main *bmain);
+
 void BKE_main_lock(struct Main *bmain);
 void BKE_main_unlock(struct Main *bmain);
 

source/blender/blenkernel/BKE_ocean.h

@@ -19,7 +19,7 @@
 #include <stdbool.h>
 
 /** \file
- * \ingroup bli
+ * \ingroup bke
  */
 
 #ifdef __cplusplus

source/blender/blenkernel/BKE_rigidbody.h

@@ -18,7 +18,7 @@
  */
 
 /** \file
- * \ingroup blenkernel
+ * \ingroup bke
  * \brief API for Blender-side Rigid Body stuff
  */
 

source/blender/blenkernel/BKE_vfontdata.h

@@ -20,7 +20,7 @@
 #pragma once
 
 /** \file
- * \ingroup bli
+ * \ingroup bke
  * \brief A structure to represent vector fonts,
  *   and to load them from PostScript fonts.
  */
@@ -49,11 +49,11 @@ typedef struct VChar {
   float width;
 } VChar;
 
-VFontData *BLI_vfontdata_from_freetypefont(struct PackedFile *pf);
-VFontData *BLI_vfontdata_copy(const VFontData *vfont_src, const int flag);
+VFontData *BKE_vfontdata_from_freetypefont(struct PackedFile *pf);
+VFontData *BKE_vfontdata_copy(const VFontData *vfont_src, const int flag);
 
-VChar *BLI_vfontchar_from_freetypefont(struct VFont *vfont, unsigned long character);
-VChar *BLI_vfontchar_copy(const VChar *vchar_src, const int flag);
+VChar *BKE_vfontdata_char_from_freetypefont(struct VFont *vfont, unsigned long character);
+VChar *BKE_vfontdata_char_copy(const VChar *vchar_src);
 
 #ifdef __cplusplus
 }

source/blender/blenkernel/CMakeLists.txt

@@ -60,6 +60,9 @@ set(INC
 
 set(INC_SYS
   ${ZLIB_INCLUDE_DIRS}
+
+  # For `vfontdata_freetype.c`.
+  ${FREETYPE_INCLUDE_DIRS}
 )
 
 set(SRC
@@ -140,7 +143,6 @@ set(SRC
   intern/fcurve_driver.c
   intern/fluid.c
   intern/fmodifier.c
-  intern/font.c
   intern/freestyle.c
   intern/geometry_component_curve.cc
   intern/geometry_component_instances.cc
@@ -286,6 +288,8 @@ set(SRC
   intern/tracking_util.c
   intern/undo_system.c
   intern/unit.c
+  intern/vfont.c
+  intern/vfontdata_freetype.c
   intern/volume.cc
   intern/volume_render.cc
   intern/volume_to_mesh.cc
@@ -360,7 +364,6 @@ set(SRC
   BKE_fcurve.h
   BKE_fcurve_driver.h
   BKE_fluid.h
-  BKE_font.h
   BKE_freestyle.h
   BKE_geometry_set.h
   BKE_geometry_set.hh
@@ -453,6 +456,8 @@ set(SRC
   BKE_tracking.h
   BKE_undo_system.h
   BKE_unit.h
+  BKE_vfont.h
+  BKE_vfontdata.h
   BKE_volume.h
   BKE_volume_render.h
   BKE_volume_to_mesh.hh
@@ -502,6 +507,9 @@ set(LIB
   bf_rna
   bf_shader_fx
   bf_simulation
+
+  # For `vfontdata_freetype.c`.
+  ${FREETYPE_LIBRARY}
 )
 
 if(WITH_BINRELOC)

source/blender/blenkernel/intern/action_mirror.c

@@ -327,10 +327,10 @@ static void action_flip_pchan(Object *ob_arm,
      * the X-axis, it turns into a 180 degree rotation over the Y-axis.
      * This has only been observed with bones that can't be flipped,
      * hence the check for `pchan_flip`. */
-    const float unit_x[4] = {1.0f, 0.0f, 0.0f, 0.0f};
-    const bool is_problematic = pchan_flip == NULL &&
-                                fabsf(dot_v4v4(pchan->bone->arm_mat[0], unit_x)) <= 1e-6;
-    if (is_problematic) {
+    const float unit_x[3] = {1.0f, 0.0f, 0.0f};
+    const bool is_x_axis_orthogonal = (pchan_flip == NULL) &&
+                                      (fabsf(dot_v3v3(pchan->bone->arm_mat[0], unit_x)) <= 1e-6f);
+    if (is_x_axis_orthogonal) {
       /* Matrix needs to flip both the X and Z axes to come out right. */
       float extra_mat[4][4] = {
           {-1.0f, 0.0f, 0.0f, 0.0f},

source/blender/blenkernel/intern/asset_catalog.cc

@@ -327,30 +327,26 @@ CatalogFilePath AssetCatalogService::find_suitable_cdf_path_for_writing(
                  "A non-empty .blend file path is required to be able to determine where the "
                  "catalog definition file should be put");
 
+  /* Ask the asset library API for an appropriate location.  */
+  char suitable_root_path[PATH_MAX];
+  const bool asset_lib_root_found = BKE_asset_library_find_suitable_root_path_from_path(
+      blend_file_path.c_str(), suitable_root_path);
+  if (asset_lib_root_found) {
+    char asset_lib_cdf_path[PATH_MAX];
+    BLI_path_join(asset_lib_cdf_path,
+                  sizeof(asset_lib_cdf_path),
+                  suitable_root_path,
+                  DEFAULT_CATALOG_FILENAME.c_str(),
+                  NULL);
+    return asset_lib_cdf_path;
+  }
+
   /* Determine the default CDF path in the same directory of the blend file. */
   char blend_dir_path[PATH_MAX];
   BLI_split_dir_part(blend_file_path.c_str(), blend_dir_path, sizeof(blend_dir_path));
   const CatalogFilePath cdf_path_next_to_blend = asset_definition_default_file_path_from_dir(
       blend_dir_path);
-
-  if (BLI_exists(cdf_path_next_to_blend.c_str())) {
-    /* - The directory containing the blend file has a blender_assets.cats.txt file?
-     *    -> Merge with & write to that file. */
-    return cdf_path_next_to_blend;
-  }
-
-  /* - There's no definition file next to the .blend file.
-   *    -> Ask the asset library API for an appropriate location.  */
-  char suitable_root_path[PATH_MAX];
-  BKE_asset_library_find_suitable_root_path_from_path(blend_file_path.c_str(), suitable_root_path);
-  char asset_lib_cdf_path[PATH_MAX];
-  BLI_path_join(asset_lib_cdf_path,
-                sizeof(asset_lib_cdf_path),
-                suitable_root_path,
-                DEFAULT_CATALOG_FILENAME.c_str(),
-                NULL);
-
-  return asset_lib_cdf_path;
+  return cdf_path_next_to_blend;
 }
 
 std::unique_ptr<AssetCatalogDefinitionFile> AssetCatalogService::construct_cdf_in_memory(

source/blender/blenkernel/intern/asset_catalog_test.cc

@@ -91,6 +91,14 @@ class AssetCatalogTest : public testing::Test {
     temp_library_path_ = "";
   }
 
+  void TearDown() override
+  {
+    if (!temp_library_path_.empty()) {
+      BLI_delete(temp_library_path_.c_str(), true, true);
+      temp_library_path_ = "";
+    }
+  }
+
   /* Register a temporary path, which will be removed at the end of the test.
    * The returned path ends in a slash. */
   CatalogFilePath use_temp_path()
@@ -177,12 +185,74 @@ class AssetCatalogTest : public testing::Test {
     });
   }
 
-  void TearDown() override
+  /* Used by on_blendfile_save__from_memory_into_existing_asset_lib* test functions. */
+  void save_from_memory_into_existing_asset_lib(const bool should_top_level_cdf_exist)
   {
-    if (!temp_library_path_.empty()) {
-      BLI_delete(temp_library_path_.c_str(), true, true);
-      temp_library_path_ = "";
+    const CatalogFilePath target_dir = create_temp_path(); /* Has trailing slash. */
+    const CatalogFilePath original_cdf_file = asset_library_root_ + "/blender_assets.cats.txt";
+    const CatalogFilePath registered_asset_lib = target_dir + "my_asset_library/";
+    const CatalogFilePath asset_lib_subdir = registered_asset_lib + "subdir/";
+    CatalogFilePath cdf_toplevel = registered_asset_lib +
+                                   AssetCatalogService::DEFAULT_CATALOG_FILENAME;
+    CatalogFilePath cdf_in_subdir = asset_lib_subdir +
+                                    AssetCatalogService::DEFAULT_CATALOG_FILENAME;
+    BLI_path_slash_native(cdf_toplevel.data());
+    BLI_path_slash_native(cdf_in_subdir.data());
+
+    /* Set up a temporary asset library for testing. */
+    bUserAssetLibrary *asset_lib_pref = BKE_preferences_asset_library_add(
+        &U, "Test", registered_asset_lib.c_str());
+    ASSERT_NE(nullptr, asset_lib_pref);
+    ASSERT_TRUE(BLI_dir_create_recursive(asset_lib_subdir.c_str()));
+
+    if (should_top_level_cdf_exist) {
+      ASSERT_EQ(0, BLI_copy(original_cdf_file.c_str(), cdf_toplevel.c_str()));
     }
+
+    /* Create an empty CDF to add complexity. It should not save to this, but to the top-level
+     * one.*/
+    ASSERT_TRUE(BLI_file_touch(cdf_in_subdir.c_str()));
+    ASSERT_EQ(0, BLI_file_size(cdf_in_subdir.c_str()));
+
+    /* Create the catalog service without loading the already-existing CDF. */
+    TestableAssetCatalogService service;
+    const CatalogFilePath blendfilename = asset_lib_subdir + "some_file.blend";
+    const AssetCatalog *cat = service.create_catalog("some/catalog/path");
+
+    /* Mock that the blend file is written to the directory already containing a CDF. */
+    ASSERT_TRUE(service.write_to_disk_on_blendfile_save(blendfilename));
+
+    /* Test that the CDF still exists in the expected location. */
+    EXPECT_TRUE(BLI_exists(cdf_toplevel.c_str()));
+    const CatalogFilePath backup_filename = cdf_toplevel + "~";
+    const bool backup_exists = BLI_exists(backup_filename.c_str());
+    EXPECT_EQ(should_top_level_cdf_exist, backup_exists)
+        << "Overwritten CDF should have been backed up.";
+
+    /* Test that the in-memory CDF has the expected file path. */
+    AssetCatalogDefinitionFile *cdf = service.get_catalog_definition_file();
+    BLI_path_slash_native(cdf->file_path.data());
+    EXPECT_EQ(cdf_toplevel, cdf->file_path);
+
+    /* Test that the in-memory catalogs have been merged with the on-disk one. */
+    AssetCatalogService loaded_service(cdf_toplevel);
+    loaded_service.load_from_disk();
+    EXPECT_NE(nullptr, loaded_service.find_catalog(cat->catalog_id));
+
+    /* This catalog comes from a pre-existing CDF that should have been merged.
+     * However, if the file doesn't exist, so does the catalog. */
+    AssetCatalog *poses_ellie_catalog = loaded_service.find_catalog(UUID_POSES_ELLIE);
+    if (should_top_level_cdf_exist) {
+      EXPECT_NE(nullptr, poses_ellie_catalog);
+    }
+    else {
+      EXPECT_EQ(nullptr, poses_ellie_catalog);
+    }
+
+    /* Test that the "red herring" CDF has not been touched. */
+    EXPECT_EQ(0, BLI_file_size(cdf_in_subdir.c_str()));
+
+    BKE_preferences_asset_library_remove(&U, asset_lib_pref);
   }
 };
 
@@ -525,51 +595,21 @@ TEST_F(AssetCatalogTest, on_blendfile_save__from_memory_into_existing_cdf_and_me
   EXPECT_NE(nullptr, loaded_service.find_catalog(UUID_POSES_ELLIE));
 }
 
-/* Create some catalogs in memory, save to subdirectory of a registered asset library. */
+/* Create some catalogs in memory, save to subdirectory of a registered asset library, where the
+ * subdirectory also contains a CDF. This should still write to the top-level dir of the asset
+ * library. */
+TEST_F(AssetCatalogTest,
+       on_blendfile_save__from_memory_into_existing_asset_lib_without_top_level_cdf)
+{
+  save_from_memory_into_existing_asset_lib(true);
+}
+
+/* Create some catalogs in memory, save to subdirectory of a registered asset library, where the
+ * subdirectory contains a CDF, but the top-level directory does not. This should still write to
+ * the top-level dir of the asset library. */
 TEST_F(AssetCatalogTest, on_blendfile_save__from_memory_into_existing_asset_lib)
 {
-  const CatalogFilePath target_dir = create_temp_path(); /* Has trailing slash. */
-  const CatalogFilePath original_cdf_file = asset_library_root_ + "/blender_assets.cats.txt";
-  const CatalogFilePath registered_asset_lib = target_dir + "my_asset_library/";
-  CatalogFilePath writable_cdf_file = registered_asset_lib +
-                                      AssetCatalogService::DEFAULT_CATALOG_FILENAME;
-  BLI_path_slash_native(writable_cdf_file.data());
-
-  /* Set up a temporary asset library for testing. */
-  bUserAssetLibrary *asset_lib_pref = BKE_preferences_asset_library_add(
-      &U, "Test", registered_asset_lib.c_str());
-  ASSERT_NE(nullptr, asset_lib_pref);
-  ASSERT_TRUE(BLI_dir_create_recursive(registered_asset_lib.c_str()));
-  ASSERT_EQ(0, BLI_copy(original_cdf_file.c_str(), writable_cdf_file.c_str()));
-
-  /* Create the catalog service without loading the already-existing CDF. */
-  TestableAssetCatalogService service;
-  const CatalogFilePath blenddirname = registered_asset_lib + "subdirectory/";
-  const CatalogFilePath blendfilename = blenddirname + "some_file.blend";
-  ASSERT_TRUE(BLI_dir_create_recursive(blenddirname.c_str()));
-  const AssetCatalog *cat = service.create_catalog("some/catalog/path");
-
-  /* Mock that the blend file is written to the directory already containing a CDF. */
-  ASSERT_TRUE(service.write_to_disk_on_blendfile_save(blendfilename));
-
-  /* Test that the CDF still exists in the expected location. */
-  EXPECT_TRUE(BLI_exists(writable_cdf_file.c_str()));
-  const CatalogFilePath backup_filename = writable_cdf_file + "~";
-  EXPECT_TRUE(BLI_exists(backup_filename.c_str()))
-      << "Overwritten CDF should have been backed up.";
-
-  /* Test that the in-memory CDF has the expected file path. */
-  AssetCatalogDefinitionFile *cdf = service.get_catalog_definition_file();
-  BLI_path_slash_native(cdf->file_path.data());
-  EXPECT_EQ(writable_cdf_file, cdf->file_path);
-
-  /* Test that the in-memory catalogs have been merged with the on-disk one. */
-  AssetCatalogService loaded_service(writable_cdf_file);
-  loaded_service.load_from_disk();
-  EXPECT_NE(nullptr, loaded_service.find_catalog(cat->catalog_id));
-  EXPECT_NE(nullptr, loaded_service.find_catalog(UUID_POSES_ELLIE));
-
-  BKE_preferences_asset_library_remove(&U, asset_lib_pref);
+  save_from_memory_into_existing_asset_lib(false);
 }
 
 TEST_F(AssetCatalogTest, create_first_catalog_from_scratch)

source/blender/blenkernel/intern/blendfile.c

@@ -347,7 +347,7 @@ static void setup_app_data(bContext *C,
   /* FIXME: Same as above, readfile's `do_version` do not allow to create new IDs. */
   /* TODO: Once this is definitively validated for 3.0 and option to not do it is removed, add a
    * version bump and check here. */
-  if (!USER_EXPERIMENTAL_TEST(&U, no_proxy_to_override_conversion)) {
+  if (mode != LOAD_UNDO && !USER_EXPERIMENTAL_TEST(&U, no_proxy_to_override_conversion)) {
     BKE_lib_override_library_main_proxy_convert(bmain, reports);
   }
 

source/blender/blenkernel/intern/bpath.c

@@ -15,7 +15,7 @@
  */
 
 /** \file
- * \ingroup bli
+ * \ingroup bke
  */
 
 /* TODO:
@@ -66,13 +66,13 @@
 #include "BLI_blenlib.h"
 #include "BLI_utildefines.h"
 
-#include "BKE_font.h"
 #include "BKE_image.h"
 #include "BKE_lib_id.h"
 #include "BKE_library.h"
 #include "BKE_main.h"
 #include "BKE_node.h"
 #include "BKE_report.h"
+#include "BKE_vfont.h"
 
 #include "BKE_bpath.h" /* own include */
 

source/blender/blenkernel/intern/cdderivedmesh.c

@@ -103,24 +103,6 @@ static int cdDM_getNumPolys(DerivedMesh *dm)
   return dm->numPolyData;
 }
 
-static void cdDM_getVert(DerivedMesh *dm, int index, MVert *r_vert)
-{
-  CDDerivedMesh *cddm = (CDDerivedMesh *)dm;
-  *r_vert = cddm->mvert[index];
-}
-
-static void cdDM_getEdge(DerivedMesh *dm, int index, MEdge *r_edge)
-{
-  CDDerivedMesh *cddm = (CDDerivedMesh *)dm;
-  *r_edge = cddm->medge[index];
-}
-
-static void cdDM_getTessFace(DerivedMesh *dm, int index, MFace *r_face)
-{
-  CDDerivedMesh *cddm = (CDDerivedMesh *)dm;
-  *r_face = cddm->mface[index];
-}
-
 static void cdDM_copyVertArray(DerivedMesh *dm, MVert *r_vert)
 {
   CDDerivedMesh *cddm = (CDDerivedMesh *)dm;
@@ -231,10 +213,6 @@ static CDDerivedMesh *cdDM_create(const char *desc)
   dm->getNumLoops = cdDM_getNumLoops;
   dm->getNumPolys = cdDM_getNumPolys;
 
-  dm->getVert = cdDM_getVert;
-  dm->getEdge = cdDM_getEdge;
-  dm->getTessFace = cdDM_getTessFace;
-
   dm->copyVertArray = cdDM_copyVertArray;
   dm->copyEdgeArray = cdDM_copyEdgeArray;
   dm->copyTessFaceArray = cdDM_copyTessFaceArray;

source/blender/blenkernel/intern/curve.c

@@ -52,13 +52,13 @@
 #include "BKE_curve.h"
 #include "BKE_curveprofile.h"
 #include "BKE_displist.h"
-#include "BKE_font.h"
 #include "BKE_idtype.h"
 #include "BKE_key.h"
 #include "BKE_lib_id.h"
 #include "BKE_lib_query.h"
 #include "BKE_main.h"
 #include "BKE_object.h"
+#include "BKE_vfont.h"
 
 #include "DEG_depsgraph.h"
 #include "DEG_depsgraph_query.h"

source/blender/blenkernel/intern/curve_convert.c

@@ -26,9 +26,9 @@
 
 #include "BKE_curve.h"
 #include "BKE_displist.h"
-#include "BKE_font.h"
 #include "BKE_lib_id.h"
 #include "BKE_modifier.h"
+#include "BKE_vfont.h"
 
 #include "DEG_depsgraph.h"
 #include "DEG_depsgraph_query.h"

source/blender/blenkernel/intern/displist.cc

@@ -47,7 +47,6 @@
 #include "BKE_anim_path.h"
 #include "BKE_curve.h"
 #include "BKE_displist.h"
-#include "BKE_font.h"
 #include "BKE_geometry_set.hh"
 #include "BKE_key.h"
 #include "BKE_lattice.h"
@@ -58,6 +57,7 @@
 #include "BKE_modifier.h"
 #include "BKE_object.h"
 #include "BKE_spline.hh"
+#include "BKE_vfont.h"
 
 #include "BLI_sys_types.h" /* For #intptr_t support. */
 

source/blender/blenkernel/intern/editmesh_bvh.c

@@ -221,7 +221,7 @@ static void bmbvh_tri_from_face(const float *cos[3],
   }
 }
 
-/* taken from bvhutils.c */
+/* Taken from `bvhutils.c`. */
 
 /* -------------------------------------------------------------------- */
 /* BKE_bmbvh_ray_cast */

source/blender/blenkernel/intern/gpencil.c

@@ -1044,7 +1044,6 @@ bGPDlayer *BKE_gpencil_layer_duplicate(const bGPDlayer *gpl_src,
                                        const bool dup_frames,
                                        const bool dup_strokes)
 {
-  const bGPDframe *gpf_src;
   bGPDframe *gpf_dst;
   bGPDlayer *gpl_dst;
 
@@ -1063,7 +1062,7 @@ bGPDlayer *BKE_gpencil_layer_duplicate(const bGPDlayer *gpl_src,
   /* copy frames */
   BLI_listbase_clear(&gpl_dst->frames);
   if (dup_frames) {
-    for (gpf_src = gpl_src->frames.first; gpf_src; gpf_src = gpf_src->next) {
+    LISTBASE_FOREACH (bGPDframe *, gpf_src, &gpl_src->frames) {
       /* make a copy of source frame */
       gpf_dst = BKE_gpencil_frame_duplicate(gpf_src, dup_strokes);
       BLI_addtail(&gpl_dst->frames, gpf_dst);

source/blender/blenkernel/intern/image.c

@@ -1560,9 +1560,9 @@ bool BKE_imtype_requires_linear_float(const char imtype)
 
 char BKE_imtype_valid_channels(const char imtype, bool write_file)
 {
-  char chan_flag = IMA_CHAN_FLAG_RGB; /* assume all support rgb */
+  char chan_flag = IMA_CHAN_FLAG_RGB; /* Assume all support RGB. */
 
-  /* alpha */
+  /* Alpha. */
   switch (imtype) {
     case R_IMF_IMTYPE_BMP:
       if (write_file) {
@@ -1583,7 +1583,7 @@ char BKE_imtype_valid_channels(const char imtype, bool write_file)
       break;
   }
 
-  /* bw */
+  /* BW. */
   switch (imtype) {
     case R_IMF_IMTYPE_BMP:
     case R_IMF_IMTYPE_PNG:
@@ -3078,8 +3078,7 @@ int BKE_imbuf_write_as(ImBuf *ibuf, const char *name, ImageFormatData *imf, cons
   ImBuf ibuf_back = *ibuf;
   int ok;
 
-  /* all data is rgba anyway,
-   * this just controls how to save for some formats */
+  /* All data is RGBA anyway, this just controls how to save for some formats. */
   ibuf->planes = imf->planes;
 
   ok = BKE_imbuf_write(ibuf, name, imf);
@@ -4611,7 +4610,7 @@ static ImBuf *load_image_single(Image *ima,
       image_init_after_load(ima, iuser, ibuf);
       *r_assign = true;
 
-      /* make packed file for autopack */
+      /* Make packed file for auto-pack. */
       if ((has_packed == false) && (G.fileflags & G_FILE_AUTOPACK)) {
         ImagePackedFile *imapf = MEM_mallocN(sizeof(ImagePackedFile), "Image Pack-file");
         BLI_addtail(&ima->packedfiles, imapf);