build_files/build_environment/cmake/versions.cmake

@@ -11,7 +11,7 @@
 # dependencies have one assigned.
 
 set(ZLIB_VERSION 1.2.13)
-set(ZLIB_URI https://zlib.net/zlib-${ZLIB_VERSION}.tar.gz)
+set(ZLIB_URI https://github.com/madler/zlib/releases/download/v${ZLIB_VERSION}/zlib-${ZLIB_VERSION}.tar.gz)
 set(ZLIB_HASH 9b8aa094c4e5765dabf4da391f00d15c)
 set(ZLIB_HASH_TYPE MD5)
 set(ZLIB_FILE zlib-${ZLIB_VERSION}.tar.gz)

doc/python_api/rst/info_gotcha.rst

@@ -272,7 +272,7 @@ Editing is where the three data types vary most.
 - Polygons are very limited for editing,
   changing materials and options like smooth works, but for anything else
   they are too inflexible and are only intended for storage.
-- Tessfaces should not be used for editing geometry because doing so will cause existing n-gons to be tessellated.
+- Loop-triangles should not be used for editing geometry because doing so will cause existing n-gons to be tessellated.
 - BMesh-faces are by far the best way to manipulate geometry.
 
 
@@ -283,7 +283,7 @@ All three data types can be used for exporting,
 the choice mostly depends on whether the target format supports n-gons or not.
 
 - Polygons are the most direct and efficient way to export providing they convert into the output format easily enough.
-- Tessfaces work well for exporting to formats which don't support n-gons,
+- Loop-triangles work well for exporting to formats which don't support n-gons,
   in fact this is the only place where their use is encouraged.
 - BMesh-Faces can work for exporting too but may not be necessary if polygons can be used
   since using BMesh gives some overhead because it's not the native storage format in Object-Mode.

intern/cycles/blender/addon/properties.py

@@ -1681,9 +1681,13 @@ class CyclesPreferences(bpy.types.AddonPreferences):
 
         import _cycles
         has_peer_memory = 0
+        has_rt_api_support = False
         for device in _cycles.available_devices(compute_device_type):
             if device[3] and self.find_existing_device_entry(device).use:
                 has_peer_memory += 1
+            if device[4] and self.find_existing_device_entry(device).use:
+                has_rt_api_support = True
+
         if has_peer_memory > 1:
             row = layout.row()
             row.use_property_split = True
@@ -1700,13 +1704,14 @@ class CyclesPreferences(bpy.types.AddonPreferences):
 
             # MetalRT only works on Apple Silicon and Navi2.
             is_arm64 = platform.machine() == 'arm64'
-            if is_arm64 or is_navi_2:
+            if is_arm64 or (is_navi_2 and has_rt_api_support):
                 col = layout.column()
                 col.use_property_split = True
                 # Kernel specialization is only supported on Apple Silicon
                 if is_arm64:
                     col.prop(self, "kernel_optimization_level")
-                col.prop(self, "use_metalrt")
+                if has_rt_api_support:
+                    col.prop(self, "use_metalrt")
 
         if compute_device_type == 'HIP':
             has_cuda, has_optix, has_hip, has_metal, has_oneapi, has_hiprt = _cycles.get_device_types()

intern/cycles/blender/python.cpp

@@ -410,11 +410,12 @@ static PyObject *available_devices_func(PyObject * /*self*/, PyObject *args)
   for (size_t i = 0; i < devices.size(); i++) {
     DeviceInfo &device = devices[i];
     string type_name = Device::string_from_type(device.type);
-    PyObject *device_tuple = PyTuple_New(4);
+    PyObject *device_tuple = PyTuple_New(5);
     PyTuple_SET_ITEM(device_tuple, 0, pyunicode_from_string(device.description.c_str()));
     PyTuple_SET_ITEM(device_tuple, 1, pyunicode_from_string(type_name.c_str()));
     PyTuple_SET_ITEM(device_tuple, 2, pyunicode_from_string(device.id.c_str()));
     PyTuple_SET_ITEM(device_tuple, 3, PyBool_FromLong(device.has_peer_memory));
+    PyTuple_SET_ITEM(device_tuple, 4, PyBool_FromLong(device.use_hardware_raytracing));
     PyTuple_SET_ITEM(ret, i, device_tuple);
   }
 

intern/cycles/device/metal/bvh.h

@@ -22,7 +22,9 @@ class BVHMetal : public BVH {
 
   API_AVAILABLE(macos(11.0))
   vector<id<MTLAccelerationStructure>> blas_array;
-  vector<uint32_t> blas_lookup;
+
+  API_AVAILABLE(macos(11.0))
+  vector<id<MTLAccelerationStructure>> unique_blas_array;
 
   bool motion_blur = false;
 

intern/cycles/device/metal/bvh.mm

@@ -132,6 +132,7 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
       geomDescMotion.indexType = MTLIndexTypeUInt32;
       geomDescMotion.triangleCount = num_indices / 3;
       geomDescMotion.intersectionFunctionTableOffset = 0;
+      geomDescMotion.opaque = true;
 
       geomDesc = geomDescMotion;
     }
@@ -146,6 +147,7 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
       geomDescNoMotion.indexType = MTLIndexTypeUInt32;
       geomDescNoMotion.triangleCount = num_indices / 3;
       geomDescNoMotion.intersectionFunctionTableOffset = 0;
+      geomDescNoMotion.opaque = true;
 
       geomDesc = geomDescNoMotion;
     }
@@ -165,6 +167,7 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
       accelDesc.motionEndBorderMode = MTLMotionBorderModeClamp;
       accelDesc.motionKeyframeCount = num_motion_steps;
     }
+    accelDesc.usage |= MTLAccelerationStructureUsageExtendedLimits;
 
     if (!use_fast_trace_bvh) {
       accelDesc.usage |= (MTLAccelerationStructureUsageRefit |
@@ -255,7 +258,8 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
                                Geometry *const geom,
                                bool refit)
 {
-  if (@available(macos 12.0, *)) {
+#  if defined(MAC_OS_VERSION_14_0)
+  if (@available(macos 14.0, *)) {
     /* Build BLAS for hair curves */
     Hair *hair = static_cast<Hair *>(geom);
     if (hair->num_curves() == 0) {
@@ -268,7 +272,6 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
     /*------------------------------------------------*/
 
     const bool use_fast_trace_bvh = (params.bvh_type == BVH_TYPE_STATIC);
-    const size_t num_segments = hair->num_segments();
 
     size_t num_motion_steps = 1;
     Attribute *motion_keys = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
@@ -276,8 +279,6 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
       num_motion_steps = hair->get_motion_steps();
     }
 
-    const size_t num_aabbs = num_segments * num_motion_steps;
-
     MTLResourceOptions storage_mode;
     if (device.hasUnifiedMemory) {
       storage_mode = MTLResourceStorageModeShared;
@@ -286,91 +287,197 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
       storage_mode = MTLResourceStorageModeManaged;
     }
 
-    /* Allocate a GPU buffer for the AABB data and populate it */
+    id<MTLBuffer> cpBuffer = nil;
-    id<MTLBuffer> aabbBuf = [device
+    id<MTLBuffer> radiusBuffer = nil;
-        newBufferWithLength:num_aabbs * sizeof(MTLAxisAlignedBoundingBox)
+    id<MTLBuffer> idxBuffer = nil;
-                    options:storage_mode];
-    MTLAxisAlignedBoundingBox *aabb_data = (MTLAxisAlignedBoundingBox *)[aabbBuf contents];
-
-    /* Get AABBs for each motion step */
-    size_t center_step = (num_motion_steps - 1) / 2;
-    for (size_t step = 0; step < num_motion_steps; ++step) {
-      /* The center step for motion vertices is not stored in the attribute */
-      const float3 *keys = hair->get_curve_keys().data();
-      if (step != center_step) {
-        size_t attr_offset = (step > center_step) ? step - 1 : step;
-        /* Technically this is a float4 array, but sizeof(float3) == sizeof(float4) */
-        keys = motion_keys->data_float3() + attr_offset * hair->get_curve_keys().size();
-      }
-
-      for (size_t j = 0, i = 0; j < hair->num_curves(); ++j) {
-        const Hair::Curve curve = hair->get_curve(j);
-
-        for (int segment = 0; segment < curve.num_segments(); ++segment, ++i) {
-          {
-            BoundBox bounds = BoundBox::empty;
-            curve.bounds_grow(segment, keys, hair->get_curve_radius().data(), bounds);
-
-            const size_t index = step * num_segments + i;
-            aabb_data[index].min = (MTLPackedFloat3 &)bounds.min;
-            aabb_data[index].max = (MTLPackedFloat3 &)bounds.max;
-          }
-        }
-      }
-    }
-
-    if (storage_mode == MTLResourceStorageModeManaged) {
-      [aabbBuf didModifyRange:NSMakeRange(0, aabbBuf.length)];
-    }
-
-#  if 0
-    for (size_t i=0; i<num_aabbs && i < 400; i++) {
-      MTLAxisAlignedBoundingBox& bb = aabb_data[i];
-      printf("  %d:   %.1f,%.1f,%.1f -- %.1f,%.1f,%.1f\n", int(i), bb.min.x, bb.min.y, bb.min.z, bb.max.x, bb.max.y, bb.max.z);
-    }
-#  endif
 
     MTLAccelerationStructureGeometryDescriptor *geomDesc;
     if (motion_blur) {
-      std::vector<MTLMotionKeyframeData *> aabb_ptrs;
+      MTLAccelerationStructureMotionCurveGeometryDescriptor *geomDescCrv =
-      aabb_ptrs.reserve(num_motion_steps);
+          [MTLAccelerationStructureMotionCurveGeometryDescriptor descriptor];
+
+      uint64_t numKeys = hair->num_keys();
+      uint64_t numCurves = hair->num_curves();
+      const array<float> &radiuses = hair->get_curve_radius();
+
+      /* Gather the curve geometry. */
+      std::vector<float3> cpData;
+      std::vector<int> idxData;
+      std::vector<float> radiusData;
+      cpData.reserve(numKeys);
+      radiusData.reserve(numKeys);
+
+      std::vector<int> step_offsets;
       for (size_t step = 0; step < num_motion_steps; ++step) {
-        MTLMotionKeyframeData *k = [MTLMotionKeyframeData data];
+
-        k.buffer = aabbBuf;
+        /* The center step for motion vertices is not stored in the attribute. */
-        k.offset = step * num_segments * sizeof(MTLAxisAlignedBoundingBox);
+        const float3 *keys = hair->get_curve_keys().data();
-        aabb_ptrs.push_back(k);
+        size_t center_step = (num_motion_steps - 1) / 2;
+        if (step != center_step) {
+          size_t attr_offset = (step > center_step) ? step - 1 : step;
+          /* Technically this is a float4 array, but sizeof(float3) == sizeof(float4). */
+          keys = motion_keys->data_float3() + attr_offset * numKeys;
+        }
+
+        step_offsets.push_back(cpData.size());
+
+        for (int c = 0; c < numCurves; ++c) {
+          const Hair::Curve curve = hair->get_curve(c);
+          int segCount = curve.num_segments();
+          int firstKey = curve.first_key;
+          uint64_t idxBase = cpData.size();
+          cpData.push_back(keys[firstKey]);
+          radiusData.push_back(radiuses[firstKey]);
+          for (int s = 0; s < segCount; ++s) {
+            if (step == 0) {
+              idxData.push_back(idxBase + s);
+            }
+            cpData.push_back(keys[firstKey + s]);
+            radiusData.push_back(radiuses[firstKey + s]);
+          }
+          cpData.push_back(keys[firstKey + curve.num_keys - 1]);
+          cpData.push_back(keys[firstKey + curve.num_keys - 1]);
+          radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]);
+          radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]);
+        }
       }
 
-      MTLAccelerationStructureMotionBoundingBoxGeometryDescriptor *geomDescMotion =
+      /* Allocate and populate MTLBuffers for geometry. */
-          [MTLAccelerationStructureMotionBoundingBoxGeometryDescriptor descriptor];
+      idxBuffer = [device newBufferWithBytes:idxData.data()
-      geomDescMotion.boundingBoxBuffers = [NSArray arrayWithObjects:aabb_ptrs.data()
+                                      length:idxData.size() * sizeof(int)
-                                                              count:aabb_ptrs.size()];
+                                     options:storage_mode];
-      geomDescMotion.boundingBoxCount = num_segments;
+
-      geomDescMotion.boundingBoxStride = sizeof(aabb_data[0]);
+      cpBuffer = [device newBufferWithBytes:cpData.data()
-      geomDescMotion.intersectionFunctionTableOffset = 1;
+                                     length:cpData.size() * sizeof(float3)
+                                    options:storage_mode];
+
+      radiusBuffer = [device newBufferWithBytes:radiusData.data()
+                                         length:radiusData.size() * sizeof(float)
+                                        options:storage_mode];
+
+      std::vector<MTLMotionKeyframeData *> cp_ptrs;
+      std::vector<MTLMotionKeyframeData *> radius_ptrs;
+      cp_ptrs.reserve(num_motion_steps);
+      radius_ptrs.reserve(num_motion_steps);
+
+      for (size_t step = 0; step < num_motion_steps; ++step) {
+        MTLMotionKeyframeData *k = [MTLMotionKeyframeData data];
+        k.buffer = cpBuffer;
+        k.offset = step_offsets[step] * sizeof(float3);
+        cp_ptrs.push_back(k);
+
+        k = [MTLMotionKeyframeData data];
+        k.buffer = radiusBuffer;
+        k.offset = step_offsets[step] * sizeof(float);
+        radius_ptrs.push_back(k);
+      }
+
+      if (storage_mode == MTLResourceStorageModeManaged) {
+        [cpBuffer didModifyRange:NSMakeRange(0, cpBuffer.length)];
+        [idxBuffer didModifyRange:NSMakeRange(0, idxBuffer.length)];
+        [radiusBuffer didModifyRange:NSMakeRange(0, radiusBuffer.length)];
+      }
+
+      geomDescCrv.controlPointBuffers = [NSArray arrayWithObjects:cp_ptrs.data()
+                                                            count:cp_ptrs.size()];
+      geomDescCrv.radiusBuffers = [NSArray arrayWithObjects:radius_ptrs.data()
+                                                      count:radius_ptrs.size()];
+
+      geomDescCrv.controlPointCount = cpData.size();
+      geomDescCrv.controlPointStride = sizeof(float3);
+      geomDescCrv.controlPointFormat = MTLAttributeFormatFloat3;
+      geomDescCrv.radiusStride = sizeof(float);
+      geomDescCrv.radiusFormat = MTLAttributeFormatFloat;
+      geomDescCrv.segmentCount = idxData.size();
+      geomDescCrv.segmentControlPointCount = 4;
+      geomDescCrv.curveType = (hair->curve_shape == CURVE_RIBBON) ? MTLCurveTypeFlat :
+                                                                    MTLCurveTypeRound;
+      geomDescCrv.curveBasis = MTLCurveBasisCatmullRom;
+      geomDescCrv.curveEndCaps = MTLCurveEndCapsDisk;
+      geomDescCrv.indexType = MTLIndexTypeUInt32;
+      geomDescCrv.indexBuffer = idxBuffer;
+      geomDescCrv.intersectionFunctionTableOffset = 1;
 
       /* Force a single any-hit call, so shadow record-all behavior works correctly */
       /* (Match optix behavior: unsigned int build_flags =
        * OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;) */
-      geomDescMotion.allowDuplicateIntersectionFunctionInvocation = false;
+      geomDescCrv.allowDuplicateIntersectionFunctionInvocation = false;
-      geomDescMotion.opaque = true;
+      geomDescCrv.opaque = true;
-      geomDesc = geomDescMotion;
+      geomDesc = geomDescCrv;
     }
     else {
-      MTLAccelerationStructureBoundingBoxGeometryDescriptor *geomDescNoMotion =
+      MTLAccelerationStructureCurveGeometryDescriptor *geomDescCrv =
-          [MTLAccelerationStructureBoundingBoxGeometryDescriptor descriptor];
+          [MTLAccelerationStructureCurveGeometryDescriptor descriptor];
-      geomDescNoMotion.boundingBoxBuffer = aabbBuf;
+
-      geomDescNoMotion.boundingBoxBufferOffset = 0;
+      uint64_t numKeys = hair->num_keys();
-      geomDescNoMotion.boundingBoxCount = int(num_aabbs);
+      uint64_t numCurves = hair->num_curves();
-      geomDescNoMotion.boundingBoxStride = sizeof(aabb_data[0]);
+      const array<float> &radiuses = hair->get_curve_radius();
-      geomDescNoMotion.intersectionFunctionTableOffset = 1;
+
+      /* Gather the curve geometry. */
+      std::vector<float3> cpData;
+      std::vector<int> idxData;
+      std::vector<float> radiusData;
+      cpData.reserve(numKeys);
+      radiusData.reserve(numKeys);
+      auto keys = hair->get_curve_keys();
+      for (int c = 0; c < numCurves; ++c) {
+        const Hair::Curve curve = hair->get_curve(c);
+        int segCount = curve.num_segments();
+        int firstKey = curve.first_key;
+        radiusData.push_back(radiuses[firstKey]);
+        uint64_t idxBase = cpData.size();
+        cpData.push_back(keys[firstKey]);
+        for (int s = 0; s < segCount; ++s) {
+          idxData.push_back(idxBase + s);
+          cpData.push_back(keys[firstKey + s]);
+          radiusData.push_back(radiuses[firstKey + s]);
+        }
+        cpData.push_back(keys[firstKey + curve.num_keys - 1]);
+        cpData.push_back(keys[firstKey + curve.num_keys - 1]);
+        radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]);
+        radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]);
+      }
+
+      /* Allocate and populate MTLBuffers for geometry. */
+      idxBuffer = [device newBufferWithBytes:idxData.data()
+                                      length:idxData.size() * sizeof(int)
+                                     options:storage_mode];
+
+      cpBuffer = [device newBufferWithBytes:cpData.data()
+                                     length:cpData.size() * sizeof(float3)
+                                    options:storage_mode];
+
+      radiusBuffer = [device newBufferWithBytes:radiusData.data()
+                                         length:radiusData.size() * sizeof(float)
+                                        options:storage_mode];
+
+      if (storage_mode == MTLResourceStorageModeManaged) {
+        [cpBuffer didModifyRange:NSMakeRange(0, cpBuffer.length)];
+        [idxBuffer didModifyRange:NSMakeRange(0, idxBuffer.length)];
+        [radiusBuffer didModifyRange:NSMakeRange(0, radiusBuffer.length)];
+      }
+      geomDescCrv.controlPointBuffer = cpBuffer;
+      geomDescCrv.radiusBuffer = radiusBuffer;
+      geomDescCrv.controlPointCount = cpData.size();
+      geomDescCrv.controlPointStride = sizeof(float3);
+      geomDescCrv.controlPointFormat = MTLAttributeFormatFloat3;
+      geomDescCrv.controlPointBufferOffset = 0;
+      geomDescCrv.segmentCount = idxData.size();
+      geomDescCrv.segmentControlPointCount = 4;
+      geomDescCrv.curveType = (hair->curve_shape == CURVE_RIBBON) ? MTLCurveTypeFlat :
+                                                                    MTLCurveTypeRound;
+      geomDescCrv.curveBasis = MTLCurveBasisCatmullRom;
+      geomDescCrv.curveEndCaps = MTLCurveEndCapsDisk;
+      geomDescCrv.indexType = MTLIndexTypeUInt32;
+      geomDescCrv.indexBuffer = idxBuffer;
+      geomDescCrv.intersectionFunctionTableOffset = 1;
 
       /* Force a single any-hit call, so shadow record-all behavior works correctly */
       /* (Match optix behavior: unsigned int build_flags =
        * OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;) */
-      geomDescNoMotion.allowDuplicateIntersectionFunctionInvocation = false;
+      geomDescCrv.allowDuplicateIntersectionFunctionInvocation = false;
-      geomDescNoMotion.opaque = true;
+      geomDescCrv.opaque = true;
-      geomDesc = geomDescNoMotion;
+      geomDesc = geomDescCrv;
     }
 
     MTLPrimitiveAccelerationStructureDescriptor *accelDesc =
@@ -389,6 +496,7 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
       accelDesc.usage |= (MTLAccelerationStructureUsageRefit |
                           MTLAccelerationStructureUsagePreferFastBuild);
     }
+    accelDesc.usage |= MTLAccelerationStructureUsageExtendedLimits;
 
     MTLAccelerationStructureSizes accelSizes = [device
         accelerationStructureSizesWithDescriptor:accelDesc];
@@ -423,10 +531,11 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
     [accelCommands addCompletedHandler:^(id<MTLCommandBuffer> /*command_buffer*/) {
       /* free temp resources */
       [scratchBuf release];
-      [aabbBuf release];
+      [cpBuffer release];
+      [radiusBuffer release];
+      [idxBuffer release];
 
       if (use_fast_trace_bvh) {
-        /* Compact the accel structure */
         uint64_t compressed_size = *(uint64_t *)sizeBuf.contents;
 
         dispatch_async(dispatch_get_global_queue(DISPATCH_QUEUE_PRIORITY_DEFAULT, 0), ^{
@@ -461,8 +570,10 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
 
     accel_struct_building = true;
     [accelCommands commit];
+
     return true;
   }
+#  endif /* MAC_OS_VERSION_14_0 */
   return false;
 }
 
@@ -605,10 +716,11 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
     if (motion_blur) {
       accelDesc.motionStartTime = 0.0f;
       accelDesc.motionEndTime = 1.0f;
-      accelDesc.motionStartBorderMode = MTLMotionBorderModeVanish;
+      //      accelDesc.motionStartBorderMode = MTLMotionBorderModeVanish;
-      accelDesc.motionEndBorderMode = MTLMotionBorderModeVanish;
+      //      accelDesc.motionEndBorderMode = MTLMotionBorderModeVanish;
       accelDesc.motionKeyframeCount = num_motion_steps;
     }
+    accelDesc.usage |= MTLAccelerationStructureUsageExtendedLimits;
 
     if (!use_fast_trace_bvh) {
       accelDesc.usage |= (MTLAccelerationStructureUsageRefit |
@@ -756,10 +868,11 @@ bool BVHMetal::build_TLAS(Progress &progress,
     uint32_t num_instances = 0;
     uint32_t num_motion_transforms = 0;
     for (Object *ob : objects) {
-      /* Skip non-traceable objects */
+      num_instances++;
+
+      /* Skip motion for non-traceable objects */
       if (!ob->is_traceable())
         continue;
-      num_instances++;
 
       if (ob->use_motion()) {
         num_motion_transforms += max((size_t)1, ob->get_motion().size());
@@ -829,28 +942,40 @@ bool BVHMetal::build_TLAS(Progress &progress,
     uint32_t instance_index = 0;
     uint32_t motion_transform_index = 0;
 
-    // allocate look up buffer for wost case scenario
+    blas_array.clear();
-    uint64_t count = objects.size();
+    blas_array.reserve(num_instances);
-    blas_lookup.resize(count);
 
     for (Object *ob : objects) {
       /* Skip non-traceable objects */
-      if (!ob->is_traceable())
-        continue;
-
       Geometry const *geom = ob->get_geometry();
-
       BVHMetal const *blas = static_cast<BVHMetal const *>(geom->bvh);
+      if (!blas || !blas->accel_struct) {
+        /* Place a degenerate instance, to ensure [[instance_id]] equals ob->get_device_index()
+         * in our intersection functions */
+        if (motion_blur) {
+          MTLAccelerationStructureMotionInstanceDescriptor *instances =
+              (MTLAccelerationStructureMotionInstanceDescriptor *)[instanceBuf contents];
+          MTLAccelerationStructureMotionInstanceDescriptor &desc = instances[instance_index++];
+          memset(&desc, 0x00, sizeof(desc));
+        }
+        else {
+          MTLAccelerationStructureUserIDInstanceDescriptor *instances =
+              (MTLAccelerationStructureUserIDInstanceDescriptor *)[instanceBuf contents];
+          MTLAccelerationStructureUserIDInstanceDescriptor &desc = instances[instance_index++];
+          memset(&desc, 0x00, sizeof(desc));
+        }
+        blas_array.push_back(nil);
+        continue;
+      }
+      blas_array.push_back(blas->accel_struct);
+
       uint32_t accel_struct_index = get_blas_index(blas);
 
       /* Add some of the object visibility bits to the mask.
        * __prim_visibility contains the combined visibility bits of all instances, so is not
        * reliable if they differ between instances.
-       *
-       * METAL_WIP: OptiX visibility mask can only contain 8 bits, so have to trade-off here
-       * and select just a few important ones.
        */
-      uint32_t mask = ob->visibility_for_tracing() & 0xFF;
+      uint32_t mask = ob->visibility_for_tracing();
 
       /* Have to have at least one bit in the mask, or else instance would always be culled. */
       if (0 == mask) {
@@ -858,11 +983,25 @@ bool BVHMetal::build_TLAS(Progress &progress,
       }
 
       /* Set user instance ID to object index */
-      int object_index = ob->get_device_index();
+      uint32_t primitive_offset = 0;
-      uint32_t user_id = uint32_t(object_index);
       int currIndex = instance_index++;
-      assert(user_id < blas_lookup.size());
+
-      blas_lookup[user_id] = accel_struct_index;
+      if (geom->geometry_type == Geometry::HAIR) {
+        /* Build BLAS for curve primitives. */
+        Hair *const hair = static_cast<Hair *const>(const_cast<Geometry *>(geom));
+        primitive_offset = uint32_t(hair->curve_segment_offset);
+      }
+      else if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
+        /* Build BLAS for triangle primitives. */
+        Mesh *const mesh = static_cast<Mesh *const>(const_cast<Geometry *>(geom));
+        primitive_offset = uint32_t(mesh->prim_offset);
+      }
+      else if (geom->geometry_type == Geometry::POINTCLOUD) {
+        /* Build BLAS for points primitives. */
+        PointCloud *const pointcloud = static_cast<PointCloud *const>(
+            const_cast<Geometry *>(geom));
+        primitive_offset = uint32_t(pointcloud->prim_offset);
+      }
 
       /* Bake into the appropriate descriptor */
       if (motion_blur) {
@@ -871,7 +1010,7 @@ bool BVHMetal::build_TLAS(Progress &progress,
         MTLAccelerationStructureMotionInstanceDescriptor &desc = instances[currIndex];
 
         desc.accelerationStructureIndex = accel_struct_index;
-        desc.userID = user_id;
+        desc.userID = primitive_offset;
         desc.mask = mask;
         desc.motionStartTime = 0.0f;
         desc.motionEndTime = 1.0f;
@@ -917,9 +1056,10 @@ bool BVHMetal::build_TLAS(Progress &progress,
         MTLAccelerationStructureUserIDInstanceDescriptor &desc = instances[currIndex];
 
         desc.accelerationStructureIndex = accel_struct_index;
-        desc.userID = user_id;
+        desc.userID = primitive_offset;
         desc.mask = mask;
         desc.intersectionFunctionTableOffset = 0;
+        desc.options = MTLAccelerationStructureInstanceOptionOpaque;
 
         float *t = (float *)&desc.transformationMatrix;
         if (ob->get_geometry()->is_instanced()) {
@@ -959,6 +1099,7 @@ bool BVHMetal::build_TLAS(Progress &progress,
       accelDesc.motionTransformCount = num_motion_transforms;
     }
 
+    accelDesc.usage |= MTLAccelerationStructureUsageExtendedLimits;
     if (!use_fast_trace_bvh) {
       accelDesc.usage |= (MTLAccelerationStructureUsageRefit |
                           MTLAccelerationStructureUsagePreferFastBuild);
@@ -1001,11 +1142,13 @@ bool BVHMetal::build_TLAS(Progress &progress,
 
     /* Cache top and bottom-level acceleration structs */
     accel_struct = accel;
-    blas_array.clear();
+
-    blas_array.reserve(all_blas.count);
+    unique_blas_array.clear();
-    for (id<MTLAccelerationStructure> blas in all_blas) {
+    unique_blas_array.reserve(all_blas.count);
-      blas_array.push_back(blas);
+    [all_blas enumerateObjectsUsingBlock:^(
-    }
+                  id<MTLAccelerationStructure> blas, NSUInteger, BOOL *) {
+      unique_blas_array.push_back(blas);
+    }];
 
     return true;
   }

intern/cycles/device/metal/device.mm

@@ -62,12 +62,17 @@ void device_metal_info(vector<DeviceInfo> &devices)
     info.has_light_tree = vendor != METAL_GPU_AMD;
     info.has_mnee = vendor != METAL_GPU_AMD;
 
-    info.use_hardware_raytracing = vendor != METAL_GPU_INTEL;
+    info.use_hardware_raytracing = false;
-    if (info.use_hardware_raytracing) {
+
-      if (@available(macos 11.0, *)) {
+    /* MetalRT now uses features exposed in Xcode versions corresponding to macOS 14+, so don't
+     * expose it in builds from older Xcode versions. */
+#  if defined(MAC_OS_VERSION_14_0)
+    if (vendor != METAL_GPU_INTEL) {
+      if (@available(macos 14.0, *)) {
         info.use_hardware_raytracing = device.supportsRaytracing;
       }
     }
+#  endif
 
     devices.push_back(info);
     device_index++;

intern/cycles/device/metal/device_impl.h

@@ -82,7 +82,6 @@ class MetalDevice : public Device {
   /* BLAS encoding & lookup */
   id<MTLArgumentEncoder> mtlBlasArgEncoder = nil;
   id<MTLBuffer> blas_buffer = nil;
-  id<MTLBuffer> blas_lookup_buffer = nil;
 
   bool use_metalrt = false;
   MetalPipelineType kernel_specialization_level = PSO_GENERIC;

intern/cycles/device/metal/device_impl.mm

@@ -81,7 +81,7 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
   mtlDevice = usable_devices[mtlDevId];
   device_vendor = MetalInfo::get_device_vendor(mtlDevice);
   assert(device_vendor != METAL_GPU_UNKNOWN);
-  metal_printf("Creating new Cycles device for Metal: %s\n", info.description.c_str());
+  metal_printf("Creating new Cycles Metal device: %s\n", info.description.c_str());
 
   /* determine default storage mode based on whether UMA is supported */
 
@@ -549,9 +549,14 @@ void MetalDevice::compile_and_load(int device_id, MetalPipelineType pso_type)
 #  endif
 
   options.fastMathEnabled = YES;
-  if (@available(macOS 12.0, *)) {
+  if (@available(macos 12.0, *)) {
     options.languageVersion = MTLLanguageVersion2_4;
   }
+#  if defined(MAC_OS_VERSION_14_0)
+  if (@available(macos 14.0, *)) {
+    options.languageVersion = MTLLanguageVersion3_1;
+  }
+#  endif
 
   if (getenv("CYCLES_METAL_PROFILING") || getenv("CYCLES_METAL_DEBUG")) {
     path_write_text(path_cache_get(string_printf("%s.metal", kernel_type_as_string(pso_type))),
@@ -1372,24 +1377,14 @@ void MetalDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         stats.mem_alloc(blas_buffer.allocatedSize);
 
         for (uint64_t i = 0; i < count; ++i) {
-          [mtlBlasArgEncoder setArgumentBuffer:blas_buffer
+          if (bvhMetalRT->blas_array[i]) {
-                                        offset:i * mtlBlasArgEncoder.encodedLength];
+            [mtlBlasArgEncoder setArgumentBuffer:blas_buffer
-          [mtlBlasArgEncoder setAccelerationStructure:bvhMetalRT->blas_array[i] atIndex:0];
+                                          offset:i * mtlBlasArgEncoder.encodedLength];
+            [mtlBlasArgEncoder setAccelerationStructure:bvhMetalRT->blas_array[i] atIndex:0];
+          }
         }
-
-        count = bvhMetalRT->blas_lookup.size();
-        bufferSize = sizeof(uint32_t) * count;
-        blas_lookup_buffer = [mtlDevice newBufferWithLength:bufferSize
-                                                    options:default_storage_mode];
-        stats.mem_alloc(blas_lookup_buffer.allocatedSize);
-
-        memcpy([blas_lookup_buffer contents],
-               bvhMetalRT -> blas_lookup.data(),
-               blas_lookup_buffer.allocatedSize);
-
         if (default_storage_mode == MTLResourceStorageModeManaged) {
           [blas_buffer didModifyRange:NSMakeRange(0, blas_buffer.length)];
-          [blas_lookup_buffer didModifyRange:NSMakeRange(0, blas_lookup_buffer.length)];
         }
       }
     }

intern/cycles/device/metal/kernel.h

@@ -22,10 +22,8 @@ enum {
   METALRT_FUNC_LOCAL_BOX,
   METALRT_FUNC_LOCAL_TRI_PRIM,
   METALRT_FUNC_LOCAL_BOX_PRIM,
-  METALRT_FUNC_CURVE_RIBBON,
+  METALRT_FUNC_CURVE,
-  METALRT_FUNC_CURVE_RIBBON_SHADOW,
+  METALRT_FUNC_CURVE_SHADOW,
-  METALRT_FUNC_CURVE_ALL,
-  METALRT_FUNC_CURVE_ALL_SHADOW,
   METALRT_FUNC_POINT,
   METALRT_FUNC_POINT_SHADOW,
   METALRT_FUNC_NUM

intern/cycles/device/metal/kernel.mm

@@ -493,10 +493,8 @@ void MetalKernelPipeline::compile()
           "__anyhit__cycles_metalrt_local_hit_box",
           "__anyhit__cycles_metalrt_local_hit_tri_prim",
           "__anyhit__cycles_metalrt_local_hit_box_prim",
-          "__intersection__curve_ribbon",
+          "__intersection__curve",
-          "__intersection__curve_ribbon_shadow",
+          "__intersection__curve_shadow",
-          "__intersection__curve_all",
-          "__intersection__curve_all_shadow",
           "__intersection__point",
           "__intersection__point_shadow",
       };
@@ -540,17 +538,8 @@ void MetalKernelPipeline::compile()
     id<MTLFunction> point_intersect_default = nil;
     id<MTLFunction> point_intersect_shadow = nil;
     if (kernel_features & KERNEL_FEATURE_HAIR) {
-      /* Add curve intersection programs. */
+      curve_intersect_default = rt_intersection_function[METALRT_FUNC_CURVE];
-      if (kernel_features & KERNEL_FEATURE_HAIR_THICK) {
+      curve_intersect_shadow = rt_intersection_function[METALRT_FUNC_CURVE_SHADOW];
-        /* Slower programs for thick hair since that also slows down ribbons.
-         * Ideally this should not be needed. */
-        curve_intersect_default = rt_intersection_function[METALRT_FUNC_CURVE_ALL];
-        curve_intersect_shadow = rt_intersection_function[METALRT_FUNC_CURVE_ALL_SHADOW];
-      }
-      else {
-        curve_intersect_default = rt_intersection_function[METALRT_FUNC_CURVE_RIBBON];
-        curve_intersect_shadow = rt_intersection_function[METALRT_FUNC_CURVE_RIBBON_SHADOW];
-      }
     }
     if (kernel_features & KERNEL_FEATURE_POINTCLOUD) {
       point_intersect_default = rt_intersection_function[METALRT_FUNC_POINT];
@@ -585,8 +574,8 @@ void MetalKernelPipeline::compile()
                          rt_intersection_function[METALRT_FUNC_LOCAL_BOX_PRIM],
                          nil];
 
-    NSMutableSet *unique_functions = [NSMutableSet
+    NSMutableSet *unique_functions = [[NSMutableSet alloc] init];
-        setWithArray:table_functions[METALRT_TABLE_DEFAULT]];
+    [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_DEFAULT]];
     [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_SHADOW]];
     [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL]];
     [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL_PRIM]];

intern/cycles/device/metal/queue.mm

@@ -490,9 +490,6 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
         [metal_device_->mtlAncillaryArgEncoder setBuffer:metal_device_->blas_buffer
                                                   offset:0
                                                  atIndex:8];
-        [metal_device_->mtlAncillaryArgEncoder setBuffer:metal_device_->blas_lookup_buffer
-                                                  offset:0
-                                                 atIndex:9];
       }
 
       for (int table = 0; table < METALRT_TABLE_NUM; table++) {
@@ -546,10 +543,8 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
         [mtlComputeCommandEncoder useResource:bvhMetalRT->accel_struct usage:MTLResourceUsageRead];
         [mtlComputeCommandEncoder useResource:metal_device_->blas_buffer
                                         usage:MTLResourceUsageRead];
-        [mtlComputeCommandEncoder useResource:metal_device_->blas_lookup_buffer
+        [mtlComputeCommandEncoder useResources:bvhMetalRT->unique_blas_array.data()
-                                        usage:MTLResourceUsageRead];
+                                         count:bvhMetalRT->unique_blas_array.size()
-        [mtlComputeCommandEncoder useResources:bvhMetalRT->blas_array.data()
-                                         count:bvhMetalRT->blas_array.size()
                                          usage:MTLResourceUsageRead];
       }
     }

intern/cycles/hydra/material.cpp

@@ -139,7 +139,7 @@ class UsdToCycles {
           {TfToken("diffuseColor"), ustring("base_color")},
           {TfToken("emissiveColor"), ustring("emission")},
           {TfToken("specularColor"), ustring("specular")},
-          {TfToken("clearcoatRoughness"), ustring("clearcoat_roughness")},
+          {TfToken("clearcoatRoughness"), ustring("coat_roughness")},
           {TfToken("opacity"), ustring("alpha")},
           // opacityThreshold
           // occlusion

intern/cycles/kernel/closure/bsdf.h

@@ -153,7 +153,6 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_MICROFACET_GGX_ID:
-    case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_GLASS_ID:
       label = bsdf_microfacet_ggx_sample(
@@ -284,7 +283,6 @@ ccl_device_inline void bsdf_roughness_eta(const KernelGlobals kg,
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_MICROFACET_GGX_ID:
-    case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_GLASS_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
@@ -385,7 +383,6 @@ ccl_device_inline int bsdf_label(const KernelGlobals kg,
       label = LABEL_TRANSMIT | LABEL_TRANSPARENT;
       break;
     case CLOSURE_BSDF_MICROFACET_GGX_ID:
-    case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID:
@@ -484,7 +481,6 @@ ccl_device_inline
       eval = bsdf_transparent_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_MICROFACET_GGX_ID:
-    case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_GLASS_ID:
       eval = bsdf_microfacet_ggx_eval(sc, Ng, sd->wi, wo, pdf);
@@ -555,7 +551,6 @@ ccl_device void bsdf_blur(KernelGlobals kg, ccl_private ShaderClosure *sc, float
 #if defined(__SVM__) || defined(__OSL__)
   switch (sc->type) {
     case CLOSURE_BSDF_MICROFACET_GGX_ID:
-    case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_GLASS_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:

intern/cycles/kernel/closure/bsdf_microfacet.h

@@ -377,17 +377,6 @@ ccl_device Spectrum bsdf_microfacet_estimate_albedo(KernelGlobals kg,
   return albedo;
 }
 
-/* Generalized Trowbridge-Reitz for clearcoat. */
-ccl_device_forceinline float bsdf_clearcoat_D(float alpha2, float cos_NH)
-{
-  if (alpha2 >= 1.0f) {
-    return M_1_PI_F;
-  }
-
-  const float t = 1.0f + (alpha2 - 1.0f) * cos_NH * cos_NH;
-  return (alpha2 - 1.0f) / (M_PI_F * logf(alpha2) * t);
-}
-
 /* Smith shadowing-masking term, here in the non-separable form.
  * For details, see:
  * Understanding the Masking-Shadowing Function in Microfacet-Based BRDFs.
@@ -529,18 +518,7 @@ ccl_device Spectrum bsdf_microfacet_eval(ccl_private const ShaderClosure *sc,
    * harder to compute. */
   if (alpha_x == alpha_y || is_transmission) { /* Isotropic. */
     float alpha2 = alpha_x * alpha_y;
-
+    D = bsdf_D<m_type>(alpha2, cos_NH);
-    if (bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID) {
-      D = bsdf_clearcoat_D(alpha2, cos_NH);
-
-      /* The masking-shadowing term for clearcoat has a fixed alpha of 0.25
-       * => alpha2 = 0.25 * 0.25 */
-      alpha2 = 0.0625f;
-    }
-    else {
-      D = bsdf_D<m_type>(alpha2, cos_NH);
-    }
-
     lambdaI = bsdf_lambda<m_type>(alpha2, cos_NI);
     lambdaO = bsdf_lambda<m_type>(alpha2, cos_NO);
   }
@@ -687,17 +665,7 @@ ccl_device int bsdf_microfacet_sample(ccl_private const ShaderClosure *sc,
       const float cos_NH = dot(N, H);
       const float cos_NO = dot(N, *wo);
 
-      if (bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID) {
+      D = bsdf_D<m_type>(alpha2, cos_NH);
-        D = bsdf_clearcoat_D(alpha2, cos_NH);
-
-        /* The masking-shadowing term for clearcoat has a fixed alpha of 0.25
-         * => alpha2 = 0.25 * 0.25 */
-        alpha2 = 0.0625f;
-      }
-      else {
-        D = bsdf_D<m_type>(alpha2, cos_NH);
-      }
-
       lambdaO = bsdf_lambda<m_type>(alpha2, cos_NO);
       lambdaI = bsdf_lambda<m_type>(alpha2, cos_NI);
     }
@@ -831,6 +799,14 @@ ccl_device void bsdf_microfacet_setup_fresnel_constant(KernelGlobals kg,
   microfacet_ggx_preserve_energy(kg, bsdf, sd, color);
 }
 
+ccl_device void bsdf_microfacet_setup_fresnel_dielectric(KernelGlobals kg,
+                                                         ccl_private MicrofacetBsdf *bsdf,
+                                                         ccl_private const ShaderData *sd)
+{
+  bsdf->fresnel_type = MicrofacetFresnel::DIELECTRIC;
+  bsdf->sample_weight *= average(bsdf_microfacet_estimate_albedo(kg, sd, bsdf, true, true));
+}
+
 /* GGX microfacet with Smith shadow-masking from:
  *
  * Microfacet Models for Refraction through Rough Surfaces
@@ -856,21 +832,6 @@ ccl_device int bsdf_microfacet_ggx_setup(ccl_private MicrofacetBsdf *bsdf)
   return SD_BSDF | bsdf_microfacet_eval_flag(bsdf);
 }
 
-ccl_device int bsdf_microfacet_ggx_clearcoat_setup(KernelGlobals kg,
-                                                   ccl_private MicrofacetBsdf *bsdf,
-                                                   ccl_private const ShaderData *sd)
-{
-  bsdf->alpha_x = saturatef(bsdf->alpha_x);
-  bsdf->alpha_y = bsdf->alpha_x;
-
-  bsdf->fresnel_type = MicrofacetFresnel::DIELECTRIC;
-  bsdf->energy_scale = 1.0f;
-  bsdf->type = CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID;
-  bsdf->sample_weight *= average(bsdf_microfacet_estimate_albedo(kg, sd, bsdf, true, true));
-
-  return SD_BSDF | bsdf_microfacet_eval_flag(bsdf);
-}
-
 ccl_device int bsdf_microfacet_ggx_refraction_setup(ccl_private MicrofacetBsdf *bsdf)
 {
   bsdf->alpha_x = saturatef(bsdf->alpha_x);

intern/cycles/kernel/closure/bssrdf.h

@@ -11,8 +11,11 @@ typedef struct Bssrdf {
 
   Spectrum radius;
   Spectrum albedo;
-  float roughness;
   float anisotropy;
+
+  /* Parameters for refractive entry bounce. */
+  float ior;
+  float alpha;
 } Bssrdf;
 
 static_assert(sizeof(ShaderClosure) >= sizeof(Bssrdf), "Bssrdf is too large!");
@@ -54,9 +57,7 @@ ccl_device float bssrdf_dipole_compute_alpha_prime(float rd, float fourthirdA)
   return xmid;
 }
 
-ccl_device void bssrdf_setup_radius(ccl_private Bssrdf *bssrdf,
+ccl_device void bssrdf_setup_radius(ccl_private Bssrdf *bssrdf, const ClosureType type)
-                                    const ClosureType type,
-                                    const float eta)
 {
   if (type == CLOSURE_BSSRDF_BURLEY_ID || type == CLOSURE_BSSRDF_RANDOM_WALK_FIXED_RADIUS_ID) {
     /* Scale mean free path length so it gives similar looking result to older
@@ -65,8 +66,8 @@ ccl_device void bssrdf_setup_radius(ccl_private Bssrdf *bssrdf,
   }
   else {
     /* Adjust radius based on IOR and albedo. */
-    const float inv_eta = 1.0f / eta;
+    const float inv_eta = 1.0f / bssrdf->ior;
-    const float F_dr = inv_eta * (-1.440f * inv_eta + 0.710f) + 0.668f + 0.0636f * eta;
+    const float F_dr = inv_eta * (-1.440f * inv_eta + 0.710f) + 0.668f + 0.0636f * bssrdf->ior;
     const float fourthirdA = (4.0f / 3.0f) * (1.0f + F_dr) /
                              (1.0f - F_dr); /* From Jensen's `Fdr` ratio formula. */
 
@@ -281,17 +282,34 @@ ccl_device_inline ccl_private Bssrdf *bssrdf_alloc(ccl_private ShaderData *sd, S
 
 ccl_device int bssrdf_setup(ccl_private ShaderData *sd,
                             ccl_private Bssrdf *bssrdf,
-                            ClosureType type,
+                            int path_flag,
-                            const float ior)
+                            ClosureType type)
 {
+  /* Clamps protecting against bad/extreme and non physical values. */
+  bssrdf->anisotropy = clamp(bssrdf->anisotropy, 0.0f, 0.9f);
+  bssrdf->ior = clamp(bssrdf->ior, 1.01f, 3.8f);
+
   int flag = 0;
 
+  if (type == CLOSURE_BSSRDF_RANDOM_WALK_ID) {
+    /* CLOSURE_BSSRDF_RANDOM_WALK_ID uses a fixed roughness. */
+    bssrdf->alpha = 1.0f;
+  }
+
   /* Verify if the radii are large enough to sample without precision issues. */
   int bssrdf_channels = SPECTRUM_CHANNELS;
   Spectrum diffuse_weight = zero_spectrum();
 
+  /* Fall back to diffuse if the radius is smaller than a quarter pixel. */
+  float min_radius = max(0.25f * sd->dP, BSSRDF_MIN_RADIUS);
+  if (path_flag & PATH_RAY_DIFFUSE_ANCESTOR) {
+    /* Always fall back to diffuse after a diffuse ancestor. Can't see it well then and adds
+     * considerable noise due to probabilities of continuing the path getting lower and lower. */
+    min_radius = FLT_MAX;
+  }
+
   FOREACH_SPECTRUM_CHANNEL (i) {
-    if (GET_SPECTRUM_CHANNEL(bssrdf->radius, i) < BSSRDF_MIN_RADIUS) {
+    if (GET_SPECTRUM_CHANNEL(bssrdf->radius, i) < min_radius) {
       GET_SPECTRUM_CHANNEL(diffuse_weight, i) = GET_SPECTRUM_CHANNEL(bssrdf->weight, i);
       GET_SPECTRUM_CHANNEL(bssrdf->weight, i) = 0.0f;
       GET_SPECTRUM_CHANNEL(bssrdf->radius, i) = 0.0f;
@@ -315,12 +333,12 @@ ccl_device int bssrdf_setup(ccl_private ShaderData *sd,
     bssrdf->type = type;
     bssrdf->sample_weight = fabsf(average(bssrdf->weight)) * bssrdf_channels;
 
-    bssrdf_setup_radius(bssrdf, type, ior);
+    bssrdf_setup_radius(bssrdf, type);
 
     flag |= SD_BSSRDF;
   }
   else {
-    bssrdf->type = type;
+    bssrdf->type = CLOSURE_NONE_ID;
     bssrdf->sample_weight = 0.0f;
   }
 

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

@@ -16,12 +16,6 @@ CCL_NAMESPACE_BEGIN
 struct MetalRTIntersectionPayload {
   RaySelfPrimitives self;
   uint visibility;
-  float u, v;
-  int prim;
-  int type;
-#if defined(__METALRT_MOTION__)
-  float time;
-#endif
 };
 
 struct MetalRTIntersectionLocalPayload {
@@ -37,9 +31,6 @@ struct MetalRTIntersectionLocalPayload {
 struct MetalRTIntersectionShadowPayload {
   RaySelfPrimitives self;
   uint visibility;
-#if defined(__METALRT_MOTION__)
-  float time;
-#endif
   int state;
   float throughput;
   short max_hits;
@@ -48,6 +39,98 @@ struct MetalRTIntersectionShadowPayload {
   bool result;
 };
 
+ccl_device_forceinline bool curve_ribbon_accept(
+    KernelGlobals kg, float u, float t, ccl_private const Ray *ray, int object, int prim, int type)
+{
+  KernelCurve kcurve = kernel_data_fetch(curves, prim);
+
+  int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
+  int k1 = k0 + 1;
+  int ka = max(k0 - 1, kcurve.first_key);
+  int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
+
+  /* We can ignore motion blur here because we don't need the positions, and it doesn't affect the
+   * radius. */
+  float radius[4];
+  radius[0] = kernel_data_fetch(curve_keys, ka).w;
+  radius[1] = kernel_data_fetch(curve_keys, k0).w;
+  radius[2] = kernel_data_fetch(curve_keys, k1).w;
+  radius[3] = kernel_data_fetch(curve_keys, kb).w;
+  const float r = metal::catmull_rom(u, radius[0], radius[1], radius[2], radius[3]);
+
+  /* MPJ TODO: Can we ignore motion and/or object transforms here? Depends on scaling? */
+  float3 ray_P = ray->P;
+  float3 ray_D = ray->D;
+  if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
+    float3 idir;
+#if defined(__METALRT_MOTION__)
+    bvh_instance_motion_push(NULL, object, ray, &ray_P, &ray_D, &idir);
+#else
+    bvh_instance_push(NULL, object, ray, &ray_P, &ray_D, &idir);
+#endif
+  }
+
+  /* ignore self intersections */
+  const float avoidance_factor = 2.0f;
+  return t * len(ray_D) > avoidance_factor * r;
+}
+
+ccl_device_forceinline float curve_ribbon_v(
+    KernelGlobals kg, float u, float t, ccl_private const Ray *ray, int object, int prim, int type)
+{
+#if defined(__METALRT_MOTION__)
+  float time = ray->time;
+#else
+  float time = 0.0f;
+#endif
+
+  const bool is_motion = (type & PRIMITIVE_MOTION);
+
+  KernelCurve kcurve = kernel_data_fetch(curves, prim);
+
+  int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
+  int k1 = k0 + 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) {
+    curve[0] = kernel_data_fetch(curve_keys, ka);
+    curve[1] = kernel_data_fetch(curve_keys, k0);
+    curve[2] = kernel_data_fetch(curve_keys, k1);
+    curve[3] = kernel_data_fetch(curve_keys, kb);
+  }
+  else {
+    motion_curve_keys(kg, object, prim, time, ka, k0, k1, kb, curve);
+  }
+
+  float3 ray_P = ray->P;
+  float3 ray_D = ray->D;
+  if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
+    float3 idir;
+#if defined(__METALRT_MOTION__)
+    bvh_instance_motion_push(NULL, object, ray, &ray_P, &ray_D, &idir);
+#else
+    bvh_instance_push(NULL, object, ray, &ray_P, &ray_D, &idir);
+#endif
+  }
+
+  const float4 P_curve4 = metal::catmull_rom(u, curve[0], curve[1], curve[2], curve[3]);
+  const float r_curve = P_curve4.w;
+
+  float3 P = ray_P + ray_D * t;
+  const float3 P_curve = float4_to_float3(P_curve4);
+
+  const float4 dPdu4 = metal::catmull_rom_derivative(u, curve[0], curve[1], curve[2], curve[3]);
+  const float3 dPdu = float4_to_float3(dPdu4);
+
+  const float3 tangent = normalize(dPdu);
+  const float3 bitangent = normalize(cross(tangent, -ray_D));
+
+  float v = dot(P - P_curve, bitangent) / r_curve;
+  return clamp(v, -1.0, 1.0f);
+}
+
 /* Scene intersection. */
 
 ccl_device_intersect bool scene_intersect(KernelGlobals kg,
@@ -79,41 +162,34 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
 
   metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
   metalrt_intersector_type metalrt_intersect;
+  metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
+  metalrt_intersect.assume_geometry_type(
+      metal::raytracing::geometry_type::triangle |
+      (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
+                                     metal::raytracing::geometry_type::none) |
+      (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
+                                     metal::raytracing::geometry_type::none));
 
-  bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
+  if (visibility & PATH_RAY_SHADOW_OPAQUE) {
-  if (triangle_only) {
+    metalrt_intersect.accept_any_intersection(true);
-    metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
   }
 
   MetalRTIntersectionPayload payload;
   payload.self = ray->self;
-  payload.u = 0.0f;
-  payload.v = 0.0f;
   payload.visibility = visibility;
 
   typename metalrt_intersector_type::result_type intersection;
 
-  uint ray_mask = visibility & 0xFF;
-  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
-    ray_mask = 0xFF;
-    /* No further intersector setup required: Default MetalRT behavior is any-hit. */
-  }
-  else if (visibility & PATH_RAY_SHADOW_OPAQUE) {
-    /* No further intersector setup required: Shadow ray early termination is controlled by the
-     * intersection handler */
-  }
-
 #if defined(__METALRT_MOTION__)
-  payload.time = ray->time;
   intersection = metalrt_intersect.intersect(r,
                                              metal_ancillaries->accel_struct,
-                                             ray_mask,
+                                             visibility,
                                              ray->time,
                                              metal_ancillaries->ift_default,
                                              payload);
 #else
   intersection = metalrt_intersect.intersect(
-      r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload);
+      r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_default, payload);
 #endif
 
   if (intersection.type == intersection_type::none) {
@@ -123,23 +199,71 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
     return false;
   }
 
-  isect->t = intersection.distance;
+  isect->object = intersection.instance_id;
-
-  isect->prim = payload.prim;
-  isect->type = payload.type;
-  isect->object = intersection.user_instance_id;
-
   isect->t = intersection.distance;
   if (intersection.type == intersection_type::triangle) {
+    isect->prim = intersection.primitive_id + intersection.user_instance_id;
+    isect->type = kernel_data_fetch(objects, intersection.instance_id).primitive_type;
     isect->u = intersection.triangle_barycentric_coord.x;
     isect->v = intersection.triangle_barycentric_coord.y;
   }
-  else {
+  else if (kernel_data.bvh.have_curves && intersection.type == intersection_type::curve) {
-    isect->u = payload.u;
+    int prim = intersection.primitive_id + intersection.user_instance_id;
-    isect->v = payload.v;
+    const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
+    isect->prim = segment.prim;
+    isect->type = segment.type;
+    isect->u = intersection.curve_parameter;
+
+    if (segment.type & PRIMITIVE_CURVE_RIBBON) {
+      isect->v = curve_ribbon_v(kg,
+                                intersection.curve_parameter,
+                                intersection.distance,
+                                ray,
+                                intersection.instance_id,
+                                segment.prim,
+                                segment.type);
+    }
+    else {
+      isect->v = 0.0f;
+    }
+  }
+  else if (kernel_data.bvh.have_points && intersection.type == intersection_type::bounding_box) {
+    const int object = intersection.instance_id;
+    const uint prim = intersection.primitive_id + intersection.user_instance_id;
+    const int prim_type = kernel_data_fetch(objects, object).primitive_type;
+
+    if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
+      float3 idir;
+#if defined(__METALRT_MOTION__)
+      bvh_instance_motion_push(NULL, object, ray, &r.origin, &r.direction, &idir);
+#else
+      bvh_instance_push(NULL, object, ray, &r.origin, &r.direction, &idir);
+#endif
+    }
+
+    if (prim_type & PRIMITIVE_POINT) {
+      if (!point_intersect(NULL,
+                           isect,
+                           r.origin,
+                           r.direction,
+                           ray->tmin,
+                           ray->tmax,
+                           object,
+                           prim,
+                           ray->time,
+                           prim_type))
+      {
+        /* Shouldn't get here */
+        kernel_assert(!"Intersection mismatch");
+        isect->t = ray->tmax;
+        isect->type = PRIMITIVE_NONE;
+        return false;
+      }
+      return true;
+    }
   }
 
-  return isect->type != PRIMITIVE_NONE;
+  return true;
 }
 
 #ifdef __BVH_LOCAL__
@@ -198,25 +322,18 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
 #  if defined(__METALRT_MOTION__)
   metalrt_intersector_type metalrt_intersect;
   typename metalrt_intersector_type::result_type intersection;
-
-  metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
-  bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
-  if (triangle_only) {
-    metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
-  }
-
-  intersection = metalrt_intersect.intersect(
-      r, metal_ancillaries->accel_struct, 0xFF, ray->time, metal_ancillaries->ift_local, payload);
 #  else
-
   metalrt_blas_intersector_type metalrt_intersect;
   typename metalrt_blas_intersector_type::result_type intersection;
+#  endif
 
   metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
-  bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
+  metalrt_intersect.assume_geometry_type(
-  if (triangle_only) {
+      metal::raytracing::geometry_type::triangle |
-    metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
+      (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
-  }
+                                     metal::raytracing::geometry_type::none) |
+      (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
+                                     metal::raytracing::geometry_type::none));
 
   // if we know we are going to get max one hit, like for random-sss-walk we can
   // optimize and accept the first hit
@@ -224,8 +341,10 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
     metalrt_intersect.accept_any_intersection(true);
   }
 
-  int blas_index = metal_ancillaries->blas_userID_to_index_lookUp[local_object];
+#  if defined(__METALRT_MOTION__)
-
+  intersection = metalrt_intersect.intersect(
+      r, metal_ancillaries->accel_struct, ~0, ray->time, metal_ancillaries->ift_local, payload);
+#  else
   if (!(kernel_data_fetch(object_flag, local_object) & SD_OBJECT_TRANSFORM_APPLIED)) {
     // transform the ray into object's local space
     Transform itfm = kernel_data_fetch(objects, local_object).itfm;
@@ -235,7 +354,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
 
   intersection = metalrt_intersect.intersect(
       r,
-      metal_ancillaries->blas_accel_structs[blas_index].blas,
+      metal_ancillaries->blas_accel_structs[local_object].blas,
       metal_ancillaries->ift_local_prim,
       payload);
 #  endif
@@ -278,13 +397,13 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
 
   metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
   metalrt_intersector_type metalrt_intersect;
-
   metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
-
+  metalrt_intersect.assume_geometry_type(
-  bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
+      metal::raytracing::geometry_type::triangle |
-  if (triangle_only) {
+      (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
-    metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
+                                     metal::raytracing::geometry_type::none) |
-  }
+      (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
+                                     metal::raytracing::geometry_type::none));
 
   MetalRTIntersectionShadowPayload payload;
   payload.self = ray->self;
@@ -296,24 +415,18 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
   payload.result = false;
   payload.state = state;
 
-  uint ray_mask = visibility & 0xFF;
-  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
-    ray_mask = 0xFF;
-  }
-
   typename metalrt_intersector_type::result_type intersection;
 
 #  if defined(__METALRT_MOTION__)
-  payload.time = ray->time;
   intersection = metalrt_intersect.intersect(r,
                                              metal_ancillaries->accel_struct,
-                                             ray_mask,
+                                             visibility,
                                              ray->time,
                                              metal_ancillaries->ift_shadow,
                                              payload);
 #  else
   intersection = metalrt_intersect.intersect(
-      r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_shadow, payload);
+      r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_shadow, payload);
 #  endif
 
   *num_recorded_hits = payload.num_recorded_hits;
@@ -347,13 +460,13 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
 
   metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
   metalrt_intersector_type metalrt_intersect;
-
   metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
-
+  metalrt_intersect.assume_geometry_type(
-  bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
+      metal::raytracing::geometry_type::triangle |
-  if (triangle_only) {
+      (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
-    metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
+                                     metal::raytracing::geometry_type::none) |
-  }
+      (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
+                                     metal::raytracing::geometry_type::none));
 
   MetalRTIntersectionPayload payload;
   payload.self = ray->self;
@@ -361,43 +474,86 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
 
   typename metalrt_intersector_type::result_type intersection;
 
-  uint ray_mask = visibility & 0xFF;
-  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
-    ray_mask = 0xFF;
-  }
-
 #  if defined(__METALRT_MOTION__)
-  payload.time = ray->time;
   intersection = metalrt_intersect.intersect(r,
                                              metal_ancillaries->accel_struct,
-                                             ray_mask,
+                                             visibility,
                                              ray->time,
                                              metal_ancillaries->ift_default,
                                              payload);
 #  else
   intersection = metalrt_intersect.intersect(
-      r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload);
+      r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_default, payload);
 #  endif
 
   if (intersection.type == intersection_type::none) {
     return false;
   }
-
+  else if (intersection.type == intersection_type::triangle) {
-  isect->prim = payload.prim;
+    isect->prim = intersection.primitive_id + intersection.user_instance_id;
-  isect->type = payload.type;
+    isect->type = kernel_data_fetch(objects, intersection.instance_id).primitive_type;
-  isect->object = intersection.user_instance_id;
-
-  isect->t = intersection.distance;
-  if (intersection.type == intersection_type::triangle) {
     isect->u = intersection.triangle_barycentric_coord.x;
     isect->v = intersection.triangle_barycentric_coord.y;
+    isect->object = intersection.instance_id;
+    isect->t = intersection.distance;
   }
-  else {
+  else if (kernel_data.bvh.have_curves && intersection.type == intersection_type::curve) {
-    isect->u = payload.u;
+    int prim = intersection.primitive_id + intersection.user_instance_id;
-    isect->v = payload.v;
+    const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
+    isect->prim = segment.prim;
+    isect->type = segment.type;
+    isect->u = intersection.curve_parameter;
+
+    if (segment.type & PRIMITIVE_CURVE_RIBBON) {
+      isect->v = curve_ribbon_v(kg,
+                                intersection.curve_parameter,
+                                intersection.distance,
+                                ray,
+                                intersection.instance_id,
+                                segment.prim,
+                                segment.type);
+    }
+    else {
+      isect->v = 0.0f;
+    }
+  }
+  else if (kernel_data.bvh.have_points && intersection.type == intersection_type::bounding_box) {
+    const int object = intersection.instance_id;
+    const uint prim = intersection.primitive_id + intersection.user_instance_id;
+    const int prim_type = kernel_data_fetch(objects, intersection.instance_id).primitive_type;
+
+    isect->object = object;
+
+    if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
+      float3 idir;
+#  if defined(__METALRT_MOTION__)
+      bvh_instance_motion_push(NULL, object, ray, &r.origin, &r.direction, &idir);
+#  else
+      bvh_instance_push(NULL, object, ray, &r.origin, &r.direction, &idir);
+#  endif
+    }
+
+    if (prim_type & PRIMITIVE_POINT) {
+      if (!point_intersect(NULL,
+                           isect,
+                           r.origin,
+                           r.direction,
+                           ray->tmin,
+                           ray->tmax,
+                           intersection.instance_id,
+                           prim,
+                           ray->time,
+                           prim_type))
+      {
+        /* Shouldn't get here */
+        kernel_assert(!"Intersection mismatch");
+        return false;
+      }
+      return true;
+    }
   }
 
-  return isect->type != PRIMITIVE_NONE;
+  return true;
 }
 #endif
 

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

@@ -27,6 +27,8 @@ using namespace metal::raytracing;
 #pragma clang diagnostic ignored "-Wunused-variable"
 #pragma clang diagnostic ignored "-Wsign-compare"
 #pragma clang diagnostic ignored "-Wuninitialized"
+#pragma clang diagnostic ignored "-Wc++17-extensions"
+#pragma clang diagnostic ignored "-Wmacro-redefined"
 
 /* Qualifiers */
 
@@ -280,17 +282,23 @@ ccl_device_forceinline uchar4 make_uchar4(const uchar x,
 #  endif /* __METALRT_MOTION__ */
 
 typedef acceleration_structure<METALRT_TAGS> metalrt_as_type;
-typedef intersection_function_table<triangle_data, METALRT_TAGS> metalrt_ift_type;
+typedef intersection_function_table<triangle_data, curve_data, METALRT_TAGS, extended_limits>
-typedef metal::raytracing::intersector<triangle_data, METALRT_TAGS> metalrt_intersector_type;
+    metalrt_ift_type;
+typedef metal::raytracing::intersector<triangle_data, curve_data, METALRT_TAGS, extended_limits>
+    metalrt_intersector_type;
 #  if defined(__METALRT_MOTION__)
 typedef acceleration_structure<primitive_motion> metalrt_blas_as_type;
-typedef intersection_function_table<triangle_data, primitive_motion> metalrt_blas_ift_type;
+typedef intersection_function_table<triangle_data, curve_data, primitive_motion, extended_limits>
-typedef metal::raytracing::intersector<triangle_data, primitive_motion>
+    metalrt_blas_ift_type;
-    metalrt_blas_intersector_type;
+typedef metal::raytracing::
+    intersector<triangle_data, curve_data, primitive_motion, extended_limits>
+        metalrt_blas_intersector_type;
 #  else
 typedef acceleration_structure<> metalrt_blas_as_type;
-typedef intersection_function_table<triangle_data> metalrt_blas_ift_type;
+typedef intersection_function_table<triangle_data, curve_data, extended_limits>
-typedef metal::raytracing::intersector<triangle_data> metalrt_blas_intersector_type;
+    metalrt_blas_ift_type;
+typedef metal::raytracing::intersector<triangle_data, curve_data, extended_limits>
+    metalrt_blas_intersector_type;
 #  endif
 
 #endif /* __METALRT__ */
@@ -326,7 +334,6 @@ struct MetalAncillaries {
   metalrt_ift_type ift_local;
   metalrt_blas_ift_type ift_local_prim;
   constant MetalRTBlasWrapper *blas_accel_structs;
-  constant int *blas_userID_to_index_lookUp;
 #endif
 };
 

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

@@ -26,13 +26,13 @@ struct BoundingBoxIntersectionResult {
   float distance [[distance]];
 };
 
-/* For a triangle intersection function. */
+/* For a primitive intersection function. */
-struct TriangleIntersectionResult {
+struct PrimitiveIntersectionResult {
   bool accept [[accept_intersection]];
   bool continue_search [[continue_search]];
 };
 
-enum { METALRT_HIT_TRIANGLE, METALRT_HIT_BOUNDING_BOX };
+enum { METALRT_HIT_TRIANGLE, METALRT_HIT_CURVE, METALRT_HIT_BOUNDING_BOX };
 
 /* Hit functions. */
 
@@ -40,20 +40,17 @@ template<typename TReturn, uint intersection_type>
 TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
                           ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload,
                           const uint object,
-                          const uint primitive_id,
+                          const uint prim,
                           const float2 barycentrics,
                           const float ray_tmax)
 {
   TReturn result;
 
-#  ifdef __BVH_LOCAL__
+#ifdef __BVH_LOCAL__
-  uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
-
   MetalKernelContext context(launch_params_metal);
 
-  if ((object != payload.local_object) || context.intersection_skip_self_local(payload.self, prim))
+  if ((object != payload.local_object) || context.intersection_skip_self_local(payload.self, prim)) {
-  {
+    /* Only intersect with matching object and skip self-intersecton. */
-    /* Only intersect with matching object and skip self-intersection. */
     result.accept = false;
     result.continue_search = true;
     return result;
@@ -124,7 +121,7 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
 #  endif
 }
 
-[[intersection(triangle, triangle_data)]] TriangleIntersectionResult
+[[intersection(triangle, triangle_data, curve_data)]] PrimitiveIntersectionResult
 __anyhit__cycles_metalrt_local_hit_tri_prim(
     constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
     ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload [[payload]],
@@ -132,27 +129,30 @@ __anyhit__cycles_metalrt_local_hit_tri_prim(
     float2 barycentrics [[barycentric_coord]],
     float ray_tmax [[distance]])
 {
+  uint prim = primitive_id + kernel_data_fetch(object_prim_offset, payload.local_object);
+
   /* instance_id, aka the user_id has been removed. If we take this function we optimized the
    * SSS for starting traversal from a primitive acceleration structure instead of the root of the
    * global AS. this means we will always be intersecting the correct object no need for the
    * user-id to check */
-  return metalrt_local_hit<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>(
+  return metalrt_local_hit<PrimitiveIntersectionResult, METALRT_HIT_TRIANGLE>(
-      launch_params_metal, payload, payload.local_object, primitive_id, barycentrics, ray_tmax);
+      launch_params_metal, payload, payload.local_object, prim, barycentrics, ray_tmax);
 }
-[[intersection(triangle, triangle_data, METALRT_TAGS)]] TriangleIntersectionResult
+[[intersection(triangle, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
 __anyhit__cycles_metalrt_local_hit_tri(
     constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
     ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload [[payload]],
-    uint instance_id [[user_instance_id]],
+    uint instance_id [[instance_id]],
     uint primitive_id [[primitive_id]],
+    uint primitive_id_offset [[user_instance_id]],
     float2 barycentrics [[barycentric_coord]],
     float ray_tmax [[distance]])
 {
-  return metalrt_local_hit<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>(
+  return metalrt_local_hit<PrimitiveIntersectionResult, METALRT_HIT_TRIANGLE>(
-      launch_params_metal, payload, instance_id, primitive_id, barycentrics, ray_tmax);
+      launch_params_metal, payload, instance_id, primitive_id + primitive_id_offset, barycentrics, ray_tmax);
 }
 
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
+[[intersection(bounding_box, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] BoundingBoxIntersectionResult
 __anyhit__cycles_metalrt_local_hit_box(const float ray_tmax [[max_distance]])
 {
   /* unused function */
@@ -163,7 +163,7 @@ __anyhit__cycles_metalrt_local_hit_box(const float ray_tmax [[max_distance]])
   return result;
 }
 
-[[intersection(bounding_box, triangle_data)]] BoundingBoxIntersectionResult
+[[intersection(bounding_box, triangle_data, curve_data)]] BoundingBoxIntersectionResult
 __anyhit__cycles_metalrt_local_hit_box_prim(const float ray_tmax [[max_distance]])
 {
   /* unused function */
@@ -180,30 +180,32 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
                             uint object,
                             uint prim,
                             const float2 barycentrics,
-                            const float ray_tmax)
+                            const float ray_tmax,
+                            const float t = 0.0f,
+                            ccl_private const Ray *ray = NULL
+                            )
 {
-#  ifdef __SHADOW_RECORD_ALL__
+#ifdef __SHADOW_RECORD_ALL__
-#    ifdef __VISIBILITY_FLAG__
+  float u = barycentrics.x;
-  const uint visibility = payload.visibility;
+  float v = barycentrics.y;
-  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
-    /* continue search */
-    return true;
-  }
-#    endif
-
-  const float u = barycentrics.x;
-  const float v = barycentrics.y;
   const int prim_type = kernel_data_fetch(objects, object).primitive_type;
-  int type = prim_type;
+  int type;
-#    ifdef __HAIR__
-  if (intersection_type != METALRT_HIT_TRIANGLE) {
-    if ((prim_type == PRIMITIVE_CURVE_THICK || prim_type == PRIMITIVE_CURVE_RIBBON)) {
-      const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
-      type = segment.type;
-      prim = segment.prim;
 
-      /* Filter out curve end-caps. */
+#  ifdef __HAIR__
-      if (u == 0.0f || u == 1.0f) {
+  if constexpr (intersection_type == METALRT_HIT_CURVE) {
+    const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
+    type = segment.type;
+    prim = segment.prim;
+    
+    /* Filter out curve end-caps. */
+    if (u == 0.0f || u == 1.0f) {
+      /* continue search */
+      return true;
+    }
+
+    if (type & PRIMITIVE_CURVE_RIBBON) {
+      MetalKernelContext context(launch_params_metal);
+      if (!context.curve_ribbon_accept(NULL, u, t, ray, object, prim, type)) {
         /* continue search */
         return true;
       }
@@ -211,6 +213,17 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
   }
 #    endif
 
+  if constexpr (intersection_type == METALRT_HIT_BOUNDING_BOX) {
+    /* Point. */
+    type = kernel_data_fetch(objects, object).primitive_type;
+    u = 0.0f;
+    v = 0.0f;
+  }
+
+  if constexpr (intersection_type == METALRT_HIT_TRIANGLE) {
+    type = prim_type;
+  }
+
   MetalKernelContext context(launch_params_metal);
 
   if (context.intersection_skip_self_shadow(payload.self, object, prim)) {
@@ -244,8 +257,9 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
     return false;
   }
 
+#  ifdef __HAIR__
   /* Always use baked shadow transparency for curves. */
-  if (type & PRIMITIVE_CURVE) {
+  if constexpr (intersection_type == METALRT_HIT_CURVE) {
     float throughput = payload.throughput;
     throughput *= context.intersection_curve_shadow_transparency(nullptr, object, prim, type, u);
     payload.throughput = throughput;
@@ -260,6 +274,7 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
       return true;
     }
   }
+#  endif
 
   payload.num_hits += 1;
   payload.num_recorded_hits += 1;
@@ -305,25 +320,26 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
   return true;
 }
 
-[[intersection(triangle, triangle_data, METALRT_TAGS)]] TriangleIntersectionResult
+[[intersection(triangle, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
 __anyhit__cycles_metalrt_shadow_all_hit_tri(
     constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
     ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]],
-    unsigned int object [[user_instance_id]],
+    const unsigned int object [[instance_id]],
-    unsigned int primitive_id [[primitive_id]],
+    const unsigned int primitive_id [[primitive_id]],
-    float2 barycentrics [[barycentric_coord]],
+    const uint primitive_id_offset [[user_instance_id]],
-    float ray_tmax [[distance]])
+    const float2 barycentrics [[barycentric_coord]],
+    const float ray_tmax [[distance]])
 {
-  uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
+  uint prim = primitive_id + primitive_id_offset;
 
-  TriangleIntersectionResult result;
+  PrimitiveIntersectionResult result;
   result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_TRIANGLE>(
       launch_params_metal, payload, object, prim, barycentrics, ray_tmax);
   result.accept = !result.continue_search;
   return result;
 }
 
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
+[[intersection(bounding_box, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] BoundingBoxIntersectionResult
 __anyhit__cycles_metalrt_shadow_all_hit_box(const float ray_tmax [[max_distance]])
 {
   /* unused function */
@@ -340,40 +356,38 @@ inline TReturnType metalrt_visibility_test(
     ray_data MetalKernelContext::MetalRTIntersectionPayload &payload,
     const uint object,
     uint prim,
-    const float u)
+    const float u,
+    const float t = 0.0f,
+    ccl_private const Ray *ray = NULL
+    )
 {
   TReturnType result;
 
-#  ifdef __HAIR__
+#ifdef __HAIR__
-  const int type = kernel_data_fetch(objects, object).primitive_type;
+  if constexpr (intersection_type == METALRT_HIT_CURVE) {
-  if (intersection_type == METALRT_HIT_BOUNDING_BOX &&
-      (type == PRIMITIVE_CURVE_THICK || type == PRIMITIVE_CURVE_RIBBON))
-  {
     /* Filter out curve end-caps. */
     if (u == 0.0f || u == 1.0f) {
       result.accept = false;
       result.continue_search = true;
       return result;
     }
+    
+    const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
+    int type = segment.type;
+    prim = segment.prim;
+
+    if (type & PRIMITIVE_CURVE_RIBBON) {
+      MetalKernelContext context(launch_params_metal);
+      if (!context.curve_ribbon_accept(NULL, u, t, ray, object, prim, type)) {
+        result.accept = false;
+        result.continue_search = true;
+        return result;
+      }
+    }
   }
 #  endif
 
   uint visibility = payload.visibility;
-#  ifdef __VISIBILITY_FLAG__
-  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
-    result.accept = false;
-    result.continue_search = true;
-    return result;
-  }
-#  endif
-
-  if (intersection_type == METALRT_HIT_TRIANGLE) {
-  }
-#  ifdef __HAIR__
-  else {
-    prim = kernel_data_fetch(curve_segments, prim).prim;
-  }
-#  endif
 
   MetalKernelContext context(launch_params_metal);
 
@@ -411,25 +425,22 @@ inline TReturnType metalrt_visibility_test(
   return result;
 }
 
-[[intersection(triangle, triangle_data, METALRT_TAGS)]] TriangleIntersectionResult
+[[intersection(triangle, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
 __anyhit__cycles_metalrt_visibility_test_tri(
     constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
     ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]],
-    unsigned int object [[user_instance_id]],
+    const unsigned int object [[instance_id]],
-    unsigned int primitive_id [[primitive_id]])
+    const uint primitive_id_offset [[user_instance_id]],
+    const unsigned int primitive_id [[primitive_id]])
 {
-  uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
+  uint prim = primitive_id + primitive_id_offset;
-  TriangleIntersectionResult result =
+  PrimitiveIntersectionResult result =
-      metalrt_visibility_test<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>(
+      metalrt_visibility_test<PrimitiveIntersectionResult, METALRT_HIT_TRIANGLE>(
           launch_params_metal, payload, object, prim, 0.0f);
-  if (result.accept) {
-    payload.prim = prim;
-    payload.type = kernel_data_fetch(objects, object).primitive_type;
-  }
   return result;
 }
 
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
+[[intersection(bounding_box, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] BoundingBoxIntersectionResult
 __anyhit__cycles_metalrt_visibility_test_box(const float ray_tmax [[max_distance]])
 {
   /* Unused function */
@@ -442,230 +453,72 @@ __anyhit__cycles_metalrt_visibility_test_box(const float ray_tmax [[max_distance
 
 /* Primitive intersection functions. */
 
-#  ifdef __HAIR__
+#ifdef __HAIR__
-ccl_device_inline void metalrt_intersection_curve(
+[[intersection(curve, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
-    constant KernelParamsMetal &launch_params_metal,
+__intersection__curve(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
-    ray_data MetalKernelContext::MetalRTIntersectionPayload &payload,
+                      ray_data MetalKernelContext::MetalRTIntersectionPayload &payload
-    const uint object,
+                      [[payload]],
-    const uint prim,
+                      const uint object [[instance_id]],
-    const uint type,
+                      const uint primitive_id [[primitive_id]],
-    const float3 ray_P,
+                      const uint primitive_id_offset [[user_instance_id]],
-    const float3 ray_D,
+                      float distance              [[distance]],
-    float time,
+                      const float3 ray_P [[origin]],
-    const float ray_tmin,
+                      const float3 ray_D [[direction]],
-    const float ray_tmax,
+                      float u [[curve_parameter]],
-    thread BoundingBoxIntersectionResult &result)
+                      const float ray_tmin [[min_distance]],
+                      const float ray_tmax [[max_distance]]
+#  if defined(__METALRT_MOTION__)
+                      ,const float time [[time]]
+#  endif                             
+                      )
 {
-#    ifdef __VISIBILITY_FLAG__
+  uint prim = primitive_id + primitive_id_offset;
-  const uint visibility = payload.visibility;
-  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
-    return;
-  }
-#    endif
 
-  Intersection isect;
+  Ray ray;
-  isect.t = ray_tmax;
+  ray.P = ray_P;
+  ray.D = ray_D;
+#if defined(__METALRT_MOTION__)
+  ray.time = time;
+#endif
 
-  MetalKernelContext context(launch_params_metal);
+  PrimitiveIntersectionResult result =
-  if (context.curve_intersect(
+    metalrt_visibility_test<PrimitiveIntersectionResult, METALRT_HIT_CURVE>(
-          NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type))
+      launch_params_metal, payload, object, prim, u, distance, &ray);
-  {
-    result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
-        launch_params_metal, payload, object, prim, isect.u);
-    if (result.accept) {
-      result.distance = isect.t;
-      payload.u = isect.u;
-      payload.v = isect.v;
-      payload.prim = prim;
-      payload.type = type;
-    }
-  }
-}
-
-ccl_device_inline void metalrt_intersection_curve_shadow(
-    constant KernelParamsMetal &launch_params_metal,
-    ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload,
-    const uint object,
-    const uint prim,
-    const uint type,
-    const float3 ray_P,
-    const float3 ray_D,
-    float time,
-    const float ray_tmin,
-    const float ray_tmax,
-    thread BoundingBoxIntersectionResult &result)
-{
-#    ifdef __VISIBILITY_FLAG__
-  const uint visibility = payload.visibility;
-  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
-    return;
-  }
-#    endif
-
-  Intersection isect;
-  isect.t = ray_tmax;
-
-  MetalKernelContext context(launch_params_metal);
-  if (context.curve_intersect(
-          NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type))
-  {
-    result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>(
-        launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax);
-    result.accept = !result.continue_search;
-  }
-}
-
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
-__intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
-                             ray_data MetalKernelContext::MetalRTIntersectionPayload &payload
-                             [[payload]],
-                             const uint object [[user_instance_id]],
-                             const uint primitive_id [[primitive_id]],
-                             const float3 ray_P [[origin]],
-                             const float3 ray_D [[direction]],
-                             const float ray_tmin [[min_distance]],
-                             const float ray_tmax [[max_distance]])
-{
-  uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
-  const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
-
-  BoundingBoxIntersectionResult result;
-  result.accept = false;
-  result.continue_search = true;
-  result.distance = ray_tmax;
-
-  if (segment.type & PRIMITIVE_CURVE_RIBBON) {
-    metalrt_intersection_curve(launch_params_metal,
-                               payload,
-                               object,
-                               segment.prim,
-                               segment.type,
-                               ray_P,
-                               ray_D,
-#    if defined(__METALRT_MOTION__)
-                               payload.time,
-#    else
-                               0.0f,
-#    endif
-                               ray_tmin,
-                               ray_tmax,
-                               result);
-  }
 
   return result;
 }
 
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
+[[intersection(curve, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
-__intersection__curve_ribbon_shadow(
+__intersection__curve_shadow(
     constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
     ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]],
-    const uint object [[user_instance_id]],
+    const uint object [[instance_id]],
     const uint primitive_id [[primitive_id]],
+    const uint primitive_id_offset [[user_instance_id]],
     const float3 ray_P [[origin]],
     const float3 ray_D [[direction]],
+    float u [[curve_parameter]],
+    float t [[distance]],
+#  if defined(__METALRT_MOTION__)
+    const float time [[time]],
+#  endif
     const float ray_tmin [[min_distance]],
     const float ray_tmax [[max_distance]])
 {
-  uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
+  uint prim = primitive_id + primitive_id_offset;
-  const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
 
-  BoundingBoxIntersectionResult result;
+  PrimitiveIntersectionResult result;
-  result.accept = false;
-  result.continue_search = true;
-  result.distance = ray_tmax;
 
-  if (segment.type & PRIMITIVE_CURVE_RIBBON) {
+  Ray ray;
-    metalrt_intersection_curve_shadow(launch_params_metal,
+  ray.P = ray_P;
-                                      payload,
+  ray.D = ray_D;
-                                      object,
+#if defined(__METALRT_MOTION__)
-                                      segment.prim,
+  ray.time = time;
-                                      segment.type,
+#endif
-                                      ray_P,
-                                      ray_D,
-#    if defined(__METALRT_MOTION__)
-                                      payload.time,
-#    else
-                                      0.0f,
-#    endif
-                                      ray_tmin,
-                                      ray_tmax,
-                                      result);
-  }
 
-  return result;
+  result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_CURVE>(
-}
+      launch_params_metal, payload, object, prim, float2(u, 0), ray_tmax, t, &ray);
-
+  result.accept = !result.continue_search;
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
-__intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
-                          ray_data MetalKernelContext::MetalRTIntersectionPayload &payload
-                          [[payload]],
-                          const uint object [[user_instance_id]],
-                          const uint primitive_id [[primitive_id]],
-                          const float3 ray_P [[origin]],
-                          const float3 ray_D [[direction]],
-                          const float ray_tmin [[min_distance]],
-                          const float ray_tmax [[max_distance]])
-{
-  uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
-  const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
-
-  BoundingBoxIntersectionResult result;
-  result.accept = false;
-  result.continue_search = true;
-  result.distance = ray_tmax;
-  metalrt_intersection_curve(launch_params_metal,
-                             payload,
-                             object,
-                             segment.prim,
-                             segment.type,
-                             ray_P,
-                             ray_D,
-#    if defined(__METALRT_MOTION__)
-                             payload.time,
-#    else
-                             0.0f,
-#    endif
-                             ray_tmin,
-                             ray_tmax,
-                             result);
-
-  return result;
-}
-
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
-__intersection__curve_all_shadow(
-    constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
-    ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]],
-    const uint object [[user_instance_id]],
-    const uint primitive_id [[primitive_id]],
-    const float3 ray_P [[origin]],
-    const float3 ray_D [[direction]],
-    const float ray_tmin [[min_distance]],
-    const float ray_tmax [[max_distance]])
-{
-  uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
-  const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
-
-  BoundingBoxIntersectionResult result;
-  result.accept = false;
-  result.continue_search = true;
-  result.distance = ray_tmax;
-
-  metalrt_intersection_curve_shadow(launch_params_metal,
-                                    payload,
-                                    object,
-                                    segment.prim,
-                                    segment.type,
-                                    ray_P,
-                                    ray_D,
-#    if defined(__METALRT_MOTION__)
-                                    payload.time,
-#    else
-                                    0.0f,
-#    endif
-                                    ray_tmin,
-                                    ray_tmax,
-                                    result);
 
   return result;
 }
@@ -685,13 +538,6 @@ ccl_device_inline void metalrt_intersection_point(
     const float ray_tmax,
     thread BoundingBoxIntersectionResult &result)
 {
-#    ifdef __VISIBILITY_FLAG__
-  const uint visibility = payload.visibility;
-  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
-    return;
-  }
-#    endif
-
   Intersection isect;
   isect.t = ray_tmax;
 
@@ -703,10 +549,6 @@ ccl_device_inline void metalrt_intersection_point(
         launch_params_metal, payload, object, prim, isect.u);
     if (result.accept) {
       result.distance = isect.t;
-      payload.u = isect.u;
-      payload.v = isect.v;
-      payload.prim = prim;
-      payload.type = type;
     }
   }
 }
@@ -724,13 +566,6 @@ ccl_device_inline void metalrt_intersection_point_shadow(
     const float ray_tmax,
     thread BoundingBoxIntersectionResult &result)
 {
-#    ifdef __VISIBILITY_FLAG__
-  const uint visibility = payload.visibility;
-  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
-    return;
-  }
-#    endif
-
   Intersection isect;
   isect.t = ray_tmax;
 
@@ -748,17 +583,21 @@ ccl_device_inline void metalrt_intersection_point_shadow(
   }
 }
 
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
+[[intersection(bounding_box, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] BoundingBoxIntersectionResult
 __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
                       ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]],
-                      const uint object [[user_instance_id]],
+                      const uint object [[instance_id]],
                       const uint primitive_id [[primitive_id]],
+                      const uint primitive_id_offset [[user_instance_id]],
                       const float3 ray_origin [[origin]],
                       const float3 ray_direction [[direction]],
+#  if defined(__METALRT_MOTION__)
+                      const float time [[time]],
+#  endif
                       const float ray_tmin [[min_distance]],
                       const float ray_tmax [[max_distance]])
 {
-  const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
+  const uint prim = primitive_id + primitive_id_offset;
   const int type = kernel_data_fetch(objects, object).primitive_type;
 
   BoundingBoxIntersectionResult result;
@@ -774,7 +613,7 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1
                              ray_origin,
                              ray_direction,
 #    if defined(__METALRT_MOTION__)
-                             payload.time,
+                             time,
 #    else
                              0.0f,
 #    endif
@@ -785,18 +624,22 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1
   return result;
 }
 
-[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
+[[intersection(bounding_box, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] BoundingBoxIntersectionResult
 __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
                              ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload
                              [[payload]],
-                             const uint object [[user_instance_id]],
+                             const uint object [[instance_id]],
                              const uint primitive_id [[primitive_id]],
+                             const uint primitive_id_offset [[user_instance_id]],
                              const float3 ray_origin [[origin]],
                              const float3 ray_direction [[direction]],
+#  if defined(__METALRT_MOTION__)
+                             const float time [[time]],
+#  endif
                              const float ray_tmin [[min_distance]],
                              const float ray_tmax [[max_distance]])
 {
-  const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
+  const uint prim = primitive_id + primitive_id_offset;
   const int type = kernel_data_fetch(objects, object).primitive_type;
 
   BoundingBoxIntersectionResult result;
@@ -812,7 +655,7 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b
                                     ray_origin,
                                     ray_direction,
 #    if defined(__METALRT_MOTION__)
-                                    payload.time,
+                                    time,
 #    else
                                     0.0f,
 #    endif

intern/cycles/kernel/integrator/state_template.h

@@ -95,7 +95,7 @@ KERNEL_STRUCT_BEGIN(subsurface)
 KERNEL_STRUCT_MEMBER(subsurface, PackedSpectrum, albedo, KERNEL_FEATURE_SUBSURFACE)
 KERNEL_STRUCT_MEMBER(subsurface, PackedSpectrum, radius, KERNEL_FEATURE_SUBSURFACE)
 KERNEL_STRUCT_MEMBER(subsurface, float, anisotropy, KERNEL_FEATURE_SUBSURFACE)
-KERNEL_STRUCT_MEMBER(subsurface, packed_float3, Ng, KERNEL_FEATURE_SUBSURFACE)
+KERNEL_STRUCT_MEMBER(subsurface, packed_float3, N, KERNEL_FEATURE_SUBSURFACE)
 KERNEL_STRUCT_END(subsurface)
 
 /********************************** Volume Stack ******************************/

intern/cycles/kernel/integrator/subsurface.h

@@ -23,6 +23,60 @@ CCL_NAMESPACE_BEGIN
 
 #ifdef __SUBSURFACE__
 
+ccl_device_inline bool subsurface_entry_bounce(KernelGlobals kg,
+                                               ccl_private const Bssrdf *bssrdf,
+                                               ccl_private ShaderData *sd,
+                                               ccl_private RNGState *rng_state,
+                                               ccl_private float3 *wo)
+{
+  float2 rand_bsdf = path_state_rng_2D(kg, rng_state, PRNG_SUBSURFACE_BSDF);
+
+  if (bssrdf->type == CLOSURE_BSSRDF_RANDOM_WALK_ID) {
+    /* CLOSURE_BSSRDF_RANDOM_WALK_ID has a 50% chance to sample a diffuse entry bounce.
+     * Also, for the refractive entry, it uses a fixed roughness of 1.0. */
+    if (rand_bsdf.x < 0.5f) {
+      rand_bsdf.x *= 2.0f;
+      float pdf;
+      sample_cos_hemisphere(-bssrdf->N, rand_bsdf, wo, &pdf);
+      return true;
+    }
+    rand_bsdf.x = 2.0f * (rand_bsdf.x - 0.5f);
+  }
+
+  const float cos_NI = dot(bssrdf->N, sd->wi);
+  if (cos_NI <= 0.0f) {
+    return false;
+  }
+
+  float3 X, Y, Z = bssrdf->N;
+  make_orthonormals(Z, &X, &Y);
+
+  const float alpha = bssrdf->alpha;
+  const float neta = 1.0f / bssrdf->ior;
+
+  /* Sample microfacet normal by transforming to/from local coordinates. */
+  const float3 local_I = make_float3(dot(X, sd->wi), dot(Y, sd->wi), cos_NI);
+  const float3 local_H = microfacet_ggx_sample_vndf(local_I, alpha, alpha, rand_bsdf);
+  const float3 H = X * local_H.x + Y * local_H.y + Z * local_H.z;
+
+  const float cos_HI = dot(H, sd->wi);
+  const float arg = 1.0f - (sqr(neta) * (1.0f - sqr(cos_HI)));
+  /* We clamp subsurface IOR to be above 1, so there should never be TIR. */
+  kernel_assert(arg >= 0.0f);
+
+  const float dnp = max(sqrtf(arg), 1e-7f);
+  const float nK = (neta * cos_HI) - dnp;
+  *wo = -(neta * sd->wi) + (nK * H);
+  return true;
+  /* Note: For a proper refractive GGX interface, we should be computing lambdaI and lambdaO
+   * and multiplying the throughput by BSDF/pdf, which for VNDF sampling works out to
+   * (1 + lambdaI) / (1 + lambdaI + lambdaO).
+   * However, this causes darkening due to the single-scattering approximation, which we'd
+   * then have to correct with a lookup table.
+   * Since we only really care about the directional distribution here, it's much easier to
+   * just skip all that instead. */
+}
+
 ccl_device int subsurface_bounce(KernelGlobals kg,
                                  IntegratorState state,
                                  ccl_private ShaderData *sd,
@@ -37,32 +91,44 @@ ccl_device int subsurface_bounce(KernelGlobals kg,
 
   /* Setup ray into surface. */
   INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
-  INTEGRATOR_STATE_WRITE(state, ray, D) = bssrdf->N;
   INTEGRATOR_STATE_WRITE(state, ray, tmin) = 0.0f;
   INTEGRATOR_STATE_WRITE(state, ray, tmax) = FLT_MAX;
   INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
   INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_zero_compact();
 
-  /* Pass along object info, reusing isect to save memory. */
+  /* Advance random number offset for bounce. */
-  INTEGRATOR_STATE_WRITE(state, subsurface, Ng) = sd->Ng;
+  INTEGRATOR_STATE_WRITE(state, path, rng_offset) += PRNG_BOUNCE_NUM;
-
-  uint32_t path_flag = (INTEGRATOR_STATE(state, path, flag) & ~PATH_RAY_CAMERA) |
-                       ((sc->type == CLOSURE_BSSRDF_BURLEY_ID) ? PATH_RAY_SUBSURFACE_DISK :
-                                                                 PATH_RAY_SUBSURFACE_RANDOM_WALK);
 
   /* Compute weight, optionally including Fresnel from entry point. */
   Spectrum weight = surface_shader_bssrdf_sample_weight(sd, sc);
+  INTEGRATOR_STATE_WRITE(state, path, throughput) *= weight;
+
+  uint32_t path_flag = (INTEGRATOR_STATE(state, path, flag) & ~PATH_RAY_CAMERA);
+  if (sc->type == CLOSURE_BSSRDF_BURLEY_ID) {
+    path_flag |= PATH_RAY_SUBSURFACE_DISK;
+    INTEGRATOR_STATE_WRITE(state, subsurface, N) = sd->Ng;
+  }
+  else {
+    path_flag |= PATH_RAY_SUBSURFACE_RANDOM_WALK;
+
+    /* Sample entry bounce into the material. */
+    RNGState rng_state;
+    path_state_rng_load(state, &rng_state);
+    float3 wo;
+    if (!subsurface_entry_bounce(kg, bssrdf, sd, &rng_state, &wo) || dot(sd->Ng, wo) >= 0.0f) {
+      /* Sampling failed, give up on this bounce. */
+      return LABEL_NONE;
+    }
+    INTEGRATOR_STATE_WRITE(state, ray, D) = wo;
+    INTEGRATOR_STATE_WRITE(state, subsurface, N) = sd->N;
+  }
 
   if (sd->flag & SD_BACKFACING) {
     path_flag |= PATH_RAY_SUBSURFACE_BACKFACING;
   }
 
-  INTEGRATOR_STATE_WRITE(state, path, throughput) *= weight;
   INTEGRATOR_STATE_WRITE(state, path, flag) = path_flag;
 
-  /* Advance random number offset for bounce. */
-  INTEGRATOR_STATE_WRITE(state, path, rng_offset) += PRNG_BOUNCE_NUM;
-
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
     if (INTEGRATOR_STATE(state, path, bounce) == 0) {
       INTEGRATOR_STATE_WRITE(state, path, pass_diffuse_weight) = one_spectrum();

intern/cycles/kernel/integrator/subsurface_disk.h

@@ -34,7 +34,7 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg,
   const float3 P = INTEGRATOR_STATE(state, ray, P);
   const float ray_dP = INTEGRATOR_STATE(state, ray, dP);
   const float time = INTEGRATOR_STATE(state, ray, time);
-  const float3 Ng = INTEGRATOR_STATE(state, subsurface, Ng);
+  const float3 Ng = INTEGRATOR_STATE(state, subsurface, N);
   const int object = INTEGRATOR_STATE(state, isect, object);
   const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
 

intern/cycles/kernel/integrator/subsurface_random_walk.h

@@ -168,24 +168,14 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
                                               ccl_private Ray &ray,
                                               ccl_private LocalIntersection &ss_isect)
 {
-  const float2 rand_bsdf = path_state_rng_2D(kg, &rng_state, PRNG_SUBSURFACE_BSDF);
-
   const float3 P = INTEGRATOR_STATE(state, ray, P);
-  const float3 N = INTEGRATOR_STATE(state, ray, D);
+  const float3 D = INTEGRATOR_STATE(state, ray, D);
   const float ray_dP = INTEGRATOR_STATE(state, ray, dP);
   const float time = INTEGRATOR_STATE(state, ray, time);
-  const float3 Ng = INTEGRATOR_STATE(state, subsurface, Ng);
+  const float3 N = INTEGRATOR_STATE(state, subsurface, N);
   const int object = INTEGRATOR_STATE(state, isect, object);
   const int prim = INTEGRATOR_STATE(state, isect, prim);
 
-  /* Sample diffuse surface scatter into the object. */
-  float3 D;
-  float pdf;
-  sample_cos_hemisphere(-N, rand_bsdf, &D, &pdf);
-  if (dot(-Ng, D) <= 0.0f) {
-    return false;
-  }
-
   /* Setup ray. */
   ray.P = P;
   ray.D = D;
@@ -436,10 +426,11 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
   if (hit) {
     kernel_assert(isfinite_safe(throughput));
 
+    /* TODO(lukas): Which PDF should we report here? Entry bounce? The random walk? Just 1.0? */
     guiding_record_bssrdf_bounce(
         kg,
         state,
-        pdf,
+        1.0f,
         N,
         D,
         safe_divide_color(throughput, INTEGRATOR_STATE(state, path, throughput)),

intern/cycles/kernel/osl/closures_setup.h

@@ -420,10 +420,6 @@ ccl_device void osl_closure_microfacet_setup(KernelGlobals kg,
   else if (closure->distribution == make_string("ashikhmin_shirley", 11318482998918370922ull)) {
     sd->flag |= bsdf_ashikhmin_shirley_setup(bsdf);
   }
-  /* Clearcoat */
-  else if (closure->distribution == make_string("clearcoat", 3490136178980547276ull)) {
-    sd->flag |= bsdf_microfacet_ggx_clearcoat_setup(kg, bsdf, sd);
-  }
   /* GGX (either single- or multi-scattering) */
   else {
     if (closure->refract == 1) {
@@ -744,23 +740,16 @@ ccl_device void osl_closure_bssrdf_setup(KernelGlobals kg,
     return;
   }
 
-  /* disable in case of diffuse ancestor, can't see it well then and
+  bssrdf->radius = closure->radius;
-   * adds considerably noise due to probabilities of continuing path
-   * getting lower and lower */
-  if (path_flag & PATH_RAY_DIFFUSE_ANCESTOR) {
-    bssrdf->radius = zero_spectrum();
-  }
-  else {
-    bssrdf->radius = closure->radius;
-  }
 
   /* create one closure per color channel */
   bssrdf->albedo = closure->albedo;
   bssrdf->N = closure->N;
-  bssrdf->roughness = closure->roughness;
+  bssrdf->alpha = sqr(closure->roughness);
-  bssrdf->anisotropy = clamp(closure->anisotropy, 0.0f, 0.9f);
+  bssrdf->ior = closure->ior;
+  bssrdf->anisotropy = closure->anisotropy;
 
-  sd->flag |= bssrdf_setup(sd, bssrdf, type, clamp(closure->ior, 1.01f, 3.8f));
+  sd->flag |= bssrdf_setup(sd, bssrdf, path_flag, type);
 }
 
 /* Hair */

intern/cycles/kernel/osl/osl.h

@@ -140,7 +140,7 @@ ccl_device void flatten_closure_tree(KernelGlobals kg,
       if (stack_size == layer_stack_level) {
         /* We just finished processing the top layers of a Layer closure, so adjust the weight to
          * account for the layering. */
-        weight *= saturatef(1.0f - reduce_max(layer_albedo / weight));
+        weight *= saturatef(1.0f - reduce_max(safe_divide_color(layer_albedo, weight)));
         layer_stack_level = -1;
         if (is_zero(weight)) {
           /* If it's fully occluded, skip the base layer we just popped from the stack and grab

intern/cycles/kernel/osl/shaders/node_principled_bsdf.osl

@@ -9,8 +9,8 @@ shader node_principled_bsdf(string distribution = "multi_ggx",
                             string subsurface_method = "random_walk",
                             color BaseColor = color(0.8, 0.8, 0.8),
                             float Subsurface = 0.0,
+                            float SubsurfaceScale = 0.1,
                             vector SubsurfaceRadius = vector(1.0, 1.0, 1.0),
-                            color SubsurfaceColor = color(0.7, 0.1, 0.1),
                             float SubsurfaceIOR = 1.4,
                             float SubsurfaceAnisotropy = 0.0,
                             float Metallic = 0.0,
@@ -22,12 +22,17 @@ shader node_principled_bsdf(string distribution = "multi_ggx",
                             float Sheen = 0.0,
                             float SheenRoughness = 0.5,
                             color SheenTint = 0.5,
-                            float Clearcoat = 0.0,
+                            float Coat = 0.0,
-                            float ClearcoatRoughness = 0.03,
+                            float CoatRoughness = 0.03,
+                            float CoatIOR = 1.5,
+                            color CoatTint = color(1.0, 1.0, 1.0),
                             float IOR = 1.45,
                             float Transmission = 0.0,
+                            color Emission = 1.0,
+                            float EmissionStrength = 0.0,
+                            float Alpha = 1.0,
                             normal Normal = N,
-                            normal ClearcoatNormal = N,
+                            normal CoatNormal = N,
                             normal Tangent = normalize(dPdu),
                             output closure color BSDF = 0)
 {
@@ -45,21 +50,21 @@ shader node_principled_bsdf(string distribution = "multi_ggx",
   }
 
   if (Metallic < 1.0 && Transmission < 1.0) {
-    color diffuse_color = mix(BaseColor, SubsurfaceColor, Subsurface);
+    BSDF = BaseColor * diffuse(Normal);
     if (Subsurface > 1e-5) {
-      BSDF = diffuse_color * bssrdf(subsurface_method,
+      vector radius = SubsurfaceScale * SubsurfaceRadius;
-                                    Normal,
+      float subsurface_ior = (subsurface_method == "random_walk") ? SubsurfaceIOR : IOR;
-                                    Subsurface * SubsurfaceRadius,
+      closure color SubsurfBSDF = bssrdf(subsurface_method,
-                                    diffuse_color,
+                                         Normal,
-                                    "roughness",
+                                         SubsurfaceScale * SubsurfaceRadius,
-                                    Roughness,
+                                         BaseColor,
-                                    "ior",
+                                         "roughness",
-                                    SubsurfaceIOR,
+                                         Roughness,
-                                    "anisotropy",
+                                         "ior",
-                                    SubsurfaceAnisotropy);
+                                         subsurface_ior,
-    }
+                                         "anisotropy",
-    else {
+                                         SubsurfaceAnisotropy);
-      BSDF = diffuse_color * diffuse(Normal);
+      BSDF = mix(BSDF, BaseColor * SubsurfBSDF, Subsurface);
     }
 
     color f0 = color(F0_from_ior(IOR));
@@ -97,14 +102,28 @@ shader node_principled_bsdf(string distribution = "multi_ggx",
     BSDF = mix(BSDF, MetallicBSDF, clamp(Metallic, 0.0, 1.0));
   }
 
-  if (Clearcoat > 1e-5) {
+  if (EmissionStrength > 0.0 && Emission != color(0.0)) {
-    float clearcoat_r2 = ClearcoatRoughness * ClearcoatRoughness;
+    BSDF += EmissionStrength * Emission * emission();
-    closure color ClearcoatBSDF = microfacet("clearcoat", ClearcoatNormal, clearcoat_r2, 1.5, 0);
+  }
-    BSDF = layer(0.25 * Clearcoat * ClearcoatBSDF, BSDF);
+
+  if (Coat > 1e-5) {
+    float coat_ior = max(CoatIOR, 1.0);
+    if (CoatTint != color(1.0)) {
+      float coat_neta = 1.0 / coat_ior;
+      float cosNI = dot(I, CoatNormal);
+      float cosNT = sqrt(1.0 - coat_neta * coat_neta * (1 - cosNI * cosNI));
+      BSDF *= pow(CoatTint, Coat / cosNT);
+    }
+    float coat_r2 = CoatRoughness * CoatRoughness;
+    closure color CoatBSDF = dielectric_bsdf(
+        CoatNormal, vector(0.0), color(1.0), color(0.0), coat_r2, coat_r2, coat_ior, "ggx");
+    BSDF = layer(Coat * CoatBSDF, BSDF);
   }
 
   if (Sheen > 1e-5) {
     closure color SheenBSDF = sheen(Normal, SheenRoughness);
     BSDF = layer(SheenTint * Sheen * SheenBSDF, BSDF);
   }
+
+  BSDF = mix(transparent(), BSDF, Alpha);
 }

intern/cycles/kernel/svm/closure.h

@@ -54,6 +54,12 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
       return svm_node_closure_bsdf_skip(kg, offset, type);
     }
   }
+  else IF_KERNEL_NODES_FEATURE(EMISSION) {
+    if (type != CLOSURE_BSDF_PRINCIPLED_ID) {
+      /* Only principled BSDF can have emission. */
+      return svm_node_closure_bsdf_skip(kg, offset, type);
+    }
+  }
   else {
     return svm_node_closure_bsdf_skip(kg, offset, type);
   }
@@ -69,8 +75,10 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
   switch (type) {
     case CLOSURE_BSDF_PRINCIPLED_ID: {
       uint specular_offset, roughness_offset, specular_tint_offset, anisotropic_offset,
-          sheen_offset, sheen_tint_offset, clearcoat_offset, clearcoat_roughness_offset,
+          sheen_offset, sheen_tint_offset, sheen_roughness_offset, coat_offset,
-          eta_offset, transmission_offset, anisotropic_rotation_offset, pad1;
+          coat_roughness_offset, coat_ior_offset, eta_offset, transmission_offset,
+          anisotropic_rotation_offset, coat_tint_offset, coat_normal_offset, dummy, alpha_offset,
+          emission_strength_offset, emission_offset;
       uint4 data_node2 = read_node(kg, &offset);
 
       float3 T = stack_load_float3(stack, data_node.y);
@@ -79,13 +87,15 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
                              &roughness_offset,
                              &specular_tint_offset,
                              &anisotropic_offset);
-      svm_unpack_node_uchar4(data_node.w,
-                             &sheen_offset,
-                             &sheen_tint_offset,
-                             &clearcoat_offset,
-                             &clearcoat_roughness_offset);
       svm_unpack_node_uchar4(
-          data_node2.x, &eta_offset, &transmission_offset, &anisotropic_rotation_offset, &pad1);
+          data_node.w, &sheen_offset, &sheen_tint_offset, &sheen_roughness_offset, &dummy);
+      svm_unpack_node_uchar4(data_node2.x,
+                             &eta_offset,
+                             &transmission_offset,
+                             &anisotropic_rotation_offset,
+                             &coat_normal_offset);
+      svm_unpack_node_uchar4(
+          data_node2.w, &coat_offset, &coat_roughness_offset, &coat_ior_offset, &coat_tint_offset);
 
       // get Disney principled parameters
       float metallic = saturatef(param1);
@@ -96,9 +106,11 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
       float anisotropic = stack_load_float(stack, anisotropic_offset);
       float sheen = stack_load_float(stack, sheen_offset);
       float3 sheen_tint = stack_load_float3(stack, sheen_tint_offset);
-      float sheen_roughness = stack_load_float(stack, data_node2.w);
+      float sheen_roughness = stack_load_float(stack, sheen_roughness_offset);
-      float clearcoat = stack_load_float(stack, clearcoat_offset);
+      float coat = stack_load_float(stack, coat_offset);
-      float clearcoat_roughness = stack_load_float(stack, clearcoat_roughness_offset);
+      float coat_roughness = stack_load_float(stack, coat_roughness_offset);
+      float coat_ior = fmaxf(stack_load_float(stack, coat_ior_offset), 1.0f);
+      float3 coat_tint = stack_load_float3(stack, coat_tint_offset);
       float transmission = saturatef(stack_load_float(stack, transmission_offset));
       float anisotropic_rotation = stack_load_float(stack, anisotropic_rotation_offset);
       float eta = fmaxf(stack_load_float(stack, eta_offset), 1e-5f);
@@ -116,28 +128,22 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
                                           __uint_as_float(data_base_color.z),
                                           __uint_as_float(data_base_color.w));
 
-      // get the additional clearcoat normal and subsurface scattering radius
+      // get the subsurface scattering data
-      uint4 data_cn_ssr = read_node(kg, &offset);
+      uint4 data_subsurf = read_node(kg, &offset);
-      float3 clearcoat_normal = stack_valid(data_cn_ssr.x) ?
-                                    stack_load_float3(stack, data_cn_ssr.x) :
-                                    sd->N;
-      clearcoat_normal = maybe_ensure_valid_specular_reflection(sd, clearcoat_normal);
-      float3 subsurface_radius = stack_valid(data_cn_ssr.y) ?
-                                     stack_load_float3(stack, data_cn_ssr.y) :
-                                     one_float3();
-      float subsurface_ior = stack_valid(data_cn_ssr.z) ? stack_load_float(stack, data_cn_ssr.z) :
-                                                          1.4f;
-      float subsurface_anisotropy = stack_valid(data_cn_ssr.w) ?
-                                        stack_load_float(stack, data_cn_ssr.w) :
-                                        0.0f;
 
-      // get the subsurface color
+      uint4 data_alpha_emission = read_node(kg, &offset);
-      uint4 data_subsurface_color = read_node(kg, &offset);
+      svm_unpack_node_uchar4(data_alpha_emission.x,
-      float3 subsurface_color = stack_valid(data_subsurface_color.x) ?
+                             &alpha_offset,
-                                    stack_load_float3(stack, data_subsurface_color.x) :
+                             &emission_strength_offset,
-                                    make_float3(__uint_as_float(data_subsurface_color.y),
+                             &emission_offset,
-                                                __uint_as_float(data_subsurface_color.z),
+                             &dummy);
-                                                __uint_as_float(data_subsurface_color.w));
+      float alpha = stack_valid(alpha_offset) ? stack_load_float(stack, alpha_offset) :
+                                                __uint_as_float(data_alpha_emission.y);
+      float3 emission = stack_load_float3(stack, emission_offset);
+      /* Emission strength */
+      emission *= stack_valid(emission_strength_offset) ?
+                      stack_load_float(stack, emission_strength_offset) :
+                      __uint_as_float(data_alpha_emission.z);
 
       Spectrum weight = closure_weight * mix_weight;
 
@@ -161,6 +167,12 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
       const bool glass_caustics = true;
 #endif
 
+      /* Before any actual shader components, apply transparency. */
+      if (alpha < 1.0f) {
+        bsdf_transparent_setup(sd, weight * (1.0f - alpha), path_flag);
+        weight *= alpha;
+      }
+
       /* First layer: Sheen */
       if (sheen > CLOSURE_WEIGHT_CUTOFF) {
         ccl_private SheenBsdf *bsdf = (ccl_private SheenBsdf *)bsdf_alloc(
@@ -175,29 +187,68 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
 
           /* Attenuate lower layers */
           Spectrum albedo = bsdf_albedo(kg, sd, (ccl_private ShaderClosure *)bsdf, true, false);
-          weight *= 1.0f - reduce_max(albedo / weight);
+          weight *= 1.0f - reduce_max(safe_divide_color(albedo, weight));
         }
       }
 
-      /* Second layer: Clearcoat */
+      /* Second layer: Coat */
-      if (reflective_caustics && clearcoat > CLOSURE_WEIGHT_CUTOFF) {
+      if (coat > CLOSURE_WEIGHT_CUTOFF) {
-        ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
+        float3 coat_normal = stack_valid(coat_normal_offset) ?
-            sd, sizeof(MicrofacetBsdf), 0.25f * clearcoat * weight);
+                                 stack_load_float3(stack, coat_normal_offset) :
+                                 sd->N;
+        coat_normal = maybe_ensure_valid_specular_reflection(sd, coat_normal);
+        if (reflective_caustics) {
+          ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
+              sd, sizeof(MicrofacetBsdf), coat * weight);
 
-        if (bsdf) {
+          if (bsdf) {
-          bsdf->N = clearcoat_normal;
+            bsdf->N = coat_normal;
-          bsdf->T = zero_float3();
+            bsdf->T = zero_float3();
-          bsdf->ior = 1.5f;
+            bsdf->ior = coat_ior;
 
-          bsdf->alpha_x = bsdf->alpha_y = sqr(clearcoat_roughness);
+            bsdf->alpha_x = bsdf->alpha_y = sqr(coat_roughness);
 
-          /* setup bsdf */
+            /* setup bsdf */
-          sd->flag |= bsdf_microfacet_ggx_clearcoat_setup(kg, bsdf, sd);
+            sd->flag |= bsdf_microfacet_ggx_setup(bsdf);
+            bsdf_microfacet_setup_fresnel_dielectric(kg, bsdf, sd);
 
-          /* Attenuate lower layers */
+            /* Attenuate lower layers */
-          Spectrum albedo = bsdf_albedo(kg, sd, (ccl_private ShaderClosure *)bsdf, true, false);
+            Spectrum albedo = bsdf_albedo(kg, sd, (ccl_private ShaderClosure *)bsdf, true, false);
-          weight *= 1.0f - reduce_max(albedo / weight);
+            weight *= 1.0f - reduce_max(safe_divide_color(albedo, weight));
+          }
         }
+
+        if (!isequal(coat_tint, one_float3())) {
+          /* Tint is normalized to perpendicular incidence.
+           * Therefore, if we define the coat thickness as length 1, the length along the ray is
+           * t = sqrt(1+tan^2(angle(N, I))) = sqrt(1+tan^2(acos(dotNI))) = 1 / dotNI.
+           * From Beer's law, we have T = exp(-sigma_e * t).
+           * Therefore, tint = exp(-sigma_e * 1) (per def.), so -sigma_e = log(tint).
+           * From this, T = exp(log(tint) * t) = exp(log(tint)) ^ t = tint ^ t;
+           *
+           * Note that this is only an approximation - it assumes that the outgoing ray
+           * follows the same angle, and that there aren't multiple internal bounces.
+           * In particular, things that could be improved:
+           * - For transmissive materials, there should not be an outgoing path at all if the path
+           *   is transmitted.
+           * - For rough materials, we could blend towards a view-independent average path length
+           *   (e.g. 2 for diffuse reflection) for the outgoing direction.
+           * However, there's also an argument to be made for keeping parameters independent of
+           * each other for more intuitive control, in particular main roughness not affecting the
+           * coat.
+           */
+          float cosNI = dot(sd->wi, coat_normal);
+          /* Refract incoming direction into coat material.
+           * TIR is no concern here since we're always coming from the outside. */
+          float cosNT = sqrtf(1.0f - sqr(1.0f / coat_ior) * (1 - sqr(cosNI)));
+          float optical_depth = 1.0f / cosNT;
+          weight *= power(rgb_to_spectrum(coat_tint), coat * optical_depth);
+        }
+      }
+
+      /* Emission (attenuated by sheen and coat) */
+      if (!is_zero(emission)) {
+        emission_setup(sd, rgb_to_spectrum(emission) * weight);
       }
 
       /* Metallic component */
@@ -295,48 +346,42 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
 
           /* Attenuate lower layers */
           Spectrum albedo = bsdf_albedo(kg, sd, (ccl_private ShaderClosure *)bsdf, true, false);
-          weight *= 1.0f - reduce_max(albedo / weight);
+          weight *= 1.0f - reduce_max(safe_divide_color(albedo, weight));
         }
       }
 
-      /* Diffuse component */
+      /* Diffuse/Subsurface component */
-      float3 diffuse_color = mix(base_color, subsurface_color, subsurface);
 #ifdef __SUBSURFACE__
-      /* disable in case of diffuse ancestor, can't see it well then and
+      ccl_private Bssrdf *bssrdf = bssrdf_alloc(sd,
-       * adds considerably noise due to probabilities of continuing path
+                                                rgb_to_spectrum(base_color) * subsurface * weight);
-       * getting lower and lower */
+      if (bssrdf) {
-      if ((subsurface > CLOSURE_WEIGHT_CUTOFF) && !(path_flag & PATH_RAY_DIFFUSE_ANCESTOR)) {
+        float3 subsurface_radius = stack_load_float3(stack, data_subsurf.y);
-        /* Skip in case of extremely low albedo. */
+        float subsurface_scale = stack_load_float(stack, data_subsurf.z);
-        if (fabsf(average(diffuse_color)) > CLOSURE_WEIGHT_CUTOFF) {
-          ccl_private Bssrdf *bssrdf = bssrdf_alloc(sd, rgb_to_spectrum(diffuse_color) * weight);
 
-          if (bssrdf) {
+        bssrdf->radius = rgb_to_spectrum(subsurface_radius * subsurface_scale);
-            bssrdf->radius = rgb_to_spectrum(subsurface_radius * subsurface);
+        bssrdf->albedo = rgb_to_spectrum(base_color);
-            bssrdf->albedo = rgb_to_spectrum(diffuse_color);
+        bssrdf->N = N;
-            bssrdf->N = N;
+        bssrdf->alpha = sqr(roughness);
-            bssrdf->roughness = roughness;
+        bssrdf->ior = eta;
-
+        bssrdf->anisotropy = stack_load_float(stack, data_subsurf.w);
-            /* Clamps protecting against bad/extreme and non physical values. */
+        if (subsurface_method == CLOSURE_BSSRDF_RANDOM_WALK_ID) {
-            subsurface_ior = clamp(subsurface_ior, 1.01f, 3.8f);
+          bssrdf->ior = stack_load_float(stack, data_subsurf.x);
-            bssrdf->anisotropy = clamp(subsurface_anisotropy, 0.0f, 0.9f);
-
-            /* setup bsdf */
-            sd->flag |= bssrdf_setup(sd, bssrdf, subsurface_method, subsurface_ior);
-          }
         }
+
+        /* setup bsdf */
+        sd->flag |= bssrdf_setup(sd, bssrdf, path_flag, subsurface_method);
       }
-      else
+#else
+      subsurface = 0.0f;
 #endif
-      {
-        ccl_private DiffuseBsdf *bsdf = (ccl_private DiffuseBsdf *)bsdf_alloc(
-            sd, sizeof(DiffuseBsdf), rgb_to_spectrum(diffuse_color) * weight);
 
-        if (bsdf) {
+      ccl_private DiffuseBsdf *bsdf = (ccl_private DiffuseBsdf *)bsdf_alloc(
-          bsdf->N = N;
+          sd, sizeof(DiffuseBsdf), rgb_to_spectrum(base_color) * (1.0f - subsurface) * weight);
+      if (bsdf) {
+        bsdf->N = N;
 
-          /* setup bsdf */
+        /* setup bsdf */
-          sd->flag |= bsdf_diffuse_setup(bsdf);
+        sd->flag |= bsdf_diffuse_setup(bsdf);
-        }
       }
 
       break;
@@ -769,22 +814,14 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
       ccl_private Bssrdf *bssrdf = bssrdf_alloc(sd, weight);
 
       if (bssrdf) {
-        /* disable in case of diffuse ancestor, can't see it well then and
-         * adds considerably noise due to probabilities of continuing path
-         * getting lower and lower */
-        if (path_flag & PATH_RAY_DIFFUSE_ANCESTOR)
-          param1 = 0.0f;
-
         bssrdf->radius = rgb_to_spectrum(stack_load_float3(stack, data_node.z) * param1);
         bssrdf->albedo = closure_weight;
         bssrdf->N = N;
-        bssrdf->roughness = FLT_MAX;
+        bssrdf->ior = param2;
+        bssrdf->alpha = 1.0f;
+        bssrdf->anisotropy = stack_load_float(stack, data_node.w);
 
-        const float subsurface_ior = clamp(param2, 1.01f, 3.8f);
+        sd->flag |= bssrdf_setup(sd, bssrdf, path_flag, (ClosureType)type);
-        const float subsurface_anisotropy = stack_load_float(stack, data_node.w);
-        bssrdf->anisotropy = clamp(subsurface_anisotropy, 0.0f, 0.9f);
-
-        sd->flag |= bssrdf_setup(sd, bssrdf, (ClosureType)type, subsurface_ior);
       }
 
       break;

intern/cycles/kernel/svm/types.h

@@ -427,7 +427,6 @@ typedef enum ClosureType {
 
   /* Glossy */
   CLOSURE_BSDF_MICROFACET_GGX_ID,
-  CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID,
   CLOSURE_BSDF_MICROFACET_BECKMANN_ID,
   CLOSURE_BSDF_MICROFACET_MULTI_GGX_ID, /* virtual closure */
   CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID,

intern/cycles/scene/shader.cpp

@@ -133,11 +133,18 @@ static float3 output_estimate_emission(ShaderOutput *output, bool &is_constant)
     return zero_float3();
   }
   else if (node->type == EmissionNode::get_node_type() ||
-           node->type == BackgroundNode::get_node_type())
+           node->type == BackgroundNode::get_node_type() ||
+           node->type == PrincipledBsdfNode::get_node_type())
   {
+    const bool is_principled = (node->type == PrincipledBsdfNode::get_node_type());
     /* Emission and Background node. */
-    ShaderInput *color_in = node->input("Color");
+    ShaderInput *color_in = node->input(is_principled ? "Emission" : "Color");
-    ShaderInput *strength_in = node->input("Strength");
+    ShaderInput *strength_in = node->input(is_principled ? "Emission Strength" : "Strength");
+
+    if (is_principled) {
+      /* Too many parameters (coat, sheen, alpha) influence Emission for the Principled BSDF. */
+      is_constant = false;
+    }
 
     float3 estimate = one_float3();
 

intern/cycles/scene/shader_nodes.cpp

@@ -2660,10 +2660,10 @@ NODE_DEFINE(PrincipledBsdfNode)
               subsurface_method_enum,
               CLOSURE_BSSRDF_RANDOM_WALK_ID);
 
-  SOCKET_IN_COLOR(base_color, "Base Color", make_float3(0.8f, 0.8f, 0.8f));
+  SOCKET_IN_COLOR(base_color, "Base Color", make_float3(0.8f, 0.8f, 0.8f))
-  SOCKET_IN_COLOR(subsurface_color, "Subsurface Color", make_float3(0.8f, 0.8f, 0.8f));
   SOCKET_IN_FLOAT(metallic, "Metallic", 0.0f);
   SOCKET_IN_FLOAT(subsurface, "Subsurface", 0.0f);
+  SOCKET_IN_FLOAT(subsurface_scale, "Subsurface Scale", 0.1f);
   SOCKET_IN_VECTOR(subsurface_radius, "Subsurface Radius", make_float3(0.1f, 0.1f, 0.1f));
   SOCKET_IN_FLOAT(subsurface_ior, "Subsurface IOR", 1.4f);
   SOCKET_IN_FLOAT(subsurface_anisotropy, "Subsurface Anisotropy", 0.0f);
@@ -2674,16 +2674,18 @@ NODE_DEFINE(PrincipledBsdfNode)
   SOCKET_IN_FLOAT(sheen, "Sheen", 0.0f);
   SOCKET_IN_FLOAT(sheen_roughness, "Sheen Roughness", 0.5f);
   SOCKET_IN_COLOR(sheen_tint, "Sheen Tint", one_float3());
-  SOCKET_IN_FLOAT(clearcoat, "Clearcoat", 0.0f);
+  SOCKET_IN_FLOAT(coat, "Coat", 0.0f);
-  SOCKET_IN_FLOAT(clearcoat_roughness, "Clearcoat Roughness", 0.03f);
+  SOCKET_IN_FLOAT(coat_roughness, "Coat Roughness", 0.03f);
+  SOCKET_IN_FLOAT(coat_ior, "Coat IOR", 1.5f);
+  SOCKET_IN_COLOR(coat_tint, "Coat Tint", one_float3());
   SOCKET_IN_FLOAT(ior, "IOR", 0.0f);
   SOCKET_IN_FLOAT(transmission, "Transmission", 0.0f);
   SOCKET_IN_FLOAT(anisotropic_rotation, "Anisotropic Rotation", 0.0f);
-  SOCKET_IN_COLOR(emission, "Emission", zero_float3());
+  SOCKET_IN_COLOR(emission, "Emission", one_float3());
-  SOCKET_IN_FLOAT(emission_strength, "Emission Strength", 1.0f);
+  SOCKET_IN_FLOAT(emission_strength, "Emission Strength", 0.0f);
   SOCKET_IN_FLOAT(alpha, "Alpha", 1.0f);
   SOCKET_IN_NORMAL(normal, "Normal", zero_float3(), SocketType::LINK_NORMAL);
-  SOCKET_IN_NORMAL(clearcoat_normal, "Clearcoat Normal", zero_float3(), SocketType::LINK_NORMAL);
+  SOCKET_IN_NORMAL(coat_normal, "Coat Normal", zero_float3(), SocketType::LINK_NORMAL);
   SOCKET_IN_NORMAL(tangent, "Tangent", zero_float3(), SocketType::LINK_TANGENT);
   SOCKET_IN_FLOAT(surface_mix_weight, "SurfaceMixWeight", 0.0f, SocketType::SVM_INTERNAL);
 
@@ -2698,60 +2700,33 @@ PrincipledBsdfNode::PrincipledBsdfNode() : BsdfBaseNode(get_node_type())
   distribution = CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_ID;
 }
 
-void PrincipledBsdfNode::expand(ShaderGraph *graph)
+void PrincipledBsdfNode::simplify_settings(Scene * /* scene */)
 {
-  ShaderOutput *principled_out = output("BSDF");
+  if (!has_surface_emission()) {
-
+    /* Emission will be zero, so optimize away any connected emission input. */
-  ShaderInput *emission_in = input("Emission");
+    ShaderInput *emission_in = input("Emission");
-  ShaderInput *emission_strength_in = input("Emission Strength");
+    ShaderInput *strength_in = input("Emission Strength");
-  if ((emission_in->link || emission != zero_float3()) &&
-      (emission_strength_in->link || emission_strength != 0.0f))
-  {
-    /* Create add closure and emission, and relink inputs. */
-    AddClosureNode *add = graph->create_node<AddClosureNode>();
-    EmissionNode *emission_node = graph->create_node<EmissionNode>();
-    ShaderOutput *new_out = add->output("Closure");
-
-    graph->add(add);
-    graph->add(emission_node);
-
-    graph->relink(emission_strength_in, emission_node->input("Strength"));
-    graph->relink(emission_in, emission_node->input("Color"));
-    graph->relink(principled_out, new_out);
-    graph->connect(emission_node->output("Emission"), add->input("Closure1"));
-    graph->connect(principled_out, add->input("Closure2"));
-
-    principled_out = new_out;
-  }
-  else {
-    /* Disconnect unused links if the other value is zero, required before
-     * we remove the input from the node entirely. */
     if (emission_in->link) {
       emission_in->disconnect();
     }
-    if (emission_strength_in->link) {
+    if (strength_in->link) {
-      emission_strength_in->disconnect();
+      strength_in->disconnect();
     }
   }
+}
 
+bool PrincipledBsdfNode::has_surface_transparent()
+{
   ShaderInput *alpha_in = input("Alpha");
-  if (alpha_in->link || alpha != 1.0f) {
+  return (alpha_in->link != NULL || alpha < (1.0f - CLOSURE_WEIGHT_CUTOFF));
-    /* Create mix and transparent BSDF for alpha transparency. */
+}
-    MixClosureNode *mix = graph->create_node<MixClosureNode>();
-    TransparentBsdfNode *transparent = graph->create_node<TransparentBsdfNode>();
 
-    graph->add(mix);
+bool PrincipledBsdfNode::has_surface_emission()
-    graph->add(transparent);
+{
-
+  ShaderInput *emission_in = input("Emission");
-    graph->relink(alpha_in, mix->input("Fac"));
+  ShaderInput *emission_strength_in = input("Emission Strength");
-    graph->relink(principled_out, mix->output("Closure"));
+  return (emission_in->link != NULL || reduce_max(emission) > CLOSURE_WEIGHT_CUTOFF) &&
-    graph->connect(transparent->output("BSDF"), mix->input("Closure1"));
+         (emission_strength_in->link != NULL || emission_strength > CLOSURE_WEIGHT_CUTOFF);
-    graph->connect(principled_out, mix->input("Closure2"));
-  }
-
-  remove_input(emission_in);
-  remove_input(emission_strength_in);
-  remove_input(alpha_in);
 }
 
 bool PrincipledBsdfNode::has_surface_bssrdf()
@@ -2772,53 +2747,44 @@ void PrincipledBsdfNode::attributes(Shader *shader, AttributeRequestSet *attribu
   ShaderNode::attributes(shader, attributes);
 }
 
-void PrincipledBsdfNode::compile(SVMCompiler &compiler,
+void PrincipledBsdfNode::compile(SVMCompiler &compiler)
-                                 ShaderInput *p_metallic,
-                                 ShaderInput *p_subsurface,
-                                 ShaderInput *p_subsurface_radius,
-                                 ShaderInput *p_subsurface_ior,
-                                 ShaderInput *p_subsurface_anisotropy,
-                                 ShaderInput *p_specular,
-                                 ShaderInput *p_roughness,
-                                 ShaderInput *p_specular_tint,
-                                 ShaderInput *p_anisotropic,
-                                 ShaderInput *p_sheen,
-                                 ShaderInput *p_sheen_roughness,
-                                 ShaderInput *p_sheen_tint,
-                                 ShaderInput *p_clearcoat,
-                                 ShaderInput *p_clearcoat_roughness,
-                                 ShaderInput *p_ior,
-                                 ShaderInput *p_transmission,
-                                 ShaderInput *p_anisotropic_rotation)
 {
   ShaderInput *base_color_in = input("Base Color");
-  ShaderInput *subsurface_color_in = input("Subsurface Color");
+
-  ShaderInput *normal_in = input("Normal");
+  ShaderInput *p_metallic = input("Metallic");
-  ShaderInput *clearcoat_normal_in = input("Clearcoat Normal");
+  ShaderInput *p_subsurface = input("Subsurface");
-  ShaderInput *tangent_in = input("Tangent");
+
+  ShaderInput *emission_strength_in = input("Emission Strength");
+  ShaderInput *alpha_in = input("Alpha");
 
   float3 weight = one_float3();
 
   compiler.add_node(NODE_CLOSURE_SET_WEIGHT, weight);
 
-  int normal_offset = compiler.stack_assign_if_linked(normal_in);
+  int normal_offset = compiler.stack_assign_if_linked(input("Normal"));
-  int clearcoat_normal_offset = compiler.stack_assign_if_linked(clearcoat_normal_in);
+  int coat_normal_offset = compiler.stack_assign_if_linked(input("Coat Normal"));
-  int tangent_offset = compiler.stack_assign_if_linked(tangent_in);
+  int tangent_offset = compiler.stack_assign_if_linked(input("Tangent"));
-  int specular_offset = compiler.stack_assign(p_specular);
+  int specular_offset = compiler.stack_assign(input("Specular"));
-  int roughness_offset = compiler.stack_assign(p_roughness);
+  int roughness_offset = compiler.stack_assign(input("Roughness"));
-  int specular_tint_offset = compiler.stack_assign(p_specular_tint);
+  int specular_tint_offset = compiler.stack_assign(input("Specular Tint"));
-  int anisotropic_offset = compiler.stack_assign(p_anisotropic);
+  int anisotropic_offset = compiler.stack_assign(input("Anisotropic"));
-  int sheen_offset = compiler.stack_assign(p_sheen);
+  int sheen_offset = compiler.stack_assign(input("Sheen"));
-  int sheen_roughness_offset = compiler.stack_assign(p_sheen_roughness);
+  int sheen_roughness_offset = compiler.stack_assign(input("Sheen Roughness"));
-  int sheen_tint_offset = compiler.stack_assign(p_sheen_tint);
+  int sheen_tint_offset = compiler.stack_assign(input("Sheen Tint"));
-  int clearcoat_offset = compiler.stack_assign(p_clearcoat);
+  int coat_offset = compiler.stack_assign(input("Coat"));
-  int clearcoat_roughness_offset = compiler.stack_assign(p_clearcoat_roughness);
+  int coat_roughness_offset = compiler.stack_assign(input("Coat Roughness"));
-  int ior_offset = compiler.stack_assign(p_ior);
+  int coat_ior_offset = compiler.stack_assign(input("Coat IOR"));
-  int transmission_offset = compiler.stack_assign(p_transmission);
+  int coat_tint_offset = compiler.stack_assign(input("Coat Tint"));
-  int anisotropic_rotation_offset = compiler.stack_assign(p_anisotropic_rotation);
+  int ior_offset = compiler.stack_assign(input("IOR"));
-  int subsurface_radius_offset = compiler.stack_assign(p_subsurface_radius);
+  int transmission_offset = compiler.stack_assign(input("Transmission"));
-  int subsurface_ior_offset = compiler.stack_assign(p_subsurface_ior);
+  int anisotropic_rotation_offset = compiler.stack_assign(input("Anisotropic Rotation"));
-  int subsurface_anisotropy_offset = compiler.stack_assign(p_subsurface_anisotropy);
+  int subsurface_radius_offset = compiler.stack_assign(input("Subsurface Radius"));
+  int subsurface_scale_offset = compiler.stack_assign(input("Subsurface Scale"));
+  int subsurface_ior_offset = compiler.stack_assign(input("Subsurface IOR"));
+  int subsurface_anisotropy_offset = compiler.stack_assign(input("Subsurface Anisotropy"));
+  int alpha_offset = compiler.stack_assign_if_linked(alpha_in);
+  int emission_strength_offset = compiler.stack_assign_if_linked(emission_strength_in);
+  int emission_offset = compiler.stack_assign(input("Emission"));
 
   compiler.add_node(NODE_CLOSURE_BSDF,
                     compiler.encode_uchar4(closure,
@@ -2834,14 +2800,15 @@ void PrincipledBsdfNode::compile(SVMCompiler &compiler,
       compiler.encode_uchar4(
           specular_offset, roughness_offset, specular_tint_offset, anisotropic_offset),
       compiler.encode_uchar4(
-          sheen_offset, sheen_tint_offset, clearcoat_offset, clearcoat_roughness_offset));
+          sheen_offset, sheen_tint_offset, sheen_roughness_offset, SVM_STACK_INVALID));
 
   compiler.add_node(
       compiler.encode_uchar4(
-          ior_offset, transmission_offset, anisotropic_rotation_offset, SVM_STACK_INVALID),
+          ior_offset, transmission_offset, anisotropic_rotation_offset, coat_normal_offset),
       distribution,
       subsurface_method,
-      sheen_roughness_offset);
+      compiler.encode_uchar4(
+          coat_offset, coat_roughness_offset, coat_ior_offset, coat_tint_offset));
 
   float3 bc_default = get_float3(base_color_in->socket_type);
 
@@ -2851,40 +2818,17 @@ void PrincipledBsdfNode::compile(SVMCompiler &compiler,
       __float_as_int(bc_default.y),
       __float_as_int(bc_default.z));
 
-  compiler.add_node(clearcoat_normal_offset,
+  compiler.add_node(subsurface_ior_offset,
                     subsurface_radius_offset,
-                    subsurface_ior_offset,
+                    subsurface_scale_offset,
                     subsurface_anisotropy_offset);
 
-  float3 ss_default = get_float3(subsurface_color_in->socket_type);
+  compiler.add_node(
-
+      compiler.encode_uchar4(
-  compiler.add_node(((subsurface_color_in->link) ? compiler.stack_assign(subsurface_color_in) :
+          alpha_offset, emission_strength_offset, emission_offset, SVM_STACK_INVALID),
-                                                   SVM_STACK_INVALID),
+      __float_as_int(get_float(alpha_in->socket_type)),
-                    __float_as_int(ss_default.x),
+      __float_as_int(get_float(emission_strength_in->socket_type)),
-                    __float_as_int(ss_default.y),
+      SVM_STACK_INVALID);
-                    __float_as_int(ss_default.z));
-}
-
-void PrincipledBsdfNode::compile(SVMCompiler &compiler)
-{
-  compile(compiler,
-          input("Metallic"),
-          input("Subsurface"),
-          input("Subsurface Radius"),
-          input("Subsurface IOR"),
-          input("Subsurface Anisotropy"),
-          input("Specular"),
-          input("Roughness"),
-          input("Specular Tint"),
-          input("Anisotropic"),
-          input("Sheen"),
-          input("Sheen Roughness"),
-          input("Sheen Tint"),
-          input("Clearcoat"),
-          input("Clearcoat Roughness"),
-          input("IOR"),
-          input("Transmission"),
-          input("Anisotropic Rotation"));
 }
 
 void PrincipledBsdfNode::compile(OSLCompiler &compiler)

intern/cycles/scene/shader_nodes.h

@@ -514,31 +514,13 @@ class PrincipledBsdfNode : public BsdfBaseNode {
  public:
   SHADER_NODE_CLASS(PrincipledBsdfNode)
 
-  void expand(ShaderGraph *graph);
   bool has_surface_bssrdf();
   bool has_bssrdf_bump();
-  void compile(SVMCompiler &compiler,
+  void simplify_settings(Scene *scene);
-               ShaderInput *metallic,
-               ShaderInput *subsurface,
-               ShaderInput *subsurface_radius,
-               ShaderInput *subsurface_ior,
-               ShaderInput *subsurface_anisotropy,
-               ShaderInput *specular,
-               ShaderInput *roughness,
-               ShaderInput *specular_tint,
-               ShaderInput *anisotropic,
-               ShaderInput *sheen,
-               ShaderInput *sheen_roughness,
-               ShaderInput *sheen_tint,
-               ShaderInput *clearcoat,
-               ShaderInput *clearcoat_roughness,
-               ShaderInput *ior,
-               ShaderInput *transmission,
-               ShaderInput *anisotropic_rotation);
 
   NODE_SOCKET_API(float3, base_color)
-  NODE_SOCKET_API(float3, subsurface_color)
   NODE_SOCKET_API(float3, subsurface_radius)
+  NODE_SOCKET_API(float, subsurface_scale)
   NODE_SOCKET_API(float, subsurface_ior)
   NODE_SOCKET_API(float, subsurface_anisotropy)
   NODE_SOCKET_API(float, metallic)
@@ -550,13 +532,15 @@ class PrincipledBsdfNode : public BsdfBaseNode {
   NODE_SOCKET_API(float, sheen)
   NODE_SOCKET_API(float, sheen_roughness)
   NODE_SOCKET_API(float3, sheen_tint)
-  NODE_SOCKET_API(float, clearcoat)
+  NODE_SOCKET_API(float, coat)
-  NODE_SOCKET_API(float, clearcoat_roughness)
+  NODE_SOCKET_API(float, coat_roughness)
+  NODE_SOCKET_API(float, coat_ior)
+  NODE_SOCKET_API(float3, coat_tint)
   NODE_SOCKET_API(float, ior)
   NODE_SOCKET_API(float, transmission)
   NODE_SOCKET_API(float, anisotropic_rotation)
   NODE_SOCKET_API(float3, normal)
-  NODE_SOCKET_API(float3, clearcoat_normal)
+  NODE_SOCKET_API(float3, coat_normal)
   NODE_SOCKET_API(float3, tangent)
   NODE_SOCKET_API(float, surface_mix_weight)
   NODE_SOCKET_API(ClosureType, distribution)
@@ -571,6 +555,8 @@ class PrincipledBsdfNode : public BsdfBaseNode {
   {
     return true;
   }
+  bool has_surface_transparent();
+  bool has_surface_emission();
 };
 
 class TranslucentBsdfNode : public BsdfNode {

intern/ghost/CMakeLists.txt

@@ -3,7 +3,7 @@
 # SPDX-License-Identifier: GPL-2.0-or-later
 
 set(INC
-  .
+  PUBLIC .
   ../clog
   ../../source/blender/imbuf
 )
@@ -89,8 +89,8 @@ if(WITH_VULKAN_BACKEND)
   )
 
   list(APPEND INC_SYS
-    ${VULKAN_INCLUDE_DIRS}
+    PUBLIC ${VULKAN_INCLUDE_DIRS}
-    ${MOLTENVK_INCLUDE_DIRS}
+    PUBLIC ${MOLTENVK_INCLUDE_DIRS}
   )
 
   list(APPEND LIB

intern/ghost/intern/GHOST_ContextWGL.cc

@@ -65,8 +65,9 @@ GHOST_ContextWGL::~GHOST_ContextWGL()
 
       s_sharedCount--;
 
-      if (s_sharedCount == 0)
+      if (s_sharedCount == 0) {
         s_sharedHGLRC = nullptr;
+      }
 
       WIN32_CHK(::wglDeleteContext(m_hGLRC));
     }
@@ -88,10 +89,12 @@ GHOST_TSuccess GHOST_ContextWGL::swapBuffers()
 
 GHOST_TSuccess GHOST_ContextWGL::setSwapInterval(int interval)
 {
-  if (epoxy_has_wgl_extension(m_hDC, "WGL_EXT_swap_control"))
+  if (epoxy_has_wgl_extension(m_hDC, "WGL_EXT_swap_control")) {
     return WIN32_CHK(::wglSwapIntervalEXT(interval)) == TRUE ? GHOST_kSuccess : GHOST_kFailure;
-  else
+  }
+  else {
     return GHOST_kFailure;
+  }
 }
 
 GHOST_TSuccess GHOST_ContextWGL::getSwapInterval(int &intervalOut)
@@ -149,8 +152,9 @@ static int weight_pixel_format(PIXELFORMATDESCRIPTOR &pfd, PIXELFORMATDESCRIPTOR
 
   weight += pfd.cColorBits - 8;
 
-  if (preferredPFD.cAlphaBits > 0 && pfd.cAlphaBits > 0)
+  if (preferredPFD.cAlphaBits > 0 && pfd.cAlphaBits > 0) {
     weight++;
+  }
 #ifdef WIN32_COMPOSITING
   if ((preferredPFD.dwFlags & PFD_SUPPORT_COMPOSITION) && (pfd.dwFlags & PFD_SUPPORT_COMPOSITION))
     weight++;
@@ -200,8 +204,9 @@ static int choose_pixel_format_legacy(HDC hDC, PIXELFORMATDESCRIPTOR &preferredP
   }
 
   /* choose any available stereo format over a non-stereo format */
-  if (iStereoPixelFormat != 0)
+  if (iStereoPixelFormat != 0) {
     iPixelFormat = iStereoPixelFormat;
+  }
 
   if (iPixelFormat == 0) {
     fprintf(stderr, "Warning! Using result of ChoosePixelFormat.\n");
@@ -372,36 +377,44 @@ struct DummyContextWGL {
 
     dummyPixelFormat = choose_pixel_format_legacy(hDC, preferredPFD);
 
-    if (dummyPixelFormat == 0)
+    if (dummyPixelFormat == 0) {
       return;
+    }
 
     PIXELFORMATDESCRIPTOR chosenPFD;
     if (!WIN32_CHK(::DescribePixelFormat(
             hDC, dummyPixelFormat, sizeof(PIXELFORMATDESCRIPTOR), &chosenPFD)))
+    {
       return;
+    }
 
     if (hWnd) {
       dummyHWND = clone_window(hWnd, nullptr);
 
-      if (dummyHWND == nullptr)
+      if (dummyHWND == nullptr) {
         return;
+      }
 
       dummyHDC = GetDC(dummyHWND);
     }
 
-    if (!WIN32_CHK(dummyHDC != nullptr))
+    if (!WIN32_CHK(dummyHDC != nullptr)) {
       return;
+    }
 
-    if (!WIN32_CHK(::SetPixelFormat(dummyHDC, dummyPixelFormat, &chosenPFD)))
+    if (!WIN32_CHK(::SetPixelFormat(dummyHDC, dummyPixelFormat, &chosenPFD))) {
       return;
+    }
 
     dummyHGLRC = ::wglCreateContext(dummyHDC);
 
-    if (!WIN32_CHK(dummyHGLRC != nullptr))
+    if (!WIN32_CHK(dummyHGLRC != nullptr)) {
       return;
+    }
 
-    if (!WIN32_CHK(::wglMakeCurrent(dummyHDC, dummyHGLRC)))
+    if (!WIN32_CHK(::wglMakeCurrent(dummyHDC, dummyHGLRC))) {
       return;
+    }
 
     has_WGL_ARB_pixel_format = epoxy_has_wgl_extension(hDC, "WGL_ARB_pixel_format");
     has_WGL_ARB_create_context = epoxy_has_wgl_extension(hDC, "WGL_ARB_create_context");
@@ -523,8 +536,9 @@ GHOST_TSuccess GHOST_ContextWGL::initializeDrawingContext()
     if (!dummy.has_WGL_ARB_create_context || ::GetPixelFormat(m_hDC) == 0) {
       int iPixelFormat = 0;
 
-      if (dummy.has_WGL_ARB_pixel_format)
+      if (dummy.has_WGL_ARB_pixel_format) {
         iPixelFormat = choose_pixel_format_arb(m_stereoVisual, needAlpha);
+      }
 
       if (iPixelFormat == 0)
         iPixelFormat = choose_pixel_format_legacy(m_hDC, dummy.preferredPFD);
@@ -541,8 +555,9 @@ GHOST_TSuccess GHOST_ContextWGL::initializeDrawingContext()
         goto error;
       }
 
-      if (needAlpha && chosenPFD.cAlphaBits == 0)
+      if (needAlpha && chosenPFD.cAlphaBits == 0) {
         fprintf(stderr, "Warning! Unable to find a pixel format with an alpha channel.\n");
+      }
 
       if (!WIN32_CHK(::SetPixelFormat(m_hDC, iPixelFormat, &chosenPFD))) {
         goto error;
@@ -553,22 +568,27 @@ GHOST_TSuccess GHOST_ContextWGL::initializeDrawingContext()
       int profileBitCore = m_contextProfileMask & WGL_CONTEXT_CORE_PROFILE_BIT_ARB;
       int profileBitCompat = m_contextProfileMask & WGL_CONTEXT_COMPATIBILITY_PROFILE_BIT_ARB;
 
-      if (!dummy.has_WGL_ARB_create_context_profile && profileBitCore)
+      if (!dummy.has_WGL_ARB_create_context_profile && profileBitCore) {
         fprintf(stderr, "Warning! OpenGL core profile not available.\n");
+      }
 
-      if (!dummy.has_WGL_ARB_create_context_profile && profileBitCompat)
+      if (!dummy.has_WGL_ARB_create_context_profile && profileBitCompat) {
         fprintf(stderr, "Warning! OpenGL compatibility profile not available.\n");
+      }
 
       int profileMask = 0;
 
-      if (dummy.has_WGL_ARB_create_context_profile && profileBitCore)
+      if (dummy.has_WGL_ARB_create_context_profile && profileBitCore) {
         profileMask |= profileBitCore;
+      }
 
-      if (dummy.has_WGL_ARB_create_context_profile && profileBitCompat)
+      if (dummy.has_WGL_ARB_create_context_profile && profileBitCompat) {
         profileMask |= profileBitCompat;
+      }
 
-      if (profileMask != m_contextProfileMask)
+      if (profileMask != m_contextProfileMask) {
         fprintf(stderr, "Warning! Ignoring untested OpenGL context profile mask bits.");
+      }
 
       std::vector<int> iAttributes;
 

intern/ghost/intern/GHOST_DisplayManagerWin32.cc

@@ -41,8 +41,9 @@ GHOST_TSuccess GHOST_DisplayManagerWin32::getNumDisplaySettings(uint8_t display,
    * function was called with #iModeNum set to zero. */
 
   DISPLAY_DEVICE display_device;
-  if (!get_dd(display, &display_device))
+  if (!get_dd(display, &display_device)) {
     return GHOST_kFailure;
+  }
 
   numSettings = 0;
   DEVMODE dm;
@@ -57,8 +58,9 @@ GHOST_TSuccess GHOST_DisplayManagerWin32::getDisplaySetting(uint8_t display,
                                                             GHOST_DisplaySetting &setting) const
 {
   DISPLAY_DEVICE display_device;
-  if (!get_dd(display, &display_device))
+  if (!get_dd(display, &display_device)) {
     return GHOST_kFailure;
+  }
 
   GHOST_TSuccess success;
   DEVMODE dm;
@@ -102,8 +104,9 @@ GHOST_TSuccess GHOST_DisplayManagerWin32::setCurrentDisplaySetting(
     uint8_t display, const GHOST_DisplaySetting &setting)
 {
   DISPLAY_DEVICE display_device;
-  if (!get_dd(display, &display_device))
+  if (!get_dd(display, &display_device)) {
     return GHOST_kFailure;
+  }
 
   GHOST_DisplaySetting match;
   findMatch(display, setting, match);

intern/ghost/intern/GHOST_SystemPathsWin32.cc

@@ -130,8 +130,9 @@ void GHOST_SystemPathsWin32::addToSystemRecentFiles(const char *filepath) const
   IShellItem *shell_item;
 
   HRESULT hr = CoInitializeEx(nullptr, COINIT_APARTMENTTHREADED | COINIT_DISABLE_OLE1DDE);
-  if (!SUCCEEDED(hr))
+  if (!SUCCEEDED(hr)) {
     return;
+  }
 
   hr = SHCreateItemFromParsingName(filepath_16, nullptr, IID_PPV_ARGS(&shell_item));
   if (SUCCEEDED(hr)) {

intern/ghost/intern/GHOST_SystemX11.cc

@@ -2734,8 +2734,9 @@ void GHOST_SystemX11::refreshXInputDevices()
 void GHOST_SystemX11::clearXInputDevices()
 {
   for (GHOST_TabletX11 &xtablet : m_xtablets) {
-    if (xtablet.Device)
+    if (xtablet.Device) {
       XCloseDevice(m_display, xtablet.Device);
+    }
   }
 
   m_xtablets.clear();

intern/ghost/intern/GHOST_WindowWin32.cc

@@ -671,8 +671,9 @@ void GHOST_WindowWin32::updateMouseCapture(GHOST_MouseCaptureEventWin32 event)
       m_nPressedButtons++;
       break;
     case MouseReleased:
-      if (m_nPressedButtons)
+      if (m_nPressedButtons) {
         m_nPressedButtons--;
+      }
       break;
     case OperatorGrab:
       m_hasGrabMouse = true;
@@ -821,12 +822,14 @@ HCURSOR GHOST_WindowWin32::getStandardCursor(GHOST_TStandardCursor shape) const
 void GHOST_WindowWin32::loadCursor(bool visible, GHOST_TStandardCursor shape) const
 {
   if (!visible) {
-    while (::ShowCursor(FALSE) >= 0)
+    while (::ShowCursor(FALSE) >= 0) {
-      ;
+      /* Pass. */
+    }
   }
   else {
-    while (::ShowCursor(TRUE) < 0)
+    while (::ShowCursor(TRUE) < 0) {
-      ;
+      /* Pass. */
+    }
   }
 
   HCURSOR cursor = getStandardCursor(shape);

intern/ghost/intern/GHOST_WindowX11.cc

@@ -402,8 +402,9 @@ bool GHOST_WindowX11::createX11_XIC()
                     XNDestroyCallback,
                     &destroy,
                     nullptr);
-  if (!m_xic)
+  if (!m_xic) {
     return false;
+  }
 
   ulong fevent;
   XGetICValues(m_xic, XNFilterEvents, &fevent, nullptr);

intern/ghost/test/gears/GHOST_C-Test.c

@@ -397,8 +397,9 @@ bool processEvent(GHOST_EventHandle hEvent, GHOST_TUserDataPtr userData)
       break;
     case GHOST_kEventWindowUpdate: {
       GHOST_WindowHandle window2 = GHOST_GetEventWindow(hEvent);
-      if (!GHOST_ValidWindow(shSystem, window2))
+      if (!GHOST_ValidWindow(shSystem, window2)) {
         break;
+      }
       setViewPortGL(window2);
       drawGL();
       GHOST_SwapWindowBuffers(window2);

intern/ghost/test/gears/GHOST_Test.cpp

@@ -602,8 +602,9 @@ bool Application::processEvent(GHOST_IEvent *event)
 
     case GHOST_kEventWindowUpdate: {
       GHOST_IWindow *window2 = event->getWindow();
-      if (!m_system->validWindow(window2))
+      if (!m_system->validWindow(window2)) {
         break;
+      }
 
       glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
 

intern/ghost/test/multitest/MultiTest.c

@@ -78,10 +78,12 @@ void rect_bevel_side(int rect[2][2], int side, float *lt, float *dk, const float
       int x = (corner == 0 || corner == 1) ? (rect[0][0] + i) : (rect[1][0] - i - 1);
       int y = (corner == 0 || corner == 3) ? (rect[0][1] + i) : (rect[1][1] - i - 1);
 
-      if (ltidx == corner)
+      if (ltidx == corner) {
         glColor3f(col[0] * ltf, col[1] * ltf, col[2] * ltf);
-      if (dkidx == corner)
+      }
+      if (dkidx == corner) {
         glColor3f(col[0] * dkf, col[1] * dkf, col[2] * dkf);
+      }
 
       glVertex2i(lx, ly);
       glVertex2i(lx = x, ly = y);
@@ -513,8 +515,9 @@ static void loggerwindow_do_key(LoggerWindow *lw, GHOST_TKey key, int press)
 {
   switch (key) {
     case GHOST_kKeyQ:
-      if (press)
+      if (press) {
         multitestapp_exit(lw->app);
+      }
       break;
   }
 }
@@ -685,8 +688,9 @@ static void extrawindow_do_key(ExtraWindow *ew, GHOST_TKey key, int press)
 {
   switch (key) {
     case GHOST_kKeyE:
-      if (press)
+      if (press) {
         multitestapp_toggle_extra_window(ew->app);
+      }
       break;
   }
 }
@@ -862,19 +866,23 @@ MultiTestApp *multitestapp_new(void)
   GHOST_EventConsumerHandle consumer = GHOST_CreateEventConsumer(multitest_event_handler, app);
 
   app->sys = GHOST_CreateSystem();
-  if (!app->sys)
+  if (!app->sys) {
     fatal("Unable to create ghost system");
+  }
 
-  if (!GHOST_AddEventConsumer(app->sys, consumer))
+  if (!GHOST_AddEventConsumer(app->sys, consumer)) {
     fatal("Unable to add multitest event consumer ");
+  }
 
   app->main = mainwindow_new(app);
-  if (!app->main)
+  if (!app->main) {
     fatal("Unable to create main window");
+  }
 
   app->logger = loggerwindow_new(app);
-  if (!app->logger)
+  if (!app->logger) {
     fatal("Unable to create logger window");
+  }
 
   app->extra = NULL;
   app->exit = 0;

intern/guardedalloc/cpp/mallocn.cpp

@@ -34,12 +34,14 @@ void *operator new[](size_t size)
 void operator delete(void *p) throw()
 {
   /* delete NULL is valid in c++ */
-  if (p)
+  if (p) {
     MEM_freeN(p);
+  }
 }
 void operator delete[](void *p) throw()
 {
   /* delete NULL is valid in c++ */
-  if (p)
+  if (p) {
     MEM_freeN(p);
+  }
 }

intern/guardedalloc/test/simpletest/memtest.c

@@ -39,8 +39,9 @@ int main(int argc, char *argv[])
   switch (argc) {
     case 2:
       verbose = atoi(argv[1]);
-      if (verbose < 0)
+      if (verbose < 0) {
         verbose = 0;
+      }
       break;
     case 1:
     default:
@@ -59,15 +60,17 @@ int main(int argc, char *argv[])
   for (i = 0; i < NUM_BLOCKS; i++) {
     int blocksize = 10000;
     char tagstring[1000];
-    if (verbose > 1)
+    if (verbose > 1) {
       printf("|--* Allocating block %d\n", i);
+    }
     sprintf(tagstring, "Memblock no. %d : ", i);
     p[i] = MEM_callocN(blocksize, strdup(tagstring));
   }
 
   /* report on that */
-  if (verbose > 1)
+  if (verbose > 1) {
     MEM_printmemlist();
+  }
 
   /* memory is there: test it */
   error_status = MEM_consistency_check();
@@ -94,16 +97,18 @@ int main(int argc, char *argv[])
   for (i = 0; i < NUM_BLOCKS; i++) {
     int blocksize = 10000;
     char tagstring[1000];
-    if (verbose > 1)
+    if (verbose > 1) {
       printf("|--* Allocating block %d\n", i);
+    }
     sprintf(tagstring, "Memblock no. %d : ", i);
     p[i] = MEM_callocN(blocksize, strdup(tagstring));
   }
 
   /* Now corrupt a few blocks. */
   ip = (int *)p[5] - 50;
-  for (i = 0; i < 1000; i++, ip++)
+  for (i = 0; i < 1000; i++, ip++) {
     *ip = i + 1;
+  }
   ip = (int *)p[6];
   *(ip + 10005) = 0;
 

scripts/presets/keyconfig/keymap_data/blender_default.py

@@ -1893,7 +1893,7 @@ def km_graph_editor(params):
         ("graph.interpolation_type", {"type": 'T', "value": 'PRESS'}, None),
         ("graph.easing_type", {"type": 'E', "value": 'PRESS', "ctrl": True}, None),
         ("graph.smooth", {"type": 'O', "value": 'PRESS', "alt": True}, None),
-        ("graph.sample", {"type": 'O', "value": 'PRESS', "shift": True, "alt": True}, None),
+        ("graph.bake_keys", {"type": 'O', "value": 'PRESS', "shift": True, "alt": True}, None),
         ("graph.keys_to_samples", {"type": 'C', "value": 'PRESS', "alt": True}, None),
         op_menu("GRAPH_MT_delete", {"type": 'X', "value": 'PRESS'}),
         ("graph.delete", {"type": 'DEL', "value": 'PRESS'}, {"properties": [("confirm", False)]}),
@@ -2552,7 +2552,7 @@ def km_dopesheet(params):
         ("action.extrapolation_type", {"type": 'E', "value": 'PRESS', "shift": True}, None),
         ("action.easing_type", {"type": 'E', "value": 'PRESS', "ctrl": True}, None),
         ("action.keyframe_type", {"type": 'R', "value": 'PRESS'}, None),
-        ("action.sample", {"type": 'O', "value": 'PRESS', "shift": True, "alt": True}, None),
+        ("action.bake_keys", {"type": 'O', "value": 'PRESS', "shift": True, "alt": True}, None),
         op_menu("DOPESHEET_MT_delete", {"type": 'X', "value": 'PRESS'}),
         ("action.delete", {"type": 'DEL', "value": 'PRESS'}, {"properties": [("confirm", False)]}),
         ("action.duplicate_move", {"type": 'D', "value": 'PRESS', "shift": True}, None),
@@ -5922,6 +5922,7 @@ def km_edit_font(params):
         ("font.text_cut", {"type": 'X', "value": 'PRESS', "ctrl": True}, None),
         ("font.text_paste", {"type": 'V', "value": 'PRESS', "ctrl": True, "repeat": True}, None),
         ("font.line_break", {"type": 'RET', "value": 'PRESS', "repeat": True}, None),
+        ("font.line_break", {"type": 'NUMPAD_ENTER', "value": 'PRESS', "repeat": True}, None),
         ("font.text_insert", {"type": 'TEXTINPUT', "value": 'ANY', "any": True, "repeat": True}, None),
         ("font.text_insert", {"type": 'BACK_SPACE', "value": 'PRESS', "alt": True, "repeat": True},
          {"properties": [("accent", True)]}),

scripts/presets/keyconfig/keymap_data/industry_compatible_data.py

@@ -1064,7 +1064,6 @@ def km_node_generic(_params):
 
     items.extend([
         op_panel("TOPBAR_PT_name", {"type": 'RET', "value": 'PRESS'}, [("keep_open", False)]),
-        ("node.add_search", {"type": 'TAB', "value": 'PRESS'}, None),
     ])
 
     return keymap

scripts/startup/bl_operators/geometry_nodes.py

@@ -87,17 +87,16 @@ def edit_geometry_nodes_modifier_poll(context):
 
 def socket_idname_to_attribute_type(idname):
     if idname.startswith("NodeSocketInt"):
-        return "INT"
+        return 'INT'
     elif idname.startswith("NodeSocketColor"):
-        return "FLOAT_COLOR"
+        return 'FLOAT_COLOR'
     elif idname.startswith("NodeSocketVector"):
-        return "FLOAT_VECTOR"
+        return 'FLOAT_VECTOR'
     elif idname.startswith("NodeSocketBool"):
-        return "BOOLEAN"
+        return 'BOOLEAN'
     elif idname.startswith("NodeSocketFloat"):
-        return "FLOAT"
+        return 'FLOAT'
     raise ValueError("Unsupported socket type")
-    return ""
 
 
 def modifier_attribute_name_get(modifier, identifier):
@@ -179,7 +178,7 @@ class MoveModifierToNodes(Operator):
                     first_geometry_input = group_node_input
 
         if not first_geometry_input:
-            self.report({"WARNING"}, "Node group must have a geometry input")
+            self.report({'WARNING'}, "Node group must have a geometry input")
             return {'CANCELLED'}
         group.links.new(group_input_node.outputs[0], first_geometry_input)
 
@@ -226,8 +225,8 @@ class MoveModifierToNodes(Operator):
             group.links.new(store_nodes[-1].outputs["Geometry"], group_output_node.inputs[data_("Geometry")])
         else:
             if not first_geometry_output:
-                self.report({"WARNING"}, "Node group must have a geometry output")
+                self.report({'WARNING'}, "Node group must have a geometry output")
-                return {"CANCELLED"}
+                return {'CANCELLED'}
             group.links.new(first_geometry_output, group_output_node.inputs[data_("Geometry")])
 
         modifier.node_group = group

scripts/startup/bl_operators/wm.py

@@ -3446,8 +3446,10 @@ class WM_MT_region_toggle_pie(Menu):
         # Use to access the labels.
         enum_items = bpy.types.Region.bl_rna.properties["type"].enum_items_static_ui
 
-        # Remove empty items.
+        # Prefer a 4 item pie menu where possible as there are some differences
-        items[:] = [item for item in items if item != []]
+        # when a pie menu has more than 4 items, see: #112129.
+        if not items_overflow and not any(items[4:]):
+            del items[4:]
 
         for region_type_list in (items + items_overflow):
             if not region_type_list:

scripts/startup/bl_ui/properties_data_lightprobe.py

@@ -107,8 +107,8 @@ class DATA_PT_lightprobe_eevee_next(DataButtonsPanel, Panel):
 
             col.separator()
 
-            col.operator("object.lightprobe_cache_bake").subset = "ACTIVE"
+            col.operator("object.lightprobe_cache_bake").subset = 'ACTIVE'
-            col.operator("object.lightprobe_cache_free").subset = "ACTIVE"
+            col.operator("object.lightprobe_cache_free").subset = 'ACTIVE'
 
             col.separator()
 

scripts/startup/bl_ui/properties_data_mesh.py

@@ -574,8 +574,8 @@ class ColorAttributesListBase():
 
         for idx, item in enumerate(attributes):
             skip = (
-                (item.domain not in {"POINT", "CORNER"}) or
+                (item.domain not in {'POINT', 'CORNER'}) or
-                (item.data_type not in {"FLOAT_COLOR", "BYTE_COLOR"}) or
+                (item.data_type not in {'FLOAT_COLOR', 'BYTE_COLOR'}) or
                 item.is_internal
             )
             flags[idx] = 0 if skip else flags[idx]

scripts/startup/bl_ui/properties_data_modifier.py

@@ -229,7 +229,7 @@ class AddModifierMenu(Operator):
     def poll(cls, context):
         # NOTE: This operator only exists to add a poll to the add modifier shortcut in the property editor.
         space = context.space_data
-        return space and space.type == 'PROPERTIES' and space.context == "MODIFIER"
+        return space and space.type == 'PROPERTIES' and space.context == 'MODIFIER'
 
     def invoke(self, context, event):
         return bpy.ops.wm.call_menu(name="OBJECT_MT_modifier_add")

scripts/startup/bl_ui/properties_render.py

@@ -651,7 +651,7 @@ class RENDER_PT_eevee_next_raytracing_reflection(EeveeRaytracingOptionsPanel):
 
     def draw_header(self, context):
         layout = self.layout
-        if context.scene.eevee.ray_split_settings == "UNIFIED":
+        if context.scene.eevee.ray_split_settings == 'UNIFIED':
             layout.label(text="Reflection & Refraction")
         else:
             layout.label(text="Reflection")
@@ -683,7 +683,7 @@ class RENDER_PT_eevee_next_raytracing_refraction(EeveeRaytracingOptionsPanel):
 
     @classmethod
     def poll(cls, context):
-        return (context.scene.eevee.ray_split_settings == "SPLIT")
+        return (context.scene.eevee.ray_split_settings == 'SPLIT')
 
     def draw(self, context):
         self.draw_internal(context, context.scene.eevee.refraction_options)
@@ -858,8 +858,8 @@ class RENDER_PT_eevee_next_indirect_lighting(RenderButtonsPanel, Panel):
         props = scene.eevee
 
         col = layout.column()
-        col.operator("object.lightprobe_cache_bake", text="Bake Light Caches", icon='RENDER_STILL').subset = "ALL"
+        col.operator("object.lightprobe_cache_bake", text="Bake Light Caches", icon='RENDER_STILL').subset = 'ALL'
-        col.operator("object.lightprobe_cache_free", text="Delete Light Caches").subset = "ALL"
+        col.operator("object.lightprobe_cache_free", text="Delete Light Caches").subset = 'ALL'
 
         col.prop(props, "gi_irradiance_pool_size", text="Pool Size")
 

scripts/startup/bl_ui/properties_view_layer.py

@@ -78,6 +78,31 @@ class VIEWLAYER_PT_eevee_layer_passes_data(ViewLayerButtonsPanel, Panel):
         col.prop(view_layer, "use_pass_normal")
 
 
+class VIEWLAYER_PT_eevee_next_layer_passes_data(ViewLayerButtonsPanel, Panel):
+    bl_label = "Data"
+    bl_parent_id = "VIEWLAYER_PT_layer_passes"
+
+    COMPAT_ENGINES = {'BLENDER_EEVEE_NEXT'}
+
+    def draw(self, context):
+        layout = self.layout
+        layout.use_property_split = True
+        layout.use_property_decorate = False
+
+        scene = context.scene
+        view_layer = context.view_layer
+
+        col = layout.column()
+        col.prop(view_layer, "use_pass_combined")
+        col.prop(view_layer, "use_pass_z")
+        col.prop(view_layer, "use_pass_mist")
+        col.prop(view_layer, "use_pass_normal")
+        col.prop(view_layer, "use_pass_position")
+        sub = col.column()
+        sub.active = not scene.eevee.use_motion_blur
+        sub.prop(view_layer, "use_pass_vector")
+
+
 class VIEWLAYER_PT_workbench_layer_passes_data(ViewLayerButtonsPanel, Panel):
     bl_label = "Data"
     bl_parent_id = "VIEWLAYER_PT_layer_passes"
@@ -269,6 +294,7 @@ classes = (
     VIEWLAYER_PT_layer_passes,
     VIEWLAYER_PT_workbench_layer_passes_data,
     VIEWLAYER_PT_eevee_layer_passes_data,
+    VIEWLAYER_PT_eevee_next_layer_passes_data,
     VIEWLAYER_PT_eevee_layer_passes_light,
     VIEWLAYER_PT_eevee_layer_passes_effects,
     VIEWLAYER_PT_layer_passes_cryptomatte,

scripts/startup/bl_ui/space_dopesheet.py

@@ -557,7 +557,7 @@ class DOPESHEET_MT_key(Menu):
         layout.separator()
         layout.operator("action.clean").channels = False
         layout.operator("action.clean", text="Clean Channels").channels = True
-        layout.operator("action.sample")
+        layout.operator("action.bake_keys")
 
         layout.separator()
         layout.operator("graph.euler_filter", text="Discontinuity (Euler) Filter")

scripts/startup/bl_ui/space_graph.py

@@ -285,7 +285,7 @@ class GRAPH_MT_channel(Menu):
         layout.separator()
         layout.operator("graph.keys_to_samples")
         layout.operator("graph.samples_to_keys")
-        layout.operator("graph.sound_bake")
+        layout.operator("graph.sound_to_samples")
 
         layout.separator()
         layout.operator("graph.euler_filter", text="Discontinuity (Euler) Filter")
@@ -305,7 +305,7 @@ class GRAPH_MT_key_density(Menu):
         # as we do not have a modal mode for it, so just execute it.
         with operator_context(layout, 'EXEC_REGION_WIN'):
             layout.operator("graph.decimate", text="Decimate (Allowed Change)").mode = 'ERROR'
-        layout.operator("graph.sample")
+        layout.operator("graph.bake_keys")
 
         layout.separator()
         layout.operator("graph.clean").channels = False

scripts/startup/bl_ui/space_outliner.py

@@ -344,7 +344,7 @@ class OUTLINER_MT_asset(Menu):
     def draw(self, _context):
         layout = self.layout
 
-        layout.operator("asset.mark", icon="ASSET_MANAGER")
+        layout.operator("asset.mark", icon='ASSET_MANAGER')
         layout.operator("asset.clear", text="Clear Asset").set_fake_user = False
         layout.operator("asset.clear", text="Clear Asset (Set Fake User)").set_fake_user = True
 

scripts/startup/bl_ui/space_toolsystem_toolbar.py

@@ -1750,7 +1750,7 @@ class _defs_paint_grease_pencil:
             if not brush:
                 return
             layout.prop(brush.gpencil_settings, "eraser_mode", expand=True)
-            if brush.gpencil_settings.eraser_mode == "HARD":
+            if brush.gpencil_settings.eraser_mode == 'HARD':
                 layout.prop(brush.gpencil_settings, "use_keep_caps_eraser")
             layout.prop(brush.gpencil_settings, "use_active_layer_only")
         return dict(

scripts/startup/bl_ui/space_view3d.py

@@ -902,7 +902,7 @@ class VIEW3D_HT_header(Header):
                 icon = 'GROUP_VCOL' if canvas_source == 'COLOR_ATTRIBUTE' else canvas_source
                 row.popover(panel="VIEW3D_PT_slots_paint_canvas", icon=icon)
             else:
-                row.popover(panel="VIEW3D_PT_slots_color_attributes", icon="GROUP_VCOL")
+                row.popover(panel="VIEW3D_PT_slots_color_attributes", icon='GROUP_VCOL')
 
             layout.popover(
                 panel="VIEW3D_PT_sculpt_automasking",
@@ -913,12 +913,12 @@ class VIEW3D_HT_header(Header):
         elif object_mode == 'VERTEX_PAINT':
             row = layout.row()
             row.ui_units_x = 6
-            row.popover(panel="VIEW3D_PT_slots_color_attributes", icon="GROUP_VCOL")
+            row.popover(panel="VIEW3D_PT_slots_color_attributes", icon='GROUP_VCOL')
 
         elif object_mode == 'WEIGHT_PAINT':
             row = layout.row()
             row.ui_units_x = 6
-            row.popover(panel="VIEW3D_PT_slots_vertex_groups", icon="GROUP_VERTEX")
+            row.popover(panel="VIEW3D_PT_slots_vertex_groups", icon='GROUP_VERTEX')
 
         elif object_mode == 'TEXTURE_PAINT':
             tool_mode = tool_settings.image_paint.mode
@@ -927,7 +927,7 @@ class VIEW3D_HT_header(Header):
             row = layout.row()
             row.ui_units_x = 9
             row.popover(panel="VIEW3D_PT_slots_projectpaint", icon=icon)
-            row.popover(panel="VIEW3D_PT_mask", icon="MOD_MASK", text="")
+            row.popover(panel="VIEW3D_PT_mask", icon='MOD_MASK', text="")
         else:
             # Transform settings depending on tool header visibility
             VIEW3D_HT_header.draw_xform_template(layout, context)
@@ -2601,7 +2601,7 @@ class VIEW3D_MT_object(Menu):
 
         layout.separator()
 
-        layout.menu("VIEW3D_MT_object_asset", icon="ASSET_MANAGER")
+        layout.menu("VIEW3D_MT_object_asset", icon='ASSET_MANAGER')
         layout.menu("VIEW3D_MT_object_parent")
         layout.menu("VIEW3D_MT_object_collection")
         layout.menu("VIEW3D_MT_object_relations")
@@ -4863,9 +4863,15 @@ class VIEW3D_MT_edit_greasepencil_delete(Menu):
         layout = self.layout
 
         layout.operator_enum("grease_pencil.dissolve", "type")
-        
+
-        layout.operator("grease_pencil.delete_frame", text="Delete Active Keyframe (Active Layer)").type = 'ACTIVE_FRAME'
+        layout.operator(
-        layout.operator("grease_pencil.delete_frame", text="Delete Active Keyframes (All Layers)").type = 'ALL_FRAMES'
+            "grease_pencil.delete_frame",
+            text="Delete Active Keyframe (Active Layer)",
+        ).type = 'ACTIVE_FRAME'
+        layout.operator(
+            "grease_pencil.delete_frame",
+            text="Delete Active Keyframes (All Layers)",
+        ).type = 'ALL_FRAMES'
 
 
 # Edit Curve
@@ -6563,7 +6569,7 @@ class VIEW3D_PT_shading_compositor(Panel):
         row = self.layout.row()
         row.active = is_supported
         row.prop(shading, "use_compositor", expand=True)
-        if shading.use_compositor != "DISABLED" and not is_supported:
+        if shading.use_compositor != 'DISABLED' and not is_supported:
             self.layout.label(text="Compositor not supported on this platform", icon='ERROR')
 
 

source/blender/animrig/ANIM_bone_collections.hh

@@ -0,0 +1,62 @@
+/* SPDX-FileCopyrightText: 2023 Blender Foundation
+ *
+ * SPDX-License-Identifier: GPL-2.0-or-later */
+
+/** \file
+ * \ingroup animrig
+ *
+ * \brief C++ functions to deal with Armature collections (i.e. the successor of bone layers).
+ */
+
+#pragma once
+
+#ifndef __cplusplus
+#  error This is a C++ header.
+#endif
+
+#include "BLI_map.hh"
+
+#include "ANIM_bone_collections.h"
+
+namespace blender::animrig {
+
+/* --------------------------------------------------------------------
+ * The following functions are only used by edit-mode Armature undo:
+ */
+
+/**
+ * Duplicates a list of BoneCollections for edit-mode undo purposes, and
+ * returns original-to-duplicate remapping data.
+ *
+ * IMPORTANT: this discards membership data in the duplicate collections.
+ * This is because this function is only intended to be used with
+ * edit-mode Armatures, where the membership information in collections
+ * is not definitive, instead being stored in the EditBones. The
+ * assumption is that the membership information in the collections will
+ * be rebuilt from the EditBones when leaving edit mode.
+ *
+ * \param do_id_user: when true, increments the user count of IDs that
+ * the BoneCollections' custom properties point to, if any.
+ *
+ * \return a map from pointers-to-the-original-collections to
+ * pointers-to-the-duplicate-collections. This can be used to remap
+ * collection pointers in other data, such as EditBones.
+ */
+blender::Map<BoneCollection *, BoneCollection *> ANIM_bonecoll_listbase_copy_no_membership(
+    ListBase *bone_colls_dst, ListBase *bone_colls_src, bool do_id_user);
+/**
+ * Frees a list of BoneCollections.
+ *
+ * IMPORTANT: although there is nothing about this function that
+ * fundamentally prevents it from being used generally, other data
+ * structures like Armature runtime data and EditBones often store
+ * direct pointers to BoneCollections, which this function does NOT
+ * handle. Prefer using higher-level functions to remove BoneCollections
+ * from Armatures.
+ *
+ * \param do_id_user: when true, decrements the user count of IDs that
+ * the BoneCollections' custom properties point to, if any.
+ */
+void ANIM_bonecoll_listbase_free(ListBase *bcolls, bool do_id_user);
+
+}  // namespace blender::animrig

source/blender/animrig/CMakeLists.txt

@@ -18,6 +18,7 @@ set(SRC
   intern/bonecolor.cc
 
   ANIM_bone_collections.h
+  ANIM_bone_collections.hh
   ANIM_bonecolor.hh
 )
 

source/blender/animrig/intern/bone_collections.cc

@@ -7,6 +7,7 @@
  */
 
 #include "BLI_linklist.h"
+#include "BLI_map.hh"
 #include "BLI_math_color.h"
 #include "BLI_string.h"
 #include "BLI_string_utf8.h"
@@ -21,9 +22,10 @@
 
 #include "BKE_animsys.h"
 #include "BKE_idprop.h"
+#include "BKE_lib_id.h"
 
 #include "ANIM_armature_iter.hh"
-#include "ANIM_bone_collections.h"
+#include "ANIM_bone_collections.hh"
 
 #include <cstring>
 #include <string>
@@ -437,3 +439,51 @@ void ANIM_armature_bonecoll_show_from_pchan(bArmature *armature, const bPoseChan
 {
   ANIM_armature_bonecoll_show_from_bone(armature, pchan->bone);
 }
+
+/* ********* */
+/* C++ only. */
+namespace blender::animrig {
+
+/* Utility functions for Armature edit-mode undo. */
+
+blender::Map<BoneCollection *, BoneCollection *> ANIM_bonecoll_listbase_copy_no_membership(
+    ListBase *bone_colls_dst, ListBase *bone_colls_src, const bool do_id_user)
+{
+  BLI_assert(BLI_listbase_is_empty(bone_colls_dst));
+
+  blender::Map<BoneCollection *, BoneCollection *> bcoll_map{};
+  LISTBASE_FOREACH (BoneCollection *, bcoll_src, bone_colls_src) {
+    BoneCollection *bcoll_dst = static_cast<BoneCollection *>(MEM_dupallocN(bcoll_src));
+
+    /* This will be rebuilt from the edit bones, so we don't need to copy it. */
+    BLI_listbase_clear(&bcoll_dst->bones);
+
+    if (bcoll_src->prop) {
+      bcoll_dst->prop = IDP_CopyProperty_ex(bcoll_src->prop,
+                                            do_id_user ? 0 : LIB_ID_CREATE_NO_USER_REFCOUNT);
+    }
+    BLI_addtail(bone_colls_dst, bcoll_dst);
+    bcoll_map.add(bcoll_src, bcoll_dst);
+  }
+
+  return bcoll_map;
+}
+
+void ANIM_bonecoll_listbase_free(ListBase *bcolls, const bool do_id_user)
+{
+  LISTBASE_FOREACH_MUTABLE (BoneCollection *, bcoll, bcolls) {
+    if (bcoll->prop) {
+      IDP_FreeProperty_ex(bcoll->prop, do_id_user);
+    }
+
+    /* This will usually already be empty, because the passed BoneCollection
+     * list is usually from ANIM_bonecoll_listbase_copy_no_membership().
+     * However, during undo this is also used to free the BoneCollection
+     * list on the Armature itself before copying over the undo BoneCollection
+     * list, in which case this of Bone pointers may not be empty. */
+    BLI_freelistN(&bcoll->bones);
+  }
+  BLI_freelistN(bcolls);
+}
+
+}  // namespace blender::animrig

source/blender/blenkernel/BKE_armature.h

@@ -195,8 +195,11 @@ bool BKE_pose_minmax(
 bool bone_autoside_name(char name[64], int strip_number, short axis, float head, float tail);
 
 /**
- * Walk the list until the bone is found (slow!),
+ * Find the bone with the given name.
- * use #BKE_armature_bone_from_name_map for multiple lookups.
+ *
+ * When doing multiple subsequent calls to this function, consider calling
+ * #BKE_armature_bone_hash_make first to hash the bone names and speed up
+ * queries.
  */
 struct Bone *BKE_armature_find_bone_name(struct bArmature *arm, const char *name);
 

source/blender/blenkernel/BKE_blender_version.h

@@ -37,7 +37,7 @@ extern "C" {
  *
  * See https://wiki.blender.org/wiki/Process/Compatibility_Handling for details. */
 #define BLENDER_FILE_MIN_VERSION 306
-#define BLENDER_FILE_MIN_SUBVERSION 12
+#define BLENDER_FILE_MIN_SUBVERSION 13
 
 /** User readable version string. */
 const char *BKE_blender_version_string(void);

source/blender/blenkernel/BKE_curves.hh

@@ -393,9 +393,22 @@ class CurvesGeometry : public ::CurvesGeometry {
   /* --------------------------------------------------------------------
    * File Read/Write.
    */
-
   void blend_read(BlendDataReader &reader);
-  void blend_write(BlendWriter &writer, ID &id);
+  /**
+   * Helper struct for `CurvesGeometry::blend_write_*` functions.
+   */
+  struct BlendWriteData {
+    /* The point custom data layers to be written. */
+    Vector<CustomDataLayer, 16> point_layers;
+    /* The curve custom data layers to be written. */
+    Vector<CustomDataLayer, 16> curve_layers;
+  };
+  /**
+   * This function needs to be called before `blend_write` and before the `CurvesGeometry` struct
+   * is written because it can mutate the `CustomData` struct.
+   */
+  BlendWriteData blend_write_prepare();
+  void blend_write(BlendWriter &writer, ID &id, const BlendWriteData &write_data);
 };
 
 static_assert(sizeof(blender::bke::CurvesGeometry) == sizeof(::CurvesGeometry));

source/blender/blenkernel/BKE_customdata.h

@@ -429,6 +429,12 @@ void CustomData_bmesh_interp(struct CustomData *data,
  */
 void CustomData_swap_corners(struct CustomData *data, int index, const int *corner_indices);
 
+/**
+ * Custom data layers can be shared through implicit sharing (`BLI_implicit_sharing.h`). This
+ * function makes sure that the layer is unshared if it was shared, which makes it mutable.
+ */
+void CustomData_ensure_data_is_mutable(struct CustomDataLayer *layer, int totelem);
+
 /**
  * Retrieve a pointer to an element of the active layer of the given \a type, chosen by the
  * \a index, if it exists.

source/blender/blenkernel/BKE_grease_pencil.hh

@@ -260,6 +260,10 @@ struct LayerTransformData {
    * i.e. each frame that is not an implicit hold. */
   Map<int, int> frames_duration;
 
+  /* Temporary copy of duplicated frames before we decide on a place to insert them.
+   * Used in the move+duplicate operator. */
+  Map<int, GreasePencilFrame> temp_frames_buffer;
+
   FrameTransformationStatus status{TRANS_CLEAR};
 };
 

source/blender/blenkernel/BKE_nla.h

@@ -272,6 +272,12 @@ void BKE_nlatrack_solo_toggle(struct AnimData *adt, struct NlaTrack *nlt);
  * Check if there is any space in the given track to add a strip of the given length.
  */
 bool BKE_nlatrack_has_space(struct NlaTrack *nlt, float start, float end);
+
+/**
+ * Check to see if there are any NLA strips in the NLA tracks.
+ */
+bool BKE_nlatrack_has_strips(ListBase *tracks);
+
 /**
  * Rearrange the strips in the track so that they are always in order
  * (usually only needed after a strip has been moved).

source/blender/blenkernel/BKE_node.h

@@ -378,13 +378,6 @@ typedef struct bNodeType {
    */
   NodeGatherSocketLinkOperationsFunction gather_link_search_ops;
 
-  /**
-   * Add to the list of search items gathered by the add-node search. The default behavior of
-   * adding a single item with the node name is usually enough, but node types can have any number
-   * of custom search items.
-   */
-  NodeGatherAddOperationsFunction gather_add_node_search_ops;
-
   /** True when the node cannot be muted. */
   bool no_muting;
 

source/blender/blenkernel/intern/anim_sys.cc

@@ -963,8 +963,23 @@ NlaEvalStrip *nlastrips_ctime_get_strip(ListBase *list,
   /* loop over strips, checking if they fall within the range */
   LISTBASE_FOREACH (NlaStrip *, strip, strips) {
     /* Check if current time occurs within this strip. */
-    if (IN_RANGE_INCL(ctime, strip->start, strip->end) ||
+
-        (strip->flag & NLASTRIP_FLAG_NO_TIME_MAP)) {
+    /* This block leads to the Action Track and non-time-remapped tweak strip evaluation to respect
+     * the extrapolation modes. If in_range, these two tracks will always output NES_TIME_WITHIN so
+     * fcurve extrapolation isn't clamped to the keyframe bounds. */
+    bool in_range = IN_RANGE_INCL(ctime, strip->start, strip->end);
+    if (strip->flag & NLASTRIP_FLAG_NO_TIME_MAP) {
+      switch (strip->extendmode) {
+        case NLASTRIP_EXTEND_HOLD:
+          in_range = true;
+          break;
+        case NLASTRIP_EXTEND_HOLD_FORWARD:
+          in_range = ctime >= strip->start;
+          break;
+      }
+    }
+
+    if (in_range) {
       /* this strip is active, so try to use it */
       estrip = strip;
       side = NES_TIME_WITHIN;
@@ -3217,16 +3232,10 @@ static void animsys_create_action_track_strip(const AnimData *adt,
   r_action_strip->extendmode = adt->act_extendmode;
   r_action_strip->influence = adt->act_influence;
 
-  /* NOTE: must set this, or else the default setting overrides,
+  /* Must set NLASTRIP_FLAG_USR_INFLUENCE, or else the default setting overrides, and influence
-   * and this setting doesn't work. */
+   * doesn't work.
-  r_action_strip->flag |= NLASTRIP_FLAG_USR_INFLUENCE;
-
-  /* Unless `extendmode` is Nothing (might be useful for flattening NLA evaluation), disable range.
-   * Extend-mode Nothing and Hold will behave as normal. Hold Forward will behave just like Hold.
    */
-  if (r_action_strip->extendmode != NLASTRIP_EXTEND_NOTHING) {
+  r_action_strip->flag |= NLASTRIP_FLAG_USR_INFLUENCE;
-    r_action_strip->flag |= NLASTRIP_FLAG_NO_TIME_MAP;
-  }
 
   const bool tweaking = (adt->flag & ADT_NLA_EDIT_ON) != 0;
   const bool soloing = (adt->flag & ADT_NLA_SOLO_TRACK) != 0;

source/blender/blenkernel/intern/curves.cc

@@ -105,12 +105,15 @@ static void curves_blend_write(BlendWriter *writer, ID *id, const void *id_addre
 {
   Curves *curves = (Curves *)id;
 
+  blender::bke::CurvesGeometry::BlendWriteData write_data =
+      curves->geometry.wrap().blend_write_prepare();
+
   /* Write LibData */
   BLO_write_id_struct(writer, Curves, id_address, &curves->id);
   BKE_id_blend_write(writer, &curves->id);
 
   /* Direct data */
-  curves->geometry.wrap().blend_write(*writer, curves->id);
+  curves->geometry.wrap().blend_write(*writer, curves->id, write_data);
 
   BLO_write_string(writer, curves->surface_uv_map);
 

source/blender/blenkernel/intern/curves_geometry.cc

@@ -1477,17 +1477,22 @@ void CurvesGeometry::blend_read(BlendDataReader &reader)
   this->update_curve_types();
 }
 
-void CurvesGeometry::blend_write(BlendWriter &writer, ID &id)
+CurvesGeometry::BlendWriteData CurvesGeometry::blend_write_prepare()
 {
-  Vector<CustomDataLayer, 16> point_layers;
+  CurvesGeometry::BlendWriteData write_data;
-  Vector<CustomDataLayer, 16> curve_layers;
+  CustomData_blend_write_prepare(this->point_data, write_data.point_layers);
-  CustomData_blend_write_prepare(this->point_data, point_layers);
+  CustomData_blend_write_prepare(this->curve_data, write_data.curve_layers);
-  CustomData_blend_write_prepare(this->curve_data, curve_layers);
+  return write_data;
+}
 
+void CurvesGeometry::blend_write(BlendWriter &writer,
+                                 ID &id,
+                                 const CurvesGeometry::BlendWriteData &write_data)
+{
   CustomData_blend_write(
-      &writer, &this->point_data, point_layers, this->point_num, CD_MASK_ALL, &id);
+      &writer, &this->point_data, write_data.point_layers, this->point_num, CD_MASK_ALL, &id);
   CustomData_blend_write(
-      &writer, &this->curve_data, curve_layers, this->curve_num, CD_MASK_ALL, &id);
+      &writer, &this->curve_data, write_data.curve_layers, this->curve_num, CD_MASK_ALL, &id);
 
   BLO_write_int32_array(&writer, this->curve_num + 1, this->curve_offsets);
 }

source/blender/blenkernel/intern/customdata.cc

@@ -2414,6 +2414,11 @@ static void ensure_layer_data_is_mutable(CustomDataLayer &layer, const int totel
   }
 }
 
+void CustomData_ensure_data_is_mutable(CustomDataLayer *layer, const int totelem)
+{
+  ensure_layer_data_is_mutable(*layer, totelem);
+}
+
 void CustomData_realloc(CustomData *data, const int old_size, const int new_size)
 {
   BLI_assert(new_size >= 0);

source/blender/blenkernel/intern/grease_pencil.cc

@@ -1564,26 +1564,43 @@ void GreasePencil::remove_drawings_with_no_users()
 void GreasePencil::move_frames(blender::bke::greasepencil::Layer &layer,
                                const blender::Map<int, int> &frame_number_destinations)
 {
-  using namespace blender;
+  return this->move_duplicate_frames(
+      layer, frame_number_destinations, blender::Map<int, GreasePencilFrame>());
+}
 
+void GreasePencil::move_duplicate_frames(
+    blender::bke::greasepencil::Layer &layer,
+    const blender::Map<int, int> &frame_number_destinations,
+    const blender::Map<int, GreasePencilFrame> &duplicate_frames)
+{
+  using namespace blender;
   Map<int, GreasePencilFrame> layer_frames_copy = layer.frames();
 
-  /* Remove all frames that have a mapping. */
+  /* Copy frames durations. */
-  for (const int frame_number : frame_number_destinations.keys()) {
+  Map<int, int> layer_frames_durations;
-    layer.remove_frame(frame_number);
+  for (const auto [frame_number, frame] : layer.frames().items()) {
+    if (!frame.is_implicit_hold()) {
+      layer_frames_durations.add(frame_number, layer.get_frame_duration_at(frame_number));
+    }
   }
 
-  /* Insert all frames of the transformation. */
   for (const auto [src_frame_number, dst_frame_number] : frame_number_destinations.items()) {
-    if (!layer_frames_copy.contains(src_frame_number)) {
+    const bool use_duplicate = duplicate_frames.contains(src_frame_number);
+
+    const Map<int, GreasePencilFrame> &frame_map = use_duplicate ? duplicate_frames :
+                                                                   layer_frames_copy;
+
+    if (!frame_map.contains(src_frame_number)) {
       continue;
     }
 
-    const GreasePencilFrame src_frame = layer_frames_copy.lookup(src_frame_number);
+    const GreasePencilFrame src_frame = frame_map.lookup(src_frame_number);
     const int drawing_index = src_frame.drawing_index;
-    const int duration = src_frame.is_implicit_hold() ?
+    const int duration = layer_frames_durations.lookup_default(src_frame_number, 0);
-                             0 :
+
-                             layer.get_frame_duration_at(src_frame_number);
+    if (!use_duplicate) {
+      layer.remove_frame(src_frame_number);
+    }
 
     /* Add and overwrite the frame at the destination number. */
     if (layer.frames().contains(dst_frame_number)) {
@@ -2031,14 +2048,17 @@ static void read_drawing_array(GreasePencil &grease_pencil, BlendDataReader *rea
 
 static void write_drawing_array(GreasePencil &grease_pencil, BlendWriter *writer)
 {
+  using namespace blender;
   BLO_write_pointer_array(writer, grease_pencil.drawing_array_num, grease_pencil.drawing_array);
   for (int i = 0; i < grease_pencil.drawing_array_num; i++) {
     GreasePencilDrawingBase *drawing_base = grease_pencil.drawing_array[i];
     switch (drawing_base->type) {
       case GP_DRAWING: {
         GreasePencilDrawing *drawing = reinterpret_cast<GreasePencilDrawing *>(drawing_base);
+        bke::CurvesGeometry::BlendWriteData write_data =
+            drawing->wrap().strokes_for_write().blend_write_prepare();
         BLO_write_struct(writer, GreasePencilDrawing, drawing);
-        drawing->wrap().strokes_for_write().blend_write(*writer, grease_pencil.id);
+        drawing->wrap().strokes_for_write().blend_write(*writer, grease_pencil.id, write_data);
         break;
       }
       case GP_DRAWING_REFERENCE: {

source/blender/blenkernel/intern/nla.cc

@@ -1221,6 +1221,24 @@ bool BKE_nlatrack_has_space(NlaTrack *nlt, float start, float end)
   return BKE_nlastrips_has_space(&nlt->strips, start, end);
 }
 
+bool BKE_nlatrack_has_strips(ListBase *tracks)
+{
+  /* sanity checks */
+  if (BLI_listbase_is_empty(tracks)) {
+    return false;
+  }
+
+  /* Check each track for NLA strips. */
+  LISTBASE_FOREACH (NlaTrack *, track, tracks) {
+    if (BLI_listbase_count(&track->strips) > 0) {
+      return true;
+    }
+  }
+
+  /* none found */
+  return false;
+}
+
 void BKE_nlatrack_sort_strips(NlaTrack *nlt)
 {
   /* sanity checks */

source/blender/blenkernel/intern/screen.cc

@@ -322,6 +322,7 @@ static void panel_list_copy(ListBase *newlb, const ListBase *lb)
   Panel *panel = static_cast<Panel *>(lb->first);
   for (; new_panel; new_panel = new_panel->next, panel = panel->next) {
     new_panel->activedata = nullptr;
+    new_panel->drawname = nullptr;
     memset(&new_panel->runtime, 0x0, sizeof(new_panel->runtime));
     panel_list_copy(&new_panel->children, &panel->children);
   }
@@ -489,6 +490,7 @@ void BKE_region_callback_free_gizmomap_set(void (*callback)(wmGizmoMap *))
 static void area_region_panels_free_recursive(Panel *panel)
 {
   MEM_SAFE_FREE(panel->activedata);
+  MEM_SAFE_FREE(panel->drawname);
 
   LISTBASE_FOREACH_MUTABLE (Panel *, child_panel, &panel->children) {
     area_region_panels_free_recursive(child_panel);
@@ -1097,6 +1099,7 @@ static void direct_link_panel_list(BlendDataReader *reader, ListBase *lb)
     panel->runtime_flag = 0;
     panel->activedata = nullptr;
     panel->type = nullptr;
+    panel->drawname = nullptr;
     panel->runtime.custom_data_ptr = nullptr;
     direct_link_panel_list(reader, &panel->children);
   }

source/blender/blenlib/intern/system_win32.c

@@ -227,8 +227,9 @@ static void bli_windows_system_backtrace_modules(FILE *fp)
 {
   fprintf(fp, "Loaded Modules :\n");
   HANDLE hModuleSnap = CreateToolhelp32Snapshot(TH32CS_SNAPMODULE, 0);
-  if (hModuleSnap == INVALID_HANDLE_VALUE)
+  if (hModuleSnap == INVALID_HANDLE_VALUE) {
     return;
+  }
 
   MODULEENTRY32 me32;
   me32.dwSize = sizeof(MODULEENTRY32);

source/blender/blenloader/intern/readfile.cc

@@ -573,8 +573,9 @@ void blo_readfile_invalidate(FileData *fd, Main *bmain, const char *message)
   /* Tag given `bmain`, and 'root 'local' main one (in case given one is a library one) as invalid.
    */
   bmain->is_read_invalid = true;
-  for (; bmain->prev != nullptr; bmain = bmain->prev)
+  for (; bmain->prev != nullptr; bmain = bmain->prev) {
-    ;
+    /* Pass. */
+  }
   bmain->is_read_invalid = true;
 
   BLO_reportf_wrap(fd->reports,

source/blender/blenloader/intern/versioning_300.cc

@@ -69,6 +69,7 @@
 #include "BKE_main_namemap.h"
 #include "BKE_mesh.hh"
 #include "BKE_modifier.h"
+#include "BKE_nla.h"
 #include "BKE_node.hh"
 #include "BKE_screen.h"
 #include "BKE_workspace.h"
@@ -1052,6 +1053,28 @@ static void version_geometry_nodes_extrude_smooth_propagation(bNodeTree &ntree)
   }
 }
 
+/* Change the action strip (if a NLA strip is preset) to HOLD instead of HOLD FORWARD to maintain
+ * backwards compatibility.*/
+static void version_nla_action_strip_hold(Main *bmain)
+{
+  ID *id;
+  FOREACH_MAIN_ID_BEGIN (bmain, id) {
+    AnimData *adt = BKE_animdata_from_id(id);
+    /* We only want to preserve existing behavior if there's an action and 1 or more NLA strips. */
+    if (adt == nullptr || adt->action == nullptr ||
+        adt->act_extendmode != NLASTRIP_EXTEND_HOLD_FORWARD)
+    {
+      continue;
+    }
+
+    if (BKE_nlatrack_has_strips(&adt->nla_tracks)) {
+      adt->act_extendmode = NLASTRIP_EXTEND_HOLD;
+    }
+
+    FOREACH_MAIN_ID_END;
+  }
+}
+
 void do_versions_after_linking_300(FileData * /*fd*/, Main *bmain)
 {
   if (MAIN_VERSION_FILE_ATLEAST(bmain, 300, 0) && !MAIN_VERSION_FILE_ATLEAST(bmain, 300, 1)) {
@@ -1341,6 +1364,10 @@ void do_versions_after_linking_300(FileData * /*fd*/, Main *bmain)
     }
   }
 
+  if (!MAIN_VERSION_FILE_ATLEAST(bmain, 306, 13)) {
+    version_nla_action_strip_hold(bmain);
+  }
+
   /**
    * Versioning code until next subversion bump goes here.
    *

source/blender/blenloader/intern/versioning_400.cc

@@ -580,6 +580,113 @@ static void version_principled_bsdf_sheen(bNodeTree *ntree)
   }
 }
 
+/* Convert subsurface inputs on the Principled BSDF. */
+static void version_principled_bsdf_subsurface(bNodeTree *ntree)
+{
+  /* - Create Subsurface Scale input
+   * - If a node's Subsurface input was connected or nonzero:
+   *   - Make the Base Color a mix of old Base Color and Subsurface Color,
+   *     using Subsurface as the mix factor
+   *   - Move Subsurface link and default value to the new Subsurface Scale input
+   *   - Set the Subsurface input to 1.0
+   * - Remove Subsurface Color input
+   */
+  LISTBASE_FOREACH (bNode *, node, &ntree->nodes) {
+    if (node->type != SH_NODE_BSDF_PRINCIPLED) {
+      continue;
+    }
+    if (nodeFindSocket(node, SOCK_IN, "Subsurface Scale")) {
+      /* Node is already updated. */
+      continue;
+    }
+
+    /* Add Scale input */
+    bNodeSocket *scale_in = nodeAddStaticSocket(
+        ntree, node, SOCK_IN, SOCK_FLOAT, PROP_DISTANCE, "Subsurface Scale", "Subsurface Scale");
+
+    bNodeSocket *subsurf = nodeFindSocket(node, SOCK_IN, "Subsurface");
+    float *subsurf_val = version_cycles_node_socket_float_value(subsurf);
+    *version_cycles_node_socket_float_value(scale_in) = *subsurf_val;
+
+    if (subsurf->link == nullptr && *subsurf_val == 0.0f) {
+      /* Node doesn't use Subsurf, we're done here. */
+      continue;
+    }
+
+    /* Fix up Subsurface Color input */
+    bNodeSocket *base_col = nodeFindSocket(node, SOCK_IN, "Base Color");
+    bNodeSocket *subsurf_col = nodeFindSocket(node, SOCK_IN, "Subsurface Color");
+    float *base_col_val = version_cycles_node_socket_rgba_value(base_col);
+    float *subsurf_col_val = version_cycles_node_socket_rgba_value(subsurf_col);
+    /* If any of the three inputs is dynamic, we need a Mix node. */
+    if (subsurf->link || subsurf_col->link || base_col->link) {
+      bNode *mix = nodeAddStaticNode(nullptr, ntree, SH_NODE_MIX);
+      static_cast<NodeShaderMix *>(mix->storage)->data_type = SOCK_RGBA;
+      mix->locx = node->locx - 170;
+      mix->locy = node->locy - 120;
+
+      bNodeSocket *a_in = nodeFindSocket(mix, SOCK_IN, "A_Color");
+      bNodeSocket *b_in = nodeFindSocket(mix, SOCK_IN, "B_Color");
+      bNodeSocket *fac_in = nodeFindSocket(mix, SOCK_IN, "Factor_Float");
+      bNodeSocket *result_out = nodeFindSocket(mix, SOCK_OUT, "Result_Color");
+
+      copy_v4_v4(version_cycles_node_socket_rgba_value(a_in), base_col_val);
+      copy_v4_v4(version_cycles_node_socket_rgba_value(b_in), subsurf_col_val);
+      *version_cycles_node_socket_float_value(fac_in) = *subsurf_val;
+
+      if (base_col->link) {
+        nodeAddLink(ntree, base_col->link->fromnode, base_col->link->fromsock, mix, a_in);
+        nodeRemLink(ntree, base_col->link);
+      }
+      if (subsurf_col->link) {
+        nodeAddLink(ntree, subsurf_col->link->fromnode, subsurf_col->link->fromsock, mix, b_in);
+        nodeRemLink(ntree, subsurf_col->link);
+      }
+      if (subsurf->link) {
+        nodeAddLink(ntree, subsurf->link->fromnode, subsurf->link->fromsock, mix, fac_in);
+        nodeAddLink(ntree, subsurf->link->fromnode, subsurf->link->fromsock, node, scale_in);
+        nodeRemLink(ntree, subsurf->link);
+      }
+      nodeAddLink(ntree, mix, result_out, node, base_col);
+    }
+    /* Mix the fixed values. */
+    interp_v4_v4v4(base_col_val, base_col_val, subsurf_col_val, *subsurf_val);
+
+    /* Set node to 100% subsurface, 0% diffuse. */
+    *subsurf_val = 1.0f;
+
+    /* Delete Subsurface Color input */
+    nodeRemoveSocket(ntree, node, subsurf_col);
+  }
+}
+
+/* Convert emission inputs on the Principled BSDF. */
+static void version_principled_bsdf_emission(bNodeTree *ntree)
+{
+  /* Blender 3.x and before would default to Emission = 0.0, Emission Strength = 1.0.
+   * Now we default the other way around (1.0 and 0.0), but because the Strength input was added
+   * a bit later, a file that only has the Emission socket would now end up as (1.0, 0.0) instead
+   * of (1.0, 1.0).
+   * Therefore, set strength to 1.0 for those files.
+   */
+  LISTBASE_FOREACH (bNode *, node, &ntree->nodes) {
+    if (node->type != SH_NODE_BSDF_PRINCIPLED) {
+      continue;
+    }
+    if (!nodeFindSocket(node, SOCK_IN, "Emission")) {
+      /* Old enough to have neither, new defaults are fine. */
+      continue;
+    }
+    if (nodeFindSocket(node, SOCK_IN, "Emission Strength")) {
+      /* New enough to have both, no need to do anything. */
+      continue;
+    }
+    bNodeSocket *sock = nodeAddStaticSocket(
+        ntree, node, SOCK_IN, SOCK_FLOAT, PROP_NONE, "Emission Strength", "Emission Strength");
+    *version_cycles_node_socket_float_value(sock) = 1.0f;
+  }
+}
+
 /* Replace old Principled Hair BSDF as a variant in the new Principled Hair BSDF. */
 static void version_replace_principled_hair_model(bNodeTree *ntree)
 {
@@ -655,6 +762,35 @@ static void versioning_convert_node_tree_socket_lists_to_interface(bNodeTree *nt
   }
 }
 
+/* Convert coat inputs on the Principled BSDF. */
+static void version_principled_bsdf_coat(bNodeTree *ntree)
+{
+  LISTBASE_FOREACH (bNode *, node, &ntree->nodes) {
+    if (node->type != SH_NODE_BSDF_PRINCIPLED) {
+      continue;
+    }
+    if (nodeFindSocket(node, SOCK_IN, "Coat IOR") != nullptr) {
+      continue;
+    }
+    bNodeSocket *coat_ior_input = nodeAddStaticSocket(
+        ntree, node, SOCK_IN, SOCK_FLOAT, PROP_NONE, "Coat IOR", "Coat IOR");
+
+    /* Adjust for 4x change in intensity. */
+    bNodeSocket *coat_input = nodeFindSocket(node, SOCK_IN, "Clearcoat");
+    *version_cycles_node_socket_float_value(coat_input) *= 0.25f;
+    /* When the coat input is dynamic, instead of inserting a *0.25 math node, set the Coat IOR
+     * to 1.2 instead - this also roughly quarters reflectivity compared to the 1.5 default. */
+    *version_cycles_node_socket_float_value(coat_ior_input) = (coat_input->link) ? 1.2f : 1.5f;
+  }
+
+  /* Rename sockets. */
+  version_node_input_socket_name(ntree, SH_NODE_BSDF_PRINCIPLED, "Clearcoat", "Coat");
+  version_node_input_socket_name(
+      ntree, SH_NODE_BSDF_PRINCIPLED, "Clearcoat Roughness", "Coat Roughness");
+  version_node_input_socket_name(
+      ntree, SH_NODE_BSDF_PRINCIPLED, "Clearcoat Normal", "Coat Normal");
+}
+
 void blo_do_versions_400(FileData *fd, Library * /*lib*/, Main *bmain)
 {
   if (!MAIN_VERSION_FILE_ATLEAST(bmain, 400, 1)) {
@@ -1062,5 +1198,17 @@ void blo_do_versions_400(FileData *fd, Library * /*lib*/, Main *bmain)
    */
   {
     /* Keep this block, even when empty. */
+
+    FOREACH_NODETREE_BEGIN (bmain, ntree, id) {
+      if (ntree->type == NTREE_SHADER) {
+        /* Convert coat inputs on the Principled BSDF. */
+        version_principled_bsdf_coat(ntree);
+        /* Convert subsurface inputs on the Principled BSDF. */
+        version_principled_bsdf_subsurface(ntree);
+        /* Convert emission on the Principled BSDF. */
+        version_principled_bsdf_emission(ntree);
+      }
+    }
+    FOREACH_NODETREE_END;
   }
 }

source/blender/blenloader/intern/versioning_defaults.cc

@@ -608,9 +608,12 @@ void BLO_update_defaults_startup_blend(Main *bmain, const char *app_template)
       for (bNode *node : ma->nodetree->all_nodes()) {
         if (node->type == SH_NODE_BSDF_PRINCIPLED) {
           bNodeSocket *roughness_socket = nodeFindSocket(node, SOCK_IN, "Roughness");
-          bNodeSocketValueFloat *roughness_data = static_cast<bNodeSocketValueFloat *>(
+          *version_cycles_node_socket_float_value(roughness_socket) = 0.5f;
-              roughness_socket->default_value);
+          bNodeSocket *emission = nodeFindSocket(node, SOCK_IN, "Emission");
-          roughness_data->value = 0.5f;
+          copy_v4_fl(version_cycles_node_socket_rgba_value(emission), 1.0f);
+          bNodeSocket *emission_strength = nodeFindSocket(node, SOCK_IN, "Emission Strength");
+          *version_cycles_node_socket_float_value(emission_strength) = 0.0f;
+
           node->custom1 = SHD_GLOSSY_MULTI_GGX;
           node->custom2 = SHD_SUBSURFACE_RANDOM_WALK;
           BKE_ntree_update_tag_node_property(ma->nodetree, node);

source/blender/draw/engines/eevee/shaders/closure_eval_surface_lib.glsl

@@ -230,15 +230,13 @@ Closure closure_eval(ClosureDiffuse diffuse, ClosureReflection reflection)
 
 /* Specular BSDF */
 CLOSURE_EVAL_FUNCTION_DECLARE_3(SpecularBSDF, Diffuse, Glossy, Glossy)
-Closure closure_eval(ClosureDiffuse diffuse,
+Closure closure_eval(ClosureDiffuse diffuse, ClosureReflection reflection, ClosureReflection coat)
-                     ClosureReflection reflection,
-                     ClosureReflection clearcoat)
 {
 #if defined(DO_SPLIT_CLOSURE_EVAL)
   Closure closure = closure_eval(diffuse);
   Closure closure_reflection = closure_eval(reflection);
-  Closure closure_clearcoat = closure_eval(clearcoat, false);
+  Closure closure_coat = closure_eval(coat, false);
-  closure.radiance += closure_reflection.radiance + closure_clearcoat.radiance;
+  closure.radiance += closure_reflection.radiance + closure_coat.radiance;
   return closure;
 #else
   /* Glue with the old system. */
@@ -250,8 +248,8 @@ Closure closure_eval(ClosureDiffuse diffuse,
   in_Diffuse_0.albedo = diffuse.color;
   in_Glossy_1.N = reflection.N;
   in_Glossy_1.roughness = reflection.roughness;
-  in_Glossy_2.N = clearcoat.N;
+  in_Glossy_2.N = coat.N;
-  in_Glossy_2.roughness = clearcoat.roughness;
+  in_Glossy_2.roughness = coat.roughness;
 
   CLOSURE_EVAL_FUNCTION_3(SpecularBSDF, Diffuse, Glossy, Glossy);
 
@@ -259,7 +257,7 @@ Closure closure_eval(ClosureDiffuse diffuse,
   if (!output_sss(diffuse, out_Diffuse_0)) {
     closure.radiance += out_Diffuse_0.radiance * diffuse.color * diffuse.weight;
   }
-  closure.radiance += out_Glossy_2.radiance * clearcoat.color * clearcoat.weight;
+  closure.radiance += out_Glossy_2.radiance * coat.color * coat.weight;
   if (!output_ssr(reflection)) {
     closure.radiance += out_Glossy_1.radiance * reflection.color * reflection.weight;
   }
@@ -271,15 +269,15 @@ Closure closure_eval(ClosureDiffuse diffuse,
 CLOSURE_EVAL_FUNCTION_DECLARE_4(PrincipledBSDF, Diffuse, Glossy, Glossy, Refraction)
 Closure closure_eval(ClosureDiffuse diffuse,
                      ClosureReflection reflection,
-                     ClosureReflection clearcoat,
+                     ClosureReflection coat,
                      ClosureRefraction refraction)
 {
 #if defined(DO_SPLIT_CLOSURE_EVAL)
   Closure closure = closure_eval(diffuse);
   Closure closure_reflection = closure_eval(reflection);
-  Closure closure_clearcoat = closure_eval(clearcoat, false);
+  Closure closure_coat = closure_eval(coat, false);
   Closure closure_refraction = closure_eval(refraction);
-  closure.radiance += closure_reflection.radiance + closure_clearcoat.radiance +
+  closure.radiance += closure_reflection.radiance + closure_coat.radiance +
                       closure_refraction.radiance;
   return closure;
 #else
@@ -290,8 +288,8 @@ Closure closure_eval(ClosureDiffuse diffuse,
   in_Diffuse_0.albedo = diffuse.color;
   in_Glossy_1.N = reflection.N;
   in_Glossy_1.roughness = reflection.roughness;
-  in_Glossy_2.N = clearcoat.N;
+  in_Glossy_2.N = coat.N;
-  in_Glossy_2.roughness = clearcoat.roughness;
+  in_Glossy_2.roughness = coat.roughness;
   in_Refraction_3.N = refraction.N;
   in_Refraction_3.roughness = refraction.roughness;
   in_Refraction_3.ior = refraction.ior;
@@ -299,7 +297,7 @@ Closure closure_eval(ClosureDiffuse diffuse,
   CLOSURE_EVAL_FUNCTION_4(PrincipledBSDF, Diffuse, Glossy, Glossy, Refraction);
 
   Closure closure = CLOSURE_DEFAULT;
-  closure.radiance += out_Glossy_2.radiance * clearcoat.color * clearcoat.weight;
+  closure.radiance += out_Glossy_2.radiance * coat.color * coat.weight;
   closure.radiance += out_Refraction_3.radiance * refraction.color * refraction.weight;
   if (!output_sss(diffuse, out_Diffuse_0)) {
     closure.radiance += out_Diffuse_0.radiance * diffuse.color * diffuse.weight;
@@ -312,10 +310,10 @@ Closure closure_eval(ClosureDiffuse diffuse,
 }
 
 CLOSURE_EVAL_FUNCTION_DECLARE_2(PrincipledBSDFMetalClearCoat, Glossy, Glossy)
-Closure closure_eval(ClosureReflection reflection, ClosureReflection clearcoat)
+Closure closure_eval(ClosureReflection reflection, ClosureReflection coat)
 {
 #if defined(DO_SPLIT_CLOSURE_EVAL)
-  Closure closure = closure_eval(clearcoat);
+  Closure closure = closure_eval(coat);
   Closure closure_reflection = closure_eval(reflection);
   closure.radiance += closure_reflection.radiance;
   return closure;
@@ -325,13 +323,13 @@ Closure closure_eval(ClosureReflection reflection, ClosureReflection clearcoat)
 
   in_Glossy_0.N = reflection.N;
   in_Glossy_0.roughness = reflection.roughness;
-  in_Glossy_1.N = clearcoat.N;
+  in_Glossy_1.N = coat.N;
-  in_Glossy_1.roughness = clearcoat.roughness;
+  in_Glossy_1.roughness = coat.roughness;
 
   CLOSURE_EVAL_FUNCTION_2(PrincipledBSDFMetalClearCoat, Glossy, Glossy);
 
   Closure closure = CLOSURE_DEFAULT;
-  closure.radiance += out_Glossy_1.radiance * clearcoat.color * clearcoat.weight;
+  closure.radiance += out_Glossy_1.radiance * coat.color * coat.weight;
   if (!output_ssr(reflection)) {
     closure.radiance += out_Glossy_0.radiance * reflection.color * reflection.weight;
   }

source/blender/draw/engines/eevee/shaders/closure_eval_volume_lib.glsl

@@ -42,20 +42,18 @@ Closure closure_eval(ClosureDiffuse diffuse, ClosureReflection reflection)
 {
   return CLOSURE_DEFAULT;
 }
-Closure closure_eval(ClosureDiffuse diffuse,
+Closure closure_eval(ClosureDiffuse diffuse, ClosureReflection reflection, ClosureReflection coat)
-                     ClosureReflection reflection,
-                     ClosureReflection clearcoat)
 {
   return CLOSURE_DEFAULT;
 }
 Closure closure_eval(ClosureDiffuse diffuse,
                      ClosureReflection reflection,
-                     ClosureReflection clearcoat,
+                     ClosureReflection coat,
                      ClosureRefraction refraction)
 {
   return CLOSURE_DEFAULT;
 }
-Closure closure_eval(ClosureReflection reflection, ClosureReflection clearcoat)
+Closure closure_eval(ClosureReflection reflection, ClosureReflection coat)
 {
   return CLOSURE_DEFAULT;
 }

source/blender/draw/engines/eevee/shaders/closure_type_lib.glsl

@@ -67,20 +67,18 @@ Closure closure_eval(ClosureHair hair);
 Closure closure_eval(ClosureReflection reflection, ClosureRefraction refraction);
 /* Dielectric BSDF. */
 Closure closure_eval(ClosureDiffuse diffuse, ClosureReflection reflection);
-/* ClearCoat BSDF. */
+/* Coat BSDF. */
-Closure closure_eval(ClosureReflection reflection, ClosureReflection clearcoat);
+Closure closure_eval(ClosureReflection reflection, ClosureReflection coat);
 /* Volume BSDF. */
 Closure closure_eval(ClosureVolumeScatter volume_scatter,
                      ClosureVolumeAbsorption volume_absorption,
                      ClosureEmission emission);
 /* Specular BSDF. */
-Closure closure_eval(ClosureDiffuse diffuse,
+Closure closure_eval(ClosureDiffuse diffuse, ClosureReflection reflection, ClosureReflection coat);
-                     ClosureReflection reflection,
-                     ClosureReflection clearcoat);
 /* Principled BSDF. */
 Closure closure_eval(ClosureDiffuse diffuse,
                      ClosureReflection reflection,
-                     ClosureReflection clearcoat,
+                     ClosureReflection coat,
                      ClosureRefraction refraction);
 
 Closure closure_add(inout Closure cl1, inout Closure cl2);

source/blender/draw/engines/eevee/shaders/infos/eevee_legacy_lightprobe_info.hh

@@ -227,7 +227,6 @@ GPU_SHADER_INTERFACE_INFO(eevee_legacy_probe_planar_downsample_geom_frag_iface,
 GPU_SHADER_CREATE_INFO(eevee_legacy_lightprobe_planar_downsample_common)
     .vertex_source("lightprobe_planar_downsample_vert.glsl")
     .fragment_source("lightprobe_planar_downsample_frag.glsl")
-    .builtins(BuiltinBits::LAYER)
     .vertex_out(eevee_legacy_probe_planar_downsample_vert_geom_iface)
     .vertex_out(eevee_legacy_probe_planar_downsample_vert_geom_flat_iface)
     .sampler(0, ImageType::FLOAT_2D_ARRAY, "source")
@@ -246,6 +245,7 @@ GPU_SHADER_CREATE_INFO(eevee_legacy_lightprobe_planar_downsample)
 #ifdef WITH_METAL_BACKEND
 GPU_SHADER_CREATE_INFO(eevee_legacy_lightprobe_planar_downsample_no_geom)
     .additional_info("eevee_legacy_lightprobe_planar_downsample_common")
+    .builtins(BuiltinBits::LAYER)
     .vertex_out(eevee_legacy_probe_planar_downsample_geom_frag_iface)
     .metal_backend_only(true)
     .do_static_compilation(true)

source/blender/draw/engines/eevee/shaders/infos/eevee_legacy_volume_info.hh

@@ -18,7 +18,6 @@ GPU_SHADER_CREATE_INFO(eevee_legacy_volumes_clear)
     .define("STANDALONE")
     .define("VOLUMETRICS")
     .define("CLEAR")
-    .builtins(BuiltinBits::LAYER)
     .additional_info("eevee_legacy_common_lib")
     .additional_info("draw_view")
     .additional_info("draw_resource_id_varying")
@@ -68,7 +67,6 @@ GPU_SHADER_CREATE_INFO(eevee_legacy_volumes_scatter_common)
     .define("STANDALONE")
     .define("VOLUMETRICS")
     .define("VOLUME_SHADOW")
-    .builtins(BuiltinBits::LAYER)
     .additional_info("eevee_legacy_common_lib")
     .additional_info("draw_view")
     .additional_info("draw_resource_id_varying")
@@ -100,6 +98,7 @@ GPU_SHADER_CREATE_INFO(eevee_legacy_volumes_scatter)
 #ifdef WITH_METAL_BACKEND
 GPU_SHADER_CREATE_INFO(eevee_legacy_volumes_scatter_no_geom)
     .additional_info("eevee_legacy_volumes_scatter_common")
+    .builtins(BuiltinBits::LAYER)
     .vertex_out(legacy_volume_geom_frag_iface)
     .metal_backend_only(true)
     .do_static_compilation(true)
@@ -133,7 +132,6 @@ GPU_SHADER_CREATE_INFO(eevee_legacy_volumes_integration_common)
     .additional_info("draw_view")
     .additional_info("eevee_legacy_volumetric_lib")
     .additional_info("draw_resource_id_varying")
-    .builtins(BuiltinBits::LAYER)
     /* NOTE: Unique sampler IDs assigned for consistency between library includes,
      * and to avoid unique assignment collision validation error.
      * However, resources will be auto assigned locations within shader usage. */
@@ -161,6 +159,7 @@ GPU_SHADER_CREATE_INFO(eevee_legacy_volumes_integration_common_geom)
 #ifdef WITH_METAL_BACKEND
 GPU_SHADER_CREATE_INFO(eevee_legacy_volumes_integration_common_no_geom)
     .additional_info("eevee_legacy_volumes_integration_common")
+    .builtins(BuiltinBits::LAYER)
     .vertex_out(legacy_volume_geom_frag_iface);
 #endif
 

source/blender/draw/engines/eevee_next/eevee_film.cc

@@ -176,6 +176,7 @@ static eViewLayerEEVEEPassType enabled_passes(const ViewLayer *view_layer)
   ENABLE_FROM_LEGACY(Z, Z)
   ENABLE_FROM_LEGACY(MIST, MIST)
   ENABLE_FROM_LEGACY(NORMAL, NORMAL)
+  ENABLE_FROM_LEGACY(POSITION, POSITION)
   ENABLE_FROM_LEGACY(SHADOW, SHADOW)
   ENABLE_FROM_LEGACY(AO, AO)
   ENABLE_FROM_LEGACY(EMIT, EMIT)
@@ -283,6 +284,7 @@ void Film::init(const int2 &extent, const rcti *output_rect)
     }
 
     const eViewLayerEEVEEPassType data_passes = EEVEE_RENDER_PASS_Z | EEVEE_RENDER_PASS_NORMAL |
+                                                EEVEE_RENDER_PASS_POSITION |
                                                 EEVEE_RENDER_PASS_VECTOR;
     const eViewLayerEEVEEPassType color_passes_1 = EEVEE_RENDER_PASS_DIFFUSE_LIGHT |
                                                    EEVEE_RENDER_PASS_SPECULAR_LIGHT |
@@ -328,6 +330,7 @@ void Film::init(const int2 &extent, const rcti *output_rect)
 
     data_.mist_id = pass_index_get(EEVEE_RENDER_PASS_MIST);
     data_.normal_id = pass_index_get(EEVEE_RENDER_PASS_NORMAL);
+    data_.position_id = pass_index_get(EEVEE_RENDER_PASS_POSITION);
     data_.vector_id = pass_index_get(EEVEE_RENDER_PASS_VECTOR);
     data_.diffuse_light_id = pass_index_get(EEVEE_RENDER_PASS_DIFFUSE_LIGHT);
     data_.diffuse_color_id = pass_index_get(EEVEE_RENDER_PASS_DIFFUSE_COLOR);

source/blender/draw/engines/eevee_next/eevee_film.hh

@@ -143,7 +143,7 @@ class Film {
   static bool pass_is_float3(eViewLayerEEVEEPassType pass_type)
   {
     return pass_storage_type(pass_type) == PASS_STORAGE_COLOR &&
-           pass_type != EEVEE_RENDER_PASS_COMBINED;
+           !ELEM(pass_type, EEVEE_RENDER_PASS_COMBINED, EEVEE_RENDER_PASS_VECTOR);
   }
 
   /* Returns layer offset in the accumulation texture. -1 if the pass is not enabled. */
@@ -158,6 +158,10 @@ class Film {
         return data_.mist_id;
       case EEVEE_RENDER_PASS_NORMAL:
         return data_.normal_id;
+      case EEVEE_RENDER_PASS_POSITION:
+        return data_.position_id;
+      case EEVEE_RENDER_PASS_VECTOR:
+        return data_.vector_id;
       case EEVEE_RENDER_PASS_DIFFUSE_LIGHT:
         return data_.diffuse_light_id;
       case EEVEE_RENDER_PASS_DIFFUSE_COLOR:
@@ -182,8 +186,6 @@ class Film {
         return data_.cryptomatte_asset_id;
       case EEVEE_RENDER_PASS_CRYPTOMATTE_MATERIAL:
         return data_.cryptomatte_material_id;
-      case EEVEE_RENDER_PASS_VECTOR:
-        return data_.vector_id;
       default:
         return -1;
     }
@@ -219,6 +221,12 @@ class Film {
       case EEVEE_RENDER_PASS_NORMAL:
         result.append(RE_PASSNAME_NORMAL);
         break;
+      case EEVEE_RENDER_PASS_POSITION:
+        result.append(RE_PASSNAME_POSITION);
+        break;
+      case EEVEE_RENDER_PASS_VECTOR:
+        result.append(RE_PASSNAME_VECTOR);
+        break;
       case EEVEE_RENDER_PASS_DIFFUSE_LIGHT:
         result.append(RE_PASSNAME_DIFFUSE_DIRECT);
         break;
@@ -255,9 +263,6 @@ class Film {
       case EEVEE_RENDER_PASS_CRYPTOMATTE_MATERIAL:
         build_cryptomatte_passes(RE_PASSNAME_CRYPTOMATTE_MATERIAL);
         break;
-      case EEVEE_RENDER_PASS_VECTOR:
-        result.append(RE_PASSNAME_VECTOR);
-        break;
       default:
         BLI_assert(0);
         break;

source/blender/draw/engines/eevee_next/eevee_instance.cc

@@ -350,7 +350,8 @@ void Instance::render_sample()
 void Instance::render_read_result(RenderLayer *render_layer, const char *view_name)
 {
   eViewLayerEEVEEPassType pass_bits = film.enabled_passes_get();
-  for (auto i : IndexRange(EEVEE_RENDER_PASS_MAX_BIT)) {
+
+  for (auto i : IndexRange(EEVEE_RENDER_PASS_MAX_BIT + 1)) {
     eViewLayerEEVEEPassType pass_type = eViewLayerEEVEEPassType(pass_bits & (1 << i));
     if (pass_type == 0) {
       continue;
@@ -498,6 +499,8 @@ void Instance::update_passes(RenderEngine *engine, Scene *scene, ViewLayer *view
   CHECK_PASS_LEGACY(Z, SOCK_FLOAT, 1, "Z");
   CHECK_PASS_LEGACY(MIST, SOCK_FLOAT, 1, "Z");
   CHECK_PASS_LEGACY(NORMAL, SOCK_VECTOR, 3, "XYZ");
+  CHECK_PASS_LEGACY(POSITION, SOCK_VECTOR, 3, "XYZ");
+  CHECK_PASS_LEGACY(VECTOR, SOCK_VECTOR, 4, "XYZW");
   CHECK_PASS_LEGACY(DIFFUSE_DIRECT, SOCK_RGBA, 3, "RGB");
   CHECK_PASS_LEGACY(DIFFUSE_COLOR, SOCK_RGBA, 3, "RGB");
   CHECK_PASS_LEGACY(GLOSSY_DIRECT, SOCK_RGBA, 3, "RGB");

source/blender/draw/engines/eevee_next/eevee_renderbuffers.cc

@@ -40,6 +40,7 @@ void RenderBuffers::sync()
   };
 
   data.normal_id = pass_index_get(EEVEE_RENDER_PASS_NORMAL, EEVEE_RENDER_PASS_AO);
+  data.position_id = pass_index_get(EEVEE_RENDER_PASS_POSITION);
   data.diffuse_light_id = pass_index_get(EEVEE_RENDER_PASS_DIFFUSE_LIGHT);
   data.diffuse_color_id = pass_index_get(EEVEE_RENDER_PASS_DIFFUSE_COLOR);
   data.specular_light_id = pass_index_get(EEVEE_RENDER_PASS_SPECULAR_LIGHT);

source/blender/draw/engines/eevee_next/eevee_shader_shared.hh

@@ -267,12 +267,13 @@ struct FilmData {
   bool1 any_render_pass_2;
   /** Controlled by user in lookdev mode or by render settings. */
   float background_opacity;
-  float _pad0, _pad1, _pad2;
+  float _pad0, _pad1;
   /** Output counts per type. */
   int color_len, value_len;
   /** Index in color_accum_img or value_accum_img of each pass. -1 if pass is not enabled. */
   int mist_id;
   int normal_id;
+  int position_id;
   int vector_id;
   int diffuse_light_id;
   int diffuse_color_id;
@@ -364,6 +365,7 @@ struct RenderBuffersInfoData {
   /* Color. */
   int color_len;
   int normal_id;
+  int position_id;
   int diffuse_light_id;
   int diffuse_color_id;
   int specular_light_id;
@@ -375,6 +377,7 @@ struct RenderBuffersInfoData {
   int value_len;
   int shadow_id;
   int ambient_occlusion_id;
+  int _pad0, _pad1, _pad2;
 };
 BLI_STATIC_ASSERT_ALIGN(RenderBuffersInfoData, 16)
 

source/blender/draw/engines/eevee_next/eevee_velocity.cc

@@ -344,8 +344,9 @@ void VelocityModule::end_sync()
                           (vel.geo.len[STEP_PREVIOUS] == GPU_vertbuf_get_vertex_len(pos_buf));
     }
     else {
-      vel.geo.do_deform = (vel.geo.len[STEP_PREVIOUS] == vel.geo.len[STEP_CURRENT]) &&
+      vel.geo.do_deform = (vel.geo.len[STEP_CURRENT] != 0) &&
-                          (vel.geo.len[STEP_NEXT] == vel.geo.len[STEP_CURRENT]);
+                          (vel.geo.len[STEP_CURRENT] == vel.geo.len[STEP_PREVIOUS]) &&
+                          (vel.geo.len[STEP_CURRENT] == vel.geo.len[STEP_NEXT]);
     }
     indirection_buf[vel.obj.resource_id] = vel;
     /* Reset for next sync. */

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

@@ -647,15 +647,22 @@ void film_process_data(ivec2 texel_film, out vec4 out_color, out float out_depth
     FilmSample film_sample = film_sample_get(0, texel_film);
 
     if (uniform_buf.film.use_reprojection || film_sample.weight < film_distance) {
-      vec4 normal = texelFetch(
-          rp_color_tx, ivec3(film_sample.texel, uniform_buf.render_pass.normal_id), 0);
       float depth = texelFetch(depth_tx, film_sample.texel, 0).x;
       vec4 vector = velocity_resolve(vector_tx, film_sample.texel, depth);
-      /* Transform to pixel space. */
+      /* Transform to pixel space, matching Cycles format. */
-      vector *= vec4(vec2(uniform_buf.film.render_extent), -vec2(uniform_buf.film.render_extent));
+      vector *= vec4(vec2(uniform_buf.film.render_extent), vec2(uniform_buf.film.render_extent));
 
       film_store_depth(texel_film, depth, out_depth);
-      film_store_data(texel_film, uniform_buf.film.normal_id, normal, out_color);
+      if (uniform_buf.film.normal_id != -1) {
+        vec4 normal = texelFetch(
+            rp_color_tx, ivec3(film_sample.texel, uniform_buf.render_pass.normal_id), 0);
+        film_store_data(texel_film, uniform_buf.film.normal_id, normal, out_color);
+      }
+      if (uniform_buf.film.position_id != -1) {
+        vec4 position = texelFetch(
+            rp_color_tx, ivec3(film_sample.texel, uniform_buf.render_pass.position_id), 0);
+        film_store_data(texel_film, uniform_buf.film.position_id, position, out_color);
+      }
       film_store_data(texel_film, uniform_buf.film.vector_id, vector, out_color);
       film_store_distance(texel_film, film_sample.weight);
     }

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

@@ -166,11 +166,11 @@ Closure closure_eval(ClosureDiffuse diffuse, ClosureReflection reflection)
   return Closure(0);
 }
 
-/* ClearCoat BSDF. */
+/* Coat BSDF. */
-Closure closure_eval(ClosureReflection reflection, ClosureReflection clearcoat)
+Closure closure_eval(ClosureReflection reflection, ClosureReflection coat)
 {
   SELECT_CLOSURE(g_reflection_data, g_reflection_rand, reflection);
-  SELECT_CLOSURE(g_reflection_data, g_reflection_rand, clearcoat);
+  SELECT_CLOSURE(g_reflection_data, g_reflection_rand, coat);
   return Closure(0);
 }
 
@@ -186,25 +186,23 @@ Closure closure_eval(ClosureVolumeScatter volume_scatter,
 }
 
 /* Specular BSDF. */
-Closure closure_eval(ClosureDiffuse diffuse,
+Closure closure_eval(ClosureDiffuse diffuse, ClosureReflection reflection, ClosureReflection coat)
-                     ClosureReflection reflection,
-                     ClosureReflection clearcoat)
 {
   SELECT_CLOSURE(g_diffuse_data, g_diffuse_rand, diffuse);
   SELECT_CLOSURE(g_reflection_data, g_reflection_rand, reflection);
-  SELECT_CLOSURE(g_reflection_data, g_reflection_rand, clearcoat);
+  SELECT_CLOSURE(g_reflection_data, g_reflection_rand, coat);
   return Closure(0);
 }
 
 /* Principled BSDF. */
 Closure closure_eval(ClosureDiffuse diffuse,
                      ClosureReflection reflection,
-                     ClosureReflection clearcoat,
+                     ClosureReflection coat,
                      ClosureRefraction refraction)
 {
   SELECT_CLOSURE(g_diffuse_data, g_diffuse_rand, diffuse);
   SELECT_CLOSURE(g_reflection_data, g_reflection_rand, reflection);
-  SELECT_CLOSURE(g_reflection_data, g_reflection_rand, clearcoat);
+  SELECT_CLOSURE(g_reflection_data, g_reflection_rand, coat);
   SELECT_CLOSURE(g_refraction_data, g_refraction_rand, refraction);
   return Closure(0);
 }

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

@@ -77,6 +77,7 @@ void main()
     imageStore(rp_cryptomatte_img, out_texel, cryptomatte_output);
   }
   output_renderpass_color(uniform_buf.render_pass.normal_id, vec4(out_normal, 1.0));
+  output_renderpass_color(uniform_buf.render_pass.position_id, vec4(g_data.P, 1.0));
   output_renderpass_color(uniform_buf.render_pass.diffuse_color_id,
                           vec4(g_diffuse_data.color, 1.0));
   output_renderpass_color(uniform_buf.render_pass.specular_color_id, vec4(specular_color, 1.0));

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

@@ -121,6 +121,7 @@ void main()
     imageStore(rp_cryptomatte_img, out_texel, cryptomatte_output);
   }
   output_renderpass_color(uniform_buf.render_pass.normal_id, vec4(out_normal, 1.0));
+  output_renderpass_color(uniform_buf.render_pass.position_id, vec4(g_data.P, 1.0));
   output_renderpass_color(uniform_buf.render_pass.diffuse_color_id,
                           vec4(g_diffuse_data.color, 1.0));
   output_renderpass_color(uniform_buf.render_pass.diffuse_light_id, vec4(diffuse_light, 1.0));

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

@@ -46,6 +46,7 @@ void main()
 
   vec4 clear_color = vec4(0.0, 0.0, 0.0, 1.0);
   output_renderpass_color(uniform_buf.render_pass.normal_id, clear_color);
+  output_renderpass_color(uniform_buf.render_pass.position_id, clear_color);
   output_renderpass_color(uniform_buf.render_pass.diffuse_light_id, clear_color);
   output_renderpass_color(uniform_buf.render_pass.specular_light_id, clear_color);
   output_renderpass_color(uniform_buf.render_pass.diffuse_color_id, clear_color);

source/blender/draw/engines/overlay/shaders/infos/overlay_edit_mode_info.hh

@@ -131,7 +131,6 @@ GPU_SHADER_CREATE_INFO(overlay_edit_mesh_face)
     .define("FACE")
     .vertex_in(0, Type::VEC3, "pos")
     .vertex_in(1, Type::UVEC4, "data")
-    .vertex_in(2, Type::VEC3, "vnor")
     .vertex_out(overlay_edit_flat_color_iface)
     .fragment_source("overlay_varying_color.glsl")
     .additional_info("overlay_edit_mesh_common");

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

@@ -11,8 +11,11 @@ void main()
   int i = 0;
   for (int x = -1; x <= 1; x++) {
     for (int y = -1; y <= 1; y++, i++) {
+      vec4 color = texture(colorBuffer, uv + vec2(x, y) * texel_size);
+      /* Clamp infinite inputs (See #112211). */
+      color = clamp(color, vec4(0.0), vec4(1e10));
       /* Use log2 space to avoid highlights creating too much aliasing. */
-      vec4 color = log2(texture(colorBuffer, uv + vec2(x, y) * texel_size) + 0.5);
+      color = log2(color + 0.5);
 
       fragColor += color * samplesWeights[i];
     }

source/blender/draw/engines/workbench/workbench_engine.cc

@@ -334,12 +334,14 @@ class Instance {
           mat.base_color = batch.debug_color();
         }
 
-        draw_mesh(ob_ref,
+        ::Image *image = nullptr;
-                  mat,
+        ImageUser *iuser = nullptr;
-                  batch.batch,
+        GPUSamplerState sampler_state = GPUSamplerState::default_sampler();
-                  handle,
+        if (object_state.color_type == V3D_SHADING_TEXTURE_COLOR) {
-                  object_state.image_paint_override,
+          get_material_image(ob_ref.object, batch.material_slot, image, iuser, sampler_state);
-                  object_state.override_sampler_state);
+        }
+
+        draw_mesh(ob_ref, mat, batch.batch, handle, image, sampler_state);
       }
     }
     else {
@@ -468,6 +470,10 @@ class Instance {
                                           GPU_DEPTH24_STENCIL8,
                                           GPU_TEXTURE_USAGE_SHADER_READ |
                                               GPU_TEXTURE_USAGE_ATTACHMENT);
+
+      fb.ensure(GPU_ATTACHMENT_TEXTURE(resources.depth_in_front_tx));
+      fb.bind();
+      GPU_framebuffer_clear_depth_stencil(fb, 1.0f, 0x00);
     }
 
     opaque_ps.draw(

source/blender/draw/engines/workbench/workbench_mesh_passes.cc

@@ -172,10 +172,10 @@ void OpaquePass::draw(Manager &manager,
     opaque_fb.bind();
 
     manager.submit(gbuffer_in_front_ps_, view);
-  }
 
-  if (resources.depth_in_front_tx.is_valid()) {
+    if (resources.depth_in_front_tx.is_valid()) {
-    GPU_texture_copy(resources.depth_in_front_tx, resources.depth_tx);
+      GPU_texture_copy(resources.depth_in_front_tx, resources.depth_tx);
+    }
   }
 
   if (!gbuffer_ps_.is_empty()) {

source/blender/draw/engines/workbench/workbench_state.cc

@@ -70,20 +70,18 @@ void SceneState::init(Object *camera_ob /*= nullptr*/)
     }
   }
 
-  if (!is_render_mode) {
+  if (shading.type < OB_SOLID) {
-    if (shading.type < OB_SOLID) {
+    shading.light = V3D_LIGHTING_FLAT;
-      shading.light = V3D_LIGHTING_FLAT;
+    shading.color_type = V3D_SHADING_OBJECT_COLOR;
-      shading.color_type = V3D_SHADING_OBJECT_COLOR;
+    shading.xray_alpha = 0.0f;
-      shading.xray_alpha = 0.0f;
-    }
-    else if (SHADING_XRAY_ENABLED(shading)) {
-      shading.xray_alpha = SHADING_XRAY_ALPHA(shading);
-    }
-    else {
-      shading.xray_alpha = 1.0f;
-    }
   }
-  xray_mode = !is_render_mode && shading.xray_alpha != 1.0f;
+  else if (SHADING_XRAY_ENABLED(shading)) {
+    shading.xray_alpha = SHADING_XRAY_ALPHA(shading);
+  }
+  else {
+    shading.xray_alpha = 1.0f;
+  }
+  xray_mode = shading.xray_alpha != 1.0f;
 
   if (SHADING_XRAY_FLAG_ENABLED(shading)) {
     /* Disable shading options that aren't supported in transparency mode. */