archive/blender-archive
Archived
1
1
Fork 0
Code Pull Requests Activity

Merge remote-tracking branch 'origin' into sculpt-dev

Browse Source
This commit is contained in:
Joseph Eagar
2023-02-07 17:45:51 -08:00
parent 4e85e0b98e 0a3df611e7
commit e9151dc96a
1293 changed files with 61217 additions and 39863 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
.gitea/default_merge_message/REBASE_TEMPLATE.md Normal file
View File

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

3
.gitea/default_merge_message/SQUASH_TEMPLATE.md Normal file
View File

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

45
.gitea/issue_template.yaml
View File

@@ -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).

44
.gitea/issue_template/bug.yaml Normal file
View File

@@ -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!

1
.gitea/issue_template/config.yaml Normal file
View File

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

10
.gitea/issue_template/design.yaml Normal file
View File

@@ -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

10
.gitea/issue_template/todo.yaml Normal file
View File

@@ -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

4
.gitea/pull_request_template.md
View File

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

17
.gitea/pull_request_template.yaml Normal file
View File

@@ -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

104
CMakeLists.txt
View File

@@ -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)
# In future, can be used with `quicklookthumbnailing/qlthumbnailreply` to create file
# thumbnails for say Finder. Turn it off for now.
option(WITH_BLENDER_THUMBNAILER "Build \"blender-thumbnailer\" thumbnail extraction utility" OFF)
elseif(WIN32)
option(WITH_BLENDER_THUMBNAILER "Build \"BlendThumb.dll\" helper for Windows explorer integration" ON)
if(WIN32)
option(WITH_BLENDER_THUMBNAILER "\
Build \"BlendThumb.dll\" helper for Windows explorer integration to support extracting \
thumbnails from `.blend` files."
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)
mark_as_advanced(WITH_XR_OPENXR)
else()
option(WITH_XR_OPENXR "Enable VR features through the OpenXR specification" ON)
set(_option_default OFF)
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,12 +370,13 @@ 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)
else()
option(WITH_JACK "Enable JACK Support (http://www.jackaudio.org)" ON)
set(_option_default OFF)
endif()
option(WITH_JACK_DYNLOAD "Enable runtime dynamic JACK libraries loading" OFF)
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)
endif()
@@ -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()
@@ -1002,6 +1035,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)
@@ -1215,13 +1250,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
@@ -1271,12 +1299,14 @@ endif()
# -----------------------------------------------------------------------------
# Configure Bullet
if(WITH_BULLET AND WITH_SYSTEM_BULLET)
find_package(Bullet)
set_and_warn_library_found("Bullet" BULLET_FOUND WITH_BULLET)
else()
set(BULLET_INCLUDE_DIRS "${CMAKE_SOURCE_DIR}/extern/bullet2/src")
# set(BULLET_LIBRARIES "")
if(WITH_BULLET)
if(WITH_SYSTEM_BULLET)
find_package(Bullet)
set_and_warn_library_found("Bullet" BULLET_FOUND WITH_BULLET)
else()
set(BULLET_INCLUDE_DIRS "${CMAKE_SOURCE_DIR}/extern/bullet2/src")
set(BULLET_LIBRARIES "extern_bullet")
endif()
endif()
@@ -1436,6 +1466,9 @@ if(CMAKE_COMPILER_IS_GNUCC)
add_check_c_compiler_flag(C_WARNINGS C_WARN_TYPE_LIMITS -Wtype-limits)
add_check_c_compiler_flag(C_WARNINGS C_WARN_FORMAT_SIGN -Wformat-signedness)
add_check_c_compiler_flag(C_WARNINGS C_WARN_RESTRICT -Wrestrict)
# Useful but too many false positives and inconvenient to suppress each occurrence.
add_check_c_compiler_flag(C_WARNINGS C_WARN_NO_STRINGOP_OVERREAD -Wno-stringop-overread)
add_check_c_compiler_flag(C_WARNINGS C_WARN_NO_STRINGOP_OVERFLOW -Wno-stringop-overflow)
# C-only.
add_check_c_compiler_flag(C_WARNINGS C_WARN_NO_NULL -Wnonnull)
@@ -1475,6 +1508,9 @@ if(CMAKE_COMPILER_IS_GNUCC)
add_check_cxx_compiler_flag(CXX_WARNINGS CXX_WARN_RESTRICT -Wrestrict)
add_check_cxx_compiler_flag(CXX_WARNINGS CXX_WARN_NO_SUGGEST_OVERRIDE -Wno-suggest-override)
add_check_cxx_compiler_flag(CXX_WARNINGS CXX_WARN_UNINITIALIZED -Wuninitialized)
# Useful but too many false positives and inconvenient to suppress each occurrence.
add_check_cxx_compiler_flag(CXX_WARNINGS CXX_WARN_NO_STRINGOP_OVERREAD -Wno-stringop-overread)
add_check_cxx_compiler_flag(CXX_WARNINGS CXX_WARN_NO_STRINGOP_OVERFLOW -Wno-stringop-overflow)
# causes too many warnings
if(NOT APPLE)

11
GNUmakefile
View File

@@ -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) \

38
README.md Normal file
View File

@@ -0,0 +1,38 @@
<!--
Keep this document short & concise,
linking to external resources instead of including content in-line.
See 'release/text/readme.html' for the end user read-me.
-->
Blender
=======
Blender is the free and open source 3D creation suite.
It supports the entirety of the 3D pipeline-modeling, rigging, animation, simulation, rendering, compositing,
motion tracking and video editing.
![Blender screenshot](https://code.blender.org/wp-content/uploads/2018/12/springrg.jpg "Blender screenshot")
Project Pages
-------------
- [Main Website](http://www.blender.org)
- [Reference Manual](https://docs.blender.org/manual/en/latest/index.html)
- [User Community](https://www.blender.org/community/)
Development
-----------
- [Build Instructions](https://wiki.blender.org/wiki/Building_Blender)
- [Code Review & Bug Tracker](https://developer.blender.org)
- [Developer Forum](https://devtalk.blender.org)
- [Developer Documentation](https://wiki.blender.org)
License
-------
Blender as a whole is licensed under the GNU General Public License, Version 3.
Individual files may have a different, but compatible license.
See [blender.org/about/license](https://www.blender.org/about/license) for details.

2
build_files/build_environment/cmake/dpcpp.cmake
View File

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

2
build_files/build_environment/cmake/llvm.cmake
View File

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

8
build_files/build_environment/cmake/ssl.cmake
View File

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

2
build_files/build_environment/cmake/usd.cmake
View File

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

2
build_files/build_environment/cmake/xml2.cmake
View File

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

6
build_files/cmake/Modules/FindMoltenVK.cmake
View File

@@ -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

6
build_files/cmake/Modules/FindOptiX.cmake
View File

@@ -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

2
build_files/cmake/Modules/FindPythonLibsUnix.cmake
View File

@@ -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}"

63
build_files/cmake/Modules/FindShaderC.cmake Normal file
View File

@@ -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)

63
build_files/cmake/Modules/FindVulkan.cmake Normal file
View File

@@ -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)

78
build_files/cmake/cmake_print_build_options.py
View File

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

10
build_files/cmake/macros.cmake
View File

@@ -550,7 +550,9 @@ function(setup_platform_linker_libs
endif()
if(WIN32 AND NOT UNIX)
target_link_libraries(${target} ${PTHREADS_LIBRARIES})
if(DEFINED PTHREADS_LIBRARIES)
target_link_libraries(${target} ${PTHREADS_LIBRARIES})
endif()
endif()
# target_link_libraries(${target} ${PLATFORM_LINKLIBS} ${CMAKE_DL_LIBS})
@@ -1115,7 +1117,7 @@ function(find_python_package
# endif()
# Not set, so initialize.
else()
string(REPLACE "." ";" _PY_VER_SPLIT "${PYTHON_VERSION}")
string(REPLACE "." ";" _PY_VER_SPLIT "${PYTHON_VERSION}")
list(GET _PY_VER_SPLIT 0 _PY_VER_MAJOR)
# re-cache
@@ -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)
@@ -1275,7 +1277,7 @@ macro(add_bundled_libraries library_dir)
list(APPEND PLATFORM_BUNDLED_LIBRARY_DIRS ${_library_dir})
unset(_all_library_versions)
unset(_library_dir)
endif()
endif()
endmacro()
macro(windows_install_shared_manifest)

16
build_files/cmake/platform/platform_apple.cmake
View File

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

7
build_files/cmake/platform/platform_old_libs_update.cmake
View File

@@ -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)

97
build_files/cmake/platform/platform_unix.cmake
View File

@@ -4,38 +4,52 @@
# Libraries configuration for any *nix system including Linux and Unix (excluding APPLE).
# Detect precompiled library directory
if(NOT DEFINED LIBDIR)
# Path to a locally compiled libraries.
set(LIBDIR_NAME ${CMAKE_SYSTEM_NAME}_${CMAKE_SYSTEM_PROCESSOR})
string(TOLOWER ${LIBDIR_NAME} LIBDIR_NAME)
set(LIBDIR_NATIVE_ABI ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_NAME})
# Path to precompiled libraries with known glibc 2.28 ABI.
set(LIBDIR_GLIBC228_ABI ${CMAKE_SOURCE_DIR}/../lib/linux_x86_64_glibc_228)
if(NOT WITH_LIBS_PRECOMPILED)
unset(LIBDIR)
else()
if(NOT DEFINED LIBDIR)
# Path to a locally compiled libraries.
set(LIBDIR_NAME ${CMAKE_SYSTEM_NAME}_${CMAKE_SYSTEM_PROCESSOR})
string(TOLOWER ${LIBDIR_NAME} LIBDIR_NAME)
set(LIBDIR_NATIVE_ABI ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_NAME})
# Choose the best suitable libraries.
if(EXISTS ${LIBDIR_NATIVE_ABI})
set(LIBDIR ${LIBDIR_NATIVE_ABI})
set(WITH_LIBC_MALLOC_HOOK_WORKAROUND True)
elseif(EXISTS ${LIBDIR_GLIBC228_ABI})
set(LIBDIR ${LIBDIR_GLIBC228_ABI})
if(WITH_MEM_JEMALLOC)
# jemalloc provides malloc hooks.
set(WITH_LIBC_MALLOC_HOOK_WORKAROUND False)
else()
# Path to precompiled libraries with known glibc 2.28 ABI.
set(LIBDIR_GLIBC228_ABI ${CMAKE_SOURCE_DIR}/../lib/linux_x86_64_glibc_228)
# Choose the best suitable libraries.
if(EXISTS ${LIBDIR_NATIVE_ABI})
set(LIBDIR ${LIBDIR_NATIVE_ABI})
set(WITH_LIBC_MALLOC_HOOK_WORKAROUND True)
elseif(EXISTS ${LIBDIR_GLIBC228_ABI})
set(LIBDIR ${LIBDIR_GLIBC228_ABI})
if(WITH_MEM_JEMALLOC)
# jemalloc provides malloc hooks.
set(WITH_LIBC_MALLOC_HOOK_WORKAROUND False)
else()
set(WITH_LIBC_MALLOC_HOOK_WORKAROUND True)
endif()
endif()
# Avoid namespace pollustion.
unset(LIBDIR_NATIVE_ABI)
unset(LIBDIR_GLIBC228_ABI)
endif()
# Avoid namespace pollustion.
unset(LIBDIR_NATIVE_ABI)
unset(LIBDIR_GLIBC228_ABI)
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})
# used by other platforms
set(PTHREADS_LIBRARIES ${CMAKE_THREAD_LIBS_INIT})
# `FindThreads` documentation notes that this may be empty
# with the system libraries provide threading functionality.
if(CMAKE_THREAD_LIBS_INIT)
list(APPEND PLATFORM_LINKLIBS ${CMAKE_THREAD_LIBS_INIT})
# used by other platforms
set(PTHREADS_LIBRARIES ${CMAKE_THREAD_LIBS_INIT})
endif()
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)

6
build_files/config/pipeline_config.yaml
View File

@@ -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

4
build_files/utils/make_bpy_wheel.py
View File

@@ -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

4
build_files/utils/make_test.py
View File

@@ -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")

19
build_files/utils/make_update.py
View File

@@ -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("."):

4
doc/python_api/sphinx_doc_gen.py
View File

@@ -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()

4
extern/audaspace/CMakeLists.txt vendored
View File

@@ -513,17 +513,19 @@ if(WITH_FFTW)
src/fx/Convolver.cpp
src/fx/ConvolverReader.cpp
src/fx/ConvolverSound.cpp
src/fx/Equalizer.cpp
src/fx/FFTConvolver.cpp
src/fx/HRTF.cpp
src/fx/ImpulseResponse.cpp
src/util/FFTPlan.cpp
)
set(FFTW_HDR
include/fx/BinauralSound.h
include/fx/BinauralSound.h
include/fx/BinauralReader.h
include/fx/Convolver.h
include/fx/ConvolverReader.h
include/fx/ConvolverSound.h
include/fx/Equalizer.h
include/fx/FFTConvolver.h
include/fx/HRTF.h
include/fx/HRTFLoader.h

11
extern/audaspace/bindings/C/AUD_Sound.cpp vendored
View File

@@ -54,6 +54,7 @@
#ifdef WITH_CONVOLUTION
#include "fx/BinauralSound.h"
#include "fx/ConvolverSound.h"
#include "fx/Equalizer.h"
#endif
#include <cassert>
@@ -768,4 +769,14 @@ AUD_API AUD_Sound* AUD_Sound_Binaural(AUD_Sound* sound, AUD_HRTF* hrtfs, AUD_Sou
}
}
AUD_API AUD_Sound* AUD_Sound_equalize(AUD_Sound* sound, float *definition, int size, float maxFreqEq, int sizeConversion)
{
assert(sound);
std::shared_ptr<Buffer> buf = std::shared_ptr<Buffer>(new Buffer(sizeof(float)*size));
std::memcpy(buf->getBuffer(), definition, sizeof(float)*size);
AUD_Sound *equalizer=new AUD_Sound(new Equalizer(*sound, buf, size, maxFreqEq, sizeConversion));
return equalizer;
}
#endif

10
extern/audaspace/bindings/C/AUD_Sound.h vendored
View File

@@ -397,6 +397,16 @@ extern AUD_API AUD_Sound* AUD_Sound_mutable(AUD_Sound* sound);
#ifdef WITH_CONVOLUTION
extern AUD_API AUD_Sound* AUD_Sound_Convolver(AUD_Sound* sound, AUD_ImpulseResponse* filter, AUD_ThreadPool* threadPool);
extern AUD_API AUD_Sound* AUD_Sound_Binaural(AUD_Sound* sound, AUD_HRTF* hrtfs, AUD_Source* source, AUD_ThreadPool* threadPool);
/**
* Creates an Equalizer for the sound
* \param sound The handle of the sound
* \param definition buffer of size*sizeof(float) with the array of equalization values
* \param maxFreqEq Maximum frequency refered by the array
* \param sizeConversion Size of the transformation. Must be 2^number (for example 1024, 2048,...)
* \return A handle to the Equalizer refered to that sound
*/
extern AUD_API AUD_Sound* AUD_Sound_equalize(AUD_Sound* sound, float *definition, int size, float maxFreqEq, int sizeConversion);
#endif
#ifdef __cplusplus

1
extern/audaspace/bindings/C/AUD_Special.h vendored
View File

@@ -53,6 +53,7 @@ extern AUD_API AUD_Handle* AUD_pauseAfter(AUD_Handle* handle, double seconds);
* \param buffer The buffer to write to. Must have a size of 3*4*length.
* \param length How many samples to read from the sound.
* \param samples_per_second How many samples to read per second of the sound.
* \param interrupt Must point to a short that equals 0. If it is set to a non-zero value, the method will be interrupted and return 0.
* \return How many samples really have been read. Always <= length.
*/
extern AUD_API int AUD_readSound(AUD_Sound* sound, float* buffer, int length, int samples_per_second, short* interrupt);

6
extern/audaspace/bindings/python/setup.py.in vendored
View File

@@ -5,12 +5,12 @@ import os
import codecs
import numpy
from distutils.core import setup, Extension
from setuptools import setup, Extension
if len(sys.argv) > 2 and sys.argv[1] == '--build-docs':
import subprocess
from distutils.core import Distribution
from distutils.command.build import build
from setuptools import Distribution
from setuptools.command.build import build
dist = Distribution()
cmd = build(dist)

106
extern/audaspace/include/fx/Equalizer.h vendored Normal file
View File

@@ -0,0 +1,106 @@
/*******************************************************************************
* Copyright 2022 Marcos Perez Gonzalez
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
******************************************************************************/
#pragma once
/**
* @file Equalizer.h
* @ingroup fx
* The Equalizer class.
*/
#include <memory>
#include <vector>
#include "ISound.h"
#include "ImpulseResponse.h"
AUD_NAMESPACE_BEGIN
class Buffer;
class ImpulseResponse;
/**
* This class represents a sound that can be modified depending on a given impulse response.
*/
class AUD_API Equalizer : public ISound
{
private:
/**
* A pointer to the imput sound.
*/
std::shared_ptr<ISound> m_sound;
/**
* Local definition of Equalizer
*/
std::shared_ptr<Buffer> m_bufEQ;
/**
* A pointer to the impulse response.
*/
std::shared_ptr<ImpulseResponse> m_impulseResponse;
/**
* delete copy constructor and operator=
*/
Equalizer(const Equalizer&) = delete;
Equalizer& operator=(const Equalizer&) = delete;
/**
* Create ImpulseResponse from the definition in the Buffer,
* using at the end a minimum phase change
*/
std::shared_ptr<ImpulseResponse> createImpulseResponse();
/**
* Create an Impulse Response with minimum phase distortion using Homomorphic
* The input is an Impulse Response
*/
std::shared_ptr<Buffer> minimumPhaseFilterHomomorphic(std::shared_ptr<Buffer> original, int lOriginal, int lWork);
/**
* Create an Impulse Response with minimum phase distortion using Hilbert
* The input is an Impulse Response
*/
std::shared_ptr<Buffer> minimumPhaseFilterHilbert(std::shared_ptr<Buffer> original, int lOriginal, int lWork);
public:
/**
* Creates a new Equalizer.
* \param sound The sound that will be equalized
*/
Equalizer(std::shared_ptr<ISound> sound, std::shared_ptr<Buffer> bufEQ, int externalSizeEq, float maxFreqEq, int sizeConversion);
virtual ~Equalizer();
virtual std::shared_ptr<IReader> createReader();
/*
* Length of the external equalizer definition. It must be the number of "float" positions of the Buffer
*/
int external_size_eq;
/*
* Length of the internal equalizer definition
*/
int filter_length;
/*
* Maximum frequency used in the equalizer definition
*/
float maxFreqEq;
};
AUD_NAMESPACE_END

2
extern/audaspace/src/fx/BinauralReader.cpp vendored
View File

@@ -27,7 +27,7 @@
AUD_NAMESPACE_BEGIN
BinauralReader::BinauralReader(std::shared_ptr<IReader> reader, std::shared_ptr<HRTF> hrtfs, std::shared_ptr<Source> source, std::shared_ptr<ThreadPool> threadPool, std::shared_ptr<FFTPlan> plan) :
m_reader(reader), m_hrtfs(hrtfs), m_source(source), m_N(plan->getSize()), m_threadPool(threadPool), m_position(0), m_eosReader(false), m_eosTail(false), m_transition(false), m_transPos(CROSSFADE_SAMPLES*NUM_OUTCHANNELS)
m_position(0), m_reader(reader), m_hrtfs(hrtfs), m_source(source), m_N(plan->getSize()), m_transition(false), m_transPos(CROSSFADE_SAMPLES*NUM_OUTCHANNELS), m_eosReader(false), m_eosTail(false), m_threadPool(threadPool)
{
if(m_hrtfs->isEmpty())
AUD_THROW(StateException, "The provided HRTF object is empty");

2
extern/audaspace/src/fx/Convolver.cpp vendored
View File

@@ -23,7 +23,7 @@
AUD_NAMESPACE_BEGIN
Convolver::Convolver(std::shared_ptr<std::vector<std::shared_ptr<std::vector<std::complex<sample_t>>>>> ir, int irLength, std::shared_ptr<ThreadPool> threadPool, std::shared_ptr<FFTPlan> plan) :
m_N(plan->getSize()), m_M(plan->getSize()/2), m_L(plan->getSize()/2), m_irBuffers(ir), m_irLength(irLength), m_threadPool(threadPool), m_numThreads(std::min(threadPool->getNumOfThreads(), static_cast<unsigned int>(m_irBuffers->size() - 1))), m_tailCounter(0), m_eos(false)
m_N(plan->getSize()), m_M(plan->getSize()/2), m_L(plan->getSize()/2), m_irBuffers(ir), m_numThreads(std::min(threadPool->getNumOfThreads(), static_cast<unsigned int>(m_irBuffers->size() - 1))), m_threadPool(threadPool), m_irLength(irLength), m_tailCounter(0), m_eos(false)
{
m_resetFlag = false;

2
extern/audaspace/src/fx/ConvolverReader.cpp vendored
View File

@@ -24,7 +24,7 @@
AUD_NAMESPACE_BEGIN
ConvolverReader::ConvolverReader(std::shared_ptr<IReader> reader, std::shared_ptr<ImpulseResponse> ir, std::shared_ptr<ThreadPool> threadPool, std::shared_ptr<FFTPlan> plan) :
m_reader(reader), m_ir(ir), m_N(plan->getSize()), m_eosReader(false), m_eosTail(false), m_inChannels(reader->getSpecs().channels), m_irChannels(ir->getSpecs().channels), m_threadPool(threadPool), m_position(0)
m_position(0), m_reader(reader), m_ir(ir), m_N(plan->getSize()), m_eosReader(false), m_eosTail(false), m_inChannels(reader->getSpecs().channels), m_irChannels(ir->getSpecs().channels), m_threadPool(threadPool)
{
m_nChannelThreads = std::min((int)threadPool->getNumOfThreads(), m_inChannels);
m_futures.resize(m_nChannelThreads);

367
extern/audaspace/src/fx/Equalizer.cpp vendored Normal file
View File

@@ -0,0 +1,367 @@
/*******************************************************************************
* Copyright 2022 Marcos Perez Gonzalez
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
******************************************************************************/
#include "fx/Equalizer.h"
#include <chrono>
#include <cstring>
#include <fstream>
#include <iomanip>
#include <iostream>
#include <memory>
#include <string>
#include "Exception.h"
#include "fx/ConvolverReader.h"
#include "fx/ImpulseResponse.h"
#include "util/Buffer.h"
#include "util/FFTPlan.h"
#include "util/ThreadPool.h"
AUD_NAMESPACE_BEGIN
Equalizer::Equalizer(std::shared_ptr<ISound> sound, std::shared_ptr<Buffer> bufEQ, int externalSizeEq, float maxFreqEq, int sizeConversion) : m_sound(sound), m_bufEQ(bufEQ)
{
this->maxFreqEq = maxFreqEq;
this->external_size_eq = externalSizeEq;
filter_length = sizeConversion;
}
Equalizer::~Equalizer()
{
}
std::shared_ptr<IReader> Equalizer::createReader()
{
std::shared_ptr<FFTPlan> fp = std::shared_ptr<FFTPlan>(new FFTPlan(filter_length));
// 2 threads to start with
return std::shared_ptr<ConvolverReader>(new ConvolverReader(m_sound->createReader(), createImpulseResponse(), std::shared_ptr<ThreadPool>(new ThreadPool(2)), fp));
}
float calculateValueArray(float* data, float minX, float maxX, int length, float posX)
{
if(posX < minX)
return 1.0;
if(posX > maxX)
return data[length - 1];
float interval = (maxX - minX) / (float) length;
int idx = (int) ((posX - minX) / interval);
return data[idx];
}
void complex_prod(float a, float b, float c, float d, float* r, float* imag)
{
float prod1 = a * c;
float prod2 = b * d;
float prod3 = (a + b) * (c + d);
// Real Part
*r = prod1 - prod2;
// Imaginary Part
*imag = prod3 - (prod1 + prod2);
}
/**
* The creation of the ImpuseResponse which will be convoluted with the sound
*
* The implementation is based on scikit-signal
*/
std::shared_ptr<ImpulseResponse> Equalizer::createImpulseResponse()
{
std::shared_ptr<FFTPlan> fp = std::shared_ptr<FFTPlan>(new FFTPlan(filter_length));
fftwf_complex* buffer = (fftwf_complex*) fp->getBuffer();
std::memset(buffer, 0, filter_length * sizeof(fftwf_complex));
std::shared_ptr<IReader> soundReader = m_sound.get()->createReader();
Specs specsSound = soundReader.get()->getSpecs();
int sampleRate = specsSound.rate;
for(unsigned i = 0; i < filter_length / 2; i++)
{
double freq = (((float) i) / (float) filter_length) * (float) sampleRate;
double dbGain = calculateValueArray(m_bufEQ->getBuffer(), 0.0, maxFreqEq, external_size_eq, freq);
// gain = 10^(decibels / 20.0)
// 0 db = 1
// 20 db = 10
// 40 db = 100
float gain = (float) pow(10.0, dbGain / 20.0);
if(i == filter_length / 2 - 1)
{
gain = 0;
}
// IMPORTANT!!!! It is needed for the minimum phase step.
// Without this, the amplitude would be square rooted
//
gain *= gain;
// Calculation of exponential with std.. or "by hand"
/*
std::complex<float> preShift= std::complex<float>(0.0, -(filter_length - 1)
/ 2. * M_PI * freq / ( sampleRate/2)); std::complex<float> shift =
std::exp(preShift);
std::complex<float> cGain = gain * shift;
*/
float imaginary_shift = -(filter_length - 1) / 2. * M_PI * freq / (sampleRate / 2);
float cGain_real = gain * cos(imaginary_shift);
float cGain_imag = gain * sin(imaginary_shift);
int i2 = filter_length - i - 1;
buffer[i][0] = cGain_real; // Real
buffer[i][1] = cGain_imag; // Imag
if(i > 0 && i2 < filter_length)
{
buffer[i2][0] = cGain_real; // Real
buffer[i2][1] = cGain_imag; // Imag
}
}
// In place. From Complex to sample_t
fp->IFFT(buffer);
// Window Hamming
sample_t* pt_sample_t = (sample_t*) buffer;
float half_filter = ((float) filter_length) / 2.0;
for(int i = 0; i < filter_length; i++)
{
// Centered in filter_length/2
float window = 0.54 - 0.46 * cos((2 * M_PI * (float) i) / (float) (filter_length - 1));
pt_sample_t[i] *= window;
}
std::shared_ptr<Buffer> b2 = std::shared_ptr<Buffer>(new Buffer(filter_length * sizeof(sample_t)));
sample_t* buffer_real = (sample_t*) buffer;
sample_t* buffer2 = b2->getBuffer();
float normaliziter = (float) filter_length;
for(int i = 0; i < filter_length; i++)
{
buffer2[i] = (buffer_real[i] / normaliziter);
}
fp->freeBuffer(buffer);
//
// Here b2 is the buffer with a "valid" FIR (remember the squared amplitude
//
std::shared_ptr<Buffer> ir_minimum = minimumPhaseFilterHomomorphic(b2, filter_length, -1);
Specs specsIR;
specsIR.rate = sampleRate;
specsIR.channels = CHANNELS_MONO;
return std::shared_ptr<ImpulseResponse>(new ImpulseResponse(std::shared_ptr<StreamBuffer>(new StreamBuffer(ir_minimum, specsIR)), fp));
}
std::shared_ptr<Buffer> Equalizer::minimumPhaseFilterHomomorphic(std::shared_ptr<Buffer> original, int lOriginal, int lWork)
{
void* b_orig = original->getBuffer();
if(lWork < lOriginal || lWork < 0)
{
lWork = (int) pow(2, ceil(log2((float) (2 * (lOriginal - 1) / 0.01))));
}
std::shared_ptr<FFTPlan> fp = std::shared_ptr<FFTPlan>(new FFTPlan(lWork, 0.1));
fftwf_complex* buffer = (fftwf_complex*) fp->getBuffer();
sample_t* b_work = (sample_t*) buffer;
// Padding with 0
std::memset(b_work, 0, lWork * sizeof(sample_t));
std::memcpy(b_work, b_orig, lOriginal * sizeof(sample_t));
fp->FFT(b_work);
for(int i = 0; i < lWork / 2; i++)
{
buffer[i][0] = fabs(sqrt(buffer[i][0] * buffer[i][0] + buffer[i][1] * buffer[i][1]));
buffer[i][1] = 0.0;
int conjugate = lWork - i - 1;
buffer[conjugate][0] = buffer[i][0];
buffer[conjugate][1] = 0.0;
}
double threshold = pow(10.0, -7);
float logThreshold = (float) log(threshold);
// take 0.25*log(|H|**2) = 0.5*log(|H|)
for(int i = 0; i < lWork; i++)
{
if(buffer[i][0] < threshold)
{
buffer[i][0] = 0.5 * logThreshold;
}
else
{
buffer[i][0] = 0.5 * log(buffer[i][0]);
}
}
fp->IFFT(buffer);
// homomorphic filter
int stop = (lOriginal + 1) / 2;
b_work[0] = b_work[0] / (float) lWork;
for(int i = 1; i < stop; i++)
{
b_work[i] = b_work[i] / (float) lWork * 2.0;
}
for(int i = stop; i < lWork; i++)
{
b_work[i] = 0;
}
fp->FFT(buffer);
// EXP
// e^x = e^ (a+bi)= e^a * e^bi = e^a * (cos b + i sin b)
for(int i = 0; i < lWork / 2; i++)
{
float new_real;
float new_imag;
new_real = exp(buffer[i][0]) * cos(buffer[i][1]);
new_imag = exp(buffer[i][0]) * sin(buffer[i][1]);
buffer[i][0] = new_real;
buffer[i][1] = new_imag;
int conjugate = lWork - i - 1;
buffer[conjugate][0] = new_real;
buffer[conjugate][1] = new_imag;
}
// IFFT
fp->IFFT(buffer);
// Create new clean Buffer with only the result and normalization
int lOut = (lOriginal / 2) + lOriginal % 2;
std::shared_ptr<Buffer> bOut = std::shared_ptr<Buffer>(new Buffer(sizeof(float) * lOut));
float* bbOut = (float*) bOut->getBuffer();
// Copy and normalize
for(int i = 0; i < lOut; i++)
{
bbOut[i] = b_work[i] / (float) lWork;
}
fp->freeBuffer(buffer);
return bOut;
}
std::shared_ptr<Buffer> Equalizer::minimumPhaseFilterHilbert(std::shared_ptr<Buffer> original, int lOriginal, int lWork)
{
void* b_orig = original->getBuffer();
if(lWork < lOriginal || lWork < 0)
{
lWork = (int) pow(2, ceil(log2((float) (2 * (lOriginal - 1) / 0.01))));
}
std::shared_ptr<FFTPlan> fp = std::shared_ptr<FFTPlan>(new FFTPlan(lWork, 0.1));
fftwf_complex* buffer = (fftwf_complex*) fp->getBuffer();
sample_t* b_work = (sample_t*) buffer;
// Padding with 0
std::memset(b_work, 0, lWork * sizeof(sample_t));
std::memcpy(b_work, b_orig, lOriginal * sizeof(sample_t));
fp->FFT(b_work);
float mymax, mymin;
float n_half = (float) (lOriginal >> 1);
for(int i = 0; i < lWork; i++)
{
float w = ((float) i) * 2.0 * M_PI / (float) lWork * n_half;
float f1 = cos(w);
float f2 = sin(w);
float f3, f4;
complex_prod(buffer[i][0], buffer[i][1], f1, f2, &f3, &f4);
buffer[i][0] = f3;
buffer[i][1] = 0.0;
if(i == 0)
{
mymax = f3;
mymin = f3;
}
else
{
if(f3 < mymin)
mymin = f3;
if(f3 > mymax)
mymax = f3;
}
}
float dp = mymax - 1;
float ds = 0 - mymin;
float S = 4.0 / pow(2, (sqrt(1 + dp + ds) + sqrt(1 - dp + ds)));
for(int i = 0; i < lWork; i++)
{
buffer[i][0] = sqrt((buffer[i][0] + ds) * S) + 1.0E-10;
}
fftwf_complex* buffer_tmp = (fftwf_complex*) std::malloc(lWork * sizeof(fftwf_complex));
std::memcpy(buffer_tmp, buffer, lWork * sizeof(fftwf_complex));
//
// Hilbert transform
//
int midpt = lWork >> 1;
for(int i = 0; i < lWork; i++)
buffer[i][0] = log(buffer[i][0]);
fp->IFFT(buffer);
b_work[0] = 0.0;
for(int i = 1; i < midpt; i++)
{
b_work[i] /= (float) lWork;
}
b_work[midpt] = 0.0;
for(int i = midpt + 1; i < lWork; i++)
{
b_work[i] /= (-1.0 * lWork);
}
fp->FFT(b_work);
// Exp
for(int i = 0; i < lWork; i++)
{
float base = exp(buffer[i][0]);
buffer[i][0] = base * cos(buffer[i][1]);
buffer[i][1] = base * sin(buffer[i][1]);
complex_prod(buffer_tmp[i][0], buffer_tmp[i][1], buffer[i][0], buffer[i][1], &(buffer[i][0]), &(buffer[i][1]));
}
std::free(buffer_tmp);
fp->IFFT(buffer);
//
// Copy and normalization
//
int n_out = n_half + lOriginal % 2;
std::shared_ptr<Buffer> b_minimum = std::shared_ptr<Buffer>(new Buffer(n_out * sizeof(sample_t)));
std::memcpy(b_minimum->getBuffer(), buffer, n_out * sizeof(sample_t));
sample_t* b_final = (sample_t*) b_minimum->getBuffer();
for(int i = 0; i < n_out; i++)
{
b_final[i] /= (float) lWork;
}
return b_minimum;
}
AUD_NAMESPACE_END

2
extern/audaspace/src/fx/FFTConvolver.cpp vendored
View File

@@ -22,7 +22,7 @@
AUD_NAMESPACE_BEGIN
FFTConvolver::FFTConvolver(std::shared_ptr<std::vector<std::complex<sample_t>>> ir, std::shared_ptr<FFTPlan> plan) :
m_plan(plan), m_N(plan->getSize()), m_M(plan->getSize()/2), m_L(plan->getSize()/2), m_tailPos(0), m_irBuffer(ir)
m_plan(plan), m_N(plan->getSize()), m_M(plan->getSize()/2), m_L(plan->getSize()/2), m_irBuffer(ir), m_tailPos(0)
{
m_tail = (float*)calloc(m_M - 1, sizeof(float));
m_realBufLen = ((m_N / 2) + 1) * 2;

4
extern/audaspace/src/fx/HRTFLoaderUnix.cpp vendored
View File

@@ -75,7 +75,7 @@ void HRTFLoader::loadHRTFs(std::shared_ptr<HRTF> hrtfs, char ear, const std::str
if(ear == 'L')
azim = 360 - azim;
}
catch(std::exception& e)
catch(...)
{
AUD_THROW(FileException, "The HRTF name doesn't follow the naming scheme: " + filename);
}
@@ -86,4 +86,4 @@ void HRTFLoader::loadHRTFs(std::shared_ptr<HRTF> hrtfs, char ear, const std::str
return;
}
AUD_NAMESPACE_END
AUD_NAMESPACE_END

4
extern/audaspace/src/fx/HRTFLoaderWindows.cpp vendored
View File

@@ -78,7 +78,7 @@ void HRTFLoader::loadHRTFs(std::shared_ptr<HRTF> hrtfs, char ear, const std::str
if(ear == 'L')
azim = 360 - azim;
}
catch(std::exception& e)
catch(...)
{
AUD_THROW(FileException, "The HRTF name doesn't follow the naming scheme: " + filename);
}
@@ -90,4 +90,4 @@ void HRTFLoader::loadHRTFs(std::shared_ptr<HRTF> hrtfs, char ear, const std::str
return;
}
AUD_NAMESPACE_END
AUD_NAMESPACE_END

24
extern/mantaflow/CMakeLists.txt vendored
View File

@@ -13,10 +13,12 @@ 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(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")
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()
@@ -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,
# meaning extern_mantaflow itself actually does not have a dependency on the
# openvdb libraries, and CMAKE is free to link the vdb libs before
# extern_mantaflow causing linker errors on linux. By explicitly declaring
# a dependency here, cmake will do the right thing.
target_link_libraries(extern_mantaflow PRIVATE ${OPENVDB_LIBRARIES})
if(WITH_OPENVDB)
# The VDB libs above are only added to as INTERFACE libs by blender_add_lib,
# meaning extern_mantaflow itself actually does not have a dependency on the
# openvdb libraries, and CMAKE is free to link the vdb libs before
# extern_mantaflow causing linker errors on linux. By explicitly declaring
# a dependency here, cmake will do the right thing.
target_link_libraries(extern_mantaflow PRIVATE ${OPENVDB_LIBRARIES})
endif()

1
extern/vulkan_memory_allocator/CMakeLists.txt vendored
View File

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

15
extern/vulkan_memory_allocator/patches/remove_compilation_warning.diff vendored Normal file
View File

@@ -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;

39116
extern/vulkan_memory_allocator/vk_mem_alloc.h vendored
View File

File diff suppressed because it is too large Load Diff

25
intern/cycles/CMakeLists.txt
View File

@@ -85,15 +85,11 @@ elseif(WIN32 AND MSVC AND NOT CMAKE_CXX_COMPILER_ID MATCHES "Clang")
# there is no /arch:SSE3, but intrinsics are available anyway
if(CMAKE_CL_64)
set(CYCLES_SSE2_KERNEL_FLAGS "${CYCLES_KERNEL_FLAGS}")
set(CYCLES_SSE3_KERNEL_FLAGS "${CYCLES_KERNEL_FLAGS}")
set(CYCLES_SSE41_KERNEL_FLAGS "${CYCLES_KERNEL_FLAGS}")
set(CYCLES_AVX_KERNEL_FLAGS "${CYCLES_AVX_ARCH_FLAGS} ${CYCLES_KERNEL_FLAGS}")
set(CYCLES_AVX2_KERNEL_FLAGS "${CYCLES_AVX2_ARCH_FLAGS} ${CYCLES_KERNEL_FLAGS}")
else()
set(CYCLES_SSE2_KERNEL_FLAGS "/arch:SSE2 ${CYCLES_KERNEL_FLAGS}")
set(CYCLES_SSE3_KERNEL_FLAGS "/arch:SSE2 ${CYCLES_KERNEL_FLAGS}")
set(CYCLES_SSE41_KERNEL_FLAGS "/arch:SSE2 ${CYCLES_KERNEL_FLAGS}")
set(CYCLES_AVX_KERNEL_FLAGS "${CYCLES_AVX_ARCH_FLAGS} ${CYCLES_KERNEL_FLAGS}")
set(CYCLES_AVX2_KERNEL_FLAGS "${CYCLES_AVX2_ARCH_FLAGS} ${CYCLES_KERNEL_FLAGS}")
endif()
@@ -126,11 +122,7 @@ elseif(CMAKE_COMPILER_IS_GNUCC OR (CMAKE_CXX_COMPILER_ID MATCHES "Clang"))
endif()
set(CYCLES_SSE2_KERNEL_FLAGS "${CYCLES_KERNEL_FLAGS} -msse -msse2")
set(CYCLES_SSE3_KERNEL_FLAGS "${CYCLES_SSE2_KERNEL_FLAGS} -msse3 -mssse3")
set(CYCLES_SSE41_KERNEL_FLAGS "${CYCLES_SSE3_KERNEL_FLAGS} -msse4.1")
if(CXX_HAS_AVX)
set(CYCLES_AVX_KERNEL_FLAGS "${CYCLES_SSE41_KERNEL_FLAGS} -mavx")
endif()
set(CYCLES_SSE41_KERNEL_FLAGS "${CYCLES_SSE2_KERNEL_FLAGS} -msse3 -mssse3 -msse4.1")
if(CXX_HAS_AVX2)
set(CYCLES_AVX2_KERNEL_FLAGS "${CYCLES_SSE41_KERNEL_FLAGS} -mavx -mavx2 -mfma -mlzcnt -mbmi -mbmi2 -mf16c")
endif()
@@ -144,13 +136,8 @@ elseif(WIN32 AND CMAKE_CXX_COMPILER_ID MATCHES "Intel")
if(CXX_HAS_SSE)
set(CYCLES_SSE2_KERNEL_FLAGS "/QxSSE2")
set(CYCLES_SSE3_KERNEL_FLAGS "/QxSSSE3")
set(CYCLES_SSE41_KERNEL_FLAGS "/QxSSE4.1")
if(CXX_HAS_AVX)
set(CYCLES_AVX_KERNEL_FLAGS "/arch:AVX")
endif()
if(CXX_HAS_AVX2)
set(CYCLES_AVX2_KERNEL_FLAGS "/QxCORE-AVX2")
endif()
@@ -174,13 +161,8 @@ elseif(CMAKE_CXX_COMPILER_ID MATCHES "Intel")
set(CYCLES_SSE2_KERNEL_FLAGS "-xsse2")
endif()
set(CYCLES_SSE3_KERNEL_FLAGS "-xssse3")
set(CYCLES_SSE41_KERNEL_FLAGS "-xsse4.1")
if(CXX_HAS_AVX)
set(CYCLES_AVX_KERNEL_FLAGS "-xavx")
endif()
if(CXX_HAS_AVX2)
set(CYCLES_AVX2_KERNEL_FLAGS "-xcore-avx2")
endif()
@@ -190,15 +172,10 @@ endif()
if(CXX_HAS_SSE)
add_definitions(
-DWITH_KERNEL_SSE2
-DWITH_KERNEL_SSE3
-DWITH_KERNEL_SSE41
)
endif()
if(CXX_HAS_AVX)
add_definitions(-DWITH_KERNEL_AVX)
endif()
if(CXX_HAS_AVX2)
add_definitions(-DWITH_KERNEL_AVX2)
endif()

10
intern/cycles/blender/addon/properties.py
View File

@@ -951,9 +951,7 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
return _cycles.debug_flags_update(scene)
debug_use_cpu_avx2: BoolProperty(name="AVX2", default=True)
debug_use_cpu_avx: BoolProperty(name="AVX", default=True)
debug_use_cpu_sse41: BoolProperty(name="SSE41", default=True)
debug_use_cpu_sse3: BoolProperty(name="SSE3", default=True)
debug_use_cpu_sse2: BoolProperty(name="SSE2", default=True)
debug_bvh_layout: EnumProperty(
name="BVH Layout",
@@ -1673,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')

2
intern/cycles/blender/addon/ui.py
View File

@@ -2112,9 +2112,7 @@ class CYCLES_RENDER_PT_debug(CyclesDebugButtonsPanel, Panel):
row = col.row(align=True)
row.prop(cscene, "debug_use_cpu_sse2", toggle=True)
row.prop(cscene, "debug_use_cpu_sse3", toggle=True)
row.prop(cscene, "debug_use_cpu_sse41", toggle=True)
row.prop(cscene, "debug_use_cpu_avx", toggle=True)
row.prop(cscene, "debug_use_cpu_avx2", toggle=True)
col.prop(cscene, "debug_bvh_layout", text="BVH")

2
intern/cycles/blender/light.cpp
View File

@@ -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());

2
intern/cycles/blender/python.cpp
View File

@@ -63,9 +63,7 @@ static void debug_flags_sync_from_scene(BL::Scene b_scene)
PointerRNA cscene = RNA_pointer_get(&b_scene.ptr, "cycles");
/* Synchronize CPU flags. */
flags.cpu.avx2 = get_boolean(cscene, "debug_use_cpu_avx2");
flags.cpu.avx = get_boolean(cscene, "debug_use_cpu_avx");
flags.cpu.sse41 = get_boolean(cscene, "debug_use_cpu_sse41");
flags.cpu.sse3 = get_boolean(cscene, "debug_use_cpu_sse3");
flags.cpu.sse2 = get_boolean(cscene, "debug_use_cpu_sse2");
flags.cpu.bvh_layout = (BVHLayout)get_enum(cscene, "debug_bvh_layout");
/* Synchronize CUDA flags. */

16
intern/cycles/cmake/macros.cmake
View File

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

2
intern/cycles/device/cpu/device.cpp
View File

@@ -45,9 +45,7 @@ string device_cpu_capabilities()
{
string capabilities = "";
capabilities += system_cpu_support_sse2() ? "SSE2 " : "";
capabilities += system_cpu_support_sse3() ? "SSE3 " : "";
capabilities += system_cpu_support_sse41() ? "SSE41 " : "";
capabilities += system_cpu_support_avx() ? "AVX " : "";
capabilities += system_cpu_support_avx2() ? "AVX2" : "";
if (capabilities[capabilities.size() - 1] == ' ')
capabilities.resize(capabilities.size() - 1);

3
intern/cycles/device/cpu/kernel.cpp
View File

@@ -9,8 +9,7 @@ CCL_NAMESPACE_BEGIN
#define KERNEL_FUNCTIONS(name) \
KERNEL_NAME_EVAL(cpu, name), KERNEL_NAME_EVAL(cpu_sse2, name), \
KERNEL_NAME_EVAL(cpu_sse3, name), KERNEL_NAME_EVAL(cpu_sse41, name), \
KERNEL_NAME_EVAL(cpu_avx, name), KERNEL_NAME_EVAL(cpu_avx2, name)
KERNEL_NAME_EVAL(cpu_sse41, name), KERNEL_NAME_EVAL(cpu_avx2, name)
#define REGISTER_KERNEL(name) name(KERNEL_FUNCTIONS(name))
#define REGISTER_KERNEL_FILM_CONVERT(name) \

21
intern/cycles/device/cpu/kernel_function.h
View File

@@ -17,13 +17,10 @@ template<typename FunctionType> class CPUKernelFunction {
public:
CPUKernelFunction(FunctionType kernel_default,
FunctionType kernel_sse2,
FunctionType kernel_sse3,
FunctionType kernel_sse41,
FunctionType kernel_avx,
FunctionType kernel_avx2)
{
kernel_info_ = get_best_kernel_info(
kernel_default, kernel_sse2, kernel_sse3, kernel_sse41, kernel_avx, kernel_avx2);
kernel_info_ = get_best_kernel_info(kernel_default, kernel_sse2, kernel_sse41, kernel_avx2);
}
template<typename... Args> inline auto operator()(Args... args) const
@@ -60,16 +57,12 @@ template<typename FunctionType> class CPUKernelFunction {
KernelInfo get_best_kernel_info(FunctionType kernel_default,
FunctionType kernel_sse2,
FunctionType kernel_sse3,
FunctionType kernel_sse41,
FunctionType kernel_avx,
FunctionType kernel_avx2)
{
/* Silence warnings about unused variables when compiling without some architectures. */
(void)kernel_sse2;
(void)kernel_sse3;
(void)kernel_sse41;
(void)kernel_avx;
(void)kernel_avx2;
#ifdef WITH_CYCLES_OPTIMIZED_KERNEL_AVX2
@@ -78,24 +71,12 @@ template<typename FunctionType> class CPUKernelFunction {
}
#endif
#ifdef WITH_CYCLES_OPTIMIZED_KERNEL_AVX
if (DebugFlags().cpu.has_avx() && system_cpu_support_avx()) {
return KernelInfo("AVX", kernel_avx);
}
#endif
#ifdef WITH_CYCLES_OPTIMIZED_KERNEL_SSE41
if (DebugFlags().cpu.has_sse41() && system_cpu_support_sse41()) {
return KernelInfo("SSE4.1", kernel_sse41);
}
#endif
#ifdef WITH_CYCLES_OPTIMIZED_KERNEL_SSE3
if (DebugFlags().cpu.has_sse3() && system_cpu_support_sse3()) {
return KernelInfo("SSE3", kernel_sse3);
}
#endif
#ifdef WITH_CYCLES_OPTIMIZED_KERNEL_SSE2
if (DebugFlags().cpu.has_sse2() && system_cpu_support_sse2()) {
return KernelInfo("SSE2", kernel_sse2);

2
intern/cycles/device/hip/util.h
View File

@@ -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

4
intern/cycles/device/metal/device.mm
View File

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

15
intern/cycles/device/metal/device_impl.mm
View File

@@ -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) \
baked_constants += string(#parent "." #name "=") + \
to_string(_type(launch_params.data.parent.name)) + "\n";
if (next_member_is_specialized) { \
baked_constants += string(#parent "." #name "=") + \
to_string(_type(launch_params.data.parent.name)) + "\n"; \
} \
else { \
string_replace( \
source, "kernel_data_" #parent "_" #name, "kernel_data." #parent ".__unused_" #name); \
next_member_is_specialized = true; \
}
# include "kernel/data_template.h"

25
intern/cycles/device/metal/kernel.mm
View File

@@ -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_int type:MTLDataType_int atIndex:Kernel_DummyConstant];
[constant_values setConstantValue:&zero_data type:MTLDataType_int atIndex:Kernel_DummyConstant];
bool next_member_is_specialized = true;
# define KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE next_member_is_specialized = false;
# 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"

3
intern/cycles/device/metal/queue.mm
View File

@@ -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];

1
intern/cycles/device/metal/util.h
View File

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

2
intern/cycles/device/metal/util.mm
View File

@@ -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;
}

8
intern/cycles/device/oneapi/device_impl.cpp
View File

@@ -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. */
static const int lowest_supported_driver_version_win = 1013430;
static const int lowest_supported_driver_version_neo = 23904;
* For Windows driver 101.4032, compute-runtime version is 24931. */
static const int lowest_supported_driver_version_win = 1014032;
static const int lowest_supported_driver_version_neo = 24931;
int OneapiDevice::parse_driver_build_version(const sycl::device &device)
{

3
intern/cycles/graph/CMakeLists.txt
View File

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

3
intern/cycles/integrator/CMakeLists.txt
View File

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

43
intern/cycles/kernel/CMakeLists.txt
View File

@@ -14,9 +14,7 @@ set(INC_SYS
set(SRC_KERNEL_DEVICE_CPU
device/cpu/kernel.cpp
device/cpu/kernel_sse2.cpp
device/cpu/kernel_sse3.cpp
device/cpu/kernel_sse41.cpp
device/cpu/kernel_avx.cpp
device/cpu/kernel_avx2.cpp
)
@@ -736,25 +734,25 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
endif()
# SYCL_CPP_FLAGS is a variable that the user can set to pass extra compiler options
set(sycl_compiler_flags
${CMAKE_CURRENT_SOURCE_DIR}/${SRC_KERNEL_DEVICE_ONEAPI}
-fsycl
-fsycl-unnamed-lambda
-fdelayed-template-parsing
-mllvm -inlinedefault-threshold=250
-mllvm -inlinehint-threshold=350
-fsycl-device-code-split=per_kernel
-fsycl-max-parallel-link-jobs=${SYCL_OFFLINE_COMPILER_PARALLEL_JOBS}
-shared
-DWITH_ONEAPI
-ffast-math
-DNDEBUG
-O2
-o ${cycles_kernel_oneapi_lib}
-I${CMAKE_CURRENT_SOURCE_DIR}/..
${SYCL_CPP_FLAGS}
)
${CMAKE_CURRENT_SOURCE_DIR}/${SRC_KERNEL_DEVICE_ONEAPI}
-fsycl
-fsycl-unnamed-lambda
-fdelayed-template-parsing
-mllvm -inlinedefault-threshold=250
-mllvm -inlinehint-threshold=350
-fsycl-device-code-split=per_kernel
-fsycl-max-parallel-link-jobs=${SYCL_OFFLINE_COMPILER_PARALLEL_JOBS}
-shared
-DWITH_ONEAPI
-ffast-math
-DNDEBUG
-O2
-o ${cycles_kernel_oneapi_lib}
-I${CMAKE_CURRENT_SOURCE_DIR}/..
${SYCL_CPP_FLAGS}
)
if (WITH_CYCLES_ONEAPI_HOST_TASK_EXECUTION)
if(WITH_CYCLES_ONEAPI_HOST_TASK_EXECUTION)
list(APPEND sycl_compiler_flags -DWITH_ONEAPI_SYCL_HOST_TASK)
endif()
@@ -942,14 +940,9 @@ set_source_files_properties(device/cpu/kernel.cpp PROPERTIES COMPILE_FLAGS "${CY
if(CXX_HAS_SSE)
set_source_files_properties(device/cpu/kernel_sse2.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_SSE2_KERNEL_FLAGS}")
set_source_files_properties(device/cpu/kernel_sse3.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_SSE3_KERNEL_FLAGS}")
set_source_files_properties(device/cpu/kernel_sse41.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_SSE41_KERNEL_FLAGS}")
endif()
if(CXX_HAS_AVX)
set_source_files_properties(device/cpu/kernel_avx.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_AVX_KERNEL_FLAGS}")
endif()
if(CXX_HAS_AVX2)
set_source_files_properties(device/cpu/kernel_avx2.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_AVX2_KERNEL_FLAGS}")
endif()

5
intern/cycles/kernel/bake/bake.h
View File

@@ -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. */
const uint32_t path_flag = PATH_RAY_EMISSION;
* This is being evaluated for all BSDFs, so path flag does not contain a specific type.
* However, we want to flag the ray visibility to ignore the sun in the background map. */
const uint32_t path_flag = PATH_RAY_EMISSION | PATH_RAY_IMPORTANCE_BAKE;
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);

173
intern/cycles/kernel/closure/bsdf.h
View File

@@ -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,
const float3 omega_in)
ccl_device_inline bool bsdf_is_transmission(ccl_private const ShaderClosure *sc, const float3 wo)
{
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,
sc,
Ng,
sd->I,
randu,
randv,
eval,
omega_in,
pdf,
&sd->lcg_state,
sampled_roughness,
eta);
label = bsdf_microfacet_multi_ggx_sample(
kg, sc, Ng, sd->wi, randu, randv, eval, wo, pdf, &sd->lcg_state, sampled_roughness, eta);
break;
case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_ID:
case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_FRESNEL_ID:
label = bsdf_microfacet_multi_ggx_glass_sample(kg,
sc,
Ng,
sd->I,
randu,
randv,
eval,
omega_in,
pdf,
&sd->lcg_state,
sampled_roughness,
eta);
label = bsdf_microfacet_multi_ggx_glass_sample(
kg, sc, Ng, sd->wi, randu, randv, eval, wo, 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);
*eval *= shift_cos_in(cosNI, frequency_multiplier);
const float cosNO = dot(*wo, sc->N);
*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,8 +461,8 @@ ccl_device_inline int bsdf_label(const KernelGlobals kg,
}
case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_ID:
case CLOSURE_BSDF_MICROFACET_MULTI_GGX_GLASS_FRESNEL_ID:
label = (bsdf_is_transmission(sc, omega_in)) ? LABEL_TRANSMIT | LABEL_GLOSSY :
LABEL_REFLECT | LABEL_GLOSSY;
label = (bsdf_is_transmission(sc, wo)) ? LABEL_TRANSMIT | LABEL_GLOSSY :
LABEL_REFLECT | LABEL_GLOSSY;
break;
case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
label = LABEL_REFLECT | LABEL_GLOSSY;
@@ -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);
if (cosNI >= 0.0f) {
eval *= shift_cos_in(cosNI, frequency_multiplier);
const float cosNO = dot(wo, sc->N);
if (cosNO >= 0.0f) {
eval *= shift_cos_in(cosNO, frequency_multiplier);
}
}
@@ -682,4 +661,38 @@ 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

34
intern/cycles/kernel/closure/bsdf_ashikhmin_shirley.h
View File

@@ -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 omega_in,
const float3 wi,
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 NdotO = dot(N, omega_in); /* and consequently we use for O omaga_in ;) */
float NdotI = dot(N, wi);
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);
float HdotI = fmaxf(fabsf(dot(H, I)), 1e-6f);
float3 H = normalize(wi + wo);
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 */
*omega_in = -I + (2.0f * HdotI) * H;
/* reflect wi on H to get wo */
*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;

54
intern/cycles/kernel/closure/bsdf_ashikhmin_velvet.h
View File

@@ -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 omega_in,
const float3 wi,
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, omega_in);
if (!(cosNO > 0 && cosNI > 0)) {
float cosNI = dot(N, wi);
float cosNO = dot(N, wo);
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;
cosNHdivHO = fmaxf(cosNHdivHO, 1e-5f);
float cosNHdivHI = cosNH / cosHI;
cosNHdivHI = fmaxf(cosNHdivHI, 1e-5f);
float fac1 = 2 * fabsf(cosNHdivHO * cosNO);
float fac2 = 2 * fabsf(cosNHdivHO * cosNI);
float fac1 = 2 * fabsf(cosNHdivHI * 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 cosNO = dot(N, I);
float cosNI = dot(N, wi);
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;
cosNHdivHO = fmaxf(cosNHdivHO, 1e-5f);
float cosNHdivHI = cosNH / cosHI;
cosNHdivHI = fmaxf(cosNHdivHI, 1e-5f);
float fac1 = 2 * fabsf(cosNHdivHO * cosNO);
float fac2 = 2 * fabsf(cosNHdivHO * cosNI);
float fac1 = 2 * fabsf(cosNHdivHI * 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);

36
intern/cycles/kernel/closure/bsdf_diffuse.h
View File

@@ -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 omega_in,
const float3 wi,
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;
*pdf = cos_pi;
return make_spectrum(cos_pi);
float cosNO = fmaxf(dot(N, wo), 0.0f) * M_1_PI_F;
*pdf = cosNO;
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 omega_in,
const float3 wi,
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;
*pdf = cos_pi;
return make_spectrum(cos_pi);
float cosNO = fmaxf(-dot(N, wo), 0.0f) * M_1_PI_F;
*pdf = cosNO;
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);
if (dot(Ng, *omega_in) < 0) {
sample_cos_hemisphere(-N, randu, randv, wo, pdf);
if (dot(Ng, *wo) < 0) {
*eval = make_spectrum(*pdf);
}
else {

20
intern/cycles/kernel/closure/bsdf_diffuse_ramp.h
View File

@@ -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 omega_in,
const float3 wi,
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);
if (cos_pi >= 0.0f) {
*pdf = cos_pi * M_1_PI_F;
return rgb_to_spectrum(bsdf_diffuse_ramp_get_color(bsdf->colors, cos_pi) * M_1_PI_F);
float cosNO = fmaxf(dot(N, wo), 0.0f);
if (cosNO >= 0.0f) {
*pdf = cosNO * M_1_PI_F;
return rgb_to_spectrum(bsdf_diffuse_ramp_get_color(bsdf->colors, cosNO) * M_1_PI_F);
}
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 {

58
intern/cycles/kernel/closure/bsdf_hair.h
View File

@@ -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 omega_in,
const float3 wi,
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);
float3 locy = normalize(I - Tg * Iz);
float Iz = dot(Tg, wi);
float3 locy = normalize(wi - Tg * Iz);
float theta_r = M_PI_2_F - fast_acosf(Iz);
float omega_in_z = dot(Tg, omega_in);
float3 omega_in_y = normalize(omega_in - Tg * omega_in_z);
float wo_z = dot(Tg, wo);
float3 wo_y = normalize(wo - Tg * wo_z);
float theta_i = M_PI_2_F - fast_acosf(omega_in_z);
float cosphi_i = dot(omega_in_y, locy);
float theta_i = M_PI_2_F - fast_acosf(wo_z);
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 omega_in,
const float3 wi,
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);
float3 locy = normalize(I - Tg * Iz);
float Iz = dot(Tg, wi);
float3 locy = normalize(wi - Tg * Iz);
float theta_r = M_PI_2_F - fast_acosf(Iz);
float omega_in_z = dot(Tg, omega_in);
float3 omega_in_y = normalize(omega_in - Tg * omega_in_z);
float wo_z = dot(Tg, wo);
float3 wo_y = normalize(wo - Tg * wo_z);
float theta_i = M_PI_2_F - fast_acosf(omega_in_z);
float phi_i = fast_acosf(dot(omega_in_y, locy));
float theta_i = M_PI_2_F - fast_acosf(wo_z);
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);
float3 locy = normalize(I - Tg * Iz);
float Iz = dot(Tg, wi);
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);
float3 locy = normalize(I - Tg * Iz);
float Iz = dot(Tg, wi);
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;
}

32
intern/cycles/kernel/closure/bsdf_hair_principled.h
View File

@@ -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));
const float3 wi = make_float3(dot(omega_in, X), dot(omega_in, Y), dot(omega_in, Z));
/* local_I is the illumination direction. */
const float3 local_O = make_float3(dot(sd->wi, X), dot(sd->wi, Y), dot(sd->wi, Z));
const float3 local_I = make_float3(dot(wo, X), dot(wo, Y), dot(wo, Z));
const 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);
}

1364
intern/cycles/kernel/closure/bsdf_microfacet.h
View File

File diff suppressed because it is too large Load Diff

94
intern/cycles/kernel/closure/bsdf_microfacet_multi.h
View File

@@ -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 omega_in,
const float3 wi,
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 localO = make_float3(dot(omega_in, X), dot(omega_in, Y), dot(omega_in, Z));
float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
float3 local_O = make_float3(dot(wo, X), dot(wo, Y), dot(wo, Z));
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,
localO,
return mf_eval_glossy(local_I,
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;
if (dot(Ng, *omega_in) <= 0.0f) {
*wo = 2 * dot(Z, wi) * Z - wi;
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 localO;
float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
float3 local_O;
*eval = mf_sample_glossy(localI,
&localO,
*eval = mf_sample_glossy(local_I,
&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 omega_in,
const float3 wi,
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 localO = make_float3(dot(omega_in, X), dot(omega_in, Y), dot(omega_in, Z));
float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
float3 local_O = make_float3(dot(wo, X), dot(wo, Y), dot(wo, Z));
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,
localO,
return mf_eval_glass(local_I,
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 localO;
float3 local_I = make_float3(dot(wi, X), dot(wi, Y), dot(wi, Z));
float3 local_O;
*eval = mf_sample_glass(localI,
&localO,
*eval = mf_sample_glass(local_I,
&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;
if (localO.z * localI.z > 0.0f) {
*wo = X * local_O.x + Y * local_O.y + Z * local_O.z;
if (local_O.z * local_I.z > 0.0f) {
return LABEL_REFLECT | LABEL_GLOSSY;
}
else {

23
intern/cycles/kernel/closure/bsdf_microfacet_multi_impl.h
View File

@@ -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);
if (order == 0)
throughput = t_color;
else
throughput *= t_color;
throughput *= interpolate_fresnel_color(wi_prev, wm, eta, cspec0);
}
}
#else /* MF_MULTI_GLOSSY */
if (use_fresnel) {
Spectrum t_color = interpolate_fresnel_color(-wr, wm, eta, F0, cspec0);
if (order == 0)
throughput = t_color;
else
throughput *= t_color;
throughput *= interpolate_fresnel_color(-wr, wm, eta, cspec0);
}
wr = mf_sample_phase_glossy(-wr, &throughput, wm);
#endif

18
intern/cycles/kernel/closure/bsdf_oren_nayar.h
View File

@@ -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 omega_in,
const float3 wi,
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) {
*eval = bsdf_oren_nayar_get_intensity(sc, bsdf->N, I, *omega_in);
if (dot(Ng, *wo) > 0.0f) {
*eval = bsdf_oren_nayar_get_intensity(sc, bsdf->N, wi, *wo);
}
else {
*pdf = 0.0f;

38
intern/cycles/kernel/closure/bsdf_phong_ramp.h
View File

@@ -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 omega_in,
const float3 wi,
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 cosNO = dot(bsdf->N, I);
float cosNI = dot(bsdf->N, wi);
float cosNO = dot(bsdf->N, wo);
if (cosNI > 0 && cosNO > 0) {
// reflect the view vector
float3 R = (2 * cosNO) * bsdf->N - I;
float cosRI = dot(R, omega_in);
if (cosRI > 0) {
float cosp = powf(cosRI, m_exponent);
float3 R = (2 * cosNI) * bsdf->N - wi;
float cosRO = dot(R, wo);
if (cosRO > 0) {
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;
if (dot(Ng, *omega_in) > 0.0f) {
*wo = (cosf(phi) * sinTheta) * T + (sinf(phi) * sinTheta) * B + (cosTheta)*R;
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);
}
}

22
intern/cycles/kernel/closure/bsdf_principled_diffuse.h
View File

@@ -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 omega_in,
const float3 wi,
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) {
const float3 V = I; // outgoing
const float3 L = omega_in; // incoming
*pdf = fmaxf(dot(N, omega_in), 0.0f) * M_1_PI_F;
if (dot(N, wo) > 0.0f) {
const float3 V = wi;
const float3 L = wo;
*pdf = fmaxf(dot(N, wo), 0.0f) * M_1_PI_F;
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) {
*eval = bsdf_principled_diffuse_compute_brdf(bsdf, N, I, *omega_in, pdf);
if (dot(Ng, *wo) > 0) {
*eval = bsdf_principled_diffuse_compute_brdf(bsdf, N, wi, *wo, pdf);
}
else {
*pdf = 0.0f;

26
intern/cycles/kernel/closure/bsdf_principled_sheen.h
View File

@@ -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 omega_in,
const float3 wi,
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) {
const float3 V = I; // outgoing
const float3 L = omega_in; // incoming
if (dot(N, wo) > 0.0f) {
const float3 V = wi;
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) {
float3 H = normalize(I + *omega_in);
if (dot(Ng, *wo) > 0) {
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();

16
intern/cycles/kernel/closure/bsdf_reflection.h
View File

@@ -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 omega_in,
const float3 wi,
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);
if (cosNO > 0) {
*omega_in = (2 * cosNO) * N - I;
if (dot(Ng, *omega_in) > 0) {
float cosNI = dot(N, wi);
if (cosNI > 0) {
*wo = (2 * cosNI) * N - wi;
if (dot(Ng, *wo) > 0) {
/* Some high number for MIS. */
*pdf = 1e6f;
*eval = make_spectrum(1e6f);

12
intern/cycles/kernel/closure/bsdf_refraction.h
View File

@@ -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 omega_in,
const float3 wi,
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;

50
intern/cycles/kernel/closure/bsdf_toon.h
View File

@@ -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 omega_in,
const float3 wi,
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 omega_in,
const float3 wi,
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 cosNO = dot(bsdf->N, I);
float cosNI = dot(bsdf->N, wi);
float cosNO = dot(bsdf->N, wo);
if (cosNI > 0 && cosNO > 0) {
/* reflect the view vector */
float3 R = (2 * cosNO) * bsdf->N - I;
float cosRI = dot(R, omega_in);
float3 R = (2 * cosNI) * bsdf->N - wi;
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) {
float cosNI = dot(bsdf->N, *omega_in);
if (dot(Ng, *wo) > 0.0f) {
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 {

10
intern/cycles/kernel/closure/bsdf_transparent.h
View File

@@ -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 omega_in,
const float3 wi,
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;

24
intern/cycles/kernel/closure/bsdf_util.h
View File

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

2
intern/cycles/kernel/closure/bssrdf.h
View File

@@ -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);
}
}

19
intern/cycles/kernel/closure/emissive.h
View File

@@ -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));
return (cosNO > 0.0f) ? 1.0f : 0.0f;
float cosNI = fabsf(dot(Ng, wi));
return (cosNI > 0.0f) ? 1.0f : 0.0f;
}
ccl_device void emissive_sample(const float3 Ng,
float randu,
float randv,
ccl_private float3 *omega_out,
ccl_private float *pdf)
ccl_device void emissive_sample(
const float3 Ng, float randu, float randv, ccl_private float3 *wi, 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);
}

26
intern/cycles/kernel/closure/volume.h
View File

@@ -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,
float3 omega_in,
const float3 wi,
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 */
*omega_in = henyey_greenstrein_sample(-I, g, randu, randv, pdf);
/* note that wi points towards the viewer and so is used negated */
*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. */

11
intern/cycles/kernel/data_template.h
View File

@@ -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

6
intern/cycles/kernel/device/cpu/kernel.h
View File

@@ -35,15 +35,9 @@ void kernel_global_memory_copy(KernelGlobalsCPU *kg, const char *name, void *mem
#define KERNEL_ARCH cpu_sse2
#include "kernel/device/cpu/kernel_arch.h"
#define KERNEL_ARCH cpu_sse3
#include "kernel/device/cpu/kernel_arch.h"
#define KERNEL_ARCH cpu_sse41
#include "kernel/device/cpu/kernel_arch.h"
#define KERNEL_ARCH cpu_avx
#include "kernel/device/cpu/kernel_arch.h"
#define KERNEL_ARCH cpu_avx2
#include "kernel/device/cpu/kernel_arch.h"

26
intern/cycles/kernel/device/cpu/kernel_avx.cpp
View File

@@ -1,26 +0,0 @@
/* SPDX-License-Identifier: Apache-2.0
* Copyright 2011-2022 Blender Foundation */
/* Optimized CPU kernel entry points. This file is compiled with AVX
* optimization flags and nearly all functions inlined, while kernel.cpp
* is compiled without for other CPU's. */
#include "util/optimization.h"
#ifndef WITH_CYCLES_OPTIMIZED_KERNEL_AVX
# define KERNEL_STUB
#else
/* SSE optimization disabled for now on 32 bit, see bug T36316. */
# if !(defined(__GNUC__) && (defined(i386) || defined(_M_IX86)))
# define __KERNEL_SSE__
# define __KERNEL_SSE2__
# define __KERNEL_SSE3__
# define __KERNEL_SSSE3__
# define __KERNEL_SSE41__
# define __KERNEL_AVX__
# endif
#endif /* WITH_CYCLES_OPTIMIZED_KERNEL_AVX */
#include "kernel/device/cpu/kernel.h"
#define KERNEL_ARCH cpu_avx
#include "kernel/device/cpu/kernel_arch_impl.h"

23
intern/cycles/kernel/device/cpu/kernel_sse3.cpp
View File

@@ -1,23 +0,0 @@
/* SPDX-License-Identifier: Apache-2.0
* Copyright 2011-2022 Blender Foundation */
/* Optimized CPU kernel entry points. This file is compiled with SSE3/SSSE3
* optimization flags and nearly all functions inlined, while kernel.cpp
* is compiled without for other CPU's. */
#include "util/optimization.h"
#ifndef WITH_CYCLES_OPTIMIZED_KERNEL_SSE3
# define KERNEL_STUB
#else
/* SSE optimization disabled for now on 32 bit, see bug T36316. */
# if !(defined(__GNUC__) && (defined(i386) || defined(_M_IX86)))
# define __KERNEL_SSE2__
# define __KERNEL_SSE3__
# define __KERNEL_SSSE3__
# endif
#endif /* WITH_CYCLES_OPTIMIZED_KERNEL_SSE3 */
#include "kernel/device/cpu/kernel.h"
#define KERNEL_ARCH cpu_sse3
#include "kernel/device/cpu/kernel_arch_impl.h"

4
intern/cycles/kernel/device/metal/context_begin.h
View File

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

10
intern/cycles/kernel/device/oneapi/compat.h
View File

@@ -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)))

19
intern/cycles/kernel/film/denoising_passes.h
View File

@@ -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;
/* 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) {
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;

2
intern/cycles/kernel/geom/curve.h
View File

@@ -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 */

Some files were not shown because too many files have changed in this diff Show More