Draw: Cleanup the GLSL intersection code

Only store center and size. Skip resource finalize computation.

.gitea/default_merge_message/REBASE_TEMPLATE.md

@@ -1,5 +0,0 @@
-${CommitTitle}
-
-${CommitBody}
-
-Pull Request #${PullRequestIndex}

.gitea/default_merge_message/SQUASH_TEMPLATE.md

@@ -1,3 +0,0 @@
-${PullRequestTitle}
-
-Pull Request #${PullRequestIndex}

.gitea/issue_template.yaml

@@ -0,0 +1,45 @@
+name: Bug Report
+about: File a bug report
+labels:
+  - bug
+ref: master
+body:
+  - type: markdown
+    attributes:
+      value: |
+                ### First time bug reporting?
+                Read [these tips](https://wiki.blender.org/wiki/Process/Bug_Reports) and watch this **[How to Report a Bug](https://www.youtube.com/watch?v=JTD0OJq_rF4)** video to make a complete, valid bug report. Remember to write your bug report in **English**.
+
+                ### What not to report here
+                For feature requests, feedback, questions or issues building Blender, see [communication channels](https://wiki.blender.org/wiki/Communication/Contact#User_Feedback_and_Requests).
+
+                ### Please verify
+                * Always test with the latest official release from [blender.org](https://www.blender.org/) and daily build from [builder.blender.org](https://builder.blender.org/).
+                * Please use `Help > Report a Bug` in Blender to automatically fill system information and exact Blender version.
+                * Test [previous Blender versions](https://download.blender.org/release/) to find the latest version that was working as expected.
+                * Find steps to redo the bug consistently without any non-official add-ons, and include a **small and simple .blend file** to demonstrate the bug.
+                * If there are multiple bugs, make multiple bug reports.
+                * Sometimes, driver or software upgrades cause problems. On Windows, try a clean install of the graphics drivers.
+
+                ### Help the developers
+                Bug fixing is important, the developers will handle a report swiftly. For that reason, we need your help to carefully provide instructions that others can follow quickly. You do your half of the work, then we do our half!
+
+                If a report is tagged with Needs Information from User and it has no reply after a week, we will assume the issue is gone and close the report.
+
+  - type: textarea
+    attributes:
+      label: "Description"
+      value: |
+               **System Information**
+               Operating system:
+               Graphics card:
+
+               **Blender Version**
+               Broken: (example: 2.80, edbf15d3c044, master, 2018-11-28, as found on the splash screen)
+               Worked: (newest version of Blender that worked as expected)
+
+               **Short description of error**
+
+               **Exact steps for others to reproduce the error**
+               Based on the default startup or an attached .blend file (as simple as possible).
+

.gitea/issue_template/bug.yaml

@@ -1,44 +0,0 @@
-name: Bug Report
-about: File a bug report
-labels:
-  - "type::Report"
-  - "status::Needs Triage"
-  - "priority::Normal"
-body:
-  - type: markdown
-    attributes:
-      value: |
-                ### Instructions
-                First time reporting? See [tips](https://wiki.blender.org/wiki/Process/Bug_Reports).
-
-                * Use **Help > Report a Bug** in Blender to fill system information and exact Blender version.
-                * Test [daily builds](https://builder.blender.org/) to verify if the issue is already fixed.
-                * Test [previous versions](https://download.blender.org/release/) to find an older working version.
-                * For feature requests, feedback, questions or build issues, see [communication channels](https://wiki.blender.org/wiki/Communication/Contact#User_Feedback_and_Requests).
-                * If there are multiple bugs, make multiple bug reports.
-
-  - type: textarea
-    id: body
-    attributes:
-      label: "Description"
-      hide_label: true
-      value: |
-               **System Information**
-               Operating system:
-               Graphics card:
-
-               **Blender Version**
-               Broken: (example: 2.80, edbf15d3c044, master, 2018-11-28, as found on the splash screen)
-               Worked: (newest version of Blender that worked as expected)
-
-               **Short description of error**
-
-               **Exact steps for others to reproduce the error**
-               Based on the default startup or an attached .blend file (as simple as possible).
-
-  - type: markdown
-    attributes:
-      value: |
-                ### Help the developers
-
-                Bug fixing is important, the developers will handle reports swiftly. For that reason, carefully provide exact steps and a **small and simple .blend file** to reproduce the problem. You do your half of the work, then we do our half!

.gitea/issue_template/config.yaml

@@ -1 +0,0 @@
-blank_issues_enabled: false

.gitea/issue_template/design.yaml

@@ -1,10 +0,0 @@
-name: Design
-about: Create a design task (for developers only)
-labels:
-  - "type::Design"
-body:
-  - type: textarea
-    id: body
-    attributes:
-      label: "Description"
-      hide_label: true

.gitea/issue_template/todo.yaml

@@ -1,10 +0,0 @@
-name: To Do
-about: Create a to do task (for developers only)
-labels:
-  - "type::To Do"
-body:
-  - type: textarea
-    id: body
-    attributes:
-      label: "Description"
-      hide_label: true

.gitea/pull_request_template.md

@@ -0,0 +1,4 @@
+---
+name: Pull Request
+about: Submit a pull request
+---

.gitea/pull_request_template.yaml

@@ -1,17 +0,0 @@
-name: Pull Request
-about: Contribute code to Blender
-body:
-  - type: markdown
-    attributes:
-      value: |
-        ### Instructions
-
-        Guides to [contributing code](https://wiki.blender.org/index.php/Dev:Doc/Process/Contributing_Code) and effective [code review](https://wiki.blender.org/index.php/Dev:Doc/Tools/Code_Review).
-
-        By submitting code here, you agree that the code is (compatible with) GNU GPL v2 or later.
-
-  - type: textarea
-    id: body
-    attributes:
-      label: "Description"
-      hide_label: true

CMakeLists.txt

@@ -167,26 +167,14 @@ get_blender_version()
 option(WITH_BLENDER "Build blender (disable to build only the blender player)" ON)
 mark_as_advanced(WITH_BLENDER)
 
-if(WIN32)
-  option(WITH_BLENDER_THUMBNAILER "\
-Build \"BlendThumb.dll\" helper for Windows explorer integration to support extracting \
-thumbnails from `.blend` files."
-    ON
-  )
+if(APPLE)
+  # In future, can be used with `quicklookthumbnailing/qlthumbnailreply` to create file
+  # thumbnails for say Finder. Turn it off for now.
+  option(WITH_BLENDER_THUMBNAILER "Build \"blender-thumbnailer\" thumbnail extraction utility" OFF)
+elseif(WIN32)
+  option(WITH_BLENDER_THUMBNAILER "Build \"BlendThumb.dll\" helper for Windows explorer integration" ON)
 else()
-  set(_option_default ON)
-  if(APPLE)
-    # In future, can be used with `quicklookthumbnailing/qlthumbnailreply`
-    # to create file thumbnails for say Finder.
-    # Turn it off for now, even though it can build on APPLE, it's not likely to be useful.
-    set(_option_default OFF)
-  endif()
-  option(WITH_BLENDER_THUMBNAILER "\
-Build stand-alone \"blender-thumbnailer\" command-line thumbnail extraction utility, \
-intended for use by file-managers to extract PNG images from `.blend` files."
-    ${_option_default}
-  )
-  unset(_option_default)
+  option(WITH_BLENDER_THUMBNAILER "Build \"blender-thumbnailer\" thumbnail extraction utility" ON)
 endif()
 
 option(WITH_INTERNATIONAL "Enable I18N (International fonts and text)" ON)
@@ -226,19 +214,14 @@ option(WITH_BULLET        "Enable Bullet (Physics Engine)" ON)
 option(WITH_SYSTEM_BULLET "Use the systems bullet library (currently unsupported due to missing features in upstream!)" )
 mark_as_advanced(WITH_SYSTEM_BULLET)
 option(WITH_OPENCOLORIO   "Enable OpenColorIO color management" ON)
-
-set(_option_default ON)
 if(APPLE)
   # There's no OpenXR runtime in sight for macOS, neither is code well
   # tested there -> disable it by default.
-  set(_option_default OFF)
-endif()
-option(WITH_XR_OPENXR "Enable VR features through the OpenXR specification" ${_option_default})
-if(APPLE)
+  option(WITH_XR_OPENXR   "Enable VR features through the OpenXR specification" OFF)
   mark_as_advanced(WITH_XR_OPENXR)
+else()
+  option(WITH_XR_OPENXR   "Enable VR features through the OpenXR specification" ON)
 endif()
-unset(_option_default)
-
 option(WITH_GMP "Enable features depending on GMP (Exact Boolean)" ON)
 
 # Compositor
@@ -370,13 +353,12 @@ else()
   set(WITH_COREAUDIO OFF)
 endif()
 if(NOT WIN32)
-  set(_option_default ON)
   if(APPLE)
-    set(_option_default OFF)
+    option(WITH_JACK          "Enable JACK Support (http://www.jackaudio.org)" OFF)
+  else()
+    option(WITH_JACK          "Enable JACK Support (http://www.jackaudio.org)" ON)
   endif()
-  option(WITH_JACK "Enable JACK Support (http://www.jackaudio.org)" ${_option_default})
-  unset(_option_default)
-  option(WITH_JACK_DYNLOAD "Enable runtime dynamic JACK libraries loading" OFF)
+  option(WITH_JACK_DYNLOAD  "Enable runtime dynamic JACK libraries loading" OFF)
 else()
   set(WITH_JACK OFF)
 endif()
@@ -417,26 +399,6 @@ mark_as_advanced(WITH_SYSTEM_GLOG)
 # Freestyle
 option(WITH_FREESTYLE     "Enable Freestyle (advanced edges rendering)" ON)
 
-# Libraries.
-if(UNIX AND NOT APPLE)
-  # Optionally build without pre-compiled libraries.
-  # NOTE: this could be supported on all platforms however in practice UNIX is the only platform
-  # that has good support for detecting installed libraries.
-  option(WITH_LIBS_PRECOMPILED "\
-Detect and link against pre-compiled libraries (typically found under \"../lib/\"). \
-Disabling this option will use the system libraries although cached paths \
-that point to pre-compiled libraries will be left as-is."
-    ON
-  )
-  mark_as_advanced(WITH_LIBS_PRECOMPILED)
-
-  option(WITH_STATIC_LIBS "Try to link with static libraries, as much as possible, to make blender more portable across distributions" OFF)
-  if(WITH_STATIC_LIBS)
-    option(WITH_BOOST_ICU "Boost uses ICU library (required for linking with static Boost built with libicu)." OFF)
-    mark_as_advanced(WITH_BOOST_ICU)
-  endif()
-endif()
-
 # Misc
 if(WIN32 OR APPLE)
   option(WITH_INPUT_IME "Enable Input Method Editor (IME) for complex Asian character input" ON)
@@ -444,6 +406,11 @@ endif()
 option(WITH_INPUT_NDOF "Enable NDOF input devices (SpaceNavigator and friends)" ON)
 if(UNIX AND NOT APPLE)
   option(WITH_INSTALL_PORTABLE "Install redistributable runtime, otherwise install into CMAKE_INSTALL_PREFIX" ON)
+  option(WITH_STATIC_LIBS "Try to link with static libraries, as much as possible, to make blender more portable across distributions" OFF)
+  if(WITH_STATIC_LIBS)
+    option(WITH_BOOST_ICU "Boost uses ICU library (required for linking with static Boost built with libicu)." OFF)
+    mark_as_advanced(WITH_BOOST_ICU)
+  endif()
 endif()
 
 option(WITH_PYTHON_INSTALL       "Copy system python into the blender install folder" ON)
@@ -524,7 +491,7 @@ endif()
 if(NOT APPLE)
   option(WITH_CYCLES_DEVICE_HIP        "Enable Cycles AMD HIP support" ON)
   option(WITH_CYCLES_HIP_BINARIES      "Build Cycles AMD HIP binaries" OFF)
-  set(CYCLES_HIP_BINARIES_ARCH gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1034 gfx1035 gfx1100 gfx1101 gfx1102 CACHE STRING "AMD HIP architectures to build binaries for")
+  set(CYCLES_HIP_BINARIES_ARCH gfx900 gfx906 gfx90c gfx902 gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1034 gfx1035 gfx1100 gfx1101 gfx1102 CACHE STRING "AMD HIP architectures to build binaries for")
   mark_as_advanced(WITH_CYCLES_DEVICE_HIP)
   mark_as_advanced(CYCLES_HIP_BINARIES_ARCH)
 endif()
@@ -1026,8 +993,6 @@ set(PLATFORM_LINKLIBS "")
 # - CMAKE_EXE_LINKER_FLAGS_DEBUG
 set(PLATFORM_LINKFLAGS "")
 set(PLATFORM_LINKFLAGS_DEBUG "")
-set(PLATFORM_LINKFLAGS_RELEASE "")
-set(PLATFORM_LINKFLAGS_EXECUTABLE "")
 
 if(NOT CMAKE_BUILD_TYPE MATCHES "Release")
   if(WITH_COMPILER_ASAN)
@@ -1241,6 +1206,13 @@ if(WITH_OPENGL)
   add_definitions(-DWITH_OPENGL)
 endif()
 
+#-----------------------------------------------------------------------------
+# Configure Vulkan.
+
+if(WITH_VULKAN_BACKEND)
+  list(APPEND BLENDER_GL_LIBRARIES ${VULKAN_LIBRARIES})
+endif()
+
 # -----------------------------------------------------------------------------
 # Configure Metal
 
@@ -1290,14 +1262,12 @@ endif()
 # -----------------------------------------------------------------------------
 # Configure Bullet
 
-if(WITH_BULLET)
-  if(WITH_SYSTEM_BULLET)
-    find_package(Bullet)
-    set_and_warn_library_found("Bullet" BULLET_FOUND WITH_BULLET)
-  else()
-    set(BULLET_INCLUDE_DIRS "${CMAKE_SOURCE_DIR}/extern/bullet2/src")
-    set(BULLET_LIBRARIES "extern_bullet")
-  endif()
+if(WITH_BULLET AND WITH_SYSTEM_BULLET)
+  find_package(Bullet)
+  set_and_warn_library_found("Bullet" BULLET_FOUND WITH_BULLET)
+else()
+  set(BULLET_INCLUDE_DIRS "${CMAKE_SOURCE_DIR}/extern/bullet2/src")
+  # set(BULLET_LIBRARIES "")
 endif()
 
 

GNUmakefile

@@ -71,13 +71,6 @@ Static Source Code Checking
    * check_mypy:            Checks all Python scripts using mypy,
                             see: source/tools/check_source/check_mypy_config.py scripts which are included.
 
-Documentation Checking
-
-   * check_wiki_file_structure:
-     Check the WIKI documentation for the source-tree's file structure
-     matches Blender's source-code.
-     See: https://wiki.blender.org/wiki/Source/File_Structure
-
 Spell Checkers
    This runs the spell checker from the developer tools repositor.
 
@@ -488,10 +481,6 @@ check_smatch: .FORCE
 check_mypy: .FORCE
 	@$(PYTHON) "$(BLENDER_DIR)/source/tools/check_source/check_mypy.py"
 
-check_wiki_file_structure: .FORCE
-	@PYTHONIOENCODING=utf_8 $(PYTHON) \
-	    "$(BLENDER_DIR)/source/tools/check_wiki/check_wiki_file_structure.py"
-
 check_spelling_py: .FORCE
 	@cd "$(BUILD_DIR)" ; \
 	PYTHONIOENCODING=utf_8 $(PYTHON) \

build_files/build_environment/cmake/dpcpp.cmake

@@ -2,7 +2,7 @@
 
 # LLVM does not switch over to cpp17 until llvm 16 and building ealier versions with
 # MSVC is leading to some crashes in ISPC. Switch back to their default on all platforms
-# for now.
+# for now. 
 string(REPLACE "-DCMAKE_CXX_STANDARD=17" " " DPCPP_CMAKE_FLAGS "${DEFAULT_CMAKE_FLAGS}")
 
 if(WIN32)

build_files/build_environment/cmake/llvm.cmake

@@ -42,7 +42,7 @@ endif()
 
 # LLVM does not switch over to cpp17 until llvm 16 and building ealier versions with
 # MSVC is leading to some crashes in ISPC. Switch back to their default on all platforms
-# for now.
+# for now. 
 string(REPLACE "-DCMAKE_CXX_STANDARD=17" " " LLVM_CMAKE_FLAGS "${DEFAULT_CMAKE_FLAGS}")
 
 # short project name due to long filename issues on windows

build_files/build_environment/cmake/ssl.cmake

@@ -10,9 +10,9 @@ if(WIN32)
     DOWNLOAD_DIR ${DOWNLOAD_DIR}
     URL_HASH ${SSL_HASH_TYPE}=${SSL_HASH}
     PREFIX ${BUILD_DIR}/ssl
-    CONFIGURE_COMMAND echo "."
-    BUILD_COMMAND echo "."
-    INSTALL_COMMAND echo "."
+    CONFIGURE_COMMAND echo "." 
+    BUILD_COMMAND echo "." 
+    INSTALL_COMMAND echo "." 
     INSTALL_DIR ${LIBDIR}/ssl
   )
 else()
@@ -46,4 +46,4 @@ else()
     INSTALL_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/ssl/src/external_ssl/ && make install
     INSTALL_DIR ${LIBDIR}/ssl
   )
-endif()
+endif()

build_files/build_environment/cmake/usd.cmake

@@ -29,7 +29,7 @@ elseif(UNIX)
   set(USD_PLATFORM_FLAGS
     -DPYTHON_INCLUDE_DIR=${LIBDIR}/python/include/python${PYTHON_SHORT_VERSION}/
     -DPYTHON_LIBRARY=${LIBDIR}/tbb/lib/${LIBPREFIX}${TBB_LIBRARY}${SHAREDLIBEXT}
-  )
+   )
 
   if(APPLE)
     set(USD_SHARED_LINKER_FLAGS "-Xlinker -undefined -Xlinker dynamic_lookup")

build_files/build_environment/cmake/xml2.cmake

@@ -1,7 +1,7 @@
 # SPDX-License-Identifier: GPL-2.0-or-later
 
 if(WIN32)
-  set(XML2_EXTRA_ARGS
+  set(XML2_EXTRA_ARGS 
     -DLIBXML2_WITH_ZLIB=OFF
     -DLIBXML2_WITH_LZMA=OFF
     -DLIBXML2_WITH_PYTHON=OFF

build_files/cmake/Modules/FindMoltenVK.cmake

@@ -19,13 +19,9 @@ ENDIF()
 
 SET(_moltenvk_SEARCH_DIRS
   ${MOLTENVK_ROOT_DIR}
+  ${LIBDIR}/vulkan/MoltenVK
 )
 
-# FIXME: These finder modules typically don't use LIBDIR,
-# this should be set by `./build_files/cmake/platform/` instead.
-IF(DEFINED LIBDIR)
-  SET(_moltenvk_SEARCH_DIRS ${_moltenvk_SEARCH_DIRS} ${LIBDIR}/moltenvk)
-ENDIF()
 
 FIND_PATH(MOLTENVK_INCLUDE_DIR
   NAMES

build_files/cmake/Modules/FindOptiX.cmake

@@ -17,13 +17,9 @@ ENDIF()
 
 SET(_optix_SEARCH_DIRS
   ${OPTIX_ROOT_DIR}
+  "$ENV{PROGRAMDATA}/NVIDIA Corporation/OptiX SDK 7.3.0"
 )
 
-# TODO: Which environment uses this?
-if(DEFINED ENV{PROGRAMDATA})
-  list(APPEND _optix_SEARCH_DIRS "$ENV{PROGRAMDATA}/NVIDIA Corporation/OptiX SDK 7.3.0")
-endif()
-
 FIND_PATH(OPTIX_INCLUDE_DIR
   NAMES
     optix.h

build_files/cmake/Modules/FindPythonLibsUnix.cmake

@@ -67,8 +67,6 @@ ENDIF()
 
 STRING(REPLACE "." "" PYTHON_VERSION_NO_DOTS ${PYTHON_VERSION})
 
-SET(_PYTHON_ABI_FLAGS "")
-
 SET(_python_SEARCH_DIRS
   ${PYTHON_ROOT_DIR}
   "$ENV{HOME}/py${PYTHON_VERSION_NO_DOTS}"

build_files/cmake/Modules/FindShaderC.cmake

@@ -1,63 +0,0 @@
-# SPDX-License-Identifier: BSD-3-Clause
-# Copyright 2023 Blender Foundation.
-
-# - Find ShaderC libraries
-# Find the ShaderC includes and libraries
-# This module defines
-#  SHADERC_INCLUDE_DIRS, where to find MoltenVK headers, Set when
-#                        SHADERC_INCLUDE_DIR is found.
-#  SHADERC_LIBRARIES, libraries to link against to use ShaderC.
-#  SHADERC_ROOT_DIR, The base directory to search for ShaderC.
-#                    This can also be an environment variable.
-#  SHADERC_FOUND, If false, do not try to use ShaderC.
-#
-
-# If SHADERC_ROOT_DIR was defined in the environment, use it.
-IF(NOT SHADERC_ROOT_DIR AND NOT $ENV{SHADERC_ROOT_DIR} STREQUAL "")
-  SET(SHADERC_ROOT_DIR $ENV{SHADERC_ROOT_DIR})
-ENDIF()
-
-SET(_shaderc_SEARCH_DIRS
-  ${SHADERC_ROOT_DIR}
-)
-
-# FIXME: These finder modules typically don't use LIBDIR,
-# this should be set by `./build_files/cmake/platform/` instead.
-IF(DEFINED LIBDIR)
-  SET(_shaderc_SEARCH_DIRS ${_shaderc_SEARCH_DIRS} ${LIBDIR}/shaderc)
-ENDIF()
-
-FIND_PATH(SHADERC_INCLUDE_DIR
-  NAMES
-    shaderc/shaderc.h
-  HINTS
-    ${_shaderc_SEARCH_DIRS}
-  PATH_SUFFIXES
-    include
-)
-
-FIND_LIBRARY(SHADERC_LIBRARY
-  NAMES
-    shaderc_combined
-  HINTS
-    ${_shaderc_SEARCH_DIRS}
-  PATH_SUFFIXES
-    lib
-)
-
-# handle the QUIETLY and REQUIRED arguments and set SHADERC_FOUND to TRUE if
-# all listed variables are TRUE
-INCLUDE(FindPackageHandleStandardArgs)
-FIND_PACKAGE_HANDLE_STANDARD_ARGS(ShaderC DEFAULT_MSG SHADERC_LIBRARY SHADERC_INCLUDE_DIR)
-
-IF(SHADERC_FOUND)
-  SET(SHADERC_LIBRARIES ${SHADERC_LIBRARY})
-  SET(SHADERC_INCLUDE_DIRS ${SHADERC_INCLUDE_DIR})
-ENDIF()
-
-MARK_AS_ADVANCED(
-  SHADERC_INCLUDE_DIR
-  SHADERC_LIBRARY
-)
-
-UNSET(_shaderc_SEARCH_DIRS)

build_files/cmake/Modules/FindVulkan.cmake

@@ -1,63 +0,0 @@
-# SPDX-License-Identifier: BSD-3-Clause
-# Copyright 2023 Blender Foundation.
-
-# - Find Vulkan libraries
-# Find the Vulkan includes and libraries
-# This module defines
-#  VULKAN_INCLUDE_DIRS, where to find Vulkan headers, Set when
-#                       VULKAN_INCLUDE_DIR is found.
-#  VULKAN_LIBRARIES, libraries to link against to use Vulkan.
-#  VULKAN_ROOT_DIR, The base directory to search for Vulkan.
-#                    This can also be an environment variable.
-#  VULKAN_FOUND, If false, do not try to use Vulkan.
-#
-
-# If VULKAN_ROOT_DIR was defined in the environment, use it.
-IF(NOT VULKAN_ROOT_DIR AND NOT $ENV{VULKAN_ROOT_DIR} STREQUAL "")
-  SET(VULKAN_ROOT_DIR $ENV{VULKAN_ROOT_DIR})
-ENDIF()
-
-SET(_vulkan_SEARCH_DIRS
-  ${VULKAN_ROOT_DIR}
-)
-
-# FIXME: These finder modules typically don't use LIBDIR,
-# this should be set by `./build_files/cmake/platform/` instead.
-IF(DEFINED LIBDIR)
-  SET(_vulkan_SEARCH_DIRS ${_vulkan_SEARCH_DIRS} ${LIBDIR}/vulkan)
-ENDIF()
-
-FIND_PATH(VULKAN_INCLUDE_DIR
-  NAMES
-    vulkan/vulkan.h
-  HINTS
-    ${_vulkan_SEARCH_DIRS}
-  PATH_SUFFIXES
-    include
-)
-
-FIND_LIBRARY(VULKAN_LIBRARY
-  NAMES
-    vulkan
-  HINTS
-    ${_vulkan_SEARCH_DIRS}
-  PATH_SUFFIXES
-    lib
-)
-
-# handle the QUIETLY and REQUIRED arguments and set VULKAN_FOUND to TRUE if
-# all listed variables are TRUE
-INCLUDE(FindPackageHandleStandardArgs)
-FIND_PACKAGE_HANDLE_STANDARD_ARGS(Vulkan DEFAULT_MSG VULKAN_LIBRARY VULKAN_INCLUDE_DIR)
-
-IF(VULKAN_FOUND)
-  SET(VULKAN_LIBRARIES ${VULKAN_LIBRARY})
-  SET(VULKAN_INCLUDE_DIRS ${VULKAN_INCLUDE_DIR})
-ENDIF()
-
-MARK_AS_ADVANCED(
-  VULKAN_INCLUDE_DIR
-  VULKAN_LIBRARY
-)
-
-UNSET(_vulkan_SEARCH_DIRS)

build_files/cmake/cmake_print_build_options.py

@@ -6,80 +6,18 @@
 import re
 import sys
 
-from typing import Optional
-
 cmakelists_file = sys.argv[-1]
 
 
-def count_backslashes_before_pos(file_data: str, pos: int) -> int:
-    slash_count = 0
-    pos -= 1
-    while pos >= 0:
-        if file_data[pos] != '\\':
-            break
-        pos -= 1
-        slash_count += 1
-    return slash_count
-
-
-def extract_cmake_string_at_pos(file_data: str, pos_beg: int) -> Optional[str]:
-    assert file_data[pos_beg - 1] == '"'
-
-    pos = pos_beg
-    # Dummy assignment.
-    pos_end = pos_beg
-    while True:
-        pos_next = file_data.find('"', pos)
-        if pos_next == -1:
-            raise Exception("Un-terminated string (parse error?)")
-
-        count_slashes = count_backslashes_before_pos(file_data, pos_next)
-        if (count_slashes % 2) == 0:
-            pos_end = pos_next
-            # Found the closing quote.
-            break
-
-        # The quote was back-slash escaped, step over it.
-        pos = pos_next + 1
-        file_data[pos_next]
-
-    assert file_data[pos_end] == '"'
-
-    if pos_beg == pos_end:
-        return None
-
-    # See: https://cmake.org/cmake/help/latest/manual/cmake-language.7.html#escape-sequences
-    text = file_data[pos_beg: pos_end].replace(
-        # Handle back-slash literals.
-        "\\\\", "\\",
-    ).replace(
-        # Handle tabs.
-        "\\t", "\t",
-    ).replace(
-        # Handle escaped quotes.
-        "\\\"", "\"",
-    ).replace(
-        # Handle tabs.
-        "\\;", ";",
-    ).replace(
-        # Handle trailing newlines.
-        "\\\n", "",
-    )
-
-    return text
-
-
-def main() -> None:
+def main():
     options = []
-    with open(cmakelists_file, 'r', encoding="utf-8") as fh:
-        file_data = fh.read()
-        for m in re.finditer(r"^\s*option\s*\(\s*(WITH_[a-zA-Z0-9_]+)\s+(\")", file_data, re.MULTILINE):
-            option_name = m.group(1)
-            option_descr = extract_cmake_string_at_pos(file_data, m.span(2)[1])
-            if option_descr is None:
-                # Possibly a parsing error, at least show something.
-                option_descr = "(UNDOCUMENTED)"
-            options.append("{:s}: {:s}".format(option_name, option_descr))
+    for l in open(cmakelists_file, 'r').readlines():
+        if not l.lstrip().startswith('#'):
+            l_option = re.sub(r'.*\boption\s*\(\s*(WITH_[a-zA-Z0-9_]+)\s+\"(.*)\"\s*.*', r'\g<1> - \g<2>', l)
+            if l_option != l:
+                l_option = l_option.strip()
+                if l_option.startswith('WITH_'):
+                    options.append(l_option)
 
     print('\n'.join(options))
 

build_files/cmake/macros.cmake

@@ -550,9 +550,7 @@ function(setup_platform_linker_libs
   endif()
 
   if(WIN32 AND NOT UNIX)
-    if(DEFINED PTHREADS_LIBRARIES)
-      target_link_libraries(${target} ${PTHREADS_LIBRARIES})
-    endif()
+    target_link_libraries(${target} ${PTHREADS_LIBRARIES})
   endif()
 
   # target_link_libraries(${target} ${PLATFORM_LINKLIBS} ${CMAKE_DL_LIBS})
@@ -1117,7 +1115,7 @@ function(find_python_package
     # endif()
     # Not set, so initialize.
   else()
-    string(REPLACE "." ";" _PY_VER_SPLIT "${PYTHON_VERSION}")
+   string(REPLACE "." ";" _PY_VER_SPLIT "${PYTHON_VERSION}")
     list(GET _PY_VER_SPLIT 0 _PY_VER_MAJOR)
 
     # re-cache
@@ -1264,7 +1262,7 @@ endmacro()
 
 # Utility to gather and install precompiled shared libraries.
 macro(add_bundled_libraries library_dir)
-  if(DEFINED LIBDIR)
+  if(EXISTS ${LIBDIR})
     set(_library_dir ${LIBDIR}/${library_dir})
     if(WIN32)
       file(GLOB _all_library_versions ${_library_dir}/*\.dll)
@@ -1277,7 +1275,7 @@ macro(add_bundled_libraries library_dir)
     list(APPEND PLATFORM_BUNDLED_LIBRARY_DIRS ${_library_dir})
     unset(_all_library_versions)
     unset(_library_dir)
-  endif()
+ endif()
 endmacro()
 
 macro(windows_install_shared_manifest)

build_files/cmake/platform/platform_apple.cmake

@@ -97,8 +97,20 @@ add_bundled_libraries(materialx/lib)
 
 if(WITH_VULKAN_BACKEND)
   find_package(MoltenVK REQUIRED)
-  find_package(ShaderC REQUIRED)
-  find_package(Vulkan REQUIRED)
+
+  if(EXISTS ${LIBDIR}/vulkan)
+    set(VULKAN_FOUND On)
+    set(VULKAN_ROOT_DIR ${LIBDIR}/vulkan/macOS)
+    set(VULKAN_INCLUDE_DIR ${VULKAN_ROOT_DIR}/include)
+    set(VULKAN_LIBRARY ${VULKAN_ROOT_DIR}/lib/libvulkan.1.dylib)
+    set(SHADERC_LIBRARY ${VULKAN_ROOT_DIR}/lib/libshaderc_combined.a)
+
+    set(VULKAN_INCLUDE_DIRS ${VULKAN_INCLUDE_DIR} ${MOLTENVK_INCLUDE_DIRS})
+    set(VULKAN_LIBRARIES ${VULKAN_LIBRARY} ${SHADERC_LIBRARY} ${MOLTENVK_LIBRARIES})
+  else()
+    message(WARNING "Vulkan SDK was not found, disabling WITH_VULKAN_BACKEND")
+    set(WITH_VULKAN_BACKEND OFF)
+  endif()
 endif()
 
 if(WITH_OPENSUBDIV)

build_files/cmake/platform/platform_old_libs_update.cmake

@@ -1,12 +1,7 @@
 # SPDX-License-Identifier: GPL-2.0-or-later
 # Copyright 2022 Blender Foundation. All rights reserved.
 
-# Auto update existing CMake caches for new libraries.
-
-# Assert that `LIBDIR` is defined.
-if(NOT (DEFINED LIBDIR))
-  message(FATAL_ERROR "Logical error, expected 'LIBDIR' to be defined!")
-endif()
+# Auto update existing CMake caches for new libraries
 
 # Clear cached variables whose name matches `pattern`.
 function(unset_cache_variables pattern)

build_files/cmake/platform/platform_unix.cmake

@@ -4,52 +4,38 @@
 # Libraries configuration for any *nix system including Linux and Unix (excluding APPLE).
 
 # Detect precompiled library directory
+if(NOT DEFINED LIBDIR)
+  # Path to a locally compiled libraries.
+  set(LIBDIR_NAME ${CMAKE_SYSTEM_NAME}_${CMAKE_SYSTEM_PROCESSOR})
+  string(TOLOWER ${LIBDIR_NAME} LIBDIR_NAME)
+  set(LIBDIR_NATIVE_ABI ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_NAME})
 
-if(NOT WITH_LIBS_PRECOMPILED)
-  unset(LIBDIR)
-else()
-  if(NOT DEFINED LIBDIR)
-    # Path to a locally compiled libraries.
-    set(LIBDIR_NAME ${CMAKE_SYSTEM_NAME}_${CMAKE_SYSTEM_PROCESSOR})
-    string(TOLOWER ${LIBDIR_NAME} LIBDIR_NAME)
-    set(LIBDIR_NATIVE_ABI ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_NAME})
+  # Path to precompiled libraries with known glibc 2.28 ABI.
+  set(LIBDIR_GLIBC228_ABI ${CMAKE_SOURCE_DIR}/../lib/linux_x86_64_glibc_228)
 
-    # Path to precompiled libraries with known glibc 2.28 ABI.
-    set(LIBDIR_GLIBC228_ABI ${CMAKE_SOURCE_DIR}/../lib/linux_x86_64_glibc_228)
-
-    # Choose the best suitable libraries.
-    if(EXISTS ${LIBDIR_NATIVE_ABI})
-      set(LIBDIR ${LIBDIR_NATIVE_ABI})
+  # Choose the best suitable libraries.
+  if(EXISTS ${LIBDIR_NATIVE_ABI})
+    set(LIBDIR ${LIBDIR_NATIVE_ABI})
+    set(WITH_LIBC_MALLOC_HOOK_WORKAROUND True)
+  elseif(EXISTS ${LIBDIR_GLIBC228_ABI})
+    set(LIBDIR ${LIBDIR_GLIBC228_ABI})
+    if(WITH_MEM_JEMALLOC)
+      # jemalloc provides malloc hooks.
+      set(WITH_LIBC_MALLOC_HOOK_WORKAROUND False)
+    else()
       set(WITH_LIBC_MALLOC_HOOK_WORKAROUND True)
-    elseif(EXISTS ${LIBDIR_GLIBC228_ABI})
-      set(LIBDIR ${LIBDIR_GLIBC228_ABI})
-      if(WITH_MEM_JEMALLOC)
-        # jemalloc provides malloc hooks.
-        set(WITH_LIBC_MALLOC_HOOK_WORKAROUND False)
-      else()
-        set(WITH_LIBC_MALLOC_HOOK_WORKAROUND True)
-      endif()
     endif()
-
-    # Avoid namespace pollustion.
-    unset(LIBDIR_NATIVE_ABI)
-    unset(LIBDIR_GLIBC228_ABI)
   endif()
 
-  if(NOT (EXISTS ${LIBDIR}))
-    message(STATUS
-      "Unable to find LIBDIR: ${LIBDIR}, system libraries may be used "
-      "(disable WITH_LIBS_PRECOMPILED to suppress this message)."
-    )
-    unset(LIBDIR)
-  endif()
+  # Avoid namespace pollustion.
+  unset(LIBDIR_NATIVE_ABI)
+  unset(LIBDIR_GLIBC228_ABI)
 endif()
 
-
 # Support restoring this value once pre-compiled libraries have been handled.
 set(WITH_STATIC_LIBS_INIT ${WITH_STATIC_LIBS})
 
-if(DEFINED LIBDIR)
+if(EXISTS ${LIBDIR})
   message(STATUS "Using pre-compiled LIBDIR: ${LIBDIR}")
 
   file(GLOB LIB_SUBDIRS ${LIBDIR}/*)
@@ -99,7 +85,7 @@ endmacro()
 # These are libraries that may be precompiled. For this we disable searching in
 # the system directories so that we don't accidentally use them instead.
 
-if(DEFINED LIBDIR)
+if(EXISTS ${LIBDIR})
   without_system_libs_begin()
 endif()
 
@@ -111,7 +97,6 @@ find_package_wrapper(Epoxy REQUIRED)
 
 if(WITH_VULKAN_BACKEND)
   find_package_wrapper(Vulkan REQUIRED)
-  find_package_wrapper(ShaderC REQUIRED)
 endif()
 
 function(check_freetype_for_brotli)
@@ -129,7 +114,7 @@ endfunction()
 if(NOT WITH_SYSTEM_FREETYPE)
   # FreeType compiled with Brotli compression for woff2.
   find_package_wrapper(Freetype REQUIRED)
-  if(DEFINED LIBDIR)
+  if(EXISTS ${LIBDIR})
     find_package_wrapper(Brotli REQUIRED)
 
     # NOTE: This is done on WIN32 & APPLE but fails on some Linux systems.
@@ -156,7 +141,7 @@ if(WITH_PYTHON)
   if(WITH_PYTHON_MODULE AND NOT WITH_INSTALL_PORTABLE)
     # Installing into `site-packages`, warn when installing into `./../lib/`
     # which script authors almost certainly don't want.
-    if(DEFINED LIBDIR)
+    if(EXISTS ${LIBDIR})
       path_is_prefix(LIBDIR PYTHON_SITE_PACKAGES _is_prefix)
       if(_is_prefix)
         message(WARNING "
@@ -232,7 +217,7 @@ if(WITH_CODEC_SNDFILE)
 endif()
 
 if(WITH_CODEC_FFMPEG)
-  if(DEFINED LIBDIR)
+  if(EXISTS ${LIBDIR})
     set(FFMPEG_ROOT_DIR ${LIBDIR}/ffmpeg)
     # Override FFMPEG components to also include static library dependencies
     # included with precompiled libraries, and to ensure correct link order.
@@ -247,7 +232,7 @@ if(WITH_CODEC_FFMPEG)
       vpx
       x264
       xvidcore)
-    if((DEFINED LIBDIR) AND (EXISTS ${LIBDIR}/ffmpeg/lib/libaom.a))
+    if(EXISTS ${LIBDIR}/ffmpeg/lib/libaom.a)
       list(APPEND FFMPEG_FIND_COMPONENTS aom)
     endif()
   elseif(FFMPEG)
@@ -445,13 +430,10 @@ if(WITH_OPENIMAGEIO)
     ${PNG_LIBRARIES}
     ${JPEG_LIBRARIES}
     ${ZLIB_LIBRARIES}
+    ${BOOST_LIBRARIES}
   )
-
   set(OPENIMAGEIO_DEFINITIONS "")
 
-  if(WITH_BOOST)
-    list(APPEND OPENIMAGEIO_LIBRARIES "${BOOST_LIBRARIES}")
-  endif()
   if(WITH_IMAGE_TIFF)
     list(APPEND OPENIMAGEIO_LIBRARIES "${TIFF_LIBRARY}")
   endif()
@@ -469,7 +451,7 @@ add_bundled_libraries(openimageio/lib)
 if(WITH_OPENCOLORIO)
   find_package_wrapper(OpenColorIO 2.0.0)
 
-  set(OPENCOLORIO_DEFINITIONS "")
+  set(OPENCOLORIO_DEFINITIONS)
   set_and_warn_library_found("OpenColorIO" OPENCOLORIO_FOUND WITH_OPENCOLORIO)
 endif()
 add_bundled_libraries(opencolorio/lib)
@@ -484,7 +466,7 @@ if(WITH_OPENIMAGEDENOISE)
 endif()
 
 if(WITH_LLVM)
-  if(DEFINED LIBDIR)
+  if(EXISTS ${LIBDIR})
     set(LLVM_STATIC ON)
   endif()
 
@@ -498,7 +480,7 @@ if(WITH_LLVM)
     endif()
 
     # Symbol conflicts with same UTF library used by OpenCollada
-    if(DEFINED LIBDIR)
+    if(EXISTS ${LIBDIR})
       if(WITH_OPENCOLLADA AND (${LLVM_VERSION} VERSION_LESS "4.0.0"))
         list(REMOVE_ITEM OPENCOLLADA_LIBRARIES ${OPENCOLLADA_UTF_LIBRARY})
       endif()
@@ -554,7 +536,7 @@ if(WITH_CYCLES AND WITH_CYCLES_PATH_GUIDING)
   endif()
 endif()
 
-if(DEFINED LIBDIR)
+if(EXISTS ${LIBDIR})
   without_system_libs_end()
 endif()
 
@@ -569,14 +551,9 @@ else()
 endif()
 
 find_package(Threads REQUIRED)
-# `FindThreads` documentation notes that this may be empty
-# with the system libraries provide threading functionality.
-if(CMAKE_THREAD_LIBS_INIT)
-  list(APPEND PLATFORM_LINKLIBS ${CMAKE_THREAD_LIBS_INIT})
-  # used by other platforms
-  set(PTHREADS_LIBRARIES ${CMAKE_THREAD_LIBS_INIT})
-endif()
-
+list(APPEND PLATFORM_LINKLIBS ${CMAKE_THREAD_LIBS_INIT})
+# used by other platforms
+set(PTHREADS_LIBRARIES ${CMAKE_THREAD_LIBS_INIT})
 
 if(CMAKE_DL_LIBS)
   list(APPEND PLATFORM_LINKLIBS ${CMAKE_DL_LIBS})
@@ -598,7 +575,7 @@ add_definitions(-D_LARGEFILE_SOURCE -D_FILE_OFFSET_BITS=64 -D_LARGEFILE64_SOURCE
 #
 # Keep last, so indirectly linked libraries don't override our own pre-compiled libs.
 
-if(DEFINED LIBDIR)
+if(EXISTS ${LIBDIR})
   # Clear the prefix path as it causes the `LIBDIR` to override system locations.
   unset(CMAKE_PREFIX_PATH)
 
@@ -654,7 +631,7 @@ if(WITH_GHOST_WAYLAND)
   # When dynamically linked WAYLAND is used and `${LIBDIR}/wayland` is present,
   # there is no need to search for the libraries as they are not needed for building.
   # Only the headers are needed which can reference the known paths.
-  if((DEFINED LIBDIR) AND (EXISTS "${LIBDIR}/wayland" AND WITH_GHOST_WAYLAND_DYNLOAD))
+  if(EXISTS "${LIBDIR}/wayland" AND WITH_GHOST_WAYLAND_DYNLOAD)
     set(_use_system_wayland OFF)
   else()
     set(_use_system_wayland ON)
@@ -718,7 +695,7 @@ if(WITH_GHOST_WAYLAND)
       add_definitions(-DWITH_GHOST_WAYLAND_LIBDECOR)
     endif()
 
-    if((DEFINED LIBDIR) AND (EXISTS "${LIBDIR}/wayland/bin/wayland-scanner"))
+    if(EXISTS "${LIBDIR}/wayland/bin/wayland-scanner")
       set(WAYLAND_SCANNER "${LIBDIR}/wayland/bin/wayland-scanner")
     else()
       pkg_get_variable(WAYLAND_SCANNER wayland-scanner wayland_scanner)

build_files/config/pipeline_config.yaml

@@ -43,10 +43,6 @@ update-code:
             branch: trunk
             commit_id: HEAD
             path: lib/benchmarks
-        assets:
-            branch: trunk
-            commit_id: HEAD
-            path: lib/assets
 
 #
 # Buildbot only configs
@@ -63,7 +59,7 @@ buildbot:
     optix:
         version: '7.3.0'
     ocloc:
-        version: '101.4032'
+        version: '101.3430'
     cmake:
         default:
             version: any

build_files/utils/make_bpy_wheel.py

@@ -24,7 +24,7 @@ import os
 import re
 import platform
 import string
-import setuptools
+import setuptools  # type: ignore
 import sys
 
 from typing import (
@@ -208,7 +208,7 @@ def main() -> None:
         return paths
 
     # Ensure this wheel is marked platform specific.
-    class BinaryDistribution(setuptools.dist.Distribution):
+    class BinaryDistribution(setuptools.dist.Distribution):  # type: ignore
         def has_ext_modules(self) -> bool:
             return True
 

build_files/utils/make_test.py

@@ -13,10 +13,10 @@ import sys
 import make_utils
 from make_utils import call
 
-# Parse arguments.
+# Parse arguments
 
 
-def parse_arguments() -> argparse.Namespace:
+def parse_arguments():
     parser = argparse.ArgumentParser()
     parser.add_argument("--ctest-command", default="ctest")
     parser.add_argument("--cmake-command", default="cmake")

build_files/utils/make_update.py

@@ -104,30 +104,17 @@ def svn_update(args: argparse.Namespace, release_version: Optional[str]) -> None
             svn_url_tests = svn_url + lib_tests
             call(svn_non_interactive + ["checkout", svn_url_tests, lib_tests_dirpath])
 
-    lib_assets = "assets"
-    lib_assets_dirpath = os.path.join(lib_dirpath, lib_assets)
-
-    if not os.path.exists(lib_assets_dirpath):
-        print_stage("Checking out Assets")
-
-        if make_utils.command_missing(args.svn_command):
-            sys.stderr.write("svn not found, can't checkout assets\n")
-            sys.exit(1)
-
-        svn_url_assets = svn_url + lib_assets
-        call(svn_non_interactive + ["checkout", svn_url_assets, lib_assets_dirpath])
-
-    # Update precompiled libraries, assets and tests
+    # Update precompiled libraries and tests
 
     if not os.path.isdir(lib_dirpath):
         print("Library path: %r, not found, skipping" % lib_dirpath)
     else:
         paths_local_and_remote = []
         if os.path.exists(os.path.join(lib_dirpath, ".svn")):
-            print_stage("Updating Precompiled Libraries, Assets and Tests (one repository)")
+            print_stage("Updating Precompiled Libraries and Tests (one repository)")
             paths_local_and_remote.append((lib_dirpath, svn_url))
         else:
-            print_stage("Updating Precompiled Libraries, Assets and Tests (multiple repositories)")
+            print_stage("Updating Precompiled Libraries and Tests (multiple repositories)")
             # Separate paths checked out.
             for dirname in os.listdir(lib_dirpath):
                 if dirname.startswith("."):

doc/python_api/sphinx_doc_gen.py

@@ -2098,8 +2098,6 @@ def write_rst_types_index(basepath):
         fw(title_string("Types (bpy.types)", "="))
         fw(".. module:: bpy.types\n\n")
         fw(".. toctree::\n")
-        # Only show top-level entries (avoids unreasonably large pages).
-        fw("   :maxdepth: 1\n")
         fw("   :glob:\n\n")
         fw("   bpy.types.*\n\n")
 
@@ -2126,8 +2124,6 @@ def write_rst_ops_index(basepath):
         write_example_ref("", fw, "bpy.ops")
         fw(".. toctree::\n")
         fw("   :caption: Submodules\n")
-        # Only show top-level entries (avoids unreasonably large pages).
-        fw("   :maxdepth: 1\n")
         fw("   :glob:\n\n")
         fw("   bpy.ops.*\n\n")
         file.close()

extern/mantaflow/CMakeLists.txt

@@ -13,12 +13,10 @@ endif()
 
 # Exporting functions from the blender binary gives linker warnings on Apple arm64 systems.
 # Silence them here.
-if(APPLE)
-  if("${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64")
-    if(CMAKE_COMPILER_IS_GNUCXX OR "${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang")
-      string(APPEND CMAKE_C_FLAGS " -fvisibility=hidden")
-      string(APPEND CMAKE_CXX_FLAGS " -fvisibility=hidden")
-    endif()
+if(APPLE AND ("${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64"))
+  if(CMAKE_COMPILER_IS_GNUCXX OR "${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang")
+    string(APPEND CMAKE_C_FLAGS " -fvisibility=hidden")
+    string(APPEND CMAKE_CXX_FLAGS " -fvisibility=hidden")
   endif()
 endif()
 
@@ -263,11 +261,9 @@ set(LIB
 
 blender_add_lib(extern_mantaflow "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")
 
-if(WITH_OPENVDB)
-  # The VDB libs above are only added to as INTERFACE libs by blender_add_lib,
-  # meaning extern_mantaflow itself actually does not have a dependency on the
-  # openvdb libraries, and CMAKE is free to link the vdb libs before
-  # extern_mantaflow causing linker errors on linux. By explicitly declaring
-  # a dependency here, cmake will do the right thing.
-  target_link_libraries(extern_mantaflow PRIVATE ${OPENVDB_LIBRARIES})
-endif()
+# The VDB libs above are only added to as INTERFACE libs by blender_add_lib,
+# meaning extern_mantaflow itself actually does not have a dependency on the
+# openvdb libraries, and CMAKE is free to link the vdb libs before
+# extern_mantaflow causing linker errors on linux. By explicitly declaring
+# a dependency here, cmake will do the right thing.
+target_link_libraries(extern_mantaflow PRIVATE ${OPENVDB_LIBRARIES})

extern/vulkan_memory_allocator/CMakeLists.txt

@@ -7,7 +7,6 @@ set(INC
 
 set(INC_SYS
   ${VULKAN_INCLUDE_DIRS}
-  ${MOLTENVK_INCLUDE_DIRS}
 )
 
 set(SRC

extern/vulkan_memory_allocator/patches/remove_compilation_warning.diff

@@ -1,15 +0,0 @@
-diff --git a/extern/vulkan_memory_allocator/vk_mem_alloc.h b/extern/vulkan_memory_allocator/vk_mem_alloc.h
-index 60f572038c0..63a9994ba46 100644
---- a/extern/vulkan_memory_allocator/vk_mem_alloc.h
-+++ b/extern/vulkan_memory_allocator/vk_mem_alloc.h
-@@ -13371,8 +13371,8 @@ bool VmaDefragmentationContext_T::IncrementCounters(VkDeviceSize bytes)
-     // Early return when max found
-     if (++m_PassStats.allocationsMoved >= m_MaxPassAllocations || m_PassStats.bytesMoved >= m_MaxPassBytes)
-     {
--        VMA_ASSERT(m_PassStats.allocationsMoved == m_MaxPassAllocations ||
--            m_PassStats.bytesMoved == m_MaxPassBytes && "Exceeded maximal pass threshold!");
-+        VMA_ASSERT((m_PassStats.allocationsMoved == m_MaxPassAllocations ||
-+            m_PassStats.bytesMoved == m_MaxPassBytes) && "Exceeded maximal pass threshold!");
-         return true;
-     }
-     return false;

intern/cycles/blender/addon/properties.py

@@ -1671,19 +1671,19 @@ class CyclesPreferences(bpy.types.AddonPreferences):
             elif device_type == 'HIP':
                 import sys
                 if sys.platform[:3] == "win":
-                    col.label(text="Requires AMD GPU with RDNA architecture", icon='BLANK1')
+                    col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
                     col.label(text="and AMD Radeon Pro 21.Q4 driver or newer", icon='BLANK1')
                 elif sys.platform.startswith("linux"):
-                    col.label(text="Requires AMD GPU with RDNA architecture", icon='BLANK1')
+                    col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
                     col.label(text="and AMD driver version 22.10 or newer", icon='BLANK1')
             elif device_type == 'ONEAPI':
                 import sys
                 if sys.platform.startswith("win"):
                     col.label(text="Requires Intel GPU with Xe-HPG architecture", icon='BLANK1')
-                    col.label(text="and Windows driver version 101.4032 or newer", icon='BLANK1')
+                    col.label(text="and Windows driver version 101.3430 or newer", icon='BLANK1')
                 elif sys.platform.startswith("linux"):
                     col.label(text="Requires Intel GPU with Xe-HPG architecture and", icon='BLANK1')
-                    col.label(text="  - intel-level-zero-gpu version 1.3.24931 or newer", icon='BLANK1')
+                    col.label(text="  - intel-level-zero-gpu version 1.3.23904 or newer", icon='BLANK1')
                     col.label(text="  - oneAPI Level-Zero Loader", icon='BLANK1')
             elif device_type == 'METAL':
                 col.label(text="Requires Apple Silicon with macOS 12.2 or newer", icon='BLANK1')
@@ -1722,20 +1722,13 @@ class CyclesPreferences(bpy.types.AddonPreferences):
             row.prop(self, "peer_memory")
 
         if compute_device_type == 'METAL':
-            import platform, re
-            isNavi2 = False
-            for device in devices:
-                obj = re.search("((RX)|(Pro)|(PRO))\s+W?6\d00X",device.name)
-                if obj:
-                    isNavi2 = True
-
-            # MetalRT only works on Apple Silicon and Navi2
-            if platform.machine() == 'arm64' or isNavi2:
+            import platform
+            # MetalRT only works on Apple Silicon at present, pending argument encoding fixes on AMD
+            # Kernel specialization is only viable on Apple Silicon at present due to relative compilation speed
+            if platform.machine() == 'arm64':
                 col = layout.column()
                 col.use_property_split = True
-                # Kernel specialization is only supported on Apple Silicon
-                if platform.machine() == 'arm64':
-                    col.prop(self, "kernel_optimization_level")
+                col.prop(self, "kernel_optimization_level")
                 col.prop(self, "use_metalrt")
 
     def draw(self, context):

intern/cycles/blender/light.cpp

@@ -48,8 +48,6 @@ void BlenderSync::sync_light(BL::Object &b_parent,
     case BL::Light::type_SPOT: {
       BL::SpotLight b_spot_light(b_light);
       light->set_size(b_spot_light.shadow_soft_size());
-      light->set_axisu(transform_get_column(&tfm, 0));
-      light->set_axisv(transform_get_column(&tfm, 1));
       light->set_light_type(LIGHT_SPOT);
       light->set_spot_angle(b_spot_light.spot_size());
       light->set_spot_smooth(b_spot_light.spot_blend());

intern/cycles/cmake/macros.cmake

@@ -111,10 +111,8 @@ macro(cycles_external_libraries_append libraries)
   endif()
   if(WITH_OPENIMAGEDENOISE)
     list(APPEND ${libraries} ${OPENIMAGEDENOISE_LIBRARIES})
-    if(APPLE)
-      if("${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64")
-        list(APPEND ${libraries} "-framework Accelerate")
-      endif()
+    if(APPLE AND "${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64")
+      list(APPEND ${libraries} "-framework Accelerate")
     endif()
   endif()
   if(WITH_ALEMBIC)
@@ -138,15 +136,7 @@ macro(cycles_external_libraries_append libraries)
     ${PYTHON_LIBRARIES}
     ${ZLIB_LIBRARIES}
     ${CMAKE_DL_LIBS}
-  )
-
-  if(DEFINED PTHREADS_LIBRARIES)
-    list(APPEND ${libraries}
-      ${PTHREADS_LIBRARIES}
-    )
-  endif()
-
-  list(APPEND ${libraries}
+    ${PTHREADS_LIBRARIES}
     ${PLATFORM_LINKLIBS}
   )
 

intern/cycles/device/hip/util.h

@@ -51,7 +51,7 @@ static inline bool hipSupportsDevice(const int hipDevId)
   hipDeviceGetAttribute(&major, hipDeviceAttributeComputeCapabilityMajor, hipDevId);
   hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
 
-  return (major >= 10);
+  return (major >= 9);
 }
 
 CCL_NAMESPACE_END

intern/cycles/device/kernel.cpp

@@ -73,10 +73,6 @@ const char *device_kernel_as_string(DeviceKernel kernel)
       return "integrator_terminated_paths_array";
     case DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY:
       return "integrator_sorted_paths_array";
-    case DEVICE_KERNEL_INTEGRATOR_SORT_BUCKET_PASS:
-      return "integrator_sort_bucket_pass";
-    case DEVICE_KERNEL_INTEGRATOR_SORT_WRITE_PASS:
-      return "integrator_sort_write_pass";
     case DEVICE_KERNEL_INTEGRATOR_COMPACT_PATHS_ARRAY:
       return "integrator_compact_paths_array";
     case DEVICE_KERNEL_INTEGRATOR_COMPACT_STATES:

intern/cycles/device/metal/device.mm

@@ -55,10 +55,6 @@ void device_metal_info(vector<DeviceInfo> &devices)
     info.denoisers = DENOISER_NONE;
     info.id = id;
 
-    if (MetalInfo::get_device_vendor(device) == METAL_GPU_AMD) {
-      info.has_light_tree = false;
-    }
-
     devices.push_back(info);
     device_index++;
   }

intern/cycles/device/metal/device_impl.h

@@ -105,8 +105,6 @@ class MetalDevice : public Device {
 
   bool use_adaptive_compilation();
 
-  bool use_local_atomic_sort() const;
-
   bool make_source_and_check_if_compile_needed(MetalPipelineType pso_type);
 
   void make_source(MetalPipelineType pso_type, const uint kernel_features);

intern/cycles/device/metal/device_impl.mm

@@ -271,11 +271,6 @@ bool MetalDevice::use_adaptive_compilation()
   return DebugFlags().metal.adaptive_compile;
 }
 
-bool MetalDevice::use_local_atomic_sort() const
-{
-  return DebugFlags().metal.use_local_atomic_sort;
-}
-
 void MetalDevice::make_source(MetalPipelineType pso_type, const uint kernel_features)
 {
   string global_defines;
@@ -283,10 +278,6 @@ void MetalDevice::make_source(MetalPipelineType pso_type, const uint kernel_feat
     global_defines += "#define __KERNEL_FEATURES__ " + to_string(kernel_features) + "\n";
   }
 
-  if (use_local_atomic_sort()) {
-    global_defines += "#define __KERNEL_LOCAL_ATOMIC_SORT__\n";
-  }
-
   if (use_metalrt) {
     global_defines += "#define __METALRT__\n";
     if (motion_blur) {
@@ -336,21 +327,10 @@ void MetalDevice::make_source(MetalPipelineType pso_type, const uint kernel_feat
 #  define KERNEL_STRUCT_BEGIN(name, parent) \
     string_replace_same_length(source, "kernel_data." #parent ".", "kernel_data_" #parent "_");
 
-    bool next_member_is_specialized = true;
-
-#  define KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE next_member_is_specialized = false;
-
     /* Add constants to md5 so that 'get_best_pipeline' is able to return a suitable match. */
 #  define KERNEL_STRUCT_MEMBER(parent, _type, name) \
-    if (next_member_is_specialized) { \
-      baked_constants += string(#parent "." #name "=") + \
-                         to_string(_type(launch_params.data.parent.name)) + "\n"; \
-    } \
-    else { \
-      string_replace( \
-          source, "kernel_data_" #parent "_" #name, "kernel_data." #parent ".__unused_" #name); \
-      next_member_is_specialized = true; \
-    }
+    baked_constants += string(#parent "." #name "=") + \
+                       to_string(_type(launch_params.data.parent.name)) + "\n";
 
 #  include "kernel/data_template.h"
 

intern/cycles/device/metal/kernel.mm

@@ -49,18 +49,6 @@ struct ShaderCache {
     if (MetalInfo::get_device_vendor(mtlDevice) == METAL_GPU_APPLE) {
       switch (MetalInfo::get_apple_gpu_architecture(mtlDevice)) {
         default:
-        case APPLE_M2_BIG:
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_STATES] = {384, 128};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INIT_FROM_CAMERA] = {640, 128};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST] = {1024, 64};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW] = {704, 704};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE] = {640, 32};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY] = {896, 768};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND] = {512, 128};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW] = {32, 32};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] = {768, 576};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY] = {896, 768};
-          break;
         case APPLE_M2:
           occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_STATES] = {32, 32};
           occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INIT_FROM_CAMERA] = {832, 32};
@@ -87,9 +75,6 @@ struct ShaderCache {
           break;
       }
     }
-
-    occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORT_BUCKET_PASS] = {1024, 1024};
-    occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORT_WRITE_PASS] = {1024, 1024};
   }
   ~ShaderCache();
 
@@ -463,18 +448,13 @@ static MTLFunctionConstantValues *GetConstantValues(KernelData const *data = nul
   if (!data) {
     data = &zero_data;
   }
-  [constant_values setConstantValue:&zero_data type:MTLDataType_int atIndex:Kernel_DummyConstant];
-
-  bool next_member_is_specialized = true;
-
-#  define KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE next_member_is_specialized = false;
+  int zero_int = 0;
+  [constant_values setConstantValue:&zero_int type:MTLDataType_int atIndex:Kernel_DummyConstant];
 
 #  define KERNEL_STRUCT_MEMBER(parent, _type, name) \
-    [constant_values setConstantValue:next_member_is_specialized ? (void *)&data->parent.name : \
-                                                                   (void *)&zero_data \
+    [constant_values setConstantValue:&data->parent.name \
                                  type:MTLDataType_##_type \
-                              atIndex:KernelData_##parent##_##name]; \
-    next_member_is_specialized = true;
+                              atIndex:KernelData_##parent##_##name];
 
 #  include "kernel/data_template.h"
 

intern/cycles/device/metal/queue.h

@@ -25,7 +25,6 @@ class MetalDeviceQueue : public DeviceQueue {
   virtual int num_concurrent_states(const size_t) const override;
   virtual int num_concurrent_busy_states(const size_t) const override;
   virtual int num_sort_partition_elements() const override;
-  virtual bool supports_local_atomic_sort() const override;
 
   virtual void init_execution() override;
 

intern/cycles/device/metal/queue.mm

@@ -278,8 +278,7 @@ int MetalDeviceQueue::num_concurrent_states(const size_t state_size) const
   if (metal_device_->device_vendor == METAL_GPU_APPLE) {
     result *= 4;
 
-    /* Increasing the state count doesn't notably benefit M1-family systems.  */
-    if (MetalInfo::get_apple_gpu_architecture(metal_device_->mtlDevice) != APPLE_M1) {
+    if (MetalInfo::get_apple_gpu_architecture(metal_device_->mtlDevice) == APPLE_M2) {
       size_t system_ram = system_physical_ram();
       size_t allocated_so_far = [metal_device_->mtlDevice currentAllocatedSize];
       size_t max_recommended_working_set = [metal_device_->mtlDevice recommendedMaxWorkingSetSize];
@@ -315,11 +314,6 @@ int MetalDeviceQueue::num_sort_partition_elements() const
   return MetalInfo::optimal_sort_partition_elements(metal_device_->mtlDevice);
 }
 
-bool MetalDeviceQueue::supports_local_atomic_sort() const
-{
-  return metal_device_->use_local_atomic_sort();
-}
-
 void MetalDeviceQueue::init_execution()
 {
   /* Synchronize all textures and memory copies before executing task. */
@@ -558,24 +552,13 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
       /* See parallel_active_index.h for why this amount of shared memory is needed.
        * Rounded up to 16 bytes for Metal */
       shared_mem_bytes = (int)round_up((num_threads_per_block + 1) * sizeof(int), 16);
+      [mtlComputeCommandEncoder setThreadgroupMemoryLength:shared_mem_bytes atIndex:0];
       break;
 
-    case DEVICE_KERNEL_INTEGRATOR_SORT_BUCKET_PASS:
-    case DEVICE_KERNEL_INTEGRATOR_SORT_WRITE_PASS: {
-      int key_count = metal_device_->launch_params.data.max_shaders;
-      shared_mem_bytes = (int)round_up(key_count * sizeof(int), 16);
-      break;
-    }
-
     default:
       break;
   }
 
-  if (shared_mem_bytes) {
-    assert(shared_mem_bytes <= 32 * 1024);
-    [mtlComputeCommandEncoder setThreadgroupMemoryLength:shared_mem_bytes atIndex:0];
-  }
-
   MTLSize size_threadgroups_per_dispatch = MTLSizeMake(
       divide_up(work_size, num_threads_per_block), 1, 1);
   MTLSize size_threads_per_threadgroup = MTLSizeMake(num_threads_per_block, 1, 1);

intern/cycles/device/metal/util.h

@@ -29,7 +29,6 @@ enum AppleGPUArchitecture {
   APPLE_UNKNOWN,
   APPLE_M1,
   APPLE_M2,
-  APPLE_M2_BIG,
 };
 
 /* Contains static Metal helper functions. */

intern/cycles/device/metal/util.mm

@@ -52,7 +52,7 @@ AppleGPUArchitecture MetalInfo::get_apple_gpu_architecture(id<MTLDevice> device)
     return APPLE_M1;
   }
   else if (strstr(device_name, "M2")) {
-    return get_apple_gpu_core_count(device) <= 10 ? APPLE_M2 : APPLE_M2_BIG;
+    return APPLE_M2;
   }
   return APPLE_UNKNOWN;
 }
@@ -64,12 +64,6 @@ MetalGPUVendor MetalInfo::get_device_vendor(id<MTLDevice> device)
     return METAL_GPU_INTEL;
   }
   else if (strstr(device_name, "AMD")) {
-    /* Setting this env var hides AMD devices thus exposing any integrated Intel devices. */
-    if (auto str = getenv("CYCLES_METAL_FORCE_INTEL")) {
-      if (atoi(str)) {
-        return METAL_GPU_UNKNOWN;
-      }
-    }
     return METAL_GPU_AMD;
   }
   else if (strstr(device_name, "Apple")) {
@@ -102,15 +96,6 @@ vector<id<MTLDevice>> const &MetalInfo::get_usable_devices()
     return usable_devices;
   }
 
-  /* If the system has both an AMD GPU (discrete) and an Intel one (integrated), prefer the AMD
-   * one. This can be overriden with CYCLES_METAL_FORCE_INTEL. */
-  bool has_usable_amd_gpu = false;
-  if (@available(macos 12.3, *)) {
-    for (id<MTLDevice> device in MTLCopyAllDevices()) {
-      has_usable_amd_gpu |= (get_device_vendor(device) == METAL_GPU_AMD);
-    }
-  }
-
   metal_printf("Usable Metal devices:\n");
   for (id<MTLDevice> device in MTLCopyAllDevices()) {
     string device_name = get_device_name(device);
@@ -126,10 +111,8 @@ vector<id<MTLDevice>> const &MetalInfo::get_usable_devices()
     }
 
 #  if defined(MAC_OS_VERSION_13_0)
-    if (!has_usable_amd_gpu) {
-      if (@available(macos 13.0, *)) {
-        usable |= (vendor == METAL_GPU_INTEL);
-      }
+    if (@available(macos 13.0, *)) {
+      usable |= (vendor == METAL_GPU_INTEL);
     }
 #  endif
 

intern/cycles/device/oneapi/device_impl.cpp

@@ -377,7 +377,7 @@ void OneapiDevice::tex_alloc(device_texture &mem)
   generic_alloc(mem);
   generic_copy_to(mem);
 
-  /* Resize if needed. Also, in case of resize - allocate in advance for future allocations. */
+  /* Resize if needed. Also, in case of resize - allocate in advance for future allocs. */
   const uint slot = mem.slot;
   if (slot >= texture_info_.size()) {
     texture_info_.resize(slot + 128);
@@ -631,9 +631,9 @@ bool OneapiDevice::enqueue_kernel(KernelContext *kernel_context,
 /* Compute-runtime (ie. NEO) version is what gets returned by sycl/L0 on Windows
  * since Windows driver 101.3268. */
 /* The same min compute-runtime version is currently required across Windows and Linux.
- * For Windows driver 101.4032, compute-runtime version is 24931. */
-static const int lowest_supported_driver_version_win = 1014032;
-static const int lowest_supported_driver_version_neo = 24931;
+ * For Windows driver 101.3430, compute-runtime version is 23904. */
+static const int lowest_supported_driver_version_win = 1013430;
+static const int lowest_supported_driver_version_neo = 23904;
 
 int OneapiDevice::parse_driver_build_version(const sycl::device &device)
 {

intern/cycles/device/queue.h

@@ -112,13 +112,6 @@ class DeviceQueue {
     return 65536;
   }
 
-  /* Does device support local atomic sorting kernels (INTEGRATOR_SORT_BUCKET_PASS and
-   * INTEGRATOR_SORT_WRITE_PASS)? */
-  virtual bool supports_local_atomic_sort() const
-  {
-    return false;
-  }
-
   /* Initialize execution of kernels on this queue.
    *
    * Will, for example, load all data required by the kernels from Device to global or path state.

intern/cycles/graph/CMakeLists.txt

@@ -5,9 +5,6 @@ set(INC
   ..
 )
 
-set(INC_SYS
-)
-
 set(SRC
   node.cpp
   node_type.cpp

intern/cycles/integrator/CMakeLists.txt

@@ -5,9 +5,6 @@ set(INC
   ..
 )
 
-set(INC_SYS
-)
-
 set(SRC
   adaptive_sampling.cpp
   denoiser.cpp

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -71,8 +71,6 @@ PathTraceWorkGPU::PathTraceWorkGPU(Device *device,
           device, "integrator_shader_mnee_sort_counter", MEM_READ_WRITE),
       integrator_shader_sort_prefix_sum_(
           device, "integrator_shader_sort_prefix_sum", MEM_READ_WRITE),
-      integrator_shader_sort_partition_key_offsets_(
-          device, "integrator_shader_sort_partition_key_offsets", MEM_READ_WRITE),
       integrator_next_main_path_index_(device, "integrator_next_main_path_index", MEM_READ_WRITE),
       integrator_next_shadow_path_index_(
           device, "integrator_next_shadow_path_index", MEM_READ_WRITE),
@@ -209,45 +207,33 @@ void PathTraceWorkGPU::alloc_integrator_sorting()
   integrator_state_gpu_.sort_partition_divisor = (int)divide_up(max_num_paths_,
                                                                 num_sort_partitions_);
 
-  if (num_sort_partitions_ > 1 && queue_->supports_local_atomic_sort()) {
-    /* Allocate array for partitioned shader sorting using local atomics. */
-    const int num_offsets = (device_scene_->data.max_shaders + 1) * num_sort_partitions_;
-    if (integrator_shader_sort_partition_key_offsets_.size() < num_offsets) {
-      integrator_shader_sort_partition_key_offsets_.alloc(num_offsets);
-      integrator_shader_sort_partition_key_offsets_.zero_to_device();
-    }
-    integrator_state_gpu_.sort_partition_key_offsets =
-        (int *)integrator_shader_sort_partition_key_offsets_.device_pointer;
+  /* Allocate arrays for shader sorting. */
+  const int sort_buckets = device_scene_->data.max_shaders * num_sort_partitions_;
+  if (integrator_shader_sort_counter_.size() < sort_buckets) {
+    integrator_shader_sort_counter_.alloc(sort_buckets);
+    integrator_shader_sort_counter_.zero_to_device();
+    integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] =
+        (int *)integrator_shader_sort_counter_.device_pointer;
+
+    integrator_shader_sort_prefix_sum_.alloc(sort_buckets);
+    integrator_shader_sort_prefix_sum_.zero_to_device();
   }
-  else {
-    /* Allocate arrays for shader sorting. */
-    const int sort_buckets = device_scene_->data.max_shaders * num_sort_partitions_;
-    if (integrator_shader_sort_counter_.size() < sort_buckets) {
-      integrator_shader_sort_counter_.alloc(sort_buckets);
-      integrator_shader_sort_counter_.zero_to_device();
-      integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] =
-          (int *)integrator_shader_sort_counter_.device_pointer;
 
-      integrator_shader_sort_prefix_sum_.alloc(sort_buckets);
-      integrator_shader_sort_prefix_sum_.zero_to_device();
+  if (device_scene_->data.kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) {
+    if (integrator_shader_raytrace_sort_counter_.size() < sort_buckets) {
+      integrator_shader_raytrace_sort_counter_.alloc(sort_buckets);
+      integrator_shader_raytrace_sort_counter_.zero_to_device();
+      integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE] =
+          (int *)integrator_shader_raytrace_sort_counter_.device_pointer;
     }
+  }
 
-    if (device_scene_->data.kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) {
-      if (integrator_shader_raytrace_sort_counter_.size() < sort_buckets) {
-        integrator_shader_raytrace_sort_counter_.alloc(sort_buckets);
-        integrator_shader_raytrace_sort_counter_.zero_to_device();
-        integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE] =
-            (int *)integrator_shader_raytrace_sort_counter_.device_pointer;
-      }
-    }
-
-    if (device_scene_->data.kernel_features & KERNEL_FEATURE_MNEE) {
-      if (integrator_shader_mnee_sort_counter_.size() < sort_buckets) {
-        integrator_shader_mnee_sort_counter_.alloc(sort_buckets);
-        integrator_shader_mnee_sort_counter_.zero_to_device();
-        integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE] =
-            (int *)integrator_shader_mnee_sort_counter_.device_pointer;
-      }
+  if (device_scene_->data.kernel_features & KERNEL_FEATURE_MNEE) {
+    if (integrator_shader_mnee_sort_counter_.size() < sort_buckets) {
+      integrator_shader_mnee_sort_counter_.alloc(sort_buckets);
+      integrator_shader_mnee_sort_counter_.zero_to_device();
+      integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE] =
+          (int *)integrator_shader_mnee_sort_counter_.device_pointer;
     }
   }
 }
@@ -465,7 +451,8 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
     work_size = num_queued;
     d_path_index = queued_paths_.device_pointer;
 
-    compute_sorted_queued_paths(kernel, num_paths_limit);
+    compute_sorted_queued_paths(
+        DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY, kernel, num_paths_limit);
   }
   else if (num_queued < work_size) {
     work_size = num_queued;
@@ -524,26 +511,11 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
   }
 }
 
-void PathTraceWorkGPU::compute_sorted_queued_paths(DeviceKernel queued_kernel,
+void PathTraceWorkGPU::compute_sorted_queued_paths(DeviceKernel kernel,
+                                                   DeviceKernel queued_kernel,
                                                    const int num_paths_limit)
 {
   int d_queued_kernel = queued_kernel;
-
-  /* Launch kernel to fill the active paths arrays. */
-  if (num_sort_partitions_ > 1 && queue_->supports_local_atomic_sort()) {
-    const int work_size = kernel_max_active_main_path_index(queued_kernel);
-    device_ptr d_queued_paths = queued_paths_.device_pointer;
-
-    int partition_size = (int)integrator_state_gpu_.sort_partition_divisor;
-
-    DeviceKernelArguments args(
-        &work_size, &partition_size, &num_paths_limit, &d_queued_paths, &d_queued_kernel);
-
-    queue_->enqueue(DEVICE_KERNEL_INTEGRATOR_SORT_BUCKET_PASS, 1024 * num_sort_partitions_, args);
-    queue_->enqueue(DEVICE_KERNEL_INTEGRATOR_SORT_WRITE_PASS, 1024 * num_sort_partitions_, args);
-    return;
-  }
-
   device_ptr d_counter = (device_ptr)integrator_state_gpu_.sort_key_counter[d_queued_kernel];
   device_ptr d_prefix_sum = integrator_shader_sort_prefix_sum_.device_pointer;
   assert(d_counter != 0 && d_prefix_sum != 0);
@@ -580,7 +552,7 @@ void PathTraceWorkGPU::compute_sorted_queued_paths(DeviceKernel queued_kernel,
                                &d_prefix_sum,
                                &d_queued_kernel);
 
-    queue_->enqueue(DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY, work_size, args);
+    queue_->enqueue(kernel, work_size, args);
   }
 }
 

intern/cycles/integrator/path_trace_work_gpu.h

@@ -70,7 +70,9 @@ class PathTraceWorkGPU : public PathTraceWork {
   void enqueue_path_iteration(DeviceKernel kernel, const int num_paths_limit = INT_MAX);
 
   void compute_queued_paths(DeviceKernel kernel, DeviceKernel queued_kernel);
-  void compute_sorted_queued_paths(DeviceKernel queued_kernel, const int num_paths_limit);
+  void compute_sorted_queued_paths(DeviceKernel kernel,
+                                   DeviceKernel queued_kernel,
+                                   const int num_paths_limit);
 
   void compact_main_paths(const int num_active_paths);
   void compact_shadow_paths();
@@ -133,7 +135,6 @@ class PathTraceWorkGPU : public PathTraceWork {
   device_vector<int> integrator_shader_raytrace_sort_counter_;
   device_vector<int> integrator_shader_mnee_sort_counter_;
   device_vector<int> integrator_shader_sort_prefix_sum_;
-  device_vector<int> integrator_shader_sort_partition_key_offsets_;
   /* Path split. */
   device_vector<int> integrator_next_main_path_index_;
   device_vector<int> integrator_next_shadow_path_index_;

intern/cycles/kernel/CMakeLists.txt

@@ -732,25 +732,25 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
   endif()
   # SYCL_CPP_FLAGS is a variable that the user can set to pass extra compiler options
   set(sycl_compiler_flags
-    ${CMAKE_CURRENT_SOURCE_DIR}/${SRC_KERNEL_DEVICE_ONEAPI}
-    -fsycl
-    -fsycl-unnamed-lambda
-    -fdelayed-template-parsing
-    -mllvm -inlinedefault-threshold=250
-    -mllvm -inlinehint-threshold=350
-    -fsycl-device-code-split=per_kernel
-    -fsycl-max-parallel-link-jobs=${SYCL_OFFLINE_COMPILER_PARALLEL_JOBS}
-    -shared
-    -DWITH_ONEAPI
-    -ffast-math
-    -DNDEBUG
-    -O2
-    -o ${cycles_kernel_oneapi_lib}
-    -I${CMAKE_CURRENT_SOURCE_DIR}/..
-    ${SYCL_CPP_FLAGS}
-  )
+      ${CMAKE_CURRENT_SOURCE_DIR}/${SRC_KERNEL_DEVICE_ONEAPI}
+      -fsycl
+      -fsycl-unnamed-lambda
+      -fdelayed-template-parsing
+      -mllvm -inlinedefault-threshold=250
+      -mllvm -inlinehint-threshold=350
+      -fsycl-device-code-split=per_kernel
+      -fsycl-max-parallel-link-jobs=${SYCL_OFFLINE_COMPILER_PARALLEL_JOBS}
+      -shared
+      -DWITH_ONEAPI
+      -ffast-math
+      -DNDEBUG
+      -O2
+      -o ${cycles_kernel_oneapi_lib}
+      -I${CMAKE_CURRENT_SOURCE_DIR}/..
+      ${SYCL_CPP_FLAGS}
+      )
 
-  if(WITH_CYCLES_ONEAPI_HOST_TASK_EXECUTION)
+  if (WITH_CYCLES_ONEAPI_HOST_TASK_EXECUTION)
     list(APPEND sycl_compiler_flags -DWITH_ONEAPI_SYCL_HOST_TASK)
   endif()
 

intern/cycles/kernel/bake/bake.h

@@ -63,9 +63,8 @@ ccl_device void kernel_background_evaluate(KernelGlobals kg,
   shader_setup_from_background(kg, &sd, ray_P, ray_D, ray_time);
 
   /* Evaluate shader.
-   * This is being evaluated for all BSDFs, so path flag does not contain a specific type.
-   * However, we want to flag the ray visibility to ignore the sun in the background map. */
-  const uint32_t path_flag = PATH_RAY_EMISSION | PATH_RAY_IMPORTANCE_BAKE;
+   * This is being evaluated for all BSDFs, so path flag does not contain a specific type. */
+  const uint32_t path_flag = PATH_RAY_EMISSION;
   surface_shader_eval<KERNEL_FEATURE_NODE_MASK_SURFACE_LIGHT &
                       ~(KERNEL_FEATURE_NODE_RAYTRACE | KERNEL_FEATURE_NODE_LIGHT_PATH)>(
       kg, INTEGRATOR_STATE_NULL, &sd, NULL, path_flag);

intern/cycles/kernel/closure/bsdf.h

@@ -102,9 +102,10 @@ ccl_device_inline float shift_cos_in(float cos_in, const float frequency_multipl
   return val;
 }
 
-ccl_device_inline bool bsdf_is_transmission(ccl_private const ShaderClosure *sc, const float3 wo)
+ccl_device_inline bool bsdf_is_transmission(ccl_private const ShaderClosure *sc,
+                                            const float3 omega_in)
 {
-  return dot(sc->N, wo) < 0.0f;
+  return dot(sc->N, omega_in) < 0.0f;
 }
 
 ccl_device_inline int bsdf_sample(KernelGlobals kg,
@@ -113,7 +114,7 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
                                   float randu,
                                   float randv,
                                   ccl_private Spectrum *eval,
-                                  ccl_private float3 *wo,
+                                  ccl_private float3 *omega_in,
                                   ccl_private float *pdf,
                                   ccl_private float2 *sampled_roughness,
                                   ccl_private float *eta)
@@ -125,43 +126,43 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
 
   switch (sc->type) {
     case CLOSURE_BSDF_DIFFUSE_ID:
-      label = bsdf_diffuse_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
+      label = bsdf_diffuse_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
 #if defined(__SVM__) || defined(__OSL__)
     case CLOSURE_BSDF_OREN_NAYAR_ID:
-      label = bsdf_oren_nayar_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
+      label = bsdf_oren_nayar_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
 #  ifdef __OSL__
     case CLOSURE_BSDF_PHONG_RAMP_ID:
       label = bsdf_phong_ramp_sample(
-          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness);
+          sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_DIFFUSE_RAMP_ID:
-      label = bsdf_diffuse_ramp_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
+      label = bsdf_diffuse_ramp_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
 #  endif
     case CLOSURE_BSDF_TRANSLUCENT_ID:
-      label = bsdf_translucent_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
+      label = bsdf_translucent_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_REFLECTION_ID:
-      label = bsdf_reflection_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf, eta);
+      label = bsdf_reflection_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, eta);
       *sampled_roughness = zero_float2();
       break;
     case CLOSURE_BSDF_REFRACTION_ID:
-      label = bsdf_refraction_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf, eta);
+      label = bsdf_refraction_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, eta);
       *sampled_roughness = zero_float2();
       break;
     case CLOSURE_BSDF_TRANSPARENT_ID:
-      label = bsdf_transparent_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
+      label = bsdf_transparent_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
       *sampled_roughness = zero_float2();
       *eta = 1.0f;
       break;
@@ -170,65 +171,85 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
     case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
       label = bsdf_microfacet_ggx_sample(
-          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness, eta);
+          kg, sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, sampled_roughness, eta);
       break;
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_ID:
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_FRESNEL_ID:
-      label = bsdf_microfacet_multi_ggx_sample(
-          kg, sc, Ng, sd->wi, randu, randv, eval, wo, pdf, &sd->lcg_state, sampled_roughness, eta);
+      label = bsdf_microfacet_multi_ggx_sample(kg,
+                                               sc,
+                                               Ng,
+                                               sd->I,
+                                               randu,
+                                               randv,
+                                               eval,
+                                               omega_in,
+                                               pdf,
+                                               &sd->lcg_state,
+                                               sampled_roughness,
+                                               eta);
       break;
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_ID:
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_FRESNEL_ID:
-      label = bsdf_microfacet_multi_ggx_glass_sample(
-          kg, sc, Ng, sd->wi, randu, randv, eval, wo, pdf, &sd->lcg_state, sampled_roughness, eta);
+      label = bsdf_microfacet_multi_ggx_glass_sample(kg,
+                                                     sc,
+                                                     Ng,
+                                                     sd->I,
+                                                     randu,
+                                                     randv,
+                                                     eval,
+                                                     omega_in,
+                                                     pdf,
+                                                     &sd->lcg_state,
+                                                     sampled_roughness,
+                                                     eta);
       break;
     case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID:
       label = bsdf_microfacet_beckmann_sample(
-          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness, eta);
+          kg, sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, sampled_roughness, eta);
       break;
     case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
       label = bsdf_ashikhmin_shirley_sample(
-          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness);
+          sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_ASHIKHMIN_VELVET_ID:
-      label = bsdf_ashikhmin_velvet_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
+      label = bsdf_ashikhmin_velvet_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_DIFFUSE_TOON_ID:
-      label = bsdf_diffuse_toon_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
+      label = bsdf_diffuse_toon_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_GLOSSY_TOON_ID:
-      label = bsdf_glossy_toon_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
+      label = bsdf_glossy_toon_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
       // double check if this is valid
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_HAIR_REFLECTION_ID:
       label = bsdf_hair_reflection_sample(
-          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness);
+          sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_HAIR_TRANSMISSION_ID:
       label = bsdf_hair_transmission_sample(
-          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness);
+          sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_HAIR_PRINCIPLED_ID:
       label = bsdf_principled_hair_sample(
-          kg, sc, sd, randu, randv, eval, wo, pdf, sampled_roughness, eta);
+          kg, sc, sd, randu, randv, eval, omega_in, pdf, sampled_roughness, eta);
       break;
     case CLOSURE_BSDF_PRINCIPLED_DIFFUSE_ID:
-      label = bsdf_principled_diffuse_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
+      label = bsdf_principled_diffuse_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_PRINCIPLED_SHEEN_ID:
-      label = bsdf_principled_sheen_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
+      label = bsdf_principled_sheen_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
@@ -253,12 +274,12 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
     const float frequency_multiplier =
         kernel_data_fetch(objects, sd->object).shadow_terminator_shading_offset;
     if (frequency_multiplier > 1.0f) {
-      const float cosNO = dot(*wo, sc->N);
-      *eval *= shift_cos_in(cosNO, frequency_multiplier);
+      const float cosNI = dot(*omega_in, sc->N);
+      *eval *= shift_cos_in(cosNI, frequency_multiplier);
     }
     if (label & LABEL_DIFFUSE) {
       if (!isequal(sc->N, sd->N)) {
-        *eval *= bump_shadowing_term(sd->N, sc->N, *wo);
+        *eval *= bump_shadowing_term(sd->N, sc->N, *omega_in);
       }
     }
   }
@@ -405,7 +426,7 @@ ccl_device_inline void bsdf_roughness_eta(const KernelGlobals kg,
 
 ccl_device_inline int bsdf_label(const KernelGlobals kg,
                                  ccl_private const ShaderClosure *sc,
-                                 const float3 wo)
+                                 const float3 omega_in)
 {
   /* For curves use the smooth normal, particularly for ribbons the geometric
    * normal gives too much darkening otherwise. */
@@ -461,8 +482,8 @@ ccl_device_inline int bsdf_label(const KernelGlobals kg,
     }
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_ID:
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_FRESNEL_ID:
-      label = (bsdf_is_transmission(sc, wo)) ? LABEL_TRANSMIT | LABEL_GLOSSY :
-                                               LABEL_REFLECT | LABEL_GLOSSY;
+      label = (bsdf_is_transmission(sc, omega_in)) ? LABEL_TRANSMIT | LABEL_GLOSSY :
+                                                     LABEL_REFLECT | LABEL_GLOSSY;
       break;
     case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
       label = LABEL_REFLECT | LABEL_GLOSSY;
@@ -483,7 +504,7 @@ ccl_device_inline int bsdf_label(const KernelGlobals kg,
       label = LABEL_TRANSMIT | LABEL_GLOSSY;
       break;
     case CLOSURE_BSDF_HAIR_PRINCIPLED_ID:
-      if (bsdf_is_transmission(sc, wo))
+      if (bsdf_is_transmission(sc, omega_in))
         label = LABEL_TRANSMIT | LABEL_GLOSSY;
       else
         label = LABEL_REFLECT | LABEL_GLOSSY;
@@ -522,83 +543,83 @@ ccl_device_inline
     bsdf_eval(KernelGlobals kg,
               ccl_private ShaderData *sd,
               ccl_private const ShaderClosure *sc,
-              const float3 wo,
+              const float3 omega_in,
               ccl_private float *pdf)
 {
   Spectrum eval = zero_spectrum();
 
   switch (sc->type) {
     case CLOSURE_BSDF_DIFFUSE_ID:
-      eval = bsdf_diffuse_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_diffuse_eval(sc, sd->I, omega_in, pdf);
       break;
 #if defined(__SVM__) || defined(__OSL__)
     case CLOSURE_BSDF_OREN_NAYAR_ID:
-      eval = bsdf_oren_nayar_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_oren_nayar_eval(sc, sd->I, omega_in, pdf);
       break;
 #  ifdef __OSL__
     case CLOSURE_BSDF_PHONG_RAMP_ID:
-      eval = bsdf_phong_ramp_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_phong_ramp_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_DIFFUSE_RAMP_ID:
-      eval = bsdf_diffuse_ramp_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_diffuse_ramp_eval(sc, sd->I, omega_in, pdf);
       break;
 #  endif
     case CLOSURE_BSDF_TRANSLUCENT_ID:
-      eval = bsdf_translucent_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_translucent_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_REFLECTION_ID:
-      eval = bsdf_reflection_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_reflection_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_REFRACTION_ID:
-      eval = bsdf_refraction_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_refraction_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_TRANSPARENT_ID:
-      eval = bsdf_transparent_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_transparent_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_MICROFACET_GGX_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_FRESNEL_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
     case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
-      eval = bsdf_microfacet_ggx_eval(sc, sd->N, sd->wi, wo, pdf);
+      eval = bsdf_microfacet_ggx_eval(sc, sd->N, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_ID:
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_FRESNEL_ID:
-      eval = bsdf_microfacet_multi_ggx_eval(sc, sd->N, sd->wi, wo, pdf, &sd->lcg_state);
+      eval = bsdf_microfacet_multi_ggx_eval(sc, sd->N, sd->I, omega_in, pdf, &sd->lcg_state);
       break;
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_ID:
     case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_FRESNEL_ID:
-      eval = bsdf_microfacet_multi_ggx_glass_eval(sc, sd->wi, wo, pdf, &sd->lcg_state);
+      eval = bsdf_microfacet_multi_ggx_glass_eval(sc, sd->I, omega_in, pdf, &sd->lcg_state);
       break;
     case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
     case CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID:
-      eval = bsdf_microfacet_beckmann_eval(sc, sd->N, sd->wi, wo, pdf);
+      eval = bsdf_microfacet_beckmann_eval(sc, sd->N, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
-      eval = bsdf_ashikhmin_shirley_eval(sc, sd->N, sd->wi, wo, pdf);
+      eval = bsdf_ashikhmin_shirley_eval(sc, sd->N, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_ASHIKHMIN_VELVET_ID:
-      eval = bsdf_ashikhmin_velvet_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_ashikhmin_velvet_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_DIFFUSE_TOON_ID:
-      eval = bsdf_diffuse_toon_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_diffuse_toon_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_GLOSSY_TOON_ID:
-      eval = bsdf_glossy_toon_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_glossy_toon_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_HAIR_PRINCIPLED_ID:
-      eval = bsdf_principled_hair_eval(kg, sd, sc, wo, pdf);
+      eval = bsdf_principled_hair_eval(kg, sd, sc, omega_in, pdf);
       break;
     case CLOSURE_BSDF_HAIR_REFLECTION_ID:
-      eval = bsdf_hair_reflection_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_hair_reflection_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_HAIR_TRANSMISSION_ID:
-      eval = bsdf_hair_transmission_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_hair_transmission_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_PRINCIPLED_DIFFUSE_ID:
-      eval = bsdf_principled_diffuse_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_principled_diffuse_eval(sc, sd->I, omega_in, pdf);
       break;
     case CLOSURE_BSDF_PRINCIPLED_SHEEN_ID:
-      eval = bsdf_principled_sheen_eval(sc, sd->wi, wo, pdf);
+      eval = bsdf_principled_sheen_eval(sc, sd->I, omega_in, pdf);
       break;
 #endif
     default:
@@ -607,7 +628,7 @@ ccl_device_inline
 
   if (CLOSURE_IS_BSDF_DIFFUSE(sc->type)) {
     if (!isequal(sc->N, sd->N)) {
-      eval *= bump_shadowing_term(sd->N, sc->N, wo);
+      eval *= bump_shadowing_term(sd->N, sc->N, omega_in);
     }
   }
 
@@ -615,9 +636,9 @@ ccl_device_inline
   const float frequency_multiplier =
       kernel_data_fetch(objects, sd->object).shadow_terminator_shading_offset;
   if (frequency_multiplier > 1.0f) {
-    const float cosNO = dot(wo, sc->N);
-    if (cosNO >= 0.0f) {
-      eval *= shift_cos_in(cosNO, frequency_multiplier);
+    const float cosNI = dot(omega_in, sc->N);
+    if (cosNI >= 0.0f) {
+      eval *= shift_cos_in(cosNI, frequency_multiplier);
     }
   }
 
@@ -661,38 +682,4 @@ ccl_device void bsdf_blur(KernelGlobals kg, ccl_private ShaderClosure *sc, float
 #endif
 }
 
-ccl_device_inline Spectrum bsdf_albedo(ccl_private const ShaderData *sd,
-                                       ccl_private const ShaderClosure *sc)
-{
-  Spectrum albedo = sc->weight;
-  /* Some closures include additional components such as Fresnel terms that cause their albedo to
-   * be below 1. The point of this function is to return a best-effort estimation of their albedo,
-   * meaning the amount of reflected/refracted light that would be expected when illuminated by a
-   * uniform white background.
-   * This is used for the denoising albedo pass and diffuse/glossy/transmission color passes.
-   * NOTE: This should always match the sample_weight of the closure - as in, if there's an albedo
-   * adjustment in here, the sample_weight should also be reduced accordingly.
-   * TODO(lukas): Consider calling this function to determine the sample_weight? Would be a bit of
-   * extra overhead though. */
-#if defined(__SVM__) || defined(__OSL__)
-  switch (sc->type) {
-    case CLOSURE_BSDF_MICROFACET_GGX_FRESNEL_ID:
-    case CLOSURE_BSDF_MICROFACET_MULTI_GGX_FRESNEL_ID:
-    case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_FRESNEL_ID:
-    case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
-      albedo *= microfacet_fresnel((ccl_private const MicrofacetBsdf *)sc, sd->wi, sc->N);
-      break;
-    case CLOSURE_BSDF_PRINCIPLED_SHEEN_ID:
-      albedo *= ((ccl_private const PrincipledSheenBsdf *)sc)->avg_value;
-      break;
-    case CLOSURE_BSDF_HAIR_PRINCIPLED_ID:
-      albedo *= bsdf_principled_hair_albedo(sc);
-      break;
-    default:
-      break;
-  }
-#endif
-  return albedo;
-}
-
 CCL_NAMESPACE_END

intern/cycles/kernel/closure/bsdf_ashikhmin_shirley.h

@@ -41,20 +41,20 @@ ccl_device_inline float bsdf_ashikhmin_shirley_roughness_to_exponent(float rough
 
 ccl_device_forceinline Spectrum bsdf_ashikhmin_shirley_eval(ccl_private const ShaderClosure *sc,
                                                             const float3 Ng,
-                                                            const float3 wi,
-                                                            const float3 wo,
+                                                            const float3 I,
+                                                            const float3 omega_in,
                                                             ccl_private float *pdf)
 {
   ccl_private const MicrofacetBsdf *bsdf = (ccl_private const MicrofacetBsdf *)sc;
-  const float cosNgO = dot(Ng, wo);
+  const float cosNgI = dot(Ng, omega_in);
   float3 N = bsdf->N;
 
-  float NdotI = dot(N, wi);
-  float NdotO = dot(N, wo);
+  float NdotI = dot(N, I);        /* in Cycles/OSL convention I is omega_out */
+  float NdotO = dot(N, omega_in); /* and consequently we use for O omaga_in ;) */
 
   float out = 0.0f;
 
-  if ((cosNgO < 0.0f) || fmaxf(bsdf->alpha_x, bsdf->alpha_y) <= 1e-4f ||
+  if ((cosNgI < 0.0f) || fmaxf(bsdf->alpha_x, bsdf->alpha_y) <= 1e-4f ||
       !(NdotI > 0.0f && NdotO > 0.0f)) {
     *pdf = 0.0f;
     return zero_spectrum();
@@ -62,15 +62,15 @@ ccl_device_forceinline Spectrum bsdf_ashikhmin_shirley_eval(ccl_private const Sh
 
   NdotI = fmaxf(NdotI, 1e-6f);
   NdotO = fmaxf(NdotO, 1e-6f);
-  float3 H = normalize(wi + wo);
-  float HdotI = fmaxf(fabsf(dot(H, wi)), 1e-6f);
+  float3 H = normalize(omega_in + I);
+  float HdotI = fmaxf(fabsf(dot(H, I)), 1e-6f);
   float HdotN = fmaxf(dot(H, N), 1e-6f);
 
   /* pump from original paper
    * (first derivative disc., but cancels the HdotI in the pdf nicely) */
-  float pump = 1.0f / fmaxf(1e-6f, (HdotI * fmaxf(NdotI, NdotO)));
+  float pump = 1.0f / fmaxf(1e-6f, (HdotI * fmaxf(NdotO, NdotI)));
   /* pump from d-brdf paper */
-  /*float pump = 1.0f / fmaxf(1e-4f, ((NdotI + NdotO) * (NdotI * NdotO))); */
+  /*float pump = 1.0f / fmaxf(1e-4f, ((NdotO + NdotI) * (NdotO*NdotI))); */
 
   float n_x = bsdf_ashikhmin_shirley_roughness_to_exponent(bsdf->alpha_x);
   float n_y = bsdf_ashikhmin_shirley_roughness_to_exponent(bsdf->alpha_y);
@@ -124,11 +124,11 @@ ccl_device_inline void bsdf_ashikhmin_shirley_sample_first_quadrant(float n_x,
 
 ccl_device int bsdf_ashikhmin_shirley_sample(ccl_private const ShaderClosure *sc,
                                              float3 Ng,
-                                             float3 wi,
+                                             float3 I,
                                              float randu,
                                              float randv,
                                              ccl_private Spectrum *eval,
-                                             ccl_private float3 *wo,
+                                             ccl_private float3 *omega_in,
                                              ccl_private float *pdf,
                                              ccl_private float2 *sampled_roughness)
 {
@@ -137,7 +137,7 @@ ccl_device int bsdf_ashikhmin_shirley_sample(ccl_private const ShaderClosure *sc
   float3 N = bsdf->N;
   int label = LABEL_REFLECT | LABEL_GLOSSY;
 
-  float NdotI = dot(N, wi);
+  float NdotI = dot(N, I);
   if (!(NdotI > 0.0f)) {
     *pdf = 0.0f;
     *eval = zero_spectrum();
@@ -198,12 +198,12 @@ ccl_device int bsdf_ashikhmin_shirley_sample(ccl_private const ShaderClosure *sc
 
   /* half vector to world space */
   float3 H = h.x * X + h.y * Y + h.z * N;
-  float HdotI = dot(H, wi);
+  float HdotI = dot(H, I);
   if (HdotI < 0.0f)
     H = -H;
 
-  /* reflect wi on H to get wo */
-  *wo = -wi + (2.0f * HdotI) * H;
+  /* reflect I on H to get omega_in */
+  *omega_in = -I + (2.0f * HdotI) * H;
 
   if (fmaxf(bsdf->alpha_x, bsdf->alpha_y) <= 1e-4f) {
     /* Some high number for MIS. */
@@ -213,7 +213,7 @@ ccl_device int bsdf_ashikhmin_shirley_sample(ccl_private const ShaderClosure *sc
   }
   else {
     /* leave the rest to eval */
-    *eval = bsdf_ashikhmin_shirley_eval(sc, N, wi, *wo, pdf);
+    *eval = bsdf_ashikhmin_shirley_eval(sc, N, I, *omega_in, pdf);
   }
 
   return label;

intern/cycles/kernel/closure/bsdf_ashikhmin_velvet.h

@@ -32,35 +32,35 @@ ccl_device int bsdf_ashikhmin_velvet_setup(ccl_private VelvetBsdf *bsdf)
 }
 
 ccl_device Spectrum bsdf_ashikhmin_velvet_eval(ccl_private const ShaderClosure *sc,
-                                               const float3 wi,
-                                               const float3 wo,
+                                               const float3 I,
+                                               const float3 omega_in,
                                                ccl_private float *pdf)
 {
   ccl_private const VelvetBsdf *bsdf = (ccl_private const VelvetBsdf *)sc;
   float m_invsigma2 = bsdf->invsigma2;
   float3 N = bsdf->N;
 
-  float cosNI = dot(N, wi);
-  float cosNO = dot(N, wo);
-  if (!(cosNI > 0 && cosNO > 0)) {
+  float cosNO = dot(N, I);
+  float cosNI = dot(N, omega_in);
+  if (!(cosNO > 0 && cosNI > 0)) {
     *pdf = 0.0f;
     return zero_spectrum();
   }
 
-  float3 H = normalize(wi + wo);
+  float3 H = normalize(omega_in + I);
 
   float cosNH = dot(N, H);
-  float cosHI = fabsf(dot(wi, H));
+  float cosHO = fabsf(dot(I, H));
 
-  if (!(fabsf(cosNH) < 1.0f - 1e-5f && cosHI > 1e-5f)) {
+  if (!(fabsf(cosNH) < 1.0f - 1e-5f && cosHO > 1e-5f)) {
     *pdf = 0.0f;
     return zero_spectrum();
   }
-  float cosNHdivHI = cosNH / cosHI;
-  cosNHdivHI = fmaxf(cosNHdivHI, 1e-5f);
+  float cosNHdivHO = cosNH / cosHO;
+  cosNHdivHO = fmaxf(cosNHdivHO, 1e-5f);
 
-  float fac1 = 2 * fabsf(cosNHdivHI * cosNI);
-  float fac2 = 2 * fabsf(cosNHdivHI * cosNO);
+  float fac1 = 2 * fabsf(cosNHdivHO * cosNO);
+  float fac2 = 2 * fabsf(cosNHdivHO * cosNI);
 
   float sinNH2 = 1 - cosNH * cosNH;
   float sinNH4 = sinNH2 * sinNH2;
@@ -69,7 +69,7 @@ ccl_device Spectrum bsdf_ashikhmin_velvet_eval(ccl_private const ShaderClosure *
   float D = expf(-cotangent2 * m_invsigma2) * m_invsigma2 * M_1_PI_F / sinNH4;
   float G = fminf(1.0f, fminf(fac1, fac2));  // TODO: derive G from D analytically
 
-  float out = 0.25f * (D * G) / cosNI;
+  float out = 0.25f * (D * G) / cosNO;
 
   *pdf = 0.5f * M_1_PI_F;
   return make_spectrum(out);
@@ -77,11 +77,11 @@ ccl_device Spectrum bsdf_ashikhmin_velvet_eval(ccl_private const ShaderClosure *
 
 ccl_device int bsdf_ashikhmin_velvet_sample(ccl_private const ShaderClosure *sc,
                                             float3 Ng,
-                                            float3 wi,
+                                            float3 I,
                                             float randu,
                                             float randv,
                                             ccl_private Spectrum *eval,
-                                            ccl_private float3 *wo,
+                                            ccl_private float3 *omega_in,
                                             ccl_private float *pdf)
 {
   ccl_private const VelvetBsdf *bsdf = (ccl_private const VelvetBsdf *)sc;
@@ -90,32 +90,32 @@ ccl_device int bsdf_ashikhmin_velvet_sample(ccl_private const ShaderClosure *sc,
 
   // we are viewing the surface from above - send a ray out with uniform
   // distribution over the hemisphere
-  sample_uniform_hemisphere(N, randu, randv, wo, pdf);
+  sample_uniform_hemisphere(N, randu, randv, omega_in, pdf);
 
-  if (!(dot(Ng, *wo) > 0)) {
+  if (!(dot(Ng, *omega_in) > 0)) {
     *pdf = 0.0f;
     *eval = zero_spectrum();
     return LABEL_NONE;
   }
 
-  float3 H = normalize(wi + *wo);
+  float3 H = normalize(*omega_in + I);
 
-  float cosNI = dot(N, wi);
-  float cosNO = dot(N, *wo);
-  float cosHI = fabsf(dot(wi, H));
+  float cosNI = dot(N, *omega_in);
+  float cosNO = dot(N, I);
   float cosNH = dot(N, H);
+  float cosHO = fabsf(dot(I, H));
 
-  if (!(fabsf(cosNI) > 1e-5f && fabsf(cosNH) < 1.0f - 1e-5f && cosHI > 1e-5f)) {
+  if (!(fabsf(cosNO) > 1e-5f && fabsf(cosNH) < 1.0f - 1e-5f && cosHO > 1e-5f)) {
     *pdf = 0.0f;
     *eval = zero_spectrum();
     return LABEL_NONE;
   }
 
-  float cosNHdivHI = cosNH / cosHI;
-  cosNHdivHI = fmaxf(cosNHdivHI, 1e-5f);
+  float cosNHdivHO = cosNH / cosHO;
+  cosNHdivHO = fmaxf(cosNHdivHO, 1e-5f);
 
-  float fac1 = 2 * fabsf(cosNHdivHI * cosNI);
-  float fac2 = 2 * fabsf(cosNHdivHI * cosNO);
+  float fac1 = 2 * fabsf(cosNHdivHO * cosNO);
+  float fac2 = 2 * fabsf(cosNHdivHO * cosNI);
 
   float sinNH2 = 1 - cosNH * cosNH;
   float sinNH4 = sinNH2 * sinNH2;
@@ -124,7 +124,7 @@ ccl_device int bsdf_ashikhmin_velvet_sample(ccl_private const ShaderClosure *sc,
   float D = expf(-cotangent2 * m_invsigma2) * m_invsigma2 * M_1_PI_F / sinNH4;
   float G = fminf(1.0f, fminf(fac1, fac2));  // TODO: derive G from D analytically
 
-  float power = 0.25f * (D * G) / cosNI;
+  float power = 0.25f * (D * G) / cosNO;
 
   *eval = make_spectrum(power);
 

intern/cycles/kernel/closure/bsdf_diffuse.h

@@ -27,34 +27,34 @@ ccl_device int bsdf_diffuse_setup(ccl_private DiffuseBsdf *bsdf)
 }
 
 ccl_device Spectrum bsdf_diffuse_eval(ccl_private const ShaderClosure *sc,
-                                      const float3 wi,
-                                      const float3 wo,
+                                      const float3 I,
+                                      const float3 omega_in,
                                       ccl_private float *pdf)
 {
   ccl_private const DiffuseBsdf *bsdf = (ccl_private const DiffuseBsdf *)sc;
   float3 N = bsdf->N;
 
-  float cosNO = fmaxf(dot(N, wo), 0.0f) * M_1_PI_F;
-  *pdf = cosNO;
-  return make_spectrum(cosNO);
+  float cos_pi = fmaxf(dot(N, omega_in), 0.0f) * M_1_PI_F;
+  *pdf = cos_pi;
+  return make_spectrum(cos_pi);
 }
 
 ccl_device int bsdf_diffuse_sample(ccl_private const ShaderClosure *sc,
                                    float3 Ng,
-                                   float3 wi,
+                                   float3 I,
                                    float randu,
                                    float randv,
                                    ccl_private Spectrum *eval,
-                                   ccl_private float3 *wo,
+                                   ccl_private float3 *omega_in,
                                    ccl_private float *pdf)
 {
   ccl_private const DiffuseBsdf *bsdf = (ccl_private const DiffuseBsdf *)sc;
   float3 N = bsdf->N;
 
   // distribution over the hemisphere
-  sample_cos_hemisphere(N, randu, randv, wo, pdf);
+  sample_cos_hemisphere(N, randu, randv, omega_in, pdf);
 
-  if (dot(Ng, *wo) > 0.0f) {
+  if (dot(Ng, *omega_in) > 0.0f) {
     *eval = make_spectrum(*pdf);
   }
   else {
@@ -73,25 +73,25 @@ ccl_device int bsdf_translucent_setup(ccl_private DiffuseBsdf *bsdf)
 }
 
 ccl_device Spectrum bsdf_translucent_eval(ccl_private const ShaderClosure *sc,
-                                          const float3 wi,
-                                          const float3 wo,
+                                          const float3 I,
+                                          const float3 omega_in,
                                           ccl_private float *pdf)
 {
   ccl_private const DiffuseBsdf *bsdf = (ccl_private const DiffuseBsdf *)sc;
   float3 N = bsdf->N;
 
-  float cosNO = fmaxf(-dot(N, wo), 0.0f) * M_1_PI_F;
-  *pdf = cosNO;
-  return make_spectrum(cosNO);
+  float cos_pi = fmaxf(-dot(N, omega_in), 0.0f) * M_1_PI_F;
+  *pdf = cos_pi;
+  return make_spectrum(cos_pi);
 }
 
 ccl_device int bsdf_translucent_sample(ccl_private const ShaderClosure *sc,
                                        float3 Ng,
-                                       float3 wi,
+                                       float3 I,
                                        float randu,
                                        float randv,
                                        ccl_private Spectrum *eval,
-                                       ccl_private float3 *wo,
+                                       ccl_private float3 *omega_in,
                                        ccl_private float *pdf)
 {
   ccl_private const DiffuseBsdf *bsdf = (ccl_private const DiffuseBsdf *)sc;
@@ -99,8 +99,8 @@ ccl_device int bsdf_translucent_sample(ccl_private const ShaderClosure *sc,
 
   // we are viewing the surface from the right side - send a ray out with cosine
   // distribution over the hemisphere
-  sample_cos_hemisphere(-N, randu, randv, wo, pdf);
-  if (dot(Ng, *wo) < 0) {
+  sample_cos_hemisphere(-N, randu, randv, omega_in, pdf);
+  if (dot(Ng, *omega_in) < 0) {
     *eval = make_spectrum(*pdf);
   }
   else {

intern/cycles/kernel/closure/bsdf_diffuse_ramp.h

@@ -48,17 +48,17 @@ ccl_device void bsdf_diffuse_ramp_blur(ccl_private ShaderClosure *sc, float roug
 }
 
 ccl_device Spectrum bsdf_diffuse_ramp_eval(ccl_private const ShaderClosure *sc,
-                                           const float3 wi,
-                                           const float3 wo,
+                                           const float3 I,
+                                           const float3 omega_in,
                                            ccl_private float *pdf)
 {
   const DiffuseRampBsdf *bsdf = (const DiffuseRampBsdf *)sc;
   float3 N = bsdf->N;
 
-  float cosNO = fmaxf(dot(N, wo), 0.0f);
-  if (cosNO >= 0.0f) {
-    *pdf = cosNO * M_1_PI_F;
-    return rgb_to_spectrum(bsdf_diffuse_ramp_get_color(bsdf->colors, cosNO) * M_1_PI_F);
+  float cos_pi = fmaxf(dot(N, omega_in), 0.0f);
+  if (cos_pi >= 0.0f) {
+    *pdf = cos_pi * M_1_PI_F;
+    return rgb_to_spectrum(bsdf_diffuse_ramp_get_color(bsdf->colors, cos_pi) * M_1_PI_F);
   }
   else {
     *pdf = 0.0f;
@@ -68,20 +68,20 @@ ccl_device Spectrum bsdf_diffuse_ramp_eval(ccl_private const ShaderClosure *sc,
 
 ccl_device int bsdf_diffuse_ramp_sample(ccl_private const ShaderClosure *sc,
                                         float3 Ng,
-                                        float3 wi,
+                                        float3 I,
                                         float randu,
                                         float randv,
                                         ccl_private Spectrum *eval,
-                                        ccl_private float3 *wo,
+                                        ccl_private float3 *omega_in,
                                         ccl_private float *pdf)
 {
   const DiffuseRampBsdf *bsdf = (const DiffuseRampBsdf *)sc;
   float3 N = bsdf->N;
 
   // distribution over the hemisphere
-  sample_cos_hemisphere(N, randu, randv, wo, pdf);
+  sample_cos_hemisphere(N, randu, randv, omega_in, pdf);
 
-  if (dot(Ng, *wo) > 0.0f) {
+  if (dot(Ng, *omega_in) > 0.0f) {
     *eval = rgb_to_spectrum(bsdf_diffuse_ramp_get_color(bsdf->colors, *pdf * M_PI_F) * M_1_PI_F);
   }
   else {

intern/cycles/kernel/closure/bsdf_hair.h

@@ -38,12 +38,12 @@ ccl_device int bsdf_hair_transmission_setup(ccl_private HairBsdf *bsdf)
 }
 
 ccl_device Spectrum bsdf_hair_reflection_eval(ccl_private const ShaderClosure *sc,
-                                              const float3 wi,
-                                              const float3 wo,
+                                              const float3 I,
+                                              const float3 omega_in,
                                               ccl_private float *pdf)
 {
   ccl_private const HairBsdf *bsdf = (ccl_private const HairBsdf *)sc;
-  if (dot(bsdf->N, wo) < 0.0f) {
+  if (dot(bsdf->N, omega_in) < 0.0f) {
     *pdf = 0.0f;
     return zero_spectrum();
   }
@@ -53,16 +53,16 @@ ccl_device Spectrum bsdf_hair_reflection_eval(ccl_private const ShaderClosure *s
   float roughness1 = bsdf->roughness1;
   float roughness2 = bsdf->roughness2;
 
-  float Iz = dot(Tg, wi);
-  float3 locy = normalize(wi - Tg * Iz);
+  float Iz = dot(Tg, I);
+  float3 locy = normalize(I - Tg * Iz);
 
   float theta_r = M_PI_2_F - fast_acosf(Iz);
 
-  float wo_z = dot(Tg, wo);
-  float3 wo_y = normalize(wo - Tg * wo_z);
+  float omega_in_z = dot(Tg, omega_in);
+  float3 omega_in_y = normalize(omega_in - Tg * omega_in_z);
 
-  float theta_i = M_PI_2_F - fast_acosf(wo_z);
-  float cosphi_i = dot(wo_y, locy);
+  float theta_i = M_PI_2_F - fast_acosf(omega_in_z);
+  float cosphi_i = dot(omega_in_y, locy);
 
   if (M_PI_2_F - fabsf(theta_i) < 0.001f || cosphi_i < 0.0f) {
     *pdf = 0.0f;
@@ -90,12 +90,12 @@ ccl_device Spectrum bsdf_hair_reflection_eval(ccl_private const ShaderClosure *s
 }
 
 ccl_device Spectrum bsdf_hair_transmission_eval(ccl_private const ShaderClosure *sc,
-                                                const float3 wi,
-                                                const float3 wo,
+                                                const float3 I,
+                                                const float3 omega_in,
                                                 ccl_private float *pdf)
 {
   ccl_private const HairBsdf *bsdf = (ccl_private const HairBsdf *)sc;
-  if (dot(bsdf->N, wo) >= 0.0f) {
+  if (dot(bsdf->N, omega_in) >= 0.0f) {
     *pdf = 0.0f;
     return zero_spectrum();
   }
@@ -104,16 +104,16 @@ ccl_device Spectrum bsdf_hair_transmission_eval(ccl_private const ShaderClosure
   float3 Tg = bsdf->T;
   float roughness1 = bsdf->roughness1;
   float roughness2 = bsdf->roughness2;
-  float Iz = dot(Tg, wi);
-  float3 locy = normalize(wi - Tg * Iz);
+  float Iz = dot(Tg, I);
+  float3 locy = normalize(I - Tg * Iz);
 
   float theta_r = M_PI_2_F - fast_acosf(Iz);
 
-  float wo_z = dot(Tg, wo);
-  float3 wo_y = normalize(wo - Tg * wo_z);
+  float omega_in_z = dot(Tg, omega_in);
+  float3 omega_in_y = normalize(omega_in - Tg * omega_in_z);
 
-  float theta_i = M_PI_2_F - fast_acosf(wo_z);
-  float phi_i = fast_acosf(dot(wo_y, locy));
+  float theta_i = M_PI_2_F - fast_acosf(omega_in_z);
+  float phi_i = fast_acosf(dot(omega_in_y, locy));
 
   if (M_PI_2_F - fabsf(theta_i) < 0.001f) {
     *pdf = 0.0f;
@@ -142,11 +142,11 @@ ccl_device Spectrum bsdf_hair_transmission_eval(ccl_private const ShaderClosure
 
 ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
                                            float3 Ng,
-                                           float3 wi,
+                                           float3 I,
                                            float randu,
                                            float randv,
                                            ccl_private Spectrum *eval,
-                                           ccl_private float3 *wo,
+                                           ccl_private float3 *omega_in,
                                            ccl_private float *pdf,
                                            ccl_private float2 *sampled_roughness)
 {
@@ -156,8 +156,8 @@ ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
   float roughness1 = bsdf->roughness1;
   float roughness2 = bsdf->roughness2;
   *sampled_roughness = make_float2(roughness1, roughness2);
-  float Iz = dot(Tg, wi);
-  float3 locy = normalize(wi - Tg * Iz);
+  float Iz = dot(Tg, I);
+  float3 locy = normalize(I - Tg * Iz);
   float3 locx = cross(locy, Tg);
   float theta_r = M_PI_2_F - fast_acosf(Iz);
 
@@ -182,7 +182,7 @@ ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
 
   float sinphi, cosphi;
   fast_sincosf(phi, &sinphi, &cosphi);
-  *wo = (cosphi * costheta_i) * locy - (sinphi * costheta_i) * locx + (sintheta_i)*Tg;
+  *omega_in = (cosphi * costheta_i) * locy - (sinphi * costheta_i) * locx + (sintheta_i)*Tg;
 
   *pdf = fabsf(phi_pdf * theta_pdf);
   if (M_PI_2_F - fabsf(theta_i) < 0.001f)
@@ -195,11 +195,11 @@ ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
 
 ccl_device int bsdf_hair_transmission_sample(ccl_private const ShaderClosure *sc,
                                              float3 Ng,
-                                             float3 wi,
+                                             float3 I,
                                              float randu,
                                              float randv,
                                              ccl_private Spectrum *eval,
-                                             ccl_private float3 *wo,
+                                             ccl_private float3 *omega_in,
                                              ccl_private float *pdf,
                                              ccl_private float2 *sampled_roughness)
 {
@@ -209,8 +209,8 @@ ccl_device int bsdf_hair_transmission_sample(ccl_private const ShaderClosure *sc
   float roughness1 = bsdf->roughness1;
   float roughness2 = bsdf->roughness2;
   *sampled_roughness = make_float2(roughness1, roughness2);
-  float Iz = dot(Tg, wi);
-  float3 locy = normalize(wi - Tg * Iz);
+  float Iz = dot(Tg, I);
+  float3 locy = normalize(I - Tg * Iz);
   float3 locx = cross(locy, Tg);
   float theta_r = M_PI_2_F - fast_acosf(Iz);
 
@@ -235,7 +235,7 @@ ccl_device int bsdf_hair_transmission_sample(ccl_private const ShaderClosure *sc
 
   float sinphi, cosphi;
   fast_sincosf(phi, &sinphi, &cosphi);
-  *wo = (cosphi * costheta_i) * locy - (sinphi * costheta_i) * locx + (sintheta_i)*Tg;
+  *omega_in = (cosphi * costheta_i) * locy - (sinphi * costheta_i) * locx + (sintheta_i)*Tg;
 
   *pdf = fabsf(phi_pdf * theta_pdf);
   if (M_PI_2_F - fabsf(theta_i) < 0.001f) {
@@ -247,7 +247,7 @@ ccl_device int bsdf_hair_transmission_sample(ccl_private const ShaderClosure *sc
   /* TODO(sergey): Should always be negative, but seems some precision issue
    * is involved here.
    */
-  kernel_assert(dot(locy, *wo) < 1e-4f);
+  kernel_assert(dot(locy, *omega_in) < 1e-4f);
 
   return LABEL_TRANSMIT | LABEL_GLOSSY;
 }

intern/cycles/kernel/closure/bsdf_hair_principled.h

@@ -41,6 +41,11 @@ static_assert(sizeof(ShaderClosure) >= sizeof(PrincipledHairBSDF),
 static_assert(sizeof(ShaderClosure) >= sizeof(PrincipledHairExtra),
               "PrincipledHairExtra is too large!");
 
+ccl_device_inline float cos_from_sin(const float s)
+{
+  return safe_sqrtf(1.0f - s * s);
+}
+
 /* Gives the change in direction in the normal plane for the given angles and p-th-order
  * scattering. */
 ccl_device_inline float delta_phi(int p, float gamma_o, float gamma_t)
@@ -174,7 +179,7 @@ ccl_device int bsdf_principled_hair_setup(ccl_private ShaderData *sd,
 
   /* Compute local frame, aligned to curve tangent and ray direction. */
   float3 X = safe_normalize(sd->dPdu);
-  float3 Y = safe_normalize(cross(X, sd->wi));
+  float3 Y = safe_normalize(cross(X, sd->I));
   float3 Z = safe_normalize(cross(X, Y));
 
   /* h -1..0..1 means the rays goes from grazing the hair, to hitting it at
@@ -254,7 +259,7 @@ ccl_device_inline void hair_alpha_angles(float sin_theta_i,
 ccl_device Spectrum bsdf_principled_hair_eval(KernelGlobals kg,
                                               ccl_private const ShaderData *sd,
                                               ccl_private const ShaderClosure *sc,
-                                              const float3 wo,
+                                              const float3 omega_in,
                                               ccl_private float *pdf)
 {
   kernel_assert(isfinite_safe(sd->P) && isfinite_safe(sd->ray_length));
@@ -266,13 +271,12 @@ ccl_device Spectrum bsdf_principled_hair_eval(KernelGlobals kg,
   kernel_assert(fabsf(dot(X, Y)) < 1e-3f);
   const float3 Z = safe_normalize(cross(X, Y));
 
-  /* local_I is the illumination direction. */
-  const float3 local_O = make_float3(dot(sd->wi, X), dot(sd->wi, Y), dot(sd->wi, Z));
-  const float3 local_I = make_float3(dot(wo, X), dot(wo, Y), dot(wo, Z));
+  const float3 wo = make_float3(dot(sd->I, X), dot(sd->I, Y), dot(sd->I, Z));
+  const float3 wi = make_float3(dot(omega_in, X), dot(omega_in, Y), dot(omega_in, Z));
 
-  const float sin_theta_o = local_O.x;
+  const float sin_theta_o = wo.x;
   const float cos_theta_o = cos_from_sin(sin_theta_o);
-  const float phi_o = atan2f(local_O.z, local_O.y);
+  const float phi_o = atan2f(wo.z, wo.y);
 
   const float sin_theta_t = sin_theta_o / bsdf->eta;
   const float cos_theta_t = cos_from_sin(sin_theta_t);
@@ -291,9 +295,9 @@ ccl_device Spectrum bsdf_principled_hair_eval(KernelGlobals kg,
   hair_attenuation(
       kg, fresnel_dielectric_cos(cos_theta_o * cos_gamma_o, bsdf->eta), T, Ap, Ap_energy);
 
-  const float sin_theta_i = local_I.x;
+  const float sin_theta_i = wi.x;
   const float cos_theta_i = cos_from_sin(sin_theta_i);
-  const float phi_i = atan2f(local_I.z, local_I.y);
+  const float phi_i = atan2f(wi.z, wi.y);
 
   const float phi = phi_i - phi_o;
 
@@ -339,7 +343,7 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
                                            float randu,
                                            float randv,
                                            ccl_private Spectrum *eval,
-                                           ccl_private float3 *wo,
+                                           ccl_private float3 *omega_in,
                                            ccl_private float *pdf,
                                            ccl_private float2 *sampled_roughness,
                                            ccl_private float *eta)
@@ -355,16 +359,16 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
   kernel_assert(fabsf(dot(X, Y)) < 1e-3f);
   const float3 Z = safe_normalize(cross(X, Y));
 
-  const float3 local_O = make_float3(dot(sd->wi, X), dot(sd->wi, Y), dot(sd->wi, Z));
+  const float3 wo = make_float3(dot(sd->I, X), dot(sd->I, Y), dot(sd->I, Z));
 
   float2 u[2];
   u[0] = make_float2(randu, randv);
   u[1].x = lcg_step_float(&sd->lcg_state);
   u[1].y = lcg_step_float(&sd->lcg_state);
 
-  const float sin_theta_o = local_O.x;
+  const float sin_theta_o = wo.x;
   const float cos_theta_o = cos_from_sin(sin_theta_o);
-  const float phi_o = atan2f(local_O.z, local_O.y);
+  const float phi_o = atan2f(wo.z, wo.y);
 
   const float sin_theta_t = sin_theta_o / bsdf->eta;
   const float cos_theta_t = cos_from_sin(sin_theta_t);
@@ -454,7 +458,7 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
   *eval = F;
   *pdf = F_energy;
 
-  *wo = X * sin_theta_i + Y * cos_theta_i * cosf(phi_i) + Z * cos_theta_i * sinf(phi_i);
+  *omega_in = X * sin_theta_i + Y * cos_theta_i * cosf(phi_i) + Z * cos_theta_i * sinf(phi_i);
 
   return LABEL_GLOSSY | ((p == 0) ? LABEL_REFLECT : LABEL_TRANSMIT);
 }

intern/cycles/kernel/closure/bsdf_microfacet_multi.h

@@ -43,7 +43,7 @@ ccl_device_forceinline float2 mf_sampleP22_11(const float cosI,
     return make_float2(r * cosf(phi), r * sinf(phi));
   }
 
-  const float sinI = sin_from_cos(cosI);
+  const float sinI = safe_sqrtf(1.0f - cosI * cosI);
   const float tanI = sinI / cosI;
   const float projA = 0.5f * (cosI + 1.0f);
   if (projA < 0.0001f)
@@ -401,7 +401,7 @@ ccl_device int bsdf_microfacet_multi_ggx_fresnel_setup(ccl_private MicrofacetBsd
 
   bsdf->type = CLOSURE_BSDF_MICROFACET_MULTI_GGX_FRESNEL_ID;
 
-  bsdf_microfacet_adjust_weight(sd, bsdf);
+  bsdf_microfacet_fresnel_color(sd, bsdf);
 
   return bsdf_microfacet_multi_ggx_common_setup(bsdf);
 }
@@ -417,15 +417,15 @@ ccl_device int bsdf_microfacet_multi_ggx_refraction_setup(ccl_private Microfacet
 
 ccl_device Spectrum bsdf_microfacet_multi_ggx_eval(ccl_private const ShaderClosure *sc,
                                                    const float3 Ng,
-                                                   const float3 wi,
-                                                   const float3 wo,
+                                                   const float3 I,
+                                                   const float3 omega_in,
                                                    ccl_private float *pdf,
                                                    ccl_private uint *lcg_state)
 {
   ccl_private const MicrofacetBsdf *bsdf = (ccl_private const MicrofacetBsdf *)sc;
-  const float cosNgO = dot(Ng, wo);
+  const float cosNgI = dot(Ng, omega_in);
 
-  if ((cosNgO < 0.0f) || bsdf->alpha_x * bsdf->alpha_y < 1e-7f) {
+  if ((cosNgI < 0.0f) || bsdf->alpha_x * bsdf->alpha_y < 1e-7f) {
     *pdf = 0.0f;
     return zero_spectrum();
   }
@@ -434,7 +434,7 @@ ccl_device Spectrum bsdf_microfacet_multi_ggx_eval(ccl_private const ShaderClosu
   Z = bsdf->N;
 
   /* Ensure that the both directions are on the outside w.r.t. the shading normal. */
-  if (dot(Z, wi) <= 0.0f || dot(Z, wo) <= 0.0f) {
+  if (dot(Z, I) <= 0.0f || dot(Z, omega_in) <= 0.0f) {
     *pdf = 0.0f;
     return zero_spectrum();
   }
@@ -447,21 +447,21 @@ ccl_device Spectrum bsdf_microfacet_multi_ggx_eval(ccl_private const ShaderClosu
   else
     make_orthonormals(Z, &X, &Y);
 
-  float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
-  float3 local_O = make_float3(dot(wo, X), dot(wo, Y), dot(wo, Z));
+  float3 localI = make_float3(dot(I, X), dot(I, Y), dot(I, Z));
+  float3 localO = make_float3(dot(omega_in, X), dot(omega_in, Y), dot(omega_in, Z));
 
   if (is_aniso)
-    *pdf = mf_ggx_aniso_pdf(local_I, local_O, make_float2(bsdf->alpha_x, bsdf->alpha_y));
+    *pdf = mf_ggx_aniso_pdf(localI, localO, make_float2(bsdf->alpha_x, bsdf->alpha_y));
   else
-    *pdf = mf_ggx_pdf(local_I, local_O, bsdf->alpha_x);
+    *pdf = mf_ggx_pdf(localI, localO, bsdf->alpha_x);
 
   if (*pdf <= 0.f) {
     *pdf = 0.f;
     return make_float3(0.f, 0.f, 0.f);
   }
 
-  return mf_eval_glossy(local_I,
-                        local_O,
+  return mf_eval_glossy(localI,
+                        localO,
                         true,
                         bsdf->extra->color,
                         bsdf->alpha_x,
@@ -475,11 +475,11 @@ ccl_device Spectrum bsdf_microfacet_multi_ggx_eval(ccl_private const ShaderClosu
 ccl_device int bsdf_microfacet_multi_ggx_sample(KernelGlobals kg,
                                                 ccl_private const ShaderClosure *sc,
                                                 float3 Ng,
-                                                float3 wi,
+                                                float3 I,
                                                 float randu,
                                                 float randv,
                                                 ccl_private Spectrum *eval,
-                                                ccl_private float3 *wo,
+                                                ccl_private float3 *omega_in,
                                                 ccl_private float *pdf,
                                                 ccl_private uint *lcg_state,
                                                 ccl_private float2 *sampled_roughness,
@@ -491,7 +491,7 @@ ccl_device int bsdf_microfacet_multi_ggx_sample(KernelGlobals kg,
   Z = bsdf->N;
 
   /* Ensure that the view direction is on the outside w.r.t. the shading normal. */
-  if (dot(Z, wi) <= 0.0f) {
+  if (dot(Z, I) <= 0.0f) {
     *pdf = 0.0f;
     return LABEL_NONE;
   }
@@ -499,8 +499,8 @@ ccl_device int bsdf_microfacet_multi_ggx_sample(KernelGlobals kg,
   /* Special case: Extremely low roughness.
    * Don't bother with microfacets, just do specular reflection. */
   if (bsdf->alpha_x * bsdf->alpha_y < 1e-7f) {
-    *wo = 2 * dot(Z, wi) * Z - wi;
-    if (dot(Ng, *wo) <= 0.0f) {
+    *omega_in = 2 * dot(Z, I) * Z - I;
+    if (dot(Ng, *omega_in) <= 0.0f) {
       *pdf = 0.0f;
       return LABEL_NONE;
     }
@@ -520,11 +520,11 @@ ccl_device int bsdf_microfacet_multi_ggx_sample(KernelGlobals kg,
   else
     make_orthonormals(Z, &X, &Y);
 
-  float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
-  float3 local_O;
+  float3 localI = make_float3(dot(I, X), dot(I, Y), dot(I, Z));
+  float3 localO;
 
-  *eval = mf_sample_glossy(local_I,
-                           &local_O,
+  *eval = mf_sample_glossy(localI,
+                           &localO,
                            bsdf->extra->color,
                            bsdf->alpha_x,
                            bsdf->alpha_y,
@@ -532,18 +532,18 @@ ccl_device int bsdf_microfacet_multi_ggx_sample(KernelGlobals kg,
                            bsdf->ior,
                            use_fresnel,
                            bsdf->extra->cspec0);
-  *wo = X * local_O.x + Y * local_O.y + Z * local_O.z;
+  *omega_in = X * localO.x + Y * localO.y + Z * localO.z;
 
   /* Ensure that the light direction is on the outside w.r.t. the geometry normal. */
-  if (dot(Ng, *wo) <= 0.0f) {
+  if (dot(Ng, *omega_in) <= 0.0f) {
     *pdf = 0.0f;
     return LABEL_NONE;
   }
 
   if (is_aniso)
-    *pdf = mf_ggx_aniso_pdf(local_I, local_O, make_float2(bsdf->alpha_x, bsdf->alpha_y));
+    *pdf = mf_ggx_aniso_pdf(localI, localO, make_float2(bsdf->alpha_x, bsdf->alpha_y));
   else
-    *pdf = mf_ggx_pdf(local_I, local_O, bsdf->alpha_x);
+    *pdf = mf_ggx_pdf(localI, localO, bsdf->alpha_x);
   *pdf = fmaxf(0.f, *pdf);
   *eval *= *pdf;
 
@@ -575,14 +575,14 @@ ccl_device int bsdf_microfacet_multi_ggx_glass_fresnel_setup(ccl_private Microfa
 
   bsdf->type = CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_FRESNEL_ID;
 
-  bsdf_microfacet_adjust_weight(sd, bsdf);
+  bsdf_microfacet_fresnel_color(sd, bsdf);
 
   return SD_BSDF | SD_BSDF_HAS_EVAL | SD_BSDF_NEEDS_LCG;
 }
 
 ccl_device Spectrum bsdf_microfacet_multi_ggx_glass_eval(ccl_private const ShaderClosure *sc,
-                                                         const float3 wi,
-                                                         const float3 wo,
+                                                         const float3 I,
+                                                         const float3 omega_in,
                                                          ccl_private float *pdf,
                                                          ccl_private uint *lcg_state)
 {
@@ -597,17 +597,17 @@ ccl_device Spectrum bsdf_microfacet_multi_ggx_glass_eval(ccl_private const Shade
   Z = bsdf->N;
   make_orthonormals(Z, &X, &Y);
 
-  float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
-  float3 local_O = make_float3(dot(wo, X), dot(wo, Y), dot(wo, Z));
+  float3 localI = make_float3(dot(I, X), dot(I, Y), dot(I, Z));
+  float3 localO = make_float3(dot(omega_in, X), dot(omega_in, Y), dot(omega_in, Z));
 
-  const bool is_transmission = local_O.z < 0.0f;
+  const bool is_transmission = localO.z < 0.0f;
   const bool use_fresnel = !is_transmission &&
                            (bsdf->type == CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_FRESNEL_ID);
 
-  *pdf = mf_glass_pdf(local_I, local_O, bsdf->alpha_x, bsdf->ior);
+  *pdf = mf_glass_pdf(localI, localO, bsdf->alpha_x, bsdf->ior);
   kernel_assert(*pdf >= 0.f);
-  return mf_eval_glass(local_I,
-                       local_O,
+  return mf_eval_glass(localI,
+                       localO,
                        !is_transmission,
                        bsdf->extra->color,
                        bsdf->alpha_x,
@@ -621,11 +621,11 @@ ccl_device Spectrum bsdf_microfacet_multi_ggx_glass_eval(ccl_private const Shade
 ccl_device int bsdf_microfacet_multi_ggx_glass_sample(KernelGlobals kg,
                                                       ccl_private const ShaderClosure *sc,
                                                       float3 Ng,
-                                                      float3 wi,
+                                                      float3 I,
                                                       float randu,
                                                       float randv,
                                                       ccl_private Spectrum *eval,
-                                                      ccl_private float3 *wo,
+                                                      ccl_private float3 *omega_in,
                                                       ccl_private float *pdf,
                                                       ccl_private uint *lcg_state,
                                                       ccl_private float2 *sampled_roughness,
@@ -642,16 +642,16 @@ ccl_device int bsdf_microfacet_multi_ggx_glass_sample(KernelGlobals kg,
   if (bsdf->alpha_x * bsdf->alpha_y < 1e-7f) {
     float3 R, T;
     bool inside;
-    float fresnel = fresnel_dielectric(bsdf->ior, Z, wi, &R, &T, &inside);
+    float fresnel = fresnel_dielectric(bsdf->ior, Z, I, &R, &T, &inside);
 
     *pdf = 1e6f;
     *eval = make_spectrum(1e6f);
     if (randu < fresnel) {
-      *wo = R;
+      *omega_in = R;
       return LABEL_REFLECT | LABEL_SINGULAR;
     }
     else {
-      *wo = T;
+      *omega_in = T;
       return LABEL_TRANSMIT | LABEL_SINGULAR;
     }
   }
@@ -660,11 +660,11 @@ ccl_device int bsdf_microfacet_multi_ggx_glass_sample(KernelGlobals kg,
 
   make_orthonormals(Z, &X, &Y);
 
-  float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
-  float3 local_O;
+  float3 localI = make_float3(dot(I, X), dot(I, Y), dot(I, Z));
+  float3 localO;
 
-  *eval = mf_sample_glass(local_I,
-                          &local_O,
+  *eval = mf_sample_glass(localI,
+                          &localO,
                           bsdf->extra->color,
                           bsdf->alpha_x,
                           bsdf->alpha_y,
@@ -672,12 +672,12 @@ ccl_device int bsdf_microfacet_multi_ggx_glass_sample(KernelGlobals kg,
                           bsdf->ior,
                           use_fresnel,
                           bsdf->extra->cspec0);
-  *pdf = mf_glass_pdf(local_I, local_O, bsdf->alpha_x, bsdf->ior);
+  *pdf = mf_glass_pdf(localI, localO, bsdf->alpha_x, bsdf->ior);
   kernel_assert(*pdf >= 0.f);
   *eval *= *pdf;
 
-  *wo = X * local_O.x + Y * local_O.y + Z * local_O.z;
-  if (local_O.z * local_I.z > 0.0f) {
+  *omega_in = X * localO.x + Y * localO.y + Z * localO.z;
+  if (localO.z * localI.z > 0.0f) {
     return LABEL_REFLECT | LABEL_GLOSSY;
   }
   else {

intern/cycles/kernel/closure/bsdf_microfacet_multi_impl.h

@@ -73,8 +73,9 @@ ccl_device_forceinline Spectrum MF_FUNCTION_FULL_NAME(mf_eval)(float3 wi,
   eval = make_spectrum(val);
 #endif
 
+  float F0 = fresnel_dielectric_cos(1.0f, eta);
   if (use_fresnel) {
-    throughput = interpolate_fresnel_color(wi, wh, eta, cspec0);
+    throughput = interpolate_fresnel_color(wi, wh, eta, F0, cspec0);
 
     eval *= throughput;
   }
@@ -143,11 +144,11 @@ ccl_device_forceinline Spectrum MF_FUNCTION_FULL_NAME(mf_eval)(float3 wi,
         throughput *= color;
       }
       else if (use_fresnel && order > 0) {
-        throughput *= interpolate_fresnel_color(wi_prev, wm, eta, cspec0);
+        throughput *= interpolate_fresnel_color(wi_prev, wm, eta, F0, cspec0);
       }
 #else /* MF_MULTI_GLOSSY */
       if (use_fresnel && order > 0) {
-        throughput *= interpolate_fresnel_color(-wr, wm, eta, cspec0);
+        throughput *= interpolate_fresnel_color(-wr, wm, eta, F0, cspec0);
       }
       wr = mf_sample_phase_glossy(-wr, &throughput, wm);
 #endif
@@ -191,6 +192,8 @@ ccl_device_forceinline Spectrum MF_FUNCTION_FULL_NAME(mf_sample)(float3 wi,
   float G1_r = 0.0f;
   bool outside = true;
 
+  float F0 = fresnel_dielectric_cos(1.0f, eta);
+
   int order;
   for (order = 0; order < 10; order++) {
     /* Sample microfacet height. */
@@ -226,12 +229,22 @@ ccl_device_forceinline Spectrum MF_FUNCTION_FULL_NAME(mf_sample)(float3 wi,
         throughput *= color;
       }
       else {
-        throughput *= interpolate_fresnel_color(wi_prev, wm, eta, cspec0);
+        Spectrum t_color = interpolate_fresnel_color(wi_prev, wm, eta, F0, cspec0);
+
+        if (order == 0)
+          throughput = t_color;
+        else
+          throughput *= t_color;
       }
     }
 #else /* MF_MULTI_GLOSSY */
     if (use_fresnel) {
-      throughput *= interpolate_fresnel_color(-wr, wm, eta, cspec0);
+      Spectrum t_color = interpolate_fresnel_color(-wr, wm, eta, F0, cspec0);
+
+      if (order == 0)
+        throughput = t_color;
+      else
+        throughput *= t_color;
     }
     wr = mf_sample_phase_glossy(-wr, &throughput, wm);
 #endif

intern/cycles/kernel/closure/bsdf_oren_nayar.h

@@ -48,14 +48,14 @@ ccl_device int bsdf_oren_nayar_setup(ccl_private OrenNayarBsdf *bsdf)
 }
 
 ccl_device Spectrum bsdf_oren_nayar_eval(ccl_private const ShaderClosure *sc,
-                                         const float3 wi,
-                                         const float3 wo,
+                                         const float3 I,
+                                         const float3 omega_in,
                                          ccl_private float *pdf)
 {
   ccl_private const OrenNayarBsdf *bsdf = (ccl_private const OrenNayarBsdf *)sc;
-  if (dot(bsdf->N, wo) > 0.0f) {
+  if (dot(bsdf->N, omega_in) > 0.0f) {
     *pdf = 0.5f * M_1_PI_F;
-    return bsdf_oren_nayar_get_intensity(sc, bsdf->N, wi, wo);
+    return bsdf_oren_nayar_get_intensity(sc, bsdf->N, I, omega_in);
   }
   else {
     *pdf = 0.0f;
@@ -65,18 +65,18 @@ ccl_device Spectrum bsdf_oren_nayar_eval(ccl_private const ShaderClosure *sc,
 
 ccl_device int bsdf_oren_nayar_sample(ccl_private const ShaderClosure *sc,
                                       float3 Ng,
-                                      float3 wi,
+                                      float3 I,
                                       float randu,
                                       float randv,
                                       ccl_private Spectrum *eval,
-                                      ccl_private float3 *wo,
+                                      ccl_private float3 *omega_in,
                                       ccl_private float *pdf)
 {
   ccl_private const OrenNayarBsdf *bsdf = (ccl_private const OrenNayarBsdf *)sc;
-  sample_uniform_hemisphere(bsdf->N, randu, randv, wo, pdf);
+  sample_uniform_hemisphere(bsdf->N, randu, randv, omega_in, pdf);
 
-  if (dot(Ng, *wo) > 0.0f) {
-    *eval = bsdf_oren_nayar_get_intensity(sc, bsdf->N, wi, *wo);
+  if (dot(Ng, *omega_in) > 0.0f) {
+    *eval = bsdf_oren_nayar_get_intensity(sc, bsdf->N, I, *omega_in);
   }
   else {
     *pdf = 0.0f;

intern/cycles/kernel/closure/bsdf_phong_ramp.h

@@ -45,23 +45,23 @@ ccl_device int bsdf_phong_ramp_setup(ccl_private PhongRampBsdf *bsdf)
 }
 
 ccl_device Spectrum bsdf_phong_ramp_eval(ccl_private const ShaderClosure *sc,
-                                         const float3 wi,
-                                         const float3 wo,
+                                         const float3 I,
+                                         const float3 omega_in,
                                          ccl_private float *pdf)
 {
   ccl_private const PhongRampBsdf *bsdf = (ccl_private const PhongRampBsdf *)sc;
   float m_exponent = bsdf->exponent;
-  float cosNI = dot(bsdf->N, wi);
-  float cosNO = dot(bsdf->N, wo);
+  float cosNI = dot(bsdf->N, omega_in);
+  float cosNO = dot(bsdf->N, I);
 
   if (cosNI > 0 && cosNO > 0) {
     // reflect the view vector
-    float3 R = (2 * cosNI) * bsdf->N - wi;
-    float cosRO = dot(R, wo);
-    if (cosRO > 0) {
-      float cosp = powf(cosRO, m_exponent);
+    float3 R = (2 * cosNO) * bsdf->N - I;
+    float cosRI = dot(R, omega_in);
+    if (cosRI > 0) {
+      float cosp = powf(cosRI, m_exponent);
       float common = 0.5f * M_1_PI_F * cosp;
-      float out = cosNO * (m_exponent + 2) * common;
+      float out = cosNI * (m_exponent + 2) * common;
       *pdf = (m_exponent + 1) * common;
       return rgb_to_spectrum(bsdf_phong_ramp_get_color(bsdf->colors, cosp) * out);
     }
@@ -77,39 +77,39 @@ ccl_device_inline float phong_ramp_exponent_to_roughness(float exponent)
 
 ccl_device int bsdf_phong_ramp_sample(ccl_private const ShaderClosure *sc,
                                       float3 Ng,
-                                      float3 wi,
+                                      float3 I,
                                       float randu,
                                       float randv,
                                       ccl_private Spectrum *eval,
-                                      ccl_private float3 *wo,
+                                      ccl_private float3 *omega_in,
                                       ccl_private float *pdf,
                                       ccl_private float2 *sampled_roughness)
 {
   ccl_private const PhongRampBsdf *bsdf = (ccl_private const PhongRampBsdf *)sc;
-  float cosNI = dot(bsdf->N, wi);
+  float cosNO = dot(bsdf->N, I);
   float m_exponent = bsdf->exponent;
   const float m_roughness = phong_ramp_exponent_to_roughness(m_exponent);
   *sampled_roughness = make_float2(m_roughness, m_roughness);
 
-  if (cosNI > 0) {
+  if (cosNO > 0) {
     // reflect the view vector
-    float3 R = (2 * cosNI) * bsdf->N - wi;
+    float3 R = (2 * cosNO) * bsdf->N - I;
     float3 T, B;
     make_orthonormals(R, &T, &B);
     float phi = M_2PI_F * randu;
     float cosTheta = powf(randv, 1 / (m_exponent + 1));
     float sinTheta2 = 1 - cosTheta * cosTheta;
     float sinTheta = sinTheta2 > 0 ? sqrtf(sinTheta2) : 0;
-    *wo = (cosf(phi) * sinTheta) * T + (sinf(phi) * sinTheta) * B + (cosTheta)*R;
-    if (dot(Ng, *wo) > 0.0f) {
+    *omega_in = (cosf(phi) * sinTheta) * T + (sinf(phi) * sinTheta) * B + (cosTheta)*R;
+    if (dot(Ng, *omega_in) > 0.0f) {
       // common terms for pdf and eval
-      float cosNO = dot(bsdf->N, *wo);
+      float cosNI = dot(bsdf->N, *omega_in);
       // make sure the direction we chose is still in the right hemisphere
-      if (cosNO > 0) {
+      if (cosNI > 0) {
         float cosp = powf(cosTheta, m_exponent);
         float common = 0.5f * M_1_PI_F * cosp;
         *pdf = (m_exponent + 1) * common;
-        float out = cosNO * (m_exponent + 2) * common;
+        float out = cosNI * (m_exponent + 2) * common;
         *eval = rgb_to_spectrum(bsdf_phong_ramp_get_color(bsdf->colors, cosp) * out);
       }
     }

intern/cycles/kernel/closure/bsdf_principled_diffuse.h

@@ -110,17 +110,17 @@ ccl_device int bsdf_principled_diffuse_setup(ccl_private PrincipledDiffuseBsdf *
 }
 
 ccl_device Spectrum bsdf_principled_diffuse_eval(ccl_private const ShaderClosure *sc,
-                                                 const float3 wi,
-                                                 const float3 wo,
+                                                 const float3 I,
+                                                 const float3 omega_in,
                                                  ccl_private float *pdf)
 {
   ccl_private const PrincipledDiffuseBsdf *bsdf = (ccl_private const PrincipledDiffuseBsdf *)sc;
   const float3 N = bsdf->N;
 
-  if (dot(N, wo) > 0.0f) {
-    const float3 V = wi;
-    const float3 L = wo;
-    *pdf = fmaxf(dot(N, wo), 0.0f) * M_1_PI_F;
+  if (dot(N, omega_in) > 0.0f) {
+    const float3 V = I;         // outgoing
+    const float3 L = omega_in;  // incoming
+    *pdf = fmaxf(dot(N, omega_in), 0.0f) * M_1_PI_F;
     return bsdf_principled_diffuse_compute_brdf(bsdf, N, V, L, pdf);
   }
   else {
@@ -131,21 +131,21 @@ ccl_device Spectrum bsdf_principled_diffuse_eval(ccl_private const ShaderClosure
 
 ccl_device int bsdf_principled_diffuse_sample(ccl_private const ShaderClosure *sc,
                                               float3 Ng,
-                                              float3 wi,
+                                              float3 I,
                                               float randu,
                                               float randv,
                                               ccl_private Spectrum *eval,
-                                              ccl_private float3 *wo,
+                                              ccl_private float3 *omega_in,
                                               ccl_private float *pdf)
 {
   ccl_private const PrincipledDiffuseBsdf *bsdf = (ccl_private const PrincipledDiffuseBsdf *)sc;
 
   float3 N = bsdf->N;
 
-  sample_cos_hemisphere(N, randu, randv, wo, pdf);
+  sample_cos_hemisphere(N, randu, randv, omega_in, pdf);
 
-  if (dot(Ng, *wo) > 0) {
-    *eval = bsdf_principled_diffuse_compute_brdf(bsdf, N, wi, *wo, pdf);
+  if (dot(Ng, *omega_in) > 0) {
+    *eval = bsdf_principled_diffuse_compute_brdf(bsdf, N, I, *omega_in, pdf);
   }
   else {
     *pdf = 0.0f;

intern/cycles/kernel/closure/bsdf_principled_sheen.h

@@ -54,25 +54,25 @@ ccl_device int bsdf_principled_sheen_setup(ccl_private const ShaderData *sd,
                                            ccl_private PrincipledSheenBsdf *bsdf)
 {
   bsdf->type = CLOSURE_BSDF_PRINCIPLED_SHEEN_ID;
-  bsdf->avg_value = calculate_avg_principled_sheen_brdf(bsdf->N, sd->wi);
+  bsdf->avg_value = calculate_avg_principled_sheen_brdf(bsdf->N, sd->I);
   bsdf->sample_weight *= bsdf->avg_value;
   return SD_BSDF | SD_BSDF_HAS_EVAL;
 }
 
 ccl_device Spectrum bsdf_principled_sheen_eval(ccl_private const ShaderClosure *sc,
-                                               const float3 wi,
-                                               const float3 wo,
+                                               const float3 I,
+                                               const float3 omega_in,
                                                ccl_private float *pdf)
 {
   ccl_private const PrincipledSheenBsdf *bsdf = (ccl_private const PrincipledSheenBsdf *)sc;
   const float3 N = bsdf->N;
 
-  if (dot(N, wo) > 0.0f) {
-    const float3 V = wi;
-    const float3 L = wo;
+  if (dot(N, omega_in) > 0.0f) {
+    const float3 V = I;         // outgoing
+    const float3 L = omega_in;  // incoming
     const float3 H = normalize(L + V);
 
-    *pdf = fmaxf(dot(N, wo), 0.0f) * M_1_PI_F;
+    *pdf = fmaxf(dot(N, omega_in), 0.0f) * M_1_PI_F;
     return calculate_principled_sheen_brdf(N, V, L, H, pdf);
   }
   else {
@@ -83,23 +83,23 @@ ccl_device Spectrum bsdf_principled_sheen_eval(ccl_private const ShaderClosure *
 
 ccl_device int bsdf_principled_sheen_sample(ccl_private const ShaderClosure *sc,
                                             float3 Ng,
-                                            float3 wi,
+                                            float3 I,
                                             float randu,
                                             float randv,
                                             ccl_private Spectrum *eval,
-                                            ccl_private float3 *wo,
+                                            ccl_private float3 *omega_in,
                                             ccl_private float *pdf)
 {
   ccl_private const PrincipledSheenBsdf *bsdf = (ccl_private const PrincipledSheenBsdf *)sc;
 
   float3 N = bsdf->N;
 
-  sample_cos_hemisphere(N, randu, randv, wo, pdf);
+  sample_cos_hemisphere(N, randu, randv, omega_in, pdf);
 
-  if (dot(Ng, *wo) > 0) {
-    float3 H = normalize(wi + *wo);
+  if (dot(Ng, *omega_in) > 0) {
+    float3 H = normalize(I + *omega_in);
 
-    *eval = calculate_principled_sheen_brdf(N, wi, *wo, H, pdf);
+    *eval = calculate_principled_sheen_brdf(N, I, *omega_in, H, pdf);
   }
   else {
     *eval = zero_spectrum();

intern/cycles/kernel/closure/bsdf_reflection.h

@@ -19,8 +19,8 @@ ccl_device int bsdf_reflection_setup(ccl_private MicrofacetBsdf *bsdf)
 }
 
 ccl_device Spectrum bsdf_reflection_eval(ccl_private const ShaderClosure *sc,
-                                         const float3 wi,
-                                         const float3 wo,
+                                         const float3 I,
+                                         const float3 omega_in,
                                          ccl_private float *pdf)
 {
   *pdf = 0.0f;
@@ -29,11 +29,11 @@ ccl_device Spectrum bsdf_reflection_eval(ccl_private const ShaderClosure *sc,
 
 ccl_device int bsdf_reflection_sample(ccl_private const ShaderClosure *sc,
                                       float3 Ng,
-                                      float3 wi,
+                                      float3 I,
                                       float randu,
                                       float randv,
                                       ccl_private Spectrum *eval,
-                                      ccl_private float3 *wo,
+                                      ccl_private float3 *omega_in,
                                       ccl_private float *pdf,
                                       ccl_private float *eta)
 {
@@ -42,10 +42,10 @@ ccl_device int bsdf_reflection_sample(ccl_private const ShaderClosure *sc,
   *eta = bsdf->ior;
 
   // only one direction is possible
-  float cosNI = dot(N, wi);
-  if (cosNI > 0) {
-    *wo = (2 * cosNI) * N - wi;
-    if (dot(Ng, *wo) > 0) {
+  float cosNO = dot(N, I);
+  if (cosNO > 0) {
+    *omega_in = (2 * cosNO) * N - I;
+    if (dot(Ng, *omega_in) > 0) {
       /* Some high number for MIS. */
       *pdf = 1e6f;
       *eval = make_spectrum(1e6f);

intern/cycles/kernel/closure/bsdf_refraction.h

@@ -19,8 +19,8 @@ ccl_device int bsdf_refraction_setup(ccl_private MicrofacetBsdf *bsdf)
 }
 
 ccl_device Spectrum bsdf_refraction_eval(ccl_private const ShaderClosure *sc,
-                                         const float3 wi,
-                                         const float3 wo,
+                                         const float3 I,
+                                         const float3 omega_in,
                                          ccl_private float *pdf)
 {
   *pdf = 0.0f;
@@ -29,11 +29,11 @@ ccl_device Spectrum bsdf_refraction_eval(ccl_private const ShaderClosure *sc,
 
 ccl_device int bsdf_refraction_sample(ccl_private const ShaderClosure *sc,
                                       float3 Ng,
-                                      float3 wi,
+                                      float3 I,
                                       float randu,
                                       float randv,
                                       ccl_private Spectrum *eval,
-                                      ccl_private float3 *wo,
+                                      ccl_private float3 *omega_in,
                                       ccl_private float *pdf,
                                       ccl_private float *eta)
 {
@@ -46,13 +46,13 @@ ccl_device int bsdf_refraction_sample(ccl_private const ShaderClosure *sc,
   float3 R, T;
   bool inside;
   float fresnel;
-  fresnel = fresnel_dielectric(m_eta, N, wi, &R, &T, &inside);
+  fresnel = fresnel_dielectric(m_eta, N, I, &R, &T, &inside);
 
   if (!inside && fresnel != 1.0f) {
     /* Some high number for MIS. */
     *pdf = 1e6f;
     *eval = make_spectrum(1e6f);
-    *wo = T;
+    *omega_in = T;
   }
   else {
     *pdf = 0.0f;

intern/cycles/kernel/closure/bsdf_toon.h

@@ -50,17 +50,17 @@ ccl_device float bsdf_toon_get_sample_angle(float max_angle, float smooth)
 }
 
 ccl_device Spectrum bsdf_diffuse_toon_eval(ccl_private const ShaderClosure *sc,
-                                           const float3 wi,
-                                           const float3 wo,
+                                           const float3 I,
+                                           const float3 omega_in,
                                            ccl_private float *pdf)
 {
   ccl_private const ToonBsdf *bsdf = (ccl_private const ToonBsdf *)sc;
-  float cosNO = dot(bsdf->N, wo);
+  float cosNI = dot(bsdf->N, omega_in);
 
-  if (cosNO >= 0.0f) {
+  if (cosNI >= 0.0f) {
     float max_angle = bsdf->size * M_PI_2_F;
     float smooth = bsdf->smooth * M_PI_2_F;
-    float angle = safe_acosf(fmaxf(cosNO, 0.0f));
+    float angle = safe_acosf(fmaxf(cosNI, 0.0f));
 
     float eval = bsdf_toon_get_intensity(max_angle, smooth, angle);
 
@@ -78,11 +78,11 @@ ccl_device Spectrum bsdf_diffuse_toon_eval(ccl_private const ShaderClosure *sc,
 
 ccl_device int bsdf_diffuse_toon_sample(ccl_private const ShaderClosure *sc,
                                         float3 Ng,
-                                        float3 wi,
+                                        float3 I,
                                         float randu,
                                         float randv,
                                         ccl_private Spectrum *eval,
-                                        ccl_private float3 *wo,
+                                        ccl_private float3 *omega_in,
                                         ccl_private float *pdf)
 {
   ccl_private const ToonBsdf *bsdf = (ccl_private const ToonBsdf *)sc;
@@ -92,9 +92,9 @@ ccl_device int bsdf_diffuse_toon_sample(ccl_private const ShaderClosure *sc,
   float angle = sample_angle * randu;
 
   if (sample_angle > 0.0f) {
-    sample_uniform_cone(bsdf->N, sample_angle, randu, randv, wo, pdf);
+    sample_uniform_cone(bsdf->N, sample_angle, randu, randv, omega_in, pdf);
 
-    if (dot(Ng, *wo) > 0.0f) {
+    if (dot(Ng, *omega_in) > 0.0f) {
       *eval = make_spectrum(*pdf * bsdf_toon_get_intensity(max_angle, smooth, angle));
     }
     else {
@@ -122,22 +122,22 @@ ccl_device int bsdf_glossy_toon_setup(ccl_private ToonBsdf *bsdf)
 }
 
 ccl_device Spectrum bsdf_glossy_toon_eval(ccl_private const ShaderClosure *sc,
-                                          const float3 wi,
-                                          const float3 wo,
+                                          const float3 I,
+                                          const float3 omega_in,
                                           ccl_private float *pdf)
 {
   ccl_private const ToonBsdf *bsdf = (ccl_private const ToonBsdf *)sc;
   float max_angle = bsdf->size * M_PI_2_F;
   float smooth = bsdf->smooth * M_PI_2_F;
-  float cosNI = dot(bsdf->N, wi);
-  float cosNO = dot(bsdf->N, wo);
+  float cosNI = dot(bsdf->N, omega_in);
+  float cosNO = dot(bsdf->N, I);
 
   if (cosNI > 0 && cosNO > 0) {
     /* reflect the view vector */
-    float3 R = (2 * cosNI) * bsdf->N - wi;
-    float cosRO = dot(R, wo);
+    float3 R = (2 * cosNO) * bsdf->N - I;
+    float cosRI = dot(R, omega_in);
 
-    float angle = safe_acosf(fmaxf(cosRO, 0.0f));
+    float angle = safe_acosf(fmaxf(cosRI, 0.0f));
 
     float eval = bsdf_toon_get_intensity(max_angle, smooth, angle);
     float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
@@ -151,32 +151,32 @@ ccl_device Spectrum bsdf_glossy_toon_eval(ccl_private const ShaderClosure *sc,
 
 ccl_device int bsdf_glossy_toon_sample(ccl_private const ShaderClosure *sc,
                                        float3 Ng,
-                                       float3 wi,
+                                       float3 I,
                                        float randu,
                                        float randv,
                                        ccl_private Spectrum *eval,
-                                       ccl_private float3 *wo,
+                                       ccl_private float3 *omega_in,
                                        ccl_private float *pdf)
 {
   ccl_private const ToonBsdf *bsdf = (ccl_private const ToonBsdf *)sc;
   float max_angle = bsdf->size * M_PI_2_F;
   float smooth = bsdf->smooth * M_PI_2_F;
-  float cosNI = dot(bsdf->N, wi);
+  float cosNO = dot(bsdf->N, I);
 
-  if (cosNI > 0) {
+  if (cosNO > 0) {
     /* reflect the view vector */
-    float3 R = (2 * cosNI) * bsdf->N - wi;
+    float3 R = (2 * cosNO) * bsdf->N - I;
 
     float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
     float angle = sample_angle * randu;
 
-    sample_uniform_cone(R, sample_angle, randu, randv, wo, pdf);
+    sample_uniform_cone(R, sample_angle, randu, randv, omega_in, pdf);
 
-    if (dot(Ng, *wo) > 0.0f) {
-      float cosNO = dot(bsdf->N, *wo);
+    if (dot(Ng, *omega_in) > 0.0f) {
+      float cosNI = dot(bsdf->N, *omega_in);
 
       /* make sure the direction we chose is still in the right hemisphere */
-      if (cosNO > 0) {
+      if (cosNI > 0) {
         *eval = make_spectrum(*pdf * bsdf_toon_get_intensity(max_angle, smooth, angle));
       }
       else {

intern/cycles/kernel/closure/bsdf_transparent.h

@@ -60,8 +60,8 @@ ccl_device void bsdf_transparent_setup(ccl_private ShaderData *sd,
 }
 
 ccl_device Spectrum bsdf_transparent_eval(ccl_private const ShaderClosure *sc,
-                                          const float3 wi,
-                                          const float3 wo,
+                                          const float3 I,
+                                          const float3 omega_in,
                                           ccl_private float *pdf)
 {
   *pdf = 0.0f;
@@ -70,15 +70,15 @@ ccl_device Spectrum bsdf_transparent_eval(ccl_private const ShaderClosure *sc,
 
 ccl_device int bsdf_transparent_sample(ccl_private const ShaderClosure *sc,
                                        float3 Ng,
-                                       float3 wi,
+                                       float3 I,
                                        float randu,
                                        float randv,
                                        ccl_private Spectrum *eval,
-                                       ccl_private float3 *wo,
+                                       ccl_private float3 *omega_in,
                                        ccl_private float *pdf)
 {
   // only one direction is possible
-  *wo = -wi;
+  *omega_in = -I;
   *pdf = 1;
   *eval = one_spectrum();
   return LABEL_TRANSMIT | LABEL_TRANSPARENT;

intern/cycles/kernel/closure/bsdf_util.h

@@ -89,21 +89,19 @@ ccl_device float schlick_fresnel(float u)
   return m2 * m2 * m;  // pow(m, 5)
 }
 
-/* Calculate the fresnel color, which is a blend between white and the F0 color */
-ccl_device_forceinline Spectrum interpolate_fresnel_color(float3 L,
-                                                          float3 H,
-                                                          float ior,
-                                                          Spectrum F0)
+/* Calculate the fresnel color which is a blend between white and the F0 color (cspec0) */
+ccl_device_forceinline Spectrum
+interpolate_fresnel_color(float3 L, float3 H, float ior, float F0, Spectrum cspec0)
 {
-  /* Compute the real Fresnel term and remap it from real_F0..1 to F0..1.
-   * The reason why we use this remapping instead of directly doing the
-   * Schlick approximation lerp(F0, 1.0, (1.0-cosLH)^5) is that for cases
-   * with similar IORs (e.g. ice in water), the relative IOR can be close
-   * enough to 1.0 that the Schlick approximation becomes inaccurate. */
-  float real_F = fresnel_dielectric_cos(dot(L, H), ior);
-  float real_F0 = fresnel_dielectric_cos(1.0f, ior);
+  /* Calculate the fresnel interpolation factor
+   * The value from fresnel_dielectric_cos(...) has to be normalized because
+   * the cspec0 keeps the F0 color
+   */
+  float F0_norm = 1.0f / (1.0f - F0);
+  float FH = (fresnel_dielectric_cos(dot(L, H), ior) - F0) * F0_norm;
 
-  return mix(F0, one_spectrum(), inverse_lerp(real_F0, 1.0f, real_F));
+  /* Blend between white and a specular color with respect to the fresnel */
+  return cspec0 * (1.0f - FH) + make_spectrum(FH);
 }
 
 ccl_device float3 ensure_valid_reflection(float3 Ng, float3 I, float3 N)

intern/cycles/kernel/closure/bssrdf.h

@@ -293,7 +293,7 @@ ccl_device int bssrdf_setup(ccl_private ShaderData *sd,
 
       /* Ad-hoc weight adjustment to avoid retro-reflection taking away half the
        * samples from BSSRDF. */
-      bsdf->sample_weight *= bsdf_principled_diffuse_retro_reflection_sample_weight(bsdf, sd->wi);
+      bsdf->sample_weight *= bsdf_principled_diffuse_retro_reflection_sample_weight(bsdf, sd->I);
     }
   }
 

intern/cycles/kernel/closure/emissive.h

@@ -36,24 +36,27 @@ ccl_device void emission_setup(ccl_private ShaderData *sd, const Spectrum weight
   }
 }
 
-/* return the probability distribution function in the direction wi,
+/* return the probability distribution function in the direction I,
  * given the parameters and the light's surface normal.  This MUST match
  * the PDF computed by sample(). */
-ccl_device float emissive_pdf(const float3 Ng, const float3 wi)
+ccl_device float emissive_pdf(const float3 Ng, const float3 I)
 {
-  float cosNI = fabsf(dot(Ng, wi));
-  return (cosNI > 0.0f) ? 1.0f : 0.0f;
+  float cosNO = fabsf(dot(Ng, I));
+  return (cosNO > 0.0f) ? 1.0f : 0.0f;
 }
 
-ccl_device void emissive_sample(
-    const float3 Ng, float randu, float randv, ccl_private float3 *wi, ccl_private float *pdf)
+ccl_device void emissive_sample(const float3 Ng,
+                                float randu,
+                                float randv,
+                                ccl_private float3 *omega_out,
+                                ccl_private float *pdf)
 {
   /* todo: not implemented and used yet */
 }
 
-ccl_device Spectrum emissive_simple_eval(const float3 Ng, const float3 wi)
+ccl_device Spectrum emissive_simple_eval(const float3 Ng, const float3 I)
 {
-  float res = emissive_pdf(Ng, wi);
+  float res = emissive_pdf(Ng, I);
 
   return make_spectrum(res);
 }

intern/cycles/kernel/closure/volume.h

@@ -49,18 +49,18 @@ ccl_device int volume_henyey_greenstein_setup(ccl_private HenyeyGreensteinVolume
 }
 
 ccl_device Spectrum volume_henyey_greenstein_eval_phase(ccl_private const ShaderVolumeClosure *svc,
-                                                        const float3 wi,
-                                                        float3 wo,
+                                                        const float3 I,
+                                                        float3 omega_in,
                                                         ccl_private float *pdf)
 {
   float g = svc->g;
 
-  /* note that wi points towards the viewer */
+  /* note that I points towards the viewer */
   if (fabsf(g) < 1e-3f) {
     *pdf = M_1_PI_F * 0.25f;
   }
   else {
-    float cos_theta = dot(-wi, wo);
+    float cos_theta = dot(-I, omega_in);
     *pdf = single_peaked_henyey_greenstein(cos_theta, g);
   }
 
@@ -88,7 +88,7 @@ henyey_greenstrein_sample(float3 D, float g, float randu, float randv, ccl_priva
     }
   }
 
-  float sin_theta = sin_from_cos(cos_theta);
+  float sin_theta = safe_sqrtf(1.0f - cos_theta * cos_theta);
   float phi = M_2PI_F * randv;
   float3 dir = make_float3(sin_theta * cosf(phi), sin_theta * sinf(phi), cos_theta);
 
@@ -100,17 +100,17 @@ henyey_greenstrein_sample(float3 D, float g, float randu, float randv, ccl_priva
 }
 
 ccl_device int volume_henyey_greenstein_sample(ccl_private const ShaderVolumeClosure *svc,
-                                               float3 wi,
+                                               float3 I,
                                                float randu,
                                                float randv,
                                                ccl_private Spectrum *eval,
-                                               ccl_private float3 *wo,
+                                               ccl_private float3 *omega_in,
                                                ccl_private float *pdf)
 {
   float g = svc->g;
 
-  /* note that wi points towards the viewer and so is used negated */
-  *wo = henyey_greenstrein_sample(-wi, g, randu, randv, pdf);
+  /* note that I points towards the viewer and so is used negated */
+  *omega_in = henyey_greenstrein_sample(-I, g, randu, randv, pdf);
   *eval = make_spectrum(*pdf); /* perfect importance sampling */
 
   return LABEL_VOLUME_SCATTER;
@@ -120,10 +120,10 @@ ccl_device int volume_henyey_greenstein_sample(ccl_private const ShaderVolumeClo
 
 ccl_device Spectrum volume_phase_eval(ccl_private const ShaderData *sd,
                                       ccl_private const ShaderVolumeClosure *svc,
-                                      float3 wo,
+                                      float3 omega_in,
                                       ccl_private float *pdf)
 {
-  return volume_henyey_greenstein_eval_phase(svc, sd->wi, wo, pdf);
+  return volume_henyey_greenstein_eval_phase(svc, sd->I, omega_in, pdf);
 }
 
 ccl_device int volume_phase_sample(ccl_private const ShaderData *sd,
@@ -131,10 +131,10 @@ ccl_device int volume_phase_sample(ccl_private const ShaderData *sd,
                                    float randu,
                                    float randv,
                                    ccl_private Spectrum *eval,
-                                   ccl_private float3 *wo,
+                                   ccl_private float3 *omega_in,
                                    ccl_private float *pdf)
 {
-  return volume_henyey_greenstein_sample(svc, sd->wi, randu, randv, eval, wo, pdf);
+  return volume_henyey_greenstein_sample(svc, sd->I, randu, randv, eval, omega_in, pdf);
 }
 
 /* Volume sampling utilities. */

intern/cycles/kernel/data_template.h

@@ -10,9 +10,6 @@
 #ifndef KERNEL_STRUCT_MEMBER
 #  define KERNEL_STRUCT_MEMBER(parent, type, name)
 #endif
-#ifndef KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE
-#  define KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE
-#endif
 
 /* Background. */
 
@@ -182,12 +179,9 @@ KERNEL_STRUCT_MEMBER(integrator, float, sample_clamp_indirect)
 KERNEL_STRUCT_MEMBER(integrator, int, use_caustics)
 /* Sampling pattern. */
 KERNEL_STRUCT_MEMBER(integrator, int, sampling_pattern)
-KERNEL_STRUCT_MEMBER(integrator, float, scrambling_distance)
-/* Sobol pattern. */
-KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE
 KERNEL_STRUCT_MEMBER(integrator, int, tabulated_sobol_sequence_size)
-KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE
 KERNEL_STRUCT_MEMBER(integrator, int, sobol_index_mask)
+KERNEL_STRUCT_MEMBER(integrator, float, scrambling_distance)
 /* Volume render. */
 KERNEL_STRUCT_MEMBER(integrator, int, use_volumes)
 KERNEL_STRUCT_MEMBER(integrator, int, volume_max_steps)
@@ -222,5 +216,4 @@ KERNEL_STRUCT_END(KernelSVMUsage)
 
 #undef KERNEL_STRUCT_BEGIN
 #undef KERNEL_STRUCT_MEMBER
-#undef KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE
 #undef KERNEL_STRUCT_END

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

@@ -401,71 +401,6 @@ ccl_gpu_kernel_threads(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
 }
 ccl_gpu_kernel_postfix
 
-ccl_gpu_kernel_threads(GPU_PARALLEL_SORT_BLOCK_SIZE)
-    ccl_gpu_kernel_signature(integrator_sort_bucket_pass,
-                             int num_states,
-                             int partition_size,
-                             int num_states_limit,
-                             ccl_global int *indices,
-                             int kernel_index)
-{
-#if defined(__KERNEL_LOCAL_ATOMIC_SORT__)
-  int max_shaders = context.launch_params_metal.data.max_shaders;
-  ccl_global ushort *d_queued_kernel = (ccl_global ushort *)
-                                           kernel_integrator_state.path.queued_kernel;
-  ccl_global uint *d_shader_sort_key = (ccl_global uint *)
-                                           kernel_integrator_state.path.shader_sort_key;
-  ccl_global int *key_offsets = (ccl_global int *)
-                                    kernel_integrator_state.sort_partition_key_offsets;
-
-  gpu_parallel_sort_bucket_pass(num_states,
-                                partition_size,
-                                max_shaders,
-                                kernel_index,
-                                d_queued_kernel,
-                                d_shader_sort_key,
-                                key_offsets,
-                                (threadgroup int *)threadgroup_array,
-                                metal_local_id,
-                                metal_local_size,
-                                metal_grid_id);
-#endif
-}
-
-ccl_gpu_kernel_threads(GPU_PARALLEL_SORT_BLOCK_SIZE)
-    ccl_gpu_kernel_signature(integrator_sort_write_pass,
-                             int num_states,
-                             int partition_size,
-                             int num_states_limit,
-                             ccl_global int *indices,
-                             int kernel_index)
-{
-#if defined(__KERNEL_LOCAL_ATOMIC_SORT__)
-  int max_shaders = context.launch_params_metal.data.max_shaders;
-  ccl_global ushort *d_queued_kernel = (ccl_global ushort *)
-                                           kernel_integrator_state.path.queued_kernel;
-  ccl_global uint *d_shader_sort_key = (ccl_global uint *)
-                                           kernel_integrator_state.path.shader_sort_key;
-  ccl_global int *key_offsets = (ccl_global int *)
-                                    kernel_integrator_state.sort_partition_key_offsets;
-
-  gpu_parallel_sort_write_pass(num_states,
-                               partition_size,
-                               max_shaders,
-                               kernel_index,
-                               num_states_limit,
-                               indices,
-                               d_queued_kernel,
-                               d_shader_sort_key,
-                               key_offsets,
-                               (threadgroup int *)threadgroup_array,
-                               metal_local_id,
-                               metal_local_size,
-                               metal_grid_id);
-#endif
-}
-ccl_gpu_kernel_postfix
-
 ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
     ccl_gpu_kernel_signature(integrator_compact_paths_array,
                              int num_states,

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

@@ -178,7 +178,7 @@ __device__
                                          simd_lane_index, \
                                          simd_group_index, \
                                          num_simd_groups, \
-                                         (threadgroup int *)threadgroup_array)
+                                         simdgroup_offset)
 #elif defined(__KERNEL_ONEAPI__)
 
 #  define gpu_parallel_active_index_array(num_states, indices, num_indices, is_active_op) \

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

@@ -19,115 +19,6 @@ CCL_NAMESPACE_BEGIN
 #  define GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE 512
 #endif
 #define GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY (~0)
-#define GPU_PARALLEL_SORT_BLOCK_SIZE 1024
-
-#if defined(__KERNEL_LOCAL_ATOMIC_SORT__)
-
-#  define atomic_store_local(p, x) \
-    atomic_store_explicit((threadgroup atomic_int *)p, x, memory_order_relaxed)
-#  define atomic_load_local(p) \
-    atomic_load_explicit((threadgroup atomic_int *)p, memory_order_relaxed)
-
-ccl_device_inline void gpu_parallel_sort_bucket_pass(const uint num_states,
-                                                     const uint partition_size,
-                                                     const uint max_shaders,
-                                                     const uint queued_kernel,
-                                                     ccl_global ushort *d_queued_kernel,
-                                                     ccl_global uint *d_shader_sort_key,
-                                                     ccl_global int *partition_key_offsets,
-                                                     ccl_gpu_shared int *buckets,
-                                                     const ushort local_id,
-                                                     const ushort local_size,
-                                                     const ushort grid_id)
-{
-  /* Zero the bucket sizes. */
-  if (local_id < max_shaders) {
-    atomic_store_local(&buckets[local_id], 0);
-  }
-
-  ccl_gpu_syncthreads();
-
-  /* Determine bucket sizes within the partitions. */
-
-  const uint partition_start = partition_size * uint(grid_id);
-  const uint partition_end = min(num_states, partition_start + partition_size);
-
-  for (int state_index = partition_start + uint(local_id); state_index < partition_end;
-       state_index += uint(local_size)) {
-    ushort kernel_index = d_queued_kernel[state_index];
-    if (kernel_index == queued_kernel) {
-      uint key = d_shader_sort_key[state_index] % max_shaders;
-      atomic_fetch_and_add_uint32(&buckets[key], 1);
-    }
-  }
-
-  ccl_gpu_syncthreads();
-
-  /* Calculate the partition's local offsets from the prefix sum of bucket sizes. */
-
-  if (local_id == 0) {
-    int offset = 0;
-    for (int i = 0; i < max_shaders; i++) {
-      partition_key_offsets[i + uint(grid_id) * (max_shaders + 1)] = offset;
-      offset = offset + atomic_load_local(&buckets[i]);
-    }
-
-    /* Store the number of active states in this partition. */
-    partition_key_offsets[max_shaders + uint(grid_id) * (max_shaders + 1)] = offset;
-  }
-}
-
-ccl_device_inline void gpu_parallel_sort_write_pass(const uint num_states,
-                                                    const uint partition_size,
-                                                    const uint max_shaders,
-                                                    const uint queued_kernel,
-                                                    const int num_states_limit,
-                                                    ccl_global int *indices,
-                                                    ccl_global ushort *d_queued_kernel,
-                                                    ccl_global uint *d_shader_sort_key,
-                                                    ccl_global int *partition_key_offsets,
-                                                    ccl_gpu_shared int *local_offset,
-                                                    const ushort local_id,
-                                                    const ushort local_size,
-                                                    const ushort grid_id)
-{
-  /* Calculate each partition's global offset from the prefix sum of the active state counts per
-   * partition. */
-
-  if (local_id < max_shaders) {
-    int partition_offset = 0;
-    for (int i = 0; i < uint(grid_id); i++) {
-      int partition_key_count = partition_key_offsets[max_shaders + uint(i) * (max_shaders + 1)];
-      partition_offset += partition_key_count;
-    }
-
-    ccl_global int *key_offsets = partition_key_offsets + (uint(grid_id) * (max_shaders + 1));
-    atomic_store_local(&local_offset[local_id], key_offsets[local_id] + partition_offset);
-  }
-
-  ccl_gpu_syncthreads();
-
-  /* Write the sorted active indices. */
-
-  const uint partition_start = partition_size * uint(grid_id);
-  const uint partition_end = min(num_states, partition_start + partition_size);
-
-  ccl_global int *key_offsets = partition_key_offsets + (uint(grid_id) * max_shaders);
-
-  for (int state_index = partition_start + uint(local_id); state_index < partition_end;
-       state_index += uint(local_size)) {
-    ushort kernel_index = d_queued_kernel[state_index];
-    if (kernel_index == queued_kernel) {
-      uint key = d_shader_sort_key[state_index] % max_shaders;
-      int index = atomic_fetch_and_add_uint32(&local_offset[key], 1);
-      if (index < num_states_limit) {
-        indices[index] = state_index;
-      }
-    }
-  }
-}
-
-#endif /* __KERNEL_LOCAL_ATOMIC_SORT__ */
 
 template<typename GetKeyOp>
 __device__ void gpu_parallel_sorted_index_array(const uint state_index,

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

@@ -105,11 +105,10 @@ struct kernel_gpu_##name \
 { \
   PARAMS_MAKER(__VA_ARGS__)(__VA_ARGS__) \
   void run(thread MetalKernelContext& context, \
-           threadgroup atomic_int *threadgroup_array, \
+           threadgroup int *simdgroup_offset, \
            const uint metal_global_id, \
            const ushort metal_local_id, \
            const ushort metal_local_size, \
-           const ushort metal_grid_id, \
            uint simdgroup_size, \
            uint simd_lane_index, \
            uint simd_group_index, \
@@ -118,24 +117,22 @@ struct kernel_gpu_##name \
 kernel void cycles_metal_##name(device const kernel_gpu_##name *params_struct, \
                                 constant KernelParamsMetal &ccl_restrict   _launch_params_metal, \
                                 constant MetalAncillaries *_metal_ancillaries, \
-                                threadgroup atomic_int *threadgroup_array[[ threadgroup(0) ]], \
+                                threadgroup int *simdgroup_offset[[ threadgroup(0) ]], \
                                 const uint metal_global_id [[thread_position_in_grid]], \
                                 const ushort metal_local_id   [[thread_position_in_threadgroup]], \
                                 const ushort metal_local_size [[threads_per_threadgroup]], \
-                                const ushort metal_grid_id    [[threadgroup_position_in_grid]], \
                                 uint simdgroup_size [[threads_per_simdgroup]], \
                                 uint simd_lane_index [[thread_index_in_simdgroup]], \
                                 uint simd_group_index [[simdgroup_index_in_threadgroup]], \
                                 uint num_simd_groups [[simdgroups_per_threadgroup]]) { \
   MetalKernelContext context(_launch_params_metal, _metal_ancillaries); \
-  params_struct->run(context, threadgroup_array, metal_global_id, metal_local_id, metal_local_size, metal_grid_id, simdgroup_size, simd_lane_index, simd_group_index, num_simd_groups); \
+  params_struct->run(context, simdgroup_offset, metal_global_id, metal_local_id, metal_local_size, simdgroup_size, simd_lane_index, simd_group_index, num_simd_groups); \
 } \
 void kernel_gpu_##name::run(thread MetalKernelContext& context, \
-                  threadgroup atomic_int *threadgroup_array, \
+                  threadgroup int *simdgroup_offset, \
                   const uint metal_global_id, \
                   const ushort metal_local_id, \
                   const ushort metal_local_size, \
-                  const ushort metal_grid_id, \
                   uint simdgroup_size, \
                   uint simd_lane_index, \
                   uint simd_group_index, \

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

@@ -34,7 +34,7 @@ class MetalKernelContext {
       kernel_assert(0);
       return 0;
     }
-
+    
 #ifdef __KERNEL_METAL_INTEL__
     template<typename TextureType, typename CoordsType>
     inline __attribute__((__always_inline__))
@@ -55,7 +55,7 @@ class MetalKernelContext {
       }
     }
 #endif
-
+    
     // texture2d
     template<>
     inline __attribute__((__always_inline__))

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

@@ -195,15 +195,7 @@ using sycl::half;
 #define fmodf(x, y) sycl::fmod((x), (y))
 #define lgammaf(x) sycl::lgamma((x))
 
-/* `sycl::native::cos` precision is not sufficient and `-ffast-math` lets
- * the current DPC++ compiler overload `sycl::cos` with it.
- * We work around this issue by directly calling the SPIRV implementation which
- * provides greater precision. */
-#if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__)
-#  define cosf(x) __spirv_ocl_cos(((float)(x)))
-#else
-#  define cosf(x) sycl::cos(((float)(x)))
-#endif
+#define cosf(x) sycl::native::cos(((float)(x)))
 #define sinf(x) sycl::native::sin(((float)(x)))
 #define powf(x, y) sycl::native::powr(((float)(x)), ((float)(y)))
 #define tanf(x) sycl::native::tan(((float)(x)))

intern/cycles/kernel/film/denoising_passes.h

@@ -58,7 +58,23 @@ ccl_device_forceinline void film_write_denoising_features_surface(KernelGlobals
     normal += sc->N * sc->sample_weight;
     sum_weight += sc->sample_weight;
 
-    if (sc->type == CLOSURE_BSDF_PRINCIPLED_DIFFUSE_ID) {
+    Spectrum closure_albedo = sc->weight;
+    /* Closures that include a Fresnel term typically have weights close to 1 even though their
+     * actual contribution is significantly lower.
+     * To account for this, we scale their weight by the average fresnel factor (the same is also
+     * done for the sample weight in the BSDF setup, so we don't need to scale that here). */
+    if (CLOSURE_IS_BSDF_MICROFACET_FRESNEL(sc->type)) {
+      ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)sc;
+      closure_albedo *= bsdf->extra->fresnel_color;
+    }
+    else if (sc->type == CLOSURE_BSDF_PRINCIPLED_SHEEN_ID) {
+      ccl_private PrincipledSheenBsdf *bsdf = (ccl_private PrincipledSheenBsdf *)sc;
+      closure_albedo *= bsdf->avg_value;
+    }
+    else if (sc->type == CLOSURE_BSDF_HAIR_PRINCIPLED_ID) {
+      closure_albedo *= bsdf_principled_hair_albedo(sc);
+    }
+    else if (sc->type == CLOSURE_BSDF_PRINCIPLED_DIFFUSE_ID) {
       /* BSSRDF already accounts for weight, retro-reflection would double up. */
       ccl_private const PrincipledDiffuseBsdf *bsdf = (ccl_private const PrincipledDiffuseBsdf *)
           sc;
@@ -67,7 +83,6 @@ ccl_device_forceinline void film_write_denoising_features_surface(KernelGlobals
       }
     }
 
-    Spectrum closure_albedo = bsdf_albedo(sd, sc);
     if (bsdf_get_specular_roughness_squared(sc) > sqr(0.075f)) {
       diffuse_albedo += closure_albedo;
       sum_nonspecular_weight += sc->sample_weight;

intern/cycles/kernel/geom/curve.h

@@ -252,7 +252,7 @@ ccl_device float3 curve_tangent_normal(KernelGlobals kg, ccl_private const Shade
 
   if (sd->type & PRIMITIVE_CURVE) {
 
-    tgN = -(-sd->wi - sd->dPdu * (dot(sd->dPdu, -sd->wi) / len_squared(sd->dPdu)));
+    tgN = -(-sd->I - sd->dPdu * (dot(sd->dPdu, -sd->I) / len_squared(sd->dPdu)));
     tgN = normalize(tgN);
 
     /* need to find suitable scaled gd for corrected normal */

intern/cycles/kernel/geom/curve_intersect.h

@@ -720,7 +720,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg,
     const float3 tangent = normalize(dPdu);
     const float3 bitangent = normalize(cross(tangent, -D));
     const float sine = sd->v;
-    const float cosine = cos_from_sin(sine);
+    const float cosine = safe_sqrtf(1.0f - sine * sine);
 
     sd->N = normalize(sine * bitangent - cosine * normalize(cross(tangent, bitangent)));
 #  if 0
@@ -738,7 +738,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg,
     /* NOTE: It is possible that P will be the same as P_inside (precision issues, or very small
      * radius). In this case use the view direction to approximate the normal. */
     const float3 P_inside = float4_to_float3(catmull_rom_basis_eval(P_curve, sd->u));
-    const float3 N = (!isequal(P, P_inside)) ? normalize(P - P_inside) : -sd->wi;
+    const float3 N = (!isequal(P, P_inside)) ? normalize(P - P_inside) : -sd->I;
 
     sd->N = N;
     sd->v = 0.0f;
@@ -757,7 +757,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg,
   }
 
   sd->P = P;
-  sd->Ng = (sd->type & PRIMITIVE_CURVE_RIBBON) ? sd->wi : sd->N;
+  sd->Ng = (sd->type & PRIMITIVE_CURVE_RIBBON) ? sd->I : sd->N;
   sd->dPdv = cross(sd->dPdu, sd->Ng);
   sd->shader = kernel_data_fetch(curves, sd->prim).shader_id;
 }

intern/cycles/kernel/geom/shader_data.h

@@ -55,7 +55,7 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
 #endif
 
   /* Read ray data into shader globals. */
-  sd->wi = -ray->D;
+  sd->I = -ray->D;
 
 #ifdef __HAIR__
   if (sd->type & PRIMITIVE_CURVE) {
@@ -111,7 +111,7 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
   sd->flag = kernel_data_fetch(shaders, (sd->shader & SHADER_MASK)).flags;
 
   /* backfacing test */
-  bool backfacing = (dot(sd->Ng, sd->wi) < 0.0f);
+  bool backfacing = (dot(sd->Ng, sd->I) < 0.0f);
 
   if (backfacing) {
     sd->flag |= SD_BACKFACING;
@@ -152,7 +152,7 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals kg,
   sd->P = P;
   sd->N = Ng;
   sd->Ng = Ng;
-  sd->wi = I;
+  sd->I = I;
   sd->shader = shader;
   if (prim != PRIM_NONE)
     sd->type = PRIMITIVE_TRIANGLE;
@@ -185,7 +185,7 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals kg,
       object_position_transform_auto(kg, sd, &sd->P);
       object_normal_transform_auto(kg, sd, &sd->Ng);
       sd->N = sd->Ng;
-      object_dir_transform_auto(kg, sd, &sd->wi);
+      object_dir_transform_auto(kg, sd, &sd->I);
     }
 
     if (sd->type == PRIMITIVE_TRIANGLE) {
@@ -227,7 +227,7 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals kg,
 
   /* backfacing test */
   if (sd->prim != PRIM_NONE) {
-    bool backfacing = (dot(sd->Ng, sd->wi) < 0.0f);
+    bool backfacing = (dot(sd->Ng, sd->I) < 0.0f);
 
     if (backfacing) {
       sd->flag |= SD_BACKFACING;
@@ -341,7 +341,7 @@ ccl_device void shader_setup_from_curve(KernelGlobals kg,
   }
 
   /* No view direction, normals or bitangent. */
-  sd->wi = zero_float3();
+  sd->I = zero_float3();
   sd->N = zero_float3();
   sd->Ng = zero_float3();
 #ifdef __DPDU__
@@ -372,7 +372,7 @@ ccl_device_inline void shader_setup_from_background(KernelGlobals kg,
   sd->P = ray_D;
   sd->N = -ray_D;
   sd->Ng = -ray_D;
-  sd->wi = -ray_D;
+  sd->I = -ray_D;
   sd->shader = kernel_data.background.surface_shader;
   sd->flag = kernel_data_fetch(shaders, (sd->shader & SHADER_MASK)).flags;
   sd->object_flag = 0;
@@ -412,7 +412,7 @@ ccl_device_inline void shader_setup_from_volume(KernelGlobals kg,
   sd->P = ray->P + ray->D * ray->tmin;
   sd->N = -ray->D;
   sd->Ng = -ray->D;
-  sd->wi = -ray->D;
+  sd->I = -ray->D;
   sd->shader = SHADER_NONE;
   sd->flag = 0;
   sd->object_flag = 0;

intern/cycles/kernel/integrator/guiding.h

@@ -44,7 +44,7 @@ ccl_device_forceinline void guiding_record_surface_segment(KernelGlobals kg,
 
   state->guiding.path_segment = kg->opgl_path_segment_storage->NextSegment();
   openpgl::cpp::SetPosition(state->guiding.path_segment, guiding_point3f(sd->P));
-  openpgl::cpp::SetDirectionOut(state->guiding.path_segment, guiding_vec3f(sd->wi));
+  openpgl::cpp::SetDirectionOut(state->guiding.path_segment, guiding_vec3f(sd->I));
   openpgl::cpp::SetVolumeScatter(state->guiding.path_segment, false);
   openpgl::cpp::SetScatteredContribution(state->guiding.path_segment, zero);
   openpgl::cpp::SetDirectContribution(state->guiding.path_segment, zero);
@@ -60,7 +60,7 @@ ccl_device_forceinline void guiding_record_surface_bounce(KernelGlobals kg,
                                                           const Spectrum weight,
                                                           const float pdf,
                                                           const float3 N,
-                                                          const float3 wo,
+                                                          const float3 omega_in,
                                                           const float2 roughness,
                                                           const float eta)
 {
@@ -78,7 +78,7 @@ ccl_device_forceinline void guiding_record_surface_bounce(KernelGlobals kg,
   openpgl::cpp::SetTransmittanceWeight(state->guiding.path_segment, guiding_vec3f(one_float3()));
   openpgl::cpp::SetVolumeScatter(state->guiding.path_segment, false);
   openpgl::cpp::SetNormal(state->guiding.path_segment, guiding_vec3f(normal));
-  openpgl::cpp::SetDirectionIn(state->guiding.path_segment, guiding_vec3f(wo));
+  openpgl::cpp::SetDirectionIn(state->guiding.path_segment, guiding_vec3f(omega_in));
   openpgl::cpp::SetPDFDirectionIn(state->guiding.path_segment, pdf);
   openpgl::cpp::SetScatteringWeight(state->guiding.path_segment, guiding_vec3f(weight_rgb));
   openpgl::cpp::SetIsDelta(state->guiding.path_segment, is_delta);
@@ -113,7 +113,7 @@ ccl_device_forceinline void guiding_record_surface_emission(KernelGlobals kg,
 ccl_device_forceinline void guiding_record_bssrdf_segment(KernelGlobals kg,
                                                           IntegratorState state,
                                                           const float3 P,
-                                                          const float3 wi)
+                                                          const float3 I)
 {
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 1
   if (!kernel_data.integrator.train_guiding) {
@@ -124,7 +124,7 @@ ccl_device_forceinline void guiding_record_bssrdf_segment(KernelGlobals kg,
 
   state->guiding.path_segment = kg->opgl_path_segment_storage->NextSegment();
   openpgl::cpp::SetPosition(state->guiding.path_segment, guiding_point3f(P));
-  openpgl::cpp::SetDirectionOut(state->guiding.path_segment, guiding_vec3f(wi));
+  openpgl::cpp::SetDirectionOut(state->guiding.path_segment, guiding_vec3f(I));
   openpgl::cpp::SetVolumeScatter(state->guiding.path_segment, true);
   openpgl::cpp::SetScatteredContribution(state->guiding.path_segment, zero);
   openpgl::cpp::SetDirectContribution(state->guiding.path_segment, zero);
@@ -166,7 +166,7 @@ ccl_device_forceinline void guiding_record_bssrdf_bounce(KernelGlobals kg,
                                                          IntegratorState state,
                                                          const float pdf,
                                                          const float3 N,
-                                                         const float3 wo,
+                                                         const float3 omega_in,
                                                          const Spectrum weight,
                                                          const Spectrum albedo)
 {
@@ -181,7 +181,7 @@ ccl_device_forceinline void guiding_record_bssrdf_bounce(KernelGlobals kg,
 
   openpgl::cpp::SetVolumeScatter(state->guiding.path_segment, false);
   openpgl::cpp::SetNormal(state->guiding.path_segment, guiding_vec3f(normal));
-  openpgl::cpp::SetDirectionIn(state->guiding.path_segment, guiding_vec3f(wo));
+  openpgl::cpp::SetDirectionIn(state->guiding.path_segment, guiding_vec3f(omega_in));
   openpgl::cpp::SetPDFDirectionIn(state->guiding.path_segment, pdf);
   openpgl::cpp::SetTransmittanceWeight(state->guiding.path_segment, guiding_vec3f(weight_rgb));
 #endif
@@ -222,7 +222,7 @@ ccl_device_forceinline void guiding_record_volume_bounce(KernelGlobals kg,
                                                          ccl_private const ShaderData *sd,
                                                          const Spectrum weight,
                                                          const float pdf,
-                                                         const float3 wo,
+                                                         const float3 omega_in,
                                                          const float roughness)
 {
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
@@ -237,7 +237,7 @@ ccl_device_forceinline void guiding_record_volume_bounce(KernelGlobals kg,
   openpgl::cpp::SetVolumeScatter(state->guiding.path_segment, true);
   openpgl::cpp::SetTransmittanceWeight(state->guiding.path_segment, guiding_vec3f(one_float3()));
   openpgl::cpp::SetNormal(state->guiding.path_segment, guiding_vec3f(normal));
-  openpgl::cpp::SetDirectionIn(state->guiding.path_segment, guiding_vec3f(wo));
+  openpgl::cpp::SetDirectionIn(state->guiding.path_segment, guiding_vec3f(omega_in));
   openpgl::cpp::SetPDFDirectionIn(state->guiding.path_segment, pdf);
   openpgl::cpp::SetScatteringWeight(state->guiding.path_segment, guiding_vec3f(weight_rgb));
   openpgl::cpp::SetIsDelta(state->guiding.path_segment, false);
@@ -467,13 +467,13 @@ ccl_device_forceinline bool guiding_bsdf_init(KernelGlobals kg,
 ccl_device_forceinline float guiding_bsdf_sample(KernelGlobals kg,
                                                  IntegratorState state,
                                                  const float2 rand_bsdf,
-                                                 ccl_private float3 *wo)
+                                                 ccl_private float3 *omega_in)
 {
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
-  pgl_vec3f pgl_wo;
+  pgl_vec3f wo;
   const pgl_point2f rand = openpgl::cpp::Point2(rand_bsdf.x, rand_bsdf.y);
-  const float pdf = kg->opgl_surface_sampling_distribution->SamplePDF(rand, pgl_wo);
-  *wo = make_float3(pgl_wo.x, pgl_wo.y, pgl_wo.z);
+  const float pdf = kg->opgl_surface_sampling_distribution->SamplePDF(rand, wo);
+  *omega_in = make_float3(wo.x, wo.y, wo.z);
   return pdf;
 #else
   return 0.0f;
@@ -482,10 +482,10 @@ ccl_device_forceinline float guiding_bsdf_sample(KernelGlobals kg,
 
 ccl_device_forceinline float guiding_bsdf_pdf(KernelGlobals kg,
                                               IntegratorState state,
-                                              const float3 wo)
+                                              const float3 omega_in)
 {
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
-  return kg->opgl_surface_sampling_distribution->PDF(guiding_vec3f(wo));
+  return kg->opgl_surface_sampling_distribution->PDF(guiding_vec3f(omega_in));
 #else
   return 0.0f;
 #endif
@@ -520,13 +520,13 @@ ccl_device_forceinline bool guiding_phase_init(KernelGlobals kg,
 ccl_device_forceinline float guiding_phase_sample(KernelGlobals kg,
                                                   IntegratorState state,
                                                   const float2 rand_phase,
-                                                  ccl_private float3 *wo)
+                                                  ccl_private float3 *omega_in)
 {
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
-  pgl_vec3f pgl_wo;
+  pgl_vec3f wo;
   const pgl_point2f rand = openpgl::cpp::Point2(rand_phase.x, rand_phase.y);
-  const float pdf = kg->opgl_volume_sampling_distribution->SamplePDF(rand, pgl_wo);
-  *wo = make_float3(pgl_wo.x, pgl_wo.y, pgl_wo.z);
+  const float pdf = kg->opgl_volume_sampling_distribution->SamplePDF(rand, wo);
+  *omega_in = make_float3(wo.x, wo.y, wo.z);
   return pdf;
 #else
   return 0.0f;
@@ -535,10 +535,10 @@ ccl_device_forceinline float guiding_phase_sample(KernelGlobals kg,
 
 ccl_device_forceinline float guiding_phase_pdf(KernelGlobals kg,
                                                IntegratorState state,
-                                               const float3 wo)
+                                               const float3 omega_in)
 {
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
-  return kg->opgl_volume_sampling_distribution->PDF(guiding_vec3f(wo));
+  return kg->opgl_volume_sampling_distribution->PDF(guiding_vec3f(omega_in));
 #else
   return 0.0f;
 #endif

intern/cycles/kernel/integrator/mnee.h

@@ -607,22 +607,24 @@ ccl_device_forceinline Spectrum mnee_eval_bsdf_contribution(ccl_private ShaderCl
 {
   ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)closure;
 
-  float cosNI = dot(bsdf->N, wi);
-  float cosNO = dot(bsdf->N, wo);
+  float cosNO = dot(bsdf->N, wi);
+  float cosNI = dot(bsdf->N, wo);
 
   float3 Ht = normalize(-(bsdf->ior * wo + wi));
-  float cosHI = dot(Ht, wi);
+  float cosHO = dot(Ht, wi);
 
   float alpha2 = bsdf->alpha_x * bsdf->alpha_y;
   float cosThetaM = dot(bsdf->N, Ht);
 
-  /* Now calculate G1(i, m) and G1(o, m). */
   float G;
   if (bsdf->type == CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID) {
-    G = bsdf_G<MicrofacetType::BECKMANN>(alpha2, cosNI, cosNO);
+    /* Eq. 26, 27: now calculate G1(i,m) and G1(o,m). */
+    G = bsdf_beckmann_G1(bsdf->alpha_x, cosNO) * bsdf_beckmann_G1(bsdf->alpha_x, cosNI);
   }
   else { /* bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID assumed */
-    G = bsdf_G<MicrofacetType::GGX>(alpha2, cosNI, cosNO);
+    /* Eq. 34: now calculate G1(i,m) and G1(o,m). */
+    G = (2.f / (1.f + safe_sqrtf(1.f + alpha2 * (1.f - cosNO * cosNO) / (cosNO * cosNO)))) *
+        (2.f / (1.f + safe_sqrtf(1.f + alpha2 * (1.f - cosNI * cosNI) / (cosNI * cosNI))));
   }
 
   /*
@@ -633,7 +635,7 @@ ccl_device_forceinline Spectrum mnee_eval_bsdf_contribution(ccl_private ShaderCl
    * contribution = bsdf_do * |do/dh| * |n.wo / n.h| / pdf_dh
    *              = (1 - F) * G * |h.wi / (n.wi * n.h^2)|
    */
-  return bsdf->weight * G * fabsf(cosHI / (cosNI * sqr(cosThetaM)));
+  return bsdf->weight * G * fabsf(cosHO / (cosNO * sqr(cosThetaM)));
 }
 
 /* Compute transfer matrix determinant |T1| = |dx1/dxn| (and |dh/dx| in the process) */
@@ -704,9 +706,9 @@ ccl_device_forceinline bool mnee_compute_transfer_matrix(ccl_private const Shade
     float ilo = -eta * ilh;
 
     float cos_theta = dot(wo, m.n);
-    float sin_theta = sin_from_cos(cos_theta);
+    float sin_theta = safe_sqrtf(1.f - sqr(cos_theta));
     float cos_phi = dot(wo, s);
-    float sin_phi = sin_from_cos(cos_phi);
+    float sin_phi = safe_sqrtf(1.f - sqr(cos_phi));
 
     /* Wo = (cos_phi * sin_theta) * s + (sin_phi * sin_theta) * t + cos_theta * n. */
     float3 dH_dtheta = ilo * (cos_theta * (cos_phi * s + sin_phi * t) - sin_theta * m.n);

intern/cycles/kernel/integrator/shade_surface.h

@@ -235,6 +235,8 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
     light_sample_to_surface_shadow_ray(kg, sd, &ls, &ray);
   }
 
+  const bool is_light = light_sample_is_light(&ls);
+
   /* Branch off shadow kernel. */
   IntegratorShadowState shadow_state = integrator_shadow_path_init(
       kg, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, false);
@@ -262,6 +264,7 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
 
   /* Copy state from main path to shadow path. */
   uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
+  shadow_flag |= (is_light) ? PATH_RAY_SHADOW_FOR_LIGHT : 0;
   const Spectrum unlit_throughput = INTEGRATOR_STATE(state, path, throughput);
   const Spectrum throughput = unlit_throughput * bsdf_eval_sum(&bsdf_eval);
 
@@ -361,7 +364,7 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
   /* BSDF closure, sample direction. */
   float bsdf_pdf = 0.0f, unguided_bsdf_pdf = 0.0f;
   BsdfEval bsdf_eval ccl_optional_struct_init;
-  float3 bsdf_wo ccl_optional_struct_init;
+  float3 bsdf_omega_in ccl_optional_struct_init;
   int label;
 
   float2 bsdf_sampled_roughness = make_float2(1.0f, 1.0f);
@@ -375,7 +378,7 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
                                                       sc,
                                                       rand_bsdf,
                                                       &bsdf_eval,
-                                                      &bsdf_wo,
+                                                      &bsdf_omega_in,
                                                       &bsdf_pdf,
                                                       &unguided_bsdf_pdf,
                                                       &bsdf_sampled_roughness,
@@ -395,7 +398,7 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
                                                sc,
                                                rand_bsdf,
                                                &bsdf_eval,
-                                               &bsdf_wo,
+                                               &bsdf_omega_in,
                                                &bsdf_pdf,
                                                &bsdf_sampled_roughness,
                                                &bsdf_eta);
@@ -413,7 +416,7 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
   }
   else {
     /* Setup ray with changed origin and direction. */
-    const float3 D = normalize(bsdf_wo);
+    const float3 D = normalize(bsdf_omega_in);
     INTEGRATOR_STATE_WRITE(state, ray, P) = integrate_surface_ray_offset(kg, sd, sd->P, D);
     INTEGRATOR_STATE_WRITE(state, ray, D) = D;
     INTEGRATOR_STATE_WRITE(state, ray, tmin) = 0.0f;
@@ -452,7 +455,7 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
                                 bsdf_weight,
                                 bsdf_pdf,
                                 sd->N,
-                                normalize(bsdf_wo),
+                                normalize(bsdf_omega_in),
                                 bsdf_sampled_roughness,
                                 bsdf_eta);
 

intern/cycles/kernel/integrator/shade_volume.h

@@ -821,6 +821,7 @@ ccl_device_forceinline void integrate_volume_direct_light(
   /* Create shadow ray. */
   Ray ray ccl_optional_struct_init;
   light_sample_to_volume_shadow_ray(kg, sd, &ls, P, &ray);
+  const bool is_light = light_sample_is_light(&ls);
 
   /* Branch off shadow kernel. */
   IntegratorShadowState shadow_state = integrator_shadow_path_init(
@@ -837,6 +838,7 @@ ccl_device_forceinline void integrate_volume_direct_light(
   const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);
   const uint16_t transparent_bounce = INTEGRATOR_STATE(state, path, transparent_bounce);
   uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
+  shadow_flag |= (is_light) ? PATH_RAY_SHADOW_FOR_LIGHT : 0;
   const Spectrum throughput_phase = throughput * bsdf_eval_sum(&phase_eval);
 
   if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
@@ -910,7 +912,7 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
   /* Phase closure, sample direction. */
   float phase_pdf = 0.0f, unguided_phase_pdf = 0.0f;
   BsdfEval phase_eval ccl_optional_struct_init;
-  float3 phase_wo ccl_optional_struct_init;
+  float3 phase_omega_in ccl_optional_struct_init;
   float sampled_roughness = 1.0f;
   int label;
 
@@ -922,7 +924,7 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
                                               svc,
                                               rand_phase,
                                               &phase_eval,
-                                              &phase_wo,
+                                              &phase_omega_in,
                                               &phase_pdf,
                                               &unguided_phase_pdf,
                                               &sampled_roughness);
@@ -936,8 +938,15 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
   else
 #  endif
   {
-    label = volume_shader_phase_sample(
-        kg, sd, phases, svc, rand_phase, &phase_eval, &phase_wo, &phase_pdf, &sampled_roughness);
+    label = volume_shader_phase_sample(kg,
+                                       sd,
+                                       phases,
+                                       svc,
+                                       rand_phase,
+                                       &phase_eval,
+                                       &phase_omega_in,
+                                       &phase_pdf,
+                                       &sampled_roughness);
 
     if (phase_pdf == 0.0f || bsdf_eval_is_zero(&phase_eval)) {
       return false;
@@ -948,7 +957,7 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
 
   /* Setup ray. */
   INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
-  INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(phase_wo);
+  INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(phase_omega_in);
   INTEGRATOR_STATE_WRITE(state, ray, tmin) = 0.0f;
   INTEGRATOR_STATE_WRITE(state, ray, tmax) = FLT_MAX;
 #  ifdef __RAY_DIFFERENTIALS__
@@ -962,7 +971,7 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
 
   /* Add phase function sampling data to the path segment. */
   guiding_record_volume_bounce(
-      kg, state, sd, phase_weight, phase_pdf, normalize(phase_wo), sampled_roughness);
+      kg, state, sd, phase_weight, phase_pdf, normalize(phase_omega_in), sampled_roughness);
 
   /* Update throughput. */
   const Spectrum throughput = INTEGRATOR_STATE(state, path, throughput);
@@ -1067,7 +1076,7 @@ ccl_device VolumeIntegrateEvent volume_integrate(KernelGlobals kg,
         float3 transmittance_weight = spectrum_to_rgb(
             safe_divide_color(result.indirect_throughput, initial_throughput));
         guiding_record_volume_transmission(kg, state, transmittance_weight);
-        guiding_record_volume_segment(kg, state, direct_P, sd.wi);
+        guiding_record_volume_segment(kg, state, direct_P, sd.I);
         guiding_generated_new_segment = true;
         unlit_throughput = result.indirect_throughput / continuation_probability;
         rand_phase_guiding = path_state_rng_1D(kg, &rng_state, PRNG_VOLUME_PHASE_GUIDING_DISTANCE);
@@ -1130,7 +1139,7 @@ ccl_device VolumeIntegrateEvent volume_integrate(KernelGlobals kg,
 #  if defined(__PATH_GUIDING__)
 #    if PATH_GUIDING_LEVEL >= 1
     if (!guiding_generated_new_segment) {
-      guiding_record_volume_segment(kg, state, sd.P, sd.wi);
+      guiding_record_volume_segment(kg, state, sd.P, sd.I);
     }
 #    endif
 #    if PATH_GUIDING_LEVEL >= 4

intern/cycles/kernel/integrator/state.h

@@ -132,9 +132,6 @@ typedef struct IntegratorStateGPU {
   /* Index of main path which will be used by a next shadow catcher split.  */
   ccl_global int *next_main_path_index;
 
-  /* Partition/key offsets used when writing sorted active indices. */
-  ccl_global int *sort_partition_key_offsets;
-
   /* Divisor used to partition active indices by locality when sorting by material.  */
   uint sort_partition_divisor;
 } IntegratorStateGPU;

intern/cycles/kernel/integrator/state_flow.h

@@ -115,13 +115,6 @@ ccl_device_forceinline void integrator_path_init_sorted(KernelGlobals kg,
   atomic_fetch_and_add_uint32(&kernel_integrator_state.queue_counter->num_queued[next_kernel], 1);
   INTEGRATOR_STATE_WRITE(state, path, queued_kernel) = next_kernel;
   INTEGRATOR_STATE_WRITE(state, path, shader_sort_key) = key_;
-
-#  if defined(__KERNEL_LOCAL_ATOMIC_SORT__)
-  if (!kernel_integrator_state.sort_key_counter[next_kernel]) {
-    return;
-  }
-#  endif
-
   atomic_fetch_and_add_uint32(&kernel_integrator_state.sort_key_counter[next_kernel][key_], 1);
 }
 
@@ -137,13 +130,6 @@ ccl_device_forceinline void integrator_path_next_sorted(KernelGlobals kg,
   atomic_fetch_and_add_uint32(&kernel_integrator_state.queue_counter->num_queued[next_kernel], 1);
   INTEGRATOR_STATE_WRITE(state, path, queued_kernel) = next_kernel;
   INTEGRATOR_STATE_WRITE(state, path, shader_sort_key) = key_;
-
-#  if defined(__KERNEL_LOCAL_ATOMIC_SORT__)
-  if (!kernel_integrator_state.sort_key_counter[next_kernel]) {
-    return;
-  }
-#  endif
-
   atomic_fetch_and_add_uint32(&kernel_integrator_state.sort_key_counter[next_kernel][key_], 1);
 }
 

intern/cycles/kernel/integrator/subsurface_random_walk.h

@@ -136,7 +136,7 @@ ccl_device_forceinline float diffusion_length_dwivedi(float alpha)
 
 ccl_device_forceinline float3 direction_from_cosine(float3 D, float cos_theta, float randv)
 {
-  float sin_theta = sin_from_cos(cos_theta);
+  float sin_theta = safe_sqrtf(1.0f - cos_theta * cos_theta);
   float phi = M_2PI_F * randv;
   float3 dir = make_float3(sin_theta * cosf(phi), sin_theta * sinf(phi), cos_theta);
 

intern/cycles/kernel/integrator/surface_shader.h

@@ -174,14 +174,14 @@ ccl_device_inline void surface_shader_prepare_closures(KernelGlobals kg,
 #if 0
 ccl_device_inline void surface_shader_validate_bsdf_sample(const KernelGlobals kg,
                                                            const ShaderClosure *sc,
-                                                           const float3 wo,
+                                                           const float3 omega_in,
                                                            const int org_label,
                                                            const float2 org_roughness,
                                                            const float org_eta)
 {
   /* Validate the the bsdf_label and bsdf_roughness_eta functions
    * by estimating the values after a bsdf sample. */
-  const int comp_label = bsdf_label(kg, sc, wo);
+  const int comp_label = bsdf_label(kg, sc, omega_in);
   kernel_assert(org_label == comp_label);
 
   float2 comp_roughness;
@@ -218,7 +218,7 @@ ccl_device_forceinline bool _surface_shader_exclude(ClosureType type, uint light
 
 ccl_device_inline float _surface_shader_bsdf_eval_mis(KernelGlobals kg,
                                                       ccl_private ShaderData *sd,
-                                                      const float3 wo,
+                                                      const float3 omega_in,
                                                       ccl_private const ShaderClosure *skip_sc,
                                                       ccl_private BsdfEval *result_eval,
                                                       float sum_pdf,
@@ -237,7 +237,7 @@ ccl_device_inline float _surface_shader_bsdf_eval_mis(KernelGlobals kg,
     if (CLOSURE_IS_BSDF_OR_BSSRDF(sc->type)) {
       if (CLOSURE_IS_BSDF(sc->type) && !_surface_shader_exclude(sc->type, light_shader_flags)) {
         float bsdf_pdf = 0.0f;
-        Spectrum eval = bsdf_eval(kg, sd, sc, wo, &bsdf_pdf);
+        Spectrum eval = bsdf_eval(kg, sd, sc, omega_in, &bsdf_pdf);
 
         if (bsdf_pdf != 0.0f) {
           bsdf_eval_accum(result_eval, sc->type, eval * sc->weight);
@@ -254,7 +254,7 @@ ccl_device_inline float _surface_shader_bsdf_eval_mis(KernelGlobals kg,
 
 ccl_device_inline float surface_shader_bsdf_eval_pdfs(const KernelGlobals kg,
                                                       ccl_private ShaderData *sd,
-                                                      const float3 wo,
+                                                      const float3 omega_in,
                                                       ccl_private BsdfEval *result_eval,
                                                       ccl_private float *pdfs,
                                                       const uint light_shader_flags)
@@ -270,7 +270,7 @@ ccl_device_inline float surface_shader_bsdf_eval_pdfs(const KernelGlobals kg,
     if (CLOSURE_IS_BSDF_OR_BSSRDF(sc->type)) {
       if (CLOSURE_IS_BSDF(sc->type) && !_surface_shader_exclude(sc->type, light_shader_flags)) {
         float bsdf_pdf = 0.0f;
-        Spectrum eval = bsdf_eval(kg, sd, sc, wo, &bsdf_pdf);
+        Spectrum eval = bsdf_eval(kg, sd, sc, omega_in, &bsdf_pdf);
         kernel_assert(bsdf_pdf >= 0.0f);
         if (bsdf_pdf != 0.0f) {
           bsdf_eval_accum(result_eval, sc->type, eval * sc->weight);
@@ -310,20 +310,20 @@ ccl_device_inline
     surface_shader_bsdf_eval(KernelGlobals kg,
                              IntegratorState state,
                              ccl_private ShaderData *sd,
-                             const float3 wo,
+                             const float3 omega_in,
                              ccl_private BsdfEval *bsdf_eval,
                              const uint light_shader_flags)
 {
   bsdf_eval_init(bsdf_eval, CLOSURE_NONE_ID, zero_spectrum());
 
   float pdf = _surface_shader_bsdf_eval_mis(
-      kg, sd, wo, NULL, bsdf_eval, 0.0f, 0.0f, light_shader_flags);
+      kg, sd, omega_in, NULL, bsdf_eval, 0.0f, 0.0f, light_shader_flags);
 
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
   if (state->guiding.use_surface_guiding) {
     const float guiding_sampling_prob = state->guiding.surface_guiding_sampling_prob;
     const float bssrdf_sampling_prob = state->guiding.bssrdf_sampling_prob;
-    const float guide_pdf = guiding_bsdf_pdf(kg, state, wo);
+    const float guide_pdf = guiding_bsdf_pdf(kg, state, omega_in);
     pdf = (guiding_sampling_prob * guide_pdf * (1.0f - bssrdf_sampling_prob)) +
           (1.0f - guiding_sampling_prob) * pdf;
   }
@@ -407,7 +407,7 @@ ccl_device int surface_shader_bsdf_guided_sample_closure(KernelGlobals kg,
                                                          ccl_private const ShaderClosure *sc,
                                                          const float2 rand_bsdf,
                                                          ccl_private BsdfEval *bsdf_eval,
-                                                         ccl_private float3 *wo,
+                                                         ccl_private float3 *omega_in,
                                                          ccl_private float *bsdf_pdf,
                                                          ccl_private float *unguided_bsdf_pdf,
                                                          ccl_private float2 *sampled_rougness,
@@ -443,14 +443,14 @@ ccl_device int surface_shader_bsdf_guided_sample_closure(KernelGlobals kg,
 
   if (sample_guiding) {
     /* Sample guiding distribution. */
-    guide_pdf = guiding_bsdf_sample(kg, state, rand_bsdf, wo);
+    guide_pdf = guiding_bsdf_sample(kg, state, rand_bsdf, omega_in);
     *bsdf_pdf = 0.0f;
 
     if (guide_pdf != 0.0f) {
       float unguided_bsdf_pdfs[MAX_CLOSURE];
 
       *unguided_bsdf_pdf = surface_shader_bsdf_eval_pdfs(
-          kg, sd, *wo, bsdf_eval, unguided_bsdf_pdfs, 0);
+          kg, sd, *omega_in, bsdf_eval, unguided_bsdf_pdfs, 0);
       *bsdf_pdf = (guiding_sampling_prob * guide_pdf * (1.0f - bssrdf_sampling_prob)) +
                   ((1.0f - guiding_sampling_prob) * (*unguided_bsdf_pdf));
       float sum_pdfs = 0.0f;
@@ -471,7 +471,7 @@ ccl_device int surface_shader_bsdf_guided_sample_closure(KernelGlobals kg,
          * the sum of all unguided_bsdf_pdfs is just < 1.0f. */
         idx = (rand_bsdf_guiding > sum_pdfs) ? sd->num_closure - 1 : idx;
 
-        label = bsdf_label(kg, &sd->closure[idx], *wo);
+        label = bsdf_label(kg, &sd->closure[idx], *omega_in);
       }
     }
 
@@ -483,11 +483,19 @@ ccl_device int surface_shader_bsdf_guided_sample_closure(KernelGlobals kg,
   else {
     /* Sample BSDF. */
     *bsdf_pdf = 0.0f;
-    label = bsdf_sample(
-        kg, sd, sc, rand_bsdf.x, rand_bsdf.y, &eval, wo, unguided_bsdf_pdf, sampled_rougness, eta);
+    label = bsdf_sample(kg,
+                        sd,
+                        sc,
+                        rand_bsdf.x,
+                        rand_bsdf.y,
+                        &eval,
+                        omega_in,
+                        unguided_bsdf_pdf,
+                        sampled_rougness,
+                        eta);
 #  if 0
     if (*unguided_bsdf_pdf > 0.0f) {
-      surface_shader_validate_bsdf_sample(kg, sc, *wo, label, sampled_roughness, eta);
+      surface_shader_validate_bsdf_sample(kg, sc, *omega_in, label, sampled_roughness, eta);
     }
 #  endif
 
@@ -499,13 +507,13 @@ ccl_device int surface_shader_bsdf_guided_sample_closure(KernelGlobals kg,
       if (sd->num_closure > 1) {
         float sweight = sc->sample_weight;
         *unguided_bsdf_pdf = _surface_shader_bsdf_eval_mis(
-            kg, sd, *wo, sc, bsdf_eval, (*unguided_bsdf_pdf) * sweight, sweight, 0);
+            kg, sd, *omega_in, sc, bsdf_eval, (*unguided_bsdf_pdf) * sweight, sweight, 0);
         kernel_assert(reduce_min(bsdf_eval_sum(bsdf_eval)) >= 0.0f);
       }
       *bsdf_pdf = *unguided_bsdf_pdf;
 
       if (use_surface_guiding) {
-        guide_pdf = guiding_bsdf_pdf(kg, state, *wo);
+        guide_pdf = guiding_bsdf_pdf(kg, state, *omega_in);
         *bsdf_pdf *= 1.0f - guiding_sampling_prob;
         *bsdf_pdf += guiding_sampling_prob * guide_pdf * (1.0f - bssrdf_sampling_prob);
       }
@@ -525,7 +533,7 @@ ccl_device int surface_shader_bsdf_sample_closure(KernelGlobals kg,
                                                   ccl_private const ShaderClosure *sc,
                                                   const float2 rand_bsdf,
                                                   ccl_private BsdfEval *bsdf_eval,
-                                                  ccl_private float3 *wo,
+                                                  ccl_private float3 *omega_in,
                                                   ccl_private float *pdf,
                                                   ccl_private float2 *sampled_roughness,
                                                   ccl_private float *eta)
@@ -538,14 +546,15 @@ ccl_device int surface_shader_bsdf_sample_closure(KernelGlobals kg,
 
   *pdf = 0.0f;
   label = bsdf_sample(
-      kg, sd, sc, rand_bsdf.x, rand_bsdf.y, &eval, wo, pdf, sampled_roughness, eta);
+      kg, sd, sc, rand_bsdf.x, rand_bsdf.y, &eval, omega_in, pdf, sampled_roughness, eta);
 
   if (*pdf != 0.0f) {
     bsdf_eval_init(bsdf_eval, sc->type, eval * sc->weight);
 
     if (sd->num_closure > 1) {
       float sweight = sc->sample_weight;
-      *pdf = _surface_shader_bsdf_eval_mis(kg, sd, *wo, sc, bsdf_eval, *pdf * sweight, sweight, 0);
+      *pdf = _surface_shader_bsdf_eval_mis(
+          kg, sd, *omega_in, sc, bsdf_eval, *pdf * sweight, sweight, 0);
     }
   }
   else {
@@ -621,7 +630,7 @@ ccl_device Spectrum surface_shader_diffuse(KernelGlobals kg, ccl_private const S
     ccl_private const ShaderClosure *sc = &sd->closure[i];
 
     if (CLOSURE_IS_BSDF_DIFFUSE(sc->type) || CLOSURE_IS_BSSRDF(sc->type))
-      eval += bsdf_albedo(sd, sc);
+      eval += sc->weight;
   }
 
   return eval;
@@ -635,7 +644,7 @@ ccl_device Spectrum surface_shader_glossy(KernelGlobals kg, ccl_private const Sh
     ccl_private const ShaderClosure *sc = &sd->closure[i];
 
     if (CLOSURE_IS_BSDF_GLOSSY(sc->type))
-      eval += bsdf_albedo(sd, sc);
+      eval += sc->weight;
   }
 
   return eval;
@@ -649,7 +658,7 @@ ccl_device Spectrum surface_shader_transmission(KernelGlobals kg, ccl_private co
     ccl_private const ShaderClosure *sc = &sd->closure[i];
 
     if (CLOSURE_IS_BSDF_TRANSMISSION(sc->type))
-      eval += bsdf_albedo(sd, sc);
+      eval += sc->weight;
   }
 
   return eval;
@@ -749,7 +758,7 @@ ccl_device Spectrum surface_shader_background(ccl_private const ShaderData *sd)
 ccl_device Spectrum surface_shader_emission(ccl_private const ShaderData *sd)
 {
   if (sd->flag & SD_EMISSION) {
-    return emissive_simple_eval(sd->Ng, sd->wi) * sd->closure_emission_background;
+    return emissive_simple_eval(sd->Ng, sd->I) * sd->closure_emission_background;
   }
   else {
     return zero_spectrum();

intern/cycles/kernel/integrator/volume_shader.h

@@ -202,7 +202,7 @@ ccl_device_inline ccl_private const ShaderVolumeClosure *volume_shader_phase_pic
 
 ccl_device_inline float _volume_shader_phase_eval_mis(ccl_private const ShaderData *sd,
                                                       ccl_private const ShaderVolumePhases *phases,
-                                                      const float3 wo,
+                                                      const float3 omega_in,
                                                       int skip_phase,
                                                       ccl_private BsdfEval *result_eval,
                                                       float sum_pdf,
@@ -214,7 +214,7 @@ ccl_device_inline float _volume_shader_phase_eval_mis(ccl_private const ShaderDa
 
     ccl_private const ShaderVolumeClosure *svc = &phases->closure[i];
     float phase_pdf = 0.0f;
-    Spectrum eval = volume_phase_eval(sd, svc, wo, &phase_pdf);
+    Spectrum eval = volume_phase_eval(sd, svc, omega_in, &phase_pdf);
 
     if (phase_pdf != 0.0f) {
       bsdf_eval_accum(result_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, eval);
@@ -230,11 +230,11 @@ ccl_device_inline float _volume_shader_phase_eval_mis(ccl_private const ShaderDa
 ccl_device float volume_shader_phase_eval(KernelGlobals kg,
                                           ccl_private const ShaderData *sd,
                                           ccl_private const ShaderVolumeClosure *svc,
-                                          const float3 wo,
+                                          const float3 omega_in,
                                           ccl_private BsdfEval *phase_eval)
 {
   float phase_pdf = 0.0f;
-  Spectrum eval = volume_phase_eval(sd, svc, wo, &phase_pdf);
+  Spectrum eval = volume_phase_eval(sd, svc, omega_in, &phase_pdf);
 
   if (phase_pdf != 0.0f) {
     bsdf_eval_accum(phase_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, eval);
@@ -247,17 +247,17 @@ ccl_device float volume_shader_phase_eval(KernelGlobals kg,
                                           IntegratorState state,
                                           ccl_private const ShaderData *sd,
                                           ccl_private const ShaderVolumePhases *phases,
-                                          const float3 wo,
+                                          const float3 omega_in,
                                           ccl_private BsdfEval *phase_eval)
 {
   bsdf_eval_init(phase_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, zero_spectrum());
 
-  float pdf = _volume_shader_phase_eval_mis(sd, phases, wo, -1, phase_eval, 0.0f, 0.0f);
+  float pdf = _volume_shader_phase_eval_mis(sd, phases, omega_in, -1, phase_eval, 0.0f, 0.0f);
 
 #  if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
   if (state->guiding.use_volume_guiding) {
     const float guiding_sampling_prob = state->guiding.volume_guiding_sampling_prob;
-    const float guide_pdf = guiding_phase_pdf(kg, state, wo);
+    const float guide_pdf = guiding_phase_pdf(kg, state, omega_in);
     pdf = (guiding_sampling_prob * guide_pdf) + (1.0f - guiding_sampling_prob) * pdf;
   }
 #  endif
@@ -272,7 +272,7 @@ ccl_device int volume_shader_phase_guided_sample(KernelGlobals kg,
                                                  ccl_private const ShaderVolumeClosure *svc,
                                                  const float2 rand_phase,
                                                  ccl_private BsdfEval *phase_eval,
-                                                 ccl_private float3 *wo,
+                                                 ccl_private float3 *omega_in,
                                                  ccl_private float *phase_pdf,
                                                  ccl_private float *unguided_phase_pdf,
                                                  ccl_private float *sampled_roughness)
@@ -304,11 +304,11 @@ ccl_device int volume_shader_phase_guided_sample(KernelGlobals kg,
 
   if (sample_guiding) {
     /* Sample guiding distribution. */
-    guide_pdf = guiding_phase_sample(kg, state, rand_phase, wo);
+    guide_pdf = guiding_phase_sample(kg, state, rand_phase, omega_in);
     *phase_pdf = 0.0f;
 
     if (guide_pdf != 0.0f) {
-      *unguided_phase_pdf = volume_shader_phase_eval(kg, sd, svc, *wo, phase_eval);
+      *unguided_phase_pdf = volume_shader_phase_eval(kg, sd, svc, *omega_in, phase_eval);
       *phase_pdf = (guiding_sampling_prob * guide_pdf) +
                    ((1.0f - guiding_sampling_prob) * (*unguided_phase_pdf));
       label = LABEL_VOLUME_SCATTER;
@@ -318,14 +318,14 @@ ccl_device int volume_shader_phase_guided_sample(KernelGlobals kg,
     /* Sample phase. */
     *phase_pdf = 0.0f;
     label = volume_phase_sample(
-        sd, svc, rand_phase.x, rand_phase.y, &eval, wo, unguided_phase_pdf);
+        sd, svc, rand_phase.x, rand_phase.y, &eval, omega_in, unguided_phase_pdf);
 
     if (*unguided_phase_pdf != 0.0f) {
       bsdf_eval_init(phase_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, eval);
 
       *phase_pdf = *unguided_phase_pdf;
       if (use_volume_guiding) {
-        guide_pdf = guiding_phase_pdf(kg, state, *wo);
+        guide_pdf = guiding_phase_pdf(kg, state, *omega_in);
         *phase_pdf *= 1.0f - guiding_sampling_prob;
         *phase_pdf += guiding_sampling_prob * guide_pdf;
       }
@@ -349,7 +349,7 @@ ccl_device int volume_shader_phase_sample(KernelGlobals kg,
                                           ccl_private const ShaderVolumeClosure *svc,
                                           float2 rand_phase,
                                           ccl_private BsdfEval *phase_eval,
-                                          ccl_private float3 *wo,
+                                          ccl_private float3 *omega_in,
                                           ccl_private float *pdf,
                                           ccl_private float *sampled_roughness)
 {
@@ -357,7 +357,7 @@ ccl_device int volume_shader_phase_sample(KernelGlobals kg,
   Spectrum eval = zero_spectrum();
 
   *pdf = 0.0f;
-  int label = volume_phase_sample(sd, svc, rand_phase.x, rand_phase.y, &eval, wo, pdf);
+  int label = volume_phase_sample(sd, svc, rand_phase.x, rand_phase.y, &eval, omega_in, pdf);
 
   if (*pdf != 0.0f) {
     bsdf_eval_init(phase_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, eval);

intern/cycles/kernel/light/area.h

@@ -102,7 +102,7 @@ ccl_device float area_light_spread_attenuation(const float3 D,
     /* The factor M_PI_F comes from integrating the radiance over the hemisphere */
     return (cos_a > 0.9999997f) ? M_PI_F : 0.0f;
   }
-  const float sin_a = sin_from_cos(cos_a);
+  const float sin_a = safe_sqrtf(1.0f - sqr(cos_a));
   const float tan_a = sin_a / cos_a;
   return max((tan_half_spread - tan_a) * normalize_spread, 0.0f);
 }

intern/cycles/kernel/light/sample.h

@@ -88,6 +88,13 @@ light_sample_shader_eval(KernelGlobals kg,
   return eval;
 }
 
+/* Test if light sample is from a light or emission from geometry. */
+ccl_device_inline bool light_sample_is_light(ccl_private const LightSample *ccl_restrict ls)
+{
+  /* return if it's a lamp for shadow pass */
+  return (ls->prim == PRIM_NONE && ls->type != LIGHT_BACKGROUND);
+}
+
 /* Early path termination of shadow rays. */
 ccl_device_inline bool light_sample_terminate(KernelGlobals kg,
                                               ccl_private const LightSample *ccl_restrict ls,