Merge branch 'master' into xr-controller-support

- Use doxygen style comments
- Improve comment for XR region type
- Improve RNA descriptions for XR nav location/rotation/scale
- Split controller draw callback into separate model and aim functions

.clang-format

@@ -180,7 +180,6 @@ 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,10 +388,6 @@ 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)
@@ -410,7 +406,6 @@ 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)
@@ -1728,12 +1723,6 @@ 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.
@@ -2005,7 +1994,6 @@ 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,6 +38,7 @@ 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} charset-normalizer==${CHARSET_NORMALIZER_VERSION} urllib3==${URLLIB3_VERSION} certifi==${CERTIFI_VERSION} requests==${REQUESTS_VERSION} zstandard==${ZSTANDARD_VERSION} --no-binary :all:
+  INSTALL_COMMAND ${PYTHON_BINARY} -m pip install ${SITE_PACKAGES_EXTRA} cython==${CYTHON_VERSION} idna==${IDNA_VERSION} chardet==${CHARDET_VERSION} urllib3==${URLLIB3_VERSION} certifi==${CERTIFI_VERSION} requests==${REQUESTS_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.7)
+set(PYTHON_VERSION 3.9.2)
 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 fddb060b483bc01850a3f412eea1d954)
+set(PYTHON_HASH f0dc9000312abeb16de4eccce9a870ab)
 set(PYTHON_HASH_TYPE MD5)
 set(PYTHON_FILE Python-${PYTHON_VERSION}.tar.xz)
 
@@ -215,18 +215,17 @@ set(NANOVDB_HASH e7b9e863ec2f3b04ead171dec2322807)
 set(NANOVDB_HASH_TYPE MD5)
 set(NANOVDB_FILE nano-vdb-${NANOVDB_GIT_UID}.tar.gz)
 
-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(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(NUMPY_VERSION 1.21.2)
-set(NUMPY_SHORT_VERSION 1.21)
+set(NUMPY_VERSION 1.19.5)
+set(NUMPY_SHORT_VERSION 1.19)
 set(NUMPY_URI https://github.com/numpy/numpy/releases/download/v${NUMPY_VERSION}/numpy-${NUMPY_VERSION}.zip)
-set(NUMPY_HASH 5638d5dae3ca387be562912312db842e)
+set(NUMPY_HASH f6a1b48717c552bbc18f1adc3cc1fe0e)
 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.7"
+PYTHON_VERSION="3.9.2"
 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="3.2"
+PYTHON_IDNA_VERSION="2.9"
 PYTHON_IDNA_VERSION_MIN="2.0"
-PYTHON_IDNA_VERSION_MAX="3.2"
+PYTHON_IDNA_VERSION_MAX="3.0"
 PYTHON_IDNA_NAME="idna"
 
-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_CHARDET_VERSION="3.0.4"
+PYTHON_CHARDET_VERSION_MIN="3.0"
+PYTHON_CHARDET_VERSION_MAX="5.0"
+PYTHON_CHARDET_NAME="chardet"
 
-PYTHON_URLLIB3_VERSION="1.26.7"
+PYTHON_URLLIB3_VERSION="1.25.9"
 PYTHON_URLLIB3_VERSION_MIN="1.0"
 PYTHON_URLLIB3_VERSION_MAX="2.0"
 PYTHON_URLLIB3_NAME="urllib3"
 
-PYTHON_CERTIFI_VERSION="2021.10.8"
-PYTHON_CERTIFI_VERSION_MIN="2021.0"
-PYTHON_CERTIFI_VERSION_MAX="2023.0"
+PYTHON_CERTIFI_VERSION="2020.4.5.2"
+PYTHON_CERTIFI_VERSION_MIN="2020.0"
+PYTHON_CERTIFI_VERSION_MAX="2021.0"
 PYTHON_CERTIFI_NAME="certifi"
 
 PYTHON_REQUESTS_VERSION="2.23.0"
@@ -414,12 +414,7 @@ PYTHON_REQUESTS_VERSION_MIN="2.0"
 PYTHON_REQUESTS_VERSION_MAX="3.0"
 PYTHON_REQUESTS_NAME="requests"
 
-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="1.19.5"
 PYTHON_NUMPY_VERSION_MIN="1.14"
 PYTHON_NUMPY_VERSION_MAX="2.0"
 PYTHON_NUMPY_NAME="numpy"
@@ -427,22 +422,20 @@ 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_CHARSET_NORMALIZER_NAME $PYTHON_CHARSET_NORMALIZER_VERSION_MIN $PYTHON_CHARSET_NORMALIZER_VERSION_MAX"
+  "$PYTHON_CHARDET_NAME $PYTHON_CHARDET_VERSION_MIN $PYTHON_CHARDET_VERSION_MAX"
   "$PYTHON_URLLIB3_NAME $PYTHON_URLLIB3_VERSION_MIN $PYTHON_URLLIB3_VERSION_MAX"
   "$PYTHON_CERTIFI_NAME $PYTHON_CERTIFI_VERSION_MIN $PYTHON_CERTIFI_VERSION_MAX"
   "$PYTHON_REQUESTS_NAME $PYTHON_REQUESTS_VERSION_MIN $PYTHON_REQUESTS_VERSION_MAX"
-  "$PYTHON_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_CHARSET_NORMALIZER_NAME==$PYTHON_CHARSET_NORMALIZER_VERSION"
+  "$PYTHON_CHARDET_NAME==$PYTHON_CHARDET_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"
 )
 
@@ -1148,11 +1141,10 @@ 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_CHARSET_NORMALIZER_VERSION] (use pip).
+    ** [Chardet $PYTHON_CHARDET_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).
@@ -2021,7 +2013,7 @@ compile_OIIO() {
   fi
 
   # To be changed each time we make edits that would modify the compiled result!
-  oiio_magic=18
+  oiio_magic=17
   _init_oiio
 
   # Force having own builds for the dependencies.
@@ -2096,7 +2088,6 @@ 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

@@ -0,0 +1,27 @@
+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

@@ -1,81 +0,0 @@
-# - 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,7 +24,6 @@ import project_source_info
 import subprocess
 import sys
 import os
-import tempfile
 
 from typing import (
     Any,
@@ -36,6 +35,7 @@ USE_QUIET = (os.environ.get("QUIET", None) is not None)
 
 CHECKER_IGNORE_PREFIX = [
     "extern",
+    "intern/moto",
 ]
 
 CHECKER_BIN = "cppcheck"
@@ -47,19 +47,13 @@ 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 cppcheck() -> None:
+def main() -> None:
     source_info = project_source_info.build_info(ignore_prefix_list=CHECKER_IGNORE_PREFIX)
     source_defines = project_source_info.build_defines_as_args()
 
@@ -84,10 +78,7 @@ def cppcheck() -> None:
             percent_str = "[" + ("%.2f]" % percent).rjust(7) + " %:"
 
             sys.stdout.flush()
-            sys.stdout.write("%s %s\n" % (
-                percent_str,
-                os.path.relpath(c, project_source_info.SOURCE_DIR)
-            ))
+            sys.stdout.write("%s " % percent_str)
 
         return subprocess.Popen(cmd)
 
@@ -99,11 +90,5 @@ def cppcheck() -> 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,9 +243,7 @@ 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
     """
@@ -273,20 +271,14 @@ def queue_processes(
 
                 if len(processes) <= job_total:
                     break
-                time.sleep(sleep)
+                else:
+                    time.sleep(0.1)
 
             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: master
+        -   branch: xr-controller-support
             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", branch, branch_fallback),
+        ("release/scripts/addons", "xr-controller-support", 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,7 +1101,6 @@ 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,7 +609,6 @@ 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,7 +736,6 @@ 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,8 +574,6 @@ 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,6 +211,7 @@ 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,6 +1197,12 @@ 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,6 +792,7 @@ 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,9 +927,6 @@ 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,14 +72,15 @@ 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);
+    }
 
-  /* 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;
   }
-
-  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:
-  explicit BlenderOutputDriver(BL::RenderEngine &b_engine);
+  BlenderOutputDriver(BL::RenderEngine &b_engine);
   ~BlenderOutputDriver();
 
   virtual void write_render_tile(const Tile &tile) override;

intern/cycles/blender/blender_session.cpp

@@ -504,10 +504,6 @@ 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,6 +545,8 @@ 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);
 
@@ -602,6 +604,10 @@ 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,6 +50,10 @@ 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,20 +439,61 @@ 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;
     }
   }
@@ -481,8 +522,10 @@ 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;
@@ -492,6 +535,7 @@ 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();
     }
@@ -501,6 +545,8 @@ 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());
@@ -513,6 +559,8 @@ 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;
@@ -561,14 +609,18 @@ 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];
@@ -581,6 +633,15 @@ 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,12 +67,8 @@ BVHBuild::~BVHBuild()
 
 /* Adding References */
 
-void BVHBuild::add_reference_triangles(BoundBox &root,
-                                       BoundBox &center,
-                                       Mesh *mesh,
-                                       int object_index)
+void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *mesh, int i)
 {
-  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);
@@ -85,7 +81,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root,
       BoundBox bounds = BoundBox::empty;
       t.bounds_grow(verts, bounds);
       if (bounds.valid() && t.valid(verts)) {
-        references.push_back(BVHReference(bounds, j, object_index, primitive_type));
+        references.push_back(BVHReference(bounds, j, i, PRIMITIVE_TRIANGLE));
         root.grow(bounds);
         center.grow(bounds.center2());
       }
@@ -105,7 +101,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root,
         t.bounds_grow(vert_steps + step * num_verts, bounds);
       }
       if (bounds.valid()) {
-        references.push_back(BVHReference(bounds, j, object_index, primitive_type));
+        references.push_back(BVHReference(bounds, j, i, PRIMITIVE_MOTION_TRIANGLE));
         root.grow(bounds);
         center.grow(bounds.center2());
       }
@@ -144,7 +140,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root,
         if (bounds.valid()) {
           const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1;
           references.push_back(
-              BVHReference(bounds, j, object_index, primitive_type, prev_time, curr_time));
+              BVHReference(bounds, j, i, PRIMITIVE_MOTION_TRIANGLE, prev_time, curr_time));
           root.grow(bounds);
           center.grow(bounds.center2());
         }
@@ -157,14 +153,18 @@ void BVHBuild::add_reference_triangles(BoundBox &root,
   }
 }
 
-void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair, int object_index)
+void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair, int i)
 {
   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 = hair->primitive_type();
+  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 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, object_index, packed_type));
+          references.push_back(BVHReference(bounds, j, i, 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, object_index, packed_type));
+          references.push_back(BVHReference(bounds, j, i, packed_type));
           root.grow(bounds);
           center.grow(bounds.center2());
         }
@@ -254,8 +254,7 @@ 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, object_index, packed_type, prev_time, curr_time));
+            references.push_back(BVHReference(bounds, j, i, packed_type, prev_time, curr_time));
             root.grow(bounds);
             center.grow(bounds.center2());
           }
@@ -269,18 +268,15 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
   }
 }
 
-void BVHBuild::add_reference_geometry(BoundBox &root,
-                                      BoundBox &center,
-                                      Geometry *geom,
-                                      int object_index)
+void BVHBuild::add_reference_geometry(BoundBox &root, BoundBox &center, Geometry *geom, int i)
 {
   if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
     Mesh *mesh = static_cast<Mesh *>(geom);
-    add_reference_triangles(root, center, mesh, object_index);
+    add_reference_triangles(root, center, mesh, i);
   }
   else if (geom->geometry_type == Geometry::HAIR) {
     Hair *hair = static_cast<Hair *>(geom);
-    add_reference_curves(root, center, hair, object_index);
+    add_reference_curves(root, center, hair, i);
   }
 }
 

intern/cycles/bvh/bvh_embree.cpp

@@ -89,9 +89,20 @@ 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;
@@ -136,7 +147,10 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
       }
       else {
         kernel_embree_convert_hit(kg, ray, hit, &current_isect);
-        if (ctx->local_object_id != current_isect.object) {
+        int object = (current_isect.object == OBJECT_NONE) ?
+                         kernel_tex_fetch(__prim_object, current_isect.prim) :
+                         current_isect.object;
+        if (ctx->local_object_id != object) {
           /* This tells Embree to continue tracing. */
           *args->valid = 0;
           break;
@@ -155,49 +169,41 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
         break;
       }
 
-      LocalIntersection *local_isect = ctx->local_isect;
+      /* 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;
+        }
+      }
+
       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;
-          }
-        }
 
-        local_isect->num_hits++;
-
-        if (local_isect->num_hits <= ctx->max_hits) {
-          hit_idx = local_isect->num_hits - 1;
+        ++ctx->local_isect->num_hits;
+        if (ctx->local_isect->num_hits <= ctx->max_hits) {
+          hit_idx = ctx->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) % local_isect->num_hits;
+          hit_idx = lcg_step_uint(ctx->lcg_state) % ctx->local_isect->num_hits;
 
           if (hit_idx >= ctx->max_hits) {
             /* This tells Embree to continue tracing. */
             *args->valid = 0;
-            return;
+            break;
           }
         }
       }
       else {
-        /* 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;
+        ctx->local_isect->num_hits = 1;
       }
-
       /* record intersection */
-      local_isect->hits[hit_idx] = current_isect;
-      local_isect->Ng[hit_idx] = normalize(make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z));
+      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));
       /* This tells Embree to continue tracing. */
       *args->valid = 0;
       break;
@@ -207,11 +213,21 @@ 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;
+        uint tri_object = (isect->object == OBJECT_NONE) ?
+                              kernel_tex_fetch(__prim_object, isect->prim) :
+                              isect->object;
         int object_flag = kernel_tex_fetch(__object_flag, tri_object);
         if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
           --ctx->num_hits;
@@ -233,7 +249,7 @@ static void rtc_filter_func_thick_curve(const RTCFilterFunctionNArguments *args)
   const RTCRay *ray = (RTCRay *)args->ray;
   RTCHit *hit = (RTCHit *)args->hit;
 
-  /* Always ignore back-facing intersections. */
+  /* Always ignore backfacing 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;
@@ -246,7 +262,7 @@ static void rtc_filter_occluded_func_thick_curve(const RTCFilterFunctionNArgumen
   const RTCRay *ray = (RTCRay *)args->ray;
   RTCHit *hit = (RTCHit *)args->hit;
 
-  /* Always ignore back-facing intersections. */
+  /* Always ignore backfacing 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;
@@ -440,7 +456,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->prim_offset;
+  size_t prim_offset = mesh->optix_prim_offset;
 
   const Attribute *attr_mP = NULL;
   size_t num_motion_steps = 1;
@@ -609,7 +625,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->curve_segment_offset;
+  size_t prim_offset = hair->optix_prim_offset;
 
   const Attribute *attr_mP = NULL;
   size_t num_motion_steps = 1;
@@ -686,7 +702,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->prim_offset);
+          rtcSetGeometryUserData(geom, (void *)mesh->optix_prim_offset);
           rtcCommitGeometry(geom);
         }
       }
@@ -695,7 +711,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->curve_segment_offset);
+          rtcSetGeometryUserData(geom, (void *)hair->optix_prim_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 "Found CUDA ${CUDA_NVCC_EXECUTABLE} (${CUDA_VERSION})")
+    message(STATUS "CUDA nvcc = ${CUDA_NVCC_EXECUTABLE}")
   else()
     message(STATUS "CUDA compiler not found, disabling WITH_CYCLES_CUDA_BINARIES")
     set(WITH_CYCLES_CUDA_BINARIES OFF)
@@ -537,16 +537,6 @@ 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,13 +768,7 @@ void OptiXDevice::denoise_color_read(DenoiseContext &context, const DenoisePass
   destination.num_components = 3;
   destination.pixel_stride = context.buffer_params.pass_stride;
 
-  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);
+  pass_accessor.get_render_tile_pixels(context.render_buffers, context.buffer_params, destination);
 }
 
 bool OptiXDevice::denoise_filter_color_preprocess(DenoiseContext &context, const DenoisePass &pass)
@@ -1252,7 +1246,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->curve_segment_offset;
+        build_input.curveArray.primitiveIndexOffset = hair->optix_prim_offset;
       }
       else {
         /* Disable visibility test any-hit program, since it is already checked during
@@ -1265,7 +1259,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->curve_segment_offset;
+        build_input.customPrimitiveArray.primitiveIndexOffset = hair->optix_prim_offset;
       }
 
       if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
@@ -1334,7 +1328,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->prim_offset;
+      build_input.triangleArray.primitiveIndexOffset = mesh->optix_prim_offset;
 
       if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
         progress.set_error("Failed to build OptiX acceleration structure");
@@ -1401,8 +1395,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. */
-      instance.instanceId = ob->get_device_index();
+      /* Set user instance ID to object index (but leave low bit blank). */
+      instance.instanceId = ob->get_device_index() << 1;
 
       /* Add some of the object visibility bits to the mask.
        * __prim_visibility contains the combined visibility bits of all instances, so is not
@@ -1514,6 +1508,9 @@ 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,13 +289,7 @@ class OIDNDenoiseContext {
      * pixels. */
     const PassAccessorCPU pass_accessor(pass_access_info, 1.0f, num_samples_);
 
-    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);
+    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,6 +149,9 @@ 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,22 +99,17 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_float(
 {
   DCHECK_EQ(destination.stride, 0) << "Custom stride for float destination is not implemented.";
 
-  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 float *buffer_data = render_buffers->buffer.data();
   const int pixel_stride = destination.pixel_stride ? destination.pixel_stride :
                                                       destination.num_components;
 
-  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;
+  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;
 
-    for (int64_t x = 0; x < buffer_params.window_width;
-         ++x, buffer += pass_stride, pixel += pixel_stride) {
       processor(kfilm_convert, buffer, pixel);
     }
   });
@@ -128,28 +123,26 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_half_rgba(
     const Destination &destination,
     const Processor &processor) const
 {
-  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 float *buffer_data = render_buffers->buffer.data();
 
   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.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) {
+  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;
 
-      float pixel_rgba[4];
-      processor(kfilm_convert, buffer, pixel_rgba);
+      float pixel[4];
+      processor(kfilm_convert, buffer, pixel);
 
-      film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel_rgba);
+      film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel);
 
-      float4_store_half(&pixel->x,
-                        make_float4(pixel_rgba[0], pixel_rgba[1], pixel_rgba[2], pixel_rgba[3]));
+      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]));
     }
   });
 }

intern/cycles/integrator/pass_accessor_gpu.cpp

@@ -43,13 +43,10 @@ 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.window_width * buffer_params.window_height;
+  const int work_size = buffer_params.width * buffer_params.height;
 
   const int destination_stride = destination.stride != 0 ? destination.stride :
-                                                           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;
+                                                           buffer_params.width;
 
   if (destination.d_pixels) {
     DCHECK_EQ(destination.stride, 0) << "Custom stride for float destination is not implemented.";
@@ -58,8 +55,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.window_width),
-                    const_cast<int *>(&offset),
+                    const_cast<int *>(&buffer_params.width),
+                    const_cast<int *>(&buffer_params.offset),
                     const_cast<int *>(&buffer_params.stride),
                     const_cast<int *>(&destination.offset),
                     const_cast<int *>(&destination_stride)};
@@ -73,8 +70,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.window_width),
-                    const_cast<int *>(&offset),
+                    const_cast<int *>(&buffer_params.width),
+                    const_cast<int *>(&buffer_params.offset),
                     const_cast<int *>(&buffer_params.stride),
                     const_cast<int *>(&destination.offset),
                     const_cast<int *>(&destination_stride)};

intern/cycles/integrator/path_trace.cpp

@@ -234,53 +234,42 @@ 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 window_height = buffer_params.window_height;
+  const int height = buffer_params.height;
 
   int current_y = 0;
   for (int i = 0; i < num_works; ++i) {
     const double weight = work_balance_infos[i].weight;
-    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);
+    const int slice_height = max(lround(height * weight), 1);
 
     /* Disallow negative values to deal with situations when there are more compute devices than
      * scan-lines. */
-    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;
+    const int remaining_height = max(0, height - current_y);
 
+    BufferParams slide_params = buffer_params;
+    slide_params.full_y = buffer_params.full_y + current_y;
     if (i < num_works - 1) {
-      slice_params.window_height = min(slice_window_height, remaining_window_height);
+      slide_params.height = min(slice_height, remaining_height);
     }
     else {
-      slice_params.window_height = remaining_window_height;
+      slide_params.height = remaining_height;
     }
 
-    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);
+    slide_params.update_offset_stride();
 
-    slice_params.update_offset_stride();
+    callback(path_trace_works[i].get(), slide_params);
 
-    callback(path_trace_works[i].get(), slice_params);
-
-    current_y += slice_params.window_height;
+    current_y += slide_params.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);
@@ -293,12 +282,6 @@ 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;
@@ -317,12 +300,9 @@ 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);
@@ -1025,12 +1005,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.window_width,
-                     full_frame_state_.render_buffers->params.window_height);
+    return make_int2(full_frame_state_.render_buffers->params.width,
+                     full_frame_state_.render_buffers->params.height);
   }
 
   const Tile &tile = tile_manager_.get_current_tile();
-  return make_int2(tile.window_width, tile.window_height);
+  return make_int2(tile.width, tile.height);
 }
 
 int2 PathTrace::get_render_tile_offset() const
@@ -1040,7 +1020,7 @@ int2 PathTrace::get_render_tile_offset() const
   }
 
   const Tile &tile = tile_manager_.get_current_tile();
-  return make_int2(tile.x + tile.window_x, tile.y + tile.window_y);
+  return make_int2(tile.x, tile.y);
 }
 
 int2 PathTrace::get_render_size() const

intern/cycles/integrator/path_trace_work.cpp

@@ -134,8 +134,7 @@ 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_buffer_params_.window_y) -
-                       (effective_big_tile_params_.full_y + effective_big_tile_params_.window_y);
+  const int offset_y = effective_buffer_params_.full_y - effective_big_tile_params_.full_y;
   const int width = effective_buffer_params_.width;
 
   PassAccessor::Destination slice_destination = destination;
@@ -192,10 +191,8 @@ 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 +
-                        effective_buffer_params_.window_x;
-  const int texture_y = effective_buffer_params_.full_y - effective_full_params_.full_y +
-                        effective_buffer_params_.window_y;
+  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;
 
   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,7 +23,6 @@
 #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"
 
@@ -31,38 +30,6 @@
 
 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,
@@ -80,7 +47,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(estimate_single_state_size())),
+      max_num_paths_(queue_->num_concurrent_states(sizeof(IntegratorStateCPU))),
       min_num_active_paths_(queue_->num_concurrent_busy_states()),
       max_active_path_index_(0)
 {
@@ -129,27 +96,16 @@ 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()
@@ -756,13 +712,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_.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 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 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;
+  const int texture_x = full_x - effective_full_params_.full_x;
+  const int texture_y = full_y - effective_full_params_.full_y;
 
   /* Re-allocate display memory if needed, and make sure the device pointer is allocated.
    *

intern/cycles/kernel/CMakeLists.txt

@@ -404,27 +404,16 @@ if(WITH_CYCLES_CUDA_BINARIES)
             -cuda-toolkit-dir "${cuda_toolkit_root_dir}"
         DEPENDS ${kernel_sources} cycles_cubin_cc)
     else()
-      set(_cuda_nvcc_args
+      add_custom_command(
+        OUTPUT ${cuda_file}
+        COMMAND ${cuda_nvcc_executable}
             -arch=${arch}
             ${CUDA_NVCC_FLAGS}
             --${format}
             ${CMAKE_CURRENT_SOURCE_DIR}${cuda_kernel_src}
             --ptxas-options="-v"
-            ${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)
+            ${cuda_flags}
+        DEPENDS ${kernel_sources})
     endif()
     delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${cuda_file}" ${CYCLES_INSTALL_PATH}/lib)
     list(APPEND cuda_cubins ${cuda_file})
@@ -483,10 +472,20 @@ endif()
 
 # HIP module
 
-if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
+if(WITH_CYCLES_HIP_BINARIES)
   # 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}
@@ -543,24 +542,23 @@ if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
         -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})
-    # Compile regular kernel
-    CYCLES_HIP_KERNEL_ADD(${arch} ${prev_arch} kernel "" "${hip_sources}" FALSE)
+      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()
   endforeach()
 
   add_custom_target(cycles_kernel_hip ALL DEPENDS ${hip_fatbins})

intern/cycles/kernel/bvh/bvh_embree.h

@@ -106,6 +106,9 @@ 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) {
@@ -118,37 +121,27 @@ 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 = 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->object = OBJECT_NONE;
   }
+  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 object)
+ccl_device_inline void kernel_embree_convert_sss_hit(const KernelGlobals *kg,
+                                                     const RTCRay *ray,
+                                                     const RTCHit *hit,
+                                                     Intersection *isect,
+                                                     int local_object_id)
 {
   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, object * 2));
+      rtcGetGeometry(kernel_data.bvh.scene, local_object_id * 2));
   isect->prim = hit->primID +
                 (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID));
-  isect->object = object;
-  isect->type = kernel_tex_fetch(__objects, object).primitive_type;
+  isect->object = local_object_id;
+  isect->type = kernel_tex_fetch(__prim_type, isect->prim);
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/bvh/bvh_shadow_all.h

@@ -130,6 +130,7 @@ 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];
@@ -137,15 +138,14 @@ ccl_device_inline
 
           /* primitive intersection */
           while (prim_addr < prim_addr2) {
-            kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) ==
-                          (type & PRIMITIVE_ALL));
+            kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) == p_type);
             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 (type & PRIMITIVE_ALL) {
+            switch (p_type) {
               case PRIMITIVE_TRIANGLE: {
                 hit = triangle_intersect(
                     kg, isect, P, dir, isect_t, visibility, object, prim_addr);
@@ -163,20 +163,17 @@ ccl_device_inline
               case PRIMITIVE_MOTION_CURVE_THICK:
               case PRIMITIVE_CURVE_RIBBON:
               case PRIMITIVE_MOTION_CURVE_RIBBON: {
-                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);
-
+                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);
                 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++) {
-                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);
+                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 (hit) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)

intern/cycles/kernel/bvh/bvh_util.h

@@ -118,18 +118,19 @@ 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 = isect->prim;
+  const int prim = kernel_tex_fetch(__prim_index, isect->prim);
   int shader = 0;
 
 #ifdef __HAIR__
-  if (isect->type & PRIMITIVE_ALL_TRIANGLE)
+  if (kernel_tex_fetch(__prim_type, isect->prim) & PRIMITIVE_ALL_TRIANGLE)
 #endif
   {
     shader = kernel_tex_fetch(__tri_shader, prim);
   }
 #ifdef __HAIR__
   else {
-    shader = kernel_tex_fetch(__curves, prim).shader_id;
+    float4 str = kernel_tex_fetch(__curves, prim);
+    shader = __float_as_int(str.z);
   }
 #endif
 
@@ -137,19 +138,21 @@ 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 prim, const int isect_type)
+    const KernelGlobals *ccl_restrict kg, const int isect_prim)
 {
+  const int prim = kernel_tex_fetch(__prim_index, isect_prim);
   int shader = 0;
 
 #ifdef __HAIR__
-  if (isect_type & PRIMITIVE_ALL_TRIANGLE)
+  if (kernel_tex_fetch(__prim_type, isect_prim) & PRIMITIVE_ALL_TRIANGLE)
 #endif
   {
     shader = kernel_tex_fetch(__tri_shader, prim);
   }
 #ifdef __HAIR__
   else {
-    shader = kernel_tex_fetch(__curves, prim).shader_id;
+    float4 str = kernel_tex_fetch(__curves, prim);
+    shader = __float_as_int(str.z);
   }
 #endif
 
@@ -159,13 +162,25 @@ 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, isect->type);
+  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);
 }
 
 ccl_device_forceinline int intersection_get_object_flags(const KernelGlobals *ccl_restrict kg,
                                                          const Intersection *ccl_restrict isect)
 {
-  return kernel_tex_fetch(__object_flag, isect->object);
+  const int object = intersection_get_object(kg, isect);
+
+  return kernel_tex_fetch(__object_flag, object);
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/closure/bsdf_principled_diffuse.h

@@ -35,25 +35,21 @@ 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, float *pdf)
+    const PrincipledDiffuseBsdf *bsdf, float3 N, float3 V, float3 L, float3 H, float *pdf)
 {
   float NdotL = dot(N, L);
+  float NdotV = dot(N, V);
 
-  if (NdotL <= 0) {
+  if (NdotL <= 0 || NdotV <= 0) {
+    *pdf = 0.0f;
     return make_float3(0.0f, 0.0f, 0.0f);
   }
 
-  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 LdotH = dot(L, H);
 
   float FL = schlick_fresnel(NdotL), FV = schlick_fresnel(NdotV);
-  const float Fd90 = 0.5f + LH2 * bsdf->roughness;
-  float Fd = (1.0f - FL + Fd90 * FL) * (1.0f - FV + Fd90 * FV);
+  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);
 
   float value = M_1_PI_F * NdotL * Fd;
 
@@ -76,10 +72,11 @@ 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, pdf);
+    return calculate_principled_diffuse_brdf(bsdf, N, V, L, H, pdf);
   }
   else {
     *pdf = 0.0f;
@@ -115,7 +112,9 @@ 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) {
-    *eval = calculate_principled_diffuse_brdf(bsdf, N, I, *omega_in, pdf);
+    float3 H = normalize(I + *omega_in);
+
+    *eval = calculate_principled_diffuse_brdf(bsdf, N, I, *omega_in, H, pdf);
 
 #ifdef __RAY_DIFFERENTIALS__
     // TODO: find a better approximation for the diffuse bounce

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

@@ -424,12 +424,8 @@ ccl_device_inline void kernel_gpu_film_convert_common(const KernelFilmConvert *k
     return;
   }
 
-  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;
-
+  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;
   ccl_global float *pixel = pixels +
                             (render_pixel_index + dst_offset) * kfilm_convert->pixel_stride;
 
@@ -455,17 +451,17 @@ ccl_device_inline void kernel_gpu_film_convert_half_rgba_common_rgba(
     return;
   }
 
-  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;
+  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;
 
   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,15 +41,22 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
   return (T *)(((uint64_t)optixGetPayload_3() << 32) | optixGetPayload_2());
 }
 
-ccl_device_forceinline int get_object_id()
+template<bool always = false> ccl_device_forceinline uint 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. */
-  return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
+  uint object = optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
 #else
-  return optixGetInstanceId();
+  uint object = 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()
@@ -101,7 +108,7 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
 #endif
 
 #ifdef __BVH_LOCAL__
-  const int object = get_object_id();
+  const uint object = get_object_id<true>();
   if (object != optixGetPayload_4() /* local_object */) {
     /* Only intersect with matching object. */
     return optixIgnoreIntersection();
@@ -145,23 +152,21 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
     local_isect->num_hits = 1;
   }
 
-  const int prim = optixGetPrimitiveIndex();
-
   Intersection *isect = &local_isect->hits[hit];
   isect->t = optixGetRayTmax();
-  isect->prim = prim;
+  isect->prim = optixGetPrimitiveIndex();
   isect->object = get_object_id();
-  isect->type = kernel_tex_fetch(__objects, isect->object).primitive_type;
+  isect->type = kernel_tex_fetch(__prim_type, isect->prim);
 
   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(__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));
+  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));
   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). */
@@ -174,32 +179,25 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
 #ifdef __SHADOW_RECORD_ALL__
   bool ignore_intersection = false;
 
-  int prim = optixGetPrimitiveIndex();
-  const uint object = get_object_id();
+  const uint prim = optixGetPrimitiveIndex();
 #  ifdef __VISIBILITY_FLAG__
   const uint visibility = optixGetPayload_4();
-  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
+  if ((kernel_tex_fetch(__prim_visibility, prim) & 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;
@@ -247,8 +245,8 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
     isect->v = v;
     isect->t = optixGetRayTmax();
     isect->prim = prim;
-    isect->object = object;
-    isect->type = type;
+    isect->object = get_object_id();
+    isect->type = kernel_tex_fetch(__prim_type, prim);
 
 #  ifdef __TRANSPARENT_SHADOWS__
     /* Detect if this surface has a shader with transparent shadows. */
@@ -276,14 +274,15 @@ 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(__objects, object).visibility & visibility) == 0) {
+  if ((kernel_tex_fetch(__prim_visibility, prim) & 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();
   }
@@ -302,9 +301,9 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
 #endif
 
 #ifdef __VISIBILITY_FLAG__
-  const uint object = get_object_id();
+  const uint prim = optixGetPrimitiveIndex();
   const uint visibility = optixGetPayload_4();
-  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
+  if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
     return optixIgnoreIntersection();
   }
 
@@ -317,39 +316,28 @@ 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_4(object);
+  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()));
 
   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 int prim, const int type)
+ccl_device_inline void optix_intersection_curve(const uint prim, const uint type)
 {
-  const int object = get_object_id();
-
-#  ifdef __VISIBILITY_FLAG__
+  const uint object = get_object_id<true>();
   const uint visibility = optixGetPayload_4();
-  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
-    return;
-  }
-#  endif
 
   float3 P = optixGetObjectRayOrigin();
   float3 dir = optixGetObjectRayDirection();
@@ -370,7 +358,7 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type)
   if (isect.t != FLT_MAX)
     isect.t *= len;
 
-  if (curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
+  if (curve_intersect(NULL, &isect, P, dir, isect.t, visibility, object, prim, time, type)) {
     optixReportIntersection(isect.t / len,
                             type & PRIMITIVE_ALL,
                             __float_as_int(isect.u),  /* Attribute_0 */
@@ -380,9 +368,9 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type)
 
 extern "C" __global__ void __intersection__curve_ribbon()
 {
-  const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, optixGetPrimitiveIndex());
-  const int prim = segment.prim;
-  const int type = segment.type;
+  const uint prim = optixGetPrimitiveIndex();
+  const uint type = kernel_tex_fetch(__prim_type, prim);
+
   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)) {
-    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = __float_as_int(curvedata.x) + 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)) {
-    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = __float_as_int(curvedata.x) + 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)) {
-    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = __float_as_int(curvedata.x) + 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)) {
-    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = __float_as_int(curvedata.x) + 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) {
-    KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
-    int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+    float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
+    int k0 = __float_as_int(curvedata.x) + 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)
 {
-  KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
-  int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+  float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
+  int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
   int k1 = k0 + 1;
 
   float4 P_curve[2];

intern/cycles/kernel/geom/geom_curve_intersect.h

@@ -630,19 +630,33 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
                                             const float3 P,
                                             const float3 dir,
                                             const float tmax,
+                                            uint visibility,
                                             int object,
-                                            int prim,
+                                            int curveAddr,
                                             float time,
                                             int type)
 {
   const bool is_motion = (type & PRIMITIVE_ALL_MOTION);
 
-  KernelCurve kcurve = kernel_tex_fetch(__curves, prim);
+#  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
 
-  int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
+  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 k1 = k0 + 1;
-  int ka = max(k0 - 1, kcurve.first_key);
-  int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 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);
 
   float4 curve[4];
   if (!is_motion) {
@@ -652,14 +666,21 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
     curve[3] = kernel_tex_fetch(__curve_keys, kb);
   }
   else {
-    motion_curve_keys(kg, object, prim, time, ka, k0, k1, kb, curve);
+    int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
+    motion_curve_keys(kg, fobject, 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 = prim;
+      isect->prim = curveAddr;
       isect->object = object;
       isect->type = type;
       return true;
@@ -669,7 +690,7 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
   }
   else {
     if (curve_intersect_recursive(P, dir, tmax, curve, isect)) {
-      isect->prim = prim;
+      isect->prim = curveAddr;
       isect->object = object;
       isect->type = type;
       return true;
@@ -687,7 +708,7 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
                                           const int isect_object,
                                           const int isect_prim)
 {
-  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+  if (isect_object != OBJECT_NONE) {
     const Transform tfm = object_get_inverse_transform(kg, sd);
 
     P = transform_point(&tfm, P);
@@ -695,12 +716,14 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
     D = safe_normalize_len(D, &t);
   }
 
-  KernelCurve kcurve = kernel_tex_fetch(__curves, isect_prim);
+  int prim = kernel_tex_fetch(__prim_index, isect_prim);
+  float4 v00 = kernel_tex_fetch(__curves, prim);
 
-  int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
+  int k0 = __float_as_int(v00.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
   int k1 = k0 + 1;
-  int ka = max(k0 - 1, kcurve.first_key);
-  int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 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);
 
   float4 P_curve[4];
 
@@ -757,13 +780,15 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
   sd->dPdv = cross(dPdu, sd->Ng);
 #  endif
 
-  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+  if (isect_object != OBJECT_NONE) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }
 
   sd->P = P;
-  sd->shader = kernel_tex_fetch(__curves, sd->prim).shader_id;
+
+  float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
+  sd->shader = __float_as_int(curvedata.z);
 }
 
 #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(__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));
+    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));
   }
   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 (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+  if (isect_object != OBJECT_NONE) {
     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 (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+  if (isect_object != OBJECT_NONE) {
     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 (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+  if (isect_object != OBJECT_NONE) {
     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 (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+  if (isect_object != OBJECT_NONE) {
     const Transform tfm = object_get_transform(kg, sd);
     P = transform_point(&tfm, P);
   }
@@ -186,9 +186,8 @@ ccl_device_inline bool motion_triangle_intersect(const KernelGlobals *kg,
       isect->t = t;
       isect->u = u;
       isect->v = v;
-      isect->prim = prim;
-      isect->object = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) :
-                                                object;
+      isect->prim = prim_addr;
+      isect->object = object;
       isect->type = PRIMITIVE_MOTION_TRIANGLE;
       return true;
     }
@@ -289,8 +288,8 @@ ccl_device_inline bool motion_triangle_intersect_local(const KernelGlobals *kg,
   isect->t = t;
   isect->u = u;
   isect->v = v;
-  isect->prim = prim;
-  isect->object = local_object;
+  isect->prim = prim_addr;
+  isect->object = 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 shader-data. */
+ * global memory as we write it to shaderdata. */
 ccl_device_inline void shader_setup_from_ray(const KernelGlobals *ccl_restrict kg,
                                              ShaderData *ccl_restrict sd,
                                              const Ray *ccl_restrict ray,
@@ -52,9 +52,10 @@ 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;
+  sd->object = (isect->object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, isect->prim) :
+                                                isect->object;
   sd->object_flag = kernel_tex_fetch(__object_flag, sd->object);
-  sd->prim = isect->prim;
+  sd->prim = kernel_tex_fetch(__prim_index, isect->prim);
   sd->lamp = LAMP_NONE;
   sd->flag = 0;
 
@@ -102,7 +103,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 (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+  if (isect->object != OBJECT_NONE) {
     /* 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(__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));
+  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));
 
   /* 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(__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));
+  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));
   /* 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(__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));
+  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));
 }
 
 /* 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(__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));
+  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));
   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(__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));
+  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));
 
   /* compute derivatives of P w.r.t. uv */
   *dPdu = (p0 - p2);

intern/cycles/kernel/geom/geom_triangle_intersect.h

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

intern/cycles/kernel/integrator/integrator_init_from_bake.h

@@ -109,17 +109,9 @@ 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, 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);
-  }
-
+  triangle_point_normal(kg, kernel_data.bake.object_index, prim, u, v, &P, &Ng, &shader);
   if (kernel_data.film.pass_background != PASS_UNUSED) {
     /* Environment baking. */
 
@@ -138,13 +130,8 @@ ccl_device bool integrator_init_from_bake(INTEGRATOR_STATE_ARGS,
   }
   else {
     /* Surface baking. */
-    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));
-    }
+    const float3 N = (shader & SHADER_SMOOTH_NORMAL) ? triangle_smooth_normal(kg, Ng, prim, u, v) :
+                                                       Ng;
 
     /* Setup ray. */
     Ray ray ccl_optional_struct_init;
@@ -156,12 +143,6 @@ 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.pass_background != PASS_UNUSED && !kernel_data.background.transparent) {
+    if (kernel_data.film.use_approximate_shadow_catcher && !kernel_data.background.transparent) {
       INTEGRATOR_STATE_WRITE(path, flag) |= PATH_RAY_SHADOW_CATCHER_BACKGROUND;
 
       if (use_raytrace_kernel) {
@@ -160,7 +160,10 @@ 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 float object_ao_distance = kernel_tex_fetch(__objects, last_isect_object).ao_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;
     if (object_ao_distance != 0.0f) {
       ray.t = object_ao_distance;
     }

intern/cycles/kernel/integrator/integrator_intersect_volume_stack.h

@@ -38,13 +38,10 @@ 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 * MAX_VOLUME_STACK_SIZE + 1];
+  Intersection hits[2 * 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;
 
@@ -58,7 +55,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);
@@ -94,15 +91,12 @@ 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 * MAX_VOLUME_STACK_SIZE + 1];
+  Intersection hits[2 * 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[MAX_VOLUME_STACK_SIZE];
+    int enclosed_volumes[VOLUME_STACK_SIZE];
     Intersection *isect = hits;
 
     qsort(hits, num_hits, sizeof(Intersection), intersections_compare);
@@ -127,7 +121,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;
@@ -142,12 +136,11 @@ ccl_device void integrator_intersect_volume_stack(INTEGRATOR_STATE_ARGS)
     }
   }
 #else
-  /* CUDA does not support definition of a variable size arrays, so use the maximum possible. */
-  int enclosed_volumes[MAX_VOLUME_STACK_SIZE];
+  int enclosed_volumes[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,8 +192,7 @@ 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 isect_type = INTEGRATOR_STATE(isect, type);
-    const int shader = intersection_get_shader_from_isect_prim(kg, isect_prim, isect_type);
+    const int shader = intersection_get_shader_from_isect_prim(kg, isect_prim);
     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,6 +59,8 @@ 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
 
@@ -83,14 +85,12 @@ 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,14 +114,12 @@ 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,9 +107,7 @@ 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,
-                        KERNEL_STRUCT_VOLUME_STACK_SIZE,
-                        KERNEL_STRUCT_VOLUME_STACK_SIZE)
+KERNEL_STRUCT_END_ARRAY(volume_stack, INTEGRATOR_VOLUME_STACK_SIZE, INTEGRATOR_VOLUME_STACK_SIZE)
 
 /********************************* Shadow Path State **************************/
 
@@ -165,5 +163,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,
-                        KERNEL_STRUCT_VOLUME_STACK_SIZE,
-                        KERNEL_STRUCT_VOLUME_STACK_SIZE)
+                        INTEGRATOR_VOLUME_STACK_SIZE,
+                        INTEGRATOR_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 < kernel_data.volume_stack_size; i++) {
+    for (int i = 0; i < INTEGRATOR_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,8 +223,6 @@ 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
@@ -232,7 +230,6 @@ 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 = ss_isect.hits[0].object;
+    const int object = intersection_get_object(kg, &ss_isect.hits[0]);
     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 >= kernel_data.volume_stack_size - 1) {
+    if (i >= 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 shader-data from stack. it's mostly setup already in
-     * shader_setup_from_volume, this switching should be quick. */
+    /* setup shaderdata 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,9 +18,11 @@
 #  define KERNEL_TEX(type, name)
 #endif
 
-/* BVH2, not used for OptiX or Embree. */
+/* bvh */
 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)
@@ -44,12 +46,10 @@ 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(KernelCurve, __curves)
+KERNEL_TEX(float4, __curves)
 KERNEL_TEX(float4, __curve_keys)
-KERNEL_TEX(KernelCurveSegment, __curve_segments)
 
 /* patches */
 KERNEL_TEX(uint, __patches)

intern/cycles/kernel/kernel_types.h

@@ -61,6 +61,8 @@ CCL_NAMESPACE_BEGIN
 #define ID_NONE (0.0f)
 #define PASS_UNUSED (~0)
 
+#define VOLUME_STACK_SIZE 4
+
 /* Kernel features */
 #define __SOBOL__
 #define __DPDU__
@@ -188,7 +190,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 {
   /* --------------------------------------------------------------------
@@ -358,6 +360,7 @@ typedef enum PassType {
   PASS_MATERIAL_ID,
   PASS_MOTION,
   PASS_MOTION_WEIGHT,
+  PASS_RENDER_TIME,
   PASS_CRYPTOMATTE,
   PASS_AOV_COLOR,
   PASS_AOV_VALUE,
@@ -606,12 +609,6 @@ 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
@@ -866,6 +863,9 @@ 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 volume_stack_size;
+  uint pad;
 
   KernelCamera cam;
   KernelFilm film;
@@ -1267,25 +1267,10 @@ typedef struct KernelObject {
 
   float ao_distance;
 
-  uint visibility;
-  int primitive_type;
+  float pad1, pad2;
 } 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,7 +110,6 @@ 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");
@@ -986,18 +985,8 @@ 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 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 {
+  else
     return false;
-  }
 }
 
 bool OSLRenderServices::get_background_attribute(const KernelGlobals *kg,

intern/cycles/kernel/osl/osl_services.h

@@ -297,7 +297,6 @@ 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:normal_map_normal", ninterp))) {
+        (!is_smooth || getattribute("geom:N", 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 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;
+    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;
 
     /* 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,20 +201,23 @@ 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, isect.hits[hit].prim, sd->time, verts);
+        motion_triangle_vertices(
+            kg, sd->object, kernel_tex_fetch(__prim_index, 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;
+      int object = (isect.hits[hit].object == OBJECT_NONE) ?
+                       kernel_tex_fetch(__prim_object, isect.hits[hit].prim) :
+                       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;
@@ -222,7 +225,7 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
 
       /* Compute smooth normal. */
       float3 N = hit_Ng;
-      int prim = isect.hits[hit].prim;
+      int prim = kernel_tex_fetch(__prim_index, 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,9 +85,6 @@ 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);
@@ -169,9 +166,6 @@ 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,6 +347,8 @@ 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,11 +97,6 @@ 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);
@@ -238,31 +233,13 @@ void BufferParams::update_offset_stride()
 
 bool BufferParams::modified(const BufferParams &other) const
 {
-  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) {
+  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)) {
     return true;
   }
 

intern/cycles/render/buffers.h

@@ -82,15 +82,6 @@ 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,6 +326,8 @@ 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,6 +46,12 @@ 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);
@@ -73,6 +79,7 @@ Geometry::Geometry(const NodeType *node_type, const Type type)
 
   bvh = NULL;
   attr_map_offset = 0;
+  optix_prim_offset = 0;
   prim_offset = 0;
 }
 
@@ -700,9 +707,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->curve_key_offset;
+        offset -= hair->curvekey_offset;
       else if (element == ATTR_ELEMENT_CURVE_KEY_MOTION)
-        offset -= hair->curve_key_offset;
+        offset -= hair->curvekey_offset;
     }
   }
   else {
@@ -965,22 +972,28 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout)
   size_t vert_size = 0;
   size_t tri_size = 0;
 
-  size_t curve_size = 0;
   size_t curve_key_size = 0;
-  size_t curve_segment_size = 0;
+  size_t curve_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) {
-    bool prim_offset_changed = false;
+    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;
+    }
 
     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;
 
@@ -1004,35 +1017,27 @@ 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);
 
-      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->curvekey_offset = curve_key_size;
       hair->prim_offset = curve_size;
 
-      curve_size += hair->num_curves();
       curve_key_size += hair->get_curve_keys().size();
-      curve_segment_size += hair->num_segments();
-    }
+      curve_size += hair->num_curves();
 
-    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;
+      hair->optix_prim_offset = optix_prim_size;
+      optix_prim_size += hair->num_segments();
     }
   }
 }
 
-void GeometryManager::device_update_mesh(Device *,
-                                         DeviceScene *dscene,
-                                         Scene *scene,
-                                         Progress &progress)
+void GeometryManager::device_update_mesh(
+    Device *, DeviceScene *dscene, Scene *scene, bool for_displacement, Progress &progress)
 {
   /* Count. */
   size_t vert_size = 0;
@@ -1040,7 +1045,6 @@ void GeometryManager::device_update_mesh(Device *,
 
   size_t curve_key_size = 0;
   size_t curve_size = 0;
-  size_t curve_segment_size = 0;
 
   size_t patch_size = 0;
 
@@ -1067,7 +1071,31 @@ void GeometryManager::device_update_mesh(Device *,
 
       curve_key_size += hair->get_curve_keys().size();
       curve_size += hair->num_curves();
-      curve_segment_size += hair->num_segments();
+    }
+  }
+
+  /* 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];
+      }
     }
   }
 
@@ -1076,7 +1104,6 @@ void GeometryManager::device_update_mesh(Device *,
     /* 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);
@@ -1102,12 +1129,13 @@ void GeometryManager::device_update_mesh(Device *,
           mesh->pack_normals(&vnormal[mesh->vert_offset]);
         }
 
-        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],
+        if (mesh->triangles_is_modified() || mesh->vert_patch_uv_is_modified() || copy_all_data) {
+          mesh->pack_verts(tri_prim_index,
                            &tri_vindex[mesh->prim_offset],
                            &tri_patch[mesh->prim_offset],
-                           &tri_patch_uv[mesh->vert_offset]);
+                           &tri_patch_uv[mesh->vert_offset],
+                           mesh->vert_offset,
+                           mesh->prim_offset);
         }
 
         if (progress.get_cancel())
@@ -1118,7 +1146,6 @@ void GeometryManager::device_update_mesh(Device *,
     /* 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();
@@ -1126,16 +1153,13 @@ void GeometryManager::device_update_mesh(Device *,
     dscene->tri_patch_uv.copy_to_device_if_modified();
   }
 
-  if (curve_segment_size != 0) {
-    progress.set_status("Updating Mesh", "Copying Curves to device");
+  if (curve_size != 0) {
+    progress.set_status("Updating Mesh", "Copying Strands to device");
 
     float4 *curve_keys = dscene->curve_keys.alloc(curve_key_size);
-    KernelCurve *curves = dscene->curves.alloc(curve_size);
-    KernelCurveSegment *curve_segments = dscene->curve_segments.alloc(curve_segment_size);
+    float4 *curves = dscene->curves.alloc(curve_size);
 
-    const bool copy_all_data = dscene->curve_keys.need_realloc() ||
-                               dscene->curves.need_realloc() ||
-                               dscene->curve_segments.need_realloc();
+    const bool copy_all_data = dscene->curve_keys.need_realloc() || dscene->curves.need_realloc();
 
     foreach (Geometry *geom, scene->geometry) {
       if (geom->is_hair()) {
@@ -1151,9 +1175,9 @@ void GeometryManager::device_update_mesh(Device *,
         }
 
         hair->pack_curves(scene,
-                          &curve_keys[hair->curve_key_offset],
+                          &curve_keys[hair->curvekey_offset],
                           &curves[hair->prim_offset],
-                          &curve_segments[hair->curve_segment_offset]);
+                          hair->curvekey_offset);
         if (progress.get_cancel())
           return;
       }
@@ -1161,7 +1185,6 @@ void GeometryManager::device_update_mesh(Device *,
 
     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()) {
@@ -1172,7 +1195,10 @@ void GeometryManager::device_update_mesh(Device *,
     foreach (Geometry *geom, scene->geometry) {
       if (geom->is_mesh()) {
         Mesh *mesh = static_cast<Mesh *>(geom);
-        mesh->pack_patches(&patch_data[mesh->patch_offset]);
+        mesh->pack_patches(&patch_data[mesh->patch_offset],
+                           mesh->vert_offset,
+                           mesh->face_offset,
+                           mesh->corner_offset);
 
         if (mesh->patch_table) {
           mesh->patch_table->copy_adjusting_offsets(&patch_data[mesh->patch_table_offset],
@@ -1186,6 +1212,23 @@ void GeometryManager::device_update_mesh(Device *,
 
     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,
@@ -1213,6 +1256,16 @@ 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);
@@ -1231,7 +1284,77 @@ 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 */
@@ -1252,23 +1375,31 @@ void GeometryManager::device_update_bvh(Device *device,
     dscene->object_node.steal_data(pack.object_node);
     dscene->object_node.copy_to_device();
   }
-  if (pack.prim_type.size()) {
+  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)) {
     dscene->prim_type.steal_data(pack.prim_type);
     dscene->prim_type.copy_to_device();
   }
-  if (pack.prim_visibility.size()) {
+  if (pack.prim_visibility.size() && (dscene->prim_visibility.is_modified() || has_bvh2_layout)) {
     dscene->prim_visibility.steal_data(pack.prim_visibility);
     dscene->prim_visibility.copy_to_device();
   }
-  if (pack.prim_index.size()) {
+  if (pack.prim_index.size() && (dscene->prim_index.need_realloc() || has_bvh2_layout)) {
     dscene->prim_index.steal_data(pack.prim_index);
     dscene->prim_index.copy_to_device();
   }
-  if (pack.prim_object.size()) {
+  if (pack.prim_object.size() && (dscene->prim_object.need_realloc() || has_bvh2_layout)) {
     dscene->prim_object.steal_data(pack.prim_object);
     dscene->prim_object.copy_to_device();
   }
-  if (pack.prim_time.size()) {
+  if (pack.prim_time.size() && (dscene->prim_time.need_realloc() || has_bvh2_layout)) {
     dscene->prim_time.steal_data(pack.prim_time);
     dscene->prim_time.copy_to_device();
   }
@@ -1498,6 +1629,8 @@ 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();
@@ -1516,7 +1649,6 @@ 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();
     }
   }
 
@@ -1559,7 +1691,6 @@ 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();
   }
@@ -1567,7 +1698,6 @@ 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;
@@ -1776,7 +1906,7 @@ void GeometryManager::device_update(Device *device,
             {"device_update (displacement: copy meshes to device)", time});
       }
     });
-    device_update_mesh(device, dscene, scene, progress);
+    device_update_mesh(device, dscene, scene, true, progress);
   }
   if (progress.get_cancel()) {
     return;
@@ -1928,7 +2058,7 @@ void GeometryManager::device_update(Device *device,
             {"device_update (copy meshes to device)", time});
       }
     });
-    device_update_mesh(device, dscene, scene, progress);
+    device_update_mesh(device, dscene, scene, false, progress);
     if (progress.get_cancel()) {
       return;
     }
@@ -1961,12 +2091,13 @@ 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();
@@ -1974,7 +2105,6 @@ 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();
@@ -1988,12 +2118,13 @@ 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);
@@ -2001,7 +2132,6 @@ 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,6 +43,24 @@ 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. */
@@ -82,6 +100,7 @@ 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(). */
@@ -125,7 +144,10 @@ class Geometry : public Node {
                    int n,
                    int total);
 
-  virtual PrimitiveType primitive_type() const = 0;
+  virtual void pack_primitives(PackedBVH *pack,
+                               int object,
+                               uint visibility,
+                               PackFlags pack_flags) = 0;
 
   /* Check whether the geometry should have own BVH built separately. Briefly,
    * own BVH is needed for geometry, if:
@@ -238,7 +260,11 @@ class GeometryManager {
 
   void device_update_object(Device *device, DeviceScene *dscene, Scene *scene, Progress &progress);
 
-  void device_update_mesh(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_attributes(Device *device,
                                 DeviceScene *dscene,

intern/cycles/render/graph.cpp

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

intern/cycles/render/hair.cpp

@@ -295,8 +295,7 @@ NODE_DEFINE(Hair)
 
 Hair::Hair() : Geometry(get_node_type(), Geometry::HAIR)
 {
-  curve_key_offset = 0;
-  curve_segment_offset = 0;
+  curvekey_offset = 0;
   curve_shape = CURVE_RIBBON;
 }
 
@@ -463,8 +462,8 @@ void Hair::apply_transform(const Transform &tfm, const bool apply_to_motion)
 
 void Hair::pack_curves(Scene *scene,
                        float4 *curve_key_co,
-                       KernelCurve *curves,
-                       KernelCurveSegment *curve_segments)
+                       float4 *curve_data,
+                       size_t curvekey_offset)
 {
   size_t curve_keys_size = curve_keys.size();
 
@@ -478,10 +477,7 @@ 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);
@@ -491,24 +487,56 @@ void Hair::pack_curves(Scene *scene,
                          scene->default_surface;
     shader_id = scene->shader_manager->get_shader_id(shader, false);
 
-    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;
+    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);
+  }
+}
 
-    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)
+{
+  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;
+      }
     }
   }
 }
 
-PrimitiveType Hair::primitive_type() const
-{
-  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,8 +21,6 @@
 
 CCL_NAMESPACE_BEGIN
 
-struct KernelCurveSegment;
-
 class Hair : public Geometry {
  public:
   NODE_DECLARE
@@ -97,8 +95,7 @@ class Hair : public Geometry {
   NODE_SOCKET_API_ARRAY(array<int>, curve_shader)
 
   /* BVH */
-  size_t curve_key_offset;
-  size_t curve_segment_offset;
+  size_t curvekey_offset;
   CurveShapeType curve_shape;
 
   /* Constructor/Destructor */
@@ -147,12 +144,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,
-                   KernelCurve *curve,
-                   KernelCurveSegment *curve_segments);
+  void pack_curves(Scene *scene, float4 *curve_key_co, float4 *curve_data, size_t curvekey_offset);
 
-  PrimitiveType primitive_type() const override;
+  void pack_primitives(PackedBVH *pack,
+                       int object,
+                       uint visibility,
+                       PackFlags pack_flags) override;
 };
 
 CCL_NAMESPACE_END

intern/cycles/render/mesh.cpp

@@ -729,7 +729,12 @@ void Mesh::pack_normals(float4 *vnormal)
   }
 }
 
-void Mesh::pack_verts(float4 *tri_verts, uint4 *tri_vindex, uint *tri_patch, float2 *tri_patch_uv)
+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)
 {
   size_t verts_size = verts.size();
 
@@ -744,19 +749,17 @@ void Mesh::pack_verts(float4 *tri_verts, uint4 *tri_vindex, uint *tri_patch, flo
   size_t triangles_size = num_triangles();
 
   for (size_t i = 0; i < triangles_size; i++) {
-    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));
+    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]);
 
     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)
+void Mesh::pack_patches(uint *patch_data, uint vert_offset, uint face_offset, uint corner_offset)
 {
   size_t num_faces = get_num_subd_faces();
   int ngons = 0;
@@ -802,9 +805,53 @@ void Mesh::pack_patches(uint *patch_data)
   }
 }
 
-PrimitiveType Mesh::primitive_type() const
+void Mesh::pack_primitives(ccl::PackedBVH *pack, int object, uint visibility, PackFlags pack_flags)
 {
-  return has_motion_blur() ? PRIMITIVE_MOTION_TRIANGLE : PRIMITIVE_TRIANGLE;
+  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]]);
+    }
+  }
 }
 
 CCL_NAMESPACE_END

intern/cycles/render/mesh.h

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

intern/cycles/render/object.cpp

@@ -60,7 +60,6 @@ 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;
@@ -367,22 +366,6 @@ 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;
@@ -529,9 +512,6 @@ 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,13 +109,6 @@ 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,6 +89,7 @@ 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);
@@ -216,6 +217,10 @@ 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,12 +49,13 @@ 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),
@@ -62,7 +63,6 @@ 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,8 +527,6 @@ 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)
@@ -644,33 +642,6 @@ 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,6 +74,8 @@ 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;
@@ -81,16 +83,14 @@ 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<KernelCurve> curves;
+  device_vector<float4> curves;
   device_vector<float4> curve_keys;
-  device_vector<KernelCurveSegment> curve_segments;
 
   device_vector<uint> patches;
 
@@ -344,9 +344,6 @@ 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,13 +157,6 @@ 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
@@ -179,9 +172,10 @@ 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()) {
-        progress.set_error(device->error_message());
+        const string &error_message = device->error_message();
+        progress.set_error(error_message);
+        progress.set_cancel(error_message);
         break;
       }
     }
@@ -286,20 +280,12 @@ 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,17 +372,8 @@ 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()
@@ -408,25 +399,18 @@ 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_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;
+  const int tile_y = index / tile_state_.num_tiles_x;
+  const int tile_x = index - tile_y * tile_state_.num_tiles_x;
 
   Tile tile;
 
-  tile.x = max(0, tile_window_x - overscan_);
-  tile.y = max(0, tile_window_y - overscan_);
+  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.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_);
+  tile.width = min(tile.width, buffer_params_.width - tile.x);
+  tile.height = min(tile.height, buffer_params_.height - tile.y);
 
   return tile;
 }
@@ -499,22 +483,11 @@ 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 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;
+  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;
 
   VLOG(3) << "Write tile at " << tile_x << ", " << tile_y;
 
@@ -526,16 +499,13 @@ 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.window_width,
+                                          tile_x + tile_params.width,
                                           tile_y,
-                                          tile_y + tile_params.window_height,
+                                          tile_y + tile_params.height,
                                           0,
                                           1,
                                           TypeDesc::FLOAT,
-                                          pixels,
-                                          xstride,
-                                          ystride,
-                                          zstride)) {
+                                          pixels)) {
     LOG(ERROR) << "Error writing tile " << write_state_.tile_out->geterror();
     return false;
   }
@@ -561,15 +531,12 @@ void TileManager::finish_write_tiles()
          ++tile_index) {
       const Tile tile = get_tile_for_index(tile_index);
 
-      const int tile_x = tile.x + tile.window_x;
-      const int tile_y = tile.y + tile.window_y;
+      VLOG(3) << "Write dummy tile at " << tile.x << ", " << tile.y;
 
-      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,
+      write_state_.tile_out->write_tiles(tile.x,
+                                         tile.x + tile.width,
+                                         tile.y,
+                                         tile.y + tile.height,
                                          0,
                                          1,
                                          TypeDesc::FLOAT,

intern/cycles/render/tile.h

@@ -35,9 +35,6 @@ 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()
   {
   }
@@ -81,11 +78,6 @@ class TileManager {
     return tile_state_.num_tiles > 1;
   }
 
-  inline int get_tile_overscan() const
-  {
-    return overscan_;
-  }
-
   bool next();
   bool done();
 
@@ -136,9 +128,6 @@ 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,6 +473,7 @@ 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
@@ -482,13 +483,19 @@ 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,6 +1140,30 @@ 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 {
+typedef struct GHOST_XrPose {
   float position[3];
   /* Blender convention (w, x, y, z) */
   float orientation_quat[4];
@@ -753,8 +753,31 @@ 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,4 +1069,39 @@ 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,7 +22,6 @@
 #include <cstring>
 
 #include "GHOST_Types.h"
-
 #include "GHOST_XrException.h"
 #include "GHOST_Xr_intern.h"
 

intern/ghost/intern/GHOST_XrAction.h

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

intern/ghost/intern/GHOST_XrContext.cpp

@@ -412,11 +412,14 @@ void GHOST_XrContext::getExtensionsToEnable(
     try_ext.push_back(XR_EXT_DEBUG_UTILS_EXTENSION_NAME);
   }
 
-  /* Try enabling interaction profile extensions. */
+  /* 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

@@ -0,0 +1,629 @@
+/*
+ * 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;
+  }
+}
+
+/** \} */