Compare commits

..

17 Commits

Author SHA1 Message Date
3ffd662586 Merge branch 'master' into tmp-workbench-shader-create-infos 2022-01-26 10:21:47 +01:00
be4f593758 Merge branch 'master' into tmp-workbench-shader-create-infos 2022-01-25 18:51:46 +01:00
a9d9ec0586 Merge branch 'master' into tmp-workbench-shader-create-infos
# Conflicts:
#	source/blender/draw/intern/shaders/common_hair_lib.glsl
2022-01-25 15:16:54 +01:00
884370dd1d Port more of the draw glue layer to GPUShaderCreateInfo 2022-01-25 15:01:24 +01:00
e379839488 Merge branch 'master' into tmp-workbench-shader-create-infos 2022-01-24 21:28:00 +01:00
154e566ccb Cleanup: Share definition of world data UBO 2022-01-24 21:22:32 +01:00
988f01d96b Fix legacy modelmat path patch 2022-01-24 18:40:31 +01:00
b3292ff40a Fix legacy modelmat path 2022-01-24 15:36:30 +01:00
d46d62cb00 Fix image tansparency cutoff 2022-01-24 15:01:55 +01:00
4677b61d29 Workbench: Port shaders to use GPUShaderCreateInfo
Also adds a few things to GPUShader for easily create shaders.
Heavy usage of macros to compose the createInfo and avoid
duplications and copy paste bugs.
This makes the link between the shader request functions
(in `workbench_shader.cc`) and the actual createInfo a bit
obscure since the names are composed and not searchable.

Differential Revision: https://developer.blender.org/D13910
2022-01-24 13:36:48 +01:00
d5182edcb3 Fix shadow 2022-01-24 13:15:32 +01:00
de78a6ea6b Fix dof 2022-01-24 13:11:51 +01:00
fa12b82e47 Fix volume, udim, and shaders test 2022-01-24 12:16:45 +01:00
cd9873cadc Fix support for hairs 2022-01-19 21:56:43 +01:00
c50d50d683 Workbench: Port shaders to use GPUShaderCreateInfo 2022-01-19 21:56:33 +01:00
5ca850ff05 GPUShaderCreateInfo: Improve debugging experience
Add more meaningful reporting and a new easy way to compile
a shader from its createInfo name.
2022-01-19 15:40:05 +01:00
443d32baff GL: Fix missing sampler type support 2022-01-19 15:38:34 +01:00
957 changed files with 14166 additions and 23333 deletions

View File

@@ -157,9 +157,8 @@ option(WITH_BLENDER "Build blender (disable to build only the blender player)" O
mark_as_advanced(WITH_BLENDER)
if(APPLE)
# In future, can be used with `quicklookthumbnailing/qlthumbnailreply` to create file
# thumbnails for say Finder. Turn it off for now.
option(WITH_BLENDER_THUMBNAILER "Build \"blender-thumbnailer\" thumbnail extraction utility" OFF)
# Currently this causes a build error linking, disable.
set(WITH_BLENDER_THUMBNAILER OFF)
elseif(WIN32)
option(WITH_BLENDER_THUMBNAILER "Build \"BlendThumb.dll\" helper for Windows explorer integration" ON)
else()
@@ -686,7 +685,7 @@ if(WIN32 OR XCODE)
option(IDE_GROUP_PROJECTS_IN_FOLDERS "Organize the projects according to source folder structure." ON)
mark_as_advanced(IDE_GROUP_PROJECTS_IN_FOLDERS)
if(IDE_GROUP_PROJECTS_IN_FOLDERS)
if (IDE_GROUP_PROJECTS_IN_FOLDERS)
set_property(GLOBAL PROPERTY USE_FOLDERS ON)
endif()
endif()

View File

@@ -19,10 +19,13 @@
set(FREETYPE_EXTRA_ARGS
-DCMAKE_RELEASE_POSTFIX:STRING=2ST
-DCMAKE_DEBUG_POSTFIX:STRING=2ST_d
-DFT_DISABLE_BZIP2=ON
-DFT_DISABLE_HARFBUZZ=ON
-DFT_DISABLE_PNG=ON
-DFT_REQUIRE_BROTLI=ON
-DWITH_BZip2=OFF
-DWITH_HarfBuzz=OFF
-DFT_WITH_HARFBUZZ=OFF
-DFT_WITH_BZIP2=OFF
-DFT_WITH_BROTLI=ON
-DCMAKE_DISABLE_FIND_PACKAGE_HarfBuzz=TRUE
-DCMAKE_DISABLE_FIND_PACKAGE_BZip2=TRUE
-DPC_BROTLIDEC_INCLUDEDIR=${LIBDIR}/brotli/include
-DPC_BROTLIDEC_LIBDIR=${LIBDIR}/brotli/lib
)

View File

@@ -31,7 +31,7 @@ ExternalProject_Add(external_python_site_packages
CONFIGURE_COMMAND ${PIP_CONFIGURE_COMMAND}
BUILD_COMMAND ""
PREFIX ${BUILD_DIR}/site_packages
INSTALL_COMMAND ${PYTHON_BINARY} -m pip install --no-cache-dir ${SITE_PACKAGES_EXTRA} cython==${CYTHON_VERSION} idna==${IDNA_VERSION} charset-normalizer==${CHARSET_NORMALIZER_VERSION} urllib3==${URLLIB3_VERSION} certifi==${CERTIFI_VERSION} requests==${REQUESTS_VERSION} zstandard==${ZSTANDARD_VERSION} --no-binary :all:
INSTALL_COMMAND ${PYTHON_BINARY} -m pip install ${SITE_PACKAGES_EXTRA} cython==${CYTHON_VERSION} idna==${IDNA_VERSION} charset-normalizer==${CHARSET_NORMALIZER_VERSION} urllib3==${URLLIB3_VERSION} certifi==${CERTIFI_VERSION} requests==${REQUESTS_VERSION} zstandard==${ZSTANDARD_VERSION} --no-binary :all:
)
if(USE_PIP_NUMPY)

View File

@@ -83,9 +83,9 @@ else()
set(OPENEXR_VERSION_POSTFIX)
endif()
set(FREETYPE_VERSION 2.11.1)
set(FREETYPE_VERSION 2.11.0)
set(FREETYPE_URI http://prdownloads.sourceforge.net/freetype/freetype-${FREETYPE_VERSION}.tar.gz)
set(FREETYPE_HASH bd4e3b007474319909a6b79d50908e85)
set(FREETYPE_HASH cf09172322f6b50cf8f568bf8fe14bde)
set(FREETYPE_HASH_TYPE MD5)
set(FREETYPE_FILE freetype-${FREETYPE_VERSION}.tar.gz)
@@ -189,11 +189,11 @@ set(OSL_HASH 1abd7ce40481771a9fa937f19595d2f2)
set(OSL_HASH_TYPE MD5)
set(OSL_FILE OpenShadingLanguage-${OSL_VERSION}.tar.gz)
set(PYTHON_VERSION 3.10.2)
set(PYTHON_SHORT_VERSION 3.10)
set(PYTHON_SHORT_VERSION_NO_DOTS 310)
set(PYTHON_VERSION 3.9.7)
set(PYTHON_SHORT_VERSION 3.9)
set(PYTHON_SHORT_VERSION_NO_DOTS 39)
set(PYTHON_URI https://www.python.org/ftp/python/${PYTHON_VERSION}/Python-${PYTHON_VERSION}.tar.xz)
set(PYTHON_HASH 14e8c22458ed7779a1957b26cde01db9)
set(PYTHON_HASH fddb060b483bc01850a3f412eea1d954)
set(PYTHON_HASH_TYPE MD5)
set(PYTHON_FILE Python-${PYTHON_VERSION}.tar.xz)
@@ -215,20 +215,18 @@ set(NANOVDB_HASH e7b9e863ec2f3b04ead171dec2322807)
set(NANOVDB_HASH_TYPE MD5)
set(NANOVDB_FILE nano-vdb-${NANOVDB_GIT_UID}.tar.gz)
set(IDNA_VERSION 3.3)
set(CHARSET_NORMALIZER_VERSION 2.0.10)
set(URLLIB3_VERSION 1.26.8)
set(IDNA_VERSION 3.2)
set(CHARSET_NORMALIZER_VERSION 2.0.6)
set(URLLIB3_VERSION 1.26.7)
set(CERTIFI_VERSION 2021.10.8)
set(REQUESTS_VERSION 2.27.1)
set(CYTHON_VERSION 0.29.26)
# The version of the zstd library used to build the Python package should match ZSTD_VERSION defined below.
# At this time of writing, 0.17.0 was already released, but built against zstd 1.5.1, while we use 1.5.0.
set(ZSTANDARD_VERSION 0.16.0)
set(REQUESTS_VERSION 2.26.0)
set(CYTHON_VERSION 0.29.24)
set(ZSTANDARD_VERSION 0.15.2 )
set(NUMPY_VERSION 1.22.0)
set(NUMPY_SHORT_VERSION 1.22)
set(NUMPY_VERSION 1.21.2)
set(NUMPY_SHORT_VERSION 1.21)
set(NUMPY_URI https://github.com/numpy/numpy/releases/download/v${NUMPY_VERSION}/numpy-${NUMPY_VERSION}.zip)
set(NUMPY_HASH 252de134862a27bd66705d29622edbfe)
set(NUMPY_HASH 5638d5dae3ca387be562912312db842e)
set(NUMPY_HASH_TYPE MD5)
set(NUMPY_FILE numpy-${NUMPY_VERSION}.zip)

View File

@@ -379,27 +379,27 @@ USE_CXX11=true
CLANG_FORMAT_VERSION_MIN="6.0"
CLANG_FORMAT_VERSION_MEX="10.0"
PYTHON_VERSION="3.10.2"
PYTHON_VERSION_SHORT="3.10"
PYTHON_VERSION_MIN="3.9"
PYTHON_VERSION_MEX="3.12"
PYTHON_VERSION="3.9.7"
PYTHON_VERSION_SHORT="3.9"
PYTHON_VERSION_MIN="3.7"
PYTHON_VERSION_MEX="3.11"
PYTHON_VERSION_INSTALLED=$PYTHON_VERSION_SHORT
PYTHON_FORCE_BUILD=false
PYTHON_FORCE_REBUILD=false
PYTHON_SKIP=false
# Additional Python modules.
PYTHON_IDNA_VERSION="3.3"
PYTHON_IDNA_VERSION="3.2"
PYTHON_IDNA_VERSION_MIN="2.0"
PYTHON_IDNA_VERSION_MEX="4.0"
PYTHON_IDNA_NAME="idna"
PYTHON_CHARSET_NORMALIZER_VERSION="2.0.10"
PYTHON_CHARSET_NORMALIZER_VERSION="2.0.6"
PYTHON_CHARSET_NORMALIZER_VERSION_MIN="2.0.6"
PYTHON_CHARSET_NORMALIZER_VERSION_MEX="2.1.0" # requests uses `charset_normalizer~=2.0.0`
PYTHON_CHARSET_NORMALIZER_NAME="charset-normalizer"
PYTHON_URLLIB3_VERSION="1.26.8"
PYTHON_URLLIB3_VERSION="1.26.7"
PYTHON_URLLIB3_VERSION_MIN="1.0"
PYTHON_URLLIB3_VERSION_MEX="2.0"
PYTHON_URLLIB3_NAME="urllib3"
@@ -409,17 +409,17 @@ PYTHON_CERTIFI_VERSION_MIN="2021.0"
PYTHON_CERTIFI_VERSION_MEX="2023.0"
PYTHON_CERTIFI_NAME="certifi"
PYTHON_REQUESTS_VERSION="2.27.1"
PYTHON_REQUESTS_VERSION="2.23.0"
PYTHON_REQUESTS_VERSION_MIN="2.0"
PYTHON_REQUESTS_VERSION_MEX="3.0"
PYTHON_REQUESTS_NAME="requests"
PYTHON_ZSTANDARD_VERSION="0.16.0"
PYTHON_ZSTANDARD_VERSION="0.15.2"
PYTHON_ZSTANDARD_VERSION_MIN="0.15.2"
PYTHON_ZSTANDARD_VERSION_MEX="0.20.0"
PYTHON_ZSTANDARD_VERSION_MEX="0.16.0"
PYTHON_ZSTANDARD_NAME="zstandard"
PYTHON_NUMPY_VERSION="1.22.0"
PYTHON_NUMPY_VERSION="1.21.2"
PYTHON_NUMPY_VERSION_MIN="1.14"
PYTHON_NUMPY_VERSION_MEX="2.0"
PYTHON_NUMPY_NAME="numpy"
@@ -499,7 +499,7 @@ LLVM_FORCE_REBUILD=false
LLVM_SKIP=false
# OSL needs to be compiled for now!
OSL_VERSION="1.11.17.0"
OSL_VERSION="1.11.14.1"
OSL_VERSION_SHORT="1.11"
OSL_VERSION_MIN="1.11"
OSL_VERSION_MEX="2.0"
@@ -4036,14 +4036,14 @@ install_DEB() {
INFO "Forced Python building, as requested..."
_do_compile_python=true
else
check_package_version_ge_lt_DEB python${PYTHON_VERSION_SHORT}-dev $PYTHON_VERSION_MIN $PYTHON_VERSION_MEX
check_package_version_ge_lt_DEB python3-dev $PYTHON_VERSION_MIN $PYTHON_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_DEB python${PYTHON_VERSION_SHORT}-dev
PYTHON_VERSION_INSTALLED=$(echo `get_package_version_DEB python3-dev` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
install_packages_DEB python3-dev
clean_Python
PRINT ""
PYTHON_VERSION_INSTALLED=$(echo `get_package_version_DEB python${PYTHON_VERSION_SHORT}-dev` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
for module in "${PYTHON_MODULES_PACKAGES[@]}"
do
module=($module)
@@ -4681,11 +4681,11 @@ install_RPM() {
else
check_package_version_ge_lt_RPM python3-devel $PYTHON_VERSION_MIN $PYTHON_VERSION_MEX
if [ $? -eq 0 ]; then
PYTHON_VERSION_INSTALLED=$(echo `get_package_version_RPM python3-devel` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
install_packages_RPM python3-devel
clean_Python
PYTHON_VERSION_INSTALLED=$(echo `get_package_version_RPM python3-devel` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
for module in "${PYTHON_MODULES_PACKAGES[@]}"
do
module=($module)
@@ -5224,12 +5224,12 @@ install_ARCH() {
else
check_package_version_ge_lt_ARCH python $PYTHON_VERSION_MIN $PYTHON_VERSION_MEX
if [ $? -eq 0 ]; then
PYTHON_VERSION_INSTALLED=$(echo `get_package_version_ARCH python` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
install_packages_ARCH python
clean_Python
PRINT ""
PYTHON_VERSION_INSTALLED=$(echo `get_package_version_ARCH python` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
for module in "${PYTHON_MODULES_PACKAGES[@]}"
do
module=($module)

View File

@@ -34,7 +34,7 @@ IF(NOT PYTHON_ROOT_DIR AND NOT $ENV{PYTHON_ROOT_DIR} STREQUAL "")
SET(PYTHON_ROOT_DIR $ENV{PYTHON_ROOT_DIR})
ENDIF()
SET(PYTHON_VERSION 3.10 CACHE STRING "Python Version (major and minor only)")
SET(PYTHON_VERSION 3.9 CACHE STRING "Python Version (major and minor only)")
MARK_AS_ADVANCED(PYTHON_VERSION)

View File

@@ -1197,21 +1197,21 @@ endfunction()
macro(openmp_delayload
projectname
)
if(MSVC)
if(WITH_OPENMP)
if(MSVC_CLANG)
set(OPENMP_DLL_NAME "libomp")
elseif(MSVC_VERSION EQUAL 1800)
set(OPENMP_DLL_NAME "vcomp120")
else()
set(OPENMP_DLL_NAME "vcomp140")
if(MSVC)
if(WITH_OPENMP)
if(MSVC_CLANG)
set(OPENMP_DLL_NAME "libomp")
elseif(MSVC_VERSION EQUAL 1800)
set(OPENMP_DLL_NAME "vcomp120")
else()
set(OPENMP_DLL_NAME "vcomp140")
endif()
set_property(TARGET ${projectname} APPEND_STRING PROPERTY LINK_FLAGS_RELEASE " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
set_property(TARGET ${projectname} APPEND_STRING PROPERTY LINK_FLAGS_DEBUG " /DELAYLOAD:${OPENMP_DLL_NAME}d.dll delayimp.lib")
set_property(TARGET ${projectname} APPEND_STRING PROPERTY LINK_FLAGS_RELWITHDEBINFO " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
set_property(TARGET ${projectname} APPEND_STRING PROPERTY LINK_FLAGS_MINSIZEREL " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
endif()
set_property(TARGET ${projectname} APPEND_STRING PROPERTY LINK_FLAGS_RELEASE " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
set_property(TARGET ${projectname} APPEND_STRING PROPERTY LINK_FLAGS_DEBUG " /DELAYLOAD:${OPENMP_DLL_NAME}d.dll delayimp.lib")
set_property(TARGET ${projectname} APPEND_STRING PROPERTY LINK_FLAGS_RELWITHDEBINFO " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
set_property(TARGET ${projectname} APPEND_STRING PROPERTY LINK_FLAGS_MINSIZEREL " /DELAYLOAD:${OPENMP_DLL_NAME}.dll delayimp.lib")
endif()
endif()
endmacro()
macro(set_and_warn_dependency

View File

@@ -128,20 +128,25 @@ if(WITH_CODEC_SNDFILE)
endif()
if(WITH_PYTHON)
# Use precompiled libraries by default.
set(PYTHON_VERSION 3.10)
# we use precompiled libraries for py 3.9 and up by default
set(PYTHON_VERSION 3.9)
if(NOT WITH_PYTHON_MODULE AND NOT WITH_PYTHON_FRAMEWORK)
# Normally cached but not since we include them with blender.
# normally cached but not since we include them with blender
set(PYTHON_INCLUDE_DIR "${LIBDIR}/python/include/python${PYTHON_VERSION}")
set(PYTHON_EXECUTABLE "${LIBDIR}/python/bin/python${PYTHON_VERSION}")
set(PYTHON_LIBRARY ${LIBDIR}/python/lib/libpython${PYTHON_VERSION}.a)
set(PYTHON_LIBPATH "${LIBDIR}/python/lib/python${PYTHON_VERSION}")
# set(PYTHON_LINKFLAGS "-u _PyMac_Error") # won't build with this enabled
else()
# Module must be compiled against Python framework.
# module must be compiled against Python framework
set(_py_framework "/Library/Frameworks/Python.framework/Versions/${PYTHON_VERSION}")
set(PYTHON_INCLUDE_DIR "${_py_framework}/include/python${PYTHON_VERSION}")
set(PYTHON_EXECUTABLE "${_py_framework}/bin/python${PYTHON_VERSION}")
set(PYTHON_LIBPATH "${_py_framework}/lib/python${PYTHON_VERSION}")
# set(PYTHON_LIBRARY python${PYTHON_VERSION})
# set(PYTHON_LINKFLAGS "-u _PyMac_Error -framework Python") # won't build with this enabled
unset(_py_framework)
endif()

View File

@@ -109,14 +109,9 @@ if(NOT WITH_SYSTEM_FREETYPE)
find_package_wrapper(Freetype REQUIRED)
if(EXISTS ${LIBDIR})
find_package_wrapper(Brotli REQUIRED)
# NOTE: This is done on WIN32 & APPLE but fails on some Linux systems.
# See: https://devtalk.blender.org/t/22536
# So `BROTLI_LIBRARIES` need to be added directly after `FREETYPE_LIBRARIES`.
#
# list(APPEND FREETYPE_LIBRARIES
# ${BROTLI_LIBRARIES}
# )
list(APPEND FREETYPE_LIBRARIES
${BROTLI_LIBRARIES}
)
endif()
endif()

View File

@@ -55,10 +55,6 @@ if(CMAKE_C_COMPILER_ID MATCHES "Clang")
message(WARNING "stripped pdb not supported with clang, disabling..")
set(WITH_WINDOWS_STRIPPED_PDB OFF)
endif()
else()
if(CMAKE_CXX_COMPILER_VERSION VERSION_LESS 19.28.29921) # MSVC 2019 16.9.16
message(FATAL_ERROR "Compiler is unsupported, MSVC 2019 16.9.16 or newer is required for building blender.")
endif()
endif()
if(NOT WITH_PYTHON_MODULE)
@@ -269,6 +265,12 @@ if(NOT DEFINED LIBDIR)
elseif(MSVC_VERSION GREATER 1919)
message(STATUS "Visual Studio 2019 detected.")
set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_BASE}_vc15)
elseif(MSVC_VERSION GREATER 1909)
message(STATUS "Visual Studio 2017 detected.")
set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_BASE}_vc15)
elseif(MSVC_VERSION EQUAL 1900)
message(STATUS "Visual Studio 2015 detected.")
set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_BASE}_vc15)
endif()
else()
message(STATUS "Using pre-compiled LIBDIR: ${LIBDIR}")
@@ -463,7 +465,7 @@ if(WITH_JACK)
endif()
if(WITH_PYTHON)
set(PYTHON_VERSION 3.10) # CACHE STRING)
set(PYTHON_VERSION 3.9) # CACHE STRING)
string(REPLACE "." "" _PYTHON_VERSION_NO_DOTS ${PYTHON_VERSION})
set(PYTHON_LIBRARY ${LIBDIR}/python/${_PYTHON_VERSION_NO_DOTS}/libs/python${_PYTHON_VERSION_NO_DOTS}.lib)

View File

@@ -3,6 +3,9 @@ echo No explicit msvc version requested, autodetecting version.
call "%~dp0\detect_msvc2019.cmd"
if %ERRORLEVEL% EQU 0 goto DetectionComplete
call "%~dp0\detect_msvc2017.cmd"
if %ERRORLEVEL% EQU 0 goto DetectionComplete
call "%~dp0\detect_msvc2022.cmd"
if %ERRORLEVEL% EQU 0 goto DetectionComplete

View File

@@ -1,3 +1,4 @@
if "%BUILD_VS_YEAR%"=="2017" set BUILD_VS_LIBDIRPOST=vc15
if "%BUILD_VS_YEAR%"=="2019" set BUILD_VS_LIBDIRPOST=vc15
if "%BUILD_VS_YEAR%"=="2022" set BUILD_VS_LIBDIRPOST=vc15

View File

@@ -19,6 +19,12 @@ if "%WITH_PYDEBUG%"=="1" (
set PYDEBUG_CMAKE_ARGS=-DWINDOWS_PYTHON_DEBUG=On
)
if "%BUILD_VS_YEAR%"=="2017" (
set BUILD_GENERATOR_POST=%WINDOWS_ARCH%
) else (
set BUILD_PLATFORM_SELECT=-A %MSBUILD_PLATFORM%
)
set BUILD_CMAKE_ARGS=%BUILD_CMAKE_ARGS% -G "Visual Studio %BUILD_VS_VER% %BUILD_VS_YEAR%%BUILD_GENERATOR_POST%" %BUILD_PLATFORM_SELECT% %TESTS_CMAKE_ARGS% %CLANG_CMAKE_ARGS% %ASAN_CMAKE_ARGS% %PYDEBUG_CMAKE_ARGS%
if NOT EXIST %BUILD_DIR%\nul (

View File

@@ -37,9 +37,15 @@ set LLVM_DIR=
:DetectionComplete
set CC=%LLVM_DIR%\bin\clang-cl
set CXX=%LLVM_DIR%\bin\clang-cl
rem build and tested against 2019 16.2
set CFLAGS=-m64 -fmsc-version=1922
set CXXFLAGS=-m64 -fmsc-version=1922
if "%BUILD_VS_YEAR%" == "2019" (
rem build and tested against 2019 16.2
set CFLAGS=-m64 -fmsc-version=1922
set CXXFLAGS=-m64 -fmsc-version=1922
) else (
rem build and tested against 2017 15.7
set CFLAGS=-m64 -fmsc-version=1914
set CXXFLAGS=-m64 -fmsc-version=1914
)
)
if "%WITH_ASAN%"=="1" (

View File

@@ -0,0 +1,3 @@
set BUILD_VS_VER=15
set BUILD_VS_YEAR=2017
call "%~dp0\detect_msvc_vswhere.cmd"

View File

@@ -50,6 +50,14 @@ if NOT "%1" == "" (
goto ERR
) else if "%1" == "x64" (
set BUILD_ARCH=x64
) else if "%1" == "2017" (
set BUILD_VS_YEAR=2017
) else if "%1" == "2017pre" (
set BUILD_VS_YEAR=2017
set VSWHERE_ARGS=-prerelease
) else if "%1" == "2017b" (
set BUILD_VS_YEAR=2017
set VSWHERE_ARGS=-products Microsoft.VisualStudio.Product.BuildTools
) else if "%1" == "2019" (
set BUILD_VS_YEAR=2019
) else if "%1" == "2019pre" (

View File

@@ -24,12 +24,12 @@ echo - nobuildinfo ^(disable buildinfo^)
echo - debug ^(Build an unoptimized debuggable build^)
echo - packagename [newname] ^(override default cpack package name^)
echo - builddir [newdir] ^(override default build folder^)
echo - 2017 ^(build with visual studio 2017^)
echo - 2017pre ^(build with visual studio 2017 pre-release^)
echo - 2017b ^(build with visual studio 2017 Build Tools^)
echo - 2019 ^(build with visual studio 2019^)
echo - 2019pre ^(build with visual studio 2019 pre-release^)
echo - 2019b ^(build with visual studio 2019 Build Tools^)
echo - 2022 ^(build with visual studio 2022^)
echo - 2022pre ^(build with visual studio 2022 pre-release^)
echo - 2022b ^(build with visual studio 2022 Build Tools^)
echo.
echo Documentation Targets ^(Not associated with building^)

View File

@@ -1,3 +1,4 @@
if "%BUILD_VS_YEAR%"=="2017" set BUILD_VS_LIBDIRPOST=vc15
if "%BUILD_VS_YEAR%"=="2019" set BUILD_VS_LIBDIRPOST=vc15
if "%BUILD_VS_YEAR%"=="2022" set BUILD_VS_LIBDIRPOST=vc15

View File

@@ -38,7 +38,7 @@ PROJECT_NAME = Blender
# could be handy for archiving the generated documentation or if some version
# control system is used.
PROJECT_NUMBER = V3.2
PROJECT_NUMBER = V3.1
# Using the PROJECT_BRIEF tag one can provide an optional one line description
# for a project that appears at the top of each page and should give viewer a

View File

@@ -22,7 +22,7 @@ Data Access
===========
The most common case for using the reference API is to find out how to access data in the blend-file.
Before going any further it's best to be aware of ID data-blocks in Blender since you will often find properties
Before going any further its best to be aware of ID data-blocks in Blender since you will often find properties
relative to them.
@@ -55,9 +55,9 @@ Start by collecting the information where the data is located.
First find this setting in the interface ``Properties editor -> Object -> Transform -> Location``.
From the button context menu select *Online Python Reference*, this will link you to:
:class:`bpy.types.Object.location`.
Being an API reference, this link often gives little more information than the tooltip, though some of the pages
Being an API reference, this link often gives little more information then the tooltip, though some of the pages
include examples (normally at the top of the page).
But you now know that you have to use ``.location`` and that it's an array of three floats.
But you now know that you have to use ``.location`` and that its an array of three floats.
So the next step is to find out where to access objects, go down to the bottom of the page to the references section,
for objects there are many references, but one of the most common places to access objects is via the context.
@@ -154,7 +154,7 @@ The tooltip includes :class:`bpy.types.SubsurfModifier.levels` but you want the
Note that the text copied won't include the ``bpy.data.collection["name"].`` component since its assumed that
you won't be doing collection look-ups on every access and typically you'll want to use the context rather
than access each :class:`bpy.types.ID` instance by name.
then access each :class:`bpy.types.ID` instance by name.
Type in the ID path into a Python console :mod:`bpy.context.active_object`.
Include the trailing dot and don't execute the code, yet.
@@ -252,6 +252,6 @@ Each entry can be selected, then copied :kbd:`Ctrl-C`, usually to paste in the t
.. note::
Not all operators get registered for display,
zooming the view for example isn't so useful to repeat so it's excluded from the output.
zooming the view for example isn't so useful to repeat so its excluded from the output.
To display *every* operator that runs see :ref:`Show All Operators <info_show_all_operators>`.

View File

@@ -229,7 +229,7 @@ removing the last items first, which is faster (as explained above):
This example shows a fast way of removing items,
for use in cases where you can alter the list order without breaking the script's functionality.
for use in cases where you can alter the list order without breaking the scripts functionality.
This works by swapping two list items, so the item you remove is always last:
.. code-block:: python
@@ -278,7 +278,7 @@ Here are three ways of joining multiple strings into one string for writing.
This also applies to any area of your code that involves a lot of string joining:
String concatenation
This is the slowest option, do **not** use this if you can avoid it, especially when writing data in a loop.
This is the slowest option, do **not** use if you can avoid it, especially when writing data in a loop.
>>> file.write(str1 + " " + str2 + " " + str3 + "\n")
@@ -288,7 +288,7 @@ String formatting
>>> file.write("%s %s %s\n" % (str1, str2, str3))
String joining
Use this to join a list of strings (the list may be temporary). In the following example, the strings are joined with
Use to join a list of strings (the list may be temporary). In the following example, the strings are joined with
a space " " in between, other examples are "" or ", ".
>>> file.write(" ".join((str1, str2, str3, "\n")))

View File

@@ -12,7 +12,7 @@ that can be troublesome and avoid practices that are known to cause instability.
Using Operators
===============
Blender's operators are tools for users to access, that can be accessed with Python too which is very useful.
Blender's operators are tools for users to access, that can access with Python too which is very useful.
Still operators have limitations that can make them cumbersome to script.
The main limits are:
@@ -20,13 +20,13 @@ The main limits are:
- Can't pass data such as objects, meshes or materials to operate on (operators use the context instead).
- The return value from calling an operator is the success (if it finished or was canceled),
in some cases it would be more logical from an API perspective to return the result of the operation.
- Operators' poll function can fail where an API function would raise an exception giving details on exactly why.
- Operators poll function can fail where an API function would raise an exception giving details on exactly why.
Why does an operator's poll fail?
---------------------------------
When calling an operator it gives an error like this:
When calling an operator gives an error like this:
>>> bpy.ops.action.clean(threshold=0.001)
RuntimeError: Operator bpy.ops.action.clean.poll() failed, context is incorrect
@@ -49,9 +49,9 @@ you should be able to find the poll function with no knowledge of C.
.. note::
Blender does have the functionality for poll functions to describe why they fail,
but it's currently not used much, if you're interested to help improve the API
but its currently not used much, if you're interested to help improve the API
feel free to add calls to :class:`bpy.types.Operator.poll_message_set` (``CTX_wm_operator_poll_msg_set`` in C)
where it's not obvious why poll fails, e.g:
where its not obvious why poll fails, e.g:
>>> bpy.ops.gpencil.draw()
RuntimeError: Operator bpy.ops.gpencil.draw.poll() Failed to find Grease Pencil data to draw into
@@ -107,7 +107,7 @@ In this case you need to call :class:`bpy.types.ViewLayer.update` after modifyin
Now all dependent data (child objects, modifiers, drivers, etc.)
have been recalculated and are available to the script within the active view layer.
has been recalculated and is available to the script within active view layer.
Can I redraw during script execution?
@@ -116,13 +116,13 @@ Can I redraw during script execution?
The official answer to this is no, or... *"You don't want to do that"*.
To give some background on the topic:
While a script executes, Blender waits for it to finish and is effectively locked until it's done;
While a script executes Blender waits for it to finish and is effectively locked until its done,
while in this state Blender won't redraw or respond to user input.
Normally this is not such a problem because scripts distributed with Blender
tend not to run for an extended period of time,
nevertheless scripts *can* take a long time to complete and it would be nice to see progress in the viewport.
Tools that lock Blender in a loop redraw are highly discouraged
When tools lock Blender in a loop redraw are highly discouraged
since they conflict with Blender's ability to run multiple operators
at once and update different parts of the interface as the tool runs.
@@ -130,7 +130,7 @@ So the solution here is to write a **modal** operator, which is an operator that
See the modal operator template in the text editor.
Modal operators execute on user input or setup their own timers to run frequently,
they can handle the events or pass through to be handled by the keymap or other modal operators.
Examples of modal operators are Transform, Painting, Fly Navigation and File Select.
Examples of a modal operators are Transform, Painting, Fly Navigation and File Select.
Writing modal operators takes more effort than a simple ``for`` loop
that contains draw calls but is more flexible and integrates better with Blender's design.
@@ -240,7 +240,7 @@ Editing
Editing is where the three data types vary most.
- Polygons are very limited for editing,
changing materials and options like smooth works, but for anything else
changing materials and options like smooth works but for anything else
they are too inflexible and are only intended for storage.
- Tessfaces should not be used for editing geometry because doing so will cause existing n-gons to be tessellated.
- BMesh-faces are by far the best way to manipulate geometry.
@@ -256,7 +256,7 @@ the choice mostly depends on whether the target format supports n-gons or not.
- Tessfaces work well for exporting to formats which don't support n-gons,
in fact this is the only place where their use is encouraged.
- BMesh-Faces can work for exporting too but may not be necessary if polygons can be used
since using BMesh gives some overhead because it's not the native storage format in Object-Mode.
since using BMesh gives some overhead because its not the native storage format in Object-Mode.
Edit Bones, Pose Bones, Bone... Bones
@@ -348,7 +348,7 @@ Armature Mode Switching
While writing scripts that deal with armatures you may find you have to switch between modes,
when doing so take care when switching out of Edit-Mode not to keep references
to the edit bones or their head/tail vectors.
Further access to these will crash Blender so it's important that the script
Further access to these will crash Blender so its important the script
clearly separates sections of the code which operate in different modes.
This is mainly an issue with Edit-Mode since pose data can be manipulated without having to be in Pose-Mode,
@@ -386,11 +386,11 @@ Or with name assignment:
Data names may not match the assigned values if they exceed the maximum length, are already used or an empty string.
It's better practice not to reference objects by names at all,
Its better practice not to reference objects by names at all,
once created you can store the data in a list, dictionary, on a class, etc;
there is rarely a reason to have to keep searching for the same data by name.
If you do need to use name references, it's best to use a dictionary to maintain
If you do need to use name references, its best to use a dictionary to maintain
a mapping between the names of the imported assets and the newly created data,
this way you don't run this risk of referencing existing data from the blend-file, or worse modifying it.
@@ -414,11 +414,11 @@ Library Collisions
Blender keeps data names unique (:class:`bpy.types.ID.name`) so you can't name two objects,
meshes, scenes, etc., the same by accident.
However, when linking in library data from another blend-file naming collisions can occur,
so it's best to avoid referencing data by name at all.
so its best to avoid referencing data by name at all.
This can be tricky at times and not even Blender handles this correctly in some cases
This can be tricky at times and not even Blender handles this correctly in some case
(when selecting the modifier object for e.g. you can't select between multiple objects with the same name),
but it's still good to try avoiding these problems in this area.
but its still good to try avoiding these problems in this area.
If you need to select between local and library data, there is a feature in ``bpy.data`` members to allow for this.
.. code-block:: python
@@ -467,11 +467,11 @@ writing a script in ``latin1`` or ``iso-8859-15``.
See `PEP 263 <https://www.python.org/dev/peps/pep-0263/>`__.
However, this complicates matters for Blender's Python API because ``.blend`` files don't have an explicit encoding.
To avoid the problem for Python integration and script authors we have decided that all strings in blend-files
To avoid the problem for Python integration and script authors we have decided all strings in blend-files
**must** be ``UTF-8``, ``ASCII`` compatible.
This means assigning strings with different encodings to an object name, for instance, will raise an error.
This means assigning strings with different encodings to an object names for instance will raise an error.
Paths are an exception to this rule since the existence of non-UTF-8 paths on the user's file system cannot be ignored.
Paths are an exception to this rule since the existence of non-UTF-8 paths on user's file system cannot be ignored.
This means seemingly harmless expressions can raise errors, e.g:
>>> print(bpy.data.filepath)
@@ -505,7 +505,7 @@ to keep it short about encoding problems -- here are some suggestions:
.. note::
Sometimes it's preferable to avoid string encoding issues by using bytes instead of Python strings,
when reading some input it's less trouble to read it as binary data
when reading some input its less trouble to read it as binary data
though you will still need to decide how to treat any strings you want to use with Blender,
some importers do this.
@@ -679,7 +679,7 @@ Undo/Redo
---------
For safety, you should assume that undo and redo always invalidates all :class:`bpy.types.ID`
instances (Object, Scene, Mesh, Light, etc.), as well obviously as all of their sub-data.
instances (Object, Scene, Mesh, Light, etc.), as weel obviously as all of their sub-data.
This example shows how you can tell undo changes the memory locations:
@@ -716,7 +716,7 @@ Tools in Blender are not allowed to modify library data.
But Python does not enforce this restriction.
This can be useful in some cases, using a script to adjust material values for example.
But it's also possible to use a script to make library data point to newly created local data,
But its also possible to use a script to make library data point to newly created local data,
which is not supported since a call to undo will remove the local data
but leave the library referencing it and likely crash.

View File

@@ -81,7 +81,7 @@ but reference an external file rather than including it directly.
Executing External Scripts
--------------------------
This is the equivalent to running the script directly, referencing a script's path from a two line code block.
This is the equivalent to running the script directly, referencing a scripts path from a two line code block.
.. code-block:: python
@@ -124,7 +124,7 @@ small script which is often useful for testing different settings quickly.
The other issue with this is the script has to be in Python's module search path.
While this is not best practice -- for testing purposes you can extend the search path,
this following example adds the current blend-file's directory to the search path
this following example adds the current blend-files directory to the search path
and then loads the script as a module.
.. code-block:: python
@@ -302,7 +302,7 @@ Python Safety (Build Option)
----------------------------
Since it's possible to access data which has been removed (see :doc:`Gotchas <info_gotcha>`),
it can be hard to track down the cause of crashes.
can make it hard to track down the cause of crashes.
To raise Python exceptions on accessing freed data (rather than crashing),
enable the CMake build option ``WITH_PYTHON_SAFETY``.
This enables data tracking which makes data access about two times slower

View File

@@ -417,8 +417,7 @@ MODULE_GROUPING = {
BLENDER_REVISION = str(bpy.app.build_hash, 'utf_8')
# '2.83.0 Beta' or '2.83.0' or '2.83.1'
BLENDER_VERSION_STRING = bpy.app.version_string
BLENDER_VERSION_DOTS = "%d.%d" % (bpy.app.version[0], bpy.app.version[1])
BLENDER_VERSION_DOTS = bpy.app.version_string
if BLENDER_REVISION != "Unknown":
# SHA1 Git hash
@@ -1725,11 +1724,11 @@ def write_sphinx_conf_py(basepath):
fw("import sys, os\n\n")
fw("extensions = ['sphinx.ext.intersphinx']\n\n")
fw("intersphinx_mapping = {'blender_manual': ('https://docs.blender.org/manual/en/dev/', None)}\n\n")
fw("project = 'Blender %s Python API'\n" % BLENDER_VERSION_STRING)
fw("project = 'Blender %s Python API'\n" % BLENDER_VERSION_DOTS)
fw("master_doc = 'index'\n")
fw("copyright = u'Blender Foundation'\n")
fw("version = '%s'\n" % BLENDER_VERSION_DOTS)
fw("release = '%s'\n" % BLENDER_VERSION_DOTS)
fw("version = '%s'\n" % BLENDER_VERSION_HASH)
fw("release = '%s'\n" % BLENDER_VERSION_HASH)
# Quiet file not in table-of-contents warnings.
fw("exclude_patterns = [\n")
@@ -1750,7 +1749,6 @@ except ModuleNotFoundError:
fw("if html_theme == 'sphinx_rtd_theme':\n")
fw(" html_theme_options = {\n")
fw(" 'display_version': False,\n")
# fw(" 'analytics_id': '',\n")
# fw(" 'collapse_navigation': True,\n")
fw(" 'sticky_navigation': False,\n")
@@ -1767,15 +1765,10 @@ except ModuleNotFoundError:
fw("html_show_search_summary = True\n")
fw("html_split_index = True\n")
fw("html_static_path = ['static']\n")
fw("templates_path = ['templates']\n")
fw("html_context = {'commit': '%s'}\n" % BLENDER_VERSION_HASH)
fw("html_extra_path = ['static/favicon.ico', 'static/blender_logo.svg']\n")
fw("html_favicon = 'static/favicon.ico'\n")
fw("html_logo = 'static/blender_logo.svg'\n")
fw("html_last_updated_fmt = '%m/%d/%Y'\n\n")
fw("if html_theme == 'sphinx_rtd_theme':\n")
fw(" html_css_files = ['css/version_switch.css']\n")
fw(" html_js_files = ['js/version_switch.js']\n")
# needed for latex, pdf gen
fw("latex_elements = {\n")
@@ -2132,9 +2125,6 @@ def copy_theme_assets(basepath):
shutil.copytree(os.path.join(SCRIPT_DIR, "static"),
os.path.join(basepath, "static"),
copy_function=shutil.copy)
shutil.copytree(os.path.join(SCRIPT_DIR, "templates"),
os.path.join(basepath, "templates"),
copy_function=shutil.copy)
def rna2sphinx(basepath):

View File

@@ -1,127 +0,0 @@
/* Override RTD theme */
.rst-versions {
border-top: 0px;
overflow: visible;
}
.version-btn.vdeact {
cursor: default;
color: dimgray;
}
.version-btn.vdeact::after {
content: "";
}
#versionwrap {
display: flex;
padding-top: 2px;
font-size: 90%;
justify-content: center;
flex-wrap: wrap;
}
.version-btn {
display: inline-block;
background-color: #272525;
width: 140px;
text-align: center;
padding: 3px 10px;
margin: 0px 5px 4px;
vertical-align: middle;
color: #27AE60;
border: solid 1px #444444;
border-radius: 3px;
cursor: pointer;
z-index: 400;
transition: border-color 0.4s;
}
.version-btn::after {
content:"\f0d8";
display: inline;
font: normal normal normal 16px/1 FontAwesome;
color: #8d8c8c;
vertical-align: top;
padding-left: 0.5em;
}
.version-btn-open::after {
color: gray;
}
.version-btn:hover, .version-btn:focus {
border-color: #525252;
}
.version-btn-open {
color: gray;
border: solid 1px gray;
}
.version-btn.wait {
cursor: wait;
}
.version-btn.disabled {
cursor: not-allowed;
color: dimgray;
}
.version-dialog {
display: none;
position: absolute;
bottom: 28px;
width: 140px;
margin: 0 5px;
padding-bottom: 4px;
background-color: #0003;
border-radius: 3px;
box-shadow: 0 0 6px #000C;
z-index: 999;
max-height: calc(100vh - 30px);
overflow-y: auto;
cursor: default;
}
.version-title {
padding: 5px;
color: black;
text-align: center;
font-size: 102%;
background-color: #27ae60;
border-bottom: solid 1.5px #444;
}
.version-list {
margin-bottom: 4px;
text-align: center;
background-color: #000C;
border: solid 1px gray;
border-radius: 0px 0px 3px 3px;
}
.version-list a, .version-list span, .version-list li {
position: relative;
display: block;
font-size: 98%;
line-height: 1.15;
width: 100%;
margin: 0;
padding: 4px 0px;
color: #404040;
}
.version-list li {
background-color: #ede9e9;
color: #404040;
padding: 1px;
}
.version-list li:hover, .version-list li a:focus {
background-color: #b9cfda;
}
.version-list li.selected, .version-list li.selected:hover {
background-color: #8d8c8c;
}
.version-list li.selected span {
cursor: default;
outline-color: red;
}
.version-arrow {
position: absolute;
width: 8px;
height: 8px;
left: 50%;
bottom: 4px;
margin-left: -4px;
transform: rotate(225deg);
background: #ede9e9;
border: 1px solid gray;
border-width: 1px 0 0 1px;
}

View File

@@ -1,323 +0,0 @@
(function() { // switch: v1.2
"use strict";
var versionsFileUrl = "https://docs.blender.org/PROD/versions.json"
var all_versions;
var Popover = function() {
function Popover(id)
{
this.isOpen = false;
this.type = (id === "version-popover");
this.$btn = $('#' + id);
this.$dialog = this.$btn.next();
this.$list = this.$dialog.children("ul");
this.sel = null;
this.beforeInit();
}
Popover.prototype = {
beforeInit : function() {
var that = this;
this.$btn.on("click", function(e) {
that.init();
e.preventDefault();
e.stopPropagation();
});
this.$btn.on("keydown", function(e) {
if (that.btnKeyFilter(e)) {
that.init();
e.preventDefault();
e.stopPropagation();
}
});
},
init : function() {
this.$btn.off("click");
this.$btn.off("keydown");
if (all_versions === undefined) {
this.$btn.addClass("wait");
this.loadVL(this);
}
else {
this.afterLoad();
}
},
loadVL : function(that) {
$.getJSON(versionsFileUrl, function(data) {
all_versions = data;
that.afterLoad();
return true;
}).fail(function() {
console.log("Version Switch Error: versions.json could not be loaded.");
that.$btn.addClass("disabled");
return false;
});
},
afterLoad : function() {
var release = DOCUMENTATION_OPTIONS.VERSION;
const m = release.match(/\d\.\d+/g);
if (m) {
release = m[0];
}
this.warnOld(release, all_versions);
var version = this.getNamed(release);
var list = this.buildList(version);
this.$list.children(":first-child").remove();
this.$list.append(list);
var that = this;
this.$list.on("keydown", function(e) {
that.keyMove(e);
});
this.$btn.removeClass("wait");
this.btnOpenHandler();
this.$btn.on("mousedown", function(e) {
that.btnOpenHandler();
e.preventDefault()
});
this.$btn.on("keydown", function(e) {
if (that.btnKeyFilter(e)) {
that.btnOpenHandler();
}
});
},
warnOld : function(release, all_versions) {
// Note this is effectively disabled now, two issues must fixed:
// * versions.js does not contain a current entry, because that leads to
// duplicate version numbers in the menu. These need to be deduplicated.
// * It only shows the warning after opening the menu to switch version
// when versions.js is loaded. This is too late to be useful.
var current = all_versions.current
if (!current)
{
// console.log("Version Switch Error: no 'current' in version.json.");
return;
}
const m = current.match(/\d\.\d+/g);
if (m) {
current = parseFloat(m[0]);
}
if (release < current) {
var currentURL = window.location.pathname.replace(release, current);
var warning = $('<div class="admonition warning"> ' +
'<p class="first admonition-title">Note</p> ' +
'<p class="last"> ' +
'You are not using the most up to date version of the documentation. ' +
'<a href="#"></a> is the newest version.' +
'</p>' +
'</div>');
warning.find('a').attr('href', currentURL).text(current);
var body = $("div.body");
if (!body.length) {
body = $("div.document");
}
body.prepend(warning);
}
},
buildList : function(v) {
var url = new URL(window.location.href);
let pathSplit = [ "", "api", v ];
if (url.pathname.startsWith("/api/")) {
pathSplit.push(url.pathname.split('/').slice(3).join('/'));
}
else {
pathSplit.push(url.pathname.substring(1));
}
if (this.type) {
var dyn = all_versions;
var cur = v;
}
var buf = [];
var that = this;
$.each(dyn, function(ix, title) {
buf.push("<li");
if (ix === cur) {
buf.push(
' class="selected" tabindex="-1" role="presentation"><span tabindex="-1" role="menuitem" aria-current="page">' +
title + '</spanp></li>');
}
else {
pathSplit[2 + that.type] = ix;
var href = new URL(url);
href.pathname = pathSplit.join('/');
buf.push(' tabindex="-1" role="presentation"><a href ="' + href + '" tabindex="-1">' +
title + '</a></li>');
}
});
return buf.join('');
},
getNamed : function(v) {
$.each(all_versions, function(ix, title) {
if (ix === "master" || ix === "latest") {
var m = title.match(/\d\.\d[\w\d\.]*/)[0];
if (parseFloat(m) == v) {
v = ix;
return false;
}
}
});
return v;
},
dialogToggle : function(speed) {
var wasClose = !this.isOpen;
var that = this;
if (!this.isOpen) {
this.$btn.addClass("version-btn-open");
this.$btn.attr("aria-pressed", true);
this.$dialog.attr("aria-hidden", false);
this.$dialog.fadeIn(speed, function() {
that.$btn.parent().on("focusout", function(e) {
that.focusoutHandler();
e.stopImmediatePropagation();
})
that.$btn.parent().on("mouseleave", function(e) {
that.mouseoutHandler();
e.stopImmediatePropagation();
});
});
this.isOpen = true;
}
else {
this.$btn.removeClass("version-btn-open");
this.$btn.attr("aria-pressed", false);
this.$dialog.attr("aria-hidden", true);
this.$btn.parent().off("focusout");
this.$btn.parent().off("mouseleave");
this.$dialog.fadeOut(speed, function() {
if (this.$sel) {
this.$sel.attr("tabindex", -1);
}
that.$btn.attr("tabindex", 0);
if (document.activeElement !== null && document.activeElement !== document &&
document.activeElement !== document.body) {
that.$btn.focus();
}
});
this.isOpen = false;
}
if (wasClose) {
if (this.$sel) {
this.$sel.attr("tabindex", -1);
}
if (document.activeElement !== null && document.activeElement !== document &&
document.activeElement !== document.body) {
var $nw = this.listEnter();
$nw.attr("tabindex", 0);
$nw.focus();
this.$sel = $nw;
}
}
},
btnOpenHandler : function() {
this.dialogToggle(300);
},
focusoutHandler : function() {
var list = this.$list;
var that = this;
setTimeout(function() {
if (list.find(":focus").length === 0) {
that.dialogToggle(200);
}
}, 200);
},
mouseoutHandler : function() {
this.dialogToggle(200);
},
btnKeyFilter : function(e) {
if (e.ctrlKey || e.shiftKey) {
return false;
}
if (e.key === " " || e.key === "Enter" || (e.key === "ArrowDown" && e.altKey) ||
e.key === "ArrowDown" || e.key === "ArrowUp") {
return true;
}
return false;
},
keyMove : function(e) {
if (e.ctrlKey || e.shiftKey) {
return true;
}
var p = true;
var $nw = $(e.target);
switch (e.key) {
case "ArrowUp":
$nw = this.listPrev($nw);
break;
case "ArrowDown":
$nw = this.listNext($nw);
break;
case "Home":
$nw = this.listFirst();
break;
case "End":
$nw = this.listLast();
break;
case "Escape":
$nw = this.listExit();
break;
case "ArrowLeft":
$nw = this.listExit();
break;
case "ArrowRight":
$nw = this.listExit();
break;
default:
p = false;
}
if (p) {
$nw.attr("tabindex", 0);
$nw.focus();
if (this.$sel) {
this.$sel.attr("tabindex", -1);
}
this.$sel = $nw;
e.preventDefault();
e.stopPropagation();
}
},
listPrev : function($nw) {
if ($nw.parent().prev().length !== 0) {
return $nw.parent().prev().children(":first-child");
}
else {
return this.listLast();
}
},
listNext : function($nw) {
if ($nw.parent().next().length !== 0) {
return $nw.parent().next().children(":first-child");
}
else {
return this.listFirst();
}
},
listFirst : function() {
return this.$list.children(":first-child").children(":first-child");
},
listLast : function() {
return this.$list.children(":last-child").children(":first-child");
},
listExit : function() {
this.mouseoutHandler();
return this.$btn;
},
listEnter : function() {
return this.$list.children(":first-child").children(":first-child");
}
};
return Popover
}();
$(document).ready(function() {
var lng_popover = new Popover("version-popover");
});
})();

View File

@@ -1,17 +0,0 @@
<div class="rst-versions" data-toggle="rst-versions" role="note" aria-label="document versions">
<ul id="versionwrap" role="presentation">
<li role="presentation">
<span id="version-popover" class="version-btn" tabindex="0" role="button" aria-label="versions selector" aria-haspopup="true" aria-controls="version-vsnlist" aria-disabled="true">
{{ release }}
</span>
<div class="version-dialog" aria-hidden="true">
<div class="version-arrow" aria-hidden="true"></div>
<div class="version-title">Versions</div>
<ul id="version-vsnlist" class="version-list" role="menu" aria-labelledby="version-popover" aria-hidden="true">
<li role="presentation">Loading...</li>
</ul>
</div>
</li>
</ul>
</div>

View File

@@ -113,6 +113,6 @@ if(WITH_MOD_FLUID)
add_subdirectory(mantaflow)
endif()
if(WITH_COMPOSITOR)
if (WITH_COMPOSITOR)
add_subdirectory(smaa_areatex)
endif()

View File

@@ -74,7 +74,7 @@ enum_panorama_types = (
"Similar to most fisheye modern lens, takes sensor dimensions into consideration"),
('MIRRORBALL', "Mirror Ball", "Uses the mirror ball mapping"),
('FISHEYE_LENS_POLYNOMIAL', "Fisheye Lens Polynomial",
"Defines the lens projection as polynomial to allow real world camera lenses to be mimicked"),
"Defines the lens projection as polynomial to allow real world camera lenses to be mimicked."),
)
enum_curve_shape = (
@@ -901,27 +901,27 @@ class CyclesCameraSettings(bpy.types.PropertyGroup):
fisheye_polynomial_k0: FloatProperty(
name="Fisheye Polynomial K0",
description="Coefficient K0 of the lens polynomial",
description="Coefficient K0 of the lens polinomial",
default=camera.default_fisheye_polynomial[0], precision=6, step=0.1, subtype='ANGLE',
)
fisheye_polynomial_k1: FloatProperty(
name="Fisheye Polynomial K1",
description="Coefficient K1 of the lens polynomial",
description="Coefficient K1 of the lens polinomial",
default=camera.default_fisheye_polynomial[1], precision=6, step=0.1, subtype='ANGLE',
)
fisheye_polynomial_k2: FloatProperty(
name="Fisheye Polynomial K2",
description="Coefficient K2 of the lens polynomial",
description="Coefficient K2 of the lens polinomial",
default=camera.default_fisheye_polynomial[2], precision=6, step=0.1, subtype='ANGLE',
)
fisheye_polynomial_k3: FloatProperty(
name="Fisheye Polynomial K3",
description="Coefficient K3 of the lens polynomial",
description="Coefficient K3 of the lens polinomial",
default=camera.default_fisheye_polynomial[3], precision=6, step=0.1, subtype='ANGLE',
)
fisheye_polynomial_k4: FloatProperty(
name="Fisheye Polynomial K4",
description="Coefficient K4 of the lens polynomial",
description="Coefficient K4 of the lens polinomial",
default=camera.default_fisheye_polynomial[4], precision=6, step=0.1, subtype='ANGLE',
)

View File

@@ -669,7 +669,7 @@ class CYCLES_RENDER_PT_performance_acceleration_structure(CyclesButtonsPanel, Pa
@classmethod
def poll(cls, context):
return not use_optix(context) or use_multi_device(context)
return not use_optix(context) or has_multi_device(context)
def draw(self, context):
import _cycles

View File

@@ -14,8 +14,6 @@
* limitations under the License.
*/
#include <optional>
#include "blender/sync.h"
#include "blender/util.h"
@@ -626,36 +624,15 @@ void BlenderSync::sync_particle_hair(
}
}
#ifdef WITH_NEW_CURVES_TYPE
static std::optional<BL::FloatAttribute> find_curves_radius_attribute(BL::Curves b_curves)
#ifdef WITH_HAIR_NODES
static float4 hair_point_as_float4(BL::HairPoint b_point)
{
for (BL::Attribute &b_attribute : b_curves.attributes) {
if (b_attribute.name() != "radius") {
continue;
}
if (b_attribute.domain() != BL::Attribute::domain_POINT) {
continue;
}
if (b_attribute.data_type() != BL::Attribute::data_type_FLOAT) {
continue;
}
return BL::FloatAttribute{b_attribute};
}
return std::nullopt;
}
static float4 hair_point_as_float4(BL::Curves b_curves,
std::optional<BL::FloatAttribute> b_attr_radius,
const int index)
{
float4 mP = float3_to_float4(get_float3(b_curves.position_data[index].vector()));
mP.w = b_attr_radius ? b_attr_radius->data[index].value() : 0.0f;
float4 mP = float3_to_float4(get_float3(b_point.co()));
mP.w = b_point.radius();
return mP;
}
static float4 interpolate_hair_points(BL::Curves b_curves,
std::optional<BL::FloatAttribute> b_attr_radius,
static float4 interpolate_hair_points(BL::Hair b_hair,
const int first_point_index,
const int num_points,
const float step)
@@ -664,12 +641,12 @@ static float4 interpolate_hair_points(BL::Curves b_curves,
const int point_a = clamp((int)curve_t, 0, num_points - 1);
const int point_b = min(point_a + 1, num_points - 1);
const float t = curve_t - (float)point_a;
return lerp(hair_point_as_float4(b_curves, b_attr_radius, first_point_index + point_a),
hair_point_as_float4(b_curves, b_attr_radius, first_point_index + point_b),
return lerp(hair_point_as_float4(b_hair.points[first_point_index + point_a]),
hair_point_as_float4(b_hair.points[first_point_index + point_b]),
t);
}
static void export_hair_curves(Scene *scene, Hair *hair, BL::Curves b_curves)
static void export_hair_curves(Scene *scene, Hair *hair, BL::Hair b_hair)
{
/* TODO: optimize so we can straight memcpy arrays from Blender? */
@@ -689,19 +666,17 @@ static void export_hair_curves(Scene *scene, Hair *hair, BL::Curves b_curves)
}
/* Reserve memory. */
const int num_keys = b_curves.points.length();
const int num_curves = b_curves.curves.length();
const int num_keys = b_hair.points.length();
const int num_curves = b_hair.curves.length();
hair->reserve_curves(num_curves, num_keys);
std::optional<BL::FloatAttribute> b_attr_radius = find_curves_radius_attribute(b_curves);
/* Export curves and points. */
vector<float> points_length;
for (int i = 0; i < num_curves; i++) {
const int first_point_index = b_curves.curve_offset_data[i].value();
const int num_points = b_curves.curve_offset_data[i + 1].value() - first_point_index;
for (BL::HairCurve &b_curve : b_hair.curves) {
const int first_point_index = b_curve.first_point_index();
const int num_points = b_curve.num_points();
float3 prev_co = zero_float3();
float length = 0.0f;
@@ -712,9 +687,10 @@ static void export_hair_curves(Scene *scene, Hair *hair, BL::Curves b_curves)
/* Position and radius. */
for (int i = 0; i < num_points; i++) {
const float3 co = get_float3(b_curves.position_data[first_point_index + i].vector());
const float radius = b_attr_radius ? b_attr_radius->data[first_point_index + i].value() :
0.0f;
BL::HairPoint b_point = b_hair.points[first_point_index + i];
const float3 co = get_float3(b_point.co());
const float radius = b_point.radius();
hair->add_curve_key(co, radius);
if (attr_intercept) {
@@ -739,7 +715,7 @@ static void export_hair_curves(Scene *scene, Hair *hair, BL::Curves b_curves)
/* Random number per curve. */
if (attr_random != NULL) {
attr_random->add(hash_uint2_to_float(i, 0));
attr_random->add(hash_uint2_to_float(b_curve.index(), 0));
}
/* Curve. */
@@ -748,7 +724,7 @@ static void export_hair_curves(Scene *scene, Hair *hair, BL::Curves b_curves)
}
}
static void export_hair_curves_motion(Hair *hair, BL::Curves b_curves, int motion_step)
static void export_hair_curves_motion(Hair *hair, BL::Hair b_hair, int motion_step)
{
/* Find or add attribute. */
Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
@@ -761,17 +737,14 @@ static void export_hair_curves_motion(Hair *hair, BL::Curves b_curves, int motio
/* Export motion keys. */
const int num_keys = hair->get_curve_keys().size();
const int num_curves = b_curves.curves.length();
float4 *mP = attr_mP->data_float4() + motion_step * num_keys;
bool have_motion = false;
int num_motion_keys = 0;
int curve_index = 0;
std::optional<BL::FloatAttribute> b_attr_radius = find_curves_radius_attribute(b_curves);
for (int i = 0; i < num_curves; i++) {
const int first_point_index = b_curves.curve_offset_data[i].value();
const int num_points = b_curves.curve_offset_data[i + 1].value() - first_point_index;
for (BL::HairCurve &b_curve : b_hair.curves) {
const int first_point_index = b_curve.first_point_index();
const int num_points = b_curve.num_points();
Hair::Curve curve = hair->get_curve(curve_index);
curve_index++;
@@ -782,7 +755,7 @@ static void export_hair_curves_motion(Hair *hair, BL::Curves b_curves, int motio
int point_index = first_point_index + i;
if (point_index < num_keys) {
mP[num_motion_keys] = hair_point_as_float4(b_curves, b_attr_radius, point_index);
mP[num_motion_keys] = hair_point_as_float4(b_hair.points[point_index]);
num_motion_keys++;
if (!have_motion) {
@@ -801,8 +774,7 @@ static void export_hair_curves_motion(Hair *hair, BL::Curves b_curves, int motio
const float step_size = curve.num_keys > 1 ? 1.0f / (curve.num_keys - 1) : 0.0f;
for (int i = 0; i < curve.num_keys; i++) {
const float step = i * step_size;
mP[num_motion_keys] = interpolate_hair_points(
b_curves, b_attr_radius, first_point_index, num_points, step);
mP[num_motion_keys] = interpolate_hair_points(b_hair, first_point_index, num_points, step);
num_motion_keys++;
}
have_motion = true;
@@ -819,12 +791,12 @@ static void export_hair_curves_motion(Hair *hair, BL::Curves b_curves, int motio
void BlenderSync::sync_hair(Hair *hair, BObjectInfo &b_ob_info, bool motion, int motion_step)
{
/* Convert Blender hair to Cycles curves. */
BL::Curves b_curves(b_ob_info.object_data);
BL::Hair b_hair(b_ob_info.object_data);
if (motion) {
export_hair_curves_motion(hair, b_curves, motion_step);
export_hair_curves_motion(hair, b_hair, motion_step);
}
else {
export_hair_curves(scene, hair, b_curves);
export_hair_curves(scene, hair, b_hair);
}
}
#else
@@ -847,8 +819,8 @@ void BlenderSync::sync_hair(BL::Depsgraph b_depsgraph, BObjectInfo &b_ob_info, H
new_hair.set_used_shaders(used_shaders);
if (view_layer.use_hair) {
#ifdef WITH_NEW_CURVES_TYPE
if (b_ob_info.object_data.is_a(&RNA_Curves)) {
#ifdef WITH_HAIR_NODES
if (b_ob_info.object_data.is_a(&RNA_Hair)) {
/* Hair object. */
sync_hair(&new_hair, b_ob_info, false);
}
@@ -901,8 +873,8 @@ void BlenderSync::sync_hair_motion(BL::Depsgraph b_depsgraph,
/* Export deformed coordinates. */
if (ccl::BKE_object_is_deform_modified(b_ob_info, b_scene, preview)) {
#ifdef WITH_NEW_CURVES_TYPE
if (b_ob_info.object_data.is_a(&RNA_Curves)) {
#ifdef WITH_HAIR_NODES
if (b_ob_info.object_data.is_a(&RNA_Hair)) {
/* Hair object. */
sync_hair(hair, b_ob_info, true, motion_step);
return;

View File

@@ -32,8 +32,8 @@ CCL_NAMESPACE_BEGIN
static Geometry::Type determine_geom_type(BObjectInfo &b_ob_info, bool use_particle_hair)
{
#ifdef WITH_NEW_CURVES_TYPE
if (b_ob_info.object_data.is_a(&RNA_Curves) || use_particle_hair) {
#ifdef WITH_HAIR_NODES
if (b_ob_info.object_data.is_a(&RNA_Hair) || use_particle_hair) {
#else
if (use_particle_hair) {
#endif
@@ -231,8 +231,8 @@ void BlenderSync::sync_geometry_motion(BL::Depsgraph &b_depsgraph,
if (progress.get_cancel())
return;
#ifdef WITH_NEW_CURVES_TYPE
if (b_ob_info.object_data.is_a(&RNA_Curves) || use_particle_hair) {
#ifdef WITH_HAIR_NODES
if (b_ob_info.object_data.is_a(&RNA_Hair) || use_particle_hair) {
#else
if (use_particle_hair) {
#endif

View File

@@ -72,7 +72,7 @@ bool BlenderSync::object_is_geometry(BObjectInfo &b_ob_info)
BL::Object::type_enum type = b_ob_info.iter_object.type();
if (type == BL::Object::type_VOLUME || type == BL::Object::type_CURVES ||
if (type == BL::Object::type_VOLUME || type == BL::Object::type_HAIR ||
type == BL::Object::type_POINTCLOUD) {
/* Will be exported attached to mesh. */
return true;
@@ -97,7 +97,7 @@ bool BlenderSync::object_can_have_geometry(BL::Object &b_ob)
case BL::Object::type_SURFACE:
case BL::Object::type_META:
case BL::Object::type_FONT:
case BL::Object::type_CURVES:
case BL::Object::type_HAIR:
case BL::Object::type_POINTCLOUD:
case BL::Object::type_VOLUME:
return true;

View File

@@ -61,26 +61,6 @@ static_assert(Object::MAX_MOTION_STEPS == Geometry::MAX_MOTION_STEPS,
# define IS_HAIR(x) (x & 1)
/* This gets called by Embree at every valid ray/object intersection.
* Things like recording subsurface or shadow hits for later evaluation
* as well as filtering for volume objects happen here.
* Cycles' own BVH does that directly inside the traversal calls.
*/
static void rtc_filter_intersection_func(const RTCFilterFunctionNArguments *args)
{
/* Current implementation in Cycles assumes only single-ray intersection queries. */
assert(args->N == 1);
RTCHit *hit = (RTCHit *)args->hit;
CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt;
const KernelGlobalsCPU *kg = ctx->kg;
const Ray *cray = ctx->ray;
if (kernel_embree_is_self_intersection(kg, hit, cray)) {
*args->valid = 0;
}
}
/* This gets called by Embree at every valid ray/object intersection.
* Things like recording subsurface or shadow hits for later evaluation
* as well as filtering for volume objects happen here.
@@ -95,16 +75,12 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
RTCHit *hit = (RTCHit *)args->hit;
CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt;
const KernelGlobalsCPU *kg = ctx->kg;
const Ray *cray = ctx->ray;
switch (ctx->type) {
case CCLIntersectContext::RAY_SHADOW_ALL: {
Intersection current_isect;
kernel_embree_convert_hit(kg, ray, hit, &current_isect);
if (intersection_skip_self_shadow(cray->self, current_isect.object, current_isect.prim)) {
*args->valid = 0;
return;
}
/* If no transparent shadows or max number of hits exceeded, all light is blocked. */
const int flags = intersection_get_shader_flags(kg, current_isect.prim, current_isect.type);
if (!(flags & (SD_HAS_TRANSPARENT_SHADOW)) || ctx->num_hits >= ctx->max_hits) {
@@ -184,10 +160,6 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
break;
}
}
if (intersection_skip_self_local(cray->self, current_isect.prim)) {
*args->valid = 0;
return;
}
/* No intersection information requested, just return a hit. */
if (ctx->max_hits == 0) {
@@ -253,11 +225,6 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
if (ctx->num_hits < ctx->max_hits) {
Intersection current_isect;
kernel_embree_convert_hit(kg, ray, hit, &current_isect);
if (intersection_skip_self(cray->self, current_isect.object, current_isect.prim)) {
*args->valid = 0;
return;
}
Intersection *isect = &ctx->isect_s[ctx->num_hits];
++ctx->num_hits;
*isect = current_isect;
@@ -269,15 +236,12 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
}
/* This tells Embree to continue tracing. */
*args->valid = 0;
break;
}
break;
}
case CCLIntersectContext::RAY_REGULAR:
default:
if (kernel_embree_is_self_intersection(kg, hit, cray)) {
*args->valid = 0;
return;
}
/* Nothing to do here. */
break;
}
}
@@ -293,14 +257,6 @@ static void rtc_filter_func_backface_cull(const RTCFilterFunctionNArguments *arg
*args->valid = 0;
return;
}
CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt;
const KernelGlobalsCPU *kg = ctx->kg;
const Ray *cray = ctx->ray;
if (kernel_embree_is_self_intersection(kg, hit, cray)) {
*args->valid = 0;
}
}
static void rtc_filter_occluded_func_backface_cull(const RTCFilterFunctionNArguments *args)
@@ -549,7 +505,6 @@ void BVHEmbree::add_triangles(const Object *ob, const Mesh *mesh, int i)
rtcSetGeometryUserData(geom_id, (void *)prim_offset);
rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func);
rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_intersection_func);
rtcSetGeometryMask(geom_id, ob->visibility_for_tracing());
rtcCommitGeometry(geom_id);
@@ -812,7 +767,6 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
rtcSetGeometryUserData(geom_id, (void *)prim_offset);
if (hair->curve_shape == CURVE_RIBBON) {
rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_intersection_func);
rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func);
}
else {

View File

@@ -142,7 +142,6 @@ class BVHParams {
top_level = false;
bvh_layout = BVH_LAYOUT_BVH2;
use_compact_structure = true;
use_unaligned_nodes = false;
num_motion_curve_steps = 0;

View File

@@ -559,10 +559,10 @@ if(WITH_CYCLES_DEVICE_METAL)
find_library(METAL_LIBRARY Metal)
# This file was added in the 12.0 SDK, use it as a way to detect the version.
if(METAL_LIBRARY AND NOT EXISTS "${METAL_LIBRARY}/Headers/MTLFunctionStitching.h")
if (METAL_LIBRARY AND NOT EXISTS "${METAL_LIBRARY}/Headers/MTLFunctionStitching.h")
message(STATUS "Metal version too old, must be SDK 12.0 or newer, disabling WITH_CYCLES_DEVICE_METAL")
set(WITH_CYCLES_DEVICE_METAL OFF)
elseif(NOT METAL_LIBRARY)
elseif (NOT METAL_LIBRARY)
message(STATUS "Metal not found, disabling WITH_CYCLES_DEVICE_METAL")
set(WITH_CYCLES_DEVICE_METAL OFF)
else()

View File

@@ -905,8 +905,8 @@ void HIPDevice::tex_alloc(device_texture &mem)
address_mode = hipAddressModeClamp;
break;
case EXTENSION_CLIP:
/* TODO(@arya): setting this to Mode Clamp instead of Mode Border
* because it's unsupported in HIP. */
// TODO : (Arya) setting this to Mode Clamp instead of Mode Border because it's unsupported
// in hip
address_mode = hipAddressModeClamp;
break;
default:

View File

@@ -226,7 +226,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
pipeline_options.usesMotionBlur = false;
pipeline_options.traversableGraphFlags =
OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_LEVEL_INSTANCING;
pipeline_options.numPayloadValues = 8;
pipeline_options.numPayloadValues = 6;
pipeline_options.numAttributeValues = 2; /* u, v */
pipeline_options.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE;
pipeline_options.pipelineLaunchParamsVariableName = "__params"; /* See globals.h */

View File

@@ -125,41 +125,20 @@ static Device *find_best_device(Device *device, DenoiserType type)
return best_device;
}
static DeviceInfo find_best_denoiser_device_info(const vector<DeviceInfo> &device_infos,
DenoiserType denoiser_type)
{
for (const DeviceInfo &device_info : device_infos) {
if ((device_info.denoisers & denoiser_type) == 0) {
continue;
}
/* TODO(sergey): Use one of the already configured devices, so that OptiX denoising can happen
* on a physical CUDA device which is already used for rendering. */
/* TODO(sergey): Choose fastest device for denoising. */
return device_info;
}
DeviceInfo none_device;
none_device.type = DEVICE_NONE;
return none_device;
}
static unique_ptr<Device> create_denoiser_device(Device *path_trace_device,
const uint device_type_mask,
DenoiserType denoiser_type)
const uint device_type_mask)
{
const vector<DeviceInfo> device_infos = Device::available_devices(device_type_mask);
if (device_infos.empty()) {
return nullptr;
}
const DeviceInfo denoiser_device_info = find_best_denoiser_device_info(device_infos,
denoiser_type);
if (denoiser_device_info.type == DEVICE_NONE) {
return nullptr;
}
/* TODO(sergey): Use one of the already configured devices, so that OptiX denoising can happen on
* a physical CUDA device which is already used for rendering. */
/* TODO(sergey): Choose fastest device for denoising. */
const DeviceInfo denoiser_device_info = device_infos.front();
unique_ptr<Device> denoiser_device(
Device::create(denoiser_device_info, path_trace_device->stats, path_trace_device->profiler));
@@ -207,8 +186,7 @@ Device *Denoiser::ensure_denoiser_device(Progress *progress)
device_creation_attempted_ = true;
const uint device_type_mask = get_device_type_mask();
local_denoiser_device_ = create_denoiser_device(
path_trace_device_, device_type_mask, params_.type);
local_denoiser_device_ = create_denoiser_device(path_trace_device_, device_type_mask);
denoiser_device_ = local_denoiser_device_.get();
return denoiser_device_;

View File

@@ -37,6 +37,8 @@ OIDNDenoiser::OIDNDenoiser(Device *path_trace_device, const DenoiseParams &param
: Denoiser(path_trace_device, params)
{
DCHECK_EQ(params.type, DENOISER_OPENIMAGEDENOISE);
DCHECK(openimagedenoise_supported()) << "OpenImageDenoiser is not supported on this platform.";
}
#ifdef WITH_OPENIMAGEDENOISE
@@ -583,9 +585,6 @@ bool OIDNDenoiser::denoise_buffer(const BufferParams &buffer_params,
const int num_samples,
bool allow_inplace_modification)
{
DCHECK(openimagedenoise_supported())
<< "OpenImageDenoiser is not supported on this platform or build.";
#ifdef WITH_OPENIMAGEDENOISE
thread_scoped_lock lock(mutex_);
@@ -636,20 +635,4 @@ uint OIDNDenoiser::get_device_type_mask() const
return DEVICE_MASK_CPU;
}
Device *OIDNDenoiser::ensure_denoiser_device(Progress *progress)
{
#ifndef WITH_OPENIMAGEDENOISE
path_trace_device_->set_error("Build without OpenImageDenoiser");
return nullptr;
#else
if (!openimagedenoise_supported()) {
path_trace_device_->set_error(
"OpenImageDenoiser is not supported on this CPU: missing SSE 4.1 support");
return nullptr;
}
return Denoiser::ensure_denoiser_device(progress);
#endif
}
CCL_NAMESPACE_END

View File

@@ -38,7 +38,6 @@ class OIDNDenoiser : public Denoiser {
protected:
virtual uint get_device_type_mask() const override;
virtual Device *ensure_denoiser_device(Progress *progress) override;
/* We only perform one denoising at a time, since OpenImageDenoise itself is multithreaded.
* Use this mutex whenever images are passed to the OIDN and needs to be denoised. */

View File

@@ -157,7 +157,7 @@ bool ShaderEval::eval_gpu(Device *device,
queue->init_execution();
/* Execute work on GPU in chunk, so we can cancel.
* TODO: query appropriate size from device. */
* TODO : query appropriate size from device.*/
const int32_t chunk_size = 65536;
device_ptr d_input = input.device_pointer;

View File

@@ -173,16 +173,15 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
uint p3 = 0;
uint p4 = visibility;
uint p5 = PRIMITIVE_NONE;
uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
uint ray_mask = visibility & 0xFF;
uint ray_flags = OPTIX_RAY_FLAG_ENFORCE_ANYHIT;
uint ray_flags = OPTIX_RAY_FLAG_NONE;
if (0 == ray_mask && (visibility & ~0xFF) != 0) {
ray_mask = 0xFF;
ray_flags = OPTIX_RAY_FLAG_ENFORCE_ANYHIT;
}
else if (visibility & PATH_RAY_SHADOW_OPAQUE) {
ray_flags |= OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT;
ray_flags = OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT;
}
optixTrace(scene_intersect_valid(ray) ? kernel_data.bvh.scene : 0,
@@ -201,9 +200,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
p2,
p3,
p4,
p5,
p6,
p7);
p5);
isect->t = __uint_as_float(p0);
isect->u = __uint_as_float(p1);
@@ -245,7 +242,6 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
}
MetalRTIntersectionPayload payload;
payload.self = ray->self;
payload.u = 0.0f;
payload.v = 0.0f;
payload.visibility = visibility;
@@ -313,7 +309,6 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR);
IntersectContext rtc_ctx(&ctx);
RTCRayHit ray_hit;
ctx.ray = ray;
kernel_embree_setup_rayhit(*ray, ray_hit, visibility);
rtcIntersect1(kernel_data.bvh.scene, &rtc_ctx.context, &ray_hit);
if (ray_hit.hit.geomID != RTC_INVALID_GEOMETRY_ID &&
@@ -361,9 +356,6 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
uint p2 = pointer_pack_to_uint_0(local_isect);
uint p3 = pointer_pack_to_uint_1(local_isect);
uint p4 = local_object;
uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
/* Is set to zero on miss or if ray is aborted, so can be used as return value. */
uint p5 = max_hits;
@@ -387,9 +379,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
p2,
p3,
p4,
p5,
p6,
p7);
p5);
return p5;
# elif defined(__METALRT__)
@@ -427,7 +417,6 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
}
MetalRTIntersectionLocalPayload payload;
payload.self = ray->self;
payload.local_object = local_object;
payload.max_hits = max_hits;
payload.local_isect.num_hits = 0;
@@ -471,7 +460,6 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
kg, has_bvh ? CCLIntersectContext::RAY_SSS : CCLIntersectContext::RAY_LOCAL);
ctx.lcg_state = lcg_state;
ctx.max_hits = max_hits;
ctx.ray = ray;
ctx.local_isect = local_isect;
if (local_isect) {
local_isect->num_hits = 0;
@@ -544,8 +532,6 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
uint p3 = max_hits;
uint p4 = visibility;
uint p5 = false;
uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
uint ray_mask = visibility & 0xFF;
if (0 == ray_mask && (visibility & ~0xFF) != 0) {
@@ -569,9 +555,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
p2,
p3,
p4,
p5,
p6,
p7);
p5);
*num_recorded_hits = uint16_unpack_from_uint_0(p2);
*throughput = __uint_as_float(p1);
@@ -604,7 +588,6 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
}
MetalRTIntersectionShadowPayload payload;
payload.self = ray->self;
payload.visibility = visibility;
payload.max_hits = max_hits;
payload.num_hits = 0;
@@ -651,7 +634,6 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
Intersection *isect_array = (Intersection *)state->shadow_isect;
ctx.isect_s = isect_array;
ctx.max_hits = max_hits;
ctx.ray = ray;
IntersectContext rtc_ctx(&ctx);
RTCRay rtc_ray;
kernel_embree_setup_ray(*ray, rtc_ray, visibility);
@@ -703,8 +685,6 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
uint p3 = 0;
uint p4 = visibility;
uint p5 = PRIMITIVE_NONE;
uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
uint ray_mask = visibility & 0xFF;
if (0 == ray_mask && (visibility & ~0xFF) != 0) {
@@ -728,9 +708,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
p2,
p3,
p4,
p5,
p6,
p7);
p5);
isect->t = __uint_as_float(p0);
isect->u = __uint_as_float(p1);
@@ -766,7 +744,6 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
}
MetalRTIntersectionPayload payload;
payload.self = ray->self;
payload.visibility = visibility;
typename metalrt_intersector_type::result_type intersection;
@@ -843,7 +820,6 @@ ccl_device_intersect uint scene_intersect_volume_all(KernelGlobals kg,
ctx.isect_s = isect;
ctx.max_hits = max_hits;
ctx.num_hits = 0;
ctx.ray = ray;
IntersectContext rtc_ctx(&ctx);
RTCRay rtc_ray;
kernel_embree_setup_ray(*ray, rtc_ray, visibility);

View File

@@ -22,8 +22,6 @@
#include "kernel/device/cpu/compat.h"
#include "kernel/device/cpu/globals.h"
#include "kernel/bvh/util.h"
#include "util/vector.h"
CCL_NAMESPACE_BEGIN
@@ -40,9 +38,6 @@ struct CCLIntersectContext {
KernelGlobals kg;
RayType type;
/* For avoiding self intersections */
const Ray *ray;
/* for shadow rays */
Intersection *isect_s;
uint max_hits;
@@ -61,7 +56,6 @@ struct CCLIntersectContext {
{
kg = kg_;
type = type_;
ray = NULL;
max_hits = 1;
num_hits = 0;
num_recorded_hits = 0;
@@ -108,34 +102,7 @@ ccl_device_inline void kernel_embree_setup_rayhit(const Ray &ray,
{
kernel_embree_setup_ray(ray, rayhit.ray, visibility);
rayhit.hit.geomID = RTC_INVALID_GEOMETRY_ID;
rayhit.hit.instID[0] = RTC_INVALID_GEOMETRY_ID;
}
ccl_device_inline bool kernel_embree_is_self_intersection(const KernelGlobals kg,
const RTCHit *hit,
const Ray *ray)
{
bool status = false;
if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) {
const int oID = hit->instID[0] / 2;
if ((ray->self.object == oID) || (ray->self.light_object == oID)) {
RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData(
rtcGetGeometry(kernel_data.bvh.scene, hit->instID[0]));
const int pID = hit->primID +
(intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID));
status = intersection_skip_self_shadow(ray->self, oID, pID);
}
}
else {
const int oID = hit->geomID / 2;
if ((ray->self.object == oID) || (ray->self.light_object == oID)) {
const int pID = hit->primID + (intptr_t)rtcGetGeometryUserData(
rtcGetGeometry(kernel_data.bvh.scene, hit->geomID));
status = intersection_skip_self_shadow(ray->self, oID, pID);
}
}
return status;
rayhit.hit.primID = RTC_INVALID_GEOMETRY_ID;
}
ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg,

View File

@@ -157,11 +157,7 @@ ccl_device_inline
}
}
/* Skip self intersection. */
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
if (intersection_skip_self_local(ray->self, prim)) {
continue;
}
if (triangle_intersect_local(kg,
local_isect,
@@ -192,11 +188,7 @@ ccl_device_inline
}
}
/* Skip self intersection. */
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
if (intersection_skip_self_local(ray->self, prim)) {
continue;
}
if (motion_triangle_intersect_local(kg,
local_isect,

View File

@@ -15,7 +15,6 @@
*/
struct MetalRTIntersectionPayload {
RaySelfPrimitives self;
uint visibility;
float u, v;
int prim;
@@ -26,7 +25,6 @@ struct MetalRTIntersectionPayload {
};
struct MetalRTIntersectionLocalPayload {
RaySelfPrimitives self;
uint local_object;
uint lcg_state;
short max_hits;
@@ -36,7 +34,6 @@ struct MetalRTIntersectionLocalPayload {
};
struct MetalRTIntersectionShadowPayload {
RaySelfPrimitives self;
uint visibility;
#if defined(__METALRT_MOTION__)
float time;

View File

@@ -160,9 +160,6 @@ ccl_device_inline
kernel_tex_fetch(__prim_object, prim_addr) :
object;
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
if (intersection_skip_self_shadow(ray->self, prim_object, prim)) {
continue;
}
switch (type & PRIMITIVE_ALL) {
case PRIMITIVE_TRIANGLE: {

View File

@@ -133,29 +133,35 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
--stack_ptr;
/* primitive intersection */
for (; prim_addr < prim_addr2; prim_addr++) {
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
switch (type & PRIMITIVE_ALL) {
case PRIMITIVE_TRIANGLE: {
for (; prim_addr < prim_addr2; prim_addr++) {
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
const int prim_object = (object == OBJECT_NONE) ?
kernel_tex_fetch(__prim_object, prim_addr) :
object;
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
if (intersection_skip_self_shadow(ray->self, prim_object, prim)) {
continue;
}
const int prim_object = (object == OBJECT_NONE) ?
kernel_tex_fetch(__prim_object, prim_addr) :
object;
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
switch (type & PRIMITIVE_ALL) {
case PRIMITIVE_TRIANGLE: {
if (triangle_intersect(
kg, isect, P, dir, isect->t, visibility, prim_object, prim, prim_addr)) {
/* shadow ray early termination */
if (visibility & PATH_RAY_SHADOW_OPAQUE)
return true;
}
break;
}
break;
}
#if BVH_FEATURE(BVH_MOTION)
case PRIMITIVE_MOTION_TRIANGLE: {
case PRIMITIVE_MOTION_TRIANGLE: {
for (; prim_addr < prim_addr2; prim_addr++) {
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
const int prim_object = (object == OBJECT_NONE) ?
kernel_tex_fetch(__prim_object, prim_addr) :
object;
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
if (motion_triangle_intersect(kg,
isect,
P,
@@ -170,21 +176,28 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
if (visibility & PATH_RAY_SHADOW_OPAQUE)
return true;
}
break;
}
break;
}
#endif /* BVH_FEATURE(BVH_MOTION) */
#if BVH_FEATURE(BVH_HAIR)
case PRIMITIVE_CURVE_THICK:
case PRIMITIVE_MOTION_CURVE_THICK:
case PRIMITIVE_CURVE_RIBBON:
case PRIMITIVE_MOTION_CURVE_RIBBON: {
case PRIMITIVE_CURVE_THICK:
case PRIMITIVE_MOTION_CURVE_THICK:
case PRIMITIVE_CURVE_RIBBON:
case PRIMITIVE_MOTION_CURVE_RIBBON: {
for (; prim_addr < prim_addr2; prim_addr++) {
if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
if (ray->time < prim_time.x || ray->time > prim_time.y) {
break;
continue;
}
}
const int prim_object = (object == OBJECT_NONE) ?
kernel_tex_fetch(__prim_object, prim_addr) :
object;
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
const bool hit = curve_intersect(
kg, isect, P, dir, isect->t, prim_object, prim, ray->time, curve_type);
@@ -193,19 +206,26 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
if (visibility & PATH_RAY_SHADOW_OPAQUE)
return true;
}
break;
}
break;
}
#endif /* BVH_FEATURE(BVH_HAIR) */
#if BVH_FEATURE(BVH_POINTCLOUD)
case PRIMITIVE_POINT:
case PRIMITIVE_MOTION_POINT: {
case PRIMITIVE_POINT:
case PRIMITIVE_MOTION_POINT: {
for (; prim_addr < prim_addr2; prim_addr++) {
if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
if (ray->time < prim_time.x || ray->time > prim_time.y) {
break;
continue;
}
}
const int prim_object = (object == OBJECT_NONE) ?
kernel_tex_fetch(__prim_object, prim_addr) :
object;
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
const int point_type = kernel_tex_fetch(__prim_type, prim_addr);
const bool hit = point_intersect(
kg, isect, P, dir, isect->t, prim_object, prim, ray->time, point_type);
@@ -214,10 +234,10 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
if (visibility & PATH_RAY_SHADOW_OPAQUE)
return true;
}
break;
}
#endif /* BVH_FEATURE(BVH_POINTCLOUD) */
break;
}
#endif /* BVH_FEATURE(BVH_POINTCLOUD) */
}
}
else {

View File

@@ -21,22 +21,54 @@ CCL_NAMESPACE_BEGIN
/* Ray offset to avoid self intersection.
*
* This function should be used to compute a modified ray start position for
* rays leaving from a surface. This is from "A Fast and Robust Method for Avoiding
* Self-Intersection" see https://research.nvidia.com/publication/2019-03_A-Fast-and
*/
* rays leaving from a surface. */
ccl_device_inline float3 ray_offset(float3 P, float3 Ng)
{
const float int_scale = 256.0f;
int3 of_i = make_int3((int)(int_scale * Ng.x), (int)(int_scale * Ng.y), (int)(int_scale * Ng.z));
#ifdef __INTERSECTION_REFINE__
const float epsilon_f = 1e-5f;
/* ideally this should match epsilon_f, but instancing and motion blur
* precision makes it problematic */
const float epsilon_test = 1.0f;
const int epsilon_i = 32;
float3 p_i = make_float3(__int_as_float(__float_as_int(P.x) + ((P.x < 0) ? -of_i.x : of_i.x)),
__int_as_float(__float_as_int(P.y) + ((P.y < 0) ? -of_i.y : of_i.y)),
__int_as_float(__float_as_int(P.z) + ((P.z < 0) ? -of_i.z : of_i.z)));
const float origin = 1.0f / 32.0f;
const float float_scale = 1.0f / 65536.0f;
return make_float3(fabsf(P.x) < origin ? P.x + float_scale * Ng.x : p_i.x,
fabsf(P.y) < origin ? P.y + float_scale * Ng.y : p_i.y,
fabsf(P.z) < origin ? P.z + float_scale * Ng.z : p_i.z);
float3 res;
/* x component */
if (fabsf(P.x) < epsilon_test) {
res.x = P.x + Ng.x * epsilon_f;
}
else {
uint ix = __float_as_uint(P.x);
ix += ((ix ^ __float_as_uint(Ng.x)) >> 31) ? -epsilon_i : epsilon_i;
res.x = __uint_as_float(ix);
}
/* y component */
if (fabsf(P.y) < epsilon_test) {
res.y = P.y + Ng.y * epsilon_f;
}
else {
uint iy = __float_as_uint(P.y);
iy += ((iy ^ __float_as_uint(Ng.y)) >> 31) ? -epsilon_i : epsilon_i;
res.y = __uint_as_float(iy);
}
/* z component */
if (fabsf(P.z) < epsilon_test) {
res.z = P.z + Ng.z * epsilon_f;
}
else {
uint iz = __float_as_uint(P.z);
iz += ((iz ^ __float_as_uint(Ng.z)) >> 31) ? -epsilon_i : epsilon_i;
res.z = __uint_as_float(iz);
}
return res;
#else
const float epsilon_f = 1e-4f;
return P + epsilon_f * Ng;
#endif
}
#if defined(__KERNEL_CPU__)
@@ -195,25 +227,4 @@ ccl_device_inline float intersection_curve_shadow_transparency(KernelGlobals kg,
return (1.0f - u) * f0 + u * f1;
}
ccl_device_inline bool intersection_skip_self(ccl_private const RaySelfPrimitives &self,
const int object,
const int prim)
{
return (self.prim == prim) && (self.object == object);
}
ccl_device_inline bool intersection_skip_self_shadow(ccl_private const RaySelfPrimitives &self,
const int object,
const int prim)
{
return ((self.prim == prim) && (self.object == object)) ||
((self.light_prim == prim) && (self.light_object == object));
}
ccl_device_inline bool intersection_skip_self_local(ccl_private const RaySelfPrimitives &self,
const int prim)
{
return (self.prim == prim);
}
CCL_NAMESPACE_END

View File

@@ -144,9 +144,6 @@ ccl_device_inline
kernel_tex_fetch(__prim_object, prim_addr) :
object;
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
if (intersection_skip_self(ray->self, prim_object, prim)) {
continue;
}
int object_flag = kernel_tex_fetch(__object_flag, prim_object);
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
@@ -167,9 +164,6 @@ ccl_device_inline
kernel_tex_fetch(__prim_object, prim_addr) :
object;
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
if (intersection_skip_self(ray->self, prim_object, prim)) {
continue;
}
int object_flag = kernel_tex_fetch(__object_flag, prim_object);
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
continue;

View File

@@ -147,9 +147,6 @@ ccl_device_inline
kernel_tex_fetch(__prim_object, prim_addr) :
object;
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
if (intersection_skip_self(ray->self, prim_object, prim)) {
continue;
}
int object_flag = kernel_tex_fetch(__object_flag, prim_object);
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
continue;
@@ -191,9 +188,6 @@ ccl_device_inline
kernel_tex_fetch(__prim_object, prim_addr) :
object;
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
if (intersection_skip_self(ray->self, prim_object, prim)) {
continue;
}
int object_flag = kernel_tex_fetch(__object_flag, prim_object);
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
continue;

View File

@@ -40,27 +40,6 @@ struct TriangleIntersectionResult
enum { METALRT_HIT_TRIANGLE, METALRT_HIT_BOUNDING_BOX };
ccl_device_inline bool intersection_skip_self(ray_data const RaySelfPrimitives& self,
const int object,
const int prim)
{
return (self.prim == prim) && (self.object == object);
}
ccl_device_inline bool intersection_skip_self_shadow(ray_data const RaySelfPrimitives& self,
const int object,
const int prim)
{
return ((self.prim == prim) && (self.object == object)) ||
((self.light_prim == prim) && (self.light_object == object));
}
ccl_device_inline bool intersection_skip_self_local(ray_data const RaySelfPrimitives& self,
const int prim)
{
return (self.prim == prim);
}
template<typename TReturn, uint intersection_type>
TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload,
@@ -74,8 +53,8 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
#ifdef __BVH_LOCAL__
uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
if ((object != payload.local_object) || intersection_skip_self_local(payload.self, prim)) {
/* Only intersect with matching object and skip self-intersecton. */
if (object != payload.local_object) {
/* Only intersect with matching object */
result.accept = false;
result.continue_search = true;
return result;
@@ -187,11 +166,6 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
}
# endif
if (intersection_skip_self_shadow(payload.self, object, prim)) {
/* continue search */
return true;
}
float u = 0.0f, v = 0.0f;
int type = 0;
if (intersection_type == METALRT_HIT_TRIANGLE) {
@@ -348,35 +322,21 @@ inline TReturnType metalrt_visibility_test(constant KernelParamsMetal &launch_pa
}
# endif
uint visibility = payload.visibility;
# ifdef __VISIBILITY_FLAG__
uint visibility = payload.visibility;
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
result.accept = false;
result.continue_search = true;
return result;
}
# endif
/* Shadow ray early termination. */
if (visibility & PATH_RAY_SHADOW_OPAQUE) {
if (intersection_skip_self_shadow(payload.self, object, prim)) {
result.accept = false;
result.continue_search = true;
return result;
}
else {
result.accept = true;
result.continue_search = false;
return result;
}
}
else {
if (intersection_skip_self(payload.self, object, prim)) {
result.accept = false;
result.continue_search = true;
return result;
}
result.accept = true;
result.continue_search = false;
return result;
}
# endif
result.accept = true;
result.continue_search = true;

View File

@@ -45,11 +45,6 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
return pointer_unpack_from_uint<T>(optixGetPayload_2(), optixGetPayload_3());
}
template<typename T> ccl_device_forceinline T *get_payload_ptr_6()
{
return (T *)(((uint64_t)optixGetPayload_7() << 32) | optixGetPayload_6());
}
ccl_device_forceinline int get_object_id()
{
#ifdef __OBJECT_MOTION__
@@ -116,12 +111,6 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
return optixIgnoreIntersection();
}
const int prim = optixGetPrimitiveIndex();
ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
if (intersection_skip_self_local(ray->self, prim)) {
return optixIgnoreIntersection();
}
const uint max_hits = optixGetPayload_5();
if (max_hits == 0) {
/* Special case for when no hit information is requested, just report that something was hit */
@@ -160,6 +149,8 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
local_isect->num_hits = 1;
}
const int prim = optixGetPrimitiveIndex();
Intersection *isect = &local_isect->hits[hit];
isect->t = optixGetRayTmax();
isect->prim = prim;
@@ -194,11 +185,6 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
}
# endif
ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
if (intersection_skip_self_shadow(ray->self, object, prim)) {
return optixIgnoreIntersection();
}
float u = 0.0f, v = 0.0f;
int type = 0;
if (optixIsTriangleHit()) {
@@ -328,12 +314,6 @@ extern "C" __global__ void __anyhit__kernel_optix_volume_test()
if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
return optixIgnoreIntersection();
}
const int prim = optixGetPrimitiveIndex();
ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
if (intersection_skip_self(ray->self, object, prim)) {
return optixIgnoreIntersection();
}
}
extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
@@ -350,31 +330,18 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
# endif
#endif
#ifdef __VISIBILITY_FLAG__
const uint object = get_object_id();
const uint visibility = optixGetPayload_4();
#ifdef __VISIBILITY_FLAG__
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
return optixIgnoreIntersection();
}
#endif
const int prim = optixGetPrimitiveIndex();
ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
/* Shadow ray early termination. */
if (visibility & PATH_RAY_SHADOW_OPAQUE) {
if (intersection_skip_self_shadow(ray->self, object, prim)) {
return optixIgnoreIntersection();
}
else {
/* Shadow ray early termination. */
return optixTerminateRay();
}
}
else {
if (intersection_skip_self(ray->self, object, prim)) {
return optixIgnoreIntersection();
}
return optixTerminateRay();
}
#endif
}
extern "C" __global__ void __closesthit__kernel_optix_hit()

View File

@@ -29,19 +29,46 @@
CCL_NAMESPACE_BEGIN
/**
* Use the barycentric coordinates to get the intersection location
/* Refine triangle intersection to more precise hit point. For rays that travel
* far the precision is often not so good, this reintersects the primitive from
* a closer distance.
*/
ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg,
ccl_private ShaderData *sd,
const int isect_object,
const int isect_prim,
const float u,
const float v,
float3 verts[3])
ccl_device_inline float3 motion_triangle_refine(KernelGlobals kg,
ccl_private ShaderData *sd,
float3 P,
float3 D,
float t,
const int isect_object,
const int isect_prim,
float3 verts[3])
{
float w = 1.0f - u - v;
float3 P = u * verts[0] + v * verts[1] + w * verts[2];
#ifdef __INTERSECTION_REFINE__
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
if (UNLIKELY(t == 0.0f)) {
return P;
}
const Transform tfm = object_get_inverse_transform(kg, sd);
P = transform_point(&tfm, P);
D = transform_direction(&tfm, D * t);
D = normalize_len(D, &t);
}
P = P + D * t;
/* Compute refined intersection distance. */
const float3 e1 = verts[0] - verts[2];
const float3 e2 = verts[1] - verts[2];
const float3 s1 = cross(D, e2);
const float invdivisor = 1.0f / dot(s1, e1);
const float3 d = P - verts[2];
const float3 s2 = cross(d, e1);
float rt = dot(e2, s2) * invdivisor;
/* Compute refined position. */
P = P + D * rt;
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_transform(kg, sd);
@@ -49,8 +76,71 @@ ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg,
}
return P;
#else
return P + D * t;
#endif
}
/* Same as above, except that t is assumed to be in object space
* for instancing.
*/
#ifdef __BVH_LOCAL__
# if defined(__KERNEL_CUDA__) && (defined(i386) || defined(_M_IX86))
ccl_device_noinline
# else
ccl_device_inline
# endif
float3
motion_triangle_refine_local(KernelGlobals kg,
ccl_private ShaderData *sd,
float3 P,
float3 D,
float t,
const int isect_object,
const int isect_prim,
float3 verts[3])
{
# if defined(__KERNEL_GPU_RAYTRACING__)
/* t is always in world space with OptiX and MetalRT. */
return motion_triangle_refine(kg, sd, P, D, t, isect_object, isect_prim, verts);
# else
# ifdef __INTERSECTION_REFINE__
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_inverse_transform(kg, sd);
P = transform_point(&tfm, P);
D = transform_direction(&tfm, D);
D = normalize(D);
}
P = P + D * t;
/* compute refined intersection distance */
const float3 e1 = verts[0] - verts[2];
const float3 e2 = verts[1] - verts[2];
const float3 s1 = cross(D, e2);
const float invdivisor = 1.0f / dot(s1, e1);
const float3 d = P - verts[2];
const float3 s2 = cross(d, e1);
float rt = dot(e2, s2) * invdivisor;
P = P + D * rt;
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_transform(kg, sd);
P = transform_point(&tfm, P);
}
return P;
# else /* __INTERSECTION_REFINE__ */
return P + D * t;
# endif /* __INTERSECTION_REFINE__ */
# endif
}
#endif /* __BVH_LOCAL__ */
/* Ray intersection. We simply compute the vertex positions at the given ray
* time and do a ray intersection with the resulting triangle.
*/

View File

@@ -68,7 +68,15 @@ ccl_device_noinline void motion_triangle_shader_setup(KernelGlobals kg,
verts[1] = (1.0f - t) * verts[1] + t * next_verts[1];
verts[2] = (1.0f - t) * verts[2] + t * next_verts[2];
/* Compute refined position. */
sd->P = motion_triangle_point_from_uv(kg, sd, isect_object, isect_prim, sd->u, sd->v, verts);
#ifdef __BVH_LOCAL__
if (is_local) {
sd->P = motion_triangle_refine_local(kg, sd, P, D, ray_t, isect_object, isect_prim, verts);
}
else
#endif /* __BVH_LOCAL__*/
{
sd->P = motion_triangle_refine(kg, sd, P, D, ray_t, isect_object, isect_prim, verts);
}
/* Compute face normal. */
float3 Ng;
if (sd->object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {

View File

@@ -89,7 +89,7 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
sd->shader = kernel_tex_fetch(__tri_shader, sd->prim);
/* vectors */
sd->P = triangle_point_from_uv(kg, sd, isect->object, isect->prim, isect->u, isect->v);
sd->P = triangle_refine(kg, sd, ray->P, ray->D, isect->t, isect->object, isect->prim);
sd->Ng = Ng;
sd->N = Ng;
@@ -190,46 +190,40 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals kg,
#ifdef __OBJECT_MOTION__
shader_setup_object_transforms(kg, sd, time);
#endif
}
else if (lamp != LAMP_NONE) {
sd->lamp = lamp;
}
/* transform into world space */
if (object_space) {
object_position_transform_auto(kg, sd, &sd->P);
object_normal_transform_auto(kg, sd, &sd->Ng);
sd->N = sd->Ng;
object_dir_transform_auto(kg, sd, &sd->I);
}
/* transform into world space */
if (object_space) {
object_position_transform_auto(kg, sd, &sd->P);
object_normal_transform_auto(kg, sd, &sd->Ng);
sd->N = sd->Ng;
object_dir_transform_auto(kg, sd, &sd->I);
}
if (sd->type == PRIMITIVE_TRIANGLE) {
/* smooth normal */
if (sd->shader & SHADER_SMOOTH_NORMAL) {
sd->N = triangle_smooth_normal(kg, Ng, sd->prim, sd->u, sd->v);
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
object_normal_transform_auto(kg, sd, &sd->N);
}
}
/* dPdu/dPdv */
#ifdef __DPDU__
triangle_dPdudv(kg, sd->prim, &sd->dPdu, &sd->dPdv);
if (sd->type == PRIMITIVE_TRIANGLE) {
/* smooth normal */
if (sd->shader & SHADER_SMOOTH_NORMAL) {
sd->N = triangle_smooth_normal(kg, Ng, sd->prim, sd->u, sd->v);
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
object_dir_transform_auto(kg, sd, &sd->dPdu);
object_dir_transform_auto(kg, sd, &sd->dPdv);
object_normal_transform_auto(kg, sd, &sd->N);
}
#endif
}
else {
/* dPdu/dPdv */
#ifdef __DPDU__
sd->dPdu = zero_float3();
sd->dPdv = zero_float3();
#endif
triangle_dPdudv(kg, sd->prim, &sd->dPdu, &sd->dPdv);
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
object_dir_transform_auto(kg, sd, &sd->dPdu);
object_dir_transform_auto(kg, sd, &sd->dPdv);
}
#endif
}
else {
if (lamp != LAMP_NONE) {
sd->lamp = lamp;
}
#ifdef __DPDU__
sd->dPdu = zero_float3();
sd->dPdv = zero_float3();

View File

@@ -142,23 +142,58 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg,
}
#endif /* __BVH_LOCAL__ */
/**
* Use the barycentric coordinates to get the intersection location
/* Refine triangle intersection to more precise hit point. For rays that travel
* far the precision is often not so good, this reintersects the primitive from
* a closer distance. */
/* Reintersections uses the paper:
*
* Tomas Moeller
* Fast, minimum storage ray/triangle intersection
* http://www.cs.virginia.edu/~gfx/Courses/2003/ImageSynthesis/papers/Acceleration/Fast%20MinimumStorage%20RayTriangle%20Intersection.pdf
*/
ccl_device_inline float3 triangle_point_from_uv(KernelGlobals kg,
ccl_private ShaderData *sd,
const int isect_object,
const int isect_prim,
const float u,
const float v)
ccl_device_inline float3 triangle_refine(KernelGlobals kg,
ccl_private ShaderData *sd,
float3 P,
float3 D,
float t,
const int isect_object,
const int isect_prim)
{
#ifdef __INTERSECTION_REFINE__
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
if (UNLIKELY(t == 0.0f)) {
return P;
}
const Transform tfm = object_get_inverse_transform(kg, sd);
P = transform_point(&tfm, P);
D = transform_direction(&tfm, D * t);
D = normalize_len(D, &t);
}
P = P + D * t;
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
const packed_float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
float w = 1.0f - u - v;
float3 P = u * tri_a + v * tri_b + w * tri_c;
float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
float3 qvec = cross(tvec, edge1);
float3 pvec = cross(D, edge2);
float det = dot(edge1, pvec);
if (det != 0.0f) {
/* If determinant is zero it means ray lies in the plane of
* the triangle. It is possible in theory due to watertight
* nature of triangle intersection. For such cases we simply
* don't refine intersection hoping it'll go all fine.
*/
float rt = dot(edge2, qvec) / det;
P = P + D * rt;
}
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_transform(kg, sd);
@@ -166,6 +201,65 @@ ccl_device_inline float3 triangle_point_from_uv(KernelGlobals kg,
}
return P;
#else
return P + D * t;
#endif
}
/* Same as above, except that t is assumed to be in object space for
* instancing.
*/
ccl_device_inline float3 triangle_refine_local(KernelGlobals kg,
ccl_private ShaderData *sd,
float3 P,
float3 D,
float t,
const int isect_object,
const int isect_prim)
{
#if defined(__KERNEL_GPU_RAYTRACING__)
/* t is always in world space with OptiX and MetalRT. */
return triangle_refine(kg, sd, P, D, t, isect_object, isect_prim);
#else
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_inverse_transform(kg, sd);
P = transform_point(&tfm, P);
D = transform_direction(&tfm, D);
D = normalize(D);
}
P = P + D * t;
# ifdef __INTERSECTION_REFINE__
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
const packed_float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
float3 qvec = cross(tvec, edge1);
float3 pvec = cross(D, edge2);
float det = dot(edge1, pvec);
if (det != 0.0f) {
/* If determinant is zero it means ray lies in the plane of
* the triangle. It is possible in theory due to watertight
* nature of triangle intersection. For such cases we simply
* don't refine intersection hoping it'll go all fine.
*/
float rt = dot(edge2, qvec) / det;
P = P + D * rt;
}
# endif /* __INTERSECTION_REFINE__ */
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_transform(kg, sd);
P = transform_point(&tfm, P);
}
return P;
#endif
}
CCL_NAMESPACE_END

View File

@@ -328,12 +328,6 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg,
/* Scene Intersection. */
Intersection isect ccl_optional_struct_init;
isect.object = OBJECT_NONE;
isect.prim = PRIM_NONE;
ray.self.object = last_isect_object;
ray.self.prim = last_isect_prim;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
bool hit = scene_intersect(kg, &ray, visibility, &isect);
/* TODO: remove this and do it in the various intersection functions instead. */

View File

@@ -156,10 +156,7 @@ ccl_device void integrator_intersect_shadow(KernelGlobals kg, IntegratorShadowSt
/* Read ray from integrator state into local memory. */
Ray ray ccl_optional_struct_init;
integrator_state_read_shadow_ray(kg, state, &ray);
ray.self.object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, object);
ray.self.prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, prim);
ray.self.light_object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 1, object);
ray.self.light_prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 1, prim);
/* Compute visibility. */
const uint visibility = integrate_intersect_shadow_visibility(kg, state);

View File

@@ -38,10 +38,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
Ray volume_ray ccl_optional_struct_init;
volume_ray.P = from_P;
volume_ray.D = normalize_len(to_P - from_P, &volume_ray.t);
volume_ray.self.object = INTEGRATOR_STATE(state, isect, object);
volume_ray.self.prim = INTEGRATOR_STATE(state, isect, prim);
volume_ray.self.light_object = OBJECT_NONE;
volume_ray.self.light_prim = PRIM_NONE;
/* Store to avoid global fetches on every intersection step. */
const uint volume_stack_size = kernel_data.volume_stack_size;
@@ -71,7 +68,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
volume_stack_enter_exit(kg, state, stack_sd);
/* Move ray forward. */
volume_ray.P = stack_sd->P;
volume_ray.P = ray_offset(stack_sd->P, -stack_sd->Ng);
if (volume_ray.t != FLT_MAX) {
volume_ray.D = normalize_len(to_P - volume_ray.P, &volume_ray.t);
}
@@ -94,10 +91,6 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
* fewest hits. */
volume_ray.D = make_float3(0.0f, 0.0f, 1.0f);
volume_ray.t = FLT_MAX;
volume_ray.self.object = OBJECT_NONE;
volume_ray.self.prim = PRIM_NONE;
volume_ray.self.light_object = OBJECT_NONE;
volume_ray.self.light_prim = PRIM_NONE;
int stack_index = 0, enclosed_index = 0;
@@ -210,7 +203,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
}
/* Move ray forward. */
volume_ray.P = stack_sd->P;
volume_ray.P = ray_offset(stack_sd->P, -stack_sd->Ng);
++step;
}
#endif

View File

@@ -37,9 +37,8 @@ ccl_device_inline void integrate_light(KernelGlobals kg,
/* Advance ray beyond light. */
/* TODO: can we make this more numerically robust to avoid reintersecting the
* same light in some cases? Ray should not intersect surface anymore as the
* object and prim ids will prevent self intersection. */
const float3 new_ray_P = ray_P + ray_D * isect.t;
* same light in some cases? */
const float3 new_ray_P = ray_offset(ray_P + ray_D * isect.t, ray_D);
INTEGRATOR_STATE_WRITE(state, ray, P) = new_ray_P;
INTEGRATOR_STATE_WRITE(state, ray, t) -= isect.t;
@@ -47,7 +46,7 @@ ccl_device_inline void integrate_light(KernelGlobals kg,
const float mis_ray_t = INTEGRATOR_STATE(state, path, mis_ray_t);
ray_P -= ray_D * mis_ray_t;
isect.t += mis_ray_t;
INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = isect.t;
INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = mis_ray_t + isect.t;
LightSample ls ccl_optional_struct_init;
const bool use_light_sample = light_sample_from_intersection(kg, &isect, ray_P, ray_D, &ls);

View File

@@ -83,10 +83,7 @@ ccl_device_inline void integrate_transparent_volume_shadow(KernelGlobals kg,
/* Setup shader data. */
Ray ray ccl_optional_struct_init;
integrator_state_read_shadow_ray(kg, state, &ray);
ray.self.object = OBJECT_NONE;
ray.self.prim = PRIM_NONE;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
/* Modify ray position and length to match current segment. */
const float start_t = (hit == 0) ? 0.0f :
INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit - 1, t);
@@ -152,7 +149,7 @@ ccl_device_inline bool integrate_transparent_shadow(KernelGlobals kg,
const float last_hit_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, num_recorded_hits - 1, t);
const float3 ray_P = INTEGRATOR_STATE(state, shadow_ray, P);
const float3 ray_D = INTEGRATOR_STATE(state, shadow_ray, D);
INTEGRATOR_STATE_WRITE(state, shadow_ray, P) = ray_P + last_hit_t * ray_D;
INTEGRATOR_STATE_WRITE(state, shadow_ray, P) = ray_offset(ray_P + last_hit_t * ray_D, ray_D);
INTEGRATOR_STATE_WRITE(state, shadow_ray, t) -= last_hit_t;
}

View File

@@ -182,35 +182,23 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(kg, shadow_state, &ray);
// Save memory by storing the light and object indices in the shadow_isect
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
/* Copy state from main path to shadow path. */
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;
shadow_flag |= (shadow_flag & PATH_RAY_ANY_PASS) ? 0 : PATH_RAY_SURFACE_PASS;
const float3 throughput = INTEGRATOR_STATE(state, path, throughput) * bsdf_eval_sum(&bsdf_eval);
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
packed_float3 pass_diffuse_weight;
packed_float3 pass_glossy_weight;
if (shadow_flag & PATH_RAY_ANY_PASS) {
/* Indirect bounce, use weights from earlier surface or volume bounce. */
pass_diffuse_weight = INTEGRATOR_STATE(state, path, pass_diffuse_weight);
pass_glossy_weight = INTEGRATOR_STATE(state, path, pass_glossy_weight);
}
else {
/* Direct light, use BSDFs at this bounce. */
shadow_flag |= PATH_RAY_SURFACE_PASS;
pass_diffuse_weight = packed_float3(bsdf_eval_pass_diffuse_weight(&bsdf_eval));
pass_glossy_weight = packed_float3(bsdf_eval_pass_glossy_weight(&bsdf_eval));
}
const packed_float3 pass_diffuse_weight =
(bounce == 0) ? packed_float3(bsdf_eval_pass_diffuse_weight(&bsdf_eval)) :
INTEGRATOR_STATE(state, path, pass_diffuse_weight);
const packed_float3 pass_glossy_weight = (bounce == 0) ?
packed_float3(
bsdf_eval_pass_glossy_weight(&bsdf_eval)) :
INTEGRATOR_STATE(state, path, pass_glossy_weight);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = pass_diffuse_weight;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = pass_glossy_weight;
}
@@ -278,11 +266,13 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
}
/* Setup ray. Note that clipping works through transparent bounces. */
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
INTEGRATOR_STATE_WRITE(state, ray, P) = ray_offset(sd->P,
(label & LABEL_TRANSMIT) ? -sd->Ng : sd->Ng);
INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(bsdf_omega_in);
INTEGRATOR_STATE_WRITE(state, ray, t) = (label & LABEL_TRANSPARENT) ?
INTEGRATOR_STATE(state, ray, t) - sd->ray_length :
FLT_MAX;
#ifdef __RAY_DIFFERENTIALS__
INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(bsdf_domega_in);
@@ -326,7 +316,7 @@ ccl_device_forceinline bool integrate_surface_volume_only_bounce(IntegratorState
}
/* Setup ray position, direction stays unchanged. */
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
INTEGRATOR_STATE_WRITE(state, ray, P) = ray_offset(sd->P, -sd->Ng);
/* Clipping works through transparent. */
INTEGRATOR_STATE_WRITE(state, ray, t) -= sd->ray_length;
@@ -370,14 +360,10 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
}
Ray ray ccl_optional_struct_init;
ray.P = sd->P;
ray.P = ray_offset(sd->P, sd->Ng);
ray.D = ao_D;
ray.t = kernel_data.integrator.ao_bounces_distance;
ray.time = sd->time;
ray.self.object = sd->object;
ray.self.prim = sd->prim;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.dP = differential_zero_compact();
ray.dD = differential_zero_compact();
@@ -389,10 +375,6 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(kg, shadow_state, &ray);
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
/* Copy state from main path to shadow path. */
const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);

View File

@@ -791,36 +791,22 @@ ccl_device_forceinline void integrate_volume_direct_light(
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(kg, shadow_state, &ray);
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
/* Copy state from main path to shadow path. */
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;
shadow_flag |= (shadow_flag & PATH_RAY_ANY_PASS) ? 0 : PATH_RAY_VOLUME_PASS;
const float3 throughput_phase = throughput * bsdf_eval_sum(&phase_eval);
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
packed_float3 pass_diffuse_weight;
packed_float3 pass_glossy_weight;
if (shadow_flag & PATH_RAY_ANY_PASS) {
/* Indirect bounce, use weights from earlier surface or volume bounce. */
pass_diffuse_weight = INTEGRATOR_STATE(state, path, pass_diffuse_weight);
pass_glossy_weight = INTEGRATOR_STATE(state, path, pass_glossy_weight);
}
else {
/* Direct light, no diffuse/glossy distinction needed for volumes. */
shadow_flag |= PATH_RAY_VOLUME_PASS;
pass_diffuse_weight = packed_float3(one_float3());
pass_glossy_weight = packed_float3(zero_float3());
}
const packed_float3 pass_diffuse_weight = (bounce == 0) ?
packed_float3(one_float3()) :
INTEGRATOR_STATE(
state, path, pass_diffuse_weight);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = pass_diffuse_weight;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = pass_glossy_weight;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = zero_float3();
}
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
@@ -887,13 +873,11 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(phase_omega_in);
INTEGRATOR_STATE_WRITE(state, ray, t) = FLT_MAX;
# ifdef __RAY_DIFFERENTIALS__
INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(phase_domega_in);
# endif
// Save memory by storing last hit prim and object in isect
INTEGRATOR_STATE_WRITE(state, isect, prim) = sd->prim;
INTEGRATOR_STATE_WRITE(state, isect, object) = sd->object;
/* Update throughput. */
const float3 throughput = INTEGRATOR_STATE(state, path, throughput);

View File

@@ -61,7 +61,6 @@ KERNEL_STRUCT_MEMBER(shadow_ray, packed_float3, D, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, t, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, time, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, dP, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, int, object, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_END(shadow_ray)
/*********************** Shadow Intersection result **************************/

View File

@@ -57,6 +57,7 @@ ccl_device int subsurface_bounce(KernelGlobals kg,
/* Pass along object info, reusing isect to save memory. */
INTEGRATOR_STATE_WRITE(state, subsurface, Ng) = sd->Ng;
INTEGRATOR_STATE_WRITE(state, isect, object) = sd->object;
uint32_t path_flag = (INTEGRATOR_STATE(state, path, flag) & ~PATH_RAY_CAMERA) |
((sc->type == CLOSURE_BSSRDF_BURLEY_ID) ? PATH_RAY_SUBSURFACE_DISK :
@@ -164,8 +165,10 @@ ccl_device_inline bool subsurface_scatter(KernelGlobals kg, IntegratorState stat
if (object_flag & SD_OBJECT_INTERSECTS_VOLUME) {
float3 P = INTEGRATOR_STATE(state, ray, P);
const float3 Ng = INTEGRATOR_STATE(state, subsurface, Ng);
const float3 offset_P = ray_offset(P, -Ng);
integrator_volume_stack_update_for_subsurface(kg, state, P, ray.P);
integrator_volume_stack_update_for_subsurface(kg, state, offset_P, ray.P);
}
}
# endif /* __VOLUME__ */

View File

@@ -99,10 +99,6 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg,
ray.dP = ray_dP;
ray.dD = differential_zero_compact();
ray.time = time;
ray.self.object = OBJECT_NONE;
ray.self.prim = PRIM_NONE;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = OBJECT_NONE;
/* Intersect with the same object. if multiple intersections are found it
* will use at most BSSRDF_MAX_HITS hits, a random subset of all hits. */

View File

@@ -195,7 +195,6 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
const float time = INTEGRATOR_STATE(state, ray, time);
const float3 Ng = INTEGRATOR_STATE(state, subsurface, Ng);
const int object = INTEGRATOR_STATE(state, isect, object);
const int prim = INTEGRATOR_STATE(state, isect, prim);
/* Sample diffuse surface scatter into the object. */
float3 D;
@@ -206,16 +205,12 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
}
/* Setup ray. */
ray.P = P;
ray.P = ray_offset(P, -Ng);
ray.D = D;
ray.t = FLT_MAX;
ray.time = time;
ray.dP = ray_dP;
ray.dD = differential_zero_compact();
ray.self.object = object;
ray.self.prim = prim;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
#ifndef __KERNEL_GPU_RAYTRACING__
/* Compute or fetch object transforms. */
@@ -382,15 +377,7 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
* If yes, we will later use backwards guided sampling in order to have a decent
* chance of connecting to it.
* TODO: Maybe use less than 10 times the mean free path? */
if (bounce == 0) {
ray.t = max(t, 10.0f / (min3(sigma_t)));
}
else {
ray.t = t;
/* After the first bounce the object can intersect the same surface again */
ray.self.object = OBJECT_NONE;
ray.self.prim = PRIM_NONE;
}
ray.t = (bounce == 0) ? max(t, 10.0f / (min3(sigma_t))) : t;
scene_intersect_local(kg, &ray, &ss_isect, object, NULL, 1);
hit = (ss_isect.num_hits > 0);
@@ -421,6 +408,13 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
if (hit) {
t = ray.t;
}
else if (bounce == 0) {
/* Restore original position if nothing was hit after the first bounce,
* without the ray_offset() that was added to avoid self-intersection.
* Otherwise if that offset is relatively large compared to the scattering
* radius, we never go back up high enough to exit the surface. */
ray.P = P;
}
/* Advance to new scatter location. */
ray.P += t * ray.D;

View File

@@ -113,30 +113,22 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
ls->P = make_float3(klight->co[0], klight->co[1], klight->co[2]);
if (type == LIGHT_SPOT) {
const float3 center = make_float3(klight->co[0], klight->co[1], klight->co[2]);
const float radius = klight->spot.radius;
const float3 dir = make_float3(
klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
/* disk oriented normal */
const float3 lightN = normalize(P - center);
ls->P = center;
ls->Ng = make_float3(klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
float radius = klight->spot.radius;
if (radius > 0.0f)
/* disk light */
ls->P += disk_light_sample(lightN, randu, randv) * radius;
const float invarea = klight->spot.invarea;
ls->pdf = invarea;
/* sphere light */
ls->P += disk_light_sample(ls->Ng, randu, randv) * radius;
ls->D = normalize_len(ls->P - P, &ls->t);
/* we set the light normal to the outgoing direction to support texturing */
ls->Ng = -ls->D;
float invarea = klight->spot.invarea;
ls->eval_fac = (0.25f * M_1_PI_F) * invarea;
ls->pdf = invarea;
/* spot light attenuation */
ls->eval_fac *= spot_light_attenuation(
dir, klight->spot.spot_angle, klight->spot.spot_smooth, -ls->D);
ls->Ng, klight->spot.spot_angle, klight->spot.spot_smooth, -ls->D);
if (!in_volume_segment && ls->eval_fac == 0.0f) {
return false;
}
@@ -145,33 +137,32 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
ls->u = uv.x;
ls->v = uv.y;
ls->pdf *= lamp_light_pdf(kg, lightN, -ls->D, ls->t);
ls->pdf *= lamp_light_pdf(kg, ls->Ng, -ls->D, ls->t);
}
else if (type == LIGHT_POINT) {
float3 center = make_float3(klight->co[0], klight->co[1], klight->co[2]);
float radius = klight->spot.radius;
/* disk oriented normal */
const float3 lightN = normalize(P - center);
ls->P = center;
float pdf = 1.0;
if (radius > 0.0f) {
ls->P += disk_light_sample(lightN, randu, randv) * radius;
ls->Ng = normalize(P - center);
ls->P += disk_light_sample(ls->Ng, randu, randv) * radius;
pdf = klight->spot.invarea;
ls->D = normalize_len(ls->P - P, &ls->t);
}
else {
ls->Ng = normalize(P - center);
}
ls->pdf = klight->spot.invarea;
ls->D = normalize_len(ls->P - P, &ls->t);
/* we set the light normal to the outgoing direction to support texturing */
ls->Ng = -ls->D;
ls->pdf = pdf;
ls->eval_fac = M_1_PI_F * 0.25f * klight->spot.invarea;
if (!in_volume_segment && ls->eval_fac == 0.0f) {
return false;
}
float2 uv = map_to_sphere(ls->Ng);
ls->u = uv.x;
ls->v = uv.y;
ls->pdf *= lamp_light_pdf(kg, lightN, -ls->D, ls->t);
ls->pdf *= lamp_light_pdf(kg, ls->Ng, -ls->D, ls->t);
}
else {
/* area light */
@@ -272,16 +263,14 @@ ccl_device bool lights_intersect(KernelGlobals kg,
if (type == LIGHT_SPOT) {
/* Spot/Disk light. */
const float mis_ray_t = INTEGRATOR_STATE(state, path, mis_ray_t);
const float3 ray_P = ray->P - ray->D * mis_ray_t;
const float3 lightP = make_float3(klight->co[0], klight->co[1], klight->co[2]);
const float3 lightN = make_float3(
klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
const float radius = klight->spot.radius;
if (radius == 0.0f) {
continue;
}
/* disk oriented normal */
const float3 lightN = normalize(ray_P - lightP);
/* One sided. */
if (dot(ray->D, lightN) >= 0.0f) {
continue;
@@ -303,10 +292,9 @@ ccl_device bool lights_intersect(KernelGlobals kg,
continue;
}
/* disk oriented normal */
const float3 lightN = normalize(ray_P - lightP);
float3 P;
if (!ray_disk_intersect(ray->P, ray->D, ray->t, lightP, lightN, radius, &P, &t)) {
const float3 lsN = normalize(ray_P - lightP);
if (!ray_disk_intersect(ray->P, ray->D, ray->t, lightP, lsN, radius, &P, &t)) {
continue;
}
}
@@ -430,8 +418,8 @@ ccl_device bool light_sample_from_intersection(KernelGlobals kg,
LightType type = (LightType)klight->type;
ls->type = type;
ls->shader = klight->shader_id;
ls->object = isect->object;
ls->prim = isect->prim;
ls->object = PRIM_NONE;
ls->prim = PRIM_NONE;
ls->lamp = lamp;
/* todo: missing texture coordinates */
ls->t = isect->t;
@@ -439,12 +427,7 @@ ccl_device bool light_sample_from_intersection(KernelGlobals kg,
ls->D = ray_D;
if (type == LIGHT_SPOT) {
const float3 center = make_float3(klight->co[0], klight->co[1], klight->co[2]);
const float3 dir = make_float3(klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
/* the normal of the oriented disk */
const float3 lightN = normalize(ray_P - center);
/* we set the light normal to the outgoing direction to support texturing*/
ls->Ng = -ls->D;
ls->Ng = make_float3(klight->spot.dir[0], klight->spot.dir[1], klight->spot.dir[2]);
float invarea = klight->spot.invarea;
ls->eval_fac = (0.25f * M_1_PI_F) * invarea;
@@ -452,7 +435,7 @@ ccl_device bool light_sample_from_intersection(KernelGlobals kg,
/* spot light attenuation */
ls->eval_fac *= spot_light_attenuation(
dir, klight->spot.spot_angle, klight->spot.spot_smooth, -ls->D);
ls->Ng, klight->spot.spot_angle, klight->spot.spot_smooth, -ls->D);
if (ls->eval_fac == 0.0f) {
return false;
@@ -464,32 +447,23 @@ ccl_device bool light_sample_from_intersection(KernelGlobals kg,
/* compute pdf */
if (ls->t != FLT_MAX)
ls->pdf *= lamp_light_pdf(kg, lightN, -ls->D, ls->t);
else
ls->pdf = 0.f;
ls->pdf *= lamp_light_pdf(kg, ls->Ng, -ls->D, ls->t);
}
else if (type == LIGHT_POINT) {
const float3 center = make_float3(klight->co[0], klight->co[1], klight->co[2]);
const float3 lighN = normalize(ray_P - center);
/* we set the light normal to the outgoing direction to support texturing*/
ls->Ng = -ls->D;
float3 center = make_float3(klight->co[0], klight->co[1], klight->co[2]);
ls->Ng = normalize(ray_P - center);
float invarea = klight->spot.invarea;
ls->eval_fac = (0.25f * M_1_PI_F) * invarea;
ls->pdf = invarea;
if (ls->eval_fac == 0.0f) {
return false;
}
float2 uv = map_to_sphere(ls->Ng);
ls->u = uv.x;
ls->v = uv.y;
/* compute pdf */
if (ls->t != FLT_MAX)
ls->pdf *= lamp_light_pdf(kg, lighN, -ls->D, ls->t);
ls->pdf *= lamp_light_pdf(kg, ls->Ng, -ls->D, ls->t);
else
ls->pdf = 0.f;
}
@@ -947,4 +921,4 @@ ccl_device_inline bool light_distribution_sample_new_position(KernelGlobals kg,
}
}
CCL_NAMESPACE_END
CCL_NAMESPACE_END

View File

@@ -198,7 +198,7 @@ ccl_device_inline float3 shadow_ray_offset(KernelGlobals kg,
float NL = dot(sd->N, L);
bool transmit = (NL < 0.0f);
float3 Ng = (transmit ? -sd->Ng : sd->Ng);
float3 P = sd->P;
float3 P = ray_offset(sd->P, Ng);
if ((sd->type & PRIMITIVE_TRIANGLE) && (sd->shader & SHADER_SMOOTH_NORMAL)) {
const float offset_cutoff =
@@ -243,7 +243,7 @@ ccl_device_inline void shadow_ray_setup(ccl_private const ShaderData *ccl_restri
}
else {
/* other lights, avoid self-intersection */
ray->D = ls->P - P;
ray->D = ray_offset(ls->P, ls->Ng) - P;
ray->D = normalize_len(ray->D, &ray->t);
}
}
@@ -257,12 +257,6 @@ ccl_device_inline void shadow_ray_setup(ccl_private const ShaderData *ccl_restri
ray->dP = differential_make_compact(sd->dP);
ray->dD = differential_zero_compact();
ray->time = sd->time;
/* Fill in intersection surface and light details. */
ray->self.prim = sd->prim;
ray->self.object = sd->object;
ray->self.light_prim = ls->prim;
ray->self.light_object = ls->object;
}
/* Create shadow ray towards light sample. */

View File

@@ -70,14 +70,10 @@ ccl_device float svm_ao(
/* Create ray. */
Ray ray;
ray.P = sd->P;
ray.P = ray_offset(sd->P, N);
ray.D = D.x * T + D.y * B + D.z * N;
ray.t = max_dist;
ray.time = sd->time;
ray.self.object = sd->object;
ray.self.prim = sd->prim;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.dP = differential_zero_compact();
ray.dD = differential_zero_compact();

View File

@@ -196,10 +196,6 @@ ccl_device float3 svm_bevel(
ray.dP = differential_zero_compact();
ray.dD = differential_zero_compact();
ray.time = sd->time;
ray.self.object = OBJECT_NONE;
ray.self.prim = PRIM_NONE;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
/* Intersect with the same object. if multiple intersections are found it
* will use at most LOCAL_MAX_HITS hits, a random subset of all hits. */
@@ -211,24 +207,15 @@ ccl_device float3 svm_bevel(
/* Quickly retrieve P and Ng without setting up ShaderData. */
float3 hit_P;
if (sd->type == PRIMITIVE_TRIANGLE) {
hit_P = triangle_point_from_uv(kg,
sd,
isect.hits[hit].object,
isect.hits[hit].prim,
isect.hits[hit].u,
isect.hits[hit].v);
hit_P = triangle_refine_local(
kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim);
}
# ifdef __OBJECT_MOTION__
else if (sd->type == PRIMITIVE_MOTION_TRIANGLE) {
float3 verts[3];
motion_triangle_vertices(kg, sd->object, isect.hits[hit].prim, sd->time, verts);
hit_P = motion_triangle_point_from_uv(kg,
sd,
isect.hits[hit].object,
isect.hits[hit].prim,
isect.hits[hit].u,
isect.hits[hit].v,
verts);
hit_P = motion_triangle_refine_local(
kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim, verts);
}
# endif /* __OBJECT_MOTION__ */

View File

@@ -512,21 +512,12 @@ typedef struct differential {
/* Ray */
typedef struct RaySelfPrimitives {
int prim; /* Primitive the ray is starting from */
int object; /* Instance prim is a part of */
int light_prim; /* Light primitive */
int light_object; /* Light object */
} RaySelfPrimitives;
typedef struct Ray {
float3 P; /* origin */
float3 D; /* direction */
float t; /* length of the ray */
float time; /* time (for motion blur) */
RaySelfPrimitives self;
#ifdef __RAY_DIFFERENTIALS__
float dP;
float dD;

View File

@@ -291,7 +291,7 @@ elseif(WITH_GHOST_X11 OR WITH_GHOST_WAYLAND)
include(CheckSymbolExists)
set(CMAKE_REQUIRED_DEFINITIONS "-D_GNU_SOURCE")
check_symbol_exists(memfd_create "sys/mman.h" HAVE_MEMFD_CREATE)
if(HAVE_MEMFD_CREATE)
if (HAVE_MEMFD_CREATE)
add_definitions(-DHAVE_MEMFD_CREATE)
endif()
@@ -307,7 +307,7 @@ elseif(WITH_GHOST_X11 OR WITH_GHOST_WAYLAND)
pkg_get_variable(WAYLAND_SCANNER wayland-scanner wayland_scanner)
pkg_check_modules(wayland-protocols wayland-protocols>=1.15)
if(${wayland-protocols_FOUND})
if (${wayland-protocols_FOUND})
pkg_get_variable(WAYLAND_PROTOCOLS_DIR wayland-protocols pkgdatadir)
else()
find_path(WAYLAND_PROTOCOLS_DIR
@@ -316,7 +316,7 @@ elseif(WITH_GHOST_X11 OR WITH_GHOST_WAYLAND)
)
endif()
if(NOT EXISTS ${WAYLAND_PROTOCOLS_DIR})
if (NOT EXISTS ${WAYLAND_PROTOCOLS_DIR})
message(FATAL_ERROR "path to wayland-protocols not found")
endif()
@@ -518,11 +518,11 @@ if(WITH_XR_OPENXR)
)
elseif(UNIX AND NOT APPLE)
list(APPEND XR_PLATFORM_DEFINES -DXR_OS_LINUX)
if(WITH_GHOST_WAYLAND)
if (WITH_GHOST_WAYLAND)
list(APPEND XR_PLATFORM_DEFINES -DXR_USE_PLATFORM_WAYLAND)
endif()
if(WITH_GHOST_X11)
if(WITH_GL_EGL)
if (WITH_GHOST_X11)
if (WITH_GL_EGL)
list(APPEND XR_PLATFORM_DEFINES -DXR_USE_PLATFORM_EGL)
else()
list(APPEND XR_PLATFORM_DEFINES -DXR_USE_PLATFORM_XLIB)

View File

@@ -889,11 +889,16 @@ extern char *GHOST_getClipboard(bool selection);
extern void GHOST_putClipboard(const char *buffer, bool selection);
/**
* Set the Console State
* \param action: console state
* Toggles console
* \param action:
* - 0: Hides
* - 1: Shows
* - 2: Toggles
* - 3: Hides if it runs not from command line
* - *: Does nothing
* \return current status (1 -visible, 0 - hidden)
*/
extern int setConsoleWindowState(GHOST_TConsoleWindowState action);
extern int GHOST_toggleConsole(int action);
/**
* Use native pixel size (MacBook pro 'retina'), if supported.

View File

@@ -411,11 +411,16 @@ class GHOST_ISystem {
#endif
/**
* Set the Console State
* \param action: console state
* Toggles console
* \param action:
* - 0: Hides.
* - 1: Shows
* - 2: Toggles
* - 3: Hides if it runs not from command line
* - *: Does nothing
* \return current status (1 -visible, 0 - hidden)
*/
virtual int setConsoleWindowState(GHOST_TConsoleWindowState action) = 0;
virtual int toggleConsole(int action) = 0;
/***************************************************************************************
* Access to clipboard.

View File

@@ -140,13 +140,6 @@ typedef enum {
// GHOST_kWindowStateUnModified,
} GHOST_TWindowState;
typedef enum {
GHOST_kConsoleWindowStateHide = 0,
GHOST_kConsoleWindowStateShow,
GHOST_kConsoleWindowStateToggle,
GHOST_kConsoleWindowStateHideForNonConsoleLaunch
} GHOST_TConsoleWindowState;
typedef enum { GHOST_kWindowOrderTop = 0, GHOST_kWindowOrderBottom } GHOST_TWindowOrder;
typedef enum {

View File

@@ -809,10 +809,10 @@ void GHOST_putClipboard(const char *buffer, bool selection)
system->putClipboard(buffer, selection);
}
int setConsoleWindowState(GHOST_TConsoleWindowState action)
int GHOST_toggleConsole(int action)
{
GHOST_ISystem *system = GHOST_ISystem::getSystem();
return system->setConsoleWindowState(action);
return system->toggleConsole(action);
}
int GHOST_UseNativePixels(void)

View File

@@ -13,9 +13,11 @@
* along with this program; if not, write to the Free Software Foundation,
* Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
*
* Copyright 2001-2002 NaN Holding BV. All rights reserved.
* 1997-2001 Id Software, Inc. Video mode switching.
Ported from Quake 2 by Alex Fraser <alex@phatcore.com>.
* The Original Code is Copyright (C) 2001-2002 by NaN Holding BV.
* All rights reserved.
* Video mode switching
* Copyright (C) 1997-2001 Id Software, Inc.
* Ported from Quake 2 by Alex Fraser <alex@phatcore.com>
*/
/** \file

View File

@@ -14,6 +14,9 @@
* Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
*
* The Original Code is Copyright (c) 2010 The Chromium Authors. All rights reserved.
* All rights reserved.
*
* The Original Code is: some of this file.
*/
/** \file
@@ -93,7 +96,7 @@ bool GHOST_ImeWin32::IsEnglishMode()
!(conversion_modes_ & (IME_CMODE_NATIVE | IME_CMODE_FULLSHAPE));
}
bool GHOST_ImeWin32::IsImeKeyEvent(char ascii, GHOST_TKey key)
bool GHOST_ImeWin32::IsImeKeyEvent(char ascii)
{
if (!(IsEnglishMode())) {
/* In Chinese, Japanese, Korean, all alpha keys are processed by IME. */
@@ -103,8 +106,7 @@ bool GHOST_ImeWin32::IsImeKeyEvent(char ascii, GHOST_TKey key)
if (IsLanguage(IMELANG_JAPANESE) && (ascii >= ' ' && ascii <= '~')) {
return true;
}
else if (IsLanguage(IMELANG_CHINESE) && ascii && strchr("!\"$'(),.:;<>?[\\]^_`/", ascii) &&
!(key == GHOST_kKeyNumpadPeriod)) {
else if (IsLanguage(IMELANG_CHINESE) && ascii && strchr("!\"$'(),.:;<>?[\\]^_`/", ascii)) {
return true;
}
}

View File

@@ -14,6 +14,9 @@
* Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
*
* The Original Code is Copyright (c) 2010 The Chromium Authors. All rights reserved.
* All rights reserved.
*
* The Original Code is: some of this file.
*/
/** \file
@@ -158,7 +161,7 @@ class GHOST_ImeWin32 {
bool IsEnglishMode();
/* Checks a key whether IME has to do handling. */
bool IsImeKeyEvent(char ascii, GHOST_TKey key);
bool IsImeKeyEvent(char ascii);
/**
* Create the IME windows, and allocate required resources for them.

View File

@@ -244,7 +244,7 @@ class GHOST_SystemCocoa : public GHOST_System {
/**
* \see GHOST_ISystem
*/
int setConsoleWindowState(GHOST_TConsoleWindowState action)
int toggleConsole(int action)
{
return 0;
}

View File

@@ -890,7 +890,7 @@ bool GHOST_SystemCocoa::processEvents(bool waitForEvent)
bool anyProcessed = false;
NSEvent *event;
/* TODO: implement timer? */
// TODO : implement timer ??
#if 0
do {
GHOST_TimerManager* timerMgr = getTimerManager();

View File

@@ -40,7 +40,7 @@ class GHOST_SystemNULL : public GHOST_System {
{
return false;
}
int setConsoleWindowState(GHOST_TConsoleWindowState action)
int toggleConsole(int action)
{
return 0;
}

View File

@@ -47,7 +47,7 @@ class GHOST_SystemSDL : public GHOST_System {
bool processEvents(bool waitForEvent);
int setConsoleWindowState(GHOST_TConsoleWindowState action)
int toggleConsole(int action)
{
return 0;
}

View File

@@ -1473,7 +1473,7 @@ bool GHOST_SystemWayland::processEvents(bool waitForEvent)
return fired || (getEventManager()->getNumEvents() > 0);
}
int GHOST_SystemWayland::setConsoleWindowState(GHOST_TConsoleWindowState /*action*/)
int GHOST_SystemWayland::toggleConsole(int /*action*/)
{
return 0;
}

View File

@@ -53,7 +53,7 @@ class GHOST_SystemWayland : public GHOST_System {
bool processEvents(bool waitForEvent) override;
int setConsoleWindowState(GHOST_TConsoleWindowState action) override;
int toggleConsole(int action) override;
GHOST_TSuccess getModifierKeys(GHOST_ModifierKeys &keys) const override;

View File

@@ -105,8 +105,6 @@
*/
#define BROKEN_PEEK_TOUCHPAD
static bool isStartedFromCommandPrompt();
static void initRawInput()
{
#ifdef WITH_INPUT_NDOF
@@ -168,10 +166,7 @@ GHOST_SystemWin32::~GHOST_SystemWin32()
{
// Shutdown COM
OleUninitialize();
if (isStartedFromCommandPrompt()) {
setConsoleWindowState(GHOST_kConsoleWindowStateShow);
}
toggleConsole(1);
}
uint64_t GHOST_SystemWin32::performanceCounterToMillis(__int64 perf_ticks) const
@@ -1225,7 +1220,7 @@ GHOST_EventKey *GHOST_SystemWin32::processKeyEvent(GHOST_WindowWin32 *window, RA
}
#ifdef WITH_INPUT_IME
if (window->getImeInput()->IsImeKeyEvent(ascii, key)) {
if (window->getImeInput()->IsImeKeyEvent(ascii)) {
return NULL;
}
#endif /* WITH_INPUT_IME */
@@ -2221,30 +2216,31 @@ static bool isStartedFromCommandPrompt()
return false;
}
int GHOST_SystemWin32::setConsoleWindowState(GHOST_TConsoleWindowState action)
int GHOST_SystemWin32::toggleConsole(int action)
{
HWND wnd = GetConsoleWindow();
switch (action) {
case GHOST_kConsoleWindowStateHideForNonConsoleLaunch: {
case 3: // startup: hide if not started from command prompt
{
if (!isStartedFromCommandPrompt()) {
ShowWindow(wnd, SW_HIDE);
m_consoleStatus = 0;
}
break;
}
case GHOST_kConsoleWindowStateHide:
case 0: // hide
ShowWindow(wnd, SW_HIDE);
m_consoleStatus = 0;
break;
case GHOST_kConsoleWindowStateShow:
case 1: // show
ShowWindow(wnd, SW_SHOW);
if (!isStartedFromCommandPrompt()) {
DeleteMenu(GetSystemMenu(wnd, FALSE), SC_CLOSE, MF_BYCOMMAND);
}
m_consoleStatus = 1;
break;
case GHOST_kConsoleWindowStateToggle:
case 2: // toggle
ShowWindow(wnd, m_consoleStatus ? SW_HIDE : SW_SHOW);
m_consoleStatus = !m_consoleStatus;
if (m_consoleStatus && !isStartedFromCommandPrompt()) {

View File

@@ -436,11 +436,16 @@ class GHOST_SystemWin32 : public GHOST_System {
static LRESULT WINAPI s_wndProc(HWND hwnd, UINT msg, WPARAM wParam, LPARAM lParam);
/**
* Set the Console State
* \param action: console state
* Toggles console
* \param action:
* - 0 - Hides
* - 1 - Shows
* - 2 - Toggles
* - 3 - Hides if it runs not from command line
* - * - Does nothing
* \return current status (1 -visible, 0 - hidden)
*/
int setConsoleWindowState(GHOST_TConsoleWindowState action);
int toggleConsole(int action);
/** The current state of the modifier keys. */
GHOST_ModifierKeys m_modifierKeys;

View File

@@ -13,9 +13,10 @@
* along with this program; if not, write to the Free Software Foundation,
* Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
*
* Copyright 2001-2002 NaN Holding BV. All rights reserved.
* 2009 Nokia Corporation and/or its subsidiary(-ies).
Part of this code has been taken from Qt, under LGPL license.
* The Original Code is Copyright (C) 2001-2002 by NaN Holding BV.
* All rights reserved.
* Part of this code has been taken from Qt, under LGPL license
* Copyright (C) 2009 Nokia Corporation and/or its subsidiary(-ies).
*/
/** \file

View File

@@ -269,7 +269,7 @@ class GHOST_SystemX11 : public GHOST_System {
/**
* \see GHOST_ISystem
*/
int setConsoleWindowState(GHOST_TConsoleWindowState /*action*/)
int toggleConsole(int /*action*/)
{
return 0;
}

View File

@@ -71,8 +71,6 @@ GHOST_WindowWin32::GHOST_WindowWin32(GHOST_SystemWin32 *system,
m_mousePresent(false),
m_inLiveResize(false),
m_system(system),
m_dropTarget(NULL),
m_hWnd(0),
m_hDC(0),
m_isDialog(dialog),
m_hasMouseCaptured(false),
@@ -80,7 +78,6 @@ GHOST_WindowWin32::GHOST_WindowWin32(GHOST_SystemWin32 *system,
m_nPressedButtons(0),
m_customCursor(0),
m_wantAlphaBackground(alphaBackground),
m_Bar(NULL),
m_wintab(NULL),
m_lastPointerTabletData(GHOST_TABLET_DATA_NONE),
m_normal_state(GHOST_kWindowStateNormal),
@@ -132,24 +129,8 @@ GHOST_WindowWin32::GHOST_WindowWin32(GHOST_SystemWin32 *system,
m_hDC = ::GetDC(m_hWnd);
if (!setDrawingContextType(type)) {
const char *title = "Blender - Unsupported Graphics Card Configuration";
const char *text =
"A graphics card and driver with support for OpenGL 3.3 or higher is "
"required.\n\nInstalling the latest driver for your graphics card might resolve the "
"issue.";
if (GetSystemMetrics(SM_CMONITORS) > 1) {
text =
"A graphics card and driver with support for OpenGL 3.3 or higher is "
"required.\n\nPlugging all monitors into your primary graphics card might resolve "
"this issue. Installing the latest driver for your graphics card could also help.";
}
MessageBox(m_hWnd, text, title, MB_OK | MB_ICONERROR);
::ReleaseDC(m_hWnd, m_hDC);
::DestroyWindow(m_hWnd);
m_hWnd = NULL;
if (!parentwindow) {
exit(0);
}
return;
}
@@ -583,12 +564,19 @@ GHOST_Context *GHOST_WindowWin32::newDrawingContext(GHOST_TDrawingContextType ty
(m_debug_context ? WGL_CONTEXT_DEBUG_BIT_ARB : 0),
GHOST_OPENGL_WGL_RESET_NOTIFICATION_STRATEGY);
if (context && !context->initializeDrawingContext()) {
delete context;
context = nullptr;
if (context->initializeDrawingContext()) {
return context;
}
else {
MessageBox(m_hWnd,
"A graphics card and driver with support for OpenGL 3.3 or higher is required.\n"
"Installing the latest driver for your graphics card may resolve the issue.\n\n"
"The program will now close.",
"Blender - Unsupported Graphics Card or Driver",
MB_OK | MB_ICONERROR);
delete context;
exit(0);
}
return context;
#elif defined(WITH_GL_PROFILE_COMPAT)
// ask for 2.1 context, driver gives any GL version >= 2.1

View File

@@ -86,7 +86,6 @@ void GHOST_XrContext::initialize(const GHOST_XrContextCreateInfo *create_info)
initApiLayers();
initExtensions();
if (isDebugMode()) {
printSDKVersion();
printAvailableAPILayersAndExtensionsInfo();
}
@@ -157,16 +156,6 @@ void GHOST_XrContext::storeInstanceProperties()
/** \name Debug Printing
* \{ */
void GHOST_XrContext::printSDKVersion()
{
const XrVersion sdk_version = XR_CURRENT_API_VERSION;
printf("OpenXR SDK Version: %u.%u.%u\n",
XR_VERSION_MAJOR(sdk_version),
XR_VERSION_MINOR(sdk_version),
XR_VERSION_PATCH(sdk_version));
}
void GHOST_XrContext::printInstanceInfo()
{
assert(m_oxr->instance != XR_NULL_HANDLE);

View File

@@ -126,7 +126,6 @@ class GHOST_XrContext : public GHOST_IXrContext {
void storeInstanceProperties();
void initDebugMessenger();
void printSDKVersion();
void printInstanceInfo();
void printAvailableAPILayersAndExtensionsInfo();
void printExtensionsAndAPILayersToEnable();

View File

@@ -292,7 +292,7 @@ target_link_libraries(multitest_c
guardedalloc_lib
wcwidth_lib
${OPENGL_gl_LIBRARY}
${FREETYPE_LIBRARIES} ${BROTLI_LIBRARIES}
${FREETYPE_LIBRARIES}
${ZLIB_LIBRARIES}
${CMAKE_DL_LIBS}
${PLATFORM_LINKLIBS}

View File

@@ -39,8 +39,8 @@
* second intern/ module with MEM_ prefix, for use in c++.
*
* \subsection memdependencies Dependencies
* - `stdlib`
* - `stdio`
* - stdlib
* - stdio
*
* \subsection memdocs API Documentation
* See \ref MEM_guardedalloc.h
@@ -268,12 +268,6 @@ void MEM_use_guarded_allocator(void);
* Allocate new memory for and constructs an object of type #T.
* #MEM_delete should be used to delete the object. Just calling #MEM_freeN is not enough when #T
* is not a trivial type.
*
* Note that when no arguments are passed, C++ will do recursive member-wise value initialization.
* That is because C++ differentiates between creating an object with `T` (default initialization)
* and `T()` (value initialization), whereby this function does the latter. Value initialization
* rules are complex, but for C-style structs, memory will be zero-initialized. So this doesn't
* match a `malloc()`, but a `calloc()` call in this case. See https://stackoverflow.com/a/4982720.
*/
template<typename T, typename... Args>
inline T *MEM_new(const char *allocation_name, Args &&...args)

View File

@@ -15,6 +15,7 @@
*
* The Original Code is Copyright (C) 2001-2002 by NaN Holding BV.
* All rights reserved.
* Original author: Laurence
*/
/** \file

View File

@@ -15,6 +15,7 @@
*
* The Original Code is Copyright (C) 2001-2002 by NaN Holding BV.
* All rights reserved.
* Original author: Laurence
*/
/** \file

View File

@@ -15,6 +15,7 @@
*
* The Original Code is Copyright (C) 2001-2002 by NaN Holding BV.
* All rights reserved.
* Original author: Laurence
*/
/** \file

View File

@@ -15,6 +15,7 @@
*
* The Original Code is Copyright (C) 2001-2002 by NaN Holding BV.
* All rights reserved.
* Original author: Laurence
*/
/** \file

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