Merge branch 'master' into retopo_transform

build_files/build_environment/install_deps.sh

@@ -478,7 +478,7 @@ OCIO_FORCE_BUILD=false
 OCIO_FORCE_REBUILD=false
 OCIO_SKIP=false
 
-IMATH_VERSION="3.1.4"
+IMATH_VERSION="3.1.5"
 IMATH_VERSION_SHORT="3.1"
 IMATH_VERSION_MIN="3.0"
 IMATH_VERSION_MEX="4.0"
@@ -487,7 +487,7 @@ IMATH_FORCE_REBUILD=false
 IMATH_SKIP=false
 _with_built_imath=false
 
-OPENEXR_VERSION="3.1.4"
+OPENEXR_VERSION="3.1.5"
 OPENEXR_VERSION_SHORT="3.1"
 OPENEXR_VERSION_MIN="3.0"
 OPENEXR_VERSION_MEX="4.0"
@@ -627,6 +627,9 @@ WEBP_DEV=""
 VPX_USE=false
 VPX_VERSION_MIN=0.9.7
 VPX_DEV=""
+AOM_USE=false
+AOM_VERSION_MIN=3.3.0
+AOM_DEV=""
 OPUS_USE=false
 OPUS_VERSION_MIN=1.1.1
 OPUS_DEV=""
@@ -1209,7 +1212,7 @@ You may also want to build them yourself (optional ones are [between brackets]):
     ** [NumPy $PYTHON_NUMPY_VERSION] (use pip).
     * Boost $BOOST_VERSION (from $BOOST_SOURCE, modules: $BOOST_BUILD_MODULES).
     * TBB $TBB_VERSION (from $TBB_SOURCE).
-    * [FFMpeg $FFMPEG_VERSION (needs libvorbis, libogg, libtheora, libx264, libmp3lame, libxvidcore, libvpx, libwebp, ...)] (from $FFMPEG_SOURCE).
+    * [FFMpeg $FFMPEG_VERSION (needs libvorbis, libogg, libtheora, libx264, libmp3lame, libxvidcore, libvpx, libaom, libwebp, ...)] (from $FFMPEG_SOURCE).
     * [OpenColorIO $OCIO_VERSION] (from $OCIO_SOURCE).
     * Imath $IMATH_VERSION (from $IMATH_SOURCE).
     * OpenEXR $OPENEXR_VERSION (from $OPENEXR_SOURCE).
@@ -3000,7 +3003,7 @@ compile_ALEMBIC() {
   fi
 
   # To be changed each time we make edits that would modify the compiled result!
-  alembic_magic=2
+  alembic_magic=3
   _init_alembic
 
   # Force having own builds for the dependencies.
@@ -3048,7 +3051,7 @@ compile_ALEMBIC() {
     fi
     if [ "$_with_built_openexr" = true ]; then
       cmake_d="$cmake_d -D USE_ARNOLD=OFF"
-      cmake_d="$cmake_d -D USE_BINARIES=OFF"
+      cmake_d="$cmake_d -D USE_BINARIES=ON"  # Tests use some Alembic binaries...
       cmake_d="$cmake_d -D USE_EXAMPLES=OFF"
       cmake_d="$cmake_d -D USE_HDF5=OFF"
       cmake_d="$cmake_d -D USE_MAYA=OFF"
@@ -3634,7 +3637,7 @@ compile_FFmpeg() {
   fi
 
   # To be changed each time we make edits that would modify the compiled result!
-  ffmpeg_magic=5
+  ffmpeg_magic=6
   _init_ffmpeg
 
   # Force having own builds for the dependencies.
@@ -3687,6 +3690,10 @@ compile_FFmpeg() {
       extra="$extra --enable-libvpx"
     fi
 
+    if [ "$AOM_USE" = true ]; then
+      extra="$extra --enable-libaom"
+    fi
+
     if [ "$WEBP_USE" = true ]; then
       extra="$extra --enable-libwebp"
     fi
@@ -4140,30 +4147,34 @@ install_DEB() {
     WEBP_USE=true
   fi
 
-  if [ "$WITH_ALL" = true ]; then
-    XVID_DEV="libxvidcore-dev"
-    check_package_DEB $XVID_DEV
-    if [ $? -eq 0 ]; then
-      XVID_USE=true
-    fi
+  XVID_DEV="libxvidcore-dev"
+  check_package_DEB $XVID_DEV
+  if [ $? -eq 0 ]; then
+    XVID_USE=true
+  fi
 
-    MP3LAME_DEV="libmp3lame-dev"
-    check_package_DEB $MP3LAME_DEV
-    if [ $? -eq 0 ]; then
-      MP3LAME_USE=true
-    fi
+  MP3LAME_DEV="libmp3lame-dev"
+  check_package_DEB $MP3LAME_DEV
+  if [ $? -eq 0 ]; then
+    MP3LAME_USE=true
+  fi
 
-    VPX_DEV="libvpx-dev"
-    check_package_version_ge_DEB $VPX_DEV $VPX_VERSION_MIN
-    if [ $? -eq 0 ]; then
-      VPX_USE=true
-    fi
+  VPX_DEV="libvpx-dev"
+  check_package_version_ge_DEB $VPX_DEV $VPX_VERSION_MIN
+  if [ $? -eq 0 ]; then
+    VPX_USE=true
+  fi
 
-    OPUS_DEV="libopus-dev"
-    check_package_version_ge_DEB $OPUS_DEV $OPUS_VERSION_MIN
-    if [ $? -eq 0 ]; then
-      OPUS_USE=true
-    fi
+  AOM_DEV="libaom-dev"
+  check_package_version_ge_DEB $AOM_DEV $AOM_VERSION_MIN
+  if [ $? -eq 0 ]; then
+    AOM_USE=true
+  fi
+
+  OPUS_DEV="libopus-dev"
+  check_package_version_ge_DEB $OPUS_DEV $OPUS_VERSION_MIN
+  if [ $? -eq 0 ]; then
+    OPUS_USE=true
   fi
 
   # Check cmake version and disable features for older distros.
@@ -4546,6 +4557,9 @@ install_DEB() {
     if [ "$VPX_USE" = true ]; then
       _packages="$_packages $VPX_DEV"
     fi
+    if [ "$AOM_USE" = true ]; then
+      _packages="$_packages $AOM_DEV"
+    fi
     if [ "$OPUS_USE" = true ]; then
       _packages="$_packages $OPUS_DEV"
     fi
@@ -4846,21 +4860,27 @@ install_RPM() {
     WEBP_USE=true
   fi
 
-  if [ "$WITH_ALL" = true ]; then
-    VPX_DEV="libvpx-devel"
-    check_package_version_ge_RPM $VPX_DEV $VPX_VERSION_MIN
-    if [ $? -eq 0 ]; then
-      VPX_USE=true
-    fi
+  VPX_DEV="libvpx-devel"
+  check_package_version_ge_RPM $VPX_DEV $VPX_VERSION_MIN
+  if [ $? -eq 0 ]; then
+    VPX_USE=true
+  fi
 
+  AOM_DEV="libaom-devel"
+  check_package_version_ge_RPM $AOM_DEV $AOM_VERSION_MIN
+  if [ $? -eq 0 ]; then
+    AOM_USE=true
+  fi
+
+  OPUS_DEV="libopus-devel"
+  check_package_version_ge_RPM $OPUS_DEV $OPUS_VERSION_MIN
+  if [ $? -eq 0 ]; then
+    OPUS_USE=true
+  fi
+
+  if [ "$WITH_ALL" = true ]; then
     PRINT ""
     install_packages_RPM libspnav-devel
-
-    OPUS_DEV="libopus-devel"
-    check_package_version_ge_RPM $OPUS_DEV $OPUS_VERSION_MIN
-    if [ $? -eq 0 ]; then
-      OPUS_USE=true
-    fi
   fi
 
   PRINT ""
@@ -5245,6 +5265,9 @@ install_RPM() {
     if [ "$VPX_USE" = true ]; then
       _packages="$_packages $VPX_DEV"
     fi
+    if [ "$AOM_USE" = true ]; then
+      _packages="$_packages $AOM_DEV"
+    fi
     if [ "$OPUS_USE" = true ]; then
       _packages="$_packages $OPUS_DEV"
     fi
@@ -5434,30 +5457,34 @@ install_ARCH() {
     WEBP_USE=true
   fi
 
-  if [ "$WITH_ALL" = true ]; then
-    XVID_DEV="xvidcore"
-    check_package_ARCH $XVID_DEV
-    if [ $? -eq 0 ]; then
-      XVID_USE=true
-    fi
+  XVID_DEV="xvidcore"
+  check_package_ARCH $XVID_DEV
+  if [ $? -eq 0 ]; then
+    XVID_USE=true
+  fi
 
-    MP3LAME_DEV="lame"
-    check_package_ARCH $MP3LAME_DEV
-    if [ $? -eq 0 ]; then
-      MP3LAME_USE=true
-    fi
+  MP3LAME_DEV="lame"
+  check_package_ARCH $MP3LAME_DEV
+  if [ $? -eq 0 ]; then
+    MP3LAME_USE=true
+  fi
 
-    VPX_DEV="libvpx"
-    check_package_version_ge_ARCH $VPX_DEV $VPX_VERSION_MIN
-    if [ $? -eq 0 ]; then
-      VPX_USE=true
-    fi
+  VPX_DEV="libvpx"
+  check_package_version_ge_ARCH $VPX_DEV $VPX_VERSION_MIN
+  if [ $? -eq 0 ]; then
+    VPX_USE=true
+  fi
 
-    OPUS_DEV="opus"
-    check_package_version_ge_ARCH $OPUS_DEV $OPUS_VERSION_MIN
-    if [ $? -eq 0 ]; then
-      OPUS_USE=true
-    fi
+  AOM_DEV="libaom"
+  check_package_version_ge_ARCH $AOM_DEV $AOM_VERSION_MIN
+  if [ $? -eq 0 ]; then
+    AOM_USE=true
+  fi
+
+  OPUS_DEV="opus"
+  check_package_version_ge_ARCH $OPUS_DEV $OPUS_VERSION_MIN
+  if [ $? -eq 0 ]; then
+    OPUS_USE=true
   fi
 
 
@@ -5835,6 +5862,9 @@ install_ARCH() {
     if [ "$VPX_USE" = true ]; then
       _packages="$_packages $VPX_DEV"
     fi
+    if [ "$AOM_USE" = true ]; then
+      _packages="$_packages $AOM_DEV"
+    fi
     if [ "$OPUS_USE" = true ]; then
       _packages="$_packages $OPUS_DEV"
     fi

build_files/cmake/platform/platform_apple.cmake

@@ -470,8 +470,9 @@ string(APPEND CMAKE_CXX_FLAGS " -ftemplate-depth=1024")
 
 # Avoid conflicts with Luxrender, and other plug-ins that may use the same
 # libraries as Blender with a different version or build options.
+set(PLATFORM_SYMBOLS_MAP ${CMAKE_SOURCE_DIR}/source/creator/symbols_apple.map)
 string(APPEND PLATFORM_LINKFLAGS
-  " -Wl,-unexported_symbols_list,'${CMAKE_SOURCE_DIR}/source/creator/osx_locals.map'"
+  " -Wl,-unexported_symbols_list,'${PLATFORM_SYMBOLS_MAP}'"
 )
 
 string(APPEND CMAKE_CXX_FLAGS " -stdlib=libc++")

build_files/cmake/platform/platform_unix.cmake

@@ -888,8 +888,9 @@ unset(_IS_LINKER_DEFAULT)
 
 # Avoid conflicts with Mesa llvmpipe, Luxrender, and other plug-ins that may
 # use the same libraries as Blender with a different version or build options.
+set(PLATFORM_SYMBOLS_MAP ${CMAKE_SOURCE_DIR}/source/creator/symbols_unix.map)
 set(PLATFORM_LINKFLAGS
-  "${PLATFORM_LINKFLAGS} -Wl,--version-script='${CMAKE_SOURCE_DIR}/source/creator/blender.map'"
+  "${PLATFORM_LINKFLAGS} -Wl,--version-script='${PLATFORM_SYMBOLS_MAP}'"
 )
 
 # Don't use position independent executable for portable install since file

intern/cycles/blender/sync.h

@@ -7,6 +7,7 @@
 #include "MEM_guardedalloc.h"
 #include "RNA_access.h"
 #include "RNA_blender_cpp.h"
+#include "RNA_path.h"
 #include "RNA_types.h"
 
 #include "blender/id_map.h"

intern/cycles/kernel/CMakeLists.txt

@@ -355,8 +355,6 @@ set(SRC_UTIL_HEADERS
   ../util/types_uint4.h
   ../util/types_uint4_impl.h
   ../util/types_ushort4.h
-  ../util/types_vector3.h
-  ../util/types_vector3_impl.h
 )
 
 set(LIB

intern/cycles/kernel/bvh/bvh.h

@@ -12,6 +12,7 @@
 
 #if defined(__EMBREE__)
 #  include "kernel/device/cpu/bvh.h"
+#  define __BVH2__
 #elif defined(__METALRT__)
 #  include "kernel/device/metal/bvh.h"
 #elif defined(__KERNEL_OPTIX__)
@@ -72,6 +73,12 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
     return false;
   }
 
+#  ifdef __EMBREE__
+  if (kernel_data.device_bvh) {
+    return kernel_embree_intersect(kg, ray, visibility, isect);
+  }
+#  endif
+
 #  ifdef __OBJECT_MOTION__
   if (kernel_data.bvh.have_motion) {
 #    ifdef __HAIR__
@@ -121,6 +128,12 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
     return false;
   }
 
+#    ifdef __EMBREE__
+  if (kernel_data.device_bvh) {
+    return kernel_embree_intersect_local(kg, ray, local_isect, local_object, lcg_state, max_hits);
+  }
+#    endif
+
 #    ifdef __OBJECT_MOTION__
   if (kernel_data.bvh.have_motion) {
     return bvh_intersect_local_motion(kg, ray, local_isect, local_object, lcg_state, max_hits);
@@ -170,6 +183,13 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
     return false;
   }
 
+#    ifdef __EMBREE__
+  if (kernel_data.device_bvh) {
+    return kernel_embree_intersect_shadow_all(
+        kg, state, ray, visibility, max_hits, num_recorded_hits, throughput);
+  }
+#    endif
+
 #    ifdef __OBJECT_MOTION__
   if (kernel_data.bvh.have_motion) {
 #      ifdef __HAIR__
@@ -254,6 +274,12 @@ ccl_device_intersect uint scene_intersect_volume(KernelGlobals kg,
     return false;
   }
 
+#    ifdef __EMBREE__
+  if (kernel_data.device_bvh) {
+    return kernel_embree_intersect_volume(kg, ray, isect, max_hits, visibility);
+  }
+#    endif
+
 #    ifdef __OBJECT_MOTION__
   if (kernel_data.bvh.have_motion) {
     return bvh_intersect_volume_all_motion(kg, ray, isect, max_hits, visibility);

intern/cycles/kernel/bvh/shadow_all.h

@@ -300,11 +300,7 @@ ccl_device_inline
       kernel_assert(object != OBJECT_NONE);
 
       /* Instance pop. */
-#if BVH_FEATURE(BVH_MOTION)
-      bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
-#else
-      bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
-#endif
+      bvh_instance_pop(ray, &P, &dir, &idir);
 
       object = OBJECT_NONE;
       node_addr = traversal_stack[stack_ptr];

intern/cycles/kernel/bvh/traversal.h

@@ -237,11 +237,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
       kernel_assert(object != OBJECT_NONE);
 
       /* instance pop */
-#if BVH_FEATURE(BVH_MOTION)
-      bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
-#else
-      bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
-#endif
+      bvh_instance_pop(ray, &P, &dir, &idir);
 
       object = OBJECT_NONE;
       node_addr = traversal_stack[stack_ptr];

intern/cycles/kernel/bvh/volume.h

@@ -210,11 +210,7 @@ ccl_device_inline
       kernel_assert(object != OBJECT_NONE);
 
       /* instance pop */
-#if BVH_FEATURE(BVH_MOTION)
-      bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
-#else
-      bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
-#endif
+      bvh_instance_pop(ray, &P, &dir, &idir);
 
       object = OBJECT_NONE;
       node_addr = traversal_stack[stack_ptr];

intern/cycles/kernel/bvh/volume_all.h

@@ -242,11 +242,7 @@ ccl_device_inline
       kernel_assert(object != OBJECT_NONE);
 
       /* Instance pop. */
-#if BVH_FEATURE(BVH_MOTION)
-      bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
-#else
-      bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
-#endif
+      bvh_instance_pop(ray, &P, &dir, &idir);
 
       object = OBJECT_NONE;
       node_addr = traversal_stack[stack_ptr];

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

@@ -166,16 +166,16 @@ ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg,
   }
   else {
     isect->type = kernel_data_fetch(objects, isect->object).primitive_type;
-    isect->u = 1.0f - hit->v - hit->u;
-    isect->v = hit->u;
+    isect->u = hit->u;
+    isect->v = hit->v;
   }
 }
 
 ccl_device_inline void kernel_embree_convert_sss_hit(
     KernelGlobals kg, const RTCRay *ray, const RTCHit *hit, Intersection *isect, int object)
 {
-  isect->u = 1.0f - hit->v - hit->u;
-  isect->v = hit->u;
+  isect->u = hit->u;
+  isect->v = hit->v;
   isect->t = ray->tfar;
   RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData(
       rtcGetGeometry(kernel_data.device_bvh, object * 2));
@@ -446,19 +446,11 @@ ccl_device void kernel_embree_filter_occluded_func_backface_cull(
 
 /* Scene intersection. */
 
-ccl_device_intersect bool scene_intersect(KernelGlobals kg,
-                                          ccl_private const Ray *ray,
-                                          const uint visibility,
-                                          ccl_private Intersection *isect)
+ccl_device_intersect bool kernel_embree_intersect(KernelGlobals kg,
+                                                  ccl_private const Ray *ray,
+                                                  const uint visibility,
+                                                  ccl_private Intersection *isect)
 {
-  if (!intersection_ray_valid(ray)) {
-    return false;
-  }
-
-  if (!kernel_data.device_bvh) {
-    return false;
-  }
-
   isect->t = ray->tmax;
   CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR);
   IntersectContext rtc_ctx(&ctx);
@@ -476,24 +468,13 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
 }
 
 #ifdef __BVH_LOCAL__
-ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
-                                                ccl_private const Ray *ray,
-                                                ccl_private LocalIntersection *local_isect,
-                                                int local_object,
-                                                ccl_private uint *lcg_state,
-                                                int max_hits)
+ccl_device_intersect bool kernel_embree_intersect_local(KernelGlobals kg,
+                                                        ccl_private const Ray *ray,
+                                                        ccl_private LocalIntersection *local_isect,
+                                                        int local_object,
+                                                        ccl_private uint *lcg_state,
+                                                        int max_hits)
 {
-  if (!intersection_ray_valid(ray)) {
-    if (local_isect) {
-      local_isect->num_hits = 0;
-    }
-    return false;
-  }
-
-  if (!kernel_data.device_bvh) {
-    return false;
-  }
-
   const bool has_bvh = !(kernel_data_fetch(object_flag, local_object) &
                          SD_OBJECT_TRANSFORM_APPLIED);
   CCLIntersectContext ctx(kg,
@@ -544,24 +525,14 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
 #endif
 
 #ifdef __SHADOW_RECORD_ALL__
-ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
-                                                     IntegratorShadowStateCPU *state,
-                                                     ccl_private const Ray *ray,
-                                                     uint visibility,
-                                                     uint max_hits,
-                                                     ccl_private uint *num_recorded_hits,
-                                                     ccl_private float *throughput)
+ccl_device_intersect bool kernel_embree_intersect_shadow_all(KernelGlobals kg,
+                                                             IntegratorShadowStateCPU *state,
+                                                             ccl_private const Ray *ray,
+                                                             uint visibility,
+                                                             uint max_hits,
+                                                             ccl_private uint *num_recorded_hits,
+                                                             ccl_private float *throughput)
 {
-  if (!intersection_ray_valid(ray)) {
-    *num_recorded_hits = 0;
-    *throughput = 1.0f;
-    return false;
-  }
-
-  if (!kernel_data.device_bvh) {
-    return false;
-  }
-
   CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_SHADOW_ALL);
   Intersection *isect_array = (Intersection *)state->shadow_isect;
   ctx.isect_s = isect_array;
@@ -579,20 +550,12 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
 #endif
 
 #ifdef __VOLUME__
-ccl_device_intersect uint scene_intersect_volume(KernelGlobals kg,
-                                                 ccl_private const Ray *ray,
-                                                 ccl_private Intersection *isect,
-                                                 const uint max_hits,
-                                                 const uint visibility)
+ccl_device_intersect uint kernel_embree_intersect_volume(KernelGlobals kg,
+                                                         ccl_private const Ray *ray,
+                                                         ccl_private Intersection *isect,
+                                                         const uint max_hits,
+                                                         const uint visibility)
 {
-  if (!intersection_ray_valid(ray)) {
-    return false;
-  }
-
-  if (!kernel_data.device_bvh) {
-    return false;
-  }
-
   CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_VOLUME_ALL);
   ctx.isect_s = isect;
   ctx.max_hits = max_hits;

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

@@ -33,38 +33,4 @@ CCL_NAMESPACE_BEGIN
 
 #define kernel_assert(cond) assert(cond)
 
-/* Macros to handle different memory storage on different devices */
-
-#ifdef __KERNEL_SSE2__
-typedef vector3<sseb> sse3b;
-typedef vector3<ssef> sse3f;
-typedef vector3<ssei> sse3i;
-
-ccl_device_inline void print_sse3b(const char *label, sse3b &a)
-{
-  print_sseb(label, a.x);
-  print_sseb(label, a.y);
-  print_sseb(label, a.z);
-}
-
-ccl_device_inline void print_sse3f(const char *label, sse3f &a)
-{
-  print_ssef(label, a.x);
-  print_ssef(label, a.y);
-  print_ssef(label, a.z);
-}
-
-ccl_device_inline void print_sse3i(const char *label, sse3i &a)
-{
-  print_ssei(label, a.x);
-  print_ssei(label, a.y);
-  print_ssei(label, a.z);
-}
-
-#  if defined(__KERNEL_AVX__) || defined(__KERNEL_AVX2__)
-typedef vector3<avxf> avx3f;
-#  endif
-
-#endif
-
 CCL_NAMESPACE_END

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

@@ -129,9 +129,8 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
 
   isect->t = intersection.distance;
   if (intersection.type == intersection_type::triangle) {
-    isect->u = 1.0f - intersection.triangle_barycentric_coord.y -
-               intersection.triangle_barycentric_coord.x;
-    isect->v = intersection.triangle_barycentric_coord.x;
+    isect->u = intersection.triangle_barycentric_coord.x;
+    isect->v = intersection.triangle_barycentric_coord.y;
   }
   else {
     isect->u = payload.u;
@@ -346,9 +345,8 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
 
   isect->t = intersection.distance;
   if (intersection.type == intersection_type::triangle) {
-    isect->u = 1.0f - intersection.triangle_barycentric_coord.y -
-               intersection.triangle_barycentric_coord.x;
-    isect->v = intersection.triangle_barycentric_coord.x;
+    isect->u = intersection.triangle_barycentric_coord.x;
+    isect->v = intersection.triangle_barycentric_coord.y;
   }
   else {
     isect->u = payload.u;

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

@@ -122,8 +122,8 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
   isect->object = object;
   isect->type = kernel_data_fetch(objects, object).primitive_type;
 
-  isect->u = 1.0f - barycentrics.y - barycentrics.x;
-  isect->v = barycentrics.x;
+  isect->u = barycentrics.x;
+  isect->v = barycentrics.y;
 
   /* Record geometric normal */
   const uint tri_vindex = kernel_data_fetch(tri_vindex, isect->prim).w;
@@ -185,18 +185,14 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
     return true;
   }
 
-  float u = 0.0f, v = 0.0f;
+  const float u = barycentrics.x;
+  const float v = barycentrics.y;
   int type = 0;
   if (intersection_type == METALRT_HIT_TRIANGLE) {
-    u = 1.0f - barycentrics.y - barycentrics.x;
-    v = barycentrics.x;
     type = kernel_data_fetch(objects, object).primitive_type;
   }
 #  ifdef __HAIR__
   else {
-    u = barycentrics.x;
-    v = barycentrics.y;
-
     const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
     type = segment.type;
     prim = segment.prim;

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

@@ -149,25 +149,13 @@ void oneapi_kernel_##name(KernelGlobalsGPU *ccl_restrict kg, \
 /* clang-format on */
 
 /* Types */
+
 /* It's not possible to use sycl types like sycl::float3, sycl::int3, etc
- * because these types have different interfaces from blender version */
+ * because these types have different interfaces from blender version. */
 
 using uchar = unsigned char;
 using sycl::half;
 
-struct float3 {
-  float x, y, z;
-};
-
-ccl_always_inline float3 make_float3(float x, float y, float z)
-{
-  return {x, y, z};
-}
-ccl_always_inline float3 make_float3(float x)
-{
-  return {x, x, x};
-}
-
 /* math functions */
 #define fabsf(x) sycl::fabs((x))
 #define copysignf(x, y) sycl::copysign((x), (y))

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

@@ -116,8 +116,8 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
   isect->type = kernel_data_fetch(objects, isect->object).primitive_type;
 
   const float2 barycentrics = optixGetTriangleBarycentrics();
-  isect->u = 1.0f - barycentrics.y - barycentrics.x;
-  isect->v = barycentrics.x;
+  isect->u = barycentrics.x;
+  isect->v = barycentrics.y;
 
   /* Record geometric normal. */
   const uint tri_vindex = kernel_data_fetch(tri_vindex, prim).w;
@@ -152,8 +152,8 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
   int type = 0;
   if (optixIsTriangleHit()) {
     const float2 barycentrics = optixGetTriangleBarycentrics();
-    u = 1.0f - barycentrics.y - barycentrics.x;
-    v = barycentrics.x;
+    u = barycentrics.x;
+    v = barycentrics.y;
     type = kernel_data_fetch(objects, object).primitive_type;
   }
 #  ifdef __HAIR__
@@ -336,8 +336,8 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
 
   if (optixIsTriangleHit()) {
     const float2 barycentrics = optixGetTriangleBarycentrics();
-    optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x));
-    optixSetPayload_2(__float_as_uint(barycentrics.x));
+    optixSetPayload_1(__float_as_uint(barycentrics.x));
+    optixSetPayload_2(__float_as_uint(barycentrics.y));
     optixSetPayload_3(prim);
     optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type);
   }

intern/cycles/kernel/geom/motion_triangle_intersect.h

@@ -27,8 +27,8 @@ ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg,
                                                        const float v,
                                                        float3 verts[3])
 {
-  float w = 1.0f - u - v;
-  float3 P = u * verts[0] + v * verts[1] + w * verts[2];
+  /* This appears to give slightly better precision than interpolating with w = (1 - u - v). */
+  float3 P = verts[0] + u * (verts[1] - verts[0]) + v * (verts[2] - verts[0]);
 
   if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);

intern/cycles/kernel/geom/object.h

@@ -503,20 +503,6 @@ ccl_device_inline void bvh_instance_push(KernelGlobals kg,
   *idir = bvh_inverse_direction(*dir);
 }
 
-/* Transform ray to exit static object in BVH. */
-
-ccl_device_inline void bvh_instance_pop(KernelGlobals kg,
-                                        int object,
-                                        ccl_private const Ray *ray,
-                                        ccl_private float3 *P,
-                                        ccl_private float3 *dir,
-                                        ccl_private float3 *idir)
-{
-  *P = ray->P;
-  *dir = bvh_clamp_direction(ray->D);
-  *idir = bvh_inverse_direction(*dir);
-}
-
 #ifdef __OBJECT_MOTION__
 /* Transform ray into object space to enter motion blurred object in BVH */
 
@@ -536,22 +522,20 @@ ccl_device_inline void bvh_instance_motion_push(KernelGlobals kg,
   *idir = bvh_inverse_direction(*dir);
 }
 
-/* Transform ray to exit motion blurred object in BVH. */
+#endif
 
-ccl_device_inline void bvh_instance_motion_pop(KernelGlobals kg,
-                                               int object,
-                                               ccl_private const Ray *ray,
-                                               ccl_private float3 *P,
-                                               ccl_private float3 *dir,
-                                               ccl_private float3 *idir)
+/* Transform ray to exit static object in BVH. */
+
+ccl_device_inline void bvh_instance_pop(ccl_private const Ray *ray,
+                                        ccl_private float3 *P,
+                                        ccl_private float3 *dir,
+                                        ccl_private float3 *idir)
 {
   *P = ray->P;
   *dir = bvh_clamp_direction(ray->D);
   *idir = bvh_inverse_direction(*dir);
 }
 
-#endif
-
 /* TODO: This can be removed when we know if no devices will require explicit
  * address space qualifiers for this case. */
 

intern/cycles/kernel/geom/subd_triangle.h

@@ -94,11 +94,11 @@ ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals kg,
     float2 uv[3];
     subd_triangle_patch_uv(kg, sd, uv);
 
-    float2 dpdu = uv[0] - uv[2];
-    float2 dpdv = uv[1] - uv[2];
+    float2 dpdu = uv[1] - uv[0];
+    float2 dpdv = uv[2] - uv[0];
 
     /* p is [s, t] */
-    float2 p = dpdu * sd->u + dpdv * sd->v + uv[2];
+    float2 p = dpdu * sd->u + dpdv * sd->v + uv[0];
 
     float a, dads, dadt;
     a = patch_eval_float(kg, sd, desc.offset, patch, p.x, p.y, 0, &dads, &dadt);
@@ -165,12 +165,12 @@ ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c;
+      *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a;
     if (dy)
-      *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c;
+      *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a;
 #endif
 
-    return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c;
+    return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a;
   }
   else if (desc.element == ATTR_ELEMENT_CORNER) {
     float2 uv[3];
@@ -195,12 +195,12 @@ ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c;
+      *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a;
     if (dy)
-      *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c;
+      *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a;
 #endif
 
-    return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c;
+    return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a;
   }
   else if (desc.element == ATTR_ELEMENT_OBJECT || desc.element == ATTR_ELEMENT_MESH) {
     if (dx)
@@ -233,11 +233,11 @@ ccl_device_noinline float2 subd_triangle_attribute_float2(KernelGlobals kg,
     float2 uv[3];
     subd_triangle_patch_uv(kg, sd, uv);
 
-    float2 dpdu = uv[0] - uv[2];
-    float2 dpdv = uv[1] - uv[2];
+    float2 dpdu = uv[1] - uv[0];
+    float2 dpdv = uv[2] - uv[0];
 
     /* p is [s, t] */
-    float2 p = dpdu * sd->u + dpdv * sd->v + uv[2];
+    float2 p = dpdu * sd->u + dpdv * sd->v + uv[0];
 
     float2 a, dads, dadt;
 
@@ -305,12 +305,12 @@ ccl_device_noinline float2 subd_triangle_attribute_float2(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c;
+      *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a;
     if (dy)
-      *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c;
+      *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a;
 #endif
 
-    return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c;
+    return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a;
   }
   else if (desc.element == ATTR_ELEMENT_CORNER) {
     float2 uv[3];
@@ -337,12 +337,12 @@ ccl_device_noinline float2 subd_triangle_attribute_float2(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c;
+      *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a;
     if (dy)
-      *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c;
+      *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a;
 #endif
 
-    return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c;
+    return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a;
   }
   else if (desc.element == ATTR_ELEMENT_OBJECT || desc.element == ATTR_ELEMENT_MESH) {
     if (dx)
@@ -375,11 +375,11 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
     float2 uv[3];
     subd_triangle_patch_uv(kg, sd, uv);
 
-    float2 dpdu = uv[0] - uv[2];
-    float2 dpdv = uv[1] - uv[2];
+    float2 dpdu = uv[1] - uv[0];
+    float2 dpdv = uv[2] - uv[0];
 
     /* p is [s, t] */
-    float2 p = dpdu * sd->u + dpdv * sd->v + uv[2];
+    float2 p = dpdu * sd->u + dpdv * sd->v + uv[0];
 
     float3 a, dads, dadt;
     a = patch_eval_float3(kg, sd, desc.offset, patch, p.x, p.y, 0, &dads, &dadt);
@@ -446,12 +446,12 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c;
+      *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a;
     if (dy)
-      *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c;
+      *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a;
 #endif
 
-    return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c;
+    return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a;
   }
   else if (desc.element == ATTR_ELEMENT_CORNER) {
     float2 uv[3];
@@ -478,12 +478,12 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c;
+      *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a;
     if (dy)
-      *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c;
+      *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a;
 #endif
 
-    return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c;
+    return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a;
   }
   else if (desc.element == ATTR_ELEMENT_OBJECT || desc.element == ATTR_ELEMENT_MESH) {
     if (dx)
@@ -516,11 +516,11 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
     float2 uv[3];
     subd_triangle_patch_uv(kg, sd, uv);
 
-    float2 dpdu = uv[0] - uv[2];
-    float2 dpdv = uv[1] - uv[2];
+    float2 dpdu = uv[1] - uv[0];
+    float2 dpdv = uv[2] - uv[0];
 
     /* p is [s, t] */
-    float2 p = dpdu * sd->u + dpdv * sd->v + uv[2];
+    float2 p = dpdu * sd->u + dpdv * sd->v + uv[0];
 
     float4 a, dads, dadt;
     if (desc.type == NODE_ATTR_RGBA) {
@@ -592,12 +592,12 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c;
+      *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a;
     if (dy)
-      *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c;
+      *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a;
 #endif
 
-    return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c;
+    return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a;
   }
   else if (desc.element == ATTR_ELEMENT_CORNER || desc.element == ATTR_ELEMENT_CORNER_BYTE) {
     float2 uv[3];
@@ -636,12 +636,12 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c;
+      *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a;
     if (dy)
-      *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c;
+      *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a;
 #endif
 
-    return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c;
+    return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a;
   }
   else if (desc.element == ATTR_ELEMENT_OBJECT || desc.element == ATTR_ELEMENT_MESH) {
     if (dx)

intern/cycles/kernel/geom/triangle.h

@@ -45,8 +45,8 @@ ccl_device_inline void triangle_point_normal(KernelGlobals kg,
   float3 v1 = kernel_data_fetch(tri_verts, tri_vindex.w + 1);
   float3 v2 = kernel_data_fetch(tri_verts, tri_vindex.w + 2);
   /* compute point */
-  float t = 1.0f - u - v;
-  *P = (u * v0 + v * v1 + t * v2);
+  float w = 1.0f - u - v;
+  *P = (w * v0 + u * v1 + v * v2);
   /* get object flags */
   int object_flag = kernel_data_fetch(object_flag, object);
   /* compute normal */
@@ -97,7 +97,7 @@ triangle_smooth_normal(KernelGlobals kg, float3 Ng, int prim, float u, float v)
   float3 n1 = kernel_data_fetch(tri_vnormal, tri_vindex.y);
   float3 n2 = kernel_data_fetch(tri_vnormal, tri_vindex.z);
 
-  float3 N = safe_normalize((1.0f - u - v) * n2 + u * n0 + v * n1);
+  float3 N = safe_normalize((1.0f - u - v) * n0 + u * n1 + v * n2);
 
   return is_zero(N) ? Ng : N;
 }
@@ -118,7 +118,7 @@ ccl_device_inline float3 triangle_smooth_normal_unnormalized(
     object_inverse_normal_transform(kg, sd, &n2);
   }
 
-  float3 N = (1.0f - u - v) * n2 + u * n0 + v * n1;
+  float3 N = (1.0f - u - v) * n0 + u * n1 + v * n2;
 
   return is_zero(N) ? Ng : N;
 }
@@ -137,8 +137,8 @@ ccl_device_inline void triangle_dPdudv(KernelGlobals kg,
   const float3 p2 = kernel_data_fetch(tri_verts, tri_vindex.w + 2);
 
   /* compute derivatives of P w.r.t. uv */
-  *dPdu = (p0 - p2);
-  *dPdv = (p1 - p2);
+  *dPdu = (p1 - p0);
+  *dPdv = (p2 - p0);
 }
 
 /* Reading attributes on various triangle elements */
@@ -167,12 +167,12 @@ ccl_device float triangle_attribute_float(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * f0 + sd->dv.dx * f1 - (sd->du.dx + sd->dv.dx) * f2;
+      *dx = sd->du.dx * f1 + sd->dv.dx * f2 - (sd->du.dx + sd->dv.dx) * f0;
     if (dy)
-      *dy = sd->du.dy * f0 + sd->dv.dy * f1 - (sd->du.dy + sd->dv.dy) * f2;
+      *dy = sd->du.dy * f1 + sd->dv.dy * f2 - (sd->du.dy + sd->dv.dy) * f0;
 #endif
 
-    return sd->u * f0 + sd->v * f1 + (1.0f - sd->u - sd->v) * f2;
+    return sd->u * f1 + sd->v * f2 + (1.0f - sd->u - sd->v) * f0;
   }
   else {
 #ifdef __RAY_DIFFERENTIALS__
@@ -217,12 +217,12 @@ ccl_device float2 triangle_attribute_float2(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * f0 + sd->dv.dx * f1 - (sd->du.dx + sd->dv.dx) * f2;
+      *dx = sd->du.dx * f1 + sd->dv.dx * f2 - (sd->du.dx + sd->dv.dx) * f0;
     if (dy)
-      *dy = sd->du.dy * f0 + sd->dv.dy * f1 - (sd->du.dy + sd->dv.dy) * f2;
+      *dy = sd->du.dy * f1 + sd->dv.dy * f2 - (sd->du.dy + sd->dv.dy) * f0;
 #endif
 
-    return sd->u * f0 + sd->v * f1 + (1.0f - sd->u - sd->v) * f2;
+    return sd->u * f1 + sd->v * f2 + (1.0f - sd->u - sd->v) * f0;
   }
   else {
 #ifdef __RAY_DIFFERENTIALS__
@@ -267,12 +267,12 @@ ccl_device float3 triangle_attribute_float3(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * f0 + sd->dv.dx * f1 - (sd->du.dx + sd->dv.dx) * f2;
+      *dx = sd->du.dx * f1 + sd->dv.dx * f2 - (sd->du.dx + sd->dv.dx) * f0;
     if (dy)
-      *dy = sd->du.dy * f0 + sd->dv.dy * f1 - (sd->du.dy + sd->dv.dy) * f2;
+      *dy = sd->du.dy * f1 + sd->dv.dy * f2 - (sd->du.dy + sd->dv.dy) * f0;
 #endif
 
-    return sd->u * f0 + sd->v * f1 + (1.0f - sd->u - sd->v) * f2;
+    return sd->u * f1 + sd->v * f2 + (1.0f - sd->u - sd->v) * f0;
   }
   else {
 #ifdef __RAY_DIFFERENTIALS__
@@ -328,12 +328,12 @@ ccl_device float4 triangle_attribute_float4(KernelGlobals kg,
 
 #ifdef __RAY_DIFFERENTIALS__
     if (dx)
-      *dx = sd->du.dx * f0 + sd->dv.dx * f1 - (sd->du.dx + sd->dv.dx) * f2;
+      *dx = sd->du.dx * f1 + sd->dv.dx * f2 - (sd->du.dx + sd->dv.dx) * f0;
     if (dy)
-      *dy = sd->du.dy * f0 + sd->dv.dy * f1 - (sd->du.dy + sd->dv.dy) * f2;
+      *dy = sd->du.dy * f1 + sd->dv.dy * f2 - (sd->du.dy + sd->dv.dy) * f0;
 #endif
 
-    return sd->u * f0 + sd->v * f1 + (1.0f - sd->u - sd->v) * f2;
+    return sd->u * f1 + sd->v * f2 + (1.0f - sd->u - sd->v) * f0;
   }
   else {
 #ifdef __RAY_DIFFERENTIALS__

intern/cycles/kernel/geom/triangle_intersect.h

@@ -145,9 +145,9 @@ ccl_device_inline float3 triangle_point_from_uv(KernelGlobals kg,
   const packed_float3 tri_a = kernel_data_fetch(tri_verts, tri_vindex + 0),
                       tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1),
                       tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2);
-  float w = 1.0f - u - v;
 
-  float3 P = u * tri_a + v * tri_b + w * tri_c;
+  /* This appears to give slightly better precision than interpolating with w = (1 - u - v). */
+  float3 P = tri_a + u * (tri_b - tri_a) + v * (tri_c - tri_a);
 
   if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);

intern/cycles/kernel/integrator/init_from_bake.h

@@ -155,6 +155,11 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg,
                                  1.0f - u);
   }
 
+  /* Convert from Blender to Cycles/Embree/OptiX barycentric convention. */
+  const float tmp = u;
+  u = v;
+  v = 1.0f - tmp - v;
+
   /* Position and normal on triangle. */
   const int object = kernel_data.bake.object_index;
   float3 P, Ng;

intern/cycles/kernel/integrator/mnee.h

@@ -186,7 +186,7 @@ ccl_device_forceinline void mnee_setup_manifold_vertex(KernelGlobals kg,
     triangle_vertices_and_normals(kg, sd_vtx->prim, verts, normals);
 
     /* Compute refined position (same code as in triangle_point_from_uv). */
-    sd_vtx->P = isect->u * verts[0] + isect->v * verts[1] + (1.f - isect->u - isect->v) * verts[2];
+    sd_vtx->P = (1.f - isect->u - isect->v) * verts[0] + isect->u * verts[1] + isect->v * verts[2];
     if (!(sd_vtx->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
       const Transform tfm = object_get_transform(kg, sd_vtx);
       sd_vtx->P = transform_point(&tfm, sd_vtx->P);
@@ -213,8 +213,8 @@ ccl_device_forceinline void mnee_setup_manifold_vertex(KernelGlobals kg,
   }
 
   /* Tangent space (position derivatives) WRT barycentric (u, v). */
-  float3 dp_du = verts[0] - verts[2];
-  float3 dp_dv = verts[1] - verts[2];
+  float3 dp_du = verts[1] - verts[0];
+  float3 dp_dv = verts[2] - verts[0];
 
   /* Geometric normal. */
   vtx->ng = normalize(cross(dp_du, dp_dv));
@@ -223,16 +223,16 @@ ccl_device_forceinline void mnee_setup_manifold_vertex(KernelGlobals kg,
 
   /* Shading normals: Interpolate normals between vertices. */
   float n_len;
-  vtx->n = normalize_len(normals[0] * sd_vtx->u + normals[1] * sd_vtx->v +
-                             normals[2] * (1.0f - sd_vtx->u - sd_vtx->v),
+  vtx->n = normalize_len(normals[0] * (1.0f - sd_vtx->u - sd_vtx->v) + normals[1] * sd_vtx->u +
+                             normals[2] * sd_vtx->v,
                          &n_len);
 
   /* Shading normal derivatives WRT barycentric (u, v)
    * we calculate the derivative of n = |u*n0 + v*n1 + (1-u-v)*n2| using:
    * d/du [f(u)/|f(u)|] = [d/du f(u)]/|f(u)| - f(u)/|f(u)|^3 <f(u), d/du f(u)>. */
   const float inv_n_len = 1.f / n_len;
-  float3 dn_du = inv_n_len * (normals[0] - normals[2]);
-  float3 dn_dv = inv_n_len * (normals[1] - normals[2]);
+  float3 dn_du = inv_n_len * (normals[1] - normals[0]);
+  float3 dn_dv = inv_n_len * (normals[2] - normals[0]);
   dn_du -= vtx->n * dot(vtx->n, dn_du);
   dn_dv -= vtx->n * dot(vtx->n, dn_dv);
 

intern/cycles/kernel/light/sample.h

@@ -137,8 +137,9 @@ ccl_device_inline float3 shadow_ray_smooth_surface_offset(
     triangle_vertices_and_normals(kg, sd->prim, V, N);
   }
 
-  const float u = sd->u, v = sd->v;
-  const float w = 1 - u - v;
+  const float u = 1.0f - sd->u - sd->v;
+  const float v = sd->u;
+  const float w = sd->v;
   float3 P = V[0] * u + V[1] * v + V[2] * w; /* Local space */
   float3 n = N[0] * u + N[1] * v + N[2] * w; /* We get away without normalization */
 

intern/cycles/kernel/osl/shaders/node_geometry.osl

@@ -20,7 +20,7 @@ shader node_geometry(normal NormalIn = N,
   Normal = NormalIn;
   TrueNormal = Ng;
   Incoming = I;
-  Parametric = point(u, v, 0.0);
+  Parametric = point(1.0 - u - v, u, 0.0);
   Backfacing = backfacing();
 
   if (bump_offset == "dx") {

intern/cycles/kernel/svm/geometry.h

@@ -34,7 +34,7 @@ ccl_device_noinline void svm_node_geometry(KernelGlobals kg,
       data = sd->Ng;
       break;
     case NODE_GEOM_uv:
-      data = make_float3(sd->u, sd->v, 0.0f);
+      data = make_float3(1.0f - sd->u - sd->v, sd->u, 0.0f);
       break;
     default:
       data = make_float3(0.0f, 0.0f, 0.0f);
@@ -57,7 +57,7 @@ ccl_device_noinline void svm_node_geometry_bump_dx(KernelGlobals kg,
       data = sd->P + sd->dP.dx;
       break;
     case NODE_GEOM_uv:
-      data = make_float3(sd->u + sd->du.dx, sd->v + sd->dv.dx, 0.0f);
+      data = make_float3(1.0f - sd->u - sd->du.dx - sd->v - sd->dv.dx, sd->u + sd->du.dx, 0.0f);
       break;
     default:
       svm_node_geometry(kg, sd, stack, type, out_offset);
@@ -84,7 +84,7 @@ ccl_device_noinline void svm_node_geometry_bump_dy(KernelGlobals kg,
       data = sd->P + sd->dP.dy;
       break;
     case NODE_GEOM_uv:
-      data = make_float3(sd->u + sd->du.dy, sd->v + sd->dv.dy, 0.0f);
+      data = make_float3(1.0f - sd->u - sd->du.dy - sd->v - sd->dv.dy, sd->u + sd->du.dy, 0.0f);
       break;
     default:
       svm_node_geometry(kg, sd, stack, type, out_offset);

intern/cycles/scene/mesh_displace.cpp

@@ -73,16 +73,16 @@ static int fill_shader_input(const Scene *scene,
 
       switch (j) {
         case 0:
-          u = 1.0f;
+          u = 0.0f;
           v = 0.0f;
           break;
         case 1:
-          u = 0.0f;
-          v = 1.0f;
+          u = 1.0f;
+          v = 0.0f;
           break;
         default:
           u = 0.0f;
-          v = 0.0f;
+          v = 1.0f;
           break;
       }
 

intern/cycles/util/CMakeLists.txt

@@ -129,8 +129,6 @@ set(SRC_HEADERS
   types_uint4.h
   types_uint4_impl.h
   types_ushort4.h
-  types_vector3.h
-  types_vector3_impl.h
   unique_ptr.h
   vector.h
   version.h

intern/cycles/util/math.h

@@ -953,7 +953,11 @@ ccl_device_inline uint prev_power_of_two(uint x)
 ccl_device_inline uint32_t reverse_integer_bits(uint32_t x)
 {
   /* Use a native instruction if it exists. */
-#if defined(__aarch64__) || defined(_M_ARM64)
+#if defined(__KERNEL_CUDA__)
+  return __brev(x);
+#elif defined(__KERNEL_METAL__)
+  return reverse_bits(x);
+#elif defined(__aarch64__) || defined(_M_ARM64)
   /* Assume the rbit is always available on 64bit ARM architecture. */
   __asm__("rbit %w0, %w1" : "=r"(x) : "r"(x));
   return x;
@@ -962,10 +966,6 @@ ccl_device_inline uint32_t reverse_integer_bits(uint32_t x)
    * This 32-bit Thumb instruction is available in ARMv6T2 and above. */
   __asm__("rbit %0, %1" : "=r"(x) : "r"(x));
   return x;
-#elif defined(__KERNEL_CUDA__)
-  return __brev(x);
-#elif defined(__KERNEL_METAL__)
-  return reverse_bits(x);
 #elif __has_builtin(__builtin_bitreverse32)
   return __builtin_bitreverse32(x);
 #else

intern/cycles/util/math_intersect.h

@@ -120,9 +120,9 @@ ccl_device_forceinline bool ray_triangle_intersect(const float3 ray_P,
    * in Embree. */
 
   /* Calculate vertices relative to ray origin. */
-  const float3 v0 = tri_c - ray_P;
-  const float3 v1 = tri_a - ray_P;
-  const float3 v2 = tri_b - ray_P;
+  const float3 v0 = tri_a - ray_P;
+  const float3 v1 = tri_b - ray_P;
+  const float3 v2 = tri_c - ray_P;
 
   /* Calculate triangle edges. */
   const float3 e0 = v2 - v0;
@@ -130,11 +130,12 @@ ccl_device_forceinline bool ray_triangle_intersect(const float3 ray_P,
   const float3 e2 = v1 - v2;
 
   /* Perform edge tests. */
-  const float U = dot(cross(v2 + v0, e0), ray_D);
-  const float V = dot(cross(v0 + v1, e1), ray_D);
-  const float W = dot(cross(v1 + v2, e2), ray_D);
+  const float U = dot(cross(e0, v2 + v0), ray_D);
+  const float V = dot(cross(e1, v0 + v1), ray_D);
+  const float W = dot(cross(e2, v1 + v2), ray_D);
 
-  const float eps = FLT_EPSILON * fabsf(U + V + W);
+  const float UVW = U + V + W;
+  const float eps = FLT_EPSILON * fabsf(UVW);
   const float minUVW = min(U, min(V, W));
   const float maxUVW = max(U, max(V, W));
 
@@ -158,8 +159,9 @@ ccl_device_forceinline bool ray_triangle_intersect(const float3 ray_P,
     return false;
   }
 
-  *isect_u = U / den;
-  *isect_v = V / den;
+  const float rcp_UVW = (fabsf(UVW) < 1e-18f) ? 0.0f : 1.0f / UVW;
+  *isect_u = min(U * rcp_UVW, 1.0f);
+  *isect_v = min(V * rcp_UVW, 1.0f);
   *isect_t = t;
   return true;
 }

intern/cycles/util/types.h

@@ -12,6 +12,7 @@
 
 #if !defined(__KERNEL_GPU__)
 #  include <stdint.h>
+#  include <stdio.h>
 #endif
 
 #include "util/defines.h"
@@ -70,6 +71,12 @@ ccl_device_inline bool is_power_of_two(size_t x)
 
 CCL_NAMESPACE_END
 
+/* Most GPU APIs matching native vector types, so we only need to implement them for
+ * CPU and oneAPI. */
+#if defined(__KERNEL_GPU__) && !defined(__KERNEL_ONEAPI__)
+#  define __KERNEL_NATIVE_VECTOR_TYPES__
+#endif
+
 /* Vectorized types declaration. */
 #include "util/types_uchar2.h"
 #include "util/types_uchar3.h"
@@ -90,8 +97,6 @@ CCL_NAMESPACE_END
 #include "util/types_float4.h"
 #include "util/types_float8.h"
 
-#include "util/types_vector3.h"
-
 /* Vectorized types implementation. */
 #include "util/types_uchar2_impl.h"
 #include "util/types_uchar3_impl.h"
@@ -110,8 +115,6 @@ CCL_NAMESPACE_END
 #include "util/types_float4_impl.h"
 #include "util/types_float8_impl.h"
 
-#include "util/types_vector3_impl.h"
-
 /* SSE types. */
 #ifndef __KERNEL_GPU__
 #  include "util/sseb.h"

intern/cycles/util/types_float2.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT2_H__
-#define __UTIL_TYPES_FLOAT2_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,18 +9,18 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct float2 {
   float x, y;
 
+#  ifndef __KERNEL_GPU__
   __forceinline float operator[](int i) const;
   __forceinline float &operator[](int i);
+#  endif
 };
 
 ccl_device_inline float2 make_float2(float x, float y);
 ccl_device_inline void print_float2(const char *label, const float2 &a);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT2_H__ */

intern/cycles/util/types_float2_impl.h

@@ -1,20 +1,16 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT2_IMPL_H__
-#define __UTIL_TYPES_FLOAT2_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 __forceinline float float2::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -28,6 +24,7 @@ __forceinline float &float2::operator[](int i)
   util_assert(i < 2);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline float2 make_float2(float x, float y)
 {
@@ -39,8 +36,6 @@ ccl_device_inline void print_float2(const char *label, const float2 &a)
 {
   printf("%s: %.8f %.8f\n", label, (double)a.x, (double)a.y);
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT2_IMPL_H__ */

intern/cycles/util/types_float3.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT3_H__
-#define __UTIL_TYPES_FLOAT3_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,28 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct ccl_try_align(16) float3
 {
-#  ifdef __KERNEL_SSE__
+#  ifdef __KERNEL_GPU__
+  /* Compact structure for GPU. */
+  float x, y, z;
+#  else
+  /* SIMD aligned structure for CPU. */
+#    ifdef __KERNEL_SSE__
   union {
     __m128 m128;
     struct {
       float x, y, z, w;
     };
   };
+#    else
+  float x, y, z, w;
+#    endif
+#  endif
 
+#  ifdef __KERNEL_SSE__
+  /* Convenient constructors and operators for SIMD, otherwise default is enough. */
   __forceinline float3();
   __forceinline float3(const float3 &a);
   __forceinline explicit float3(const __m128 &a);
@@ -29,18 +39,18 @@ struct ccl_try_align(16) float3
   __forceinline operator __m128 &();
 
   __forceinline float3 &operator=(const float3 &a);
-#  else  /* __KERNEL_SSE__ */
-  float x, y, z, w;
-#  endif /* __KERNEL_SSE__ */
+#  endif
 
+#  ifndef __KERNEL_GPU__
   __forceinline float operator[](int i) const;
   __forceinline float &operator[](int i);
+#  endif
 };
 
 ccl_device_inline float3 make_float3(float f);
 ccl_device_inline float3 make_float3(float x, float y, float z);
 ccl_device_inline void print_float3(const char *label, const float3 &a);
-#endif /* !defined(__KERNEL_GPU__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 /* Smaller float3 for storage. For math operations this must be converted to float3, so that on the
  * CPU SIMD instructions can be used. */
@@ -78,5 +88,3 @@ struct packed_float3 {
 static_assert(sizeof(packed_float3) == 12, "packed_float3 expected to be exactly 12 bytes");
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT3_H__ */

intern/cycles/util/types_float3_impl.h

@@ -1,20 +1,15 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT3_IMPL_H__
-#define __UTIL_TYPES_FLOAT3_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 #  ifdef __KERNEL_SSE__
 __forceinline float3::float3()
 {
@@ -45,6 +40,7 @@ __forceinline float3 &float3::operator=(const float3 &a)
 }
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
 __forceinline float float3::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -58,23 +54,32 @@ __forceinline float &float3::operator[](int i)
   util_assert(i < 3);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline float3 make_float3(float f)
 {
-#  ifdef __KERNEL_SSE__
-  float3 a(_mm_set1_ps(f));
+#  ifdef __KERNEL_GPU__
+  float3 a = {f, f, f};
 #  else
+#    ifdef __KERNEL_SSE__
+  float3 a(_mm_set1_ps(f));
+#    else
   float3 a = {f, f, f, f};
+#    endif
 #  endif
   return a;
 }
 
 ccl_device_inline float3 make_float3(float x, float y, float z)
 {
-#  ifdef __KERNEL_SSE__
-  float3 a(_mm_set_ps(0.0f, z, y, x));
+#  ifdef __KERNEL_GPU__
+  float3 a = {x, y, z};
 #  else
+#    ifdef __KERNEL_SSE__
+  float3 a(_mm_set_ps(0.0f, z, y, x));
+#    else
   float3 a = {x, y, z, 0.0f};
+#    endif
 #  endif
   return a;
 }
@@ -83,8 +88,6 @@ ccl_device_inline void print_float3(const char *label, const float3 &a)
 {
   printf("%s: %.8f %.8f %.8f\n", label, (double)a.x, (double)a.y, (double)a.z);
 }
-#endif /* !defined(__KERNEL_GPU__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT3_IMPL_H__ */

intern/cycles/util/types_float4.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT4_H__
-#define __UTIL_TYPES_FLOAT4_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct int4;
 
 struct ccl_try_align(16) float4
@@ -35,16 +34,16 @@ struct ccl_try_align(16) float4
   float x, y, z, w;
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
   __forceinline float operator[](int i) const;
   __forceinline float &operator[](int i);
+#  endif
 };
 
 ccl_device_inline float4 make_float4(float f);
 ccl_device_inline float4 make_float4(float x, float y, float z, float w);
 ccl_device_inline float4 make_float4(const int4 &i);
 ccl_device_inline void print_float4(const char *label, const float4 &a);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT4_H__ */

intern/cycles/util/types_float4_impl.h

@@ -1,20 +1,15 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT4_IMPL_H__
-#define __UTIL_TYPES_FLOAT4_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 #  ifdef __KERNEL_SSE__
 __forceinline float4::float4()
 {
@@ -41,6 +36,7 @@ __forceinline float4 &float4::operator=(const float4 &a)
 }
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
 __forceinline float float4::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -54,6 +50,7 @@ __forceinline float &float4::operator[](int i)
   util_assert(i < 4);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline float4 make_float4(float f)
 {
@@ -89,8 +86,6 @@ ccl_device_inline void print_float4(const char *label, const float4 &a)
 {
   printf("%s: %.8f %.8f %.8f %.8f\n", label, (double)a.x, (double)a.y, (double)a.z, (double)a.w);
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT4_IMPL_H__ */

intern/cycles/util/types_float8.h

@@ -2,8 +2,7 @@
  * Original code Copyright 2017, Intel Corporation
  * Modifications Copyright 2018-2022 Blender Foundation. */
 
-#ifndef __UTIL_TYPES_FLOAT8_H__
-#define __UTIL_TYPES_FLOAT8_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -12,7 +11,7 @@
 CCL_NAMESPACE_BEGIN
 
 /* float8 is a reserved type in Metal that has not been implemented. For
- * that reason this is named float8_t. */
+ * that reason this is named float8_t and not using native vector types. */
 
 #ifdef __KERNEL_GPU__
 struct float8_t
@@ -52,5 +51,3 @@ ccl_device_inline float8_t
 make_float8_t(float a, float b, float c, float d, float e, float f, float g, float h);
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT8_H__ */

intern/cycles/util/types_float8_impl.h

@@ -2,17 +2,12 @@
  * Original code Copyright 2017, Intel Corporation
  * Modifications Copyright 2018-2022 Blender Foundation. */
 
-#ifndef __UTIL_TYPES_FLOAT8_IMPL_H__
-#define __UTIL_TYPES_FLOAT8_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
 #ifdef __KERNEL_AVX2__
@@ -83,5 +78,3 @@ make_float8_t(float a, float b, float c, float d, float e, float f, float g, flo
 }
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT8_IMPL_H__ */

intern/cycles/util/types_int2.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT2_H__
-#define __UTIL_TYPES_INT2_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct int2 {
   int x, y;
 
+#  ifndef __KERNEL_GPU__
   __forceinline int operator[](int i) const;
   __forceinline int &operator[](int i);
+#  endif
 };
 
 ccl_device_inline int2 make_int2(int x, int y);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT2_H__ */

intern/cycles/util/types_int2_impl.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT2_IMPL_H__
-#define __UTIL_TYPES_INT2_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 int int2::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -24,14 +24,13 @@ int &int2::operator[](int i)
   util_assert(i < 2);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline int2 make_int2(int x, int y)
 {
   int2 a = {x, y};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT2_IMPL_H__ */

intern/cycles/util/types_int3.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT3_H__
-#define __UTIL_TYPES_INT3_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,10 +9,15 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct ccl_try_align(16) int3
 {
-#  ifdef __KERNEL_SSE__
+#  ifdef __KERNEL_GPU__
+  /* Compact structure on the GPU. */
+  int x, y, z;
+#  else
+  /* SIMD aligned structure for CPU. */
+#    ifdef __KERNEL_SSE__
   union {
     __m128i m128;
     struct {
@@ -29,19 +33,20 @@ struct ccl_try_align(16) int3
   __forceinline operator __m128i &();
 
   __forceinline int3 &operator=(const int3 &a);
-#  else  /* __KERNEL_SSE__ */
+#    else  /* __KERNEL_SSE__ */
   int x, y, z, w;
-#  endif /* __KERNEL_SSE__ */
+#    endif /* __KERNEL_SSE__ */
+#  endif
 
+#  ifndef __KERNEL_GPU__
   __forceinline int operator[](int i) const;
   __forceinline int &operator[](int i);
+#  endif
 };
 
 ccl_device_inline int3 make_int3(int i);
 ccl_device_inline int3 make_int3(int x, int y, int z);
 ccl_device_inline void print_int3(const char *label, const int3 &a);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT3_H__ */

intern/cycles/util/types_int3_impl.h

@@ -1,20 +1,15 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT3_IMPL_H__
-#define __UTIL_TYPES_INT3_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 #  ifdef __KERNEL_SSE__
 __forceinline int3::int3()
 {
@@ -45,6 +40,7 @@ __forceinline int3 &int3::operator=(const int3 &a)
 }
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
 __forceinline int int3::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -58,25 +54,33 @@ __forceinline int &int3::operator[](int i)
   util_assert(i < 3);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline int3 make_int3(int i)
 {
-#  ifdef __KERNEL_SSE__
-  int3 a(_mm_set1_epi32(i));
+#  ifdef __KERNEL_GPU__
+  int3 a = {i, i, i};
 #  else
+#    ifdef __KERNEL_SSE__
+  int3 a(_mm_set1_epi32(i));
+#    else
   int3 a = {i, i, i, i};
+#    endif
 #  endif
   return a;
 }
 
 ccl_device_inline int3 make_int3(int x, int y, int z)
 {
-#  ifdef __KERNEL_SSE__
-  int3 a(_mm_set_epi32(0, z, y, x));
+#  ifdef __KERNEL_GPU__
+  int3 a = {x, y, z};
 #  else
+#    ifdef __KERNEL_SSE__
+  int3 a(_mm_set_epi32(0, z, y, x));
+#    else
   int3 a = {x, y, z, 0};
+#    endif
 #  endif
-
   return a;
 }
 
@@ -84,8 +88,6 @@ ccl_device_inline void print_int3(const char *label, const int3 &a)
 {
   printf("%s: %d %d %d\n", label, a.x, a.y, a.z);
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT3_IMPL_H__ */

intern/cycles/util/types_int4.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT4_H__
-#define __UTIL_TYPES_INT4_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 
 struct float3;
 struct float4;
@@ -37,8 +36,10 @@ struct ccl_try_align(16) int4
   int x, y, z, w;
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
   __forceinline int operator[](int i) const;
   __forceinline int &operator[](int i);
+#  endif
 };
 
 ccl_device_inline int4 make_int4(int i);
@@ -46,8 +47,6 @@ ccl_device_inline int4 make_int4(int x, int y, int z, int w);
 ccl_device_inline int4 make_int4(const float3 &f);
 ccl_device_inline int4 make_int4(const float4 &f);
 ccl_device_inline void print_int4(const char *label, const int4 &a);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT4_H__ */

intern/cycles/util/types_int4_impl.h

@@ -1,20 +1,15 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT4_IMPL_H__
-#define __UTIL_TYPES_INT4_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 #  ifdef __KERNEL_SSE__
 __forceinline int4::int4()
 {
@@ -45,6 +40,7 @@ __forceinline int4 &int4::operator=(const int4 &a)
 }
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
 __forceinline int int4::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -58,6 +54,7 @@ __forceinline int &int4::operator[](int i)
   util_assert(i < 4);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline int4 make_int4(int i)
 {
@@ -105,8 +102,6 @@ ccl_device_inline void print_int4(const char *label, const int4 &a)
 {
   printf("%s: %d %d %d %d\n", label, a.x, a.y, a.z, a.w);
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT4_IMPL_H__ */

intern/cycles/util/types_uchar2.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UCHAR2_H__
-#define __UTIL_TYPES_UCHAR2_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uchar2 {
   uchar x, y;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uchar operator[](int i) const;
   __forceinline uchar &operator[](int i);
+#  endif
 };
 
 ccl_device_inline uchar2 make_uchar2(uchar x, uchar y);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UCHAR2_H__ */

intern/cycles/util/types_uchar2_impl.h

@@ -10,7 +10,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 uchar uchar2::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -24,13 +25,14 @@ uchar &uchar2::operator[](int i)
   util_assert(i < 2);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uchar2 make_uchar2(uchar x, uchar y)
 {
   uchar2 a = {x, y};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
 

intern/cycles/util/types_uchar3.h

@@ -10,16 +10,18 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uchar3 {
   uchar x, y, z;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uchar operator[](int i) const;
   __forceinline uchar &operator[](int i);
+#  endif
 };
 
 ccl_device_inline uchar3 make_uchar3(uchar x, uchar y, uchar z);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
 

intern/cycles/util/types_uchar3_impl.h

@@ -10,7 +10,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 uchar uchar3::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -24,13 +25,14 @@ uchar &uchar3::operator[](int i)
   util_assert(i < 3);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uchar3 make_uchar3(uchar x, uchar y, uchar z)
 {
   uchar3 a = {x, y, z};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
 

intern/cycles/util/types_uchar4.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UCHAR4_H__
-#define __UTIL_TYPES_UCHAR4_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uchar4 {
   uchar x, y, z, w;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uchar operator[](int i) const;
   __forceinline uchar &operator[](int i);
+#  endif
 };
 
 ccl_device_inline uchar4 make_uchar4(uchar x, uchar y, uchar z, uchar w);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UCHAR4_H__ */

intern/cycles/util/types_uchar4_impl.h

@@ -10,7 +10,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 uchar uchar4::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -24,13 +25,14 @@ uchar &uchar4::operator[](int i)
   util_assert(i < 4);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uchar4 make_uchar4(uchar x, uchar y, uchar z, uchar w)
 {
   uchar4 a = {x, y, z, w};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
 

intern/cycles/util/types_uint2.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT2_H__
-#define __UTIL_TYPES_UINT2_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uint2 {
   uint x, y;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uint operator[](uint i) const;
   __forceinline uint &operator[](uint i);
+#  endif
 };
 
 ccl_device_inline uint2 make_uint2(uint x, uint y);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT2_H__ */

intern/cycles/util/types_uint2_impl.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT2_IMPL_H__
-#define __UTIL_TYPES_UINT2_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 __forceinline uint uint2::operator[](uint i) const
 {
   util_assert(i < 2);
@@ -22,14 +22,13 @@ __forceinline uint &uint2::operator[](uint i)
   util_assert(i < 2);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uint2 make_uint2(uint x, uint y)
 {
   uint2 a = {x, y};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT2_IMPL_H__ */

intern/cycles/util/types_uint3.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT3_H__
-#define __UTIL_TYPES_UINT3_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uint3 {
   uint x, y, z;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uint operator[](uint i) const;
   __forceinline uint &operator[](uint i);
+#  endif
 };
 
 ccl_device_inline uint3 make_uint3(uint x, uint y, uint z);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT3_H__ */

intern/cycles/util/types_uint3_impl.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT3_IMPL_H__
-#define __UTIL_TYPES_UINT3_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 __forceinline uint uint3::operator[](uint i) const
 {
   util_assert(i < 3);
@@ -22,14 +22,13 @@ __forceinline uint &uint3::operator[](uint i)
   util_assert(i < 3);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uint3 make_uint3(uint x, uint y, uint z)
 {
   uint3 a = {x, y, z};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT3_IMPL_H__ */

intern/cycles/util/types_uint4.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT4_H__
-#define __UTIL_TYPES_UINT4_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uint4 {
   uint x, y, z, w;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uint operator[](uint i) const;
   __forceinline uint &operator[](uint i);
+#  endif
 };
 
 ccl_device_inline uint4 make_uint4(uint x, uint y, uint z, uint w);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT4_H__ */

intern/cycles/util/types_uint4_impl.h

@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT4_IMPL_H__
-#define __UTIL_TYPES_UINT4_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 __forceinline uint uint4::operator[](uint i) const
 {
   util_assert(i < 3);
@@ -22,14 +22,13 @@ __forceinline uint &uint4::operator[](uint i)
   util_assert(i < 3);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uint4 make_uint4(uint x, uint y, uint z, uint w)
 {
   uint4 a = {x, y, z, w};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT4_IMPL_H__ */

intern/cycles/util/types_ushort4.h

@@ -10,7 +10,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 
 struct ushort4 {
   uint16_t x, y, z, w;

intern/cycles/util/types_vector3.h

@@ -1,26 +0,0 @@
-/* SPDX-License-Identifier: Apache-2.0
- * Copyright 2011-2022 Blender Foundation */
-
-#ifndef __UTIL_TYPES_VECTOR3_H__
-#define __UTIL_TYPES_VECTOR3_H__
-
-#ifndef __UTIL_TYPES_H__
-#  error "Do not include this file directly, include util/types.h instead."
-#endif
-
-CCL_NAMESPACE_BEGIN
-
-#ifndef __KERNEL_GPU__
-template<typename T> class vector3 {
- public:
-  T x, y, z;
-
-  __forceinline vector3();
-  __forceinline vector3(const T &a);
-  __forceinline vector3(const T &x, const T &y, const T &z);
-};
-#endif /* __KERNEL_GPU__ */
-
-CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_VECTOR3_H__ */

intern/cycles/util/types_vector3_impl.h

@@ -1,30 +0,0 @@
-/* SPDX-License-Identifier: Apache-2.0
- * Copyright 2011-2022 Blender Foundation */
-
-#ifndef __UTIL_TYPES_VECTOR3_IMPL_H__
-#define __UTIL_TYPES_VECTOR3_IMPL_H__
-
-#ifndef __UTIL_TYPES_H__
-#  error "Do not include this file directly, include util/types.h instead."
-#endif
-
-CCL_NAMESPACE_BEGIN
-
-#ifndef __KERNEL_GPU__
-template<typename T> ccl_always_inline vector3<T>::vector3()
-{
-}
-
-template<typename T> ccl_always_inline vector3<T>::vector3(const T &a) : x(a), y(a), z(a)
-{
-}
-
-template<typename T>
-ccl_always_inline vector3<T>::vector3(const T &x, const T &y, const T &z) : x(x), y(y), z(z)
-{
-}
-#endif /* __KERNEL_GPU__ */
-
-CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_VECTOR3_IMPL_H__ */

release/scripts/startup/bl_operators/node.py

@@ -149,37 +149,6 @@ class NODE_OT_add_node(NodeAddOperator, Operator):
     bl_options = {'REGISTER', 'UNDO'}
 
 
-# Add a node and link it to an existing socket
-class NODE_OT_add_and_link_node(NodeAddOperator, Operator):
-    '''Add a node to the active tree and link to an existing socket'''
-    bl_idname = "node.add_and_link_node"
-    bl_label = "Add and Link Node"
-    bl_options = {'REGISTER', 'UNDO'}
-
-    link_socket_index: IntProperty(
-        name="Link Socket Index",
-        description="Index of the socket to link",
-    )
-
-    def execute(self, context):
-        space = context.space_data
-        ntree = space.edit_tree
-
-        node = self.create_node(context)
-        if not node:
-            return {'CANCELLED'}
-
-        to_socket = getattr(context, "link_to_socket", None)
-        if to_socket:
-            ntree.links.new(node.outputs[self.link_socket_index], to_socket)
-
-        from_socket = getattr(context, "link_from_socket", None)
-        if from_socket:
-            ntree.links.new(from_socket, node.inputs[self.link_socket_index])
-
-        return {'FINISHED'}
-
-
 class NODE_OT_add_search(NodeAddOperator, Operator):
     '''Add a node to the active tree'''
     bl_idname = "node.add_search"
@@ -306,7 +275,6 @@ class NODE_OT_tree_path_parent(Operator):
 classes = (
     NodeSetting,
 
-    NODE_OT_add_and_link_node,
     NODE_OT_add_node,
     NODE_OT_add_search,
     NODE_OT_collapse_hide_unused_toggle,

release/scripts/startup/bl_ui/properties_render.py

@@ -162,6 +162,35 @@ class RENDER_PT_eevee_motion_blur(RenderButtonsPanel, Panel):
         col.prop(props, "motion_blur_steps", text="Steps")
 
 
+class RENDER_PT_eevee_next_motion_blur(RenderButtonsPanel, Panel):
+    bl_label = "Motion Blur"
+    bl_options = {'DEFAULT_CLOSED'}
+    COMPAT_ENGINES = {'BLENDER_EEVEE_NEXT'}
+
+    @classmethod
+    def poll(cls, context):
+        return (context.engine in cls.COMPAT_ENGINES)
+
+    def draw_header(self, context):
+        scene = context.scene
+        props = scene.eevee
+        self.layout.prop(props, "use_motion_blur", text="")
+
+    def draw(self, context):
+        layout = self.layout
+        layout.use_property_split = True
+        scene = context.scene
+        props = scene.eevee
+
+        layout.active = props.use_motion_blur
+        col = layout.column()
+        col.prop(props, "motion_blur_position", text="Position")
+        col.prop(props, "motion_blur_shutter")
+        col.separator()
+        col.prop(props, "motion_blur_depth_scale")
+        col.prop(props, "motion_blur_steps", text="Steps")
+
+
 class RENDER_PT_eevee_depth_of_field(RenderButtonsPanel, Panel):
     bl_label = "Depth of Field"
     bl_options = {'DEFAULT_CLOSED'}
@@ -756,6 +785,7 @@ classes = (
     RENDER_PT_eevee_film,
 
     RENDER_PT_eevee_next_sampling,
+    RENDER_PT_eevee_next_motion_blur,
     RENDER_PT_eevee_next_film,
 
     RENDER_PT_gpencil,

release/scripts/startup/bl_ui/properties_view_layer.py

@@ -79,6 +79,7 @@ class VIEWLAYER_PT_eevee_next_layer_passes_data(ViewLayerButtonsPanel, Panel):
         layout.use_property_split = True
         layout.use_property_decorate = False
 
+        scene = context.scene
         view_layer = context.view_layer
 
         col = layout.column()
@@ -87,7 +88,9 @@ class VIEWLAYER_PT_eevee_next_layer_passes_data(ViewLayerButtonsPanel, Panel):
         col.prop(view_layer, "use_pass_mist")
         col.prop(view_layer, "use_pass_normal")
         col.prop(view_layer, "use_pass_position")
-        col.prop(view_layer, "use_pass_vector")
+        sub = col.column()
+        sub.active = not scene.eevee.use_motion_blur
+        sub.prop(view_layer, "use_pass_vector")
 
 
 class VIEWLAYER_PT_eevee_layer_passes_light(ViewLayerButtonsPanel, Panel):

source/blender/blenfont/BLF_api.h

@@ -18,10 +18,10 @@ extern "C" {
 #define BLF_DATAFILES_FONTS_DIR "fonts"
 
 /* File name of the default variable-width font. */
-#define BLF_DEFAULT_PROPORTIONAL_FONT "droidsans.ttf"
+#define BLF_DEFAULT_PROPORTIONAL_FONT "DejaVuSans.woff2"
 
 /* File name of the default fixed-pitch font. */
-#define BLF_DEFAULT_MONOSPACED_FONT "bmonofont-i18n.ttf"
+#define BLF_DEFAULT_MONOSPACED_FONT "DejaVuSansMono.woff2"
 
 /* enable this only if needed (unused circa 2016) */
 #define BLF_BLUR_ENABLE 0
@@ -351,6 +351,8 @@ enum {
   BLF_DEFAULT = 1 << 14,
   /** Must only be used as last font in the stack. */
   BLF_LAST_RESORT = 1 << 15,
+  /** Failure to load this font. Don't try again. */
+  BLF_BAD_FONT = 1 << 16,
 };
 
 #define BLF_DRAW_STR_DUMMY_MAX 1024

source/blender/blenfont/intern/blf_font.c

@@ -18,8 +18,9 @@
 
 #include FT_FREETYPE_H
 #include FT_GLYPH_H
-#include FT_TRUETYPE_TABLES_H  /* For TT_OS2 */
 #include FT_MULTIPLE_MASTERS_H /* Variable font support. */
+#include FT_TRUETYPE_IDS_H     /* Codepoint coverage constants. */
+#include FT_TRUETYPE_TABLES_H  /* For TT_OS2 */
 
 #include "MEM_guardedalloc.h"
 
@@ -28,6 +29,7 @@
 #include "BLI_listbase.h"
 #include "BLI_math.h"
 #include "BLI_math_color_blend.h"
+#include "BLI_path_util.h"
 #include "BLI_rect.h"
 #include "BLI_string.h"
 #include "BLI_string_utf8.h"
@@ -836,7 +838,7 @@ static void blf_font_boundbox_foreach_glyph_ex(FontBLF *font,
   size_t i = 0, i_curr;
   rcti gbox_px;
 
-  if (str_len == 0) {
+  if (str_len == 0 || str[0] == 0) {
     /* early output. */
     return;
   }
@@ -1157,7 +1159,7 @@ int blf_font_ascender(FontBLF *font)
 
 char *blf_display_name(FontBLF *font)
 {
-  if (!font->face->family_name) {
+  if (!blf_ensure_face(font) || !font->face->family_name) {
     return NULL;
   }
   return BLI_sprintfN("%s %s", font->face->family_name, font->face->style_name);
@@ -1244,14 +1246,28 @@ static void blf_font_fill(FontBLF *font)
   font->glyph_cache_mutex = &blf_glyph_cache_mutex;
 }
 
-FontBLF *blf_font_new(const char *name, const char *filepath)
+/**
+ * Create an FT_Face for this font if not already existing.
+ */
+bool blf_ensure_face(FontBLF *font)
 {
-  FontBLF *font;
-  FT_Error err;
-  char *mfile;
+  if (font->face) {
+    return true;
+  }
+
+  if (font->flags & BLF_BAD_FONT) {
+    return false;
+  }
+
+  FT_Error err;
+
+  if (font->filepath) {
+    err = FT_New_Face(ft_lib, font->filepath, 0, &font->face);
+  }
+  if (font->mem) {
+    err = FT_New_Memory_Face(ft_lib, font->mem, (FT_Long)font->mem_size, 0, &font->face);
+  }
 
-  font = (FontBLF *)MEM_callocN(sizeof(FontBLF), "blf_font_new");
-  err = FT_New_Face(ft_lib, filepath, 0, &font->face);
   if (err) {
     if (ELEM(err, FT_Err_Unknown_File_Format, FT_Err_Unimplemented_Feature)) {
       printf("Format of this font file is not supported\n");
@@ -1259,8 +1275,8 @@ FontBLF *blf_font_new(const char *name, const char *filepath)
     else {
       printf("Error encountered while opening font file\n");
     }
-    MEM_freeN(font);
-    return NULL;
+    font->flags |= BLF_BAD_FONT;
+    return false;
   }
 
   err = FT_Select_Charmap(font->face, FT_ENCODING_UNICODE);
@@ -1272,28 +1288,28 @@ FontBLF *blf_font_new(const char *name, const char *filepath)
   }
   if (err) {
     printf("Can't set a character map!\n");
-    FT_Done_Face(font->face);
-    MEM_freeN(font);
-    return NULL;
+    font->flags |= BLF_BAD_FONT;
+    return false;
   }
 
-  mfile = blf_dir_metrics_search(filepath);
-  if (mfile) {
-    err = FT_Attach_File(font->face, mfile);
-    if (err) {
-      fprintf(stderr, "FT_Attach_File failed to load '%s' with error %d\n", filepath, (int)err);
+  if (font->filepath) {
+    char *mfile = blf_dir_metrics_search(font->filepath);
+    if (mfile) {
+      err = FT_Attach_File(font->face, mfile);
+      if (err) {
+        fprintf(stderr,
+                "FT_Attach_File failed to load '%s' with error %d\n",
+                font->filepath,
+                (int)err);
+      }
+      MEM_freeN(mfile);
     }
-    MEM_freeN(mfile);
   }
 
   if (FT_HAS_MULTIPLE_MASTERS(font->face)) {
     FT_Get_MM_Var(font->face, &(font->variations));
   }
 
-  font->name = BLI_strdup(name);
-  font->filepath = BLI_strdup(filepath);
-  blf_font_fill(font);
-
   /* Save TrueType table with bits to quickly test most unicode block coverage. */
   TT_OS2 *os2_table = (TT_OS2 *)FT_Get_Sfnt_Table(font->face, FT_SFNT_OS2);
   if (os2_table) {
@@ -1303,17 +1319,11 @@ FontBLF *blf_font_new(const char *name, const char *filepath)
     font->UnicodeRanges[3] = (uint)os2_table->ulUnicodeRange4;
   }
 
-  /* Detect "Last resort" fonts. They have everything. Usually except last 5 bits.  */
-  if (font->UnicodeRanges[0] == 0xffffffffU && font->UnicodeRanges[1] == 0xffffffffU &&
-      font->UnicodeRanges[2] == 0xffffffffU && font->UnicodeRanges[3] >= 0x7FFFFFFU) {
-    font->flags |= BLF_LAST_RESORT;
-  }
-
   if (FT_IS_FIXED_WIDTH(font->face)) {
     font->flags |= BLF_MONOSPACED;
   }
 
-  if (FT_HAS_KERNING(font->face)) {
+  if (FT_HAS_KERNING(font->face) && !font->kerning_cache) {
     /* Create kerning cache table and fill with value indicating "unset". */
     font->kerning_cache = MEM_mallocN(sizeof(KerningCacheBLF), __func__);
     for (uint i = 0; i < KERNING_CACHE_TABLE_SIZE; i++) {
@@ -1323,49 +1333,114 @@ FontBLF *blf_font_new(const char *name, const char *filepath)
     }
   }
 
+  return true;
+}
+
+typedef struct eFaceDetails {
+  char name[50];
+  unsigned int coverage1;
+  unsigned int coverage2;
+  unsigned int coverage3;
+  unsigned int coverage4;
+} eFaceDetails;
+
+/* Details about the fallback fonts we ship, so that we can load only when needed. */
+static const eFaceDetails static_face_details[] = {
+    {"lastresort.woff2", UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX},
+    {"Noto Sans CJK Regular.woff2", 0x30000083L, 0x2BDF3C10L, 0x16L, 0},
+    {"NotoEmoji-VariableFont_wght.woff2", 0x80000003L, 0x241E4ACL, 0x14000000L, 0x4000000L},
+    {"NotoSansArabic-VariableFont_wdth,wght.woff2", TT_UCR_ARABIC, 0, 0, 0},
+    {"NotoSansArmenian-VariableFont_wdth,wght.woff2",
+     TT_UCR_ARMENIAN,
+     TT_UCR_ALPHABETIC_PRESENTATION_FORMS,
+     0,
+     0},
+    {"NotoSansBengali-VariableFont_wdth,wght.woff2", TT_UCR_BENGALI, 0, 0, 0},
+    {"NotoSansDevanagari-Regular.woff2", TT_UCR_DEVANAGARI, 0, 0, 0},
+    {"NotoSansEthiopic-Regular.woff2", 0, 0, TT_UCR_ETHIOPIC, 0},
+    {"NotoSansGeorgian-VariableFont_wdth,wght.woff2", TT_UCR_GEORGIAN, 0, 0, 0},
+    {"NotoSansGujarati-Regular.woff2", TT_UCR_GUJARATI, 0, 0, 0},
+    {"NotoSansGurmukhi-VariableFont_wdth,wght.woff2", TT_UCR_GURMUKHI, 0, 0, 0},
+    {"NotoSansHebrew-VariableFont_wdth,wght.woff2", TT_UCR_HEBREW, 0, 0, 0},
+    {"NotoSansJavanese-Regular.woff2", 0x80000003L, 0x2000L, 0, 0},
+    {"NotoSansKannada-VariableFont_wdth,wght.woff2", TT_UCR_KANNADA, 0, 0, 0},
+    {"NotoSansMalayalam-VariableFont_wdth,wght.woff2", TT_UCR_MALAYALAM, 0, 0, 0},
+    {"NotoSansMath-Regular.woff2", 0, TT_UCR_MATHEMATICAL_OPERATORS, 0, 0},
+    {"NotoSansMyanmar-Regular.woff2", 0, 0, TT_UCR_MYANMAR, 0},
+    {"NotoSansSymbols-VariableFont_wght.woff2", 0x3L, 0x200E4B4L, 0, 0},
+    {"NotoSansSymbols2-Regular.woff2", 0x80000003L, 0x200E3E4L, 0x40020L, 0x580A048L},
+    {"NotoSansTamil-VariableFont_wdth,wght.woff2", TT_UCR_TAMIL, 0, 0, 0},
+    {"NotoSansTelugu-VariableFont_wdth,wght.woff2", TT_UCR_TELUGU, 0, 0, 0},
+    {"NotoSansThai-VariableFont_wdth,wght.woff2", TT_UCR_THAI, 0, 0, 0},
+};
+
+/* Create a new font from filename OR from passed memory pointer. */
+static FontBLF *blf_font_new_ex(const char *name,
+                                const char *filepath,
+                                const unsigned char *mem,
+                                const size_t mem_size)
+{
+  FontBLF *font = (FontBLF *)MEM_callocN(sizeof(FontBLF), "blf_font_new");
+
+  font->name = BLI_strdup(name);
+  font->filepath = filepath ? BLI_strdup(filepath) : NULL;
+  if (mem) {
+    font->mem = (void *)mem;
+    font->mem_size = mem_size;
+  }
+  blf_font_fill(font);
+
+  /* If we have static details about this font we don't need to load the Face. */
+  const eFaceDetails *static_details = NULL;
+  char filename[256];
+  for (int i = 0; i < (int)ARRAY_SIZE(static_face_details); i++) {
+    BLI_split_file_part(font->filepath, filename, sizeof(filename));
+    if (STREQ(static_face_details[i].name, filename)) {
+      static_details = &static_face_details[i];
+      font->UnicodeRanges[0] = static_details->coverage1;
+      font->UnicodeRanges[1] = static_details->coverage2;
+      font->UnicodeRanges[2] = static_details->coverage3;
+      font->UnicodeRanges[3] = static_details->coverage4;
+      break;
+    }
+  }
+
+  if (!static_details) {
+    if (!blf_ensure_face(font)) {
+      blf_font_free(font);
+      return NULL;
+    }
+  }
+
+  /* Detect "Last resort" fonts. They have everything. Usually except last 5 bits.  */
+  if (font->UnicodeRanges[0] == 0xffffffffU && font->UnicodeRanges[1] == 0xffffffffU &&
+      font->UnicodeRanges[2] == 0xffffffffU && font->UnicodeRanges[3] >= 0x7FFFFFFU) {
+    font->flags |= BLF_LAST_RESORT;
+  }
+
   return font;
 }
 
-void blf_font_attach_from_mem(FontBLF *font, const unsigned char *mem, int mem_size)
+FontBLF *blf_font_new(const char *name, const char *filename)
+{
+  return blf_font_new_ex(name, filename, NULL, 0);
+}
+
+FontBLF *blf_font_new_from_mem(const char *name, const unsigned char *mem, const size_t mem_size)
+{
+  return blf_font_new_ex(name, NULL, mem, mem_size);
+}
+
+void blf_font_attach_from_mem(FontBLF *font, const unsigned char *mem, const size_t mem_size)
 {
   FT_Open_Args open;
 
   open.flags = FT_OPEN_MEMORY;
   open.memory_base = (const FT_Byte *)mem;
-  open.memory_size = mem_size;
+  open.memory_size = (FT_Long)mem_size;
   FT_Attach_Stream(font->face, &open);
 }
 
-FontBLF *blf_font_new_from_mem(const char *name, const unsigned char *mem, int mem_size)
-{
-  FontBLF *font;
-  FT_Error err;
-
-  font = (FontBLF *)MEM_callocN(sizeof(FontBLF), "blf_font_new_from_mem");
-  err = FT_New_Memory_Face(ft_lib, mem, mem_size, 0, &font->face);
-  if (err) {
-    MEM_freeN(font);
-    return NULL;
-  }
-
-  err = FT_Select_Charmap(font->face, ft_encoding_unicode);
-  if (err) {
-    printf("Can't set the unicode character map!\n");
-    FT_Done_Face(font->face);
-    MEM_freeN(font);
-    return NULL;
-  }
-
-  if (FT_HAS_MULTIPLE_MASTERS(font->face)) {
-    FT_Get_MM_Var(font->face, &(font->variations));
-  }
-
-  font->name = BLI_strdup(name);
-  font->filepath = NULL;
-  blf_font_fill(font);
-  return font;
-}
-
 void blf_font_free(FontBLF *font)
 {
   blf_glyph_cache_clear(font);
@@ -1378,7 +1453,9 @@ void blf_font_free(FontBLF *font)
     FT_Done_MM_Var(ft_lib, font->variations);
   }
 
-  FT_Done_Face(font->face);
+  if (font->face) {
+    FT_Done_Face(font->face);
+  }
   if (font->filepath) {
     MEM_freeN(font->filepath);
   }
@@ -1396,6 +1473,10 @@ void blf_font_free(FontBLF *font)
 
 bool blf_font_size(FontBLF *font, float size, unsigned int dpi)
 {
+  if (!blf_ensure_face(font)) {
+    return false;
+  }
+
   /* FreeType uses fixed-point integers in 64ths. */
   FT_F26Dot6 ft_size = lroundf(size * 64.0f);
   /* Adjust our new size to be on even 64ths. */

source/blender/blenfont/intern/blf_font_default.c

@@ -66,8 +66,6 @@ void BLF_load_font_stack()
           }
           else {
             BLF_enable(font_id, BLF_DEFAULT);
-            /* TODO: FontBLF will later load FT_Face on demand.  When this is in
-             * place we can drop this face now since we have all needed data. */
           }
         }
       }

source/blender/blenfont/intern/blf_glyph.c

@@ -584,16 +584,22 @@ static FT_UInt blf_glyph_index_from_charcode(FontBLF **font, const uint charcode
       continue;
     }
     if (coverage_bit < 0 || blf_font_has_coverage_bit(f, coverage_bit)) {
-      glyph_index = FT_Get_Char_Index(f->face, charcode);
-      if (glyph_index) {
-        *font = f;
-        return glyph_index;
+      if (blf_ensure_face(f)) {
+        glyph_index = FT_Get_Char_Index(f->face, charcode);
+        if (glyph_index) {
+          *font = f;
+          return glyph_index;
+        }
       }
     }
   }
 
+#ifdef DEBUG
+  printf("Unicode character U+%04X not found in loaded fonts. \n", charcode);
+#endif
+
   /* Not found in the stack, return from Last Resort if there is one. */
-  if (last_resort) {
+  if (last_resort && blf_ensure_face(last_resort)) {
     glyph_index = FT_Get_Char_Index(last_resort->face, charcode);
     if (glyph_index) {
       *font = last_resort;

source/blender/blenfont/intern/blf_internal.h

@@ -39,12 +39,14 @@ void blf_font_exit(void);
 
 bool blf_font_id_is_valid(int fontid);
 
+bool blf_ensure_face(struct FontBLF *font);
+
 void blf_draw_buffer__start(struct FontBLF *font);
 void blf_draw_buffer__end(void);
 
 struct FontBLF *blf_font_new(const char *name, const char *filepath);
-struct FontBLF *blf_font_new_from_mem(const char *name, const unsigned char *mem, int mem_size);
-void blf_font_attach_from_mem(struct FontBLF *font, const unsigned char *mem, int mem_size);
+struct FontBLF *blf_font_new_from_mem(const char *name, const unsigned char *mem, size_t mem_size);
+void blf_font_attach_from_mem(struct FontBLF *font, const unsigned char *mem, size_t mem_size);
 
 /**
  * Change font's output size. Returns true if successful in changing the size.

source/blender/blenfont/intern/blf_internal_types.h

@@ -240,9 +240,13 @@ typedef struct FontBLF {
   /* # of times this font was loaded */
   unsigned int reference_count;
 
-  /** File-path or NULL. */
+  /* Full path to font file or NULL if from memory. */
   char *filepath;
 
+  /* Pointer to in-memory font, or NULL if from file. */
+  void *mem;
+  size_t mem_size;
+
   /* Copied from the SFNT OS/2 table. Bit flags for unicode blocks and ranges
    * considered "functional". Cached here because face might not always exist.
    * See: https://docs.microsoft.com/en-us/typography/opentype/spec/os2#ur */

source/blender/blenkernel/BKE_idprop.h

@@ -7,6 +7,7 @@
  */
 
 #include "BLI_compiler_attrs.h"
+#include "BLI_sys_types.h"
 
 #ifdef __cplusplus
 extern "C" {
@@ -19,6 +20,7 @@ struct BlendWriter;
 struct ID;
 struct IDProperty;
 struct IDPropertyUIData;
+struct Library;
 
 typedef union IDPropertyTemplate {
   int i;
@@ -318,7 +320,7 @@ void IDP_BlendReadData_impl(struct BlendDataReader *reader,
                             struct IDProperty **prop,
                             const char *caller_func_id);
 #define IDP_BlendDataRead(reader, prop) IDP_BlendReadData_impl(reader, prop, __func__)
-void IDP_BlendReadLib(struct BlendLibReader *reader, struct IDProperty *prop);
+void IDP_BlendReadLib(struct BlendLibReader *reader, struct Library *lib, struct IDProperty *prop);
 void IDP_BlendReadExpand(struct BlendExpander *expander, struct IDProperty *prop);
 
 typedef enum eIDPropertyUIDataType {

source/blender/blenkernel/BKE_idprop.hh

@@ -45,6 +45,9 @@ std::unique_ptr<IDProperty, IDPropertyDeleter> create(StringRefNull prop_name, d
 std::unique_ptr<IDProperty, IDPropertyDeleter> create(StringRefNull prop_name,
                                                       const StringRefNull value);
 
+/** \brief Allocate a new IDProperty of type IDP_ID, set its name and value. */
+std::unique_ptr<IDProperty, IDPropertyDeleter> create(StringRefNull prop_name, ID *id);
+
 /**
  * \brief Allocate a new IDProperty of type IDP_ARRAY and subtype IDP_INT.
  *

source/blender/blenkernel/BKE_node.h

@@ -374,6 +374,9 @@ typedef struct bNodeTreeType {
   int type;        /* type identifier */
   char idname[64]; /* identifier name */
 
+  /* The ID name of group nodes for this type. */
+  char group_idname[64];
+
   char ui_name[64];
   char ui_description[256];
   int ui_icon;

source/blender/blenkernel/BKE_screen.h

@@ -291,7 +291,7 @@ enum {
 
 /* Draw an item in the uiList */
 typedef void (*uiListDrawItemFunc)(struct uiList *ui_list,
-                                   struct bContext *C,
+                                   const struct bContext *C,
                                    struct uiLayout *layout,
                                    struct PointerRNA *dataptr,
                                    struct PointerRNA *itemptr,
@@ -303,12 +303,12 @@ typedef void (*uiListDrawItemFunc)(struct uiList *ui_list,
 
 /* Draw the filtering part of an uiList */
 typedef void (*uiListDrawFilterFunc)(struct uiList *ui_list,
-                                     struct bContext *C,
+                                     const struct bContext *C,
                                      struct uiLayout *layout);
 
 /* Filter items of an uiList */
 typedef void (*uiListFilterItemsFunc)(struct uiList *ui_list,
-                                      struct bContext *C,
+                                      const struct bContext *C,
                                       struct PointerRNA *,
                                       const char *propname);