Compare commits

..

47 Commits

Author SHA1 Message Date
ea2db47a0b Use null for root tree for node tree change propagation 2023-01-10 16:49:57 -05:00
eaa17cfae2 Merge branch 'master' into temp-nodes-group-declarations 2023-01-10 16:41:44 -05:00
082449b737 Add missing tag for socket declaration node sockets update 2023-01-09 23:22:59 -05:00
b1bbd33ed6 Merge branch 'master' into temp-nodes-group-declarations 2023-01-09 17:48:39 -05:00
96609ade04 Merge branch 'master' into temp-nodes-group-declarations 2023-01-09 16:35:40 -05:00
844b3ed472 Fix various issues in socket declaration updating/building/comparison 2023-01-02 16:06:54 -05:00
392f4481a7 Remove redundant node declaration allocation 2023-01-02 16:06:37 -05:00
612c0ae3c8 Correct flag used for compact sockets 2023-01-02 14:57:32 -05:00
1847b63d27 Fix crash when creating a group 2023-01-02 14:54:04 -05:00
dd32b8fc97 Merge branch 'master' into temp-nodes-group-declarations 2023-01-02 14:03:24 -05:00
06b5ceb0c4 Fixes to group creation 2022-12-29 16:31:19 -05:00
ef89ba2b34 Merge branch 'master' into temp-nodes-group-declarations 2022-12-29 15:40:18 -05:00
399fc4af2a Merge branch 'master' into temp-nodes-group-declarations 2022-12-29 14:12:26 -05:00
92a16d0193 Merge branch 'master' into temp-nodes-group-declarations 2022-12-29 13:30:23 -05:00
e76e8fe87a Various cleanups, making the diff smaller 2022-12-27 16:40:58 -05:00
0c87d9b8f7 Merge branch 'master' into temp-nodes-group-declarations 2022-12-27 12:04:36 -05:00
d5b2979ad5 Cleanup: Fix typos 2022-12-27 10:52:44 -05:00
b60fc1c4c8 Merge branch 'master' into temp-nodes-group-declarations 2022-12-27 10:50:26 -05:00
4052267be7 Copy over "collapsed" flag 2022-12-22 16:01:24 -05:00
60210762e0 Handle property subtypes 2022-12-22 15:58:25 -05:00
13d984cfaa Add "Custom" socket declaration
I didn't test this yet
2022-12-22 15:18:07 -05:00
195a271bab Add more early return checks to geometry node group declaration 2022-12-22 15:04:27 -05:00
ddee26288d Simplify declaration functions for "extend" socket 2022-12-22 14:42:33 -05:00
af5a9b5748 Improve comments in various places 2022-12-22 14:42:17 -05:00
109f18df71 Fix clang format differences 2022-12-22 14:41:44 -05:00
450693ea6b Make function static again 2022-12-22 14:41:31 -05:00
f06635bfde Improve comment for dynamic declaration callback 2022-12-22 14:41:16 -05:00
21a8d7c69b Fix shader socket declaration
Copy and paste error. All tests pass now
2022-12-22 14:39:53 -05:00
1ffe27c438 Fix updating of input and output nodes when creating group 2022-12-22 11:51:20 -05:00
86b9b61371 Merge branch 'master' into temp-nodes-group-declarations 2022-12-22 10:42:37 -05:00
a2992774ac Only rebuild declaration when it is dynamic 2022-12-21 13:00:04 -06:00
cc57344527 Merge branch 'master' into temp-nodes-group-declarations 2022-12-21 12:58:57 -06:00
e4139476bf Merge branch 'master' into temp-nodes-group-declarations 2022-12-20 14:50:24 -06:00
80a3bed01a Don't connect between two extend sockets 2022-12-19 21:10:15 -06:00
439761d15c Remove theoretically unnecessary optimization 2022-12-19 20:51:57 -06:00
7b4dcec3ec Various fixes 2022-12-19 16:59:55 -06:00
f4d52d060b Merge branch 'master' into temp-nodes-group-declarations 2022-12-19 16:02:50 -06:00
a90642816c Some progress 2022-12-16 16:42:19 -06:00
8a27da7b3b Merge branch 'master' into temp-nodes-group-declarations 2022-12-16 14:50:08 -06:00
c174184385 Merge branch 'master' into temp-nodes-group-declarations 2022-12-16 11:32:09 -06:00
660e240acf Basic startup works 2022-12-15 17:43:08 -06:00
f509bd37b1 Fix build errors 2022-12-15 17:18:34 -06:00
542bff5e94 Declarations for group input and output nodes 2022-12-15 16:44:08 -06:00
d82166833a Merge branch 'master' into temp-nodes-group-declarations 2022-12-15 16:22:26 -06:00
64297ec8ad Cleanup 2022-12-15 16:17:16 -06:00
510155ed5c Merge branch 'master' into temp-nodes-group-declarations 2022-12-15 15:06:34 -06:00
59b1eb145d Start of refactoring to use declaration for group nodes 2022-12-14 18:42:04 -06:00
1404 changed files with 40356 additions and 62457 deletions

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

@@ -1,38 +0,0 @@
<!--
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.

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

@@ -513,7 +513,6 @@ 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
@@ -525,7 +524,6 @@ if(WITH_FFTW)
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

View File

@@ -54,7 +54,6 @@
#ifdef WITH_CONVOLUTION
#include "fx/BinauralSound.h"
#include "fx/ConvolverSound.h"
#include "fx/Equalizer.h"
#endif
#include <cassert>
@@ -769,14 +768,4 @@ 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

View File

@@ -397,16 +397,6 @@ 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

View File

@@ -53,7 +53,6 @@ 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);

View File

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

View File

@@ -1,106 +0,0 @@
/*******************************************************************************
* 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

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_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)
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)
{
if(m_hrtfs->isEmpty())
AUD_THROW(StateException, "The provided HRTF object is empty");

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_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_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_resetFlag = false;

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_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_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_nChannelThreads = std::min((int)threadPool->getNumOfThreads(), m_inChannels);
m_futures.resize(m_nChannelThreads);

View File

@@ -1,367 +0,0 @@
/*******************************************************************************
* 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

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_irBuffer(ir), m_tailPos(0)
m_plan(plan), m_N(plan->getSize()), m_M(plan->getSize()/2), m_L(plan->getSize()/2), m_tailPos(0), m_irBuffer(ir)
{
m_tail = (float*)calloc(m_M - 1, sizeof(float));
m_realBufLen = ((m_N / 2) + 1) * 2;

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(...)
catch(std::exception& e)
{
AUD_THROW(FileException, "The HRTF name doesn't follow the naming scheme: " + filename);
}

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(...)
catch(std::exception& e)
{
AUD_THROW(FileException, "The HRTF name doesn't follow the naming scheme: " + filename);
}

View File

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

View File

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

View File

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

View File

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

View File

@@ -85,11 +85,15 @@ 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()
@@ -122,7 +126,11 @@ 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_SSE41_KERNEL_FLAGS "${CYCLES_SSE2_KERNEL_FLAGS} -msse3 -mssse3 -msse4.1")
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()
if(CXX_HAS_AVX2)
set(CYCLES_AVX2_KERNEL_FLAGS "${CYCLES_SSE41_KERNEL_FLAGS} -mavx -mavx2 -mfma -mlzcnt -mbmi -mbmi2 -mf16c")
endif()
@@ -136,8 +144,13 @@ 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()
@@ -161,8 +174,13 @@ 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()
@@ -172,10 +190,15 @@ 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()

View File

@@ -951,7 +951,9 @@ 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",
@@ -1671,19 +1673,19 @@ class CyclesPreferences(bpy.types.AddonPreferences):
elif device_type == 'HIP':
import sys
if sys.platform[:3] == "win":
col.label(text="Requires AMD GPU with RDNA architecture", icon='BLANK1')
col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
col.label(text="and AMD Radeon Pro 21.Q4 driver or newer", icon='BLANK1')
elif sys.platform.startswith("linux"):
col.label(text="Requires AMD GPU with RDNA architecture", icon='BLANK1')
col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
col.label(text="and AMD driver version 22.10 or newer", icon='BLANK1')
elif device_type == 'ONEAPI':
import sys
if sys.platform.startswith("win"):
col.label(text="Requires Intel GPU with Xe-HPG architecture", icon='BLANK1')
col.label(text="and Windows driver version 101.4032 or newer", icon='BLANK1')
col.label(text="and Windows driver version 101.3430 or newer", icon='BLANK1')
elif sys.platform.startswith("linux"):
col.label(text="Requires Intel GPU with Xe-HPG architecture and", icon='BLANK1')
col.label(text=" - intel-level-zero-gpu version 1.3.24931 or newer", icon='BLANK1')
col.label(text=" - intel-level-zero-gpu version 1.3.23904 or newer", icon='BLANK1')
col.label(text=" - oneAPI Level-Zero Loader", icon='BLANK1')
elif device_type == 'METAL':
col.label(text="Requires Apple Silicon with macOS 12.2 or newer", icon='BLANK1')

View File

@@ -2112,7 +2112,9 @@ 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")

View File

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

View File

@@ -63,7 +63,9 @@ 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. */

View File

@@ -26,7 +26,7 @@
CCL_NAMESPACE_BEGIN
typedef unordered_multimap<void *, ShaderInput *> PtrInputMap;
typedef map<void *, ShaderInput *> PtrInputMap;
typedef map<void *, ShaderOutput *> PtrOutputMap;
typedef map<string, ConvertNode *> ProxyMap;
@@ -1251,9 +1251,7 @@ static void add_nodes(Scene *scene,
ConvertNode *proxy = graph->create_node<ConvertNode>(to_socket_type, to_socket_type, true);
/* Muted nodes can result in multiple Cycles input sockets mapping to the same Blender
* input socket, so this needs to be a multimap. */
input_map.emplace(b_link.from_socket().ptr.data, proxy->inputs[0]);
input_map[b_link.from_socket().ptr.data] = proxy->inputs[0];
output_map[b_link.to_socket().ptr.data] = proxy->outputs[0];
graph->add(proxy);
@@ -1288,7 +1286,7 @@ static void add_nodes(Scene *scene,
/* register the proxy node for internal binding */
group_proxy_input_map[b_input.identifier()] = proxy;
input_map.emplace(b_input.ptr.data, proxy->inputs[0]);
input_map[b_input.ptr.data] = proxy->inputs[0];
set_default_value(proxy->inputs[0], b_input, b_data, b_ntree);
}
@@ -1340,7 +1338,7 @@ static void add_nodes(Scene *scene,
if (proxy_it != proxy_output_map.end()) {
ConvertNode *proxy = proxy_it->second;
input_map.emplace(b_input.ptr.data, proxy->inputs[0]);
input_map[b_input.ptr.data] = proxy->inputs[0];
set_default_value(proxy->inputs[0], b_input, b_data, b_ntree);
}
@@ -1371,7 +1369,7 @@ static void add_nodes(Scene *scene,
/* XXX should not happen, report error? */
continue;
}
input_map.emplace(b_input.ptr.data, input);
input_map[b_input.ptr.data] = input;
set_default_value(input, b_input, b_data, b_ntree);
}
@@ -1403,25 +1401,22 @@ static void add_nodes(Scene *scene,
BL::NodeSocket b_from_sock = b_link.from_socket();
BL::NodeSocket b_to_sock = b_link.to_socket();
ShaderOutput *output = nullptr;
ShaderOutput *output = 0;
ShaderInput *input = 0;
PtrOutputMap::iterator output_it = output_map.find(b_from_sock.ptr.data);
if (output_it != output_map.end())
output = output_it->second;
PtrInputMap::iterator input_it = input_map.find(b_to_sock.ptr.data);
if (input_it != input_map.end())
input = input_it->second;
/* either socket may be NULL when the node was not exported, typically
/* either node may be NULL when the node was not exported, typically
* because the node type is not supported */
if (output != nullptr) {
ShaderOutput *output = output_it->second;
auto inputs = input_map.equal_range(b_to_sock.ptr.data);
for (PtrInputMap::iterator input_it = inputs.first; input_it != inputs.second; ++input_it) {
ShaderInput *input = input_it->second;
if (input != nullptr) {
if (output && input)
graph->connect(output, input);
}
}
}
}
}
static void add_nodes(Scene *scene,
BL::RenderEngine &b_engine,

View File

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

View File

@@ -45,7 +45,9 @@ 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);

View File

@@ -9,7 +9,8 @@ CCL_NAMESPACE_BEGIN
#define KERNEL_FUNCTIONS(name) \
KERNEL_NAME_EVAL(cpu, name), KERNEL_NAME_EVAL(cpu_sse2, name), \
KERNEL_NAME_EVAL(cpu_sse41, name), KERNEL_NAME_EVAL(cpu_avx2, 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)
#define REGISTER_KERNEL(name) name(KERNEL_FUNCTIONS(name))
#define REGISTER_KERNEL_FILM_CONVERT(name) \

View File

@@ -17,10 +17,13 @@ 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_sse41, kernel_avx2);
kernel_info_ = get_best_kernel_info(
kernel_default, kernel_sse2, kernel_sse3, kernel_sse41, kernel_avx, kernel_avx2);
}
template<typename... Args> inline auto operator()(Args... args) const
@@ -57,12 +60,16 @@ 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
@@ -71,12 +78,24 @@ 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);

View File

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

View File

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

View File

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

View File

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

View File

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

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

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

View File

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

View File

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

View File

@@ -14,7 +14,9 @@ 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
)
@@ -938,9 +940,14 @@ 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()

View File

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

View File

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

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

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

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

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

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

View File

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

File diff suppressed because it is too large Load Diff

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

View File

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

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

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

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

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

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

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

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

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 wi,
const float3 wo,
const float3 I,
const float3 omega_in,
ccl_private float *pdf)
{
*pdf = 0.0f;
@@ -70,15 +70,15 @@ ccl_device Spectrum bsdf_transparent_eval(ccl_private const ShaderClosure *sc,
ccl_device int bsdf_transparent_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float3 I,
float randu,
float randv,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float3 *omega_in,
ccl_private float *pdf)
{
// only one direction is possible
*wo = -wi;
*omega_in = -I;
*pdf = 1;
*eval = one_spectrum();
return LABEL_TRANSMIT | LABEL_TRANSPARENT;

View File

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

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->wi);
bsdf->sample_weight *= bsdf_principled_diffuse_retro_reflection_sample_weight(bsdf, sd->I);
}
}

View File

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

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

View File

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

View File

@@ -35,9 +35,15 @@ 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"

View File

@@ -0,0 +1,26 @@
/* 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"

View File

@@ -0,0 +1,23 @@
/* 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"

View File

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

View File

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

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->wi - sd->dPdu * (dot(sd->dPdu, -sd->wi) / len_squared(sd->dPdu)));
tgN = -(-sd->I - sd->dPdu * (dot(sd->dPdu, -sd->I) / len_squared(sd->dPdu)));
tgN = normalize(tgN);
/* need to find suitable scaled gd for corrected normal */

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

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