PointCloud: Initial rendering support for Workbench

Also includes outline overlays. Removes the temp overlay drawing

We make the geometry follow camera like billboards this uses less
geometry. Currently we use half octahedron for now. Goal would be
to use icospheres.

This patch also optimize the case when pointcloud has uniform radius.
However we should premultiply the radius prop by the default radius
beforehand to avoid a multiplication on CPU.

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

.clang-tidy

@@ -0,0 +1,43 @@
+Checks:  >
+  -*,
+  readability-*,
+  -readability-uppercase-literal-suffix,
+  -readability-magic-numbers,
+  -readability-isolate-declaration,
+  -readability-convert-member-functions-to-static,
+  -readability-implicit-bool-conversion,
+  -readability-avoid-const-params-in-decls,
+  -readability-simplify-boolean-expr,
+  -readability-make-member-function-const,
+
+  -readability-misleading-indentation,
+
+  -readability-else-after-return,
+  -readability-inconsistent-declaration-parameter-name,
+  -readability-redundant-preprocessor,
+  -readability-function-size,
+  -readability-function-size,
+  -readability-redundant-string-init,
+  -readability-redundant-member-init,
+  -readability-const-return-type,
+  -readability-static-accessed-through-instance,
+  -readability-redundant-declaration,
+  -readability-qualified-auto,
+
+  bugprone-*,
+  -bugprone-narrowing-conversions,
+  -bugprone-unhandled-self-assignment,
+  -bugprone-branch-clone,
+  -bugprone-macro-parentheses,
+
+  -bugprone-sizeof-expression,
+  -bugprone-integer-division,
+  -bugprone-incorrect-roundings,
+  -bugprone-suspicious-string-compare,
+  -bugprone-not-null-terminated-result,
+  -bugprone-suspicious-missing-comma,
+  -bugprone-parent-virtual-call,
+  -bugprone-infinite-loop,
+  -bugprone-copy-constructor-init,
+
+WarningsAsErrors: '*'

CMakeLists.txt

@@ -415,6 +415,11 @@ mark_as_advanced(WITH_CXX_GUARDEDALLOC)
 option(WITH_ASSERT_ABORT "Call abort() when raising an assertion through BLI_assert()" ON)
 mark_as_advanced(WITH_ASSERT_ABORT)
 
+if(UNIX AND NOT APPLE)
+  option(WITH_CLANG_TIDY "Use Clang Tidy to analyze the source code (only enable for development on Linux using Clang)" OFF)
+  mark_as_advanced(WITH_CLANG_TIDY)
+endif()
+
 option(WITH_BOOST "Enable features depending on boost" ON)
 option(WITH_TBB   "Enable features depending on TBB (OpenVDB, OpenImageDenoise, sculpt multithreading)" ON)
 

build_files/build_environment/CMakeLists.txt

@@ -30,7 +30,7 @@
 #   build_deps 2015 x64 / build_deps 2015 x86
 #
 # MAC OS X USAGE:
-#   Install with homebrew: brew install cmake autoconf automake libtool yasm nasm
+#   Install with homebrew: brew install cmake autoconf automake libtool yasm nasm bison
 #   Run "make deps" from main Blender directory
 #
 # LINUX USAGE:
@@ -76,6 +76,7 @@ include(cmake/llvm.cmake)
 include(cmake/clang.cmake)
 if(APPLE)
   include(cmake/openmp.cmake)
+  include(cmake/nasm.cmake)
 endif()
 include(cmake/openimageio.cmake)
 include(cmake/tiff.cmake)

build_files/build_environment/cmake/boost.cmake

@@ -44,7 +44,7 @@ if(WIN32)
 elseif(APPLE)
   set(BOOST_CONFIGURE_COMMAND ./bootstrap.sh)
   set(BOOST_BUILD_COMMAND ./b2)
-  set(BOOST_BUILD_OPTIONS toolset=darwin cxxflags=${PLATFORM_CXXFLAGS} linkflags=${PLATFORM_LDFLAGS} visibility=global --disable-icu boost.locale.icu=off)
+  set(BOOST_BUILD_OPTIONS toolset=clang-darwin cxxflags=${PLATFORM_CXXFLAGS} linkflags=${PLATFORM_LDFLAGS} visibility=global --disable-icu boost.locale.icu=off)
   set(BOOST_HARVEST_CMD echo .)
   set(BOOST_PATCH_COMMAND echo .)
 else()

build_files/build_environment/cmake/check_software.cmake

@@ -30,6 +30,7 @@ if(UNIX)
     nasm
     yasm
     tclsh
+    bison
   )
 
   foreach(_software ${_required_software})
@@ -40,6 +41,12 @@ if(UNIX)
     unset(_software_find CACHE)
   endforeach()
 
+  if(APPLE)
+    if(NOT EXISTS "/usr/local/opt/bison/bin/bison")
+      set(_software_missing "${_software_missing} bison")
+    endif()
+  endif()
+
   if(_software_missing)
     message(
       "\n"
@@ -50,7 +57,7 @@ if(UNIX)
       "  apt install autoconf automake libtool yasm nasm tcl\n"
       "\n"
       "On macOS (with homebrew):\n"
-      "  brew install cmake autoconf automake libtool yasm nasm\n"
+      "  brew install cmake autoconf automake libtool yasm nasm bison\n"
       "\n"
       "Other platforms:\n"
       "  Install equivalent packages.\n")

build_files/build_environment/cmake/ffmpeg.cmake

@@ -50,7 +50,8 @@ if(APPLE)
   set(FFMPEG_EXTRA_FLAGS
     ${FFMPEG_EXTRA_FLAGS}
     --target-os=darwin
-    )
+    --x86asmexe=${LIBDIR}/nasm/bin/nasm
+  )
 endif()
 
 ExternalProject_Add(external_ffmpeg
@@ -143,6 +144,12 @@ if(WIN32)
     external_zlib_mingw
   )
 endif()
+if(APPLE)
+  add_dependencies(
+    external_ffmpeg
+    external_nasm
+  )
+endif()
 
 if(BUILD_MODE STREQUAL Release AND WIN32)
   ExternalProject_Add_Step(external_ffmpeg after_install

build_files/build_environment/cmake/freetype.cmake

@@ -24,7 +24,8 @@ set(FREETYPE_EXTRA_ARGS
   -DFT_WITH_HARFBUZZ=OFF
   -DFT_WITH_BZIP2=OFF
   -DCMAKE_DISABLE_FIND_PACKAGE_HarfBuzz=TRUE
-  -DCMAKE_DISABLE_FIND_PACKAGE_BZip2=TRUE)
+  -DCMAKE_DISABLE_FIND_PACKAGE_BZip2=TRUE
+  -DCMAKE_DISABLE_FIND_PACKAGE_BrotliDec=TRUE)
 
 ExternalProject_Add(external_freetype
   URL ${FREETYPE_URI}

build_files/build_environment/cmake/ispc.cmake

@@ -22,6 +22,17 @@ if(WIN32)
     -DBISON_EXECUTABLE=${LIBDIR}/flexbison/win_bison.exe
     -DM4_EXECUTABLE=${DOWNLOAD_DIR}/mingw/mingw64/msys/1.0/bin/m4.exe
   )
+elseif(APPLE)
+  # Use bison installed via Homebrew.
+  # The one which comes which Xcode toolset is too old.
+  set(ISPC_EXTRA_ARGS_APPLE
+    -DBISON_EXECUTABLE=/usr/local/opt/bison/bin/bison
+  )
+elseif(UNIX)
+  set(ISPC_EXTRA_ARGS_UNIX
+    -DCMAKE_C_COMPILER=${LIBDIR}/clang/bin/clang
+    -DCMAKE_CXX_COMPILER=${LIBDIR}/clang/bin/clang++
+  )
 endif()
 
 set(ISPC_EXTRA_ARGS
@@ -36,6 +47,8 @@ set(ISPC_EXTRA_ARGS
     -DCLANG_LIBRARY_DIR=${LIBDIR}/clang/lib
     -DCLANG_INCLUDE_DIRS=${LIBDIR}/clang/include
     ${ISPC_EXTRA_ARGS_WIN}
+    ${ISPC_EXTRA_ARGS_APPLE}
+    ${ISPC_EXTRA_ARGS_UNIX}
 )
 
 ExternalProject_Add(external_ispc

build_files/build_environment/cmake/nasm.cmake

@@ -0,0 +1,29 @@
+# ***** BEGIN GPL LICENSE BLOCK *****
+#
+# This program is free software; you can redistribute it and/or
+# modify it under the terms of the GNU General Public License
+# as published by the Free Software Foundation; either version 2
+# of the License, or (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program; if not, write to the Free Software Foundation,
+# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+#
+# ***** END GPL LICENSE BLOCK *****
+
+ExternalProject_Add(external_nasm
+  URL ${NASM_URI}
+  DOWNLOAD_DIR ${DOWNLOAD_DIR}
+  URL_HASH SHA256=${NASM_HASH}
+  PREFIX ${BUILD_DIR}/nasm
+  PATCH_COMMAND ${PATCH_CMD} --verbose -p 1 -N -d ${BUILD_DIR}/nasm/src/external_nasm < ${PATCH_DIR}/nasm.diff
+  CONFIGURE_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/nasm/src/external_nasm/ && ${CONFIGURE_COMMAND} --prefix=${LIBDIR}/nasm
+  BUILD_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/nasm/src/external_nasm/ && make -j${MAKE_THREADS}
+  INSTALL_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/nasm/src/external_nasm/ && make install
+  INSTALL_DIR ${LIBDIR}/nasm
+)

build_files/build_environment/cmake/ogg.cmake

@@ -21,6 +21,7 @@ ExternalProject_Add(external_ogg
   DOWNLOAD_DIR ${DOWNLOAD_DIR}
   URL_HASH SHA256=${OGG_HASH}
   PREFIX ${BUILD_DIR}/ogg
+  PATCH_COMMAND ${PATCH_CMD} --verbose -p 1 -N -d ${BUILD_DIR}/ogg/src/external_ogg < ${PATCH_DIR}/ogg.diff
   CONFIGURE_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/ogg/src/external_ogg/ && ${CONFIGURE_COMMAND} --prefix=${LIBDIR}/ogg --disable-shared --enable-static
   BUILD_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/ogg/src/external_ogg/ && make -j${MAKE_THREADS}
   INSTALL_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/ogg/src/external_ogg/ && make install

build_files/build_environment/cmake/options.cmake

@@ -113,14 +113,18 @@ else()
       COMMAND xcode-select --print-path
       OUTPUT_VARIABLE XCODE_DEV_PATH OUTPUT_STRIP_TRAILING_WHITESPACE
     )
+    execute_process(
+      COMMAND xcodebuild -version -sdk macosx SDKVersion
+      OUTPUT_VARIABLE MACOSX_SDK_VERSION OUTPUT_STRIP_TRAILING_WHITESPACE)
+
     set(OSX_ARCHITECTURES x86_64)
-    set(OSX_DEPLOYMENT_TARGET 10.11)
+    set(OSX_DEPLOYMENT_TARGET 10.13)
     set(OSX_SYSROOT ${XCODE_DEV_PATH}/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk)
 
     set(PLATFORM_CFLAGS "-isysroot ${OSX_SYSROOT} -mmacosx-version-min=${OSX_DEPLOYMENT_TARGET}")
     set(PLATFORM_CXXFLAGS "-isysroot ${OSX_SYSROOT} -mmacosx-version-min=${OSX_DEPLOYMENT_TARGET} -std=c++11 -stdlib=libc++")
     set(PLATFORM_LDFLAGS "-isysroot ${OSX_SYSROOT} -mmacosx-version-min=${OSX_DEPLOYMENT_TARGET}")
-    set(PLATFORM_BUILD_TARGET --build=x86_64-apple-darwin15.0.0) # OS X 10.11
+    set(PLATFORM_BUILD_TARGET --build=x86_64-apple-darwin17.0.0) # OS X 10.13
     set(PLATFORM_CMAKE_FLAGS
       -DCMAKE_OSX_ARCHITECTURES:STRING=${OSX_ARCHITECTURES}
       -DCMAKE_OSX_DEPLOYMENT_TARGET:STRING=${OSX_DEPLOYMENT_TARGET}
@@ -155,6 +159,7 @@ else()
 
   set(CONFIGURE_ENV
     export MACOSX_DEPLOYMENT_TARGET=${OSX_DEPLOYMENT_TARGET} &&
+    export MACOSX_SDK_VERSION=${OSX_DEPLOYMENT_TARGET} &&
     export CFLAGS=${PLATFORM_CFLAGS} &&
     export CXXFLAGS=${PLATFORM_CXXFLAGS} &&
     export LDFLAGS=${PLATFORM_LDFLAGS}

build_files/build_environment/cmake/python.cmake

@@ -48,7 +48,12 @@ if(WIN32)
 
 else()
   if(APPLE)
-    # disable functions that can be in 10.13 sdk but aren't available on 10.9 target
+    # Disable functions that can be in 10.13 sdk but aren't available on 10.9 target.
+    #
+    # Disable libintl (gettext library) as it might come from Homebrew, which makes
+    # it so test program compiles, but the Python does not. This is because for Python
+    # we use isysroot, which seems to forbid using libintl.h.
+    # The gettext functionality seems to come from CoreFoundation, so should be all fine.
     set(PYTHON_FUNC_CONFIGS
       export ac_cv_func_futimens=no &&
       export ac_cv_func_utimensat=no &&
@@ -60,7 +65,10 @@ else()
       export ac_cv_func_getentropy=no &&
       export ac_cv_func_mkostemp=no &&
       export ac_cv_func_mkostemps=no &&
-      export ac_cv_func_timingsafe_bcmp=no)
+      export ac_cv_func_timingsafe_bcmp=no &&
+      export ac_cv_header_libintl_h=no &&
+      export ac_cv_lib_intl_textdomain=no
+    )
     set(PYTHON_CONFIGURE_ENV ${CONFIGURE_ENV} && ${PYTHON_FUNC_CONFIGS})
     set(PYTHON_BINARY ${BUILD_DIR}/python/src/external_python/python.exe)
   else()

build_files/build_environment/cmake/tiff.cmake

@@ -16,6 +16,12 @@
 #
 # ***** END GPL LICENSE BLOCK *****
 
+if(WITH_WEBP)
+  set(WITH_TIFF_WEBP ON)
+else()
+  set(WITH_TIFF_WEBP OFF)
+endif()
+
 set(TIFF_EXTRA_ARGS
   -DZLIB_LIBRARY=${LIBDIR}/zlib/lib/${ZLIB_LIBRARY}
   -DZLIB_INCLUDE_DIR=${LIBDIR}/zlib/include
@@ -23,6 +29,8 @@ set(TIFF_EXTRA_ARGS
   -DBUILD_SHARED_LIBS=OFF
   -Dlzma=OFF
   -Djbig=OFF
+  -Dzstd=OFF
+  -Dwebp=${WITH_TIFF_WEBP}
 )
 
 ExternalProject_Add(external_tiff

build_files/build_environment/cmake/versions.cmake

@@ -305,6 +305,10 @@ set(MESA_VERSION 18.3.1)
 set(MESA_URI ftp://ftp.freedesktop.org/pub/mesa//mesa-${MESA_VERSION}.tar.xz)
 set(MESA_HASH d60828056d77bfdbae0970f9b15fb1be)
 
+set(NASM_VERSION 2.15.02)
+set(NASM_URI https://www.nasm.us/pub/nasm/releasebuilds/${NASM_VERSION}/nasm-${NASM_VERSION}.tar.xz)
+set(NASM_HASH f4fd1329b1713e1ccd34b2fc121c4bcd278c9f91cc4cb205ae8fcd2e4728dd14)
+
 set(XR_OPENXR_SDK_VERSION 1.0.8)
 set(XR_OPENXR_SDK_URI https://github.com/KhronosGroup/OpenXR-SDK/archive/release-${XR_OPENXR_SDK_VERSION}.tar.gz)
 set(XR_OPENXR_SDK_HASH c6de63d2e0f9029aa58dfa97cad8ce07)

build_files/build_environment/cmake/x264.cmake

@@ -21,12 +21,21 @@ if(WIN32)
 endif()
 
 
+if(APPLE)
+  set(X264_CONFIGURE_ENV
+    export AS=${LIBDIR}/nasm/bin/nasm
+  )
+else()
+  set(X264_CONFIGURE_ENV echo .)
+endif()
+
 ExternalProject_Add(external_x264
   URL ${X264_URI}
   DOWNLOAD_DIR ${DOWNLOAD_DIR}
   URL_HASH SHA256=${X264_HASH}
   PREFIX ${BUILD_DIR}/x264
-  CONFIGURE_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/x264/src/external_x264/ && ${CONFIGURE_COMMAND} --prefix=${LIBDIR}/x264
+  CONFIGURE_COMMAND ${CONFIGURE_ENV} && ${X264_CONFIGURE_ENV} && cd ${BUILD_DIR}/x264/src/external_x264/ &&
+    ${CONFIGURE_COMMAND} --prefix=${LIBDIR}/x264
     --enable-static
     --enable-pic
     --disable-lavf
@@ -39,3 +48,10 @@ ExternalProject_Add(external_x264
 if(MSVC)
   set_target_properties(external_x264 PROPERTIES FOLDER Mingw)
 endif()
+
+if(APPLE)
+  add_dependencies(
+    external_x264
+    external_nasm
+  )
+endif()

build_files/build_environment/patches/ffmpeg.diff

@@ -9,3 +9,62 @@
  enabled libopenmpt        && require_pkg_config libopenmpt "libopenmpt >= 0.2.6557" libopenmpt/libopenmpt.h openmpt_module_create -lstdc++ && append libopenmpt_extralibs "-lstdc++"
  enabled libopus           && {
      enabled libopus_decoder && {
+--- a/libavcodec/cfhddata.c
++++ b/libavcodec/cfhddata.c
+@@ -276,10 +276,10 @@
+ av_cold int ff_cfhd_init_vlcs(CFHDContext *s)
+ {
+     int i, j, ret = 0;
+-    uint32_t new_cfhd_vlc_bits[NB_VLC_TABLE_18 * 2];
+-    uint8_t  new_cfhd_vlc_len[NB_VLC_TABLE_18 * 2];
+-    uint16_t new_cfhd_vlc_run[NB_VLC_TABLE_18 * 2];
+-    int16_t  new_cfhd_vlc_level[NB_VLC_TABLE_18 * 2];
++    uint32_t *new_cfhd_vlc_bits = av_calloc(sizeof(uint32_t), NB_VLC_TABLE_18 * 2);
++    uint8_t  *new_cfhd_vlc_len = av_calloc(sizeof(uint8_t), NB_VLC_TABLE_18 * 2);
++    uint16_t *new_cfhd_vlc_run = av_calloc(sizeof(uint16_t), NB_VLC_TABLE_18 * 2);
++    int16_t  *new_cfhd_vlc_level = av_calloc(sizeof(int16_t), NB_VLC_TABLE_18 * 2);
+ 
+     /** Similar to dv.c, generate signed VLC tables **/
+ 
+@@ -305,8 +305,13 @@
+ 
+     ret = init_vlc(&s->vlc_9, VLC_BITS, j, new_cfhd_vlc_len,
+                    1, 1, new_cfhd_vlc_bits, 4, 4, 0);
+-    if (ret < 0)
++    if (ret < 0) {
++        av_free(new_cfhd_vlc_bits);
++        av_free(new_cfhd_vlc_len);
++        av_free(new_cfhd_vlc_run);
++        av_free(new_cfhd_vlc_level);
+         return ret;
++    }
+     for (i = 0; i < s->vlc_9.table_size; i++) {
+         int code = s->vlc_9.table[i][0];
+         int len  = s->vlc_9.table[i][1];
+@@ -346,8 +351,14 @@
+ 
+     ret = init_vlc(&s->vlc_18, VLC_BITS, j, new_cfhd_vlc_len,
+                    1, 1, new_cfhd_vlc_bits, 4, 4, 0);
+-    if (ret < 0)
++    if (ret < 0) {
++        av_free(new_cfhd_vlc_bits);
++        av_free(new_cfhd_vlc_len);
++        av_free(new_cfhd_vlc_run);
++        av_free(new_cfhd_vlc_level);
+         return ret;
++    }
++
+     av_assert0(s->vlc_18.table_size == 4572);
+ 
+     for (i = 0; i < s->vlc_18.table_size; i++) {
+@@ -367,5 +378,10 @@
+         s->table_18_rl_vlc[i].run   = run;
+     }
+ 
++    av_free(new_cfhd_vlc_bits);
++    av_free(new_cfhd_vlc_len);
++    av_free(new_cfhd_vlc_run);
++    av_free(new_cfhd_vlc_level);
++
+     return ret;
+ }

build_files/build_environment/patches/ispc.diff

@@ -34,3 +34,52 @@ diff -Naur orig/cmake/GenerateBuiltins.cmake.txt external_ispc/cmake/GenerateBui
      elseif ("${bit}" STREQUAL "64" AND ${arch} STREQUAL "x86")
          set(target_arch "x86_64")
      elseif ("${bit}" STREQUAL "32" AND ${arch} STREQUAL "arm")
+diff --git a/CMakeLists.txt b/CMakeLists.txt
+index 46a8db8..f53beef 100644
+--- a/CMakeLists.txt
++++ b/CMakeLists.txt
+@@ -36,8 +36,12 @@
+ cmake_minimum_required(VERSION 3.13)
+ 
+ if (UNIX)
+-    set(CMAKE_C_COMPILER "clang")
+-    set(CMAKE_CXX_COMPILER "clang++")
++    if (NOT CMAKE_C_COMPILER)
++        set(CMAKE_C_COMPILER "clang")
++    endif()
++    if (NOT CMAKE_CXX_COMPILER)
++        set(CMAKE_CXX_COMPILER "clang++")
++    endif()
+ endif()
+ 
+ set(PROJECT_NAME ispc)
+@@ -412,6 +416,29 @@ else()
+     endif()
+ endif()
+ 
++# Link against libstdc++.a which must be provided to the linker after
++# LLVM and CLang libraries.
++# This is needed because some of LLVM/CLang dependencies are using
++# std::make_shared, which is defined in one of those:
++#  - libclang-cpp.so
++#  - libstdc++.a
++# Using the former one is tricky because then generated binary depends
++# on a library which is outside of the LD_LIBRARY_PATH.
++#
++# Hence, using C++ implementation from G++ which seems to work just fine.
++# In fact, from investigation seems that libclang-cpp.so itself is pulling
++# std::_Sp_make_shared_tag from G++'s libstdc++.a.
++if(UNIX AND NOT APPLE)
++    execute_process(
++        COMMAND g++ --print-file-name libstdc++.a
++        OUTPUT_VARIABLE GCC_LIBSTDCXX_A
++        OUTPUT_STRIP_TRAILING_WHITESPACE
++    )
++    if(GCC_LIBSTDCXX_A AND EXISTS ${GCC_LIBSTDCXX_A})
++        target_link_libraries(${PROJECT_NAME} ${GCC_LIBSTDCXX_A})
++    endif()
++endif()
++
+ # Build target for utility checking host ISA
+ if (ISPC_INCLUDE_UTILS)
+     add_executable(check_isa "")

build_files/build_environment/patches/nasm.diff

@@ -0,0 +1,129 @@
+diff --git a/output/macho.h b/output/macho.h
+index 538c531e..fd5e8849 100644
+--- a/output/macho.h
++++ b/output/macho.h
+@@ -60,6 +60,8 @@
+ #define LC_SEGMENT			0x1
+ #define LC_SEGMENT_64			0x19
+ #define LC_SYMTAB			0x2
++#define LC_VERSION_MIN_MACOSX		0x24
++#define LC_BUILD_VERSION		0x32
+ 
+ /* Symbol type bits */
+ #define N_STAB				0xe0
+diff --git a/output/outmacho.c b/output/outmacho.c
+index 08147883..de6ec902 100644
+--- a/output/outmacho.c
++++ b/output/outmacho.c
+@@ -38,6 +38,8 @@
+ 
+ #include "compiler.h"
+ 
++#include <stdlib.h>
++
+ #include "nctype.h"
+ 
+ #include "nasm.h"
+@@ -64,6 +66,8 @@
+ #define MACHO_SYMCMD_SIZE		24
+ #define MACHO_NLIST_SIZE		12
+ #define MACHO_RELINFO_SIZE		8
++#define MACHO_BUILD_VERSION_SIZE	24
++#define MACHO_VERSION_MIN_MACOSX_SIZE	16
+ 
+ #define MACHO_HEADER64_SIZE		32
+ #define MACHO_SEGCMD64_SIZE		72
+@@ -1224,6 +1228,46 @@ static void macho_layout_symbols (uint32_t *numsyms,
+     }
+ }
+ 
++static bool get_full_version_from_env (const char *variable_name,
++                                       int *r_major,
++                                       int *r_minor,
++                                       int *r_patch)  {
++    *r_major = 0;
++    *r_minor = 0;
++    *r_patch = 0;
++
++    const char *value = getenv(variable_name);
++    if (value == NULL || value[0] == '\0') {
++        return false;
++    }
++
++    const char *current_value = value;
++    const char *end_value = value + strlen(value);
++
++    char *endptr;
++
++    *r_major = strtol(current_value, &endptr, 10);
++    if (endptr >= end_value) {
++        return true;
++    }
++    current_value = endptr + 1;
++
++    *r_minor = strtol(current_value, &endptr, 10);
++    if (endptr >= end_value) {
++        return true;
++    }
++    current_value = endptr + 1;
++
++    *r_patch = strtol(current_value, &endptr, 10);
++
++    return true;
++}
++
++static bool need_version_min_macosx_command (void) {
++    return getenv("MACOSX_DEPLOYMENT_TARGET") &&
++           getenv("MACOSX_SDK_VERSION");
++}
++
+ /* Calculate some values we'll need for writing later.  */
+ 
+ static void macho_calculate_sizes (void)
+@@ -1270,6 +1314,12 @@ static void macho_calculate_sizes (void)
+         head_sizeofcmds += fmt.segcmd_size  + seg_nsects * fmt.sectcmd_size;
+     }
+ 
++    /* LC_VERSION_MIN_MACOSX */
++    if (need_version_min_macosx_command()) {
++	++head_ncmds;
++	head_sizeofcmds += MACHO_VERSION_MIN_MACOSX_SIZE;
++    }
++
+     if (nsyms > 0) {
+ 	++head_ncmds;
+ 	head_sizeofcmds += MACHO_SYMCMD_SIZE;
+@@ -1653,6 +1703,33 @@ static void macho_write (void)
+     else
+         nasm_warn(WARN_OTHER, "no sections?");
+ 
++#define ENCODE_BUILD_VERSION(major, minor, patch) \
++    (((major) << 16) | ((minor) << 8) | (patch))
++
++    if (0) {
++        fwriteint32_t(LC_BUILD_VERSION, ofile); /* cmd == LC_BUILD_VERSION */
++        fwriteint32_t(MACHO_BUILD_VERSION_SIZE, ofile); /* size of load command */
++        fwriteint32_t(1, ofile);    /* platform */
++        fwriteint32_t(ENCODE_BUILD_VERSION(10, 13, 0), ofile);    /* minos, X.Y.Z is encoded in nibbles xxxx.yy.zz */
++        fwriteint32_t(ENCODE_BUILD_VERSION(10, 15, 4), ofile);    /* sdk, X.Y.Z is encoded in nibbles xxxx.yy.zz */
++        fwriteint32_t(0, ofile);    /* number of tool entries following this */
++    }
++
++    if (need_version_min_macosx_command()) {
++        int sdk_major, sdk_minor, sdk_patch;
++        get_full_version_from_env("MACOSX_SDK_VERSION", &sdk_major, &sdk_minor, &sdk_patch);
++
++        int version_major, version_minor, version_patch;
++        get_full_version_from_env("MACOSX_DEPLOYMENT_TARGET", &version_major, &version_minor, &version_patch);
++
++        fwriteint32_t(LC_VERSION_MIN_MACOSX, ofile); /* cmd == LC_VERSION_MIN_MACOSX */
++        fwriteint32_t(MACHO_VERSION_MIN_MACOSX_SIZE, ofile); /* size of load command */
++        fwriteint32_t(ENCODE_BUILD_VERSION(version_major, version_minor, version_patch), ofile);    /* minos, X.Y.Z is encoded in nibbles xxxx.yy.zz */
++        fwriteint32_t(ENCODE_BUILD_VERSION(sdk_major, sdk_minor, sdk_patch), ofile);    /* sdk, X.Y.Z is encoded in nibbles xxxx.yy.zz */
++    }
++
++#undef ENCODE_BUILD_VERSION
++
+     if (nsyms > 0) {
+         /* write out symbol command */
+         fwriteint32_t(LC_SYMTAB, ofile); /* cmd == LC_SYMTAB */

build_files/build_environment/patches/ogg.diff

@@ -0,0 +1,12 @@
+diff --git a/include/ogg/os_types.h b/include/ogg/os_types.h
+index eb8a322..6f73b72 100644
+--- a/include/ogg/os_types.h
++++ b/include/ogg/os_types.h
+@@ -71,6 +71,7 @@
+ #elif (defined(__APPLE__) && defined(__MACH__)) /* MacOS X Framework build */
+ 
+ #  include <sys/types.h>
++#  include <stdint.h>
+    typedef int16_t ogg_int16_t;
+    typedef uint16_t ogg_uint16_t;
+    typedef int32_t ogg_int32_t;

build_files/cmake/Modules/FindClangTidy.cmake

@@ -0,0 +1,104 @@
+# - Find clang-tidy executable
+#
+# Find the native clang-tidy executable
+#
+# This module defines
+#  CLANG_TIDY_EXECUTABLE, the ful lpath to clang-tidy executable
+#
+#  CLANG_TIDY_VERSION, the full version of the clang-tidy in the
+#                      major,minor.patch format
+#
+# CLANG_TIDY_VERSION_MAJOR,
+# CLANG_TIDY_VERSION_MINOR,
+# CLANG_TIDY_VERSION_PATCH, individual components of the clang-tidy version.
+#
+#  CLANG_TIDY_FOUND, If false, do not try to use Eigen3.
+
+#=============================================================================
+# Copyright 2020 Blender Foundation.
+#
+# Distributed under the OSI-approved BSD License (the "License");
+# see accompanying file Copyright.txt for details.
+#
+# This software is distributed WITHOUT ANY WARRANTY; without even the
+# implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+# See the License for more information.
+#=============================================================================
+
+# If CLANG_TIDY_ROOT_DIR was defined in the environment, use it.
+if(NOT CLANG_TIDY_ROOT_DIR AND NOT $ENV{CLANG_TIDY_ROOT_DIR} STREQUAL "")
+  set(CLANG_TIDY_ROOT_DIR $ENV{CLANG_TIDY_ROOT_DIR})
+endif()
+
+set(_clang_tidy_SEARCH_DIRS
+  ${CLANG_TIDY_ROOT_DIR}
+  /usr/local/bin
+)
+
+# TODO(sergey): Find more reliable way of finding the latest clang-tidy.
+find_program(CLANG_TIDY_EXECUTABLE
+  NAMES
+    clang-tidy-10
+    clang-tidy-9
+    clang-tidy-8
+    clang-tidy-7
+    clang-tidy
+  HINTS
+    ${_clang_tidy_SEARCH_DIRS}
+)
+
+if(CLANG_TIDY_EXECUTABLE)
+  # Mark clang-tidy as found.
+  set(CLANG_TIDY_FOUND TRUE)
+
+  # Setup fallback values.
+  set(CLANG_TIDY_VERSION_MAJOR 0)
+  set(CLANG_TIDY_VERSION_MINOR 0)
+  set(CLANG_TIDY_VERSION_PATCH 0)
+
+  # Get version from the output.
+  #
+  # NOTE: Don't use name of the executable file since that only includes a
+  # major version. Also, even the major version might be missing in the
+  # executable name.
+  execute_process(COMMAND ${CLANG_TIDY_EXECUTABLE} -version
+                  OUTPUT_VARIABLE CLANG_TIDY_VERSION_RAW
+                  ERROR_QUIET
+                  OUTPUT_STRIP_TRAILING_WHITESPACE)
+
+  # Parse parts.
+  if(CLANG_TIDY_VERSION_RAW MATCHES "LLVM version .*")
+    # Strip the LLVM prefix and get list of individual version components.
+    string(REGEX REPLACE
+           ".*LLVM version ([.0-9]+).*" "\\1"
+           CLANG_SEMANTIC_VERSION "${CLANG_TIDY_VERSION_RAW}")
+    string(REPLACE "." ";" CLANG_VERSION_PARTS "${CLANG_SEMANTIC_VERSION}")
+    list(LENGTH CLANG_VERSION_PARTS NUM_CLANG_TIDY_VERSION_PARTS)
+
+    # Extract components into corresponding variables.
+    if(NUM_CLANG_TIDY_VERSION_PARTS GREATER 0)
+      list(GET CLANG_VERSION_PARTS 0 CLANG_TIDY_VERSION_MAJOR)
+    endif()
+    if(NUM_CLANG_TIDY_VERSION_PARTS GREATER 1)
+      list(GET CLANG_VERSION_PARTS 1 CLANG_TIDY_VERSION_MINOR)
+    endif()
+    if(NUM_CLANG_TIDY_VERSION_PARTS GREATER 2)
+      list(GET CLANG_VERSION_PARTS 2 CLANG_TIDY_VERSION_PATCH)
+    endif()
+
+    # Unset temp variables.
+    unset(NUM_CLANG_TIDY_VERSION_PARTS)
+    unset(CLANG_SEMANTIC_VERSION)
+    unset(CLANG_VERSION_PARTS)
+  endif()
+
+  # Construct full semantic version.
+  set(CLANG_TIDY_VERSION "${CLANG_TIDY_VERSION_MAJOR}.\
+${CLANG_TIDY_VERSION_MINOR}.\
+${CLANG_TIDY_VERSION_PATCH}")
+  unset(CLANG_TIDY_VERSION_RAW)
+
+  message(STATUS "Found clang-tidy ${CLANG_TIDY_EXECUTABLE} (${CLANG_TIDY_VERSION})")
+else()
+  set(CLANG_TIDY_FOUND FALSE)
+endif()

build_files/cmake/Modules/FindEmbree.cmake

@@ -82,7 +82,7 @@ FIND_LIBRARY(EMBREE_LIBRARY
 # handle the QUIETLY and REQUIRED arguments and set EMBREE_FOUND to TRUE if
 # all listed variables are TRUE
 INCLUDE(FindPackageHandleStandardArgs)
-FIND_PACKAGE_HANDLE_STANDARD_ARGS(EMBREE DEFAULT_MSG
+FIND_PACKAGE_HANDLE_STANDARD_ARGS(Embree DEFAULT_MSG
     _embree_LIBRARIES EMBREE_INCLUDE_DIR)
 
 IF(EMBREE_FOUND)

build_files/cmake/platform/platform_apple.cmake

@@ -20,7 +20,7 @@
 
 # Libraries configuration for Apple.
 
-set(MACOSX_DEPLOYMENT_TARGET "10.11")
+set(MACOSX_DEPLOYMENT_TARGET "10.13")
 
 macro(find_package_wrapper)
 # do nothing, just satisfy the macro

build_files/cmake/platform/platform_apple_xcode.cmake

@@ -65,13 +65,9 @@ endif()
 
 message(STATUS "Detected OS X ${OSX_SYSTEM} and Xcode ${XCODE_VERSION} at ${XCODE_BUNDLE}")
 
-# Older Xcode versions had different approach to the directory hiearchy.
+# Require a relatively recent Xcode version.
-# Require newer Xcode which is also have better chances of being able to compile with the
+if(${XCODE_VERSION} VERSION_LESS 10.0)
-# required deployment target.
+  message(FATAL_ERROR "Only Xcode version 10.0 and newer is supported")
-#
-# NOTE: Xcode version 8.2 is the latest one which runs on macOS 10.11.
-if(${XCODE_VERSION} VERSION_LESS 8.2)
-  message(FATAL_ERROR "Only Xcode version 8.2 and newer is supported")
 endif()
 
 # note: xcode-select path could be ambiguous,
@@ -133,14 +129,14 @@ if(${CMAKE_GENERATOR} MATCHES "Xcode")
 endif()
 unset(OSX_SDKROOT)
 
-# 10.11 is our min. target, if you use higher sdk, weak linking happens
+# 10.13 is our min. target, if you use higher sdk, weak linking happens
 if(CMAKE_OSX_DEPLOYMENT_TARGET)
-  if(${CMAKE_OSX_DEPLOYMENT_TARGET} VERSION_LESS 10.11)
+  if(${CMAKE_OSX_DEPLOYMENT_TARGET} VERSION_LESS 10.13)
-    message(STATUS "Setting deployment target to 10.11, lower versions are not supported")
+    message(STATUS "Setting deployment target to 10.13, lower versions are not supported")
-    set(CMAKE_OSX_DEPLOYMENT_TARGET "10.11" CACHE STRING "" FORCE)
+    set(CMAKE_OSX_DEPLOYMENT_TARGET "10.13" CACHE STRING "" FORCE)
   endif()
 else()
-  set(CMAKE_OSX_DEPLOYMENT_TARGET "10.11" CACHE STRING "" FORCE)
+  set(CMAKE_OSX_DEPLOYMENT_TARGET "10.13" CACHE STRING "" FORCE)
 endif()
 
 if(NOT ${CMAKE_GENERATOR} MATCHES "Xcode")

build_files/utils/make_test.py

@@ -40,7 +40,8 @@ if make_utils.command_missing(git_command):
 
 # Test if we are building a specific release version.
 branch = make_utils.git_branch(git_command)
-release_version = make_utils.git_branch_release_version(branch)
+tag = make_utils.git_tag(git_command)
+release_version = make_utils.git_branch_release_version(branch, tag)
 lib_tests_dirpath = os.path.join('..', 'lib', "tests")
 
 if not os.path.exists(lib_tests_dirpath):

build_files/utils/make_update.py

@@ -197,7 +197,8 @@ if __name__ == "__main__":
 
     # Test if we are building a specific release version.
     branch = make_utils.git_branch(args.git_command)
-    release_version = make_utils.git_branch_release_version(branch)
+    tag = make_utils.git_tag(args.git_command)
+    release_version = make_utils.git_branch_release_version(branch, tag)
 
     if not args.no_libraries:
         svn_update(args, release_version)

build_files/utils/make_utils.py

@@ -36,7 +36,7 @@ def check_output(cmd, exit_on_error=True):
     return output.strip()
 
 def git_branch(git_command):
-    # Test if we are building a specific release version.
+    # Get current branch name.
     try:
         branch = subprocess.check_output([git_command, "rev-parse", "--abbrev-ref", "HEAD"])
     except subprocess.CalledProcessError as e:
@@ -45,10 +45,23 @@ def git_branch(git_command):
 
     return branch.strip().decode('utf8')
 
-def git_branch_release_version(branch):
+def git_tag(git_command):
+    # Get current tag name.
+    try:
+        tag = subprocess.check_output([git_command, "describe", "--exact-match"])
+    except subprocess.CalledProcessError as e:
+        return None
+
+    return tag.strip().decode('utf8')
+
+def git_branch_release_version(branch, tag):
     release_version = re.search("^blender-v(.*)-release$", branch)
     if release_version:
         release_version = release_version.group(1)
+    elif tag:
+        release_version = re.search("^v([0-9]*\.[0-9]*).*", tag)
+        if release_version:
+            release_version = release_version.group(1)
     return release_version
 
 def svn_libraries_base_url(release_version):

doc/python_api/rst/info_overview.rst

@@ -248,7 +248,7 @@ using the ``bl_idname`` rather than the classes original name.
 .. note::
 
    There are some exceptions to this for class names which aren't guarantee to be unique.
-   In this case use: :func:`bpy.types.Struct.bl_rna_get_subclass`.
+   In this case use: :func:`bpy.types.Struct.bl_rna_get_subclass_py`.
 
 
 When loading a class, Blender performs sanity checks making sure all required properties and functions are found,

extern/mantaflow/helper/util/vectorbase.h

@@ -248,12 +248,14 @@ template<class S> class Vector3D {
  protected:
 };
 
-//! helper to check whether float/double value is non-zero
+//! helper to check whether value is non-zero
-inline bool notZero(Real f)
+template<class S> inline bool notZero(S v)
 {
-  if (std::abs(f) > VECTOR_EPSILON)
+  return (std::abs(v) > VECTOR_EPSILON);
-    return true;
+}
-  return false;
+template<class S> inline bool notZero(Vector3D<S> v)
+{
+  return (std::abs(norm(v)) > VECTOR_EPSILON);
 }
 
 //************************************************************************

extern/mantaflow/preprocessed/fileio/iomeshes.cpp

@@ -315,10 +315,14 @@ int readObjFile(const std::string &name, Mesh *mesh, bool append)
     return 0;
   }
 
+  const Real dx = mesh->getParent()->getDx();
+  const Vec3 gs = toVec3(mesh->getParent()->getGridSize());
+
   if (!append)
     mesh->clear();
   int nodebase = mesh->numNodes();
-  int cnt = nodebase;
+  int cntNodes = nodebase, cntNormals = nodebase;
+
   while (ifs.good() && !ifs.eof()) {
     string id;
     ifs >> id;
@@ -333,19 +337,23 @@ int readObjFile(const std::string &name, Mesh *mesh, bool append)
     }
     else if (id == "vn") {
       // normals
-      if (!mesh->numNodes()) {
+      if (mesh->numNodes() != cntNodes) {
         errMsg("invalid amount of nodes");
         return 0;
       }
-      Node n = mesh->nodes(cnt);
+      Node *n = &mesh->nodes(cntNormals);
-      ifs >> n.normal.x >> n.normal.y >> n.normal.z;
+      ifs >> n->normal.x >> n->normal.y >> n->normal.z;
-      cnt++;
+      cntNormals++;
     }
     else if (id == "v") {
       // vertex
       Node n;
       ifs >> n.pos.x >> n.pos.y >> n.pos.z;
+      // convert to grid space
+      n.pos /= dx;
+      n.pos += gs * 0.5;
       mesh->addNode(n);
+      cntNodes++;
     }
     else if (id == "g") {
       // group
@@ -408,7 +416,6 @@ int writeObjFile(const string &name, Mesh *mesh)
   // write normals
   for (int i = 0; i < numVerts; i++) {
     Vector3D<float> n = toVec3f(mesh->nodes(i).normal);
-    // normalize to unit cube around 0
     ofs << "vn " << n.value[0] << " " << n.value[1] << " " << n.value[2] << " "
         << "\n";
   }

extern/mantaflow/preprocessed/gitinfo.h

@@ -1,3 +1,3 @@
 
 
-#define MANTA_GIT_VERSION "commit d80d3c821de74315ab26b5efd153d41477b976c4"
+#define MANTA_GIT_VERSION "commit 7395d36e3f504edbdabe34b30edc855b422c7baa"

intern/CMakeLists.txt

@@ -30,6 +30,7 @@ add_subdirectory(opensubdiv)
 add_subdirectory(mikktspace)
 add_subdirectory(glew-mx)
 add_subdirectory(eigen)
+add_subdirectory(sky)
 
 if(WITH_AUDASPACE)
   add_subdirectory(audaspace)

intern/cycles/app/CMakeLists.txt

@@ -35,7 +35,7 @@ if(WITH_CYCLES_OSL)
 endif()
 
 if(NOT CYCLES_STANDALONE_REPOSITORY)
-  list(APPEND LIBRARIES bf_intern_glew_mx bf_intern_guardedalloc bf_intern_numaapi)
+  list(APPEND LIBRARIES bf_intern_glew_mx bf_intern_guardedalloc bf_intern_numaapi bf_intern_sky)
 endif()
 
 if(WITH_CYCLES_LOGGING)

intern/cycles/blender/addon/engine.py

@@ -224,7 +224,7 @@ def system_info():
     import _cycles
     return _cycles.system_info()
 
-def list_render_passes(srl):
+def list_render_passes(scene, srl):
     # Builtin Blender passes.
     yield ("Combined", "RGBA", 'COLOR')
 
@@ -279,14 +279,17 @@ def list_render_passes(srl):
             yield ("Denoising Normal",          "XYZ", 'VECTOR')
             yield ("Denoising Albedo",          "RGB", 'COLOR')
             yield ("Denoising Depth",           "Z",   'VALUE')
-            yield ("Denoising Shadowing",       "X",   'VALUE')
+
-            yield ("Denoising Variance",        "RGB", 'COLOR')
+            if scene.cycles.denoiser == 'NLM':
-            yield ("Denoising Intensity",       "X",   'VALUE')
+                yield ("Denoising Shadowing",       "X",   'VALUE')
-            clean_options = ("denoising_diffuse_direct", "denoising_diffuse_indirect",
+                yield ("Denoising Variance",        "RGB", 'COLOR')
-                             "denoising_glossy_direct", "denoising_glossy_indirect",
+                yield ("Denoising Intensity",       "X",   'VALUE')
-                             "denoising_transmission_direct", "denoising_transmission_indirect")
+
-            if any(getattr(crl, option) for option in clean_options):
+                clean_options = ("denoising_diffuse_direct", "denoising_diffuse_indirect",
-                yield ("Denoising Clean", "RGB", 'COLOR')
+                                 "denoising_glossy_direct", "denoising_glossy_indirect",
+                                 "denoising_transmission_direct", "denoising_transmission_indirect")
+                if any(getattr(crl, option) for option in clean_options):
+                    yield ("Denoising Clean", "RGB", 'COLOR')
 
     # Custom AOV passes.
     for aov in crl.aovs:
@@ -298,15 +301,15 @@ def list_render_passes(srl):
 def register_passes(engine, scene, view_layer):
     # Detect duplicate render pass names, first one wins.
     listed = set()
-    for name, channelids, channeltype in list_render_passes(view_layer):
+    for name, channelids, channeltype in list_render_passes(scene, view_layer):
         if name not in listed:
             engine.register_pass(scene, view_layer, name, len(channelids), channelids, channeltype)
             listed.add(name)
 
-def detect_conflicting_passes(view_layer):
+def detect_conflicting_passes(scene, view_layer):
     # Detect conflicting render pass names for UI.
     counter = {}
-    for name, _, _ in list_render_passes(view_layer):
+    for name, _, _ in list_render_passes(scene, view_layer):
         counter[name] = counter.get(name, 0) + 1
 
     for aov in view_layer.cycles.aovs:

intern/cycles/blender/addon/properties.py

@@ -182,6 +182,7 @@ enum_aov_types = (
     ('COLOR', "Color", "Write a Color pass", 1),
 )
 
+
 def enum_openimagedenoise_denoiser(self, context):
     if _cycles.with_openimagedenoise:
         return [('OPENIMAGEDENOISE', "OpenImageDenoise", "Use Intel OpenImageDenoise AI denoiser running on the CPU", 4)]
@@ -196,14 +197,11 @@ def enum_preview_denoiser(self, context):
     optix_items = enum_optix_denoiser(self, context)
     oidn_items = enum_openimagedenoise_denoiser(self, context)
 
-    if len(optix_items):
+    if len(optix_items) or len(oidn_items):
-        auto_label = "Fastest (Optix)"
+        items = [('AUTO', "Automatic", "Use the fastest available denoiser for viewport rendering (OptiX if available, OpenImageDenoise otherwise)", 0)]
-    elif len(oidn_items):
-        auto_label = "Fastest (OpenImageDenoise)"
     else:
-        auto_label = "None"
+        items = [('AUTO', "None", "Blender was compiled without a viewport denoiser", 0)]
 
-    items = [('AUTO', auto_label, "Use the fastest available denoiser for viewport rendering", 0)]
     items += optix_items
     items += oidn_items
     return items
@@ -211,14 +209,23 @@ def enum_preview_denoiser(self, context):
 def enum_denoiser(self, context):
     items = [('NLM', "NLM", "Cycles native non-local means denoiser, running on any compute device", 1)]
     items += enum_optix_denoiser(self, context)
+    items += enum_openimagedenoise_denoiser(self, context)
     return items
 
-enum_denoising_optix_input_passes = (
+enum_denoising_input_passes = (
     ('RGB', "Color", "Use only color as input", 1),
     ('RGB_ALBEDO', "Color + Albedo", "Use color and albedo data as input", 2),
     ('RGB_ALBEDO_NORMAL', "Color + Albedo + Normal", "Use color, albedo and normal data as input", 3),
 )
 
+
+def update_render_passes(self, context):
+    scene = context.scene
+    view_layer = context.view_layer
+    view_layer.update_render_passes()
+    engine.detect_conflicting_passes(scene, view_layer)
+
+
 class CyclesRenderSettings(bpy.types.PropertyGroup):
 
     device: EnumProperty(
@@ -264,9 +271,12 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
 
     denoiser: EnumProperty(
         name="Denoiser",
-        description="Denoise the image with the selected denoiser",
+        description="Denoise the image with the selected denoiser. "
+        "For denoising the image after rendering, denoising data render passes "
+        "also adapt to the selected denoiser",
         items=enum_denoiser,
         default=1,
+        update=update_render_passes,
     )
     preview_denoiser: EnumProperty(
         name="Viewport Denoiser",
@@ -821,6 +831,7 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
     debug_use_cuda_split_kernel: BoolProperty(name="Split Kernel", default=False)
 
     debug_optix_cuda_streams: IntProperty(name="CUDA Streams", default=1, min=1)
+    debug_optix_curves_api: BoolProperty(name="Native OptiX Curve Primitive", default=False)
 
     debug_opencl_kernel_type: EnumProperty(
         name="OpenCL Kernel Type",
@@ -1294,12 +1305,6 @@ class CyclesCurveRenderSettings(bpy.types.PropertyGroup):
         del bpy.types.Scene.cycles_curves
 
 
-def update_render_passes(self, context):
-    view_layer = context.view_layer
-    view_layer.update_render_passes()
-    engine.detect_conflicting_passes(view_layer)
-
-
 class CyclesAOVPass(bpy.types.PropertyGroup):
     name: StringProperty(
         name="Name",
@@ -1433,7 +1438,7 @@ class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
     )
     denoising_store_passes: BoolProperty(
         name="Store Denoising Passes",
-        description="Store the denoising feature passes and the noisy image",
+        description="Store the denoising feature passes and the noisy image. The passes adapt to the denoiser selected for rendering",
         default=False,
         update=update_render_passes,
     )
@@ -1446,11 +1451,18 @@ class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
 
     denoising_optix_input_passes: EnumProperty(
         name="Input Passes",
-        description="Passes handed over to the OptiX denoiser (this can have different effects on the denoised image)",
+        description="Passes used by the denoiser to distinguish noise from shader and geometry detail",
-        items=enum_denoising_optix_input_passes,
+        items=enum_denoising_input_passes,
         default='RGB_ALBEDO',
     )
 
+    denoising_openimagedenoise_input_passes: EnumProperty(
+        name="Input Passes",
+        description="Passes used by the denoiser to distinguish noise from shader and geometry detail",
+        items=enum_denoising_input_passes,
+        default='RGB_ALBEDO_NORMAL',
+    )
+
     use_pass_crypto_object: BoolProperty(
         name="Cryptomatte Object",
         description="Render cryptomatte object pass, for isolating objects in compositing",

intern/cycles/blender/addon/ui.py

@@ -265,7 +265,12 @@ class CYCLES_RENDER_PT_sampling_denoising(CyclesButtonsPanel, Panel):
         row = heading.row(align=True)
         row.prop(cscene, "use_denoising", text="")
         sub = row.row()
+
         sub.active = cscene.use_denoising
+        for view_layer in scene.view_layers:
+            if view_layer.cycles.denoising_store_passes:
+                sub.active = True
+
         sub.prop(cscene, "denoiser", text="")
 
         heading = layout.column(align=False, heading="Viewport")
@@ -777,10 +782,6 @@ class CYCLES_RENDER_PT_filter(CyclesButtonsPanel, Panel):
         col.prop(view_layer, "use_solid", text="Surfaces")
         col.prop(view_layer, "use_strand", text="Hair")
         col.prop(view_layer, "use_volumes", text="Volumes")
-        if with_freestyle:
-            sub = col.row(align=True)
-            sub.prop(view_layer, "use_freestyle", text="Freestyle")
-            sub.active = rd.use_freestyle
 
 
 class CYCLES_RENDER_PT_override(CyclesButtonsPanel, Panel):
@@ -1007,6 +1008,7 @@ class CYCLES_RENDER_PT_denoising(CyclesButtonsPanel, Panel):
             col.prop(cycles_view_layer, "denoising_optix_input_passes")
             return
         elif denoiser == 'OPENIMAGEDENOISE':
+            col.prop(cycles_view_layer, "denoising_openimagedenoise_input_passes")
             return
 
         col.prop(cycles_view_layer, "denoising_radius", text="Radius")
@@ -1222,6 +1224,7 @@ def has_geometry_visibility(ob):
     return ob and ((ob.type in {'MESH', 'CURVE', 'SURFACE', 'FONT', 'META', 'LIGHT'}) or
                     (ob.instance_type == 'COLLECTION' and ob.instance_collection))
 
+
 class CYCLES_OBJECT_PT_shading(CyclesButtonsPanel, Panel):
     bl_label = "Shading"
     bl_context = "object"
@@ -1244,6 +1247,7 @@ class CYCLES_OBJECT_PT_shading(CyclesButtonsPanel, Panel):
             col = flow.column()
             col.prop(cob, "shadow_terminator_offset")
 
+
 class CYCLES_OBJECT_PT_visibility(CyclesButtonsPanel, Panel):
     bl_label = "Visibility"
     bl_context = "object"
@@ -2024,6 +2028,7 @@ class CYCLES_RENDER_PT_debug(CyclesButtonsPanel, Panel):
         col = layout.column()
         col.label(text="OptiX Flags:")
         col.prop(cscene, "debug_optix_cuda_streams")
+        col.prop(cscene, "debug_optix_curves_api")
 
         col.separator()
 

intern/cycles/blender/blender_camera.cpp

@@ -709,6 +709,10 @@ static void blender_camera_from_view(BlenderCamera *bcam,
 
   /* 3d view transform */
   bcam->matrix = transform_inverse(get_transform(b_rv3d.view_matrix()));
+
+  /* dimensions */
+  bcam->full_width = width;
+  bcam->full_height = height;
 }
 
 static void blender_camera_view_subset(BL::RenderEngine &b_engine,

intern/cycles/blender/blender_id_map.h

@@ -200,7 +200,7 @@ template<typename K, typename T> class id_map {
  * To uniquely identify instances, we use the parent, object and persistent instance ID.
  * We also export separate object for a mesh and its particle hair. */
 
-enum { OBJECT_PERSISTENT_ID_SIZE = 16 };
+enum { OBJECT_PERSISTENT_ID_SIZE = 8 /* MAX_DUPLI_RECUR in Blender. */ };
 
 struct ObjectKey {
   void *parent;

intern/cycles/blender/blender_object.cpp

@@ -59,7 +59,7 @@ bool BlenderSync::BKE_object_is_modified(BL::Object &b_ob)
   return false;
 }
 
-bool BlenderSync::object_is_mesh(BL::Object &b_ob)
+bool BlenderSync::object_is_geometry(BL::Object &b_ob)
 {
   BL::ID b_ob_data = b_ob.data();
 
@@ -143,7 +143,7 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph,
   }
 
   /* only interested in object that we can create meshes from */
-  if (!object_is_mesh(b_ob)) {
+  if (!object_is_geometry(b_ob)) {
     return NULL;
   }
 

intern/cycles/blender/blender_python.cpp

@@ -92,6 +92,7 @@ bool debug_flags_sync_from_scene(BL::Scene b_scene)
   flags.cuda.split_kernel = get_boolean(cscene, "debug_use_cuda_split_kernel");
   /* Synchronize OptiX flags. */
   flags.optix.cuda_streams = get_int(cscene, "debug_optix_cuda_streams");
+  flags.optix.curves_api = get_boolean(cscene, "debug_optix_curves_api");
   /* Synchronize OpenCL device type. */
   switch (get_enum(cscene, "debug_opencl_device_type")) {
     case 0:

intern/cycles/blender/blender_shader.cpp

@@ -815,9 +815,10 @@ static ShaderNode *add_node(Scene *scene,
     sky->ground_albedo = b_sky_node.ground_albedo();
     sky->sun_disc = b_sky_node.sun_disc();
     sky->sun_size = b_sky_node.sun_size();
+    sky->sun_intensity = b_sky_node.sun_intensity();
     sky->sun_elevation = b_sky_node.sun_elevation();
     sky->sun_rotation = b_sky_node.sun_rotation();
-    sky->altitude = b_sky_node.altitude();
+    sky->altitude = 1000.0f * b_sky_node.altitude();
     sky->air_density = b_sky_node.air_density();
     sky->dust_density = b_sky_node.dust_density();
     sky->ozone_density = b_sky_node.ozone_density();

intern/cycles/blender/blender_sync.cpp

@@ -147,30 +147,43 @@ void BlenderSync::sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d
     /* Object */
     else if (b_id.is_a(&RNA_Object)) {
       BL::Object b_ob(b_id);
-      const bool updated_geometry = b_update->is_updated_geometry();
+      const bool is_geometry = object_is_geometry(b_ob);
+      const bool is_light = !is_geometry && object_is_light(b_ob);
 
-      if (b_update->is_updated_transform() || b_update->is_updated_shading()) {
+      if (is_geometry || is_light) {
-        object_map.set_recalc(b_ob);
+        const bool updated_geometry = b_update->is_updated_geometry();
-        light_map.set_recalc(b_ob);
-      }
 
-      if (object_is_mesh(b_ob)) {
+        /* Geometry (mesh, hair, volume). */
-        if (updated_geometry ||
+        if (is_geometry) {
-            (object_subdivision_type(b_ob, preview, experimental) != Mesh::SUBDIVISION_NONE)) {
+          if (b_update->is_updated_transform() || b_update->is_updated_shading()) {
-          BL::ID key = BKE_object_is_modified(b_ob) ? b_ob : b_ob.data();
+            object_map.set_recalc(b_ob);
-          geometry_map.set_recalc(key);
+          }
+
+          if (updated_geometry ||
+              (object_subdivision_type(b_ob, preview, experimental) != Mesh::SUBDIVISION_NONE)) {
+            BL::ID key = BKE_object_is_modified(b_ob) ? b_ob : b_ob.data();
+            geometry_map.set_recalc(key);
+          }
+
+          if (updated_geometry) {
+            BL::Object::particle_systems_iterator b_psys;
+            for (b_ob.particle_systems.begin(b_psys); b_psys != b_ob.particle_systems.end();
+                 ++b_psys) {
+              particle_system_map.set_recalc(b_ob);
+            }
+          }
         }
-      }
+        /* Light */
-      else if (object_is_light(b_ob)) {
+        else if (is_light) {
-        if (updated_geometry) {
+          if (b_update->is_updated_transform() || b_update->is_updated_shading()) {
-          light_map.set_recalc(b_ob);
+            object_map.set_recalc(b_ob);
-        }
+            light_map.set_recalc(b_ob);
-      }
+          }
 
-      if (updated_geometry) {
+          if (updated_geometry) {
-        BL::Object::particle_systems_iterator b_psys;
+            light_map.set_recalc(b_ob);
-        for (b_ob.particle_systems.begin(b_psys); b_psys != b_ob.particle_systems.end(); ++b_psys)
+          }
-          particle_system_map.set_recalc(b_ob);
+        }
       }
     }
     /* Mesh */
@@ -941,7 +954,13 @@ DenoiseParams BlenderSync::get_denoise_params(BL::Scene &b_scene,
       denoising.strength = get_float(clayer, "denoising_strength");
       denoising.feature_strength = get_float(clayer, "denoising_feature_strength");
       denoising.relative_pca = get_boolean(clayer, "denoising_relative_pca");
-      denoising.optix_input_passes = get_enum(clayer, "denoising_optix_input_passes");
+
+      denoising.input_passes = (DenoiserInput)get_enum(
+          clayer,
+          (denoising.type == DENOISER_OPTIX) ? "denoising_optix_input_passes" :
+                                               "denoising_openimagedenoise_input_passes",
+          DENOISER_INPUT_NUM,
+          DENOISER_INPUT_RGB_ALBEDO_NORMAL);
 
       denoising.store_passes = get_boolean(clayer, "denoising_store_passes");
     }

intern/cycles/blender/blender_sync.h

@@ -208,7 +208,7 @@ class BlenderSync {
   /* util */
   void find_shader(BL::ID &id, vector<Shader *> &used_shaders, Shader *default_shader);
   bool BKE_object_is_modified(BL::Object &b_ob);
-  bool object_is_mesh(BL::Object &b_ob);
+  bool object_is_geometry(BL::Object &b_ob);
   bool object_is_light(BL::Object &b_ob);
 
   /* variables */

intern/cycles/cmake/external_libs.cmake

@@ -133,9 +133,9 @@ if(CYCLES_STANDALONE_REPOSITORY)
   set(BOOST_DEFINITIONS "-DBOOST_ALL_NO_LIB")
 
   ####
-  # embree
+  # Embree
   if(WITH_CYCLES_EMBREE)
-    find_package(embree 3.8.0 REQUIRED)
+    find_package(Embree 3.8.0 REQUIRED)
   endif()
 
   ####

intern/cycles/device/cuda/device_cuda_impl.cpp

@@ -1760,7 +1760,7 @@ void CUDADevice::denoise(RenderTile &rtile, DenoisingTask &denoising)
   denoising.render_buffer.samples = rtile.sample;
   denoising.buffer.gpu_temporary_mem = true;
 
-  denoising.run_denoising(&rtile);
+  denoising.run_denoising(rtile);
 }
 
 void CUDADevice::adaptive_sampling_filter(uint filter_sample,

intern/cycles/device/device.cpp

@@ -209,13 +209,13 @@ bool Device::bind_fallback_display_space_shader(const float width, const float h
     glUseProgram(fallback_shader_program);
     image_texture_location = glGetUniformLocation(fallback_shader_program, "image_texture");
     if (image_texture_location < 0) {
-      LOG(ERROR) << "Shader doesn't containt the 'image_texture' uniform.";
+      LOG(ERROR) << "Shader doesn't contain the 'image_texture' uniform.";
       return false;
     }
 
     fullscreen_location = glGetUniformLocation(fallback_shader_program, "fullscreen");
     if (fullscreen_location < 0) {
-      LOG(ERROR) << "Shader doesn't containt the 'fullscreen' uniform.";
+      LOG(ERROR) << "Shader doesn't contain the 'fullscreen' uniform.";
       return false;
     }
 

intern/cycles/device/device.h

@@ -439,10 +439,10 @@ class Device {
   {
     return 0;
   }
-  virtual void map_neighbor_tiles(Device * /*sub_device*/, RenderTile * /*tiles*/)
+  virtual void map_neighbor_tiles(Device * /*sub_device*/, RenderTileNeighbors & /*neighbors*/)
   {
   }
-  virtual void unmap_neighbor_tiles(Device * /*sub_device*/, RenderTile * /*tiles*/)
+  virtual void unmap_neighbor_tiles(Device * /*sub_device*/, RenderTileNeighbors & /*neighbors*/)
   {
   }
 

intern/cycles/device/device_cpu.cpp

@@ -182,6 +182,7 @@ class CPUDevice : public Device {
   oidn::DeviceRef oidn_device;
   oidn::FilterRef oidn_filter;
 #endif
+  thread_spin_lock oidn_task_lock;
 
   bool use_split_kernel;
 
@@ -948,12 +949,25 @@ class CPUDevice : public Device {
     }
   }
 
-  void denoise_openimagedenoise(DeviceTask &task, RenderTile &rtile)
+  void denoise_openimagedenoise_buffer(DeviceTask &task,
+                                       float *buffer,
+                                       const size_t offset,
+                                       const size_t stride,
+                                       const size_t x,
+                                       const size_t y,
+                                       const size_t w,
+                                       const size_t h,
+                                       const float scale)
   {
 #ifdef WITH_OPENIMAGEDENOISE
     assert(openimagedenoise_supported());
 
-    /* Only one at a time, since OpenImageDenoise itself is multithreaded. */
+    /* Only one at a time, since OpenImageDenoise itself is multithreaded for full
+     * buffers, and for tiled rendering because creating multiple devices and filters
+     * is slow and memory hungry as well.
+     *
+     * TODO: optimize tiled rendering case, by batching together denoising of many
+     * tiles somehow? */
     static thread_mutex mutex;
     thread_scoped_lock lock(mutex);
 
@@ -964,54 +978,192 @@ class CPUDevice : public Device {
     }
     if (!oidn_filter) {
       oidn_filter = oidn_device.newFilter("RT");
+      oidn_filter.set("hdr", true);
+      oidn_filter.set("srgb", false);
     }
 
-    /* Copy pixels from compute device to CPU (no-op for CPU device). */
-    rtile.buffers->buffer.copy_from_device();
-
     /* Set images with appropriate stride for our interleaved pass storage. */
-    const struct {
+    struct {
       const char *name;
-      int offset;
+      const int offset;
-    } passes[] = {{"color", task.pass_denoising_data + DENOISING_PASS_COLOR},
+      const bool scale;
-                  {"normal", task.pass_denoising_data + DENOISING_PASS_NORMAL},
+      const bool use;
-                  {"albedo", task.pass_denoising_data + DENOISING_PASS_ALBEDO},
+      array<float> scaled_buffer;
-                  {"output", 0},
+    } passes[] = {{"color", task.pass_denoising_data + DENOISING_PASS_COLOR, false, true},
+                  {"albedo",
+                   task.pass_denoising_data + DENOISING_PASS_ALBEDO,
+                   true,
+                   task.denoising.input_passes >= DENOISER_INPUT_RGB_ALBEDO},
+                  {"normal",
+                   task.pass_denoising_data + DENOISING_PASS_NORMAL,
+                   true,
+                   task.denoising.input_passes >= DENOISER_INPUT_RGB_ALBEDO_NORMAL},
+                  {"output", 0, false, true},
                   { NULL,
                     0 }};
 
     for (int i = 0; passes[i].name; i++) {
-      const int64_t offset = rtile.offset + rtile.x + rtile.y * rtile.stride;
+      if (!passes[i].use) {
-      const int64_t buffer_offset = (offset * task.pass_stride + passes[i].offset) * sizeof(float);
+        continue;
-      const int64_t pixel_stride = task.pass_stride * sizeof(float);
+      }
-      const int64_t row_stride = rtile.stride * pixel_stride;
 
-      oidn_filter.setImage(passes[i].name,
+      const int64_t pixel_offset = offset + x + y * stride;
-                           (char *)rtile.buffer + buffer_offset,
+      const int64_t buffer_offset = (pixel_offset * task.pass_stride + passes[i].offset);
-                           oidn::Format::Float3,
+      const int64_t pixel_stride = task.pass_stride;
-                           rtile.w,
+      const int64_t row_stride = stride * pixel_stride;
-                           rtile.h,
+
-                           0,
+      if (passes[i].scale && scale != 1.0f) {
-                           pixel_stride,
+        /* Normalize albedo and normal passes as they are scaled by the number of samples.
-                           row_stride);
+         * For the color passes OIDN will perform auto-exposure making it unnecessary. */
+        array<float> &scaled_buffer = passes[i].scaled_buffer;
+        scaled_buffer.resize(w * h * 3);
+
+        for (int y = 0; y < h; y++) {
+          const float *pass_row = buffer + buffer_offset + y * row_stride;
+          float *scaled_row = scaled_buffer.data() + y * w * 3;
+
+          for (int x = 0; x < w; x++) {
+            scaled_row[x * 3 + 0] = pass_row[x * pixel_stride + 0] * scale;
+            scaled_row[x * 3 + 1] = pass_row[x * pixel_stride + 1] * scale;
+            scaled_row[x * 3 + 2] = pass_row[x * pixel_stride + 2] * scale;
+          }
+        }
+
+        oidn_filter.setImage(
+            passes[i].name, scaled_buffer.data(), oidn::Format::Float3, w, h, 0, 0, 0);
+      }
+      else {
+        oidn_filter.setImage(passes[i].name,
+                             buffer + buffer_offset,
+                             oidn::Format::Float3,
+                             w,
+                             h,
+                             0,
+                             pixel_stride * sizeof(float),
+                             row_stride * sizeof(float));
+      }
     }
 
     /* Execute filter. */
-    oidn_filter.set("hdr", true);
-    oidn_filter.set("srgb", false);
     oidn_filter.commit();
     oidn_filter.execute();
-
-    /* todo: it may be possible to avoid this copy, but we have to ensure that
-     * when other code copies data from the device it doesn't overwrite the
-     * denoiser buffers. */
-    rtile.buffers->buffer.copy_to_device();
 #else
     (void)task;
-    (void)rtile;
+    (void)buffer;
+    (void)offset;
+    (void)stride;
+    (void)x;
+    (void)y;
+    (void)w;
+    (void)h;
+    (void)scale;
 #endif
   }
 
+  void denoise_openimagedenoise(DeviceTask &task, RenderTile &rtile)
+  {
+    if (task.type == DeviceTask::DENOISE_BUFFER) {
+      /* Copy pixels from compute device to CPU (no-op for CPU device). */
+      rtile.buffers->buffer.copy_from_device();
+
+      denoise_openimagedenoise_buffer(task,
+                                      (float *)rtile.buffer,
+                                      rtile.offset,
+                                      rtile.stride,
+                                      rtile.x,
+                                      rtile.y,
+                                      rtile.w,
+                                      rtile.h,
+                                      1.0f / rtile.sample);
+
+      /* todo: it may be possible to avoid this copy, but we have to ensure that
+       * when other code copies data from the device it doesn't overwrite the
+       * denoiser buffers. */
+      rtile.buffers->buffer.copy_to_device();
+    }
+    else {
+      /* Per-tile denoising. */
+      rtile.sample = rtile.start_sample + rtile.num_samples;
+      const float scale = 1.0f / rtile.sample;
+      const float invscale = rtile.sample;
+      const size_t pass_stride = task.pass_stride;
+
+      /* Map neighboring tiles into one buffer for denoising. */
+      RenderTileNeighbors neighbors(rtile);
+      task.map_neighbor_tiles(neighbors, this);
+      RenderTile &center_tile = neighbors.tiles[RenderTileNeighbors::CENTER];
+      rtile = center_tile;
+
+      /* Calculate size of the tile to denoise (including overlap). The overlap
+       * size was chosen empirically. OpenImageDenoise specifies an overlap size
+       * of 128 but this is significantly bigger than typical tile size. */
+      const int4 rect = rect_clip(rect_expand(center_tile.bounds(), 64), neighbors.bounds());
+      const int2 rect_size = make_int2(rect.z - rect.x, rect.w - rect.y);
+
+      /* Adjacent tiles are in separate memory regions, copy into single buffer. */
+      array<float> merged(rect_size.x * rect_size.y * task.pass_stride);
+
+      for (int i = 0; i < RenderTileNeighbors::SIZE; i++) {
+        RenderTile &ntile = neighbors.tiles[i];
+        if (!ntile.buffer) {
+          continue;
+        }
+
+        const int xmin = max(ntile.x, rect.x);
+        const int ymin = max(ntile.y, rect.y);
+        const int xmax = min(ntile.x + ntile.w, rect.z);
+        const int ymax = min(ntile.y + ntile.h, rect.w);
+
+        const size_t tile_offset = ntile.offset + xmin + ymin * ntile.stride;
+        const float *tile_buffer = (float *)ntile.buffer + tile_offset * pass_stride;
+
+        const size_t merged_stride = rect_size.x;
+        const size_t merged_offset = (xmin - rect.x) + (ymin - rect.y) * merged_stride;
+        float *merged_buffer = merged.data() + merged_offset * pass_stride;
+
+        for (int y = ymin; y < ymax; y++) {
+          for (int x = 0; x < pass_stride * (xmax - xmin); x++) {
+            merged_buffer[x] = tile_buffer[x] * scale;
+          }
+          tile_buffer += ntile.stride * pass_stride;
+          merged_buffer += merged_stride * pass_stride;
+        }
+      }
+
+      /* Denoise */
+      denoise_openimagedenoise_buffer(
+          task, merged.data(), 0, rect_size.x, 0, 0, rect_size.x, rect_size.y, 1.0f);
+
+      /* Copy back result from merged buffer. */
+      RenderTile &ntile = neighbors.target;
+      if (ntile.buffer) {
+        const int xmin = max(ntile.x, rect.x);
+        const int ymin = max(ntile.y, rect.y);
+        const int xmax = min(ntile.x + ntile.w, rect.z);
+        const int ymax = min(ntile.y + ntile.h, rect.w);
+
+        const size_t tile_offset = ntile.offset + xmin + ymin * ntile.stride;
+        float *tile_buffer = (float *)ntile.buffer + tile_offset * pass_stride;
+
+        const size_t merged_stride = rect_size.x;
+        const size_t merged_offset = (xmin - rect.x) + (ymin - rect.y) * merged_stride;
+        const float *merged_buffer = merged.data() + merged_offset * pass_stride;
+
+        for (int y = ymin; y < ymax; y++) {
+          for (int x = 0; x < pass_stride * (xmax - xmin); x += pass_stride) {
+            tile_buffer[x + 0] = merged_buffer[x + 0] * invscale;
+            tile_buffer[x + 1] = merged_buffer[x + 1] * invscale;
+            tile_buffer[x + 2] = merged_buffer[x + 2] * invscale;
+          }
+          tile_buffer += ntile.stride * pass_stride;
+          merged_buffer += merged_stride * pass_stride;
+        }
+      }
+
+      task.unmap_neighbor_tiles(neighbors, this);
+    }
+  }
+
   void denoise_nlm(DenoisingTask &denoising, RenderTile &tile)
   {
     ProfilingHelper profiling(denoising.profiler, PROFILING_DENOISING);
@@ -1040,7 +1192,7 @@ class CPUDevice : public Device {
     denoising.render_buffer.samples = tile.sample;
     denoising.buffer.gpu_temporary_mem = false;
 
-    denoising.run_denoising(&tile);
+    denoising.run_denoising(tile);
   }
 
   void thread_render(DeviceTask &task)
@@ -1070,10 +1222,23 @@ class CPUDevice : public Device {
       }
     }
 
+    /* NLM denoiser. */
     DenoisingTask *denoising = NULL;
 
+    /* OpenImageDenoise: we can only denoise with one thread at a time, so to
+     * avoid waiting with mutex locks in the denoiser, we let only a single
+     * thread acquire denoising tiles. */
+    uint tile_types = task.tile_types;
+    bool hold_denoise_lock = false;
+    if ((tile_types & RenderTile::DENOISE) && task.denoising.type == DENOISER_OPENIMAGEDENOISE) {
+      if (!oidn_task_lock.try_lock()) {
+        tile_types &= ~RenderTile::DENOISE;
+        hold_denoise_lock = true;
+      }
+    }
+
     RenderTile tile;
-    while (task.acquire_tile(this, tile, task.tile_types)) {
+    while (task.acquire_tile(this, tile, tile_types)) {
       if (tile.task == RenderTile::PATH_TRACE) {
         if (use_split_kernel) {
           device_only_memory<uchar> void_buffer(this, "void_buffer");
@@ -1108,6 +1273,10 @@ class CPUDevice : public Device {
       }
     }
 
+    if (hold_denoise_lock) {
+      oidn_task_lock.unlock();
+    }
+
     profiler.remove_state(&kg->profiler);
 
     thread_kernel_globals_free((KernelGlobals *)kgbuffer.device_pointer);

intern/cycles/device/device_denoising.cpp

@@ -71,29 +71,30 @@ DenoisingTask::~DenoisingTask()
   tile_info_mem.free();
 }
 
-void DenoisingTask::set_render_buffer(RenderTile *rtiles)
+void DenoisingTask::set_render_buffer(RenderTileNeighbors &neighbors)
 {
-  for (int i = 0; i < 9; i++) {
+  for (int i = 0; i < RenderTileNeighbors::SIZE; i++) {
-    tile_info->offsets[i] = rtiles[i].offset;
+    RenderTile &rtile = neighbors.tiles[i];
-    tile_info->strides[i] = rtiles[i].stride;
+    tile_info->offsets[i] = rtile.offset;
-    tile_info->buffers[i] = rtiles[i].buffer;
+    tile_info->strides[i] = rtile.stride;
+    tile_info->buffers[i] = rtile.buffer;
   }
-  tile_info->x[0] = rtiles[3].x;
+  tile_info->x[0] = neighbors.tiles[3].x;
-  tile_info->x[1] = rtiles[4].x;
+  tile_info->x[1] = neighbors.tiles[4].x;
-  tile_info->x[2] = rtiles[5].x;
+  tile_info->x[2] = neighbors.tiles[5].x;
-  tile_info->x[3] = rtiles[5].x + rtiles[5].w;
+  tile_info->x[3] = neighbors.tiles[5].x + neighbors.tiles[5].w;
-  tile_info->y[0] = rtiles[1].y;
+  tile_info->y[0] = neighbors.tiles[1].y;
-  tile_info->y[1] = rtiles[4].y;
+  tile_info->y[1] = neighbors.tiles[4].y;
-  tile_info->y[2] = rtiles[7].y;
+  tile_info->y[2] = neighbors.tiles[7].y;
-  tile_info->y[3] = rtiles[7].y + rtiles[7].h;
+  tile_info->y[3] = neighbors.tiles[7].y + neighbors.tiles[7].h;
 
-  target_buffer.offset = rtiles[9].offset;
+  target_buffer.offset = neighbors.target.offset;
-  target_buffer.stride = rtiles[9].stride;
+  target_buffer.stride = neighbors.target.stride;
-  target_buffer.ptr = rtiles[9].buffer;
+  target_buffer.ptr = neighbors.target.buffer;
 
-  if (do_prefilter && rtiles[9].buffers) {
+  if (do_prefilter && neighbors.target.buffers) {
     target_buffer.denoising_output_offset =
-        rtiles[9].buffers->params.get_denoising_prefiltered_offset();
+        neighbors.target.buffers->params.get_denoising_prefiltered_offset();
   }
   else {
     target_buffer.denoising_output_offset = 0;
@@ -320,12 +321,11 @@ void DenoisingTask::reconstruct()
   functions.solve(target_buffer.ptr);
 }
 
-void DenoisingTask::run_denoising(RenderTile *tile)
+void DenoisingTask::run_denoising(RenderTile &tile)
 {
-  RenderTile rtiles[10];
+  RenderTileNeighbors neighbors(tile);
-  rtiles[4] = *tile;
+  functions.map_neighbor_tiles(neighbors);
-  functions.map_neighbor_tiles(rtiles);
+  set_render_buffer(neighbors);
-  set_render_buffer(rtiles);
 
   setup_denoising_buffer();
 
@@ -347,7 +347,7 @@ void DenoisingTask::run_denoising(RenderTile *tile)
     write_buffer();
   }
 
-  functions.unmap_neighbor_tiles(rtiles);
+  functions.unmap_neighbor_tiles(neighbors);
 }
 
 CCL_NAMESPACE_END

intern/cycles/device/device_denoising.h

@@ -102,8 +102,8 @@ class DenoisingTask {
                   device_ptr output_ptr)>
         detect_outliers;
     function<bool(int out_offset, device_ptr frop_ptr, device_ptr buffer_ptr)> write_feature;
-    function<void(RenderTile *rtiles)> map_neighbor_tiles;
+    function<void(RenderTileNeighbors &neighbors)> map_neighbor_tiles;
-    function<void(RenderTile *rtiles)> unmap_neighbor_tiles;
+    function<void(RenderTileNeighbors &neighbors)> unmap_neighbor_tiles;
   } functions;
 
   /* Stores state of the current Reconstruction operation,
@@ -154,7 +154,7 @@ class DenoisingTask {
   DenoisingTask(Device *device, const DeviceTask &task);
   ~DenoisingTask();
 
-  void run_denoising(RenderTile *tile);
+  void run_denoising(RenderTile &tile);
 
   struct DenoiseBuffers {
     int pass_stride;
@@ -179,7 +179,7 @@ class DenoisingTask {
  protected:
   Device *device;
 
-  void set_render_buffer(RenderTile *rtiles);
+  void set_render_buffer(RenderTileNeighbors &neighbors);
   void setup_denoising_buffer();
   void prefilter_shadowing();
   void prefilter_features();

intern/cycles/device/device_multi.cpp

@@ -177,8 +177,11 @@ class MultiDevice : public Device {
         return false;
 
     if (requested_features.use_denoising) {
+      /* Only need denoising feature, everything else is unused. */
+      DeviceRequestedFeatures denoising_features;
+      denoising_features.use_denoising = true;
       foreach (SubDevice &sub, denoising_devices)
-        if (!sub.device->load_kernels(requested_features))
+        if (!sub.device->load_kernels(denoising_features))
           return false;
     }
 
@@ -581,20 +584,22 @@ class MultiDevice : public Device {
     return -1;
   }
 
-  void map_neighbor_tiles(Device *sub_device, RenderTile *tiles)
+  void map_neighbor_tiles(Device *sub_device, RenderTileNeighbors &neighbors)
   {
-    for (int i = 0; i < 9; i++) {
+    for (int i = 0; i < RenderTileNeighbors::SIZE; i++) {
-      if (!tiles[i].buffers) {
+      RenderTile &tile = neighbors.tiles[i];
+
+      if (!tile.buffers) {
         continue;
       }
 
-      device_vector<float> &mem = tiles[i].buffers->buffer;
+      device_vector<float> &mem = tile.buffers->buffer;
-      tiles[i].buffer = mem.device_pointer;
+      tile.buffer = mem.device_pointer;
 
       if (mem.device == this && matching_rendering_and_denoising_devices) {
         /* Skip unnecessary copies in viewport mode (buffer covers the
          * whole image), but still need to fix up the tile device pointer. */
-        map_tile(sub_device, tiles[i]);
+        map_tile(sub_device, tile);
         continue;
       }
 
@@ -607,15 +612,15 @@ class MultiDevice : public Device {
          * also required for the case where a CPU thread is denoising
          * a tile rendered on the GPU. In that case we have to avoid
          * overwriting the buffer being de-noised by the CPU thread. */
-        if (!tiles[i].buffers->map_neighbor_copied) {
+        if (!tile.buffers->map_neighbor_copied) {
-          tiles[i].buffers->map_neighbor_copied = true;
+          tile.buffers->map_neighbor_copied = true;
           mem.copy_from_device();
         }
 
         if (mem.device == this) {
           /* Can re-use memory if tile is already allocated on the sub device. */
-          map_tile(sub_device, tiles[i]);
+          map_tile(sub_device, tile);
-          mem.swap_device(sub_device, mem.device_size, tiles[i].buffer);
+          mem.swap_device(sub_device, mem.device_size, tile.buffer);
         }
         else {
           mem.swap_device(sub_device, 0, 0);
@@ -623,40 +628,42 @@ class MultiDevice : public Device {
 
         mem.copy_to_device();
 
-        tiles[i].buffer = mem.device_pointer;
+        tile.buffer = mem.device_pointer;
-        tiles[i].device_size = mem.device_size;
+        tile.device_size = mem.device_size;
 
         mem.restore_device();
       }
     }
   }
 
-  void unmap_neighbor_tiles(Device *sub_device, RenderTile *tiles)
+  void unmap_neighbor_tiles(Device *sub_device, RenderTileNeighbors &neighbors)
   {
-    device_vector<float> &mem = tiles[9].buffers->buffer;
+    RenderTile &target_tile = neighbors.target;
+    device_vector<float> &mem = target_tile.buffers->buffer;
 
     if (mem.device == this && matching_rendering_and_denoising_devices) {
       return;
     }
 
     /* Copy denoised result back to the host. */
-    mem.swap_device(sub_device, tiles[9].device_size, tiles[9].buffer);
+    mem.swap_device(sub_device, target_tile.device_size, target_tile.buffer);
     mem.copy_from_device();
     mem.restore_device();
 
     /* Copy denoised result to the original device. */
     mem.copy_to_device();
 
-    for (int i = 0; i < 9; i++) {
+    for (int i = 0; i < RenderTileNeighbors::SIZE; i++) {
-      if (!tiles[i].buffers) {
+      RenderTile &tile = neighbors.tiles[i];
+      if (!tile.buffers) {
         continue;
       }
 
-      device_vector<float> &mem = tiles[i].buffers->buffer;
+      device_vector<float> &mem = tile.buffers->buffer;
 
       if (mem.device != sub_device && mem.device != this) {
         /* Free up memory again if it was allocated for the copy above. */
-        mem.swap_device(sub_device, tiles[i].device_size, tiles[i].buffer);
+        mem.swap_device(sub_device, tile.device_size, tile.buffer);
         sub_device->mem_free(mem);
         mem.restore_device();
       }

intern/cycles/device/device_optix.cpp

@@ -131,8 +131,12 @@ class OptiXDevice : public CUDADevice {
     PG_RGEN,
     PG_MISS,
     PG_HITD,  // Default hit group
-    PG_HITL,  // __BVH_LOCAL__ hit group
     PG_HITS,  // __SHADOW_RECORD_ALL__ hit group
+    PG_HITL,  // __BVH_LOCAL__ hit group (only used for triangles)
+#  if OPTIX_ABI_VERSION >= 36
+    PG_HITD_MOTION,
+    PG_HITS_MOTION,
+#  endif
 #  ifdef WITH_CYCLES_DEBUG
     PG_EXCP,
 #  endif
@@ -177,6 +181,7 @@ class OptiXDevice : public CUDADevice {
   OptixDeviceContext context = NULL;
 
   OptixModule optix_module = NULL;  // All necessary OptiX kernels are in one module
+  OptixModule builtin_modules[2] = {};
   OptixPipeline pipelines[NUM_PIPELINES] = {};
 
   bool motion_blur = false;
@@ -264,6 +269,9 @@ class OptiXDevice : public CUDADevice {
     // Unload modules
     if (optix_module != NULL)
       optixModuleDestroy(optix_module);
+    for (unsigned int i = 0; i < 2; ++i)
+      if (builtin_modules[i] != NULL)
+        optixModuleDestroy(builtin_modules[i]);
     for (unsigned int i = 0; i < NUM_PIPELINES; ++i)
       if (pipelines[i] != NULL)
         optixPipelineDestroy(pipelines[i]);
@@ -338,6 +346,12 @@ class OptiXDevice : public CUDADevice {
       optixModuleDestroy(optix_module);
       optix_module = NULL;
     }
+    for (unsigned int i = 0; i < 2; ++i) {
+      if (builtin_modules[i] != NULL) {
+        optixModuleDestroy(builtin_modules[i]);
+        builtin_modules[i] = NULL;
+      }
+    }
     for (unsigned int i = 0; i < NUM_PIPELINES; ++i) {
       if (pipelines[i] != NULL) {
         optixPipelineDestroy(pipelines[i]);
@@ -369,6 +383,18 @@ class OptiXDevice : public CUDADevice {
 #  endif
     pipeline_options.pipelineLaunchParamsVariableName = "__params";  // See kernel_globals.h
 
+#  if OPTIX_ABI_VERSION >= 36
+    pipeline_options.usesPrimitiveTypeFlags = OPTIX_PRIMITIVE_TYPE_FLAGS_TRIANGLE;
+    if (requested_features.use_hair) {
+      if (DebugFlags().optix.curves_api && requested_features.use_hair_thick) {
+        pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_CUBIC_BSPLINE;
+      }
+      else {
+        pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_CUSTOM;
+      }
+    }
+#  endif
+
     // Keep track of whether motion blur is enabled, so to enable/disable motion in BVH builds
     // This is necessary since objects may be reported to have motion if the Vector pass is
     // active, but may still need to be rendered without motion blur if that isn't active as well
@@ -442,6 +468,34 @@ class OptiXDevice : public CUDADevice {
         group_descs[PG_HITD].hitgroup.entryFunctionNameIS = "__intersection__curve_ribbon";
         group_descs[PG_HITS].hitgroup.entryFunctionNameIS = "__intersection__curve_ribbon";
       }
+
+#  if OPTIX_ABI_VERSION >= 36
+      if (DebugFlags().optix.curves_api && requested_features.use_hair_thick) {
+        OptixBuiltinISOptions builtin_options;
+        builtin_options.builtinISModuleType = OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE;
+        builtin_options.usesMotionBlur = false;
+
+        check_result_optix_ret(optixBuiltinISModuleGet(
+            context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[0]));
+
+        group_descs[PG_HITD].hitgroup.moduleIS = builtin_modules[0];
+        group_descs[PG_HITD].hitgroup.entryFunctionNameIS = nullptr;
+        group_descs[PG_HITS].hitgroup.moduleIS = builtin_modules[0];
+        group_descs[PG_HITS].hitgroup.entryFunctionNameIS = nullptr;
+
+        if (motion_blur) {
+          builtin_options.usesMotionBlur = true;
+
+          check_result_optix_ret(optixBuiltinISModuleGet(
+              context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[1]));
+
+          group_descs[PG_HITD_MOTION] = group_descs[PG_HITD];
+          group_descs[PG_HITD_MOTION].hitgroup.moduleIS = builtin_modules[1];
+          group_descs[PG_HITS_MOTION] = group_descs[PG_HITS];
+          group_descs[PG_HITS_MOTION].hitgroup.moduleIS = builtin_modules[1];
+        }
+      }
+#  endif
     }
 
     if (requested_features.use_subsurface || requested_features.use_shader_raytrace) {
@@ -493,8 +547,14 @@ class OptiXDevice : public CUDADevice {
     unsigned int trace_css = stack_size[PG_HITD].cssCH;
     // This is based on the maximum of closest-hit and any-hit/intersection programs
     trace_css = std::max(trace_css, stack_size[PG_HITD].cssIS + stack_size[PG_HITD].cssAH);
-    trace_css = std::max(trace_css, stack_size[PG_HITL].cssIS + stack_size[PG_HITL].cssAH);
     trace_css = std::max(trace_css, stack_size[PG_HITS].cssIS + stack_size[PG_HITS].cssAH);
+    trace_css = std::max(trace_css, stack_size[PG_HITL].cssIS + stack_size[PG_HITL].cssAH);
+#  if OPTIX_ABI_VERSION >= 36
+    trace_css = std::max(trace_css,
+                         stack_size[PG_HITD_MOTION].cssIS + stack_size[PG_HITD_MOTION].cssAH);
+    trace_css = std::max(trace_css,
+                         stack_size[PG_HITS_MOTION].cssIS + stack_size[PG_HITS_MOTION].cssAH);
+#  endif
 
     OptixPipelineLinkOptions link_options;
     link_options.maxTraceDepth = 1;
@@ -503,17 +563,23 @@ class OptiXDevice : public CUDADevice {
 #  else
     link_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_LINEINFO;
 #  endif
-    link_options.overrideUsesMotionBlur = pipeline_options.usesMotionBlur;
+#  if OPTIX_ABI_VERSION < 24
+    link_options.overrideUsesMotionBlur = motion_blur;
+#  endif
 
     {  // Create path tracing pipeline
       OptixProgramGroup pipeline_groups[] = {
-          groups[PG_RGEN],
+        groups[PG_RGEN],
-          groups[PG_MISS],
+        groups[PG_MISS],
-          groups[PG_HITD],
+        groups[PG_HITD],
-          groups[PG_HITS],
+        groups[PG_HITS],
-          groups[PG_HITL],
+        groups[PG_HITL],
+#  if OPTIX_ABI_VERSION >= 36
+        groups[PG_HITD_MOTION],
+        groups[PG_HITS_MOTION],
+#  endif
 #  ifdef WITH_CYCLES_DEBUG
-          groups[PG_EXCP],
+        groups[PG_EXCP],
 #  endif
       };
       check_result_optix_ret(
@@ -530,8 +596,8 @@ class OptiXDevice : public CUDADevice {
       const unsigned int css = stack_size[PG_RGEN].cssRG + link_options.maxTraceDepth * trace_css;
 
       // Set stack size depending on pipeline options
-      check_result_optix_ret(optixPipelineSetStackSize(
+      check_result_optix_ret(
-          pipelines[PIP_PATH_TRACE], 0, 0, css, (pipeline_options.usesMotionBlur ? 3 : 2)));
+          optixPipelineSetStackSize(pipelines[PIP_PATH_TRACE], 0, 0, css, (motion_blur ? 3 : 2)));
     }
 
     // Only need to create shader evaluation pipeline if one of these features is used:
@@ -541,15 +607,19 @@ class OptiXDevice : public CUDADevice {
 
     if (use_shader_eval_pipeline) {  // Create shader evaluation pipeline
       OptixProgramGroup pipeline_groups[] = {
-          groups[PG_BAKE],
+        groups[PG_BAKE],
-          groups[PG_DISP],
+        groups[PG_DISP],
-          groups[PG_BACK],
+        groups[PG_BACK],
-          groups[PG_MISS],
+        groups[PG_MISS],
-          groups[PG_HITD],
+        groups[PG_HITD],
-          groups[PG_HITS],
+        groups[PG_HITS],
-          groups[PG_HITL],
+        groups[PG_HITL],
+#  if OPTIX_ABI_VERSION >= 36
+        groups[PG_HITD_MOTION],
+        groups[PG_HITS_MOTION],
+#  endif
 #  ifdef WITH_CYCLES_DEBUG
-          groups[PG_EXCP],
+        groups[PG_EXCP],
 #  endif
       };
       check_result_optix_ret(
@@ -672,7 +742,11 @@ class OptiXDevice : public CUDADevice {
       sbt_params.missRecordCount = 1;
       sbt_params.hitgroupRecordBase = sbt_data.device_pointer + PG_HITD * sizeof(SbtRecord);
       sbt_params.hitgroupRecordStrideInBytes = sizeof(SbtRecord);
-      sbt_params.hitgroupRecordCount = 3;  // PG_HITD, PG_HITL, PG_HITS
+#  if OPTIX_ABI_VERSION >= 36
+      sbt_params.hitgroupRecordCount = 5;  // PG_HITD(_MOTION), PG_HITS(_MOTION), PG_HITL
+#  else
+      sbt_params.hitgroupRecordCount = 3;  // PG_HITD, PG_HITS, PG_HITL
+#  endif
 
       // Launch the ray generation program
       check_result_optix(optixLaunch(pipelines[PIP_PATH_TRACE],
@@ -727,19 +801,18 @@ class OptiXDevice : public CUDADevice {
       //   0 1 2
       //   3 4 5
       //   6 7 8  9
-      RenderTile rtiles[10];
+      RenderTileNeighbors neighbors(rtile);
-      rtiles[4] = rtile;
+      task.map_neighbor_tiles(neighbors, this);
-      task.map_neighbor_tiles(rtiles, this);
+      RenderTile &center_tile = neighbors.tiles[RenderTileNeighbors::CENTER];
-      rtile = rtiles[4];  // Tile may have been modified by mapping code
+      RenderTile &target_tile = neighbors.target;
+      rtile = center_tile;  // Tile may have been modified by mapping code
 
       // Calculate size of the tile to denoise (including overlap)
-      int4 rect = make_int4(
+      int4 rect = center_tile.bounds();
-          rtiles[4].x, rtiles[4].y, rtiles[4].x + rtiles[4].w, rtiles[4].y + rtiles[4].h);
       // Overlap between tiles has to be at least 64 pixels
       // TODO(pmours): Query this value from OptiX
       rect = rect_expand(rect, 64);
-      int4 clip_rect = make_int4(
+      int4 clip_rect = neighbors.bounds();
-          rtiles[3].x, rtiles[1].y, rtiles[5].x + rtiles[5].w, rtiles[7].y + rtiles[7].h);
       rect = rect_clip(rect, clip_rect);
       int2 rect_size = make_int2(rect.z - rect.x, rect.w - rect.y);
       int2 overlap_offset = make_int2(rtile.x - rect.x, rtile.y - rect.y);
@@ -760,14 +833,14 @@ class OptiXDevice : public CUDADevice {
       device_only_memory<float> input(this, "denoiser input");
       device_vector<TileInfo> tile_info_mem(this, "denoiser tile info", MEM_READ_WRITE);
 
-      if ((!rtiles[0].buffer || rtiles[0].buffer == rtile.buffer) &&
+      bool contiguous_memory = true;
-          (!rtiles[1].buffer || rtiles[1].buffer == rtile.buffer) &&
+      for (int i = 0; i < RenderTileNeighbors::SIZE; i++) {
-          (!rtiles[2].buffer || rtiles[2].buffer == rtile.buffer) &&
+        if (neighbors.tiles[i].buffer && neighbors.tiles[i].buffer != rtile.buffer) {
-          (!rtiles[3].buffer || rtiles[3].buffer == rtile.buffer) &&
+          contiguous_memory = false;
-          (!rtiles[5].buffer || rtiles[5].buffer == rtile.buffer) &&
+        }
-          (!rtiles[6].buffer || rtiles[6].buffer == rtile.buffer) &&
+      }
-          (!rtiles[7].buffer || rtiles[7].buffer == rtile.buffer) &&
+
-          (!rtiles[8].buffer || rtiles[8].buffer == rtile.buffer)) {
+      if (contiguous_memory) {
         // Tiles are in continous memory, so can just subtract overlap offset
         input_ptr -= (overlap_offset.x + overlap_offset.y * rtile.stride) * pixel_stride;
         // Stride covers the whole width of the image and not just a single tile
@@ -782,19 +855,19 @@ class OptiXDevice : public CUDADevice {
         input_stride *= rect_size.x;
 
         TileInfo *tile_info = tile_info_mem.alloc(1);
-        for (int i = 0; i < 9; i++) {
+        for (int i = 0; i < RenderTileNeighbors::SIZE; i++) {
-          tile_info->offsets[i] = rtiles[i].offset;
+          tile_info->offsets[i] = neighbors.tiles[i].offset;
-          tile_info->strides[i] = rtiles[i].stride;
+          tile_info->strides[i] = neighbors.tiles[i].stride;
-          tile_info->buffers[i] = rtiles[i].buffer;
+          tile_info->buffers[i] = neighbors.tiles[i].buffer;
         }
-        tile_info->x[0] = rtiles[3].x;
+        tile_info->x[0] = neighbors.tiles[3].x;
-        tile_info->x[1] = rtiles[4].x;
+        tile_info->x[1] = neighbors.tiles[4].x;
-        tile_info->x[2] = rtiles[5].x;
+        tile_info->x[2] = neighbors.tiles[5].x;
-        tile_info->x[3] = rtiles[5].x + rtiles[5].w;
+        tile_info->x[3] = neighbors.tiles[5].x + neighbors.tiles[5].w;
-        tile_info->y[0] = rtiles[1].y;
+        tile_info->y[0] = neighbors.tiles[1].y;
-        tile_info->y[1] = rtiles[4].y;
+        tile_info->y[1] = neighbors.tiles[4].y;
-        tile_info->y[2] = rtiles[7].y;
+        tile_info->y[2] = neighbors.tiles[7].y;
-        tile_info->y[3] = rtiles[7].y + rtiles[7].h;
+        tile_info->y[3] = neighbors.tiles[7].y + neighbors.tiles[7].h;
         tile_info_mem.copy_to_device();
 
         void *args[] = {
@@ -804,7 +877,7 @@ class OptiXDevice : public CUDADevice {
 
 #  if OPTIX_DENOISER_NO_PIXEL_STRIDE
       device_only_memory<float> input_rgb(this, "denoiser input rgb");
-      input_rgb.alloc_to_device(rect_size.x * rect_size.y * 3 * task.denoising.optix_input_passes);
+      input_rgb.alloc_to_device(rect_size.x * rect_size.y * 3 * task.denoising.input_passes);
 
       void *input_args[] = {&input_rgb.device_pointer,
                             &input_ptr,
@@ -813,7 +886,7 @@ class OptiXDevice : public CUDADevice {
                             &input_stride,
                             &task.pass_stride,
                             const_cast<int *>(pass_offset),
-                            &task.denoising.optix_input_passes,
+                            &task.denoising.input_passes,
                             &rtile.sample};
       launch_filter_kernel(
           "kernel_cuda_filter_convert_to_rgb", rect_size.x, rect_size.y, input_args);
@@ -824,7 +897,7 @@ class OptiXDevice : public CUDADevice {
 #  endif
 
       const bool recreate_denoiser = (denoiser == NULL) ||
-                                     (task.denoising.optix_input_passes != denoiser_input_passes);
+                                     (task.denoising.input_passes != denoiser_input_passes);
       if (recreate_denoiser) {
         // Destroy existing handle before creating new one
         if (denoiser != NULL) {
@@ -833,23 +906,29 @@ class OptiXDevice : public CUDADevice {
 
         // Create OptiX denoiser handle on demand when it is first used
         OptixDenoiserOptions denoiser_options;
-        assert(task.denoising.optix_input_passes >= 1 && task.denoising.optix_input_passes <= 3);
+        assert(task.denoising.input_passes >= 1 && task.denoising.input_passes <= 3);
         denoiser_options.inputKind = static_cast<OptixDenoiserInputKind>(
-            OPTIX_DENOISER_INPUT_RGB + (task.denoising.optix_input_passes - 1));
+            OPTIX_DENOISER_INPUT_RGB + (task.denoising.input_passes - 1));
+#  if OPTIX_ABI_VERSION < 28
         denoiser_options.pixelFormat = OPTIX_PIXEL_FORMAT_FLOAT3;
+#  endif
         check_result_optix_ret(optixDenoiserCreate(context, &denoiser_options, &denoiser));
         check_result_optix_ret(
             optixDenoiserSetModel(denoiser, OPTIX_DENOISER_MODEL_KIND_HDR, NULL, 0));
 
         // OptiX denoiser handle was created with the requested number of input passes
-        denoiser_input_passes = task.denoising.optix_input_passes;
+        denoiser_input_passes = task.denoising.input_passes;
       }
 
       OptixDenoiserSizes sizes = {};
       check_result_optix_ret(
           optixDenoiserComputeMemoryResources(denoiser, rect_size.x, rect_size.y, &sizes));
 
+#  if OPTIX_ABI_VERSION < 28
       const size_t scratch_size = sizes.recommendedScratchSizeInBytes;
+#  else
+      const size_t scratch_size = sizes.withOverlapScratchSizeInBytes;
+#  endif
       const size_t scratch_offset = sizes.stateSizeInBytes;
 
       // Allocate denoiser state if tile size has changed since last setup
@@ -897,10 +976,10 @@ class OptiXDevice : public CUDADevice {
       int2 output_offset = overlap_offset;
       overlap_offset = make_int2(0, 0);  // Not supported by denoiser API, so apply manually
 #  else
-      output_layers[0].data = rtiles[9].buffer + pixel_offset;
+      output_layers[0].data = target_tile.buffer + pixel_offset;
-      output_layers[0].width = rtiles[9].w;
+      output_layers[0].width = target_tile.w;
-      output_layers[0].height = rtiles[9].h;
+      output_layers[0].height = target_tile.h;
-      output_layers[0].rowStrideInBytes = rtiles[9].stride * pixel_stride;
+      output_layers[0].rowStrideInBytes = target_tile.stride * pixel_stride;
       output_layers[0].pixelStrideInBytes = pixel_stride;
 #  endif
       output_layers[0].format = OPTIX_PIXEL_FORMAT_FLOAT3;
@@ -913,7 +992,7 @@ class OptiXDevice : public CUDADevice {
                                                  denoiser_state.device_pointer,
                                                  scratch_offset,
                                                  input_layers,
-                                                 task.denoising.optix_input_passes,
+                                                 task.denoising.input_passes,
                                                  overlap_offset.x,
                                                  overlap_offset.y,
                                                  output_layers,
@@ -922,26 +1001,26 @@ class OptiXDevice : public CUDADevice {
 
 #  if OPTIX_DENOISER_NO_PIXEL_STRIDE
       void *output_args[] = {&input_ptr,
-                             &rtiles[9].buffer,
+                             &target_tile.buffer,
                              &output_offset.x,
                              &output_offset.y,
                              &rect_size.x,
                              &rect_size.y,
-                             &rtiles[9].x,
+                             &target_tile.x,
-                             &rtiles[9].y,
+                             &target_tile.y,
-                             &rtiles[9].w,
+                             &target_tile.w,
-                             &rtiles[9].h,
+                             &target_tile.h,
-                             &rtiles[9].offset,
+                             &target_tile.offset,
-                             &rtiles[9].stride,
+                             &target_tile.stride,
                              &task.pass_stride,
                              &rtile.sample};
       launch_filter_kernel(
-          "kernel_cuda_filter_convert_from_rgb", rtiles[9].w, rtiles[9].h, output_args);
+          "kernel_cuda_filter_convert_from_rgb", target_tile.w, target_tile.h, output_args);
 #  endif
 
       check_result_cuda_ret(cuStreamSynchronize(0));
 
-      task.unmap_neighbor_tiles(rtiles, this);
+      task.unmap_neighbor_tiles(neighbors, this);
     }
     else {
       // Run CUDA denoising kernels
@@ -993,7 +1072,11 @@ class OptiXDevice : public CUDADevice {
       sbt_params.missRecordCount = 1;
       sbt_params.hitgroupRecordBase = sbt_data.device_pointer + PG_HITD * sizeof(SbtRecord);
       sbt_params.hitgroupRecordStrideInBytes = sizeof(SbtRecord);
-      sbt_params.hitgroupRecordCount = 3;  // PG_HITD, PG_HITL, PG_HITS
+#  if OPTIX_ABI_VERSION >= 36
+      sbt_params.hitgroupRecordCount = 5;  // PG_HITD(_MOTION), PG_HITS(_MOTION), PG_HITL
+#  else
+      sbt_params.hitgroupRecordCount = 3;  // PG_HITD, PG_HITS, PG_HITL
+#  endif
 
       check_result_optix(optixLaunch(pipelines[PIP_SHADER_EVAL],
                                      cuda_stream[thread_index],
@@ -1070,7 +1153,7 @@ class OptiXDevice : public CUDADevice {
                                            &build_input,
                                            1,
                                            temp_mem.device_pointer,
-                                           temp_mem.device_size,
+                                           sizes.tempSizeInBytes,
                                            out_data,
                                            sizes.outputSizeInBytes,
                                            &out_handle,
@@ -1142,7 +1225,6 @@ class OptiXDevice : public CUDADevice {
           continue;
         }
 
-        const size_t num_curves = hair->num_curves();
         const size_t num_segments = hair->num_segments();
 
         size_t num_motion_steps = 1;
@@ -1152,7 +1234,18 @@ class OptiXDevice : public CUDADevice {
         }
 
         device_vector<OptixAabb> aabb_data(this, "temp_aabb_data", MEM_READ_ONLY);
-        aabb_data.alloc(num_segments * num_motion_steps);
+#  if OPTIX_ABI_VERSION >= 36
+        device_vector<int> index_data(this, "temp_index_data", MEM_READ_ONLY);
+        device_vector<float4> vertex_data(this, "temp_vertex_data", MEM_READ_ONLY);
+        // Four control points for each curve segment
+        const size_t num_vertices = num_segments * 4;
+        if (DebugFlags().optix.curves_api && hair->curve_shape == CURVE_THICK) {
+          index_data.alloc(num_segments);
+          vertex_data.alloc(num_vertices * num_motion_steps);
+        }
+        else
+#  endif
+          aabb_data.alloc(num_segments * num_motion_steps);
 
         // Get AABBs for each motion step
         for (size_t step = 0; step < num_motion_steps; ++step) {
@@ -1165,44 +1258,127 @@ class OptiXDevice : public CUDADevice {
             keys = motion_keys->data_float3() + attr_offset * hair->curve_keys.size();
           }
 
-          size_t i = step * num_segments;
+          for (size_t j = 0, i = 0; j < hair->num_curves(); ++j) {
-          for (size_t j = 0; j < num_curves; ++j) {
+            const Hair::Curve curve = hair->get_curve(j);
-            const Hair::Curve c = hair->get_curve(j);
 
-            for (size_t k = 0; k < c.num_segments(); ++i, ++k) {
+            for (int segment = 0; segment < curve.num_segments(); ++segment, ++i) {
-              BoundBox bounds = BoundBox::empty;
+#  if OPTIX_ABI_VERSION >= 36
-              c.bounds_grow(k, keys, hair->curve_radius.data(), bounds);
+              if (DebugFlags().optix.curves_api && hair->curve_shape == CURVE_THICK) {
+                int k0 = curve.first_key + segment;
+                int k1 = k0 + 1;
+                int ka = max(k0 - 1, curve.first_key);
+                int kb = min(k1 + 1, curve.first_key + curve.num_keys - 1);
 
-              aabb_data[i].minX = bounds.min.x;
+                const float4 px = make_float4(keys[ka].x, keys[k0].x, keys[k1].x, keys[kb].x);
-              aabb_data[i].minY = bounds.min.y;
+                const float4 py = make_float4(keys[ka].y, keys[k0].y, keys[k1].y, keys[kb].y);
-              aabb_data[i].minZ = bounds.min.z;
+                const float4 pz = make_float4(keys[ka].z, keys[k0].z, keys[k1].z, keys[kb].z);
-              aabb_data[i].maxX = bounds.max.x;
+                const float4 pw = make_float4(hair->curve_radius[ka],
-              aabb_data[i].maxY = bounds.max.y;
+                                              hair->curve_radius[k0],
-              aabb_data[i].maxZ = bounds.max.z;
+                                              hair->curve_radius[k1],
+                                              hair->curve_radius[kb]);
+
+                // Convert Catmull-Rom data to Bezier spline
+                static const float4 cr2bsp0 = make_float4(+7, -4, +5, -2) / 6.f;
+                static const float4 cr2bsp1 = make_float4(-2, 11, -4, +1) / 6.f;
+                static const float4 cr2bsp2 = make_float4(+1, -4, 11, -2) / 6.f;
+                static const float4 cr2bsp3 = make_float4(-2, +5, -4, +7) / 6.f;
+
+                index_data[i] = i * 4;
+                float4 *const v = vertex_data.data() + step * num_vertices + index_data[i];
+                v[0] = make_float4(
+                    dot(cr2bsp0, px), dot(cr2bsp0, py), dot(cr2bsp0, pz), dot(cr2bsp0, pw));
+                v[1] = make_float4(
+                    dot(cr2bsp1, px), dot(cr2bsp1, py), dot(cr2bsp1, pz), dot(cr2bsp1, pw));
+                v[2] = make_float4(
+                    dot(cr2bsp2, px), dot(cr2bsp2, py), dot(cr2bsp2, pz), dot(cr2bsp2, pw));
+                v[3] = make_float4(
+                    dot(cr2bsp3, px), dot(cr2bsp3, py), dot(cr2bsp3, pz), dot(cr2bsp3, pw));
+              }
+              else
+#  endif
+              {
+                BoundBox bounds = BoundBox::empty;
+                curve.bounds_grow(segment, keys, hair->curve_radius.data(), bounds);
+
+                const size_t index = step * num_segments + i;
+                aabb_data[index].minX = bounds.min.x;
+                aabb_data[index].minY = bounds.min.y;
+                aabb_data[index].minZ = bounds.min.z;
+                aabb_data[index].maxX = bounds.max.x;
+                aabb_data[index].maxY = bounds.max.y;
+                aabb_data[index].maxZ = bounds.max.z;
+              }
             }
           }
         }
 
         // Upload AABB data to GPU
         aabb_data.copy_to_device();
+#  if OPTIX_ABI_VERSION >= 36
+        index_data.copy_to_device();
+        vertex_data.copy_to_device();
+#  endif
 
         vector<device_ptr> aabb_ptrs;
         aabb_ptrs.reserve(num_motion_steps);
+#  if OPTIX_ABI_VERSION >= 36
+        vector<device_ptr> width_ptrs;
+        vector<device_ptr> vertex_ptrs;
+        width_ptrs.reserve(num_motion_steps);
+        vertex_ptrs.reserve(num_motion_steps);
+#  endif
         for (size_t step = 0; step < num_motion_steps; ++step) {
           aabb_ptrs.push_back(aabb_data.device_pointer + step * num_segments * sizeof(OptixAabb));
+#  if OPTIX_ABI_VERSION >= 36
+          const device_ptr base_ptr = vertex_data.device_pointer +
+                                      step * num_vertices * sizeof(float4);
+          width_ptrs.push_back(base_ptr + 3 * sizeof(float));  // Offset by vertex size
+          vertex_ptrs.push_back(base_ptr);
+#  endif
         }
 
-        // Disable visibility test anyhit program, since it is already checked during intersection
+        // Force a single any-hit call, so shadow record-all behavior works correctly
-        // Those trace calls that require anyhit can force it with OPTIX_RAY_FLAG_ENFORCE_ANYHIT
+        unsigned int build_flags = OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;
-        unsigned int build_flags = OPTIX_GEOMETRY_FLAG_DISABLE_ANYHIT;
         OptixBuildInput build_input = {};
-        build_input.type = OPTIX_BUILD_INPUT_TYPE_CUSTOM_PRIMITIVES;
+#  if OPTIX_ABI_VERSION >= 36
-        build_input.aabbArray.aabbBuffers = (CUdeviceptr *)aabb_ptrs.data();
+        if (DebugFlags().optix.curves_api && hair->curve_shape == CURVE_THICK) {
-        build_input.aabbArray.numPrimitives = num_segments;
+          build_input.type = OPTIX_BUILD_INPUT_TYPE_CURVES;
-        build_input.aabbArray.strideInBytes = sizeof(OptixAabb);
+          build_input.curveArray.curveType = OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE;
-        build_input.aabbArray.flags = &build_flags;
+          build_input.curveArray.numPrimitives = num_segments;
-        build_input.aabbArray.numSbtRecords = 1;
+          build_input.curveArray.vertexBuffers = (CUdeviceptr *)vertex_ptrs.data();
-        build_input.aabbArray.primitiveIndexOffset = hair->optix_prim_offset;
+          build_input.curveArray.numVertices = num_vertices;
+          build_input.curveArray.vertexStrideInBytes = sizeof(float4);
+          build_input.curveArray.widthBuffers = (CUdeviceptr *)width_ptrs.data();
+          build_input.curveArray.widthStrideInBytes = sizeof(float4);
+          build_input.curveArray.indexBuffer = (CUdeviceptr)index_data.device_pointer;
+          build_input.curveArray.indexStrideInBytes = sizeof(int);
+          build_input.curveArray.flag = build_flags;
+          build_input.curveArray.primitiveIndexOffset = hair->optix_prim_offset;
+        }
+        else
+#  endif
+        {
+          // Disable visibility test any-hit program, since it is already checked during
+          // intersection. Those trace calls that require anyhit can force it with a ray flag.
+          build_flags |= OPTIX_GEOMETRY_FLAG_DISABLE_ANYHIT;
+
+          build_input.type = OPTIX_BUILD_INPUT_TYPE_CUSTOM_PRIMITIVES;
+#  if OPTIX_ABI_VERSION < 23
+          build_input.aabbArray.aabbBuffers = (CUdeviceptr *)aabb_ptrs.data();
+          build_input.aabbArray.numPrimitives = num_segments;
+          build_input.aabbArray.strideInBytes = sizeof(OptixAabb);
+          build_input.aabbArray.flags = &build_flags;
+          build_input.aabbArray.numSbtRecords = 1;
+          build_input.aabbArray.primitiveIndexOffset = hair->optix_prim_offset;
+#  else
+          build_input.customPrimitiveArray.aabbBuffers = (CUdeviceptr *)aabb_ptrs.data();
+          build_input.customPrimitiveArray.numPrimitives = num_segments;
+          build_input.customPrimitiveArray.strideInBytes = sizeof(OptixAabb);
+          build_input.customPrimitiveArray.flags = &build_flags;
+          build_input.customPrimitiveArray.numSbtRecords = 1;
+          build_input.customPrimitiveArray.primitiveIndexOffset = hair->optix_prim_offset;
+#  endif
+        }
 
         // Allocate memory for new BLAS and build it
         OptixTraversableHandle handle;
@@ -1257,8 +1433,8 @@ class OptiXDevice : public CUDADevice {
           vertex_ptrs.push_back(vertex_data.device_pointer + num_verts * step * sizeof(float3));
         }
 
-        // No special build flags for triangle primitives
+        // Force a single any-hit call, so shadow record-all behavior works correctly
-        unsigned int build_flags = OPTIX_GEOMETRY_FLAG_NONE;
+        unsigned int build_flags = OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;
         OptixBuildInput build_input = {};
         build_input.type = OPTIX_BUILD_INPUT_TYPE_TRIANGLES;
         build_input.triangleArray.vertexBuffers = (CUdeviceptr *)vertex_ptrs.data();
@@ -1324,9 +1500,26 @@ class OptiXDevice : public CUDADevice {
       // Set user instance ID to object index
       instance.instanceId = ob->get_device_index();
 
-      // Volumes have a special bit set in the visibility mask so a trace can mask only volumes
+      // Have to have at least one bit in the mask, or else instance would always be culled
-      // See 'scene_intersect_volume' in bvh.h
+      instance.visibilityMask = 1;
-      instance.visibilityMask = (ob->geometry->has_volume ? 3 : 1);
+
+      if (ob->geometry->has_volume) {
+        // Volumes have a special bit set in the visibility mask so a trace can mask only volumes
+        instance.visibilityMask |= 2;
+      }
+
+      if (ob->geometry->type == Geometry::HAIR) {
+        // Same applies to curves (so they can be skipped in local trace calls)
+        instance.visibilityMask |= 4;
+
+#  if OPTIX_ABI_VERSION >= 36
+        if (motion_blur && ob->geometry->has_motion_blur() && DebugFlags().optix.curves_api &&
+            static_cast<const Hair *>(ob->geometry)->curve_shape == CURVE_THICK) {
+          // Select between motion blur and non-motion blur built-in intersection module
+          instance.sbtOffset = PG_HITD_MOTION - PG_HITD;
+        }
+#  endif
+      }
 
       // Insert motion traversable if object has motion
       if (motion_blur && ob->use_motion()) {

intern/cycles/device/device_task.h

@@ -29,6 +29,7 @@ CCL_NAMESPACE_BEGIN
 class Device;
 class RenderBuffers;
 class RenderTile;
+class RenderTileNeighbors;
 class Tile;
 
 enum DenoiserType {
@@ -41,6 +42,14 @@ enum DenoiserType {
   DENOISER_ALL = ~0,
 };
 
+enum DenoiserInput {
+  DENOISER_INPUT_RGB = 1,
+  DENOISER_INPUT_RGB_ALBEDO = 2,
+  DENOISER_INPUT_RGB_ALBEDO_NORMAL = 3,
+
+  DENOISER_INPUT_NUM,
+};
+
 typedef int DenoiserTypeMask;
 
 class DenoiseParams {
@@ -72,10 +81,10 @@ class DenoiseParams {
   /* Clamp the input to the range of +-1e8. Should be enough for any legitimate data. */
   bool clamp_input;
 
-  /** Optix Denoiser **/
+  /** OIDN/Optix Denoiser **/
 
-  /* Passes handed over to the OptiX denoiser (default to color + albedo). */
+  /* Passes handed over to the OIDN/OptiX denoiser (default to color + albedo). */
-  int optix_input_passes;
+  DenoiserInput input_passes;
 
   DenoiseParams()
   {
@@ -91,7 +100,7 @@ class DenoiseParams {
     neighbor_frames = 2;
     clamp_input = true;
 
-    optix_input_passes = 2;
+    input_passes = DENOISER_INPUT_RGB_ALBEDO_NORMAL;
 
     start_sample = 0;
   }
@@ -150,8 +159,8 @@ class DeviceTask {
   function<void(RenderTile &)> update_tile_sample;
   function<void(RenderTile &)> release_tile;
   function<bool()> get_cancel;
-  function<void(RenderTile *, Device *)> map_neighbor_tiles;
+  function<void(RenderTileNeighbors &, Device *)> map_neighbor_tiles;
-  function<void(RenderTile *, Device *)> unmap_neighbor_tiles;
+  function<void(RenderTileNeighbors &, Device *)> unmap_neighbor_tiles;
 
   uint tile_types;
   DenoiseParams denoising;

intern/cycles/device/opencl/device_opencl_impl.cpp

@@ -1850,7 +1850,7 @@ void OpenCLDevice::denoise(RenderTile &rtile, DenoisingTask &denoising)
   denoising.render_buffer.samples = rtile.sample;
   denoising.buffer.gpu_temporary_mem = true;
 
-  denoising.run_denoising(&rtile);
+  denoising.run_denoising(rtile);
 }
 
 void OpenCLDevice::shader(DeviceTask &task)

intern/cycles/device/opencl/memory_manager.cpp

@@ -64,6 +64,9 @@ void MemoryManager::DeviceBuffer::update_device_memory(OpenCLDevice *device)
     total_size += alloc_size;
   }
 
+  /* Always allocate non-empty buffer, NULL pointers cause problems with some drivers. */
+  total_size = max(total_size, 16);
+
   if (need_realloc) {
     cl_ulong max_buffer_size;
     clGetDeviceInfo(
@@ -251,7 +254,7 @@ void MemoryManager::set_kernel_arg_buffers(cl_kernel kernel, cl_uint *narg)
       device->kernel_set_args(kernel, (*narg)++, *device_buffer.buffer);
     }
     else {
-      device->kernel_set_args(kernel, (*narg)++, 0);
+      device->kernel_set_args(kernel, (*narg)++);
     }
   }
 }

intern/cycles/kernel/bvh/bvh.h

@@ -172,11 +172,11 @@ ccl_device_intersect bool scene_intersect(KernelGlobals *kg,
              0.0f,
              ray->t,
              ray->time,
-             0xFF,
+             0xF,
              OPTIX_RAY_FLAG_NONE,
-             0,
-             0,
              0,  // SBT offset for PG_HITD
+             0,
+             0,
              p0,
              p1,
              p2,
@@ -264,12 +264,13 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals *kg,
              0.0f,
              ray->t,
              ray->time,
+             // Skip curves
+             0x3,
              // Need to always call into __anyhit__kernel_optix_local_hit
-             0xFF,
              OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
-             1,
+             2,  // SBT offset for PG_HITL
+             0,
              0,
-             0,  // SBT offset for PG_HITL
              p0,
              p1,
              p2,
@@ -374,12 +375,12 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals *kg,
              0.0f,
              ray->t,
              ray->time,
+             0xF,
              // Need to always call into __anyhit__kernel_optix_shadow_all_hit
-             0xFF,
              OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
-             2,
+             1,  // SBT offset for PG_HITS
+             0,
              0,
-             0,  // SBT offset for PG_HITS
              p0,
              p1,
              *num_hits,
@@ -458,12 +459,12 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals *kg,
              0.0f,
              ray->t,
              ray->time,
-             // Visibility mask set to only intersect objects with volumes
+             // Skip everything but volumes
-             0x02,
+             0x2,
              OPTIX_RAY_FLAG_NONE,
-             0,
-             0,
              0,  // SBT offset for PG_HITD
+             0,
+             0,
              p0,
              p1,
              p2,

intern/cycles/kernel/geom/geom_curve_intersect.h

@@ -734,7 +734,6 @@ ccl_device_inline void curve_shader_setup(KernelGlobals *kg,
   }
 
   sd->u = isect->u;
-  sd->v = isect->v;
 
   P = P + D * t;
 
@@ -750,6 +749,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals *kg,
 
     sd->N = normalize(sine * bitangent - cosine * normalize(cross(tangent, bitangent)));
     sd->Ng = -D;
+    sd->v = isect->v;
 
 #  if 0
     /* This approximates the position and geometric normal of a thick curve too,
@@ -764,8 +764,11 @@ ccl_device_inline void curve_shader_setup(KernelGlobals *kg,
      * This could be optimized by recording the normal in the intersection,
      * however for Optix this would go beyond the size of the payload. */
     const float3 P_inside = float4_to_float3(catmull_rom_basis_eval(P_curve, isect->u));
-    sd->Ng = normalize(P - P_inside);
+    const float3 Ng = normalize(P - P_inside);
-    sd->N = sd->Ng;
+
+    sd->N = Ng;
+    sd->Ng = Ng;
+    sd->v = 0.0f;
   }
 
 #  ifdef __DPDU__

intern/cycles/kernel/kernel_camera.h

@@ -379,7 +379,7 @@ ccl_device_inline void camera_sample(KernelGlobals *kg,
     const int shutter_table_offset = kernel_data.cam.shutter_table_offset;
     ray->time = lookup_table_read(kg, time, shutter_table_offset, SHUTTER_TABLE_SIZE);
     /* TODO(sergey): Currently single rolling shutter effect type only
-     * where scan-lines are acquired from top to bottom and whole scanline
+     * where scan-lines are acquired from top to bottom and whole scan-line
      * is acquired at once (no delay in acquisition happens between pixels
      * of single scan-line).
      *

intern/cycles/kernel/kernel_path.h

@@ -284,19 +284,11 @@ ccl_device_forceinline bool kernel_path_shader_apply(KernelGlobals *kg,
 #ifdef __HOLDOUT__
   if (((sd->flag & SD_HOLDOUT) || (sd->object_flag & SD_OBJECT_HOLDOUT_MASK)) &&
       (state->flag & PATH_RAY_TRANSPARENT_BACKGROUND)) {
+    const float3 holdout_weight = shader_holdout_apply(kg, sd);
     if (kernel_data.background.transparent) {
-      float3 holdout_weight;
-      if (sd->object_flag & SD_OBJECT_HOLDOUT_MASK) {
-        holdout_weight = make_float3(1.0f, 1.0f, 1.0f);
-      }
-      else {
-        holdout_weight = shader_holdout_eval(kg, sd);
-      }
-      /* any throughput is ok, should all be identical here */
       L->transparent += average(holdout_weight * throughput);
     }
-
+    if (isequal_float3(holdout_weight, make_float3(1.0f, 1.0f, 1.0f))) {
-    if (sd->object_flag & SD_OBJECT_HOLDOUT_MASK) {
       return false;
     }
   }

intern/cycles/kernel/kernel_shader.h

@@ -1017,15 +1017,36 @@ ccl_device float3 shader_emissive_eval(ShaderData *sd)
 
 /* Holdout */
 
-ccl_device float3 shader_holdout_eval(KernelGlobals *kg, ShaderData *sd)
+ccl_device float3 shader_holdout_apply(KernelGlobals *kg, ShaderData *sd)
 {
   float3 weight = make_float3(0.0f, 0.0f, 0.0f);
 
-  for (int i = 0; i < sd->num_closure; i++) {
+  /* For objects marked as holdout, preserve transparency and remove all other
-    ShaderClosure *sc = &sd->closure[i];
+   * closures, replacing them with a holdout weight. */
+  if (sd->object_flag & SD_OBJECT_HOLDOUT_MASK) {
+    if ((sd->flag & SD_TRANSPARENT) && !(sd->flag & SD_HAS_ONLY_VOLUME)) {
+      weight = make_float3(1.0f, 1.0f, 1.0f) - sd->closure_transparent_extinction;
 
-    if (CLOSURE_IS_HOLDOUT(sc->type))
+      for (int i = 0; i < sd->num_closure; i++) {
-      weight += sc->weight;
+        ShaderClosure *sc = &sd->closure[i];
+        if (!CLOSURE_IS_BSDF_TRANSPARENT(sc->type)) {
+          sc->type = NBUILTIN_CLOSURES;
+        }
+      }
+
+      sd->flag &= ~(SD_CLOSURE_FLAGS - (SD_TRANSPARENT | SD_BSDF));
+    }
+    else {
+      weight = make_float3(1.0f, 1.0f, 1.0f);
+    }
+  }
+  else {
+    for (int i = 0; i < sd->num_closure; i++) {
+      ShaderClosure *sc = &sd->closure[i];
+      if (CLOSURE_IS_HOLDOUT(sc->type)) {
+        weight += sc->weight;
+      }
+    }
   }
 
   return weight;

intern/cycles/kernel/kernels/optix/kernel_optix.cu

@@ -15,6 +15,7 @@
  * limitations under the License.
  */
 
+// clang-format off
 #include "kernel/kernel_compat_optix.h"
 #include "util/util_atomic.h"
 #include "kernel/kernel_types.h"
@@ -23,6 +24,7 @@
 
 #include "kernel/kernel_path.h"
 #include "kernel/kernel_bake.h"
+// clang-format on
 
 template<typename T> ccl_device_forceinline T *get_payload_ptr_0()
 {
@@ -139,8 +141,8 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
   }
   else {
     if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) {
-      // Record closest intersection only (do not terminate ray here, since there is no guarantee
+      // Record closest intersection only
-      // about distance ordering in anyhit)
+      // Do not terminate ray here, since there is no guarantee about distance ordering in any-hit
       return optixIgnoreIntersection();
     }
 
@@ -153,15 +155,9 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
   isect->object = get_object_id();
   isect->type = kernel_tex_fetch(__prim_type, isect->prim);
 
-  if (optixIsTriangleHit()) {
+  const float2 barycentrics = optixGetTriangleBarycentrics();
-    const float2 barycentrics = optixGetTriangleBarycentrics();
+  isect->u = 1.0f - barycentrics.y - barycentrics.x;
-    isect->u = 1.0f - barycentrics.y - barycentrics.x;
+  isect->v = barycentrics.x;
-    isect->v = barycentrics.x;
-  }
-  else {
-    isect->u = __uint_as_float(optixGetAttribute_0());
-    isect->v = __uint_as_float(optixGetAttribute_1());
-  }
 
   // Record geometric normal
   const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim);
@@ -198,10 +194,18 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
     isect->u = 1.0f - barycentrics.y - barycentrics.x;
     isect->v = barycentrics.x;
   }
+#  ifdef __HAIR__
   else {
-    isect->u = __uint_as_float(optixGetAttribute_0());
+    const float u = __uint_as_float(optixGetAttribute_0());
+    isect->u = u;
     isect->v = __uint_as_float(optixGetAttribute_1());
+
+    // Filter out curve endcaps
+    if (u == 0.0f || u == 1.0f) {
+      return optixIgnoreIntersection();
+    }
   }
+#  endif
 
 #  ifdef __TRANSPARENT_SHADOWS__
   // Detect if this surface has a shader with transparent shadows
@@ -213,7 +217,6 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
 #  ifdef __TRANSPARENT_SHADOWS__
   }
 
-  // TODO(pmours): Do we need REQUIRE_UNIQUE_ANYHIT for this to work?
   optixSetPayload_2(optixGetPayload_2() + 1);  // num_hits++
 
   // Continue tracing
@@ -227,13 +230,25 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
   uint visibility = optixGetPayload_4();
 #ifdef __VISIBILITY_FLAG__
   const uint prim = optixGetPrimitiveIndex();
-  if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0)
+  if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
     return optixIgnoreIntersection();
+  }
+#endif
+
+#ifdef __HAIR__
+  if (!optixIsTriangleHit()) {
+    // Filter out curve endcaps
+    const float u = __uint_as_float(optixGetAttribute_0());
+    if (u == 0.0f || u == 1.0f) {
+      return optixIgnoreIntersection();
+    }
+  }
 #endif
 
   // Shadow ray early termination
-  if (visibility & PATH_RAY_SHADOW_OPAQUE)
+  if (visibility & PATH_RAY_SHADOW_OPAQUE) {
     return optixTerminateRay();
+  }
 }
 
 extern "C" __global__ void __closesthit__kernel_optix_hit()
@@ -250,7 +265,7 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
     optixSetPayload_2(__float_as_uint(barycentrics.x));
   }
   else {
-    optixSetPayload_1(optixGetAttribute_0());
+    optixSetPayload_1(optixGetAttribute_0());  // Same as 'optixGetCurveParameter()'
     optixSetPayload_2(optixGetAttribute_1());
   }
 }
@@ -286,7 +301,6 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
                             __float_as_int(isect.u),   // Attribute_0
                             __float_as_int(isect.v));  // Attribute_1
   }
-
 }
 
 extern "C" __global__ void __intersection__curve_ribbon()

intern/cycles/kernel/shaders/node_sky_texture.osl

@@ -122,12 +122,18 @@ vector geographical_to_direction(float lat, float lon)
   return vector(cos(lat) * cos(lon), cos(lat) * sin(lon), sin(lat));
 }
 
-color sky_radiance_nishita(vector dir, float nishita_data[9], string filename)
+float precise_angle(vector a, vector b)
+{
+  return 2.0 * atan2(length(a - b), length(a + b));
+}
+
+color sky_radiance_nishita(vector dir, float nishita_data[10], string filename)
 {
   /* definitions */
   float sun_elevation = nishita_data[6];
   float sun_rotation = nishita_data[7];
   float angular_diameter = nishita_data[8];
+  float sun_intensity = nishita_data[9];
   int sun_disc = angular_diameter > 0;
   float alpha = 1.0;
   color xyz;
@@ -138,13 +144,13 @@ color sky_radiance_nishita(vector dir, float nishita_data[9], string filename)
   if (dir[2] >= 0.0) {
     /* definitions */
     vector sun_dir = geographical_to_direction(sun_elevation, sun_rotation + M_PI_2);
-    float sun_dir_angle = acos(dot(dir, sun_dir));
+    float sun_dir_angle = precise_angle(dir, sun_dir);
     float half_angular = angular_diameter / 2.0;
     float dir_elevation = M_PI_2 - direction[0];
 
     /* if ray inside sun disc render it, otherwise render sky */
     if (sun_dir_angle < half_angular && sun_disc == 1) {
-      /* get 3 pixels data */
+      /* get 2 pixels data */
       color pixel_bottom = color(nishita_data[0], nishita_data[1], nishita_data[2]);
       color pixel_top = color(nishita_data[3], nishita_data[4], nishita_data[5]);
       float y;
@@ -153,13 +159,13 @@ color sky_radiance_nishita(vector dir, float nishita_data[9], string filename)
       if (sun_elevation - half_angular > 0.0) {
         if ((sun_elevation + half_angular) > 0.0) {
           y = ((dir_elevation - sun_elevation) / angular_diameter) + 0.5;
-          xyz = mix(pixel_bottom, pixel_top, y);
+          xyz = mix(pixel_bottom, pixel_top, y) * sun_intensity;
         }
       }
       else {
         if (sun_elevation + half_angular > 0.0) {
           y = dir_elevation / (sun_elevation + half_angular);
-          xyz = mix(pixel_bottom, pixel_top, y);
+          xyz = mix(pixel_bottom, pixel_top, y) * sun_intensity;
         }
       }
       /* limb darkening, coefficient is 0.6f */
@@ -171,7 +177,8 @@ color sky_radiance_nishita(vector dir, float nishita_data[9], string filename)
     else {
       /* sky interpolation */
       float x = (direction[1] + M_PI + sun_rotation) / M_2PI;
-      float y = 1.0 - (dir_elevation / M_PI_2);
+      /* more pixels toward horizon compensation */
+      float y = 1.0 - sqrt(dir_elevation / M_PI_2);
       if (x > 1.0) {
         x = x - 1.0;
       }
@@ -201,19 +208,20 @@ color sky_radiance_nishita(vector dir, float nishita_data[9], string filename)
   return xyz_to_rgb(xyz[0], xyz[1], xyz[2]) * 120000.0;
 }
 
-shader node_sky_texture(int use_mapping = 0,
+shader node_sky_texture(
-                        matrix mapping = matrix(0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0),
+    int use_mapping = 0,
-                        vector Vector = P,
+    matrix mapping = matrix(0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0),
-                        string type = "hosek_wilkie",
+    vector Vector = P,
-                        float theta = 0.0,
+    string type = "hosek_wilkie",
-                        float phi = 0.0,
+    float theta = 0.0,
-                        string filename = "",
+    float phi = 0.0,
-                        color radiance = color(0.0, 0.0, 0.0),
+    string filename = "",
-                        float config_x[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
+    color radiance = color(0.0, 0.0, 0.0),
-                        float config_y[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
+    float config_x[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
-                        float config_z[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
+    float config_y[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
-                        float nishita_data[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
+    float config_z[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
-                        output color Color = color(0.0, 0.0, 0.0))
+    float nishita_data[10] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
+    output color Color = color(0.0, 0.0, 0.0))
 {
   vector p = Vector;
 

intern/cycles/kernel/svm/svm_sky.h

@@ -136,6 +136,7 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals *kg,
   float sun_elevation = nishita_data[6];
   float sun_rotation = nishita_data[7];
   float angular_diameter = nishita_data[8];
+  float sun_intensity = nishita_data[9];
   bool sun_disc = (angular_diameter > 0.0f);
   float3 xyz;
   /* convert dir to spherical coordinates */
@@ -145,13 +146,13 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals *kg,
   if (dir.z >= 0.0f) {
     /* definitions */
     float3 sun_dir = geographical_to_direction(sun_elevation, sun_rotation + M_PI_2_F);
-    float sun_dir_angle = acos(dot(dir, sun_dir));
+    float sun_dir_angle = precise_angle(dir, sun_dir);
     float half_angular = angular_diameter / 2.0f;
     float dir_elevation = M_PI_2_F - direction.x;
 
     /* if ray inside sun disc render it, otherwise render sky */
     if (sun_disc && sun_dir_angle < half_angular) {
-      /* get 3 pixels data */
+      /* get 2 pixels data */
       float3 pixel_bottom = make_float3(nishita_data[0], nishita_data[1], nishita_data[2]);
       float3 pixel_top = make_float3(nishita_data[3], nishita_data[4], nishita_data[5]);
       float y;
@@ -160,13 +161,13 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals *kg,
       if (sun_elevation - half_angular > 0.0f) {
         if (sun_elevation + half_angular > 0.0f) {
           y = ((dir_elevation - sun_elevation) / angular_diameter) + 0.5f;
-          xyz = interp(pixel_bottom, pixel_top, y);
+          xyz = interp(pixel_bottom, pixel_top, y) * sun_intensity;
         }
       }
       else {
         if (sun_elevation + half_angular > 0.0f) {
           y = dir_elevation / (sun_elevation + half_angular);
-          xyz = interp(pixel_bottom, pixel_top, y);
+          xyz = interp(pixel_bottom, pixel_top, y) * sun_intensity;
         }
       }
       /* limb darkening, coefficient is 0.6f */
@@ -178,7 +179,8 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals *kg,
     else {
       /* sky interpolation */
       float x = (direction.y + M_PI_F + sun_rotation) / M_2PI_F;
-      float y = dir_elevation / M_PI_2_F;
+      /* more pixels toward horizon compensation */
+      float y = safe_sqrtf(dir_elevation / M_PI_2_F);
       if (x > 1.0f) {
         x -= 1.0f;
       }
@@ -301,7 +303,7 @@ ccl_device void svm_node_tex_sky(
   /* Nishita */
   else {
     /* Define variables */
-    float nishita_data[9];
+    float nishita_data[10];
 
     float4 data = read_node_float(kg, offset);
     nishita_data[0] = data.x;
@@ -317,7 +319,8 @@ ccl_device void svm_node_tex_sky(
 
     data = read_node_float(kg, offset);
     nishita_data[8] = data.x;
-    uint texture_id = __float_as_uint(data.y);
+    nishita_data[9] = data.y;
+    uint texture_id = __float_as_uint(data.z);
 
     /* Compute Sky */
     f = sky_radiance_nishita(kg, dir, nishita_data, texture_id);

intern/cycles/render/CMakeLists.txt

@@ -2,6 +2,7 @@
 set(INC
   ..
   ../../glew-mx
+  ../../sky/include
 )
 
 set(INC_SYS
@@ -92,6 +93,7 @@ set(LIB
   cycles_device
   cycles_subd
   cycles_util
+  bf_intern_sky
 )
 
 if(WITH_CYCLES_OSL)

intern/cycles/render/buffers.h

@@ -52,7 +52,7 @@ class BufferParams {
   /* passes */
   vector<Pass> passes;
   bool denoising_data_pass;
-  /* If only some light path types should be denoised, an additional pass is needed. */
+  /* If only some light path types should be target, an additional pass is needed. */
   bool denoising_clean_pass;
   /* When we're prefiltering the passes during rendering, we need to keep both the
    * original and the prefiltered data around because neighboring tiles might still
@@ -149,6 +149,50 @@ class RenderTile {
   RenderBuffers *buffers;
 
   RenderTile();
+
+  int4 bounds() const
+  {
+    return make_int4(x,      /* xmin */
+                     y,      /* ymin */
+                     x + w,  /* xmax */
+                     y + h); /* ymax */
+  }
+};
+
+/* Render Tile Neighbors
+ * Set of neighboring tiles used for denoising. Tile order:
+ *  0 1 2
+ *  3 4 5
+ *  6 7 8 */
+
+class RenderTileNeighbors {
+ public:
+  static const int SIZE = 9;
+  static const int CENTER = 4;
+
+  RenderTile tiles[SIZE];
+  RenderTile target;
+
+  RenderTileNeighbors(const RenderTile &center)
+  {
+    tiles[CENTER] = center;
+  }
+
+  int4 bounds() const
+  {
+    return make_int4(tiles[3].x,               /* xmin */
+                     tiles[1].y,               /* ymin */
+                     tiles[5].x + tiles[5].w,  /* xmax */
+                     tiles[7].y + tiles[7].h); /* ymax */
+  }
+
+  void set_bounds_from_center()
+  {
+    tiles[3].x = tiles[CENTER].x;
+    tiles[1].y = tiles[CENTER].y;
+    tiles[5].x = tiles[CENTER].x + tiles[CENTER].w;
+    tiles[7].y = tiles[CENTER].y + tiles[CENTER].h;
+  }
 };
 
 CCL_NAMESPACE_END

intern/cycles/render/camera.cpp

@@ -26,6 +26,7 @@
 #include "util/util_function.h"
 #include "util/util_logging.h"
 #include "util/util_math_cdf.h"
+#include "util/util_task.h"
 #include "util/util_vector.h"
 
 /* needed for calculating differentials */
@@ -496,20 +497,35 @@ void Camera::device_update_volume(Device * /*device*/, DeviceScene *dscene, Scen
   if (!need_device_update && !need_flags_update) {
     return;
   }
-  KernelCamera *kcam = &dscene->data.cam;
+
-  BoundBox viewplane_boundbox = viewplane_bounds_get();
+  KernelIntegrator *kintegrator = &dscene->data.integrator;
-  for (size_t i = 0; i < scene->objects.size(); ++i) {
+  if (kintegrator->use_volumes) {
-    Object *object = scene->objects[i];
+    KernelCamera *kcam = &dscene->data.cam;
-    if (object->geometry->has_volume && viewplane_boundbox.intersects(object->bounds)) {
+    BoundBox viewplane_boundbox = viewplane_bounds_get();
-      /* TODO(sergey): Consider adding more grained check. */
+
-      VLOG(1) << "Detected camera inside volume.";
+    /* Parallel object update, with grain size to avoid too much threading overhead
-      kcam->is_inside_volume = 1;
+     * for individual objects. */
-      break;
+    static const int OBJECTS_PER_TASK = 32;
+    parallel_for(blocked_range<size_t>(0, scene->objects.size(), OBJECTS_PER_TASK),
+                 [&](const blocked_range<size_t> &r) {
+                   for (size_t i = r.begin(); i != r.end(); i++) {
+                     Object *object = scene->objects[i];
+                     if (object->geometry->has_volume &&
+                         viewplane_boundbox.intersects(object->bounds)) {
+                       /* TODO(sergey): Consider adding more grained check. */
+                       VLOG(1) << "Detected camera inside volume.";
+                       kcam->is_inside_volume = 1;
+                       parallel_for_cancel();
+                       break;
+                     }
+                   }
+                 });
+
+    if (!kcam->is_inside_volume) {
+      VLOG(1) << "Camera is outside of the volume.";
     }
   }
-  if (!kcam->is_inside_volume) {
+
-    VLOG(1) << "Camera is outside of the volume.";
-  }
   need_device_update = false;
   need_flags_update = false;
 }

intern/cycles/render/colorspace.cpp

@@ -266,7 +266,7 @@ inline void processor_apply_pixels(const OCIO::Processor *processor, T *pixels,
 {
   /* TODO: implement faster version for when we know the conversion
    * is a simple matrix transform between linear spaces. In that case
-   * unpremultiply is not needed. */
+   * un-premultiply is not needed. */
 
   /* Process large images in chunks to keep temporary memory requirement down. */
   const size_t chunk_size = std::min((size_t)(16 * 1024 * 1024), num_pixels);
@@ -354,7 +354,7 @@ void ColorSpaceManager::to_scene_linear(ColorSpaceProcessor *processor_,
         processor->applyRGB(pixel);
       }
       else {
-        /* Unassociate and associate alpha since color management should not
+        /* Un-associate and associate alpha since color management should not
          * be affected by transparency. */
         float alpha = pixel[3];
         float inv_alpha = 1.0f / alpha;

intern/cycles/render/denoising.cpp

@@ -271,42 +271,45 @@ bool DenoiseTask::acquire_tile(Device *device, Device *tile_device, RenderTile &
  *
  * However, since there is only one large memory, the denoised result has to be written to
  * a different buffer to avoid having to copy an entire horizontal slice of the image. */
-void DenoiseTask::map_neighboring_tiles(RenderTile *tiles, Device *tile_device)
+void DenoiseTask::map_neighboring_tiles(RenderTileNeighbors &neighbors, Device *tile_device)
 {
+  RenderTile &center_tile = neighbors.tiles[RenderTileNeighbors::CENTER];
+  RenderTile &target_tile = neighbors.target;
+
   /* Fill tile information. */
-  for (int i = 0; i < 9; i++) {
+  for (int i = 0; i < RenderTileNeighbors::SIZE; i++) {
-    if (i == 4) {
+    if (i == RenderTileNeighbors::CENTER) {
       continue;
     }
 
+    RenderTile &tile = neighbors.tiles[i];
     int dx = (i % 3) - 1;
     int dy = (i / 3) - 1;
-    tiles[i].x = clamp(tiles[4].x + dx * denoiser->tile_size.x, 0, image.width);
+    tile.x = clamp(center_tile.x + dx * denoiser->tile_size.x, 0, image.width);
-    tiles[i].w = clamp(tiles[4].x + (dx + 1) * denoiser->tile_size.x, 0, image.width) - tiles[i].x;
+    tile.w = clamp(center_tile.x + (dx + 1) * denoiser->tile_size.x, 0, image.width) - tile.x;
-    tiles[i].y = clamp(tiles[4].y + dy * denoiser->tile_size.y, 0, image.height);
+    tile.y = clamp(center_tile.y + dy * denoiser->tile_size.y, 0, image.height);
-    tiles[i].h = clamp(tiles[4].y + (dy + 1) * denoiser->tile_size.y, 0, image.height) -
+    tile.h = clamp(center_tile.y + (dy + 1) * denoiser->tile_size.y, 0, image.height) - tile.y;
-                 tiles[i].y;
 
-    tiles[i].buffer = tiles[4].buffer;
+    tile.buffer = center_tile.buffer;
-    tiles[i].offset = tiles[4].offset;
+    tile.offset = center_tile.offset;
-    tiles[i].stride = image.width;
+    tile.stride = image.width;
   }
 
   /* Allocate output buffer. */
   device_vector<float> *output_mem = new device_vector<float>(
       tile_device, "denoising_output", MEM_READ_WRITE);
-  output_mem->alloc(OUTPUT_NUM_CHANNELS * tiles[4].w * tiles[4].h);
+  output_mem->alloc(OUTPUT_NUM_CHANNELS * center_tile.w * center_tile.h);
 
   /* Fill output buffer with noisy image, assumed by kernel_filter_finalize
    * when skipping denoising of some pixels. */
   float *result = output_mem->data();
-  float *in = &image.pixels[image.num_channels * (tiles[4].y * image.width + tiles[4].x)];
+  float *in = &image.pixels[image.num_channels * (center_tile.y * image.width + center_tile.x)];
 
   const DenoiseImageLayer &layer = image.layers[current_layer];
   const int *input_to_image_channel = layer.input_to_image_channel.data();
 
-  for (int y = 0; y < tiles[4].h; y++) {
+  for (int y = 0; y < center_tile.h; y++) {
-    for (int x = 0; x < tiles[4].w; x++, result += OUTPUT_NUM_CHANNELS) {
+    for (int x = 0; x < center_tile.w; x++, result += OUTPUT_NUM_CHANNELS) {
       for (int i = 0; i < OUTPUT_NUM_CHANNELS; i++) {
         result[i] = in[image.num_channels * x + input_to_image_channel[INPUT_NOISY_IMAGE + i]];
       }
@@ -317,35 +320,38 @@ void DenoiseTask::map_neighboring_tiles(RenderTile *tiles, Device *tile_device)
   output_mem->copy_to_device();
 
   /* Fill output tile info. */
-  tiles[9] = tiles[4];
+  target_tile = center_tile;
-  tiles[9].buffer = output_mem->device_pointer;
+  target_tile.buffer = output_mem->device_pointer;
-  tiles[9].stride = tiles[9].w;
+  target_tile.stride = target_tile.w;
-  tiles[9].offset -= tiles[9].x + tiles[9].y * tiles[9].stride;
+  target_tile.offset -= target_tile.x + target_tile.y * target_tile.stride;
 
   thread_scoped_lock output_lock(output_mutex);
-  assert(output_pixels.count(tiles[4].tile_index) == 0);
+  assert(output_pixels.count(center_tile.tile_index) == 0);
-  output_pixels[tiles[9].tile_index] = output_mem;
+  output_pixels[target_tile.tile_index] = output_mem;
 }
 
-void DenoiseTask::unmap_neighboring_tiles(RenderTile *tiles)
+void DenoiseTask::unmap_neighboring_tiles(RenderTileNeighbors &neighbors)
 {
+  RenderTile &center_tile = neighbors.tiles[RenderTileNeighbors::CENTER];
+  RenderTile &target_tile = neighbors.target;
+
   thread_scoped_lock output_lock(output_mutex);
-  assert(output_pixels.count(tiles[4].tile_index) == 1);
+  assert(output_pixels.count(center_tile.tile_index) == 1);
-  device_vector<float> *output_mem = output_pixels[tiles[9].tile_index];
+  device_vector<float> *output_mem = output_pixels[target_tile.tile_index];
-  output_pixels.erase(tiles[4].tile_index);
+  output_pixels.erase(center_tile.tile_index);
   output_lock.unlock();
 
   /* Copy denoised pixels from device. */
-  output_mem->copy_from_device(0, OUTPUT_NUM_CHANNELS * tiles[9].w, tiles[9].h);
+  output_mem->copy_from_device(0, OUTPUT_NUM_CHANNELS * target_tile.w, target_tile.h);
 
   float *result = output_mem->data();
-  float *out = &image.pixels[image.num_channels * (tiles[9].y * image.width + tiles[9].x)];
+  float *out = &image.pixels[image.num_channels * (target_tile.y * image.width + target_tile.x)];
 
   const DenoiseImageLayer &layer = image.layers[current_layer];
   const int *output_to_image_channel = layer.output_to_image_channel.data();
 
-  for (int y = 0; y < tiles[9].h; y++) {
+  for (int y = 0; y < target_tile.h; y++) {
-    for (int x = 0; x < tiles[9].w; x++, result += OUTPUT_NUM_CHANNELS) {
+    for (int x = 0; x < target_tile.w; x++, result += OUTPUT_NUM_CHANNELS) {
       for (int i = 0; i < OUTPUT_NUM_CHANNELS; i++) {
         out[image.num_channels * x + output_to_image_channel[i]] = result[i];
       }

intern/cycles/render/denoising.h

@@ -196,8 +196,8 @@ class DenoiseTask {
 
   /* Device task callbacks */
   bool acquire_tile(Device *device, Device *tile_device, RenderTile &tile);
-  void map_neighboring_tiles(RenderTile *tiles, Device *tile_device);
+  void map_neighboring_tiles(RenderTileNeighbors &neighbors, Device *tile_device);
-  void unmap_neighboring_tiles(RenderTile *tiles);
+  void unmap_neighboring_tiles(RenderTileNeighbors &neighbors);
   void release_tile();
   bool get_cancel();
 };

intern/cycles/render/graph.cpp

@@ -942,12 +942,12 @@ void ShaderGraph::refine_bump_nodes()
       foreach (NodePair &pair, nodes_dy)
         add(pair.second);
 
-      /* connect what is connected is bump to samplecenter input*/
+      /* Connect what is connected is bump to sample-center input. */
       connect(out, node->input("SampleCenter"));
 
-      /* bump input is just for connectivity purpose for the graph input,
+      /* Bump input is just for connectivity purpose for the graph input,
-       * we re-connected this input to samplecenter, so lets disconnect it
+       * we re-connected this input to sample-center, so lets disconnect it
-       * from bump input */
+       * from bump input. */
       disconnect(bump_input);
     }
   }

intern/cycles/render/image_sky.cpp

@@ -16,16 +16,20 @@
 
 #include "render/image_sky.h"
 
+#include "sky_model.h"
+
 #include "util/util_image.h"
 #include "util/util_logging.h"
 #include "util/util_path.h"
-#include "util/util_sky_model.h"
 #include "util/util_task.h"
 
 CCL_NAMESPACE_BEGIN
 
-SkyLoader::SkyLoader(
+SkyLoader::SkyLoader(float sun_elevation,
-    float sun_elevation, int altitude, float air_density, float dust_density, float ozone_density)
+                     float altitude,
+                     float air_density,
+                     float dust_density,
+                     float ozone_density)
     : sun_elevation(sun_elevation),
       altitude(altitude),
       air_density(air_density),
@@ -56,23 +60,22 @@ bool SkyLoader::load_pixels(const ImageMetaData &metadata,
   int width = metadata.width;
   int height = metadata.height;
   float *pixel_data = (float *)pixels;
-  float altitude_f = (float)altitude;
 
   /* precompute sky texture */
   const int rows_per_task = divide_up(1024, width);
   parallel_for(blocked_range<size_t>(0, height, rows_per_task),
                [&](const blocked_range<size_t> &r) {
-                 nishita_skymodel_precompute_texture(pixel_data,
+                 SKY_nishita_skymodel_precompute_texture(pixel_data,
-                                                     metadata.channels,
+                                                         metadata.channels,
-                                                     r.begin(),
+                                                         r.begin(),
-                                                     r.end(),
+                                                         r.end(),
-                                                     width,
+                                                         width,
-                                                     height,
+                                                         height,
-                                                     sun_elevation,
+                                                         sun_elevation,
-                                                     altitude_f,
+                                                         altitude,
-                                                     air_density,
+                                                         air_density,
-                                                     dust_density,
+                                                         dust_density,
-                                                     ozone_density);
+                                                         ozone_density);
                });
 
   return true;

intern/cycles/render/image_sky.h

@@ -21,14 +21,14 @@ CCL_NAMESPACE_BEGIN
 class SkyLoader : public ImageLoader {
  private:
   float sun_elevation;
-  int altitude;
+  float altitude;
   float air_density;
   float dust_density;
   float ozone_density;
 
  public:
   SkyLoader(float sun_elevation,
-            int altitude,
+            float altitude,
             float air_density,
             float dust_density,
             float ozone_density);

intern/cycles/render/light.cpp

@@ -625,13 +625,19 @@ void LightManager::device_update_background(Device *device,
           }
         }
 
+        /* Determine sun direction from lat/long and texture mapping. */
         float latitude = sky->sun_elevation;
         float longitude = M_2PI_F - sky->sun_rotation + M_PI_2_F;
+        float3 sun_direction = make_float3(
+            cosf(latitude) * cosf(longitude), cosf(latitude) * sinf(longitude), sinf(latitude));
+        Transform sky_transform = transform_inverse(sky->tex_mapping.compute_transform());
+        sun_direction = transform_direction(&sky_transform, sun_direction);
+
+        /* Pack sun direction and size. */
         float half_angle = sky->sun_size * 0.5f;
-        kbackground->sun = make_float4(cosf(latitude) * cosf(longitude),
+        kbackground->sun = make_float4(
-                                       cosf(latitude) * sinf(longitude),
+            sun_direction.x, sun_direction.y, sun_direction.z, half_angle);
-                                       sinf(latitude),
+
-                                       half_angle);
         kbackground->sun_weight = 4.0f;
         environment_res.x = max(environment_res.x, 512);
         environment_res.y = max(environment_res.y, 256);

intern/cycles/render/nodes.cpp

@@ -27,9 +27,10 @@
 #include "render/scene.h"
 #include "render/svm.h"
 
+#include "sky_model.h"
+
 #include "util/util_foreach.h"
 #include "util/util_logging.h"
-#include "util/util_sky_model.h"
 #include "util/util_transform.h"
 
 #include "kernel/svm/svm_color_util.h"
@@ -631,7 +632,7 @@ typedef struct SunSky {
 
   /* Parameter */
   float radiance_x, radiance_y, radiance_z;
-  float config_x[9], config_y[9], config_z[9], nishita_data[9];
+  float config_x[9], config_y[9], config_z[9], nishita_data[10];
 } SunSky;
 
 /* Preetham model */
@@ -726,8 +727,8 @@ static void sky_texture_precompute_hosek(SunSky *sunsky,
   float solarElevation = M_PI_2_F - theta;
 
   /* Initialize Sky Model */
-  ArHosekSkyModelState *sky_state;
+  SKY_ArHosekSkyModelState *sky_state;
-  sky_state = arhosek_xyz_skymodelstate_alloc_init(
+  sky_state = SKY_arhosek_xyz_skymodelstate_alloc_init(
       (double)turbidity, (double)ground_albedo, (double)solarElevation);
 
   /* Copy values from sky_state to SunSky */
@@ -741,25 +742,31 @@ static void sky_texture_precompute_hosek(SunSky *sunsky,
   sunsky->radiance_z = (float)sky_state->radiances[2];
 
   /* Free sky_state */
-  arhosekskymodelstate_free(sky_state);
+  SKY_arhosekskymodelstate_free(sky_state);
 }
 
 /* Nishita improved */
 static void sky_texture_precompute_nishita(SunSky *sunsky,
                                            bool sun_disc,
                                            float sun_size,
+                                           float sun_intensity,
                                            float sun_elevation,
                                            float sun_rotation,
-                                           int altitude,
+                                           float altitude,
                                            float air_density,
                                            float dust_density)
 {
   /* sample 2 sun pixels */
   float pixel_bottom[3];
   float pixel_top[3];
-  float altitude_f = (float)altitude;
+  SKY_nishita_skymodel_precompute_sun(
-  nishita_skymodel_precompute_sun(
+      sun_elevation, sun_size, altitude, air_density, dust_density, pixel_bottom, pixel_top);
-      sun_elevation, sun_size, altitude_f, air_density, dust_density, pixel_bottom, pixel_top);
+  /* limit sun rotation between 0 and 360 degrees */
+  sun_rotation = fmodf(sun_rotation, M_2PI_F);
+  if (sun_rotation < 0.0f) {
+    sun_rotation += M_2PI_F;
+  }
+  sun_rotation = M_2PI_F - sun_rotation;
   /* send data to svm_sky */
   sunsky->nishita_data[0] = pixel_bottom[0];
   sunsky->nishita_data[1] = pixel_bottom[1];
@@ -768,8 +775,9 @@ static void sky_texture_precompute_nishita(SunSky *sunsky,
   sunsky->nishita_data[4] = pixel_top[1];
   sunsky->nishita_data[5] = pixel_top[2];
   sunsky->nishita_data[6] = sun_elevation;
-  sunsky->nishita_data[7] = M_2PI_F - sun_rotation;
+  sunsky->nishita_data[7] = sun_rotation;
   sunsky->nishita_data[8] = sun_disc ? sun_size : 0.0f;
+  sunsky->nishita_data[9] = sun_intensity;
 }
 
 NODE_DEFINE(SkyTextureNode)
@@ -789,9 +797,10 @@ NODE_DEFINE(SkyTextureNode)
   SOCKET_FLOAT(ground_albedo, "Ground Albedo", 0.3f);
   SOCKET_BOOLEAN(sun_disc, "Sun Disc", true);
   SOCKET_FLOAT(sun_size, "Sun Size", 0.009512f);
+  SOCKET_FLOAT(sun_intensity, "Sun Intensity", 1.0f);
   SOCKET_FLOAT(sun_elevation, "Sun Elevation", M_PI_2_F);
   SOCKET_FLOAT(sun_rotation, "Sun Rotation", 0.0f);
-  SOCKET_INT(altitude, "Altitude", 0);
+  SOCKET_FLOAT(altitude, "Altitude", 1.0f);
   SOCKET_FLOAT(air_density, "Air", 1.0f);
   SOCKET_FLOAT(dust_density, "Dust", 1.0f);
   SOCKET_FLOAT(ozone_density, "Ozone", 1.0f);
@@ -819,12 +828,17 @@ void SkyTextureNode::compile(SVMCompiler &compiler)
   else if (type == NODE_SKY_HOSEK)
     sky_texture_precompute_hosek(&sunsky, sun_direction, turbidity, ground_albedo);
   else if (type == NODE_SKY_NISHITA) {
+    /* Clamp altitude to reasonable values.
+     * Below 1m causes numerical issues and above 60km is space. */
+    float clamped_altitude = clamp(altitude, 1.0f, 59999.0f);
+
     sky_texture_precompute_nishita(&sunsky,
                                    sun_disc,
                                    sun_size,
+                                   sun_intensity,
                                    sun_elevation,
                                    sun_rotation,
-                                   altitude,
+                                   clamped_altitude,
                                    air_density,
                                    dust_density);
     /* precomputed texture image parameters */
@@ -836,7 +850,7 @@ void SkyTextureNode::compile(SVMCompiler &compiler)
     /* precompute sky texture */
     if (handle.empty()) {
       SkyLoader *loader = new SkyLoader(
-          sun_elevation, altitude, air_density, dust_density, ozone_density);
+          sun_elevation, clamped_altitude, air_density, dust_density, ozone_density);
       handle = image_manager->add_image(loader, impar);
     }
   }
@@ -891,7 +905,10 @@ void SkyTextureNode::compile(SVMCompiler &compiler)
                       __float_as_uint(sunsky.nishita_data[5]),
                       __float_as_uint(sunsky.nishita_data[6]),
                       __float_as_uint(sunsky.nishita_data[7]));
-    compiler.add_node(__float_as_uint(sunsky.nishita_data[8]), handle.svm_slot(), 0, 0);
+    compiler.add_node(__float_as_uint(sunsky.nishita_data[8]),
+                      __float_as_uint(sunsky.nishita_data[9]),
+                      handle.svm_slot(),
+                      0);
   }
 
   tex_mapping.compile_end(compiler, vector_in, vector_offset);
@@ -907,12 +924,17 @@ void SkyTextureNode::compile(OSLCompiler &compiler)
   else if (type == NODE_SKY_HOSEK)
     sky_texture_precompute_hosek(&sunsky, sun_direction, turbidity, ground_albedo);
   else if (type == NODE_SKY_NISHITA) {
+    /* Clamp altitude to reasonable values.
+     * Below 1m causes numerical issues and above 60km is space. */
+    float clamped_altitude = clamp(altitude, 1.0f, 59999.0f);
+
     sky_texture_precompute_nishita(&sunsky,
                                    sun_disc,
                                    sun_size,
+                                   sun_intensity,
                                    sun_elevation,
                                    sun_rotation,
-                                   altitude,
+                                   clamped_altitude,
                                    air_density,
                                    dust_density);
     /* precomputed texture image parameters */
@@ -924,7 +946,7 @@ void SkyTextureNode::compile(OSLCompiler &compiler)
     /* precompute sky texture */
     if (handle.empty()) {
       SkyLoader *loader = new SkyLoader(
-          sun_elevation, altitude, air_density, dust_density, ozone_density);
+          sun_elevation, clamped_altitude, air_density, dust_density, ozone_density);
       handle = image_manager->add_image(loader, impar);
     }
   }
@@ -939,7 +961,7 @@ void SkyTextureNode::compile(OSLCompiler &compiler)
   compiler.parameter_array("config_x", sunsky.config_x, 9);
   compiler.parameter_array("config_y", sunsky.config_y, 9);
   compiler.parameter_array("config_z", sunsky.config_z, 9);
-  compiler.parameter_array("nishita_data", sunsky.nishita_data, 9);
+  compiler.parameter_array("nishita_data", sunsky.nishita_data, 10);
   /* nishita texture */
   if (type == NODE_SKY_NISHITA) {
     compiler.parameter_texture("filename", handle.svm_slot());

intern/cycles/render/nodes.h

@@ -170,9 +170,10 @@ class SkyTextureNode : public TextureNode {
   float ground_albedo;
   bool sun_disc;
   float sun_size;
+  float sun_intensity;
   float sun_elevation;
   float sun_rotation;
-  int altitude;
+  float altitude;
   float air_density;
   float dust_density;
   float ozone_density;

intern/cycles/render/object.cpp

@@ -823,6 +823,12 @@ void ObjectManager::apply_static_transforms(DeviceScene *dscene, Scene *scene, P
       Mesh *mesh = static_cast<Mesh *>(geom);
       apply = apply && mesh->subdivision_type == Mesh::SUBDIVISION_NONE;
     }
+    else if (geom->type == Geometry::HAIR) {
+      /* Can't apply non-uniform scale to curves, this can't be represented by
+       * control points and radius alone. */
+      float scale;
+      apply = apply && transform_uniform_scale(object->tfm, scale);
+    }
 
     if (apply) {
       if (!(motion_blur && object->use_motion())) {

intern/cycles/render/session.cpp

@@ -536,7 +536,7 @@ void Session::release_tile(RenderTile &rtile, const bool need_denoise)
   denoising_cond.notify_all();
 }
 
-void Session::map_neighbor_tiles(RenderTile *tiles, Device *tile_device)
+void Session::map_neighbor_tiles(RenderTileNeighbors &neighbors, Device *tile_device)
 {
   thread_scoped_lock tile_lock(tile_mutex);
 
@@ -546,75 +546,77 @@ void Session::map_neighbor_tiles(RenderTile *tiles, Device *tile_device)
       tile_manager.state.buffer.full_x + tile_manager.state.buffer.width,
       tile_manager.state.buffer.full_y + tile_manager.state.buffer.height);
 
+  RenderTile &center_tile = neighbors.tiles[RenderTileNeighbors::CENTER];
+
   if (!tile_manager.schedule_denoising) {
     /* Fix up tile slices with overlap. */
     if (tile_manager.slice_overlap != 0) {
-      int y = max(tiles[4].y - tile_manager.slice_overlap, image_region.y);
+      int y = max(center_tile.y - tile_manager.slice_overlap, image_region.y);
-      tiles[4].h = min(tiles[4].y + tiles[4].h + tile_manager.slice_overlap, image_region.w) - y;
+      center_tile.h = min(center_tile.y + center_tile.h + tile_manager.slice_overlap,
-      tiles[4].y = y;
+                          image_region.w) -
+                      y;
+      center_tile.y = y;
     }
 
     /* Tiles are not being denoised individually, which means the entire image is processed. */
-    tiles[3].x = tiles[4].x;
+    neighbors.set_bounds_from_center();
-    tiles[1].y = tiles[4].y;
-    tiles[5].x = tiles[4].x + tiles[4].w;
-    tiles[7].y = tiles[4].y + tiles[4].h;
   }
   else {
-    int center_idx = tiles[4].tile_index;
+    int center_idx = center_tile.tile_index;
     assert(tile_manager.state.tiles[center_idx].state == Tile::DENOISE);
 
     for (int dy = -1, i = 0; dy <= 1; dy++) {
       for (int dx = -1; dx <= 1; dx++, i++) {
+        RenderTile &rtile = neighbors.tiles[i];
         int nindex = tile_manager.get_neighbor_index(center_idx, i);
         if (nindex >= 0) {
           Tile *tile = &tile_manager.state.tiles[nindex];
 
-          tiles[i].x = image_region.x + tile->x;
+          rtile.x = image_region.x + tile->x;
-          tiles[i].y = image_region.y + tile->y;
+          rtile.y = image_region.y + tile->y;
-          tiles[i].w = tile->w;
+          rtile.w = tile->w;
-          tiles[i].h = tile->h;
+          rtile.h = tile->h;
 
           if (buffers) {
-            tile_manager.state.buffer.get_offset_stride(tiles[i].offset, tiles[i].stride);
+            tile_manager.state.buffer.get_offset_stride(rtile.offset, rtile.stride);
 
-            tiles[i].buffer = buffers->buffer.device_pointer;
+            rtile.buffer = buffers->buffer.device_pointer;
-            tiles[i].buffers = buffers;
+            rtile.buffers = buffers;
           }
           else {
             assert(tile->buffers);
-            tile->buffers->params.get_offset_stride(tiles[i].offset, tiles[i].stride);
+            tile->buffers->params.get_offset_stride(rtile.offset, rtile.stride);
 
-            tiles[i].buffer = tile->buffers->buffer.device_pointer;
+            rtile.buffer = tile->buffers->buffer.device_pointer;
-            tiles[i].buffers = tile->buffers;
+            rtile.buffers = tile->buffers;
           }
         }
         else {
-          int px = tiles[4].x + dx * params.tile_size.x;
+          int px = center_tile.x + dx * params.tile_size.x;
-          int py = tiles[4].y + dy * params.tile_size.y;
+          int py = center_tile.y + dy * params.tile_size.y;
 
-          tiles[i].x = clamp(px, image_region.x, image_region.z);
+          rtile.x = clamp(px, image_region.x, image_region.z);
-          tiles[i].y = clamp(py, image_region.y, image_region.w);
+          rtile.y = clamp(py, image_region.y, image_region.w);
-          tiles[i].w = tiles[i].h = 0;
+          rtile.w = rtile.h = 0;
 
-          tiles[i].buffer = (device_ptr)NULL;
+          rtile.buffer = (device_ptr)NULL;
-          tiles[i].buffers = NULL;
+          rtile.buffers = NULL;
         }
       }
     }
   }
 
-  assert(tiles[4].buffers);
+  assert(center_tile.buffers);
-  device->map_neighbor_tiles(tile_device, tiles);
+  device->map_neighbor_tiles(tile_device, neighbors);
 
   /* The denoised result is written back to the original tile. */
-  tiles[9] = tiles[4];
+  neighbors.target = center_tile;
 }
 
-void Session::unmap_neighbor_tiles(RenderTile *tiles, Device *tile_device)
+void Session::unmap_neighbor_tiles(RenderTileNeighbors &neighbors, Device *tile_device)
 {
   thread_scoped_lock tile_lock(tile_mutex);
-  device->unmap_neighbor_tiles(tile_device, tiles);
+  device->unmap_neighbor_tiles(tile_device, neighbors);
 }
 
 void Session::run_cpu()
@@ -1003,7 +1005,7 @@ bool Session::update_scene()
   int height = tile_manager.state.buffer.full_height;
   int resolution = tile_manager.state.resolution_divider;
 
-  if (width != cam->width || height != cam->height) {
+  if (width != cam->width || height != cam->height || resolution != cam->resolution) {
     cam->width = width;
     cam->height = height;
     cam->resolution = resolution;
@@ -1126,6 +1128,11 @@ bool Session::render_need_denoise(bool &delayed)
 {
   delayed = false;
 
+  /* Not supported yet for baking. */
+  if (read_bake_tile_cb) {
+    return false;
+  }
+
   /* Denoising enabled? */
   if (!params.denoising.need_denoising_task()) {
     return false;

intern/cycles/render/session.h

@@ -198,8 +198,8 @@ class Session {
   void update_tile_sample(RenderTile &tile);
   void release_tile(RenderTile &tile, const bool need_denoise);
 
-  void map_neighbor_tiles(RenderTile *tiles, Device *tile_device);
+  void map_neighbor_tiles(RenderTileNeighbors &neighbors, Device *tile_device);
-  void unmap_neighbor_tiles(RenderTile *tiles, Device *tile_device);
+  void unmap_neighbor_tiles(RenderTileNeighbors &neighbors, Device *tile_device);
 
   bool device_use_gl;
 

intern/cycles/util/CMakeLists.txt

@@ -98,10 +98,6 @@ set(SRC_HEADERS
   util_rect.h
   util_set.h
   util_simd.h
-  util_sky_model.cpp
-  util_sky_model.h
-  util_sky_model_data.h
-  util_sky_nishita.cpp
   util_avxf.h
   util_avxb.h
   util_semaphore.h

intern/cycles/util/util_debug.cpp

@@ -83,6 +83,7 @@ DebugFlags::OptiX::OptiX()
 void DebugFlags::OptiX::reset()
 {
   cuda_streams = 1;
+  curves_api = false;
 }
 
 DebugFlags::OpenCL::OpenCL() : device_type(DebugFlags::OpenCL::DEVICE_ALL), debug(false)

intern/cycles/util/util_debug.h

@@ -108,6 +108,9 @@ class DebugFlags {
 
     /* Number of CUDA streams to launch kernels concurrently from. */
     int cuda_streams;
+
+    /* Use OptiX curves API for hair instead of custom implementation. */
+    bool curves_api;
   };
 
   /* Descriptor of OpenCL feature-set to be used. */

intern/cycles/util/util_math.h

@@ -787,6 +787,16 @@ ccl_device_inline float compare_floats(float a, float b, float abs_diff, int ulp
   return (abs(__float_as_int(a) - __float_as_int(b)) < ulp_diff);
 }
 
+/* Calculate the angle between the two vectors a and b.
+ * The usual approach acos(dot(a, b)) has severe precision issues for small angles,
+ * which are avoided by this method.
+ * Based on "Mangled Angles" from https://people.eecs.berkeley.edu/~wkahan/Mindless.pdf
+ */
+ccl_device_inline float precise_angle(float3 a, float3 b)
+{
+  return 2.0f * atan2f(len(a - b), len(a + b));
+}
+
 CCL_NAMESPACE_END
 
 #endif /* __UTIL_MATH_H__ */

intern/cycles/util/util_tbb.h

@@ -34,6 +34,11 @@ using tbb::blocked_range;
 using tbb::enumerable_thread_specific;
 using tbb::parallel_for;
 
+static inline void parallel_for_cancel()
+{
+  tbb::task::self().cancel_group_execution();
+}
+
 CCL_NAMESPACE_END
 
 #endif /* __UTIL_TBB_H__ */

intern/cycles/util/util_thread.h

@@ -29,9 +29,9 @@
 #  include <pthread.h>
 #endif
 
-#ifdef __APPLE__
+/* NOTE: Use tbb/spin_mutex.h instead of util_tbb.h because some of the TBB
-#  include <libkern/OSAtomic.h>
+ * functionality requires RTTI, which is disabled for OSL kernel. */
-#endif
+#include <tbb/spin_mutex.h>
 
 #include "util/util_function.h"
 
@@ -65,76 +65,7 @@ class thread {
   int node_;
 };
 
-/* Own wrapper around pthread's spin lock to make it's use easier. */
+using thread_spin_lock = tbb::spin_mutex;
-
-class thread_spin_lock {
- public:
-#ifdef __APPLE__
-  inline thread_spin_lock()
-  {
-    spin_ = OS_SPINLOCK_INIT;
-  }
-
-  inline void lock()
-  {
-    OSSpinLockLock(&spin_);
-  }
-
-  inline void unlock()
-  {
-    OSSpinLockUnlock(&spin_);
-  }
-#elif defined(_WIN32)
-  inline thread_spin_lock()
-  {
-    const DWORD SPIN_COUNT = 50000;
-    InitializeCriticalSectionAndSpinCount(&cs_, SPIN_COUNT);
-  }
-
-  inline ~thread_spin_lock()
-  {
-    DeleteCriticalSection(&cs_);
-  }
-
-  inline void lock()
-  {
-    EnterCriticalSection(&cs_);
-  }
-
-  inline void unlock()
-  {
-    LeaveCriticalSection(&cs_);
-  }
-#else
-  inline thread_spin_lock()
-  {
-    pthread_spin_init(&spin_, 0);
-  }
-
-  inline ~thread_spin_lock()
-  {
-    pthread_spin_destroy(&spin_);
-  }
-
-  inline void lock()
-  {
-    pthread_spin_lock(&spin_);
-  }
-
-  inline void unlock()
-  {
-    pthread_spin_unlock(&spin_);
-  }
-#endif
- protected:
-#ifdef __APPLE__
-  OSSpinLock spin_;
-#elif defined(_WIN32)
-  CRITICAL_SECTION cs_;
-#else
-  pthread_spinlock_t spin_;
-#endif
-};
 
 class thread_scoped_spin_lock {
  public:

intern/ghost/GHOST_C-api.h

@@ -362,8 +362,8 @@ extern GHOST_TSuccess GHOST_HasCursorShape(GHOST_WindowHandle windowhandle,
  * \param mask The mask data for the cursor.
  * \param sizex The width of the cursor
  * \param sizey The height of the cursor
- * \param hotX The X coordinate of the cursor hotspot.
+ * \param hotX The X coordinate of the cursor hot-spot.
- * \param hotY The Y coordinate of the cursor hotspot.
+ * \param hotY The Y coordinate of the cursor hot-spot.
  * \param canInvertColor Let macOS invert cursor color to match platform convention.
  * \return Indication of success.
  */

intern/ghost/GHOST_IWindow.h

@@ -287,8 +287,8 @@ class GHOST_IWindow {
    * Set the shape of the cursor to a custom cursor.
    * \param   bitmap  The bitmap data for the cursor.
    * \param   mask    The mask data for the cursor.
-   * \param   hotX    The X coordinate of the cursor hotspot.
+   * \param   hotX    The X coordinate of the cursor hot-spot.
-   * \param   hotY    The Y coordinate of the cursor hotspot.
+   * \param   hotY    The Y coordinate of the cursor hot-spot.
    * \return  Indication of success.
    */
   virtual GHOST_TSuccess setCustomCursorShape(GHOST_TUns8 *bitmap,

intern/ghost/intern/GHOST_SystemX11.cpp

@@ -1324,7 +1324,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
         int revert_to;
 
         /* as ICCCM say, we need reply this event
-         * with a SetInputFocus, the data[1] have
+         * with a #SetInputFocus, the data[1] have
          * the valid timestamp (send by the wm).
          *
          * Some WM send this event before the
@@ -1345,7 +1345,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
       else {
 #ifdef WITH_XDND
         /* try to handle drag event
-         * (if there's no such events, GHOST_HandleClientMessage will return zero) */
+         * (if there's no such events, #GHOST_HandleClientMessage will return zero) */
         if (window->getDropTarget()->GHOST_HandleClientMessage(xe) == false) {
           /* Unknown client message, ignore */
         }
@@ -1366,12 +1366,12 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
 
     case EnterNotify:
     case LeaveNotify: {
-      /* XCrossingEvents pointer leave enter window.
+      /* #XCrossingEvents pointer leave enter window.
-       * also do cursor move here, MotionNotify only
+       * also do cursor move here, #MotionNotify only
        * happens when motion starts & ends inside window.
        * we only do moves when the crossing mode is 'normal'
-       * (really crossing between windows) since some windowmanagers
+       * (really crossing between windows) since some window-managers
-       * also send grab/ungrab crossings for mousewheel events.
+       * also send grab/un-grab crossings for mouse-wheel events.
        */
       XCrossingEvent &xce = xe->xcrossing;
       if (xce.mode == NotifyNormal) {
@@ -1396,11 +1396,11 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
     case MapNotify:
       /*
        * From ICCCM:
-       * [ Clients can select for StructureNotify on their
+       * [ Clients can select for #StructureNotify on their
        *   top-level windows to track transition between
-       *   Normal and Iconic states. Receipt of a MapNotify
+       *   Normal and Iconic states. Receipt of a #MapNotify
        *   event will indicate a transition to the Normal
-       *   state, and receipt of an UnmapNotify event will
+       *   state, and receipt of an #UnmapNotify event will
        *   indicate a transition to the Iconic state. ]
        */
       if (window->m_post_init == True) {
@@ -1441,7 +1441,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
       nxe.xselection.target = xse->target;
       nxe.xselection.time = xse->time;
 
-      /* Check to see if the requestor is asking for String */
+      /* Check to see if the requester is asking for String */
       if (xse->target == utf8_string || xse->target == string || xse->target == compound_text ||
           xse->target == c_string) {
         if (xse->selection == XInternAtom(m_display, "PRIMARY", False)) {
@@ -1487,7 +1487,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
         nxe.xselection.property = None;
       }
 
-      /* Send the event to the client 0 0 == False, SelectionNotify */
+      /* Send the event to the client 0 0 == False, #SelectionNotify */
       XSendEvent(m_display, xse->requestor, 0, 0, &nxe);
       XFlush(m_display);
       break;
@@ -1513,7 +1513,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
 
           /* Note: This event might be generated with incomplete data-set
            * (don't exactly know why, looks like in some cases, if the value does not change,
-           * it is not included in subsequent XDeviceMotionEvent events).
+           * it is not included in subsequent #XDeviceMotionEvent events).
            * So we have to check which values this event actually contains!
            */
 
@@ -1569,14 +1569,13 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
 GHOST_TSuccess GHOST_SystemX11::getModifierKeys(GHOST_ModifierKeys &keys) const
 {
 
-  /* Analyze the masks returned from XQueryPointer. */
+  /* Analyze the masks returned from #XQueryPointer. */
 
   memset((void *)m_keyboard_vector, 0, sizeof(m_keyboard_vector));
 
   XQueryKeymap(m_display, (char *)m_keyboard_vector);
 
-  /* now translate key symbols into keycodes and
+  /* Now translate key symbols into key-codes and test with vector. */
-   * test with vector. */
 
   const static KeyCode shift_l = XKeysymToKeycode(m_display, XK_Shift_L);
   const static KeyCode shift_r = XKeysymToKeycode(m_display, XK_Shift_R);
@@ -1671,7 +1670,7 @@ GHOST_TSuccess GHOST_SystemX11::setCursorPosition(GHOST_TInt32 x, GHOST_TInt32 y
 {
 
   /* This is a brute force move in screen coordinates
-   * XWarpPointer does relative moves so first determine the
+   * #XWarpPointer does relative moves so first determine the
    * current pointer position. */
 
   int cx, cy;

intern/ghost/intern/GHOST_Window.h

@@ -34,7 +34,7 @@ class GHOST_Context;
  * Dimensions are given in screen coordinates that are relative to the
  * upper-left corner of the screen.
  * Implements part of the GHOST_IWindow interface and adds some methods to
- * be implemented by childs of this class.
+ * be implemented by sub-classes of this class.
  */
 class GHOST_Window : public GHOST_IWindow {
  public:
@@ -124,8 +124,8 @@ class GHOST_Window : public GHOST_IWindow {
    * Set the shape of the cursor to a custom cursor.
    * \param   bitmap  The bitmap data for the cursor.
    * \param   mask    The mask data for the cursor.
-   * \param   hotX    The X coordinate of the cursor hotspot.
+   * \param   hotX    The X coordinate of the cursor hot-spot.
-   * \param   hotY    The Y coordinate of the cursor hotspot.
+   * \param   hotY    The Y coordinate of the cursor hot-spot.
    * \return  Indication of success.
    */
   GHOST_TSuccess setCustomCursorShape(GHOST_TUns8 *bitmap,

intern/guardedalloc/MEM_guardedalloc.h

@@ -239,6 +239,15 @@ void MEM_use_guarded_allocator(void);
     { \
       if (mem) \
         MEM_freeN(mem); \
+    } \
+    void *operator new(size_t /*count*/, void *ptr) \
+    { \
+      return ptr; \
+    } \
+    /* This is the matching delete operator to the placement-new operator above. Both parameters \
+     * will have the same value. Without this, we get the warning C4291 on windows. */ \
+    void operator delete(void * /*ptr_to_free*/, void * /*ptr*/) \
+    { \
     }
 
 /* Needed when type includes a namespace, then the namespace should not be

intern/guardedalloc/intern/mallocn_guarded_impl.c

@@ -198,10 +198,12 @@ print_error(const char *str, ...)
   va_end(ap);
   buf[sizeof(buf) - 1] = '\0';
 
-  if (error_callback)
+  if (error_callback) {
     error_callback(buf);
-  else
+  }
+  else {
     fputs(buf, stderr);
+  }
 }
 
 static pthread_mutex_t thread_lock = PTHREAD_MUTEX_INITIALIZER;
@@ -261,13 +263,16 @@ void *MEM_guarded_dupallocN(const void *vmemh)
     memh--;
 
 #ifndef DEBUG_MEMDUPLINAME
-    if (LIKELY(memh->alignment == 0))
+    if (LIKELY(memh->alignment == 0)) {
       newp = MEM_guarded_mallocN(memh->len, "dupli_alloc");
-    else
+    }
+    else {
       newp = MEM_guarded_mallocN_aligned(memh->len, (size_t)memh->alignment, "dupli_alloc");
+    }
 
-    if (newp == NULL)
+    if (newp == NULL) {
       return NULL;
+    }
 #else
     {
       MemHead *nmemh;
@@ -450,8 +455,9 @@ void *MEM_guarded_mallocN(size_t len, const char *str)
 
   if (LIKELY(memh)) {
     make_memhead_header(memh, len, str);
-    if (UNLIKELY(malloc_debug_memset && len))
+    if (UNLIKELY(malloc_debug_memset && len)) {
       memset(memh + 1, 255, len);
+    }
 
 #ifdef DEBUG_MEMCOUNTER
     if (_mallocn_count == DEBUG_MEMCOUNTER_ERROR_VAL)
@@ -522,8 +528,9 @@ void *MEM_guarded_mallocN_aligned(size_t len, size_t alignment, const char *str)
 
     make_memhead_header(memh, len, str);
     memh->alignment = (short)alignment;
-    if (UNLIKELY(malloc_debug_memset && len))
+    if (UNLIKELY(malloc_debug_memset && len)) {
       memset(memh + 1, 255, len);
+    }
 
 #ifdef DEBUG_MEMCOUNTER
     if (_mallocn_count == DEBUG_MEMCOUNTER_ERROR_VAL)
@@ -601,12 +608,15 @@ static int compare_len(const void *p1, const void *p2)
   const MemPrintBlock *pb1 = (const MemPrintBlock *)p1;
   const MemPrintBlock *pb2 = (const MemPrintBlock *)p2;
 
-  if (pb1->len < pb2->len)
+  if (pb1->len < pb2->len) {
     return 1;
-  else if (pb1->len == pb2->len)
+  }
+  else if (pb1->len == pb2->len) {
     return 0;
-  else
+  }
+  else {
     return -1;
+  }
 }
 
 void MEM_guarded_printmemlist_stats(void)
@@ -636,8 +646,9 @@ void MEM_guarded_printmemlist_stats(void)
   totpb = 0;
 
   membl = membase->first;
-  if (membl)
+  if (membl) {
     membl = MEMNEXT(membl);
+  }
 
   while (membl && pb) {
     pb->name = membl->name;
@@ -654,10 +665,12 @@ void MEM_guarded_printmemlist_stats(void)
     }
 #endif
 
-    if (membl->next)
+    if (membl->next) {
       membl = MEMNEXT(membl->next);
-    else
+    }
+    else {
       break;
+    }
   }
 
   /* sort by name and add together blocks with the same name */
@@ -737,8 +750,9 @@ static void MEM_guarded_printmemlist_internal(int pydict)
   mem_lock_thread();
 
   membl = membase->first;
-  if (membl)
+  if (membl) {
     membl = MEMNEXT(membl);
+  }
 
   if (pydict) {
     print_error("# membase_debug.py\n");
@@ -771,10 +785,12 @@ static void MEM_guarded_printmemlist_internal(int pydict)
       print_memhead_backtrace(membl);
 #endif
     }
-    if (membl->next)
+    if (membl->next) {
       membl = MEMNEXT(membl->next);
-    else
+    }
+    else {
       break;
+    }
   }
   if (pydict) {
     print_error("]\n\n");
@@ -791,15 +807,18 @@ void MEM_guarded_callbackmemlist(void (*func)(void *))
   mem_lock_thread();
 
   membl = membase->first;
-  if (membl)
+  if (membl) {
     membl = MEMNEXT(membl);
+  }
 
   while (membl) {
     func(membl + 1);
-    if (membl->next)
+    if (membl->next) {
       membl = MEMNEXT(membl->next);
-    else
+    }
+    else {
       break;
+    }
   }
 
   mem_unlock_thread();
@@ -890,24 +909,25 @@ void MEM_guarded_freeN(void *vmemh)
     MemorY_ErroR(memh->name, "end corrupt");
     name = check_memlist(memh);
     if (name != NULL) {
-      if (name != memh->name)
+      if (name != memh->name) {
         MemorY_ErroR(name, "is also corrupt");
+      }
     }
   }
   else {
     mem_lock_thread();
     name = check_memlist(memh);
     mem_unlock_thread();
-    if (name == NULL)
+    if (name == NULL) {
       MemorY_ErroR("free", "pointer not in memlist");
-    else
+    }
+    else {
       MemorY_ErroR(name, "error in header");
+    }
   }
 
   totblock--;
   /* here a DUMP should happen */
-
-  return;
 }
 
 /* --------------------------------------------------------------------- */
@@ -930,10 +950,12 @@ static void addtail(volatile localListBase *listbase, void *vlink)
   link->next = NULL;
   link->prev = listbase->last;
 
-  if (listbase->last)
+  if (listbase->last) {
     ((struct localLink *)listbase->last)->next = link;
-  if (listbase->first == NULL)
+  }
+  if (listbase->first == NULL) {
     listbase->first = link;
+  }
   listbase->last = link;
 }
 
@@ -950,15 +972,19 @@ static void remlink(volatile localListBase *listbase, void *vlink)
     return;
 #endif
 
-  if (link->next)
+  if (link->next) {
     link->next->prev = link->prev;
-  if (link->prev)
+  }
+  if (link->prev) {
     link->prev->next = link->next;
+  }
 
-  if (listbase->last == link)
+  if (listbase->last == link) {
     listbase->last = link->prev;
-  if (listbase->first == link)
+  }
+  if (listbase->first == link) {
     listbase->first = link->next;
+  }
 }
 
 static void rem_memblock(MemHead *memh)
@@ -966,10 +992,12 @@ static void rem_memblock(MemHead *memh)
   mem_lock_thread();
   remlink(membase, &memh->next);
   if (memh->prev) {
-    if (memh->next)
+    if (memh->next) {
       MEMNEXT(memh->prev)->nextname = MEMNEXT(memh->next)->name;
-    else
+    }
+    else {
       MEMNEXT(memh->prev)->nextname = NULL;
+    }
   }
   mem_unlock_thread();
 
@@ -981,8 +1009,9 @@ static void rem_memblock(MemHead *memh)
     free((char *)memh->name);
 #endif
 
-  if (UNLIKELY(malloc_debug_memset && memh->len))
+  if (UNLIKELY(malloc_debug_memset && memh->len)) {
     memset(memh + 1, 255, memh->len);
+  }
   if (LIKELY(memh->alignment == 0)) {
     free(memh);
   }
@@ -1006,78 +1035,100 @@ static const char *check_memlist(MemHead *memh)
   const char *name;
 
   forw = membase->first;
-  if (forw)
+  if (forw) {
     forw = MEMNEXT(forw);
+  }
   forwok = NULL;
   while (forw) {
-    if (forw->tag1 != MEMTAG1 || forw->tag2 != MEMTAG2)
+    if (forw->tag1 != MEMTAG1 || forw->tag2 != MEMTAG2) {
       break;
+    }
     forwok = forw;
-    if (forw->next)
+    if (forw->next) {
       forw = MEMNEXT(forw->next);
-    else
+    }
+    else {
       forw = NULL;
+    }
   }
 
   back = (MemHead *)membase->last;
-  if (back)
+  if (back) {
     back = MEMNEXT(back);
+  }
   backok = NULL;
   while (back) {
-    if (back->tag1 != MEMTAG1 || back->tag2 != MEMTAG2)
+    if (back->tag1 != MEMTAG1 || back->tag2 != MEMTAG2) {
       break;
+    }
     backok = back;
-    if (back->prev)
+    if (back->prev) {
       back = MEMNEXT(back->prev);
-    else
+    }
+    else {
       back = NULL;
+    }
   }
 
-  if (forw != back)
+  if (forw != back) {
     return ("MORE THAN 1 MEMORYBLOCK CORRUPT");
+  }
 
   if (forw == NULL && back == NULL) {
     /* no wrong headers found then but in search of memblock */
 
     forw = membase->first;
-    if (forw)
+    if (forw) {
       forw = MEMNEXT(forw);
+    }
     forwok = NULL;
     while (forw) {
-      if (forw == memh)
+      if (forw == memh) {
         break;
-      if (forw->tag1 != MEMTAG1 || forw->tag2 != MEMTAG2)
+      }
+      if (forw->tag1 != MEMTAG1 || forw->tag2 != MEMTAG2) {
         break;
+      }
       forwok = forw;
-      if (forw->next)
+      if (forw->next) {
         forw = MEMNEXT(forw->next);
-      else
+      }
+      else {
         forw = NULL;
+      }
     }
-    if (forw == NULL)
+    if (forw == NULL) {
       return NULL;
+    }
 
     back = (MemHead *)membase->last;
-    if (back)
+    if (back) {
       back = MEMNEXT(back);
+    }
     backok = NULL;
     while (back) {
-      if (back == memh)
+      if (back == memh) {
         break;
-      if (back->tag1 != MEMTAG1 || back->tag2 != MEMTAG2)
+      }
+      if (back->tag1 != MEMTAG1 || back->tag2 != MEMTAG2) {
         break;
+      }
       backok = back;
-      if (back->prev)
+      if (back->prev) {
         back = MEMNEXT(back->prev);
-      else
+      }
+      else {
         back = NULL;
+      }
     }
   }
 
-  if (forwok)
+  if (forwok) {
     name = forwok->nextname;
-  else
+  }
+  else {
     name = "No name found";
+  }
 
   if (forw == memh) {
     /* to be sure but this block is removed from the list */

intern/guardedalloc/test/simpletest/memtest.c

@@ -148,7 +148,7 @@ int main(int argc, char *argv[])
       fprintf(stderr, "|\n|--* Errors were detected\n");
     }
     else {
-      fprintf(stderr, "|\n|--* Test exited succesfully\n");
+      fprintf(stderr, "|\n|--* Test exited successfully\n");
     }
 
     fprintf(stderr, "|\n*** Finished test\n\n");

intern/libc_compat/libc_compat.c

@@ -28,6 +28,7 @@
 #  if defined(__GLIBC_PREREQ) && __GLIBC_PREREQ(2, 31)
 
 double __exp_finite(double x);
+double __exp2_finite(double x);
 double __acos_finite(double x);
 double __asin_finite(double x);
 double __log2_finite(double x);
@@ -35,6 +36,7 @@ double __log10_finite(double x);
 double __log_finite(double x);
 double __pow_finite(double x, double y);
 float __expf_finite(float x);
+float __exp2f_finite(float x);
 float __acosf_finite(float x);
 float __asinf_finite(float x);
 float __log2f_finite(float x);
@@ -47,6 +49,11 @@ double __exp_finite(double x)
   return exp(x);
 }
 
+double __exp2_finite(double x)
+{
+  return exp2(x);
+}
+
 double __acos_finite(double x)
 {
   return acos(x);
@@ -82,6 +89,11 @@ float __expf_finite(float x)
   return expf(x);
 }
 
+float __exp2f_finite(float x)
+{
+  return exp2f(x);
+}
+
 float __acosf_finite(float x)
 {
   return acosf(x);

intern/libmv/libmv/simple_pipeline/pipeline.cc

@@ -316,8 +316,8 @@ double InternalReprojectionError(
   }
   LG << "Skipped " << num_skipped << " markers.";
   LG << "Reprojected " << num_reprojected << " markers.";
-  LG << "Total error: " << total_error;
+  LG << "Total error: " << total_error << " px";
-  LG << "Average error: " << (total_error / num_reprojected) << " [pixels].";
+  LG << "Average error: " << (total_error / num_reprojected) << " px";
   return total_error / num_reprojected;
 }
 

intern/mantaflow/extern/manta_fluid_API.h

@@ -31,52 +31,52 @@ extern "C" {
 struct MANTA;
 
 /* Fluid functions */
-struct MANTA *manta_init(int *res, struct FluidModifierData *mmd);
+struct MANTA *manta_init(int *res, struct FluidModifierData *fmd);
 void manta_free(struct MANTA *fluid);
-void manta_ensure_obstacle(struct MANTA *fluid, struct FluidModifierData *mmd);
+void manta_ensure_obstacle(struct MANTA *fluid, struct FluidModifierData *fmd);
-void manta_ensure_guiding(struct MANTA *fluid, struct FluidModifierData *mmd);
+void manta_ensure_guiding(struct MANTA *fluid, struct FluidModifierData *fmd);
-void manta_ensure_invelocity(struct MANTA *fluid, struct FluidModifierData *mmd);
+void manta_ensure_invelocity(struct MANTA *fluid, struct FluidModifierData *fmd);
-void manta_ensure_outflow(struct MANTA *fluid, struct FluidModifierData *mmd);
+void manta_ensure_outflow(struct MANTA *fluid, struct FluidModifierData *fmd);
-int manta_write_config(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_write_config(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_write_data(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_write_data(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_write_noise(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_write_noise(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_read_config(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_read_config(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
 int manta_read_data(struct MANTA *fluid,
-                    struct FluidModifierData *mmd,
+                    struct FluidModifierData *fmd,
                     int framenr,
                     bool resumable);
 int manta_read_noise(struct MANTA *fluid,
-                     struct FluidModifierData *mmd,
+                     struct FluidModifierData *fmd,
                      int framenr,
                      bool resumable);
-int manta_read_mesh(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_read_mesh(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
 int manta_read_particles(struct MANTA *fluid,
-                         struct FluidModifierData *mmd,
+                         struct FluidModifierData *fmd,
                          int framenr,
                          bool resumable);
 int manta_read_guiding(struct MANTA *fluid,
-                       struct FluidModifierData *mmd,
+                       struct FluidModifierData *fmd,
                        int framenr,
                        bool sourceDomain);
-int manta_bake_data(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_bake_data(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_bake_noise(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_bake_noise(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_bake_mesh(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_bake_mesh(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_bake_particles(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_bake_particles(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_bake_guiding(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_bake_guiding(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_has_data(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_has_data(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_has_noise(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_has_noise(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_has_mesh(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_has_mesh(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
-int manta_has_particles(struct MANTA *fluid, struct FluidModifierData *mmd, int framenr);
+int manta_has_particles(struct MANTA *fluid, struct FluidModifierData *fmd, int framenr);
 int manta_has_guiding(struct MANTA *fluid,
-                      struct FluidModifierData *mmd,
+                      struct FluidModifierData *fmd,
                       int framenr,
                       bool domain);
 
-void manta_update_variables(struct MANTA *fluid, struct FluidModifierData *mmd);
+void manta_update_variables(struct MANTA *fluid, struct FluidModifierData *fmd);
 int manta_get_frame(struct MANTA *fluid);
 float manta_get_timestep(struct MANTA *fluid);
 void manta_adapt_timestep(struct MANTA *fluid);
-bool manta_needs_realloc(struct MANTA *fluid, struct FluidModifierData *mmd);
+bool manta_needs_realloc(struct MANTA *fluid, struct FluidModifierData *fmd);
 
 /* Fluid accessors */
 size_t manta_get_index(int x, int max_x, int y, int max_y, int z /*, int max_z */);
@@ -110,7 +110,7 @@ float *manta_get_phiout_in(struct MANTA *fluid);
 float *manta_get_phioutstatic_in(struct MANTA *fluid);
 
 /* Smoke functions */
-void manta_smoke_export_script(struct MANTA *smoke, struct FluidModifierData *mmd);
+void manta_smoke_export_script(struct MANTA *smoke, struct FluidModifierData *fmd);
 void manta_smoke_export(struct MANTA *smoke,
                         float *dt,
                         float *dx,
@@ -151,9 +151,9 @@ void manta_smoke_turbulence_get_rgba_fixed_color(struct MANTA *smoke,
                                                  float color[3],
                                                  float *data,
                                                  int sequential);
-void manta_smoke_ensure_heat(struct MANTA *smoke, struct FluidModifierData *mmd);
+void manta_smoke_ensure_heat(struct MANTA *smoke, struct FluidModifierData *fmd);
-void manta_smoke_ensure_fire(struct MANTA *smoke, struct FluidModifierData *mmd);
+void manta_smoke_ensure_fire(struct MANTA *smoke, struct FluidModifierData *fmd);
-void manta_smoke_ensure_colors(struct MANTA *smoke, struct FluidModifierData *mmd);
+void manta_smoke_ensure_colors(struct MANTA *smoke, struct FluidModifierData *fmd);
 
 /* Smoke accessors */
 float *manta_smoke_get_density(struct MANTA *smoke);
@@ -190,8 +190,8 @@ void manta_smoke_turbulence_get_res(struct MANTA *smoke, int *res);
 int manta_smoke_turbulence_get_cells(struct MANTA *smoke);
 
 /* Liquid functions */
-void manta_liquid_export_script(struct MANTA *smoke, struct FluidModifierData *mmd);
+void manta_liquid_export_script(struct MANTA *smoke, struct FluidModifierData *fmd);
-void manta_liquid_ensure_sndparts(struct MANTA *fluid, struct FluidModifierData *mmd);
+void manta_liquid_ensure_sndparts(struct MANTA *fluid, struct FluidModifierData *fmd);
 
 /* Liquid accessors */
 int manta_liquid_get_particle_res_x(struct MANTA *liquid);

intern/mantaflow/intern/MANTA_main.h

@@ -37,7 +37,7 @@ using std::vector;
 
 struct MANTA {
  public:
-  MANTA(int *res, struct FluidModifierData *mmd);
+  MANTA(int *res, struct FluidModifierData *fmd);
   MANTA(){};
   virtual ~MANTA();
 
@@ -61,68 +61,68 @@ struct MANTA {
   } Triangle;
 
   // Manta step, handling everything
-  void step(struct FluidModifierData *mmd, int startFrame);
+  void step(struct FluidModifierData *fmd, int startFrame);
 
   // Grid initialization functions
-  void initHeat(struct FluidModifierData *mmd = NULL);
+  void initHeat(struct FluidModifierData *fmd = NULL);
-  void initFire(struct FluidModifierData *mmd = NULL);
+  void initFire(struct FluidModifierData *fmd = NULL);
-  void initColors(struct FluidModifierData *mmd = NULL);
+  void initColors(struct FluidModifierData *fmd = NULL);
-  void initFireHigh(struct FluidModifierData *mmd = NULL);
+  void initFireHigh(struct FluidModifierData *fmd = NULL);
-  void initColorsHigh(struct FluidModifierData *mmd = NULL);
+  void initColorsHigh(struct FluidModifierData *fmd = NULL);
-  void initLiquid(FluidModifierData *mmd = NULL);
+  void initLiquid(FluidModifierData *fmd = NULL);
-  void initLiquidMesh(FluidModifierData *mmd = NULL);
+  void initLiquidMesh(FluidModifierData *fmd = NULL);
-  void initObstacle(FluidModifierData *mmd = NULL);
+  void initObstacle(FluidModifierData *fmd = NULL);
-  void initCurvature(FluidModifierData *mmd = NULL);
+  void initCurvature(FluidModifierData *fmd = NULL);
-  void initGuiding(FluidModifierData *mmd = NULL);
+  void initGuiding(FluidModifierData *fmd = NULL);
-  void initFractions(FluidModifierData *mmd = NULL);
+  void initFractions(FluidModifierData *fmd = NULL);
-  void initInVelocity(FluidModifierData *mmd = NULL);
+  void initInVelocity(FluidModifierData *fmd = NULL);
-  void initOutflow(FluidModifierData *mmd = NULL);
+  void initOutflow(FluidModifierData *fmd = NULL);
-  void initSndParts(FluidModifierData *mmd = NULL);
+  void initSndParts(FluidModifierData *fmd = NULL);
-  void initLiquidSndParts(FluidModifierData *mmd = NULL);
+  void initLiquidSndParts(FluidModifierData *fmd = NULL);
 
   // Pointer transfer: Mantaflow -> Blender
   void updatePointers();
 
   // Write cache
-  bool writeConfiguration(FluidModifierData *mmd, int framenr);
+  bool writeConfiguration(FluidModifierData *fmd, int framenr);
-  bool writeData(FluidModifierData *mmd, int framenr);
+  bool writeData(FluidModifierData *fmd, int framenr);
-  bool writeNoise(FluidModifierData *mmd, int framenr);
+  bool writeNoise(FluidModifierData *fmd, int framenr);
   // write calls for mesh and particles were left in bake calls for now
 
   // Read cache (via Manta save/load)
-  bool readConfiguration(FluidModifierData *mmd, int framenr);
+  bool readConfiguration(FluidModifierData *fmd, int framenr);
-  bool readData(FluidModifierData *mmd, int framenr, bool resumable);
+  bool readData(FluidModifierData *fmd, int framenr, bool resumable);
-  bool readNoise(FluidModifierData *mmd, int framenr, bool resumable);
+  bool readNoise(FluidModifierData *fmd, int framenr, bool resumable);
-  bool readMesh(FluidModifierData *mmd, int framenr);
+  bool readMesh(FluidModifierData *fmd, int framenr);
-  bool readParticles(FluidModifierData *mmd, int framenr, bool resumable);
+  bool readParticles(FluidModifierData *fmd, int framenr, bool resumable);
-  bool readGuiding(FluidModifierData *mmd, int framenr, bool sourceDomain);
+  bool readGuiding(FluidModifierData *fmd, int framenr, bool sourceDomain);
 
   // Read cache (via file read functions in MANTA - e.g. read .bobj.gz meshes, .uni particles)
-  bool updateMeshStructures(FluidModifierData *mmd, int framenr);
+  bool updateMeshStructures(FluidModifierData *fmd, int framenr);
-  bool updateFlipStructures(FluidModifierData *mmd, int framenr);
+  bool updateFlipStructures(FluidModifierData *fmd, int framenr);
-  bool updateParticleStructures(FluidModifierData *mmd, int framenr);
+  bool updateParticleStructures(FluidModifierData *fmd, int framenr);
-  bool updateSmokeStructures(FluidModifierData *mmd, int framenr);
+  bool updateSmokeStructures(FluidModifierData *fmd, int framenr);
-  bool updateNoiseStructures(FluidModifierData *mmd, int framenr);
+  bool updateNoiseStructures(FluidModifierData *fmd, int framenr);
-  bool updateVariables(FluidModifierData *mmd);
+  bool updateVariables(FluidModifierData *fmd);
 
   // Bake cache
-  bool bakeData(FluidModifierData *mmd, int framenr);
+  bool bakeData(FluidModifierData *fmd, int framenr);
-  bool bakeNoise(FluidModifierData *mmd, int framenr);
+  bool bakeNoise(FluidModifierData *fmd, int framenr);
-  bool bakeMesh(FluidModifierData *mmd, int framenr);
+  bool bakeMesh(FluidModifierData *fmd, int framenr);
-  bool bakeParticles(FluidModifierData *mmd, int framenr);
+  bool bakeParticles(FluidModifierData *fmd, int framenr);
-  bool bakeGuiding(FluidModifierData *mmd, int framenr);
+  bool bakeGuiding(FluidModifierData *fmd, int framenr);
 
   // IO for Mantaflow scene script
-  void exportSmokeScript(struct FluidModifierData *mmd);
+  void exportSmokeScript(struct FluidModifierData *fmd);
-  void exportLiquidScript(struct FluidModifierData *mmd);
+  void exportLiquidScript(struct FluidModifierData *fmd);
 
   // Check cache status by frame
-  bool hasConfig(FluidModifierData *mmd, int framenr);
+  bool hasConfig(FluidModifierData *fmd, int framenr);
-  bool hasData(FluidModifierData *mmd, int framenr);
+  bool hasData(FluidModifierData *fmd, int framenr);
-  bool hasNoise(FluidModifierData *mmd, int framenr);
+  bool hasNoise(FluidModifierData *fmd, int framenr);
-  bool hasMesh(FluidModifierData *mmd, int framenr);
+  bool hasMesh(FluidModifierData *fmd, int framenr);
-  bool hasParticles(FluidModifierData *mmd, int framenr);
+  bool hasParticles(FluidModifierData *fmd, int framenr);
-  bool hasGuiding(FluidModifierData *mmd, int framenr, bool sourceDomain);
+  bool hasGuiding(FluidModifierData *fmd, int framenr, bool sourceDomain);
 
   inline size_t getTotalCells()
   {
@@ -739,7 +739,7 @@ struct MANTA {
   float getTimestep();
   void adaptTimestep();
 
-  bool needsRealloc(FluidModifierData *mmd);
+  bool needsRealloc(FluidModifierData *fmd);
 
  private:
   // simulation constants
@@ -874,20 +874,20 @@ struct MANTA {
   vector<pVel> *mSndParticleVelocity;
   vector<float> *mSndParticleLife;
 
-  void initializeRNAMap(struct FluidModifierData *mmd = NULL);
+  void initializeRNAMap(struct FluidModifierData *fmd = NULL);
-  void initDomain(struct FluidModifierData *mmd = NULL);
+  void initDomain(struct FluidModifierData *fmd = NULL);
-  void initNoise(struct FluidModifierData *mmd = NULL);
+  void initNoise(struct FluidModifierData *fmd = NULL);
-  void initMesh(struct FluidModifierData *mmd = NULL);
+  void initMesh(struct FluidModifierData *fmd = NULL);
-  void initSmoke(struct FluidModifierData *mmd = NULL);
+  void initSmoke(struct FluidModifierData *fmd = NULL);
-  void initSmokeNoise(struct FluidModifierData *mmd = NULL);
+  void initSmokeNoise(struct FluidModifierData *fmd = NULL);
   void initializeMantaflow();
   void terminateMantaflow();
   bool runPythonString(vector<string> commands);
   string getRealValue(const string &varName);
   string parseLine(const string &line);
-  string parseScript(const string &setup_string, FluidModifierData *mmd = NULL);
+  string parseScript(const string &setup_string, FluidModifierData *fmd = NULL);
-  string getDirectory(struct FluidModifierData *mmd, string subdirectory);
+  string getDirectory(struct FluidModifierData *fmd, string subdirectory);
-  string getFile(struct FluidModifierData *mmd,
+  string getFile(struct FluidModifierData *fmd,
                  string subdirectory,
                  string fname,
                  string extension,

intern/mantaflow/intern/manta_fluid_API.cpp

@@ -27,9 +27,9 @@
 #include "manta_fluid_API.h"
 
 /* Fluid functions */
-MANTA *manta_init(int *res, struct FluidModifierData *mmd)
+MANTA *manta_init(int *res, struct FluidModifierData *fmd)
 {
-  return new MANTA(res, mmd);
+  return new MANTA(res, fmd);
 }
 void manta_free(MANTA *fluid)
 {
@@ -37,173 +37,173 @@ void manta_free(MANTA *fluid)
   fluid = nullptr;
 }
 
-void manta_ensure_obstacle(MANTA *fluid, struct FluidModifierData *mmd)
+void manta_ensure_obstacle(MANTA *fluid, struct FluidModifierData *fmd)
 {
   if (!fluid)
     return;
-  fluid->initObstacle(mmd);
+  fluid->initObstacle(fmd);
   fluid->updatePointers();
 }
-void manta_ensure_guiding(MANTA *fluid, struct FluidModifierData *mmd)
+void manta_ensure_guiding(MANTA *fluid, struct FluidModifierData *fmd)
 {
   if (!fluid)
     return;
-  fluid->initGuiding(mmd);
+  fluid->initGuiding(fmd);
   fluid->updatePointers();
 }
-void manta_ensure_invelocity(MANTA *fluid, struct FluidModifierData *mmd)
+void manta_ensure_invelocity(MANTA *fluid, struct FluidModifierData *fmd)
 {
   if (!fluid)
     return;
-  fluid->initInVelocity(mmd);
+  fluid->initInVelocity(fmd);
   fluid->updatePointers();
 }
-void manta_ensure_outflow(MANTA *fluid, struct FluidModifierData *mmd)
+void manta_ensure_outflow(MANTA *fluid, struct FluidModifierData *fmd)
 {
   if (!fluid)
     return;
-  fluid->initOutflow(mmd);
+  fluid->initOutflow(fmd);
   fluid->updatePointers();
 }
 
-int manta_write_config(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_write_config(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->writeConfiguration(mmd, framenr);
+  return fluid->writeConfiguration(fmd, framenr);
 }
 
-int manta_write_data(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_write_data(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->writeData(mmd, framenr);
+  return fluid->writeData(fmd, framenr);
 }
 
-int manta_write_noise(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_write_noise(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->writeNoise(mmd, framenr);
+  return fluid->writeNoise(fmd, framenr);
 }
 
-int manta_read_config(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_read_config(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->readConfiguration(mmd, framenr);
+  return fluid->readConfiguration(fmd, framenr);
 }
 
-int manta_read_data(MANTA *fluid, FluidModifierData *mmd, int framenr, bool resumable)
+int manta_read_data(MANTA *fluid, FluidModifierData *fmd, int framenr, bool resumable)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->readData(mmd, framenr, resumable);
+  return fluid->readData(fmd, framenr, resumable);
 }
 
-int manta_read_noise(MANTA *fluid, FluidModifierData *mmd, int framenr, bool resumable)
+int manta_read_noise(MANTA *fluid, FluidModifierData *fmd, int framenr, bool resumable)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->readNoise(mmd, framenr, resumable);
+  return fluid->readNoise(fmd, framenr, resumable);
 }
 
-int manta_read_mesh(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_read_mesh(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->readMesh(mmd, framenr);
+  return fluid->readMesh(fmd, framenr);
 }
 
-int manta_read_particles(MANTA *fluid, FluidModifierData *mmd, int framenr, bool resumable)
+int manta_read_particles(MANTA *fluid, FluidModifierData *fmd, int framenr, bool resumable)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->readParticles(mmd, framenr, resumable);
+  return fluid->readParticles(fmd, framenr, resumable);
 }
 
-int manta_read_guiding(MANTA *fluid, FluidModifierData *mmd, int framenr, bool sourceDomain)
+int manta_read_guiding(MANTA *fluid, FluidModifierData *fmd, int framenr, bool sourceDomain)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->readGuiding(mmd, framenr, sourceDomain);
+  return fluid->readGuiding(fmd, framenr, sourceDomain);
 }
 
-int manta_bake_data(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_bake_data(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->bakeData(mmd, framenr);
+  return fluid->bakeData(fmd, framenr);
 }
 
-int manta_bake_noise(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_bake_noise(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->bakeNoise(mmd, framenr);
+  return fluid->bakeNoise(fmd, framenr);
 }
 
-int manta_bake_mesh(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_bake_mesh(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->bakeMesh(mmd, framenr);
+  return fluid->bakeMesh(fmd, framenr);
 }
 
-int manta_bake_particles(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_bake_particles(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->bakeParticles(mmd, framenr);
+  return fluid->bakeParticles(fmd, framenr);
 }
 
-int manta_bake_guiding(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_bake_guiding(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->bakeGuiding(mmd, framenr);
+  return fluid->bakeGuiding(fmd, framenr);
 }
 
-int manta_has_data(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_has_data(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->hasData(mmd, framenr);
+  return fluid->hasData(fmd, framenr);
 }
 
-int manta_has_noise(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_has_noise(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->hasNoise(mmd, framenr);
+  return fluid->hasNoise(fmd, framenr);
 }
 
-int manta_has_mesh(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_has_mesh(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->hasMesh(mmd, framenr);
+  return fluid->hasMesh(fmd, framenr);
 }
 
-int manta_has_particles(MANTA *fluid, FluidModifierData *mmd, int framenr)
+int manta_has_particles(MANTA *fluid, FluidModifierData *fmd, int framenr)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->hasParticles(mmd, framenr);
+  return fluid->hasParticles(fmd, framenr);
 }
 
-int manta_has_guiding(MANTA *fluid, FluidModifierData *mmd, int framenr, bool domain)
+int manta_has_guiding(MANTA *fluid, FluidModifierData *fmd, int framenr, bool domain)
 {
-  if (!fluid || !mmd)
+  if (!fluid || !fmd)
     return 0;
-  return fluid->hasGuiding(mmd, framenr, domain);
+  return fluid->hasGuiding(fmd, framenr, domain);
 }
 
-void manta_update_variables(MANTA *fluid, FluidModifierData *mmd)
+void manta_update_variables(MANTA *fluid, FluidModifierData *fmd)
 {
   if (!fluid)
     return;
-  fluid->updateVariables(mmd);
+  fluid->updateVariables(fmd);
 }
 
 int manta_get_frame(MANTA *fluid)
@@ -227,11 +227,11 @@ void manta_adapt_timestep(MANTA *fluid)
   fluid->adaptTimestep();
 }
 
-bool manta_needs_realloc(MANTA *fluid, FluidModifierData *mmd)
+bool manta_needs_realloc(MANTA *fluid, FluidModifierData *fmd)
 {
   if (!fluid)
     return false;
-  return fluid->needsRealloc(mmd);
+  return fluid->needsRealloc(fmd);
 }
 
 /* Fluid accessors */
@@ -361,11 +361,11 @@ float *manta_get_phioutstatic_in(MANTA *fluid)
 }
 
 /* Smoke functions */
-void manta_smoke_export_script(MANTA *smoke, FluidModifierData *mmd)
+void manta_smoke_export_script(MANTA *smoke, FluidModifierData *fmd)
 {
-  if (!smoke || !mmd)
+  if (!smoke || !fmd)
     return;
-  smoke->exportSmokeScript(mmd);
+  smoke->exportSmokeScript(fmd);
 }
 
 void manta_smoke_export(MANTA *smoke,
@@ -527,31 +527,31 @@ void manta_smoke_turbulence_get_rgba_fixed_color(MANTA *smoke,
   get_rgba_fixed_color(color, smoke->getTotalCellsHigh(), data, sequential);
 }
 
-void manta_smoke_ensure_heat(MANTA *smoke, struct FluidModifierData *mmd)
+void manta_smoke_ensure_heat(MANTA *smoke, struct FluidModifierData *fmd)
 {
   if (smoke) {
-    smoke->initHeat(mmd);
+    smoke->initHeat(fmd);
     smoke->updatePointers();
   }
 }
 
-void manta_smoke_ensure_fire(MANTA *smoke, struct FluidModifierData *mmd)
+void manta_smoke_ensure_fire(MANTA *smoke, struct FluidModifierData *fmd)
 {
   if (smoke) {
-    smoke->initFire(mmd);
+    smoke->initFire(fmd);
     if (smoke->usingNoise()) {
-      smoke->initFireHigh(mmd);
+      smoke->initFireHigh(fmd);
     }
     smoke->updatePointers();
   }
 }
 
-void manta_smoke_ensure_colors(MANTA *smoke, struct FluidModifierData *mmd)
+void manta_smoke_ensure_colors(MANTA *smoke, struct FluidModifierData *fmd)
 {
   if (smoke) {
-    smoke->initColors(mmd);
+    smoke->initColors(fmd);
     if (smoke->usingNoise()) {
-      smoke->initColorsHigh(mmd);
+      smoke->initColorsHigh(fmd);
     }
     smoke->updatePointers();
   }
@@ -700,17 +700,17 @@ int manta_smoke_turbulence_get_cells(MANTA *smoke)
 }
 
 /* Liquid functions */
-void manta_liquid_export_script(MANTA *liquid, FluidModifierData *mmd)
+void manta_liquid_export_script(MANTA *liquid, FluidModifierData *fmd)
 {
-  if (!liquid || !mmd)
+  if (!liquid || !fmd)
     return;
-  liquid->exportLiquidScript(mmd);
+  liquid->exportLiquidScript(fmd);
 }
 
-void manta_liquid_ensure_sndparts(MANTA *liquid, struct FluidModifierData *mmd)
+void manta_liquid_ensure_sndparts(MANTA *liquid, struct FluidModifierData *fmd)
 {
   if (liquid) {
-    liquid->initLiquidSndParts(mmd);
+    liquid->initLiquidSndParts(fmd);
     liquid->updatePointers();
   }
 }

intern/mantaflow/intern/strings/fluid_script.h

@@ -146,19 +146,19 @@ mantaMsg('1 Mantaflow cell is ' + str(ratioMetersToRes_s$ID$) + ' Blender length
 ratioResToBLength_s$ID$ = float(res_s$ID$) / float(domainSize_s$ID$) # [cells / blength] (blength: cm, m, or km, ... )\n\
 mantaMsg('1 Blender length unit is ' + str(ratioResToBLength_s$ID$) + ' Mantaflow cells long.')\n\
 \n\
-ratioBTimeToTimstep_s$ID$ = float(1) / float(frameLengthRaw_s$ID$) # the time within 1 blender time unit, see also fluid.c\n\
+ratioBTimeToTimestep_s$ID$ = float(1) / float(frameLengthRaw_s$ID$) # the time within 1 blender time unit, see also fluid.c\n\
-mantaMsg('1 Blender time unit is ' + str(ratioBTimeToTimstep_s$ID$) + ' Mantaflow time units long.')\n\
+mantaMsg('1 Blender time unit is ' + str(ratioBTimeToTimestep_s$ID$) + ' Mantaflow time units long.')\n\
 \n\
 ratioFrameToFramelength_s$ID$ = float(1) / float(frameLengthUnscaled_s$ID$ ) # the time within 1 frame\n\
 mantaMsg('frame / frameLength is ' + str(ratioFrameToFramelength_s$ID$) + ' Mantaflow time units long.')\n\
 \n\
-scaleAcceleration_s$ID$ = ratioResToBLength_s$ID$ * (ratioBTimeToTimstep_s$ID$**2)# [meters/btime^2] to [cells/timestep^2] (btime: sec, min, or h, ...)\n\
+scaleAcceleration_s$ID$ = ratioResToBLength_s$ID$ * (ratioBTimeToTimestep_s$ID$**2)# [meters/btime^2] to [cells/timestep^2] (btime: sec, min, or h, ...)\n\
 mantaMsg('scaleAcceleration is ' + str(scaleAcceleration_s$ID$))\n\
 \n\
 scaleSpeedFrames_s$ID$ = ratioResToBLength_s$ID$ * ratioFrameToFramelength_s$ID$ # [blength/frame] to [cells/frameLength]\n\
 mantaMsg('scaleSpeed is ' + str(scaleSpeedFrames_s$ID$))\n\
 \n\
-scaleSpeedTime_s$ID$ = ratioResToBLength_s$ID$ * ratioBTimeToTimstep_s$ID$ # [blength/btime] to [cells/frameLength]\n\
+scaleSpeedTime_s$ID$ = ratioResToBLength_s$ID$ * ratioBTimeToTimestep_s$ID$ # [blength/btime] to [cells/frameLength]\n\
 mantaMsg('scaleSpeedTime is ' + str(scaleSpeedTime_s$ID$))\n\
 \n\
 gravity_s$ID$ *= scaleAcceleration_s$ID$ # scale from world acceleration to cell based acceleration\n\
@@ -461,14 +461,14 @@ if 'fluid_data_dict_resume_s$ID$' in globals(): fluid_data_dict_resume_s$ID$.cle
 if 'fluid_guiding_dict_s$ID$' in globals(): fluid_guiding_dict_s$ID$.clear()\n\
 if 'fluid_vel_dict_s$ID$' in globals(): fluid_vel_dict_s$ID$.clear()\n\
 \n\
-# Delete all childs from objects (e.g. pdata for particles)\n\
+# Delete all children from objects (e.g. pdata for particles)\n\
-mantaMsg('Release solver childs childs')\n\
+mantaMsg('Release solver childrens children')\n\
 for var in list(globals()):\n\
     if var.endswith('_pp$ID$') or var.endswith('_mesh$ID$'):\n\
         del globals()[var]\n\
 \n\
-# Now delete childs from solver objects\n\
+# Now delete children from solver objects\n\
-mantaMsg('Release solver childs')\n\
+mantaMsg('Release solver children')\n\
 for var in list(globals()):\n\
     if var.endswith('_s$ID$') or var.endswith('_sn$ID$') or var.endswith('_sm$ID$') or var.endswith('_sp$ID$') or var.endswith('_sg$ID$'):\n\
         del globals()[var]\n\
@@ -502,10 +502,12 @@ gc.collect()\n";
 // BAKE
 //////////////////////////////////////////////////////////////////////
 
+/* This has to match the behavior of BLI_path_frame,
+ * for positive and negative frame numbers. */
 const std::string fluid_cache_helper =
     "\n\
 def fluid_cache_get_framenr_formatted_$ID$(framenr):\n\
-    return str(framenr).zfill(4) # framenr with leading zeroes\n";
+    return str(framenr).zfill(4) if framenr >= 0 else str(framenr).zfill(5)\n";
 
 const std::string fluid_bake_multiprocessing =
     "\n\

intern/sky/CMakeLists.txt

@@ -0,0 +1,35 @@
+# ***** BEGIN GPL LICENSE BLOCK *****
+#
+# This program is free software; you can redistribute it and/or
+# modify it under the terms of the GNU General Public License
+# as published by the Free Software Foundation; either version 2
+# of the License, or (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program; if not, write to the Free Software Foundation,
+# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+#
+# ***** END GPL LICENSE BLOCK *****
+
+set(INC
+  include
+)
+
+set(INC_SYS
+
+)
+
+set(SRC
+  source/sky_model.cpp
+  source/sky_nishita.cpp
+)
+
+set(LIB
+)
+
+blender_add_lib(bf_intern_sky "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")

intern/sky/include/sky_model.h

@@ -298,14 +298,14 @@ HINT #1:   if you want to model the sky of an earth-like planet that orbits
            previous paragraph.
 */
 
-#include "util/util_types.h"
+#ifndef __SKY_MODEL_H__
+#define __SKY_MODEL_H__
 
-CCL_NAMESPACE_BEGIN
+#ifdef __cplusplus
+extern "C" {
+#endif
 
-#ifndef _SKY_MODEL_H_
+typedef double SKY_ArHosekSkyModelConfiguration[9];
-#  define _SKY_MODEL_H_
-
-typedef double ArHosekSkyModelConfiguration[9];
 
 //   Spectral version of the model
 
@@ -335,8 +335,8 @@ typedef double ArHosekSkyModelConfiguration[9];
 
 ---------------------------------------------------------------------------- */
 
-typedef struct ArHosekSkyModelState {
+typedef struct SKY_ArHosekSkyModelState {
-  ArHosekSkyModelConfiguration configs[11];
+  SKY_ArHosekSkyModelConfiguration configs[11];
   double radiances[11];
   double turbidity;
   double solar_radius;
@@ -344,7 +344,7 @@ typedef struct ArHosekSkyModelState {
   double emission_correction_factor_sun[11];
   double albedo;
   double elevation;
-} ArHosekSkyModelState;
+} SKY_ArHosekSkyModelState;
 
 /* ----------------------------------------------------------------------------
 
@@ -355,9 +355,9 @@ typedef struct ArHosekSkyModelState {
 
 ---------------------------------------------------------------------------- */
 
-ArHosekSkyModelState *arhosekskymodelstate_alloc_init(const double solar_elevation,
+SKY_ArHosekSkyModelState *SKY_arhosekskymodelstate_alloc_init(const double solar_elevation,
-                                                      const double atmospheric_turbidity,
+                                                              const double atmospheric_turbidity,
-                                                      const double ground_albedo);
+                                                              const double ground_albedo);
 
 /* ----------------------------------------------------------------------------
 
@@ -388,66 +388,68 @@ ArHosekSkyModelState *arhosekskymodelstate_alloc_init(const double solar_elevati
 
 ---------------------------------------------------------------------------- */
 
-ArHosekSkyModelState *arhosekskymodelstate_alienworld_alloc_init(
+SKY_ArHosekSkyModelState *SKY_arhosekskymodelstate_alienworld_alloc_init(
     const double solar_elevation,
     const double solar_intensity,
     const double solar_surface_temperature_kelvin,
     const double atmospheric_turbidity,
     const double ground_albedo);
 
-void arhosekskymodelstate_free(ArHosekSkyModelState *state);
+void SKY_arhosekskymodelstate_free(SKY_ArHosekSkyModelState *state);
 
-double arhosekskymodel_radiance(ArHosekSkyModelState *state,
+double SKY_arhosekskymodel_radiance(SKY_ArHosekSkyModelState *state,
-                                double theta,
+                                    double theta,
-                                double gamma,
+                                    double gamma,
-                                double wavelength);
+                                    double wavelength);
 
 // CIE XYZ and RGB versions
 
-ArHosekSkyModelState *arhosek_xyz_skymodelstate_alloc_init(const double turbidity,
+SKY_ArHosekSkyModelState *SKY_arhosek_xyz_skymodelstate_alloc_init(const double turbidity,
-                                                           const double albedo,
+                                                                   const double albedo,
-                                                           const double elevation);
+                                                                   const double elevation);
 
-ArHosekSkyModelState *arhosek_rgb_skymodelstate_alloc_init(const double turbidity,
+SKY_ArHosekSkyModelState *SKY_arhosek_rgb_skymodelstate_alloc_init(const double turbidity,
-                                                           const double albedo,
+                                                                   const double albedo,
-                                                           const double elevation);
+                                                                   const double elevation);
 
-double arhosek_tristim_skymodel_radiance(ArHosekSkyModelState *state,
+double SKY_arhosek_tristim_skymodel_radiance(SKY_ArHosekSkyModelState *state,
-                                         double theta,
+                                             double theta,
-                                         double gamma,
+                                             double gamma,
-                                         int channel);
+                                             int channel);
 
 //   Delivers the complete function: sky + sun, including limb darkening.
 //   Please read the above description before using this - there are several
 //   caveats!
 
-double arhosekskymodel_solar_radiance(ArHosekSkyModelState *state,
+double SKY_arhosekskymodel_solar_radiance(SKY_ArHosekSkyModelState *state,
-                                      double theta,
+                                          double theta,
-                                      double gamma,
+                                          double gamma,
-                                      double wavelength);
+                                          double wavelength);
-
-#endif  // _SKY_MODEL_H_
 
 /* Nishita improved sky model */
 
-void nishita_skymodel_precompute_texture(float *pixels,
+void SKY_nishita_skymodel_precompute_texture(float *pixels,
-                                         int stride,
+                                             int stride,
-                                         int start_y,
+                                             int start_y,
-                                         int end_y,
+                                             int end_y,
-                                         int width,
+                                             int width,
-                                         int height,
+                                             int height,
-                                         float sun_elevation,
+                                             float sun_elevation,
+                                             float altitude,
+                                             float air_density,
+                                             float dust_density,
+                                             float ozone_density);
+
+void SKY_nishita_skymodel_precompute_sun(float sun_elevation,
+                                         float angular_diameter,
                                          float altitude,
                                          float air_density,
                                          float dust_density,
-                                         float ozone_density);
+                                         float *r_pixel_bottom,
+                                         float *r_pixel_top);
 
-void nishita_skymodel_precompute_sun(float sun_elevation,
+#ifdef __cplusplus
-                                     float angular_diameter,
+}
-                                     float altitude,
+#endif
-                                     float air_density,
-                                     float dust_density,
-                                     float *pixel_bottom,
-                                     float *pixel_top);
 
-CCL_NAMESPACE_END
+#endif  // __SKY_MODEL_H__

intern/sky/source/sky_float3.h

@@ -0,0 +1,157 @@
+/*
+ * 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.
+ */
+
+#ifndef __SKY_FLOAT3_H__
+#define __SKY_FLOAT3_H__
+
+// minimal float3 + util_math.h implementation for nishita sky model
+
+#include <math.h>
+
+#ifndef M_PI_F
+#  define M_PI_F (3.1415926535897932f) /* pi */
+#endif
+#ifndef M_PI_2_F
+#  define M_PI_2_F (1.5707963267948966f) /* pi/2 */
+#endif
+#ifndef M_2PI_F
+#  define M_2PI_F (6.2831853071795864f) /* 2*pi */
+#endif
+
+struct float3 {
+  float x, y, z;
+
+  float3() = default;
+
+  float3(const float *ptr) : x{ptr[0]}, y{ptr[1]}, z{ptr[2]}
+  {
+  }
+
+  float3(const float (*ptr)[3]) : float3((const float *)ptr)
+  {
+  }
+
+  explicit float3(float value) : x(value), y(value), z(value)
+  {
+  }
+
+  explicit float3(int value) : x(value), y(value), z(value)
+  {
+  }
+
+  float3(float x, float y, float z) : x{x}, y{y}, z{z}
+  {
+  }
+
+  operator const float *() const
+  {
+    return &x;
+  }
+
+  operator float *()
+  {
+    return &x;
+  }
+
+  friend float3 operator*(const float3 &a, float b)
+  {
+    return {a.x * b, a.y * b, a.z * b};
+  }
+
+  friend float3 operator*(float b, const float3 &a)
+  {
+    return {a.x * b, a.y * b, a.z * b};
+  }
+
+  friend float3 operator-(const float3 &a, const float3 &b)
+  {
+    return {a.x - b.x, a.y - b.y, a.z - b.z};
+  }
+
+  friend float3 operator-(const float3 &a)
+  {
+    return {-a.x, -a.y, -a.z};
+  }
+
+  float length_squared() const
+  {
+    return x * x + y * y + z * z;
+  }
+
+  float length() const
+  {
+    return sqrt(length_squared());
+  }
+
+  static float distance(const float3 &a, const float3 &b)
+  {
+    return (a - b).length();
+  }
+
+  friend float3 operator+(const float3 &a, const float3 &b)
+  {
+    return {a.x + b.x, a.y + b.y, a.z + b.z};
+  }
+
+  void operator+=(const float3 &b)
+  {
+    this->x += b.x;
+    this->y += b.y;
+    this->z += b.z;
+  }
+
+  friend float3 operator*(const float3 &a, const float3 &b)
+  {
+    return {a.x * b.x, a.y * b.y, a.z * b.z};
+  }
+};
+
+inline float sqr(float a)
+{
+  return a * a;
+}
+
+inline float3 make_float3(float x, float y, float z)
+{
+  return float3(x, y, z);
+}
+
+inline float dot(const float3 &a, const float3 &b)
+{
+  return a.x * b.x + a.y * b.y + a.z * b.z;
+}
+
+inline float distance(const float3 &a, const float3 &b)
+{
+  return float3::distance(a, b);
+}
+
+inline float len_squared(float3 f)
+{
+  return f.length_squared();
+}
+
+inline float len(float3 f)
+{
+  return f.length();
+}
+
+inline float reduce_add(float3 f)
+{
+  return f.x + f.y + f.z;
+}
+
+#endif /* __SKY_FLOAT3_H__ */