Merge remote-tracking branch 'b/main'

UPBGE · Feb 6, 2025 · 5c70c40 · 5c70c40
2 parents 23aa3f3 + f3603b4
commit 5c70c40
Show file tree

Hide file tree

Showing 178 changed files with 2,379 additions and 1,519 deletions.
diff --git a/build_files/build_environment/cmake/versions.cmake b/build_files/build_environment/cmake/versions.cmake
@@ -1152,15 +1152,15 @@ set(UNIFIED_MEMORY_FRAMEWORK_COPYRIGHT "Copyright (C) 2023-2024 Intel Corporatio
 # compiler, the versions used are taken from the following location
 # https://github.com/intel/intel-graphics-compiler/releases
 
-set(IGC_VERSION 1.0.17384.29)
-set(IGC_URI https://github.com/intel/intel-graphics-compiler/archive/refs/tags/igc-${IGC_VERSION}.tar.gz)
-set(IGC_HASH de7b1ba9cb1369f9aa26343bb4ad1ac7e5cb5f9f4517071f25d853e46cae6195)
+set(IGC_VERSION 2.1.14)
+set(IGC_URI https://github.com/intel/intel-graphics-compiler/archive/refs/tags/v${IGC_VERSION}.tar.gz)
+set(IGC_HASH ebc39f8238dec2983c227d9db204db333b439ba5cd100e0a819a0208ac8066ca)
 set(IGC_HASH_TYPE SHA256)
-set(IGC_FILE igc-${IGC_VERSION}.tar.gz)
+set(IGC_FILE ${IGC_VERSION}.tar.gz)
 set(IGC_NAME IGC)
 set(IGC_HOMEPAGE https://github.com/intel/intel-graphics-compiler)
 set(IGC_LICENSE SPDX:MIT)
-set(IGC_COPYRIGHT "Copyright (C) 2019-2021 Intel Corporation")
+set(IGC_COPYRIGHT "Copyright (C) 2019-2024 Intel Corporation")
 
 set(IGC_LLVM_VERSION llvmorg-14.0.5)
 set(IGC_LLVM_URI https://github.com/llvm/llvm-project/archive/refs/tags/${IGC_LLVM_VERSION}.tar.gz)
@@ -1181,24 +1181,24 @@ set(IGC_LLVM_COPYRIGHT "Copyright (c) 2003-2019 University of Illinois at Urbana
 #
 # WARNING WARNING WARNING
 
-set(IGC_OPENCL_CLANG_VERSION 470cf0018e1ef6fc92eda1356f5f31f7da452abc)
+set(IGC_OPENCL_CLANG_VERSION v14.0.1)
 set(IGC_OPENCL_CLANG_URI https://github.com/intel/opencl-clang/archive/${IGC_OPENCL_CLANG_VERSION}.tar.gz)
-set(IGC_OPENCL_CLANG_HASH fa410e0b4cc5b3fc3262e3b6aaace3543207a20ecd004f48dfec9a970f1fe4e2)
+set(IGC_OPENCL_CLANG_HASH 77a023e311f825bbdc242fe9b397bd68bac23a88e68142b4919fa322c6c70392)
 set(IGC_OPENCL_CLANG_HASH_TYPE SHA256)
 set(IGC_OPENCL_CLANG_FILE opencl-clang-${IGC_OPENCL_CLANG_VERSION}.tar.gz)
 set(IGC_OPENCL_CLANG_HOMEPAGE https://github.com/intel/opencl-clang/)
 set(IGC_OPENCL_CLANG_LICENSE SPDX:Apache-2.0 WITH LLVM-exception)
 set(IGC_OPENCL_CLANG_COPYRIGHT "Copyright (c) Intel Corporation (2009-2017).")
 
-set(IGC_VCINTRINSICS_VERSION v0.19.0)
-set(IGC_VCINTRINSICS_URI https://github.com/intel/vc-intrinsics/archive/refs/tags/${IGC_VCINTRINSICS_VERSION}.tar.gz)
-set(IGC_VCINTRINSICS_HASH b708df2fddc9fcb2cac5d6f26870f2e105f8395c0208ecd8acc38cbf175aee52)
+set(IGC_VCINTRINSICS_VERSION 0.20.1)
+set(IGC_VCINTRINSICS_URI https://github.com/intel/vc-intrinsics/archive/refs/tags/v${IGC_VCINTRINSICS_VERSION}.tar.gz)
+set(IGC_VCINTRINSICS_HASH 01a68a3c8299075ad2b094264bef5109d660f1d53faf614f4f95745ad05f4e3e)
 set(IGC_VCINTRINSICS_HASH_TYPE SHA256)
 set(IGC_VCINTRINSICS_FILE vc-intrinsics-${IGC_VCINTRINSICS_VERSION}.tar.gz)
 set(IGC_VCINTRINSICS_NAME "VC Intrinsics")
 set(IGC_VCINTRINSICS_HOMEPAGE https://github.com/intel/vc-intrinsics)
 set(IGC_VCINTRINSICS_LICENSE SPDX:MIT)
-set(IGC_VCINTRINSICS_COPYRIGHT "Copyright (c) 2019 Intel Corporation")
+set(IGC_VCINTRINSICS_COPYRIGHT "Copyright (C) 2020-2021 Intel Corporation")
 
 set(IGC_SPIRV_HEADERS_VERSION vulkan-sdk-1.3.275.0)
 set(IGC_SPIRV_HEADERS_URI https://github.com/KhronosGroup/SPIRV-Headers/archive/refs/tags/${IGC_SPIRV_HEADERS_VERSION}.tar.gz)
@@ -1220,9 +1220,9 @@ set(IGC_SPIRV_TOOLS_HOMEPAGE https://github.com/KhronosGroup/SPIRV-Tools/)
 set(IGC_SPIRV_TOOLS_LICENSE SPDX:Apache-2.0)
 set(IGC_SPIRV_TOOLS_COPYRIGHT "Copyright (c) 2015-2016 The Khronos Group Inc.")
 
-set(IGC_SPIRV_TRANSLATOR_VERSION 2823e7052b7999c10fff63bc8089e5aa205716f4)
+set(IGC_SPIRV_TRANSLATOR_VERSION 1cfcf1b5ec2ec9c52fd56dc7cb47dcf12125e8db)
 set(IGC_SPIRV_TRANSLATOR_URI https://github.com/KhronosGroup/SPIRV-LLVM-Translator/archive/${IGC_SPIRV_TRANSLATOR_VERSION}.tar.gz)
-set(IGC_SPIRV_TRANSLATOR_HASH 39eb3e033a0a1f5c69622d6a0b87e296b4e090d2f613f1014ee6fedcc2d3ca83)
+set(IGC_SPIRV_TRANSLATOR_HASH 5384344b47b67e3811419599c138284de5490304030096f3bc10271915a8ad76)
 set(IGC_SPIRV_TRANSLATOR_HASH_TYPE SHA256)
 set(IGC_SPIRV_TRANSLATOR_FILE SPIR-V-Translator-${IGC_SPIRV_TRANSLATOR_VERSION}.tar.gz)
 set(IGC_SPIRV_TRANSLATOR_NAME "LLVM/SPIR-V Bi-Directional Translator")
@@ -1243,9 +1243,9 @@ http://llvm.org
 ### Intel Graphics Compiler DEPS END ###
 ########################################
 
-set(GMMLIB_VERSION intel-gmmlib-22.4.1)
+set(GMMLIB_VERSION intel-gmmlib-22.5.3)
 set(GMMLIB_URI https://github.com/intel/gmmlib/archive/refs/tags/${GMMLIB_VERSION}.tar.gz)
-set(GMMLIB_HASH 451fbe2eac26533a896ca0da0356354ecc38680f273fce7d121c6a22251ed21e)
+set(GMMLIB_HASH a8ccba4a71319b66a901df016998a76745ce9624032b6555517607c5bbb2b4e8)
 set(GMMLIB_HASH_TYPE SHA256)
 set(GMMLIB_NAME "Intel(R) Graphics Memory Management Library")
 set(GMMLIB_FILE ${GMMLIB_VERSION}.tar.gz)
@@ -1257,14 +1257,14 @@ Copyright (c) 2016 Gabi Melman.
 Copyright 2008, Google Inc. All rights reserved.
 ]=])
 
-set(OCLOC_VERSION 24.31.30508.9)
+set(OCLOC_VERSION 24.45.31740.15)
 set(OCLOC_URI https://github.com/intel/compute-runtime/archive/refs/tags/${OCLOC_VERSION}.tar.gz)
-set(OCLOC_HASH 7c2b5708e996fc9e61997f1821d9be1e0fd43c9f29cfe3fea383a01d9aa92868)
+set(OCLOC_HASH f0e3b3b038b585753530a8f1b00a8e71b045ee007b8392aaa986ccad92cf5123)
 set(OCLOC_HASH_TYPE SHA256)
 set(OCLOC_FILE ocloc-${OCLOC_VERSION}.tar.gz)
 set(OCLOC_HOMEPAGE https://github.com/intel/compute-runtime)
 set(OCLOC_LICENSE SPDX:MIT)
-set(OCLOC_COPYRIGHT "Copyright (C) 2021 Intel Corporation")
+set(OCLOC_COPYRIGHT "Copyright (C) 2018-2023 Intel Corporation")
 
 set(AOM_VERSION 3.4.0)
 set(AOM_URI https://storage.googleapis.com/aom-releases/libaom-${AOM_VERSION}.tar.gz)

diff --git a/build_files/build_environment/patches/ocloc.diff b/build_files/build_environment/patches/ocloc.diff
@@ -1,14 +1,13 @@
 diff --git a/shared/offline_compiler/source/CMakeLists.txt b/shared/offline_compiler/source/CMakeLists.txt
-index 39af3af46c..2e0b37e35f 100644
+index 9d16ca0e736..88e236acd1a 100644
 --- a/shared/offline_compiler/source/CMakeLists.txt
 +++ b/shared/offline_compiler/source/CMakeLists.txt
-@@ -450,7 +450,7 @@ endif()
- if(WIN32)
-   set(LIBOCLOC_OUTPUT_NAME_SUFFIX ${OCLOC_BITNESS_SUFFIX})
- elseif(UNIX)
--  set(OCLOC_OUTPUT_NAME_SUFFIX "-${NEO_OCLOC_VERSION_MAJOR}.${NEO_OCLOC_VERSION_MINOR}")
-+#  set(OCLOC_OUTPUT_NAME_SUFFIX "-${NEO_OCLOC_VERSION_MAJOR}.${NEO_OCLOC_VERSION_MINOR}")
-   if(NOT NEO_CURRENT_PLATFORMS_SUPPORT)
-     set(LIBOCLOC_OUTPUT_NAME_SUFFIX "-${NEO_OCLOC_VERSION_MAJOR}.${NEO_OCLOC_VERSION_MINOR}")
+@@ -463,7 +463,7 @@ elseif(UNIX)
+     set(NEO_OCLOC_FORMER_LIB_NAME ${CMAKE_SHARED_LIBRARY_PREFIX}${OCLOC_NAME}${NEO__LEGACY_NAME_SUFFIX}${CMAKE_SHARED_LIBRARY_SUFFIX})
    endif()
-
+
+-  set(OCLOC_OUTPUT_NAME_SUFFIX "-${NEO_OCL_VERSION_MAJOR}.${NEO_OCL_VERSION_MINOR}.${NEO_OCLOC_VERSION_MODE}")
++  #set(OCLOC_OUTPUT_NAME_SUFFIX "-${NEO_OCL_VERSION_MAJOR}.${NEO_OCL_VERSION_MINOR}.${NEO_OCLOC_VERSION_MODE}")
+ endif()
+ set_target_properties(${OCLOC_NAME} PROPERTIES OUTPUT_NAME ${OCLOC_NAME}${OCLOC_OUTPUT_NAME_SUFFIX})
+ set_target_properties(${OCLOC_NAME}_lib PROPERTIES OUTPUT_NAME ${OCLOC_NAME}${LIBOCLOC_OUTPUT_NAME_SUFFIX})
diff --git a/doc/doxygen/Doxyfile b/doc/doxygen/Doxyfile
@@ -38,7 +38,7 @@ PROJECT_NAME           = Blender
 # could be handy for archiving the generated documentation or if some version
 # control system is used.
 
-PROJECT_NUMBER         = V4.4
+PROJECT_NUMBER         = V4.5
 
 # Using the PROJECT_BRIEF tag one can provide an optional one line description
 # for a project that appears at the top of each page and should give viewer a

diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py
@@ -1792,7 +1792,7 @@ def _draw_devices(self, layout, device_type, devices):
                     col.label(text=rpt_("and Windows driver version %s or newer") % driver_version,
                               icon='BLANK1', translate=False)
                 elif sys.platform.startswith("linux"):
-                    driver_version = "XX.XX.29735.20"
+                    driver_version = "XX.XX.31740.15"
                     col.label(
                         text=rpt_("Requires Intel GPU with Xe-HPG architecture and"),
                         icon='BLANK1',

diff --git a/intern/cycles/blender/sync.cpp b/intern/cycles/blender/sync.cpp
@@ -187,7 +187,7 @@ void BlenderSync::sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d
             const map<void *, set<BL::ID>>::const_iterator instance_geometries =
                 instance_geometries_by_object.find(b_ob.ptr.data);
             if (instance_geometries != instance_geometries_by_object.end()) {
-              for (BL::ID const geometry : instance_geometries->second) {
+              for (BL::ID const &geometry : instance_geometries->second) {
                 geometry_map.set_recalc(geometry);
               }
             }

diff --git a/intern/cycles/device/oneapi/device_impl.cpp b/intern/cycles/device/oneapi/device_impl.cpp
@@ -1067,7 +1067,7 @@ static const int lowest_supported_driver_version_win = 1015730;
  * This information is returned by `ocloc query OCL_DRIVER_VERSION`.*/
 static const int lowest_supported_driver_version_neo = 29550;
 #  else
-static const int lowest_supported_driver_version_neo = 29735;
+static const int lowest_supported_driver_version_neo = 31740;
 #  endif
 
 int parse_driver_build_version(const sycl::device &device)

diff --git a/intern/cycles/kernel/closure/bsdf_principled_hair_huang.h b/intern/cycles/kernel/closure/bsdf_principled_hair_huang.h
@@ -38,7 +38,7 @@ struct HuangHairExtra {
   float pixel_coverage;
 
   /* Valid integration interval. */
-  float gamma_m_min, gamma_m_max;
+  Interval<float> gamma_m;
 };
 
 struct HuangHairBSDF {
@@ -73,6 +73,8 @@ static_assert(sizeof(ShaderClosure) >= sizeof(HuangHairExtra), "HuangHairExtra i
 /** \name Hair coordinate system utils.
  * \{ */
 
+/* TODO(weizhen): add a namespace `huanghair` for the below utility functions. */
+
 /* Returns `sin(theta)` of the given direction. */
 ccl_device_inline float sin_theta(const float3 w)
 {
@@ -124,10 +126,16 @@ ccl_device_inline float2 dir_sph(const float3 w)
   return make_float2(dir_theta(w), dir_phi(w));
 }
 
+/* Check whether the hair cross-section is circular. */
+ccl_device_inline bool is_circular(const float b)
+{
+  return b == 1.0f;
+}
+
 /* Conversion between `gamma` and `phi`. Notations see Figure 5 in the paper. */
 ccl_device_inline float to_phi(const float gamma, const float b)
 {
-  if (b == 1.0f) {
+  if (is_circular(b)) {
     return gamma;
   }
   float sin_gamma;
@@ -138,7 +146,7 @@ ccl_device_inline float to_phi(const float gamma, const float b)
 
 ccl_device_inline float to_gamma(const float phi, const float b)
 {
-  if (b == 1.0f) {
+  if (is_circular(b)) {
     return phi;
   }
   float sin_phi;
@@ -152,7 +160,7 @@ ccl_device_inline float to_gamma(const float phi, const float b)
  * Also, make use of `r = sqrt(sqr(cos_phi_i) + sqr(b * sin_phi_i))` to pre-map `h` to [-1, 1]. */
 ccl_device_inline float h_to_gamma(const float h_div_r, const float b, const float3 wi)
 {
-  return (b == 1.0f) ? -asinf(h_div_r) : atan2f(wi.z, -b * wi.x) - acosf(-h_div_r);
+  return is_circular(b) ? -asinf(h_div_r) : atan2f(wi.z, -b * wi.x) - acosf(-h_div_r);
 }
 
 /* Compute the coordinate on the ellipse, given `gamma` and the aspect ratio between the minor axis
@@ -178,7 +186,7 @@ ccl_device_inline float3 sphg_dir(const float theta, const float gamma, const fl
   fast_sincosf(theta, &sin_theta, &cos_theta);
   fast_sincosf(gamma, &sin_gamma, &cos_gamma);
 
-  if (b == 1.0f || fabsf(cos_gamma) < 1e-6f) {
+  if (is_circular(b) || fabsf(cos_gamma) < 1e-6f) {
     sin_phi = sin_gamma;
     cos_phi = cos_gamma;
   }
@@ -349,7 +357,7 @@ ccl_device Spectrum bsdf_hair_huang_eval_r(KernelGlobals kg,
   /* Maximal sample resolution. */
   float res = roughness * 0.7f;
 
-  const float gamma_m_range = bsdf->extra->gamma_m_max - bsdf->extra->gamma_m_min;
+  const float gamma_m_range = bsdf->extra->gamma_m.length();
 
   /* Number of intervals should be even. */
   const size_t intervals = 2 * (size_t)ceilf(gamma_m_range / res * 0.5f);
@@ -360,7 +368,7 @@ ccl_device Spectrum bsdf_hair_huang_eval_r(KernelGlobals kg,
   /* Integrate using Composite Simpson's 1/3 rule. */
   float integral = 0.0f;
   for (size_t i = 0; i <= intervals; i++) {
-    const float gamma_m = bsdf->extra->gamma_m_min + i * res;
+    const float gamma_m = bsdf->extra->gamma_m.min + i * res;
     const float3 wm = sphg_dir(bsdf->tilt, gamma_m, b);
 
     if (microfacet_visible(wi, wo, make_float3(wm.x, 0.0f, wm.z), wh)) {
@@ -407,7 +415,6 @@ ccl_device Spectrum bsdf_hair_huang_eval_residual(KernelGlobals kg,
 
   /* Get minor axis, assuming major axis is 1. */
   const float b = bsdf->aspect_ratio;
-  const bool is_circular = (b == 1.0f);
 
   const Spectrum mu_a = bsdf->sigma;
   const float eta = bsdf->eta;
@@ -418,7 +425,7 @@ ccl_device Spectrum bsdf_hair_huang_eval_residual(KernelGlobals kg,
   const float sqrt_roughness = sqrtf(roughness);
 
   float res = roughness * 0.8f;
-  const float gamma_m_range = bsdf->extra->gamma_m_max - bsdf->extra->gamma_m_min;
+  const float gamma_m_range = bsdf->extra->gamma_m.length();
   const size_t intervals = 2 * (size_t)ceilf(gamma_m_range / res * 0.5f);
   res = gamma_m_range / intervals;
 
@@ -427,7 +434,7 @@ ccl_device Spectrum bsdf_hair_huang_eval_residual(KernelGlobals kg,
   Spectrum S_trrt = zero_spectrum();
   for (size_t i = 0; i <= intervals; i++) {
 
-    const float gamma_mi = bsdf->extra->gamma_m_min + i * res;
+    const float gamma_mi = bsdf->extra->gamma_m.min + i * res;
 
     const float3 wmi = sphg_dir(bsdf->tilt, gamma_mi, b);
     const float3 wmi_ = sphg_dir(0.0f, gamma_mi, b);
@@ -464,7 +471,7 @@ ccl_device Spectrum bsdf_hair_huang_eval_residual(KernelGlobals kg,
     const float weight = (i == 0 || i == intervals) ? 0.5f : (i % 2 + 1);
 
     const Spectrum A_t = exp(mu_a / cos_theta(wt) *
-                             (is_circular ?
+                             (is_circular(b) ?
                                   2.0f * cosf(gamma_mi - phi_t) :
                                   -len(to_point(gamma_mi, b) - to_point(gamma_mt + M_PI_F, b))));
 
@@ -544,7 +551,7 @@ ccl_device Spectrum bsdf_hair_huang_eval_residual(KernelGlobals kg,
       const float D3 = bsdf_D<MicrofacetType::GGX>(roughness2, cos_mh3);
 
       const Spectrum A_tr = exp(mu_a / cos_theta(wtr) *
-                                -(is_circular ?
+                                -(is_circular(b) ?
                                       2.0f * fabsf(cosf(phi_tr - gamma_mt)) :
                                       len(to_point(gamma_mtr, b) - to_point(gamma_mt, b))));
 
@@ -606,7 +613,6 @@ ccl_device int bsdf_hair_huang_sample(const KernelGlobals kg,
 
   /* Get minor axis, assuming major axis is 1. */
   const float b = bsdf->aspect_ratio;
-  const bool is_circular = (b == 1.0f);
 
   /* Sample `h` for farfield model, as the computed intersection might have numerical issues. */
   const float h_div_r = is_nearfield(bsdf) ? bsdf->h / bsdf->extra->radius :
@@ -673,7 +679,7 @@ ccl_device int bsdf_hair_huang_sample(const KernelGlobals kg,
   if (cos_mi2 > 0.0f && microfacet_visible(-wt, wmi_, wh1) && microfacet_visible(-wt, wmt_, wh2)) {
     const Spectrum mu_a = bsdf->sigma;
     const Spectrum A_t = exp(mu_a / cos_theta(wt) *
-                             (is_circular ?
+                             (is_circular(b) ?
                                   2.0f * cosf(phi_t - gamma_mi) :
                                   -len(to_point(gamma_mi, b) - to_point(gamma_mt + M_PI_F, b))));
 
@@ -704,7 +710,7 @@ ccl_device int bsdf_hair_huang_sample(const KernelGlobals kg,
     const float cos_mi3 = dot(wmtr, wtr);
     if (cos_mi3 > 0.0f) {
       const Spectrum A_tr = exp(mu_a / cos_theta(wtr) *
-                                -(is_circular ?
+                                -(is_circular(b) ?
                                       2.0f * fabsf(cosf(phi_tr - gamma_mt)) :
                                       len(to_point(gamma_mt, b) - to_point(gamma_mtr, b))));
 
@@ -826,16 +832,16 @@ ccl_device Spectrum bsdf_hair_huang_eval(KernelGlobals kg,
   }
   const float r = bsdf->extra->radius;
   const float b = bsdf->aspect_ratio;
-  const float phi_i = (b == 1.0f) ? 0.0f : dir_phi(local_I);
-  float gamma_m_min = to_gamma(phi_i - half_span, b);
-  float gamma_m_max = to_gamma(phi_i + half_span, b);
-  if (gamma_m_max < gamma_m_min) {
-    gamma_m_max += M_2PI_F;
+  const float phi_i = is_circular(b) ? 0.0f : dir_phi(local_I);
+
+  Interval<float> gamma_m = {to_gamma(phi_i - half_span, b), to_gamma(phi_i + half_span, b)};
+  if (gamma_m.max < gamma_m.min) {
+    gamma_m.max += M_2PI_F;
   }
 
   /* Prevent numerical issues at the boundary. */
-  gamma_m_min += 1e-3f;
-  gamma_m_max -= 1e-3f;
+  gamma_m.min += 1e-3f;
+  gamma_m.max -= 1e-3f;
 
   /* Length of the integral interval. */
   float dh = 2.0f * r;
@@ -845,42 +851,41 @@ ccl_device Spectrum bsdf_hair_huang_eval(KernelGlobals kg,
      * Inspired by [An Efficient and Practical Near and Far Field Fur Reflectance Model]
      * (https://sites.cs.ucsb.edu/~lingqi/publications/paper_fur2.pdf) by Ling-Qi Yan, Henrik Wann
      * Jensen and Ravi Ramamoorthi. */
-    const float h_max = min(bsdf->h + bsdf->extra->pixel_coverage, r);
-    const float h_min = max(bsdf->h - bsdf->extra->pixel_coverage, -r);
+    const float half_pixel = bsdf->extra->pixel_coverage;
+    const Interval<float> h = intervals_intersection(Interval<float>{-r, r},
+                                                     {bsdf->h - half_pixel, bsdf->h + half_pixel});
 
     /* At the boundaries the hair might not cover the whole pixel. */
-    dh = h_max - h_min;
+    dh = h.length();
 
-    float nearfield_gamma_min = h_to_gamma(h_max / r, bsdf->aspect_ratio, local_I);
-    float nearfield_gamma_max = h_to_gamma(h_min / r, bsdf->aspect_ratio, local_I);
+    Interval<float> nearfield_gamma = {h_to_gamma(h.max / r, bsdf->aspect_ratio, local_I),
+                                       h_to_gamma(h.min / r, bsdf->aspect_ratio, local_I)};
 
-    if (nearfield_gamma_max < nearfield_gamma_min) {
-      nearfield_gamma_max += M_2PI_F;
+    if (nearfield_gamma.max < nearfield_gamma.min) {
+      nearfield_gamma.max += M_2PI_F;
     }
 
     /* Wrap range to compute the intersection. */
-    if ((gamma_m_max - nearfield_gamma_min) > M_2PI_F) {
-      gamma_m_min -= M_2PI_F;
-      gamma_m_max -= M_2PI_F;
+    if ((gamma_m.max - nearfield_gamma.min) > M_2PI_F) {
+      gamma_m.min -= M_2PI_F;
+      gamma_m.max -= M_2PI_F;
     }
-    else if ((nearfield_gamma_max - gamma_m_min) > M_2PI_F) {
-      nearfield_gamma_min -= M_2PI_F;
-      nearfield_gamma_max -= M_2PI_F;
+    else if ((nearfield_gamma.max - gamma_m.min) > M_2PI_F) {
+      nearfield_gamma.min -= M_2PI_F;
+      nearfield_gamma.max -= M_2PI_F;
     }
 
-    gamma_m_min = fmaxf(gamma_m_min, nearfield_gamma_min);
-    gamma_m_max = fminf(gamma_m_max, nearfield_gamma_max);
+    gamma_m = intervals_intersection(gamma_m, nearfield_gamma);
   }
 
-  bsdf->extra->gamma_m_min = gamma_m_min;
-  bsdf->extra->gamma_m_max = gamma_m_max;
-
-  if (!(bsdf->extra->gamma_m_min < bsdf->extra->gamma_m_max)) {
+  if (gamma_m.is_empty()) {
     /* No overlap between the valid range and the visible range. Can happen at grazing `theta`
      * angles. */
     return zero_spectrum();
   }
 
+  bsdf->extra->gamma_m = gamma_m;
+
   const float projected_area = cos_theta(local_I) * dh;
 
   return (bsdf_hair_huang_eval_r(kg, sc, local_I, local_O) +