Fix mistakes with background drawing

Played with a more complex and more efficient version, but that made
things pretty complicated. Now this is a simple, but less optimized, solution.

CMakeLists.txt

@@ -497,11 +497,10 @@ endif()
 # We default options to whatever default standard in the current compiler.
 if(CMAKE_COMPILER_IS_GNUCC AND (NOT "${CMAKE_C_COMPILER_VERSION}" VERSION_LESS "6.0") AND (NOT WITH_CXX11))
 	set(_c11_init ON)
-	set(_cxx11_init ON)
 else()
 	set(_c11_init OFF)
-	set(_cxx11_init OFF)
 endif()
+set(_cxx11_init ON)
 
 option(WITH_C11 "Build with C11 standard enabled, for development use only!" ${_c11_init})
 mark_as_advanced(WITH_C11)
@@ -628,6 +627,12 @@ if(APPLE)
 			# to silence sdk not found warning, just overrides CMAKE_OSX_SYSROOT
 			set(CMAKE_XCODE_ATTRIBUTE_SDKROOT macosx${OSX_SYSTEM})
 		endif()
+
+		# QuickTime framework is no longer available in SDK 10.12+
+		if(WITH_CODEC_QUICKTIME AND ${OSX_SYSTEM} VERSION_GREATER 10.11)
+			set(WITH_CODEC_QUICKTIME OFF)
+			message(STATUS "QuickTime not supported by SDK ${OSX_SYSTEM}, disabling WITH_CODEC_QUICKTIME")
+		endif()
 	endif()
 
 	if(OSX_SYSTEM MATCHES 10.9)
@@ -861,7 +866,7 @@ endif()
 # linux only, not cached
 set(WITH_BINRELOC OFF)
 
-# MAXOSX only, set to avoid uninitialized
+# MACOSX only, set to avoid uninitialized
 set(EXETYPE "")
 
 # C/C++ flags
@@ -1567,7 +1572,7 @@ if(WITH_CXX11)
 	if(CMAKE_COMPILER_IS_GNUCC OR CMAKE_C_COMPILER_ID MATCHES "Clang")
 		# TODO(sergey): Do we want c++11 or gnu-c++11 here?
 		set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11")
-	elseif(MSVC12)
+	elseif(MSVC)
 		# Nothing special is needed, C++11 features are available by default.
 	else()
 		message(FATAL_ERROR "Compiler ${CMAKE_C_COMPILER_ID} is not supported for C++11 build yet")

build_files/build_environment/install_deps.sh

@@ -284,7 +284,7 @@ NO_BUILD=false
 NO_CONFIRM=false
 USE_CXX11=true  # Mandatory in blender2.8
 
-PYTHON_VERSION="3.5.1"
+PYTHON_VERSION="3.5.2"
 PYTHON_VERSION_MIN="3.5"
 PYTHON_FORCE_BUILD=false
 PYTHON_FORCE_REBUILD=false
@@ -317,7 +317,7 @@ OPENEXR_FORCE_REBUILD=false
 OPENEXR_SKIP=false
 _with_built_openexr=false
 
-OIIO_VERSION="1.6.9"
+OIIO_VERSION="1.7.8"
 OIIO_VERSION_MIN="1.6.0"
 OIIO_VERSION_MAX="1.9.0"  # UNKNOWN currently # Not supported by current OSL...
 OIIO_FORCE_BUILD=false
@@ -332,14 +332,14 @@ LLVM_FORCE_REBUILD=false
 LLVM_SKIP=false
 
 # OSL needs to be compiled for now!
-OSL_VERSION="1.7.3"
+OSL_VERSION="1.7.5"
 OSL_VERSION_MIN=$OSL_VERSION
 OSL_FORCE_BUILD=false
 OSL_FORCE_REBUILD=false
 OSL_SKIP=false
 
 # OpenSubdiv needs to be compiled for now
-OSD_VERSION="3.0.5"
+OSD_VERSION="3.1.1"
 OSD_VERSION_MIN=$OSD_VERSION
 OSD_FORCE_BUILD=false
 OSD_FORCE_REBUILD=false
@@ -367,7 +367,7 @@ OPENCOLLADA_FORCE_BUILD=false
 OPENCOLLADA_FORCE_REBUILD=false
 OPENCOLLADA_SKIP=false
 
-FFMPEG_VERSION="2.8.4"
+FFMPEG_VERSION="3.2.1"
 FFMPEG_VERSION_MIN="2.8.4"
 FFMPEG_FORCE_BUILD=false
 FFMPEG_FORCE_REBUILD=false
@@ -705,6 +705,21 @@ if [ "$WITH_ALL" = true -a "$OPENCOLLADA_SKIP" = false ]; then
 fi
 
 
+WARNING "****WARNING****"
+PRINT "If you are experiencing issues building Blender, _*TRY A FRESH, CLEAN BUILD FIRST*_!"
+PRINT "The same goes for install_deps itself, if you encounter issues, please first erase everything in $SRC and $INST"
+PRINT "(provided obviously you did not add anything yourself in those dirs!), and run install_deps.sh again!"
+PRINT "Often, changes in the libs built by this script, or in your distro package, cannot be handled simply, so..."
+PRINT ""
+PRINT "You may also try to use the '--build-foo' options to bypass your distribution's packages"
+PRINT "for some troublesome/buggy libraries..."
+PRINT ""
+PRINT ""
+PRINT "Ran with:"
+PRINT "    install_deps.sh $COMMANDLINE"
+PRINT ""
+PRINT ""
+
 
 # This has to be done here, because user might force some versions...
 PYTHON_SOURCE=( "https://www.python.org/ftp/python/$PYTHON_VERSION/Python-$PYTHON_VERSION.tgz" )
@@ -777,6 +792,8 @@ However, if you are experiencing linking errors (also when building Blender itse
 
 Please note that until the transition to C++11-built libraries if completed in your distribution, situation will
 remain fuzzy and incompatibilities may happen..."
+  PRINT ""
+  PRINT ""
   CXXFLAGS="$CXXFLAGS -std=c++11"
   export CXXFLAGS
 fi
@@ -2455,7 +2472,7 @@ compile_FFmpeg() {
         --enable-avfilter --disable-vdpau \
         --disable-bzlib --disable-libgsm --disable-libspeex \
         --enable-pthreads --enable-zlib --enable-stripping --enable-runtime-cpudetect \
-        --disable-vaapi --disable-libfaac --disable-nonfree --enable-gpl \
+        --disable-vaapi --disable-nonfree --enable-gpl \
         --disable-postproc --disable-librtmp --disable-libopencore-amrnb \
         --disable-libopencore-amrwb --disable-libdc1394 --disable-version3 --disable-outdev=sdl \
         --disable-libxcb \
@@ -4154,16 +4171,6 @@ print_info_ffmpeglink() {
 }
 
 print_info() {
-  PRINT ""
-  PRINT ""
-  WARNING "****WARNING****"
-  PRINT "If you are experiencing issues building Blender, _*TRY A FRESH, CLEAN BUILD FIRST*_!"
-  PRINT "The same goes for install_deps itself, if you encounter issues, please first erase everything in $SRC and $INST"
-  PRINT "(provided obviously you did not add anything yourself in those dirs!), and run install_deps.sh again!"
-  PRINT "Often, changes in the libs built by this script, or in your distro package, cannot be handled simply, so..."
-  PRINT ""
-  PRINT "You may also try to use the '--build-foo' options to bypass your distribution's packages"
-  PRINT "for some troublesome/buggy libraries..."
   PRINT ""
   PRINT ""
   PRINT "Ran with:"

build_files/buildbot/master.cfg

@@ -297,8 +297,8 @@ def generic_builder(id, libdir='', branch='', rsync=False):
 # Builders
 
 add_builder(c, 'mac_x86_64_10_6_cmake', 'darwin-9.x.universal', generic_builder, hour=5)
-add_builder(c, 'linux_glibc211_i686_cmake', '', generic_builder, hour=1)
-add_builder(c, 'linux_glibc211_x86_64_cmake', '', generic_builder, hour=2)
+# add_builder(c, 'linux_glibc211_i686_cmake', '', generic_builder, hour=1)
+# add_builder(c, 'linux_glibc211_x86_64_cmake', '', generic_builder, hour=2)
 add_builder(c, 'linux_glibc219_i686_cmake', '', generic_builder, hour=3)
 add_builder(c, 'linux_glibc219_x86_64_cmake', '', generic_builder, hour=4)
 add_builder(c, 'win32_cmake_vc2013', 'windows_vc12', generic_builder, hour=3)

build_files/cmake/platform/platform_apple.cmake

@@ -158,7 +158,7 @@ if(WITH_CODEC_FFMPEG)
 		mp3lame swscale x264 xvidcore theora theoradec theoraenc vorbis vorbisenc vorbisfile ogg
 	)
 	if(WITH_CXX11)
-		set(FFMPEG_LIBRARIES ${FFMPEG_LIBRARIES} schroedinger orc vpx)
+		set(FFMPEG_LIBRARIES ${FFMPEG_LIBRARIES} schroedinger orc vpx webp swresample)
 	endif()
 	set(FFMPEG_LIBPATH ${FFMPEG}/lib)
 endif()
@@ -316,6 +316,9 @@ if(WITH_OPENIMAGEIO)
 		${OPENEXR_LIBRARIES}
 		${ZLIB_LIBRARIES}
 	)
+	if(WITH_CXX11)
+		set(OPENIMAGEIO_LIBRARIES ${OPENIMAGEIO_LIBRARIES} ${LIBDIR}/ffmpeg/lib/libwebp.a)
+	endif()
 	set(OPENIMAGEIO_LIBPATH
 		${OPENIMAGEIO}/lib
 		${JPEG_LIBPATH}

build_files/cmake/platform/platform_win32_msvc.cmake

@@ -112,7 +112,7 @@ set(CMAKE_STATIC_LINKER_FLAGS "${CMAKE_STATIC_LINKER_FLAGS} /ignore:4221")
 
 # MSVC only, Mingw doesnt need
 if(CMAKE_CL_64)
-	set(PLATFORM_LINKFLAGS "/MACHINE:X64 /OPT:NOREF ${PLATFORM_LINKFLAGS}")
+	set(PLATFORM_LINKFLAGS "/MACHINE:X64 ${PLATFORM_LINKFLAGS}")
 else()
 	set(PLATFORM_LINKFLAGS "/MACHINE:IX86 /LARGEADDRESSAWARE ${PLATFORM_LINKFLAGS}")
 endif()
@@ -238,14 +238,14 @@ if(WITH_CODEC_FFMPEG)
 	windows_find_package(FFMPEG)
 	if(NOT FFMPEG_FOUND)
 		warn_hardcoded_paths(ffmpeg)
-		set(FFMPEG_LIBRARY_VERSION 55)
-		set(FFMPEG_LIBRARY_VERSION_AVU 52)
+		set(FFMPEG_LIBRARY_VERSION 57)
+		set(FFMPEG_LIBRARY_VERSION_AVU 55)
 		set(FFMPEG_LIBRARIES
-			${LIBDIR}/ffmpeg/lib/avcodec-${FFMPEG_LIBRARY_VERSION}.lib
-			${LIBDIR}/ffmpeg/lib/avformat-${FFMPEG_LIBRARY_VERSION}.lib
-			${LIBDIR}/ffmpeg/lib/avdevice-${FFMPEG_LIBRARY_VERSION}.lib
-			${LIBDIR}/ffmpeg/lib/avutil-${FFMPEG_LIBRARY_VERSION_AVU}.lib
-			${LIBDIR}/ffmpeg/lib/swscale-2.lib
+			${LIBDIR}/ffmpeg/lib/avcodec.lib
+			${LIBDIR}/ffmpeg/lib/avformat.lib
+			${LIBDIR}/ffmpeg/lib/avdevice.lib
+			${LIBDIR}/ffmpeg/lib/avutil.lib
+			${LIBDIR}/ffmpeg/lib/swscale.lib
 			)
 	endif()
 endif()
@@ -380,6 +380,7 @@ if(WITH_OPENIMAGEIO)
 	set(OPENCOLORIO_DEFINITIONS "-DOCIO_STATIC_BUILD")
 	set(OPENIMAGEIO_IDIFF "${OPENIMAGEIO}/bin/idiff.exe")
 	add_definitions(-DOIIO_STATIC_BUILD)
+	add_definitions(-DOIIO_NO_SSE=1)
 endif()
 
 if(WITH_LLVM)

doc/python_api/rst/bge_types/bge.types.KX_GameObject.rst

@@ -405,7 +405,7 @@ base class --- :class:`SCA_IObject`
 
       .. note::
 
-         This attribute is experemental and may be removed (but probably wont be).
+         This attribute is experimental and may be removed (but probably wont be).
 
       .. note::
 
@@ -419,7 +419,7 @@ base class --- :class:`SCA_IObject`
 
       .. note::
 
-         This attribute is experemental and may be removed (but probably wont be).
+         This attribute is experimental and may be removed (but probably wont be).
 
       .. note::
 
@@ -453,7 +453,7 @@ base class --- :class:`SCA_IObject`
 
    .. attribute:: childrenRecursive
 
-      all children of this object including childrens children, (read-only).
+      all children of this object including children's children, (read-only).
 
       :type: :class:`CListValue` of :class:`KX_GameObject`'s
 
@@ -536,7 +536,7 @@ base class --- :class:`SCA_IObject`
 
    .. method:: getAxisVect(vect)
 
-      Returns the axis vector rotates by the objects worldspace orientation.
+      Returns the axis vector rotates by the object's worldspace orientation.
       This is the equivalent of multiplying the vector by the orientation matrix.
 
       :arg vect: a vector to align the axis.
@@ -596,7 +596,7 @@ base class --- :class:`SCA_IObject`
 
       Gets the game object's linear velocity.
 
-      This method returns the game object's velocity through it's centre of mass, ie no angular velocity component.
+      This method returns the game object's velocity through it's center of mass, ie no angular velocity component.
 
       :arg local:
          * False: you get the "global" velocity ie: relative to world orientation.
@@ -609,7 +609,7 @@ base class --- :class:`SCA_IObject`
 
       Sets the game object's linear velocity.
 
-      This method sets game object's velocity through it's centre of mass,
+      This method sets game object's velocity through it's center of mass,
       ie no angular velocity component.
 
       This requires a dynamic object.
@@ -814,7 +814,7 @@ base class --- :class:`SCA_IObject`
             # do something
             pass
 
-      The face paremeter determines the orientation of the normal.
+      The face parameter determines the orientation of the normal.
 
       * 0 => hit normal is always oriented towards the ray origin (as if you casted the ray from outside)
       * 1 => hit normal is the real face normal (only for mesh object, otherwise face has no effect)
@@ -911,7 +911,7 @@ base class --- :class:`SCA_IObject`
 
       .. note::
 
-         The gameObject argument has an advantage that it can convert from a mesh with modifiers applied (such as subsurf).
+         The gameObject argument has an advantage that it can convert from a mesh with modifiers applied (such as the Subdivision Surface modifier).
 
       .. warning::
 
@@ -919,7 +919,7 @@ base class --- :class:`SCA_IObject`
 
       .. warning::
 
-         If the object is a part of a combound object it will fail (parent or child)
+         If the object is a part of a compound object it will fail (parent or child)
 
       .. warning::
 

doc/python_api/rst/bgl.rst

@@ -12,7 +12,7 @@ contents: dir(bgl).  A simple search on the web can point to more
 than enough material to teach OpenGL programming, from books to many
 collections of tutorials.
 
-Here is a comprehensive `list of books <https://www.opengl.org/documentation/books/>`__ (non free).
+Here is a comprehensive `list of books <https://www.khronos.org/developers/books/>`__ (non free).
 The `arcsynthesis tutorials <https://web.archive.org/web/20150225192611/http://www.arcsynthesis.org/gltut/index.html>`__
 is one of the best resources to learn modern OpenGL and
 `g-truc <http://www.g-truc.net/post-opengl-samples.html#menu>`__
@@ -2067,7 +2067,7 @@ offers a set of extensive examples, including advanced features.
    :arg length: Returns the length of the string returned in source (excluding the null terminator).
    :type source: :class:`bgl.Buffer` char.
    :arg source: Specifies an array of characters that is used to return the source code string.
-   
+
 
 .. function:: glShaderSource(shader, shader_string):
 

doc/python_api/rst/info_api_reference.rst

@@ -204,7 +204,7 @@ Lets say we want to access the texture of a brush via Python, to adjust its ``co
 
 - Start in the default scene and enable 'Sculpt' mode from the 3D-View header.
 - From the toolbar expand the **Texture** panel and add a new texture.
-  *Notice the texture button its self doesn't have very useful links (you can check the tool-tips).*
+  *Notice the texture button its self doesn't have very useful links (you can check the tooltips).*
 - The contrast setting isn't exposed in the sculpt toolbar, so view the texture in the properties panel...
 
   - In the properties button select the Texture context.

doc/python_api/rst/info_overview.rst

@@ -19,7 +19,7 @@ This is a typical Python environment so tutorials on how to write Python scripts
 will work running the scripts in Blender too.
 Blender provides the :mod:`bpy` module to the Python interpreter.
 This module can be imported in a script and gives access to Blender data, classes, and functions.
-Scripts that deal with Blender data will need to import this module. 
+Scripts that deal with Blender data will need to import this module.
 
 Here is a simple example of moving a vertex of the object named **Cube**:
 
@@ -80,7 +80,7 @@ To run as modules:
 
 
 Add-ons
-------
+-------
 
 Some of Blenders functionality is best kept optional,
 alongside scripts loaded at startup we have add-ons which are kept in their own directory ``scripts/addons``,
@@ -213,7 +213,7 @@ A simple Blender/Python module can look like this:
        bpy.utils.register_class(SimpleOperator)
 
    def unregister():
-       bpy.utils.unregister_class(SimpleOperator)    
+       bpy.utils.unregister_class(SimpleOperator)
 
    if __name__ == "__main__":
        register()
@@ -327,7 +327,7 @@ Say you want to store material settings for a custom engine.
 .. note::
 
    *The class must be registered before being used in a property, failing to do so will raise an error:*
-   
+
    ``ValueError: bpy_struct "Material" registration error: my_custom_props could not register``
 
 
@@ -429,4 +429,3 @@ Calling these operators:
    >>> bpy.ops.object.operator_2()
    Hello World OBJECT_OT_operator_2
    {'FINISHED'}
-

doc/python_api/sphinx_doc_gen.py

@@ -1565,9 +1565,9 @@ def pyrna2sphinx(basepath):
 
     # operators
     def write_ops():
-        API_BASEURL = "https://developer.blender.org/diffusion/B/browse/master/release/scripts/ "
-        API_BASEURL_ADDON = "https://developer.blender.org/diffusion/BA/"
-        API_BASEURL_ADDON_CONTRIB = "https://developer.blender.org/diffusion/BAC/"
+        API_BASEURL = "https://developer.blender.org/diffusion/B/browse/master/release/scripts "
+        API_BASEURL_ADDON = "https://developer.blender.org/diffusion/BA"
+        API_BASEURL_ADDON_CONTRIB = "https://developer.blender.org/diffusion/BAC"
 
         op_modules = {}
         for op in ops.values():
@@ -1632,13 +1632,9 @@ def write_sphinx_conf_py(basepath):
     file = open(filepath, "w", encoding="utf-8")
     fw = file.write
 
-    fw("import sys, os\n")
-    fw("\n")
-    fw("extensions = ['sphinx.ext.intersphinx']\n")
-    fw("\n")
-    fw("intersphinx_mapping = {'blender_manual': ('https://www.blender.org/manual/', None)}\n")
-    fw("\n")
-
+    fw("import sys, os\n\n")
+    fw("extensions = ['sphinx.ext.intersphinx']\n\n")
+    fw("intersphinx_mapping = {'blender_manual': ('https://docs.blender.org/manual/en/dev/', None)}\n\n")
     fw("project = 'Blender'\n")
     # fw("master_doc = 'index'\n")
     fw("copyright = u'Blender Foundation'\n")
@@ -1655,12 +1651,16 @@ def write_sphinx_conf_py(basepath):
 
     # not helpful since the source is generated, adds to upload size.
     fw("html_copy_source = False\n")
+    fw("html_show_sphinx = False\n")
     fw("html_split_index = True\n")
     fw("\n")
 
     # needed for latex, pdf gen
+    fw("latex_elements = {\n")
+    fw("  'papersize': 'a4paper',\n")
+    fw("}\n\n")
+
     fw("latex_documents = [ ('contents', 'contents.tex', 'Blender Index', 'Blender Foundation', 'manual'), ]\n")
-    fw("latex_paper_size = 'a4paper'\n")
     file.close()
 
 

doc/python_api/sphinx_doc_update.py

@@ -41,9 +41,9 @@ import tempfile
 import zipfile
 
 
-DEFAULT_RSYNC_SERVER = "www.blender.org"
+DEFAULT_RSYNC_SERVER = "docs.blender.org"
 DEFAULT_RSYNC_ROOT = "/api/"
-DEFAULT_SYMLINK_ROOT = "/data/www/vhosts/www.blender.org/api"
+DEFAULT_SYMLINK_ROOT = "/data/www/vhosts/docs.blender.org/api"
 
 
 def argparse_create():

extern/Eigen3/Eigen/src/StlSupport/StdVector.h

@@ -77,7 +77,7 @@ namespace std {
   void resize(size_type new_size)
   { resize(new_size, T()); }
 
-#if defined(_VECTOR_)
+#if defined(_VECTOR_) && (_MSC_VER<1910)
   // workaround MSVC std::vector implementation
   void resize(size_type new_size, const value_type& x)
   {
@@ -110,7 +110,7 @@ namespace std {
       vector_base::insert(vector_base::end(), new_size - vector_base::size(), x);
   }
 #else
-  // either GCC 4.1 or non-GCC
+  // either GCC 4.1, MSVC2017 or non-GCC
   // default implementation which should always work.
   void resize(size_type new_size, const value_type& x)
   {

extern/clew/include/clew.h

@@ -369,7 +369,8 @@ typedef unsigned int cl_GLenum;
 #endif
 
 /* Define basic vector types */
-#if defined( __VEC__ )
+/* WOrkaround for ppc64el platform: conflicts with bool from C++. */
+#if defined( __VEC__ ) && !(defined(__PPC64__) && defined(__LITTLE_ENDIAN__))
    #include <altivec.h>   /* may be omitted depending on compiler. AltiVec spec provides no way to detect whether the header is required. */
    typedef vector unsigned char     __cl_uchar16;
    typedef vector signed char       __cl_char16;

extern/curve_fit_nd/curve_fit_nd.h

@@ -137,7 +137,7 @@ int curve_fit_cubic_to_points_refit_db(
         const double          error_threshold,
         const unsigned int    calc_flag,
         const unsigned int   *corners,
-        unsigned int          corners_len,
+        const unsigned int    corners_len,
         const double          corner_angle,
 
         double **r_cubic_array, unsigned int *r_cubic_array_len,

extern/gflags/README.blender

@@ -18,6 +18,8 @@ Local modifications:
 - Applied some modifications from fork https://github.com/Nazg-Gul/gflags.git
   (see https://github.com/gflags/gflags/pull/129)
 
-- Avoid attemot of acquiring mutex lock in FlagRegistry::GlobalRegistry when
+- Avoid attempt of acquiring mutex lock in FlagRegistry::GlobalRegistry when
   doing static flags initialization. See d81dd2d in Blender repository.
 
+- Made `google::{anonymous}::FlagValue::ValueSize() const` inlined, so it does
+  not trigger strict compiler warning.

extern/gflags/src/gflags.cc

@@ -218,7 +218,7 @@ class FlagValue {
   bool Equal(const FlagValue& x) const;
   FlagValue* New() const;   // creates a new one with default value
   void CopyFrom(const FlagValue& x);
-  int ValueSize() const;
+  inline int ValueSize() const;
 
   // Calls the given validate-fn on value_buffer_, and returns
   // whatever it returns.  But first casts validate_fn_proto to a
@@ -443,7 +443,7 @@ void FlagValue::CopyFrom(const FlagValue& x) {
   }
 }
 
-int FlagValue::ValueSize() const {
+inline int FlagValue::ValueSize() const {
   if (type_ > FV_MAX_INDEX) {
     assert(false);  // unknown type
     return 0;

extern/glog/src/raw_logging.cc

@@ -59,7 +59,8 @@
 # include <unistd.h>
 #endif
 
-#if defined(HAVE_SYSCALL_H) || defined(HAVE_SYS_SYSCALL_H)
+// Hurd does not have SYS_write.
+#if (defined(HAVE_SYSCALL_H) || defined(HAVE_SYS_SYSCALL_H)) && !defined(__GNU__)
 # define safe_write(fd, s, len)  syscall(SYS_write, fd, s, len)
 #else
   // Not so safe, but what can you do?

extern/glog/src/stacktrace_powerpc-inl.h

@@ -111,7 +111,7 @@ int GetStackTrace(void** result, int max_depth, int skip_count) {
       result[n++] = *(sp+2);
 #elif defined(_CALL_SYSV)
       result[n++] = *(sp+1);
-#elif defined(__APPLE__) || (defined(__linux) && defined(__PPC64__))
+#elif defined(__APPLE__) || ((defined(__linux) || defined(__linux__)) && defined(__PPC64__))
       // This check is in case the compiler doesn't define _CALL_AIX/etc.
       result[n++] = *(sp+2);
 #elif defined(__linux)

extern/gtest/include/gtest/internal/gtest-internal.h

@@ -60,6 +60,10 @@
 #include <string>
 #include <vector>
 
+#if (__cplusplus > 199711L) || (defined(_MSC_VER) && _MSC_VER >= 1800)
+# include <type_traits>
+#endif
+
 #include "gtest/gtest-message.h"
 #include "gtest/internal/gtest-string.h"
 #include "gtest/internal/gtest-filepath.h"
@@ -854,6 +858,7 @@ struct AddReference<T&> { typedef T& type; };  // NOLINT
 template <typename From, typename To>
 class ImplicitlyConvertible {
  private:
+#if !((__cplusplus > 199711L) || (defined(_MSC_VER) && _MSC_VER >= 1800))
   // We need the following helper functions only for their types.
   // They have no implementations.
 
@@ -874,6 +879,7 @@ class ImplicitlyConvertible {
   // implicitly converted to type To.
   static char Helper(To);
   static char (&Helper(...))[2];  // NOLINT
+#endif
 
   // We have to put the 'public' section after the 'private' section,
   // or MSVC refuses to compile the code.
@@ -883,6 +889,8 @@ class ImplicitlyConvertible {
   // instantiation.  The simplest workaround is to use its C++0x type traits
   // functions (C++Builder 2009 and above only).
   static const bool value = __is_convertible(From, To);
+#elif (__cplusplus > 199711L) || (defined(_MSC_VER) && _MSC_VER >= 1800)
+  static const bool value = std::is_convertible<From, To>::value;
 #else
   // MSVC warns about implicitly converting from double to int for
   // possible loss of data, so we need to temporarily disable the

extern/rangetree/intern/range_tree.c

@@ -808,7 +808,7 @@ bool range_tree_uint_retake(RangeTreeUInt *rt, const uint value)
 
 uint range_tree_uint_take_any(RangeTreeUInt *rt)
 {
-	Node *node = node = rt->list.first;
+	Node *node = rt->list.first;
 	uint value = node->min;
 	if (value == node->max) {
 		rt_node_remove(rt, node);

intern/audaspace/FX/AUD_LimiterReader.cpp

@@ -110,10 +110,10 @@ void AUD_LimiterReader::read(int& length, bool& eos, sample_t* buffer)
 			eos = true;
 		}
 
-		if(position < m_start * rate)
+		if(position < int(m_start * rate))
 		{
 			int len2 = length;
-			for(int len = m_start * rate - position;
+			for(int len = int(m_start * rate) - position;
 				len2 == length && !eos;
 				len -= length)
 			{

intern/cycles/CMakeLists.txt

@@ -74,7 +74,6 @@ elseif(CMAKE_COMPILER_IS_GNUCC)
 	if(CXX_HAS_AVX2)
 		set(CYCLES_AVX2_KERNEL_FLAGS "-ffast-math -msse -msse2 -msse3 -mssse3 -msse4.1 -mavx -mavx2 -mfma -mlzcnt -mbmi -mbmi2 -mf16c -mfpmath=sse")
 	endif()
-	set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -ffast-math")
 elseif(CMAKE_CXX_COMPILER_ID MATCHES "Clang")
 	check_cxx_compiler_flag(-msse CXX_HAS_SSE)
 	check_cxx_compiler_flag(-mavx CXX_HAS_AVX)
@@ -90,7 +89,6 @@ elseif(CMAKE_CXX_COMPILER_ID MATCHES "Clang")
 	if(CXX_HAS_AVX2)
 		set(CYCLES_AVX2_KERNEL_FLAGS "-ffast-math -msse -msse2 -msse3 -mssse3 -msse4.1 -mavx -mavx2 -mfma -mlzcnt -mbmi -mbmi2 -mf16c")
 	endif()
-	set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -ffast-math")
 endif()
 
 if(CXX_HAS_SSE)

intern/cycles/app/cycles_standalone.cpp

@@ -72,20 +72,17 @@ static void session_print(const string& str)
 
 static void session_print_status()
 {
-	int sample, tile;
-	double total_time, sample_time, render_time;
 	string status, substatus;
 
 	/* get status */
-	sample = options.session->progress.get_sample();
-	options.session->progress.get_tile(tile, total_time, sample_time, render_time);
+	float progress = options.session->progress.get_progress();
 	options.session->progress.get_status(status, substatus);
 
 	if(substatus != "")
 		status += ": " + substatus;
 
 	/* print status */
-	status = string_printf("Sample %d   %s", sample, status.c_str());
+	status = string_printf("Progress %05.2f   %s", (double) progress*100, status.c_str());
 	session_print(status);
 }
 
@@ -167,13 +164,12 @@ static void display_info(Progress& progress)
 	latency = (elapsed - last);
 	last = elapsed;
 
-	int sample, tile;
-	double total_time, sample_time, render_time;
+	double total_time, sample_time;
 	string status, substatus;
 
-	sample = progress.get_sample();
-	progress.get_tile(tile, total_time, sample_time, render_time);
+	progress.get_time(total_time, sample_time);
 	progress.get_status(status, substatus);
+	float progress_val = progress.get_progress();
 
 	if(substatus != "")
 		status += ": " + substatus;
@@ -184,10 +180,10 @@ static void display_info(Progress& progress)
 	        "%s"
 	        "        Time: %.2f"
 	        "        Latency: %.4f"
-	        "        Sample: %d"
+	        "        Progress: %05.2f"
 	        "        Average: %.4f"
 	        "        Interactive: %s",
-	        status.c_str(), total_time, latency, sample, sample_time, interactive.c_str());
+	        status.c_str(), total_time, latency, (double) progress_val*100, sample_time, interactive.c_str());
 
 	view_display_info(str.c_str());
 

intern/cycles/app/cycles_xml.cpp

@@ -523,7 +523,7 @@ static void xml_read_mesh(const XMLReadState& state, pugi::xml_node node)
 
 	/* we don't yet support arbitrary attributes, for now add vertex
 	 * coordinates as generated coordinates if requested */
-	if (mesh->need_attribute(state.scene, ATTR_STD_GENERATED)) {
+	if(mesh->need_attribute(state.scene, ATTR_STD_GENERATED)) {
 		Attribute *attr = mesh->attributes.add(ATTR_STD_GENERATED);
 		memcpy(attr->data_float3(), mesh->verts.data(), sizeof(float3)*mesh->verts.size());
 	}

intern/cycles/blender/CMakeLists.txt

@@ -25,6 +25,8 @@ set(SRC
 	blender_camera.cpp
 	blender_mesh.cpp
 	blender_object.cpp
+	blender_object_cull.cpp
+	blender_particles.cpp
 	blender_curves.cpp
 	blender_logging.cpp
 	blender_python.cpp
@@ -34,6 +36,7 @@ set(SRC
 	blender_texture.cpp
 
 	CCL_api.h
+	blender_object_cull.h
 	blender_sync.h
 	blender_session.h
 	blender_texture.h

intern/cycles/blender/addon/__init__.py

@@ -23,11 +23,25 @@ bl_info = {
     "location": "Info header, render engine menu",
     "description": "Cycles Render Engine integration",
     "warning": "",
-    "wiki_url": "https://www.blender.org/manual/render/cycles/index.html",
+    "wiki_url": "https://docs.blender.org/manual/en/dev/render/cycles/",
     "tracker_url": "",
     "support": 'OFFICIAL',
     "category": "Render"}
 
+# Support 'reload' case.
+if "bpy" in locals():
+    import importlib
+    if "engine" in locals():
+        importlib.reload(engine)
+    if "version_update" in locals():
+        importlib.reload(version_update)
+    if "ui" in locals():
+        importlib.reload(ui)
+    if "properties" in locals():
+        importlib.reload(properties)
+    if "presets" in locals():
+        importlib.reload(presets)
+
 import bpy
 
 from . import (

intern/cycles/blender/addon/properties.py

@@ -30,7 +30,7 @@ import _cycles
 
 enum_devices = (
     ('CPU', "CPU", "Use CPU for rendering"),
-    ('GPU', "GPU Compute", "Use GPU compute device for rendering, configured in user preferences"),
+    ('GPU', "GPU Compute", "Use GPU compute device for rendering, configured in the system tab in the user preferences"),
     )
 
 if _cycles.with_network:
@@ -129,6 +129,16 @@ enum_device_type = (
     ('OPENCL', "OpenCL", "OpenCL", 2)
     )
 
+enum_texture_limit = (
+    ('OFF', "No Limit", "No texture size limit", 0),
+    ('128', "128", "Limit texture size to 128 pixels", 1),
+    ('256', "256", "Limit texture size to 256 pixels", 2),
+    ('512', "512", "Limit texture size to 512 pixels", 3),
+    ('1024', "1024", "Limit texture size to 1024 pixels", 4),
+    ('2048', "2048", "Limit texture size to 2048 pixels", 5),
+    ('4096', "4096", "Limit texture size to 4096 pixels", 6),
+    ('8192', "8192", "Limit texture size to 8192 pixels", 7),
+    )
 
 class CyclesRenderSettings(bpy.types.PropertyGroup):
     @classmethod
@@ -278,7 +288,7 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
                 description="Probabilistically terminate light samples when the light contribution is below this threshold (more noise but faster rendering). "
                             "Zero disables the test and never ignores lights",
                 min=0.0, max=1.0,
-                default=0.05,
+                default=0.01,
                 )
 
         cls.caustics_reflective = BoolProperty(
@@ -518,6 +528,12 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
                 description="Use special type BVH optimized for hair (uses more ram but renders faster)",
                 default=True,
                 )
+        cls.debug_bvh_time_steps = IntProperty(
+                name="BVH Time Steps",
+                description="Split BVH primitives by this number of time steps to speed up render time in cost of memory",
+                default=0,
+                min=0, max=16,
+                )
         cls.tile_order = EnumProperty(
                 name="Tile Order",
                 description="Tile order for rendering",
@@ -566,6 +582,19 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
                 min=0.0, max=5.0
                 )
 
+        cls.use_distance_cull = BoolProperty(
+                name="Use Distance Cull",
+                description="Allow objects to be culled based on the distance from camera",
+                default=False,
+                )
+
+        cls.distance_cull_margin = FloatProperty(
+                name="Cull Distance",
+                description="Cull objects which are further away from camera than this distance",
+                default=50,
+                min=0.0
+                )
+
         cls.motion_blur_position = EnumProperty(
             name="Motion Blur Position",
             default='CENTER',
@@ -595,6 +624,20 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
             min=0.0, max=1.0,
             )
 
+        cls.texture_limit = EnumProperty(
+            name="Viewport Texture Limit",
+            default='OFF',
+            description="Limit texture size used by viewport rendering",
+            items=enum_texture_limit
+            )
+
+        cls.texture_limit_render = EnumProperty(
+            name="Render Texture Limit",
+            default='OFF',
+            description="Limit texture size used by final rendering",
+            items=enum_texture_limit
+            )
+
         # Various fine-tuning debug flags
 
         def devices_update_callback(self, context):
@@ -1016,6 +1059,12 @@ class CyclesObjectSettings(bpy.types.PropertyGroup):
                 default=False,
                 )
 
+        cls.use_distance_cull = BoolProperty(
+                name="Use Distance Cull",
+                description="Allow this object and its duplicators to be culled by distance from camera",
+                default=False,
+                )
+
         cls.use_adaptive_subdivision = BoolProperty(
                 name="Use Adaptive Subdivision",
                 description="Use adaptive render time subdivision",
@@ -1094,6 +1143,49 @@ class CyclesCurveRenderSettings(bpy.types.PropertyGroup):
         del bpy.types.Scene.cycles_curves
 
 
+class CyclesCurveSettings(bpy.types.PropertyGroup):
+    @classmethod
+    def register(cls):
+        bpy.types.ParticleSettings.cycles = PointerProperty(
+                name="Cycles Hair Settings",
+                description="Cycles hair settings",
+                type=cls,
+                )
+        cls.radius_scale = FloatProperty(
+                name="Radius Scaling",
+                description="Multiplier of width properties",
+                min=0.0, max=1000.0,
+                default=0.01,
+                )
+        cls.root_width = FloatProperty(
+                name="Root Size",
+                description="Strand's width at root",
+                min=0.0, max=1000.0,
+                default=1.0,
+                )
+        cls.tip_width = FloatProperty(
+                name="Tip Multiplier",
+                description="Strand's width at tip",
+                min=0.0, max=1000.0,
+                default=0.0,
+                )
+        cls.shape = FloatProperty(
+                name="Strand Shape",
+                description="Strand shape parameter",
+                min=-1.0, max=1.0,
+                default=0.0,
+                )
+        cls.use_closetip = BoolProperty(
+                name="Close tip",
+                description="Set tip radius to zero",
+                default=True,
+                )
+
+    @classmethod
+    def unregister(cls):
+        del bpy.types.ParticleSettings.cycles
+
+
 class CyclesDeviceSettings(bpy.types.PropertyGroup):
     @classmethod
     def register(cls):
@@ -1126,7 +1218,7 @@ class CyclesPreferences(bpy.types.AddonPreferences):
 
     def get_devices(self):
         import _cycles
-        # Layout of the device tuples: (Name, Type, Internal ID, Persistent ID)
+        # Layout of the device tuples: (Name, Type, Persistent ID)
         device_list = _cycles.available_devices()
 
         cuda_devices = []
@@ -1174,21 +1266,19 @@ class CyclesPreferences(bpy.types.AddonPreferences):
 
 
     def draw_impl(self, layout, context):
-        layout.label(text="Compute Device:")
+        layout.label(text="Cycles Compute Device:")
         layout.row().prop(self, "compute_device_type", expand=True)
 
         cuda_devices, opencl_devices = self.get_devices()
         row = layout.row()
 
-        if cuda_devices:
+        if self.compute_device_type == 'CUDA' and cuda_devices:
             col = row.column(align=True)
-            col.label(text="CUDA devices:")
             for device in cuda_devices:
                 col.prop(device, "use", text=device.name, toggle=True)
 
-        if opencl_devices:
+        if self.compute_device_type == 'OPENCL' and opencl_devices:
             col = row.column(align=True)
-            col.label(text="OpenCL devices:")
             for device in opencl_devices:
                 col.prop(device, "use", text=device.name, toggle=True)
 
@@ -1207,6 +1297,7 @@ def register():
     bpy.utils.register_class(CyclesMeshSettings)
     bpy.utils.register_class(CyclesObjectSettings)
     bpy.utils.register_class(CyclesCurveRenderSettings)
+    bpy.utils.register_class(CyclesCurveSettings)
     bpy.utils.register_class(CyclesDeviceSettings)
     bpy.utils.register_class(CyclesPreferences)
 
@@ -1221,5 +1312,6 @@ def unregister():
     bpy.utils.unregister_class(CyclesObjectSettings)
     bpy.utils.unregister_class(CyclesVisibilitySettings)
     bpy.utils.unregister_class(CyclesCurveRenderSettings)
+    bpy.utils.unregister_class(CyclesCurveSettings)
     bpy.utils.unregister_class(CyclesDeviceSettings)
     bpy.utils.unregister_class(CyclesPreferences)

intern/cycles/blender/addon/ui.py

@@ -217,7 +217,7 @@ class CyclesRender_PT_sampling(CyclesButtonsPanel, Panel):
         draw_samples_info(layout, context)
 
 
-class CyclesRender_PT_geometery(CyclesButtonsPanel, Panel):
+class CyclesRender_PT_geometry(CyclesButtonsPanel, Panel):
     bl_label = "Geometry"
     bl_options = {'DEFAULT_CLOSED'}
 
@@ -226,6 +226,7 @@ class CyclesRender_PT_geometery(CyclesButtonsPanel, Panel):
 
         scene = context.scene
         cscene = scene.cycles
+        ccscene = scene.cycles_curves
 
         if cscene.feature_set == 'EXPERIMENTAL':
             split = layout.split()
@@ -252,6 +253,25 @@ class CyclesRender_PT_geometery(CyclesButtonsPanel, Panel):
             row.prop(cscene, "volume_step_size")
             row.prop(cscene, "volume_max_steps")
 
+        layout.prop(ccscene, "use_curves", text="Use Hair")
+        col = layout.column()
+        col.active = ccscene.use_curves
+
+        col.prop(ccscene, "primitive", text="Primitive")
+        col.prop(ccscene, "shape", text="Shape")
+
+        if not (ccscene.primitive in {'CURVE_SEGMENTS', 'LINE_SEGMENTS'} and ccscene.shape == 'RIBBONS'):
+            col.prop(ccscene, "cull_backfacing", text="Cull back-faces")
+
+        if ccscene.primitive == 'TRIANGLES' and ccscene.shape == 'THICK':
+            col.prop(ccscene, "resolution", text="Resolution")
+        elif ccscene.primitive == 'CURVE_SEGMENTS':
+            col.prop(ccscene, "subdivisions", text="Curve subdivisions")
+
+        row = col.row()
+        row.prop(ccscene, "minimum_width", text="Min Pixels")
+        row.prop(ccscene, "maximum_width", text="Max Ext.")
+
 
 class CyclesRender_PT_light_paths(CyclesButtonsPanel, Panel):
     bl_label = "Light Paths"
@@ -412,6 +432,10 @@ class CyclesRender_PT_performance(CyclesButtonsPanel, Panel):
         col.prop(cscene, "debug_use_spatial_splits")
         col.prop(cscene, "debug_use_hair_bvh")
 
+        row = col.row()
+        row.active = not cscene.debug_use_spatial_splits
+        row.prop(cscene, "debug_bvh_time_steps")
+
 
 class CyclesRender_PT_layer_options(CyclesButtonsPanel, Panel):
     bl_label = "Layer"
@@ -767,8 +791,13 @@ class CyclesObject_PT_cycles_settings(CyclesButtonsPanel, Panel):
         col = layout.column()
         col.label(text="Performance:")
         row = col.row()
-        row.active = scene.render.use_simplify and cscene.use_camera_cull
-        row.prop(cob, "use_camera_cull")
+        sub = row.row()
+        sub.active = scene.render.use_simplify and cscene.use_camera_cull
+        sub.prop(cob, "use_camera_cull")
+
+        sub = row.row()
+        sub.active = scene.render.use_simplify and cscene.use_distance_cull
+        sub.prop(cob, "use_distance_cull")
 
 
 class CYCLES_OT_use_shading_nodes(Operator):
@@ -1358,41 +1387,35 @@ class CyclesTexture_PT_colors(CyclesButtonsPanel, Panel):
             layout.template_color_ramp(mapping, "color_ramp", expand=True)
 
 
-class CyclesRender_PT_CurveRendering(CyclesButtonsPanel, Panel):
-    bl_label = "Cycles Hair Rendering"
+class CyclesParticle_PT_textures(CyclesButtonsPanel, Panel):
+    bl_label = "Textures"
     bl_context = "particle"
+    bl_options = {'DEFAULT_CLOSED'}
 
     @classmethod
     def poll(cls, context):
         psys = context.particle_system
-        return CyclesButtonsPanel.poll(context) and psys and psys.settings.type == 'HAIR'
-
-    def draw_header(self, context):
-        ccscene = context.scene.cycles_curves
-        self.layout.prop(ccscene, "use_curves", text="")
+        return psys and CyclesButtonsPanel.poll(context)
 
     def draw(self, context):
         layout = self.layout
 
-        scene = context.scene
-        ccscene = scene.cycles_curves
-
-        layout.active = ccscene.use_curves
-
-        layout.prop(ccscene, "primitive", text="Primitive")
-        layout.prop(ccscene, "shape", text="Shape")
-
-        if not (ccscene.primitive in {'CURVE_SEGMENTS', 'LINE_SEGMENTS'} and ccscene.shape == 'RIBBONS'):
-            layout.prop(ccscene, "cull_backfacing", text="Cull back-faces")
-
-        if ccscene.primitive == 'TRIANGLES' and ccscene.shape == 'THICK':
-            layout.prop(ccscene, "resolution", text="Resolution")
-        elif ccscene.primitive == 'CURVE_SEGMENTS':
-            layout.prop(ccscene, "subdivisions", text="Curve subdivisions")
+        psys = context.particle_system
+        part = psys.settings
 
         row = layout.row()
-        row.prop(ccscene, "minimum_width", text="Min Pixels")
-        row.prop(ccscene, "maximum_width", text="Max Ext.")
+        row.template_list("TEXTURE_UL_texslots", "", part, "texture_slots", part, "active_texture_index", rows=2)
+
+        col = row.column(align=True)
+        col.operator("texture.slot_move", text="", icon='TRIA_UP').type = 'UP'
+        col.operator("texture.slot_move", text="", icon='TRIA_DOWN').type = 'DOWN'
+        col.menu("TEXTURE_MT_specials", icon='DOWNARROW_HLT', text="")
+
+        if not part.active_texture:
+            layout.template_ID(part, "active_texture", new="texture.new")
+        else:
+            slot = part.texture_slots[part.active_texture_index]
+            layout.template_ID(slot, "texture", new="texture.new")
 
 
 class CyclesRender_PT_bake(CyclesButtonsPanel, Panel):
@@ -1506,6 +1529,37 @@ class CyclesRender_PT_debug(CyclesButtonsPanel, Panel):
         col.prop(cscene, "debug_use_opencl_debug", text="Debug")
 
 
+class CyclesParticle_PT_CurveSettings(CyclesButtonsPanel, Panel):
+    bl_label = "Cycles Hair Settings"
+    bl_context = "particle"
+
+    @classmethod
+    def poll(cls, context):
+        scene = context.scene
+        ccscene = scene.cycles_curves
+        psys = context.particle_system
+        use_curves = ccscene.use_curves and psys
+        return CyclesButtonsPanel.poll(context) and use_curves and psys.settings.type == 'HAIR'
+
+    def draw(self, context):
+        layout = self.layout
+
+        psys = context.particle_settings
+        cpsys = psys.cycles
+
+        row = layout.row()
+        row.prop(cpsys, "shape", text="Shape")
+
+        layout.label(text="Thickness:")
+        row = layout.row()
+        row.prop(cpsys, "root_width", text="Root")
+        row.prop(cpsys, "tip_width", text="Tip")
+
+        row = layout.row()
+        row.prop(cpsys, "radius_scale", text="Scaling")
+        row.prop(cpsys, "use_closetip", text="Close tip")
+
+
 class CyclesScene_PT_simplify(CyclesButtonsPanel, Panel):
     bl_label = "Simplify"
     bl_context = "scene"
@@ -1523,22 +1577,41 @@ class CyclesScene_PT_simplify(CyclesButtonsPanel, Panel):
         cscene = scene.cycles
 
         layout.active = rd.use_simplify
+
+        col = layout.column(align=True)
+        col.label(text="Subdivision")
+        row = col.row(align=True)
+        row.prop(rd, "simplify_subdivision", text="Viewport")
+        row.prop(rd, "simplify_subdivision_render", text="Render")
+
+
+        col = layout.column(align=True)
+        col.label(text="Child Particles")
+        row = col.row(align=True)
+        row.prop(rd, "simplify_child_particles", text="Viewport")
+        row.prop(rd, "simplify_child_particles_render", text="Render")
+
+        col = layout.column(align=True)
+        split = col.split()
+        sub = split.column()
+        sub.label(text="Texture Limit Viewport")
+        sub.prop(cscene, "texture_limit", text="")
+        sub = split.column()
+        sub.label(text="Texture Limit Render")
+        sub.prop(cscene, "texture_limit_render", text="")
+
         split = layout.split()
-
         col = split.column()
-        col.label(text="Viewport:")
-        col.prop(rd, "simplify_subdivision", text="Subdivision")
-
-        col = split.column()
-        col.label(text="Render:")
-        col.prop(rd, "simplify_subdivision_render", text="Subdivision")
-
-        col = layout.column()
         col.prop(cscene, "use_camera_cull")
-        subsub = col.column()
-        subsub.active = cscene.use_camera_cull
-        subsub.prop(cscene, "camera_cull_margin")
+        row = col.row()
+        row.active = cscene.use_camera_cull
+        row.prop(cscene, "camera_cull_margin")
 
+        col = split.column()
+        col.prop(cscene, "use_distance_cull")
+        row = col.row()
+        row.active = cscene.use_distance_cull
+        row.prop(cscene, "distance_cull_margin", text="Distance")
 
 def draw_device(self, context):
     scene = context.scene
@@ -1552,11 +1625,9 @@ def draw_device(self, context):
 
         split = layout.split(percentage=1/3)
         split.label("Device:")
-        row = split.row(align=True)
-        sub = row.split(align=True)
-        sub.active = show_device_selection(context)
-        sub.prop(cscene, "device", text="")
-        row.operator("wm.addon_userpref_show", text="", icon='PREFERENCES').module = __package__
+        row = split.row()
+        row.active = show_device_selection(context)
+        row.prop(cscene, "device", text="")
 
         if engine.with_osl() and use_cpu(context):
             layout.prop(cscene, "shading_system")

intern/cycles/blender/addon/version_update.py

@@ -172,6 +172,24 @@ def custom_bake_remap(scene):
 
 @persistent
 def do_versions(self):
+    if bpy.context.user_preferences.version <= (2, 78, 1):
+        prop = bpy.context.user_preferences.addons[__package__].preferences
+        system = bpy.context.user_preferences.system
+        if not prop.is_property_set("compute_device_type"):
+            # Device might not currently be available so this can fail
+            try:
+                if system.legacy_compute_device_type == 1:
+                    prop.compute_device_type = 'OPENCL'
+                elif system.legacy_compute_device_type == 2:
+                    prop.compute_device_type = 'CUDA'
+                else:
+                    prop.compute_device_type = 'NONE'
+            except:
+                pass
+
+            # Init device list for UI
+            prop.get_devices()
+
     # We don't modify startup file because it assumes to
     # have all the default values only.
     if not bpy.data.is_saved:

intern/cycles/blender/blender_curves.cpp

@@ -29,21 +29,6 @@
 
 CCL_NAMESPACE_BEGIN
 
-/* Utilities */
-
-/* Hair curve functions */
-
-void curveinterp_v3_v3v3v3v3(float3 *p, float3 *v1, float3 *v2, float3 *v3, float3 *v4, const float w[4]);
-void interp_weights(float t, float data[4]);
-float shaperadius(float shape, float root, float tip, float time);
-void InterpolateKeySegments(int seg, int segno, int key, int curve, float3 *keyloc, float *time, ParticleCurveData *CData);
-void ExportCurveSegments(Scene *scene, Mesh *mesh, ParticleCurveData *CData);
-void ExportCurveTrianglePlanes(Mesh *mesh, ParticleCurveData *CData,
-                               float3 RotCam, bool is_ortho);
-void ExportCurveTriangleGeometry(Mesh *mesh, ParticleCurveData *CData, int resolution);
-void ExportCurveTriangleUV(ParticleCurveData *CData, int vert_offset, int resol, float3 *uvdata);
-void ExportCurveTriangleVcol(ParticleCurveData *CData, int vert_offset, int resol, uchar4 *cdata);
-
 ParticleCurveData::ParticleCurveData()
 {
 }
@@ -52,7 +37,7 @@ ParticleCurveData::~ParticleCurveData()
 {
 }
 
-void interp_weights(float t, float data[4])
+static void interp_weights(float t, float data[4])
 {
 	/* Cardinal curve interpolation */
 	float t2 = t * t;
@@ -65,17 +50,19 @@ void interp_weights(float t, float data[4])
 	data[3] =  fc          * t3  - fc * t2;
 }
 
-void curveinterp_v3_v3v3v3v3(float3 *p, float3 *v1, float3 *v2, float3 *v3, float3 *v4, const float w[4])
+static void curveinterp_v3_v3v3v3v3(float3 *p,
+                                    float3 *v1, float3 *v2, float3 *v3, float3 *v4,
+                                    const float w[4])
 {
 	p->x = v1->x * w[0] + v2->x * w[1] + v3->x * w[2] + v4->x * w[3];
 	p->y = v1->y * w[0] + v2->y * w[1] + v3->y * w[2] + v4->y * w[3];
 	p->z = v1->z * w[0] + v2->z * w[1] + v3->z * w[2] + v4->z * w[3];
 }
 
-float shaperadius(float shape, float root, float tip, float time)
+static float shaperadius(float shape, float root, float tip, float time)
 {
 	float radius = 1.0f - time;
-	
+
 	if(shape != 0.0f) {
 		if(shape < 0.0f)
 			radius = powf(radius, 1.0f + shape);
@@ -87,7 +74,13 @@ float shaperadius(float shape, float root, float tip, float time)
 
 /* curve functions */
 
-void InterpolateKeySegments(int seg, int segno, int key, int curve, float3 *keyloc, float *time, ParticleCurveData *CData)
+static void InterpolateKeySegments(int seg,
+                                   int segno,
+                                   int key,
+                                   int curve,
+                                   float3 *keyloc,
+                                   float *time,
+                                   ParticleCurveData *CData)
 {
 	float3 ckey_loc1 = CData->curvekey_co[key];
 	float3 ckey_loc2 = ckey_loc1;
@@ -116,16 +109,244 @@ void InterpolateKeySegments(int seg, int segno, int key, int curve, float3 *keyl
 		curveinterp_v3_v3v3v3v3(keyloc, &ckey_loc1, &ckey_loc2, &ckey_loc3, &ckey_loc4, t);
 }
 
-void ExportCurveTrianglePlanes(Mesh *mesh, ParticleCurveData *CData,
-                               float3 RotCam, bool is_ortho)
+static bool ObtainCacheParticleData(Mesh *mesh,
+                                    BL::Mesh *b_mesh,
+                                    BL::Object *b_ob,
+                                    ParticleCurveData *CData,
+                                    bool background)
+{
+	int curvenum = 0;
+	int keyno = 0;
+
+	if(!(mesh && b_mesh && b_ob && CData))
+		return false;
+
+	Transform tfm = get_transform(b_ob->matrix_world());
+	Transform itfm = transform_quick_inverse(tfm);
+
+	BL::Object::modifiers_iterator b_mod;
+	for(b_ob->modifiers.begin(b_mod); b_mod != b_ob->modifiers.end(); ++b_mod) {
+		if((b_mod->type() == b_mod->type_PARTICLE_SYSTEM) && (background ? b_mod->show_render() : b_mod->show_viewport())) {
+			BL::ParticleSystemModifier psmd((const PointerRNA)b_mod->ptr);
+			BL::ParticleSystem b_psys((const PointerRNA)psmd.particle_system().ptr);
+			BL::ParticleSettings b_part((const PointerRNA)b_psys.settings().ptr);
+
+			if((b_part.render_type() == BL::ParticleSettings::render_type_PATH) && (b_part.type() == BL::ParticleSettings::type_HAIR)) {
+				int shader = clamp(b_part.material()-1, 0, mesh->used_shaders.size()-1);
+				int draw_step = background ? b_part.render_step() : b_part.draw_step();
+				int totparts = b_psys.particles.length();
+				int totchild = background ? b_psys.child_particles.length() : (int)((float)b_psys.child_particles.length() * (float)b_part.draw_percentage() / 100.0f);
+				int totcurves = totchild;
+
+				if(b_part.child_type() == 0 || totchild == 0)
+					totcurves += totparts;
+
+				if(totcurves == 0)
+					continue;
+
+				int ren_step = (1 << draw_step) + 1;
+				if(b_part.kink() == BL::ParticleSettings::kink_SPIRAL)
+					ren_step += b_part.kink_extra_steps();
+
+				PointerRNA cpsys = RNA_pointer_get(&b_part.ptr, "cycles");
+
+				CData->psys_firstcurve.push_back_slow(curvenum);
+				CData->psys_curvenum.push_back_slow(totcurves);
+				CData->psys_shader.push_back_slow(shader);
+
+				float radius = get_float(cpsys, "radius_scale") * 0.5f;
+
+				CData->psys_rootradius.push_back_slow(radius * get_float(cpsys, "root_width"));
+				CData->psys_tipradius.push_back_slow(radius * get_float(cpsys, "tip_width"));
+				CData->psys_shape.push_back_slow(get_float(cpsys, "shape"));
+				CData->psys_closetip.push_back_slow(get_boolean(cpsys, "use_closetip"));
+
+				int pa_no = 0;
+				if(!(b_part.child_type() == 0) && totchild != 0)
+					pa_no = totparts;
+
+				int num_add = (totparts+totchild - pa_no);
+				CData->curve_firstkey.reserve(CData->curve_firstkey.size() + num_add);
+				CData->curve_keynum.reserve(CData->curve_keynum.size() + num_add);
+				CData->curve_length.reserve(CData->curve_length.size() + num_add);
+				CData->curvekey_co.reserve(CData->curvekey_co.size() + num_add*ren_step);
+				CData->curvekey_time.reserve(CData->curvekey_time.size() + num_add*ren_step);
+
+				for(; pa_no < totparts+totchild; pa_no++) {
+					int keynum = 0;
+					CData->curve_firstkey.push_back_slow(keyno);
+
+					float curve_length = 0.0f;
+					float3 pcKey;
+					for(int step_no = 0; step_no < ren_step; step_no++) {
+						float nco[3];
+						b_psys.co_hair(*b_ob, pa_no, step_no, nco);
+						float3 cKey = make_float3(nco[0], nco[1], nco[2]);
+						cKey = transform_point(&itfm, cKey);
+						if(step_no > 0) {
+							float step_length = len(cKey - pcKey);
+							if(step_length == 0.0f)
+								continue;
+							curve_length += step_length;
+						}
+						CData->curvekey_co.push_back_slow(cKey);
+						CData->curvekey_time.push_back_slow(curve_length);
+						pcKey = cKey;
+						keynum++;
+					}
+					keyno += keynum;
+
+					CData->curve_keynum.push_back_slow(keynum);
+					CData->curve_length.push_back_slow(curve_length);
+					curvenum++;
+				}
+			}
+		}
+	}
+
+	return true;
+}
+
+static bool ObtainCacheParticleUV(Mesh *mesh,
+                                  BL::Mesh *b_mesh,
+                                  BL::Object *b_ob,
+                                  ParticleCurveData *CData,
+                                  bool background,
+                                  int uv_num)
+{
+	if(!(mesh && b_mesh && b_ob && CData))
+		return false;
+
+	CData->curve_uv.clear();
+
+	BL::Object::modifiers_iterator b_mod;
+	for(b_ob->modifiers.begin(b_mod); b_mod != b_ob->modifiers.end(); ++b_mod) {
+		if((b_mod->type() == b_mod->type_PARTICLE_SYSTEM) && (background ? b_mod->show_render() : b_mod->show_viewport())) {
+			BL::ParticleSystemModifier psmd((const PointerRNA)b_mod->ptr);
+			BL::ParticleSystem b_psys((const PointerRNA)psmd.particle_system().ptr);
+			BL::ParticleSettings b_part((const PointerRNA)b_psys.settings().ptr);
+
+			if((b_part.render_type() == BL::ParticleSettings::render_type_PATH) && (b_part.type() == BL::ParticleSettings::type_HAIR)) {
+				int totparts = b_psys.particles.length();
+				int totchild = background ? b_psys.child_particles.length() : (int)((float)b_psys.child_particles.length() * (float)b_part.draw_percentage() / 100.0f);
+				int totcurves = totchild;
+
+				if(b_part.child_type() == 0 || totchild == 0)
+					totcurves += totparts;
+
+				if(totcurves == 0)
+					continue;
+
+				int pa_no = 0;
+				if(!(b_part.child_type() == 0) && totchild != 0)
+					pa_no = totparts;
+
+				int num_add = (totparts+totchild - pa_no);
+				CData->curve_uv.reserve(CData->curve_uv.size() + num_add);
+
+				BL::ParticleSystem::particles_iterator b_pa;
+				b_psys.particles.begin(b_pa);
+				for(; pa_no < totparts+totchild; pa_no++) {
+					/* Add UVs */
+					BL::Mesh::tessface_uv_textures_iterator l;
+					b_mesh->tessface_uv_textures.begin(l);
+
+					float3 uv = make_float3(0.0f, 0.0f, 0.0f);
+					if(b_mesh->tessface_uv_textures.length())
+						b_psys.uv_on_emitter(psmd, *b_pa, pa_no, uv_num, &uv.x);
+					CData->curve_uv.push_back_slow(uv);
+
+					if(pa_no < totparts && b_pa != b_psys.particles.end())
+						++b_pa;
+				}
+			}
+		}
+	}
+
+	return true;
+}
+
+static bool ObtainCacheParticleVcol(Mesh *mesh,
+                                    BL::Mesh *b_mesh,
+                                    BL::Object *b_ob,
+                                    ParticleCurveData *CData,
+                                    bool background,
+                                    int vcol_num)
+{
+	if(!(mesh && b_mesh && b_ob && CData))
+		return false;
+
+	CData->curve_vcol.clear();
+
+	BL::Object::modifiers_iterator b_mod;
+	for(b_ob->modifiers.begin(b_mod); b_mod != b_ob->modifiers.end(); ++b_mod) {
+		if((b_mod->type() == b_mod->type_PARTICLE_SYSTEM) && (background ? b_mod->show_render() : b_mod->show_viewport())) {
+			BL::ParticleSystemModifier psmd((const PointerRNA)b_mod->ptr);
+			BL::ParticleSystem b_psys((const PointerRNA)psmd.particle_system().ptr);
+			BL::ParticleSettings b_part((const PointerRNA)b_psys.settings().ptr);
+
+			if((b_part.render_type() == BL::ParticleSettings::render_type_PATH) && (b_part.type() == BL::ParticleSettings::type_HAIR)) {
+				int totparts = b_psys.particles.length();
+				int totchild = background ? b_psys.child_particles.length() : (int)((float)b_psys.child_particles.length() * (float)b_part.draw_percentage() / 100.0f);
+				int totcurves = totchild;
+
+				if(b_part.child_type() == 0 || totchild == 0)
+					totcurves += totparts;
+
+				if(totcurves == 0)
+					continue;
+
+				int pa_no = 0;
+				if(!(b_part.child_type() == 0) && totchild != 0)
+					pa_no = totparts;
+
+				int num_add = (totparts+totchild - pa_no);
+				CData->curve_vcol.reserve(CData->curve_vcol.size() + num_add);
+
+				BL::ParticleSystem::particles_iterator b_pa;
+				b_psys.particles.begin(b_pa);
+				for(; pa_no < totparts+totchild; pa_no++) {
+					/* Add vertex colors */
+					BL::Mesh::tessface_vertex_colors_iterator l;
+					b_mesh->tessface_vertex_colors.begin(l);
+
+					float3 vcol = make_float3(0.0f, 0.0f, 0.0f);
+					if(b_mesh->tessface_vertex_colors.length())
+						b_psys.mcol_on_emitter(psmd, *b_pa, pa_no, vcol_num, &vcol.x);
+					CData->curve_vcol.push_back_slow(vcol);
+
+					if(pa_no < totparts && b_pa != b_psys.particles.end())
+						++b_pa;
+				}
+			}
+		}
+	}
+
+	return true;
+}
+
+static void set_resolution(BL::Object *b_ob, BL::Scene *scene, bool render)
+{
+	BL::Object::modifiers_iterator b_mod;
+	for(b_ob->modifiers.begin(b_mod); b_mod != b_ob->modifiers.end(); ++b_mod) {
+		if((b_mod->type() == b_mod->type_PARTICLE_SYSTEM) && ((b_mod->show_viewport()) || (b_mod->show_render()))) {
+			BL::ParticleSystemModifier psmd((const PointerRNA)b_mod->ptr);
+			BL::ParticleSystem b_psys((const PointerRNA)psmd.particle_system().ptr);
+			b_psys.set_resolution(*scene, *b_ob, (render)? 2: 1);
+		}
+	}
+}
+
+static void ExportCurveTrianglePlanes(Mesh *mesh, ParticleCurveData *CData,
+                                      float3 RotCam, bool is_ortho)
 {
 	int vertexno = mesh->verts.size();
 	int vertexindex = vertexno;
 	int numverts = 0, numtris = 0;
 
 	/* compute and reserve size of arrays */
-	for(int sys = 0; sys < CData->psys_firstcurve.size() ; sys++) {
-		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys] ; curve++) {
+	for(int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
+		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys]; curve++) {
 			if(CData->curve_keynum[curve] <= 1 || CData->curve_length[curve] == 0.0f)
 				continue;
 
@@ -137,8 +358,8 @@ void ExportCurveTrianglePlanes(Mesh *mesh, ParticleCurveData *CData,
 	mesh->reserve_mesh(mesh->verts.size() + numverts, mesh->num_triangles() + numtris);
 
 	/* actually export */
-	for(int sys = 0; sys < CData->psys_firstcurve.size() ; sys++) {
-		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys] ; curve++) {
+	for(int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
+		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys]; curve++) {
 			if(CData->curve_keynum[curve] <= 1 || CData->curve_length[curve] == 0.0f)
 				continue;
 
@@ -163,7 +384,7 @@ void ExportCurveTrianglePlanes(Mesh *mesh, ParticleCurveData *CData,
 
 				if(curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1)
 					v1 = CData->curvekey_co[curvekey] - CData->curvekey_co[max(curvekey - 1, CData->curve_firstkey[curve])];
-				else 
+				else
 					v1 = CData->curvekey_co[curvekey + 1] - CData->curvekey_co[curvekey - 1];
 
 				time = CData->curvekey_time[curvekey]/CData->curve_length[curve];
@@ -199,15 +420,17 @@ void ExportCurveTrianglePlanes(Mesh *mesh, ParticleCurveData *CData,
 	/* texture coords still needed */
 }
 
-void ExportCurveTriangleGeometry(Mesh *mesh, ParticleCurveData *CData, int resolution)
+static void ExportCurveTriangleGeometry(Mesh *mesh,
+                                        ParticleCurveData *CData,
+                                        int resolution)
 {
 	int vertexno = mesh->verts.size();
 	int vertexindex = vertexno;
 	int numverts = 0, numtris = 0;
 
 	/* compute and reserve size of arrays */
-	for(int sys = 0; sys < CData->psys_firstcurve.size() ; sys++) {
-		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys] ; curve++) {
+	for(int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
+		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys]; curve++) {
 			if(CData->curve_keynum[curve] <= 1 || CData->curve_length[curve] == 0.0f)
 				continue;
 
@@ -219,8 +442,8 @@ void ExportCurveTriangleGeometry(Mesh *mesh, ParticleCurveData *CData, int resol
 	mesh->reserve_mesh(mesh->verts.size() + numverts, mesh->num_triangles() + numtris);
 
 	/* actually export */
-	for(int sys = 0; sys < CData->psys_firstcurve.size() ; sys++) {
-		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys] ; curve++) {
+	for(int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
+		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys]; curve++) {
 			if(CData->curve_keynum[curve] <= 1 || CData->curve_length[curve] == 0.0f)
 				continue;
 
@@ -331,7 +554,7 @@ void ExportCurveTriangleGeometry(Mesh *mesh, ParticleCurveData *CData, int resol
 	/* texture coords still needed */
 }
 
-void ExportCurveSegments(Scene *scene, Mesh *mesh, ParticleCurveData *CData)
+static void ExportCurveSegments(Scene *scene, Mesh *mesh, ParticleCurveData *CData)
 {
 	int num_keys = 0;
 	int num_curves = 0;
@@ -340,13 +563,13 @@ void ExportCurveSegments(Scene *scene, Mesh *mesh, ParticleCurveData *CData)
 		return;
 
 	Attribute *attr_intercept = NULL;
-	
+
 	if(mesh->need_attribute(scene, ATTR_STD_CURVE_INTERCEPT))
 		attr_intercept = mesh->curve_attributes.add(ATTR_STD_CURVE_INTERCEPT);
 
 	/* compute and reserve size of arrays */
-	for(int sys = 0; sys < CData->psys_firstcurve.size() ; sys++) {
-		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys] ; curve++) {
+	for(int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
+		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys]; curve++) {
 			if(CData->curve_keynum[curve] <= 1 || CData->curve_length[curve] == 0.0f)
 				continue;
 
@@ -365,8 +588,8 @@ void ExportCurveSegments(Scene *scene, Mesh *mesh, ParticleCurveData *CData)
 	num_curves = 0;
 
 	/* actually export */
-	for(int sys = 0; sys < CData->psys_firstcurve.size() ; sys++) {
-		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys] ; curve++) {
+	for(int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
+		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys]; curve++) {
 			if(CData->curve_keynum[curve] <= 1 || CData->curve_length[curve] == 0.0f)
 				continue;
 
@@ -460,8 +683,13 @@ static void ExportCurveSegmentsMotion(Mesh *mesh, ParticleCurveData *CData, int
 	/* in case of new attribute, we verify if there really was any motion */
 	if(new_attribute) {
 		if(i != numkeys || !have_motion) {
-			/* no motion, remove attributes again */
-			VLOG(1) << "No motion, removing attribute";
+			/* No motion or hair "topology" changed, remove attributes again. */
+			if(i != numkeys) {
+				VLOG(1) << "Hair topology changed, removing attribute.";
+			}
+			else {
+				VLOG(1) << "No motion, removing attribute.";
+			}
 			mesh->curve_attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION);
 		}
 		else if(time_index > 0) {
@@ -481,7 +709,10 @@ static void ExportCurveSegmentsMotion(Mesh *mesh, ParticleCurveData *CData, int
 	}
 }
 
-void ExportCurveTriangleUV(ParticleCurveData *CData, int vert_offset, int resol, float3 *uvdata)
+static void ExportCurveTriangleUV(ParticleCurveData *CData,
+                                  int vert_offset,
+                                  int resol,
+                                  float3 *uvdata)
 {
 	if(uvdata == NULL)
 		return;
@@ -491,8 +722,8 @@ void ExportCurveTriangleUV(ParticleCurveData *CData, int vert_offset, int resol,
 
 	int vertexindex = vert_offset;
 
-	for(int sys = 0; sys < CData->psys_firstcurve.size() ; sys++) {
-		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys] ; curve++) {
+	for(int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
+		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys]; curve++) {
 			if(CData->curve_keynum[curve] <= 1 || CData->curve_length[curve] == 0.0f)
 				continue;
 
@@ -526,15 +757,18 @@ void ExportCurveTriangleUV(ParticleCurveData *CData, int vert_offset, int resol,
 	}
 }
 
-void ExportCurveTriangleVcol(ParticleCurveData *CData, int vert_offset, int resol, uchar4 *cdata)
+static void ExportCurveTriangleVcol(ParticleCurveData *CData,
+                                    int vert_offset,
+                                    int resol,
+                                    uchar4 *cdata)
 {
 	if(cdata == NULL)
 		return;
 
 	int vertexindex = vert_offset;
 
-	for(int sys = 0; sys < CData->psys_firstcurve.size() ; sys++) {
-		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys] ; curve++) {
+	for(int sys = 0; sys < CData->psys_firstcurve.size(); sys++) {
+		for(int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys]; curve++) {
 			if(CData->curve_keynum[curve] <= 1 || CData->curve_length[curve] == 0.0f)
 				continue;
 
@@ -620,6 +854,20 @@ void BlenderSync::sync_curve_settings()
 	}
 
 	if(curve_system_manager->modified_mesh(prev_curve_system_manager)) {
+		BL::BlendData::objects_iterator b_ob;
+
+		for(b_data.objects.begin(b_ob); b_ob != b_data.objects.end(); ++b_ob) {
+			if(object_is_mesh(*b_ob)) {
+				BL::Object::particle_systems_iterator b_psys;
+				for(b_ob->particle_systems.begin(b_psys); b_psys != b_ob->particle_systems.end(); ++b_psys) {
+					if((b_psys->settings().render_type()==BL::ParticleSettings::render_type_PATH)&&(b_psys->settings().type()==BL::ParticleSettings::type_HAIR)) {
+						BL::ID key = BKE_object_is_modified(*b_ob)? *b_ob: b_ob->data();
+						mesh_map.set_recalc(key);
+						object_map.set_recalc(*b_ob);
+					}
+				}
+			}
+		}
 	}
 
 	if(curve_system_manager->modified(prev_curve_system_manager))
@@ -644,7 +892,7 @@ void BlenderSync::sync_curves(Mesh *mesh,
 	/* obtain general settings */
 	bool use_curves = scene->curve_system_manager->use_curves;
 
-	if(!use_curves) {
+	if(!(use_curves && b_ob.mode() != b_ob.mode_PARTICLE_EDIT)) {
 		if(!motion)
 			mesh->compute_bounds();
 		return;
@@ -661,6 +909,11 @@ void BlenderSync::sync_curves(Mesh *mesh,
 
 	ParticleCurveData CData;
 
+	if(!preview)
+		set_resolution(&b_ob, &b_scene, true);
+
+	ObtainCacheParticleData(mesh, &b_mesh, &b_ob, &CData, !preview);
+
 	/* add hair geometry to mesh */
 	if(primitive == CURVE_TRIANGLES) {
 		if(triangle_method == CURVE_CAMERA_TRIANGLES) {
@@ -728,6 +981,8 @@ void BlenderSync::sync_curves(Mesh *mesh,
 			if(!mesh->need_attribute(scene, ustring(l->name().c_str())))
 				continue;
 
+			ObtainCacheParticleVcol(mesh, &b_mesh, &b_ob, &CData, !preview, vcol_num);
+
 			if(primitive == CURVE_TRIANGLES) {
 				Attribute *attr_vcol = mesh->attributes.add(
 					ustring(l->name().c_str()), TypeDesc::TypeColor, ATTR_ELEMENT_CORNER_BYTE);
@@ -767,6 +1022,8 @@ void BlenderSync::sync_curves(Mesh *mesh,
 			if(mesh->need_attribute(scene, name) || mesh->need_attribute(scene, std)) {
 				Attribute *attr_uv;
 
+				ObtainCacheParticleUV(mesh, &b_mesh, &b_ob, &CData, !preview, uv_num);
+
 				if(primitive == CURVE_TRIANGLES) {
 					if(active_render)
 						attr_uv = mesh->attributes.add(std, name);
@@ -797,8 +1054,10 @@ void BlenderSync::sync_curves(Mesh *mesh,
 		}
 	}
 
+	if(!preview)
+		set_resolution(&b_ob, &b_scene, false);
+
 	mesh->compute_bounds();
 }
 
 CCL_NAMESPACE_END
-

intern/cycles/blender/blender_mesh.cpp

@@ -597,8 +597,8 @@ static void create_mesh(Scene *scene,
                         Mesh *mesh,
                         BL::Mesh& b_mesh,
                         const vector<Shader*>& used_shaders,
-                        bool subdivision=false,
-                        bool subdivide_uvs=true)
+                        bool subdivision = false,
+                        bool subdivide_uvs = true)
 {
 	/* count vertices and faces */
 	int numverts = b_mesh.vertices.length();
@@ -671,28 +671,10 @@ static void create_mesh(Scene *scene,
 			int shader = clamp(f->material_index(), 0, used_shaders.size()-1);
 			bool smooth = f->use_smooth() || use_loop_normals;
 
-			/* split vertices if normal is different
+			/* Create triangles.
 			 *
-			 * note all vertex attributes must have been set here so we can split
-			 * and copy attributes in split_vertex without remapping later */
-			if(use_loop_normals) {
-				BL::Array<float, 12> loop_normals = f->split_normals();
-
-				for(int i = 0; i < n; i++) {
-					float3 loop_N = make_float3(loop_normals[i * 3], loop_normals[i * 3 + 1], loop_normals[i * 3 + 2]);
-
-					if(N[vi[i]] != loop_N) {
-						int new_vi = mesh->split_vertex(vi[i]);
-
-						/* set new normal and vertex index */
-						N = attr_N->data_float3();
-						N[new_vi] = loop_N;
-						vi[i] = new_vi;
-					}
-				}
-			}
-
-			/* create triangles */
+			 * NOTE: Autosmooth is already taken care about.
+			 */
 			if(n == 4) {
 				if(is_zero(cross(mesh->verts[vi[1]] - mesh->verts[vi[0]], mesh->verts[vi[2]] - mesh->verts[vi[0]])) ||
 				   is_zero(cross(mesh->verts[vi[2]] - mesh->verts[vi[0]], mesh->verts[vi[3]] - mesh->verts[vi[0]])))
@@ -724,24 +706,8 @@ static void create_mesh(Scene *scene,
 
 			vi.reserve(n);
 			for(int i = 0; i < n; i++) {
+				/* NOTE: Autosmooth is already taken care about. */
 				vi[i] = b_mesh.loops[p->loop_start() + i].vertex_index();
-
-				/* split vertices if normal is different
-				 *
-				 * note all vertex attributes must have been set here so we can split
-				 * and copy attributes in split_vertex without remapping later */
-				if(use_loop_normals) {
-					float3 loop_N = get_float3(b_mesh.loops[p->loop_start() + i].normal());
-
-					if(N[vi[i]] != loop_N) {
-						int new_vi = mesh->split_vertex(vi[i]);
-
-						/* set new normal and vertex index */
-						N = attr_N->data_float3();
-						N[new_vi] = loop_N;
-						vi[i] = new_vi;
-					}
-				}
 			}
 
 			/* create subd faces */
@@ -961,7 +927,13 @@ Mesh *BlenderSync::sync_mesh(BL::Object& b_ob,
 
 		mesh->subdivision_type = object_subdivision_type(b_ob, preview, experimental);
 
-		BL::Mesh b_mesh = object_to_mesh(b_data, b_ob, b_scene, true, !preview, need_undeformed, mesh->subdivision_type);
+		BL::Mesh b_mesh = object_to_mesh(b_data,
+		                                 b_ob,
+		                                 b_scene,
+		                                 true,
+		                                 !preview,
+		                                 need_undeformed,
+		                                 mesh->subdivision_type);
 
 		if(b_mesh) {
 			if(render_layer.use_surfaces && !hide_tris) {
@@ -1086,7 +1058,13 @@ void BlenderSync::sync_mesh_motion(BL::Object& b_ob,
 
 	if(ccl::BKE_object_is_deform_modified(b_ob, b_scene, preview)) {
 		/* get derived mesh */
-		b_mesh = object_to_mesh(b_data, b_ob, b_scene, true, !preview, false, false);
+		b_mesh = object_to_mesh(b_data,
+		                        b_ob,
+		                        b_scene,
+		                        true,
+		                        !preview,
+		                        false,
+		                        Mesh::SUBDIVISION_NONE);
 	}
 
 	if(!b_mesh) {
@@ -1157,10 +1135,12 @@ void BlenderSync::sync_mesh_motion(BL::Object& b_ob,
 			{
 				/* no motion, remove attributes again */
 				if(b_mesh.vertices.length() != numverts) {
-					VLOG(1) << "Topology differs, disabling motion blur.";
+					VLOG(1) << "Topology differs, disabling motion blur for object "
+					        << b_ob.name();
 				}
 				else {
-					VLOG(1) << "No actual deformation motion for object " << b_ob.name();
+					VLOG(1) << "No actual deformation motion for object "
+					        << b_ob.name();
 				}
 				mesh->attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION);
 				if(attr_mN)

intern/cycles/blender/blender_object.cpp

@@ -25,6 +25,7 @@
 #include "particles.h"
 #include "shader.h"
 
+#include "blender_object_cull.h"
 #include "blender_sync.h"
 #include "blender_util.h"
 
@@ -235,55 +236,6 @@ void BlenderSync::sync_background_light(bool use_portal)
 
 /* Object */
 
-/* TODO(sergey): Not really optimal, consider approaches based on k-DOP in order
- * to reduce number of objects which are wrongly considered visible.
- */
-static bool object_boundbox_clip(Scene *scene,
-                                 BL::Object& b_ob,
-                                 Transform& tfm,
-                                 float margin)
-{
-	Camera *cam = scene->camera;
-	Transform& worldtondc = cam->worldtondc;
-	BL::Array<float, 24> boundbox = b_ob.bound_box();
-	float3 bb_min = make_float3(FLT_MAX, FLT_MAX, FLT_MAX),
-	       bb_max = make_float3(-FLT_MAX, -FLT_MAX, -FLT_MAX);
-	bool all_behind = true;
-	for(int i = 0; i < 8; ++i) {
-		float3 p = make_float3(boundbox[3 * i + 0],
-		                       boundbox[3 * i + 1],
-		                       boundbox[3 * i + 2]);
-		p = transform_point(&tfm, p);
-
-		float4 b = make_float4(p.x, p.y, p.z, 1.0f);
-		float4 c = make_float4(dot(worldtondc.x, b),
-		                       dot(worldtondc.y, b),
-		                       dot(worldtondc.z, b),
-		                       dot(worldtondc.w, b));
-		p = float4_to_float3(c / c.w);
-		if(c.z < 0.0f) {
-			p.x = 1.0f - p.x;
-			p.y = 1.0f - p.y;
-		}
-		if(c.z >= -margin) {
-			all_behind = false;
-		}
-		bb_min = min(bb_min, p);
-		bb_max = max(bb_max, p);
-	}
-	if(!all_behind) {
-		if(bb_min.x >= 1.0f + margin ||
-		   bb_min.y >= 1.0f + margin ||
-		   bb_max.x <= -margin ||
-		   bb_max.y <= -margin)
-		{
-			return true;
-		}
-		return false;
-	}
-	return true;
-}
-
 Object *BlenderSync::sync_object(BL::Object& b_parent,
                                  int persistent_id[OBJECT_PERSISTENT_ID_SIZE],
                                  BL::DupliObject& b_dupli_ob,
@@ -291,8 +243,7 @@ Object *BlenderSync::sync_object(BL::Object& b_parent,
                                  uint layer_flag,
                                  float motion_time,
                                  bool hide_tris,
-                                 bool use_camera_cull,
-                                 float camera_cull_margin,
+                                 BlenderObjectCulling& culling,
                                  bool *use_portal)
 {
 	BL::Object b_ob = (b_dupli_ob ? b_dupli_ob.object() : b_parent);
@@ -308,11 +259,12 @@ Object *BlenderSync::sync_object(BL::Object& b_parent,
 	}
 
 	/* only interested in object that we can create meshes from */
-	if(!object_is_mesh(b_ob))
+	if(!object_is_mesh(b_ob)) {
 		return NULL;
+	}
 
-	/* Perform camera space culling. */
-	if(use_camera_cull && object_boundbox_clip(scene, b_ob, tfm, camera_cull_margin)) {
+	/* Perform object culling. */
+	if(culling.test(scene, b_ob, tfm)) {
 		return NULL;
 	}
 
@@ -466,9 +418,29 @@ static bool object_render_hide(BL::Object& b_ob,
                                bool parent_hide,
                                bool& hide_triangles)
 {
+	/* check if we should render or hide particle emitter */
+	BL::Object::particle_systems_iterator b_psys;
+
+	bool hair_present = false;
+	bool show_emitter = false;
+	bool hide_emitter = false;
 	bool hide_as_dupli_parent = false;
 	bool hide_as_dupli_child_original = false;
 
+	for(b_ob.particle_systems.begin(b_psys); b_psys != b_ob.particle_systems.end(); ++b_psys) {
+		if((b_psys->settings().render_type() == BL::ParticleSettings::render_type_PATH) &&
+		   (b_psys->settings().type()==BL::ParticleSettings::type_HAIR))
+			hair_present = true;
+
+		if(b_psys->settings().use_render_emitter())
+			show_emitter = true;
+		else
+			hide_emitter = true;
+	}
+
+	if(show_emitter)
+		hide_emitter = false;
+
 	/* duplicators hidden by default, except dupliframes which duplicate self */
 	if(b_ob.is_duplicator())
 		if(top_level || b_ob.dupli_type() != BL::Object::dupli_type_FRAMES)
@@ -488,9 +460,17 @@ static bool object_render_hide(BL::Object& b_ob,
 		parent = parent.parent();
 	}
 	
-	hide_triangles = false;
+	hide_triangles = hide_emitter;
 
-	return (hide_as_dupli_parent || hide_as_dupli_child_original);
+	if(show_emitter) {
+		return false;
+	}
+	else if(hair_present) {
+		return hide_as_dupli_child_original;
+	}
+	else {
+		return (hide_as_dupli_parent || hide_as_dupli_child_original);
+	}
 }
 
 static bool object_render_hide_duplis(BL::Object& b_ob)
@@ -513,23 +493,15 @@ void BlenderSync::sync_objects(BL::SpaceView3D& b_v3d, float motion_time)
 		light_map.pre_sync();
 		mesh_map.pre_sync();
 		object_map.pre_sync();
+		particle_system_map.pre_sync();
 		motion_times.clear();
 	}
 	else {
 		mesh_motion_synced.clear();
 	}
 
-	bool allow_camera_cull = false;
-	float camera_cull_margin = 0.0f;
-	if(b_scene.render().use_simplify()) {
-		PointerRNA cscene = RNA_pointer_get(&b_scene.ptr, "cycles");
-		allow_camera_cull = scene->camera->type != CAMERA_PANORAMA &&
-		                    !b_scene.render().use_multiview() &&
-		                    get_boolean(cscene, "use_camera_cull");
-		if(allow_camera_cull) {
-			camera_cull_margin = get_float(cscene, "camera_cull_margin");
-		}
-	}
+	/* initialize culling */
+	BlenderObjectCulling culling(scene, b_scene);
 
 	/* object loop */
 	BL::Scene::object_bases_iterator b_base;
@@ -561,12 +533,9 @@ void BlenderSync::sync_objects(BL::SpaceView3D& b_v3d, float motion_time)
 			if(!hide) {
 				progress.set_sync_status("Synchronizing object", b_ob.name());
 
-				PointerRNA cobject = RNA_pointer_get(&b_ob.ptr, "cycles");
-				bool use_camera_cull = allow_camera_cull && get_boolean(cobject, "use_camera_cull");
-				if(use_camera_cull) {
-					/* Need to have proper projection matrix. */
-					scene->camera->update();
-				}
+				/* load per-object culling data */
+				culling.init_object(scene, b_ob);
+
 				if(b_ob.is_duplicator() && !object_render_hide_duplis(b_ob)) {
 					/* dupli objects */
 					b_ob.dupli_list_create(b_scene, dupli_settings);
@@ -586,16 +555,22 @@ void BlenderSync::sync_objects(BL::SpaceView3D& b_v3d, float motion_time)
 							BL::Array<int, OBJECT_PERSISTENT_ID_SIZE> persistent_id = b_dup->persistent_id();
 
 							/* sync object and mesh or light data */
-							sync_object(b_ob,
-							            persistent_id.data,
-							            *b_dup,
-							            tfm,
-							            ob_layer,
-							            motion_time,
-							            hide_tris,
-							            use_camera_cull,
-							            camera_cull_margin,
-							            &use_portal);
+							Object *object = sync_object(b_ob,
+							                             persistent_id.data,
+							                             *b_dup,
+							                             tfm,
+							                             ob_layer,
+							                             motion_time,
+							                             hide_tris,
+							                             culling,
+							                             &use_portal);
+
+							/* sync possible particle data, note particle_id
+							 * starts counting at 1, first is dummy particle */
+							if(!motion && object) {
+								sync_dupli_particle(b_ob, *b_dup, object);
+							}
+
 						}
 					}
 
@@ -616,8 +591,7 @@ void BlenderSync::sync_objects(BL::SpaceView3D& b_v3d, float motion_time)
 					            ob_layer,
 					            motion_time,
 					            hide_tris,
-					            use_camera_cull,
-					            camera_cull_margin,
+					            culling,
 					            &use_portal);
 				}
 			}
@@ -638,6 +612,8 @@ void BlenderSync::sync_objects(BL::SpaceView3D& b_v3d, float motion_time)
 			scene->mesh_manager->tag_update(scene);
 		if(object_map.post_sync())
 			scene->object_manager->tag_update(scene);
+		if(particle_system_map.post_sync())
+			scene->particle_system_manager->tag_update(scene);
 	}
 
 	if(motion)

intern/cycles/blender/blender_object_cull.cpp

@@ -0,0 +1,149 @@
+/*
+ * Copyright 2011-2016 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.
+ */
+
+#include <cstdlib>
+
+#include "camera.h"
+
+#include "blender_object_cull.h"
+
+CCL_NAMESPACE_BEGIN
+
+BlenderObjectCulling::BlenderObjectCulling(Scene *scene, BL::Scene& b_scene)
+        : use_scene_camera_cull_(false),
+          use_camera_cull_(false),
+          camera_cull_margin_(0.0f),
+          use_scene_distance_cull_(false),
+          use_distance_cull_(false),
+          distance_cull_margin_(0.0f)
+{
+	if(b_scene.render().use_simplify()) {
+		PointerRNA cscene = RNA_pointer_get(&b_scene.ptr, "cycles");
+
+		use_scene_camera_cull_ = scene->camera->type != CAMERA_PANORAMA &&
+		                         !b_scene.render().use_multiview() &&
+		                         get_boolean(cscene, "use_camera_cull");
+		use_scene_distance_cull_ = scene->camera->type != CAMERA_PANORAMA &&
+		                           !b_scene.render().use_multiview() &&
+		                           get_boolean(cscene, "use_distance_cull");
+
+		camera_cull_margin_ = get_float(cscene, "camera_cull_margin");
+		distance_cull_margin_ = get_float(cscene, "distance_cull_margin");
+
+		if(distance_cull_margin_ == 0.0f) {
+			use_scene_distance_cull_ = false;
+		}
+	}
+}
+
+void BlenderObjectCulling::init_object(Scene *scene, BL::Object& b_ob)
+{
+	if(!use_scene_camera_cull_ && !use_scene_distance_cull_) {
+		return;
+	}
+
+	PointerRNA cobject = RNA_pointer_get(&b_ob.ptr, "cycles");
+
+	use_camera_cull_ = use_scene_camera_cull_ && get_boolean(cobject, "use_camera_cull");
+	use_distance_cull_ = use_scene_distance_cull_ && get_boolean(cobject, "use_distance_cull");
+
+	if(use_camera_cull_ || use_distance_cull_) {
+		/* Need to have proper projection matrix. */
+		scene->camera->update();
+	}
+}
+
+bool BlenderObjectCulling::test(Scene *scene, BL::Object& b_ob, Transform& tfm)
+{
+	if(!use_camera_cull_ && !use_distance_cull_) {
+		return false;
+	}
+
+	/* Compute world space bounding box corners. */
+	float3 bb[8];
+	BL::Array<float, 24> boundbox = b_ob.bound_box();
+	for(int i = 0; i < 8; ++i) {
+		float3 p = make_float3(boundbox[3 * i + 0],
+		                       boundbox[3 * i + 1],
+		                       boundbox[3 * i + 2]);
+		bb[i] = transform_point(&tfm, p);
+	}
+
+	bool camera_culled = use_camera_cull_ && test_camera(scene, bb);
+	bool distance_culled = use_distance_cull_ && test_distance(scene, bb);
+
+	return ((camera_culled && distance_culled) ||
+	        (camera_culled && !use_distance_cull_) ||
+	        (distance_culled && !use_camera_cull_));
+}
+
+/* TODO(sergey): Not really optimal, consider approaches based on k-DOP in order
+ * to reduce number of objects which are wrongly considered visible.
+ */
+bool BlenderObjectCulling::test_camera(Scene *scene, float3 bb[8])
+{
+	Camera *cam = scene->camera;
+	Transform& worldtondc = cam->worldtondc;
+	float3 bb_min = make_float3(FLT_MAX, FLT_MAX, FLT_MAX),
+	       bb_max = make_float3(-FLT_MAX, -FLT_MAX, -FLT_MAX);
+	bool all_behind = true;
+	for(int i = 0; i < 8; ++i) {
+		float3 p = bb[i];
+		float4 b = make_float4(p.x, p.y, p.z, 1.0f);
+		float4 c = make_float4(dot(worldtondc.x, b),
+		                       dot(worldtondc.y, b),
+		                       dot(worldtondc.z, b),
+		                       dot(worldtondc.w, b));
+		p = float4_to_float3(c / c.w);
+		if(c.z < 0.0f) {
+			p.x = 1.0f - p.x;
+			p.y = 1.0f - p.y;
+		}
+		if(c.z >= -camera_cull_margin_) {
+			all_behind = false;
+		}
+		bb_min = min(bb_min, p);
+		bb_max = max(bb_max, p);
+	}
+	if(all_behind) {
+		return true;
+	}
+	return (bb_min.x >= 1.0f + camera_cull_margin_ ||
+	        bb_min.y >= 1.0f + camera_cull_margin_ ||
+	        bb_max.x <= -camera_cull_margin_ ||
+	        bb_max.y <= -camera_cull_margin_);
+}
+
+bool BlenderObjectCulling::test_distance(Scene *scene, float3 bb[8])
+{
+	float3 camera_position = transform_get_column(&scene->camera->matrix, 3);
+	float3 bb_min = make_float3(FLT_MAX, FLT_MAX, FLT_MAX),
+	       bb_max = make_float3(-FLT_MAX, -FLT_MAX, -FLT_MAX);
+
+	/* Find min & max points for x & y & z on bounding box */
+	for(int i = 0; i < 8; ++i) {
+		float3 p = bb[i];
+		bb_min = min(bb_min, p);
+		bb_max = max(bb_max, p);
+	}
+
+	float3 closest_point = max(min(bb_max,camera_position),bb_min);
+	return (len_squared(camera_position - closest_point) >
+	        distance_cull_margin_ * distance_cull_margin_);
+}
+
+CCL_NAMESPACE_END
+

intern/cycles/blender/blender_object_cull.h

@@ -0,0 +1,49 @@
+/*
+ * Copyright 2011-2016 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifndef __BLENDER_OBJECT_CULL_H__
+#define __BLENDER_OBJECT_CULL_H__
+
+#include "blender_sync.h"
+#include "util_types.h"
+
+CCL_NAMESPACE_BEGIN
+
+class Scene;
+
+class BlenderObjectCulling
+{
+public:
+	BlenderObjectCulling(Scene *scene, BL::Scene& b_scene);
+
+	void init_object(Scene *scene, BL::Object& b_ob);
+	bool test(Scene *scene, BL::Object& b_ob, Transform& tfm);
+
+private:
+	bool test_camera(Scene *scene, float3 bb[8]);
+	bool test_distance(Scene *scene, float3 bb[8]);
+
+	bool use_scene_camera_cull_;
+	bool use_camera_cull_;
+	float camera_cull_margin_;
+	bool use_scene_distance_cull_;
+	bool use_distance_cull_;
+	float distance_cull_margin_;
+};
+
+CCL_NAMESPACE_END
+
+#endif /* __BLENDER_OBJECT_CULL_H__ */

intern/cycles/blender/blender_particles.cpp

@@ -0,0 +1,92 @@
+/*
+ * Copyright 2011-2013 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.
+ */
+
+#include "mesh.h"
+#include "object.h"
+#include "particles.h"
+
+#include "blender_sync.h"
+#include "blender_util.h"
+
+#include "util_foreach.h"
+
+CCL_NAMESPACE_BEGIN
+
+/* Utilities */
+
+bool BlenderSync::sync_dupli_particle(BL::Object& b_ob,
+                                      BL::DupliObject& b_dup,
+                                      Object *object)
+{
+	/* test if this dupli was generated from a particle sytem */
+	BL::ParticleSystem b_psys = b_dup.particle_system();
+	if(!b_psys)
+		return false;
+
+	object->hide_on_missing_motion = true;
+
+	/* test if we need particle data */
+	if(!object->mesh->need_attribute(scene, ATTR_STD_PARTICLE))
+		return false;
+
+	/* don't handle child particles yet */
+	BL::Array<int, OBJECT_PERSISTENT_ID_SIZE> persistent_id = b_dup.persistent_id();
+
+	if(persistent_id[0] >= b_psys.particles.length())
+		return false;
+
+	/* find particle system */
+	ParticleSystemKey key(b_ob, persistent_id);
+	ParticleSystem *psys;
+
+	bool first_use = !particle_system_map.is_used(key);
+	bool need_update = particle_system_map.sync(&psys, b_ob, b_dup.object(), key);
+
+	/* no update needed? */
+	if(!need_update && !object->mesh->need_update && !scene->object_manager->need_update)
+		return true;
+
+	/* first time used in this sync loop? clear and tag update */
+	if(first_use) {
+		psys->particles.clear();
+		psys->tag_update(scene);
+	}
+
+	/* add particle */
+	BL::Particle b_pa = b_psys.particles[persistent_id[0]];
+	Particle pa;
+	
+	pa.index = persistent_id[0];
+	pa.age = b_scene.frame_current() - b_pa.birth_time();
+	pa.lifetime = b_pa.lifetime();
+	pa.location = get_float3(b_pa.location());
+	pa.rotation = get_float4(b_pa.rotation());
+	pa.size = b_pa.size();
+	pa.velocity = get_float3(b_pa.velocity());
+	pa.angular_velocity = get_float3(b_pa.angular_velocity());
+
+	psys->particles.push_back_slow(pa);
+
+	if(object->particle_index != psys->particles.size() - 1)
+		scene->object_manager->tag_update(scene);
+	object->particle_system = psys;
+	object->particle_index = psys->particles.size() - 1;
+
+	/* return that this object has particle data */
+	return true;
+}
+
+CCL_NAMESPACE_END

intern/cycles/blender/blender_session.cpp

@@ -126,8 +126,8 @@ void BlenderSession::create_session()
 
 	/* setup callbacks for builtin image support */
 	scene->image_manager->builtin_image_info_cb = function_bind(&BlenderSession::builtin_image_info, this, _1, _2, _3, _4, _5, _6, _7);
-	scene->image_manager->builtin_image_pixels_cb = function_bind(&BlenderSession::builtin_image_pixels, this, _1, _2, _3);
-	scene->image_manager->builtin_image_float_pixels_cb = function_bind(&BlenderSession::builtin_image_float_pixels, this, _1, _2, _3);
+	scene->image_manager->builtin_image_pixels_cb = function_bind(&BlenderSession::builtin_image_pixels, this, _1, _2, _3, _4);
+	scene->image_manager->builtin_image_float_pixels_cb = function_bind(&BlenderSession::builtin_image_float_pixels, this, _1, _2, _3, _4);
 
 	/* create session */
 	session = new Session(session_params);
@@ -305,12 +305,16 @@ static PassType get_pass_type(BL::RenderPass& b_pass)
 #ifdef WITH_CYCLES_DEBUG
 		case BL::RenderPass::type_DEBUG:
 		{
-			if(b_pass.debug_type() == BL::RenderPass::debug_type_BVH_TRAVERSAL_STEPS)
-				return PASS_BVH_TRAVERSAL_STEPS;
-			if(b_pass.debug_type() == BL::RenderPass::debug_type_BVH_TRAVERSED_INSTANCES)
-				return PASS_BVH_TRAVERSED_INSTANCES;
-			if(b_pass.debug_type() == BL::RenderPass::debug_type_RAY_BOUNCES)
-				return PASS_RAY_BOUNCES;
+			switch(b_pass.debug_type()) {
+				case BL::RenderPass::debug_type_BVH_TRAVERSED_NODES:
+					return PASS_BVH_TRAVERSED_NODES;
+				case BL::RenderPass::debug_type_BVH_TRAVERSED_INSTANCES:
+					return PASS_BVH_TRAVERSED_INSTANCES;
+				case BL::RenderPass::debug_type_BVH_INTERSECTIONS:
+					return PASS_BVH_INTERSECTIONS;
+				case BL::RenderPass::debug_type_RAY_BOUNCES:
+					return PASS_RAY_BOUNCES;
+			}
 			break;
 		}
 #endif
@@ -580,7 +584,7 @@ static void populate_bake_data(BakeData *data, const
 	BL::BakePixel bp = pixel_array;
 
 	int i;
-	for(i=0; i < num_pixels; i++) {
+	for(i = 0; i < num_pixels; i++) {
 		if(bp.object_id() == object_id) {
 			data->set(i, bp.primitive_id(), bp.uv(), bp.du_dx(), bp.du_dy(), bp.dv_dx(), bp.dv_dy());
 		} else {
@@ -930,38 +934,13 @@ void BlenderSession::get_status(string& status, string& substatus)
 
 void BlenderSession::get_progress(float& progress, double& total_time, double& render_time)
 {
-	double tile_time;
-	int tile, sample, samples_per_tile;
-	int tile_total = session->tile_manager.state.num_tiles;
-	int samples = session->tile_manager.state.sample + 1;
-	int total_samples = session->tile_manager.get_num_effective_samples();
-
-	session->progress.get_tile(tile, total_time, render_time, tile_time);
-
-	sample = session->progress.get_sample();
-	samples_per_tile = session->tile_manager.get_num_effective_samples();
-
-	if(background && samples_per_tile && tile_total)
-		progress = ((float)sample / (float)(tile_total * samples_per_tile));
-	else if(!background && samples > 0 && total_samples != INT_MAX)
-		progress = ((float)samples) / total_samples;
-	else
-		progress = 0.0;
+	session->progress.get_time(total_time, render_time);
+	progress = session->progress.get_progress();
 }
 
 void BlenderSession::update_bake_progress()
 {
-	float progress;
-	int sample, samples_per_task, parts_total;
-
-	sample = session->progress.get_sample();
-	samples_per_task = scene->bake_manager->num_samples;
-	parts_total = scene->bake_manager->num_parts;
-
-	if(samples_per_task)
-		progress = ((float)sample / (float)(parts_total * samples_per_task));
-	else
-		progress = 0.0;
+	float progress = session->progress.get_progress();
 
 	if(progress != last_progress) {
 		b_engine.update_progress(progress);
@@ -1080,7 +1059,13 @@ int BlenderSession::builtin_image_frame(const string &builtin_name)
 	return atoi(builtin_name.substr(last + 1, builtin_name.size() - last - 1).c_str());
 }
 
-void BlenderSession::builtin_image_info(const string &builtin_name, void *builtin_data, bool &is_float, int &width, int &height, int &depth, int &channels)
+void BlenderSession::builtin_image_info(const string &builtin_name,
+                                        void *builtin_data,
+                                        bool &is_float,
+                                        int &width,
+                                        int &height,
+                                        int &depth,
+                                        int &channels)
 {
 	/* empty image */
 	is_float = false;
@@ -1158,60 +1143,67 @@ void BlenderSession::builtin_image_info(const string &builtin_name, void *builti
 	}
 }
 
-bool BlenderSession::builtin_image_pixels(const string &builtin_name, void *builtin_data, unsigned char *pixels)
+bool BlenderSession::builtin_image_pixels(const string &builtin_name,
+                                          void *builtin_data,
+                                          unsigned char *pixels,
+                                          const size_t pixels_size)
 {
-	if(!builtin_data)
+	if(!builtin_data) {
 		return false;
+	}
 
-	int frame = builtin_image_frame(builtin_name);
+	const int frame = builtin_image_frame(builtin_name);
 
 	PointerRNA ptr;
 	RNA_id_pointer_create((ID*)builtin_data, &ptr);
 	BL::Image b_image(ptr);
 
-	int width = b_image.size()[0];
-	int height = b_image.size()[1];
-	int channels = b_image.channels();
+	const int width = b_image.size()[0];
+	const int height = b_image.size()[1];
+	const int channels = b_image.channels();
 
-	unsigned char *image_pixels;
-	image_pixels = image_get_pixels_for_frame(b_image, frame);
-	size_t num_pixels = ((size_t)width) * height;
+	unsigned char *image_pixels = image_get_pixels_for_frame(b_image, frame);
+	const size_t num_pixels = ((size_t)width) * height;
 
-	if(image_pixels) {
-		memcpy(pixels, image_pixels, num_pixels * channels * sizeof(unsigned char));
+	if(image_pixels && num_pixels * channels == pixels_size) {
+		memcpy(pixels, image_pixels, pixels_size * sizeof(unsigned char));
 		MEM_freeN(image_pixels);
 	}
 	else {
 		if(channels == 1) {
-			memset(pixels, 0, num_pixels * sizeof(unsigned char));
+			memset(pixels, 0, pixels_size * sizeof(unsigned char));
 		}
 		else {
+			const size_t num_pixels_safe = pixels_size / channels;
 			unsigned char *cp = pixels;
-			for(size_t i = 0; i < num_pixels; i++, cp += channels) {
+			for(size_t i = 0; i < num_pixels_safe; i++, cp += channels) {
 				cp[0] = 255;
 				cp[1] = 0;
 				cp[2] = 255;
-				if(channels == 4)
+				if(channels == 4) {
 					cp[3] = 255;
+				}
 			}
 		}
 	}
-
-	/* premultiply, byte images are always straight for blender */
+	/* Premultiply, byte images are always straight for Blender. */
 	unsigned char *cp = pixels;
 	for(size_t i = 0; i < num_pixels; i++, cp += channels) {
 		cp[0] = (cp[0] * cp[3]) >> 8;
 		cp[1] = (cp[1] * cp[3]) >> 8;
 		cp[2] = (cp[2] * cp[3]) >> 8;
 	}
-
 	return true;
 }
 
-bool BlenderSession::builtin_image_float_pixels(const string &builtin_name, void *builtin_data, float *pixels)
+bool BlenderSession::builtin_image_float_pixels(const string &builtin_name,
+                                                void *builtin_data,
+                                                float *pixels,
+                                                const size_t pixels_size)
 {
-	if(!builtin_data)
+	if(!builtin_data) {
 		return false;
+	}
 
 	PointerRNA ptr;
 	RNA_id_pointer_create((ID*)builtin_data, &ptr);
@@ -1222,16 +1214,16 @@ bool BlenderSession::builtin_image_float_pixels(const string &builtin_name, void
 		BL::Image b_image(b_id);
 		int frame = builtin_image_frame(builtin_name);
 
-		int width = b_image.size()[0];
-		int height = b_image.size()[1];
-		int channels = b_image.channels();
+		const int width = b_image.size()[0];
+		const int height = b_image.size()[1];
+		const int channels = b_image.channels();
 
 		float *image_pixels;
 		image_pixels = image_get_float_pixels_for_frame(b_image, frame);
-		size_t num_pixels = ((size_t)width) * height;
+		const size_t num_pixels = ((size_t)width) * height;
 
-		if(image_pixels) {
-			memcpy(pixels, image_pixels, num_pixels * channels * sizeof(float));
+		if(image_pixels && num_pixels * channels == pixels_size) {
+			memcpy(pixels, image_pixels, pixels_size * sizeof(float));
 			MEM_freeN(image_pixels);
 		}
 		else {
@@ -1239,13 +1231,15 @@ bool BlenderSession::builtin_image_float_pixels(const string &builtin_name, void
 				memset(pixels, 0, num_pixels * sizeof(float));
 			}
 			else {
+				const size_t num_pixels_safe = pixels_size / channels;
 				float *fp = pixels;
-				for(int i = 0; i < num_pixels; i++, fp += channels) {
+				for(int i = 0; i < num_pixels_safe; i++, fp += channels) {
 					fp[0] = 1.0f;
 					fp[1] = 0.0f;
 					fp[2] = 1.0f;
-					if(channels == 4)
+					if(channels == 4) {
 						fp[3] = 1.0f;
+					}
 				}
 			}
 		}
@@ -1257,8 +1251,9 @@ bool BlenderSession::builtin_image_float_pixels(const string &builtin_name, void
 		BL::Object b_ob(b_id);
 		BL::SmokeDomainSettings b_domain = object_smoke_domain_find(b_ob);
 
-		if(!b_domain)
+		if(!b_domain) {
 			return false;
+		}
 
 		int3 resolution = get_int3(b_domain.domain_resolution());
 		int length, amplify = (b_domain.use_high_resolution())? b_domain.amplify() + 1: 1;
@@ -1270,10 +1265,10 @@ bool BlenderSession::builtin_image_float_pixels(const string &builtin_name, void
 			amplify = 1;
 		}
 
-		int width = resolution.x * amplify;
-		int height = resolution.y * amplify;
-		int depth = resolution.z * amplify;
-		size_t num_pixels = ((size_t)width) * height * depth;
+		const int width = resolution.x * amplify;
+		const int height = resolution.y * amplify;
+		const int depth = resolution.z * amplify;
+		const size_t num_pixels = ((size_t)width) * height * depth;
 
 		if(builtin_name == Attribute::standard_name(ATTR_STD_VOLUME_DENSITY)) {
 			SmokeDomainSettings_density_grid_get_length(&b_domain.ptr, &length);

intern/cycles/blender/blender_session.h

@@ -145,9 +145,21 @@ protected:
 	void do_write_update_render_tile(RenderTile& rtile, bool do_update_only);
 
 	int builtin_image_frame(const string &builtin_name);
-	void builtin_image_info(const string &builtin_name, void *builtin_data, bool &is_float, int &width, int &height, int &depth, int &channels);
-	bool builtin_image_pixels(const string &builtin_name, void *builtin_data, unsigned char *pixels);
-	bool builtin_image_float_pixels(const string &builtin_name, void *builtin_data, float *pixels);
+	void builtin_image_info(const string &builtin_name,
+	                        void *builtin_data,
+	                        bool &is_float,
+	                        int &width,
+	                        int &height,
+	                        int &depth,
+	                        int &channels);
+	bool builtin_image_pixels(const string &builtin_name,
+	                          void *builtin_data,
+	                          unsigned char *pixels,
+	                          const size_t pixels_size);
+	bool builtin_image_float_pixels(const string &builtin_name,
+	                                void *builtin_data,
+	                                float *pixels,
+	                                const size_t pixels_size);
 
 	/* Update tile manager to reflect resumable render settings. */
 	void update_resumable_tile_manager(int num_samples);

intern/cycles/blender/blender_sync.cpp

@@ -56,6 +56,7 @@ BlenderSync::BlenderSync(BL::RenderEngine& b_engine,
   object_map(&scene->objects),
   mesh_map(&scene->meshes),
   light_map(&scene->lights),
+  particle_system_map(&scene->particle_systems),
   world_map(NULL),
   world_recalc(false),
   scene(scene),
@@ -143,6 +144,12 @@ bool BlenderSync::sync_recalc()
 			if(b_ob->is_updated_data() || b_ob->data().is_updated())
 				light_map.set_recalc(*b_ob);
 		}
+		
+		if(b_ob->is_updated_data()) {
+			BL::Object::particle_systems_iterator b_psys;
+			for(b_ob->particle_systems.begin(b_psys); b_psys != b_ob->particle_systems.end(); ++b_psys)
+				particle_system_map.set_recalc(*b_ob);
+		}
 	}
 
 	BL::BlendData::meshes_iterator b_mesh;
@@ -176,6 +183,7 @@ bool BlenderSync::sync_recalc()
 		object_map.has_recalc() ||
 		light_map.has_recalc() ||
 		mesh_map.has_recalc() ||
+		particle_system_map.has_recalc() ||
 		BlendDataObjects_is_updated_get(&b_data.ptr) ||
 		world_recalc;
 
@@ -490,12 +498,27 @@ SceneParams BlenderSync::get_scene_params(BL::Scene& b_scene,
 
 	params.use_bvh_spatial_split = RNA_boolean_get(&cscene, "debug_use_spatial_splits");
 	params.use_bvh_unaligned_nodes = RNA_boolean_get(&cscene, "debug_use_hair_bvh");
+	params.num_bvh_time_steps = RNA_int_get(&cscene, "debug_bvh_time_steps");
 
 	if(background && params.shadingsystem != SHADINGSYSTEM_OSL)
 		params.persistent_data = r.use_persistent_data();
 	else
 		params.persistent_data = false;
 
+	int texture_limit;
+	if(background) {
+		texture_limit = RNA_enum_get(&cscene, "texture_limit_render");
+	}
+	else {
+		texture_limit = RNA_enum_get(&cscene, "texture_limit");
+	}
+	if(texture_limit > 0 && b_scene.render().use_simplify()) {
+		params.texture_limit = 1 << (texture_limit + 6);
+	}
+	else {
+		params.texture_limit = 0;
+	}
+
 #if !(defined(__GNUC__) && (defined(i386) || defined(_M_IX86)))
 	if(is_cpu) {
 		params.use_qbvh = DebugFlags().cpu.qbvh && system_cpu_support_sse2();

intern/cycles/blender/blender_sync.h

@@ -35,6 +35,7 @@
 CCL_NAMESPACE_BEGIN
 
 class Background;
+class BlenderObjectCulling;
 class Camera;
 class Film;
 class Light;
@@ -122,8 +123,7 @@ private:
 	                    uint layer_flag,
 	                    float motion_time,
 	                    bool hide_tris,
-	                    bool use_camera_cull,
-	                    float camera_cull_margin,
+	                    BlenderObjectCulling& culling,
 	                    bool *use_portal);
 	void sync_light(BL::Object& b_parent,
 	                int persistent_id[OBJECT_PERSISTENT_ID_SIZE],
@@ -139,6 +139,11 @@ private:
 	                        int width, int height,
 	                        float motion_time);
 
+	/* particles */
+	bool sync_dupli_particle(BL::Object& b_ob,
+	                         BL::DupliObject& b_dup,
+	                         Object *object);
+
 	/* Images. */
 	void sync_images();
 
@@ -157,6 +162,7 @@ private:
 	id_map<ObjectKey, Object> object_map;
 	id_map<void*, Mesh> mesh_map;
 	id_map<ObjectKey, Light> light_map;
+	id_map<ParticleSystemKey, ParticleSystem> particle_system_map;
 	set<Mesh*> mesh_synced;
 	set<Mesh*> mesh_motion_synced;
 	set<float> motion_times;

intern/cycles/blender/blender_util.h

@@ -48,12 +48,12 @@ static inline BL::Mesh object_to_mesh(BL::BlendData& data,
                                       bool apply_modifiers,
                                       bool render,
                                       bool calc_undeformed,
-                                      bool subdivision)
+                                      Mesh::SubdivisionType subdivision_type)
 {
 	bool subsurf_mod_show_render;
 	bool subsurf_mod_show_viewport;
 
-	if(subdivision) {
+	if(subdivision_type != Mesh::SUBDIVISION_NONE) {
 		BL::Modifier subsurf_mod = object.modifiers[object.modifiers.length()-1];
 
 		subsurf_mod_show_render = subsurf_mod.show_render();
@@ -65,7 +65,7 @@ static inline BL::Mesh object_to_mesh(BL::BlendData& data,
 
 	BL::Mesh me = data.meshes.new_from_object(scene, object, apply_modifiers, (render)? 2: 1, false, calc_undeformed);
 
-	if(subdivision) {
+	if(subdivision_type != Mesh::SUBDIVISION_NONE) {
 		BL::Modifier subsurf_mod = object.modifiers[object.modifiers.length()-1];
 
 		subsurf_mod.show_render(subsurf_mod_show_render);
@@ -74,9 +74,14 @@ static inline BL::Mesh object_to_mesh(BL::BlendData& data,
 
 	if((bool)me) {
 		if(me.use_auto_smooth()) {
-			me.calc_normals_split();
+			if(subdivision_type == Mesh::SUBDIVISION_CATMULL_CLARK) {
+				me.calc_normals_split();
+			}
+			else {
+				me.split_faces();
+			}
 		}
-		if(!subdivision) {
+		if(subdivision_type == Mesh::SUBDIVISION_NONE) {
 			me.calc_tessface(true);
 		}
 	}
@@ -754,6 +759,33 @@ struct ObjectKey {
 	}
 };
 
+/* Particle System Key */
+
+struct ParticleSystemKey {
+	void *ob;
+	int id[OBJECT_PERSISTENT_ID_SIZE];
+
+	ParticleSystemKey(void *ob_, int id_[OBJECT_PERSISTENT_ID_SIZE])
+	: ob(ob_)
+	{
+		if(id_)
+			memcpy(id, id_, sizeof(id));
+		else
+			memset(id, 0, sizeof(id));
+	}
+
+	bool operator<(const ParticleSystemKey& k) const
+	{
+		/* first id is particle index, we don't compare that */
+		if(ob < k.ob)
+			return true;
+		else if(ob == k.ob)
+			return memcmp(id+1, k.id+1, sizeof(int)*(OBJECT_PERSISTENT_ID_SIZE-1)) < 0;
+
+		return false;
+	}
+};
+
 CCL_NAMESPACE_END
 
 #endif /* __BLENDER_UTIL_H__ */

intern/cycles/bvh/bvh.cpp

@@ -845,6 +845,8 @@ void QBVH::pack_aligned_inner(const BVHStackEntry& e,
 	                  bounds,
 	                  child,
 	                  e.node->m_visibility,
+	                  e.node->m_time_from,
+	                  e.node->m_time_to,
 	                  num);
 }
 
@@ -852,12 +854,17 @@ void QBVH::pack_aligned_node(int idx,
                              const BoundBox *bounds,
                              const int *child,
                              const uint visibility,
+                             const float time_from,
+                             const float time_to,
                              const int num)
 {
 	float4 data[BVH_QNODE_SIZE];
 	memset(data, 0, sizeof(data));
 
 	data[0].x = __uint_as_float(visibility & ~PATH_RAY_NODE_UNALIGNED);
+	data[0].y = time_from;
+	data[0].z = time_to;
+
 	for(int i = 0; i < num; i++) {
 		float3 bb_min = bounds[i].min;
 		float3 bb_max = bounds[i].max;
@@ -908,6 +915,8 @@ void QBVH::pack_unaligned_inner(const BVHStackEntry& e,
 	                    bounds,
 	                    child,
 	                    e.node->m_visibility,
+	                    e.node->m_time_from,
+	                    e.node->m_time_to,
 	                    num);
 }
 
@@ -916,12 +925,16 @@ void QBVH::pack_unaligned_node(int idx,
                                const BoundBox *bounds,
                                const int *child,
                                const uint visibility,
+                               const float time_from,
+                               const float time_to,
                                const int num)
 {
 	float4 data[BVH_UNALIGNED_QNODE_SIZE];
 	memset(data, 0, sizeof(data));
 
 	data[0].x = __uint_as_float(visibility | PATH_RAY_NODE_UNALIGNED);
+	data[0].y = time_from;
+	data[0].z = time_to;
 
 	for(int i = 0; i < num; i++) {
 		Transform space = BVHUnaligned::compute_node_transform(
@@ -1207,6 +1220,8 @@ void QBVH::refit_node(int idx, bool leaf, BoundBox& bbox, uint& visibility)
 			                    child_bbox,
 			                    &c[0],
 			                    visibility,
+			                    0.0f,
+			                    1.0f,
 			                    4);
 		}
 		else {
@@ -1214,6 +1229,8 @@ void QBVH::refit_node(int idx, bool leaf, BoundBox& bbox, uint& visibility)
 			                  child_bbox,
 			                  &c[0],
 			                  visibility,
+			                  0.0f,
+			                  1.0f,
 			                  4);
 		}
 	}

intern/cycles/bvh/bvh.h

@@ -175,6 +175,8 @@ protected:
 	                       const BoundBox *bounds,
 	                       const int *child,
 	                       const uint visibility,
+	                       const float time_from,
+	                       const float time_to,
 	                       const int num);
 
 	void pack_unaligned_inner(const BVHStackEntry& e,
@@ -185,6 +187,8 @@ protected:
 	                         const BoundBox *bounds,
 	                         const int *child,
 	                         const uint visibility,
+	                         const float time_from,
+	                         const float time_to,
 	                         const int num);
 
 	/* refit */

intern/cycles/bvh/bvh_build.cpp

@@ -112,81 +112,237 @@ BVHBuild::~BVHBuild()
 
 /* Adding References */
 
-void BVHBuild::add_reference_mesh(BoundBox& root, BoundBox& center, Mesh *mesh, int i)
+void BVHBuild::add_reference_triangles(BoundBox& root, BoundBox& center, Mesh *mesh, int i)
 {
-	if(params.primitive_mask & PRIMITIVE_ALL_TRIANGLE) {
-		Attribute *attr_mP = NULL;
-
-		if(mesh->has_motion_blur())
-			attr_mP = mesh->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
-
-		size_t num_triangles = mesh->num_triangles();
-		for(uint j = 0; j < num_triangles; j++) {
-			Mesh::Triangle t = mesh->get_triangle(j);
+	const Attribute *attr_mP = NULL;
+	if(mesh->has_motion_blur()) {
+		attr_mP = mesh->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+	}
+	const size_t num_triangles = mesh->num_triangles();
+	for(uint j = 0; j < num_triangles; j++) {
+		Mesh::Triangle t = mesh->get_triangle(j);
+		const float3 *verts = &mesh->verts[0];
+		if(attr_mP == NULL) {
 			BoundBox bounds = BoundBox::empty;
-			PrimitiveType type = PRIMITIVE_TRIANGLE;
-
-			t.bounds_grow(&mesh->verts[0], bounds);
-
-			/* motion triangles */
-			if(attr_mP) {
-				size_t mesh_size = mesh->verts.size();
-				size_t steps = mesh->motion_steps - 1;
-				float3 *vert_steps = attr_mP->data_float3();
-
-				for(size_t i = 0; i < steps; i++)
-					t.bounds_grow(vert_steps + i*mesh_size, bounds);
-
-				type = PRIMITIVE_MOTION_TRIANGLE;
-			}
-
+			t.bounds_grow(verts, bounds);
 			if(bounds.valid()) {
-				references.push_back(BVHReference(bounds, j, i, type));
+				references.push_back(BVHReference(bounds,
+				                                  j,
+				                                  i,
+				                                  PRIMITIVE_TRIANGLE));
 				root.grow(bounds);
 				center.grow(bounds.center2());
 			}
 		}
-	}
-
-	if(params.primitive_mask & PRIMITIVE_ALL_CURVE) {
-		Attribute *curve_attr_mP = NULL;
-
-		if(mesh->has_motion_blur())
-			curve_attr_mP = mesh->curve_attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
-
-		size_t num_curves = mesh->num_curves();
-		for(uint j = 0; j < num_curves; j++) {
-			Mesh::Curve curve = mesh->get_curve(j);
-			PrimitiveType type = PRIMITIVE_CURVE;
-
-			for(int k = 0; k < curve.num_keys - 1; k++) {
-				BoundBox bounds = BoundBox::empty;
-				curve.bounds_grow(k, &mesh->curve_keys[0], &mesh->curve_radius[0], bounds);
-
-				/* motion curve */
-				if(curve_attr_mP) {
-					size_t mesh_size = mesh->curve_keys.size();
-					size_t steps = mesh->motion_steps - 1;
-					float3 *key_steps = curve_attr_mP->data_float3();
-
-					for(size_t i = 0; i < steps; i++)
-						curve.bounds_grow(k, key_steps + i*mesh_size, &mesh->curve_radius[0], bounds);
-
-					type = PRIMITIVE_MOTION_CURVE;
-				}
-
+		else if(params.num_motion_triangle_steps == 0 || params.use_spatial_split) {
+			/* Motion triangles, simple case: single node for the whole
+			 * primitive. Lowest memory footprint and faster BVH build but
+			 * least optimal ray-tracing.
+			 */
+			/* TODO(sergey): Support motion steps for spatially split BVH. */
+			const size_t num_verts = mesh->verts.size();
+			const size_t num_steps = mesh->motion_steps;
+			const float3 *vert_steps = attr_mP->data_float3();
+			BoundBox bounds = BoundBox::empty;
+			t.bounds_grow(verts, bounds);
+			for(size_t step = 0; step < num_steps - 1; step++) {
+				t.bounds_grow(vert_steps + step*num_verts, bounds);
+			}
+			if(bounds.valid()) {
+				references.push_back(
+				        BVHReference(bounds,
+				                     j,
+				                     i,
+				                     PRIMITIVE_MOTION_TRIANGLE));
+				root.grow(bounds);
+				center.grow(bounds.center2());
+			}
+		}
+		else {
+			/* Motion triangles, trace optimized case:  we split triangle
+			 * primitives into separate nodes for each of the time steps.
+			 * This way we minimize overlap of neighbor curve primitives.
+			 */
+			const int num_bvh_steps = params.num_motion_curve_steps * 2 + 1;
+			const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1);
+			const size_t num_verts = mesh->verts.size();
+			const size_t num_steps = mesh->motion_steps;
+			const float3 *vert_steps = attr_mP->data_float3();
+			/* Calculate bounding box of the previous time step.
+			 * Will be reused later to avoid duplicated work on
+			 * calculating BVH time step boundbox.
+			 */
+			float3 prev_verts[3];
+			t.motion_verts(verts,
+			               vert_steps,
+			               num_verts,
+			               num_steps,
+			               0.0f,
+			               prev_verts);
+			BoundBox prev_bounds = BoundBox::empty;
+			prev_bounds.grow(prev_verts[0]);
+			prev_bounds.grow(prev_verts[1]);
+			prev_bounds.grow(prev_verts[2]);
+			/* Create all primitive time steps, */
+			for(int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
+				const float curr_time = (float)(bvh_step) * num_bvh_steps_inv_1;
+				float3 curr_verts[3];
+				t.motion_verts(verts,
+				               vert_steps,
+				               num_verts,
+				               num_steps,
+				               curr_time,
+				               curr_verts);
+				BoundBox curr_bounds = BoundBox::empty;
+				curr_bounds.grow(curr_verts[0]);
+				curr_bounds.grow(curr_verts[1]);
+				curr_bounds.grow(curr_verts[2]);
+				BoundBox bounds = prev_bounds;
+				bounds.grow(curr_bounds);
 				if(bounds.valid()) {
-					int packed_type = PRIMITIVE_PACK_SEGMENT(type, k);
+					const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1;
+					references.push_back(
+					        BVHReference(bounds,
+					                     j,
+					                     i,
+					                     PRIMITIVE_MOTION_TRIANGLE,
+					                     prev_time,
+					                     curr_time));
+					root.grow(bounds);
+					center.grow(bounds.center2());
+				}
+				/* Current time boundbox becomes previous one for the
+				 * next time step.
+				 */
+				prev_bounds = curr_bounds;
+			}
+		}
+	}
+}
 
+void BVHBuild::add_reference_curves(BoundBox& root, BoundBox& center, Mesh *mesh, int i)
+{
+	const Attribute *curve_attr_mP = NULL;
+	if(mesh->has_motion_blur()) {
+		curve_attr_mP = mesh->curve_attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+	}
+	const size_t num_curves = mesh->num_curves();
+	for(uint j = 0; j < num_curves; j++) {
+		const Mesh::Curve curve = mesh->get_curve(j);
+		const float *curve_radius = &mesh->curve_radius[0];
+		for(int k = 0; k < curve.num_keys - 1; k++) {
+			if(curve_attr_mP == NULL) {
+				/* Really simple logic for static hair. */
+				BoundBox bounds = BoundBox::empty;
+				curve.bounds_grow(k, &mesh->curve_keys[0], curve_radius, bounds);
+				if(bounds.valid()) {
+					int packed_type = PRIMITIVE_PACK_SEGMENT(PRIMITIVE_CURVE, k);
 					references.push_back(BVHReference(bounds, j, i, packed_type));
 					root.grow(bounds);
 					center.grow(bounds.center2());
 				}
 			}
+			else if(params.num_motion_curve_steps == 0 || params.use_spatial_split) {
+				/* Simple case of motion curves: single node for the while
+				 * shutter time. Lowest memory usage but less optimal
+				 * rendering.
+				 */
+				/* TODO(sergey): Support motion steps for spatially split BVH. */
+				BoundBox bounds = BoundBox::empty;
+				curve.bounds_grow(k, &mesh->curve_keys[0], curve_radius, bounds);
+				const size_t num_keys = mesh->curve_keys.size();
+				const size_t num_steps = mesh->motion_steps;
+				const float3 *key_steps = curve_attr_mP->data_float3();
+				for(size_t step = 0; step < num_steps - 1; step++) {
+					curve.bounds_grow(k,
+					                  key_steps + step*num_keys,
+					                  curve_radius,
+					                  bounds);
+				}
+				if(bounds.valid()) {
+					int packed_type = PRIMITIVE_PACK_SEGMENT(PRIMITIVE_MOTION_CURVE, k);
+					references.push_back(BVHReference(bounds,
+					                                  j,
+					                                  i,
+					                                  packed_type));
+					root.grow(bounds);
+					center.grow(bounds.center2());
+				}
+			}
+			else {
+				/* Motion curves, trace optimized case:  we split curve keys
+				 * primitives into separate nodes for each of the time steps.
+				 * This way we minimize overlap of neighbor curve primitives.
+				 */
+				const int num_bvh_steps = params.num_motion_curve_steps * 2 + 1;
+				const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1);
+				const size_t num_steps = mesh->motion_steps;
+				const float3 *curve_keys = &mesh->curve_keys[0];
+				const float3 *key_steps = curve_attr_mP->data_float3();
+				const size_t num_keys = mesh->curve_keys.size();
+				/* Calculate bounding box of the previous time step.
+				 * Will be reused later to avoid duplicated work on
+				 * calculating BVH time step boundbox.
+				 */
+				float4 prev_keys[4];
+				curve.cardinal_motion_keys(curve_keys,
+				                           curve_radius,
+				                           key_steps,
+				                           num_keys,
+				                           num_steps,
+				                           0.0f,
+				                           k - 1, k, k + 1, k + 2,
+				                           prev_keys);
+				BoundBox prev_bounds = BoundBox::empty;
+				curve.bounds_grow(prev_keys, prev_bounds);
+				/* Create all primitive time steps, */
+				for(int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
+					const float curr_time = (float)(bvh_step) * num_bvh_steps_inv_1;
+					float4 curr_keys[4];
+					curve.cardinal_motion_keys(curve_keys,
+					                           curve_radius,
+					                           key_steps,
+					                           num_keys,
+					                           num_steps,
+					                           curr_time,
+					                           k - 1, k, k + 1, k + 2,
+					                           curr_keys);
+					BoundBox curr_bounds = BoundBox::empty;
+					curve.bounds_grow(curr_keys, curr_bounds);
+					BoundBox bounds = prev_bounds;
+					bounds.grow(curr_bounds);
+					if(bounds.valid()) {
+						const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1;
+						int packed_type = PRIMITIVE_PACK_SEGMENT(PRIMITIVE_MOTION_CURVE, k);
+						references.push_back(BVHReference(bounds,
+						                                  j,
+						                                  i,
+						                                  packed_type,
+						                                  prev_time,
+						                                  curr_time));
+						root.grow(bounds);
+						center.grow(bounds.center2());
+					}
+					/* Current time boundbox becomes previous one for the
+					 * next time step.
+					 */
+					prev_bounds = curr_bounds;
+				}
+			}
 		}
 	}
 }
 
+void BVHBuild::add_reference_mesh(BoundBox& root, BoundBox& center, Mesh *mesh, int i)
+{
+	if(params.primitive_mask & PRIMITIVE_ALL_TRIANGLE) {
+		add_reference_triangles(root, center, mesh, i);
+	}
+	if(params.primitive_mask & PRIMITIVE_ALL_CURVE) {
+		add_reference_curves(root, center, mesh, i);
+	}
+}
+
 void BVHBuild::add_reference_object(BoundBox& root, BoundBox& center, Object *ob, int i)
 {
 	references.push_back(BVHReference(ob->bounds, -1, i, 0));
@@ -200,7 +356,7 @@ static size_t count_curve_segments(Mesh *mesh)
 
 	for(size_t i = 0; i < num_curves; i++)
 		num += mesh->get_curve(i).num_keys - 1;
-	
+
 	return num;
 }
 
@@ -344,6 +500,7 @@ BVHNode* BVHBuild::run()
 		else {
 			/*rotate(rootnode, 4, 5);*/
 			rootnode->update_visibility();
+			rootnode->update_time();
 		}
 		if(rootnode != NULL) {
 			VLOG(1) << "BVH build statistics:\n"
@@ -371,7 +528,7 @@ void BVHBuild::progress_update()
 {
 	if(time_dt() - progress_start_time < 0.25)
 		return;
-	
+
 	double progress_start = (double)progress_count/(double)progress_total;
 	double duplicates = (double)(progress_total - progress_original_total)/(double)progress_total;
 
@@ -379,7 +536,7 @@ void BVHBuild::progress_update()
 	                           progress_start * 100.0, duplicates * 100.0);
 
 	progress.set_substatus(msg);
-	progress_start_time = time_dt(); 
+	progress_start_time = time_dt();
 }
 
 void BVHBuild::thread_build_node(InnerNode *inner,
@@ -435,6 +592,7 @@ bool BVHBuild::range_within_max_leaf_size(const BVHRange& range,
 		return false;
 
 	size_t num_triangles = 0;
+	size_t num_motion_triangles = 0;
 	size_t num_curves = 0;
 	size_t num_motion_curves = 0;
 
@@ -445,13 +603,16 @@ bool BVHBuild::range_within_max_leaf_size(const BVHRange& range,
 			num_curves++;
 		if(ref.prim_type() & PRIMITIVE_MOTION_CURVE)
 			num_motion_curves++;
-		else if(ref.prim_type() & PRIMITIVE_ALL_TRIANGLE)
+		else if(ref.prim_type() & PRIMITIVE_TRIANGLE)
 			num_triangles++;
+		else if(ref.prim_type() & PRIMITIVE_MOTION_TRIANGLE)
+			num_motion_triangles++;
 	}
 
-	return (num_triangles < params.max_triangle_leaf_size) &&
-	       (num_curves < params.max_curve_leaf_size) &&
-	       (num_motion_curves < params.max_curve_leaf_size);
+	return (num_triangles <= params.max_triangle_leaf_size) &&
+	       (num_motion_triangles <= params.max_motion_triangle_leaf_size) &&
+	       (num_curves <= params.max_curve_leaf_size) &&
+	       (num_motion_curves <= params.max_motion_curve_leaf_size);
 }
 
 /* multithreaded binning builder */
@@ -689,18 +850,24 @@ BVHNode *BVHBuild::create_object_leaf_nodes(const BVHReference *ref, int start,
 		prim_object[start] = ref->prim_object();
 
 		uint visibility = objects[ref->prim_object()]->visibility;
-		return new LeafNode(ref->bounds(), visibility, start, start+1);
+		BVHNode *leaf_node =  new LeafNode(ref->bounds(), visibility, start, start+1);
+		leaf_node->m_time_from = ref->time_from();
+		leaf_node->m_time_to = ref->time_to();
+		return leaf_node;
 	}
 	else {
 		int mid = num/2;
-		BVHNode *leaf0 = create_object_leaf_nodes(ref, start, mid); 
-		BVHNode *leaf1 = create_object_leaf_nodes(ref+mid, start+mid, num-mid); 
+		BVHNode *leaf0 = create_object_leaf_nodes(ref, start, mid);
+		BVHNode *leaf1 = create_object_leaf_nodes(ref+mid, start+mid, num-mid);
 
 		BoundBox bounds = BoundBox::empty;
 		bounds.grow(leaf0->m_bounds);
 		bounds.grow(leaf1->m_bounds);
 
-		return new InnerNode(bounds, leaf0, leaf1);
+		BVHNode *inner_node = new InnerNode(bounds, leaf0, leaf1);
+		inner_node->m_time_from = min(leaf0->m_time_from, leaf1->m_time_from);
+		inner_node->m_time_to = max(leaf0->m_time_to, leaf1->m_time_to);
+		return inner_node;
 	}
 }
 
@@ -804,6 +971,16 @@ BVHNode* BVHBuild::create_leaf_node(const BVHRange& range,
 			                                   visibility[i],
 			                                   start_index,
 			                                   start_index + num);
+			if(true) {
+				float time_from = 1.0f, time_to = 0.0f;
+				for(int j = 0; j < num; ++j) {
+					const BVHReference &ref = p_ref[i][j];
+					time_from = min(time_from, ref.time_from());
+					time_to = max(time_to, ref.time_to());
+				}
+				leaf_node->m_time_from = time_from;
+				leaf_node->m_time_to = time_to;
+			}
 			if(alignment_found) {
 				/* Need to recalculate leaf bounds with new alignment. */
 				leaf_node->m_bounds = BoundBox::empty;
@@ -918,7 +1095,7 @@ BVHNode* BVHBuild::create_leaf_node(const BVHRange& range,
 		BVHNode *inner = new InnerNode(inner_bounds, leaves[1], leaves[2]);
 		return new InnerNode(range.bounds(), leaves[0], inner);
 	} else {
-		/* Shpuld be doing more branches if more primitive types added. */
+		/* Should be doing more branches if more primitive types added. */
 		assert(num_leaves <= 5);
 		BoundBox inner_bounds_a = merge(leaves[0]->m_bounds, leaves[1]->m_bounds);
 		BoundBox inner_bounds_b = merge(leaves[2]->m_bounds, leaves[3]->m_bounds);
@@ -951,7 +1128,7 @@ void BVHBuild::rotate(BVHNode *node, int max_depth)
 	/* nothing to rotate if we reached a leaf node. */
 	if(node->is_leaf() || max_depth < 0)
 		return;
-	
+
 	InnerNode *parent = (InnerNode*)node;
 
 	/* rotate all children first */

intern/cycles/bvh/bvh_build.h

@@ -63,6 +63,8 @@ protected:
 	friend class BVHObjectBinning;
 
 	/* Adding references. */
+	void add_reference_triangles(BoundBox& root, BoundBox& center, Mesh *mesh, int i);
+	void add_reference_curves(BoundBox& root, BoundBox& center, Mesh *mesh, int i);
 	void add_reference_mesh(BoundBox& root, BoundBox& center, Mesh *mesh, int i);
 	void add_reference_object(BoundBox& root, BoundBox& center, Object *ob, int i);
 	void add_references(BVHRange& root);

intern/cycles/bvh/bvh_node.cpp

@@ -176,6 +176,19 @@ uint BVHNode::update_visibility()
 	return m_visibility;
 }
 
+void BVHNode::update_time()
+{
+	if(!is_leaf()) {
+		InnerNode *inner = (InnerNode*)this;
+		BVHNode *child0 = inner->children[0];
+		BVHNode *child1 = inner->children[1];
+		child0->update_time();
+		child1->update_time();
+		m_time_from = min(child0->m_time_from, child1->m_time_from);
+		m_time_to =  max(child0->m_time_to, child1->m_time_to);
+	}
+}
+
 /* Inner Node */
 
 void InnerNode::print(int depth) const

intern/cycles/bvh/bvh_node.h

@@ -47,7 +47,9 @@ class BVHNode
 {
 public:
 	BVHNode() : m_is_unaligned(false),
-	            m_aligned_space(NULL)
+	            m_aligned_space(NULL),
+	            m_time_from(0.0f),
+	            m_time_to(1.0f)
 	{
 	}
 
@@ -91,12 +93,15 @@ public:
 	void deleteSubtree();
 
 	uint update_visibility();
+	void update_time();
 
 	bool m_is_unaligned;
 
 	// TODO(sergey): Can be stored as 3x3 matrix, but better to have some
 	// utilities and type defines in util_transform first.
 	Transform *m_aligned_space;
+
+	float m_time_from, m_time_to;
 };
 
 class InnerNode : public BVHNode

intern/cycles/bvh/bvh_params.h

@@ -43,7 +43,9 @@ public:
 	/* number of primitives in leaf */
 	int min_leaf_size;
 	int max_triangle_leaf_size;
+	int max_motion_triangle_leaf_size;
 	int max_curve_leaf_size;
+	int max_motion_curve_leaf_size;
 
 	/* object or mesh level bvh */
 	bool top_level;
@@ -59,6 +61,17 @@ public:
 	 */
 	bool use_unaligned_nodes;
 
+	/* Split time range to this number of steps and create leaf node for each
+	 * of this time steps.
+	 *
+	 * Speeds up rendering of motion curve primitives in the cost of higher
+	 * memory usage.
+	 */
+	int num_motion_curve_steps;
+
+	/* Same as above, but for triangle primitives. */
+	int num_motion_triangle_steps;
+
 	/* fixed parameters */
 	enum {
 		MAX_DEPTH = 64,
@@ -80,13 +93,17 @@ public:
 
 		min_leaf_size = 1;
 		max_triangle_leaf_size = 8;
-		max_curve_leaf_size = 2;
+		max_motion_triangle_leaf_size = 8;
+		max_curve_leaf_size = 1;
+		max_motion_curve_leaf_size = 4;
 
 		top_level = false;
 		use_qbvh = false;
 		use_unaligned_nodes = false;
 
 		primitive_mask = PRIMITIVE_ALL;
+
+		num_motion_curve_steps = 0;
 	}
 
 	/* SAH costs */
@@ -113,8 +130,15 @@ class BVHReference
 public:
 	__forceinline BVHReference() {}
 
-	__forceinline BVHReference(const BoundBox& bounds_, int prim_index_, int prim_object_, int prim_type)
-	: rbounds(bounds_)
+	__forceinline BVHReference(const BoundBox& bounds_,
+	                           int prim_index_,
+	                           int prim_object_,
+	                           int prim_type,
+	                           float time_from = 0.0f,
+	                           float time_to = 1.0f)
+	        : rbounds(bounds_),
+	          time_from_(time_from),
+	          time_to_(time_to)
 	{
 		rbounds.min.w = __int_as_float(prim_index_);
 		rbounds.max.w = __int_as_float(prim_object_);
@@ -125,6 +149,9 @@ public:
 	__forceinline int prim_index() const { return __float_as_int(rbounds.min.w); }
 	__forceinline int prim_object() const { return __float_as_int(rbounds.max.w); }
 	__forceinline int prim_type() const { return type; }
+	__forceinline float time_from() const { return time_from_; }
+	__forceinline float time_to() const { return time_to_; }
+
 
 	BVHReference& operator=(const BVHReference &arg) {
 		if(&arg != this) {
@@ -133,9 +160,11 @@ public:
 		return *this;
 	}
 
+
 protected:
 	BoundBox rbounds;
 	uint type;
+	float time_from_, time_to_;
 };
 
 /* BVH Range

intern/cycles/device/device.cpp

@@ -64,6 +64,8 @@ std::ostream& operator <<(std::ostream &os,
 	   << string_from_bool(requested_features.use_integrator_branched) << std::endl;
 	os << "Use Patch Evaluation: "
 	   << string_from_bool(requested_features.use_patch_evaluation) << std::endl;
+	os << "Use Transparent Shadows: "
+	   << string_from_bool(requested_features.use_transparent) << std::endl;
 	return os;
 }
 

intern/cycles/device/device.h

@@ -117,6 +117,9 @@ public:
 
 	/* Use OpenSubdiv patch evaluation */
 	bool use_patch_evaluation;
+	
+	/* Use Transparent shadows */
+	bool use_transparent;
 
 	DeviceRequestedFeatures()
 	{
@@ -133,6 +136,7 @@ public:
 		use_volume = false;
 		use_integrator_branched = false;
 		use_patch_evaluation = false;
+		use_transparent = false;
 	}
 
 	bool modified(const DeviceRequestedFeatures& requested_features)
@@ -148,7 +152,8 @@ public:
 		         use_subsurface == requested_features.use_subsurface &&
 		         use_volume == requested_features.use_volume &&
 		         use_integrator_branched == requested_features.use_integrator_branched &&
-		         use_patch_evaluation == requested_features.use_patch_evaluation);
+		         use_patch_evaluation == requested_features.use_patch_evaluation &&
+		         use_transparent == requested_features.use_transparent);
 	}
 
 	/* Convert the requested features structure to a build options,
@@ -189,6 +194,9 @@ public:
 		if(!use_patch_evaluation) {
 			build_options += " -D__NO_PATCH_EVAL__";
 		}
+		if(!use_transparent) {
+			build_options += " -D__NO_TRANSPARENT__";
+		}
 		return build_options;
 	}
 };
@@ -220,6 +228,7 @@ public:
 	DeviceInfo info;
 	virtual const string& error_message() { return error_msg; }
 	bool have_error() { return !error_message().empty(); }
+	virtual bool show_samples() const { return false; }
 
 	/* statistics */
 	Stats &stats;

intern/cycles/device/device_cpu.cpp

@@ -112,6 +112,11 @@ public:
 		task_pool.stop();
 	}
 
+	virtual bool show_samples() const
+	{
+		return (TaskScheduler::num_threads() == 1);
+	}
+
 	void mem_alloc(device_memory& mem, MemoryType /*type*/)
 	{
 		mem.device_pointer = mem.data_pointer;
@@ -275,7 +280,7 @@ public:
 
 				tile.sample = sample + 1;
 
-				task.update_progress(&tile);
+				task.update_progress(&tile, tile.w*tile.h);
 			}
 
 			task.release_tile(tile);

intern/cycles/device/device_cuda.cpp

@@ -115,6 +115,12 @@ public:
 		return path_exists(cubins_path);
 	}
 
+	virtual bool show_samples() const
+	{
+		/* The CUDADevice only processes one tile at a time, so showing samples is fine. */
+		return true;
+	}
+
 /*#ifdef NDEBUG
 #define cuda_abort()
 #else
@@ -124,7 +130,7 @@ public:
 	{
 		if(first_error) {
 			fprintf(stderr, "\nRefer to the Cycles GPU rendering documentation for possible solutions:\n");
-			fprintf(stderr, "http://www.blender.org/manual/render/cycles/gpu_rendering.html\n\n");
+			fprintf(stderr, "https://docs.blender.org/manual/en/dev/render/cycles/gpu_rendering.html\n\n");
 			first_error = false;
 		}
 	}
@@ -1267,7 +1273,7 @@ public:
 
 					tile.sample = sample + 1;
 
-					task->update_progress(&tile);
+					task->update_progress(&tile, tile.w*tile.h);
 				}
 
 				task->release_tile(tile);
@@ -1418,7 +1424,11 @@ void device_cuda_info(vector<DeviceInfo>& devices)
 		cuDeviceGetAttribute(&pci_location[0], CU_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID, num);
 		cuDeviceGetAttribute(&pci_location[1], CU_DEVICE_ATTRIBUTE_PCI_BUS_ID, num);
 		cuDeviceGetAttribute(&pci_location[2], CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID, num);
-		info.id = string_printf("CUDA_%s_%04x:%02x:%02x", name, pci_location[0], pci_location[1], pci_location[2]);
+		info.id = string_printf("CUDA_%s_%04x:%02x:%02x",
+		                        name,
+		                        (unsigned int)pci_location[0],
+		                        (unsigned int)pci_location[1],
+		                        (unsigned int)pci_location[2]);
 
 		/* if device has a kernel timeout, assume it is used for display */
 		if(cuDeviceGetAttribute(&attr, CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT, num) == CUDA_SUCCESS && attr == 1) {

intern/cycles/device/device_multi.cpp

@@ -89,6 +89,14 @@ public:
 		return error_msg;
 	}
 
+	virtual bool show_samples() const
+	{
+		if(devices.size() > 1) {
+			return false;
+		}
+		return devices.front().device->show_samples();
+	}
+
 	bool load_kernels(const DeviceRequestedFeatures& requested_features)
 	{
 		foreach(SubDevice& sub, devices)

intern/cycles/device/device_network.cpp

@@ -51,6 +51,11 @@ public:
 
 	thread_mutex rpc_lock;
 
+	virtual bool show_samples() const
+	{
+		return false;
+	}
+
 	NetworkDevice(DeviceInfo& info, Stats &stats, const char *address)
 	: Device(info, stats, true), socket(io_service)
 	{

intern/cycles/device/device_task.cpp

@@ -19,6 +19,8 @@
 
 #include "device_task.h"
 
+#include "buffers.h"
+
 #include "util_algorithm.h"
 #include "util_time.h"
 
@@ -99,14 +101,18 @@ void DeviceTask::split(list<DeviceTask>& tasks, int num, int max_size)
 	}
 }
 
-void DeviceTask::update_progress(RenderTile *rtile)
+void DeviceTask::update_progress(RenderTile *rtile, int pixel_samples)
 {
 	if((type != PATH_TRACE) &&
 	   (type != SHADER))
 		return;
 
-	if(update_progress_sample)
-		update_progress_sample();
+	if(update_progress_sample) {
+		if(pixel_samples == -1) {
+			pixel_samples = shader_w;
+		}
+		update_progress_sample(pixel_samples, rtile? rtile->sample : 0);
+	}
 
 	if(update_tile_sample) {
 		double current_time = time_dt();

intern/cycles/device/device_task.h

@@ -56,10 +56,10 @@ public:
 	int get_subtask_count(int num, int max_size = 0);
 	void split(list<DeviceTask>& tasks, int num, int max_size = 0);
 
-	void update_progress(RenderTile *rtile);
+	void update_progress(RenderTile *rtile, int pixel_samples = -1);
 
 	function<bool(Device *device, RenderTile&)> acquire_tile;
-	function<void(void)> update_progress_sample;
+	function<void(long, int)> update_progress_sample;
 	function<void(RenderTile&)> update_tile_sample;
 	function<void(RenderTile&)> release_tile;
 	function<bool(void)> get_cancel;

intern/cycles/device/opencl/opencl.h

@@ -16,14 +16,14 @@
 
 #ifdef WITH_OPENCL
 
-#include "clew.h"
-
 #include "device.h"
 
 #include "util_map.h"
 #include "util_param.h"
 #include "util_string.h"
 
+#include "clew.h"
+
 CCL_NAMESPACE_BEGIN
 
 #define CL_MEM_PTR(p) ((cl_mem)(uintptr_t)(p))

intern/cycles/device/opencl/opencl_mega.cpp

@@ -39,6 +39,10 @@ public:
 	{
 	}
 
+	virtual bool show_samples() const {
+		return true;
+	}
+
 	virtual void load_kernels(const DeviceRequestedFeatures& /*requested_features*/,
 	                          vector<OpenCLProgram*> &programs)
 	{
@@ -120,7 +124,7 @@ public:
 
 					tile.sample = sample + 1;
 
-					task->update_progress(&tile);
+					task->update_progress(&tile, tile.w*tile.h);
 				}
 
 				/* Complete kernel execution before release tile */

intern/cycles/device/opencl/opencl_split.cpp

@@ -247,6 +247,10 @@ public:
 		}
 	}
 
+	virtual bool show_samples() const {
+		return false;
+	}
+
 	/* Split kernel utility functions. */
 	size_t get_tex_size(const char *tex_name)
 	{

intern/cycles/device/opencl/opencl_util.cpp

@@ -667,7 +667,10 @@ string OpenCLInfo::get_hardware_id(string platform_name, cl_device_id device_id)
 		/* Use cl_amd_device_topology extension. */
 		cl_char topology[24];
 		if(clGetDeviceInfo(device_id, 0x4037, sizeof(topology), topology, NULL) == CL_SUCCESS && topology[0] == 1) {
-			return string_printf("%02x:%02x.%01x", topology[21], topology[22], topology[23]);
+			return string_printf("%02x:%02x.%01x",
+			                     (unsigned int)topology[21],
+			                     (unsigned int)topology[22],
+			                     (unsigned int)topology[23]);
 		}
 	}
 	else if(platform_name == "NVIDIA CUDA") {
@@ -675,7 +678,10 @@ string OpenCLInfo::get_hardware_id(string platform_name, cl_device_id device_id)
 		cl_int bus_id, slot_id;
 		if(clGetDeviceInfo(device_id, 0x4008, sizeof(cl_int), &bus_id,  NULL) == CL_SUCCESS &&
 		   clGetDeviceInfo(device_id, 0x4009, sizeof(cl_int), &slot_id, NULL) == CL_SUCCESS) {
-			return string_printf("%02x:%02x.%01x", bus_id, slot_id>>3, slot_id & 0x7);
+			return string_printf("%02x:%02x.%01x",
+			                     (unsigned int)(bus_id),
+			                     (unsigned int)(slot_id >> 3),
+			                     (unsigned int)(slot_id & 0x7));
 		}
 	}
 	/* No general way to get a hardware ID from OpenCL => give up. */

intern/cycles/kernel/CMakeLists.txt

@@ -164,6 +164,8 @@ set(SRC_GEOM_HEADERS
 	geom/geom_curve.h
 	geom/geom_motion_curve.h
 	geom/geom_motion_triangle.h
+	geom/geom_motion_triangle_intersect.h
+	geom/geom_motion_triangle_shader.h
 	geom/geom_object.h
 	geom/geom_patch.h
 	geom/geom_primitive.h

intern/cycles/kernel/bvh/bvh_shadow_all.h

@@ -187,7 +187,7 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 
 					/* primitive intersection */
 					while(prim_addr < prim_addr2) {
-						kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
+						kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) == p_type);
 
 						bool hit;
 
@@ -222,6 +222,7 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 #if BVH_FEATURE(BVH_HAIR)
 							case PRIMITIVE_CURVE:
 							case PRIMITIVE_MOTION_CURVE: {
+								const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
 								if(kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) {
 									hit = bvh_cardinal_curve_intersect(kg,
 									                                   isect_array,
@@ -231,7 +232,7 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 									                                   object,
 									                                   prim_addr,
 									                                   ray->time,
-									                                   type,
+									                                   curve_type,
 									                                   NULL,
 									                                   0, 0);
 								}
@@ -244,7 +245,7 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 									                          object,
 									                          prim_addr,
 									                          ray->time,
-									                          type,
+									                          curve_type,
 									                          NULL,
 									                          0, 0);
 								}

intern/cycles/kernel/bvh/bvh_subsurface.h

@@ -72,7 +72,7 @@ void BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 	ss_isect->num_hits = 0;
 
 	const int object_flag = kernel_tex_fetch(__object_flag, subsurface_object);
-	if(!(object_flag & SD_TRANSFORM_APPLIED)) {
+	if(!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
 #if BVH_FEATURE(BVH_MOTION)
 		Transform ob_itfm;
 		bvh_instance_motion_push(kg,

intern/cycles/kernel/bvh/bvh_traversal.h

@@ -213,7 +213,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 						--stack_ptr;
 					}
 				}
-				BVH_DEBUG_NEXT_STEP();
+				BVH_DEBUG_NEXT_NODE();
 			}
 
 			/* if node is leaf, fetch triangle list */
@@ -235,7 +235,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 					switch(type & PRIMITIVE_ALL) {
 						case PRIMITIVE_TRIANGLE: {
 							for(; prim_addr < prim_addr2; prim_addr++) {
-								BVH_DEBUG_NEXT_STEP();
+								BVH_DEBUG_NEXT_INTERSECTION();
 								kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
 								if(triangle_intersect(kg,
 								                      &isect_precalc,
@@ -264,7 +264,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 #if BVH_FEATURE(BVH_MOTION)
 						case PRIMITIVE_MOTION_TRIANGLE: {
 							for(; prim_addr < prim_addr2; prim_addr++) {
-								BVH_DEBUG_NEXT_STEP();
+								BVH_DEBUG_NEXT_INTERSECTION();
 								kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
 								if(motion_triangle_intersect(kg,
 								                             isect,
@@ -296,8 +296,9 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 						case PRIMITIVE_CURVE:
 						case PRIMITIVE_MOTION_CURVE: {
 							for(; prim_addr < prim_addr2; prim_addr++) {
-								BVH_DEBUG_NEXT_STEP();
-								kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
+								BVH_DEBUG_NEXT_INTERSECTION();
+								const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
+								kernel_assert((curve_type & PRIMITIVE_ALL) == (type & PRIMITIVE_ALL));
 								bool hit;
 								if(kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) {
 									hit = bvh_cardinal_curve_intersect(kg,
@@ -308,7 +309,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 									                                   object,
 									                                   prim_addr,
 									                                   ray->time,
-									                                   type,
+									                                   curve_type,
 									                                   lcg_state,
 									                                   difl,
 									                                   extmax);
@@ -322,7 +323,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 									                          object,
 									                          prim_addr,
 									                          ray->time,
-									                          type,
+									                          curve_type,
 									                          lcg_state,
 									                          difl,
 									                          extmax);

intern/cycles/kernel/bvh/bvh_types.h

@@ -50,12 +50,17 @@ CCL_NAMESPACE_BEGIN
 #ifdef __KERNEL_DEBUG__
 #  define BVH_DEBUG_INIT() \
 	do { \
-		isect->num_traversal_steps = 0; \
+		isect->num_traversed_nodes = 0; \
 		isect->num_traversed_instances = 0; \
+		isect->num_intersections = 0; \
 	} while(0)
-#  define BVH_DEBUG_NEXT_STEP() \
+#  define BVH_DEBUG_NEXT_NODE() \
 	do { \
-		++isect->num_traversal_steps; \
+		++isect->num_traversed_nodes; \
+	} while(0)
+#  define BVH_DEBUG_NEXT_INTERSECTION() \
+	do { \
+		++isect->num_intersections; \
 	} while(0)
 #  define BVH_DEBUG_NEXT_INSTANCE() \
 	do { \
@@ -63,7 +68,8 @@ CCL_NAMESPACE_BEGIN
 	} while(0)
 #else  /* __KERNEL_DEBUG__ */
 #  define BVH_DEBUG_INIT()
-#  define BVH_DEBUG_NEXT_STEP()
+#  define BVH_DEBUG_NEXT_NODE()
+#  define BVH_DEBUG_NEXT_INTERSECTION()
 #  define BVH_DEBUG_NEXT_INSTANCE()
 #endif  /* __KERNEL_DEBUG__ */
 

intern/cycles/kernel/bvh/bvh_volume.h

@@ -236,9 +236,7 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 					/* instance push */
 					object = kernel_tex_fetch(__prim_object, -prim_addr-1);
 					int object_flag = kernel_tex_fetch(__object_flag, object);
-
 					if(object_flag & SD_OBJECT_HAS_VOLUME) {
-
 #  if BVH_FEATURE(BVH_MOTION)
 						bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &isect->t, &ob_itfm);
 #  else

intern/cycles/kernel/bvh/bvh_volume_all.h

@@ -287,7 +287,6 @@ uint BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
 					/* instance push */
 					object = kernel_tex_fetch(__prim_object, -prim_addr-1);
 					int object_flag = kernel_tex_fetch(__object_flag, object);
-
 					if(object_flag & SD_OBJECT_HAS_VOLUME) {
 
 #  if BVH_FEATURE(BVH_MOTION)

intern/cycles/kernel/bvh/qbvh_shadow_all.h

@@ -106,14 +106,20 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 			while(node_addr >= 0 && node_addr != ENTRYPOINT_SENTINEL) {
 				float4 inodes = kernel_tex_fetch(__bvh_nodes, node_addr+0);
 
+				if(false
 #ifdef __VISIBILITY_FLAG__
-				if((__float_as_uint(inodes.x) & PATH_RAY_SHADOW) == 0) {
+				   || ((__float_as_uint(inodes.x) & PATH_RAY_SHADOW) == 0)
+#endif
+#if BVH_FEATURE(BVH_MOTION)
+				   || UNLIKELY(ray->time < inodes.y)
+				   || UNLIKELY(ray->time > inodes.z)
+#endif
+				) {
 					/* Pop. */
 					node_addr = traversal_stack[stack_ptr].addr;
 					--stack_ptr;
 					continue;
 				}
-#endif
 
 				ssef dist;
 				int child_mask = NODE_INTERSECT(kg,
@@ -262,7 +268,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 
 					/* Primitive intersection. */
 					while(prim_addr < prim_addr2) {
-						kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
+						kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) == p_type);
 
 						bool hit;
 
@@ -297,6 +303,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 #if BVH_FEATURE(BVH_HAIR)
 							case PRIMITIVE_CURVE:
 							case PRIMITIVE_MOTION_CURVE: {
+								const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
 								if(kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) {
 									hit = bvh_cardinal_curve_intersect(kg,
 									                                   isect_array,
@@ -306,7 +313,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 									                                   object,
 									                                   prim_addr,
 									                                   ray->time,
-									                                   type,
+									                                   curve_type,
 									                                   NULL,
 									                                   0, 0);
 								}
@@ -319,7 +326,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 									                          object,
 									                          prim_addr,
 									                          ray->time,
-									                          type,
+									                          curve_type,
 									                          NULL,
 									                          0, 0);
 								}

intern/cycles/kernel/bvh/qbvh_subsurface.h

@@ -61,7 +61,7 @@ ccl_device void BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 	ss_isect->num_hits = 0;
 
 	const int object_flag = kernel_tex_fetch(__object_flag, subsurface_object);
-	if(!(object_flag & SD_TRANSFORM_APPLIED)) {
+	if(!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
 #if BVH_FEATURE(BVH_MOTION)
 		Transform ob_itfm;
 		bvh_instance_motion_push(kg,

intern/cycles/kernel/bvh/qbvh_traversal.h

@@ -117,6 +117,10 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 				float4 inodes = kernel_tex_fetch(__bvh_nodes, node_addr+0);
 
 				if(UNLIKELY(node_dist > isect->t)
+#if BVH_FEATURE(BVH_MOTION)
+				   || UNLIKELY(ray->time < inodes.y)
+				   || UNLIKELY(ray->time > inodes.z)
+#endif
 #ifdef __VISIBILITY_FLAG__
 				   || (__float_as_uint(inodes.x) & visibility) == 0)
 #endif
@@ -131,7 +135,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 				int child_mask;
 				ssef dist;
 
-				BVH_DEBUG_NEXT_STEP();
+				BVH_DEBUG_NEXT_NODE();
 
 #if BVH_FEATURE(BVH_HAIR_MINIMUM_WIDTH)
 				if(difl != 0.0f) {
@@ -326,7 +330,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 					switch(type & PRIMITIVE_ALL) {
 						case PRIMITIVE_TRIANGLE: {
 							for(; prim_addr < prim_addr2; prim_addr++) {
-								BVH_DEBUG_NEXT_STEP();
+								BVH_DEBUG_NEXT_INTERSECTION();
 								kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
 								if(triangle_intersect(kg,
 								                      &isect_precalc,
@@ -347,7 +351,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 #if BVH_FEATURE(BVH_MOTION)
 						case PRIMITIVE_MOTION_TRIANGLE: {
 							for(; prim_addr < prim_addr2; prim_addr++) {
-								BVH_DEBUG_NEXT_STEP();
+								BVH_DEBUG_NEXT_INTERSECTION();
 								kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
 								if(motion_triangle_intersect(kg,
 								                             isect,
@@ -371,8 +375,9 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 						case PRIMITIVE_CURVE:
 						case PRIMITIVE_MOTION_CURVE: {
 							for(; prim_addr < prim_addr2; prim_addr++) {
-								BVH_DEBUG_NEXT_STEP();
-								kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
+								BVH_DEBUG_NEXT_INTERSECTION();
+								const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
+								kernel_assert((curve_type & PRIMITIVE_ALL) == (type & PRIMITIVE_ALL));
 								bool hit;
 								if(kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) {
 									hit = bvh_cardinal_curve_intersect(kg,
@@ -383,7 +388,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 									                                   object,
 									                                   prim_addr,
 									                                   ray->time,
-									                                   type,
+									                                   curve_type,
 									                                   lcg_state,
 									                                   difl,
 									                                   extmax);
@@ -397,7 +402,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 									                          object,
 									                          prim_addr,
 									                          ray->time,
-									                          type,
+									                          curve_type,
 									                          lcg_state,
 									                          difl,
 									                          extmax);

intern/cycles/kernel/bvh/qbvh_volume.h

@@ -293,9 +293,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 					/* Instance push. */
 					object = kernel_tex_fetch(__prim_object, -prim_addr-1);
 					int object_flag = kernel_tex_fetch(__object_flag, object);
-
 					if(object_flag & SD_OBJECT_HAS_VOLUME) {
-
 #  if BVH_FEATURE(BVH_MOTION)
 						bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &isect->t, &ob_itfm);
 #  else

intern/cycles/kernel/bvh/qbvh_volume_all.h

@@ -344,9 +344,7 @@ ccl_device uint BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
 					/* Instance push. */
 					object = kernel_tex_fetch(__prim_object, -prim_addr-1);
 					int object_flag = kernel_tex_fetch(__object_flag, object);
-
 					if(object_flag & SD_OBJECT_HAS_VOLUME) {
-
 #  if BVH_FEATURE(BVH_MOTION)
 						bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &isect_t, &ob_itfm);
 #  else

intern/cycles/kernel/closure/bsdf_ashikhmin_shirley.h

@@ -143,6 +143,7 @@ ccl_device int bsdf_ashikhmin_shirley_sample(const ShaderClosure *sc, float3 Ng,
 {
 	const MicrofacetBsdf *bsdf = (const MicrofacetBsdf*)sc;
 	float3 N = bsdf->N;
+	int label = LABEL_REFLECT | LABEL_GLOSSY;
 
 	float NdotI = dot(N, I);
 	if(NdotI > 0.0f) {
@@ -211,6 +212,7 @@ ccl_device int bsdf_ashikhmin_shirley_sample(const ShaderClosure *sc, float3 Ng,
 			/* Some high number for MIS. */
 			*pdf = 1e6f;
 			*eval = make_float3(1e6f, 1e6f, 1e6f);
+			label = LABEL_REFLECT | LABEL_SINGULAR;
 		}
 		else {
 			/* leave the rest to eval_reflect */
@@ -224,7 +226,7 @@ ccl_device int bsdf_ashikhmin_shirley_sample(const ShaderClosure *sc, float3 Ng,
 #endif
 	}
 
-	return LABEL_REFLECT|LABEL_GLOSSY;
+	return label;
 }
 
 

intern/cycles/kernel/closure/bsdf_hair.h

@@ -267,7 +267,10 @@ ccl_device int bsdf_hair_transmission_sample(const ShaderClosure *sc, float3 Ng,
 
 	*eval = make_float3(*pdf, *pdf, *pdf);
 
-	kernel_assert(dot(locy, *omega_in) < 0.0f);
+	/* TODO(sergey): Should always be negative, but seems some precision issue
+	 * is involved here.
+	 */
+	kernel_assert(dot(locy, *omega_in) < 1e-4f);
 
 	return LABEL_TRANSMIT|LABEL_GLOSSY;
 }

intern/cycles/kernel/closure/bsdf_microfacet.h

@@ -452,6 +452,7 @@ ccl_device int bsdf_microfacet_ggx_sample(KernelGlobals *kg, const ShaderClosure
 	float alpha_y = bsdf->alpha_y;
 	bool m_refractive = bsdf->type == CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID;
 	float3 N = bsdf->N;
+	int label;
 
 	float cosNO = dot(N, I);
 	if(cosNO > 0) {
@@ -477,6 +478,7 @@ ccl_device int bsdf_microfacet_ggx_sample(KernelGlobals *kg, const ShaderClosure
 		/* reflection or refraction? */
 		if(!m_refractive) {
 			float cosMO = dot(m, I);
+			label = LABEL_REFLECT | LABEL_GLOSSY;
 
 			if(cosMO > 0) {
 				/* eq. 39 - compute actual reflected direction */
@@ -487,6 +489,7 @@ ccl_device int bsdf_microfacet_ggx_sample(KernelGlobals *kg, const ShaderClosure
 						/* some high number for MIS */
 						*pdf = 1e6f;
 						*eval = make_float3(1e6f, 1e6f, 1e6f);
+						label = LABEL_REFLECT | LABEL_SINGULAR;
 					}
 					else {
 						/* microfacet normal is visible to this ray */
@@ -549,6 +552,8 @@ ccl_device int bsdf_microfacet_ggx_sample(KernelGlobals *kg, const ShaderClosure
 			}
 		}
 		else {
+			label = LABEL_TRANSMIT | LABEL_GLOSSY;
+
 			/* CAUTION: the i and o variables are inverted relative to the paper
 			 * eq. 39 - compute actual refractive direction */
 			float3 R, T;
@@ -576,6 +581,7 @@ ccl_device int bsdf_microfacet_ggx_sample(KernelGlobals *kg, const ShaderClosure
 					/* some high number for MIS */
 					*pdf = 1e6f;
 					*eval = make_float3(1e6f, 1e6f, 1e6f);
+					label = LABEL_TRANSMIT | LABEL_SINGULAR;
 				}
 				else {
 					/* eq. 33 */
@@ -607,7 +613,10 @@ ccl_device int bsdf_microfacet_ggx_sample(KernelGlobals *kg, const ShaderClosure
 			}
 		}
 	}
-	return (m_refractive) ? LABEL_TRANSMIT|LABEL_GLOSSY : LABEL_REFLECT|LABEL_GLOSSY;
+	else {
+		label = (m_refractive) ? LABEL_TRANSMIT|LABEL_GLOSSY : LABEL_REFLECT|LABEL_GLOSSY;
+	}
+	return label;
 }
 
 /* Beckmann microfacet with Smith shadow-masking from:
@@ -815,6 +824,7 @@ ccl_device int bsdf_microfacet_beckmann_sample(KernelGlobals *kg, const ShaderCl
 	float alpha_y = bsdf->alpha_y;
 	bool m_refractive = bsdf->type == CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID;
 	float3 N = bsdf->N;
+	int label;
 
 	float cosNO = dot(N, I);
 	if(cosNO > 0) {
@@ -839,6 +849,7 @@ ccl_device int bsdf_microfacet_beckmann_sample(KernelGlobals *kg, const ShaderCl
 
 		/* reflection or refraction? */
 		if(!m_refractive) {
+			label = LABEL_REFLECT | LABEL_GLOSSY;
 			float cosMO = dot(m, I);
 
 			if(cosMO > 0) {
@@ -850,6 +861,7 @@ ccl_device int bsdf_microfacet_beckmann_sample(KernelGlobals *kg, const ShaderCl
 						/* some high number for MIS */
 						*pdf = 1e6f;
 						*eval = make_float3(1e6f, 1e6f, 1e6f);
+						label = LABEL_REFLECT | LABEL_SINGULAR;
 					}
 					else {
 						/* microfacet normal is visible to this ray
@@ -904,6 +916,8 @@ ccl_device int bsdf_microfacet_beckmann_sample(KernelGlobals *kg, const ShaderCl
 			}
 		}
 		else {
+			label = LABEL_TRANSMIT | LABEL_GLOSSY;
+
 			/* CAUTION: the i and o variables are inverted relative to the paper
 			 * eq. 39 - compute actual refractive direction */
 			float3 R, T;
@@ -931,6 +945,7 @@ ccl_device int bsdf_microfacet_beckmann_sample(KernelGlobals *kg, const ShaderCl
 					/* some high number for MIS */
 					*pdf = 1e6f;
 					*eval = make_float3(1e6f, 1e6f, 1e6f);
+					label = LABEL_TRANSMIT | LABEL_SINGULAR;
 				}
 				else {
 					/* eq. 33 */
@@ -963,7 +978,10 @@ ccl_device int bsdf_microfacet_beckmann_sample(KernelGlobals *kg, const ShaderCl
 			}
 		}
 	}
-	return (m_refractive) ? LABEL_TRANSMIT|LABEL_GLOSSY : LABEL_REFLECT|LABEL_GLOSSY;
+	else {
+		label = (m_refractive) ? LABEL_TRANSMIT|LABEL_GLOSSY : LABEL_REFLECT|LABEL_GLOSSY;
+	}
+	return label;
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/geom/geom.h

@@ -23,6 +23,8 @@
 #include "geom_subd_triangle.h"
 #include "geom_triangle_intersect.h"
 #include "geom_motion_triangle.h"
+#include "geom_motion_triangle_intersect.h"
+#include "geom_motion_triangle_shader.h"
 #include "geom_motion_curve.h"
 #include "geom_curve.h"
 #include "geom_volume.h"

intern/cycles/kernel/geom/geom_curve.h

@@ -255,6 +255,17 @@ ccl_device_forceinline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Inte
 		int ka = max(k0 - 1, v00.x);
 		int kb = min(k1 + 1, v00.x + v00.y - 1);
 
+#ifdef __KERNEL_AVX2__
+		avxf P_curve_0_1, P_curve_2_3;
+		if(type & PRIMITIVE_CURVE) {
+			P_curve_0_1 = _mm256_loadu2_m128(&kg->__curve_keys.data[k0].x, &kg->__curve_keys.data[ka].x);
+			P_curve_2_3 = _mm256_loadu2_m128(&kg->__curve_keys.data[kb].x, &kg->__curve_keys.data[k1].x);
+		}
+		else {
+			int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
+			motion_cardinal_curve_keys_avx(kg, fobject, prim, time, ka, k0, k1, kb, &P_curve_0_1,&P_curve_2_3);
+		}
+#else  /* __KERNEL_AVX2__ */
 		ssef P_curve[4];
 
 		if(type & PRIMITIVE_CURVE) {
@@ -267,6 +278,7 @@ ccl_device_forceinline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Inte
 			int fobject = (object == OBJECT_NONE)? kernel_tex_fetch(__prim_object, curveAddr): object;
 			motion_cardinal_curve_keys(kg, fobject, prim, time, ka, k0, k1, kb, (float4*)&P_curve);
 		}
+#endif  /* __KERNEL_AVX2__ */
 
 		ssef rd_sgn = set_sign_bit<0, 1, 1, 1>(shuffle<0>(rd_ss));
 		ssef mul_zxxy = shuffle<2, 0, 0, 1>(vdir) * rd_sgn;
@@ -278,6 +290,33 @@ ccl_device_forceinline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Inte
 		ssef htfm1 = shuffle<1, 0, 1, 3>(load1f_first(extract<0>(d_ss)), vdir0);
 		ssef htfm2 = shuffle<1, 3, 2, 3>(mul_shuf, vdir0);
 
+#ifdef __KERNEL_AVX2__
+		const avxf vPP = _mm256_broadcast_ps(&P.m128);
+		const avxf htfm00 = avxf(htfm0.m128, htfm0.m128);
+		const avxf htfm11 = avxf(htfm1.m128, htfm1.m128);
+		const avxf htfm22 = avxf(htfm2.m128, htfm2.m128);
+
+		const avxf p01 = madd(shuffle<0>(P_curve_0_1 - vPP),
+		                      htfm00,
+		                      madd(shuffle<1>(P_curve_0_1 - vPP),
+		                           htfm11,
+		                           shuffle<2>(P_curve_0_1 - vPP) * htfm22));
+		const avxf p23 = madd(shuffle<0>(P_curve_2_3 - vPP),
+		                      htfm00,
+		                      madd(shuffle<1>(P_curve_2_3 - vPP),
+		                           htfm11,
+		                           shuffle<2>(P_curve_2_3 - vPP)*htfm22));
+
+		const ssef p0 = _mm256_castps256_ps128(p01);
+		const ssef p1 = _mm256_extractf128_ps(p01, 1);
+		const ssef p2 = _mm256_castps256_ps128(p23);
+		const ssef p3 = _mm256_extractf128_ps(p23, 1);
+
+		const ssef P_curve_1 = _mm256_extractf128_ps(P_curve_0_1, 1);
+		r_st = ((float4 &)P_curve_1).w;
+		const ssef P_curve_2 = _mm256_castps256_ps128(P_curve_2_3);
+		r_en = ((float4 &)P_curve_2).w;
+#else  /* __KERNEL_AVX2__ */
 		ssef htfm[] = { htfm0, htfm1, htfm2 };
 		ssef vP = load4f(P);
 		ssef p0 = transform_point_T3(htfm, P_curve[0] - vP);
@@ -285,6 +324,10 @@ ccl_device_forceinline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Inte
 		ssef p2 = transform_point_T3(htfm, P_curve[2] - vP);
 		ssef p3 = transform_point_T3(htfm, P_curve[3] - vP);
 
+		r_st = ((float4 &)P_curve[1]).w;
+		r_en = ((float4 &)P_curve[2]).w;
+#endif  /* __KERNEL_AVX2__ */
+
 		float fc = 0.71f;
 		ssef vfc = ssef(fc);
 		ssef vfcxp3 = vfc * p3;
@@ -294,8 +337,6 @@ ccl_device_forceinline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Inte
 		vcurve_coef[2] = madd(ssef(fc * 2.0f), p0, madd(ssef(fc - 3.0f), p1, msub(ssef(3.0f - 2.0f * fc), p2, vfcxp3)));
 		vcurve_coef[3] = msub(ssef(fc - 2.0f), p2 - p1, msub(vfc, p0, vfcxp3));
 
-		r_st = ((float4 &)P_curve[1]).w;
-		r_en = ((float4 &)P_curve[2]).w;
 	}
 #else
 	float3 curve_coef[4];
@@ -383,8 +424,9 @@ ccl_device_forceinline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Inte
 
 	/* begin loop */
 	while(!(tree >> (depth))) {
-		float i_st = tree * resol;
-		float i_en = i_st + (level * resol);
+		const float i_st = tree * resol;
+		const float i_en = i_st + (level * resol);
+
 #ifdef __KERNEL_SSE2__
 		ssef vi_st = ssef(i_st), vi_en = ssef(i_en);
 		ssef vp_st = madd(madd(madd(vcurve_coef[3], vi_st, vcurve_coef[2]), vi_st, vcurve_coef[1]), vi_st, vcurve_coef[0]);
@@ -458,13 +500,23 @@ ccl_device_forceinline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Inte
 
 			if(flags & CURVE_KN_RIBBONS) {
 				float3 tg = (p_en - p_st);
+#ifdef __KERNEL_SSE__
+				const float3 tg_sq = tg * tg;
+				float w = tg_sq.x + tg_sq.y;
+#else
 				float w = tg.x * tg.x + tg.y * tg.y;
+#endif
 				if(w == 0) {
 					tree++;
 					level = tree & -tree;
 					continue;
 				}
+#ifdef __KERNEL_SSE__
+				const float3 p_sttg = p_st * tg;
+				w = -(p_sttg.x + p_sttg.y) / w;
+#else
 				w = -(p_st.x * tg.x + p_st.y * tg.y) / w;
+#endif
 				w = saturate(w);
 
 				/* compute u on the curve segment */
@@ -496,7 +548,13 @@ ccl_device_forceinline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Inte
 				if(difl != 0.0f) {
 					mw_extension = min(difl * fabsf(bmaxz), extmax);
 					r_ext = mw_extension + r_curr;
+#ifdef __KERNEL_SSE__
+					const float3 p_curr_sq = p_curr * p_curr;
+					const float3 dxxx = _mm_sqrt_ss(_mm_hadd_ps(p_curr_sq.m128, p_curr_sq.m128));
+					float d = dxxx.x;
+#else
 					float d = sqrtf(p_curr.x * p_curr.x + p_curr.y * p_curr.y);
+#endif
 					float d0 = d - r_curr;
 					float d1 = d + r_curr;
 					float inv_mw_extension = 1.0f/mw_extension;
@@ -853,7 +911,7 @@ ccl_device_forceinline bool bvh_curve_intersect(KernelGlobals *kg, Intersection
 #  undef len3_squared
 #  undef len3
 #  undef dot3
-#  endif
+#endif
 }
 
 ccl_device_inline float3 curvetangent(float t, float3 p0, float3 p1, float3 p2, float3 p3)

intern/cycles/kernel/geom/geom_motion_curve.h

@@ -50,12 +50,12 @@ ccl_device_inline int find_attribute_curve_motion(KernelGlobals *kg, int object,
 ccl_device_inline void motion_curve_keys_for_step(KernelGlobals *kg, int offset, int numkeys, int numsteps, int step, int k0, int k1, float4 keys[2])
 {
 	if(step == numsteps) {
-		/* center step: regular vertex location */
+		/* center step: regular key location */
 		keys[0] = kernel_tex_fetch(__curve_keys, k0);
 		keys[1] = kernel_tex_fetch(__curve_keys, k1);
 	}
 	else {
-		/* center step not stored in this array */
+		/* center step is not stored in this array */
 		if(step > numsteps)
 			step--;
 
@@ -97,14 +97,14 @@ ccl_device_inline void motion_curve_keys(KernelGlobals *kg, int object, int prim
 ccl_device_inline void motion_cardinal_curve_keys_for_step(KernelGlobals *kg, int offset, int numkeys, int numsteps, int step, int k0, int k1, int k2, int k3, float4 keys[4])
 {
 	if(step == numsteps) {
-		/* center step: regular vertex location */
+		/* center step: regular key location */
 		keys[0] = kernel_tex_fetch(__curve_keys, k0);
 		keys[1] = kernel_tex_fetch(__curve_keys, k1);
 		keys[2] = kernel_tex_fetch(__curve_keys, k2);
 		keys[3] = kernel_tex_fetch(__curve_keys, k3);
 	}
 	else {
-		/* center step not store in this array */
+		/* center step is not stored in this array */
 		if(step > numsteps)
 			step--;
 
@@ -118,7 +118,12 @@ ccl_device_inline void motion_cardinal_curve_keys_for_step(KernelGlobals *kg, in
 }
 
 /* return 2 curve key locations */
-ccl_device_inline void motion_cardinal_curve_keys(KernelGlobals *kg, int object, int prim, float time, int k0, int k1, int k2, int k3, float4 keys[4])
+ccl_device_inline void motion_cardinal_curve_keys(KernelGlobals *kg,
+                                                  int object,
+                                                  int prim,
+                                                  float time,
+                                                  int k0, int k1, int k2, int k3,
+                                                  float4 keys[4])
 {
 	/* get motion info */
 	int numsteps, numkeys;
@@ -147,6 +152,65 @@ ccl_device_inline void motion_cardinal_curve_keys(KernelGlobals *kg, int object,
 	keys[3] = (1.0f - t)*keys[3] + t*next_keys[3];
 }
 
+#ifdef __KERNEL_AVX2__
+/* Similar to above, but returns keys as pair of two AVX registers with each
+ * holding two float4.
+ */
+ccl_device_inline void motion_cardinal_curve_keys_avx(KernelGlobals *kg,
+                                                      int object,
+                                                      int prim,
+                                                      float time,
+                                                      int k0, int k1,
+                                                      int k2, int k3,
+                                                      avxf *out_keys_0_1,
+                                                      avxf *out_keys_2_3)
+{
+	/* Get motion info. */
+	int numsteps, numkeys;
+	object_motion_info(kg, object, &numsteps, NULL, &numkeys);
+
+	/* Figure out which steps we need to fetch and their interpolation factor. */
+	int maxstep = numsteps * 2;
+	int step = min((int)(time*maxstep), maxstep - 1);
+	float t = time*maxstep - step;
+
+	/* Find attribute. */
+	AttributeElement elem;
+	int offset = find_attribute_curve_motion(kg,
+	                                         object,
+	                                         ATTR_STD_MOTION_VERTEX_POSITION,
+	                                         &elem);
+	kernel_assert(offset != ATTR_STD_NOT_FOUND);
+
+	/* Fetch key coordinates. */
+	float4 next_keys[4];
+	float4 keys[4];
+	motion_cardinal_curve_keys_for_step(kg,
+	                                    offset,
+	                                    numkeys,
+	                                    numsteps,
+	                                    step,
+	                                    k0, k1, k2, k3,
+	                                    keys);
+	motion_cardinal_curve_keys_for_step(kg,
+	                                    offset,
+	                                    numkeys,
+	                                    numsteps,
+	                                    step + 1,
+	                                    k0, k1, k2, k3,
+	                                    next_keys);
+
+	const avxf keys_0_1 = avxf(keys[0].m128, keys[1].m128);
+	const avxf keys_2_3 = avxf(keys[2].m128, keys[3].m128);
+	const avxf next_keys_0_1 = avxf(next_keys[0].m128, next_keys[1].m128);
+	const avxf next_keys_2_3 = avxf(next_keys[2].m128, next_keys[3].m128);
+
+	/* Interpolate between steps. */
+	*out_keys_0_1 = (1.0f - t) * keys_0_1 + t*next_keys_0_1;
+	*out_keys_2_3 = (1.0f - t) * keys_2_3 + t*next_keys_2_3;
+}
+#endif
+
 #endif
 
 CCL_NAMESPACE_END

intern/cycles/kernel/geom/geom_motion_triangle.h

@@ -76,7 +76,7 @@ ccl_device_inline void motion_triangle_normals_for_step(KernelGlobals *kg, uint4
 		normals[2] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.z));
 	}
 	else {
-		/* center step not stored in this array */
+		/* center step is not stored in this array */
 		if(step > numsteps)
 			step--;
 
@@ -117,312 +117,4 @@ ccl_device_inline void motion_triangle_vertices(KernelGlobals *kg, int object, i
 	verts[2] = (1.0f - t)*verts[2] + t*next_verts[2];
 }
 
-/* Refine triangle intersection to more precise hit point. For rays that travel
- * far the precision is often not so good, this reintersects the primitive from
- * a closer distance. */
-
-ccl_device_inline float3 motion_triangle_refine(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray, float3 verts[3])
-{
-	float3 P = ray->P;
-	float3 D = ray->D;
-	float t = isect->t;
-
-#ifdef __INTERSECTION_REFINE__
-	if(isect->object != OBJECT_NONE) {
-		if(UNLIKELY(t == 0.0f)) {
-			return P;
-		}
-#  ifdef __OBJECT_MOTION__
-		Transform tfm = ccl_fetch(sd, ob_itfm);
-#  else
-		Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM);
-#  endif
-
-		P = transform_point(&tfm, P);
-		D = transform_direction(&tfm, D*t);
-		D = normalize_len(D, &t);
-	}
-
-	P = P + D*t;
-
-	/* compute refined intersection distance */
-	const float3 e1 = verts[0] - verts[2];
-	const float3 e2 = verts[1] - verts[2];
-	const float3 s1 = cross(D, e2);
-
-	const float invdivisor = 1.0f/dot(s1, e1);
-	const float3 d = P - verts[2];
-	const float3 s2 = cross(d, e1);
-	float rt = dot(e2, s2)*invdivisor;
-
-	/* compute refined position */
-	P = P + D*rt;
-
-	if(isect->object != OBJECT_NONE) {
-#  ifdef __OBJECT_MOTION__
-		Transform tfm = ccl_fetch(sd, ob_tfm);
-#  else
-		Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM);
-#  endif
-
-		P = transform_point(&tfm, P);
-	}
-
-	return P;
-#else
-	return P + D*t;
-#endif
-}
-
-/* Same as above, except that isect->t is assumed to be in object space for instancing */
-
-#ifdef __SUBSURFACE__
-#  if defined(__KERNEL_CUDA__) && (defined(i386) || defined(_M_IX86))
-ccl_device_noinline
-#  else
-ccl_device_inline
-#  endif
-float3 motion_triangle_refine_subsurface(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray, float3 verts[3])
-{
-	float3 P = ray->P;
-	float3 D = ray->D;
-	float t = isect->t;
-
-#  ifdef __INTERSECTION_REFINE__
-	if(isect->object != OBJECT_NONE) {
-#    ifdef __OBJECT_MOTION__
-		Transform tfm = ccl_fetch(sd, ob_itfm);
-#    else
-		Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM);
-#    endif
-
-		P = transform_point(&tfm, P);
-		D = transform_direction(&tfm, D);
-		D = normalize(D);
-	}
-
-	P = P + D*t;
-
-	/* compute refined intersection distance */
-	const float3 e1 = verts[0] - verts[2];
-	const float3 e2 = verts[1] - verts[2];
-	const float3 s1 = cross(D, e2);
-
-	const float invdivisor = 1.0f/dot(s1, e1);
-	const float3 d = P - verts[2];
-	const float3 s2 = cross(d, e1);
-	float rt = dot(e2, s2)*invdivisor;
-
-	P = P + D*rt;
-
-	if(isect->object != OBJECT_NONE) {
-#    ifdef __OBJECT_MOTION__
-		Transform tfm = ccl_fetch(sd, ob_tfm);
-#    else
-		Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM);
-#    endif
-
-		P = transform_point(&tfm, P);
-	}
-
-	return P;
-#  else
-	return P + D*t;
-#  endif
-}
-#endif
-
-/* Setup of motion triangle specific parts of ShaderData, moved into this one
- * function to more easily share computation of interpolated positions and
- * normals */
-
-/* return 3 triangle vertex normals */
-ccl_device_noinline void motion_triangle_shader_setup(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray, bool subsurface)
-{
-	/* get shader */
-	ccl_fetch(sd, shader) = kernel_tex_fetch(__tri_shader, ccl_fetch(sd, prim));
-
-	/* get motion info */
-	int numsteps, numverts;
-	object_motion_info(kg, ccl_fetch(sd, object), &numsteps, &numverts, NULL);
-
-	/* figure out which steps we need to fetch and their interpolation factor */
-	int maxstep = numsteps*2;
-	int step = min((int)(ccl_fetch(sd, time)*maxstep), maxstep-1);
-	float t = ccl_fetch(sd, time)*maxstep - step;
-
-	/* find attribute */
-	AttributeElement elem;
-	int offset = find_attribute_motion(kg, ccl_fetch(sd, object), ATTR_STD_MOTION_VERTEX_POSITION, &elem);
-	kernel_assert(offset != ATTR_STD_NOT_FOUND);
-
-	/* fetch vertex coordinates */
-	float3 verts[3], next_verts[3];
-	uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, ccl_fetch(sd, prim));
-
-	motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step, verts);
-	motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step+1, next_verts);
-
-	/* interpolate between steps */
-	verts[0] = (1.0f - t)*verts[0] + t*next_verts[0];
-	verts[1] = (1.0f - t)*verts[1] + t*next_verts[1];
-	verts[2] = (1.0f - t)*verts[2] + t*next_verts[2];
-
-	/* compute refined position */
-#ifdef __SUBSURFACE__
-	if(!subsurface)
-#endif
-		ccl_fetch(sd, P) = motion_triangle_refine(kg, sd, isect, ray, verts);
-#ifdef __SUBSURFACE__
-	else
-		ccl_fetch(sd, P) = motion_triangle_refine_subsurface(kg, sd, isect, ray, verts);
-#endif
-
-	/* compute face normal */
-	float3 Ng;
-	if(ccl_fetch(sd, flag) & SD_NEGATIVE_SCALE_APPLIED)
-		Ng = normalize(cross(verts[2] - verts[0], verts[1] - verts[0]));
-	else
-		Ng = normalize(cross(verts[1] - verts[0], verts[2] - verts[0]));
-
-	ccl_fetch(sd, Ng) = Ng;
-	ccl_fetch(sd, N) = Ng;
-
-	/* compute derivatives of P w.r.t. uv */
-#ifdef __DPDU__
-	ccl_fetch(sd, dPdu) = (verts[0] - verts[2]);
-	ccl_fetch(sd, dPdv) = (verts[1] - verts[2]);
-#endif
-
-	/* compute smooth normal */
-	if(ccl_fetch(sd, shader) & SHADER_SMOOTH_NORMAL) {
-		/* find attribute */
-		AttributeElement elem;
-		int offset = find_attribute_motion(kg, ccl_fetch(sd, object), ATTR_STD_MOTION_VERTEX_NORMAL, &elem);
-		kernel_assert(offset != ATTR_STD_NOT_FOUND);
-
-		/* fetch vertex coordinates */
-		float3 normals[3], next_normals[3];
-		motion_triangle_normals_for_step(kg, tri_vindex, offset, numverts, numsteps, step, normals);
-		motion_triangle_normals_for_step(kg, tri_vindex, offset, numverts, numsteps, step+1, next_normals);
-
-		/* interpolate between steps */
-		normals[0] = (1.0f - t)*normals[0] + t*next_normals[0];
-		normals[1] = (1.0f - t)*normals[1] + t*next_normals[1];
-		normals[2] = (1.0f - t)*normals[2] + t*next_normals[2];
-
-		/* interpolate between vertices */
-		float u = ccl_fetch(sd, u);
-		float v = ccl_fetch(sd, v);
-		float w = 1.0f - u - v;
-		ccl_fetch(sd, N) = (u*normals[0] + v*normals[1] + w*normals[2]);
-	}
-}
-
-/* Ray intersection. We simply compute the vertex positions at the given ray
- * time and do a ray intersection with the resulting triangle */
-
-ccl_device_inline bool motion_triangle_intersect(KernelGlobals *kg, Intersection *isect,
-	float3 P, float3 dir, float time, uint visibility, int object, int triAddr)
-{
-	/* primitive index for vertex location lookup */
-	int prim = kernel_tex_fetch(__prim_index, triAddr);
-	int fobject = (object == OBJECT_NONE)? kernel_tex_fetch(__prim_object, triAddr): object;
-
-	/* get vertex locations for intersection */
-	float3 verts[3];
-	motion_triangle_vertices(kg, fobject, prim, time, verts);
-
-	/* ray-triangle intersection, unoptimized */
-	float t, u, v;
-
-	if(ray_triangle_intersect_uv(P, dir, isect->t, verts[2], verts[0], verts[1], &u, &v, &t)) {
-#ifdef __VISIBILITY_FLAG__
-		/* visibility flag test. we do it here under the assumption
-		 * that most triangles are culled by node flags */
-		if(kernel_tex_fetch(__prim_visibility, triAddr) & visibility)
-#endif
-		{
-			isect->t = t;
-			isect->u = u;
-			isect->v = v;
-			isect->prim = triAddr;
-			isect->object = object;
-			isect->type = PRIMITIVE_MOTION_TRIANGLE;
-		
-			return true;
-		}
-	}
-
-	return false;
-}
-
-/* Special ray intersection routines for subsurface scattering. In that case we
- * only want to intersect with primitives in the same object, and if case of
- * multiple hits we pick a single random primitive as the intersection point. */
-
-#ifdef __SUBSURFACE__
-ccl_device_inline void motion_triangle_intersect_subsurface(
-        KernelGlobals *kg,
-        SubsurfaceIntersection *ss_isect,
-        float3 P,
-        float3 dir,
-        float time,
-        int object,
-        int triAddr,
-        float tmax,
-        uint *lcg_state,
-        int max_hits)
-{
-	/* primitive index for vertex location lookup */
-	int prim = kernel_tex_fetch(__prim_index, triAddr);
-	int fobject = (object == OBJECT_NONE)? kernel_tex_fetch(__prim_object, triAddr): object;
-
-	/* get vertex locations for intersection */
-	float3 verts[3];
-	motion_triangle_vertices(kg, fobject, prim, time, verts);
-
-	/* ray-triangle intersection, unoptimized */
-	float t, u, v;
-
-	if(ray_triangle_intersect_uv(P, dir, tmax, verts[2], verts[0], verts[1], &u, &v, &t)) {
-		for(int i = min(max_hits, ss_isect->num_hits) - 1; i >= 0; --i) {
-			if(ss_isect->hits[i].t == t) {
-				return;
-			}
-		}
-
-		ss_isect->num_hits++;
-
-		int hit;
-
-		if(ss_isect->num_hits <= max_hits) {
-			hit = ss_isect->num_hits - 1;
-		}
-		else {
-			/* reservoir sampling: if we are at the maximum number of
-			 * hits, randomly replace element or skip it */
-			hit = lcg_step_uint(lcg_state) % ss_isect->num_hits;
-
-			if(hit >= max_hits)
-				return;
-		}
-
-		/* record intersection */
-		Intersection *isect = &ss_isect->hits[hit];
-		isect->t = t;
-		isect->u = u;
-		isect->v = v;
-		isect->prim = triAddr;
-		isect->object = object;
-		isect->type = PRIMITIVE_MOTION_TRIANGLE;
-
-		/* Record geometric normal. */
-		ss_isect->Ng[hit] = normalize(cross(verts[1] - verts[0],
-		                                    verts[2] - verts[0]));
-	}
-}
-#endif
-
 CCL_NAMESPACE_END
-

intern/cycles/kernel/geom/geom_motion_triangle_intersect.h

@@ -0,0 +1,280 @@
+/*
+ * Copyright 2011-2016 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.
+ */
+
+/* Motion Triangle Primitive
+ *
+ * These are stored as regular triangles, plus extra positions and normals at
+ * times other than the frame center. Computing the triangle vertex positions
+ * or normals at a given ray time is a matter of interpolation of the two steps
+ * between which the ray time lies.
+ *
+ * The extra positions and normals are stored as ATTR_STD_MOTION_VERTEX_POSITION
+ * and ATTR_STD_MOTION_VERTEX_NORMAL mesh attributes.
+ */
+
+CCL_NAMESPACE_BEGIN
+
+/* Refine triangle intersection to more precise hit point. For rays that travel
+ * far the precision is often not so good, this reintersects the primitive from
+ * a closer distance.
+ */
+
+ccl_device_inline float3 motion_triangle_refine(KernelGlobals *kg,
+                                                ShaderData *sd,
+                                                const Intersection *isect,
+                                                const Ray *ray,
+                                                float3 verts[3])
+{
+	float3 P = ray->P;
+	float3 D = ray->D;
+	float t = isect->t;
+
+#ifdef __INTERSECTION_REFINE__
+	if(isect->object != OBJECT_NONE) {
+		if(UNLIKELY(t == 0.0f)) {
+			return P;
+		}
+#  ifdef __OBJECT_MOTION__
+		Transform tfm = ccl_fetch(sd, ob_itfm);
+#  else
+		Transform tfm = object_fetch_transform(kg,
+		                                       isect->object,
+		                                       OBJECT_INVERSE_TRANSFORM);
+#  endif
+
+		P = transform_point(&tfm, P);
+		D = transform_direction(&tfm, D*t);
+		D = normalize_len(D, &t);
+	}
+
+	P = P + D*t;
+
+	/* Compute refined intersection distance. */
+	const float3 e1 = verts[0] - verts[2];
+	const float3 e2 = verts[1] - verts[2];
+	const float3 s1 = cross(D, e2);
+
+	const float invdivisor = 1.0f/dot(s1, e1);
+	const float3 d = P - verts[2];
+	const float3 s2 = cross(d, e1);
+	float rt = dot(e2, s2)*invdivisor;
+
+	/* Compute refined position. */
+	P = P + D*rt;
+
+	if(isect->object != OBJECT_NONE) {
+#  ifdef __OBJECT_MOTION__
+		Transform tfm = ccl_fetch(sd, ob_tfm);
+#  else
+		Transform tfm = object_fetch_transform(kg,
+		                                       isect->object,
+		                                       OBJECT_TRANSFORM);
+#  endif
+
+		P = transform_point(&tfm, P);
+	}
+
+	return P;
+#else
+	return P + D*t;
+#endif
+}
+
+/* Same as above, except that isect->t is assumed to be in object space
+ * for instancing.
+ */
+
+#ifdef __SUBSURFACE__
+#  if defined(__KERNEL_CUDA__) && (defined(i386) || defined(_M_IX86))
+ccl_device_noinline
+#  else
+ccl_device_inline
+#  endif
+float3 motion_triangle_refine_subsurface(KernelGlobals *kg,
+                                         ShaderData *sd,
+                                         const Intersection *isect,
+                                         const Ray *ray,
+                                         float3 verts[3])
+{
+	float3 P = ray->P;
+	float3 D = ray->D;
+	float t = isect->t;
+
+#  ifdef __INTERSECTION_REFINE__
+	if(isect->object != OBJECT_NONE) {
+#    ifdef __OBJECT_MOTION__
+		Transform tfm = ccl_fetch(sd, ob_itfm);
+#    else
+		Transform tfm = object_fetch_transform(kg,
+		                                       isect->object,
+		                                       OBJECT_INVERSE_TRANSFORM);
+#    endif
+
+		P = transform_point(&tfm, P);
+		D = transform_direction(&tfm, D);
+		D = normalize(D);
+	}
+
+	P = P + D*t;
+
+	/* compute refined intersection distance */
+	const float3 e1 = verts[0] - verts[2];
+	const float3 e2 = verts[1] - verts[2];
+	const float3 s1 = cross(D, e2);
+
+	const float invdivisor = 1.0f/dot(s1, e1);
+	const float3 d = P - verts[2];
+	const float3 s2 = cross(d, e1);
+	float rt = dot(e2, s2)*invdivisor;
+
+	P = P + D*rt;
+
+	if(isect->object != OBJECT_NONE) {
+#    ifdef __OBJECT_MOTION__
+		Transform tfm = ccl_fetch(sd, ob_tfm);
+#    else
+		Transform tfm = object_fetch_transform(kg,
+		                                       isect->object,
+		                                       OBJECT_TRANSFORM);
+#    endif
+
+		P = transform_point(&tfm, P);
+	}
+
+	return P;
+#  else  /* __INTERSECTION_REFINE__ */
+	return P + D*t;
+#  endif  /* __INTERSECTION_REFINE__ */
+}
+#endif  /* __SUBSURFACE__ */
+
+
+/* Ray intersection. We simply compute the vertex positions at the given ray
+ * time and do a ray intersection with the resulting triangle.
+ */
+
+ccl_device_inline bool motion_triangle_intersect(KernelGlobals *kg,
+                                                 Intersection *isect,
+                                                 float3 P,
+                                                 float3 dir,
+                                                 float time,
+                                                 uint visibility,
+                                                 int object,
+                                                 int prim_addr)
+{
+	/* Primitive index for vertex location lookup. */
+	int prim = kernel_tex_fetch(__prim_index, prim_addr);
+	int fobject = (object == OBJECT_NONE)
+	                  ? kernel_tex_fetch(__prim_object, prim_addr)
+	                  : object;
+	/* Get vertex locations for intersection. */
+	float3 verts[3];
+	motion_triangle_vertices(kg, fobject, prim, time, verts);
+	/* Ray-triangle intersection, unoptimized. */
+	float t, u, v;
+	if(ray_triangle_intersect_uv(P,
+	                             dir,
+	                             isect->t,
+	                             verts[2], verts[0], verts[1],
+	                             &u, &v, &t))
+	{
+#ifdef __VISIBILITY_FLAG__
+		/* Visibility flag test. we do it here under the assumption
+		 * that most triangles are culled by node flags.
+		 */
+		if(kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
+#endif
+		{
+			isect->t = t;
+			isect->u = u;
+			isect->v = v;
+			isect->prim = prim_addr;
+			isect->object = object;
+			isect->type = PRIMITIVE_MOTION_TRIANGLE;
+			return true;
+		}
+	}
+	return false;
+}
+
+/* Special ray intersection routines for subsurface scattering. In that case we
+ * only want to intersect with primitives in the same object, and if case of
+ * multiple hits we pick a single random primitive as the intersection point.
+ */
+#ifdef __SUBSURFACE__
+ccl_device_inline void motion_triangle_intersect_subsurface(
+        KernelGlobals *kg,
+        SubsurfaceIntersection *ss_isect,
+        float3 P,
+        float3 dir,
+        float time,
+        int object,
+        int prim_addr,
+        float tmax,
+        uint *lcg_state,
+        int max_hits)
+{
+	/* Primitive index for vertex location lookup. */
+	int prim = kernel_tex_fetch(__prim_index, prim_addr);
+	int fobject = (object == OBJECT_NONE)
+	                  ? kernel_tex_fetch(__prim_object, prim_addr)
+	                  : object;
+	/* Get vertex locations for intersection. */
+	float3 verts[3];
+	motion_triangle_vertices(kg, fobject, prim, time, verts);
+	/* Ray-triangle intersection, unoptimized. */
+	float t, u, v;
+	if(ray_triangle_intersect_uv(P,
+	                             dir,
+	                             tmax,
+	                             verts[2], verts[0], verts[1],
+	                             &u, &v, &t))
+	{
+		for(int i = min(max_hits, ss_isect->num_hits) - 1; i >= 0; --i) {
+			if(ss_isect->hits[i].t == t) {
+				return;
+			}
+		}
+		ss_isect->num_hits++;
+		int hit;
+		if(ss_isect->num_hits <= max_hits) {
+			hit = ss_isect->num_hits - 1;
+		}
+		else {
+			/* Reservoir sampling: if we are at the maximum number of
+			 * hits, randomly replace element or skip it.
+			 */
+			hit = lcg_step_uint(lcg_state) % ss_isect->num_hits;
+
+			if(hit >= max_hits)
+				return;
+		}
+		/* Record intersection. */
+		Intersection *isect = &ss_isect->hits[hit];
+		isect->t = t;
+		isect->u = u;
+		isect->v = v;
+		isect->prim = prim_addr;
+		isect->object = object;
+		isect->type = PRIMITIVE_MOTION_TRIANGLE;
+		/* Record geometric normal. */
+		ss_isect->Ng[hit] = normalize(cross(verts[1] - verts[0],
+		                                    verts[2] - verts[0]));
+	}
+}
+#endif  /* __SUBSURFACE__ */
+
+CCL_NAMESPACE_END

intern/cycles/kernel/geom/geom_motion_triangle_shader.h

@@ -0,0 +1,123 @@
+/*
+ * Copyright 2011-2016 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.
+ */
+
+/* Motion Triangle Primitive
+ *
+ * These are stored as regular triangles, plus extra positions and normals at
+ * times other than the frame center. Computing the triangle vertex positions
+ * or normals at a given ray time is a matter of interpolation of the two steps
+ * between which the ray time lies.
+ *
+ * The extra positions and normals are stored as ATTR_STD_MOTION_VERTEX_POSITION
+ * and ATTR_STD_MOTION_VERTEX_NORMAL mesh attributes.
+ */
+
+CCL_NAMESPACE_BEGIN
+
+/* Setup of motion triangle specific parts of ShaderData, moved into this one
+ * function to more easily share computation of interpolated positions and
+ * normals */
+
+/* return 3 triangle vertex normals */
+ccl_device_noinline void motion_triangle_shader_setup(KernelGlobals *kg,
+                                                      ShaderData *sd, const
+                                                      Intersection *isect,
+                                                      const Ray *ray,
+                                                      bool subsurface)
+{
+	/* Get shader. */
+	ccl_fetch(sd, shader) = kernel_tex_fetch(__tri_shader, ccl_fetch(sd, prim));
+	/* Get motion info. */
+	/* TODO(sergey): This logic is really similar to motion_triangle_vertices(),
+	 * can we de-duplicate something here?
+	 */
+	int numsteps, numverts;
+	object_motion_info(kg, ccl_fetch(sd, object), &numsteps, &numverts, NULL);
+	/* Figure out which steps we need to fetch and their interpolation factor. */
+	int maxstep = numsteps*2;
+	int step = min((int)(ccl_fetch(sd, time)*maxstep), maxstep-1);
+	float t = ccl_fetch(sd, time)*maxstep - step;
+	/* Find attribute. */
+	AttributeElement elem;
+	int offset = find_attribute_motion(kg, ccl_fetch(sd, object),
+	                                   ATTR_STD_MOTION_VERTEX_POSITION,
+	                                   &elem);
+	kernel_assert(offset != ATTR_STD_NOT_FOUND);
+	/* Fetch vertex coordinates. */
+	float3 verts[3], next_verts[3];
+	uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, ccl_fetch(sd, prim));
+	motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step, verts);
+	motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step+1, next_verts);
+	/* Interpolate between steps. */
+	verts[0] = (1.0f - t)*verts[0] + t*next_verts[0];
+	verts[1] = (1.0f - t)*verts[1] + t*next_verts[1];
+	verts[2] = (1.0f - t)*verts[2] + t*next_verts[2];
+	/* Compute refined position. */
+#ifdef __SUBSURFACE__
+	if(subsurface) {
+		ccl_fetch(sd, P) = motion_triangle_refine_subsurface(kg,
+		                                                     sd,
+		                                                     isect,
+		                                                     ray,
+		                                                     verts);
+	}
+	else
+#endif  /*  __SUBSURFACE__*/
+	{
+		ccl_fetch(sd, P) = motion_triangle_refine(kg, sd, isect, ray, verts);
+	}
+	/* Compute face normal. */
+	float3 Ng;
+	if(ccl_fetch(sd, object_flag) & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
+		Ng = normalize(cross(verts[2] - verts[0], verts[1] - verts[0]));
+	}
+	else {
+		Ng = normalize(cross(verts[1] - verts[0], verts[2] - verts[0]));
+	}
+	ccl_fetch(sd, Ng) = Ng;
+	ccl_fetch(sd, N) = Ng;
+	/* Compute derivatives of P w.r.t. uv. */
+#ifdef __DPDU__
+	ccl_fetch(sd, dPdu) = (verts[0] - verts[2]);
+	ccl_fetch(sd, dPdv) = (verts[1] - verts[2]);
+#endif
+	/* Compute smooth normal. */
+	if(ccl_fetch(sd, shader) & SHADER_SMOOTH_NORMAL) {
+		/* Find attribute. */
+		AttributeElement elem;
+		int offset = find_attribute_motion(kg,
+		                                   ccl_fetch(sd, object),
+		                                   ATTR_STD_MOTION_VERTEX_NORMAL,
+		                                   &elem);
+		kernel_assert(offset != ATTR_STD_NOT_FOUND);
+		/* Fetch vertex coordinates. */
+		float3 normals[3], next_normals[3];
+		motion_triangle_normals_for_step(kg, tri_vindex, offset, numverts, numsteps, step, normals);
+		motion_triangle_normals_for_step(kg, tri_vindex, offset, numverts, numsteps, step+1, next_normals);
+		/* Interpolate between steps. */
+		normals[0] = (1.0f - t)*normals[0] + t*next_normals[0];
+		normals[1] = (1.0f - t)*normals[1] + t*next_normals[1];
+		normals[2] = (1.0f - t)*normals[2] + t*next_normals[2];
+		/* Interpolate between vertices. */
+		float u = ccl_fetch(sd, u);
+		float v = ccl_fetch(sd, v);
+		float w = 1.0f - u - v;
+		ccl_fetch(sd, N) = (u*normals[0] + v*normals[1] + w*normals[2]);
+	}
+}
+
+CCL_NAMESPACE_END
+

intern/cycles/kernel/geom/geom_object.h

@@ -113,7 +113,6 @@ ccl_device_inline Transform object_fetch_transform_motion(KernelGlobals *kg, int
 ccl_device_inline Transform object_fetch_transform_motion_test(KernelGlobals *kg, int object, float time, Transform *itfm)
 {
 	int object_flag = kernel_tex_fetch(__object_flag, object);
-
 	if(object_flag & SD_OBJECT_MOTION) {
 		/* if we do motion blur */
 		Transform tfm = object_fetch_transform_motion(kg, object, time);

intern/cycles/kernel/geom/geom_primitive.h

@@ -157,8 +157,9 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals *kg, ShaderData *
 	if(is_curve_primitive) {
 		center = curve_motion_center_location(kg, sd);
 
-		if(!(ccl_fetch(sd, flag) & SD_TRANSFORM_APPLIED))
+		if(!(ccl_fetch(sd, object_flag) & SD_OBJECT_TRANSFORM_APPLIED)) {
 			object_position_transform(kg, sd, &center);
+		}
 	}
 	else
 #endif
@@ -181,7 +182,7 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals *kg, ShaderData *
 		motion_post = primitive_attribute_float3(kg, sd, desc, NULL, NULL);
 
 #ifdef __HAIR__
-		if(is_curve_primitive && (ccl_fetch(sd, flag) & SD_OBJECT_HAS_VERTEX_MOTION) == 0) {
+		if(is_curve_primitive && (ccl_fetch(sd, object_flag) & SD_OBJECT_HAS_VERTEX_MOTION) == 0) {
 			object_position_transform(kg, sd, &motion_pre);
 			object_position_transform(kg, sd, &motion_post);
 		}

intern/cycles/kernel/geom/geom_triangle.h

@@ -32,10 +32,12 @@ ccl_device_inline float3 triangle_normal(KernelGlobals *kg, ShaderData *sd)
 	const float3 v2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w+2));
 
 	/* return normal */
-	if(ccl_fetch(sd, flag) & SD_NEGATIVE_SCALE_APPLIED)
+	if(ccl_fetch(sd, object_flag) & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
 		return normalize(cross(v2 - v0, v1 - v0));
-	else
+	}
+	else {
 		return normalize(cross(v1 - v0, v2 - v0));
+	}
 }
 
 /* point and normal on triangle  */
@@ -46,20 +48,18 @@ ccl_device_inline void triangle_point_normal(KernelGlobals *kg, int object, int
 	float3 v0 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w+0));
 	float3 v1 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w+1));
 	float3 v2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w+2));
-
 	/* compute point */
 	float t = 1.0f - u - v;
 	*P = (u*v0 + v*v1 + t*v2);
-
 	/* get object flags */
 	int object_flag = kernel_tex_fetch(__object_flag, object);
-
 	/* compute normal */
-	if(object_flag & SD_NEGATIVE_SCALE_APPLIED)
+	if(object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
 		*Ng = normalize(cross(v2 - v0, v1 - v0));
-	else
+	}
+	else {
 		*Ng = normalize(cross(v1 - v0, v2 - v0));
-
+	}
 	/* shader`*/
 	*shader = kernel_tex_fetch(__tri_shader, prim);
 }

intern/cycles/kernel/geom/geom_triangle_intersect.h

@@ -108,7 +108,7 @@ ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
                                           float3 P,
                                           uint visibility,
                                           int object,
-                                          int triAddr)
+                                          int prim_addr)
 {
 	const int kx = isect_precalc->kx;
 	const int ky = isect_precalc->ky;
@@ -118,7 +118,7 @@ ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
 	const float Sz = isect_precalc->Sz;
 
 	/* Calculate vertices relative to ray origin. */
-	const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, triAddr);
+	const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
 
 #if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
 	const avxf avxf_P(P.m128, P.m128);
@@ -129,10 +129,10 @@ ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
 	const avxf AB = tri_ab - avxf_P;
 	const avxf BC = tri_bc - avxf_P;
 
-	const __m256i permuteMask = _mm256_set_epi32(0x3, kz, ky, kx, 0x3, kz, ky, kx);
+	const __m256i permute_mask = _mm256_set_epi32(0x3, kz, ky, kx, 0x3, kz, ky, kx);
 
-	const avxf AB_k = shuffle(AB, permuteMask);
-	const avxf BC_k = shuffle(BC, permuteMask);
+	const avxf AB_k = shuffle(AB, permute_mask);
+	const avxf BC_k = shuffle(BC, permute_mask);
 
 	/* Akz, Akz, Bkz, Bkz, Bkz, Bkz, Ckz, Ckz */
 	const avxf ABBC_kz = shuffle<2>(AB_k, BC_k);
@@ -155,14 +155,14 @@ ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
 	/* By, Bx, Cy, Cx, By, Bx, Ay, Ax */
 	const avxf BCBA_yx = permute<3,2,7,6,3,2,1,0>(ABBC_xy);
 
-	const avxf negMask(0,0,0,0,0x80000000, 0x80000000, 0x80000000, 0x80000000);
+	const avxf neg_mask(0,0,0,0,0x80000000, 0x80000000, 0x80000000, 0x80000000);
 
 	/* W           U                             V
 	 * (AxBy-AyBx) (BxCy-ByCx) XX XX (BxBy-ByBx) (CxAy-CyAx) XX XX
 	 */
-	const avxf WUxxxxVxx_neg = _mm256_hsub_ps(ABBC_xy * BCBA_yx, negMask /* Dont care */);
+	const avxf WUxxxxVxx_neg = _mm256_hsub_ps(ABBC_xy * BCBA_yx, neg_mask /* Dont care */);
 
-	const avxf WUVWnegWUVW = permute<0,1,5,0,0,1,5,0>(WUxxxxVxx_neg) ^ negMask;
+	const avxf WUVWnegWUVW = permute<0,1,5,0,0,1,5,0>(WUxxxxVxx_neg) ^ neg_mask;
 
 	/* Calculate scaled barycentric coordinates. */
 	float WUVW_array[4];
@@ -231,7 +231,7 @@ ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
 #ifdef __VISIBILITY_FLAG__
 	/* visibility flag test. we do it here under the assumption
 	 * that most triangles are culled by node flags */
-	if(kernel_tex_fetch(__prim_visibility, triAddr) & visibility)
+	if(kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
 #endif
 	{
 #ifdef __KERNEL_CUDA__
@@ -241,7 +241,7 @@ ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
 #endif
 		/* Normalize U, V, W, and T. */
 		const float inv_det = 1.0f / det;
-		isect->prim = triAddr;
+		isect->prim = prim_addr;
 		isect->object = object;
 		isect->type = PRIMITIVE_TRIANGLE;
 		isect->u = U * inv_det;
@@ -264,7 +264,7 @@ ccl_device_inline void triangle_intersect_subsurface(
         SubsurfaceIntersection *ss_isect,
         float3 P,
         int object,
-        int triAddr,
+        int prim_addr,
         float tmax,
         uint *lcg_state,
         int max_hits)
@@ -277,7 +277,7 @@ ccl_device_inline void triangle_intersect_subsurface(
 	const float Sz = isect_precalc->Sz;
 
 	/* Calculate vertices relative to ray origin. */
-	const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, triAddr);
+	const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
 	const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex+0),
 	             tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex+1),
 	             tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex+2);
@@ -415,7 +415,7 @@ ccl_device_inline void triangle_intersect_subsurface(
 
 	/* record intersection */
 	Intersection *isect = &ss_isect->hits[hit];
-	isect->prim = triAddr;
+	isect->prim = prim_addr;
 	isect->object = object;
 	isect->type = PRIMITIVE_TRIANGLE;
 	isect->u = U * inv_det;

intern/cycles/kernel/kernel_bake.h

@@ -320,7 +320,7 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
 	                         P, Ng, Ng,
 	                         shader, object, prim,
 	                         u, v, 1.0f, 0.5f,
-	                         !(kernel_tex_fetch(__object_flag, object) & SD_TRANSFORM_APPLIED),
+	                         !(kernel_tex_fetch(__object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED),
 	                         LAMP_NONE);
 	sd.I = sd.N;
 

intern/cycles/kernel/kernel_debug.h

@@ -18,8 +18,9 @@ CCL_NAMESPACE_BEGIN
 
 ccl_device_inline void debug_data_init(DebugData *debug_data)
 {
-	debug_data->num_bvh_traversal_steps = 0;
+	debug_data->num_bvh_traversed_nodes = 0;
 	debug_data->num_bvh_traversed_instances = 0;
+	debug_data->num_bvh_intersections = 0;
 	debug_data->num_ray_bounces = 0;
 }
 
@@ -30,16 +31,21 @@ ccl_device_inline void kernel_write_debug_passes(KernelGlobals *kg,
                                                  int sample)
 {
 	int flag = kernel_data.film.pass_flag;
-	if(flag & PASS_BVH_TRAVERSAL_STEPS) {
-		kernel_write_pass_float(buffer + kernel_data.film.pass_bvh_traversal_steps,
+	if(flag & PASS_BVH_TRAVERSED_NODES) {
+		kernel_write_pass_float(buffer + kernel_data.film.pass_bvh_traversed_nodes,
 		                        sample,
-		                        debug_data->num_bvh_traversal_steps);
+		                        debug_data->num_bvh_traversed_nodes);
 	}
 	if(flag & PASS_BVH_TRAVERSED_INSTANCES) {
 		kernel_write_pass_float(buffer + kernel_data.film.pass_bvh_traversed_instances,
 		                        sample,
 		                        debug_data->num_bvh_traversed_instances);
 	}
+	if(flag & PASS_BVH_INTERSECTIONS) {
+		kernel_write_pass_float(buffer + kernel_data.film.pass_bvh_intersections,
+		                        sample,
+		                        debug_data->num_bvh_intersections);
+	}
 	if(flag & PASS_RAY_BOUNCES) {
 		kernel_write_pass_float(buffer + kernel_data.film.pass_ray_bounces,
 		                        sample,

intern/cycles/kernel/kernel_jitter.h

@@ -149,6 +149,15 @@ ccl_device_inline uint cmj_hash(uint i, uint p)
 	return i;
 }
 
+ccl_device_inline uint cmj_hash_simple(uint i, uint p)
+{
+	i = (i ^ 61) ^ p;
+	i += i << 3;
+	i ^= i >> 4;
+	i *= 0x27d4eb2d;
+	return i;
+}
+
 ccl_device_inline float cmj_randfloat(uint i, uint p)
 {
 	return cmj_hash(i, p) * (1.0f / 4294967808.0f);

intern/cycles/kernel/kernel_light.h

@@ -767,7 +767,7 @@ ccl_device void object_transform_light_sample(KernelGlobals *kg, LightSample *ls
 {
 #ifdef __INSTANCING__
 	/* instance transform */
-	if(!(kernel_tex_fetch(__object_flag, object) & SD_TRANSFORM_APPLIED)) {
+	if(!(kernel_tex_fetch(__object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
 #  ifdef __OBJECT_MOTION__
 		Transform itfm;
 		Transform tfm = object_fetch_transform_motion_test(kg, object, time, &itfm);

intern/cycles/kernel/kernel_path.h

@@ -141,6 +141,10 @@ ccl_device void kernel_path_indirect(KernelGlobals *kg,
 #endif  /* __LAMP_MIS__ */
 
 #ifdef __VOLUME__
+		/* Sanitize volume stack. */
+		if(!hit) {
+			kernel_volume_clean_stack(kg, state->volume_stack);
+		}
 		/* volume attenuation, emission, scatter */
 		if(state->volume_stack[0].shader != SHADER_NONE) {
 			Ray volume_ray = *ray;
@@ -448,7 +452,7 @@ bool kernel_path_subsurface_scatter(
 #  ifdef __VOLUME__
 		ss_indirect->need_update_volume_stack =
 		        kernel_data.integrator.use_volumes &&
-		        ccl_fetch(sd, flag) & SD_OBJECT_INTERSECTS_VOLUME;
+		        ccl_fetch(sd, object_flag) & SD_OBJECT_INTERSECTS_VOLUME;
 #  endif  /* __VOLUME__ */
 
 		/* compute lighting with the BSDF closure */
@@ -630,8 +634,9 @@ ccl_device_inline float4 kernel_path_integrate(KernelGlobals *kg,
 
 #ifdef __KERNEL_DEBUG__
 		if(state.flag & PATH_RAY_CAMERA) {
-			debug_data.num_bvh_traversal_steps += isect.num_traversal_steps;
+			debug_data.num_bvh_traversed_nodes += isect.num_traversed_nodes;
 			debug_data.num_bvh_traversed_instances += isect.num_traversed_instances;
+			debug_data.num_bvh_intersections += isect.num_intersections;
 		}
 		debug_data.num_ray_bounces++;
 #endif  /* __KERNEL_DEBUG__ */
@@ -658,6 +663,10 @@ ccl_device_inline float4 kernel_path_integrate(KernelGlobals *kg,
 #endif  /* __LAMP_MIS__ */
 
 #ifdef __VOLUME__
+		/* Sanitize volume stack. */
+		if(!hit) {
+			kernel_volume_clean_stack(kg, state.volume_stack);
+		}
 		/* volume attenuation, emission, scatter */
 		if(state.volume_stack[0].shader != SHADER_NONE) {
 			Ray volume_ray = ray;
@@ -768,21 +777,25 @@ ccl_device_inline float4 kernel_path_integrate(KernelGlobals *kg,
 
 		/* holdout */
 #ifdef __HOLDOUT__
-		if((sd.flag & (SD_HOLDOUT|SD_HOLDOUT_MASK)) && (state.flag & PATH_RAY_CAMERA)) {
+		if(((sd.flag & SD_HOLDOUT) ||
+		    (sd.object_flag & SD_OBJECT_HOLDOUT_MASK)) &&
+		   (state.flag & PATH_RAY_CAMERA))
+		{
 			if(kernel_data.background.transparent) {
 				float3 holdout_weight;
-				
-				if(sd.flag & SD_HOLDOUT_MASK)
+				if(sd.object_flag & SD_OBJECT_HOLDOUT_MASK) {
 					holdout_weight = make_float3(1.0f, 1.0f, 1.0f);
-				else
+				}
+				else {
 					holdout_weight = shader_holdout_eval(kg, &sd);
-
+				}
 				/* any throughput is ok, should all be identical here */
 				L_transparent += average(holdout_weight*throughput);
 			}
 
-			if(sd.flag & SD_HOLDOUT_MASK)
+			if(sd.object_flag & SD_OBJECT_HOLDOUT_MASK) {
 				break;
+			}
 		}
 #endif  /* __HOLDOUT__ */
 

intern/cycles/kernel/kernel_path_branched.h

@@ -167,8 +167,9 @@ ccl_device void kernel_branched_path_subsurface_scatter(KernelGlobals *kg,
 			                                                  true);
 #ifdef __VOLUME__
 			Ray volume_ray = *ray;
-			bool need_update_volume_stack = kernel_data.integrator.use_volumes &&
-			                                ccl_fetch(sd, flag) & SD_OBJECT_INTERSECTS_VOLUME;
+			bool need_update_volume_stack =
+			        kernel_data.integrator.use_volumes &&
+			        ccl_fetch(sd, object_flag) & SD_OBJECT_INTERSECTS_VOLUME;
 #endif  /* __VOLUME__ */
 
 			/* compute lighting with the BSDF closure */
@@ -288,12 +289,17 @@ ccl_device float4 kernel_branched_path_integrate(KernelGlobals *kg, RNG *rng, in
 #endif  /* __HAIR__ */
 
 #ifdef __KERNEL_DEBUG__
-		debug_data.num_bvh_traversal_steps += isect.num_traversal_steps;
+		debug_data.num_bvh_traversed_nodes += isect.num_traversed_nodes;
 		debug_data.num_bvh_traversed_instances += isect.num_traversed_instances;
+		debug_data.num_bvh_intersections += isect.num_intersections;
 		debug_data.num_ray_bounces++;
 #endif  /* __KERNEL_DEBUG__ */
 
 #ifdef __VOLUME__
+		/* Sanitize volume stack. */
+		if(!hit) {
+			kernel_volume_clean_stack(kg, state.volume_stack);
+		}
 		/* volume attenuation, emission, scatter */
 		if(state.volume_stack[0].shader != SHADER_NONE) {
 			Ray volume_ray = ray;
@@ -468,21 +474,21 @@ ccl_device float4 kernel_branched_path_integrate(KernelGlobals *kg, RNG *rng, in
 
 		/* holdout */
 #ifdef __HOLDOUT__
-		if(sd.flag & (SD_HOLDOUT|SD_HOLDOUT_MASK)) {
+		if((sd.flag & SD_HOLDOUT) || (sd.object_flag & SD_OBJECT_HOLDOUT_MASK)) {
 			if(kernel_data.background.transparent) {
 				float3 holdout_weight;
-				
-				if(sd.flag & SD_HOLDOUT_MASK)
+				if(sd.object_flag & SD_OBJECT_HOLDOUT_MASK) {
 					holdout_weight = make_float3(1.0f, 1.0f, 1.0f);
-				else
+				}
+				else {
 					holdout_weight = shader_holdout_eval(kg, &sd);
-
+				}
 				/* any throughput is ok, should all be identical here */
 				L_transparent += average(holdout_weight*throughput);
 			}
-
-			if(sd.flag & SD_HOLDOUT_MASK)
+			if(sd.object_flag & SD_OBJECT_HOLDOUT_MASK) {
 				break;
+			}
 		}
 #endif  /* __HOLDOUT__ */
 

intern/cycles/kernel/kernel_random.h

@@ -120,13 +120,11 @@ ccl_device_forceinline float path_rng_1D(KernelGlobals *kg, ccl_addr_space RNG *
 	/* Cranly-Patterson rotation using rng seed */
 	float shift;
 
-	/* using the same *rng value to offset seems to give correlation issues,
-	 * we could hash it with the dimension but this has a performance impact,
-	 * we need to find a solution for this */
-	if(dimension & 1)
-		shift = (*rng >> 16) * (1.0f/(float)0xFFFF);
-	else
-		shift = (*rng & 0xFFFF) * (1.0f/(float)0xFFFF);
+	/* Hash rng with dimension to solve correlation issues.
+	 * See T38710, T50116.
+	 */
+	RNG tmp_rng = cmj_hash_simple(dimension, *rng);
+	shift = tmp_rng * (1.0f/(float)0xFFFFFFFF);
 
 	return r + shift - floorf(r + shift);
 #endif

intern/cycles/kernel/kernel_shader.h

@@ -38,7 +38,7 @@ CCL_NAMESPACE_BEGIN
 #ifdef __OBJECT_MOTION__
 ccl_device void shader_setup_object_transforms(KernelGlobals *kg, ShaderData *sd, float time)
 {
-	if(ccl_fetch(sd, flag) & SD_OBJECT_MOTION) {
+	if(ccl_fetch(sd, object_flag) & SD_OBJECT_MOTION) {
 		ccl_fetch(sd, ob_tfm) = object_fetch_transform_motion(kg, ccl_fetch(sd, object), time);
 		ccl_fetch(sd, ob_itfm) = transform_quick_inverse(ccl_fetch(sd, ob_tfm));
 	}
@@ -59,7 +59,9 @@ ccl_device_noinline void shader_setup_from_ray(KernelGlobals *kg,
 #endif
 
 	ccl_fetch(sd, type) = isect->type;
-	ccl_fetch(sd, flag) = kernel_tex_fetch(__object_flag, ccl_fetch(sd, object));
+	ccl_fetch(sd, flag) = 0;
+	ccl_fetch(sd, object_flag) = kernel_tex_fetch(__object_flag,
+	                                              ccl_fetch(sd, object));
 
 	/* matrices and time */
 #ifdef __OBJECT_MOTION__
@@ -160,10 +162,11 @@ void shader_setup_from_subsurface(
         const Intersection *isect,
         const Ray *ray)
 {
-	bool backfacing = sd->flag & SD_BACKFACING;
+	const bool backfacing = sd->flag & SD_BACKFACING;
 
 	/* object, matrices, time, ray_length stay the same */
-	sd->flag = kernel_tex_fetch(__object_flag, sd->object);
+	sd->flag = 0;
+	sd->object_flag = kernel_tex_fetch(__object_flag, sd->object);
 	sd->prim = kernel_tex_fetch(__prim_index, isect->prim);
 	sd->type = isect->type;
 
@@ -271,8 +274,10 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals *kg,
 	ccl_fetch(sd, ray_length) = t;
 
 	ccl_fetch(sd, flag) = kernel_tex_fetch(__shader_flag, (ccl_fetch(sd, shader) & SHADER_MASK)*SHADER_SIZE);
+	ccl_fetch(sd, object_flag) = 0;
 	if(ccl_fetch(sd, object) != OBJECT_NONE) {
-		ccl_fetch(sd, flag) |= kernel_tex_fetch(__object_flag, ccl_fetch(sd, object));
+		ccl_fetch(sd, object_flag) |= kernel_tex_fetch(__object_flag,
+		                                               ccl_fetch(sd, object));
 
 #ifdef __OBJECT_MOTION__
 		shader_setup_object_transforms(kg, sd, time);
@@ -298,7 +303,7 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals *kg,
 			ccl_fetch(sd, N) = triangle_smooth_normal(kg, ccl_fetch(sd, prim), ccl_fetch(sd, u), ccl_fetch(sd, v));
 
 #ifdef __INSTANCING__
-			if(!(ccl_fetch(sd, flag) & SD_TRANSFORM_APPLIED)) {
+			if(!(ccl_fetch(sd, object_flag) & SD_OBJECT_TRANSFORM_APPLIED)) {
 				object_normal_transform_auto(kg, sd, &ccl_fetch(sd, N));
 			}
 #endif
@@ -309,7 +314,7 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals *kg,
 		triangle_dPdudv(kg, ccl_fetch(sd, prim), &ccl_fetch(sd, dPdu), &ccl_fetch(sd, dPdv));
 
 #  ifdef __INSTANCING__
-		if(!(ccl_fetch(sd, flag) & SD_TRANSFORM_APPLIED)) {
+		if(!(ccl_fetch(sd, object_flag) & SD_OBJECT_TRANSFORM_APPLIED)) {
 			object_dir_transform_auto(kg, sd, &ccl_fetch(sd, dPdu));
 			object_dir_transform_auto(kg, sd, &ccl_fetch(sd, dPdv));
 		}
@@ -364,7 +369,7 @@ ccl_device void shader_setup_from_displace(KernelGlobals *kg, ShaderData *sd,
 	                         P, Ng, I,
 	                         shader, object, prim,
 	                         u, v, 0.0f, 0.5f,
-	                         !(kernel_tex_fetch(__object_flag, object) & SD_TRANSFORM_APPLIED),
+	                         !(kernel_tex_fetch(__object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED),
 	                         LAMP_NONE);
 }
 
@@ -379,6 +384,7 @@ ccl_device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderDat
 	ccl_fetch(sd, I) = -ray->D;
 	ccl_fetch(sd, shader) = kernel_data.background.surface_shader;
 	ccl_fetch(sd, flag) = kernel_tex_fetch(__shader_flag, (ccl_fetch(sd, shader) & SHADER_MASK)*SHADER_SIZE);
+	ccl_fetch(sd, object_flag) = 0;
 #ifdef __OBJECT_MOTION__
 	ccl_fetch(sd, time) = ray->time;
 #endif
@@ -420,6 +426,7 @@ ccl_device_inline void shader_setup_from_volume(KernelGlobals *kg, ShaderData *s
 	sd->I = -ray->D;
 	sd->shader = SHADER_NONE;
 	sd->flag = 0;
+	sd->object_flag = 0;
 #ifdef __OBJECT_MOTION__
 	sd->time = ray->time;
 #endif
@@ -1027,6 +1034,7 @@ ccl_device_inline void shader_eval_volume(KernelGlobals *kg,
 	sd->num_closure = 0;
 	sd->num_closure_extra = 0;
 	sd->flag = 0;
+	sd->object_flag = 0;
 
 	for(int i = 0; stack[i].shader != SHADER_NONE; i++) {
 		/* setup shaderdata from stack. it's mostly setup already in
@@ -1034,11 +1042,12 @@ ccl_device_inline void shader_eval_volume(KernelGlobals *kg,
 		sd->object = stack[i].object;
 		sd->shader = stack[i].shader;
 
-		sd->flag &= ~(SD_SHADER_FLAGS|SD_OBJECT_FLAGS);
+		sd->flag &= ~SD_SHADER_FLAGS;
 		sd->flag |= kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*SHADER_SIZE);
+		sd->object_flag &= ~SD_OBJECT_FLAGS;
 
 		if(sd->object != OBJECT_NONE) {
-			sd->flag |= kernel_tex_fetch(__object_flag, sd->object);
+			sd->object_flag |= kernel_tex_fetch(__object_flag, sd->object);
 
 #ifdef __OBJECT_MOTION__
 			/* todo: this is inefficient for motion blur, we should be

intern/cycles/kernel/kernel_types.h

@@ -192,6 +192,9 @@ CCL_NAMESPACE_BEGIN
 #ifdef __NO_PATCH_EVAL__
 #  undef __PATCH_EVAL__
 #endif
+#ifdef __NO_TRANSPARENT__
+#  undef __TRANSPARENT_SHADOWS__
+#endif
 
 /* Random Numbers */
 
@@ -342,9 +345,10 @@ typedef enum PassType {
 	PASS_SUBSURFACE_COLOR = (1 << 24),
 	PASS_LIGHT = (1 << 25), /* no real pass, used to force use_light_pass */
 #ifdef __KERNEL_DEBUG__
-	PASS_BVH_TRAVERSAL_STEPS = (1 << 26),
+	PASS_BVH_TRAVERSED_NODES = (1 << 26),
 	PASS_BVH_TRAVERSED_INSTANCES = (1 << 27),
-	PASS_RAY_BOUNCES = (1 << 28),
+	PASS_BVH_INTERSECTIONS = (1 << 28),
+	PASS_RAY_BOUNCES = (1 << 29),
 #endif
 } PassType;
 
@@ -539,35 +543,38 @@ typedef ccl_addr_space struct Intersection {
 	int type;
 
 #ifdef __KERNEL_DEBUG__
-	int num_traversal_steps;
+	int num_traversed_nodes;
 	int num_traversed_instances;
+	int num_intersections;
 #endif
 } Intersection;
 
 /* Primitives */
 
 typedef enum PrimitiveType {
-	PRIMITIVE_NONE = 0,
-	PRIMITIVE_TRIANGLE = 1,
-	PRIMITIVE_MOTION_TRIANGLE = 2,
-	PRIMITIVE_CURVE = 4,
-	PRIMITIVE_MOTION_CURVE = 8,
-	/* Lamp primitive is not included below on purpose, since it is no real traceable primitive */
-	PRIMITIVE_LAMP = 16,
+	PRIMITIVE_NONE            = 0,
+	PRIMITIVE_TRIANGLE        = (1 << 0),
+	PRIMITIVE_MOTION_TRIANGLE = (1 << 1),
+	PRIMITIVE_CURVE           = (1 << 2),
+	PRIMITIVE_MOTION_CURVE    = (1 << 3),
+	/* Lamp primitive is not included below on purpose,
+	 * since it is no real traceable primitive.
+	 */
+	PRIMITIVE_LAMP            = (1 << 4),
 
 	PRIMITIVE_ALL_TRIANGLE = (PRIMITIVE_TRIANGLE|PRIMITIVE_MOTION_TRIANGLE),
 	PRIMITIVE_ALL_CURVE = (PRIMITIVE_CURVE|PRIMITIVE_MOTION_CURVE),
 	PRIMITIVE_ALL_MOTION = (PRIMITIVE_MOTION_TRIANGLE|PRIMITIVE_MOTION_CURVE),
 	PRIMITIVE_ALL = (PRIMITIVE_ALL_TRIANGLE|PRIMITIVE_ALL_CURVE),
 
-	/* Total number of different primitives.
+	/* Total number of different traceable primitives.
 	 * NOTE: This is an actual value, not a bitflag.
 	 */
 	PRIMITIVE_NUM_TOTAL = 4,
 } PrimitiveType;
 
-#define PRIMITIVE_PACK_SEGMENT(type, segment) ((segment << 16) | type)
-#define PRIMITIVE_UNPACK_SEGMENT(type) (type >> 16)
+#define PRIMITIVE_PACK_SEGMENT(type, segment) ((segment << PRIMITIVE_NUM_TOTAL) | (type))
+#define PRIMITIVE_UNPACK_SEGMENT(type) (type >> PRIMITIVE_NUM_TOTAL)
 
 /* Attributes */
 
@@ -685,56 +692,108 @@ typedef enum ShaderContext {
 /* Shader Data
  *
  * Main shader state at a point on the surface or in a volume. All coordinates
- * are in world space. */
+ * are in world space.
+ */
 
 enum ShaderDataFlag {
-	/* runtime flags */
-	SD_BACKFACING      = (1 << 0),   /* backside of surface? */
-	SD_EMISSION        = (1 << 1),   /* have emissive closure? */
-	SD_BSDF            = (1 << 2),   /* have bsdf closure? */
-	SD_BSDF_HAS_EVAL   = (1 << 3),   /* have non-singular bsdf closure? */
-	SD_BSSRDF          = (1 << 4),   /* have bssrdf */
-	SD_HOLDOUT         = (1 << 5),   /* have holdout closure? */
-	SD_ABSORPTION      = (1 << 6),   /* have volume absorption closure? */
-	SD_SCATTER         = (1 << 7),   /* have volume phase closure? */
-	SD_AO              = (1 << 8),   /* have ao closure? */
-	SD_TRANSPARENT     = (1 << 9),  /* have transparent closure? */
+	/* Runtime flags. */
+
+	/* Set when ray hits backside of surface. */
+	SD_BACKFACING      = (1 << 0),
+	/* Shader has emissive closure. */
+	SD_EMISSION        = (1 << 1),
+	/* Shader has BSDF closure. */
+	SD_BSDF            = (1 << 2),
+	/* Shader has non-singular BSDF closure. */
+	SD_BSDF_HAS_EVAL   = (1 << 3),
+	/* Shader has BSSRDF closure. */
+	SD_BSSRDF          = (1 << 4),
+	/* Shader has holdout closure. */
+	SD_HOLDOUT         = (1 << 5),
+	/* Shader has volume absorption closure. */
+	SD_ABSORPTION      = (1 << 6),
+	/* Shader has have volume phase (scatter) closure. */
+	SD_SCATTER         = (1 << 7),
+	/* Shader has AO closure. */
+	SD_AO              = (1 << 8),
+	/* Shader has transparent closure. */
+	SD_TRANSPARENT     = (1 << 9),
+	/* BSDF requires LCG for evaluation. */
 	SD_BSDF_NEEDS_LCG  = (1 << 10),
 
-	SD_CLOSURE_FLAGS = (SD_EMISSION|SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSSRDF|
-	                    SD_HOLDOUT|SD_ABSORPTION|SD_SCATTER|SD_AO|
+	SD_CLOSURE_FLAGS = (SD_EMISSION |
+	                    SD_BSDF |
+	                    SD_BSDF_HAS_EVAL |
+	                    SD_BSSRDF |
+	                    SD_HOLDOUT |
+	                    SD_ABSORPTION |
+	                    SD_SCATTER |
+	                    SD_AO |
 	                    SD_BSDF_NEEDS_LCG),
 
-	/* shader flags */
-	SD_USE_MIS                = (1 << 12),  /* direct light sample */
-	SD_HAS_TRANSPARENT_SHADOW = (1 << 13),  /* has transparent shadow */
-	SD_HAS_VOLUME             = (1 << 14),  /* has volume shader */
-	SD_HAS_ONLY_VOLUME        = (1 << 15),  /* has only volume shader, no surface */
-	SD_HETEROGENEOUS_VOLUME   = (1 << 16),  /* has heterogeneous volume */
-	SD_HAS_BSSRDF_BUMP        = (1 << 17),  /* bssrdf normal uses bump */
-	SD_VOLUME_EQUIANGULAR     = (1 << 18),  /* use equiangular sampling */
-	SD_VOLUME_MIS             = (1 << 19),  /* use multiple importance sampling */
-	SD_VOLUME_CUBIC           = (1 << 20),  /* use cubic interpolation for voxels */
-	SD_HAS_BUMP               = (1 << 21),  /* has data connected to the displacement input */
-	SD_HAS_DISPLACEMENT       = (1 << 22),  /* has true displacement */
-	SD_HAS_CONSTANT_EMISSION  = (1 << 23),  /* has constant emission (value stored in __shader_flag) */
+	/* Shader flags. */
 
-	SD_SHADER_FLAGS = (SD_USE_MIS|SD_HAS_TRANSPARENT_SHADOW|SD_HAS_VOLUME|
-	                   SD_HAS_ONLY_VOLUME|SD_HETEROGENEOUS_VOLUME|
-	                   SD_HAS_BSSRDF_BUMP|SD_VOLUME_EQUIANGULAR|SD_VOLUME_MIS|
-	                   SD_VOLUME_CUBIC|SD_HAS_BUMP|SD_HAS_DISPLACEMENT|SD_HAS_CONSTANT_EMISSION),
+	/* direct light sample */
+	SD_USE_MIS                = (1 << 16),
+	/* Has transparent shadow. */
+	SD_HAS_TRANSPARENT_SHADOW = (1 << 17),
+	/* Has volume shader. */
+	SD_HAS_VOLUME             = (1 << 18),
+	/* Has only volume shader, no surface. */
+	SD_HAS_ONLY_VOLUME        = (1 << 19),
+	/* Has heterogeneous volume. */
+	SD_HETEROGENEOUS_VOLUME   = (1 << 20),
+	/* BSSRDF normal uses bump. */
+	SD_HAS_BSSRDF_BUMP        = (1 << 21),
+	/* Use equiangular volume sampling */
+	SD_VOLUME_EQUIANGULAR     = (1 << 22),
+	/* Use multiple importance volume sampling. */
+	SD_VOLUME_MIS             = (1 << 23),
+	/* Use cubic interpolation for voxels. */
+	SD_VOLUME_CUBIC           = (1 << 24),
+	/* Has data connected to the displacement input. */
+	SD_HAS_BUMP               = (1 << 25),
+	/* Has true displacement. */
+	SD_HAS_DISPLACEMENT       = (1 << 26),
+	/* Has constant emission (value stored in __shader_flag) */
+	SD_HAS_CONSTANT_EMISSION  = (1 << 27),
 
-	/* object flags */
-	SD_HOLDOUT_MASK             = (1 << 24),  /* holdout for camera rays */
-	SD_OBJECT_MOTION            = (1 << 25),  /* has object motion blur */
-	SD_TRANSFORM_APPLIED        = (1 << 26),  /* vertices have transform applied */
-	SD_NEGATIVE_SCALE_APPLIED   = (1 << 27),  /* vertices have negative scale applied */
-	SD_OBJECT_HAS_VOLUME        = (1 << 28),  /* object has a volume shader */
-	SD_OBJECT_INTERSECTS_VOLUME = (1 << 29),  /* object intersects AABB of an object with volume shader */
-	SD_OBJECT_HAS_VERTEX_MOTION = (1 << 30),  /* has position for motion vertices */
+	SD_SHADER_FLAGS = (SD_USE_MIS |
+	                   SD_HAS_TRANSPARENT_SHADOW |
+	                   SD_HAS_VOLUME |
+	                   SD_HAS_ONLY_VOLUME |
+	                   SD_HETEROGENEOUS_VOLUME|
+	                   SD_HAS_BSSRDF_BUMP |
+	                   SD_VOLUME_EQUIANGULAR |
+	                   SD_VOLUME_MIS |
+	                   SD_VOLUME_CUBIC |
+	                   SD_HAS_BUMP |
+	                   SD_HAS_DISPLACEMENT |
+	                   SD_HAS_CONSTANT_EMISSION)
+};
 
-	SD_OBJECT_FLAGS = (SD_HOLDOUT_MASK|SD_OBJECT_MOTION|SD_TRANSFORM_APPLIED|
-	                   SD_NEGATIVE_SCALE_APPLIED|SD_OBJECT_HAS_VOLUME|
+	/* Object flags. */
+enum ShaderDataObjectFlag {
+	/* Holdout for camera rays. */
+	SD_OBJECT_HOLDOUT_MASK           = (1 << 0),
+	/* Has object motion blur. */
+	SD_OBJECT_MOTION                 = (1 << 1),
+	/* Vertices have transform applied. */
+	SD_OBJECT_TRANSFORM_APPLIED      = (1 << 2),
+	/* Vertices have negative scale applied. */
+	SD_OBJECT_NEGATIVE_SCALE_APPLIED = (1 << 3),
+	/* Object has a volume shader. */
+	SD_OBJECT_HAS_VOLUME             = (1 << 4),
+	/* Object intersects AABB of an object with volume shader. */
+	SD_OBJECT_INTERSECTS_VOLUME      = (1 << 5),
+	/* Has position for motion vertices. */
+	SD_OBJECT_HAS_VERTEX_MOTION      = (1 << 6),
+
+	SD_OBJECT_FLAGS = (SD_OBJECT_HOLDOUT_MASK |
+	                   SD_OBJECT_MOTION |
+	                   SD_OBJECT_TRANSFORM_APPLIED |
+	                   SD_OBJECT_NEGATIVE_SCALE_APPLIED |
+	                   SD_OBJECT_HAS_VOLUME |
 	                   SD_OBJECT_INTERSECTS_VOLUME)
 };
 
@@ -773,6 +832,8 @@ typedef ccl_addr_space struct ShaderData {
 	ccl_soa_member(int, shader);
 	/* booleans describing shader, see ShaderDataFlag */
 	ccl_soa_member(int, flag);
+	/* booleans describing object of the shader, see ShaderDataObjectFlag */
+	ccl_soa_member(int, object_flag);
 
 	/* primitive id if there is one, ~0 otherwise */
 	ccl_soa_member(int, prim);
@@ -1035,10 +1096,10 @@ typedef struct KernelFilm {
 	float mist_falloff;
 
 #ifdef __KERNEL_DEBUG__
-	int pass_bvh_traversal_steps;
+	int pass_bvh_traversed_nodes;
 	int pass_bvh_traversed_instances;
+	int pass_bvh_intersections;
 	int pass_ray_bounces;
-	int pass_pad3;
 #endif
 } KernelFilm;
 static_assert_align(KernelFilm, 16);
@@ -1183,10 +1244,9 @@ static_assert_align(KernelData, 16);
  * really important here.
  */
 typedef ccl_addr_space struct DebugData {
-	// Total number of BVH node traversal steps and primitives intersections
-	// for the camera rays.
-	int num_bvh_traversal_steps;
+	int num_bvh_traversed_nodes;
 	int num_bvh_traversed_instances;
+	int num_bvh_intersections;
 	int num_ray_bounces;
 } DebugData;
 #endif

intern/cycles/kernel/kernel_volume.h

@@ -245,11 +245,18 @@ ccl_device float kernel_volume_equiangular_sample(Ray *ray, float3 light_P, floa
 	float t = ray->t;
 
 	float delta = dot((light_P - ray->P) , ray->D);
-	float D = sqrtf(len_squared(light_P - ray->P) - delta * delta);
+	float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta);
+	if(UNLIKELY(D == 0.0f)) {
+		*pdf = 0.0f;
+		return 0.0f;
+	}
 	float theta_a = -atan2f(delta, D);
 	float theta_b = atan2f(t - delta, D);
 	float t_ = D * tanf((xi * theta_b) + (1 - xi) * theta_a);
-
+	if(UNLIKELY(theta_b == theta_a)) {
+		*pdf = 0.0f;
+		return 0.0f;
+	}
 	*pdf = D / ((theta_b - theta_a) * (D * D + t_ * t_));
 
 	return min(t, delta + t_); /* min is only for float precision errors */
@@ -258,13 +265,19 @@ ccl_device float kernel_volume_equiangular_sample(Ray *ray, float3 light_P, floa
 ccl_device float kernel_volume_equiangular_pdf(Ray *ray, float3 light_P, float sample_t)
 {
 	float delta = dot((light_P - ray->P) , ray->D);
-	float D = sqrtf(len_squared(light_P - ray->P) - delta * delta);
+	float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta);
+	if(UNLIKELY(D == 0.0f)) {
+		return 0.0f;
+	}
 
 	float t = ray->t;
 	float t_ = sample_t - delta;
 
 	float theta_a = -atan2f(delta, D);
 	float theta_b = atan2f(t - delta, D);
+	if(UNLIKELY(theta_b == theta_a)) {
+		return 0.0f;
+	}
 
 	float pdf = D / ((theta_b - theta_a) * (D * D + t_ * t_));
 
@@ -569,17 +582,12 @@ ccl_device VolumeIntegrateResult kernel_volume_integrate_heterogeneous_distance(
 ccl_device_noinline VolumeIntegrateResult kernel_volume_integrate(KernelGlobals *kg,
 	PathState *state, ShaderData *sd, Ray *ray, PathRadiance *L, float3 *throughput, RNG *rng, bool heterogeneous)
 {
-	/* workaround to fix correlation bug in T38710, can find better solution
-	 * in random number generator later, for now this is done here to not impact
-	 * performance of rendering without volumes */
-	RNG tmp_rng = cmj_hash(*rng, state->rng_offset);
-
 	shader_setup_from_volume(kg, sd, ray);
 
 	if(heterogeneous)
-		return kernel_volume_integrate_heterogeneous_distance(kg, state, ray, sd, L, throughput, &tmp_rng);
+		return kernel_volume_integrate_heterogeneous_distance(kg, state, ray, sd, L, throughput, rng);
 	else
-		return kernel_volume_integrate_homogeneous(kg, state, ray, sd, L, throughput, &tmp_rng, true);
+		return kernel_volume_integrate_homogeneous(kg, state, ray, sd, L, throughput, rng, true);
 }
 
 /* Decoupled Volume Sampling
@@ -958,6 +966,9 @@ ccl_device VolumeIntegrateResult kernel_volume_decoupled_scatter(
 			mis_weight = 2.0f*power_heuristic(pdf, distance_pdf);
 		}
 	}
+	if(sample_t < 1e-6f) {
+		return VOLUME_PATH_SCATTERED;
+	}
 
 	/* compute transmittance up to this step */
 	if(step != segment->steps)
@@ -1251,4 +1262,30 @@ ccl_device void kernel_volume_stack_update_for_subsurface(KernelGlobals *kg,
 }
 #endif
 
+/* Clean stack after the last bounce.
+ *
+ * It is expected that all volumes are closed manifolds, so at the time when ray
+ * hits nothing (for example, it is a last bounce which goes to environment) the
+ * only expected volume in the stack is the world's one. All the rest volume
+ * entries should have been exited already.
+ *
+ * This isn't always true because of ray intersection precision issues, which
+ * could lead us to an infinite non-world volume in the stack, causing render
+ * artifacts.
+ *
+ * Use this function after the last bounce to get rid of all volumes apart from
+ * the world's one after the last bounce to avoid render artifacts.
+ */
+ccl_device_inline void kernel_volume_clean_stack(KernelGlobals *kg,
+                                                 VolumeStack *volume_stack)
+{
+	if(kernel_data.background.volume_shader != SHADER_NONE) {
+		/* Keep the world's volume in stack. */
+		volume_stack[1].shader = SHADER_NONE;
+	}
+	else {
+		volume_stack[0].shader = SHADER_NONE;
+	}
+}
+
 CCL_NAMESPACE_END

intern/cycles/kernel/osl/osl_services.cpp

@@ -102,6 +102,8 @@ ustring OSLRenderServices::u_curve_tangent_normal("geom:curve_tangent_normal");
 #endif
 ustring OSLRenderServices::u_path_ray_length("path:ray_length");
 ustring OSLRenderServices::u_path_ray_depth("path:ray_depth");
+ustring OSLRenderServices::u_path_diffuse_depth("path:diffuse_depth");
+ustring OSLRenderServices::u_path_glossy_depth("path:glossy_depth");
 ustring OSLRenderServices::u_path_transparent_depth("path:transparent_depth");
 ustring OSLRenderServices::u_path_transmission_depth("path:transmission_depth");
 ustring OSLRenderServices::u_trace("trace");
@@ -710,7 +712,7 @@ bool OSLRenderServices::get_object_standard_attribute(KernelGlobals *kg, ShaderD
 		else
 			motion_triangle_vertices(kg, sd->object, sd->prim, sd->time, P);
 
-		if(!(sd->flag & SD_TRANSFORM_APPLIED)) {
+		if(!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
 			object_position_transform(kg, sd, &P[0]);
 			object_position_transform(kg, sd, &P[1]);
 			object_position_transform(kg, sd, &P[2]);
@@ -759,6 +761,24 @@ bool OSLRenderServices::get_background_attribute(KernelGlobals *kg, ShaderData *
 		int f = state->bounce;
 		return set_attribute_int(f, type, derivatives, val);
 	}
+	else if(name == u_path_diffuse_depth) {
+		/* Diffuse Ray Depth */
+		PathState *state = sd->osl_path_state;
+		int f = state->diffuse_bounce;
+		return set_attribute_int(f, type, derivatives, val);
+	}
+	else if(name == u_path_glossy_depth) {
+		/* Glossy Ray Depth */
+		PathState *state = sd->osl_path_state;
+		int f = state->glossy_bounce;
+		return set_attribute_int(f, type, derivatives, val);
+	}
+	else if(name == u_path_transmission_depth) {
+		/* Transmission Ray Depth */
+		PathState *state = sd->osl_path_state;
+		int f = state->transmission_bounce;
+		return set_attribute_int(f, type, derivatives, val);
+	}
 	else if(name == u_path_transparent_depth) {
 		/* Transparent Ray Depth */
 		PathState *state = sd->osl_path_state;

intern/cycles/kernel/osl/osl_services.h

@@ -165,6 +165,8 @@ public:
 	static ustring u_curve_tangent_normal;
 	static ustring u_path_ray_length;
 	static ustring u_path_ray_depth;
+	static ustring u_path_diffuse_depth;
+	static ustring u_path_glossy_depth;
 	static ustring u_path_transparent_depth;
 	static ustring u_path_transmission_depth;
 	static ustring u_trace;

intern/cycles/kernel/shaders/node_light_path.osl

@@ -27,6 +27,8 @@ shader node_light_path(
 	output float IsVolumeScatterRay = 0.0,
 	output float RayLength = 0.0,
 	output float RayDepth = 0.0,
+	output float DiffuseDepth = 0.0,
+	output float GlossyDepth = 0.0,
 	output float TransparentDepth = 0.0,
 	output float TransmissionDepth = 0.0)
 {
@@ -45,6 +47,14 @@ shader node_light_path(
 	getattribute("path:ray_depth", ray_depth);
 	RayDepth = (float)ray_depth;
 
+	int diffuse_depth;
+	getattribute("path:diffuse_depth", diffuse_depth);
+	DiffuseDepth = (float)diffuse_depth;
+
+	int glossy_depth;
+	getattribute("path:glossy_depth", glossy_depth);
+	GlossyDepth = (float)glossy_depth;
+
 	int transparent_depth;
 	getattribute("path:transparent_depth", transparent_depth);
 	TransparentDepth = (float)transparent_depth;

intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h

@@ -137,22 +137,22 @@ ccl_device void kernel_holdout_emission_blurring_pathtermination_ao(
 
 		/* holdout */
 #ifdef __HOLDOUT__
-		if((ccl_fetch(sd, flag) & (SD_HOLDOUT|SD_HOLDOUT_MASK)) &&
+		if(((ccl_fetch(sd, flag) & SD_HOLDOUT) ||
+		    (ccl_fetch(sd, object_flag) & SD_OBJECT_HOLDOUT_MASK)) &&
 		   (state->flag & PATH_RAY_CAMERA))
 		{
 			if(kernel_data.background.transparent) {
 				float3 holdout_weight;
-
-				if(ccl_fetch(sd, flag) & SD_HOLDOUT_MASK)
+				if(ccl_fetch(sd, object_flag) & SD_OBJECT_HOLDOUT_MASK) {
 					holdout_weight = make_float3(1.0f, 1.0f, 1.0f);
-				else
+				}
+				else {
 					holdout_weight = shader_holdout_eval(kg, sd);
-
+				}
 				/* any throughput is ok, should all be identical here */
 				L_transparent_coop[ray_index] += average(holdout_weight*throughput);
 			}
-
-			if(ccl_fetch(sd, flag) & SD_HOLDOUT_MASK) {
+			if(ccl_fetch(sd, object_flag) & SD_OBJECT_HOLDOUT_MASK) {
 				ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER);
 				*enqueue_flag = 1;
 			}

intern/cycles/kernel/split/kernel_scene_intersect.h

@@ -116,8 +116,9 @@ ccl_device void kernel_scene_intersect(
 
 #ifdef __KERNEL_DEBUG__
 	if(state.flag & PATH_RAY_CAMERA) {
-		debug_data->num_bvh_traversal_steps += isect->num_traversal_steps;
+		debug_data->num_bvh_traversed_nodes += isect->num_traversed_nodes;
 		debug_data->num_bvh_traversed_instances += isect->num_traversed_instances;
+		debug_data->num_bvh_intersections += isect->num_intersections;
 	}
 	debug_data->num_ray_bounces++;
 #endif

intern/cycles/kernel/svm/svm_light_path.h

@@ -34,6 +34,8 @@ ccl_device void svm_node_light_path(ShaderData *sd, ccl_addr_space PathState *st
 		case NODE_LP_backfacing: info = (ccl_fetch(sd, flag) & SD_BACKFACING)? 1.0f: 0.0f; break;
 		case NODE_LP_ray_length: info = ccl_fetch(sd, ray_length); break;
 		case NODE_LP_ray_depth: info = (float)state->bounce; break;
+		case NODE_LP_ray_diffuse: info = (float)state->diffuse_bounce; break;
+		case NODE_LP_ray_glossy: info = (float)state->glossy_bounce; break;
 		case NODE_LP_ray_transparent: info = (float)state->transparent_bounce; break;
 		case NODE_LP_ray_transmission: info = (float)state->transmission_bounce; break;
 	}