blender
/
blender
136
406
Fork 589
Code Issues 6.2k Pull Requests 750 Projects 19 Wiki Activity

Merge branch 'master' into sculpt-dev

This commit is contained in:
Joseph Eagar 2021-10-23 23:15:36 -07:00
parent d21737189e 8c21667f3c
commit d294084432
281 changed files with 5549 additions and 2480 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
.github/pull_request_template.md vendored Normal file
View File

@ -0,0 +1,5 @@
This repository is only used as a mirror of git.blender.org. Blender development happens on
https://developer.blender.org.
To get started with contributing code, please see:
https://wiki.blender.org/wiki/Process/Contributing_Code

65
CMakeLists.txt
View File

@ -398,45 +398,55 @@ if(WITH_PYTHON_INSTALL)
set(PYTHON_REQUESTS_PATH "" CACHE PATH "Path to python site-packages or dist-packages containing 'requests' module")
mark_as_advanced(PYTHON_REQUESTS_PATH)
endif()
option(WITH_PYTHON_INSTALL_ZSTANDARD "Copy zstandard into the blender install folder" ON)
set(PYTHON_ZSTANDARD_PATH "" CACHE PATH "Path to python site-packages or dist-packages containing 'zstandard' module")
mark_as_advanced(PYTHON_ZSTANDARD_PATH)
endif()
option(WITH_CPU_SIMD "Enable SIMD instruction if they're detected on the host machine" ON)
mark_as_advanced(WITH_CPU_SIMD)
# Cycles
option(WITH_CYCLES "Enable Cycles Render Engine" ON)
option(WITH_CYCLES_STANDALONE "Build Cycles standalone application" OFF)
option(WITH_CYCLES_STANDALONE_GUI "Build Cycles standalone with GUI" OFF)
option(WITH_CYCLES_OSL "Build Cycles with OSL support" ON)
option(WITH_CYCLES_EMBREE "Build Cycles with Embree support" ON)
option(WITH_CYCLES_CUDA_BINARIES "Build Cycles CUDA binaries" OFF)
option(WITH_CYCLES_CUBIN_COMPILER "Build cubins with nvrtc based compiler instead of nvcc" OFF)
option(WITH_CYCLES_CUDA_BUILD_SERIAL "Build cubins one after another (useful on machines with limited RAM)" OFF)
mark_as_advanced(WITH_CYCLES_CUDA_BUILD_SERIAL)
set(CYCLES_TEST_DEVICES CPU CACHE STRING "Run regression tests on the specified device types (CPU CUDA OPTIX)" )
set(CYCLES_CUDA_BINARIES_ARCH sm_30 sm_35 sm_37 sm_50 sm_52 sm_60 sm_61 sm_70 sm_75 sm_86 compute_75 CACHE STRING "CUDA architectures to build binaries for")
mark_as_advanced(CYCLES_CUDA_BINARIES_ARCH)
option(WITH_CYCLES_HIP_BINARIES "Build Cycles HIP binaries" OFF)
unset(PLATFORM_DEFAULT)
option(WITH_CYCLES_LOGGING "Build Cycles with logging support" ON)
option(WITH_CYCLES_DEBUG_NAN "Build Cycles with additional asserts for detecting NaNs and invalid values" OFF)
option(WITH_CYCLES_NATIVE_ONLY "Build Cycles with native kernel only (which fits current CPU, use for development only)" OFF)
option(WITH_CYCLES_KERNEL_ASAN "Build Cycles kernels with address sanitizer when WITH_COMPILER_ASAN is on, even if it's very slow" OFF)
option(WITH_CYCLES "Enable Cycles Render Engine" ON)
option(WITH_CYCLES_OSL "Build Cycles with OpenShadingLanguage support" ON)
option(WITH_CYCLES_EMBREE "Build Cycles with Embree support" ON)
option(WITH_CYCLES_LOGGING "Build Cycles with logging support" ON)
option(WITH_CYCLES_STANDALONE "Build Cycles standalone application" OFF)
option(WITH_CYCLES_STANDALONE_GUI "Build Cycles standalone with GUI" OFF)
option(WITH_CYCLES_DEBUG_NAN "Build Cycles with additional asserts for detecting NaNs and invalid values" OFF)
option(WITH_CYCLES_NATIVE_ONLY "Build Cycles with native kernel only (which fits current CPU, use for development only)" OFF)
option(WITH_CYCLES_KERNEL_ASAN "Build Cycles kernels with address sanitizer when WITH_COMPILER_ASAN is on, even if it's very slow" OFF)
set(CYCLES_TEST_DEVICES CPU CACHE STRING "Run regression tests on the specified device types (CPU CUDA OPTIX HIP)" )
mark_as_advanced(WITH_CYCLES_KERNEL_ASAN)
mark_as_advanced(WITH_CYCLES_CUBIN_COMPILER)
mark_as_advanced(WITH_CYCLES_LOGGING)
mark_as_advanced(WITH_CYCLES_DEBUG_NAN)
mark_as_advanced(WITH_CYCLES_NATIVE_ONLY)
option(WITH_CYCLES_DEVICE_CUDA "Enable Cycles CUDA compute support" ON)
option(WITH_CYCLES_DEVICE_OPTIX "Enable Cycles OptiX support" ON)
option(WITH_CYCLES_DEVICE_HIP "Enable Cycles HIP support" ON)
mark_as_advanced(WITH_CYCLES_DEVICE_HIP)
# NVIDIA CUDA & OptiX
option(WITH_CYCLES_DEVICE_CUDA "Enable Cycles NVIDIA CUDA compute support" ON)
option(WITH_CYCLES_DEVICE_OPTIX "Enable Cycles NVIDIA OptiX support" ON)
mark_as_advanced(WITH_CYCLES_DEVICE_CUDA)
option(WITH_CUDA_DYNLOAD "Dynamically load CUDA libraries at runtime" ON)
option(WITH_CYCLES_CUDA_BINARIES "Build Cycles NVIDIA CUDA binaries" OFF)
set(CYCLES_CUDA_BINARIES_ARCH sm_30 sm_35 sm_37 sm_50 sm_52 sm_60 sm_61 sm_70 sm_75 sm_86 compute_75 CACHE STRING "CUDA architectures to build binaries for")
option(WITH_CYCLES_CUBIN_COMPILER "Build cubins with nvrtc based compiler instead of nvcc" OFF)
option(WITH_CYCLES_CUDA_BUILD_SERIAL "Build cubins one after another (useful on machines with limited RAM)" OFF)
option(WITH_CUDA_DYNLOAD "Dynamically load CUDA libraries at runtime (for developers, makes cuda-gdb work)" ON)
mark_as_advanced(CYCLES_CUDA_BINARIES_ARCH)
mark_as_advanced(WITH_CYCLES_CUBIN_COMPILER)
mark_as_advanced(WITH_CYCLES_CUDA_BUILD_SERIAL)
mark_as_advanced(WITH_CUDA_DYNLOAD)
# AMD HIP
option(WITH_CYCLES_DEVICE_HIP "Enable Cycles AMD HIP support" OFF)
option(WITH_CYCLES_HIP_BINARIES "Build Cycles AMD HIP binaries" OFF)
set(CYCLES_HIP_BINARIES_ARCH gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1034 CACHE STRING "AMD HIP architectures to build binaries for")
mark_as_advanced(WITH_CYCLES_DEVICE_HIP)
mark_as_advanced(CYCLES_HIP_BINARIES_ARCH)
# Draw Manager
option(WITH_DRAW_DEBUG "Add extra debug capabilities to Draw Manager" OFF)
mark_as_advanced(WITH_DRAW_DEBUG)
@ -1741,6 +1751,12 @@ if(WITH_PYTHON)
elseif(WITH_PYTHON_INSTALL_REQUESTS)
find_python_package(requests "")
endif()
if(WIN32 OR APPLE)
# pass, we have this in lib/python/site-packages
elseif(WITH_PYTHON_INSTALL_ZSTANDARD)
find_python_package(zstandard "")
endif()
endif()
# Select C++17 as the standard for C++ projects.
@ -2012,6 +2028,7 @@ if(FIRST_RUN)
endif()
info_cfg_option(WITH_PYTHON_INSTALL)
info_cfg_option(WITH_PYTHON_INSTALL_NUMPY)
info_cfg_option(WITH_PYTHON_INSTALL_ZSTANDARD)
info_cfg_option(WITH_PYTHON_MODULE)
info_cfg_option(WITH_PYTHON_SAFETY)

1
build_files/build_environment/cmake/numpy.cmake
View File

@ -38,7 +38,6 @@ ExternalProject_Add(external_numpy
PREFIX ${BUILD_DIR}/numpy
PATCH_COMMAND ${NUMPY_PATCH}
CONFIGURE_COMMAND ""
PATCH_COMMAND COMMAND ${PATCH_CMD} -p 1 -d ${BUILD_DIR}/numpy/src/external_numpy < ${PATCH_DIR}/numpy.diff
LOG_BUILD 1
BUILD_COMMAND ${PYTHON_BINARY} ${BUILD_DIR}/numpy/src/external_numpy/setup.py build ${NUMPY_BUILD_OPTION} install --old-and-unmanageable
INSTALL_COMMAND ""

10
build_files/build_environment/cmake/python_site_packages.cmake
View File

@ -18,14 +18,20 @@
if(WIN32 AND BUILD_MODE STREQUAL Debug)
set(SITE_PACKAGES_EXTRA --global-option build --global-option --debug)
# zstandard is determined to build and link release mode libs in a debug
# configuration, the only way to make it happy is to bend to its will
# and give it a library to link with.
set(PIP_CONFIGURE_COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/python/libs/python${PYTHON_SHORT_VERSION_NO_DOTS}_d.lib ${LIBDIR}/python/libs/python${PYTHON_SHORT_VERSION_NO_DOTS}.lib)
else()
set(PIP_CONFIGURE_COMMAND echo ".")
endif()
ExternalProject_Add(external_python_site_packages
DOWNLOAD_COMMAND ""
CONFIGURE_COMMAND ""
CONFIGURE_COMMAND ${PIP_CONFIGURE_COMMAND}
BUILD_COMMAND ""
PREFIX ${BUILD_DIR}/site_packages
INSTALL_COMMAND ${PYTHON_BINARY} -m pip install ${SITE_PACKAGES_EXTRA} cython==${CYTHON_VERSION} idna==${IDNA_VERSION} chardet==${CHARDET_VERSION} urllib3==${URLLIB3_VERSION} certifi==${CERTIFI_VERSION} requests==${REQUESTS_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)

23
build_files/build_environment/cmake/versions.cmake
View File

@ -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.9.2)
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 f0dc9000312abeb16de4eccce9a870ab)
set(PYTHON_HASH fddb060b483bc01850a3f412eea1d954)
set(PYTHON_HASH_TYPE MD5)
set(PYTHON_FILE Python-${PYTHON_VERSION}.tar.xz)
@ -215,17 +215,18 @@ set(NANOVDB_HASH e7b9e863ec2f3b04ead171dec2322807)
set(NANOVDB_HASH_TYPE MD5)
set(NANOVDB_FILE nano-vdb-${NANOVDB_GIT_UID}.tar.gz)
set(IDNA_VERSION 2.10)
set(CHARDET_VERSION 4.0.0)
set(URLLIB3_VERSION 1.26.3)
set(CERTIFI_VERSION 2020.12.5)
set(REQUESTS_VERSION 2.25.1)
set(CYTHON_VERSION 0.29.21)
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.26.0)
set(CYTHON_VERSION 0.29.24)
set(ZSTANDARD_VERSION 0.15.2 )
set(NUMPY_VERSION 1.19.5)
set(NUMPY_SHORT_VERSION 1.19)
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 f6a1b48717c552bbc18f1adc3cc1fe0e)
set(NUMPY_HASH 5638d5dae3ca387be562912312db842e)
set(NUMPY_HASH_TYPE MD5)
set(NUMPY_FILE numpy-${NUMPY_VERSION}.zip)

170
build_files/build_environment/install_deps.sh
View File

@ -371,71 +371,78 @@ NO_BUILD=false
NO_CONFIRM=false
USE_CXX11=true
# Note about versions: Min is inclusive, Max is exclusive (i.e. XXX_VERSION_MIN <= ACTUAL_VERSION < XXX_VERSION_MAX)
# Note about versions: Min is inclusive, Mex is 'minimum exclusive' (i.e. XXX_VERSION_MIN <= ACTUAL_VERSION < XXX_VERSION_MEX)
# XXX_VERSION is officially supported/used version in official builds.
# XXX_VERSION_SHORT is used for various things, like preferred version (when distribution provides several of them),
# and to name shortcuts to built libraries' installation directories...
CLANG_FORMAT_VERSION_MIN="6.0"
CLANG_FORMAT_VERSION_MAX="10.0"
CLANG_FORMAT_VERSION_MEX="10.0"
PYTHON_VERSION="3.9.2"
PYTHON_VERSION="3.9.7"
PYTHON_VERSION_SHORT="3.9"
PYTHON_VERSION_MIN="3.7"
PYTHON_VERSION_MAX="3.11"
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="2.9"
PYTHON_IDNA_VERSION="3.2"
PYTHON_IDNA_VERSION_MIN="2.0"
PYTHON_IDNA_VERSION_MAX="3.0"
PYTHON_IDNA_VERSION_MEX="4.0"
PYTHON_IDNA_NAME="idna"
PYTHON_CHARDET_VERSION="3.0.4"
PYTHON_CHARDET_VERSION_MIN="3.0"
PYTHON_CHARDET_VERSION_MAX="5.0"
PYTHON_CHARDET_NAME="chardet"
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.25.9"
PYTHON_URLLIB3_VERSION="1.26.7"
PYTHON_URLLIB3_VERSION_MIN="1.0"
PYTHON_URLLIB3_VERSION_MAX="2.0"
PYTHON_URLLIB3_VERSION_MEX="2.0"
PYTHON_URLLIB3_NAME="urllib3"
PYTHON_CERTIFI_VERSION="2020.4.5.2"
PYTHON_CERTIFI_VERSION_MIN="2020.0"
PYTHON_CERTIFI_VERSION_MAX="2021.0"
PYTHON_CERTIFI_VERSION="2021.10.8"
PYTHON_CERTIFI_VERSION_MIN="2021.0"
PYTHON_CERTIFI_VERSION_MEX="2023.0"
PYTHON_CERTIFI_NAME="certifi"
PYTHON_REQUESTS_VERSION="2.23.0"
PYTHON_REQUESTS_VERSION_MIN="2.0"
PYTHON_REQUESTS_VERSION_MAX="3.0"
PYTHON_REQUESTS_VERSION_MEX="3.0"
PYTHON_REQUESTS_NAME="requests"
PYTHON_NUMPY_VERSION="1.19.5"
PYTHON_ZSTANDARD_VERSION="0.15.2"
PYTHON_ZSTANDARD_VERSION_MIN="0.15.2"
PYTHON_ZSTANDARD_VERSION_MEX="0.16.0"
PYTHON_ZSTANDARD_NAME="zstandard"
PYTHON_NUMPY_VERSION="1.21.2"
PYTHON_NUMPY_VERSION_MIN="1.14"
PYTHON_NUMPY_VERSION_MAX="2.0"
PYTHON_NUMPY_VERSION_MEX="2.0"
PYTHON_NUMPY_NAME="numpy"
# As package-ready parameters (only used with distro packages).
PYTHON_MODULES_PACKAGES=(
"$PYTHON_IDNA_NAME $PYTHON_IDNA_VERSION_MIN $PYTHON_IDNA_VERSION_MAX"
"$PYTHON_CHARDET_NAME $PYTHON_CHARDET_VERSION_MIN $PYTHON_CHARDET_VERSION_MAX"
"$PYTHON_URLLIB3_NAME $PYTHON_URLLIB3_VERSION_MIN $PYTHON_URLLIB3_VERSION_MAX"
"$PYTHON_CERTIFI_NAME $PYTHON_CERTIFI_VERSION_MIN $PYTHON_CERTIFI_VERSION_MAX"
"$PYTHON_REQUESTS_NAME $PYTHON_REQUESTS_VERSION_MIN $PYTHON_REQUESTS_VERSION_MAX"
"$PYTHON_NUMPY_NAME $PYTHON_NUMPY_VERSION_MIN $PYTHON_NUMPY_VERSION_MAX"
"$PYTHON_IDNA_NAME $PYTHON_IDNA_VERSION_MIN $PYTHON_IDNA_VERSION_MEX"
"$PYTHON_CHARSET_NORMALIZER_NAME $PYTHON_CHARSET_NORMALIZER_VERSION_MIN $PYTHON_CHARSET_NORMALIZER_VERSION_MEX"
"$PYTHON_URLLIB3_NAME $PYTHON_URLLIB3_VERSION_MIN $PYTHON_URLLIB3_VERSION_MEX"
"$PYTHON_CERTIFI_NAME $PYTHON_CERTIFI_VERSION_MIN $PYTHON_CERTIFI_VERSION_MEX"
"$PYTHON_REQUESTS_NAME $PYTHON_REQUESTS_VERSION_MIN $PYTHON_REQUESTS_VERSION_MEX"
"$PYTHON_ZSTANDARD_NAME $PYTHON_ZSTANDARD_VERSION_MIN $PYTHON_ZSTANDARD_VERSION_MEX"
"$PYTHON_NUMPY_NAME $PYTHON_NUMPY_VERSION_MIN $PYTHON_NUMPY_VERSION_MEX"
)
# As pip-ready parameters (only used when building python).
PYTHON_MODULES_PIP=(
"$PYTHON_IDNA_NAME==$PYTHON_IDNA_VERSION"
"$PYTHON_CHARDET_NAME==$PYTHON_CHARDET_VERSION"
"$PYTHON_CHARSET_NORMALIZER_NAME==$PYTHON_CHARSET_NORMALIZER_VERSION"
"$PYTHON_URLLIB3_NAME==$PYTHON_URLLIB3_VERSION"
"$PYTHON_CERTIFI_NAME==$PYTHON_CERTIFI_VERSION"
"$PYTHON_REQUESTS_NAME==$PYTHON_REQUESTS_VERSION"
"$PYTHON_ZSTANDARD_NAME==$PYTHON_ZSTANDARD_VERSION"
"$PYTHON_NUMPY_NAME==$PYTHON_NUMPY_VERSION"
)
@ -443,7 +450,7 @@ PYTHON_MODULES_PIP=(
BOOST_VERSION="1.73.0"
BOOST_VERSION_SHORT="1.73"
BOOST_VERSION_MIN="1.49"
BOOST_VERSION_MAX="2.0"
BOOST_VERSION_MEX="2.0"
BOOST_FORCE_BUILD=false
BOOST_FORCE_REBUILD=false
BOOST_SKIP=false
@ -452,7 +459,7 @@ TBB_VERSION="2020"
TBB_VERSION_SHORT="2020"
TBB_VERSION_UPDATE="_U2" # Used for source packages...
TBB_VERSION_MIN="2018"
TBB_VERSION_MAX="2022"
TBB_VERSION_MEX="2022"
TBB_FORCE_BUILD=false
TBB_FORCE_REBUILD=false
TBB_SKIP=false
@ -460,7 +467,7 @@ TBB_SKIP=false
OCIO_VERSION="2.0.0"
OCIO_VERSION_SHORT="2.0"
OCIO_VERSION_MIN="2.0"
OCIO_VERSION_MAX="3.0"
OCIO_VERSION_MEX="3.0"
OCIO_FORCE_BUILD=false
OCIO_FORCE_REBUILD=false
OCIO_SKIP=false
@ -468,7 +475,7 @@ OCIO_SKIP=false
OPENEXR_VERSION="2.5.5"
OPENEXR_VERSION_SHORT="2.5"
OPENEXR_VERSION_MIN="2.4"
OPENEXR_VERSION_MAX="3.0"
OPENEXR_VERSION_MEX="3.0"
OPENEXR_FORCE_BUILD=false
OPENEXR_FORCE_REBUILD=false
OPENEXR_SKIP=false
@ -477,7 +484,7 @@ _with_built_openexr=false
OIIO_VERSION="2.2.15.1"
OIIO_VERSION_SHORT="2.2"
OIIO_VERSION_MIN="2.1.12"
OIIO_VERSION_MAX="2.3.0"
OIIO_VERSION_MEX="2.3.0"
OIIO_FORCE_BUILD=false
OIIO_FORCE_REBUILD=false
OIIO_SKIP=false
@ -485,7 +492,7 @@ OIIO_SKIP=false
LLVM_VERSION="12.0.0"
LLVM_VERSION_SHORT="12.0"
LLVM_VERSION_MIN="11.0"
LLVM_VERSION_MAX="13.0"
LLVM_VERSION_MEX="13.0"
LLVM_VERSION_FOUND=""
LLVM_FORCE_BUILD=false
LLVM_FORCE_REBUILD=false
@ -495,7 +502,7 @@ LLVM_SKIP=false
OSL_VERSION="1.11.14.1"
OSL_VERSION_SHORT="1.11"
OSL_VERSION_MIN="1.11"
OSL_VERSION_MAX="2.0"
OSL_VERSION_MEX="2.0"
OSL_FORCE_BUILD=false
OSL_FORCE_REBUILD=false
OSL_SKIP=false
@ -504,7 +511,7 @@ OSL_SKIP=false
OSD_VERSION="3.4.3"
OSD_VERSION_SHORT="3.4"
OSD_VERSION_MIN="3.4"
OSD_VERSION_MAX="4.0"
OSD_VERSION_MEX="4.0"
OSD_FORCE_BUILD=false
OSD_FORCE_REBUILD=false
OSD_SKIP=false
@ -515,7 +522,7 @@ OPENVDB_BLOSC_VERSION="1.5.0"
OPENVDB_VERSION="8.0.1"
OPENVDB_VERSION_SHORT="8.0"
OPENVDB_VERSION_MIN="8.0"
OPENVDB_VERSION_MAX="8.1"
OPENVDB_VERSION_MEX="8.1"
OPENVDB_FORCE_BUILD=false
OPENVDB_FORCE_REBUILD=false
OPENVDB_SKIP=false
@ -524,7 +531,7 @@ OPENVDB_SKIP=false
ALEMBIC_VERSION="1.7.16"
ALEMBIC_VERSION_SHORT="1.7"
ALEMBIC_VERSION_MIN="1.7"
ALEMBIC_VERSION_MAX="2.0"
ALEMBIC_VERSION_MEX="2.0"
ALEMBIC_FORCE_BUILD=false
ALEMBIC_FORCE_REBUILD=false
ALEMBIC_SKIP=false
@ -532,7 +539,7 @@ ALEMBIC_SKIP=false
USD_VERSION="21.02"
USD_VERSION_SHORT="21.02"
USD_VERSION_MIN="20.05"
USD_VERSION_MAX="22.00"
USD_VERSION_MEX="22.00"
USD_FORCE_BUILD=false
USD_FORCE_REBUILD=false
USD_SKIP=false
@ -540,7 +547,7 @@ USD_SKIP=false
OPENCOLLADA_VERSION="1.6.68"
OPENCOLLADA_VERSION_SHORT="1.6"
OPENCOLLADA_VERSION_MIN="1.6.68"
OPENCOLLADA_VERSION_MAX="1.7"
OPENCOLLADA_VERSION_MEX="1.7"
OPENCOLLADA_FORCE_BUILD=false
OPENCOLLADA_FORCE_REBUILD=false
OPENCOLLADA_SKIP=false
@ -548,7 +555,7 @@ OPENCOLLADA_SKIP=false
EMBREE_VERSION="3.10.0"
EMBREE_VERSION_SHORT="3.10"
EMBREE_VERSION_MIN="3.10"
EMBREE_VERSION_MAX="4.0"
EMBREE_VERSION_MEX="4.0"
EMBREE_FORCE_BUILD=false
EMBREE_FORCE_REBUILD=false
EMBREE_SKIP=false
@ -556,7 +563,7 @@ EMBREE_SKIP=false
OIDN_VERSION="1.4.1"
OIDN_VERSION_SHORT="1.4"
OIDN_VERSION_MIN="1.4.0"
OIDN_VERSION_MAX="1.5"
OIDN_VERSION_MEX="1.5"
OIDN_FORCE_BUILD=false
OIDN_FORCE_REBUILD=false
OIDN_SKIP=false
@ -566,7 +573,7 @@ ISPC_VERSION="1.16.0"
FFMPEG_VERSION="4.4"
FFMPEG_VERSION_SHORT="4.4"
FFMPEG_VERSION_MIN="3.0"
FFMPEG_VERSION_MAX="5.0"
FFMPEG_VERSION_MEX="5.0"
FFMPEG_FORCE_BUILD=false
FFMPEG_FORCE_REBUILD=false
FFMPEG_SKIP=false
@ -575,7 +582,7 @@ _ffmpeg_list_sep=";"
XR_OPENXR_VERSION="1.0.17"
XR_OPENXR_VERSION_SHORT="1.0"
XR_OPENXR_VERSION_MIN="1.0.8"
XR_OPENXR_VERSION_MAX="2.0"
XR_OPENXR_VERSION_MEX="2.0"
XR_OPENXR_FORCE_BUILD=false
XR_OPENXR_FORCE_REBUILD=false
XR_OPENXR_SKIP=false
@ -1141,10 +1148,11 @@ You may also want to build them yourself (optional ones are [between brackets]):
* Python $PYTHON_VERSION (from $PYTHON_SOURCE).
** [IDNA $PYTHON_IDNA_VERSION] (use pip).
** [Chardet $PYTHON_CHARDET_VERSION] (use pip).
** [Charset Normalizer $PYTHON_CHARSET_NORMALIZER_VERSION] (use pip).
** [Urllib3 $PYTHON_URLLIB3_VERSION] (use pip).
** [Certifi $PYTHON_CERTIFI_VERSION] (use pip).
** [Requests $PYTHON_REQUESTS_VERSION] (use pip).
** [ZStandard $PYTHON_ZSTANDARD_VERSION] (use pip).
** [NumPy $PYTHON_NUMPY_VERSION] (use pip).
* Boost $BOOST_VERSION (from $BOOST_SOURCE, modules: $BOOST_BUILD_MODULES).
* TBB $TBB_VERSION (from $TBB_SOURCE).
@ -4028,7 +4036,7 @@ install_DEB() {
INFO "Forced Python building, as requested..."
_do_compile_python=true
else
check_package_version_ge_lt_DEB python3-dev $PYTHON_VERSION_MIN $PYTHON_VERSION_MAX
check_package_version_ge_lt_DEB python3-dev $PYTHON_VERSION_MIN $PYTHON_VERSION_MEX
if [ $? -eq 0 ]; then
PYTHON_VERSION_INSTALLED=$(echo `get_package_version_DEB python3-dev` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
@ -4041,8 +4049,8 @@ install_DEB() {
module=($module)
package="python3-${module[0]}"
package_vmin=${module[1]}
package_vmax=${module[2]}
check_package_version_ge_lt_DEB "$package" $package_vmin $package_vmax
package_vmex=${module[2]}
check_package_version_ge_lt_DEB "$package" $package_vmin $package_vmex
if [ $? -eq 0 ]; then
install_packages_DEB "$package"
else
@ -4068,7 +4076,7 @@ install_DEB() {
INFO "Forced Boost building, as requested..."
compile_Boost
else
check_package_version_ge_lt_DEB libboost-dev $BOOST_VERSION_MIN $BOOST_VERSION_MAX
check_package_version_ge_lt_DEB libboost-dev $BOOST_VERSION_MIN $BOOST_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_DEB libboost-dev
@ -4089,7 +4097,7 @@ install_DEB() {
INFO "Forced TBB building, as requested..."
compile_TBB
else
check_package_version_ge_lt_DEB libtbb-dev $TBB_VERSION_MIN $TBB_VERSION_MAX
check_package_version_ge_lt_DEB libtbb-dev $TBB_VERSION_MIN $TBB_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_DEB libtbb-dev
clean_TBB
@ -4106,7 +4114,7 @@ install_DEB() {
INFO "Forced OpenColorIO building, as requested..."
compile_OCIO
else
check_package_version_ge_lt_DEB libopencolorio-dev $OCIO_VERSION_MIN $OCIO_VERSION_MAX
check_package_version_ge_lt_DEB libopencolorio-dev $OCIO_VERSION_MIN $OCIO_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_DEB libopencolorio-dev
clean_OCIO
@ -4123,7 +4131,7 @@ install_DEB() {
INFO "Forced ILMBase/OpenEXR building, as requested..."
compile_OPENEXR
else
check_package_version_ge_lt_DEB libopenexr-dev $OPENEXR_VERSION_MIN $OPENEXR_VERSION_MAX
check_package_version_ge_lt_DEB libopenexr-dev $OPENEXR_VERSION_MIN $OPENEXR_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_DEB libopenexr-dev
OPENEXR_VERSION=`get_package_version_DEB libopenexr-dev`
@ -4144,7 +4152,7 @@ install_DEB() {
INFO "Forced OpenImageIO building, as requested..."
compile_OIIO
else
check_package_version_ge_lt_DEB libopenimageio-dev $OIIO_VERSION_MIN $OIIO_VERSION_MAX
check_package_version_ge_lt_DEB libopenimageio-dev $OIIO_VERSION_MIN $OIIO_VERSION_MEX
if [ $? -eq 0 -a "$_with_built_openexr" = false ]; then
install_packages_DEB libopenimageio-dev openimageio-tools
clean_OIIO
@ -4164,7 +4172,7 @@ install_DEB() {
INFO "Forced LLVM building, as requested..."
_do_compile_llvm=true
else
check_package_version_ge_lt_DEB llvm-dev $LLVM_VERSION_MIN $LLVM_VERSION_MAX
check_package_version_ge_lt_DEB llvm-dev $LLVM_VERSION_MIN $LLVM_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_DEB llvm-dev clang libclang-dev
have_llvm=true
@ -4195,7 +4203,7 @@ install_DEB() {
INFO "Forced OpenShadingLanguage building, as requested..."
_do_compile_osl=true
else
check_package_version_ge_lt_DEB libopenshadinglanguage-dev $OSL_VERSION_MIN $OSL_VERSION_MAX
check_package_version_ge_lt_DEB libopenshadinglanguage-dev $OSL_VERSION_MIN $OSL_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_DEB libopenshadinglanguage-dev
clean_OSL
@ -4233,7 +4241,7 @@ install_DEB() {
INFO "Forced OpenVDB building, as requested..."
compile_OPENVDB
else
check_package_version_ge_lt_DEB libopenvdb-dev $OPENVDB_VERSION_MIN $OPENVDB_VERSION_MAX
check_package_version_ge_lt_DEB libopenvdb-dev $OPENVDB_VERSION_MIN $OPENVDB_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_DEB libopenvdb-dev libblosc-dev
clean_OPENVDB
@ -4295,7 +4303,7 @@ install_DEB() {
_do_compile_embree=true
else
# There is a package, but it does not provide everything that Blender needs...
#~ check_package_version_ge_lt_DEB libembree-dev $EMBREE_VERSION_MIN $EMBREE_VERSION_MAX
#~ check_package_version_ge_lt_DEB libembree-dev $EMBREE_VERSION_MIN $EMBREE_VERSION_MEX
#~ if [ $? -eq 0 ]; then
#~ install_packages_DEB libembree-dev
#~ clean_Embree
@ -4337,7 +4345,7 @@ install_DEB() {
# XXX Debian Testing / Ubuntu 16.04 finally includes FFmpeg, so check as usual
check_package_DEB ffmpeg
if [ $? -eq 0 ]; then
check_package_version_ge_lt_DEB ffmpeg $FFMPEG_VERSION_MIN $FFMPEG_VERSION_MAX
check_package_version_ge_lt_DEB ffmpeg $FFMPEG_VERSION_MIN $FFMPEG_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_DEB libavdevice-dev
clean_FFmpeg
@ -4671,7 +4679,7 @@ install_RPM() {
INFO "Forced Python building, as requested..."
_do_compile_python=true
else
check_package_version_ge_lt_RPM python3-devel $PYTHON_VERSION_MIN $PYTHON_VERSION_MAX
check_package_version_ge_lt_RPM python3-devel $PYTHON_VERSION_MIN $PYTHON_VERSION_MEX
if [ $? -eq 0 ]; then
PYTHON_VERSION_INSTALLED=$(echo `get_package_version_RPM python3-devel` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
@ -4683,8 +4691,8 @@ install_RPM() {
module=($module)
package="python3-${module[0]}"
package_vmin=${module[1]}
package_vmax=${module[2]}
check_package_version_ge_lt_RPM "$package" $package_vmin $package_vmax
package_vmex=${module[2]}
check_package_version_ge_lt_RPM "$package" $package_vmin $package_vmex
if [ $? -eq 0 ]; then
install_packages_RPM "$package"
else
@ -4711,7 +4719,7 @@ install_RPM() {
INFO "Forced Boost building, as requested..."
_do_compile_boost=true
else
check_package_version_ge_lt_RPM boost-devel $BOOST_VERSION_MIN $BOOST_VERSION_MAX
check_package_version_ge_lt_RPM boost-devel $BOOST_VERSION_MIN $BOOST_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_RPM boost-devel
clean_Boost
@ -4738,7 +4746,7 @@ install_RPM() {
INFO "Forced TBB building, as requested..."
compile_TBB
else
check_package_version_ge_lt_RPM tbb-devel $TBB_VERSION_MIN $TBB_VERSION_MAX
check_package_version_ge_lt_RPM tbb-devel $TBB_VERSION_MIN $TBB_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_RPM tbb-devel
clean_TBB
@ -4756,7 +4764,7 @@ install_RPM() {
compile_OCIO
else
if [ "$RPM" = "SUSE" ]; then
check_package_version_ge_lt_RPM OpenColorIO-devel $OCIO_VERSION_MIN $OCIO_VERSION_MAX
check_package_version_ge_lt_RPM OpenColorIO-devel $OCIO_VERSION_MIN $OCIO_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_RPM OpenColorIO-devel
clean_OCIO
@ -4776,7 +4784,7 @@ install_RPM() {
INFO "Forced ILMBase/OpenEXR building, as requested..."
compile_OPENEXR
else
check_package_version_ge_lt_RPM openexr-devel $OPENEXR_VERSION_MIN $OPENEXR_VERSION_MAX
check_package_version_ge_lt_RPM openexr-devel $OPENEXR_VERSION_MIN $OPENEXR_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_RPM openexr-devel
OPENEXR_VERSION=`get_package_version_RPM openexr-devel`
@ -4794,7 +4802,7 @@ install_RPM() {
INFO "Forced OpenImageIO building, as requested..."
compile_OIIO
else
check_package_version_ge_lt_RPM OpenImageIO-devel $OIIO_VERSION_MIN $OIIO_VERSION_MAX
check_package_version_ge_lt_RPM OpenImageIO-devel $OIIO_VERSION_MIN $OIIO_VERSION_MEX
if [ $? -eq 0 -a $_with_built_openexr == false ]; then
install_packages_RPM OpenImageIO-devel OpenImageIO-utils
clean_OIIO
@ -4819,7 +4827,7 @@ install_RPM() {
else
CLANG_DEV="clang-devel"
fi
check_package_version_ge_lt_RPM llvm-devel $LLVM_VERSION_MIN $LLVM_VERSION_MAX
check_package_version_ge_lt_RPM llvm-devel $LLVM_VERSION_MIN $LLVM_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_RPM llvm-devel $CLANG_DEV
have_llvm=true
@ -4855,7 +4863,7 @@ install_RPM() {
else
OSL_DEV="openshadinglanguage-devel"
fi
check_package_version_ge_lt_RPM $OSL_DEV $OSL_VERSION_MIN $OSL_VERSION_MAX
check_package_version_ge_lt_RPM $OSL_DEV $OSL_VERSION_MIN $OSL_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_RPM $OSL_DEV
clean_OSL
@ -4950,7 +4958,7 @@ install_RPM() {
_do_compile_embree=true
else
# There is a package, but it does not provide everything that Blender needs...
#~ check_package_version_ge_lt_RPM embree-devel $EMBREE_VERSION_MIN $EMBREE_VERSION_MAX
#~ check_package_version_ge_lt_RPM embree-devel $EMBREE_VERSION_MIN $EMBREE_VERSION_MEX
#~ if [ $? -eq 0 ]; then
#~ install_packages_RPM embree-devel
#~ clean_Embree
@ -4989,7 +4997,7 @@ install_RPM() {
INFO "Forced FFMpeg building, as requested..."
compile_FFmpeg
else
check_package_version_ge_lt_RPM ffmpeg-devel $FFMPEG_VERSION_MIN $FFMPEG_VERSION_MAX
check_package_version_ge_lt_RPM ffmpeg-devel $FFMPEG_VERSION_MIN $FFMPEG_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_RPM ffmpeg ffmpeg-devel
clean_FFmpeg
@ -5214,7 +5222,7 @@ install_ARCH() {
INFO "Forced Python building, as requested..."
_do_compile_python=true
else
check_package_version_ge_lt_ARCH python $PYTHON_VERSION_MIN $PYTHON_VERSION_MAX
check_package_version_ge_lt_ARCH python $PYTHON_VERSION_MIN $PYTHON_VERSION_MEX
if [ $? -eq 0 ]; then
PYTHON_VERSION_INSTALLED=$(echo `get_package_version_ARCH python` | sed -r 's/^([0-9]+\.[0-9]+).*/\1/')
@ -5227,8 +5235,8 @@ install_ARCH() {
module=($module)
package="python-${module[0]}"
package_vmin=${module[1]}
package_vmax=${module[2]}
check_package_version_ge_lt_ARCH "$package" $package_vmin $package_vmax
package_vmex=${module[2]}
check_package_version_ge_lt_ARCH "$package" $package_vmin $package_vmex
if [ $? -eq 0 ]; then
install_packages_ARCH "$package"
else
@ -5254,7 +5262,7 @@ install_ARCH() {
INFO "Forced Boost building, as requested..."
compile_Boost
else
check_package_version_ge_lt_ARCH boost $BOOST_VERSION_MIN $BOOST_VERSION_MAX
check_package_version_ge_lt_ARCH boost $BOOST_VERSION_MIN $BOOST_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_ARCH boost
clean_Boost
@ -5271,7 +5279,7 @@ install_ARCH() {
INFO "Forced TBB building, as requested..."
compile_TBB
else
check_package_version_ge_lt_ARCH intel-tbb $TBB_VERSION_MIN $TBB_VERSION_MAX
check_package_version_ge_lt_ARCH intel-tbb $TBB_VERSION_MIN $TBB_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_ARCH intel-tbb
clean_TBB
@ -5288,7 +5296,7 @@ install_ARCH() {
INFO "Forced OpenColorIO building, as requested..."
compile_OCIO
else
check_package_version_ge_lt_ARCH opencolorio $OCIO_VERSION_MIN $OCIO_VERSION_MAX
check_package_version_ge_lt_ARCH opencolorio $OCIO_VERSION_MIN $OCIO_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_ARCH opencolorio
clean_OCIO
@ -5305,7 +5313,7 @@ install_ARCH() {
INFO "Forced ILMBase/OpenEXR building, as requested..."
compile_OPENEXR
else
check_package_version_ge_lt_ARCH openexr $OPENEXR_VERSION_MIN $OPENEXR_VERSION_MAX
check_package_version_ge_lt_ARCH openexr $OPENEXR_VERSION_MIN $OPENEXR_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_ARCH openexr
OPENEXR_VERSION=`get_package_version_ARCH openexr`
@ -5324,7 +5332,7 @@ install_ARCH() {
INFO "Forced OpenImageIO building, as requested..."
compile_OIIO
else
check_package_version_ge_lt_ARCH openimageio $OIIO_VERSION_MIN $OIIO_VERSION_MAX
check_package_version_ge_lt_ARCH openimageio $OIIO_VERSION_MIN $OIIO_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_ARCH openimageio
clean_OIIO
@ -5344,7 +5352,7 @@ install_ARCH() {
INFO "Forced LLVM building, as requested..."
_do_compile_llvm=true
else
check_package_version_ge_lt_ARCH llvm $LLVM_VERSION_MIN $LLVM_VERSION_MAX
check_package_version_ge_lt_ARCH llvm $LLVM_VERSION_MIN $LLVM_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_ARCH llvm clang
have_llvm=true
@ -5375,7 +5383,7 @@ install_ARCH() {
INFO "Forced OpenShadingLanguage building, as requested..."
_do_compile_osl=true
else
check_package_version_ge_lt_ARCH openshadinglanguage $OSL_VERSION_MIN $OSL_VERSION_MAX
check_package_version_ge_lt_ARCH openshadinglanguage $OSL_VERSION_MIN $OSL_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_ARCH openshadinglanguage
clean_OSL
@ -5401,7 +5409,7 @@ install_ARCH() {
INFO "Forced OpenSubdiv building, as requested..."
compile_OSD
else
check_package_version_ge_lt_ARCH opensubdiv $OSD_VERSION_MIN $OSD_VERSION_MAX
check_package_version_ge_lt_ARCH opensubdiv $OSD_VERSION_MIN $OSD_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_ARCH opensubdiv
clean_OSD
@ -5418,7 +5426,7 @@ install_ARCH() {
INFO "Forced OpenVDB building, as requested..."
compile_OPENVDB
else
check_package_version_ge_lt_ARCH openvdb $OPENVDB_VERSION_MIN $OPENVDB_VERSION_MAX
check_package_version_ge_lt_ARCH openvdb $OPENVDB_VERSION_MIN $OPENVDB_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_ARCH openvdb
clean_OPENVDB
@ -5484,7 +5492,7 @@ install_ARCH() {
_do_compile_embree=true
else
# There is a package, but it does not provide everything that Blender needs...
#~ check_package_version_ge_lt_ARCH embree $EMBREE_VERSION_MIN $EMBREE_VERSION_MAX
#~ check_package_version_ge_lt_ARCH embree $EMBREE_VERSION_MIN $EMBREE_VERSION_MEX
#~ if [ $? -eq 0 ]; then
#~ install_packages_ARCH embree
#~ clean_Embree
@ -5523,7 +5531,7 @@ install_ARCH() {
INFO "Forced FFMpeg building, as requested..."
compile_FFmpeg
else
check_package_version_ge_lt_ARCH ffmpeg $FFMPEG_VERSION_MIN $FFMPEG_VERSION_MAX
check_package_version_ge_lt_ARCH ffmpeg $FFMPEG_VERSION_MIN $FFMPEG_VERSION_MEX
if [ $? -eq 0 ]; then
install_packages_ARCH ffmpeg
clean_FFmpeg

27
build_files/build_environment/patches/numpy.diff
View File

@ -1,27 +0,0 @@
diff --git a/numpy/distutils/system_info.py b/numpy/distutils/system_info.py
index ba2b1f4..b10f7df 100644
--- a/numpy/distutils/system_info.py
+++ b/numpy/distutils/system_info.py
@@ -2164,8 +2164,8 @@ class accelerate_info(system_info):
'accelerate' in libraries):
if intel:
args.extend(['-msse3'])
- else:
- args.extend(['-faltivec'])
+# else:
+# args.extend(['-faltivec'])
args.extend([
'-I/System/Library/Frameworks/vecLib.framework/Headers'])
link_args.extend(['-Wl,-framework', '-Wl,Accelerate'])
@@ -2174,8 +2174,8 @@ class accelerate_info(system_info):
'veclib' in libraries):
if intel:
args.extend(['-msse3'])
- else:
- args.extend(['-faltivec'])
+# else:
+# args.extend(['-faltivec'])
args.extend([
'-I/System/Library/Frameworks/vecLib.framework/Headers'])
link_args.extend(['-Wl,-framework', '-Wl,vecLib'])

3
build_files/build_environment/windows/build_deps.cmd
View File

@ -79,6 +79,9 @@ set STAGING=%BUILD_DIR%\S
rem for python module build
set MSSdk=1
set DISTUTILS_USE_SDK=1
rem if you let pip pick its own build dirs, it'll stick it somewhere deep inside the user profile
rem and cython will refuse to link due to a path that gets too long.
set TMPDIR=c:\t\
rem for python externals source to be shared between the various archs and compilers
mkdir %BUILD_DIR%\downloads\externals

2
build_files/windows/find_dependencies.cmd
View File

@ -3,7 +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)
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\37\bin\python.exe
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
if NOT "%verbose%" == "" (
echo svn : "%SVN%"
echo cmake : "%CMAKE%"

2
build_files/windows/format.cmd
View File

@ -10,7 +10,7 @@ exit /b 1
echo found clang-format in %CF_PATH%
if EXIST %PYTHON% (
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\37\bin\python.exe
set PYTHON=%BLENDER_DIR%\..\lib\win64_vc15\python\39\bin\python.exe
goto detect_python_done
)

2
extern/audaspace/bindings/python/PySound.cpp vendored
View File

@ -290,7 +290,7 @@ PyDoc_STRVAR(M_aud_Sound_buffer_doc,
".. classmethod:: buffer(data, rate)\n\n"
" Creates a sound from a data buffer.\n\n"
" :arg data: The data as two dimensional numpy array.\n"
" :type data: numpy.ndarray\n"
" :type data: :class:`numpy.ndarray`\n"
" :arg rate: The sample rate.\n"
" :type rate: double\n"
" :return: The created :class:`Sound` object.\n"

141
extern/hipew/include/hipew.h vendored
View File

@ -425,6 +425,105 @@ typedef struct HIPdevprop_st {
int textureAlign;
} HIPdevprop;
typedef struct {
// 32-bit Atomics
unsigned hasGlobalInt32Atomics : 1; ///< 32-bit integer atomics for global memory.
unsigned hasGlobalFloatAtomicExch : 1; ///< 32-bit float atomic exch for global memory.
unsigned hasSharedInt32Atomics : 1; ///< 32-bit integer atomics for shared memory.
unsigned hasSharedFloatAtomicExch : 1; ///< 32-bit float atomic exch for shared memory.
unsigned hasFloatAtomicAdd : 1; ///< 32-bit float atomic add in global and shared memory.
// 64-bit Atomics
unsigned hasGlobalInt64Atomics : 1; ///< 64-bit integer atomics for global memory.
unsigned hasSharedInt64Atomics : 1; ///< 64-bit integer atomics for shared memory.
// Doubles
unsigned hasDoubles : 1; ///< Double-precision floating point.
// Warp cross-lane operations
unsigned hasWarpVote : 1; ///< Warp vote instructions (__any, __all).
unsigned hasWarpBallot : 1; ///< Warp ballot instructions (__ballot).
unsigned hasWarpShuffle : 1; ///< Warp shuffle operations. (__shfl_*).
unsigned hasFunnelShift : 1; ///< Funnel two words into one with shift&mask caps.
// Sync
unsigned hasThreadFenceSystem : 1; ///< __threadfence_system.
unsigned hasSyncThreadsExt : 1; ///< __syncthreads_count, syncthreads_and, syncthreads_or.
// Misc
unsigned hasSurfaceFuncs : 1; ///< Surface functions.
unsigned has3dGrid : 1; ///< Grid and group dims are 3D (rather than 2D).
unsigned hasDynamicParallelism : 1; ///< Dynamic parallelism.
} hipDeviceArch_t;
typedef struct hipDeviceProp_t {
char name[256]; ///< Device name.
size_t totalGlobalMem; ///< Size of global memory region (in bytes).
size_t sharedMemPerBlock; ///< Size of shared memory region (in bytes).
int regsPerBlock; ///< Registers per block.
int warpSize; ///< Warp size.
int maxThreadsPerBlock; ///< Max work items per work group or workgroup max size.
int maxThreadsDim[3]; ///< Max number of threads in each dimension (XYZ) of a block.
int maxGridSize[3]; ///< Max grid dimensions (XYZ).
int clockRate; ///< Max clock frequency of the multiProcessors in khz.
int memoryClockRate; ///< Max global memory clock frequency in khz.
int memoryBusWidth; ///< Global memory bus width in bits.
size_t totalConstMem; ///< Size of shared memory region (in bytes).
int major; ///< Major compute capability. On HCC, this is an approximation and features may
///< differ from CUDA CC. See the arch feature flags for portable ways to query
///< feature caps.
int minor; ///< Minor compute capability. On HCC, this is an approximation and features may
///< differ from CUDA CC. See the arch feature flags for portable ways to query
///< feature caps.
int multiProcessorCount; ///< Number of multi-processors (compute units).
int l2CacheSize; ///< L2 cache size.
int maxThreadsPerMultiProcessor; ///< Maximum resident threads per multi-processor.
int computeMode; ///< Compute mode.
int clockInstructionRate; ///< Frequency in khz of the timer used by the device-side "clock*"
///< instructions. New for HIP.
hipDeviceArch_t arch; ///< Architectural feature flags. New for HIP.
int concurrentKernels; ///< Device can possibly execute multiple kernels concurrently.
int pciDomainID; ///< PCI Domain ID
int pciBusID; ///< PCI Bus ID.
int pciDeviceID; ///< PCI Device ID.
size_t maxSharedMemoryPerMultiProcessor; ///< Maximum Shared Memory Per Multiprocessor.
int isMultiGpuBoard; ///< 1 if device is on a multi-GPU board, 0 if not.
int canMapHostMemory; ///< Check whether HIP can map host memory
int gcnArch; ///< DEPRECATED: use gcnArchName instead
char gcnArchName[256]; ///< AMD GCN Arch Name.
int integrated; ///< APU vs dGPU
int cooperativeLaunch; ///< HIP device supports cooperative launch
int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices
int maxTexture1DLinear; ///< Maximum size for 1D textures bound to linear memory
int maxTexture1D; ///< Maximum number of elements in 1D images
int maxTexture2D[2]; ///< Maximum dimensions (width, height) of 2D images, in image elements
int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image elements
unsigned int* hdpMemFlushCntl; ///< Addres of HDP_MEM_COHERENCY_FLUSH_CNTL register
unsigned int* hdpRegFlushCntl; ///< Addres of HDP_REG_COHERENCY_FLUSH_CNTL register
size_t memPitch; ///<Maximum pitch in bytes allowed by memory copies
size_t textureAlignment; ///<Alignment requirement for textures
size_t texturePitchAlignment; ///<Pitch alignment requirement for texture references bound to pitched memory
int kernelExecTimeoutEnabled; ///<Run time limit for kernels executed on the device
int ECCEnabled; ///<Device has ECC support enabled
int tccDriver; ///< 1:If device is Tesla device using TCC driver, else 0
int cooperativeMultiDeviceUnmatchedFunc; ///< HIP device supports cooperative launch on multiple
///devices with unmatched functions
int cooperativeMultiDeviceUnmatchedGridDim; ///< HIP device supports cooperative launch on multiple
///devices with unmatched grid dimensions
int cooperativeMultiDeviceUnmatchedBlockDim; ///< HIP device supports cooperative launch on multiple
///devices with unmatched block dimensions
int cooperativeMultiDeviceUnmatchedSharedMem; ///< HIP device supports cooperative launch on multiple
///devices with unmatched shared memories
int isLargeBar; ///< 1: if it is a large PCI bar device, else 0
int asicRevision; ///< Revision of the GPU in this device
int managedMemory; ///< Device supports allocating managed memory on this system
int directManagedMemAccessFromHost; ///< Host can directly access managed memory on the device without migration
int concurrentManagedAccess; ///< Device can coherently access managed memory concurrently with the CPU
int pageableMemoryAccess; ///< Device supports coherently accessing pageable memory
///< without calling hipHostRegister on it
int pageableMemoryAccessUsesHostPageTables; ///< Device accesses pageable memory via the host's page tables
} hipDeviceProp_t;
typedef enum HIPpointer_attribute_enum {
HIP_POINTER_ATTRIBUTE_CONTEXT = 1,
HIP_POINTER_ATTRIBUTE_MEMORY_TYPE = 2,
@ -951,6 +1050,25 @@ typedef enum HIPGLmap_flags_enum {
HIP_GL_MAP_RESOURCE_FLAGS_WRITE_DISCARD = 0x02,
} HIPGLmap_flags;
/**
* hipRTC related
*/
typedef struct _hiprtcProgram* hiprtcProgram;
typedef enum hiprtcResult {
HIPRTC_SUCCESS = 0,
HIPRTC_ERROR_OUT_OF_MEMORY = 1,
HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2,
HIPRTC_ERROR_INVALID_INPUT = 3,
HIPRTC_ERROR_INVALID_PROGRAM = 4,
HIPRTC_ERROR_INVALID_OPTION = 5,
HIPRTC_ERROR_COMPILATION = 6,
HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7,
HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8,
HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9,
HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10,
HIPRTC_ERROR_INTERNAL_ERROR = 11
} hiprtcResult;
/* Function types. */
typedef hipError_t HIPAPI thipGetErrorName(hipError_t error, const char** pStr);
@ -958,6 +1076,7 @@ typedef hipError_t HIPAPI thipInit(unsigned int Flags);
typedef hipError_t HIPAPI thipDriverGetVersion(int* driverVersion);
typedef hipError_t HIPAPI thipGetDevice(hipDevice_t* device, int ordinal);
typedef hipError_t HIPAPI thipGetDeviceCount(int* count);
typedef hipError_t HIPAPI thipGetDeviceProperties(hipDeviceProp_t* props, int deviceId);
typedef hipError_t HIPAPI thipDeviceGetName(char* name, int len, hipDevice_t dev);
typedef hipError_t HIPAPI thipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attrib, hipDevice_t dev);
typedef hipError_t HIPAPI thipDeviceComputeCapability(int* major, int* minor, hipDevice_t dev);
@ -1071,6 +1190,16 @@ typedef hipError_t HIPAPI thipGraphicsMapResources(unsigned int count, hipGraphi
typedef hipError_t HIPAPI thipGraphicsUnmapResources(unsigned int count, hipGraphicsResource* resources, hipStream_t hStream);
typedef hipError_t HIPAPI thipGraphicsGLRegisterBuffer(hipGraphicsResource* pCudaResource, GLuint buffer, unsigned int Flags);
typedef hipError_t HIPAPI thipGLGetDevices(unsigned int* pHipDeviceCount, int* pHipDevices, unsigned int hipDeviceCount, hipGLDeviceList deviceList);
typedef hiprtcResult HIPAPI thiprtcGetErrorString(hiprtcResult result);
typedef hiprtcResult HIPAPI thiprtcAddNameExpression(hiprtcProgram prog, const char* name_expression);
typedef hiprtcResult HIPAPI thiprtcCompileProgram(hiprtcProgram prog, int numOptions, const char** options);
typedef hiprtcResult HIPAPI thiprtcCreateProgram(hiprtcProgram* prog, const char* src, const char* name, int numHeaders, const char** headers, const char** includeNames);
typedef hiprtcResult HIPAPI thiprtcDestroyProgram(hiprtcProgram* prog);
typedef hiprtcResult HIPAPI thiprtcGetLoweredName(hiprtcProgram prog, const char* name_expression, const char** lowered_name);
typedef hiprtcResult HIPAPI thiprtcGetProgramLog(hiprtcProgram prog, char* log);
typedef hiprtcResult HIPAPI thiprtcGetProgramLogSize(hiprtcProgram prog, size_t* logSizeRet);
typedef hiprtcResult HIPAPI thiprtcGetCode(hiprtcProgram prog, char* code);
typedef hiprtcResult HIPAPI thiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet);
/* Function declarations. */
@ -1079,6 +1208,7 @@ extern thipInit *hipInit;
extern thipDriverGetVersion *hipDriverGetVersion;
extern thipGetDevice *hipGetDevice;
extern thipGetDeviceCount *hipGetDeviceCount;
extern thipGetDeviceProperties *hipGetDeviceProperties;
extern thipDeviceGetName *hipDeviceGetName;
extern thipDeviceGetAttribute *hipDeviceGetAttribute;
extern thipDeviceComputeCapability *hipDeviceComputeCapability;
@ -1187,6 +1317,17 @@ extern thipGraphicsUnmapResources *hipGraphicsUnmapResources;
extern thipGraphicsGLRegisterBuffer *hipGraphicsGLRegisterBuffer;
extern thipGLGetDevices *hipGLGetDevices;
extern thiprtcGetErrorString* hiprtcGetErrorString;
extern thiprtcAddNameExpression* hiprtcAddNameExpression;
extern thiprtcCompileProgram* hiprtcCompileProgram;
extern thiprtcCreateProgram* hiprtcCreateProgram;
extern thiprtcDestroyProgram* hiprtcDestroyProgram;
extern thiprtcGetLoweredName* hiprtcGetLoweredName;
extern thiprtcGetProgramLog* hiprtcGetProgramLog;
extern thiprtcGetProgramLogSize* hiprtcGetProgramLogSize;
extern thiprtcGetCode* hiprtcGetCode;
extern thiprtcGetCodeSize* hiprtcGetCodeSize;
enum {
HIPEW_SUCCESS = 0,

23
extern/hipew/src/hipew.c vendored
View File

@ -70,6 +70,7 @@ thipInit *hipInit;
thipDriverGetVersion *hipDriverGetVersion;
thipGetDevice *hipGetDevice;
thipGetDeviceCount *hipGetDeviceCount;
thipGetDeviceProperties *hipGetDeviceProperties;
thipDeviceGetName *hipDeviceGetName;
thipDeviceGetAttribute *hipDeviceGetAttribute;
thipDeviceComputeCapability *hipDeviceComputeCapability;
@ -178,6 +179,17 @@ thipGraphicsResourceGetMappedPointer *hipGraphicsResourceGetMappedPointer;
thipGraphicsGLRegisterBuffer *hipGraphicsGLRegisterBuffer;
thipGLGetDevices *hipGLGetDevices;
thiprtcGetErrorString* hiprtcGetErrorString;
thiprtcAddNameExpression* hiprtcAddNameExpression;
thiprtcCompileProgram* hiprtcCompileProgram;
thiprtcCreateProgram* hiprtcCreateProgram;
thiprtcDestroyProgram* hiprtcDestroyProgram;
thiprtcGetLoweredName* hiprtcGetLoweredName;
thiprtcGetProgramLog* hiprtcGetProgramLog;
thiprtcGetProgramLogSize* hiprtcGetProgramLogSize;
thiprtcGetCode* hiprtcGetCode;
thiprtcGetCodeSize* hiprtcGetCodeSize;
static DynamicLibrary dynamic_library_open_find(const char **paths) {
@ -242,6 +254,7 @@ static int hipewHipInit(void) {
HIP_LIBRARY_FIND_CHECKED(hipDriverGetVersion);
HIP_LIBRARY_FIND_CHECKED(hipGetDevice);
HIP_LIBRARY_FIND_CHECKED(hipGetDeviceCount);
HIP_LIBRARY_FIND_CHECKED(hipGetDeviceProperties);
HIP_LIBRARY_FIND_CHECKED(hipDeviceGetName);
HIP_LIBRARY_FIND_CHECKED(hipDeviceGetAttribute);
HIP_LIBRARY_FIND_CHECKED(hipDeviceComputeCapability);
@ -346,6 +359,16 @@ static int hipewHipInit(void) {
HIP_LIBRARY_FIND_CHECKED(hipGraphicsGLRegisterBuffer);
HIP_LIBRARY_FIND_CHECKED(hipGLGetDevices);
#endif
HIP_LIBRARY_FIND_CHECKED(hiprtcGetErrorString);
HIP_LIBRARY_FIND_CHECKED(hiprtcAddNameExpression);
HIP_LIBRARY_FIND_CHECKED(hiprtcCompileProgram);
HIP_LIBRARY_FIND_CHECKED(hiprtcCreateProgram);
HIP_LIBRARY_FIND_CHECKED(hiprtcDestroyProgram);
HIP_LIBRARY_FIND_CHECKED(hiprtcGetLoweredName);
HIP_LIBRARY_FIND_CHECKED(hiprtcGetProgramLog);
HIP_LIBRARY_FIND_CHECKED(hiprtcGetProgramLogSize);
HIP_LIBRARY_FIND_CHECKED(hiprtcGetCode);
HIP_LIBRARY_FIND_CHECKED(hiprtcGetCodeSize);
result = HIPEW_SUCCESS;
return result;
}

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

@ -1329,7 +1329,7 @@ class CyclesPreferences(bpy.types.AddonPreferences):
elif entry.type == 'CPU':
cpu_devices.append(entry)
# Extend all GPU devices with CPU.
if compute_device_type != 'CPU' and compute_device_type != 'HIP':
if compute_device_type != 'CPU':
devices.extend(cpu_devices)
return devices
@ -1373,8 +1373,18 @@ class CyclesPreferences(bpy.types.AddonPreferences):
if not found_device:
col = box.column(align=True)
col.label(text="No compatible GPUs found for path tracing", icon='INFO')
col.label(text="Cycles will render on the CPU", icon='BLANK1')
col.label(text="No compatible GPUs found for Cycles", icon='INFO')
if device_type == 'CUDA':
col.label(text="Requires NVIDIA GPU with compute capability 3.0", icon='BLANK1')
elif device_type == 'OPTIX':
col.label(text="Requires NVIDIA GPU with compute capability 5.0", icon='BLANK1')
col.label(text="and NVIDIA driver version 470 or newer", icon='BLANK1')
elif device_type == 'HIP':
import sys
col.label(text="Requires discrete AMD GPU with ??? architecture", icon='BLANK1')
if sys.platform[:3] == "win":
col.label(text="and AMD driver version ??? or newer", icon='BLANK1')
return
for device in devices:

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

@ -50,9 +50,7 @@ CPUKernels::CPUKernels()
REGISTER_KERNEL(adaptive_sampling_filter_x),
REGISTER_KERNEL(adaptive_sampling_filter_y),
/* Cryptomatte. */
REGISTER_KERNEL(cryptomatte_postprocess),
/* Bake. */
REGISTER_KERNEL(bake)
REGISTER_KERNEL(cryptomatte_postprocess)
{
}

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

@ -102,10 +102,6 @@ class CPUKernels {
CryptomattePostprocessFunction cryptomatte_postprocess;
/* Bake. */
CPUKernelFunction<void (*)(const KernelGlobalsCPU *, float *, int, int, int, int, int)> bake;
CPUKernels();
};

6
intern/cycles/device/cuda/device_impl.cpp
View File

@ -454,7 +454,7 @@ bool CUDADevice::load_kernels(const uint kernel_features)
return (result == CUDA_SUCCESS);
}
void CUDADevice::reserve_local_memory(const uint /* kernel_features */)
void CUDADevice::reserve_local_memory(const uint kernel_features)
{
/* Together with CU_CTX_LMEM_RESIZE_TO_MAX, this reserves local memory
* needed for kernel launches, so that we can reliably figure out when
@ -468,7 +468,9 @@ void CUDADevice::reserve_local_memory(const uint /* kernel_features */)
{
/* Use the biggest kernel for estimation. */
const DeviceKernel test_kernel = DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE;
const DeviceKernel test_kernel = (kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) ?
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE :
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE;
/* Launch kernel, using just 1 block appears sufficient to reserve memory for all
* multiprocessors. It would be good to do this in parallel for the multi GPU case

16
intern/cycles/device/cuda/queue.cpp
View File

@ -41,13 +41,19 @@ CUDADeviceQueue::~CUDADeviceQueue()
int CUDADeviceQueue::num_concurrent_states(const size_t state_size) const
{
int num_states = max(cuda_device_->get_num_multiprocessors() *
cuda_device_->get_max_num_threads_per_multiprocessor() * 16,
1048576);
const int max_num_threads = cuda_device_->get_num_multiprocessors() *
cuda_device_->get_max_num_threads_per_multiprocessor();
int num_states = max(max_num_threads, 65536) * 16;
const char *factor_str = getenv("CYCLES_CONCURRENT_STATES_FACTOR");
if (factor_str) {
num_states = max((int)(num_states * atof(factor_str)), 1024);
const float factor = (float)atof(factor_str);
if (factor != 0.0f) {
num_states = max((int)(num_states * factor), 1024);
}
else {
VLOG(3) << "CYCLES_CONCURRENT_STATES_FACTOR evaluated to 0";
}
}
VLOG(3) << "GPU queue concurrent states: " << num_states << ", using up to "
@ -107,6 +113,8 @@ bool CUDADeviceQueue::enqueue(DeviceKernel kernel, const int work_size, void *ar
case DEVICE_KERNEL_INTEGRATOR_TERMINATED_PATHS_ARRAY:
case DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY:
case DEVICE_KERNEL_INTEGRATOR_COMPACT_PATHS_ARRAY:
case DEVICE_KERNEL_INTEGRATOR_TERMINATED_SHADOW_PATHS_ARRAY:
case DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_PATHS_ARRAY:
/* See parall_active_index.h for why this amount of shared memory is needed. */
shared_mem_bytes = (num_threads_per_block + 1) * sizeof(int);
break;

6
intern/cycles/device/device_kernel.cpp
View File

@ -64,6 +64,12 @@ const char *device_kernel_as_string(DeviceKernel kernel)
return "integrator_compact_paths_array";
case DEVICE_KERNEL_INTEGRATOR_COMPACT_STATES:
return "integrator_compact_states";
case DEVICE_KERNEL_INTEGRATOR_TERMINATED_SHADOW_PATHS_ARRAY:
return "integrator_terminated_shadow_paths_array";
case DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_PATHS_ARRAY:
return "integrator_compact_shadow_paths_array";
case DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_STATES:
return "integrator_compact_shadow_states";
case DEVICE_KERNEL_INTEGRATOR_RESET:
return "integrator_reset";
case DEVICE_KERNEL_INTEGRATOR_SHADOW_CATCHER_COUNT_POSSIBLE_SPLITS:

30
intern/cycles/device/hip/device_impl.cpp
View File

@ -208,7 +208,7 @@ bool HIPDevice::use_adaptive_compilation()
return DebugFlags().hip.adaptive_compile;
}
/* Common NVCC flags which stays the same regardless of shading model,
/* Common HIPCC flags which stays the same regardless of shading model,
* kernel sources md5 and only depends on compiler or compilation settings.
*/
string HIPDevice::compile_kernel_get_common_cflags(const uint kernel_features)
@ -239,11 +239,13 @@ string HIPDevice::compile_kernel(const uint kernel_features,
int major, minor;
hipDeviceGetAttribute(&major, hipDeviceAttributeComputeCapabilityMajor, hipDevId);
hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
hipDeviceProp_t props;
hipGetDeviceProperties(&props, hipDevId);
/* Attempt to use kernel provided with Blender. */
if (!use_adaptive_compilation()) {
if (!force_ptx) {
const string fatbin = path_get(string_printf("lib/%s_sm_%d%d.cubin", name, major, minor));
const string fatbin = path_get(string_printf("lib/%s_%s.fatbin", name, props.gcnArchName));
VLOG(1) << "Testing for pre-compiled kernel " << fatbin << ".";
if (path_exists(fatbin)) {
VLOG(1) << "Using precompiled kernel.";
@ -283,17 +285,21 @@ string HIPDevice::compile_kernel(const uint kernel_features,
const string kernel_md5 = util_md5_string(source_md5 + common_cflags);
const char *const kernel_ext = "genco";
std::string options;
# ifdef _WIN32
const char *const options =
"save-temps -Wno-parentheses-equality -Wno-unused-value --hipcc-func-supp";
options.append("Wno-parentheses-equality -Wno-unused-value --hipcc-func-supp -ffast-math");
# else
const char *const options =
"save-temps -Wno-parentheses-equality -Wno-unused-value --hipcc-func-supp -O3 -ggdb";
options.append("Wno-parentheses-equality -Wno-unused-value --hipcc-func-supp -O3 -ffast-math");
# endif
# ifdef _DEBUG
options.append(" -save-temps");
# endif
options.append(" --amdgpu-target=").append(props.gcnArchName);
const string include_path = source_path;
const char *const kernel_arch = force_ptx ? "compute" : "sm";
const char *const kernel_arch = props.gcnArchName;
const string fatbin_file = string_printf(
"cycles_%s_%s_%d%d_%s", name, kernel_arch, major, minor, kernel_md5.c_str());
"cycles_%s_%s_%s", name, kernel_arch, kernel_md5.c_str());
const string fatbin = path_cache_get(path_join("kernels", fatbin_file));
VLOG(1) << "Testing for locally compiled kernel " << fatbin << ".";
if (path_exists(fatbin)) {
@ -350,7 +356,7 @@ string HIPDevice::compile_kernel(const uint kernel_features,
string command = string_printf("%s -%s -I %s --%s %s -o \"%s\"",
hipcc,
options,
options.c_str(),
include_path.c_str(),
kernel_ext,
source_path.c_str(),
@ -430,7 +436,7 @@ bool HIPDevice::load_kernels(const uint kernel_features)
return (result == hipSuccess);
}
void HIPDevice::reserve_local_memory(const uint)
void HIPDevice::reserve_local_memory(const uint kernel_features)
{
/* Together with hipDeviceLmemResizeToMax, this reserves local memory
* needed for kernel launches, so that we can reliably figure out when
@ -444,7 +450,9 @@ void HIPDevice::reserve_local_memory(const uint)
{
/* Use the biggest kernel for estimation. */
const DeviceKernel test_kernel = DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE;
const DeviceKernel test_kernel = (kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) ?
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE :
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE;
/* Launch kernel, using just 1 block appears sufficient to reserve memory for all
* multiprocessors. It would be good to do this in parallel for the multi GPU case

19
intern/cycles/device/hip/queue.cpp
View File

@ -41,22 +41,19 @@ HIPDeviceQueue::~HIPDeviceQueue()
int HIPDeviceQueue::num_concurrent_states(const size_t state_size) const
{
int num_states = 0;
const int max_num_threads = hip_device_->get_num_multiprocessors() *
hip_device_->get_max_num_threads_per_multiprocessor();
if (max_num_threads == 0) {
num_states = 1048576; // 65536 * 16
}
else {
num_states = max_num_threads * 16;
}
int num_states = ((max_num_threads == 0) ? 65536 : max_num_threads) * 16;
const char *factor_str = getenv("CYCLES_CONCURRENT_STATES_FACTOR");
if (factor_str) {
float factor = atof(factor_str);
if (!factor)
const float factor = (float)atof(factor_str);
if (factor != 0.0f) {
num_states = max((int)(num_states * factor), 1024);
}
else {
VLOG(3) << "CYCLES_CONCURRENT_STATES_FACTOR evaluated to 0";
num_states = max((int)(num_states * factor), 1024);
}
}
VLOG(3) << "GPU queue concurrent states: " << num_states << ", using up to "
@ -116,6 +113,8 @@ bool HIPDeviceQueue::enqueue(DeviceKernel kernel, const int work_size, void *arg
case DEVICE_KERNEL_INTEGRATOR_TERMINATED_PATHS_ARRAY:
case DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY:
case DEVICE_KERNEL_INTEGRATOR_COMPACT_PATHS_ARRAY:
case DEVICE_KERNEL_INTEGRATOR_TERMINATED_SHADOW_PATHS_ARRAY:
case DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_PATHS_ARRAY:
/* See parall_active_index.h for why this amount of shared memory is needed. */
shared_mem_bytes = (num_threads_per_block + 1) * sizeof(int);
break;

3
intern/cycles/device/optix/device_impl.cpp
View File

@ -377,9 +377,6 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
group_descs[PG_CALL_SVM_BEVEL].callables.moduleDC = optix_module;
group_descs[PG_CALL_SVM_BEVEL].callables.entryFunctionNameDC =
"__direct_callable__svm_node_bevel";
group_descs[PG_CALL_AO_PASS].kind = OPTIX_PROGRAM_GROUP_KIND_CALLABLES;
group_descs[PG_CALL_AO_PASS].callables.moduleDC = optix_module;
group_descs[PG_CALL_AO_PASS].callables.entryFunctionNameDC = "__direct_callable__ao_pass";
}
optix_assert(optixProgramGroupCreate(

1
intern/cycles/device/optix/device_impl.h
View File

@ -45,7 +45,6 @@ enum {
PG_HITS_MOTION,
PG_CALL_SVM_AO,
PG_CALL_SVM_BEVEL,
PG_CALL_AO_PASS,
NUM_PROGRAM_GROUPS
};

2
intern/cycles/integrator/pass_accessor.cpp
View File

@ -115,7 +115,7 @@ static void pad_pixels(const BufferParams &buffer_params,
}
if (destination.pixels_half_rgba) {
const half one = float_to_half(1.0f);
const half one = float_to_half_display(1.0f);
half4 *pixel = destination.pixels_half_rgba + destination.offset;
for (size_t i = 0; i < size; i++, pixel++) {

4
intern/cycles/integrator/pass_accessor_cpu.cpp
View File

@ -148,8 +148,8 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_half_rgba(
film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel_rgba);
float4_store_half(&pixel->x,
make_float4(pixel_rgba[0], pixel_rgba[1], pixel_rgba[2], pixel_rgba[3]));
*pixel = float4_to_half4_display(
make_float4(pixel_rgba[0], pixel_rgba[1], pixel_rgba[2], pixel_rgba[3]));
}
});
}

215
intern/cycles/integrator/path_trace_work_gpu.cpp
View File

@ -78,24 +78,25 @@ PathTraceWorkGPU::PathTraceWorkGPU(Device *device,
integrator_shader_sort_counter_(device, "integrator_shader_sort_counter", MEM_READ_WRITE),
integrator_shader_raytrace_sort_counter_(
device, "integrator_shader_raytrace_sort_counter", MEM_READ_WRITE),
integrator_shader_sort_prefix_sum_(
device, "integrator_shader_sort_prefix_sum", MEM_READ_WRITE),
integrator_next_main_path_index_(device, "integrator_next_main_path_index", MEM_READ_WRITE),
integrator_next_shadow_path_index_(
device, "integrator_next_shadow_path_index", MEM_READ_WRITE),
integrator_next_shadow_catcher_path_index_(
device, "integrator_next_shadow_catcher_path_index", MEM_READ_WRITE),
queued_paths_(device, "queued_paths", MEM_READ_WRITE),
num_queued_paths_(device, "num_queued_paths", MEM_READ_WRITE),
work_tiles_(device, "work_tiles", MEM_READ_WRITE),
display_rgba_half_(device, "display buffer half", MEM_READ_WRITE),
max_num_paths_(queue_->num_concurrent_states(estimate_single_state_size())),
min_num_active_paths_(queue_->num_concurrent_busy_states()),
max_active_path_index_(0)
min_num_active_main_paths_(queue_->num_concurrent_busy_states()),
max_active_main_path_index_(0)
{
memset(&integrator_state_gpu_, 0, sizeof(integrator_state_gpu_));
/* Limit number of active paths to the half of the overall state. This is due to the logic in the
* path compaction which relies on the fact that regeneration does not happen sooner than half of
* the states are available again. */
min_num_active_paths_ = min(min_num_active_paths_, max_num_paths_ / 2);
min_num_active_main_paths_ = min(min_num_active_main_paths_, max_num_paths_ / 2);
}
void PathTraceWorkGPU::alloc_integrator_soa()
@ -200,6 +201,9 @@ void PathTraceWorkGPU::alloc_integrator_sorting()
integrator_shader_raytrace_sort_counter_.alloc(max_shaders);
integrator_shader_raytrace_sort_counter_.zero_to_device();
integrator_shader_sort_prefix_sum_.alloc(max_shaders);
integrator_shader_sort_prefix_sum_.zero_to_device();
integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] =
(int *)integrator_shader_sort_counter_.device_pointer;
integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE] =
@ -217,13 +221,13 @@ void PathTraceWorkGPU::alloc_integrator_path_split()
(int *)integrator_next_shadow_path_index_.device_pointer;
}
if (integrator_next_shadow_catcher_path_index_.size() == 0) {
integrator_next_shadow_catcher_path_index_.alloc(1);
if (integrator_next_main_path_index_.size() == 0) {
integrator_next_main_path_index_.alloc(1);
integrator_next_shadow_path_index_.data()[0] = 0;
integrator_next_shadow_catcher_path_index_.zero_to_device();
integrator_next_main_path_index_.zero_to_device();
integrator_state_gpu_.next_shadow_catcher_path_index =
(int *)integrator_next_shadow_catcher_path_index_.device_pointer;
integrator_state_gpu_.next_main_path_index =
(int *)integrator_next_main_path_index_.device_pointer;
}
}
@ -298,7 +302,7 @@ void PathTraceWorkGPU::render_samples(RenderStatistics &statistics,
break;
}
num_busy_accum += get_num_active_paths();
num_busy_accum += num_active_main_paths_paths();
++num_iterations;
}
@ -357,26 +361,16 @@ bool PathTraceWorkGPU::enqueue_path_iteration()
return false;
}
/* If the number of shadow kernels dropped to zero, set the next shadow path
* index to zero as well.
*
* TODO: use shadow path compaction to lower it more often instead of letting
* it fill up entirely? */
const int num_queued_shadow =
queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW] +
queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW];
if (num_queued_shadow == 0) {
if (integrator_next_shadow_path_index_.data()[0] != 0) {
integrator_next_shadow_path_index_.data()[0] = 0;
queue_->copy_to_device(integrator_next_shadow_path_index_);
}
}
/* For kernels that add shadow paths, check if there is enough space available.
* If not, schedule shadow kernels first to clear out the shadow paths. */
int num_paths_limit = INT_MAX;
if (kernel_creates_shadow_paths(kernel)) {
if (max_num_paths_ - integrator_next_shadow_path_index_.data()[0] <
queue_counter->num_queued[kernel]) {
compact_shadow_paths();
const int available_shadow_paths = max_num_paths_ -
integrator_next_shadow_path_index_.data()[0];
if (available_shadow_paths < queue_counter->num_queued[kernel]) {
if (queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW]) {
enqueue_path_iteration(DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW);
return true;
@ -386,10 +380,14 @@ bool PathTraceWorkGPU::enqueue_path_iteration()
return true;
}
}
else if (kernel_creates_ao_paths(kernel)) {
/* AO kernel creates two shadow paths, so limit number of states to schedule. */
num_paths_limit = available_shadow_paths / 2;
}
}
/* Schedule kernel with maximum number of queued items. */
enqueue_path_iteration(kernel);
enqueue_path_iteration(kernel, num_paths_limit);
/* Update next shadow path index for kernels that can add shadow paths. */
if (kernel_creates_shadow_paths(kernel)) {
@ -399,12 +397,12 @@ bool PathTraceWorkGPU::enqueue_path_iteration()
return true;
}
void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel)
void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num_paths_limit)
{
void *d_path_index = (void *)NULL;
/* Create array of path indices for which this kernel is queued to be executed. */
int work_size = kernel_max_active_path_index(kernel);
int work_size = kernel_max_active_main_path_index(kernel);
IntegratorQueueCounter *queue_counter = integrator_queue_counter_.data();
int num_queued = queue_counter->num_queued[kernel];
@ -414,7 +412,8 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel)
work_size = num_queued;
d_path_index = (void *)queued_paths_.device_pointer;
compute_sorted_queued_paths(DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY, kernel);
compute_sorted_queued_paths(
DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY, kernel, num_paths_limit);
}
else if (num_queued < work_size) {
work_size = num_queued;
@ -430,6 +429,8 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel)
}
}
work_size = min(work_size, num_paths_limit);
DCHECK_LE(work_size, max_num_paths_);
switch (kernel) {
@ -464,17 +465,20 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel)
}
}
void PathTraceWorkGPU::compute_sorted_queued_paths(DeviceKernel kernel, DeviceKernel queued_kernel)
void PathTraceWorkGPU::compute_sorted_queued_paths(DeviceKernel kernel,
DeviceKernel queued_kernel,
const int num_paths_limit)
{
int d_queued_kernel = queued_kernel;
void *d_counter = integrator_state_gpu_.sort_key_counter[d_queued_kernel];
assert(d_counter != nullptr);
void *d_prefix_sum = (void *)integrator_shader_sort_prefix_sum_.device_pointer;
assert(d_counter != nullptr && d_prefix_sum != nullptr);
/* Compute prefix sum of number of active paths with each shader. */
{
const int work_size = 1;
int max_shaders = device_scene_->data.max_shaders;
void *args[] = {&d_counter, &max_shaders};
void *args[] = {&d_counter, &d_prefix_sum, &max_shaders};
queue_->enqueue(DEVICE_KERNEL_PREFIX_SUM, work_size, args);
}
@ -483,29 +487,24 @@ void PathTraceWorkGPU::compute_sorted_queued_paths(DeviceKernel kernel, DeviceKe
/* Launch kernel to fill the active paths arrays. */
{
/* TODO: this could be smaller for terminated paths based on amount of work we want
* to schedule. */
const int work_size = kernel_max_active_path_index(queued_kernel);
* to schedule, and also based on num_paths_limit.
*
* Also, when the number paths is limited it may be better to prefer paths from the
* end of the array since compaction would need to do less work. */
const int work_size = kernel_max_active_main_path_index(queued_kernel);
void *d_queued_paths = (void *)queued_paths_.device_pointer;
void *d_num_queued_paths = (void *)num_queued_paths_.device_pointer;
void *args[] = {const_cast<int *>(&work_size),
const_cast<int *>(&num_paths_limit),
&d_queued_paths,
&d_num_queued_paths,
&d_counter,
&d_prefix_sum,
&d_queued_kernel};
queue_->enqueue(kernel, work_size, args);
}
if (queued_kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE) {
queue_->zero_to_device(integrator_shader_sort_counter_);
}
else if (queued_kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE) {
queue_->zero_to_device(integrator_shader_raytrace_sort_counter_);
}
else {
assert(0);
}
}
void PathTraceWorkGPU::compute_queued_paths(DeviceKernel kernel, DeviceKernel queued_kernel)
@ -513,7 +512,7 @@ void PathTraceWorkGPU::compute_queued_paths(DeviceKernel kernel, DeviceKernel qu
int d_queued_kernel = queued_kernel;
/* Launch kernel to fill the active paths arrays. */
const int work_size = kernel_max_active_path_index(queued_kernel);
const int work_size = kernel_max_active_main_path_index(queued_kernel);
void *d_queued_paths = (void *)queued_paths_.device_pointer;
void *d_num_queued_paths = (void *)num_queued_paths_.device_pointer;
void *args[] = {
@ -523,18 +522,76 @@ void PathTraceWorkGPU::compute_queued_paths(DeviceKernel kernel, DeviceKernel qu
queue_->enqueue(kernel, work_size, args);
}
void PathTraceWorkGPU::compact_states(const int num_active_paths)
void PathTraceWorkGPU::compact_main_paths(const int num_active_paths)
{
/* Early out if there is nothing that needs to be compacted. */
if (num_active_paths == 0) {
max_active_path_index_ = 0;
}
/* Compact fragmented path states into the start of the array, moving any paths
* with index higher than the number of active paths into the gaps. */
if (max_active_path_index_ == num_active_paths) {
max_active_main_path_index_ = 0;
return;
}
const int min_compact_paths = 32;
if (max_active_main_path_index_ == num_active_paths ||
max_active_main_path_index_ < min_compact_paths) {
return;
}
/* Compact. */
compact_paths(num_active_paths,
max_active_main_path_index_,
DEVICE_KERNEL_INTEGRATOR_TERMINATED_PATHS_ARRAY,
DEVICE_KERNEL_INTEGRATOR_COMPACT_PATHS_ARRAY,
DEVICE_KERNEL_INTEGRATOR_COMPACT_STATES);
/* Adjust max active path index now we know which part of the array is actually used. */
max_active_main_path_index_ = num_active_paths;
}
void PathTraceWorkGPU::compact_shadow_paths()
{
IntegratorQueueCounter *queue_counter = integrator_queue_counter_.data();
const int num_active_paths =
queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW] +
queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW];
/* Early out if there is nothing that needs to be compacted. */
if (num_active_paths == 0) {
if (integrator_next_shadow_path_index_.data()[0] != 0) {
integrator_next_shadow_path_index_.data()[0] = 0;
queue_->copy_to_device(integrator_next_shadow_path_index_);
}
return;
}
/* Compact if we can reduce the space used by half. Not always since
* compaction has a cost. */
const float shadow_compact_ratio = 0.5f;
const int min_compact_paths = 32;
if (integrator_next_shadow_path_index_.data()[0] < num_active_paths * shadow_compact_ratio ||
integrator_next_shadow_path_index_.data()[0] < min_compact_paths) {
return;
}
/* Compact. */
compact_paths(num_active_paths,
integrator_next_shadow_path_index_.data()[0],
DEVICE_KERNEL_INTEGRATOR_TERMINATED_SHADOW_PATHS_ARRAY,
DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_PATHS_ARRAY,
DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_STATES);
/* Adjust max active path index now we know which part of the array is actually used. */
integrator_next_shadow_path_index_.data()[0] = num_active_paths;
queue_->copy_to_device(integrator_next_shadow_path_index_);
}
void PathTraceWorkGPU::compact_paths(const int num_active_paths,
const int max_active_path_index,
DeviceKernel terminated_paths_kernel,
DeviceKernel compact_paths_kernel,
DeviceKernel compact_kernel)
{
/* Compact fragmented path states into the start of the array, moving any paths
* with index higher than the number of active paths into the gaps. */
void *d_compact_paths = (void *)queued_paths_.device_pointer;
void *d_num_queued_paths = (void *)num_queued_paths_.device_pointer;
@ -545,17 +602,17 @@ void PathTraceWorkGPU::compact_states(const int num_active_paths)
int work_size = num_active_paths;
void *args[] = {&work_size, &d_compact_paths, &d_num_queued_paths, &offset};
queue_->zero_to_device(num_queued_paths_);
queue_->enqueue(DEVICE_KERNEL_INTEGRATOR_TERMINATED_PATHS_ARRAY, work_size, args);
queue_->enqueue(terminated_paths_kernel, work_size, args);
}
/* Create array of paths that we need to compact, where the path index is bigger
* than the number of active paths. */
{
int work_size = max_active_path_index_;
int work_size = max_active_path_index;
void *args[] = {
&work_size, &d_compact_paths, &d_num_queued_paths, const_cast<int *>(&num_active_paths)};
queue_->zero_to_device(num_queued_paths_);
queue_->enqueue(DEVICE_KERNEL_INTEGRATOR_COMPACT_PATHS_ARRAY, work_size, args);
queue_->enqueue(compact_paths_kernel, work_size, args);
}
queue_->copy_from_device(num_queued_paths_);
@ -570,13 +627,8 @@ void PathTraceWorkGPU::compact_states(const int num_active_paths)
int terminated_states_offset = num_active_paths;
void *args[] = {
&d_compact_paths, &active_states_offset, &terminated_states_offset, &work_size};
queue_->enqueue(DEVICE_KERNEL_INTEGRATOR_COMPACT_STATES, work_size, args);
queue_->enqueue(compact_kernel, work_size, args);
}
queue_->synchronize();
/* Adjust max active path index now we know which part of the array is actually used. */
max_active_path_index_ = num_active_paths;
}
bool PathTraceWorkGPU::enqueue_work_tiles(bool &finished)
@ -590,7 +642,7 @@ bool PathTraceWorkGPU::enqueue_work_tiles(bool &finished)
return false;
}
int num_active_paths = get_num_active_paths();
int num_active_paths = num_active_main_paths_paths();
/* Don't schedule more work if canceling. */
if (is_cancel_requested()) {
@ -630,7 +682,7 @@ bool PathTraceWorkGPU::enqueue_work_tiles(bool &finished)
/* Schedule when we're out of paths or there are too few paths to keep the
* device occupied. */
int num_paths = num_active_paths;
if (num_paths == 0 || num_paths < min_num_active_paths_) {
if (num_paths == 0 || num_paths < min_num_active_main_paths_) {
/* Get work tiles until the maximum number of path is reached. */
while (num_paths < max_num_camera_paths) {
KernelWorkTile work_tile;
@ -657,11 +709,11 @@ bool PathTraceWorkGPU::enqueue_work_tiles(bool &finished)
/* Compact state array when number of paths becomes small relative to the
* known maximum path index, which makes computing active index arrays slow. */
compact_states(num_active_paths);
compact_main_paths(num_active_paths);
if (has_shadow_catcher()) {
integrator_next_shadow_catcher_path_index_.data()[0] = num_paths;
queue_->copy_to_device(integrator_next_shadow_catcher_path_index_);
integrator_next_main_path_index_.data()[0] = num_paths;
queue_->copy_to_device(integrator_next_main_path_index_);
}
enqueue_work_tiles((device_scene_->data.bake.use) ? DEVICE_KERNEL_INTEGRATOR_INIT_FROM_BAKE :
@ -714,12 +766,11 @@ void PathTraceWorkGPU::enqueue_work_tiles(DeviceKernel kernel,
queue_->enqueue(kernel, max_tile_work_size * num_work_tiles, args);
max_active_path_index_ = path_index_offset + num_predicted_splits;
max_active_main_path_index_ = path_index_offset + num_predicted_splits;
}
int PathTraceWorkGPU::get_num_active_paths()
int PathTraceWorkGPU::num_active_main_paths_paths()
{
/* TODO: this is wrong, does not account for duplicates with shadow! */
IntegratorQueueCounter *queue_counter = integrator_queue_counter_.data();
int num_paths = 0;
@ -727,7 +778,10 @@ int PathTraceWorkGPU::get_num_active_paths()
DCHECK_GE(queue_counter->num_queued[i], 0)
<< "Invalid number of queued states for kernel "
<< device_kernel_as_string(static_cast<DeviceKernel>(i));
num_paths += queue_counter->num_queued[i];
if (!kernel_is_shadow_path((DeviceKernel)i)) {
num_paths += queue_counter->num_queued[i];
}
}
return num_paths;
@ -992,7 +1046,7 @@ bool PathTraceWorkGPU::has_shadow_catcher() const
int PathTraceWorkGPU::shadow_catcher_count_possible_splits()
{
if (max_active_path_index_ == 0) {
if (max_active_main_path_index_ == 0) {
return 0;
}
@ -1002,7 +1056,7 @@ int PathTraceWorkGPU::shadow_catcher_count_possible_splits()
queue_->zero_to_device(num_queued_paths_);
const int work_size = max_active_path_index_;
const int work_size = max_active_main_path_index_;
void *d_num_queued_paths = (void *)num_queued_paths_.device_pointer;
void *args[] = {const_cast<int *>(&work_size), &d_num_queued_paths};
@ -1026,16 +1080,23 @@ bool PathTraceWorkGPU::kernel_creates_shadow_paths(DeviceKernel kernel)
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME);
}
bool PathTraceWorkGPU::kernel_creates_ao_paths(DeviceKernel kernel)
{
return (device_scene_->data.film.pass_ao != PASS_UNUSED) &&
(kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE);
}
bool PathTraceWorkGPU::kernel_is_shadow_path(DeviceKernel kernel)
{
return (kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW);
}
int PathTraceWorkGPU::kernel_max_active_path_index(DeviceKernel kernel)
int PathTraceWorkGPU::kernel_max_active_main_path_index(DeviceKernel kernel)
{
return (kernel_is_shadow_path(kernel)) ? integrator_next_shadow_path_index_.data()[0] :
max_active_path_index_;
max_active_main_path_index_;
}
CCL_NAMESPACE_END

26
intern/cycles/integrator/path_trace_work_gpu.h
View File

@ -79,14 +79,22 @@ class PathTraceWorkGPU : public PathTraceWork {
const int num_predicted_splits);
bool enqueue_path_iteration();
void enqueue_path_iteration(DeviceKernel kernel);
void enqueue_path_iteration(DeviceKernel kernel, const int num_paths_limit = INT_MAX);
void compute_queued_paths(DeviceKernel kernel, DeviceKernel queued_kernel);
void compute_sorted_queued_paths(DeviceKernel kernel, DeviceKernel queued_kernel);
void compute_sorted_queued_paths(DeviceKernel kernel,
DeviceKernel queued_kernel,
const int num_paths_limit);
void compact_states(const int num_active_paths);
void compact_main_paths(const int num_active_paths);
void compact_shadow_paths();
void compact_paths(const int num_active_paths,
const int max_active_path_index,
DeviceKernel terminated_paths_kernel,
DeviceKernel compact_paths_kernel,
DeviceKernel compact_kernel);
int get_num_active_paths();
int num_active_main_paths_paths();
/* Check whether graphics interop can be used for the PathTraceDisplay update. */
bool should_use_graphics_interop();
@ -116,8 +124,9 @@ class PathTraceWorkGPU : public PathTraceWork {
/* Kernel properties. */
bool kernel_uses_sorting(DeviceKernel kernel);
bool kernel_creates_shadow_paths(DeviceKernel kernel);
bool kernel_creates_ao_paths(DeviceKernel kernel);
bool kernel_is_shadow_path(DeviceKernel kernel);
int kernel_max_active_path_index(DeviceKernel kernel);
int kernel_max_active_main_path_index(DeviceKernel kernel);
/* Integrator queue. */
unique_ptr<DeviceQueue> queue_;
@ -136,9 +145,10 @@ class PathTraceWorkGPU : public PathTraceWork {
/* Shader sorting. */
device_vector<int> integrator_shader_sort_counter_;
device_vector<int> integrator_shader_raytrace_sort_counter_;
device_vector<int> integrator_shader_sort_prefix_sum_;
/* Path split. */
device_vector<int> integrator_next_main_path_index_;
device_vector<int> integrator_next_shadow_path_index_;
device_vector<int> integrator_next_shadow_catcher_path_index_;
/* Temporary buffer to get an array of queued path for a particular kernel. */
device_vector<int> queued_paths_;
@ -162,12 +172,12 @@ class PathTraceWorkGPU : public PathTraceWork {
/* Minimum number of paths which keeps the device bust. If the actual number of paths falls below
* this value more work will be scheduled. */
int min_num_active_paths_;
int min_num_active_main_paths_;
/* Maximum path index, effective number of paths used may be smaller than
* the size of the integrator_state_ buffer so can avoid iterating over the
* full buffer. */
int max_active_path_index_;
int max_active_main_path_index_;
};
CCL_NAMESPACE_END

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

@ -487,9 +487,6 @@ endif()
# HIP module
if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
# 64 bit only
set(HIP_BITS 64)
# build for each arch
set(hip_sources device/hip/kernel.cpp
${SRC_HEADERS}
@ -503,33 +500,39 @@ if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
)
set(hip_fatbins)
macro(CYCLES_HIP_KERNEL_ADD arch prev_arch name flags sources experimental)
if(${arch} MATCHES "compute_.*")
set(format "ptx")
else()
set(format "fatbin")
endif()
macro(CYCLES_HIP_KERNEL_ADD arch name flags sources experimental)
set(format "fatbin")
set(hip_file ${name}_${arch}.${format})
set(kernel_sources ${sources})
if(NOT ${prev_arch} STREQUAL "none")
if(${prev_arch} MATCHES "compute_.*")
set(kernel_sources ${kernel_sources} ${name}_${prev_arch}.ptx)
else()
set(kernel_sources ${kernel_sources} ${name}_${prev_arch}.fatbin)
endif()
endif()
set(hip_kernel_src "/device/hip/${name}.cpp")
set(hip_flags ${flags}
if(WIN32)
set(hip_command ${CMAKE_COMMAND})
set(hip_flags
-E env "HIP_PATH=${HIP_ROOT_DIR}" "PATH=${HIP_PERL_DIR}"
${HIP_HIPCC_EXECUTABLE}.bat)
else()
set(hip_command ${HIP_HIPCC_EXECUTABLE})
set(hip_flags)
endif()
set(hip_flags
${hip_flags}
--amdgpu-target=${arch}
${HIP_HIPCC_FLAGS}
--genco
${CMAKE_CURRENT_SOURCE_DIR}${hip_kernel_src}
${flags}
-D CCL_NAMESPACE_BEGIN=
-D CCL_NAMESPACE_END=
-D HIPCC
-m ${HIP_BITS}
-I ${CMAKE_CURRENT_SOURCE_DIR}/..
-I ${CMAKE_CURRENT_SOURCE_DIR}/device/hip
--use_fast_math
-Wno-parentheses-equality
-Wno-unused-value
--hipcc-func-supp
-ffast-math
-o ${CMAKE_CURRENT_BINARY_DIR}/${hip_file})
if(${experimental})
@ -541,29 +544,17 @@ if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
set(hip_flags ${hip_flags} -D __KERNEL_DEBUG__)
endif()
if(WITH_NANOVDB)
set(hip_flags ${hip_flags}
-D WITH_NANOVDB
-I "${NANOVDB_INCLUDE_DIR}")
endif()
add_custom_command(
OUTPUT ${hip_file}
COMMAND ${HIP_HIPCC_EXECUTABLE}
-arch=${arch}
${HIP_HIPCC_FLAGS}
--${format}
${CMAKE_CURRENT_SOURCE_DIR}${hip_kernel_src}
${hip_flags}
COMMAND ${hip_command} ${hip_flags}
DEPENDS ${kernel_sources})
delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${hip_file}" ${CYCLES_INSTALL_PATH}/lib)
list(APPEND hip_fatbins ${hip_file})
endmacro()
set(prev_arch "none")
foreach(arch ${CYCLES_HIP_BINARIES_ARCH})
# Compile regular kernel
CYCLES_HIP_KERNEL_ADD(${arch} ${prev_arch} kernel "" "${hip_sources}" FALSE)
CYCLES_HIP_KERNEL_ADD(${arch} kernel "" "${hip_sources}" FALSE)
endforeach()
add_custom_target(cycles_kernel_hip ALL DEPENDS ${hip_fatbins})

4
intern/cycles/kernel/device/cpu/image.h
View File

@ -72,12 +72,12 @@ template<typename T> struct TextureInterpolator {
static ccl_always_inline float4 read(half4 r)
{
return half4_to_float4(r);
return half4_to_float4_image(r);
}
static ccl_always_inline float4 read(half r)
{
float f = half_to_float(r);
float f = half_to_float_image(r);
return make_float4(f, f, f, 1.0f);
}

8
intern/cycles/kernel/device/cpu/kernel_arch.h
View File

@ -107,12 +107,4 @@ void KERNEL_FUNCTION_FULL_NAME(cryptomatte_postprocess)(const KernelGlobalsCPU *
ccl_global float *render_buffer,
int pixel_index);
/* --------------------------------------------------------------------
* Bake.
*/
/* TODO(sergey): Needs to be re-implemented. Or not? Brecht did it already :) */
void KERNEL_FUNCTION_FULL_NAME(bake)(
const KernelGlobalsCPU *kg, float *buffer, int sample, int x, int y, int offset, int stride);
#undef KERNEL_ARCH

19
intern/cycles/kernel/device/cpu/kernel_arch_impl.h
View File

@ -231,25 +231,6 @@ void KERNEL_FUNCTION_FULL_NAME(cryptomatte_postprocess)(const KernelGlobalsCPU *
#endif
}
/* --------------------------------------------------------------------
* Bake.
*/
/* TODO(sergey): Needs to be re-implemented. Or not? Brecht did it already :) */
void KERNEL_FUNCTION_FULL_NAME(bake)(
const KernelGlobalsCPU *kg, float *buffer, int sample, int x, int y, int offset, int stride)
{
#if 0
# ifdef KERNEL_STUB
STUB_ASSERT(KERNEL_ARCH, bake);
# else
# ifdef __BAKING__
kernel_bake_evaluate(kg, buffer, sample, x, y, offset, stride);
# endif
# endif /* KERNEL_STUB */
#endif
}
#undef KERNEL_INVOKE
#undef DEFINE_INTEGRATOR_KERNEL
#undef DEFINE_INTEGRATOR_SHADE_KERNEL

7
intern/cycles/kernel/device/cuda/compat.h
View File

@ -128,6 +128,13 @@ __device__ half __float2half(const float f)
return val;
}
__device__ float __half2float(const half h)
{
float val;
asm("{ cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(h));
return val;
}
/* Types */
#include "util/util_half.h"

67
intern/cycles/kernel/device/gpu/kernel.h
View File

@ -281,12 +281,35 @@ extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_B
});
}
extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
kernel_gpu_integrator_terminated_shadow_paths_array(int num_states,
int *indices,
int *num_indices,
int indices_offset)
{
gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
num_states, indices + indices_offset, num_indices, [](const int state) {
return (INTEGRATOR_STATE(state, shadow_path, queued_kernel) == 0);
});
}
extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
kernel_gpu_integrator_sorted_paths_array(
int num_states, int *indices, int *num_indices, int *key_prefix_sum, int kernel)
kernel_gpu_integrator_sorted_paths_array(int num_states,
int num_states_limit,
int *indices,
int *num_indices,
int *key_counter,
int *key_prefix_sum,
int kernel)
{
gpu_parallel_sorted_index_array<GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE>(
num_states, indices, num_indices, key_prefix_sum, [kernel](const int state) {
num_states,
num_states_limit,
indices,
num_indices,
key_counter,
key_prefix_sum,
[kernel](const int state) {
return (INTEGRATOR_STATE(state, path, queued_kernel) == kernel) ?
INTEGRATOR_STATE(state, path, shader_sort_key) :
GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY;
@ -321,10 +344,40 @@ extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_SORTED_INDEX_DEFAULT_B
}
}
extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE)
kernel_gpu_prefix_sum(int *values, int num_values)
extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
kernel_gpu_integrator_compact_shadow_paths_array(int num_states,
int *indices,
int *num_indices,
int num_active_paths)
{
gpu_parallel_prefix_sum<GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE>(values, num_values);
gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>(
num_states, indices, num_indices, [num_active_paths](const int state) {
return (state >= num_active_paths) &&
(INTEGRATOR_STATE(state, shadow_path, queued_kernel) != 0);
});
}
extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE)
kernel_gpu_integrator_compact_shadow_states(const int *active_terminated_states,
const int active_states_offset,
const int terminated_states_offset,
const int work_size)
{
const int global_index = ccl_gpu_global_id_x();
if (global_index < work_size) {
const int from_state = active_terminated_states[active_states_offset + global_index];
const int to_state = active_terminated_states[terminated_states_offset + global_index];
integrator_shadow_state_move(NULL, to_state, from_state);
}
}
extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE)
kernel_gpu_prefix_sum(int *counter, int *prefix_sum, int num_values)
{
gpu_parallel_prefix_sum<GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE>(
counter, prefix_sum, num_values);
}
/* --------------------------------------------------------------------
@ -463,7 +516,7 @@ ccl_device_inline void kernel_gpu_film_convert_half_rgba_common_rgba(
film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel);
ccl_global half4 *out = ((ccl_global half4 *)rgba) + rgba_offset + y * rgba_stride + x;
float4_store_half((ccl_global half *)out, make_float4(pixel[0], pixel[1], pixel[2], pixel[3]));
*out = float4_to_half4_display(make_float4(pixel[0], pixel[1], pixel[2], pixel[3]));
}
/* Common implementation for half4 destination and 3-channel input pass. */

8
intern/cycles/kernel/device/gpu/parallel_prefix_sum.h
View File

@ -33,7 +33,8 @@ CCL_NAMESPACE_BEGIN
# define GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE 512
#endif
template<uint blocksize> __device__ void gpu_parallel_prefix_sum(int *values, const int num_values)
template<uint blocksize>
__device__ void gpu_parallel_prefix_sum(int *counter, int *prefix_sum, const int num_values)
{
if (!(ccl_gpu_block_idx_x == 0 && ccl_gpu_thread_idx_x == 0)) {
return;
@ -41,8 +42,9 @@ template<uint blocksize> __device__ void gpu_parallel_prefix_sum(int *values, co
int offset = 0;
for (int i = 0; i < num_values; i++) {
const int new_offset = offset + values[i];
values[i] = offset;
const int new_offset = offset + counter[i];
prefix_sum[i] = offset;
counter[i] = 0;
offset = new_offset;
}
}

12
intern/cycles/kernel/device/gpu/parallel_sorted_index.h
View File

@ -35,8 +35,10 @@ CCL_NAMESPACE_BEGIN
template<uint blocksize, typename GetKeyOp>
__device__ void gpu_parallel_sorted_index_array(const uint num_states,
const int num_states_limit,
int *indices,
int *num_indices,
int *key_counter,
int *key_prefix_sum,
GetKeyOp get_key_op)
{
@ -46,7 +48,15 @@ __device__ void gpu_parallel_sorted_index_array(const uint num_states,
if (key != GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY) {
const uint index = atomic_fetch_and_add_uint32(&key_prefix_sum[key], 1);
indices[index] = state_index;
if (index < num_states_limit) {
/* Assign state index. */
indices[index] = state_index;
}
else {
/* Can't process this state now, increase the counter again so that
* it will be handled in another iteration. */
atomic_fetch_and_add_uint32(&key_counter[key], 1);
}
}
}

4
intern/cycles/kernel/device/hip/globals.h
View File

@ -27,10 +27,10 @@ CCL_NAMESPACE_BEGIN
/* Not actually used, just a NULL pointer that gets passed everywhere, which we
* hope gets optimized out by the compiler. */
struct KernelGlobals {
/* NOTE: Keep the size in sync with SHADOW_STACK_MAX_HITS. */
struct KernelGlobalsGPU {
int unused[1];
};
typedef ccl_global const KernelGlobalsGPU *ccl_restrict KernelGlobals;
/* Global scene data and textures */
__constant__ KernelData __data;

7
intern/cycles/kernel/device/optix/compat.h
View File

@ -120,6 +120,13 @@ __device__ half __float2half(const float f)
return val;
}
__device__ float __half2float(const half h)
{
float val;
asm("{ cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(h));
return val;
}
/* Types */
#include "util/util_half.h"

2
intern/cycles/kernel/integrator/integrator_init_from_bake.h
View File

@ -185,7 +185,7 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg,
/* Setup next kernel to execute. */
const int shader_index = shader & SHADER_MASK;
const int shader_flags = kernel_tex_fetch(__shaders, shader_index).flags;
if ((shader_flags & SD_HAS_RAYTRACE) || (kernel_data.film.pass_ao != PASS_UNUSED)) {
if (shader_flags & SD_HAS_RAYTRACE) {
INTEGRATOR_PATH_INIT_SORTED(DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE, shader_index);
}
else {

3
intern/cycles/kernel/integrator/integrator_intersect_closest.h
View File

@ -111,8 +111,7 @@ ccl_device_forceinline void integrator_intersect_shader_next_kernel(
* Note that the splitting leaves kernel and sorting counters as-is, so use INIT semantic for
* the matte path. */
const bool use_raytrace_kernel = ((shader_flags & SD_HAS_RAYTRACE) ||
(kernel_data.film.pass_ao != PASS_UNUSED));
const bool use_raytrace_kernel = (shader_flags & SD_HAS_RAYTRACE);
if (use_raytrace_kernel) {
INTEGRATOR_PATH_NEXT_SORTED(

40
intern/cycles/kernel/integrator/integrator_megakernel.h
View File

@ -34,16 +34,12 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
ccl_global float *ccl_restrict render_buffer)
{
/* Each kernel indicates the next kernel to execute, so here we simply
* have to check what that kernel is and execute it.
*
* TODO: investigate if we can use device side enqueue for GPUs to avoid
* having to compile this big kernel. */
* have to check what that kernel is and execute it. */
while (true) {
/* Handle any shadow paths before we potentially create more shadow paths. */
const uint32_t shadow_queued_kernel = INTEGRATOR_STATE(
&state->shadow, shadow_path, queued_kernel);
if (shadow_queued_kernel) {
/* First handle any shadow paths before we potentially create more shadow paths. */
switch (shadow_queued_kernel) {
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
integrator_intersect_shadow(kg, &state->shadow);
@ -55,10 +51,30 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
kernel_assert(0);
break;
}
continue;
}
else if (INTEGRATOR_STATE(state, path, queued_kernel)) {
/* Then handle regular path kernels. */
switch (INTEGRATOR_STATE(state, path, queued_kernel)) {
/* Handle any AO paths before we potentially create more AO paths. */
const uint32_t ao_queued_kernel = INTEGRATOR_STATE(&state->ao, shadow_path, queued_kernel);
if (ao_queued_kernel) {
switch (ao_queued_kernel) {
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
integrator_intersect_shadow(kg, &state->ao);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW:
integrator_shade_shadow(kg, &state->ao, render_buffer);
break;
default:
kernel_assert(0);
break;
}
continue;
}
/* Then handle regular path kernels. */
const uint32_t queued_kernel = INTEGRATOR_STATE(state, path, queued_kernel);
if (queued_kernel) {
switch (queued_kernel) {
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST:
integrator_intersect_closest(kg, state);
break;
@ -87,10 +103,10 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
kernel_assert(0);
break;
}
continue;
}
else {
break;
}
break;
}
}

2
intern/cycles/kernel/integrator/integrator_shade_background.h
View File

@ -198,7 +198,7 @@ ccl_device void integrator_shade_background(KernelGlobals kg,
const int shader = intersection_get_shader_from_isect_prim(kg, isect_prim, isect_type);
const int shader_flags = kernel_tex_fetch(__shaders, shader).flags;
if ((shader_flags & SD_HAS_RAYTRACE) || (kernel_data.film.pass_ao != PASS_UNUSED)) {
if (shader_flags & SD_HAS_RAYTRACE) {
INTEGRATOR_PATH_NEXT_SORTED(DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND,
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE,
shader);

1
intern/cycles/kernel/integrator/integrator_shade_shadow.h
View File

@ -136,6 +136,7 @@ ccl_device_inline bool integrate_transparent_shadow(KernelGlobals kg,
INTEGRATOR_STATE_WRITE(state, shadow_path, throughput) = throughput;
INTEGRATOR_STATE_WRITE(state, shadow_path, transparent_bounce) += 1;
INTEGRATOR_STATE_WRITE(state, shadow_path, rng_offset) += PRNG_BOUNCE_NUM;
}
/* Note we do not need to check max_transparent_bounce here, the number

91
intern/cycles/kernel/integrator/integrator_shade_surface.h
View File

@ -168,7 +168,8 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
const bool is_light = light_sample_is_light(&ls);
/* Branch off shadow kernel. */
INTEGRATOR_SHADOW_PATH_INIT(shadow_state, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW);
INTEGRATOR_SHADOW_PATH_INIT(
shadow_state, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, shadow);
/* Copy volume stack and enter/exit volume. */
integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
@ -199,9 +200,8 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
state, path, render_pixel_index);
INTEGRATOR_STATE_WRITE(
shadow_state, shadow_path, rng_offset) = INTEGRATOR_STATE(state, path, rng_offset) -
PRNG_BOUNCE_NUM * transparent_bounce;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_offset) = INTEGRATOR_STATE(
state, path, rng_offset);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_hash) = INTEGRATOR_STATE(
state, path, rng_hash);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, sample) = INTEGRATOR_STATE(
@ -324,26 +324,14 @@ ccl_device_forceinline bool integrate_surface_volume_only_bounce(IntegratorState
}
#endif
#if defined(__AO__) && defined(__SHADER_RAYTRACE__)
#if defined(__AO__)
ccl_device_forceinline void integrate_surface_ao_pass(
KernelGlobals kg,
ConstIntegratorState state,
IntegratorState state,
ccl_private const ShaderData *ccl_restrict sd,
ccl_private const RNGState *ccl_restrict rng_state,
ccl_global float *ccl_restrict render_buffer)
{
# ifdef __KERNEL_OPTIX__
optixDirectCall<void>(2, kg, state, sd, rng_state, render_buffer);
}
extern "C" __device__ void __direct_callable__ao_pass(
KernelGlobals kg,
ConstIntegratorState state,
ccl_private const ShaderData *ccl_restrict sd,
ccl_private const RNGState *ccl_restrict rng_state,
ccl_global float *ccl_restrict render_buffer)
{
# endif /* __KERNEL_OPTIX__ */
float bsdf_u, bsdf_v;
path_state_rng_2D(kg, rng_state, PRNG_BSDF_U, &bsdf_u, &bsdf_v);
@ -352,24 +340,47 @@ extern "C" __device__ void __direct_callable__ao_pass(
float ao_pdf;
sample_cos_hemisphere(ao_N, bsdf_u, bsdf_v, &ao_D, &ao_pdf);
if (dot(sd->Ng, ao_D) > 0.0f && ao_pdf != 0.0f) {
Ray ray ccl_optional_struct_init;
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.dP = differential_zero_compact();
ray.dD = differential_zero_compact();
Intersection isect ccl_optional_struct_init;
if (!scene_intersect(kg, &ray, PATH_RAY_SHADOW_OPAQUE, &isect)) {
ccl_global float *buffer = kernel_pass_pixel_render_buffer(kg, state, render_buffer);
const float3 throughput = INTEGRATOR_STATE(state, path, throughput);
kernel_write_pass_float3(buffer + kernel_data.film.pass_ao, throughput);
}
if (!(dot(sd->Ng, ao_D) > 0.0f && ao_pdf != 0.0f)) {
return;
}
Ray ray ccl_optional_struct_init;
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.dP = differential_zero_compact();
ray.dD = differential_zero_compact();
/* Branch off shadow kernel. */
INTEGRATOR_SHADOW_PATH_INIT(shadow_state, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, ao);
/* Copy volume stack and enter/exit volume. */
integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(kg, shadow_state, &ray);
/* 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) | PATH_RAY_SHADOW_FOR_AO;
const float3 throughput = INTEGRATOR_STATE(state, path, throughput) * shader_bsdf_alpha(kg, sd);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
state, path, render_pixel_index);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_offset) = INTEGRATOR_STATE(
state, path, rng_offset);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_hash) = INTEGRATOR_STATE(
state, path, rng_hash);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, sample) = INTEGRATOR_STATE(
state, path, sample);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, flag) = shadow_flag;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, bounce) = bounce;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transparent_bounce) = transparent_bounce;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, throughput) = throughput;
}
#endif /* defined(__AO__) && defined(__SHADER_RAYTRACE__) */
#endif /* defined(__AO__) */
template<uint node_feature_mask>
ccl_device bool integrate_surface(KernelGlobals kg,
@ -474,14 +485,12 @@ ccl_device bool integrate_surface(KernelGlobals kg,
PROFILING_EVENT(PROFILING_SHADE_SURFACE_DIRECT_LIGHT);
integrate_surface_direct_light(kg, state, &sd, &rng_state);
#if defined(__AO__) && defined(__SHADER_RAYTRACE__)
#if defined(__AO__)
/* Ambient occlusion pass. */
if (node_feature_mask & KERNEL_FEATURE_NODE_RAYTRACE) {
if ((kernel_data.film.pass_ao != PASS_UNUSED) &&
(INTEGRATOR_STATE(state, path, flag) & PATH_RAY_CAMERA)) {
PROFILING_EVENT(PROFILING_SHADE_SURFACE_AO);
integrate_surface_ao_pass(kg, state, &sd, &rng_state, render_buffer);
}
if ((kernel_data.film.pass_ao != PASS_UNUSED) &&
(INTEGRATOR_STATE(state, path, flag) & PATH_RAY_CAMERA)) {
PROFILING_EVENT(PROFILING_SHADE_SURFACE_AO);
integrate_surface_ao_pass(kg, state, &sd, &rng_state, render_buffer);
}
#endif

8
intern/cycles/kernel/integrator/integrator_shade_volume.h
View File

@ -776,7 +776,8 @@ ccl_device_forceinline void integrate_volume_direct_light(
const bool is_light = light_sample_is_light(ls);
/* Branch off shadow kernel. */
INTEGRATOR_SHADOW_PATH_INIT(shadow_state, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW);
INTEGRATOR_SHADOW_PATH_INIT(
shadow_state, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, shadow);
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(kg, shadow_state, &ray);
@ -798,9 +799,8 @@ ccl_device_forceinline void integrate_volume_direct_light(
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
state, path, render_pixel_index);
INTEGRATOR_STATE_WRITE(
shadow_state, shadow_path, rng_offset) = INTEGRATOR_STATE(state, path, rng_offset) -
PRNG_BOUNCE_NUM * transparent_bounce;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_offset) = INTEGRATOR_STATE(
state, path, rng_offset);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_hash) = INTEGRATOR_STATE(
state, path, rng_hash);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, sample) = INTEGRATOR_STATE(

3
intern/cycles/kernel/integrator/integrator_state.h
View File

@ -92,6 +92,7 @@ typedef struct IntegratorStateCPU {
#undef KERNEL_STRUCT_VOLUME_STACK_SIZE
IntegratorShadowStateCPU shadow;
IntegratorShadowStateCPU ao;
} IntegratorStateCPU;
/* Path Queue
@ -138,7 +139,7 @@ typedef struct IntegratorStateGPU {
ccl_global int *next_shadow_path_index;
/* Index of main path which will be used by a next shadow catcher split. */
ccl_global int *next_shadow_catcher_path_index;
ccl_global int *next_main_path_index;
} IntegratorStateGPU;
/* Abstraction

6
intern/cycles/kernel/integrator/integrator_state_flow.h
View File

@ -63,7 +63,7 @@ CCL_NAMESPACE_BEGIN
&kernel_integrator_state.queue_counter->num_queued[current_kernel], 1); \
INTEGRATOR_STATE_WRITE(state, path, queued_kernel) = 0;
# define INTEGRATOR_SHADOW_PATH_INIT(shadow_state, state, next_kernel) \
# define INTEGRATOR_SHADOW_PATH_INIT(shadow_state, state, next_kernel, shadow_type) \
IntegratorShadowState shadow_state = atomic_fetch_and_add_uint32( \
&kernel_integrator_state.next_shadow_path_index[0], 1); \
atomic_fetch_and_add_uint32(&kernel_integrator_state.queue_counter->num_queued[next_kernel], \
@ -129,8 +129,8 @@ CCL_NAMESPACE_BEGIN
(void)current_kernel; \
}
# define INTEGRATOR_SHADOW_PATH_INIT(shadow_state, state, next_kernel) \
IntegratorShadowState shadow_state = &state->shadow; \
# define INTEGRATOR_SHADOW_PATH_INIT(shadow_state, state, next_kernel, shadow_type) \
IntegratorShadowState shadow_state = &state->shadow_type; \
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, queued_kernel) = next_kernel;
# define INTEGRATOR_SHADOW_PATH_NEXT(current_kernel, next_kernel) \
{ \

58
intern/cycles/kernel/integrator/integrator_state_util.h
View File

@ -265,6 +265,62 @@ ccl_device_inline void integrator_state_move(KernelGlobals kg,
INTEGRATOR_STATE_WRITE(state, path, queued_kernel) = 0;
}
ccl_device_inline void integrator_shadow_state_copy_only(KernelGlobals kg,
ConstIntegratorShadowState to_state,
ConstIntegratorShadowState state)
{
int index;
/* Rely on the compiler to optimize out unused assignments and `while(false)`'s. */
# define KERNEL_STRUCT_BEGIN(name) \
index = 0; \
do {
# define KERNEL_STRUCT_MEMBER(parent_struct, type, name, feature) \
if (kernel_integrator_state.parent_struct.name != nullptr) { \
kernel_integrator_state.parent_struct.name[to_state] = \
kernel_integrator_state.parent_struct.name[state]; \
}
# define KERNEL_STRUCT_ARRAY_MEMBER(parent_struct, type, name, feature) \
if (kernel_integrator_state.parent_struct[index].name != nullptr) { \
kernel_integrator_state.parent_struct[index].name[to_state] = \
kernel_integrator_state.parent_struct[index].name[state]; \
}
# define KERNEL_STRUCT_END(name) \
} \
while (false) \
;
# define KERNEL_STRUCT_END_ARRAY(name, cpu_array_size, gpu_array_size) \
++index; \
} \
while (index < gpu_array_size) \
;
# define KERNEL_STRUCT_VOLUME_STACK_SIZE kernel_data.volume_stack_size
# include "kernel/integrator/integrator_shadow_state_template.h"
# undef KERNEL_STRUCT_BEGIN
# undef KERNEL_STRUCT_MEMBER
# undef KERNEL_STRUCT_ARRAY_MEMBER
# undef KERNEL_STRUCT_END
# undef KERNEL_STRUCT_END_ARRAY
# undef KERNEL_STRUCT_VOLUME_STACK_SIZE
}
ccl_device_inline void integrator_shadow_state_move(KernelGlobals kg,
ConstIntegratorState to_state,
ConstIntegratorState state)
{
integrator_shadow_state_copy_only(kg, to_state, state);
INTEGRATOR_STATE_WRITE(state, shadow_path, queued_kernel) = 0;
}
#endif
/* NOTE: Leaves kernel scheduling information untouched. Use INIT semantic for one of the paths
@ -274,7 +330,7 @@ ccl_device_inline void integrator_state_shadow_catcher_split(KernelGlobals kg,
{
#if defined(__KERNEL_GPU__)
ConstIntegratorState to_state = atomic_fetch_and_add_uint32(
&kernel_integrator_state.next_shadow_catcher_path_index[0], 1);
&kernel_integrator_state.next_main_path_index[0], 1);
integrator_state_copy_only(kg, to_state, state);
#else

2
intern/cycles/kernel/integrator/integrator_subsurface.h
View File

@ -182,7 +182,7 @@ ccl_device_inline bool subsurface_scatter(KernelGlobals kg, IntegratorState stat
const int shader = intersection_get_shader(kg, &ss_isect.hits[0]);
const int shader_flags = kernel_tex_fetch(__shaders, shader).flags;
if ((shader_flags & SD_HAS_RAYTRACE) || (kernel_data.film.pass_ao != PASS_UNUSED)) {
if (shader_flags & SD_HAS_RAYTRACE) {
INTEGRATOR_PATH_NEXT_SORTED(DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE,
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE,
shader);

7
intern/cycles/kernel/kernel_accumulate.h
View File

@ -408,6 +408,13 @@ ccl_device_inline void kernel_accum_light(KernelGlobals kg,
const uint32_t path_flag = INTEGRATOR_STATE(state, shadow_path, flag);
const int sample = INTEGRATOR_STATE(state, shadow_path, sample);
/* Ambient occlusion. */
if (path_flag & PATH_RAY_SHADOW_FOR_AO) {
kernel_write_pass_float3(buffer + kernel_data.film.pass_ao, contribution);
return;
}
/* Direct light shadow. */
kernel_accum_combined_pass(kg, path_flag, sample, contribution, buffer);
#ifdef __PASSES__

6
intern/cycles/kernel/kernel_path_state.h
View File

@ -28,6 +28,7 @@ ccl_device_inline void path_state_init_queues(IntegratorState state)
INTEGRATOR_STATE_WRITE(state, path, queued_kernel) = 0;
#ifdef __KERNEL_CPU__
INTEGRATOR_STATE_WRITE(&state->shadow, shadow_path, queued_kernel) = 0;
INTEGRATOR_STATE_WRITE(&state->ao, shadow_path, queued_kernel) = 0;
#endif
}
@ -298,11 +299,8 @@ ccl_device_inline void path_state_rng_load(ConstIntegratorState state,
ccl_device_inline void shadow_path_state_rng_load(ConstIntegratorShadowState state,
ccl_private RNGState *rng_state)
{
const uint shadow_bounces = INTEGRATOR_STATE(state, shadow_path, transparent_bounce);
rng_state->rng_hash = INTEGRATOR_STATE(state, shadow_path, rng_hash);
rng_state->rng_offset = INTEGRATOR_STATE(state, shadow_path, rng_offset) +
PRNG_BOUNCE_NUM * shadow_bounces;
rng_state->rng_offset = INTEGRATOR_STATE(state, shadow_path, rng_offset);
rng_state->sample = INTEGRATOR_STATE(state, shadow_path, sample);
}

12
intern/cycles/kernel/kernel_types.h
View File

@ -285,21 +285,22 @@ enum PathRayFlag {
PATH_RAY_VOLUME_PASS = (1U << 26U),
PATH_RAY_ANY_PASS = (PATH_RAY_REFLECT_PASS | PATH_RAY_TRANSMISSION_PASS | PATH_RAY_VOLUME_PASS),
/* Shadow ray is for a light or surface. */
/* Shadow ray is for a light or surface, or AO. */
PATH_RAY_SHADOW_FOR_LIGHT = (1U << 27U),
PATH_RAY_SHADOW_FOR_AO = (1U << 28U),
/* A shadow catcher object was hit and the path was split into two. */
PATH_RAY_SHADOW_CATCHER_HIT = (1U << 28U),
PATH_RAY_SHADOW_CATCHER_HIT = (1U << 29U),
/* A shadow catcher object was hit and this path traces only shadow catchers, writing them into
* their dedicated pass for later division.
*
* NOTE: Is not covered with `PATH_RAY_ANY_PASS` because shadow catcher does special handling
* which is separate from the light passes. */
PATH_RAY_SHADOW_CATCHER_PASS = (1U << 29U),
PATH_RAY_SHADOW_CATCHER_PASS = (1U << 30U),
/* Path is evaluating background for an approximate shadow catcher with non-transparent film. */
PATH_RAY_SHADOW_CATCHER_BACKGROUND = (1U << 30U),
PATH_RAY_SHADOW_CATCHER_BACKGROUND = (1U << 31U),
};
/* Configure ray visibility bits for rays and objects respectively,
@ -1457,6 +1458,9 @@ typedef enum DeviceKernel {
DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY,
DEVICE_KERNEL_INTEGRATOR_COMPACT_PATHS_ARRAY,
DEVICE_KERNEL_INTEGRATOR_COMPACT_STATES,
DEVICE_KERNEL_INTEGRATOR_TERMINATED_SHADOW_PATHS_ARRAY,
DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_PATHS_ARRAY,
DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_STATES,
DEVICE_KERNEL_INTEGRATOR_RESET,
DEVICE_KERNEL_INTEGRATOR_SHADOW_CATCHER_COUNT_POSSIBLE_SPLITS,

4
intern/cycles/render/film.cpp
View File

@ -677,10 +677,6 @@ uint Film::get_kernel_features(const Scene *scene) const
kernel_features |= KERNEL_FEATURE_SHADOW_PASS;
}
}
if (pass_type == PASS_AO) {
kernel_features |= KERNEL_FEATURE_NODE_RAYTRACE;
}
}
return kernel_features;

172
intern/cycles/util/util_half.h
View File

@ -59,99 +59,16 @@ struct half4 {
half x, y, z, w;
};
/* Conversion to/from half float for image textures
*
* Simplified float to half for fast sampling on processor without a native
* instruction, and eliminating any NaN and inf values. */
ccl_device_inline half float_to_half_image(float f)
{
#if defined(__KERNEL_CUDA__) || defined(__KERNEL_HIP__)
ccl_device_inline void float4_store_half(ccl_private half *h, float4 f)
{
h[0] = __float2half(f.x);
h[1] = __float2half(f.y);
h[2] = __float2half(f.z);
h[3] = __float2half(f.w);
}
return __float2half(f);
#else
ccl_device_inline void float4_store_half(ccl_private half *h, float4 f)
{
# ifndef __KERNEL_SSE2__
for (int i = 0; i < 4; i++) {
/* optimized float to half for pixels:
* assumes no negative, no nan, no inf, and sets denormal to 0 */
union {
uint i;
float f;
} in;
in.f = (f[i] > 0.0f) ? ((f[i] < 65504.0f) ? f[i] : 65504.0f) : 0.0f;
int x = in.i;
int absolute = x & 0x7FFFFFFF;
int Z = absolute + 0xC8000000;
int result = (absolute < 0x38800000) ? 0 : Z;
int rshift = (result >> 13);
h[i] = (rshift & 0x7FFF);
}
# else
/* same as above with SSE */
ssef x = min(max(load4f(f), 0.0f), 65504.0f);
# ifdef __KERNEL_AVX2__
ssei rpack = _mm_cvtps_ph(x, 0);
# else
ssei absolute = cast(x) & 0x7FFFFFFF;
ssei Z = absolute + 0xC8000000;
ssei result = andnot(absolute < 0x38800000, Z);
ssei rshift = (result >> 13) & 0x7FFF;
ssei rpack = _mm_packs_epi32(rshift, rshift);
# endif
_mm_storel_pi((__m64 *)h, _mm_castsi128_ps(rpack));
# endif
}
# ifndef __KERNEL_HIP__
ccl_device_inline float half_to_float(half h)
{
float f;
*((int *)&f) = ((h & 0x8000) << 16) | (((h & 0x7c00) + 0x1C000) << 13) | ((h & 0x03FF) << 13);
return f;
}
# else
ccl_device_inline float half_to_float(std::uint32_t a) noexcept
{
std::uint32_t u = ((a << 13) + 0x70000000U) & 0x8fffe000U;
std::uint32_t v = __float_as_uint(__uint_as_float(u) *
__uint_as_float(0x77800000U) /*0x1.0p+112f*/) +
0x38000000U;
u = (a & 0x7fff) != 0 ? v : u;
return __uint_as_float(u) * __uint_as_float(0x07800000U) /*0x1.0p-112f*/;
}
# endif /* __KERNEL_HIP__ */
ccl_device_inline float4 half4_to_float4(half4 h)
{
float4 f;
f.x = half_to_float(h.x);
f.y = half_to_float(h.y);
f.z = half_to_float(h.z);
f.w = half_to_float(h.w);
return f;
}
ccl_device_inline half float_to_half(float f)
{
const uint u = __float_as_uint(f);
/* Sign bit, shifted to its position. */
uint sign_bit = u & 0x80000000;
@ -170,9 +87,82 @@ ccl_device_inline half float_to_half(float f)
value_bits = (exponent_bits == 0 ? 0 : value_bits);
/* Re-insert sign bit and return. */
return (value_bits | sign_bit);
#endif
}
ccl_device_inline float half_to_float_image(half h)
{
#if defined(__KERNEL_CUDA__) || defined(__KERNEL_HIP__)
return __half2float(h);
#else
const int x = ((h & 0x8000) << 16) | (((h & 0x7c00) + 0x1C000) << 13) | ((h & 0x03FF) << 13);
return __int_as_float(x);
#endif
}
ccl_device_inline float4 half4_to_float4_image(const half4 h)
{
/* Unable to use because it gives different results half_to_float_image, can we
* modify float_to_half_image so the conversion results are identical? */
#if 0 /* defined(__KERNEL_AVX2__) */
/* CPU: AVX. */
__m128i x = _mm_castpd_si128(_mm_load_sd((const double *)&h));
return float4(_mm_cvtph_ps(x));
#endif
const float4 f = make_float4(half_to_float_image(h.x),
half_to_float_image(h.y),
half_to_float_image(h.z),
half_to_float_image(h.w));
return f;
}
/* Conversion to half float texture for display.
*
* Simplified float to half for fast display texture conversion on processors
* without a native instruction. Assumes no negative, no NaN, no inf, and sets
* denormal to 0. */
ccl_device_inline half float_to_half_display(const float f)
{
#if defined(__KERNEL_CUDA__) || defined(__KERNEL_HIP__)
return __float2half(f);
#else
const int x = __float_as_int((f > 0.0f) ? ((f < 65504.0f) ? f : 65504.0f) : 0.0f);
const int absolute = x & 0x7FFFFFFF;
const int Z = absolute + 0xC8000000;
const int result = (absolute < 0x38800000) ? 0 : Z;
const int rshift = (result >> 13);
return (rshift & 0x7FFF);
#endif
}
ccl_device_inline half4 float4_to_half4_display(const float4 f)
{
#ifdef __KERNEL_SSE2__
/* CPU: SSE and AVX. */
ssef x = min(max(load4f(f), 0.0f), 65504.0f);
# ifdef __KERNEL_AVX2__
ssei rpack = _mm_cvtps_ph(x, 0);
# else
ssei absolute = cast(x) & 0x7FFFFFFF;
ssei Z = absolute + 0xC8000000;
ssei result = andnot(absolute < 0x38800000, Z);
ssei rshift = (result >> 13) & 0x7FFF;
ssei rpack = _mm_packs_epi32(rshift, rshift);
# endif
half4 h;
_mm_storel_pi((__m64 *)&h, _mm_castsi128_ps(rpack));
return h;
#else
/* GPU and scalar fallback. */
const half4 h = {float_to_half_display(f.x),
float_to_half_display(f.y),
float_to_half_display(f.z),
float_to_half_display(f.w)};
return h;
#endif
}
CCL_NAMESPACE_END

4
intern/cycles/util/util_image.h
View File

@ -56,7 +56,7 @@ template<> inline float util_image_cast_to_float(uint16_t value)
}
template<> inline float util_image_cast_to_float(half value)
{
return half_to_float(value);
return half_to_float_image(value);
}
/* Cast float value to output pixel type. */
@ -88,7 +88,7 @@ template<> inline uint16_t util_image_cast_from_float(float value)
}
template<> inline half util_image_cast_from_float(float value)
{
return float_to_half(value);
return float_to_half_image(value);
}
CCL_NAMESPACE_END

3
intern/ghost/intern/GHOST_SystemCocoa.mm
View File

@ -116,8 +116,6 @@ static GHOST_TKey convertKey(int rawCode, unichar recvChar, UInt16 keyAction)
case kVK_ANSI_Z: return GHOST_kKeyZ;
#endif
/* Numbers keys: mapped to handle some int'l keyboard (e.g. French). */
case kVK_ISO_Section:
return GHOST_kKeyUnknown;
case kVK_ANSI_1:
return GHOST_kKey1;
case kVK_ANSI_2:
@ -257,6 +255,7 @@ static GHOST_TKey convertKey(int rawCode, unichar recvChar, UInt16 keyAction)
case kVK_ANSI_LeftBracket: return GHOST_kKeyLeftBracket;
case kVK_ANSI_RightBracket: return GHOST_kKeyRightBracket;
case kVK_ANSI_Grave: return GHOST_kKeyAccentGrave;
case kVK_ISO_Section: return GHOST_kKeyUnknown;
#endif
case kVK_VolumeUp:
case kVK_VolumeDown:

2
intern/libmv/CMakeLists.txt
View File

@ -156,6 +156,7 @@ if(WITH_LIBMV)
libmv/base/scoped_ptr.h
libmv/base/vector.h
libmv/base/vector_utils.h
libmv/build/build_config.h
libmv/image/array_nd.h
libmv/image/convolve.h
libmv/image/correlation.h
@ -199,6 +200,7 @@ if(WITH_LIBMV)
libmv/simple_pipeline/reconstruction_scale.h
libmv/simple_pipeline/resect.h
libmv/simple_pipeline/tracks.h
libmv/threading/threading.h
libmv/tracking/brute_region_tracker.h
libmv/tracking/hybrid_region_tracker.h
libmv/tracking/kalman_filter.h

946
intern/libmv/ChangeLog
View File

File diff suppressed because it is too large Load Diff

24
intern/libmv/bundle.sh
View File

@ -9,7 +9,8 @@ fi
BRANCH="master"
repo="git://git.blender.org/libmv.git"
# repo="git://git.blender.org/libmv.git"
repo="/home/sergey/Developer/libmv"
tmp=`mktemp -d`
git clone -b $BRANCH $repo $tmp/libmv
@ -26,16 +27,16 @@ done
rm -rf $tmp
sources=`find ./libmv -type f -iname '*.cc' -or -iname '*.cpp' -or -iname '*.c' | grep -v _test.cc | grep -v test_data_sets | sed -r 's/^\.\//\t\t/' | sort -d`
headers=`find ./libmv -type f -iname '*.h' | grep -v test_data_sets | sed -r 's/^\.\//\t\t/' | sort -d`
sources=`find ./libmv -type f -iname '*.cc' -or -iname '*.cpp' -or -iname '*.c' | grep -v _test.cc | grep -v test_data_sets | sed -r 's/^\.\// /' | sort -d`
headers=`find ./libmv -type f -iname '*.h' | grep -v test_data_sets | sed -r 's/^\.\// /' | sort -d`
third_sources=`find ./third_party -type f -iname '*.cc' -or -iname '*.cpp' -or -iname '*.c' | sed -r 's/^\.\//\t\t/' | sort -d`
third_headers=`find ./third_party -type f -iname '*.h' | sed -r 's/^\.\//\t\t/' | sort -d`
third_sources=`find ./third_party -type f -iname '*.cc' -or -iname '*.cpp' -or -iname '*.c' | sed -r 's/^\.\// /' | sort -d`
third_headers=`find ./third_party -type f -iname '*.h' | sed -r 's/^\.\// /' | sort -d`
tests=`find ./libmv -type f -iname '*_test.cc' | sort -d | awk ' { name=gensub(".*/([A-Za-z_]+)_test.cc", "\\\\1", $1); printf("\t\tBLENDER_SRC_GTEST(\"libmv_%s\" \"%s\" \"libmv_test_dataset;bf_intern_libmv;extern_ceres\")\n", name, $1) } '`
tests=`find ./libmv -type f -iname '*_test.cc' | sort -d | awk ' { name=gensub(".*/([A-Za-z_]+)_test.cc", "\\\\1", "g", $1); printf(" blender_add_test_executable(\"libmv_%s\" \"%s\" \"\${INC}\" \"\${INC_SYS}\" \"libmv_test_dataset;bf_intern_libmv;extern_ceres\")\n", name, $1) } '`
src_dir=`find ./libmv -type f -iname '*.cc' -exec dirname {} \; -or -iname '*.cpp' -exec dirname {} \; -or -iname '*.c' -exec dirname {} \; | sed -r 's/^\.\//\t\t/' | sort -d | uniq`
src_third_dir=`find ./third_party -type f -iname '*.cc' -exec dirname {} \; -or -iname '*.cpp' -exec dirname {} \; -or -iname '*.c' -exec dirname {} \; | sed -r 's/^\.\//\t\t/' | sort -d | uniq`
src_dir=`find ./libmv -type f -iname '*.cc' -exec dirname {} \; -or -iname '*.cpp' -exec dirname {} \; -or -iname '*.c' -exec dirname {} \; | sed -r 's/^\.\// /' | sort -d | uniq`
src_third_dir=`find ./third_party -type f -iname '*.cc' -exec dirname {} \; -or -iname '*.cpp' -exec dirname {} \; -or -iname '*.c' -exec dirname {} \; | sed -r 's/^\.\// /' | sort -d | uniq`
src=""
win_src=""
for x in $src_dir $src_third_dir; do
@ -119,6 +120,9 @@ set(LIB
if(WITH_LIBMV)
setup_libdirs()
if(WIN32)
add_definitions(-D_USE_MATH_DEFINES)
endif()
add_definitions(\${GFLAGS_DEFINES})
add_definitions(\${GLOG_DEFINES})
add_definitions(\${CERES_DEFINES})
@ -186,7 +190,9 @@ ${third_headers}
if(WITH_GTESTS)
blender_add_lib(libmv_test_dataset "./libmv/multiview/test_data_sets.cc" "${INC}" "${INC_SYS}" "")
include(GTestTesting)
blender_add_lib(libmv_test_dataset "./libmv/multiview/test_data_sets.cc" "\${INC}" "\${INC_SYS}" "")
${tests}
endif()

2
intern/libmv/files.txt
View File

@ -119,6 +119,7 @@ libmv/simple_pipeline/resect.h
libmv/simple_pipeline/resect_test.cc
libmv/simple_pipeline/tracks.cc
libmv/simple_pipeline/tracks.h
libmv/threading/threading.h
libmv/tracking/brute_region_tracker.cc
libmv/tracking/brute_region_tracker.h
libmv/tracking/brute_region_tracker_test.cc
@ -138,6 +139,7 @@ libmv/tracking/track_region.cc
libmv/tracking/track_region.h
libmv/tracking/trklt_region_tracker.cc
libmv/tracking/trklt_region_tracker.h
third_party/.clang-format
third_party/msinttypes/inttypes.h
third_party/msinttypes/README.libmv
third_party/msinttypes/stdint.h

2
intern/libmv/mkfiles.sh
View File

@ -1,6 +1,6 @@
#!/bin/sh
find ./libmv/ -type f | sed -r 's/^\.\///' | sort > files.txt
find ./third_party/ -mindepth 2 -type f | \
find ./third_party/ -type f | \
grep -v third_party/ceres | \
sed -r 's/^\.\///' | sort >> files.txt

21
intern/libmv/third_party/msinttypes/inttypes.h vendored
View File

@ -1,35 +1,32 @@
/* No need to format 3rd-party compatibility headers. */
/* clang-format off */
// ISO C9x compliant inttypes.h for Microsoft Visual Studio
// Based on ISO/IEC 9899:TC2 Committee draft (May 6, 2005) WG14/N1124
//
// Based on ISO/IEC 9899:TC2 Committee draft (May 6, 2005) WG14/N1124
//
// Copyright (c) 2006 Alexander Chemeris
//
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
//
// 1. Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
//
// 2. Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimer in the
// documentation and/or other materials provided with the distribution.
//
//
// 3. The name of the author may be used to endorse or promote products
// derived from this software without specific prior written permission.
//
//
// THIS SOFTWARE IS PROVIDED BY THE AUTHOR ``AS IS'' AND ANY EXPRESS OR IMPLIED
// WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF
// MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO
// EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
// SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
// OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
// OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
// WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR
// OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF
// ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
//
///////////////////////////////////////////////////////////////////////////////
#ifndef _MSC_VER // [

21
intern/libmv/third_party/msinttypes/stdint.h vendored
View File

@ -1,35 +1,32 @@
/* No need to format 3rd-party compatibility headers. */
/* clang-format off */
// ISO C9x compliant stdint.h for Microsoft Visual Studio
// Based on ISO/IEC 9899:TC2 Committee draft (May 6, 2005) WG14/N1124
//
// Based on ISO/IEC 9899:TC2 Committee draft (May 6, 2005) WG14/N1124
//
// Copyright (c) 2006-2008 Alexander Chemeris
//
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
//
// 1. Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
//
// 2. Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimer in the
// documentation and/or other materials provided with the distribution.
//
//
// 3. The name of the author may be used to endorse or promote products
// derived from this software without specific prior written permission.
//
//
// THIS SOFTWARE IS PROVIDED BY THE AUTHOR ``AS IS'' AND ANY EXPRESS OR IMPLIED
// WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF
// MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO
// EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
// SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
// OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
// OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
// WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR
// OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF
// ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
//
///////////////////////////////////////////////////////////////////////////////
#ifndef _MSC_VER // [

54
release/datafiles/blender_icons.svg
View File

@ -7719,7 +7719,7 @@
<g
style="display:inline;fill:#ffffff;enable-background:new"
id="g13475"
transform="translate(-21.000002,4.4999696e-6)">
transform="translate(-20)">
<path
sodipodi:nodetypes="ccccccccccccccccccccccccccccccccccccccccccccccccccccccc"
mask="none"
@ -7733,7 +7733,7 @@
style="opacity:0.6;fill:#ffffff">
<path
id="path12187"
style="display:inline;opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;stroke:none;stroke-width:0.99999994;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;enable-background:new"
style="display:inline;opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;enable-background:new"
d="m 48,753 v 3.5 c 0,0.82843 0.67157,1.5 1.5,1.5 0.82843,0 1.5,-0.67157 1.5,-1.5 V 756 c 0,-0.5 0.53412,-1 1,-1 0.55229,0 1,0.44772 1,1 v 2.5 c 0,0.82843 0.67157,1.5 1.5,1.5 0.82843,0 1.5,-0.67157 1.5,-1.5 V 757 c 0,-0.55228 0.44772,-1 1,-1 0.55229,0 1,-0.44771 1,-1 v -2 h -2 v 0 h -2 v 0 h -2 v 0 h -2 v 0 z"
inkscape:connector-curvature="0"
sodipodi:nodetypes="csssssssssssscccccccccc" />
@ -7741,7 +7741,7 @@
</g>
<g
style="display:inline;fill:#ffffff;enable-background:new"
transform="matrix(-1,0,0,1,88.999998,-231)"
transform="matrix(-1,0,0,1,90,-231)"
id="g12197">
<path
inkscape:connector-curvature="0"
@ -8589,7 +8589,7 @@
<g
style="display:inline;fill:#ffffff;enable-background:new"
id="g12259"
transform="matrix(0,-1,-1,0,312,-85.999995)"
transform="matrix(0,-1,-1,0,313,-85.999995)"
inkscape:export-filename="C:\Users\Andrzej Ambroż\Desktop\mtrx.png"
inkscape:export-xdpi="96"
inkscape:export-ydpi="96">
@ -8843,14 +8843,14 @@
<g
style="display:inline;fill:#ffffff;enable-background:new"
id="g12761"
transform="translate(71.999998,4.4999696e-6)"
transform="translate(72.999998)"
inkscape:export-filename="C:\Users\Andrzej Ambroż\Desktop\mtrx.png"
inkscape:export-xdpi="96"
inkscape:export-ydpi="96">
<path
style="display:inline;opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;enable-background:new"
d="m 51.5,515 a 0.50005,0.50005 0 0 0 -0.353516,0.14648 l -3,3 A 0.50005,0.50005 0 0 0 48,518.5 v 10 a 0.50005,0.50005 0 0 0 0.5,0.5 h 10 a 0.50005,0.50005 0 0 0 0.353516,-0.14648 l 3,-3 A 0.50005,0.50005 0 0 0 62,525.5 v -10 A 0.50005,0.50005 0 0 0 61.5,515 Z m 0.207031,1 H 61 v 9.29297 L 58.292969,528 H 49 v -9.29297 z M 52,519 v 2 h 2 v -2 z m 2,2 v 2 h 2 v -2 z m 2,0 h 2 v -2 h -2 z m 0,2 v 2 h 2 v -2 z m 0,2 h -2 v 2 h 2 z m -2,0 v -2 h -2 v 2 z m -2,0 h -2 v 2 h 2 z m 0,-2 v -2 h -2 v 2 z"
transform="translate(-71.999998,-4.4999696e-6)"
transform="translate(-71.999998)"
id="path12740"
inkscape:connector-curvature="0" />
</g>
@ -10425,13 +10425,13 @@
sodipodi:nodetypes="sssss" />
</g>
<g
transform="translate(-1.8536743e-6,4.4999696e-6)"
transform="translate(0.97000661)"
style="display:inline;fill:#ffffff;enable-background:new"
id="g14346">
<path
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:accumulate"
d="m 184.27734,515 c -0.3909,-0.007 -0.822,0.0708 -1.2832,0.25781 -1.22985,0.49878 -2.7569,1.66874 -4.85742,3.90039 -2.09205,2.22266 -3.24657,3.73559 -3.77539,4.91602 -0.26441,0.59022 -0.37396,1.11195 -0.31641,1.58984 0.0576,0.4779 0.29454,0.88243 0.60156,1.18946 0.013,0.0131 0.0267,0.0255 0.041,0.0371 l 2.5,2 c 0.49277,0.39559 1.11361,-0.29808 0.66602,-0.74414 l -2,-2 c -0.36948,-0.36947 -0.29946,-0.99351 0,-1.29296 0.3216,-0.32161 0.86101,-0.43196 1.29296,0 l 2,2 a 0.50005,0.50005 0 1 0 0.70704,-0.70704 l -2,-2 c -0.36948,-0.36947 -0.29946,-0.99351 0,-1.29296 0.3216,-0.32161 0.86101,-0.43196 1.29296,0 l 2,2 a 0.50005,0.50005 0 1 0 0.70704,-0.70704 l -1.97657,-1.9746 c -0.008,-0.009 -0.0153,-0.0172 -0.0234,-0.0254 -0.36948,-0.36947 -0.29946,-0.99351 0,-1.29296 0.3216,-0.32161 0.86101,-0.43196 1.29296,0 l 2,2 a 0.50005,0.50005 0 1 0 0.70704,-0.70704 l -1.97657,-1.9746 c -0.008,-0.009 -0.0153,-0.0172 -0.0234,-0.0254 -0.36948,-0.36947 -0.29946,-0.99351 0,-1.29296 0.3216,-0.32161 0.86101,-0.43196 1.29296,0 l 2,2 a 0.50005,0.50005 0 1 0 0.70704,-0.70704 l -2,-2 c -0.36948,-0.36947 -0.29946,-0.99351 0,-1.29296 0.3216,-0.32161 0.86101,-0.43196 1.29296,0 l 2,2 c 0.44606,0.44759 1.13973,-0.17325 0.74414,-0.66602 l -2,-2.5 c -0.0116,-0.0143 -0.024,-0.028 -0.0371,-0.041 -0.30703,-0.30702 -0.7149,-0.52843 -1.19922,-0.61132 -0.12109,-0.0207 -0.24665,-0.0327 -0.37696,-0.0352 z"
transform="translate(1.8536743e-6,-4.4999696e-6)"
transform="translate(1.8536743e-6)"
id="path14236"
inkscape:connector-curvature="0" />
</g>
@ -11015,7 +11015,7 @@
inkscape:connector-curvature="0" />
<g
style="display:inline;fill:#ffffff;enable-background:new"
transform="translate(-189,168)"
transform="translate(-188,168)"
id="g21028">
<path
inkscape:connector-curvature="0"
@ -11084,7 +11084,7 @@
<g
style="display:inline;fill:#ffffff;enable-background:new"
id="g19653"
transform="translate(-628,42.000005)">
transform="translate(-626.99987,42.000005)">
<path
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:accumulate"
d="m 143.13867,515.0332 c -0.49613,-0.0451 -1.03862,0.15972 -1.49219,0.61328 l -1.75,1.75 c -0.19519,0.19527 -0.19519,0.51177 0,0.70704 l 3,3 c 0.19527,0.19519 0.51177,0.19519 0.70704,0 l 1.75,-1.75 c 0.45356,-0.45357 0.65838,-0.99606 0.61328,-1.49219 -0.0451,-0.49613 -0.30436,-0.90592 -0.61328,-1.21485 l -1,-1 c -0.30893,-0.30892 -0.71872,-0.56817 -1.21485,-0.61328 z m -4.39648,3.7168 a 0.50005,0.50005 0 0 0 -0.34571,0.14648 l -0.25,0.25 a 0.50005,0.50005 0 0 0 0,0.70704 l 3,3 a 0.50005,0.50005 0 0 0 0.70704,0 l 0.25,-0.25 a 0.50005,0.50005 0 0 0 0,-0.70704 l -3,-3 A 0.50005,0.50005 0 0 0 138.74219,518.75 Z m -1.25196,2.24609 a 0.50005,0.50005 0 0 0 -0.34375,0.15039 l -4.25,4.25 c -0.83838,0.83839 -1.06879,1.7573 -0.80273,2.4668 C 132.35981,528.57278 133.04167,529 133.75,529 H 138 c 1.00042,0 2,-0.79793 2,-2 v -3.5 a 0.50005,0.50005 0 1 0 -1,0 v 3.5 c 0,0.66505 -0.50442,1 -1,1 h -4.25 c -0.29167,0 -0.60981,-0.19778 -0.71875,-0.48828 -0.10894,-0.2905 -0.0893,-0.74659 0.57227,-1.4082 l 4.25,-4.25 a 0.50005,0.50005 0 0 0 -0.36329,-0.85743 z"
@ -11096,7 +11096,7 @@
inkscape:export-ydpi="96"
inkscape:export-xdpi="96"
inkscape:export-filename="C:\Users\Andrzej Ambroż\Desktop\mtrx.png"
transform="translate(418,-4.7064591e-6)"
transform="translate(418)"
id="g19647" />
</g>
<path
@ -11416,7 +11416,7 @@
<g
style="display:inline;fill:#ffffff;enable-background:new"
id="g22292"
transform="translate(-88.000002,-170)">
transform="translate(-87.000002,-170)">
<g
transform="translate(20,10)"
id="g22287"
@ -13064,7 +13064,7 @@
inkscape:export-ydpi="96"
inkscape:export-xdpi="96"
inkscape:export-filename="C:\Users\Andrzej Ambroż\Desktop\mtrx.png"
transform="translate(-21.000047,168)"
transform="translate(-20,168)"
id="g15951-6"
style="display:inline;opacity:1;fill:#ffffff;enable-background:new">
<path
@ -14462,8 +14462,8 @@
</g>
<path
id="path16385"
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:0.99999994px;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
d="m 515.5,432 c -0.1326,2e-5 -0.25976,0.0527 -0.35352,0.14647 L 513.29297,434 H 510.5 c -0.27613,3e-5 -0.49997,0.22387 -0.5,0.5 v 8 c 3e-5,0.27613 0.22387,0.49997 0.5,0.5 h 6.45312 0.004 c 0.0131,-8.2e-4 0.0261,-0.002 0.0391,-0.004 0.004,5e-5 0.008,5e-5 0.0117,0 0.0143,0.002 0.0286,0.003 0.043,0.004 h 0.006 5.44336 c 0.27613,-3e-5 0.49997,-0.22387 0.5,-0.5 v -8 c -3e-5,-0.27613 -0.22387,-0.49997 -0.5,-0.5 h -1.79297 l -1.85351,-1.85353 C 518.76,432.05268 518.6327,431.99995 518.5,432 Z m 1.49781,1.99804 c 2.21589,-3.2e-4 4.00185,1.7855 4.00195,4.00196 -9e-4,2.19736 -1.75841,3.97457 -3.95508,4 -0.0131,8.2e-4 -0.0261,0.002 -0.0391,0.004 -0.0149,-0.002 -0.0299,-0.003 -0.0449,-0.004 -10e-4,0 -0.003,0 -0.004,0 -2.19824,-0.0228 -3.95803,-1.80057 -3.95898,-4 9e-5,-2.21578 1.78479,-4.00131 4,-4.00196 z M 516.99805,435 C 515.34742,435 514,436.34922 514,438 c 0,1.65079 1.34742,3 2.99805,3 1.65063,0 3,-1.34921 3,-3 0,-1.65078 -1.34937,-3 -3,-3 z m 0,1 c 1.11009,0 2,0.88955 2,2 0,1.11045 -0.88991,2 -2,2 C 515.88796,440 515,439.11045 515,438 c 0,-1.11045 0.88796,-2 1.99805,-2 z"
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1px;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
d="m 516.5,432 c -0.1326,2e-5 -0.25976,0.0527 -0.35352,0.14647 L 514.29297,434 H 511.5 c -0.27613,3e-5 -0.49997,0.22387 -0.5,0.5 v 8 c 3e-5,0.27613 0.22387,0.49997 0.5,0.5 h 6.45312 0.004 c 0.0131,-8.2e-4 0.0261,-0.002 0.0391,-0.004 0.004,5e-5 0.008,5e-5 0.0117,0 0.0143,0.002 0.0286,0.003 0.043,0.004 h 0.006 5.44336 c 0.27613,-3e-5 0.49997,-0.22387 0.5,-0.5 v -8 c -3e-5,-0.27613 -0.22387,-0.49997 -0.5,-0.5 h -1.79297 l -1.85351,-1.85353 C 519.76,432.05268 519.6327,431.99995 519.5,432 Z m 1.49781,1.99804 c 2.21589,-3.2e-4 4.00185,1.7855 4.00195,4.00196 -9e-4,2.19736 -1.75841,3.97457 -3.95508,4 -0.0131,8.2e-4 -0.0261,0.002 -0.0391,0.004 -0.0149,-0.002 -0.0299,-0.003 -0.0449,-0.004 -10e-4,0 -0.003,0 -0.004,0 -2.19824,-0.0228 -3.95803,-1.80057 -3.95898,-4 9e-5,-2.21578 1.78479,-4.00131 4,-4.00196 z M 517.99805,435 C 516.34742,435 515,436.34922 515,438 c 0,1.65079 1.34742,3 2.99805,3 1.65063,0 3,-1.34921 3,-3 0,-1.65078 -1.34937,-3 -3,-3 z m 0,1 c 1.11009,0 2,0.88955 2,2 0,1.11045 -0.88991,2 -2,2 C 516.88796,440 516,439.11045 516,438 c 0,-1.11045 0.88796,-2 1.99805,-2 z"
inkscape:connector-curvature="0" />
<path
inkscape:connector-curvature="0"
@ -14473,18 +14473,18 @@
<path
inkscape:connector-curvature="0"
id="path22635-0-9"
d="m 494.5,432 a 0.50005,0.50005 0 0 0 -0.35352,0.14648 L 492.29297,434 H 489.5 a 0.50005,0.50005 0 0 0 -0.5,0.5 v 8 a 0.50005,0.50005 0 0 0 0.5,0.5 h 12 a 0.50005,0.50005 0 0 0 0.5,-0.5 v -8 a 0.50005,0.50005 0 0 0 -0.5,-0.5 h -1.79297 l -1.85351,-1.85352 A 0.50005,0.50005 0 0 0 497.5,432 Z m 0.20703,1 h 2.58594 l 1.85351,1.85352 A 0.50005,0.50005 0 0 0 499.5,435 h 1.5 v 7 h -11 v -7 h 2.5 a 0.50005,0.50005 0 0 0 0.35352,-0.14648 z m -0.46094,2.74414 a 0.50005,0.50005 0 0 0 -0.34961,0.85938 l 1.39649,1.39648 -1.39649,1.39648 a 0.50005,0.50005 0 1 0 0.70704,0.70704 L 496,438.70703 l 1.39648,1.39649 a 0.50005,0.50005 0 1 0 0.70704,-0.70704 L 496.70703,438 l 1.39649,-1.39648 a 0.50005,0.50005 0 1 0 -0.70704,-0.70704 L 496,437.29297 l -1.39648,-1.39649 a 0.50005,0.50005 0 0 0 -0.35743,-0.15234 z"
d="m 495.5,432 a 0.50005,0.50005 0 0 0 -0.35352,0.14648 L 493.29297,434 H 490.5 a 0.50005,0.50005 0 0 0 -0.5,0.5 v 8 a 0.50005,0.50005 0 0 0 0.5,0.5 h 12 a 0.50005,0.50005 0 0 0 0.5,-0.5 v -8 a 0.50005,0.50005 0 0 0 -0.5,-0.5 h -1.79297 l -1.85351,-1.85352 A 0.50005,0.50005 0 0 0 498.5,432 Z m 0.20703,1 h 2.58594 l 1.85351,1.85352 A 0.50005,0.50005 0 0 0 500.5,435 h 1.5 v 7 h -11 v -7 h 2.5 a 0.50005,0.50005 0 0 0 0.35352,-0.14648 z m -0.46094,2.74414 a 0.50005,0.50005 0 0 0 -0.34961,0.85938 l 1.39649,1.39648 -1.39649,1.39648 a 0.50005,0.50005 0 1 0 0.70704,0.70704 L 497,438.70703 l 1.39648,1.39649 a 0.50005,0.50005 0 1 0 0.70704,-0.70704 L 497.70703,438 l 1.39649,-1.39648 a 0.50005,0.50005 0 1 0 -0.70704,-0.70704 L 497,437.29297 l -1.39648,-1.39649 a 0.50005,0.50005 0 0 0 -0.35743,-0.15234 z"
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:0.6;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new" />
<path
sodipodi:nodetypes="ccccccccc"
inkscape:connector-curvature="0"
id="path19969-0"
d="m 469.49023,433 c -0.3497,0.006 -0.58488,0.36077 -0.45507,0.68555 l 4,10.00195 c 0.1779,0.45034 0.82806,0.4102 0.94922,-0.0586 l 1.17578,-4.46875 4.46679,-1.17578 c 0.46499,-0.12321 0.50486,-0.76769 0.0586,-0.94727 l -10,-4.00195 c -0.0621,-0.0247 -0.12852,-0.0366 -0.19532,-0.0351 z"
d="m 471.49125,433 c -0.3497,0.006 -0.58488,0.36077 -0.45507,0.68555 l 4,10.00195 c 0.1779,0.45034 0.82806,0.4102 0.94922,-0.0586 l 1.17578,-4.46875 4.46679,-1.17578 c 0.46499,-0.12321 0.50486,-0.76769 0.0586,-0.94727 l -10,-4.00195 c -0.0621,-0.0247 -0.12852,-0.0366 -0.19532,-0.0351 z"
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1px;stroke-linecap:butt;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new" />
<path
inkscape:connector-curvature="0"
id="path19971-4"
d="m 448.49023,433 a 0.50005,0.50005 0 0 0 -0.45507,0.68555 l 4,10.00195 a 0.50050292,0.50050292 0 1 0 0.92968,-0.37109 l -3.5664,-8.91797 8.91601,3.5664 a 0.50005,0.50005 0 1 0 0.3711,-0.92773 l -10,-4.00195 A 0.50005,0.50005 0 0 0 448.49023,433 Z"
d="m 450.49078,433 a 0.50005,0.50005 0 0 0 -0.45507,0.68555 l 4,10.00195 a 0.50050292,0.50050292 0 1 0 0.92968,-0.37109 l -3.5664,-8.91797 8.91601,3.5664 a 0.50005,0.50005 0 1 0 0.3711,-0.92773 l -10,-4.00195 A 0.50005,0.50005 0 0 0 450.49078,433 Z"
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:0.6;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1px;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new" />
<g
transform="translate(37.00977,-1667.9941)"
@ -15577,22 +15577,22 @@
</g>
<path
id="path22885-3-7"
d="m 537.4922,557 c -0.1299,0 -0.2539,0.055 -0.3457,0.1465 -1.4018,1.4018 -2.9571,1.8535 -5.6465,1.8535 -0.2761,0 -0.5,0.2239 -0.5,0.5 v 3 c 0,2.4627 0.6805,4.0682 1.7871,5.2754 1.1066,1.2072 2.5736,2.0242 4.1777,3.1348 0.084,0.058 0.1832,0.09 0.2852,0.09 h 0.5 c 0.102,-2e-4 0.2015,-0.032 0.2852,-0.09 1.6041,-1.1106 3.0711,-1.9276 4.1777,-3.1348 C 543.3195,566.5682 544,564.9627 544,562.5 v -3 c 0,-0.2761 -0.2239,-0.5 -0.5,-0.5 -2.6894,0 -4.2447,-0.4517 -5.6465,-1.8535 -0.096,-0.096 -0.226,-0.1486 -0.3613,-0.1465 z m 3.4766,3.9902 a 1.0001,1.0001 0 0 1 0.8124,1.6348 l -4,5 a 1.0001,1.0001 0 0 1 -1.4882,0.082 l -2,-2 a 1.0001,1.0001 0 1 1 1.414,-1.414 l 1.209,1.209 3.3028,-4.127 a 1.0001,1.0001 0 0 1 0.75,-0.3848 z"
d="m 538.4922,557 c -0.1299,0 -0.2539,0.055 -0.3457,0.1465 -1.4018,1.4018 -2.9571,1.8535 -5.6465,1.8535 -0.2761,0 -0.5,0.2239 -0.5,0.5 v 3 c 0,2.4627 0.6805,4.0682 1.7871,5.2754 1.1066,1.2072 2.5736,2.0242 4.1777,3.1348 0.084,0.058 0.1832,0.09 0.2852,0.09 h 0.5 c 0.102,-2e-4 0.2015,-0.032 0.2852,-0.09 1.6041,-1.1106 3.0711,-1.9276 4.1777,-3.1348 C 544.3195,566.5682 545,564.9627 545,562.5 v -3 c 0,-0.2761 -0.2239,-0.5 -0.5,-0.5 -2.6894,0 -4.2447,-0.4517 -5.6465,-1.8535 -0.096,-0.096 -0.226,-0.1486 -0.3613,-0.1465 z m 3.4766,3.9902 a 1.0001,1.0001 0 0 1 0.8124,1.6348 l -4,5 a 1.0001,1.0001 0 0 1 -1.4882,0.082 l -2,-2 a 1.0001,1.0001 0 1 1 1.414,-1.414 l 1.209,1.209 3.3028,-4.127 a 1.0001,1.0001 0 0 1 0.75,-0.3848 z"
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:1;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1px;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
inkscape:connector-curvature="0" />
<path
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:0.6;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
d="m 516.49219,557 a 0.50005,0.50005 0 0 0 -0.34571,0.14648 C 514.74469,558.54828 513.18939,559 510.5,559 a 0.50005,0.50005 0 0 0 -0.5,0.5 v 3 c 0,2.46272 0.6805,4.06818 1.78711,5.27539 1.10661,1.20722 2.57356,2.02419 4.17773,3.13477 A 0.50005,0.50005 0 0 0 516.25,571 h 0.5 a 0.50005,0.50005 0 0 0 0.28516,-0.0898 c 1.60417,-1.11058 3.07112,-1.92755 4.17773,-3.13477 C 522.3195,566.56818 523,564.96272 523,562.5 v -3 a 0.50005,0.50005 0 0 0 -0.5,-0.5 c -2.68939,0 -4.24469,-0.45172 -5.64648,-1.85352 A 0.50005,0.50005 0 0 0 516.49219,557 Z m 0.008,1.07031 c 1.42533,1.26825 3.16641,1.79779 5.5,1.86524 V 562.5 c 0,2.28728 -0.5695,3.55682 -1.52539,4.59961 -0.92707,1.01135 -2.30655,1.81425 -3.88867,2.90039 h -0.17188 c -1.58212,-1.08614 -2.9616,-1.88904 -3.88867,-2.90039 C 511.5695,566.05682 511,564.78728 511,562.5 v -2.56445 c 2.33359,-0.0675 4.07467,-0.59699 5.5,-1.86524 z"
d="m 517.49219,557 a 0.50005,0.50005 0 0 0 -0.34571,0.14648 C 515.74469,558.54828 514.18939,559 511.5,559 a 0.50005,0.50005 0 0 0 -0.5,0.5 v 3 c 0,2.46272 0.6805,4.06818 1.78711,5.27539 1.10661,1.20722 2.57356,2.02419 4.17773,3.13477 A 0.50005,0.50005 0 0 0 517.25,571 h 0.5 a 0.50005,0.50005 0 0 0 0.28516,-0.0898 c 1.60417,-1.11058 3.07112,-1.92755 4.17773,-3.13477 C 523.3195,566.56818 524,564.96272 524,562.5 v -3 a 0.50005,0.50005 0 0 0 -0.5,-0.5 c -2.68939,0 -4.24469,-0.45172 -5.64648,-1.85352 A 0.50005,0.50005 0 0 0 517.49219,557 Z m 0.008,1.07031 c 1.42533,1.26825 3.16641,1.79779 5.5,1.86524 V 562.5 c 0,2.28728 -0.5695,3.55682 -1.52539,4.59961 -0.92707,1.01135 -2.30655,1.81425 -3.88867,2.90039 h -0.17188 c -1.58212,-1.08614 -2.9616,-1.88904 -3.88867,-2.90039 C 512.5695,566.05682 512,564.78728 512,562.5 v -2.56445 c 2.33359,-0.0675 4.07467,-0.59699 5.5,-1.86524 z"
id="path22877-22-0"
inkscape:connector-curvature="0" />
<path
inkscape:connector-curvature="0"
d="m 511.99933,414.00018 h 9 v 5 h -9 z m -1.5,-2 c -0.27613,3e-5 -0.49997,0.22387 -0.5,0.5 v 8 c 3e-5,0.27613 0.22387,0.49997 0.5,0.5 h 5.5 v 1 h -2.5 c -0.67616,-0.01 -0.67616,1.00956 0,1 h 6 c 0.6573,-0.009 0.6573,-0.9907 0,-1 h -2.5 v -1 h 5.5 c 0.27613,-3e-5 0.49997,-0.22387 0.5,-0.5 v -8 c -3e-5,-0.27613 -0.22387,-0.49997 -0.5,-0.5 z m 0.5,1 h 11 v 7 h -11 z"
d="m 513,414.00018 h 9 v 5 h -9 z m -1.5,-2 c -0.27613,3e-5 -0.49997,0.22387 -0.5,0.5 v 8 c 3e-5,0.27613 0.22387,0.49997 0.5,0.5 h 5.5 v 1 h -2.5 c -0.67616,-0.01 -0.67616,1.00956 0,1 h 6 c 0.6573,-0.009 0.6573,-0.9907 0,-1 H 518 v -1 h 5.5 c 0.27613,-3e-5 0.49997,-0.22387 0.5,-0.5 v -8 c -3e-5,-0.27613 -0.22387,-0.49997 -0.5,-0.5 z m 0.5,1 h 11 v 7 h -11 z"
style="opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker:none;paint-order:fill markers stroke"
id="rect22324-2" />
<path
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:0.6;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:evenodd;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:round;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
d="m 489.5,412 c -0.27613,3e-5 -0.49997,0.22387 -0.5,0.5 v 8 c 3e-5,0.27613 0.22387,0.49997 0.5,0.5 h 5.5 v 1 h -2.5 c -0.67616,-0.01 -0.67616,1.00956 0,1 h 6 c 0.6573,-0.009 0.6573,-0.9907 0,-1 H 496 v -1 h 5.5 c 0.27613,-3e-5 0.49997,-0.22387 0.5,-0.5 v -8 c -3e-5,-0.27613 -0.22387,-0.49997 -0.5,-0.5 z m 0.5,1 h 11 v 7 h -11 z"
d="m 490.5,412 c -0.27613,3e-5 -0.49997,0.22387 -0.5,0.5 v 8 c 3e-5,0.27613 0.22387,0.49997 0.5,0.5 h 5.5 v 1 h -2.5 c -0.67616,-0.01 -0.67616,1.00956 0,1 h 6 c 0.6573,-0.009 0.6573,-0.9907 0,-1 H 497 v -1 h 5.5 c 0.27613,-3e-5 0.49997,-0.22387 0.5,-0.5 v -8 c -3e-5,-0.27613 -0.22387,-0.49997 -0.5,-0.5 z m 0.5,1 h 11 v 7 h -11 z"
id="path22338-8"
inkscape:connector-curvature="0"
sodipodi:nodetypes="cccccccccccccccccccccc" />
@ -17302,11 +17302,11 @@
inkscape:connector-curvature="0" />
</g>
<path
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:0.98999999;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:miter;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
d="m 418.85321,140.04954 -2.82,-2.82 a 0.62,0.62 0 0 0 -0.4,-0.18 0.6,0.6 0 0 0 -0.6,0.6 0.62,0.62 0 0 0 0.18,0.43 l 1,1 -9.18,9.12 -1,-1 a 0.62,0.62 0 0 0 -0.4,-0.15 0.6,0.6 0 0 0 -0.6,0.6 0.62,0.62 0 0 0 0.18,0.4 l 2.82,2.82 a 0.6,0.6 0 0 0 0.82,-0.82 l -1,-1 9.18,-9.15 1,1 a 0.6,0.6 0 0 0 0.82,-0.85 z"
id="path3261"
inkscape:connector-curvature="0" />
<g
style="color:#000000;font-style:normal;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:medium;line-height:normal;font-family:sans-serif;font-variant-ligatures:normal;font-variant-position:normal;font-variant-caps:normal;font-variant-numeric:normal;font-variant-alternates:normal;font-feature-settings:normal;text-indent:0;text-align:start;text-decoration:none;text-decoration-line:none;text-decoration-style:solid;text-decoration-color:#000000;letter-spacing:normal;word-spacing:normal;text-transform:none;writing-mode:lr-tb;direction:ltr;text-orientation:mixed;dominant-baseline:auto;baseline-shift:baseline;text-anchor:start;white-space:normal;shape-padding:0;clip-rule:nonzero;display:inline;overflow:visible;visibility:visible;opacity:0.98999999;isolation:auto;mix-blend-mode:normal;color-interpolation:sRGB;color-interpolation-filters:linearRGB;solid-color:#000000;solid-opacity:1;vector-effect:none;fill:#ffffff;fill-opacity:1;fill-rule:nonzero;stroke:none;stroke-width:1;stroke-linecap:round;stroke-linejoin:miter;stroke-miterlimit:4;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;color-rendering:auto;image-rendering:auto;shape-rendering:auto;text-rendering:auto;enable-background:new"
d="m 418.85321,140.04954 -2.82,-2.82 a 0.62,0.62 0 0 0 -0.4,-0.18 0.6,0.6 0 0 0 -0.6,0.6 0.62,0.62 0 0 0 0.18,0.43 l 1,1 -9.18,9.12 -1,-1 a 0.62,0.62 0 0 0 -0.4,-0.15 0.6,0.6 0 0 0 -0.6,0.6 0.62,0.62 0 0 0 0.18,0.4 l 2.82,2.82 a 0.6,0.6 0 0 0 0.82,-0.82 l -1,-1 9.18,-9.15 1,1 a 0.6,0.6 0 0 0 0.82,-0.85 z"
id="path3261"
inkscape:connector-curvature="0" />
<g
style="display:inline;enable-background:new"
transform="translate(-0.35812,42.294299)"
id="g4073">

Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 2.5 MiB

After

Width:  |  Height:  |  Size: 2.5 MiB

BIN
release/datafiles/blender_icons16/icon16_editmode_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_fake_user_off.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_fake_user_on.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_force_boid.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_object_datamode.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_particlemode.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_pose_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_restrict_render_off.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_restrict_render_on.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_restrict_select_off.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_restrict_select_on.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_restrict_view_off.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_restrict_view_on.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_sculptmode_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_tpaint_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_uv_data.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_vpaint_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons16/icon16_wpaint_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_editmode_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_fake_user_off.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_fake_user_on.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_force_boid.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_object_datamode.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_particlemode.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_pose_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_restrict_render_off.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_restrict_render_on.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_restrict_select_off.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_restrict_select_on.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_restrict_view_off.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_restrict_view_on.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_sculptmode_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_tpaint_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_uv_data.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_vpaint_hlt.dat
View File

Binary file not shown.

BIN
release/datafiles/blender_icons32/icon32_wpaint_hlt.dat
View File

Binary file not shown.

6
release/datafiles/userdef/userdef_default.c
View File

@ -35,9 +35,9 @@ const UserDef U_default = {
.subversionfile = BLENDER_FILE_SUBVERSION,
.flag = (USER_AUTOSAVE | USER_TOOLTIPS | USER_RELPATHS | USER_RELEASECONFIRM |
USER_SCRIPT_AUTOEXEC_DISABLE | USER_NONEGFRAMES),
.dupflag = USER_DUP_MESH | USER_DUP_CURVE | USER_DUP_SURF | USER_DUP_FONT | USER_DUP_MBALL |
USER_DUP_LAMP | USER_DUP_ARM | USER_DUP_ACT | USER_DUP_LIGHTPROBE |
USER_DUP_GPENCIL,
.dupflag = USER_DUP_MESH | USER_DUP_CURVE | USER_DUP_SURF | USER_DUP_LATTICE | USER_DUP_FONT |
USER_DUP_MBALL | USER_DUP_LAMP | USER_DUP_ARM | USER_DUP_CAMERA | USER_DUP_SPEAKER |
USER_DUP_ACT | USER_DUP_LIGHTPROBE | USER_DUP_GPENCIL,
.pref_flag = USER_PREF_FLAG_SAVE,
.savetime = 2,
.tempdir = "",

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