Merge branch 'master' into tmp-eevee-shadow-commit

# Conflicts:
#	source/blender/gpu/shaders/infos/gpu_shader_test_info.hh

.gitea/default_merge_message/REBASE_TEMPLATE.md

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

.gitea/default_merge_message/SQUASH_TEMPLATE.md

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

.gitea/issue_template.yaml

@@ -1,45 +0,0 @@
-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

@@ -0,0 +1,44 @@
+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

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

.gitea/issue_template/design.yaml

@@ -0,0 +1,10 @@
+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

@@ -0,0 +1,10 @@
+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

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

.gitea/pull_request_template.yaml

@@ -0,0 +1,17 @@
+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,14 +167,26 @@ get_blender_version()
 option(WITH_BLENDER "Build blender (disable to build only the blender player)" ON)
 mark_as_advanced(WITH_BLENDER)
 
-if(APPLE)
+if(WIN32)
-  # In future, can be used with `quicklookthumbnailing/qlthumbnailreply` to create file
+  option(WITH_BLENDER_THUMBNAILER "\
-  # thumbnails for say Finder. Turn it off for now.
+Build \"BlendThumb.dll\" helper for Windows explorer integration to support extracting \
-  option(WITH_BLENDER_THUMBNAILER "Build \"blender-thumbnailer\" thumbnail extraction utility" OFF)
+thumbnails from `.blend` files."
-elseif(WIN32)
+    ON
-  option(WITH_BLENDER_THUMBNAILER "Build \"BlendThumb.dll\" helper for Windows explorer integration" ON)
+  )
 else()
-  option(WITH_BLENDER_THUMBNAILER "Build \"blender-thumbnailer\" thumbnail extraction utility" ON)
+  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)
 endif()
 
 option(WITH_INTERNATIONAL "Enable I18N (International fonts and text)" ON)
@@ -214,14 +226,19 @@ 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.
-  option(WITH_XR_OPENXR   "Enable VR features through the OpenXR specification" OFF)
+  set(_option_default OFF)
-  mark_as_advanced(WITH_XR_OPENXR)
-else()
-  option(WITH_XR_OPENXR   "Enable VR features through the OpenXR specification" ON)
 endif()
+option(WITH_XR_OPENXR "Enable VR features through the OpenXR specification" ${_option_default})
+if(APPLE)
+  mark_as_advanced(WITH_XR_OPENXR)
+endif()
+unset(_option_default)
+
 option(WITH_GMP "Enable features depending on GMP (Exact Boolean)" ON)
 
 # Compositor
@@ -353,11 +370,12 @@ else()
   set(WITH_COREAUDIO OFF)
 endif()
 if(NOT WIN32)
+  set(_option_default ON)
   if(APPLE)
-    option(WITH_JACK          "Enable JACK Support (http://www.jackaudio.org)" OFF)
+    set(_option_default 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)
 else()
   set(WITH_JACK OFF)
@@ -399,6 +417,26 @@ 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)
@@ -406,11 +444,6 @@ 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)
@@ -491,7 +524,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 gfx900 gfx906 gfx90c gfx902 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 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()
@@ -993,6 +1026,8 @@ 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)
@@ -1206,13 +1241,6 @@ if(WITH_OPENGL)
   add_definitions(-DWITH_OPENGL)
 endif()
 
-#-----------------------------------------------------------------------------
-# Configure Vulkan.
-
-if(WITH_VULKAN_BACKEND)
-  list(APPEND BLENDER_GL_LIBRARIES ${VULKAN_LIBRARIES})
-endif()
-
 # -----------------------------------------------------------------------------
 # Configure Metal
 
@@ -1262,12 +1290,14 @@ endif()
 # -----------------------------------------------------------------------------
 # Configure Bullet
 
-if(WITH_BULLET AND WITH_SYSTEM_BULLET)
+if(WITH_BULLET)
+  if(WITH_SYSTEM_BULLET)
     find_package(Bullet)
     set_and_warn_library_found("Bullet" BULLET_FOUND WITH_BULLET)
-else()
+  else()
     set(BULLET_INCLUDE_DIRS "${CMAKE_SOURCE_DIR}/extern/bullet2/src")
-  # set(BULLET_LIBRARIES "")
+    set(BULLET_LIBRARIES "extern_bullet")
+  endif()
 endif()
 
 

GNUmakefile

@@ -71,6 +71,13 @@ 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.
 
@@ -481,6 +488,10 @@ 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/cmake/Modules/FindMoltenVK.cmake

@@ -19,9 +19,13 @@ 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,9 +17,13 @@ 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,6 +67,8 @@ 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

@@ -0,0 +1,63 @@
+# 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

@@ -0,0 +1,63 @@
+# 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,18 +6,80 @@
 import re
 import sys
 
+from typing import Optional
+
 cmakelists_file = sys.argv[-1]
 
 
-def main():
+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:
     options = []
-    for l in open(cmakelists_file, 'r').readlines():
+    with open(cmakelists_file, 'r', encoding="utf-8") as fh:
-        if not l.lstrip().startswith('#'):
+        file_data = fh.read()
-            l_option = re.sub(r'.*\boption\s*\(\s*(WITH_[a-zA-Z0-9_]+)\s+\"(.*)\"\s*.*', r'\g<1> - \g<2>', l)
+        for m in re.finditer(r"^\s*option\s*\(\s*(WITH_[a-zA-Z0-9_]+)\s+(\")", file_data, re.MULTILINE):
-            if l_option != l:
+            option_name = m.group(1)
-                l_option = l_option.strip()
+            option_descr = extract_cmake_string_at_pos(file_data, m.span(2)[1])
-                if l_option.startswith('WITH_'):
+            if option_descr is None:
-                    options.append(l_option)
+                # Possibly a parsing error, at least show something.
+                option_descr = "(UNDOCUMENTED)"
+            options.append("{:s}: {:s}".format(option_name, option_descr))
 
     print('\n'.join(options))
 

build_files/cmake/macros.cmake

@@ -550,8 +550,10 @@ function(setup_platform_linker_libs
   endif()
 
   if(WIN32 AND NOT UNIX)
+    if(DEFINED PTHREADS_LIBRARIES)
       target_link_libraries(${target} ${PTHREADS_LIBRARIES})
     endif()
+  endif()
 
   # target_link_libraries(${target} ${PLATFORM_LINKLIBS} ${CMAKE_DL_LIBS})
   target_link_libraries(${target} ${PLATFORM_LINKLIBS})
@@ -1262,7 +1264,7 @@ endmacro()
 
 # Utility to gather and install precompiled shared libraries.
 macro(add_bundled_libraries library_dir)
-  if(EXISTS ${LIBDIR})
+  if(DEFINED LIBDIR)
     set(_library_dir ${LIBDIR}/${library_dir})
     if(WIN32)
       file(GLOB _all_library_versions ${_library_dir}/*\.dll)

build_files/cmake/platform/platform_apple.cmake

@@ -97,20 +97,8 @@ add_bundled_libraries(materialx/lib)
 
 if(WITH_VULKAN_BACKEND)
   find_package(MoltenVK REQUIRED)
-
+  find_package(ShaderC REQUIRED)
-  if(EXISTS ${LIBDIR}/vulkan)
+  find_package(Vulkan REQUIRED)
-    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,7 +1,12 @@
 # SPDX-License-Identifier: GPL-2.0-or-later
 # Copyright 2022 Blender Foundation. All rights reserved.
 
-# Auto update existing CMake caches for new libraries
+# 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()
 
 # Clear cached variables whose name matches `pattern`.
 function(unset_cache_variables pattern)

build_files/cmake/platform/platform_unix.cmake

@@ -4,7 +4,11 @@
 # Libraries configuration for any *nix system including Linux and Unix (excluding APPLE).
 
 # Detect precompiled library directory
-if(NOT DEFINED LIBDIR)
+
+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)
@@ -30,12 +34,22 @@ if(NOT DEFINED LIBDIR)
     # 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()
 endif()
 
+
 # Support restoring this value once pre-compiled libraries have been handled.
 set(WITH_STATIC_LIBS_INIT ${WITH_STATIC_LIBS})
 
-if(EXISTS ${LIBDIR})
+if(DEFINED LIBDIR)
   message(STATUS "Using pre-compiled LIBDIR: ${LIBDIR}")
 
   file(GLOB LIB_SUBDIRS ${LIBDIR}/*)
@@ -85,7 +99,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(EXISTS ${LIBDIR})
+if(DEFINED LIBDIR)
   without_system_libs_begin()
 endif()
 
@@ -97,6 +111,7 @@ 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)
@@ -114,7 +129,7 @@ endfunction()
 if(NOT WITH_SYSTEM_FREETYPE)
   # FreeType compiled with Brotli compression for woff2.
   find_package_wrapper(Freetype REQUIRED)
-  if(EXISTS ${LIBDIR})
+  if(DEFINED LIBDIR)
     find_package_wrapper(Brotli REQUIRED)
 
     # NOTE: This is done on WIN32 & APPLE but fails on some Linux systems.
@@ -141,7 +156,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(EXISTS ${LIBDIR})
+    if(DEFINED LIBDIR)
       path_is_prefix(LIBDIR PYTHON_SITE_PACKAGES _is_prefix)
       if(_is_prefix)
         message(WARNING "
@@ -217,7 +232,7 @@ if(WITH_CODEC_SNDFILE)
 endif()
 
 if(WITH_CODEC_FFMPEG)
-  if(EXISTS ${LIBDIR})
+  if(DEFINED 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.
@@ -232,7 +247,7 @@ if(WITH_CODEC_FFMPEG)
       vpx
       x264
       xvidcore)
-    if(EXISTS ${LIBDIR}/ffmpeg/lib/libaom.a)
+    if((DEFINED LIBDIR) AND (EXISTS ${LIBDIR}/ffmpeg/lib/libaom.a))
       list(APPEND FFMPEG_FIND_COMPONENTS aom)
     endif()
   elseif(FFMPEG)
@@ -430,10 +445,13 @@ 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()
@@ -451,7 +469,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)
@@ -466,7 +484,7 @@ if(WITH_OPENIMAGEDENOISE)
 endif()
 
 if(WITH_LLVM)
-  if(EXISTS ${LIBDIR})
+  if(DEFINED LIBDIR)
     set(LLVM_STATIC ON)
   endif()
 
@@ -480,7 +498,7 @@ if(WITH_LLVM)
     endif()
 
     # Symbol conflicts with same UTF library used by OpenCollada
-    if(EXISTS ${LIBDIR})
+    if(DEFINED LIBDIR)
       if(WITH_OPENCOLLADA AND (${LLVM_VERSION} VERSION_LESS "4.0.0"))
         list(REMOVE_ITEM OPENCOLLADA_LIBRARIES ${OPENCOLLADA_UTF_LIBRARY})
       endif()
@@ -536,7 +554,7 @@ if(WITH_CYCLES AND WITH_CYCLES_PATH_GUIDING)
   endif()
 endif()
 
-if(EXISTS ${LIBDIR})
+if(DEFINED LIBDIR)
   without_system_libs_end()
 endif()
 
@@ -551,9 +569,14 @@ else()
 endif()
 
 find_package(Threads REQUIRED)
-list(APPEND PLATFORM_LINKLIBS ${CMAKE_THREAD_LIBS_INIT})
+# `FindThreads` documentation notes that this may be empty
-# used by other platforms
+# with the system libraries provide threading functionality.
-set(PTHREADS_LIBRARIES ${CMAKE_THREAD_LIBS_INIT})
+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()
+
 
 if(CMAKE_DL_LIBS)
   list(APPEND PLATFORM_LINKLIBS ${CMAKE_DL_LIBS})
@@ -575,7 +598,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(EXISTS ${LIBDIR})
+if(DEFINED LIBDIR)
   # Clear the prefix path as it causes the `LIBDIR` to override system locations.
   unset(CMAKE_PREFIX_PATH)
 
@@ -631,7 +654,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(EXISTS "${LIBDIR}/wayland" AND WITH_GHOST_WAYLAND_DYNLOAD)
+  if((DEFINED LIBDIR) AND (EXISTS "${LIBDIR}/wayland" AND WITH_GHOST_WAYLAND_DYNLOAD))
     set(_use_system_wayland OFF)
   else()
     set(_use_system_wayland ON)
@@ -695,7 +718,7 @@ if(WITH_GHOST_WAYLAND)
       add_definitions(-DWITH_GHOST_WAYLAND_LIBDECOR)
     endif()
 
-    if(EXISTS "${LIBDIR}/wayland/bin/wayland-scanner")
+    if((DEFINED LIBDIR) AND (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,6 +43,10 @@ update-code:
             branch: trunk
             commit_id: HEAD
             path: lib/benchmarks
+        assets:
+            branch: trunk
+            commit_id: HEAD
+            path: lib/assets
 
 #
 # Buildbot only configs
@@ -59,7 +63,7 @@ buildbot:
     optix:
         version: '7.3.0'
     ocloc:
-        version: '101.3430'
+        version: '101.4032'
     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  # type: ignore
+import setuptools
 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):  # type: ignore
+    class BinaryDistribution(setuptools.dist.Distribution):
         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():
+def parse_arguments() -> argparse.Namespace:
     parser = argparse.ArgumentParser()
     parser.add_argument("--ctest-command", default="ctest")
     parser.add_argument("--cmake-command", default="cmake")

build_files/utils/make_update.py

@@ -104,17 +104,30 @@ 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])
 
-    # Update precompiled libraries and tests
+    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
 
     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 and Tests (one repository)")
+            print_stage("Updating Precompiled Libraries, Assets and Tests (one repository)")
             paths_local_and_remote.append((lib_dirpath, svn_url))
         else:
-            print_stage("Updating Precompiled Libraries and Tests (multiple repositories)")
+            print_stage("Updating Precompiled Libraries, Assets 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,6 +2098,8 @@ 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")
 
@@ -2124,6 +2126,8 @@ 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,11 +13,13 @@ endif()
 
 # Exporting functions from the blender binary gives linker warnings on Apple arm64 systems.
 # Silence them here.
-if(APPLE AND ("${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64"))
+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()
+  endif()
 endif()
 
 set(MANTAVERSION "0.13")
@@ -261,9 +263,11 @@ set(LIB
 
 blender_add_lib(extern_mantaflow "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")
 
-# The VDB libs above are only added to as INTERFACE libs by blender_add_lib,
+if(WITH_OPENVDB)
-# meaning extern_mantaflow itself actually does not have a dependency on the
+  # The VDB libs above are only added to as INTERFACE libs by blender_add_lib,
-# openvdb libraries, and CMAKE is free to link the vdb libs before
+  # meaning extern_mantaflow itself actually does not have a dependency on the
-# extern_mantaflow causing linker errors on linux. By explicitly declaring
+  # openvdb libraries, and CMAKE is free to link the vdb libs before
-# a dependency here, cmake will do the right thing.
+  # extern_mantaflow causing linker errors on linux. By explicitly declaring
-target_link_libraries(extern_mantaflow PRIVATE ${OPENVDB_LIBRARIES})
+  # a dependency here, cmake will do the right thing.
+  target_link_libraries(extern_mantaflow PRIVATE ${OPENVDB_LIBRARIES})
+endif()

extern/vulkan_memory_allocator/CMakeLists.txt

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

extern/vulkan_memory_allocator/patches/remove_compilation_warning.diff

@@ -0,0 +1,15 @@
+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;

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 ||
+        VMA_ASSERT((m_PassStats.allocationsMoved == m_MaxPassAllocations ||
-            m_PassStats.bytesMoved == m_MaxPassBytes && "Exceeded maximal pass threshold!");
+            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 Vega or RDNA architecture", icon='BLANK1')
+                    col.label(text="Requires AMD GPU with 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 Vega or RDNA architecture", icon='BLANK1')
+                    col.label(text="Requires AMD GPU with 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.3430 or newer", icon='BLANK1')
+                    col.label(text="and Windows driver version 101.4032 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.23904 or newer", icon='BLANK1')
+                    col.label(text="  - intel-level-zero-gpu version 1.3.24931 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')

intern/cycles/blender/light.cpp

@@ -48,6 +48,8 @@ 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,12 @@ macro(cycles_external_libraries_append libraries)
   endif()
   if(WITH_OPENIMAGEDENOISE)
     list(APPEND ${libraries} ${OPENIMAGEDENOISE_LIBRARIES})
-    if(APPLE AND "${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64")
+    if(APPLE)
+      if("${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64")
         list(APPEND ${libraries} "-framework Accelerate")
       endif()
     endif()
+  endif()
   if(WITH_ALEMBIC)
     list(APPEND ${libraries} ${ALEMBIC_LIBRARIES})
   endif()
@@ -136,7 +138,15 @@ 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}
     ${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 >= 9);
+  return (major >= 10);
 }
 
 CCL_NAMESPACE_END

intern/cycles/device/metal/device_impl.mm

@@ -327,10 +327,21 @@ 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";
+                         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; \
+    }
 
 #  include "kernel/data_template.h"
 

intern/cycles/device/metal/kernel.mm

@@ -49,6 +49,18 @@ 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};
@@ -448,13 +460,18 @@ static MTLFunctionConstantValues *GetConstantValues(KernelData const *data = nul
   if (!data) {
     data = &zero_data;
   }
-  int zero_int = 0;
+  [constant_values setConstantValue:&zero_data type:MTLDataType_int atIndex:Kernel_DummyConstant];
-  [constant_values setConstantValue:&zero_int type:MTLDataType_int atIndex:Kernel_DummyConstant];
+
+  bool next_member_is_specialized = true;
+
+#  define KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE next_member_is_specialized = false;
 
 #  define KERNEL_STRUCT_MEMBER(parent, _type, name) \
-    [constant_values setConstantValue:&data->parent.name \
+    [constant_values setConstantValue:next_member_is_specialized ? (void *)&data->parent.name : \
+                                                                   (void *)&zero_data \
                                  type:MTLDataType_##_type \
-                              atIndex:KernelData_##parent##_##name];
+                              atIndex:KernelData_##parent##_##name]; \
+    next_member_is_specialized = true;
 
 #  include "kernel/data_template.h"
 

intern/cycles/device/metal/queue.mm

@@ -278,7 +278,8 @@ int MetalDeviceQueue::num_concurrent_states(const size_t state_size) const
   if (metal_device_->device_vendor == METAL_GPU_APPLE) {
     result *= 4;
 
-    if (MetalInfo::get_apple_gpu_architecture(metal_device_->mtlDevice) == APPLE_M2) {
+    /* Increasing the state count doesn't notably benefit M1-family systems.  */
+    if (MetalInfo::get_apple_gpu_architecture(metal_device_->mtlDevice) != APPLE_M1) {
       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];

intern/cycles/device/metal/util.h

@@ -29,6 +29,7 @@ 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 APPLE_M2;
+    return get_apple_gpu_core_count(device) <= 10 ? APPLE_M2 : APPLE_M2_BIG;
   }
   return APPLE_UNKNOWN;
 }

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 allocs. */
+  /* Resize if needed. Also, in case of resize - allocate in advance for future allocations. */
   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.3430, compute-runtime version is 23904. */
+ * For Windows driver 101.4032, compute-runtime version is 24931. */
-static const int lowest_supported_driver_version_win = 1013430;
+static const int lowest_supported_driver_version_win = 1014032;
-static const int lowest_supported_driver_version_neo = 23904;
+static const int lowest_supported_driver_version_neo = 24931;
 
 int OneapiDevice::parse_driver_build_version(const sycl::device &device)
 {

intern/cycles/graph/CMakeLists.txt

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

intern/cycles/integrator/CMakeLists.txt

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

intern/cycles/kernel/CMakeLists.txt

@@ -750,7 +750,7 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
     ${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,8 +63,9 @@ 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. */
+   * This is being evaluated for all BSDFs, so path flag does not contain a specific type.
-  const uint32_t path_flag = PATH_RAY_EMISSION;
+   * 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;
   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,10 +102,9 @@ 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,
+ccl_device_inline bool bsdf_is_transmission(ccl_private const ShaderClosure *sc, const float3 wo)
-                                            const float3 omega_in)
 {
-  return dot(sc->N, omega_in) < 0.0f;
+  return dot(sc->N, wo) < 0.0f;
 }
 
 ccl_device_inline int bsdf_sample(KernelGlobals kg,
@@ -114,7 +113,7 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
                                   float randu,
                                   float randv,
                                   ccl_private Spectrum *eval,
-                                  ccl_private float3 *omega_in,
+                                  ccl_private float3 *wo,
                                   ccl_private float *pdf,
                                   ccl_private float2 *sampled_roughness,
                                   ccl_private float *eta)
@@ -126,43 +125,43 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
 
   switch (sc->type) {
     case CLOSURE_BSDF_DIFFUSE_ID:
-      label = bsdf_diffuse_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
+      label = bsdf_diffuse_sample(sc, Ng, sd->wi, randu, randv, eval, wo, 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->I, randu, randv, eval, omega_in, pdf);
+      label = bsdf_oren_nayar_sample(sc, Ng, sd->wi, randu, randv, eval, wo, 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->I, randu, randv, eval, omega_in, pdf, sampled_roughness);
+          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_DIFFUSE_RAMP_ID:
-      label = bsdf_diffuse_ramp_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
+      label = bsdf_diffuse_ramp_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
 #  endif
     case CLOSURE_BSDF_TRANSLUCENT_ID:
-      label = bsdf_translucent_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
+      label = bsdf_translucent_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_REFLECTION_ID:
-      label = bsdf_reflection_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, eta);
+      label = bsdf_reflection_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf, eta);
       *sampled_roughness = zero_float2();
       break;
     case CLOSURE_BSDF_REFRACTION_ID:
-      label = bsdf_refraction_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, eta);
+      label = bsdf_refraction_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf, eta);
       *sampled_roughness = zero_float2();
       break;
     case CLOSURE_BSDF_TRANSPARENT_ID:
-      label = bsdf_transparent_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
+      label = bsdf_transparent_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
       *sampled_roughness = zero_float2();
       *eta = 1.0f;
       break;
@@ -171,85 +170,65 @@ 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(
-          kg, sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, sampled_roughness, eta);
+          sc, Ng, sd->wi, randu, randv, eval, wo, 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,
+      label = bsdf_microfacet_multi_ggx_sample(
-                                               sc,
+          kg, sc, Ng, sd->wi, randu, randv, eval, wo, pdf, &sd->lcg_state, sampled_roughness, eta);
-                                               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,
+      label = bsdf_microfacet_multi_ggx_glass_sample(
-                                                     sc,
+          kg, sc, Ng, sd->wi, randu, randv, eval, wo, pdf, &sd->lcg_state, sampled_roughness, eta);
-                                                     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(
-          kg, sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, sampled_roughness, eta);
+          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness, eta);
       break;
     case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
       label = bsdf_ashikhmin_shirley_sample(
-          sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, sampled_roughness);
+          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_ASHIKHMIN_VELVET_ID:
-      label = bsdf_ashikhmin_velvet_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
+      label = bsdf_ashikhmin_velvet_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_DIFFUSE_TOON_ID:
-      label = bsdf_diffuse_toon_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
+      label = bsdf_diffuse_toon_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_GLOSSY_TOON_ID:
-      label = bsdf_glossy_toon_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
+      label = bsdf_glossy_toon_sample(sc, Ng, sd->wi, randu, randv, eval, wo, 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->I, randu, randv, eval, omega_in, pdf, sampled_roughness);
+          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_HAIR_TRANSMISSION_ID:
       label = bsdf_hair_transmission_sample(
-          sc, Ng, sd->I, randu, randv, eval, omega_in, pdf, sampled_roughness);
+          sc, Ng, sd->wi, randu, randv, eval, wo, pdf, sampled_roughness);
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_HAIR_PRINCIPLED_ID:
       label = bsdf_principled_hair_sample(
-          kg, sc, sd, randu, randv, eval, omega_in, pdf, sampled_roughness, eta);
+          kg, sc, sd, randu, randv, eval, wo, pdf, sampled_roughness, eta);
       break;
     case CLOSURE_BSDF_PRINCIPLED_DIFFUSE_ID:
-      label = bsdf_principled_diffuse_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
+      label = bsdf_principled_diffuse_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
     case CLOSURE_BSDF_PRINCIPLED_SHEEN_ID:
-      label = bsdf_principled_sheen_sample(sc, Ng, sd->I, randu, randv, eval, omega_in, pdf);
+      label = bsdf_principled_sheen_sample(sc, Ng, sd->wi, randu, randv, eval, wo, pdf);
       *sampled_roughness = one_float2();
       *eta = 1.0f;
       break;
@@ -274,12 +253,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 cosNI = dot(*omega_in, sc->N);
+      const float cosNO = dot(*wo, sc->N);
-      *eval *= shift_cos_in(cosNI, frequency_multiplier);
+      *eval *= shift_cos_in(cosNO, frequency_multiplier);
     }
     if (label & LABEL_DIFFUSE) {
       if (!isequal(sc->N, sd->N)) {
-        *eval *= bump_shadowing_term(sd->N, sc->N, *omega_in);
+        *eval *= bump_shadowing_term(sd->N, sc->N, *wo);
       }
     }
   }
@@ -426,7 +405,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 omega_in)
+                                 const float3 wo)
 {
   /* For curves use the smooth normal, particularly for ribbons the geometric
    * normal gives too much darkening otherwise. */
@@ -482,7 +461,7 @@ 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, omega_in)) ? LABEL_TRANSMIT | LABEL_GLOSSY :
+      label = (bsdf_is_transmission(sc, wo)) ? LABEL_TRANSMIT | LABEL_GLOSSY :
                                                LABEL_REFLECT | LABEL_GLOSSY;
       break;
     case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
@@ -504,7 +483,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, omega_in))
+      if (bsdf_is_transmission(sc, wo))
         label = LABEL_TRANSMIT | LABEL_GLOSSY;
       else
         label = LABEL_REFLECT | LABEL_GLOSSY;
@@ -543,83 +522,83 @@ ccl_device_inline
     bsdf_eval(KernelGlobals kg,
               ccl_private ShaderData *sd,
               ccl_private const ShaderClosure *sc,
-              const float3 omega_in,
+              const float3 wo,
               ccl_private float *pdf)
 {
   Spectrum eval = zero_spectrum();
 
   switch (sc->type) {
     case CLOSURE_BSDF_DIFFUSE_ID:
-      eval = bsdf_diffuse_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_diffuse_eval(sc, sd->wi, wo, pdf);
       break;
 #if defined(__SVM__) || defined(__OSL__)
     case CLOSURE_BSDF_OREN_NAYAR_ID:
-      eval = bsdf_oren_nayar_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_oren_nayar_eval(sc, sd->wi, wo, pdf);
       break;
 #  ifdef __OSL__
     case CLOSURE_BSDF_PHONG_RAMP_ID:
-      eval = bsdf_phong_ramp_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_phong_ramp_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_DIFFUSE_RAMP_ID:
-      eval = bsdf_diffuse_ramp_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_diffuse_ramp_eval(sc, sd->wi, wo, pdf);
       break;
 #  endif
     case CLOSURE_BSDF_TRANSLUCENT_ID:
-      eval = bsdf_translucent_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_translucent_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_REFLECTION_ID:
-      eval = bsdf_reflection_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_reflection_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_REFRACTION_ID:
-      eval = bsdf_refraction_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_refraction_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_TRANSPARENT_ID:
-      eval = bsdf_transparent_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_transparent_eval(sc, sd->wi, wo, 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->I, omega_in, pdf);
+      eval = bsdf_microfacet_ggx_eval(sc, sd->N, sd->wi, wo, 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->I, omega_in, pdf, &sd->lcg_state);
+      eval = bsdf_microfacet_multi_ggx_eval(sc, sd->N, sd->wi, wo, 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->I, omega_in, pdf, &sd->lcg_state);
+      eval = bsdf_microfacet_multi_ggx_glass_eval(sc, sd->wi, wo, 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->I, omega_in, pdf);
+      eval = bsdf_microfacet_beckmann_eval(sc, sd->N, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
-      eval = bsdf_ashikhmin_shirley_eval(sc, sd->N, sd->I, omega_in, pdf);
+      eval = bsdf_ashikhmin_shirley_eval(sc, sd->N, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_ASHIKHMIN_VELVET_ID:
-      eval = bsdf_ashikhmin_velvet_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_ashikhmin_velvet_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_DIFFUSE_TOON_ID:
-      eval = bsdf_diffuse_toon_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_diffuse_toon_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_GLOSSY_TOON_ID:
-      eval = bsdf_glossy_toon_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_glossy_toon_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_HAIR_PRINCIPLED_ID:
-      eval = bsdf_principled_hair_eval(kg, sd, sc, omega_in, pdf);
+      eval = bsdf_principled_hair_eval(kg, sd, sc, wo, pdf);
       break;
     case CLOSURE_BSDF_HAIR_REFLECTION_ID:
-      eval = bsdf_hair_reflection_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_hair_reflection_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_HAIR_TRANSMISSION_ID:
-      eval = bsdf_hair_transmission_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_hair_transmission_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_PRINCIPLED_DIFFUSE_ID:
-      eval = bsdf_principled_diffuse_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_principled_diffuse_eval(sc, sd->wi, wo, pdf);
       break;
     case CLOSURE_BSDF_PRINCIPLED_SHEEN_ID:
-      eval = bsdf_principled_sheen_eval(sc, sd->I, omega_in, pdf);
+      eval = bsdf_principled_sheen_eval(sc, sd->wi, wo, pdf);
       break;
 #endif
     default:
@@ -628,7 +607,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, omega_in);
+      eval *= bump_shadowing_term(sd->N, sc->N, wo);
     }
   }
 
@@ -636,9 +615,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 cosNI = dot(omega_in, sc->N);
+    const float cosNO = dot(wo, sc->N);
-    if (cosNI >= 0.0f) {
+    if (cosNO >= 0.0f) {
-      eval *= shift_cos_in(cosNI, frequency_multiplier);
+      eval *= shift_cos_in(cosNO, frequency_multiplier);
     }
   }
 
@@ -682,4 +661,37 @@ 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 I,
+                                                            const float3 wi,
-                                                            const float3 omega_in,
+                                                            const float3 wo,
                                                             ccl_private float *pdf)
 {
   ccl_private const MicrofacetBsdf *bsdf = (ccl_private const MicrofacetBsdf *)sc;
-  const float cosNgI = dot(Ng, omega_in);
+  const float cosNgO = dot(Ng, wo);
   float3 N = bsdf->N;
 
-  float NdotI = dot(N, I);        /* in Cycles/OSL convention I is omega_out */
+  float NdotI = dot(N, wi);
-  float NdotO = dot(N, omega_in); /* and consequently we use for O omaga_in ;) */
+  float NdotO = dot(N, wo);
 
   float out = 0.0f;
 
-  if ((cosNgI < 0.0f) || fmaxf(bsdf->alpha_x, bsdf->alpha_y) <= 1e-4f ||
+  if ((cosNgO < 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(omega_in + I);
+  float3 H = normalize(wi + wo);
-  float HdotI = fmaxf(fabsf(dot(H, I)), 1e-6f);
+  float HdotI = fmaxf(fabsf(dot(H, wi)), 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(NdotO, NdotI)));
+  float pump = 1.0f / fmaxf(1e-6f, (HdotI * fmaxf(NdotI, NdotO)));
   /* pump from d-brdf paper */
-  /*float pump = 1.0f / fmaxf(1e-4f, ((NdotO + NdotI) * (NdotO*NdotI))); */
+  /*float pump = 1.0f / fmaxf(1e-4f, ((NdotI + NdotO) * (NdotI * NdotO))); */
 
   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 I,
+                                             float3 wi,
                                              float randu,
                                              float randv,
                                              ccl_private Spectrum *eval,
-                                             ccl_private float3 *omega_in,
+                                             ccl_private float3 *wo,
                                              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, I);
+  float NdotI = dot(N, wi);
   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, I);
+  float HdotI = dot(H, wi);
   if (HdotI < 0.0f)
     H = -H;
 
-  /* reflect I on H to get omega_in */
+  /* reflect wi on H to get wo */
-  *omega_in = -I + (2.0f * HdotI) * H;
+  *wo = -wi + (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, I, *omega_in, pdf);
+    *eval = bsdf_ashikhmin_shirley_eval(sc, N, wi, *wo, 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 I,
+                                               const float3 wi,
-                                               const float3 omega_in,
+                                               const float3 wo,
                                                ccl_private float *pdf)
 {
   ccl_private const VelvetBsdf *bsdf = (ccl_private const VelvetBsdf *)sc;
   float m_invsigma2 = bsdf->invsigma2;
   float3 N = bsdf->N;
 
-  float cosNO = dot(N, I);
+  float cosNI = dot(N, wi);
-  float cosNI = dot(N, omega_in);
+  float cosNO = dot(N, wo);
-  if (!(cosNO > 0 && cosNI > 0)) {
+  if (!(cosNI > 0 && cosNO > 0)) {
     *pdf = 0.0f;
     return zero_spectrum();
   }
 
-  float3 H = normalize(omega_in + I);
+  float3 H = normalize(wi + wo);
 
   float cosNH = dot(N, H);
-  float cosHO = fabsf(dot(I, H));
+  float cosHI = fabsf(dot(wi, H));
 
-  if (!(fabsf(cosNH) < 1.0f - 1e-5f && cosHO > 1e-5f)) {
+  if (!(fabsf(cosNH) < 1.0f - 1e-5f && cosHI > 1e-5f)) {
     *pdf = 0.0f;
     return zero_spectrum();
   }
-  float cosNHdivHO = cosNH / cosHO;
+  float cosNHdivHI = cosNH / cosHI;
-  cosNHdivHO = fmaxf(cosNHdivHO, 1e-5f);
+  cosNHdivHI = fmaxf(cosNHdivHI, 1e-5f);
 
-  float fac1 = 2 * fabsf(cosNHdivHO * cosNO);
+  float fac1 = 2 * fabsf(cosNHdivHI * cosNI);
-  float fac2 = 2 * fabsf(cosNHdivHO * cosNI);
+  float fac2 = 2 * fabsf(cosNHdivHI * cosNO);
 
   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) / cosNO;
+  float out = 0.25f * (D * G) / cosNI;
 
   *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 I,
+                                            float3 wi,
                                             float randu,
                                             float randv,
                                             ccl_private Spectrum *eval,
-                                            ccl_private float3 *omega_in,
+                                            ccl_private float3 *wo,
                                             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, omega_in, pdf);
+  sample_uniform_hemisphere(N, randu, randv, wo, pdf);
 
-  if (!(dot(Ng, *omega_in) > 0)) {
+  if (!(dot(Ng, *wo) > 0)) {
     *pdf = 0.0f;
     *eval = zero_spectrum();
     return LABEL_NONE;
   }
 
-  float3 H = normalize(*omega_in + I);
+  float3 H = normalize(wi + *wo);
 
-  float cosNI = dot(N, *omega_in);
+  float cosNI = dot(N, wi);
-  float cosNO = dot(N, I);
+  float cosNO = dot(N, *wo);
+  float cosHI = fabsf(dot(wi, H));
   float cosNH = dot(N, H);
-  float cosHO = fabsf(dot(I, H));
 
-  if (!(fabsf(cosNO) > 1e-5f && fabsf(cosNH) < 1.0f - 1e-5f && cosHO > 1e-5f)) {
+  if (!(fabsf(cosNI) > 1e-5f && fabsf(cosNH) < 1.0f - 1e-5f && cosHI > 1e-5f)) {
     *pdf = 0.0f;
     *eval = zero_spectrum();
     return LABEL_NONE;
   }
 
-  float cosNHdivHO = cosNH / cosHO;
+  float cosNHdivHI = cosNH / cosHI;
-  cosNHdivHO = fmaxf(cosNHdivHO, 1e-5f);
+  cosNHdivHI = fmaxf(cosNHdivHI, 1e-5f);
 
-  float fac1 = 2 * fabsf(cosNHdivHO * cosNO);
+  float fac1 = 2 * fabsf(cosNHdivHI * cosNI);
-  float fac2 = 2 * fabsf(cosNHdivHO * cosNI);
+  float fac2 = 2 * fabsf(cosNHdivHI * cosNO);
 
   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) / cosNO;
+  float power = 0.25f * (D * G) / cosNI;
 
   *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 I,
+                                      const float3 wi,
-                                      const float3 omega_in,
+                                      const float3 wo,
                                       ccl_private float *pdf)
 {
   ccl_private const DiffuseBsdf *bsdf = (ccl_private const DiffuseBsdf *)sc;
   float3 N = bsdf->N;
 
-  float cos_pi = fmaxf(dot(N, omega_in), 0.0f) * M_1_PI_F;
+  float cosNO = fmaxf(dot(N, wo), 0.0f) * M_1_PI_F;
-  *pdf = cos_pi;
+  *pdf = cosNO;
-  return make_spectrum(cos_pi);
+  return make_spectrum(cosNO);
 }
 
 ccl_device int bsdf_diffuse_sample(ccl_private const ShaderClosure *sc,
                                    float3 Ng,
-                                   float3 I,
+                                   float3 wi,
                                    float randu,
                                    float randv,
                                    ccl_private Spectrum *eval,
-                                   ccl_private float3 *omega_in,
+                                   ccl_private float3 *wo,
                                    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, omega_in, pdf);
+  sample_cos_hemisphere(N, randu, randv, wo, pdf);
 
-  if (dot(Ng, *omega_in) > 0.0f) {
+  if (dot(Ng, *wo) > 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 I,
+                                          const float3 wi,
-                                          const float3 omega_in,
+                                          const float3 wo,
                                           ccl_private float *pdf)
 {
   ccl_private const DiffuseBsdf *bsdf = (ccl_private const DiffuseBsdf *)sc;
   float3 N = bsdf->N;
 
-  float cos_pi = fmaxf(-dot(N, omega_in), 0.0f) * M_1_PI_F;
+  float cosNO = fmaxf(-dot(N, wo), 0.0f) * M_1_PI_F;
-  *pdf = cos_pi;
+  *pdf = cosNO;
-  return make_spectrum(cos_pi);
+  return make_spectrum(cosNO);
 }
 
 ccl_device int bsdf_translucent_sample(ccl_private const ShaderClosure *sc,
                                        float3 Ng,
-                                       float3 I,
+                                       float3 wi,
                                        float randu,
                                        float randv,
                                        ccl_private Spectrum *eval,
-                                       ccl_private float3 *omega_in,
+                                       ccl_private float3 *wo,
                                        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, omega_in, pdf);
+  sample_cos_hemisphere(-N, randu, randv, wo, pdf);
-  if (dot(Ng, *omega_in) < 0) {
+  if (dot(Ng, *wo) < 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 I,
+                                           const float3 wi,
-                                           const float3 omega_in,
+                                           const float3 wo,
                                            ccl_private float *pdf)
 {
   const DiffuseRampBsdf *bsdf = (const DiffuseRampBsdf *)sc;
   float3 N = bsdf->N;
 
-  float cos_pi = fmaxf(dot(N, omega_in), 0.0f);
+  float cosNO = fmaxf(dot(N, wo), 0.0f);
-  if (cos_pi >= 0.0f) {
+  if (cosNO >= 0.0f) {
-    *pdf = cos_pi * M_1_PI_F;
+    *pdf = cosNO * M_1_PI_F;
-    return rgb_to_spectrum(bsdf_diffuse_ramp_get_color(bsdf->colors, cos_pi) * M_1_PI_F);
+    return rgb_to_spectrum(bsdf_diffuse_ramp_get_color(bsdf->colors, cosNO) * 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 I,
+                                        float3 wi,
                                         float randu,
                                         float randv,
                                         ccl_private Spectrum *eval,
-                                        ccl_private float3 *omega_in,
+                                        ccl_private float3 *wo,
                                         ccl_private float *pdf)
 {
   const DiffuseRampBsdf *bsdf = (const DiffuseRampBsdf *)sc;
   float3 N = bsdf->N;
 
   // distribution over the hemisphere
-  sample_cos_hemisphere(N, randu, randv, omega_in, pdf);
+  sample_cos_hemisphere(N, randu, randv, wo, pdf);
 
-  if (dot(Ng, *omega_in) > 0.0f) {
+  if (dot(Ng, *wo) > 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 I,
+                                              const float3 wi,
-                                              const float3 omega_in,
+                                              const float3 wo,
                                               ccl_private float *pdf)
 {
   ccl_private const HairBsdf *bsdf = (ccl_private const HairBsdf *)sc;
-  if (dot(bsdf->N, omega_in) < 0.0f) {
+  if (dot(bsdf->N, wo) < 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, I);
+  float Iz = dot(Tg, wi);
-  float3 locy = normalize(I - Tg * Iz);
+  float3 locy = normalize(wi - Tg * Iz);
 
   float theta_r = M_PI_2_F - fast_acosf(Iz);
 
-  float omega_in_z = dot(Tg, omega_in);
+  float wo_z = dot(Tg, wo);
-  float3 omega_in_y = normalize(omega_in - Tg * omega_in_z);
+  float3 wo_y = normalize(wo - Tg * wo_z);
 
-  float theta_i = M_PI_2_F - fast_acosf(omega_in_z);
+  float theta_i = M_PI_2_F - fast_acosf(wo_z);
-  float cosphi_i = dot(omega_in_y, locy);
+  float cosphi_i = dot(wo_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 I,
+                                                const float3 wi,
-                                                const float3 omega_in,
+                                                const float3 wo,
                                                 ccl_private float *pdf)
 {
   ccl_private const HairBsdf *bsdf = (ccl_private const HairBsdf *)sc;
-  if (dot(bsdf->N, omega_in) >= 0.0f) {
+  if (dot(bsdf->N, wo) >= 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, I);
+  float Iz = dot(Tg, wi);
-  float3 locy = normalize(I - Tg * Iz);
+  float3 locy = normalize(wi - Tg * Iz);
 
   float theta_r = M_PI_2_F - fast_acosf(Iz);
 
-  float omega_in_z = dot(Tg, omega_in);
+  float wo_z = dot(Tg, wo);
-  float3 omega_in_y = normalize(omega_in - Tg * omega_in_z);
+  float3 wo_y = normalize(wo - Tg * wo_z);
 
-  float theta_i = M_PI_2_F - fast_acosf(omega_in_z);
+  float theta_i = M_PI_2_F - fast_acosf(wo_z);
-  float phi_i = fast_acosf(dot(omega_in_y, locy));
+  float phi_i = fast_acosf(dot(wo_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 I,
+                                           float3 wi,
                                            float randu,
                                            float randv,
                                            ccl_private Spectrum *eval,
-                                           ccl_private float3 *omega_in,
+                                           ccl_private float3 *wo,
                                            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, I);
+  float Iz = dot(Tg, wi);
-  float3 locy = normalize(I - Tg * Iz);
+  float3 locy = normalize(wi - 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);
-  *omega_in = (cosphi * costheta_i) * locy - (sinphi * costheta_i) * locx + (sintheta_i)*Tg;
+  *wo = (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 I,
+                                             float3 wi,
                                              float randu,
                                              float randv,
                                              ccl_private Spectrum *eval,
-                                             ccl_private float3 *omega_in,
+                                             ccl_private float3 *wo,
                                              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, I);
+  float Iz = dot(Tg, wi);
-  float3 locy = normalize(I - Tg * Iz);
+  float3 locy = normalize(wi - 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);
-  *omega_in = (cosphi * costheta_i) * locy - (sinphi * costheta_i) * locx + (sintheta_i)*Tg;
+  *wo = (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, *omega_in) < 1e-4f);
+  kernel_assert(dot(locy, *wo) < 1e-4f);
 
   return LABEL_TRANSMIT | LABEL_GLOSSY;
 }

intern/cycles/kernel/closure/bsdf_hair_principled.h

@@ -41,11 +41,6 @@ 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)
@@ -179,7 +174,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->I));
+  float3 Y = safe_normalize(cross(X, sd->wi));
   float3 Z = safe_normalize(cross(X, Y));
 
   /* h -1..0..1 means the rays goes from grazing the hair, to hitting it at
@@ -259,7 +254,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 omega_in,
+                                              const float3 wo,
                                               ccl_private float *pdf)
 {
   kernel_assert(isfinite_safe(sd->P) && isfinite_safe(sd->ray_length));
@@ -271,12 +266,13 @@ 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));
 
-  const float3 wo = make_float3(dot(sd->I, X), dot(sd->I, Y), dot(sd->I, Z));
+  /* local_I is the illumination direction. */
-  const float3 wi = make_float3(dot(omega_in, X), dot(omega_in, Y), dot(omega_in, Z));
+  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 float sin_theta_o = wo.x;
+  const float sin_theta_o = local_O.x;
   const float cos_theta_o = cos_from_sin(sin_theta_o);
-  const float phi_o = atan2f(wo.z, wo.y);
+  const float phi_o = atan2f(local_O.z, local_O.y);
 
   const float sin_theta_t = sin_theta_o / bsdf->eta;
   const float cos_theta_t = cos_from_sin(sin_theta_t);
@@ -295,9 +291,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 = wi.x;
+  const float sin_theta_i = local_I.x;
   const float cos_theta_i = cos_from_sin(sin_theta_i);
-  const float phi_i = atan2f(wi.z, wi.y);
+  const float phi_i = atan2f(local_I.z, local_I.y);
 
   const float phi = phi_i - phi_o;
 
@@ -343,7 +339,7 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
                                            float randu,
                                            float randv,
                                            ccl_private Spectrum *eval,
-                                           ccl_private float3 *omega_in,
+                                           ccl_private float3 *wo,
                                            ccl_private float *pdf,
                                            ccl_private float2 *sampled_roughness,
                                            ccl_private float *eta)
@@ -359,16 +355,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 wo = make_float3(dot(sd->I, X), dot(sd->I, Y), dot(sd->I, Z));
+  const float3 local_O = make_float3(dot(sd->wi, X), dot(sd->wi, Y), dot(sd->wi, 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 = wo.x;
+  const float sin_theta_o = local_O.x;
   const float cos_theta_o = cos_from_sin(sin_theta_o);
-  const float phi_o = atan2f(wo.z, wo.y);
+  const float phi_o = atan2f(local_O.z, local_O.y);
 
   const float sin_theta_t = sin_theta_o / bsdf->eta;
   const float cos_theta_t = cos_from_sin(sin_theta_t);
@@ -458,7 +454,7 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
   *eval = F;
   *pdf = F_energy;
 
-  *omega_in = X * sin_theta_i + Y * cos_theta_i * cosf(phi_i) + Z * cos_theta_i * sinf(phi_i);
+  *wo = 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 = safe_sqrtf(1.0f - cosI * cosI);
+  const float sinI = sin_from_cos(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_fresnel_color(sd, bsdf);
+  bsdf_microfacet_adjust_weight(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 I,
+                                                   const float3 wi,
-                                                   const float3 omega_in,
+                                                   const float3 wo,
                                                    ccl_private float *pdf,
                                                    ccl_private uint *lcg_state)
 {
   ccl_private const MicrofacetBsdf *bsdf = (ccl_private const MicrofacetBsdf *)sc;
-  const float cosNgI = dot(Ng, omega_in);
+  const float cosNgO = dot(Ng, wo);
 
-  if ((cosNgI < 0.0f) || bsdf->alpha_x * bsdf->alpha_y < 1e-7f) {
+  if ((cosNgO < 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, I) <= 0.0f || dot(Z, omega_in) <= 0.0f) {
+  if (dot(Z, wi) <= 0.0f || dot(Z, wo) <= 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 localI = make_float3(dot(I, X), dot(I, Y), dot(I, Z));
+  float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
-  float3 localO = make_float3(dot(omega_in, X), dot(omega_in, Y), dot(omega_in, Z));
+  float3 local_O = make_float3(dot(wo, X), dot(wo, Y), dot(wo, Z));
 
   if (is_aniso)
-    *pdf = mf_ggx_aniso_pdf(localI, localO, make_float2(bsdf->alpha_x, bsdf->alpha_y));
+    *pdf = mf_ggx_aniso_pdf(local_I, local_O, make_float2(bsdf->alpha_x, bsdf->alpha_y));
   else
-    *pdf = mf_ggx_pdf(localI, localO, bsdf->alpha_x);
+    *pdf = mf_ggx_pdf(local_I, local_O, bsdf->alpha_x);
 
   if (*pdf <= 0.f) {
     *pdf = 0.f;
     return make_float3(0.f, 0.f, 0.f);
   }
 
-  return mf_eval_glossy(localI,
+  return mf_eval_glossy(local_I,
-                        localO,
+                        local_O,
                         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 I,
+                                                float3 wi,
                                                 float randu,
                                                 float randv,
                                                 ccl_private Spectrum *eval,
-                                                ccl_private float3 *omega_in,
+                                                ccl_private float3 *wo,
                                                 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, I) <= 0.0f) {
+  if (dot(Z, wi) <= 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) {
-    *omega_in = 2 * dot(Z, I) * Z - I;
+    *wo = 2 * dot(Z, wi) * Z - wi;
-    if (dot(Ng, *omega_in) <= 0.0f) {
+    if (dot(Ng, *wo) <= 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 localI = make_float3(dot(I, X), dot(I, Y), dot(I, Z));
+  float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
-  float3 localO;
+  float3 local_O;
 
-  *eval = mf_sample_glossy(localI,
+  *eval = mf_sample_glossy(local_I,
-                           &localO,
+                           &local_O,
                            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);
-  *omega_in = X * localO.x + Y * localO.y + Z * localO.z;
+  *wo = X * local_O.x + Y * local_O.y + Z * local_O.z;
 
   /* Ensure that the light direction is on the outside w.r.t. the geometry normal. */
-  if (dot(Ng, *omega_in) <= 0.0f) {
+  if (dot(Ng, *wo) <= 0.0f) {
     *pdf = 0.0f;
     return LABEL_NONE;
   }
 
   if (is_aniso)
-    *pdf = mf_ggx_aniso_pdf(localI, localO, make_float2(bsdf->alpha_x, bsdf->alpha_y));
+    *pdf = mf_ggx_aniso_pdf(local_I, local_O, make_float2(bsdf->alpha_x, bsdf->alpha_y));
   else
-    *pdf = mf_ggx_pdf(localI, localO, bsdf->alpha_x);
+    *pdf = mf_ggx_pdf(local_I, local_O, 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_fresnel_color(sd, bsdf);
+  bsdf_microfacet_adjust_weight(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 I,
+                                                         const float3 wi,
-                                                         const float3 omega_in,
+                                                         const float3 wo,
                                                          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 localI = make_float3(dot(I, X), dot(I, Y), dot(I, Z));
+  float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
-  float3 localO = make_float3(dot(omega_in, X), dot(omega_in, Y), dot(omega_in, Z));
+  float3 local_O = make_float3(dot(wo, X), dot(wo, Y), dot(wo, Z));
 
-  const bool is_transmission = localO.z < 0.0f;
+  const bool is_transmission = local_O.z < 0.0f;
   const bool use_fresnel = !is_transmission &&
                            (bsdf->type == CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_FRESNEL_ID);
 
-  *pdf = mf_glass_pdf(localI, localO, bsdf->alpha_x, bsdf->ior);
+  *pdf = mf_glass_pdf(local_I, local_O, bsdf->alpha_x, bsdf->ior);
   kernel_assert(*pdf >= 0.f);
-  return mf_eval_glass(localI,
+  return mf_eval_glass(local_I,
-                       localO,
+                       local_O,
                        !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 I,
+                                                      float3 wi,
                                                       float randu,
                                                       float randv,
                                                       ccl_private Spectrum *eval,
-                                                      ccl_private float3 *omega_in,
+                                                      ccl_private float3 *wo,
                                                       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, I, &R, &T, &inside);
+    float fresnel = fresnel_dielectric(bsdf->ior, Z, wi, &R, &T, &inside);
 
     *pdf = 1e6f;
     *eval = make_spectrum(1e6f);
     if (randu < fresnel) {
-      *omega_in = R;
+      *wo = R;
       return LABEL_REFLECT | LABEL_SINGULAR;
     }
     else {
-      *omega_in = T;
+      *wo = 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 localI = make_float3(dot(I, X), dot(I, Y), dot(I, Z));
+  float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
-  float3 localO;
+  float3 local_O;
 
-  *eval = mf_sample_glass(localI,
+  *eval = mf_sample_glass(local_I,
-                          &localO,
+                          &local_O,
                           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(localI, localO, bsdf->alpha_x, bsdf->ior);
+  *pdf = mf_glass_pdf(local_I, local_O, bsdf->alpha_x, bsdf->ior);
   kernel_assert(*pdf >= 0.f);
   *eval *= *pdf;
 
-  *omega_in = X * localO.x + Y * localO.y + Z * localO.z;
+  *wo = X * local_O.x + Y * local_O.y + Z * local_O.z;
-  if (localO.z * localI.z > 0.0f) {
+  if (local_O.z * local_I.z > 0.0f) {
     return LABEL_REFLECT | LABEL_GLOSSY;
   }
   else {

intern/cycles/kernel/closure/bsdf_microfacet_multi_impl.h

@@ -73,9 +73,8 @@ 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, F0, cspec0);
+    throughput = interpolate_fresnel_color(wi, wh, eta, cspec0);
 
     eval *= throughput;
   }
@@ -144,11 +143,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, F0, cspec0);
+        throughput *= interpolate_fresnel_color(wi_prev, wm, eta, cspec0);
       }
 #else /* MF_MULTI_GLOSSY */
       if (use_fresnel && order > 0) {
-        throughput *= interpolate_fresnel_color(-wr, wm, eta, F0, cspec0);
+        throughput *= interpolate_fresnel_color(-wr, wm, eta, cspec0);
       }
       wr = mf_sample_phase_glossy(-wr, &throughput, wm);
 #endif
@@ -192,8 +191,6 @@ 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. */
@@ -229,22 +226,12 @@ ccl_device_forceinline Spectrum MF_FUNCTION_FULL_NAME(mf_sample)(float3 wi,
         throughput *= color;
       }
       else {
-        Spectrum t_color = interpolate_fresnel_color(wi_prev, wm, eta, F0, cspec0);
+        throughput *= interpolate_fresnel_color(wi_prev, wm, eta, cspec0);
-
-        if (order == 0)
-          throughput = t_color;
-        else
-          throughput *= t_color;
       }
     }
 #else /* MF_MULTI_GLOSSY */
     if (use_fresnel) {
-      Spectrum t_color = interpolate_fresnel_color(-wr, wm, eta, F0, cspec0);
+      throughput *= interpolate_fresnel_color(-wr, wm, eta, 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 I,
+                                         const float3 wi,
-                                         const float3 omega_in,
+                                         const float3 wo,
                                          ccl_private float *pdf)
 {
   ccl_private const OrenNayarBsdf *bsdf = (ccl_private const OrenNayarBsdf *)sc;
-  if (dot(bsdf->N, omega_in) > 0.0f) {
+  if (dot(bsdf->N, wo) > 0.0f) {
     *pdf = 0.5f * M_1_PI_F;
-    return bsdf_oren_nayar_get_intensity(sc, bsdf->N, I, omega_in);
+    return bsdf_oren_nayar_get_intensity(sc, bsdf->N, wi, wo);
   }
   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 I,
+                                      float3 wi,
                                       float randu,
                                       float randv,
                                       ccl_private Spectrum *eval,
-                                      ccl_private float3 *omega_in,
+                                      ccl_private float3 *wo,
                                       ccl_private float *pdf)
 {
   ccl_private const OrenNayarBsdf *bsdf = (ccl_private const OrenNayarBsdf *)sc;
-  sample_uniform_hemisphere(bsdf->N, randu, randv, omega_in, pdf);
+  sample_uniform_hemisphere(bsdf->N, randu, randv, wo, pdf);
 
-  if (dot(Ng, *omega_in) > 0.0f) {
+  if (dot(Ng, *wo) > 0.0f) {
-    *eval = bsdf_oren_nayar_get_intensity(sc, bsdf->N, I, *omega_in);
+    *eval = bsdf_oren_nayar_get_intensity(sc, bsdf->N, wi, *wo);
   }
   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 I,
+                                         const float3 wi,
-                                         const float3 omega_in,
+                                         const float3 wo,
                                          ccl_private float *pdf)
 {
   ccl_private const PhongRampBsdf *bsdf = (ccl_private const PhongRampBsdf *)sc;
   float m_exponent = bsdf->exponent;
-  float cosNI = dot(bsdf->N, omega_in);
+  float cosNI = dot(bsdf->N, wi);
-  float cosNO = dot(bsdf->N, I);
+  float cosNO = dot(bsdf->N, wo);
 
   if (cosNI > 0 && cosNO > 0) {
     // reflect the view vector
-    float3 R = (2 * cosNO) * bsdf->N - I;
+    float3 R = (2 * cosNI) * bsdf->N - wi;
-    float cosRI = dot(R, omega_in);
+    float cosRO = dot(R, wo);
-    if (cosRI > 0) {
+    if (cosRO > 0) {
-      float cosp = powf(cosRI, m_exponent);
+      float cosp = powf(cosRO, m_exponent);
       float common = 0.5f * M_1_PI_F * cosp;
-      float out = cosNI * (m_exponent + 2) * common;
+      float out = cosNO * (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 I,
+                                      float3 wi,
                                       float randu,
                                       float randv,
                                       ccl_private Spectrum *eval,
-                                      ccl_private float3 *omega_in,
+                                      ccl_private float3 *wo,
                                       ccl_private float *pdf,
                                       ccl_private float2 *sampled_roughness)
 {
   ccl_private const PhongRampBsdf *bsdf = (ccl_private const PhongRampBsdf *)sc;
-  float cosNO = dot(bsdf->N, I);
+  float cosNI = dot(bsdf->N, wi);
   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 (cosNO > 0) {
+  if (cosNI > 0) {
     // reflect the view vector
-    float3 R = (2 * cosNO) * bsdf->N - I;
+    float3 R = (2 * cosNI) * bsdf->N - wi;
     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;
-    *omega_in = (cosf(phi) * sinTheta) * T + (sinf(phi) * sinTheta) * B + (cosTheta)*R;
+    *wo = (cosf(phi) * sinTheta) * T + (sinf(phi) * sinTheta) * B + (cosTheta)*R;
-    if (dot(Ng, *omega_in) > 0.0f) {
+    if (dot(Ng, *wo) > 0.0f) {
       // common terms for pdf and eval
-      float cosNI = dot(bsdf->N, *omega_in);
+      float cosNO = dot(bsdf->N, *wo);
       // make sure the direction we chose is still in the right hemisphere
-      if (cosNI > 0) {
+      if (cosNO > 0) {
         float cosp = powf(cosTheta, m_exponent);
         float common = 0.5f * M_1_PI_F * cosp;
         *pdf = (m_exponent + 1) * common;
-        float out = cosNI * (m_exponent + 2) * common;
+        float out = cosNO * (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 I,
+                                                 const float3 wi,
-                                                 const float3 omega_in,
+                                                 const float3 wo,
                                                  ccl_private float *pdf)
 {
   ccl_private const PrincipledDiffuseBsdf *bsdf = (ccl_private const PrincipledDiffuseBsdf *)sc;
   const float3 N = bsdf->N;
 
-  if (dot(N, omega_in) > 0.0f) {
+  if (dot(N, wo) > 0.0f) {
-    const float3 V = I;         // outgoing
+    const float3 V = wi;
-    const float3 L = omega_in;  // incoming
+    const float3 L = wo;
-    *pdf = fmaxf(dot(N, omega_in), 0.0f) * M_1_PI_F;
+    *pdf = fmaxf(dot(N, wo), 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 I,
+                                              float3 wi,
                                               float randu,
                                               float randv,
                                               ccl_private Spectrum *eval,
-                                              ccl_private float3 *omega_in,
+                                              ccl_private float3 *wo,
                                               ccl_private float *pdf)
 {
   ccl_private const PrincipledDiffuseBsdf *bsdf = (ccl_private const PrincipledDiffuseBsdf *)sc;
 
   float3 N = bsdf->N;
 
-  sample_cos_hemisphere(N, randu, randv, omega_in, pdf);
+  sample_cos_hemisphere(N, randu, randv, wo, pdf);
 
-  if (dot(Ng, *omega_in) > 0) {
+  if (dot(Ng, *wo) > 0) {
-    *eval = bsdf_principled_diffuse_compute_brdf(bsdf, N, I, *omega_in, pdf);
+    *eval = bsdf_principled_diffuse_compute_brdf(bsdf, N, wi, *wo, 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->I);
+  bsdf->avg_value = calculate_avg_principled_sheen_brdf(bsdf->N, sd->wi);
   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 I,
+                                               const float3 wi,
-                                               const float3 omega_in,
+                                               const float3 wo,
                                                ccl_private float *pdf)
 {
   ccl_private const PrincipledSheenBsdf *bsdf = (ccl_private const PrincipledSheenBsdf *)sc;
   const float3 N = bsdf->N;
 
-  if (dot(N, omega_in) > 0.0f) {
+  if (dot(N, wo) > 0.0f) {
-    const float3 V = I;         // outgoing
+    const float3 V = wi;
-    const float3 L = omega_in;  // incoming
+    const float3 L = wo;
     const float3 H = normalize(L + V);
 
-    *pdf = fmaxf(dot(N, omega_in), 0.0f) * M_1_PI_F;
+    *pdf = fmaxf(dot(N, wo), 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 I,
+                                            float3 wi,
                                             float randu,
                                             float randv,
                                             ccl_private Spectrum *eval,
-                                            ccl_private float3 *omega_in,
+                                            ccl_private float3 *wo,
                                             ccl_private float *pdf)
 {
   ccl_private const PrincipledSheenBsdf *bsdf = (ccl_private const PrincipledSheenBsdf *)sc;
 
   float3 N = bsdf->N;
 
-  sample_cos_hemisphere(N, randu, randv, omega_in, pdf);
+  sample_cos_hemisphere(N, randu, randv, wo, pdf);
 
-  if (dot(Ng, *omega_in) > 0) {
+  if (dot(Ng, *wo) > 0) {
-    float3 H = normalize(I + *omega_in);
+    float3 H = normalize(wi + *wo);
 
-    *eval = calculate_principled_sheen_brdf(N, I, *omega_in, H, pdf);
+    *eval = calculate_principled_sheen_brdf(N, wi, *wo, 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 I,
+                                         const float3 wi,
-                                         const float3 omega_in,
+                                         const float3 wo,
                                          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 I,
+                                      float3 wi,
                                       float randu,
                                       float randv,
                                       ccl_private Spectrum *eval,
-                                      ccl_private float3 *omega_in,
+                                      ccl_private float3 *wo,
                                       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 cosNO = dot(N, I);
+  float cosNI = dot(N, wi);
-  if (cosNO > 0) {
+  if (cosNI > 0) {
-    *omega_in = (2 * cosNO) * N - I;
+    *wo = (2 * cosNI) * N - wi;
-    if (dot(Ng, *omega_in) > 0) {
+    if (dot(Ng, *wo) > 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 I,
+                                         const float3 wi,
-                                         const float3 omega_in,
+                                         const float3 wo,
                                          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 I,
+                                      float3 wi,
                                       float randu,
                                       float randv,
                                       ccl_private Spectrum *eval,
-                                      ccl_private float3 *omega_in,
+                                      ccl_private float3 *wo,
                                       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, I, &R, &T, &inside);
+  fresnel = fresnel_dielectric(m_eta, N, wi, &R, &T, &inside);
 
   if (!inside && fresnel != 1.0f) {
     /* Some high number for MIS. */
     *pdf = 1e6f;
     *eval = make_spectrum(1e6f);
-    *omega_in = T;
+    *wo = 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 I,
+                                           const float3 wi,
-                                           const float3 omega_in,
+                                           const float3 wo,
                                            ccl_private float *pdf)
 {
   ccl_private const ToonBsdf *bsdf = (ccl_private const ToonBsdf *)sc;
-  float cosNI = dot(bsdf->N, omega_in);
+  float cosNO = dot(bsdf->N, wo);
 
-  if (cosNI >= 0.0f) {
+  if (cosNO >= 0.0f) {
     float max_angle = bsdf->size * M_PI_2_F;
     float smooth = bsdf->smooth * M_PI_2_F;
-    float angle = safe_acosf(fmaxf(cosNI, 0.0f));
+    float angle = safe_acosf(fmaxf(cosNO, 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 I,
+                                        float3 wi,
                                         float randu,
                                         float randv,
                                         ccl_private Spectrum *eval,
-                                        ccl_private float3 *omega_in,
+                                        ccl_private float3 *wo,
                                         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, omega_in, pdf);
+    sample_uniform_cone(bsdf->N, sample_angle, randu, randv, wo, pdf);
 
-    if (dot(Ng, *omega_in) > 0.0f) {
+    if (dot(Ng, *wo) > 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 I,
+                                          const float3 wi,
-                                          const float3 omega_in,
+                                          const float3 wo,
                                           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, omega_in);
+  float cosNI = dot(bsdf->N, wi);
-  float cosNO = dot(bsdf->N, I);
+  float cosNO = dot(bsdf->N, wo);
 
   if (cosNI > 0 && cosNO > 0) {
     /* reflect the view vector */
-    float3 R = (2 * cosNO) * bsdf->N - I;
+    float3 R = (2 * cosNI) * bsdf->N - wi;
-    float cosRI = dot(R, omega_in);
+    float cosRO = dot(R, wo);
 
-    float angle = safe_acosf(fmaxf(cosRI, 0.0f));
+    float angle = safe_acosf(fmaxf(cosRO, 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 I,
+                                       float3 wi,
                                        float randu,
                                        float randv,
                                        ccl_private Spectrum *eval,
-                                       ccl_private float3 *omega_in,
+                                       ccl_private float3 *wo,
                                        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 cosNO = dot(bsdf->N, I);
+  float cosNI = dot(bsdf->N, wi);
 
-  if (cosNO > 0) {
+  if (cosNI > 0) {
     /* reflect the view vector */
-    float3 R = (2 * cosNO) * bsdf->N - I;
+    float3 R = (2 * cosNI) * bsdf->N - wi;
 
     float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
     float angle = sample_angle * randu;
 
-    sample_uniform_cone(R, sample_angle, randu, randv, omega_in, pdf);
+    sample_uniform_cone(R, sample_angle, randu, randv, wo, pdf);
 
-    if (dot(Ng, *omega_in) > 0.0f) {
+    if (dot(Ng, *wo) > 0.0f) {
-      float cosNI = dot(bsdf->N, *omega_in);
+      float cosNO = dot(bsdf->N, *wo);
 
       /* make sure the direction we chose is still in the right hemisphere */
-      if (cosNI > 0) {
+      if (cosNO > 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 I,
+                                          const float3 wi,
-                                          const float3 omega_in,
+                                          const float3 wo,
                                           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 I,
+                                       float3 wi,
                                        float randu,
                                        float randv,
                                        ccl_private Spectrum *eval,
-                                       ccl_private float3 *omega_in,
+                                       ccl_private float3 *wo,
                                        ccl_private float *pdf)
 {
   // only one direction is possible
-  *omega_in = -I;
+  *wo = -wi;
   *pdf = 1;
   *eval = one_spectrum();
   return LABEL_TRANSMIT | LABEL_TRANSPARENT;

intern/cycles/kernel/closure/bsdf_util.h

@@ -89,19 +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 (cspec0) */
+/* 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, float F0, Spectrum cspec0)
+interpolate_fresnel_color(float3 L, float3 H, float ior, Spectrum F0)
 {
-  /* Calculate the fresnel interpolation factor
+  /* Compute the real Fresnel term and remap it from real_F0..1 to F0..1.
-   * The value from fresnel_dielectric_cos(...) has to be normalized because
+   * The reason why we use this remapping instead of directly doing the
-   * the cspec0 keeps the F0 color
+   * 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
-  float F0_norm = 1.0f / (1.0f - F0);
+   * enough to 1.0 that the Schlick approximation becomes inaccurate. */
-  float FH = (fresnel_dielectric_cos(dot(L, H), ior) - F0) * F0_norm;
+  float real_F = fresnel_dielectric_cos(dot(L, H), ior);
+  float real_F0 = fresnel_dielectric_cos(1.0f, ior);
 
-  /* Blend between white and a specular color with respect to the fresnel */
+  return mix(F0, one_spectrum(), inverse_lerp(real_F0, 1.0f, real_F));
-  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->I);
+      bsdf->sample_weight *= bsdf_principled_diffuse_retro_reflection_sample_weight(bsdf, sd->wi);
     }
   }
 

intern/cycles/kernel/closure/emissive.h

@@ -36,27 +36,24 @@ ccl_device void emission_setup(ccl_private ShaderData *sd, const Spectrum weight
   }
 }
 
-/* return the probability distribution function in the direction I,
+/* return the probability distribution function in the direction wi,
  * 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 I)
+ccl_device float emissive_pdf(const float3 Ng, const float3 wi)
 {
-  float cosNO = fabsf(dot(Ng, I));
+  float cosNI = fabsf(dot(Ng, wi));
-  return (cosNO > 0.0f) ? 1.0f : 0.0f;
+  return (cosNI > 0.0f) ? 1.0f : 0.0f;
 }
 
-ccl_device void emissive_sample(const float3 Ng,
+ccl_device void emissive_sample(
-                                float randu,
+    const float3 Ng, float randu, float randv, ccl_private float3 *wi, ccl_private float *pdf)
-                                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 I)
+ccl_device Spectrum emissive_simple_eval(const float3 Ng, const float3 wi)
 {
-  float res = emissive_pdf(Ng, I);
+  float res = emissive_pdf(Ng, wi);
 
   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 I,
+                                                        const float3 wi,
-                                                        float3 omega_in,
+                                                        float3 wo,
                                                         ccl_private float *pdf)
 {
   float g = svc->g;
 
-  /* note that I points towards the viewer */
+  /* note that wi points towards the viewer */
   if (fabsf(g) < 1e-3f) {
     *pdf = M_1_PI_F * 0.25f;
   }
   else {
-    float cos_theta = dot(-I, omega_in);
+    float cos_theta = dot(-wi, wo);
     *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 = safe_sqrtf(1.0f - cos_theta * cos_theta);
+  float sin_theta = sin_from_cos(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 I,
+                                               float3 wi,
                                                float randu,
                                                float randv,
                                                ccl_private Spectrum *eval,
-                                               ccl_private float3 *omega_in,
+                                               ccl_private float3 *wo,
                                                ccl_private float *pdf)
 {
   float g = svc->g;
 
-  /* note that I points towards the viewer and so is used negated */
+  /* note that wi points towards the viewer and so is used negated */
-  *omega_in = henyey_greenstrein_sample(-I, g, randu, randv, pdf);
+  *wo = henyey_greenstrein_sample(-wi, 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 omega_in,
+                                      float3 wo,
                                       ccl_private float *pdf)
 {
-  return volume_henyey_greenstein_eval_phase(svc, sd->I, omega_in, pdf);
+  return volume_henyey_greenstein_eval_phase(svc, sd->wi, wo, 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 *omega_in,
+                                   ccl_private float3 *wo,
                                    ccl_private float *pdf)
 {
-  return volume_henyey_greenstein_sample(svc, sd->I, randu, randv, eval, omega_in, pdf);
+  return volume_henyey_greenstein_sample(svc, sd->wi, randu, randv, eval, wo, pdf);
 }
 
 /* Volume sampling utilities. */

intern/cycles/kernel/data_template.h

@@ -10,6 +10,9 @@
 #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. */
 
@@ -179,9 +182,12 @@ 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, int, tabulated_sobol_sequence_size)
-KERNEL_STRUCT_MEMBER(integrator, int, sobol_index_mask)
 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)
 /* Volume render. */
 KERNEL_STRUCT_MEMBER(integrator, int, use_volumes)
 KERNEL_STRUCT_MEMBER(integrator, int, volume_max_steps)
@@ -216,4 +222,5 @@ 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/oneapi/compat.h

@@ -195,7 +195,15 @@ using sycl::half;
 #define fmodf(x, y) sycl::fmod((x), (y))
 #define lgammaf(x) sycl::lgamma((x))
 
-#define cosf(x) sycl::native::cos(((float)(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 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,23 +58,7 @@ ccl_device_forceinline void film_write_denoising_features_surface(KernelGlobals
     normal += sc->N * sc->sample_weight;
     sum_weight += sc->sample_weight;
 
-    Spectrum closure_albedo = sc->weight;
+    if (sc->type == CLOSURE_BSDF_PRINCIPLED_DIFFUSE_ID) {
-    /* 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;
@@ -83,6 +67,7 @@ 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->I - sd->dPdu * (dot(sd->dPdu, -sd->I) / len_squared(sd->dPdu)));
+    tgN = -(-sd->wi - sd->dPdu * (dot(sd->dPdu, -sd->wi) / 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 = safe_sqrtf(1.0f - sine * sine);
+    const float cosine = cos_from_sin(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->I;
+    const float3 N = (!isequal(P, P_inside)) ? normalize(P - P_inside) : -sd->wi;
 
     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->I : sd->N;
+  sd->Ng = (sd->type & PRIMITIVE_CURVE_RIBBON) ? sd->wi : 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->I = -ray->D;
+  sd->wi = -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->I) < 0.0f);
+  bool backfacing = (dot(sd->Ng, sd->wi) < 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->I = I;
+  sd->wi = 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->I);
+      object_dir_transform_auto(kg, sd, &sd->wi);
     }
 
     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->I) < 0.0f);
+    bool backfacing = (dot(sd->Ng, sd->wi) < 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->I = zero_float3();
+  sd->wi = 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->I = -ray_D;
+  sd->wi = -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->I = -ray->D;
+  sd->wi = -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->I));
+  openpgl::cpp::SetDirectionOut(state->guiding.path_segment, guiding_vec3f(sd->wi));
   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 omega_in,
+                                                          const float3 wo,
                                                           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(omega_in));
+  openpgl::cpp::SetDirectionIn(state->guiding.path_segment, guiding_vec3f(wo));
   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 I)
+                                                          const float3 wi)
 {
 #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(I));
+  openpgl::cpp::SetDirectionOut(state->guiding.path_segment, guiding_vec3f(wi));
   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 omega_in,
+                                                         const float3 wo,
                                                          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(omega_in));
+  openpgl::cpp::SetDirectionIn(state->guiding.path_segment, guiding_vec3f(wo));
   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 omega_in,
+                                                         const float3 wo,
                                                          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(omega_in));
+  openpgl::cpp::SetDirectionIn(state->guiding.path_segment, guiding_vec3f(wo));
   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 *omega_in)
+                                                 ccl_private float3 *wo)
 {
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
-  pgl_vec3f wo;
+  pgl_vec3f pgl_wo;
   const pgl_point2f rand = openpgl::cpp::Point2(rand_bsdf.x, rand_bsdf.y);
-  const float pdf = kg->opgl_surface_sampling_distribution->SamplePDF(rand, wo);
+  const float pdf = kg->opgl_surface_sampling_distribution->SamplePDF(rand, pgl_wo);
-  *omega_in = make_float3(wo.x, wo.y, wo.z);
+  *wo = make_float3(pgl_wo.x, pgl_wo.y, pgl_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 omega_in)
+                                              const float3 wo)
 {
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
-  return kg->opgl_surface_sampling_distribution->PDF(guiding_vec3f(omega_in));
+  return kg->opgl_surface_sampling_distribution->PDF(guiding_vec3f(wo));
 #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 *omega_in)
+                                                  ccl_private float3 *wo)
 {
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
-  pgl_vec3f wo;
+  pgl_vec3f pgl_wo;
   const pgl_point2f rand = openpgl::cpp::Point2(rand_phase.x, rand_phase.y);
-  const float pdf = kg->opgl_volume_sampling_distribution->SamplePDF(rand, wo);
+  const float pdf = kg->opgl_volume_sampling_distribution->SamplePDF(rand, pgl_wo);
-  *omega_in = make_float3(wo.x, wo.y, wo.z);
+  *wo = make_float3(pgl_wo.x, pgl_wo.y, pgl_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 omega_in)
+                                               const float3 wo)
 {
 #if defined(__PATH_GUIDING__) && PATH_GUIDING_LEVEL >= 4
-  return kg->opgl_volume_sampling_distribution->PDF(guiding_vec3f(omega_in));
+  return kg->opgl_volume_sampling_distribution->PDF(guiding_vec3f(wo));
 #else
   return 0.0f;
 #endif

intern/cycles/kernel/integrator/mnee.h

@@ -607,24 +607,22 @@ ccl_device_forceinline Spectrum mnee_eval_bsdf_contribution(ccl_private ShaderCl
 {
   ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)closure;
 
-  float cosNO = dot(bsdf->N, wi);
+  float cosNI = dot(bsdf->N, wi);
-  float cosNI = dot(bsdf->N, wo);
+  float cosNO = dot(bsdf->N, wo);
 
   float3 Ht = normalize(-(bsdf->ior * wo + wi));
-  float cosHO = dot(Ht, wi);
+  float cosHI = 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) {
-    /* Eq. 26, 27: now calculate G1(i,m) and G1(o,m). */
+    G = bsdf_G<MicrofacetType::BECKMANN>(alpha2, cosNI, cosNO);
-    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 */
-    /* Eq. 34: now calculate G1(i,m) and G1(o,m). */
+    G = bsdf_G<MicrofacetType::GGX>(alpha2, cosNI, cosNO);
-    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))));
   }
 
   /*
@@ -635,7 +633,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(cosHO / (cosNO * sqr(cosThetaM)));
+  return bsdf->weight * G * fabsf(cosHI / (cosNI * sqr(cosThetaM)));
 }
 
 /* Compute transfer matrix determinant |T1| = |dx1/dxn| (and |dh/dx| in the process) */
@@ -706,9 +704,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 = safe_sqrtf(1.f - sqr(cos_theta));
+    float sin_theta = sin_from_cos(cos_theta);
     float cos_phi = dot(wo, s);
-    float sin_phi = safe_sqrtf(1.f - sqr(cos_phi));
+    float sin_phi = sin_from_cos(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,8 +235,6 @@ 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);
@@ -264,7 +262,6 @@ 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);
 
@@ -364,7 +361,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_omega_in ccl_optional_struct_init;
+  float3 bsdf_wo ccl_optional_struct_init;
   int label;
 
   float2 bsdf_sampled_roughness = make_float2(1.0f, 1.0f);
@@ -378,7 +375,7 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
                                                       sc,
                                                       rand_bsdf,
                                                       &bsdf_eval,
-                                                      &bsdf_omega_in,
+                                                      &bsdf_wo,
                                                       &bsdf_pdf,
                                                       &unguided_bsdf_pdf,
                                                       &bsdf_sampled_roughness,
@@ -398,7 +395,7 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
                                                sc,
                                                rand_bsdf,
                                                &bsdf_eval,
-                                               &bsdf_omega_in,
+                                               &bsdf_wo,
                                                &bsdf_pdf,
                                                &bsdf_sampled_roughness,
                                                &bsdf_eta);
@@ -416,7 +413,7 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
   }
   else {
     /* Setup ray with changed origin and direction. */
-    const float3 D = normalize(bsdf_omega_in);
+    const float3 D = normalize(bsdf_wo);
     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;
@@ -455,7 +452,7 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
                                 bsdf_weight,
                                 bsdf_pdf,
                                 sd->N,
-                                normalize(bsdf_omega_in),
+                                normalize(bsdf_wo),
                                 bsdf_sampled_roughness,
                                 bsdf_eta);
 

intern/cycles/kernel/integrator/shade_volume.h

@@ -821,7 +821,6 @@ 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(
@@ -838,7 +837,6 @@ 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) {
@@ -912,7 +910,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_omega_in ccl_optional_struct_init;
+  float3 phase_wo ccl_optional_struct_init;
   float sampled_roughness = 1.0f;
   int label;
 
@@ -924,7 +922,7 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
                                               svc,
                                               rand_phase,
                                               &phase_eval,
-                                              &phase_omega_in,
+                                              &phase_wo,
                                               &phase_pdf,
                                               &unguided_phase_pdf,
                                               &sampled_roughness);
@@ -938,15 +936,8 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
   else
 #  endif
   {
-    label = volume_shader_phase_sample(kg,
+    label = volume_shader_phase_sample(
-                                       sd,
+        kg, sd, phases, svc, rand_phase, &phase_eval, &phase_wo, &phase_pdf, &sampled_roughness);
-                                       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;
@@ -957,7 +948,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_omega_in);
+  INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(phase_wo);
   INTEGRATOR_STATE_WRITE(state, ray, tmin) = 0.0f;
   INTEGRATOR_STATE_WRITE(state, ray, tmax) = FLT_MAX;
 #  ifdef __RAY_DIFFERENTIALS__
@@ -971,7 +962,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_omega_in), sampled_roughness);
+      kg, state, sd, phase_weight, phase_pdf, normalize(phase_wo), sampled_roughness);
 
   /* Update throughput. */
   const Spectrum throughput = INTEGRATOR_STATE(state, path, throughput);
@@ -1076,7 +1067,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.I);
+        guiding_record_volume_segment(kg, state, direct_P, sd.wi);
         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);
@@ -1139,7 +1130,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.I);
+      guiding_record_volume_segment(kg, state, sd.P, sd.wi);
     }
 #    endif
 #    if PATH_GUIDING_LEVEL >= 4

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 = safe_sqrtf(1.0f - cos_theta * cos_theta);
+  float sin_theta = sin_from_cos(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 omega_in,
+                                                           const float3 wo,
                                                            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, omega_in);
+  const int comp_label = bsdf_label(kg, sc, wo);
   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 omega_in,
+                                                      const float3 wo,
                                                       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, omega_in, &bsdf_pdf);
+        Spectrum eval = bsdf_eval(kg, sd, sc, wo, &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 omega_in,
+                                                      const float3 wo,
                                                       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, omega_in, &bsdf_pdf);
+        Spectrum eval = bsdf_eval(kg, sd, sc, wo, &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 omega_in,
+                             const float3 wo,
                              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, omega_in, NULL, bsdf_eval, 0.0f, 0.0f, light_shader_flags);
+      kg, sd, wo, 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, omega_in);
+    const float guide_pdf = guiding_bsdf_pdf(kg, state, wo);
     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 *omega_in,
+                                                         ccl_private float3 *wo,
                                                          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, omega_in);
+    guide_pdf = guiding_bsdf_sample(kg, state, rand_bsdf, wo);
     *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, *omega_in, bsdf_eval, unguided_bsdf_pdfs, 0);
+          kg, sd, *wo, 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], *omega_in);
+        label = bsdf_label(kg, &sd->closure[idx], *wo);
       }
     }
 
@@ -483,19 +483,11 @@ ccl_device int surface_shader_bsdf_guided_sample_closure(KernelGlobals kg,
   else {
     /* Sample BSDF. */
     *bsdf_pdf = 0.0f;
-    label = bsdf_sample(kg,
+    label = bsdf_sample(
-                        sd,
+        kg, sd, sc, rand_bsdf.x, rand_bsdf.y, &eval, wo, unguided_bsdf_pdf, sampled_rougness, eta);
-                        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, *omega_in, label, sampled_roughness, eta);
+      surface_shader_validate_bsdf_sample(kg, sc, *wo, label, sampled_roughness, eta);
     }
 #  endif
 
@@ -507,13 +499,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, *omega_in, sc, bsdf_eval, (*unguided_bsdf_pdf) * sweight, sweight, 0);
+            kg, sd, *wo, 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, *omega_in);
+        guide_pdf = guiding_bsdf_pdf(kg, state, *wo);
         *bsdf_pdf *= 1.0f - guiding_sampling_prob;
         *bsdf_pdf += guiding_sampling_prob * guide_pdf * (1.0f - bssrdf_sampling_prob);
       }
@@ -533,7 +525,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 *omega_in,
+                                                  ccl_private float3 *wo,
                                                   ccl_private float *pdf,
                                                   ccl_private float2 *sampled_roughness,
                                                   ccl_private float *eta)
@@ -546,15 +538,14 @@ 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, omega_in, pdf, sampled_roughness, eta);
+      kg, sd, sc, rand_bsdf.x, rand_bsdf.y, &eval, wo, 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(
+      *pdf = _surface_shader_bsdf_eval_mis(kg, sd, *wo, sc, bsdf_eval, *pdf * sweight, sweight, 0);
-          kg, sd, *omega_in, sc, bsdf_eval, *pdf * sweight, sweight, 0);
     }
   }
   else {
@@ -630,7 +621,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 += sc->weight;
+      eval += bsdf_albedo(sd, sc);
   }
 
   return eval;
@@ -644,7 +635,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 += sc->weight;
+      eval += bsdf_albedo(sd, sc);
   }
 
   return eval;
@@ -658,7 +649,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 += sc->weight;
+      eval += bsdf_albedo(sd, sc);
   }
 
   return eval;
@@ -758,7 +749,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->I) * sd->closure_emission_background;
+    return emissive_simple_eval(sd->Ng, sd->wi) * 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 omega_in,
+                                                      const float3 wo,
                                                       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, omega_in, &phase_pdf);
+    Spectrum eval = volume_phase_eval(sd, svc, wo, &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 omega_in,
+                                          const float3 wo,
                                           ccl_private BsdfEval *phase_eval)
 {
   float phase_pdf = 0.0f;
-  Spectrum eval = volume_phase_eval(sd, svc, omega_in, &phase_pdf);
+  Spectrum eval = volume_phase_eval(sd, svc, wo, &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 omega_in,
+                                          const float3 wo,
                                           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, omega_in, -1, phase_eval, 0.0f, 0.0f);
+  float pdf = _volume_shader_phase_eval_mis(sd, phases, wo, -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, omega_in);
+    const float guide_pdf = guiding_phase_pdf(kg, state, wo);
     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 *omega_in,
+                                                 ccl_private float3 *wo,
                                                  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, omega_in);
+    guide_pdf = guiding_phase_sample(kg, state, rand_phase, wo);
     *phase_pdf = 0.0f;
 
     if (guide_pdf != 0.0f) {
-      *unguided_phase_pdf = volume_shader_phase_eval(kg, sd, svc, *omega_in, phase_eval);
+      *unguided_phase_pdf = volume_shader_phase_eval(kg, sd, svc, *wo, 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, omega_in, unguided_phase_pdf);
+        sd, svc, rand_phase.x, rand_phase.y, &eval, wo, 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, *omega_in);
+        guide_pdf = guiding_phase_pdf(kg, state, *wo);
         *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 *omega_in,
+                                          ccl_private float3 *wo,
                                           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, omega_in, pdf);
+  int label = volume_phase_sample(sd, svc, rand_phase.x, rand_phase.y, &eval, wo, 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 = safe_sqrtf(1.0f - sqr(cos_a));
+  const float sin_a = sin_from_cos(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,13 +88,6 @@ 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,

intern/cycles/kernel/light/spot.h

@@ -7,24 +7,13 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device float spot_light_attenuation(float3 dir,
+ccl_device float spot_light_attenuation(const ccl_global KernelSpotLight *spot, float3 ray)
-                                        float cos_half_spot_angle,
-                                        float spot_smooth,
-                                        float3 N)
 {
-  float attenuation = dot(dir, N);
+  const float3 scaled_ray = safe_normalize(
+      make_float3(dot(ray, spot->axis_u), dot(ray, spot->axis_v), dot(ray, spot->dir)) /
+      spot->len);
 
-  if (attenuation <= cos_half_spot_angle) {
+  return smoothstepf((scaled_ray.z - spot->cos_half_spot_angle) / spot->spot_smooth);
-    attenuation = 0.0f;
-  }
-  else {
-    float t = attenuation - cos_half_spot_angle;
-
-    if (t < spot_smooth && spot_smooth != 0.0f)
-      attenuation *= smoothstepf(t / spot_smooth);
-  }
-
-  return attenuation;
 }
 
 template<bool in_volume_segment>
@@ -57,8 +46,7 @@ ccl_device_inline bool spot_light_sample(const ccl_global KernelLight *klight,
   ls->eval_fac = (0.25f * M_1_PI_F) * invarea;
 
   /* spot light attenuation */
-  ls->eval_fac *= spot_light_attenuation(
+  ls->eval_fac *= spot_light_attenuation(&klight->spot, -ls->D);
-      klight->spot.dir, klight->spot.cos_half_spot_angle, klight->spot.spot_smooth, -ls->D);
   if (!in_volume_segment && ls->eval_fac == 0.0f) {
     return false;
   }
@@ -87,8 +75,7 @@ ccl_device_forceinline void spot_light_update_position(const ccl_global KernelLi
   ls->pdf = invarea;
 
   /* spot light attenuation */
-  ls->eval_fac *= spot_light_attenuation(
+  ls->eval_fac *= spot_light_attenuation(&klight->spot, ls->Ng);
-      klight->spot.dir, klight->spot.cos_half_spot_angle, klight->spot.spot_smooth, ls->Ng);
 }
 
 ccl_device_inline bool spot_light_intersect(const ccl_global KernelLight *klight,
@@ -129,8 +116,7 @@ ccl_device_inline bool spot_light_sample_from_intersection(
   ls->pdf = invarea;
 
   /* spot light attenuation */
-  ls->eval_fac *= spot_light_attenuation(
+  ls->eval_fac *= spot_light_attenuation(&klight->spot, -ls->D);
-      klight->spot.dir, klight->spot.cos_half_spot_angle, klight->spot.spot_smooth, -ls->D);
 
   if (ls->eval_fac == 0.0f) {
     return false;

intern/cycles/kernel/light/tree.h

@@ -47,11 +47,6 @@ ccl_device float light_tree_cos_bounding_box_angle(const BoundingBox bbox,
   return cos_theta_u;
 }
 
-ccl_device_forceinline float sin_from_cos(const float c)
-{
-  return safe_sqrtf(1.0f - sqr(c));
-}
-
 /* Compute vector v as in Fig .8. P_v is the corresponding point along the ray. */
 ccl_device float3 compute_v(
     const float3 centroid, const float3 P, const float3 D, const float3 bcone_axis, const float t)

intern/cycles/kernel/light/triangle.h

@@ -63,7 +63,7 @@ ccl_device_forceinline float triangle_light_pdf(KernelGlobals kg,
   const float3 e2 = V[2] - V[1];
   const float longest_edge_squared = max(len_squared(e0), max(len_squared(e1), len_squared(e2)));
   const float3 N = cross(e0, e1);
-  const float distance_to_plane = fabsf(dot(N, sd->I * t)) / dot(N, N);
+  const float distance_to_plane = fabsf(dot(N, sd->wi * t)) / dot(N, N);
   const float area = 0.5f * len(N);
 
   float pdf;
@@ -71,7 +71,7 @@ ccl_device_forceinline float triangle_light_pdf(KernelGlobals kg,
   if (longest_edge_squared > distance_to_plane * distance_to_plane) {
     /* sd contains the point on the light source
      * calculate Px, the point that we're shading */
-    const float3 Px = sd->P + sd->I * t;
+    const float3 Px = sd->P + sd->wi * t;
     const float3 v0_p = V[0] - Px;
     const float3 v1_p = V[1] - Px;
     const float3 v2_p = V[2] - Px;
@@ -99,7 +99,7 @@ ccl_device_forceinline float triangle_light_pdf(KernelGlobals kg,
       return 0.0f;
     }
 
-    pdf = triangle_light_pdf_area_sampling(sd->Ng, sd->I, t) / area;
+    pdf = triangle_light_pdf_area_sampling(sd->Ng, sd->wi, t) / area;
   }
 
   /* Belongs in distribution.h but can reuse computations here. */
@@ -218,7 +218,7 @@ ccl_device_forceinline bool triangle_light_sample(KernelGlobals kg,
     /* Finally, select a random point along the edge of the new triangle
      * That point on the spherical triangle is the sampled ray direction */
     const float z = 1.0f - randv * (1.0f - dot(C_, B));
-    ls->D = z * B + safe_sqrtf(1.0f - z * z) * safe_normalize(C_ - dot(C_, B) * B);
+    ls->D = z * B + sin_from_cos(z) * safe_normalize(C_ - dot(C_, B) * B);
 
     /* calculate intersection with the planar triangle */
     if (!ray_triangle_intersect(

intern/cycles/kernel/osl/closures_setup.h

@@ -80,7 +80,7 @@ ccl_device void osl_closure_diffuse_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
 
   sd->flag |= bsdf_diffuse_setup(bsdf);
 }
@@ -101,7 +101,7 @@ ccl_device void osl_closure_oren_nayar_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->roughness = closure->roughness;
 
   sd->flag |= bsdf_oren_nayar_setup(bsdf);
@@ -123,7 +123,7 @@ ccl_device void osl_closure_translucent_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
 
   sd->flag |= bsdf_translucent_setup(bsdf);
 }
@@ -144,7 +144,7 @@ ccl_device void osl_closure_reflection_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
 
   sd->flag |= bsdf_reflection_setup(bsdf);
 }
@@ -165,7 +165,7 @@ ccl_device void osl_closure_refraction_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->ior = closure->ior;
 
   sd->flag |= bsdf_refraction_setup(bsdf);
@@ -199,7 +199,7 @@ ccl_device void osl_closure_microfacet_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = closure->alpha_y;
   bsdf->ior = closure->ior;
@@ -257,7 +257,7 @@ ccl_device void osl_closure_microfacet_ggx_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
 
   sd->flag |= bsdf_microfacet_ggx_isotropic_setup(bsdf);
@@ -280,7 +280,7 @@ ccl_device void osl_closure_microfacet_ggx_aniso_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = closure->alpha_y;
   bsdf->T = closure->T;
@@ -305,7 +305,7 @@ ccl_device void osl_closure_microfacet_ggx_refraction_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->ior = closure->ior;
 
@@ -337,7 +337,7 @@ ccl_device void osl_closure_microfacet_ggx_fresnel_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = bsdf->alpha_x;
   bsdf->ior = closure->ior;
@@ -345,7 +345,6 @@ ccl_device void osl_closure_microfacet_ggx_fresnel_setup(
   bsdf->extra = extra;
   bsdf->extra->color = rgb_to_spectrum(closure->color);
   bsdf->extra->cspec0 = rgb_to_spectrum(closure->cspec0);
-  bsdf->extra->clearcoat = 0.0f;
 
   bsdf->T = zero_float3();
 
@@ -375,7 +374,7 @@ ccl_device void osl_closure_microfacet_ggx_aniso_fresnel_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = closure->alpha_y;
   bsdf->ior = closure->ior;
@@ -383,7 +382,6 @@ ccl_device void osl_closure_microfacet_ggx_aniso_fresnel_setup(
   bsdf->extra = extra;
   bsdf->extra->color = rgb_to_spectrum(closure->color);
   bsdf->extra->cspec0 = rgb_to_spectrum(closure->cspec0);
-  bsdf->extra->clearcoat = 0.0f;
 
   bsdf->T = closure->T;
 
@@ -418,7 +416,7 @@ ccl_device void osl_closure_microfacet_multi_ggx_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = bsdf->alpha_x;
   bsdf->ior = 1.0f;
@@ -426,7 +424,6 @@ ccl_device void osl_closure_microfacet_multi_ggx_setup(
   bsdf->extra = extra;
   bsdf->extra->color = rgb_to_spectrum(closure->color);
   bsdf->extra->cspec0 = zero_spectrum();
-  bsdf->extra->clearcoat = 0.0f;
 
   bsdf->T = zero_float3();
 
@@ -459,7 +456,7 @@ ccl_device void osl_closure_microfacet_multi_ggx_glass_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = bsdf->alpha_x;
   bsdf->ior = closure->ior;
@@ -467,7 +464,6 @@ ccl_device void osl_closure_microfacet_multi_ggx_glass_setup(
   bsdf->extra = extra;
   bsdf->extra->color = rgb_to_spectrum(closure->color);
   bsdf->extra->cspec0 = zero_spectrum();
-  bsdf->extra->clearcoat = 0.0f;
 
   bsdf->T = zero_float3();
 
@@ -500,7 +496,7 @@ ccl_device void osl_closure_microfacet_multi_ggx_aniso_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = closure->alpha_y;
   bsdf->ior = 1.0f;
@@ -508,7 +504,6 @@ ccl_device void osl_closure_microfacet_multi_ggx_aniso_setup(
   bsdf->extra = extra;
   bsdf->extra->color = rgb_to_spectrum(closure->color);
   bsdf->extra->cspec0 = zero_spectrum();
-  bsdf->extra->clearcoat = 0.0f;
 
   bsdf->T = closure->T;
 
@@ -543,7 +538,7 @@ ccl_device void osl_closure_microfacet_multi_ggx_fresnel_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = bsdf->alpha_x;
   bsdf->ior = closure->ior;
@@ -551,7 +546,6 @@ ccl_device void osl_closure_microfacet_multi_ggx_fresnel_setup(
   bsdf->extra = extra;
   bsdf->extra->color = rgb_to_spectrum(closure->color);
   bsdf->extra->cspec0 = rgb_to_spectrum(closure->cspec0);
-  bsdf->extra->clearcoat = 0.0f;
 
   bsdf->T = zero_float3();
 
@@ -584,7 +578,7 @@ ccl_device void osl_closure_microfacet_multi_ggx_glass_fresnel_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = bsdf->alpha_x;
   bsdf->ior = closure->ior;
@@ -592,7 +586,6 @@ ccl_device void osl_closure_microfacet_multi_ggx_glass_fresnel_setup(
   bsdf->extra = extra;
   bsdf->extra->color = rgb_to_spectrum(closure->color);
   bsdf->extra->cspec0 = rgb_to_spectrum(closure->cspec0);
-  bsdf->extra->clearcoat = 0.0f;
 
   bsdf->T = zero_float3();
 
@@ -625,7 +618,7 @@ ccl_device void osl_closure_microfacet_multi_ggx_aniso_fresnel_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = closure->alpha_y;
   bsdf->ior = closure->ior;
@@ -633,7 +626,6 @@ ccl_device void osl_closure_microfacet_multi_ggx_aniso_fresnel_setup(
   bsdf->extra = extra;
   bsdf->extra->color = rgb_to_spectrum(closure->color);
   bsdf->extra->cspec0 = rgb_to_spectrum(closure->cspec0);
-  bsdf->extra->clearcoat = 0.0f;
 
   bsdf->T = closure->T;
 
@@ -659,7 +651,7 @@ ccl_device void osl_closure_microfacet_beckmann_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
 
   sd->flag |= bsdf_microfacet_beckmann_isotropic_setup(bsdf);
@@ -682,7 +674,7 @@ ccl_device void osl_closure_microfacet_beckmann_aniso_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = closure->alpha_y;
   bsdf->T = closure->T;
@@ -707,7 +699,7 @@ ccl_device void osl_closure_microfacet_beckmann_refraction_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->ior = closure->ior;
 
@@ -733,7 +725,7 @@ ccl_device void osl_closure_ashikhmin_velvet_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->sigma = closure->sigma;
 
   sd->flag |= bsdf_ashikhmin_velvet_setup(bsdf);
@@ -756,7 +748,7 @@ ccl_device void osl_closure_ashikhmin_shirley_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->alpha_x = closure->alpha_x;
   bsdf->alpha_y = closure->alpha_y;
   bsdf->T = closure->T;
@@ -780,7 +772,7 @@ ccl_device void osl_closure_diffuse_toon_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->size = closure->size;
   bsdf->smooth = closure->smooth;
 
@@ -803,7 +795,7 @@ ccl_device void osl_closure_glossy_toon_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->size = closure->size;
   bsdf->smooth = closure->smooth;
 
@@ -829,7 +821,7 @@ ccl_device void osl_closure_principled_diffuse_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->roughness = closure->roughness;
 
   sd->flag |= bsdf_principled_diffuse_setup(bsdf);
@@ -852,7 +844,7 @@ ccl_device void osl_closure_principled_sheen_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->avg_value = 0.0f;
 
   sd->flag |= bsdf_principled_sheen_setup(sd, bsdf);
@@ -865,27 +857,18 @@ ccl_device void osl_closure_principled_clearcoat_setup(
     float3 weight,
     ccl_private const PrincipledClearcoatClosure *closure)
 {
+  weight *= 0.25f * closure->clearcoat;
   ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
       sd, sizeof(MicrofacetBsdf), rgb_to_spectrum(weight));
   if (!bsdf) {
     return;
   }
 
-  MicrofacetExtra *extra = (MicrofacetExtra *)closure_alloc_extra(sd, sizeof(MicrofacetExtra));
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
-  if (!extra) {
-    return;
-  }
-
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
   bsdf->alpha_x = closure->clearcoat_roughness;
   bsdf->alpha_y = closure->clearcoat_roughness;
   bsdf->ior = 1.5f;
 
-  bsdf->extra = extra;
-  bsdf->extra->color = zero_spectrum();
-  bsdf->extra->cspec0 = make_spectrum(0.04f);
-  bsdf->extra->clearcoat = closure->clearcoat;
-
   bsdf->T = zero_float3();
 
   sd->flag |= bsdf_microfacet_ggx_clearcoat_setup(bsdf, sd);
@@ -948,7 +931,7 @@ ccl_device void osl_closure_diffuse_ramp_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
 
   bsdf->colors = (float3 *)closure_alloc_extra(sd, sizeof(float3) * 8);
   if (!bsdf->colors) {
@@ -973,7 +956,7 @@ ccl_device void osl_closure_phong_ramp_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->exponent = closure->exponent;
 
   bsdf->colors = (float3 *)closure_alloc_extra(sd, sizeof(float3) * 8);
@@ -1024,7 +1007,7 @@ ccl_device void osl_closure_bssrdf_setup(KernelGlobals kg,
 
   /* create one closure per color channel */
   bssrdf->albedo = closure->albedo;
-  bssrdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bssrdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bssrdf->roughness = closure->roughness;
   bssrdf->anisotropy = clamp(closure->anisotropy, 0.0f, 0.9f);
 
@@ -1049,7 +1032,7 @@ ccl_device void osl_closure_hair_reflection_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->T = closure->T;
   bsdf->roughness1 = closure->roughness1;
   bsdf->roughness2 = closure->roughness2;
@@ -1075,7 +1058,7 @@ ccl_device void osl_closure_hair_transmission_setup(
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->T = closure->T;
   bsdf->roughness1 = closure->roughness1;
   bsdf->roughness2 = closure->roughness2;
@@ -1107,7 +1090,7 @@ ccl_device void osl_closure_principled_hair_setup(KernelGlobals kg,
     return;
   }
 
-  bsdf->N = ensure_valid_reflection(sd->Ng, sd->I, closure->N);
+  bsdf->N = ensure_valid_reflection(sd->Ng, sd->wi, closure->N);
   bsdf->sigma = closure->sigma;
   bsdf->v = closure->v;
   bsdf->s = closure->s;

intern/cycles/kernel/osl/osl.h

@@ -25,13 +25,13 @@ ccl_device_inline void shaderdata_to_shaderglobals(KernelGlobals kg,
                                                    ccl_private ShaderGlobals *globals)
 {
   const differential3 dP = differential_from_compact(sd->Ng, sd->dP);
-  const differential3 dI = differential_from_compact(sd->I, sd->dI);
+  const differential3 dI = differential_from_compact(sd->wi, sd->dI);
 
   /* copy from shader data to shader globals */
   globals->P = sd->P;
   globals->dPdx = dP.dx;
   globals->dPdy = dP.dy;
-  globals->I = sd->I;
+  globals->I = sd->wi;
   globals->dIdx = dI.dx;
   globals->dIdy = dI.dy;
   globals->N = sd->N;
@@ -161,7 +161,10 @@ ccl_device_inline void osl_eval_nodes(KernelGlobals kg,
                         /* shadeindex = */ 0);
 #  endif
 
-  if (globals.Ci) {
+  if constexpr (type == SHADER_TYPE_DISPLACEMENT) {
+    sd->P = globals.P;
+  }
+  else if (globals.Ci) {
     flatten_closure_tree(kg, sd, path_flag, globals.Ci);
   }
 }

intern/cycles/kernel/osl/services.cpp

@@ -20,6 +20,7 @@
 
 #include "kernel/osl/globals.h"
 #include "kernel/osl/services.h"
+#include "kernel/osl/types.h"
 
 #include "util/foreach.h"
 #include "util/log.h"
@@ -119,6 +120,8 @@ ustring OSLRenderServices::u_u("u");
 ustring OSLRenderServices::u_v("v");
 ustring OSLRenderServices::u_empty;
 
+ImageManager *OSLRenderServices::image_manager = nullptr;
+
 OSLRenderServices::OSLRenderServices(OSL::TextureSystem *texture_system, int device_type)
     : OSL::RendererServices(texture_system), device_type_(device_type)
 {
@@ -1154,7 +1157,7 @@ TextureSystem::TextureHandle *OSLRenderServices::get_texture_handle(ustring file
     /* For non-OIIO textures, just return a pointer to our own OSLTextureHandle. */
     if (it != textures.end()) {
       if (it->second->type != OSLTextureHandle::OIIO) {
-        return (TextureSystem::TextureHandle *)it->second.get();
+        return reinterpret_cast<TextureSystem::TextureHandle *>(it->second.get());
       }
     }
 
@@ -1173,17 +1176,54 @@ TextureSystem::TextureHandle *OSLRenderServices::get_texture_handle(ustring file
 
     /* Assign OIIO texture handle and return. */
     it->second->oiio_handle = handle;
-    return (TextureSystem::TextureHandle *)it->second.get();
+    return reinterpret_cast<TextureSystem::TextureHandle *>(it->second.get());
   }
   else {
-    if (it != textures.end() && it->second->type == OSLTextureHandle::SVM &&
+    /* Construct GPU texture handle for existing textures. */
-        it->second->svm_slots[0].w == -1) {
+    if (it != textures.end()) {
+      switch (it->second->type) {
+        case OSLTextureHandle::OIIO:
+          return NULL;
+        case OSLTextureHandle::SVM:
+          if (!it->second->handle.empty() && it->second->handle.get_manager() != image_manager) {
+            it.clear();
+            break;
+          }
+          return reinterpret_cast<TextureSystem::TextureHandle *>(OSL_TEXTURE_HANDLE_TYPE_SVM |
+                                                                  it->second->svm_slots[0].y);
+        case OSLTextureHandle::IES:
+          if (!it->second->handle.empty() && it->second->handle.get_manager() != image_manager) {
+            it.clear();
+            break;
+          }
+          return reinterpret_cast<TextureSystem::TextureHandle *>(OSL_TEXTURE_HANDLE_TYPE_IES |
+                                                                  it->second->svm_slots[0].y);
+        case OSLTextureHandle::AO:
           return reinterpret_cast<TextureSystem::TextureHandle *>(
-          static_cast<uintptr_t>(it->second->svm_slots[0].y + 1));
+              OSL_TEXTURE_HANDLE_TYPE_AO_OR_BEVEL | 1);
+        case OSLTextureHandle::BEVEL:
+          return reinterpret_cast<TextureSystem::TextureHandle *>(
+              OSL_TEXTURE_HANDLE_TYPE_AO_OR_BEVEL | 2);
+      }
     }
 
+    if (!image_manager) {
       return NULL;
     }
+
+    /* Load new textures using SVM image manager. */
+    ImageHandle handle = image_manager->add_image(filename.string(), ImageParams());
+    if (handle.empty()) {
+      return NULL;
+    }
+
+    if (!textures.insert(filename, new OSLTextureHandle(handle))) {
+      return NULL;
+    }
+
+    return reinterpret_cast<TextureSystem::TextureHandle *>(OSL_TEXTURE_HANDLE_TYPE_SVM |
+                                                            handle.svm_slot());
+  }
 }
 
 bool OSLRenderServices::good(TextureSystem::TextureHandle *texture_handle)
@@ -1720,8 +1760,8 @@ bool OSLRenderServices::getmessage(OSL::ShaderGlobals *sg,
           return set_attribute_float3(f, type, derivatives, val);
         }
         else if (name == u_I) {
-          const differential3 dI = differential_from_compact(sd->I, sd->dI);
+          const differential3 dI = differential_from_compact(sd->wi, sd->dI);
-          float3 f[3] = {sd->I, dI.dx, dI.dy};
+          float3 f[3] = {sd->wi, dI.dx, dI.dy};
           return set_attribute_float3(f, type, derivatives, val);
         }
         else if (name == u_u) {

intern/cycles/kernel/osl/services.h

@@ -16,6 +16,8 @@
 #include <OSL/oslexec.h>
 #include <OSL/rendererservices.h>
 
+#include "scene/image.h"
+
 #ifdef WITH_PTEX
 class PtexCache;
 #endif
@@ -54,10 +56,20 @@ struct OSLTextureHandle : public OIIO::RefCnt {
   {
   }
 
+  OSLTextureHandle(const ImageHandle &handle)
+      : type(SVM),
+        svm_slots(handle.get_svm_slots()),
+        oiio_handle(nullptr),
+        processor(nullptr),
+        handle(handle)
+  {
+  }
+
   Type type;
   vector<int4> svm_slots;
   OSL::TextureSystem::TextureHandle *oiio_handle;
   ColorSpaceProcessor *processor;
+  ImageHandle handle;
 };
 
 typedef OIIO::intrusive_ptr<OSLTextureHandle> OSLTextureHandleRef;
@@ -324,6 +336,8 @@ class OSLRenderServices : public OSL::RendererServices {
    * shading system. */
   OSLTextureHandleMap textures;
 
+  static ImageManager *image_manager;
+
  private:
   int device_type_;
 };

intern/cycles/kernel/osl/services_gpu.h

@@ -1443,6 +1443,8 @@ OSL_NOISE_IMPL(osl_snoise, snoise)
 
 /* Texturing */
 
+#include "kernel/svm/ies.h"
+
 ccl_device_extern ccl_private OSLTextureOptions *osl_get_texture_options(
     ccl_private ShaderGlobals *sg)
 {
@@ -1548,15 +1550,12 @@ ccl_device_extern bool osl_texture(ccl_private ShaderGlobals *sg,
                                    ccl_private float *dalphady,
                                    ccl_private void *errormessage)
 {
-  if (!texture_handle) {
+  const unsigned int type = OSL_TEXTURE_HANDLE_TYPE(texture_handle);
-    return false;
+  const unsigned int slot = OSL_TEXTURE_HANDLE_SLOT(texture_handle);
-  }
-
-  /* Only SVM textures are supported. */
-  int id = static_cast<int>(reinterpret_cast<size_t>(texture_handle) - 1);
-
-  const float4 rgba = kernel_tex_image_interp(nullptr, id, s, 1.0f - t);
 
+  switch (type) {
+    case OSL_TEXTURE_HANDLE_TYPE_SVM: {
+      const float4 rgba = kernel_tex_image_interp(nullptr, slot, s, 1.0f - t);
       if (nchannels > 0)
         result[0] = rgba.x;
       if (nchannels > 1)
@@ -1565,8 +1564,17 @@ ccl_device_extern bool osl_texture(ccl_private ShaderGlobals *sg,
         result[2] = rgba.z;
       if (alpha)
         *alpha = rgba.w;
-
       return true;
+    }
+    case OSL_TEXTURE_HANDLE_TYPE_IES: {
+      if (nchannels > 0)
+        result[0] = kernel_ies_interp(nullptr, slot, s, t);
+      return true;
+    }
+    default: {
+      return false;
+    }
+  }
 }
 
 ccl_device_extern bool osl_texture3d(ccl_private ShaderGlobals *sg,
@@ -1586,15 +1594,12 @@ ccl_device_extern bool osl_texture3d(ccl_private ShaderGlobals *sg,
                                      ccl_private float *dalphady,
                                      ccl_private void *errormessage)
 {
-  if (!texture_handle) {
+  const unsigned int type = OSL_TEXTURE_HANDLE_TYPE(texture_handle);
-    return false;
+  const unsigned int slot = OSL_TEXTURE_HANDLE_SLOT(texture_handle);
-  }
-
-  /* Only SVM textures are supported. */
-  int id = static_cast<int>(reinterpret_cast<size_t>(texture_handle) - 1);
-
-  const float4 rgba = kernel_tex_image_interp_3d(nullptr, id, *P, INTERPOLATION_NONE);
 
+  switch (type) {
+    case OSL_TEXTURE_HANDLE_TYPE_SVM: {
+      const float4 rgba = kernel_tex_image_interp_3d(nullptr, slot, *P, INTERPOLATION_NONE);
       if (nchannels > 0)
         result[0] = rgba.x;
       if (nchannels > 1)
@@ -1603,8 +1608,12 @@ ccl_device_extern bool osl_texture3d(ccl_private ShaderGlobals *sg,
         result[2] = rgba.z;
       if (alpha)
         *alpha = rgba.w;
-
       return true;
+    }
+    default: {
+      return false;
+    }
+  }
 }
 
 ccl_device_extern bool osl_environment(ccl_private ShaderGlobals *sg,

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

@@ -111,8 +111,8 @@ shader node_principled_bsdf(string distribution = "Multiscatter GGX",
     float eta = backfacing() ? 1.0 / f : f;
 
     if (distribution == "GGX" || Roughness <= 5e-2) {
-      float cosNO = dot(Normal, I);
+      float cosNI = dot(Normal, I);
-      float Fr = fresnel_dielectric_cos(cosNO, eta);
+      float Fr = fresnel_dielectric_cos(cosNI, eta);
 
       float refl_roughness = Roughness;
       if (Roughness <= 1e-2)

intern/cycles/kernel/osl/shaders/node_sky_texture.osl

@@ -135,8 +135,9 @@ color sky_radiance_nishita(vector dir, float nishita_data[10], string filename)
     float half_angular = angular_diameter / 2.0;
     float dir_elevation = M_PI_2 - direction[0];
 
-    /* if ray inside sun disc render it, otherwise render sky */
+    /* if ray inside sun disc render it, otherwise render sky.
-    if (sun_dir_angle < half_angular && sun_disc == 1) {
+     * alternatively, ignore the sun if we're evaluating the background texture. */
+    if (sun_dir_angle < half_angular && sun_disc == 1 && raytype("importance_bake") != 1) {
       /* get 2 pixels data */
       color pixel_bottom = color(nishita_data[0], nishita_data[1], nishita_data[2]);
       color pixel_top = color(nishita_data[3], nishita_data[4], nishita_data[5]);

intern/cycles/kernel/osl/types.h

@@ -96,4 +96,13 @@ struct OSLNoiseOptions {
 struct OSLTextureOptions {
 };
 
+#define OSL_TEXTURE_HANDLE_TYPE_IES ((uintptr_t)0x2 << 30)
+#define OSL_TEXTURE_HANDLE_TYPE_SVM ((uintptr_t)0x1 << 30)
+#define OSL_TEXTURE_HANDLE_TYPE_AO_OR_BEVEL ((uintptr_t)0x3 << 30)
+
+#define OSL_TEXTURE_HANDLE_TYPE(handle) \
+  ((unsigned int)((uintptr_t)(handle) & ((uintptr_t)0x3 << 30)))
+#define OSL_TEXTURE_HANDLE_SLOT(handle) \
+  ((unsigned int)((uintptr_t)(handle) & ((uintptr_t)0x3FFFFFFF)))
+
 CCL_NAMESPACE_END

intern/cycles/kernel/sample/mapping.h

@@ -33,19 +33,19 @@ ccl_device void make_orthonormals_tangent(const float3 N,
 
 /* sample direction with cosine weighted distributed in hemisphere */
 ccl_device_inline void sample_cos_hemisphere(
-    const float3 N, float randu, float randv, ccl_private float3 *omega_in, ccl_private float *pdf)
+    const float3 N, float randu, float randv, ccl_private float3 *wo, ccl_private float *pdf)
 {
   to_unit_disk(&randu, &randv);
   float costheta = sqrtf(max(1.0f - randu * randu - randv * randv, 0.0f));
   float3 T, B;
   make_orthonormals(N, &T, &B);
-  *omega_in = randu * T + randv * B + costheta * N;
+  *wo = randu * T + randv * B + costheta * N;
   *pdf = costheta * M_1_PI_F;
 }
 
 /* sample direction uniformly distributed in hemisphere */
 ccl_device_inline void sample_uniform_hemisphere(
-    const float3 N, float randu, float randv, ccl_private float3 *omega_in, ccl_private float *pdf)
+    const float3 N, float randu, float randv, ccl_private float3 *wo, ccl_private float *pdf)
 {
   float z = randu;
   float r = sqrtf(max(0.0f, 1.0f - z * z));
@@ -55,7 +55,7 @@ ccl_device_inline void sample_uniform_hemisphere(
 
   float3 T, B;
   make_orthonormals(N, &T, &B);
-  *omega_in = x * T + y * B + z * N;
+  *wo = x * T + y * B + z * N;
   *pdf = 0.5f * M_1_PI_F;
 }
 
@@ -64,20 +64,21 @@ ccl_device_inline void sample_uniform_cone(const float3 N,
                                            float angle,
                                            float randu,
                                            float randv,
-                                           ccl_private float3 *omega_in,
+                                           ccl_private float3 *wo,
                                            ccl_private float *pdf)
 {
-  float zMin = cosf(angle);
+  const float cosThetaMin = cosf(angle);
-  float z = zMin - zMin * randu + randu;
+  const float cosTheta = mix(cosThetaMin, 1.0f, randu);
-  float r = safe_sqrtf(1.0f - sqr(z));
+  const float sinTheta = sin_from_cos(cosTheta);
-  float phi = M_2PI_F * randv;
+  const float phi = M_2PI_F * randv;
-  float x = r * cosf(phi);
+  const float x = sinTheta * cosf(phi);
-  float y = r * sinf(phi);
+  const float y = sinTheta * sinf(phi);
+  const float z = cosTheta;
 
   float3 T, B;
   make_orthonormals(N, &T, &B);
-  *omega_in = x * T + y * B + z * N;
+  *wo = x * T + y * B + z * N;
-  *pdf = M_1_2PI_F / (1.0f - zMin);
+  *pdf = M_1_2PI_F / (1.0f - cosThetaMin);
 }
 
 ccl_device_inline float pdf_uniform_cone(const float3 N, float3 D, float angle)

intern/cycles/kernel/svm/brick.h

@@ -46,17 +46,8 @@ ccl_device_noinline_cpu float2 svm_brick(float3 p,
   float tint = saturatef((brick_noise((rownum << 16) + (bricknum & 0xFFFF)) + bias));
   float min_dist = min(min(x, y), min(brick_width - x, row_height - y));
 
-  float mortar;
-  if (min_dist >= mortar_size) {
-    mortar = 0.0f;
-  }
-  else if (mortar_smooth == 0.0f) {
-    mortar = 1.0f;
-  }
-  else {
   min_dist = 1.0f - min_dist / mortar_size;
-    mortar = (min_dist < mortar_smooth) ? smoothstepf(min_dist / mortar_smooth) : 1.0f;
+  float mortar = smoothstepf(min_dist / mortar_smooth);
-  }
 
   return make_float2(tint, mortar);
 }

intern/cycles/kernel/svm/closure.h

@@ -102,7 +102,7 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
 
   float3 N = stack_valid(data_node.x) ? stack_load_float3(stack, data_node.x) : sd->N;
   if (!(sd->type & PRIMITIVE_CURVE)) {
-    N = ensure_valid_reflection(sd->Ng, sd->I, N);
+    N = ensure_valid_reflection(sd->Ng, sd->wi, N);
   }
 
   float param1 = (stack_valid(param1_offset)) ? stack_load_float(stack, param1_offset) :
@@ -162,8 +162,8 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
       float ior = (sd->flag & SD_BACKFACING) ? 1.0f / eta : eta;
 
       // calculate fresnel for refraction
-      float cosNO = dot(N, sd->I);
+      float cosNI = dot(N, sd->wi);
-      float fresnel = fresnel_dielectric_cos(cosNO, ior);
+      float fresnel = fresnel_dielectric_cos(cosNI, ior);
 
       // calculate weights of the diffuse and specular part
       float diffuse_weight = (1.0f - saturatef(metallic)) * (1.0f - saturatef(transmission));
@@ -185,7 +185,7 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
                                     stack_load_float3(stack, data_cn_ssr.x) :
                                     sd->N;
       if (!(sd->type & PRIMITIVE_CURVE)) {
-        clearcoat_normal = ensure_valid_reflection(sd->Ng, sd->I, clearcoat_normal);
+        clearcoat_normal = ensure_valid_reflection(sd->Ng, sd->wi, clearcoat_normal);
       }
       float3 subsurface_radius = stack_valid(data_cn_ssr.y) ?
                                      stack_load_float3(stack, data_cn_ssr.y) :
@@ -333,7 +333,6 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
             bsdf->extra->cspec0 = rgb_to_spectrum(
                 (specular * 0.08f * tmp_col) * (1.0f - metallic) + base_color * metallic);
             bsdf->extra->color = rgb_to_spectrum(base_color);
-            bsdf->extra->clearcoat = 0.0f;
 
             /* setup bsdf */
             if (distribution == CLOSURE_BSDF_MICROFACET_GGX_GLASS_ID ||
@@ -383,7 +382,6 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
 
                 bsdf->extra->color = rgb_to_spectrum(base_color);
                 bsdf->extra->cspec0 = rgb_to_spectrum(cspec0);
-                bsdf->extra->clearcoat = 0.0f;
 
                 /* setup bsdf */
                 sd->flag |= bsdf_microfacet_ggx_fresnel_setup(bsdf, sd);
@@ -440,7 +438,6 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
 
               bsdf->extra->color = rgb_to_spectrum(base_color);
               bsdf->extra->cspec0 = rgb_to_spectrum(cspec0);
-              bsdf->extra->clearcoat = 0.0f;
 
               /* setup bsdf */
               sd->flag |= bsdf_microfacet_multi_ggx_glass_fresnel_setup(bsdf, sd);
@@ -455,31 +452,21 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
 #ifdef __CAUSTICS_TRICKS__
       if (kernel_data.integrator.caustics_reflective || (path_flag & PATH_RAY_DIFFUSE) == 0) {
 #endif
-        if (clearcoat > CLOSURE_WEIGHT_CUTOFF) {
+        Spectrum clearcoat_weight = 0.25f * clearcoat * weight;
         ccl_private MicrofacetBsdf *bsdf = (ccl_private MicrofacetBsdf *)bsdf_alloc(
-              sd, sizeof(MicrofacetBsdf), weight);
+            sd, sizeof(MicrofacetBsdf), clearcoat_weight);
-          ccl_private MicrofacetExtra *extra =
-              (bsdf != NULL) ?
-                  (ccl_private MicrofacetExtra *)closure_alloc_extra(sd, sizeof(MicrofacetExtra)) :
-                  NULL;
 
-          if (bsdf && extra) {
+        if (bsdf) {
           bsdf->N = clearcoat_normal;
           bsdf->T = zero_float3();
           bsdf->ior = 1.5f;
-            bsdf->extra = extra;
 
           bsdf->alpha_x = clearcoat_roughness * clearcoat_roughness;
           bsdf->alpha_y = clearcoat_roughness * clearcoat_roughness;
 
-            bsdf->extra->color = zero_spectrum();
-            bsdf->extra->cspec0 = make_spectrum(0.04f);
-            bsdf->extra->clearcoat = clearcoat;
-
           /* setup bsdf */
           sd->flag |= bsdf_microfacet_ggx_clearcoat_setup(bsdf, sd);
         }
-        }
 #ifdef __CAUSTICS_TRICKS__
       }
 #endif
@@ -584,7 +571,6 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
         if (bsdf->extra) {
           bsdf->extra->color = rgb_to_spectrum(stack_load_float3(stack, data_node.w));
           bsdf->extra->cspec0 = zero_spectrum();
-          bsdf->extra->clearcoat = 0.0f;
           sd->flag |= bsdf_microfacet_multi_ggx_setup(bsdf);
         }
       }
@@ -652,8 +638,8 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
       eta = (sd->flag & SD_BACKFACING) ? 1.0f / eta : eta;
 
       /* fresnel */
-      float cosNO = dot(N, sd->I);
+      float cosNI = dot(N, sd->wi);
-      float fresnel = fresnel_dielectric_cos(cosNO, eta);
+      float fresnel = fresnel_dielectric_cos(cosNI, eta);
       float roughness = sqr(param1);
 
       /* reflection */
@@ -724,7 +710,6 @@ ccl_device_noinline int svm_node_closure_bsdf(KernelGlobals kg,
       kernel_assert(stack_valid(data_node.z));
       bsdf->extra->color = rgb_to_spectrum(stack_load_float3(stack, data_node.z));
       bsdf->extra->cspec0 = zero_spectrum();
-      bsdf->extra->clearcoat = 0.0f;
 
       /* setup bsdf */
       sd->flag |= bsdf_microfacet_multi_ggx_glass_setup(bsdf);

intern/cycles/kernel/svm/displace.h

@@ -71,7 +71,7 @@ ccl_device_noinline void svm_node_set_bump(KernelGlobals kg,
       object_normal_transform(kg, sd, &normal_out);
     }
 
-    normal_out = ensure_valid_reflection(sd->Ng, sd->I, normal_out);
+    normal_out = ensure_valid_reflection(sd->Ng, sd->wi, normal_out);
     stack_store_float3(stack, node.w, normal_out);
   }
   else

intern/cycles/kernel/svm/fresnel.h

@@ -22,7 +22,7 @@ ccl_device_noinline void svm_node_fresnel(ccl_private ShaderData *sd,
   eta = fmaxf(eta, 1e-5f);
   eta = (sd->flag & SD_BACKFACING) ? 1.0f / eta : eta;
 
-  float f = fresnel_dielectric_cos(dot(sd->I, normal_in), eta);
+  float f = fresnel_dielectric_cos(dot(sd->wi, normal_in), eta);
 
   stack_store_float(stack, out_offset, f);
 }
@@ -50,10 +50,10 @@ ccl_device_noinline void svm_node_layer_weight(ccl_private ShaderData *sd,
     float eta = fmaxf(1.0f - blend, 1e-5f);
     eta = (sd->flag & SD_BACKFACING) ? eta : 1.0f / eta;
 
-    f = fresnel_dielectric_cos(dot(sd->I, normal_in), eta);
+    f = fresnel_dielectric_cos(dot(sd->wi, normal_in), eta);
   }
   else {
-    f = fabsf(dot(sd->I, normal_in));
+    f = fabsf(dot(sd->wi, normal_in));
 
     if (blend != 0.5f) {
       blend = clamp(blend, 0.0f, 1.0f - 1e-5f);

intern/cycles/kernel/svm/geometry.h

@@ -28,7 +28,7 @@ ccl_device_noinline void svm_node_geometry(KernelGlobals kg,
       break;
 #endif
     case NODE_GEOM_I:
-      data = sd->I;
+      data = sd->wi;
       break;
     case NODE_GEOM_Ng:
       data = sd->Ng;

intern/cycles/kernel/svm/ies.h

@@ -84,6 +84,7 @@ ccl_device_inline float kernel_ies_interp(KernelGlobals kg, int slot, float h_an
   return max(cubic_interp(a, b, c, d, h_frac), 0.0f);
 }
 
+#ifdef __SVM__
 ccl_device_noinline void svm_node_ies(KernelGlobals kg,
                                       ccl_private ShaderData *sd,
                                       ccl_private float *stack,
@@ -105,5 +106,6 @@ ccl_device_noinline void svm_node_ies(KernelGlobals kg,
     stack_store_float(stack, fac_offset, fac);
   }
 }
+#endif
 
 CCL_NAMESPACE_END

intern/cycles/kernel/svm/sky.h

@@ -118,6 +118,7 @@ ccl_device float3 geographical_to_direction(float lat, float lon)
 
 ccl_device float3 sky_radiance_nishita(KernelGlobals kg,
                                        float3 dir,
+                                       uint32_t path_flag,
                                        float3 pixel_bottom,
                                        float3 pixel_top,
                                        ccl_private float *nishita_data,
@@ -140,8 +141,9 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals kg,
     float half_angular = angular_diameter / 2.0f;
     float dir_elevation = M_PI_2_F - direction.x;
 
-    /* if ray inside sun disc render it, otherwise render sky */
+    /* if ray inside sun disc render it, otherwise render sky.
-    if (sun_disc && sun_dir_angle < half_angular) {
+     * alternatively, ignore the sun if we're evaluating the background texture. */
+    if (sun_disc && sun_dir_angle < half_angular && !(path_flag & PATH_RAY_IMPORTANCE_BAKE)) {
       /* get 2 pixels data */
       float y;
 
@@ -197,8 +199,12 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals kg,
   return xyz_to_rgb_clamped(kg, xyz);
 }
 
-ccl_device_noinline int svm_node_tex_sky(
+ccl_device_noinline int svm_node_tex_sky(KernelGlobals kg,
-    KernelGlobals kg, ccl_private ShaderData *sd, ccl_private float *stack, uint4 node, int offset)
+                                         ccl_private ShaderData *sd,
+                                         uint32_t path_flag,
+                                         ccl_private float *stack,
+                                         uint4 node,
+                                         int offset)
 {
   /* Load data */
   uint dir_offset = node.y;
@@ -310,7 +316,8 @@ ccl_device_noinline int svm_node_tex_sky(
     uint texture_id = __float_as_uint(data.z);
 
     /* Compute Sky */
-    f = sky_radiance_nishita(kg, dir, pixel_bottom, pixel_top, nishita_data, texture_id);
+    f = sky_radiance_nishita(
+        kg, dir, path_flag, pixel_bottom, pixel_top, nishita_data, texture_id);
   }
 
   stack_store_float3(stack, out_offset, f);