Compare commits
1 Commits
tmp-eevee-
...
tmp-gpu-po
Author | SHA1 | Date | |
---|---|---|---|
7f5219caa4 |
@@ -264,9 +264,6 @@ ForEachMacros:
|
||||
- SET_SLOT_PROBING_BEGIN
|
||||
- MAP_SLOT_PROBING_BEGIN
|
||||
- VECTOR_SET_SLOT_PROBING_BEGIN
|
||||
- LIGHT_FOREACH_BEGIN_DIRECTIONAL
|
||||
- LIGHT_FOREACH_BEGIN_LOCAL
|
||||
- LIGHT_FOREACH_BEGIN_LOCAL_NO_CULL
|
||||
|
||||
StatementMacros:
|
||||
- PyObject_HEAD
|
||||
|
@@ -273,13 +273,11 @@ endif()
|
||||
|
||||
if(UNIX AND NOT APPLE)
|
||||
option(WITH_SYSTEM_GLEW "Use GLEW OpenGL wrapper library provided by the operating system" OFF)
|
||||
option(WITH_SYSTEM_GLEW "Use GLEW OpenGL wrapper library provided by the operating system" OFF)
|
||||
option(WITH_SYSTEM_FREETYPE "Use the freetype library provided by the operating system" OFF)
|
||||
option(WITH_SYSTEM_GLES "Use OpenGL ES library provided by the operating system" ON)
|
||||
else()
|
||||
# not an option for other OS's
|
||||
set(WITH_SYSTEM_GLEW OFF)
|
||||
set(WITH_SYSTEM_GLES OFF)
|
||||
set(WITH_SYSTEM_FREETYPE OFF)
|
||||
endif()
|
||||
|
||||
|
||||
@@ -710,12 +708,9 @@ if(UNIX AND NOT APPLE)
|
||||
endif()
|
||||
|
||||
# Installation process.
|
||||
set(POSTINSTALL_SCRIPT "" CACHE FILEPATH "Run given CMake script after installation process")
|
||||
option(POSTINSTALL_SCRIPT "Run given CMake script after installation process" OFF)
|
||||
mark_as_advanced(POSTINSTALL_SCRIPT)
|
||||
|
||||
set(POSTCONFIGURE_SCRIPT "" CACHE FILEPATH "Run given CMake script as the last step of CMake configuration")
|
||||
mark_as_advanced(POSTCONFIGURE_SCRIPT)
|
||||
|
||||
# end option(...)
|
||||
|
||||
|
||||
@@ -2079,8 +2074,3 @@ endif()
|
||||
if(0)
|
||||
print_all_vars()
|
||||
endif()
|
||||
|
||||
# Should be the last step of configuration.
|
||||
if(POSTCONFIGURE_SCRIPT)
|
||||
include(${POSTCONFIGURE_SCRIPT})
|
||||
endif()
|
||||
|
@@ -63,7 +63,6 @@ include(cmake/jpeg.cmake)
|
||||
include(cmake/blosc.cmake)
|
||||
include(cmake/pthreads.cmake)
|
||||
include(cmake/openexr.cmake)
|
||||
include(cmake/brotli.cmake)
|
||||
include(cmake/freetype.cmake)
|
||||
include(cmake/freeglut.cmake)
|
||||
include(cmake/glew.cmake)
|
||||
|
@@ -25,13 +25,8 @@ else()
|
||||
endif()
|
||||
|
||||
if(WIN32)
|
||||
if(MSVC_VERSION GREATER_EQUAL 1920) # 2019
|
||||
set(BOOST_TOOLSET toolset=msvc-14.2)
|
||||
set(BOOST_COMPILER_STRING -vc142)
|
||||
else() # 2017
|
||||
set(BOOST_TOOLSET toolset=msvc-14.1)
|
||||
set(BOOST_COMPILER_STRING -vc141)
|
||||
endif()
|
||||
set(BOOST_TOOLSET toolset=msvc-14.1)
|
||||
set(BOOST_COMPILER_STRING -vc141)
|
||||
|
||||
set(BOOST_CONFIGURE_COMMAND bootstrap.bat)
|
||||
set(BOOST_BUILD_COMMAND b2)
|
||||
|
@@ -1,38 +0,0 @@
|
||||
# ***** BEGIN GPL LICENSE BLOCK *****
|
||||
#
|
||||
# This program is free software; you can redistribute it and/or
|
||||
# modify it under the terms of the GNU General Public License
|
||||
# as published by the Free Software Foundation; either version 2
|
||||
# of the License, or (at your option) any later version.
|
||||
#
|
||||
# This program is distributed in the hope that it will be useful,
|
||||
# but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
# GNU General Public License for more details.
|
||||
#
|
||||
# You should have received a copy of the GNU General Public License
|
||||
# along with this program; if not, write to the Free Software Foundation,
|
||||
# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
|
||||
#
|
||||
# ***** END GPL LICENSE BLOCK *****
|
||||
|
||||
set(BROTLI_EXTRA_ARGS
|
||||
)
|
||||
|
||||
ExternalProject_Add(external_brotli
|
||||
URL file://${PACKAGE_DIR}/${BROTLI_FILE}
|
||||
DOWNLOAD_DIR ${DOWNLOAD_DIR}
|
||||
URL_HASH ${BROTLI_HASH_TYPE}=${BROTLI_HASH}
|
||||
PREFIX ${BUILD_DIR}/brotli
|
||||
CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/brotli ${DEFAULT_CMAKE_FLAGS} ${BROTLI_EXTRA_ARGS}
|
||||
INSTALL_DIR ${LIBDIR}/brotli
|
||||
)
|
||||
|
||||
if(BUILD_MODE STREQUAL Release AND WIN32)
|
||||
ExternalProject_Add_Step(external_brotli after_install
|
||||
COMMAND ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/brotli/include ${HARVEST_TARGET}/brotli/include
|
||||
COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/brotli/lib/brotlidec-static${LIBEXT} ${HARVEST_TARGET}/brotli/lib/brotlidec-static${LIBEXT}
|
||||
COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/brotli/lib/brotlicommon-static${LIBEXT} ${HARVEST_TARGET}/brotli/lib/brotlicommon-static${LIBEXT}
|
||||
DEPENDEES install
|
||||
)
|
||||
endif()
|
@@ -94,4 +94,3 @@ download_source(POTRACE)
|
||||
download_source(HARU)
|
||||
download_source(ZSTD)
|
||||
download_source(FLEX)
|
||||
download_source(BROTLI)
|
||||
|
@@ -23,12 +23,9 @@ set(FREETYPE_EXTRA_ARGS
|
||||
-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
|
||||
)
|
||||
-DCMAKE_DISABLE_FIND_PACKAGE_BrotliDec=TRUE)
|
||||
|
||||
ExternalProject_Add(external_freetype
|
||||
URL file://${PACKAGE_DIR}/${FREETYPE_FILE}
|
||||
@@ -39,11 +36,6 @@ ExternalProject_Add(external_freetype
|
||||
INSTALL_DIR ${LIBDIR}/freetype
|
||||
)
|
||||
|
||||
add_dependencies(
|
||||
external_freetype
|
||||
external_brotli
|
||||
)
|
||||
|
||||
if(BUILD_MODE STREQUAL Release AND WIN32)
|
||||
ExternalProject_Add_Step(external_freetype after_install
|
||||
COMMAND ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/freetype ${HARVEST_TARGET}/freetype
|
||||
|
@@ -79,8 +79,6 @@ endfunction()
|
||||
harvest(alembic/include alembic/include "*.h")
|
||||
harvest(alembic/lib/libAlembic.a alembic/lib/libAlembic.a)
|
||||
harvest(alembic/bin alembic/bin "*")
|
||||
harvest(brotli/include brotli/include "*.h")
|
||||
harvest(brotli/lib brotli/lib "*.a")
|
||||
harvest(boost/include boost/include "*")
|
||||
harvest(boost/lib boost/lib "*.a")
|
||||
harvest(ffmpeg/include ffmpeg/include "*.h")
|
||||
|
@@ -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)
|
||||
|
@@ -83,9 +83,9 @@ else()
|
||||
set(OPENEXR_VERSION_POSTFIX)
|
||||
endif()
|
||||
|
||||
set(FREETYPE_VERSION 2.11.0)
|
||||
set(FREETYPE_VERSION 2.10.2)
|
||||
set(FREETYPE_URI http://prdownloads.sourceforge.net/freetype/freetype-${FREETYPE_VERSION}.tar.gz)
|
||||
set(FREETYPE_HASH cf09172322f6b50cf8f568bf8fe14bde)
|
||||
set(FREETYPE_HASH b1cb620e4c875cd4d1bfa04945400945)
|
||||
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)
|
||||
|
||||
@@ -502,9 +500,3 @@ set(ZSTD_FILE zstd-${ZSTD_VERSION}.tar.gz)
|
||||
|
||||
set(SSE2NEON_GIT https://github.com/DLTcollab/sse2neon.git)
|
||||
set(SSE2NEON_GIT_HASH fe5ff00bb8d19b327714a3c290f3e2ce81ba3525)
|
||||
|
||||
set(BROTLI_VERSION v1.0.9)
|
||||
set(BROTLI_URI https://github.com/google/brotli/archive/refs/tags/${BROTLI_VERSION}.tar.gz)
|
||||
set(BROTLI_HASH f9e8d81d0405ba66d181529af42a3354f838c939095ff99930da6aa9cdf6fe46)
|
||||
set(BROTLI_HASH_TYPE SHA256)
|
||||
set(BROTLI_FILE brotli-${BROTLI_VERSION}.tar.gz)
|
||||
|
@@ -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"
|
||||
|
@@ -1,14 +1,21 @@
|
||||
@echo off
|
||||
if NOT "%1" == "" (
|
||||
if "%1" == "2017" (
|
||||
echo "Building for VS2017"
|
||||
set VSVER=15.0
|
||||
set VSVER_SHORT=15
|
||||
set BuildDir=VS15
|
||||
if "%1" == "2013" (
|
||||
echo "Building for VS2013"
|
||||
set VSVER=12.0
|
||||
set VSVER_SHORT=12
|
||||
set BuildDir=VS12
|
||||
goto par2
|
||||
)
|
||||
if "%1" == "2019" (
|
||||
echo "Building for VS2019"
|
||||
if "%1" == "2015" (
|
||||
echo "Building for VS2015"
|
||||
set VSVER=14.0
|
||||
set VSVER_SHORT=14
|
||||
set BuildDir=VS14
|
||||
goto par2
|
||||
)
|
||||
if "%1" == "2017" (
|
||||
echo "Building for VS2017"
|
||||
set VSVER=15.0
|
||||
set VSVER_SHORT=15
|
||||
set BuildDir=VS15
|
||||
@@ -18,22 +25,40 @@ if NOT "%1" == "" (
|
||||
)
|
||||
:usage
|
||||
|
||||
Echo Usage build_deps 2017/2019 x64
|
||||
Echo Usage build_deps 2013/2015/2017 x64/x86
|
||||
goto exit
|
||||
:par2
|
||||
if NOT "%2" == "" (
|
||||
if "%2" == "x86" (
|
||||
echo "Building for x86"
|
||||
set HARVESTROOT=Windows_vc
|
||||
set ARCH=86
|
||||
if "%1" == "2013" (
|
||||
set CMAKE_BUILDER=Visual Studio 12 2013
|
||||
)
|
||||
if "%1" == "2015" (
|
||||
set CMAKE_BUILDER=Visual Studio 14 2015
|
||||
)
|
||||
if "%1" == "2017" (
|
||||
set CMAKE_BUILDER=Visual Studio 15 2017
|
||||
)
|
||||
|
||||
goto start
|
||||
)
|
||||
if "%2" == "x64" (
|
||||
echo "Building for x64"
|
||||
set HARVESTROOT=Win64_vc
|
||||
set ARCH=64
|
||||
if "%1" == "2019" (
|
||||
set CMAKE_BUILDER=Visual Studio 16 2019
|
||||
set CMAKE_BUILD_ARCH=-A x64
|
||||
if "%1" == "2013" (
|
||||
set CMAKE_BUILDER=Visual Studio 12 2013 Win64
|
||||
)
|
||||
if "%1" == "2015" (
|
||||
set CMAKE_BUILDER=Visual Studio 14 2015 Win64
|
||||
)
|
||||
if "%1" == "2017" (
|
||||
set CMAKE_BUILDER=Visual Studio 15 2017 Win64
|
||||
set CMAKE_BUILD_ARCH=
|
||||
)
|
||||
|
||||
goto start
|
||||
)
|
||||
)
|
||||
@@ -95,7 +120,7 @@ set path=%BUILD_DIR%\downloads\mingw\mingw64\msys\1.0\bin\;%BUILD_DIR%\downloads
|
||||
mkdir %STAGING%\%BuildDir%%ARCH%R
|
||||
cd %Staging%\%BuildDir%%ARCH%R
|
||||
echo %DATE% %TIME% : Start > %StatusFile%
|
||||
cmake -G "%CMAKE_BUILDER%" %CMAKE_BUILD_ARCH% -Thost=x64 %SOURCE_DIR% -DPACKAGE_DIR=%BUILD_DIR%/packages -DDOWNLOAD_DIR=%BUILD_DIR%/downloads -DBUILD_MODE=Release -DHARVEST_TARGET=%HARVEST_DIR%/%HARVESTROOT%%VSVER_SHORT%/
|
||||
cmake -G "%CMAKE_BUILDER%" -Thost=x64 %SOURCE_DIR% -DPACKAGE_DIR=%BUILD_DIR%/packages -DDOWNLOAD_DIR=%BUILD_DIR%/downloads -DBUILD_MODE=Release -DHARVEST_TARGET=%HARVEST_DIR%/%HARVESTROOT%%VSVER_SHORT%/
|
||||
echo %DATE% %TIME% : Release Configuration done >> %StatusFile%
|
||||
if "%dobuild%" == "1" (
|
||||
msbuild /m "ll.vcxproj" /p:Configuration=Release /fl /flp:logfile=BlenderDeps_llvm.log;Verbosity=normal
|
||||
@@ -108,7 +133,7 @@ if "%NODEBUG%" == "1" goto exit
|
||||
cd %BUILD_DIR%
|
||||
mkdir %STAGING%\%BuildDir%%ARCH%D
|
||||
cd %Staging%\%BuildDir%%ARCH%D
|
||||
cmake -G "%CMAKE_BUILDER%" %CMAKE_BUILD_ARCH% -Thost=x64 %SOURCE_DIR% -DPACKAGE_DIR=%BUILD_DIR%/packages -DDOWNLOAD_DIR=%BUILD_DIR%/downloads -DCMAKE_BUILD_TYPE=Debug -DBUILD_MODE=Debug -DHARVEST_TARGET=%HARVEST_DIR%/%HARVESTROOT%%VSVER_SHORT%/ %CMAKE_DEBUG_OPTIONS%
|
||||
cmake -G "%CMAKE_BUILDER%" -Thost=x64 %SOURCE_DIR% -DPACKAGE_DIR=%BUILD_DIR%/packages -DDOWNLOAD_DIR=%BUILD_DIR%/downloads -DCMAKE_BUILD_TYPE=Debug -DBUILD_MODE=Debug -DHARVEST_TARGET=%HARVEST_DIR%/%HARVESTROOT%%VSVER_SHORT%/ %CMAKE_DEBUG_OPTIONS%
|
||||
echo %DATE% %TIME% : Debug Configuration done >> %StatusFile%
|
||||
if "%dobuild%" == "1" (
|
||||
msbuild /m "ll.vcxproj" /p:Configuration=Debug /fl /flp:logfile=BlenderDeps_llvm.log;;Verbosity=normal
|
||||
|
@@ -1,83 +0,0 @@
|
||||
# - Find Brotli library (compression for freetype/woff2).
|
||||
# This module defines
|
||||
# BROTLI_INCLUDE_DIRS, where to find Brotli headers, Set when
|
||||
# BROTLI_INCLUDE_DIR is found.
|
||||
# BROTLI_LIBRARIES, libraries to link against to use Brotli.
|
||||
# BROTLI_ROOT_DIR, The base directory to search for Brotli.
|
||||
# This can also be an environment variable.
|
||||
# BROTLI_FOUND, If false, do not try to use Brotli.
|
||||
#
|
||||
|
||||
#=============================================================================
|
||||
# Copyright 2022 Blender Foundation.
|
||||
#
|
||||
# Distributed under the OSI-approved BSD 3-Clause License,
|
||||
# see accompanying file BSD-3-Clause-license.txt for details.
|
||||
#=============================================================================
|
||||
|
||||
# If BROTLI_ROOT_DIR was defined in the environment, use it.
|
||||
IF(NOT BROTLI_ROOT_DIR AND NOT $ENV{BROTLI_ROOT_DIR} STREQUAL "")
|
||||
SET(BROTLI_ROOT_DIR $ENV{BROTLI_ROOT_DIR})
|
||||
ENDIF()
|
||||
|
||||
SET(_BROTLI_SEARCH_DIRS
|
||||
${BROTLI_ROOT_DIR}
|
||||
)
|
||||
|
||||
FIND_PATH(BROTLI_INCLUDE_DIR
|
||||
NAMES
|
||||
brotli/decode.h
|
||||
HINTS
|
||||
${_BROTLI_SEARCH_DIRS}
|
||||
PATH_SUFFIXES
|
||||
include
|
||||
DOC "Brotli header files"
|
||||
)
|
||||
|
||||
FIND_LIBRARY(BROTLI_LIBRARY_COMMON
|
||||
NAMES
|
||||
# Some builds use a special `-static` postfix in their static libraries names.
|
||||
brotlicommon-static
|
||||
brotlicommon
|
||||
HINTS
|
||||
${_BROTLI_SEARCH_DIRS}
|
||||
PATH_SUFFIXES
|
||||
lib64 lib lib/static
|
||||
DOC "Brotli static common library"
|
||||
)
|
||||
FIND_LIBRARY(BROTLI_LIBRARY_DEC
|
||||
NAMES
|
||||
# Some builds use a special `-static` postfix in their static libraries names.
|
||||
brotlidec-static
|
||||
brotlidec
|
||||
HINTS
|
||||
${_BROTLI_SEARCH_DIRS}
|
||||
PATH_SUFFIXES
|
||||
lib64 lib lib/static
|
||||
DOC "Brotli static decode library"
|
||||
)
|
||||
|
||||
|
||||
IF(${BROTLI_LIBRARY_COMMON_NOTFOUND} or ${BROTLI_LIBRARY_DEC_NOTFOUND})
|
||||
set(BROTLI_FOUND FALSE)
|
||||
ELSE()
|
||||
# handle the QUIETLY and REQUIRED arguments and set BROTLI_FOUND to TRUE if
|
||||
# all listed variables are TRUE
|
||||
INCLUDE(FindPackageHandleStandardArgs)
|
||||
FIND_PACKAGE_HANDLE_STANDARD_ARGS(Brotli DEFAULT_MSG BROTLI_LIBRARY_COMMON BROTLI_LIBRARY_DEC BROTLI_INCLUDE_DIR)
|
||||
|
||||
IF(BROTLI_FOUND)
|
||||
get_filename_component(BROTLI_LIBRARY_DIR ${BROTLI_LIBRARY_COMMON} DIRECTORY)
|
||||
SET(BROTLI_INCLUDE_DIRS ${BROTLI_INCLUDE_DIR})
|
||||
SET(BROTLI_LIBRARIES ${BROTLI_LIBRARY_DEC} ${BROTLI_LIBRARY_COMMON})
|
||||
ENDIF()
|
||||
ENDIF()
|
||||
|
||||
MARK_AS_ADVANCED(
|
||||
BROTLI_INCLUDE_DIR
|
||||
BROTLI_LIBRARY_COMMON
|
||||
BROTLI_LIBRARY_DEC
|
||||
BROTLI_LIBRARY_DIR
|
||||
)
|
||||
|
||||
UNSET(_BROTLI_SEARCH_DIRS)
|
@@ -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()
|
||||
|
||||
@@ -161,11 +166,7 @@ if(WITH_FFTW3)
|
||||
find_package(Fftw3)
|
||||
endif()
|
||||
|
||||
# FreeType compiled with Brotli compression for woff2.
|
||||
find_package(Freetype REQUIRED)
|
||||
list(APPEND FREETYPE_LIBRARIES
|
||||
${LIBDIR}/brotli/lib/libbrotlicommon-static.a
|
||||
${LIBDIR}/brotli/lib/libbrotlidec-static.a)
|
||||
|
||||
if(WITH_IMAGE_OPENEXR)
|
||||
find_package(OpenEXR)
|
||||
|
@@ -48,9 +48,6 @@ if(NOT DEFINED LIBDIR)
|
||||
unset(LIBDIR_CENTOS7_ABI)
|
||||
endif()
|
||||
|
||||
# Support restoring this value once pre-compiled libraries have been handled.
|
||||
set(WITH_STATIC_LIBS_INIT ${WITH_STATIC_LIBS})
|
||||
|
||||
if(EXISTS ${LIBDIR})
|
||||
message(STATUS "Using pre-compiled LIBDIR: ${LIBDIR}")
|
||||
|
||||
@@ -103,22 +100,7 @@ find_package_wrapper(JPEG REQUIRED)
|
||||
find_package_wrapper(PNG REQUIRED)
|
||||
find_package_wrapper(ZLIB REQUIRED)
|
||||
find_package_wrapper(Zstd REQUIRED)
|
||||
|
||||
if(NOT WITH_SYSTEM_FREETYPE)
|
||||
# FreeType compiled with Brotli compression for woff2.
|
||||
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 `FREETYPE_LIBRARIES`.
|
||||
#
|
||||
# list(APPEND FREETYPE_LIBRARIES
|
||||
# ${BROTLI_LIBRARIES}
|
||||
# )
|
||||
endif()
|
||||
endif()
|
||||
find_package_wrapper(Freetype REQUIRED)
|
||||
|
||||
if(WITH_PYTHON)
|
||||
# No way to set py35, remove for now.
|
||||
@@ -554,21 +536,6 @@ add_definitions(-D_LARGEFILE_SOURCE -D_FILE_OFFSET_BITS=64 -D_LARGEFILE64_SOURCE
|
||||
#
|
||||
# Keep last, so indirectly linked libraries don't override our own pre-compiled libs.
|
||||
|
||||
if(EXISTS ${LIBDIR})
|
||||
# Clear the prefix path as it causes the `LIBDIR` to override system locations.
|
||||
unset(CMAKE_PREFIX_PATH)
|
||||
|
||||
# Since the pre-compiled `LIBDIR` directories have been handled, don't prefer static libraries.
|
||||
set(WITH_STATIC_LIBS ${WITH_STATIC_LIBS_INIT})
|
||||
endif()
|
||||
|
||||
if(WITH_SYSTEM_FREETYPE)
|
||||
find_package_wrapper(Freetype)
|
||||
if(NOT FREETYPE_FOUND)
|
||||
message(FATAL_ERROR "Failed finding system FreeType version!")
|
||||
endif()
|
||||
endif()
|
||||
|
||||
if(WITH_LZO AND WITH_SYSTEM_LZO)
|
||||
find_package_wrapper(LZO)
|
||||
if(NOT LZO_FOUND)
|
||||
|
@@ -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}")
|
||||
@@ -345,11 +347,7 @@ set(FREETYPE_INCLUDE_DIRS
|
||||
${LIBDIR}/freetype/include
|
||||
${LIBDIR}/freetype/include/freetype2
|
||||
)
|
||||
set(FREETYPE_LIBRARIES
|
||||
${LIBDIR}/freetype/lib/freetype2ST.lib
|
||||
${LIBDIR}/brotli/lib/brotlidec-static.lib
|
||||
${LIBDIR}/brotli/lib/brotlicommon-static.lib
|
||||
)
|
||||
set(FREETYPE_LIBRARY ${LIBDIR}/freetype/lib/freetype2ST.lib)
|
||||
windows_find_package(freetype REQUIRED)
|
||||
|
||||
if(WITH_FFTW3)
|
||||
@@ -463,7 +461,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)
|
||||
|
@@ -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
|
||||
|
||||
|
@@ -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
|
||||
|
||||
|
@@ -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 (
|
||||
|
@@ -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" (
|
||||
|
3
build_files/windows/detect_msvc2017.cmd
Normal file
3
build_files/windows/detect_msvc2017.cmd
Normal file
@@ -0,0 +1,3 @@
|
||||
set BUILD_VS_VER=15
|
||||
set BUILD_VS_YEAR=2017
|
||||
call "%~dp0\detect_msvc_vswhere.cmd"
|
@@ -3,32 +3,7 @@ for %%X in (svn.exe) do (set SVN=%%~$PATH:X)
|
||||
for %%X in (cmake.exe) do (set CMAKE=%%~$PATH:X)
|
||||
for %%X in (ctest.exe) do (set CTEST=%%~$PATH:X)
|
||||
for %%X in (git.exe) do (set GIT=%%~$PATH:X)
|
||||
REM For python, default on 39 but if that does not exist also check
|
||||
REM the 310,311 and 312 folders to see if those are there, it checks
|
||||
REM this far ahead to ensure good lib folder compatiblity in the future.
|
||||
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
|
||||
if EXIST %PYTHON% (
|
||||
goto detect_python_done
|
||||
)
|
||||
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\310\bin\python.exe
|
||||
if EXIST %PYTHON% (
|
||||
goto detect_python_done
|
||||
)
|
||||
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\311\bin\python.exe
|
||||
if EXIST %PYTHON% (
|
||||
goto detect_python_done
|
||||
)
|
||||
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\312\bin\python.exe
|
||||
if EXIST %PYTHON% (
|
||||
goto detect_python_done
|
||||
)
|
||||
|
||||
if NOT EXIST %PYTHON% (
|
||||
echo Warning: Python not found, there is likely an issue with the library folder
|
||||
set PYTHON=""
|
||||
)
|
||||
|
||||
:detect_python_done
|
||||
if NOT "%verbose%" == "" (
|
||||
echo svn : "%SVN%"
|
||||
echo cmake : "%CMAKE%"
|
||||
@@ -36,3 +11,7 @@ if NOT "%verbose%" == "" (
|
||||
echo git : "%GIT%"
|
||||
echo python : "%PYTHON%"
|
||||
)
|
||||
if "%CMAKE%" == "" (
|
||||
echo Cmake not found in path, required for building, exiting...
|
||||
exit /b 1
|
||||
)
|
||||
|
@@ -9,11 +9,17 @@ exit /b 1
|
||||
:detect_done
|
||||
echo found clang-format in %CF_PATH%
|
||||
|
||||
if NOT EXIST %PYTHON% (
|
||||
echo python not found, required for this operation
|
||||
exit /b 1
|
||||
if EXIST %PYTHON% (
|
||||
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
|
||||
goto detect_python_done
|
||||
)
|
||||
|
||||
echo python not found in lib folder
|
||||
exit /b 1
|
||||
|
||||
:detect_python_done
|
||||
echo found python (%PYTHON%)
|
||||
|
||||
set FORMAT_PATHS=%BLENDER_DIR%\source\tools\utils_maintenance\clang_format_paths.py
|
||||
|
||||
REM The formatting script expects clang-format to be in the current PATH.
|
||||
|
@@ -1,8 +1,18 @@
|
||||
if NOT EXIST %PYTHON% (
|
||||
echo python not found, required for this operation
|
||||
exit /b 1
|
||||
if EXIST "%PYTHON%" (
|
||||
goto detect_python_done
|
||||
)
|
||||
|
||||
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
|
||||
if EXIST %PYTHON% (
|
||||
goto detect_python_done
|
||||
)
|
||||
|
||||
echo python not found at %PYTHON%
|
||||
exit /b 1
|
||||
|
||||
:detect_python_done
|
||||
echo found python (%PYTHON%)
|
||||
|
||||
call "%~dp0\find_inkscape.cmd"
|
||||
|
||||
if EXIST "%INKSCAPE_BIN%" (
|
||||
|
@@ -1,8 +1,18 @@
|
||||
if NOT EXIST %PYTHON% (
|
||||
echo python not found, required for this operation
|
||||
exit /b 1
|
||||
if EXIST %PYTHON% (
|
||||
goto detect_python_done
|
||||
)
|
||||
|
||||
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
|
||||
if EXIST %PYTHON% (
|
||||
goto detect_python_done
|
||||
)
|
||||
|
||||
echo python not found at %PYTHON%
|
||||
exit /b 1
|
||||
|
||||
:detect_python_done
|
||||
echo found python (%PYTHON%)
|
||||
|
||||
call "%~dp0\find_blender.cmd"
|
||||
|
||||
if EXIST "%BLENDER_BIN%" (
|
||||
|
@@ -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" (
|
||||
|
@@ -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^)
|
||||
|
@@ -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
|
||||
|
||||
|
@@ -1,7 +1,10 @@
|
||||
if NOT EXIST %PYTHON% (
|
||||
echo python not found, required for this operation
|
||||
exit /b 1
|
||||
if EXIST %PYTHON% (
|
||||
goto detect_python_done
|
||||
)
|
||||
|
||||
echo python not found in lib folder
|
||||
exit /b 1
|
||||
|
||||
:detect_python_done
|
||||
|
||||
REM Use -B to avoid writing __pycache__ in lib directory and causing update conflicts.
|
||||
|
@@ -743,7 +743,7 @@ will re-allocate objects data,
|
||||
any references to a meshes vertices/polygons/UVs, armatures bones,
|
||||
curves points, etc. cannot be accessed after switching mode.
|
||||
|
||||
Only the reference to the data itself can be re-accessed, the following example will crash.
|
||||
Only the reference to the data its self can be re-accessed, the following example will crash.
|
||||
|
||||
.. code-block:: python
|
||||
|
||||
|
@@ -1762,7 +1762,6 @@ except ModuleNotFoundError:
|
||||
fw("html_show_sphinx = False\n")
|
||||
fw("html_baseurl = 'https://docs.blender.org/api/current/'\n")
|
||||
fw("html_use_opensearch = 'https://docs.blender.org/api/current'\n")
|
||||
fw("html_show_search_summary = True\n")
|
||||
fw("html_split_index = True\n")
|
||||
fw("html_static_path = ['static']\n")
|
||||
fw("html_extra_path = ['static/favicon.ico', 'static/blender_logo.svg']\n")
|
||||
|
@@ -667,11 +667,6 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
|
||||
description="Use special type BVH optimized for hair (uses more ram but renders faster)",
|
||||
default=True,
|
||||
)
|
||||
debug_use_compact_bvh: BoolProperty(
|
||||
name="Use Compact BVH",
|
||||
description="Use compact BVH structure (uses less ram but renders slower)",
|
||||
default=True,
|
||||
)
|
||||
debug_bvh_time_steps: IntProperty(
|
||||
name="BVH Time Steps",
|
||||
description="Split BVH primitives by this number of time steps to speed up render time in cost of memory",
|
||||
@@ -1452,19 +1447,6 @@ class CyclesPreferences(bpy.types.AddonPreferences):
|
||||
num += 1
|
||||
return num
|
||||
|
||||
def has_multi_device(self):
|
||||
import _cycles
|
||||
compute_device_type = self.get_compute_device_type()
|
||||
device_list = _cycles.available_devices(compute_device_type)
|
||||
for device in device_list:
|
||||
if device[1] == compute_device_type:
|
||||
continue
|
||||
for dev in self.devices:
|
||||
if dev.use and dev.id == device[2]:
|
||||
return True
|
||||
|
||||
return False
|
||||
|
||||
def has_active_device(self):
|
||||
return self.get_num_gpu_devices() > 0
|
||||
|
||||
|
@@ -118,11 +118,11 @@ def use_optix(context):
|
||||
|
||||
return (get_device_type(context) == 'OPTIX' and cscene.device == 'GPU')
|
||||
|
||||
def use_multi_device(context):
|
||||
|
||||
def use_sample_all_lights(context):
|
||||
cscene = context.scene.cycles
|
||||
if cscene.device != 'GPU':
|
||||
return False
|
||||
return context.preferences.addons[__package__].preferences.has_multi_device()
|
||||
|
||||
return cscene.sample_all_lights_direct or cscene.sample_all_lights_indirect
|
||||
|
||||
|
||||
def show_device_active(context):
|
||||
@@ -667,10 +667,6 @@ class CYCLES_RENDER_PT_performance_acceleration_structure(CyclesButtonsPanel, Pa
|
||||
bl_label = "Acceleration Structure"
|
||||
bl_parent_id = "CYCLES_RENDER_PT_performance"
|
||||
|
||||
@classmethod
|
||||
def poll(cls, context):
|
||||
return not use_optix(context) or has_multi_device(context)
|
||||
|
||||
def draw(self, context):
|
||||
import _cycles
|
||||
|
||||
@@ -683,33 +679,21 @@ class CYCLES_RENDER_PT_performance_acceleration_structure(CyclesButtonsPanel, Pa
|
||||
|
||||
col = layout.column()
|
||||
|
||||
use_embree = _cycles.with_embree
|
||||
|
||||
use_embree = False
|
||||
if use_cpu(context):
|
||||
col.prop(cscene, "debug_use_spatial_splits")
|
||||
if use_embree:
|
||||
col.prop(cscene, "debug_use_compact_bvh")
|
||||
else:
|
||||
sub = col.column()
|
||||
sub.active = not cscene.debug_use_spatial_splits
|
||||
sub.prop(cscene, "debug_bvh_time_steps")
|
||||
|
||||
col.prop(cscene, "debug_use_hair_bvh")
|
||||
|
||||
use_embree = _cycles.with_embree
|
||||
if not use_embree:
|
||||
sub = col.column(align=True)
|
||||
sub.label(text="Cycles built without Embree support")
|
||||
sub.label(text="CPU raytracing performance will be poor")
|
||||
else:
|
||||
col.prop(cscene, "debug_use_spatial_splits")
|
||||
sub = col.column()
|
||||
sub.active = not cscene.debug_use_spatial_splits
|
||||
sub.prop(cscene, "debug_bvh_time_steps")
|
||||
|
||||
col.prop(cscene, "debug_use_hair_bvh")
|
||||
|
||||
# CPU is used in addition to a GPU
|
||||
if use_multi_device(context) and use_embree:
|
||||
col.prop(cscene, "debug_use_compact_bvh")
|
||||
col.prop(cscene, "debug_use_spatial_splits")
|
||||
sub = col.column()
|
||||
sub.active = not use_embree
|
||||
sub.prop(cscene, "debug_use_hair_bvh")
|
||||
sub = col.column()
|
||||
sub.active = not cscene.debug_use_spatial_splits and not use_embree
|
||||
sub.prop(cscene, "debug_bvh_time_steps")
|
||||
|
||||
|
||||
class CYCLES_RENDER_PT_performance_final_render(CyclesButtonsPanel, Panel):
|
||||
|
@@ -1071,15 +1071,7 @@ static void create_subd_mesh(Scene *scene,
|
||||
|
||||
for (BL::MeshEdge &e : b_mesh.edges) {
|
||||
if (e.crease() != 0.0f) {
|
||||
mesh->add_edge_crease(e.vertices()[0], e.vertices()[1], e.crease());
|
||||
}
|
||||
}
|
||||
|
||||
for (BL::MeshVertexCreaseLayer &c : b_mesh.vertex_creases) {
|
||||
for (int i = 0; i < c.data.length(); ++i) {
|
||||
if (c.data[i].value() != 0.0f) {
|
||||
mesh->add_vertex_crease(i, c.data[i].value());
|
||||
}
|
||||
mesh->add_crease(e.vertices()[0], e.vertices()[1], e.crease());
|
||||
}
|
||||
}
|
||||
|
||||
|
@@ -51,6 +51,8 @@ bool BlenderOutputDriver::read_render_tile(const Tile &tile)
|
||||
|
||||
BL::RenderLayer b_rlay = *b_single_rlay;
|
||||
|
||||
vector<float> pixels(tile.size.x * tile.size.y * 4);
|
||||
|
||||
/* Copy each pass.
|
||||
* TODO:copy only the required ones for better performance? */
|
||||
for (BL::RenderPass &b_pass : b_rlay.passes) {
|
||||
@@ -107,7 +109,7 @@ void BlenderOutputDriver::write_render_tile(const Tile &tile)
|
||||
|
||||
BL::RenderLayer b_rlay = *b_single_rlay;
|
||||
|
||||
vector<float> pixels(static_cast<size_t>(tile.size.x) * tile.size.y * 4);
|
||||
vector<float> pixels(tile.size.x * tile.size.y * 4);
|
||||
|
||||
/* Copy each pass. */
|
||||
for (BL::RenderPass &b_pass : b_rlay.passes) {
|
||||
|
@@ -689,9 +689,6 @@ static ShaderNode *add_node(Scene *scene,
|
||||
else if (b_node.is_a(&RNA_ShaderNodeHairInfo)) {
|
||||
node = graph->create_node<HairInfoNode>();
|
||||
}
|
||||
else if (b_node.is_a(&RNA_ShaderNodePointInfo)) {
|
||||
node = graph->create_node<PointInfoNode>();
|
||||
}
|
||||
else if (b_node.is_a(&RNA_ShaderNodeVolumeInfo)) {
|
||||
node = graph->create_node<VolumeInfoNode>();
|
||||
}
|
||||
|
@@ -787,7 +787,6 @@ SceneParams BlenderSync::get_scene_params(BL::Scene &b_scene, bool background)
|
||||
params.bvh_type = BVH_TYPE_DYNAMIC;
|
||||
|
||||
params.use_bvh_spatial_split = RNA_boolean_get(&cscene, "debug_use_spatial_splits");
|
||||
params.use_bvh_compact_structure = RNA_boolean_get(&cscene, "debug_use_compact_bvh");
|
||||
params.use_bvh_unaligned_nodes = RNA_boolean_get(&cscene, "debug_use_hair_bvh");
|
||||
params.num_bvh_time_steps = RNA_int_get(&cscene, "debug_bvh_time_steps");
|
||||
|
||||
|
@@ -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, ¤t_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, ¤t_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)
|
||||
@@ -399,12 +355,10 @@ void BVHEmbree::build(Progress &progress, Stats *stats, RTCDevice rtc_device_)
|
||||
}
|
||||
|
||||
const bool dynamic = params.bvh_type == BVH_TYPE_DYNAMIC;
|
||||
const bool compact = params.use_compact_structure;
|
||||
|
||||
scene = rtcNewScene(rtc_device);
|
||||
const RTCSceneFlags scene_flags = (dynamic ? RTC_SCENE_FLAG_DYNAMIC : RTC_SCENE_FLAG_NONE) |
|
||||
(compact ? RTC_SCENE_FLAG_COMPACT : RTC_SCENE_FLAG_NONE) |
|
||||
RTC_SCENE_FLAG_ROBUST;
|
||||
RTC_SCENE_FLAG_COMPACT | RTC_SCENE_FLAG_ROBUST;
|
||||
rtcSetSceneFlags(scene, scene_flags);
|
||||
build_quality = dynamic ? RTC_BUILD_QUALITY_LOW :
|
||||
(params.use_spatial_split ? RTC_BUILD_QUALITY_HIGH :
|
||||
@@ -549,7 +503,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 +765,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 {
|
||||
|
@@ -97,9 +97,6 @@ class BVHParams {
|
||||
*/
|
||||
bool use_unaligned_nodes;
|
||||
|
||||
/* Use compact acceleration structure (Embree)*/
|
||||
bool use_compact_structure;
|
||||
|
||||
/* Split time range to this number of steps and create leaf node for each
|
||||
* of this time steps.
|
||||
*
|
||||
|
@@ -58,11 +58,6 @@ class BVHMetal : public BVH {
|
||||
id<MTLCommandQueue> queue,
|
||||
Geometry *const geom,
|
||||
bool refit);
|
||||
bool build_BLAS_pointcloud(Progress &progress,
|
||||
id<MTLDevice> device,
|
||||
id<MTLCommandQueue> queue,
|
||||
Geometry *const geom,
|
||||
bool refit);
|
||||
bool build_TLAS(Progress &progress, id<MTLDevice> device, id<MTLCommandQueue> queue, bool refit);
|
||||
};
|
||||
|
||||
|
@@ -19,7 +19,6 @@
|
||||
# include "scene/hair.h"
|
||||
# include "scene/mesh.h"
|
||||
# include "scene/object.h"
|
||||
# include "scene/pointcloud.h"
|
||||
|
||||
# include "util/progress.h"
|
||||
|
||||
@@ -476,220 +475,6 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
|
||||
return false;
|
||||
}
|
||||
|
||||
bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
|
||||
id<MTLDevice> device,
|
||||
id<MTLCommandQueue> queue,
|
||||
Geometry *const geom,
|
||||
bool refit)
|
||||
{
|
||||
if (@available(macos 12.0, *)) {
|
||||
/* Build BLAS for point cloud */
|
||||
PointCloud *pointcloud = static_cast<PointCloud *>(geom);
|
||||
if (pointcloud->num_points() == 0) {
|
||||
return false;
|
||||
}
|
||||
|
||||
/*------------------------------------------------*/
|
||||
BVH_status("Building pointcloud BLAS | %7d points | %s",
|
||||
(int)pointcloud->num_points(),
|
||||
geom->name.c_str());
|
||||
/*------------------------------------------------*/
|
||||
|
||||
const size_t num_points = pointcloud->get_points().size();
|
||||
const float3 *points = pointcloud->get_points().data();
|
||||
const float *radius = pointcloud->get_radius().data();
|
||||
|
||||
const bool use_fast_trace_bvh = (params.bvh_type == BVH_TYPE_STATIC);
|
||||
|
||||
size_t num_motion_steps = 1;
|
||||
Attribute *motion_keys = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
|
||||
if (motion_blur && pointcloud->get_use_motion_blur() && motion_keys) {
|
||||
num_motion_steps = pointcloud->get_motion_steps();
|
||||
}
|
||||
|
||||
const size_t num_aabbs = num_motion_steps;
|
||||
|
||||
MTLResourceOptions storage_mode;
|
||||
if (device.hasUnifiedMemory) {
|
||||
storage_mode = MTLResourceStorageModeShared;
|
||||
}
|
||||
else {
|
||||
storage_mode = MTLResourceStorageModeManaged;
|
||||
}
|
||||
|
||||
/* Allocate a GPU buffer for the AABB data and populate it */
|
||||
id<MTLBuffer> aabbBuf = [device
|
||||
newBufferWithLength:num_aabbs * sizeof(MTLAxisAlignedBoundingBox)
|
||||
options:storage_mode];
|
||||
MTLAxisAlignedBoundingBox *aabb_data = (MTLAxisAlignedBoundingBox *)[aabbBuf contents];
|
||||
|
||||
/* Get AABBs for each motion step */
|
||||
size_t center_step = (num_motion_steps - 1) / 2;
|
||||
for (size_t step = 0; step < num_motion_steps; ++step) {
|
||||
/* The center step for motion vertices is not stored in the attribute */
|
||||
if (step != center_step) {
|
||||
size_t attr_offset = (step > center_step) ? step - 1 : step;
|
||||
points = motion_keys->data_float3() + attr_offset * num_points;
|
||||
}
|
||||
|
||||
for (size_t j = 0; j < num_points; ++j) {
|
||||
const PointCloud::Point point = pointcloud->get_point(j);
|
||||
BoundBox bounds = BoundBox::empty;
|
||||
point.bounds_grow(points, radius, bounds);
|
||||
|
||||
const size_t index = step * num_points + j;
|
||||
aabb_data[index].min = (MTLPackedFloat3 &)bounds.min;
|
||||
aabb_data[index].max = (MTLPackedFloat3 &)bounds.max;
|
||||
}
|
||||
}
|
||||
|
||||
if (storage_mode == MTLResourceStorageModeManaged) {
|
||||
[aabbBuf didModifyRange:NSMakeRange(0, aabbBuf.length)];
|
||||
}
|
||||
|
||||
# if 0
|
||||
for (size_t i=0; i<num_aabbs && i < 400; i++) {
|
||||
MTLAxisAlignedBoundingBox& bb = aabb_data[i];
|
||||
printf(" %d: %.1f,%.1f,%.1f -- %.1f,%.1f,%.1f\n", int(i), bb.min.x, bb.min.y, bb.min.z, bb.max.x, bb.max.y, bb.max.z);
|
||||
}
|
||||
# endif
|
||||
|
||||
MTLAccelerationStructureGeometryDescriptor *geomDesc;
|
||||
if (motion_blur) {
|
||||
std::vector<MTLMotionKeyframeData *> aabb_ptrs;
|
||||
aabb_ptrs.reserve(num_motion_steps);
|
||||
for (size_t step = 0; step < num_motion_steps; ++step) {
|
||||
MTLMotionKeyframeData *k = [MTLMotionKeyframeData data];
|
||||
k.buffer = aabbBuf;
|
||||
k.offset = step * num_points * sizeof(MTLAxisAlignedBoundingBox);
|
||||
aabb_ptrs.push_back(k);
|
||||
}
|
||||
|
||||
MTLAccelerationStructureMotionBoundingBoxGeometryDescriptor *geomDescMotion =
|
||||
[MTLAccelerationStructureMotionBoundingBoxGeometryDescriptor descriptor];
|
||||
geomDescMotion.boundingBoxBuffers = [NSArray arrayWithObjects:aabb_ptrs.data()
|
||||
count:aabb_ptrs.size()];
|
||||
geomDescMotion.boundingBoxCount = num_points;
|
||||
geomDescMotion.boundingBoxStride = sizeof(aabb_data[0]);
|
||||
geomDescMotion.intersectionFunctionTableOffset = 2;
|
||||
|
||||
/* Force a single any-hit call, so shadow record-all behavior works correctly */
|
||||
/* (Match optix behavior: unsigned int build_flags =
|
||||
* OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;) */
|
||||
geomDescMotion.allowDuplicateIntersectionFunctionInvocation = false;
|
||||
geomDescMotion.opaque = true;
|
||||
geomDesc = geomDescMotion;
|
||||
}
|
||||
else {
|
||||
MTLAccelerationStructureBoundingBoxGeometryDescriptor *geomDescNoMotion =
|
||||
[MTLAccelerationStructureBoundingBoxGeometryDescriptor descriptor];
|
||||
geomDescNoMotion.boundingBoxBuffer = aabbBuf;
|
||||
geomDescNoMotion.boundingBoxBufferOffset = 0;
|
||||
geomDescNoMotion.boundingBoxCount = int(num_aabbs);
|
||||
geomDescNoMotion.boundingBoxStride = sizeof(aabb_data[0]);
|
||||
geomDescNoMotion.intersectionFunctionTableOffset = 2;
|
||||
|
||||
/* Force a single any-hit call, so shadow record-all behavior works correctly */
|
||||
/* (Match optix behavior: unsigned int build_flags =
|
||||
* OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;) */
|
||||
geomDescNoMotion.allowDuplicateIntersectionFunctionInvocation = false;
|
||||
geomDescNoMotion.opaque = true;
|
||||
geomDesc = geomDescNoMotion;
|
||||
}
|
||||
|
||||
MTLPrimitiveAccelerationStructureDescriptor *accelDesc =
|
||||
[MTLPrimitiveAccelerationStructureDescriptor descriptor];
|
||||
accelDesc.geometryDescriptors = @[ geomDesc ];
|
||||
|
||||
if (motion_blur) {
|
||||
accelDesc.motionStartTime = 0.0f;
|
||||
accelDesc.motionEndTime = 1.0f;
|
||||
accelDesc.motionStartBorderMode = MTLMotionBorderModeVanish;
|
||||
accelDesc.motionEndBorderMode = MTLMotionBorderModeVanish;
|
||||
accelDesc.motionKeyframeCount = num_motion_steps;
|
||||
}
|
||||
|
||||
if (!use_fast_trace_bvh) {
|
||||
accelDesc.usage |= (MTLAccelerationStructureUsageRefit |
|
||||
MTLAccelerationStructureUsagePreferFastBuild);
|
||||
}
|
||||
|
||||
MTLAccelerationStructureSizes accelSizes = [device
|
||||
accelerationStructureSizesWithDescriptor:accelDesc];
|
||||
id<MTLAccelerationStructure> accel_uncompressed = [device
|
||||
newAccelerationStructureWithSize:accelSizes.accelerationStructureSize];
|
||||
id<MTLBuffer> scratchBuf = [device newBufferWithLength:accelSizes.buildScratchBufferSize
|
||||
options:MTLResourceStorageModePrivate];
|
||||
id<MTLBuffer> sizeBuf = [device newBufferWithLength:8 options:MTLResourceStorageModeShared];
|
||||
id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
|
||||
id<MTLAccelerationStructureCommandEncoder> accelEnc =
|
||||
[accelCommands accelerationStructureCommandEncoder];
|
||||
if (refit) {
|
||||
[accelEnc refitAccelerationStructure:accel_struct
|
||||
descriptor:accelDesc
|
||||
destination:accel_uncompressed
|
||||
scratchBuffer:scratchBuf
|
||||
scratchBufferOffset:0];
|
||||
}
|
||||
else {
|
||||
[accelEnc buildAccelerationStructure:accel_uncompressed
|
||||
descriptor:accelDesc
|
||||
scratchBuffer:scratchBuf
|
||||
scratchBufferOffset:0];
|
||||
}
|
||||
if (use_fast_trace_bvh) {
|
||||
[accelEnc writeCompactedAccelerationStructureSize:accel_uncompressed
|
||||
toBuffer:sizeBuf
|
||||
offset:0
|
||||
sizeDataType:MTLDataTypeULong];
|
||||
}
|
||||
[accelEnc endEncoding];
|
||||
[accelCommands addCompletedHandler:^(id<MTLCommandBuffer> command_buffer) {
|
||||
/* free temp resources */
|
||||
[scratchBuf release];
|
||||
[aabbBuf release];
|
||||
|
||||
if (use_fast_trace_bvh) {
|
||||
/* Compact the accel structure */
|
||||
uint64_t compressed_size = *(uint64_t *)sizeBuf.contents;
|
||||
|
||||
dispatch_async(dispatch_get_global_queue(DISPATCH_QUEUE_PRIORITY_DEFAULT, 0), ^{
|
||||
id<MTLCommandBuffer> accelCommands = [queue commandBuffer];
|
||||
id<MTLAccelerationStructureCommandEncoder> accelEnc =
|
||||
[accelCommands accelerationStructureCommandEncoder];
|
||||
id<MTLAccelerationStructure> accel = [device
|
||||
newAccelerationStructureWithSize:compressed_size];
|
||||
[accelEnc copyAndCompactAccelerationStructure:accel_uncompressed
|
||||
toAccelerationStructure:accel];
|
||||
[accelEnc endEncoding];
|
||||
[accelCommands addCompletedHandler:^(id<MTLCommandBuffer> command_buffer) {
|
||||
uint64_t allocated_size = [accel allocatedSize];
|
||||
stats.mem_alloc(allocated_size);
|
||||
accel_struct = accel;
|
||||
[accel_uncompressed release];
|
||||
accel_struct_building = false;
|
||||
}];
|
||||
[accelCommands commit];
|
||||
});
|
||||
}
|
||||
else {
|
||||
/* set our acceleration structure to the uncompressed structure */
|
||||
accel_struct = accel_uncompressed;
|
||||
|
||||
uint64_t allocated_size = [accel_struct allocatedSize];
|
||||
stats.mem_alloc(allocated_size);
|
||||
accel_struct_building = false;
|
||||
}
|
||||
[sizeBuf release];
|
||||
}];
|
||||
|
||||
accel_struct_building = true;
|
||||
[accelCommands commit];
|
||||
return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
bool BVHMetal::build_BLAS(Progress &progress,
|
||||
id<MTLDevice> device,
|
||||
id<MTLCommandQueue> queue,
|
||||
@@ -706,8 +491,6 @@ bool BVHMetal::build_BLAS(Progress &progress,
|
||||
return build_BLAS_mesh(progress, device, queue, geom, refit);
|
||||
case Geometry::HAIR:
|
||||
return build_BLAS_hair(progress, device, queue, geom, refit);
|
||||
case Geometry::POINTCLOUD:
|
||||
return build_BLAS_pointcloud(progress, device, queue, geom, refit);
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
|
@@ -115,8 +115,6 @@ class MetalDevice : public Device {
|
||||
|
||||
void load_texture_info();
|
||||
|
||||
void erase_allocation(device_memory &mem);
|
||||
|
||||
virtual bool should_use_graphics_interop() override;
|
||||
|
||||
virtual unique_ptr<DeviceQueue> gpu_queue_create() override;
|
||||
|
@@ -87,14 +87,17 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
|
||||
default:
|
||||
break;
|
||||
case METAL_GPU_INTEL: {
|
||||
use_metalrt = false;
|
||||
max_threads_per_threadgroup = 64;
|
||||
break;
|
||||
}
|
||||
case METAL_GPU_AMD: {
|
||||
use_metalrt = false;
|
||||
max_threads_per_threadgroup = 128;
|
||||
break;
|
||||
}
|
||||
case METAL_GPU_APPLE: {
|
||||
use_metalrt = true;
|
||||
max_threads_per_threadgroup = 512;
|
||||
break;
|
||||
}
|
||||
@@ -429,25 +432,6 @@ void MetalDevice::load_texture_info()
|
||||
}
|
||||
}
|
||||
|
||||
void MetalDevice::erase_allocation(device_memory &mem)
|
||||
{
|
||||
stats.mem_free(mem.device_size);
|
||||
mem.device_pointer = 0;
|
||||
mem.device_size = 0;
|
||||
|
||||
auto it = metal_mem_map.find(&mem);
|
||||
if (it != metal_mem_map.end()) {
|
||||
MetalMem *mmem = it->second.get();
|
||||
|
||||
/* blank out reference to MetalMem* in the launch params (fixes crash T94736) */
|
||||
if (mmem->pointer_index >= 0) {
|
||||
device_ptr *pointers = (device_ptr *)&launch_params;
|
||||
pointers[mmem->pointer_index] = 0;
|
||||
}
|
||||
metal_mem_map.erase(it);
|
||||
}
|
||||
}
|
||||
|
||||
MetalDevice::MetalMem *MetalDevice::generic_alloc(device_memory &mem)
|
||||
{
|
||||
size_t size = mem.memory_size();
|
||||
@@ -577,7 +561,11 @@ void MetalDevice::generic_free(device_memory &mem)
|
||||
mmem.mtlBuffer = nil;
|
||||
}
|
||||
|
||||
erase_allocation(mem);
|
||||
stats.mem_free(mem.device_size);
|
||||
mem.device_pointer = 0;
|
||||
mem.device_size = 0;
|
||||
|
||||
metal_mem_map.erase(&mem);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -966,7 +954,10 @@ void MetalDevice::tex_free(device_texture &mem)
|
||||
delayed_free_list.push_back(mmem.mtlTexture);
|
||||
mmem.mtlTexture = nil;
|
||||
}
|
||||
erase_allocation(mem);
|
||||
stats.mem_free(mem.device_size);
|
||||
mem.device_pointer = 0;
|
||||
mem.device_size = 0;
|
||||
metal_mem_map.erase(&mem);
|
||||
}
|
||||
}
|
||||
|
||||
|
@@ -36,8 +36,6 @@ enum {
|
||||
METALRT_FUNC_CURVE_RIBBON_SHADOW,
|
||||
METALRT_FUNC_CURVE_ALL,
|
||||
METALRT_FUNC_CURVE_ALL_SHADOW,
|
||||
METALRT_FUNC_POINT,
|
||||
METALRT_FUNC_POINT_SHADOW,
|
||||
METALRT_FUNC_NUM
|
||||
};
|
||||
|
||||
|
@@ -358,8 +358,6 @@ bool MetalDeviceKernels::load(MetalDevice *device, int kernel_type)
|
||||
"__intersection__curve_ribbon_shadow",
|
||||
"__intersection__curve_all",
|
||||
"__intersection__curve_all_shadow",
|
||||
"__intersection__point",
|
||||
"__intersection__point_shadow",
|
||||
};
|
||||
assert(sizeof(function_names) / sizeof(function_names[0]) == METALRT_FUNC_NUM);
|
||||
|
||||
@@ -402,50 +400,36 @@ bool MetalDeviceKernels::load(MetalDevice *device, int kernel_type)
|
||||
NSArray *function_list = nil;
|
||||
|
||||
if (device->use_metalrt) {
|
||||
id<MTLFunction> curve_intersect_default = nil;
|
||||
id<MTLFunction> curve_intersect_shadow = nil;
|
||||
id<MTLFunction> point_intersect_default = nil;
|
||||
id<MTLFunction> point_intersect_shadow = nil;
|
||||
id<MTLFunction> box_intersect_default = nil;
|
||||
id<MTLFunction> box_intersect_shadow = nil;
|
||||
if (device->kernel_features & KERNEL_FEATURE_HAIR) {
|
||||
/* Add curve intersection programs. */
|
||||
if (device->kernel_features & KERNEL_FEATURE_HAIR_THICK) {
|
||||
/* Slower programs for thick hair since that also slows down ribbons.
|
||||
* Ideally this should not be needed. */
|
||||
curve_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL];
|
||||
curve_intersect_shadow =
|
||||
rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL_SHADOW];
|
||||
box_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL];
|
||||
box_intersect_shadow = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_ALL_SHADOW];
|
||||
}
|
||||
else {
|
||||
curve_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_RIBBON];
|
||||
curve_intersect_shadow =
|
||||
box_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_RIBBON];
|
||||
box_intersect_shadow =
|
||||
rt_intersection_funcs[kernel_type][METALRT_FUNC_CURVE_RIBBON_SHADOW];
|
||||
}
|
||||
}
|
||||
if (device->kernel_features & KERNEL_FEATURE_POINTCLOUD) {
|
||||
point_intersect_default = rt_intersection_funcs[kernel_type][METALRT_FUNC_POINT];
|
||||
point_intersect_shadow = rt_intersection_funcs[kernel_type][METALRT_FUNC_POINT_SHADOW];
|
||||
}
|
||||
table_functions[METALRT_TABLE_DEFAULT] = [NSArray
|
||||
arrayWithObjects:rt_intersection_funcs[kernel_type][METALRT_FUNC_DEFAULT_TRI],
|
||||
curve_intersect_default ?
|
||||
curve_intersect_default :
|
||||
rt_intersection_funcs[kernel_type][METALRT_FUNC_DEFAULT_BOX],
|
||||
point_intersect_default ?
|
||||
point_intersect_default :
|
||||
box_intersect_default ?
|
||||
box_intersect_default :
|
||||
rt_intersection_funcs[kernel_type][METALRT_FUNC_DEFAULT_BOX],
|
||||
nil];
|
||||
table_functions[METALRT_TABLE_SHADOW] = [NSArray
|
||||
arrayWithObjects:rt_intersection_funcs[kernel_type][METALRT_FUNC_SHADOW_TRI],
|
||||
curve_intersect_shadow ?
|
||||
curve_intersect_shadow :
|
||||
rt_intersection_funcs[kernel_type][METALRT_FUNC_SHADOW_BOX],
|
||||
point_intersect_shadow ?
|
||||
point_intersect_shadow :
|
||||
box_intersect_shadow ?
|
||||
box_intersect_shadow :
|
||||
rt_intersection_funcs[kernel_type][METALRT_FUNC_SHADOW_BOX],
|
||||
nil];
|
||||
table_functions[METALRT_TABLE_LOCAL] = [NSArray
|
||||
arrayWithObjects:rt_intersection_funcs[kernel_type][METALRT_FUNC_LOCAL_TRI],
|
||||
rt_intersection_funcs[kernel_type][METALRT_FUNC_LOCAL_BOX],
|
||||
rt_intersection_funcs[kernel_type][METALRT_FUNC_LOCAL_BOX],
|
||||
nil];
|
||||
|
||||
|
@@ -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 */
|
||||
|
@@ -141,7 +141,6 @@ bool PassAccessor::get_render_tile_pixels(const RenderBuffers *render_buffers,
|
||||
const PassType type = pass_access_info_.type;
|
||||
const PassMode mode = pass_access_info_.mode;
|
||||
const PassInfo pass_info = Pass::get_info(type, pass_access_info_.include_albedo);
|
||||
int num_written_components = pass_info.num_components;
|
||||
|
||||
if (pass_info.num_components == 1) {
|
||||
/* Single channel passes. */
|
||||
@@ -189,10 +188,8 @@ bool PassAccessor::get_render_tile_pixels(const RenderBuffers *render_buffers,
|
||||
else if ((pass_info.divide_type != PASS_NONE || pass_info.direct_type != PASS_NONE ||
|
||||
pass_info.indirect_type != PASS_NONE) &&
|
||||
mode != PassMode::DENOISED) {
|
||||
/* RGB lighting passes that need to divide out color and/or sum direct and indirect.
|
||||
* These can also optionally write alpha like the combined pass. */
|
||||
/* RGB lighting passes that need to divide out color and/or sum direct and indirect. */
|
||||
get_pass_light_path(render_buffers, buffer_params, destination);
|
||||
num_written_components = 4;
|
||||
}
|
||||
else {
|
||||
/* Passes that need no special computation, or denoised passes that already
|
||||
@@ -218,7 +215,7 @@ bool PassAccessor::get_render_tile_pixels(const RenderBuffers *render_buffers,
|
||||
}
|
||||
}
|
||||
|
||||
pad_pixels(buffer_params, destination, num_written_components);
|
||||
pad_pixels(buffer_params, destination, pass_info.num_components);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
@@ -820,15 +820,8 @@ void PathTrace::tile_buffer_read()
|
||||
return;
|
||||
}
|
||||
|
||||
/* Read buffers back from device. */
|
||||
tbb::parallel_for_each(path_trace_works_, [&](unique_ptr<PathTraceWork> &path_trace_work) {
|
||||
path_trace_work->copy_render_buffers_from_device();
|
||||
});
|
||||
|
||||
/* Read (subset of) passes from output driver. */
|
||||
PathTraceTile tile(*this);
|
||||
if (output_driver_->read_render_tile(tile)) {
|
||||
/* Copy buffers to device again. */
|
||||
tbb::parallel_for_each(path_trace_works_, [](unique_ptr<PathTraceWork> &path_trace_work) {
|
||||
path_trace_work->copy_render_buffers_to_device();
|
||||
});
|
||||
|
@@ -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);
|
||||
|
@@ -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,
|
||||
|
@@ -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,
|
||||
|
@@ -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;
|
||||
|
@@ -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: {
|
||||
|
@@ -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 {
|
||||
|
@@ -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
|
||||
|
@@ -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;
|
||||
|
@@ -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;
|
||||
|
@@ -199,18 +199,22 @@ ccl_device int volume_sample_channel(float3 albedo,
|
||||
* Tracing". Matt Jen-Yuan Chiang, Peter Kutz, Brent Burley. SIGGRAPH 2016. */
|
||||
float3 weights = fabs(throughput * albedo);
|
||||
float sum_weights = weights.x + weights.y + weights.z;
|
||||
float3 weights_pdf;
|
||||
|
||||
if (sum_weights > 0.0f) {
|
||||
*pdf = weights / sum_weights;
|
||||
weights_pdf = weights / sum_weights;
|
||||
}
|
||||
else {
|
||||
*pdf = make_float3(1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f);
|
||||
weights_pdf = make_float3(1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f);
|
||||
}
|
||||
|
||||
if (rand < pdf->x) {
|
||||
*pdf = weights_pdf;
|
||||
|
||||
/* OpenCL does not support -> on float3, so don't use pdf->x. */
|
||||
if (rand < weights_pdf.x) {
|
||||
return 0;
|
||||
}
|
||||
else if (rand < pdf->x + pdf->y) {
|
||||
else if (rand < weights_pdf.x + weights_pdf.y) {
|
||||
return 1;
|
||||
}
|
||||
else {
|
||||
|
@@ -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;
|
||||
@@ -616,150 +576,6 @@ __intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
#endif /* __HAIR__ */
|
||||
|
||||
#ifdef __POINTCLOUD__
|
||||
ccl_device_inline
|
||||
void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal,
|
||||
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload,
|
||||
const uint object,
|
||||
const uint prim,
|
||||
const uint type,
|
||||
const float3 ray_origin,
|
||||
const float3 ray_direction,
|
||||
float time,
|
||||
const float ray_tmax,
|
||||
thread BoundingBoxIntersectionResult &result)
|
||||
{
|
||||
# ifdef __VISIBILITY_FLAG__
|
||||
const uint visibility = payload.visibility;
|
||||
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
|
||||
return;
|
||||
}
|
||||
# endif
|
||||
|
||||
float3 P = ray_origin;
|
||||
float3 dir = ray_direction;
|
||||
|
||||
/* The direction is not normalized by default, but the point intersection routine expects that */
|
||||
float len;
|
||||
dir = normalize_len(dir, &len);
|
||||
|
||||
Intersection isect;
|
||||
isect.t = ray_tmax;
|
||||
/* Transform maximum distance into object space. */
|
||||
if (isect.t != FLT_MAX)
|
||||
isect.t *= len;
|
||||
|
||||
MetalKernelContext context(launch_params_metal);
|
||||
if (context.point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
|
||||
result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
|
||||
launch_params_metal, payload, object, prim, isect.u);
|
||||
if (result.accept) {
|
||||
result.distance = isect.t / len;
|
||||
payload.u = isect.u;
|
||||
payload.v = isect.v;
|
||||
payload.prim = prim;
|
||||
payload.type = type;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
ccl_device_inline
|
||||
void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params_metal,
|
||||
ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload,
|
||||
const uint object,
|
||||
const uint prim,
|
||||
const uint type,
|
||||
const float3 ray_origin,
|
||||
const float3 ray_direction,
|
||||
float time,
|
||||
const float ray_tmax,
|
||||
thread BoundingBoxIntersectionResult &result)
|
||||
{
|
||||
const uint visibility = payload.visibility;
|
||||
|
||||
float3 P = ray_origin;
|
||||
float3 dir = ray_direction;
|
||||
|
||||
/* The direction is not normalized by default, but the point intersection routine expects that */
|
||||
float len;
|
||||
dir = normalize_len(dir, &len);
|
||||
|
||||
Intersection isect;
|
||||
isect.t = ray_tmax;
|
||||
/* Transform maximum distance into object space */
|
||||
if (isect.t != FLT_MAX)
|
||||
isect.t *= len;
|
||||
|
||||
MetalKernelContext context(launch_params_metal);
|
||||
if (context.point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
|
||||
result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>(
|
||||
launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax);
|
||||
result.accept = !result.continue_search;
|
||||
|
||||
if (result.accept) {
|
||||
result.distance = isect.t / len;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]]
|
||||
BoundingBoxIntersectionResult
|
||||
__intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
|
||||
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]],
|
||||
const uint object [[user_instance_id]],
|
||||
const uint primitive_id [[primitive_id]],
|
||||
const float3 ray_origin [[origin]],
|
||||
const float3 ray_direction [[direction]],
|
||||
const float ray_tmax [[max_distance]])
|
||||
{
|
||||
const uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
const int type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
|
||||
BoundingBoxIntersectionResult result;
|
||||
result.accept = false;
|
||||
result.continue_search = true;
|
||||
result.distance = ray_tmax;
|
||||
|
||||
metalrt_intersection_point(launch_params_metal, payload, object, prim, type, ray_origin, ray_direction,
|
||||
# if defined(__METALRT_MOTION__)
|
||||
payload.time,
|
||||
# else
|
||||
0.0f,
|
||||
# endif
|
||||
ray_tmax, result);
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]]
|
||||
BoundingBoxIntersectionResult
|
||||
__intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
|
||||
ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]],
|
||||
const uint object [[user_instance_id]],
|
||||
const uint primitive_id [[primitive_id]],
|
||||
const float3 ray_origin [[origin]],
|
||||
const float3 ray_direction [[direction]],
|
||||
const float ray_tmax [[max_distance]])
|
||||
{
|
||||
const uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
const int type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
|
||||
BoundingBoxIntersectionResult result;
|
||||
result.accept = false;
|
||||
result.continue_search = true;
|
||||
result.distance = ray_tmax;
|
||||
|
||||
metalrt_intersection_point_shadow(launch_params_metal, payload, object, prim, type, ray_origin, ray_direction,
|
||||
# if defined(__METALRT_MOTION__)
|
||||
payload.time,
|
||||
# else
|
||||
0.0f,
|
||||
# endif
|
||||
ray_tmax, result);
|
||||
|
||||
return result;
|
||||
}
|
||||
#endif /* __POINTCLOUD__ */
|
||||
#endif /* __METALRT__ */
|
||||
|
@@ -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()
|
||||
|
@@ -214,21 +214,6 @@ ccl_device_inline void film_get_pass_pixel_light_path(
|
||||
pixel[0] = f.x;
|
||||
pixel[1] = f.y;
|
||||
pixel[2] = f.z;
|
||||
|
||||
/* Optional alpha channel. */
|
||||
if (kfilm_convert->num_components >= 4) {
|
||||
if (kfilm_convert->pass_combined != PASS_UNUSED) {
|
||||
float scale, scale_exposure;
|
||||
film_get_scale_and_scale_exposure(kfilm_convert, buffer, &scale, &scale_exposure);
|
||||
|
||||
ccl_global const float *in_combined = buffer + kfilm_convert->pass_combined;
|
||||
const float alpha = in_combined[3] * scale;
|
||||
pixel[3] = film_transparency_to_alpha(alpha);
|
||||
}
|
||||
else {
|
||||
pixel[3] = 1.0f;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
ccl_device_inline void film_get_pass_pixel_float3(ccl_global const KernelFilmConvert *ccl_restrict
|
||||
|
@@ -226,18 +226,6 @@ ccl_device float curve_thickness(KernelGlobals kg, ccl_private const ShaderData
|
||||
return r * 2.0f;
|
||||
}
|
||||
|
||||
/* Curve random */
|
||||
|
||||
ccl_device float curve_random(KernelGlobals kg, ccl_private const ShaderData *sd)
|
||||
{
|
||||
if (sd->type & PRIMITIVE_CURVE) {
|
||||
const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_CURVE_RANDOM);
|
||||
return (desc.offset != ATTR_STD_NOT_FOUND) ? curve_attribute_float(kg, sd, desc, NULL, NULL) :
|
||||
0.0f;
|
||||
}
|
||||
return 0.0f;
|
||||
}
|
||||
|
||||
/* Curve location for motion pass, linear interpolation between keys and
|
||||
* ignoring radius because we do the same for the motion keys */
|
||||
|
||||
|
@@ -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.
|
||||
*/
|
||||
|
@@ -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) {
|
||||
|
@@ -81,7 +81,7 @@ ccl_device float3 point_attribute_float3(KernelGlobals kg,
|
||||
# endif
|
||||
|
||||
if (desc.element == ATTR_ELEMENT_VERTEX) {
|
||||
return kernel_tex_fetch(__attributes_float3, desc.offset + sd->prim);
|
||||
return float4_to_float3(kernel_tex_fetch(__attributes_float4, desc.offset + sd->prim));
|
||||
}
|
||||
else {
|
||||
return make_float3(0.0f, 0.0f, 0.0f);
|
||||
@@ -109,59 +109,17 @@ ccl_device float4 point_attribute_float4(KernelGlobals kg,
|
||||
}
|
||||
}
|
||||
|
||||
/* Point position */
|
||||
|
||||
ccl_device float3 point_position(KernelGlobals kg, ccl_private const ShaderData *sd)
|
||||
{
|
||||
if (sd->type & PRIMITIVE_POINT) {
|
||||
/* World space center. */
|
||||
float3 P = (sd->type & PRIMITIVE_MOTION) ?
|
||||
float4_to_float3(motion_point(kg, sd->object, sd->prim, sd->time)) :
|
||||
float4_to_float3(kernel_tex_fetch(__points, sd->prim));
|
||||
|
||||
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
|
||||
object_position_transform(kg, sd, &P);
|
||||
}
|
||||
|
||||
return P;
|
||||
}
|
||||
|
||||
return zero_float3();
|
||||
}
|
||||
|
||||
/* Point radius */
|
||||
|
||||
ccl_device float point_radius(KernelGlobals kg, ccl_private const ShaderData *sd)
|
||||
{
|
||||
if (sd->type & PRIMITIVE_POINT) {
|
||||
/* World space radius. */
|
||||
const float r = kernel_tex_fetch(__points, sd->prim).w;
|
||||
|
||||
if (sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED) {
|
||||
return r;
|
||||
}
|
||||
else {
|
||||
float3 dir = make_float3(r, r, r);
|
||||
object_dir_transform(kg, sd, &dir);
|
||||
return average(dir);
|
||||
}
|
||||
return kernel_tex_fetch(__points, sd->prim).w;
|
||||
}
|
||||
|
||||
return 0.0f;
|
||||
}
|
||||
|
||||
/* Point random */
|
||||
|
||||
ccl_device float point_random(KernelGlobals kg, ccl_private const ShaderData *sd)
|
||||
{
|
||||
if (sd->type & PRIMITIVE_POINT) {
|
||||
const AttributeDescriptor desc = find_attribute(kg, sd, ATTR_STD_POINT_RANDOM);
|
||||
return (desc.offset != ATTR_STD_NOT_FOUND) ? point_attribute_float(kg, sd, desc, NULL, NULL) :
|
||||
0.0f;
|
||||
}
|
||||
return 0.0f;
|
||||
}
|
||||
|
||||
/* Point location for motion pass, linear interpolation between keys and
|
||||
* ignoring radius because we do the same for the motion keys */
|
||||
|
||||
|
@@ -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;
|
||||
|
||||
|
@@ -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
|
||||
|
@@ -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. */
|
||||
|
@@ -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);
|
||||
|
||||
|
@@ -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
|
||||
|
@@ -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);
|
||||
|
@@ -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;
|
||||
}
|
||||
|
||||
|
@@ -182,18 +182,13 @@ 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;
|
||||
shadow_flag |= 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) {
|
||||
@@ -271,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);
|
||||
@@ -319,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;
|
||||
@@ -363,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();
|
||||
|
||||
@@ -382,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);
|
||||
|
@@ -791,17 +791,13 @@ 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;
|
||||
shadow_flag |= PATH_RAY_VOLUME_PASS;
|
||||
const float3 throughput_phase = throughput * bsdf_eval_sum(&phase_eval);
|
||||
|
||||
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
|
||||
@@ -877,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);
|
||||
|
@@ -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 **************************/
|
||||
|
@@ -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__ */
|
||||
|
@@ -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. */
|
||||
|
@@ -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;
|
||||
|
@@ -418,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;
|
||||
|
@@ -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. */
|
||||
|
@@ -116,8 +116,6 @@ ustring OSLRenderServices::u_curve_tangent_normal("geom:curve_tangent_normal");
|
||||
ustring OSLRenderServices::u_curve_random("geom:curve_random");
|
||||
ustring OSLRenderServices::u_is_point("geom:is_point");
|
||||
ustring OSLRenderServices::u_point_radius("geom:point_radius");
|
||||
ustring OSLRenderServices::u_point_position("geom:point_position");
|
||||
ustring OSLRenderServices::u_point_random("geom:point_random");
|
||||
ustring OSLRenderServices::u_normal_map_normal("geom:normal_map_normal");
|
||||
ustring OSLRenderServices::u_path_ray_length("path:ray_length");
|
||||
ustring OSLRenderServices::u_path_ray_depth("path:ray_depth");
|
||||
@@ -1001,10 +999,6 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg
|
||||
float3 f = curve_tangent_normal(kg, sd);
|
||||
return set_attribute_float3(f, type, derivatives, val);
|
||||
}
|
||||
else if (name == u_curve_random) {
|
||||
float f = curve_random(kg, sd);
|
||||
return set_attribute_float(f, type, derivatives, val);
|
||||
}
|
||||
/* point attributes */
|
||||
else if (name == u_is_point) {
|
||||
float f = (sd->type & PRIMITIVE_POINT) != 0;
|
||||
@@ -1014,14 +1008,6 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg
|
||||
float f = point_radius(kg, sd);
|
||||
return set_attribute_float(f, type, derivatives, val);
|
||||
}
|
||||
else if (name == u_point_position) {
|
||||
float3 f = point_position(kg, sd);
|
||||
return set_attribute_float3(f, type, derivatives, val);
|
||||
}
|
||||
else if (name == u_point_random) {
|
||||
float f = point_random(kg, sd);
|
||||
return set_attribute_float(f, type, derivatives, val);
|
||||
}
|
||||
else if (name == u_normal_map_normal) {
|
||||
if (sd->type & PRIMITIVE_TRIANGLE) {
|
||||
float3 f = triangle_smooth_normal_unnormalized(kg, sd, sd->Ng, sd->prim, sd->u, sd->v);
|
||||
|
@@ -298,9 +298,7 @@ class OSLRenderServices : public OSL::RendererServices {
|
||||
static ustring u_curve_tangent_normal;
|
||||
static ustring u_curve_random;
|
||||
static ustring u_is_point;
|
||||
static ustring u_point_position;
|
||||
static ustring u_point_radius;
|
||||
static ustring u_point_random;
|
||||
static ustring u_normal_map_normal;
|
||||
static ustring u_path_ray_length;
|
||||
static ustring u_path_ray_depth;
|
||||
|
@@ -49,7 +49,6 @@ set(SRC_OSL
|
||||
node_glossy_bsdf.osl
|
||||
node_gradient_texture.osl
|
||||
node_hair_info.osl
|
||||
node_point_info.osl
|
||||
node_scatter_volume.osl
|
||||
node_absorption_volume.osl
|
||||
node_principled_volume.osl
|
||||
|
@@ -1,26 +0,0 @@
|
||||
/*
|
||||
* Copyright 2011-2022 Blender Foundation
|
||||
*
|
||||
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||
* you may not use this file except in compliance with the License.
|
||||
* You may obtain a copy of the License at
|
||||
*
|
||||
* http://www.apache.org/licenses/LICENSE-2.0
|
||||
*
|
||||
* Unless required by applicable law or agreed to in writing, software
|
||||
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
* See the License for the specific language governing permissions and
|
||||
* limitations under the License.
|
||||
*/
|
||||
|
||||
#include "stdcycles.h"
|
||||
|
||||
shader node_point_info(output point Position = point(0.0, 0.0, 0.0),
|
||||
output float Radius = 0.0,
|
||||
output float Random = 0.0)
|
||||
{
|
||||
getattribute("geom:point_position", Position);
|
||||
getattribute("geom:point_radius", Radius);
|
||||
getattribute("geom:point_random", Random);
|
||||
}
|
@@ -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();
|
||||
|
||||
|
@@ -87,9 +87,7 @@ ccl_device_noinline void svm_node_attr(KernelGlobals kg,
|
||||
if (node.y == ATTR_STD_GENERATED && desc.element == ATTR_ELEMENT_NONE) {
|
||||
/* No generated attribute, fall back to object coordinates. */
|
||||
float3 f = sd->P;
|
||||
if (sd->object != OBJECT_NONE) {
|
||||
object_inverse_position_transform(kg, sd, &f);
|
||||
}
|
||||
object_inverse_position_transform(kg, sd, &f);
|
||||
if (type == NODE_ATTR_OUTPUT_FLOAT) {
|
||||
stack_store_float(stack, out_offset, average(f));
|
||||
}
|
||||
@@ -181,9 +179,7 @@ ccl_device_noinline void svm_node_attr_bump_dx(KernelGlobals kg,
|
||||
if (node.y == ATTR_STD_GENERATED && desc.element == ATTR_ELEMENT_NONE) {
|
||||
/* No generated attribute, fall back to object coordinates. */
|
||||
float3 f = sd->P + sd->dP.dx;
|
||||
if (sd->object != OBJECT_NONE) {
|
||||
object_inverse_position_transform(kg, sd, &f);
|
||||
}
|
||||
object_inverse_position_transform(kg, sd, &f);
|
||||
if (type == NODE_ATTR_OUTPUT_FLOAT) {
|
||||
stack_store_float(stack, out_offset, average(f));
|
||||
}
|
||||
@@ -279,9 +275,7 @@ ccl_device_noinline void svm_node_attr_bump_dy(KernelGlobals kg,
|
||||
if (node.y == ATTR_STD_GENERATED && desc.element == ATTR_ELEMENT_NONE) {
|
||||
/* No generated attribute, fall back to object coordinates. */
|
||||
float3 f = sd->P + sd->dP.dy;
|
||||
if (sd->object != OBJECT_NONE) {
|
||||
object_inverse_position_transform(kg, sd, &f);
|
||||
}
|
||||
object_inverse_position_transform(kg, sd, &f);
|
||||
if (type == NODE_ATTR_OUTPUT_FLOAT) {
|
||||
stack_store_float(stack, out_offset, average(f));
|
||||
}
|
||||
|
@@ -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__ */
|
||||
|
||||
|
@@ -242,6 +242,13 @@ ccl_device_noinline void svm_node_hair_info(KernelGlobals kg,
|
||||
stack_store_float(stack, out_offset, data);
|
||||
break;
|
||||
}
|
||||
# if 0
|
||||
case NODE_INFO_CURVE_FADE: {
|
||||
data = sd->curve_transparency;
|
||||
stack_store_float(stack, out_offset, data);
|
||||
break;
|
||||
}
|
||||
# endif
|
||||
case NODE_INFO_CURVE_TANGENT_NORMAL: {
|
||||
data3 = curve_tangent_normal(kg, sd);
|
||||
stack_store_float3(stack, out_offset, data3);
|
||||
@@ -251,28 +258,4 @@ ccl_device_noinline void svm_node_hair_info(KernelGlobals kg,
|
||||
}
|
||||
#endif
|
||||
|
||||
#ifdef __POINTCLOUD__
|
||||
|
||||
/* Point Info */
|
||||
|
||||
ccl_device_noinline void svm_node_point_info(KernelGlobals kg,
|
||||
ccl_private ShaderData *sd,
|
||||
ccl_private float *stack,
|
||||
uint type,
|
||||
uint out_offset)
|
||||
{
|
||||
switch (type) {
|
||||
case NODE_INFO_POINT_POSITION:
|
||||
stack_store_float3(stack, out_offset, point_position(kg, sd));
|
||||
break;
|
||||
case NODE_INFO_POINT_RADIUS:
|
||||
stack_store_float(stack, out_offset, point_radius(kg, sd));
|
||||
break;
|
||||
case NODE_INFO_POINT_RANDOM:
|
||||
break; /* handled as attribute */
|
||||
}
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
@@ -454,14 +454,13 @@ ccl_device void svm_eval_nodes(KernelGlobals kg,
|
||||
break;
|
||||
#if defined(__HAIR__)
|
||||
case NODE_HAIR_INFO:
|
||||
svm_node_hair_info(kg, sd, stack, node.y, node.z);
|
||||
break;
|
||||
#endif
|
||||
#if defined(__POINTCLOUD__)
|
||||
case NODE_POINT_INFO:
|
||||
svm_node_point_info(kg, sd, stack, node.y, node.z);
|
||||
IF_KERNEL_NODES_FEATURE(HAIR)
|
||||
{
|
||||
svm_node_hair_info(kg, sd, stack, node.y, node.z);
|
||||
}
|
||||
break;
|
||||
#endif
|
||||
|
||||
case NODE_TEXTURE_MAPPING:
|
||||
offset = svm_node_texture_mapping(kg, sd, stack, node.y, node.z, offset);
|
||||
break;
|
||||
|
@@ -81,7 +81,6 @@ typedef enum ShaderNodeType {
|
||||
NODE_OBJECT_INFO,
|
||||
NODE_PARTICLE_INFO,
|
||||
NODE_HAIR_INFO,
|
||||
NODE_POINT_INFO,
|
||||
NODE_TEXTURE_MAPPING,
|
||||
NODE_MAPPING,
|
||||
NODE_MIN_MAX,
|
||||
@@ -177,16 +176,12 @@ typedef enum NodeHairInfo {
|
||||
NODE_INFO_CURVE_INTERCEPT,
|
||||
NODE_INFO_CURVE_LENGTH,
|
||||
NODE_INFO_CURVE_THICKNESS,
|
||||
/* Fade for minimum hair width transiency. */
|
||||
// NODE_INFO_CURVE_FADE,
|
||||
NODE_INFO_CURVE_TANGENT_NORMAL,
|
||||
NODE_INFO_CURVE_RANDOM,
|
||||
} NodeHairInfo;
|
||||
|
||||
typedef enum NodePointInfo {
|
||||
NODE_INFO_POINT_POSITION,
|
||||
NODE_INFO_POINT_RADIUS,
|
||||
NODE_INFO_POINT_RANDOM,
|
||||
} NodePointInfo;
|
||||
|
||||
typedef enum NodeLightPath {
|
||||
NODE_LP_camera = 0,
|
||||
NODE_LP_shadow,
|
||||
|
@@ -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;
|
||||
@@ -1574,21 +1565,21 @@ enum KernelFeatureFlag : uint32_t {
|
||||
KERNEL_FEATURE_NODE_BSDF = (1U << 0U),
|
||||
KERNEL_FEATURE_NODE_EMISSION = (1U << 1U),
|
||||
KERNEL_FEATURE_NODE_VOLUME = (1U << 2U),
|
||||
KERNEL_FEATURE_NODE_BUMP = (1U << 3U),
|
||||
KERNEL_FEATURE_NODE_BUMP_STATE = (1U << 4U),
|
||||
KERNEL_FEATURE_NODE_VORONOI_EXTRA = (1U << 5U),
|
||||
KERNEL_FEATURE_NODE_RAYTRACE = (1U << 6U),
|
||||
KERNEL_FEATURE_NODE_AOV = (1U << 7U),
|
||||
KERNEL_FEATURE_NODE_LIGHT_PATH = (1U << 8U),
|
||||
KERNEL_FEATURE_NODE_HAIR = (1U << 3U),
|
||||
KERNEL_FEATURE_NODE_BUMP = (1U << 4U),
|
||||
KERNEL_FEATURE_NODE_BUMP_STATE = (1U << 5U),
|
||||
KERNEL_FEATURE_NODE_VORONOI_EXTRA = (1U << 6U),
|
||||
KERNEL_FEATURE_NODE_RAYTRACE = (1U << 7U),
|
||||
KERNEL_FEATURE_NODE_AOV = (1U << 8U),
|
||||
KERNEL_FEATURE_NODE_LIGHT_PATH = (1U << 9U),
|
||||
|
||||
/* Use denoising kernels and output denoising passes. */
|
||||
KERNEL_FEATURE_DENOISING = (1U << 9U),
|
||||
KERNEL_FEATURE_DENOISING = (1U << 10U),
|
||||
|
||||
/* Use path tracing kernels. */
|
||||
KERNEL_FEATURE_PATH_TRACING = (1U << 10U),
|
||||
KERNEL_FEATURE_PATH_TRACING = (1U << 11U),
|
||||
|
||||
/* BVH/sampling kernel features. */
|
||||
KERNEL_FEATURE_POINTCLOUD = (1U << 11U),
|
||||
KERNEL_FEATURE_HAIR = (1U << 12U),
|
||||
KERNEL_FEATURE_HAIR_THICK = (1U << 13U),
|
||||
KERNEL_FEATURE_OBJECT_MOTION = (1U << 14U),
|
||||
@@ -1625,6 +1616,9 @@ enum KernelFeatureFlag : uint32_t {
|
||||
KERNEL_FEATURE_AO_PASS = (1U << 25U),
|
||||
KERNEL_FEATURE_AO_ADDITIVE = (1U << 26U),
|
||||
KERNEL_FEATURE_AO = (KERNEL_FEATURE_AO_PASS | KERNEL_FEATURE_AO_ADDITIVE),
|
||||
|
||||
/* Point clouds. */
|
||||
KERNEL_FEATURE_POINTCLOUD = (1U << 27U),
|
||||
};
|
||||
|
||||
/* Shader node feature mask, to specialize shader evaluation for kernels. */
|
||||
@@ -1634,7 +1628,7 @@ enum KernelFeatureFlag : uint32_t {
|
||||
KERNEL_FEATURE_NODE_LIGHT_PATH)
|
||||
#define KERNEL_FEATURE_NODE_MASK_SURFACE_SHADOW \
|
||||
(KERNEL_FEATURE_NODE_BSDF | KERNEL_FEATURE_NODE_EMISSION | KERNEL_FEATURE_NODE_VOLUME | \
|
||||
KERNEL_FEATURE_NODE_BUMP | KERNEL_FEATURE_NODE_BUMP_STATE | \
|
||||
KERNEL_FEATURE_NODE_HAIR | KERNEL_FEATURE_NODE_BUMP | KERNEL_FEATURE_NODE_BUMP_STATE | \
|
||||
KERNEL_FEATURE_NODE_VORONOI_EXTRA | KERNEL_FEATURE_NODE_LIGHT_PATH)
|
||||
#define KERNEL_FEATURE_NODE_MASK_SURFACE \
|
||||
(KERNEL_FEATURE_NODE_MASK_SURFACE_SHADOW | KERNEL_FEATURE_NODE_RAYTRACE | \
|
||||
|
@@ -1178,12 +1178,6 @@ void AlembicProcedural::read_subd(AlembicObject *abc_object, Abc::chrono_t frame
|
||||
cached_data.subd_creases_weight.copy_to_socket(
|
||||
frame_time, mesh, mesh->get_subd_creases_weight_socket());
|
||||
|
||||
cached_data.subd_vertex_crease_indices.copy_to_socket(
|
||||
frame_time, mesh, mesh->get_subd_vert_creases_socket());
|
||||
|
||||
cached_data.subd_vertex_crease_weights.copy_to_socket(
|
||||
frame_time, mesh, mesh->get_subd_vert_creases_weight_socket());
|
||||
|
||||
mesh->set_num_subd_faces(mesh->get_subd_shader().size());
|
||||
|
||||
/* Update attributes. */
|
||||
|
@@ -320,8 +320,6 @@ struct CachedData {
|
||||
DataStore<int> num_ngons;
|
||||
DataStore<array<int>> subd_creases_edge;
|
||||
DataStore<array<float>> subd_creases_weight;
|
||||
DataStore<array<int>> subd_vertex_crease_indices;
|
||||
DataStore<array<float>> subd_vertex_crease_weights;
|
||||
|
||||
/* hair data */
|
||||
DataStore<array<float3>> curve_keys;
|
||||
|
@@ -478,9 +478,7 @@ static void add_subd_polygons(CachedData &cached_data, const SubDSchemaData &dat
|
||||
cached_data.uv_loops.add_data(uv_loops, time);
|
||||
}
|
||||
|
||||
static void add_subd_edge_creases(CachedData &cached_data,
|
||||
const SubDSchemaData &data,
|
||||
chrono_t time)
|
||||
static void add_subd_creases(CachedData &cached_data, const SubDSchemaData &data, chrono_t time)
|
||||
{
|
||||
if (!(data.crease_indices.valid() && data.crease_indices.valid() &&
|
||||
data.crease_sharpnesses.valid())) {
|
||||
@@ -519,37 +517,6 @@ static void add_subd_edge_creases(CachedData &cached_data,
|
||||
}
|
||||
}
|
||||
|
||||
static void add_subd_vertex_creases(CachedData &cached_data,
|
||||
const SubDSchemaData &data,
|
||||
chrono_t time)
|
||||
{
|
||||
if (!(data.corner_indices.valid() && data.crease_sharpnesses.valid())) {
|
||||
return;
|
||||
}
|
||||
|
||||
const ISampleSelector iss = ISampleSelector(time);
|
||||
const Int32ArraySamplePtr creases_indices = data.crease_indices.getValue(iss);
|
||||
const FloatArraySamplePtr creases_sharpnesses = data.crease_sharpnesses.getValue(iss);
|
||||
|
||||
if (!(creases_indices && creases_sharpnesses) ||
|
||||
creases_indices->size() != creases_sharpnesses->size()) {
|
||||
return;
|
||||
}
|
||||
|
||||
array<float> sharpnesses;
|
||||
sharpnesses.reserve(creases_indices->size());
|
||||
array<int> indices;
|
||||
indices.reserve(creases_indices->size());
|
||||
|
||||
for (size_t i = 0; i < creases_indices->size(); i++) {
|
||||
indices.push_back_reserved((*creases_indices)[i]);
|
||||
sharpnesses.push_back_reserved((*creases_sharpnesses)[i]);
|
||||
}
|
||||
|
||||
cached_data.subd_vertex_crease_indices.add_data(indices, time);
|
||||
cached_data.subd_vertex_crease_weights.add_data(sharpnesses, time);
|
||||
}
|
||||
|
||||
static void read_subd_geometry(CachedData &cached_data, const SubDSchemaData &data, chrono_t time)
|
||||
{
|
||||
const ISampleSelector iss = ISampleSelector(time);
|
||||
@@ -558,8 +525,7 @@ static void read_subd_geometry(CachedData &cached_data, const SubDSchemaData &da
|
||||
|
||||
if (data.topology_variance != kHomogenousTopology || cached_data.shader.size() == 0) {
|
||||
add_subd_polygons(cached_data, data, time);
|
||||
add_subd_edge_creases(cached_data, data, time);
|
||||
add_subd_vertex_creases(cached_data, data, time);
|
||||
add_subd_creases(cached_data, data, time);
|
||||
}
|
||||
}
|
||||
|
||||
|
@@ -76,10 +76,9 @@ struct SubDSchemaData {
|
||||
|
||||
vector<FaceSetShaderIndexPair> shader_face_sets;
|
||||
|
||||
// Those are unsupported for now.
|
||||
Alembic::AbcGeom::IInt32ArrayProperty corner_indices;
|
||||
Alembic::AbcGeom::IFloatArrayProperty corner_sharpnesses;
|
||||
|
||||
// Those are unsupported for now.
|
||||
Alembic::AbcGeom::IInt32Property face_varying_interpolate_boundary;
|
||||
Alembic::AbcGeom::IInt32Property face_varying_propagate_corners;
|
||||
Alembic::AbcGeom::IInt32Property interpolate_boundary;
|
||||
|
@@ -263,9 +263,7 @@ template<typename T> inline void cast_from_float4(T *data, float4 value)
|
||||
|
||||
/* Slower versions for other all data types, which needs to convert to float and back. */
|
||||
template<typename T, bool compress_as_srgb = false>
|
||||
inline void processor_apply_pixels_rgba(const OCIO::Processor *processor,
|
||||
T *pixels,
|
||||
size_t num_pixels)
|
||||
inline void processor_apply_pixels(const OCIO::Processor *processor, T *pixels, size_t num_pixels)
|
||||
{
|
||||
/* TODO: implement faster version for when we know the conversion
|
||||
* is a simple matrix transform between linear spaces. In that case
|
||||
@@ -312,79 +310,25 @@ inline void processor_apply_pixels_rgba(const OCIO::Processor *processor,
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, bool compress_as_srgb = false>
|
||||
inline void processor_apply_pixels_grayscale(const OCIO::Processor *processor,
|
||||
T *pixels,
|
||||
size_t num_pixels)
|
||||
{
|
||||
OCIO::ConstCPUProcessorRcPtr device_processor = processor->getDefaultCPUProcessor();
|
||||
|
||||
/* Process large images in chunks to keep temporary memory requirement down. */
|
||||
const size_t chunk_size = std::min((size_t)(16 * 1024 * 1024), num_pixels);
|
||||
vector<float> float_pixels(chunk_size * 3);
|
||||
|
||||
for (size_t j = 0; j < num_pixels; j += chunk_size) {
|
||||
size_t width = std::min(chunk_size, num_pixels - j);
|
||||
|
||||
/* Convert to 3 channels, since that's the minimum required by OpenColorIO. */
|
||||
{
|
||||
const T *pixel = pixels + j;
|
||||
float *fpixel = float_pixels.data();
|
||||
for (size_t i = 0; i < width; i++, pixel++, fpixel += 3) {
|
||||
const float f = util_image_cast_to_float<T>(*pixel);
|
||||
fpixel[0] = f;
|
||||
fpixel[1] = f;
|
||||
fpixel[2] = f;
|
||||
}
|
||||
}
|
||||
|
||||
OCIO::PackedImageDesc desc((float *)float_pixels.data(), width, 1, 3);
|
||||
device_processor->apply(desc);
|
||||
|
||||
{
|
||||
T *pixel = pixels + j;
|
||||
const float *fpixel = float_pixels.data();
|
||||
for (size_t i = 0; i < width; i++, pixel++, fpixel += 3) {
|
||||
float f = average(make_float3(fpixel[0], fpixel[1], fpixel[2]));
|
||||
if (compress_as_srgb) {
|
||||
f = color_linear_to_srgb(f);
|
||||
}
|
||||
*pixel = util_image_cast_from_float<T>(f);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
template<typename T>
|
||||
void ColorSpaceManager::to_scene_linear(
|
||||
ustring colorspace, T *pixels, size_t num_pixels, bool is_rgba, bool compress_as_srgb)
|
||||
void ColorSpaceManager::to_scene_linear(ustring colorspace,
|
||||
T *pixels,
|
||||
size_t num_pixels,
|
||||
bool compress_as_srgb)
|
||||
{
|
||||
#ifdef WITH_OCIO
|
||||
const OCIO::Processor *processor = (const OCIO::Processor *)get_processor(colorspace);
|
||||
|
||||
if (processor) {
|
||||
if (is_rgba) {
|
||||
if (compress_as_srgb) {
|
||||
/* Compress output as sRGB. */
|
||||
processor_apply_pixels_rgba<T, true>(processor, pixels, num_pixels);
|
||||
}
|
||||
else {
|
||||
/* Write output as scene linear directly. */
|
||||
processor_apply_pixels_rgba<T>(processor, pixels, num_pixels);
|
||||
}
|
||||
if (compress_as_srgb) {
|
||||
/* Compress output as sRGB. */
|
||||
processor_apply_pixels<T, true>(processor, pixels, num_pixels);
|
||||
}
|
||||
else {
|
||||
if (compress_as_srgb) {
|
||||
/* Compress output as sRGB. */
|
||||
processor_apply_pixels_grayscale<T, true>(processor, pixels, num_pixels);
|
||||
}
|
||||
else {
|
||||
/* Write output as scene linear directly. */
|
||||
processor_apply_pixels_grayscale<T>(processor, pixels, num_pixels);
|
||||
}
|
||||
/* Write output as scene linear directly. */
|
||||
processor_apply_pixels<T>(processor, pixels, num_pixels);
|
||||
}
|
||||
}
|
||||
#else
|
||||
@@ -404,11 +348,6 @@ void ColorSpaceManager::to_scene_linear(ColorSpaceProcessor *processor_,
|
||||
|
||||
if (processor) {
|
||||
OCIO::ConstCPUProcessorRcPtr device_processor = processor->getDefaultCPUProcessor();
|
||||
if (channels == 1) {
|
||||
float3 rgb = make_float3(pixel[0], pixel[0], pixel[0]);
|
||||
device_processor->applyRGB(&rgb.x);
|
||||
pixel[0] = average(rgb);
|
||||
}
|
||||
if (channels == 3) {
|
||||
device_processor->applyRGB(pixel);
|
||||
}
|
||||
@@ -451,9 +390,9 @@ void ColorSpaceManager::free_memory()
|
||||
}
|
||||
|
||||
/* Template instantiations so we don't have to inline functions. */
|
||||
template void ColorSpaceManager::to_scene_linear(ustring, uchar *, size_t, bool, bool);
|
||||
template void ColorSpaceManager::to_scene_linear(ustring, ushort *, size_t, bool, bool);
|
||||
template void ColorSpaceManager::to_scene_linear(ustring, half *, size_t, bool, bool);
|
||||
template void ColorSpaceManager::to_scene_linear(ustring, float *, size_t, bool, bool);
|
||||
template void ColorSpaceManager::to_scene_linear(ustring, uchar *, size_t, bool);
|
||||
template void ColorSpaceManager::to_scene_linear(ustring, ushort *, size_t, bool);
|
||||
template void ColorSpaceManager::to_scene_linear(ustring, half *, size_t, bool);
|
||||
template void ColorSpaceManager::to_scene_linear(ustring, float *, size_t, bool);
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
@@ -43,8 +43,10 @@ class ColorSpaceManager {
|
||||
/* Convert pixels in the specified colorspace to scene linear color for
|
||||
* rendering. Must be a colorspace returned from detect_known_colorspace. */
|
||||
template<typename T>
|
||||
static void to_scene_linear(
|
||||
ustring colorspace, T *pixels, size_t num_pixels, bool is_rgba, bool compress_as_srgb);
|
||||
static void to_scene_linear(ustring colorspace,
|
||||
T *pixels,
|
||||
size_t num_pixels,
|
||||
bool compress_as_srgb);
|
||||
|
||||
/* Efficiently convert pixels to scene linear colorspace at render time,
|
||||
* for OSL where the image texture cache contains original pixels. The
|
||||
|
@@ -441,13 +441,9 @@ void ConstantFolder::fold_mapping(NodeMappingType type) const
|
||||
if (is_zero(scale_in)) {
|
||||
make_zero();
|
||||
}
|
||||
else if (
|
||||
/* Can't constant fold since we always need to normalize the output. */
|
||||
(type != NODE_MAPPING_TYPE_NORMAL) &&
|
||||
/* Check all use values are zero, note location is not used by vector and normal types. */
|
||||
(is_zero(location_in) || type == NODE_MAPPING_TYPE_VECTOR ||
|
||||
type == NODE_MAPPING_TYPE_NORMAL) &&
|
||||
is_zero(rotation_in) && is_one(scale_in)) {
|
||||
else if ((is_zero(location_in) || type == NODE_MAPPING_TYPE_VECTOR ||
|
||||
type == NODE_MAPPING_TYPE_NORMAL) &&
|
||||
is_zero(rotation_in) && is_one(scale_in)) {
|
||||
try_bypass_or_make_constant(vector_in);
|
||||
}
|
||||
}
|
||||
|
@@ -236,7 +236,6 @@ void Geometry::compute_bvh(
|
||||
|
||||
BVHParams bparams;
|
||||
bparams.use_spatial_split = params->use_bvh_spatial_split;
|
||||
bparams.use_compact_structure = params->use_bvh_compact_structure;
|
||||
bparams.bvh_layout = bvh_layout;
|
||||
bparams.use_unaligned_nodes = dscene->data.bvh.have_curves &&
|
||||
params->use_bvh_unaligned_nodes;
|
||||
|
@@ -576,13 +576,13 @@ bool ImageManager::file_load_image(Image *img, int texture_limit)
|
||||
pixels[i * 4 + 3] = one;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (img->metadata.colorspace != u_colorspace_raw &&
|
||||
img->metadata.colorspace != u_colorspace_srgb) {
|
||||
/* Convert to scene linear. */
|
||||
ColorSpaceManager::to_scene_linear(
|
||||
img->metadata.colorspace, pixels, num_pixels, is_rgba, img->metadata.compress_as_srgb);
|
||||
if (img->metadata.colorspace != u_colorspace_raw &&
|
||||
img->metadata.colorspace != u_colorspace_srgb) {
|
||||
/* Convert to scene linear. */
|
||||
ColorSpaceManager::to_scene_linear(
|
||||
img->metadata.colorspace, pixels, num_pixels, img->metadata.compress_as_srgb);
|
||||
}
|
||||
}
|
||||
|
||||
/* Make sure we don't have buggy values. */
|
||||
@@ -891,10 +891,6 @@ void ImageManager::device_free(Device *device)
|
||||
void ImageManager::collect_statistics(RenderStats *stats)
|
||||
{
|
||||
foreach (const Image *image, images) {
|
||||
if (!image) {
|
||||
/* Image may have been freed due to lack of users. */
|
||||
continue;
|
||||
}
|
||||
stats->image.textures.add_entry(
|
||||
NamedSizeEntry(image->loader->name(), image->mem->memory_size()));
|
||||
}
|
||||
|
Some files were not shown because too many files have changed in this diff Show More
Reference in New Issue
Block a user