remove unnecessary mutex

.clang-format

@@ -191,7 +191,6 @@ ForEachMacros:
   - FOREACH_MAIN_ID_BEGIN
   - FOREACH_MAIN_LISTBASE_BEGIN
   - FOREACH_MAIN_LISTBASE_ID_BEGIN
-  - FOREACH_MESH_BUFFER_CACHE
   - FOREACH_NODETREE_BEGIN
   - FOREACH_OBJECT_BEGIN
   - FOREACH_OBJECT_FLAG_BEGIN
@@ -214,7 +213,6 @@ ForEachMacros:
   - GHASH_ITER_INDEX
   - GPU_SELECT_LOAD_IF_PICKSEL_LIST
   - GP_EDITABLE_STROKES_BEGIN
-  - GP_EVALUATED_STROKES_BEGIN
   - GSET_FOREACH_BEGIN
   - GSET_ITER
   - GSET_ITER_INDEX
@@ -238,6 +236,7 @@ ForEachMacros:
   - LISTBASE_FOREACH_BACKWARD
   - LISTBASE_FOREACH_MUTABLE
   - LISTBASE_FOREACH_BACKWARD_MUTABLE
+  - MAN2D_ITER_AXES_BEGIN
   - MAN_ITER_AXES_BEGIN
   - NODE_INSTANCE_HASH_ITER
   - NODE_SOCKET_TYPES_BEGIN
@@ -245,16 +244,12 @@ ForEachMacros:
   - NODE_TYPES_BEGIN
   - PIXEL_LOOPER_BEGIN
   - PIXEL_LOOPER_BEGIN_CHANNELS
-  - RENDER_PASS_ITER_BEGIN
   - RNA_BEGIN
   - RNA_PROP_BEGIN
   - RNA_STRUCT_BEGIN
   - RNA_STRUCT_BEGIN_SKIP_RNA_TYPE
-  - SCULPT_VERTEX_DUPLICATES_AND_NEIGHBORS_ITER_BEGIN
-  - SCULPT_VERTEX_NEIGHBORS_ITER_BEGIN
   - SEQP_BEGIN
   - SEQ_BEGIN
-  - SURFACE_QUAD_ITER_BEGIN
   - foreach
 
 # Use once we bump the minimum version to version 8.

CMakeLists.txt

@@ -323,10 +323,6 @@ option(WITH_FREESTYLE     "Enable Freestyle (advanced edges rendering)" ON)
 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)
 
-# New simulation data block
-option(WITH_NEW_SIMULATION_TYPE "Enable simulation data block (use for development only, don't save in files)" OFF)
-mark_as_advanced(WITH_NEW_SIMULATION_TYPE)
-
 # Misc
 if(WIN32)
   option(WITH_INPUT_IME "Enable Input Method Editor (IME) for complex Asian character input" ON)
@@ -637,10 +633,9 @@ set_and_warn_dependency(WITH_BOOST WITH_OPENVDB        OFF)
 set_and_warn_dependency(WITH_BOOST WITH_OPENCOLORIO    OFF)
 set_and_warn_dependency(WITH_BOOST WITH_QUADRIFLOW     OFF)
 set_and_warn_dependency(WITH_BOOST WITH_USD            OFF)
-set_and_warn_dependency(WITH_BOOST WITH_ALEMBIC        OFF)
 
 if(WITH_BOOST AND NOT (WITH_CYCLES OR WITH_OPENIMAGEIO OR WITH_INTERNATIONAL OR
-   WITH_OPENVDB OR WITH_OPENCOLORIO OR WITH_USD OR WITH_ALEMBIC))
+   WITH_OPENVDB OR WITH_OPENCOLORIO OR WITH_USD))
   message(STATUS "No dependencies need 'WITH_BOOST' forcing WITH_BOOST=OFF")
   set(WITH_BOOST OFF)
 endif()

build_files/build_environment/cmake/versions.cmake

@@ -146,8 +146,8 @@ set(PYTHON_URI https://www.python.org/ftp/python/${PYTHON_VERSION}/Python-${PYTH
 set(PYTHON_HASH d33e4aae66097051c2eca45ee3604803)
 
 set(TBB_VERSION 2019_U9)
-set(TBB_URI https://github.com/oneapi-src/oneTBB/archive/${TBB_VERSION}.tar.gz)
-set(TBB_HASH 26263622e9187212ec240dcf01b66207)
+set(TBB_URI https://github.com/01org/tbb/archive/${TBB_VERSION}.tar.gz)
+set(TBB_HASH 584edbec127c508f2cd5b6e79ad200fc)
 
 set(OPENVDB_VERSION 7.0.0)
 set(OPENVDB_URI https://github.com/dreamworksanimation/openvdb/archive/v${OPENVDB_VERSION}.tar.gz)

build_files/build_environment/cmake/xr_openxr.cmake

@@ -29,7 +29,6 @@ if(UNIX AND NOT APPLE)
     -DBUILD_WITH_WAYLAND_HEADERS=OFF
     -DBUILD_WITH_XCB_HEADERS=OFF
     -DBUILD_WITH_XLIB_HEADERS=ON
-    -DBUILD_WITH_SYSTEM_JSONCPP=OFF
     -DCMAKE_CXX_FLAGS=-DDISABLE_STD_FILESYSTEM=1
   )
 endif()

build_files/build_environment/install_deps.sh

@@ -63,8 +63,8 @@ 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, \
+skip-xr-openxr\
 -- "$@" \
 )
 
@@ -1289,7 +1289,7 @@ compile_Python() {
 
     ./configure --prefix=$_inst --libdir=$_inst/lib --enable-ipv6 \
         --enable-loadable-sqlite-extensions --with-dbmliborder=bdb \
-        --with-computed-gotos --with-pymalloc --enable-shared
+        --with-computed-gotos --with-pymalloc
 
     make -j$THREADS && make install
     make clean
@@ -1310,8 +1310,6 @@ compile_Python() {
     INFO "Own Python-$PYTHON_VERSION is up to date, nothing to do!"
     INFO "If you want to force rebuild of this lib, use the --force-python option."
   fi
-
-  run_ldconfig "python-$PYTHON_VERSION_MIN"
 }
 
 # ----------------------------------------------------------------------------
@@ -3130,7 +3128,7 @@ compile_XR_OpenXR_SDK() {
   fi
 
   # To be changed each time we make edits that would modify the compiled result!
-  xr_openxr_magic=2
+  xr_openxr_magic=1
   _init_xr_openxr_sdk
 
   # Clean install if needed!
@@ -3187,7 +3185,6 @@ compile_XR_OpenXR_SDK() {
     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_d="$cmake_d -D BUILD_WITH_SYSTEM_JSONCPP=OFF"
 
     cmake $cmake_d "-DCMAKE_CXX_FLAGS=-DDISABLE_STD_FILESYSTEM=1" ..
 
@@ -5191,7 +5188,7 @@ print_info() {
   PRINT ""
   PRINT "If you're using CMake add this to your configuration flags:"
 
-  _buildargs="-U *SNDFILE* -U PYTHON* -U *BOOST* -U *Boost*"
+  _buildargs="-U *SNDFILE* -U *PYTHON* -U *BOOST* -U *Boost*"
   _buildargs="$_buildargs -U *OPENCOLORIO* -U *OPENEXR* -U *OPENIMAGEIO* -U *LLVM* -U *CYCLES*"
   _buildargs="$_buildargs -U *OPENSUBDIV* -U *OPENVDB* -U *COLLADA* -U *FFMPEG* -U *ALEMBIC* -U *USD*"
 

build_files/buildbot/config/blender_linux.cmake

@@ -7,6 +7,9 @@ message(STATUS "Building in CentOS 7 64bit environment")
 set(LIBDIR_NAME "linux_centos7_x86_64")
 set(WITH_CXX11_ABI           OFF CACHE BOOL "" FORCE)
 
+# Default to only build Blender
+set(WITH_BLENDER             ON  CACHE BOOL "" FORCE)
+
 # ######## Linux-specific build options ########
 # Options which are specific to Linux-only platforms
 
@@ -17,6 +20,12 @@ set(WITH_DOC_MANPAGE         OFF CACHE BOOL "" FORCE)
 
 set(WITH_JACK_DYNLOAD        ON  CACHE BOOL "" FORCE)
 set(WITH_SDL_DYNLOAD         ON  CACHE BOOL "" FORCE)
+set(WITH_SYSTEM_GLEW         OFF CACHE BOOL "" FORCE)
+
+set(WITH_OPENMP_STATIC       ON  CACHE BOOL "" FORCE)
+
+set(WITH_PYTHON_INSTALL_NUMPY    ON CACHE BOOL "" FORCE)
+set(WITH_PYTHON_INSTALL_REQUESTS ON CACHE BOOL "" FORCE)
 
 # ######## Release environment specific settings ########
 
@@ -24,5 +33,13 @@ set(LIBDIR "${CMAKE_CURRENT_LIST_DIR}/../../../../lib/${LIBDIR_NAME}" CACHE STRI
 
 # Platform specific configuration, to ensure static linking against everything.
 
+set(Boost_USE_STATIC_LIBS    ON CACHE BOOL "" FORCE)
+
+# We need to link OpenCOLLADA against PCRE library. Even though it is not installed
+# on /usr, we do not really care -- all we care is PCRE_FOUND be TRUE and its
+# library pointing to a valid one.
+set(PCRE_INCLUDE_DIR          "/usr/include"                        CACHE STRING "" FORCE)
+set(PCRE_LIBRARY              "${LIBDIR}/opencollada/lib/libpcre.a" CACHE STRING "" FORCE)
+
 # Additional linking libraries
 set(CMAKE_EXE_LINKER_FLAGS   "-lrt -static-libstdc++ -no-pie"  CACHE STRING "" FORCE)

build_files/cmake/Modules/FindOpenImageDenoise.cmake

@@ -95,7 +95,7 @@ FIND_LIBRARY(OPENIMAGEDENOISE_LIBRARY
 # handle the QUIETLY and REQUIRED arguments and set OPENIMAGEDENOISE_FOUND to TRUE if
 # all listed variables are TRUE
 INCLUDE(FindPackageHandleStandardArgs)
-FIND_PACKAGE_HANDLE_STANDARD_ARGS(OpenImageDenoise DEFAULT_MSG
+FIND_PACKAGE_HANDLE_STANDARD_ARGS(OPENIMAGEDENOISE DEFAULT_MSG
     OPENIMAGEDENOISE_LIBRARY OPENIMAGEDENOISE_INCLUDE_DIR)
 
 IF(OPENIMAGEDENOISE_FOUND)

build_files/cmake/platform/platform_apple.cmake

@@ -222,10 +222,12 @@ if(WITH_OPENCOLLADA)
     -lMathMLSolver
     -lGeneratedSaxParser
     -lbuffer -lftoa -lUTF
+    ${OPENCOLLADA_LIBPATH}/libxml2.a
   )
-  # PCRE and XML2 are bundled with OpenCollada.
+  # PCRE is bundled with openCollada
+  # set(PCRE ${LIBDIR}/pcre)
+  # set(PCRE_LIBPATH ${PCRE}/lib)
   set(PCRE_LIBRARIES pcre)
-  set(XML2_LIBRARIES xml2)
 endif()
 
 if(WITH_SDL)
@@ -447,9 +449,7 @@ if(${XCODE_VERSION} VERSION_EQUAL 5 OR ${XCODE_VERSION} VERSION_GREATER 5)
   # Xcode 5 is always using CLANG, which has too low template depth of 128 for libmv
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -ftemplate-depth=1024")
 endif()
-
-# Avoid conflicts with Luxrender, and other plug-ins that may use the same
-# libraries as Blender with a different version or build options.
+# Get rid of eventually clashes, we export some symbols explicitly as local
 set(PLATFORM_LINKFLAGS
   "${PLATFORM_LINKFLAGS} -Xlinker -unexported_symbols_list -Xlinker '${CMAKE_SOURCE_DIR}/source/creator/osx_locals.map'"
 )

build_files/cmake/platform/platform_apple_xcode.cmake

@@ -20,6 +20,10 @@
 
 # Xcode and system configuration for Apple.
 
+# require newer cmake on osx because of version handling,
+# older cmake cannot handle 2 digit subversion!
+cmake_minimum_required(VERSION 3.0.0)
+
 if(NOT CMAKE_OSX_ARCHITECTURES)
   set(CMAKE_OSX_ARCHITECTURES x86_64 CACHE STRING
     "Choose the architecture you want to build Blender for: i386, x86_64 or ppc"
@@ -41,98 +45,54 @@ execute_process(
     OUTPUT_VARIABLE XCODE_CHECK OUTPUT_STRIP_TRAILING_WHITESPACE)
 string(REPLACE "/Contents/Developer" "" XCODE_BUNDLE ${XCODE_CHECK}) # truncate to bundlepath in any case
 
-if(NOT ${CMAKE_GENERATOR} MATCHES "Xcode")
-  # Unix makefile generator does not fill XCODE_VERSION var, so we get it with a command.
-  # Note that `xcodebuild -version` gives output in two lines: first line will include
-  # Xcode version, second one will include build number. We are only interested in the
-  # former one. Here is an example of the output:
-  #   Xcode 11.4
-  #   Build version 11E146
-  # The expected XCODE_VERSION in this case is 11.4.
+if(${CMAKE_GENERATOR} MATCHES "Xcode")
 
+  # earlier xcode has no bundled developer dir, no sense in getting xcode path from
+  if(${XCODE_VERSION} VERSION_GREATER 4.2)
+    # reduce to XCode name without dp extension
+    string(SUBSTRING "${XCODE_CHECK}" 14 6 DP_NAME)
+    if(${DP_NAME} MATCHES Xcode5)
+      set(XCODE_VERSION 5)
+    endif()
+  endif()
+
+  ##### cmake incompatibility with xcode  4.3 and higher #####
+  if(${XCODE_VERSION} MATCHES '') # cmake fails due looking for xcode in the wrong path, thus will be empty var
+    message(FATAL_ERROR "Xcode 4.3 and higher must be used with cmake 2.8-8 or higher")
+  endif()
+  ### end cmake incompatibility with xcode 4.3 and higher ###
+
+  if(${XCODE_VERSION} VERSION_EQUAL 4 OR ${XCODE_VERSION} VERSION_GREATER 4 AND ${XCODE_VERSION} VERSION_LESS 4.3)
+    # Xcode 4 defaults to the Apple LLVM Compiler.
+    # Override the default compiler selection because Blender only compiles with gcc up to xcode 4.2
+    set(CMAKE_XCODE_ATTRIBUTE_GCC_VERSION "com.apple.compilers.llvmgcc42")
+    message(STATUS "Setting compiler to: " ${CMAKE_XCODE_ATTRIBUTE_GCC_VERSION})
+  endif()
+else() # unix makefile generator does not fill XCODE_VERSION var, so we get it with a command
   execute_process(COMMAND xcodebuild -version OUTPUT_VARIABLE XCODE_VERS_BUILD_NR)
-
-  # Convert output to a single line by replacling newlines with spaces.
-  # This is needed because regex replace can not operate through the newline character
-  # and applies substitutions for each individual lines.
-  string(REPLACE "\n" " " XCODE_VERS_BUILD_NR_SINGLE_LINE "${XCODE_VERS_BUILD_NR}")
-
-  string(REGEX REPLACE "(.*)Xcode ([0-9\\.]+).*" "\\2" XCODE_VERSION "${XCODE_VERS_BUILD_NR_SINGLE_LINE}")
-
+  string(SUBSTRING "${XCODE_VERS_BUILD_NR}" 6 3 XCODE_VERSION) # truncate away build-nr
   unset(XCODE_VERS_BUILD_NR)
-  unset(XCODE_VERS_BUILD_NR_SINGLE_LINE)
 endif()
 
 message(STATUS "Detected OS X ${OSX_SYSTEM} and Xcode ${XCODE_VERSION} at ${XCODE_BUNDLE}")
 
-# Older Xcode versions had different approach to the directory hiearchy.
-# Require newer Xcode which is also have better chances of being able to compile with the
-# required deployment target.
-#
-# NOTE: Xcode version 8.2 is the latest one which runs on macOS 10.11.
-if(${XCODE_VERSION} VERSION_LESS 8.2)
-  message(FATAL_ERROR "Only Xcode version 8.2 and newer is supported")
-endif()
-
-# note: xcode-select path could be ambiguous,
-# cause /Applications/Xcode.app/Contents/Developer or /Applications/Xcode.app would be allowed
-# so i use a selfcomposed bundlepath here
-set(OSX_SYSROOT_PREFIX ${XCODE_BUNDLE}/Contents/Developer/Platforms/MacOSX.platform)
-message(STATUS "OSX_SYSROOT_PREFIX: " ${OSX_SYSROOT_PREFIX})
-
-# Collect list of OSX system versions which will be used to detect path to corresponding SDK.
-# Start with macOS SDK version reported by xcodebuild and include possible extra ones.
-#
-# The reason for need of extra ones is because it's possible that xcodebuild will report
-# SDK version in the full manner (aka major.minor.patch), but the actual path will only
-# include major.minor.
-#
-# This happens, for example, on macOS Catalina 10.15.4 and Xcode 11.4: xcodebuild on this
-# system outputs "10.15.4", but the actual SDK path is MacOSX10.15.sdk.
-#
-# This should be safe from picking wrong SDK version because (a) xcodebuild reports full semantic
-# SDK version, so such SDK does exist on the system. And if it doesn't exist with full version
-# in the path, what SDK is in the major.minor folder then.
-set(OSX_SDK_TEST_VERSIONS ${OSX_SYSTEM})
-if(OSX_SYSTEM MATCHES "([0-9]+)\\.([0-9]+)\\.([0-9]+)")
-  string(REGEX REPLACE "([0-9]+)\\.([0-9]+)\\.([0-9]+)" "\\1.\\2" OSX_SYSTEM_NO_PATCH "${OSX_SYSTEM}")
-  list(APPEND OSX_SDK_TEST_VERSIONS ${OSX_SYSTEM_NO_PATCH})
-  unset(OSX_SYSTEM_NO_PATCH)
-endif()
-
-# Loop through all possible versions and pick the first one which resolves to a valid SDK path.
-set(OSX_SDK_PATH)
-set(OSX_SDK_FOUND FALSE)
-set(OSX_SDK_PREFIX ${OSX_SYSROOT_PREFIX}/Developer/SDKs)
-set(OSX_SDKROOT)
-foreach(OSX_SDK_VERSION ${OSX_SDK_TEST_VERSIONS})
-  set(CURRENT_OSX_SDK_PATH "${OSX_SDK_PREFIX}/MacOSX${OSX_SDK_VERSION}.sdk")
-  if(EXISTS ${CURRENT_OSX_SDK_PATH})
-    set(OSX_SDK_PATH "${CURRENT_OSX_SDK_PATH}")
-    set(OSX_SDKROOT macosx${OSX_SDK_VERSION})
-    set(OSX_SDK_FOUND TRUE)
-    break()
+if(${XCODE_VERSION} VERSION_LESS 4.3)
+  # use guaranteed existing sdk
+  set(CMAKE_OSX_SYSROOT /Developer/SDKs/MacOSX${OSX_SYSTEM}.sdk CACHE PATH "" FORCE)
+else()
+  # note: xcode-select path could be ambiguous,
+  # cause /Applications/Xcode.app/Contents/Developer or /Applications/Xcode.app would be allowed
+  # so i use a selfcomposed bundlepath here
+  set(OSX_SYSROOT_PREFIX ${XCODE_BUNDLE}/Contents/Developer/Platforms/MacOSX.platform)
+  message(STATUS "OSX_SYSROOT_PREFIX: " ${OSX_SYSROOT_PREFIX})
+  set(OSX_DEVELOPER_PREFIX /Developer/SDKs/MacOSX${OSX_SYSTEM}.sdk) # use guaranteed existing sdk
+  set(CMAKE_OSX_SYSROOT ${OSX_SYSROOT_PREFIX}/${OSX_DEVELOPER_PREFIX} CACHE PATH "" FORCE)
+  if(${CMAKE_GENERATOR} MATCHES "Xcode")
+    # to silence sdk not found warning, just overrides CMAKE_OSX_SYSROOT
+    set(CMAKE_XCODE_ATTRIBUTE_SDKROOT macosx${OSX_SYSTEM})
   endif()
-endforeach()
-unset(OSX_SDK_PREFIX)
-unset(OSX_SDK_TEST_VERSIONS)
-
-if(NOT OSX_SDK_FOUND)
-  message(FATAL_ERROR "Unable to find SDK for macOS version ${OSX_SYSTEM}")
 endif()
 
-message(STATUS "Detected OSX_SYSROOT: ${OSX_SDK_PATH}")
-
-set(CMAKE_OSX_SYSROOT ${OSX_SDK_PATH} CACHE PATH "" FORCE)
-unset(OSX_SDK_PATH)
-unset(OSX_SDK_FOUND)
-
-if(${CMAKE_GENERATOR} MATCHES "Xcode")
-  # to silence sdk not found warning, just overrides CMAKE_OSX_SYSROOT
-  set(CMAKE_XCODE_ATTRIBUTE_SDKROOT ${OSX_SDKROOT})
-endif()
-unset(OSX_SDKROOT)
-
 # 10.11 is our min. target, if you use higher sdk, weak linking happens
 if(CMAKE_OSX_DEPLOYMENT_TARGET)
   if(${CMAKE_OSX_DEPLOYMENT_TARGET} VERSION_LESS 10.11)

build_files/cmake/platform/platform_unix.cmake

@@ -195,14 +195,8 @@ endif()
 if(WITH_OPENCOLLADA)
   find_package_wrapper(OpenCOLLADA)
   if(OPENCOLLADA_FOUND)
-    if(WITH_STATIC_LIBS)
-      # PCRE is bundled with OpenCollada without headers, so can't use
-      # find_package reliably to detect it.
-      set(PCRE_LIBRARIES ${LIBDIR}/opencollada/lib/libpcre.a)
-    else()
-      find_package_wrapper(PCRE)
-    endif()
     find_package_wrapper(XML2)
+    find_package_wrapper(PCRE)
   else()
     set(WITH_OPENCOLLADA OFF)
   endif()
@@ -411,6 +405,13 @@ if(WITH_LLVM)
   endif()
 endif()
 
+if(WITH_LLVM OR WITH_SDL_DYNLOAD)
+  # Fix for conflict with Mesa llvmpipe
+  set(PLATFORM_LINKFLAGS
+    "${PLATFORM_LINKFLAGS} -Wl,--version-script='${CMAKE_SOURCE_DIR}/source/creator/blender.map'"
+  )
+endif()
+
 if(WITH_OPENSUBDIV)
   find_package_wrapper(OpenSubdiv)
 
@@ -600,16 +601,3 @@ elseif(CMAKE_C_COMPILER_ID MATCHES "Intel")
   set(PLATFORM_CFLAGS "-pipe -fPIC -funsigned-char -fno-strict-aliasing")
   set(PLATFORM_LINKFLAGS "${PLATFORM_LINKFLAGS} -static-intel")
 endif()
-
-# Avoid conflicts with Mesa llvmpipe, Luxrender, and other plug-ins that may
-# use the same libraries as Blender with a different version or build options.
-set(PLATFORM_LINKFLAGS
-  "${PLATFORM_LINKFLAGS} -Wl,--version-script='${CMAKE_SOURCE_DIR}/source/creator/blender.map'"
-)
-
-# Don't use position independent executable for portable install since file
-# browsers can't properly detect blender as an executable then. Still enabled
-# for non-portable installs as typically used by Linux distributions.
-if(WITH_INSTALL_PORTABLE)
-  set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -no-pie")
-endif()

build_files/cmake/platform/platform_win32.cmake

@@ -669,7 +669,7 @@ if(WITH_USD)
     set(USD_INCLUDE_DIRS ${LIBDIR}/usd/include)
     set(USD_RELEASE_LIB ${LIBDIR}/usd/lib/libusd_m.lib)
     set(USD_DEBUG_LIB ${LIBDIR}/usd/lib/libusd_m_d.lib)
-    set(USD_LIBRARY_DIR ${LIBDIR}/usd/lib)
+    set(USD_LIBRARY_DIR ${LIBDIR}/usd/lib/usd)
     set(USD_LIBRARIES
       debug ${USD_DEBUG_LIB}
       optimized ${USD_RELEASE_LIB}

build_files/windows/autodetect_msvc.cmd

@@ -6,6 +6,9 @@ if %ERRORLEVEL% EQU 0 goto DetectionComplete
 call "%~dp0\detect_msvc2019.cmd"
 if %ERRORLEVEL% EQU 0 goto DetectionComplete
 
+call "%~dp0\detect_msvc2015.cmd"
+if %ERRORLEVEL% EQU 0 goto DetectionComplete
+
 echo Compiler Detection failed. Use verbose switch for more information. 
 exit /b 1
 

build_files/windows/detect_msvc2015.cmd

@@ -0,0 +1,3 @@
+set BUILD_VS_VER=14
+set BUILD_VS_YEAR=2015
+call "%~dp0\detect_msvc_classic.cmd"

build_files/windows/detect_msvc_classic.cmd

@@ -0,0 +1,69 @@
+if NOT "%verbose%" == "" (
+	echo Detecting msvc %BUILD_VS_YEAR%
+)
+set KEY_NAME="HKEY_LOCAL_MACHINE\SOFTWARE\Wow6432Node\Microsoft\VisualStudio\%BUILD_VS_VER%.0\Setup\VC"
+for /F "usebackq skip=2 tokens=1-2*" %%A IN (`REG QUERY %KEY_NAME% /v ProductDir 2^>nul`) DO set MSVC_VC_DIR=%%C
+if DEFINED MSVC_VC_DIR (
+	if NOT "%verbose%" == "" (
+		echo Visual Studio %BUILD_VS_YEAR% on Win64 detected at "%MSVC_VC_DIR%"
+	)
+	goto msvc_detect_finally
+)
+
+REM Check 32 bits
+set KEY_NAME="HKEY_LOCAL_MACHINE\SOFTWARE\Microsoft\VisualStudio\%BUILD_VS_VER%.0\Setup\VC"
+for /F "usebackq skip=2 tokens=1-2*" %%A IN (`REG QUERY %KEY_NAME% /v ProductDir 2^>nul`) DO set MSVC_VC_DIR=%%C
+if DEFINED MSVC_VC_DIR (
+	if NOT "%verbose%" == "" (
+		echo Visual Studio %BUILD_VS_YEAR% on Win32 detected at "%MSVC_VC_DIR%"
+	)
+	goto msvc_detect_finally
+)
+if NOT "%verbose%" == "" (
+	echo Visual Studio %BUILD_VS_YEAR% not found. 
+)
+goto FAIL
+:msvc_detect_finally
+set VCVARS=%MSVC_VC_DIR%\vcvarsall.bat
+if not exist "%VCVARS%" (
+	echo "%VCVARS%" not found.
+	goto FAIL
+)
+
+call "%vcvars%" %BUILD_ARCH%
+
+rem try msbuild
+msbuild /version > NUL 
+if errorlevel 1 (
+	if NOT "%verbose%" == "" (
+		echo Visual Studio %BUILD_VS_YEAR% msbuild not found
+	)
+	goto FAIL
+)
+
+if NOT "%verbose%" == "" (
+		echo Visual Studio %BUILD_VS_YEAR% msbuild found 
+)
+
+REM try the c++ compiler
+cl 2> NUL 1>&2
+if errorlevel 1 (
+	if NOT "%verbose%" == "" (
+		echo Visual Studio %BUILD_VS_YEAR% C/C++ Compiler not found
+	)
+	goto FAIL
+)
+
+if NOT "%verbose%" == "" (
+		echo Visual Studio %BUILD_VS_YEAR% C/C++ Compiler found
+)
+goto DetectionComplete
+
+:FAIL
+exit /b 1
+
+:DetectionComplete
+if NOT "%verbose%" == "" (
+		echo Visual Studio %BUILD_VS_YEAR% Detected successfully 
+)
+exit /b 0

build_files/windows/detect_msvc_vswhere.cmd

@@ -27,13 +27,7 @@ if NOT "%verbose%" == "" (
 
 if "%VS_InstallDir%"=="" (
 	if NOT "%verbose%" == "" (
-		echo.
-		echo Visual Studio is detected but no suitable installation was found. 
-		echo.
-		echo Check the "Desktop development with C++" workload has been installed. 
-		echo. 
-		echo If you are attempting to use either Visual Studio Preview version or the Visual C++ Build tools, Please see 'make help' on how to opt in to those toolsets.
-		echo. 
+		echo Visual Studio is detected but the "Desktop development with C++" workload has not been instlled
 		goto FAIL
 	)
 )

build_files/windows/parse_arguments.cmd

@@ -66,6 +66,8 @@ if NOT "%1" == "" (
 	) else if "%1" == "2019b" (
 		set BUILD_VS_YEAR=2019
 		set VSWHERE_ARGS=-products Microsoft.VisualStudio.Product.BuildTools
+	) else if "%1" == "2015" (
+		set BUILD_VS_YEAR=2015
 	) else if "%1" == "packagename" (
 		set BUILD_CMAKE_ARGS=%BUILD_CMAKE_ARGS% -DCPACK_OVERRIDE_PACKAGENAME="%2"
 		shift /1

build_files/windows/show_help.cmd

@@ -27,13 +27,11 @@ echo - builddir [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^)
-echo - 2019 ^(build with visual studio 2019^)
-echo - 2019pre ^(build with visual studio 2019 pre-release^)
-echo - 2019b ^(build with visual studio 2019 Build Tools^)
 
 echo.
 echo Experimental options
 echo - with_opengl_tests ^(enable both the render and draw opengl test suites^)
+echo - 2015 ^(build with visual studio 2015^)
 echo - clang ^(enable building with clang^)
 echo - asan ^(enable asan when building with clang^)
 echo - ninja ^(enable building with ninja instead of msbuild^)

doc/python_api/examples/aud.py

@@ -2,20 +2,20 @@
 Basic Sound Playback
 ++++++++++++++++++++
 
-This script shows how to use the classes: :class:`Device`, :class:`Sound` and
+This script shows how to use the classes: :class:`Device`, :class:`Factory` and
 :class:`Handle`.
 """
 import aud
 
-device = aud.Device()
+device = aud.device()
 # load sound file (it can be a video file with audio)
-sound = aud.Sound('music.ogg')
+factory = aud.Factory('music.ogg')
 
 # play the audio, this return a handle to control play/pause
-handle = device.play(sound)
+handle = device.play(factory)
 # if the audio is not too big and will be used often you can buffer it
-sound_buffered = aud.Sound.buffer(sound)
-handle_buffered = device.play(sound_buffered)
+factory_buffered = aud.Factory.buffer(factory)
+handle_buffered = device.play(factory_buffered)
 
 # stop the sounds (otherwise they play until their ends)
 handle.stop()

doc/python_api/rst/info_gotcha.rst

@@ -205,15 +205,15 @@ Support Overview
 
    * - Usage
      - :class:`bpy.types.MeshPolygon`
-     - :class:`bpy.types.MeshLoopTriangle`
+     - :class:`bpy.types.MeshTessFace`
      - :class:`bmesh.types.BMFace`
    * - Import/Create
      - Poor *(inflexible)*
-     - Unusable *(read-only)*.
+     - Good *(supported as upgrade path)*
      - Best
    * - Manipulate
      - Poor *(inflexible)*
-     - Unusable *(read-only)*.
+     - Poor *(loses ngons)*
      - Best
    * - Export/Output
      - Good *(ngon support)*

doc/python_api/sphinx_doc_gen.py

@@ -1052,7 +1052,6 @@ context_type_map = {
     "selected_editable_fcurves": ("FCurve", True),
     "selected_editable_objects": ("Object", True),
     "selected_editable_sequences": ("Sequence", True),
-    "selected_nla_strips": ("NlaStrip", True),
     "selected_nodes": ("Node", True),
     "selected_objects": ("Object", True),
     "selected_pose_bones": ("PoseBone", True),
@@ -1074,7 +1073,6 @@ context_type_map = {
     "visible_pose_bones": ("PoseBone", True),
     "visible_fcurves": ("FCurve", True),
     "weight_paint_object": ("Object", False),
-    "volume": ("Volume", False),
     "world": ("World", False),
 }
 

extern/audaspace/bindings/doc/tutorials.rst

@@ -4,51 +4,35 @@ Tutorials
 Introduction
 ------------
 
-The C and Python binding for audaspace were designed with simplicity in mind.
-This means however that to use the full capabilities of audaspace,
-there is no way around the C++ library.
+The C and Python binding for audaspace were designed with simplicity in mind. This means however that to use the full capabilities of audaspace, there is no way around the C++ library.
 
 Simple Demo
 -----------
 
-The **simple.py** example program contains all the basic
-building blocks for an application using audaspace.
-These building blocks are basically the classes :class:`aud.Device`,
-:class:`aud.Sound` and :class:`aud.Handle`.
+The **simple.py** example program contains all the basic building blocks for an application using audaspace. These building blocks are basically the classes :class:`aud.Device`, :class:`aud.Sound` and :class:`aud.Handle`.
 
-We start with importing :mod:`aud` and :mod:`time`
-as the modules we need for our simple example.
+We start with importing :mod:`aud` and :mod:`time` as the modules we need for our simple example.
 
 .. code-block:: python
 
    #!/usr/bin/python
    import aud, time
 
-The first step now is to open an output device and this
-can simply be done by allocating a :class:`aud.Device` object.
+The first step now is to open an output device and this can simply be done by allocating a :class:`aud.Device` object.
 
 .. code-block:: python
 
    device = aud.Device()
 
-To create a sound we can choose to load one from a :func:`aud.Sound.file`,
-or we use one of our signal generators. We decide to do the latter
-and create a :func:`aud.Sound.sine` signal with a frequency of 440 Hz.
+To create a sound we can choose to load one from a :func:`aud.Sound.file`, or we use one of our signal generators. We decide to do the latter and create a :func:`aud.Sound.sine` signal with a frequency of 440 Hz.
 
 .. code-block:: python
 
    sine = aud.Sound.sine(440)
 
-.. note:: At this point nothing is playing back yet,
-:class:`aud.Sound` objects are just descriptions of sounds.
+.. note:: At this point nothing is playing back yet, :class:`aud.Sound` objects are just descriptions of sounds.
 
-However instead of a sine wave, we would like to have a square wave
-to produce a more retro gaming sound. We could of course use the
-:func:`aud.Sound.square` generator instead of sine,
-but we want to show how to apply effects,
-so we apply a :func:`aud.Sound.threshold`
-which makes a square wave out of our sine too,
-even if less efficient than directly generating the square wave.
+However instead of a sine wave, we would like to have a square wave to produce a more retro gaming sound. We could of course use the :func:`aud.Sound.square` generator instead of sine, but we want to show how to apply effects, so we apply a :func:`aud.Sound.threshold` which makes a square wave out of our sine too, even if less efficient than directly generating the square wave.
 
 .. code-block:: python
 
@@ -56,19 +40,13 @@ even if less efficient than directly generating the square wave.
 
 .. note:: The :class:`aud.Sound` class offers generator and effect functions.
 
-The we can play our sound by calling the
-:func:`aud.Device.play` method of our device.
-This method returns a :class:`aud.Handle`
-which is used to control the playback of the sound.
+The we can play our sound by calling the :func:`aud.Device.play` method of our device. This method returns a :class:`aud.Handle` which is used to control the playback of the sound.
 
 .. code-block:: python
 
    handle = device.play(square)
 
-Now if we do nothing else anymore the application will quit immediately,
-so we won't hear much of our square wave,
-so we decide to wait for three seconds before
-quitting the application by calling :func:`time.sleep`.
+Now if we do nothing else anymore the application will quit immediately, so we won't hear much of our square wave, so we decide to wait for three seconds before quitting the application by calling :func:`time.sleep`.
 
 .. code-block:: python
 
@@ -77,47 +55,29 @@ quitting the application by calling :func:`time.sleep`.
 Audioplayer
 -----------
 
-Now that we know the basics of audaspace,
-we can build our own music player easily
-by just slightly changing the previous program.
-The **player.py** example does exactly that,
-let's have a short look at the differences:
+Now that we know the basics of audaspace, we can build our own music player easily by just slightly changing the previous program. The **player.py** example does exactly that, let's have a short look at the differences:
 
-Instead of creating a sine signal and thresholding it,
-we in fact use the :func:`aud.Sound.file` function to load a sound from a file.
-The filename we pass is the first command line argument our application got.
+Instead of creating a sine signal and thresholding it, we in fact use the :func:`aud.Sound.file` function to load a sound from a file. The filename we pass is the first command line argument our application got.
 
 .. code-block:: python
 
    sound = aud.Sound.file(sys.argv[1])
 
-When the sound gets played back we now want to wait until
-the whole file has been played, so we use the :data:`aud.Handle.status`
-property to determine whether the sound finished playing.
+When the sound gets played back we now want to wait until the whole file has been played, so we use the :data:`aud.Handle.status` property to determine whether the sound finished playing.
 
 .. code-block:: python
 
    while handle.status:
    	time.sleep(0.1)
 
-We don't make any error checks if the user actually added a command
-line argument. As an exercise you could extend this program to play
-any number of command line supplied files in sequence.
+We don't make any error checks if the user actually added a command line argument. As an exercise you could extend this program to play any number of command line supplied files in sequence.
 
 Siren
 -----
 
-Let's get a little bit more complex. The **siren.py** example
-plays a generated siren sound that circles around your head.
-Depending on how many speakers you have and if the output
-device used supports the speaker setup, you will hear this effect.
-With stereo speakers you should at least hear some left-right-panning.
+Let's get a little bit more complex. The **siren.py** example plays a generated siren sound that circles around your head. Depending on how many speakers you have and if the output device used supports the speaker setup, you will hear this effect. With stereo speakers you should at least hear some left-right-panning.
 
-We start off again with importing the modules we need and
-we also define some properties of our siren sound.
-We want it to consist of two sine sounds with different frequencies.
-We define a length for the sine sounds and how long a fade in/out should take.
-We also know already how to open a device.
+We start off again with importing the modules we need and we also define some properties of our siren sound. We want it to consist of two sine sounds with different frequencies. We define a length for the sine sounds and how long a fade in/out should take. We also know already how to open a device.
 
 .. code-block:: python
 
@@ -128,35 +88,27 @@ We also know already how to open a device.
 
    device = aud.Device()
 
-The next thing to do is to define our sine waves and apply all the required effects.
-As each of the effect functions returns the corresponding sound,
-we can easily chain those calls together.
+The next thing to do is to define our sine waves and apply all the required effects. As each of the effect functions returns the corresponding sound, we can easily chain those calls together.
 
 .. code-block:: python
 
    high = aud.Sound.sine(880).limit(0, length).fadein(0, fadelength).fadeout(length - fadelength, length)
    low = aud.Sound.sine(700).limit(0, length).fadein(0, fadelength).fadeout(length - fadelength, length).volume(0.6)
 
-The next step is to connect the two sines,
-which we do using the :func:`aud.Sound.join` function.
+The next step is to connect the two sines, which we do using the :func:`aud.Sound.join` function.
 
 .. code-block:: python
 
    sound = high.join(low)
 
-The generated siren sound can now be played back and what we also do is to loop it.
-Therefore we set the :data:`aud.Handle.loop_count` to a negative value to loop forever.
+The generated siren sound can now be played back and what we also do is to loop it. Therefore we set the :data:`aud.Handle.loop_count` to a negative value to loop forever.
 
 .. code-block:: python
 
    handle = device.play(sound)
    handle.loop_count = -1
 
-Now we use some timing code to make sure our demo runs for 10 seconds,
-but we also use the time to update the location of our playing sound,
-with the :data:`aud.Handle.location` property, which is a three dimensional vector.
-The trigonometic calculation based on the running time of the program keeps
-the sound on the XZ plane letting it follow a circle around us.
+Now we use some timing code to make sure our demo runs for 10 seconds, but we also use the time to update the location of our playing sound, with the :data:`aud.Handle.location` property, which is a three dimensional vector. The trigonometic calculation based on the running time of the program keeps the sound on the XZ plane letting it follow a circle around us.
 
 .. code-block:: python
 
@@ -167,54 +119,33 @@ the sound on the XZ plane letting it follow a circle around us.
 
    	handle.location = [math.sin(angle), 0, -math.cos(angle)]
 
-As an exercise you could try to let the sound come from the far left
-and go to the far right and a little bit in front of you within the
-10 second runtime of the program. With this change you should be able
-to hear the volume of the sound change, depending on how far it is away from you.
-Updating the :data:`aud.Handle.velocity` property properly also enables the doppler effect.
-Compare your solution to the **siren2.py** demo.
+As an exercise you could try to let the sound come from the far left and go to the far right and a little bit in front of you within the 10 second runtime of the program. With this change you should be able to hear the volume of the sound change, depending on how far it is away from you. Updating the :data:`aud.Handle.velocity` property properly also enables the doppler effect. Compare your solution to the **siren2.py** demo.
 
 Tetris
 ------
 
-The **tetris.py** demo application shows an even more
-complex application which generates retro tetris music.
-Looking at the source code there should be nothing new here,
-again the functions used from audaspace are the same as in the previous examples.
-In the :func:`parseNote` function all single notes get joined which leads
-to a very long chain of sounds. If you think of :func:`aud.Sound.join`
-as a function that creates a binary tree with the two joined sounds as
-leaves then the :func:`parseNote` function creates a very unbalanced tree.
+The **tetris.py** demo application shows an even more complex application which generates retro tetris music. Looking at the source code there should be nothing new here, again the functions used from audaspace are the same as in the previous examples. In the :func:`parseNote` function all single notes get joined which leads to a very long chain of sounds. If you think of :func:`aud.Sound.join` as a function that creates a binary tree with the two joined sounds as leaves then the :func:`parseNote` function creates a very unbalanced tree.
 
-Insted we could rewrite the code to use two other classes:
-:class:`aud.Sequence` and :class:`aud.SequenceEntry` to sequence the notes.
-The **tetris2.py** application does exactly that.
-Before the while loop we add a variable that stores the current position
-in the score and create a new :class:`aud.Sequence` object.
+Insted we could rewrite the code to use two other classes: :class:`aud.Sequence` and :class:`aud.SequenceEntry` to sequence the notes. The **tetris2.py** application does exactly that. Before the while loop we add a variable that stores the current position in the score and create a new :class:`aud.Sequence` object.
 
 .. code-block:: python
 
    position = 0
    sequence = aud.Sequence()
 
-Then in the loop we can create the note simply by chaining the
-:func:`aud.Sound.square` generator and :func:`aud.Sound.fadein`
-and :func:`aud.Sound.fadeout` effects.
+Then in the loop we can create the note simply by chaining the :func:`aud.Sound.square` generator and :func:`aud.Sound.fadein` and :func:`aud.Sound.fadeout` effects.
 
 .. code-block:: python
 
    note = aud.Sound.square(freq, rate).fadein(0, fadelength).fadeout(length - fadelength, fadelength)
 
-Now instead of using :func:`aud.Sound.limit` and :func:`aud.Sound.join`
-we simply add the sound to the sequence.
+Now instead of using :func:`aud.Sound.limit` and :func:`aud.Sound.join` we simply add the sound to the sequence.
 
 .. code-block:: python
 
    entry = sequence.add(note, position, position + length, 0)
 
-The entry returned from the :func:`aud.Sequence.add`
-function is an object of the :class:`aud.SequenceEntry` class.
-We can use this entry to mute the note in case it's actually a pause.
+The entry returned from the :func:`aud.Sequence.add` function is an object of the :class:`aud.SequenceEntry` class. We can use this entry to mute the note in case it's actually a pause.
 
 .. code-block:: python
 
@@ -227,14 +158,9 @@ Lastly we have to update our position variable.
 
    position += length
 
-Now in **tetris2.py** we used the :data:`aud.SequenceEntry.muted`
-property to show how the :class:`aud.SequenceEntry` class can be used,
-but it would actually be smarter to not even create a note for pauses and just skip them.
-You can try to implement this as an exercise and then check out the solution in **tetris3.py**.
+Now in **tetris2.py** we used the :data:`aud.SequenceEntry.muted` property to show how the :class:`aud.SequenceEntry` class can be used, but it would actually be smarter to not even create a note for pauses and just skip them. You can try to implement this as an exercise and then check out the solution in **tetris3.py**.
 
 Conclusion
 ----------
 
-We introduced all five currently available classes in the audaspace Python API.
-Of course all classes offer a lot more functions than have been used in these demo applications,
-check out the specific class documentation for more details.
+We introduced all five currently available classes in the audaspace Python API. Of course all classes offer a lot more functions than have been used in these demo applications, check out the specific class documentation for more details.

extern/audaspace/bindings/python/PyDevice.cpp

@@ -124,17 +124,15 @@ Device_new(PyTypeObject* type, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Device_lock_doc,
-			 ".. classmethod:: lock()\n\n"
-			 "   Locks the device so that it's guaranteed, that no samples are\n"
-			 "   read from the streams until :meth:`unlock` is called.\n"
-			 "   This is useful if you want to do start/stop/pause/resume some\n"
-			 "   sounds at the same time.\n\n"
-			 "   .. note::\n\n"
-			 "      The device has to be unlocked as often as locked to be\n"
-			 "      able to continue playback.\n\n"
-			 "   .. warning::\n\n"
-			 "      Make sure the time between locking and unlocking is\n"
-			 "      as short as possible to avoid clicks.");
+			 "lock()\n\n"
+			 "Locks the device so that it's guaranteed, that no samples are "
+			 "read from the streams until :meth:`unlock` is called.\n"
+			 "This is useful if you want to do start/stop/pause/resume some "
+			 "sounds at the same time.\n\n"
+			 ".. note:: The device has to be unlocked as often as locked to be "
+			 "able to continue playback.\n\n"
+			 ".. warning:: Make sure the time between locking and unlocking is "
+			 "as short as possible to avoid clicks.");
 
 static PyObject *
 Device_lock(Device* self)
@@ -152,15 +150,15 @@ Device_lock(Device* self)
 }
 
 PyDoc_STRVAR(M_aud_Device_play_doc,
-			 ".. classmethod:: play(sound, keep=False)\n\n"
-			 "   Plays a sound.\n\n"
-			 "   :arg sound: The sound to play.\n"
-			 "   :type sound: :class:`Sound`\n"
-			 "   :arg keep: See :attr:`Handle.keep`.\n"
-			 "   :type keep: bool\n"
-			 "   :return: The playback handle with which playback can be\n"
-			 "      controlled with.\n"
-			 "   :rtype: :class:`Handle`");
+			 "play(sound, keep=False)\n\n"
+			 "Plays a sound.\n\n"
+			 ":arg sound: The sound to play.\n"
+			 ":type sound: :class:`Sound`\n"
+			 ":arg keep: See :attr:`Handle.keep`.\n"
+			 ":type keep: bool\n"
+			 ":return: The playback handle with which playback can be "
+			 "controlled with.\n"
+			 ":rtype: :class:`Handle`");
 
 static PyObject *
 Device_play(Device* self, PyObject* args, PyObject* kwds)
@@ -212,8 +210,8 @@ Device_play(Device* self, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Device_stopAll_doc,
-			 ".. classmethod:: stopAll()\n\n"
-			 "   Stops all playing and paused sounds.");
+			 "stopAll()\n\n"
+			 "Stops all playing and paused sounds.");
 
 static PyObject *
 Device_stopAll(Device* self)
@@ -231,9 +229,9 @@ Device_stopAll(Device* self)
 }
 
 PyDoc_STRVAR(M_aud_Device_unlock_doc,
-			 ".. classmethod:: unlock()\n\n"
-			 "   Unlocks the device after a lock call, see :meth:`lock` for\n"
-			 "   details.");
+			 "unlock()\n\n"
+			 "Unlocks the device after a lock call, see :meth:`lock` for "
+			 "details.");
 
 static PyObject *
 Device_unlock(Device* self)
@@ -286,7 +284,7 @@ Device_get_channels(Device* self, void* nothing)
 
 PyDoc_STRVAR(M_aud_Device_distance_model_doc,
 			 "The distance model of the device.\n\n"
-			 ".. seealso:: `OpenAL Documentation <https://www.openal.org/documentation/>`__");
+			 ".. seealso:: http://connect.creativelabs.com/openal/Documentation/OpenAL%201.1%20Specification.htm#_Toc199835864");
 
 static PyObject *
 Device_get_distance_model(Device* self, void* nothing)

extern/audaspace/bindings/python/PyDynamicMusic.cpp

@@ -60,12 +60,12 @@ DynamicMusic_dealloc(DynamicMusicP* self)
 }
 
 PyDoc_STRVAR(M_aud_DynamicMusic_addScene_doc,
-	".. classmethod:: addScene(scene)\n\n"
-	"   Adds a new scene.\n\n"
-	"   :arg scene: The scene sound.\n"
-	"   :type scene: :class:`Sound`\n"
-	"   :return: The new scene id.\n"
-	"   :rtype: int");
+	"addScene(scene)\n\n"
+	"Adds a new scene.\n\n"
+	":arg scene: The scene sound.\n"
+	":type scene: :class:`Sound`\n"
+	":return: The new scene id.\n"
+	":rtype: int");
 
 static PyObject *
 DynamicMusic_addScene(DynamicMusicP* self, PyObject* args)
@@ -90,16 +90,16 @@ DynamicMusic_addScene(DynamicMusicP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_DynamicMusic_addTransition_doc,
-	".. classmethod:: addTransition(ini, end, transition)\n\n"
-	"   Adds a new scene.\n\n"
-	"   :arg ini: the initial scene foor the transition.\n"
-	"   :type ini: int\n"
-	"   :arg end: The final scene for the transition.\n"
-	"   :type end: int\n"
-	"   :arg transition: The transition sound.\n"
-	"   :type transition: :class:`Sound`\n"
-	"   :return: false if the ini or end scenes don't exist, true othrwise.\n"
-	"   :rtype: bool");
+	"addTransition(ini, end, transition)\n\n"
+	"Adds a new scene.\n\n"
+	":arg ini: the initial scene foor the transition.\n"
+	":type ini: int\n"
+	":arg end: The final scene for the transition.\n"
+	":type end: int\n"
+	":arg transition: The transition sound.\n"
+	":type transition: :class:`Sound`\n"
+	":return: false if the ini or end scenes don't exist, true othrwise.\n"
+	":rtype: bool");
 
 static PyObject *
 DynamicMusic_addTransition(DynamicMusicP* self, PyObject* args)
@@ -125,10 +125,10 @@ DynamicMusic_addTransition(DynamicMusicP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_DynamicMusic_resume_doc,
-	".. classmethod:: resume()\n\n"
-	"   Resumes playback of the scene.\n\n"
-	"   :return: Whether the action succeeded.\n"
-	"   :rtype: bool");
+	"resume()\n\n"
+	"Resumes playback of the scene.\n\n"
+	":return: Whether the action succeeded.\n"
+	":rtype: bool");
 
 static PyObject *
 DynamicMusic_resume(DynamicMusicP* self)
@@ -145,10 +145,10 @@ DynamicMusic_resume(DynamicMusicP* self)
 }
 
 PyDoc_STRVAR(M_aud_DynamicMusic_pause_doc,
-	".. classmethod:: pause()\n\n"
-	"   Pauses playback of the scene.\n\n"
-	"   :return: Whether the action succeeded.\n"
-	"   :rtype: bool");
+	"pause()\n\n"
+	"Pauses playback of the scene.\n\n"
+	":return: Whether the action succeeded.\n"
+	":rtype: bool");
 
 static PyObject *
 DynamicMusic_pause(DynamicMusicP* self)
@@ -165,10 +165,10 @@ DynamicMusic_pause(DynamicMusicP* self)
 }
 
 PyDoc_STRVAR(M_aud_DynamicMusic_stop_doc,
-	".. classmethod:: stop()\n\n"
-	"   Stops playback of the scene.\n\n"
-	"   :return: Whether the action succeeded.\n"
-	"   :rtype: bool\n\n");
+	"stop()\n\n"
+	"Stops playback of the scene.\n\n"
+	":return: Whether the action succeeded.\n"
+	":rtype: bool\n\n");
 
 static PyObject *
 DynamicMusic_stop(DynamicMusicP* self)
@@ -464,4 +464,4 @@ void addDynamicMusicToModule(PyObject* module)
 {
 	Py_INCREF(&DynamicMusicType);
 	PyModule_AddObject(module, "DynamicMusic", (PyObject *)&DynamicMusicType);
-}
+}

extern/audaspace/bindings/python/PyHRTF.cpp

@@ -54,16 +54,16 @@ HRTF_dealloc(HRTFP* self)
 }
 
 PyDoc_STRVAR(M_aud_HRTF_addImpulseResponse_doc,
-	".. classmethod:: addImpulseResponseFromSound(sound, azimuth, elevation)\n\n"
-	"   Adds a new hrtf to the HRTF object\n\n"
-	"   :arg sound: The sound that contains the hrtf.\n"
-	"   :type sound: :class:`Sound`\n"
-	"   :arg azimuth: The azimuth angle of the hrtf.\n"
-	"   :type azimuth: float\n"
-	"   :arg elevation: The elevation angle of the hrtf.\n"
-	"   :type elevation: float\n"
-	"   :return: Whether the action succeeded.\n"
-	"   :rtype: bool");
+	"addImpulseResponseFromSound(sound, azimuth, elevation)\n\n"
+	"Adds a new hrtf to the HRTF object\n\n"
+	":arg sound: The sound that contains the hrtf.\n"
+	":type sound: :class:`Sound`\n"
+	":arg azimuth: The azimuth angle of the hrtf.\n"
+	":type azimuth: float\n"
+	":arg elevation: The elevation angle of the hrtf.\n"
+	":type elevation: float\n"
+	":return: Whether the action succeeded.\n"
+	":rtype: bool");
 
 static PyObject *
 HRTF_addImpulseResponseFromSound(HRTFP* self, PyObject* args)
@@ -90,14 +90,14 @@ HRTF_addImpulseResponseFromSound(HRTFP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_HRTF_loadLeftHrtfSet_doc,
-	".. classmethod:: loadLeftHrtfSet(extension, directory)\n\n"
-	"   Loads all HRTFs from a directory.\n\n"
-	"   :arg extension: The file extension of the hrtfs.\n"
-	"   :type extension: string\n"
-	"   :arg directory: The path to where the HRTF files are located.\n"
-	"   :type extension: string\n"
-	"   :return: The loaded :class:`HRTF` object.\n"
-	"   :rtype: :class:`HRTF`\n\n");
+	"loadLeftHrtfSet(extension, directory)\n\n"
+	"Loads all HRTFs from a directory.\n\n"
+	":arg extension: The file extension of the hrtfs.\n"
+	":type extension: string\n"
+	":arg directory: The path to where the HRTF files are located.\n"
+	":type extension: string\n"
+	":return: The loaded :class:`HRTF` object.\n"
+	":rtype: :class:`HRTF`\n\n");
 
 static PyObject *
 HRTF_loadLeftHrtfSet(PyTypeObject* type, PyObject* args)
@@ -125,14 +125,14 @@ HRTF_loadLeftHrtfSet(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_HRTF_loadRightHrtfSet_doc,
-	".. classmethod:: loadLeftHrtfSet(extension, directory)\n\n"
-	"   Loads all HRTFs from a directory.\n\n"
-	"   :arg extension: The file extension of the hrtfs.\n"
-	"   :type extension: string\n"
-	"   :arg directory: The path to where the HRTF files are located.\n"
-	"   :type extension: string\n"
-	"   :return: The loaded :class:`HRTF` object.\n"
-	"   :rtype: :class:`HRTF`\n\n");
+	"loadLeftHrtfSet(extension, directory)\n\n"
+	"Loads all HRTFs from a directory.\n\n"
+	":arg extension: The file extension of the hrtfs.\n"
+	":type extension: string\n"
+	":arg directory: The path to where the HRTF files are located.\n"
+	":type extension: string\n"
+	":return: The loaded :class:`HRTF` object.\n"
+	":rtype: :class:`HRTF`\n\n");
 
 static PyObject *
 HRTF_loadRightHrtfSet(PyTypeObject* type, PyObject* args)

extern/audaspace/bindings/python/PyHandle.cpp

@@ -38,10 +38,10 @@ Handle_dealloc(Handle* self)
 }
 
 PyDoc_STRVAR(M_aud_Handle_pause_doc,
-			 ".. classmethod:: pause()\n\n"
-			 "   Pauses playback.\n\n"
-			 "   :return: Whether the action succeeded.\n"
-			 "   :rtype: bool");
+			 "pause()\n\n"
+			 "Pauses playback.\n\n"
+			 ":return: Whether the action succeeded.\n"
+			 ":rtype: bool");
 
 static PyObject *
 Handle_pause(Handle* self)
@@ -58,10 +58,10 @@ Handle_pause(Handle* self)
 }
 
 PyDoc_STRVAR(M_aud_Handle_resume_doc,
-			 ".. classmethod:: resume()\n\n"
-			 "   Resumes playback.\n\n"
-			 "   :return: Whether the action succeeded.\n"
-			 "   :rtype: bool");
+			 "resume()\n\n"
+			 "Resumes playback.\n\n"
+			 ":return: Whether the action succeeded.\n"
+			 ":rtype: bool");
 
 static PyObject *
 Handle_resume(Handle* self)
@@ -78,11 +78,11 @@ Handle_resume(Handle* self)
 }
 
 PyDoc_STRVAR(M_aud_Handle_stop_doc,
-			 ".. classmethod:: stop()\n\n"
-			 "   Stops playback.\n\n"
-			 "   :return: Whether the action succeeded.\n"
-			 "   :rtype: bool\n\n"
-			 "   .. note:: This makes the handle invalid.");
+			 "stop()\n\n"
+			 "Stops playback.\n\n"
+			 ":return: Whether the action succeeded.\n"
+			 ":rtype: bool\n\n"
+			 ".. note:: This makes the handle invalid.");
 
 static PyObject *
 Handle_stop(Handle* self)
@@ -1122,3 +1122,5 @@ void addHandleToModule(PyObject* module)
 	Py_INCREF(&HandleType);
 	PyModule_AddObject(module, "Handle", (PyObject *)&HandleType);
 }
+
+

extern/audaspace/bindings/python/PyPlaybackManager.cpp

@@ -60,15 +60,14 @@ PlaybackManager_dealloc(PlaybackManagerP* self)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_play_doc,
-	".. classmethod:: setVolume(sound, catKey)\n\n"
-	"   Plays a sound through the playback manager and assigns it to a category.\n\n"
-	"   :arg sound: The sound to play.\n"
-	"   :type sound: :class:`Sound`\n"
-	"   :arg catKey: the key of the category in which the sound will be added,\n"
-	"      if it doesn't exist, a new one will be created.\n"
-	"   :type catKey: int\n"
-	"   :return: The playback handle with which playback can be controlled with.\n"
-	"   :rtype: :class:`Handle`");
+	"setVolume(sound, catKey)\n\n"
+	"Plays a sound through the playback manager and assigns it to a category.\n\n"
+	":arg sound: The sound to play.\n"
+	":type sound: :class:`Sound`\n"
+	":arg catKey: the key of the category in which the sound will be added, if it doesn't exist, a new one will be created.\n"
+	":type catKey: int\n"
+	":return: The playback handle with which playback can be controlled with.\n"
+	":rtype: :class:`Handle`");
 
 static PyObject *
 PlaybackManager_play(PlaybackManagerP* self, PyObject* args)
@@ -104,12 +103,12 @@ PlaybackManager_play(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_resume_doc,
-	".. classmethod:: resume(catKey)\n\n"
-	"   Resumes playback of the catgory.\n\n"
-	"   :arg catKey: the key of the category.\n"
-	"   :type catKey: int\n"
-	"   :return: Whether the action succeeded.\n"
-	"   :rtype: bool");
+	"resume(catKey)\n\n"
+	"Resumes playback of the catgory.\n\n"
+	":arg catKey: the key of the category.\n"
+	":type catKey: int\n"
+	":return: Whether the action succeeded.\n"
+	":rtype: bool");
 
 static PyObject *
 PlaybackManager_resume(PlaybackManagerP* self, PyObject* args)
@@ -131,12 +130,12 @@ PlaybackManager_resume(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_pause_doc,
-	".. classmethod:: pause(catKey)\n\n"
-	"   Pauses playback of the category.\n\n"
-	"   :arg catKey: the key of the category.\n"
-	"   :type catKey: int\n"
-	"   :return: Whether the action succeeded.\n"
-	"   :rtype: bool");
+	"pause(catKey)\n\n"
+	"Pauses playback of the category.\n\n"
+	":arg catKey: the key of the category.\n"
+	":type catKey: int\n"
+	":return: Whether the action succeeded.\n"
+	":rtype: bool");
 
 static PyObject *
 PlaybackManager_pause(PlaybackManagerP* self, PyObject* args)
@@ -158,12 +157,12 @@ PlaybackManager_pause(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_add_category_doc,
-	".. classmethod:: addCategory(volume)\n\n"
-	"   Adds a category with a custom volume.\n\n"
-	"   :arg volume: The volume for ther new category.\n"
-	"   :type volume: float\n"
-	"   :return: The key of the new category.\n"
-	"   :rtype: int\n\n");
+	"addCategory(volume)\n\n"
+	"Adds a category with a custom volume.\n\n"
+	":arg volume: The volume for ther new category.\n"
+	":type volume: float\n"
+	":return: The key of the new category.\n"
+	":rtype: int\n\n");
 
 static PyObject *
 PlaybackManager_add_category(PlaybackManagerP* self, PyObject* args)
@@ -185,12 +184,12 @@ PlaybackManager_add_category(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_get_volume_doc,
-	".. classmethod:: getVolume(catKey)\n\n"
-	"   Retrieves the volume of a category.\n\n"
-	"   :arg catKey: the key of the category.\n"
-	"   :type catKey: int\n"
-	"   :return: The volume of the cateogry.\n"
-	"   :rtype: float\n\n");
+	"getVolume(catKey)\n\n"
+	"Retrieves the volume of a category.\n\n"
+	":arg catKey: the key of the category.\n"
+	":type catKey: int\n"
+	":return: The volume of the cateogry.\n"
+	":rtype: float\n\n");
 
 static PyObject *
 PlaybackManager_get_volume(PlaybackManagerP* self, PyObject* args)
@@ -212,14 +211,14 @@ PlaybackManager_get_volume(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_set_volume_doc,
-	".. classmethod:: setVolume(volume, catKey)\n\n"
-	"   Changes the volume of a category.\n\n"
-	"   :arg volume: the new volume value.\n"
-	"   :type volume: float\n"
-	"   :arg catKey: the key of the category.\n"
-	"   :type catKey: int\n"
-	"   :return: Whether the action succeeded.\n"
-	"   :rtype: int\n\n");
+	"setVolume(volume, catKey)\n\n"
+	"Changes the volume of a category.\n\n"
+	":arg volume: the new volume value.\n"
+	":type volume: float\n"
+	":arg catKey: the key of the category.\n"
+	":type catKey: int\n"
+	":return: Whether the action succeeded.\n"
+	":rtype: int\n\n");
 
 static PyObject *
 PlaybackManager_set_volume(PlaybackManagerP* self, PyObject* args)
@@ -242,12 +241,12 @@ PlaybackManager_set_volume(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_stop_doc,
-	".. classmethod:: stop(catKey)\n\n"
-	"   Stops playback of the category.\n\n"
-	"   :arg catKey: the key of the category.\n"
-	"   :type catKey: int\n"
-	"   :return: Whether the action succeeded.\n"
-	"   :rtype: bool\n\n");
+	"stop(catKey)\n\n"
+	"Stops playback of the category.\n\n"
+	":arg catKey: the key of the category.\n"
+	":type catKey: int\n"
+	":return: Whether the action succeeded.\n"
+	":rtype: bool\n\n");
 
 static PyObject *
 PlaybackManager_stop(PlaybackManagerP* self, PyObject* args)
@@ -269,8 +268,8 @@ PlaybackManager_stop(PlaybackManagerP* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_PlaybackManager_clean_doc,
-	".. classmethod:: clean()\n\n"
-	"   Cleans all the invalid and finished sound from the playback manager.\n\n");
+	"clean()\n\n"
+	"Cleans all the invalid and finished sound from the playback manager.\n\n");
 
 static PyObject *
 PlaybackManager_clean(PlaybackManagerP* self)

extern/audaspace/bindings/python/PySequence.cpp

@@ -99,18 +99,18 @@ Sequence_new(PyTypeObject* type, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Sequence_add_doc,
-			 ".. classmethod:: add()\n\n"
-			 "   Adds a new entry to the sequence.\n\n"
-			 "   :arg sound: The sound this entry should play.\n"
-			 "   :type sound: :class:`Sound`\n"
-			 "   :arg begin: The start time.\n"
-			 "   :type begin: float\n"
-			 "   :arg end: The end time or a negative value if determined by the sound.\n"
-			 "   :type end: float\n"
-			 "   :arg skip: How much seconds should be skipped at the beginning.\n"
-			 "   :type skip: float\n"
-			 "   :return: The entry added.\n"
-			 "   :rtype: :class:`SequenceEntry`");
+			 "add()\n\n"
+			 "Adds a new entry to the sequence.\n\n"
+			 ":arg sound: The sound this entry should play.\n"
+			 ":type sound: :class:`Sound`\n"
+			 ":arg begin: The start time.\n"
+			 ":type begin: float\n"
+			 ":arg end: The end time or a negative value if determined by the sound.\n"
+			 ":type end: float\n"
+			 ":arg skip: How much seconds should be skipped at the beginning.\n"
+			 ":type skip: float\n"
+			 ":return: The entry added.\n"
+			 ":rtype: :class:`SequenceEntry`");
 
 static PyObject *
 Sequence_add(Sequence* self, PyObject* args, PyObject* kwds)
@@ -151,10 +151,10 @@ Sequence_add(Sequence* self, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Sequence_remove_doc,
-			 ".. classmethod:: remove()\n\n"
-			 "   Removes an entry from the sequence.\n\n"
-			 "   :arg entry: The entry to remove.\n"
-			 "   :type entry: :class:`SequenceEntry`\n");
+			 "remove()\n\n"
+			 "Removes an entry from the sequence.\n\n"
+			 ":arg entry: The entry to remove.\n"
+			 ":type entry: :class:`SequenceEntry`\n");
 
 static PyObject *
 Sequence_remove(Sequence* self, PyObject* args)
@@ -183,16 +183,16 @@ Sequence_remove(Sequence* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sequence_setAnimationData_doc,
-			 ".. classmethod:: setAnimationData()\n\n"
-			 "   Writes animation data to a sequence.\n\n"
-			 "   :arg type: The type of animation data.\n"
-			 "   :type type: int\n"
-			 "   :arg frame: The frame this data is for.\n"
-			 "   :type frame: int\n"
-			 "   :arg data: The data to write.\n"
-			 "   :type data: sequence of float\n"
-			 "   :arg animated: Whether the attribute is animated.\n"
-			 "   :type animated: bool");
+			 "setAnimationData()\n\n"
+			 "Writes animation data to a sequence.\n\n"
+			 ":arg type: The type of animation data.\n"
+			 ":type type: int\n"
+			 ":arg frame: The frame this data is for.\n"
+			 ":type frame: int\n"
+			 ":arg data: The data to write.\n"
+			 ":type data: sequence of float\n"
+			 ":arg animated: Whether the attribute is animated.\n"
+			 ":type animated: bool");
 
 static PyObject *
 Sequence_setAnimationData(Sequence* self, PyObject* args)
@@ -325,7 +325,7 @@ Sequence_set_channels(Sequence* self, PyObject* args, void* nothing)
 
 PyDoc_STRVAR(M_aud_Sequence_distance_model_doc,
 			 "The distance model of the sequence.\n\n"
-			  ".. seealso:: `OpenAL Documentation <https://www.openal.org/documentation/>`__");
+			 ".. seealso:: http://connect.creativelabs.com/openal/Documentation/OpenAL%201.1%20Specification.htm#_Toc199835864");
 
 static PyObject *
 Sequence_get_distance_model(Sequence* self, void* nothing)

extern/audaspace/bindings/python/PySequenceEntry.cpp

@@ -43,14 +43,14 @@ SequenceEntry_dealloc(SequenceEntry* self)
 }
 
 PyDoc_STRVAR(M_aud_SequenceEntry_move_doc,
-			 ".. classmethod:: move()\n\n"
-			 "   Moves the entry.\n\n"
-			 "   :arg begin: The new start time.\n"
-			 "   :type begin: float\n"
-			 "   :arg end: The new end time or a negative value if unknown.\n"
-			 "   :type end: float\n"
-			 "   :arg skip: How many seconds to skip at the beginning.\n"
-			 "   :type skip: float\n");
+			 "move()\n\n"
+			 "Moves the entry.\n\n"
+			 ":arg begin: The new start time.\n"
+			 ":type begin: float\n"
+			 ":arg end: The new end time or a negative value if unknown.\n"
+			 ":type end: float\n"
+			 ":arg skip: How many seconds to skip at the beginning.\n"
+			 ":type skip: float\n");
 
 static PyObject *
 SequenceEntry_move(SequenceEntry* self, PyObject* args)
@@ -73,16 +73,16 @@ SequenceEntry_move(SequenceEntry* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_SequenceEntry_setAnimationData_doc,
-			 ".. classmethod:: setAnimationData()\n\n"
-			 "   Writes animation data to a sequenced entry.\n\n"
-			 "   :arg type: The type of animation data.\n"
-			 "   :type type: int\n"
-			 "   :arg frame: The frame this data is for.\n"
-			 "   :type frame: int\n"
-			 "   :arg data: The data to write.\n"
-			 "   :type data: sequence of float\n"
-			 "   :arg animated: Whether the attribute is animated.\n"
-			 "   :type animated: bool");
+			 "setAnimationData()\n\n"
+			 "Writes animation data to a sequenced entry.\n\n"
+			 ":arg type: The type of animation data.\n"
+			 ":type type: int\n"
+			 ":arg frame: The frame this data is for.\n"
+			 ":type frame: int\n"
+			 ":arg data: The data to write.\n"
+			 ":type data: sequence of float\n"
+			 ":arg animated: Whether the attribute is animated.\n"
+			 ":type animated: bool");
 
 static PyObject *
 SequenceEntry_setAnimationData(SequenceEntry* self, PyObject* args)

extern/audaspace/bindings/python/PySound.cpp

@@ -114,11 +114,11 @@ Sound_new(PyTypeObject* type, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Sound_data_doc,
-			 ".. classmethod:: data()\n\n"
-			 "   Retrieves the data of the sound as numpy array.\n\n"
-			 "   :return: A two dimensional numpy float array.\n"
-			 "   :rtype: :class:`numpy.ndarray`\n\n"
-			 "   .. note:: Best efficiency with cached sounds.");
+			 "data()\n\n"
+			 "Retrieves the data of the sound as numpy array.\n\n"
+			 ":return: A two dimensional numpy float array.\n"
+			 ":rtype: :class:`numpy.ndarray`\n\n"
+			 ".. note:: Best efficiency with cached sounds.");
 
 static PyObject *
 Sound_data(Sound* self)
@@ -145,24 +145,24 @@ Sound_data(Sound* self)
 }
 
 PyDoc_STRVAR(M_aud_Sound_write_doc,
-			 ".. classmethod:: write(filename, rate, channels, format, container, codec, bitrate, buffersize)\n\n"
-			 "   Writes the sound to a file.\n\n"
-			 "   :arg filename: The path to write to.\n"
-			 "   :type filename: string\n"
-			 "   :arg rate: The sample rate to write with.\n"
-			 "   :type rate: int\n"
-			 "   :arg channels: The number of channels to write with.\n"
-			 "   :type channels: int\n"
-			 "   :arg format: The sample format to write with.\n"
-			 "   :type format: int\n"
-			 "   :arg container: The container format for the file.\n"
-			 "   :type container: int\n"
-			 "   :arg codec: The codec to use in the file.\n"
-			 "   :type codec: int\n"
-			 "   :arg bitrate: The bitrate to write with.\n"
-			 "   :type bitrate: int\n"
-			 "   :arg buffersize: The size of the writing buffer.\n"
-			 "   :type buffersize: int");
+			 "write(filename, rate, channels, format, container, codec, bitrate, buffersize)\n\n"
+			 "Writes the sound to a file.\n\n"
+			 ":arg filename: The path to write to.\n"
+			 ":type filename: string\n"
+			 ":arg rate: The sample rate to write with.\n"
+			 ":type rate: int\n"
+			 ":arg channels: The number of channels to write with.\n"
+			 ":type channels: int\n"
+			 ":arg format: The sample format to write with.\n"
+			 ":type format: int\n"
+			 ":arg container: The container format for the file.\n"
+			 ":type container: int\n"
+			 ":arg codec: The codec to use in the file.\n"
+			 ":type codec: int\n"
+			 ":arg bitrate: The bitrate to write with.\n"
+			 ":type bitrate: int\n"
+			 ":arg buffersize: The size of the writing buffer.\n"
+			 ":type buffersize: int\n");
 
 static PyObject *
 Sound_write(Sound* self, PyObject* args, PyObject* kwds)
@@ -286,14 +286,14 @@ Sound_write(Sound* self, PyObject* args, PyObject* kwds)
 }
 
 PyDoc_STRVAR(M_aud_Sound_buffer_doc,
-			 ".. classmethod:: buffer(data, rate)\n\n"
-			 "   Creates a sound from a data buffer.\n\n"
-			 "   :arg data: The data as two dimensional numpy array.\n"
-			 "   :type data: numpy.ndarray\n"
-			 "   :arg rate: The sample rate.\n"
-			 "   :type rate: double\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "buffer(data, rate)\n\n"
+			 "Creates a sound from a data buffer.\n\n"
+			 ":arg data: The data as two dimensional numpy array.\n"
+			 ":type data: numpy.ndarray\n"
+			 ":arg rate: The sample rate.\n"
+			 ":type rate: double\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_buffer(PyTypeObject* type, PyObject* args)
@@ -356,17 +356,16 @@ Sound_buffer(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_cache_doc,
-			 ".. classmethod:: cache()\n\n"
-			 "   Caches a sound into RAM.\n\n"
-			 "   This saves CPU usage needed for decoding and file access if the\n"
-			 "   underlying sound reads from a file on the harddisk,\n"
-			 "   but it consumes a lot of memory.\n\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. note:: Only known-length factories can be buffered.\n\n"
-			 "   .. warning::\n\n"
-			 "      Raw PCM data needs a lot of space, only buffer\n"
-			 "      short factories.");
+			 "cache()\n\n"
+			 "Caches a sound into RAM.\n"
+			 "This saves CPU usage needed for decoding and file access if the "
+			 "underlying sound reads from a file on the harddisk, but it "
+			 "consumes a lot of memory.\n\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. note:: Only known-length factories can be buffered.\n\n"
+			 ".. warning:: Raw PCM data needs a lot of space, only buffer "
+			 "short factories.");
 
 static PyObject *
 Sound_cache(Sound* self)
@@ -392,16 +391,15 @@ Sound_cache(Sound* self)
 }
 
 PyDoc_STRVAR(M_aud_Sound_file_doc,
-			 ".. classmethod:: file(filename)\n\n"
-			 "   Creates a sound object of a sound file.\n\n"
-			 "   :arg filename: Path of the file.\n"
-			 "   :type filename: string\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. warning::\n\n"
-			 "      If the file doesn't exist or can't be read you will\n"
-			 "      not get an exception immediately, but when you try to start\n"
-			 "      playback of that sound.\n");
+			 "file(filename)\n\n"
+			 "Creates a sound object of a sound file.\n\n"
+			 ":arg filename: Path of the file.\n"
+			 ":type filename: string\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. warning:: If the file doesn't exist or can't be read you will "
+			 "not get an exception immediately, but when you try to start "
+			 "playback of that sound.");
 
 static PyObject *
 Sound_file(PyTypeObject* type, PyObject* args)
@@ -432,15 +430,15 @@ Sound_file(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_sawtooth_doc,
-			 ".. classmethod:: sawtooth(frequency, rate=48000)\n\n"
-			 "   Creates a sawtooth sound which plays a sawtooth wave.\n\n"
-			 "   :arg frequency: The frequency of the sawtooth wave in Hz.\n"
-			 "   :type frequency: float\n"
-			 "   :arg rate: The sampling rate in Hz. It's recommended to set this\n"
-			 "      value to the playback device's samling rate to avoid resamping.\n"
-			 "   :type rate: int\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "sawtooth(frequency, rate=48000)\n\n"
+			 "Creates a sawtooth sound which plays a sawtooth wave.\n\n"
+			 ":arg frequency: The frequency of the sawtooth wave in Hz.\n"
+			 ":type frequency: float\n"
+			 ":arg rate: The sampling rate in Hz. It's recommended to set this "
+			 "value to the playback device's samling rate to avoid resamping.\n"
+			 ":type rate: int\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_sawtooth(PyTypeObject* type, PyObject* args)
@@ -472,13 +470,13 @@ Sound_sawtooth(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_silence_doc,
-			 ".. classmethod:: silence(rate=48000)\n\n"
-			 "   Creates a silence sound which plays simple silence.\n\n"
-			 "   :arg rate: The sampling rate in Hz. It's recommended to set this\n"
-			 "      value to the playback device's samling rate to avoid resamping.\n"
-			 "   :type rate: int\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "silence(rate=48000)\n\n"
+			 "Creates a silence sound which plays simple silence.\n\n"
+			 ":arg rate: The sampling rate in Hz. It's recommended to set this "
+			 "value to the playback device's samling rate to avoid resamping.\n"
+			 ":type rate: int\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_silence(PyTypeObject* type, PyObject* args)
@@ -509,15 +507,15 @@ Sound_silence(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_sine_doc,
-			 ".. classmethod:: sine(frequency, rate=48000)\n\n"
-			 "   Creates a sine sound which plays a sine wave.\n\n"
-			 "   :arg frequency: The frequency of the sine wave in Hz.\n"
-			 "   :type frequency: float\n"
-			 "   :arg rate: The sampling rate in Hz. It's recommended to set this\n"
-			 "      value to the playback device's samling rate to avoid resamping.\n"
-			 "   :type rate: int\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "sine(frequency, rate=48000)\n\n"
+			 "Creates a sine sound which plays a sine wave.\n\n"
+			 ":arg frequency: The frequency of the sine wave in Hz.\n"
+			 ":type frequency: float\n"
+			 ":arg rate: The sampling rate in Hz. It's recommended to set this "
+			 "value to the playback device's samling rate to avoid resamping.\n"
+			 ":type rate: int\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_sine(PyTypeObject* type, PyObject* args)
@@ -549,15 +547,15 @@ Sound_sine(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_square_doc,
-			 ".. classmethod:: square(frequency, rate=48000)\n\n"
-			 "   Creates a square sound which plays a square wave.\n\n"
-			 "   :arg frequency: The frequency of the square wave in Hz.\n"
-			 "   :type frequency: float\n"
-			 "   :arg rate: The sampling rate in Hz. It's recommended to set this\n"
-			 "      value to the playback device's samling rate to avoid resamping.\n"
-			 "   :type rate: int\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "square(frequency, rate=48000)\n\n"
+			 "Creates a square sound which plays a square wave.\n\n"
+			 ":arg frequency: The frequency of the square wave in Hz.\n"
+			 ":type frequency: float\n"
+			 ":arg rate: The sampling rate in Hz. It's recommended to set this "
+			 "value to the playback device's samling rate to avoid resamping.\n"
+			 ":type rate: int\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_square(PyTypeObject* type, PyObject* args)
@@ -589,15 +587,15 @@ Sound_square(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_triangle_doc,
-			 ".. classmethod:: triangle(frequency, rate=48000)\n\n"
-			 "   Creates a triangle sound which plays a triangle wave.\n\n"
-			 "   :arg frequency: The frequency of the triangle wave in Hz.\n"
-			 "   :type frequency: float\n"
-			 "   :arg rate: The sampling rate in Hz. It's recommended to set this\n"
-			 "      value to the playback device's samling rate to avoid resamping.\n"
-			 "   :type rate: int\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "triangle(frequency, rate=48000)\n\n"
+			 "Creates a triangle sound which plays a triangle wave.\n\n"
+			 ":arg frequency: The frequency of the triangle wave in Hz.\n"
+			 ":type frequency: float\n"
+			 ":arg rate: The sampling rate in Hz. It's recommended to set this "
+			 "value to the playback device's samling rate to avoid resamping.\n"
+			 ":type rate: int\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_triangle(PyTypeObject* type, PyObject* args)
@@ -629,16 +627,14 @@ Sound_triangle(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_accumulate_doc,
-			 ".. classmethod:: accumulate(additive=False)\n\n"
-			 "   Accumulates a sound by summing over positive input\n"
-			 "   differences thus generating a monotonic sigal.\n"
-			 "   If additivity is set to true negative input differences get added too,\n"
-			 "   but positive ones with a factor of two.\n\n"
-			 "   Note that with additivity the signal is not monotonic anymore.\n\n"
-			 "   :arg additive: Whether the accumulation should be additive or not.\n"
-			 "   :type time: bool\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "accumulate(additive=False)\n\n"
+			 "Accumulates a sound by summing over positive input differences thus generating a monotonic sigal. "
+			 "If additivity is set to true negative input differences get added too, but positive ones with a factor of two. "
+			 "Note that with additivity the signal is not monotonic anymore.\n\n"
+			 ":arg additive: Whether the accumulation should be additive or not.\n"
+			 ":type time: bool\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_accumulate(Sound* self, PyObject* args)
@@ -681,19 +677,19 @@ Sound_accumulate(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_ADSR_doc,
-			 ".. classmethod:: ADSR(attack, decay, sustain, release)\n\n"
-			 "   Attack-Decay-Sustain-Release envelopes the volume of a sound.\n"
-			 "   Note: there is currently no way to trigger the release with this API.\n\n"
-			 "   :arg attack: The attack time in seconds.\n"
-			 "   :type attack: float\n"
-			 "   :arg decay: The decay time in seconds.\n"
-			 "   :type decay: float\n"
-			 "   :arg sustain: The sustain level.\n"
-			 "   :type sustain: float\n"
-			 "   :arg release: The release level.\n"
-			 "   :type release: float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "ADSR(attack,decay,sustain,release)\n\n"
+			 "Attack-Decay-Sustain-Release envelopes the volume of a sound. "
+			 "Note: there is currently no way to trigger the release with this API.\n\n"
+			 ":arg attack: The attack time in seconds.\n"
+			 ":type attack: float\n"
+			 ":arg decay: The decay time in seconds.\n"
+			 ":type decay: float\n"
+			 ":arg sustain: The sustain level.\n"
+			 ":type sustain: float\n"
+			 ":arg release: The release level.\n"
+			 ":type release: float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_ADSR(Sound* self, PyObject* args)
@@ -724,12 +720,14 @@ Sound_ADSR(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_delay_doc,
-			 ".. classmethod:: delay(time)\n\n"
-			 "   Delays by playing adding silence in front of the other sound's data.\n\n"
-			 "   :arg time: How many seconds of silence should be added before the sound.\n"
-			 "   :type time: float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "delay(time)\n\n"
+			 "Delays by playing adding silence in front of the other sound's "
+			 "data.\n\n"
+			 ":arg time: How many seconds of silence should be added before "
+			 "the sound.\n"
+			 ":type time: float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_delay(Sound* self, PyObject* args)
@@ -760,18 +758,19 @@ Sound_delay(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_envelope_doc,
-			 ".. classmethod:: envelope(attack, release, threshold, arthreshold)\n\n"
-			 "   Delays by playing adding silence in front of the other sound's data.\n\n"
-			 "   :arg attack: The attack factor.\n"
-			 "   :type attack: float\n"
-			 "   :arg release: The release factor.\n"
-			 "   :type release: float\n"
-			 "   :arg threshold: The general threshold value.\n"
-			 "   :type threshold: float\n"
-			 "   :arg arthreshold: The attack/release threshold value.\n"
-			 "   :type arthreshold: float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "envelope(attack, release, threshold, arthreshold)\n\n"
+			 "Delays by playing adding silence in front of the other sound's "
+			 "data.\n\n"
+			 ":arg attack: The attack factor.\n"
+			 ":type attack: float\n"
+			 ":arg release: The release factor.\n"
+			 ":type release: float\n"
+			 ":arg threshold: The general threshold value.\n"
+			 ":type threshold: float\n"
+			 ":arg arthreshold: The attack/release threshold value.\n"
+			 ":type arthreshold: float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_envelope(Sound* self, PyObject* args)
@@ -802,16 +801,16 @@ Sound_envelope(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_fadein_doc,
-			 ".. classmethod:: fadein(start, length)\n\n"
-			 "   Fades a sound in by raising the volume linearly in the given\n"
-			 "   time interval.\n\n"
-			 "   :arg start: Time in seconds when the fading should start.\n"
-			 "   :type start: float\n"
-			 "   :arg length: Time in seconds how long the fading should last.\n"
-			 "   :type length: float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. note:: Before the fade starts it plays silence.");
+			 "fadein(start, length)\n\n"
+			 "Fades a sound in by raising the volume linearly in the given "
+			 "time interval.\n\n"
+			 ":arg start: Time in seconds when the fading should start.\n"
+			 ":type start: float\n"
+			 ":arg length: Time in seconds how long the fading should last.\n"
+			 ":type length: float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. note:: Before the fade starts it plays silence.");
 
 static PyObject *
 Sound_fadein(Sound* self, PyObject* args)
@@ -842,18 +841,17 @@ Sound_fadein(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_fadeout_doc,
-			 ".. classmethod:: fadeout(start, length)\n\n"
-			 "   Fades a sound in by lowering the volume linearly in the given\n"
-			 "   time interval.\n\n"
-			 "   :arg start: Time in seconds when the fading should start.\n"
-			 "   :type start: float\n"
-			 "   :arg length: Time in seconds how long the fading should last.\n"
-			 "   :type length: float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. note::\n\n"
-			 "      After the fade this sound plays silence, so that\n"
-			 "      the length of the sound is not altered.");
+			 "fadeout(start, length)\n\n"
+			 "Fades a sound in by lowering the volume linearly in the given "
+			 "time interval.\n\n"
+			 ":arg start: Time in seconds when the fading should start.\n"
+			 ":type start: float\n"
+			 ":arg length: Time in seconds how long the fading should last.\n"
+			 ":type length: float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. note:: After the fade this sound plays silence, so that "
+			 "the length of the sound is not altered.");
 
 static PyObject *
 Sound_fadeout(Sound* self, PyObject* args)
@@ -884,20 +882,20 @@ Sound_fadeout(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_filter_doc,
-			 ".. classmethod:: filter(b, a = (1))\n\n"
-			 "   Filters a sound with the supplied IIR filter coefficients.\n"
-			 "   Without the second parameter you'll get a FIR filter.\n\n"
-			 "   If the first value of the a sequence is 0,\n"
-			 "   it will be set to 1 automatically.\n"
-			 "   If the first value of the a sequence is neither 0 nor 1, all\n"
-			 "   filter coefficients will be scaled by this value so that it is 1\n"
-			 "   in the end, you don't have to scale yourself.\n\n"
-			 "   :arg b: The nominator filter coefficients.\n"
-			 "   :type b: sequence of float\n"
-			 "   :arg a: The denominator filter coefficients.\n"
-			 "   :type a: sequence of float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "filter(b, a = (1))\n\n"
+			 "Filters a sound with the supplied IIR filter coefficients.\n"
+			 "Without the second parameter you'll get a FIR filter.\n"
+			 "If the first value of the a sequence is 0 it will be set to 1 "
+			 "automatically.\n"
+			 "If the first value of the a sequence is neither 0 nor 1, all "
+			 "filter coefficients will be scaled by this value so that it is 1 "
+			 "in the end, you don't have to scale yourself.\n\n"
+			 ":arg b: The nominator filter coefficients.\n"
+			 ":type b: sequence of float\n"
+			 ":arg a: The denominator filter coefficients.\n"
+			 ":type a: sequence of float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_filter(Sound* self, PyObject* args)
@@ -984,15 +982,15 @@ Sound_filter(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_highpass_doc,
-			 ".. classmethod:: highpass(frequency, Q=0.5)\n\n"
-			 "   Creates a second order highpass filter based on the transfer\n"
-			 "   function :math:`H(s) = s^2 / (s^2 + s/Q + 1)`\n\n"
-			 "   :arg frequency: The cut off trequency of the highpass.\n"
-			 "   :type frequency: float\n"
-			 "   :arg Q: Q factor of the lowpass.\n"
-			 "   :type Q: float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "highpass(frequency, Q=0.5)\n\n"
+			 "Creates a second order highpass filter based on the transfer "
+			 "function H(s) = s^2 / (s^2 + s/Q + 1)\n\n"
+			 ":arg frequency: The cut off trequency of the highpass.\n"
+			 ":type frequency: float\n"
+			 ":arg Q: Q factor of the lowpass.\n"
+			 ":type Q: float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_highpass(Sound* self, PyObject* args)
@@ -1024,14 +1022,14 @@ Sound_highpass(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_limit_doc,
-			 ".. classmethod:: limit(start, end)\n\n"
-			 "   Limits a sound within a specific start and end time.\n\n"
-			 "   :arg start: Start time in seconds.\n"
-			 "   :type start: float\n"
-			 "   :arg end: End time in seconds.\n"
-			 "   :type end: float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "limit(start, end)\n\n"
+			 "Limits a sound within a specific start and end time.\n\n"
+			 ":arg start: Start time in seconds.\n"
+			 ":type start: float\n"
+			 ":arg end: End time in seconds.\n"
+			 ":type end: float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_limit(Sound* self, PyObject* args)
@@ -1062,16 +1060,15 @@ Sound_limit(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_loop_doc,
-			 ".. classmethod:: loop(count)\n\n"
-			 "   Loops a sound.\n\n"
-			 "   :arg count: How often the sound should be looped.\n"
-			 "      Negative values mean endlessly.\n"
-			 "   :type count: integer\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. note::\n\n"
-			 "      This is a filter function, you might consider using\n"
-			 "      :attr:`Handle.loop_count` instead.");
+			 "loop(count)\n\n"
+			 "Loops a sound.\n\n"
+			 ":arg count: How often the sound should be looped. "
+			 "Negative values mean endlessly.\n"
+			 ":type count: integer\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. note:: This is a filter function, you might consider using "
+			 ":attr:`Handle.loop_count` instead.");
 
 static PyObject *
 Sound_loop(Sound* self, PyObject* args)
@@ -1102,15 +1099,15 @@ Sound_loop(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_lowpass_doc,
-			 ".. classmethod:: lowpass(frequency, Q=0.5)\n\n"
-			 "   Creates a second order lowpass filter based on the transfer "
-			 "   function :math:`H(s) = 1 / (s^2 + s/Q + 1)`\n\n"
-			 "   :arg frequency: The cut off trequency of the lowpass.\n"
-			 "   :type frequency: float\n"
-			 "   :arg Q: Q factor of the lowpass.\n"
-			 "   :type Q: float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "lowpass(frequency, Q=0.5)\n\n"
+			 "Creates a second order lowpass filter based on the transfer "
+			 "function H(s) = 1 / (s^2 + s/Q + 1)\n\n"
+			 ":arg frequency: The cut off trequency of the lowpass.\n"
+			 ":type frequency: float\n"
+			 ":arg Q: Q factor of the lowpass.\n"
+			 ":type Q: float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_lowpass(Sound* self, PyObject* args)
@@ -1142,15 +1139,14 @@ Sound_lowpass(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_modulate_doc,
-			 ".. classmethod:: modulate(sound)\n\n"
-			 "   Modulates two factories.\n\n"
-			 "   :arg sound: The sound to modulate over the other.\n"
-			 "   :type sound: :class:`Sound`\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. note::\n\n"
-			 "      The two factories have to have the same specifications\n"
-			 "      (channels and samplerate).");
+			 "modulate(sound)\n\n"
+			 "Modulates two factories.\n\n"
+			 ":arg sound: The sound to modulate over the other.\n"
+			 ":type sound: :class:`Sound`\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. note:: The two factories have to have the same specifications "
+			 "(channels and samplerate).");
 
 static PyObject *
 Sound_modulate(Sound* self, PyObject* object)
@@ -1184,19 +1180,17 @@ Sound_modulate(Sound* self, PyObject* object)
 }
 
 PyDoc_STRVAR(M_aud_Sound_pitch_doc,
-			 ".. classmethod:: pitch(factor)\n\n"
-			 "   Changes the pitch of a sound with a specific factor.\n\n"
-			 "   :arg factor: The factor to change the pitch with.\n"
-			 "   :type factor: float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. note::\n\n"
-			 "      This is done by changing the sample rate of the\n"
-			 "      underlying sound, which has to be an integer, so the factor\n"
-			 "      value rounded and the factor may not be 100 % accurate.\n\n"
-			 "   .. note::\n\n"
-			 "      This is a filter function, you might consider using\n"
-			 "      :attr:`Handle.pitch` instead.");
+			 "pitch(factor)\n\n"
+			 "Changes the pitch of a sound with a specific factor.\n\n"
+			 ":arg factor: The factor to change the pitch with.\n"
+			 ":type factor: float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. note:: This is done by changing the sample rate of the "
+			 "underlying sound, which has to be an integer, so the factor "
+			 "value rounded and the factor may not be 100 % accurate.\n\n"
+			 ".. note:: This is a filter function, you might consider using "
+			 ":attr:`Handle.pitch` instead.");
 
 static PyObject *
 Sound_pitch(Sound* self, PyObject* args)
@@ -1227,12 +1221,12 @@ Sound_pitch(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_rechannel_doc,
-			 ".. classmethod:: rechannel(channels)\n\n"
-			 "   Rechannels the sound.\n\n"
-			 "   :arg channels: The new channel configuration.\n"
-			 "   :type channels: int\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "rechannel(channels)\n\n"
+			 "Rechannels the sound.\n\n"
+			 ":arg channels: The new channel configuration.\n"
+			 ":type channels: int\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_rechannel(Sound* self, PyObject* args)
@@ -1267,14 +1261,14 @@ Sound_rechannel(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_resample_doc,
-			 ".. classmethod:: resample(rate, high_quality)\n\n"
-			 "   Resamples the sound.\n\n"
-			 "   :arg rate: The new sample rate.\n"
-			 "   :type rate: double\n"
-			 "   :arg high_quality: When true use a higher quality but slower resampler.\n"
-			 "   :type high_quality: bool\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "resample(rate, high_quality)\n\n"
+			 "Resamples the sound.\n\n"
+			 ":arg rate: The new sample rate.\n"
+			 ":type rate: double\n"
+			 ":arg high_quality: When true use a higher quality but slower resampler.\n"
+			 ":type high_quality: bool\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_resample(Sound* self, PyObject* args)
@@ -1322,19 +1316,17 @@ Sound_resample(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_reverse_doc,
-			 ".. classmethod:: reverse()\n\n"
-			 "   Plays a sound reversed.\n\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. note::\n\n"
-			 "      The sound has to have a finite length and has to be seekable.\n"
-			 "      It's recommended to use this only with factories with\n"
-			 "      fast and accurate seeking, which is not true for encoded audio\n"
-			 "      files, such ones should be buffered using :meth:`cache` before\n"
-			 "      being played reversed.\n\n"
-			 "   .. warning::\n\n"
-			 "      If seeking is not accurate in the underlying sound\n"
-			 "      you'll likely hear skips/jumps/cracks.");
+			 "reverse()\n\n"
+			 "Plays a sound reversed.\n\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. note:: The sound has to have a finite length and has to be "
+			 "seekable. It's recommended to use this only with factories	 with "
+			 "fast and accurate seeking, which is not true for encoded audio "
+			 "files, such ones should be buffered using :meth:`cache` before "
+			 "being played reversed.\n\n"
+			 ".. warning:: If seeking is not accurate in the underlying sound "
+			 "you'll likely hear skips/jumps/cracks.");
 
 static PyObject *
 Sound_reverse(Sound* self)
@@ -1360,10 +1352,10 @@ Sound_reverse(Sound* self)
 }
 
 PyDoc_STRVAR(M_aud_Sound_sum_doc,
-			 ".. classmethod:: sum()\n\n"
-			 "   Sums the samples of a sound.\n\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "sum()\n\n"
+			 "Sums the samples of a sound.\n\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_sum(Sound* self)
@@ -1389,12 +1381,12 @@ Sound_sum(Sound* self)
 }
 
 PyDoc_STRVAR(M_aud_Sound_threshold_doc,
-			 ".. classmethod:: threshold(threshold = 0)\n\n"
-			 "   Makes a threshold wave out of an audio wave by setting all samples\n"
-			 "   with a amplitude >= threshold to 1, all <= -threshold to -1 and\n"
-			 "   all between to 0.\n\n"
-			 "   :arg threshold: Threshold value over which an amplitude counts\n"
-			 "      non-zero.\n"
+			 "threshold(threshold = 0)\n\n"
+			 "Makes a threshold wave out of an audio wave by setting all samples "
+			 "with a amplitude >= threshold to 1, all <= -threshold to -1 and "
+			 "all between to 0.\n\n"
+			 ":arg threshold: Threshold value over which an amplitude counts "
+			 "non-zero.\n"
 			 ":type threshold: float\n"
 			 ":return: The created :class:`Sound` object.\n"
 			 ":rtype: :class:`Sound`");
@@ -1428,16 +1420,15 @@ Sound_threshold(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_volume_doc,
-			 ".. classmethod:: volume(volume)\n\n"
-			 "   Changes the volume of a sound.\n\n"
-			 "   :arg volume: The new volume..\n"
-			 "   :type volume: float\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. note:: Should be in the range [0, 1] to avoid clipping.\n\n"
-			 "   .. note::\n\n"
-			 "      This is a filter function, you might consider using\n"
-			 "      :attr:`Handle.volume` instead.");
+			 "volume(volume)\n\n"
+			 "Changes the volume of a sound.\n\n"
+			 ":arg volume: The new volume..\n"
+			 ":type volume: float\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. note:: Should be in the range [0, 1] to avoid clipping.\n\n"
+			 ".. note:: This is a filter function, you might consider using "
+			 ":attr:`Handle.volume` instead.");
 
 static PyObject *
 Sound_volume(Sound* self, PyObject* args)
@@ -1468,15 +1459,14 @@ Sound_volume(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_join_doc,
-			 ".. classmethod:: join(sound)\n\n"
-			 "   Plays two factories in sequence.\n\n"
-			 "   :arg sound: The sound to play second.\n"
-			 "   :type sound: :class:`Sound`\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. note::\n\n"
-			 "   The two factories have to have the same specifications\n"
-			 "   (channels and samplerate).");
+			 "join(sound)\n\n"
+			 "Plays two factories in sequence.\n\n"
+			 ":arg sound: The sound to play second.\n"
+			 ":type sound: :class:`Sound`\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. note:: The two factories have to have the same specifications "
+			 "(channels and samplerate).");
 
 static PyObject *
 Sound_join(Sound* self, PyObject* object)
@@ -1511,15 +1501,14 @@ Sound_join(Sound* self, PyObject* object)
 }
 
 PyDoc_STRVAR(M_aud_Sound_mix_doc,
-			 ".. classmethod:: mix(sound)\n\n"
-			 "   Mixes two factories.\n\n"
-			 "   :arg sound: The sound to mix over the other.\n"
-			 "   :type sound: :class:`Sound`\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`\n\n"
-			 "   .. note::\n\n"
-			 "      The two factories have to have the same specifications\n"
-			 "      (channels and samplerate).");
+			 "mix(sound)\n\n"
+			 "Mixes two factories.\n\n"
+			 ":arg sound: The sound to mix over the other.\n"
+			 ":type sound: :class:`Sound`\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`\n\n"
+			 ".. note:: The two factories have to have the same specifications "
+			 "(channels and samplerate).");
 
 static PyObject *
 Sound_mix(Sound* self, PyObject* object)
@@ -1553,11 +1542,11 @@ Sound_mix(Sound* self, PyObject* object)
 }
 
 PyDoc_STRVAR(M_aud_Sound_pingpong_doc,
-			 ".. classmethod:: pingpong()\n\n"
-			 "   Plays a sound forward and then backward.\n"
-			 "   This is like joining a sound with its reverse.\n\n"
-			 "   :return: The created :class:`Sound` object.\n"
-			 "   :rtype: :class:`Sound`");
+			 "pingpong()\n\n"
+			 "Plays a sound forward and then backward.\n"
+			 "This is like joining a sound with its reverse.\n\n"
+			 ":return: The created :class:`Sound` object.\n"
+			 ":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_pingpong(Sound* self)
@@ -1583,12 +1572,12 @@ Sound_pingpong(Sound* self)
 }
 
 PyDoc_STRVAR(M_aud_Sound_list_doc,
-	".. classmethod:: list()\n\n"
-	"   Creates an empty sound list that can contain several sounds.\n\n"
-	"   :arg random: whether the playback will be random or not.\n"
-	"   :type random: int\n"
-	"   :return: The created :class:`Sound` object.\n"
-	"   :rtype: :class:`Sound`");
+	"list()\n\n"
+	"Creates an empty sound list that can contain several sounds.\n\n"
+	":arg random: wether the playback will be random or not.\n"
+	":type random: int\n"
+	":return: The created :class:`Sound` object.\n"
+	":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_list(PyTypeObject* type, PyObject* args)
@@ -1619,11 +1608,11 @@ Sound_list(PyTypeObject* type, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_mutable_doc,
-	".. classmethod:: mutable()\n\n"
-	"   Creates a sound that will be restarted when sought backwards.\n"
-	"   If the original sound is a sound list, the playing sound can change.\n\n"
-	"   :return: The created :class:`Sound` object.\n"
-	"   :rtype: :class:`Sound`");
+	"mutable()\n\n"
+	"Creates a sound that will be restarted when sought backwards.\n"
+	"If the original sound is a sound list, the playing sound can change.\n\n"
+	":return: The created :class:`Sound` object.\n"
+	":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_mutable(Sound* self)
@@ -1649,12 +1638,12 @@ Sound_mutable(Sound* self)
 }
 
 PyDoc_STRVAR(M_aud_Sound_list_addSound_doc,
-	".. classmethod:: addSound(sound)\n\n"
-	"   Adds a new sound to a sound list.\n\n"
-	"   :arg sound: The sound that will be added to the list.\n"
-	"   :type sound: :class:`Sound`\n\n"
-	"   .. note:: You can only add a sound to a sound list.");
-
+	"addSound(sound)\n\n"
+	"Adds a new sound to a sound list.\n\n"
+	":arg sound: The sound that will be added to the list.\n"
+	":type sound: :class:`Sound`\n\n"
+	".. note:: You can only add a sound to a sound list.");
+	
 static PyObject *
 Sound_list_addSound(Sound* self, PyObject* object)
 {
@@ -1682,14 +1671,14 @@ Sound_list_addSound(Sound* self, PyObject* object)
 #ifdef WITH_CONVOLUTION
 
 PyDoc_STRVAR(M_aud_Sound_convolver_doc,
-	".. classmethod:: convolver()\n\n"
-	"   Creates a sound that will apply convolution to another sound.\n\n"
-	"   :arg impulseResponse: The filter with which convolve the sound.\n"
-	"   :type impulseResponse: :class:`ImpulseResponse`\n"
-	"   :arg threadPool: A thread pool used to parallelize convolution.\n"
-	"   :type threadPool: :class:`ThreadPool`\n"
-	"   :return: The created :class:`Sound` object.\n"
-	"   :rtype: :class:`Sound`");
+	"convolver()\n\n"
+	"Creates a sound that will apply convolution to another sound.\n\n"
+	":arg impulseResponse: The filter with which convolve the sound.\n"
+	":type impulseResponse: :class:`ImpulseResponse`\n"
+	":arg threadPool: A thread pool used to parallelize convolution.\n"
+	":type threadPool: :class:`ThreadPool`\n"
+	":return: The created :class:`Sound` object.\n"
+	":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_convolver(Sound* self, PyObject* args)
@@ -1731,16 +1720,16 @@ Sound_convolver(Sound* self, PyObject* args)
 }
 
 PyDoc_STRVAR(M_aud_Sound_binaural_doc,
-	".. classmethod:: convolver()\n\n"
-	"   Creates a binaural sound using another sound as source. The original sound must be mono\n\n"
-	"   :arg hrtfs: An HRTF set.\n"
-	"   :type hrtf: :class:`HRTF`\n"
-	"   :arg source: An object representing the source position of the sound.\n"
-	"   :type source: :class:`Source`\n"
-	"   :arg threadPool: A thread pool used to parallelize convolution.\n"
-	"   :type threadPool: :class:`ThreadPool`\n"
-	"   :return: The created :class:`Sound` object.\n"
-	"   :rtype: :class:`Sound`");
+	"convolver()\n\n"
+	"Creates a binaural sound using another sound as source. The original sound must be mono\n\n"
+	":arg hrtfs: An HRTF set.\n"
+	":type hrtf: :class:`HRTF`\n"
+	":arg source: An object representing the source position of the sound.\n"
+	":type source: :class:`Source`\n"
+	":arg threadPool: A thread pool used to parallelize convolution.\n"
+	":type threadPool: :class:`ThreadPool`\n"
+	":return: The created :class:`Sound` object.\n"
+	":rtype: :class:`Sound`");
 
 static PyObject *
 Sound_binaural(Sound* self, PyObject* args)

extern/mantaflow/UPDATE.sh

@@ -71,7 +71,7 @@ rm $BLENDER_INSTALLATION/blender/tmp/preprocessed/fileio/*.reg
 cd $BLENDER_INSTALLATION/blender/tmp/
 
 echo "Applying clang format to Mantaflow source files"
-find . -iname *.h -o -iname *.cpp | xargs clang-format --verbose -i -style=file -sort-includes=0
+find . -iname *.h -o -iname *.cpp | xargs clang-format --verbose -i -style=file
 find . -iname *.h -o -iname *.cpp | xargs dos2unix --verbose
 
 # ==================== 5) MOVE MANTAFLOW FILES TO EXTERN/ ================================

extern/mantaflow/preprocessed/gitinfo.h

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

extern/mantaflow/preprocessed/grid.cpp

@@ -853,147 +853,6 @@ template<class T> struct knPermuteAxes : public KernelBase {
   int axis2;
 };
 
-struct knJoinVec : public KernelBase {
-  knJoinVec(Grid<Vec3> &a, const Grid<Vec3> &b, bool keepMax)
-      : KernelBase(&a, 0), a(a), b(b), keepMax(keepMax)
-  {
-    runMessage();
-    run();
-  }
-  inline void op(IndexInt idx, Grid<Vec3> &a, const Grid<Vec3> &b, bool keepMax) const
-  {
-    Real a1 = normSquare(a[idx]);
-    Real b1 = normSquare(b[idx]);
-    a[idx] = (keepMax) ? max(a1, b1) : min(a1, b1);
-  }
-  inline Grid<Vec3> &getArg0()
-  {
-    return a;
-  }
-  typedef Grid<Vec3> type0;
-  inline const Grid<Vec3> &getArg1()
-  {
-    return b;
-  }
-  typedef Grid<Vec3> type1;
-  inline bool &getArg2()
-  {
-    return keepMax;
-  }
-  typedef bool type2;
-  void runMessage()
-  {
-    debMsg("Executing kernel knJoinVec ", 3);
-    debMsg("Kernel range"
-               << " x " << maxX << " y " << maxY << " z " << minZ << " - " << maxZ << " ",
-           4);
-  };
-  void operator()(const tbb::blocked_range<IndexInt> &__r) const
-  {
-    for (IndexInt idx = __r.begin(); idx != (IndexInt)__r.end(); idx++)
-      op(idx, a, b, keepMax);
-  }
-  void run()
-  {
-    tbb::parallel_for(tbb::blocked_range<IndexInt>(0, size), *this);
-  }
-  Grid<Vec3> &a;
-  const Grid<Vec3> &b;
-  bool keepMax;
-};
-struct knJoinInt : public KernelBase {
-  knJoinInt(Grid<int> &a, const Grid<int> &b, bool keepMax)
-      : KernelBase(&a, 0), a(a), b(b), keepMax(keepMax)
-  {
-    runMessage();
-    run();
-  }
-  inline void op(IndexInt idx, Grid<int> &a, const Grid<int> &b, bool keepMax) const
-  {
-    a[idx] = (keepMax) ? max(a[idx], b[idx]) : min(a[idx], b[idx]);
-  }
-  inline Grid<int> &getArg0()
-  {
-    return a;
-  }
-  typedef Grid<int> type0;
-  inline const Grid<int> &getArg1()
-  {
-    return b;
-  }
-  typedef Grid<int> type1;
-  inline bool &getArg2()
-  {
-    return keepMax;
-  }
-  typedef bool type2;
-  void runMessage()
-  {
-    debMsg("Executing kernel knJoinInt ", 3);
-    debMsg("Kernel range"
-               << " x " << maxX << " y " << maxY << " z " << minZ << " - " << maxZ << " ",
-           4);
-  };
-  void operator()(const tbb::blocked_range<IndexInt> &__r) const
-  {
-    for (IndexInt idx = __r.begin(); idx != (IndexInt)__r.end(); idx++)
-      op(idx, a, b, keepMax);
-  }
-  void run()
-  {
-    tbb::parallel_for(tbb::blocked_range<IndexInt>(0, size), *this);
-  }
-  Grid<int> &a;
-  const Grid<int> &b;
-  bool keepMax;
-};
-struct knJoinReal : public KernelBase {
-  knJoinReal(Grid<Real> &a, const Grid<Real> &b, bool keepMax)
-      : KernelBase(&a, 0), a(a), b(b), keepMax(keepMax)
-  {
-    runMessage();
-    run();
-  }
-  inline void op(IndexInt idx, Grid<Real> &a, const Grid<Real> &b, bool keepMax) const
-  {
-    a[idx] = (keepMax) ? max(a[idx], b[idx]) : min(a[idx], b[idx]);
-  }
-  inline Grid<Real> &getArg0()
-  {
-    return a;
-  }
-  typedef Grid<Real> type0;
-  inline const Grid<Real> &getArg1()
-  {
-    return b;
-  }
-  typedef Grid<Real> type1;
-  inline bool &getArg2()
-  {
-    return keepMax;
-  }
-  typedef bool type2;
-  void runMessage()
-  {
-    debMsg("Executing kernel knJoinReal ", 3);
-    debMsg("Kernel range"
-               << " x " << maxX << " y " << maxY << " z " << minZ << " - " << maxZ << " ",
-           4);
-  };
-  void operator()(const tbb::blocked_range<IndexInt> &__r) const
-  {
-    for (IndexInt idx = __r.begin(); idx != (IndexInt)__r.end(); idx++)
-      op(idx, a, b, keepMax);
-  }
-  void run()
-  {
-    tbb::parallel_for(tbb::blocked_range<IndexInt>(0, size), *this);
-  }
-  Grid<Real> &a;
-  const Grid<Real> &b;
-  bool keepMax;
-};
-
 template<class T> Grid<T> &Grid<T>::safeDivide(const Grid<T> &a)
 {
   knGridSafeDiv<T>(*this, a);
@@ -1069,18 +928,6 @@ void Grid<T>::permuteAxesCopyToGrid(int axis0, int axis1, int axis2, Grid<T> &ou
             "Permuted grids must have the same dimensions!");
   knPermuteAxes<T>(*this, out, axis0, axis1, axis2);
 }
-template<> void Grid<Vec3>::join(const Grid<Vec3> &a, bool keepMax)
-{
-  knJoinVec(*this, a, keepMax);
-}
-template<> void Grid<int>::join(const Grid<int> &a, bool keepMax)
-{
-  knJoinInt(*this, a, keepMax);
-}
-template<> void Grid<Real>::join(const Grid<Real> &a, bool keepMax)
-{
-  knJoinReal(*this, a, keepMax);
-}
 
 template<> Real Grid<Real>::getMax() const
 {

extern/mantaflow/preprocessed/grid.h

@@ -966,38 +966,10 @@ template<class T> class Grid : public GridBase {
     }
   }
 
-  //! join other grid by either keeping min or max value at cell
-  void join(const Grid<T> &a, bool keepMax = true);
-  static PyObject *_W_27(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
-  {
-    try {
-      PbArgs _args(_linargs, _kwds);
-      Grid *pbo = dynamic_cast<Grid *>(Pb::objFromPy(_self));
-      bool noTiming = _args.getOpt<bool>("notiming", -1, 0);
-      pbPreparePlugin(pbo->getParent(), "Grid::join", !noTiming);
-      PyObject *_retval = 0;
-      {
-        ArgLocker _lock;
-        const Grid<T> &a = *_args.getPtr<Grid<T>>("a", 0, &_lock);
-        bool keepMax = _args.getOpt<bool>("keepMax", 1, true, &_lock);
-        pbo->_args.copy(_args);
-        _retval = getPyNone();
-        pbo->join(a, keepMax);
-        pbo->_args.check();
-      }
-      pbFinalizePlugin(pbo->getParent(), "Grid::join", !noTiming);
-      return _retval;
-    }
-    catch (std::exception &e) {
-      pbSetError("Grid::join", e.what());
-      return 0;
-    }
-  }
-
   // common compound operators
   //! get absolute max value in grid
   Real getMaxAbs() const;
-  static PyObject *_W_28(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_27(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1022,7 +994,7 @@ template<class T> class Grid : public GridBase {
 
   //! get max value in grid
   Real getMax() const;
-  static PyObject *_W_29(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_28(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1047,7 +1019,7 @@ template<class T> class Grid : public GridBase {
 
   //! get min value in grid
   Real getMin() const;
-  static PyObject *_W_30(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_29(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1072,7 +1044,7 @@ template<class T> class Grid : public GridBase {
 
   //! calculate L1 norm of grid content
   Real getL1(int bnd = 0);
-  static PyObject *_W_31(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_30(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1098,7 +1070,7 @@ template<class T> class Grid : public GridBase {
 
   //! calculate L2 norm of grid content
   Real getL2(int bnd = 0);
-  static PyObject *_W_32(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_31(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1124,7 +1096,7 @@ template<class T> class Grid : public GridBase {
 
   //! set all boundary cells to constant value (Dirichlet)
   void setBound(T value, int boundaryWidth = 1);
-  static PyObject *_W_33(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_32(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1152,7 +1124,7 @@ template<class T> class Grid : public GridBase {
 
   //! set all boundary cells to last inner value (Neumann)
   void setBoundNeumann(int boundaryWidth = 1);
-  static PyObject *_W_34(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_33(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1179,7 +1151,7 @@ template<class T> class Grid : public GridBase {
 
   //! get data pointer of grid
   std::string getDataPointer();
-  static PyObject *_W_35(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_34(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1204,7 +1176,7 @@ template<class T> class Grid : public GridBase {
 
   //! debugging helper, print grid from python. skip boundary of width bnd
   void printGrid(int zSlice = -1, bool printIndex = false, int bnd = 1);
-  static PyObject *_W_36(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_35(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1272,7 +1244,7 @@ class MACGrid : public Grid<Vec3> {
   {
     mType = (GridType)(TypeMAC | TypeVec3);
   }
-  static int _W_37(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static int _W_36(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     PbClass *obj = Pb::objFromPy(_self);
     if (obj)
@@ -1354,7 +1326,7 @@ class MACGrid : public Grid<Vec3> {
   //! set all boundary cells of a MAC grid to certain value (Dirchlet). Respects staggered grid
   //! locations optionally, only set normal components
   void setBoundMAC(Vec3 value, int boundaryWidth, bool normalOnly = false);
-  static PyObject *_W_38(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_37(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1395,7 +1367,7 @@ class FlagGrid : public Grid<int> {
   {
     mType = (GridType)(TypeFlags | TypeInt);
   }
-  static int _W_39(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static int _W_38(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     PbClass *obj = Pb::objFromPy(_self);
     if (obj)
@@ -1575,7 +1547,7 @@ class FlagGrid : public Grid<int> {
                   const std::string &inflow = "      ",
                   const std::string &outflow = "      ",
                   Grid<Real> *phiWalls = 0x00);
-  static PyObject *_W_40(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_39(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1609,7 +1581,7 @@ class FlagGrid : public Grid<int> {
 
   //! set fluid flags inside levelset (liquids)
   void updateFromLevelset(LevelsetGrid &levelset);
-  static PyObject *_W_41(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_40(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1636,7 +1608,7 @@ class FlagGrid : public Grid<int> {
 
   //! set all cells (except obs/in/outflow) to type (fluid by default)
   void fillGrid(int type = TypeFluid);
-  static PyObject *_W_42(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_41(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);
@@ -1665,7 +1637,7 @@ class FlagGrid : public Grid<int> {
   //! warning for large grids! only regular int returned (due to python interface)
   //! optionally creates mask in RealGrid (1 where flag matches, 0 otherwise)
   int countCells(int flag, int bnd = 0, Grid<Real> *mask = NULL);
-  static PyObject *_W_43(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
+  static PyObject *_W_42(PyObject *_self, PyObject *_linargs, PyObject *_kwds)
   {
     try {
       PbArgs _args(_linargs, _kwds);

extern/mantaflow/preprocessed/grid.h.reg.cpp

@@ -8,11 +8,11 @@ namespace Manta {
 #ifdef _C_FlagGrid
 static const Pb::Register _R_26("FlagGrid", "FlagGrid", "Grid<int>");
 template<> const char *Namify<FlagGrid>::S = "FlagGrid";
-static const Pb::Register _R_27("FlagGrid", "FlagGrid", FlagGrid::_W_39);
-static const Pb::Register _R_28("FlagGrid", "initDomain", FlagGrid::_W_40);
-static const Pb::Register _R_29("FlagGrid", "updateFromLevelset", FlagGrid::_W_41);
-static const Pb::Register _R_30("FlagGrid", "fillGrid", FlagGrid::_W_42);
-static const Pb::Register _R_31("FlagGrid", "countCells", FlagGrid::_W_43);
+static const Pb::Register _R_27("FlagGrid", "FlagGrid", FlagGrid::_W_38);
+static const Pb::Register _R_28("FlagGrid", "initDomain", FlagGrid::_W_39);
+static const Pb::Register _R_29("FlagGrid", "updateFromLevelset", FlagGrid::_W_40);
+static const Pb::Register _R_30("FlagGrid", "fillGrid", FlagGrid::_W_41);
+static const Pb::Register _R_31("FlagGrid", "countCells", FlagGrid::_W_42);
 #endif
 #ifdef _C_Grid
 static const Pb::Register _R_32("Grid<int>", "Grid<int>", "GridBase");
@@ -35,95 +35,92 @@ static const Pb::Register _R_47("Grid<int>", "clamp", Grid<int>::_W_23);
 static const Pb::Register _R_48("Grid<int>", "stomp", Grid<int>::_W_24);
 static const Pb::Register _R_49("Grid<int>", "permuteAxes", Grid<int>::_W_25);
 static const Pb::Register _R_50("Grid<int>", "permuteAxesCopyToGrid", Grid<int>::_W_26);
-static const Pb::Register _R_51("Grid<int>", "join", Grid<int>::_W_27);
-static const Pb::Register _R_52("Grid<int>", "getMaxAbs", Grid<int>::_W_28);
-static const Pb::Register _R_53("Grid<int>", "getMax", Grid<int>::_W_29);
-static const Pb::Register _R_54("Grid<int>", "getMin", Grid<int>::_W_30);
-static const Pb::Register _R_55("Grid<int>", "getL1", Grid<int>::_W_31);
-static const Pb::Register _R_56("Grid<int>", "getL2", Grid<int>::_W_32);
-static const Pb::Register _R_57("Grid<int>", "setBound", Grid<int>::_W_33);
-static const Pb::Register _R_58("Grid<int>", "setBoundNeumann", Grid<int>::_W_34);
-static const Pb::Register _R_59("Grid<int>", "getDataPointer", Grid<int>::_W_35);
-static const Pb::Register _R_60("Grid<int>", "printGrid", Grid<int>::_W_36);
-static const Pb::Register _R_61("Grid<Real>", "Grid<Real>", "GridBase");
+static const Pb::Register _R_51("Grid<int>", "getMaxAbs", Grid<int>::_W_27);
+static const Pb::Register _R_52("Grid<int>", "getMax", Grid<int>::_W_28);
+static const Pb::Register _R_53("Grid<int>", "getMin", Grid<int>::_W_29);
+static const Pb::Register _R_54("Grid<int>", "getL1", Grid<int>::_W_30);
+static const Pb::Register _R_55("Grid<int>", "getL2", Grid<int>::_W_31);
+static const Pb::Register _R_56("Grid<int>", "setBound", Grid<int>::_W_32);
+static const Pb::Register _R_57("Grid<int>", "setBoundNeumann", Grid<int>::_W_33);
+static const Pb::Register _R_58("Grid<int>", "getDataPointer", Grid<int>::_W_34);
+static const Pb::Register _R_59("Grid<int>", "printGrid", Grid<int>::_W_35);
+static const Pb::Register _R_60("Grid<Real>", "Grid<Real>", "GridBase");
 template<> const char *Namify<Grid<Real>>::S = "Grid<Real>";
-static const Pb::Register _R_62("Grid<Real>", "Grid", Grid<Real>::_W_9);
-static const Pb::Register _R_63("Grid<Real>", "save", Grid<Real>::_W_10);
-static const Pb::Register _R_64("Grid<Real>", "load", Grid<Real>::_W_11);
-static const Pb::Register _R_65("Grid<Real>", "clear", Grid<Real>::_W_12);
-static const Pb::Register _R_66("Grid<Real>", "copyFrom", Grid<Real>::_W_13);
-static const Pb::Register _R_67("Grid<Real>", "getGridType", Grid<Real>::_W_14);
-static const Pb::Register _R_68("Grid<Real>", "add", Grid<Real>::_W_15);
-static const Pb::Register _R_69("Grid<Real>", "sub", Grid<Real>::_W_16);
-static const Pb::Register _R_70("Grid<Real>", "setConst", Grid<Real>::_W_17);
-static const Pb::Register _R_71("Grid<Real>", "addConst", Grid<Real>::_W_18);
-static const Pb::Register _R_72("Grid<Real>", "addScaled", Grid<Real>::_W_19);
-static const Pb::Register _R_73("Grid<Real>", "mult", Grid<Real>::_W_20);
-static const Pb::Register _R_74("Grid<Real>", "multConst", Grid<Real>::_W_21);
-static const Pb::Register _R_75("Grid<Real>", "safeDivide", Grid<Real>::_W_22);
-static const Pb::Register _R_76("Grid<Real>", "clamp", Grid<Real>::_W_23);
-static const Pb::Register _R_77("Grid<Real>", "stomp", Grid<Real>::_W_24);
-static const Pb::Register _R_78("Grid<Real>", "permuteAxes", Grid<Real>::_W_25);
-static const Pb::Register _R_79("Grid<Real>", "permuteAxesCopyToGrid", Grid<Real>::_W_26);
-static const Pb::Register _R_80("Grid<Real>", "join", Grid<Real>::_W_27);
-static const Pb::Register _R_81("Grid<Real>", "getMaxAbs", Grid<Real>::_W_28);
-static const Pb::Register _R_82("Grid<Real>", "getMax", Grid<Real>::_W_29);
-static const Pb::Register _R_83("Grid<Real>", "getMin", Grid<Real>::_W_30);
-static const Pb::Register _R_84("Grid<Real>", "getL1", Grid<Real>::_W_31);
-static const Pb::Register _R_85("Grid<Real>", "getL2", Grid<Real>::_W_32);
-static const Pb::Register _R_86("Grid<Real>", "setBound", Grid<Real>::_W_33);
-static const Pb::Register _R_87("Grid<Real>", "setBoundNeumann", Grid<Real>::_W_34);
-static const Pb::Register _R_88("Grid<Real>", "getDataPointer", Grid<Real>::_W_35);
-static const Pb::Register _R_89("Grid<Real>", "printGrid", Grid<Real>::_W_36);
-static const Pb::Register _R_90("Grid<Vec3>", "Grid<Vec3>", "GridBase");
+static const Pb::Register _R_61("Grid<Real>", "Grid", Grid<Real>::_W_9);
+static const Pb::Register _R_62("Grid<Real>", "save", Grid<Real>::_W_10);
+static const Pb::Register _R_63("Grid<Real>", "load", Grid<Real>::_W_11);
+static const Pb::Register _R_64("Grid<Real>", "clear", Grid<Real>::_W_12);
+static const Pb::Register _R_65("Grid<Real>", "copyFrom", Grid<Real>::_W_13);
+static const Pb::Register _R_66("Grid<Real>", "getGridType", Grid<Real>::_W_14);
+static const Pb::Register _R_67("Grid<Real>", "add", Grid<Real>::_W_15);
+static const Pb::Register _R_68("Grid<Real>", "sub", Grid<Real>::_W_16);
+static const Pb::Register _R_69("Grid<Real>", "setConst", Grid<Real>::_W_17);
+static const Pb::Register _R_70("Grid<Real>", "addConst", Grid<Real>::_W_18);
+static const Pb::Register _R_71("Grid<Real>", "addScaled", Grid<Real>::_W_19);
+static const Pb::Register _R_72("Grid<Real>", "mult", Grid<Real>::_W_20);
+static const Pb::Register _R_73("Grid<Real>", "multConst", Grid<Real>::_W_21);
+static const Pb::Register _R_74("Grid<Real>", "safeDivide", Grid<Real>::_W_22);
+static const Pb::Register _R_75("Grid<Real>", "clamp", Grid<Real>::_W_23);
+static const Pb::Register _R_76("Grid<Real>", "stomp", Grid<Real>::_W_24);
+static const Pb::Register _R_77("Grid<Real>", "permuteAxes", Grid<Real>::_W_25);
+static const Pb::Register _R_78("Grid<Real>", "permuteAxesCopyToGrid", Grid<Real>::_W_26);
+static const Pb::Register _R_79("Grid<Real>", "getMaxAbs", Grid<Real>::_W_27);
+static const Pb::Register _R_80("Grid<Real>", "getMax", Grid<Real>::_W_28);
+static const Pb::Register _R_81("Grid<Real>", "getMin", Grid<Real>::_W_29);
+static const Pb::Register _R_82("Grid<Real>", "getL1", Grid<Real>::_W_30);
+static const Pb::Register _R_83("Grid<Real>", "getL2", Grid<Real>::_W_31);
+static const Pb::Register _R_84("Grid<Real>", "setBound", Grid<Real>::_W_32);
+static const Pb::Register _R_85("Grid<Real>", "setBoundNeumann", Grid<Real>::_W_33);
+static const Pb::Register _R_86("Grid<Real>", "getDataPointer", Grid<Real>::_W_34);
+static const Pb::Register _R_87("Grid<Real>", "printGrid", Grid<Real>::_W_35);
+static const Pb::Register _R_88("Grid<Vec3>", "Grid<Vec3>", "GridBase");
 template<> const char *Namify<Grid<Vec3>>::S = "Grid<Vec3>";
-static const Pb::Register _R_91("Grid<Vec3>", "Grid", Grid<Vec3>::_W_9);
-static const Pb::Register _R_92("Grid<Vec3>", "save", Grid<Vec3>::_W_10);
-static const Pb::Register _R_93("Grid<Vec3>", "load", Grid<Vec3>::_W_11);
-static const Pb::Register _R_94("Grid<Vec3>", "clear", Grid<Vec3>::_W_12);
-static const Pb::Register _R_95("Grid<Vec3>", "copyFrom", Grid<Vec3>::_W_13);
-static const Pb::Register _R_96("Grid<Vec3>", "getGridType", Grid<Vec3>::_W_14);
-static const Pb::Register _R_97("Grid<Vec3>", "add", Grid<Vec3>::_W_15);
-static const Pb::Register _R_98("Grid<Vec3>", "sub", Grid<Vec3>::_W_16);
-static const Pb::Register _R_99("Grid<Vec3>", "setConst", Grid<Vec3>::_W_17);
-static const Pb::Register _R_100("Grid<Vec3>", "addConst", Grid<Vec3>::_W_18);
-static const Pb::Register _R_101("Grid<Vec3>", "addScaled", Grid<Vec3>::_W_19);
-static const Pb::Register _R_102("Grid<Vec3>", "mult", Grid<Vec3>::_W_20);
-static const Pb::Register _R_103("Grid<Vec3>", "multConst", Grid<Vec3>::_W_21);
-static const Pb::Register _R_104("Grid<Vec3>", "safeDivide", Grid<Vec3>::_W_22);
-static const Pb::Register _R_105("Grid<Vec3>", "clamp", Grid<Vec3>::_W_23);
-static const Pb::Register _R_106("Grid<Vec3>", "stomp", Grid<Vec3>::_W_24);
-static const Pb::Register _R_107("Grid<Vec3>", "permuteAxes", Grid<Vec3>::_W_25);
-static const Pb::Register _R_108("Grid<Vec3>", "permuteAxesCopyToGrid", Grid<Vec3>::_W_26);
-static const Pb::Register _R_109("Grid<Vec3>", "join", Grid<Vec3>::_W_27);
-static const Pb::Register _R_110("Grid<Vec3>", "getMaxAbs", Grid<Vec3>::_W_28);
-static const Pb::Register _R_111("Grid<Vec3>", "getMax", Grid<Vec3>::_W_29);
-static const Pb::Register _R_112("Grid<Vec3>", "getMin", Grid<Vec3>::_W_30);
-static const Pb::Register _R_113("Grid<Vec3>", "getL1", Grid<Vec3>::_W_31);
-static const Pb::Register _R_114("Grid<Vec3>", "getL2", Grid<Vec3>::_W_32);
-static const Pb::Register _R_115("Grid<Vec3>", "setBound", Grid<Vec3>::_W_33);
-static const Pb::Register _R_116("Grid<Vec3>", "setBoundNeumann", Grid<Vec3>::_W_34);
-static const Pb::Register _R_117("Grid<Vec3>", "getDataPointer", Grid<Vec3>::_W_35);
-static const Pb::Register _R_118("Grid<Vec3>", "printGrid", Grid<Vec3>::_W_36);
+static const Pb::Register _R_89("Grid<Vec3>", "Grid", Grid<Vec3>::_W_9);
+static const Pb::Register _R_90("Grid<Vec3>", "save", Grid<Vec3>::_W_10);
+static const Pb::Register _R_91("Grid<Vec3>", "load", Grid<Vec3>::_W_11);
+static const Pb::Register _R_92("Grid<Vec3>", "clear", Grid<Vec3>::_W_12);
+static const Pb::Register _R_93("Grid<Vec3>", "copyFrom", Grid<Vec3>::_W_13);
+static const Pb::Register _R_94("Grid<Vec3>", "getGridType", Grid<Vec3>::_W_14);
+static const Pb::Register _R_95("Grid<Vec3>", "add", Grid<Vec3>::_W_15);
+static const Pb::Register _R_96("Grid<Vec3>", "sub", Grid<Vec3>::_W_16);
+static const Pb::Register _R_97("Grid<Vec3>", "setConst", Grid<Vec3>::_W_17);
+static const Pb::Register _R_98("Grid<Vec3>", "addConst", Grid<Vec3>::_W_18);
+static const Pb::Register _R_99("Grid<Vec3>", "addScaled", Grid<Vec3>::_W_19);
+static const Pb::Register _R_100("Grid<Vec3>", "mult", Grid<Vec3>::_W_20);
+static const Pb::Register _R_101("Grid<Vec3>", "multConst", Grid<Vec3>::_W_21);
+static const Pb::Register _R_102("Grid<Vec3>", "safeDivide", Grid<Vec3>::_W_22);
+static const Pb::Register _R_103("Grid<Vec3>", "clamp", Grid<Vec3>::_W_23);
+static const Pb::Register _R_104("Grid<Vec3>", "stomp", Grid<Vec3>::_W_24);
+static const Pb::Register _R_105("Grid<Vec3>", "permuteAxes", Grid<Vec3>::_W_25);
+static const Pb::Register _R_106("Grid<Vec3>", "permuteAxesCopyToGrid", Grid<Vec3>::_W_26);
+static const Pb::Register _R_107("Grid<Vec3>", "getMaxAbs", Grid<Vec3>::_W_27);
+static const Pb::Register _R_108("Grid<Vec3>", "getMax", Grid<Vec3>::_W_28);
+static const Pb::Register _R_109("Grid<Vec3>", "getMin", Grid<Vec3>::_W_29);
+static const Pb::Register _R_110("Grid<Vec3>", "getL1", Grid<Vec3>::_W_30);
+static const Pb::Register _R_111("Grid<Vec3>", "getL2", Grid<Vec3>::_W_31);
+static const Pb::Register _R_112("Grid<Vec3>", "setBound", Grid<Vec3>::_W_32);
+static const Pb::Register _R_113("Grid<Vec3>", "setBoundNeumann", Grid<Vec3>::_W_33);
+static const Pb::Register _R_114("Grid<Vec3>", "getDataPointer", Grid<Vec3>::_W_34);
+static const Pb::Register _R_115("Grid<Vec3>", "printGrid", Grid<Vec3>::_W_35);
 #endif
 #ifdef _C_GridBase
-static const Pb::Register _R_119("GridBase", "GridBase", "PbClass");
+static const Pb::Register _R_116("GridBase", "GridBase", "PbClass");
 template<> const char *Namify<GridBase>::S = "GridBase";
-static const Pb::Register _R_120("GridBase", "GridBase", GridBase::_W_0);
-static const Pb::Register _R_121("GridBase", "getSizeX", GridBase::_W_1);
-static const Pb::Register _R_122("GridBase", "getSizeY", GridBase::_W_2);
-static const Pb::Register _R_123("GridBase", "getSizeZ", GridBase::_W_3);
-static const Pb::Register _R_124("GridBase", "getSize", GridBase::_W_4);
-static const Pb::Register _R_125("GridBase", "is3D", GridBase::_W_5);
-static const Pb::Register _R_126("GridBase", "is4D", GridBase::_W_6);
-static const Pb::Register _R_127("GridBase", "getSizeT", GridBase::_W_7);
-static const Pb::Register _R_128("GridBase", "getStrideT", GridBase::_W_8);
+static const Pb::Register _R_117("GridBase", "GridBase", GridBase::_W_0);
+static const Pb::Register _R_118("GridBase", "getSizeX", GridBase::_W_1);
+static const Pb::Register _R_119("GridBase", "getSizeY", GridBase::_W_2);
+static const Pb::Register _R_120("GridBase", "getSizeZ", GridBase::_W_3);
+static const Pb::Register _R_121("GridBase", "getSize", GridBase::_W_4);
+static const Pb::Register _R_122("GridBase", "is3D", GridBase::_W_5);
+static const Pb::Register _R_123("GridBase", "is4D", GridBase::_W_6);
+static const Pb::Register _R_124("GridBase", "getSizeT", GridBase::_W_7);
+static const Pb::Register _R_125("GridBase", "getStrideT", GridBase::_W_8);
 #endif
 #ifdef _C_MACGrid
-static const Pb::Register _R_129("MACGrid", "MACGrid", "Grid<Vec3>");
+static const Pb::Register _R_126("MACGrid", "MACGrid", "Grid<Vec3>");
 template<> const char *Namify<MACGrid>::S = "MACGrid";
-static const Pb::Register _R_130("MACGrid", "MACGrid", MACGrid::_W_37);
-static const Pb::Register _R_131("MACGrid", "setBoundMAC", MACGrid::_W_38);
+static const Pb::Register _R_127("MACGrid", "MACGrid", MACGrid::_W_36);
+static const Pb::Register _R_128("MACGrid", "setBoundMAC", MACGrid::_W_37);
 #endif
 static const Pb::Register _R_7("GridType_TypeNone", 0);
 static const Pb::Register _R_8("GridType_TypeReal", 1);
@@ -250,9 +247,6 @@ void PbRegister_file_7()
   KEEP_UNUSED(_R_126);
   KEEP_UNUSED(_R_127);
   KEEP_UNUSED(_R_128);
-  KEEP_UNUSED(_R_129);
-  KEEP_UNUSED(_R_130);
-  KEEP_UNUSED(_R_131);
 }
 }
 }  // namespace Manta

extern/mantaflow/preprocessed/plugin/secondaryparticles.cpp

@@ -525,7 +525,7 @@ struct knFlipSampleSecondaryParticlesMoreCylinders : public KernelBase {
     if (!(flags(i, j, k) & itype))
       return;
 
-    static RandomStream mRand(9832);
+    RandomStream mRand(9832);
     Real radius =
         0.25;  // diameter=0.5 => sampling with two cylinders in each dimension since cell size=1
     for (Real x = i - radius; x <= i + radius; x += 2 * radius) {
@@ -791,9 +791,11 @@ struct knFlipSampleSecondaryParticles : public KernelBase {
     const int n = KE * (k_ta * TA + k_wc * WC) * dt;  // number of secondary particles
     if (n == 0)
       return;
-    static RandomStream mRand(9832);
+    RandomStream mRand(9832);
 
-    Vec3 xi = Vec3(i, j, k) + mRand.getVec3();  // randomized offset uniform in cell
+    Vec3 xi = Vec3(i + mRand.getReal(),
+                   j + mRand.getReal(),
+                   k + mRand.getReal());  // randomized offset uniform in cell
     Vec3 vi = v.getInterpolated(xi);
     Vec3 dir = dt * vi;                               // direction of movement of current particle
     Vec3 e1 = getNormalized(Vec3(dir.z, 0, -dir.x));  // perpendicular to dir
@@ -1832,7 +1834,7 @@ struct knFlipDeleteParticlesInObstacle : public KernelBase {
     }
     int gridIndex = flags.index(xidx);
     // remove particles that penetrate obstacles
-    if (flags.isObstacle(gridIndex) || flags.isOutflow(gridIndex)) {
+    if (flags[gridIndex] == FlagGrid::TypeObstacle || flags[gridIndex] == FlagGrid::TypeOutflow) {
       pts.kill(idx);
     }
   }

intern/cycles/CMakeLists.txt

@@ -313,7 +313,9 @@ if(WITH_CYCLES_CUDA_BINARIES AND (NOT WITH_CYCLES_CUBIN_COMPILER))
       set(MAX_MSVC 1910)
     elseif(${CUDA_VERSION} EQUAL "9.1")
       set(MAX_MSVC 1911)
-    elseif(${CUDA_VERSION} LESS "11.0")
+    elseif(${CUDA_VERSION} EQUAL "10.0")
+      set(MAX_MSVC 1999)
+    elseif(${CUDA_VERSION} EQUAL "10.1")
       set(MAX_MSVC 1999)
     endif()
     if(NOT MSVC_VERSION LESS ${MAX_MSVC} OR CMAKE_C_COMPILER_ID MATCHES "Clang")
@@ -330,7 +332,7 @@ if(WITH_CYCLES_CUDA_BINARIES AND (NOT WITH_CYCLES_CUBIN_COMPILER))
 endif()
 
 # NVRTC gives wrong rendering result in CUDA 10.0, so we must use NVCC.
-if(WITH_CYCLES_CUDA_BINARIES AND WITH_CYCLES_CUBIN_COMPILER AND NOT WITH_CYCLES_CUBIN_COMPILER_OVERRRIDE)
+if(WITH_CYCLES_CUDA_BINARIES AND WITH_CYCLES_CUBIN_COMPILER)
   if(NOT (${CUDA_VERSION} VERSION_LESS 10.0))
     message(STATUS "cycles_cubin_cc not supported for CUDA 10.0+, using nvcc instead.")
     set(WITH_CYCLES_CUBIN_COMPILER OFF)

intern/cycles/app/cycles_cubin_cc.cpp

@@ -43,8 +43,7 @@ template<typename T> std::string to_string(const T &n)
 
 class CompilationSettings {
  public:
-  CompilationSettings()
-      : target_arch(0), bits(64), verbose(false), fast_math(false), ptx_only(false)
+  CompilationSettings() : target_arch(0), bits(64), verbose(false), fast_math(false)
   {
   }
 
@@ -58,13 +57,12 @@ class CompilationSettings {
   int bits;
   bool verbose;
   bool fast_math;
-  bool ptx_only;
 };
 
 static bool compile_cuda(CompilationSettings &settings)
 {
-  const char *headers[] = {"stdlib.h", "float.h", "math.h", "stdio.h", "stddef.h"};
-  const char *header_content[] = {"\n", "\n", "\n", "\n", "\n"};
+  const char *headers[] = {"stdlib.h", "float.h", "math.h", "stdio.h"};
+  const char *header_content[] = {"\n", "\n", "\n", "\n"};
 
   printf("Building %s\n", settings.input_file.c_str());
 
@@ -85,8 +83,6 @@ static bool compile_cuda(CompilationSettings &settings)
   options.push_back("-D__KERNEL_CUDA_VERSION__=" + std::to_string(cuewNvrtcVersion()));
   options.push_back("-arch=compute_" + std::to_string(settings.target_arch));
   options.push_back("--device-as-default-execution-space");
-  options.push_back("-DCYCLES_CUBIN_CC");
-  options.push_back("--std=c++11");
   if (settings.fast_math)
     options.push_back("--use_fast_math");
 
@@ -138,14 +134,10 @@ static bool compile_cuda(CompilationSettings &settings)
     fprintf(stderr, "Error: nvrtcGetPTX failed (%d)\n\n", (int)result);
     return false;
   }
-  if (settings.ptx_only) {
-    settings.ptx_file = settings.output_file;
-  }
-  else {
-    /* Write a file in the temp folder with the ptx code. */
-    settings.ptx_file = OIIO::Filesystem::temp_directory_path() + "/" +
-                        OIIO::Filesystem::unique_path();
-  }
+
+  /* Write a file in the temp folder with the ptx code. */
+  settings.ptx_file = OIIO::Filesystem::temp_directory_path() + "/" +
+                      OIIO::Filesystem::unique_path();
   FILE *f = fopen(settings.ptx_file.c_str(), "wb");
   fwrite(&ptx_code[0], 1, ptx_size, f);
   fclose(f);
@@ -257,9 +249,6 @@ static bool parse_parameters(int argc, const char **argv, CompilationSettings &s
              "-D %L",
              &settings.defines,
              "Add additional defines",
-             "-ptx",
-             &settings.ptx_only,
-             "emit PTX code",
              "-v",
              &settings.verbose,
              "Use verbose logging",
@@ -314,10 +303,8 @@ int main(int argc, const char **argv)
     exit(EXIT_FAILURE);
   }
 
-  if (!settings.ptx_only) {
-    if (!link_ptxas(settings)) {
-      exit(EXIT_FAILURE);
-    }
+  if (!link_ptxas(settings)) {
+    exit(EXIT_FAILURE);
   }
 
   return 0;

intern/cycles/blender/addon/engine.py

@@ -260,16 +260,15 @@ def list_render_passes(srl):
     if crl.use_pass_volume_indirect:           yield ("VolumeInd",                     "RGB", 'COLOR')
 
     # Cryptomatte passes.
-    crypto_depth = (crl.pass_crypto_depth + 1) // 2
     if crl.use_pass_crypto_object:
-        for i in range(0, crypto_depth):
-            yield ("CryptoObject" + '{:02d}'.format(i), "RGBA", 'COLOR')
+        for i in range(0, crl.pass_crypto_depth, 2):
+            yield ("CryptoObject" + '{:02d}'.format(i//2), "RGBA", 'COLOR')
     if crl.use_pass_crypto_material:
-        for i in range(0, crypto_depth):
-            yield ("CryptoMaterial" + '{:02d}'.format(i), "RGBA", 'COLOR')
+        for i in range(0, crl.pass_crypto_depth, 2):
+            yield ("CryptoMaterial" + '{:02d}'.format(i//2), "RGBA", 'COLOR')
     if srl.cycles.use_pass_crypto_asset:
-        for i in range(0, crypto_depth):
-            yield ("CryptoAsset" + '{:02d}'.format(i), "RGBA", 'COLOR')
+        for i in range(0, srl.cycles.pass_crypto_depth, 2):
+            yield ("CryptoAsset" + '{:02d}'.format(i//2), "RGBA", 'COLOR')
 
     # Denoising passes.
     if crl.use_denoising or crl.denoising_store_passes:

intern/cycles/blender/addon/ui.py

@@ -1917,6 +1917,7 @@ class CYCLES_RENDER_PT_bake_influence(CyclesButtonsPanel, Panel):
             flow.prop(cbk, "use_pass_diffuse")
             flow.prop(cbk, "use_pass_glossy")
             flow.prop(cbk, "use_pass_transmission")
+            flow.prop(cbk, "use_pass_subsurface")
             flow.prop(cbk, "use_pass_ambient_occlusion")
             flow.prop(cbk, "use_pass_emit")
 
@@ -2086,7 +2087,7 @@ class CYCLES_RENDER_PT_simplify_viewport(CyclesButtonsPanel, Panel):
         col.prop(rd, "simplify_child_particles", text="Child Particles")
         col.prop(cscene, "texture_limit", text="Texture Limit")
         col.prop(cscene, "ao_bounces", text="AO Bounces")
-
+        col.prop(rd, "use_simplify_smoke_highres")
 
 class CYCLES_RENDER_PT_simplify_render(CyclesButtonsPanel, Panel):
     bl_label = "Render"

intern/cycles/blender/blender_sync.cpp

@@ -633,11 +633,11 @@ vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer &b_rlay,
 
   /* Cryptomatte stores two ID/weight pairs per RGBA layer.
    * User facing parameter is the number of pairs. */
-  int crypto_depth = divide_up(min(16, get_int(crp, "pass_crypto_depth")), 2);
+  int crypto_depth = min(16, get_int(crp, "pass_crypto_depth")) / 2;
   scene->film->cryptomatte_depth = crypto_depth;
   scene->film->cryptomatte_passes = CRYPT_NONE;
   if (get_boolean(crp, "use_pass_crypto_object")) {
-    for (int i = 0; i < crypto_depth; i++) {
+    for (int i = 0; i < crypto_depth; ++i) {
       string passname = cryptomatte_prefix + string_printf("Object%02d", i);
       b_engine.add_pass(passname.c_str(), 4, "RGBA", b_view_layer.name().c_str());
       Pass::add(PASS_CRYPTOMATTE, passes, passname.c_str());
@@ -646,7 +646,7 @@ vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer &b_rlay,
                                                         CRYPT_OBJECT);
   }
   if (get_boolean(crp, "use_pass_crypto_material")) {
-    for (int i = 0; i < crypto_depth; i++) {
+    for (int i = 0; i < crypto_depth; ++i) {
       string passname = cryptomatte_prefix + string_printf("Material%02d", i);
       b_engine.add_pass(passname.c_str(), 4, "RGBA", b_view_layer.name().c_str());
       Pass::add(PASS_CRYPTOMATTE, passes, passname.c_str());
@@ -655,7 +655,7 @@ vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer &b_rlay,
                                                         CRYPT_MATERIAL);
   }
   if (get_boolean(crp, "use_pass_crypto_asset")) {
-    for (int i = 0; i < crypto_depth; i++) {
+    for (int i = 0; i < crypto_depth; ++i) {
       string passname = cryptomatte_prefix + string_printf("Asset%02d", i);
       b_engine.add_pass(passname.c_str(), 4, "RGBA", b_view_layer.name().c_str());
       Pass::add(PASS_CRYPTOMATTE, passes, passname.c_str());

intern/cycles/device/device.cpp

@@ -597,7 +597,6 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
 
   info.has_half_images = true;
   info.has_volume_decoupled = true;
-  info.has_adaptive_stop_per_sample = true;
   info.has_osl = true;
   info.has_profiling = true;
 
@@ -640,7 +639,6 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
     /* Accumulate device info. */
     info.has_half_images &= device.has_half_images;
     info.has_volume_decoupled &= device.has_volume_decoupled;
-    info.has_adaptive_stop_per_sample &= device.has_adaptive_stop_per_sample;
     info.has_osl &= device.has_osl;
     info.has_profiling &= device.has_profiling;
   }

intern/cycles/device/device.h

@@ -75,13 +75,12 @@ class DeviceInfo {
   string description;
   string id; /* used for user preferences, should stay fixed with changing hardware config */
   int num;
-  bool display_device;               /* GPU is used as a display device. */
-  bool has_half_images;              /* Support half-float textures. */
-  bool has_volume_decoupled;         /* Decoupled volume shading. */
-  bool has_adaptive_stop_per_sample; /* Per-sample adaptive sampling stopping. */
-  bool has_osl;                      /* Support Open Shading Language. */
-  bool use_split_kernel;             /* Use split or mega kernel. */
-  bool has_profiling;                /* Supports runtime collection of profiling info. */
+  bool display_device;       /* GPU is used as a display device. */
+  bool has_half_images;      /* Support half-float textures. */
+  bool has_volume_decoupled; /* Decoupled volume shading. */
+  bool has_osl;              /* Support Open Shading Language. */
+  bool use_split_kernel;     /* Use split or mega kernel. */
+  bool has_profiling;        /* Supports runtime collection of profiling info. */
   int cpu_threads;
   vector<DeviceInfo> multi_devices;
   vector<DeviceInfo> denoising_devices;
@@ -95,7 +94,6 @@ class DeviceInfo {
     display_device = false;
     has_half_images = false;
     has_volume_decoupled = false;
-    has_adaptive_stop_per_sample = false;
     has_osl = false;
     use_split_kernel = false;
     has_profiling = false;

intern/cycles/device/device_cpu.cpp

@@ -839,7 +839,7 @@ class CPUDevice : public Device {
     return true;
   }
 
-  bool adaptive_sampling_filter(KernelGlobals *kg, RenderTile &tile, int sample)
+  bool adaptive_sampling_filter(KernelGlobals *kg, RenderTile &tile)
   {
     WorkTile wtile;
     wtile.x = tile.x;
@@ -850,24 +850,11 @@ class CPUDevice : public Device {
     wtile.stride = tile.stride;
     wtile.buffer = (float *)tile.buffer;
 
-    /* For CPU we do adaptive stopping per sample so we can stop earlier, but
-     * for combined CPU + GPU rendering we match the GPU and do it per tile
-     * after a given number of sample steps. */
-    if (!kernel_data.integrator.adaptive_stop_per_sample) {
-      for (int y = wtile.y; y < wtile.y + wtile.h; ++y) {
-        for (int x = wtile.x; x < wtile.x + wtile.w; ++x) {
-          const int index = wtile.offset + x + y * wtile.stride;
-          float *buffer = wtile.buffer + index * kernel_data.film.pass_stride;
-          kernel_do_adaptive_stopping(kg, buffer, sample);
-        }
-      }
-    }
-
     bool any = false;
-    for (int y = wtile.y; y < wtile.y + wtile.h; ++y) {
+    for (int y = tile.y; y < tile.y + tile.h; ++y) {
       any |= kernel_do_adaptive_filter_x(kg, y, &wtile);
     }
-    for (int x = wtile.x; x < wtile.x + wtile.w; ++x) {
+    for (int x = tile.x; x < tile.x + tile.w; ++x) {
       any |= kernel_do_adaptive_filter_y(kg, x, &wtile);
     }
     return (!any);
@@ -930,7 +917,7 @@ class CPUDevice : public Device {
       tile.sample = sample + 1;
 
       if (task.adaptive_sampling.use && task.adaptive_sampling.need_filter(sample)) {
-        const bool stop = adaptive_sampling_filter(kg, tile, sample);
+        const bool stop = adaptive_sampling_filter(kg, tile);
         if (stop) {
           const int num_progress_samples = end_sample - sample;
           tile.sample = end_sample;
@@ -1340,7 +1327,6 @@ void device_cpu_info(vector<DeviceInfo> &devices)
   info.id = "CPU";
   info.num = 0;
   info.has_volume_decoupled = true;
-  info.has_adaptive_stop_per_sample = true;
   info.has_osl = true;
   info.has_half_images = true;
   info.has_profiling = true;

intern/cycles/device/device_cuda.cpp

@@ -129,7 +129,6 @@ void device_cuda_info(vector<DeviceInfo> &devices)
 
     info.has_half_images = (major >= 3);
     info.has_volume_decoupled = false;
-    info.has_adaptive_stop_per_sample = false;
 
     int pci_location[3] = {0, 0, 0};
     cuDeviceGetAttribute(&pci_location[0], CU_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID, num);

intern/cycles/device/device_network.cpp

@@ -311,7 +311,6 @@ void device_network_info(vector<DeviceInfo> &devices)
 
   /* todo: get this info from device */
   info.has_volume_decoupled = false;
-  info.has_adaptive_stop_per_sample = false;
   info.has_osl = false;
 
   devices.push_back(info);

intern/cycles/device/device_opencl.cpp

@@ -119,7 +119,6 @@ void device_opencl_info(vector<DeviceInfo> &devices)
     info.display_device = true;
     info.use_split_kernel = true;
     info.has_volume_decoupled = false;
-    info.has_adaptive_stop_per_sample = false;
     info.id = id;
 
     /* Check OpenCL extensions */

intern/cycles/device/device_optix.cpp

@@ -1558,7 +1558,7 @@ void device_optix_info(vector<DeviceInfo> &devices)
     }
 
     // Only add devices with RTX support
-    if (rtcore_version == 0 && !getenv("CYCLES_OPTIX_TEST"))
+    if (rtcore_version == 0)
       it = cuda_devices.erase(it);
     else
       ++it;

intern/cycles/device/device_task.cpp

@@ -138,7 +138,8 @@ void DeviceTask::update_progress(RenderTile *rtile, int pixel_samples)
 
 /* Adaptive Sampling */
 
-AdaptiveSampling::AdaptiveSampling() : use(true), adaptive_step(0), min_samples(0)
+AdaptiveSampling::AdaptiveSampling()
+    : use(true), adaptive_step(ADAPTIVE_SAMPLE_STEP), min_samples(0)
 {
 }
 

intern/cycles/device/opencl/device_opencl.h

@@ -88,12 +88,9 @@ class OpenCLInfo {
   static bool device_supported(const string &platform_name, const cl_device_id device_id);
   static bool platform_version_check(cl_platform_id platform, string *error = NULL);
   static bool device_version_check(cl_device_id device, string *error = NULL);
-  static bool get_device_version(cl_device_id device,
-                                 int *r_major,
-                                 int *r_minor,
-                                 string *error = NULL);
   static string get_hardware_id(const string &platform_name, cl_device_id device_id);
-  static void get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices);
+  static void get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices,
+                                 bool force_all = false);
 
   /* ** Some handy shortcuts to low level cl*GetInfo() functions. ** */
 

intern/cycles/device/opencl/device_opencl_impl.cpp

@@ -257,16 +257,16 @@ void OpenCLDevice::OpenCLSplitPrograms::load_kernels(
 
     /* Ordered with most complex kernels first, to reduce overall compile time. */
     ADD_SPLIT_KERNEL_PROGRAM(subsurface_scatter);
-    ADD_SPLIT_KERNEL_PROGRAM(direct_lighting);
-    ADD_SPLIT_KERNEL_PROGRAM(indirect_background);
     if (requested_features.use_volume || is_preview) {
       ADD_SPLIT_KERNEL_PROGRAM(do_volume);
     }
-    ADD_SPLIT_KERNEL_PROGRAM(shader_eval);
-    ADD_SPLIT_KERNEL_PROGRAM(lamp_emission);
-    ADD_SPLIT_KERNEL_PROGRAM(holdout_emission_blurring_pathtermination_ao);
     ADD_SPLIT_KERNEL_PROGRAM(shadow_blocked_dl);
     ADD_SPLIT_KERNEL_PROGRAM(shadow_blocked_ao);
+    ADD_SPLIT_KERNEL_PROGRAM(holdout_emission_blurring_pathtermination_ao);
+    ADD_SPLIT_KERNEL_PROGRAM(lamp_emission);
+    ADD_SPLIT_KERNEL_PROGRAM(direct_lighting);
+    ADD_SPLIT_KERNEL_PROGRAM(indirect_background);
+    ADD_SPLIT_KERNEL_PROGRAM(shader_eval);
 
     /* Quick kernels bundled in a single program to reduce overhead of starting
      * Blender processes. */
@@ -1896,17 +1896,6 @@ string OpenCLDevice::kernel_build_options(const string *debug_src)
 {
   string build_options = "-cl-no-signed-zeros -cl-mad-enable ";
 
-  /* Build with OpenCL 2.0 if available, this improves performance
-   * with AMD OpenCL drivers on Windows and Linux (legacy drivers).
-   * Note that OpenCL selects the highest 1.x version by default,
-   * only for 2.0 do we need the explicit compiler flag. */
-  int version_major, version_minor;
-  if (OpenCLInfo::get_device_version(cdDevice, &version_major, &version_minor)) {
-    if (version_major >= 2) {
-      build_options += "-cl-std=CL2.0 ";
-    }
-  }
-
   if (platform_name == "NVIDIA CUDA") {
     build_options +=
         "-D__KERNEL_OPENCL_NVIDIA__ "

intern/cycles/device/opencl/opencl_util.cpp

@@ -23,7 +23,6 @@
 #  include "util/util_logging.h"
 #  include "util/util_md5.h"
 #  include "util/util_path.h"
-#  include "util/util_semaphore.h"
 #  include "util/util_system.h"
 #  include "util/util_time.h"
 
@@ -391,27 +390,8 @@ static void escape_python_string(string &str)
   string_replace(str, "'", "\'");
 }
 
-static int opencl_compile_process_limit()
-{
-  /* Limit number of concurrent processes compiling, with a heuristic based
-   * on total physical RAM and estimate of memory usage needed when compiling
-   * with all Cycles features enabled.
-   *
-   * This is somewhat arbitrary as we don't know the actual available RAM or
-   * how much the kernel compilation will needed depending on the features, but
-   * better than not limiting at all. */
-  static const int64_t GB = 1024LL * 1024LL * 1024LL;
-  static const int64_t process_memory = 2 * GB;
-  static const int64_t base_memory = 2 * GB;
-  static const int64_t system_memory = system_physical_ram();
-  static const int64_t process_limit = (system_memory - base_memory) / process_memory;
-
-  return max((int)process_limit, 1);
-}
-
 bool OpenCLDevice::OpenCLProgram::compile_separate(const string &clbin)
 {
-  /* Construct arguments. */
   vector<string> args;
   args.push_back("--background");
   args.push_back("--factory-startup");
@@ -439,23 +419,14 @@ bool OpenCLDevice::OpenCLProgram::compile_separate(const string &clbin)
       kernel_file_escaped.c_str(),
       clbin_escaped.c_str()));
 
-  /* Limit number of concurrent processes compiling. */
-  static thread_counting_semaphore semaphore(opencl_compile_process_limit());
-  semaphore.acquire();
-
-  /* Compile. */
-  const double starttime = time_dt();
+  double starttime = time_dt();
   add_log(string("Cycles: compiling OpenCL program ") + program_name + "...", false);
   add_log(string("Build flags: ") + kernel_build_options, true);
-  const bool success = system_call_self(args);
-  const double elapsed = time_dt() - starttime;
-
-  semaphore.release();
-
-  if (!success || !path_exists(clbin)) {
+  if (!system_call_self(args) || !path_exists(clbin)) {
     return false;
   }
 
+  double elapsed = time_dt() - starttime;
   add_log(
       string_printf("Kernel compilation of %s finished in %.2lfs.", program_name.c_str(), elapsed),
       false);
@@ -776,10 +747,6 @@ bool OpenCLInfo::device_supported(const string &platform_name, const cl_device_i
   }
   VLOG(3) << "OpenCL driver version " << driver_major << "." << driver_minor;
 
-  if (getenv("CYCLES_OPENCL_TEST")) {
-    return true;
-  }
-
   /* It is possible to have Iris GPU on AMD/Apple OpenCL framework
    * (aka, it will not be on Intel framework). This isn't supported
    * and needs an explicit blacklist.
@@ -839,30 +806,18 @@ bool OpenCLInfo::platform_version_check(cl_platform_id platform, string *error)
   return true;
 }
 
-bool OpenCLInfo::get_device_version(cl_device_id device, int *r_major, int *r_minor, string *error)
+bool OpenCLInfo::device_version_check(cl_device_id device, string *error)
 {
+  const int req_major = 1, req_minor = 1;
+  int major, minor;
   char version[256];
   clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION, sizeof(version), &version, NULL);
-  if (sscanf(version, "OpenCL C %d.%d", r_major, r_minor) < 2) {
+  if (sscanf(version, "OpenCL C %d.%d", &major, &minor) < 2) {
     if (error != NULL) {
       *error = string_printf("OpenCL: failed to parse OpenCL C version string (%s).", version);
     }
     return false;
   }
-  if (error != NULL) {
-    *error = "";
-  }
-  return true;
-}
-
-bool OpenCLInfo::device_version_check(cl_device_id device, string *error)
-{
-  const int req_major = 1, req_minor = 1;
-  int major, minor;
-  if (!get_device_version(device, &major, &minor, error)) {
-    return false;
-  }
-
   if (!((major == req_major && minor >= req_minor) || (major > req_major))) {
     if (error != NULL) {
       *error = string_printf("OpenCL: C version 1.1 or later required, found %d.%d", major, minor);
@@ -903,7 +858,7 @@ string OpenCLInfo::get_hardware_id(const string &platform_name, cl_device_id dev
   return "";
 }
 
-void OpenCLInfo::get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices)
+void OpenCLInfo::get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices, bool force_all)
 {
   const cl_device_type device_type = OpenCLInfo::device_type();
   static bool first_time = true;
@@ -969,7 +924,7 @@ void OpenCLInfo::get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices
         FIRST_VLOG(2) << "Ignoring device " << device_name << " due to old compiler version.";
         continue;
       }
-      if (device_supported(platform_name, device_id)) {
+      if (force_all || device_supported(platform_name, device_id)) {
         cl_device_type device_type;
         if (!get_device_type(device_id, &device_type, &error)) {
           FIRST_VLOG(2) << "Ignoring device " << device_name

intern/cycles/kernel/CMakeLists.txt

@@ -452,7 +452,7 @@ if(WITH_CYCLES_CUDA_BINARIES)
       endif()
 
       add_custom_command(
-        OUTPUT ${cuda_file}
+        OUTPUT ${cuda_cubin}
         COMMAND ${CUBIN_CC_ENV}
             "$<TARGET_FILE:cycles_cubin_cc>"
             -target ${CUDA_ARCH}
@@ -461,6 +461,7 @@ if(WITH_CYCLES_CUDA_BINARIES)
             -v
             -cuda-toolkit-dir "${CUDA_TOOLKIT_ROOT_DIR}"
         DEPENDS ${kernel_sources} cycles_cubin_cc)
+      set(cuda_file ${cuda_cubin})
     else()
       add_custom_command(
         OUTPUT ${cuda_file}
@@ -516,6 +517,7 @@ if(WITH_CYCLES_DEVICE_OPTIX)
       -I "${OPTIX_INCLUDE_DIR}"
       -I "${CMAKE_CURRENT_SOURCE_DIR}/.."
       -I "${CMAKE_CURRENT_SOURCE_DIR}/kernels/cuda"
+      -arch=sm_30
       --use_fast_math
       -o ${output})
 
@@ -523,62 +525,25 @@ if(WITH_CYCLES_DEVICE_OPTIX)
       set(cuda_flags ${cuda_flags}
         -D __KERNEL_DEBUG__)
     endif()
-    if(WITH_CYCLES_CUBIN_COMPILER)
 
-      # Needed to find libnvrtc-builtins.so. Can't do it from inside
-      # cycles_cubin_cc since the env variable is read before main()
-      if(APPLE)
-        set(CUBIN_CC_ENV ${CMAKE_COMMAND}
-          -E env DYLD_LIBRARY_PATH="${CUDA_TOOLKIT_ROOT_DIR}/lib")
-      elseif(UNIX)
-        set(CUBIN_CC_ENV ${CMAKE_COMMAND}
-          -E env LD_LIBRARY_PATH="${CUDA_TOOLKIT_ROOT_DIR}/lib64")
-      endif()
+    add_custom_command(
+      OUTPUT
+        ${output}
+      DEPENDS
+        ${input}
+        ${SRC_HEADERS}
+        ${SRC_KERNELS_CUDA_HEADERS}
+        ${SRC_KERNELS_OPTIX_HEADERS}
+        ${SRC_BVH_HEADERS}
+        ${SRC_SVM_HEADERS}
+        ${SRC_GEOM_HEADERS}
+        ${SRC_CLOSURE_HEADERS}
+        ${SRC_UTIL_HEADERS}
+      COMMAND
+        ${CUDA_NVCC_EXECUTABLE} --ptx ${cuda_flags} ${input}
+      WORKING_DIRECTORY
+        "${CMAKE_CURRENT_SOURCE_DIR}")
 
-      add_custom_command(
-        OUTPUT ${output}
-        DEPENDS
-          ${input}
-          ${SRC_HEADERS}
-          ${SRC_KERNELS_CUDA_HEADERS}
-          ${SRC_KERNELS_OPTIX_HEADERS}
-          ${SRC_BVH_HEADERS}
-          ${SRC_SVM_HEADERS}
-          ${SRC_GEOM_HEADERS}
-          ${SRC_CLOSURE_HEADERS}
-          ${SRC_UTIL_HEADERS}
-        COMMAND ${CUBIN_CC_ENV}
-            "$<TARGET_FILE:cycles_cubin_cc>"
-            -target 30
-            -ptx
-            -i ${CMAKE_CURRENT_SOURCE_DIR}/${input}
-            ${cuda_flags}
-            -v
-            -cuda-toolkit-dir "${CUDA_TOOLKIT_ROOT_DIR}"
-        DEPENDS ${kernel_sources} cycles_cubin_cc)
-    else()
-      add_custom_command(
-        OUTPUT
-          ${output}
-        DEPENDS
-          ${input}
-          ${SRC_HEADERS}
-          ${SRC_KERNELS_CUDA_HEADERS}
-          ${SRC_KERNELS_OPTIX_HEADERS}
-          ${SRC_BVH_HEADERS}
-          ${SRC_SVM_HEADERS}
-          ${SRC_GEOM_HEADERS}
-          ${SRC_CLOSURE_HEADERS}
-          ${SRC_UTIL_HEADERS}
-        COMMAND
-          ${CUDA_NVCC_EXECUTABLE}
-          --ptx
-          -arch=sm_30
-          ${cuda_flags}
-          ${input}
-        WORKING_DIRECTORY
-          "${CMAKE_CURRENT_SOURCE_DIR}")
-    endif()
     list(APPEND optix_ptx ${output})
 
     delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${output}" ${CYCLES_INSTALL_PATH}/lib)

intern/cycles/kernel/closure/bsdf_hair_principled.h

@@ -493,36 +493,6 @@ ccl_device void bsdf_principled_hair_blur(ShaderClosure *sc, float roughness)
   bsdf->m0_roughness = fmaxf(roughness, bsdf->m0_roughness);
 }
 
-/* Hair Albedo */
-
-ccl_device_inline float bsdf_principled_hair_albedo_roughness_scale(
-    const float azimuthal_roughness)
-{
-  const float x = azimuthal_roughness;
-  return (((((0.245f * x) + 5.574f) * x - 10.73f) * x + 2.532f) * x - 0.215f) * x + 5.969f;
-}
-
-ccl_device float3 bsdf_principled_hair_albedo(ShaderClosure *sc)
-{
-  PrincipledHairBSDF *bsdf = (PrincipledHairBSDF *)sc;
-  return exp3(-sqrt(bsdf->sigma) * bsdf_principled_hair_albedo_roughness_scale(bsdf->v));
-}
-
-ccl_device_inline float3
-bsdf_principled_hair_sigma_from_reflectance(const float3 color, const float azimuthal_roughness)
-{
-  const float3 sigma = log3(color) /
-                       bsdf_principled_hair_albedo_roughness_scale(azimuthal_roughness);
-  return sigma * sigma;
-}
-
-ccl_device_inline float3 bsdf_principled_hair_sigma_from_concentration(const float eumelanin,
-                                                                       const float pheomelanin)
-{
-  return eumelanin * make_float3(0.506f, 0.841f, 1.653f) +
-         pheomelanin * make_float3(0.343f, 0.733f, 1.924f);
-}
-
 CCL_NAMESPACE_END
 
 #endif /* __BSDF_HAIR_PRINCIPLED_H__ */

intern/cycles/kernel/geom/geom_motion_triangle_intersect.h

@@ -103,21 +103,17 @@ ccl_device_inline
                                  const Ray *ray,
                                  float3 verts[3])
 {
-#  ifdef __KERNEL_OPTIX__
-  /* isect->t is always in world space with OptiX. */
-  return motion_triangle_refine(kg, sd, isect, ray, verts);
-#  else
   float3 P = ray->P;
   float3 D = ray->D;
   float t = isect->t;
 
-#    ifdef __INTERSECTION_REFINE__
+#  ifdef __INTERSECTION_REFINE__
   if (isect->object != OBJECT_NONE) {
-#      ifdef __OBJECT_MOTION__
+#    ifdef __OBJECT_MOTION__
     Transform tfm = sd->ob_itfm;
-#      else
+#    else
     Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM);
-#      endif
+#    endif
 
     P = transform_point(&tfm, P);
     D = transform_direction(&tfm, D);
@@ -139,20 +135,19 @@ ccl_device_inline
   P = P + D * rt;
 
   if (isect->object != OBJECT_NONE) {
-#      ifdef __OBJECT_MOTION__
+#    ifdef __OBJECT_MOTION__
     Transform tfm = sd->ob_tfm;
-#      else
+#    else
     Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM);
-#      endif
+#    endif
 
     P = transform_point(&tfm, P);
   }
 
   return P;
-#    else  /* __INTERSECTION_REFINE__ */
+#  else  /* __INTERSECTION_REFINE__ */
   return P + D * t;
-#    endif /* __INTERSECTION_REFINE__ */
-#  endif
+#  endif /* __INTERSECTION_REFINE__ */
 }
 #endif /* __BVH_LOCAL__ */
 

intern/cycles/kernel/geom/geom_triangle_intersect.h

@@ -690,20 +690,16 @@ ccl_device_inline float3 triangle_refine_local(KernelGlobals *kg,
                                                const Intersection *isect,
                                                const Ray *ray)
 {
-#ifdef __KERNEL_OPTIX__
-  /* isect->t is always in world space with OptiX. */
-  return triangle_refine(kg, sd, isect, ray);
-#else
   float3 P = ray->P;
   float3 D = ray->D;
   float t = isect->t;
 
   if (isect->object != OBJECT_NONE) {
-#  ifdef __OBJECT_MOTION__
+#ifdef __OBJECT_MOTION__
     Transform tfm = sd->ob_itfm;
-#  else
+#else
     Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM);
-#  endif
+#endif
 
     P = transform_point(&tfm, P);
     D = transform_direction(&tfm, D);
@@ -712,7 +708,7 @@ ccl_device_inline float3 triangle_refine_local(KernelGlobals *kg,
 
   P = P + D * t;
 
-#  ifdef __INTERSECTION_REFINE__
+#ifdef __INTERSECTION_REFINE__
   const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim);
   const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0),
                tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1),
@@ -732,20 +728,19 @@ ccl_device_inline float3 triangle_refine_local(KernelGlobals *kg,
     float rt = dot(edge2, qvec) / det;
     P = P + D * rt;
   }
-#  endif /* __INTERSECTION_REFINE__ */
+#endif /* __INTERSECTION_REFINE__ */
 
   if (isect->object != OBJECT_NONE) {
-#  ifdef __OBJECT_MOTION__
+#ifdef __OBJECT_MOTION__
     Transform tfm = sd->ob_tfm;
-#  else
+#else
     Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM);
-#  endif
+#endif
 
     P = transform_point(&tfm, P);
   }
 
   return P;
-#endif
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/kernel_adaptive_sampling.h

@@ -150,7 +150,6 @@ ccl_device void kernel_adaptive_post_adjust(KernelGlobals *kg,
   }
 #endif /* __DENOISING_FEATURES__ */
 
-  /* Cryptomatte. */
   if (kernel_data.film.cryptomatte_passes) {
     int num_slots = 0;
     num_slots += (kernel_data.film.cryptomatte_passes & CRYPT_OBJECT) ? 1 : 0;
@@ -163,14 +162,6 @@ ccl_device void kernel_adaptive_post_adjust(KernelGlobals *kg,
       id_buffer[slot].y *= sample_multiplier;
     }
   }
-
-  /* AOVs. */
-  for (int i = 0; i < kernel_data.film.pass_aov_value_num; i++) {
-    *(buffer + kernel_data.film.pass_aov_value + i) *= sample_multiplier;
-  }
-  for (int i = 0; i < kernel_data.film.pass_aov_color_num; i++) {
-    *((ccl_global float4 *)(buffer + kernel_data.film.pass_aov_color) + i) *= sample_multiplier;
-  }
 }
 
 /* This is a simple box filter in two passes.

intern/cycles/kernel/kernel_compat_cuda.h

@@ -37,11 +37,8 @@ typedef unsigned long long uint64_t;
 typedef unsigned short half;
 typedef unsigned long long CUtexObject;
 
-#ifdef CYCLES_CUBIN_CC
-#  define FLT_MIN 1.175494350822287507969e-38f
-#  define FLT_MAX 340282346638528859811704183484516925440.0f
-#  define FLT_EPSILON 1.192092896e-07F
-#endif
+#define FLT_MIN 1.175494350822287507969e-38f
+#define FLT_MAX 340282346638528859811704183484516925440.0f
 
 __device__ half __float2half(const float f)
 {

intern/cycles/kernel/kernel_compat_optix.h

@@ -35,11 +35,9 @@ typedef unsigned int uint32_t;
 typedef unsigned long long uint64_t;
 typedef unsigned short half;
 typedef unsigned long long CUtexObject;
-#ifdef CYCLES_CUBIN_CC
-#  define FLT_MIN 1.175494350822287507969e-38f
-#  define FLT_MAX 340282346638528859811704183484516925440.0f
-#  define FLT_EPSILON 1.192092896e-07F
-#endif
+
+#define FLT_MIN 1.175494350822287507969e-38f
+#define FLT_MAX 340282346638528859811704183484516925440.0f
 
 __device__ half __float2half(const float f)
 {

intern/cycles/kernel/kernel_passes.h

@@ -87,9 +87,6 @@ ccl_device_inline void kernel_update_denoising_features(KernelGlobals *kg,
       PrincipledSheenBsdf *bsdf = (PrincipledSheenBsdf *)sc;
       closure_albedo *= bsdf->avg_value;
     }
-    else if (sc->type == CLOSURE_BSDF_HAIR_PRINCIPLED_ID) {
-      closure_albedo *= bsdf_principled_hair_albedo(sc);
-    }
 
     if (bsdf_get_specular_roughness_squared(sc) > sqr(0.075f)) {
       diffuse_albedo += closure_albedo;
@@ -403,13 +400,9 @@ ccl_device_inline void kernel_write_result(KernelGlobals *kg,
                                make_float4(L_sum.x * 2.0f, L_sum.y * 2.0f, L_sum.z * 2.0f, 0.0f));
     }
 #ifdef __KERNEL_CPU__
-    if ((sample > kernel_data.integrator.adaptive_min_samples) &&
-        kernel_data.integrator.adaptive_stop_per_sample) {
-      const int step = kernel_data.integrator.adaptive_step;
-
-      if ((sample & (step - 1)) == (step - 1)) {
-        kernel_do_adaptive_stopping(kg, buffer, sample);
-      }
+    if (sample > kernel_data.integrator.adaptive_min_samples &&
+        (sample & (ADAPTIVE_SAMPLE_STEP - 1)) == (ADAPTIVE_SAMPLE_STEP - 1)) {
+      kernel_do_adaptive_stopping(kg, buffer, sample);
     }
 #endif
   }

intern/cycles/kernel/kernel_subsurface.h

@@ -428,17 +428,12 @@ ccl_device_noinline
     hit = (ss_isect->num_hits > 0);
 
     if (hit) {
-#ifdef __KERNEL_OPTIX__
-      /* t is always in world space with OptiX. */
-      t = ss_isect->hits[0].t;
-#else
       /* Compute world space distance to surface hit. */
       float3 D = ray->D;
       object_inverse_dir_transform(kg, sd, &D);
       D = normalize(D) * ss_isect->hits[0].t;
       object_dir_transform(kg, sd, &D);
       t = len(D);
-#endif
     }
 
     /* Advance to new scatter location. */

intern/cycles/kernel/kernel_types.h

@@ -63,6 +63,11 @@ CCL_NAMESPACE_BEGIN
 
 #define VOLUME_STACK_SIZE 32
 
+/* Adaptive sampling constants */
+#define ADAPTIVE_SAMPLE_STEP 4
+static_assert((ADAPTIVE_SAMPLE_STEP & (ADAPTIVE_SAMPLE_STEP - 1)) == 0,
+              "ADAPTIVE_SAMPLE_STEP must be power of two for bitwise operations to work");
+
 /* Split kernel constants */
 #define WORK_POOL_SIZE_GPU 64
 #define WORK_POOL_SIZE_CPU 1
@@ -106,6 +111,8 @@ CCL_NAMESPACE_BEGIN
 #ifndef __KERNEL_AO_PREVIEW__
 #  define __SVM__
 #  define __EMISSION__
+#  define __TEXTURES__
+#  define __EXTRA_NODES__
 #  define __HOLDOUT__
 #  define __MULTI_CLOSURE__
 #  define __TRANSPARENT_SHADOWS__
@@ -1237,9 +1244,7 @@ typedef struct KernelFilm {
 
   int pass_aov_color;
   int pass_aov_value;
-  int pass_aov_color_num;
-  int pass_aov_value_num;
-  int pad1, pad2, pad3;
+  int pad1;
 
   /* XYZ to rendering color space transform. float4 instead of float3 to
    * ensure consistent padding/alignment across devices. */
@@ -1262,7 +1267,7 @@ typedef struct KernelFilm {
   int use_display_exposure;
   int use_display_pass_alpha;
 
-  int pad4, pad5, pad6;
+  int pad3, pad4, pad5;
 } KernelFilm;
 static_assert_align(KernelFilm, 16);
 
@@ -1345,8 +1350,6 @@ typedef struct KernelIntegrator {
   int sampling_pattern;
   int aa_samples;
   int adaptive_min_samples;
-  int adaptive_step;
-  int adaptive_stop_per_sample;
   float adaptive_threshold;
 
   /* volume render */
@@ -1359,7 +1362,7 @@ typedef struct KernelIntegrator {
 
   int max_closures;
 
-  int pad1;
+  int pad1, pad2, pad3;
 } KernelIntegrator;
 static_assert_align(KernelIntegrator, 16);
 

intern/cycles/kernel/svm/svm.h

@@ -231,8 +231,6 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
     uint4 node = read_node(kg, &offset);
 
     switch (node.x) {
-      case NODE_END:
-        return;
 #if NODES_GROUP(NODE_GROUP_LEVEL_0)
       case NODE_SHADER_JUMP: {
         if (type == SHADER_TYPE_SURFACE)
@@ -312,6 +310,7 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
         svm_node_vector_displacement(kg, sd, stack, node, &offset);
         break;
 #  endif /* NODES_FEATURE(NODE_FEATURE_BUMP) */
+#  ifdef __TEXTURES__
       case NODE_TEX_IMAGE:
         svm_node_tex_image(kg, sd, stack, node, &offset);
         break;
@@ -321,7 +320,9 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
       case NODE_TEX_NOISE:
         svm_node_tex_noise(kg, sd, stack, node.y, node.z, node.w, &offset);
         break;
-#  if NODES_FEATURE(NODE_FEATURE_BUMP)
+#  endif /* __TEXTURES__ */
+#  ifdef __EXTRA_NODES__
+#    if NODES_FEATURE(NODE_FEATURE_BUMP)
       case NODE_SET_BUMP:
         svm_node_set_bump(kg, sd, stack, node);
         break;
@@ -346,19 +347,20 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
       case NODE_CLOSURE_SET_NORMAL:
         svm_node_set_normal(kg, sd, stack, node.y, node.z);
         break;
-#    if NODES_FEATURE(NODE_FEATURE_BUMP_STATE)
+#      if NODES_FEATURE(NODE_FEATURE_BUMP_STATE)
       case NODE_ENTER_BUMP_EVAL:
         svm_node_enter_bump_eval(kg, sd, stack, node.y);
         break;
       case NODE_LEAVE_BUMP_EVAL:
         svm_node_leave_bump_eval(kg, sd, stack, node.y);
         break;
-#    endif /* NODES_FEATURE(NODE_FEATURE_BUMP_STATE) */
-#  endif   /* NODES_FEATURE(NODE_FEATURE_BUMP) */
+#      endif /* NODES_FEATURE(NODE_FEATURE_BUMP_STATE) */
+#    endif   /* NODES_FEATURE(NODE_FEATURE_BUMP) */
       case NODE_HSV:
         svm_node_hsv(kg, sd, stack, node, &offset);
         break;
-#endif /* NODES_GROUP(NODE_GROUP_LEVEL_0) */
+#  endif /* __EXTRA_NODES__ */
+#endif   /* NODES_GROUP(NODE_GROUP_LEVEL_0) */
 
 #if NODES_GROUP(NODE_GROUP_LEVEL_1)
       case NODE_CLOSURE_HOLDOUT:
@@ -378,6 +380,7 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
         svm_node_principled_volume(kg, sd, stack, node, type, path_flag, &offset);
         break;
 #  endif /* NODES_FEATURE(NODE_FEATURE_VOLUME) */
+#  ifdef __EXTRA_NODES__
       case NODE_MATH:
         svm_node_math(kg, sd, stack, node.y, node.z, node.w, &offset);
         break;
@@ -402,12 +405,15 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
       case NODE_PARTICLE_INFO:
         svm_node_particle_info(kg, sd, stack, node.y, node.z);
         break;
-#  if defined(__HAIR__) && NODES_FEATURE(NODE_FEATURE_HAIR)
+#    ifdef __HAIR__
+#      if NODES_FEATURE(NODE_FEATURE_HAIR)
       case NODE_HAIR_INFO:
         svm_node_hair_info(kg, sd, stack, node.y, node.z);
         break;
-#  endif /* NODES_FEATURE(NODE_FEATURE_HAIR) */
-#endif   /* NODES_GROUP(NODE_GROUP_LEVEL_1) */
+#      endif /* NODES_FEATURE(NODE_FEATURE_HAIR) */
+#    endif   /* __HAIR__ */
+#  endif     /* __EXTRA_NODES__ */
+#endif       /* NODES_GROUP(NODE_GROUP_LEVEL_1) */
 
 #if NODES_GROUP(NODE_GROUP_LEVEL_2)
       case NODE_TEXTURE_MAPPING:
@@ -422,6 +428,7 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
       case NODE_CAMERA:
         svm_node_camera(kg, sd, stack, node.y, node.z, node.w);
         break;
+#  ifdef __TEXTURES__
       case NODE_TEX_ENVIRONMENT:
         svm_node_tex_environment(kg, sd, stack, node);
         break;
@@ -452,6 +459,8 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
       case NODE_TEX_WHITE_NOISE:
         svm_node_tex_white_noise(kg, sd, stack, node.y, node.z, node.w, &offset);
         break;
+#  endif /* __TEXTURES__ */
+#  ifdef __EXTRA_NODES__
       case NODE_NORMAL:
         svm_node_normal(kg, sd, stack, node.y, node.z, node.w, &offset);
         break;
@@ -461,7 +470,19 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
       case NODE_IES:
         svm_node_ies(kg, sd, stack, node, &offset);
         break;
-#endif /* NODES_GROUP(NODE_GROUP_LEVEL_2) */
+      case NODE_AOV_START:
+        if (!svm_node_aov_check(state, buffer)) {
+          return;
+        }
+        break;
+      case NODE_AOV_COLOR:
+        svm_node_aov_color(kg, sd, stack, node, buffer);
+        break;
+      case NODE_AOV_VALUE:
+        svm_node_aov_value(kg, sd, stack, node, buffer);
+        break;
+#  endif /* __EXTRA_NODES__ */
+#endif   /* NODES_GROUP(NODE_GROUP_LEVEL_2) */
 
 #if NODES_GROUP(NODE_GROUP_LEVEL_3)
       case NODE_RGB_CURVES:
@@ -474,6 +495,7 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
       case NODE_NORMAL_MAP:
         svm_node_normal_map(kg, sd, stack, node);
         break;
+#  ifdef __EXTRA_NODES__
       case NODE_INVERT:
         svm_node_invert(sd, stack, node.y, node.z, node.w);
         break;
@@ -513,6 +535,12 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
       case NODE_CLAMP:
         svm_node_clamp(kg, sd, stack, node.y, node.z, node.w, &offset);
         break;
+#  endif /* __EXTRA_NODES__ */
+#  if NODES_FEATURE(NODE_FEATURE_VOLUME)
+      case NODE_TEX_VOXEL:
+        svm_node_tex_voxel(kg, sd, stack, node, &offset);
+        break;
+#  endif /* NODES_FEATURE(NODE_FEATURE_VOLUME) */
 #  ifdef __SHADER_RAYTRACE__
       case NODE_BEVEL:
         svm_node_bevel(kg, sd, state, stack, node);
@@ -522,25 +550,8 @@ ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg,
         break;
 #  endif /* __SHADER_RAYTRACE__ */
 #endif   /* NODES_GROUP(NODE_GROUP_LEVEL_3) */
-
-#if NODES_GROUP(NODE_GROUP_LEVEL_4)
-#  if NODES_FEATURE(NODE_FEATURE_VOLUME)
-      case NODE_TEX_VOXEL:
-        svm_node_tex_voxel(kg, sd, stack, node, &offset);
-        break;
-#  endif /* NODES_FEATURE(NODE_FEATURE_VOLUME) */
-      case NODE_AOV_START:
-        if (!svm_node_aov_check(state, buffer)) {
-          return;
-        }
-        break;
-      case NODE_AOV_COLOR:
-        svm_node_aov_color(kg, sd, stack, node, buffer);
-        break;
-      case NODE_AOV_VALUE:
-        svm_node_aov_value(kg, sd, stack, node, buffer);
-        break;
-#endif /* NODES_GROUP(NODE_GROUP_LEVEL_4) */
+      case NODE_END:
+        return;
       default:
         kernel_assert(!"Unknown node type was passed to the SVM machine");
         return;

intern/cycles/kernel/svm/svm_closure.h

@@ -16,6 +16,23 @@
 
 CCL_NAMESPACE_BEGIN
 
+/* Hair Melanin */
+
+ccl_device_inline float3 sigma_from_concentration(float eumelanin, float pheomelanin)
+{
+  return eumelanin * make_float3(0.506f, 0.841f, 1.653f) +
+         pheomelanin * make_float3(0.343f, 0.733f, 1.924f);
+}
+
+ccl_device_inline float3 sigma_from_reflectance(float3 color, float azimuthal_roughness)
+{
+  float x = azimuthal_roughness;
+  float roughness_fac = (((((0.245f * x) + 5.574f) * x - 10.73f) * x + 2.532f) * x - 0.215f) * x +
+                        5.969f;
+  float3 sigma = log3(color) / roughness_fac;
+  return sigma * sigma;
+}
+
 /* Closure Nodes */
 
 ccl_device void svm_node_glass_setup(
@@ -851,26 +868,24 @@ ccl_device void svm_node_closure_bsdf(KernelGlobals *kg,
             /* Benedikt Bitterli's melanin ratio remapping. */
             float eumelanin = melanin * (1.0f - melanin_redness);
             float pheomelanin = melanin * melanin_redness;
-            float3 melanin_sigma = bsdf_principled_hair_sigma_from_concentration(eumelanin,
-                                                                                 pheomelanin);
+            float3 melanin_sigma = sigma_from_concentration(eumelanin, pheomelanin);
 
             /* Optional tint. */
             float3 tint = stack_load_float3(stack, tint_ofs);
-            float3 tint_sigma = bsdf_principled_hair_sigma_from_reflectance(tint,
-                                                                            radial_roughness);
+            float3 tint_sigma = sigma_from_reflectance(tint, radial_roughness);
 
             bsdf->sigma = melanin_sigma + tint_sigma;
             break;
           }
           case NODE_PRINCIPLED_HAIR_REFLECTANCE: {
             float3 color = stack_load_float3(stack, color_ofs);
-            bsdf->sigma = bsdf_principled_hair_sigma_from_reflectance(color, radial_roughness);
+            bsdf->sigma = sigma_from_reflectance(color, radial_roughness);
             break;
           }
           default: {
             /* Fallback to brownish hair, same as defaults for melanin. */
             kernel_assert(!"Invalid Principled Hair parametrization!");
-            bsdf->sigma = bsdf_principled_hair_sigma_from_concentration(0.0f, 0.8054375f);
+            bsdf->sigma = sigma_from_concentration(0.0f, 0.8054375f);
             break;
           }
         }

intern/cycles/kernel/svm/svm_image.h

@@ -16,6 +16,8 @@
 
 CCL_NAMESPACE_BEGIN
 
+#ifdef __TEXTURES__
+
 ccl_device float4 svm_image_texture(KernelGlobals *kg, int id, float x, float y, uint flags)
 {
   if (id == -1) {
@@ -244,4 +246,6 @@ ccl_device void svm_node_tex_environment(KernelGlobals *kg,
     stack_store_float(stack, alpha_offset, f.w);
 }
 
+#endif /* __TEXTURES__ */
+
 CCL_NAMESPACE_END

intern/cycles/kernel/svm/svm_math.h

@@ -51,7 +51,7 @@ ccl_device void svm_node_vector_math(KernelGlobals *kg,
 
   float3 a = stack_load_float3(stack, a_stack_offset);
   float3 b = stack_load_float3(stack, b_stack_offset);
-  float3 c = make_float3(0.0f, 0.0f, 0.0f);
+  float3 c;
   float scale = stack_load_float(stack, scale_stack_offset);
 
   float value;

intern/cycles/kernel/svm/svm_noise.h

@@ -65,7 +65,7 @@ ccl_device_noinline_cpu float perlin_1d(float x)
  * supported, we do a standard implementation, but if it is supported, we
  * do an implementation using SSE intrinsics.
  */
-#if !defined(__KERNEL_SSE2__)
+#ifndef __KERNEL_SSE2__
 
 /* ** Standard Implementation ** */
 
@@ -266,7 +266,7 @@ ccl_device_noinline_cpu float perlin_4d(float x, float y, float z, float w)
   return r;
 }
 
-#else /* SSE is supported. */
+#else
 
 /* ** SSE Implementation ** */
 
@@ -300,57 +300,6 @@ ccl_device_inline ssef bi_mix(ssef p, ssef f)
   return mix(g, shuffle<1>(g), shuffle<1>(f));
 }
 
-ccl_device_inline ssef fade(const ssef &t)
-{
-  ssef a = madd(t, 6.0f, -15.0f);
-  ssef b = madd(t, a, 10.0f);
-  return (t * t) * (t * b);
-}
-
-/* Negate val if the nth bit of h is 1. */
-#  define negate_if_nth_bit(val, h, n) ((val) ^ cast(((h) & (1 << (n))) << (31 - (n))))
-
-ccl_device_inline ssef grad(const ssei &hash, const ssef &x, const ssef &y)
-{
-  ssei h = hash & 7;
-  ssef u = select(h < 4, x, y);
-  ssef v = 2.0f * select(h < 4, y, x);
-  return negate_if_nth_bit(u, h, 0) + negate_if_nth_bit(v, h, 1);
-}
-
-/* We use SSE to compute and interpolate 4 gradients at once:
- *
- *    Point  Offset from v0
- *     v0       (0, 0)
- *     v1       (0, 1)
- *     v2       (1, 0)    (0, 1, 0, 1) = shuffle<0, 2, 0, 2>(shuffle<1, 1, 1, 1>(V, V + 1))
- *     v3       (1, 1)         ^
- *               |  |__________|       (0, 0, 1, 1) = shuffle<0, 0, 0, 0>(V, V + 1)
- *               |                          ^
- *               |__________________________|
- *
- */
-ccl_device_noinline float perlin_2d(float x, float y)
-{
-  ssei XY;
-  ssef fxy = floorfrac(ssef(x, y, 0.0f, 0.0f), &XY);
-  ssef uv = fade(fxy);
-
-  ssei XY1 = XY + 1;
-  ssei X = shuffle<0, 0, 0, 0>(XY, XY1);
-  ssei Y = shuffle<0, 2, 0, 2>(shuffle<1, 1, 1, 1>(XY, XY1));
-
-  ssei h = hash_ssei2(X, Y);
-
-  ssef fxy1 = fxy - 1.0f;
-  ssef fx = shuffle<0, 0, 0, 0>(fxy, fxy1);
-  ssef fy = shuffle<0, 2, 0, 2>(shuffle<1, 1, 1, 1>(fxy, fxy1));
-
-  ssef g = grad(h, fx, fy);
-
-  return extract<0>(bi_mix(g, uv));
-}
-
 /* SSE Trilinear Interpolation:
  *
  * The function takes three ssef inputs:
@@ -391,12 +340,34 @@ ccl_device_inline ssef tri_mix(ssef p, ssef q, ssef f)
   return mix(g, shuffle<1>(g), shuffle<2>(f));
 }
 
-/* 3D and 4D noise can be accelerated using AVX, so we first check if AVX
- * is supported, that is, if __KERNEL_AVX__ is defined. If it is not
- * supported, we do an SSE implementation, but if it is supported,
- * we do an implementation using AVX intrinsics.
+/* SSE Quadrilinear Interpolation:
+ *
+ * Quadrilinear interpolation is as simple as a linear interpolation
+ * between two trilinear interpolations.
+ *
  */
-#  if !defined(__KERNEL_AVX__)
+ccl_device_inline ssef quad_mix(ssef p, ssef q, ssef r, ssef s, ssef f)
+{
+  return mix(tri_mix(p, q, f), tri_mix(r, s, f), shuffle<3>(f));
+}
+
+ccl_device_inline ssef fade(const ssef &t)
+{
+  ssef a = madd(t, 6.0f, -15.0f);
+  ssef b = madd(t, a, 10.0f);
+  return (t * t) * (t * b);
+}
+
+/* Negate val if the nth bit of h is 1. */
+#  define negate_if_nth_bit(val, h, n) ((val) ^ cast(((h) & (1 << (n))) << (31 - (n))))
+
+ccl_device_inline ssef grad(const ssei &hash, const ssef &x, const ssef &y)
+{
+  ssei h = hash & 7;
+  ssef u = select(h < 4, x, y);
+  ssef v = 2.0f * select(h < 4, y, x);
+  return negate_if_nth_bit(u, h, 0) + negate_if_nth_bit(v, h, 1);
+}
 
 ccl_device_inline ssef grad(const ssei &hash, const ssef &x, const ssef &y, const ssef &z)
 {
@@ -417,15 +388,37 @@ grad(const ssei &hash, const ssef &x, const ssef &y, const ssef &z, const ssef &
   return negate_if_nth_bit(u, h, 0) + negate_if_nth_bit(v, h, 1) + negate_if_nth_bit(s, h, 2);
 }
 
-/* SSE Quadrilinear Interpolation:
+/* We use SSE to compute and interpolate 4 gradients at once:
  *
- * Quadrilinear interpolation is as simple as a linear interpolation
- * between two trilinear interpolations.
+ *    Point  Offset from v0
+ *     v0       (0, 0)
+ *     v1       (0, 1)
+ *     v2       (1, 0)    (0, 1, 0, 1) = shuffle<0, 2, 0, 2>(shuffle<1, 1, 1, 1>(V, V + 1))
+ *     v3       (1, 1)         ^
+ *               |  |__________|       (0, 0, 1, 1) = shuffle<0, 0, 0, 0>(V, V + 1)
+ *               |                          ^
+ *               |__________________________|
  *
  */
-ccl_device_inline ssef quad_mix(ssef p, ssef q, ssef r, ssef s, ssef f)
+ccl_device_noinline float perlin_2d(float x, float y)
 {
-  return mix(tri_mix(p, q, f), tri_mix(r, s, f), shuffle<3>(f));
+  ssei XY;
+  ssef fxy = floorfrac(ssef(x, y, 0.0f, 0.0f), &XY);
+  ssef uv = fade(fxy);
+
+  ssei XY1 = XY + 1;
+  ssei X = shuffle<0, 0, 0, 0>(XY, XY1);
+  ssei Y = shuffle<0, 2, 0, 2>(shuffle<1, 1, 1, 1>(XY, XY1));
+
+  ssei h = hash_ssei2(X, Y);
+
+  ssef fxy1 = fxy - 1.0f;
+  ssef fx = shuffle<0, 0, 0, 0>(fxy, fxy1);
+  ssef fy = shuffle<0, 2, 0, 2>(shuffle<1, 1, 1, 1>(fxy, fxy1));
+
+  ssef g = grad(h, fx, fy);
+
+  return extract<0>(bi_mix(g, uv));
 }
 
 /* We use SSE to compute and interpolate 4 gradients at once. Since we have 8
@@ -529,148 +522,6 @@ ccl_device_noinline float perlin_4d(float x, float y, float z, float w)
 
   return extract<0>(quad_mix(g1, g2, g3, g4, uvws));
 }
-
-#  else /* AVX is supported. */
-
-/* AVX Implementation */
-
-ccl_device_inline avxf grad(const avxi &hash, const avxf &x, const avxf &y, const avxf &z)
-{
-  avxi h = hash & 15;
-  avxf u = select(h < 8, x, y);
-  avxf vt = select((h == 12) | (h == 14), x, z);
-  avxf v = select(h < 4, y, vt);
-  return negate_if_nth_bit(u, h, 0) + negate_if_nth_bit(v, h, 1);
-}
-
-ccl_device_inline avxf
-grad(const avxi &hash, const avxf &x, const avxf &y, const avxf &z, const avxf &w)
-{
-  avxi h = hash & 31;
-  avxf u = select(h < 24, x, y);
-  avxf v = select(h < 16, y, z);
-  avxf s = select(h < 8, z, w);
-  return negate_if_nth_bit(u, h, 0) + negate_if_nth_bit(v, h, 1) + negate_if_nth_bit(s, h, 2);
-}
-
-/* SSE Quadrilinear Interpolation:
- *
- * The interpolation is done in two steps:
- * 1. Interpolate p and q along the w axis to get s.
- * 2. Trilinearly interpolate (s0, s1, s2, s3) and (s4, s5, s6, s7) to get the final
- *    value. (s0, s1, s2, s3) and (s4, s5, s6, s7) are generated by extracting the
- *    low and high ssef from s.
- *
- */
-ccl_device_inline ssef quad_mix(avxf p, avxf q, ssef f)
-{
-  ssef fv = shuffle<3>(f);
-  avxf s = mix(p, q, avxf(fv, fv));
-  return tri_mix(low(s), high(s), f);
-}
-
-/* We use AVX to compute and interpolate 8 gradients at once.
- *
- *    Point  Offset from v0
- *     v0      (0, 0, 0)
- *     v1      (0, 0, 1)    The full avx type is computed by inserting the following
- *     v2      (0, 1, 0)    sse types into both the low and high parts of the avx.
- *     v3      (0, 1, 1)
- *     v4      (1, 0, 0)
- *     v5      (1, 0, 1)    (0, 1, 0, 1) = shuffle<0, 2, 0, 2>(shuffle<2, 2, 2, 2>(V, V + 1))
- *     v6      (1, 1, 0)         ^
- *     v7      (1, 1, 1)         |
- *                 |  |__________|       (0, 0, 1, 1) = shuffle<1, 1, 1, 1>(V, V + 1)
- *                 |                          ^
- *                 |__________________________|
- *
- */
-ccl_device_noinline float perlin_3d(float x, float y, float z)
-{
-  ssei XYZ;
-  ssef fxyz = floorfrac(ssef(x, y, z, 0.0f), &XYZ);
-  ssef uvw = fade(fxyz);
-
-  ssei XYZ1 = XYZ + 1;
-  ssei X = shuffle<0>(XYZ);
-  ssei X1 = shuffle<0>(XYZ1);
-  ssei Y = shuffle<1, 1, 1, 1>(XYZ, XYZ1);
-  ssei Z = shuffle<0, 2, 0, 2>(shuffle<2, 2, 2, 2>(XYZ, XYZ1));
-
-  avxi h = hash_avxi3(avxi(X, X1), avxi(Y, Y), avxi(Z, Z));
-
-  ssef fxyz1 = fxyz - 1.0f;
-  ssef fx = shuffle<0>(fxyz);
-  ssef fx1 = shuffle<0>(fxyz1);
-  ssef fy = shuffle<1, 1, 1, 1>(fxyz, fxyz1);
-  ssef fz = shuffle<0, 2, 0, 2>(shuffle<2, 2, 2, 2>(fxyz, fxyz1));
-
-  avxf g = grad(h, avxf(fx, fx1), avxf(fy, fy), avxf(fz, fz));
-
-  return extract<0>(tri_mix(low(g), high(g), uvw));
-}
-
-/* We use AVX to compute and interpolate 8 gradients at once. Since we have 16
- * gradients in 4D, we need to compute two sets of gradients at the points:
- *
- *    Point  Offset from v0
- *     v0     (0, 0, 0, 0)
- *     v1     (0, 0, 1, 0)  The full avx type is computed by inserting the following
- *     v2     (0, 1, 0, 0)  sse types into both the low and high parts of the avx.
- *     v3     (0, 1, 1, 0)
- *     v4     (1, 0, 0, 0)
- *     v5     (1, 0, 1, 0)  (0, 1, 0, 1) = shuffle<0, 2, 0, 2>(shuffle<2, 2, 2, 2>(V, V + 1))
- *     v6     (1, 1, 0, 0)    ^
- *     v7     (1, 1, 1, 0)    |
- *                |  |________|    (0, 0, 1, 1) = shuffle<1, 1, 1, 1>(V, V + 1)
- *                |                       ^
- *                |_______________________|
- *
- *    Point  Offset from v0
- *     v8     (0, 0, 0, 1)
- *     v9     (0, 0, 1, 1)
- *     v10    (0, 1, 0, 1)
- *     v11    (0, 1, 1, 1)
- *     v12    (1, 0, 0, 1)
- *     v13    (1, 0, 1, 1)
- *     v14    (1, 1, 0, 1)
- *     v15    (1, 1, 1, 1)
- *
- */
-ccl_device_noinline float perlin_4d(float x, float y, float z, float w)
-{
-  ssei XYZW;
-  ssef fxyzw = floorfrac(ssef(x, y, z, w), &XYZW);
-  ssef uvws = fade(fxyzw);
-
-  ssei XYZW1 = XYZW + 1;
-  ssei X = shuffle<0>(XYZW);
-  ssei X1 = shuffle<0>(XYZW1);
-  ssei Y = shuffle<1, 1, 1, 1>(XYZW, XYZW1);
-  ssei Z = shuffle<0, 2, 0, 2>(shuffle<2, 2, 2, 2>(XYZW, XYZW1));
-  ssei W = shuffle<3>(XYZW);
-  ssei W1 = shuffle<3>(XYZW1);
-
-  avxi h1 = hash_avxi4(avxi(X, X1), avxi(Y, Y), avxi(Z, Z), avxi(W, W));
-  avxi h2 = hash_avxi4(avxi(X, X1), avxi(Y, Y), avxi(Z, Z), avxi(W1, W1));
-
-  ssef fxyzw1 = fxyzw - 1.0f;
-  ssef fx = shuffle<0>(fxyzw);
-  ssef fx1 = shuffle<0>(fxyzw1);
-  ssef fy = shuffle<1, 1, 1, 1>(fxyzw, fxyzw1);
-  ssef fz = shuffle<0, 2, 0, 2>(shuffle<2, 2, 2, 2>(fxyzw, fxyzw1));
-  ssef fw = shuffle<3>(fxyzw);
-  ssef fw1 = shuffle<3>(fxyzw1);
-
-  avxf g1 = grad(h1, avxf(fx, fx1), avxf(fy, fy), avxf(fz, fz), avxf(fw, fw));
-  avxf g2 = grad(h2, avxf(fx, fx1), avxf(fy, fy), avxf(fz, fz), avxf(fw1, fw1));
-
-  return extract<0>(quad_mix(g1, g2, uvws));
-}
-#  endif
-
-#  undef negate_if_nth_bit
-
 #endif
 
 /* Remap the output of noise to a predictable range [-1, 1].

intern/cycles/kernel/svm/svm_types.h

@@ -42,8 +42,7 @@ CCL_NAMESPACE_BEGIN
 #define NODE_GROUP_LEVEL_1 1
 #define NODE_GROUP_LEVEL_2 2
 #define NODE_GROUP_LEVEL_3 3
-#define NODE_GROUP_LEVEL_4 4
-#define NODE_GROUP_LEVEL_MAX NODE_GROUP_LEVEL_4
+#define NODE_GROUP_LEVEL_MAX NODE_GROUP_LEVEL_3
 
 #define NODE_FEATURE_VOLUME (1 << 0)
 #define NODE_FEATURE_HAIR (1 << 1)
@@ -63,98 +62,98 @@ CCL_NAMESPACE_BEGIN
 
 typedef enum ShaderNodeType {
   NODE_END = 0,
-  NODE_SHADER_JUMP,
   NODE_CLOSURE_BSDF,
   NODE_CLOSURE_EMISSION,
   NODE_CLOSURE_BACKGROUND,
   NODE_CLOSURE_SET_WEIGHT,
   NODE_CLOSURE_WEIGHT,
-  NODE_EMISSION_WEIGHT,
   NODE_MIX_CLOSURE,
   NODE_JUMP_IF_ZERO,
   NODE_JUMP_IF_ONE,
-  NODE_GEOMETRY,
-  NODE_CONVERT,
-  NODE_TEX_COORD,
-  NODE_VALUE_F,
-  NODE_VALUE_V,
-  NODE_ATTR,
-  NODE_VERTEX_COLOR,
-  NODE_GEOMETRY_BUMP_DX,
-  NODE_GEOMETRY_BUMP_DY,
-  NODE_SET_DISPLACEMENT,
-  NODE_DISPLACEMENT,
-  NODE_VECTOR_DISPLACEMENT,
   NODE_TEX_IMAGE,
   NODE_TEX_IMAGE_BOX,
-  NODE_TEX_NOISE,
-  NODE_SET_BUMP,
-  NODE_ATTR_BUMP_DX,
-  NODE_ATTR_BUMP_DY,
-  NODE_VERTEX_COLOR_BUMP_DX,
-  NODE_VERTEX_COLOR_BUMP_DY,
-  NODE_TEX_COORD_BUMP_DX,
-  NODE_TEX_COORD_BUMP_DY,
-  NODE_CLOSURE_SET_NORMAL,
-  NODE_ENTER_BUMP_EVAL,
-  NODE_LEAVE_BUMP_EVAL,
-  NODE_HSV,
-  NODE_CLOSURE_HOLDOUT,
-  NODE_FRESNEL,
-  NODE_LAYER_WEIGHT,
-  NODE_CLOSURE_VOLUME,
-  NODE_PRINCIPLED_VOLUME,
-  NODE_MATH,
-  NODE_VECTOR_MATH,
-  NODE_RGB_RAMP,
-  NODE_GAMMA,
-  NODE_BRIGHTCONTRAST,
-  NODE_LIGHT_PATH,
-  NODE_OBJECT_INFO,
-  NODE_PARTICLE_INFO,
-  NODE_HAIR_INFO,
-  NODE_TEXTURE_MAPPING,
-  NODE_MAPPING,
-  NODE_MIN_MAX,
-  NODE_CAMERA,
-  NODE_TEX_ENVIRONMENT,
   NODE_TEX_SKY,
+  NODE_GEOMETRY,
+  NODE_GEOMETRY_DUPLI,
+  NODE_LIGHT_PATH,
+  NODE_VALUE_F,
+  NODE_VALUE_V,
+  NODE_MIX,
+  NODE_ATTR,
+  NODE_CONVERT,
+  NODE_FRESNEL,
+  NODE_WIREFRAME,
+  NODE_WAVELENGTH,
+  NODE_BLACKBODY,
+  NODE_EMISSION_WEIGHT,
   NODE_TEX_GRADIENT,
   NODE_TEX_VORONOI,
   NODE_TEX_MUSGRAVE,
   NODE_TEX_WAVE,
   NODE_TEX_MAGIC,
-  NODE_TEX_CHECKER,
-  NODE_TEX_BRICK,
-  NODE_TEX_WHITE_NOISE,
-  NODE_NORMAL,
-  NODE_LIGHT_FALLOFF,
-  NODE_IES,
-  NODE_RGB_CURVES,
-  NODE_VECTOR_CURVES,
-  NODE_TANGENT,
-  NODE_NORMAL_MAP,
-  NODE_INVERT,
-  NODE_MIX,
+  NODE_TEX_NOISE,
+  NODE_SHADER_JUMP,
+  NODE_SET_DISPLACEMENT,
+  NODE_GEOMETRY_BUMP_DX,
+  NODE_GEOMETRY_BUMP_DY,
+  NODE_SET_BUMP,
+  NODE_MATH,
+  NODE_VECTOR_MATH,
+  NODE_VECTOR_TRANSFORM,
+  NODE_MAPPING,
+  NODE_TEX_COORD,
+  NODE_TEX_COORD_BUMP_DX,
+  NODE_TEX_COORD_BUMP_DY,
+  NODE_ATTR_BUMP_DX,
+  NODE_ATTR_BUMP_DY,
+  NODE_TEX_ENVIRONMENT,
+  NODE_CLOSURE_HOLDOUT,
+  NODE_LAYER_WEIGHT,
+  NODE_CLOSURE_VOLUME,
   NODE_SEPARATE_VECTOR,
   NODE_COMBINE_VECTOR,
   NODE_SEPARATE_HSV,
   NODE_COMBINE_HSV,
-  NODE_VECTOR_ROTATE,
-  NODE_VECTOR_TRANSFORM,
-  NODE_WIREFRAME,
-  NODE_WAVELENGTH,
-  NODE_BLACKBODY,
+  NODE_HSV,
+  NODE_CAMERA,
+  NODE_INVERT,
+  NODE_NORMAL,
+  NODE_GAMMA,
+  NODE_TEX_CHECKER,
+  NODE_BRIGHTCONTRAST,
+  NODE_RGB_RAMP,
+  NODE_RGB_CURVES,
+  NODE_VECTOR_CURVES,
+  NODE_MIN_MAX,
+  NODE_LIGHT_FALLOFF,
+  NODE_OBJECT_INFO,
+  NODE_PARTICLE_INFO,
+  NODE_TEX_BRICK,
+  NODE_CLOSURE_SET_NORMAL,
+  NODE_AMBIENT_OCCLUSION,
+  NODE_TANGENT,
+  NODE_NORMAL_MAP,
+  NODE_HAIR_INFO,
+  NODE_UVMAP,
+  NODE_TEX_VOXEL,
+  NODE_ENTER_BUMP_EVAL,
+  NODE_LEAVE_BUMP_EVAL,
+  NODE_BEVEL,
+  NODE_DISPLACEMENT,
+  NODE_VECTOR_DISPLACEMENT,
+  NODE_PRINCIPLED_VOLUME,
+  NODE_IES,
   NODE_MAP_RANGE,
   NODE_CLAMP,
-  NODE_BEVEL,
-  NODE_AMBIENT_OCCLUSION,
-  NODE_TEX_VOXEL,
+  NODE_TEXTURE_MAPPING,
+  NODE_TEX_WHITE_NOISE,
+  NODE_VERTEX_COLOR,
+  NODE_VERTEX_COLOR_BUMP_DX,
+  NODE_VERTEX_COLOR_BUMP_DY,
   NODE_AOV_START,
-  NODE_AOV_COLOR,
   NODE_AOV_VALUE,
-  /* NOTE: for best OpenCL performance, item definition in the enum must
-   * match the switch case order in svm.h. */
+  NODE_AOV_COLOR,
+  NODE_VECTOR_ROTATE,
 } ShaderNodeType;
 
 typedef enum NodeAttributeType {

intern/cycles/render/attribute.cpp

@@ -48,7 +48,7 @@ Attribute::Attribute(
 Attribute::~Attribute()
 {
   /* For voxel data, we need to free the image handle. */
-  if (element == ATTR_ELEMENT_VOXEL && buffer.size()) {
+  if (element == ATTR_ELEMENT_VOXEL) {
     ImageHandle &handle = data_voxel();
     handle.~ImageHandle();
   }

intern/cycles/render/attribute.h

@@ -59,11 +59,7 @@ class Attribute {
             AttributeElement element,
             Geometry *geom,
             AttributePrimitive prim);
-  Attribute(Attribute &&other) = default;
-  Attribute(const Attribute &other) = delete;
-  Attribute &operator=(const Attribute &other) = delete;
   ~Attribute();
-
   void set(ustring name, TypeDesc type, AttributeElement element);
   void resize(Geometry *geom, AttributePrimitive prim, bool reserve_only);
   void resize(size_t num_elements);

intern/cycles/render/film.cpp

@@ -203,10 +203,9 @@ void Pass::add(PassType type, vector<Pass> &passes, const char *name)
 
   passes.push_back(pass);
 
-  /* Order from by components, to ensure alignment so passes with size 4
-   * come first and then passes with size 1. Note this must use stable sort
-   * so cryptomatte passes remain in the right order. */
-  stable_sort(&passes[0], &passes[0] + passes.size(), compare_pass_order);
+  /* order from by components, to ensure alignment so passes with size 4
+   * come first and then passes with size 1 */
+  sort(&passes[0], &passes[0] + passes.size(), compare_pass_order);
 
   if (pass.divide_type != PASS_NONE)
     Pass::add(pass.divide_type, passes);
@@ -362,10 +361,8 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
   kfilm->light_pass_flag = 0;
   kfilm->pass_stride = 0;
   kfilm->use_light_pass = use_light_visibility;
-  kfilm->pass_aov_value_num = 0;
-  kfilm->pass_aov_color_num = 0;
 
-  bool have_cryptomatte = false;
+  bool have_cryptomatte = false, have_aov_color = false, have_aov_value = false;
 
   for (size_t i = 0; i < passes.size(); i++) {
     Pass &pass = passes[i];
@@ -500,16 +497,16 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
         kfilm->pass_sample_count = kfilm->pass_stride;
         break;
       case PASS_AOV_COLOR:
-        if (kfilm->pass_aov_color_num == 0) {
+        if (!have_aov_color) {
           kfilm->pass_aov_color = kfilm->pass_stride;
+          have_aov_color = true;
         }
-        kfilm->pass_aov_color_num++;
         break;
       case PASS_AOV_VALUE:
-        if (kfilm->pass_aov_value_num == 0) {
+        if (!have_aov_value) {
           kfilm->pass_aov_value = kfilm->pass_stride;
+          have_aov_value = true;
         }
-        kfilm->pass_aov_value_num++;
         break;
       default:
         assert(false);

intern/cycles/render/image.cpp

@@ -119,9 +119,6 @@ void ImageHandle::clear()
   foreach (const int slot, tile_slots) {
     manager->remove_image_user(slot);
   }
-
-  tile_slots.clear();
-  manager = NULL;
 }
 
 bool ImageHandle::empty()
@@ -283,7 +280,6 @@ void ImageManager::set_osl_texture_system(void *texture_system)
 bool ImageManager::set_animation_frame_update(int frame)
 {
   if (frame != animation_frame) {
-    thread_scoped_lock device_lock(images_mutex);
     animation_frame = frame;
 
     for (size_t slot = 0; slot < images.size(); slot++) {
@@ -378,7 +374,7 @@ int ImageManager::add_image_slot(ImageLoader *loader,
   Image *img;
   size_t slot;
 
-  thread_scoped_lock device_lock(images_mutex);
+  thread_scoped_lock device_lock(device_mutex);
 
   /* Fnd existing image. */
   for (slot = 0; slot < images.size(); slot++) {
@@ -419,7 +415,6 @@ int ImageManager::add_image_slot(ImageLoader *loader,
 
 void ImageManager::add_image_user(int slot)
 {
-  thread_scoped_lock device_lock(images_mutex);
   Image *image = images[slot];
   assert(image && image->users >= 1);
 
@@ -428,7 +423,6 @@ void ImageManager::add_image_user(int slot)
 
 void ImageManager::remove_image_user(int slot)
 {
-  thread_scoped_lock device_lock(images_mutex);
   Image *image = images[slot];
   assert(image && image->users >= 1);
 

intern/cycles/render/image.h

@@ -206,7 +206,6 @@ class ImageManager {
   bool has_half_images;
 
   thread_mutex device_mutex;
-  thread_mutex images_mutex;
   int animation_frame;
 
   vector<Image *> images;

intern/cycles/render/integrator.cpp

@@ -190,13 +190,6 @@ void Integrator::device_update(Device *device, DeviceScene *dscene, Scene *scene
   else {
     kintegrator->adaptive_min_samples = max(4, adaptive_min_samples);
   }
-
-  kintegrator->adaptive_step = 4;
-  kintegrator->adaptive_stop_per_sample = device->info.has_adaptive_stop_per_sample;
-
-  /* Adaptive step must be a power of two for bitwise operations to work. */
-  assert((kintegrator->adaptive_step & (kintegrator->adaptive_step - 1)) == 0);
-
   if (aa_samples > 0 && adaptive_threshold == 0.0f) {
     kintegrator->adaptive_threshold = max(0.001f, 1.0f / (float)aa_samples);
     VLOG(1) << "Cycles adaptive sampling: automatic threshold = "

intern/cycles/render/mesh_volume.cpp

@@ -25,18 +25,18 @@
 
 CCL_NAMESPACE_BEGIN
 
-const int64_t VOXEL_INDEX_NONE = -1;
-
-static int64_t compute_voxel_index(const int3 &resolution, int64_t x, int64_t y, int64_t z)
+static size_t compute_voxel_index(const int3 &resolution, size_t x, size_t y, size_t z)
 {
-  if (x < 0 || x >= resolution.x) {
-    return VOXEL_INDEX_NONE;
+  if (x == -1 || x >= resolution.x) {
+    return -1;
   }
-  else if (y < 0 || y >= resolution.y) {
-    return VOXEL_INDEX_NONE;
+
+  if (y == -1 || y >= resolution.y) {
+    return -1;
   }
-  else if (z < 0 || z >= resolution.z) {
-    return VOXEL_INDEX_NONE;
+
+  if (z == -1 || z >= resolution.z) {
+    return -1;
   }
 
   return x + y * resolution.x + z * resolution.x * resolution.y;
@@ -184,15 +184,15 @@ VolumeMeshBuilder::VolumeMeshBuilder(VolumeParams *volume_params)
   params = volume_params;
   number_of_nodes = 0;
 
-  const int64_t x = divide_up(params->resolution.x, CUBE_SIZE);
-  const int64_t y = divide_up(params->resolution.y, CUBE_SIZE);
-  const int64_t z = divide_up(params->resolution.z, CUBE_SIZE);
+  const size_t x = divide_up(params->resolution.x, CUBE_SIZE);
+  const size_t y = divide_up(params->resolution.y, CUBE_SIZE);
+  const size_t z = divide_up(params->resolution.z, CUBE_SIZE);
 
   /* Adding 2*pad_size since we pad in both positive and negative directions
    * along the axis. */
-  const int64_t px = divide_up(params->resolution.x + 2 * params->pad_size, CUBE_SIZE);
-  const int64_t py = divide_up(params->resolution.y + 2 * params->pad_size, CUBE_SIZE);
-  const int64_t pz = divide_up(params->resolution.z + 2 * params->pad_size, CUBE_SIZE);
+  const size_t px = divide_up(params->resolution.x + 2 * params->pad_size, CUBE_SIZE);
+  const size_t py = divide_up(params->resolution.y + 2 * params->pad_size, CUBE_SIZE);
+  const size_t pz = divide_up(params->resolution.z + 2 * params->pad_size, CUBE_SIZE);
 
   res = make_int3(px, py, pz);
   pad_offset = make_int3(px - x, py - y, pz - z);
@@ -209,10 +209,7 @@ void VolumeMeshBuilder::add_node(int x, int y, int z)
 
   assert((index_x >= 0) && (index_y >= 0) && (index_z >= 0));
 
-  const int64_t index = compute_voxel_index(res, index_x, index_y, index_z);
-  if (index == VOXEL_INDEX_NONE) {
-    return;
-  }
+  const size_t index = compute_voxel_index(res, index_x, index_y, index_z);
 
   /* We already have a node here. */
   if (grid[index] == 1) {
@@ -259,7 +256,7 @@ void VolumeMeshBuilder::generate_vertices_and_quads(vector<ccl::int3> &vertices_
   for (int z = 0; z < res.z; ++z) {
     for (int y = 0; y < res.y; ++y) {
       for (int x = 0; x < res.x; ++x) {
-        int64_t voxel_index = compute_voxel_index(res, x, y, z);
+        size_t voxel_index = compute_voxel_index(res, x, y, z);
         if (grid[voxel_index] == 0) {
           continue;
         }
@@ -288,32 +285,32 @@ void VolumeMeshBuilder::generate_vertices_and_quads(vector<ccl::int3> &vertices_
          */
 
         voxel_index = compute_voxel_index(res, x - 1, y, z);
-        if (voxel_index == VOXEL_INDEX_NONE || grid[voxel_index] == 0) {
+        if (voxel_index == -1 || grid[voxel_index] == 0) {
           create_quad(corners, vertices_is, quads, res, used_verts, QUAD_X_MIN);
         }
 
         voxel_index = compute_voxel_index(res, x + 1, y, z);
-        if (voxel_index == VOXEL_INDEX_NONE || grid[voxel_index] == 0) {
+        if (voxel_index == -1 || grid[voxel_index] == 0) {
           create_quad(corners, vertices_is, quads, res, used_verts, QUAD_X_MAX);
         }
 
         voxel_index = compute_voxel_index(res, x, y - 1, z);
-        if (voxel_index == VOXEL_INDEX_NONE || grid[voxel_index] == 0) {
+        if (voxel_index == -1 || grid[voxel_index] == 0) {
           create_quad(corners, vertices_is, quads, res, used_verts, QUAD_Y_MIN);
         }
 
         voxel_index = compute_voxel_index(res, x, y + 1, z);
-        if (voxel_index == VOXEL_INDEX_NONE || grid[voxel_index] == 0) {
+        if (voxel_index == -1 || grid[voxel_index] == 0) {
           create_quad(corners, vertices_is, quads, res, used_verts, QUAD_Y_MAX);
         }
 
         voxel_index = compute_voxel_index(res, x, y, z - 1);
-        if (voxel_index == VOXEL_INDEX_NONE || grid[voxel_index] == 0) {
+        if (voxel_index == -1 || grid[voxel_index] == 0) {
           create_quad(corners, vertices_is, quads, res, used_verts, QUAD_Z_MIN);
         }
 
         voxel_index = compute_voxel_index(res, x, y, z + 1);
-        if (voxel_index == VOXEL_INDEX_NONE || grid[voxel_index] == 0) {
+        if (voxel_index == -1 || grid[voxel_index] == 0) {
           create_quad(corners, vertices_is, quads, res, used_verts, QUAD_Z_MAX);
         }
       }
@@ -458,7 +455,7 @@ void GeometryManager::create_volume_mesh(Mesh *mesh, Progress &progress)
   for (int z = 0; z < resolution.z; ++z) {
     for (int y = 0; y < resolution.y; ++y) {
       for (int x = 0; x < resolution.x; ++x) {
-        int64_t voxel_index = compute_voxel_index(resolution, x, y, z);
+        size_t voxel_index = compute_voxel_index(resolution, x, y, z);
 
         for (size_t i = 0; i < voxel_grids.size(); ++i) {
           const VoxelAttributeGrid &voxel_grid = voxel_grids[i];

intern/cycles/render/nodes.h

@@ -197,11 +197,6 @@ class OutputAOVNode : public ShaderNode {
 
   ustring name;
 
-  virtual int get_group()
-  {
-    return NODE_GROUP_LEVEL_4;
-  }
-
   /* Don't allow output node de-duplication. */
   virtual bool equals(const ShaderNode & /*other*/)
   {
@@ -344,7 +339,7 @@ class PointDensityTextureNode : public ShaderNode {
   SHADER_NODE_NO_CLONE_CLASS(PointDensityTextureNode)
   virtual int get_group()
   {
-    return NODE_GROUP_LEVEL_4;
+    return NODE_GROUP_LEVEL_3;
   }
 
   ~PointDensityTextureNode();

intern/cycles/render/session.cpp

@@ -1110,7 +1110,6 @@ void Session::render(bool with_denoising)
   task.adaptive_sampling.use = (scene->integrator->sampling_pattern == SAMPLING_PATTERN_PMJ) &&
                                scene->dscene.data.film.pass_adaptive_aux_buffer;
   task.adaptive_sampling.min_samples = scene->dscene.data.integrator.adaptive_min_samples;
-  task.adaptive_sampling.adaptive_step = scene->dscene.data.integrator.adaptive_step;
 
   /* Acquire render tiles by default. */
   task.tile_types = RenderTile::PATH_TRACE;

intern/cycles/util/CMakeLists.txt

@@ -102,7 +102,6 @@ set(SRC_HEADERS
   util_sky_model_data.h
   util_avxf.h
   util_avxb.h
-  util_semaphore.h
   util_sseb.h
   util_ssef.h
   util_ssei.h

intern/cycles/util/util_algorithm.h

@@ -25,7 +25,6 @@ using std::max;
 using std::min;
 using std::remove;
 using std::sort;
-using std::stable_sort;
 using std::swap;
 
 CCL_NAMESPACE_END

intern/cycles/util/util_avxb.h

@@ -16,7 +16,7 @@
  */
 
 #ifndef __UTIL_AVXB_H__
-#define __UTIL_AVXB_H__
+#  define __UTIL_AVXB_H__
 
 CCL_NAMESPACE_BEGIN
 
@@ -53,10 +53,6 @@ struct avxb {
   __forceinline avxb(const __m256 input) : m256(input)
   {
   }
-  __forceinline avxb(const __m128 &a, const __m128 &b)
-      : m256(_mm256_insertf128_ps(_mm256_castps128_ps256(a), b, 1))
-  {
-  }
   __forceinline operator const __m256 &(void)const
   {
     return m256;
@@ -150,9 +146,9 @@ __forceinline const avxb operator!=(const avxb &a, const avxb &b)
 }
 __forceinline const avxb operator==(const avxb &a, const avxb &b)
 {
-#ifdef __KERNEL_AVX2__
+#  ifdef __KERNEL_AVX2__
   return _mm256_castsi256_ps(_mm256_cmpeq_epi32(a, b));
-#else
+#  else
   __m128i a_lo = _mm_castps_si128(_mm256_extractf128_ps(a, 0));
   __m128i a_hi = _mm_castps_si128(_mm256_extractf128_ps(a, 1));
   __m128i b_lo = _mm_castps_si128(_mm256_extractf128_ps(b, 0));
@@ -161,16 +157,16 @@ __forceinline const avxb operator==(const avxb &a, const avxb &b)
   __m128i c_hi = _mm_cmpeq_epi32(a_hi, b_hi);
   __m256i result = _mm256_insertf128_si256(_mm256_castsi128_si256(c_lo), c_hi, 1);
   return _mm256_castsi256_ps(result);
-#endif
+#  endif
 }
 
 __forceinline const avxb select(const avxb &m, const avxb &t, const avxb &f)
 {
-#if defined(__KERNEL_SSE41__)
+#  if defined(__KERNEL_SSE41__)
   return _mm256_blendv_ps(f, t, m);
-#else
+#  else
   return _mm256_or_ps(_mm256_and_ps(m, t), _mm256_andnot_ps(m, f));
-#endif
+#  endif
 }
 
 ////////////////////////////////////////////////////////////////////////////////
@@ -190,18 +186,18 @@ __forceinline const avxb unpackhi(const avxb &a, const avxb &b)
 /// Reduction Operations
 ////////////////////////////////////////////////////////////////////////////////
 
-#if defined(__KERNEL_SSE41__)
+#  if defined(__KERNEL_SSE41__)
 __forceinline size_t popcnt(const avxb &a)
 {
   return __popcnt(_mm256_movemask_ps(a));
 }
-#else
+#  else
 __forceinline size_t popcnt(const avxb &a)
 {
   return bool(a[0]) + bool(a[1]) + bool(a[2]) + bool(a[3]) + bool(a[4]) + bool(a[5]) + bool(a[6]) +
          bool(a[7]);
 }
-#endif
+#  endif
 
 __forceinline bool reduce_and(const avxb &a)
 {
@@ -238,6 +234,8 @@ ccl_device_inline void print_avxb(const char *label, const avxb &a)
   printf("%s: %d %d %d %d %d %d %d %d\n", label, a[0], a[1], a[2], a[3], a[4], a[5], a[6], a[7]);
 }
 
+#endif
+
 CCL_NAMESPACE_END
 
-#endif
+//#endif

intern/cycles/util/util_avxf.h

@@ -15,7 +15,7 @@
  */
 
 #ifndef __UTIL_AVXF_H__
-#define __UTIL_AVXF_H__
+#  define __UTIL_AVXF_H__
 
 CCL_NAMESPACE_BEGIN
 
@@ -140,11 +140,6 @@ __forceinline void dot3(const avxf &a, const avxf &b, float &den, float &den2)
 /// Unary Operators
 ////////////////////////////////////////////////////////////////////////////////
 
-__forceinline const avxf cast(const __m256i &a)
-{
-  return _mm256_castsi256_ps(a);
-}
-
 __forceinline const avxf mm256_sqrt(const avxf &a)
 {
   return _mm256_sqrt_ps(a.m256);
@@ -264,35 +259,16 @@ template<size_t i0> __forceinline const avxf shuffle(const avxf &a)
   return shuffle<i0>(a, a);
 }
 
-template<size_t i> __forceinline float extract(const avxf &a)
-{
-  __m256 b = shuffle<i, i, i, i>(a).m256;
-  return _mm256_cvtss_f32(b);
-}
-template<> __forceinline float extract<0>(const avxf &a)
-{
-  return _mm256_cvtss_f32(a.m256);
-}
-
-__forceinline ssef low(const avxf &a)
-{
-  return _mm256_extractf128_ps(a.m256, 0);
-}
-__forceinline ssef high(const avxf &a)
-{
-  return _mm256_extractf128_ps(a.m256, 1);
-}
-
 template<int i0, int i1, int i2, int i3, int i4, int i5, int i6, int i7>
 __forceinline const avxf permute(const avxf &a)
 {
-#ifdef __KERNEL_AVX2__
+#  ifdef __KERNEL_AVX2__
   return _mm256_permutevar8x32_ps(a, _mm256_set_epi32(i7, i6, i5, i4, i3, i2, i1, i0));
-#else
+#  else
   float temp[8];
   _mm256_storeu_ps((float *)&temp, a);
   return avxf(temp[i7], temp[i6], temp[i5], temp[i4], temp[i3], temp[i2], temp[i1], temp[i0]);
-#endif
+#  endif
 }
 
 template<int S0, int S1, int S2, int S3, int S4, int S5, int S6, int S7>
@@ -333,51 +309,39 @@ __forceinline avxf mini(const avxf &a, const avxf &b)
 ////////////////////////////////////////////////////////////////////////////////
 __forceinline const avxf madd(const avxf &a, const avxf &b, const avxf &c)
 {
-#ifdef __KERNEL_AVX2__
+#  ifdef __KERNEL_AVX2__
   return _mm256_fmadd_ps(a, b, c);
-#else
+#  else
   return c + (a * b);
-#endif
+#  endif
 }
 
 __forceinline const avxf nmadd(const avxf &a, const avxf &b, const avxf &c)
 {
-#ifdef __KERNEL_AVX2__
+#  ifdef __KERNEL_AVX2__
   return _mm256_fnmadd_ps(a, b, c);
-#else
+#  else
   return c - (a * b);
-#endif
+#  endif
 }
 __forceinline const avxf msub(const avxf &a, const avxf &b, const avxf &c)
 {
-#ifdef __KERNEL_AVX2__
+#  ifdef __KERNEL_AVX2__
   return _mm256_fmsub_ps(a, b, c);
-#else
+#  else
   return (a * b) - c;
-#endif
+#  endif
 }
 
 ////////////////////////////////////////////////////////////////////////////////
-/// Comparison Operators + Select
+/// Comparison Operators
 ////////////////////////////////////////////////////////////////////////////////
 __forceinline const avxb operator<=(const avxf &a, const avxf &b)
 {
   return _mm256_cmp_ps(a.m256, b.m256, _CMP_LE_OS);
 }
 
-__forceinline const avxf select(const avxb &m, const avxf &t, const avxf &f)
-{
-  return _mm256_blendv_ps(f, t, m);
-}
-
-////////////////////////////////////////////////////////////////////////////////
-/// Common Functions
-////////////////////////////////////////////////////////////////////////////////
-
-__forceinline avxf mix(const avxf &a, const avxf &b, const avxf &t)
-{
-  return madd(t, b, (avxf(1.0f) - t) * a);
-}
+#endif
 
 #ifndef _mm256_set_m128
 #  define _mm256_set_m128(/* __m128 */ hi, /* __m128 */ lo) \
@@ -388,5 +352,3 @@ __forceinline avxf mix(const avxf &a, const avxf &b, const avxf &t)
   _mm256_set_m128(_mm_loadu_ps(hiaddr), _mm_loadu_ps(loaddr))
 
 CCL_NAMESPACE_END
-
-#endif

intern/cycles/util/util_avxi.h

@@ -1,745 +0,0 @@
-/*
- * Copyright 2009-2013 Intel Corporation
- *
- * 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 __UTIL_AVXI_H__
-#define __UTIL_AVXI_H__
-
-CCL_NAMESPACE_BEGIN
-
-struct avxb;
-
-struct avxi {
-  typedef avxb Mask;  // mask type for us
-  enum { size = 8 };  // number of SIMD elements
-  union {             // data
-    __m256i m256;
-#if !defined(__KERNEL_AVX2__)
-    struct {
-      __m128i l, h;
-    };
-#endif
-    int32_t v[8];
-  };
-
-  ////////////////////////////////////////////////////////////////////////////////
-  /// Constructors, Assignment & Cast Operators
-  ////////////////////////////////////////////////////////////////////////////////
-
-  __forceinline avxi()
-  {
-  }
-  __forceinline avxi(const avxi &a)
-  {
-    m256 = a.m256;
-  }
-  __forceinline avxi &operator=(const avxi &a)
-  {
-    m256 = a.m256;
-    return *this;
-  }
-
-  __forceinline avxi(const __m256i a) : m256(a)
-  {
-  }
-  __forceinline operator const __m256i &(void)const
-  {
-    return m256;
-  }
-  __forceinline operator __m256i &(void)
-  {
-    return m256;
-  }
-
-  __forceinline explicit avxi(const ssei &a)
-      : m256(_mm256_insertf128_si256(_mm256_castsi128_si256(a), a, 1))
-  {
-  }
-  __forceinline avxi(const ssei &a, const ssei &b)
-      : m256(_mm256_insertf128_si256(_mm256_castsi128_si256(a), b, 1))
-  {
-  }
-#if defined(__KERNEL_AVX2__)
-  __forceinline avxi(const __m128i &a, const __m128i &b)
-      : m256(_mm256_insertf128_si256(_mm256_castsi128_si256(a), b, 1))
-  {
-  }
-#else
-  __forceinline avxi(const __m128i &a, const __m128i &b) : l(a), h(b)
-  {
-  }
-#endif
-  __forceinline explicit avxi(const int32_t *const a)
-      : m256(_mm256_castps_si256(_mm256_loadu_ps((const float *)a)))
-  {
-  }
-  __forceinline avxi(int32_t a) : m256(_mm256_set1_epi32(a))
-  {
-  }
-  __forceinline avxi(int32_t a, int32_t b) : m256(_mm256_set_epi32(b, a, b, a, b, a, b, a))
-  {
-  }
-  __forceinline avxi(int32_t a, int32_t b, int32_t c, int32_t d)
-      : m256(_mm256_set_epi32(d, c, b, a, d, c, b, a))
-  {
-  }
-  __forceinline avxi(
-      int32_t a, int32_t b, int32_t c, int32_t d, int32_t e, int32_t f, int32_t g, int32_t h)
-      : m256(_mm256_set_epi32(h, g, f, e, d, c, b, a))
-  {
-  }
-
-  __forceinline explicit avxi(const __m256 a) : m256(_mm256_cvtps_epi32(a))
-  {
-  }
-
-  ////////////////////////////////////////////////////////////////////////////////
-  /// Constants
-  ////////////////////////////////////////////////////////////////////////////////
-
-  __forceinline avxi(ZeroTy) : m256(_mm256_setzero_si256())
-  {
-  }
-#if defined(__KERNEL_AVX2__)
-  __forceinline avxi(OneTy) : m256(_mm256_set1_epi32(1))
-  {
-  }
-  __forceinline avxi(PosInfTy) : m256(_mm256_set1_epi32(pos_inf))
-  {
-  }
-  __forceinline avxi(NegInfTy) : m256(_mm256_set1_epi32(neg_inf))
-  {
-  }
-#else
-  __forceinline avxi(OneTy) : m256(_mm256_set_epi32(1, 1, 1, 1, 1, 1, 1, 1))
-  {
-  }
-  __forceinline avxi(PosInfTy)
-      : m256(_mm256_set_epi32(
-            pos_inf, pos_inf, pos_inf, pos_inf, pos_inf, pos_inf, pos_inf, pos_inf))
-  {
-  }
-  __forceinline avxi(NegInfTy)
-      : m256(_mm256_set_epi32(
-            neg_inf, neg_inf, neg_inf, neg_inf, neg_inf, neg_inf, neg_inf, neg_inf))
-  {
-  }
-#endif
-  __forceinline avxi(StepTy) : m256(_mm256_set_epi32(7, 6, 5, 4, 3, 2, 1, 0))
-  {
-  }
-
-  ////////////////////////////////////////////////////////////////////////////////
-  /// Array Access
-  ////////////////////////////////////////////////////////////////////////////////
-
-  __forceinline const int32_t &operator[](const size_t i) const
-  {
-    assert(i < 8);
-    return v[i];
-  }
-  __forceinline int32_t &operator[](const size_t i)
-  {
-    assert(i < 8);
-    return v[i];
-  }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-/// Unary Operators
-////////////////////////////////////////////////////////////////////////////////
-
-__forceinline const avxi cast(const __m256 &a)
-{
-  return _mm256_castps_si256(a);
-}
-__forceinline const avxi operator+(const avxi &a)
-{
-  return a;
-}
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxi operator-(const avxi &a)
-{
-  return _mm256_sub_epi32(_mm256_setzero_si256(), a.m256);
-}
-__forceinline const avxi abs(const avxi &a)
-{
-  return _mm256_abs_epi32(a.m256);
-}
-#else
-__forceinline const avxi operator-(const avxi &a)
-{
-  return avxi(_mm_sub_epi32(_mm_setzero_si128(), a.l), _mm_sub_epi32(_mm_setzero_si128(), a.h));
-}
-__forceinline const avxi abs(const avxi &a)
-{
-  return avxi(_mm_abs_epi32(a.l), _mm_abs_epi32(a.h));
-}
-#endif
-
-////////////////////////////////////////////////////////////////////////////////
-/// Binary Operators
-////////////////////////////////////////////////////////////////////////////////
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxi operator+(const avxi &a, const avxi &b)
-{
-  return _mm256_add_epi32(a.m256, b.m256);
-}
-#else
-__forceinline const avxi operator+(const avxi &a, const avxi &b)
-{
-  return avxi(_mm_add_epi32(a.l, b.l), _mm_add_epi32(a.h, b.h));
-}
-#endif
-__forceinline const avxi operator+(const avxi &a, const int32_t b)
-{
-  return a + avxi(b);
-}
-__forceinline const avxi operator+(const int32_t a, const avxi &b)
-{
-  return avxi(a) + b;
-}
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxi operator-(const avxi &a, const avxi &b)
-{
-  return _mm256_sub_epi32(a.m256, b.m256);
-}
-#else
-__forceinline const avxi operator-(const avxi &a, const avxi &b)
-{
-  return avxi(_mm_sub_epi32(a.l, b.l), _mm_sub_epi32(a.h, b.h));
-}
-#endif
-__forceinline const avxi operator-(const avxi &a, const int32_t b)
-{
-  return a - avxi(b);
-}
-__forceinline const avxi operator-(const int32_t a, const avxi &b)
-{
-  return avxi(a) - b;
-}
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxi operator*(const avxi &a, const avxi &b)
-{
-  return _mm256_mullo_epi32(a.m256, b.m256);
-}
-#else
-__forceinline const avxi operator*(const avxi &a, const avxi &b)
-{
-  return avxi(_mm_mullo_epi32(a.l, b.l), _mm_mullo_epi32(a.h, b.h));
-}
-#endif
-__forceinline const avxi operator*(const avxi &a, const int32_t b)
-{
-  return a * avxi(b);
-}
-__forceinline const avxi operator*(const int32_t a, const avxi &b)
-{
-  return avxi(a) * b;
-}
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxi operator&(const avxi &a, const avxi &b)
-{
-  return _mm256_and_si256(a.m256, b.m256);
-}
-#else
-__forceinline const avxi operator&(const avxi &a, const avxi &b)
-{
-  return _mm256_castps_si256(_mm256_and_ps(_mm256_castsi256_ps(a), _mm256_castsi256_ps(b)));
-}
-#endif
-__forceinline const avxi operator&(const avxi &a, const int32_t b)
-{
-  return a & avxi(b);
-}
-__forceinline const avxi operator&(const int32_t a, const avxi &b)
-{
-  return avxi(a) & b;
-}
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxi operator|(const avxi &a, const avxi &b)
-{
-  return _mm256_or_si256(a.m256, b.m256);
-}
-#else
-__forceinline const avxi operator|(const avxi &a, const avxi &b)
-{
-  return _mm256_castps_si256(_mm256_or_ps(_mm256_castsi256_ps(a), _mm256_castsi256_ps(b)));
-}
-#endif
-__forceinline const avxi operator|(const avxi &a, const int32_t b)
-{
-  return a | avxi(b);
-}
-__forceinline const avxi operator|(const int32_t a, const avxi &b)
-{
-  return avxi(a) | b;
-}
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxi operator^(const avxi &a, const avxi &b)
-{
-  return _mm256_xor_si256(a.m256, b.m256);
-}
-#else
-__forceinline const avxi operator^(const avxi &a, const avxi &b)
-{
-  return _mm256_castps_si256(_mm256_xor_ps(_mm256_castsi256_ps(a), _mm256_castsi256_ps(b)));
-}
-#endif
-__forceinline const avxi operator^(const avxi &a, const int32_t b)
-{
-  return a ^ avxi(b);
-}
-__forceinline const avxi operator^(const int32_t a, const avxi &b)
-{
-  return avxi(a) ^ b;
-}
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxi operator<<(const avxi &a, const int32_t n)
-{
-  return _mm256_slli_epi32(a.m256, n);
-}
-__forceinline const avxi operator>>(const avxi &a, const int32_t n)
-{
-  return _mm256_srai_epi32(a.m256, n);
-}
-
-__forceinline const avxi sra(const avxi &a, const int32_t b)
-{
-  return _mm256_srai_epi32(a.m256, b);
-}
-__forceinline const avxi srl(const avxi &a, const int32_t b)
-{
-  return _mm256_srli_epi32(a.m256, b);
-}
-#else
-__forceinline const avxi operator<<(const avxi &a, const int32_t n)
-{
-  return avxi(_mm_slli_epi32(a.l, n), _mm_slli_epi32(a.h, n));
-}
-__forceinline const avxi operator>>(const avxi &a, const int32_t n)
-{
-  return avxi(_mm_srai_epi32(a.l, n), _mm_srai_epi32(a.h, n));
-}
-
-__forceinline const avxi sra(const avxi &a, const int32_t b)
-{
-  return avxi(_mm_srai_epi32(a.l, b), _mm_srai_epi32(a.h, b));
-}
-__forceinline const avxi srl(const avxi &a, const int32_t b)
-{
-  return avxi(_mm_srli_epi32(a.l, b), _mm_srli_epi32(a.h, b));
-}
-#endif
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxi min(const avxi &a, const avxi &b)
-{
-  return _mm256_min_epi32(a.m256, b.m256);
-}
-#else
-__forceinline const avxi min(const avxi &a, const avxi &b)
-{
-  return avxi(_mm_min_epi32(a.l, b.l), _mm_min_epi32(a.h, b.h));
-}
-#endif
-__forceinline const avxi min(const avxi &a, const int32_t b)
-{
-  return min(a, avxi(b));
-}
-__forceinline const avxi min(const int32_t a, const avxi &b)
-{
-  return min(avxi(a), b);
-}
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxi max(const avxi &a, const avxi &b)
-{
-  return _mm256_max_epi32(a.m256, b.m256);
-}
-#else
-__forceinline const avxi max(const avxi &a, const avxi &b)
-{
-  return avxi(_mm_max_epi32(a.l, b.l), _mm_max_epi32(a.h, b.h));
-}
-#endif
-__forceinline const avxi max(const avxi &a, const int32_t b)
-{
-  return max(a, avxi(b));
-}
-__forceinline const avxi max(const int32_t a, const avxi &b)
-{
-  return max(avxi(a), b);
-}
-
-////////////////////////////////////////////////////////////////////////////////
-/// Assignment Operators
-////////////////////////////////////////////////////////////////////////////////
-
-__forceinline avxi &operator+=(avxi &a, const avxi &b)
-{
-  return a = a + b;
-}
-__forceinline avxi &operator+=(avxi &a, const int32_t b)
-{
-  return a = a + b;
-}
-
-__forceinline avxi &operator-=(avxi &a, const avxi &b)
-{
-  return a = a - b;
-}
-__forceinline avxi &operator-=(avxi &a, const int32_t b)
-{
-  return a = a - b;
-}
-
-__forceinline avxi &operator*=(avxi &a, const avxi &b)
-{
-  return a = a * b;
-}
-__forceinline avxi &operator*=(avxi &a, const int32_t b)
-{
-  return a = a * b;
-}
-
-__forceinline avxi &operator&=(avxi &a, const avxi &b)
-{
-  return a = a & b;
-}
-__forceinline avxi &operator&=(avxi &a, const int32_t b)
-{
-  return a = a & b;
-}
-
-__forceinline avxi &operator|=(avxi &a, const avxi &b)
-{
-  return a = a | b;
-}
-__forceinline avxi &operator|=(avxi &a, const int32_t b)
-{
-  return a = a | b;
-}
-
-__forceinline avxi &operator^=(avxi &a, const avxi &b)
-{
-  return a = a ^ b;
-}
-__forceinline avxi &operator^=(avxi &a, const int32_t b)
-{
-  return a = a ^ b;
-}
-
-__forceinline avxi &operator<<=(avxi &a, const int32_t b)
-{
-  return a = a << b;
-}
-__forceinline avxi &operator>>=(avxi &a, const int32_t b)
-{
-  return a = a >> b;
-}
-
-////////////////////////////////////////////////////////////////////////////////
-/// Comparison Operators + Select
-////////////////////////////////////////////////////////////////////////////////
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxb operator==(const avxi &a, const avxi &b)
-{
-  return _mm256_castsi256_ps(_mm256_cmpeq_epi32(a.m256, b.m256));
-}
-#else
-__forceinline const avxb operator==(const avxi &a, const avxi &b)
-{
-  return avxb(_mm_castsi128_ps(_mm_cmpeq_epi32(a.l, b.l)),
-              _mm_castsi128_ps(_mm_cmpeq_epi32(a.h, b.h)));
-}
-#endif
-__forceinline const avxb operator==(const avxi &a, const int32_t b)
-{
-  return a == avxi(b);
-}
-__forceinline const avxb operator==(const int32_t a, const avxi &b)
-{
-  return avxi(a) == b;
-}
-
-__forceinline const avxb operator!=(const avxi &a, const avxi &b)
-{
-  return !(a == b);
-}
-__forceinline const avxb operator!=(const avxi &a, const int32_t b)
-{
-  return a != avxi(b);
-}
-__forceinline const avxb operator!=(const int32_t a, const avxi &b)
-{
-  return avxi(a) != b;
-}
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxb operator<(const avxi &a, const avxi &b)
-{
-  return _mm256_castsi256_ps(_mm256_cmpgt_epi32(b.m256, a.m256));
-}
-#else
-__forceinline const avxb operator<(const avxi &a, const avxi &b)
-{
-  return avxb(_mm_castsi128_ps(_mm_cmplt_epi32(a.l, b.l)),
-              _mm_castsi128_ps(_mm_cmplt_epi32(a.h, b.h)));
-}
-#endif
-__forceinline const avxb operator<(const avxi &a, const int32_t b)
-{
-  return a < avxi(b);
-}
-__forceinline const avxb operator<(const int32_t a, const avxi &b)
-{
-  return avxi(a) < b;
-}
-
-__forceinline const avxb operator>=(const avxi &a, const avxi &b)
-{
-  return !(a < b);
-}
-__forceinline const avxb operator>=(const avxi &a, const int32_t b)
-{
-  return a >= avxi(b);
-}
-__forceinline const avxb operator>=(const int32_t a, const avxi &b)
-{
-  return avxi(a) >= b;
-}
-
-#if defined(__KERNEL_AVX2__)
-__forceinline const avxb operator>(const avxi &a, const avxi &b)
-{
-  return _mm256_castsi256_ps(_mm256_cmpgt_epi32(a.m256, b.m256));
-}
-#else
-__forceinline const avxb operator>(const avxi &a, const avxi &b)
-{
-  return avxb(_mm_castsi128_ps(_mm_cmpgt_epi32(a.l, b.l)),
-              _mm_castsi128_ps(_mm_cmpgt_epi32(a.h, b.h)));
-}
-#endif
-__forceinline const avxb operator>(const avxi &a, const int32_t b)
-{
-  return a > avxi(b);
-}
-__forceinline const avxb operator>(const int32_t a, const avxi &b)
-{
-  return avxi(a) > b;
-}
-
-__forceinline const avxb operator<=(const avxi &a, const avxi &b)
-{
-  return !(a > b);
-}
-__forceinline const avxb operator<=(const avxi &a, const int32_t b)
-{
-  return a <= avxi(b);
-}
-__forceinline const avxb operator<=(const int32_t a, const avxi &b)
-{
-  return avxi(a) <= b;
-}
-
-__forceinline const avxi select(const avxb &m, const avxi &t, const avxi &f)
-{
-  return _mm256_castps_si256(_mm256_blendv_ps(_mm256_castsi256_ps(f), _mm256_castsi256_ps(t), m));
-}
-
-////////////////////////////////////////////////////////////////////////////////
-/// Movement/Shifting/Shuffling Functions
-////////////////////////////////////////////////////////////////////////////////
-
-#if defined(__KERNEL_AVX2__)
-__forceinline avxi unpacklo(const avxi &a, const avxi &b)
-{
-  return _mm256_unpacklo_epi32(a.m256, b.m256);
-}
-__forceinline avxi unpackhi(const avxi &a, const avxi &b)
-{
-  return _mm256_unpackhi_epi32(a.m256, b.m256);
-}
-#else
-__forceinline avxi unpacklo(const avxi &a, const avxi &b)
-{
-  return _mm256_castps_si256(_mm256_unpacklo_ps(_mm256_castsi256_ps(a), _mm256_castsi256_ps(b)));
-}
-__forceinline avxi unpackhi(const avxi &a, const avxi &b)
-{
-  return _mm256_castps_si256(_mm256_unpackhi_ps(_mm256_castsi256_ps(a), _mm256_castsi256_ps(b)));
-}
-#endif
-
-template<size_t i> __forceinline const avxi shuffle(const avxi &a)
-{
-  return _mm256_castps_si256(_mm256_permute_ps(_mm256_castsi256_ps(a), _MM_SHUFFLE(i, i, i, i)));
-}
-
-template<size_t i0, size_t i1> __forceinline const avxi shuffle(const avxi &a)
-{
-  return _mm256_permute2f128_si256(a, a, (i1 << 4) | (i0 << 0));
-}
-
-template<size_t i0, size_t i1> __forceinline const avxi shuffle(const avxi &a, const avxi &b)
-{
-  return _mm256_permute2f128_si256(a, b, (i1 << 4) | (i0 << 0));
-}
-
-template<size_t i0, size_t i1, size_t i2, size_t i3>
-__forceinline const avxi shuffle(const avxi &a)
-{
-  return _mm256_castps_si256(
-      _mm256_permute_ps(_mm256_castsi256_ps(a), _MM_SHUFFLE(i3, i2, i1, i0)));
-}
-
-template<size_t i0, size_t i1, size_t i2, size_t i3>
-__forceinline const avxi shuffle(const avxi &a, const avxi &b)
-{
-  return _mm256_castps_si256(_mm256_shuffle_ps(
-      _mm256_castsi256_ps(a), _mm256_castsi256_ps(b), _MM_SHUFFLE(i3, i2, i1, i0)));
-}
-
-template<> __forceinline const avxi shuffle<0, 0, 2, 2>(const avxi &b)
-{
-  return _mm256_castps_si256(_mm256_moveldup_ps(_mm256_castsi256_ps(b)));
-}
-template<> __forceinline const avxi shuffle<1, 1, 3, 3>(const avxi &b)
-{
-  return _mm256_castps_si256(_mm256_movehdup_ps(_mm256_castsi256_ps(b)));
-}
-template<> __forceinline const avxi shuffle<0, 1, 0, 1>(const avxi &b)
-{
-  return _mm256_castps_si256(
-      _mm256_castpd_ps(_mm256_movedup_pd(_mm256_castps_pd(_mm256_castsi256_ps(b)))));
-}
-
-__forceinline const avxi broadcast(const int *ptr)
-{
-  return _mm256_castps_si256(_mm256_broadcast_ss((const float *)ptr));
-}
-template<size_t i> __forceinline const avxi insert(const avxi &a, const ssei &b)
-{
-  return _mm256_insertf128_si256(a, b, i);
-}
-template<size_t i> __forceinline const ssei extract(const avxi &a)
-{
-  return _mm256_extractf128_si256(a, i);
-}
-
-////////////////////////////////////////////////////////////////////////////////
-/// Reductions
-////////////////////////////////////////////////////////////////////////////////
-
-__forceinline const avxi vreduce_min2(const avxi &v)
-{
-  return min(v, shuffle<1, 0, 3, 2>(v));
-}
-__forceinline const avxi vreduce_min4(const avxi &v)
-{
-  avxi v1 = vreduce_min2(v);
-  return min(v1, shuffle<2, 3, 0, 1>(v1));
-}
-__forceinline const avxi vreduce_min(const avxi &v)
-{
-  avxi v1 = vreduce_min4(v);
-  return min(v1, shuffle<1, 0>(v1));
-}
-
-__forceinline const avxi vreduce_max2(const avxi &v)
-{
-  return max(v, shuffle<1, 0, 3, 2>(v));
-}
-__forceinline const avxi vreduce_max4(const avxi &v)
-{
-  avxi v1 = vreduce_max2(v);
-  return max(v1, shuffle<2, 3, 0, 1>(v1));
-}
-__forceinline const avxi vreduce_max(const avxi &v)
-{
-  avxi v1 = vreduce_max4(v);
-  return max(v1, shuffle<1, 0>(v1));
-}
-
-__forceinline const avxi vreduce_add2(const avxi &v)
-{
-  return v + shuffle<1, 0, 3, 2>(v);
-}
-__forceinline const avxi vreduce_add4(const avxi &v)
-{
-  avxi v1 = vreduce_add2(v);
-  return v1 + shuffle<2, 3, 0, 1>(v1);
-}
-__forceinline const avxi vreduce_add(const avxi &v)
-{
-  avxi v1 = vreduce_add4(v);
-  return v1 + shuffle<1, 0>(v1);
-}
-
-__forceinline int reduce_min(const avxi &v)
-{
-  return extract<0>(extract<0>(vreduce_min(v)));
-}
-__forceinline int reduce_max(const avxi &v)
-{
-  return extract<0>(extract<0>(vreduce_max(v)));
-}
-__forceinline int reduce_add(const avxi &v)
-{
-  return extract<0>(extract<0>(vreduce_add(v)));
-}
-
-__forceinline size_t select_min(const avxi &v)
-{
-  return __bsf(movemask(v == vreduce_min(v)));
-}
-__forceinline size_t select_max(const avxi &v)
-{
-  return __bsf(movemask(v == vreduce_max(v)));
-}
-
-__forceinline size_t select_min(const avxb &valid, const avxi &v)
-{
-  const avxi a = select(valid, v, avxi(pos_inf));
-  return __bsf(movemask(valid & (a == vreduce_min(a))));
-}
-__forceinline size_t select_max(const avxb &valid, const avxi &v)
-{
-  const avxi a = select(valid, v, avxi(neg_inf));
-  return __bsf(movemask(valid & (a == vreduce_max(a))));
-}
-
-////////////////////////////////////////////////////////////////////////////////
-/// Output Operators
-////////////////////////////////////////////////////////////////////////////////
-
-ccl_device_inline void print_avxi(const char *label, const avxi &a)
-{
-  printf("%s: %d %d %d %d %d %d %d %d\n", label, a[0], a[1], a[2], a[3], a[4], a[5], a[6], a[7]);
-}
-
-CCL_NAMESPACE_END
-
-#endif

intern/cycles/util/util_hash.h

@@ -312,60 +312,6 @@ ccl_device_inline ssei hash_ssei4(ssei kx, ssei ky, ssei kz, ssei kw)
   return c;
 }
 
-#  if defined(__KERNEL_AVX__)
-ccl_device_inline avxi hash_avxi(avxi kx)
-{
-  avxi a, b, c;
-  a = b = c = avxi(0xdeadbeef + (1 << 2) + 13);
-
-  a += kx;
-  final(a, b, c);
-
-  return c;
-}
-
-ccl_device_inline avxi hash_avxi2(avxi kx, avxi ky)
-{
-  avxi a, b, c;
-  a = b = c = avxi(0xdeadbeef + (2 << 2) + 13);
-
-  b += ky;
-  a += kx;
-  final(a, b, c);
-
-  return c;
-}
-
-ccl_device_inline avxi hash_avxi3(avxi kx, avxi ky, avxi kz)
-{
-  avxi a, b, c;
-  a = b = c = avxi(0xdeadbeef + (3 << 2) + 13);
-
-  c += kz;
-  b += ky;
-  a += kx;
-  final(a, b, c);
-
-  return c;
-}
-
-ccl_device_inline avxi hash_avxi4(avxi kx, avxi ky, avxi kz, avxi kw)
-{
-  avxi a, b, c;
-  a = b = c = avxi(0xdeadbeef + (4 << 2) + 13);
-
-  a += kx;
-  b += ky;
-  c += kz;
-  mix(a, b, c);
-
-  a += kw;
-  final(a, b, c);
-
-  return c;
-}
-#  endif
-
 #  undef rot
 #  undef final
 #  undef mix

intern/cycles/util/util_semaphore.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 __UTIL_SEMAPHORE_H__
-#define __UTIL_SEMAPHORE_H__
-
-#include "util/util_thread.h"
-
-CCL_NAMESPACE_BEGIN
-
-/* Counting Semaphore
- *
- * To restrict concurrent access to a resource to a specified number
- * of threads. Similar to std::counting_semaphore from C++20. */
-
-class thread_counting_semaphore {
- public:
-  explicit thread_counting_semaphore(const int count) : count(count)
-  {
-  }
-
-  thread_counting_semaphore(const thread_counting_semaphore &) = delete;
-
-  void acquire()
-  {
-    thread_scoped_lock lock(mutex);
-    while (count == 0) {
-      condition.wait(lock);
-    }
-    count--;
-  }
-
-  void release()
-  {
-    thread_scoped_lock lock(mutex);
-    count++;
-    condition.notify_one();
-  }
-
- protected:
-  thread_mutex mutex;
-  thread_condition_variable condition;
-  int count;
-};
-
-CCL_NAMESPACE_END
-
-#endif /* __UTIL_SEMAPHORE_H__ */

intern/cycles/util/util_simd.h

@@ -75,28 +75,6 @@ static struct FalseTy {
   }
 } False ccl_maybe_unused;
 
-static struct ZeroTy {
-  __forceinline operator float() const
-  {
-    return 0;
-  }
-  __forceinline operator int() const
-  {
-    return 0;
-  }
-} zero ccl_maybe_unused;
-
-static struct OneTy {
-  __forceinline operator float() const
-  {
-    return 1;
-  }
-  __forceinline operator int() const
-  {
-    return 1;
-  }
-} one ccl_maybe_unused;
-
 static struct NegInfTy {
   __forceinline operator float() const
   {
@@ -119,9 +97,6 @@ static struct PosInfTy {
   }
 } inf ccl_maybe_unused, pos_inf ccl_maybe_unused;
 
-static struct StepTy {
-} step ccl_maybe_unused;
-
 /* Intrinsics Functions */
 
 #    if defined(__BMI__) && defined(__GNUC__)
@@ -588,13 +563,6 @@ __forceinline __m128 _mm_round_ps_emu(__m128 value, const int flags)
 
 #    endif /* !(defined(__KERNEL_SSE41__) || defined(__SSE4_1__) || defined(__SSE4_2__)) */
 
-/* Older GCC versions do not have _mm256_cvtss_f32 yet, so define it ourselves.
- * _mm256_castps256_ps128 generates no instructions so this is just as efficient. */
-#    if defined(__KERNEL_AVX__) || defined(__KERNEL_AVX2__)
-#      undef _mm256_cvtss_f32
-#      define _mm256_cvtss_f32(a) (_mm_cvtss_f32(_mm256_castps256_ps128(a)))
-#    endif
-
 #  else /* __KERNEL_SSE2__ */
 
 /* This section is for utility functions which operates on non-register data

intern/cycles/util/util_static_assert.h

@@ -24,9 +24,9 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if defined(__KERNEL_OPENCL__) || defined(CYCLES_CUBIN_CC)
+#ifdef __KERNEL_OPENCL__
 #  define static_assert(statement, message)
-#endif /* __KERNEL_OPENCL__ */
+#endif /* __KERNEL_GPU__ */
 
 #define static_assert_align(st, align) \
   static_assert((sizeof(st) % (align) == 0), "Structure must be strictly aligned")  // NOLINT

intern/cycles/util/util_types.h

@@ -158,7 +158,6 @@ CCL_NAMESPACE_END
 #  if defined(__KERNEL_AVX__) || defined(__KERNEL_AVX2__)
 #    include "util/util_avxb.h"
 #    include "util/util_avxf.h"
-#    include "util/util_avxi.h"
 #  endif
 #endif
 

intern/dualcon/intern/ModelReader.h

@@ -19,10 +19,6 @@
 
 #include "GeoCommon.h"
 
-#ifdef WITH_CXX_GUARDEDALLOC
-#  include "MEM_guardedalloc.h"
-#endif
-
 /*
  * Virtual class for input file readers
  *

intern/ghost/CMakeLists.txt

@@ -368,11 +368,11 @@ if(WITH_XR_OPENXR)
 
     GHOST_IXrContext.h
     intern/GHOST_IXrGraphicsBinding.h
+    intern/GHOST_Xr_intern.h
+    intern/GHOST_Xr_openxr_includes.h
     intern/GHOST_XrContext.h
     intern/GHOST_XrSession.h
     intern/GHOST_XrSwapchain.h
-    intern/GHOST_Xr_intern.h
-    intern/GHOST_Xr_openxr_includes.h
   )
   list(APPEND INC_SYS
     ${XR_OPENXR_SDK_INCLUDE_DIR}

intern/ghost/GHOST_C-api.h

@@ -214,6 +214,25 @@ extern GHOST_ContextHandle GHOST_CreateOpenGLContext(GHOST_SystemHandle systemha
 extern GHOST_TSuccess GHOST_DisposeOpenGLContext(GHOST_SystemHandle systemhandle,
                                                  GHOST_ContextHandle contexthandle);
 
+#ifdef WIN32
+/**
+ * Create a new offscreen context.
+ * Never explicitly delete the context, use disposeContext() instead.
+ * \param systemhandle The handle to the system
+ * \return A handle to the new context ( == NULL if creation failed).
+ */
+GHOST_ContextHandle GHOST_CreateDirectXContext(GHOST_SystemHandle systemhandle);
+
+/**
+ * Dispose of a context.
+ * \param systemhandle The handle to the system
+ * \param contexthandle Handle to the context to be disposed.
+ * \return Indication of success.
+ */
+GHOST_TSuccess GHOST_DisposeDirectXContext(GHOST_SystemHandle systemhandle,
+                                           GHOST_ContextHandle contexthandle);
+#endif
+
 /**
  * Returns the window user data.
  * \param windowhandle The handle to the window
@@ -738,7 +757,7 @@ extern GHOST_TSuccess GHOST_ActivateOpenGLContext(GHOST_ContextHandle contexthan
 extern GHOST_TSuccess GHOST_ReleaseOpenGLContext(GHOST_ContextHandle contexthandle);
 
 /**
- * Get the OpenGL frame-buffer handle that serves as a default frame-buffer.
+ * Get the OpenGL framebuffer handle that serves as a default framebuffer.
  */
 extern unsigned int GHOST_GetContextDefaultOpenGLFramebuffer(GHOST_ContextHandle contexthandle);
 
@@ -750,7 +769,7 @@ extern unsigned int GHOST_GetContextDefaultOpenGLFramebuffer(GHOST_ContextHandle
 extern int GHOST_isUpsideDownContext(GHOST_ContextHandle contexthandle);
 
 /**
- * Get the OpenGL frame-buffer handle that serves as a default frame-buffer.
+ * Get the OpenGL framebuffer handle that serves as a default framebuffer.
  */
 extern unsigned int GHOST_GetDefaultOpenGLFramebuffer(GHOST_WindowHandle windwHandle);
 
@@ -761,6 +780,13 @@ extern unsigned int GHOST_GetDefaultOpenGLFramebuffer(GHOST_WindowHandle windwHa
  */
 extern void GHOST_SetTabletAPI(GHOST_SystemHandle systemhandle, GHOST_TTabletAPI api);
 
+/**
+ * Returns the status of the tablet
+ * \param windowhandle The handle to the window
+ * \return Status of tablet
+ */
+extern const GHOST_TabletData *GHOST_GetTabletData(GHOST_WindowHandle windowhandle);
+
 /**
  * Access to rectangle width.
  * \param rectanglehandle The handle to the rectangle
@@ -1052,13 +1078,6 @@ void GHOST_XrSessionDrawViews(GHOST_XrContextHandle xr_context, void *customdata
  */
 int GHOST_XrSessionIsRunning(const GHOST_XrContextHandle xr_context);
 
-/**
- * Check if \a xr_context has a session that requires an upside-down frame-buffer (compared to
- * OpenGL). If true, the render result should be flipped vertically for correct output.
- * \note: Only to be called after session start, may otherwise result in a false negative.
- */
-int GHOST_XrSessionNeedsUpsideDownDrawing(const GHOST_XrContextHandle xr_context);
-
 /* events */
 /**
  * Invoke handling of all OpenXR events for \a xr_context. Should be called on every main-loop

intern/ghost/GHOST_IContext.h

@@ -60,6 +60,11 @@ class GHOST_IContext {
 
   virtual GHOST_TSuccess swapBuffers() = 0;
 
+  /**
+   * Returns if the window is rendered upside down compared to OpenGL.
+   */
+  virtual bool isUpsideDown() const = 0;
+
 #ifdef WITH_CXX_GUARDEDALLOC
   MEM_CXX_CLASS_ALLOC_FUNCS("GHOST:GHOST_IContext")
 #endif

intern/ghost/GHOST_ISystem.h

@@ -266,6 +266,12 @@ class GHOST_ISystem {
    */
   virtual GHOST_IContext *createOffscreenContext() = 0;
 
+  /**
+   * Overload to allow requesting a different context type. By default only OpenGL is supported.
+   * However by explicitly overloading this a system may add support for others.
+   */
+  virtual GHOST_IContext *createOffscreenContext(GHOST_TDrawingContextType type) = 0;
+
   /**
    * Dispose of a context.
    * \param   context Pointer to the context to be disposed.

intern/ghost/GHOST_IWindow.h

@@ -243,6 +243,16 @@ class GHOST_IWindow {
 
   virtual bool isDialog() const = 0;
 
+  /**
+   * Returns the tablet data (pressure etc).
+   * \return The tablet data (pressure etc).
+   */
+  virtual const GHOST_TabletData &GetTabletData()
+  {
+    /* Default state when no tablet is used, for systems without tablet support. */
+    return GHOST_TABLET_DATA_DEFAULT;
+  }
+
   /***************************************************************************************
    * Progress bar functionality
    ***************************************************************************************/

intern/ghost/GHOST_IXrContext.h

@@ -37,8 +37,6 @@ class GHOST_IXrContext {
   virtual void setGraphicsContextBindFuncs(GHOST_XrGraphicsContextBindFn bind_fn,
                                            GHOST_XrGraphicsContextUnbindFn unbind_fn) = 0;
   virtual void setDrawViewFunc(GHOST_XrDrawViewFn draw_view_fn) = 0;
-
-  virtual bool needsUpsideDownDrawing() const = 0;
 };
 
 #endif  // __GHOST_IXRCONTEXT_H__

intern/ghost/GHOST_Types.h

@@ -117,8 +117,8 @@ typedef struct GHOST_TabletData {
   float Ytilt; /* as above */
 } GHOST_TabletData;
 
-static const GHOST_TabletData GHOST_TABLET_DATA_NONE = {
-    GHOST_kTabletModeNone, /* No cursor in range */
+static const GHOST_TabletData GHOST_TABLET_DATA_DEFAULT = {
+    GHOST_kTabletModeNone, /* No tablet connected. */
     1.0f,                  /* Pressure */
     0.0f,                  /* Xtilt */
     0.0f};                 /* Ytilt */
@@ -598,8 +598,8 @@ typedef void (*GHOST_TimerProcPtr)(struct GHOST_TimerTaskHandle__ *task, GHOST_T
 
 #ifdef WITH_XR_OPENXR
 
-struct GHOST_XrDrawViewInfo;
 struct GHOST_XrError;
+struct GHOST_XrDrawViewInfo;
 /**
  * The XR view (i.e. the OpenXR runtime) may require a different graphics library than OpenGL. An
  * offscreen texture of the viewport will then be drawn into using OpenGL, but the final texture
@@ -621,8 +621,9 @@ typedef void (*GHOST_XrErrorHandlerFn)(const struct GHOST_XrError *);
 
 typedef void (*GHOST_XrSessionExitFn)(void *customdata);
 
-typedef void *(*GHOST_XrGraphicsContextBindFn)(void);
-typedef void (*GHOST_XrGraphicsContextUnbindFn)(GHOST_ContextHandle graphics_context);
+typedef void *(*GHOST_XrGraphicsContextBindFn)(enum GHOST_TXrGraphicsBinding graphics_lib);
+typedef void (*GHOST_XrGraphicsContextUnbindFn)(enum GHOST_TXrGraphicsBinding graphics_lib,
+                                                GHOST_ContextHandle graphics_context);
 typedef void (*GHOST_XrDrawViewFn)(const struct GHOST_XrDrawViewInfo *draw_view, void *customdata);
 
 /* An array of GHOST_TXrGraphicsBinding items defining the candidate bindings to use. The first

intern/ghost/intern/GHOST_C-api.cpp

@@ -150,6 +150,25 @@ GHOST_TSuccess GHOST_DisposeOpenGLContext(GHOST_SystemHandle systemhandle,
   return system->disposeContext(context);
 }
 
+#ifdef WIN32
+GHOST_ContextHandle GHOST_CreateDirectXContext(GHOST_SystemHandle systemhandle)
+{
+  GHOST_ISystem *system = (GHOST_ISystem *)systemhandle;
+
+  return (GHOST_ContextHandle)system->createOffscreenContext(GHOST_kDrawingContextTypeD3D);
+}
+
+GHOST_TSuccess GHOST_DisposeDirectXContext(GHOST_SystemHandle systemhandle,
+                                           GHOST_ContextHandle contexthandle)
+{
+  GHOST_ISystem *system = (GHOST_ISystem *)systemhandle;
+  GHOST_IContext *context = (GHOST_IContext *)contexthandle;
+
+  return system->disposeContext(context);
+}
+
+#endif
+
 GHOST_WindowHandle GHOST_CreateWindow(GHOST_SystemHandle systemhandle,
                                       const char *title,
                                       GHOST_TInt32 left,
@@ -697,6 +716,13 @@ unsigned int GHOST_GetContextDefaultOpenGLFramebuffer(GHOST_ContextHandle contex
   return context->getDefaultFramebuffer();
 }
 
+int GHOST_isUpsideDownContext(GHOST_ContextHandle contexthandle)
+{
+  GHOST_IContext *context = (GHOST_IContext *)contexthandle;
+
+  return context->isUpsideDown();
+}
+
 unsigned int GHOST_GetDefaultOpenGLFramebuffer(GHOST_WindowHandle windowhandle)
 {
   GHOST_IWindow *window = (GHOST_IWindow *)windowhandle;
@@ -717,6 +743,11 @@ void GHOST_SetTabletAPI(GHOST_SystemHandle systemhandle, GHOST_TTabletAPI api)
   system->setTabletAPI(api);
 }
 
+const GHOST_TabletData *GHOST_GetTabletData(GHOST_WindowHandle windowhandle)
+{
+  return &((GHOST_IWindow *)windowhandle)->GetTabletData();
+}
+
 GHOST_TInt32 GHOST_GetWidthRectangle(GHOST_RectangleHandle rectanglehandle)
 {
   return ((GHOST_Rect *)rectanglehandle)->getWidth();
@@ -960,12 +991,4 @@ void GHOST_XrDrawViewFunc(GHOST_XrContextHandle xr_contexthandle, GHOST_XrDrawVi
   GHOST_XR_CAPI_CALL(xr_context->setDrawViewFunc(draw_view_fn), xr_context);
 }
 
-int GHOST_XrSessionNeedsUpsideDownDrawing(const GHOST_XrContextHandle xr_contexthandle)
-{
-  GHOST_IXrContext *xr_context = (GHOST_IXrContext *)xr_contexthandle;
-
-  GHOST_XR_CAPI_CALL_RET(xr_context->needsUpsideDownDrawing(), xr_context);
-  return 0; /* Only reached if exception is thrown. */
-}
-
 #endif

intern/ghost/intern/GHOST_Context.h

@@ -120,9 +120,9 @@ class GHOST_Context : public GHOST_IContext {
   }
 
   /**
-   * Returns if the context is rendered upside down compared to OpenGL.
+   * Returns if the window is rendered upside down compared to OpenGL.
    */
-  virtual inline bool isUpsideDown() const
+  inline bool isUpsideDown() const
   {
     return false;
   }

intern/ghost/intern/GHOST_ContextCGL.mm

@@ -18,7 +18,7 @@
  */
 
 /** \file
- * \ingroup GHOST
+ *  \ingroup GHOST
  *
  * Definition of GHOST_ContextCGL class.
  */

intern/ghost/intern/GHOST_EventButton.h

@@ -37,19 +37,17 @@ class GHOST_EventButton : public GHOST_Event {
    * Constructor.
    * \param time      The time this event was generated.
    * \param type      The type of this event.
-   * \param window    The window of this event.
-   * \param button    The state of the buttons were at the time of the event.
-   * \param tablet    The tablet data associated with this event.
+   * \param window: The window of this event.
+   * \param button: The state of the buttons were at the time of the event.
    */
   GHOST_EventButton(GHOST_TUns64 time,
                     GHOST_TEventType type,
                     GHOST_IWindow *window,
-                    GHOST_TButtonMask button,
-                    const GHOST_TabletData &tablet)
+                    GHOST_TButtonMask button)
       : GHOST_Event(time, type, window)
   {
     m_buttonEventData.button = button;
-    m_buttonEventData.tablet = tablet;
+    m_buttonEventData.tablet = window->GetTabletData();
     m_data = &m_buttonEventData;
   }
 

intern/ghost/intern/GHOST_EventCursor.h

@@ -38,19 +38,17 @@ class GHOST_EventCursor : public GHOST_Event {
    * \param type      The type of this event.
    * \param x         The x-coordinate of the location the cursor was at the time of the event.
    * \param y         The y-coordinate of the location the cursor was at the time of the event.
-   * \param tablet    The tablet data associated with this event.
    */
   GHOST_EventCursor(GHOST_TUns64 msec,
                     GHOST_TEventType type,
                     GHOST_IWindow *window,
                     GHOST_TInt32 x,
-                    GHOST_TInt32 y,
-                    const GHOST_TabletData &tablet)
+                    GHOST_TInt32 y)
       : GHOST_Event(msec, type, window)
   {
     m_cursorEventData.x = x;
     m_cursorEventData.y = y;
-    m_cursorEventData.tablet = tablet;
+    m_cursorEventData.tablet = window->GetTabletData();
     m_data = &m_cursorEventData;
   }
 

intern/ghost/intern/GHOST_IXrGraphicsBinding.h

@@ -61,7 +61,6 @@ class GHOST_IXrGraphicsBinding {
       uint32_t image_count) = 0;
   virtual void submitToSwapchainImage(XrSwapchainImageBaseHeader *swapchain_image,
                                       const GHOST_XrDrawViewInfo *draw_info) = 0;
-  virtual bool needsUpsideDownDrawing(GHOST_Context &ghost_ctx) const = 0;
 
  protected:
   /* Use GHOST_XrGraphicsBindingCreateFromType! */
@@ -69,6 +68,6 @@ class GHOST_IXrGraphicsBinding {
 };
 
 std::unique_ptr<GHOST_IXrGraphicsBinding> GHOST_XrGraphicsBindingCreateFromType(
-    GHOST_TXrGraphicsBinding type, GHOST_Context *ghost_ctx);
+    GHOST_TXrGraphicsBinding type);
 
 #endif /* __GHOST_IXRGRAPHICSBINDING_H__ */