Workbench: Refactor: Antialiasing

This new object makes it easier to resolve shader code dependency.

.arcconfig

@@ -1,7 +1,6 @@
 {
 	"project_id" : "Blender",
 	"conduit_uri" : "https://developer.blender.org/",
-	"phabricator.uri" : "https://developer.blender.org/",
 	"git.default-relative-commit" : "origin/master",
 	"arc.land.update.default" : "rebase",
 	"arc.land.onto.default" : "master"

.clang-format

@@ -132,7 +132,9 @@ PenaltyBreakAssignment: 100
 
 AllowShortFunctionsOnASingleLine: None
 
-SortIncludes: true
+# Disable for now since it complicates initial migration tests,
+# TODO: look into enabling this in the future.
+SortIncludes: false
 
 # Don't right align escaped newlines to the right because we have a wide default
 AlignEscapedNewlines: DontAlign

CMakeLists.txt

@@ -180,14 +180,6 @@ option(WITH_BULLET        "Enable Bullet (Physics Engine)" ON)
 option(WITH_SYSTEM_BULLET "Use the systems bullet library (currently unsupported due to missing features in upstream!)" )
 mark_as_advanced(WITH_SYSTEM_BULLET)
 option(WITH_OPENCOLORIO   "Enable OpenColorIO color management" ON)
-if(APPLE)
-  # There's no OpenXR runtime in sight for macOS, neither is code well
-  # tested there -> disable it by default.
-  option(WITH_XR_OPENXR   "Enable VR features through the OpenXR specification" OFF)
-  mark_as_advanced(WITH_XR_OPENXR)
-else()
-  option(WITH_XR_OPENXR   "Enable VR features through the OpenXR specification" ON)
-endif()
 
 # Compositor
 option(WITH_COMPOSITOR         "Enable the tile based nodal compositor" ON)
@@ -319,10 +311,6 @@ mark_as_advanced(WITH_SYSTEM_GLOG)
 # Freestyle
 option(WITH_FREESTYLE     "Enable Freestyle (advanced edges rendering)" ON)
 
-# New object types
-option(WITH_NEW_OBJECT_TYPES "Enable new hair and pointcloud objects (use for development only, don't save in files)" OFF)
-mark_as_advanced(WITH_NEW_OBJECT_TYPES)
-
 # Misc
 if(WIN32)
   option(WITH_INPUT_IME "Enable Input Method Editor (IME) for complex Asian character input" ON)
@@ -645,9 +633,6 @@ set_and_warn_dependency(WITH_TBB WITH_OPENIMAGEDENOISE  OFF)
 set_and_warn_dependency(WITH_TBB WITH_OPENVDB           OFF)
 set_and_warn_dependency(WITH_TBB WITH_MOD_FLUID         OFF)
 
-# OpenVDB uses 'half' type from OpenEXR & fails to link without OpenEXR enabled.
-set_and_warn_dependency(WITH_IMAGE_OPENEXR WITH_OPENVDB OFF)
-
 # auto enable openimageio for cycles
 if(WITH_CYCLES)
   set(WITH_OPENIMAGEIO ON)
@@ -684,7 +669,6 @@ if(WITH_GHOST_SDL OR WITH_HEADLESS)
   set(WITH_X11_ALPHA     OFF)
   set(WITH_GHOST_XDND    OFF)
   set(WITH_INPUT_IME     OFF)
-  set(WITH_XR_OPENXR     OFF)
 endif()
 
 if(WITH_CPU_SSE)
@@ -1711,7 +1695,6 @@ if(FIRST_RUN)
   info_cfg_option(WITH_CYCLES)
   info_cfg_option(WITH_FREESTYLE)
   info_cfg_option(WITH_OPENCOLORIO)
-  info_cfg_option(WITH_XR_OPENXR)
   info_cfg_option(WITH_OPENIMAGEDENOISE)
   info_cfg_option(WITH_OPENVDB)
   info_cfg_option(WITH_ALEMBIC)

build_files/build_environment/CMakeLists.txt

@@ -99,7 +99,6 @@ else()
 endif()
 include(cmake/openimagedenoise.cmake)
 include(cmake/embree.cmake)
-include(cmake/xr_openxr.cmake)
 
 if(WITH_WEBP)
   include(cmake/webp.cmake)

build_files/build_environment/cmake/harvest.cmake

@@ -161,8 +161,6 @@ harvest(opensubdiv/include opensubdiv/include "*.h")
 harvest(opensubdiv/lib opensubdiv/lib "*.a")
 harvest(openvdb/include/openvdb openvdb/include/openvdb "*.h")
 harvest(openvdb/lib openvdb/lib "*.a")
-harvest(xr_openxr_sdk/include/openxr xr_openxr_sdk/include/openxr "*.h")
-harvest(xr_openxr_sdk/lib xr_openxr_sdk/lib "*.a")
 harvest(osl/bin osl/bin "oslc")
 harvest(osl/include osl/include "*.h")
 harvest(osl/lib osl/lib "*.a")

build_files/build_environment/cmake/versions.cmake

@@ -318,7 +318,3 @@ set(LIBGLU_HASH 151aef599b8259efe9acd599c96ea2a3)
 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(XR_OPENXR_SDK_VERSION 1.0.6)
-set(XR_OPENXR_SDK_URI https://github.com/KhronosGroup/OpenXR-SDK/archive/release-${XR_OPENXR_SDK_VERSION}.tar.gz)
-set(XR_OPENXR_SDK_HASH 21daea7c3bfec365298d779a0e19caa1)

build_files/build_environment/cmake/xr_openxr.cmake

@@ -1,59 +0,0 @@
-# ***** BEGIN GPL LICENSE BLOCK *****
-#
-# This program is free software; you can redistribute it and/or
-# modify it under the terms of the GNU General Public License
-# as published by the Free Software Foundation; either version 2
-# of the License, or (at your option) any later version.
-#
-# This program is distributed in the hope that it will be useful,
-# but WITHOUT ANY WARRANTY; without even the implied warranty of
-# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-# GNU General Public License for more details.
-#
-# You should have received a copy of the GNU General Public License
-# along with this program; if not, write to the Free Software Foundation,
-# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-#
-# ***** END GPL LICENSE BLOCK *****
-
-
-# Keep flags in sync with install_deps.sh ones in compile_XR_OpenXR_SDK()
-set(XR_OPENXR_SDK_EXTRA_ARGS
-  -DBUILD_FORCE_GENERATION=OFF
-  -DBUILD_LOADER=ON
-  -DDYNAMIC_LOADER=OFF
-)
-
-if(UNIX AND NOT APPLE)
-  list(APPEND XR_OPENXR_SDK_EXTRA_ARGS
-    -DBUILD_WITH_WAYLAND_HEADERS=OFF
-    -DBUILD_WITH_XCB_HEADERS=OFF
-    -DBUILD_WITH_XLIB_HEADERS=ON
-    -DCMAKE_CXX_FLAGS=-DDISABLE_STD_FILESYSTEM=1
-  )
-endif()
-
-ExternalProject_Add(external_xr_openxr_sdk
-  URL ${XR_OPENXR_SDK_URI}
-  DOWNLOAD_DIR ${DOWNLOAD_DIR}
-  URL_HASH MD5=${XR_OPENXR_SDK_HASH}
-  PREFIX ${BUILD_DIR}/xr_openxr_sdk
-  CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${LIBDIR}/xr_openxr_sdk ${DEFAULT_CMAKE_FLAGS} ${XR_OPENXR_SDK_EXTRA_ARGS}
-  INSTALL_DIR ${LIBDIR}/xr_openxr_sdk
-)
-
-if(WIN32)
-  if(BUILD_MODE STREQUAL Release)
-    ExternalProject_Add_Step(external_xr_openxr_sdk after_install
-      COMMAND ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/xr_openxr_sdk/include/openxr ${HARVEST_TARGET}/xr_openxr_sdk/include/openxr
-      COMMAND ${CMAKE_COMMAND} -E copy_directory ${LIBDIR}/xr_openxr_sdk/lib ${HARVEST_TARGET}/xr_openxr_sdk/lib
-      DEPENDEES install
-    )
-  endif()
-  if(BUILD_MODE STREQUAL Debug)
-    ExternalProject_Add_Step(external_xr_openxr_sdk after_install
-      COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/xr_openxr_sdk/lib/openxr_loader.lib ${HARVEST_TARGET}/xr_openxr_sdk/lib/openxr_loader_d.lib
-      DEPENDEES install
-    )
-  endif()
-endif()

build_files/build_environment/install_deps.sh

@@ -52,19 +52,16 @@ getopt \
 -o s:i:t:h \
 --long source:,install:,tmp:,info:,threads:,help,show-deps,no-sudo,no-build,no-confirm,\
 with-all,with-opencollada,with-jack,with-embree,with-oidn,\
-ver-ocio:,ver-oiio:,ver-llvm:,ver-osl:,ver-osd:,ver-openvdb:,ver-xr-openxr:,\
+ver-ocio:,ver-oiio:,ver-llvm:,ver-osl:,ver-osd:,ver-openvdb:,\
 force-all,force-python,force-numpy,force-boost,\
 force-ocio,force-openexr,force-oiio,force-llvm,force-osl,force-osd,force-openvdb,\
 force-ffmpeg,force-opencollada,force-alembic,force-embree,force-oidn,force-usd,\
-force-xr-openxr,\
 build-all,build-python,build-numpy,build-boost,\
 build-ocio,build-openexr,build-oiio,build-llvm,build-osl,build-osd,build-openvdb,\
 build-ffmpeg,build-opencollada,build-alembic,build-embree,build-oidn,build-usd,\
-build-xr-openxr,\
 skip-python,skip-numpy,skip-boost,\
 skip-ocio,skip-openexr,skip-oiio,skip-llvm,skip-osl,skip-osd,skip-openvdb,\
-skip-ffmpeg,skip-opencollada,skip-alembic,skip-embree,skip-oidn,skip-usd, \
-skip-xr-openxr\
+skip-ffmpeg,skip-opencollada,skip-alembic,skip-embree,skip-oidn,skip-usd \
 -- "$@" \
 )
 
@@ -172,9 +169,6 @@ ARGUMENTS_INFO="\"COMMAND LINE ARGUMENTS:
     --ver-openvdb=<ver>
         Force version of OpenVDB library.
 
-    --ver-xr-openxr=<ver>
-        Force version of OpenXR-SDK.
-
     Note about the --ver-foo options:
         It may not always work as expected (some libs are actually checked out from a git rev...), yet it might help
         to fix some build issues (like LLVM mismatch with the version used by your graphic system).
@@ -230,9 +224,6 @@ ARGUMENTS_INFO="\"COMMAND LINE ARGUMENTS:
     --build-usd
         Force the build of Universal Scene Description.
 
-    --build-xr-openxr
-        Force the build of OpenXR-SDK.
-
     Note about the --build-foo options:
         * They force the script to prefer building dependencies rather than using available packages.
           This may make things simpler and allow working around some distribution bugs, but on the other hand it will
@@ -294,9 +285,6 @@ ARGUMENTS_INFO="\"COMMAND LINE ARGUMENTS:
     --force-usd
         Force the rebuild of Universal Scene Description.
 
-    --force-xr-openxr
-        Force the rebuild of OpenXR-SDK.
-
     Note about the --force-foo options:
         * They obviously only have an effect if those libraries are built by this script
           (i.e. if there is no available and satisfactory package)!
@@ -349,10 +337,7 @@ ARGUMENTS_INFO="\"COMMAND LINE ARGUMENTS:
         Unconditionally skip FFMpeg installation/building.
 
     --skip-usd
-        Unconditionally skip Universal Scene Description installation/building.
-
-    --skip-xr-openxr
-        Unconditionally skip OpenXR-SDK installation/building.\""
+        Unconditionally skip Universal Scene Description installation/building.\""
 
 # ----------------------------------------------------------------------------
 # Main Vars
@@ -469,11 +454,6 @@ FFMPEG_FORCE_REBUILD=false
 FFMPEG_SKIP=false
 _ffmpeg_list_sep=";"
 
-XR_OPENXR_VERSION="1.0.6"
-XR_OPENXR_FORCE_BUILD=false
-XR_OPENXR_FORCE_REBUILD=false
-XR_OPENXR_SKIP=false
-
 # FFMPEG optional libs.
 VORBIS_USE=false
 VORBIS_DEV=""
@@ -644,11 +624,6 @@ while true; do
       OPENVDB_VERSION_MIN=$OPENVDB_VERSION
       shift; shift; continue
     ;;
-    --ver-xr-openxr)
-      XR_OPENXR_VERSION="$2"
-      XR_OPENXR_VERSION_MIN=$XR_OPENXR_VERSION
-      shift; shift; continue
-    ;;
     --build-all)
       PYTHON_FORCE_BUILD=true
       NUMPY_FORCE_BUILD=true
@@ -666,7 +641,6 @@ while true; do
       FFMPEG_FORCE_BUILD=true
       ALEMBIC_FORCE_BUILD=true
       USD_FORCE_BUILD=true
-      XR_OPENXR_FORCE_BUILD=true
       shift; continue
     ;;
     --build-python)
@@ -721,9 +695,6 @@ while true; do
     --build-usd)
       USD_FORCE_BUILD=true; shift; continue
     ;;
-    --build-xr-openxr)
-      XR_OPENXR_FORCE_BUILD=true; shift; continue
-    ;;
     --force-all)
       PYTHON_FORCE_REBUILD=true
       NUMPY_FORCE_REBUILD=true
@@ -741,7 +712,6 @@ while true; do
       FFMPEG_FORCE_REBUILD=true
       ALEMBIC_FORCE_REBUILD=true
       USD_FORCE_REBUILD=true
-      XR_OPENXR_FORCE_REBUILD=true
       shift; continue
     ;;
     --force-python)
@@ -794,9 +764,6 @@ while true; do
     --force-usd)
       USD_FORCE_REBUILD=true; shift; continue
     ;;
-    --force-xr-openxr)
-      XR_OPENXR_FORCE_REBUILD=true; shift; continue
-    ;;
     --skip-python)
       PYTHON_SKIP=true; shift; continue
     ;;
@@ -845,9 +812,6 @@ while true; do
     --skip-usd)
       USD_SKIP=true; shift; continue
     ;;
-    --skip-xr-openxr)
-      XR_OPENXR_SKIP=true; shift; continue
-    ;;
     --)
       # no more arguments to parse
       break
@@ -976,12 +940,6 @@ OIDN_SOURCE=( "https://github.com/OpenImageDenoise/oidn/releases/download/v${OID
 
 FFMPEG_SOURCE=( "http://ffmpeg.org/releases/ffmpeg-$FFMPEG_VERSION.tar.bz2" )
 
-XR_OPENXR_USE_REPO=false
-XR_OPENXR_SOURCE=("https://github.com/KhronosGroup/OpenXR-SDK/archive/release-${XR_OPENXR_VERSION}.tar.gz")
-#~ XR_OPENXR_SOURCE_REPO=("https://github.com/KhronosGroup/OpenXR-SDK-Source.git")
-#~ XR_OPENXR_REPO_UID="5292e57fda47561e672fba0a4b6e545c0f25dd8d"
-#~ XR_OPENXR_REPO_BRANCH="master"
-
 # C++11 is required now
 CXXFLAGS_BACK=$CXXFLAGS
 CXXFLAGS="$CXXFLAGS -std=c++11"
@@ -1024,8 +982,7 @@ You may also want to build them yourself (optional ones are [between brackets]):
     * [Embree $EMBREE_VERSION] (from $EMBREE_SOURCE).
     * [OpenImageDenoise $OIDN_VERSION] (from $OIDN_SOURCE).
     * [Alembic $ALEMBIC_VERSION] (from $ALEMBIC_SOURCE).
-    * [Universal Scene Description $USD_VERSION] (from $USD_SOURCE).
-    * [OpenXR-SDK $XR_OPENXR_VERSION] (from $XR_OPENXR_SOURCE).\""
+    * [Universal Scene Description $USD_VERSION] (from $USD_SOURCE).\""
 
 if [ "$DO_SHOW_DEPS" = true ]; then
   PRINT ""
@@ -3101,116 +3058,6 @@ compile_FFmpeg() {
   fi
 }
 
-# ----------------------------------------------------------------------------
-# Build OpenXR SDK
-
-_init_xr_openxr_sdk() {
-  _src=$SRC/XR-OpenXR-SDK-$XR_OPENXR_VERSION
-  _git=true
-  _inst=$INST/xr-openxr-sdk-$XR_OPENXR_VERSION
-  _inst_shortcut=$INST/xr-openxr-sdk
-}
-
-_update_deps_xr_openxr_sdk() {
-  :
-}
-
-clean_XR_OpenXR_SDK() {
-  _init_xr_openxr_sdk
-  _clean
-  _update_deps_xr_openxr_sdk
-}
-
-compile_XR_OpenXR_SDK() {
-  if [ "$NO_BUILD" = true ]; then
-    WARNING "--no-build enabled, OpenXR will not be compiled!"
-    return
-  fi
-
-  # To be changed each time we make edits that would modify the compiled result!
-  xr_openxr_magic=1
-  _init_xr_openxr_sdk
-
-  # Clean install if needed!
-  magic_compile_check xr-openxr-$OPENXR_VERSION $xr_openxr_magic
-  if [ $? -eq 1 -o "$XR_OPENXR_FORCE_REBUILD" = true ]; then
-    clean_XR_OpenXR_SDK
-  fi
-
-  if [ ! -d $_inst ]; then
-    INFO "Building XR-OpenXR-SDK-$XR_OPENXR_VERSION"
-    _is_building=true
-
-    # Rebuild dependencies as well!
-    _update_deps_xr_openxr_sdk
-
-    prepare_opt
-
-    if [ ! -d $_src ]; then
-      mkdir -p $SRC
-
-      if [ "$XR_OPENXR_USE_REPO" = true ]; then
-        git clone $XR_OPENXR_SOURCE_REPO $_src
-      else
-        download XR_OPENXR_SOURCE[@] "$_src.tar.gz"
-        INFO "Unpacking XR-OpenXR-SDK-$XR_OPENXR_VERSION"
-        tar -C $SRC --transform "s,(.*/?)OpenXR-SDK-[^/]*(.*),\1XR-OpenXR-SDK-$XR_OPENXR_VERSION\2,x" \
-            -xf $_src.tar.gz
-      fi
-    fi
-
-    cd $_src
-
-    if [ "$XR_OPENXR_USE_REPO" = true ]; then
-      git pull origin $XR_OPENXR_REPO_BRANCH
-
-      # Stick to same rev as windows' libs...
-      git checkout $XR_OPENXR_REPO_UID
-      git reset --hard
-    fi
-
-    # Always refresh the whole build!
-    if [ -d build ]; then
-      rm -rf build
-    fi
-    mkdir build
-    cd build
-
-    # Keep flags in sync with XR_OPENXR_SDK_EXTRA_ARGS in xr_openxr.cmake!
-    cmake_d="-D CMAKE_BUILD_TYPE=Release"
-    cmake_d="$cmake_d -D CMAKE_INSTALL_PREFIX=$_inst"
-    cmake_d="$cmake_d -D BUILD_FORCE_GENERATION=OFF"
-    cmake_d="$cmake_d -D BUILD_LOADER=ON"
-    cmake_d="$cmake_d -D DYNAMIC_LOADER=OFF"
-    cmake_d="$cmake_d -D BUILD_WITH_WAYLAND_HEADERS=OFF"
-    cmake_d="$cmake_d -D BUILD_WITH_XCB_HEADERS=OFF"
-    cmake_d="$cmake_d -D BUILD_WITH_XLIB_HEADERS=ON"
-
-    cmake $cmake_d "-DCMAKE_CXX_FLAGS=-DDISABLE_STD_FILESYSTEM=1" ..
-
-    make -j$THREADS && make install
-    make clean
-
-    if [ -d $_inst ]; then
-      _create_inst_shortcut
-    else
-      ERROR "XR-OpenXR-SDK-$XR_OPENXR_VERSION failed to compile, exiting"
-      exit 1
-    fi
-
-    magic_compile_set xr-openxr-$XR_OPENXR_VERSION $xr_openxr_magic
-
-    cd $CWD
-    INFO "Done compiling XR-OpenXR-SDK-$XR_OPENXR_VERSION!"
-    _is_building=false
-  else
-    INFO "Own XR-OpenXR-SDK-$XR_OPENXR_VERSION is up to date, nothing to do!"
-    INFO "If you want to force rebuild of this lib, use the --force-xr-openxr option."
-  fi
-
-  run_ldconfig "xr-openxr-sdk"
-}
-
 
 # ----------------------------------------------------------------------------
 # Install on DEB-like
@@ -3755,18 +3602,6 @@ install_DEB() {
       compile_FFmpeg
     fi
   fi
-
-  PRINT ""
-  if [ "$XR_OPENXR_SKIP" = true ]; then
-    WARNING "Skipping OpenXR-SDK installation, as requested..."
-  elif [ "$XR_OPENXR_FORCE_BUILD" = true ]; then
-    INFO "Forced OpenXR-SDK building, as requested..."
-    compile_XR_OpenXR_SDK
-  else
-    # No package currently!
-    PRINT ""
-    compile_XR_OpenXR_SDK
-  fi
 }
 
 
@@ -4373,17 +4208,6 @@ install_RPM() {
       compile_FFmpeg
     fi
   fi
-
-  PRINT ""
-  if [ "$XR_OPENXR_SKIP" = true ]; then
-    WARNING "Skipping OpenXR-SDK installation, as requested..."
-  elif [ "$XR_OPENXR_FORCE_BUILD" = true ]; then
-    INFO "Forced OpenXR-SDK building, as requested..."
-    compile_XR_OpenXR_SDK
-  else
-    # No package currently!
-    compile_XR_OpenXR_SDK
-  fi
 }
 
 
@@ -4885,17 +4709,6 @@ install_ARCH() {
       compile_FFmpeg
     fi
   fi
-
-  PRINT ""
-  if [ "$XR_OPENXR_SKIP" = true ]; then
-    WARNING "Skipping OpenXR-SDK installation, as requested..."
-  elif [ "$XR_OPENXR_FORCE_BUILD" = true ]; then
-    INFO "Forced OpenXR-SDK building, as requested..."
-    compile_XR_OpenXR_SDK
-  else
-    # No package currently!
-    compile_XR_OpenXR_SDK
-  fi
 }
 
 
@@ -5093,17 +4906,6 @@ install_OTHER() {
     INFO "Forced FFMpeg building, as requested..."
     compile_FFmpeg
   fi
-
-  PRINT ""
-  if [ "$XR_OPENXR_SKIP" = true ]; then
-    WARNING "Skipping OpenXR-SDK installation, as requested..."
-  elif [ "$XR_OPENXR_FORCE_BUILD" = true ]; then
-    INFO "Forced OpenXR-SDK building, as requested..."
-    compile_XR_OpenXR_SDK
-  else
-    # No package currently!
-    compile_XR_OpenXR_SDK
-  fi
 }
 
 # ----------------------------------------------------------------------------
@@ -5372,17 +5174,6 @@ print_info() {
     fi
   fi
 
-  if [ "$XR_OPENXR_SKIP" = false ]; then
-    _1="-D WITH_XR_OPENXR=ON"
-    PRINT "  $_1"
-    _buildargs="$_buildargs $_1"
-    if [ -d $INST/xr-openxr-sdk ]; then
-      _1="-D XR_OPENXR_SDK_ROOT_DIR=$INST/xr-openxr-sdk"
-      PRINT "  $_1"
-      _buildargs="$_buildargs $_1"
-    fi
-  fi
-
   PRINT ""
   PRINT "Or even simpler, just run (in your blender-source dir):"
   PRINT "  make -j$THREADS BUILD_CMAKE_ARGS=\"$_buildargs\""

build_files/buildbot/buildbot_utils.py

@@ -98,7 +98,7 @@ class VersionInfo:
             self.is_development_build = False
         else:
             # Development build
-            self.full_version = self.version + self.version_char + '-' + self.hash
+            self.full_version = self.version + '-' + self.hash
             self.is_development_build = True
 
     def _parse_header_file(self, filename, define):

build_files/cmake/Modules/FindXR-OpenXR-SDK.cmake

@@ -1,73 +0,0 @@
-# - Find OpenXR-SDK libraries
-# Find the native OpenXR-SDK includes and libraries
-#
-# Note that there is a distinction between the OpenXR standard and the SDK. The
-# latter provides utilities to use the standard but is not part of it. Most
-# importantly, it contains C headers and a loader library, which manages
-# dynamic linking to OpenXR runtimes like Monado, Windows Mixed Reality or
-# Oculus. See the repository for more details:
-# https://github.com/KhronosGroup/OpenXR-SDK
-#
-# This module defines
-#  XR_OPENXR_SDK_INCLUDE_DIRS, where to find OpenXR-SDK headers, Set when
-#                           XR_OPENXR_SDK_INCLUDE_DIR is found.
-#  XR_OPENXR_SDK_LIBRARIES, libraries to link against to use OpenXR.
-#  XR_OPENXR_SDK_ROOT_DIR, the base directory to search for OpenXR-SDK.
-#                        This can also be an environment variable.
-#  XR_OPENXR_SDK_FOUND, if false, do not try to use OpenXR-SDK.
-#
-# also defined, but not for general use are
-#  XR_OPENXR_SDK_LOADER_LIBRARY, where to find the OpenXR-SDK loader library.
-
-#=============================================================================
-# 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 XR_OPENXR_SDK_ROOT_DIR was defined in the environment, use it.
-IF(NOT XR_OPENXR_SDK_ROOT_DIR AND NOT $ENV{XR_OPENXR_SDK_ROOT_DIR} STREQUAL "")
-  SET(XR_OPENXR_SDK_ROOT_DIR $ENV{XR_OPENXR_SDK_ROOT_DIR})
-ENDIF()
-
-SET(_xr_openxr_sdk_SEARCH_DIRS
-  ${XR_OPENXR_SDK_ROOT_DIR}
-  /opt/lib/xr-openxr-sdk
-)
-
-FIND_PATH(XR_OPENXR_SDK_INCLUDE_DIR
-  NAMES
-    openxr/openxr.h
-  HINTS
-    ${_xr_openxr_sdk_SEARCH_DIRS}
-  PATH_SUFFIXES
-    include
-)
-
-FIND_LIBRARY(XR_OPENXR_SDK_LOADER_LIBRARY
-  NAMES
-    openxr_loader
-  HINTS
-    ${_xr_openxr_sdk_SEARCH_DIRS}
-  PATH_SUFFIXES
-    lib64 lib
-)
-
-# handle the QUIETLY and REQUIRED arguments and set XR_OPENXR_SDK_FOUND to TRUE if
-# all listed variables are TRUE
-INCLUDE(FindPackageHandleStandardArgs)
-FIND_PACKAGE_HANDLE_STANDARD_ARGS(XR_OPENXR_SDK DEFAULT_MSG
-    XR_OPENXR_SDK_LOADER_LIBRARY XR_OPENXR_SDK_INCLUDE_DIR)
-
-IF(XR_OPENXR_SDK_FOUND)
-  SET(XR_OPENXR_SDK_LIBRARIES ${XR_OPENXR_SDK_LOADER_LIBRARY})
-  SET(XR_OPENXR_SDK_INCLUDE_DIRS ${XR_OPENXR_SDK_INCLUDE_DIR})
-ENDIF(XR_OPENXR_SDK_FOUND)
-
-MARK_AS_ADVANCED(
-  XR_OPENXR_SDK_INCLUDE_DIR
-  XR_OPENXR_SDK_LOADER_LIBRARY
-)

build_files/cmake/config/blender_full.cmake

@@ -61,6 +61,3 @@ if(UNIX AND NOT APPLE)
   set(WITH_X11_XINPUT          ON  CACHE BOOL "" FORCE)
   set(WITH_X11_XF86VMODE       ON  CACHE BOOL "" FORCE)
 endif()
-if(NOT APPLE)
-  set(WITH_XR_OPENXR           ON  CACHE BOOL "" FORCE)
-endif()

build_files/cmake/config/blender_lite.cmake

@@ -6,6 +6,7 @@
 #
 
 set(WITH_INSTALL_PORTABLE    ON  CACHE BOOL "" FORCE)
+set(WITH_SYSTEM_GLEW         ON  CACHE BOOL "" FORCE)
 
 set(WITH_ALEMBIC             OFF CACHE BOOL "" FORCE)
 set(WITH_BOOST               OFF CACHE BOOL "" FORCE)
@@ -44,7 +45,6 @@ set(WITH_OPENAL              OFF CACHE BOOL "" FORCE)
 set(WITH_OPENCOLLADA         OFF CACHE BOOL "" FORCE)
 set(WITH_OPENCOLORIO         OFF CACHE BOOL "" FORCE)
 set(WITH_OPENIMAGEDENOISE    OFF CACHE BOOL "" FORCE)
-set(WITH_XR_OPENXR           OFF CACHE BOOL "" FORCE)
 set(WITH_OPENIMAGEIO         OFF CACHE BOOL "" FORCE)
 set(WITH_OPENMP              OFF CACHE BOOL "" FORCE)
 set(WITH_OPENSUBDIV          OFF CACHE BOOL "" FORCE)

build_files/cmake/config/blender_release.cmake

@@ -65,6 +65,3 @@ if(UNIX AND NOT APPLE)
   set(WITH_X11_XINPUT          ON  CACHE BOOL "" FORCE)
   set(WITH_X11_XF86VMODE       ON  CACHE BOOL "" FORCE)
 endif()
-if(NOT APPLE)
-  set(WITH_XR_OPENXR              ON  CACHE BOOL "" FORCE)
-endif()

build_files/cmake/macros.cmake

@@ -122,7 +122,7 @@ function(target_link_libraries_optimized
   )
 
   foreach(_LIB ${LIBS})
-    target_link_libraries(${TARGET} INTERFACE optimized "${_LIB}")
+    target_link_libraries(${TARGET} optimized "${_LIB}")
   endforeach()
 endfunction()
 
@@ -132,7 +132,7 @@ function(target_link_libraries_debug
   )
 
   foreach(_LIB ${LIBS})
-    target_link_libraries(${TARGET} INTERFACE debug "${_LIB}")
+    target_link_libraries(${TARGET} debug "${_LIB}")
   endforeach()
 endfunction()
 
@@ -170,7 +170,6 @@ function(blender_include_dirs_sys
 endfunction()
 
 function(blender_source_group
-  name
   sources
   )
 
@@ -206,13 +205,6 @@ function(blender_source_group
       source_group("${GROUP_ID}" FILES ${_SRC})
     endforeach()
   endif()
-
-  # if enabled, set the FOLDER property for visual studio projects
-  if(WINDOWS_USE_VISUAL_STUDIO_PROJECT_FOLDERS)
-    get_filename_component(FolderDir ${CMAKE_CURRENT_SOURCE_DIR} DIRECTORY)
-    string(REPLACE ${CMAKE_SOURCE_DIR} "" FolderDir ${FolderDir})
-    set_target_properties(${name} PROPERTIES FOLDER ${FolderDir})
-  endif()
 endfunction()
 
 
@@ -303,11 +295,11 @@ function(blender_add_lib__impl
         set(next_library_mode "${library_lower}")
       else()
         if("${next_library_mode}" STREQUAL "optimized")
-          target_link_libraries(${name} INTERFACE optimized ${library})
+          target_link_libraries(${name} optimized ${library})
         elseif("${next_library_mode}" STREQUAL "debug")
-          target_link_libraries(${name} INTERFACE debug ${library})
+          target_link_libraries(${name} debug ${library})
         else()
-          target_link_libraries(${name} INTERFACE ${library})
+          target_link_libraries(${name} ${library})
         endif()
         set(next_library_mode "")
       endif()
@@ -316,7 +308,14 @@ function(blender_add_lib__impl
 
   # works fine without having the includes
   # listed is helpful for IDE's (QtCreator/MSVC)
-  blender_source_group("${name}" "${sources}")
+  blender_source_group("${sources}")
+
+  # if enabled, set the FOLDER property for visual studio projects
+  if(WINDOWS_USE_VISUAL_STUDIO_PROJECT_FOLDERS)
+    get_filename_component(FolderDir ${CMAKE_CURRENT_SOURCE_DIR} DIRECTORY)
+    string(REPLACE ${CMAKE_SOURCE_DIR} "" FolderDir ${FolderDir})
+    set_target_properties(${name} PROPERTIES FOLDER ${FolderDir})
+  endif()
 
   list_assert_duplicates("${sources}")
   list_assert_duplicates("${includes}")
@@ -947,7 +946,7 @@ function(data_to_c_simple
   set_source_files_properties(${_file_to} PROPERTIES GENERATED TRUE)
 endfunction()
 
-# Function for converting pixmap directory to a '.png' and then a '.c' file.
+# macro for converting pixmap directory to a png and then a c file
 function(data_to_c_simple_icons
   path_from icon_prefix icon_names
   list_to_add
@@ -1160,12 +1159,12 @@ macro(blender_precompile_headers target cpp header)
 endmacro()
 
 macro(set_and_warn_dependency
-  _dependency _setting _val)
-  # when $_dependency is disabled, forces $_setting = $_val
-  if(NOT ${${_dependency}} AND ${${_setting}})
-    message(STATUS "'${_dependency}' is disabled: forcing 'set(${_setting} ${_val})'")
-    set(${_setting} ${_val})
-  endif()
+    _dependency _setting _val)
+    # when $_dependency is disabled, forces $_setting = $_val
+    if(NOT ${${_dependency}} AND ${${_setting}})
+      message(STATUS "'${_dependency}' is disabled: forcing 'set(${_setting} ${_val})'")
+      set(${_setting} ${_val})
+    endif()
 endmacro()
 
 macro(without_system_libs_begin)

build_files/cmake/platform/platform_apple.cmake

@@ -411,21 +411,11 @@ if(WITH_OPENMP)
 
     # Copy libomp.dylib to allow executables like datatoc and tests to work.
     execute_process(
-      COMMAND mkdir -p ${CMAKE_BINARY_DIR}/Resources/lib
-      COMMAND cp -p ${LIBDIR}/openmp/lib/libomp.dylib ${CMAKE_BINARY_DIR}/Resources/lib/libomp.dylib
-    )
+        COMMAND mkdir -p ${CMAKE_BINARY_DIR}/Resources/lib
+        COMMAND cp -p ${LIBDIR}/openmp/lib/libomp.dylib ${CMAKE_BINARY_DIR}/Resources/lib/libomp.dylib)
     execute_process(
-      COMMAND mkdir -p ${CMAKE_BINARY_DIR}/bin/Resources/lib
-      COMMAND cp -p ${LIBDIR}/openmp/lib/libomp.dylib ${CMAKE_BINARY_DIR}/bin/Resources/lib/libomp.dylib
-    )
-  endif()
-endif()
-
-if(WITH_XR_OPENXR)
-  find_package(OpenXR-SDK)
-  if(NOT OPENXR_SDK_FOUND)
-    message(WARNING "OpenXR-SDK was not found, disabling WITH_XR_OPENXR")
-    set(WITH_XR_OPENXR OFF)
+        COMMAND mkdir -p ${CMAKE_BINARY_DIR}/bin/Resources/lib
+        COMMAND cp -p ${LIBDIR}/openmp/lib/libomp.dylib ${CMAKE_BINARY_DIR}/bin/Resources/lib/libomp.dylib)
   endif()
 endif()
 

build_files/cmake/platform/platform_unix.cmake

@@ -428,14 +428,6 @@ if(WITH_TBB)
   find_package_wrapper(TBB)
 endif()
 
-if(WITH_XR_OPENXR)
-  find_package(XR-OpenXR-SDK)
-  if(NOT XR_OPENXR_SDK_FOUND)
-    message(WARNING "OpenXR-SDK not found, disabling WITH_XR_OPENXR")
-    set(WITH_XR_OPENXR OFF)
-  endif()
-endif()
-
 if(EXISTS ${LIBDIR})
   without_system_libs_end()
 endif()

build_files/cmake/platform/platform_win32.cmake

@@ -671,8 +671,8 @@ if(WITH_USD)
     set(USD_DEBUG_LIB ${LIBDIR}/usd/lib/libusd_m_d.lib)
     set(USD_LIBRARY_DIR ${LIBDIR}/usd/lib/usd)
     set(USD_LIBRARIES
-      debug ${USD_DEBUG_LIB}
-      optimized ${USD_RELEASE_LIB}
+        debug ${USD_DEBUG_LIB}
+        optimized ${USD_RELEASE_LIB}
     )
   endif()
 endif()
@@ -713,15 +713,3 @@ if(WINDOWS_PYTHON_DEBUG)
 </Project>")
   endif()
 endif()
-
-if(WITH_XR_OPENXR)
-  if(EXISTS ${LIBDIR}/xr_openxr_sdk)
-    set(XR_OPENXR_SDK ${LIBDIR}/xr_openxr_sdk)
-    set(XR_OPENXR_SDK_LIBPATH ${LIBDIR}/xr_openxr_sdk/lib)
-    set(XR_OPENXR_SDK_INCLUDE_DIR ${XR_OPENXR_SDK}/include)
-    set(XR_OPENXR_SDK_LIBRARIES optimized ${XR_OPENXR_SDK_LIBPATH}/openxr_loader.lib debug ${XR_OPENXR_SDK_LIBPATH}/openxr_loader_d.lib)
-  else()
-    message(WARNING "OpenXR-SDK was not found, disabling WITH_XR_OPENXR")
-    set(WITH_XR_OPENXR OFF)
-  endif()
-endif()

build_files/windows/show_help.cmd

@@ -23,7 +23,7 @@ echo - with_tests ^(enable building unit tests^)
 echo - nobuildinfo ^(disable buildinfo^)
 echo - debug ^(Build an unoptimized debuggable build^)
 echo - packagename [newname] ^(override default cpack package name^)
-echo - builddir [newdir] ^(override default build folder^)
+echo - buildir [newdir] ^(override default build folder^)
 echo - 2017 ^(build with visual studio 2017^)
 echo - 2017pre ^(build with visual studio 2017 pre-release^)
 echo - 2017b ^(build with visual studio 2017 Build Tools^)

doc/python_api/examples/bpy.msgbus.1.py

@@ -1,44 +0,0 @@
-"""
-The message bus system can be used to receive notifications when properties of
-Blender datablocks are changed via the data API.
-
-
-Limitations
------------
-
-The message bus system is triggered by updates via the RNA system. This means
-that the following updates will result in a notification on the message bus:
-
-- Changes via the Python API, for example ``some_object.location.x += 3``.
-- Changes via the sliders, fields, and buttons in the user interface.
-
-The following updates do **not** trigger message bus notifications:
-
-- Moving objects in the 3D Viewport.
-- Changes performed by the animation system.
-
-
-Example Use
------------
-
-Below is an example of subscription to changes in the active object's location.
-"""
-
-import bpy
-
-# Any Python object can act as the subscription's owner.
-owner = object()
-
-subscribe_to = bpy.context.object.location
-
-def msgbus_callback(*args):
-    # This will print:
-    # Something changed! (1, 2, 3)
-    print("Something changed!", args)
-
-bpy.msgbus.subscribe_rna(
-    key=subscribe_to,
-    owner=owner,
-    args=(1, 2, 3),
-    notify=msgbus_callback,
-)

doc/python_api/examples/bpy.msgbus.2.py

@@ -1,6 +0,0 @@
-"""
-Some properties are converted to Python objects when you retrieve them. This
-needs to be avoided in order to create the subscription, by using
-``datablock.path_resolve("property_name", False)``:
-"""
-subscribe_to = bpy.context.object.path_resolve("name", False)

doc/python_api/examples/bpy.msgbus.3.py

@@ -1,5 +0,0 @@
-"""
-It is also possible to create subscriptions on a property of all instances of a
-certain type:
-"""
-subscribe_to = (bpy.types.Object, "location")

doc/python_api/examples/bpy.types.AddonPreferences.1.py

@@ -6,7 +6,7 @@ bl_info = {
     "location": "SpaceBar Search -> Add-on Preferences Example",
     "description": "Example Add-on",
     "warning": "",
-    "doc_url": "",
+    "wiki_url": "",
     "tracker_url": "",
     "category": "Object",
 }

doc/python_api/rst/info_quickstart.rst

@@ -255,9 +255,9 @@ Examples:
    >>> bpy.ops.object.scale_apply()
    {'FINISHED'}
 
-.. tip::
+.. note::
 
-   The :ref:`Operator Cheat Sheet <blender_manual:bpy.ops.wm.operator_cheat_sheet>`.
+   The menu item: :menuselection:`Help --> Operator Cheat Sheet`
    gives a list of all operators and their default values in Python syntax, along with the generated docs.
    This is a good way to get an overview of all Blender's operators.
 

doc/python_api/sphinx_doc_gen.py

@@ -1756,7 +1756,6 @@ def write_rst_contents(basepath):
     app_modules = (
         "bpy.context",  # note: not actually a module
         "bpy.data",     # note: not actually a module
-        "bpy.msgbus",   # note: not actually a module
         "bpy.ops",
         "bpy.types",
 
@@ -1847,29 +1846,6 @@ def write_rst_ops_index(basepath):
         file.close()
 
 
-def write_rst_msgbus(basepath):
-    """
-    Write the rst files of bpy.msgbus module
-    """
-    if 'bpy.msgbus' in EXCLUDE_MODULES:
-        return
-
-    # Write the index.
-    filepath = os.path.join(basepath, "bpy.msgbus.rst")
-    file = open(filepath, "w", encoding="utf-8")
-    fw = file.write
-    fw(title_string("Message Bus (bpy.msgbus)", "="))
-    write_example_ref("", fw, "bpy.msgbus")
-    fw(".. toctree::\n")
-    fw("   :glob:\n\n")
-    fw("   bpy.msgbus.*\n\n")
-    file.close()
-
-    # Write the contents.
-    pymodule2sphinx(basepath, 'bpy.msgbus', bpy.msgbus, 'Message Bus')
-    EXAMPLE_SET_USED.add("bpy.msgbus")
-
-
 def write_rst_data(basepath):
     '''
     Write the rst file of bpy.data module
@@ -2024,7 +2000,6 @@ def rna2sphinx(basepath):
     write_rst_bpy(basepath)                 # bpy, disabled by default
     write_rst_types_index(basepath)         # bpy.types
     write_rst_ops_index(basepath)           # bpy.ops
-    write_rst_msgbus(basepath)              # bpy.msgbus
     pyrna2sphinx(basepath)                  # bpy.types.* and bpy.ops.*
     write_rst_data(basepath)                # bpy.data
     write_rst_importable_modules(basepath)

extern/bullet2/CMakeLists.txt

@@ -415,11 +415,4 @@ if(CMAKE_COMPILER_IS_GNUCXX)
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fpermissive")
 endif()
 
-if(MSVC)
-  # bullet is responsible for quite a few silly warnings
-  # suppress all of them. Not great, but they really needed
-  # to sort that out themselves.
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /W0")
-endif()
-
 blender_add_lib(extern_bullet "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")

extern/draco/CMakeLists.txt

@@ -24,21 +24,6 @@ set(CMAKE_CXX_STANDARD 14)
 add_subdirectory(dracoenc)
 
 # Build blender-draco-exporter module.
-set(SRC
-  src/draco-compressor.cpp
-  src/draco-compressor.h
-)
-
-set(INC
-  dracoenc/src
-)
-
-set(LIB
-  dracoenc
-)
-
-add_library(extern_draco SHARED "${SRC}")
-target_include_directories(extern_draco PUBLIC "${INC}")
-target_link_libraries(extern_draco PUBLIC "${LIB}")
-
-blender_source_group(extern_draco "${SRC}")
+add_library(extern_draco SHARED src/draco-compressor.cpp src/draco-compressor.h)
+target_include_directories(extern_draco PUBLIC dracoenc/src)
+target_link_libraries(extern_draco PUBLIC dracoenc)

extern/glog/CMakeLists.txt

@@ -77,11 +77,6 @@ if(WIN32)
   list(APPEND INC
     src/windows
   )
-  if(MSVC)
-    # Suppress warning about google::LogMessageFatal::~LogMessageFatal
-    # not returning.
-    add_definitions("/wd4722")
-  endif()
 else()
   list(APPEND INC
     include

extern/mantaflow/CMakeLists.txt

@@ -49,11 +49,6 @@ if(WITH_TBB)
   add_definitions(-DTBB=1)
 endif()
 
-if(WITH_OPENVDB)
-  add_definitions(-DOPENVDB=1)
-  add_definitions(-DOPENVDB_STATICLIB)
-endif()
-
 if(WIN32)
   add_definitions(-D_USE_MATH_DEFINES)
 endif()
@@ -83,6 +78,7 @@ if(WITH_TBB)
 endif()
 
 if(WITH_OPENVDB)
+  add_definitions(-DOPENVDB=1 ${OPENVDB_DEFINITIONS})
   list(APPEND INC_SYS
     ${OPENVDB_INCLUDE_DIRS}
   )

extern/mantaflow/helper/pwrapper/registry.cpp

@@ -192,7 +192,7 @@ int cbDisableConstructor(PyObject *self, PyObject *args, PyObject *kwds)
   return -1;
 }
 
-PyMODINIT_FUNC PyInit_manta_main(void)
+PyMODINIT_FUNC PyInit_Main(void)
 {
   MantaEnsureRegistration();
 #if PY_MAJOR_VERSION >= 3
@@ -567,7 +567,7 @@ void WrapperRegistry::construct(const string &scriptname, const vector<string> &
   registerDummyTypes();
 
   // work around for certain gcc versions, cast to char*
-  PyImport_AppendInittab((char *)gDefaultModuleName.c_str(), PyInit_manta_main);
+  PyImport_AppendInittab((char *)gDefaultModuleName.c_str(), PyInit_Main);
 }
 
 inline PyObject *castPy(PyTypeObject *p)

extern/mantaflow/helper/pwrapper/registry.h

@@ -62,7 +62,7 @@ void MantaEnsureRegistration();
 
 #ifdef BLENDER
 #  ifdef PyMODINIT_FUNC
-PyMODINIT_FUNC PyInit_manta_main(void);
+PyMODINIT_FUNC PyInit_Main(void);
 #  endif
 #endif
 

extern/mantaflow/helper/util/rcmatrix.h

@@ -17,13 +17,16 @@
 
 // link to omp & tbb for now
 #if OPENMP == 1 || TBB == 1
-#  define MANTA_ENABLE_PARALLEL 1
+#  define MANTA_ENABLE_PARALLEL 0
 // allow the preconditioner to be computed in parallel? (can lead to slightly non-deterministic
 // results)
 #  define MANTA_ENABLE_PARALLEL_PC 0
+// use c++11 code?
+#  define MANTA_USE_CPP11 1
 #else
 #  define MANTA_ENABLE_PARALLEL 0
 #  define MANTA_ENABLE_PARALLEL_PC 0
+#  define MANTA_USE_CPP11 0
 #endif
 
 #if MANTA_ENABLE_PARALLEL == 1
@@ -500,7 +503,11 @@ template<class N, class T> struct RCMatrix {
       for (Iterator it = row_begin(i); it; ++it)
         result.fix_element(it.index(), i, it.value());
     }
+#if MANTA_USE_CPP11 == 1
+    return std::move(result);
+#else
     return result;
+#endif
   }
 
   RCMatrix getKtK() const
@@ -525,7 +532,12 @@ template<class N, class T> struct RCMatrix {
         }
       }
     }
-    parallel_end return result;
+    parallel_end
+#if MANTA_USE_CPP11 == 1
+        return std::move(result);
+#else
+        return result;
+#endif
   }
 
   RCMatrix operator*(const RCMatrix &m) const
@@ -549,7 +561,12 @@ template<class N, class T> struct RCMatrix {
         }
       }
     }
-    parallel_end return result;
+    parallel_end
+#if MANTA_USE_CPP11 == 1
+        return std::move(result);
+#else
+        return result;
+#endif
   }
 
   RCMatrix sqrt() const
@@ -564,7 +581,12 @@ template<class N, class T> struct RCMatrix {
         result.set_element(i, j, std::sqrt(it_A.value()));
       }
     }
-    parallel_end return result;
+    parallel_end
+#if MANTA_USE_CPP11 == 1
+        return std::move(result);
+#else
+        return result;
+#endif
   }
 
   RCMatrix operator*(const double k) const
@@ -579,7 +601,12 @@ template<class N, class T> struct RCMatrix {
         result.add_to_element(i, j, it_A.value() * k);
       }
     }
-    parallel_end return result;
+    parallel_end
+#if MANTA_USE_CPP11 == 1
+        return std::move(result);
+#else
+        return result;
+#endif
   }
 
   RCMatrix applyKernel(const RCMatrix &kernel, const int nx, const int ny) const
@@ -613,7 +640,12 @@ template<class N, class T> struct RCMatrix {
         }
       }
     }
-    parallel_end return result;
+    parallel_end
+#if MANTA_USE_CPP11 == 1
+        return std::move(result);
+#else
+        return result;
+#endif
   }
 
   RCMatrix applyHorizontalKernel(const RCMatrix &kernel, const int nx, const int ny) const
@@ -647,7 +679,12 @@ template<class N, class T> struct RCMatrix {
         }
       }
     }
-    parallel_end return result;
+    parallel_end
+#if MANTA_USE_CPP11 == 1
+        return std::move(result);
+#else
+        return result;
+#endif
   }
 
   RCMatrix applyVerticalKernel(const RCMatrix &kernel, const int nx, const int ny) const
@@ -681,7 +718,12 @@ template<class N, class T> struct RCMatrix {
         }
       }
     }
-    parallel_end return result;
+    parallel_end
+#if MANTA_USE_CPP11 == 1
+        return std::move(result);
+#else
+        return result;
+#endif
   }
 
   RCMatrix applySeparableKernel(const RCMatrix &kernelH,
@@ -705,7 +747,11 @@ template<class N, class T> struct RCMatrix {
   {
     std::vector<T> result(n, 0.0);
     multiply(rhs, result);
+#if MANTA_USE_CPP11 == 1
+    return std::move(result);
+#else
     return result;
+#endif
   }
   void multiply(const std::vector<T> &rhs, std::vector<T> &result) const
   {
@@ -786,7 +832,11 @@ template<class N, class T> struct RCMatrix {
     for (N i = 0; i < result.size(); i++) {
       result[i] = std::abs(result[i] - rhs[i]);
     }
+#if MANTA_USE_CPP11 == 1
+    return std::move(result);
+#else
     return result;
+#endif
   }
   T norm() const
   {
@@ -974,7 +1024,11 @@ template<class N, class T> struct RCFixedMatrix {
   {
     std::vector<T> result(n, 0.0);
     multiply(rhs, result);
+#if MANTA_USE_CPP11 == 1
+    return std::move(result);
+#else
     return result;
+#endif
   }
   void multiply(const std::vector<T> &rhs, std::vector<T> &result) const
   {
@@ -1010,7 +1064,12 @@ template<class N, class T> struct RCFixedMatrix {
         }
       }
     }
-    parallel_end return result;
+    parallel_end
+#if MANTA_USE_CPP11 == 1
+        return std::move(result);
+#else
+        return result;
+#endif
   }
 
   RCMatrix<N, T> toRCMatrix() const
@@ -1028,7 +1087,12 @@ template<class N, class T> struct RCFixedMatrix {
         result.matrix[i]->value[j] = value[rowstart[i] + j];
       }
     }
-    parallel_end return result;
+    parallel_end
+#if MANTA_USE_CPP11 == 1
+        return std::move(result);
+#else
+        return result;
+#endif
   }
 };
 

extern/mantaflow/helper/util/vectorbase.h

@@ -15,7 +15,7 @@
 #define _VECTORBASE_H
 
 // get rid of windos min/max defines
-#if (defined(WIN32) || defined(_WIN32)) && !defined(NOMINMAX)
+#if defined(WIN32) || defined(_WIN32)
 #  define NOMINMAX
 #endif
 

extern/mantaflow/preprocessed/gitinfo.h

@@ -1,3 +1,3 @@
 
 
-#define MANTA_GIT_VERSION "commit 5fbd3d04381b21afce4a593d1fe2d9bc7bef5424"
+#define MANTA_GIT_VERSION "commit ce000bcbd7004e6549ac2f118755fcdc1f679bc3"

extern/mantaflow/preprocessed/plugin/initplugins.cpp

@@ -1829,15 +1829,13 @@ struct KnUpdateFlagsObs : public KernelBase {
                    const MACGrid *fractions,
                    const Grid<Real> &phiObs,
                    const Grid<Real> *phiOut,
-                   const Grid<Real> *phiIn,
-                   int boundaryWidth)
-      : KernelBase(&flags, boundaryWidth),
+                   const Grid<Real> *phiIn)
+      : KernelBase(&flags, 1),
         flags(flags),
         fractions(fractions),
         phiObs(phiObs),
         phiOut(phiOut),
-        phiIn(phiIn),
-        boundaryWidth(boundaryWidth)
+        phiIn(phiIn)
   {
     runMessage();
     run();
@@ -1849,8 +1847,7 @@ struct KnUpdateFlagsObs : public KernelBase {
                  const MACGrid *fractions,
                  const Grid<Real> &phiObs,
                  const Grid<Real> *phiOut,
-                 const Grid<Real> *phiIn,
-                 int boundaryWidth) const
+                 const Grid<Real> *phiIn) const
   {
 
     bool isObs = false;
@@ -1913,11 +1910,6 @@ struct KnUpdateFlagsObs : public KernelBase {
     return phiIn;
   }
   typedef Grid<Real> type4;
-  inline int &getArg5()
-  {
-    return boundaryWidth;
-  }
-  typedef int type5;
   void runMessage()
   {
     debMsg("Executing kernel KnUpdateFlagsObs ", 3);
@@ -1931,15 +1923,15 @@ struct KnUpdateFlagsObs : public KernelBase {
     const int _maxY = maxY;
     if (maxZ > 1) {
       for (int k = __r.begin(); k != (int)__r.end(); k++)
-        for (int j = boundaryWidth; j < _maxY; j++)
-          for (int i = boundaryWidth; i < _maxX; i++)
-            op(i, j, k, flags, fractions, phiObs, phiOut, phiIn, boundaryWidth);
+        for (int j = 1; j < _maxY; j++)
+          for (int i = 1; i < _maxX; i++)
+            op(i, j, k, flags, fractions, phiObs, phiOut, phiIn);
     }
     else {
       const int k = 0;
       for (int j = __r.begin(); j != (int)__r.end(); j++)
-        for (int i = boundaryWidth; i < _maxX; i++)
-          op(i, j, k, flags, fractions, phiObs, phiOut, phiIn, boundaryWidth);
+        for (int i = 1; i < _maxX; i++)
+          op(i, j, k, flags, fractions, phiObs, phiOut, phiIn);
     }
   }
   void run()
@@ -1947,14 +1939,13 @@ struct KnUpdateFlagsObs : public KernelBase {
     if (maxZ > 1)
       tbb::parallel_for(tbb::blocked_range<IndexInt>(minZ, maxZ), *this);
     else
-      tbb::parallel_for(tbb::blocked_range<IndexInt>(boundaryWidth, maxY), *this);
+      tbb::parallel_for(tbb::blocked_range<IndexInt>(1, maxY), *this);
   }
   FlagGrid &flags;
   const MACGrid *fractions;
   const Grid<Real> &phiObs;
   const Grid<Real> *phiOut;
   const Grid<Real> *phiIn;
-  int boundaryWidth;
 };
 
 //! update obstacle and outflow flags from levelsets
@@ -1963,10 +1954,9 @@ void setObstacleFlags(FlagGrid &flags,
                       const Grid<Real> &phiObs,
                       const MACGrid *fractions = NULL,
                       const Grid<Real> *phiOut = NULL,
-                      const Grid<Real> *phiIn = NULL,
-                      int boundaryWidth = 1)
+                      const Grid<Real> *phiIn = NULL)
 {
-  KnUpdateFlagsObs(flags, fractions, phiObs, phiOut, phiIn, boundaryWidth);
+  KnUpdateFlagsObs(flags, fractions, phiObs, phiOut, phiIn);
 }
 static PyObject *_W_18(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
 {
@@ -1983,9 +1973,8 @@ static PyObject *_W_18(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
       const MACGrid *fractions = _args.getPtrOpt<MACGrid>("fractions", 2, NULL, &_lock);
       const Grid<Real> *phiOut = _args.getPtrOpt<Grid<Real>>("phiOut", 3, NULL, &_lock);
       const Grid<Real> *phiIn = _args.getPtrOpt<Grid<Real>>("phiIn", 4, NULL, &_lock);
-      int boundaryWidth = _args.getOpt<int>("boundaryWidth", 5, 1, &_lock);
       _retval = getPyNone();
-      setObstacleFlags(flags, phiObs, fractions, phiOut, phiIn, boundaryWidth);
+      setObstacleFlags(flags, phiObs, fractions, phiOut, phiIn);
       _args.check();
     }
     pbFinalizePlugin(parent, "setObstacleFlags", !noTiming);

extern/mantaflow/preprocessed/plugin/pressure.cpp

@@ -1171,11 +1171,6 @@ void solvePressureSystem(Grid<Real> &rhs,
     maxIter = 100;
 
     pmg = gMapMG[parent];
-    // Release MG from previous step if present (e.g. if previous solve was with MGStatic)
-    if (pmg && preconditioner == PcMGDynamic) {
-      releaseMG(parent);
-      pmg = nullptr;
-    }
     if (!pmg) {
       pmg = new GridMg(pressure.getSize());
       gMapMG[parent] = pmg;

intern/atomic/intern/atomic_ops_msvc.h

@@ -37,8 +37,8 @@
 #endif
 #define WIN32_LEAN_AND_MEAN
 
-#include <intrin.h>
 #include <windows.h>
+#include <intrin.h>
 
 #if defined(__clang__)
 #  pragma GCC diagnostic push

intern/atomic/intern/atomic_ops_utils.h

@@ -47,10 +47,10 @@
 #ifndef __ATOMIC_OPS_UTILS_H__
 #define __ATOMIC_OPS_UTILS_H__
 
-#include <limits.h>
-#include <stddef.h>
 #include <stdint.h>
+#include <stddef.h>
 #include <stdlib.h>
+#include <limits.h>
 
 #include <assert.h>
 

intern/audaspace/intern/AUD_PyInit.cpp

@@ -25,8 +25,8 @@
 #include "AUD_PyInit.h"
 
 #include <AUD_Sound.h>
-#include <python/PyAPI.h>
 #include <python/PySound.h>
+#include <python/PyAPI.h>
 
 extern "C" {
 extern void *BKE_sound_get_factory(void *sound);

intern/clog/clog.c

@@ -18,23 +18,23 @@
  * \ingroup clog
  */
 
-#include <assert.h>
 #include <stdarg.h>
-#include <stdint.h>
 #include <stdlib.h>
 #include <string.h>
+#include <stdint.h>
+#include <assert.h>
 
 /* Disable for small single threaded programs
  * to avoid having to link with pthreads. */
 #ifdef WITH_CLOG_PTHREADS
-#  include "atomic_ops.h"
 #  include <pthread.h>
+#  include "atomic_ops.h"
 #endif
 
 /* For 'isatty' to check for color. */
 #if defined(__unix__) || defined(__APPLE__) || defined(__HAIKU__)
-#  include <sys/time.h>
 #  include <unistd.h>
+#  include <sys/time.h>
 #endif
 
 #if defined(_MSC_VER)

intern/cycles/CMakeLists.txt

@@ -177,11 +177,14 @@ if(CXX_HAS_AVX2)
   add_definitions(-DWITH_KERNEL_AVX2)
 endif()
 
-# LLVM and OSL need to build without RTTI
-if(WIN32 AND MSVC)
-  set(RTTI_DISABLE_FLAGS "/GR- -DBOOST_NO_RTTI -DBOOST_NO_TYPEID")
-elseif(CMAKE_COMPILER_IS_GNUCC OR (CMAKE_C_COMPILER_ID MATCHES "Clang"))
-  set(RTTI_DISABLE_FLAGS "-fno-rtti -DBOOST_NO_RTTI -DBOOST_NO_TYPEID")
+if(WITH_CYCLES_OSL)
+  # LLVM and OSL need to build without RTTI
+  if(WIN32 AND MSVC)
+    set(RTTI_DISABLE_FLAGS "/GR- -DBOOST_NO_RTTI -DBOOST_NO_TYPEID")
+  elseif(CMAKE_COMPILER_IS_GNUCC OR (CMAKE_C_COMPILER_ID MATCHES "Clang"))
+    set(RTTI_DISABLE_FLAGS "-fno-rtti -DBOOST_NO_RTTI -DBOOST_NO_TYPEID")
+  endif()
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${RTTI_DISABLE_FLAGS}")
 endif()
 
 # Definitions and Includes
@@ -350,6 +353,17 @@ if(WITH_CYCLES_NETWORK)
   add_definitions(-DWITH_NETWORK)
 endif()
 
+if(WITH_OPENCOLORIO)
+  add_definitions(-DWITH_OCIO)
+  include_directories(
+    SYSTEM
+    ${OPENCOLORIO_INCLUDE_DIRS}
+  )
+  if(WIN32)
+    add_definitions(-DOpenColorIO_STATIC)
+  endif()
+endif()
+
 if(WITH_CYCLES_STANDALONE OR WITH_CYCLES_NETWORK OR WITH_CYCLES_CUBIN_COMPILER)
   add_subdirectory(app)
 endif()

intern/cycles/app/CMakeLists.txt

@@ -51,17 +51,14 @@ endif()
 
 # Common configuration.
 
-link_directories(
-  ${OPENIMAGEIO_LIBPATH}
-  ${BOOST_LIBPATH}
-  ${PNG_LIBPATH}
-  ${JPEG_LIBPATH}
-  ${ZLIB_LIBPATH}
-  ${TIFF_LIBPATH}
-  ${OPENEXR_LIBPATH}
-  ${OPENJPEG_LIBPATH}
-  ${OPENVDB_LIBPATH}
-)
+link_directories(${OPENIMAGEIO_LIBPATH}
+                 ${BOOST_LIBPATH}
+                 ${PNG_LIBPATH}
+                 ${JPEG_LIBPATH}
+                 ${ZLIB_LIBPATH}
+                 ${TIFF_LIBPATH}
+                 ${OPENEXR_LIBPATH}
+                 ${OPENJPEG_LIBPATH})
 
 if(WITH_OPENCOLORIO)
   link_directories(${OPENCOLORIO_LIBPATH})

intern/cycles/app/cycles_cubin_cc.cpp

@@ -14,8 +14,8 @@
  * limitations under the License.
  */
 
-#include <stdint.h>
 #include <stdio.h>
+#include <stdint.h>
 
 #include <string>
 #include <vector>

intern/cycles/app/cycles_server.cpp

@@ -20,11 +20,11 @@
 
 #include "util/util_args.h"
 #include "util/util_foreach.h"
-#include "util/util_logging.h"
 #include "util/util_path.h"
 #include "util/util_stats.h"
 #include "util/util_string.h"
 #include "util/util_task.h"
+#include "util/util_logging.h"
 
 using namespace ccl;
 

intern/cycles/app/cycles_standalone.cpp

@@ -16,12 +16,12 @@
 
 #include <stdio.h>
 
-#include "device/device.h"
 #include "render/buffers.h"
 #include "render/camera.h"
-#include "render/integrator.h"
+#include "device/device.h"
 #include "render/scene.h"
 #include "render/session.h"
+#include "render/integrator.h"
 
 #include "util/util_args.h"
 #include "util/util_foreach.h"

intern/cycles/app/cycles_xml.cpp

@@ -16,9 +16,9 @@
 
 #include <stdio.h>
 
+#include <sstream>
 #include <algorithm>
 #include <iterator>
-#include <sstream>
 
 #include "graph/node_xml.h"
 
@@ -32,8 +32,8 @@
 #include "render/nodes.h"
 #include "render/object.h"
 #include "render/osl.h"
-#include "render/scene.h"
 #include "render/shader.h"
+#include "render/scene.h"
 
 #include "subd/subd_patch.h"
 #include "subd/subd_split.h"
@@ -292,7 +292,7 @@ static void xml_read_shader_graph(XMLReadState &state, Shader *shader, xml_node
             filepath = path_join(state.base, filepath);
           }
 
-          snode = OSLShaderManager::osl_node(manager, filepath);
+          snode = ((OSLShaderManager *)manager)->osl_node(filepath);
 
           if (!snode) {
             fprintf(stderr, "Failed to create OSL node from \"%s\".\n", filepath.c_str());

intern/cycles/blender/CMakeLists.txt

@@ -38,7 +38,6 @@ set(SRC
   CCL_api.h
   blender_device.h
   blender_id_map.h
-  blender_image.h
   blender_object_cull.h
   blender_sync.h
   blender_session.h
@@ -92,20 +91,6 @@ if(WITH_MOD_FLUID)
   add_definitions(-DWITH_FLUID)
 endif()
 
-if(WITH_NEW_OBJECT_TYPES)
-  add_definitions(-DWITH_NEW_OBJECT_TYPES)
-endif()
-
-if(WITH_OPENVDB)
-  add_definitions(-DWITH_OPENVDB ${OPENVDB_DEFINITIONS})
-  list(APPEND INC_SYS
-    ${OPENVDB_INCLUDE_DIRS}
-  )
-  list(APPEND LIB
-    ${OPENVDB_LIBRARIES}
-  )
-endif()
-
 blender_add_lib(bf_intern_cycles "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")
 
 # avoid link failure with clang 3.4 debug

intern/cycles/blender/addon/__init__.py

@@ -22,7 +22,7 @@ bl_info = {
     "blender": (2, 80, 0),
     "description": "Cycles renderer integration",
     "warning": "",
-    "doc_url": "https://docs.blender.org/manual/en/latest/render/cycles/",
+    "wiki_url": "https://docs.blender.org/manual/en/latest/render/cycles/",
     "tracker_url": "",
     "support": 'OFFICIAL',
     "category": "Render"}

intern/cycles/blender/addon/engine.py

@@ -33,7 +33,7 @@ def _is_using_buggy_driver():
             # in the version string, but those cards do not quite work and
             # causing crashes.
             return True
-        regex = re.compile(".*Compatibility Profile Context ([0-9]+(\\.[0-9]+)+)$")
+        regex = re.compile(".*Compatibility Profile Context ([0-9]+(\.[0-9]+)+)$")
         if not regex.match(version):
             # Skip cards like FireGL
             return False
@@ -255,7 +255,6 @@ def list_render_passes(srl):
     if crl.pass_debug_bvh_traversed_instances: yield ("Debug BVH Traversed Instances", "X",   'VALUE')
     if crl.pass_debug_bvh_intersections:       yield ("Debug BVH Intersections",       "X",   'VALUE')
     if crl.pass_debug_ray_bounces:             yield ("Debug Ray Bounces",             "X",   'VALUE')
-    if crl.pass_debug_sample_count:            yield ("Debug Sample Count",            "X",   'VALUE')
     if crl.use_pass_volume_direct:             yield ("VolumeDir",                     "RGB", 'COLOR')
     if crl.use_pass_volume_indirect:           yield ("VolumeInd",                     "RGB", 'COLOR')
 

intern/cycles/blender/addon/properties.py

@@ -112,7 +112,6 @@ enum_use_layer_samples = (
 enum_sampling_pattern = (
     ('SOBOL', "Sobol", "Use Sobol random sampling pattern"),
     ('CORRELATED_MUTI_JITTER', "Correlated Multi-Jitter", "Use Correlated Multi-Jitter random sampling pattern"),
-    ('PROGRESSIVE_MUTI_JITTER', "Progressive Multi-Jitter", "Use Progressive Multi-Jitter random sampling pattern"),
 )
 
 enum_integrator = (
@@ -350,26 +349,6 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         default=0.01,
     )
 
-    use_adaptive_sampling: BoolProperty(
-        name="Use Adaptive Sampling",
-        description="Automatically reduce the number of samples per pixel based on estimated noise level",
-        default=False,
-    )
-
-    adaptive_threshold: FloatProperty(
-        name="Adaptive Sampling Threshold",
-        description="Noise level step to stop sampling at, lower values reduce noise the cost of render time. Zero for automatic setting based on number of AA samples",
-        min=0.0, max=1.0,
-        default=0.0,
-        precision=4,
-    )
-    adaptive_min_samples: IntProperty(
-        name="Adaptive Min Samples",
-        description="Minimum AA samples for adaptive sampling, to discover noisy features before stopping sampling. Zero for automatic setting based on number of AA samples",
-        min=0, max=4096,
-        default=0,
-    )
-
     min_light_bounces: IntProperty(
             name="Min Light Bounces",
             description="Minimum number of light bounces. Setting this higher reduces noise in the first bounces, "
@@ -444,20 +423,13 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
         default=8,
     )
 
-    volume_step_rate: FloatProperty(
-        name="Step Rate",
-        description="Globally adjust detail for volume rendering, on top of automatically estimated step size. "
-                    "Higher values reduce render time, lower values render with more detail",
-        default=1.0,
-        min=0.01, max=100.0, soft_min=0.1, soft_max=10.0, precision=2
-    )
-
-    volume_preview_step_rate: FloatProperty(
-        name="Step Rate",
-        description="Globally adjust detail for volume rendering, on top of automatically estimated step size. "
-                    "Higher values reduce render time, lower values render with more detail",
-        default=1.0,
-        min=0.01, max=100.0, soft_min=0.1, soft_max=10.0, precision=2
+    volume_step_size: FloatProperty(
+        name="Step Size",
+        description="Distance between volume shader samples when rendering the volume "
+        "(lower values give more accurate and detailed results, but also increased render time)",
+        default=0.1,
+        min=0.0000001, max=100000.0, soft_min=0.01, soft_max=1.0, precision=4,
+        unit='LENGTH'
     )
 
     volume_max_steps: IntProperty(
@@ -941,14 +913,6 @@ class CyclesMaterialSettings(bpy.types.PropertyGroup):
         default='LINEAR',
     )
 
-    volume_step_rate: FloatProperty(
-        name="Step Rate",
-        description="Scale the distance between volume shader samples when rendering the volume "
-                    "(lower values give more accurate and detailed results, but also increased render time)",
-        default=1.0,
-        min=0.001, max=1000.0, soft_min=0.1, soft_max=10.0, precision=4
-    )
-
     displacement_method: EnumProperty(
         name="Displacement Method",
         description="Method to use for the displacement",
@@ -1059,13 +1023,6 @@ class CyclesWorldSettings(bpy.types.PropertyGroup):
         items=enum_volume_interpolation,
         default='LINEAR',
     )
-    volume_step_size: FloatProperty(
-        name="Step Size",
-        description="Distance between volume shader samples when rendering the volume "
-                    "(lower values give more accurate and detailed results, but also increased render time)",
-        default=1.0,
-        min=0.0000001, max=100000.0, soft_min=0.1, soft_max=100.0, precision=4
-    )
 
     @classmethod
     def register(cls):
@@ -1340,12 +1297,7 @@ class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
         default=False,
         update=update_render_passes,
     )
-    pass_debug_sample_count: BoolProperty(
-        name="Debug Sample Count",
-        description="Number of samples/camera rays per pixel",
-        default=False,
-        update=update_render_passes,
-    )
+
     use_pass_volume_direct: BoolProperty(
         name="Volume Direct",
         description="Deliver direct volumetric scattering pass",

intern/cycles/blender/addon/ui.py

@@ -230,32 +230,6 @@ class CYCLES_RENDER_PT_sampling_sub_samples(CyclesButtonsPanel, Panel):
         draw_samples_info(layout, context)
 
 
-class CYCLES_RENDER_PT_sampling_adaptive(CyclesButtonsPanel, Panel):
-    bl_label = "Adaptive Sampling"
-    bl_parent_id = "CYCLES_RENDER_PT_sampling"
-    bl_options = {'DEFAULT_CLOSED'}
-
-    def draw_header(self, context):
-        layout = self.layout
-        scene = context.scene
-        cscene = scene.cycles
-
-        layout.prop(cscene, "use_adaptive_sampling", text="")
-
-    def draw(self, context):
-        layout = self.layout
-        layout.use_property_split = True
-        layout.use_property_decorate = False
-
-        scene = context.scene
-        cscene = scene.cycles
-
-        layout.active = cscene.use_adaptive_sampling
-
-        col = layout.column(align=True)
-        col.prop(cscene, "adaptive_min_samples", text="Min Samples")
-        col.prop(cscene, "adaptive_threshold", text="Noise Threshold")
-
 class CYCLES_RENDER_PT_sampling_advanced(CyclesButtonsPanel, Panel):
     bl_label = "Advanced"
     bl_parent_id = "CYCLES_RENDER_PT_sampling"
@@ -273,9 +247,7 @@ class CYCLES_RENDER_PT_sampling_advanced(CyclesButtonsPanel, Panel):
         row.prop(cscene, "seed")
         row.prop(cscene, "use_animated_seed", text="", icon='TIME')
 
-        col = layout.column(align=True)
-        col.active = not(cscene.use_adaptive_sampling)
-        col.prop(cscene, "sampling_pattern", text="Pattern")
+        layout.prop(cscene, "sampling_pattern", text="Pattern")
 
         layout.prop(cscene, "use_square_samples")
 
@@ -373,7 +345,7 @@ class CYCLES_RENDER_PT_subdivision(CyclesButtonsPanel, Panel):
         col = layout.column()
         sub = col.column(align=True)
         sub.prop(cscene, "dicing_rate", text="Dicing Rate Render")
-        sub.prop(cscene, "preview_dicing_rate", text="Viewport")
+        sub.prop(cscene, "preview_dicing_rate", text="Preview")
 
         col.separator()
 
@@ -428,11 +400,9 @@ class CYCLES_RENDER_PT_volumes(CyclesButtonsPanel, Panel):
         scene = context.scene
         cscene = scene.cycles
 
-        col = layout.column(align=True)
-        col.prop(cscene, "volume_step_rate", text="Step Rate Render")
-        col.prop(cscene, "volume_preview_step_rate", text="Viewport")
-
-        layout.prop(cscene, "volume_max_steps", text="Max Steps")
+        col = layout.column()
+        col.prop(cscene, "volume_step_size", text="Step Size")
+        col.prop(cscene, "volume_max_steps", text="Max Steps")
 
 
 class CYCLES_RENDER_PT_light_paths(CyclesButtonsPanel, Panel):
@@ -772,8 +742,6 @@ class CYCLES_RENDER_PT_filter(CyclesButtonsPanel, Panel):
         col.prop(view_layer, "use_solid", text="Surfaces")
         col = flow.column()
         col.prop(view_layer, "use_strand", text="Hair")
-        col = flow.column()
-        col.prop(view_layer, "use_volumes", text="Volumes")
         if with_freestyle:
             col = flow.column()
             col.prop(view_layer, "use_freestyle", text="Freestyle")
@@ -845,8 +813,6 @@ class CYCLES_RENDER_PT_passes_data(CyclesButtonsPanel, Panel):
         col.prop(cycles_view_layer, "denoising_store_passes", text="Denoising Data")
         col = flow.column()
         col.prop(cycles_view_layer, "pass_debug_render_time", text="Render Time")
-        col = flow.column()
-        col.prop(cycles_view_layer, "pass_debug_sample_count", text="Sample Count")
 
         layout.separator()
 
@@ -1417,6 +1383,8 @@ class CYCLES_LIGHT_PT_light(CyclesButtonsPanel, Panel):
         light = context.light
         clamp = light.cycles
 
+        layout.use_property_decorate = False
+
         if self.bl_space_type == 'PROPERTIES':
             layout.row().prop(light, "type", expand=True)
             layout.use_property_split = True
@@ -1698,9 +1666,6 @@ class CYCLES_WORLD_PT_settings_volume(CyclesButtonsPanel, Panel):
         sub.prop(cworld, "volume_sampling", text="Sampling")
         col.prop(cworld, "volume_interpolation", text="Interpolation")
         col.prop(cworld, "homogeneous_volume", text="Homogeneous")
-        sub = col.column()
-        sub.active = not cworld.homogeneous_volume
-        sub.prop(cworld, "volume_step_size")
 
 
 class CYCLES_MATERIAL_PT_preview(CyclesButtonsPanel, Panel):
@@ -1832,9 +1797,6 @@ class CYCLES_MATERIAL_PT_settings_volume(CyclesButtonsPanel, Panel):
         sub.prop(cmat, "volume_sampling", text="Sampling")
         col.prop(cmat, "volume_interpolation", text="Interpolation")
         col.prop(cmat, "homogeneous_volume", text="Homogeneous")
-        sub = col.column()
-        sub.active = not cmat.homogeneous_volume
-        sub.prop(cmat, "volume_step_rate")
 
     def draw(self, context):
         self.draw_shared(self, context, context.material)
@@ -2276,7 +2238,6 @@ classes = (
     CYCLES_PT_integrator_presets,
     CYCLES_RENDER_PT_sampling,
     CYCLES_RENDER_PT_sampling_sub_samples,
-    CYCLES_RENDER_PT_sampling_adaptive,
     CYCLES_RENDER_PT_sampling_advanced,
     CYCLES_RENDER_PT_light_paths,
     CYCLES_RENDER_PT_light_paths_max_bounces,

intern/cycles/blender/blender_curves.cpp

@@ -595,6 +595,40 @@ static void ExportCurveTriangleGeometry(Mesh *mesh, ParticleCurveData *CData, in
   /* texture coords still needed */
 }
 
+static void export_hair_motion_validate_attribute(Hair *hair,
+                                                  int motion_step,
+                                                  int num_motion_keys,
+                                                  bool have_motion)
+{
+  Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+  const int num_keys = hair->curve_keys.size();
+
+  if (num_motion_keys != num_keys || !have_motion) {
+    /* No motion or hair "topology" changed, remove attributes again. */
+    if (num_motion_keys != num_keys) {
+      VLOG(1) << "Hair topology changed, removing attribute.";
+    }
+    else {
+      VLOG(1) << "No motion, removing attribute.";
+    }
+    hair->attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION);
+  }
+  else if (motion_step > 0) {
+    VLOG(1) << "Filling in new motion vertex position for motion_step " << motion_step;
+
+    /* Motion, fill up previous steps that we might have skipped because
+     * they had no motion, but we need them anyway now. */
+    for (int step = 0; step < motion_step; step++) {
+      float4 *mP = attr_mP->data_float4() + step * num_keys;
+
+      for (int key = 0; key < num_keys; key++) {
+        mP[key] = float3_to_float4(hair->curve_keys[key]);
+        mP[key].w = hair->curve_radius[key];
+      }
+    }
+  }
+}
+
 static void ExportCurveSegments(Scene *scene, Hair *hair, ParticleCurveData *CData)
 {
   int num_keys = 0;
@@ -714,40 +748,6 @@ static float4 LerpCurveSegmentMotionCV(ParticleCurveData *CData, int sys, int cu
   return lerp(mP, mP2, remainder);
 }
 
-static void export_hair_motion_validate_attribute(Hair *hair,
-                                                  int motion_step,
-                                                  int num_motion_keys,
-                                                  bool have_motion)
-{
-  Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
-  const int num_keys = hair->curve_keys.size();
-
-  if (num_motion_keys != num_keys || !have_motion) {
-    /* No motion or hair "topology" changed, remove attributes again. */
-    if (num_motion_keys != num_keys) {
-      VLOG(1) << "Hair topology changed, removing attribute.";
-    }
-    else {
-      VLOG(1) << "No motion, removing attribute.";
-    }
-    hair->attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION);
-  }
-  else if (motion_step > 0) {
-    VLOG(1) << "Filling in new motion vertex position for motion_step " << motion_step;
-
-    /* Motion, fill up previous steps that we might have skipped because
-     * they had no motion, but we need them anyway now. */
-    for (int step = 0; step < motion_step; step++) {
-      float4 *mP = attr_mP->data_float4() + step * num_keys;
-
-      for (int key = 0; key < num_keys; key++) {
-        mP[key] = float3_to_float4(hair->curve_keys[key]);
-        mP[key].w = hair->curve_radius[key];
-      }
-    }
-  }
-}
-
 static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, int motion_step)
 {
   VLOG(1) << "Exporting curve motion segments for hair " << hair->name << ", motion step "
@@ -817,7 +817,7 @@ static void ExportCurveSegmentsMotion(Hair *hair, ParticleCurveData *CData, int
     }
   }
 
-  /* In case of new attribute, we verify if there really was any motion. */
+  /* in case of new attribute, we verify if there really was any motion */
   if (new_attribute) {
     export_hair_motion_validate_attribute(hair, motion_step, i, have_motion);
   }
@@ -1154,194 +1154,6 @@ void BlenderSync::sync_particle_hair(
   }
 }
 
-#ifdef WITH_NEW_OBJECT_TYPES
-static float4 hair_point_as_float4(BL::HairPoint b_point)
-{
-  float4 mP = float3_to_float4(get_float3(b_point.co()));
-  mP.w = b_point.radius();
-  return mP;
-}
-
-static float4 interpolate_hair_points(BL::Hair b_hair,
-                                      const int first_point_index,
-                                      const int num_points,
-                                      const float step)
-{
-  const float curve_t = step * (num_points - 1);
-  const int point_a = clamp((int)curve_t, 0, num_points - 1);
-  const int point_b = min(point_a + 1, num_points - 1);
-  const float t = curve_t - (float)point_a;
-  return lerp(hair_point_as_float4(b_hair.points[first_point_index + point_a]),
-              hair_point_as_float4(b_hair.points[first_point_index + point_b]),
-              t);
-}
-
-static void export_hair_curves(Scene *scene, Hair *hair, BL::Hair b_hair)
-{
-  /* TODO: optimize so we can straight memcpy arrays from Blender? */
-
-  /* Add requested attributes. */
-  Attribute *attr_intercept = NULL;
-  Attribute *attr_random = NULL;
-
-  if (hair->need_attribute(scene, ATTR_STD_CURVE_INTERCEPT)) {
-    attr_intercept = hair->attributes.add(ATTR_STD_CURVE_INTERCEPT);
-  }
-  if (hair->need_attribute(scene, ATTR_STD_CURVE_RANDOM)) {
-    attr_random = hair->attributes.add(ATTR_STD_CURVE_RANDOM);
-  }
-
-  /* Reserve memory. */
-  const int num_keys = b_hair.points.length();
-  const int num_curves = b_hair.curves.length();
-
-  if (num_curves > 0) {
-    VLOG(1) << "Exporting curve segments for hair " << hair->name;
-  }
-
-  hair->reserve_curves(num_curves, num_keys);
-
-  /* Export curves and points. */
-  vector<float> points_length;
-
-  BL::Hair::curves_iterator b_curve_iter;
-  for (b_hair.curves.begin(b_curve_iter); b_curve_iter != b_hair.curves.end(); ++b_curve_iter) {
-    BL::HairCurve b_curve = *b_curve_iter;
-    const int first_point_index = b_curve.first_point_index();
-    const int num_points = b_curve.num_points();
-
-    float3 prev_co = make_float3(0.0f, 0.0f, 0.0f);
-    float length = 0.0f;
-    if (attr_intercept) {
-      points_length.clear();
-      points_length.reserve(num_points);
-    }
-
-    /* Position and radius. */
-    for (int i = 0; i < num_points; i++) {
-      BL::HairPoint b_point = b_hair.points[first_point_index + i];
-
-      const float3 co = get_float3(b_point.co());
-      const float radius = b_point.radius();
-      hair->add_curve_key(co, radius);
-
-      if (attr_intercept) {
-        if (i > 0) {
-          length += len(co - prev_co);
-          points_length.push_back(length);
-        }
-        prev_co = co;
-      }
-    }
-
-    /* Normalized 0..1 attribute along curve. */
-    if (attr_intercept) {
-      for (int i = 0; i < num_points; i++) {
-        attr_intercept->add((length == 0.0f) ? 0.0f : points_length[i] / length);
-      }
-    }
-
-    /* Random number per curve. */
-    if (attr_random != NULL) {
-      attr_random->add(hash_uint2_to_float(b_curve.index(), 0));
-    }
-
-    /* Curve. */
-    const int shader_index = 0;
-    hair->add_curve(first_point_index, shader_index);
-  }
-}
-
-static void export_hair_curves_motion(Hair *hair, BL::Hair b_hair, int motion_step)
-{
-  VLOG(1) << "Exporting curve motion segments for hair " << hair->name << ", motion step "
-          << motion_step;
-
-  /* Find or add attribute. */
-  Attribute *attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
-  bool new_attribute = false;
-
-  if (!attr_mP) {
-    VLOG(1) << "Creating new motion vertex position attribute";
-    attr_mP = hair->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
-    new_attribute = true;
-  }
-
-  /* Export motion keys. */
-  const int num_keys = hair->curve_keys.size();
-  float4 *mP = attr_mP->data_float4() + motion_step * num_keys;
-  bool have_motion = false;
-  int num_motion_keys = 0;
-  int curve_index = 0;
-
-  BL::Hair::curves_iterator b_curve_iter;
-  for (b_hair.curves.begin(b_curve_iter); b_curve_iter != b_hair.curves.end(); ++b_curve_iter) {
-    BL::HairCurve b_curve = *b_curve_iter;
-    const int first_point_index = b_curve.first_point_index();
-    const int num_points = b_curve.num_points();
-
-    Hair::Curve curve = hair->get_curve(curve_index);
-    curve_index++;
-
-    if (num_points == curve.num_keys) {
-      /* Number of keys matches. */
-      for (int i = 0; i < num_points; i++) {
-        int point_index = first_point_index + i;
-
-        if (point_index < num_keys) {
-          mP[num_motion_keys] = hair_point_as_float4(b_hair.points[point_index]);
-          num_motion_keys++;
-
-          if (!have_motion) {
-            /* TODO: use epsilon for comparison? Was needed for particles due to
-             * transform, but ideally should not happen anymore. */
-            float4 curve_key = float3_to_float4(hair->curve_keys[i]);
-            curve_key.w = hair->curve_radius[i];
-            have_motion = !(mP[i] == curve_key);
-          }
-        }
-      }
-    }
-    else {
-      /* Number of keys has changed. Generate an interpolated version
-       * to preserve motion blur. */
-      const float step_size = curve.num_keys > 1 ? 1.0f / (curve.num_keys - 1) : 0.0f;
-      for (int i = 0; i < curve.num_keys; i++) {
-        const float step = i * step_size;
-        mP[num_motion_keys] = interpolate_hair_points(b_hair, first_point_index, num_points, step);
-        num_motion_keys++;
-      }
-      have_motion = true;
-    }
-  }
-
-  /* In case of new attribute, we verify if there really was any motion. */
-  if (new_attribute) {
-    export_hair_motion_validate_attribute(hair, motion_step, num_motion_keys, have_motion);
-  }
-}
-#endif /* WITH_NEW_OBJECT_TYPES */
-
-/* Hair object. */
-void BlenderSync::sync_hair(Hair *hair, BL::Object &b_ob, bool motion, int motion_step)
-{
-#ifdef WITH_NEW_OBJECT_TYPES
-  /* Convert Blender hair to Cycles curves. */
-  BL::Hair b_hair(b_ob.data());
-  if (motion) {
-    export_hair_curves_motion(hair, b_hair, motion_step);
-  }
-  else {
-    export_hair_curves(scene, hair, b_hair);
-  }
-#else
-  (void)hair;
-  (void)b_ob;
-  (void)motion;
-  (void)motion_step;
-#endif /* WITH_NEW_OBJECT_TYPES */
-}
-
 void BlenderSync::sync_hair(BL::Depsgraph b_depsgraph,
                             BL::Object b_ob,
                             Geometry *geom,
@@ -1367,24 +1179,14 @@ void BlenderSync::sync_hair(BL::Depsgraph b_depsgraph,
   geom->used_shaders = used_shaders;
 
   if (view_layer.use_hair && scene->curve_system_manager->use_curves) {
-#ifdef WITH_NEW_OBJECT_TYPES
-    if (b_ob.type() == BL::Object::type_HAIR) {
-      /* Hair object. */
-      sync_hair(hair, b_ob, false);
-      assert(mesh == NULL);
-    }
-    else
-#endif
-    {
-      /* Particle hair. */
-      bool need_undeformed = geom->need_attribute(scene, ATTR_STD_GENERATED);
-      BL::Mesh b_mesh = object_to_mesh(
-          b_data, b_ob, b_depsgraph, need_undeformed, Mesh::SUBDIVISION_NONE);
+    /* Particle hair. */
+    bool need_undeformed = geom->need_attribute(scene, ATTR_STD_GENERATED);
+    BL::Mesh b_mesh = object_to_mesh(
+        b_data, b_ob, b_depsgraph, need_undeformed, Mesh::SUBDIVISION_NONE);
 
-      if (b_mesh) {
-        sync_particle_hair(geom, b_mesh, b_ob, false);
-        free_object_to_mesh(b_data, b_ob, b_mesh);
-      }
+    if (b_mesh) {
+      sync_particle_hair(geom, b_mesh, b_ob, false);
+      free_object_to_mesh(b_data, b_ob, b_mesh);
     }
   }
 
@@ -1411,24 +1213,13 @@ void BlenderSync::sync_hair_motion(BL::Depsgraph b_depsgraph,
 
   /* Export deformed coordinates. */
   if (ccl::BKE_object_is_deform_modified(b_ob, b_scene, preview)) {
-#ifdef WITH_NEW_OBJECT_TYPES
-    if (b_ob.type() == BL::Object::type_HAIR) {
-      /* Hair object. */
-      sync_hair(hair, b_ob, true, motion_step);
-      assert(mesh == NULL);
+    /* Particle hair. */
+    BL::Mesh b_mesh = object_to_mesh(b_data, b_ob, b_depsgraph, false, Mesh::SUBDIVISION_NONE);
+    if (b_mesh) {
+      sync_particle_hair(geom, b_mesh, b_ob, true, motion_step);
+      free_object_to_mesh(b_data, b_ob, b_mesh);
       return;
     }
-    else
-#endif
-    {
-      /* Particle hair. */
-      BL::Mesh b_mesh = object_to_mesh(b_data, b_ob, b_depsgraph, false, Mesh::SUBDIVISION_NONE);
-      if (b_mesh) {
-        sync_particle_hair(geom, b_mesh, b_ob, true, motion_step);
-        free_object_to_mesh(b_data, b_ob, b_mesh);
-        return;
-      }
-    }
   }
 
   /* No deformation on this frame, copy coordinates if other frames did have it. */

intern/cycles/blender/blender_device.cpp

@@ -17,8 +17,6 @@
 #include "blender/blender_device.h"
 #include "blender/blender_util.h"
 
-#include "util/util_foreach.h"
-
 CCL_NAMESPACE_BEGIN
 
 enum DenoiserType {

intern/cycles/blender/blender_device.h

@@ -18,9 +18,9 @@
 #define __BLENDER_DEVICE_H__
 
 #include "MEM_guardedalloc.h"
+#include "RNA_types.h"
 #include "RNA_access.h"
 #include "RNA_blender_cpp.h"
-#include "RNA_types.h"
 
 #include "device/device.h"
 

intern/cycles/blender/blender_geometry.cpp

@@ -23,8 +23,6 @@
 #include "blender/blender_sync.h"
 #include "blender/blender_util.h"
 
-#include "util/util_foreach.h"
-
 CCL_NAMESPACE_BEGIN
 
 Geometry *BlenderSync::sync_geometry(BL::Depsgraph &b_depsgraph,
@@ -38,19 +36,11 @@ Geometry *BlenderSync::sync_geometry(BL::Depsgraph &b_depsgraph,
   BL::ID b_key_id = (BKE_object_is_modified(b_ob)) ? b_ob_instance : b_ob_data;
   GeometryKey key(b_key_id.ptr.data, use_particle_hair);
   BL::Material material_override = view_layer.material_override;
-  Shader *default_shader = (b_ob.type() == BL::Object::type_VOLUME) ? scene->default_volume :
-                                                                      scene->default_surface;
-#ifdef WITH_NEW_OBJECT_TYPES
-  Geometry::Type geom_type = ((b_ob.type() == BL::Object::type_HAIR || use_particle_hair) &&
+  Shader *default_shader = scene->default_surface;
+  Geometry::Type geom_type = (use_particle_hair &&
                               (scene->curve_system_manager->primitive != CURVE_TRIANGLES)) ?
                                  Geometry::HAIR :
                                  Geometry::MESH;
-#else
-  Geometry::Type geom_type = ((use_particle_hair) &&
-                              (scene->curve_system_manager->primitive != CURVE_TRIANGLES)) ?
-                                 Geometry::HAIR :
-                                 Geometry::MESH;
-#endif
 
   /* Find shader indices. */
   vector<Shader *> used_shaders;
@@ -129,14 +119,10 @@ Geometry *BlenderSync::sync_geometry(BL::Depsgraph &b_depsgraph,
 
   geom->name = ustring(b_ob_data.name().c_str());
 
-#ifdef WITH_NEW_OBJECT_TYPES
-  if (b_ob.type() == BL::Object::type_HAIR || use_particle_hair) {
-#else
   if (use_particle_hair) {
-#endif
     sync_hair(b_depsgraph, b_ob, geom, used_shaders);
   }
-  else if (b_ob.type() == BL::Object::type_VOLUME || object_fluid_gas_domain_find(b_ob)) {
+  else if (object_fluid_gas_domain_find(b_ob)) {
     Mesh *mesh = static_cast<Mesh *>(geom);
     sync_volume(b_ob, mesh, used_shaders);
   }
@@ -173,14 +159,10 @@ void BlenderSync::sync_geometry_motion(BL::Depsgraph &b_depsgraph,
     return;
   }
 
-#ifdef WITH_NEW_OBJECT_TYPES
-  if (b_ob.type() == BL::Object::type_HAIR || use_particle_hair) {
-#else
   if (use_particle_hair) {
-#endif
     sync_hair_motion(b_depsgraph, b_ob, geom, motion_step);
   }
-  else if (b_ob.type() == BL::Object::type_VOLUME || object_fluid_gas_domain_find(b_ob)) {
+  else if (object_fluid_gas_domain_find(b_ob)) {
     /* No volume motion blur support yet. */
   }
   else {

intern/cycles/blender/blender_image.cpp

@@ -14,71 +14,206 @@
  * limitations under the License.
  */
 
-#include "MEM_guardedalloc.h"
+#include "render/image.h"
 
-#include "blender/blender_image.h"
+#include "blender/blender_sync.h"
 #include "blender/blender_session.h"
 #include "blender/blender_util.h"
 
 CCL_NAMESPACE_BEGIN
 
-/* Packed Images */
-
-BlenderImageLoader::BlenderImageLoader(BL::Image b_image, int frame)
-    : b_image(b_image), frame(frame), free_cache(!b_image.has_data())
+/* builtin image file name is actually an image datablock name with
+ * absolute sequence frame number concatenated via '@' character
+ *
+ * this function splits frame from builtin name
+ */
+int BlenderSession::builtin_image_frame(const string &builtin_name)
 {
+  int last = builtin_name.find_last_of('@');
+  return atoi(builtin_name.substr(last + 1, builtin_name.size() - last - 1).c_str());
 }
 
-bool BlenderImageLoader::load_metadata(ImageMetaData &metadata)
+void BlenderSession::builtin_image_info(const string &builtin_name,
+                                        void *builtin_data,
+                                        ImageMetaData &metadata)
 {
-  metadata.width = b_image.size()[0];
-  metadata.height = b_image.size()[1];
-  metadata.depth = 1;
-  metadata.channels = b_image.channels();
+  /* empty image */
+  metadata.width = 1;
+  metadata.height = 1;
 
-  if (b_image.is_float()) {
-    if (metadata.channels == 1) {
-      metadata.type = IMAGE_DATA_TYPE_FLOAT;
+  if (!builtin_data)
+    return;
+
+  /* recover ID pointer */
+  PointerRNA ptr;
+  RNA_id_pointer_create((ID *)builtin_data, &ptr);
+  BL::ID b_id(ptr);
+
+  if (b_id.is_a(&RNA_Image)) {
+    /* image data */
+    BL::Image b_image(b_id);
+
+    metadata.builtin_free_cache = !b_image.has_data();
+    metadata.is_float = b_image.is_float();
+    metadata.width = b_image.size()[0];
+    metadata.height = b_image.size()[1];
+    metadata.depth = 1;
+    metadata.channels = b_image.channels();
+
+    if (metadata.is_float) {
+      /* Float images are already converted on the Blender side,
+       * no need to do anything in Cycles. */
+      metadata.colorspace = u_colorspace_raw;
     }
-    else if (metadata.channels == 4) {
-      metadata.type = IMAGE_DATA_TYPE_FLOAT4;
-    }
-    else {
-      return false;
+  }
+  else if (b_id.is_a(&RNA_Object)) {
+    /* smoke volume data */
+    BL::Object b_ob(b_id);
+    BL::FluidDomainSettings b_domain = object_fluid_gas_domain_find(b_ob);
+
+    metadata.is_float = true;
+    metadata.depth = 1;
+    metadata.channels = 1;
+
+    if (!b_domain)
+      return;
+
+    if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_DENSITY) ||
+        builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_FLAME) ||
+        builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_HEAT) ||
+        builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_TEMPERATURE))
+      metadata.channels = 1;
+    else if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_COLOR))
+      metadata.channels = 4;
+    else if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_VELOCITY))
+      metadata.channels = 3;
+    else
+      return;
+
+    int3 resolution = get_int3(b_domain.domain_resolution());
+    int amplify = (b_domain.use_noise()) ? b_domain.noise_scale() : 1;
+
+    /* Velocity and heat data is always low-resolution. */
+    if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_VELOCITY) ||
+        builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_HEAT)) {
+      amplify = 1;
     }
 
-    /* Float images are already converted on the Blender side,
-     * no need to do anything in Cycles. */
-    metadata.colorspace = u_colorspace_raw;
+    metadata.width = resolution.x * amplify;
+    metadata.height = resolution.y * amplify;
+    metadata.depth = resolution.z * amplify;
   }
   else {
-    if (metadata.channels == 1) {
-      metadata.type = IMAGE_DATA_TYPE_BYTE;
+    /* TODO(sergey): Check we're indeed in shader node tree. */
+    PointerRNA ptr;
+    RNA_pointer_create(NULL, &RNA_Node, builtin_data, &ptr);
+    BL::Node b_node(ptr);
+    if (b_node.is_a(&RNA_ShaderNodeTexPointDensity)) {
+      BL::ShaderNodeTexPointDensity b_point_density_node(b_node);
+      metadata.channels = 4;
+      metadata.width = b_point_density_node.resolution();
+      metadata.height = metadata.width;
+      metadata.depth = metadata.width;
+      metadata.is_float = true;
     }
-    else if (metadata.channels == 4) {
-      metadata.type = IMAGE_DATA_TYPE_BYTE4;
+  }
+}
+
+bool BlenderSession::builtin_image_pixels(const string &builtin_name,
+                                          void *builtin_data,
+                                          int tile,
+                                          unsigned char *pixels,
+                                          const size_t pixels_size,
+                                          const bool associate_alpha,
+                                          const bool free_cache)
+{
+  if (!builtin_data) {
+    return false;
+  }
+
+  const int frame = builtin_image_frame(builtin_name);
+
+  PointerRNA ptr;
+  RNA_id_pointer_create((ID *)builtin_data, &ptr);
+  BL::Image b_image(ptr);
+
+  const int width = b_image.size()[0];
+  const int height = b_image.size()[1];
+  const int channels = b_image.channels();
+
+  unsigned char *image_pixels = image_get_pixels_for_frame(b_image, frame, tile);
+  const size_t num_pixels = ((size_t)width) * height;
+
+  if (image_pixels && num_pixels * channels == pixels_size) {
+    memcpy(pixels, image_pixels, pixels_size * sizeof(unsigned char));
+  }
+  else {
+    if (channels == 1) {
+      memset(pixels, 0, pixels_size * sizeof(unsigned char));
     }
     else {
-      return false;
+      const size_t num_pixels_safe = pixels_size / channels;
+      unsigned char *cp = pixels;
+      for (size_t i = 0; i < num_pixels_safe; i++, cp += channels) {
+        cp[0] = 255;
+        cp[1] = 0;
+        cp[2] = 255;
+        if (channels == 4) {
+          cp[3] = 255;
+        }
+      }
     }
   }
 
+  if (image_pixels) {
+    MEM_freeN(image_pixels);
+  }
+
+  /* Free image buffers to save memory during render. */
+  if (free_cache) {
+    b_image.buffers_free();
+  }
+
+  if (associate_alpha) {
+    /* Premultiply, byte images are always straight for Blender. */
+    unsigned char *cp = pixels;
+    for (size_t i = 0; i < num_pixels; i++, cp += channels) {
+      cp[0] = (cp[0] * cp[3]) >> 8;
+      cp[1] = (cp[1] * cp[3]) >> 8;
+      cp[2] = (cp[2] * cp[3]) >> 8;
+    }
+  }
   return true;
 }
 
-bool BlenderImageLoader::load_pixels(const ImageMetaData &metadata,
-                                     void *pixels,
-                                     const size_t pixels_size,
-                                     const bool associate_alpha)
+bool BlenderSession::builtin_image_float_pixels(const string &builtin_name,
+                                                void *builtin_data,
+                                                int tile,
+                                                float *pixels,
+                                                const size_t pixels_size,
+                                                const bool,
+                                                const bool free_cache)
 {
-  const size_t num_pixels = ((size_t)metadata.width) * metadata.height;
-  const int channels = metadata.channels;
-  const int tile = 0; /* TODO(lukas): Support tiles here? */
+  if (!builtin_data) {
+    return false;
+  }
 
-  if (b_image.is_float()) {
+  PointerRNA ptr;
+  RNA_id_pointer_create((ID *)builtin_data, &ptr);
+  BL::ID b_id(ptr);
+
+  if (b_id.is_a(&RNA_Image)) {
     /* image data */
+    BL::Image b_image(b_id);
+    int frame = builtin_image_frame(builtin_name);
+
+    const int width = b_image.size()[0];
+    const int height = b_image.size()[1];
+    const int channels = b_image.channels();
+
     float *image_pixels;
     image_pixels = image_get_float_pixels_for_frame(b_image, frame, tile);
+    const size_t num_pixels = ((size_t)width) * height;
 
     if (image_pixels && num_pixels * channels == pixels_size) {
       memcpy(pixels, image_pixels, pixels_size * sizeof(float));
@@ -89,7 +224,7 @@ bool BlenderImageLoader::load_pixels(const ImageMetaData &metadata,
       }
       else {
         const size_t num_pixels_safe = pixels_size / channels;
-        float *fp = (float *)pixels;
+        float *fp = pixels;
         for (int i = 0; i < num_pixels_safe; i++, fp += channels) {
           fp[0] = 1.0f;
           fp[1] = 0.0f;
@@ -104,91 +239,107 @@ bool BlenderImageLoader::load_pixels(const ImageMetaData &metadata,
     if (image_pixels) {
       MEM_freeN(image_pixels);
     }
-  }
-  else {
-    unsigned char *image_pixels = image_get_pixels_for_frame(b_image, frame, tile);
 
-    if (image_pixels && num_pixels * channels == pixels_size) {
-      memcpy(pixels, image_pixels, pixels_size * sizeof(unsigned char));
+    /* Free image buffers to save memory during render. */
+    if (free_cache) {
+      b_image.buffers_free();
+    }
+
+    return true;
+  }
+  else if (b_id.is_a(&RNA_Object)) {
+    /* smoke volume data */
+    BL::Object b_ob(b_id);
+    BL::FluidDomainSettings b_domain = object_fluid_gas_domain_find(b_ob);
+
+    if (!b_domain) {
+      return false;
+    }
+#ifdef WITH_FLUID
+    int3 resolution = get_int3(b_domain.domain_resolution());
+    int length, amplify = (b_domain.use_noise()) ? b_domain.noise_scale() : 1;
+
+    /* Velocity and heat data is always low-resolution. */
+    if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_VELOCITY) ||
+        builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_HEAT)) {
+      amplify = 1;
+    }
+
+    const int width = resolution.x * amplify;
+    const int height = resolution.y * amplify;
+    const int depth = resolution.z * amplify;
+    const size_t num_pixels = ((size_t)width) * height * depth;
+
+    if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_DENSITY)) {
+      FluidDomainSettings_density_grid_get_length(&b_domain.ptr, &length);
+      if (length == num_pixels) {
+        FluidDomainSettings_density_grid_get(&b_domain.ptr, pixels);
+        return true;
+      }
+    }
+    else if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_FLAME)) {
+      /* this is in range 0..1, and interpreted by the OpenGL smoke viewer
+       * as 1500..3000 K with the first part faded to zero density */
+      FluidDomainSettings_flame_grid_get_length(&b_domain.ptr, &length);
+      if (length == num_pixels) {
+        FluidDomainSettings_flame_grid_get(&b_domain.ptr, pixels);
+        return true;
+      }
+    }
+    else if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_COLOR)) {
+      /* the RGB is "premultiplied" by density for better interpolation results */
+      FluidDomainSettings_color_grid_get_length(&b_domain.ptr, &length);
+      if (length == num_pixels * 4) {
+        FluidDomainSettings_color_grid_get(&b_domain.ptr, pixels);
+        return true;
+      }
+    }
+    else if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_VELOCITY)) {
+      FluidDomainSettings_velocity_grid_get_length(&b_domain.ptr, &length);
+      if (length == num_pixels * 3) {
+        FluidDomainSettings_velocity_grid_get(&b_domain.ptr, pixels);
+        return true;
+      }
+    }
+    else if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_HEAT)) {
+      FluidDomainSettings_heat_grid_get_length(&b_domain.ptr, &length);
+      if (length == num_pixels) {
+        FluidDomainSettings_heat_grid_get(&b_domain.ptr, pixels);
+        return true;
+      }
+    }
+    else if (builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_TEMPERATURE)) {
+      FluidDomainSettings_temperature_grid_get_length(&b_domain.ptr, &length);
+      if (length == num_pixels) {
+        FluidDomainSettings_temperature_grid_get(&b_domain.ptr, pixels);
+        return true;
+      }
     }
     else {
-      if (channels == 1) {
-        memset(pixels, 0, pixels_size * sizeof(unsigned char));
-      }
-      else {
-        const size_t num_pixels_safe = pixels_size / channels;
-        unsigned char *cp = (unsigned char *)pixels;
-        for (size_t i = 0; i < num_pixels_safe; i++, cp += channels) {
-          cp[0] = 255;
-          cp[1] = 0;
-          cp[2] = 255;
-          if (channels == 4) {
-            cp[3] = 255;
-          }
-        }
-      }
+      fprintf(
+          stderr, "Cycles error: unknown volume attribute %s, skipping\n", builtin_name.c_str());
+      pixels[0] = 0.0f;
+      return false;
     }
-
-    if (image_pixels) {
-      MEM_freeN(image_pixels);
-    }
-
-    if (associate_alpha) {
-      /* Premultiply, byte images are always straight for Blender. */
-      unsigned char *cp = (unsigned char *)pixels;
-      for (size_t i = 0; i < num_pixels; i++, cp += channels) {
-        cp[0] = (cp[0] * cp[3]) >> 8;
-        cp[1] = (cp[1] * cp[3]) >> 8;
-        cp[2] = (cp[2] * cp[3]) >> 8;
-      }
+#endif
+    fprintf(stderr, "Cycles error: unexpected smoke volume resolution, skipping\n");
+  }
+  else {
+    /* We originally were passing view_layer here but in reality we need a
+     * a depsgraph to pass to the RE_point_density_minmax() function.
+     */
+    /* TODO(sergey): Check we're indeed in shader node tree. */
+    PointerRNA ptr;
+    RNA_pointer_create(NULL, &RNA_Node, builtin_data, &ptr);
+    BL::Node b_node(ptr);
+    if (b_node.is_a(&RNA_ShaderNodeTexPointDensity)) {
+      BL::ShaderNodeTexPointDensity b_point_density_node(b_node);
+      int length;
+      b_point_density_node.calc_point_density(b_depsgraph, &length, &pixels);
     }
   }
 
-  /* Free image buffers to save memory during render. */
-  if (free_cache) {
-    b_image.buffers_free();
-  }
-
-  return true;
-}
-
-string BlenderImageLoader::name() const
-{
-  return BL::Image(b_image).name();
-}
-
-bool BlenderImageLoader::equals(const ImageLoader &other) const
-{
-  const BlenderImageLoader &other_loader = (const BlenderImageLoader &)other;
-  return b_image == other_loader.b_image && frame == other_loader.frame;
-}
-
-/* Point Density */
-
-BlenderPointDensityLoader::BlenderPointDensityLoader(BL::Depsgraph b_depsgraph,
-                                                     BL::ShaderNodeTexPointDensity b_node)
-    : b_depsgraph(b_depsgraph), b_node(b_node)
-{
-}
-
-bool BlenderPointDensityLoader::load_metadata(ImageMetaData &metadata)
-{
-  metadata.channels = 4;
-  metadata.width = b_node.resolution();
-  metadata.height = metadata.width;
-  metadata.depth = metadata.width;
-  metadata.type = IMAGE_DATA_TYPE_FLOAT4;
-  return true;
-}
-
-bool BlenderPointDensityLoader::load_pixels(const ImageMetaData &,
-                                            void *pixels,
-                                            const size_t,
-                                            const bool)
-{
-  int length;
-  b_node.calc_point_density(b_depsgraph, &length, (float **)&pixels);
-  return true;
+  return false;
 }
 
 void BlenderSession::builtin_images_load()
@@ -206,15 +357,4 @@ void BlenderSession::builtin_images_load()
   manager->device_load_builtin(device, session->scene, session->progress);
 }
 
-string BlenderPointDensityLoader::name() const
-{
-  return BL::ShaderNodeTexPointDensity(b_node).name();
-}
-
-bool BlenderPointDensityLoader::equals(const ImageLoader &other) const
-{
-  const BlenderPointDensityLoader &other_loader = (const BlenderPointDensityLoader &)other;
-  return b_node == other_loader.b_node && b_depsgraph == other_loader.b_depsgraph;
-}
-
 CCL_NAMESPACE_END

intern/cycles/blender/blender_image.h

@@ -1,61 +0,0 @@
-/*
- * Copyright 2011-2020 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#ifndef __BLENDER_IMAGE_H__
-#define __BLENDER_IMAGE_H__
-
-#include "RNA_blender_cpp.h"
-
-#include "render/image.h"
-
-CCL_NAMESPACE_BEGIN
-
-class BlenderImageLoader : public ImageLoader {
- public:
-  BlenderImageLoader(BL::Image b_image, int frame);
-
-  bool load_metadata(ImageMetaData &metadata) override;
-  bool load_pixels(const ImageMetaData &metadata,
-                   void *pixels,
-                   const size_t pixels_size,
-                   const bool associate_alpha) override;
-  string name() const override;
-  bool equals(const ImageLoader &other) const override;
-
-  BL::Image b_image;
-  int frame;
-  bool free_cache;
-};
-
-class BlenderPointDensityLoader : public ImageLoader {
- public:
-  BlenderPointDensityLoader(BL::Depsgraph depsgraph, BL::ShaderNodeTexPointDensity b_node);
-
-  bool load_metadata(ImageMetaData &metadata) override;
-  bool load_pixels(const ImageMetaData &metadata,
-                   void *pixels,
-                   const size_t pixels_size,
-                   const bool associate_alpha) override;
-  string name() const override;
-  bool equals(const ImageLoader &other) const override;
-
-  BL::Depsgraph b_depsgraph;
-  BL::ShaderNodeTexPointDensity b_node;
-};
-
-CCL_NAMESPACE_END
-
-#endif /* __BLENDER_IMAGE_H__ */

intern/cycles/blender/blender_mesh.cpp

@@ -14,25 +14,25 @@
  * limitations under the License.
  */
 
-#include "render/camera.h"
 #include "render/colorspace.h"
 #include "render/mesh.h"
 #include "render/object.h"
 #include "render/scene.h"
+#include "render/camera.h"
 
-#include "blender/blender_session.h"
 #include "blender/blender_sync.h"
+#include "blender/blender_session.h"
 #include "blender/blender_util.h"
 
 #include "subd/subd_patch.h"
 #include "subd/subd_split.h"
 
 #include "util/util_algorithm.h"
-#include "util/util_disjoint_set.h"
 #include "util/util_foreach.h"
 #include "util/util_hash.h"
 #include "util/util_logging.h"
 #include "util/util_math.h"
+#include "util/util_disjoint_set.h"
 
 #include "mikktspace.h"
 

intern/cycles/blender/blender_object.cpp

@@ -15,14 +15,14 @@
  */
 
 #include "render/camera.h"
-#include "render/graph.h"
 #include "render/integrator.h"
+#include "render/graph.h"
 #include "render/light.h"
 #include "render/mesh.h"
-#include "render/nodes.h"
 #include "render/object.h"
-#include "render/particles.h"
 #include "render/scene.h"
+#include "render/nodes.h"
+#include "render/particles.h"
 #include "render/shader.h"
 
 #include "blender/blender_object_cull.h"
@@ -67,20 +67,10 @@ bool BlenderSync::object_is_mesh(BL::Object &b_ob)
     return false;
   }
 
-  BL::Object::type_enum type = b_ob.type();
-
-#ifdef WITH_NEW_OBJECT_TYPES
-  if (type == BL::Object::type_VOLUME || type == BL::Object::type_HAIR) {
-#else
-  if (type == BL::Object::type_VOLUME) {
-#endif
-    /* Will be exported attached to mesh. */
-    return true;
-  }
-  else if (type == BL::Object::type_CURVE) {
+  if (b_ob.type() == BL::Object::type_CURVE) {
     /* Skip exporting curves without faces, overhead can be
      * significant if there are many for path animation. */
-    BL::Curve b_curve(b_ob_data);
+    BL::Curve b_curve(b_ob.data());
 
     return (b_curve.bevel_object() || b_curve.extrude() != 0.0f || b_curve.bevel_depth() != 0.0f ||
             b_curve.dimensions() == BL::Curve::dimensions_2D || b_ob.modifiers.length());

intern/cycles/blender/blender_python.cpp

@@ -19,8 +19,8 @@
 #include "blender/CCL_api.h"
 
 #include "blender/blender_device.h"
-#include "blender/blender_session.h"
 #include "blender/blender_sync.h"
+#include "blender/blender_session.h"
 #include "blender/blender_util.h"
 
 #include "render/denoising.h"
@@ -38,8 +38,8 @@
 #ifdef WITH_OSL
 #  include "render/osl.h"
 
-#  include <OSL/oslconfig.h>
 #  include <OSL/oslquery.h>
+#  include <OSL/oslconfig.h>
 #endif
 
 #ifdef WITH_OPENCL

intern/cycles/blender/blender_session.cpp

@@ -41,8 +41,8 @@
 #include "util/util_progress.h"
 #include "util/util_time.h"
 
-#include "blender/blender_session.h"
 #include "blender/blender_sync.h"
+#include "blender/blender_session.h"
 #include "blender/blender_util.h"
 
 CCL_NAMESPACE_BEGIN
@@ -138,6 +138,14 @@ void BlenderSession::create_session()
   scene = new Scene(scene_params, session->device);
   scene->name = b_scene.name();
 
+  /* setup callbacks for builtin image support */
+  scene->image_manager->builtin_image_info_cb = function_bind(
+      &BlenderSession::builtin_image_info, this, _1, _2, _3);
+  scene->image_manager->builtin_image_pixels_cb = function_bind(
+      &BlenderSession::builtin_image_pixels, this, _1, _2, _3, _4, _5, _6, _7);
+  scene->image_manager->builtin_image_float_pixels_cb = function_bind(
+      &BlenderSession::builtin_image_float_pixels, this, _1, _2, _3, _4, _5, _6, _7);
+
   session->scene = scene;
 
   /* There is no single depsgraph to use for the entire render.
@@ -462,8 +470,7 @@ void BlenderSession::render(BL::Depsgraph &b_depsgraph_)
   b_rlay_name = b_view_layer.name();
 
   /* add passes */
-  vector<Pass> passes = sync->sync_render_passes(
-      b_rlay, b_view_layer, session_params.adaptive_sampling);
+  vector<Pass> passes = sync->sync_render_passes(b_rlay, b_view_layer);
   buffer_params.passes = passes;
 
   PointerRNA crl = RNA_pointer_get(&b_view_layer.ptr, "cycles");

intern/cycles/blender/blender_session.h

@@ -17,19 +17,15 @@
 #ifndef __BLENDER_SESSION_H__
 #define __BLENDER_SESSION_H__
 
-#include "RNA_blender_cpp.h"
-
 #include "device/device.h"
-
-#include "render/bake.h"
 #include "render/scene.h"
 #include "render/session.h"
+#include "render/bake.h"
 
 #include "util/util_vector.h"
 
 CCL_NAMESPACE_BEGIN
 
-class BlenderSync;
 class ImageMetaData;
 class Scene;
 class Session;
@@ -157,6 +153,22 @@ class BlenderSession {
                                      bool do_update_only);
   void do_write_update_render_tile(RenderTile &rtile, bool do_update_only, bool highlight);
 
+  int builtin_image_frame(const string &builtin_name);
+  void builtin_image_info(const string &builtin_name, void *builtin_data, ImageMetaData &metadata);
+  bool builtin_image_pixels(const string &builtin_name,
+                            void *builtin_data,
+                            int tile,
+                            unsigned char *pixels,
+                            const size_t pixels_size,
+                            const bool associate_alpha,
+                            const bool free_cache);
+  bool builtin_image_float_pixels(const string &builtin_name,
+                                  void *builtin_data,
+                                  int tile,
+                                  float *pixels,
+                                  const size_t pixels_size,
+                                  const bool associate_alpha,
+                                  const bool free_cache);
   void builtin_images_load();
 
   /* Update tile manager to reflect resumable render settings. */

intern/cycles/blender/blender_shader.cpp

@@ -23,15 +23,14 @@
 #include "render/scene.h"
 #include "render/shader.h"
 
-#include "blender/blender_image.h"
-#include "blender/blender_sync.h"
 #include "blender/blender_texture.h"
+#include "blender/blender_sync.h"
 #include "blender/blender_util.h"
 
 #include "util/util_debug.h"
 #include "util/util_foreach.h"
-#include "util/util_set.h"
 #include "util/util_string.h"
+#include "util/util_set.h"
 #include "util/util_task.h"
 
 CCL_NAMESPACE_BEGIN
@@ -328,7 +327,6 @@ static ShaderNode *add_node(Scene *scene,
     BL::ShaderNodeVectorRotate b_vector_rotate_node(b_node);
     VectorRotateNode *vector_rotate_node = new VectorRotateNode();
     vector_rotate_node->type = (NodeVectorRotateType)b_vector_rotate_node.rotation_type();
-    vector_rotate_node->invert = b_vector_rotate_node.invert();
     node = vector_rotate_node;
   }
   else if (b_node.is_a(&RNA_ShaderNodeVectorTransform)) {
@@ -620,16 +618,16 @@ static ShaderNode *add_node(Scene *scene,
       /* create script node */
       BL::ShaderNodeScript b_script_node(b_node);
 
-      ShaderManager *manager = scene->shader_manager;
+      OSLShaderManager *manager = (OSLShaderManager *)scene->shader_manager;
       string bytecode_hash = b_script_node.bytecode_hash();
 
       if (!bytecode_hash.empty()) {
-        node = OSLShaderManager::osl_node(manager, "", bytecode_hash, b_script_node.bytecode());
+        node = manager->osl_node("", bytecode_hash, b_script_node.bytecode());
       }
       else {
         string absolute_filepath = blender_absolute_path(
             b_data, b_ntree, b_script_node.filepath());
-        node = OSLShaderManager::osl_node(manager, absolute_filepath, "");
+        node = manager->osl_node(absolute_filepath, "");
       }
     }
 #else
@@ -651,18 +649,6 @@ static ShaderNode *add_node(Scene *scene,
     get_tex_mapping(&image->tex_mapping, b_texture_mapping);
 
     if (b_image) {
-      PointerRNA colorspace_ptr = b_image.colorspace_settings().ptr;
-      image->colorspace = get_enum_identifier(colorspace_ptr, "name");
-
-      image->animated = b_image_node.image_user().use_auto_refresh();
-      image->alpha_type = get_image_alpha_type(b_image);
-
-      image->tiles.clear();
-      BL::Image::tiles_iterator b_iter;
-      for (b_image.tiles.begin(b_iter); b_iter != b_image.tiles.end(); ++b_iter) {
-        image->tiles.push_back(b_iter->number());
-      }
-
       /* builtin images will use callback-based reading because
        * they could only be loaded correct from blender side
        */
@@ -679,13 +665,34 @@ static ShaderNode *add_node(Scene *scene,
          */
         int scene_frame = b_scene.frame_current();
         int image_frame = image_user_frame_number(b_image_user, scene_frame);
-        image->handle = scene->image_manager->add_image(
-            new BlenderImageLoader(b_image, image_frame), image->image_params());
+        image->filename = b_image.name() + "@" + string_printf("%d", image_frame);
+        image->builtin_data = b_image.ptr.data;
       }
       else {
         image->filename = image_user_file_path(
             b_image_user, b_image, b_scene.frame_current(), true);
+        image->builtin_data = NULL;
       }
+
+      PointerRNA colorspace_ptr = b_image.colorspace_settings().ptr;
+      image->colorspace = get_enum_identifier(colorspace_ptr, "name");
+
+      image->animated = b_image_node.image_user().use_auto_refresh();
+      image->alpha_type = get_image_alpha_type(b_image);
+
+      image->tiles.clear();
+      BL::Image::tiles_iterator b_iter;
+      for (b_image.tiles.begin(b_iter); b_iter != b_image.tiles.end(); ++b_iter) {
+        image->tiles.push_back(b_iter->number());
+      }
+
+      /* TODO: restore */
+      /* TODO(sergey): Does not work properly when we change builtin type. */
+#if 0
+      if (b_image.is_updated()) {
+        scene->image_manager->tag_reload_image(image->image_key());
+      }
+#endif
     }
     node = image;
   }
@@ -701,12 +708,6 @@ static ShaderNode *add_node(Scene *scene,
     get_tex_mapping(&env->tex_mapping, b_texture_mapping);
 
     if (b_image) {
-      PointerRNA colorspace_ptr = b_image.colorspace_settings().ptr;
-      env->colorspace = get_enum_identifier(colorspace_ptr, "name");
-
-      env->animated = b_env_node.image_user().use_auto_refresh();
-      env->alpha_type = get_image_alpha_type(b_image);
-
       bool is_builtin = b_image.packed_file() || b_image.source() == BL::Image::source_GENERATED ||
                         b_image.source() == BL::Image::source_MOVIE ||
                         (b_engine.is_preview() && b_image.source() != BL::Image::source_SEQUENCE);
@@ -714,13 +715,28 @@ static ShaderNode *add_node(Scene *scene,
       if (is_builtin) {
         int scene_frame = b_scene.frame_current();
         int image_frame = image_user_frame_number(b_image_user, scene_frame);
-        env->handle = scene->image_manager->add_image(new BlenderImageLoader(b_image, image_frame),
-                                                      env->image_params());
+        env->filename = b_image.name() + "@" + string_printf("%d", image_frame);
+        env->builtin_data = b_image.ptr.data;
       }
       else {
         env->filename = image_user_file_path(
             b_image_user, b_image, b_scene.frame_current(), false);
+        env->builtin_data = NULL;
       }
+
+      PointerRNA colorspace_ptr = b_image.colorspace_settings().ptr;
+      env->colorspace = get_enum_identifier(colorspace_ptr, "name");
+
+      env->animated = b_env_node.image_user().use_auto_refresh();
+      env->alpha_type = get_image_alpha_type(b_image);
+
+      /* TODO: restore */
+      /* TODO(sergey): Does not work properly when we change builtin type. */
+#if 0
+      if (b_image.is_updated()) {
+        scene->image_manager->tag_reload_image(env->image_key());
+      }
+#endif
     }
     node = env;
   }
@@ -864,13 +880,18 @@ static ShaderNode *add_node(Scene *scene,
   else if (b_node.is_a(&RNA_ShaderNodeTexPointDensity)) {
     BL::ShaderNodeTexPointDensity b_point_density_node(b_node);
     PointDensityTextureNode *point_density = new PointDensityTextureNode();
+    point_density->filename = b_point_density_node.name();
     point_density->space = (NodeTexVoxelSpace)b_point_density_node.space();
     point_density->interpolation = get_image_interpolation(b_point_density_node);
-    point_density->handle = scene->image_manager->add_image(
-        new BlenderPointDensityLoader(b_depsgraph, b_point_density_node),
-        point_density->image_params());
+    point_density->builtin_data = b_point_density_node.ptr.data;
+    point_density->image_manager = scene->image_manager;
 
-    b_point_density_node.cache_point_density(b_depsgraph);
+    /* TODO(sergey): Use more proper update flag. */
+    if (true) {
+      point_density->add_image();
+      b_point_density_node.cache_point_density(b_depsgraph);
+      scene->image_manager->tag_reload_image(point_density->image_key());
+    }
     node = point_density;
 
     /* Transformation form world space to texture space.
@@ -1260,7 +1281,6 @@ void BlenderSync::sync_materials(BL::Depsgraph &b_depsgraph, bool update_all)
       shader->heterogeneous_volume = !get_boolean(cmat, "homogeneous_volume");
       shader->volume_sampling_method = get_volume_sampling(cmat);
       shader->volume_interpolation_method = get_volume_interpolation(cmat);
-      shader->volume_step_rate = get_float(cmat, "volume_step_rate");
       shader->displacement_method = get_displacement_method(cmat);
 
       shader->set_graph(graph);
@@ -1325,7 +1345,6 @@ void BlenderSync::sync_world(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d,
       shader->heterogeneous_volume = !get_boolean(cworld, "homogeneous_volume");
       shader->volume_sampling_method = get_volume_sampling(cworld);
       shader->volume_interpolation_method = get_volume_interpolation(cworld);
-      shader->volume_step_rate = get_float(cworld, "volume_step_size");
     }
     else if (new_viewport_parameters.use_scene_world && b_world) {
       BackgroundNode *background = new BackgroundNode();

intern/cycles/blender/blender_sync.cpp

@@ -16,7 +16,6 @@
 
 #include "render/background.h"
 #include "render/camera.h"
-#include "render/curves.h"
 #include "render/film.h"
 #include "render/graph.h"
 #include "render/integrator.h"
@@ -26,18 +25,19 @@
 #include "render/object.h"
 #include "render/scene.h"
 #include "render/shader.h"
+#include "render/curves.h"
 
 #include "device/device.h"
 
 #include "blender/blender_device.h"
-#include "blender/blender_session.h"
 #include "blender/blender_sync.h"
+#include "blender/blender_session.h"
 #include "blender/blender_util.h"
 
 #include "util/util_debug.h"
 #include "util/util_foreach.h"
-#include "util/util_hash.h"
 #include "util/util_opengl.h"
+#include "util/util_hash.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -178,11 +178,6 @@ void BlenderSync::sync_recalc(BL::Depsgraph &b_depsgraph, BL::SpaceView3D &b_v3d
         world_recalc = true;
       }
     }
-    /* Volume */
-    else if (b_id.is_a(&RNA_Volume)) {
-      BL::Volume b_volume(b_id);
-      geometry_map.set_recalc(b_volume);
-    }
   }
 
   BlenderViewportParameters new_viewport_parameters(b_v3d);
@@ -262,8 +257,7 @@ void BlenderSync::sync_integrator()
   integrator->transparent_max_bounce = get_int(cscene, "transparent_max_bounces");
 
   integrator->volume_max_steps = get_int(cscene, "volume_max_steps");
-  integrator->volume_step_rate = (preview) ? get_float(cscene, "volume_preview_step_rate") :
-                                             get_float(cscene, "volume_step_rate");
+  integrator->volume_step_size = get_float(cscene, "volume_step_size");
 
   integrator->caustics_reflective = get_boolean(cscene, "caustics_reflective");
   integrator->caustics_refractive = get_boolean(cscene, "caustics_refractive");
@@ -302,16 +296,6 @@ void BlenderSync::sync_integrator()
   integrator->sample_all_lights_indirect = get_boolean(cscene, "sample_all_lights_indirect");
   integrator->light_sampling_threshold = get_float(cscene, "light_sampling_threshold");
 
-  if (RNA_boolean_get(&cscene, "use_adaptive_sampling")) {
-    integrator->sampling_pattern = SAMPLING_PATTERN_PMJ;
-    integrator->adaptive_min_samples = get_int(cscene, "adaptive_min_samples");
-    integrator->adaptive_threshold = get_float(cscene, "adaptive_threshold");
-  }
-  else {
-    integrator->adaptive_min_samples = INT_MAX;
-    integrator->adaptive_threshold = 0.0f;
-  }
-
   int diffuse_samples = get_int(cscene, "diffuse_samples");
   int glossy_samples = get_int(cscene, "glossy_samples");
   int transmission_samples = get_int(cscene, "transmission_samples");
@@ -328,8 +312,6 @@ void BlenderSync::sync_integrator()
     integrator->mesh_light_samples = mesh_light_samples * mesh_light_samples;
     integrator->subsurface_samples = subsurface_samples * subsurface_samples;
     integrator->volume_samples = volume_samples * volume_samples;
-    integrator->adaptive_min_samples = min(
-        integrator->adaptive_min_samples * integrator->adaptive_min_samples, INT_MAX);
   }
   else {
     integrator->diffuse_samples = diffuse_samples;
@@ -411,7 +393,6 @@ void BlenderSync::sync_view_layer(BL::SpaceView3D & /*b_v3d*/, BL::ViewLayer &b_
   view_layer.use_background_ao = b_view_layer.use_ao();
   view_layer.use_surfaces = b_view_layer.use_solid();
   view_layer.use_hair = b_view_layer.use_strand();
-  view_layer.use_volumes = b_view_layer.use_volumes();
 
   /* Material override. */
   view_layer.material_override = b_view_layer.material_override();
@@ -503,8 +484,6 @@ PassType BlenderSync::get_pass_type(BL::RenderPass &b_pass)
   MAP_PASS("Debug Ray Bounces", PASS_RAY_BOUNCES);
 #endif
   MAP_PASS("Debug Render Time", PASS_RENDER_TIME);
-  MAP_PASS("AdaptiveAuxBuffer", PASS_ADAPTIVE_AUX_BUFFER);
-  MAP_PASS("Debug Sample Count", PASS_SAMPLE_COUNT);
   if (string_startswith(name, cryptomatte_prefix)) {
     return PASS_CRYPTOMATTE;
   }
@@ -540,9 +519,7 @@ int BlenderSync::get_denoising_pass(BL::RenderPass &b_pass)
   return -1;
 }
 
-vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer &b_rlay,
-                                             BL::ViewLayer &b_view_layer,
-                                             bool adaptive_sampling)
+vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer &b_rlay, BL::ViewLayer &b_view_layer)
 {
   vector<Pass> passes;
 
@@ -618,10 +595,6 @@ vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer &b_rlay,
     b_engine.add_pass("Debug Render Time", 1, "X", b_view_layer.name().c_str());
     Pass::add(PASS_RENDER_TIME, passes, "Debug Render Time");
   }
-  if (get_boolean(crp, "pass_debug_sample_count")) {
-    b_engine.add_pass("Debug Sample Count", 1, "X", b_view_layer.name().c_str());
-    Pass::add(PASS_SAMPLE_COUNT, passes, "Debug Sample Count");
-  }
   if (get_boolean(crp, "use_pass_volume_direct")) {
     b_engine.add_pass("VolumeDir", 3, "RGB", b_view_layer.name().c_str());
     Pass::add(PASS_VOLUME_DIRECT, passes, "VolumeDir");
@@ -668,13 +641,6 @@ vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer &b_rlay,
                                                         CRYPT_ACCURATE);
   }
 
-  if (adaptive_sampling) {
-    Pass::add(PASS_ADAPTIVE_AUX_BUFFER, passes);
-    if (!get_boolean(crp, "pass_debug_sample_count")) {
-      Pass::add(PASS_SAMPLE_COUNT, passes);
-    }
-  }
-
   RNA_BEGIN (&crp, b_aov, "aovs") {
     bool is_color = (get_enum(b_aov, "type") == 1);
     string name = get_string(b_aov, "name");
@@ -914,8 +880,6 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine &b_engine,
   params.use_profiling = params.device.has_profiling && !b_engine.is_preview() && background &&
                          BlenderSession::print_render_stats;
 
-  params.adaptive_sampling = RNA_boolean_get(&cscene, "use_adaptive_sampling");
-
   return params;
 }
 

intern/cycles/blender/blender_sync.h

@@ -18,9 +18,9 @@
 #define __BLENDER_SYNC_H__
 
 #include "MEM_guardedalloc.h"
+#include "RNA_types.h"
 #include "RNA_access.h"
 #include "RNA_blender_cpp.h"
-#include "RNA_types.h"
 
 #include "blender/blender_id_map.h"
 #include "blender/blender_viewport.h"
@@ -71,9 +71,7 @@ class BlenderSync {
                  int height,
                  void **python_thread_state);
   void sync_view_layer(BL::SpaceView3D &b_v3d, BL::ViewLayer &b_view_layer);
-  vector<Pass> sync_render_passes(BL::RenderLayer &b_render_layer,
-                                  BL::ViewLayer &b_view_layer,
-                                  bool adaptive_sampling);
+  vector<Pass> sync_render_passes(BL::RenderLayer &b_render_layer, BL::ViewLayer &b_view_layer);
   void sync_integrator();
   void sync_camera(BL::RenderSettings &b_render,
                    BL::Object &b_override,
@@ -157,7 +155,6 @@ class BlenderSync {
                         BL::Object b_ob,
                         Geometry *geom,
                         int motion_step);
-  void sync_hair(Hair *hair, BL::Object &b_ob, bool motion, int motion_step = 0);
   void sync_particle_hair(
       Geometry *geom, BL::Mesh &b_mesh, BL::Object &b_ob, bool motion, int motion_step = 0);
   void sync_curve_settings();
@@ -237,7 +234,6 @@ class BlenderSync {
           use_background_ao(true),
           use_surfaces(true),
           use_hair(true),
-          use_volumes(true),
           samples(0),
           bound_samples(false)
     {
@@ -249,7 +245,6 @@ class BlenderSync {
     bool use_background_ao;
     bool use_surfaces;
     bool use_hair;
-    bool use_volumes;
     int samples;
     bool bound_samples;
   } view_layer;

intern/cycles/blender/blender_texture.h

@@ -17,8 +17,8 @@
 #ifndef __BLENDER_TEXTURE_H__
 #define __BLENDER_TEXTURE_H__
 
-#include "blender/blender_sync.h"
 #include <stdlib.h>
+#include "blender/blender_sync.h"
 
 CCL_NAMESPACE_BEGIN
 

intern/cycles/blender/blender_viewport.h

@@ -18,9 +18,9 @@
 #define __BLENDER_VIEWPORT_H__
 
 #include "MEM_guardedalloc.h"
+#include "RNA_types.h"
 #include "RNA_access.h"
 #include "RNA_blender_cpp.h"
-#include "RNA_types.h"
 
 #include "render/film.h"
 #include "util/util_param.h"

intern/cycles/blender/blender_volume.cpp

@@ -1,3 +1,4 @@
+
 /*
  * Copyright 2011-2013 Blender Foundation
  *
@@ -15,177 +16,14 @@
  */
 
 #include "render/colorspace.h"
-#include "render/image.h"
-#include "render/image_vdb.h"
 #include "render/mesh.h"
 #include "render/object.h"
 
 #include "blender/blender_sync.h"
 #include "blender/blender_util.h"
 
-#ifdef WITH_OPENVDB
-#  include <openvdb/openvdb.h>
-openvdb::GridBase::ConstPtr BKE_volume_grid_openvdb_for_read(const struct Volume *volume,
-                                                             struct VolumeGrid *grid);
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-/* TODO: verify this is not loading unnecessary attributes. */
-class BlenderSmokeLoader : public ImageLoader {
- public:
-  BlenderSmokeLoader(const BL::Object &b_ob, AttributeStandard attribute)
-      : b_ob(b_ob), attribute(attribute)
-  {
-  }
-
-  bool load_metadata(ImageMetaData &metadata) override
-  {
-    BL::FluidDomainSettings b_domain = object_fluid_gas_domain_find(b_ob);
-
-    if (!b_domain) {
-      return false;
-    }
-
-    if (attribute == ATTR_STD_VOLUME_DENSITY || attribute == ATTR_STD_VOLUME_FLAME ||
-        attribute == ATTR_STD_VOLUME_HEAT || attribute == ATTR_STD_VOLUME_TEMPERATURE) {
-      metadata.type = IMAGE_DATA_TYPE_FLOAT;
-      metadata.channels = 1;
-    }
-    else if (attribute == ATTR_STD_VOLUME_COLOR) {
-      metadata.type = IMAGE_DATA_TYPE_FLOAT4;
-      metadata.channels = 4;
-    }
-    else if (attribute == ATTR_STD_VOLUME_VELOCITY) {
-      metadata.type = IMAGE_DATA_TYPE_FLOAT4;
-      metadata.channels = 3;
-    }
-    else {
-      return false;
-    }
-
-    int3 resolution = get_int3(b_domain.domain_resolution());
-    int amplify = (b_domain.use_noise()) ? b_domain.noise_scale() : 1;
-
-    /* Velocity and heat data is always low-resolution. */
-    if (attribute == ATTR_STD_VOLUME_VELOCITY || attribute == ATTR_STD_VOLUME_HEAT) {
-      amplify = 1;
-    }
-
-    metadata.width = resolution.x * amplify;
-    metadata.height = resolution.y * amplify;
-    metadata.depth = resolution.z * amplify;
-
-    /* Create a matrix to transform from object space to mesh texture space.
-     * This does not work with deformations but that can probably only be done
-     * well with a volume grid mapping of coordinates. */
-    BL::Mesh b_mesh(b_ob.data());
-    float3 loc, size;
-    mesh_texture_space(b_mesh, loc, size);
-    metadata.transform_3d = transform_translate(-loc) * transform_scale(size);
-    metadata.use_transform_3d = true;
-
-    return true;
-  }
-
-  bool load_pixels(const ImageMetaData &, void *pixels, const size_t, const bool) override
-  {
-    /* smoke volume data */
-    BL::FluidDomainSettings b_domain = object_fluid_gas_domain_find(b_ob);
-
-    if (!b_domain) {
-      return false;
-    }
-#ifdef WITH_FLUID
-    int3 resolution = get_int3(b_domain.domain_resolution());
-    int length, amplify = (b_domain.use_noise()) ? b_domain.noise_scale() : 1;
-
-    /* Velocity and heat data is always low-resolution. */
-    if (attribute == ATTR_STD_VOLUME_VELOCITY || attribute == ATTR_STD_VOLUME_HEAT) {
-      amplify = 1;
-    }
-
-    const int width = resolution.x * amplify;
-    const int height = resolution.y * amplify;
-    const int depth = resolution.z * amplify;
-    const size_t num_pixels = ((size_t)width) * height * depth;
-
-    float *fpixels = (float *)pixels;
-
-    if (attribute == ATTR_STD_VOLUME_DENSITY) {
-      FluidDomainSettings_density_grid_get_length(&b_domain.ptr, &length);
-      if (length == num_pixels) {
-        FluidDomainSettings_density_grid_get(&b_domain.ptr, fpixels);
-        return true;
-      }
-    }
-    else if (attribute == ATTR_STD_VOLUME_FLAME) {
-      /* this is in range 0..1, and interpreted by the OpenGL smoke viewer
-       * as 1500..3000 K with the first part faded to zero density */
-      FluidDomainSettings_flame_grid_get_length(&b_domain.ptr, &length);
-      if (length == num_pixels) {
-        FluidDomainSettings_flame_grid_get(&b_domain.ptr, fpixels);
-        return true;
-      }
-    }
-    else if (attribute == ATTR_STD_VOLUME_COLOR) {
-      /* the RGB is "premultiplied" by density for better interpolation results */
-      FluidDomainSettings_color_grid_get_length(&b_domain.ptr, &length);
-      if (length == num_pixels * 4) {
-        FluidDomainSettings_color_grid_get(&b_domain.ptr, fpixels);
-        return true;
-      }
-    }
-    else if (attribute == ATTR_STD_VOLUME_VELOCITY) {
-      FluidDomainSettings_velocity_grid_get_length(&b_domain.ptr, &length);
-      if (length == num_pixels * 3) {
-        FluidDomainSettings_velocity_grid_get(&b_domain.ptr, fpixels);
-        return true;
-      }
-    }
-    else if (attribute == ATTR_STD_VOLUME_HEAT) {
-      FluidDomainSettings_heat_grid_get_length(&b_domain.ptr, &length);
-      if (length == num_pixels) {
-        FluidDomainSettings_heat_grid_get(&b_domain.ptr, fpixels);
-        return true;
-      }
-    }
-    else if (attribute == ATTR_STD_VOLUME_TEMPERATURE) {
-      FluidDomainSettings_temperature_grid_get_length(&b_domain.ptr, &length);
-      if (length == num_pixels) {
-        FluidDomainSettings_temperature_grid_get(&b_domain.ptr, fpixels);
-        return true;
-      }
-    }
-    else {
-      fprintf(stderr,
-              "Cycles error: unknown volume attribute %s, skipping\n",
-              Attribute::standard_name(attribute));
-      fpixels[0] = 0.0f;
-      return false;
-    }
-#else
-    (void)pixels;
-#endif
-    fprintf(stderr, "Cycles error: unexpected smoke volume resolution, skipping\n");
-    return false;
-  }
-
-  string name() const override
-  {
-    return Attribute::standard_name(attribute);
-  }
-
-  bool equals(const ImageLoader &other) const override
-  {
-    const BlenderSmokeLoader &other_loader = (const BlenderSmokeLoader &)other;
-    return b_ob == other_loader.b_ob && attribute == other_loader.attribute;
-  }
-
-  BL::Object b_ob;
-  AttributeStandard attribute;
-};
-
 static void sync_smoke_volume(Scene *scene, BL::Object &b_ob, Mesh *mesh, float frame)
 {
   BL::FluidDomainSettings b_domain = object_fluid_gas_domain_find(b_ob);
@@ -193,6 +31,7 @@ static void sync_smoke_volume(Scene *scene, BL::Object &b_ob, Mesh *mesh, float
     return;
   }
 
+  ImageManager *image_manager = scene->image_manager;
   AttributeStandard attributes[] = {ATTR_STD_VOLUME_DENSITY,
                                     ATTR_STD_VOLUME_COLOR,
                                     ATTR_STD_VOLUME_FLAME,
@@ -207,172 +46,47 @@ static void sync_smoke_volume(Scene *scene, BL::Object &b_ob, Mesh *mesh, float
       continue;
     }
 
-    mesh->volume_clipping = b_domain.clipping();
+    mesh->volume_isovalue = b_domain.clipping();
 
     Attribute *attr = mesh->attributes.add(std);
+    VoxelAttribute *volume_data = attr->data_voxel();
+    ImageMetaData metadata;
 
-    ImageLoader *loader = new BlenderSmokeLoader(b_ob, std);
-    ImageParams params;
-    params.frame = frame;
+    ImageKey key;
+    key.filename = Attribute::standard_name(std);
+    key.builtin_data = b_ob.ptr.data;
 
-    attr->data_voxel() = scene->image_manager->add_image(loader, params);
-  }
-}
-
-class BlenderVolumeLoader : public VDBImageLoader {
- public:
-  BlenderVolumeLoader(BL::Volume b_volume, const string &grid_name)
-      : VDBImageLoader(grid_name),
-        b_volume(b_volume),
-        b_volume_grid(PointerRNA_NULL),
-        unload(false)
-  {
-#ifdef WITH_OPENVDB
-    /* Find grid with matching name. */
-    BL::Volume::grids_iterator b_grid_iter;
-    for (b_volume.grids.begin(b_grid_iter); b_grid_iter != b_volume.grids.end(); ++b_grid_iter) {
-      if (b_grid_iter->name() == grid_name) {
-        b_volume_grid = *b_grid_iter;
-      }
-    }
-#endif
+    volume_data->manager = image_manager;
+    volume_data->slot = image_manager->add_image(key, frame, metadata);
   }
 
-  bool load_metadata(ImageMetaData &metadata) override
-  {
-    if (!b_volume_grid) {
-      return false;
-    }
+  /* Create a matrix to transform from object space to mesh texture space.
+   * This does not work with deformations but that can probably only be done
+   * well with a volume grid mapping of coordinates. */
+  if (mesh->need_attribute(scene, ATTR_STD_GENERATED_TRANSFORM)) {
+    Attribute *attr = mesh->attributes.add(ATTR_STD_GENERATED_TRANSFORM);
+    Transform *tfm = attr->data_transform();
 
-    unload = !b_volume_grid.is_loaded();
+    BL::Mesh b_mesh(b_ob.data());
+    float3 loc, size;
+    mesh_texture_space(b_mesh, loc, size);
 
-#ifdef WITH_OPENVDB
-    Volume *volume = (Volume *)b_volume.ptr.data;
-    VolumeGrid *volume_grid = (VolumeGrid *)b_volume_grid.ptr.data;
-    grid = BKE_volume_grid_openvdb_for_read(volume, volume_grid);
-#endif
-
-    return VDBImageLoader::load_metadata(metadata);
+    *tfm = transform_translate(-loc) * transform_scale(size);
   }
-
-  bool load_pixels(const ImageMetaData &metadata,
-                   void *pixels,
-                   const size_t pixel_size,
-                   const bool associate_alpha) override
-  {
-    if (!b_volume_grid) {
-      return false;
-    }
-
-    return VDBImageLoader::load_pixels(metadata, pixels, pixel_size, associate_alpha);
-  }
-
-  bool equals(const ImageLoader &other) const override
-  {
-    /* TODO: detect multiple volume datablocks with the same filepath. */
-    const BlenderVolumeLoader &other_loader = (const BlenderVolumeLoader &)other;
-    return b_volume == other_loader.b_volume && b_volume_grid == other_loader.b_volume_grid;
-  }
-
-  void cleanup() override
-  {
-    VDBImageLoader::cleanup();
-    if (b_volume_grid && unload) {
-      b_volume_grid.unload();
-    }
-  }
-
-  BL::Volume b_volume;
-  BL::VolumeGrid b_volume_grid;
-  bool unload;
-};
-
-static void sync_volume_object(BL::BlendData &b_data, BL::Object &b_ob, Scene *scene, Mesh *mesh)
-{
-  BL::Volume b_volume(b_ob.data());
-  b_volume.grids.load(b_data.ptr.data);
-
-  BL::VolumeRender b_render(b_volume.render());
-
-  mesh->volume_clipping = b_render.clipping();
-  mesh->volume_step_size = b_render.step_size();
-  mesh->volume_object_space = (b_render.space() == BL::VolumeRender::space_OBJECT);
-
-  /* Find grid with matching name. */
-  BL::Volume::grids_iterator b_grid_iter;
-  for (b_volume.grids.begin(b_grid_iter); b_grid_iter != b_volume.grids.end(); ++b_grid_iter) {
-    BL::VolumeGrid b_grid = *b_grid_iter;
-    ustring name = ustring(b_grid.name());
-    AttributeStandard std = ATTR_STD_NONE;
-
-    if (name == Attribute::standard_name(ATTR_STD_VOLUME_DENSITY)) {
-      std = ATTR_STD_VOLUME_DENSITY;
-    }
-    else if (name == Attribute::standard_name(ATTR_STD_VOLUME_COLOR)) {
-      std = ATTR_STD_VOLUME_COLOR;
-    }
-    else if (name == Attribute::standard_name(ATTR_STD_VOLUME_FLAME)) {
-      std = ATTR_STD_VOLUME_FLAME;
-    }
-    else if (name == Attribute::standard_name(ATTR_STD_VOLUME_HEAT)) {
-      std = ATTR_STD_VOLUME_HEAT;
-    }
-    else if (name == Attribute::standard_name(ATTR_STD_VOLUME_TEMPERATURE)) {
-      std = ATTR_STD_VOLUME_TEMPERATURE;
-    }
-    else if (name == Attribute::standard_name(ATTR_STD_VOLUME_VELOCITY)) {
-      std = ATTR_STD_VOLUME_VELOCITY;
-    }
-
-    if ((std != ATTR_STD_NONE && mesh->need_attribute(scene, std)) ||
-        mesh->need_attribute(scene, name)) {
-      Attribute *attr = (std != ATTR_STD_NONE) ?
-                            mesh->attributes.add(std) :
-                            mesh->attributes.add(name, TypeDesc::TypeFloat, ATTR_ELEMENT_VOXEL);
-
-      ImageLoader *loader = new BlenderVolumeLoader(b_volume, name.string());
-      ImageParams params;
-      params.frame = b_volume.grids.frame();
-
-      attr->data_voxel() = scene->image_manager->add_image(loader, params);
-    }
-  }
-}
-
-/* If the voxel attributes change, we need to rebuild the bounding mesh. */
-static vector<int> get_voxel_image_slots(Mesh *mesh)
-{
-  vector<int> slots;
-  for (const Attribute &attr : mesh->attributes.attributes) {
-    if (attr.element == ATTR_ELEMENT_VOXEL) {
-      slots.push_back(attr.data_voxel().svm_slot());
-    }
-  }
-
-  return slots;
 }
 
 void BlenderSync::sync_volume(BL::Object &b_ob, Mesh *mesh, const vector<Shader *> &used_shaders)
 {
-  vector<int> old_voxel_slots = get_voxel_image_slots(mesh);
+  bool old_has_voxel_attributes = mesh->has_voxel_attributes();
 
   mesh->clear();
   mesh->used_shaders = used_shaders;
 
-  if (view_layer.use_volumes) {
-    if (b_ob.type() == BL::Object::type_VOLUME) {
-      /* Volume object. Create only attributes, bounding mesh will then
-       * be automatically generated later. */
-      sync_volume_object(b_data, b_ob, scene, mesh);
-    }
-    else {
-      /* Smoke domain. */
-      sync_smoke_volume(scene, b_ob, mesh, b_scene.frame_current());
-    }
-  }
+  /* Smoke domain. */
+  sync_smoke_volume(scene, b_ob, mesh, b_scene.frame_current());
 
   /* Tag update. */
-  bool rebuild = (old_voxel_slots != get_voxel_image_slots(mesh));
+  bool rebuild = (old_has_voxel_attributes != mesh->has_voxel_attributes());
   mesh->tag_update(scene, rebuild);
 }
 

intern/cycles/bvh/bvh.cpp

@@ -535,9 +535,8 @@ void BVH::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
 
         /* Modify offsets into arrays */
         int4 data = bvh_nodes[i + nsize_bbox];
-
+        int4 data1 = bvh_nodes[i + nsize_bbox - 1];
         if (use_obvh) {
-          int4 data1 = bvh_nodes[i + nsize_bbox - 1];
           data.z += (data.z < 0) ? -noffset_leaf : noffset;
           data.w += (data.w < 0) ? -noffset_leaf : noffset;
           data.x += (data.x < 0) ? -noffset_leaf : noffset;
@@ -546,8 +545,6 @@ void BVH::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
           data1.w += (data1.w < 0) ? -noffset_leaf : noffset;
           data1.x += (data1.x < 0) ? -noffset_leaf : noffset;
           data1.y += (data1.y < 0) ? -noffset_leaf : noffset;
-          pack_nodes[pack_nodes_offset + nsize_bbox] = data;
-          pack_nodes[pack_nodes_offset + nsize_bbox - 1] = data1;
         }
         else {
           data.z += (data.z < 0) ? -noffset_leaf : noffset;
@@ -556,7 +553,10 @@ void BVH::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
             data.x += (data.x < 0) ? -noffset_leaf : noffset;
             data.y += (data.y < 0) ? -noffset_leaf : noffset;
           }
-          pack_nodes[pack_nodes_offset + nsize_bbox] = data;
+        }
+        pack_nodes[pack_nodes_offset + nsize_bbox] = data;
+        if (use_obvh) {
+          pack_nodes[pack_nodes_offset + nsize_bbox - 1] = data1;
         }
 
         /* Usually this copies nothing, but we better

intern/cycles/bvh/bvh_build.cpp

@@ -22,20 +22,20 @@
 #include "bvh/bvh_params.h"
 #include "bvh_split.h"
 
-#include "render/curves.h"
 #include "render/hair.h"
 #include "render/mesh.h"
 #include "render/object.h"
 #include "render/scene.h"
+#include "render/curves.h"
 
 #include "util/util_algorithm.h"
 #include "util/util_foreach.h"
 #include "util/util_logging.h"
 #include "util/util_progress.h"
-#include "util/util_queue.h"
-#include "util/util_simd.h"
 #include "util/util_stack_allocator.h"
+#include "util/util_simd.h"
 #include "util/util_time.h"
+#include "util/util_queue.h"
 
 CCL_NAMESPACE_BEGIN
 

intern/cycles/bvh/bvh_embree.cpp

@@ -35,9 +35,9 @@
 
 #ifdef WITH_EMBREE
 
-#  include <embree3/rtcore_geometry.h>
 #  include <pmmintrin.h>
 #  include <xmmintrin.h>
+#  include <embree3/rtcore_geometry.h>
 
 #  include "bvh/bvh_embree.h"
 
@@ -45,9 +45,9 @@
  */
 #  include "kernel/bvh/bvh_embree.h"
 #  include "kernel/kernel_compat_cpu.h"
+#  include "kernel/split/kernel_split_data_types.h"
 #  include "kernel/kernel_globals.h"
 #  include "kernel/kernel_random.h"
-#  include "kernel/split/kernel_split_data_types.h"
 
 #  include "render/hair.h"
 #  include "render/mesh.h"

intern/cycles/bvh/bvh_optix.cpp

@@ -18,11 +18,10 @@
 #ifdef WITH_OPTIX
 
 #  include "bvh/bvh_optix.h"
-#  include "render/geometry.h"
 #  include "render/hair.h"
+#  include "render/geometry.h"
 #  include "render/mesh.h"
 #  include "render/object.h"
-#  include "util/util_foreach.h"
 #  include "util/util_logging.h"
 #  include "util/util_progress.h"
 

intern/cycles/device/cuda/device_cuda.h

@@ -82,17 +82,6 @@ class CUDADevice : public Device {
   device_vector<TextureInfo> texture_info;
   bool need_texture_info;
 
-  /* Kernels */
-  struct {
-    bool loaded;
-
-    CUfunction adaptive_stopping;
-    CUfunction adaptive_filter_x;
-    CUfunction adaptive_filter_y;
-    CUfunction adaptive_scale_samples;
-    int adaptive_num_threads_per_block;
-  } functions;
-
   static bool have_precompiled_kernels();
 
   virtual bool show_samples() const;
@@ -125,8 +114,6 @@ class CUDADevice : public Device {
 
   virtual bool load_kernels(const DeviceRequestedFeatures &requested_features);
 
-  void load_functions();
-
   void reserve_local_memory(const DeviceRequestedFeatures &requested_features);
 
   void init_host_memory();
@@ -155,13 +142,9 @@ class CUDADevice : public Device {
 
   virtual void const_copy_to(const char *name, void *host, size_t size);
 
-  void global_alloc(device_memory &mem);
+  void tex_alloc(device_memory &mem);
 
-  void global_free(device_memory &mem);
-
-  void tex_alloc(device_texture &mem);
-
-  void tex_free(device_texture &mem);
+  void tex_free(device_memory &mem);
 
   bool denoising_non_local_means(device_ptr image_ptr,
                                  device_ptr guide_ptr,
@@ -214,15 +197,6 @@ class CUDADevice : public Device {
 
   void denoise(RenderTile &rtile, DenoisingTask &denoising);
 
-  void adaptive_sampling_filter(uint filter_sample,
-                                WorkTile *wtile,
-                                CUdeviceptr d_wtile,
-                                CUstream stream = 0);
-  void adaptive_sampling_post(RenderTile &rtile,
-                              WorkTile *wtile,
-                              CUdeviceptr d_wtile,
-                              CUstream stream = 0);
-
   void path_trace(DeviceTask &task, RenderTile &rtile, device_vector<WorkTile> &work_tiles);
 
   void film_convert(DeviceTask &task,

intern/cycles/device/cuda/device_cuda_impl.cpp

@@ -39,8 +39,8 @@
 #  include "util/util_path.h"
 #  include "util/util_string.h"
 #  include "util/util_system.h"
-#  include "util/util_time.h"
 #  include "util/util_types.h"
+#  include "util/util_time.h"
 #  include "util/util_windows.h"
 
 #  include "kernel/split/kernel_split_data_types.h"
@@ -185,7 +185,7 @@ void CUDADevice::cuda_error_message(const string &message)
 }
 
 CUDADevice::CUDADevice(DeviceInfo &info, Stats &stats, Profiler &profiler, bool background_)
-    : Device(info, stats, profiler, background_), texture_info(this, "__texture_info", MEM_GLOBAL)
+    : Device(info, stats, profiler, background_), texture_info(this, "__texture_info", MEM_TEXTURE)
 {
   first_error = true;
   background = background_;
@@ -208,8 +208,6 @@ CUDADevice::CUDADevice(DeviceInfo &info, Stats &stats, Profiler &profiler, bool
   map_host_used = 0;
   can_map_host = 0;
 
-  functions.loaded = false;
-
   /* Intialize CUDA. */
   if (cuda_error(cuInit(0)))
     return;
@@ -533,42 +531,9 @@ bool CUDADevice::load_kernels(const DeviceRequestedFeatures &requested_features)
     reserve_local_memory(requested_features);
   }
 
-  load_functions();
-
   return (result == CUDA_SUCCESS);
 }
 
-void CUDADevice::load_functions()
-{
-  /* TODO: load all functions here. */
-  if (functions.loaded) {
-    return;
-  }
-  functions.loaded = true;
-
-  cuda_assert(cuModuleGetFunction(
-      &functions.adaptive_stopping, cuModule, "kernel_cuda_adaptive_stopping"));
-  cuda_assert(cuModuleGetFunction(
-      &functions.adaptive_filter_x, cuModule, "kernel_cuda_adaptive_filter_x"));
-  cuda_assert(cuModuleGetFunction(
-      &functions.adaptive_filter_y, cuModule, "kernel_cuda_adaptive_filter_y"));
-  cuda_assert(cuModuleGetFunction(
-      &functions.adaptive_scale_samples, cuModule, "kernel_cuda_adaptive_scale_samples"));
-
-  cuda_assert(cuFuncSetCacheConfig(functions.adaptive_stopping, CU_FUNC_CACHE_PREFER_L1));
-  cuda_assert(cuFuncSetCacheConfig(functions.adaptive_filter_x, CU_FUNC_CACHE_PREFER_L1));
-  cuda_assert(cuFuncSetCacheConfig(functions.adaptive_filter_y, CU_FUNC_CACHE_PREFER_L1));
-  cuda_assert(cuFuncSetCacheConfig(functions.adaptive_scale_samples, CU_FUNC_CACHE_PREFER_L1));
-
-  int unused_min_blocks;
-  cuda_assert(cuOccupancyMaxPotentialBlockSize(&unused_min_blocks,
-                                               &functions.adaptive_num_threads_per_block,
-                                               functions.adaptive_scale_samples,
-                                               NULL,
-                                               0,
-                                               0));
-}
-
 void CUDADevice::reserve_local_memory(const DeviceRequestedFeatures &requested_features)
 {
   if (use_split_kernel()) {
@@ -684,8 +649,7 @@ void CUDADevice::move_textures_to_host(size_t size, bool for_texture)
       device_memory &mem = *pair.first;
       CUDAMem *cmem = &pair.second;
 
-      bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) &&
-                        (&mem != &texture_info);
+      bool is_texture = (mem.type == MEM_TEXTURE) && (&mem != &texture_info);
       bool is_image = is_texture && (mem.data_height > 1);
 
       /* Can't move this type of memory. */
@@ -725,7 +689,8 @@ void CUDADevice::move_textures_to_host(size_t size, bool for_texture)
       device_ptr prev_pointer = max_mem->device_pointer;
       size_t prev_size = max_mem->device_size;
 
-      mem_copy_to(*max_mem);
+      tex_free(*max_mem);
+      tex_alloc(*max_mem);
       size = (max_size >= size) ? 0 : size - max_size;
 
       max_mem->device_pointer = prev_pointer;
@@ -759,7 +724,7 @@ CUDADevice::CUDAMem *CUDADevice::generic_alloc(device_memory &mem, size_t pitch_
    * If there is not enough room for working memory, we will try to move
    * textures to host memory, assuming the performance impact would have
    * been worse for working memory. */
-  bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) && (&mem != &texture_info);
+  bool is_texture = (mem.type == MEM_TEXTURE) && (&mem != &texture_info);
   bool is_image = is_texture && (mem.data_height > 1);
 
   size_t headroom = (is_texture) ? device_texture_headroom : device_working_headroom;
@@ -922,9 +887,6 @@ void CUDADevice::mem_alloc(device_memory &mem)
   else if (mem.type == MEM_TEXTURE) {
     assert(!"mem_alloc not supported for textures.");
   }
-  else if (mem.type == MEM_GLOBAL) {
-    assert(!"mem_alloc not supported for global memory.");
-  }
   else {
     generic_alloc(mem);
   }
@@ -935,13 +897,9 @@ void CUDADevice::mem_copy_to(device_memory &mem)
   if (mem.type == MEM_PIXELS) {
     assert(!"mem_copy_to not supported for pixels.");
   }
-  else if (mem.type == MEM_GLOBAL) {
-    global_free(mem);
-    global_alloc(mem);
-  }
   else if (mem.type == MEM_TEXTURE) {
-    tex_free((device_texture &)mem);
-    tex_alloc((device_texture &)mem);
+    tex_free(mem);
+    tex_alloc(mem);
   }
   else {
     if (!mem.device_pointer) {
@@ -957,7 +915,7 @@ void CUDADevice::mem_copy_from(device_memory &mem, int y, int w, int h, int elem
   if (mem.type == MEM_PIXELS && !background) {
     pixels_copy_from(mem, y, w, h);
   }
-  else if (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) {
+  else if (mem.type == MEM_TEXTURE) {
     assert(!"mem_copy_from not supported for textures.");
   }
   else if (mem.host_pointer) {
@@ -1000,11 +958,8 @@ void CUDADevice::mem_free(device_memory &mem)
   if (mem.type == MEM_PIXELS && !background) {
     pixels_free(mem);
   }
-  else if (mem.type == MEM_GLOBAL) {
-    global_free(mem);
-  }
   else if (mem.type == MEM_TEXTURE) {
-    tex_free((device_texture &)mem);
+    tex_free(mem);
   }
   else {
     generic_free(mem);
@@ -1027,25 +982,7 @@ void CUDADevice::const_copy_to(const char *name, void *host, size_t size)
   cuda_assert(cuMemcpyHtoD(mem, host, size));
 }
 
-void CUDADevice::global_alloc(device_memory &mem)
-{
-  CUDAContextScope scope(this);
-
-  generic_alloc(mem);
-  generic_copy_to(mem);
-
-  const_copy_to(mem.name, &mem.device_pointer, sizeof(mem.device_pointer));
-}
-
-void CUDADevice::global_free(device_memory &mem)
-{
-  if (mem.device_pointer) {
-    CUDAContextScope scope(this);
-    generic_free(mem);
-  }
-}
-
-void CUDADevice::tex_alloc(device_texture &mem)
+void CUDADevice::tex_alloc(device_memory &mem)
 {
   CUDAContextScope scope(this);
 
@@ -1055,7 +992,7 @@ void CUDADevice::tex_alloc(device_texture &mem)
   size_t size = mem.memory_size();
 
   CUaddress_mode address_mode = CU_TR_ADDRESS_MODE_WRAP;
-  switch (mem.info.extension) {
+  switch (mem.extension) {
     case EXTENSION_REPEAT:
       address_mode = CU_TR_ADDRESS_MODE_WRAP;
       break;
@@ -1071,13 +1008,22 @@ void CUDADevice::tex_alloc(device_texture &mem)
   }
 
   CUfilter_mode filter_mode;
-  if (mem.info.interpolation == INTERPOLATION_CLOSEST) {
+  if (mem.interpolation == INTERPOLATION_CLOSEST) {
     filter_mode = CU_TR_FILTER_MODE_POINT;
   }
   else {
     filter_mode = CU_TR_FILTER_MODE_LINEAR;
   }
 
+  /* Data Storage */
+  if (mem.interpolation == INTERPOLATION_NONE) {
+    generic_alloc(mem);
+    generic_copy_to(mem);
+
+    const_copy_to(bind_name.c_str(), &mem.device_pointer, sizeof(mem.device_pointer));
+    return;
+  }
+
   /* Image Texture Storage */
   CUarray_format_enum format;
   switch (mem.data_type) {
@@ -1188,6 +1134,15 @@ void CUDADevice::tex_alloc(device_texture &mem)
   }
 
   /* Kepler+, bindless textures. */
+  int flat_slot = 0;
+  if (string_startswith(mem.name, "__tex_image")) {
+    int pos = string(mem.name).rfind("_");
+    flat_slot = atoi(mem.name + pos + 1);
+  }
+  else {
+    assert(0);
+  }
+
   CUDA_RESOURCE_DESC resDesc;
   memset(&resDesc, 0, sizeof(resDesc));
 
@@ -1224,20 +1179,25 @@ void CUDADevice::tex_alloc(device_texture &mem)
   cuda_assert(cuTexObjectCreate(&cmem->texobject, &resDesc, &texDesc, NULL));
 
   /* Resize once */
-  const uint slot = mem.slot;
-  if (slot >= texture_info.size()) {
+  if (flat_slot >= texture_info.size()) {
     /* Allocate some slots in advance, to reduce amount
      * of re-allocations. */
-    texture_info.resize(slot + 128);
+    texture_info.resize(flat_slot + 128);
   }
 
   /* Set Mapping and tag that we need to (re-)upload to device */
-  texture_info[slot] = mem.info;
-  texture_info[slot].data = (uint64_t)cmem->texobject;
+  TextureInfo &info = texture_info[flat_slot];
+  info.data = (uint64_t)cmem->texobject;
+  info.cl_buffer = 0;
+  info.interpolation = mem.interpolation;
+  info.extension = mem.extension;
+  info.width = mem.data_width;
+  info.height = mem.data_height;
+  info.depth = mem.data_depth;
   need_texture_info = true;
 }
 
-void CUDADevice::tex_free(device_texture &mem)
+void CUDADevice::tex_free(device_memory &mem)
 {
   if (mem.device_pointer) {
     CUDAContextScope scope(this);
@@ -1706,80 +1666,6 @@ void CUDADevice::denoise(RenderTile &rtile, DenoisingTask &denoising)
   denoising.run_denoising(&rtile);
 }
 
-void CUDADevice::adaptive_sampling_filter(uint filter_sample,
-                                          WorkTile *wtile,
-                                          CUdeviceptr d_wtile,
-                                          CUstream stream)
-{
-  const int num_threads_per_block = functions.adaptive_num_threads_per_block;
-
-  /* These are a series of tiny kernels because there is no grid synchronization
-   * from within a kernel, so multiple kernel launches it is. */
-  uint total_work_size = wtile->h * wtile->w;
-  void *args2[] = {&d_wtile, &filter_sample, &total_work_size};
-  uint num_blocks = divide_up(total_work_size, num_threads_per_block);
-  cuda_assert(cuLaunchKernel(functions.adaptive_stopping,
-                             num_blocks,
-                             1,
-                             1,
-                             num_threads_per_block,
-                             1,
-                             1,
-                             0,
-                             stream,
-                             args2,
-                             0));
-  total_work_size = wtile->h;
-  num_blocks = divide_up(total_work_size, num_threads_per_block);
-  cuda_assert(cuLaunchKernel(functions.adaptive_filter_x,
-                             num_blocks,
-                             1,
-                             1,
-                             num_threads_per_block,
-                             1,
-                             1,
-                             0,
-                             stream,
-                             args2,
-                             0));
-  total_work_size = wtile->w;
-  num_blocks = divide_up(total_work_size, num_threads_per_block);
-  cuda_assert(cuLaunchKernel(functions.adaptive_filter_y,
-                             num_blocks,
-                             1,
-                             1,
-                             num_threads_per_block,
-                             1,
-                             1,
-                             0,
-                             stream,
-                             args2,
-                             0));
-}
-
-void CUDADevice::adaptive_sampling_post(RenderTile &rtile,
-                                        WorkTile *wtile,
-                                        CUdeviceptr d_wtile,
-                                        CUstream stream)
-{
-  const int num_threads_per_block = functions.adaptive_num_threads_per_block;
-  uint total_work_size = wtile->h * wtile->w;
-
-  void *args[] = {&d_wtile, &rtile.start_sample, &rtile.sample, &total_work_size};
-  uint num_blocks = divide_up(total_work_size, num_threads_per_block);
-  cuda_assert(cuLaunchKernel(functions.adaptive_scale_samples,
-                             num_blocks,
-                             1,
-                             1,
-                             num_threads_per_block,
-                             1,
-                             1,
-                             0,
-                             stream,
-                             args,
-                             0));
-}
-
 void CUDADevice::path_trace(DeviceTask &task,
                             RenderTile &rtile,
                             device_vector<WorkTile> &work_tiles)
@@ -1829,9 +1715,6 @@ void CUDADevice::path_trace(DeviceTask &task,
   }
 
   uint step_samples = divide_up(min_blocks * num_threads_per_block, wtile->w * wtile->h);
-  if (task.adaptive_sampling.use) {
-    step_samples = task.adaptive_sampling.align_static_samples(step_samples);
-  }
 
   /* Render all samples. */
   int start_sample = rtile.start_sample;
@@ -1853,12 +1736,6 @@ void CUDADevice::path_trace(DeviceTask &task,
     cuda_assert(
         cuLaunchKernel(cuPathTrace, num_blocks, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
 
-    /* Run the adaptive sampling kernels at selected samples aligned to step samples. */
-    uint filter_sample = sample + wtile->num_samples - 1;
-    if (task.adaptive_sampling.use && task.adaptive_sampling.need_filter(filter_sample)) {
-      adaptive_sampling_filter(filter_sample, wtile, d_work_tiles);
-    }
-
     cuda_assert(cuCtxSynchronize());
 
     /* Update progress. */
@@ -1870,14 +1747,6 @@ void CUDADevice::path_trace(DeviceTask &task,
         break;
     }
   }
-
-  /* Finalize adaptive sampling. */
-  if (task.adaptive_sampling.use) {
-    CUdeviceptr d_work_tiles = (CUdeviceptr)work_tiles.device_pointer;
-    adaptive_sampling_post(rtile, wtile, d_work_tiles);
-    cuda_assert(cuCtxSynchronize());
-    task.update_progress(&rtile, rtile.w * rtile.h * wtile->num_samples);
-  }
 }
 
 void CUDADevice::film_convert(DeviceTask &task,
@@ -2275,7 +2144,7 @@ void CUDADevice::thread_run(DeviceTask *task)
 {
   CUDAContextScope scope(this);
 
-  if (task->type == DeviceTask::RENDER) {
+  if (task->type == DeviceTask::RENDER || task->type == DeviceTask::DENOISE) {
     DeviceRequestedFeatures requested_features;
     if (use_split_kernel()) {
       if (split_kernel == NULL) {
@@ -2290,7 +2159,7 @@ void CUDADevice::thread_run(DeviceTask *task)
     RenderTile tile;
     DenoisingTask denoising(this, *task);
 
-    while (task->acquire_tile(this, tile, task->tile_types)) {
+    while (task->acquire_tile(this, tile)) {
       if (tile.task == RenderTile::PATH_TRACE) {
         if (use_split_kernel()) {
           device_only_memory<uchar> void_buffer(this, "void_buffer");

intern/cycles/device/device.cpp

@@ -25,11 +25,11 @@
 #include "util/util_logging.h"
 #include "util/util_math.h"
 #include "util/util_opengl.h"
-#include "util/util_string.h"
-#include "util/util_system.h"
 #include "util/util_time.h"
+#include "util/util_system.h"
 #include "util/util_types.h"
 #include "util/util_vector.h"
+#include "util/util_string.h"
 
 CCL_NAMESPACE_BEGIN
 

intern/cycles/device/device.h

@@ -27,8 +27,8 @@
 #include "util/util_list.h"
 #include "util/util_stats.h"
 #include "util/util_string.h"
-#include "util/util_texture.h"
 #include "util/util_thread.h"
+#include "util/util_texture.h"
 #include "util/util_types.h"
 #include "util/util_vector.h"
 

intern/cycles/device/device_cpu.cpp

@@ -29,19 +29,16 @@
 #include "device/device_intern.h"
 #include "device/device_split_kernel.h"
 
-// clang-format off
 #include "kernel/kernel.h"
 #include "kernel/kernel_compat_cpu.h"
 #include "kernel/kernel_types.h"
 #include "kernel/split/kernel_split_data.h"
 #include "kernel/kernel_globals.h"
-#include "kernel/kernel_adaptive_sampling.h"
 
 #include "kernel/filter/filter.h"
 
 #include "kernel/osl/osl_shader.h"
 #include "kernel/osl/osl_globals.h"
-// clang-format on
 
 #include "render/buffers.h"
 #include "render/coverage.h"
@@ -264,7 +261,7 @@ class CPUDevice : public Device {
 
   CPUDevice(DeviceInfo &info_, Stats &stats_, Profiler &profiler_, bool background_)
       : Device(info_, stats_, profiler_, background_),
-        texture_info(this, "__texture_info", MEM_GLOBAL),
+        texture_info(this, "__texture_info", MEM_TEXTURE),
 #define REGISTER_KERNEL(name) name##_kernel(KERNEL_FUNCTIONS(name))
         REGISTER_KERNEL(path_trace),
         REGISTER_KERNEL(convert_to_half_float),
@@ -320,10 +317,6 @@ class CPUDevice : public Device {
     REGISTER_SPLIT_KERNEL(next_iteration_setup);
     REGISTER_SPLIT_KERNEL(indirect_subsurface);
     REGISTER_SPLIT_KERNEL(buffer_update);
-    REGISTER_SPLIT_KERNEL(adaptive_stopping);
-    REGISTER_SPLIT_KERNEL(adaptive_filter_x);
-    REGISTER_SPLIT_KERNEL(adaptive_filter_y);
-    REGISTER_SPLIT_KERNEL(adaptive_adjust_samples);
 #undef REGISTER_SPLIT_KERNEL
 #undef KERNEL_FUNCTIONS
   }
@@ -372,9 +365,6 @@ class CPUDevice : public Device {
     if (mem.type == MEM_TEXTURE) {
       assert(!"mem_alloc not supported for textures.");
     }
-    else if (mem.type == MEM_GLOBAL) {
-      assert(!"mem_alloc not supported for global memory.");
-    }
     else {
       if (mem.name) {
         VLOG(1) << "Buffer allocate: " << mem.name << ", "
@@ -399,13 +389,9 @@ class CPUDevice : public Device {
 
   void mem_copy_to(device_memory &mem)
   {
-    if (mem.type == MEM_GLOBAL) {
-      global_free(mem);
-      global_alloc(mem);
-    }
-    else if (mem.type == MEM_TEXTURE) {
-      tex_free((device_texture &)mem);
-      tex_alloc((device_texture &)mem);
+    if (mem.type == MEM_TEXTURE) {
+      tex_free(mem);
+      tex_alloc(mem);
     }
     else if (mem.type == MEM_PIXELS) {
       assert(!"mem_copy_to not supported for pixels.");
@@ -437,11 +423,8 @@ class CPUDevice : public Device {
 
   void mem_free(device_memory &mem)
   {
-    if (mem.type == MEM_GLOBAL) {
-      global_free(mem);
-    }
-    else if (mem.type == MEM_TEXTURE) {
-      tex_free((device_texture &)mem);
+    if (mem.type == MEM_TEXTURE) {
+      tex_free(mem);
     }
     else if (mem.device_pointer) {
       if (mem.type == MEM_DEVICE_ONLY) {
@@ -463,50 +446,51 @@ class CPUDevice : public Device {
     kernel_const_copy(&kernel_globals, name, host, size);
   }
 
-  void global_alloc(device_memory &mem)
-  {
-    VLOG(1) << "Global memory allocate: " << mem.name << ", "
-            << string_human_readable_number(mem.memory_size()) << " bytes. ("
-            << string_human_readable_size(mem.memory_size()) << ")";
-
-    kernel_global_memory_copy(&kernel_globals, mem.name, mem.host_pointer, mem.data_size);
-
-    mem.device_pointer = (device_ptr)mem.host_pointer;
-    mem.device_size = mem.memory_size();
-    stats.mem_alloc(mem.device_size);
-  }
-
-  void global_free(device_memory &mem)
-  {
-    if (mem.device_pointer) {
-      mem.device_pointer = 0;
-      stats.mem_free(mem.device_size);
-      mem.device_size = 0;
-    }
-  }
-
-  void tex_alloc(device_texture &mem)
+  void tex_alloc(device_memory &mem)
   {
     VLOG(1) << "Texture allocate: " << mem.name << ", "
             << string_human_readable_number(mem.memory_size()) << " bytes. ("
             << string_human_readable_size(mem.memory_size()) << ")";
 
+    if (mem.interpolation == INTERPOLATION_NONE) {
+      /* Data texture. */
+      kernel_tex_copy(&kernel_globals, mem.name, mem.host_pointer, mem.data_size);
+    }
+    else {
+      /* Image Texture. */
+      int flat_slot = 0;
+      if (string_startswith(mem.name, "__tex_image")) {
+        int pos = string(mem.name).rfind("_");
+        flat_slot = atoi(mem.name + pos + 1);
+      }
+      else {
+        assert(0);
+      }
+
+      if (flat_slot >= texture_info.size()) {
+        /* Allocate some slots in advance, to reduce amount
+         * of re-allocations. */
+        texture_info.resize(flat_slot + 128);
+      }
+
+      TextureInfo &info = texture_info[flat_slot];
+      info.data = (uint64_t)mem.host_pointer;
+      info.cl_buffer = 0;
+      info.interpolation = mem.interpolation;
+      info.extension = mem.extension;
+      info.width = mem.data_width;
+      info.height = mem.data_height;
+      info.depth = mem.data_depth;
+
+      need_texture_info = true;
+    }
+
     mem.device_pointer = (device_ptr)mem.host_pointer;
     mem.device_size = mem.memory_size();
     stats.mem_alloc(mem.device_size);
-
-    const uint slot = mem.slot;
-    if (slot >= texture_info.size()) {
-      /* Allocate some slots in advance, to reduce amount of re-allocations. */
-      texture_info.resize(slot + 128);
-    }
-
-    texture_info[slot] = mem.info;
-    texture_info[slot].data = (uint64_t)mem.host_pointer;
-    need_texture_info = true;
   }
 
-  void tex_free(device_texture &mem)
+  void tex_free(device_memory &mem)
   {
     if (mem.device_pointer) {
       mem.device_pointer = 0;
@@ -527,7 +511,7 @@ class CPUDevice : public Device {
 
   void thread_run(DeviceTask *task)
   {
-    if (task->type == DeviceTask::RENDER)
+    if (task->type == DeviceTask::RENDER || task->type == DeviceTask::DENOISE)
       thread_render(*task);
     else if (task->type == DeviceTask::SHADER)
       thread_shader(*task);
@@ -839,49 +823,6 @@ class CPUDevice : public Device {
     return true;
   }
 
-  bool adaptive_sampling_filter(KernelGlobals *kg, RenderTile &tile)
-  {
-    WorkTile wtile;
-    wtile.x = tile.x;
-    wtile.y = tile.y;
-    wtile.w = tile.w;
-    wtile.h = tile.h;
-    wtile.offset = tile.offset;
-    wtile.stride = tile.stride;
-    wtile.buffer = (float *)tile.buffer;
-
-    bool any = false;
-    for (int y = tile.y; y < tile.y + tile.h; ++y) {
-      any |= kernel_do_adaptive_filter_x(kg, y, &wtile);
-    }
-    for (int x = tile.x; x < tile.x + tile.w; ++x) {
-      any |= kernel_do_adaptive_filter_y(kg, x, &wtile);
-    }
-    return (!any);
-  }
-
-  void adaptive_sampling_post(const RenderTile &tile, KernelGlobals *kg)
-  {
-    float *render_buffer = (float *)tile.buffer;
-    for (int y = tile.y; y < tile.y + tile.h; y++) {
-      for (int x = tile.x; x < tile.x + tile.w; x++) {
-        int index = tile.offset + x + y * tile.stride;
-        ccl_global float *buffer = render_buffer + index * kernel_data.film.pass_stride;
-        if (buffer[kernel_data.film.pass_sample_count] < 0.0f) {
-          buffer[kernel_data.film.pass_sample_count] = -buffer[kernel_data.film.pass_sample_count];
-          float sample_multiplier = tile.sample / max((float)tile.start_sample + 1.0f,
-                                                      buffer[kernel_data.film.pass_sample_count]);
-          if (sample_multiplier != 1.0f) {
-            kernel_adaptive_post_adjust(kg, buffer, sample_multiplier);
-          }
-        }
-        else {
-          kernel_adaptive_post_adjust(kg, buffer, tile.sample / (tile.sample - 1.0f));
-        }
-      }
-    }
-  }
-
   void path_trace(DeviceTask &task, RenderTile &tile, KernelGlobals *kg)
   {
     const bool use_coverage = kernel_data.film.cryptomatte_passes & CRYPT_ACCURATE;
@@ -914,27 +855,14 @@ class CPUDevice : public Device {
           path_trace_kernel()(kg, render_buffer, sample, x, y, tile.offset, tile.stride);
         }
       }
-      tile.sample = sample + 1;
 
-      if (task.adaptive_sampling.use && task.adaptive_sampling.need_filter(sample)) {
-        const bool stop = adaptive_sampling_filter(kg, tile);
-        if (stop) {
-          const int num_progress_samples = end_sample - sample;
-          tile.sample = end_sample;
-          task.update_progress(&tile, tile.w * tile.h * num_progress_samples);
-          break;
-        }
-      }
+      tile.sample = sample + 1;
 
       task.update_progress(&tile, tile.w * tile.h);
     }
     if (use_coverage) {
       coverage.finalize();
     }
-
-    if (task.adaptive_sampling.use) {
-      adaptive_sampling_post(tile, kg);
-    }
   }
 
   void denoise(DenoisingTask &denoising, RenderTile &tile)
@@ -999,7 +927,7 @@ class CPUDevice : public Device {
     DenoisingTask denoising(this, task);
     denoising.profiler = &kg->profiler;
 
-    while (task.acquire_tile(this, tile, task.tile_types)) {
+    while (task.acquire_tile(this, tile)) {
       if (tile.task == RenderTile::PATH_TRACE) {
         if (use_split_kernel) {
           device_only_memory<uchar> void_buffer(this, "void_buffer");

intern/cycles/device/device_cuda.cpp

@@ -16,9 +16,9 @@
 
 #ifdef WITH_CUDA
 
-#  include "device/cuda/device_cuda.h"
 #  include "device/device.h"
 #  include "device/device_intern.h"
+#  include "device/cuda/device_cuda.h"
 
 #  include "util/util_logging.h"
 #  include "util/util_string.h"

intern/cycles/device/device_intern.h

@@ -17,15 +17,9 @@
 #ifndef __DEVICE_INTERN_H__
 #define __DEVICE_INTERN_H__
 
-#include "util/util_string.h"
-#include "util/util_vector.h"
-
 CCL_NAMESPACE_BEGIN
 
 class Device;
-class DeviceInfo;
-class Profiler;
-class Stats;
 
 Device *device_cpu_create(DeviceInfo &info, Stats &stats, Profiler &profiler, bool background);
 bool device_opencl_init();

intern/cycles/device/device_memory.cpp

@@ -14,8 +14,8 @@
  * limitations under the License.
  */
 
-#include "device/device_memory.h"
 #include "device/device.h"
+#include "device/device_memory.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -31,6 +31,8 @@ device_memory::device_memory(Device *device, const char *name, MemoryType type)
       data_depth(0),
       type(type),
       name(name),
+      interpolation(INTERPOLATION_NONE),
+      extension(EXTENSION_REPEAT),
       device(device),
       device_pointer(0),
       host_pointer(0),
@@ -74,7 +76,7 @@ void device_memory::host_free()
 
 void device_memory::device_alloc()
 {
-  assert(!device_pointer && type != MEM_TEXTURE && type != MEM_GLOBAL);
+  assert(!device_pointer && type != MEM_TEXTURE);
   device->mem_alloc(*this);
 }
 
@@ -94,7 +96,7 @@ void device_memory::device_copy_to()
 
 void device_memory::device_copy_from(int y, int w, int h, int elem)
 {
-  assert(type != MEM_TEXTURE && type != MEM_READ_ONLY && type != MEM_GLOBAL);
+  assert(type != MEM_TEXTURE && type != MEM_READ_ONLY);
   device->mem_copy_from(*this, y, w, h, elem);
 }
 
@@ -137,93 +139,4 @@ device_sub_ptr::~device_sub_ptr()
   device->mem_free_sub_ptr(ptr);
 }
 
-/* Device Texture */
-
-device_texture::device_texture(Device *device,
-                               const char *name,
-                               const uint slot,
-                               ImageDataType image_data_type,
-                               InterpolationType interpolation,
-                               ExtensionType extension)
-    : device_memory(device, name, MEM_TEXTURE), slot(slot)
-{
-  switch (image_data_type) {
-    case IMAGE_DATA_TYPE_FLOAT4:
-      data_type = TYPE_FLOAT;
-      data_elements = 4;
-      break;
-    case IMAGE_DATA_TYPE_FLOAT:
-      data_type = TYPE_FLOAT;
-      data_elements = 1;
-      break;
-    case IMAGE_DATA_TYPE_BYTE4:
-      data_type = TYPE_UCHAR;
-      data_elements = 4;
-      break;
-    case IMAGE_DATA_TYPE_BYTE:
-      data_type = TYPE_UCHAR;
-      data_elements = 1;
-      break;
-    case IMAGE_DATA_TYPE_HALF4:
-      data_type = TYPE_HALF;
-      data_elements = 4;
-      break;
-    case IMAGE_DATA_TYPE_HALF:
-      data_type = TYPE_HALF;
-      data_elements = 1;
-      break;
-    case IMAGE_DATA_TYPE_USHORT4:
-      data_type = TYPE_UINT16;
-      data_elements = 4;
-      break;
-    case IMAGE_DATA_TYPE_USHORT:
-      data_type = TYPE_UINT16;
-      data_elements = 1;
-      break;
-    case IMAGE_DATA_NUM_TYPES:
-      assert(0);
-      return;
-  }
-
-  memset(&info, 0, sizeof(info));
-  info.data_type = image_data_type;
-  info.interpolation = interpolation;
-  info.extension = extension;
-}
-
-device_texture::~device_texture()
-{
-  device_free();
-  host_free();
-}
-
-/* Host memory allocation. */
-void *device_texture::alloc(const size_t width, const size_t height, const size_t depth)
-{
-  const size_t new_size = size(width, height, depth);
-
-  if (new_size != data_size) {
-    device_free();
-    host_free();
-    host_pointer = host_alloc(data_elements * datatype_size(data_type) * new_size);
-    assert(device_pointer == 0);
-  }
-
-  data_size = new_size;
-  data_width = width;
-  data_height = height;
-  data_depth = depth;
-
-  info.width = width;
-  info.height = height;
-  info.depth = depth;
-
-  return host_pointer;
-}
-
-void device_texture::copy_to_device()
-{
-  device_copy_to();
-}
-
 CCL_NAMESPACE_END

intern/cycles/device/device_memory.h

@@ -23,7 +23,6 @@
 
 #include "util/util_array.h"
 #include "util/util_half.h"
-#include "util/util_string.h"
 #include "util/util_texture.h"
 #include "util/util_types.h"
 #include "util/util_vector.h"
@@ -32,14 +31,7 @@ CCL_NAMESPACE_BEGIN
 
 class Device;
 
-enum MemoryType {
-  MEM_READ_ONLY,
-  MEM_READ_WRITE,
-  MEM_DEVICE_ONLY,
-  MEM_GLOBAL,
-  MEM_TEXTURE,
-  MEM_PIXELS
-};
+enum MemoryType { MEM_READ_ONLY, MEM_READ_WRITE, MEM_DEVICE_ONLY, MEM_TEXTURE, MEM_PIXELS };
 
 /* Supported Data Types */
 
@@ -216,6 +208,8 @@ class device_memory {
   size_t data_depth;
   MemoryType type;
   const char *name;
+  InterpolationType interpolation;
+  ExtensionType extension;
 
   /* Pointers. */
   Device *device;
@@ -316,7 +310,7 @@ template<typename T> class device_only_memory : public device_memory {
  * in and copied to the device with copy_to_device(). Or alternatively
  * allocated and set to zero on the device with zero_to_device().
  *
- * When using memory type MEM_GLOBAL, a pointer to this memory will be
+ * When using memory type MEM_TEXTURE, a pointer to this memory will be
  * automatically attached to kernel globals, using the provided name
  * matching an entry in kernel_textures.h. */
 
@@ -509,33 +503,6 @@ class device_sub_ptr {
   device_ptr ptr;
 };
 
-/* Device Texture
- *
- * 2D or 3D image texture memory. */
-
-class device_texture : public device_memory {
- public:
-  device_texture(Device *device,
-                 const char *name,
-                 const uint slot,
-                 ImageDataType image_data_type,
-                 InterpolationType interpolation,
-                 ExtensionType extension);
-  ~device_texture();
-
-  void *alloc(const size_t width, const size_t height, const size_t depth = 0);
-  void copy_to_device();
-
-  uint slot;
-  TextureInfo info;
-
- protected:
-  size_t size(const size_t width, const size_t height, const size_t depth)
-  {
-    return width * ((height == 0) ? 1 : height) * ((depth == 0) ? 1 : depth);
-  }
-};
-
 CCL_NAMESPACE_END
 
 #endif /* __DEVICE_MEMORY_H__ */

intern/cycles/device/device_multi.cpp

@@ -14,8 +14,8 @@
  * limitations under the License.
  */
 
-#include <sstream>
 #include <stdlib.h>
+#include <sstream>
 
 #include "device/device.h"
 #include "device/device_intern.h"
@@ -183,14 +183,6 @@ class MultiDevice : public Device {
     return true;
   }
 
-  virtual void *osl_memory()
-  {
-    if (devices.size() > 1) {
-      return NULL;
-    }
-    return devices.front().device->osl_memory();
-  }
-
   void mem_alloc(device_memory &mem)
   {
     device_ptr key = unique_key++;
@@ -490,24 +482,11 @@ class MultiDevice : public Device {
 
   void task_add(DeviceTask &task)
   {
-    list<SubDevice> task_devices = devices;
-    if (!denoising_devices.empty()) {
-      if (task.type == DeviceTask::DENOISE_BUFFER) {
-        /* Denoising tasks should be redirected to the denoising devices entirely. */
-        task_devices = denoising_devices;
-      }
-      else if (task.type == DeviceTask::RENDER && (task.tile_types & RenderTile::DENOISE)) {
-        const uint tile_types = task.tile_types;
-        /* For normal rendering tasks only redirect the denoising part to the denoising devices.
-         * Do not need to split the task here, since they all run through 'acquire_tile'. */
-        task.tile_types = RenderTile::DENOISE;
-        foreach (SubDevice &sub, denoising_devices) {
-          sub.device->task_add(task);
-        }
-        /* Rendering itself should still be executed on the rendering devices. */
-        task.tile_types = tile_types ^ RenderTile::DENOISE;
-      }
-    }
+    list<SubDevice> &task_devices = denoising_devices.empty() ||
+                                            (task.type != DeviceTask::DENOISE &&
+                                             task.type != DeviceTask::DENOISE_BUFFER) ?
+                                        devices :
+                                        denoising_devices;
 
     list<DeviceTask> tasks;
     task.split(tasks, task_devices.size());

intern/cycles/device/device_network.cpp

@@ -14,9 +14,9 @@
  * limitations under the License.
  */
 
-#include "device/device_network.h"
 #include "device/device.h"
 #include "device/device_intern.h"
+#include "device/device_network.h"
 
 #include "util/util_foreach.h"
 #include "util/util_logging.h"

intern/cycles/device/device_network.h

@@ -19,19 +19,19 @@
 
 #ifdef WITH_NETWORK
 
-#  include <boost/archive/binary_iarchive.hpp>
-#  include <boost/archive/binary_oarchive.hpp>
 #  include <boost/archive/text_iarchive.hpp>
 #  include <boost/archive/text_oarchive.hpp>
+#  include <boost/archive/binary_iarchive.hpp>
+#  include <boost/archive/binary_oarchive.hpp>
 #  include <boost/array.hpp>
 #  include <boost/asio.hpp>
 #  include <boost/bind.hpp>
 #  include <boost/serialization/vector.hpp>
 #  include <boost/thread.hpp>
 
-#  include <deque>
 #  include <iostream>
 #  include <sstream>
+#  include <deque>
 
 #  include "render/buffers.h"
 

intern/cycles/device/device_opencl.cpp

@@ -16,9 +16,9 @@
 
 #ifdef WITH_OPENCL
 
-#  include "device/opencl/device_opencl.h"
 #  include "device/device.h"
 #  include "device/device_intern.h"
+#  include "device/opencl/device_opencl.h"
 
 #  include "util/util_foreach.h"
 #  include "util/util_logging.h"

intern/cycles/device/device_optix.cpp

@@ -17,28 +17,28 @@
 
 #ifdef WITH_OPTIX
 
-#  include "bvh/bvh.h"
 #  include "device/cuda/device_cuda.h"
-#  include "device/device_denoising.h"
 #  include "device/device_intern.h"
-#  include "render/buffers.h"
+#  include "device/device_denoising.h"
+#  include "bvh/bvh.h"
+#  include "render/scene.h"
 #  include "render/hair.h"
 #  include "render/mesh.h"
 #  include "render/object.h"
-#  include "render/scene.h"
-#  include "util/util_debug.h"
-#  include "util/util_logging.h"
+#  include "render/buffers.h"
 #  include "util/util_md5.h"
 #  include "util/util_path.h"
 #  include "util/util_time.h"
+#  include "util/util_debug.h"
+#  include "util/util_logging.h"
 
 #  ifdef WITH_CUDA_DYNLOAD
 #    include <cuew.h>
 // Do not use CUDA SDK headers when using CUEW
 #    define OPTIX_DONT_INCLUDE_CUDA
 #  endif
-#  include <optix_function_table_definition.h>
 #  include <optix_stubs.h>
+#  include <optix_function_table_definition.h>
 
 // TODO(pmours): Disable this once drivers have native support
 #  define OPTIX_DENOISER_NO_PIXEL_STRIDE 1
@@ -186,15 +186,14 @@ class OptiXDevice : public CUDADevice {
   OptixTraversableHandle tlas_handle = 0;
 
   OptixDenoiser denoiser = NULL;
-  device_only_memory<unsigned char> denoiser_state;
+  pair<int2, CUdeviceptr> denoiser_state = {};
   int denoiser_input_passes = 0;
 
  public:
   OptiXDevice(DeviceInfo &info_, Stats &stats_, Profiler &profiler_, bool background_)
       : CUDADevice(info_, stats_, profiler_, background_),
         sbt_data(this, "__sbt", MEM_READ_ONLY),
-        launch_params(this, "__params"),
-        denoiser_state(this, "__denoiser_state")
+        launch_params(this, "__params")
   {
     // Store number of CUDA streams in device info
     info.cpu_threads = DebugFlags().optix.cuda_streams;
@@ -256,10 +255,13 @@ class OptiXDevice : public CUDADevice {
       cuMemFree(mem);
     }
 
+    if (denoiser_state.second) {
+      cuMemFree(denoiser_state.second);
+    }
+
     sbt_data.free();
     texture_info.free();
     launch_params.free();
-    denoiser_state.free();
 
     // Unload modules
     if (optix_module != NULL)
@@ -477,9 +479,9 @@ class OptiXDevice : public CUDADevice {
     // Calculate maximum trace continuation stack size
     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 = max(trace_css, stack_size[PG_HITD].cssIS + stack_size[PG_HITD].cssAH);
+    trace_css = max(trace_css, stack_size[PG_HITL].cssIS + stack_size[PG_HITL].cssAH);
+    trace_css = max(trace_css, stack_size[PG_HITS].cssIS + stack_size[PG_HITS].cssAH);
 
     OptixPipelineLinkOptions link_options;
     link_options.maxTraceDepth = 1;
@@ -548,9 +550,8 @@ class OptiXDevice : public CUDADevice {
                               &pipelines[PIP_SHADER_EVAL]));
 
       // Calculate continuation stack size based on the maximum of all ray generation stack sizes
-      const unsigned int css = std::max(stack_size[PG_BAKE].cssRG,
-                                        std::max(stack_size[PG_DISP].cssRG,
-                                                 stack_size[PG_BACK].cssRG)) +
+      const unsigned int css = max(stack_size[PG_BAKE].cssRG,
+                                   max(stack_size[PG_DISP].cssRG, stack_size[PG_BACK].cssRG)) +
                                link_options.maxTraceDepth * trace_css;
 
       check_result_optix_ret(optixPipelineSetStackSize(
@@ -570,14 +571,9 @@ class OptiXDevice : public CUDADevice {
     if (have_error())
       return;  // Abort early if there was an error previously
 
-    if (task.type == DeviceTask::RENDER) {
-      if (thread_index != 0) {
-        // Only execute denoising in a single thread (see also 'task_add')
-        task.tile_types &= ~RenderTile::DENOISE;
-      }
-
+    if (task.type == DeviceTask::RENDER || task.type == DeviceTask::DENOISE) {
       RenderTile tile;
-      while (task.acquire_tile(this, tile, task.tile_types)) {
+      while (task.acquire_tile(this, tile)) {
         if (tile.task == RenderTile::PATH_TRACE)
           launch_render(task, tile, thread_index);
         else if (tile.task == RenderTile::DENOISE)
@@ -628,11 +624,7 @@ class OptiXDevice : public CUDADevice {
 
     const int end_sample = rtile.start_sample + rtile.num_samples;
     // Keep this number reasonable to avoid running into TDRs
-    int step_samples = (info.display_device ? 8 : 32);
-    if (task.adaptive_sampling.use) {
-      step_samples = task.adaptive_sampling.align_static_samples(step_samples);
-    }
-
+    const int step_samples = (info.display_device ? 8 : 32);
     // Offset into launch params buffer so that streams use separate data
     device_ptr launch_params_ptr = launch_params.device_pointer +
                                    thread_index * launch_params.data_elements;
@@ -643,9 +635,10 @@ class OptiXDevice : public CUDADevice {
       // Copy work tile information to device
       wtile.num_samples = min(step_samples, end_sample - sample);
       wtile.start_sample = sample;
-      device_ptr d_wtile_ptr = launch_params_ptr + offsetof(KernelParams, tile);
-      check_result_cuda(
-          cuMemcpyHtoDAsync(d_wtile_ptr, &wtile, sizeof(wtile), cuda_stream[thread_index]));
+      check_result_cuda(cuMemcpyHtoDAsync(launch_params_ptr + offsetof(KernelParams, tile),
+                                          &wtile,
+                                          sizeof(wtile),
+                                          cuda_stream[thread_index]));
 
       OptixShaderBindingTable sbt_params = {};
       sbt_params.raygenRecord = sbt_data.device_pointer + PG_RGEN * sizeof(SbtRecord);
@@ -670,12 +663,6 @@ class OptiXDevice : public CUDADevice {
                                      wtile.h,
                                      1));
 
-      // Run the adaptive sampling kernels at selected samples aligned to step samples.
-      uint filter_sample = wtile.start_sample + wtile.num_samples - 1;
-      if (task.adaptive_sampling.use && task.adaptive_sampling.need_filter(filter_sample)) {
-        adaptive_sampling_filter(filter_sample, &wtile, d_wtile_ptr, cuda_stream[thread_index]);
-      }
-
       // Wait for launch to finish
       check_result_cuda(cuStreamSynchronize(cuda_stream[thread_index]));
 
@@ -687,14 +674,6 @@ class OptiXDevice : public CUDADevice {
       if (task.get_cancel() && !task.need_finish_queue)
         return;  // Cancel rendering
     }
-
-    // Finalize adaptive sampling
-    if (task.adaptive_sampling.use) {
-      device_ptr d_wtile_ptr = launch_params_ptr + offsetof(KernelParams, tile);
-      adaptive_sampling_post(rtile, &wtile, d_wtile_ptr, cuda_stream[thread_index]);
-      check_result_cuda(cuStreamSynchronize(cuda_stream[thread_index]));
-      task.update_progress(&rtile, rtile.w * rtile.h * wtile.num_samples);
-    }
   }
 
   bool launch_denoise(DeviceTask &task, RenderTile &rtile)
@@ -834,26 +813,32 @@ class OptiXDevice : public CUDADevice {
       check_result_optix_ret(
           optixDenoiserComputeMemoryResources(denoiser, rect_size.x, rect_size.y, &sizes));
 
+      auto &state = denoiser_state.second;
+      auto &state_size = denoiser_state.first;
       const size_t scratch_size = sizes.recommendedScratchSizeInBytes;
       const size_t scratch_offset = sizes.stateSizeInBytes;
 
       // Allocate denoiser state if tile size has changed since last setup
-      if (recreate_denoiser || (denoiser_state.data_width != rect_size.x ||
-                                denoiser_state.data_height != rect_size.y)) {
-        denoiser_state.alloc_to_device(scratch_offset + scratch_size);
+      if (state_size.x != rect_size.x || state_size.y != rect_size.y || recreate_denoiser) {
+        // Free existing state before allocating new one
+        if (state) {
+          cuMemFree(state);
+          state = 0;
+        }
+
+        check_result_cuda_ret(cuMemAlloc(&state, scratch_offset + scratch_size));
 
         // Initialize denoiser state for the current tile size
         check_result_optix_ret(optixDenoiserSetup(denoiser,
                                                   0,
                                                   rect_size.x,
                                                   rect_size.y,
-                                                  denoiser_state.device_pointer,
+                                                  state,
                                                   scratch_offset,
-                                                  denoiser_state.device_pointer + scratch_offset,
+                                                  state + scratch_offset,
                                                   scratch_size));
 
-        denoiser_state.data_width = rect_size.x;
-        denoiser_state.data_height = rect_size.y;
+        state_size = rect_size;
       }
 
       // Set up input and output layer information
@@ -895,14 +880,14 @@ class OptiXDevice : public CUDADevice {
       check_result_optix_ret(optixDenoiserInvoke(denoiser,
                                                  0,
                                                  &params,
-                                                 denoiser_state.device_pointer,
+                                                 state,
                                                  scratch_offset,
                                                  input_layers,
                                                  task.denoising.optix_input_passes,
                                                  overlap_offset.x,
                                                  overlap_offset.y,
                                                  output_layers,
-                                                 denoiser_state.device_pointer + scratch_offset,
+                                                 state + scratch_offset,
                                                  scratch_size));
 
 #  if OPTIX_DENOISER_NO_PIXEL_STRIDE
@@ -1474,7 +1459,7 @@ class OptiXDevice : public CUDADevice {
       return;
     }
 
-    if (task.type == DeviceTask::DENOISE_BUFFER) {
+    if (task.type == DeviceTask::DENOISE || task.type == DeviceTask::DENOISE_BUFFER) {
       // Execute denoising in a single thread (e.g. to avoid race conditions during creation)
       task_pool.push(new OptiXDeviceTask(this, task, 0));
       return;

intern/cycles/device/device_split_kernel.cpp

@@ -55,10 +55,6 @@ DeviceSplitKernel::DeviceSplitKernel(Device *device)
   kernel_next_iteration_setup = NULL;
   kernel_indirect_subsurface = NULL;
   kernel_buffer_update = NULL;
-  kernel_adaptive_stopping = NULL;
-  kernel_adaptive_filter_x = NULL;
-  kernel_adaptive_filter_y = NULL;
-  kernel_adaptive_adjust_samples = NULL;
 }
 
 DeviceSplitKernel::~DeviceSplitKernel()
@@ -87,10 +83,6 @@ DeviceSplitKernel::~DeviceSplitKernel()
   delete kernel_next_iteration_setup;
   delete kernel_indirect_subsurface;
   delete kernel_buffer_update;
-  delete kernel_adaptive_stopping;
-  delete kernel_adaptive_filter_x;
-  delete kernel_adaptive_filter_y;
-  delete kernel_adaptive_adjust_samples;
 }
 
 bool DeviceSplitKernel::load_kernels(const DeviceRequestedFeatures &requested_features)
@@ -122,10 +114,6 @@ bool DeviceSplitKernel::load_kernels(const DeviceRequestedFeatures &requested_fe
   LOAD_KERNEL(next_iteration_setup);
   LOAD_KERNEL(indirect_subsurface);
   LOAD_KERNEL(buffer_update);
-  LOAD_KERNEL(adaptive_stopping);
-  LOAD_KERNEL(adaptive_filter_x);
-  LOAD_KERNEL(adaptive_filter_y);
-  LOAD_KERNEL(adaptive_adjust_samples);
 
 #undef LOAD_KERNEL
 
@@ -214,21 +202,13 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
     /* initial guess to start rolling average */
     const int initial_num_samples = 1;
     /* approx number of samples per second */
-    const int samples_per_second = (avg_time_per_sample > 0.0) ?
-                                       int(double(time_multiplier) / avg_time_per_sample) + 1 :
-                                       initial_num_samples;
+    int samples_per_second = (avg_time_per_sample > 0.0) ?
+                                 int(double(time_multiplier) / avg_time_per_sample) + 1 :
+                                 initial_num_samples;
 
     RenderTile subtile = tile;
     subtile.start_sample = tile.sample;
-    subtile.num_samples = samples_per_second;
-
-    if (task->adaptive_sampling.use) {
-      subtile.num_samples = task->adaptive_sampling.align_dynamic_samples(subtile.start_sample,
-                                                                          subtile.num_samples);
-    }
-
-    /* Don't go beyond requested number of samples. */
-    subtile.num_samples = min(subtile.num_samples,
+    subtile.num_samples = min(samples_per_second,
                               tile.start_sample + tile.num_samples - tile.sample);
 
     if (device->have_error()) {
@@ -322,23 +302,6 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
       }
     }
 
-    int filter_sample = tile.sample + subtile.num_samples - 1;
-    if (task->adaptive_sampling.use && task->adaptive_sampling.need_filter(filter_sample)) {
-      size_t buffer_size[2];
-      buffer_size[0] = round_up(tile.w, local_size[0]);
-      buffer_size[1] = round_up(tile.h, local_size[1]);
-      kernel_adaptive_stopping->enqueue(
-          KernelDimensions(buffer_size, local_size), kgbuffer, kernel_data);
-      buffer_size[0] = round_up(tile.h, local_size[0]);
-      buffer_size[1] = round_up(1, local_size[1]);
-      kernel_adaptive_filter_x->enqueue(
-          KernelDimensions(buffer_size, local_size), kgbuffer, kernel_data);
-      buffer_size[0] = round_up(tile.w, local_size[0]);
-      buffer_size[1] = round_up(1, local_size[1]);
-      kernel_adaptive_filter_y->enqueue(
-          KernelDimensions(buffer_size, local_size), kgbuffer, kernel_data);
-    }
-
     double time_per_sample = ((time_dt() - start_time) / subtile.num_samples);
 
     if (avg_time_per_sample == 0.0) {
@@ -361,28 +324,6 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
     }
   }
 
-  if (task->adaptive_sampling.use) {
-    /* Reset the start samples. */
-    RenderTile subtile = tile;
-    subtile.start_sample = tile.start_sample;
-    subtile.num_samples = tile.sample - tile.start_sample;
-    enqueue_split_kernel_data_init(KernelDimensions(global_size, local_size),
-                                   subtile,
-                                   num_global_elements,
-                                   kgbuffer,
-                                   kernel_data,
-                                   split_data,
-                                   ray_state,
-                                   queue_index,
-                                   use_queues_flag,
-                                   work_pool_wgs);
-    size_t buffer_size[2];
-    buffer_size[0] = round_up(tile.w, local_size[0]);
-    buffer_size[1] = round_up(tile.h, local_size[1]);
-    kernel_adaptive_adjust_samples->enqueue(
-        KernelDimensions(buffer_size, local_size), kgbuffer, kernel_data);
-  }
-
   return true;
 }
 

intern/cycles/device/device_split_kernel.h

@@ -75,10 +75,6 @@ class DeviceSplitKernel {
   SplitKernelFunction *kernel_next_iteration_setup;
   SplitKernelFunction *kernel_indirect_subsurface;
   SplitKernelFunction *kernel_buffer_update;
-  SplitKernelFunction *kernel_adaptive_stopping;
-  SplitKernelFunction *kernel_adaptive_filter_x;
-  SplitKernelFunction *kernel_adaptive_filter_y;
-  SplitKernelFunction *kernel_adaptive_adjust_samples;
 
   /* Global memory variables [porting]; These memory is used for
    * co-operation between different kernels; Data written by one

intern/cycles/device/device_task.cpp

@@ -68,7 +68,7 @@ int DeviceTask::get_subtask_count(int num, int max_size)
   if (type == SHADER) {
     num = min(shader_w, num);
   }
-  else if (type == RENDER) {
+  else if (type == RENDER || type == DENOISE) {
   }
   else {
     num = min(h, num);
@@ -94,7 +94,7 @@ void DeviceTask::split(list<DeviceTask> &tasks, int num, int max_size)
       tasks.push_back(task);
     }
   }
-  else if (type == RENDER) {
+  else if (type == RENDER || type == DENOISE) {
     for (int i = 0; i < num; i++)
       tasks.push_back(*this);
   }
@@ -136,59 +136,4 @@ void DeviceTask::update_progress(RenderTile *rtile, int pixel_samples)
   }
 }
 
-/* Adaptive Sampling */
-
-AdaptiveSampling::AdaptiveSampling()
-    : use(true), adaptive_step(ADAPTIVE_SAMPLE_STEP), min_samples(0)
-{
-}
-
-/* Render samples in steps that align with the adaptive filtering. */
-int AdaptiveSampling::align_static_samples(int samples) const
-{
-  if (samples > adaptive_step) {
-    /* Make multiple of adaptive_step. */
-    while (samples % adaptive_step != 0) {
-      samples--;
-    }
-  }
-  else if (samples < adaptive_step) {
-    /* Make divisor of adaptive_step. */
-    while (adaptive_step % samples != 0) {
-      samples--;
-    }
-  }
-
-  return max(samples, 1);
-}
-
-/* Render samples in steps that align with the adaptive filtering, with the
- * suggested number of samples dynamically changing. */
-int AdaptiveSampling::align_dynamic_samples(int offset, int samples) const
-{
-  /* Round so that we end up on multiples of adaptive_samples. */
-  samples += offset;
-
-  if (samples > adaptive_step) {
-    /* Make multiple of adaptive_step. */
-    while (samples % adaptive_step != 0) {
-      samples--;
-    }
-  }
-
-  samples -= offset;
-
-  return max(samples, 1);
-}
-
-bool AdaptiveSampling::need_filter(int sample) const
-{
-  if (sample > min_samples) {
-    return (sample & (adaptive_step - 1)) == (adaptive_step - 1);
-  }
-  else {
-    return false;
-  }
-}
-
 CCL_NAMESPACE_END

intern/cycles/device/device_task.h

@@ -62,22 +62,9 @@ class DenoiseParams {
   }
 };
 
-class AdaptiveSampling {
- public:
-  AdaptiveSampling();
-
-  int align_static_samples(int samples) const;
-  int align_dynamic_samples(int offset, int samples) const;
-  bool need_filter(int sample) const;
-
-  bool use;
-  int adaptive_step;
-  int min_samples;
-};
-
 class DeviceTask : public Task {
  public:
-  typedef enum { RENDER, FILM_CONVERT, SHADER, DENOISE_BUFFER } Type;
+  typedef enum { RENDER, DENOISE, DENOISE_BUFFER, FILM_CONVERT, SHADER } Type;
   Type type;
 
   int x, y, w, h;
@@ -103,7 +90,7 @@ class DeviceTask : public Task {
 
   void update_progress(RenderTile *rtile, int pixel_samples = -1);
 
-  function<bool(Device *device, RenderTile &, uint)> acquire_tile;
+  function<bool(Device *device, RenderTile &)> acquire_tile;
   function<void(long, int)> update_progress_sample;
   function<void(RenderTile &)> update_tile_sample;
   function<void(RenderTile &)> release_tile;
@@ -111,7 +98,6 @@ class DeviceTask : public Task {
   function<void(RenderTile *, Device *)> map_neighbor_tiles;
   function<void(RenderTile *, Device *)> unmap_neighbor_tiles;
 
-  uint tile_types;
   DenoiseParams denoising;
   bool denoising_from_render;
   vector<int> denoising_frames;
@@ -128,7 +114,6 @@ class DeviceTask : public Task {
 
   bool need_finish_queue;
   bool integrator_branched;
-  AdaptiveSampling adaptive_sampling;
 
  protected:
   double last_update_time;

intern/cycles/device/opencl/device_opencl.h

@@ -428,10 +428,8 @@ class OpenCLDevice : public Device {
   int mem_sub_ptr_alignment();
 
   void const_copy_to(const char *name, void *host, size_t size);
-  void global_alloc(device_memory &mem);
-  void global_free(device_memory &mem);
-  void tex_alloc(device_texture &mem);
-  void tex_free(device_texture &mem);
+  void tex_alloc(device_memory &mem);
+  void tex_free(device_memory &mem);
 
   size_t global_size_round_up(int group_size, int global_size);
   void enqueue_kernel(cl_kernel kernel,
@@ -447,7 +445,6 @@ class OpenCLDevice : public Device {
                     device_ptr rgba_byte,
                     device_ptr rgba_half);
   void shader(DeviceTask &task);
-  void update_adaptive(DeviceTask &task, RenderTile &tile, int sample);
 
   void denoise(RenderTile &tile, DenoisingTask &denoising);
 

intern/cycles/device/opencl/device_opencl_impl.cpp

@@ -56,11 +56,7 @@ static const string SPLIT_BUNDLE_KERNELS =
     "enqueue_inactive "
     "next_iteration_setup "
     "indirect_subsurface "
-    "buffer_update "
-    "adaptive_stopping "
-    "adaptive_filter_x "
-    "adaptive_filter_y "
-    "adaptive_adjust_samples";
+    "buffer_update";
 
 const string OpenCLDevice::get_opencl_program_name(const string &kernel_name)
 {
@@ -287,10 +283,6 @@ void OpenCLDevice::OpenCLSplitPrograms::load_kernels(
     ADD_SPLIT_KERNEL_BUNDLE_PROGRAM(next_iteration_setup);
     ADD_SPLIT_KERNEL_BUNDLE_PROGRAM(indirect_subsurface);
     ADD_SPLIT_KERNEL_BUNDLE_PROGRAM(buffer_update);
-    ADD_SPLIT_KERNEL_BUNDLE_PROGRAM(adaptive_stopping);
-    ADD_SPLIT_KERNEL_BUNDLE_PROGRAM(adaptive_filter_x);
-    ADD_SPLIT_KERNEL_BUNDLE_PROGRAM(adaptive_filter_y);
-    ADD_SPLIT_KERNEL_BUNDLE_PROGRAM(adaptive_adjust_samples);
     programs.push_back(&program_split);
 
 #  undef ADD_SPLIT_KERNEL_PROGRAM
@@ -613,7 +605,7 @@ OpenCLDevice::OpenCLDevice(DeviceInfo &info, Stats &stats, Profiler &profiler, b
       kernel_programs(this),
       preview_programs(this),
       memory_manager(this),
-      texture_info(this, "__texture_info", MEM_GLOBAL)
+      texture_info(this, "__texture_info", MEM_TEXTURE)
 {
   cpPlatform = NULL;
   cdDevice = NULL;
@@ -945,7 +937,7 @@ void OpenCLDevice::mem_alloc(device_memory &mem)
   cl_mem_flags mem_flag;
   void *mem_ptr = NULL;
 
-  if (mem.type == MEM_READ_ONLY || mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL)
+  if (mem.type == MEM_READ_ONLY || mem.type == MEM_TEXTURE)
     mem_flag = CL_MEM_READ_ONLY;
   else
     mem_flag = CL_MEM_READ_WRITE;
@@ -969,13 +961,9 @@ void OpenCLDevice::mem_alloc(device_memory &mem)
 
 void OpenCLDevice::mem_copy_to(device_memory &mem)
 {
-  if (mem.type == MEM_GLOBAL) {
-    global_free(mem);
-    global_alloc(mem);
-  }
-  else if (mem.type == MEM_TEXTURE) {
-    tex_free((device_texture &)mem);
-    tex_alloc((device_texture &)mem);
+  if (mem.type == MEM_TEXTURE) {
+    tex_free(mem);
+    tex_alloc(mem);
   }
   else {
     if (!mem.device_pointer) {
@@ -1081,11 +1069,8 @@ void OpenCLDevice::mem_zero(device_memory &mem)
 
 void OpenCLDevice::mem_free(device_memory &mem)
 {
-  if (mem.type == MEM_GLOBAL) {
-    global_free(mem);
-  }
-  else if (mem.type == MEM_TEXTURE) {
-    tex_free((device_texture &)mem);
+  if (mem.type == MEM_TEXTURE) {
+    tex_free(mem);
   }
   else {
     if (mem.device_pointer) {
@@ -1108,7 +1093,7 @@ int OpenCLDevice::mem_sub_ptr_alignment()
 device_ptr OpenCLDevice::mem_alloc_sub_ptr(device_memory &mem, int offset, int size)
 {
   cl_mem_flags mem_flag;
-  if (mem.type == MEM_READ_ONLY || mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL)
+  if (mem.type == MEM_READ_ONLY || mem.type == MEM_TEXTURE)
     mem_flag = CL_MEM_READ_ONLY;
   else
     mem_flag = CL_MEM_READ_WRITE;
@@ -1148,9 +1133,9 @@ void OpenCLDevice::const_copy_to(const char *name, void *host, size_t size)
   data->copy_to_device();
 }
 
-void OpenCLDevice::global_alloc(device_memory &mem)
+void OpenCLDevice::tex_alloc(device_memory &mem)
 {
-  VLOG(1) << "Global memory allocate: " << mem.name << ", "
+  VLOG(1) << "Texture allocate: " << mem.name << ", "
           << string_human_readable_number(mem.memory_size()) << " bytes. ("
           << string_human_readable_size(mem.memory_size()) << ")";
 
@@ -1162,7 +1147,7 @@ void OpenCLDevice::global_alloc(device_memory &mem)
   textures_need_update = true;
 }
 
-void OpenCLDevice::global_free(device_memory &mem)
+void OpenCLDevice::tex_free(device_memory &mem)
 {
   if (mem.device_pointer) {
     mem.device_pointer = 0;
@@ -1180,25 +1165,6 @@ void OpenCLDevice::global_free(device_memory &mem)
   }
 }
 
-void OpenCLDevice::tex_alloc(device_texture &mem)
-{
-  VLOG(1) << "Texture allocate: " << mem.name << ", "
-          << string_human_readable_number(mem.memory_size()) << " bytes. ("
-          << string_human_readable_size(mem.memory_size()) << ")";
-
-  memory_manager.alloc(mem.name, mem);
-  /* Set the pointer to non-null to keep code that inspects its value from thinking its
-   * unallocated. */
-  mem.device_pointer = 1;
-  textures[mem.name] = &mem;
-  textures_need_update = true;
-}
-
-void OpenCLDevice::tex_free(device_texture &mem)
-{
-  global_free(mem);
-}
-
 size_t OpenCLDevice::global_size_round_up(int group_size, int global_size)
 {
   int r = global_size % group_size;
@@ -1299,10 +1265,10 @@ void OpenCLDevice::flush_texture_buffers()
 
   foreach (TexturesMap::value_type &tex, textures) {
     string name = tex.first;
-    device_memory *mem = tex.second;
 
-    if (mem->type == MEM_TEXTURE) {
-      const uint id = ((device_texture *)mem)->slot;
+    if (string_startswith(name, "__tex_image")) {
+      int pos = name.rfind("_");
+      int id = atoi(name.data() + pos + 1);
       texture_slots.push_back(texture_slot_t(name, num_data_slots + id));
       num_slots = max(num_slots, num_data_slots + id + 1);
     }
@@ -1315,20 +1281,22 @@ void OpenCLDevice::flush_texture_buffers()
 
   /* Fill in descriptors */
   foreach (texture_slot_t &slot, texture_slots) {
-    device_memory *mem = textures[slot.name];
     TextureInfo &info = texture_info[slot.slot];
 
     MemoryManager::BufferDescriptor desc = memory_manager.get_descriptor(slot.name);
-
-    if (mem->type == MEM_TEXTURE) {
-      info = ((device_texture *)mem)->info;
-    }
-    else {
-      memset(&info, 0, sizeof(TextureInfo));
-    }
-
     info.data = desc.offset;
     info.cl_buffer = desc.device_buffer;
+
+    if (string_startswith(slot.name, "__tex_image")) {
+      device_memory *mem = textures[slot.name];
+
+      info.width = mem->data_width;
+      info.height = mem->data_height;
+      info.depth = mem->data_depth;
+
+      info.interpolation = mem->interpolation;
+      info.extension = mem->extension;
+    }
   }
 
   /* Force write of descriptors. */
@@ -1340,7 +1308,7 @@ void OpenCLDevice::thread_run(DeviceTask *task)
 {
   flush_texture_buffers();
 
-  if (task->type == DeviceTask::RENDER) {
+  if (task->type == DeviceTask::RENDER || task->type == DeviceTask::DENOISE) {
     RenderTile tile;
     DenoisingTask denoising(this, *task);
 
@@ -1349,7 +1317,7 @@ void OpenCLDevice::thread_run(DeviceTask *task)
     kgbuffer.alloc_to_device(1);
 
     /* Keep rendering tiles until done. */
-    while (task->acquire_tile(this, tile, task->tile_types)) {
+    while (task->acquire_tile(this, tile)) {
       if (tile.task == RenderTile::PATH_TRACE) {
         assert(tile.task == RenderTile::PATH_TRACE);
         scoped_timer timer(&tile.buffers->render_time);

intern/cycles/device/opencl/memory_manager.h

@@ -19,8 +19,8 @@
 #include "device/device.h"
 
 #include "util/util_map.h"
-#include "util/util_string.h"
 #include "util/util_vector.h"
+#include "util/util_string.h"
 
 #include "clew.h"
 

intern/cycles/device/opencl/opencl_util.cpp

@@ -16,15 +16,15 @@
 
 #ifdef WITH_OPENCL
 
-#  include "device/device_intern.h"
 #  include "device/opencl/device_opencl.h"
+#  include "device/device_intern.h"
 
 #  include "util/util_debug.h"
 #  include "util/util_logging.h"
 #  include "util/util_md5.h"
 #  include "util/util_path.h"
-#  include "util/util_system.h"
 #  include "util/util_time.h"
+#  include "util/util_system.h"
 
 using std::cerr;
 using std::endl;

intern/cycles/kernel/CMakeLists.txt

@@ -36,10 +36,6 @@ set(SRC_CUDA_KERNELS
 )
 
 set(SRC_OPENCL_KERNELS
-  kernels/opencl/kernel_adaptive_stopping.cl
-  kernels/opencl/kernel_adaptive_filter_x.cl
-  kernels/opencl/kernel_adaptive_filter_y.cl
-  kernels/opencl/kernel_adaptive_adjust_samples.cl
   kernels/opencl/kernel_bake.cl
   kernels/opencl/kernel_base.cl
   kernels/opencl/kernel_displace.cl
@@ -98,7 +94,6 @@ set(SRC_BVH_HEADERS
 
 set(SRC_HEADERS
   kernel_accumulate.h
-  kernel_adaptive_sampling.h
   kernel_bake.h
   kernel_camera.h
   kernel_color.h
@@ -329,10 +324,6 @@ set(SRC_UTIL_HEADERS
 )
 
 set(SRC_SPLIT_HEADERS
-  split/kernel_adaptive_adjust_samples.h
-  split/kernel_adaptive_filter_x.h
-  split/kernel_adaptive_filter_y.h
-  split/kernel_adaptive_stopping.h
   split/kernel_branched.h
   split/kernel_buffer_update.h
   split/kernel_data_init.h

intern/cycles/kernel/bvh/bvh_embree.h

@@ -17,12 +17,9 @@
 #include <embree3/rtcore_ray.h>
 #include <embree3/rtcore_scene.h>
 
-// clang-format off
 #include "kernel/kernel_compat_cpu.h"
 #include "kernel/split/kernel_split_data_types.h"
 #include "kernel/kernel_globals.h"
-// clang-format on
-
 #include "util/util_vector.h"
 
 CCL_NAMESPACE_BEGIN

intern/cycles/kernel/closure/bsdf.h

@@ -14,7 +14,6 @@
  * limitations under the License.
  */
 
-// clang-format off
 #include "kernel/closure/bsdf_ashikhmin_velvet.h"
 #include "kernel/closure/bsdf_diffuse.h"
 #include "kernel/closure/bsdf_oren_nayar.h"
@@ -33,7 +32,6 @@
 #include "kernel/closure/bsdf_principled_sheen.h"
 #include "kernel/closure/bssrdf.h"
 #include "kernel/closure/volume.h"
-// clang-format on
 
 CCL_NAMESPACE_BEGIN
 

intern/cycles/kernel/geom/geom.h

@@ -14,7 +14,6 @@
  * limitations under the License.
  */
 
-// clang-format off
 #include "kernel/geom/geom_attribute.h"
 #include "kernel/geom/geom_object.h"
 #ifdef __PATCH_EVAL__
@@ -31,4 +30,3 @@
 #include "kernel/geom/geom_curve_intersect.h"
 #include "kernel/geom/geom_volume.h"
 #include "kernel/geom/geom_primitive.h"
-// clang-format on

intern/cycles/kernel/geom/geom_object.h

@@ -320,26 +320,6 @@ ccl_device_inline uint object_patch_map_offset(KernelGlobals *kg, int object)
   return kernel_tex_fetch(__objects, object).patch_map_offset;
 }
 
-/* Volume step size */
-
-ccl_device_inline float object_volume_density(KernelGlobals *kg, int object)
-{
-  if (object == OBJECT_NONE) {
-    return 1.0f;
-  }
-
-  return kernel_tex_fetch(__objects, object).surface_area;
-}
-
-ccl_device_inline float object_volume_step_size(KernelGlobals *kg, int object)
-{
-  if (object == OBJECT_NONE) {
-    return kernel_data.background.volume_step_size;
-  }
-
-  return kernel_tex_fetch(__object_volume_step, object);
-}
-
 /* Pass ID for shader */
 
 ccl_device int shader_pass_id(KernelGlobals *kg, const ShaderData *sd)

intern/cycles/kernel/geom/geom_volume.h

@@ -51,14 +51,10 @@ ccl_device float volume_attribute_float(KernelGlobals *kg,
                                         const ShaderData *sd,
                                         const AttributeDescriptor desc)
 {
-  /* todo: optimize this so we don't have to transform both here and in
-   * kernel_tex_image_interp_3d when possible. Also could optimize for the
-   * common case where transform is translation/scale only. */
-  float3 P = sd->P;
-  object_inverse_position_transform(kg, sd, &P);
+  float3 P = volume_normalized_position(kg, sd, sd->P);
   InterpolationType interp = (sd->flag & SD_VOLUME_CUBIC) ? INTERPOLATION_CUBIC :
                                                             INTERPOLATION_NONE;
-  float4 r = kernel_tex_image_interp_3d(kg, desc.offset, P, interp);
+  float4 r = kernel_tex_image_interp_3d(kg, desc.offset, P.x, P.y, P.z, interp);
   return average(float4_to_float3(r));
 }
 
@@ -66,11 +62,10 @@ ccl_device float3 volume_attribute_float3(KernelGlobals *kg,
                                           const ShaderData *sd,
                                           const AttributeDescriptor desc)
 {
-  float3 P = sd->P;
-  object_inverse_position_transform(kg, sd, &P);
+  float3 P = volume_normalized_position(kg, sd, sd->P);
   InterpolationType interp = (sd->flag & SD_VOLUME_CUBIC) ? INTERPOLATION_CUBIC :
                                                             INTERPOLATION_NONE;
-  float4 r = kernel_tex_image_interp_3d(kg, desc.offset, P, interp);
+  float4 r = kernel_tex_image_interp_3d(kg, desc.offset, P.x, P.y, P.z, interp);
 
   if (r.w > 1e-6f && r.w != 1.0f) {
     /* For RGBA colors, unpremultiply after interpolation. */

intern/cycles/kernel/kernel.h

@@ -19,8 +19,8 @@
 
 /* CPU Kernel Interface */
 
-#include "kernel/kernel_types.h"
 #include "util/util_types.h"
+#include "kernel/kernel_types.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -38,7 +38,7 @@ void *kernel_osl_memory(KernelGlobals *kg);
 bool kernel_osl_use(KernelGlobals *kg);
 
 void kernel_const_copy(KernelGlobals *kg, const char *name, void *host, size_t size);
-void kernel_global_memory_copy(KernelGlobals *kg, const char *name, void *mem, size_t size);
+void kernel_tex_copy(KernelGlobals *kg, const char *name, void *mem, size_t size);
 
 #define KERNEL_ARCH cpu
 #include "kernel/kernels/cpu/kernel_cpu.h"