Cycles: Support automatic differentiation of shader nodes in SVM

This is an internal change preparing for the texture cache. Only implemented
for surfaces, and currently supports the following nodes:

- Geometry
- Tangent
- Mapping
- Attribute
- Texture Coordinate
- Environment Texture
- Image Texture
- Vector math
- UV Map
- Combine/Separate XYZ
- Bump

This has some impact on GPU rendering performance. Various changes were
made to optimize this, but rendering can still be a few % slower on some
GPUs. Some of the optimizations done:

* Use different node types enum for derivative nodes, consecutive to ensure the
  jump table works.
* Template various SVM derivative nodes to separate them from the
  non-derivative case, and avoid using dual types in those implementations.
* Use template on return type for stack_store and stack_load to make the above
  easier to implement.
* Use template on return type of primitive attribute reading to make derivative
  and non-derivative variations.
* Unify derivative and bump dx/dy nodes. Now it's a single derivative node
  that handles both cases.
* Derivative nodes are disabled in volume shaders for now.
* Tweak inlining on a few functions.

Co-authored-by: Brecht Van Lommel <brecht@blender.org>

Pull Request: https://projects.blender.org/blender/blender/pulls/155706

intern/cycles/kernel/bake/bake.h

@@ -66,7 +66,7 @@ ccl_device void kernel_background_evaluate(KernelGlobals kg,
 
   /* Setup shader data. */
   ShaderData sd;
-  shader_setup_from_background(kg, &sd, ray_P, ray_D, ray_time);
+  shader_setup_from_background(kg, &sd, ray_P, ray_D, 0.0f, ray_time);
 
   /* Evaluate shader.
    * This is being evaluated for all BSDFs, so path flag does not contain a specific type.

intern/cycles/kernel/camera/projection.h

@@ -28,6 +28,18 @@ ccl_device float2 direction_to_equirectangular_range(const float3 dir, const flo
   return make_float2(u, v);
 }
 
+ccl_device dual2 direction_to_equirectangular_range(const dual3 dir, const float4 range)
+{
+  if (is_zero(dir)) {
+    return make_zero<dual2>();
+  }
+
+  const dual1 u = (atan2(dir.y(), dir.x()) - range.y) / range.x;
+  const dual1 v = (acos(dir.z() / len(dir)) - range.w) / range.z;
+
+  return make_float2(u, v);
+}
+
 ccl_device float3 equirectangular_range_to_direction(const float u,
                                                      const float v,
                                                      const float4 range)
@@ -42,6 +54,11 @@ ccl_device float2 direction_to_equirectangular(const float3 dir)
   return direction_to_equirectangular_range(dir, make_float4(-M_2PI_F, M_PI_F, -M_PI_F, M_PI_F));
 }
 
+ccl_device dual2 direction_to_equirectangular(const dual3 dir)
+{
+  return direction_to_equirectangular_range(dir, make_float4(-M_2PI_F, M_PI_F, -M_PI_F, M_PI_F));
+}
+
 ccl_device float3 equirectangular_to_direction(const float u, const float v)
 {
   return equirectangular_range_to_direction(u, v, make_float4(-M_2PI_F, M_PI_F, -M_PI_F, M_PI_F));
@@ -234,6 +251,19 @@ ccl_device float2 direction_to_mirrorball(float3 dir)
   return make_float2(u, v);
 }
 
+ccl_device dual2 direction_to_mirrorball(dual3 dir)
+{
+  /* inverse of mirrorball_to_direction */
+  dir.val.y -= 1.0f;
+
+  dir = dir * 0.5f * inversesqrt(-0.5f * dir.y());
+
+  const dual1 u = 0.5f * (dir.x() + 1.0f);
+  const dual1 v = 0.5f * (dir.z() + 1.0f);
+
+  return make_float2(u, v);
+}
+
 /* Single face of a equiangular cube map projection as described in
  * https://blog.google/products/google-ar-vr/bringing-pixels-front-and-center-vr-video/ */
 ccl_device float3 equiangular_cubemap_face_to_direction(float u, float v)

intern/cycles/kernel/data_template.h

@@ -236,6 +236,9 @@ KERNEL_STRUCT_END(KernelIntegrator)
 
 KERNEL_STRUCT_BEGIN(KernelSVMUsage, svm_usage)
 #define SHADER_NODE_TYPE(type) KERNEL_STRUCT_MEMBER(svm_usage, int, type)
+#define SHADER_NODE_TYPE_DERIVATIVE(type) \
+  SHADER_NODE_TYPE(type) \
+  SHADER_NODE_TYPE(type##_DERIVATIVE)
 #include "kernel/svm/node_types_template.h"
 KERNEL_STRUCT_END(KernelSVMUsage)
 

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

@@ -8,6 +8,7 @@
 #include "kernel/device/cpu/globals.h"
 #include "kernel/util/image_2d.h"
 
+#include "util/defines.h"
 #include "util/half.h"
 #include "util/types_image.h"
 
@@ -325,9 +326,9 @@ template<typename TexT, typename OutT = float4> struct ImageInterpolator {
 #undef SET_CUBIC_SPLINE_WEIGHTS
 
 ccl_device float4 kernel_image_interp(KernelGlobals kg,
+                                      ccl_private ShaderData * /*sd*/,
                                       const int image_texture_id,
-                                      const float x,
-                                      float y)
+                                      dual2 uv)
 {
   if (image_texture_id == KERNEL_IMAGE_NONE) {
     return IMAGE_MISSING_RGBA;
@@ -337,6 +338,8 @@ ccl_device float4 kernel_image_interp(KernelGlobals kg,
     return IMAGE_MISSING_RGBA;
   }
   const KernelImageInfo &info = kernel_data_fetch(image_info, tex.image_info_id);
+  const float x = uv.val.x;
+  const float y = uv.val.y;
 
   if (UNLIKELY(!info.data)) {
     return zero_float4();
@@ -374,16 +377,16 @@ ccl_device float4 kernel_image_interp(KernelGlobals kg,
 }
 
 ccl_device_forceinline float4 kernel_image_interp_with_udim(KernelGlobals kg,
-                                                            ShaderData * /*sd*/,
+                                                            ShaderData *sd,
                                                             const int udim_id,
-                                                            float2 uv)
+                                                            dual2 uv)
 {
-  const int image_texture_id = kernel_image_udim_map(kg, udim_id, uv);
+  const int image_texture_id = kernel_image_udim_map(kg, udim_id, uv.val);
   if (image_texture_id == KERNEL_IMAGE_NONE) {
     return IMAGE_MISSING_RGBA;
   }
 
-  return kernel_image_interp(kg, image_texture_id, uv.x, uv.y);
+  return kernel_image_interp(kg, sd, image_texture_id, uv);
 }
 
 } /* Namespace. */

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

@@ -6,6 +6,7 @@
 
 #include "kernel/globals.h"
 #include "kernel/util/image_2d.h"
+#include "util/defines.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -85,9 +86,9 @@ ccl_device_noinline T kernel_image_interp_bicubic(const ccl_global KernelImageIn
 }
 
 ccl_device float4 kernel_image_interp(KernelGlobals kg,
+                                      ccl_private ShaderData * /*sd*/,
                                       const int image_texture_id,
-                                      const float x,
-                                      float y)
+                                      const dual2 uv)
 {
   if (image_texture_id == KERNEL_IMAGE_NONE) {
     return IMAGE_MISSING_RGBA;
@@ -97,6 +98,8 @@ ccl_device float4 kernel_image_interp(KernelGlobals kg,
     return IMAGE_MISSING_RGBA;
   }
   const ccl_global KernelImageInfo &info = kernel_data_fetch(image_info, tex.image_info_id);
+  const float x = uv.val.x;
+  const float y = uv.val.y;
 
   /* float4, byte4, ushort4 and half4 */
   const int image_type = info.data_type;
@@ -128,16 +131,16 @@ ccl_device float4 kernel_image_interp(KernelGlobals kg,
 }
 
 ccl_device_forceinline float4 kernel_image_interp_with_udim(KernelGlobals kg,
-                                                            ccl_private ShaderData * /*sd*/,
+                                                            ccl_private ShaderData *sd,
                                                             const int udim_id,
-                                                            float2 uv)
+                                                            dual2 uv)
 {
-  const int image_texture_id = kernel_image_udim_map(kg, udim_id, uv);
+  const int image_texture_id = kernel_image_udim_map(kg, udim_id, uv.val);
   if (image_texture_id == KERNEL_IMAGE_NONE) {
     return IMAGE_MISSING_RGBA;
   }
 
-  return kernel_image_interp(kg, image_texture_id, uv.x, uv.y);
+  return kernel_image_interp(kg, sd, image_texture_id, uv);
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/features.h

@@ -110,6 +110,8 @@ CCL_NAMESPACE_BEGIN
   if constexpr ((node_feature_mask & (KERNEL_FEATURE_##feature)) != 0U)
 #define IF_KERNEL_NODES_FEATURE(feature) \
   if constexpr ((node_feature_mask & (KERNEL_FEATURE_NODE_##feature)) != 0U)
+#define IF_NOT_KERNEL_NODES_FEATURE(feature) \
+  if constexpr ((node_feature_mask & (KERNEL_FEATURE_NODE_##feature)) == 0U)
 
 /* Kernel Feature Guards
  *

intern/cycles/kernel/geom/curve.h

@@ -44,35 +44,36 @@ ccl_device_inline T curve_attribute_dfdy(const ccl_private differential &du,
   return du.dy * (f1 - f0);
 }
 
-/* Read attributes on various curve elements, and compute the partial derivatives if requested. */
+/* Read attributes on various curve elements. T is the return type, which can be a plain type
+ * or a dual type to include derivatives. */
 
 template<typename T>
-ccl_device dual<T> curve_attribute(KernelGlobals kg,
-                                   const ccl_private ShaderData *sd,
-                                   const AttributeDescriptor desc,
-                                   const bool dx = false,
-                                   const bool dy = false)
+ccl_device T curve_attribute(KernelGlobals kg,
+                             const ccl_private ShaderData *sd,
+                             const AttributeDescriptor desc)
 {
-  dual<T> result;
+  using BaseT = dual_base_t<T>;
+
   if (desc.element & ATTR_ELEMENT_CURVE_KEY) {
     const KernelCurve curve = kernel_data_fetch(curves, sd->prim);
     const int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     const int k1 = k0 + 1;
 
-    const T f0 = attribute_data_fetch<T>(kg, desc.element, desc.offset + k0);
-    const T f1 = attribute_data_fetch<T>(kg, desc.element, desc.offset + k1);
+    const BaseT f0 = attribute_data_fetch<BaseT>(kg, desc.element, desc.offset + k0);
+    const BaseT f1 = attribute_data_fetch<BaseT>(kg, desc.element, desc.offset + k1);
 
+    if constexpr (is_dual_v<T>) {
+      T result;
+      result.val = mix(f0, f1, sd->u);
 #  ifdef __RAY_DIFFERENTIALS__
-    if (dx) {
       result.dx = curve_attribute_dfdx(sd->du, f0, f1);
-    }
-    if (dy) {
       result.dy = curve_attribute_dfdy(sd->du, f0, f1);
-    }
 #  endif
-
-    result.val = mix(f0, f1, sd->u);
-    return result;
+      return result;
+    }
+    else {
+      return mix(f0, f1, sd->u);
+    }
   }
 
   /* idea: we can't derive any useful differentials here, but for tiled
@@ -81,9 +82,9 @@ ccl_device dual<T> curve_attribute(KernelGlobals kg,
    * could be computed somehow? */
 
   if (desc.element & ATTR_ELEMENT_CURVE) {
-    return dual<T>(attribute_data_fetch<T>(kg, desc.element, desc.offset + sd->prim));
+    return T(attribute_data_fetch<BaseT>(kg, desc.element, desc.offset + sd->prim));
   }
-  return make_zero<dual<T>>();
+  return make_zero<T>();
 }
 
 /* Curve thickness */
@@ -129,7 +130,7 @@ ccl_device float curve_random(KernelGlobals kg, const ccl_private ShaderData *sd
 {
   if (sd->type & PRIMITIVE_CURVE) {
     const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_CURVE_RANDOM);
-    return (desc.offset != ATTR_STD_NOT_FOUND) ? curve_attribute<float>(kg, sd, desc).val : 0.0f;
+    return (desc.offset != ATTR_STD_NOT_FOUND) ? curve_attribute<float>(kg, sd, desc) : 0.0f;
   }
   return 0.0f;
 }

intern/cycles/kernel/geom/object.h

@@ -148,9 +148,10 @@ ccl_device_inline void object_position_transform(KernelGlobals kg,
 
 /* Transform position from world to object space */
 
+template<class T>
 ccl_device_inline void object_inverse_position_transform(KernelGlobals kg,
                                                          const ccl_private ShaderData *sd,
-                                                         ccl_private float3 *P)
+                                                         ccl_private T *P)
 {
 #ifdef __OBJECT_MOTION__
   if (sd->object_flag & SD_OBJECT_MOTION) {
@@ -163,6 +164,17 @@ ccl_device_inline void object_inverse_position_transform(KernelGlobals kg,
   *P = transform_point(&tfm, *P);
 }
 
+/* Convenience wrapper that checks for OBJECT_NONE before transforming.
+ * Works with both plain types (float3) and dual types (dual3). */
+template<class Float3Type>
+ccl_device_inline void object_inverse_position_transform_if_object(
+    KernelGlobals kg, const ccl_private ShaderData *sd, ccl_private Float3Type *P)
+{
+  if (sd->object != OBJECT_NONE) {
+    object_inverse_position_transform(kg, sd, P);
+  }
+}
+
 /* Transform normal from world to object space */
 
 ccl_device_inline void object_inverse_normal_transform(KernelGlobals kg,
@@ -185,10 +197,10 @@ ccl_device_inline void object_inverse_normal_transform(KernelGlobals kg,
 }
 
 /* Transform normal from object to world space */
-
+template<class T>
 ccl_device_inline void object_normal_transform(KernelGlobals kg,
                                                const ccl_private ShaderData *sd,
-                                               ccl_private float3 *N)
+                                               ccl_private T *N)
 {
 #ifdef __OBJECT_MOTION__
   if (sd->object_flag & SD_OBJECT_MOTION) {

intern/cycles/kernel/geom/point.h

@@ -22,16 +22,14 @@ CCL_NAMESPACE_BEGIN
 /* Reading attributes on various point elements */
 
 template<typename T>
-ccl_device dual<T> point_attribute(KernelGlobals kg,
-                                   const ccl_private ShaderData *sd,
-                                   const AttributeDescriptor desc,
-                                   const bool /* dx */ = false,
-                                   const bool /* dy */ = false)
+ccl_device T point_attribute(KernelGlobals kg,
+                             const ccl_private ShaderData *sd,
+                             const AttributeDescriptor desc)
 {
   if (desc.element & ATTR_ELEMENT_VERTEX) {
-    return dual<T>(attribute_data_fetch<T>(kg, desc.element, desc.offset + sd->prim));
+    return T(attribute_data_fetch<dual_base_t<T>>(kg, desc.element, desc.offset + sd->prim));
   }
-  return make_zero<dual<T>>();
+  return make_zero<T>();
 }
 
 /* Point position */
@@ -81,7 +79,7 @@ ccl_device float point_random(KernelGlobals kg, const ccl_private ShaderData *sd
 {
   if (sd->type & PRIMITIVE_POINT) {
     const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_POINT_RANDOM);
-    return (desc.offset != ATTR_STD_NOT_FOUND) ? point_attribute<float>(kg, sd, desc).val : 0.0f;
+    return (desc.offset != ATTR_STD_NOT_FOUND) ? point_attribute<float>(kg, sd, desc) : 0.0f;
   }
   return 0.0f;
 }

intern/cycles/kernel/geom/primitive.h

@@ -29,31 +29,31 @@ CCL_NAMESPACE_BEGIN
  * heavy volume interpolation code. */
 
 template<typename T>
-ccl_device_forceinline dual<T> primitive_surface_attribute(KernelGlobals kg,
-                                                           const ccl_private ShaderData *sd,
-                                                           const AttributeDescriptor desc,
-                                                           const bool dx = false,
-                                                           const bool dy = false)
+ccl_device_forceinline T primitive_surface_attribute(KernelGlobals kg,
+                                                     const ccl_private ShaderData *sd,
+                                                     const AttributeDescriptor desc)
 {
+  using BaseT = dual_base_t<T>;
+
   if (desc.element & (ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
-    return dual<T>(attribute_data_fetch<T>(kg, desc.element, desc.offset));
+    return T(attribute_data_fetch<BaseT>(kg, desc.element, desc.offset));
   }
 
   if (sd->type & PRIMITIVE_TRIANGLE) {
-    return triangle_attribute<T>(kg, sd, desc, dx, dy);
+    return triangle_attribute<T>(kg, sd, desc);
   }
 #ifdef __HAIR__
   if (sd->type & PRIMITIVE_CURVE) {
-    return curve_attribute<T>(kg, sd, desc, dx, dy);
+    return curve_attribute<T>(kg, sd, desc);
   }
 #endif
 #ifdef __POINTCLOUD__
   else if (sd->type & PRIMITIVE_POINT) {
-    return point_attribute<T>(kg, sd, desc, dx, dy);
+    return point_attribute<T>(kg, sd, desc);
   }
 #endif
   else {
-    return make_zero<dual<T>>();
+    return make_zero<T>();
   }
 }
 
@@ -71,7 +71,7 @@ ccl_device void primitive_normal_set_undisplaced(KernelGlobals kg,
     if (ndesc.offset == ATTR_STD_NOT_FOUND) {
       return;
     }
-    N = safe_normalize(primitive_surface_attribute<float3>(kg, sd, ndesc, false, false).val);
+    N = safe_normalize(primitive_surface_attribute<float3>(kg, sd, ndesc));
   }
   else {
     N = triangle_face_normal_undisplaced(kg, sd, position_undisplaced_offset);
@@ -116,7 +116,7 @@ ccl_device_forceinline float3 primitive_uv(KernelGlobals kg, const ccl_private S
     return make_float3(0.0f, 0.0f, 0.0f);
   }
 
-  const float2 uv = primitive_surface_attribute<float2>(kg, sd, desc).val;
+  const float2 uv = primitive_surface_attribute<float2>(kg, sd, desc);
   return make_float3(uv.x, uv.y, 1.0f);
 }
 
@@ -135,8 +135,8 @@ ccl_device bool primitive_ptex(KernelGlobals kg,
     return false;
   }
 
-  const float3 uv3 = primitive_surface_attribute<float3>(kg, sd, desc_uv).val;
-  const float face_id_f = primitive_surface_attribute<float>(kg, sd, desc_face_id).val;
+  const float3 uv3 = primitive_surface_attribute<float3>(kg, sd, desc_uv);
+  const float face_id_f = primitive_surface_attribute<float>(kg, sd, desc_face_id);
 
   *uv = make_float2(uv3.x, uv3.y);
   *face_id = (int)face_id_f;
@@ -146,15 +146,16 @@ ccl_device bool primitive_ptex(KernelGlobals kg,
 
 /* Surface tangent */
 
-ccl_device float3 primitive_tangent(KernelGlobals kg, ccl_private ShaderData *sd)
+template<typename Float3Type>
+ccl_device Float3Type primitive_tangent(KernelGlobals kg, ccl_private ShaderData *sd)
 {
 #if defined(__HAIR__) || defined(__POINTCLOUD__)
   if (sd->type & (PRIMITIVE_CURVE | PRIMITIVE_POINT)) {
 #  ifdef __DPDU__
-    return normalize(sd->dPdu);
+    return Float3Type(normalize(sd->dPdu));
   }
 #  else
-    return make_float3(0.0f, 0.0f, 0.0f);
+    return make_zero<Float3Type>();
 #  endif
 #endif
 
@@ -162,16 +163,24 @@ ccl_device float3 primitive_tangent(KernelGlobals kg, ccl_private ShaderData *sd
   const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_GENERATED);
 
   if (desc.offset != ATTR_STD_NOT_FOUND) {
-    float3 data = primitive_surface_attribute<float3>(kg, sd, desc).val;
-    data = make_float3(-(data.y - 0.5f), (data.x - 0.5f), 0.0f);
-    object_normal_transform(kg, sd, &data);
-    return cross(sd->N, normalize(cross(data, sd->N)));
+    if constexpr (is_dual_v<Float3Type>) {
+      dual3 data = primitive_surface_attribute<dual3>(kg, sd, desc);
+      data = make_float3(-(data.y() - 0.5f), (data.x() - 0.5f), dual1());
+      object_normal_transform(kg, sd, &data);
+      return cross(sd->N, normalize(cross(data, sd->N)));
+    }
+    else {
+      float3 data = primitive_surface_attribute<float3>(kg, sd, desc);
+      data = make_float3(-(data.y - 0.5f), (data.x - 0.5f), 0.0f);
+      object_normal_transform(kg, sd, &data);
+      return cross(sd->N, normalize(cross(data, sd->N)));
+    }
   }
   /* otherwise use surface derivatives */
 #ifdef __DPDU__
-  return normalize(sd->dPdu);
+  return Float3Type(normalize(sd->dPdu));
 #else
-  return make_float3(0.0f, 0.0f, 0.0f);
+  return make_zero<Float3Type>();
 #endif
 }
 
@@ -221,9 +230,9 @@ ccl_device_forceinline float4 primitive_motion_vector(KernelGlobals kg,
 
 #if defined(__HAIR__) || defined(__POINTCLOUD__)
     if (is_curve_or_point) {
-      motion_pre = make_float3(primitive_surface_attribute<float4>(kg, sd, desc).val);
+      motion_pre = make_float3(primitive_surface_attribute<float4>(kg, sd, desc));
       desc.offset += numverts;
-      motion_post = make_float3(primitive_surface_attribute<float4>(kg, sd, desc).val);
+      motion_post = make_float3(primitive_surface_attribute<float4>(kg, sd, desc));
 
       /* Curve */
       if ((sd->object_flag & SD_OBJECT_HAS_VERTEX_MOTION) == 0) {
@@ -236,9 +245,9 @@ ccl_device_forceinline float4 primitive_motion_vector(KernelGlobals kg,
         if (sd->type & PRIMITIVE_TRIANGLE)
     {
       /* Triangle */
-      motion_pre = triangle_attribute<float3>(kg, sd, desc).val;
+      motion_pre = triangle_attribute<float3>(kg, sd, desc);
       desc.offset += numverts;
-      motion_post = triangle_attribute<float3>(kg, sd, desc).val;
+      motion_post = triangle_attribute<float3>(kg, sd, desc);
     }
   }
 

intern/cycles/kernel/geom/shader_data.h

@@ -36,10 +36,10 @@ ccl_device void shader_setup_object_transforms(KernelGlobals kg,
 
 /* TODO: break this up if it helps reduce register pressure to load data from
  * global memory as we write it to shader-data. */
-ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
-                                             ccl_private ShaderData *ccl_restrict sd,
-                                             const ccl_private Ray *ccl_restrict ray,
-                                             const ccl_private Intersection *ccl_restrict isect)
+ccl_device_noinline void shader_setup_from_ray(KernelGlobals kg,
+                                               ccl_private ShaderData *ccl_restrict sd,
+                                               const ccl_private Ray *ccl_restrict ray,
+                                               const ccl_private Intersection *ccl_restrict isect)
 {
   /* Read intersection data into shader globals.
    *
@@ -367,6 +367,7 @@ ccl_device_inline void shader_setup_from_background(KernelGlobals kg,
                                                     ccl_private ShaderData *ccl_restrict sd,
                                                     const float3 ray_P,
                                                     const float3 ray_D,
+                                                    const float ray_dD,
                                                     const float ray_time)
 {
   /* for NDC coordinates */
@@ -391,16 +392,17 @@ ccl_device_inline void shader_setup_from_background(KernelGlobals kg,
 
 #ifdef __DPDU__
   /* dPdu/dPdv */
-  sd->dPdu = zero_float3();
-  sd->dPdv = zero_float3();
+  /* Construct arbitrary local coordinate system. */
+  make_orthonormals(sd->Ng, &sd->dPdu, &sd->dPdv);
 #endif
 
 #ifdef __RAY_DIFFERENTIALS__
   /* differentials */
-  sd->dP = differential_zero_compact(); /* TODO: ray->dP */
-  sd->dI = differential_zero_compact();
-  sd->du = differential_zero();
-  sd->dv = differential_zero();
+  sd->dP = ray_dD;
+  sd->dI = differential_incoming_compact(ray_dD);
+  /* Make the uv coordinate system match the constructed local coordinate system. */
+  sd->du.dx = sd->dv.dy = sd->dP;
+  sd->du.dy = sd->dv.dx = 0.0f;
 #endif
 }
 

intern/cycles/kernel/geom/triangle.h

@@ -279,16 +279,15 @@ ccl_device_inline T triangle_attribute_dfdy(const ccl_private differential &du,
   return du.dy * f1 + dv.dy * f2 - (du.dy + dv.dy) * f0;
 }
 
-/* Read attributes on various triangle elements, and compute the partial derivatives if requested.
- */
+/* Read attributes on various triangle elements. T is the return type, which can be a plain type
+ * (float, float3, etc.) or a dual type (dual1, dual3, etc.) to include derivatives. */
 template<typename T>
-ccl_device dual<T> triangle_attribute(KernelGlobals kg,
-                                      const ccl_private ShaderData *sd,
-                                      const AttributeDescriptor desc,
-                                      const bool dx = false,
-                                      const bool dy = false)
+ccl_device T triangle_attribute(KernelGlobals kg,
+                                const ccl_private ShaderData *sd,
+                                const AttributeDescriptor desc)
 {
-  dual<T> result;
+  using BaseT = dual_base_t<T>;
+
   if (desc.element & (ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_CORNER)) {
     int i0, i1, i2;
 
@@ -306,26 +305,27 @@ ccl_device dual<T> triangle_attribute(KernelGlobals kg,
       i2 = tri + 2;
     }
 
-    T f[3];
-    attribute_data_fetch_3<T>(kg, desc.element, desc.offset, i0, i1, i2, f);
+    BaseT f[3];
+    attribute_data_fetch_3<BaseT>(kg, desc.element, desc.offset, i0, i1, i2, f);
 
+    if constexpr (is_dual_v<T>) {
+      T result;
+      result.val = triangle_interpolate(sd->u, sd->v, f[0], f[1], f[2]);
 #ifdef __RAY_DIFFERENTIALS__
-    if (dx) {
       result.dx = triangle_attribute_dfdx(sd->du, sd->dv, f[0], f[1], f[2]);
-    }
-    if (dy) {
       result.dy = triangle_attribute_dfdy(sd->du, sd->dv, f[0], f[1], f[2]);
-    }
 #endif
-
-    result.val = triangle_interpolate(sd->u, sd->v, f[0], f[1], f[2]);
-    return result;
+      return result;
+    }
+    else {
+      return triangle_interpolate(sd->u, sd->v, f[0], f[1], f[2]);
+    }
   }
 
   if (desc.element & ATTR_ELEMENT_FACE) {
-    return dual<T>(attribute_data_fetch<T>(kg, desc.element, desc.offset + sd->prim));
+    return T(attribute_data_fetch<BaseT>(kg, desc.element, desc.offset + sd->prim));
   }
-  return make_zero<dual<T>>();
+  return make_zero<T>();
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/geom/volume.h

@@ -26,15 +26,14 @@ CCL_NAMESPACE_BEGIN
 
 /* Return position normalized to 0..1 in mesh bounds */
 
-ccl_device_inline float3 volume_normalized_position(KernelGlobals kg,
-                                                    const ccl_private ShaderData *sd,
-                                                    float3 P)
+template<typename Float3Type>
+ccl_device_inline Float3Type volume_normalized_position(KernelGlobals kg,
+                                                        const ccl_private ShaderData *sd,
+                                                        Float3Type P)
 {
-  /* todo: optimize this so it's just a single matrix multiplication when
-   * possible (not motion blur), or perhaps even just translation + scale */
   const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_GENERATED_TRANSFORM);
 
-  object_inverse_position_transform(kg, sd, &P);
+  object_inverse_position_transform_if_object(kg, sd, &P);
 
   if (desc.offset != ATTR_STD_NOT_FOUND) {
     const Transform tfm = primitive_attribute_matrix(kg, desc);

intern/cycles/kernel/integrator/shade_background.h

@@ -51,6 +51,7 @@ ccl_device Spectrum integrator_eval_background_shader(KernelGlobals kg,
                                emission_sd,
                                INTEGRATOR_STATE(state, ray, P),
                                INTEGRATOR_STATE(state, ray, D),
+                               INTEGRATOR_STATE(state, ray, dD),
                                INTEGRATOR_STATE(state, ray, time));
 
   PROFILING_SHADER(emission_sd->object, emission_sd->shader);

intern/cycles/kernel/integrator/shade_light.h

@@ -134,7 +134,12 @@ ccl_device bool integrate_light_nee(KernelGlobals kg, IntegratorShadowState stat
 
     if (light_type == LIGHT_BACKGROUND) {
       /* Background light. */
-      shader_setup_from_background(kg, emission_sd, ray.P, ray.D, ray.time);
+#ifdef __RAY_DIFFERENTIALS__
+      const float ray_dD = ray.dD;
+#else
+      const float ray_dD = 0.0f;
+#endif
+      shader_setup_from_background(kg, emission_sd, ray.P, ray.D, ray_dD, ray.time);
       is_background = true;
     }
     else {

intern/cycles/kernel/osl/closures.cpp

@@ -137,7 +137,7 @@ void osl_eval_nodes_surface(const ThreadKernelGlobalsCPU *kg,
         const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_POSITION_UNDISPLACED);
         kernel_assert(desc.offset != ATTR_STD_NOT_FOUND);
 
-        dual3 P = primitive_surface_attribute<float3>(kg, sd, desc, true, true);
+        dual3 P = primitive_surface_attribute<dual3>(kg, sd, desc);
         object_position_transform(kg, sd, &P);
 
         sd->P = P.val;

intern/cycles/kernel/osl/osl.h

@@ -220,7 +220,7 @@ ccl_device_inline void osl_eval_nodes(KernelGlobals kg,
         const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_POSITION_UNDISPLACED);
         kernel_assert(desc.offset != ATTR_STD_NOT_FOUND);
 
-        dual3 P = primitive_surface_attribute<float3>(kg, sd, desc, true, true);
+        dual3 P = primitive_surface_attribute<dual3>(kg, sd, desc);
 
         object_position_transform(kg, sd, &P);
 

intern/cycles/kernel/osl/services.cpp

@@ -16,6 +16,7 @@
 #include "util/colorspace.h"
 #include "util/log.h"
 #include "util/string.h"
+#include "util/types_image.h"
 
 #include "kernel/device/cpu/image.h"
 
@@ -426,7 +427,7 @@ bool OSLRenderServices::texture(OSLUStringHash filename,
     }
     case OSLTextureHandle::SVM: {
       const float4 rgba = kernel_image_interp_with_udim(
-          kernel_globals, sd, handle->id, make_float2(s, 1.0f - t));
+          kernel_globals, sd, handle->id, dual2(make_float2(s, 1.0f - t)));
 
       result[0] = rgba[0];
       if (nchannels > 1) {

intern/cycles/kernel/osl/services_gpu.h

@@ -307,7 +307,7 @@ ccl_device_extern bool rs_texture(ccl_private ShaderGlobals *sg,
     case OSL_TEXTURE_HANDLE_TYPE_SVM: {
       ccl_private ShaderData *sd = sg->sd;
       const float4 rgba = kernel_image_interp_with_udim(
-          nullptr, sd, image_texture_id, make_float2(s, 1.0f - t));
+          nullptr, sd, image_texture_id, dual2(make_float2(s, 1.0f - t)));
       if (nchannels > 0) {
         result[0] = rgba.x;
       }

intern/cycles/kernel/osl/services_shared.h

@@ -592,7 +592,12 @@ ccl_device_inline bool osl_shared_get_object_attribute_impl(KernelGlobals kg,
   else
 #endif
   {
-    data = primitive_surface_attribute<T>(kg, sd, desc, derivatives, derivatives);
+    if (derivatives) {
+      data = primitive_surface_attribute<dual<T>>(kg, sd, desc);
+    }
+    else {
+      data = dual<T>(primitive_surface_attribute<T>(kg, sd, desc));
+    }
   }
   return set_attribute(data, type, derivatives, val);
 }
@@ -702,10 +707,10 @@ ccl_device_inline bool osl_shared_get_background_attribute(KernelGlobals kg,
     if ((sg->raytype & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE &&
         kernel_data.cam.type == CAMERA_ORTHOGRAPHIC)
     {
-      ndc.val = camera_world_to_ndc(kg, sd, sd->ray_P);
+      ndc = dual3(camera_world_to_ndc(kg, sd, sd->ray_P));
     }
     else {
-      ndc.val = camera_world_to_ndc(kg, sd, sd->P);
+      ndc = dual3(camera_world_to_ndc(kg, sd, sd->P));
 
       if (derivatives) {
         const differential3 dP = differential_from_compact(sd->Ng, sd->dP);

intern/cycles/kernel/svm/attribute.h

@@ -13,8 +13,6 @@
 
 #include "kernel/svm/util.h"
 
-#include "kernel/util/differential.h"
-
 CCL_NAMESPACE_BEGIN
 
 /* Attribute Node */
@@ -49,193 +47,118 @@ ccl_device AttributeDescriptor svm_node_attr_init(KernelGlobals kg,
   return desc;
 }
 
+/* Store attribute to the stack. Float3Type is float3 or dual3. */
+
+template<typename Float3Type>
 ccl_device_inline void svm_node_attr_store(const NodeAttributeOutputType type,
                                            ccl_private float *stack,
                                            const uint out_offset,
-                                           const float f)
+                                           const ccl_private Float3Type &f)
 {
-  if (type == NODE_ATTR_OUTPUT_FLOAT) {
-    stack_store_float(stack, out_offset, f);
-  }
-  else if (type == NODE_ATTR_OUTPUT_FLOAT3) {
-    stack_store_float3(stack, out_offset, make_float3(f));
+  using FloatType = dual_scalar_t<Float3Type>;
+  if (type == NODE_ATTR_OUTPUT_FLOAT3) {
+    stack_store(stack, out_offset, f);
   }
   else {
-    stack_store_float(stack, out_offset, 1.0f);
+    stack_store(stack, out_offset, FloatType(average(f)));
   }
 }
 
-ccl_device_inline void svm_node_attr_store(const NodeAttributeOutputType type,
-                                           ccl_private float *stack,
-                                           const uint out_offset,
-                                           const ccl_private float2 &f)
+/* Core surface attribute evaluation, returning Float3Type = float3 or dual3.
+ * Fetches the attribute, applies output type conversion (float3 or scalar-as-float3),
+ * and computes derivatives when Float3Type is a dual type. */
+
+template<typename Float3Type>
+ccl_device_inline Float3Type svm_node_attr_surface_eval(KernelGlobals kg,
+                                                        ccl_private ShaderData *sd,
+                                                        const uint4 node,
+                                                        const NodeAttributeOutputType type,
+                                                        const AttributeDescriptor desc)
 {
-  if (type == NODE_ATTR_OUTPUT_FLOAT) {
-    stack_store_float(stack, out_offset, f.x);
-  }
-  else if (type == NODE_ATTR_OUTPUT_FLOAT3) {
-    stack_store_float3(stack, out_offset, make_float3(f));
-  }
-  else {
-    stack_store_float(stack, out_offset, 1.0f);
-  }
-}
-
-ccl_device_inline void svm_node_attr_store(const NodeAttributeOutputType type,
-                                           ccl_private float *stack,
-                                           const uint out_offset,
-                                           const ccl_private float3 &f)
-{
-  if (type == NODE_ATTR_OUTPUT_FLOAT) {
-    stack_store_float(stack, out_offset, average(f));
-  }
-  else if (type == NODE_ATTR_OUTPUT_FLOAT3) {
-    stack_store_float3(stack, out_offset, f);
-  }
-  else {
-    stack_store_float(stack, out_offset, 1.0f);
-  }
-}
-
-ccl_device_inline void svm_node_attr_store(const NodeAttributeOutputType type,
-                                           ccl_private float *stack,
-                                           const uint out_offset,
-                                           const ccl_private float4 &f)
-{
-  if (type == NODE_ATTR_OUTPUT_FLOAT) {
-    stack_store_float(stack, out_offset, average(make_float3(f)));
-  }
-  else if (type == NODE_ATTR_OUTPUT_FLOAT3) {
-    stack_store_float3(stack, out_offset, make_float3(f));
-  }
-  else {
-    kernel_assert(type == NODE_ATTR_OUTPUT_FLOAT_ALPHA);
-    stack_store_float(stack, out_offset, f.w);
-  }
-}
-
-template<class T>
-ccl_device_inline void svm_surface_attr(KernelGlobals kg,
-                                        const ccl_private ShaderData *sd,
-                                        const AttributeDescriptor desc,
-                                        const NodeAttributeOutputType type,
-                                        ccl_private float *stack,
-                                        const uint out_offset)
-{
-  T f = primitive_surface_attribute<T>(kg, sd, desc).val;
-  svm_node_attr_store(type, stack, out_offset, f);
-}
-
-template<class T>
-ccl_device_inline void svm_surface_attr_dx(KernelGlobals kg,
-                                           const ccl_private ShaderData *sd,
-                                           const AttributeDescriptor desc,
-                                           const float bump_filter_width,
-                                           const NodeAttributeOutputType type,
-                                           ccl_private float *stack,
-                                           const uint out_offset)
-{
-  dual<T> f = primitive_surface_attribute<T>(kg, sd, desc, true, false);
-  f.val += f.dx * bump_filter_width;
-  svm_node_attr_store(type, stack, out_offset, f.val);
-}
-
-template<class T>
-ccl_device_inline void svm_surface_attr_dy(KernelGlobals kg,
-                                           const ccl_private ShaderData *sd,
-                                           const AttributeDescriptor desc,
-                                           const float bump_filter_width,
-                                           const NodeAttributeOutputType type,
-                                           ccl_private float *stack,
-                                           const uint out_offset)
-{
-  dual<T> f = primitive_surface_attribute<T>(kg, sd, desc, false, true);
-  f.val += f.dy * bump_filter_width;
-  svm_node_attr_store(type, stack, out_offset, f.val);
-}
-
-template<uint node_feature_mask>
-ccl_device_noinline void svm_node_attr(KernelGlobals kg,
-                                       ccl_private ShaderData *sd,
-                                       ccl_private float *stack,
-                                       const uint4 node)
-{
-  NodeAttributeOutputType type = NODE_ATTR_OUTPUT_FLOAT;
-  uint out_offset = 0;
-  const AttributeDescriptor desc = svm_node_attr_init(kg, sd, node, &type, &out_offset);
-
-#ifdef __VOLUME__
-  IF_KERNEL_NODES_FEATURE(VOLUME)
-  {
-    /* Volumes
-     * NOTE: moving this into its own node type might help improve performance. */
-    if (primitive_is_volume_attribute(sd)) {
-      const bool stochastic_sample = node.w;
-      const float4 value = volume_attribute_float4(kg, sd, desc, stochastic_sample);
-
-      if (type == NODE_ATTR_OUTPUT_FLOAT) {
-        const float f = volume_attribute_value<float>(value);
-        stack_store_float(stack, out_offset, f);
-      }
-      else if (type == NODE_ATTR_OUTPUT_FLOAT3) {
-        const float3 f = volume_attribute_value<float3>(value);
-        stack_store_float3(stack, out_offset, f);
-      }
-      else {
-        const float f = volume_attribute_alpha(value);
-        stack_store_float(stack, out_offset, f);
-      }
-      return;
-    }
-  }
-#endif
+  using FloatType = dual_scalar_t<Float3Type>;
 
   if (sd->type == PRIMITIVE_LAMP && node.y == ATTR_STD_UV) {
-    stack_store_float3(stack, out_offset, make_float3(1.0f - sd->u - sd->v, sd->u, 0.0f));
-    return;
+    Float3Type uv(make_float3(1.0f - sd->u - sd->v, sd->u, 0.0f));
+    if constexpr (is_dual_v<Float3Type>) {
+      uv.dx = make_float3(-sd->du.dx - sd->dv.dx, sd->du.dx, 0.0f);
+      uv.dy = make_float3(-sd->du.dy - sd->dv.dy, sd->du.dy, 0.0f);
+    }
+    return uv;
   }
 
   if (node.y == ATTR_STD_GENERATED && desc.element == ATTR_ELEMENT_NONE) {
-    /* No generated attribute, fall back to object coordinates. */
-    float3 f = sd->P;
-    if (sd->object != OBJECT_NONE) {
-      object_inverse_position_transform(kg, sd, &f);
-    }
-    svm_node_attr_store(type, stack, out_offset, f);
-    return;
+    Float3Type f = shading_position<Float3Type>(sd);
+    object_inverse_position_transform_if_object(kg, sd, &f);
+    return f;
   }
 
-  /* Surface. */
+  /* Surface attribute fetch with output type conversion. */
   if (desc.type == NODE_ATTR_FLOAT) {
-    svm_surface_attr<float>(kg, sd, desc, type, stack, out_offset);
+    FloatType f = primitive_surface_attribute<FloatType>(kg, sd, desc);
+    if (type == NODE_ATTR_OUTPUT_FLOAT_ALPHA) {
+      return make_float3(FloatType(1.0f));
+    }
+    return make_float3(f, f, f);
   }
-  else if (desc.type == NODE_ATTR_FLOAT2) {
-    svm_surface_attr<float2>(kg, sd, desc, type, stack, out_offset);
+
+  if (desc.type == NODE_ATTR_FLOAT2) {
+    if constexpr (is_dual_v<Float3Type>) {
+      dual2 f = primitive_surface_attribute<dual2>(kg, sd, desc);
+      if (type == NODE_ATTR_OUTPUT_FLOAT) {
+        return make_float3(f.x());
+      }
+      if (type == NODE_ATTR_OUTPUT_FLOAT_ALPHA) {
+        return make_float3(FloatType(1.0f));
+      }
+      return make_float3(f);
+    }
+    else {
+      float2 f = primitive_surface_attribute<float2>(kg, sd, desc);
+      if (type == NODE_ATTR_OUTPUT_FLOAT) {
+        return make_float3(f.x);
+      }
+      if (type == NODE_ATTR_OUTPUT_FLOAT_ALPHA) {
+        return make_float3(FloatType(1.0f));
+      }
+      return make_float3(f);
+    }
   }
-  else if (desc.type == NODE_ATTR_FLOAT4 || desc.type == NODE_ATTR_RGBA) {
-    svm_surface_attr<float4>(kg, sd, desc, type, stack, out_offset);
+
+  if (desc.type == NODE_ATTR_FLOAT4 || desc.type == NODE_ATTR_RGBA) {
+    if constexpr (is_dual_v<Float3Type>) {
+      dual4 f = primitive_surface_attribute<dual4>(kg, sd, desc);
+      if (type == NODE_ATTR_OUTPUT_FLOAT) {
+        return make_float3(average(make_float3(f)));
+      }
+      if (type == NODE_ATTR_OUTPUT_FLOAT_ALPHA) {
+        return make_float3(f.w());
+      }
+      return make_float3(f);
+    }
+    else {
+      float4 f = primitive_surface_attribute<float4>(kg, sd, desc);
+      if (type == NODE_ATTR_OUTPUT_FLOAT) {
+        return make_float3(average(make_float3(f)));
+      }
+      if (type == NODE_ATTR_OUTPUT_FLOAT_ALPHA) {
+        return make_float3(f.w);
+      }
+      return make_float3(f);
+    }
   }
-  else {
-    svm_surface_attr<float3>(kg, sd, desc, type, stack, out_offset);
+
+  Float3Type f = primitive_surface_attribute<Float3Type>(kg, sd, desc);
+  if (type == NODE_ATTR_OUTPUT_FLOAT) {
+    return make_float3(average(f));
   }
+  if (type == NODE_ATTR_OUTPUT_FLOAT_ALPHA) {
+    return make_float3(FloatType(1.0f));
+  }
+  return f;
 }
 
-/* Position offsetted in x direction. */
-ccl_device_forceinline float3 svm_node_bump_P_dx(const ccl_private ShaderData *sd,
-                                                 const float bump_filter_width)
-{
-  return sd->P + dPdx(sd) * bump_filter_width;
-}
-
-/* Position offsetted in y direction. */
-ccl_device_forceinline float3 svm_node_bump_P_dy(const ccl_private ShaderData *sd,
-                                                 const float bump_filter_width)
-{
-  return sd->P + dPdy(sd) * bump_filter_width;
-}
-
-/* Evaluate attributes at a position shifted in x direction. */
-ccl_device_noinline void svm_node_attr_bump_dx(KernelGlobals kg,
+/* Surface attribute node. */
+ccl_device_noinline void svm_node_attr_surface(KernelGlobals kg,
                                                ccl_private ShaderData *sd,
                                                ccl_private float *stack,
                                                const uint4 node)
@@ -243,83 +166,66 @@ ccl_device_noinline void svm_node_attr_bump_dx(KernelGlobals kg,
   NodeAttributeOutputType type = NODE_ATTR_OUTPUT_FLOAT;
   uint out_offset = 0;
   const AttributeDescriptor desc = svm_node_attr_init(kg, sd, node, &type, &out_offset);
-  const float bump_filter_width = __uint_as_float(node.w);
 
-#ifdef __VOLUME__
-  /* Volume */
-  if (primitive_is_volume_attribute(sd)) {
-    svm_node_attr_store(type, stack, out_offset, 0.0f);
-    return;
-  }
-#endif
-
-  if (node.y == ATTR_STD_GENERATED && desc.element == ATTR_ELEMENT_NONE) {
-    /* No generated attribute, fall back to object coordinates. */
-    float3 f_x = svm_node_bump_P_dx(sd, bump_filter_width);
-    if (sd->object != OBJECT_NONE) {
-      object_inverse_position_transform(kg, sd, &f_x);
-    }
-    svm_node_attr_store(type, stack, out_offset, f_x);
-    return;
-  }
-
-  /* Surface */
-  if (desc.type == NODE_ATTR_FLOAT) {
-    svm_surface_attr_dx<float>(kg, sd, desc, bump_filter_width, type, stack, out_offset);
-  }
-  else if (desc.type == NODE_ATTR_FLOAT2) {
-    svm_surface_attr_dx<float2>(kg, sd, desc, bump_filter_width, type, stack, out_offset);
-  }
-  else if (desc.type == NODE_ATTR_FLOAT4 || desc.type == NODE_ATTR_RGBA) {
-    svm_surface_attr_dx<float4>(kg, sd, desc, bump_filter_width, type, stack, out_offset);
-  }
-  else {
-    svm_surface_attr_dx<float3>(kg, sd, desc, bump_filter_width, type, stack, out_offset);
-  }
+  float3 data = svm_node_attr_surface_eval<float3>(kg, sd, node, type, desc);
+  svm_node_attr_store(type, stack, out_offset, data);
 }
 
-/* Evaluate attributes at a position shifted in y direction. */
-ccl_device_noinline void svm_node_attr_bump_dy(KernelGlobals kg,
-                                               ccl_private ShaderData *sd,
-                                               ccl_private float *stack,
-                                               const uint4 node)
+/* Evaluate surface attributes with derivatives and optional bump offset.
+ * Used for derivative tracking and bump mapping. */
+
+ccl_device_noinline void svm_node_attr_derivative(KernelGlobals kg,
+                                                  ccl_private ShaderData *sd,
+                                                  ccl_private float *stack,
+                                                  const uint4 node)
 {
   NodeAttributeOutputType type = NODE_ATTR_OUTPUT_FLOAT;
   uint out_offset = 0;
   const AttributeDescriptor desc = svm_node_attr_init(kg, sd, node, &type, &out_offset);
+
+  uint unused1, unused2, bump_offset, store_derivatives;
+  svm_unpack_node_uchar4(node.z, &unused1, &unused2, &bump_offset, &store_derivatives);
   const float bump_filter_width = __uint_as_float(node.w);
 
-#ifdef __VOLUME__
-  /* Volume */
-  if (primitive_is_volume_attribute(sd)) {
-    svm_node_attr_store(type, stack, out_offset, 0.0f);
-    return;
+  dual3 data = svm_node_attr_surface_eval<dual3>(kg, sd, node, type, desc);
+  if (bump_offset == NODE_BUMP_OFFSET_DX) {
+    data.val += data.dx * bump_filter_width;
   }
-#endif
-
-  if (node.y == ATTR_STD_GENERATED && desc.element == ATTR_ELEMENT_NONE) {
-    /* No generated attribute, fall back to object coordinates. */
-    /* TODO:(weizhen) */
-    float3 f_y = svm_node_bump_P_dy(sd, bump_filter_width);
-    if (sd->object != OBJECT_NONE) {
-      object_inverse_position_transform(kg, sd, &f_y);
-    }
-    svm_node_attr_store(type, stack, out_offset, f_y);
-    return;
+  else if (bump_offset == NODE_BUMP_OFFSET_DY) {
+    data.val += data.dy * bump_filter_width;
   }
 
-  /* Surface */
-  if (desc.type == NODE_ATTR_FLOAT) {
-    svm_surface_attr_dy<float>(kg, sd, desc, bump_filter_width, type, stack, out_offset);
-  }
-  else if (desc.type == NODE_ATTR_FLOAT2) {
-    svm_surface_attr_dy<float2>(kg, sd, desc, bump_filter_width, type, stack, out_offset);
-  }
-  else if (desc.type == NODE_ATTR_FLOAT4 || desc.type == NODE_ATTR_RGBA) {
-    svm_surface_attr_dy<float4>(kg, sd, desc, bump_filter_width, type, stack, out_offset);
+  if (store_derivatives) {
+    svm_node_attr_store(type, stack, out_offset, data);
   }
   else {
-    svm_surface_attr_dy<float3>(kg, sd, desc, bump_filter_width, type, stack, out_offset);
+    svm_node_attr_store(type, stack, out_offset, float3(data.val));
+  }
+}
+
+/* Volume attribute node. Volumes have no derivatives or bump. */
+ccl_device_noinline void svm_node_attr_volume(KernelGlobals kg,
+                                              ccl_private ShaderData *sd,
+                                              ccl_private float *stack,
+                                              const uint4 node)
+{
+  kernel_assert(primitive_is_volume_attribute(sd));
+
+  NodeAttributeOutputType type = NODE_ATTR_OUTPUT_FLOAT;
+  uint out_offset = 0;
+  const AttributeDescriptor desc = svm_node_attr_init(kg, sd, node, &type, &out_offset);
+
+  const bool stochastic_sample = node.w;
+  const float4 value = volume_attribute_float4(kg, sd, desc, stochastic_sample);
+
+  if (type == NODE_ATTR_OUTPUT_FLOAT) {
+    stack_store_float(stack, out_offset, volume_attribute_value<float>(value));
+  }
+  else if (type == NODE_ATTR_OUTPUT_FLOAT3) {
+    stack_store_float3(stack, out_offset, volume_attribute_value<float3>(value));
+  }
+  else {
+    stack_store_float(stack, out_offset, volume_attribute_alpha(value));
   }
 }
 

intern/cycles/kernel/svm/bump.h

@@ -30,7 +30,7 @@ ccl_device_noinline void svm_node_enter_bump_eval(KernelGlobals kg,
   /* Set position as if undisplaced. */
   const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_POSITION_UNDISPLACED);
   if (desc.offset != ATTR_STD_NOT_FOUND) {
-    dual3 attr = primitive_surface_attribute<float3>(kg, sd, desc, true, true);
+    dual3 attr = primitive_surface_attribute<dual3>(kg, sd, desc);
     object_position_transform(kg, sd, &attr);
 
     sd->P = attr.val;

intern/cycles/kernel/svm/closure.h

@@ -921,7 +921,7 @@ ccl_device
       const AttributeDescriptor attr_descr_random = find_attribute(kg, sd, data_node2.y);
       float random = 0.0f;
       if (attr_descr_random.offset != ATTR_STD_NOT_FOUND) {
-        random = primitive_surface_attribute<float>(kg, sd, attr_descr_random).val;
+        random = primitive_surface_attribute<float>(kg, sd, attr_descr_random);
       }
       else {
         random = stack_load_float_default(stack, random_ofs, data_node3.y);
@@ -1053,7 +1053,7 @@ ccl_device
           if (bsdf->aspect_ratio != 1.0f) {
             /* Align ellipse major axis with the curve normal direction. */
             const AttributeDescriptor attr_descr_normal = find_attribute(kg, sd, shared_ofs2);
-            bsdf->N = curve_attribute<float3>(kg, sd, attr_descr_normal).val;
+            bsdf->N = curve_attribute<float3>(kg, sd, attr_descr_normal);
           }
 
           bsdf->roughness = roughness;

intern/cycles/kernel/svm/convert.h

@@ -12,24 +12,27 @@ CCL_NAMESPACE_BEGIN
 
 /* Conversion Nodes */
 
+template<typename FloatType, typename Float3Type>
 ccl_device_noinline void svm_node_convert(
     KernelGlobals kg, ccl_private float *stack, const uint type, const uint from, const uint to)
 {
+
   switch ((NodeConvert)type) {
     case NODE_CONVERT_FI: {
+      /* TODO(weizhen): should actually store 0 for int, but none of the nodes that we compute
+       * derivatives for has int inputs, so seems fine. */
       const float f = stack_load_float(stack, from);
       stack_store_int(stack, to, float_to_int(f));
       break;
     }
     case NODE_CONVERT_FV: {
-      const float f = stack_load_float(stack, from);
-      stack_store_float3(stack, to, make_float3(f, f, f));
+      const FloatType f = stack_load<FloatType>(stack, from);
+      stack_store(stack, to, make_float3(f, f, f));
       break;
     }
     case NODE_CONVERT_CF: {
-      const float3 f = stack_load_float3(stack, from);
-      const float g = linear_rgb_to_gray(kg, f);
-      stack_store_float(stack, to, g);
+      const Float3Type f = stack_load<Float3Type>(stack, from);
+      stack_store(stack, to, linear_rgb_to_gray(kg, f));
       break;
     }
     case NODE_CONVERT_CI: {
@@ -39,9 +42,8 @@ ccl_device_noinline void svm_node_convert(
       break;
     }
     case NODE_CONVERT_VF: {
-      const float3 f = stack_load_float3(stack, from);
-      const float g = average(f);
-      stack_store_float(stack, to, g);
+      const Float3Type f = stack_load<Float3Type>(stack, from);
+      stack_store(stack, to, average(f));
       break;
     }
     case NODE_CONVERT_VI: {
@@ -52,14 +54,16 @@ ccl_device_noinline void svm_node_convert(
     }
     case NODE_CONVERT_IF: {
       const float f = (float)stack_load_int(stack, from);
-      stack_store_float(stack, to, f);
+      stack_store(stack, to, FloatType(f));
       break;
     }
     case NODE_CONVERT_IV: {
       const float f = (float)stack_load_int(stack, from);
-      stack_store_float3(stack, to, make_float3(f, f, f));
+      stack_store(stack, to, Float3Type(make_float3(f, f, f)));
       break;
     }
+    default:
+      assert(false);
   }
 }
 

intern/cycles/kernel/svm/displace.h

@@ -206,7 +206,7 @@ ccl_device_noinline int svm_node_vector_displacement(KernelGlobals kg,
       const AttributeDescriptor attr = find_attribute(kg, sd, node.z);
       float3 tangent;
       if (attr.offset != ATTR_STD_NOT_FOUND) {
-        tangent = primitive_surface_attribute<float3>(kg, sd, attr).val;
+        tangent = primitive_surface_attribute<float3>(kg, sd, attr);
       }
       else {
         tangent = normalize(sd->dPdu);
@@ -215,7 +215,7 @@ ccl_device_noinline int svm_node_vector_displacement(KernelGlobals kg,
       float3 bitangent = safe_normalize(cross(normal, tangent));
       const AttributeDescriptor attr_sign = find_attribute(kg, sd, node.w);
       if (attr_sign.offset != ATTR_STD_NOT_FOUND) {
-        const float sign = primitive_surface_attribute<float>(kg, sd, attr_sign).val;
+        const float sign = primitive_surface_attribute<float>(kg, sd, attr_sign);
         bitangent *= sign;
       }
 

intern/cycles/kernel/svm/geometry.h

@@ -16,102 +16,78 @@ CCL_NAMESPACE_BEGIN
 
 /* Geometry Node */
 
-ccl_device_noinline void svm_node_geometry(KernelGlobals kg,
-                                           ccl_private ShaderData *sd,
-                                           ccl_private float *stack,
-                                           const uint type,
-                                           const uint out_offset)
+template<typename Float3Type>
+ccl_device_inline Float3Type svm_node_geometry_eval(KernelGlobals kg,
+                                                    ccl_private ShaderData *sd,
+                                                    const uint type)
 {
-  float3 data;
+  Float3Type data;
 
   switch (type) {
     case NODE_GEOM_P:
-      data = sd->P;
+      data = shading_position<Float3Type>(sd);
       break;
     case NODE_GEOM_N:
-      data = sd->N;
+      data = Float3Type(sd->N);
       break;
 #ifdef __DPDU__
     case NODE_GEOM_T:
-      data = primitive_tangent(kg, sd);
+      data = primitive_tangent<Float3Type>(kg, sd);
       break;
 #endif
     case NODE_GEOM_I:
-      data = sd->wi;
+      data = shading_incoming<Float3Type>(sd);
       break;
     case NODE_GEOM_Ng:
-      data = sd->Ng;
+      data = Float3Type(sd->Ng);
       break;
     case NODE_GEOM_uv:
-      data = make_float3(1.0f - sd->u - sd->v, sd->u, 0.0f);
+      data = Float3Type(make_float3(1.0f - sd->u - sd->v, sd->u, 0.0f));
+      if constexpr (is_dual_v<Float3Type>) {
+        data.dx = make_float3(-sd->du.dx - sd->dv.dx, sd->du.dx, 0.0f);
+        data.dy = make_float3(-sd->du.dy - sd->dv.dy, sd->du.dy, 0.0f);
+      }
       break;
     default:
-      data = make_float3(0.0f, 0.0f, 0.0f);
+      data = Float3Type(make_float3(0.0f, 0.0f, 0.0f));
   }
 
-  stack_store_float3(stack, out_offset, data);
+  return data;
 }
 
-ccl_device_noinline void svm_node_geometry_bump_dx(KernelGlobals kg,
-                                                   ccl_private ShaderData *sd,
-                                                   ccl_private float *stack,
-                                                   const uint type,
-                                                   const uint out_offset,
-                                                   const float bump_filter_width)
+ccl_device_noinline void svm_node_geometry(KernelGlobals kg,
+                                           ccl_private ShaderData *sd,
+                                           ccl_private float *stack,
+                                           const uint4 node)
 {
-#ifdef __RAY_DIFFERENTIALS__
-  float3 data;
-
-  switch (type) {
-    case NODE_GEOM_P:
-      data = svm_node_bump_P_dx(sd, bump_filter_width);
-      break;
-    case NODE_GEOM_uv: {
-      const float u_x = sd->u + sd->du.dx * bump_filter_width;
-      const float v_x = sd->v + sd->dv.dx * bump_filter_width;
-      data = make_float3(1.0f - u_x - v_x, u_x, 0.0f);
-      break;
-    }
-    default:
-      svm_node_geometry(kg, sd, stack, type, out_offset);
-      return;
-  }
-
-  stack_store_float3(stack, out_offset, data);
-#else
-  svm_node_geometry(kg, sd, stack, type, out_offset);
-#endif
+  uint type, unused1, unused2;
+  svm_unpack_node_uchar3(node.y, &type, &unused1, &unused2);
+  stack_store(stack, node.z, svm_node_geometry_eval<float3>(kg, sd, type));
 }
 
-ccl_device_noinline void svm_node_geometry_bump_dy(KernelGlobals kg,
-                                                   ccl_private ShaderData *sd,
-                                                   ccl_private float *stack,
-                                                   const uint type,
-                                                   const uint out_offset,
-                                                   const float bump_filter_width)
+ccl_device_noinline void svm_node_geometry_derivative(KernelGlobals kg,
+                                                      ccl_private ShaderData *sd,
+                                                      ccl_private float *stack,
+                                                      const uint4 node)
 {
-#ifdef __RAY_DIFFERENTIALS__
-  float3 data;
+  uint type, bump_offset, store_derivatives;
+  svm_unpack_node_uchar3(node.y, &type, &bump_offset, &store_derivatives);
 
-  switch (type) {
-    case NODE_GEOM_P:
-      data = svm_node_bump_P_dy(sd, bump_filter_width);
-      break;
-    case NODE_GEOM_uv: {
-      const float u_y = sd->u + sd->du.dy * bump_filter_width;
-      const float v_y = sd->v + sd->dv.dy * bump_filter_width;
-      data = make_float3(1.0f - u_y - v_y, u_y, 0.0f);
-      break;
-    }
-    default:
-      svm_node_geometry(kg, sd, stack, type, out_offset);
-      return;
+  /* Compute with derivatives, then apply first-order bump offset. */
+  dual3 data = svm_node_geometry_eval<dual3>(kg, sd, type);
+  const float bump_filter_width = __uint_as_float(node.w);
+  if (bump_offset == NODE_BUMP_OFFSET_DX) {
+    data.val += data.dx * bump_filter_width;
+  }
+  else if (bump_offset == NODE_BUMP_OFFSET_DY) {
+    data.val += data.dy * bump_filter_width;
+  }
+  if (store_derivatives) {
+    stack_store(stack, node.z, data);
+  }
+  else {
+    stack_store(stack, node.z, data.val);
   }
-
-  stack_store_float3(stack, out_offset, data);
-#else
-  svm_node_geometry(kg, sd, stack, type, out_offset);
-#endif
 }
 
 /* Object Info */

intern/cycles/kernel/svm/image.h

@@ -19,7 +19,7 @@
 CCL_NAMESPACE_BEGIN
 
 ccl_device float4 svm_image_texture(
-    KernelGlobals kg, ccl_private ShaderData *sd, const int id, const float2 uv, const uint flags)
+    KernelGlobals kg, ccl_private ShaderData *sd, const int id, const dual2 uv, const uint flags)
 {
   float4 r = kernel_image_interp_with_udim(kg, sd, id, uv);
   const float alpha = r.w;
@@ -37,16 +37,29 @@ ccl_device float4 svm_image_texture(
 }
 
 /* Remap coordinate from 0..1 box to -1..-1 */
-ccl_device_inline float3 texco_remap_square(const float3 co)
+template<class Float3Type> ccl_device_inline Float3Type texco_remap_square(const Float3Type co)
 {
   return (co - make_float3(0.5f, 0.5f, 0.5f)) * 2.0f;
 }
 
-ccl_device_noinline int svm_node_tex_image(KernelGlobals kg,
-                                           ccl_private ShaderData *sd,
-                                           ccl_private float *stack,
-                                           const uint4 node,
-                                           int offset)
+template<class Float3Type>
+ccl_device_inline auto svm_node_tex_image_mapping(const Float3Type co, const uint proj)
+{
+  if (proj == NODE_IMAGE_PROJ_SPHERE) {
+    return map_to_sphere(texco_remap_square(co));
+  }
+  if (proj == NODE_IMAGE_PROJ_TUBE) {
+    return map_to_tube(texco_remap_square(co));
+  }
+
+  return make_float2(co);
+}
+
+ccl_device_noinline void svm_node_tex_image(KernelGlobals kg,
+                                            ccl_private ShaderData *sd,
+                                            ccl_private float *stack,
+                                            const uint4 node,
+                                            const bool derivative)
 {
   uint co_offset;
   uint out_offset;
@@ -55,36 +68,32 @@ ccl_device_noinline int svm_node_tex_image(KernelGlobals kg,
 
   svm_unpack_node_uchar4(node.z, &co_offset, &out_offset, &alpha_offset, &flags);
 
-  float3 co = stack_load_float3(stack, co_offset);
-  float2 tex_co;
-  if (node.w == NODE_IMAGE_PROJ_SPHERE) {
-    co = texco_remap_square(co);
-    tex_co = map_to_sphere(co);
-  }
-  else if (node.w == NODE_IMAGE_PROJ_TUBE) {
-    co = texco_remap_square(co);
-    tex_co = map_to_tube(co);
+  dual2 tex_co;
+  if (derivative) {
+    const dual3 co = stack_load<dual3>(stack, co_offset);
+    tex_co = svm_node_tex_image_mapping(co, node.w);
   }
   else {
-    tex_co = make_float2(co.x, co.y);
+    const float3 co = stack_load_float3(stack, co_offset);
+    tex_co = dual2(svm_node_tex_image_mapping(co, node.w));
   }
 
   const int id = node.y;
   const float4 f = svm_image_texture(kg, sd, id, tex_co, flags);
 
   if (stack_valid(out_offset)) {
-    stack_store_float3(stack, out_offset, make_float3(f.x, f.y, f.z));
+    stack_store_float3(stack, out_offset, make_float3(f));
   }
   if (stack_valid(alpha_offset)) {
     stack_store_float(stack, alpha_offset, f.w);
   }
-  return offset;
 }
 
 ccl_device_noinline void svm_node_tex_image_box(KernelGlobals kg,
                                                 ccl_private ShaderData *sd,
                                                 ccl_private float *stack,
-                                                const uint4 node)
+                                                const uint4 node,
+                                                const bool derivative)
 {
   /* get object space normal */
   float3 N = sd->N;
@@ -158,22 +167,23 @@ ccl_device_noinline void svm_node_tex_image_box(KernelGlobals kg,
   uint flags;
   svm_unpack_node_uchar4(node.z, &co_offset, &out_offset, &alpha_offset, &flags);
 
-  const float3 co = stack_load_float3(stack, co_offset);
   const uint id = node.y;
-
   float4 f = zero_float4();
 
+  const dual3 co = (derivative) ? stack_load<dual3>(stack, co_offset) :
+                                  dual3(stack_load_float3(stack, co_offset));
+
   /* Map so that no textures are flipped, rotation is somewhat arbitrary. */
   if (weight.x > 0.0f) {
-    const float2 uv = make_float2((signed_N.x < 0.0f) ? 1.0f - co.y : co.y, co.z);
+    const dual2 uv = make_float2((signed_N.x < 0.0f) ? 1.0f - co.y() : co.y(), co.z());
     f += weight.x * svm_image_texture(kg, sd, id, uv, flags);
   }
   if (weight.y > 0.0f) {
-    const float2 uv = make_float2((signed_N.y > 0.0f) ? 1.0f - co.x : co.x, co.z);
+    const dual2 uv = make_float2((signed_N.y > 0.0f) ? 1.0f - co.x() : co.x(), co.z());
     f += weight.y * svm_image_texture(kg, sd, id, uv, flags);
   }
   if (weight.z > 0.0f) {
-    const float2 uv = make_float2((signed_N.z > 0.0f) ? 1.0f - co.y : co.y, co.x);
+    const dual2 uv = make_float2((signed_N.z > 0.0f) ? 1.0f - co.y() : co.y(), co.x());
     f += weight.z * svm_image_texture(kg, sd, id, uv, flags);
   }
 
@@ -185,30 +195,38 @@ ccl_device_noinline void svm_node_tex_image_box(KernelGlobals kg,
   }
 }
 
+template<class Float3Type>
+ccl_device_inline auto svm_node_tex_environment_projection(Float3Type co, const uint proj)
+{
+  co = safe_normalize(co);
+  if (proj == 0) {
+    return direction_to_equirectangular(co);
+  }
+  return direction_to_mirrorball(co);
+}
+
 ccl_device_noinline void svm_node_tex_environment(KernelGlobals kg,
                                                   ccl_private ShaderData *sd,
                                                   ccl_private float *stack,
-                                                  const uint4 node)
+                                                  const uint4 node,
+                                                  const bool derivative)
 {
   const uint id = node.y;
   uint co_offset;
   uint out_offset;
   uint alpha_offset;
   uint flags;
-  const uint projection = node.w;
 
   svm_unpack_node_uchar4(node.z, &co_offset, &out_offset, &alpha_offset, &flags);
 
-  float3 co = stack_load_float3(stack, co_offset);
-  float2 uv;
-
-  co = safe_normalize(co);
-
-  if (projection == 0) {
-    uv = direction_to_equirectangular(co);
+  dual2 uv;
+  if (derivative) {
+    const dual3 co = stack_load<dual3>(stack, co_offset);
+    uv = svm_node_tex_environment_projection(co, node.w);
   }
   else {
-    uv = direction_to_mirrorball(co);
+    const float3 co = stack_load_float3(stack, co_offset);
+    uv = dual2(svm_node_tex_environment_projection(co, node.w));
   }
 
   const float4 f = svm_image_texture(kg, sd, id, uv, flags);

intern/cycles/kernel/svm/mapping.h

@@ -11,6 +11,7 @@ CCL_NAMESPACE_BEGIN
 
 /* Mapping Node */
 
+template<typename Float3Type>
 ccl_device_noinline void svm_node_mapping(ccl_private float *stack,
                                           const uint type,
                                           const uint inputs_stack_offsets,
@@ -26,13 +27,13 @@ ccl_device_noinline void svm_node_mapping(ccl_private float *stack,
                          &rotation_stack_offset,
                          &scale_stack_offset);
 
-  const float3 vector = stack_load_float3(stack, vector_stack_offset);
   const float3 location = stack_load_float3(stack, location_stack_offset);
   const float3 rotation = stack_load_float3(stack, rotation_stack_offset);
   const float3 scale = stack_load_float3(stack, scale_stack_offset);
 
-  const float3 result = svm_mapping((NodeMappingType)type, vector, location, rotation, scale);
-  stack_store_float3(stack, result_stack_offset, result);
+  const Float3Type vector = stack_load<Float3Type>(stack, vector_stack_offset);
+  const Float3Type result = svm_mapping((NodeMappingType)type, vector, location, rotation, scale);
+  stack_store(stack, result_stack_offset, result);
 }
 
 /* Texture Mapping */

intern/cycles/kernel/svm/mapping_util.h

@@ -12,11 +12,12 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device float3 svm_mapping(NodeMappingType type,
-                              const float3 vector,
-                              const float3 location,
-                              const float3 rotation,
-                              const float3 scale)
+template<class Float3Type>
+ccl_device Float3Type svm_mapping(NodeMappingType type,
+                                  const Float3Type vector,
+                                  const float3 location,
+                                  const float3 rotation,
+                                  const float3 scale)
 {
   const Transform rotationTransform = euler_to_transform(rotation);
   switch (type) {
@@ -30,7 +31,7 @@ ccl_device float3 svm_mapping(NodeMappingType type,
     case NODE_MAPPING_TYPE_NORMAL:
       return safe_normalize(transform_direction(&rotationTransform, safe_divide(vector, scale)));
     default:
-      return make_float3(0.0f, 0.0f, 0.0f);
+      return Float3Type(zero_float3());
   }
 }
 

intern/cycles/kernel/svm/math.h

@@ -27,6 +27,7 @@ ccl_device_noinline void svm_node_math(ccl_private float *stack,
   stack_store_float(stack, result_stack_offset, result);
 }
 
+template<typename Float3Type>
 ccl_device_noinline int svm_node_vector_math(KernelGlobals kg,
                                              ccl_private float *stack,
                                              const uint type,
@@ -34,6 +35,8 @@ ccl_device_noinline int svm_node_vector_math(KernelGlobals kg,
                                              const uint outputs_stack_offsets,
                                              int offset)
 {
+  using FloatType = dual_scalar_t<Float3Type>;
+
   uint value_stack_offset;
   uint vector_stack_offset;
   uint a_stack_offset;
@@ -43,30 +46,30 @@ ccl_device_noinline int svm_node_vector_math(KernelGlobals kg,
       inputs_stack_offsets, &a_stack_offset, &b_stack_offset, &param1_stack_offset);
   svm_unpack_node_uchar2(outputs_stack_offsets, &value_stack_offset, &vector_stack_offset);
 
-  const float3 a = stack_load_float3(stack, a_stack_offset);
-  const float3 b = stack_load_float3(stack, b_stack_offset);
-  float3 c = make_float3(0.0f, 0.0f, 0.0f);
-  const float param1 = stack_load_float(stack, param1_stack_offset);
-
-  float value;
-  float3 vector;
+  const Float3Type a = stack_load<Float3Type>(stack, a_stack_offset);
+  const Float3Type b = stack_load<Float3Type>(stack, b_stack_offset);
+  Float3Type c = make_zero<Float3Type>();
+  const FloatType param1 = stack_load<FloatType>(stack, param1_stack_offset);
 
   /* 3 Vector Operators */
   if (type == NODE_VECTOR_MATH_WRAP || type == NODE_VECTOR_MATH_FACEFORWARD ||
       type == NODE_VECTOR_MATH_MULTIPLY_ADD)
   {
     const uint4 extra_node = read_node(kg, &offset);
-    c = stack_load_float3(stack, extra_node.x);
+    c = stack_load<Float3Type>(stack, extra_node.x);
   }
 
+  FloatType value = make_zero<FloatType>();
+  Float3Type vector = make_zero<Float3Type>();
   svm_vector_math(&value, &vector, (NodeVectorMathType)type, a, b, c, param1);
 
   if (stack_valid(value_stack_offset)) {
-    stack_store_float(stack, value_stack_offset, value);
+    stack_store(stack, value_stack_offset, value);
   }
   if (stack_valid(vector_stack_offset)) {
-    stack_store_float3(stack, vector_stack_offset, vector);
+    stack_store(stack, vector_stack_offset, vector);
   }
+
   return offset;
 }
 

intern/cycles/kernel/svm/math_util.h

@@ -12,13 +12,14 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device void svm_vector_math(ccl_private float *value,
-                                ccl_private float3 *vector,
+template<class Float3Type, class FloatType>
+ccl_device void svm_vector_math(ccl_private FloatType *value,
+                                ccl_private Float3Type *vector,
                                 NodeVectorMathType type,
-                                const float3 a,
-                                const float3 b,
-                                const float3 c,
-                                float param1)
+                                const Float3Type a,
+                                const Float3Type b,
+                                const Float3Type c,
+                                const FloatType param1)
 {
   switch (type) {
     case NODE_VECTOR_MATH_ADD:
@@ -112,8 +113,8 @@ ccl_device void svm_vector_math(ccl_private float *value,
       *vector = tan(a);
       break;
     default:
-      *vector = zero_float3();
-      *value = 0.0f;
+      *vector = Float3Type(zero_float3());
+      *value = FloatType(0.0f);
   }
 }
 

intern/cycles/kernel/svm/node_types_template.h

@@ -5,9 +5,13 @@
 #ifndef SHADER_NODE_TYPE
 #  define SHADER_NODE_TYPE(name)
 #endif
+#ifndef SHADER_NODE_TYPE_DERIVATIVE
+#  define SHADER_NODE_TYPE_DERIVATIVE(name) SHADER_NODE_TYPE(name)
+#endif
 
-/* NOTE: for best OpenCL performance, item definition in the enum must
- * match the switch case order in `svm.h`. */
+/* NOTE: For good performance with jump tables on some GPU backends, the enum must
+ * match the switch order in `svm.h`. It is also assumed the derivative variation
+ * directly follows the regular node type. */
 
 SHADER_NODE_TYPE(NODE_END)
 SHADER_NODE_TYPE(NODE_SHADER_JUMP)
@@ -20,28 +24,20 @@ SHADER_NODE_TYPE(NODE_EMISSION_WEIGHT)
 SHADER_NODE_TYPE(NODE_MIX_CLOSURE)
 SHADER_NODE_TYPE(NODE_JUMP_IF_ZERO)
 SHADER_NODE_TYPE(NODE_JUMP_IF_ONE)
-SHADER_NODE_TYPE(NODE_GEOMETRY)
-SHADER_NODE_TYPE(NODE_CONVERT)
-SHADER_NODE_TYPE(NODE_TEX_COORD)
-SHADER_NODE_TYPE(NODE_VALUE_F)
-SHADER_NODE_TYPE(NODE_VALUE_V)
-SHADER_NODE_TYPE(NODE_ATTR)
-SHADER_NODE_TYPE(NODE_VERTEX_COLOR)
-SHADER_NODE_TYPE(NODE_GEOMETRY_BUMP_DX)
-SHADER_NODE_TYPE(NODE_GEOMETRY_BUMP_DY)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_GEOMETRY)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_CONVERT)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_TEX_COORD)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_VALUE_F)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_VALUE_V)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_ATTR)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_VERTEX_COLOR)
 SHADER_NODE_TYPE(NODE_SET_DISPLACEMENT)
 SHADER_NODE_TYPE(NODE_DISPLACEMENT)
 SHADER_NODE_TYPE(NODE_VECTOR_DISPLACEMENT)
-SHADER_NODE_TYPE(NODE_TEX_IMAGE)
-SHADER_NODE_TYPE(NODE_TEX_IMAGE_BOX)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_TEX_IMAGE)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_TEX_IMAGE_BOX)
 SHADER_NODE_TYPE(NODE_TEX_NOISE)
 SHADER_NODE_TYPE(NODE_SET_BUMP)
-SHADER_NODE_TYPE(NODE_ATTR_BUMP_DX)
-SHADER_NODE_TYPE(NODE_ATTR_BUMP_DY)
-SHADER_NODE_TYPE(NODE_VERTEX_COLOR_BUMP_DX)
-SHADER_NODE_TYPE(NODE_VERTEX_COLOR_BUMP_DY)
-SHADER_NODE_TYPE(NODE_TEX_COORD_BUMP_DX)
-SHADER_NODE_TYPE(NODE_TEX_COORD_BUMP_DY)
 SHADER_NODE_TYPE(NODE_CLOSURE_SET_NORMAL)
 SHADER_NODE_TYPE(NODE_ENTER_BUMP_EVAL)
 SHADER_NODE_TYPE(NODE_LEAVE_BUMP_EVAL)
@@ -53,7 +49,7 @@ SHADER_NODE_TYPE(NODE_CLOSURE_VOLUME)
 SHADER_NODE_TYPE(NODE_VOLUME_COEFFICIENTS)
 SHADER_NODE_TYPE(NODE_PRINCIPLED_VOLUME)
 SHADER_NODE_TYPE(NODE_MATH)
-SHADER_NODE_TYPE(NODE_VECTOR_MATH)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_VECTOR_MATH)
 SHADER_NODE_TYPE(NODE_RGB_RAMP)
 SHADER_NODE_TYPE(NODE_GAMMA)
 SHADER_NODE_TYPE(NODE_BRIGHTCONTRAST)
@@ -63,10 +59,10 @@ SHADER_NODE_TYPE(NODE_PARTICLE_INFO)
 SHADER_NODE_TYPE(NODE_HAIR_INFO)
 SHADER_NODE_TYPE(NODE_POINT_INFO)
 SHADER_NODE_TYPE(NODE_TEXTURE_MAPPING)
-SHADER_NODE_TYPE(NODE_MAPPING)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_MAPPING)
 SHADER_NODE_TYPE(NODE_MIN_MAX)
 SHADER_NODE_TYPE(NODE_CAMERA)
-SHADER_NODE_TYPE(NODE_TEX_ENVIRONMENT)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_TEX_ENVIRONMENT)
 SHADER_NODE_TYPE(NODE_TEX_SKY)
 SHADER_NODE_TYPE(NODE_TEX_GRADIENT)
 SHADER_NODE_TYPE(NODE_TEX_VORONOI)
@@ -80,15 +76,15 @@ SHADER_NODE_TYPE(NODE_NORMAL)
 SHADER_NODE_TYPE(NODE_LIGHT_FALLOFF)
 SHADER_NODE_TYPE(NODE_IES)
 SHADER_NODE_TYPE(NODE_CURVES)
-SHADER_NODE_TYPE(NODE_TANGENT)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_TANGENT)
 SHADER_NODE_TYPE(NODE_NORMAL_MAP)
 SHADER_NODE_TYPE(NODE_RADIAL_TILING)
 SHADER_NODE_TYPE(NODE_INVERT)
 SHADER_NODE_TYPE(NODE_MIX)
 SHADER_NODE_TYPE(NODE_SEPARATE_COLOR)
 SHADER_NODE_TYPE(NODE_COMBINE_COLOR)
-SHADER_NODE_TYPE(NODE_SEPARATE_VECTOR)
-SHADER_NODE_TYPE(NODE_COMBINE_VECTOR)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_SEPARATE_VECTOR)
+SHADER_NODE_TYPE_DERIVATIVE(NODE_COMBINE_VECTOR)
 SHADER_NODE_TYPE(NODE_VECTOR_ROTATE)
 SHADER_NODE_TYPE(NODE_VECTOR_TRANSFORM)
 SHADER_NODE_TYPE(NODE_WIREFRAME)
@@ -108,8 +104,10 @@ SHADER_NODE_TYPE(NODE_MIX_COLOR)
 SHADER_NODE_TYPE(NODE_MIX_FLOAT)
 SHADER_NODE_TYPE(NODE_MIX_VECTOR)
 SHADER_NODE_TYPE(NODE_MIX_VECTOR_NON_UNIFORM)
+SHADER_NODE_TYPE(NODE_NONE)
 
 /* Padding for struct alignment. */
 SHADER_NODE_TYPE(NODE_PAD1)
 
 #undef SHADER_NODE_TYPE
+#undef SHADER_NODE_TYPE_DERIVATIVE

intern/cycles/kernel/svm/sepcomb_vector.h

@@ -10,34 +10,57 @@ CCL_NAMESPACE_BEGIN
 
 /* Vector combine / separate, used for the RGB and XYZ nodes */
 
+template<typename Float3Type>
 ccl_device void svm_node_combine_vector(ccl_private float *stack,
                                         const uint in_offset,
                                         const uint vector_index,
                                         const uint out_offset)
 {
-  const float vector = stack_load_float(stack, in_offset);
+  using FloatType = dual_scalar_t<Float3Type>;
+  const FloatType value = stack_load<FloatType>(stack, in_offset);
 
   if (stack_valid(out_offset)) {
-    stack_store_float(stack, out_offset + vector_index, vector);
+    if constexpr (is_dual_v<Float3Type>) {
+      stack_store_float(stack, out_offset + vector_index, value.val);
+      stack_store_float(stack, out_offset + vector_index + 3, value.dx);
+      stack_store_float(stack, out_offset + vector_index + 6, value.dy);
+    }
+    else {
+      stack_store_float(stack, out_offset + vector_index, value);
+    }
   }
 }
 
+template<typename Float3Type>
 ccl_device void svm_node_separate_vector(ccl_private float *stack,
                                          const uint ivector_offset,
                                          const uint vector_index,
                                          const uint out_offset)
 {
-  const float3 vector = stack_load_float3(stack, ivector_offset);
+  const Float3Type vector = stack_load<Float3Type>(stack, ivector_offset);
 
   if (stack_valid(out_offset)) {
-    if (vector_index == 0) {
-      stack_store_float(stack, out_offset, vector.x);
-    }
-    else if (vector_index == 1) {
-      stack_store_float(stack, out_offset, vector.y);
+    if constexpr (is_dual_v<Float3Type>) {
+      if (vector_index == 0) {
+        stack_store(stack, out_offset, vector.x());
+      }
+      else if (vector_index == 1) {
+        stack_store(stack, out_offset, vector.y());
+      }
+      else {
+        stack_store(stack, out_offset, vector.z());
+      }
     }
     else {
-      stack_store_float(stack, out_offset, vector.z);
+      if (vector_index == 0) {
+        stack_store(stack, out_offset, vector.x);
+      }
+      else if (vector_index == 1) {
+        stack_store(stack, out_offset, vector.y);
+      }
+      else {
+        stack_store(stack, out_offset, vector.z);
+      }
     }
   }
 }

intern/cycles/kernel/svm/sky.h

@@ -10,8 +10,10 @@
 #include "kernel/svm/util.h"
 
 #include "kernel/util/colorspace.h"
+#include "kernel/util/differential.h"
 
 #include "util/color.h"
+#include "util/defines.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -134,6 +136,7 @@ ccl_device float3 geographical_to_direction(const float lat, const float lon)
 }
 
 ccl_device float3 sky_radiance_nishita(KernelGlobals kg,
+                                       ccl_private ShaderData *sd,
                                        const float3 dir,
                                        const uint32_t path_flag,
                                        const float3 pixel_bottom,
@@ -172,13 +175,14 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals kg,
   const float x = fractf((-direction.y - M_PI_2_F + sun_rotation) * M_1_2PI_F);
   /* Undo the non-linear transformation from the sky LUT */
   const float y = copysignf(sqrtf(fabsf(dir_elevation) * M_2_PI_F), dir_elevation) * 0.5f + 0.5f;
-  xyz += make_float3(kernel_image_interp(kg, texture_id, x, y));
+  xyz += make_float3(kernel_image_interp(kg, sd, texture_id, dual2(make_float2(x, y))));
 
   /* Convert to RGB */
   return xyz_to_rgb_clamped(kg, xyz);
 }
 
 ccl_device_noinline int svm_node_tex_sky(KernelGlobals kg,
+                                         ccl_private ShaderData *sd,
                                          const uint32_t path_flag,
                                          ccl_private float *stack,
                                          const uint4 node,
@@ -299,7 +303,8 @@ ccl_device_noinline int svm_node_tex_sky(KernelGlobals kg,
     const uint texture_id = __float_as_uint(data.w);
 
     /* Compute Sky */
-    rgb = sky_radiance_nishita(kg, dir, path_flag, pixel_bottom, pixel_top, sky_data, texture_id);
+    rgb = sky_radiance_nishita(
+        kg, sd, dir, path_flag, pixel_bottom, pixel_top, sky_data, texture_id);
   }
 
   stack_store_float3(stack, out_offset, rgb);

intern/cycles/kernel/svm/svm.h

@@ -172,36 +172,72 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
       }
       break;
       SVM_CASE(NODE_GEOMETRY)
-      svm_node_geometry(kg, sd, stack, node.y, node.z);
+      svm_node_geometry(kg, sd, stack, node);
+      break;
+      SVM_CASE(NODE_GEOMETRY_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_geometry_derivative(kg, sd, stack, node);
+      }
       break;
       SVM_CASE(NODE_CONVERT)
-      svm_node_convert(kg, stack, node.y, node.z, node.w);
+      svm_node_convert<float, float3>(kg, stack, node.y, node.z, node.w);
+      break;
+      SVM_CASE(NODE_CONVERT_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_convert<dual1, dual3>(kg, stack, node.y, node.z, node.w);
+      }
       break;
       SVM_CASE(NODE_TEX_COORD)
       offset = svm_node_tex_coord(kg, sd, path_flag, stack, node, offset);
       break;
+      SVM_CASE(NODE_TEX_COORD_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        offset = svm_node_tex_coord_derivative(kg, sd, path_flag, stack, node, offset);
+      }
+      break;
       SVM_CASE(NODE_VALUE_F)
-      svm_node_value_f(stack, node.y, node.z);
+      svm_node_value_f<float>(stack, node.y, node.z);
+      break;
+      SVM_CASE(NODE_VALUE_F_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_value_f<dual1>(stack, node.y, node.z);
+      }
       break;
       SVM_CASE(NODE_VALUE_V)
-      offset = svm_node_value_v(kg, stack, node.y, offset);
+      offset = svm_node_value_v<float3>(kg, stack, node.y, offset);
+      break;
+      SVM_CASE(NODE_VALUE_V_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        offset = svm_node_value_v<dual3>(kg, stack, node.y, offset);
+      }
       break;
       SVM_CASE(NODE_ATTR)
-      svm_node_attr<node_feature_mask>(kg, sd, stack, node);
+      IF_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_attr_volume(kg, sd, stack, node);
+      }
+      else {
+        svm_node_attr_surface(kg, sd, stack, node);
+      }
+      break;
+      SVM_CASE(NODE_ATTR_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_attr_derivative(kg, sd, stack, node);
+      }
       break;
       SVM_CASE(NODE_VERTEX_COLOR)
       svm_node_vertex_color(kg, sd, stack, node);
       break;
-      SVM_CASE(NODE_GEOMETRY_BUMP_DX)
-      IF_KERNEL_NODES_FEATURE(BUMP)
+      SVM_CASE(NODE_VERTEX_COLOR_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
       {
-        svm_node_geometry_bump_dx(kg, sd, stack, node.y, node.z, __uint_as_float(node.w));
-      }
-      break;
-      SVM_CASE(NODE_GEOMETRY_BUMP_DY)
-      IF_KERNEL_NODES_FEATURE(BUMP)
-      {
-        svm_node_geometry_bump_dy(kg, sd, stack, node.y, node.z, __uint_as_float(node.w));
+        svm_node_vertex_color_derivative(kg, sd, stack, node);
       }
       break;
       SVM_CASE(NODE_SET_DISPLACEMENT)
@@ -214,10 +250,22 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
       offset = svm_node_vector_displacement<node_feature_mask>(kg, sd, stack, node, offset);
       break;
       SVM_CASE(NODE_TEX_IMAGE)
-      offset = svm_node_tex_image(kg, sd, stack, node, offset);
+      svm_node_tex_image(kg, sd, stack, node, false);
+      break;
+      SVM_CASE(NODE_TEX_IMAGE_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_tex_image(kg, sd, stack, node, true);
+      }
       break;
       SVM_CASE(NODE_TEX_IMAGE_BOX)
-      svm_node_tex_image_box(kg, sd, stack, node);
+      svm_node_tex_image_box(kg, sd, stack, node, false);
+      break;
+      SVM_CASE(NODE_TEX_IMAGE_BOX_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_tex_image_box(kg, sd, stack, node, true);
+      }
       break;
       SVM_CASE(NODE_TEX_NOISE)
       offset = svm_node_tex_noise(kg, stack, node.y, node.z, node.w, offset);
@@ -225,42 +273,6 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
       SVM_CASE(NODE_SET_BUMP)
       offset = svm_node_set_bump<node_feature_mask>(kg, sd, stack, node, offset);
       break;
-      SVM_CASE(NODE_ATTR_BUMP_DX)
-      IF_KERNEL_NODES_FEATURE(BUMP)
-      {
-        svm_node_attr_bump_dx(kg, sd, stack, node);
-      }
-      break;
-      SVM_CASE(NODE_ATTR_BUMP_DY)
-      IF_KERNEL_NODES_FEATURE(BUMP)
-      {
-        svm_node_attr_bump_dy(kg, sd, stack, node);
-      }
-      break;
-      SVM_CASE(NODE_VERTEX_COLOR_BUMP_DX)
-      IF_KERNEL_NODES_FEATURE(BUMP)
-      {
-        svm_node_vertex_color_bump_dx(kg, sd, stack, node);
-      }
-      break;
-      SVM_CASE(NODE_VERTEX_COLOR_BUMP_DY)
-      IF_KERNEL_NODES_FEATURE(BUMP)
-      {
-        svm_node_vertex_color_bump_dy(kg, sd, stack, node);
-      }
-      break;
-      SVM_CASE(NODE_TEX_COORD_BUMP_DX)
-      IF_KERNEL_NODES_FEATURE(BUMP)
-      {
-        offset = svm_node_tex_coord_bump_dx(kg, sd, path_flag, stack, node, offset);
-      }
-      break;
-      SVM_CASE(NODE_TEX_COORD_BUMP_DY)
-      IF_KERNEL_NODES_FEATURE(BUMP)
-      {
-        offset = svm_node_tex_coord_bump_dy(kg, sd, path_flag, stack, node, offset);
-      }
-      break;
       SVM_CASE(NODE_CLOSURE_SET_NORMAL)
       IF_KERNEL_NODES_FEATURE(BUMP)
       {
@@ -314,7 +326,13 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
       svm_node_math(stack, node.y, node.z, node.w);
       break;
       SVM_CASE(NODE_VECTOR_MATH)
-      offset = svm_node_vector_math(kg, stack, node.y, node.z, node.w, offset);
+      offset = svm_node_vector_math<float3>(kg, stack, node.y, node.z, node.w, offset);
+      break;
+      SVM_CASE(NODE_VECTOR_MATH_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        offset = svm_node_vector_math<dual3>(kg, stack, node.y, node.z, node.w, offset);
+      }
       break;
       SVM_CASE(NODE_RGB_RAMP)
       offset = svm_node_rgb_ramp(kg, stack, node, offset);
@@ -348,7 +366,13 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
       offset = svm_node_texture_mapping(kg, stack, node.y, node.z, offset);
       break;
       SVM_CASE(NODE_MAPPING)
-      svm_node_mapping(stack, node.y, node.z, node.w);
+      svm_node_mapping<float3>(stack, node.y, node.z, node.w);
+      break;
+      SVM_CASE(NODE_MAPPING_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_mapping<dual3>(stack, node.y, node.z, node.w);
+      }
       break;
       SVM_CASE(NODE_MIN_MAX)
       offset = svm_node_min_max(kg, stack, node.y, node.z, offset);
@@ -357,10 +381,16 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
       svm_node_camera(kg, sd, stack, node.y, node.z, node.w);
       break;
       SVM_CASE(NODE_TEX_ENVIRONMENT)
-      svm_node_tex_environment(kg, sd, stack, node);
+      svm_node_tex_environment(kg, sd, stack, node, false);
+      break;
+      SVM_CASE(NODE_TEX_ENVIRONMENT_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_tex_environment(kg, sd, stack, node, true);
+      }
       break;
       SVM_CASE(NODE_TEX_SKY)
-      offset = svm_node_tex_sky(kg, path_flag, stack, node, offset);
+      offset = svm_node_tex_sky(kg, sd, path_flag, stack, node, offset);
       break;
       SVM_CASE(NODE_TEX_GRADIENT)
       svm_node_tex_gradient(stack, node);
@@ -399,7 +429,13 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
       offset = svm_node_curves(kg, stack, node, offset);
       break;
       SVM_CASE(NODE_TANGENT)
-      svm_node_tangent(kg, sd, stack, node);
+      svm_node_tangent<float3>(kg, sd, stack, node);
+      break;
+      SVM_CASE(NODE_TANGENT_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_tangent<dual3>(kg, sd, stack, node);
+      }
       break;
       SVM_CASE(NODE_NORMAL_MAP)
       svm_node_normal_map(kg, sd, stack, node);
@@ -420,10 +456,22 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
       svm_node_combine_color(stack, node.y, node.z, node.w);
       break;
       SVM_CASE(NODE_SEPARATE_VECTOR)
-      svm_node_separate_vector(stack, node.y, node.z, node.w);
+      svm_node_separate_vector<float3>(stack, node.y, node.z, node.w);
+      break;
+      SVM_CASE(NODE_SEPARATE_VECTOR_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_separate_vector<dual3>(stack, node.y, node.z, node.w);
+      }
       break;
       SVM_CASE(NODE_COMBINE_VECTOR)
-      svm_node_combine_vector(stack, node.y, node.z, node.w);
+      svm_node_combine_vector<float3>(stack, node.y, node.z, node.w);
+      break;
+      SVM_CASE(NODE_COMBINE_VECTOR_DERIVATIVE)
+      IF_NOT_KERNEL_NODES_FEATURE(VOLUME)
+      {
+        svm_node_combine_vector<dual3>(stack, node.y, node.z, node.w);
+      }
       break;
       SVM_CASE(NODE_VECTOR_ROTATE)
       svm_node_vector_rotate(stack, node.y, node.z, node.w);

intern/cycles/kernel/svm/tex_coord.h

@@ -17,30 +17,167 @@
 
 CCL_NAMESPACE_BEGIN
 
+/* Smooth normal with screen-space derivatives for texture coordinate use.
+ * Returns the interpolated normal in object space, with dx/dy representing
+ * the per-pixel change from ray differentials. */
+ccl_device_inline dual3 svm_texco_smooth_normal(KernelGlobals kg, const ccl_private ShaderData *sd)
+{
+  if ((sd->type & PRIMITIVE_TRIANGLE) && (sd->shader & SHADER_SMOOTH_NORMAL)) {
+    float3 N_x, N_y;
+    float3 N;
+    if (sd->type == PRIMITIVE_TRIANGLE) {
+      N = triangle_smooth_normal(kg,
+                                 sd->Ng,
+                                 sd->object,
+                                 sd->object_flag,
+                                 sd->prim,
+                                 sd->u,
+                                 sd->v,
+                                 sd->du,
+                                 sd->dv,
+                                 N_x,
+                                 N_y);
+    }
+    else {
+      N = motion_triangle_smooth_normal(
+          kg, sd->Ng, sd->object, sd->prim, sd->time, sd->u, sd->v, sd->du, sd->dv, N_x, N_y);
+    }
+    if (sd->flag & SD_BACKFACING) {
+      N = -N;
+      N_x = -N_x;
+      N_y = -N_y;
+    }
+    if (sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED) {
+      object_inverse_normal_transform(kg, sd, &N);
+      object_inverse_normal_transform(kg, sd, &N_x);
+      object_inverse_normal_transform(kg, sd, &N_y);
+    }
+    return dual3(N, N_x - N, N_y - N);
+  }
+
+  /* Flat normal or non-triangle: no derivative. */
+  float3 N = sd->N;
+  object_inverse_normal_transform(kg, sd, &N);
+  return dual3(N);
+}
+
 /* Texture Coordinate Node */
 
-ccl_device_inline float3 svm_texco_reflection(const ccl_private ShaderData *sd)
+template<typename Float3Type>
+ccl_device_inline Float3Type svm_texco_reflection(const ccl_private ShaderData *sd)
 {
-  float3 data = sd->wi;
+  Float3Type data = shading_incoming<Float3Type>(sd);
   if (sd->object != OBJECT_NONE) {
     data = -reflect(data, sd->N);
   }
   return data;
 }
 
-ccl_device_inline float3 svm_texco_camera(KernelGlobals kg,
-                                          const ccl_private ShaderData *sd,
-                                          const ccl_private float3 &P)
+template<typename Float3Type>
+ccl_device_inline Float3Type svm_texco_camera(KernelGlobals kg,
+                                              const ccl_private ShaderData *sd,
+                                              const ccl_private Float3Type &P)
 {
-  float3 data = P;
+  Float3Type data(P);
   const Transform tfm = kernel_data.cam.worldtocamera;
   if (sd->object == OBJECT_NONE) {
-    data += camera_position(kg);
+    data = data + camera_position(kg);
   }
   data = transform_point(&tfm, data);
   return data;
 }
 
+template<typename Float3Type>
+ccl_device_noinline Float3Type svm_node_tex_coord_eval(KernelGlobals kg,
+                                                       ccl_private ShaderData *sd,
+                                                       const uint32_t path_flag,
+                                                       const uint type,
+                                                       ccl_private int *offset)
+{
+  Float3Type data;
+
+  switch ((NodeTexCoord)type) {
+    case NODE_TEXCO_OBJECT:
+    case NODE_TEXCO_OBJECT_WITH_TRANSFORM: {
+      data = shading_position<Float3Type>(sd);
+      if (type == NODE_TEXCO_OBJECT) {
+        object_inverse_position_transform_if_object(kg, sd, &data);
+      }
+      else {
+        Transform tfm;
+        tfm.x = read_node_float(kg, offset);
+        tfm.y = read_node_float(kg, offset);
+        tfm.z = read_node_float(kg, offset);
+        data = transform_point(&tfm, data);
+      }
+      break;
+    }
+    case NODE_TEXCO_NORMAL: {
+      if constexpr (is_dual_v<Float3Type>) {
+        data = svm_texco_smooth_normal(kg, sd);
+      }
+      else {
+        data = sd->N;
+        object_inverse_normal_transform(kg, sd, &data);
+      }
+      break;
+    }
+    case NODE_TEXCO_CAMERA: {
+      const Float3Type P = shading_position<Float3Type>(sd);
+      data = svm_texco_camera<Float3Type>(kg, sd, P);
+      break;
+    }
+    case NODE_TEXCO_WINDOW: {
+      if ((path_flag & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE &&
+          kernel_data.cam.type == CAMERA_ORTHOGRAPHIC)
+      {
+        data = Float3Type(camera_world_to_ndc(kg, sd, sd->ray_P));
+      }
+      else {
+        data = Float3Type(camera_world_to_ndc(kg, sd, sd->P));
+        if constexpr (is_dual_v<Float3Type>) {
+          data.dx.x = 1.0f / kernel_data.cam.width;
+          data.dy.y = 1.0f / kernel_data.cam.height;
+        }
+      }
+      if constexpr (is_dual_v<Float3Type>) {
+        data.val.z = 0.0f;
+      }
+      else {
+        data.z = 0.0f;
+      }
+      break;
+    }
+    case NODE_TEXCO_REFLECTION: {
+      data = svm_texco_reflection<Float3Type>(sd);
+      break;
+    }
+    case NODE_TEXCO_DUPLI_GENERATED: {
+      data = Float3Type(object_dupli_generated(kg, sd->object));
+      break;
+    }
+    case NODE_TEXCO_DUPLI_UV: {
+      data = Float3Type(object_dupli_uv(kg, sd->object));
+      break;
+    }
+    case NODE_TEXCO_VOLUME_GENERATED: {
+      data = shading_position<Float3Type>(sd);
+
+#ifdef __VOLUME__
+      if (sd->object != OBJECT_NONE) {
+        data = volume_normalized_position<Float3Type>(kg, sd, data);
+      }
+#endif
+      break;
+    }
+    default:
+      data = make_zero<Float3Type>();
+      break;
+  }
+
+  return data;
+}
+
 ccl_device_noinline int svm_node_tex_coord(KernelGlobals kg,
                                            ccl_private ShaderData *sd,
                                            const uint32_t path_flag,
@@ -48,277 +185,42 @@ ccl_device_noinline int svm_node_tex_coord(KernelGlobals kg,
                                            const uint4 node,
                                            int offset)
 {
-  float3 data = zero_float3();
-  const uint type = node.y;
-  const uint out_offset = node.z;
-
-  switch ((NodeTexCoord)type) {
-    case NODE_TEXCO_OBJECT:
-    case NODE_TEXCO_OBJECT_WITH_TRANSFORM: {
-      data = sd->P;
-      if (type == NODE_TEXCO_OBJECT) {
-        if (sd->object != OBJECT_NONE) {
-          object_inverse_position_transform(kg, sd, &data);
-        }
-      }
-      else {
-        Transform tfm;
-        tfm.x = read_node_float(kg, &offset);
-        tfm.y = read_node_float(kg, &offset);
-        tfm.z = read_node_float(kg, &offset);
-        data = transform_point(&tfm, data);
-      }
-      break;
-    }
-    case NODE_TEXCO_NORMAL: {
-      data = sd->N;
-      object_inverse_normal_transform(kg, sd, &data);
-      break;
-    }
-    case NODE_TEXCO_CAMERA: {
-      const float3 P = sd->P;
-      data = svm_texco_camera(kg, sd, P);
-      break;
-    }
-    case NODE_TEXCO_WINDOW: {
-      if ((path_flag & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE &&
-          kernel_data.cam.type == CAMERA_ORTHOGRAPHIC)
-      {
-        data = camera_world_to_ndc(kg, sd, sd->ray_P);
-      }
-      else {
-        data = camera_world_to_ndc(kg, sd, sd->P);
-      }
-      data.z = 0.0f;
-      break;
-    }
-    case NODE_TEXCO_REFLECTION: {
-      data = svm_texco_reflection(sd);
-      break;
-    }
-    case NODE_TEXCO_DUPLI_GENERATED: {
-      data = object_dupli_generated(kg, sd->object);
-      break;
-    }
-    case NODE_TEXCO_DUPLI_UV: {
-      data = object_dupli_uv(kg, sd->object);
-      break;
-    }
-    case NODE_TEXCO_VOLUME_GENERATED: {
-      data = sd->P;
-
-#ifdef __VOLUME__
-      if (sd->object != OBJECT_NONE) {
-        data = volume_normalized_position(kg, sd, data);
-      }
-#endif
-      break;
-    }
-  }
-
-  stack_store_float3(stack, out_offset, data);
+  uint type, unused1, unused2;
+  svm_unpack_node_uchar3(node.y, &type, &unused1, &unused2);
+  const float3 data = svm_node_tex_coord_eval<float3>(kg, sd, path_flag, type, &offset);
+  stack_store(stack, node.z, data);
   return offset;
 }
 
-ccl_device_inline float3 texco_normal_from_uv(KernelGlobals kg,
-                                              ccl_private ShaderData *sd,
-                                              const float u,
-                                              const float v)
+ccl_device_noinline int svm_node_tex_coord_derivative(KernelGlobals kg,
+                                                      ccl_private ShaderData *sd,
+                                                      const uint32_t path_flag,
+                                                      ccl_private float *stack,
+                                                      const uint4 node,
+                                                      int offset)
 {
-  float3 N;
-  if ((sd->type & PRIMITIVE_TRIANGLE) && (sd->shader & SHADER_SMOOTH_NORMAL)) {
-    N = (sd->type == PRIMITIVE_TRIANGLE) ?
-            triangle_smooth_normal(
-                kg, zero_float3(), sd->object, sd->object_flag, sd->prim, u, v) :
-            motion_triangle_smooth_normal(kg, zero_float3(), sd->object, sd->prim, u, v, sd->time);
-    if (is_zero(N)) {
-      N = sd->Ng;
-      object_inverse_normal_transform(kg, sd, &N);
-    }
-    else {
-      if (sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED) {
-        /* Transform to local space. */
-        object_inverse_normal_transform(kg, sd, &N);
-      }
-      if (sd->flag & SD_BACKFACING) {
-        N = -N;
-      }
-    }
+  uint type, bump_offset, store_derivatives;
+  svm_unpack_node_uchar3(node.y, &type, &bump_offset, &store_derivatives);
+
+  dual3 data = svm_node_tex_coord_eval<dual3>(kg, sd, path_flag, type, &offset);
+  const float bump_filter_width = __uint_as_float(node.w);
+  if (bump_offset == NODE_BUMP_OFFSET_DX) {
+    data.val += data.dx * bump_filter_width;
+  }
+  else if (bump_offset == NODE_BUMP_OFFSET_DY) {
+    data.val += data.dy * bump_filter_width;
+  }
+  /* Normal texture coordinate must be normalized after bump offset, matching OSL. */
+  if (type == NODE_TEXCO_NORMAL) {
+    data = safe_normalize(data);
+  }
+  if (store_derivatives) {
+    stack_store(stack, node.z, data);
   }
   else {
-    /* TODO: implement for curve. */
-    N = sd->N;
-    object_inverse_normal_transform(kg, sd, &N);
+    stack_store(stack, node.z, data.val);
   }
-  return N;
-}
-
-ccl_device_noinline int svm_node_tex_coord_bump_dx(KernelGlobals kg,
-                                                   ccl_private ShaderData *sd,
-                                                   const uint32_t path_flag,
-                                                   ccl_private float *stack,
-                                                   const uint4 node,
-                                                   int offset)
-{
-#ifdef __RAY_DIFFERENTIALS__
-  float3 data = zero_float3();
-  const uint type = node.y;
-  const uint out_offset = node.z;
-  const float bump_filter_width = __uint_as_float(node.w);
-
-  switch ((NodeTexCoord)type) {
-    case NODE_TEXCO_OBJECT:
-    case NODE_TEXCO_OBJECT_WITH_TRANSFORM: {
-      data = svm_node_bump_P_dx(sd, bump_filter_width);
-      if (type == NODE_TEXCO_OBJECT) {
-        if (sd->object != OBJECT_NONE) {
-          object_inverse_position_transform(kg, sd, &data);
-        }
-      }
-      else {
-        Transform tfm;
-        tfm.x = read_node_float(kg, &offset);
-        tfm.y = read_node_float(kg, &offset);
-        tfm.z = read_node_float(kg, &offset);
-        data = transform_point(&tfm, data);
-      }
-      break;
-    }
-    case NODE_TEXCO_NORMAL: {
-      data = texco_normal_from_uv(
-          kg, sd, sd->u + sd->du.dx * bump_filter_width, sd->v + sd->dv.dx * bump_filter_width);
-      break;
-    }
-    case NODE_TEXCO_CAMERA: {
-      const float3 P = svm_node_bump_P_dx(sd, bump_filter_width);
-      data = svm_texco_camera(kg, sd, P);
-      break;
-    }
-    case NODE_TEXCO_WINDOW: {
-      if ((path_flag & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE &&
-          kernel_data.cam.type == CAMERA_ORTHOGRAPHIC)
-      {
-        data = camera_world_to_ndc(kg, sd, sd->ray_P);
-      }
-      else {
-        data = camera_world_to_ndc(kg, sd, svm_node_bump_P_dx(sd, bump_filter_width));
-      }
-      data.z = 0.0f;
-      break;
-    }
-    case NODE_TEXCO_REFLECTION: {
-      data = svm_texco_reflection(sd);
-      break;
-    }
-    case NODE_TEXCO_DUPLI_GENERATED: {
-      data = object_dupli_generated(kg, sd->object);
-      break;
-    }
-    case NODE_TEXCO_DUPLI_UV: {
-      data = object_dupli_uv(kg, sd->object);
-      break;
-    }
-    case NODE_TEXCO_VOLUME_GENERATED: {
-      data = svm_node_bump_P_dx(sd, bump_filter_width);
-
-#  ifdef __VOLUME__
-      if (sd->object != OBJECT_NONE) {
-        data = volume_normalized_position(kg, sd, data);
-      }
-#  endif
-      break;
-    }
-  }
-
-  stack_store_float3(stack, out_offset, data);
   return offset;
-#else
-  return svm_node_tex_coord(kg, sd, path_flag, stack, node, offset);
-#endif
-}
-
-ccl_device_noinline int svm_node_tex_coord_bump_dy(KernelGlobals kg,
-                                                   ccl_private ShaderData *sd,
-                                                   const uint32_t path_flag,
-                                                   ccl_private float *stack,
-                                                   const uint4 node,
-                                                   int offset)
-{
-#ifdef __RAY_DIFFERENTIALS__
-  float3 data = zero_float3();
-  const uint type = node.y;
-  const uint out_offset = node.z;
-  const float bump_filter_width = __uint_as_float(node.w);
-
-  switch ((NodeTexCoord)type) {
-    case NODE_TEXCO_OBJECT:
-    case NODE_TEXCO_OBJECT_WITH_TRANSFORM: {
-      data = svm_node_bump_P_dy(sd, bump_filter_width);
-      if (type == NODE_TEXCO_OBJECT) {
-        if (sd->object != OBJECT_NONE) {
-          object_inverse_position_transform(kg, sd, &data);
-        }
-      }
-      else {
-        Transform tfm;
-        tfm.x = read_node_float(kg, &offset);
-        tfm.y = read_node_float(kg, &offset);
-        tfm.z = read_node_float(kg, &offset);
-        data = transform_point(&tfm, data);
-      }
-      break;
-    }
-    case NODE_TEXCO_NORMAL: {
-      data = texco_normal_from_uv(
-          kg, sd, sd->u + sd->du.dy * bump_filter_width, sd->v + sd->dv.dy * bump_filter_width);
-      break;
-    }
-    case NODE_TEXCO_CAMERA: {
-      const float3 P = svm_node_bump_P_dy(sd, bump_filter_width);
-      data = svm_texco_camera(kg, sd, P);
-      break;
-    }
-    case NODE_TEXCO_WINDOW: {
-      if ((path_flag & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE &&
-          kernel_data.cam.type == CAMERA_ORTHOGRAPHIC)
-      {
-        data = camera_world_to_ndc(kg, sd, sd->ray_P);
-      }
-      else {
-        data = camera_world_to_ndc(kg, sd, svm_node_bump_P_dy(sd, bump_filter_width));
-      }
-      data.z = 0.0f;
-      break;
-    }
-    case NODE_TEXCO_REFLECTION: {
-      data = svm_texco_reflection(sd);
-      break;
-    }
-    case NODE_TEXCO_DUPLI_GENERATED: {
-      data = object_dupli_generated(kg, sd->object);
-      break;
-    }
-    case NODE_TEXCO_DUPLI_UV: {
-      data = object_dupli_uv(kg, sd->object);
-      break;
-    }
-    case NODE_TEXCO_VOLUME_GENERATED: {
-      data = svm_node_bump_P_dy(sd, bump_filter_width);
-
-#  ifdef __VOLUME__
-      if (sd->object != OBJECT_NONE) {
-        data = volume_normalized_position(kg, sd, data);
-      }
-#  endif
-      break;
-    }
-  }
-
-  stack_store_float3(stack, out_offset, data);
-  return offset;
-#else
-  return svm_node_tex_coord(kg, sd, path_flag, stack, node, offset);
-#endif
 }
 
 ccl_device_noinline void svm_node_normal_map(KernelGlobals kg,
@@ -367,8 +269,8 @@ ccl_device_noinline void svm_node_normal_map(KernelGlobals kg,
     }
 
     /* get _unnormalized_ interpolated normal and tangent */
-    const float3 tangent = primitive_surface_attribute<float3>(kg, sd, attr).val;
-    const float sign = primitive_surface_attribute<float>(kg, sd, attr_sign).val;
+    const float3 tangent = primitive_surface_attribute<float3>(kg, sd, attr);
+    const float sign = primitive_surface_attribute<float>(kg, sd, attr_sign);
     float3 normal;
 
     if (sd->shader & SHADER_SMOOTH_NORMAL) {
@@ -377,8 +279,7 @@ ccl_device_noinline void svm_node_normal_map(KernelGlobals kg,
               find_attribute(kg, sd->object, sd->prim, ATTR_STD_NORMAL_UNDISPLACED) :
               AttributeDescriptor{ATTR_ELEMENT_NONE, NODE_ATTR_FLOAT3, ATTR_STD_NOT_FOUND};
       if (attr_undisplaced_normal.offset != ATTR_STD_NOT_FOUND) {
-        normal =
-            primitive_surface_attribute<float3>(kg, sd, attr_undisplaced_normal, false, false).val;
+        normal = primitive_surface_attribute<float3>(kg, sd, attr_undisplaced_normal);
         /* Can't interpolate in tangent space as the displaced normal is not used
          * for the tangent frame. */
         linear_interpolate_strength = true;
@@ -454,6 +355,7 @@ ccl_device_noinline void svm_node_normal_map(KernelGlobals kg,
   stack_store_float3(stack, normal_offset, N);
 }
 
+template<typename Float3Type>
 ccl_device_noinline void svm_node_tangent(KernelGlobals kg,
                                           ccl_private ShaderData *sd,
                                           ccl_private float *stack,
@@ -464,54 +366,73 @@ ccl_device_noinline void svm_node_tangent(KernelGlobals kg,
   uint axis;
   svm_unpack_node_uchar3(node.y, &tangent_offset, &direction_type, &axis);
 
-  float3 tangent;
-  float3 attribute_value;
   const AttributeDescriptor desc = find_attribute(kg, sd, node.z);
-  if (desc.offset != ATTR_STD_NOT_FOUND) {
-    if (desc.type == NODE_ATTR_FLOAT2) {
-      const float2 value = primitive_surface_attribute<float2>(kg, sd, desc).val;
-      attribute_value.x = value.x;
-      attribute_value.y = value.y;
-      attribute_value.z = 0.0f;
-    }
-    else {
-      attribute_value = primitive_surface_attribute<float3>(kg, sd, desc).val;
-    }
-  }
 
+  Float3Type tangent;
   if (direction_type == NODE_TANGENT_UVMAP) {
     /* UV map */
     if (desc.offset == ATTR_STD_NOT_FOUND) {
-      stack_store_float3(stack, tangent_offset, zero_float3());
+      stack_store(stack, tangent_offset, Float3Type());
       return;
     }
-    tangent = attribute_value;
+    if (desc.type == NODE_ATTR_FLOAT2) {
+      if constexpr (is_dual_v<Float3Type>) {
+        tangent = make_float3(primitive_surface_attribute<dual2>(kg, sd, desc));
+      }
+      else {
+        tangent = make_float3(primitive_surface_attribute<float2>(kg, sd, desc));
+      }
+    }
+    else {
+      tangent = primitive_surface_attribute<Float3Type>(kg, sd, desc);
+    }
   }
   else {
     /* radial */
-    float3 generated;
-
+    Float3Type generated;
     if (desc.offset == ATTR_STD_NOT_FOUND) {
-      generated = sd->P;
+      generated = shading_position<Float3Type>(sd);
+    }
+    else if (desc.type == NODE_ATTR_FLOAT2) {
+      if constexpr (is_dual_v<Float3Type>) {
+        generated = make_float3(primitive_surface_attribute<dual2>(kg, sd, desc));
+      }
+      else {
+        generated = make_float3(primitive_surface_attribute<float2>(kg, sd, desc));
+      }
     }
     else {
-      generated = attribute_value;
+      generated = primitive_surface_attribute<Float3Type>(kg, sd, desc);
     }
 
-    if (axis == NODE_TANGENT_AXIS_X) {
-      tangent = make_float3(0.0f, -(generated.z - 0.5f), (generated.y - 0.5f));
-    }
-    else if (axis == NODE_TANGENT_AXIS_Y) {
-      tangent = make_float3(-(generated.z - 0.5f), 0.0f, (generated.x - 0.5f));
+    if constexpr (is_dual_v<Float3Type>) {
+      using FloatType = dual_scalar_t<Float3Type>;
+      if (axis == NODE_TANGENT_AXIS_X) {
+        tangent = make_float3(FloatType(), -(generated.z() - 0.5f), (generated.y() - 0.5f));
+      }
+      else if (axis == NODE_TANGENT_AXIS_Y) {
+        tangent = make_float3(-(generated.z() - 0.5f), FloatType(), (generated.x() - 0.5f));
+      }
+      else {
+        tangent = make_float3(-(generated.y() - 0.5f), (generated.x() - 0.5f), FloatType());
+      }
     }
     else {
-      tangent = make_float3(-(generated.y - 0.5f), (generated.x - 0.5f), 0.0f);
+      if (axis == NODE_TANGENT_AXIS_X) {
+        tangent = make_float3(0.0f, -(generated.z - 0.5f), (generated.y - 0.5f));
+      }
+      else if (axis == NODE_TANGENT_AXIS_Y) {
+        tangent = make_float3(-(generated.z - 0.5f), 0.0f, (generated.x - 0.5f));
+      }
+      else {
+        tangent = make_float3(-(generated.y - 0.5f), (generated.x - 0.5f), 0.0f);
+      }
     }
   }
 
   object_normal_transform(kg, sd, &tangent);
   tangent = cross(sd->N, normalize(cross(tangent, sd->N)));
-  stack_store_float3(stack, tangent_offset, tangent);
+  stack_store(stack, tangent_offset, tangent);
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/svm/types.h

@@ -21,6 +21,7 @@ CCL_NAMESPACE_BEGIN
 
 enum ShaderNodeType {
 #define SHADER_NODE_TYPE(name) name,
+#define SHADER_NODE_TYPE_DERIVATIVE(name) name, name##_DERIVATIVE,
 #include "node_types_template.h"
 
   NODE_NUM
@@ -273,7 +274,8 @@ enum NodeConvert {
   NODE_CONVERT_VF,
   NODE_CONVERT_VI,
   NODE_CONVERT_IF,
-  NODE_CONVERT_IV
+  NODE_CONVERT_IV,
+  NODE_CONVERT_NONE,
 };
 
 enum NodeNoiseType {

intern/cycles/kernel/svm/util.h

@@ -62,6 +62,61 @@ ccl_device_inline void stack_store_float(ccl_private float *stack, const uint a,
   stack[a] = f;
 }
 
+/* Type-based stack load. T can be float, float3, dual1, or dual3.
+ * When T is a dual type, derivatives are loaded from adjacent stack slots. */
+
+template<typename T> ccl_device_inline T stack_load(const ccl_private float *stack, const uint a);
+
+ccl_device_template_spec float stack_load(const ccl_private float *stack, const uint a)
+{
+  return stack_load_float(stack, a);
+}
+
+ccl_device_template_spec float3 stack_load(const ccl_private float *stack, const uint a)
+{
+  return stack_load_float3(stack, a);
+}
+
+ccl_device_template_spec dual1 stack_load(const ccl_private float *stack, const uint a)
+{
+  return {
+      stack_load_float(stack, a), stack_load_float(stack, a + 1), stack_load_float(stack, a + 2)};
+}
+
+ccl_device_template_spec dual3 stack_load(const ccl_private float *stack, const uint a)
+{
+  return {stack_load_float3(stack, a),
+          stack_load_float3(stack, a + 3),
+          stack_load_float3(stack, a + 6)};
+}
+
+/* Type-based stack store. Overloaded for plain and dual types.
+ * For dual types, derivatives are stored in adjacent stack slots. */
+
+ccl_device_inline void stack_store(ccl_private float *stack, const uint a, const float f)
+{
+  stack_store_float(stack, a, f);
+}
+
+ccl_device_inline void stack_store(ccl_private float *stack, const uint a, const float3 f)
+{
+  stack_store_float3(stack, a, f);
+}
+
+ccl_device_inline void stack_store(ccl_private float *stack, const uint a, const dual1 f)
+{
+  stack_store_float(stack, a, f.val);
+  stack_store_float(stack, a + 1, f.dx);
+  stack_store_float(stack, a + 2, f.dy);
+}
+
+ccl_device_inline void stack_store(ccl_private float *stack, const uint a, const dual3 f)
+{
+  stack_store_float3(stack, a, f.val);
+  stack_store_float3(stack, a + 3, f.dx);
+  stack_store_float3(stack, a + 6, f.dy);
+}
+
 ccl_device_inline int stack_load_int(const ccl_private float *stack, const uint a)
 {
   kernel_assert(a < SVM_STACK_SIZE);
@@ -157,4 +212,38 @@ ccl_device_forceinline float3 dPdy(const ccl_private ShaderData *sd)
   return sd->dPdu * sd->du.dy + sd->dPdv * sd->dv.dy;
 }
 
+/* Shading position, returns Float3Type = float3 (no derivatives) or dual3 (with derivatives). */
+
+template<typename Float3Type>
+ccl_device_inline Float3Type shading_position(const ccl_private ShaderData *sd)
+{
+  if constexpr (is_dual_v<Float3Type>) {
+    dual3 P(sd->P);
+    P.dx = dPdx(sd);
+    P.dy = dPdy(sd);
+    return P;
+  }
+  else {
+    return sd->P;
+  }
+}
+
+/* Shading incoming direction, returns Float3Type = float3 or dual3. */
+
+template<typename Float3Type>
+ccl_device_inline Float3Type shading_incoming(const ccl_private ShaderData *sd)
+{
+  if constexpr (is_dual_v<Float3Type>) {
+    dual3 I(sd->wi);
+    float3 dIdx, dIdy;
+    make_orthonormals(sd->wi, &dIdx, &dIdy);
+    I.dx = sd->dI * dIdx;
+    I.dy = sd->dI * dIdy;
+    return I;
+  }
+  else {
+    return sd->wi;
+  }
+}
+
 CCL_NAMESPACE_END

intern/cycles/kernel/svm/value.h

@@ -10,13 +10,16 @@ CCL_NAMESPACE_BEGIN
 
 /* Value Nodes */
 
+template<typename FloatType>
 ccl_device void svm_node_value_f(ccl_private float *stack,
                                  const uint ivalue,
                                  const uint out_offset)
 {
-  stack_store_float(stack, out_offset, __uint_as_float(ivalue));
+  /* Derivative of a constant is zero. */
+  stack_store(stack, out_offset, FloatType(__uint_as_float(ivalue)));
 }
 
+template<typename Float3Type>
 ccl_device int svm_node_value_v(KernelGlobals kg,
                                 ccl_private float *stack,
                                 const uint out_offset,
@@ -27,7 +30,8 @@ ccl_device int svm_node_value_v(KernelGlobals kg,
   const float3 p = make_float3(
       __uint_as_float(node1.y), __uint_as_float(node1.z), __uint_as_float(node1.w));
 
-  stack_store_float3(stack, out_offset, p);
+  /* Derivative of a constant is zero. */
+  stack_store(stack, out_offset, Float3Type(p));
   return offset;
 }
 

intern/cycles/kernel/svm/vertex_color.h

@@ -24,12 +24,12 @@ ccl_device_noinline void svm_node_vertex_color(KernelGlobals kg,
   const AttributeDescriptor descriptor = find_attribute(kg, sd, layer_id);
   if (descriptor.offset != ATTR_STD_NOT_FOUND) {
     if (descriptor.type == NODE_ATTR_FLOAT4 || descriptor.type == NODE_ATTR_RGBA) {
-      const float4 vertex_color = primitive_surface_attribute<float4>(kg, sd, descriptor).val;
+      const float4 vertex_color = primitive_surface_attribute<float4>(kg, sd, descriptor);
       stack_store_float3(stack, color_offset, make_float3(vertex_color));
       stack_store_float(stack, alpha_offset, vertex_color.w);
     }
     else {
-      const float3 vertex_color = primitive_surface_attribute<float3>(kg, sd, descriptor).val;
+      const float3 vertex_color = primitive_surface_attribute<float3>(kg, sd, descriptor);
       stack_store_float3(stack, color_offset, vertex_color);
       stack_store_float(stack, alpha_offset, 1.0f);
     }
@@ -40,60 +40,39 @@ ccl_device_noinline void svm_node_vertex_color(KernelGlobals kg,
   }
 }
 
-ccl_device_noinline void svm_node_vertex_color_bump_dx(KernelGlobals kg,
-                                                       ccl_private ShaderData *sd,
-                                                       ccl_private float *stack,
-                                                       const uint4 node)
+ccl_device_noinline void svm_node_vertex_color_derivative(KernelGlobals kg,
+                                                          ccl_private ShaderData *sd,
+                                                          ccl_private float *stack,
+                                                          const uint4 node)
 {
   uint layer_id;
   uint color_offset;
   uint alpha_offset;
-  svm_unpack_node_uchar3(node.y, &layer_id, &color_offset, &alpha_offset);
+  uint bump_offset;
+  svm_unpack_node_uchar4(node.y, &layer_id, &color_offset, &alpha_offset, &bump_offset);
   const float bump_filter_width = __uint_as_float(node.z);
 
   const AttributeDescriptor descriptor = find_attribute(kg, sd, layer_id);
   if (descriptor.offset != ATTR_STD_NOT_FOUND) {
     if (descriptor.type == NODE_ATTR_FLOAT4 || descriptor.type == NODE_ATTR_RGBA) {
-      dual4 vertex_color = primitive_surface_attribute<float4>(kg, sd, descriptor, true, false);
-      vertex_color.val += vertex_color.dx * bump_filter_width;
+      dual4 vertex_color = primitive_surface_attribute<dual4>(kg, sd, descriptor);
+      if (bump_offset == NODE_BUMP_OFFSET_DX) {
+        vertex_color.val += vertex_color.dx * bump_filter_width;
+      }
+      else if (bump_offset == NODE_BUMP_OFFSET_DY) {
+        vertex_color.val += vertex_color.dy * bump_filter_width;
+      }
       stack_store_float3(stack, color_offset, make_float3(vertex_color.val));
       stack_store_float(stack, alpha_offset, vertex_color.val.w);
     }
     else {
-      dual3 vertex_color = primitive_surface_attribute<float3>(kg, sd, descriptor, true, false);
-      vertex_color.val += vertex_color.dx * bump_filter_width;
-      stack_store_float3(stack, color_offset, vertex_color.val);
-      stack_store_float(stack, alpha_offset, 1.0f);
-    }
-  }
-  else {
-    stack_store_float3(stack, color_offset, make_float3(0.0f, 0.0f, 0.0f));
-    stack_store_float(stack, alpha_offset, 0.0f);
-  }
-}
-
-ccl_device_noinline void svm_node_vertex_color_bump_dy(KernelGlobals kg,
-                                                       ccl_private ShaderData *sd,
-                                                       ccl_private float *stack,
-                                                       const uint4 node)
-{
-  uint layer_id;
-  uint color_offset;
-  uint alpha_offset;
-  svm_unpack_node_uchar3(node.y, &layer_id, &color_offset, &alpha_offset);
-  const float bump_filter_width = __uint_as_float(node.z);
-
-  const AttributeDescriptor descriptor = find_attribute(kg, sd, layer_id);
-  if (descriptor.offset != ATTR_STD_NOT_FOUND) {
-    if (descriptor.type == NODE_ATTR_FLOAT4 || descriptor.type == NODE_ATTR_RGBA) {
-      dual4 vertex_color = primitive_surface_attribute<float4>(kg, sd, descriptor, false, true);
-      vertex_color.val += vertex_color.dy * bump_filter_width;
-      stack_store_float3(stack, color_offset, make_float3(vertex_color.val));
-      stack_store_float(stack, alpha_offset, vertex_color.val.w);
-    }
-    else {
-      dual3 vertex_color = primitive_surface_attribute<float3>(kg, sd, descriptor, false, true);
-      vertex_color.val += vertex_color.dy * bump_filter_width;
+      dual3 vertex_color = primitive_surface_attribute<dual3>(kg, sd, descriptor);
+      if (bump_offset == NODE_BUMP_OFFSET_DX) {
+        vertex_color.val += vertex_color.dx * bump_filter_width;
+      }
+      else if (bump_offset == NODE_BUMP_OFFSET_DY) {
+        vertex_color.val += vertex_color.dy * bump_filter_width;
+      }
       stack_store_float3(stack, color_offset, vertex_color.val);
       stack_store_float(stack, alpha_offset, 1.0f);
     }

intern/cycles/kernel/util/colorspace.h

@@ -31,7 +31,7 @@ ccl_device float3 rec709_to_rgb(KernelGlobals kg, const float3 rec709)
                          dot(make_float3(kernel_data.film.rec709_to_b), rec709));
 }
 
-ccl_device float linear_rgb_to_gray(KernelGlobals kg, const float3 c)
+template<class T> ccl_device auto linear_rgb_to_gray(KernelGlobals kg, const T c)
 {
   return dot(c, make_float3(kernel_data.film.rgb_to_y));
 }

intern/cycles/scene/image.h

@@ -176,6 +176,8 @@ class ImageManager {
 
   bool need_update() const;
 
+  bool get_use_texture_cache() const;
+
  private:
   bool need_update_;
 

intern/cycles/scene/shader_graph.h

@@ -7,6 +7,7 @@
 #include "graph/node.h"
 #include "graph/node_type.h"
 
+#include "kernel/svm/types.h"
 #include "kernel/types.h"
 
 #include "util/map.h"
@@ -59,15 +60,11 @@ enum ShaderNodeSpecialType {
   SHADER_SPECIAL_TYPE_LIGHT_PATH,
 };
 
-/* Input
- *
- * Input socket for a shader node. May be linked to an output or not. If not
- * linked, it will either get a fixed default value, or e.g. a texture
- * coordinate. */
+/* Base class for ShaderInput and ShaderOutput. */
 
-class ShaderInput {
+class ShaderIO {
  public:
-  ShaderInput(const SocketType &socket_type_, ShaderNode *parent_)
+  ShaderIO(const SocketType &socket_type_, ShaderNode *parent_)
       : socket_type(socket_type_), parent(parent_)
 
   {
@@ -77,15 +74,33 @@ class ShaderInput {
   {
     return socket_type.ui_name;
   }
-  int flags() const
-  {
-    return socket_type.flags;
-  }
   SocketType::Type type() const
   {
     return socket_type.type;
   }
 
+  const SocketType &socket_type;
+  ShaderNode *parent;
+
+  virtual void disconnect() {};
+
+  int stack_offset = SVM_STACK_INVALID; /* for SVM compiler */
+};
+
+/* Input
+ *
+ * Input socket for a shader node. May be linked to an output or not. If not
+ * linked, it will either get a fixed default value, or e.g. a texture
+ * coordinate. */
+
+class ShaderInput final : public ShaderIO {
+ public:
+  using ShaderIO::ShaderIO;
+
+  int flags() const
+  {
+    return socket_type.flags;
+  }
   void set(const float f)
   {
     ((Node *)parent)->set(socket_type, f);
@@ -99,12 +114,9 @@ class ShaderInput {
     ((Node *)parent)->set(socket_type, f);
   }
 
-  void disconnect();
+  void disconnect() override;
 
-  const SocketType &socket_type;
-  ShaderNode *parent;
   ShaderOutput *link = nullptr;
-  int stack_offset = SVM_STACK_INVALID; /* for SVM compiler */
 
   /* Keeps track of whether a constant was folded in this socket, to avoid over-optimizing when the
    * link is null. */
@@ -115,28 +127,13 @@ class ShaderInput {
  *
  * Output socket for a shader node. */
 
-class ShaderOutput {
+class ShaderOutput final : public ShaderIO {
  public:
-  ShaderOutput(const SocketType &socket_type_, ShaderNode *parent_)
-      : socket_type(socket_type_), parent(parent_)
-  {
-  }
+  using ShaderIO::ShaderIO;
 
-  ustring name() const
-  {
-    return socket_type.ui_name;
-  }
-  SocketType::Type type() const
-  {
-    return socket_type.type;
-  }
+  void disconnect() override;
 
-  void disconnect();
-
-  const SocketType &socket_type;
-  ShaderNode *parent;
   vector<ShaderInput *> links;
-  int stack_offset = SVM_STACK_INVALID; /* for SVM compiler */
 };
 
 /* Node
@@ -211,6 +208,22 @@ class ShaderNode : public Node {
   {
     return false;
   }
+  virtual ShaderNodeType shader_node_type() const
+  {
+    return NODE_NONE;
+  }
+  virtual bool is_texture_node_and_needs_derivatives(const SVMCompiler & /*compiler*/)
+  {
+    return false;
+  }
+  bool need_derivatives() const
+  {
+    return need_derivatives_;
+  }
+  void set_need_derivatives(const bool need_derivatives = true)
+  {
+    need_derivatives_ = need_derivatives;
+  }
 
   unique_ptr_vector<ShaderInput> inputs;
   unique_ptr_vector<ShaderOutput> outputs;
@@ -256,6 +269,8 @@ class ShaderNode : public Node {
   virtual bool equals(const ShaderNode &other);
 
  protected:
+  bool need_derivatives_ = false;
+
   /* Disconnect the input with the given name if it is connected.
    * Used to optimize away unused inputs. */
   void disconnect_unused_input(const char *name);

intern/cycles/scene/shader_nodes.cpp

@@ -185,7 +185,8 @@ int TextureMapping::compile_begin(SVMCompiler &compiler, ShaderInput *vector_in)
 {
   if (!skip()) {
     const int offset_in = compiler.stack_assign(vector_in);
-    const int offset_out = compiler.stack_find_offset(SocketType::VECTOR);
+    assert(vector_in->type() == SocketType::VECTOR || vector_in->type() == SocketType::POINT);
+    const int offset_out = compiler.stack_find_offset(vector_in);
 
     compile(compiler, offset_in, offset_out);
 
@@ -200,7 +201,7 @@ void TextureMapping::compile_end(SVMCompiler &compiler,
                                  const int vector_offset)
 {
   if (!skip()) {
-    compiler.stack_clear_offset(vector_in->type(), vector_offset);
+    compiler.stack_clear_offset(vector_in, vector_offset);
   }
 }
 
@@ -368,11 +369,19 @@ void ImageTextureNode::attributes(Shader *shader, AttributeRequestSet *attribute
   ShaderNode::attributes(shader, attributes);
 }
 
+ShaderNodeType ImageTextureNode::shader_node_type() const
+{
+  if (projection != NODE_IMAGE_PROJ_BOX) {
+    return NODE_TEX_IMAGE;
+  }
+  return NODE_TEX_IMAGE_BOX;
+}
+
 void ImageTextureNode::update_images(const SVMCompiler &compiler)
 {
   if (handle.empty()) {
-    cull_tiles(compiler.scene, compiler.current_graph);
     ImageManager *image_manager = compiler.scene->image_manager.get();
+    cull_tiles(compiler.scene, compiler.current_graph);
     handle = image_manager->add_image(filename.string(), image_params(), tiles);
   }
 }
@@ -406,7 +415,7 @@ void ImageTextureNode::compile(SVMCompiler &compiler)
   }
 
   if (projection != NODE_IMAGE_PROJ_BOX) {
-    compiler.add_node(NODE_TEX_IMAGE,
+    compiler.add_node(this,
                       handle.kernel_id(),
                       compiler.encode_uchar4(vector_offset,
                                              compiler.stack_assign_if_linked(color_out),
@@ -415,7 +424,7 @@ void ImageTextureNode::compile(SVMCompiler &compiler)
                       projection);
   }
   else {
-    compiler.add_node(NODE_TEX_IMAGE_BOX,
+    compiler.add_node(this,
                       handle.kernel_id(),
                       compiler.encode_uchar4(vector_offset,
                                              compiler.stack_assign_if_linked(color_out),
@@ -575,7 +584,7 @@ void EnvironmentTextureNode::compile(SVMCompiler &compiler)
     flags |= NODE_IMAGE_COMPRESS_AS_SRGB;
   }
 
-  compiler.add_node(NODE_TEX_ENVIRONMENT,
+  compiler.add_node(this,
                     handle.kernel_id(),
                     compiler.encode_uchar4(vector_offset,
                                            compiler.stack_assign_if_linked(color_out),
@@ -1945,7 +1954,7 @@ void MappingNode::compile(SVMCompiler &compiler)
   const int result_stack_offset = compiler.stack_assign(vector_out);
 
   compiler.add_node(
-      NODE_MAPPING,
+      this,
       mapping_type,
       compiler.encode_uchar4(
           vector_stack_offset, location_stack_offset, rotation_stack_offset, scale_stack_offset),
@@ -1982,10 +1991,8 @@ void RGBToBWNode::constant_fold(const ConstantFolder &folder)
 
 void RGBToBWNode::compile(SVMCompiler &compiler)
 {
-  compiler.add_node(NODE_CONVERT,
-                    NODE_CONVERT_CF,
-                    compiler.stack_assign(inputs[0]),
-                    compiler.stack_assign(outputs[0]));
+  compiler.add_node(
+      this, NODE_CONVERT_CF, compiler.stack_assign(inputs[0]), compiler.stack_assign(outputs[0]));
 }
 
 void RGBToBWNode::compile(OSLCompiler &compiler)
@@ -2135,6 +2142,43 @@ void ConvertNode::constant_fold(const ConstantFolder &folder)
   }
 }
 
+NodeConvert ConvertNode::convert_type()
+{
+  if (from == SocketType::FLOAT) {
+    if (to == SocketType::INT) {
+      /* float to int */
+      return NODE_CONVERT_FI;
+    }
+    /* float to float3 */
+    return NODE_CONVERT_FV;
+  }
+  if (from == SocketType::INT) {
+    if (to == SocketType::FLOAT) {
+      /* int to float */
+      return NODE_CONVERT_IF;
+    }
+    /* int to vector/point/normal */
+    return NODE_CONVERT_IV;
+  }
+  if (to == SocketType::FLOAT) {
+    if (from == SocketType::COLOR) {
+      /* color to float */
+      return NODE_CONVERT_CF;
+    }
+    /* vector/point/normal to float */
+    return NODE_CONVERT_VF;
+  }
+  if (to == SocketType::INT) {
+    if (from == SocketType::COLOR) {
+      /* color to int */
+      return NODE_CONVERT_CI;
+    }
+    /* vector/point/normal to int */
+    return NODE_CONVERT_VI;
+  }
+  return NODE_CONVERT_NONE;
+}
+
 void ConvertNode::compile(SVMCompiler &compiler)
 {
   /* proxy nodes should have been removed at this point */
@@ -2143,53 +2187,9 @@ void ConvertNode::compile(SVMCompiler &compiler)
   ShaderInput *in = inputs[0];
   ShaderOutput *out = outputs[0];
 
-  if (from == SocketType::FLOAT) {
-    if (to == SocketType::INT) {
-      /* float to int */
-      compiler.add_node(
-          NODE_CONVERT, NODE_CONVERT_FI, compiler.stack_assign(in), compiler.stack_assign(out));
-    }
-    else {
-      /* float to float3 */
-      compiler.add_node(
-          NODE_CONVERT, NODE_CONVERT_FV, compiler.stack_assign(in), compiler.stack_assign(out));
-    }
-  }
-  else if (from == SocketType::INT) {
-    if (to == SocketType::FLOAT) {
-      /* int to float */
-      compiler.add_node(
-          NODE_CONVERT, NODE_CONVERT_IF, compiler.stack_assign(in), compiler.stack_assign(out));
-    }
-    else {
-      /* int to vector/point/normal */
-      compiler.add_node(
-          NODE_CONVERT, NODE_CONVERT_IV, compiler.stack_assign(in), compiler.stack_assign(out));
-    }
-  }
-  else if (to == SocketType::FLOAT) {
-    if (from == SocketType::COLOR) {
-      /* color to float */
-      compiler.add_node(
-          NODE_CONVERT, NODE_CONVERT_CF, compiler.stack_assign(in), compiler.stack_assign(out));
-    }
-    else {
-      /* vector/point/normal to float */
-      compiler.add_node(
-          NODE_CONVERT, NODE_CONVERT_VF, compiler.stack_assign(in), compiler.stack_assign(out));
-    }
-  }
-  else if (to == SocketType::INT) {
-    if (from == SocketType::COLOR) {
-      /* color to int */
-      compiler.add_node(
-          NODE_CONVERT, NODE_CONVERT_CI, compiler.stack_assign(in), compiler.stack_assign(out));
-    }
-    else {
-      /* vector/point/normal to int */
-      compiler.add_node(
-          NODE_CONVERT, NODE_CONVERT_VI, compiler.stack_assign(in), compiler.stack_assign(out));
-    }
+  const NodeConvert type = convert_type();
+  if (type != NODE_CONVERT_NONE) {
+    compiler.add_node(this, type, compiler.stack_assign(in), compiler.stack_assign(out));
   }
   else {
     /* float3 to float3 */
@@ -2199,8 +2199,7 @@ void ConvertNode::compile(SVMCompiler &compiler)
     }
     else {
       /* set 0,0,0 value */
-      compiler.add_node(NODE_VALUE_V, compiler.stack_assign(out));
-      compiler.add_node(NODE_VALUE_V, value_color);
+      compiler.add_value_node(this, value_color, compiler.stack_assign(out));
     }
   }
 }
@@ -2275,7 +2274,7 @@ void BsdfNode::compile(SVMCompiler &compiler,
   const int data_z_offset = (data_z) ? compiler.stack_assign(data_z) : SVM_STACK_INVALID;
   const int data_w_offset = (data_w) ? compiler.stack_assign(data_w) : SVM_STACK_INVALID;
 
-  compiler.add_node(NODE_CLOSURE_BSDF,
+  compiler.add_node(this,
                     compiler.encode_uchar4(
                         closure,
                         (bsdf_y) ? compiler.stack_assign_if_linked(bsdf_y) : SVM_STACK_INVALID,
@@ -2390,7 +2389,7 @@ void MetallicBsdfNode::compile(SVMCompiler &compiler)
   const int tangent_offset = compiler.stack_assign_if_linked(input("Tangent"));
   const int rotation_offset = compiler.stack_assign(input("Rotation"));
 
-  compiler.add_node(NODE_CLOSURE_BSDF,
+  compiler.add_node(this,
                     compiler.encode_uchar4(fresnel_type,
                                            compiler.stack_assign_if_linked(roughness_in),
                                            compiler.stack_assign_if_linked(anisotropy_in),
@@ -2947,7 +2946,7 @@ void PrincipledBsdfNode::compile(SVMCompiler &compiler)
   }
 
   compiler.add_node(
-      NODE_CLOSURE_BSDF,
+      this,
       compiler.encode_uchar4(
           closure, ior_offset, roughness_offset, compiler.closure_mix_weight_offset()),
       __float_as_int(get_float(input("IOR")->socket_type)),
@@ -3838,7 +3837,7 @@ void PrincipledHairBsdfNode::compile(SVMCompiler &compiler)
   /* Encode all parameters into data nodes. */
   /* node */
   compiler.add_node(
-      NODE_CLOSURE_BSDF,
+      this,
       /* Socket IDs can be packed 4 at a time into a single data packet */
       compiler.encode_uchar4(
           closure, roughness_ofs, random_roughness_ofs, compiler.closure_mix_weight_offset()),
@@ -3981,55 +3980,107 @@ void GeometryNode::attributes(Shader *shader, AttributeRequestSet *attributes)
   ShaderNode::attributes(shader, attributes);
 }
 
+ShaderNodeType GeometryNode::shader_node_type() const
+{
+  return NODE_GEOMETRY;
+}
+
+static uint shader_bump_to_node_bump_offset(ShaderBump bump)
+{
+  switch (bump) {
+    case SHADER_BUMP_DX:
+      return NODE_BUMP_OFFSET_DX;
+    case SHADER_BUMP_DY:
+      return NODE_BUMP_OFFSET_DY;
+    default:
+      return NODE_BUMP_OFFSET_CENTER;
+  }
+}
+
+/* Construct a temporary AttributeNode to get the type and derivative info for SVM. */
+static AttributeNode attr_node_copy_from(const ShaderNode *node)
+{
+  AttributeNode attr_node;
+  attr_node.bump = node->bump;
+  attr_node.set_need_derivatives(node->need_derivatives());
+  return attr_node;
+}
+
+/* Construct a temporary GeometryNode to get the type and derivative info for SVM. */
+static GeometryNode geom_node_copy_from(const ShaderNode *node)
+{
+  GeometryNode geom_node;
+  geom_node.bump = node->bump;
+  geom_node.set_need_derivatives(node->need_derivatives());
+  return geom_node;
+}
+
 void GeometryNode::compile(SVMCompiler &compiler)
 {
+  const uint bump_offset = shader_bump_to_node_bump_offset(bump);
+  const bool use_derivative = need_derivatives() || (bump != SHADER_BUMP_NONE);
+  const bool store_derivatives = need_derivatives();
   ShaderOutput *out;
-  ShaderNodeType geom_node = NODE_GEOMETRY;
-  ShaderNodeType attr_node = NODE_ATTR;
-
-  if (bump == SHADER_BUMP_DX) {
-    geom_node = NODE_GEOMETRY_BUMP_DX;
-    attr_node = NODE_ATTR_BUMP_DX;
-  }
-  else if (bump == SHADER_BUMP_DY) {
-    geom_node = NODE_GEOMETRY_BUMP_DY;
-    attr_node = NODE_ATTR_BUMP_DY;
-  }
 
   out = output("Position");
   if (!out->links.empty()) {
-    compiler.add_node(
-        geom_node, NODE_GEOM_P, compiler.stack_assign(out), __float_as_uint(bump_filter_width));
+    compiler.add_node_derivative(
+        NODE_GEOMETRY,
+        use_derivative,
+        compiler.encode_uchar4(NODE_GEOM_P, bump_offset, store_derivatives),
+        compiler.stack_assign(out),
+        __float_as_uint(bump_filter_width));
   }
 
+  /* Currently no bump offset is supported for Normal, Tangent, True Normal, and Incoming. */
   out = output("Normal");
   if (!out->links.empty()) {
-    compiler.add_node(
-        geom_node, NODE_GEOM_N, compiler.stack_assign(out), __float_as_uint(bump_filter_width));
+    compiler.add_node_derivative(
+        NODE_GEOMETRY,
+        use_derivative,
+        compiler.encode_uchar4(NODE_GEOM_N, NODE_BUMP_OFFSET_CENTER, store_derivatives),
+        compiler.stack_assign(out),
+        __float_as_uint(bump_filter_width));
   }
 
   out = output("Tangent");
   if (!out->links.empty()) {
-    compiler.add_node(
-        geom_node, NODE_GEOM_T, compiler.stack_assign(out), __float_as_uint(bump_filter_width));
+    compiler.add_node_derivative(
+        NODE_GEOMETRY,
+        use_derivative,
+        compiler.encode_uchar4(NODE_GEOM_T, NODE_BUMP_OFFSET_CENTER, store_derivatives),
+        compiler.stack_assign(out),
+        __float_as_uint(bump_filter_width));
   }
 
   out = output("True Normal");
   if (!out->links.empty()) {
-    compiler.add_node(
-        geom_node, NODE_GEOM_Ng, compiler.stack_assign(out), __float_as_uint(bump_filter_width));
+    compiler.add_node_derivative(
+        NODE_GEOMETRY,
+        use_derivative,
+        compiler.encode_uchar4(NODE_GEOM_Ng, NODE_BUMP_OFFSET_CENTER, store_derivatives),
+        compiler.stack_assign(out),
+        __float_as_uint(bump_filter_width));
   }
 
   out = output("Incoming");
   if (!out->links.empty()) {
-    compiler.add_node(
-        geom_node, NODE_GEOM_I, compiler.stack_assign(out), __float_as_uint(bump_filter_width));
+    compiler.add_node_derivative(
+        NODE_GEOMETRY,
+        use_derivative,
+        compiler.encode_uchar4(NODE_GEOM_I, NODE_BUMP_OFFSET_CENTER, store_derivatives),
+        compiler.stack_assign(out),
+        __float_as_uint(bump_filter_width));
   }
 
   out = output("Parametric");
   if (!out->links.empty()) {
-    compiler.add_node(
-        geom_node, NODE_GEOM_uv, compiler.stack_assign(out), __float_as_uint(bump_filter_width));
+    compiler.add_node_derivative(
+        NODE_GEOMETRY,
+        use_derivative,
+        compiler.encode_uchar4(NODE_GEOM_uv, bump_offset, store_derivatives),
+        compiler.stack_assign(out),
+        __float_as_uint(bump_filter_width));
   }
 
   out = output("Backfacing");
@@ -4037,29 +4088,31 @@ void GeometryNode::compile(SVMCompiler &compiler)
     compiler.add_node(NODE_LIGHT_PATH, NODE_LP_backfacing, compiler.stack_assign(out));
   }
 
+  const AttributeNode attr_node = attr_node_copy_from(this);
+
   out = output("Pointiness");
   if (!out->links.empty()) {
     if (compiler.output_type() != SHADER_TYPE_VOLUME) {
-      compiler.add_node(attr_node,
+      compiler.add_node(&attr_node,
                         ATTR_STD_POINTINESS,
                         compiler.encode_uchar4(compiler.stack_assign(out), NODE_ATTR_OUTPUT_FLOAT),
                         __float_as_uint(bump_filter_width));
     }
     else {
-      compiler.add_node(NODE_VALUE_F, __float_as_int(0.0f), compiler.stack_assign(out));
+      compiler.add_value_node(this, __float_as_int(0.0f), compiler.stack_assign(out));
     }
   }
 
   out = output("Random Per Island");
   if (!out->links.empty()) {
     if (compiler.output_type() != SHADER_TYPE_VOLUME) {
-      compiler.add_node(attr_node,
+      compiler.add_node(&attr_node,
                         ATTR_STD_RANDOM_PER_ISLAND,
                         compiler.encode_uchar4(compiler.stack_assign(out), NODE_ATTR_OUTPUT_FLOAT),
                         __float_as_uint(bump_filter_width));
     }
     else {
-      compiler.add_node(NODE_VALUE_F, __float_as_int(0.0f), compiler.stack_assign(out));
+      compiler.add_value_node(this, __float_as_int(0.0f), compiler.stack_assign(out));
     }
   }
 }
@@ -4127,86 +4180,108 @@ void TextureCoordinateNode::attributes(Shader *shader, AttributeRequestSet *attr
   ShaderNode::attributes(shader, attributes);
 }
 
+ShaderNodeType TextureCoordinateNode::shader_node_type() const
+{
+  return NODE_TEX_COORD;
+}
+
 void TextureCoordinateNode::compile(SVMCompiler &compiler)
 {
+  const uint bump_offset = shader_bump_to_node_bump_offset(bump);
+  const bool use_derivative = need_derivatives() || (bump != SHADER_BUMP_NONE);
+  const bool store_derivatives = need_derivatives();
   ShaderOutput *out;
-  ShaderNodeType texco_node = NODE_TEX_COORD;
-  ShaderNodeType attr_node = NODE_ATTR;
-  ShaderNodeType geom_node = NODE_GEOMETRY;
 
-  if (bump == SHADER_BUMP_DX) {
-    texco_node = NODE_TEX_COORD_BUMP_DX;
-    attr_node = NODE_ATTR_BUMP_DX;
-    geom_node = NODE_GEOMETRY_BUMP_DX;
-  }
-  else if (bump == SHADER_BUMP_DY) {
-    texco_node = NODE_TEX_COORD_BUMP_DY;
-    attr_node = NODE_ATTR_BUMP_DY;
-    geom_node = NODE_GEOMETRY_BUMP_DY;
-  }
+  const AttributeNode attr_node = attr_node_copy_from(this);
+  const GeometryNode geom_node = geom_node_copy_from(this);
 
   out = output("Generated");
   if (!out->links.empty()) {
     if (compiler.background) {
-      compiler.add_node(
-          geom_node, NODE_GEOM_P, compiler.stack_assign(out), __float_as_uint(bump_filter_width));
+      compiler.add_node_derivative(
+          NODE_GEOMETRY,
+          use_derivative,
+          compiler.encode_uchar4(NODE_GEOM_P, bump_offset, store_derivatives),
+          compiler.stack_assign(out),
+          __float_as_uint(bump_filter_width));
     }
     else {
       if (from_dupli) {
-        compiler.add_node(texco_node,
-                          NODE_TEXCO_DUPLI_GENERATED,
-                          compiler.stack_assign(out),
-                          __float_as_uint(bump_filter_width));
+        /* Dupli generated coordinates are constant, no bump offset. */
+        compiler.add_node_derivative(NODE_TEX_COORD,
+                                     use_derivative,
+                                     compiler.encode_uchar4(NODE_TEXCO_DUPLI_GENERATED,
+                                                            NODE_BUMP_OFFSET_CENTER,
+                                                            store_derivatives),
+                                     compiler.stack_assign(out),
+                                     __float_as_uint(bump_filter_width));
       }
       else if (compiler.output_type() == SHADER_TYPE_VOLUME) {
-        compiler.add_node(texco_node,
-                          NODE_TEXCO_VOLUME_GENERATED,
-                          compiler.stack_assign(out),
-                          __float_as_uint(bump_filter_width));
+        compiler.add_node_derivative(
+            NODE_TEX_COORD,
+            use_derivative,
+            compiler.encode_uchar4(NODE_TEXCO_VOLUME_GENERATED, bump_offset, store_derivatives),
+            compiler.stack_assign(out),
+            __float_as_uint(bump_filter_width));
       }
       else {
         const int attr = compiler.attribute(ATTR_STD_GENERATED);
-        compiler.add_node(
-            attr_node,
-            attr,
-            compiler.encode_uchar4(compiler.stack_assign(out), NODE_ATTR_OUTPUT_FLOAT3),
-            __float_as_uint(bump_filter_width));
+        compiler.add_node_derivative(NODE_ATTR,
+                                     use_derivative,
+                                     attr,
+                                     compiler.encode_uchar4(compiler.stack_assign(out),
+                                                            NODE_ATTR_OUTPUT_FLOAT3,
+                                                            bump_offset,
+                                                            store_derivatives),
+                                     __float_as_uint(bump_filter_width));
       }
     }
   }
 
   out = output("Normal");
   if (!out->links.empty()) {
-    compiler.add_node(texco_node,
-                      NODE_TEXCO_NORMAL,
-                      compiler.stack_assign(out),
-                      __float_as_uint(bump_filter_width));
+    compiler.add_node_derivative(
+        NODE_TEX_COORD,
+        use_derivative,
+        compiler.encode_uchar4(NODE_TEXCO_NORMAL, bump_offset, store_derivatives),
+        compiler.stack_assign(out),
+        __float_as_uint(bump_filter_width));
   }
 
   out = output("UV");
   if (!out->links.empty()) {
     if (from_dupli) {
-      compiler.add_node(texco_node,
-                        NODE_TEXCO_DUPLI_UV,
-                        compiler.stack_assign(out),
-                        __float_as_uint(bump_filter_width));
+      /* Dupli UV coordinates arent constant, no bump offset. */
+      compiler.add_node_derivative(
+          NODE_TEX_COORD,
+          use_derivative,
+          compiler.encode_uchar4(NODE_TEXCO_DUPLI_UV, NODE_BUMP_OFFSET_CENTER, store_derivatives),
+          compiler.stack_assign(out),
+          __float_as_uint(bump_filter_width));
     }
     else {
       const int attr = compiler.attribute(ATTR_STD_UV);
-      compiler.add_node(
-          attr_node,
+      compiler.add_node_derivative(
+          NODE_ATTR,
+          use_derivative,
           attr,
-          compiler.encode_uchar4(compiler.stack_assign(out), NODE_ATTR_OUTPUT_FLOAT3),
+          compiler.encode_uchar4(
+              compiler.stack_assign(out), NODE_ATTR_OUTPUT_FLOAT3, bump_offset, store_derivatives),
           __float_as_uint(bump_filter_width));
     }
   }
 
   out = output("Object");
   if (!out->links.empty()) {
-    compiler.add_node(texco_node,
-                      (use_transform) ? NODE_TEXCO_OBJECT_WITH_TRANSFORM : NODE_TEXCO_OBJECT,
-                      compiler.stack_assign(out),
-                      __float_as_uint(bump_filter_width));
+    compiler.add_node_derivative(NODE_TEX_COORD,
+                                 use_derivative,
+                                 compiler.encode_uchar4((use_transform) ?
+                                                            NODE_TEXCO_OBJECT_WITH_TRANSFORM :
+                                                            NODE_TEXCO_OBJECT,
+                                                        bump_offset,
+                                                        store_derivatives),
+                                 compiler.stack_assign(out),
+                                 __float_as_uint(bump_filter_width));
     if (use_transform) {
       const Transform ob_itfm = transform_inverse(ob_tfm);
       compiler.add_node(ob_itfm.x);
@@ -4217,31 +4292,43 @@ void TextureCoordinateNode::compile(SVMCompiler &compiler)
 
   out = output("Camera");
   if (!out->links.empty()) {
-    compiler.add_node(texco_node,
-                      NODE_TEXCO_CAMERA,
-                      compiler.stack_assign(out),
-                      __float_as_uint(bump_filter_width));
+    compiler.add_node_derivative(
+        NODE_TEX_COORD,
+        use_derivative,
+        compiler.encode_uchar4(NODE_TEXCO_CAMERA, bump_offset, store_derivatives),
+        compiler.stack_assign(out),
+        __float_as_uint(bump_filter_width));
   }
 
   out = output("Window");
   if (!out->links.empty()) {
-    compiler.add_node(texco_node,
-                      NODE_TEXCO_WINDOW,
-                      compiler.stack_assign(out),
-                      __float_as_uint(bump_filter_width));
+    compiler.add_node_derivative(
+        NODE_TEX_COORD,
+        use_derivative,
+        compiler.encode_uchar4(NODE_TEXCO_WINDOW, bump_offset, store_derivatives),
+        compiler.stack_assign(out),
+        __float_as_uint(bump_filter_width));
   }
 
+  /* Reflection currently does not support bump offset. */
   out = output("Reflection");
   if (!out->links.empty()) {
     if (compiler.background) {
-      compiler.add_node(
-          geom_node, NODE_GEOM_I, compiler.stack_assign(out), __float_as_uint(bump_filter_width));
+      compiler.add_node_derivative(
+          NODE_GEOMETRY,
+          use_derivative,
+          compiler.encode_uchar4(NODE_GEOM_I, NODE_BUMP_OFFSET_CENTER, store_derivatives),
+          compiler.stack_assign(out),
+          __float_as_uint(bump_filter_width));
     }
     else {
-      compiler.add_node(texco_node,
-                        NODE_TEXCO_REFLECTION,
-                        compiler.stack_assign(out),
-                        __float_as_uint(bump_filter_width));
+      compiler.add_node_derivative(NODE_TEX_COORD,
+                                   use_derivative,
+                                   compiler.encode_uchar4(NODE_TEXCO_REFLECTION,
+                                                          NODE_BUMP_OFFSET_CENTER,
+                                                          store_derivatives),
+                                   compiler.stack_assign(out),
+                                   __float_as_uint(bump_filter_width));
     }
   }
 }
@@ -4308,30 +4395,30 @@ void UVMapNode::attributes(Shader *shader, AttributeRequestSet *attributes)
   ShaderNode::attributes(shader, attributes);
 }
 
+ShaderNodeType UVMapNode::shader_node_type() const
+{
+  return NODE_TEX_COORD;
+}
+
 void UVMapNode::compile(SVMCompiler &compiler)
 {
+  const uint bump_offset = shader_bump_to_node_bump_offset(bump);
+  const bool use_derivative = need_derivatives() || (bump != SHADER_BUMP_NONE);
+  const bool store_derivatives = need_derivatives();
   ShaderOutput *out = output("UV");
-  ShaderNodeType texco_node = NODE_TEX_COORD;
-  ShaderNodeType attr_node = NODE_ATTR;
-  int attr;
-
-  if (bump == SHADER_BUMP_DX) {
-    texco_node = NODE_TEX_COORD_BUMP_DX;
-    attr_node = NODE_ATTR_BUMP_DX;
-  }
-  else if (bump == SHADER_BUMP_DY) {
-    texco_node = NODE_TEX_COORD_BUMP_DY;
-    attr_node = NODE_ATTR_BUMP_DY;
-  }
 
   if (!out->links.empty()) {
     if (from_dupli) {
-      compiler.add_node(texco_node,
-                        NODE_TEXCO_DUPLI_UV,
-                        compiler.stack_assign(out),
-                        __float_as_uint(bump_filter_width));
+      /* Dupli UV coordinates are constant, no bump offset. */
+      compiler.add_node_derivative(
+          NODE_TEX_COORD,
+          use_derivative,
+          compiler.encode_uchar4(NODE_TEXCO_DUPLI_UV, NODE_BUMP_OFFSET_CENTER, store_derivatives),
+          compiler.stack_assign(out),
+          __float_as_uint(bump_filter_width));
     }
     else {
+      int attr;
       if (!attribute.empty()) {
         attr = compiler.attribute(attribute);
       }
@@ -4339,10 +4426,12 @@ void UVMapNode::compile(SVMCompiler &compiler)
         attr = compiler.attribute(ATTR_STD_UV);
       }
 
-      compiler.add_node(
-          attr_node,
+      compiler.add_node_derivative(
+          NODE_ATTR,
+          use_derivative,
           attr,
-          compiler.encode_uchar4(compiler.stack_assign(out), NODE_ATTR_OUTPUT_FLOAT3),
+          compiler.encode_uchar4(
+              compiler.stack_assign(out), NODE_ATTR_OUTPUT_FLOAT3, bump_offset, store_derivatives),
           __float_as_uint(bump_filter_width));
     }
   }
@@ -4956,8 +5045,15 @@ void VertexColorNode::attributes(Shader *shader, AttributeRequestSet *attributes
   ShaderNode::attributes(shader, attributes);
 }
 
+ShaderNodeType VertexColorNode::shader_node_type() const
+{
+  return NODE_VERTEX_COLOR;
+}
+
 void VertexColorNode::compile(SVMCompiler &compiler)
 {
+  const uint bump_offset = shader_bump_to_node_bump_offset(bump);
+  const bool use_derivative = need_derivatives() || (bump != SHADER_BUMP_NONE);
   ShaderOutput *color_out = output("Color");
   ShaderOutput *alpha_out = output("Alpha");
   int layer_id = 0;
@@ -4969,23 +5065,13 @@ void VertexColorNode::compile(SVMCompiler &compiler)
     layer_id = compiler.attribute(ATTR_STD_VERTEX_COLOR);
   }
 
-  ShaderNodeType node;
-
-  if (bump == SHADER_BUMP_DX) {
-    node = NODE_VERTEX_COLOR_BUMP_DX;
-  }
-  else if (bump == SHADER_BUMP_DY) {
-    node = NODE_VERTEX_COLOR_BUMP_DY;
-  }
-  else {
-    node = NODE_VERTEX_COLOR;
-  }
-
-  compiler.add_node(node,
-                    compiler.encode_uchar4(layer_id,
-                                           compiler.stack_assign(color_out),
-                                           compiler.stack_assign(alpha_out)),
-                    __float_as_uint(bump_filter_width));
+  compiler.add_node_derivative(NODE_VERTEX_COLOR,
+                               use_derivative,
+                               compiler.encode_uchar4(layer_id,
+                                                      compiler.stack_assign(color_out),
+                                                      compiler.stack_assign(alpha_out),
+                                                      bump_offset),
+                               __float_as_uint(bump_filter_width));
 }
 
 void VertexColorNode::compile(OSLCompiler &compiler)
@@ -5039,7 +5125,7 @@ void ValueNode::compile(SVMCompiler &compiler)
 {
   ShaderOutput *val_out = output("Value");
 
-  compiler.add_node(NODE_VALUE_F, __float_as_int(value), compiler.stack_assign(val_out));
+  compiler.add_value_node(this, __float_as_int(value), compiler.stack_assign(val_out));
 }
 
 void ValueNode::compile(OSLCompiler &compiler)
@@ -5072,8 +5158,7 @@ void ColorNode::compile(SVMCompiler &compiler)
   ShaderOutput *color_out = output("Color");
 
   if (!color_out->links.empty()) {
-    compiler.add_node(NODE_VALUE_V, compiler.stack_assign(color_out));
-    compiler.add_node(NODE_VALUE_V, value);
+    compiler.add_value_node(this, value, compiler.stack_assign(color_out));
   }
 }
 
@@ -5706,19 +5791,11 @@ void CombineXYZNode::constant_fold(const ConstantFolder &folder)
 
 void CombineXYZNode::compile(SVMCompiler &compiler)
 {
-  ShaderInput *x_in = input("X");
-  ShaderInput *y_in = input("Y");
-  ShaderInput *z_in = input("Z");
   ShaderOutput *vector_out = output("Vector");
 
-  compiler.add_node(
-      NODE_COMBINE_VECTOR, compiler.stack_assign(x_in), 0, compiler.stack_assign(vector_out));
-
-  compiler.add_node(
-      NODE_COMBINE_VECTOR, compiler.stack_assign(y_in), 1, compiler.stack_assign(vector_out));
-
-  compiler.add_node(
-      NODE_COMBINE_VECTOR, compiler.stack_assign(z_in), 2, compiler.stack_assign(vector_out));
+  compiler.add_node(this, compiler.stack_assign(input("X")), 0, compiler.stack_assign(vector_out));
+  compiler.add_node(this, compiler.stack_assign(input("Y")), 1, compiler.stack_assign(vector_out));
+  compiler.add_node(this, compiler.stack_assign(input("Z")), 2, compiler.stack_assign(vector_out));
 }
 
 void CombineXYZNode::compile(OSLCompiler &compiler)
@@ -5915,18 +5992,10 @@ void SeparateXYZNode::constant_fold(const ConstantFolder &folder)
 void SeparateXYZNode::compile(SVMCompiler &compiler)
 {
   ShaderInput *vector_in = input("Vector");
-  ShaderOutput *x_out = output("X");
-  ShaderOutput *y_out = output("Y");
-  ShaderOutput *z_out = output("Z");
 
-  compiler.add_node(
-      NODE_SEPARATE_VECTOR, compiler.stack_assign(vector_in), 0, compiler.stack_assign(x_out));
-
-  compiler.add_node(
-      NODE_SEPARATE_VECTOR, compiler.stack_assign(vector_in), 1, compiler.stack_assign(y_out));
-
-  compiler.add_node(
-      NODE_SEPARATE_VECTOR, compiler.stack_assign(vector_in), 2, compiler.stack_assign(z_out));
+  compiler.add_node(this, compiler.stack_assign(vector_in), 0, compiler.stack_assign(output("X")));
+  compiler.add_node(this, compiler.stack_assign(vector_in), 1, compiler.stack_assign(output("Y")));
+  compiler.add_node(this, compiler.stack_assign(vector_in), 2, compiler.stack_assign(output("Z")));
 }
 
 void SeparateXYZNode::compile(OSLCompiler &compiler)
@@ -6032,56 +6101,68 @@ void AttributeNode::attributes(Shader *shader, AttributeRequestSet *attributes)
   ShaderNode::attributes(shader, attributes);
 }
 
+ShaderNodeType AttributeNode::shader_node_type() const
+{
+  return NODE_ATTR;
+}
+
 void AttributeNode::compile(SVMCompiler &compiler)
 {
+  const uint bump_offset = shader_bump_to_node_bump_offset(bump);
+  const bool use_derivative = need_derivatives() || (bump != SHADER_BUMP_NONE);
+  const bool store_derivatives = need_derivatives();
   ShaderOutput *color_out = output("Color");
   ShaderOutput *vector_out = output("Vector");
   ShaderOutput *fac_out = output("Fac");
   ShaderOutput *alpha_out = output("Alpha");
-  ShaderNodeType attr_node = NODE_ATTR;
   const int attr = compiler.attribute_standard(attribute);
   const uint bump_filter_or_stochastic = (compiler.output_type() == SHADER_TYPE_VOLUME) ?
                                              stochastic_sample :
                                              __float_as_uint(bump_filter_width);
 
-  if (bump == SHADER_BUMP_DX) {
-    attr_node = NODE_ATTR_BUMP_DX;
-  }
-  else if (bump == SHADER_BUMP_DY) {
-    attr_node = NODE_ATTR_BUMP_DY;
-  }
-
   if (!color_out->links.empty() || !vector_out->links.empty()) {
     if (!color_out->links.empty()) {
-      compiler.add_node(
-          attr_node,
-          attr,
-          compiler.encode_uchar4(compiler.stack_assign(color_out), NODE_ATTR_OUTPUT_FLOAT3),
-          bump_filter_or_stochastic);
+      compiler.add_node_derivative(NODE_ATTR,
+                                   use_derivative,
+                                   attr,
+                                   compiler.encode_uchar4(compiler.stack_assign(color_out),
+                                                          NODE_ATTR_OUTPUT_FLOAT3,
+                                                          bump_offset,
+                                                          store_derivatives),
+                                   bump_filter_or_stochastic);
     }
     if (!vector_out->links.empty()) {
-      compiler.add_node(
-          attr_node,
-          attr,
-          compiler.encode_uchar4(compiler.stack_assign(vector_out), NODE_ATTR_OUTPUT_FLOAT3),
-          bump_filter_or_stochastic);
+      compiler.add_node_derivative(NODE_ATTR,
+                                   use_derivative,
+                                   attr,
+                                   compiler.encode_uchar4(compiler.stack_assign(vector_out),
+                                                          NODE_ATTR_OUTPUT_FLOAT3,
+                                                          bump_offset,
+                                                          store_derivatives),
+                                   bump_filter_or_stochastic);
     }
   }
 
   if (!fac_out->links.empty()) {
-    compiler.add_node(
-        attr_node,
-        attr,
-        compiler.encode_uchar4(compiler.stack_assign(fac_out), NODE_ATTR_OUTPUT_FLOAT),
-        bump_filter_or_stochastic);
+    compiler.add_node_derivative(NODE_ATTR,
+                                 use_derivative,
+                                 attr,
+                                 compiler.encode_uchar4(compiler.stack_assign(fac_out),
+                                                        NODE_ATTR_OUTPUT_FLOAT,
+                                                        bump_offset,
+                                                        store_derivatives),
+                                 bump_filter_or_stochastic);
   }
 
   if (!alpha_out->links.empty()) {
-    compiler.add_node(
-        attr_node,
-        attr,
-        compiler.encode_uchar4(compiler.stack_assign(alpha_out), NODE_ATTR_OUTPUT_FLOAT_ALPHA),
-        bump_filter_or_stochastic);
+    compiler.add_node_derivative(NODE_ATTR,
+                                 use_derivative,
+                                 attr,
+                                 compiler.encode_uchar4(compiler.stack_assign(alpha_out),
+                                                        NODE_ATTR_OUTPUT_FLOAT_ALPHA,
+                                                        bump_offset,
+                                                        store_derivatives),
+                                 bump_filter_or_stochastic);
   }
 }
 
@@ -6917,26 +6998,20 @@ void VectorMathNode::compile(SVMCompiler &compiler)
   const int value_stack_offset = compiler.stack_assign_if_linked(value_out);
   const int vector_stack_offset = compiler.stack_assign_if_linked(vector_out);
 
+  compiler.add_node(
+      this,
+      math_type,
+      compiler.encode_uchar4(vector1_stack_offset, vector2_stack_offset, param1_stack_offset),
+      compiler.encode_uchar4(value_stack_offset, vector_stack_offset));
+
   /* 3 Vector Operators */
   if (math_type == NODE_VECTOR_MATH_WRAP || math_type == NODE_VECTOR_MATH_FACEFORWARD ||
       math_type == NODE_VECTOR_MATH_MULTIPLY_ADD)
   {
     ShaderInput *vector3_in = input("Vector3");
     const int vector3_stack_offset = compiler.stack_assign(vector3_in);
-    compiler.add_node(
-        NODE_VECTOR_MATH,
-        math_type,
-        compiler.encode_uchar4(vector1_stack_offset, vector2_stack_offset, param1_stack_offset),
-        compiler.encode_uchar4(value_stack_offset, vector_stack_offset));
     compiler.add_node(vector3_stack_offset);
   }
-  else {
-    compiler.add_node(
-        NODE_VECTOR_MATH,
-        math_type,
-        compiler.encode_uchar4(vector1_stack_offset, vector2_stack_offset, param1_stack_offset),
-        compiler.encode_uchar4(value_stack_offset, vector_stack_offset));
-  }
 }
 
 void VectorMathNode::compile(OSLCompiler &compiler)
@@ -6982,7 +7057,7 @@ void VectorRotateNode::compile(SVMCompiler &compiler)
   ShaderInput *angle_in = input("Angle");
   ShaderOutput *vector_out = output("Vector");
 
-  compiler.add_node(NODE_VECTOR_ROTATE,
+  compiler.add_node(this,
                     compiler.encode_uchar4(rotate_type,
                                            compiler.stack_assign(vector_in),
                                            compiler.stack_assign(rotation_in),
@@ -7033,7 +7108,7 @@ void VectorTransformNode::compile(SVMCompiler &compiler)
   ShaderOutput *vector_out = output("Vector");
 
   compiler.add_node(
-      NODE_VECTOR_TRANSFORM,
+      this,
       compiler.encode_uchar4(transform_type, convert_from, convert_to),
       compiler.encode_uchar4(compiler.stack_assign(vector_in), compiler.stack_assign(vector_out)));
 }
@@ -7157,10 +7232,7 @@ void CurvesNode::constant_fold(const ConstantFolder &folder, ShaderInput *value_
   }
 }
 
-void CurvesNode::compile(SVMCompiler &compiler,
-                         const int type,
-                         ShaderInput *value_in,
-                         ShaderOutput *value_out)
+void CurvesNode::compile(SVMCompiler &compiler, ShaderInput *value_in, ShaderOutput *value_out)
 {
   if (curves.size() == 0) {
     return;
@@ -7168,7 +7240,7 @@ void CurvesNode::compile(SVMCompiler &compiler,
 
   ShaderInput *fac_in = input("Fac");
 
-  compiler.add_node(ShaderNodeType(type),
+  compiler.add_node(this,
                     compiler.encode_uchar4(compiler.stack_assign(fac_in),
                                            compiler.stack_assign(value_in),
                                            compiler.stack_assign(value_out),
@@ -7233,7 +7305,7 @@ void RGBCurvesNode::constant_fold(const ConstantFolder &folder)
 
 void RGBCurvesNode::compile(SVMCompiler &compiler)
 {
-  CurvesNode::compile(compiler, NODE_CURVES, input("Color"), output("Color"));
+  CurvesNode::compile(compiler, input("Color"), output("Color"));
 }
 
 void RGBCurvesNode::compile(OSLCompiler &compiler)
@@ -7269,7 +7341,7 @@ void VectorCurvesNode::constant_fold(const ConstantFolder &folder)
 
 void VectorCurvesNode::compile(SVMCompiler &compiler)
 {
-  CurvesNode::compile(compiler, NODE_CURVES, input("Vector"), output("Vector"));
+  CurvesNode::compile(compiler, input("Vector"), output("Vector"));
 }
 
 void VectorCurvesNode::compile(OSLCompiler &compiler)
@@ -7459,9 +7531,7 @@ void SetNormalNode::compile(SVMCompiler &compiler)
   ShaderInput *direction_in = input("Direction");
   ShaderOutput *normal_out = output("Normal");
 
-  compiler.add_node(NODE_CLOSURE_SET_NORMAL,
-                    compiler.stack_assign(direction_in),
-                    compiler.stack_assign(normal_out));
+  compiler.add_node(this, compiler.stack_assign(direction_in), compiler.stack_assign(normal_out));
 }
 
 void SetNormalNode::compile(OSLCompiler &compiler)
@@ -7837,7 +7907,7 @@ void TangentNode::compile(SVMCompiler &compiler)
   }
 
   compiler.add_node(
-      NODE_TANGENT,
+      this,
       compiler.encode_uchar4(compiler.stack_assign(tangent_out), direction_type, axis),
       attr);
 }
@@ -8090,7 +8160,7 @@ void RaycastNode::compile(SVMCompiler &compiler)
   ShaderOutput *hit_position_out = output("Hit Position");
   ShaderOutput *hit_normal_out = output("Hit Normal");
 
-  compiler.add_node(NODE_RAYCAST,
+  compiler.add_node(this,
                     compiler.encode_uchar4(compiler.stack_assign(position_in),
                                            compiler.stack_assign(direction_in),
                                            compiler.stack_assign(length_in),

intern/cycles/scene/shader_nodes.h

@@ -86,6 +86,12 @@ class ImageSlotTextureNode : public TextureNode {
 
   virtual void update_images(const SVMCompiler &compiler) = 0;
 
+  bool is_texture_node_and_needs_derivatives(const SVMCompiler &compiler) override
+  {
+    update_images(compiler);
+    return need_derivatives();
+  }
+
   ImageHandle handle;
 };
 
@@ -109,6 +115,8 @@ class ImageTextureNode : public ImageSlotTextureNode {
 
   void update_images(const SVMCompiler &compiler) override;
 
+  ShaderNodeType shader_node_type() const override;
+
   /* Parameters. */
   NODE_SOCKET_API(ustring, filename)
   NODE_SOCKET_API(ustring, colorspace)
@@ -145,6 +153,11 @@ class EnvironmentTextureNode : public ImageSlotTextureNode {
 
   void update_images(const SVMCompiler &compiler) override;
 
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_TEX_ENVIRONMENT;
+  }
+
   /* Parameters. */
   NODE_SOCKET_API(ustring, filename)
   NODE_SOCKET_API(ustring, colorspace)
@@ -394,6 +407,10 @@ class MappingNode : public ShaderNode {
  public:
   SHADER_NODE_CLASS(MappingNode)
   void constant_fold(const ConstantFolder &folder) override;
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_MAPPING;
+  }
 
   NODE_SOCKET_API(float3, vector)
   NODE_SOCKET_API(float3, location)
@@ -410,6 +427,10 @@ class RGBToBWNode : public ShaderNode {
   {
     return true;
   }
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_CONVERT;
+  }
 
   NODE_SOCKET_API(float3, color)
 };
@@ -421,11 +442,15 @@ class ConvertNode : public ShaderNode {
   SHADER_NODE_BASE_CLASS(ConvertNode)
 
   void constant_fold(const ConstantFolder &folder) override;
-
   bool is_linear_operation() override
   {
     return true;
   }
+  NodeConvert convert_type();
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_CONVERT;
+  }
 
  private:
   SocketType::Type from, to;
@@ -474,6 +499,11 @@ class BsdfBaseNode : public ShaderNode {
     return ShaderNode::get_feature() | KERNEL_FEATURE_NODE_BSDF;
   }
 
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_CLOSURE_BSDF;
+  }
+
  protected:
   ClosureType closure;
 };
@@ -987,6 +1017,7 @@ class GeometryNode : public ShaderNode {
     return true;
   }
   int get_group();
+  ShaderNodeType shader_node_type() const override;
 };
 
 class TextureCoordinateNode : public ShaderNode {
@@ -1001,6 +1032,7 @@ class TextureCoordinateNode : public ShaderNode {
   {
     return true;
   }
+  ShaderNodeType shader_node_type() const override;
 
   NODE_SOCKET_API(bool, from_dupli)
   NODE_SOCKET_API(bool, use_transform)
@@ -1019,6 +1051,7 @@ class UVMapNode : public ShaderNode {
   {
     return true;
   }
+  ShaderNodeType shader_node_type() const override;
 
   NODE_SOCKET_API(ustring, attribute)
   NODE_SOCKET_API(bool, from_dupli)
@@ -1113,6 +1146,7 @@ class VertexColorNode : public ShaderNode {
   {
     return true;
   }
+  ShaderNodeType shader_node_type() const override;
 
   NODE_SOCKET_API(ustring, layer_name)
 };
@@ -1277,6 +1311,10 @@ class CombineXYZNode : public ShaderNode {
   {
     return true;
   }
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_COMBINE_VECTOR;
+  }
 
   NODE_SOCKET_API(float, x)
   NODE_SOCKET_API(float, y)
@@ -1323,6 +1361,10 @@ class SeparateXYZNode : public ShaderNode {
   {
     return true;
   }
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_SEPARATE_VECTOR;
+  }
 
   NODE_SOCKET_API(float3, vector)
 };
@@ -1350,6 +1392,7 @@ class AttributeNode : public ShaderNode {
   {
     return true;
   }
+  ShaderNodeType shader_node_type() const override;
 
   NODE_SOCKET_API(ustring, attribute)
 
@@ -1489,6 +1532,10 @@ class VectorMathNode : public ShaderNode {
   SHADER_NODE_CLASS(VectorMathNode)
   void constant_fold(const ConstantFolder &folder) override;
   bool is_linear_operation() override;
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_VECTOR_MATH;
+  }
 
   NODE_SOCKET_API(float3, vector1)
   NODE_SOCKET_API(float3, vector2)
@@ -1508,6 +1555,11 @@ class VectorRotateNode : public ShaderNode {
   NODE_SOCKET_API(float3, axis)
   NODE_SOCKET_API(float, angle)
   NODE_SOCKET_API(float3, rotation)
+
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_VECTOR_ROTATE;
+  }
 };
 
 class VectorTransformNode : public ShaderNode {
@@ -1518,6 +1570,11 @@ class VectorTransformNode : public ShaderNode {
   NODE_SOCKET_API(NodeVectorTransformConvertSpace, convert_from)
   NODE_SOCKET_API(NodeVectorTransformConvertSpace, convert_to)
   NODE_SOCKET_API(float3, vector)
+
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_VECTOR_TRANSFORM;
+  }
 };
 
 class BumpNode : public ShaderNode {
@@ -1560,11 +1617,12 @@ class CurvesNode : public ShaderNode {
  protected:
   using ShaderNode::constant_fold;
   void constant_fold(const ConstantFolder &folder, ShaderInput *value_in);
-  void compile(SVMCompiler &compiler,
-               const int type,
-               ShaderInput *value_in,
-               ShaderOutput *value_out);
+  void compile(SVMCompiler &compiler, ShaderInput *value_in, ShaderOutput *value_out);
   void compile(OSLCompiler &compiler, const char *name);
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_CURVES;
+  }
 };
 
 class RGBCurvesNode : public CurvesNode {
@@ -1607,6 +1665,11 @@ class SetNormalNode : public ShaderNode {
  public:
   SHADER_NODE_CLASS(SetNormalNode)
   NODE_SOCKET_API(float3, direction)
+
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_CLOSURE_SET_NORMAL;
+  }
 };
 
 class OSLNode final : public ShaderNode {
@@ -1710,6 +1773,10 @@ class TangentNode : public ShaderNode {
   {
     return true;
   }
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_TANGENT;
+  }
 
   NODE_SOCKET_API(NodeTangentDirectionType, direction_type)
   NODE_SOCKET_API(NodeTangentAxis, axis)
@@ -1782,6 +1849,10 @@ class RaycastNode : public ShaderNode {
   {
     return KERNEL_FEATURE_NODE_RAYTRACE;
   }
+  ShaderNodeType shader_node_type() const override
+  {
+    return NODE_RAYCAST;
+  }
 
   NODE_SOCKET_API(float3, position)
   NODE_SOCKET_API(float3, direction)

intern/cycles/scene/svm.cpp

@@ -18,6 +18,7 @@
 
 #include "util/log.h"
 #include "util/progress.h"
+#include "util/queue.h"
 #include "util/task.h"
 
 CCL_NAMESPACE_BEGIN
@@ -190,6 +191,14 @@ int SVMCompiler::stack_size(SocketType::Type type)
   return size;
 }
 
+int SVMCompiler::stack_size(const ShaderIO *io)
+{
+  const SocketType::Type type = io->type();
+  const bool derivative = io->parent->need_derivatives();
+
+  return derivative ? stack_size(type) * 3 : stack_size(type);
+}
+
 int SVMCompiler::stack_find_offset(const int size)
 {
   int offset = -1;
@@ -224,14 +233,14 @@ int SVMCompiler::stack_find_offset(const int size)
   return 0;
 }
 
-int SVMCompiler::stack_find_offset(SocketType::Type type)
+int SVMCompiler::stack_find_offset(const ShaderIO *io)
 {
-  return stack_find_offset(stack_size(type));
+  return stack_find_offset(stack_size(io));
 }
 
-void SVMCompiler::stack_clear_offset(SocketType::Type type, const int offset)
+void SVMCompiler::stack_clear_offset(const ShaderIO *io, const int offset)
 {
-  const int size = stack_size(type);
+  const int size = stack_size(io);
 
   for (int i = 0; i < size; i++) {
     active_stack.users[offset + i]--;
@@ -248,25 +257,22 @@ int SVMCompiler::stack_assign(ShaderInput *input)
       input->stack_offset = input->link->stack_offset;
     }
     else {
-      Node *node = input->parent;
+      const ShaderNode *node = input->parent;
 
       /* not linked to output -> add nodes to load default value */
-      input->stack_offset = stack_find_offset(input->type());
+      input->stack_offset = stack_find_offset(input);
 
       if (input->type() == SocketType::FLOAT) {
-        add_node(NODE_VALUE_F,
-                 __float_as_int(node->get_float(input->socket_type)),
-                 input->stack_offset);
+        add_value_node(
+            node, __float_as_int(node->get_float(input->socket_type)), input->stack_offset);
       }
       else if (input->type() == SocketType::INT) {
-        add_node(NODE_VALUE_F, node->get_int(input->socket_type), input->stack_offset);
+        add_value_node(node, node->get_int(input->socket_type), input->stack_offset);
       }
       else if (input->type() == SocketType::VECTOR || input->type() == SocketType::NORMAL ||
                input->type() == SocketType::POINT || input->type() == SocketType::COLOR)
       {
-
-        add_node(NODE_VALUE_V, input->stack_offset);
-        add_node(NODE_VALUE_V, node->get_float3(input->socket_type));
+        add_value_node(node, node->get_float3(input->socket_type), input->stack_offset);
       }
       else { /* should not get called for closure */
         assert(0);
@@ -281,7 +287,7 @@ int SVMCompiler::stack_assign(ShaderOutput *output)
 {
   /* if no stack offset assigned yet, find one */
   if (output->stack_offset == SVM_STACK_INVALID) {
-    output->stack_offset = stack_find_offset(output->type());
+    output->stack_offset = stack_find_offset(output);
   }
 
   return output->stack_offset;
@@ -335,8 +341,7 @@ void SVMCompiler::stack_link(ShaderInput *input, ShaderOutput *output)
     assert(stack_size(output->type()) == stack_size(input->link->type()));
 
     output->stack_offset = input->link->stack_offset;
-
-    const int size = stack_size(output->type());
+    const int size = stack_size(output);
 
     for (int i = 0; i < size; i++) {
       active_stack.users[output->stack_offset + i]++;
@@ -367,7 +372,7 @@ void SVMCompiler::stack_clear_users(ShaderNode *node, ShaderNodeSet &done)
       }
 
       if (all_done) {
-        stack_clear_offset(output->type(), output->stack_offset);
+        stack_clear_offset(output, output->stack_offset);
         output->stack_offset = SVM_STACK_INVALID;
 
         for (ShaderInput *in : output->links) {
@@ -382,7 +387,7 @@ void SVMCompiler::stack_clear_temporary(ShaderNode *node)
 {
   for (ShaderInput *input : node->inputs) {
     if (!input->link && input->stack_offset != SVM_STACK_INVALID) {
-      stack_clear_offset(input->type(), input->stack_offset);
+      stack_clear_offset(input, input->stack_offset);
       input->stack_offset = SVM_STACK_INVALID;
     }
   }
@@ -408,12 +413,42 @@ void SVMCompiler::add_node(ShaderNodeType type, const int a, int b, const int c)
   svm_node_types_used[type] = true;
   current_svm_nodes.push_back_slow(make_int4(type, a, b, c));
 }
-
-void SVMCompiler::add_node(ShaderNodeType type, const float3 &f)
+static ShaderNodeType svm_node_type_with_derivatives(ShaderNodeType type)
 {
-  svm_node_types_used[type] = true;
-  current_svm_nodes.push_back_slow(
-      make_int4(type, __float_as_int(f.x), __float_as_int(f.y), __float_as_int(f.z)));
+  switch (type) {
+#define SHADER_NODE_TYPE_DERIVATIVE(name) \
+  case name: \
+    return name##_DERIVATIVE;
+#include "kernel/svm/node_types_template.h"
+    default:
+      break;
+  }
+
+  return type;
+}
+
+void SVMCompiler::add_node_derivative(
+    const ShaderNodeType type, const bool need_derivatives, const int a, const int b, const int c)
+{
+  /* Only support derivatives for surface for now. */
+  const ShaderNodeType node_x = (need_derivatives && current_type != SHADER_TYPE_VOLUME) ?
+                                    svm_node_type_with_derivatives(type) :
+                                    type;
+  svm_node_types_used[node_x] = true;
+  add_node(node_x, a, b, c);
+}
+
+void SVMCompiler::add_node(const ShaderNode *node, const int a, int b, const int c)
+{
+  const ShaderNodeType type = node->shader_node_type();
+  assert(type != NODE_NONE);
+  add_node_derivative(type, node->need_derivatives(), a, b, c);
+}
+
+void SVMCompiler::add_node(const ShaderNodeType type, const float3 &f, const bool need_derivatives)
+{
+  add_node_derivative(
+      type, need_derivatives, __float_as_int(f.x), __float_as_int(f.y), __float_as_int(f.z));
 }
 
 void SVMCompiler::add_node(const float4 &f)
@@ -422,6 +457,19 @@ void SVMCompiler::add_node(const float4 &f)
       __float_as_int(f.x), __float_as_int(f.y), __float_as_int(f.z), __float_as_int(f.w)));
 }
 
+void SVMCompiler::add_value_node(const ShaderNode *node, const int value, const int stack_offset)
+{
+  add_node_derivative(NODE_VALUE_F, node->need_derivatives(), value, stack_offset);
+}
+
+void SVMCompiler::add_value_node(const ShaderNode *node,
+                                 const float3 &value,
+                                 const int stack_offset)
+{
+  add_node_derivative(NODE_VALUE_V, node->need_derivatives(), stack_offset);
+  add_node(NODE_VALUE_V, value, node->need_derivatives());
+}
+
 uint SVMCompiler::attribute(ustring name)
 {
   return scene->shader_manager->get_attribute_id(name);
@@ -750,6 +798,43 @@ void SVMCompiler::generate_multi_closure(ShaderNode *root_node,
   state->nodes_done_flag[node->id] = true;
 }
 
+static void mark_nodes_requiring_derivatives(const SVMCompiler &compiler,
+                                             ShaderGraph *graph,
+                                             const ShaderType type)
+{
+  if (type == SHADER_TYPE_VOLUME) {
+    /* Only support derivatives for surface for now. */
+    return;
+  }
+  queue<ShaderNode *> traverse_queue;
+  ShaderNodeSet scheduled;
+  /* Check if texture nodes need derivatives. */
+  for (ShaderNode *node : graph->nodes) {
+    if (node->is_texture_node_and_needs_derivatives(compiler)) {
+      traverse_queue.push(node);
+      scheduled.insert(node);
+    }
+  }
+  /* Mark all ancestors of texture nodes as requiring derivatives, if the texture nodes themselves
+   * need derivatives. */
+  while (!traverse_queue.empty()) {
+    ShaderNode *node = traverse_queue.front();
+    traverse_queue.pop();
+    node->set_need_derivatives();
+    LOG_DEBUG << "Marking " << node->name << " as requiring derivatives";
+    for (ShaderInput *input : node->inputs) {
+      if (input->link == nullptr) {
+        continue;
+      }
+      if (scheduled.find(input->link->parent) != scheduled.end()) {
+        continue;
+      }
+      traverse_queue.push(input->link->parent);
+      scheduled.insert(input->link->parent);
+    }
+  }
+}
+
 void SVMCompiler::compile_type(Shader *shader, ShaderGraph *graph, ShaderType type)
 {
   /* Converting a shader graph into svm_nodes that can be executed
@@ -806,6 +891,8 @@ void SVMCompiler::compile_type(Shader *shader, ShaderGraph *graph, ShaderType ty
     }
   }
 
+  mark_nodes_requiring_derivatives(*this, graph, type);
+
   /* for the bump shader we need add a node to store the shader state */
   const bool need_bump_state = (type == SHADER_TYPE_BUMP) &&
                                (shader->get_displacement_method() == DISPLACE_BOTH);

intern/cycles/scene/svm.h

@@ -89,14 +89,22 @@ class SVMCompiler {
   int stack_assign_if_not_equal(ShaderInput *input, const float value);
   int stack_assign_if_not_equal(ShaderInput *input, const float3 value);
   int stack_find_offset(const int size);
-  int stack_find_offset(SocketType::Type type);
-  void stack_clear_offset(SocketType::Type type, const int offset);
+  int stack_find_offset(const ShaderIO *io);
+  void stack_clear_offset(const ShaderIO *io, const int offset);
   void stack_link(ShaderInput *input, ShaderOutput *output);
 
   void add_node(ShaderNodeType type, const int a = 0, const int b = 0, const int c = 0);
+  void add_node_derivative(const ShaderNodeType type,
+                           const bool need_derivatives,
+                           const int a = 0,
+                           const int b = 0,
+                           const int c = 0);
   void add_node(const int a = 0, const int b = 0, const int c = 0, const int d = 0);
-  void add_node(ShaderNodeType type, const float3 &f);
+  void add_node(const ShaderNode *node, const int a = 0, const int b = 0, const int c = 0);
+  void add_node(ShaderNodeType type, const float3 &f, const bool need_derivatives = 0);
   void add_node(const float4 &f);
+  void add_value_node(const ShaderNode *node, const int value, const int stack_offset);
+  void add_value_node(const ShaderNode *node, const float3 &value, const int stack_offset);
   uint attribute(ustring name);
   uint attribute(AttributeStandard std);
   uint attribute_standard(ustring name);
@@ -194,6 +202,7 @@ class SVMCompiler {
 
   void stack_clear_temporary(ShaderNode *node);
   int stack_size(SocketType::Type type);
+  int stack_size(const ShaderIO *io);
   void stack_clear_users(ShaderNode *node, ShaderNodeSet &done);
 
   /* single closure */

intern/cycles/util/math_base.h

@@ -678,6 +678,11 @@ ccl_device float bits_to_01(const uint bits)
   return bits * (1.0f / (float)0xFFFFFFFF);
 }
 
+ccl_device_inline bool is_zero(const float a)
+{
+  return a == 0.0f;
+}
+
 #if !defined(__KERNEL_GPU__)
 #  if defined(__GNUC__)
 ccl_device_inline uint popcount(const uint x)

intern/cycles/util/math_dual.h

@@ -5,7 +5,9 @@
 #pragma once
 
 #include "util/math_base.h"
+#include "util/math_float3.h"
 #include "util/types_dual.h"
+#include "util/types_float3.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -29,18 +31,212 @@ ccl_device_template_spec dual4 make_zero()
   return dual4();
 }
 
+template<class T> ccl_device_inline bool is_zero(const ccl_private dual<T> &a)
+{
+  return is_zero(a.val);
+}
+
+ccl_device_inline bool operator<(const ccl_private dual1 &a, const float b)
+{
+  return a.val < b;
+}
+
 /* Multiplication of dual by scalar. */
 template<class T1, class T2> ccl_device_inline dual<T1> operator*(const dual<T1> a, T2 b)
 {
   return {a.val * b, a.dx * b, a.dy * b};
 }
 
+/* Multiplication of scalar by dual. */
+template<class T> ccl_device_inline dual<T> operator*(const T a, const ccl_private dual<T> &b)
+{
+  return {a * b.val, a * b.dx, a * b.dy};
+}
+
+/* Multiplication of duals.
+ * (uv)' = uv' + u'v. */
+template<class T1, class T2>
+ccl_device_inline dual<T1> operator*(const ccl_private dual<T1> &u, const ccl_private dual<T2> &v)
+{
+  return {u.val * v.val, u.val * v.dx + u.dx * v.val, u.val * v.dy + u.dy * v.val};
+}
+
+/* Division of dual by scalar. */
+template<class T> ccl_device_inline dual<T> operator/(const dual<T> a, T b)
+{
+  const T inv_b = 1.0f / b;
+  return {a.val * inv_b, a.dx * inv_b, a.dy * inv_b};
+}
+
+/* Division of dual by dual.
+ * (u/v)' = (u' - v' * u/v) / v. */
+template<class T1, class T2>
+ccl_device_inline dual<T1> operator/(const ccl_private dual<T1> &u, const ccl_private dual<T2> &v)
+{
+  const T2 inv_v = 1.0f / v.val;
+  /* NOTE: Numerically `u/v != u*inv_v`, for compatibility we compute `u/v`. */
+  const T1 u_v = u.val / v.val;
+  return {u_v, (u.dx - u_v * v.dx) * inv_v, (u.dy - u_v * v.dy) * inv_v};
+}
+
+template<class T1, class T2>
+ccl_device_inline dual<T1> operator/=(ccl_private dual<T1> &a, const ccl_private dual<T2> &b)
+{
+  return a = a / b;
+}
+
+/* Addition of duals. */
+template<class T> ccl_device_inline dual<T> operator+(const dual<T> a, const dual<T> b)
+{
+  return {a.val + b.val, a.dx + b.dx, a.dy + b.dy};
+}
+
+/* Addition of dual and scalar. */
+template<class T1, class T2> ccl_device_inline dual<T1> operator+(const dual<T1> a, T2 b)
+{
+  return {a.val + b, a.dx, a.dy};
+}
+
+/* Addition of scalar and dual. */
+template<class T1, class T2> ccl_device_inline dual<T2> operator+(const T1 a, const dual<T2> b)
+{
+  return {a + b.val, b.dx, b.dy};
+}
+
+/* Subtraction of dual by scalar. */
+template<class T1, class T2> ccl_device_inline dual<T1> operator-(const dual<T1> a, T2 b)
+{
+  return {a.val - b, a.dx, a.dy};
+}
+
+/* Subtraction of scalar by dual. */
+template<class T1, class T2> ccl_device_inline dual<T2> operator-(const T1 a, const dual<T2> b)
+{
+  return {a - b.val, -b.dx, -b.dy};
+}
+
+/* Subtraction of duals. */
+template<class T>
+ccl_device_inline dual<T> operator-(const ccl_private dual<T> &a, const ccl_private dual<T> &b)
+{
+  return {a.val - b.val, a.dx - b.dx, a.dy - b.dy};
+}
+
 /* Negation. */
 template<class T> ccl_device_inline dual<T> operator-(const ccl_private dual<T> &a)
 {
   return {-a.val, -a.dx, -a.dy};
 }
 
+/* dfdx = dfdu * dudx */
+template<class T>
+ccl_device_inline dual<T> chain_rule(const ccl_private dual<T> &u,
+                                     const ccl_private T &f,
+                                     const ccl_private T &dfdu)
+{
+  return {f, dfdu * u.dx, dfdu * u.dy};
+}
+
+/* dfdx = dfdu * dudx + dfdv * dvdx. */
+template<class T>
+ccl_device_inline dual<T> chain_rule(const ccl_private dual<T> &u,
+                                     const ccl_private dual<T> &v,
+                                     const ccl_private T &f,
+                                     const ccl_private T &dfdu,
+                                     const ccl_private T &dfdv)
+{
+  return {f, dfdu * u.dx + dfdv * v.dx, dfdu * u.dy + dfdv * v.dy};
+}
+
+template<class MaskType>
+ccl_device_inline dual3 select(const MaskType mask, const dual3 a, const dual3 b)
+{
+#if defined(__KERNEL_METAL__)
+  const bool3 mask_ = bool3(mask);
+  return {metal::select(b.val, a.val, mask_),
+          metal::select(b.dx, a.dx, mask_),
+          metal::select(b.dy, a.dy, mask_)};
+#elif defined(__KERNEL_SSE__)
+#  ifdef __KERNEL_SSE42__
+  const auto mask_ = _mm_castsi128_ps(mask.m128);
+  return {float3(_mm_blendv_ps(b.val.m128, a.val.m128, mask_)),
+          float3(_mm_blendv_ps(b.dx.m128, a.dx.m128, mask_)),
+          float3(_mm_blendv_ps(b.dy.m128, a.dy.m128, mask_))};
+#  else
+  const auto mask_ = _mm_castsi128_ps(mask);
+  return {float3(_mm_or_ps(_mm_and_ps(mask_, a.val), _mm_andnot_ps(mask_, b.val))),
+          float3(_mm_or_ps(_mm_and_ps(mask_, a.dx), _mm_andnot_ps(mask_, b.dx))),
+          float3(_mm_or_ps(_mm_and_ps(mask_, a.dy), _mm_andnot_ps(mask_, b.dy)))};
+#  endif
+#else
+  return make_float3(mask.x ? a.x() : b.x(), mask.y ? a.y() : b.y(), mask.z ? a.z() : b.z());
+#endif
+}
+
+/* Functions with zero derivatives. */
+template<class T> ccl_device_inline dual<T> floor(const ccl_private dual<T> &a)
+{
+  return dual<T>(floor(a.val));
+}
+
+template<class T> ccl_device_inline dual<T> ceil(const ccl_private dual<T> &a)
+{
+  return dual<T>(ceil(a.val));
+}
+
+template<class T> ccl_device_inline dual<T> compatible_sign(const ccl_private dual<T> &u)
+{
+  return dual<T>(compatible_sign(u.val));
+}
+
+/* f = u - round(u / v) * v, f' = u'. */
+ccl_device_inline dual3 safe_fmod(const dual3 u, const dual3 v)
+{
+  return {safe_fmod(u.val, v.val), u.dx, u.dy};
+}
+
+ccl_device_inline dual3 safe_floored_fmod(const dual3 a, const dual3 b)
+{
+  return select(component_is_zero(b.val), make_zero<dual3>(), a - floor(a.val / b.val) * b);
+}
+
+template<class T> ccl_device_inline dual<T> safe_divide(const dual<T> f, const T g)
+{
+  return select(component_is_zero(g), make_zero<dual<T>>(), f / g);
+}
+
+template<class T>
+ccl_device_inline dual<T> safe_divide(const ccl_private dual<T> &f, const ccl_private dual<T> &g)
+{
+  return select(component_is_zero(g.val), make_zero<dual<T>>(), f / g);
+}
+
+/* Adapted from GODOT-engine math_funcs.h. */
+ccl_device_inline dual3 wrap(const dual3 value, const dual3 max, const dual3 min)
+{
+  return safe_floored_fmod(value - min, max - min) + min;
+}
+
+ccl_device_inline dual3 min(const ccl_private dual3 &a, const ccl_private dual3 &b)
+{
+  return select(a.val < b.val, a, b);
+}
+
+ccl_device_inline dual3 max(const ccl_private dual3 &a, const ccl_private dual3 &b)
+{
+  return select(a.val > b.val, a, b);
+}
+
+ccl_device_inline dual1 max(const ccl_private dual1 &a, const ccl_private dual1 &b)
+{
+  return a.val > b.val ? a : b;
+}
+
+ccl_device_inline dual3 fabs(const ccl_private dual3 &a)
+{
+  return select(a.val > zero_float3(), a, -a);
+}
+
 template<class T> ccl_device_inline dual1 average(const dual<T> a)
 {
   return {average(a.val), average(a.dx), average(a.dy)};
@@ -51,9 +247,185 @@ template<class T> ccl_device_inline dual1 reduce_add(const dual<T> a)
   return {reduce_add(a.val), reduce_add(a.dx), reduce_add(a.dy)};
 }
 
+/* f(u) = sqrt(u), dfdu = 1 / (2 * sqrt(u)). */
+ccl_device_inline dual1 sqrt(const ccl_private dual1 &u)
+{
+  const float f = sqrtf(u.val);
+  return chain_rule(u, f, 0.5f / f);
+}
+
+template<class T> ccl_device_inline dual1 len(const ccl_private dual<T> &a)
+{
+  return sqrt(dot(a, a));
+}
+
 template<class T1, class T2> ccl_device_inline dual1 dot(const dual<T1> a, const T2 b)
 {
   return reduce_add(a * b);
 }
 
+template<class T> ccl_device_inline dual1 len_squared(const ccl_private dual<T> &a)
+{
+  return dot(a, a);
+}
+
+template<class T>
+ccl_device_inline dual1 distance(const ccl_private dual<T> &a, const ccl_private dual<T> &b)
+{
+  return len(a - b);
+}
+
+ccl_device_inline dual3 cross(const ccl_private dual3 &a, const ccl_private dual3 &b)
+{
+  return {cross(a.val, b.val),
+          cross(a.val, b.dx) + cross(a.dx, b.val),
+          cross(a.val, b.dy) + cross(a.dy, b.val)};
+}
+
+ccl_device_inline dual3 cross(const ccl_private dual3 &a, const ccl_private float3 &b)
+{
+  return {cross(a.val, b), cross(a.dx, b), cross(a.dy, b)};
+}
+
+ccl_device_inline dual3 cross(const ccl_private float3 &a, const ccl_private dual3 &b)
+{
+  return -cross(b, a);
+}
+
+/* f(u) = 1 / sqrt(u), dfdu = -1 / (2 * u^(3/2)). */
+ccl_device_inline dual1 inversesqrt(const ccl_private dual1 &u)
+{
+  const float f = inversesqrtf(u.val);
+  return chain_rule(u, f, -0.5f * safe_divide(f, u.val));
+}
+
+template<class T> ccl_device_inline dual<T> normalize(const ccl_private dual<T> &a)
+{
+  return a * inversesqrt(len_squared(a));
+}
+
+template<class T> ccl_device_inline dual<T> safe_normalize(const ccl_private dual<T> &a)
+{
+  const dual1 len_sq = len_squared(a);
+  return is_zero(len_sq) ? make_zero<dual<T>>() : a * inversesqrt(len_sq);
+}
+
+/* f(y, x) = atan2(y, x),
+ * dfdx = -y / (x^2 + y^2),
+ * dfdy = x / (x^2 + y^2) */
+ccl_device_inline dual1 atan2(const ccl_private dual1 &y, const ccl_private dual1 &x)
+{
+  const float inv_len = safe_divide(1.0f, sqr(x.val) + sqr(y.val));
+  const float dfdx = -y.val * inv_len;
+  const float dfdy = x.val * inv_len;
+  return chain_rule(x, y, atan2f(y.val, x.val), dfdx, dfdy);
+}
+
+/* f(u) = acos(u), dfdu = -1 / sqrt(1 - u^2). */
+ccl_device_inline dual1 acos(const ccl_private dual1 &u)
+{
+  return chain_rule(u, acosf(u.val), -inversesqrtf(1.0f - sqr(u.val)));
+}
+
+ccl_device_inline dual1 safe_acos(const ccl_private dual1 &u)
+{
+  const float dfdu = (fabsf(u.val) >= 1.0f) ? 0.0f : -inversesqrtf(1.0f - sqr(u.val));
+  return chain_rule(u, safe_acosf(u.val), dfdu);
+}
+
+template<class T> ccl_device_inline dual3 reflect(const dual3 incident, const T unit_normal)
+{
+  return incident - unit_normal * make_float3(dot(incident, unit_normal)) * 2.0f;
+}
+
+ccl_device_inline dual3 refract(const dual3 incident, const dual3 normal, const dual1 eta)
+{
+  const dual1 NI = dot(incident, normal);
+  const dual1 k = 1.0f - eta * eta * (1.0f - NI * NI);
+  if (k.val < 0.0f) {
+    return dual3();
+  }
+  return incident * eta - normal * (eta * NI + sqrt(k));
+}
+
+ccl_device_inline dual3 faceforward(const dual3 vector,
+                                    const dual3 incident,
+                                    const dual3 reference)
+{
+  return (dot(reference, incident) < 0.0f) ? vector : -vector;
+}
+
+ccl_device_inline dual3 project(const dual3 v, const dual3 v_proj)
+{
+  const dual1 len_squared = dot(v_proj, v_proj);
+  return (len_squared.val != 0.0f) ? v_proj * (dot(v, v_proj) / len_squared) : dual3();
+}
+
+template<class T> ccl_device_inline dual<T> sin(const ccl_private dual<T> &x)
+{
+  T sinx, cosx;
+  sincos(x.val, &sinx, &cosx);
+  return chain_rule(x, sinx, cosx);
+}
+
+template<class T> ccl_device_inline dual<T> cos(const ccl_private dual<T> &x)
+{
+  T sinx, cosx;
+  sincos(x.val, &sinx, &cosx);
+  return chain_rule(x, cosx, -sinx);
+}
+
+ccl_device_inline dual3 tan(const ccl_private dual3 &x)
+{
+  const float3 tanx = tan(x.val);
+  const float3 secx = safe_divide(one_float3(), cos(x.val));
+  return chain_rule(x, tanx, sqr(secx));
+}
+
+/* f(u, v) = u^v, dfdu = v u^(v-1), dfdv = u^v ln(u). */
+template<class T>
+ccl_device_inline dual<T> safe_pow(const ccl_private dual<T> &u, const ccl_private dual<T> &v)
+{
+  /* u^(v-1). */
+  const T u_v_minus_1 = safe_pow(u.val, v.val - 1.0f);
+  /* u^v = u * u^(v-1). */
+  /* NOTE: numerically `u^v != u*u^(v-1)`, but the current behaviour matches OSL. */
+  const T f = u.val * u_v_minus_1;
+  return chain_rule(u, v, f, v.val * u_v_minus_1, f * safe_log(u.val));
+}
+
+/* Projections. */
+ccl_device_inline dual2 map_to_tube(const dual3 co)
+{
+  dual1 u, v;
+  const dual1 length = len(make_float2(co));
+  if (length.val > 0.0f) {
+    u = (1.0f - (atan2(co.x(), co.y()) / M_PI_F)) * 0.5f;
+    v = (co.z() + 1.0f) * 0.5f;
+  }
+  else {
+    u = v = make_zero<dual1>();
+  }
+  return make_float2(u, v);
+}
+
+ccl_device_inline dual2 map_to_sphere(const dual3 co)
+{
+  const dual1 l = dot(co, co);
+  dual1 u, v;
+  if (l.val > 0.0f) {
+    if (UNLIKELY(co.val.x == 0.0f && co.val.y == 0.0f)) {
+      u = make_zero<dual1>(); /* Otherwise domain error. */
+    }
+    else {
+      u = (0.5f - atan2(co.x(), co.y()) * M_1_2PI_F);
+    }
+    v = 1.0f - safe_acos(co.z() * inversesqrt(l)) * M_1_PI_F;
+  }
+  else {
+    u = v = make_zero<dual1>();
+  }
+  return make_float2(u, v);
+}
+
 CCL_NAMESPACE_END

intern/cycles/util/math_float3.h

@@ -621,7 +621,7 @@ ccl_device_inline float3 select(const MaskType mask, const float3 a, const float
 #  ifdef __KERNEL_SSE42__
   return float3(_mm_blendv_ps(b.m128, a.m128, _mm_castsi128_ps(mask.m128)));
 #  else
-  return float4(
+  return float3(
       _mm_or_ps(_mm_and_ps(_mm_castsi128_ps(mask), a), _mm_andnot_ps(_mm_castsi128_ps(mask), b)));
 #  endif
 #else
@@ -646,7 +646,22 @@ ccl_device_inline float3 safe_pow(const float3 a, const float3 b)
   return make_float3(safe_powf(a.x, b.x), safe_powf(a.y, b.y), safe_powf(a.z, b.z));
 }
 
-ccl_device_inline auto isequal_mask(const float3 a, const float3 b)
+ccl_device_inline float3 safe_log(const float3 v)
+{
+  return select(v > zero_float3(), log(v), zero_float3());
+}
+
+ccl_device_inline void sincos(const float3 x, ccl_private float3 *sine, ccl_private float3 *cosine)
+{
+#if defined(__KERNEL_METAL__)
+  *sine = sincos(x, *cosine);
+#else
+  *sine = sin(x);
+  *cosine = cos(x);
+#endif
+}
+
+ccl_device_inline auto component_wise_equal(const float3 a, const float3 b)
 {
 #if defined(__KERNEL_METAL__)
   return a == b;
@@ -659,14 +674,14 @@ ccl_device_inline auto isequal_mask(const float3 a, const float3 b)
 #endif
 }
 
-ccl_device_inline auto is_zero_mask(const float3 a)
+ccl_device_inline auto component_is_zero(const float3 a)
 {
-  return isequal_mask(a, zero_float3());
+  return component_wise_equal(a, zero_float3());
 }
 
 ccl_device_inline float3 safe_floored_fmod(const float3 a, const float3 b)
 {
-  return select(is_zero_mask(b), zero_float3(), a - floor(a / b) * b);
+  return select(component_is_zero(b), zero_float3(), a - floor(a / b) * b);
 }
 
 ccl_device_inline float3 wrap(const float3 value, const float3 max, const float3 min)
@@ -676,7 +691,7 @@ ccl_device_inline float3 wrap(const float3 value, const float3 max, const float3
 
 ccl_device_inline float3 safe_fmod(const float3 a, const float3 b)
 {
-  return select(is_zero_mask(b), zero_float3(), fmod(a, b));
+  return select(component_is_zero(b), zero_float3(), fmod(a, b));
 }
 
 ccl_device_inline float3 compatible_sign(const float3 v)

intern/cycles/util/projection.h

@@ -32,30 +32,6 @@ ccl_device_inline float2 polar_to_cartesian(const float r, const float phi)
   return make_float2(r * cosf(phi), r * sinf(phi));
 }
 
-/* Transform p from a local coordinate system (spanned by X and Y) into global coordinates. */
-template<class T> ccl_device_inline T to_global(const float2 p, const T X, const T Y)
-{
-  return p.x * X + p.y * Y;
-}
-
-/* Transform p from a local coordinate system (spanned by X, Y and Z) into global coordinates. */
-template<class T> ccl_device_inline T to_global(const float3 p, const T X, const T Y, const T Z)
-{
-  return p.x * X + p.y * Y + p.z * Z;
-}
-
-/* Transform p from global coordinates into a local coordinate system (spanned by X and Y). */
-template<class T> ccl_device_inline float2 to_local(const T p, const T X, const T Y)
-{
-  return make_float2(dot(p, X), dot(p, Y));
-}
-
-/* Transform p from global coordinates into a local coordinate system (spanned by X, Y and Z). */
-template<class T> ccl_device_inline float3 to_local(const T p, const T X, const T Y, const T Z)
-{
-  return make_float3(dot(p, X), dot(p, Y), dot(p, Z));
-}
-
 ccl_device_inline float3 disk_to_hemisphere(const float2 p)
 {
   return make_float3(p.x, p.y, safe_sqrtf(1.0f - len_squared(p)));

intern/cycles/util/transform.h

@@ -17,6 +17,36 @@
 
 CCL_NAMESPACE_BEGIN
 
+/* Transform p from a local coordinate system (spanned by X and Y) into global coordinates. */
+template<class T> ccl_device_inline T to_global(const float2 p, const T X, const T Y)
+{
+  return p.x * X + p.y * Y;
+}
+
+/* Transform p from a local coordinate system (spanned by X, Y and Z) into global coordinates. */
+template<class T> ccl_device_inline T to_global(const float3 p, const T X, const T Y, const T Z)
+{
+  return p.x * X + p.y * Y + p.z * Z;
+}
+
+/* Transform p from global coordinates into a local coordinate system (spanned by X and Y). */
+template<class T> ccl_device_inline float2 to_local(const T p, const T X, const T Y)
+{
+  return make_float2(dot(p, X), dot(p, Y));
+}
+
+/* Transform p from global coordinates into a local coordinate system (spanned by X, Y and Z). */
+template<class T> ccl_device_inline float3 to_local(const T p, const T X, const T Y, const T Z)
+{
+  return make_float3(dot(p, X), dot(p, Y), dot(p, Z));
+}
+
+template<class T>
+ccl_device_inline dual3 to_local(const dual<T> p, const T X, const T Y, const T Z)
+{
+  return make_float3(dot(p, X), dot(p, Y), dot(p, Z));
+}
+
 /* Affine transformation, stored as 4x3 matrix. */
 
 struct Transform {
@@ -146,14 +176,19 @@ ccl_device_inline float3 transform_direction(const ccl_private Transform *t, con
 #endif
 }
 
-ccl_device_inline float3 transform_direction_transposed(const ccl_private Transform *t,
-                                                        const float3 a)
+ccl_device_inline dual3 transform_direction(const ccl_private Transform *t, const dual3 a)
+{
+  return to_local(a, make_float3(t->x), make_float3(t->y), make_float3(t->z));
+}
+
+template<class T>
+ccl_device_inline T transform_direction_transposed(const ccl_private Transform *t, const T a)
 {
   const float3 x = make_float3(t->x.x, t->y.x, t->z.x);
   const float3 y = make_float3(t->x.y, t->y.y, t->z.y);
   const float3 z = make_float3(t->x.z, t->y.z, t->z.z);
 
-  return make_float3(dot(x, a), dot(y, a), dot(z, a));
+  return to_local(a, x, y, z);
 }
 
 ccl_device_inline Transform make_transform(const float a,

intern/cycles/util/types_dual.h

@@ -17,6 +17,18 @@ template<class T> struct dual {
   ccl_device_inline_method dual(const T val, const T dx, const T dy) : val(val), dx(dx), dy(dy) {}
 };
 
+template<> struct dual<float> {
+  float val = 0.0f;
+  float dx = 0.0f;
+  float dy = 0.0f;
+  dual() = default;
+  ccl_device_inline_method explicit dual(const float val) : val(val) {}
+  ccl_device_inline_method dual(const float val, const float dx, const float dy)
+      : val(val), dx(dx), dy(dy)
+  {
+  }
+};
+
 template<> struct dual<float2> {
   float2 val = make_float2(0.0f);
   float2 dx = make_float2(0.0f);
@@ -27,6 +39,14 @@ template<> struct dual<float2> {
       : val(val), dx(dx), dy(dy)
   {
   }
+  ccl_device_inline_method dual<float> x() const
+  {
+    return {val.x, dx.x, dy.x};
+  }
+  ccl_device_inline_method dual<float> y() const
+  {
+    return {val.y, dx.y, dy.y};
+  }
 };
 
 template<> struct dual<float3> {
@@ -39,6 +59,18 @@ template<> struct dual<float3> {
       : val(val), dx(dx), dy(dy)
   {
   }
+  ccl_device_inline_method dual<float> x() const
+  {
+    return {val.x, dx.x, dy.x};
+  }
+  ccl_device_inline_method dual<float> y() const
+  {
+    return {val.y, dx.y, dy.y};
+  }
+  ccl_device_inline_method dual<float> z() const
+  {
+    return {val.z, dx.z, dy.z};
+  }
 };
 
 template<> struct dual<float4> {
@@ -51,6 +83,22 @@ template<> struct dual<float4> {
       : val(val), dx(dx), dy(dy)
   {
   }
+  ccl_device_inline_method dual<float> x() const
+  {
+    return {val.x, dx.x, dy.x};
+  }
+  ccl_device_inline_method dual<float> y() const
+  {
+    return {val.y, dx.y, dy.y};
+  }
+  ccl_device_inline_method dual<float> z() const
+  {
+    return {val.z, dx.z, dy.z};
+  }
+  ccl_device_inline_method dual<float> w() const
+  {
+    return {val.w, dx.w, dy.w};
+  }
 };
 
 using dual1 = dual<float>;
@@ -58,9 +106,48 @@ using dual2 = dual<float2>;
 using dual3 = dual<float3>;
 using dual4 = dual<float4>;
 
-template<class T> ccl_device_inline dual3 make_float3(const ccl_private dual<T> &a)
+/* Dual type traits. */
+
+template<typename T> struct is_dual {
+  ccl_static_constexpr bool value = false;
+};
+template<typename U> struct is_dual<dual<U>> {
+  ccl_static_constexpr bool value = true;
+};
+
+template<typename T> ccl_static_constexpr bool is_dual_v = is_dual<T>::value;
+
+/* Base (non-dual) type. E.g. dual_base_t<dual3> = float3, dual_base_t<float3> = float3. */
+
+template<typename T> struct dual_base_type {
+  using type = T;
+};
+template<typename U> struct dual_base_type<dual<U>> {
+  using type = U;
+};
+template<typename T> using dual_base_t = typename dual_base_type<T>::type;
+
+/* Scalar type corresponding to a vector type. */
+
+template<typename T> struct dual_scalar_type {
+  using type = T;
+};
+template<> struct dual_scalar_type<float3> {
+  using type = float;
+};
+template<> struct dual_scalar_type<dual3> {
+  using type = dual1;
+};
+template<typename T> using dual_scalar_t = typename dual_scalar_type<T>::type;
+
+ccl_device_inline dual2 make_float2(const dual3 a)
 {
-  return {make_float3(a.val), make_float3(a.dx), make_float3(a.dy)};
+  return {make_float2(a.val), make_float2(a.dx), make_float2(a.dy)};
+}
+
+ccl_device_inline dual2 make_float2(const dual1 a, const dual1 b)
+{
+  return {make_float2(a.val, b.val), make_float2(a.dx, b.dx), make_float2(a.dy, b.dy)};
 }
 
 ccl_device_inline dual3 make_float3(const dual1 a, const dual1 b, const dual1 c)
@@ -70,6 +157,18 @@ ccl_device_inline dual3 make_float3(const dual1 a, const dual1 b, const dual1 c)
           make_float3(a.dy, b.dy, c.dy)};
 }
 
+template<class T> ccl_device_inline dual3 make_float3(const ccl_private dual<T> &a)
+{
+  return {make_float3(a.val), make_float3(a.dx), make_float3(a.dy)};
+}
+
+ccl_device_inline dual4 make_float4(const dual1 a, const dual1 b, const dual1 c, const dual1 d)
+{
+  return {make_float4(a.val, b.val, c.val, d.val),
+          make_float4(a.dx, b.dx, c.dx, d.dx),
+          make_float4(a.dy, b.dy, c.dy, d.dy)};
+}
+
 ccl_device_inline dual4 make_float4(const dual3 a)
 {
   return {make_float4(a.val), make_float4(a.dx, 0.0f), make_float4(a.dy, 0.0f)};