Merge branch 'master' into blender2.8

Conflicts:
	source/blender/blenkernel/intern/library.c
	source/blender/blenkernel/intern/particle.c

59 files changed, 1952 insertions, 354 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index d2916283091..281aa52cae3 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -397,6 +397,7 @@ option(WITH_CYCLES					"Enable Cycles Render Engine" ON)
 option(WITH_CYCLES_STANDALONE		"Build Cycles standalone application" OFF)
 option(WITH_CYCLES_STANDALONE_GUI	"Build Cycles standalone with GUI" OFF)
 option(WITH_CYCLES_OSL				"Build Cycles with OSL support" ${_init_CYCLES_OSL})
+option(WITH_CYCLES_OPENSUBDIV		"Build Cycles with OpenSubdiv support" ON)
 option(WITH_CYCLES_CUDA_BINARIES	"Build Cycles CUDA binaries" OFF)
 set(CYCLES_CUDA_BINARIES_ARCH sm_20 sm_21 sm_30 sm_35 sm_37 sm_50 sm_52 CACHE STRING "CUDA architectures to build binaries for")
 mark_as_advanced(CYCLES_CUDA_BINARIES_ARCH)
@@ -1101,7 +1102,6 @@ if(UNIX AND NOT APPLE)
 	endif()
 
 	if(WITH_ALEMBIC)
-		set(ALEMBIC_ROOT_DIR ${LIBDIR}/alembic)
 		find_package_wrapper(Alembic)
 
 		if(WITH_ALEMBIC_HDF5)
@@ -1350,10 +1350,8 @@ elseif(WIN32)
 		# MSVC11 needs _ALLOW_KEYWORD_MACROS to build
 		add_definitions(-D_ALLOW_KEYWORD_MACROS)
 
-		if(CMAKE_CL_64)
-			# We want to support Vista level ABI for x64
-			add_definitions(-D_WIN32_WINNT=0x600)
-		endif()
+		# We want to support Vista level ABI 
+		add_definitions(-D_WIN32_WINNT=0x600)
 
 		# Make cmake find the msvc redistributables
 		set(CMAKE_INSTALL_SYSTEM_RUNTIME_LIBS_SKIP TRUE)
@@ -1685,18 +1683,11 @@ elseif(WIN32)
 		endif()
 
 		if(WITH_ALEMBIC)
-			set(ALEMBIC_ROOT_DIR ${LIBDIR}/alembic)
-			find_package(Alembic)
-
-			if(WITH_ALEMBIC_HDF5)
-				set(HDF5_ROOT_DIR ${LIBDIR}/hdf5)
-				find_package(HDF5)
-			endif()
-
-			if(NOT ALEMBIC_FOUND OR (WITH_ALEMBIC_HDF5 AND NOT HDF5_FOUND))
-				set(WITH_ALEMBIC OFF)
-				set(WITH_ALEMBIC_HDF5 OFF)
-			endif()
+			set(ALEMBIC ${LIBDIR}/alembic)
+			set(ALEMBIC_INCLUDE_DIR ${ALEMBIC}/include)
+			set(ALEMBIC_INCLUDE_DIRS ${ALEMBIC_INCLUDE_DIR})
+			set(ALEMBIC_LIBPATH ${ALEMBIC}/lib)
+			set(ALEMBIC_LIBRARIES optimized alembic debug alembic_d)
 		endif()
 
 		if(WITH_MOD_CLOTH_ELTOPO)
@@ -1999,14 +1990,16 @@ elseif(WIN32)
 		endif()
 
 		if(WITH_ALEMBIC)
-			set(ALEMBIC_ROOT_DIR ${LIBDIR}/alembic)
+			# TODO(sergey): For until someone drops by and compiles libraries for
+			# MinGW we allow users to compile their own Alembic library and use
+			# that via find_package(),
+			#
+			# Once precompiled libraries are there we'll use hardcoded locations.
 			find_package_wrapper(Alembic)
-
 			if(WITH_ALEMBIC_HDF5)
 				set(HDF5_ROOT_DIR ${LIBDIR}/hdf5)
 				find_package_wrapper(HDF5)
 			endif()
-
 			if(NOT ALEMBIC_FOUND OR (WITH_ALEMBIC_HDF5 AND NOT HDF5_FOUND))
 				set(WITH_ALEMBIC OFF)
 				set(WITH_ALEMBIC_HDF5 OFF)
@@ -2098,17 +2091,11 @@ elseif(APPLE)
 	endif()
 
 	if(WITH_ALEMBIC)
-		set(ALEMBIC_ROOT_DIR ${LIBDIR}/alembic)
-		find_package(Alembic)
-		if(WITH_ALEMBIC_HDF5)
-			set(HDF5_ROOT_DIR ${LIBDIR}/hdf5)
-			find_package(HDF5)
-		endif()
-
-		if(NOT ALEMBIC_FOUND OR (WITH_ALEMBIC_HDF5 AND NOT HDF5_FOUND))
-			set(WITH_ALEMBIC OFF)
-			set(WITH_ALEMBIC_HDF5 OFF)
-		endif()
+		set(ALEMBIC ${LIBDIR}/alembic)
+		set(ALEMBIC_INCLUDE_DIR ${ALEMBIC}/include)
+		set(ALEMBIC_INCLUDE_DIRS ${ALEMBIC_INCLUDE_DIR})
+		set(ALEMBIC_LIBPATH ${ALEMBIC}/lib)
+		set(ALEMBIC_LIBRARIES Alembic)
 	endif()
 
 	if(WITH_OPENSUBDIV)
@@ -2520,6 +2507,11 @@ if(WITH_CYCLES)
 			)
 		endif()
 	endif()
+
+	if(WITH_CYCLES_OPENSUBDIV AND NOT WITH_OPENSUBDIV)
+		message(STATUS "WITH_CYCLES_OPENSUBDIV requires WITH_OPENSUBDIV to be ON, turning OFF")
+		set(WITH_CYCLES_OPENSUBDIV OFF)
+	endif()
 endif()
 
 if(WITH_INTERNATIONAL)
@@ -2572,8 +2564,6 @@ else()
 	endif()
 	unset(_SYSTEM_BIG_ENDIAN)
 endif()
-
-
 if(WITH_IMAGE_OPENJPEG)
 	if(WITH_SYSTEM_OPENJPEG)
 		# dealt with above
diff --git a/intern/cycles/CMakeLists.txt b/intern/cycles/CMakeLists.txt
index 3b410b2a1e1..97854a88e84 100644
--- a/intern/cycles/CMakeLists.txt
+++ b/intern/cycles/CMakeLists.txt
@@ -146,6 +146,14 @@ if(WITH_CYCLES_OSL)
 	)
 endif()
 
+if(WITH_CYCLES_OPENSUBDIV)
+	add_definitions(-DWITH_OPENSUBDIV)
+	include_directories(
+		SYSTEM
+		${OPENSUBDIV_INCLUDE_DIR}
+	)
+endif()
+
 set(WITH_CYCLES_DEVICE_OPENCL TRUE)
 set(WITH_CYCLES_DEVICE_CUDA TRUE)
 set(WITH_CYCLES_DEVICE_MULTI TRUE)
diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py
index ab798373d9d..17defc5c0fa 100644
--- a/intern/cycles/blender/addon/properties.py
+++ b/intern/cycles/blender/addon/properties.py
@@ -775,6 +775,13 @@ class CyclesMaterialSettings(bpy.types.PropertyGroup):
                 default='LINEAR',
                 )
 
+        cls.displacement_method = EnumProperty(
+                name="Displacement Method",
+                description="Method to use for the displacement",
+                items=enum_displacement_methods,
+                default='BUMP',
+                )
+
     @classmethod
     def unregister(cls):
         del bpy.types.Material.cycles
@@ -952,13 +959,6 @@ class CyclesMeshSettings(bpy.types.PropertyGroup):
                 type=cls,
                 )
 
-        cls.displacement_method = EnumProperty(
-                name="Displacement Method",
-                description="Method to use for the displacement",
-                items=enum_displacement_methods,
-                default='BUMP',
-                )
-
     @classmethod
     def unregister(cls):
         del bpy.types.Mesh.cycles
diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index 65937908a67..3de309a34ea 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -674,40 +674,6 @@ class Cycles_PT_context_material(CyclesButtonsPanel, Panel):
             split.separator()
 
 
-class Cycles_PT_mesh_displacement(CyclesButtonsPanel, Panel):
-    bl_label = "Displacement"
-    bl_context = "data"
-
-    @classmethod
-    def poll(cls, context):
-        if CyclesButtonsPanel.poll(context):
-            if context.mesh or context.curve or context.meta_ball:
-                if context.scene.cycles.feature_set == 'EXPERIMENTAL':
-                    return True
-
-        return False
-
-    def draw(self, context):
-        layout = self.layout
-
-        mesh = context.mesh
-        curve = context.curve
-        mball = context.meta_ball
-
-        if mesh:
-            cdata = mesh.cycles
-        elif curve:
-            cdata = curve.cycles
-        elif mball:
-            cdata = mball.cycles
-
-        split = layout.split()
-
-        col = split.column()
-        sub = col.column(align=True)
-        sub.label(text="Displacement:")
-        sub.prop(cdata, "displacement_method", text="")
-
 class CyclesObject_PT_motion_blur(CyclesButtonsPanel, Panel):
     bl_label = "Motion Blur"
     bl_context = "object"
@@ -1219,6 +1185,11 @@ class CyclesMaterial_PT_settings(CyclesButtonsPanel, Panel):
         col.prop(cmat, "sample_as_light", text="Multiple Importance")
         col.prop(cmat, "use_transparent_shadow")
 
+        if context.scene.cycles.feature_set == 'EXPERIMENTAL':
+            col.separator()
+            col.label(text="Displacement:")
+            col.prop(cmat, "displacement_method", text="")
+
         col = split.column()
         col.label(text="Volume:")
         sub = col.column()
diff --git a/intern/cycles/blender/blender_mesh.cpp b/intern/cycles/blender/blender_mesh.cpp
index 74fd4cb44a0..6dc26c2981b 100644
--- a/intern/cycles/blender/blender_mesh.cpp
+++ b/intern/cycles/blender/blender_mesh.cpp
@@ -409,7 +409,8 @@ static void attr_create_uv_map(Scene *scene,
                                BL::Mesh& b_mesh,
                                const vector<int>& nverts,
                                const vector<int>& face_flags,
-                               bool subdivision)
+                               bool subdivision,
+                               bool subdivide_uvs)
 {
 	if(subdivision) {
 		BL::Mesh::uv_layers_iterator l;
@@ -429,6 +430,10 @@ static void attr_create_uv_map(Scene *scene,
 				else
 					attr = mesh->subd_attributes.add(name, TypeDesc::TypePoint, ATTR_ELEMENT_CORNER);
 
+				if(subdivide_uvs) {
+					attr->flags |= ATTR_SUBDIVIDED;
+				}
+
 				BL::Mesh::polygons_iterator p;
 				float3 *fdata = attr->data_float3();
 
@@ -592,7 +597,8 @@ static void create_mesh(Scene *scene,
                         Mesh *mesh,
                         BL::Mesh& b_mesh,
                         const vector<Shader*>& used_shaders,
-                        bool subdivision=false)
+                        bool subdivision=false,
+                        bool subdivide_uvs=true)
 {
 	/* count vertices and faces */
 	int numverts = b_mesh.vertices.length();
@@ -638,6 +644,7 @@ static void create_mesh(Scene *scene,
 	/* create generated coordinates from undeformed coordinates */
 	if(mesh->need_attribute(scene, ATTR_STD_GENERATED)) {
 		Attribute *attr = attributes.add(ATTR_STD_GENERATED);
+		attr->flags |= ATTR_SUBDIVIDED;
 
 		float3 loc, size;
 		mesh_texture_space(b_mesh, loc, size);
@@ -746,7 +753,7 @@ static void create_mesh(Scene *scene,
 	 * The calculate functions will check whether they're needed or not.
 	 */
 	attr_create_vertex_color(scene, mesh, b_mesh, nverts, face_flags, subdivision);
-	attr_create_uv_map(scene, mesh, b_mesh, nverts, face_flags, subdivision);
+	attr_create_uv_map(scene, mesh, b_mesh, nverts, face_flags, subdivision, subdivide_uvs);
 
 	/* for volume objects, create a matrix to transform from object space to
 	 * mesh texture space. this does not work with deformations but that can
@@ -770,8 +777,35 @@ static void create_subd_mesh(Scene *scene,
                              float dicing_rate,
                              int max_subdivisions)
 {
-	create_mesh(scene, mesh, b_mesh, used_shaders, true);
+	BL::SubsurfModifier subsurf_mod(b_ob.modifiers[b_ob.modifiers.length()-1]);
+	bool subdivide_uvs = subsurf_mod.use_subsurf_uv();
+
+	create_mesh(scene, mesh, b_mesh, used_shaders, true, subdivide_uvs);
+
+	/* export creases */
+	size_t num_creases = 0;
+	BL::Mesh::edges_iterator e;
+
+	for(b_mesh.edges.begin(e); e != b_mesh.edges.end(); ++e) {
+		if(e->crease() != 0.0f) {
+			num_creases++;
+		}
+	}
+
+	mesh->subd_creases.resize(num_creases);
 
+	Mesh::SubdEdgeCrease* crease = mesh->subd_creases.data();
+	for(b_mesh.edges.begin(e); e != b_mesh.edges.end(); ++e) {
+		if(e->crease() != 0.0f) {
+			crease->v[0] = e->vertices()[0];
+			crease->v[1] = e->vertices()[1];
+			crease->crease = e->crease();
+
+			crease++;
+		}
+	}
+
+	/* set subd params */
 	SubdParams sdparams(mesh);
 
 	PointerRNA cobj = RNA_pointer_get(&b_ob.ptr, "cycles");
@@ -903,8 +937,6 @@ Mesh *BlenderSync::sync_mesh(BL::Object& b_ob,
 	mesh_synced.insert(mesh);
 
 	/* create derived mesh */
-	PointerRNA cmesh = RNA_pointer_get(&b_ob_data.ptr, "cycles");
-
 	array<int> oldtriangle = mesh->triangles;
 	
 	/* compares curve_keys rather than strands in order to handle quick hair
@@ -936,7 +968,7 @@ Mesh *BlenderSync::sync_mesh(BL::Object& b_ob,
 			BL::Modifier mod = b_ob.modifiers[b_ob.modifiers.length()-1];
 			bool enabled = preview ? mod.show_viewport() : mod.show_render();
 
-			if(enabled && mod.type() == BL::Modifier::type_SUBSURF && RNA_int_get(&cobj, "use_adaptive_subdivision")) {
+			if(enabled && mod.type() == BL::Modifier::type_SUBSURF && RNA_boolean_get(&cobj, "use_adaptive_subdivision")) {
 				BL::SubsurfModifier subsurf(mod);
 
 				if(subsurf.subdivision_type() == BL::SubsurfModifier::subdivision_type_CATMULL_CLARK) {
@@ -974,21 +1006,6 @@ Mesh *BlenderSync::sync_mesh(BL::Object& b_ob,
 	}
 	mesh->geometry_flags = requested_geometry_flags;
 
-	/* displacement method */
-	if(cmesh.data) {
-		const int method = get_enum(cmesh,
-		                            "displacement_method",
-		                            Mesh::DISPLACE_NUM_METHODS,
-		                            Mesh::DISPLACE_BUMP);
-
-		if(method == 0 || !experimental)
-			mesh->displacement_method = Mesh::DISPLACE_BUMP;
-		else if(method == 1)
-			mesh->displacement_method = Mesh::DISPLACE_TRUE;
-		else
-			mesh->displacement_method = Mesh::DISPLACE_BOTH;
-	}
-
 	/* fluid motion */
 	sync_mesh_fluid_motion(b_ob, scene, mesh);
 
diff --git a/intern/cycles/blender/blender_object.cpp b/intern/cycles/blender/blender_object.cpp
index 9d91d9b0a63..8e798d18647 100644
--- a/intern/cycles/blender/blender_object.cpp
+++ b/intern/cycles/blender/blender_object.cpp
@@ -329,16 +329,18 @@ Object *BlenderSync::sync_object(BL::Object& b_parent,
 			/* object transformation */
 			if(tfm != object->tfm) {
 				VLOG(1) << "Object " << b_ob.name() << " motion detected.";
-				if(motion_time == -1.0f) {
-					object->motion.pre = tfm;
-					object->use_motion = true;
-				}
-				else if(motion_time == 1.0f) {
-					object->motion.post = tfm;
+				if(motion_time == -1.0f || motion_time == 1.0f) {
 					object->use_motion = true;
 				}
 			}
 
+			if(motion_time == -1.0f) {
+				object->motion.pre = tfm;
+			}
+			else if(motion_time == 1.0f) {
+				object->motion.post = tfm;
+			}
+
 			/* mesh deformation */
 			if(object->mesh)
 				sync_mesh_motion(b_ob, object, motion_time);
diff --git a/intern/cycles/blender/blender_shader.cpp b/intern/cycles/blender/blender_shader.cpp
index 2fe8ee90334..534bc6cc897 100644
--- a/intern/cycles/blender/blender_shader.cpp
+++ b/intern/cycles/blender/blender_shader.cpp
@@ -64,6 +64,14 @@ static VolumeInterpolation get_volume_interpolation(PointerRNA& ptr)
 	                                     VOLUME_INTERPOLATION_LINEAR);
 }
 
+static DisplacementMethod get_displacement_method(PointerRNA& ptr)
+{
+	return (DisplacementMethod)get_enum(ptr,
+	                                    "displacement_method",
+	                                    DISPLACE_NUM_METHODS,
+	                                    DISPLACE_BUMP);
+}
+
 static int validate_enum_value(int value, int num_values, int default_value)
 {
 	if(value >= num_values) {
@@ -1182,6 +1190,7 @@ void BlenderSync::sync_materials(bool update_all)
 			shader->heterogeneous_volume = !get_boolean(cmat, "homogeneous_volume");
 			shader->volume_sampling_method = get_volume_sampling(cmat);
 			shader->volume_interpolation_method = get_volume_interpolation(cmat);
+			shader->displacement_method = (experimental) ? get_displacement_method(cmat) : DISPLACE_BUMP;
 
 			shader->set_graph(graph);
 			shader->tag_update(scene);
diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp
index 6a511ea7316..a85f34082db 100644
--- a/intern/cycles/device/device_cuda.cpp
+++ b/intern/cycles/device/device_cuda.cpp
@@ -297,7 +297,7 @@ public:
 			cuda_error_message("CUDA nvcc compiler version could not be parsed.");
 			return false;
 		}
-		if(cuda_version < 60) {
+		if(cuda_version < 75) {
 			printf("Unsupported CUDA version %d.%d detected, "
 			       "you need CUDA 7.5 or newer.\n",
 			       major, minor);
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index 42298a0811d..7bef247d3bd 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -162,6 +162,7 @@ set(SRC_GEOM_HEADERS
 	geom/geom_motion_curve.h
 	geom/geom_motion_triangle.h
 	geom/geom_object.h
+	geom/geom_patch.h
 	geom/geom_primitive.h
 	geom/geom_subd_triangle.h
 	geom/geom_triangle.h
diff --git a/intern/cycles/kernel/geom/geom.h b/intern/cycles/kernel/geom/geom.h
index 493afdc4f62..11548324e18 100644
--- a/intern/cycles/kernel/geom/geom.h
+++ b/intern/cycles/kernel/geom/geom.h
@@ -17,6 +17,7 @@
 
 #include "geom_attribute.h"
 #include "geom_object.h"
+#include "geom_patch.h"
 #include "geom_triangle.h"
 #include "geom_subd_triangle.h"
 #include "geom_triangle_intersect.h"
diff --git a/intern/cycles/kernel/geom/geom_attribute.h b/intern/cycles/kernel/geom/geom_attribute.h
index e036c2752d7..8604d30ad34 100644
--- a/intern/cycles/kernel/geom/geom_attribute.h
+++ b/intern/cycles/kernel/geom/geom_attribute.h
@@ -45,7 +45,7 @@ ccl_device_inline uint attribute_primitive_type(KernelGlobals *kg, const ShaderD
 
 ccl_device_inline AttributeDescriptor attribute_not_found()
 {
-	const AttributeDescriptor desc = {ATTR_ELEMENT_NONE, (NodeAttributeType)0, ATTR_STD_NOT_FOUND};
+	const AttributeDescriptor desc = {ATTR_ELEMENT_NONE, (NodeAttributeType)0, 0, ATTR_STD_NOT_FOUND};
 	return desc;
 }
 
@@ -79,7 +79,8 @@ ccl_device_inline AttributeDescriptor find_attribute(KernelGlobals *kg, const Sh
 
 	/* return result */
 	desc.offset = (attr_map.y == ATTR_ELEMENT_NONE) ? (int)ATTR_STD_NOT_FOUND : (int)attr_map.z;
-	desc.type = (NodeAttributeType)attr_map.w;
+	desc.type = (NodeAttributeType)(attr_map.w & 0xff);
+	desc.flags = (AttributeFlag)(attr_map.w >> 8);
 
 	return desc;
 }
diff --git a/intern/cycles/kernel/geom/geom_object.h b/intern/cycles/kernel/geom/geom_object.h
index c0d15a95954..883c5dc100d 100644
--- a/intern/cycles/kernel/geom/geom_object.h
+++ b/intern/cycles/kernel/geom/geom_object.h
@@ -292,6 +292,18 @@ ccl_device_inline void object_motion_info(KernelGlobals *kg, int object, int *nu
 		*numverts = __float_as_int(f.w);
 }
 
+/* Offset to an objects patch map */
+
+ccl_device_inline uint object_patch_map_offset(KernelGlobals *kg, int object)
+{
+	if(object == OBJECT_NONE)
+		return 0;
+
+	int offset = object*OBJECT_SIZE + 11;
+	float4 f = kernel_tex_fetch(__objects, offset);
+	return __float_as_uint(f.x);
+}
+
 /* Pass ID for shader */
 
 ccl_device int shader_pass_id(KernelGlobals *kg, const ShaderData *sd)
diff --git a/intern/cycles/kernel/geom/geom_patch.h b/intern/cycles/kernel/geom/geom_patch.h
new file mode 100644
index 00000000000..6a0ff5a4a04
--- /dev/null
+++ b/intern/cycles/kernel/geom/geom_patch.h
@@ -0,0 +1,343 @@
+/*
+ * Based on code from OpenSubdiv released under this license:
+ *
+ * Copyright 2013 Pixar
+ *
+ * Licensed under the Apache License, Version 2.0 (the "Apache License")
+ * with the following modification; you may not use this file except in
+ * compliance with the Apache License and the following modification to it:
+ * Section 6. Trademarks. is deleted and replaced with:
+ *
+ * 6. Trademarks. This License does not grant permission to use the trade
+ *   names, trademarks, service marks, or product names of the Licensor
+ *   and its affiliates, except as required to comply with Section 4(c) of
+ *   the License and to reproduce the content of the NOTICE file.
+ *
+ * You may obtain a copy of the Apache License at
+ *
+ *    http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the Apache License with the above modification is
+ * distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the Apache License for the specific
+ * language governing permissions and limitations under the Apache License.
+ *
+ */
+
+CCL_NAMESPACE_BEGIN
+
+typedef struct PatchHandle {
+	int array_index, patch_index, vert_index;
+} PatchHandle;
+
+ccl_device_inline int patch_map_resolve_quadrant(float median, float *u, float *v)
+{
+	int quadrant = -1;
+
+	if(*u < median) {
+		if(*v < median) {
+			quadrant = 0;
+		}
+		else {
+			quadrant = 1;
+			*v -= median;
+		}
+	}
+	else {
+		if(*v < median) {
+			quadrant = 3;
+		}
+		else {
+			quadrant = 2;
+			*v -= median;
+		}
+		*u -= median;
+	}
+
+	return quadrant;
+}
+
+/* retrieve PatchHandle from patch coords */
+
+ccl_device_inline PatchHandle patch_map_find_patch(KernelGlobals *kg, int object, int patch, float u, float v)
+{
+	PatchHandle handle;
+
+	kernel_assert((u >= 0.0f) && (u <= 1.0f) && (v >= 0.0f) && (v <= 1.0f));
+
+	int node = (object_patch_map_offset(kg, object) + patch)/2;
+	float median = 0.5f;
+
+	for(int depth = 0; depth < 0xff; depth++) {
+		float delta = median * 0.5f;
+
+		int quadrant = patch_map_resolve_quadrant(median, &u, &v);
+		kernel_assert(quadrant >= 0);
+
+		uint child = kernel_tex_fetch(__patches, node + quadrant);
+
+		/* is the quadrant a hole? */
+		if(!(child & PATCH_MAP_NODE_IS_SET)) {
+			handle.array_index = -1;
+			return handle;
+		}
+
+		uint index = child & PATCH_MAP_NODE_INDEX_MASK;
+
+		if(child & PATCH_MAP_NODE_IS_LEAF) {
+			handle.array_index = kernel_tex_fetch(__patches, index + 0);
+			handle.patch_index = kernel_tex_fetch(__patches, index + 1);
+			handle.vert_index = kernel_tex_fetch(__patches, index + 2);
+
+			return handle;
+		} else {
+			node = index;
+		}
+
+		median = delta;
+	}
+
+	/* no leaf found */
+	kernel_assert(0);
+
+	handle.array_index = -1;
+	return handle;
+}
+
+ccl_device_inline void patch_eval_bspline_weights(float t, float *point, float *deriv)
+{
+	/* The four uniform cubic B-Spline basis functions evaluated at t */
+	float inv_6 = 1.0f / 6.0f;
+
+	float t2 = t * t;
+	float t3 = t * t2;
+
+	point[0] = inv_6 * (1.0f - 3.0f*(t - t2) - t3);
+	point[1] = inv_6 * (4.0f - 6.0f*t2 + 3.0f*t3);
+	point[2] = inv_6 * (1.0f + 3.0f*(t + t2 - t3));
+	point[3] = inv_6 * t3;
+
+	/* Derivatives of the above four basis functions at t */
+	deriv[0] = -0.5f*t2 + t - 0.5f;
+	deriv[1] =  1.5f*t2 - 2.0f*t;
+	deriv[2] = -1.5f*t2 + t + 0.5f;
+	deriv[3] =  0.5f*t2;
+}
+
+ccl_device_inline void patch_eval_adjust_boundary_weights(uint bits, float *s, float *t)
+{
+	int boundary = ((bits >> 8) & 0xf);
+
+	if(boundary & 1) {
+		t[2] -= t[0];
+		t[1] += 2*t[0];
+		t[0] = 0;
+	}
+
+	if(boundary & 2) {
+		s[1] -= s[3];
+		s[2] += 2*s[3];
+		s[3] = 0;
+	}
+
+	if(boundary & 4) {
+		t[1] -= t[3];
+		t[2] += 2*t[3];
+		t[3] = 0;
+	}
+
+	if(boundary & 8) {
+		s[2] -= s[0];
+		s[1] += 2*s[0];
+		s[0] = 0;
+	}
+}
+
+ccl_device_inline int patch_eval_depth(uint patch_bits)
+{
+	return (patch_bits & 0xf);
+}
+
+ccl_device_inline float patch_eval_param_fraction(uint patch_bits)
+{
+	bool non_quad_root = (patch_bits >> 4) & 0x1;
+	int depth = patch_eval_depth(patch_bits);
+
+	if(non_quad_root) {
+		return 1.0f / (float)(1 << (depth-1));
+	}
+	else {
+		return 1.0f / (float)(1 << depth);
+	}
+}
+
+ccl_device_inline void patch_eval_normalize_coords(uint patch_bits, float *u, float *v)
+{
+	float frac = patch_eval_param_fraction(patch_bits);
+
+	int iu = (patch_bits >> 22) & 0x3ff;
+	int iv = (patch_bits >> 12) & 0x3ff;
+
+	/* top left corner */
+	float pu = (float)iu*frac;
+	float pv = (float)iv*frac;
+
+	/* normalize uv coordinates */
+	*u = (*u - pu) / frac;
+	*v = (*v - pv) / frac;
+}
+
+/* retrieve patch control indices */
+
+ccl_device_inline int patch_eval_indices(KernelGlobals *kg, const PatchHandle *handle, int channel,
+                                         int indices[PATCH_MAX_CONTROL_VERTS])
+{
+	int index_base = kernel_tex_fetch(__patches, handle->array_index + 2) + handle->vert_index;
+
+	/* XXX: regular patches only */
+	for(int i = 0; i < 16; i++) {
+		indices[i] = kernel_tex_fetch(__patches, index_base + i);
+	}
+
+	return 16;
+}
+
+/* evaluate patch basis functions */
+
+ccl_device_inline void patch_eval_basis(KernelGlobals *kg, const PatchHandle *handle, float u, float v,
+                                float weights[PATCH_MAX_CONTROL_VERTS],
+                                float weights_du[PATCH_MAX_CONTROL_VERTS],
+                                float weights_dv[PATCH_MAX_CONTROL_VERTS])
+{
+	uint patch_bits = kernel_tex_fetch(__patches, handle->patch_index + 1); /* read patch param */
+	float d_scale = 1 << patch_eval_depth(patch_bits);
+
+	bool non_quad_root = (patch_bits >> 4) & 0x1;
+	if(non_quad_root) {
+		d_scale *= 0.5f;
+	}
+
+	patch_eval_normalize_coords(patch_bits, &u, &v);
+
+	/* XXX: regular patches only for now. */
+
+	float s[4], t[4], ds[4], dt[4];
+
+	patch_eval_bspline_weights(u, s, ds);
+	patch_eval_bspline_weights(v, t, dt);
+
+	patch_eval_adjust_boundary_weights(patch_bits, s, t);
+	patch_eval_adjust_boundary_weights(patch_bits, ds, dt);
+
+	for(int k = 0; k < 4; k++) {
+		for(int l = 0; l < 4; l++) {
+			weights[4*k+l] = s[l] * t[k];
+			weights_du[4*k+l] = ds[l] * t[k] * d_scale;
+			weights_dv[4*k+l] = s[l] * dt[k] * d_scale;
+		}
+	}
+}
+
+/* generic function for evaluating indices and weights from patch coords */
+
+ccl_device_inline int patch_eval_control_verts(KernelGlobals *kg, int object, int patch, float u, float v, int channel,
+                                        int indices[PATCH_MAX_CONTROL_VERTS],
+                                        float weights[PATCH_MAX_CONTROL_VERTS],
+                                        float weights_du[PATCH_MAX_CONTROL_VERTS],
+                                        float weights_dv[PATCH_MAX_CONTROL_VERTS])
+{
+	PatchHandle handle = patch_map_find_patch(kg, object, patch, u, v);
+	kernel_assert(handle.array_index >= 0);
+
+	int num_control = patch_eval_indices(kg, &handle, channel, indices);
+	patch_eval_basis(kg, &handle, u, v, weights, weights_du, weights_dv);
+
+	return num_control;
+}
+
+/* functions for evaluating attributes on patches */
+
+ccl_device float patch_eval_float(KernelGlobals *kg, const ShaderData *sd, int offset,
+                                  int patch, float u, float v, int channel,
+                                  float *du, float* dv)
+{
+	int indices[PATCH_MAX_CONTROL_VERTS];
+	float weights[PATCH_MAX_CONTROL_VERTS];
+	float weights_du[PATCH_MAX_CONTROL_VERTS];
+	float weights_dv[PATCH_MAX_CONTROL_VERTS];
+
+	int num_control = patch_eval_control_verts(kg, ccl_fetch(sd, object), patch, u, v, channel,
+	                                           indices, weights, weights_du, weights_dv);
+
+	float val = 0.0f;
+	if(du) *du = 0.0f;
+	if(dv) *dv = 0.0f;
+
+	for(int i = 0; i < num_control; i++) {
+		float v = kernel_tex_fetch(__attributes_float, offset + indices[i]);
+
+		val += v * weights[i];
+		if(du) *du += v * weights_du[i];
+		if(dv) *dv += v * weights_dv[i];
+	}
+
+	return val;
+}
+
+ccl_device float3 patch_eval_float3(KernelGlobals *kg, const ShaderData *sd, int offset,
+                                    int patch, float u, float v, int channel,
+                                    float3 *du, float3 *dv)
+{
+	int indices[PATCH_MAX_CONTROL_VERTS];
+	float weights[PATCH_MAX_CONTROL_VERTS];
+	float weights_du[PATCH_MAX_CONTROL_VERTS];
+	float weights_dv[PATCH_MAX_CONTROL_VERTS];
+
+	int num_control = patch_eval_control_verts(kg, ccl_fetch(sd, object), patch, u, v, channel,
+	                                           indices, weights, weights_du, weights_dv);
+
+	float3 val = make_float3(0.0f, 0.0f, 0.0f);
+	if(du) *du = make_float3(0.0f, 0.0f, 0.0f);
+	if(dv) *dv = make_float3(0.0f, 0.0f, 0.0f);
+
+	for(int i = 0; i < num_control; i++) {
+		float3 v = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + indices[i]));
+
+		val += v * weights[i];
+		if(du) *du += v * weights_du[i];
+		if(dv) *dv += v * weights_dv[i];
+	}
+
+	return val;
+}
+
+ccl_device float3 patch_eval_uchar4(KernelGlobals *kg, const ShaderData *sd, int offset,
+                                    int patch, float u, float v, int channel,
+                                    float3 *du, float3 *dv)
+{
+	int indices[PATCH_MAX_CONTROL_VERTS];
+	float weights[PATCH_MAX_CONTROL_VERTS];
+	float weights_du[PATCH_MAX_CONTROL_VERTS];
+	float weights_dv[PATCH_MAX_CONTROL_VERTS];
+
+	int num_control = patch_eval_control_verts(kg, ccl_fetch(sd, object), patch, u, v, channel,
+	                                           indices, weights, weights_du, weights_dv);
+
+	float3 val = make_float3(0.0f, 0.0f, 0.0f);
+	if(du) *du = make_float3(0.0f, 0.0f, 0.0f);
+	if(dv) *dv = make_float3(0.0f, 0.0f, 0.0f);
+
+	for(int i = 0; i < num_control; i++) {
+		float3 v = color_byte_to_float(kernel_tex_fetch(__attributes_uchar4, offset + indices[i]));
+
+		val += v * weights[i];
+		if(du) *du += v * weights_du[i];
+		if(dv) *dv += v * weights_dv[i];
+	}
+
+	return val;
+}
+
+CCL_NAMESPACE_END
+
diff --git a/intern/cycles/kernel/geom/geom_subd_triangle.h b/intern/cycles/kernel/geom/geom_subd_triangle.h
index 79a2ce5cc70..fccacf435f9 100644
--- a/intern/cycles/kernel/geom/geom_subd_triangle.h
+++ b/intern/cycles/kernel/geom/geom_subd_triangle.h
@@ -97,11 +97,54 @@ ccl_device_inline void subd_triangle_patch_corners(KernelGlobals *kg, int patch,
 
 /* Reading attributes on various subdivision triangle elements */
 
-ccl_device float subd_triangle_attribute_float(KernelGlobals *kg, const ShaderData *sd, const AttributeDescriptor desc, float *dx, float *dy)
+ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals *kg, const ShaderData *sd, const AttributeDescriptor desc, float *dx, float *dy)
 {
 	int patch = subd_triangle_patch(kg, sd);
 
-	if(desc.element == ATTR_ELEMENT_FACE) {
+	if(desc.flags & ATTR_SUBDIVIDED) {
+		float2 uv[3];
+		subd_triangle_patch_uv(kg, sd, uv);
+
+		float2 dpdu = uv[0] - uv[2];
+		float2 dpdv = uv[1] - uv[2];
+
+		/* p is [s, t] */
+		float2 p = dpdu * ccl_fetch(sd, u) + dpdv * ccl_fetch(sd, v) + uv[2];
+
+		float a, dads, dadt;
+		a = patch_eval_float(kg, sd, desc.offset, patch, p.x, p.y, 0, &dads, &dadt);
+
+#ifdef __RAY_DIFFERENTIALS__
+		if(dx || dy) {
+			float dsdu = dpdu.x;
+			float dtdu = dpdu.y;
+			float dsdv = dpdv.x;
+			float dtdv = dpdv.y;
+
+			if(dx) {
+				float dudx = ccl_fetch(sd, du).dx;
+				float dvdx = ccl_fetch(sd, dv).dx;
+
+				float dsdx = dsdu*dudx + dsdv*dvdx;
+				float dtdx = dtdu*dudx + dtdv*dvdx;
+
+				*dx = dads*dsdx + dadt*dtdx;
+			}
+			if(dy) {
+				float dudy = ccl_fetch(sd, du).dy;
+				float dvdy = ccl_fetch(sd, dv).dy;
+
+				float dsdy = dsdu*dudy + dsdv*dvdy;
+				float dtdy = dtdu*dudy + dtdv*dvdy;
+
+				*dy = dads*dsdy + dadt*dtdy;
+			}
+		}
+#endif
+
+		return a;
+	}
+	else if(desc.element == ATTR_ELEMENT_FACE) {
 		if(dx) *dx = 0.0f;
 		if(dy) *dy = 0.0f;
 
@@ -110,9 +153,8 @@ ccl_device float subd_triangle_attribute_float(KernelGlobals *kg, const ShaderDa
 	else if(desc.element == ATTR_ELEMENT_VERTEX || desc.element == ATTR_ELEMENT_VERTEX_MOTION) {
 		float2 uv[3];
 		subd_triangle_patch_uv(kg, sd, uv);
-		uint4 v = subd_triangle_patch_indices(kg, patch);
 
-		float a, b, c;
+		uint4 v = subd_triangle_patch_indices(kg, patch);
 
 		float f0 = kernel_tex_fetch(__attributes_float, desc.offset + v.x);
 		float f1 = kernel_tex_fetch(__attributes_float, desc.offset + v.y);
@@ -124,9 +166,9 @@ ccl_device float subd_triangle_attribute_float(KernelGlobals *kg, const ShaderDa
 			f3 = (f3+f0)*0.5f;
 		}
 
-		a = mix(mix(f0, f1, uv[0].x), mix(f3, f2, uv[0].x), uv[0].y);
-		b = mix(mix(f0, f1, uv[1].x), mix(f3, f2, uv[1].x), uv[1].y);
-		c = mix(mix(f0, f1, uv[2].x), mix(f3, f2, uv[2].x), uv[2].y);
+		float a = mix(mix(f0, f1, uv[0].x), mix(f3, f2, uv[0].x), uv[0].y);
+		float b = mix(mix(f0, f1, uv[1].x), mix(f3, f2, uv[1].x), uv[1].y);
+		float c = mix(mix(f0, f1, uv[2].x), mix(f3, f2, uv[2].x), uv[2].y);
 
 #ifdef __RAY_DIFFERENTIALS__
 		if(dx) *dx = ccl_fetch(sd, du).dx*a + ccl_fetch(sd, dv).dx*b - (ccl_fetch(sd, du).dx + ccl_fetch(sd, dv).dx)*c;
@@ -136,13 +178,11 @@ ccl_device float subd_triangle_attribute_float(KernelGlobals *kg, const ShaderDa
 		return ccl_fetch(sd, u)*a + ccl_fetch(sd, v)*b + (1.0f - ccl_fetch(sd, u) - ccl_fetch(sd, v))*c;
 	}
 	else if(desc.element == ATTR_ELEMENT_CORNER) {
-		int corners[4];
-		subd_triangle_patch_corners(kg, patch, corners);
-
 		float2 uv[3];
 		subd_triangle_patch_uv(kg, sd, uv);
 
-		float a, b, c;
+		int corners[4];
+		subd_triangle_patch_corners(kg, patch, corners);
 
 		float f0 = kernel_tex_fetch(__attributes_float, corners[0] + desc.offset);
 		float f1 = kernel_tex_fetch(__attributes_float, corners[1] + desc.offset);
@@ -154,9 +194,9 @@ ccl_device float subd_triangle_attribute_float(KernelGlobals *kg, const ShaderDa
 			f3 = (f3+f0)*0.5f;
 		}
 
-		a = mix(mix(f0, f1, uv[0].x), mix(f3, f2, uv[0].x), uv[0].y);
-		b = mix(mix(f0, f1, uv[1].x), mix(f3, f2, uv[1].x), uv[1].y);
-		c = mix(mix(f0, f1, uv[2].x), mix(f3, f2, uv[2].x), uv[2].y);
+		float a = mix(mix(f0, f1, uv[0].x), mix(f3, f2, uv[0].x), uv[0].y);
+		float b = mix(mix(f0, f1, uv[1].x), mix(f3, f2, uv[1].x), uv[1].y);
+		float c = mix(mix(f0, f1, uv[2].x), mix(f3, f2, uv[2].x), uv[2].y);
 
 #ifdef __RAY_DIFFERENTIALS__
 		if(dx) *dx = ccl_fetch(sd, du).dx*a + ccl_fetch(sd, dv).dx*b - (ccl_fetch(sd, du).dx + ccl_fetch(sd, dv).dx)*c;
@@ -173,11 +213,60 @@ ccl_device float subd_triangle_attribute_float(KernelGlobals *kg, const ShaderDa
 	}
 }
 
-ccl_device float3 subd_triangle_attribute_float3(KernelGlobals *kg, const ShaderData *sd, const AttributeDescriptor desc, float3 *dx, float3 *dy)
+ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals *kg, const ShaderData *sd, const AttributeDescriptor desc, float3 *dx, float3 *dy)
 {
 	int patch = subd_triangle_patch(kg, sd);
 
-	if(desc.element == ATTR_ELEMENT_FACE) {
+	if(desc.flags & ATTR_SUBDIVIDED) {
+		float2 uv[3];
+		subd_triangle_patch_uv(kg, sd, uv);
+
+		float2 dpdu = uv[0] - uv[2];
+		float2 dpdv = uv[1] - uv[2];
+
+		/* p is [s, t] */
+		float2 p = dpdu * ccl_fetch(sd, u) + dpdv * ccl_fetch(sd, v) + uv[2];
+
+		float3 a, dads, dadt;
+
+		if(desc.element == ATTR_ELEMENT_CORNER_BYTE) {
+			a = patch_eval_uchar4(kg, sd, desc.offset, patch, p.x, p.y, 0, &dads, &dadt);
+		}
+		else {
+			a = patch_eval_float3(kg, sd, desc.offset, patch, p.x, p.y, 0, &dads, &dadt);
+		}
+
+#ifdef __RAY_DIFFERENTIALS__
+		if(dx || dy) {
+			float dsdu = dpdu.x;
+			float dtdu = dpdu.y;
+			float dsdv = dpdv.x;
+			float dtdv = dpdv.y;
+
+			if(dx) {
+				float dudx = ccl_fetch(sd, du).dx;
+				float dvdx = ccl_fetch(sd, dv).dx;
+
+				float dsdx = dsdu*dudx + dsdv*dvdx;
+				float dtdx = dtdu*dudx + dtdv*dvdx;
+
+				*dx = dads*dsdx + dadt*dtdx;
+			}
+			if(dy) {
+				float dudy = ccl_fetch(sd, du).dy;
+				float dvdy = ccl_fetch(sd, dv).dy;
+
+				float dsdy = dsdu*dudy + dsdv*dvdy;
+				float dtdy = dtdu*dudy + dtdv*dvdy;
+
+				*dy = dads*dsdy + dadt*dtdy;
+			}
+		}
+#endif
+
+		return a;
+	}
+	else if(desc.element == ATTR_ELEMENT_FACE) {
 		if(dx) *dx = make_float3(0.0f, 0.0f, 0.0f);
 		if(dy) *dy = make_float3(0.0f, 0.0f, 0.0f);
 
@@ -186,9 +275,8 @@ ccl_device float3 subd_triangle_attribute_float3(KernelGlobals *kg, const Shader
 	else if(desc.element == ATTR_ELEMENT_VERTEX || desc.element == ATTR_ELEMENT_VERTEX_MOTION) {
 		float2 uv[3];
 		subd_triangle_patch_uv(kg, sd, uv);
-		uint4 v = subd_triangle_patch_indices(kg, patch);
 
-		float3 a, b, c;
+		uint4 v = subd_triangle_patch_indices(kg, patch);
 
 		float3 f0 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + v.x));
 		float3 f1 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + v.y));
@@ -200,9 +288,9 @@ ccl_device float3 subd_triangle_attribute_float3(KernelGlobals *kg, const Shader
 			f3 = (f3+f0)*0.5f;
 		}
 
-		a = mix(mix(f0, f1, uv[0].x), mix(f3, f2, uv[0].x), uv[0].y);
-		b = mix(mix(f0, f1, uv[1].x), mix(f3, f2, uv[1].x), uv[1].y);
-		c = mix(mix(f0, f1, uv[2].x), mix(f3, f2, uv[2].x), uv[2].y);
+		float3 a = mix(mix(f0, f1, uv[0].x), mix(f3, f2, uv[0].x), uv[0].y);
+		float3 b = mix(mix(f0, f1, uv[1].x), mix(f3, f2, uv[1].x), uv[1].y);
+		float3 c = mix(mix(f0, f1, uv[2].x), mix(f3, f2, uv[2].x), uv[2].y);
 
 #ifdef __RAY_DIFFERENTIALS__
 		if(dx) *dx = ccl_fetch(sd, du).dx*a + ccl_fetch(sd, dv).dx*b - (ccl_fetch(sd, du).dx + ccl_fetch(sd, dv).dx)*c;
@@ -212,13 +300,12 @@ ccl_device float3 subd_triangle_attribute_float3(KernelGlobals *kg, const Shader
 		return ccl_fetch(sd, u)*a + ccl_fetch(sd, v)*b + (1.0f - ccl_fetch(sd, u) - ccl_fetch(sd, v))*c;
 	}
 	else if(desc.element == ATTR_ELEMENT_CORNER || desc.element == ATTR_ELEMENT_CORNER_BYTE) {
-		int corners[4];
-		subd_triangle_patch_corners(kg, patch, corners);
-
 		float2 uv[3];
 		subd_triangle_patch_uv(kg, sd, uv);
 
-		float3 a, b, c;
+		int corners[4];
+		subd_triangle_patch_corners(kg, patch, corners);
+
 		float3 f0, f1, f2, f3;
 
 		if(desc.element == ATTR_ELEMENT_CORNER) {
@@ -239,9 +326,9 @@ ccl_device float3 subd_triangle_attribute_float3(KernelGlobals *kg, const Shader
 			f3 = (f3+f0)*0.5f;
 		}
 
-		a = mix(mix(f0, f1, uv[0].x), mix(f3, f2, uv[0].x), uv[0].y);
-		b = mix(mix(f0, f1, uv[1].x), mix(f3, f2, uv[1].x), uv[1].y);
-		c = mix(mix(f0, f1, uv[2].x), mix(f3, f2, uv[2].x), uv[2].y);
+		float3 a = mix(mix(f0, f1, uv[0].x), mix(f3, f2, uv[0].x), uv[0].y);
+		float3 b = mix(mix(f0, f1, uv[1].x), mix(f3, f2, uv[1].x), uv[1].y);
+		float3 c = mix(mix(f0, f1, uv[2].x), mix(f3, f2, uv[2].x), uv[2].y);
 
 #ifdef __RAY_DIFFERENTIALS__
 		if(dx) *dx = ccl_fetch(sd, du).dx*a + ccl_fetch(sd, dv).dx*b - (ccl_fetch(sd, du).dx + ccl_fetch(sd, dv).dx)*c;
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index a039b414006..063220b542e 100644
--- a/intern/cycles/kernel/kernel_compat_cuda.h
+++ b/intern/cycles/kernel/kernel_compat_cuda.h
@@ -47,6 +47,7 @@
 #define ccl_may_alias
 #define ccl_addr_space
 #define ccl_restrict __restrict__
+#define ccl_align(n) __align__(n)
 
 /* No assert supported for CUDA */
 
diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h
index 8505cb85576..2ae89dde7c4 100644
--- a/intern/cycles/kernel/kernel_compat_opencl.h
+++ b/intern/cycles/kernel/kernel_compat_opencl.h
@@ -40,6 +40,7 @@
 #define ccl_local __local
 #define ccl_private __private
 #define ccl_restrict restrict
+#define ccl_align(n) __attribute__((aligned(n)))
 
 #ifdef __SPLIT_KERNEL__
 #  define ccl_addr_space __global
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index 57b116e1dd7..f3b10c21b9d 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -34,7 +34,7 @@
 CCL_NAMESPACE_BEGIN
 
 /* constants */
-#define OBJECT_SIZE 		11
+#define OBJECT_SIZE 		12
 #define OBJECT_VECTOR_SIZE	6
 #define LIGHT_SIZE			5
 #define FILTER_TABLE_SIZE	1024
@@ -624,9 +624,15 @@ typedef enum AttributeStandard {
 	ATTR_STD_NOT_FOUND = ~0
 } AttributeStandard;
 
+typedef enum AttributeFlag {
+	ATTR_FINAL_SIZE = (1 << 0),
+	ATTR_SUBDIVIDED = (1 << 1),
+} AttributeFlag;
+
 typedef struct AttributeDescriptor {
 	AttributeElement element;
 	NodeAttributeType type;
+	uint flags; /* see enum AttributeFlag */
 	int offset;
 } AttributeDescriptor;
 
@@ -650,23 +656,18 @@ typedef struct AttributeDescriptor {
  * ShaderClosure has a fixed size, and any extra space must be allocated
  * with closure_alloc_extra().
  *
- * float3 is 12 bytes on CUDA and 16 bytes on CPU/OpenCL, we set the data
- * size to ensure ShaderClosure is 80 bytes total everywhere. */
+ * We pad the struct to 80 bytes and ensure it is aligned to 16 bytes, which
+ * we assume to be the maximum required alignment for any struct. */
 
 #define SHADER_CLOSURE_BASE \
 	float3 weight; \
 	ClosureType type; \
 	float sample_weight \
 
-typedef ccl_addr_space struct ShaderClosure {
+typedef ccl_addr_space struct ccl_align(16) ShaderClosure {
 	SHADER_CLOSURE_BASE;
 
-	/* pad to 80 bytes, data types are aligned to own size */
-#ifdef __KERNEL_CUDA__
-	float data[15];
-#else
-	float data[14];
-#endif
+	float data[14]; /* pad to 80 bytes */
 } ShaderClosure;
 
 /* Shader Context
@@ -741,7 +742,7 @@ enum ShaderDataFlag {
 #  define SD_THREAD (get_global_id(1) * get_global_size(0) + get_global_id(0))
 #  if defined(__SPLIT_KERNEL_AOS__)
      /* ShaderData is stored as an Array-of-Structures */
-#    define ccl_soa_member(type, name) type soa_##name;
+#    define ccl_soa_member(type, name) type soa_##name
 #    define ccl_fetch(s, t) (s[SD_THREAD].soa_##t)
 #    define ccl_fetch_array(s, t, index) (&s[SD_THREAD].soa_##t[index])
 #  else
@@ -749,7 +750,7 @@ enum ShaderDataFlag {
 #    define SD_GLOBAL_SIZE (get_global_size(0) * get_global_size(1))
 #    define SD_FIELD_SIZE(t) sizeof(((struct ShaderData*)0)->t)
 #    define SD_OFFSETOF(t) ((char*)(&((struct ShaderData*)0)->t) - (char*)0)
-#    define ccl_soa_member(type, name) type soa_##name;
+#    define ccl_soa_member(type, name) type soa_##name
 #    define ccl_fetch(s, t) (((ShaderData*)((ccl_addr_space char*)s + SD_GLOBAL_SIZE * SD_OFFSETOF(soa_##t) +  SD_FIELD_SIZE(soa_##t) * SD_THREAD - SD_OFFSETOF(soa_##t)))->soa_##t)
 #    define ccl_fetch_array(s, t, index) (&ccl_fetch(s, t)[index])
 #  endif
@@ -1245,6 +1246,16 @@ enum RayState {
 #define REMOVE_RAY_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] = (ray_state[ray_index] & (~flag)))
 #define IS_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] & flag)
 
+/* Patches */
+
+#define PATCH_MAX_CONTROL_VERTS 16
+
+/* Patch map node flags */
+
+#define PATCH_MAP_NODE_IS_SET (1 << 30)
+#define PATCH_MAP_NODE_IS_LEAF (1 << 31)
+#define PATCH_MAP_NODE_INDEX_MASK (~(PATCH_MAP_NODE_IS_SET | PATCH_MAP_NODE_IS_LEAF))
+
 CCL_NAMESPACE_END
 
 #endif /*  __KERNEL_TYPES_H__ */
diff --git a/intern/cycles/kernel/svm/svm_math_util.h b/intern/cycles/kernel/svm/svm_math_util.h
index 3f7d18a02fe..6d13a0d8e02 100644
--- a/intern/cycles/kernel/svm/svm_math_util.h
+++ b/intern/cycles/kernel/svm/svm_math_util.h
@@ -32,21 +32,17 @@ ccl_device void svm_vector_math(float *Fac, float3 *Vector, NodeVectorMath type,
 		*Fac = average_fac(*Vector);
 	}
 	else if(type == NODE_VECTOR_MATH_AVERAGE) {
-		*Fac = len(Vector1 + Vector2);
-		*Vector = normalize(Vector1 + Vector2);
+		*Vector = safe_normalize_len(Vector1 + Vector2, Fac);
 	}
 	else if(type == NODE_VECTOR_MATH_DOT_PRODUCT) {
 		*Fac = dot(Vector1, Vector2);
 		*Vector = make_float3(0.0f, 0.0f, 0.0f);
 	}
 	else if(type == NODE_VECTOR_MATH_CROSS_PRODUCT) {
-		float3 c = cross(Vector1, Vector2);
-		*Fac = len(c);
-		*Vector = normalize(c);
+		*Vector = safe_normalize_len(cross(Vector1, Vector2), Fac);
 	}
 	else if(type == NODE_VECTOR_MATH_NORMALIZE) {
-		*Fac = len(Vector1);
-		*Vector = normalize(Vector1);
+		*Vector = safe_normalize_len(Vector1, Fac);
 	}
 	else {
 		*Fac = 0.0f;
diff --git a/intern/cycles/render/attribute.cpp b/intern/cycles/render/attribute.cpp
index a77ae1121a1..c0d429a583c 100644
--- a/intern/cycles/render/attribute.cpp
+++ b/intern/cycles/render/attribute.cpp
@@ -44,6 +44,7 @@ void Attribute::set(ustring name_, TypeDesc type_, AttributeElement element_)
 	type = type_;
 	element = element_;
 	std = ATTR_STD_NONE;
+	flags = 0;
 
 	/* string and matrix not supported! */
 	assert(type == TypeDesc::TypeFloat || type == TypeDesc::TypeColor ||
@@ -61,6 +62,11 @@ void Attribute::resize(Mesh *mesh, AttributePrimitive prim, bool reserve_only)
 	}
 }
 
+void Attribute::resize(size_t num_elements)
+{
+	buffer.resize(num_elements * data_sizeof(), 0);
+}
+
 void Attribute::add(const float& f)
 {
 	char *data = (char*)&f;
@@ -130,6 +136,10 @@ size_t Attribute::data_sizeof() const
 
 size_t Attribute::element_size(Mesh *mesh, AttributePrimitive prim) const
 {
+	if(flags & ATTR_FINAL_SIZE) {
+		return buffer.size() / data_sizeof();
+	}
+
 	size_t size;
 
 	switch(element) {
diff --git a/intern/cycles/render/attribute.h b/intern/cycles/render/attribute.h
index 67d067d60c9..f4538c76369 100644
--- a/intern/cycles/render/attribute.h
+++ b/intern/cycles/render/attribute.h
@@ -54,11 +54,13 @@ public:
 	TypeDesc type;
 	vector<char> buffer;
 	AttributeElement element;
+	uint flags; /* enum AttributeFlag */
 
 	Attribute() {}
 	~Attribute();
 	void set(ustring name, TypeDesc type, AttributeElement element);
 	void resize(Mesh *mesh, AttributePrimitive prim, bool reserve_only);
+	void resize(size_t num_elements);
 
 	size_t data_sizeof() const;
 	size_t element_size(Mesh *mesh, AttributePrimitive prim) const;
diff --git a/intern/cycles/render/mesh.cpp b/intern/cycles/render/mesh.cpp
index 2edec40b131..f90c19a11c8 100644
--- a/intern/cycles/render/mesh.cpp
+++ b/intern/cycles/render/mesh.cpp
@@ -30,6 +30,8 @@
 
 #include "osl_globals.h"
 
+#include "subd_patch_table.h"
+
 #include "util_foreach.h"
 #include "util_logging.h"
 #include "util_progress.h"
@@ -112,19 +114,12 @@ float3 Mesh::SubdFace::normal(const Mesh *mesh) const
 	return safe_normalize(cross(v1 - v0, v2 - v0));
 }
 
-
 /* Mesh */
 
 NODE_DEFINE(Mesh)
 {
 	NodeType* type = NodeType::add("mesh", create);
 
-	static NodeEnum displacement_method_enum;
-	displacement_method_enum.insert("bump", DISPLACE_BUMP);
-	displacement_method_enum.insert("true", DISPLACE_TRUE);
-	displacement_method_enum.insert("both", DISPLACE_BOTH);
-	SOCKET_ENUM(displacement_method, "Displacement Method", displacement_method_enum, DISPLACE_BUMP);
-
 	SOCKET_UINT(motion_steps, "Motion Steps", 3);
 	SOCKET_BOOLEAN(use_motion_blur, "Use Motion Blur", false);
 
@@ -177,11 +172,14 @@ Mesh::Mesh()
 	num_ngons = 0;
 
 	subdivision_type = SUBDIVISION_NONE;
+
+	patch_table = NULL;
 }
 
 Mesh::~Mesh()
 {
 	delete bvh;
+	delete patch_table;
 }
 
 void Mesh::resize_mesh(int numverts, int numtris)
@@ -274,6 +272,8 @@ void Mesh::clear()
 
 	num_subd_verts = 0;
 
+	subd_creases.clear();
+
 	attributes.clear();
 	curve_attributes.clear();
 	subd_attributes.clear();
@@ -283,6 +283,9 @@ void Mesh::clear()
 	transform_negative_scaled = false;
 	transform_normal = transform_identity();
 	geometry_flags = GEOMETRY_NONE;
+
+	delete patch_table;
+	patch_table = NULL;
 }
 
 int Mesh::split_vertex(int vertex)
@@ -705,7 +708,6 @@ void Mesh::pack_patches(uint *patch_data, uint vert_offset, uint face_offset, ui
 	}
 }
 
-
 void Mesh::compute_bvh(DeviceScene *dscene,
                        SceneParams *params,
                        Progress *progress,
@@ -779,6 +781,17 @@ bool Mesh::has_motion_blur() const
 	         curve_attributes.find(ATTR_STD_MOTION_VERTEX_POSITION)));
 }
 
+bool Mesh::has_true_displacement() const
+{
+	foreach(Shader *shader, used_shaders) {
+		if(shader->has_displacement && shader->displacement_method != DISPLACE_BUMP) {
+			return true;
+		}
+	}
+
+	return false;
+}
+
 bool Mesh::need_build_bvh() const
 {
 	return !transform_applied || has_surface_bssrdf;
@@ -834,6 +847,7 @@ void MeshManager::update_osl_attributes(Device *device, Scene *scene, vector<Att
 			osl_attr.desc.element = ATTR_ELEMENT_OBJECT;
 			osl_attr.value = attr;
 			osl_attr.desc.offset = 0;
+			osl_attr.desc.flags = 0;
 
 			og->attribute_map[i*ATTR_PRIM_TYPES + ATTR_PRIM_TRIANGLE][attr.name()] = osl_attr;
 			og->attribute_map[i*ATTR_PRIM_TYPES + ATTR_PRIM_CURVE][attr.name()] = osl_attr;
@@ -977,6 +991,8 @@ void MeshManager::update_svm_attributes(Device *device, DeviceScene *dscene, Sce
 					attr_map[index].w = NODE_ATTR_MATRIX;
 				else
 					attr_map[index].w = NODE_ATTR_FLOAT3;
+
+				attr_map[index].w |= req.triangle_desc.flags << 8;
 			}
 
 			index++;
@@ -992,6 +1008,8 @@ void MeshManager::update_svm_attributes(Device *device, DeviceScene *dscene, Sce
 					attr_map[index].w = NODE_ATTR_MATRIX;
 				else
 					attr_map[index].w = NODE_ATTR_FLOAT3;
+
+				attr_map[index].w |= req.curve_desc.flags << 8;
 			}
 
 			index++;
@@ -1007,6 +1025,8 @@ void MeshManager::update_svm_attributes(Device *device, DeviceScene *dscene, Sce
 					attr_map[index].w = NODE_ATTR_MATRIX;
 				else
 					attr_map[index].w = NODE_ATTR_FLOAT3;
+
+				attr_map[index].w |= req.subd_desc.flags << 8;
 			}
 
 			index++;
@@ -1071,6 +1091,7 @@ static void update_attribute_element_offset(Mesh *mesh,
 	if(mattr) {
 		/* store element and type */
 		desc.element = mattr->element;
+		desc.flags = mattr->flags;
 		type = mattr->type;
 
 		/* store attribute data in arrays */
@@ -1127,7 +1148,11 @@ static void update_attribute_element_offset(Mesh *mesh,
 
 		/* mesh vertex/curve index is global, not per object, so we sneak
 		 * a correction for that in here */
-		if(element == ATTR_ELEMENT_VERTEX)
+		if(mesh->subdivision_type == Mesh::SUBDIVISION_CATMULL_CLARK && desc.flags & ATTR_SUBDIVIDED) {
+			/* indices for subdivided attributes are retrieved
+			 * from patch table so no need for correction here*/
+		}
+		else if(element == ATTR_ELEMENT_VERTEX)
 			offset -= mesh->vert_offset;
 		else if(element == ATTR_ELEMENT_VERTEX_MOTION)
 			offset -= mesh->vert_offset;
@@ -1323,6 +1348,12 @@ void MeshManager::mesh_calc_offset(Scene *scene)
 		if(mesh->subd_faces.size()) {
 			Mesh::SubdFace& last = mesh->subd_faces[mesh->subd_faces.size()-1];
 			patch_size += (last.ptex_offset + last.num_ptex_faces()) * 8;
+
+			/* patch tables are stored in same array so include them in patch_size */
+			if(mesh->patch_table) {
+				mesh->patch_table_offset = patch_size;
+				patch_size += mesh->patch_table->total_size();
+			}
 		}
 		face_size += mesh->subd_faces.size();
 		corner_size += mesh->subd_face_corners.size();
@@ -1354,6 +1385,12 @@ void MeshManager::device_update_mesh(Device *device,
 		if(mesh->subd_faces.size()) {
 			Mesh::SubdFace& last = mesh->subd_faces[mesh->subd_faces.size()-1];
 			patch_size += (last.ptex_offset + last.num_ptex_faces()) * 8;
+
+			/* patch tables are stored in same array so include them in patch_size */
+			if(mesh->patch_table) {
+				mesh->patch_table_offset = patch_size;
+				patch_size += mesh->patch_table->total_size();
+			}
 		}
 	}
 
@@ -1436,6 +1473,11 @@ void MeshManager::device_update_mesh(Device *device,
 
 		foreach(Mesh *mesh, scene->meshes) {
 			mesh->pack_patches(&patch_data[mesh->patch_offset], mesh->vert_offset, mesh->face_offset, mesh->corner_offset);
+
+			if(mesh->patch_table) {
+				mesh->patch_table->copy_adjusting_offsets(&patch_data[mesh->patch_table_offset], mesh->patch_table_offset);
+			}
+
 			if(progress.get_cancel()) return;
 		}
 
@@ -1622,7 +1664,7 @@ void MeshManager::device_update(Device *device, DeviceScene *dscene, Scene *scen
 	bool old_need_object_flags_update = false;
 	foreach(Mesh *mesh, scene->meshes) {
 		if(mesh->need_update &&
-		   mesh->displacement_method != Mesh::DISPLACE_BUMP)
+		   mesh->has_true_displacement())
 		{
 			true_displacement_used = true;
 			break;
@@ -1648,6 +1690,10 @@ void MeshManager::device_update(Device *device, DeviceScene *dscene, Scene *scen
 	}
 	if(progress.get_cancel()) return;
 
+	/* after mesh data has been copied to device memory we need to update
+	 * offsets for patch tables as this can't be known before hand */
+	scene->object_manager->device_update_patch_map_offsets(device, dscene, scene);
+
 	device_update_attributes(device, dscene, scene, progress);
 	if(progress.get_cancel()) return;
 
diff --git a/intern/cycles/render/mesh.h b/intern/cycles/render/mesh.h
index c9ae9aab888..eff5c50e635 100644
--- a/intern/cycles/render/mesh.h
+++ b/intern/cycles/render/mesh.h
@@ -40,6 +40,7 @@ class Scene;
 class SceneParams;
 class AttributeRequest;
 class DiagSplit;
+struct PackedPatchTable;
 
 /* Mesh */
 
@@ -110,13 +111,9 @@ public:
 		int num_ptex_faces() const { return num_corners == 4 ? 1 : num_corners; }
 	};
 
-	/* Displacement */
-	enum DisplacementMethod {
-		DISPLACE_BUMP = 0,
-		DISPLACE_TRUE = 1,
-		DISPLACE_BOTH = 2,
-
-		DISPLACE_NUM_METHODS,
+	struct SubdEdgeCrease {
+		int v[2];
+		float crease;
 	};
 
 	enum SubdivisionType {
@@ -157,6 +154,8 @@ public:
 	array<int> subd_face_corners;
 	int num_ngons;
 
+	array<SubdEdgeCrease> subd_creases;
+
 	vector<Shader*> used_shaders;
 	AttributeSet attributes;
 	AttributeSet curve_attributes;
@@ -166,7 +165,8 @@ public:
 	bool transform_applied;
 	bool transform_negative_scaled;
 	Transform transform_normal;
-	DisplacementMethod displacement_method;
+
+	PackedPatchTable *patch_table;
 
 	uint motion_steps;
 	bool use_motion_blur;
@@ -184,6 +184,7 @@ public:
 	size_t curvekey_offset;
 
 	size_t patch_offset;
+	size_t patch_table_offset;
 	size_t face_offset;
 	size_t corner_offset;
 
@@ -234,6 +235,7 @@ public:
 	void tag_update(Scene *scene, bool rebuild);
 
 	bool has_motion_blur() const;
+	bool has_true_displacement() const;
 
 	/* Check whether the mesh should have own BVH built separately. Briefly,
 	 * own BVH is needed for mesh, if:
diff --git a/intern/cycles/render/mesh_displace.cpp b/intern/cycles/render/mesh_displace.cpp
index 95f46ff02a2..ef9cfedd412 100644
--- a/intern/cycles/render/mesh_displace.cpp
+++ b/intern/cycles/render/mesh_displace.cpp
@@ -26,19 +26,27 @@
 
 CCL_NAMESPACE_BEGIN
 
+static float3 compute_face_normal(const Mesh::Triangle& t, float3 *verts)
+{
+	float3 v0 = verts[t.v[0]];
+	float3 v1 = verts[t.v[1]];
+	float3 v2 = verts[t.v[2]];
+
+	float3 norm = cross(v1 - v0, v2 - v0);
+	float normlen = len(norm);
+
+	if(normlen == 0.0f)
+		return make_float3(1.0f, 0.0f, 0.0f);
+
+	return norm / normlen;
+}
+
 bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Mesh *mesh, Progress& progress)
 {
 	/* verify if we have a displacement shader */
-	bool has_displacement = false;
-
-	if(mesh->displacement_method != Mesh::DISPLACE_BUMP) {
-		foreach(Shader *shader, mesh->used_shaders)
-			if(shader->has_displacement)
-				has_displacement = true;
-	}
-	
-	if(!has_displacement)
+	if(!mesh->has_true_displacement()) {
 		return false;
+	}
 
 	string msg = string_printf("Computing Displacement %s", mesh->name.c_str());
 	progress.set_status("Updating Mesh", msg);
@@ -67,8 +75,9 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
 		Shader *shader = (shader_index < mesh->used_shaders.size()) ?
 			mesh->used_shaders[shader_index] : scene->default_surface;
 
-		if(!shader->has_displacement)
+		if(!shader->has_displacement || shader->displacement_method == DISPLACE_BUMP) {
 			continue;
+		}
 
 		for(int j = 0; j < 3; j++) {
 			if(done[t.v[j]])
@@ -153,8 +162,9 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
 		Shader *shader = (shader_index < mesh->used_shaders.size()) ?
 			mesh->used_shaders[shader_index] : scene->default_surface;
 
-		if(!shader->has_displacement)
+		if(!shader->has_displacement || shader->displacement_method == DISPLACE_BUMP) {
 			continue;
+		}
 
 		for(int j = 0; j < 3; j++) {
 			if(!done[t.v[j]]) {
@@ -178,9 +188,131 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
 	mesh->attributes.remove(ATTR_STD_FACE_NORMAL);
 	mesh->add_face_normals();
 
-	if(mesh->displacement_method == Mesh::DISPLACE_TRUE) {
-		mesh->attributes.remove(ATTR_STD_VERTEX_NORMAL);
-		mesh->add_vertex_normals();
+	bool need_recompute_vertex_normals = false;
+
+	foreach(Shader *shader, mesh->used_shaders) {
+		if(shader->has_displacement && shader->displacement_method == DISPLACE_TRUE) {
+			need_recompute_vertex_normals = true;
+			break;
+		}
+	}
+
+	if(need_recompute_vertex_normals) {
+		bool flip = mesh->transform_negative_scaled;
+		vector<bool> tri_has_true_disp(num_triangles, false);
+
+		for(size_t i = 0; i < num_triangles; i++) {
+			int shader_index = mesh->shader[i];
+			Shader *shader = (shader_index < mesh->used_shaders.size()) ?
+				mesh->used_shaders[shader_index] : scene->default_surface;
+
+			tri_has_true_disp[i] = shader->has_displacement && shader->displacement_method == DISPLACE_TRUE;
+		}
+
+		/* static vertex normals */
+
+		/* get attributes */
+		Attribute *attr_fN = mesh->attributes.find(ATTR_STD_FACE_NORMAL);
+		Attribute *attr_vN = mesh->attributes.find(ATTR_STD_VERTEX_NORMAL);
+
+		float3 *fN = attr_fN->data_float3();
+		float3 *vN = attr_vN->data_float3();
+
+		/* compute vertex normals */
+
+		/* zero vertex normals on triangles with true displacement */
+		for(size_t i = 0; i < num_triangles; i++) {
+			if(tri_has_true_disp[i]) {
+				for(size_t j = 0; j < 3; j++) {
+					vN[mesh->get_triangle(i).v[j]] = make_float3(0.0f, 0.0f, 0.0f);
+				}
+			}
+		}
+
+		/* add face normals to vertex normals */
+		for(size_t i = 0; i < num_triangles; i++) {
+			if(tri_has_true_disp[i]) {
+				for(size_t j = 0; j < 3; j++) {
+					vN[mesh->get_triangle(i).v[j]] += fN[i];
+				}
+			}
+		}
+
+		/* normalize vertex normals */
+		done.clear();
+		done.resize(num_verts, false);
+
+		for(size_t i = 0; i < num_triangles; i++) {
+			if(tri_has_true_disp[i]) {
+				for(size_t j = 0; j < 3; j++) {
+					int vert = mesh->get_triangle(i).v[j];
+
+					if(done[vert]) {
+						continue;
+					}
+
+					vN[vert] = normalize(vN[vert]);
+					if(flip)
+						vN[vert] = -vN[vert];
+
+					done[vert] = true;
+				}
+			}
+		}
+
+		/* motion vertex normals */
+		Attribute *attr_mP = mesh->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+		Attribute *attr_mN = mesh->attributes.find(ATTR_STD_MOTION_VERTEX_NORMAL);
+
+		if(mesh->has_motion_blur() && attr_mP && attr_mN) {
+			for(int step = 0; step < mesh->motion_steps - 1; step++) {
+				float3 *mP = attr_mP->data_float3() + step*mesh->verts.size();
+				float3 *mN = attr_mN->data_float3() + step*mesh->verts.size();
+
+				/* compute */
+
+				/* zero vertex normals on triangles with true displacement */
+				for(size_t i = 0; i < num_triangles; i++) {
+					if(tri_has_true_disp[i]) {
+						for(size_t j = 0; j < 3; j++) {
+							mN[mesh->get_triangle(i).v[j]] = make_float3(0.0f, 0.0f, 0.0f);
+						}
+					}
+				}
+
+				/* add face normals to vertex normals */
+				for(size_t i = 0; i < num_triangles; i++) {
+					if(tri_has_true_disp[i]) {
+						for(size_t j = 0; j < 3; j++) {
+							float3 fN = compute_face_normal(mesh->get_triangle(i), mP);
+							mN[mesh->get_triangle(i).v[j]] += fN;
+						}
+					}
+				}
+
+				/* normalize vertex normals */
+				done.clear();
+				done.resize(num_verts, false);
+
+				for(size_t i = 0; i < num_triangles; i++) {
+					if(tri_has_true_disp[i]) {
+						for(size_t j = 0; j < 3; j++) {
+							int vert = mesh->get_triangle(i).v[j];
+
+							if(done[vert]) {
+								continue;
+							}
+
+							mN[vert] = normalize(mN[vert]);
+							if(flip)
+								mN[vert] = -mN[vert];
+
+							done[vert] = true;
+						}
+					}
+				}
+			}
+		}
 	}
 
 	return true;
diff --git a/intern/cycles/render/mesh_subdivision.cpp b/intern/cycles/render/mesh_subdivision.cpp
index fe8e41e8d35..efb40efbb79 100644
--- a/intern/cycles/render/mesh_subdivision.cpp
+++ b/intern/cycles/render/mesh_subdivision.cpp
@@ -19,13 +19,302 @@
 
 #include "subd_split.h"
 #include "subd_patch.h"
+#include "subd_patch_table.h"
 
 #include "util_foreach.h"
 
 CCL_NAMESPACE_BEGIN
 
+#ifdef WITH_OPENSUBDIV
+
+CCL_NAMESPACE_END
+
+#include <opensubdiv/far/topologyRefinerFactory.h>
+#include <opensubdiv/far/primvarRefiner.h>
+#include <opensubdiv/far/patchTableFactory.h>
+#include <opensubdiv/far/patchMap.h>
+
+/* specializations of TopologyRefinerFactory for ccl::Mesh */
+
+namespace OpenSubdiv {
+namespace OPENSUBDIV_VERSION {
+namespace Far {
+	template<>
+	bool TopologyRefinerFactory<ccl::Mesh>::resizeComponentTopology(TopologyRefiner& refiner, ccl::Mesh const& mesh)
+	{
+		setNumBaseVertices(refiner, mesh.verts.size());
+		setNumBaseFaces(refiner, mesh.subd_faces.size());
+
+		ccl::Mesh::SubdFace* face = &mesh.subd_faces[0];
+
+		for(int i = 0; i < mesh.subd_faces.size(); i++, face++) {
+			setNumBaseFaceVertices(refiner, i, face->num_corners);
+		}
+
+		return true;
+	}
+
+	template<>
+	bool TopologyRefinerFactory<ccl::Mesh>::assignComponentTopology(TopologyRefiner& refiner, ccl::Mesh const& mesh)
+	{
+		ccl::Mesh::SubdFace* face = &mesh.subd_faces[0];
+
+		for(int i = 0; i < mesh.subd_faces.size(); i++, face++) {
+			IndexArray face_verts = getBaseFaceVertices(refiner, i);
+
+			int* corner = &mesh.subd_face_corners[face->start_corner];
+
+			for(int j = 0; j < face->num_corners; j++, corner++) {
+				face_verts[j] = *corner;
+			}
+		}
+
+		return true;
+	}
+
+	template<>
+	bool TopologyRefinerFactory<ccl::Mesh>::assignComponentTags(TopologyRefiner& refiner, ccl::Mesh const& mesh)
+	{
+		const ccl::Mesh::SubdEdgeCrease* crease = mesh.subd_creases.data();
+
+		for(int i = 0; i < mesh.subd_creases.size(); i++, crease++) {
+			Index edge = findBaseEdge(refiner, crease->v[0], crease->v[1]);
+
+			if(edge != INDEX_INVALID) {
+				setBaseEdgeSharpness(refiner, edge, crease->crease * 10.0f);
+			}
+		}
+
+		for(int i = 0; i < mesh.verts.size(); i++) {
+			ConstIndexArray vert_edges = getBaseVertexEdges(refiner, i);
+
+			if(vert_edges.size() == 2) {
+				float sharpness = refiner.getLevel(0).getEdgeSharpness(vert_edges[0]);
+				sharpness = std::min(sharpness, refiner.getLevel(0).getEdgeSharpness(vert_edges[1]));
+
+				setBaseVertexSharpness(refiner, i, sharpness);
+			}
+		}
+
+		return true;
+	}
+
+	template<>
+	bool TopologyRefinerFactory<ccl::Mesh>::assignFaceVaryingTopology(TopologyRefiner& /*refiner*/, ccl::Mesh const& /*mesh*/)
+	{
+		return true;
+	}
+
+	template<>
+	void TopologyRefinerFactory<ccl::Mesh>::reportInvalidTopology(TopologyError /*err_code*/,
+		char const */*msg*/, ccl::Mesh const& /*mesh*/)
+	{
+	}
+} /* namespace Far */
+} /* namespace OPENSUBDIV_VERSION */
+} /* namespace OpenSubdiv */
+
+CCL_NAMESPACE_BEGIN
+
+using namespace OpenSubdiv;
+
+/* struct that implements OpenSubdiv's vertex interface */
+
+template<typename T>
+struct OsdValue {
+	T value;
+
+	OsdValue() {}
+
+	void Clear(void* = 0) {
+		memset(&value, 0, sizeof(T));
+	}
+
+	void AddWithWeight(OsdValue<T> const& src, float weight) {
+		value += src.value * weight;
+	}
+};
+
+template<>
+void OsdValue<uchar4>::AddWithWeight(OsdValue<uchar4> const& src, float weight)
+{
+	for(int i = 0; i < 4; i++) {
+		value[i] += (uchar)(src.value[i] * weight);
+	}
+}
+
+/* class for holding OpenSubdiv data used during tessellation */
+
+class OsdData {
+	Mesh* mesh;
+	vector<OsdValue<float3> > verts;
+	Far::TopologyRefiner* refiner;
+	Far::PatchTable* patch_table;
+	Far::PatchMap* patch_map;
+
+public:
+	OsdData() : mesh(NULL), refiner(NULL), patch_table(NULL), patch_map(NULL) {}
+
+	~OsdData()
+	{
+		delete refiner;
+		delete patch_table;
+		delete patch_map;
+	}
+
+	void build_from_mesh(Mesh* mesh_)
+	{
+		mesh = mesh_;
+
+		/* type and options */
+		Sdc::SchemeType type = Sdc::SCHEME_CATMARK;
+
+		Sdc::Options options;
+		options.SetVtxBoundaryInterpolation(Sdc::Options::VTX_BOUNDARY_EDGE_ONLY);
+
+		/* create refiner */
+		refiner = Far::TopologyRefinerFactory<Mesh>::Create(*mesh,
+				Far::TopologyRefinerFactory<Mesh>::Options(type, options));
+
+		/* adaptive refinement */
+		int max_isolation = 10;
+		refiner->RefineAdaptive(Far::TopologyRefiner::AdaptiveOptions(max_isolation));
+
+		/* create patch table */
+		Far::PatchTableFactory::Options patch_options;
+		patch_options.endCapType = Far::PatchTableFactory::Options::ENDCAP_GREGORY_BASIS;
+
+		patch_table = Far::PatchTableFactory::Create(*refiner, patch_options);
+
+		/* interpolate verts */
+		int num_refiner_verts = refiner->GetNumVerticesTotal();
+		int num_local_points = patch_table->GetNumLocalPoints();
+
+		verts.resize(num_refiner_verts + num_local_points);
+		for(int i = 0; i < mesh->verts.size(); i++) {
+			verts[i].value = mesh->verts[i];
+		}
+
+		OsdValue<float3>* src = &verts[0];
+		for(int i = 0; i < refiner->GetMaxLevel(); i++) {
+			OsdValue<float3>* dest = src + refiner->GetLevel(i).GetNumVertices();
+			Far::PrimvarRefiner(*refiner).Interpolate(i+1, src, dest);
+			src = dest;
+		}
+
+		patch_table->ComputeLocalPointValues(&verts[0], &verts[num_refiner_verts]);
+
+		/* create patch map */
+		patch_map = new Far::PatchMap(*patch_table);
+	}
+
+	void subdivide_attribute(Attribute& attr)
+	{
+		Far::PrimvarRefiner primvar_refiner(*refiner);
+
+		if(attr.element == ATTR_ELEMENT_VERTEX) {
+			int num_refiner_verts = refiner->GetNumVerticesTotal();
+			int num_local_points = patch_table->GetNumLocalPoints();
+
+			attr.resize(num_refiner_verts + num_local_points);
+			attr.flags |= ATTR_FINAL_SIZE;
+
+			char* src = &attr.buffer[0];
+
+			for(int i = 0; i < refiner->GetMaxLevel(); i++) {
+				char* dest = src + refiner->GetLevel(i).GetNumVertices() * attr.data_sizeof();
+
+				if(attr.same_storage(attr.type, TypeDesc::TypeFloat)) {
+					primvar_refiner.Interpolate(i+1, (OsdValue<float>*)src, (OsdValue<float>*&)dest);
+				}
+				else {
+					primvar_refiner.Interpolate(i+1, (OsdValue<float4>*)src, (OsdValue<float4>*&)dest);
+				}
+
+				src = dest;
+			}
+
+			if(attr.same_storage(attr.type, TypeDesc::TypeFloat)) {
+				patch_table->ComputeLocalPointValues((OsdValue<float>*)&attr.buffer[0],
+					                                 (OsdValue<float>*)&attr.buffer[num_refiner_verts * attr.data_sizeof()]);
+			}
+			else {
+				patch_table->ComputeLocalPointValues((OsdValue<float4>*)&attr.buffer[0],
+					                                 (OsdValue<float4>*)&attr.buffer[num_refiner_verts * attr.data_sizeof()]);
+			}
+		}
+		else if(attr.element == ATTR_ELEMENT_CORNER || attr.element == ATTR_ELEMENT_CORNER_BYTE) {
+			// TODO(mai): fvar interpolation
+		}
+	}
+
+	friend struct OsdPatch;
+	friend class Mesh;
+};
+
+/* ccl::Patch implementation that uses OpenSubdiv for eval */
+
+struct OsdPatch : Patch {
+	OsdData* osd_data;
+
+	OsdPatch(OsdData* data) : osd_data(data) {}
+
+	void eval(float3 *P, float3 *dPdu, float3 *dPdv, float3 *N, float u, float v)
+	{
+		const Far::PatchTable::PatchHandle* handle = osd_data->patch_map->FindPatch(patch_index, u, v);
+		assert(handle);
+
+		float p_weights[20], du_weights[20], dv_weights[20];
+		osd_data->patch_table->EvaluateBasis(*handle, u, v, p_weights, du_weights, dv_weights);
+
+		Far::ConstIndexArray cv = osd_data->patch_table->GetPatchVertices(*handle);
+
+		float3 du, dv;
+		if(P) *P = make_float3(0.0f, 0.0f, 0.0f);
+		du = make_float3(0.0f, 0.0f, 0.0f);
+		dv = make_float3(0.0f, 0.0f, 0.0f);
+
+		for(int i = 0; i < cv.size(); i++) {
+			float3 p = osd_data->verts[cv[i]].value;
+
+			if(P) *P += p * p_weights[i];
+			du += p * du_weights[i];
+			dv += p * dv_weights[i];
+		}
+
+		if(dPdu) *dPdu = du;
+		if(dPdv) *dPdv = dv;
+		if(N) *N = normalize(cross(du, dv));
+	}
+
+	BoundBox bound() { return BoundBox::empty; }
+};
+
+#endif
+
 void Mesh::tessellate(DiagSplit *split)
 {
+#ifdef WITH_OPENSUBDIV
+	OsdData osd_data;
+	bool need_packed_patch_table = false;
+
+	if(subdivision_type == SUBDIVISION_CATMULL_CLARK) {
+		osd_data.build_from_mesh(this);
+	}
+	else
+#endif
+	{
+		/* force linear subdivision if OpenSubdiv is unavailable to avoid
+		 * falling into catmull-clark code paths by accident
+		 */
+		subdivision_type = SUBDIVISION_LINEAR;
+
+		/* force disable attribute subdivision for same reason as above */
+		foreach(Attribute& attr, subd_attributes.attributes) {
+			attr.flags &= ~ATTR_SUBDIVIDED;
+		}
+	}
+
 	int num_faces = subd_faces.size();
 
 	Attribute *attr_vN = subd_attributes.find(ATTR_STD_VERTEX_NORMAL);
@@ -36,113 +325,158 @@ void Mesh::tessellate(DiagSplit *split)
 
 		if(face.is_quad()) {
 			/* quad */
-			LinearQuadPatch patch;
-			float3 *hull = patch.hull;
-			float3 *normals = patch.normals;
+			QuadDice::SubPatch subpatch;
 
-			patch.patch_index = face.ptex_offset;
-			patch.shader = face.shader;
+			LinearQuadPatch quad_patch;
+#ifdef WITH_OPENSUBDIV
+			OsdPatch osd_patch(&osd_data);
 
-			for(int i = 0; i < 4; i++) {
-				hull[i] = verts[subd_face_corners[face.start_corner+i]];
+			if(subdivision_type == SUBDIVISION_CATMULL_CLARK) {
+				osd_patch.patch_index = face.ptex_offset;
+
+				subpatch.patch = &osd_patch;
 			}
+			else
+#endif
+			{
+				float3 *hull = quad_patch.hull;
+				float3 *normals = quad_patch.normals;
+
+				quad_patch.patch_index = face.ptex_offset;
 
-			if(face.smooth) {
 				for(int i = 0; i < 4; i++) {
-					normals[i] = vN[subd_face_corners[face.start_corner+i]];
+					hull[i] = verts[subd_face_corners[face.start_corner+i]];
 				}
-			}
-			else {
-				float3 N = face.normal(this);
-				for(int i = 0; i < 4; i++) {
-					normals[i] = N;
+
+				if(face.smooth) {
+					for(int i = 0; i < 4; i++) {
+						normals[i] = vN[subd_face_corners[face.start_corner+i]];
+					}
+				}
+				else {
+					float3 N = face.normal(this);
+					for(int i = 0; i < 4; i++) {
+						normals[i] = N;
+					}
 				}
+
+				swap(hull[2], hull[3]);
+				swap(normals[2], normals[3]);
+
+				subpatch.patch = &quad_patch;
 			}
 
-			swap(hull[2], hull[3]);
-			swap(normals[2], normals[3]);
+			subpatch.patch->shader = face.shader;
 
 			/* Quad faces need to be split at least once to line up with split ngons, we do this
 			 * here in this manner because if we do it later edge factors may end up slightly off.
 			 */
-			QuadDice::SubPatch subpatch;
-			subpatch.patch = &patch;
-
 			subpatch.P00 = make_float2(0.0f, 0.0f);
 			subpatch.P10 = make_float2(0.5f, 0.0f);
 			subpatch.P01 = make_float2(0.0f, 0.5f);
 			subpatch.P11 = make_float2(0.5f, 0.5f);
-			split->split_quad(&patch, &subpatch);
+			split->split_quad(subpatch.patch, &subpatch);
 
 			subpatch.P00 = make_float2(0.5f, 0.0f);
 			subpatch.P10 = make_float2(1.0f, 0.0f);
 			subpatch.P01 = make_float2(0.5f, 0.5f);
 			subpatch.P11 = make_float2(1.0f, 0.5f);
-			split->split_quad(&patch, &subpatch);
+			split->split_quad(subpatch.patch, &subpatch);
 
 			subpatch.P00 = make_float2(0.0f, 0.5f);
 			subpatch.P10 = make_float2(0.5f, 0.5f);
 			subpatch.P01 = make_float2(0.0f, 1.0f);
 			subpatch.P11 = make_float2(0.5f, 1.0f);
-			split->split_quad(&patch, &subpatch);
+			split->split_quad(subpatch.patch, &subpatch);
 
 			subpatch.P00 = make_float2(0.5f, 0.5f);
 			subpatch.P10 = make_float2(1.0f, 0.5f);
 			subpatch.P01 = make_float2(0.5f, 1.0f);
 			subpatch.P11 = make_float2(1.0f, 1.0f);
-			split->split_quad(&patch, &subpatch);
+			split->split_quad(subpatch.patch, &subpatch);
 		}
 		else {
 			/* ngon */
-			float3 center_vert = make_float3(0.0f, 0.0f, 0.0f);
-			float3 center_normal = make_float3(0.0f, 0.0f, 0.0f);
+#ifdef WITH_OPENSUBDIV
+			if(subdivision_type == SUBDIVISION_CATMULL_CLARK) {
+				OsdPatch patch(&osd_data);
+
+				patch.shader = face.shader;
 
-			float inv_num_corners = 1.0f/float(face.num_corners);
-			for(int corner = 0; corner < face.num_corners; corner++) {
-				center_vert += verts[subd_face_corners[face.start_corner + corner]] * inv_num_corners;
-				center_normal += vN[subd_face_corners[face.start_corner + corner]] * inv_num_corners;
+				for(int corner = 0; corner < face.num_corners; corner++) {
+					patch.patch_index = face.ptex_offset + corner;
+
+					split->split_quad(&patch);
+				}
 			}
+			else
+#endif
+			{
+				float3 center_vert = make_float3(0.0f, 0.0f, 0.0f);
+				float3 center_normal = make_float3(0.0f, 0.0f, 0.0f);
+
+				float inv_num_corners = 1.0f/float(face.num_corners);
+				for(int corner = 0; corner < face.num_corners; corner++) {
+					center_vert += verts[subd_face_corners[face.start_corner + corner]] * inv_num_corners;
+					center_normal += vN[subd_face_corners[face.start_corner + corner]] * inv_num_corners;
+				}
 
-			for(int corner = 0; corner < face.num_corners; corner++) {
-				LinearQuadPatch patch;
-				float3 *hull = patch.hull;
-				float3 *normals = patch.normals;
+				for(int corner = 0; corner < face.num_corners; corner++) {
+					LinearQuadPatch patch;
+					float3 *hull = patch.hull;
+					float3 *normals = patch.normals;
 
-				patch.patch_index = face.ptex_offset + corner;
+					patch.patch_index = face.ptex_offset + corner;
 
-				patch.shader = face.shader;
+					patch.shader = face.shader;
 
-				hull[0] = verts[subd_face_corners[face.start_corner + mod(corner + 0, face.num_corners)]];
-				hull[1] = verts[subd_face_corners[face.start_corner + mod(corner + 1, face.num_corners)]];
-				hull[2] = verts[subd_face_corners[face.start_corner + mod(corner - 1, face.num_corners)]];
-				hull[3] = center_vert;
+					hull[0] = verts[subd_face_corners[face.start_corner + mod(corner + 0, face.num_corners)]];
+					hull[1] = verts[subd_face_corners[face.start_corner + mod(corner + 1, face.num_corners)]];
+					hull[2] = verts[subd_face_corners[face.start_corner + mod(corner - 1, face.num_corners)]];
+					hull[3] = center_vert;
 
-				hull[1] = (hull[1] + hull[0]) * 0.5;
-				hull[2] = (hull[2] + hull[0]) * 0.5;
+					hull[1] = (hull[1] + hull[0]) * 0.5;
+					hull[2] = (hull[2] + hull[0]) * 0.5;
 
-				if(face.smooth) {
-					normals[0] = vN[subd_face_corners[face.start_corner + mod(corner + 0, face.num_corners)]];
-					normals[1] = vN[subd_face_corners[face.start_corner + mod(corner + 1, face.num_corners)]];
-					normals[2] = vN[subd_face_corners[face.start_corner + mod(corner - 1, face.num_corners)]];
-					normals[3] = center_normal;
+					if(face.smooth) {
+						normals[0] = vN[subd_face_corners[face.start_corner + mod(corner + 0, face.num_corners)]];
+						normals[1] = vN[subd_face_corners[face.start_corner + mod(corner + 1, face.num_corners)]];
+						normals[2] = vN[subd_face_corners[face.start_corner + mod(corner - 1, face.num_corners)]];
+						normals[3] = center_normal;
 
-					normals[1] = (normals[1] + normals[0]) * 0.5;
-					normals[2] = (normals[2] + normals[0]) * 0.5;
-				}
-				else {
-					float3 N = face.normal(this);
-					for(int i = 0; i < 4; i++) {
-						normals[i] = N;
+						normals[1] = (normals[1] + normals[0]) * 0.5;
+						normals[2] = (normals[2] + normals[0]) * 0.5;
+					}
+					else {
+						float3 N = face.normal(this);
+						for(int i = 0; i < 4; i++) {
+							normals[i] = N;
+						}
 					}
-				}
 
-				split->split_quad(&patch);
+					split->split_quad(&patch);
+				}
 			}
 		}
 	}
 
 	/* interpolate center points for attributes */
 	foreach(Attribute& attr, subd_attributes.attributes) {
+#ifdef WITH_OPENSUBDIV
+		if(subdivision_type == SUBDIVISION_CATMULL_CLARK && attr.flags & ATTR_SUBDIVIDED) {
+			if(attr.element == ATTR_ELEMENT_CORNER || attr.element == ATTR_ELEMENT_CORNER_BYTE) {
+				/* keep subdivision for corner attributes disabled for now */
+				attr.flags &= ~ATTR_SUBDIVIDED;
+			}
+			else {
+				osd_data.subdivide_attribute(attr);
+
+				need_packed_patch_table = true;
+				continue;
+			}
+		}
+#endif
+
 		char* data = attr.data();
 		size_t stride = attr.data_sizeof();
 		int ngons = 0;
@@ -218,6 +552,15 @@ void Mesh::tessellate(DiagSplit *split)
 			default: break;
 		}
 	}
+
+#ifdef WITH_OPENSUBDIV
+	/* pack patch tables */
+	if(need_packed_patch_table) {
+		delete patch_table;
+		patch_table = new PackedPatchTable;
+		patch_table->pack(osd_data.patch_table);
+	}
+#endif
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/render/object.cpp b/intern/cycles/render/object.cpp
index 5dbe2e4d6a1..62076f3a865 100644
--- a/intern/cycles/render/object.cpp
+++ b/intern/cycles/render/object.cpp
@@ -29,6 +29,8 @@
 #include "util_progress.h"
 #include "util_vector.h"
 
+#include "subd_patch_table.h"
+
 CCL_NAMESPACE_BEGIN
 
 /* Object */
@@ -607,6 +609,40 @@ void ObjectManager::device_update_flags(Device *device,
 	device->tex_alloc("__object_flag", dscene->object_flag);
 }
 
+void ObjectManager::device_update_patch_map_offsets(Device *device, DeviceScene *dscene, Scene *scene)
+{
+	if (scene->objects.size() == 0)
+		return;
+
+	uint4* objects = (uint4*)dscene->objects.get_data();
+
+	bool update = false;
+
+	int object_index = 0;
+	foreach(Object *object, scene->objects) {
+		int offset = object_index*OBJECT_SIZE + 11;
+
+		Mesh* mesh = object->mesh;
+
+		if(mesh->patch_table) {
+			uint patch_map_offset = 2*(mesh->patch_table_offset + mesh->patch_table->total_size() -
+			                           mesh->patch_table->num_nodes * PATCH_NODE_SIZE) - mesh->patch_offset;
+
+			if(objects[offset].x != patch_map_offset) {
+				objects[offset].x = patch_map_offset;
+				update = true;
+			}
+		}
+
+		object_index++;
+	}
+
+	if(update) {
+		device->tex_free(dscene->objects);
+		device->tex_alloc("__objects", dscene->objects);
+	}
+}
+
 void ObjectManager::device_free(Device *device, DeviceScene *dscene)
 {
 	device->tex_free(dscene->objects);
@@ -656,7 +692,7 @@ void ObjectManager::apply_static_transforms(DeviceScene *dscene, Scene *scene, u
 		 * Could be solved by moving reference counter to Mesh.
 		 */
 		if((mesh_users[object->mesh] == 1 && !object->mesh->has_surface_bssrdf) &&
-		   object->mesh->displacement_method == Mesh::DISPLACE_BUMP)
+		   !object->mesh->has_true_displacement())
 		{
 			if(!(motion_blur && object->use_motion)) {
 				if(!object->mesh->transform_applied) {
diff --git a/intern/cycles/render/object.h b/intern/cycles/render/object.h
index 7ab73f3c91a..2e5837f672f 100644
--- a/intern/cycles/render/object.h
+++ b/intern/cycles/render/object.h
@@ -97,6 +97,8 @@ public:
 	                         Scene *scene,
 	                         Progress& progress,
 	                         bool bounds_valid = true);
+	void device_update_patch_map_offsets(Device *device, DeviceScene *dscene, Scene *scene);
+
 	void device_free(Device *device, DeviceScene *dscene);
 
 	void tag_update(Scene *scene);
diff --git a/intern/cycles/render/shader.cpp b/intern/cycles/render/shader.cpp
index 4cdb878df45..d000cca5a45 100644
--- a/intern/cycles/render/shader.cpp
+++ b/intern/cycles/render/shader.cpp
@@ -150,6 +150,12 @@ NODE_DEFINE(Shader)
 	volume_interpolation_method_enum.insert("cubic", VOLUME_INTERPOLATION_CUBIC);
 	SOCKET_ENUM(volume_interpolation_method, "Volume Interpolation Method", volume_interpolation_method_enum, VOLUME_INTERPOLATION_LINEAR);
 
+	static NodeEnum displacement_method_enum;
+	displacement_method_enum.insert("bump", DISPLACE_BUMP);
+	displacement_method_enum.insert("true", DISPLACE_TRUE);
+	displacement_method_enum.insert("both", DISPLACE_BOTH);
+	SOCKET_ENUM(displacement_method, "Displacement Method", displacement_method_enum, DISPLACE_BUMP);
+
 	return type;
 }
 
@@ -173,6 +179,8 @@ Shader::Shader()
 	has_object_dependency = false;
 	has_integrator_dependency = false;
 
+	displacement_method = DISPLACE_BUMP;
+
 	id = -1;
 	used = false;
 
@@ -310,7 +318,7 @@ int ShaderManager::get_shader_id(Shader *shader, Mesh *mesh, bool smooth)
 	int id = shader->id*2;
 	
 	/* index depends bump since this setting is not in the shader */
-	if(mesh && mesh->displacement_method != Mesh::DISPLACE_TRUE)
+	if(mesh && shader->displacement_method != DISPLACE_TRUE)
 		id += 1;
 	/* smooth flag */
 	if(smooth)
diff --git a/intern/cycles/render/shader.h b/intern/cycles/render/shader.h
index dc57ed4e4eb..060ad7056bc 100644
--- a/intern/cycles/render/shader.h
+++ b/intern/cycles/render/shader.h
@@ -66,6 +66,14 @@ enum VolumeInterpolation {
 	VOLUME_NUM_INTERPOLATION,
 };
 
+enum DisplacementMethod {
+	DISPLACE_BUMP = 0,
+	DISPLACE_TRUE = 1,
+	DISPLACE_BOTH = 2,
+
+	DISPLACE_NUM_METHODS,
+};
+
 /* Shader describing the appearance of a Mesh, Light or Background.
  *
  * While there is only a single shader graph, it has three outputs: surface,
@@ -110,6 +118,9 @@ public:
 	bool has_object_dependency;
 	bool has_integrator_dependency;
 
+	/* displacement */
+	DisplacementMethod displacement_method;
+
 	/* requested mesh attributes */
 	AttributeRequestSet attributes;
 
diff --git a/intern/cycles/subd/CMakeLists.txt b/intern/cycles/subd/CMakeLists.txt
index db497013693..9265299e82b 100644
--- a/intern/cycles/subd/CMakeLists.txt
+++ b/intern/cycles/subd/CMakeLists.txt
@@ -16,6 +16,7 @@ set(SRC
 	subd_dice.cpp
 	subd_patch.cpp
 	subd_split.cpp
+	subd_patch_table.cpp
 )
 
 set(SRC_HEADERS
@@ -24,10 +25,6 @@ set(SRC_HEADERS
 	subd_split.h
 )
 
-if(WITH_CYCLES_OPENSUBDIV)
-	add_definitions(-DWITH_OPENSUBDIV)
-endif()
-
 include_directories(${INC})
 include_directories(SYSTEM ${INC_SYS})
 
diff --git a/intern/cycles/subd/subd_patch_table.cpp b/intern/cycles/subd/subd_patch_table.cpp
new file mode 100644
index 00000000000..68ec1b2c6a6
--- /dev/null
+++ b/intern/cycles/subd/subd_patch_table.cpp
@@ -0,0 +1,297 @@
+/*
+ * Based on code from OpenSubdiv released under this license:
+ *
+ * Copyright 2014 DreamWorks Animation LLC.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "Apache License")
+ * with the following modification; you may not use this file except in
+ * compliance with the Apache License and the following modification to it:
+ * Section 6. Trademarks. is deleted and replaced with:
+ *
+ * 6. Trademarks. This License does not grant permission to use the trade
+ *   names, trademarks, service marks, or product names of the Licensor
+ *   and its affiliates, except as required to comply with Section 4(c) of
+ *   the License and to reproduce the content of the NOTICE file.
+ *
+ * You may obtain a copy of the Apache License at
+ *
+ *    http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the Apache License with the above modification is
+ * distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the Apache License for the specific
+ * language governing permissions and limitations under the Apache License.
+ *
+ */
+
+#include "subd_patch_table.h"
+#include "kernel_types.h"
+
+#include "util_math.h"
+
+#ifdef WITH_OPENSUBDIV
+#include <opensubdiv/far/patchTable.h>
+#endif
+
+CCL_NAMESPACE_BEGIN
+
+#ifdef WITH_OPENSUBDIV
+
+using namespace OpenSubdiv;
+
+/* functions for building patch maps */
+
+struct PatchMapQuadNode {
+	/* sets all the children to point to the patch of index */
+	void set_child(int index)
+	{
+		for (int i = 0; i < 4; i++) {
+			children[i] = index | PATCH_MAP_NODE_IS_SET | PATCH_MAP_NODE_IS_LEAF;
+		}
+	}
+
+	/* sets the child in quadrant to point to the node or patch of the given index */
+	void set_child(unsigned char quadrant, int index, bool is_leaf=true)
+	{
+		assert(quadrant < 4);
+		children[quadrant] = index | PATCH_MAP_NODE_IS_SET | (is_leaf ? PATCH_MAP_NODE_IS_LEAF : 0);
+	}
+
+	uint children[4];
+};
+
+template<class T>
+static int resolve_quadrant(T& median, T& u, T& v)
+{
+	int quadrant = -1;
+
+	if(u < median) {
+		if(v < median) {
+			quadrant = 0;
+		}
+		else {
+			quadrant = 1;
+			v -= median;
+		}
+	}
+	else {
+		if(v < median) {
+			quadrant = 3;
+		}
+		else {
+			quadrant = 2;
+			v -= median;
+		}
+		u -= median;
+	}
+
+	return quadrant;
+}
+
+static void build_patch_map(PackedPatchTable& table, OpenSubdiv::Far::PatchTable* patch_table, int offset)
+{
+	int num_faces = 0;
+
+	for(int array = 0; array < table.num_arrays; array++) {
+		Far::ConstPatchParamArray params = patch_table->GetPatchParams(array);
+
+		for(int j = 0; j < patch_table->GetNumPatches(array); j++) {
+			num_faces = max(num_faces, (int)params[j].GetFaceId());
+		}
+	}
+	num_faces++;
+
+	vector<PatchMapQuadNode> quadtree;
+	quadtree.reserve(num_faces + table.num_patches);
+	quadtree.resize(num_faces);
+
+	/* adjust offsets to make indices relative to the table */
+	int handle_index = -(table.num_patches * PATCH_HANDLE_SIZE);
+	offset += table.total_size();
+
+	/* populate the quadtree from the FarPatchArrays sub-patches */
+	for(int array = 0; array < table.num_arrays; array++) {
+		Far::ConstPatchParamArray params = patch_table->GetPatchParams(array);
+
+		for(int i = 0; i < patch_table->GetNumPatches(array); i++, handle_index += PATCH_HANDLE_SIZE) {
+			const Far::PatchParam& param = params[i];
+			unsigned short depth = param.GetDepth();
+
+			PatchMapQuadNode* node = &quadtree[params[i].GetFaceId()];
+
+			if(depth == (param.NonQuadRoot() ? 1 : 0)) {
+				/* special case : regular BSpline face w/ no sub-patches */
+				node->set_child(handle_index + offset);
+				continue;
+			}
+
+			int u = param.GetU();
+			int v = param.GetV();
+			int pdepth = param.NonQuadRoot() ? depth-2 : depth-1;
+			int half = 1 << pdepth;
+
+			for(int j = 0; j < depth; j++) {
+				int delta = half >> 1;
+
+				int quadrant = resolve_quadrant(half, u, v);
+				assert(quadrant >= 0);
+
+				half = delta;
+
+				if(j == pdepth) {
+					/* we have reached the depth of the sub-patch : add a leaf */
+					assert(!(node->children[quadrant] & PATCH_MAP_NODE_IS_SET));
+					node->set_child(quadrant, handle_index + offset, true);
+					break;
+				}
+				else {
+					/* travel down the child node of the corresponding quadrant */
+					if(!(node->children[quadrant] & PATCH_MAP_NODE_IS_SET)) {
+						/* create a new branch in the quadrant */
+						quadtree.push_back(PatchMapQuadNode());
+
+						int idx = (int)quadtree.size() - 1;
+						node->set_child(quadrant, idx*4 + offset, false);
+
+						node = &quadtree[idx];
+					}
+					else {
+						/* travel down an existing branch */
+						uint idx = node->children[quadrant] & PATCH_MAP_NODE_INDEX_MASK;
+						node = &(quadtree[(idx - offset)/4]);
+					}
+				}
+			}
+		}
+	}
+
+	/* copy into table */
+	assert(table.table.size() == table.total_size());
+	uint map_offset = table.total_size();
+
+	table.num_nodes = quadtree.size() * 4;
+	table.table.resize(table.total_size());
+
+	uint* data = &table.table[map_offset];
+
+	for(int i = 0; i < quadtree.size(); i++) {
+		for(int j = 0; j < 4; j++) {
+			assert(quadtree[i].children[j] & PATCH_MAP_NODE_IS_SET);
+			*(data++) = quadtree[i].children[j];
+		}
+	}
+}
+
+#endif
+
+/* packed patch table functions */
+
+size_t PackedPatchTable::total_size()
+{
+	return num_arrays * PATCH_ARRAY_SIZE +
+		   num_indices +
+		   num_patches * (PATCH_PARAM_SIZE + PATCH_HANDLE_SIZE) +
+		   num_nodes * PATCH_NODE_SIZE;
+}
+
+void PackedPatchTable::pack(Far::PatchTable* patch_table, int offset)
+{
+	num_arrays = 0;
+	num_patches = 0;
+	num_indices = 0;
+	num_nodes = 0;
+
+#ifdef WITH_OPENSUBDIV
+	num_arrays = patch_table->GetNumPatchArrays();
+
+	for(int i = 0; i < num_arrays; i++) {
+		int patches = patch_table->GetNumPatches(i);
+		int num_control = patch_table->GetPatchArrayDescriptor(i).GetNumControlVertices();
+
+		num_patches += patches;
+		num_indices += patches * num_control;
+	}
+
+	table.resize(total_size());
+	uint* data = &table[0];
+
+	uint* array = data;
+	uint* index = array + num_arrays * PATCH_ARRAY_SIZE;
+	uint* param = index + num_indices;
+	uint* handle = param + num_patches * PATCH_PARAM_SIZE;
+
+	uint current_param = 0;
+
+	for(int i = 0; i < num_arrays; i++) {
+		*(array++) = patch_table->GetPatchArrayDescriptor(i).GetType();
+		*(array++) = patch_table->GetNumPatches(i);
+		*(array++) = (index - data) + offset;
+		*(array++) = (param - data) + offset;
+
+		Far::ConstIndexArray indices = patch_table->GetPatchArrayVertices(i);
+
+		for(int j = 0; j < indices.size(); j++) {
+			*(index++) = indices[j];
+		}
+
+		const Far::PatchParamTable& param_table = patch_table->GetPatchParamTable();
+
+		int num_control = patch_table->GetPatchArrayDescriptor(i).GetNumControlVertices();
+		int patches = patch_table->GetNumPatches(i);
+
+		for(int j = 0; j < patches; j++, current_param++) {
+			*(param++) = param_table[current_param].field0;
+			*(param++) = param_table[current_param].field1;
+
+			*(handle++) = (array - data) - PATCH_ARRAY_SIZE + offset;
+			*(handle++) = (param - data) - PATCH_PARAM_SIZE + offset;
+			*(handle++) = j * num_control;
+		}
+	}
+
+	build_patch_map(*this, patch_table, offset);
+#else
+	(void)patch_table;
+	(void)offset;
+#endif
+}
+
+void PackedPatchTable::copy_adjusting_offsets(uint* dest, int doffset)
+{
+	uint* src = &table[0];
+
+	/* arrays */
+	for(int i = 0; i < num_arrays; i++) {
+		*(dest++) = *(src++);
+		*(dest++) = *(src++);
+		*(dest++) = *(src++) + doffset;
+		*(dest++) = *(src++) + doffset;
+	}
+
+	/* indices */
+	for(int i = 0; i < num_indices; i++) {
+		*(dest++) = *(src++);
+	}
+
+	/* params */
+	for(int i = 0; i < num_patches; i++) {
+		*(dest++) = *(src++);
+		*(dest++) = *(src++);
+	}
+
+	/* handles */
+	for(int i = 0; i < num_patches; i++) {
+		*(dest++) = *(src++) + doffset;
+		*(dest++) = *(src++) + doffset;
+		*(dest++) = *(src++);
+	}
+
+	/* nodes */
+	for(int i = 0; i < num_nodes; i++) {
+		*(dest++) = *(src++) + doffset;
+	}
+}
+
+CCL_NAMESPACE_END
+
diff --git a/intern/cycles/subd/subd_patch_table.h b/intern/cycles/subd/subd_patch_table.h
new file mode 100644
index 00000000000..c8c7ecf9e47
--- /dev/null
+++ b/intern/cycles/subd/subd_patch_table.h
@@ -0,0 +1,63 @@
+/*
+ * 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 __SUBD_PATCH_TABLE_H__
+#define __SUBD_PATCH_TABLE_H__
+
+#include "util_types.h"
+#include "util_vector.h"
+
+#ifdef WITH_OPENSUBDIV
+#ifdef _MSC_VER
+#  include "iso646.h"
+#endif
+
+#include <opensubdiv/far/patchTable.h>
+#endif
+
+CCL_NAMESPACE_BEGIN
+
+#ifdef WITH_OPENSUBDIV
+using namespace OpenSubdiv;
+#else
+/* forward declare for when OpenSubdiv is unavailable */
+namespace Far { struct PatchTable; }
+#endif
+
+#define PATCH_ARRAY_SIZE 4
+#define PATCH_PARAM_SIZE 2
+#define PATCH_HANDLE_SIZE 3
+#define PATCH_NODE_SIZE 1
+
+struct PackedPatchTable {
+	vector<uint> table;
+
+	size_t num_arrays;
+	size_t num_indices;
+	size_t num_patches;
+	size_t num_nodes;
+
+	/* calculated size from num_* members */
+	size_t total_size();
+
+	void pack(Far::PatchTable* patch_table, int offset = 0);
+	void copy_adjusting_offsets(uint* dest, int doffset);
+};
+
+CCL_NAMESPACE_END
+
+#endif /* __SUBD_PATCH_TABLE_H__ */
+
diff --git a/intern/cycles/test/CMakeLists.txt b/intern/cycles/test/CMakeLists.txt
index 80fe893826a..9af777fb9dd 100644
--- a/intern/cycles/test/CMakeLists.txt
+++ b/intern/cycles/test/CMakeLists.txt
@@ -26,6 +26,7 @@ set(ALL_CYCLES_LIBRARIES
 	cycles_device
 	cycles_bvh
 	cycles_graph
+	cycles_subd
 	cycles_util
 	${OPENIMAGEIO_LIBRARIES}
 )
@@ -41,6 +42,16 @@ if(WITH_IMAGE_OPENJPEG AND NOT WITH_SYSTEM_OPENJPEG)
 		extern_openjpeg
 	)
 endif()
+if(WITH_CYCLES_OPENSUBDIV)
+	add_definitions(-DWITH_OPENSUBDIV)
+	include_directories(
+		SYSTEM
+		${OPENSUBDIV_INCLUDE_DIR}
+	)
+	list(APPEND ALL_CYCLES_LIBRARIES
+		${OPENSUBDIV_LIBRARIES}
+	)
+endif()
 list(APPEND ALL_CYCLES_LIBRARIES
 	${BOOST_LIBRARIES}
 )
diff --git a/intern/cycles/util/util_math.h b/intern/cycles/util/util_math.h
index 13aba0646d2..89a882d9b9d 100644
--- a/intern/cycles/util/util_math.h
+++ b/intern/cycles/util/util_math.h
@@ -572,6 +572,12 @@ ccl_device_inline float3 safe_normalize(const float3 a)
 	return (t != 0.0f)? a/t: a;
 }
 
+ccl_device_inline float3 safe_normalize_len(const float3 a, float *t)
+{
+	*t = len(a);
+	return (*t != 0.0f)? a/(*t): a;
+}
+
 #ifndef __KERNEL_OPENCL__
 
 ccl_device_inline bool operator==(const float3 a, const float3 b)
diff --git a/intern/cycles/util/util_vector.h b/intern/cycles/util/util_vector.h
index 6f8c3f6f3de..546b17570bb 100644
--- a/intern/cycles/util/util_vector.h
+++ b/intern/cycles/util/util_vector.h
@@ -222,6 +222,11 @@ public:
 		return datasize_;
 	}
 
+	T* data()
+	{
+		return data_;
+	}
+
 	const T* data() const
 	{
 		return data_;
diff --git a/release/scripts/startup/bl_ui/properties_grease_pencil_common.py b/release/scripts/startup/bl_ui/properties_grease_pencil_common.py
index c80c5ca7ddd..90cf410036a 100644
--- a/release/scripts/startup/bl_ui/properties_grease_pencil_common.py
+++ b/release/scripts/startup/bl_ui/properties_grease_pencil_common.py
@@ -80,6 +80,7 @@ class GreasePencilDrawingToolsPanel:
         sub = col.column(align=True)
         sub.prop(context.tool_settings, "use_gpencil_additive_drawing", text="Additive Drawing")
         sub.prop(context.tool_settings, "use_gpencil_continuous_drawing", text="Continuous Drawing")
+        sub.prop(context.tool_settings, "use_gpencil_draw_onback", text="Draw on Back")
 
         col.separator()
         col.separator()
@@ -658,7 +659,7 @@ class GPENCIL_UL_palettecolor(UIList):
             row = split.row(align=True)
             row.prop(palcolor, "color", text="", emboss=palcolor.is_stroke_visible)
             row.prop(palcolor, "fill_color", text="", emboss=palcolor.is_fill_visible)
-            split.prop(palcolor, "info", text="", emboss=False)
+            split.prop(palcolor, "name", text="", emboss=False)
 
             row = layout.row(align=True)
             row.prop(palcolor, "lock", text="", emboss=False)
diff --git a/source/blender/alembic/CMakeLists.txt b/source/blender/alembic/CMakeLists.txt
index f579a3b200a..42bd6a9c340 100644
--- a/source/blender/alembic/CMakeLists.txt
+++ b/source/blender/alembic/CMakeLists.txt
@@ -40,7 +40,7 @@ set(INC_SYS
 	${HDF5_INCLUDE_DIRS}
 	${OPENEXR_INCLUDE_DIRS}
 )
-if(APPLE)
+if(APPLE OR WIN32)
 	list(APPEND INC_SYS
 		${BOOST_INCLUDE_DIR}
 	)
diff --git a/source/blender/blenkernel/BKE_idcode.h b/source/blender/blenkernel/BKE_idcode.h
index 6de0efe2709..964a49435f1 100644
--- a/source/blender/blenkernel/BKE_idcode.h
+++ b/source/blender/blenkernel/BKE_idcode.h
@@ -42,6 +42,8 @@ bool        BKE_idcode_is_valid(short idcode);
 int         BKE_idcode_to_idfilter(const short idcode);
 short       BKE_idcode_from_idfilter(const int idfilter);
 
+int BKE_idcode_to_index(const short idcode);
+
 /**
  * Return an ID code and steps the index forward 1.
  *
diff --git a/source/blender/blenkernel/BKE_main.h b/source/blender/blenkernel/BKE_main.h
index 6487954844b..7eba01e6d5a 100644
--- a/source/blender/blenkernel/BKE_main.h
+++ b/source/blender/blenkernel/BKE_main.h
@@ -42,6 +42,8 @@
  */
 #include "DNA_listBase.h"
 
+#include "BKE_library.h"
+
 #ifdef __cplusplus
 extern "C" {
 #endif
@@ -103,7 +105,7 @@ typedef struct Main {
 	ListBase linestyle;
 	ListBase cachefiles;
 
-	char id_tag_update[256];
+	char id_tag_update[MAX_LIBARRAY];
 
 	/* Evaluation context used by viewport */
 	struct EvaluationContext *eval_ctx;
diff --git a/source/blender/blenkernel/BKE_sequencer.h b/source/blender/blenkernel/BKE_sequencer.h
index 30bb6954019..811e9136fc9 100644
--- a/source/blender/blenkernel/BKE_sequencer.h
+++ b/source/blender/blenkernel/BKE_sequencer.h
@@ -477,4 +477,6 @@ struct ImBuf *BKE_sequencer_render_mask_input(
         int cfra, int fra_offset, bool make_float);
 void BKE_sequencer_color_balance_apply(struct StripColorBalance *cb, struct ImBuf *ibuf, float mul, bool make_float, struct ImBuf *mask_input);
 
+void BKE_sequencer_all_free_anim_ibufs(int cfra);
+
 #endif  /* __BKE_SEQUENCER_H__ */
diff --git a/source/blender/blenkernel/intern/depsgraph.c b/source/blender/blenkernel/intern/depsgraph.c
index 85ecf495805..b6bfe336183 100644
--- a/source/blender/blenkernel/intern/depsgraph.c
+++ b/source/blender/blenkernel/intern/depsgraph.c
@@ -69,6 +69,7 @@
 #include "BKE_effect.h"
 #include "BKE_fcurve.h"
 #include "BKE_global.h"
+#include "BKE_idcode.h"
 #include "BKE_image.h"
 #include "BKE_key.h"
 #include "BKE_library.h"
@@ -2602,9 +2603,7 @@ void DAG_ids_flush_tagged(Main *bmain)
 		ListBase *lb = lbarray[a];
 		ID *id = lb->first;
 
-		/* we tag based on first ID type character to avoid 
-		 * looping over all ID's in case there are no tags */
-		if (id && bmain->id_tag_update[id->name[0]]) {
+		if (id && bmain->id_tag_update[BKE_idcode_to_index(GS(id->name))]) {
 			for (; id; id = id->next) {
 				if (id->tag & (LIB_TAG_ID_RECALC | LIB_TAG_ID_RECALC_DATA)) {
 					
@@ -2644,9 +2643,7 @@ void DAG_ids_check_recalc(Main *bmain, Scene *scene, bool time)
 		ListBase *lb = lbarray[a];
 		ID *id = lb->first;
 
-		/* we tag based on first ID type character to avoid 
-		 * looping over all ID's in case there are no tags */
-		if (id && bmain->id_tag_update[id->name[0]]) {
+		if (id && bmain->id_tag_update[BKE_idcode_to_index(GS(id->name))]) {
 			updated = true;
 			break;
 		}
@@ -2727,9 +2724,7 @@ void DAG_ids_clear_recalc(Main *bmain)
 		ListBase *lb = lbarray[a];
 		ID *id = lb->first;
 
-		/* we tag based on first ID type character to avoid 
-		 * looping over all ID's in case there are no tags */
-		if (id && bmain->id_tag_update[id->name[0]]) {
+		if (id && bmain->id_tag_update[BKE_idcode_to_index(GS(id->name))]) {
 			for (; id; id = id->next) {
 				if (id->tag & (LIB_TAG_ID_RECALC | LIB_TAG_ID_RECALC_DATA))
 					id->tag &= ~(LIB_TAG_ID_RECALC | LIB_TAG_ID_RECALC_DATA);
@@ -2836,12 +2831,12 @@ void DAG_id_type_tag(Main *bmain, short idtype)
 		DAG_id_type_tag(bmain, ID_SCE);
 	}
 
-	bmain->id_tag_update[((char *)&idtype)[0]] = 1;
+	bmain->id_tag_update[BKE_idcode_to_index(idtype)] = 1;
 }
 
 int DAG_id_type_tagged(Main *bmain, short idtype)
 {
-	return bmain->id_tag_update[((char *)&idtype)[0]];
+	return bmain->id_tag_update[BKE_idcode_to_index(idtype)];
 }
 
 #if 0 // UNUSED
diff --git a/source/blender/blenkernel/intern/idcode.c b/source/blender/blenkernel/intern/idcode.c
index ddc4fb37848..c76d072cb64 100644
--- a/source/blender/blenkernel/intern/idcode.c
+++ b/source/blender/blenkernel/intern/idcode.c
@@ -259,6 +259,55 @@ short BKE_idcode_from_idfilter(const int idfilter)
 }
 
 /**
+ * Convert an idcode into an index (e.g. ID_OB -> INDEX_ID_OB).
+ */
+int BKE_idcode_to_index(const short idcode)
+{
+#define CASE_IDINDEX(_id) case ID_##_id: return INDEX_ID_##_id
+
+	switch ((ID_Type)idcode) {
+		CASE_IDINDEX(AC);
+		CASE_IDINDEX(AR);
+		CASE_IDINDEX(BR);
+		CASE_IDINDEX(CA);
+		CASE_IDINDEX(CF);
+		CASE_IDINDEX(CU);
+		CASE_IDINDEX(GD);
+		CASE_IDINDEX(GR);
+		CASE_IDINDEX(IM);
+		CASE_IDINDEX(KE);
+		CASE_IDINDEX(IP);
+		CASE_IDINDEX(LA);
+		CASE_IDINDEX(LI);
+		CASE_IDINDEX(LS);
+		CASE_IDINDEX(LT);
+		CASE_IDINDEX(MA);
+		CASE_IDINDEX(MB);
+		CASE_IDINDEX(MC);
+		CASE_IDINDEX(ME);
+		CASE_IDINDEX(MSK);
+		CASE_IDINDEX(NT);
+		CASE_IDINDEX(OB);
+		CASE_IDINDEX(PAL);
+		CASE_IDINDEX(PC);
+		CASE_IDINDEX(SCE);
+		CASE_IDINDEX(SCR);
+		CASE_IDINDEX(SPK);
+		CASE_IDINDEX(SO);
+		CASE_IDINDEX(TE);
+		CASE_IDINDEX(TXT);
+		CASE_IDINDEX(VF);
+		CASE_IDINDEX(WM);
+		CASE_IDINDEX(WO);
+	}
+
+	BLI_assert(0);
+	return -1;
+
+#undef CASE_IDINDEX
+}
+
+/**
  * Convert an idcode into a name (plural).
  *
  * \param idcode: The code to convert.
diff --git a/source/blender/blenkernel/intern/library.c b/source/blender/blenkernel/intern/library.c
index 933622eed83..ca98f5d04b7 100644
--- a/source/blender/blenkernel/intern/library.c
+++ b/source/blender/blenkernel/intern/library.c
@@ -744,59 +744,55 @@ void BKE_main_lib_objects_recalc_all(Main *bmain)
  * \note MAX_LIBARRAY define should match this code */
 int set_listbasepointers(Main *main, ListBase **lb)
 {
-	int a = 0;
-
 	/* BACKWARDS! also watch order of free-ing! (mesh<->mat), first items freed last.
 	 * This is important because freeing data decreases usercounts of other datablocks,
 	 * if this data is its self freed it can crash. */
-	lb[a++] = &(main->library);  /* Libraries may be accessed from pretty much any other ID... */
-	lb[a++] = &(main->ipo);
-	lb[a++] = &(main->action); /* moved here to avoid problems when freeing with animato (aligorith) */
-	lb[a++] = &(main->key);
-	lb[a++] = &(main->gpencil); /* referenced by nodes, objects, view, scene etc, before to free after. */
-	lb[a++] = &(main->nodetree);
-	lb[a++] = &(main->image);
-	lb[a++] = &(main->tex);
-	lb[a++] = &(main->mat);
-	lb[a++] = &(main->vfont);
+	lb[INDEX_ID_LI] = &(main->library);  /* Libraries may be accessed from pretty much any other ID... */
+	lb[INDEX_ID_IP] = &(main->ipo);
+	lb[INDEX_ID_AC] = &(main->action); /* moved here to avoid problems when freeing with animato (aligorith) */
+	lb[INDEX_ID_KE] = &(main->key);
+	lb[INDEX_ID_GD] = &(main->gpencil); /* referenced by nodes, objects, view, scene etc, before to free after. */
+	lb[INDEX_ID_NT] = &(main->nodetree);
+	lb[INDEX_ID_IM] = &(main->image);
+	lb[INDEX_ID_TE] = &(main->tex);
+	lb[INDEX_ID_MA] = &(main->mat);
+	lb[INDEX_ID_VF] = &(main->vfont);
 	
 	/* Important!: When adding a new object type,
 	 * the specific data should be inserted here 
 	 */
 
-	lb[a++] = &(main->armature);
-
-	lb[a++] = &(main->cachefiles);
-	lb[a++] = &(main->mesh);
-	lb[a++] = &(main->curve);
-	lb[a++] = &(main->mball);
-
-	lb[a++] = &(main->latt);
-	lb[a++] = &(main->lamp);
-	lb[a++] = &(main->camera);
-
-	lb[a++] = &(main->text);
-	lb[a++] = &(main->sound);
-	lb[a++] = &(main->group);
-	lb[a++] = &(main->palettes);
-	lb[a++] = &(main->paintcurves);
-	lb[a++] = &(main->brush);
-	lb[a++] = &(main->speaker);
-
-	lb[a++] = &(main->world);
-	lb[a++] = &(main->movieclip);
-	lb[a++] = &(main->screen);
-	lb[a++] = &(main->object);
-	lb[a++] = &(main->linestyle); /* referenced by scenes */
-	lb[a++] = &(main->scene);
-	lb[a++] = &(main->wm);
-	lb[a++] = &(main->mask);
+	lb[INDEX_ID_AR] = &(main->armature);
+
+	lb[INDEX_ID_CF] = &(main->cachefiles);
+	lb[INDEX_ID_ME] = &(main->mesh);
+	lb[INDEX_ID_CU] = &(main->curve);
+	lb[INDEX_ID_MB] = &(main->mball);
+
+	lb[INDEX_ID_LT] = &(main->latt);
+	lb[INDEX_ID_LA] = &(main->lamp);
+	lb[INDEX_ID_CA] = &(main->camera);
+
+	lb[INDEX_ID_TXT] = &(main->text);
+	lb[INDEX_ID_SO]  = &(main->sound);
+	lb[INDEX_ID_GR]  = &(main->group);
+	lb[INDEX_ID_PAL] = &(main->palettes);
+	lb[INDEX_ID_PC]  = &(main->paintcurves);
+	lb[INDEX_ID_BR]  = &(main->brush);
+	lb[INDEX_ID_SPK] = &(main->speaker);
+
+	lb[INDEX_ID_WO]  = &(main->world);
+	lb[INDEX_ID_MC]  = &(main->movieclip);
+	lb[INDEX_ID_SCR] = &(main->screen);
+	lb[INDEX_ID_OB]  = &(main->object);
+	lb[INDEX_ID_LS]  = &(main->linestyle); /* referenced by scenes */
+	lb[INDEX_ID_SCE] = &(main->scene);
+	lb[INDEX_ID_WM]  = &(main->wm);
+	lb[INDEX_ID_MSK] = &(main->mask);
 	
-	lb[a] = NULL;
-
-	BLI_assert(a + 1 == MAX_LIBARRAY);
+	lb[INDEX_ID_NULL] = NULL;
 
-	return a;
+	return (MAX_LIBARRAY - 1);
 }
 
 /* *********** ALLOC AND FREE *****************
@@ -1645,12 +1641,12 @@ void BKE_library_make_local(Main *bmain, const Library *lib, const bool untagged
 
 		/* Do not explicitly make local non-linkable IDs (shapekeys, in fact), they are assumed to be handled
 		 * by real datablocks responsible of them. */
-		const bool do_skip = (id && BKE_idcode_is_linkable(GS(id->name)));
+		const bool do_skip = (id && !BKE_idcode_is_linkable(GS(id->name)));
 
 		for (; id; id = id_next) {
 			id->newid = NULL;
 			id_next = id->next;  /* id is possibly being inserted again */
-			
+
 			/* The check on the second line (LIB_TAG_PRE_EXISTING) is done so its
 			 * possible to tag data you don't want to be made local, used for
 			 * appending data, so any libdata already linked wont become local
diff --git a/source/blender/blenkernel/intern/sequencer.c b/source/blender/blenkernel/intern/sequencer.c
index 6067a8b2d9b..c240aa27343 100644
--- a/source/blender/blenkernel/intern/sequencer.c
+++ b/source/blender/blenkernel/intern/sequencer.c
@@ -1656,6 +1656,9 @@ static bool seq_proxy_get_fname(Editing *ed, Sequence *seq, int cfra, int render
 	else if ((proxy->storage & SEQ_STORAGE_PROXY_CUSTOM_DIR) && (proxy->storage & SEQ_STORAGE_PROXY_CUSTOM_FILE)) {
 		BLI_strncpy(dir, seq->strip->proxy->dir, sizeof(dir));
 	}
+	else if (proxy->storage & SEQ_STORAGE_PROXY_CUSTOM_FILE) {
+		BLI_strncpy(dir, seq->strip->proxy->dir, sizeof(dir));
+	}
 	else if (sanim && sanim->anim && (proxy->storage & SEQ_STORAGE_PROXY_CUSTOM_DIR)) {
 		char fname[FILE_MAXFILE];
 		BLI_strncpy(dir, seq->strip->proxy->dir, sizeof(dir));
@@ -1675,13 +1678,21 @@ static bool seq_proxy_get_fname(Editing *ed, Sequence *seq, int cfra, int render
 	if (view_id > 0)
 		BLI_snprintf(suffix, sizeof(suffix), "_%d", view_id);
 
-	if (proxy->storage & SEQ_STORAGE_PROXY_CUSTOM_FILE && sanim && sanim->anim &&
+	if (proxy->storage & SEQ_STORAGE_PROXY_CUSTOM_FILE &&
 	    ed->proxy_storage != SEQ_EDIT_PROXY_DIR_STORAGE)
 	{
-		BLI_join_dirfile(name, PROXY_MAXFILE,
-		                 dir, proxy->file);
-		BLI_path_abs(name, G.main->name);
-		BLI_snprintf(name, PROXY_MAXFILE, "%s_%s", name, suffix);
+		char fname[FILE_MAXFILE];
+		BLI_join_dirfile(fname, PROXY_MAXFILE, dir, proxy->file);
+		BLI_path_abs(fname, G.main->name);
+		if (suffix[0] != '\0') {
+			/* TODO(sergey): This will actually append suffix after extension
+			 * which is weird but how was originally coded in multiview branch.
+			 */
+			BLI_snprintf(name, PROXY_MAXFILE, "%s_%s", fname, suffix);
+		}
+		else {
+			BLI_strncpy(name, fname, PROXY_MAXFILE);
+		}
 
 		return true;
 	}
@@ -5595,3 +5606,31 @@ int BKE_sequencer_find_next_prev_edit(
 
 	return best_frame;
 }
+
+static void sequencer_all_free_anim_ibufs(ListBase *seqbase, int cfra)
+{
+	for (Sequence *seq = seqbase->first; seq != NULL; seq = seq->next) {
+		if (seq->enddisp < cfra || seq->startdisp > cfra) {
+			BKE_sequence_free_anim(seq);
+		}
+		if (seq->type == SEQ_TYPE_META) {
+			sequencer_all_free_anim_ibufs(&seq->seqbase, cfra);
+		}
+	}
+}
+
+void BKE_sequencer_all_free_anim_ibufs(int cfra)
+{
+	BKE_sequencer_cache_cleanup();
+	for (Scene *scene = G.main->scene.first;
+	     scene != NULL;
+	     scene = scene->id.next)
+	{
+		Editing *ed = BKE_sequencer_editing_get(scene, false);
+		if (ed == NULL) {
+			/* Ignore scenes without sequencer. */
+			continue;
+		}
+		sequencer_all_free_anim_ibufs(&ed->seqbase, cfra);
+	}
+}
diff --git a/source/blender/depsgraph/intern/depsgraph_query.cc b/source/blender/depsgraph/intern/depsgraph_query.cc
index cac4eaae215..7f2f6a65f5e 100644
--- a/source/blender/depsgraph/intern/depsgraph_query.cc
+++ b/source/blender/depsgraph/intern/depsgraph_query.cc
@@ -33,6 +33,7 @@
 #include "MEM_guardedalloc.h"
 
 extern "C" {
+#include "BKE_idcode.h"
 #include "BKE_main.h"
 
 #include "DEG_depsgraph_query.h"
@@ -42,7 +43,7 @@ extern "C" {
 
 bool DEG_id_type_tagged(Main *bmain, short idtype)
 {
-	return bmain->id_tag_update[((unsigned char *)&idtype)[0]] != 0;
+	return bmain->id_tag_update[BKE_idcode_to_index(idtype)] != 0;
 }
 
 short DEG_get_eval_flags_for_id(Depsgraph *graph, ID *id)
diff --git a/source/blender/depsgraph/intern/depsgraph_tag.cc b/source/blender/depsgraph/intern/depsgraph_tag.cc
index 304e3e14bae..75fb55a1b4d 100644
--- a/source/blender/depsgraph/intern/depsgraph_tag.cc
+++ b/source/blender/depsgraph/intern/depsgraph_tag.cc
@@ -43,6 +43,7 @@ extern "C" {
 
 #include "BLI_task.h"
 
+#include "BKE_idcode.h"
 #include "BKE_library.h"
 #include "BKE_main.h"
 #include "BKE_node.h"
@@ -225,10 +226,8 @@ void DEG_id_type_tag(Main *bmain, short idtype)
 		DEG_id_type_tag(bmain, ID_WO);
 		DEG_id_type_tag(bmain, ID_SCE);
 	}
-	/* We tag based on first ID type character to avoid
-	 * looping over all ID's in case there are no tags.
-	 */
-	bmain->id_tag_update[((unsigned char *)&idtype)[0]] = 1;
+
+	bmain->id_tag_update[BKE_idcode_to_index(idtype)] = 1;
 }
 
 /* Recursively push updates out to all nodes dependent on this,
@@ -339,10 +338,7 @@ void DEG_ids_check_recalc(Main *bmain, Scene *scene, bool time)
 		ListBase *lb = lbarray[a];
 		ID *id = (ID *)lb->first;
 
-		/* We tag based on first ID type character to avoid
-		 * looping over all ID's in case there are no tags.
-		 */
-		if (id && bmain->id_tag_update[(unsigned char)id->name[0]]) {
+		if (id && bmain->id_tag_update[BKE_idcode_to_index(GS(id->name))]) {
 			updated = true;
 			break;
 		}
@@ -367,10 +363,7 @@ void DEG_ids_clear_recalc(Main *bmain)
 		ListBase *lb = lbarray[a];
 		ID *id = (ID *)lb->first;
 
-		/* We tag based on first ID type character to avoid
-		 * looping over all ID's in case there are no tags.
-		 */
-		if (id && bmain->id_tag_update[(unsigned char)id->name[0]]) {
+		if (id && bmain->id_tag_update[BKE_idcode_to_index(GS(id->name))]) {
 			for (; id; id = (ID *)id->next) {
 				id->tag &= ~(LIB_TAG_ID_RECALC | LIB_TAG_ID_RECALC_DATA);
 
diff --git a/source/blender/editors/gpencil/gpencil_data.c b/source/blender/editors/gpencil/gpencil_data.c
index 6f2ebbed7d7..876f873e5fa 100644
--- a/source/blender/editors/gpencil/gpencil_data.c
+++ b/source/blender/editors/gpencil/gpencil_data.c
@@ -747,7 +747,7 @@ static int gp_stroke_arrange_exec(bContext *C, wmOperator *op)
 	bGPDframe *gpf = gpl->actframe;
 	/* temp listbase to store selected strokes */
 	ListBase selected = {NULL};
-	const int direction = RNA_enum_get(op->ptr, "type");
+	const int direction = RNA_enum_get(op->ptr, "direction");
 
 	/* verify if any selected stroke is in the extreme of the stack and select to move */
 	for (gps = gpf->strokes.first; gps; gps = gps->next) {
diff --git a/source/blender/editors/gpencil/gpencil_paint.c b/source/blender/editors/gpencil/gpencil_paint.c
index eefdee70304..e7e39a85792 100644
--- a/source/blender/editors/gpencil/gpencil_paint.c
+++ b/source/blender/editors/gpencil/gpencil_paint.c
@@ -692,6 +692,7 @@ static void gp_stroke_newfrombuffer(tGPsdata *p)
 	bGPDspoint *pt;
 	tGPspoint *ptc;
 	bGPDbrush *brush = p->brush;
+	ToolSettings *ts = p->scene->toolsettings;
 	
 	int i, totelem;
 	/* since strokes are so fine, when using their depth we need a margin otherwise they might get missed */
@@ -925,8 +926,16 @@ static void gp_stroke_newfrombuffer(tGPsdata *p)
 	gps->palcolor = palcolor;
 	strcpy(gps->colorname, palcolor->info);
 
-	/* add stroke to frame */
-	BLI_addtail(&p->gpf->strokes, gps);
+	/* add stroke to frame, usually on tail of the listbase, but if on back is enabled the stroke is added on listbase head 
+	* because the drawing order is inverse and the head stroke is the first to draw. This is very useful for artist
+	* when drawing the background
+	*/
+	if ((ts->gpencil_flags & GP_TOOL_FLAG_PAINT_ONBACK) && (p->paintmode != GP_PAINTMODE_DRAW_POLY)) {
+		BLI_addhead(&p->gpf->strokes, gps);
+	}
+	else {
+		BLI_addtail(&p->gpf->strokes, gps);
+	}
 	gp_stroke_added_enable(p);
 }
 
@@ -2278,6 +2287,28 @@ static void gpencil_stroke_end(wmOperator *op)
 	p->gpf = NULL;
 }
 
+/* Move last stroke in the listbase to the head to be drawn below all previous strokes in the layer */
+static void gpencil_move_last_stroke_to_back(bContext *C)
+{
+	/* move last stroke (the polygon) to head of the listbase stroke to draw on back of all previous strokes */
+	bGPdata *gpd = ED_gpencil_data_get_active(C);
+	bGPDlayer *gpl = BKE_gpencil_layer_getactive(gpd);
+
+	/* sanity checks */
+	if (ELEM(NULL, gpd, gpl, gpl->actframe)) {
+		return;
+	}
+
+	bGPDframe *gpf = gpl->actframe;
+	bGPDstroke *gps = gpf->strokes.last;
+	if (ELEM(NULL, gps)) {
+		return;
+	}
+
+	BLI_remlink(&gpf->strokes, gps);
+	BLI_insertlinkbefore(&gpf->strokes, gpf->strokes.first, gps);
+}
+
 /* events handling during interactive drawing part of operator */
 static int gpencil_draw_modal(bContext *C, wmOperator *op, const wmEvent *event)
 {
@@ -2334,6 +2365,12 @@ static int gpencil_draw_modal(bContext *C, wmOperator *op, const wmEvent *event)
 	if (ELEM(event->type, RETKEY, PADENTER, ESCKEY, SPACEKEY, EKEY)) {
 		/* exit() ends the current stroke before cleaning up */
 		/* printf("\t\tGP - end of paint op + end of stroke\n"); */
+		/* if drawing polygon and enable on back, must move stroke */
+		if ((p->scene->toolsettings->gpencil_flags & GP_TOOL_FLAG_PAINT_ONBACK) && (p->paintmode == GP_PAINTMODE_DRAW_POLY)) {
+			if (p->flags & GP_PAINTFLAG_STROKEADDED) {
+				gpencil_move_last_stroke_to_back(C);
+			}
+		}
 		p->status = GP_STATUS_DONE;
 		estate = OPERATOR_FINISHED;
 	}
@@ -2390,9 +2427,12 @@ static int gpencil_draw_modal(bContext *C, wmOperator *op, const wmEvent *event)
 			}
 			else {
 				/* printf("\t\tGP - end of stroke + op\n"); */
-				/* disable paint session */
-				p->scene->toolsettings->gpencil_flags &= ~GP_TOOL_FLAG_PAINTSESSIONS_ON;
-
+				/* if drawing polygon and enable on back, must move stroke */
+				if ((p->scene->toolsettings->gpencil_flags & GP_TOOL_FLAG_PAINT_ONBACK) && (p->paintmode == GP_PAINTMODE_DRAW_POLY)) {
+					if (p->flags & GP_PAINTFLAG_STROKEADDED) {
+						gpencil_move_last_stroke_to_back(C);
+					}
+				}
 				p->status = GP_STATUS_DONE;
 				estate = OPERATOR_FINISHED;
 			}
@@ -2422,9 +2462,6 @@ static int gpencil_draw_modal(bContext *C, wmOperator *op, const wmEvent *event)
 					in_bounds = true;
 				}
 				else {
-					/* disable paint session */
-					p->scene->toolsettings->gpencil_flags &= ~GP_TOOL_FLAG_PAINTSESSIONS_ON;
-
 					/* Out of bounds, or invalid in some other way */
 					p->status = GP_STATUS_ERROR;
 					estate = OPERATOR_CANCELLED;
@@ -2441,9 +2478,6 @@ static int gpencil_draw_modal(bContext *C, wmOperator *op, const wmEvent *event)
 				in_bounds = BLI_rcti_isect_pt_v(&region_rect, event->mval);
 			}
 			else {
-				/* disable paint session */
-				p->scene->toolsettings->gpencil_flags &= ~GP_TOOL_FLAG_PAINTSESSIONS_ON;
-
 				/* No region */
 				p->status = GP_STATUS_ERROR;
 				estate = OPERATOR_CANCELLED;
@@ -2471,9 +2505,6 @@ static int gpencil_draw_modal(bContext *C, wmOperator *op, const wmEvent *event)
 				p = gpencil_stroke_begin(C, op);
 				
 				if (p->status == GP_STATUS_ERROR) {
-					/* disable paint session */
-					p->scene->toolsettings->gpencil_flags &= ~GP_TOOL_FLAG_PAINTSESSIONS_ON;
-
 					estate = OPERATOR_CANCELLED;
 				}
 			}
@@ -2482,9 +2513,12 @@ static int gpencil_draw_modal(bContext *C, wmOperator *op, const wmEvent *event)
 				 * NOTE: Don't eter this case if an error occurred while finding the
 				 *       region (as above)
 				 */
-				/* disable paint session */
-				p->scene->toolsettings->gpencil_flags &= ~GP_TOOL_FLAG_PAINTSESSIONS_ON;
-
+				/* if drawing polygon and enable on back, must move stroke */
+				if ((p->scene->toolsettings->gpencil_flags & GP_TOOL_FLAG_PAINT_ONBACK) && (p->paintmode == GP_PAINTMODE_DRAW_POLY)) {
+					if (p->flags & GP_PAINTFLAG_STROKEADDED) {
+						gpencil_move_last_stroke_to_back(C);
+					}
+				}
 				p->status = GP_STATUS_DONE;
 				estate = OPERATOR_FINISHED;
 			}
diff --git a/source/blender/editors/interface/interface_templates.c b/source/blender/editors/interface/interface_templates.c
index b3f2192c26c..bb9da1a141b 100644
--- a/source/blender/editors/interface/interface_templates.c
+++ b/source/blender/editors/interface/interface_templates.c
@@ -857,8 +857,8 @@ static uiLayout *draw_modifier(
 		
 		/* mode enabling buttons */
 		UI_block_align_begin(block);
-		/* Softbody not allowed in this situation, enforce! */
-		if (((md->type != eModifierType_Softbody && md->type != eModifierType_Collision) || !(ob->pd && ob->pd->deflect)) &&
+		/* Collision and Surface are always enabled, hide buttons! */
+		if (((md->type != eModifierType_Collision) || !(ob->pd && ob->pd->deflect)) &&
 		    (md->type != eModifierType_Surface) )
 		{
 			uiItemR(row, &ptr, "show_render", 0, "", ICON_NONE);
diff --git a/source/blender/editors/object/object_relations.c b/source/blender/editors/object/object_relations.c
index 6a9f5a9bfaf..fb067ccca81 100644
--- a/source/blender/editors/object/object_relations.c
+++ b/source/blender/editors/object/object_relations.c
@@ -1381,7 +1381,7 @@ static int move_to_layer_exec(bContext *C, wmOperator *op)
 			/* upper byte is used for local view */
 			local = base->lay & 0xFF000000;
 			base->lay = lay + local;
-			base->object->lay = lay;
+			base->object->lay = base->lay;
 			/* if (base->object->type == OB_LAMP) is_lamp = true; */
 		}
 		CTX_DATA_END;
diff --git a/source/blender/editors/space_graph/space_graph.c b/source/blender/editors/space_graph/space_graph.c
index 2582ba4be8d..b35d1b2c777 100644
--- a/source/blender/editors/space_graph/space_graph.c
+++ b/source/blender/editors/space_graph/space_graph.c
@@ -687,7 +687,7 @@ static void graph_id_remap(ScrArea *UNUSED(sa), SpaceLink *slink, ID *old_id, ID
 		return;
 	}
 
-	if ((ID *)sgraph->ads->filter_grp == old_id) {
+	if (sgraph->ads && (ID *)sgraph->ads->filter_grp == old_id) {
 		sgraph->ads->filter_grp = (Group *)new_id;
 	}
 }
diff --git a/source/blender/editors/space_view3d/drawvolume.c b/source/blender/editors/space_view3d/drawvolume.c
index a667b46d676..29c06388cf6 100644
--- a/source/blender/editors/space_view3d/drawvolume.c
+++ b/source/blender/editors/space_view3d/drawvolume.c
@@ -378,11 +378,13 @@ void draw_smoke_volume(SmokeDomainSettings *sds, Object *ob,
 
 	/* setup buffer and draw */
 
-	int gl_depth = 0, gl_blend = 0;
+	int gl_depth = 0, gl_blend = 0, gl_depth_write = 0;
 	glGetBooleanv(GL_BLEND, (GLboolean *)&gl_blend);
 	glGetBooleanv(GL_DEPTH_TEST, (GLboolean *)&gl_depth);
+	glGetBooleanv(GL_DEPTH_WRITEMASK, (GLboolean *)&gl_depth_write);
 
 	glEnable(GL_DEPTH_TEST);
+	glDepthMask(GL_FALSE);
 	glEnable(GL_BLEND);
 	glBlendFunc(GL_ONE, GL_ONE_MINUS_SRC_ALPHA);
 
@@ -421,6 +423,8 @@ void draw_smoke_volume(SmokeDomainSettings *sds, Object *ob,
 
 	GPU_shader_unbind();
 
+	glDepthMask(gl_depth_write);
+
 	if (!gl_blend) {
 		glDisable(GL_BLEND);
 	}
diff --git a/source/blender/gpu/intern/gpu_texture.c b/source/blender/gpu/intern/gpu_texture.c
index 827c52c9a5f..54f0003c086 100644
--- a/source/blender/gpu/intern/gpu_texture.c
+++ b/source/blender/gpu/intern/gpu_texture.c
@@ -55,7 +55,8 @@ struct GPUTexture {
 	int number;         /* number for multitexture binding */
 	int refcount;       /* reference count */
 	GLenum target;      /* GL_TEXTURE_* */
-	GLenum target_base; /* same as target, (but no multisample) */
+	GLenum target_base; /* same as target, (but no multisample)
+	                     * use it for unbinding */
 	GLuint bindcode;    /* opengl identifier for texture */
 	int fromblender;    /* we got the texture from Blender */
 
@@ -374,6 +375,9 @@ GPUTexture *GPU_texture_from_blender(Image *ima, ImageUser *iuser, int textarget
 	GLint bindcode = GPU_verify_image(ima, iuser, textarget, 0, 0, mipmap, is_data);
 	GPU_update_image_time(ima, time);
 
+	/* see GPUInput::textarget: it can take two values - GL_TEXTURE_2D and GL_TEXTURE_CUBE_MAP
+	 * these values are correct for glDisable, so textarget can be safely used in
+	 * GPU_texture_bind/GPU_texture_unbind through tex->target_base */
 	if (textarget == GL_TEXTURE_2D)
 		gputt = TEXTARGET_TEXTURE_2D;
 	else
@@ -390,7 +394,7 @@ GPUTexture *GPU_texture_from_blender(Image *ima, ImageUser *iuser, int textarget
 	tex->number = -1;
 	tex->refcount = 1;
 	tex->target = textarget;
-	tex->target_base = GL_TEXTURE_2D;
+	tex->target_base = textarget;
 	tex->fromblender = 1;
 
 	ima->gputexture[gputt] = tex;
@@ -626,11 +630,11 @@ void GPU_texture_bind(GPUTexture *tex, int number)
 	GLenum arbnumber = (GLenum)((GLuint)GL_TEXTURE0 + number);
 	if (number != 0) glActiveTexture(arbnumber);
 	if (tex->bindcode != 0) {
-		glBindTexture(tex->target, tex->bindcode);
+		glBindTexture(tex->target_base, tex->bindcode);
 	}
 	else
-		GPU_invalid_tex_bind(tex->target);
-	glEnable(tex->target);
+		GPU_invalid_tex_bind(tex->target_base);
+	glEnable(tex->target_base);
 	if (number != 0) glActiveTexture(GL_TEXTURE0);
 
 	tex->number = number;
@@ -652,8 +656,6 @@ void GPU_texture_unbind(GPUTexture *tex)
 
 	GLenum arbnumber = (GLenum)((GLuint)GL_TEXTURE0 + tex->number);
 	if (tex->number != 0) glActiveTexture(arbnumber);
-	glBindTexture(tex->target, 0);
-	glDisable(tex->target);
 	glBindTexture(tex->target_base, 0);
 	glDisable(tex->target_base);
 	if (tex->number != 0) glActiveTexture(GL_TEXTURE0);
diff --git a/source/blender/makesdna/DNA_ID.h b/source/blender/makesdna/DNA_ID.h
index 725eefe5155..45eb31235f5 100644
--- a/source/blender/makesdna/DNA_ID.h
+++ b/source/blender/makesdna/DNA_ID.h
@@ -379,6 +379,45 @@ enum {
 	FILTER_ID_CF        = (1 << 27),
 };
 
+/* IMPORTANT: this enum matches the order currently use in set_lisbasepointers,
+ * keep them in sync! */
+enum {
+	INDEX_ID_LI = 0,
+	INDEX_ID_IP,
+	INDEX_ID_AC,
+	INDEX_ID_KE,
+	INDEX_ID_GD,
+	INDEX_ID_NT,
+	INDEX_ID_IM,
+	INDEX_ID_TE,
+	INDEX_ID_MA,
+	INDEX_ID_VF,
+	INDEX_ID_AR,
+	INDEX_ID_CF,
+	INDEX_ID_ME,
+	INDEX_ID_CU,
+	INDEX_ID_MB,
+	INDEX_ID_LT,
+	INDEX_ID_LA,
+	INDEX_ID_CA,
+	INDEX_ID_TXT,
+	INDEX_ID_SO,
+	INDEX_ID_GR,
+	INDEX_ID_PAL,
+	INDEX_ID_PC,
+	INDEX_ID_BR,
+	INDEX_ID_SPK,
+	INDEX_ID_WO,
+	INDEX_ID_MC,
+	INDEX_ID_SCR,
+	INDEX_ID_OB,
+	INDEX_ID_LS,
+	INDEX_ID_SCE,
+	INDEX_ID_WM,
+	INDEX_ID_MSK,
+	INDEX_ID_NULL,
+};
+
 #ifdef __cplusplus
 }
 #endif
diff --git a/source/blender/makesdna/DNA_scene_types.h b/source/blender/makesdna/DNA_scene_types.h
index 66b79282533..d4edc8a1c20 100644
--- a/source/blender/makesdna/DNA_scene_types.h
+++ b/source/blender/makesdna/DNA_scene_types.h
@@ -2052,6 +2052,8 @@ typedef enum eGPencil_Flags {
 	GP_TOOL_FLAG_PAINTSESSIONS_ON       = (1 << 0),
 	/* When creating new frames, the last frame gets used as the basis for the new one */
 	GP_TOOL_FLAG_RETAIN_LAST            = (1 << 1),
+	/* Add the strokes below all strokes in the layer */
+	GP_TOOL_FLAG_PAINT_ONBACK = (1 << 2)
 } eGPencil_Flags;
 
 /* toolsettings->gpencil_src */
diff --git a/source/blender/makesrna/intern/rna_gpencil.c b/source/blender/makesrna/intern/rna_gpencil.c
index 80f4d5d05b6..7424c190501 100644
--- a/source/blender/makesrna/intern/rna_gpencil.c
+++ b/source/blender/makesrna/intern/rna_gpencil.c
@@ -1381,8 +1381,9 @@ static void rna_def_gpencil_palettecolor(BlenderRNA *brna)
 	RNA_def_property_update(prop, NC_GPENCIL | ND_DATA, "rna_GPencil_update");
 
 	/* Name */
-	prop = RNA_def_property(srna, "info", PROP_STRING, PROP_NONE);
-	RNA_def_property_ui_text(prop, "Info", "Color name");
+	prop = RNA_def_property(srna, "name", PROP_STRING, PROP_NONE);
+	RNA_def_property_string_sdna(prop, NULL, "info");
+	RNA_def_property_ui_text(prop, "Name", "Color name");
 	RNA_def_property_string_funcs(prop, NULL, NULL, "rna_GPencilPaletteColor_info_set");
 	RNA_def_struct_name_property(srna, prop);
 	RNA_def_property_update(prop, NC_GPENCIL | ND_DATA, "rna_GPencil_update");
diff --git a/source/blender/makesrna/intern/rna_scene.c b/source/blender/makesrna/intern/rna_scene.c
index f39e7dc10b7..06b5bd87dd4 100644
--- a/source/blender/makesrna/intern/rna_scene.c
+++ b/source/blender/makesrna/intern/rna_scene.c
@@ -2593,6 +2593,12 @@ static void rna_def_tool_settings(BlenderRNA  *brna)
 	                         "are included as the basis for the new one");
 	RNA_def_property_update(prop, NC_SCENE | ND_TOOLSETTINGS, NULL);
 	
+	prop = RNA_def_property(srna, "use_gpencil_draw_onback", PROP_BOOLEAN, PROP_NONE);
+	RNA_def_property_boolean_sdna(prop, NULL, "gpencil_flags", GP_TOOL_FLAG_PAINT_ONBACK);
+	RNA_def_property_ui_text(prop, "Draw Strokes on Back",
+		"When draw new strokes, the new stroke is drawn below of all strokes in the layer");
+	RNA_def_property_update(prop, NC_SCENE | ND_TOOLSETTINGS, NULL);
+
 	prop = RNA_def_property(srna, "grease_pencil_source", PROP_ENUM, PROP_NONE);
 	RNA_def_property_enum_bitflag_sdna(prop, NULL, "gpencil_src");
 	RNA_def_property_enum_items(prop, gpencil_source_3d_items);
diff --git a/source/blender/render/intern/raytrace/rayobject_rtbuild.cpp b/source/blender/render/intern/raytrace/rayobject_rtbuild.cpp
index 02a49fc3c8f..724a809077e 100644
--- a/source/blender/render/intern/raytrace/rayobject_rtbuild.cpp
+++ b/source/blender/render/intern/raytrace/rayobject_rtbuild.cpp
@@ -42,6 +42,10 @@
 #include "BLI_math.h"
 #include "BLI_utildefines.h"
 
+#if __cplusplus >= 201103L
+using std::isfinite;
+#endif
+
 static bool selected_node(RTBuilder::Object *node)
 {
 	return node->selected;
diff --git a/source/blender/render/intern/source/pipeline.c b/source/blender/render/intern/source/pipeline.c
index fc99bedfa0f..363267bb363 100644
--- a/source/blender/render/intern/source/pipeline.c
+++ b/source/blender/render/intern/source/pipeline.c
@@ -2784,6 +2784,7 @@ static void do_render_all_options(Render *re)
 
 	/* ensure no images are in memory from previous animated sequences */
 	BKE_image_all_free_anim_ibufs(re->r.cfra);
+	BKE_sequencer_all_free_anim_ibufs(re->r.cfra);
 
 	if (RE_engine_render(re, 1)) {
 		/* in this case external render overrides all */
diff --git a/source/gameengine/GamePlayer/ghost/GPG_ghost.cpp b/source/gameengine/GamePlayer/ghost/GPG_ghost.cpp
index d033afacc08..edbbf93bf7a 100644
--- a/source/gameengine/GamePlayer/ghost/GPG_ghost.cpp
+++ b/source/gameengine/GamePlayer/ghost/GPG_ghost.cpp
@@ -61,6 +61,7 @@ extern "C"
 
 #include "DNA_scene_types.h"
 #include "DNA_userdef_types.h"
+#include "DNA_genfile.h"
 
 #include "BLO_readfile.h"
 #include "BLO_runtime.h"
@@ -492,6 +493,8 @@ int main(
 	// freeing up GPU_Textures works correctly.
 	BLI_threadapi_init();
 
+	DNA_sdna_current_init();
+
 	RNA_init();
 
 	init_nodesystem();