Correct numpy hash

These changes will be committed in two separate consecutive commits.
It's easier to review & build as a whole, hence the single diff.

----------------------

*Bump Python 3.9.2 → 3.9.7*

If we're going to add a Python package to Blender, this usually means
rebuilding all of Python, so we might just as well bump it to the
latest 3.9 release.

----------------------

*Bundle the `zstandard` Python package*

This package allows Python scripts to handle compressed blend files.
This is for example needed by #blender_asset_tracer to send files to a
Flamenco render farm.

This change includes a new `WITH_PYTHON_INSTALL_ZSTANDARD` build-time
option, to control whether to actually install the package. For this
the already-existing approach for Requests was copied. I'm not too
happy with the way Python packages and their dependents (Audaspace &
Mantaflow) are handled, but that's for a different cleanup commit.

Differential Revision: https://developer.blender.org/D12777

.clang-format

@@ -180,6 +180,7 @@ ForEachMacros:
   - CTX_DATA_BEGIN_WITH_ID
   - DEG_OBJECT_ITER_BEGIN
   - DEG_OBJECT_ITER_FOR_RENDER_ENGINE_BEGIN
+  - DRW_ENABLED_ENGINE_ITER
   - DRIVER_TARGETS_LOOPER_BEGIN
   - DRIVER_TARGETS_USED_LOOPER_BEGIN
   - FOREACH_BASE_IN_EDIT_MODE_BEGIN

CMakeLists.txt

@@ -388,6 +388,10 @@ 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)
@@ -406,6 +410,7 @@ 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)
@@ -1723,6 +1728,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.
@@ -1994,6 +2005,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)
 

build_files/build_environment/cmake/numpy.cmake

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

build_files/build_environment/cmake/python_site_packages.cmake

@@ -25,7 +25,7 @@ ExternalProject_Add(external_python_site_packages
   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)

build_files/build_environment/cmake/versions.cmake

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

build_files/build_environment/install_deps.sh

@@ -379,7 +379,7 @@ USE_CXX11=true
 CLANG_FORMAT_VERSION_MIN="6.0"
 CLANG_FORMAT_VERSION_MAX="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"
@@ -389,24 +389,24 @@ 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_MAX="3.2"
 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_MAX="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_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_MAX="2023.0"
 PYTHON_CERTIFI_NAME="certifi"
 
 PYTHON_REQUESTS_VERSION="2.23.0"
@@ -414,7 +414,12 @@ PYTHON_REQUESTS_VERSION_MIN="2.0"
 PYTHON_REQUESTS_VERSION_MAX="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_MAX="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_NAME="numpy"
@@ -422,20 +427,22 @@ 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_CHARSET_NORMALIZER_NAME $PYTHON_CHARSET_NORMALIZER_VERSION_MIN $PYTHON_CHARSET_NORMALIZER_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_ZSTANDARD_NAME $PYTHON_ZSTANDARD_VERSION_MIN $PYTHON_ZSTANDARD_VERSION_MAX"
   "$PYTHON_NUMPY_NAME $PYTHON_NUMPY_VERSION_MIN $PYTHON_NUMPY_VERSION_MAX"
 )
 
 # 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"
 )
 
@@ -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).
+    ** [Chardet $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).
@@ -2013,7 +2021,7 @@ compile_OIIO() {
   fi
 
   # To be changed each time we make edits that would modify the compiled result!
-  oiio_magic=17
+  oiio_magic=18
   _init_oiio
 
   # Force having own builds for the dependencies.
@@ -2088,6 +2096,7 @@ compile_OIIO() {
     cmake_d="$cmake_d -D USE_PYTHON=OFF"
     cmake_d="$cmake_d -D USE_FFMPEG=OFF"
     cmake_d="$cmake_d -D USE_OPENCV=OFF"
+    cmake_d="$cmake_d -D USE_OPENVDB=OFF"
     cmake_d="$cmake_d -D BUILD_TESTING=OFF"
     cmake_d="$cmake_d -D OIIO_BUILD_TESTS=OFF"
     cmake_d="$cmake_d -D OIIO_BUILD_TOOLS=OFF"

build_files/build_environment/patches/numpy.diff

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

build_files/cmake/Modules/FindHIP.cmake

@@ -0,0 +1,81 @@
+# - Find HIP compiler
+#
+# This module defines
+#  HIP_HIPCC_EXECUTABLE, the full path to the hipcc executable
+#  HIP_VERSION, the HIP compiler version
+#
+#  HIP_FOUND, if the HIP toolkit is found.
+
+#=============================================================================
+# Copyright 2021 Blender Foundation.
+#
+# Distributed under the OSI-approved BSD 3-Clause License,
+# see accompanying file BSD-3-Clause-license.txt for details.
+#=============================================================================
+
+# If HIP_ROOT_DIR was defined in the environment, use it.
+if(NOT HIP_ROOT_DIR AND NOT $ENV{HIP_ROOT_DIR} STREQUAL "")
+  set(HIP_ROOT_DIR $ENV{HIP_ROOT_DIR})
+endif()
+
+set(_hip_SEARCH_DIRS
+  ${HIP_ROOT_DIR}
+)
+
+find_program(HIP_HIPCC_EXECUTABLE
+  NAMES
+    hipcc
+  HINTS
+    ${_hip_SEARCH_DIRS}
+  PATH_SUFFIXES
+    bin
+)
+
+if(HIP_HIPCC_EXECUTABLE AND NOT EXISTS ${HIP_HIPCC_EXECUTABLE})
+  message(WARNING "Cached or directly specified hipcc executable does not exist.")
+  set(HIP_FOUND FALSE)
+elseif(HIP_HIPCC_EXECUTABLE)
+  set(HIP_FOUND TRUE)
+
+  set(HIP_VERSION_MAJOR 0)
+  set(HIP_VERSION_MINOR 0)
+  set(HIP_VERSION_PATCH 0)
+
+  # Get version from the output.
+  execute_process(COMMAND ${HIP_HIPCC_EXECUTABLE} --version
+                  OUTPUT_VARIABLE HIP_VERSION_RAW
+                  ERROR_QUIET
+                  OUTPUT_STRIP_TRAILING_WHITESPACE)
+
+  # Parse parts.
+  if(HIP_VERSION_RAW MATCHES "HIP version: .*")
+    # Strip the HIP prefix and get list of individual version components.
+    string(REGEX REPLACE
+           ".*HIP version: ([.0-9]+).*" "\\1"
+           HIP_SEMANTIC_VERSION "${HIP_VERSION_RAW}")
+    string(REPLACE "." ";" HIP_VERSION_PARTS "${HIP_SEMANTIC_VERSION}")
+    list(LENGTH HIP_VERSION_PARTS NUM_HIP_VERSION_PARTS)
+
+    # Extract components into corresponding variables.
+    if(NUM_HIP_VERSION_PARTS GREATER 0)
+      list(GET HIP_VERSION_PARTS 0 HIP_VERSION_MAJOR)
+    endif()
+    if(NUM_HIP_VERSION_PARTS GREATER 1)
+      list(GET HIP_VERSION_PARTS 1 HIP_VERSION_MINOR)
+    endif()
+    if(NUM_HIP_VERSION_PARTS GREATER 2)
+      list(GET HIP_VERSION_PARTS 2 HIP_VERSION_PATCH)
+    endif()
+
+    # Unset temp variables.
+    unset(NUM_HIP_VERSION_PARTS)
+    unset(HIP_SEMANTIC_VERSION)
+    unset(HIP_VERSION_PARTS)
+  endif()
+
+  # Construct full semantic version.
+  set(HIP_VERSION "${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_PATCH}")
+  unset(HIP_VERSION_RAW)
+else()
+  set(HIP_FOUND FALSE)
+endif()

build_files/cmake/cmake_static_check_cppcheck.py

@@ -24,6 +24,7 @@ import project_source_info
 import subprocess
 import sys
 import os
+import tempfile
 
 from typing import (
     Any,
@@ -35,7 +36,6 @@ USE_QUIET = (os.environ.get("QUIET", None) is not None)
 
 CHECKER_IGNORE_PREFIX = [
     "extern",
-    "intern/moto",
 ]
 
 CHECKER_BIN = "cppcheck"
@@ -47,13 +47,19 @@ CHECKER_ARGS = [
     "--max-configs=1",  # speeds up execution
     #  "--check-config", # when includes are missing
     "--enable=all",  # if you want sixty hundred pedantic suggestions
+
+    # Quiet output, otherwise all defines/includes are printed (overly verbose).
+    # Only enable this for troubleshooting (if defines are not set as expected for example).
+    "--quiet",
+
+    # NOTE: `--cppcheck-build-dir=<dir>` is added later as a temporary directory.
 ]
 
 if USE_QUIET:
     CHECKER_ARGS.append("--quiet")
 
 
-def main() -> None:
+def cppcheck() -> None:
     source_info = project_source_info.build_info(ignore_prefix_list=CHECKER_IGNORE_PREFIX)
     source_defines = project_source_info.build_defines_as_args()
 
@@ -78,7 +84,10 @@ def main() -> None:
             percent_str = "[" + ("%.2f]" % percent).rjust(7) + " %:"
 
             sys.stdout.flush()
-            sys.stdout.write("%s " % percent_str)
+            sys.stdout.write("%s %s\n" % (
+                percent_str,
+                os.path.relpath(c, project_source_info.SOURCE_DIR)
+            ))
 
         return subprocess.Popen(cmd)
 
@@ -90,5 +99,11 @@ def main() -> None:
     print("Finished!")
 
 
+def main() -> None:
+    with tempfile.TemporaryDirectory() as temp_dir:
+        CHECKER_ARGS.append("--cppcheck-build-dir=" + temp_dir)
+        cppcheck()
+
+
 if __name__ == "__main__":
     main()

build_files/cmake/project_source_info.py

@@ -243,7 +243,9 @@ def build_defines_as_args() -> List[str]:
 # use this module.
 def queue_processes(
         process_funcs: Sequence[Tuple[Callable[..., subprocess.Popen[Any]], Tuple[Any, ...]]],
+        *,
         job_total: int =-1,
+        sleep: float = 0.1,
 ) -> None:
     """ Takes a list of function arg pairs, each function must return a process
     """
@@ -271,14 +273,20 @@ def queue_processes(
 
                 if len(processes) <= job_total:
                     break
-                else:
-                    time.sleep(0.1)
+                time.sleep(sleep)
 
             sys.stdout.flush()
             sys.stderr.flush()
 
             processes.append(func(*args))
 
+        # Don't return until all jobs have finished.
+        while 1:
+            processes[:] = [p for p in processes if p.poll() is None]
+            if not processes:
+                break
+            time.sleep(sleep)
+
 
 def main() -> None:
     if not os.path.exists(join(CMAKE_DIR, "CMakeCache.txt")):

build_files/config/pipeline_config.yaml

@@ -5,7 +5,7 @@
 update-code:
     git:
         submodules:
-        -   branch: xr-controller-support
+        -   branch: master
             commit_id: HEAD
             path: release/scripts/addons
         -   branch: master

build_files/utils/make_update.py

@@ -178,7 +178,7 @@ def submodules_update(args, release_version, branch):
         branch = branch_fallback
 
     submodules = [
-        ("release/scripts/addons", "xr-controller-support", branch_fallback),
+        ("release/scripts/addons", branch, branch_fallback),
         ("release/scripts/addons_contrib", branch, branch_fallback),
         ("release/datafiles/locale", branch, branch_fallback),
         ("source/tools", branch, branch_fallback),

doc/python_api/sphinx_doc_gen.py

@@ -1101,6 +1101,7 @@ context_type_map = {
     "scene": ("Scene", False),
     "sculpt_object": ("Object", False),
     "selectable_objects": ("Object", True),
+    "selected_asset_files": ("FileSelectEntry", True),
     "selected_bones": ("EditBone", True),
     "selected_editable_bones": ("EditBone", True),
     "selected_editable_fcurves": ("FCurve", True),

extern/cuew/include/cuew.h

@@ -609,6 +609,7 @@ typedef enum cudaError_enum {
   CUDA_ERROR_INVALID_GRAPHICS_CONTEXT = 219,
   CUDA_ERROR_NVLINK_UNCORRECTABLE = 220,
   CUDA_ERROR_JIT_COMPILER_NOT_FOUND = 221,
+  CUDA_ERROR_UNSUPPORTED_PTX_VERSION = 222,
   CUDA_ERROR_INVALID_SOURCE = 300,
   CUDA_ERROR_FILE_NOT_FOUND = 301,
   CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND = 302,

extern/cuew/src/cuew.c

@@ -736,6 +736,7 @@ const char *cuewErrorString(CUresult result) {
     case CUDA_ERROR_INVALID_GRAPHICS_CONTEXT: return "Invalid graphics context";
     case CUDA_ERROR_NVLINK_UNCORRECTABLE: return "Nvlink uncorrectable";
     case CUDA_ERROR_JIT_COMPILER_NOT_FOUND: return "Jit compiler not found";
+    case CUDA_ERROR_UNSUPPORTED_PTX_VERSION: return "Unsupported PTX version";
     case CUDA_ERROR_INVALID_SOURCE: return "Invalid source";
     case CUDA_ERROR_FILE_NOT_FOUND: return "File not found";
     case CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND: return "Link to a shared object failed to resolve";

extern/smaa_areatex/smaa_areatex.cpp

@@ -574,6 +574,8 @@ Dbl2 AreaDiag::area(Dbl2 p1, Dbl2 p2, int left)
 	Dbl2 d = p2 - p1;
 	if (d.x == 0.0)
 		return Dbl2(0.0, 1.0);
+	if (d.y == 0.0)
+		return Dbl2(1.0, 0.0);
 
 	double x1 = (double)(1 + left);
 	double x2 = x1 + 1.0;

intern/cycles/blender/addon/engine.py

@@ -211,7 +211,6 @@ def list_render_passes(scene, srl):
     if crl.use_pass_shadow_catcher:        yield ("Shadow Catcher",      "RGB",  'COLOR')
 
     # Debug passes.
-    if crl.pass_debug_render_time:             yield ("Debug Render Time",             "X",   'VALUE')
     if crl.pass_debug_sample_count:            yield ("Debug Sample Count",            "X",   'VALUE')
 
     # Cryptomatte passes.

intern/cycles/blender/addon/properties.py

@@ -1197,12 +1197,6 @@ class CyclesCurveRenderSettings(bpy.types.PropertyGroup):
 
 class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
 
-    pass_debug_render_time: BoolProperty(
-        name="Debug Render Time",
-        description="Render time in milliseconds per sample and pixel",
-        default=False,
-        update=update_render_passes,
-    )
     pass_debug_sample_count: BoolProperty(
         name="Debug Sample Count",
         description="Number of samples/camera rays per pixel",

intern/cycles/blender/addon/ui.py

@@ -792,7 +792,6 @@ class CYCLES_RENDER_PT_passes_data(CyclesButtonsPanel, Panel):
         col.prop(view_layer, "use_pass_material_index")
 
         col = layout.column(heading="Debug", align=True)
-        col.prop(cycles_view_layer, "pass_debug_render_time", text="Render Time")
         col.prop(cycles_view_layer, "pass_debug_sample_count", text="Sample Count")
 
         layout.prop(view_layer, "pass_alpha_threshold")

intern/cycles/blender/blender_camera.cpp

@@ -927,6 +927,9 @@ BufferParams BlenderSync::get_buffer_params(
     params.height = height;
   }
 
+  params.window_width = params.width;
+  params.window_height = params.height;
+
   return params;
 }
 

intern/cycles/blender/blender_output_driver.cpp

@@ -72,15 +72,14 @@ bool BlenderOutputDriver::update_render_tile(const Tile &tile)
     write_render_tile(tile);
     return true;
   }
-  else {
-    /* Don't highlight full-frame tile. */
-    if (!(tile.size == tile.full_size)) {
-      b_engine_.tile_highlight_clear_all();
-      b_engine_.tile_highlight_set(tile.offset.x, tile.offset.y, tile.size.x, tile.size.y, true);
-    }
 
-    return false;
+  /* Don't highlight full-frame tile. */
+  if (!(tile.size == tile.full_size)) {
+    b_engine_.tile_highlight_clear_all();
+    b_engine_.tile_highlight_set(tile.offset.x, tile.offset.y, tile.size.x, tile.size.y, true);
   }
+
+  return false;
 }
 
 void BlenderOutputDriver::write_render_tile(const Tile &tile)

intern/cycles/blender/blender_output_driver.h

@@ -26,7 +26,7 @@ CCL_NAMESPACE_BEGIN
 
 class BlenderOutputDriver : public OutputDriver {
  public:
-  BlenderOutputDriver(BL::RenderEngine &b_engine);
+  explicit BlenderOutputDriver(BL::RenderEngine &b_engine);
   ~BlenderOutputDriver();
 
   virtual void write_render_tile(const Tile &tile) override;

intern/cycles/blender/blender_session.cpp

@@ -504,6 +504,10 @@ void BlenderSession::render_frame_finish()
   /* Clear driver. */
   session->set_output_driver(nullptr);
   session->full_buffer_written_cb = function_null;
+
+  /* All the files are handled.
+   * Clear the list so that this session can be re-used by Persistent Data. */
+  full_buffer_files_.clear();
 }
 
 static PassType bake_type_to_pass(const string &bake_type_str, const int bake_filter)

intern/cycles/blender/blender_sync.cpp

@@ -545,8 +545,6 @@ static PassType get_blender_pass_type(BL::RenderPass &b_pass)
   MAP_PASS("Shadow Catcher", PASS_SHADOW_CATCHER);
   MAP_PASS("Noisy Shadow Catcher", PASS_SHADOW_CATCHER);
 
-  MAP_PASS("Debug Render Time", PASS_RENDER_TIME);
-
   MAP_PASS("AdaptiveAuxBuffer", PASS_ADAPTIVE_AUX_BUFFER);
   MAP_PASS("Debug Sample Count", PASS_SAMPLE_COUNT);
 
@@ -604,10 +602,6 @@ void BlenderSync::sync_render_passes(BL::RenderLayer &b_rlay, BL::ViewLayer &b_v
   PointerRNA crl = RNA_pointer_get(&b_view_layer.ptr, "cycles");
 
   /* Debug passes. */
-  if (get_boolean(crl, "pass_debug_render_time")) {
-    b_engine.add_pass("Debug Render Time", 1, "X", b_view_layer.name().c_str());
-    pass_add(scene, PASS_RENDER_TIME, "Debug Render Time");
-  }
   if (get_boolean(crl, "pass_debug_sample_count")) {
     b_engine.add_pass("Debug Sample Count", 1, "X", b_view_layer.name().c_str());
     pass_add(scene, PASS_SAMPLE_COUNT, "Debug Sample Count");

intern/cycles/bvh/bvh.h

@@ -50,10 +50,6 @@ struct PackedBVH {
   array<int4> leaf_nodes;
   /* object index to BVH node index mapping for instances */
   array<int> object_node;
-  /* Mapping from primitive index to index in triangle array. */
-  array<uint> prim_tri_index;
-  /* Continuous storage of triangle vertices. */
-  array<float4> prim_tri_verts;
   /* primitive type - triangle or strand */
   array<int> prim_type;
   /* visibility visibilitys for primitives */

intern/cycles/bvh/bvh2.cpp

@@ -439,61 +439,20 @@ void BVH2::refit_primitives(int start, int end, BoundBox &bbox, uint &visibility
 
 /* Triangles */
 
-void BVH2::pack_triangle(int idx, float4 tri_verts[3])
-{
-  int tob = pack.prim_object[idx];
-  assert(tob >= 0 && tob < objects.size());
-  const Mesh *mesh = static_cast<const Mesh *>(objects[tob]->get_geometry());
-
-  int tidx = pack.prim_index[idx];
-  Mesh::Triangle t = mesh->get_triangle(tidx);
-  const float3 *vpos = &mesh->verts[0];
-  float3 v0 = vpos[t.v[0]];
-  float3 v1 = vpos[t.v[1]];
-  float3 v2 = vpos[t.v[2]];
-
-  tri_verts[0] = float3_to_float4(v0);
-  tri_verts[1] = float3_to_float4(v1);
-  tri_verts[2] = float3_to_float4(v2);
-}
-
 void BVH2::pack_primitives()
 {
   const size_t tidx_size = pack.prim_index.size();
-  size_t num_prim_triangles = 0;
-  /* Count number of triangles primitives in BVH. */
-  for (unsigned int i = 0; i < tidx_size; i++) {
-    if ((pack.prim_index[i] != -1)) {
-      if ((pack.prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
-        ++num_prim_triangles;
-      }
-    }
-  }
   /* Reserve size for arrays. */
-  pack.prim_tri_index.clear();
-  pack.prim_tri_index.resize(tidx_size);
-  pack.prim_tri_verts.clear();
-  pack.prim_tri_verts.resize(num_prim_triangles * 3);
   pack.prim_visibility.clear();
   pack.prim_visibility.resize(tidx_size);
   /* Fill in all the arrays. */
-  size_t prim_triangle_index = 0;
   for (unsigned int i = 0; i < tidx_size; i++) {
     if (pack.prim_index[i] != -1) {
       int tob = pack.prim_object[i];
       Object *ob = objects[tob];
-      if ((pack.prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
-        pack_triangle(i, (float4 *)&pack.prim_tri_verts[3 * prim_triangle_index]);
-        pack.prim_tri_index[i] = 3 * prim_triangle_index;
-        ++prim_triangle_index;
-      }
-      else {
-        pack.prim_tri_index[i] = -1;
-      }
       pack.prim_visibility[i] = ob->visibility_for_tracing();
     }
     else {
-      pack.prim_tri_index[i] = -1;
       pack.prim_visibility[i] = 0;
     }
   }
@@ -522,10 +481,8 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
 
   /* reserve */
   size_t prim_index_size = pack.prim_index.size();
-  size_t prim_tri_verts_size = pack.prim_tri_verts.size();
 
   size_t pack_prim_index_offset = prim_index_size;
-  size_t pack_prim_tri_verts_offset = prim_tri_verts_size;
   size_t pack_nodes_offset = nodes_size;
   size_t pack_leaf_nodes_offset = leaf_nodes_size;
   size_t object_offset = 0;
@@ -535,7 +492,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
 
     if (geom->need_build_bvh(params.bvh_layout)) {
       prim_index_size += bvh->pack.prim_index.size();
-      prim_tri_verts_size += bvh->pack.prim_tri_verts.size();
       nodes_size += bvh->pack.nodes.size();
       leaf_nodes_size += bvh->pack.leaf_nodes.size();
     }
@@ -545,8 +501,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
   pack.prim_type.resize(prim_index_size);
   pack.prim_object.resize(prim_index_size);
   pack.prim_visibility.resize(prim_index_size);
-  pack.prim_tri_verts.resize(prim_tri_verts_size);
-  pack.prim_tri_index.resize(prim_index_size);
   pack.nodes.resize(nodes_size);
   pack.leaf_nodes.resize(leaf_nodes_size);
   pack.object_node.resize(objects.size());
@@ -559,8 +513,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
   int *pack_prim_type = (pack.prim_type.size()) ? &pack.prim_type[0] : NULL;
   int *pack_prim_object = (pack.prim_object.size()) ? &pack.prim_object[0] : NULL;
   uint *pack_prim_visibility = (pack.prim_visibility.size()) ? &pack.prim_visibility[0] : NULL;
-  float4 *pack_prim_tri_verts = (pack.prim_tri_verts.size()) ? &pack.prim_tri_verts[0] : NULL;
-  uint *pack_prim_tri_index = (pack.prim_tri_index.size()) ? &pack.prim_tri_index[0] : NULL;
   int4 *pack_nodes = (pack.nodes.size()) ? &pack.nodes[0] : NULL;
   int4 *pack_leaf_nodes = (pack.leaf_nodes.size()) ? &pack.leaf_nodes[0] : NULL;
   float2 *pack_prim_time = (pack.prim_time.size()) ? &pack.prim_time[0] : NULL;
@@ -609,18 +561,14 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
       int *bvh_prim_index = &bvh->pack.prim_index[0];
       int *bvh_prim_type = &bvh->pack.prim_type[0];
       uint *bvh_prim_visibility = &bvh->pack.prim_visibility[0];
-      uint *bvh_prim_tri_index = &bvh->pack.prim_tri_index[0];
       float2 *bvh_prim_time = bvh->pack.prim_time.size() ? &bvh->pack.prim_time[0] : NULL;
 
       for (size_t i = 0; i < bvh_prim_index_size; i++) {
         if (bvh->pack.prim_type[i] & PRIMITIVE_ALL_CURVE) {
           pack_prim_index[pack_prim_index_offset] = bvh_prim_index[i] + geom_prim_offset;
-          pack_prim_tri_index[pack_prim_index_offset] = -1;
         }
         else {
           pack_prim_index[pack_prim_index_offset] = bvh_prim_index[i] + geom_prim_offset;
-          pack_prim_tri_index[pack_prim_index_offset] = bvh_prim_tri_index[i] +
-                                                        pack_prim_tri_verts_offset;
         }
 
         pack_prim_type[pack_prim_index_offset] = bvh_prim_type[i];
@@ -633,15 +581,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
       }
     }
 
-    /* Merge triangle vertices data. */
-    if (bvh->pack.prim_tri_verts.size()) {
-      const size_t prim_tri_size = bvh->pack.prim_tri_verts.size();
-      memcpy(pack_prim_tri_verts + pack_prim_tri_verts_offset,
-             &bvh->pack.prim_tri_verts[0],
-             prim_tri_size * sizeof(float4));
-      pack_prim_tri_verts_offset += prim_tri_size;
-    }
-
     /* merge nodes */
     if (bvh->pack.leaf_nodes.size()) {
       int4 *leaf_nodes_offset = &bvh->pack.leaf_nodes[0];

intern/cycles/bvh/bvh_build.cpp

@@ -67,8 +67,12 @@ BVHBuild::~BVHBuild()
 
 /* Adding References */
 
-void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *mesh, int i)
+void BVHBuild::add_reference_triangles(BoundBox &root,
+                                       BoundBox &center,
+                                       Mesh *mesh,
+                                       int object_index)
 {
+  const PrimitiveType primitive_type = mesh->primitive_type();
   const Attribute *attr_mP = NULL;
   if (mesh->has_motion_blur()) {
     attr_mP = mesh->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
@@ -81,7 +85,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
       BoundBox bounds = BoundBox::empty;
       t.bounds_grow(verts, bounds);
       if (bounds.valid() && t.valid(verts)) {
-        references.push_back(BVHReference(bounds, j, i, PRIMITIVE_TRIANGLE));
+        references.push_back(BVHReference(bounds, j, object_index, primitive_type));
         root.grow(bounds);
         center.grow(bounds.center2());
       }
@@ -101,7 +105,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
         t.bounds_grow(vert_steps + step * num_verts, bounds);
       }
       if (bounds.valid()) {
-        references.push_back(BVHReference(bounds, j, i, PRIMITIVE_MOTION_TRIANGLE));
+        references.push_back(BVHReference(bounds, j, object_index, primitive_type));
         root.grow(bounds);
         center.grow(bounds.center2());
       }
@@ -140,7 +144,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
         if (bounds.valid()) {
           const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1;
           references.push_back(
-              BVHReference(bounds, j, i, PRIMITIVE_MOTION_TRIANGLE, prev_time, curr_time));
+              BVHReference(bounds, j, object_index, primitive_type, prev_time, curr_time));
           root.grow(bounds);
           center.grow(bounds.center2());
         }
@@ -153,18 +157,14 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
   }
 }
 
-void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair, int i)
+void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair, int object_index)
 {
   const Attribute *curve_attr_mP = NULL;
   if (hair->has_motion_blur()) {
     curve_attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
   }
 
-  const PrimitiveType primitive_type =
-      (curve_attr_mP != NULL) ?
-          ((hair->curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON :
-                                                 PRIMITIVE_MOTION_CURVE_THICK) :
-          ((hair->curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON : PRIMITIVE_CURVE_THICK);
+  const PrimitiveType primitive_type = hair->primitive_type();
 
   const size_t num_curves = hair->num_curves();
   for (uint j = 0; j < num_curves; j++) {
@@ -177,7 +177,7 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
         curve.bounds_grow(k, &hair->get_curve_keys()[0], curve_radius, bounds);
         if (bounds.valid()) {
           int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
-          references.push_back(BVHReference(bounds, j, i, packed_type));
+          references.push_back(BVHReference(bounds, j, object_index, packed_type));
           root.grow(bounds);
           center.grow(bounds.center2());
         }
@@ -198,7 +198,7 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
         }
         if (bounds.valid()) {
           int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
-          references.push_back(BVHReference(bounds, j, i, packed_type));
+          references.push_back(BVHReference(bounds, j, object_index, packed_type));
           root.grow(bounds);
           center.grow(bounds.center2());
         }
@@ -254,7 +254,8 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
           if (bounds.valid()) {
             const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1;
             int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
-            references.push_back(BVHReference(bounds, j, i, packed_type, prev_time, curr_time));
+            references.push_back(
+                BVHReference(bounds, j, object_index, packed_type, prev_time, curr_time));
             root.grow(bounds);
             center.grow(bounds.center2());
           }
@@ -268,15 +269,18 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
   }
 }
 
-void BVHBuild::add_reference_geometry(BoundBox &root, BoundBox &center, Geometry *geom, int i)
+void BVHBuild::add_reference_geometry(BoundBox &root,
+                                      BoundBox &center,
+                                      Geometry *geom,
+                                      int object_index)
 {
   if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
     Mesh *mesh = static_cast<Mesh *>(geom);
-    add_reference_triangles(root, center, mesh, i);
+    add_reference_triangles(root, center, mesh, object_index);
   }
   else if (geom->geometry_type == Geometry::HAIR) {
     Hair *hair = static_cast<Hair *>(geom);
-    add_reference_curves(root, center, hair, i);
+    add_reference_curves(root, center, hair, object_index);
   }
 }
 

intern/cycles/bvh/bvh_embree.cpp

@@ -89,20 +89,9 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
 
       /* Test if we need to record this transparent intersection. */
       if (ctx->num_hits < ctx->max_hits || ray->tfar < ctx->max_t) {
-        /* Skip already recorded intersections. */
-        int num_recorded_hits = min(ctx->num_hits, ctx->max_hits);
-
-        for (int i = 0; i < num_recorded_hits; ++i) {
-          if (current_isect.object == ctx->isect_s[i].object &&
-              current_isect.prim == ctx->isect_s[i].prim && current_isect.t == ctx->isect_s[i].t) {
-            /* This intersection was already recorded, skip it. */
-            *args->valid = 0;
-            return;
-          }
-        }
-
         /* If maximum number of hits was reached, replace the intersection with the
          * highest distance. We want to find the N closest intersections. */
+        const int num_recorded_hits = min(ctx->num_hits, ctx->max_hits);
         int isect_index = num_recorded_hits;
         if (num_recorded_hits + 1 >= ctx->max_hits) {
           float max_t = ctx->isect_s[0].t;
@@ -147,10 +136,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
       }
       else {
         kernel_embree_convert_hit(kg, ray, hit, &current_isect);
-        int object = (current_isect.object == OBJECT_NONE) ?
-                         kernel_tex_fetch(__prim_object, current_isect.prim) :
-                         current_isect.object;
-        if (ctx->local_object_id != object) {
+        if (ctx->local_object_id != current_isect.object) {
           /* This tells Embree to continue tracing. */
           *args->valid = 0;
           break;
@@ -169,41 +155,49 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
         break;
       }
 
-      /* See triangle_intersect_subsurface() for the native equivalent. */
-      for (int i = min(ctx->max_hits, ctx->local_isect->num_hits) - 1; i >= 0; --i) {
-        if (ctx->local_isect->hits[i].t == ray->tfar) {
-          /* This tells Embree to continue tracing. */
-          *args->valid = 0;
-          break;
-        }
-      }
-
+      LocalIntersection *local_isect = ctx->local_isect;
       int hit_idx = 0;
 
       if (ctx->lcg_state) {
+        /* See triangle_intersect_subsurface() for the native equivalent. */
+        for (int i = min(ctx->max_hits, local_isect->num_hits) - 1; i >= 0; --i) {
+          if (local_isect->hits[i].t == ray->tfar) {
+            /* This tells Embree to continue tracing. */
+            *args->valid = 0;
+            return;
+          }
+        }
 
-        ++ctx->local_isect->num_hits;
-        if (ctx->local_isect->num_hits <= ctx->max_hits) {
-          hit_idx = ctx->local_isect->num_hits - 1;
+        local_isect->num_hits++;
+
+        if (local_isect->num_hits <= ctx->max_hits) {
+          hit_idx = local_isect->num_hits - 1;
         }
         else {
           /* reservoir sampling: if we are at the maximum number of
            * hits, randomly replace element or skip it */
-          hit_idx = lcg_step_uint(ctx->lcg_state) % ctx->local_isect->num_hits;
+          hit_idx = lcg_step_uint(ctx->lcg_state) % local_isect->num_hits;
 
           if (hit_idx >= ctx->max_hits) {
             /* This tells Embree to continue tracing. */
             *args->valid = 0;
-            break;
+            return;
           }
         }
       }
       else {
-        ctx->local_isect->num_hits = 1;
+        /* Record closest intersection only. */
+        if (local_isect->num_hits && current_isect.t > local_isect->hits[0].t) {
+          *args->valid = 0;
+          return;
+        }
+
+        local_isect->num_hits = 1;
       }
+
       /* record intersection */
-      ctx->local_isect->hits[hit_idx] = current_isect;
-      ctx->local_isect->Ng[hit_idx] = normalize(make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z));
+      local_isect->hits[hit_idx] = current_isect;
+      local_isect->Ng[hit_idx] = normalize(make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z));
       /* This tells Embree to continue tracing. */
       *args->valid = 0;
       break;
@@ -213,21 +207,11 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
       if (ctx->num_hits < ctx->max_hits) {
         Intersection current_isect;
         kernel_embree_convert_hit(kg, ray, hit, &current_isect);
-        for (size_t i = 0; i < ctx->num_hits; ++i) {
-          if (current_isect.object == ctx->isect_s[i].object &&
-              current_isect.prim == ctx->isect_s[i].prim && current_isect.t == ctx->isect_s[i].t) {
-            /* This intersection was already recorded, skip it. */
-            *args->valid = 0;
-            break;
-          }
-        }
         Intersection *isect = &ctx->isect_s[ctx->num_hits];
         ++ctx->num_hits;
         *isect = current_isect;
         /* Only primitives from volume object. */
-        uint tri_object = (isect->object == OBJECT_NONE) ?
-                              kernel_tex_fetch(__prim_object, isect->prim) :
-                              isect->object;
+        uint tri_object = isect->object;
         int object_flag = kernel_tex_fetch(__object_flag, tri_object);
         if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
           --ctx->num_hits;
@@ -249,7 +233,7 @@ static void rtc_filter_func_thick_curve(const RTCFilterFunctionNArguments *args)
   const RTCRay *ray = (RTCRay *)args->ray;
   RTCHit *hit = (RTCHit *)args->hit;
 
-  /* Always ignore backfacing intersections. */
+  /* Always ignore back-facing intersections. */
   if (dot(make_float3(ray->dir_x, ray->dir_y, ray->dir_z),
           make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)) > 0.0f) {
     *args->valid = 0;
@@ -262,7 +246,7 @@ static void rtc_filter_occluded_func_thick_curve(const RTCFilterFunctionNArgumen
   const RTCRay *ray = (RTCRay *)args->ray;
   RTCHit *hit = (RTCHit *)args->hit;
 
-  /* Always ignore backfacing intersections. */
+  /* Always ignore back-facing intersections. */
   if (dot(make_float3(ray->dir_x, ray->dir_y, ray->dir_z),
           make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)) > 0.0f) {
     *args->valid = 0;
@@ -456,7 +440,7 @@ void BVHEmbree::add_instance(Object *ob, int i)
 
 void BVHEmbree::add_triangles(const Object *ob, const Mesh *mesh, int i)
 {
-  size_t prim_offset = mesh->optix_prim_offset;
+  size_t prim_offset = mesh->prim_offset;
 
   const Attribute *attr_mP = NULL;
   size_t num_motion_steps = 1;
@@ -625,7 +609,7 @@ void BVHEmbree::set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, c
 
 void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
 {
-  size_t prim_offset = hair->optix_prim_offset;
+  size_t prim_offset = hair->curve_segment_offset;
 
   const Attribute *attr_mP = NULL;
   size_t num_motion_steps = 1;
@@ -702,7 +686,7 @@ void BVHEmbree::refit(Progress &progress)
         if (mesh->num_triangles() > 0) {
           RTCGeometry geom = rtcGetGeometry(scene, geom_id);
           set_tri_vertex_buffer(geom, mesh, true);
-          rtcSetGeometryUserData(geom, (void *)mesh->optix_prim_offset);
+          rtcSetGeometryUserData(geom, (void *)mesh->prim_offset);
           rtcCommitGeometry(geom);
         }
       }
@@ -711,7 +695,7 @@ void BVHEmbree::refit(Progress &progress)
         if (hair->num_curves() > 0) {
           RTCGeometry geom = rtcGetGeometry(scene, geom_id + 1);
           set_curve_vertex_buffer(geom, hair, true);
-          rtcSetGeometryUserData(geom, (void *)hair->optix_prim_offset);
+          rtcSetGeometryUserData(geom, (void *)hair->curve_segment_offset);
           rtcCommitGeometry(geom);
         }
       }

intern/cycles/cmake/external_libs.cmake

@@ -521,7 +521,7 @@ endif()
 if(WITH_CYCLES_CUDA_BINARIES OR NOT WITH_CUDA_DYNLOAD)
   find_package(CUDA) # Try to auto locate CUDA toolkit
   if(CUDA_FOUND)
-    message(STATUS "CUDA nvcc = ${CUDA_NVCC_EXECUTABLE}")
+    message(STATUS "Found CUDA ${CUDA_NVCC_EXECUTABLE} (${CUDA_VERSION})")
   else()
     message(STATUS "CUDA compiler not found, disabling WITH_CYCLES_CUDA_BINARIES")
     set(WITH_CYCLES_CUDA_BINARIES OFF)
@@ -537,6 +537,16 @@ endif()
 # HIP
 ###########################################################################
 
+if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
+  find_package(HIP)
+  if(HIP_FOUND)
+    message(STATUS "Found HIP ${HIP_HIPCC_EXECUTABLE} (${HIP_VERSION})")
+  else()
+    message(STATUS "HIP compiler not found, disabling WITH_CYCLES_HIP_BINARIES")
+    set(WITH_CYCLES_HIP_BINARIES OFF)
+  endif()
+endif()
+
 if(NOT WITH_HIP_DYNLOAD)
   set(WITH_HIP_DYNLOAD ON)
 endif()

intern/cycles/device/optix/device_impl.cpp

@@ -768,7 +768,13 @@ void OptiXDevice::denoise_color_read(DenoiseContext &context, const DenoisePass
   destination.num_components = 3;
   destination.pixel_stride = context.buffer_params.pass_stride;
 
-  pass_accessor.get_render_tile_pixels(context.render_buffers, context.buffer_params, destination);
+  BufferParams buffer_params = context.buffer_params;
+  buffer_params.window_x = 0;
+  buffer_params.window_y = 0;
+  buffer_params.window_width = buffer_params.width;
+  buffer_params.window_height = buffer_params.height;
+
+  pass_accessor.get_render_tile_pixels(context.render_buffers, buffer_params, destination);
 }
 
 bool OptiXDevice::denoise_filter_color_preprocess(DenoiseContext &context, const DenoisePass &pass)
@@ -1246,7 +1252,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         build_input.curveArray.indexBuffer = (CUdeviceptr)index_data.device_pointer;
         build_input.curveArray.indexStrideInBytes = sizeof(int);
         build_input.curveArray.flag = build_flags;
-        build_input.curveArray.primitiveIndexOffset = hair->optix_prim_offset;
+        build_input.curveArray.primitiveIndexOffset = hair->curve_segment_offset;
       }
       else {
         /* Disable visibility test any-hit program, since it is already checked during
@@ -1259,7 +1265,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         build_input.customPrimitiveArray.strideInBytes = sizeof(OptixAabb);
         build_input.customPrimitiveArray.flags = &build_flags;
         build_input.customPrimitiveArray.numSbtRecords = 1;
-        build_input.customPrimitiveArray.primitiveIndexOffset = hair->optix_prim_offset;
+        build_input.customPrimitiveArray.primitiveIndexOffset = hair->curve_segment_offset;
       }
 
       if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
@@ -1328,7 +1334,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
        * buffers for that purpose. OptiX does not allow this to be zero though, so just pass in
        * one and rely on that having the same meaning in this case. */
       build_input.triangleArray.numSbtRecords = 1;
-      build_input.triangleArray.primitiveIndexOffset = mesh->optix_prim_offset;
+      build_input.triangleArray.primitiveIndexOffset = mesh->prim_offset;
 
       if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
         progress.set_error("Failed to build OptiX acceleration structure");
@@ -1395,8 +1401,8 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
       instance.transform[5] = 1.0f;
       instance.transform[10] = 1.0f;
 
-      /* Set user instance ID to object index (but leave low bit blank). */
-      instance.instanceId = ob->get_device_index() << 1;
+      /* Set user instance ID to object index. */
+      instance.instanceId = ob->get_device_index();
 
       /* Add some of the object visibility bits to the mask.
        * __prim_visibility contains the combined visibility bits of all instances, so is not
@@ -1508,9 +1514,6 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         else {
           /* Disable instance transform if geometry already has it applied to vertex data. */
           instance.flags |= OPTIX_INSTANCE_FLAG_DISABLE_TRANSFORM;
-          /* Non-instanced objects read ID from 'prim_object', so distinguish
-           * them from instanced objects with the low bit set. */
-          instance.instanceId |= 1;
         }
       }
     }

intern/cycles/integrator/denoiser_oidn.cpp

@@ -289,7 +289,13 @@ class OIDNDenoiseContext {
      * pixels. */
     const PassAccessorCPU pass_accessor(pass_access_info, 1.0f, num_samples_);
 
-    pass_accessor.get_render_tile_pixels(render_buffers_, buffer_params_, destination);
+    BufferParams buffer_params = buffer_params_;
+    buffer_params.window_x = 0;
+    buffer_params.window_y = 0;
+    buffer_params.window_width = buffer_params.width;
+    buffer_params.window_height = buffer_params.height;
+
+    pass_accessor.get_render_tile_pixels(render_buffers_, buffer_params, destination);
   }
 
   /* Read pass pixels using PassAccessor into a temporary buffer which is owned by the pass.. */

intern/cycles/integrator/pass_accessor.cpp

@@ -149,9 +149,6 @@ bool PassAccessor::get_render_tile_pixels(const RenderBuffers *render_buffers,
       /* Denoised passes store their final pixels, no need in special calculation. */
       get_pass_float(render_buffers, buffer_params, destination);
     }
-    else if (type == PASS_RENDER_TIME) {
-      /* TODO(sergey): Needs implementation. */
-    }
     else if (type == PASS_DEPTH) {
       get_pass_depth(render_buffers, buffer_params, destination);
     }

intern/cycles/integrator/pass_accessor_cpu.cpp

@@ -99,17 +99,22 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_float(
 {
   DCHECK_EQ(destination.stride, 0) << "Custom stride for float destination is not implemented.";
 
-  const float *buffer_data = render_buffers->buffer.data();
+  const int64_t pass_stride = buffer_params.pass_stride;
+  const int64_t buffer_row_stride = buffer_params.stride * buffer_params.pass_stride;
+
+  const float *window_data = render_buffers->buffer.data() + buffer_params.window_x * pass_stride +
+                             buffer_params.window_y * buffer_row_stride;
+
   const int pixel_stride = destination.pixel_stride ? destination.pixel_stride :
                                                       destination.num_components;
 
-  tbb::parallel_for(0, buffer_params.height, [&](int64_t y) {
-    int64_t pixel_index = y * buffer_params.width;
-    for (int64_t x = 0; x < buffer_params.width; ++x, ++pixel_index) {
-      const int64_t input_pixel_offset = pixel_index * buffer_params.pass_stride;
-      const float *buffer = buffer_data + input_pixel_offset;
-      float *pixel = destination.pixels + (pixel_index + destination.offset) * pixel_stride;
+  tbb::parallel_for(0, buffer_params.window_height, [&](int64_t y) {
+    const float *buffer = window_data + y * buffer_row_stride;
+    float *pixel = destination.pixels +
+                   (y * buffer_params.width + destination.offset) * pixel_stride;
 
+    for (int64_t x = 0; x < buffer_params.window_width;
+         ++x, buffer += pass_stride, pixel += pixel_stride) {
       processor(kfilm_convert, buffer, pixel);
     }
   });
@@ -123,26 +128,28 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_half_rgba(
     const Destination &destination,
     const Processor &processor) const
 {
-  const float *buffer_data = render_buffers->buffer.data();
+  const int64_t pass_stride = buffer_params.pass_stride;
+  const int64_t buffer_row_stride = buffer_params.stride * buffer_params.pass_stride;
+
+  const float *window_data = render_buffers->buffer.data() + buffer_params.window_x * pass_stride +
+                             buffer_params.window_y * buffer_row_stride;
 
   half4 *dst_start = destination.pixels_half_rgba + destination.offset;
   const int destination_stride = destination.stride != 0 ? destination.stride :
                                                            buffer_params.width;
 
-  tbb::parallel_for(0, buffer_params.height, [&](int64_t y) {
-    int64_t pixel_index = y * buffer_params.width;
-    half4 *dst_row_start = dst_start + y * destination_stride;
-    for (int64_t x = 0; x < buffer_params.width; ++x, ++pixel_index) {
-      const int64_t input_pixel_offset = pixel_index * buffer_params.pass_stride;
-      const float *buffer = buffer_data + input_pixel_offset;
+  tbb::parallel_for(0, buffer_params.window_height, [&](int64_t y) {
+    const float *buffer = window_data + y * buffer_row_stride;
+    half4 *pixel = dst_start + y * destination_stride;
+    for (int64_t x = 0; x < buffer_params.window_width; ++x, buffer += pass_stride, ++pixel) {
 
-      float pixel[4];
-      processor(kfilm_convert, buffer, pixel);
+      float pixel_rgba[4];
+      processor(kfilm_convert, buffer, pixel_rgba);
 
-      film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel);
+      film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel_rgba);
 
-      half4 *pixel_half_rgba = dst_row_start + x;
-      float4_store_half(&pixel_half_rgba->x, make_float4(pixel[0], pixel[1], pixel[2], pixel[3]));
+      float4_store_half(&pixel->x,
+                        make_float4(pixel_rgba[0], pixel_rgba[1], pixel_rgba[2], pixel_rgba[3]));
     }
   });
 }

intern/cycles/integrator/pass_accessor_gpu.cpp

@@ -43,10 +43,13 @@ void PassAccessorGPU::run_film_convert_kernels(DeviceKernel kernel,
   KernelFilmConvert kfilm_convert;
   init_kernel_film_convert(&kfilm_convert, buffer_params, destination);
 
-  const int work_size = buffer_params.width * buffer_params.height;
+  const int work_size = buffer_params.window_width * buffer_params.window_height;
 
   const int destination_stride = destination.stride != 0 ? destination.stride :
-                                                           buffer_params.width;
+                                                           buffer_params.window_width;
+
+  const int offset = buffer_params.window_x * buffer_params.pass_stride +
+                     buffer_params.window_y * buffer_params.stride * buffer_params.pass_stride;
 
   if (destination.d_pixels) {
     DCHECK_EQ(destination.stride, 0) << "Custom stride for float destination is not implemented.";
@@ -55,8 +58,8 @@ void PassAccessorGPU::run_film_convert_kernels(DeviceKernel kernel,
                     const_cast<device_ptr *>(&destination.d_pixels),
                     const_cast<device_ptr *>(&render_buffers->buffer.device_pointer),
                     const_cast<int *>(&work_size),
-                    const_cast<int *>(&buffer_params.width),
-                    const_cast<int *>(&buffer_params.offset),
+                    const_cast<int *>(&buffer_params.window_width),
+                    const_cast<int *>(&offset),
                     const_cast<int *>(&buffer_params.stride),
                     const_cast<int *>(&destination.offset),
                     const_cast<int *>(&destination_stride)};
@@ -70,8 +73,8 @@ void PassAccessorGPU::run_film_convert_kernels(DeviceKernel kernel,
                     const_cast<device_ptr *>(&destination.d_pixels_half_rgba),
                     const_cast<device_ptr *>(&render_buffers->buffer.device_pointer),
                     const_cast<int *>(&work_size),
-                    const_cast<int *>(&buffer_params.width),
-                    const_cast<int *>(&buffer_params.offset),
+                    const_cast<int *>(&buffer_params.window_width),
+                    const_cast<int *>(&offset),
                     const_cast<int *>(&buffer_params.stride),
                     const_cast<int *>(&destination.offset),
                     const_cast<int *>(&destination_stride)};

intern/cycles/integrator/path_trace.cpp

@@ -234,42 +234,53 @@ template<typename Callback>
 static void foreach_sliced_buffer_params(const vector<unique_ptr<PathTraceWork>> &path_trace_works,
                                          const vector<WorkBalanceInfo> &work_balance_infos,
                                          const BufferParams &buffer_params,
+                                         const int overscan,
                                          const Callback &callback)
 {
   const int num_works = path_trace_works.size();
-  const int height = buffer_params.height;
+  const int window_height = buffer_params.window_height;
 
   int current_y = 0;
   for (int i = 0; i < num_works; ++i) {
     const double weight = work_balance_infos[i].weight;
-    const int slice_height = max(lround(height * weight), 1);
+    const int slice_window_full_y = buffer_params.full_y + buffer_params.window_y + current_y;
+    const int slice_window_height = max(lround(window_height * weight), 1);
 
     /* Disallow negative values to deal with situations when there are more compute devices than
      * scan-lines. */
-    const int remaining_height = max(0, height - current_y);
+    const int remaining_window_height = max(0, window_height - current_y);
+
+    BufferParams slice_params = buffer_params;
+
+    slice_params.full_y = max(slice_window_full_y - overscan, buffer_params.full_y);
+    slice_params.window_y = slice_window_full_y - slice_params.full_y;
 
-    BufferParams slide_params = buffer_params;
-    slide_params.full_y = buffer_params.full_y + current_y;
     if (i < num_works - 1) {
-      slide_params.height = min(slice_height, remaining_height);
+      slice_params.window_height = min(slice_window_height, remaining_window_height);
     }
     else {
-      slide_params.height = remaining_height;
+      slice_params.window_height = remaining_window_height;
     }
 
-    slide_params.update_offset_stride();
+    slice_params.height = slice_params.window_y + slice_params.window_height + overscan;
+    slice_params.height = min(slice_params.height,
+                              buffer_params.height + buffer_params.full_y - slice_params.full_y);
 
-    callback(path_trace_works[i].get(), slide_params);
+    slice_params.update_offset_stride();
 
-    current_y += slide_params.height;
+    callback(path_trace_works[i].get(), slice_params);
+
+    current_y += slice_params.window_height;
   }
 }
 
 void PathTrace::update_allocated_work_buffer_params()
 {
+  const int overscan = tile_manager_.get_tile_overscan();
   foreach_sliced_buffer_params(path_trace_works_,
                                work_balance_infos_,
                                big_tile_params_,
+                               overscan,
                                [](PathTraceWork *path_trace_work, const BufferParams &params) {
                                  RenderBuffers *buffers = path_trace_work->get_render_buffers();
                                  buffers->reset(params);
@@ -282,6 +293,12 @@ static BufferParams scale_buffer_params(const BufferParams &params, int resoluti
 
   scaled_params.width = max(1, params.width / resolution_divider);
   scaled_params.height = max(1, params.height / resolution_divider);
+
+  scaled_params.window_x = params.window_x / resolution_divider;
+  scaled_params.window_y = params.window_y / resolution_divider;
+  scaled_params.window_width = params.window_width / resolution_divider;
+  scaled_params.window_height = params.window_height / resolution_divider;
+
   scaled_params.full_x = params.full_x / resolution_divider;
   scaled_params.full_y = params.full_y / resolution_divider;
   scaled_params.full_width = params.full_width / resolution_divider;
@@ -300,9 +317,12 @@ void PathTrace::update_effective_work_buffer_params(const RenderWork &render_wor
   const BufferParams scaled_big_tile_params = scale_buffer_params(big_tile_params_,
                                                                   resolution_divider);
 
+  const int overscan = tile_manager_.get_tile_overscan();
+
   foreach_sliced_buffer_params(path_trace_works_,
                                work_balance_infos_,
                                scaled_big_tile_params,
+                               overscan,
                                [&](PathTraceWork *path_trace_work, const BufferParams params) {
                                  path_trace_work->set_effective_buffer_params(
                                      scaled_full_params, scaled_big_tile_params, params);
@@ -1005,12 +1025,12 @@ bool PathTrace::set_render_tile_pixels(PassAccessor &pass_accessor,
 int2 PathTrace::get_render_tile_size() const
 {
   if (full_frame_state_.render_buffers) {
-    return make_int2(full_frame_state_.render_buffers->params.width,
-                     full_frame_state_.render_buffers->params.height);
+    return make_int2(full_frame_state_.render_buffers->params.window_width,
+                     full_frame_state_.render_buffers->params.window_height);
   }
 
   const Tile &tile = tile_manager_.get_current_tile();
-  return make_int2(tile.width, tile.height);
+  return make_int2(tile.window_width, tile.window_height);
 }
 
 int2 PathTrace::get_render_tile_offset() const
@@ -1020,7 +1040,7 @@ int2 PathTrace::get_render_tile_offset() const
   }
 
   const Tile &tile = tile_manager_.get_current_tile();
-  return make_int2(tile.x, tile.y);
+  return make_int2(tile.x + tile.window_x, tile.y + tile.window_y);
 }
 
 int2 PathTrace::get_render_size() const

intern/cycles/integrator/path_trace_work.cpp

@@ -134,7 +134,8 @@ void PathTraceWork::copy_from_denoised_render_buffers(const RenderBuffers *rende
 bool PathTraceWork::get_render_tile_pixels(const PassAccessor &pass_accessor,
                                            const PassAccessor::Destination &destination)
 {
-  const int offset_y = effective_buffer_params_.full_y - effective_big_tile_params_.full_y;
+  const int offset_y = (effective_buffer_params_.full_y + effective_buffer_params_.window_y) -
+                       (effective_big_tile_params_.full_y + effective_big_tile_params_.window_y);
   const int width = effective_buffer_params_.width;
 
   PassAccessor::Destination slice_destination = destination;
@@ -191,8 +192,10 @@ PassAccessor::Destination PathTraceWork::get_display_destination_template(
   PassAccessor::Destination destination(film_->get_display_pass());
 
   const int2 display_texture_size = display->get_texture_size();
-  const int texture_x = effective_buffer_params_.full_x - effective_full_params_.full_x;
-  const int texture_y = effective_buffer_params_.full_y - effective_full_params_.full_y;
+  const int texture_x = effective_buffer_params_.full_x - effective_full_params_.full_x +
+                        effective_buffer_params_.window_x;
+  const int texture_y = effective_buffer_params_.full_y - effective_full_params_.full_y +
+                        effective_buffer_params_.window_y;
 
   destination.offset = texture_y * display_texture_size.x + texture_x;
   destination.stride = display_texture_size.x;

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -23,6 +23,7 @@
 #include "render/buffers.h"
 #include "render/scene.h"
 #include "util/util_logging.h"
+#include "util/util_string.h"
 #include "util/util_tbb.h"
 #include "util/util_time.h"
 
@@ -30,6 +31,38 @@
 
 CCL_NAMESPACE_BEGIN
 
+static size_t estimate_single_state_size()
+{
+  size_t state_size = 0;
+
+#define KERNEL_STRUCT_BEGIN(name) for (int array_index = 0;; array_index++) {
+#define KERNEL_STRUCT_MEMBER(parent_struct, type, name, feature) state_size += sizeof(type);
+#define KERNEL_STRUCT_ARRAY_MEMBER(parent_struct, type, name, feature) state_size += sizeof(type);
+#define KERNEL_STRUCT_END(name) \
+  break; \
+  }
+#define KERNEL_STRUCT_END_ARRAY(name, cpu_array_size, gpu_array_size) \
+  if (array_index >= gpu_array_size - 1) { \
+    break; \
+  } \
+  }
+/* TODO(sergey): Look into better estimation for fields which depend on scene features. Maybe
+ * maximum state calculation should happen as `alloc_work_memory()`, so that we can react to an
+ * updated scene state here.
+ * For until then use common value. Currently this size is only used for logging, but is weak to
+ * rely on this. */
+#define KERNEL_STRUCT_VOLUME_STACK_SIZE 4
+#include "kernel/integrator/integrator_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
+
+  return state_size;
+}
+
 PathTraceWorkGPU::PathTraceWorkGPU(Device *device,
                                    Film *film,
                                    DeviceScene *device_scene,
@@ -47,7 +80,7 @@ PathTraceWorkGPU::PathTraceWorkGPU(Device *device,
       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(sizeof(IntegratorStateCPU))),
+      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)
 {
@@ -96,16 +129,27 @@ void PathTraceWorkGPU::alloc_integrator_soa()
   break; \
   }
 #define KERNEL_STRUCT_END_ARRAY(name, cpu_array_size, gpu_array_size) \
-  if (array_index == gpu_array_size - 1) { \
+  if (array_index >= gpu_array_size - 1) { \
     break; \
   } \
   }
+#define KERNEL_STRUCT_VOLUME_STACK_SIZE (device_scene_->data.volume_stack_size)
 #include "kernel/integrator/integrator_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
+
+  if (VLOG_IS_ON(3)) {
+    size_t total_soa_size = 0;
+    for (auto &&soa_memory : integrator_state_soa_) {
+      total_soa_size += soa_memory->memory_size();
+    }
+
+    VLOG(3) << "GPU SoA state size: " << string_human_readable_size(total_soa_size);
+  }
 }
 
 void PathTraceWorkGPU::alloc_integrator_queue()
@@ -712,13 +756,13 @@ void PathTraceWorkGPU::copy_to_display_naive(PathTraceDisplay *display,
 {
   const int full_x = effective_buffer_params_.full_x;
   const int full_y = effective_buffer_params_.full_y;
-  const int width = effective_buffer_params_.width;
-  const int height = effective_buffer_params_.height;
-  const int final_width = buffers_->params.width;
-  const int final_height = buffers_->params.height;
+  const int width = effective_buffer_params_.window_width;
+  const int height = effective_buffer_params_.window_height;
+  const int final_width = buffers_->params.window_width;
+  const int final_height = buffers_->params.window_height;
 
-  const int texture_x = full_x - effective_full_params_.full_x;
-  const int texture_y = full_y - effective_full_params_.full_y;
+  const int texture_x = full_x - effective_full_params_.full_x + effective_buffer_params_.window_x;
+  const int texture_y = full_y - effective_full_params_.full_y + effective_buffer_params_.window_y;
 
   /* Re-allocate display memory if needed, and make sure the device pointer is allocated.
    *

intern/cycles/kernel/CMakeLists.txt

@@ -404,16 +404,27 @@ if(WITH_CYCLES_CUDA_BINARIES)
             -cuda-toolkit-dir "${cuda_toolkit_root_dir}"
         DEPENDS ${kernel_sources} cycles_cubin_cc)
     else()
-      add_custom_command(
-        OUTPUT ${cuda_file}
-        COMMAND ${cuda_nvcc_executable}
+      set(_cuda_nvcc_args
             -arch=${arch}
             ${CUDA_NVCC_FLAGS}
             --${format}
             ${CMAKE_CURRENT_SOURCE_DIR}${cuda_kernel_src}
             --ptxas-options="-v"
-            ${cuda_flags}
-        DEPENDS ${kernel_sources})
+            ${cuda_flags})
+
+      if(WITH_COMPILER_CCACHE AND CCACHE_PROGRAM)
+        add_custom_command(
+          OUTPUT ${cuda_file}
+          COMMAND ${CCACHE_PROGRAM} ${cuda_nvcc_executable} ${_cuda_nvcc_args}
+          DEPENDS ${kernel_sources})
+      else()
+        add_custom_command(
+          OUTPUT ${cuda_file}
+          COMMAND ${cuda_nvcc_executable} ${_cuda_nvcc_args}
+          DEPENDS ${kernel_sources})
+      endif()
+
+      unset(_cuda_nvcc_args)
     endif()
     delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${cuda_file}" ${CYCLES_INSTALL_PATH}/lib)
     list(APPEND cuda_cubins ${cuda_file})
@@ -472,20 +483,10 @@ endif()
 
 # HIP module
 
-if(WITH_CYCLES_HIP_BINARIES)
+if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
   # 64 bit only
   set(HIP_BITS 64)
 
-  # HIP version
-  execute_process(COMMAND ${HIP_HIPCC_EXECUTABLE} "--version" OUTPUT_VARIABLE HIPCC_OUT)
-  string(REGEX REPLACE ".*release ([0-9]+)\\.([0-9]+).*" "\\1" HIP_VERSION_MAJOR "${HIPCC_OUT}")
-  string(REGEX REPLACE ".*release ([0-9]+)\\.([0-9]+).*" "\\2" HIP_VERSION_MINOR "${HIPCC_OUT}")
-  set(HIP_VERSION "${HIP_VERSION_MAJOR}${HIP_VERSION_MINOR}")
-
-
-  message(WARNING
-    "HIP version ${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR} detected")
-
   # build for each arch
   set(hip_sources device/hip/kernel.cpp
     ${SRC_HEADERS}
@@ -542,23 +543,24 @@ if(WITH_CYCLES_HIP_BINARIES)
         -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}
+      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})
-      set(hip_hipcc_executable ${HIP_HIPCC_EXECUTABLE})
-      set(hip_toolkit_root_dir ${HIP_TOOLKIT_ROOT_DIR})
-    if(DEFINED hip_hipcc_executable AND DEFINED hip_toolkit_root_dir)
-      # Compile regular kernel
-      CYCLES_HIP_KERNEL_ADD(${arch} ${prev_arch} kernel "" "${hip_sources}" FALSE)
-
-      if(WITH_CYCLES_HIP_BUILD_SERIAL)
-        set(prev_arch ${arch})
-      endif()
-
-      unset(hip_hipcc_executable)
-      unset(hip_toolkit_root_dir)
-    endif()
+    # Compile regular kernel
+    CYCLES_HIP_KERNEL_ADD(${arch} ${prev_arch} kernel "" "${hip_sources}" FALSE)
   endforeach()
 
   add_custom_target(cycles_kernel_hip ALL DEPENDS ${hip_fatbins})

intern/cycles/kernel/bvh/bvh_embree.h

@@ -106,9 +106,6 @@ ccl_device_inline void kernel_embree_convert_hit(const KernelGlobals *kg,
                                                  const RTCHit *hit,
                                                  Intersection *isect)
 {
-  bool is_hair = hit->geomID & 1;
-  isect->u = is_hair ? hit->u : 1.0f - hit->v - hit->u;
-  isect->v = is_hair ? hit->v : hit->u;
   isect->t = ray->tfar;
   isect->Ng = make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z);
   if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) {
@@ -121,27 +118,37 @@ ccl_device_inline void kernel_embree_convert_hit(const KernelGlobals *kg,
   else {
     isect->prim = hit->primID + (intptr_t)rtcGetGeometryUserData(
                                     rtcGetGeometry(kernel_data.bvh.scene, hit->geomID));
-    isect->object = OBJECT_NONE;
+    isect->object = hit->geomID / 2;
+  }
+
+  const bool is_hair = hit->geomID & 1;
+  if (is_hair) {
+    const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, isect->prim);
+    isect->type = segment.type;
+    isect->prim = segment.prim;
+    isect->u = hit->u;
+    isect->v = hit->v;
+  }
+  else {
+    isect->type = kernel_tex_fetch(__objects, isect->object).primitive_type;
+    isect->u = 1.0f - hit->v - hit->u;
+    isect->v = hit->u;
   }
-  isect->type = kernel_tex_fetch(__prim_type, isect->prim);
 }
 
-ccl_device_inline void kernel_embree_convert_sss_hit(const KernelGlobals *kg,
-                                                     const RTCRay *ray,
-                                                     const RTCHit *hit,
-                                                     Intersection *isect,
-                                                     int local_object_id)
+ccl_device_inline void kernel_embree_convert_sss_hit(
+    const KernelGlobals *kg, const RTCRay *ray, const RTCHit *hit, Intersection *isect, int object)
 {
   isect->u = 1.0f - hit->v - hit->u;
   isect->v = hit->u;
   isect->t = ray->tfar;
   isect->Ng = make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z);
   RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData(
-      rtcGetGeometry(kernel_data.bvh.scene, local_object_id * 2));
+      rtcGetGeometry(kernel_data.bvh.scene, object * 2));
   isect->prim = hit->primID +
                 (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID));
-  isect->object = local_object_id;
-  isect->type = kernel_tex_fetch(__prim_type, isect->prim);
+  isect->object = object;
+  isect->type = kernel_tex_fetch(__objects, object).primitive_type;
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/bvh/bvh_shadow_all.h

@@ -130,7 +130,6 @@ ccl_device_inline
         if (prim_addr >= 0) {
           const int prim_addr2 = __float_as_int(leaf.y);
           const uint type = __float_as_int(leaf.w);
-          const uint p_type = type & PRIMITIVE_ALL;
 
           /* pop */
           node_addr = traversal_stack[stack_ptr];
@@ -138,14 +137,15 @@ ccl_device_inline
 
           /* primitive intersection */
           while (prim_addr < prim_addr2) {
-            kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) == p_type);
+            kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) ==
+                          (type & PRIMITIVE_ALL));
             bool hit;
 
             /* todo: specialized intersect functions which don't fill in
              * isect unless needed and check SD_HAS_TRANSPARENT_SHADOW?
              * might give a few % performance improvement */
 
-            switch (p_type) {
+            switch (type & PRIMITIVE_ALL) {
               case PRIMITIVE_TRIANGLE: {
                 hit = triangle_intersect(
                     kg, isect, P, dir, isect_t, visibility, object, prim_addr);
@@ -163,17 +163,20 @@ ccl_device_inline
               case PRIMITIVE_MOTION_CURVE_THICK:
               case PRIMITIVE_CURVE_RIBBON:
               case PRIMITIVE_MOTION_CURVE_RIBBON: {
-                const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
-                hit = curve_intersect(kg,
-                                      isect,
-                                      P,
-                                      dir,
-                                      isect_t,
-                                      visibility,
-                                      object,
-                                      prim_addr,
-                                      ray->time,
-                                      curve_type);
+                if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) {
+                  const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
+                  if (ray->time < prim_time.x || ray->time > prim_time.y) {
+                    hit = false;
+                    break;
+                  }
+                }
+
+                const int curve_object = kernel_tex_fetch(__prim_object, prim_addr);
+                const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
+                const int curve_prim = kernel_tex_fetch(__prim_index, prim_addr);
+                hit = curve_intersect(
+                    kg, isect, P, dir, isect_t, curve_object, curve_prim, ray->time, curve_type);
+
                 break;
               }
 #endif

intern/cycles/kernel/bvh/bvh_traversal.h

@@ -165,18 +165,18 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(const KernelGlobals *kg,
             case PRIMITIVE_CURVE_RIBBON:
             case PRIMITIVE_MOTION_CURVE_RIBBON: {
               for (; prim_addr < prim_addr2; prim_addr++) {
-                const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
-                kernel_assert((curve_type & PRIMITIVE_ALL) == (type & PRIMITIVE_ALL));
-                const bool hit = curve_intersect(kg,
-                                                 isect,
-                                                 P,
-                                                 dir,
-                                                 isect->t,
-                                                 visibility,
-                                                 object,
-                                                 prim_addr,
-                                                 ray->time,
-                                                 curve_type);
+                if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) {
+                  const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
+                  if (ray->time < prim_time.x || ray->time > prim_time.y) {
+                    continue;
+                  }
+                }
+
+                const int curve_object = kernel_tex_fetch(__prim_object, prim_addr);
+                const int curve_prim = kernel_tex_fetch(__prim_index, prim_addr);
+                const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
+                const bool hit = curve_intersect(
+                    kg, isect, P, dir, isect->t, curve_object, curve_prim, ray->time, curve_type);
                 if (hit) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)

intern/cycles/kernel/bvh/bvh_util.h

@@ -118,19 +118,18 @@ ccl_device_inline void sort_intersections(Intersection *hits, uint num_hits)
 ccl_device_forceinline int intersection_get_shader_flags(const KernelGlobals *ccl_restrict kg,
                                                          const Intersection *ccl_restrict isect)
 {
-  const int prim = kernel_tex_fetch(__prim_index, isect->prim);
+  const int prim = isect->prim;
   int shader = 0;
 
 #ifdef __HAIR__
-  if (kernel_tex_fetch(__prim_type, isect->prim) & PRIMITIVE_ALL_TRIANGLE)
+  if (isect->type & PRIMITIVE_ALL_TRIANGLE)
 #endif
   {
     shader = kernel_tex_fetch(__tri_shader, prim);
   }
 #ifdef __HAIR__
   else {
-    float4 str = kernel_tex_fetch(__curves, prim);
-    shader = __float_as_int(str.z);
+    shader = kernel_tex_fetch(__curves, prim).shader_id;
   }
 #endif
 
@@ -138,21 +137,19 @@ ccl_device_forceinline int intersection_get_shader_flags(const KernelGlobals *cc
 }
 
 ccl_device_forceinline int intersection_get_shader_from_isect_prim(
-    const KernelGlobals *ccl_restrict kg, const int isect_prim)
+    const KernelGlobals *ccl_restrict kg, const int prim, const int isect_type)
 {
-  const int prim = kernel_tex_fetch(__prim_index, isect_prim);
   int shader = 0;
 
 #ifdef __HAIR__
-  if (kernel_tex_fetch(__prim_type, isect_prim) & PRIMITIVE_ALL_TRIANGLE)
+  if (isect_type & PRIMITIVE_ALL_TRIANGLE)
 #endif
   {
     shader = kernel_tex_fetch(__tri_shader, prim);
   }
 #ifdef __HAIR__
   else {
-    float4 str = kernel_tex_fetch(__curves, prim);
-    shader = __float_as_int(str.z);
+    shader = kernel_tex_fetch(__curves, prim).shader_id;
   }
 #endif
 
@@ -162,25 +159,13 @@ ccl_device_forceinline int intersection_get_shader_from_isect_prim(
 ccl_device_forceinline int intersection_get_shader(const KernelGlobals *ccl_restrict kg,
                                                    const Intersection *ccl_restrict isect)
 {
-  return intersection_get_shader_from_isect_prim(kg, isect->prim);
-}
-
-ccl_device_forceinline int intersection_get_object(const KernelGlobals *ccl_restrict kg,
-                                                   const Intersection *ccl_restrict isect)
-{
-  if (isect->object != OBJECT_NONE) {
-    return isect->object;
-  }
-
-  return kernel_tex_fetch(__prim_object, isect->prim);
+  return intersection_get_shader_from_isect_prim(kg, isect->prim, isect->type);
 }
 
 ccl_device_forceinline int intersection_get_object_flags(const KernelGlobals *ccl_restrict kg,
                                                          const Intersection *ccl_restrict isect)
 {
-  const int object = intersection_get_object(kg, isect);
-
-  return kernel_tex_fetch(__object_flag, object);
+  return kernel_tex_fetch(__object_flag, isect->object);
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/closure/bsdf_principled_diffuse.h

@@ -35,21 +35,25 @@ static_assert(sizeof(ShaderClosure) >= sizeof(PrincipledDiffuseBsdf),
               "PrincipledDiffuseBsdf is too large!");
 
 ccl_device float3 calculate_principled_diffuse_brdf(
-    const PrincipledDiffuseBsdf *bsdf, float3 N, float3 V, float3 L, float3 H, float *pdf)
+    const PrincipledDiffuseBsdf *bsdf, float3 N, float3 V, float3 L, float *pdf)
 {
   float NdotL = dot(N, L);
-  float NdotV = dot(N, V);
 
-  if (NdotL <= 0 || NdotV <= 0) {
-    *pdf = 0.0f;
+  if (NdotL <= 0) {
     return make_float3(0.0f, 0.0f, 0.0f);
   }
 
-  float LdotH = dot(L, H);
+  float NdotV = dot(N, V);
+
+  /* H = normalize(L + V);  // Bisector of an angle between L and V.
+   * LH2 = 2 * dot(L, H)^2 = 2cos(x)^2 = cos(2x) + 1 = dot(L, V) + 1,
+   * half-angle x between L and V is at most 90 deg
+   */
+  float LH2 = dot(L, V) + 1;
 
   float FL = schlick_fresnel(NdotL), FV = schlick_fresnel(NdotV);
-  const float Fd90 = 0.5f + 2.0f * LdotH * LdotH * bsdf->roughness;
-  float Fd = (1.0f * (1.0f - FL) + Fd90 * FL) * (1.0f * (1.0f - FV) + Fd90 * FV);
+  const float Fd90 = 0.5f + LH2 * bsdf->roughness;
+  float Fd = (1.0f - FL + Fd90 * FL) * (1.0f - FV + Fd90 * FV);
 
   float value = M_1_PI_F * NdotL * Fd;
 
@@ -72,11 +76,10 @@ ccl_device float3 bsdf_principled_diffuse_eval_reflect(const ShaderClosure *sc,
   float3 N = bsdf->N;
   float3 V = I;         // outgoing
   float3 L = omega_in;  // incoming
-  float3 H = normalize(L + V);
 
   if (dot(N, omega_in) > 0.0f) {
     *pdf = fmaxf(dot(N, omega_in), 0.0f) * M_1_PI_F;
-    return calculate_principled_diffuse_brdf(bsdf, N, V, L, H, pdf);
+    return calculate_principled_diffuse_brdf(bsdf, N, V, L, pdf);
   }
   else {
     *pdf = 0.0f;
@@ -112,9 +115,7 @@ ccl_device int bsdf_principled_diffuse_sample(const ShaderClosure *sc,
   sample_cos_hemisphere(N, randu, randv, omega_in, pdf);
 
   if (dot(Ng, *omega_in) > 0) {
-    float3 H = normalize(I + *omega_in);
-
-    *eval = calculate_principled_diffuse_brdf(bsdf, N, I, *omega_in, H, pdf);
+    *eval = calculate_principled_diffuse_brdf(bsdf, N, I, *omega_in, pdf);
 
 #ifdef __RAY_DIFFERENTIALS__
     // TODO: find a better approximation for the diffuse bounce

intern/cycles/kernel/device/gpu/kernel.h

@@ -424,8 +424,12 @@ ccl_device_inline void kernel_gpu_film_convert_common(const KernelFilmConvert *k
     return;
   }
 
-  const uint64_t render_buffer_offset = (uint64_t)render_pixel_index * kfilm_convert->pass_stride;
-  ccl_global const float *buffer = render_buffer + render_buffer_offset;
+  const int x = render_pixel_index % width;
+  const int y = render_pixel_index / width;
+
+  ccl_global const float *buffer = render_buffer + offset + x * kfilm_convert->pass_stride +
+                                   y * stride * kfilm_convert->pass_stride;
+
   ccl_global float *pixel = pixels +
                             (render_pixel_index + dst_offset) * kfilm_convert->pixel_stride;
 
@@ -451,17 +455,17 @@ ccl_device_inline void kernel_gpu_film_convert_half_rgba_common_rgba(
     return;
   }
 
-  const uint64_t render_buffer_offset = (uint64_t)render_pixel_index * kfilm_convert->pass_stride;
-  ccl_global const float *buffer = render_buffer + render_buffer_offset;
+  const int x = render_pixel_index % width;
+  const int y = render_pixel_index / width;
+
+  ccl_global const float *buffer = render_buffer + offset + x * kfilm_convert->pass_stride +
+                                   y * stride * kfilm_convert->pass_stride;
 
   float pixel[4];
   processor(kfilm_convert, buffer, pixel);
 
   film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel);
 
-  const int x = render_pixel_index % width;
-  const int y = render_pixel_index / width;
-
   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]));
 }

intern/cycles/kernel/device/optix/kernel.cu

@@ -41,22 +41,15 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
   return (T *)(((uint64_t)optixGetPayload_3() << 32) | optixGetPayload_2());
 }
 
-template<bool always = false> ccl_device_forceinline uint get_object_id()
+ccl_device_forceinline int get_object_id()
 {
 #ifdef __OBJECT_MOTION__
-  /* Always get the the instance ID from the TLAS.
+  /* Always get the the instance ID from the TLAS
    * There might be a motion transform node between TLAS and BLAS which does not have one. */
-  uint object = optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
+  return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
 #else
-  uint object = optixGetInstanceId();
+  return optixGetInstanceId();
 #endif
-  /* Choose between always returning object ID or only for instances. */
-  if (always || (object & 1) == 0)
-    /* Can just remove the low bit since instance always contains object ID. */
-    return object >> 1;
-  else
-    /* Set to OBJECT_NONE if this is not an instanced object. */
-    return OBJECT_NONE;
 }
 
 extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()
@@ -108,7 +101,7 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
 #endif
 
 #ifdef __BVH_LOCAL__
-  const uint object = get_object_id<true>();
+  const int object = get_object_id();
   if (object != optixGetPayload_4() /* local_object */) {
     /* Only intersect with matching object. */
     return optixIgnoreIntersection();
@@ -152,21 +145,23 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
     local_isect->num_hits = 1;
   }
 
+  const int prim = optixGetPrimitiveIndex();
+
   Intersection *isect = &local_isect->hits[hit];
   isect->t = optixGetRayTmax();
-  isect->prim = optixGetPrimitiveIndex();
+  isect->prim = prim;
   isect->object = get_object_id();
-  isect->type = kernel_tex_fetch(__prim_type, isect->prim);
+  isect->type = kernel_tex_fetch(__objects, isect->object).primitive_type;
 
   const float2 barycentrics = optixGetTriangleBarycentrics();
   isect->u = 1.0f - barycentrics.y - barycentrics.x;
   isect->v = barycentrics.x;
 
   /* Record geometric normal. */
-  const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim);
-  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0));
-  const float3 tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1));
-  const float3 tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));
+  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
+  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0));
+  const float3 tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1));
+  const float3 tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
   local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
 
   /* Continue tracing (without this the trace call would return after the first hit). */
@@ -179,25 +174,32 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
 #ifdef __SHADOW_RECORD_ALL__
   bool ignore_intersection = false;
 
-  const uint prim = optixGetPrimitiveIndex();
+  int prim = optixGetPrimitiveIndex();
+  const uint object = get_object_id();
 #  ifdef __VISIBILITY_FLAG__
   const uint visibility = optixGetPayload_4();
-  if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
+  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
     ignore_intersection = true;
   }
 #  endif
 
   float u = 0.0f, v = 0.0f;
+  int type = 0;
   if (optixIsTriangleHit()) {
     const float2 barycentrics = optixGetTriangleBarycentrics();
     u = 1.0f - barycentrics.y - barycentrics.x;
     v = barycentrics.x;
+    type = kernel_tex_fetch(__objects, object).primitive_type;
   }
 #  ifdef __HAIR__
   else {
     u = __uint_as_float(optixGetAttribute_0());
     v = __uint_as_float(optixGetAttribute_1());
 
+    const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
+    type = segment.type;
+    prim = segment.prim;
+
     /* Filter out curve endcaps. */
     if (u == 0.0f || u == 1.0f) {
       ignore_intersection = true;
@@ -245,8 +247,8 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
     isect->v = v;
     isect->t = optixGetRayTmax();
     isect->prim = prim;
-    isect->object = get_object_id();
-    isect->type = kernel_tex_fetch(__prim_type, prim);
+    isect->object = object;
+    isect->type = type;
 
 #  ifdef __TRANSPARENT_SHADOWS__
     /* Detect if this surface has a shader with transparent shadows. */
@@ -274,15 +276,14 @@ extern "C" __global__ void __anyhit__kernel_optix_volume_test()
   }
 #endif
 
+  const uint object = get_object_id();
 #ifdef __VISIBILITY_FLAG__
-  const uint prim = optixGetPrimitiveIndex();
   const uint visibility = optixGetPayload_4();
-  if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
+  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
     return optixIgnoreIntersection();
   }
 #endif
 
-  const uint object = get_object_id<true>();
   if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
     return optixIgnoreIntersection();
   }
@@ -301,9 +302,9 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
 #endif
 
 #ifdef __VISIBILITY_FLAG__
-  const uint prim = optixGetPrimitiveIndex();
+  const uint object = get_object_id();
   const uint visibility = optixGetPayload_4();
-  if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
+  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
     return optixIgnoreIntersection();
   }
 
@@ -316,28 +317,39 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
 
 extern "C" __global__ void __closesthit__kernel_optix_hit()
 {
+  const int object = get_object_id();
+  const int prim = optixGetPrimitiveIndex();
+
   optixSetPayload_0(__float_as_uint(optixGetRayTmax())); /* Intersection distance */
-  optixSetPayload_3(optixGetPrimitiveIndex());
-  optixSetPayload_4(get_object_id());
-  /* Can be PRIMITIVE_TRIANGLE and PRIMITIVE_MOTION_TRIANGLE or curve type and segment index. */
-  optixSetPayload_5(kernel_tex_fetch(__prim_type, optixGetPrimitiveIndex()));
+  optixSetPayload_4(object);
 
   if (optixIsTriangleHit()) {
     const float2 barycentrics = optixGetTriangleBarycentrics();
     optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x));
     optixSetPayload_2(__float_as_uint(barycentrics.x));
+    optixSetPayload_3(prim);
+    optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type);
   }
   else {
+    const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
     optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */
     optixSetPayload_2(optixGetAttribute_1());
+    optixSetPayload_3(segment.prim);
+    optixSetPayload_5(segment.type);
   }
 }
 
 #ifdef __HAIR__
-ccl_device_inline void optix_intersection_curve(const uint prim, const uint type)
+ccl_device_inline void optix_intersection_curve(const int prim, const int type)
 {
-  const uint object = get_object_id<true>();
+  const int object = get_object_id();
+
+#  ifdef __VISIBILITY_FLAG__
   const uint visibility = optixGetPayload_4();
+  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
+    return;
+  }
+#  endif
 
   float3 P = optixGetObjectRayOrigin();
   float3 dir = optixGetObjectRayDirection();
@@ -358,7 +370,7 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
   if (isect.t != FLT_MAX)
     isect.t *= len;
 
-  if (curve_intersect(NULL, &isect, P, dir, isect.t, visibility, object, prim, time, type)) {
+  if (curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
     optixReportIntersection(isect.t / len,
                             type & PRIMITIVE_ALL,
                             __float_as_int(isect.u),  /* Attribute_0 */
@@ -368,9 +380,9 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
 
 extern "C" __global__ void __intersection__curve_ribbon()
 {
-  const uint prim = optixGetPrimitiveIndex();
-  const uint type = kernel_tex_fetch(__prim_type, prim);
-
+  const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, optixGetPrimitiveIndex());
+  const int prim = segment.prim;
+  const int type = segment.type;
   if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) {
     optix_intersection_curve(prim, type);
   }

intern/cycles/kernel/geom/geom_curve.h

@@ -34,8 +34,8 @@ ccl_device float curve_attribute_float(const KernelGlobals *kg,
                                        float *dy)
 {
   if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
-    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     int k1 = k0 + 1;
 
     float f0 = kernel_tex_fetch(__attributes_float, desc.offset + k0);
@@ -76,8 +76,8 @@ ccl_device float2 curve_attribute_float2(const KernelGlobals *kg,
                                          float2 *dy)
 {
   if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
-    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     int k1 = k0 + 1;
 
     float2 f0 = kernel_tex_fetch(__attributes_float2, desc.offset + k0);
@@ -122,8 +122,8 @@ ccl_device float3 curve_attribute_float3(const KernelGlobals *kg,
                                          float3 *dy)
 {
   if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
-    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     int k1 = k0 + 1;
 
     float3 f0 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + k0));
@@ -164,8 +164,8 @@ ccl_device float4 curve_attribute_float4(const KernelGlobals *kg,
                                          float4 *dy)
 {
   if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
-    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     int k1 = k0 + 1;
 
     float4 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + k0);
@@ -206,8 +206,8 @@ ccl_device float curve_thickness(const KernelGlobals *kg, const ShaderData *sd)
   float r = 0.0f;
 
   if (sd->type & PRIMITIVE_ALL_CURVE) {
-    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
     int k1 = k0 + 1;
 
     float4 P_curve[2];
@@ -231,8 +231,8 @@ ccl_device float curve_thickness(const KernelGlobals *kg, const ShaderData *sd)
 
 ccl_device float3 curve_motion_center_location(const KernelGlobals *kg, const ShaderData *sd)
 {
-  float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-  int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+  KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
+  int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
   int k1 = k0 + 1;
 
   float4 P_curve[2];

intern/cycles/kernel/geom/geom_curve_intersect.h

@@ -630,33 +630,19 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
                                             const float3 P,
                                             const float3 dir,
                                             const float tmax,
-                                            uint visibility,
                                             int object,
-                                            int curveAddr,
+                                            int prim,
                                             float time,
                                             int type)
 {
   const bool is_motion = (type & PRIMITIVE_ALL_MOTION);
 
-#  ifndef __KERNEL_OPTIX__ /* See OptiX motion flag OPTIX_MOTION_FLAG_[START|END]_VANISH */
-  if (is_motion && kernel_data.bvh.use_bvh_steps) {
-    const float2 prim_time = kernel_tex_fetch(__prim_time, curveAddr);
-    if (time < prim_time.x || time > prim_time.y) {
-      return false;
-    }
-  }
-#  endif
+  KernelCurve kcurve = kernel_tex_fetch(__curves, prim);
 
-  int segment = PRIMITIVE_UNPACK_SEGMENT(type);
-  int prim = kernel_tex_fetch(__prim_index, curveAddr);
-
-  float4 v00 = kernel_tex_fetch(__curves, prim);
-
-  int k0 = __float_as_int(v00.x) + segment;
+  int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
   int k1 = k0 + 1;
-
-  int ka = max(k0 - 1, __float_as_int(v00.x));
-  int kb = min(k1 + 1, __float_as_int(v00.x) + __float_as_int(v00.y) - 1);
+  int ka = max(k0 - 1, kcurve.first_key);
+  int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
 
   float4 curve[4];
   if (!is_motion) {
@@ -666,21 +652,14 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
     curve[3] = kernel_tex_fetch(__curve_keys, kb);
   }
   else {
-    int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
-    motion_curve_keys(kg, fobject, prim, time, ka, k0, k1, kb, curve);
+    motion_curve_keys(kg, object, prim, time, ka, k0, k1, kb, curve);
   }
 
-#  ifdef __VISIBILITY_FLAG__
-  if (!(kernel_tex_fetch(__prim_visibility, curveAddr) & visibility)) {
-    return false;
-  }
-#  endif
-
   if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) {
     /* todo: adaptive number of subdivisions could help performance here. */
     const int subdivisions = kernel_data.bvh.curve_subdivisions;
     if (ribbon_intersect(P, dir, tmax, subdivisions, curve, isect)) {
-      isect->prim = curveAddr;
+      isect->prim = prim;
       isect->object = object;
       isect->type = type;
       return true;
@@ -690,7 +669,7 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
   }
   else {
     if (curve_intersect_recursive(P, dir, tmax, curve, isect)) {
-      isect->prim = curveAddr;
+      isect->prim = prim;
       isect->object = object;
       isect->type = type;
       return true;
@@ -708,7 +687,7 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
                                           const int isect_object,
                                           const int isect_prim)
 {
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_inverse_transform(kg, sd);
 
     P = transform_point(&tfm, P);
@@ -716,14 +695,12 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
     D = safe_normalize_len(D, &t);
   }
 
-  int prim = kernel_tex_fetch(__prim_index, isect_prim);
-  float4 v00 = kernel_tex_fetch(__curves, prim);
+  KernelCurve kcurve = kernel_tex_fetch(__curves, isect_prim);
 
-  int k0 = __float_as_int(v00.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+  int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
   int k1 = k0 + 1;
-
-  int ka = max(k0 - 1, __float_as_int(v00.x));
-  int kb = min(k1 + 1, __float_as_int(v00.x) + __float_as_int(v00.y) - 1);
+  int ka = max(k0 - 1, kcurve.first_key);
+  int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
 
   float4 P_curve[4];
 
@@ -780,15 +757,13 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
   sd->dPdv = cross(dPdu, sd->Ng);
 #  endif
 
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }
 
   sd->P = P;
-
-  float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
-  sd->shader = __float_as_int(curvedata.z);
+  sd->shader = kernel_tex_fetch(__curves, sd->prim).shader_id;
 }
 
 #endif

intern/cycles/kernel/geom/geom_motion_triangle.h

@@ -72,9 +72,9 @@ ccl_device_inline void motion_triangle_verts_for_step(const KernelGlobals *kg,
 {
   if (step == numsteps) {
     /* center step: regular vertex location */
-    verts[0] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-    verts[1] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-    verts[2] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+    verts[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+    verts[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+    verts[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
   }
   else {
     /* center step not store in this array */

intern/cycles/kernel/geom/geom_motion_triangle_intersect.h

@@ -44,7 +44,7 @@ ccl_device_inline float3 motion_triangle_refine(const KernelGlobals *kg,
                                                 float3 verts[3])
 {
 #ifdef __INTERSECTION_REFINE__
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     if (UNLIKELY(t == 0.0f)) {
       return P;
     }
@@ -70,7 +70,7 @@ ccl_device_inline float3 motion_triangle_refine(const KernelGlobals *kg,
   /* Compute refined position. */
   P = P + D * rt;
 
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }
@@ -106,7 +106,7 @@ ccl_device_inline
   return motion_triangle_refine(kg, sd, P, D, t, isect_object, isect_prim, verts);
 #  else
 #    ifdef __INTERSECTION_REFINE__
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_inverse_transform(kg, sd);
 
     P = transform_point(&tfm, P);
@@ -128,7 +128,7 @@ ccl_device_inline
 
   P = P + D * rt;
 
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }
@@ -186,8 +186,9 @@ ccl_device_inline bool motion_triangle_intersect(const KernelGlobals *kg,
       isect->t = t;
       isect->u = u;
       isect->v = v;
-      isect->prim = prim_addr;
-      isect->object = object;
+      isect->prim = prim;
+      isect->object = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) :
+                                                object;
       isect->type = PRIMITIVE_MOTION_TRIANGLE;
       return true;
     }
@@ -288,8 +289,8 @@ ccl_device_inline bool motion_triangle_intersect_local(const KernelGlobals *kg,
   isect->t = t;
   isect->u = u;
   isect->v = v;
-  isect->prim = prim_addr;
-  isect->object = object;
+  isect->prim = prim;
+  isect->object = local_object;
   isect->type = PRIMITIVE_MOTION_TRIANGLE;
 
   /* Record geometric normal. */

intern/cycles/kernel/geom/geom_shader_data.h

@@ -37,7 +37,7 @@ ccl_device void shader_setup_object_transforms(const KernelGlobals *ccl_restrict
 #endif
 
 /* TODO: break this up if it helps reduce register pressure to load data from
- * global memory as we write it to shaderdata. */
+ * global memory as we write it to shader-data. */
 ccl_device_inline void shader_setup_from_ray(const KernelGlobals *ccl_restrict kg,
                                              ShaderData *ccl_restrict sd,
                                              const Ray *ccl_restrict ray,
@@ -52,10 +52,9 @@ ccl_device_inline void shader_setup_from_ray(const KernelGlobals *ccl_restrict k
   sd->v = isect->v;
   sd->ray_length = isect->t;
   sd->type = isect->type;
-  sd->object = (isect->object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, isect->prim) :
-                                                isect->object;
+  sd->object = isect->object;
   sd->object_flag = kernel_tex_fetch(__object_flag, sd->object);
-  sd->prim = kernel_tex_fetch(__prim_index, isect->prim);
+  sd->prim = isect->prim;
   sd->lamp = LAMP_NONE;
   sd->flag = 0;
 
@@ -103,7 +102,7 @@ ccl_device_inline void shader_setup_from_ray(const KernelGlobals *ccl_restrict k
 
   sd->flag |= kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
 
-  if (isect->object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     /* instance transform */
     object_normal_transform_auto(kg, sd, &sd->N);
     object_normal_transform_auto(kg, sd, &sd->Ng);

intern/cycles/kernel/geom/geom_triangle.h

@@ -29,9 +29,9 @@ ccl_device_inline float3 triangle_normal(const KernelGlobals *kg, ShaderData *sd
 {
   /* load triangle vertices */
   const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
-  const float3 v0 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-  const float3 v1 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-  const float3 v2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+  const float3 v0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+  const float3 v1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+  const float3 v2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
 
   /* return normal */
   if (sd->object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
@@ -54,9 +54,9 @@ ccl_device_inline void triangle_point_normal(const KernelGlobals *kg,
 {
   /* load triangle vertices */
   const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
-  float3 v0 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-  float3 v1 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-  float3 v2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+  float3 v0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+  float3 v1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+  float3 v2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
   /* compute point */
   float t = 1.0f - u - v;
   *P = (u * v0 + v * v1 + t * v2);
@@ -78,9 +78,9 @@ ccl_device_inline void triangle_point_normal(const KernelGlobals *kg,
 ccl_device_inline void triangle_vertices(const KernelGlobals *kg, int prim, float3 P[3])
 {
   const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
-  P[0] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-  P[1] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-  P[2] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+  P[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+  P[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+  P[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
 }
 
 /* Triangle vertex locations and vertex normals */
@@ -91,9 +91,9 @@ ccl_device_inline void triangle_vertices_and_normals(const KernelGlobals *kg,
                                                      float3 N[3])
 {
   const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
-  P[0] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-  P[1] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-  P[2] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+  P[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+  P[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+  P[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
   N[0] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.x));
   N[1] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.y));
   N[2] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.z));
@@ -145,9 +145,9 @@ ccl_device_inline void triangle_dPdudv(const KernelGlobals *kg,
 {
   /* fetch triangle vertex coordinates */
   const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
-  const float3 p0 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
-  const float3 p1 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
-  const float3 p2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
+  const float3 p0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
+  const float3 p1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
+  const float3 p2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
 
   /* compute derivatives of P w.r.t. uv */
   *dPdu = (p0 - p2);

intern/cycles/kernel/geom/geom_triangle_intersect.h

@@ -35,13 +35,14 @@ ccl_device_inline bool triangle_intersect(const KernelGlobals *kg,
                                           int object,
                                           int prim_addr)
 {
-  const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
+  const int prim = kernel_tex_fetch(__prim_index, prim_addr);
+  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
 #if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
-  const ssef *ssef_verts = (ssef *)&kg->__prim_tri_verts.data[tri_vindex];
+  const ssef *ssef_verts = (ssef *)&kg->__tri_verts.data[tri_vindex];
 #else
-  const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0),
-               tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1),
-               tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2);
+  const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
+               tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
+               tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
 #endif
   float t, u, v;
   if (ray_triangle_intersect(P,
@@ -64,8 +65,9 @@ ccl_device_inline bool triangle_intersect(const KernelGlobals *kg,
     if (kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
 #endif
     {
-      isect->prim = prim_addr;
-      isect->object = object;
+      isect->object = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) :
+                                                object;
+      isect->prim = prim;
       isect->type = PRIMITIVE_TRIANGLE;
       isect->u = u;
       isect->v = v;
@@ -102,13 +104,14 @@ ccl_device_inline bool triangle_intersect_local(const KernelGlobals *kg,
     }
   }
 
-  const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
+  const int prim = kernel_tex_fetch(__prim_index, prim_addr);
+  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
 #  if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
-  const ssef *ssef_verts = (ssef *)&kg->__prim_tri_verts.data[tri_vindex];
+  const ssef *ssef_verts = (ssef *)&kg->__tri_verts.data[tri_vindex];
 #  else
-  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0)),
-               tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1)),
-               tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));
+  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0)),
+               tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1)),
+               tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
 #  endif
   float t, u, v;
   if (!ray_triangle_intersect(P,
@@ -167,8 +170,8 @@ ccl_device_inline bool triangle_intersect_local(const KernelGlobals *kg,
 
   /* Record intersection. */
   Intersection *isect = &local_isect->hits[hit];
-  isect->prim = prim_addr;
-  isect->object = object;
+  isect->prim = prim;
+  isect->object = local_object;
   isect->type = PRIMITIVE_TRIANGLE;
   isect->u = u;
   isect->v = v;
@@ -176,9 +179,9 @@ ccl_device_inline bool triangle_intersect_local(const KernelGlobals *kg,
 
   /* Record geometric normal. */
 #  if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
-  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0)),
-               tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1)),
-               tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));
+  const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0)),
+               tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1)),
+               tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
 #  endif
   local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
 
@@ -206,7 +209,7 @@ ccl_device_inline float3 triangle_refine(const KernelGlobals *kg,
                                          const int isect_prim)
 {
 #ifdef __INTERSECTION_REFINE__
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     if (UNLIKELY(t == 0.0f)) {
       return P;
     }
@@ -219,10 +222,10 @@ ccl_device_inline float3 triangle_refine(const KernelGlobals *kg,
 
   P = P + D * t;
 
-  const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect_prim);
-  const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0),
-               tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1),
-               tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2);
+  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
+  const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
+               tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
+               tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
   float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
   float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
   float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
@@ -239,7 +242,7 @@ ccl_device_inline float3 triangle_refine(const KernelGlobals *kg,
     P = P + D * rt;
   }
 
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }
@@ -265,7 +268,7 @@ ccl_device_inline float3 triangle_refine_local(const KernelGlobals *kg,
   /* t is always in world space with OptiX. */
   return triangle_refine(kg, sd, P, D, t, isect_object, isect_prim);
 #else
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_inverse_transform(kg, sd);
 
     P = transform_point(&tfm, P);
@@ -276,10 +279,10 @@ ccl_device_inline float3 triangle_refine_local(const KernelGlobals *kg,
   P = P + D * t;
 
 #  ifdef __INTERSECTION_REFINE__
-  const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect_prim);
-  const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0),
-               tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1),
-               tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2);
+  const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
+  const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
+               tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
+               tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
   float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
   float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
   float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
@@ -297,7 +300,7 @@ ccl_device_inline float3 triangle_refine_local(const KernelGlobals *kg,
   }
 #  endif /* __INTERSECTION_REFINE__ */
 
-  if (isect_object != OBJECT_NONE) {
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }

intern/cycles/kernel/integrator/integrator_init_from_bake.h

@@ -109,9 +109,17 @@ ccl_device bool integrator_init_from_bake(INTEGRATOR_STATE_ARGS,
   }
 
   /* Position and normal on triangle. */
+  const int object = kernel_data.bake.object_index;
   float3 P, Ng;
   int shader;
-  triangle_point_normal(kg, kernel_data.bake.object_index, prim, u, v, &P, &Ng, &shader);
+  triangle_point_normal(kg, object, prim, u, v, &P, &Ng, &shader);
+
+  const int object_flag = kernel_tex_fetch(__object_flag, object);
+  if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+    Transform tfm = object_fetch_transform(kg, object, OBJECT_TRANSFORM);
+    P = transform_point_auto(&tfm, P);
+  }
+
   if (kernel_data.film.pass_background != PASS_UNUSED) {
     /* Environment baking. */
 
@@ -130,8 +138,13 @@ ccl_device bool integrator_init_from_bake(INTEGRATOR_STATE_ARGS,
   }
   else {
     /* Surface baking. */
-    const float3 N = (shader & SHADER_SMOOTH_NORMAL) ? triangle_smooth_normal(kg, Ng, prim, u, v) :
-                                                       Ng;
+    float3 N = (shader & SHADER_SMOOTH_NORMAL) ? triangle_smooth_normal(kg, Ng, prim, u, v) : Ng;
+
+    if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+      Transform itfm = object_fetch_transform(kg, object, OBJECT_INVERSE_TRANSFORM);
+      N = normalize(transform_direction_transposed(&itfm, N));
+      Ng = normalize(transform_direction_transposed(&itfm, Ng));
+    }
 
     /* Setup ray. */
     Ray ray ccl_optional_struct_init;
@@ -143,6 +156,12 @@ ccl_device bool integrator_init_from_bake(INTEGRATOR_STATE_ARGS,
     /* Setup differentials. */
     float3 dPdu, dPdv;
     triangle_dPdudv(kg, prim, &dPdu, &dPdv);
+    if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+      Transform tfm = object_fetch_transform(kg, object, OBJECT_TRANSFORM);
+      dPdu = transform_direction(&tfm, dPdu);
+      dPdv = transform_direction(&tfm, dPdv);
+    }
+
     differential3 dP;
     dP.dx = dPdu * dudx + dPdv * dvdx;
     dP.dy = dPdu * dudy + dPdv * dvdy;

intern/cycles/kernel/integrator/integrator_intersect_closest.h

@@ -123,7 +123,7 @@ ccl_device_forceinline void integrator_intersect_shader_next_kernel(
 #ifdef __SHADOW_CATCHER__
   const int object_flags = intersection_get_object_flags(kg, isect);
   if (kernel_shadow_catcher_split(INTEGRATOR_STATE_PASS, object_flags)) {
-    if (kernel_data.film.use_approximate_shadow_catcher && !kernel_data.background.transparent) {
+    if (kernel_data.film.pass_background != PASS_UNUSED && !kernel_data.background.transparent) {
       INTEGRATOR_STATE_WRITE(path, flag) |= PATH_RAY_SHADOW_CATCHER_BACKGROUND;
 
       if (use_raytrace_kernel) {
@@ -160,10 +160,7 @@ ccl_device void integrator_intersect_closest(INTEGRATOR_STATE_ARGS)
   if (path_state_ao_bounce(INTEGRATOR_STATE_PASS)) {
     ray.t = kernel_data.integrator.ao_bounces_distance;
 
-    const int last_object = last_isect_object != OBJECT_NONE ?
-                                last_isect_object :
-                                kernel_tex_fetch(__prim_object, last_isect_prim);
-    const float object_ao_distance = kernel_tex_fetch(__objects, last_object).ao_distance;
+    const float object_ao_distance = kernel_tex_fetch(__objects, last_isect_object).ao_distance;
     if (object_ao_distance != 0.0f) {
       ray.t = object_ao_distance;
     }

intern/cycles/kernel/integrator/integrator_intersect_volume_stack.h

@@ -38,10 +38,13 @@ ccl_device void integrator_volume_stack_update_for_subsurface(INTEGRATOR_STATE_A
   volume_ray.P = from_P;
   volume_ray.D = normalize_len(to_P - from_P, &volume_ray.t);
 
+  /* Store to avoid global fetches on every intersection step. */
+  const uint volume_stack_size = kernel_data.volume_stack_size;
+
 #ifdef __VOLUME_RECORD_ALL__
-  Intersection hits[2 * VOLUME_STACK_SIZE + 1];
+  Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1];
   uint num_hits = scene_intersect_volume_all(
-      kg, &volume_ray, hits, 2 * VOLUME_STACK_SIZE, PATH_RAY_ALL_VISIBILITY);
+      kg, &volume_ray, hits, 2 * volume_stack_size, PATH_RAY_ALL_VISIBILITY);
   if (num_hits > 0) {
     Intersection *isect = hits;
 
@@ -55,7 +58,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(INTEGRATOR_STATE_A
 #else
   Intersection isect;
   int step = 0;
-  while (step < 2 * VOLUME_STACK_SIZE &&
+  while (step < 2 * volume_stack_size &&
          scene_intersect_volume(kg, &volume_ray, &isect, PATH_RAY_ALL_VISIBILITY)) {
     shader_setup_from_ray(kg, stack_sd, &volume_ray, &isect);
     volume_stack_enter_exit(INTEGRATOR_STATE_PASS, stack_sd);
@@ -91,12 +94,15 @@ ccl_device void integrator_intersect_volume_stack(INTEGRATOR_STATE_ARGS)
     stack_index++;
   }
 
+  /* Store to avoid global fetches on every intersection step. */
+  const uint volume_stack_size = kernel_data.volume_stack_size;
+
 #ifdef __VOLUME_RECORD_ALL__
-  Intersection hits[2 * VOLUME_STACK_SIZE + 1];
+  Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1];
   uint num_hits = scene_intersect_volume_all(
-      kg, &volume_ray, hits, 2 * VOLUME_STACK_SIZE, visibility);
+      kg, &volume_ray, hits, 2 * volume_stack_size, visibility);
   if (num_hits > 0) {
-    int enclosed_volumes[VOLUME_STACK_SIZE];
+    int enclosed_volumes[MAX_VOLUME_STACK_SIZE];
     Intersection *isect = hits;
 
     qsort(hits, num_hits, sizeof(Intersection), intersections_compare);
@@ -121,7 +127,7 @@ ccl_device void integrator_intersect_volume_stack(INTEGRATOR_STATE_ARGS)
             break;
           }
         }
-        if (need_add && stack_index < VOLUME_STACK_SIZE - 1) {
+        if (need_add && stack_index < volume_stack_size - 1) {
           const VolumeStack new_entry = {stack_sd->object, stack_sd->shader};
           integrator_state_write_volume_stack(INTEGRATOR_STATE_PASS, stack_index, new_entry);
           ++stack_index;
@@ -136,11 +142,12 @@ ccl_device void integrator_intersect_volume_stack(INTEGRATOR_STATE_ARGS)
     }
   }
 #else
-  int enclosed_volumes[VOLUME_STACK_SIZE];
+  /* CUDA does not support definition of a variable size arrays, so use the maximum possible. */
+  int enclosed_volumes[MAX_VOLUME_STACK_SIZE];
   int step = 0;
 
-  while (stack_index < VOLUME_STACK_SIZE - 1 && enclosed_index < VOLUME_STACK_SIZE - 1 &&
-         step < 2 * VOLUME_STACK_SIZE) {
+  while (stack_index < volume_stack_size - 1 && enclosed_index < volume_stack_size - 1 &&
+         step < 2 * volume_stack_size) {
     Intersection isect;
     if (!scene_intersect_volume(kg, &volume_ray, &isect, visibility)) {
       break;

intern/cycles/kernel/integrator/integrator_shade_background.h

@@ -192,7 +192,8 @@ ccl_device void integrator_shade_background(INTEGRATOR_STATE_ARGS,
     INTEGRATOR_STATE_WRITE(path, flag) &= ~PATH_RAY_SHADOW_CATCHER_BACKGROUND;
 
     const int isect_prim = INTEGRATOR_STATE(isect, prim);
-    const int shader = intersection_get_shader_from_isect_prim(kg, isect_prim);
+    const int isect_type = INTEGRATOR_STATE(isect, type);
+    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)) {

intern/cycles/kernel/integrator/integrator_state.h

@@ -59,8 +59,6 @@ CCL_NAMESPACE_BEGIN
  *
  * TODO: these could be made dynamic depending on the features used in the scene. */
 
-#define INTEGRATOR_VOLUME_STACK_SIZE VOLUME_STACK_SIZE
-
 #define INTEGRATOR_SHADOW_ISECT_SIZE_CPU 1024
 #define INTEGRATOR_SHADOW_ISECT_SIZE_GPU 4
 
@@ -85,12 +83,14 @@ typedef struct IntegratorStateCPU {
 #define KERNEL_STRUCT_END_ARRAY(name, cpu_size, gpu_size) \
   } \
   name[cpu_size];
+#define KERNEL_STRUCT_VOLUME_STACK_SIZE MAX_VOLUME_STACK_SIZE
 #include "kernel/integrator/integrator_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
 } IntegratorStateCPU;
 
 /* Path Queue
@@ -114,12 +114,14 @@ typedef struct IntegratorStateGPU {
 #define KERNEL_STRUCT_END_ARRAY(name, cpu_size, gpu_size) \
   } \
   name[gpu_size];
+#define KERNEL_STRUCT_VOLUME_STACK_SIZE MAX_VOLUME_STACK_SIZE
 #include "kernel/integrator/integrator_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
 
   /* Count number of queued kernels. */
   IntegratorQueueCounter *queue_counter;

intern/cycles/kernel/integrator/integrator_state_template.h

@@ -107,7 +107,9 @@ KERNEL_STRUCT_END(subsurface)
 KERNEL_STRUCT_BEGIN(volume_stack)
 KERNEL_STRUCT_ARRAY_MEMBER(volume_stack, int, object, KERNEL_FEATURE_VOLUME)
 KERNEL_STRUCT_ARRAY_MEMBER(volume_stack, int, shader, KERNEL_FEATURE_VOLUME)
-KERNEL_STRUCT_END_ARRAY(volume_stack, INTEGRATOR_VOLUME_STACK_SIZE, INTEGRATOR_VOLUME_STACK_SIZE)
+KERNEL_STRUCT_END_ARRAY(volume_stack,
+                        KERNEL_STRUCT_VOLUME_STACK_SIZE,
+                        KERNEL_STRUCT_VOLUME_STACK_SIZE)
 
 /********************************* Shadow Path State **************************/
 
@@ -163,5 +165,5 @@ KERNEL_STRUCT_BEGIN(shadow_volume_stack)
 KERNEL_STRUCT_ARRAY_MEMBER(shadow_volume_stack, int, object, KERNEL_FEATURE_VOLUME)
 KERNEL_STRUCT_ARRAY_MEMBER(shadow_volume_stack, int, shader, KERNEL_FEATURE_VOLUME)
 KERNEL_STRUCT_END_ARRAY(shadow_volume_stack,
-                        INTEGRATOR_VOLUME_STACK_SIZE,
-                        INTEGRATOR_VOLUME_STACK_SIZE)
+                        KERNEL_STRUCT_VOLUME_STACK_SIZE,
+                        KERNEL_STRUCT_VOLUME_STACK_SIZE)

intern/cycles/kernel/integrator/integrator_state_util.h

@@ -155,7 +155,7 @@ ccl_device_forceinline void integrator_state_read_shadow_isect(INTEGRATOR_STATE_
 ccl_device_forceinline void integrator_state_copy_volume_stack_to_shadow(INTEGRATOR_STATE_ARGS)
 {
   if (kernel_data.kernel_features & KERNEL_FEATURE_VOLUME) {
-    for (int i = 0; i < INTEGRATOR_VOLUME_STACK_SIZE; i++) {
+    for (int i = 0; i < kernel_data.volume_stack_size; i++) {
       INTEGRATOR_STATE_ARRAY_WRITE(shadow_volume_stack, i, object) = INTEGRATOR_STATE_ARRAY(
           volume_stack, i, object);
       INTEGRATOR_STATE_ARRAY_WRITE(shadow_volume_stack, i, shader) = INTEGRATOR_STATE_ARRAY(
@@ -223,6 +223,8 @@ ccl_device_inline void integrator_state_copy_only(const IntegratorState to_state
     while (index < gpu_array_size) \
       ;
 
+#  define KERNEL_STRUCT_VOLUME_STACK_SIZE kernel_data.volume_stack_size
+
 #  include "kernel/integrator/integrator_state_template.h"
 
 #  undef KERNEL_STRUCT_BEGIN
@@ -230,6 +232,7 @@ ccl_device_inline void integrator_state_copy_only(const IntegratorState to_state
 #  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_state_move(const IntegratorState to_state,

intern/cycles/kernel/integrator/integrator_subsurface.h

@@ -577,7 +577,7 @@ ccl_device_inline bool subsurface_scatter(INTEGRATOR_STATE_ARGS)
 #  ifdef __VOLUME__
   /* Update volume stack if needed. */
   if (kernel_data.integrator.use_volumes) {
-    const int object = intersection_get_object(kg, &ss_isect.hits[0]);
+    const int object = ss_isect.hits[0].object;
     const int object_flag = kernel_tex_fetch(__object_flag, object);
 
     if (object_flag & SD_OBJECT_INTERSECTS_VOLUME) {

intern/cycles/kernel/integrator/integrator_volume_stack.h

@@ -72,7 +72,7 @@ ccl_device void volume_stack_enter_exit(INTEGRATOR_STATE_ARGS,
     }
 
     /* If we exceed the stack limit, ignore. */
-    if (i >= VOLUME_STACK_SIZE - 1) {
+    if (i >= kernel_data.volume_stack_size - 1) {
       return;
     }
 

intern/cycles/kernel/kernel_shader.h

@@ -780,8 +780,8 @@ ccl_device_inline void shader_eval_volume(INTEGRATOR_STATE_CONST_ARGS,
       break;
     }
 
-    /* setup shaderdata from stack. it's mostly setup already in
-     * shader_setup_from_volume, this switching should be quick */
+    /* Setup shader-data from stack. it's mostly setup already in
+     * shader_setup_from_volume, this switching should be quick. */
     sd->object = entry.object;
     sd->lamp = LAMP_NONE;
     sd->shader = entry.shader;

intern/cycles/kernel/kernel_textures.h

@@ -18,11 +18,9 @@
 #  define KERNEL_TEX(type, name)
 #endif
 
-/* bvh */
+/* BVH2, not used for OptiX or Embree. */
 KERNEL_TEX(float4, __bvh_nodes)
 KERNEL_TEX(float4, __bvh_leaf_nodes)
-KERNEL_TEX(float4, __prim_tri_verts)
-KERNEL_TEX(uint, __prim_tri_index)
 KERNEL_TEX(uint, __prim_type)
 KERNEL_TEX(uint, __prim_visibility)
 KERNEL_TEX(uint, __prim_index)
@@ -46,10 +44,12 @@ KERNEL_TEX(float4, __tri_vnormal)
 KERNEL_TEX(uint4, __tri_vindex)
 KERNEL_TEX(uint, __tri_patch)
 KERNEL_TEX(float2, __tri_patch_uv)
+KERNEL_TEX(float4, __tri_verts)
 
 /* curves */
-KERNEL_TEX(float4, __curves)
+KERNEL_TEX(KernelCurve, __curves)
 KERNEL_TEX(float4, __curve_keys)
+KERNEL_TEX(KernelCurveSegment, __curve_segments)
 
 /* patches */
 KERNEL_TEX(uint, __patches)

intern/cycles/kernel/kernel_types.h

@@ -61,8 +61,6 @@ CCL_NAMESPACE_BEGIN
 #define ID_NONE (0.0f)
 #define PASS_UNUSED (~0)
 
-#define VOLUME_STACK_SIZE 4
-
 /* Kernel features */
 #define __SOBOL__
 #define __DPDU__
@@ -190,7 +188,7 @@ enum SamplingPattern {
   SAMPLING_NUM_PATTERNS,
 };
 
-/* these flags values correspond to raytypes in osl.cpp, so keep them in sync! */
+/* These flags values correspond to `raytypes` in `osl.cpp`, so keep them in sync! */
 
 enum PathRayFlag {
   /* --------------------------------------------------------------------
@@ -360,7 +358,6 @@ typedef enum PassType {
   PASS_MATERIAL_ID,
   PASS_MOTION,
   PASS_MOTION_WEIGHT,
-  PASS_RENDER_TIME,
   PASS_CRYPTOMATTE,
   PASS_AOV_COLOR,
   PASS_AOV_VALUE,
@@ -609,6 +606,12 @@ typedef struct AttributeDescriptor {
 #  define MAX_CLOSURE __MAX_CLOSURE__
 #endif
 
+#ifndef __MAX_VOLUME_STACK_SIZE__
+#  define MAX_VOLUME_STACK_SIZE 32
+#else
+#  define MAX_VOLUME_STACK_SIZE __MAX_VOLUME_STACK_SIZE__
+#endif
+
 #define MAX_VOLUME_CLOSURE 8
 
 /* This struct is the base class for all closures. The common members are
@@ -863,9 +866,6 @@ typedef struct VolumeStack {
 
 /* Struct to gather multiple nearby intersections. */
 typedef struct LocalIntersection {
-  Ray ray;
-  float3 weight[LOCAL_MAX_HITS];
-
   int num_hits;
   struct Intersection hits[LOCAL_MAX_HITS];
   float3 Ng[LOCAL_MAX_HITS];
@@ -1224,7 +1224,7 @@ typedef struct KernelData {
   uint kernel_features;
   uint max_closures;
   uint max_shaders;
-  uint pad;
+  uint volume_stack_size;
 
   KernelCamera cam;
   KernelFilm film;
@@ -1267,10 +1267,25 @@ typedef struct KernelObject {
 
   float ao_distance;
 
-  float pad1, pad2;
+  uint visibility;
+  int primitive_type;
 } KernelObject;
 static_assert_align(KernelObject, 16);
 
+typedef struct KernelCurve {
+  int shader_id;
+  int first_key;
+  int num_keys;
+  int type;
+} KernelCurve;
+static_assert_align(KernelCurve, 16);
+
+typedef struct KernelCurveSegment {
+  int prim;
+  int type;
+} KernelCurveSegment;
+static_assert_align(KernelCurveSegment, 8);
+
 typedef struct KernelSpotLight {
   float radius;
   float invarea;

intern/cycles/kernel/osl/osl_services.cpp

@@ -110,6 +110,7 @@ ustring OSLRenderServices::u_curve_thickness("geom:curve_thickness");
 ustring OSLRenderServices::u_curve_length("geom:curve_length");
 ustring OSLRenderServices::u_curve_tangent_normal("geom:curve_tangent_normal");
 ustring OSLRenderServices::u_curve_random("geom:curve_random");
+ustring OSLRenderServices::u_normal_map_normal("geom:normal_map_normal");
 ustring OSLRenderServices::u_path_ray_length("path:ray_length");
 ustring OSLRenderServices::u_path_ray_depth("path:ray_depth");
 ustring OSLRenderServices::u_path_diffuse_depth("path:diffuse_depth");
@@ -985,8 +986,18 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobals *kg,
     float3 f = curve_tangent_normal(kg, sd);
     return set_attribute_float3(f, type, derivatives, val);
   }
-  else
+  else if (name == u_normal_map_normal) {
+    if (sd->type & PRIMITIVE_ALL_TRIANGLE) {
+      float3 f = triangle_smooth_normal_unnormalized(kg, sd, sd->Ng, sd->prim, sd->u, sd->v);
+      return set_attribute_float3(f, type, derivatives, val);
+    }
+    else {
+      return false;
+    }
+  }
+  else {
     return false;
+  }
 }
 
 bool OSLRenderServices::get_background_attribute(const KernelGlobals *kg,

intern/cycles/kernel/osl/osl_services.h

@@ -297,6 +297,7 @@ class OSLRenderServices : public OSL::RendererServices {
   static ustring u_curve_length;
   static ustring u_curve_tangent_normal;
   static ustring u_curve_random;
+  static ustring u_normal_map_normal;
   static ustring u_path_ray_length;
   static ustring u_path_ray_depth;
   static ustring u_path_diffuse_depth;

intern/cycles/kernel/shaders/node_normal_map.osl

@@ -45,7 +45,7 @@ shader node_normal_map(normal NormalIn = N,
 
     // get _unnormalized_ interpolated normal and tangent
     if (getattribute(attr_name, tangent) && getattribute(attr_sign_name, tangent_sign) &&
-        (!is_smooth || getattribute("geom:N", ninterp))) {
+        (!is_smooth || getattribute("geom:normal_map_normal", ninterp))) {
       // apply normal map
       vector B = tangent_sign * cross(ninterp, tangent);
       Normal = normalize(mcolor[0] * tangent + mcolor[1] * B + mcolor[2] * ninterp);

intern/cycles/kernel/svm/svm_bevel.h

@@ -182,17 +182,17 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
     float3 disk_P = (disk_r * cosf(phi)) * disk_T + (disk_r * sinf(phi)) * disk_B;
 
     /* Create ray. */
-    Ray *ray = &isect.ray;
-    ray->P = sd->P + disk_N * disk_height + disk_P;
-    ray->D = -disk_N;
-    ray->t = 2.0f * disk_height;
-    ray->dP = differential_zero_compact();
-    ray->dD = differential_zero_compact();
-    ray->time = sd->time;
+    Ray ray ccl_optional_struct_init;
+    ray.P = sd->P + disk_N * disk_height + disk_P;
+    ray.D = -disk_N;
+    ray.t = 2.0f * disk_height;
+    ray.dP = differential_zero_compact();
+    ray.dD = differential_zero_compact();
+    ray.time = sd->time;
 
     /* Intersect with the same object. if multiple intersections are found it
      * will use at most LOCAL_MAX_HITS hits, a random subset of all hits. */
-    scene_intersect_local(kg, ray, &isect, sd->object, &lcg_state, LOCAL_MAX_HITS);
+    scene_intersect_local(kg, &ray, &isect, sd->object, &lcg_state, LOCAL_MAX_HITS);
 
     int num_eval_hits = min(isect.num_hits, LOCAL_MAX_HITS);
 
@@ -201,23 +201,20 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
       float3 hit_P;
       if (sd->type & PRIMITIVE_TRIANGLE) {
         hit_P = triangle_refine_local(
-            kg, sd, ray->P, ray->D, ray->t, isect.hits[hit].object, isect.hits[hit].prim);
+            kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim);
       }
 #  ifdef __OBJECT_MOTION__
       else if (sd->type & PRIMITIVE_MOTION_TRIANGLE) {
         float3 verts[3];
-        motion_triangle_vertices(
-            kg, sd->object, kernel_tex_fetch(__prim_index, isect.hits[hit].prim), sd->time, verts);
+        motion_triangle_vertices(kg, sd->object, isect.hits[hit].prim, sd->time, verts);
         hit_P = motion_triangle_refine_local(
-            kg, sd, ray->P, ray->D, ray->t, isect.hits[hit].object, isect.hits[hit].prim, verts);
+            kg, sd, ray.P, ray.D, ray.t, isect.hits[hit].object, isect.hits[hit].prim, verts);
       }
 #  endif /* __OBJECT_MOTION__ */
 
       /* Get geometric normal. */
       float3 hit_Ng = isect.Ng[hit];
-      int object = (isect.hits[hit].object == OBJECT_NONE) ?
-                       kernel_tex_fetch(__prim_object, isect.hits[hit].prim) :
-                       isect.hits[hit].object;
+      int object = isect.hits[hit].object;
       int object_flag = kernel_tex_fetch(__object_flag, object);
       if (object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
         hit_Ng = -hit_Ng;
@@ -225,7 +222,7 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
 
       /* Compute smooth normal. */
       float3 N = hit_Ng;
-      int prim = kernel_tex_fetch(__prim_index, isect.hits[hit].prim);
+      int prim = isect.hits[hit].prim;
       int shader = kernel_tex_fetch(__tri_shader, prim);
 
       if (shader & SHADER_SMOOTH_NORMAL) {

intern/cycles/kernel/svm/svm_closure.h

@@ -85,6 +85,9 @@ ccl_device_noinline int svm_node_closure_bsdf(
   }
 
   float3 N = stack_valid(data_node.x) ? stack_load_float3(stack, data_node.x) : sd->N;
+  if (!(sd->type & PRIMITIVE_ALL_CURVE)) {
+    N = ensure_valid_reflection(sd->Ng, sd->I, N);
+  }
 
   float param1 = (stack_valid(param1_offset)) ? stack_load_float(stack, param1_offset) :
                                                 __uint_as_float(node.z);
@@ -166,6 +169,9 @@ ccl_device_noinline int svm_node_closure_bsdf(
       float3 clearcoat_normal = stack_valid(data_cn_ssr.x) ?
                                     stack_load_float3(stack, data_cn_ssr.x) :
                                     sd->N;
+      if (!(sd->type & PRIMITIVE_ALL_CURVE)) {
+        clearcoat_normal = ensure_valid_reflection(sd->Ng, sd->I, clearcoat_normal);
+      }
       float3 subsurface_radius = stack_valid(data_cn_ssr.y) ?
                                      stack_load_float3(stack, data_cn_ssr.y) :
                                      make_float3(1.0f, 1.0f, 1.0f);

intern/cycles/kernel/svm/svm_math_util.h

@@ -200,43 +200,43 @@ ccl_device float svm_math(NodeMathType type, float a, float b, float c)
   }
 }
 
-/* Calculate color in range 800..12000 using an approximation
- * a/x+bx+c for R and G and ((at + b)t + c)t + d) for B
- * Max absolute error for RGB is (0.00095, 0.00077, 0.00057),
- * which is enough to get the same 8 bit/channel color.
- */
-
-ccl_static_constant float blackbody_table_r[6][3] = {
-    {2.52432244e+03f, -1.06185848e-03f, 3.11067539e+00f},
-    {3.37763626e+03f, -4.34581697e-04f, 1.64843306e+00f},
-    {4.10671449e+03f, -8.61949938e-05f, 6.41423749e-01f},
-    {4.66849800e+03f, 2.85655028e-05f, 1.29075375e-01f},
-    {4.60124770e+03f, 2.89727618e-05f, 1.48001316e-01f},
-    {3.78765709e+03f, 9.36026367e-06f, 3.98995841e-01f},
-};
-
-ccl_static_constant float blackbody_table_g[6][3] = {
-    {-7.50343014e+02f, 3.15679613e-04f, 4.73464526e-01f},
-    {-1.00402363e+03f, 1.29189794e-04f, 9.08181524e-01f},
-    {-1.22075471e+03f, 2.56245413e-05f, 1.20753416e+00f},
-    {-1.42546105e+03f, -4.01730887e-05f, 1.44002695e+00f},
-    {-1.18134453e+03f, -2.18913373e-05f, 1.30656109e+00f},
-    {-5.00279505e+02f, -4.59745390e-06f, 1.09090465e+00f},
-};
-
-ccl_static_constant float blackbody_table_b[6][4] = {
-    {0.0f, 0.0f, 0.0f, 0.0f}, /* zeros should be optimized by compiler */
-    {0.0f, 0.0f, 0.0f, 0.0f},
-    {0.0f, 0.0f, 0.0f, 0.0f},
-    {-2.02524603e-11f, 1.79435860e-07f, -2.60561875e-04f, -1.41761141e-02f},
-    {-2.22463426e-13f, -1.55078698e-08f, 3.81675160e-04f, -7.30646033e-01f},
-    {6.72595954e-13f, -2.73059993e-08f, 4.24068546e-04f, -7.52204323e-01f},
-};
-
 ccl_device float3 svm_math_blackbody_color(float t)
 {
   /* TODO(lukas): Reimplement in XYZ. */
 
+  /* Calculate color in range 800..12000 using an approximation
+   * a/x+bx+c for R and G and ((at + b)t + c)t + d) for B
+   * Max absolute error for RGB is (0.00095, 0.00077, 0.00057),
+   * which is enough to get the same 8 bit/channel color.
+   */
+
+  const float blackbody_table_r[6][3] = {
+      {2.52432244e+03f, -1.06185848e-03f, 3.11067539e+00f},
+      {3.37763626e+03f, -4.34581697e-04f, 1.64843306e+00f},
+      {4.10671449e+03f, -8.61949938e-05f, 6.41423749e-01f},
+      {4.66849800e+03f, 2.85655028e-05f, 1.29075375e-01f},
+      {4.60124770e+03f, 2.89727618e-05f, 1.48001316e-01f},
+      {3.78765709e+03f, 9.36026367e-06f, 3.98995841e-01f},
+  };
+
+  const float blackbody_table_g[6][3] = {
+      {-7.50343014e+02f, 3.15679613e-04f, 4.73464526e-01f},
+      {-1.00402363e+03f, 1.29189794e-04f, 9.08181524e-01f},
+      {-1.22075471e+03f, 2.56245413e-05f, 1.20753416e+00f},
+      {-1.42546105e+03f, -4.01730887e-05f, 1.44002695e+00f},
+      {-1.18134453e+03f, -2.18913373e-05f, 1.30656109e+00f},
+      {-5.00279505e+02f, -4.59745390e-06f, 1.09090465e+00f},
+  };
+
+  const float blackbody_table_b[6][4] = {
+      {0.0f, 0.0f, 0.0f, 0.0f}, /* zeros should be optimized by compiler */
+      {0.0f, 0.0f, 0.0f, 0.0f},
+      {0.0f, 0.0f, 0.0f, 0.0f},
+      {-2.02524603e-11f, 1.79435860e-07f, -2.60561875e-04f, -1.41761141e-02f},
+      {-2.22463426e-13f, -1.55078698e-08f, 3.81675160e-04f, -7.30646033e-01f},
+      {6.72595954e-13f, -2.73059993e-08f, 4.24068546e-04f, -7.52204323e-01f},
+  };
+
   if (t >= 12000.0f) {
     return make_float3(0.826270103f, 0.994478524f, 1.56626022f);
   }

intern/cycles/kernel/svm/svm_tex_coord.h

@@ -347,8 +347,6 @@ ccl_device_noinline void svm_node_normal_map(const KernelGlobals *kg,
     N = safe_normalize(sd->N + (N - sd->N) * strength);
   }
 
-  N = ensure_valid_reflection(sd->Ng, sd->I, N);
-
   if (is_zero(N)) {
     N = sd->N;
   }

intern/cycles/kernel/svm/svm_wavelength.h

@@ -34,44 +34,44 @@ CCL_NAMESPACE_BEGIN
 
 /* Wavelength to RGB */
 
-// CIE colour matching functions xBar, yBar, and zBar for
-//   wavelengths from 380 through 780 nanometers, every 5
-//   nanometers.  For a wavelength lambda in this range:
-//        cie_colour_match[(lambda - 380) / 5][0] = xBar
-//        cie_colour_match[(lambda - 380) / 5][1] = yBar
-//        cie_colour_match[(lambda - 380) / 5][2] = zBar
-ccl_static_constant float cie_colour_match[81][3] = {
-    {0.0014f, 0.0000f, 0.0065f}, {0.0022f, 0.0001f, 0.0105f}, {0.0042f, 0.0001f, 0.0201f},
-    {0.0076f, 0.0002f, 0.0362f}, {0.0143f, 0.0004f, 0.0679f}, {0.0232f, 0.0006f, 0.1102f},
-    {0.0435f, 0.0012f, 0.2074f}, {0.0776f, 0.0022f, 0.3713f}, {0.1344f, 0.0040f, 0.6456f},
-    {0.2148f, 0.0073f, 1.0391f}, {0.2839f, 0.0116f, 1.3856f}, {0.3285f, 0.0168f, 1.6230f},
-    {0.3483f, 0.0230f, 1.7471f}, {0.3481f, 0.0298f, 1.7826f}, {0.3362f, 0.0380f, 1.7721f},
-    {0.3187f, 0.0480f, 1.7441f}, {0.2908f, 0.0600f, 1.6692f}, {0.2511f, 0.0739f, 1.5281f},
-    {0.1954f, 0.0910f, 1.2876f}, {0.1421f, 0.1126f, 1.0419f}, {0.0956f, 0.1390f, 0.8130f},
-    {0.0580f, 0.1693f, 0.6162f}, {0.0320f, 0.2080f, 0.4652f}, {0.0147f, 0.2586f, 0.3533f},
-    {0.0049f, 0.3230f, 0.2720f}, {0.0024f, 0.4073f, 0.2123f}, {0.0093f, 0.5030f, 0.1582f},
-    {0.0291f, 0.6082f, 0.1117f}, {0.0633f, 0.7100f, 0.0782f}, {0.1096f, 0.7932f, 0.0573f},
-    {0.1655f, 0.8620f, 0.0422f}, {0.2257f, 0.9149f, 0.0298f}, {0.2904f, 0.9540f, 0.0203f},
-    {0.3597f, 0.9803f, 0.0134f}, {0.4334f, 0.9950f, 0.0087f}, {0.5121f, 1.0000f, 0.0057f},
-    {0.5945f, 0.9950f, 0.0039f}, {0.6784f, 0.9786f, 0.0027f}, {0.7621f, 0.9520f, 0.0021f},
-    {0.8425f, 0.9154f, 0.0018f}, {0.9163f, 0.8700f, 0.0017f}, {0.9786f, 0.8163f, 0.0014f},
-    {1.0263f, 0.7570f, 0.0011f}, {1.0567f, 0.6949f, 0.0010f}, {1.0622f, 0.6310f, 0.0008f},
-    {1.0456f, 0.5668f, 0.0006f}, {1.0026f, 0.5030f, 0.0003f}, {0.9384f, 0.4412f, 0.0002f},
-    {0.8544f, 0.3810f, 0.0002f}, {0.7514f, 0.3210f, 0.0001f}, {0.6424f, 0.2650f, 0.0000f},
-    {0.5419f, 0.2170f, 0.0000f}, {0.4479f, 0.1750f, 0.0000f}, {0.3608f, 0.1382f, 0.0000f},
-    {0.2835f, 0.1070f, 0.0000f}, {0.2187f, 0.0816f, 0.0000f}, {0.1649f, 0.0610f, 0.0000f},
-    {0.1212f, 0.0446f, 0.0000f}, {0.0874f, 0.0320f, 0.0000f}, {0.0636f, 0.0232f, 0.0000f},
-    {0.0468f, 0.0170f, 0.0000f}, {0.0329f, 0.0119f, 0.0000f}, {0.0227f, 0.0082f, 0.0000f},
-    {0.0158f, 0.0057f, 0.0000f}, {0.0114f, 0.0041f, 0.0000f}, {0.0081f, 0.0029f, 0.0000f},
-    {0.0058f, 0.0021f, 0.0000f}, {0.0041f, 0.0015f, 0.0000f}, {0.0029f, 0.0010f, 0.0000f},
-    {0.0020f, 0.0007f, 0.0000f}, {0.0014f, 0.0005f, 0.0000f}, {0.0010f, 0.0004f, 0.0000f},
-    {0.0007f, 0.0002f, 0.0000f}, {0.0005f, 0.0002f, 0.0000f}, {0.0003f, 0.0001f, 0.0000f},
-    {0.0002f, 0.0001f, 0.0000f}, {0.0002f, 0.0001f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f},
-    {0.0001f, 0.0000f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f}, {0.0000f, 0.0000f, 0.0000f}};
-
 ccl_device_noinline void svm_node_wavelength(
     const KernelGlobals *kg, ShaderData *sd, float *stack, uint wavelength, uint color_out)
 {
+  // CIE colour matching functions xBar, yBar, and zBar for
+  //   wavelengths from 380 through 780 nanometers, every 5
+  //   nanometers.  For a wavelength lambda in this range:
+  //        cie_colour_match[(lambda - 380) / 5][0] = xBar
+  //        cie_colour_match[(lambda - 380) / 5][1] = yBar
+  //        cie_colour_match[(lambda - 380) / 5][2] = zBar
+  const float cie_colour_match[81][3] = {
+      {0.0014f, 0.0000f, 0.0065f}, {0.0022f, 0.0001f, 0.0105f}, {0.0042f, 0.0001f, 0.0201f},
+      {0.0076f, 0.0002f, 0.0362f}, {0.0143f, 0.0004f, 0.0679f}, {0.0232f, 0.0006f, 0.1102f},
+      {0.0435f, 0.0012f, 0.2074f}, {0.0776f, 0.0022f, 0.3713f}, {0.1344f, 0.0040f, 0.6456f},
+      {0.2148f, 0.0073f, 1.0391f}, {0.2839f, 0.0116f, 1.3856f}, {0.3285f, 0.0168f, 1.6230f},
+      {0.3483f, 0.0230f, 1.7471f}, {0.3481f, 0.0298f, 1.7826f}, {0.3362f, 0.0380f, 1.7721f},
+      {0.3187f, 0.0480f, 1.7441f}, {0.2908f, 0.0600f, 1.6692f}, {0.2511f, 0.0739f, 1.5281f},
+      {0.1954f, 0.0910f, 1.2876f}, {0.1421f, 0.1126f, 1.0419f}, {0.0956f, 0.1390f, 0.8130f},
+      {0.0580f, 0.1693f, 0.6162f}, {0.0320f, 0.2080f, 0.4652f}, {0.0147f, 0.2586f, 0.3533f},
+      {0.0049f, 0.3230f, 0.2720f}, {0.0024f, 0.4073f, 0.2123f}, {0.0093f, 0.5030f, 0.1582f},
+      {0.0291f, 0.6082f, 0.1117f}, {0.0633f, 0.7100f, 0.0782f}, {0.1096f, 0.7932f, 0.0573f},
+      {0.1655f, 0.8620f, 0.0422f}, {0.2257f, 0.9149f, 0.0298f}, {0.2904f, 0.9540f, 0.0203f},
+      {0.3597f, 0.9803f, 0.0134f}, {0.4334f, 0.9950f, 0.0087f}, {0.5121f, 1.0000f, 0.0057f},
+      {0.5945f, 0.9950f, 0.0039f}, {0.6784f, 0.9786f, 0.0027f}, {0.7621f, 0.9520f, 0.0021f},
+      {0.8425f, 0.9154f, 0.0018f}, {0.9163f, 0.8700f, 0.0017f}, {0.9786f, 0.8163f, 0.0014f},
+      {1.0263f, 0.7570f, 0.0011f}, {1.0567f, 0.6949f, 0.0010f}, {1.0622f, 0.6310f, 0.0008f},
+      {1.0456f, 0.5668f, 0.0006f}, {1.0026f, 0.5030f, 0.0003f}, {0.9384f, 0.4412f, 0.0002f},
+      {0.8544f, 0.3810f, 0.0002f}, {0.7514f, 0.3210f, 0.0001f}, {0.6424f, 0.2650f, 0.0000f},
+      {0.5419f, 0.2170f, 0.0000f}, {0.4479f, 0.1750f, 0.0000f}, {0.3608f, 0.1382f, 0.0000f},
+      {0.2835f, 0.1070f, 0.0000f}, {0.2187f, 0.0816f, 0.0000f}, {0.1649f, 0.0610f, 0.0000f},
+      {0.1212f, 0.0446f, 0.0000f}, {0.0874f, 0.0320f, 0.0000f}, {0.0636f, 0.0232f, 0.0000f},
+      {0.0468f, 0.0170f, 0.0000f}, {0.0329f, 0.0119f, 0.0000f}, {0.0227f, 0.0082f, 0.0000f},
+      {0.0158f, 0.0057f, 0.0000f}, {0.0114f, 0.0041f, 0.0000f}, {0.0081f, 0.0029f, 0.0000f},
+      {0.0058f, 0.0021f, 0.0000f}, {0.0041f, 0.0015f, 0.0000f}, {0.0029f, 0.0010f, 0.0000f},
+      {0.0020f, 0.0007f, 0.0000f}, {0.0014f, 0.0005f, 0.0000f}, {0.0010f, 0.0004f, 0.0000f},
+      {0.0007f, 0.0002f, 0.0000f}, {0.0005f, 0.0002f, 0.0000f}, {0.0003f, 0.0001f, 0.0000f},
+      {0.0002f, 0.0001f, 0.0000f}, {0.0002f, 0.0001f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f},
+      {0.0001f, 0.0000f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f}, {0.0000f, 0.0000f, 0.0000f}};
+
   float lambda_nm = stack_load_float(stack, wavelength);
   float ii = (lambda_nm - 380.0f) * (1.0f / 5.0f);  // scaled 0..80
   int i = float_to_int(ii);

intern/cycles/render/buffers.cpp

@@ -97,6 +97,11 @@ NODE_DEFINE(BufferParams)
   SOCKET_INT(width, "Width", 0);
   SOCKET_INT(height, "Height", 0);
 
+  SOCKET_INT(window_x, "Window X", 0);
+  SOCKET_INT(window_y, "Window Y", 0);
+  SOCKET_INT(window_width, "Window Width", 0);
+  SOCKET_INT(window_height, "Window Height", 0);
+
   SOCKET_INT(full_x, "Full X", 0);
   SOCKET_INT(full_y, "Full Y", 0);
   SOCKET_INT(full_width, "Full Width", 0);
@@ -233,13 +238,31 @@ void BufferParams::update_offset_stride()
 
 bool BufferParams::modified(const BufferParams &other) const
 {
-  if (!(width == other.width && height == other.height && full_x == other.full_x &&
-        full_y == other.full_y && full_width == other.full_width &&
-        full_height == other.full_height && offset == other.offset && stride == other.stride &&
-        pass_stride == other.pass_stride && layer == other.layer && view == other.view &&
-        exposure == other.exposure &&
-        use_approximate_shadow_catcher == other.use_approximate_shadow_catcher &&
-        use_transparent_background == other.use_transparent_background)) {
+  if (width != other.width || height != other.height) {
+    return true;
+  }
+
+  if (full_x != other.full_x || full_y != other.full_y || full_width != other.full_width ||
+      full_height != other.full_height) {
+    return true;
+  }
+
+  if (window_x != other.window_x || window_y != other.window_y ||
+      window_width != other.window_width || window_height != other.window_height) {
+    return true;
+  }
+
+  if (offset != other.offset || stride != other.stride || pass_stride != other.pass_stride) {
+    return true;
+  }
+
+  if (layer != other.layer || view != other.view) {
+    return false;
+  }
+
+  if (exposure != other.exposure ||
+      use_approximate_shadow_catcher != other.use_approximate_shadow_catcher ||
+      use_transparent_background != other.use_transparent_background) {
     return true;
   }
 

intern/cycles/render/buffers.h

@@ -82,6 +82,15 @@ class BufferParams : public Node {
   int width = 0;
   int height = 0;
 
+  /* Windows defines which part of the buffers is visible. The part outside of the window is
+   * considered an "overscan".
+   *
+   * Window X and Y are relative to the position of the buffer in the full buffer. */
+  int window_x = 0;
+  int window_y = 0;
+  int window_width = 0;
+  int window_height = 0;
+
   /* Offset into and width/height of the full buffer. */
   int full_x = 0;
   int full_y = 0;

intern/cycles/render/film.cpp

@@ -326,8 +326,6 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
         kfilm->pass_bake_differential = kfilm->pass_stride;
         break;
 
-      case PASS_RENDER_TIME:
-        break;
       case PASS_CRYPTOMATTE:
         kfilm->pass_cryptomatte = have_cryptomatte ?
                                       min(kfilm->pass_cryptomatte, kfilm->pass_stride) :

intern/cycles/render/geometry.cpp

@@ -46,12 +46,6 @@ CCL_NAMESPACE_BEGIN
 
 /* Geometry */
 
-PackFlags operator|=(PackFlags &pack_flags, uint32_t value)
-{
-  pack_flags = (PackFlags)((uint32_t)pack_flags | value);
-  return pack_flags;
-}
-
 NODE_ABSTRACT_DEFINE(Geometry)
 {
   NodeType *type = NodeType::add("geometry_base", NULL);
@@ -79,7 +73,6 @@ Geometry::Geometry(const NodeType *node_type, const Type type)
 
   bvh = NULL;
   attr_map_offset = 0;
-  optix_prim_offset = 0;
   prim_offset = 0;
 }
 
@@ -707,9 +700,9 @@ void GeometryManager::update_attribute_element_offset(Geometry *geom,
       if (element == ATTR_ELEMENT_CURVE)
         offset -= hair->prim_offset;
       else if (element == ATTR_ELEMENT_CURVE_KEY)
-        offset -= hair->curvekey_offset;
+        offset -= hair->curve_key_offset;
       else if (element == ATTR_ELEMENT_CURVE_KEY_MOTION)
-        offset -= hair->curvekey_offset;
+        offset -= hair->curve_key_offset;
     }
   }
   else {
@@ -972,28 +965,22 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout)
   size_t vert_size = 0;
   size_t tri_size = 0;
 
-  size_t curve_key_size = 0;
   size_t curve_size = 0;
+  size_t curve_key_size = 0;
+  size_t curve_segment_size = 0;
 
   size_t patch_size = 0;
   size_t face_size = 0;
   size_t corner_size = 0;
 
-  size_t optix_prim_size = 0;
-
   foreach (Geometry *geom, scene->geometry) {
-    if (geom->optix_prim_offset != optix_prim_size) {
-      /* Need to rebuild BVH in OptiX, since refit only allows modified mesh data there */
-      const bool has_optix_bvh = bvh_layout == BVH_LAYOUT_OPTIX ||
-                                 bvh_layout == BVH_LAYOUT_MULTI_OPTIX ||
-                                 bvh_layout == BVH_LAYOUT_MULTI_OPTIX_EMBREE;
-      geom->need_update_rebuild |= has_optix_bvh;
-      geom->need_update_bvh_for_offset = true;
-    }
+    bool prim_offset_changed = false;
 
     if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
       Mesh *mesh = static_cast<Mesh *>(geom);
 
+      prim_offset_changed = (mesh->prim_offset != tri_size);
+
       mesh->vert_offset = vert_size;
       mesh->prim_offset = tri_size;
 
@@ -1017,27 +1004,35 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout)
 
       face_size += mesh->get_num_subd_faces();
       corner_size += mesh->subd_face_corners.size();
-
-      mesh->optix_prim_offset = optix_prim_size;
-      optix_prim_size += mesh->num_triangles();
     }
     else if (geom->is_hair()) {
       Hair *hair = static_cast<Hair *>(geom);
 
-      hair->curvekey_offset = curve_key_size;
+      prim_offset_changed = (hair->curve_segment_offset != curve_segment_size);
+      hair->curve_key_offset = curve_key_size;
+      hair->curve_segment_offset = curve_segment_size;
       hair->prim_offset = curve_size;
 
-      curve_key_size += hair->get_curve_keys().size();
       curve_size += hair->num_curves();
+      curve_key_size += hair->get_curve_keys().size();
+      curve_segment_size += hair->num_segments();
+    }
 
-      hair->optix_prim_offset = optix_prim_size;
-      optix_prim_size += hair->num_segments();
+    if (prim_offset_changed) {
+      /* Need to rebuild BVH in OptiX, since refit only allows modified mesh data there */
+      const bool has_optix_bvh = bvh_layout == BVH_LAYOUT_OPTIX ||
+                                 bvh_layout == BVH_LAYOUT_MULTI_OPTIX ||
+                                 bvh_layout == BVH_LAYOUT_MULTI_OPTIX_EMBREE;
+      geom->need_update_rebuild |= has_optix_bvh;
+      geom->need_update_bvh_for_offset = true;
     }
   }
 }
 
-void GeometryManager::device_update_mesh(
-    Device *, DeviceScene *dscene, Scene *scene, bool for_displacement, Progress &progress)
+void GeometryManager::device_update_mesh(Device *,
+                                         DeviceScene *dscene,
+                                         Scene *scene,
+                                         Progress &progress)
 {
   /* Count. */
   size_t vert_size = 0;
@@ -1045,6 +1040,7 @@ void GeometryManager::device_update_mesh(
 
   size_t curve_key_size = 0;
   size_t curve_size = 0;
+  size_t curve_segment_size = 0;
 
   size_t patch_size = 0;
 
@@ -1071,31 +1067,7 @@ void GeometryManager::device_update_mesh(
 
       curve_key_size += hair->get_curve_keys().size();
       curve_size += hair->num_curves();
-    }
-  }
-
-  /* Create mapping from triangle to primitive triangle array. */
-  vector<uint> tri_prim_index(tri_size);
-  if (for_displacement) {
-    /* For displacement kernels we do some trickery to make them believe
-     * we've got all required data ready. However, that data is different
-     * from final render kernels since we don't have BVH yet, so can't
-     * really use same semantic of arrays.
-     */
-    foreach (Geometry *geom, scene->geometry) {
-      if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
-        Mesh *mesh = static_cast<Mesh *>(geom);
-        for (size_t i = 0; i < mesh->num_triangles(); ++i) {
-          tri_prim_index[i + mesh->prim_offset] = 3 * (i + mesh->prim_offset);
-        }
-      }
-    }
-  }
-  else {
-    for (size_t i = 0; i < dscene->prim_index.size(); ++i) {
-      if ((dscene->prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
-        tri_prim_index[dscene->prim_index[i]] = dscene->prim_tri_index[i];
-      }
+      curve_segment_size += hair->num_segments();
     }
   }
 
@@ -1104,6 +1076,7 @@ void GeometryManager::device_update_mesh(
     /* normals */
     progress.set_status("Updating Mesh", "Computing normals");
 
+    float4 *tri_verts = dscene->tri_verts.alloc(tri_size * 3);
     uint *tri_shader = dscene->tri_shader.alloc(tri_size);
     float4 *vnormal = dscene->tri_vnormal.alloc(vert_size);
     uint4 *tri_vindex = dscene->tri_vindex.alloc(tri_size);
@@ -1129,13 +1102,12 @@ void GeometryManager::device_update_mesh(
           mesh->pack_normals(&vnormal[mesh->vert_offset]);
         }
 
-        if (mesh->triangles_is_modified() || mesh->vert_patch_uv_is_modified() || copy_all_data) {
-          mesh->pack_verts(tri_prim_index,
+        if (mesh->verts_is_modified() || mesh->triangles_is_modified() ||
+            mesh->vert_patch_uv_is_modified() || copy_all_data) {
+          mesh->pack_verts(&tri_verts[mesh->prim_offset * 3],
                            &tri_vindex[mesh->prim_offset],
                            &tri_patch[mesh->prim_offset],
-                           &tri_patch_uv[mesh->vert_offset],
-                           mesh->vert_offset,
-                           mesh->prim_offset);
+                           &tri_patch_uv[mesh->vert_offset]);
         }
 
         if (progress.get_cancel())
@@ -1146,6 +1118,7 @@ void GeometryManager::device_update_mesh(
     /* vertex coordinates */
     progress.set_status("Updating Mesh", "Copying Mesh to device");
 
+    dscene->tri_verts.copy_to_device_if_modified();
     dscene->tri_shader.copy_to_device_if_modified();
     dscene->tri_vnormal.copy_to_device_if_modified();
     dscene->tri_vindex.copy_to_device_if_modified();
@@ -1153,13 +1126,16 @@ void GeometryManager::device_update_mesh(
     dscene->tri_patch_uv.copy_to_device_if_modified();
   }
 
-  if (curve_size != 0) {
-    progress.set_status("Updating Mesh", "Copying Strands to device");
+  if (curve_segment_size != 0) {
+    progress.set_status("Updating Mesh", "Copying Curves to device");
 
     float4 *curve_keys = dscene->curve_keys.alloc(curve_key_size);
-    float4 *curves = dscene->curves.alloc(curve_size);
+    KernelCurve *curves = dscene->curves.alloc(curve_size);
+    KernelCurveSegment *curve_segments = dscene->curve_segments.alloc(curve_segment_size);
 
-    const bool copy_all_data = dscene->curve_keys.need_realloc() || dscene->curves.need_realloc();
+    const bool copy_all_data = dscene->curve_keys.need_realloc() ||
+                               dscene->curves.need_realloc() ||
+                               dscene->curve_segments.need_realloc();
 
     foreach (Geometry *geom, scene->geometry) {
       if (geom->is_hair()) {
@@ -1175,9 +1151,9 @@ void GeometryManager::device_update_mesh(
         }
 
         hair->pack_curves(scene,
-                          &curve_keys[hair->curvekey_offset],
+                          &curve_keys[hair->curve_key_offset],
                           &curves[hair->prim_offset],
-                          hair->curvekey_offset);
+                          &curve_segments[hair->curve_segment_offset]);
         if (progress.get_cancel())
           return;
       }
@@ -1185,6 +1161,7 @@ void GeometryManager::device_update_mesh(
 
     dscene->curve_keys.copy_to_device_if_modified();
     dscene->curves.copy_to_device_if_modified();
+    dscene->curve_segments.copy_to_device_if_modified();
   }
 
   if (patch_size != 0 && dscene->patches.need_realloc()) {
@@ -1195,10 +1172,7 @@ void GeometryManager::device_update_mesh(
     foreach (Geometry *geom, scene->geometry) {
       if (geom->is_mesh()) {
         Mesh *mesh = static_cast<Mesh *>(geom);
-        mesh->pack_patches(&patch_data[mesh->patch_offset],
-                           mesh->vert_offset,
-                           mesh->face_offset,
-                           mesh->corner_offset);
+        mesh->pack_patches(&patch_data[mesh->patch_offset]);
 
         if (mesh->patch_table) {
           mesh->patch_table->copy_adjusting_offsets(&patch_data[mesh->patch_table_offset],
@@ -1212,23 +1186,6 @@ void GeometryManager::device_update_mesh(
 
     dscene->patches.copy_to_device();
   }
-
-  if (for_displacement) {
-    float4 *prim_tri_verts = dscene->prim_tri_verts.alloc(tri_size * 3);
-    foreach (Geometry *geom, scene->geometry) {
-      if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
-        Mesh *mesh = static_cast<Mesh *>(geom);
-        for (size_t i = 0; i < mesh->num_triangles(); ++i) {
-          Mesh::Triangle t = mesh->get_triangle(i);
-          size_t offset = 3 * (i + mesh->prim_offset);
-          prim_tri_verts[offset + 0] = float3_to_float4(mesh->verts[t.v[0]]);
-          prim_tri_verts[offset + 1] = float3_to_float4(mesh->verts[t.v[1]]);
-          prim_tri_verts[offset + 2] = float3_to_float4(mesh->verts[t.v[2]]);
-        }
-      }
-    }
-    dscene->prim_tri_verts.copy_to_device();
-  }
 }
 
 void GeometryManager::device_update_bvh(Device *device,
@@ -1256,16 +1213,6 @@ void GeometryManager::device_update_bvh(Device *device,
   const bool can_refit = scene->bvh != nullptr &&
                          (bparams.bvh_layout == BVHLayout::BVH_LAYOUT_OPTIX);
 
-  PackFlags pack_flags = PackFlags::PACK_NONE;
-
-  if (scene->bvh == nullptr) {
-    pack_flags |= PackFlags::PACK_ALL;
-  }
-
-  if (dscene->prim_visibility.is_modified()) {
-    pack_flags |= PackFlags::PACK_VISIBILITY;
-  }
-
   BVH *bvh = scene->bvh;
   if (!scene->bvh) {
     bvh = scene->bvh = BVH::create(bparams, scene->geometry, scene->objects, device);
@@ -1284,77 +1231,7 @@ void GeometryManager::device_update_bvh(Device *device,
     pack = std::move(static_cast<BVH2 *>(bvh)->pack);
   }
   else {
-    progress.set_status("Updating Scene BVH", "Packing BVH primitives");
-
-    size_t num_prims = 0;
-    size_t num_tri_verts = 0;
-    foreach (Geometry *geom, scene->geometry) {
-      if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
-        Mesh *mesh = static_cast<Mesh *>(geom);
-        num_prims += mesh->num_triangles();
-        num_tri_verts += 3 * mesh->num_triangles();
-      }
-      else if (geom->is_hair()) {
-        Hair *hair = static_cast<Hair *>(geom);
-        num_prims += hair->num_segments();
-      }
-    }
-
     pack.root_index = -1;
-
-    if (pack_flags != PackFlags::PACK_ALL) {
-      /* if we do not need to recreate the BVH, then only the vertices are updated, so we can
-       * safely retake the memory */
-      dscene->prim_tri_verts.give_data(pack.prim_tri_verts);
-
-      if ((pack_flags & PackFlags::PACK_VISIBILITY) != 0) {
-        dscene->prim_visibility.give_data(pack.prim_visibility);
-      }
-    }
-    else {
-      /* It is not strictly necessary to skip those resizes we if do not have to repack, as the OS
-       * will not allocate pages if we do not touch them, however it does help catching bugs. */
-      pack.prim_tri_index.resize(num_prims);
-      pack.prim_tri_verts.resize(num_tri_verts);
-      pack.prim_type.resize(num_prims);
-      pack.prim_index.resize(num_prims);
-      pack.prim_object.resize(num_prims);
-      pack.prim_visibility.resize(num_prims);
-    }
-
-    // Merge visibility flags of all objects and find object index for non-instanced geometry
-    unordered_map<const Geometry *, pair<int, uint>> geometry_to_object_info;
-    geometry_to_object_info.reserve(scene->geometry.size());
-    foreach (Object *ob, scene->objects) {
-      const Geometry *const geom = ob->get_geometry();
-      pair<int, uint> &info = geometry_to_object_info[geom];
-      info.second |= ob->visibility_for_tracing();
-      if (!geom->is_instanced()) {
-        info.first = ob->get_device_index();
-      }
-    }
-
-    TaskPool pool;
-    // Iterate over scene mesh list instead of objects, since 'optix_prim_offset' was calculated
-    // based on that list, which may be ordered differently from the object list.
-    foreach (Geometry *geom, scene->geometry) {
-      /* Make a copy of the pack_flags so the current geometry's flags do not pollute the others'.
-       */
-      PackFlags geom_pack_flags = pack_flags;
-
-      if (geom->is_modified()) {
-        geom_pack_flags |= PackFlags::PACK_VERTICES;
-      }
-
-      if (geom_pack_flags == PACK_NONE) {
-        continue;
-      }
-
-      const pair<int, uint> &info = geometry_to_object_info[geom];
-      pool.push(function_bind(
-          &Geometry::pack_primitives, geom, &pack, info.first, info.second, geom_pack_flags));
-    }
-    pool.wait_work();
   }
 
   /* copy to device */
@@ -1375,31 +1252,23 @@ void GeometryManager::device_update_bvh(Device *device,
     dscene->object_node.steal_data(pack.object_node);
     dscene->object_node.copy_to_device();
   }
-  if (pack.prim_tri_index.size() && (dscene->prim_tri_index.need_realloc() || has_bvh2_layout)) {
-    dscene->prim_tri_index.steal_data(pack.prim_tri_index);
-    dscene->prim_tri_index.copy_to_device();
-  }
-  if (pack.prim_tri_verts.size()) {
-    dscene->prim_tri_verts.steal_data(pack.prim_tri_verts);
-    dscene->prim_tri_verts.copy_to_device();
-  }
-  if (pack.prim_type.size() && (dscene->prim_type.need_realloc() || has_bvh2_layout)) {
+  if (pack.prim_type.size()) {
     dscene->prim_type.steal_data(pack.prim_type);
     dscene->prim_type.copy_to_device();
   }
-  if (pack.prim_visibility.size() && (dscene->prim_visibility.is_modified() || has_bvh2_layout)) {
+  if (pack.prim_visibility.size()) {
     dscene->prim_visibility.steal_data(pack.prim_visibility);
     dscene->prim_visibility.copy_to_device();
   }
-  if (pack.prim_index.size() && (dscene->prim_index.need_realloc() || has_bvh2_layout)) {
+  if (pack.prim_index.size()) {
     dscene->prim_index.steal_data(pack.prim_index);
     dscene->prim_index.copy_to_device();
   }
-  if (pack.prim_object.size() && (dscene->prim_object.need_realloc() || has_bvh2_layout)) {
+  if (pack.prim_object.size()) {
     dscene->prim_object.steal_data(pack.prim_object);
     dscene->prim_object.copy_to_device();
   }
-  if (pack.prim_time.size() && (dscene->prim_time.need_realloc() || has_bvh2_layout)) {
+  if (pack.prim_time.size()) {
     dscene->prim_time.steal_data(pack.prim_time);
     dscene->prim_time.copy_to_device();
   }
@@ -1629,8 +1498,6 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
     dscene->bvh_nodes.tag_realloc();
     dscene->bvh_leaf_nodes.tag_realloc();
     dscene->object_node.tag_realloc();
-    dscene->prim_tri_verts.tag_realloc();
-    dscene->prim_tri_index.tag_realloc();
     dscene->prim_type.tag_realloc();
     dscene->prim_visibility.tag_realloc();
     dscene->prim_index.tag_realloc();
@@ -1649,6 +1516,7 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
     if (device_update_flags & DEVICE_CURVE_DATA_NEEDS_REALLOC) {
       dscene->curves.tag_realloc();
       dscene->curve_keys.tag_realloc();
+      dscene->curve_segments.tag_realloc();
     }
   }
 
@@ -1691,6 +1559,7 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
   if (device_update_flags & DEVICE_MESH_DATA_MODIFIED) {
     /* if anything else than vertices or shaders are modified, we would need to reallocate, so
      * these are the only arrays that can be updated */
+    dscene->tri_verts.tag_modified();
     dscene->tri_vnormal.tag_modified();
     dscene->tri_shader.tag_modified();
   }
@@ -1698,6 +1567,7 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
   if (device_update_flags & DEVICE_CURVE_DATA_MODIFIED) {
     dscene->curve_keys.tag_modified();
     dscene->curves.tag_modified();
+    dscene->curve_segments.tag_modified();
   }
 
   need_flags_update = false;
@@ -1906,7 +1776,7 @@ void GeometryManager::device_update(Device *device,
             {"device_update (displacement: copy meshes to device)", time});
       }
     });
-    device_update_mesh(device, dscene, scene, true, progress);
+    device_update_mesh(device, dscene, scene, progress);
   }
   if (progress.get_cancel()) {
     return;
@@ -2058,7 +1928,7 @@ void GeometryManager::device_update(Device *device,
             {"device_update (copy meshes to device)", time});
       }
     });
-    device_update_mesh(device, dscene, scene, false, progress);
+    device_update_mesh(device, dscene, scene, progress);
     if (progress.get_cancel()) {
       return;
     }
@@ -2091,13 +1961,12 @@ void GeometryManager::device_update(Device *device,
   dscene->bvh_nodes.clear_modified();
   dscene->bvh_leaf_nodes.clear_modified();
   dscene->object_node.clear_modified();
-  dscene->prim_tri_verts.clear_modified();
-  dscene->prim_tri_index.clear_modified();
   dscene->prim_type.clear_modified();
   dscene->prim_visibility.clear_modified();
   dscene->prim_index.clear_modified();
   dscene->prim_object.clear_modified();
   dscene->prim_time.clear_modified();
+  dscene->tri_verts.clear_modified();
   dscene->tri_shader.clear_modified();
   dscene->tri_vindex.clear_modified();
   dscene->tri_patch.clear_modified();
@@ -2105,6 +1974,7 @@ void GeometryManager::device_update(Device *device,
   dscene->tri_patch_uv.clear_modified();
   dscene->curves.clear_modified();
   dscene->curve_keys.clear_modified();
+  dscene->curve_segments.clear_modified();
   dscene->patches.clear_modified();
   dscene->attributes_map.clear_modified();
   dscene->attributes_float.clear_modified();
@@ -2118,13 +1988,12 @@ void GeometryManager::device_free(Device *device, DeviceScene *dscene, bool forc
   dscene->bvh_nodes.free_if_need_realloc(force_free);
   dscene->bvh_leaf_nodes.free_if_need_realloc(force_free);
   dscene->object_node.free_if_need_realloc(force_free);
-  dscene->prim_tri_verts.free_if_need_realloc(force_free);
-  dscene->prim_tri_index.free_if_need_realloc(force_free);
   dscene->prim_type.free_if_need_realloc(force_free);
   dscene->prim_visibility.free_if_need_realloc(force_free);
   dscene->prim_index.free_if_need_realloc(force_free);
   dscene->prim_object.free_if_need_realloc(force_free);
   dscene->prim_time.free_if_need_realloc(force_free);
+  dscene->tri_verts.free_if_need_realloc(force_free);
   dscene->tri_shader.free_if_need_realloc(force_free);
   dscene->tri_vnormal.free_if_need_realloc(force_free);
   dscene->tri_vindex.free_if_need_realloc(force_free);
@@ -2132,6 +2001,7 @@ void GeometryManager::device_free(Device *device, DeviceScene *dscene, bool forc
   dscene->tri_patch_uv.free_if_need_realloc(force_free);
   dscene->curves.free_if_need_realloc(force_free);
   dscene->curve_keys.free_if_need_realloc(force_free);
+  dscene->curve_segments.free_if_need_realloc(force_free);
   dscene->patches.free_if_need_realloc(force_free);
   dscene->attributes_map.free_if_need_realloc(force_free);
   dscene->attributes_float.free_if_need_realloc(force_free);

intern/cycles/render/geometry.h

@@ -43,24 +43,6 @@ class Shader;
 class Volume;
 struct PackedBVH;
 
-/* Flags used to determine which geometry data need to be packed. */
-enum PackFlags : uint32_t {
-  PACK_NONE = 0u,
-
-  /* Pack the geometry information (e.g. triangle or curve keys indices). */
-  PACK_GEOMETRY = (1u << 0),
-
-  /* Pack the vertices, for Meshes and Volumes' bounding meshes. */
-  PACK_VERTICES = (1u << 1),
-
-  /* Pack the visibility flags for each triangle or curve. */
-  PACK_VISIBILITY = (1u << 2),
-
-  PACK_ALL = (PACK_GEOMETRY | PACK_VERTICES | PACK_VISIBILITY),
-};
-
-PackFlags operator|=(PackFlags &pack_flags, uint32_t value);
-
 /* Geometry
  *
  * Base class for geometric types like Mesh and Hair. */
@@ -100,7 +82,6 @@ class Geometry : public Node {
   BVH *bvh;
   size_t attr_map_offset;
   size_t prim_offset;
-  size_t optix_prim_offset;
 
   /* Shader Properties */
   bool has_volume;         /* Set in the device_update_flags(). */
@@ -144,10 +125,7 @@ class Geometry : public Node {
                    int n,
                    int total);
 
-  virtual void pack_primitives(PackedBVH *pack,
-                               int object,
-                               uint visibility,
-                               PackFlags pack_flags) = 0;
+  virtual PrimitiveType primitive_type() const = 0;
 
   /* Check whether the geometry should have own BVH built separately. Briefly,
    * own BVH is needed for geometry, if:
@@ -260,11 +238,7 @@ class GeometryManager {
 
   void device_update_object(Device *device, DeviceScene *dscene, Scene *scene, Progress &progress);
 
-  void device_update_mesh(Device *device,
-                          DeviceScene *dscene,
-                          Scene *scene,
-                          bool for_displacement,
-                          Progress &progress);
+  void device_update_mesh(Device *device, DeviceScene *dscene, Scene *scene, Progress &progress);
 
   void device_update_attributes(Device *device,
                                 DeviceScene *dscene,

intern/cycles/render/graph.cpp

@@ -1149,7 +1149,9 @@ int ShaderGraph::get_num_closures()
       num_closures += 8;
     }
     else if (CLOSURE_IS_VOLUME(closure_type)) {
-      num_closures += VOLUME_STACK_SIZE;
+      /* TODO(sergey): Verify this is still needed, since we have special minimized volume storage
+       * for the volume steps. */
+      num_closures += MAX_VOLUME_STACK_SIZE;
     }
     else if (closure_type == CLOSURE_BSDF_HAIR_PRINCIPLED_ID) {
       num_closures += 4;

intern/cycles/render/hair.cpp

@@ -295,7 +295,8 @@ NODE_DEFINE(Hair)
 
 Hair::Hair() : Geometry(get_node_type(), Geometry::HAIR)
 {
-  curvekey_offset = 0;
+  curve_key_offset = 0;
+  curve_segment_offset = 0;
   curve_shape = CURVE_RIBBON;
 }
 
@@ -462,8 +463,8 @@ void Hair::apply_transform(const Transform &tfm, const bool apply_to_motion)
 
 void Hair::pack_curves(Scene *scene,
                        float4 *curve_key_co,
-                       float4 *curve_data,
-                       size_t curvekey_offset)
+                       KernelCurve *curves,
+                       KernelCurveSegment *curve_segments)
 {
   size_t curve_keys_size = curve_keys.size();
 
@@ -477,7 +478,10 @@ void Hair::pack_curves(Scene *scene,
   }
 
   /* pack curve segments */
+  const PrimitiveType type = primitive_type();
+
   size_t curve_num = num_curves();
+  size_t index = 0;
 
   for (size_t i = 0; i < curve_num; i++) {
     Curve curve = get_curve(i);
@@ -487,56 +491,24 @@ void Hair::pack_curves(Scene *scene,
                          scene->default_surface;
     shader_id = scene->shader_manager->get_shader_id(shader, false);
 
-    curve_data[i] = make_float4(__int_as_float(curve.first_key + curvekey_offset),
-                                __int_as_float(curve.num_keys),
-                                __int_as_float(shader_id),
-                                0.0f);
+    curves[i].shader_id = shader_id;
+    curves[i].first_key = curve_key_offset + curve.first_key;
+    curves[i].num_keys = curve.num_keys;
+    curves[i].type = type;
+
+    for (int k = 0; k < curve.num_segments(); ++k, ++index) {
+      curve_segments[index].prim = prim_offset + i;
+      curve_segments[index].type = PRIMITIVE_PACK_SEGMENT(type, k);
+    }
   }
 }
 
-void Hair::pack_primitives(PackedBVH *pack, int object, uint visibility, PackFlags pack_flags)
+PrimitiveType Hair::primitive_type() const
 {
-  if (curve_first_key.empty())
-    return;
-
-  /* Separate loop as other arrays are not initialized if their packing is not required. */
-  if ((pack_flags & PACK_VISIBILITY) != 0) {
-    unsigned int *prim_visibility = &pack->prim_visibility[optix_prim_offset];
-
-    size_t index = 0;
-    for (size_t j = 0; j < num_curves(); ++j) {
-      Curve curve = get_curve(j);
-      for (size_t k = 0; k < curve.num_segments(); ++k, ++index) {
-        prim_visibility[index] = visibility;
-      }
-    }
-  }
-
-  if ((pack_flags & PACK_GEOMETRY) != 0) {
-    unsigned int *prim_tri_index = &pack->prim_tri_index[optix_prim_offset];
-    int *prim_type = &pack->prim_type[optix_prim_offset];
-    int *prim_index = &pack->prim_index[optix_prim_offset];
-    int *prim_object = &pack->prim_object[optix_prim_offset];
-    // 'pack->prim_time' is unused by Embree and OptiX
-
-    uint type = has_motion_blur() ?
-                    ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON :
-                                                     PRIMITIVE_MOTION_CURVE_THICK) :
-                    ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON :
-                                                     PRIMITIVE_CURVE_THICK);
-
-    size_t index = 0;
-    for (size_t j = 0; j < num_curves(); ++j) {
-      Curve curve = get_curve(j);
-      for (size_t k = 0; k < curve.num_segments(); ++k, ++index) {
-        prim_tri_index[index] = -1;
-        prim_type[index] = PRIMITIVE_PACK_SEGMENT(type, k);
-        // Each curve segment points back to its curve index
-        prim_index[index] = j + prim_offset;
-        prim_object[index] = object;
-      }
-    }
-  }
+  return has_motion_blur() ?
+             ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON :
+                                              PRIMITIVE_MOTION_CURVE_THICK) :
+             ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON : PRIMITIVE_CURVE_THICK);
 }
 
 CCL_NAMESPACE_END

intern/cycles/render/hair.h

@@ -21,6 +21,8 @@
 
 CCL_NAMESPACE_BEGIN
 
+struct KernelCurveSegment;
+
 class Hair : public Geometry {
  public:
   NODE_DECLARE
@@ -95,7 +97,8 @@ class Hair : public Geometry {
   NODE_SOCKET_API_ARRAY(array<int>, curve_shader)
 
   /* BVH */
-  size_t curvekey_offset;
+  size_t curve_key_offset;
+  size_t curve_segment_offset;
   CurveShapeType curve_shape;
 
   /* Constructor/Destructor */
@@ -144,12 +147,12 @@ class Hair : public Geometry {
   void get_uv_tiles(ustring map, unordered_set<int> &tiles) override;
 
   /* BVH */
-  void pack_curves(Scene *scene, float4 *curve_key_co, float4 *curve_data, size_t curvekey_offset);
+  void pack_curves(Scene *scene,
+                   float4 *curve_key_co,
+                   KernelCurve *curve,
+                   KernelCurveSegment *curve_segments);
 
-  void pack_primitives(PackedBVH *pack,
-                       int object,
-                       uint visibility,
-                       PackFlags pack_flags) override;
+  PrimitiveType primitive_type() const override;
 };
 
 CCL_NAMESPACE_END

intern/cycles/render/mesh.cpp

@@ -729,12 +729,7 @@ void Mesh::pack_normals(float4 *vnormal)
   }
 }
 
-void Mesh::pack_verts(const vector<uint> &tri_prim_index,
-                      uint4 *tri_vindex,
-                      uint *tri_patch,
-                      float2 *tri_patch_uv,
-                      size_t vert_offset,
-                      size_t tri_offset)
+void Mesh::pack_verts(float4 *tri_verts, uint4 *tri_vindex, uint *tri_patch, float2 *tri_patch_uv)
 {
   size_t verts_size = verts.size();
 
@@ -749,17 +744,19 @@ void Mesh::pack_verts(const vector<uint> &tri_prim_index,
   size_t triangles_size = num_triangles();
 
   for (size_t i = 0; i < triangles_size; i++) {
-    Triangle t = get_triangle(i);
-    tri_vindex[i] = make_uint4(t.v[0] + vert_offset,
-                               t.v[1] + vert_offset,
-                               t.v[2] + vert_offset,
-                               tri_prim_index[i + tri_offset]);
+    const Triangle t = get_triangle(i);
+    tri_vindex[i] = make_uint4(
+        t.v[0] + vert_offset, t.v[1] + vert_offset, t.v[2] + vert_offset, 3 * (prim_offset + i));
 
     tri_patch[i] = (!get_num_subd_faces()) ? -1 : (triangle_patch[i] * 8 + patch_offset);
+
+    tri_verts[i * 3] = float3_to_float4(verts[t.v[0]]);
+    tri_verts[i * 3 + 1] = float3_to_float4(verts[t.v[1]]);
+    tri_verts[i * 3 + 2] = float3_to_float4(verts[t.v[2]]);
   }
 }
 
-void Mesh::pack_patches(uint *patch_data, uint vert_offset, uint face_offset, uint corner_offset)
+void Mesh::pack_patches(uint *patch_data)
 {
   size_t num_faces = get_num_subd_faces();
   int ngons = 0;
@@ -805,53 +802,9 @@ void Mesh::pack_patches(uint *patch_data, uint vert_offset, uint face_offset, ui
   }
 }
 
-void Mesh::pack_primitives(ccl::PackedBVH *pack, int object, uint visibility, PackFlags pack_flags)
+PrimitiveType Mesh::primitive_type() const
 {
-  if (triangles.empty())
-    return;
-
-  const size_t num_prims = num_triangles();
-
-  /* Use prim_offset for indexing as it is computed per geometry type, and prim_tri_verts does not
-   * contain data for Hair geometries. */
-  float4 *prim_tri_verts = &pack->prim_tri_verts[prim_offset * 3];
-  // 'pack->prim_time' is unused by Embree and OptiX
-
-  uint type = has_motion_blur() ? PRIMITIVE_MOTION_TRIANGLE : PRIMITIVE_TRIANGLE;
-
-  /* Separate loop as other arrays are not initialized if their packing is not required. */
-  if ((pack_flags & PackFlags::PACK_VISIBILITY) != 0) {
-    unsigned int *prim_visibility = &pack->prim_visibility[optix_prim_offset];
-    for (size_t k = 0; k < num_prims; ++k) {
-      prim_visibility[k] = visibility;
-    }
-  }
-
-  if ((pack_flags & PackFlags::PACK_GEOMETRY) != 0) {
-    /* Use optix_prim_offset for indexing as those arrays also contain data for Hair geometries. */
-    unsigned int *prim_tri_index = &pack->prim_tri_index[optix_prim_offset];
-    int *prim_type = &pack->prim_type[optix_prim_offset];
-    int *prim_index = &pack->prim_index[optix_prim_offset];
-    int *prim_object = &pack->prim_object[optix_prim_offset];
-
-    for (size_t k = 0; k < num_prims; ++k) {
-      if ((pack_flags & PackFlags::PACK_GEOMETRY) != 0) {
-        prim_tri_index[k] = (prim_offset + k) * 3;
-        prim_type[k] = type;
-        prim_index[k] = prim_offset + k;
-        prim_object[k] = object;
-      }
-    }
-  }
-
-  if ((pack_flags & PackFlags::PACK_VERTICES) != 0) {
-    for (size_t k = 0; k < num_prims; ++k) {
-      const Mesh::Triangle t = get_triangle(k);
-      prim_tri_verts[k * 3] = float3_to_float4(verts[t.v[0]]);
-      prim_tri_verts[k * 3 + 1] = float3_to_float4(verts[t.v[1]]);
-      prim_tri_verts[k * 3 + 2] = float3_to_float4(verts[t.v[2]]);
-    }
-  }
+  return has_motion_blur() ? PRIMITIVE_MOTION_TRIANGLE : PRIMITIVE_TRIANGLE;
 }
 
 CCL_NAMESPACE_END

intern/cycles/render/mesh.h

@@ -224,18 +224,10 @@ class Mesh : public Geometry {
 
   void pack_shaders(Scene *scene, uint *shader);
   void pack_normals(float4 *vnormal);
-  void pack_verts(const vector<uint> &tri_prim_index,
-                  uint4 *tri_vindex,
-                  uint *tri_patch,
-                  float2 *tri_patch_uv,
-                  size_t vert_offset,
-                  size_t tri_offset);
-  void pack_patches(uint *patch_data, uint vert_offset, uint face_offset, uint corner_offset);
+  void pack_verts(float4 *tri_verts, uint4 *tri_vindex, uint *tri_patch, float2 *tri_patch_uv);
+  void pack_patches(uint *patch_data);
 
-  void pack_primitives(PackedBVH *pack,
-                       int object,
-                       uint visibility,
-                       PackFlags pack_flags) override;
+  PrimitiveType primitive_type() const override;
 
   void tessellate(DiagSplit *split);
 

intern/cycles/render/object.cpp

@@ -60,6 +60,7 @@ struct UpdateObjectTransformState {
 
   /* Packed object arrays. Those will be filled in. */
   uint *object_flag;
+  uint *object_visibility;
   KernelObject *objects;
   Transform *object_motion_pass;
   DecomposedTransform *object_motion;
@@ -366,6 +367,22 @@ float Object::compute_volume_step_size() const
   return step_size;
 }
 
+bool Object::check_is_volume() const
+{
+  if (geometry->geometry_type == Geometry::VOLUME) {
+    return true;
+  }
+
+  for (Node *node : get_geometry()->get_used_shaders()) {
+    const Shader *shader = static_cast<const Shader *>(node);
+    if (shader->has_volume_connected) {
+      return true;
+    }
+  }
+
+  return false;
+}
+
 int Object::get_device_index() const
 {
   return index;
@@ -512,6 +529,9 @@ void ObjectManager::device_update_object_transform(UpdateObjectTransformState *s
                                              (1.0f - 0.5f * ob->shadow_terminator_shading_offset);
   kobject.shadow_terminator_geometry_offset = ob->shadow_terminator_geometry_offset;
 
+  kobject.visibility = ob->visibility_for_tracing();
+  kobject.primitive_type = geom->primitive_type();
+
   /* Object flag. */
   if (ob->use_holdout) {
     flag |= SD_OBJECT_HOLDOUT_MASK;

intern/cycles/render/object.h

@@ -109,6 +109,13 @@ class Object : public Node {
   /* Compute step size from attributes, shaders, transforms. */
   float compute_volume_step_size() const;
 
+  /* Check whether this object requires volume sampling (and hence might require space in the
+   * volume stack).
+   *
+   * Note that this is a naive iteration over shaders, which allows to access information prior
+   * to `scene_update()`. */
+  bool check_is_volume() const;
+
  protected:
   /* Specifies the position of the object in scene->objects and
    * in the device vectors. Gets set in device_update. */

intern/cycles/render/pass.cpp

@@ -89,7 +89,6 @@ const NodeEnum *Pass::get_type_enum()
     pass_type_enum.insert("material_id", PASS_MATERIAL_ID);
     pass_type_enum.insert("motion", PASS_MOTION);
     pass_type_enum.insert("motion_weight", PASS_MOTION_WEIGHT);
-    pass_type_enum.insert("render_time", PASS_RENDER_TIME);
     pass_type_enum.insert("cryptomatte", PASS_CRYPTOMATTE);
     pass_type_enum.insert("aov_color", PASS_AOV_COLOR);
     pass_type_enum.insert("aov_value", PASS_AOV_VALUE);
@@ -217,10 +216,6 @@ PassInfo Pass::get_info(const PassType type, const bool include_albedo)
       pass_info.num_components = 3;
       pass_info.use_exposure = false;
       break;
-    case PASS_RENDER_TIME:
-      /* This pass is handled entirely on the host side. */
-      pass_info.num_components = 0;
-      break;
 
     case PASS_DIFFUSE_COLOR:
     case PASS_GLOSSY_COLOR:

intern/cycles/render/scene.cpp

@@ -49,13 +49,12 @@ DeviceScene::DeviceScene(Device *device)
     : bvh_nodes(device, "__bvh_nodes", MEM_GLOBAL),
       bvh_leaf_nodes(device, "__bvh_leaf_nodes", MEM_GLOBAL),
       object_node(device, "__object_node", MEM_GLOBAL),
-      prim_tri_index(device, "__prim_tri_index", MEM_GLOBAL),
-      prim_tri_verts(device, "__prim_tri_verts", MEM_GLOBAL),
       prim_type(device, "__prim_type", MEM_GLOBAL),
       prim_visibility(device, "__prim_visibility", MEM_GLOBAL),
       prim_index(device, "__prim_index", MEM_GLOBAL),
       prim_object(device, "__prim_object", MEM_GLOBAL),
       prim_time(device, "__prim_time", MEM_GLOBAL),
+      tri_verts(device, "__tri_verts", MEM_GLOBAL),
       tri_shader(device, "__tri_shader", MEM_GLOBAL),
       tri_vnormal(device, "__tri_vnormal", MEM_GLOBAL),
       tri_vindex(device, "__tri_vindex", MEM_GLOBAL),
@@ -63,6 +62,7 @@ DeviceScene::DeviceScene(Device *device)
       tri_patch_uv(device, "__tri_patch_uv", MEM_GLOBAL),
       curves(device, "__curves", MEM_GLOBAL),
       curve_keys(device, "__curve_keys", MEM_GLOBAL),
+      curve_segments(device, "__curve_segments", MEM_GLOBAL),
       patches(device, "__patches", MEM_GLOBAL),
       objects(device, "__objects", MEM_GLOBAL),
       object_motion_pass(device, "__object_motion_pass", MEM_GLOBAL),
@@ -527,6 +527,8 @@ void Scene::update_kernel_features()
   const uint max_closures = (params.background) ? get_max_closure_count() : MAX_CLOSURE;
   dscene.data.max_closures = max_closures;
   dscene.data.max_shaders = shaders.size();
+
+  dscene.data.volume_stack_size = get_volume_stack_size();
 }
 
 bool Scene::update(Progress &progress)
@@ -642,6 +644,33 @@ int Scene::get_max_closure_count()
   return max_closure_global;
 }
 
+int Scene::get_volume_stack_size() const
+{
+  /* Quick non-expensive check. Can over-estimate maximum possible nested level, but does not
+   * require expensive calculation during pre-processing. */
+  int num_volume_objects = 0;
+  for (const Object *object : objects) {
+    if (object->check_is_volume()) {
+      ++num_volume_objects;
+    }
+
+    if (num_volume_objects == MAX_VOLUME_STACK_SIZE) {
+      break;
+    }
+  }
+
+  /* Count background world for the stack. */
+  const Shader *background_shader = background->get_shader(this);
+  if (background_shader && background_shader->has_volume_connected) {
+    ++num_volume_objects;
+  }
+
+  /* Space for terminator. */
+  ++num_volume_objects;
+
+  return min(num_volume_objects, MAX_VOLUME_STACK_SIZE);
+}
+
 bool Scene::has_shadow_catcher()
 {
   if (shadow_catcher_modified_) {

intern/cycles/render/scene.h

@@ -74,8 +74,6 @@ class DeviceScene {
   device_vector<int4> bvh_nodes;
   device_vector<int4> bvh_leaf_nodes;
   device_vector<int> object_node;
-  device_vector<uint> prim_tri_index;
-  device_vector<float4> prim_tri_verts;
   device_vector<int> prim_type;
   device_vector<uint> prim_visibility;
   device_vector<int> prim_index;
@@ -83,14 +81,16 @@ class DeviceScene {
   device_vector<float2> prim_time;
 
   /* mesh */
+  device_vector<float4> tri_verts;
   device_vector<uint> tri_shader;
   device_vector<float4> tri_vnormal;
   device_vector<uint4> tri_vindex;
   device_vector<uint> tri_patch;
   device_vector<float2> tri_patch_uv;
 
-  device_vector<float4> curves;
+  device_vector<KernelCurve> curves;
   device_vector<float4> curve_keys;
+  device_vector<KernelCurveSegment> curve_segments;
 
   device_vector<uint> patches;
 
@@ -344,6 +344,9 @@ class Scene : public NodeOwner {
   /* Get maximum number of closures to be used in kernel. */
   int get_max_closure_count();
 
+  /* Get size of a volume stack needed to render this scene.  */
+  int get_volume_stack_size() const;
+
   template<typename T> void delete_node_impl(T *node)
   {
     delete node;

intern/cycles/render/session.cpp

@@ -157,6 +157,13 @@ void Session::run_main_render_loop()
       continue;
     }
 
+    /* Stop rendering if error happened during scene update or other step of preparing scene
+     * for render. */
+    if (device->have_error()) {
+      progress.set_error(device->error_message());
+      break;
+    }
+
     {
       /* buffers mutex is locked entirely while rendering each
        * sample, and released/reacquired on each iteration to allow
@@ -172,10 +179,9 @@ void Session::run_main_render_loop()
       /* update status and timing */
       update_status_time();
 
+      /* Stop rendering if error happened during path tracing. */
       if (device->have_error()) {
-        const string &error_message = device->error_message();
-        progress.set_error(error_message);
-        progress.set_cancel(error_message);
+        progress.set_error(device->error_message());
         break;
       }
     }
@@ -280,12 +286,20 @@ RenderWork Session::run_update_for_next_iteration()
       BufferParams tile_params = buffer_params_;
 
       const Tile &tile = tile_manager_.get_current_tile();
+
       tile_params.width = tile.width;
       tile_params.height = tile.height;
+
+      tile_params.window_x = tile.window_x;
+      tile_params.window_y = tile.window_y;
+      tile_params.window_width = tile.window_width;
+      tile_params.window_height = tile.window_height;
+
       tile_params.full_x = tile.x + buffer_params_.full_x;
       tile_params.full_y = tile.y + buffer_params_.full_y;
       tile_params.full_width = buffer_params_.full_width;
       tile_params.full_height = buffer_params_.full_height;
+
       tile_params.update_offset_stride();
 
       path_trace_->reset(buffer_params_, tile_params);

intern/cycles/render/tile.cpp

@@ -372,8 +372,17 @@ void TileManager::update(const BufferParams &params, const Scene *scene)
   configure_image_spec_from_buffer(&write_state_.image_spec, buffer_params_, tile_size_);
 
   const DenoiseParams denoise_params = scene->integrator->get_denoise_params();
+  const AdaptiveSampling adaptive_sampling = scene->integrator->get_adaptive_sampling();
+
   node_to_image_spec_atttributes(
       &write_state_.image_spec, &denoise_params, ATTR_DENOISE_SOCKET_PREFIX);
+
+  if (adaptive_sampling.use) {
+    overscan_ = 4;
+  }
+  else {
+    overscan_ = 0;
+  }
 }
 
 bool TileManager::done()
@@ -399,18 +408,25 @@ Tile TileManager::get_tile_for_index(int index) const
   /* TODO(sergey): Consider using hilbert spiral, or. maybe, even configurable. Not sure this
    * brings a lot of value since this is only applicable to BIG tiles. */
 
-  const int tile_y = index / tile_state_.num_tiles_x;
-  const int tile_x = index - tile_y * tile_state_.num_tiles_x;
+  const int tile_index_y = index / tile_state_.num_tiles_x;
+  const int tile_index_x = index - tile_index_y * tile_state_.num_tiles_x;
+
+  const int tile_window_x = tile_index_x * tile_size_.x;
+  const int tile_window_y = tile_index_y * tile_size_.y;
 
   Tile tile;
 
-  tile.x = tile_x * tile_size_.x;
-  tile.y = tile_y * tile_size_.y;
-  tile.width = tile_size_.x;
-  tile.height = tile_size_.y;
+  tile.x = max(0, tile_window_x - overscan_);
+  tile.y = max(0, tile_window_y - overscan_);
 
-  tile.width = min(tile.width, buffer_params_.width - tile.x);
-  tile.height = min(tile.height, buffer_params_.height - tile.y);
+  tile.window_x = tile_window_x - tile.x;
+  tile.window_y = tile_window_y - tile.y;
+  tile.window_width = min(tile_size_.x, buffer_params_.width - tile_window_x);
+  tile.window_height = min(tile_size_.y, buffer_params_.height - tile_window_y);
+
+  tile.width = min(buffer_params_.width - tile.x, tile.window_x + tile.window_width + overscan_);
+  tile.height = min(buffer_params_.height - tile.y,
+                    tile.window_y + tile.window_height + overscan_);
 
   return tile;
 }
@@ -483,11 +499,22 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
 
   DCHECK_EQ(tile_buffers.params.pass_stride, buffer_params_.pass_stride);
 
+  vector<float> pixel_storage;
+
   const BufferParams &tile_params = tile_buffers.params;
 
-  const float *pixels = tile_buffers.buffer.data();
-  const int tile_x = tile_params.full_x - buffer_params_.full_x;
-  const int tile_y = tile_params.full_y - buffer_params_.full_y;
+  const int tile_x = tile_params.full_x - buffer_params_.full_x + tile_params.window_x;
+  const int tile_y = tile_params.full_y - buffer_params_.full_y + tile_params.window_y;
+
+  const int64_t pass_stride = tile_params.pass_stride;
+  const int64_t tile_row_stride = tile_params.width * pass_stride;
+
+  const int64_t xstride = pass_stride * sizeof(float);
+  const int64_t ystride = xstride * tile_params.width;
+  const int64_t zstride = ystride * tile_params.height;
+
+  const float *pixels = tile_buffers.buffer.data() + tile_params.window_x * pass_stride +
+                        tile_params.window_y * tile_row_stride;
 
   VLOG(3) << "Write tile at " << tile_x << ", " << tile_y;
 
@@ -499,13 +526,16 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
    * The only thing we have to ensure is that the tile_x and tile_y are a multiple of the
    * image tile size, which happens in compute_render_tile_size. */
   if (!write_state_.tile_out->write_tiles(tile_x,
-                                          tile_x + tile_params.width,
+                                          tile_x + tile_params.window_width,
                                           tile_y,
-                                          tile_y + tile_params.height,
+                                          tile_y + tile_params.window_height,
                                           0,
                                           1,
                                           TypeDesc::FLOAT,
-                                          pixels)) {
+                                          pixels,
+                                          xstride,
+                                          ystride,
+                                          zstride)) {
     LOG(ERROR) << "Error writing tile " << write_state_.tile_out->geterror();
     return false;
   }
@@ -531,12 +561,15 @@ void TileManager::finish_write_tiles()
          ++tile_index) {
       const Tile tile = get_tile_for_index(tile_index);
 
-      VLOG(3) << "Write dummy tile at " << tile.x << ", " << tile.y;
+      const int tile_x = tile.x + tile.window_x;
+      const int tile_y = tile.y + tile.window_y;
 
-      write_state_.tile_out->write_tiles(tile.x,
-                                         tile.x + tile.width,
-                                         tile.y,
-                                         tile.y + tile.height,
+      VLOG(3) << "Write dummy tile at " << tile_x << ", " << tile_y;
+
+      write_state_.tile_out->write_tiles(tile_x,
+                                         tile_x + tile.window_width,
+                                         tile_y,
+                                         tile_y + tile.window_height,
                                          0,
                                          1,
                                          TypeDesc::FLOAT,

intern/cycles/render/tile.h

@@ -35,6 +35,9 @@ class Tile {
   int x = 0, y = 0;
   int width = 0, height = 0;
 
+  int window_x = 0, window_y = 0;
+  int window_width = 0, window_height = 0;
+
   Tile()
   {
   }
@@ -78,6 +81,11 @@ class TileManager {
     return tile_state_.num_tiles > 1;
   }
 
+  inline int get_tile_overscan() const
+  {
+    return overscan_;
+  }
+
   bool next();
   bool done();
 
@@ -128,6 +136,9 @@ class TileManager {
 
   int2 tile_size_ = make_int2(0, 0);
 
+  /* Number of extra pixels around the actual tile to render. */
+  int overscan_ = 0;
+
   BufferParams buffer_params_;
 
   /* Tile scheduling state. */

intern/ghost/CMakeLists.txt

@@ -473,7 +473,6 @@ if(WITH_XR_OPENXR)
     intern/GHOST_Xr.cpp
     intern/GHOST_XrAction.cpp
     intern/GHOST_XrContext.cpp
-    intern/GHOST_XrControllerModel.cpp
     intern/GHOST_XrEvent.cpp
     intern/GHOST_XrGraphicsBinding.cpp
     intern/GHOST_XrSession.cpp
@@ -483,19 +482,13 @@ if(WITH_XR_OPENXR)
     intern/GHOST_IXrGraphicsBinding.h
     intern/GHOST_XrAction.h
     intern/GHOST_XrContext.h
-    intern/GHOST_XrControllerModel.h
     intern/GHOST_XrException.h
     intern/GHOST_XrSession.h
     intern/GHOST_XrSwapchain.h
     intern/GHOST_Xr_intern.h
     intern/GHOST_Xr_openxr_includes.h
   )
-  list(APPEND INC
-    ../../extern/json/include
-    ../../extern/tinygltf
-  )
   list(APPEND INC_SYS
-    ${EIGEN3_INCLUDE_DIRS}
     ${XR_OPENXR_SDK_INCLUDE_DIR}
   )
   list(APPEND LIB

intern/ghost/GHOST_C-api.h

@@ -1140,30 +1140,6 @@ void GHOST_XrGetActionCustomdataArray(GHOST_XrContextHandle xr_context,
                                       const char *action_set_name,
                                       void **r_customdata_array);
 
-/* controller model */
-/**
- * Load the OpenXR controller model.
- */
-int GHOST_XrLoadControllerModel(GHOST_XrContextHandle xr_context, const char *subaction_path);
-
-/**
- * Unload the OpenXR controller model.
- */
-void GHOST_XrUnloadControllerModel(GHOST_XrContextHandle xr_context, const char *subaction_path);
-
-/**
- * Update component transforms for the OpenXR controller model.
- */
-int GHOST_XrUpdateControllerModelComponents(GHOST_XrContextHandle xr_context,
-                                            const char *subaction_path);
-
-/**
- * Get vertex data for the OpenXR controller model.
- */
-int GHOST_XrGetControllerModelData(GHOST_XrContextHandle xr_context,
-                                   const char *subaction_path,
-                                   GHOST_XrControllerModelData *r_data);
-
 #endif /* WITH_XR_OPENXR */
 
 #ifdef __cplusplus

intern/ghost/GHOST_Types.h

@@ -643,7 +643,7 @@ typedef void (*GHOST_XrDrawViewFn)(const struct GHOST_XrDrawViewInfo *draw_view,
  */
 typedef const GHOST_TXrGraphicsBinding *GHOST_XrGraphicsBindingCandidates;
 
-typedef struct GHOST_XrPose {
+typedef struct {
   float position[3];
   /* Blender convention (w, x, y, z) */
   float orientation_quat[4];
@@ -753,31 +753,8 @@ typedef struct GHOST_XrActionProfileInfo {
   const char *profile_path;
   uint32_t count_subaction_paths;
   const char **subaction_paths;
-  /** Bindings for each subaction path. */
+  /* Bindings for each subaction path. */
   const GHOST_XrActionBindingInfo *bindings;
 } GHOST_XrActionProfileInfo;
 
-typedef struct GHOST_XrControllerModelVertex {
-  float position[3];
-  float normal[3];
-} GHOST_XrControllerModelVertex;
-
-typedef struct GHOST_XrControllerModelComponent {
-  /** World space transform. */
-  float transform[4][4];
-  uint32_t vertex_offset;
-  uint32_t vertex_count;
-  uint32_t index_offset;
-  uint32_t index_count;
-} GHOST_XrControllerModelComponent;
-
-typedef struct GHOST_XrControllerModelData {
-  uint32_t count_vertices;
-  const GHOST_XrControllerModelVertex *vertices;
-  uint32_t count_indices;
-  const uint32_t *indices;
-  uint32_t count_components;
-  const GHOST_XrControllerModelComponent *components;
-} GHOST_XrControllerModelData;
-
 #endif /* WITH_XR_OPENXR */

intern/ghost/intern/GHOST_C-api.cpp

@@ -1069,39 +1069,4 @@ void GHOST_XrGetActionCustomdataArray(GHOST_XrContextHandle xr_contexthandle,
                      xr_context);
 }
 
-int GHOST_XrLoadControllerModel(GHOST_XrContextHandle xr_contexthandle, const char *subaction_path)
-{
-  GHOST_IXrContext *xr_context = (GHOST_IXrContext *)xr_contexthandle;
-  GHOST_XrSession *xr_session = xr_context->getSession();
-  GHOST_XR_CAPI_CALL_RET(xr_session->loadControllerModel(subaction_path), xr_context);
-  return 0;
-}
-
-void GHOST_XrUnloadControllerModel(GHOST_XrContextHandle xr_contexthandle,
-                                   const char *subaction_path)
-{
-  GHOST_IXrContext *xr_context = (GHOST_IXrContext *)xr_contexthandle;
-  GHOST_XrSession *xr_session = xr_context->getSession();
-  GHOST_XR_CAPI_CALL(xr_session->unloadControllerModel(subaction_path), xr_context);
-}
-
-int GHOST_XrUpdateControllerModelComponents(GHOST_XrContextHandle xr_contexthandle,
-                                            const char *subaction_path)
-{
-  GHOST_IXrContext *xr_context = (GHOST_IXrContext *)xr_contexthandle;
-  GHOST_XrSession *xr_session = xr_context->getSession();
-  GHOST_XR_CAPI_CALL_RET(xr_session->updateControllerModelComponents(subaction_path), xr_context);
-  return 0;
-}
-
-int GHOST_XrGetControllerModelData(GHOST_XrContextHandle xr_contexthandle,
-                                   const char *subaction_path,
-                                   GHOST_XrControllerModelData *r_data)
-{
-  GHOST_IXrContext *xr_context = (GHOST_IXrContext *)xr_contexthandle;
-  GHOST_XrSession *xr_session = xr_context->getSession();
-  GHOST_XR_CAPI_CALL_RET(xr_session->getControllerModelData(subaction_path, *r_data), xr_context);
-  return 0;
-}
-
 #endif /* WITH_XR_OPENXR */

intern/ghost/intern/GHOST_SystemPathsUnix.cpp

@@ -142,7 +142,7 @@ const char *GHOST_SystemPathsUnix::getUserSpecialDir(GHOST_TUserSpecialDirTypes
   }
 
   static string path = "";
-  /* Pipe stderr to /dev/null to avoid error prints. We will fail gracefully still. */
+  /* Pipe `stderr` to `/dev/null` to avoid error prints. We will fail gracefully still. */
   string command = string("xdg-user-dir ") + type_str + " 2> /dev/null";
 
   FILE *fstream = popen(command.c_str(), "r");
@@ -152,7 +152,7 @@ const char *GHOST_SystemPathsUnix::getUserSpecialDir(GHOST_TUserSpecialDirTypes
   std::stringstream path_stream;
   while (!feof(fstream)) {
     char c = fgetc(fstream);
-    /* xdg-user-dir ends the path with '\n'. */
+    /* `xdg-user-dir` ends the path with '\n'. */
     if (c == '\n') {
       break;
     }

intern/ghost/intern/GHOST_SystemX11.cpp

@@ -1044,7 +1044,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
 
 #ifdef USE_NON_LATIN_KB_WORKAROUND
       /* XXX: Code below is kinda awfully convoluted... Issues are:
-       * - In keyboards like latin ones, numbers need a 'Shift' to be accessed but key_sym
+       * - In keyboards like Latin ones, numbers need a 'Shift' to be accessed but key_sym
        *   is unmodified (or anyone swapping the keys with `xmodmap`).
        * - #XLookupKeysym seems to always use first defined key-map (see T47228), which generates
        *   key-codes unusable by ghost_key_from_keysym for non-Latin-compatible key-maps.
@@ -1131,7 +1131,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
         }
       }
 #else
-      /* In keyboards like latin ones,
+      /* In keyboards like Latin ones,
        * numbers needs a 'Shift' to be accessed but key_sym
        * is unmodified (or anyone swapping the keys with xmodmap).
        *

intern/ghost/intern/GHOST_XrAction.cpp

@@ -22,6 +22,7 @@
 #include <cstring>
 
 #include "GHOST_Types.h"
+
 #include "GHOST_XrException.h"
 #include "GHOST_Xr_intern.h"
 

intern/ghost/intern/GHOST_XrAction.h

@@ -26,7 +26,6 @@
 #include <map>
 #include <memory>
 #include <string>
-#include <vector>
 
 #include "GHOST_Util.h"
 

intern/ghost/intern/GHOST_XrContext.cpp

@@ -412,14 +412,11 @@ void GHOST_XrContext::getExtensionsToEnable(
     try_ext.push_back(XR_EXT_DEBUG_UTILS_EXTENSION_NAME);
   }
 
-  /* Interaction profile extensions. */
+  /* Try enabling interaction profile extensions. */
   try_ext.push_back(XR_EXT_HP_MIXED_REALITY_CONTROLLER_EXTENSION_NAME);
   try_ext.push_back(XR_HTC_VIVE_COSMOS_CONTROLLER_INTERACTION_EXTENSION_NAME);
   try_ext.push_back(XR_HUAWEI_CONTROLLER_INTERACTION_EXTENSION_NAME);
 
-  /* Controller model extension. */
-  try_ext.push_back(XR_MSFT_CONTROLLER_MODEL_EXTENSION_NAME);
-
   /* Varjo quad view extension. */
   try_ext.push_back(XR_VARJO_QUAD_VIEWS_EXTENSION_NAME);
 

intern/ghost/intern/GHOST_XrControllerModel.cpp

@@ -1,629 +0,0 @@
-/*
- * This program is free software; you can redistribute it and/or
- * modify it under the terms of the GNU General Public License
- * as published by the Free Software Foundation; either version 2
- * of the License, or (at your option) any later version.
- *
- * This program is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
- * GNU General Public License for more details.
- *
- * You should have received a copy of the GNU General Public License
- * along with this program; if not, write to the Free Software Foundation,
- * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
- */
-
-/** \file
- * \ingroup GHOST
- */
-
-#include <cassert>
-
-#include <Eigen/Core>
-#include <Eigen/Geometry>
-
-#include "GHOST_Types.h"
-#include "GHOST_XrException.h"
-#include "GHOST_Xr_intern.h"
-
-#include "GHOST_XrControllerModel.h"
-
-#define TINYGLTF_IMPLEMENTATION
-#define TINYGLTF_NO_STB_IMAGE
-#define TINYGLTF_NO_STB_IMAGE_WRITE
-#define STBIWDEF static inline
-#include "tiny_gltf.h"
-
-struct GHOST_XrControllerModelNode {
-  int32_t parent_idx = -1;
-  int32_t component_idx = -1;
-  float local_transform[4][4];
-};
-
-/* -------------------------------------------------------------------- */
-/** \name glTF Utilities
- *
- * Adapted from Microsoft OpenXR-Mixed Reality Samples (MIT License):
- * https://github.com/microsoft/OpenXR-MixedReality
- * \{ */
-
-struct GHOST_XrPrimitive {
-  std::vector<GHOST_XrControllerModelVertex> vertices;
-  std::vector<uint32_t> indices;
-};
-
-/**
- * Validate that an accessor does not go out of bounds of the buffer view that it references and
- * that the buffer view does not exceed the bounds of the buffer that it references
- */
-static void validate_accessor(const tinygltf::Accessor &accessor,
-                              const tinygltf::BufferView &buffer_view,
-                              const tinygltf::Buffer &buffer,
-                              size_t byte_stride,
-                              size_t element_size)
-{
-  /* Make sure the accessor does not go out of range of the buffer view. */
-  if (accessor.byteOffset + (accessor.count - 1) * byte_stride + element_size >
-      buffer_view.byteLength) {
-    throw GHOST_XrException("glTF: Accessor goes out of range of bufferview.");
-  }
-
-  /* Make sure the buffer view does not go out of range of the buffer. */
-  if (buffer_view.byteOffset + buffer_view.byteLength > buffer.data.size()) {
-    throw GHOST_XrException("glTF: BufferView goes out of range of buffer.");
-  }
-}
-
-template<float (GHOST_XrControllerModelVertex::*field)[3]>
-static void read_vertices(const tinygltf::Accessor &accessor,
-                          const tinygltf::BufferView &buffer_view,
-                          const tinygltf::Buffer &buffer,
-                          GHOST_XrPrimitive &primitive)
-{
-  if (accessor.type != TINYGLTF_TYPE_VEC3) {
-    throw GHOST_XrException(
-        "glTF: Accessor for primitive attribute has incorrect type (VEC3 expected).");
-  }
-
-  if (accessor.componentType != TINYGLTF_COMPONENT_TYPE_FLOAT) {
-    throw GHOST_XrException(
-        "glTF: Accessor for primitive attribute has incorrect component type (FLOAT expected).");
-  }
-
-  /* If stride is not specified, it is tightly packed. */
-  constexpr size_t packed_size = sizeof(float) * 3;
-  const size_t stride = buffer_view.byteStride == 0 ? packed_size : buffer_view.byteStride;
-  validate_accessor(accessor, buffer_view, buffer, stride, packed_size);
-
-  /* Resize the vertices vector, if necessary, to include room for the attribute data.
-     If there are multiple attributes for a primitive, the first one will resize, and the
-     subsequent will not need to. */
-  primitive.vertices.resize(accessor.count);
-
-  /* Copy the attribute value over from the glTF buffer into the appropriate vertex field. */
-  const uint8_t *buffer_ptr = buffer.data.data() + buffer_view.byteOffset + accessor.byteOffset;
-  for (size_t i = 0; i < accessor.count; i++, buffer_ptr += stride) {
-    memcpy(primitive.vertices[i].*field, buffer_ptr, stride);
-  }
-}
-
-static void load_attribute_accessor(const tinygltf::Model &gltf_model,
-                                    const std::string &attribute_name,
-                                    int accessor_id,
-                                    GHOST_XrPrimitive &primitive)
-{
-  const auto &accessor = gltf_model.accessors.at(accessor_id);
-
-  if (accessor.bufferView == -1) {
-    throw GHOST_XrException("glTF: Accessor for primitive attribute specifies no bufferview.");
-  }
-
-  const tinygltf::BufferView &buffer_view = gltf_model.bufferViews.at(accessor.bufferView);
-  if (buffer_view.target != TINYGLTF_TARGET_ARRAY_BUFFER && buffer_view.target != 0) {
-    throw GHOST_XrException(
-        "glTF: Accessor for primitive attribute uses bufferview with invalid 'target' type.");
-  }
-
-  const tinygltf::Buffer &buffer = gltf_model.buffers.at(buffer_view.buffer);
-
-  if (attribute_name.compare("POSITION") == 0) {
-    read_vertices<&GHOST_XrControllerModelVertex::position>(
-        accessor, buffer_view, buffer, primitive);
-  }
-  else if (attribute_name.compare("NORMAL") == 0) {
-    read_vertices<&GHOST_XrControllerModelVertex::normal>(
-        accessor, buffer_view, buffer, primitive);
-  }
-}
-
-/**
- * Reads index data from a glTF primitive into a GHOST_XrPrimitive. glTF indices may be 8bit, 16bit
- * or 32bit integers. This will coalesce indices from the source type(s) into a 32bit integer.
- */
-template<typename TSrcIndex>
-static void read_indices(const tinygltf::Accessor &accessor,
-                         const tinygltf::BufferView &buffer_view,
-                         const tinygltf::Buffer &buffer,
-                         GHOST_XrPrimitive &primitive)
-{
-  if (buffer_view.target != TINYGLTF_TARGET_ELEMENT_ARRAY_BUFFER &&
-      buffer_view.target != 0) { /* Allow 0 (not specified) even though spec doesn't seem to allow
-                                    this (BoomBox GLB fails). */
-    throw GHOST_XrException(
-        "glTF: Accessor for indices uses bufferview with invalid 'target' type.");
-  }
-
-  constexpr size_t component_size_bytes = sizeof(TSrcIndex);
-  if (buffer_view.byteStride != 0 &&
-      buffer_view.byteStride !=
-          component_size_bytes) { /* Index buffer must be packed per glTF spec. */
-    throw GHOST_XrException(
-        "glTF: Accessor for indices uses bufferview with invalid 'byteStride'.");
-  }
-
-  validate_accessor(accessor, buffer_view, buffer, component_size_bytes, component_size_bytes);
-
-  if ((accessor.count % 3) != 0) { /* Since only triangles are supported, enforce that the number
-                                      of indices is divisible by 3. */
-    throw GHOST_XrException("glTF: Unexpected number of indices for triangle primitive");
-  }
-
-  const TSrcIndex *index_buffer = reinterpret_cast<const TSrcIndex *>(
-      buffer.data.data() + buffer_view.byteOffset + accessor.byteOffset);
-  for (uint32_t i = 0; i < accessor.count; i++) {
-    primitive.indices.push_back(*(index_buffer + i));
-  }
-}
-
-/**
- * Reads index data from a glTF primitive into a GHOST_XrPrimitive.
- */
-static void load_index_accessor(const tinygltf::Model &gltf_model,
-                                const tinygltf::Accessor &accessor,
-                                GHOST_XrPrimitive &primitive)
-{
-  if (accessor.type != TINYGLTF_TYPE_SCALAR) {
-    throw GHOST_XrException("glTF: Accessor for indices specifies invalid 'type'.");
-  }
-
-  if (accessor.bufferView == -1) {
-    throw GHOST_XrException("glTF: Index accessor without bufferView is currently not supported.");
-  }
-
-  const tinygltf::BufferView &buffer_view = gltf_model.bufferViews.at(accessor.bufferView);
-  const tinygltf::Buffer &buffer = gltf_model.buffers.at(buffer_view.buffer);
-
-  if (accessor.componentType == TINYGLTF_COMPONENT_TYPE_UNSIGNED_BYTE) {
-    read_indices<uint8_t>(accessor, buffer_view, buffer, primitive);
-  }
-  else if (accessor.componentType == TINYGLTF_COMPONENT_TYPE_UNSIGNED_SHORT) {
-    read_indices<uint16_t>(accessor, buffer_view, buffer, primitive);
-  }
-  else if (accessor.componentType == TINYGLTF_COMPONENT_TYPE_UNSIGNED_INT) {
-    read_indices<uint32_t>(accessor, buffer_view, buffer, primitive);
-  }
-  else {
-    throw GHOST_XrException("glTF: Accessor for indices specifies invalid 'componentType'.");
-  }
-}
-
-static GHOST_XrPrimitive read_primitive(const tinygltf::Model &gltf_model,
-                                        const tinygltf::Primitive &gltf_primitive)
-{
-  if (gltf_primitive.mode != TINYGLTF_MODE_TRIANGLES) {
-    throw GHOST_XrException(
-        "glTF: Unsupported primitive mode. Only TINYGLTF_MODE_TRIANGLES is supported.");
-  }
-
-  GHOST_XrPrimitive primitive;
-
-  /* glTF vertex data is stored in an attribute dictionary.Loop through each attribute and insert
-   * it into the GHOST_XrPrimitive. */
-  for (const auto &[attr_name, accessor_idx] : gltf_primitive.attributes) {
-    load_attribute_accessor(gltf_model, attr_name, accessor_idx, primitive);
-  }
-
-  if (gltf_primitive.indices != -1) {
-    /* If indices are specified for the glTF primitive, read them into the GHOST_XrPrimitive. */
-    load_index_accessor(gltf_model, gltf_model.accessors.at(gltf_primitive.indices), primitive);
-  }
-
-  return primitive;
-}
-
-/**
- * Calculate node local and world transforms.
- */
-static void calc_node_transforms(const tinygltf::Node &gltf_node,
-                                 const float parent_transform[4][4],
-                                 float r_local_transform[4][4],
-                                 float r_world_transform[4][4])
-{
-  /* A node may specify either a 4x4 matrix or TRS (Translation - Rotation - Scale) values, but not
-   * both. */
-  if (gltf_node.matrix.size() == 16) {
-    const std::vector<double> &dm = gltf_node.matrix;
-    float m[4][4] = {(float)dm[0],
-                     (float)dm[1],
-                     (float)dm[2],
-                     (float)dm[3],
-                     (float)dm[4],
-                     (float)dm[5],
-                     (float)dm[6],
-                     (float)dm[7],
-                     (float)dm[8],
-                     (float)dm[9],
-                     (float)dm[10],
-                     (float)dm[11],
-                     (float)dm[12],
-                     (float)dm[13],
-                     (float)dm[14],
-                     (float)dm[15]};
-    memcpy(r_local_transform, m, sizeof(float) * 16);
-  }
-  else {
-    /* No matrix is present, so construct a matrix from the TRS values (each one is optional). */
-    std::vector<double> translation = gltf_node.translation;
-    std::vector<double> rotation = gltf_node.rotation;
-    std::vector<double> scale = gltf_node.scale;
-    Eigen::Matrix4f &m = *(Eigen::Matrix4f *)r_local_transform;
-    Eigen::Quaternionf q;
-    Eigen::Matrix3f scalemat;
-
-    if (translation.size() != 3) {
-      translation.resize(3);
-      translation[0] = translation[1] = translation[2] = 0.0;
-    }
-    if (rotation.size() != 4) {
-      rotation.resize(4);
-      rotation[0] = rotation[1] = rotation[2] = 0.0;
-      rotation[3] = 1.0;
-    }
-    if (scale.size() != 3) {
-      scale.resize(3);
-      scale[0] = scale[1] = scale[2] = 1.0;
-    }
-
-    q.w() = (float)rotation[3];
-    q.x() = (float)rotation[0];
-    q.y() = (float)rotation[1];
-    q.z() = (float)rotation[2];
-    q.normalize();
-
-    scalemat.setIdentity();
-    scalemat(0, 0) = (float)scale[0];
-    scalemat(1, 1) = (float)scale[1];
-    scalemat(2, 2) = (float)scale[2];
-
-    m.setIdentity();
-    m.block<3, 3>(0, 0) = q.toRotationMatrix() * scalemat;
-    m.block<3, 1>(0, 3) = Eigen::Vector3f(
-        (float)translation[0], (float)translation[1], (float)translation[2]);
-  }
-
-  *(Eigen::Matrix4f *)r_world_transform = *(Eigen::Matrix4f *)parent_transform *
-                                          *(Eigen::Matrix4f *)r_local_transform;
-}
-
-static void load_node(const tinygltf::Model &gltf_model,
-                      int gltf_node_id,
-                      int32_t parent_idx,
-                      const float parent_transform[4][4],
-                      const std::string &parent_name,
-                      const std::vector<XrControllerModelNodePropertiesMSFT> &node_properties,
-                      std::vector<GHOST_XrControllerModelVertex> &vertices,
-                      std::vector<uint32_t> &indices,
-                      std::vector<GHOST_XrControllerModelComponent> &components,
-                      std::vector<GHOST_XrControllerModelNode> &nodes,
-                      std::vector<int32_t> &node_state_indices)
-{
-  const tinygltf::Node &gltf_node = gltf_model.nodes.at(gltf_node_id);
-  float world_transform[4][4];
-
-  GHOST_XrControllerModelNode &node = nodes.emplace_back();
-  const int32_t node_idx = (int32_t)(nodes.size() - 1);
-  node.parent_idx = parent_idx;
-  calc_node_transforms(gltf_node, parent_transform, node.local_transform, world_transform);
-
-  for (size_t i = 0; i < node_properties.size(); ++i) {
-    if ((node_state_indices[i] < 0) && (parent_name == node_properties[i].parentNodeName) &&
-        (gltf_node.name == node_properties[i].nodeName)) {
-      node_state_indices[i] = node_idx;
-      break;
-    }
-  }
-
-  if (gltf_node.mesh != -1) {
-    const tinygltf::Mesh &gltf_mesh = gltf_model.meshes.at(gltf_node.mesh);
-
-    GHOST_XrControllerModelComponent &component = components.emplace_back();
-    node.component_idx = components.size() - 1;
-    memcpy(component.transform, world_transform, sizeof(component.transform));
-    component.vertex_offset = vertices.size();
-    component.index_offset = indices.size();
-
-    for (const tinygltf::Primitive &gltf_primitive : gltf_mesh.primitives) {
-      /* Read the primitive data from the glTF buffers. */
-      const GHOST_XrPrimitive primitive = read_primitive(gltf_model, gltf_primitive);
-
-      const size_t start_vertex = vertices.size();
-      size_t offset = start_vertex;
-      size_t count = primitive.vertices.size();
-      vertices.resize(offset + count);
-      memcpy(vertices.data() + offset,
-             primitive.vertices.data(),
-             count * sizeof(decltype(primitive.vertices)::value_type));
-
-      offset = indices.size();
-      count = primitive.indices.size();
-      indices.resize(offset + count);
-      for (size_t i = 0; i < count; i += 3) {
-        indices[offset + i + 0] = start_vertex + primitive.indices[i + 0];
-        indices[offset + i + 1] = start_vertex + primitive.indices[i + 2];
-        indices[offset + i + 2] = start_vertex + primitive.indices[i + 1];
-      }
-    }
-
-    component.vertex_count = vertices.size() - component.vertex_offset;
-    component.index_count = indices.size() - component.index_offset;
-  }
-
-  /* Recursively load all children. */
-  for (const int child_node_id : gltf_node.children) {
-    load_node(gltf_model,
-              child_node_id,
-              node_idx,
-              world_transform,
-              gltf_node.name,
-              node_properties,
-              vertices,
-              indices,
-              components,
-              nodes,
-              node_state_indices);
-  }
-}
-
-/** \} */
-
-/* -------------------------------------------------------------------- */
-/** \name OpenXR Extension Functions
- *
- * \{ */
-
-static PFN_xrGetControllerModelKeyMSFT g_xrGetControllerModelKeyMSFT = nullptr;
-static PFN_xrLoadControllerModelMSFT g_xrLoadControllerModelMSFT = nullptr;
-static PFN_xrGetControllerModelPropertiesMSFT g_xrGetControllerModelPropertiesMSFT = nullptr;
-static PFN_xrGetControllerModelStateMSFT g_xrGetControllerModelStateMSFT = nullptr;
-static XrInstance g_instance = XR_NULL_HANDLE;
-
-#define INIT_EXTENSION_FUNCTION(name) \
-  CHECK_XR( \
-      xrGetInstanceProcAddr(instance, #name, reinterpret_cast<PFN_xrVoidFunction *>(&g_##name)), \
-      "Failed to get pointer to extension function: " #name);
-
-static void init_controller_model_extension_functions(XrInstance instance)
-{
-  if (instance != g_instance) {
-    g_instance = instance;
-    g_xrGetControllerModelKeyMSFT = nullptr;
-    g_xrLoadControllerModelMSFT = nullptr;
-    g_xrGetControllerModelPropertiesMSFT = nullptr;
-    g_xrGetControllerModelStateMSFT = nullptr;
-  }
-
-  if (g_xrGetControllerModelKeyMSFT == nullptr) {
-    INIT_EXTENSION_FUNCTION(xrGetControllerModelKeyMSFT);
-  }
-  if (g_xrLoadControllerModelMSFT == nullptr) {
-    INIT_EXTENSION_FUNCTION(xrLoadControllerModelMSFT);
-  }
-  if (g_xrGetControllerModelPropertiesMSFT == nullptr) {
-    INIT_EXTENSION_FUNCTION(xrGetControllerModelPropertiesMSFT);
-  }
-  if (g_xrGetControllerModelStateMSFT == nullptr) {
-    INIT_EXTENSION_FUNCTION(xrGetControllerModelStateMSFT);
-  }
-}
-
-/** \} */
-
-/* -------------------------------------------------------------------- */
-/** \name GHOST_XrControllerModel
- *
- * \{ */
-
-GHOST_XrControllerModel::GHOST_XrControllerModel(XrInstance instance,
-                                                 const char *subaction_path_str)
-{
-  init_controller_model_extension_functions(instance);
-
-  CHECK_XR(xrStringToPath(instance, subaction_path_str, &m_subaction_path),
-           (std::string("Failed to get user path \"") + subaction_path_str + "\".").data());
-}
-
-GHOST_XrControllerModel::~GHOST_XrControllerModel()
-{
-  if (m_load_task.valid()) {
-    m_load_task.wait();
-  }
-}
-
-void GHOST_XrControllerModel::load(XrSession session)
-{
-  if (m_data_loaded || m_load_task.valid()) {
-    return;
-  }
-
-  /* Get model key. */
-  XrControllerModelKeyStateMSFT key_state{XR_TYPE_CONTROLLER_MODEL_KEY_STATE_MSFT};
-  CHECK_XR(g_xrGetControllerModelKeyMSFT(session, m_subaction_path, &key_state),
-           "Failed to get controller model key state.");
-
-  if (key_state.modelKey != XR_NULL_CONTROLLER_MODEL_KEY_MSFT) {
-    m_model_key = key_state.modelKey;
-    /* Load asynchronously. */
-    m_load_task = std::async(std::launch::async,
-                             [&, session = session]() { return loadControllerModel(session); });
-  }
-}
-
-void GHOST_XrControllerModel::loadControllerModel(XrSession session)
-{
-  /* Load binary buffers. */
-  uint32_t buf_size = 0;
-  CHECK_XR(g_xrLoadControllerModelMSFT(session, m_model_key, 0, &buf_size, nullptr),
-           "Failed to get controller model buffer size.");
-
-  std::vector<uint8_t> buf((size_t)buf_size);
-  CHECK_XR(g_xrLoadControllerModelMSFT(session, m_model_key, buf_size, &buf_size, buf.data()),
-           "Failed to load controller model binary buffers.");
-
-  /* Convert to glTF model. */
-  tinygltf::TinyGLTF gltf_loader;
-  tinygltf::Model gltf_model;
-  std::string err_msg;
-  {
-    /* Workaround for TINYGLTF_NO_STB_IMAGE define. Set custom image loader to prevent failure when
-     * parsing image data. */
-    auto load_img_func = [](tinygltf::Image *img,
-                            const int p0,
-                            std::string *p1,
-                            std::string *p2,
-                            int p3,
-                            int p4,
-                            const unsigned char *p5,
-                            int p6,
-                            void *user_pointer) -> bool {
-      (void)img;
-      (void)p0;
-      (void)p1;
-      (void)p2;
-      (void)p3;
-      (void)p4;
-      (void)p5;
-      (void)p6;
-      (void)user_pointer;
-      return true;
-    };
-    gltf_loader.SetImageLoader(load_img_func, nullptr);
-  }
-
-  if (!gltf_loader.LoadBinaryFromMemory(&gltf_model, &err_msg, nullptr, buf.data(), buf_size)) {
-    throw GHOST_XrException(("Failed to load glTF controller model: " + err_msg).c_str());
-  }
-
-  /* Get node properties. */
-  XrControllerModelPropertiesMSFT model_properties{XR_TYPE_CONTROLLER_MODEL_PROPERTIES_MSFT};
-  model_properties.nodeCapacityInput = 0;
-  CHECK_XR(g_xrGetControllerModelPropertiesMSFT(session, m_model_key, &model_properties),
-           "Failed to get controller model node properties count.");
-
-  std::vector<XrControllerModelNodePropertiesMSFT> node_properties(
-      model_properties.nodeCountOutput, {XR_TYPE_CONTROLLER_MODEL_NODE_PROPERTIES_MSFT});
-  model_properties.nodeCapacityInput = (uint32_t)node_properties.size();
-  model_properties.nodeProperties = node_properties.data();
-  CHECK_XR(g_xrGetControllerModelPropertiesMSFT(session, m_model_key, &model_properties),
-           "Failed to get controller model node properties.");
-
-  m_node_state_indices.resize(node_properties.size(), -1);
-
-  /* Get mesh vertex data. */
-  const tinygltf::Scene &default_scene = gltf_model.scenes.at(
-      (gltf_model.defaultScene == -1) ? 0 : gltf_model.defaultScene);
-  const int32_t root_idx = -1;
-  const std::string root_name = "";
-  float root_transform[4][4] = {0};
-  root_transform[0][0] = root_transform[1][1] = root_transform[2][2] = root_transform[3][3] = 1.0f;
-
-  for (const int node_id : default_scene.nodes) {
-    load_node(gltf_model,
-              node_id,
-              root_idx,
-              root_transform,
-              root_name,
-              node_properties,
-              m_vertices,
-              m_indices,
-              m_components,
-              m_nodes,
-              m_node_state_indices);
-  }
-
-  m_data_loaded = true;
-}
-
-void GHOST_XrControllerModel::updateComponents(XrSession session)
-{
-  if (!m_data_loaded) {
-    return;
-  }
-
-  /* Get node states. */
-  XrControllerModelStateMSFT model_state{XR_TYPE_CONTROLLER_MODEL_STATE_MSFT};
-  model_state.nodeCapacityInput = 0;
-  CHECK_XR(g_xrGetControllerModelStateMSFT(session, m_model_key, &model_state),
-           "Failed to get controller model node state count.");
-
-  const uint32_t count = model_state.nodeCountOutput;
-  std::vector<XrControllerModelNodeStateMSFT> node_states(
-      count, {XR_TYPE_CONTROLLER_MODEL_NODE_STATE_MSFT});
-  model_state.nodeCapacityInput = count;
-  model_state.nodeStates = node_states.data();
-  CHECK_XR(g_xrGetControllerModelStateMSFT(session, m_model_key, &model_state),
-           "Failed to get controller model node states.");
-
-  /* Update node local transforms. */
-  assert(m_node_state_indices.size() == count);
-
-  for (uint32_t state_idx = 0; state_idx < count; ++state_idx) {
-    const int32_t &node_idx = m_node_state_indices[state_idx];
-    if (node_idx >= 0) {
-      const XrPosef &pose = node_states[state_idx].nodePose;
-      Eigen::Matrix4f &m = *(Eigen::Matrix4f *)m_nodes[node_idx].local_transform;
-      Eigen::Quaternionf q(
-          pose.orientation.w, pose.orientation.x, pose.orientation.y, pose.orientation.z);
-      m.setIdentity();
-      m.block<3, 3>(0, 0) = q.toRotationMatrix();
-      m.block<3, 1>(0, 3) = Eigen::Vector3f(pose.position.x, pose.position.y, pose.position.z);
-    }
-  }
-
-  /* Calculate component transforms (in world space). */
-  std::vector<Eigen::Matrix4f> world_transforms(m_nodes.size());
-  uint32_t i = 0;
-  for (const GHOST_XrControllerModelNode &node : m_nodes) {
-    world_transforms[i] = (node.parent_idx >= 0) ? world_transforms[node.parent_idx] *
-                                                       *(Eigen::Matrix4f *)node.local_transform :
-                                                   *(Eigen::Matrix4f *)node.local_transform;
-    if (node.component_idx >= 0) {
-      memcpy(m_components[node.component_idx].transform,
-             world_transforms[i].data(),
-             sizeof(m_components[node.component_idx].transform));
-    }
-    ++i;
-  }
-}
-
-void GHOST_XrControllerModel::getData(GHOST_XrControllerModelData &r_data)
-{
-  if (m_data_loaded) {
-    r_data.count_vertices = (uint32_t)m_vertices.size();
-    r_data.vertices = m_vertices.data();
-    r_data.count_indices = (uint32_t)m_indices.size();
-    r_data.indices = m_indices.data();
-    r_data.count_components = (uint32_t)m_components.size();
-    r_data.components = m_components.data();
-  }
-  else {
-    r_data.count_vertices = 0;
-    r_data.vertices = nullptr;
-    r_data.count_indices = 0;
-    r_data.indices = nullptr;
-    r_data.count_components = 0;
-    r_data.components = nullptr;
-  }
-}
-
-/** \} */