1
1

Compare commits

..

1 Commits

Author SHA1 Message Date
7f5219caa4 Migrating polyline shaders to create info.
GPU_SHADER_3D_POLYLINE_(FLAT/SMOOTH/UNIFORM/CLIPPED_UNIFORM)_COLOR

TODO:
* GPU_SHADER_3D_CLIPPED_UNIFORM_COLOR
* Check all usages of clipped uniform shaders.

{T95016}
2022-01-19 14:46:24 +01:00
1036 changed files with 36252 additions and 45968 deletions

View File

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

View File

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

View File

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

View File

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

View File

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

View File

@@ -94,4 +94,3 @@ download_source(POTRACE)
download_source(HARU)
download_source(ZSTD)
download_source(FLEX)
download_source(BROTLI)

View File

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

View File

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

View File

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

View File

@@ -83,9 +83,9 @@ else()
set(OPENEXR_VERSION_POSTFIX)
endif()
set(FREETYPE_VERSION 2.11.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)

View File

@@ -379,27 +379,27 @@ USE_CXX11=true
CLANG_FORMAT_VERSION_MIN="6.0"
CLANG_FORMAT_VERSION_MEX="10.0"
PYTHON_VERSION="3.10.2"
PYTHON_VERSION_SHORT="3.10"
PYTHON_VERSION_MIN="3.9"
PYTHON_VERSION_MEX="3.12"
PYTHON_VERSION="3.9.7"
PYTHON_VERSION_SHORT="3.9"
PYTHON_VERSION_MIN="3.7"
PYTHON_VERSION_MEX="3.11"
PYTHON_VERSION_INSTALLED=$PYTHON_VERSION_SHORT
PYTHON_FORCE_BUILD=false
PYTHON_FORCE_REBUILD=false
PYTHON_SKIP=false
# Additional Python modules.
PYTHON_IDNA_VERSION="3.3"
PYTHON_IDNA_VERSION="3.2"
PYTHON_IDNA_VERSION_MIN="2.0"
PYTHON_IDNA_VERSION_MEX="4.0"
PYTHON_IDNA_NAME="idna"
PYTHON_CHARSET_NORMALIZER_VERSION="2.0.10"
PYTHON_CHARSET_NORMALIZER_VERSION="2.0.6"
PYTHON_CHARSET_NORMALIZER_VERSION_MIN="2.0.6"
PYTHON_CHARSET_NORMALIZER_VERSION_MEX="2.1.0" # requests uses `charset_normalizer~=2.0.0`
PYTHON_CHARSET_NORMALIZER_NAME="charset-normalizer"
PYTHON_URLLIB3_VERSION="1.26.8"
PYTHON_URLLIB3_VERSION="1.26.7"
PYTHON_URLLIB3_VERSION_MIN="1.0"
PYTHON_URLLIB3_VERSION_MEX="2.0"
PYTHON_URLLIB3_NAME="urllib3"
@@ -409,17 +409,17 @@ PYTHON_CERTIFI_VERSION_MIN="2021.0"
PYTHON_CERTIFI_VERSION_MEX="2023.0"
PYTHON_CERTIFI_NAME="certifi"
PYTHON_REQUESTS_VERSION="2.27.1"
PYTHON_REQUESTS_VERSION="2.23.0"
PYTHON_REQUESTS_VERSION_MIN="2.0"
PYTHON_REQUESTS_VERSION_MEX="3.0"
PYTHON_REQUESTS_NAME="requests"
PYTHON_ZSTANDARD_VERSION="0.16.0"
PYTHON_ZSTANDARD_VERSION="0.15.2"
PYTHON_ZSTANDARD_VERSION_MIN="0.15.2"
PYTHON_ZSTANDARD_VERSION_MEX="0.20.0"
PYTHON_ZSTANDARD_VERSION_MEX="0.16.0"
PYTHON_ZSTANDARD_NAME="zstandard"
PYTHON_NUMPY_VERSION="1.22.0"
PYTHON_NUMPY_VERSION="1.21.2"
PYTHON_NUMPY_VERSION_MIN="1.14"
PYTHON_NUMPY_VERSION_MEX="2.0"
PYTHON_NUMPY_NAME="numpy"
@@ -499,7 +499,7 @@ LLVM_FORCE_REBUILD=false
LLVM_SKIP=false
# OSL needs to be compiled for now!
OSL_VERSION="1.11.17.0"
OSL_VERSION="1.11.14.1"
OSL_VERSION_SHORT="1.11"
OSL_VERSION_MIN="1.11"
OSL_VERSION_MEX="2.0"

View File

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

View File

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

View File

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

View File

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

View File

@@ -55,10 +55,6 @@ if(CMAKE_C_COMPILER_ID MATCHES "Clang")
message(WARNING "stripped pdb not supported with clang, disabling..")
set(WITH_WINDOWS_STRIPPED_PDB OFF)
endif()
else()
if(CMAKE_CXX_COMPILER_VERSION VERSION_LESS 19.28.29921) # MSVC 2019 16.9.16
message(FATAL_ERROR "Compiler is unsupported, MSVC 2019 16.9.16 or newer is required for building blender.")
endif()
endif()
if(NOT WITH_PYTHON_MODULE)
@@ -269,6 +265,12 @@ if(NOT DEFINED LIBDIR)
elseif(MSVC_VERSION GREATER 1919)
message(STATUS "Visual Studio 2019 detected.")
set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_BASE}_vc15)
elseif(MSVC_VERSION GREATER 1909)
message(STATUS "Visual Studio 2017 detected.")
set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_BASE}_vc15)
elseif(MSVC_VERSION EQUAL 1900)
message(STATUS "Visual Studio 2015 detected.")
set(LIBDIR ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_BASE}_vc15)
endif()
else()
message(STATUS "Using pre-compiled LIBDIR: ${LIBDIR}")
@@ -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)

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

@@ -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%" (

View File

@@ -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%" (

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

@@ -89,7 +89,7 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
sd->shader = kernel_tex_fetch(__tri_shader, sd->prim);
/* vectors */
sd->P = triangle_point_from_uv(kg, sd, isect->object, isect->prim, isect->u, isect->v);
sd->P = triangle_refine(kg, sd, ray->P, ray->D, isect->t, isect->object, isect->prim);
sd->Ng = Ng;
sd->N = Ng;

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

@@ -182,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);

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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

View File

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