Merged changes in the trunk up to revision 43748.

Conflicts resolved:
source/blender/editors/include/UI_resources.h

doc/python_api/examples/bge.constraints.py

@@ -1,6 +1,6 @@
 """
 Basic Physics Constraint
-++++++++++++++++++++++
+++++++++++++++++++++++++
 Example of how to create a hinge Physics Constraint between two objects.
 """
 from bge import logic

doc/python_api/examples/bge.texture.1.py

@@ -1,6 +1,6 @@
 """
-Texture replacement
+Texture Replacement
-++++++++++++++++++++++
++++++++++++++++++++
 Example of how to replace a texture in game with an external image.
 createTexture() and removeTexture() are to be called from a module Python
 Controller.

doc/python_api/examples/bge.texture.py

@@ -1,6 +1,6 @@
 """
 Basic Video Playback
-++++++++++++++++++++++
+++++++++++++++++++++
 Example of how to replace a texture in game with a video. It needs to run everyframe
 """
 import bge

doc/python_api/rst/bge.constraints.rst

@@ -5,6 +5,11 @@ Physics Constraints (bge.constraints)
 .. module:: bge.constraints
 
 .. literalinclude:: ../examples/bge.constraints.py
+   :language: rest
+   :lines: 2-4
+
+.. literalinclude:: ../examples/bge.constraints.py
+   :lines: 6-
 
 .. function:: createConstraint(physicsid, physicsid2, constrainttype, [pivotX, pivotY, pivotZ, [axisX, axisY, axisZ, [flag]]]])
 

doc/python_api/rst/bge.render.rst

@@ -77,6 +77,14 @@ Functions
    
    :rtype: integer
 
+.. function:: setWindowSize(width, height)
+
+   Set the width and height of the window (in pixels). This also works for fullscreen applications.
+   
+   :type width: integer
+   :type height: integer
+
+
 .. function:: makeScreenshot(filename)
 
    Writes a screenshot to the given filename.

doc/python_api/rst/bge.texture.rst

@@ -37,8 +37,18 @@ When the texture object is deleted, the new texture is deleted and the old textu
 .. module:: bge.texture
 
 .. literalinclude:: ../examples/bge.texture.py
+   :language: rest
+   :lines: 2-4
+
+.. literalinclude:: ../examples/bge.texture.py
+   :lines: 6-
 
 .. literalinclude:: ../examples/bge.texture.1.py
+   :language: rest
+   :lines: 2-6
+
+.. literalinclude:: ../examples/bge.texture.1.py
+   :lines: 8-
 
 .. class:: VideoFFmpeg(file [, capture=-1, rate=25.0, width=0, height=0])
 

doc/python_api/rst/bge.types.rst

@@ -66,6 +66,12 @@ Game Types (bge.types)
 
       :type: dictionary {:ref:`keycode<keyboard-keys>`::ref:`status<input-status>`, ...}
 
+   .. attribute:: active_events
+
+      A dictionary containing the status of only the active keyboard events or keys. (read-only).
+
+      :type: dictionary {:ref:`keycode<keyboard-keys>`::ref:`status<input-status>`, ...}
+
 .. class:: SCA_PythonMouse(PyObjectPlus)
 
    The current mouse.
@@ -75,6 +81,12 @@ Game Types (bge.types)
       a dictionary containing the status of each mouse event. (read-only).
 
       :type: dictionary {:ref:`keycode<mouse-keys>`::ref:`status<input-status>`, ...}
+
+   .. attribute:: active_events
+
+      a dictionary containing the status of only the active mouse events. (read-only).
+
+      :type: dictionary {:ref:`keycode<mouse-keys>`::ref:`status<input-status>`, ...}
       
    .. attribute:: position
 
@@ -942,7 +954,7 @@ Game Types (bge.types)
 
       .. deprecated:: use :data:`localPosition` and :data:`worldPosition`.
 
-      :type: :class:`mathurils.Vector`
+      :type: :class:`mathutils.Vector`
 
    .. attribute:: orientation
 
@@ -980,7 +992,7 @@ Game Types (bge.types)
 
    .. attribute:: worldScale
 
-      The object's world scaling factor. Read-only. [sx, sy, sz]
+      The object's world scaling factor. [sx, sy, sz]
 
       :type: :class:`mathutils.Vector`
 
@@ -995,6 +1007,18 @@ Game Types (bge.types)
       The object's world position. [x, y, z]
 
       :type: :class:`mathutils.Vector`
+
+   .. attribute:: localTransform
+
+      The object's local space transform matrix. 4x4 Matrix.
+
+      :type: :class:`mathutils.Matrix`
+
+   .. attribute:: worldTransform
+
+      The object's world space transform matrix. 4x4 Matrix.
+
+      :type: :class:`mathutils.Matrix`
 	  
    .. attribute:: localLinearVelocity
       

extern/libmv/third_party/glog/src/config.h

@@ -2,14 +2,14 @@
 /* src/config.h.in.  Generated from configure.ac by autoheader.  */
 
 /* Namespace for Google classes */
-#ifdef __APPLE__
+#if defined(__APPLE__)
  #include "config_mac.h"
-#elif __FreeBSD__
+#elif defined(__FreeBSD__) || defined(__FreeBSD_kernel__)
  #include "config_freebsd.h"
-#elif __MINGW32__
+#elif defined(__MINGW32__)
  #include "windows/config.h"
-#elif __GNUC__
+#elif defined(__linux__)
  #include "config_linux.h"
-#elif _MSC_VER
+#elif defined(_MSC_VER)
  #include "windows/config.h"
 #endif

extern/libmv/third_party/glog/src/config_linux.h

@@ -133,8 +133,10 @@
 /* How to access the PC from a struct ucontext */
 #if defined(_M_X64) || defined(__amd64__) || defined(__x86_64__)
   #define PC_FROM_UCONTEXT uc_mcontext.gregs[REG_RIP]
-#else
+#elif defined(_M_IX86) || defined(__i386__) || defined(__i486__) || defined(__i586__) || defined(__i686__)
   #define PC_FROM_UCONTEXT uc_mcontext.gregs[REG_EIP]
+#else
+  #undef PC_FROM_UCONTEXT
 #endif
 
 /* Define to necessary symbol if this constant uses a non-standard name on

intern/boolop/SConscript

@@ -8,6 +8,7 @@ incs += ' ../../source/blender/blenlib'
 defs = []
 
 if not env['WITH_BF_CARVE']:
+    import os
     sources = env.Glob('intern/*.cpp')
     sources.remove('intern' + os.sep + 'BOP_CarveInterface.cpp')
 else:

intern/boolop/intern/BOP_CarveInterface.cpp

@@ -15,12 +15,13 @@
  * along with this program; if not, write to the Free Software Foundation,
  * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
  *
- * The Original Code is Copyright (C) 2001-2002 by NaN Holding BV.
+ * The Original Code is Copyright (C) 2010 by the Blender Foundation.
  * All rights reserved.
  *
  * The Original Code is: all of this file.
  *
- * Contributor(s): none yet.
+ * Contributor(s): Ken Hughes,
+ *                 Sergey Sharybin.
  *
  * ***** END GPL LICENSE BLOCK *****
  */
@@ -36,6 +37,9 @@
 #include <carve/interpolator.hpp>
 #include <carve/rescale.hpp>
 
+#include <iostream>
+
+using namespace carve::mesh;
 typedef unsigned int uint;
 
 #define MAX(x,y) ((x)>(y)?(x):(y))
@@ -61,10 +65,168 @@ static int isFacePlanar(CSG_IFace &face, std::vector<carve::geom3d::Vector> &ver
 	return 1;
 }
 
-static carve::mesh::MeshSet<3> *Carve_addMesh(CSG_FaceIteratorDescriptor& face_it,
+static MeshSet<3> *Carve_meshSetFromMeshes(std::vector<MeshSet<3>::mesh_t*> &meshes)
-                                              CSG_VertexIteratorDescriptor& vertex_it,
+{
-                                              carve::interpolate::FaceAttr<uint> &oface_num,
+	std::vector<MeshSet<3>::mesh_t*> new_meshes;
-                                              uint &num_origfaces )
+
+	std::vector<MeshSet<3>::mesh_t*>::iterator it = meshes.begin();
+	for(; it!=meshes.end(); it++) {
+		MeshSet<3>::mesh_t *mesh = *it;
+		MeshSet<3>::mesh_t *new_mesh = new MeshSet<3>::mesh_t(mesh->faces);
+
+		new_meshes.push_back(new_mesh);
+	}
+
+	return new MeshSet<3>(new_meshes);
+}
+
+static void Carve_getIntersectedOperandMeshes(std::vector<MeshSet<3>::mesh_t*> &meshes,
+                                              std::vector<MeshSet<3>::aabb_t> &precomputedAABB,
+                                              MeshSet<3>::aabb_t &otherAABB,
+                                              std::vector<MeshSet<3>::mesh_t*> &operandMeshes)
+{
+	std::vector<MeshSet<3>::mesh_t*>::iterator it = meshes.begin();
+	std::vector<MeshSet<3>::aabb_t>::iterator aabb_it = precomputedAABB.begin();
+	std::vector<MeshSet<3>::aabb_t> usedAABB;
+
+	while(it != meshes.end()) {
+		MeshSet<3>::mesh_t *mesh = *it;
+		MeshSet<3>::aabb_t aabb = mesh->getAABB();
+		bool isIntersect = false;
+
+		std::vector<MeshSet<3>::aabb_t>::iterator used_it = usedAABB.begin();
+		for(; used_it!=usedAABB.end(); used_it++) {
+			MeshSet<3>::aabb_t usedAABB = *used_it;
+
+			if(usedAABB.intersects(aabb) && usedAABB.intersects(otherAABB)) {
+				isIntersect = true;
+				break;
+			}
+		}
+
+		if(!isIntersect) {
+			operandMeshes.push_back(mesh);
+			usedAABB.push_back(aabb);
+
+			it = meshes.erase(it);
+			aabb_it = precomputedAABB.erase(aabb_it);
+		}
+		else {
+			it++;
+			aabb_it++;
+		}
+	}
+}
+
+static MeshSet<3> *Carve_getIntersectedOperand(std::vector<MeshSet<3>::mesh_t*> &meshes,
+                                               std::vector<MeshSet<3>::aabb_t> &precomputedAABB,
+                                               MeshSet<3>::aabb_t &otherAABB)
+{
+	std::vector<MeshSet<3>::mesh_t*> operandMeshes;
+	Carve_getIntersectedOperandMeshes(meshes, precomputedAABB, otherAABB, operandMeshes);
+
+	return Carve_meshSetFromMeshes(operandMeshes);
+}
+
+static MeshSet<3> *Carve_unionIntersectingMeshes(MeshSet<3> *poly,
+                                                 std::vector<MeshSet<3>::aabb_t> &precomputedAABB,
+                                                 MeshSet<3>::aabb_t &otherAABB,
+                                                 carve::interpolate::FaceAttr<uint> &oface_num)
+{
+	if(poly->meshes.size()<=1)
+		return poly;
+
+	carve::csg::CSG csg;
+
+	oface_num.installHooks(csg);
+	csg.hooks.registerHook(new carve::csg::CarveTriangulator, carve::csg::CSG::Hooks::PROCESS_OUTPUT_FACE_BIT);
+
+	std::vector<MeshSet<3>::mesh_t*> orig_meshes =
+			std::vector<MeshSet<3>::mesh_t*>(poly->meshes.begin(), poly->meshes.end());
+
+	MeshSet<3> *left = Carve_getIntersectedOperand(orig_meshes, precomputedAABB, otherAABB);
+
+	while(orig_meshes.size()) {
+		MeshSet<3> *right = Carve_getIntersectedOperand(orig_meshes, precomputedAABB, otherAABB);
+
+		try {
+			if(left->meshes.size()==0) {
+				delete left;
+
+				left = right;
+			}
+			else {
+				MeshSet<3> *result = csg.compute(left, right, carve::csg::CSG::UNION, NULL, carve::csg::CSG::CLASSIFY_EDGE);
+
+				delete left;
+				delete right;
+
+				left = result;
+			}
+		}
+		catch(carve::exception e) {
+			std::cerr << "CSG failed, exception " << e.str() << std::endl;
+
+			delete right;
+		}
+		catch(...) {
+			delete left;
+			delete right;
+
+			throw "Unknown error in Carve library";
+		}
+	}
+
+	return left;
+}
+
+static MeshSet<3>::aabb_t Carve_computeAABB(MeshSet<3> *poly,
+                                            std::vector<MeshSet<3>::aabb_t> &precomputedAABB)
+{
+	MeshSet<3>::aabb_t overallAABB;
+	std::vector<MeshSet<3>::mesh_t*>::iterator it = poly->meshes.begin();
+
+	for(; it!=poly->meshes.end(); it++) {
+		MeshSet<3>::aabb_t aabb;
+		MeshSet<3>::mesh_t *mesh = *it;
+
+		aabb = mesh->getAABB();
+		precomputedAABB.push_back(aabb);
+
+		overallAABB.unionAABB(aabb);
+	}
+
+	return overallAABB;
+}
+
+static void Carve_prepareOperands(MeshSet<3> **left_r, MeshSet<3> **right_r,
+                                  carve::interpolate::FaceAttr<uint> &oface_num)
+{
+	MeshSet<3> *left, *right;
+
+	std::vector<MeshSet<3>::aabb_t> left_precomputedAABB;
+	std::vector<MeshSet<3>::aabb_t> right_precomputedAABB;
+
+	MeshSet<3>::aabb_t leftAABB = Carve_computeAABB(*left_r, left_precomputedAABB);
+	MeshSet<3>::aabb_t rightAABB = Carve_computeAABB(*right_r, right_precomputedAABB);
+
+	left = Carve_unionIntersectingMeshes(*left_r, left_precomputedAABB, rightAABB, oface_num);
+	right = Carve_unionIntersectingMeshes(*right_r, right_precomputedAABB, leftAABB, oface_num);
+
+	if(left != *left_r)
+		delete *left_r;
+
+	if(right != *right_r)
+		delete *right_r;
+
+	*left_r = left;
+	*right_r = right;
+}
+
+static MeshSet<3> *Carve_addMesh(CSG_FaceIteratorDescriptor &face_it,
+                                 CSG_VertexIteratorDescriptor &vertex_it,
+                                 carve::interpolate::FaceAttr<uint> &oface_num,
+                                 uint &num_origfaces)
 {
 	CSG_IVertex vertex;
 	std::vector<carve::geom3d::Vector> vertices;
@@ -126,12 +288,12 @@ static carve::mesh::MeshSet<3> *Carve_addMesh(CSG_FaceIteratorDescriptor& face_i
 		}
 	}
 
-	carve::mesh::MeshSet<3> *poly = new carve::mesh::MeshSet<3> (vertices, numfaces, f);
+	MeshSet<3> *poly = new MeshSet<3> (vertices, numfaces, f);
 
 	uint i;
-	carve::mesh::MeshSet<3>::face_iter face_iter = poly->faceBegin();
+	MeshSet<3>::face_iter face_iter = poly->faceBegin();
 	for (i = 0; face_iter != poly->faceEnd(); ++face_iter, ++i) {
-		carve::mesh::MeshSet<3>::face_t *face = *face_iter;
+		MeshSet<3>::face_t *face = *face_iter;
 		oface_num.setAttribute(face, forig[i]);
 	}
 
@@ -139,8 +301,8 @@ static carve::mesh::MeshSet<3> *Carve_addMesh(CSG_FaceIteratorDescriptor& face_i
 }
 
 // check whether two faces share an edge, and if so merge them
-static uint quadMerge(std::map<carve::mesh::MeshSet<3>::vertex_t*, uint> *vertexToIndex_map,
+static uint quadMerge(std::map<MeshSet<3>::vertex_t*, uint> *vertexToIndex_map,
-                      carve::mesh::MeshSet<3>::face_t *f1, carve::mesh::MeshSet<3>::face_t *f2,
+                      MeshSet<3>::face_t *f1, MeshSet<3>::face_t *f2,
                       uint v, uint quad[4])
 {
 	uint current, n1, p1, n2, p2;
@@ -188,23 +350,23 @@ static uint quadMerge(std::map<carve::mesh::MeshSet<3>::vertex_t*, uint> *vertex
 	return 0;
 }
 
-static BSP_CSGMesh* Carve_exportMesh(carve::mesh::MeshSet<3>* &poly, carve::interpolate::FaceAttr<uint> &oface_num,
+static BSP_CSGMesh *Carve_exportMesh(MeshSet<3>* &poly, carve::interpolate::FaceAttr<uint> &oface_num,
                                      uint num_origfaces)
 {
 	uint i;
-	BSP_CSGMesh* outputMesh = BSP_CSGMesh::New();
+	BSP_CSGMesh *outputMesh = BSP_CSGMesh::New();
 
 	if (outputMesh == NULL)
 		return NULL;
 
-	std::vector<BSP_MVertex>* vertices = new std::vector<BSP_MVertex>;
+	std::vector<BSP_MVertex> *vertices = new std::vector<BSP_MVertex>;
 
 	outputMesh->SetVertices(vertices);
 
-	std::map<carve::mesh::MeshSet<3>::vertex_t*, uint> vertexToIndex_map;
+	std::map<MeshSet<3>::vertex_t*, uint> vertexToIndex_map;
-	std::vector<carve::mesh::MeshSet<3>::vertex_t>::iterator it = poly->vertex_storage.begin();
+	std::vector<MeshSet<3>::vertex_t>::iterator it = poly->vertex_storage.begin();
 	for (i = 0; it != poly->vertex_storage.end(); ++i, ++it) {
-		carve::mesh::MeshSet<3>::vertex_t *vertex = &(*it);
+		MeshSet<3>::vertex_t *vertex = &(*it);
 		vertexToIndex_map[vertex] = i;
 	}
 
@@ -217,13 +379,13 @@ static BSP_CSGMesh* Carve_exportMesh(carve::mesh::MeshSet<3>* &poly, carve::inte
 	}
 
 	// build vectors of faces for each original face and each vertex 
-	std::vector< std::vector<uint> > vi(poly->vertex_storage.size());
+	std::vector<std::vector<uint> > vi(poly->vertex_storage.size());
-	std::vector< std::vector<uint> > ofaces(num_origfaces);
+	std::vector<std::vector<uint> > ofaces(num_origfaces);
-	carve::mesh::MeshSet<3>::face_iter face_iter = poly->faceBegin();
+	MeshSet<3>::face_iter face_iter = poly->faceBegin();
 	for (i = 0; face_iter != poly->faceEnd(); ++face_iter, ++i) {
-		carve::mesh::MeshSet<3>::face_t *f = *face_iter;
+		MeshSet<3>::face_t *f = *face_iter;
 		ofaces[oface_num.getAttribute(f)].push_back(i);
-		carve::mesh::MeshSet<3>::face_t::edge_iter_t edge_iter = f->begin();
+		MeshSet<3>::face_t::edge_iter_t edge_iter = f->begin();
 		for (; edge_iter != f->end(); ++edge_iter) {
 			int index = vertexToIndex_map[edge_iter->vert];
 			vi[index].push_back(i);
@@ -242,7 +404,7 @@ static BSP_CSGMesh* Carve_exportMesh(carve::mesh::MeshSet<3>* &poly, carve::inte
 			uint findex = fl.back();
 			fl.pop_back();
 
-			carve::mesh::MeshSet<3>::face_t *f = *(poly->faceBegin() + findex);
+			MeshSet<3>::face_t *f = *(poly->faceBegin() + findex);
 
 			// add all information except vertices to the output mesh
 			outputMesh->FaceSet().push_back(BSP_MFace());
@@ -256,7 +418,7 @@ static BSP_CSGMesh* Carve_exportMesh(carve::mesh::MeshSet<3>* &poly, carve::inte
 			// the original face
 			uint result = 0;
 
-			carve::mesh::MeshSet<3>::face_t::edge_iter_t edge_iter = f->begin();
+			MeshSet<3>::face_t::edge_iter_t edge_iter = f->begin();
 			for (; edge_iter != f->end(); ++edge_iter) {
 				int v = vertexToIndex_map[edge_iter->vert];
 				for (uint pos2=0; !result && pos2 < vi[v].size();pos2++) {
@@ -266,7 +428,7 @@ static BSP_CSGMesh* Carve_exportMesh(carve::mesh::MeshSet<3>* &poly, carve::inte
 					if (findex == otherf)
 						continue;
 
-					carve::mesh::MeshSet<3>::face_t *f2 = *(poly->faceBegin() + otherf);
+					MeshSet<3>::face_t *f2 = *(poly->faceBegin() + otherf);
 
 					// if other face doesn't have the same original face,
 					// ignore it also
@@ -302,7 +464,7 @@ static BSP_CSGMesh* Carve_exportMesh(carve::mesh::MeshSet<3>* &poly, carve::inte
 				outFace.m_verts.push_back(quadverts[2]);
 				outFace.m_verts.push_back(quadverts[3]);
 			} else {
-				carve::mesh::MeshSet<3>::face_t::edge_iter_t edge_iter = f->begin();
+				MeshSet<3>::face_t::edge_iter_t edge_iter = f->begin();
 				for (; edge_iter != f->end(); ++edge_iter) {
 					//int index = ofacevert_num.getAttribute(f, edge_iter.idx());
 					int index = vertexToIndex_map[edge_iter->vert];
@@ -337,7 +499,7 @@ BoolOpState BOP_performBooleanOperation(BoolOpType                    opType,
                                         CSG_VertexIteratorDescriptor  obBVertices)
 {
 	carve::csg::CSG::OP op;
-	carve::mesh::MeshSet<3> *left, *right, *output;
+	MeshSet<3> *left, *right, *output = NULL;
 	carve::csg::CSG csg;
 	carve::geom3d::Vector min, max;
 	carve::interpolate::FaceAttr<uint> oface_num;
@@ -360,6 +522,19 @@ BoolOpState BOP_performBooleanOperation(BoolOpType                    opType,
 	left = Carve_addMesh(obAFaces, obAVertices, oface_num, num_origfaces );
 	right = Carve_addMesh(obBFaces, obBVertices, oface_num, num_origfaces );
 
+	Carve_prepareOperands(&left, &right, oface_num);
+
+	if(left->meshes.size() == 0 || right->meshes.size()==0) {
+		// normally sohuldn't happen (zero-faces objects are handled by modifier itself), but
+		// unioning intersecting meshes which doesn't have consistent normals might lead to
+		// empty result which wouldn't work here
+
+		delete left;
+		delete right;
+
+		return BOP_ERROR;
+	}
+
 	min.x = max.x = left->vertex_storage[0].v.x;
 	min.y = max.y = left->vertex_storage[0].v.y;
 	min.z = max.z = left->vertex_storage[0].v.z;
@@ -390,13 +565,29 @@ BoolOpState BOP_performBooleanOperation(BoolOpType                    opType,
 	csg.hooks.registerHook(new carve::csg::CarveTriangulator, carve::csg::CSG::Hooks::PROCESS_OUTPUT_FACE_BIT);
 
 	oface_num.installHooks(csg);
-	output = csg.compute( left, right, op, NULL, carve::csg::CSG::CLASSIFY_EDGE);
+
+	try {
+		output = csg.compute(left, right, op, NULL, carve::csg::CSG::CLASSIFY_EDGE);
+	}
+	catch(carve::exception e) {
+		std::cerr << "CSG failed, exception " << e.str() << std::endl;
+	}
+	catch(...) {
+		delete left;
+		delete right;
+
+		throw "Unknown error in Carve library";
+	}
+
 	delete left;
 	delete right;
 
+	if(!output)
+		return BOP_ERROR;
+
 	output->transform(rev_r);
 
-	*outputMesh = Carve_exportMesh( output, oface_num, num_origfaces);
+	*outputMesh = Carve_exportMesh(output, oface_num, num_origfaces);
 	delete output;
 
 	return BOP_OK;

intern/cycles/app/cycles_xml.cpp

@@ -462,6 +462,9 @@ static void xml_read_shader_graph(const XMLReadState& state, Shader *shader, pug
 		else if(string_iequals(node.name(), "gamma")) {
 			snode = new GammaNode();
 		}
+		else if(string_iequals(node.name(), "brightness")) {
+			snode = new BrightContrastNode();
+		}
 		else if(string_iequals(node.name(), "combine_rgb")) {
 			snode = new CombineRGBNode();
 		}

intern/cycles/blender/addon/ui.py

@@ -183,6 +183,38 @@ class CyclesRender_PT_layers(CyclesButtonsPanel, Panel):
 
         layout.separator()
 
+        split = layout.split()
+
+        col = split.column()
+        col.label(text="Passes:")
+        col.prop(rl, "use_pass_combined")
+        col.prop(rl, "use_pass_z")
+        col.prop(rl, "use_pass_normal")
+        col.prop(rl, "use_pass_object_index")
+        col.prop(rl, "use_pass_material_index")
+        col.prop(rl, "use_pass_emit")
+        col.prop(rl, "use_pass_environment")
+
+        col = split.column()
+        col.label()
+        col.label(text="Diffuse:")
+        row = col.row(align=True)
+        row.prop(rl, "use_pass_diffuse_direct", text="Direct", toggle=True)
+        row.prop(rl, "use_pass_diffuse_indirect", text="Indirect", toggle=True)
+        row.prop(rl, "use_pass_diffuse_color", text="Color", toggle=True)
+        col.label(text="Glossy:")
+        row = col.row(align=True)
+        row.prop(rl, "use_pass_glossy_direct", text="Direct", toggle=True)
+        row.prop(rl, "use_pass_glossy_indirect", text="Indirect", toggle=True)
+        row.prop(rl, "use_pass_glossy_color", text="Color", toggle=True)
+        col.label(text="Transmission:")
+        row = col.row(align=True)
+        row.prop(rl, "use_pass_transmission_direct", text="Direct", toggle=True)
+        row.prop(rl, "use_pass_transmission_indirect", text="Indirect", toggle=True)
+        row.prop(rl, "use_pass_transmission_color", text="Color", toggle=True)
+
+        layout.separator()
+
         rl = rd.layers[0]
         layout.prop(rl, "material_override", text="Material")
 
@@ -241,7 +273,7 @@ class CyclesCamera_PT_dof(CyclesButtonsPanel, Panel):
 
 
 class Cycles_PT_context_material(CyclesButtonsPanel, Panel):
-    bl_label = "Surface"
+    bl_label = ""
     bl_context = "material"
     bl_options = {'HIDE_HEADER'}
 
@@ -337,16 +369,13 @@ class CyclesObject_PT_ray_visibility(CyclesButtonsPanel, Panel):
         ob = context.object
         visibility = ob.cycles_visibility
 
-        split = layout.split()
+        flow = layout.column_flow()
-
+        
-        col = split.column()
+        flow.prop(visibility, "camera")
-        col.prop(visibility, "camera")
+        flow.prop(visibility, "diffuse")
-        col.prop(visibility, "diffuse")
+        flow.prop(visibility, "glossy")
-        col.prop(visibility, "glossy")
+        flow.prop(visibility, "transmission")
-
+        flow.prop(visibility, "shadow")
-        col = split.column()
-        col.prop(visibility, "transmission")
-        col.prop(visibility, "shadow")
 
 
 def find_node(material, nodetype):
@@ -473,17 +502,13 @@ class CyclesWorld_PT_settings(CyclesButtonsPanel, Panel):
         world = context.world
         cworld = world.cycles
 
-        split = layout.split()
+        col = layout.column()
-        col = split.column()
 
         col.prop(cworld, "sample_as_light")
         row = col.row()
         row.active = cworld.sample_as_light
         row.prop(cworld, "sample_map_resolution")
 
-        col = split.column()
-        col.label()
-
 
 class CyclesWorld_PT_volume(CyclesButtonsPanel, Panel):
     bl_label = "Volume"
@@ -617,14 +642,9 @@ class CyclesTexture_PT_context(CyclesButtonsPanel, Panel):
                 col.template_ID(user, "texture", new="texture.new")
 
             if tex:
-                row = split.row()
+                split = layout.split(percentage=0.2)
-                row.prop(tex, "use_nodes", icon="NODETREE", text="")
+                split.label(text="Type:")
-                row.label()
+                split.prop(tex, "type", text="")
-
-                if not tex.use_nodes:
-                    split = layout.split(percentage=0.2)
-                    split.label(text="Type:")
-                    split.prop(tex, "type", text="")
 
 
 class CyclesTexture_PT_nodes(CyclesButtonsPanel, Panel):

intern/cycles/blender/blender_object.cpp

@@ -159,24 +159,26 @@ void BlenderSync::sync_background_light()
 {
 	BL::World b_world = b_scene.world();
 
-	PointerRNA cworld = RNA_pointer_get(&b_world.ptr, "cycles");
+	if(b_world) {
-	bool sample_as_light = get_boolean(cworld, "sample_as_light");
+		PointerRNA cworld = RNA_pointer_get(&b_world.ptr, "cycles");
+		bool sample_as_light = get_boolean(cworld, "sample_as_light");
 
-	if(sample_as_light) {
+		if(sample_as_light) {
-		/* test if we need to sync */
+			/* test if we need to sync */
-		Light *light;
+			Light *light;
-		ObjectKey key(b_world, 0, b_world);
+			ObjectKey key(b_world, 0, b_world);
 
-		if(light_map.sync(&light, b_world, b_world, key) ||
+			if(light_map.sync(&light, b_world, b_world, key) ||
-		   world_recalc ||
+			   world_recalc ||
-		   b_world.ptr.data != world_map)
+			   b_world.ptr.data != world_map)
-		{
+			{
-			light->type = LIGHT_BACKGROUND;
+				light->type = LIGHT_BACKGROUND;
-			light->map_resolution  = get_int(cworld, "sample_map_resolution");
+				light->map_resolution  = get_int(cworld, "sample_map_resolution");
-			light->shader = scene->default_background;
+				light->shader = scene->default_background;
 
-			light->tag_update(scene);
+				light->tag_update(scene);
-			light_map.set_recalc(b_world);
+				light_map.set_recalc(b_world);
+			}
 		}
 	}
 
@@ -212,6 +214,7 @@ void BlenderSync::sync_object(BL::Object b_parent, int b_index, BL::Object b_ob,
 	/* object sync */
 	if(object_updated || (object->mesh && object->mesh->need_update)) {
 		object->name = b_ob.name().c_str();
+		object->pass_id = b_ob.pass_index();
 		object->tfm = tfm;
 
 		/* visibility flags for both parent */

intern/cycles/blender/blender_session.cpp

@@ -116,9 +116,68 @@ void BlenderSession::free_session()
 	delete session;
 }
 
+static PassType get_pass_type(BL::RenderPass b_pass)
+{
+	switch(b_pass.type()) {
+		case BL::RenderPass::type_COMBINED:
+			return PASS_COMBINED;
+
+		case BL::RenderPass::type_Z:
+			return PASS_DEPTH;
+		case BL::RenderPass::type_NORMAL:
+			return PASS_NORMAL;
+		case BL::RenderPass::type_OBJECT_INDEX:
+			return PASS_OBJECT_ID;
+		case BL::RenderPass::type_UV:
+			return PASS_UV;
+		case BL::RenderPass::type_MATERIAL_INDEX:
+			return PASS_MATERIAL_ID;
+
+		case BL::RenderPass::type_DIFFUSE_DIRECT:
+			return PASS_DIFFUSE_DIRECT;
+		case BL::RenderPass::type_GLOSSY_DIRECT:
+			return PASS_GLOSSY_DIRECT;
+		case BL::RenderPass::type_TRANSMISSION_DIRECT:
+			return PASS_TRANSMISSION_DIRECT;
+
+		case BL::RenderPass::type_DIFFUSE_INDIRECT:
+			return PASS_DIFFUSE_INDIRECT;
+		case BL::RenderPass::type_GLOSSY_INDIRECT:
+			return PASS_GLOSSY_INDIRECT;
+		case BL::RenderPass::type_TRANSMISSION_INDIRECT:
+			return PASS_TRANSMISSION_INDIRECT;
+
+		case BL::RenderPass::type_DIFFUSE_COLOR:
+			return PASS_DIFFUSE_COLOR;
+		case BL::RenderPass::type_GLOSSY_COLOR:
+			return PASS_GLOSSY_COLOR;
+		case BL::RenderPass::type_TRANSMISSION_COLOR:
+			return PASS_TRANSMISSION_COLOR;
+
+		case BL::RenderPass::type_EMIT:
+			return PASS_EMISSION;
+		case BL::RenderPass::type_ENVIRONMENT:
+			return PASS_BACKGROUND;
+
+		case BL::RenderPass::type_DIFFUSE:
+		case BL::RenderPass::type_SHADOW:
+		case BL::RenderPass::type_AO:
+		case BL::RenderPass::type_COLOR:
+		case BL::RenderPass::type_REFRACTION:
+		case BL::RenderPass::type_SPECULAR:
+		case BL::RenderPass::type_REFLECTION:
+		case BL::RenderPass::type_VECTOR:
+		case BL::RenderPass::type_MIST:
+			return PASS_NONE;
+	}
+	
+	return PASS_NONE;
+}
+
 void BlenderSession::render()
 {
 	/* get buffer parameters */
+	SessionParams session_params = BlenderSync::get_session_params(b_userpref, b_scene, background);
 	BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
 	int w = buffer_params.width, h = buffer_params.height;
 
@@ -143,6 +202,29 @@ void BlenderSession::render()
 		/* set layer */
 		b_rlay = *b_iter;
 
+		/* add passes */
+		vector<Pass> passes;
+		Pass::add(PASS_COMBINED, passes);
+
+		if(session_params.device.advanced_shading) {
+			BL::RenderLayer::passes_iterator b_pass_iter;
+			
+			for(b_rlay.passes.begin(b_pass_iter); b_pass_iter != b_rlay.passes.end(); ++b_pass_iter) {
+				BL::RenderPass b_pass(*b_pass_iter);
+				PassType pass_type = get_pass_type(b_pass);
+
+				if(pass_type != PASS_NONE)
+					Pass::add(pass_type, passes);
+			}
+		}
+
+		buffer_params.passes = passes;
+		scene->film->passes = passes;
+		scene->film->tag_update(scene);
+
+		/* update session */
+		session->reset(buffer_params, session_params.samples);
+
 		/* update scene */
 		sync->sync_data(b_v3d, active);
 
@@ -165,22 +247,41 @@ void BlenderSession::write_render_result()
 {
 	/* get state */
 	RenderBuffers *buffers = session->buffers;
+
+	/* copy data from device */
+	if(!buffers->copy_from_device())
+		return;
+
+	BufferParams& params = buffers->params;
 	float exposure = scene->film->exposure;
 	double total_time, sample_time;
 	int sample;
+
 	session->progress.get_sample(sample, total_time, sample_time);
 
-	/* get pixels */
+	vector<float> pixels(params.width*params.height*4);
-	float4 *pixels = buffers->copy_from_device(exposure, sample);
 
-	if(!pixels)
+	/* copy each pass */
-		return;
+	BL::RenderLayer::passes_iterator b_iter;
+	
+	for(b_rlay.passes.begin(b_iter); b_iter != b_rlay.passes.end(); ++b_iter) {
+		BL::RenderPass b_pass(*b_iter);
 
-	/* write pixels */
+		/* find matching pass type */
-	rna_RenderLayer_rect_set(&b_rlay.ptr, (float*)pixels);
+		PassType pass_type = get_pass_type(b_pass);
+		int components = b_pass.channels();
+
+		/* copy pixels */
+		if(buffers->get_pass(pass_type, exposure, sample, components, &pixels[0]))
+			rna_RenderPass_rect_set(&b_pass.ptr, &pixels[0]);
+	}
+
+	/* copy combined pass */
+	if(buffers->get_pass(PASS_COMBINED, exposure, sample, 4, &pixels[0]))
+		rna_RenderLayer_rect_set(&b_rlay.ptr, &pixels[0]);
+
+	/* tag result as updated */
 	RE_engine_update_result((RenderEngine*)b_engine.ptr.data, (RenderResult*)b_rr.ptr.data);
-
-	delete [] pixels;
 }
 
 void BlenderSession::synchronize()

intern/cycles/blender/blender_shader.cpp

@@ -142,6 +142,10 @@ static ShaderNode *add_node(BL::BlendData b_data, ShaderGraph *graph, BL::Shader
 			node = new GammaNode();
 			break;
 		}
+		case BL::ShaderNode::type_BRIGHTCONTRAST: {
+			node = new BrightContrastNode();
+			break;
+		}
 		case BL::ShaderNode::type_MIX_RGB: {
 			BL::ShaderNodeMixRGB b_mix_node(b_node);
 			MixNode *mix = new MixNode();
@@ -632,6 +636,7 @@ void BlenderSync::sync_materials()
 			ShaderGraph *graph = new ShaderGraph();
 
 			shader->name = b_mat->name().c_str();
+			shader->pass_id = b_mat->pass_index();
 
 			/* create nodes */
 			if(b_mat->use_nodes() && b_mat->node_tree()) {

intern/cycles/device/device.cpp

@@ -183,6 +183,9 @@ Device *Device::create(DeviceInfo& info, bool background, int threads)
 			return NULL;
 	}
 
+	if(device)
+		device->info = info;
+
 	return device;
 }
 

intern/cycles/device/device.h

@@ -51,6 +51,7 @@ public:
 	string id;
 	int num;
 	bool display_device;
+	bool advanced_shading;
 	vector<DeviceInfo> multi_devices;
 
 	DeviceInfo()
@@ -59,6 +60,7 @@ public:
 		id = "CPU";
 		num = 0;
 		display_device = false;
+		advanced_shading = true;
 	}
 };
 
@@ -101,10 +103,8 @@ protected:
 public:
 	virtual ~Device() {}
 
-	virtual bool support_full_kernel() = 0;
-
 	/* info */
-	virtual string description() = 0;
+	DeviceInfo info;
 	virtual const string& error_message() { return error_msg; }
 
 	/* regular memory */

intern/cycles/device/device_cpu.cpp

@@ -72,16 +72,11 @@ public:
 		kernel_globals_free(kg);
 	}
 
-	bool support_full_kernel()
+	bool support_advanced_shading()
 	{
 		return true;
 	}
 
-	string description()
-	{
-		return system_cpu_brand_string();
-	}
-
 	void mem_alloc(device_memory& mem, MemoryType type)
 	{
 		mem.device_pointer = mem.data_pointer;
@@ -162,7 +157,7 @@ public:
 		if(system_cpu_support_optimized()) {
 			for(int y = task.y; y < task.y + task.h; y++) {
 				for(int x = task.x; x < task.x + task.w; x++)
-					kernel_cpu_optimized_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state,
+					kernel_cpu_optimized_path_trace(kg, (float*)task.buffer, (unsigned int*)task.rng_state,
 						task.sample, x, y, task.offset, task.stride);
 
 				if(tasks.worker_cancel())
@@ -174,7 +169,7 @@ public:
 		{
 			for(int y = task.y; y < task.y + task.h; y++) {
 				for(int x = task.x; x < task.x + task.w; x++)
-					kernel_cpu_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state,
+					kernel_cpu_path_trace(kg, (float*)task.buffer, (unsigned int*)task.rng_state,
 						task.sample, x, y, task.offset, task.stride);
 
 				if(tasks.worker_cancel())
@@ -194,7 +189,7 @@ public:
 		if(system_cpu_support_optimized()) {
 			for(int y = task.y; y < task.y + task.h; y++)
 				for(int x = task.x; x < task.x + task.w; x++)
-					kernel_cpu_optimized_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer,
+					kernel_cpu_optimized_tonemap(kg, (uchar4*)task.rgba, (float*)task.buffer,
 						task.sample, task.resolution, x, y, task.offset, task.stride);
 		}
 		else
@@ -202,7 +197,7 @@ public:
 		{
 			for(int y = task.y; y < task.y + task.h; y++)
 				for(int x = task.x; x < task.x + task.w; x++)
-					kernel_cpu_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer,
+					kernel_cpu_tonemap(kg, (uchar4*)task.rgba, (float*)task.buffer,
 						task.sample, task.resolution, x, y, task.offset, task.stride);
 		}
 	}
@@ -271,6 +266,7 @@ void device_cpu_info(vector<DeviceInfo>& devices)
 	info.description = system_cpu_brand_string();
 	info.id = "CPU";
 	info.num = 0;
+	info.advanced_shading = true;
 
 	devices.insert(devices.begin(), info);
 }

intern/cycles/device/device_cuda.cpp

@@ -106,11 +106,6 @@ public:
 		}
 	}
 
-	static int cuda_align_up(int& offset, int alignment)
-	{
-		return (offset + alignment - 1) & ~(alignment - 1);
-	}
-
 #ifdef NDEBUG
 #define cuda_abort()
 #else
@@ -194,26 +189,6 @@ public:
 		cuda_assert(cuCtxDetach(cuContext))
 	}
 
-	bool support_full_kernel()
-	{
-		int major, minor;
-		cuDeviceComputeCapability(&major, &minor, cuDevId);
-
-		return (major >= 2);
-	}
-
-	string description()
-	{
-		/* print device information */
-		char deviceName[256];
-
-		cuda_push_context();
-		cuDeviceGetName(deviceName, 256, cuDevId);
-		cuda_pop_context();
-
-		return string("CUDA ") + deviceName;
-	}
-
 	bool support_device(bool experimental)
 	{
 		if(!experimental) {
@@ -505,7 +480,7 @@ public:
 		offset += sizeof(d_rng_state);
 
 		int sample = task.sample;
-		offset = cuda_align_up(offset, __alignof(sample));
+		offset = align_up(offset, __alignof(sample));
 
 		cuda_assert(cuParamSeti(cuPathTrace, offset, task.sample))
 		offset += sizeof(task.sample);
@@ -569,7 +544,7 @@ public:
 		offset += sizeof(d_buffer);
 
 		int sample = task.sample;
-		offset = cuda_align_up(offset, __alignof(sample));
+		offset = align_up(offset, __alignof(sample));
 
 		cuda_assert(cuParamSeti(cuFilmConvert, offset, task.sample))
 		offset += sizeof(task.sample);
@@ -638,7 +613,7 @@ public:
 		offset += sizeof(d_offset);
 
 		int shader_eval_type = task.shader_eval_type;
-		offset = cuda_align_up(offset, __alignof(shader_eval_type));
+		offset = align_up(offset, __alignof(shader_eval_type));
 
 		cuda_assert(cuParamSeti(cuDisplace, offset, task.shader_eval_type))
 		offset += sizeof(task.shader_eval_type);
@@ -881,6 +856,10 @@ void device_cuda_info(vector<DeviceInfo>& devices)
 		info.id = string_printf("CUDA_%d", num);
 		info.num = num;
 
+		int major, minor;
+		cuDeviceComputeCapability(&major, &minor, num);
+		info.advanced_shading = (major >= 2);
+
 		/* if device has a kernel timeout, assume it is used for display */
 		if(cuDeviceGetAttribute(&attr, CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT, num) == CUDA_SUCCESS && attr == 1) {
 			info.display_device = true;

intern/cycles/device/device_multi.cpp

@@ -76,16 +76,6 @@ public:
 			delete sub.device;
 	}
 
-	bool support_full_kernel()
-	{
-		foreach(SubDevice& sub, devices) {
-			if(!sub.device->support_full_kernel())
-				return false;
-		}
-
-		return true;
-	}
-
 	const string& error_message()
 	{
 		foreach(SubDevice& sub, devices) {
@@ -99,38 +89,6 @@ public:
 		return error_msg;
 	}
 
-	string description()
-	{
-		/* create map to find duplicate descriptions */
-		map<string, int> dupli_map;
-		map<string, int>::iterator dt;
-
-		foreach(SubDevice& sub, devices) {
-			string key = sub.device->description();
-
-			if(dupli_map.find(key) == dupli_map.end())
-				dupli_map[key] = 1;
-			else
-				dupli_map[key]++;
-		}
-
-		/* generate string */
-		stringstream desc;
-		bool first = true;
-
-		for(dt = dupli_map.begin(); dt != dupli_map.end(); dt++) {
-			if(!first) desc << ", ";
-			first = false;
-
-			if(dt->second > 1)
-				desc << dt->second << "x " << dt->first;
-			else
-				desc << dt->first;
-		}
-
-		return desc.str();
-	}
-
 	bool load_kernels(bool experimental)
 	{
 		foreach(SubDevice& sub, devices)
@@ -344,6 +302,8 @@ static void device_multi_add(vector<DeviceInfo>& devices, DeviceType type, bool
 	map<string, int>::iterator dt;
 	int num_added = 0, num_display = 0;
 
+	info.advanced_shading = true;
+
 	foreach(DeviceInfo& subinfo, devices) {
 		if(subinfo.type == type) {
 			if(subinfo.display_device) {
@@ -363,6 +323,8 @@ static void device_multi_add(vector<DeviceInfo>& devices, DeviceType type, bool
 			info.multi_devices.push_back(subinfo);
 			if(subinfo.display_device)
 				info.display_device = true;
+			if(!subinfo.advanced_shading)
+				info.advanced_shading = false;
 			num_added++;
 		}
 	}

intern/cycles/device/device_network.cpp

@@ -57,24 +57,6 @@ public:
 	{
 	}
 
-	bool support_full_kernel()
-	{
-		return false;
-	}
-
-	string description()
-	{
-		RPCSend snd(socket, "description");
-		snd.write();
-
-		RPCReceive rcv(socket);
-		string desc_string;
-
-		*rcv.archive & desc_string;
-
-		return desc_string + " (remote)";
-	}
-
 	void mem_alloc(device_memory& mem, MemoryType type)
 	{
 #if 0

intern/cycles/device/device_opencl.cpp

@@ -453,20 +453,6 @@ public:
 			clReleaseContext(cxContext);
 	}
 
-	bool support_full_kernel()
-	{
-		return false;
-	}
-
-	string description()
-	{
-		char name[1024];
-
-		clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(name), &name, NULL);
-
-		return string("OpenCL ") + name;
-	}
-
 	void mem_alloc(device_memory& mem, MemoryType type)
 	{
 		size_t size = mem.memory_size();
@@ -750,6 +736,7 @@ void device_opencl_info(vector<DeviceInfo>& devices)
 		info.num = num;
 		/* we don't know if it's used for display, but assume it is */
 		info.display_device = true;
+		info.advanced_shading = false;
 
 		devices.push_back(info);
 	}

intern/cycles/kernel/CMakeLists.txt

@@ -15,6 +15,7 @@ set(SRC
 
 set(SRC_HEADERS
 	kernel.h
+	kernel_accumulate.h
 	kernel_bvh.h
 	kernel_camera.h
 	kernel_compat_cpu.h
@@ -30,6 +31,7 @@ set(SRC_HEADERS
 	kernel_mbvh.h
 	kernel_montecarlo.h
 	kernel_object.h
+	kernel_passes.h
 	kernel_path.h
 	kernel_qbvh.h
 	kernel_random.h
@@ -61,6 +63,7 @@ set(SRC_SVM_HEADERS
 	svm/svm_displace.h
 	svm/svm_fresnel.h
 	svm/svm_gamma.h
+	svm/svm_brightness.h
 	svm/svm_geometry.h
 	svm/svm_gradient.h
 	svm/svm_hsv.h

intern/cycles/kernel/kernel.cl

@@ -28,7 +28,7 @@
 
 __kernel void kernel_ocl_path_trace(
 	__constant KernelData *data,
-	__global float4 *buffer,
+	__global float *buffer,
 	__global uint *rng_state,
 
 #define KERNEL_TEX(type, ttype, name) \
@@ -56,7 +56,7 @@ __kernel void kernel_ocl_path_trace(
 __kernel void kernel_ocl_tonemap(
 	__constant KernelData *data,
 	__global uchar4 *rgba,
-	__global float4 *buffer,
+	__global float *buffer,
 
 #define KERNEL_TEX(type, ttype, name) \
 	__global type *name,

intern/cycles/kernel/kernel.cpp

@@ -204,14 +204,14 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t
 
 /* Path Tracing */
 
-void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
+void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
 {
 	kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */
 
-void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
+void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride)
 {
 	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }

intern/cycles/kernel/kernel.cu

@@ -26,7 +26,7 @@
 #include "kernel_path.h"
 #include "kernel_displace.h"
 
-extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
+extern "C" __global__ void kernel_cuda_path_trace(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
 {
 	int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
 	int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
@@ -35,7 +35,7 @@ extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_stat
 		kernel_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride);
 }
 
-extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float4 *buffer, int sample, int resolution, int sx, int sy, int sw, int sh, int offset, int stride)
+extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float *buffer, int sample, int resolution, int sx, int sy, int sw, int sh, int offset, int stride)
 {
 	int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
 	int y = sy + blockDim.y*blockIdx.y + threadIdx.y;

intern/cycles/kernel/kernel.h

@@ -36,17 +36,17 @@ bool kernel_osl_use(KernelGlobals *kg);
 void kernel_const_copy(KernelGlobals *kg, const char *name, void *host, size_t size);
 void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t width, size_t height);
 
-void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state,
+void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state,
 	int sample, int x, int y, int offset, int stride);
-void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer,
+void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer,
 	int sample, int resolution, int x, int y, int offset, int stride);
 void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float4 *output,
 	int type, int i);
 
 #ifdef WITH_OPTIMIZED_KERNEL
-void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state,
+void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state,
 	int sample, int x, int y, int offset, int stride);
-void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer,
+void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer,
 	int sample, int resolution, int x, int y, int offset, int stride);
 void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float4 *output,
 	int type, int i);

intern/cycles/kernel/kernel_accumulate.h

@@ -0,0 +1,283 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ */
+
+CCL_NAMESPACE_BEGIN
+
+/* BSDF Eval
+ *
+ * BSDF evaluation result, split per BSDF type. This is used to accumulate
+ * render passes separately. */
+
+__device_inline void bsdf_eval_init(BsdfEval *eval, ClosureType type, float3 value, int use_light_pass)
+{
+#ifdef __PASSES__
+	eval->use_light_pass = use_light_pass;
+
+	if(eval->use_light_pass) {
+		eval->diffuse = make_float3(0.0f, 0.0f, 0.0f);
+		eval->glossy = make_float3(0.0f, 0.0f, 0.0f);
+		eval->transmission = make_float3(0.0f, 0.0f, 0.0f);
+		eval->transparent = make_float3(0.0f, 0.0f, 0.0f);
+
+		if(type == CLOSURE_BSDF_TRANSPARENT_ID)
+			eval->transparent = value;
+		else if(CLOSURE_IS_BSDF_DIFFUSE(type))
+			eval->diffuse = value;
+		else if(CLOSURE_IS_BSDF_GLOSSY(type))
+			eval->glossy = value;
+		else
+			eval->transmission = value;
+	}
+	else
+		eval->diffuse = value;
+#else
+	*eval = value;
+#endif
+}
+
+__device_inline void bsdf_eval_accum(BsdfEval *eval, ClosureType type, float3 value)
+{
+#ifdef __PASSES__
+	if(eval->use_light_pass) {
+		if(CLOSURE_IS_BSDF_DIFFUSE(type))
+			eval->diffuse += value;
+		else if(CLOSURE_IS_BSDF_GLOSSY(type))
+			eval->glossy += value;
+		else
+			eval->transmission += value;
+
+		/* skipping transparent, this function is used by for eval(), will be zero then */
+	}
+	else
+		eval->diffuse += value;
+#else
+	*eval += value;
+#endif
+}
+
+__device_inline bool bsdf_eval_is_zero(BsdfEval *eval)
+{
+#ifdef __PASSES__
+	if(eval->use_light_pass) {
+		return is_zero(eval->diffuse)
+			&& is_zero(eval->glossy)
+			&& is_zero(eval->transmission)
+			&& is_zero(eval->transparent);
+	}
+	else
+		return is_zero(eval->diffuse);
+#else
+	return is_zero(*eval);
+#endif
+}
+
+__device_inline void bsdf_eval_mul(BsdfEval *eval, float3 value)
+{
+#ifdef __PASSES__
+	if(eval->use_light_pass) {
+		eval->diffuse *= value;
+		eval->glossy *= value;
+		eval->transmission *= value;
+
+		/* skipping transparent, this function is used by for eval(), will be zero then */
+	}
+	else
+		eval->diffuse *= value;
+#else
+	*eval *= value;
+#endif
+}
+
+/* Path Radiance
+ *
+ * We accumulate different render passes separately. After summing at the end
+ * to get the combined result, it should be identical. We definte directly
+ * visible as the first non-transparent hit, while indirectly visible are the
+ * bounces after that. */
+
+__device_inline void path_radiance_init(PathRadiance *L, int use_light_pass)
+{
+	/* clear all */
+#ifdef __PASSES__
+	L->use_light_pass = use_light_pass;
+
+	if(use_light_pass) {
+		L->indirect = make_float3(0.0f, 0.0f, 0.0f);
+		L->direct_throughput = make_float3(0.0f, 0.0f, 0.0f);
+		L->direct_emission = make_float3(0.0f, 0.0f, 0.0f);
+
+		L->color_diffuse = make_float3(0.0f, 0.0f, 0.0f);
+		L->color_glossy = make_float3(0.0f, 0.0f, 0.0f);
+		L->color_transmission = make_float3(0.0f, 0.0f, 0.0f);
+
+		L->direct_diffuse = make_float3(0.0f, 0.0f, 0.0f);
+		L->direct_glossy = make_float3(0.0f, 0.0f, 0.0f);
+		L->direct_transmission = make_float3(0.0f, 0.0f, 0.0f);
+
+		L->indirect_diffuse = make_float3(0.0f, 0.0f, 0.0f);
+		L->indirect_glossy = make_float3(0.0f, 0.0f, 0.0f);
+		L->indirect_transmission = make_float3(0.0f, 0.0f, 0.0f);
+
+		L->emission = make_float3(0.0f, 0.0f, 0.0f);
+		L->background = make_float3(0.0f, 0.0f, 0.0f);
+	}
+	else
+		L->emission = make_float3(0.0f, 0.0f, 0.0f);
+#else
+	*L = make_float3(0.0f, 0.0f, 0.0f);
+#endif
+}
+
+__device_inline void path_radiance_bsdf_bounce(PathRadiance *L, float3 *throughput,
+	BsdfEval *bsdf_eval, float bsdf_pdf, int bounce, int bsdf_label)
+{
+	float inverse_pdf = 1.0f/bsdf_pdf;
+
+#ifdef __PASSES__
+	if(L->use_light_pass) {
+		if(bounce == 0) {
+			if(bsdf_label & LABEL_TRANSPARENT) {
+				/* transparent bounce before first hit */
+				*throughput *= bsdf_eval->transparent*inverse_pdf;
+			}
+			else {
+				/* first on directly visible surface */
+				float3 value = *throughput*inverse_pdf;
+
+				L->indirect_diffuse = bsdf_eval->diffuse*value;
+				L->indirect_glossy = bsdf_eval->glossy*value;
+				L->indirect_transmission = bsdf_eval->transmission*value;
+
+				*throughput = L->indirect_diffuse + L->indirect_glossy + L->indirect_transmission;
+				
+				L->direct_throughput = *throughput;
+			}
+		}
+		else {
+			/* indirectly visible through BSDF */
+			float3 sum = (bsdf_eval->diffuse + bsdf_eval->glossy + bsdf_eval->transmission + bsdf_eval->transparent)*inverse_pdf;
+			*throughput *= sum;
+		}
+	}
+	else
+		*throughput *= bsdf_eval->diffuse*inverse_pdf;
+#else
+	*throughput *= *bsdf_eval*inverse_pdf;
+#endif
+}
+
+__device_inline void path_radiance_accum_emission(PathRadiance *L, float3 throughput, float3 value, int bounce)
+{
+#ifdef __PASSES__
+	if(L->use_light_pass) {
+		if(bounce == 0)
+			L->emission += throughput*value;
+		else if(bounce == 1)
+			L->direct_emission += throughput*value;
+		else
+			L->indirect += throughput*value;
+	}
+	else
+		L->emission += throughput*value;
+#else
+	*L += throughput*value;
+#endif
+}
+
+__device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughput, BsdfEval *bsdf_eval, int bounce)
+{
+#ifdef __PASSES__
+	if(L->use_light_pass) {
+		if(bounce == 0) {
+			/* directly visible lighting */
+			L->direct_diffuse += throughput*bsdf_eval->diffuse;
+			L->direct_glossy += throughput*bsdf_eval->glossy;
+			L->direct_transmission += throughput*bsdf_eval->transmission;
+		}
+		else {
+			/* indirectly visible lighting after BSDF bounce */
+			float3 sum = bsdf_eval->diffuse + bsdf_eval->glossy + bsdf_eval->transmission;
+			L->indirect += throughput*sum;
+		}
+	}
+	else
+		L->emission += throughput*bsdf_eval->diffuse;
+#else
+	*L += throughput*(*bsdf_eval);
+#endif
+}
+
+__device_inline void path_radiance_accum_background(PathRadiance *L, float3 throughput, float3 value, int bounce)
+{
+#ifdef __PASSES__
+	if(L->use_light_pass) {
+		if(bounce == 0)
+			L->background += throughput*value;
+		else if(bounce == 1)
+			L->direct_emission += throughput*value;
+		else
+			L->indirect += throughput*value;
+	}
+	else
+		L->emission += throughput*value;
+#else
+	*L += throughput*value;
+#endif
+}
+
+__device_inline float3 safe_divide_color(float3 a, float3 b)
+{
+	float x, y, z;
+
+	x = (b.x != 0.0f)? a.x/b.x: 0.0f;
+	y = (b.y != 0.0f)? a.y/b.y: 0.0f;
+	z = (b.z != 0.0f)? a.z/b.z: 0.0f;
+
+	return make_float3(x, y, z);
+}
+
+__device_inline float3 path_radiance_sum(PathRadiance *L)
+{
+#ifdef __PASSES__
+	if(L->use_light_pass) {
+		/* this division is a bit ugly, but means we only have to keep track of
+		   only a single throughput further along the path, here we recover just
+		   the indirect parth that is not influenced by any particular BSDF type */
+		L->direct_emission = safe_divide_color(L->direct_emission, L->direct_throughput);
+		L->direct_diffuse += L->indirect_diffuse*L->direct_emission;
+		L->direct_glossy += L->indirect_glossy*L->direct_emission;
+		L->direct_transmission += L->indirect_transmission*L->direct_emission;
+
+		L->indirect = safe_divide_color(L->indirect, L->direct_throughput);
+		L->indirect_diffuse *= L->indirect;
+		L->indirect_glossy *= L->indirect;
+		L->indirect_transmission *= L->indirect;
+
+		return L->emission + L->background
+			+ L->direct_diffuse + L->direct_glossy + L->direct_transmission
+			+ L->indirect_diffuse + L->indirect_glossy + L->indirect_transmission;
+	}
+	else
+		return L->emission;
+#else
+	return *L;
+#endif
+}
+
+CCL_NAMESPACE_END
+

intern/cycles/kernel/kernel_emission.h

@@ -27,6 +27,7 @@ __device float3 direct_emissive_eval(KernelGlobals *kg, float rando,
 	ShaderData sd;
 	float3 eval;
 
+#ifdef __BACKGROUND_MIS__
 	if(ls->type == LIGHT_BACKGROUND) {
 		Ray ray;
 		ray.D = ls->D;
@@ -36,7 +37,9 @@ __device float3 direct_emissive_eval(KernelGlobals *kg, float rando,
 		shader_setup_from_background(kg, &sd, &ray);
 		eval = shader_eval_background(kg, &sd, 0);
 	}
-	else {
+	else
+#endif
+	{
 		shader_setup_from_sample(kg, &sd, ls->P, ls->Ng, I, ls->shader, ls->object, ls->prim, u, v);
 		ls->Ng = sd.Ng;
 
@@ -57,7 +60,7 @@ __device float3 direct_emissive_eval(KernelGlobals *kg, float rando,
 }
 
 __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex,
-	float randt, float rando, float randu, float randv, Ray *ray, float3 *eval)
+	float randt, float rando, float randu, float randv, Ray *ray, BsdfEval *eval)
 {
 	LightSample ls;
 
@@ -83,32 +86,33 @@ __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex,
 		return false;
 
 	/* evaluate closure */
-	*eval = direct_emissive_eval(kg, rando, &ls, randu, randv, -ls.D);
+	float3 light_eval = direct_emissive_eval(kg, rando, &ls, randu, randv, -ls.D);
 
-	if(is_zero(*eval))
+	if(is_zero(light_eval))
 		return false;
 
 	/* todo: use visbility flag to skip lights */
 
 	/* evaluate BSDF at shading point */
 	float bsdf_pdf;
-	float3 bsdf_eval = shader_bsdf_eval(kg, sd, ls.D, &bsdf_pdf);
 
-	*eval *= bsdf_eval/pdf;
+	shader_bsdf_eval(kg, sd, ls.D, eval, &bsdf_pdf);
-
-	if(is_zero(*eval))
-		return false;
 
 	if(ls.prim != ~0 || ls.type == LIGHT_BACKGROUND) {
 		/* multiple importance sampling */
 		float mis_weight = power_heuristic(pdf, bsdf_pdf);
-		*eval *= mis_weight;
+		light_eval *= mis_weight;
 	}
 	/* todo: clean up these weights */
 	else if(ls.shader & SHADER_AREA_LIGHT)
-		*eval *= 0.25f; /* area lamp */
+		light_eval *= 0.25f; /* area lamp */
 	else if(ls.t != FLT_MAX)
-		*eval *= 0.25f*M_1_PI_F; /* point lamp */
+		light_eval *= 0.25f*M_1_PI_F; /* point lamp */
+	
+	bsdf_eval_mul(eval, light_eval/pdf);
+
+	if(bsdf_eval_is_zero(eval))
+		return false;
 
 	if(ls.shader & SHADER_CAST_SHADOW) {
 		/* setup ray */
@@ -163,6 +167,7 @@ __device float3 indirect_background(KernelGlobals *kg, Ray *ray, int path_flag,
 	float3 L = shader_eval_background(kg, &sd, path_flag);
 	shader_release(kg, &sd);
 
+#ifdef __BACKGROUND_MIS__
 	/* check if background light exists or if we should skip pdf */
 	int res = kernel_data.integrator.pdf_background_res;
 
@@ -174,6 +179,7 @@ __device float3 indirect_background(KernelGlobals *kg, Ray *ray, int path_flag,
 
 		return L*mis_weight;
 	}
+#endif
 
 	return L;
 #else

intern/cycles/kernel/kernel_film.h

@@ -48,15 +48,22 @@ __device uchar4 film_float_to_byte(float4 color)
 	return result;
 }
 
-__device void kernel_film_tonemap(KernelGlobals *kg, __global uchar4 *rgba, __global float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
+__device void kernel_film_tonemap(KernelGlobals *kg,
+	__global uchar4 *rgba, __global float *buffer,
+	int sample, int resolution, int x, int y, int offset, int stride)
 {
+	/* buffer offset */
 	int index = offset + x + y*stride;
-	float4 irradiance = buffer[index];
 
+	rgba += index;
+	buffer += index*kernel_data.film.pass_stride;
+
+	/* map colors */
+	float4 irradiance = *((__global float4*)buffer);
 	float4 float_result = film_map(kg, irradiance, sample);
 	uchar4 byte_result = film_float_to_byte(float_result);
 
-	rgba[index] = byte_result;
+	*rgba = byte_result;
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/kernel_light.h

@@ -192,6 +192,7 @@ __device void regular_light_sample(KernelGlobals *kg, int point,
 		ls->D = -D;
 		ls->t = FLT_MAX;
 	}
+#ifdef __BACKGROUND_MIS__
 	else if(type == LIGHT_BACKGROUND) {
 		/* infinite area light (e.g. light dome or env light) */
 		float3 D = background_light_sample(kg, randu, randv, pdf);
@@ -201,6 +202,7 @@ __device void regular_light_sample(KernelGlobals *kg, int point,
 		ls->D = -D;
 		ls->t = FLT_MAX;
 	}
+#endif
 	else {
 		ls->P = make_float3(data0.y, data0.z, data0.w);
 

intern/cycles/kernel/kernel_object.h

@@ -64,5 +64,15 @@ __device_inline float object_surface_area(KernelGlobals *kg, int object)
 	return f.x;
 }
 
+__device_inline float object_pass_id(KernelGlobals *kg, int object)
+{
+	if(object == ~0)
+		return 0.0f;
+
+	int offset = object*OBJECT_SIZE + OBJECT_PROPERTIES;
+	float4 f = kernel_tex_fetch(__objects, offset);
+	return f.y;
+}
+
 CCL_NAMESPACE_END
 

intern/cycles/kernel/kernel_optimized.cpp

@@ -35,14 +35,14 @@ CCL_NAMESPACE_BEGIN
 
 /* Path Tracing */
 
-void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
+void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
 {
 	kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */
 
-void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
+void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride)
 {
 	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }

intern/cycles/kernel/kernel_passes.h

@@ -0,0 +1,137 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ */
+
+CCL_NAMESPACE_BEGIN
+
+__device_inline void kernel_write_pass_float(__global float *buffer, int sample, float value)
+{
+	__global float *buf = buffer;
+	*buf = (sample == 0)? value: *buf + value;
+}
+
+__device_inline void kernel_write_pass_float3(__global float *buffer, int sample, float3 value)
+{
+	__global float3 *buf = (__global float3*)buffer;
+	*buf = (sample == 0)? value: *buf + value;
+}
+
+__device_inline void kernel_write_pass_float4(__global float *buffer, int sample, float4 value)
+{
+	__global float4 *buf = (__global float4*)buffer;
+	*buf = (sample == 0)? value: *buf + value;
+}
+
+__device_inline void kernel_write_data_passes(KernelGlobals *kg, __global float *buffer, PathRadiance *L,
+	ShaderData *sd, int sample, int path_flag, float3 throughput)
+{
+#ifdef __PASSES__
+	if(!(path_flag & PATH_RAY_CAMERA))
+		return;
+
+	int flag = kernel_data.film.pass_flag;
+
+	if(!(flag & PASS_ALL))
+		return;
+	
+	/* todo: add alpha treshold */
+	if(!(path_flag & PATH_RAY_TRANSPARENT)) {
+		if(sample == 0) {
+			if(flag & PASS_DEPTH) {
+				Transform tfm = kernel_data.cam.worldtocamera;
+				float depth = len(transform(&tfm, sd->P));
+
+				kernel_write_pass_float(buffer + kernel_data.film.pass_depth, sample, depth);
+			}
+			if(flag & PASS_OBJECT_ID) {
+				float id = object_pass_id(kg, sd->object);
+				kernel_write_pass_float(buffer + kernel_data.film.pass_object_id, sample, id);
+			}
+			if(flag & PASS_MATERIAL_ID) {
+				float id = shader_pass_id(kg, sd);
+				kernel_write_pass_float(buffer + kernel_data.film.pass_material_id, sample, id);
+			}
+		}
+
+		if(flag & PASS_NORMAL) {
+			float3 normal = sd->N;
+			kernel_write_pass_float3(buffer + kernel_data.film.pass_normal, sample, normal);
+		}
+		if(flag & PASS_UV) {
+			float3 uv = make_float3(0.0f, 0.0f, 0.0f); /* todo: request and lookup */
+			kernel_write_pass_float3(buffer + kernel_data.film.pass_uv, sample, uv);
+		}
+	}
+
+	if(flag & (PASS_DIFFUSE_INDIRECT|PASS_DIFFUSE_COLOR|PASS_DIFFUSE_DIRECT))
+		L->color_diffuse += shader_bsdf_diffuse(kg, sd)*throughput;
+	if(flag & (PASS_GLOSSY_INDIRECT|PASS_GLOSSY_COLOR|PASS_GLOSSY_DIRECT))
+		L->color_glossy += shader_bsdf_glossy(kg, sd)*throughput;
+	if(flag & (PASS_TRANSMISSION_INDIRECT|PASS_TRANSMISSION_COLOR|PASS_TRANSMISSION_DIRECT))
+		L->color_transmission += shader_bsdf_transmission(kg, sd)*throughput;
+#endif
+}
+
+__device_inline void kernel_write_light_passes(KernelGlobals *kg, __global float *buffer, PathRadiance *L, int sample)
+{
+#ifdef __PASSES__
+	int flag = kernel_data.film.pass_flag;
+
+	if(!kernel_data.film.use_light_pass)
+		return;
+	
+	if(flag & PASS_DIFFUSE_INDIRECT) {
+		float3 color = safe_divide_color(L->indirect_diffuse, L->color_diffuse);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_indirect, sample, color);
+	}
+	if(flag & PASS_GLOSSY_INDIRECT) {
+		float3 color = safe_divide_color(L->indirect_glossy, L->color_glossy);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_indirect, sample, color);
+	}
+	if(flag & PASS_TRANSMISSION_INDIRECT) {
+		float3 color = safe_divide_color(L->indirect_transmission, L->color_transmission);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_indirect, sample, color);
+	}
+	if(flag & PASS_DIFFUSE_DIRECT) {
+		float3 color = safe_divide_color(L->direct_diffuse, L->color_diffuse);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_direct, sample, color);
+	}
+	if(flag & PASS_GLOSSY_DIRECT) {
+		float3 color = safe_divide_color(L->direct_glossy, L->color_glossy);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_direct, sample, color);
+	}
+	if(flag & PASS_TRANSMISSION_DIRECT) {
+		float3 color = safe_divide_color(L->direct_transmission, L->color_transmission);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_direct, sample, color);
+	}
+
+	if(flag & PASS_EMISSION)
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_emission, sample, L->emission);
+	if(flag & PASS_BACKGROUND)
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_background, sample, L->background);
+
+	if(flag & PASS_DIFFUSE_COLOR)
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_color, sample, L->color_diffuse);
+	if(flag & PASS_GLOSSY_COLOR)
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_color, sample, L->color_glossy);
+	if(flag & PASS_TRANSMISSION_COLOR)
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_color, sample, L->color_transmission);
+#endif
+}
+
+CCL_NAMESPACE_END
+

intern/cycles/kernel/kernel_path.h

@@ -25,39 +25,16 @@
 #else
 #include "kernel_bvh.h"
 #endif
+#include "kernel_accumulate.h"
 #include "kernel_camera.h"
 #include "kernel_shader.h"
 #include "kernel_light.h"
 #include "kernel_emission.h"
 #include "kernel_random.h"
+#include "kernel_passes.h"
 
 CCL_NAMESPACE_BEGIN
 
-#ifdef __MODIFY_TP__
-__device float3 path_terminate_modified_throughput(KernelGlobals *kg, __global float3 *buffer, int x, int y, int offset, int stride, int sample)
-{
-	/* modify throughput to influence path termination probability, to avoid
-	   darker regions receiving fewer samples than lighter regions. also RGB
-	   are weighted differently. proper validation still remains to be done. */
-	const float3 weights = make_float3(1.0f, 1.33f, 0.66f);
-	const float3 one = make_float3(1.0f, 1.0f, 1.0f);
-	const int minsample = 5;
-	const float minL = 0.1f;
-
-	if(sample >= minsample) {
-		float3 L = buffer[offset + x + y*stride];
-		float3 Lmin = make_float3(minL, minL, minL);
-		float correct = (float)(sample+1)/(float)sample;
-
-		L = film_map(L*correct, sample);
-
-		return weights/clamp(L, Lmin, one);
-	}
-
-	return weights;
-}
-#endif
-
 typedef struct PathState {
 	uint flag;
 	int bounce;
@@ -168,7 +145,7 @@ __device_inline float path_state_terminate_probability(KernelGlobals *kg, PathSt
 	return average(throughput);
 }
 
-__device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ray, Intersection *isect, float3 *light_L)
+__device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ray, Intersection *isect, BsdfEval *L_light)
 {
 	if(ray->t == 0.0f)
 		return false;
@@ -208,7 +185,7 @@ __device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ra
 				}
 
 				if(!scene_intersect(kg, ray, PATH_RAY_SHADOW_TRANSPARENT, isect)) {
-					*light_L *= throughput;
+					bsdf_eval_mul(L_light, throughput);
 					return false;
 				}
 
@@ -234,13 +211,16 @@ __device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ra
 	return result;
 }
 
-__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, float3 throughput)
+__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer)
 {
 	/* initialize */
-	float3 L = make_float3(0.0f, 0.0f, 0.0f);
+	PathRadiance L;
-	float Ltransparent = 0.0f;
+	float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
+	float L_transparent = 0.0f;
 
-#ifdef __EMISSION__
+	path_radiance_init(&L, kernel_data.film.use_light_pass);
+
+#if defined(__EMISSION__) || defined(__BACKGROUND__)
 	float ray_pdf = 0.0f;
 #endif
 	PathState state;
@@ -257,13 +237,15 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 		if(!scene_intersect(kg, &ray, visibility, &isect)) {
 			/* eval background shader if nothing hit */
 			if(kernel_data.background.transparent && (state.flag & PATH_RAY_CAMERA)) {
-				Ltransparent += average(throughput);
+				L_transparent += average(throughput);
 			}
+#ifdef __BACKGROUND__
 			else {
 				/* sample background shader */
-				float3 background_L = indirect_background(kg, &ray, state.flag, ray_pdf);
+				float3 L_background = indirect_background(kg, &ray, state.flag, ray_pdf);
-				L += throughput*background_L;
+				path_radiance_accum_background(&L, throughput, L_background, state.bounce);
 			}
+#endif
 
 			break;
 		}
@@ -274,19 +256,24 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 		float rbsdf = path_rng(kg, rng, sample, rng_offset + PRNG_BSDF);
 		shader_eval_surface(kg, &sd, rbsdf, state.flag);
 
+		kernel_write_data_passes(kg, buffer, &L, &sd, sample, state.flag, throughput);
+
 #ifdef __HOLDOUT__
 		if((sd.flag & SD_HOLDOUT) && (state.flag & PATH_RAY_CAMERA)) {
 			float3 holdout_weight = shader_holdout_eval(kg, &sd);
 
 			if(kernel_data.background.transparent)
-				Ltransparent += average(holdout_weight*throughput);
+				/* any throughput is ok, should all be identical here */
+				L_transparent += average(holdout_weight*throughput);
 		}
 #endif
 
 #ifdef __EMISSION__
 		/* emission */
-		if(sd.flag & SD_EMISSION)
+		if(sd.flag & SD_EMISSION) {
-			L += throughput*indirect_emission(kg, &sd, isect.t, state.flag, ray_pdf);
+			float3 emission = indirect_emission(kg, &sd, isect.t, state.flag, ray_pdf);
+			path_radiance_accum_emission(&L, throughput, emission, state.bounce);
+		}
 #endif
 
 		/* path termination. this is a strange place to put the termination, it's
@@ -310,7 +297,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 				float light_v = path_rng(kg, rng, sample, rng_offset + PRNG_LIGHT_V);
 
 				Ray light_ray;
-				float3 light_L;
+				BsdfEval L_light;
 
 #ifdef __MULTI_LIGHT__
 				/* index -1 means randomly sample from distribution */
@@ -320,10 +307,10 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 #else
 				const int i = -1;
 #endif
-					if(direct_emission(kg, &sd, i, light_t, light_o, light_u, light_v, &light_ray, &light_L)) {
+					if(direct_emission(kg, &sd, i, light_t, light_o, light_u, light_v, &light_ray, &L_light)) {
 						/* trace shadow ray */
-						if(!shadow_blocked(kg, &state, &light_ray, &isect, &light_L))
+						if(!shadow_blocked(kg, &state, &light_ray, &isect, &L_light))
-							L += throughput*light_L;
+							path_radiance_accum_light(&L, throughput, &L_light, state.bounce);
 					}
 #ifdef __MULTI_LIGHT__
 				}
@@ -338,7 +325,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 
 		/* sample BSDF */
 		float bsdf_pdf;
-		float3 bsdf_eval;
+		BsdfEval bsdf_eval;
 		float3 bsdf_omega_in;
 		differential3 bsdf_domega_in;
 		float bsdf_u = path_rng(kg, rng, sample, rng_offset + PRNG_BSDF_U);
@@ -350,11 +337,11 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 
 		shader_release(kg, &sd);
 
-		if(bsdf_pdf == 0.0f || is_zero(bsdf_eval))
+		if(bsdf_pdf == 0.0f || bsdf_eval_is_zero(&bsdf_eval))
 			break;
 
 		/* modify throughput */
-		throughput *= bsdf_eval/bsdf_pdf;
+		path_radiance_bsdf_bounce(&L, &throughput, &bsdf_eval, bsdf_pdf, state.bounce, label);
 
 		/* set labels */
 #if defined(__EMISSION__) || defined(__BACKGROUND__)
@@ -374,18 +361,31 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 #endif
 	}
 
-	return make_float4(L.x, L.y, L.z, 1.0f - Ltransparent);
+	float3 L_sum = path_radiance_sum(&L);
+
+	kernel_write_light_passes(kg, buffer, &L, sample);
+
+	return make_float4(L_sum.x, L_sum.y, L_sum.z, 1.0f - L_transparent);
 }
 
-__device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int sample, int x, int y, int offset, int stride)
+__device void kernel_path_trace(KernelGlobals *kg,
+	__global float *buffer, __global uint *rng_state,
+	int sample, int x, int y, int offset, int stride)
 {
+	/* buffer offset */
+	int index = offset + x + y*stride;
+	int pass_stride = kernel_data.film.pass_stride;
+
+	rng_state += index;
+	buffer += index*pass_stride;
+
 	/* initialize random numbers */
 	RNG rng;
 
 	float filter_u;
 	float filter_v;
 
-	path_rng_init(kg, rng_state, sample, &rng, x, y, offset, stride, &filter_u, &filter_v);
+	path_rng_init(kg, rng_state, sample, &rng, x, y, &filter_u, &filter_v);
 
 	/* sample camera ray */
 	Ray ray;
@@ -396,23 +396,12 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl
 	camera_sample(kg, x, y, filter_u, filter_v, lens_u, lens_v, &ray);
 
 	/* integrate */
-#ifdef __MODIFY_TP__
+	float4 L = kernel_path_integrate(kg, &rng, sample, ray, buffer);
-	float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, offset, stride, sample);
-	float4 L = kernel_path_integrate(kg, &rng, sample, ray, throughput)/throughput;
-#else
-	float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
-	float4 L = kernel_path_integrate(kg, &rng, sample, ray, throughput);
-#endif
 
 	/* accumulate result in output buffer */
-	int index = offset + x + y*stride;
+	kernel_write_pass_float4(buffer, sample, L);
 
-	if(sample == 0)
+	path_rng_end(kg, rng_state, rng);
-		buffer[index] = L;
-	else
-		buffer[index] += L;
-
-	path_rng_end(kg, rng_state, rng, x, y, offset, stride);
 }
 
 CCL_NAMESPACE_END

intern/cycles/kernel/kernel_random.h

@@ -123,7 +123,7 @@ __device_inline float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dime
 #endif
 }
 
-__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy)
+__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy)
 {
 #ifdef __SOBOL_FULL_SCREEN__
 	uint px, py;
@@ -135,19 +135,31 @@ __device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state,
 
 	*rng ^= kernel_data.integrator.seed;
 
-	*fx = size * (float)px * (1.0f/(float)0xFFFFFFFF) - x;
+	if(sample == 0) {
-	*fy = size * (float)py * (1.0f/(float)0xFFFFFFFF) - y;
+		*fx = 0.5f;
+		*fy = 0.5f;
+	}
+	else {
+		*fx = size * (float)px * (1.0f/(float)0xFFFFFFFF) - x;
+		*fy = size * (float)py * (1.0f/(float)0xFFFFFFFF) - y;
+	}
 #else
-	*rng = rng_state[offset + x + y*stride];
+	*rng = *rng_state;
 
 	*rng ^= kernel_data.integrator.seed;
 
-	*fx = path_rng(kg, rng, sample, PRNG_FILTER_U);
+	if(sample == 0) {
-	*fy = path_rng(kg, rng, sample, PRNG_FILTER_V);
+		*fx = 0.5f;
+		*fy = 0.5f;
+	}
+	else {
+		*fx = path_rng(kg, rng, sample, PRNG_FILTER_U);
+		*fy = path_rng(kg, rng, sample, PRNG_FILTER_V);
+	}
 #endif
 }
 
-__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride)
+__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng)
 {
 	/* nothing to do */
 }
@@ -163,21 +175,27 @@ __device float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dimension)
 	return (float)*rng * (1.0f/(float)0xFFFFFFFF);
 }
 
-__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy)
+__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy)
 {
 	/* load state */
-	*rng = rng_state[offset + x + y*stride];
+	*rng = *rng_state;
 
 	*rng ^= kernel_data.integrator.seed;
 
-	*fx = path_rng(kg, rng, sample, PRNG_FILTER_U);
+	if(sample == 0) {
-	*fy = path_rng(kg, rng, sample, PRNG_FILTER_V);
+		*fx = 0.5f;
+		*fy = 0.5f;
+	}
+	else {
+		*fx = path_rng(kg, rng, sample, PRNG_FILTER_U);
+		*fy = path_rng(kg, rng, sample, PRNG_FILTER_V);
+	}
 }
 
-__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride)
+__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng)
 {
 	/* store state for next sample */
-	rng_state[offset + x + y*stride] = rng;
+	*rng_state = rng;
 }
 
 #endif

intern/cycles/kernel/kernel_shader.h

@@ -26,7 +26,6 @@
  *
  */
 
-
 #ifdef __OSL__
 
 #include "osl_shader.h"
@@ -75,7 +74,7 @@ __device_inline void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd,
 	if(sd->shader & SHADER_SMOOTH_NORMAL)
 		sd->N = triangle_smooth_normal(kg, sd->prim, sd->u, sd->v);
 
-	sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK);
+	sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2);
 
 #ifdef __DPDU__
 	/* dPdu/dPdv */
@@ -166,7 +165,7 @@ __device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd,
 #endif
 	}
 
-	sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK);
+	sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2);
 
 #ifdef __DPDU__
 	/* dPdu/dPdv */
@@ -243,7 +242,7 @@ __device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData
 	sd->Ng = -sd->P;
 	sd->I = -sd->P;
 	sd->shader = kernel_data.background.shader;
-	sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK);
+	sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2);
 
 #ifdef __INSTANCING__
 	sd->object = ~0;
@@ -275,8 +274,8 @@ __device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData
 
 #ifdef __MULTI_CLOSURE__
 
-__device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float3 omega_in, float *pdf,
+__device_inline void _shader_bsdf_multi_eval(const ShaderData *sd, const float3 omega_in, float *pdf,
-	int skip_bsdf, float3 sum_eval, float sum_pdf, float sum_sample_weight)
+	int skip_bsdf, BsdfEval *bsdf_eval, float sum_pdf, float sum_sample_weight)
 {
 	for(int i = 0; i< sd->num_closure; i++) {
 		if(i == skip_bsdf)
@@ -293,7 +292,7 @@ __device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float
 #endif
 
 			if(bsdf_pdf != 0.0f) {
-				sum_eval += eval*sc->weight;
+				bsdf_eval_accum(bsdf_eval, sc->type, eval*sc->weight);
 				sum_pdf += bsdf_pdf*sc->sample_weight;
 			}
 
@@ -302,25 +301,27 @@ __device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float
 	}
 
 	*pdf = sum_pdf/sum_sample_weight;
-	return sum_eval;
 }
 
 #endif
 
-__device float3 shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd,
+__device void shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd,
-	const float3 omega_in, float *pdf)
+	const float3 omega_in, BsdfEval *eval, float *pdf)
 {
 #ifdef __MULTI_CLOSURE__
-	return _shader_bsdf_multi_eval(sd, omega_in, pdf, -1, make_float3(0.0f, 0.0f, 0.0f), 0.0f, 0.0f);
+	bsdf_eval_init(eval, NBUILTIN_CLOSURES, make_float3(0.0f, 0.0f, 0.0f), kernel_data.film.use_light_pass);
+
+	return _shader_bsdf_multi_eval(sd, omega_in, pdf, -1, eval, 0.0f, 0.0f);
 #else
 	const ShaderClosure *sc = &sd->closure;
+
 	*pdf = 0.0f;
-	return svm_bsdf_eval(sd, sc, omega_in, pdf)*sc->weight;
+	*eval = svm_bsdf_eval(sd, sc, omega_in, pdf)*sc->weight;
 #endif
 }
 
 __device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd,
-	float randu, float randv, float3 *eval,
+	float randu, float randv, BsdfEval *bsdf_eval,
 	float3 *omega_in, differential3 *domega_in, float *pdf)
 {
 #ifdef __MULTI_CLOSURE__
@@ -359,27 +360,28 @@ __device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd,
 
 	const ShaderClosure *sc = &sd->closure[sampled];
 	int label;
+	float3 eval;
 
 	*pdf = 0.0f;
 #ifdef __OSL__
-	label = OSLShader::bsdf_sample(sd, sc, randu, randv, *eval, *omega_in, *domega_in, *pdf);
+	label = OSLShader::bsdf_sample(sd, sc, randu, randv, eval, *omega_in, *domega_in, *pdf);
 #else
-	label = svm_bsdf_sample(sd, sc, randu, randv, eval, omega_in, domega_in, pdf);
+	label = svm_bsdf_sample(sd, sc, randu, randv, &eval, omega_in, domega_in, pdf);
 #endif
 
-	*eval *= sc->weight;
+	bsdf_eval_init(bsdf_eval, sc->type, eval*sc->weight, kernel_data.film.use_light_pass);
 
 	if(sd->num_closure > 1 && *pdf != 0.0f) {
 		float sweight = sc->sample_weight;
-		*eval = _shader_bsdf_multi_eval(sd, *omega_in, pdf, sampled, *eval, *pdf*sweight, sweight);
+		_shader_bsdf_multi_eval(sd, *omega_in, pdf, sampled, bsdf_eval, *pdf*sweight, sweight);
 	}
 
 	return label;
 #else
 	/* sample the single closure that we picked */
 	*pdf = 0.0f;
-	int label = svm_bsdf_sample(sd, &sd->closure, randu, randv, eval, omega_in, domega_in, pdf);
+	int label = svm_bsdf_sample(sd, &sd->closure, randu, randv, bsdf_eval, omega_in, domega_in, pdf);
-	*eval *= sd->closure.weight;
+	*bsdf_eval *= sd->closure.weight;
 	return label;
 #endif
 }
@@ -421,6 +423,68 @@ __device float3 shader_bsdf_transparency(KernelGlobals *kg, ShaderData *sd)
 #endif
 }
 
+__device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd)
+{
+#ifdef __MULTI_CLOSURE__
+	float3 eval = make_float3(0.0f, 0.0f, 0.0f);
+
+	for(int i = 0; i< sd->num_closure; i++) {
+		ShaderClosure *sc = &sd->closure[i];
+
+		if(CLOSURE_IS_BSDF_DIFFUSE(sc->type))
+			eval += sc->weight;
+	}
+
+	return eval;
+#else
+	if(CLOSURE_IS_BSDF_DIFFUSE(sd->closure.type))
+		return sd->closure.weight;
+	else
+		return make_float3(0.0f, 0.0f, 0.0f);
+#endif
+}
+
+__device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd)
+{
+#ifdef __MULTI_CLOSURE__
+	float3 eval = make_float3(0.0f, 0.0f, 0.0f);
+
+	for(int i = 0; i< sd->num_closure; i++) {
+		ShaderClosure *sc = &sd->closure[i];
+
+		if(CLOSURE_IS_BSDF_GLOSSY(sc->type))
+			eval += sc->weight;
+	}
+
+	return eval;
+#else
+	if(CLOSURE_IS_BSDF_GLOSSY(sd->closure.type))
+		return sd->closure.weight;
+	else
+		return make_float3(0.0f, 0.0f, 0.0f);
+#endif
+}
+
+__device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd)
+{
+#ifdef __MULTI_CLOSURE__
+	float3 eval = make_float3(0.0f, 0.0f, 0.0f);
+
+	for(int i = 0; i< sd->num_closure; i++) {
+		ShaderClosure *sc = &sd->closure[i];
+
+		if(CLOSURE_IS_BSDF_TRANSMISSION(sc->type))
+			eval += sc->weight;
+	}
+
+	return eval;
+#else
+	if(CLOSURE_IS_BSDF_TRANSMISSION(sd->closure.type))
+		return sd->closure.weight;
+	else
+		return make_float3(0.0f, 0.0f, 0.0f);
+#endif
+}
 
 /* Emission */
 
@@ -588,12 +652,17 @@ __device bool shader_transparent_shadow(KernelGlobals *kg, Intersection *isect)
 	int prim = kernel_tex_fetch(__prim_index, isect->prim);
 	float4 Ns = kernel_tex_fetch(__tri_normal, prim);
 	int shader = __float_as_int(Ns.w);
-	int flag = kernel_tex_fetch(__shader_flag, shader & SHADER_MASK);
+	int flag = kernel_tex_fetch(__shader_flag, (shader & SHADER_MASK)*2);
 
 	return (flag & SD_HAS_SURFACE_TRANSPARENT) != 0;
 }
 #endif
 
+__device int shader_pass_id(KernelGlobals *kg, ShaderData *sd)
+{
+	return kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2 + 1);
+}
+
 /* Free ShaderData */
 
 __device void shader_release(KernelGlobals *kg, ShaderData *sd)

intern/cycles/kernel/kernel_types.h

@@ -70,6 +70,8 @@ CCL_NAMESPACE_BEGIN
 #ifdef __KERNEL_ADV_SHADING__
 #define __MULTI_CLOSURE__
 #define __TRANSPARENT_SHADOWS__
+#define __PASSES__
+#define __BACKGROUND_MIS__
 #endif
 
 //#define __MULTI_LIGHT__
@@ -150,6 +152,75 @@ typedef enum ClosureLabel {
 	LABEL_STOP = 2048
 } ClosureLabel;
 
+/* Render Passes */
+
+typedef enum PassType {
+	PASS_NONE = 0,
+	PASS_COMBINED = 1,
+	PASS_DEPTH = 2,
+	PASS_NORMAL = 8,
+	PASS_UV = 16,
+	PASS_OBJECT_ID = 32,
+	PASS_MATERIAL_ID = 64,
+	PASS_DIFFUSE_COLOR = 128,
+	PASS_GLOSSY_COLOR = 256,
+	PASS_TRANSMISSION_COLOR = 512,
+	PASS_DIFFUSE_INDIRECT = 1024,
+	PASS_GLOSSY_INDIRECT = 2048,
+	PASS_TRANSMISSION_INDIRECT = 4096,
+	PASS_DIFFUSE_DIRECT = 8192,
+	PASS_GLOSSY_DIRECT = 16384,
+	PASS_TRANSMISSION_DIRECT = 32768,
+	PASS_EMISSION = 65536,
+	PASS_BACKGROUND = 131072
+} PassType;
+
+#define PASS_ALL (~0)
+
+#ifdef __PASSES__
+
+typedef float3 PathThroughput;
+
+struct PathRadiance {
+	int use_light_pass;
+
+	float3 emission;
+	float3 background;
+
+	float3 indirect;
+	float3 direct_throughput;
+	float3 direct_emission;
+
+	float3 color_diffuse;
+	float3 color_glossy;
+	float3 color_transmission;
+
+	float3 direct_diffuse;
+	float3 direct_glossy;
+	float3 direct_transmission;
+
+	float3 indirect_diffuse;
+	float3 indirect_glossy;
+	float3 indirect_transmission;
+};
+
+struct BsdfEval {
+	int use_light_pass;
+
+	float3 diffuse;
+	float3 glossy;
+	float3 transmission;
+	float3 transparent;
+};
+
+#else
+
+typedef float3 PathThroughput;
+typedef float3 PathRadiance;
+typedef float3 BsdfEval;
+
+#endif
+
 /* Shader Flag */
 
 typedef enum ShaderFlag {
@@ -353,7 +424,32 @@ typedef struct KernelCamera {
 
 typedef struct KernelFilm {
 	float exposure;
-	int pad1, pad2, pad3;
+	int pass_flag;
+	int pass_stride;
+	int use_light_pass;
+
+	int pass_combined;
+	int pass_depth;
+	int pass_normal;
+	int pass_pad;
+
+	int pass_uv;
+	int pass_object_id;
+	int pass_material_id;
+	int pass_diffuse_color;
+
+	int pass_glossy_color;
+	int pass_transmission_color;
+	int pass_diffuse_indirect;
+	int pass_glossy_indirect;
+
+	int pass_transmission_indirect;
+	int pass_diffuse_direct;
+	int pass_glossy_direct;
+	int pass_transmission_direct;
+
+	int pass_emission;
+	int pass_background;
 } KernelFilm;
 
 typedef struct KernelBackground {

intern/cycles/kernel/osl/nodes/CMakeLists.txt

@@ -20,6 +20,7 @@ set(SRC_OSL
 	node_environment_texture.osl
 	node_fresnel.osl
 	node_gamma.osl
+	node_brightness.osl
 	node_geometry.osl
 	node_glass_bsdf.osl
 	node_glossy_bsdf.osl

intern/cycles/kernel/osl/nodes/node_brightness.osl

@@ -0,0 +1,50 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ */
+
+#include "stdosl.h"
+
+shader node_brightness(
+	color ColorIn = color(0.8, 0.8, 0.8),
+	float Bright = 0.0,
+	float Contrast = 0.0,
+	output ColorOut = color(0.8, 0.8, 0.8)
+{
+	float delta = Contrast * (1.0/200.0);
+	float a = 1.0 - delta * 2.0;
+	float b;
+
+	Bright *= 1.0/100.0;
+
+	/*
+	* The algorithm is by Werner D. Streidt
+	* (http://visca.com/ffactory/archives/5-99/msg00021.html)
+	* Extracted of OpenCV demhist.c
+	*/
+
+	if (Contrast > 0.0) {
+		a = (a < 0.0 ? 1.0/a : 0.0);
+		b = a * (Brightness - delta);
+	}
+	else {
+		delta *= -1.0;
+		b = a * (Brightness + delta);
+	}
+
+	ColorOut = a * ColorIn + b;
+}
+

intern/cycles/kernel/svm/svm.h

@@ -131,6 +131,7 @@ CCL_NAMESPACE_END
 #include "svm_hsv.h"
 #include "svm_image.h"
 #include "svm_gamma.h"
+#include "svm_brightness.h"
 #include "svm_invert.h"
 #include "svm_light_path.h"
 #include "svm_magic.h"
@@ -270,6 +271,9 @@ __device_noinline void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderT
 			case NODE_GAMMA:
 				svm_node_gamma(sd, stack, node.y, node.z, node.w);
 				break;
+			case NODE_BRIGHTCONTRAST:
+				svm_node_brightness(sd, stack, node.y, node.z, node.w);
+				break;
 			case NODE_MIX:
 				svm_node_mix(kg, sd, stack, node.y, node.z, node.w, &offset);
 				break;

intern/cycles/kernel/svm/svm_brightness.h

@@ -0,0 +1,57 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ */
+
+CCL_NAMESPACE_BEGIN
+
+__device void svm_node_brightness(ShaderData *sd, float *stack, uint in_color, uint out_color, uint node)
+{
+	uint bright_offset, contrast_offset;
+	float3 color = stack_load_float3(stack, in_color);
+
+	decode_node_uchar4(node, &bright_offset, &contrast_offset, NULL, NULL);
+	float brightness = stack_load_float(stack, bright_offset);
+	float contrast  = stack_load_float(stack, contrast_offset);
+
+	brightness *= 1.0f/100.0f;
+	float delta = contrast * (1.0f/200.0f);
+	float a = 1.0f - delta * 2.0f;
+	float b;
+
+	/*
+	* The algorithm is by Werner D. Streidt
+	* (http://visca.com/ffactory/archives/5-99/msg00021.html)
+	* Extracted of OpenCV demhist.c
+	*/
+	if (contrast > 0.0f) {
+		a = (a > 0.0f? (1.0f / a): 0.0f);
+		b = a * (brightness - delta);
+	}
+	else {
+		delta *= -1.0f;
+		b = a * (brightness + delta);
+	}
+
+	color.x = a*color.x + b;
+	color.y = a*color.y + b;
+	color.z = a*color.z + b;
+
+	if (stack_valid(out_color))
+		stack_store_float3(stack, out_color, color);
+}
+
+CCL_NAMESPACE_END

intern/cycles/kernel/svm/svm_types.h

@@ -87,7 +87,8 @@ typedef enum NodeType {
 	NODE_INVERT = 5400,
 	NODE_NORMAL = 5500,
 	NODE_GAMMA = 5600,
-	NODE_TEX_CHECKER = 5700
+	NODE_TEX_CHECKER = 5700,
+	NODE_BRIGHTCONTRAST = 5800
 } NodeType;
 
 typedef enum NodeAttributeType {
@@ -265,22 +266,26 @@ typedef enum ShaderType {
 
 typedef enum ClosureType {
 	CLOSURE_BSDF_ID,
+
 	CLOSURE_BSDF_DIFFUSE_ID,
 	CLOSURE_BSDF_OREN_NAYAR_ID,
-	CLOSURE_BSDF_TRANSLUCENT_ID,
+
 	CLOSURE_BSDF_REFLECTION_ID,
-	CLOSURE_BSDF_REFRACTION_ID,
-	CLOSURE_BSDF_GLASS_ID,
-	CLOSURE_BSDF_TRANSPARENT_ID,
 	CLOSURE_BSDF_MICROFACET_GGX_ID,
-	CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID,
 	CLOSURE_BSDF_MICROFACET_BECKMANN_ID,
-	CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID,
 	CLOSURE_BSDF_WARD_ID,
-	CLOSURE_BSDF_ASHIKHMIN_VELVET_ID,
-	CLOSURE_BSDF_WESTIN_BACKSCATTER_ID,
 	CLOSURE_BSDF_WESTIN_SHEEN_ID,
 
+	CLOSURE_BSDF_TRANSLUCENT_ID,
+	CLOSURE_BSDF_REFRACTION_ID,
+	CLOSURE_BSDF_WESTIN_BACKSCATTER_ID,
+	CLOSURE_BSDF_ASHIKHMIN_VELVET_ID,
+	CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID,
+	CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID,
+	CLOSURE_BSDF_GLASS_ID,
+
+	CLOSURE_BSDF_TRANSPARENT_ID,
+
 	CLOSURE_BSSRDF_CUBIC_ID,
 	CLOSURE_EMISSION_ID,
 	CLOSURE_DEBUG_ID,
@@ -296,7 +301,10 @@ typedef enum ClosureType {
 } ClosureType;
 
 /* watch this, being lazy with memory usage */
-#define CLOSURE_IS_BSDF(type) (type <= CLOSURE_BSDF_WESTIN_SHEEN_ID)
+#define CLOSURE_IS_BSDF(type) (type <= CLOSURE_BSDF_TRANSPARENT_ID)
+#define CLOSURE_IS_BSDF_DIFFUSE(type) (type >= CLOSURE_BSDF_DIFFUSE_ID && type <= CLOSURE_BSDF_OREN_NAYAR_ID)
+#define CLOSURE_IS_BSDF_GLOSSY(type) (type >= CLOSURE_BSDF_REFLECTION_ID && type <= CLOSURE_BSDF_WESTIN_SHEEN_ID)
+#define CLOSURE_IS_BSDF_TRANSMISSION(type) (type >= CLOSURE_BSDF_TRANSLUCENT_ID && type <= CLOSURE_BSDF_GLASS_ID)
 #define CLOSURE_IS_VOLUME(type) (type >= CLOSURE_VOLUME_ID && type <= CLOSURE_VOLUME_ISOTROPIC_ID)
 #define CLOSURE_IS_EMISSION(type) (type == CLOSURE_EMISSION_ID)
 #define CLOSURE_IS_HOLDOUT(type) (type == CLOSURE_HOLDOUT_ID)

intern/cycles/render/buffers.cpp

@@ -22,6 +22,7 @@
 #include "device.h"
 
 #include "util_debug.h"
+#include "util_foreach.h"
 #include "util_hash.h"
 #include "util_image.h"
 #include "util_math.h"
@@ -31,6 +32,48 @@
 
 CCL_NAMESPACE_BEGIN
 
+/* Buffer Params */
+
+BufferParams::BufferParams()
+{
+	width = 0;
+	height = 0;
+
+	full_x = 0;
+	full_y = 0;
+	full_width = 0;
+	full_height = 0;
+
+	Pass::add(PASS_COMBINED, passes);
+}
+
+void BufferParams::get_offset_stride(int& offset, int& stride)
+{
+	offset = -(full_x + full_y*width);
+	stride = width;
+}
+
+bool BufferParams::modified(const BufferParams& params)
+{
+	return !(full_x == params.full_x
+		&& full_y == params.full_y
+		&& width == params.width
+		&& height == params.height
+		&& full_width == params.full_width
+		&& full_height == params.full_height
+		&& Pass::equals(passes, params.passes));
+}
+
+int BufferParams::get_passes_size()
+{
+	int size = 0;
+
+	foreach(Pass& pass, passes)
+		size += pass.components;
+	
+	return align_up(size, 4);
+}
+
 /* Render Buffers */
 
 RenderBuffers::RenderBuffers(Device *device_)
@@ -64,7 +107,7 @@ void RenderBuffers::reset(Device *device, BufferParams& params_)
 	device_free();
 	
 	/* allocate buffer */
-	buffer.resize(params.width, params.height);
+	buffer.resize(params.width*params.height*params.get_passes_size());
 	device->mem_alloc(buffer, MEM_READ_WRITE);
 	device->mem_zero(buffer);
 
@@ -82,31 +125,83 @@ void RenderBuffers::reset(Device *device, BufferParams& params_)
 	device->mem_copy_to(rng_state);
 }
 
-float4 *RenderBuffers::copy_from_device(float exposure, int sample)
+bool RenderBuffers::copy_from_device()
 {
 	if(!buffer.device_pointer)
-		return NULL;
+		return false;
 
-	device->mem_copy_from(buffer, 0, params.width, params.height, sizeof(float4));
+	device->mem_copy_from(buffer, 0, params.width, params.height, params.get_passes_size()*sizeof(float));
 
-	float4 *out = new float4[params.width*params.height];
+	return true;
-	float4 *in = (float4*)buffer.data_pointer;
+}
-	float scale = 1.0f/(float)sample;
-	
-	for(int i = params.width*params.height - 1; i >= 0; i--) {
-		float4 rgba = in[i]*scale;
 
-		rgba.x = rgba.x*exposure;
+bool RenderBuffers::get_pass(PassType type, float exposure, int sample, int components, float *pixels)
-		rgba.y = rgba.y*exposure;
+{
-		rgba.z = rgba.z*exposure;
+	int pass_offset = 0;
 
-		/* clamp since alpha might be > 1.0 due to russian roulette */
+	foreach(Pass& pass, params.passes) {
-		rgba.w = clamp(rgba.w, 0.0f, 1.0f);
+		if(pass.type != type) {
+			pass_offset += pass.components;
+			continue;
+		}
 
-		out[i] = rgba;
+		float *in = (float*)buffer.data_pointer + pass_offset;
+		int pass_stride = params.get_passes_size();
+
+		float scale = (pass.filter)? 1.0f/(float)sample: 1.0f;
+		float scale_exposure = (pass.exposure)? scale*exposure: scale;
+
+		int size = params.width*params.height;
+
+		if(components == 1) {
+			assert(pass.components == components);
+
+			/* scalar */
+			if(type == PASS_DEPTH) {
+				for(int i = 0; i < size; i++, in += pass_stride, pixels++) {
+					float f = *in;
+					pixels[0] = (f == 0.0f)? 1e10f: f*scale_exposure;
+				}
+			}
+			else {
+				for(int i = 0; i < size; i++, in += pass_stride, pixels++) {
+					float f = *in;
+					pixels[0] = f*scale_exposure;
+				}
+			}
+		}
+		else if(components == 3) {
+			assert(pass.components == 4);
+
+			/* RGB/vector */
+			for(int i = 0; i < size; i++, in += pass_stride, pixels += 3) {
+				float3 f = make_float3(in[0], in[1], in[2]);
+
+				pixels[0] = f.x*scale_exposure;
+				pixels[1] = f.y*scale_exposure;
+				pixels[2] = f.z*scale_exposure;
+			}
+		}
+		else if(components == 4) {
+			assert(pass.components == components);
+
+			/* RGBA */
+			for(int i = 0; i < size; i++, in += pass_stride, pixels += 4) {
+				float4 f = make_float4(in[0], in[1], in[2], in[3]);
+
+				pixels[0] = f.x*scale_exposure;
+				pixels[1] = f.y*scale_exposure;
+				pixels[2] = f.z*scale_exposure;
+
+				/* clamp since alpha might be > 1.0 due to russian roulette */
+				pixels[3] = clamp(f.w*scale, 0.0f, 1.0f);
+			}
+		}
+
+		return true;
 	}
 
-	return out;
+	return false;
 }
 
 /* Display Buffer */

intern/cycles/render/buffers.h

@@ -21,6 +21,10 @@
 
 #include "device_memory.h"
 
+#include "film.h"
+
+#include "kernel_types.h"
+
 #include "util_string.h"
 #include "util_thread.h"
 #include "util_types.h"
@@ -45,32 +49,16 @@ public:
 	int full_width;
 	int full_height;
 
-	BufferParams()
+	/* passes */
-	{
+	vector<Pass> passes;
-		width = 0;
-		height = 0;
 
-		full_x = 0;
+	/* functions */
-		full_y = 0;
+	BufferParams();
-		full_width = 0;
-		full_height = 0;
-	}
 
-	void get_offset_stride(int& offset, int& stride)
+	void get_offset_stride(int& offset, int& stride);
-	{
+	bool modified(const BufferParams& params);
-		offset = -(full_x + full_y*width);
+	void add_pass(PassType type);
-		stride = width;
+	int get_passes_size();
-	}
-
-	bool modified(const BufferParams& params)
-	{
-		return !(full_x == params.full_x
-			&& full_y == params.full_y
-			&& width == params.width
-			&& height == params.height
-			&& full_width == params.full_width
-			&& full_height == params.full_height);
-	}
 };
 
 /* Render Buffers */
@@ -80,7 +68,7 @@ public:
 	/* buffer parameters */
 	BufferParams params;
 	/* float buffer */
-	device_vector<float4> buffer;
+	device_vector<float> buffer;
 	/* random number generator state */
 	device_vector<uint> rng_state;
 	/* mutex, must be locked manually by callers */
@@ -90,7 +78,9 @@ public:
 	~RenderBuffers();
 
 	void reset(Device *device, BufferParams& params);
-	float4 *copy_from_device(float exposure, int sample);
+
+	bool copy_from_device();
+	bool get_pass(PassType type, float exposure, int sample, int components, float *pixels);
 
 protected:
 	void device_free();

intern/cycles/render/film.cpp

@@ -21,11 +21,123 @@
 #include "film.h"
 #include "scene.h"
 
+#include "util_algorithm.h"
+#include "util_foreach.h"
+
 CCL_NAMESPACE_BEGIN
 
+/* Pass */
+
+static bool compare_pass_order(const Pass& a, const Pass& b)
+{
+	if(a.components == b.components)
+		return (a.type < b.type);
+	return (a.components > b.components);
+}
+
+void Pass::add(PassType type, vector<Pass>& passes)
+{
+	Pass pass;
+
+	pass.type = type;
+	pass.filter = true;
+	pass.exposure = false;
+
+	switch(type) {
+		case PASS_NONE:
+			pass.components = 0;
+			break;
+		case PASS_COMBINED:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_DEPTH:
+			pass.components = 1;
+			pass.filter = false;
+			break;
+		case PASS_NORMAL:
+			pass.components = 4;
+			break;
+		case PASS_UV:
+			pass.components = 4;
+			break;
+		case PASS_OBJECT_ID:
+			pass.components = 1;
+			pass.filter = false;
+			break;
+		case PASS_MATERIAL_ID:
+			pass.components = 1;
+			pass.filter = false;
+			break;
+		case PASS_DIFFUSE_COLOR:
+			pass.components = 4;
+			break;
+		case PASS_GLOSSY_COLOR:
+			pass.components = 4;
+			break;
+		case PASS_TRANSMISSION_COLOR:
+			pass.components = 4;
+			break;
+		case PASS_DIFFUSE_INDIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_GLOSSY_INDIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_TRANSMISSION_INDIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_DIFFUSE_DIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_GLOSSY_DIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_TRANSMISSION_DIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+
+		case PASS_EMISSION:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_BACKGROUND:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+	}
+
+	passes.push_back(pass);
+
+	/* order from by components, to ensure alignment so passes with size 4
+	   come first and then passes with size 1 */
+	sort(passes.begin(), passes.end(), compare_pass_order);
+}
+
+bool Pass::equals(const vector<Pass>& A, const vector<Pass>& B)
+{
+	if(A.size() != B.size())
+		return false;
+	
+	for(int i = 0; i < A.size(); i++)
+		if(A[i].type != B[i].type)
+			return false;
+	
+	return true;
+}
+
+/* Film */
+
 Film::Film()
 {
 	exposure = 0.8f;
+	Pass::add(PASS_COMBINED, passes);
 	need_update = true;
 }
 
@@ -42,6 +154,84 @@ void Film::device_update(Device *device, DeviceScene *dscene)
 
 	/* update __data */
 	kfilm->exposure = exposure;
+	kfilm->pass_flag = 0;
+	kfilm->pass_stride = 0;
+	kfilm->use_light_pass = 0;
+
+	foreach(Pass& pass, passes) {
+		kfilm->pass_flag |= pass.type;
+
+		switch(pass.type) {
+			case PASS_COMBINED:
+				kfilm->pass_combined = kfilm->pass_stride;
+				break;
+			case PASS_DEPTH:
+				kfilm->pass_depth = kfilm->pass_stride;
+				break;
+			case PASS_NORMAL:
+				kfilm->pass_normal = kfilm->pass_stride;
+				break;
+			case PASS_UV:
+				kfilm->pass_uv = kfilm->pass_stride;
+				break;
+			case PASS_OBJECT_ID:
+				kfilm->pass_object_id = kfilm->pass_stride;
+				break;
+			case PASS_MATERIAL_ID:
+				kfilm->pass_material_id = kfilm->pass_stride;
+				break;
+			case PASS_DIFFUSE_COLOR:
+				kfilm->pass_diffuse_color = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_GLOSSY_COLOR:
+				kfilm->pass_glossy_color = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_TRANSMISSION_COLOR:
+				kfilm->pass_transmission_color = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_DIFFUSE_INDIRECT:
+				kfilm->pass_diffuse_indirect = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_GLOSSY_INDIRECT:
+				kfilm->pass_glossy_indirect = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_TRANSMISSION_INDIRECT:
+				kfilm->pass_transmission_indirect = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_DIFFUSE_DIRECT:
+				kfilm->pass_diffuse_direct = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_GLOSSY_DIRECT:
+				kfilm->pass_glossy_direct = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_TRANSMISSION_DIRECT:
+				kfilm->pass_transmission_direct = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+
+			case PASS_EMISSION:
+				kfilm->pass_emission = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_BACKGROUND:
+				kfilm->pass_background = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+			case PASS_NONE:
+				break;
+		}
+
+		kfilm->pass_stride += pass.components;
+	}
+
+	kfilm->pass_stride = align_up(kfilm->pass_stride, 4);
 
 	need_update = false;
 }
@@ -52,7 +242,8 @@ void Film::device_free(Device *device, DeviceScene *dscene)
 
 bool Film::modified(const Film& film)
 {
-	return !(exposure == film.exposure);
+	return !(exposure == film.exposure
+		&& Pass::equals(passes, film.passes));
 }
 
 void Film::tag_update(Scene *scene)

intern/cycles/render/film.h

@@ -20,6 +20,9 @@
 #define __FILM_H__
 
 #include "util_string.h"
+#include "util_vector.h"
+
+#include "kernel_types.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -27,9 +30,21 @@ class Device;
 class DeviceScene;
 class Scene;
 
+class Pass {
+public:
+	PassType type;
+	int components;
+	bool filter;
+	bool exposure;
+
+	static void add(PassType type, vector<Pass>& passes);
+	static bool equals(const vector<Pass>& A, const vector<Pass>& B);
+};
+
 class Film {
 public:
 	float exposure;
+	vector<Pass> passes;
 	bool need_update;
 
 	Film();

intern/cycles/render/light.cpp

@@ -402,6 +402,16 @@ void LightManager::device_update_points(Device *device, DeviceScene *dscene, Sce
 
 	float4 *light_data = dscene->light_data.resize(scene->lights.size()*LIGHT_SIZE);
 
+	if(!device->info.advanced_shading) {
+		/* remove unsupported light */
+		foreach(Light *light, scene->lights) {
+			if(light->type == LIGHT_BACKGROUND) {
+				scene->lights.erase(std::remove(scene->lights.begin(), scene->lights.end(), light), scene->lights.end());
+				break;
+			}
+		}
+	}
+
 	for(size_t i = 0; i < scene->lights.size(); i++) {
 		Light *light = scene->lights[i];
 		float3 co = light->co;
@@ -431,7 +441,7 @@ void LightManager::device_update_points(Device *device, DeviceScene *dscene, Sce
 			shader_id &= ~SHADER_AREA_LIGHT;
 
 			light_data[i*LIGHT_SIZE + 0] = make_float4(__int_as_float(light->type), 0.0f, 0.0f, 0.0f);
-			light_data[i*LIGHT_SIZE + 1] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
+			light_data[i*LIGHT_SIZE + 1] = make_float4(__int_as_float(shader_id), 0.0f, 0.0f, 0.0f);
 			light_data[i*LIGHT_SIZE + 2] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 			light_data[i*LIGHT_SIZE + 3] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 		}

intern/cycles/render/nodes.cpp

@@ -1820,6 +1820,38 @@ void GammaNode::compile(OSLCompiler& compiler)
 	compiler.add(this, "node_gamma");
 }
 
+/* Bright Contrast */
+BrightContrastNode::BrightContrastNode()
+: ShaderNode("brightness")
+{
+	add_input("Color", SHADER_SOCKET_COLOR);
+	add_input("Bright", SHADER_SOCKET_FLOAT);
+	add_input("Contrast", SHADER_SOCKET_FLOAT);
+	add_output("Color", SHADER_SOCKET_COLOR);
+}
+
+void BrightContrastNode::compile(SVMCompiler& compiler)
+{
+	ShaderInput *color_in = input("Color");
+	ShaderInput *bright_in = input("Bright");
+	ShaderInput *contrast_in = input("Contrast");
+	ShaderOutput *color_out = output("Color");
+
+	compiler.stack_assign(color_in);
+	compiler.stack_assign(bright_in);
+	compiler.stack_assign(contrast_in);
+	compiler.stack_assign(color_out);
+
+	compiler.add_node(NODE_BRIGHTCONTRAST,
+		color_in->stack_offset, color_out->stack_offset,
+		compiler.encode_uchar4(bright_in->stack_offset, contrast_in->stack_offset));
+}
+
+void BrightContrastNode::compile(OSLCompiler& compiler)
+{
+	compiler.add(this, "node_brightness");
+}
+
 /* Separate RGB */
 SeparateRGBNode::SeparateRGBNode()
 : ShaderNode("separate_rgb")

intern/cycles/render/nodes.h

@@ -320,6 +320,11 @@ public:
 	SHADER_NODE_CLASS(GammaNode)
 };
 
+class BrightContrastNode : public ShaderNode {
+public:
+	SHADER_NODE_CLASS(BrightContrastNode)
+};
+
 class SeparateRGBNode : public ShaderNode {
 public:
 	SHADER_NODE_CLASS(SeparateRGBNode)

intern/cycles/render/object.cpp

@@ -36,6 +36,7 @@ Object::Object()
 	mesh = NULL;
 	tfm = transform_identity();
 	visibility = ~0;
+	pass_id = 0;
 }
 
 Object::~Object()
@@ -135,6 +136,7 @@ void ObjectManager::device_update_transforms(Device *device, DeviceScene *dscene
 		/* todo: correct for displacement, and move to a better place */
 		float uniform_scale;
 		float surface_area = 0.0f;
+		float pass_id = ob->pass_id;
 		
 		if(transform_uniform_scale(tfm, uniform_scale)) {
 			map<Mesh*, float>::iterator it = surface_area_map.find(mesh);
@@ -171,7 +173,7 @@ void ObjectManager::device_update_transforms(Device *device, DeviceScene *dscene
 		memcpy(&objects[offset], &tfm, sizeof(float4)*4);
 		memcpy(&objects[offset+4], &itfm, sizeof(float4)*4);
 		memcpy(&objects[offset+8], &ntfm, sizeof(float4)*4);
-		objects[offset+12] = make_float4(surface_area, 0.0f, 0.0f, 0.0f);
+		objects[offset+12] = make_float4(surface_area, pass_id, 0.0f, 0.0f);
 
 		i++;
 

intern/cycles/render/object.h

@@ -41,6 +41,7 @@ public:
 	Transform tfm;
 	BoundBox bounds;
 	ustring name;
+	int pass_id;
 	vector<ParamValue> attributes;
 	uint visibility;
 

intern/cycles/render/shader.cpp

@@ -35,6 +35,7 @@ CCL_NAMESPACE_BEGIN
 Shader::Shader()
 {
 	name = "";
+	pass_id = 0;
 
 	graph = NULL;
 	graph_bump = NULL;
@@ -167,7 +168,7 @@ void ShaderManager::device_update_common(Device *device, DeviceScene *dscene, Sc
 	if(scene->shaders.size() == 0)
 		return;
 
-	uint shader_flag_size = scene->shaders.size()*2;
+	uint shader_flag_size = scene->shaders.size()*4;
 	uint *shader_flag = dscene->shader_flag.resize(shader_flag_size);
 	uint i = 0;
 
@@ -184,7 +185,9 @@ void ShaderManager::device_update_common(Device *device, DeviceScene *dscene, Sc
 			flag |= SD_HOMOGENEOUS_VOLUME;
 
 		shader_flag[i++] = flag;
+		shader_flag[i++] = shader->pass_id;
 		shader_flag[i++] = flag;
+		shader_flag[i++] = shader->pass_id;
 	}
 
 	device->tex_alloc("__shader_flag", dscene->shader_flag);

intern/cycles/render/shader.h

@@ -47,6 +47,7 @@ class Shader {
 public:
 	/* name */
 	string name;
+	int pass_id;
 
 	/* shader graph */
 	ShaderGraph *graph;

intern/cycles/render/svm.cpp

@@ -58,7 +58,7 @@ void SVMShaderManager::device_update(Device *device, DeviceScene *dscene, Scene
 	}
 	
 	bool sunsky_done = false;
-	bool use_multi_closure = device->support_full_kernel();
+	bool use_multi_closure = device->info.advanced_shading;
 
 	for(i = 0; i < scene->shaders.size(); i++) {
 		Shader *shader = scene->shaders[i];

intern/cycles/util/util_progress.h

@@ -96,7 +96,7 @@ public:
 	{
 		thread_scoped_lock lock(progress_mutex);
 
-		start_time = start_time;
+		start_time = start_time_;
 	}
 
 	void set_sample(int sample_, double sample_time_)

intern/cycles/util/util_types.h

@@ -277,6 +277,11 @@ __device float4 make_float4(float x, float y, float z, float w)
 	return a;
 }
 
+__device int align_up(int offset, int alignment)
+{
+	return (offset + alignment - 1) & ~(alignment - 1);
+}
+
 #endif
 
 CCL_NAMESPACE_END

intern/ffmpeg/ffmpeg_compat.h

@@ -35,6 +35,7 @@
 
 #include <libavcodec/avcodec.h>
 #include <libavutil/rational.h>
+#include <libavutil/opt.h>
 
 #if (LIBAVFORMAT_VERSION_MAJOR > 52) || ((LIBAVFORMAT_VERSION_MAJOR >= 52) && (LIBAVFORMAT_VERSION_MINOR >= 101))
 #define FFMPEG_HAVE_PARSE_UTILS 1

intern/ghost/GHOST_ISystem.h

@@ -271,6 +271,15 @@ public:
 	 */
 	virtual GHOST_TSuccess beginFullScreen(const GHOST_DisplaySetting& setting, GHOST_IWindow** window,
 		const bool stereoVisual, const GHOST_TUns16 numOfAASamples=0) = 0;
+	
+	/**
+	 * Updates the resolution while in fullscreen mode.
+	 * @param setting	The new setting of the display.
+	 * @param window	Window displayed in full screen.
+	 *
+	 * @return	Indication of success.
+	 */
+	virtual GHOST_TSuccess updateFullScreen(const GHOST_DisplaySetting& setting, GHOST_IWindow** window) = 0;
 
 	/**
 	 * Ends full screen mode.

intern/ghost/intern/GHOST_NDOFManager.cpp

@@ -77,6 +77,12 @@ static const char* ndof_button_names[] = {
 	"NDOF_BUTTON_8",
 	"NDOF_BUTTON_9",
 	"NDOF_BUTTON_10",
+	// more general-purpose buttons
+	"NDOF_BUTTON_A",
+	"NDOF_BUTTON_B",
+	"NDOF_BUTTON_C",
+	// the end
+	"NDOF_BUTTON_LAST"
 };
 #endif
 
@@ -169,7 +175,7 @@ static const NDOF_ButtonT SpaceMousePro_HID_map[] = {
 };
 
 /* this is the older SpacePilot (sans Pro)
- * thanks to polosson for the info in this table */
+ * thanks to polosson for info about this device */
 static const NDOF_ButtonT SpacePilot_HID_map[] = {
 	NDOF_BUTTON_1,
 	NDOF_BUTTON_2,
@@ -194,6 +200,23 @@ static const NDOF_ButtonT SpacePilot_HID_map[] = {
 	NDOF_BUTTON_NONE // the CONFIG button -- what does it do?
 };
 
+/* this is the older Spaceball 5000 USB
+ * thanks to Tehrasha Darkon for info about this device */
+static const NDOF_ButtonT Spaceball5000_HID_map[] = {
+	NDOF_BUTTON_1,
+	NDOF_BUTTON_2,
+	NDOF_BUTTON_3,
+	NDOF_BUTTON_4,
+	NDOF_BUTTON_5,
+	NDOF_BUTTON_6,
+	NDOF_BUTTON_7,
+	NDOF_BUTTON_8,
+	NDOF_BUTTON_9,
+	NDOF_BUTTON_A,
+	NDOF_BUTTON_B,
+	NDOF_BUTTON_C
+};
+
 GHOST_NDOFManager::GHOST_NDOFManager(GHOST_System& sys)
 	: m_system(sys)
 	, m_deviceType(NDOF_UnknownDevice) // each platform has its own device detection code
@@ -237,12 +260,12 @@ bool GHOST_NDOFManager::setDevice(unsigned short vendor_id, unsigned short produ
 					m_buttonCount = 15;
 					break;
 				case 0xC629:
-					puts("ndof: using SpacePilotPro");
+					puts("ndof: using SpacePilot Pro");
 					m_deviceType = NDOF_SpacePilotPro;
 					m_buttonCount = 31;
 					break;
 				case 0xC62B:
-					puts("ndof: using SpaceMousePro");
+					puts("ndof: using SpaceMouse Pro");
 					m_deviceType = NDOF_SpaceMousePro;
 					m_buttonCount = 27;
 					// ^^ actually has 15 buttons, but their HID codes range from 0 to 26
@@ -255,6 +278,12 @@ bool GHOST_NDOFManager::setDevice(unsigned short vendor_id, unsigned short produ
 					m_buttonCount = 21;
 					break;
 
+				case 0xC621:
+					puts("ndof: using Spaceball 5000");
+					m_deviceType = NDOF_Spaceball5000;
+					m_buttonCount = 12;
+					break;
+
 				case 0xC623:
 					puts("ndof: SpaceTraveler not supported, please file a bug report");
 					m_buttonCount = 8;
@@ -385,8 +414,21 @@ void GHOST_NDOFManager::updateButton(int button_number, bool press, GHOST_TUns64
 				default: sendButtonEvent(SpacePilot_HID_map[button_number], press, time, window);
 			}
 			break;
+		case NDOF_Spaceball5000:
+			// has no special 'keyboard' buttons
+			sendButtonEvent(Spaceball5000_HID_map[button_number], press, time, window);
+			break;
 		case NDOF_UnknownDevice:
-			printf("ndof: button %d on unknown device (ignoring)\n", button_number);
+			printf("ndof: button %d on unknown device (", button_number);
+			// map to the 'general purpose' buttons
+			// this is mainly for old serial devices
+			if (button_number < NDOF_BUTTON_LAST - NDOF_BUTTON_1) {
+				printf("sending)\n");
+				sendButtonEvent((NDOF_ButtonT)(NDOF_BUTTON_1 + button_number), press, time, window);
+			}
+			else {
+				printf("discarding)\n");
+			}
 	}
 
 	int mask = 1 << button_number;

intern/ghost/intern/GHOST_NDOFManager.h

@@ -40,7 +40,8 @@ typedef enum {
 	NDOF_SpaceMousePro,
 
 	// older devices
-	NDOF_SpacePilot
+	NDOF_SpacePilot,
+	NDOF_Spaceball5000
 
 	} NDOF_DeviceT;
 
@@ -87,7 +88,12 @@ typedef enum {
 	NDOF_BUTTON_8,
 	NDOF_BUTTON_9,
 	NDOF_BUTTON_10,
-
+	// more general-purpose buttons
+	NDOF_BUTTON_A,
+	NDOF_BUTTON_B,
+	NDOF_BUTTON_C,
+	// the end
+	NDOF_BUTTON_LAST
 	} NDOF_ButtonT;
 
 class GHOST_NDOFManager

intern/ghost/intern/GHOST_System.cpp

@@ -168,6 +168,19 @@ GHOST_TSuccess GHOST_System::beginFullScreen(const GHOST_DisplaySetting& setting
 }
 
 
+GHOST_TSuccess GHOST_System::updateFullScreen(const GHOST_DisplaySetting& setting, GHOST_IWindow** window)
+{
+	GHOST_TSuccess success = GHOST_kFailure;
+	GHOST_ASSERT(m_windowManager, "GHOST_System::updateFullScreen(): invalid window manager");
+	if(m_displayManager) {
+		if (m_windowManager->getFullScreen()) {
+			success = m_displayManager->setCurrentDisplaySetting(GHOST_DisplayManager::kMainDisplay, setting);
+		}
+	}
+
+	return success;
+}
+
 GHOST_TSuccess GHOST_System::endFullScreen(void)
 {
 	GHOST_TSuccess success = GHOST_kFailure;

intern/ghost/intern/GHOST_System.h

@@ -145,6 +145,15 @@ public:
 	 */
 	virtual GHOST_TSuccess beginFullScreen(const GHOST_DisplaySetting& setting, GHOST_IWindow** window,
 		const bool stereoVisual, const GHOST_TUns16 numOfAASamples=0);
+		
+	/**
+	 * Updates the resolution while in fullscreen mode.
+	 * @param setting	The new setting of the display.
+	 * @param window	Window displayed in full screen.
+	 *
+	 * @return	Indication of success.
+	 */
+	virtual GHOST_TSuccess updateFullScreen(const GHOST_DisplaySetting& setting, GHOST_IWindow** window);
 
 	/**
 	 * Ends full screen mode.

release/scripts/presets/keyconfig/maya.py

@@ -219,7 +219,8 @@ kmi.properties.value_1 = 'DISABLED'
 kmi.properties.value_2 = 'ENABLED'
 kmi = km.keymap_items.new('view3d.game_start', 'P', 'PRESS')
 kmi = km.keymap_items.new('object.select_all', 'A', 'PRESS')
-kmi = km.keymap_items.new('object.select_all', 'I', 'PRESS', ctrl=True).action = 'INVERT'
+kmi = km.keymap_items.new('object.select_all', 'I', 'PRESS', ctrl=True)
+kmi.properties.action = 'INVERT'
 kmi = km.keymap_items.new('object.select_linked', 'L', 'PRESS', shift=True)
 kmi = km.keymap_items.new('object.select_grouped', 'G', 'PRESS', shift=True)
 kmi = km.keymap_items.new('object.select_mirror', 'M', 'PRESS', shift=True, ctrl=True)

release/scripts/presets/tracking_settings/blurry_footage.py

@@ -9,3 +9,6 @@ settings.default_search_size = 100
 settings.default_frames_limit = 0
 settings.default_pattern_match = 'PREV_FRAME'
 settings.default_margin = 0
+settings.use_default_red_channel = True
+settings.use_default_green_channel = True
+settings.use_default_blue_channel = True

release/scripts/presets/tracking_settings/default.py

@@ -9,3 +9,6 @@ settings.default_search_size = 61
 settings.default_frames_limit = 0
 settings.default_pattern_match = 'KEYFRAME'
 settings.default_margin = 0
+settings.use_default_red_channel = True
+settings.use_default_green_channel = True
+settings.use_default_blue_channel = True

release/scripts/presets/tracking_settings/fast_motion.py

@@ -9,3 +9,6 @@ settings.default_search_size = 300
 settings.default_frames_limit = 0
 settings.default_pattern_match = 'PREV_FRAME'
 settings.default_margin = 5
+settings.use_default_red_channel = True
+settings.use_default_green_channel = True
+settings.use_default_blue_channel = True

release/scripts/startup/bl_operators/clip.py

@@ -90,6 +90,34 @@ def CLIP_track_view_selected(sc, track):
     return False
 
 
+def CLIP_default_settings_from_track(clip, track):
+    settings = clip.tracking.settings
+
+    width = clip.size[0]
+    height = clip.size[1]
+
+    pattern = track.pattern_max - track.pattern_min
+    search = track.search_max - track.search_min
+
+    pattern[0] = pattern[0] * clip.size[0]
+    pattern[1] = pattern[1] * clip.size[1]
+
+    search[0] = search[0] * clip.size[0]
+    search[1] = search[1] * clip.size[1]
+
+    settings.default_tracker = track.tracker
+    settings.default_pyramid_levels = track.pyramid_levels
+    settings.default_correlation_min = track.correlation_min
+    settings.default_pattern_size = max(pattern[0], pattern[1])
+    settings.default_search_size = max(search[0], search[1])
+    settings.default_frames_limit = track.frames_limit
+    settings.default_pattern_match = track.pattern_match
+    settings.default_margin = track.margin
+    settings.use_default_red_channel = track.use_red_channel
+    settings.use_default_green_channel = track.use_green_channel
+    settings.use_default_blue_channel = track.use_blue_channel
+
+
 class CLIP_OT_track_to_empty(Operator):
     """Create an Empty object which will be copying movement of active track"""
 
@@ -805,3 +833,29 @@ class CLIP_OT_setup_tracking_scene(Operator):
         self._setupObjects(context)
 
         return {'FINISHED'}
+
+class CLIP_OT_track_settings_as_default(Operator):
+    """Copy trackign settings from active track to default settings"""
+
+    bl_idname = "clip.track_settings_as_default"
+    bl_label = "Track Settings As Default"
+    bl_options = {'UNDO', 'REGISTER'}
+
+    @classmethod
+    def poll(cls, context):
+        sc = context.space_data
+
+        if sc.type != 'CLIP_EDITOR':
+            return False
+
+        clip = sc.clip
+
+        return clip and clip.tracking.tracks.active
+
+    def execute(self, context):
+        sc = context.space_data
+        clip = sc.clip
+
+        CLIP_default_settings_from_track(clip, clip.tracking.tracks.active)
+
+        return {'FINISHED'}

release/scripts/startup/bl_operators/presets.py

@@ -409,7 +409,10 @@ class AddPresetTrackingSettings(AddPresetBase, Operator):
         "settings.default_search_size",
         "settings.default_frames_limit",
         "settings.default_pattern_match",
-        "settings.default_margin"
+        "settings.default_margin",
+        "settings.use_default_red_channel",
+        "settings.use_default_green_channel",
+        "settings.use_default_blue_channel"
     ]
 
     preset_subdir = "tracking_settings"

release/scripts/startup/bl_ui/properties_object_constraint.py

@@ -434,25 +434,28 @@ class ConstraintButtonsPanel():
     def ACTION(self, context, layout, con):
         self.target_template(layout, con)
 
-        layout.prop(con, "action")
-
-        layout.prop(con, "transform_channel")
-
         split = layout.split()
-
+        
-        col = split.column(align=True)
+        col = split.column()
-        col.label(text="Action Length:")
+        col.label(text="From Target:")
-        col.prop(con, "frame_start", text="Start")
+        col.prop(con, "transform_channel", text="")
-        col.prop(con, "frame_end", text="End")
+        col.prop(con, "target_space", text="")
-
+        
+        col = split.column()
+        col.label(text="To Action:")
+        col.prop(con, "action", text="")
+        
+        split = layout.split()
+        
         col = split.column(align=True)
         col.label(text="Target Range:")
         col.prop(con, "min", text="Min")
         col.prop(con, "max", text="Max")
-
+        
-        row = layout.row()
+        col = split.column(align=True)
-        row.label(text="Convert:")
+        col.label(text="Action Range:")
-        row.prop(con, "target_space", text="")
+        col.prop(con, "frame_start", text="Start")
+        col.prop(con, "frame_end", text="End")
 
     def LOCKED_TRACK(self, context, layout, con):
         self.target_template(layout, con)

release/scripts/startup/bl_ui/properties_particle.py

@@ -911,7 +911,7 @@ class PARTICLE_PT_render(ParticleButtonsPanel, Panel):
                 col.prop_search(psys, "billboard_time_index_uv", ob.data, "uv_textures")
 
             split = layout.split(percentage=0.33)
-            split.label(text="Split uv's:")
+            split.label(text="Split UVs:")
             split.prop(part, "billboard_uv_split", text="Number of splits")
 
             if psys:

release/scripts/startup/bl_ui/properties_physics_cloth.py

@@ -87,6 +87,7 @@ class PHYSICS_PT_cloth(PhysicButtonsPanel, Panel):
         col.label(text="Damping:")
         col.prop(cloth, "spring_damping", text="Spring")
         col.prop(cloth, "air_damping", text="Air")
+        col.prop(cloth, "vel_damping", text="Velocity")
 
         col.prop(cloth, "use_pin_cloth", text="Pinning")
         sub = col.column()

release/scripts/startup/bl_ui/properties_physics_fluid.py

@@ -79,6 +79,7 @@ class PHYSICS_PT_fluid(PhysicButtonsPanel, Panel):
             sub = col.column(align=True)
             sub.prop(fluid, "start_time", text="Start")
             sub.prop(fluid, "end_time", text="End")
+            col.prop(fluid, "simulation_rate", text="Speed")
 
             col = split.column()
             col.label()
@@ -230,6 +231,10 @@ class PHYSICS_PT_domain_gravity(PhysicButtonsPanel, Panel):
         if fluid.viscosity_preset == 'MANUAL':
             sub.prop(fluid, "viscosity_base", text="Base")
             sub.prop(fluid, "viscosity_exponent", text="Exponent", slider=True)
+        else:
+            # just for padding to prevent jumping around
+            sub.separator()
+            sub.separator()
 
         col.label(text="Optimization:")
         col.prop(fluid, "grid_levels", slider=True)

release/scripts/startup/bl_ui/space_clip.py

@@ -127,6 +127,13 @@ class CLIP_PT_tools_marker(Panel):
 
             col.separator()
 
+            row = col.row(align=True)
+            row.prop(settings, "use_default_red_channel", text="R", toggle=True)
+            row.prop(settings, "use_default_green_channel", text="G", toggle=True)
+            row.prop(settings, "use_default_blue_channel", text="B", toggle=True)
+
+            col.separator()
+
             sub = col.column(align=True)
             sub.prop(settings, "default_pattern_size")
             sub.prop(settings, "default_search_size")
@@ -147,6 +154,9 @@ class CLIP_PT_tools_marker(Panel):
             col.label(text="Match:")
             col.prop(settings, "default_pattern_match", text="")
 
+            col.separator()
+            col.operator('clip.track_settings_as_default', text="Copy From Active Track")
+
 
 class CLIP_PT_tools_tracking(Panel):
     bl_space_type = 'CLIP_EDITOR'
@@ -172,6 +182,7 @@ class CLIP_PT_tools_tracking(Panel):
         props.backwards = True
         props.sequence = True
         props = row.operator("clip.track_markers", text="", icon='PLAY')
+        props.backwards = False
         props.sequence = True
         row.operator("clip.track_markers", text="", icon='FRAME_NEXT')
 
@@ -558,6 +569,9 @@ class CLIP_PT_display(Panel):
 
         col.prop(sc, "lock_selection")
 
+        if sc.view == 'GRAPH':
+            col.prop(sc, "lock_time_cursor")
+
         clip = sc.clip
         if clip:
             col.label(text="Display Aspect Ratio:")

release/scripts/startup/bl_ui/space_logic.py

@@ -68,7 +68,7 @@ class LOGIC_PT_properties(Panel):
             row = box.row()
             row.prop(prop, "name", text="")
             row.prop(prop, "type", text="")
-            row.prop(prop, "value", text="", toggle=True)  # we don't care about the type. rna will display correctly
+            row.prop(prop, "value", text="")
             row.prop(prop, "show_debug", text="", toggle=True, icon='INFO')
             row.operator("object.game_property_remove", text="", icon='X', emboss=False).index = i
 

release/scripts/startup/bl_ui/space_node.py

@@ -58,6 +58,11 @@ class NODE_HT_header(Header):
                 # Don't show "Use Nodes" Button when Engine is BI for Lamps
                 if snode_id and not (scene.render.use_shading_nodes == 0 and ob.type == 'LAMP'):
                     layout.prop(snode_id, "use_nodes")
+                    
+            if snode.shader_type == 'WORLD':
+                layout.template_ID(scene, "world", new="world.new")
+                if snode_id:
+                    layout.prop(snode_id, "use_nodes")
 
         elif snode.tree_type == 'TEXTURE':
             layout.prop(snode, "texture_type", text="", expand=True)

release/scripts/startup/bl_ui/space_view3d_toolbar.py

@@ -526,9 +526,9 @@ class VIEW3D_PT_tools_brush(Panel, View3DPaintPanel):
 
                 if brush.use_space and tool != 'SMOOTH':
                     if brush.use_space_atten:
-                        row.prop(brush, "use_space_atten", toggle=True, icon='LOCKED')
+                        row.prop(brush, "use_space_atten", toggle=True, text="", icon='LOCKED')
                     else:
-                        row.prop(brush, "use_space_atten", toggle=True, icon='UNLOCKED')
+                        row.prop(brush, "use_space_atten", toggle=True, text="", icon='UNLOCKED')
 
                 self.prop_unified_strength(row, context, brush, "strength", text="Strength")
                 self.prop_unified_strength(row, context, brush, "use_pressure_strength")

source/blender/avi/intern/endian.c

@@ -42,6 +42,10 @@
 #include "endian.h"
 #include "avi_intern.h"
 
+#ifdef __BIG_ENDIAN__
+#include "MEM_guardedalloc.h"
+#endif
+
 #ifdef __BIG_ENDIAN__
 static void invert (int *num)
 {

source/blender/blenkernel/BKE_DerivedMesh.h

@@ -538,6 +538,15 @@ int sculpt_get_deform_matrices(struct Scene *scene, struct Object *ob,
 								float (**deformmats)[3][3], float (**deformcos)[3]);
 
 void weight_to_rgb(float r_rgb[3], const float weight);
+/* Update the weight MCOL preview layer.
+ * If weights are NULL, use object's active vgroup(s).
+ * Else, weights must be an array of weight float values.
+ *     If indices is NULL, it must be of numVerts length.
+ *     Else, it must be of num length, as indices, which contains vertices' idx to apply weights to.
+ *         (other vertices are assumed zero weight).
+ */
+void DM_update_weight_mcol(struct Object *ob, struct DerivedMesh *dm, int const draw_flag,
+                           float *weights, int num, const int *indices);
 
 /* convert layers requested by a GLSL material to actually available layers in
  * the DerivedMesh, with both a pointer for arrays and an offset for editmesh */

source/blender/blenkernel/BKE_armature.h

@@ -86,7 +86,7 @@ int bone_autoside_name (char name[64], int strip_number, short axis, float head,
 
 struct Bone *get_named_bone (struct bArmature *arm, const char *name);
 
-float distfactor_to_bone (float vec[3], float b1[3], float b2[3], float rad1, float rad2, float rdist);
+float distfactor_to_bone(const float vec[3], const float b1[3], const float b2[3], float rad1, float rad2, float rdist);
 
 void where_is_armature (struct bArmature *arm);
 void where_is_armature_bone(struct Bone *bone, struct Bone *prevbone);

source/blender/blenkernel/BKE_blender.h

@@ -42,7 +42,7 @@ extern "C" {
  * and keep comment above the defines.
  * Use STRINGIFY() rather than defining with quotes */
 #define BLENDER_VERSION			261
-#define BLENDER_SUBVERSION		3
+#define BLENDER_SUBVERSION		4
 
 #define BLENDER_MINVERSION		250
 #define BLENDER_MINSUBVERSION	0
@@ -51,7 +51,7 @@ extern "C" {
 		/* can be left blank, otherwise a,b,c... etc with no quotes */
 #define BLENDER_VERSION_CHAR	
 		/* alpha/beta/rc/release, docs use this */
-#define BLENDER_VERSION_CYCLE	alpha
+#define BLENDER_VERSION_CYCLE	beta
 
 extern char versionstr[]; /* from blender.c */
 

source/blender/blenkernel/BKE_cloth.h

@@ -216,7 +216,7 @@ void clmdSetInterruptCallBack ( int ( *f ) ( void ) );
 void cloth_free_modifier_extern ( struct ClothModifierData *clmd );
 void cloth_free_modifier ( struct ClothModifierData *clmd );
 void cloth_init ( struct ClothModifierData *clmd );
-struct DerivedMesh *clothModifier_do ( struct ClothModifierData *clmd, struct Scene *scene, struct Object *ob, struct DerivedMesh *dm);
+void clothModifier_do ( struct ClothModifierData *clmd, struct Scene *scene, struct Object *ob, struct DerivedMesh *dm, float (*vertexCos)[3]);
 
 void cloth_update_normals ( ClothVertex *verts, int nVerts, struct MFace *face, int totface );
 int cloth_uses_vgroup(struct ClothModifierData *clmd);

source/blender/blenkernel/BKE_customdata.h

@@ -67,7 +67,7 @@ extern const CustomDataMask CD_MASK_FACECORNERS;
 #define CD_DUPLICATE 4  /* do a full copy of all layers, only allowed if source
 						   has same number of elements */
 
-#define CD_TYPE_AS_MASK(_type) (CustomDataMask)(1 << (CustomDataMask)(_type))
+#define CD_TYPE_AS_MASK(_type) (CustomDataMask)((CustomDataMask)1 << (CustomDataMask)(_type))
 
 /* initialises a CustomData object with the same layer setup as source.
  * mask is a bitfield where (mask & (1 << (layer type))) indicates

source/blender/blenkernel/BKE_modifier.h

@@ -99,7 +99,10 @@ typedef enum {
 	eModifierTypeFlag_Single = (1<<7),
 
 	/* Some modifier can't be added manually by user */
-	eModifierTypeFlag_NoUserAdd = (1<<8)
+	eModifierTypeFlag_NoUserAdd = (1<<8),
+
+	/* For modifiers that use CD_WEIGHT_MCOL for preview. */
+	eModifierTypeFlag_UsesPreview = (1<<9)
 } ModifierTypeFlag;
 
 typedef void (*ObjectWalkFunc)(void *userData, struct Object *ob, struct Object **obpoin);
@@ -323,6 +326,7 @@ void          modifier_setError(struct ModifierData *md, const char *format, ...
 __attribute__ ((format (printf, 2, 3)))
 #endif
 ;
+int           modifier_isPreview(struct ModifierData *md);
 
 void          modifiers_foreachObjectLink(struct Object *ob,
 										  ObjectWalkFunc walk,
@@ -349,6 +353,7 @@ struct Object *modifiers_isDeformedByLattice(struct Object *ob);
 int           modifiers_usesArmature(struct Object *ob, struct bArmature *arm);
 int           modifiers_isCorrectableDeformed(struct Object *ob);
 void          modifier_freeTemporaryData(struct ModifierData *md);
+int           modifiers_isPreview(struct Object *ob);
 
 int           modifiers_indexInObject(struct Object *ob, struct ModifierData *md);
 
@@ -362,6 +367,9 @@ struct LinkNode *modifiers_calcDataMasks(struct Scene *scene,
 										 struct ModifierData *md,
 										 CustomDataMask dataMask,
 										 int required_mode);
+struct ModifierData *modifiers_getLastPreview(struct Scene *scene,
+                                              struct ModifierData *md,
+                                              int required_mode);
 struct ModifierData  *modifiers_getVirtualModifierList(struct Object *ob);
 
 /* ensure modifier correctness when changing ob->data */

source/blender/blenkernel/BKE_node.h

@@ -523,6 +523,7 @@ struct ShadeResult;
 #define SH_NODE_VOLUME_ISOTROPIC		162
 #define SH_NODE_GAMMA				163
 #define SH_NODE_TEX_CHECKER			164
+#define SH_NODE_BRIGHTCONTRAST			165
 
 /* custom defines options for Material node */
 #define SH_NODE_MAT_DIFF   1
@@ -556,25 +557,34 @@ void			ntreeGPUMaterialNodes(struct bNodeTree *ntree, struct GPUMaterial *mat);
 /* ************** COMPOSITE NODES *************** */
 
 /* output socket defines */
-#define RRES_OUT_IMAGE		0
+#define RRES_OUT_IMAGE				0
-#define RRES_OUT_ALPHA		1
+#define RRES_OUT_ALPHA				1
-#define RRES_OUT_Z			2
+#define RRES_OUT_Z					2
-#define RRES_OUT_NORMAL		3
+#define RRES_OUT_NORMAL				3
-#define RRES_OUT_UV			4
+#define RRES_OUT_UV					4
-#define RRES_OUT_VEC		5
+#define RRES_OUT_VEC				5
-#define RRES_OUT_RGBA		6
+#define RRES_OUT_RGBA				6
-#define RRES_OUT_DIFF		7
+#define RRES_OUT_DIFF				7
-#define RRES_OUT_SPEC		8
+#define RRES_OUT_SPEC				8
-#define RRES_OUT_SHADOW		9
+#define RRES_OUT_SHADOW				9
-#define RRES_OUT_AO			10
+#define RRES_OUT_AO					10
-#define RRES_OUT_REFLECT	11
+#define RRES_OUT_REFLECT			11
-#define RRES_OUT_REFRACT	12
+#define RRES_OUT_REFRACT			12
-#define RRES_OUT_INDIRECT	13
+#define RRES_OUT_INDIRECT			13
-#define RRES_OUT_INDEXOB	14
+#define RRES_OUT_INDEXOB			14
-#define RRES_OUT_INDEXMA	15
+#define RRES_OUT_INDEXMA			15
-#define RRES_OUT_MIST		16
+#define RRES_OUT_MIST				16
-#define RRES_OUT_EMIT		17
+#define RRES_OUT_EMIT				17
-#define RRES_OUT_ENV		18
+#define RRES_OUT_ENV				18
+#define RRES_OUT_DIFF_DIRECT		19
+#define RRES_OUT_DIFF_INDIRECT		20
+#define RRES_OUT_DIFF_COLOR			21
+#define RRES_OUT_GLOSSY_DIRECT		22
+#define RRES_OUT_GLOSSY_INDIRECT	23
+#define RRES_OUT_GLOSSY_COLOR		24
+#define RRES_OUT_TRANSM_DIRECT		25
+#define RRES_OUT_TRANSM_INDIRECT	26
+#define RRES_OUT_TRANSM_COLOR		27
 
 /* note: types are needed to restore callbacks, don't change values */
 #define CMP_NODE_VIEWER		201

source/blender/blenkernel/BKE_tracking.h

@@ -61,6 +61,7 @@ struct MovieTrackingMarker *BKE_tracking_get_marker(struct MovieTrackingTrack *t
 struct MovieTrackingMarker *BKE_tracking_ensure_marker(struct MovieTrackingTrack *track, int framenr);
 struct MovieTrackingMarker *BKE_tracking_exact_marker(struct MovieTrackingTrack *track, int framenr);
 int BKE_tracking_has_marker(struct MovieTrackingTrack *track, int framenr);
+int BKE_tracking_has_enabled_marker(struct MovieTrackingTrack *track, int framenr);
 
 void BKE_tracking_free_track(struct MovieTrackingTrack *track);
 

source/blender/blenkernel/intern/DerivedMesh.c

@@ -722,24 +722,23 @@ void vDM_ColorBand_store(ColorBand *coba)
  * note that we could save some memory and allocate RGB only but then we'd need to
  * re-arrange the colors when copying to the face since MCol has odd ordering,
  * so leave this as is - campbell */
-static unsigned char *calc_weightpaint_vert_array(Object *ob, int const draw_flag, ColorBand *coba)
+static unsigned char *calc_weightpaint_vert_array(Object *ob, DerivedMesh *dm, int const draw_flag, ColorBand *coba)
 {
-	Mesh *me = ob->data;
+	MDeformVert *dv = DM_get_vert_data_layer(dm, CD_MDEFORMVERT);
-	unsigned char *wtcol_v = MEM_mallocN (sizeof(unsigned char) * me->totvert * 4, "weightmap_v");
+	int numVerts = dm->getNumVerts(dm);
+	unsigned char *wtcol_v = MEM_mallocN (sizeof(unsigned char) * numVerts * 4, "weightmap_v");
 
-	if (me->dvert) {
+	if (dv) {
 		unsigned char *wc = wtcol_v;
-		MDeformVert *dv= me->dvert;
 		unsigned int i;
 
-		/* varisbles for multipaint */
+		/* variables for multipaint */
 		const int defbase_tot = BLI_countlist(&ob->defbase);
 		const int defbase_act = ob->actdef-1;
 		char *dg_flags = MEM_mallocN(defbase_tot * sizeof(char), __func__);
 		const int selected = get_selected_defgroups(ob, dg_flags, defbase_tot);
-		/* const int unselected = defbase_tot - selected; */ /* UNUSED */
 
-		for (i = me->totvert; i != 0; i--, wc += 4, dv++) {
+		for (i = numVerts; i != 0; i--, wc += 4, dv++) {
 			calc_weightpaint_vert_color(wc, dv, coba, defbase_tot, defbase_act, dg_flags, selected, draw_flag);
 		}
 
@@ -748,48 +747,94 @@ static unsigned char *calc_weightpaint_vert_array(Object *ob, int const draw_fla
 	else {
 		int col_i;
 		weightpaint_color((unsigned char *)&col_i, coba, 0.0f);
-		fill_vn_i((int *)wtcol_v, me->totvert, col_i);
+		fill_vn_i((int *)wtcol_v, numVerts, col_i);
 	}
 
 	return wtcol_v;
 }
 
-static void add_weight_mcol_dm(Object *ob, DerivedMesh *dm, int const draw_flag)
+/* return an array of vertex weight colors from given weights, caller must free.
+ *
+ * note that we could save some memory and allocate RGB only but then we'd need to
+ * re-arrange the colors when copying to the face since MCol has odd ordering,
+ * so leave this as is - campbell */
+static unsigned char *calc_colors_from_weights_array(const int num, float *weights)
+{
+	unsigned char *wtcol_v = MEM_mallocN(sizeof(unsigned char) * num * 4, "weightmap_v");
+	unsigned char *wc = wtcol_v;
+	int i;
+
+	for (i = 0; i < num; i++, wc += 4, weights++)
+		weightpaint_color((unsigned char *) wc, NULL, *weights);
+
+	return wtcol_v;
+}
+
+void DM_update_weight_mcol(Object *ob, DerivedMesh *dm, int const draw_flag,
+                           float *weights, int num, const int *indices)
 {
 	ColorBand *coba= stored_cb;	/* warning, not a local var */
 
-	Mesh *me = ob->data;
+	MFace *mf = dm->getFaceArray(dm);
-	unsigned char *wtcol_v = calc_weightpaint_vert_array(ob, draw_flag, coba);
+	int numFaces = dm->getNumFaces(dm);
-	unsigned char *wtcol_f = MEM_mallocN (sizeof(unsigned char) * me->totface*4*4, "weightmap_f");
+	int numVerts = dm->getNumVerts(dm);
-	unsigned char *wtcol_f_step = wtcol_f;
+	unsigned char *wtcol_v;
-
+	unsigned char *wtcol_f = dm->getFaceDataArray(dm, CD_WEIGHT_MCOL);
-	MFace *mf = me->mface;
 	int i;
 
-	for (i=0; i<me->totface; i++, mf++, wtcol_f_step += (4 * 4)) {
+	/* If no CD_WEIGHT_MCOL existed yet, add a new one! */
+	if (!wtcol_f)
+		wtcol_f = CustomData_add_layer(&dm->faceData, CD_WEIGHT_MCOL, CD_CALLOC, NULL, numFaces);
+
+	if (wtcol_f) {
+		unsigned char *wtcol_f_step = wtcol_f;
+
+		/* Weights are given by caller. */
+		if (weights) {
+			float *w = weights;
+			/* If indices is not NULL, it means we do not have weights for all vertices,
+			 * so we must create them (and set them to zero)... */
+			if(indices) {
+				w = MEM_callocN(sizeof(float)*numVerts, "Temp weight array DM_update_weight_mcol");
+				i = num;
+				while(i--)
+					w[indices[i]] = weights[i];
+			}
+
+			/* Convert float weights to colors. */
+			wtcol_v = calc_colors_from_weights_array(numVerts, w);
+
+			if(indices)
+				MEM_freeN(w);
+		}
+
+		/* No weights given, take them from active vgroup(s). */
+		else
+			wtcol_v = calc_weightpaint_vert_array(ob, dm, draw_flag, coba);
+
+		/* Now copy colors in all face verts. */
+		for (i = 0; i < numFaces; i++, mf++, wtcol_f_step += (4 * 4)) {
 #if 0
-		unsigned int fidx= mf->v4 ? 3:2;
+			unsigned int fidx= mf->v4 ? 3:2;
 
 #else	/* better zero out triangles 4th component. else valgrind complains when the buffer's copied */
-		unsigned int fidx;
+			unsigned int fidx;
-		if (mf->v4) {
+			if (mf->v4) {
-			fidx = 3;
+				fidx = 3;
-		}
+			}
-		else {
+			else {
-			fidx = 2;
+				fidx = 2;
-			*(int *)(&wtcol_f_step[3 * 4]) = 0;
+				*(int *)(&wtcol_f_step[3 * 4]) = 0;
-		}
+			}
 #endif
 
-		do {
+			do {
-			copy_v4_v4_char((char *)&wtcol_f_step[fidx * 4],
+				copy_v4_v4_char((char *)&wtcol_f_step[fidx * 4],
-			                (char *)&wtcol_v[4 * (*(&mf->v1 + fidx))]);
+						        (char *)&wtcol_v[4 * (*(&mf->v1 + fidx))]);
-		} while (fidx--);
+			} while (fidx--);
+		}
+		MEM_freeN(wtcol_v);
 	}
-
-	MEM_freeN(wtcol_v);
-
-	CustomData_add_layer(&dm->faceData, CD_WEIGHT_MCOL, CD_ASSIGN, wtcol_f, dm->numFaceData);
 }
 
 /* new value for useDeform -1  (hack for the gameengine):
@@ -803,7 +848,7 @@ static void mesh_calc_modifiers(Scene *scene, Object *ob, float (*inputVertexCos
 								int needMapping, CustomDataMask dataMask, int index, int useCache)
 {
 	Mesh *me = ob->data;
-	ModifierData *firstmd, *md;
+	ModifierData *firstmd, *md, *previewmd = NULL;
 	LinkNode *datamasks, *curr;
 	CustomDataMask mask, nextmask, append_mask = 0;
 	float (*deformedVerts)[3] = NULL;
@@ -816,8 +861,17 @@ static void mesh_calc_modifiers(Scene *scene, Object *ob, float (*inputVertexCos
 	int has_multires = mmd != NULL, multires_applied = 0;
 	int sculpt_mode = ob->mode & OB_MODE_SCULPT && ob->sculpt;
 
-	int draw_flag= ((scene->toolsettings->multipaint ? CALC_WP_MULTIPAINT : 0) |
+	const int draw_flag= ((scene->toolsettings->multipaint ? CALC_WP_MULTIPAINT : 0) |
-	                (scene->toolsettings->auto_normalize ? CALC_WP_AUTO_NORMALIZE : 0));
+	                      (scene->toolsettings->auto_normalize ? CALC_WP_AUTO_NORMALIZE : 0));
+	/* Generic preview only in object mode! */
+	const int do_mod_mcol = (ob->mode == OB_MODE_OBJECT);
+#if 0 /* XXX Will re-enable this when we have global mod stack options. */
+	const int do_final_wmcol = (scene->toolsettings->weights_preview == WP_WPREVIEW_FINAL) && do_wmcol;
+#endif
+	const int do_final_wmcol = FALSE;
+	int do_init_wmcol = ((dataMask & CD_MASK_WEIGHT_MCOL) && (ob->mode & OB_MODE_WEIGHT_PAINT) && !do_final_wmcol);
+	/* XXX Same as above... For now, only weights preview in WPaint mode. */
+	const int do_mod_wmcol = do_init_wmcol;
 
 	if(mmd && !mmd->sculptlvl)
 		has_multires = 0;
@@ -842,6 +896,14 @@ static void mesh_calc_modifiers(Scene *scene, Object *ob, float (*inputVertexCos
 	datamasks = modifiers_calcDataMasks(scene, ob, md, dataMask, required_mode);
 	curr = datamasks;
 
+	if(do_mod_wmcol || do_mod_mcol) {
+		/* Find the last active modifier generating a preview, or NULL if none. */
+		/* XXX Currently, DPaint modifier just ignores this.
+		 *     Needs a stupid hack...
+		 *     The whole "modifier preview" thing has to be (re?)designed, anyway! */
+		previewmd = modifiers_getLastPreview(scene, md, required_mode);
+	}
+
 	if(deform_r) *deform_r = NULL;
 	*final_r = NULL;
 
@@ -997,8 +1059,8 @@ static void mesh_calc_modifiers(Scene *scene, Object *ob, float (*inputVertexCos
 					CDDM_calc_normals(dm);
 				}
 
-				if((dataMask & CD_MASK_WEIGHT_MCOL) && (ob->mode & OB_MODE_WEIGHT_PAINT))
+				if(do_init_wmcol)
-					add_weight_mcol_dm(ob, dm, draw_flag); 
+					DM_update_weight_mcol(ob, dm, draw_flag, NULL, 0, NULL);
 
 				/* Constructive modifiers need to have an origindex
 				 * otherwise they wont have anywhere to copy the data from.
@@ -1085,8 +1147,14 @@ static void mesh_calc_modifiers(Scene *scene, Object *ob, float (*inputVertexCos
 			}
 
 			/* in case of dynamic paint, make sure preview mask remains for following modifiers */
+			/* XXX Temp and hackish solution! */
 			if (md->type == eModifierType_DynamicPaint)
 				append_mask |= CD_MASK_WEIGHT_MCOL;
+			/* In case of active preview modifier, make sure preview mask remains for following modifiers. */
+			else if ((md == previewmd) && (do_mod_wmcol)) {
+				DM_update_weight_mcol(ob, dm, draw_flag, NULL, 0, NULL);
+				append_mask |= CD_MASK_WEIGHT_MCOL;
+			}
 		}
 
 		isPrevDeform= (mti->type == eModifierTypeType_OnlyDeform);
@@ -1114,10 +1182,19 @@ static void mesh_calc_modifiers(Scene *scene, Object *ob, float (*inputVertexCos
 		CDDM_apply_vert_coords(finaldm, deformedVerts);
 		CDDM_calc_normals(finaldm);
 
-		if((dataMask & CD_MASK_WEIGHT_MCOL) && (ob->mode & OB_MODE_WEIGHT_PAINT))
+#if 0 /* For later nice mod preview! */
-			add_weight_mcol_dm(ob, finaldm, draw_flag);
+		/* In case we need modified weights in CD_WEIGHT_MCOL, we have to re-compute it. */
+		if(do_final_wmcol)
+			DM_update_weight_mcol(ob, finaldm, draw_flag, NULL, 0, NULL);
+#endif
 	} else if(dm) {
 		finaldm = dm;
+
+#if 0 /* For later nice mod preview! */
+		/* In case we need modified weights in CD_WEIGHT_MCOL, we have to re-compute it. */
+		if(do_final_wmcol)
+			DM_update_weight_mcol(ob, finaldm, draw_flag, NULL, 0, NULL);
+#endif
 	} else {
 		finaldm = CDDM_from_mesh(me, ob);
 
@@ -1126,8 +1203,9 @@ static void mesh_calc_modifiers(Scene *scene, Object *ob, float (*inputVertexCos
 			CDDM_calc_normals(finaldm);
 		}
 
-		if((dataMask & CD_MASK_WEIGHT_MCOL) && (ob->mode & OB_MODE_WEIGHT_PAINT))
+		/* In this case, we should never have weight-modifying modifiers in stack... */
-			add_weight_mcol_dm(ob, finaldm, draw_flag);
+		if(do_init_wmcol)
+			DM_update_weight_mcol(ob, finaldm, draw_flag, NULL, 0, NULL);
 	}
 
 	/* add an orco layer if needed */

source/blender/blenkernel/intern/armature.c

@@ -665,7 +665,7 @@ static void b_bone_deform(bPoseChanDeform *pdef_info, Bone *bone, float *co, Dua
 }
 
 /* using vec with dist to bone b1 - b2 */
-float distfactor_to_bone (float vec[3], float b1[3], float b2[3], float rad1, float rad2, float rdist)
+float distfactor_to_bone(const float vec[3], const float b1[3], const float b2[3], float rad1, float rad2, float rdist)
 {
 	float dist=0.0f; 
 	float bdelta[3];
@@ -677,18 +677,18 @@ float distfactor_to_bone (float vec[3], float b1[3], float b2[3], float rad1, fl
 	
 	sub_v3_v3v3(pdelta, vec, b1);
 	
-	a = bdelta[0]*pdelta[0] + bdelta[1]*pdelta[1] + bdelta[2]*pdelta[2];
+	a = dot_v3v3(bdelta, pdelta);
-	hsqr = ((pdelta[0]*pdelta[0]) + (pdelta[1]*pdelta[1]) + (pdelta[2]*pdelta[2]));
+	hsqr = dot_v3v3(pdelta, pdelta);
 	
-	if (a < 0.0F){
+	if (a < 0.0f) {
 		/* If we're past the end of the bone, do a spherical field attenuation thing */
-		dist= ((b1[0]-vec[0])*(b1[0]-vec[0]) +(b1[1]-vec[1])*(b1[1]-vec[1]) +(b1[2]-vec[2])*(b1[2]-vec[2])) ;
+		dist = len_squared_v3v3(b1, vec);
 		rad= rad1;
 	}
-	else if (a > l){
+	else if (a > l) {
 		/* If we're past the end of the bone, do a spherical field attenuation thing */
-		dist= ((b2[0]-vec[0])*(b2[0]-vec[0]) +(b2[1]-vec[1])*(b2[1]-vec[1]) +(b2[2]-vec[2])*(b2[2]-vec[2])) ;
+		dist = len_squared_v3v3(b2, vec);
-		rad= rad2;
+		rad = rad2;
 	}
 	else {
 		dist= (hsqr - (a*a));
@@ -709,7 +709,7 @@ float distfactor_to_bone (float vec[3], float b1[3], float b2[3], float rad1, fl
 		if(rdist==0.0f || dist >= l) 
 			return 0.0f;
 		else {
-			a= (float)sqrt(dist)-rad;
+			a = sqrtf(dist)-rad;
 			return 1.0f-( a*a )/( rdist*rdist );
 		}
 	}

source/blender/blenkernel/intern/cloth.c

@@ -90,7 +90,7 @@ static CM_SOLVER_DEF	solvers [] =
 /* ********** cloth engine ******* */
 /* Prototypes for internal functions.
 */
-static void cloth_to_object (Object *ob,  ClothModifierData *clmd, DerivedMesh *dm);
+static void cloth_to_object (Object *ob,  ClothModifierData *clmd, float (*vertexCos)[3]);
 static void cloth_from_mesh ( ClothModifierData *clmd, DerivedMesh *dm );
 static int cloth_from_object(Object *ob, ClothModifierData *clmd, DerivedMesh *dm, float framenr, int first);
 static int cloth_build_springs ( ClothModifierData *clmd, DerivedMesh *dm );
@@ -130,6 +130,7 @@ void cloth_init ( ClothModifierData *clmd )
 	clmd->sim_parms->presets = 2; /* cotton as start setting */
 	clmd->sim_parms->timescale = 1.0f; /* speed factor, describes how fast cloth moves */
 	clmd->sim_parms->reset = 0;
+	clmd->sim_parms->vel_damping = 1.0f; /* 1.0 = no damping, 0.0 = fully dampened */
 	
 	clmd->coll_parms->self_friction = 5.0;
 	clmd->coll_parms->friction = 5.0;
@@ -428,9 +429,8 @@ static int do_step_cloth(Object *ob, ClothModifierData *clmd, DerivedMesh *resul
 /************************************************
  * clothModifier_do - main simulation function
 ************************************************/
-DerivedMesh *clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob, DerivedMesh *dm)
+void clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob, DerivedMesh *dm, float (*vertexCos)[3])
 {
-	DerivedMesh *result;
 	PointCache *cache;
 	PTCacheID pid;
 	float timescale;
@@ -440,20 +440,14 @@ DerivedMesh *clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob,
 	clmd->scene= scene;	/* nice to pass on later :) */
 	framenr= (int)scene->r.cfra;
 	cache= clmd->point_cache;
-	result = CDDM_copy(dm);
 
 	BKE_ptcache_id_from_cloth(&pid, ob, clmd);
 	BKE_ptcache_id_time(&pid, scene, framenr, &startframe, &endframe, &timescale);
 	clmd->sim_parms->timescale= timescale;
 
-	if(!result) {
-		BKE_ptcache_invalidate(cache);
-		return dm;
-	}
-
 	if(clmd->sim_parms->reset
 		|| (framenr == (startframe - clmd->sim_parms->preroll) && clmd->sim_parms->preroll != 0)
-		|| (clmd->clothObject && result->getNumVerts(result) != clmd->clothObject->numverts))
+		|| (clmd->clothObject && dm->getNumVerts(dm) != clmd->clothObject->numverts))
 	{
 		clmd->sim_parms->reset = 0;
 		cache->flag |= PTCACHE_OUTDATED;
@@ -461,7 +455,7 @@ DerivedMesh *clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob,
 		BKE_ptcache_validate(cache, 0);
 		cache->last_exact= 0;
 		cache->flag &= ~PTCACHE_REDO_NEEDED;
-		return result;
+		return;
 	}
 	
 	// unused in the moment, calculated separately in implicit.c
@@ -473,20 +467,20 @@ DerivedMesh *clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob,
 
 		/* do simulation */
 		if(!do_init_cloth(ob, clmd, dm, framenr))
-			return result;
+			return;
 
 		do_step_cloth(ob, clmd, dm, framenr);
-		cloth_to_object(ob, clmd, result);
+		cloth_to_object(ob, clmd, vertexCos);
 
 		clmd->clothObject->last_frame= framenr;
 
-		return result;
+		return;
 	}
 
 	/* simulation is only active during a specific period */
 	if(framenr < startframe) {
 		BKE_ptcache_invalidate(cache);
-		return result;
+		return;
 	}
 	else if(framenr > endframe) {
 		framenr= endframe;
@@ -494,7 +488,7 @@ DerivedMesh *clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob,
 
 	/* initialize simulation data if it didn't exist already */
 	if(!do_init_cloth(ob, clmd, dm, framenr))
-		return result;
+		return;
 
 	if((framenr == startframe) && (clmd->sim_parms->preroll == 0)) {
 		BKE_ptcache_id_reset(scene, &pid, PTCACHE_RESET_OUTDATED);
@@ -502,7 +496,7 @@ DerivedMesh *clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob,
 		BKE_ptcache_validate(cache, framenr);
 		cache->flag &= ~PTCACHE_REDO_NEEDED;
 		clmd->clothObject->last_frame= framenr;
-		return result;
+		return;
 	}
 
 	/* try to read from cache */
@@ -510,7 +504,7 @@ DerivedMesh *clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob,
 
 	if(cache_result == PTCACHE_READ_EXACT || cache_result == PTCACHE_READ_INTERPOLATED) {
 		implicit_set_positions(clmd);
-		cloth_to_object (ob, clmd, result);
+		cloth_to_object (ob, clmd, vertexCos);
 
 		BKE_ptcache_validate(cache, framenr);
 
@@ -519,7 +513,7 @@ DerivedMesh *clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob,
 
 		clmd->clothObject->last_frame= framenr;
 
-		return result;
+		return;
 	}
 	else if(cache_result==PTCACHE_READ_OLD) {
 		implicit_set_positions(clmd);
@@ -527,11 +521,11 @@ DerivedMesh *clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob,
 	else if( /*ob->id.lib ||*/ (cache->flag & PTCACHE_BAKED)) { /* 2.4x disabled lib, but this can be used in some cases, testing further - campbell */
 		/* if baked and nothing in cache, do nothing */
 		BKE_ptcache_invalidate(cache);
-		return result;
+		return;
 	}
 
 	if(framenr!=clmd->clothObject->last_frame+1)
-		return result;
+		return;
 
 	/* if on second frame, write cache for first frame */
 	if(cache->simframe == startframe && (cache->flag & PTCACHE_OUTDATED || cache->last_exact==0))
@@ -548,10 +542,8 @@ DerivedMesh *clothModifier_do(ClothModifierData *clmd, Scene *scene, Object *ob,
 	else
 		BKE_ptcache_write(&pid, framenr);
 
-	cloth_to_object (ob, clmd, result);
+	cloth_to_object (ob, clmd, vertexCos);
 	clmd->clothObject->last_frame= framenr;
-
-	return result;
 }
 
 /* frees all */
@@ -706,24 +698,19 @@ void cloth_free_modifier_extern ( ClothModifierData *clmd )
  * cloth_to_object - copies the deformed vertices to the object.
  *
  **/
-static void cloth_to_object (Object *ob,  ClothModifierData *clmd, DerivedMesh *dm)
+static void cloth_to_object (Object *ob,  ClothModifierData *clmd, float (*vertexCos)[3])
 {
 	unsigned int	i = 0;
-	MVert *mvert = NULL;
-	unsigned int numverts;
 	Cloth *cloth = clmd->clothObject;
 
 	if (clmd->clothObject) {
 		/* inverse matrix is not uptodate... */
 		invert_m4_m4(ob->imat, ob->obmat);
 
-		mvert = CDDM_get_verts(dm);
+		for (i = 0; i < cloth->numverts; i++)
-		numverts = dm->getNumVerts(dm);
-
-		for (i = 0; i < numverts; i++)
 		{
-			copy_v3_v3 (mvert[i].co, cloth->verts[i].x);
+			copy_v3_v3 (vertexCos[i], cloth->verts[i].x);
-			mul_m4_v3(ob->imat, mvert[i].co);	/* cloth is in global coords */
+			mul_m4_v3(ob->imat, vertexCos[i]);	/* cloth is in global coords */
 		}
 	}
 }

source/blender/blenkernel/intern/constraint.c

@@ -2775,8 +2775,7 @@ static void stretchto_evaluate (bConstraint *con, bConstraintOb *cob, ListBase *
 		/* store Z orientation before destroying obmat */
 		normalize_v3_v3(zz, cob->matrix[2]);
 		
-		dist = len_v3v3(cob->matrix[3], ct->matrix[3]);
+		/* XXX That makes the constraint buggy with asymmetrically scaled objects, see #29940. */
-		/* XXX What was all that for??? Makes the constraint buggy with scaled objects, see #29940. */
 /*		sub_v3_v3v3(vec, cob->matrix[3], ct->matrix[3]);*/
 /*		vec[0] /= size[0];*/
 /*		vec[1] /= size[1];*/
@@ -2784,10 +2783,14 @@ static void stretchto_evaluate (bConstraint *con, bConstraintOb *cob, ListBase *
 		
 /*		dist = normalize_v3(vec);*/
 		
+		dist = len_v3v3(cob->matrix[3], ct->matrix[3]);
+		/* Only Y constrained object axis scale should be used, to keep same length when scaling it. */
+		dist /= size[1];
+		
 		/* data->orglength==0 occurs on first run, and after 'R' button is clicked */
-		if (data->orglength == 0)  
+		if (data->orglength == 0)
 			data->orglength = dist;
-		if (data->bulge == 0) 
+		if (data->bulge == 0)
 			data->bulge = 1.0;
 		
 		scale[1] = dist/data->orglength;

source/blender/blenkernel/intern/displist.c

@@ -833,6 +833,7 @@ static void curve_calc_modifiers_post(Scene *scene, Object *ob, ListBase *dispba
 	int editmode = (!forRender && cu->editnurb);
 	DerivedMesh *dm= NULL, *ndm;
 	float (*vertCos)[3] = NULL;
+	int useCache = !forRender;
 
 	if(forRender) required_mode = eModifierMode_Render;
 	else required_mode = eModifierMode_Realtime;
@@ -911,7 +912,7 @@ static void curve_calc_modifiers_post(Scene *scene, Object *ob, ListBase *dispba
 				vertCos= NULL;
 			}
 
-			ndm = mti->applyModifier(md, ob, dm, forRender, editmode);
+			ndm = mti->applyModifier(md, ob, dm, forRender, useCache);
 
 			if (ndm) {
 				/* Modifier returned a new derived mesh */

source/blender/blenkernel/intern/dynamicpaint.c

@@ -97,6 +97,10 @@ static int neighY[8] = {0,1,1, 1, 0,-1,-1,-1};
 /* brush mesh raycast status */
 #define HIT_VOLUME 1
 #define HIT_PROXIMITY 2
+/* dynamicPaint_findNeighbourPixel() return codes */
+#define NOT_FOUND -1
+#define ON_MESH_EDGE -2
+#define OUT_OF_TEXTURE -3
 /* paint effect default movement per frame in global units */
 #define EFF_MOVEMENT_PER_FRAME 0.05f
 /* initial wave time factor */
@@ -134,10 +138,10 @@ typedef struct Vec3f {
 	float v[3];
 } Vec3f;
 
-typedef struct BakeNeighPoint {
+typedef struct BakeAdjPoint {
 	float dir[3];	/* vector pointing towards this neighbour */
 	float dist;		/* distance to */
-} BakeNeighPoint;
+} BakeAdjPoint;
 
 /* Surface data used while processing a frame	*/
 typedef struct PaintBakeNormal {
@@ -156,7 +160,7 @@ typedef struct PaintBakeData {
 	Bounds3D mesh_bounds;
 
 	/* adjacency info */
-	BakeNeighPoint *bNeighs; /* current global neighbour distances and directions, if required */
+	BakeAdjPoint *bNeighs; /* current global neighbour distances and directions, if required */
 	double average_dist;
 	/* space partitioning */
 	VolumeGrid *grid;		/* space partitioning grid to optimize brush checks */
@@ -188,13 +192,6 @@ typedef struct ImgSeqFormatData {
 	Vec3f *barycentricWeights;		/* b-weights for all pixel samples */
 } ImgSeqFormatData;
 
-#if 0 /* UNUSED */
-typedef struct EffVelPoint {
-	float previous_pos[3];
-	float previous_vel[3];
-} EffVelPoint;
-#endif
-
 /* adjacency data flags */
 #define ADJ_ON_MESH_EDGE (1<<0)
 
@@ -470,19 +467,25 @@ static void object_cacheIgnoreClear(Object *ob, int state)
 	BLI_freelistN(&pidlist);
 }
 
-static void subframe_updateObject(Scene *scene, Object *ob, int flags, float frame)
+static int subframe_updateObject(Scene *scene, Object *ob, int flags, float frame)
 {
 	DynamicPaintModifierData *pmd = (DynamicPaintModifierData *)modifiers_findByType(ob, eModifierType_DynamicPaint);
 	bConstraint *con;
 
 	/* if other is dynamic paint canvas, dont update */
 	if (pmd && pmd->canvas)
-		return;
+		return 1;
 
 	/* if object has parents, update them too */
 	if (flags & UPDATE_PARENTS) {
-		if (ob->parent) subframe_updateObject(scene, ob->parent, 0, frame);
+		int is_canvas = 0;
-		if (ob->track) subframe_updateObject(scene, ob->track, 0, frame);
+		if (ob->parent) is_canvas += subframe_updateObject(scene, ob->parent, 0, frame);
+		if (ob->track) is_canvas += subframe_updateObject(scene, ob->track, 0, frame);
+
+		/* skip subframe if object is parented
+		*  to vertex of a dynamic paint canvas */
+		if (is_canvas && (ob->partype == PARVERT1 || ob->partype == PARVERT3))
+			return 0;
 
 		/* also update constraint targets */
 		for (con = ob->constraints.first; con; con=con->next) {
@@ -519,6 +522,8 @@ static void subframe_updateObject(Scene *scene, Object *ob, int flags, float fra
 	}
 	else
 		where_is_object_time(scene, ob, frame);
+
+	return 0;
 }
 
 static void scene_setSubframe(Scene *scene, float subframe)
@@ -1222,7 +1227,7 @@ static int surface_usesAdjData(DynamicPaintSurface *surface)
 static void dynamicPaint_initAdjacencyData(DynamicPaintSurface *surface, int force_init)
 {
 	PaintSurfaceData *sData = surface->data;
-	PaintAdjData *ed;
+	PaintAdjData *ad;
 	int *temp_data;
 	int neigh_points = 0;
 
@@ -1238,17 +1243,17 @@ static void dynamicPaint_initAdjacencyData(DynamicPaintSurface *surface, int for
 	if (!neigh_points) return;
 
 	/* allocate memory */
-	ed = sData->adj_data = MEM_callocN(sizeof(PaintAdjData), "Surface Adj Data");
+	ad = sData->adj_data = MEM_callocN(sizeof(PaintAdjData), "Surface Adj Data");
-	if (!ed) return;
+	if (!ad) return;
-	ed->n_index = MEM_callocN(sizeof(int)*sData->total_points, "Surface Adj Index");
+	ad->n_index = MEM_callocN(sizeof(int)*sData->total_points, "Surface Adj Index");
-	ed->n_num = MEM_callocN(sizeof(int)*sData->total_points, "Surface Adj Counts");
+	ad->n_num = MEM_callocN(sizeof(int)*sData->total_points, "Surface Adj Counts");
 	temp_data = MEM_callocN(sizeof(int)*sData->total_points, "Temp Adj Data");
-	ed->n_target = MEM_callocN(sizeof(int)*neigh_points, "Surface Adj Targets");
+	ad->n_target = MEM_callocN(sizeof(int)*neigh_points, "Surface Adj Targets");
-	ed->flags = MEM_callocN(sizeof(int)*sData->total_points, "Surface Adj Flags");
+	ad->flags = MEM_callocN(sizeof(int)*sData->total_points, "Surface Adj Flags");
-	ed->total_targets = neigh_points;
+	ad->total_targets = neigh_points;
 
 	/* in case of allocation error, free memory */
-	if (!ed->n_index || !ed->n_num || !ed->n_target || !temp_data) {
+	if (!ad->n_index || !ad->n_num || !ad->n_target || !temp_data) {
 		dynamicPaint_freeAdjData(sData);
 		if (temp_data) MEM_freeN(temp_data);
 		setError(surface->canvas, "Not enough free memory.");
@@ -1267,14 +1272,15 @@ static void dynamicPaint_initAdjacencyData(DynamicPaintSurface *surface, int for
 
 		/* count number of edges per vertex */
 		for (i=0; i<numOfEdges; i++) {
-			ed->n_num[edge[i].v1]++;
+			ad->n_num[edge[i].v1]++;
-			ed->n_num[edge[i].v2]++;
+			ad->n_num[edge[i].v2]++;
 
 			temp_data[edge[i].v1]++;
 			temp_data[edge[i].v2]++;
 		}
 
-		/* to locate points on "mesh edge" */
+		/* also add number of vertices to temp_data
+		*  to locate points on "mesh edge" */
 		for (i=0; i<numOfFaces; i++) {
 			temp_data[face[i].v1]++;
 			temp_data[face[i].v2]++;
@@ -1288,7 +1294,7 @@ static void dynamicPaint_initAdjacencyData(DynamicPaintSurface *surface, int for
 		for (i=0; i<sData->total_points; i++) {
 			if ((temp_data[i]%2) ||
 				temp_data[i] < 4)
-				ed->flags[i] |= ADJ_ON_MESH_EDGE;
+				ad->flags[i] |= ADJ_ON_MESH_EDGE;
 				
 			/* reset temp data */ 
 			temp_data[i] = 0;
@@ -1297,22 +1303,22 @@ static void dynamicPaint_initAdjacencyData(DynamicPaintSurface *surface, int for
 		/* order n_index array */
 		n_pos = 0;
 		for (i=0; i<sData->total_points; i++) {
-			ed->n_index[i] = n_pos;
+			ad->n_index[i] = n_pos;
-			n_pos += ed->n_num[i];
+			n_pos += ad->n_num[i];
 		}
 
 		/* and now add neighbour data using that info */
 		for (i=0; i<numOfEdges; i++) {
 			/* first vertex */
 			int index = edge[i].v1;
-			n_pos = ed->n_index[index]+temp_data[index];
+			n_pos = ad->n_index[index]+temp_data[index];
-			ed->n_target[n_pos] = edge[i].v2;
+			ad->n_target[n_pos] = edge[i].v2;
 			temp_data[index]++;
 
 			/* second vertex */
 			index = edge[i].v2;
-			n_pos = ed->n_index[index]+temp_data[index];
+			n_pos = ad->n_index[index]+temp_data[index];
-			ed->n_target[n_pos] = edge[i].v1;
+			ad->n_target[n_pos] = edge[i].v1;
 			temp_data[index]++;
 		}
 	}
@@ -1500,10 +1506,11 @@ void dynamicPaint_clearSurface(DynamicPaintSurface *surface)
 int dynamicPaint_resetSurface(DynamicPaintSurface *surface)
 {
 	int numOfPoints = dynamicPaint_surfaceNumOfPoints(surface);
-	/* dont touch image sequence types. they get handled only on bake */
+	/* free existing data */
-	if (surface->format == MOD_DPAINT_SURFACE_F_IMAGESEQ) return 1;
-
 	if (surface->data) dynamicPaint_freeSurfaceData(surface);
+
+	/* dont reallocate for image sequence types. they get handled only on bake */
+	if (surface->format == MOD_DPAINT_SURFACE_F_IMAGESEQ) return 1;
 	if (numOfPoints < 1) return 0;
 
 	/* allocate memory */
@@ -1574,7 +1581,6 @@ static struct DerivedMesh *dynamicPaint_Modifier_apply(DynamicPaintModifierData
 
 		DynamicPaintSurface *surface = pmd->canvas->surfaces.first;
 		int update_normals = 0;
-		pmd->canvas->flags &= ~MOD_DPAINT_PREVIEW_READY;
 
 		/* loop through surfaces */
 		for (; surface; surface=surface->next) {
@@ -1651,7 +1657,6 @@ static struct DerivedMesh *dynamicPaint_Modifier_apply(DynamicPaintModifierData
 										}
 									}
 								}
-								pmd->canvas->flags |= MOD_DPAINT_PREVIEW_READY;
 							}
 						}
 
@@ -1711,29 +1716,7 @@ static struct DerivedMesh *dynamicPaint_Modifier_apply(DynamicPaintModifierData
 						if (surface->flags & MOD_DPAINT_PREVIEW) {
 							/* Save preview results to weight layer, to be
 							*   able to share same drawing methods */
-							MFace *mface = result->getFaceArray(result);
+							DM_update_weight_mcol(ob, result, 0, weight, 0, NULL);
-							int numOfFaces = result->getNumFaces(result);
-							int i;
-							MCol *col = result->getFaceDataArray(result, CD_WEIGHT_MCOL);
-							if (!col) col = CustomData_add_layer(&result->faceData, CD_WEIGHT_MCOL, CD_CALLOC, NULL, numOfFaces);
-
-							if (col) {
-								#pragma omp parallel for schedule(static)
-								for (i=0; i<numOfFaces; i++) {
-									float temp_color[3];
-									int j = (mface[i].v4) ? 4 : 3;
-									while (j--) {
-										int index = *((&mface[i].v1)+j);
-
-										weight_to_rgb(temp_color, weight[index]);
-										col[i*4+j].r = FTOCHAR(temp_color[2]);
-										col[i*4+j].g = FTOCHAR(temp_color[1]);
-										col[i*4+j].b = FTOCHAR(temp_color[0]);
-										col[i*4+j].a = 255;
-									}
-								}
-								pmd->canvas->flags |= MOD_DPAINT_PREVIEW_READY;
-							}
 						}
 
 						/* apply weights into a vertex group, if doesnt exists add a new layer */
@@ -1923,8 +1906,8 @@ static int dynamicPaint_findNeighbourPixel(PaintUVPoint *tempPoints, DerivedMesh
 	x = px + neighX[n_index];
 	y = py + neighY[n_index];
 
-	if (x<0 || x>=w) return -1;
+	if (x<0 || x>=w) return OUT_OF_TEXTURE;
-	if (y<0 || y>=h) return -1;
+	if (y<0 || y>=h) return OUT_OF_TEXTURE;
 
 	tPoint = &tempPoints[x+w*y];		/* UV neighbour */
 	cPoint = &tempPoints[px+w*py];		/* Origin point */
@@ -2037,8 +2020,8 @@ static int dynamicPaint_findNeighbourPixel(PaintUVPoint *tempPoints, DerivedMesh
 				}
 			}
 
-			/* If none found return -1	*/
+			/* If none found pixel is on mesh edge	*/
-			if (target_face == -1) return -1;
+			if (target_face == -1) return ON_MESH_EDGE;
 
 			/*
 			*	If target face is connected in UV space as well, just use original index
@@ -2076,15 +2059,15 @@ static int dynamicPaint_findNeighbourPixel(PaintUVPoint *tempPoints, DerivedMesh
 			final_pixel[1] = (int)floor(pixel[1]);
 
 			/* If current pixel uv is outside of texture	*/
-			if (final_pixel[0] < 0 || final_pixel[0] >= w) return -1;
+			if (final_pixel[0] < 0 || final_pixel[0] >= w) return OUT_OF_TEXTURE;
-			if (final_pixel[1] < 0 || final_pixel[1] >= h) return -1;
+			if (final_pixel[1] < 0 || final_pixel[1] >= h) return OUT_OF_TEXTURE;
 
 			final_index = final_pixel[0] + w * final_pixel[1];
 
 			/* If we ended up to our origin point ( mesh has smaller than pixel sized faces)	*/
-			if (final_index == (px+w*py)) return -1;
+			if (final_index == (px+w*py)) return NOT_FOUND;
 			/* If found pixel still lies on wrong face ( mesh has smaller than pixel sized faces)	*/
-			if (tempPoints[final_index].face_index != target_face) return -1;
+			if (tempPoints[final_index].face_index != target_face) return NOT_FOUND;
 
 			/*
 			*	If final point is an "edge pixel", use it's "real" neighbour instead
@@ -2466,11 +2449,14 @@ int dynamicPaint_createUVSurface(DynamicPaintSurface *surface)
 								*  If not found, -1 is returned */
 								int n_target = dynamicPaint_findNeighbourPixel(tempPoints, dm, uvname, w, h, tx, ty, i);
 
-								if (n_target != -1) {
+								if (n_target >= 0) {
 									ed->n_target[n_pos] = final_index[n_target];
 									ed->n_num[final_index[index]]++;
 									n_pos++;
 								}
+								else if (n_target == ON_MESH_EDGE || n_target == OUT_OF_TEXTURE) {
+									ed->flags[final_index[index]] |= ADJ_ON_MESH_EDGE;
+								}
 							}
 						}
 					}
@@ -3171,8 +3157,8 @@ static int dynamicPaint_paintMesh(DynamicPaintSurface *surface,
 			mul_m4_v3(brushOb->obmat, mvert[ii].co);
 			boundInsert(&mesh_bb, mvert[ii].co);
 
-			/* for project brush calculate average normal */
+			/* for proximity project calculate average normal */
-			if (brush->flags & MOD_DPAINT_PROX_PROJECT) {
+			if (brush->flags & MOD_DPAINT_PROX_PROJECT && brush->collision != MOD_DPAINT_COL_VOLUME) {
 				float nor[3];
 				normal_short_to_float_v3(nor, mvert[ii].no);
 				mul_mat3_m4_v3(brushOb->obmat, nor);
@@ -3182,7 +3168,7 @@ static int dynamicPaint_paintMesh(DynamicPaintSurface *surface,
 			}
 		}
 
-		if (brush->flags & MOD_DPAINT_PROX_PROJECT) {
+		if (brush->flags & MOD_DPAINT_PROX_PROJECT && brush->collision != MOD_DPAINT_COL_VOLUME) {
 			mul_v3_fl(avg_brushNor, 1.0f/(float)numOfVerts);
 			/* instead of null vector use positive z */
 			if (!(MIN3(avg_brushNor[0],avg_brushNor[1],avg_brushNor[2])))
@@ -3566,7 +3552,7 @@ static int dynamicPaint_paintParticles(DynamicPaintSurface *surface,
 		/* Proceed only if particle is active	*/
 		if(pa->alive == PARS_UNBORN && (part->flag & PART_UNBORN)==0) continue;									
 		else if(pa->alive == PARS_DEAD && (part->flag & PART_DIED)==0) continue;									
-		else if(pa->flag & PARS_NO_DISP || pa->flag & PARS_UNEXIST) continue;
+		else if(pa->flag & PARS_UNEXIST) continue;
 
 		/*	for debug purposes check if any NAN particle proceeds
 		*	For some reason they get past activity check, this should rule most of them out	*/
@@ -3879,14 +3865,13 @@ static int dynamicPaint_paintSinglePoint(DynamicPaintSurface *surface, float *po
 /***************************** Dynamic Paint Step / Baking ******************************/
 
 /*
-*	Calculate current frame neighbouring point distances
+*	Calculate current frame distances and directions for adjacency data
-*	and direction vectors
 */
-static void dynamicPaint_prepareNeighbourData(DynamicPaintSurface *surface, int force_init)
+static void dynamicPaint_prepareAdjacencyData(DynamicPaintSurface *surface, int force_init)
 {
 	PaintSurfaceData *sData = surface->data;
 	PaintBakeData *bData = sData->bData;
-	BakeNeighPoint *bNeighs;
+	BakeAdjPoint *bNeighs;
 	PaintAdjData *adj_data = sData->adj_data;
 	Vec3f *realCoord = bData->realCoord;
 	int index;
@@ -3894,7 +3879,7 @@ static void dynamicPaint_prepareNeighbourData(DynamicPaintSurface *surface, int
 	if ((!surface_usesAdjDistance(surface) && !force_init) || !sData->adj_data) return;
 
 	if (bData->bNeighs) MEM_freeN(bData->bNeighs);
-	bNeighs = bData->bNeighs = MEM_mallocN(sData->adj_data->total_targets*sizeof(struct BakeNeighPoint),"PaintEffectBake");
+	bNeighs = bData->bNeighs = MEM_mallocN(sData->adj_data->total_targets*sizeof(struct BakeAdjPoint),"PaintEffectBake");
 	if (!bNeighs) return;
 
 	#pragma omp parallel for schedule(static)
@@ -3933,7 +3918,7 @@ static void dynamicPaint_prepareNeighbourData(DynamicPaintSurface *surface, int
 /* find two adjacency points (closest_id) and influence (closest_d) to move paint towards when affected by a force  */
 void surface_determineForceTargetPoints(PaintSurfaceData *sData, int index, float force[3], float closest_d[2], int closest_id[2])
 {
-	BakeNeighPoint *bNeighs = sData->bData->bNeighs;
+	BakeAdjPoint *bNeighs = sData->bData->bNeighs;
 	int numOfNeighs = sData->adj_data->n_num[index];
 	int i;
 
@@ -4002,7 +3987,7 @@ static void dynamicPaint_doSmudge(DynamicPaintSurface *surface, DynamicPaintBrus
 {
 	PaintSurfaceData *sData = surface->data;
 	PaintBakeData *bData = sData->bData;
-	BakeNeighPoint *bNeighs = sData->bData->bNeighs;
+	BakeAdjPoint *bNeighs = sData->bData->bNeighs;
 	int index, steps, step;
 	float eff_scale, max_velocity = 0.0f;
 
@@ -4161,7 +4146,7 @@ static int dynamicPaint_prepareEffectStep(DynamicPaintSurface *surface, Scene *s
 static void dynamicPaint_doEffectStep(DynamicPaintSurface *surface, float *force, PaintPoint *prevPoint, float timescale, float steps)
 {
 	PaintSurfaceData *sData = surface->data;
-	BakeNeighPoint *bNeighs = sData->bData->bNeighs;
+	BakeAdjPoint *bNeighs = sData->bData->bNeighs;
 	float distance_scale = getSurfaceDimension(sData)/CANVAS_REL_SIZE;
 	int index;
 	timescale /= steps;
@@ -4190,7 +4175,7 @@ static void dynamicPaint_doEffectStep(DynamicPaintSurface *surface, float *force
 			/*	Loop through neighbouring points	*/
 			for (i=0; i<numOfNeighs; i++) {
 				int n_index = sData->adj_data->n_index[index]+i;
-				float w_factor /* , p_alpha = pPoint->e_alpha */ /* UNUSED */;
+				float w_factor;
 				PaintPoint *ePoint = &prevPoint[sData->adj_data->n_target[n_index]];
 				float speed_scale = (bNeighs[n_index].dist<eff_scale) ? 1.0f : eff_scale/bNeighs[n_index].dist;
 				float color_mix = (MIN3(ePoint->wetness, pPoint->wetness, 1.0f))*0.25f*surface->color_spread_speed;
@@ -4327,7 +4312,7 @@ static void dynamicPaint_doEffectStep(DynamicPaintSurface *surface, float *force
 void dynamicPaint_doWaveStep(DynamicPaintSurface *surface, float timescale)
 {
 	PaintSurfaceData *sData = surface->data;
-	BakeNeighPoint *bNeighs = sData->bData->bNeighs;
+	BakeAdjPoint *bNeighs = sData->bData->bNeighs;
 	int index;
 	int steps, ss;
 	float dt, min_dist, damp_factor;
@@ -4760,8 +4745,8 @@ static int dynamicPaint_generateBakeData(DynamicPaintSurface *surface, Scene *sc
 
 	/* generate surface space partitioning grid */
 	surfaceGenerateGrid(surface);
-	/* calculate current frame neighbouring point distances and global dirs */
+	/* calculate current frame adjacency point distances and global dirs */
-	dynamicPaint_prepareNeighbourData(surface, 0);
+	dynamicPaint_prepareAdjacencyData(surface, 0);
 
 	/* Copy current frame vertices to check against in next frame */
 	copy_m4_m4(bData->prev_obmat, ob->obmat);
@@ -4844,7 +4829,7 @@ static int dynamicPaint_doStep(Scene *scene, Object *ob, DynamicPaintSurface *su
 						if (!sData->adj_data)
 							dynamicPaint_initAdjacencyData(surface, 1);
 						if (!bData->bNeighs)
-							dynamicPaint_prepareNeighbourData(surface, 1);
+							dynamicPaint_prepareAdjacencyData(surface, 1);
 					}
 
 					/* update object data on this subframe */

source/blender/blenkernel/intern/implicit.c

@@ -1852,6 +1852,9 @@ int implicit_solver (Object *ob, float frame, ClothModifierData *clmd, ListBase
 	
 	while(step < tf)
 	{	
+		// damping velocity for artistic reasons
+		mul_lfvectorS(id->V, id->V, clmd->sim_parms->vel_damping, numverts);
+
 		// calculate forces
 		cloth_calc_force(clmd, frame, id->F, id->X, id->V, id->dFdV, id->dFdX, effectors, step, id->M);
 		

source/blender/blenkernel/intern/modifier.c

@@ -142,6 +142,19 @@ int modifier_supportsMapping(ModifierData *md)
 			(mti->flags & eModifierTypeFlag_SupportsMapping));
 }
 
+int modifier_isPreview(ModifierData *md)
+{
+	ModifierTypeInfo *mti = modifierType_getInfo(md->type);
+
+	if (!(mti->flags & eModifierTypeFlag_UsesPreview))
+		return FALSE;
+
+	if (md->mode & eModifierMode_Realtime)
+		return TRUE;
+
+	return FALSE;
+}
+
 ModifierData *modifiers_findByType(Object *ob, ModifierType type)
 {
 	ModifierData *md = ob->modifiers.first;
@@ -385,6 +398,21 @@ LinkNode *modifiers_calcDataMasks(struct Scene *scene, Object *ob, ModifierData
 	return dataMasks;
 }
 
+ModifierData *modifiers_getLastPreview(struct Scene *scene, ModifierData *md, int required_mode)
+{
+	ModifierData *tmp_md = NULL;
+
+	if (required_mode != eModifierMode_Realtime)
+		return tmp_md;
+
+	/* Find the latest modifier in stack generating preview. */
+	for(; md; md = md->next) {
+		if(modifier_isEnabled(scene, md, required_mode) && modifier_isPreview(md))
+			tmp_md = md;
+	}
+	return tmp_md;
+}
+
 ModifierData *modifiers_getVirtualModifierList(Object *ob)
 {
 		/* Kinda hacky, but should be fine since we are never
@@ -545,6 +573,20 @@ int modifiers_isCorrectableDeformed(Object *ob)
 	return 0;
 }
 
+/* Check whether the given object has a modifier in its stack that uses WEIGHT_MCOL CD layer
+ * to preview something... Used by DynamicPaint and WeightVG currently. */
+int modifiers_isPreview(Object *ob)
+{
+	ModifierData *md = ob->modifiers.first;
+
+	for (; md; md = md->next) {
+		if (modifier_isPreview(md))
+			return TRUE;
+	}
+
+	return FALSE;
+}
+
 int modifiers_indexInObject(Object *ob, ModifierData *md_seek)
 {
 	int i= 0;

source/blender/blenkernel/intern/nla.c

@@ -1448,9 +1448,17 @@ void BKE_nla_validate_state (AnimData *adt)
 			 */
 			// TODO: 1 solution is to tie this in with auto-blending...
 			if (strip->extendmode != NLASTRIP_EXTEND_NOTHING) {
+				/* 1) First strip must be set to extend hold, otherwise, stuff before acts dodgy
+				 * 2) Only overwrite extend mode if *not* changing it will most probably result in 
+				 * occlusion problems, which will occur iff
+				 *	- blendmode = REPLACE
+				 *	- all channels the same (this is fiddly to test, so is currently assumed)
+				 *
+				 * Should fix problems such as [#29869]
+				 */
 				if (strip == fstrip)
 					strip->extendmode= NLASTRIP_EXTEND_HOLD;
-				else
+				else if (strip->blendmode == NLASTRIP_MODE_REPLACE)
 					strip->extendmode= NLASTRIP_EXTEND_HOLD_FORWARD;
 			}
 		}
@@ -1542,6 +1550,34 @@ short BKE_nla_tweakmode_enter (AnimData *adt)
 			break;
 		}	
 	}
+	
+	/* There are situations where we may have multiple strips selected and we want to enter tweakmode on all 
+	 * of those at once. Usually in those cases, it will usually just be a single strip per AnimData. 
+	 * In such cases, compromise and take the last selected track and/or last selected strip [#28468] 
+	 */
+	if (activeTrack == NULL) {
+		/* try last selected track for active strip */
+		for (nlt = adt->nla_tracks.last; nlt; nlt = nlt->prev) {
+			if (nlt->flag & NLATRACK_SELECTED) {
+				/* assume this is the active track */
+				activeTrack= nlt;
+				
+				/* try to find active strip */
+				activeStrip= BKE_nlastrip_find_active(nlt);
+				break;
+			}
+		}	
+	}
+	if ((activeTrack) && (activeStrip == NULL)) {
+		/* no active strip in active or last selected track; compromise for first selected (assuming only single)... */
+		for (strip = activeTrack->strips.first; strip; strip= strip->next) {
+			if (strip->flag & (NLASTRIP_FLAG_SELECT|NLASTRIP_FLAG_ACTIVE)) {
+				activeStrip = strip;
+				break;
+			}
+		}
+	}
+	
 	if ELEM3(NULL, activeTrack, activeStrip, activeStrip->act) {
 		if (G.f & G_DEBUG) {
 			printf("NLA tweakmode enter - neither active requirement found \n");

source/blender/blenkernel/intern/node.c

@@ -1881,6 +1881,7 @@ static void registerShaderNodes(bNodeTreeType *ttype)
 	register_node_type_sh_material(ttype);
 	register_node_type_sh_camera(ttype);
 	register_node_type_sh_gamma(ttype);
+	register_node_type_sh_brightcontrast(ttype);
 	register_node_type_sh_value(ttype);
 	register_node_type_sh_rgb(ttype);
 	register_node_type_sh_mix_rgb(ttype);

source/blender/blenkernel/intern/particle_system.c

@@ -1126,6 +1126,9 @@ static int distribute_threads_init_data(ParticleThread *threads, Scene *scene, D
 			fprintf(stderr,"Particle distribution error: Nothing to emit from!\n");
 
 		if(dm != finaldm) dm->release(dm);
+
+		BLI_kdtree_free(tree);
+
 		return 0;
 	}
 
@@ -3477,6 +3480,7 @@ static void do_hair_dynamics(ParticleSimulationData *sim)
 	int totedge;
 	int k;
 	float hairmat[4][4];
+	float (*deformedVerts)[3];
 
 	if(!psys->clmd) {
 		psys->clmd = (ClothModifierData*)modifier_new(eModifierType_Cloth);
@@ -3570,7 +3574,15 @@ static void do_hair_dynamics(ParticleSimulationData *sim)
 	psys->clmd->point_cache = psys->pointcache;
 	psys->clmd->sim_parms->effector_weights = psys->part->effector_weights;
 
-	psys->hair_out_dm = clothModifier_do(psys->clmd, sim->scene, sim->ob, dm);
+	deformedVerts = MEM_callocN(sizeof(*deformedVerts)*dm->getNumVerts(dm), "do_hair_dynamics vertexCos");
+	psys->hair_out_dm = CDDM_copy(dm);
+	psys->hair_out_dm->getVertCos(psys->hair_out_dm, deformedVerts);
+
+	clothModifier_do(psys->clmd, sim->scene, sim->ob, dm, deformedVerts);
+
+	CDDM_apply_vert_coords(psys->hair_out_dm, deformedVerts);
+
+	MEM_freeN(deformedVerts);
 
 	psys->clmd->sim_parms->effector_weights = NULL;
 }

source/blender/blenkernel/intern/scene.c

@@ -538,6 +538,8 @@ Scene *add_scene(const char *name)
 	sce->gm.recastData.detailsampledist = 6.0f;
 	sce->gm.recastData.detailsamplemaxerror = 1.0f;
 
+	sce->gm.exitkey = 218; // Blender key code for ESC
+
 	sound_create_scene(sce);
 
 	return sce;

source/blender/blenkernel/intern/smoke.c

@@ -77,6 +77,9 @@
 
 #include "BKE_smoke.h"
 
+/* UNUSED so far, may be enabled later */
+/* #define USE_SMOKE_COLLISION_DM */
+
 #ifdef WITH_SMOKE
 
 #ifdef _WIN32
@@ -617,9 +620,11 @@ static void smokeModifier_freeCollision(SmokeModifierData *smd)
 			smd->coll->bvhtree = NULL;
 		}
 
+#ifdef USE_SMOKE_COLLISION_DM
 		if(smd->coll->dm)
 			smd->coll->dm->release(smd->coll->dm);
 		smd->coll->dm = NULL;
+#endif
 
 		MEM_freeN(smd->coll);
 		smd->coll = NULL;
@@ -682,9 +687,11 @@ void smokeModifier_reset(struct SmokeModifierData *smd)
 				smd->coll->bvhtree = NULL;
 			}
 
+#ifdef USE_SMOKE_COLLISION_DM
 			if(smd->coll->dm)
 				smd->coll->dm->release(smd->coll->dm);
 			smd->coll->dm = NULL;
+#endif
 
 		}
 	}
@@ -772,7 +779,10 @@ void smokeModifier_createType(struct SmokeModifierData *smd)
 			smd->coll->points = NULL;
 			smd->coll->numpoints = 0;
 			smd->coll->bvhtree = NULL;
+
+#ifdef USE_SMOKE_COLLISION_DM
 			smd->coll->dm = NULL;
+#endif
 		}
 	}
 }
@@ -1339,11 +1349,13 @@ void smokeModifier_do(SmokeModifierData *smd, Scene *scene, Object *ob, DerivedM
 		{
 			// XXX TODO
 			smd->time = scene->r.cfra;
-			
+
+#ifdef USE_SMOKE_COLLISION_DM
 			if(smd->coll->dm)
 				smd->coll->dm->release(smd->coll->dm);
 
 			smd->coll->dm = CDDM_copy(dm);
+#endif
 
 			// rigid movement support
 			copy_m4_m4(smd->coll->mat_old, smd->coll->mat);