== Multires/VPaint ==

Enable multires painting.

* Added operator/RNA/UI to toggle multires for vertex colors. (UI is not great right now, just a button in mesh data properties.)
* Added layer names to GridKey. These are used (in combination with CD type) to identify the source layer from CustomData.
* Changed default mcol to white with alpha=0
* Renamed CD_FACEGRID to CD_GRIDS, removed CD_DISP
* Added a new CustomDataMultires type (CD_GRIDS), which stores layered data like CustomData. However, it only stores float types, and is very simplified.
* Reworked PaintMask to use the new CD multires stuff, also used for vertex colors.
* Started changing vpaint to internally use floats.
* Layering support for multires vpaint. Layer alpha is used to combine the output.

TODO:
* Doesn't handle layer renaming yet (so if you rename your mcol layer, it'll disassociate from the multires data)
* Layers for non-multires vpaint
* Default mcol layer in the startup blend has full alpha, so for testing layers you have to delete the default mcol layer and add new ones.
* Multires level in vpaint is controlled by the Preview level

19 files changed, 786 insertions, 227 deletions

diff --git a/release/scripts/ui/properties_data_mesh.py b/release/scripts/ui/properties_data_mesh.py
index 44d3d19e793..2e5fbe7e37d 100644
--- a/release/scripts/ui/properties_data_mesh.py
+++ b/release/scripts/ui/properties_data_mesh.py
@@ -366,6 +366,10 @@ class DATA_PT_vertex_colors(DataButtonsPanel):
         lay = me.active_vertex_color
         if lay:
             layout.prop(lay, "name")
+            if lay.multiresolution:
+                layout.operator("mesh.vertex_color_multiresolution_toggle", text="Remove Multires")
+            else:
+                layout.operator("mesh.vertex_color_multiresolution_toggle", text="Add Multires")
 
 
 classes = [
diff --git a/source/blender/blenkernel/BKE_customdata.h b/source/blender/blenkernel/BKE_customdata.h
index ff8d2dc300c..f8f56ef03fd 100644
--- a/source/blender/blenkernel/BKE_customdata.h
+++ b/source/blender/blenkernel/BKE_customdata.h
@@ -35,6 +35,7 @@
 struct ID;
 struct CustomData;
 struct CustomDataLayer;
+struct CustomDataMultires;
 typedef unsigned int CustomDataMask;
 
 extern const CustomDataMask CD_MASK_BAREMESH;
@@ -234,6 +235,8 @@ int CustomData_get_render_layer(const struct CustomData *data, int type);
 int CustomData_get_clone_layer(const struct CustomData *data, int type);
 int CustomData_get_stencil_layer(const struct CustomData *data, int type);
 
+char *CustomData_get_layer_name_at_offset(const struct CustomData *data, int type, int offset);
+
 /* copies the data from source to the data element at index in the first
  * layer of type
  * no effect if there is no layer of type
@@ -271,6 +274,7 @@ void CustomData_set_layer_stencil_index(struct CustomData *data, int type, int n
 
 /* adds flag to the layer flags */
 void CustomData_set_layer_flag(struct CustomData *data, int type, int flag);
+void CustomData_set_layer_offset_flag(struct CustomData *data, int type, int offset, int flag);
 
 /* alloc/free a block of custom data attached to one element in editmode */
 void CustomData_em_set_default(struct CustomData *data, void **block);
@@ -311,7 +315,38 @@ void CustomData_from_bmeshpoly(struct CustomData *fdata, struct CustomData *pdat
 void CustomData_bmesh_init_pool(struct CustomData *data, int allocsize);
 
 /* Subsurf grids */
-void CustomData_set_num_grid_elements(struct CustomData *data, int grid_elems);
+
+/* return the number of layers of type that have multires data */
+int CustomData_get_multires_count(struct CustomData *cd, int type);
+
+/* allocates a list of names of layers that have multires data */
+void *CustomData_get_multires_names(struct CustomData *cd, int type);
+
+/* number of floats used per-element for the multires of a customdata type */
+int CustomData_multires_type_totfloat(int type);
+
+/* returns the multires data for a layer matching name and type,
+   or NULL if no such layer found */
+float *CustomData_multires_get_data(struct CustomDataMultires *cdm, int type,
+				    char *name);
+
+/* if layer matching type and name exists, free and replace its griddata
+   otherwise create the layer and set its griddata */
+void CustomData_multires_assign_data(struct CustomDataMultires *cdm, int type,
+				     char *name, float *data);
+
+/* insert a multires layer of the specified type */
+void CustomData_multires_add_layer(struct CustomDataMultires *cdm, int type,
+				   char *name, float *data);
+
+/* remove the multires layer with matching source name
+   returns 1 if succesful, 0 otherwise */
+int CustomData_multires_remove_layer(struct CustomDataMultires *cdm, int type,
+				     char *name);
+
+/* rename a layer matching type and old_name */
+void CustomData_multires_rename(struct CustomDataMultires *cdm, int type,
+				char *old_name, char *name);
 
 /* External file storage */
 
diff --git a/source/blender/blenkernel/BKE_dmgrid.h b/source/blender/blenkernel/BKE_dmgrid.h
index c676fdca5f2..fef5df2017f 100644
--- a/source/blender/blenkernel/BKE_dmgrid.h
+++ b/source/blender/blenkernel/BKE_dmgrid.h
@@ -4,37 +4,45 @@
 /* Each grid element can contain zero or more layers of coordinates,
    paint masks, and normals; these numbers are stored in the GridKey
 
-   For now, co and no can have only zero or one layers, only mask is
-   really variable.
+   The name arrays are the unique names of the source customdata layer
 */
 typedef struct GridKey {
 	int co;
 	int color;
 	int mask;
 	int no;
+
+	/* key to identify the source layer */
+	char (*color_names)[32];
+	char (*mask_names)[32];
 } GridKey;
 
-#define GRIDELEM_KEY_INIT(_gridkey, _totco, _totcolor, _totmask, _totno) \
-	((_gridkey)->co = _totco, (_gridkey)->color = _totcolor, (_gridkey)->mask = _totmask, (_gridkey)->no = _totno)
+#define GRIDELEM_KEY_INIT(_key, _totco, _totcolor, _totmask, _totno) \
+	((_key)->co = _totco, (_key)->color = _totcolor,	     \
+	 (_key)->mask = _totmask, (_key)->no = _totno,		     \
+	 (_key)->color_names = NULL, (_key)->mask_names = NULL)
 
-#define GRIDELEM_SIZE(_key) ((3*(_key)->co + 3*(_key)->color + (_key)->mask + 3*(_key)->no) * sizeof(float))
-#define GRIDELEM_INTERP_COUNT(_key) (3*(_key)->co + 3*(_key)->color + (_key)->mask)
+#define GRIDELEM_SIZE(_key) ((3*(_key)->co + 4*(_key)->color + (_key)->mask + 3*(_key)->no) * sizeof(float))
+#define GRIDELEM_INTERP_COUNT(_key) (3*(_key)->co + 4*(_key)->color + (_key)->mask)
 
 #define GRIDELEM_COLOR_OFFSET(_key) (3*(_key)->co*sizeof(float))
-#define GRIDELEM_MASK_OFFSET(_key) (GRIDELEM_COLOR_OFFSET(_key) + 3*(_key)->color*sizeof(float))
+#define GRIDELEM_MASK_OFFSET(_key) (GRIDELEM_COLOR_OFFSET(_key) + 4*(_key)->color*sizeof(float))
 #define GRIDELEM_NO_OFFSET(_key) (GRIDELEM_MASK_OFFSET(_key) + (_key)->mask*sizeof(float))
 
 #define GRIDELEM_AT(_grid, _elem, _key) ((struct DMGridData*)(((char*)(_grid)) + (_elem) * GRIDELEM_SIZE(_key)))
 #define GRIDELEM_INC(_grid, _inc, _key) ((_grid) = GRIDELEM_AT(_grid, _inc, _key))
 
+ /* I can't figure out how to cast this type without a typedef,
+    having the array length is useful to directly index layers */
+typedef float (*gridelem_f4)[4];
 #define GRIDELEM_CO(_grid, _key) ((float*)(_grid))
-#define GRIDELEM_COLOR(_grid, _key) ((float*)((char*)(_grid) + GRIDELEM_COLOR_OFFSET(_key)))
-#define GRIDELEM_NO(_grid, _key) ((float*)((char*)(_grid) + GRIDELEM_NO_OFFSET(_key)))
+#define GRIDELEM_COLOR(_grid, _key) ((gridelem_f4)((char*)(_grid) + GRIDELEM_COLOR_OFFSET(_key)))
 #define GRIDELEM_MASK(_grid, _key) ((float*)((char*)(_grid) + GRIDELEM_MASK_OFFSET(_key)))
+#define GRIDELEM_NO(_grid, _key) ((float*)((char*)(_grid) + GRIDELEM_NO_OFFSET(_key)))
 
 #define GRIDELEM_CO_AT(_grid, _elem, _key) GRIDELEM_CO(GRIDELEM_AT(_grid, _elem, _key), _key)
 #define GRIDELEM_COLOR_AT(_grid, _elem, _key) GRIDELEM_COLOR(GRIDELEM_AT(_grid, _elem, _key), _key)
-#define GRIDELEM_NO_AT(_grid, _elem, _key) GRIDELEM_NO(GRIDELEM_AT(_grid, _elem, _key), _key)
 #define GRIDELEM_MASK_AT(_grid, _elem, _key) GRIDELEM_MASK(GRIDELEM_AT(_grid, _elem, _key), _key)
+#define GRIDELEM_NO_AT(_grid, _elem, _key) GRIDELEM_NO(GRIDELEM_AT(_grid, _elem, _key), _key)
 
 #endif
diff --git a/source/blender/blenkernel/intern/CCGSubSurf.c b/source/blender/blenkernel/intern/CCGSubSurf.c
index c3ccde1625b..131db2b843a 100644
--- a/source/blender/blenkernel/intern/CCGSubSurf.c
+++ b/source/blender/blenkernel/intern/CCGSubSurf.c
@@ -749,6 +749,15 @@ void ccgSubSurf_free(CCGSubSurf *ss) {
 	_ehash_free(ss->eMap, (EHEntryFreeFP) _edge_free, ss);
 	_ehash_free(ss->vMap, (EHEntryFreeFP) _vert_free, ss);
 
+	/* free gridkey */
+	if(ss->meshIFC.gridkey.color_names)
+		MEM_freeN(ss->meshIFC.gridkey.color_names);
+	if(ss->meshIFC.gridkey.mask_names)
+		MEM_freeN(ss->meshIFC.gridkey.mask_names);
+
+	ss->meshIFC.gridkey.color_names = NULL;
+	ss->meshIFC.gridkey.mask_names = NULL;
+
 	CCGSUBSURF_free(ss, ss);
 
 	if (allocatorIFC.release) {
diff --git a/source/blender/blenkernel/intern/customdata.c b/source/blender/blenkernel/intern/customdata.c
index 96ba8c71494..9bccbc7f33e 100644
--- a/source/blender/blenkernel/intern/customdata.c
+++ b/source/blender/blenkernel/intern/customdata.c
@@ -749,7 +749,7 @@ static void layerSwap_mcol(void *data, const int *corner_indices)
 
 static void layerDefault_mcol(void *data, int count)
 {
-	static MCol default_mcol = {255, 255, 255, 255};
+	static MCol default_mcol = {0, 255, 255, 255}; /* abgr */
 	MCol *mcol = (MCol*)data;
 	int i;
 
@@ -759,26 +759,41 @@ static void layerDefault_mcol(void *data, int count)
 
 /* Grid */
 
-void layerCopy_grid(const void *source_v, void *dest_v, int count)
+static void layerCopy_grid(const void *source_v, void *dest_v, int count)
 {
-	const CustomData *source = source_v;
-	CustomData *dest = dest_v;
-	int i;
+	const CustomDataMultires *source = source_v;
+	CustomDataMultires *dest = dest_v;
+	int i, j;
 
 	for(i = 0; i < count; ++i) {
-		CustomData_copy(source + i, dest + i, ~0, CD_DUPLICATE,
-				source[i].grid_elems);
-		dest[i].grid_elems = source[i].grid_elems;
+		dest[i].totlayer = source[i].totlayer;
+		dest[i].layers = MEM_callocN(sizeof(CustomDataMultiresLayer) *
+					     dest[i].totlayer,
+					     "CustomDataMultiresLayers");
+
+		for(j = 0; j < source[i].totlayer; ++j) {
+			CustomDataMultiresLayer *dl = dest[i].layers + j;
+			CustomDataMultiresLayer *sl = source[i].layers + j;
+
+			dl->type = sl->type;
+			dl->griddata = MEM_dupallocN(sl->griddata);
+			BLI_strncpy(dl->name, sl->name, sizeof(dl->name));
+		}
 	}
 }
 
-void layerFree_grid(void *data, int count, int size)
+static void layerFree_grid(void *data_v, int count, int size)
 {
-	CustomData *cd = data;
-	int i;
+	CustomDataMultires *data = data_v;
+	int i, j;
 
 	for(i = 0; i < count; ++i) {
-		CustomData_free(cd + i, cd[i].grid_elems);
+		for(j = 0; j < data[i].totlayer; ++j) {
+			if(data[i].layers[j].griddata)
+				MEM_freeN(data[i].layers[j].griddata);
+		}
+		if(data[i].layers)
+			MEM_freeN(data[i].layers);
 	}
 }
 
@@ -822,9 +837,11 @@ const LayerTypeInfo LAYERTYPEINFO[CD_NUMTYPES] = {
 	 layerSwap_mcol, layerDefault_mcol},
 	{sizeof(MCol)*4, "MCol", 4, "TexturedCol", NULL, NULL, layerInterp_mcol,
 	 layerSwap_mcol, layerDefault_mcol},
+	/* CD_CLOTH_ORCO */
 	{sizeof(float)*3, "", 0, NULL, NULL, NULL, NULL, NULL, NULL},
-	{sizeof(CustomData), "CustomData", 1, "Face Grid", layerCopy_grid, layerFree_grid, NULL, NULL, NULL},
-	{sizeof(float)*3, "", 0, NULL, NULL, NULL, NULL, NULL, NULL},
+	/* CD_GRID */
+	{sizeof(CustomData), "CustomDataMultires", 1, "Grid", layerCopy_grid, layerFree_grid, NULL, NULL, NULL},
+	/* CD_PAINTMASK */
 	{sizeof(float), "", 0, "Mask", NULL, NULL, NULL, NULL, NULL},
 };
 
@@ -833,8 +850,8 @@ const char *LAYERTYPENAMES[CD_NUMTYPES] = {
 	/*   5-9 */ "CDMTFace", "CDMCol", "CDOrigIndex", "CDNormal", "CDFlags",
 	/* 10-14 */ "CDMFloatProperty", "CDMIntProperty","CDMStringProperty", "CDOrigSpace", "CDOrco",
 	/* 15-19 */ "CDMTexPoly", "CDMLoopUV", "CDMloopCol", "CDTangent", "CDMDisps",
-	/* 20-24 */ "CDWeightMCol", "CDIDMCol", "CDTextureMCol", "CDClothOrco", "CDFaceGrid",
-	/* 25-26 */ "CDDisp", "CDPaintMask"
+	/* 20-24 */ "CDWeightMCol", "CDIDMCol", "CDTextureMCol", "CDClothOrco", "CDGrid",
+	/* 25-26 */ "CDPaintMask"
 };
 
 const CustomDataMask CD_MASK_BAREMESH =
@@ -843,11 +860,11 @@ const CustomDataMask CD_MASK_MESH =
 	CD_MASK_MVERT | CD_MASK_MEDGE | CD_MASK_MFACE |
 	CD_MASK_MSTICKY | CD_MASK_MDEFORMVERT | CD_MASK_MTFACE | CD_MASK_MCOL |
 	CD_MASK_PROP_FLT | CD_MASK_PROP_INT | CD_MASK_PROP_STR | CD_MASK_MDISPS |
-	CD_MASK_FACEGRID | CD_MASK_PAINTMASK;
+	CD_MASK_PAINTMASK;
 const CustomDataMask CD_MASK_EDITMESH =
 	CD_MASK_MSTICKY | CD_MASK_MDEFORMVERT | CD_MASK_MTFACE |
 	CD_MASK_MCOL|CD_MASK_PROP_FLT | CD_MASK_PROP_INT | CD_MASK_PROP_STR |
-	CD_MASK_MDISPS | CD_MASK_FACEGRID | CD_MASK_PAINTMASK;
+	CD_MASK_MDISPS | CD_MASK_PAINTMASK;
 const CustomDataMask CD_MASK_DERIVEDMESH =
 	CD_MASK_MSTICKY | CD_MASK_MDEFORMVERT | CD_MASK_MTFACE |
 	CD_MASK_MCOL | CD_MASK_ORIGINDEX | CD_MASK_PROP_FLT | CD_MASK_PROP_INT | CD_MASK_CLOTH_ORCO |
@@ -1100,6 +1117,20 @@ int CustomData_get_stencil_layer(const CustomData *data, int type)
 	return -1;
 }
 
+char *CustomData_get_layer_name_at_offset(const CustomData *data, int type, int offset)
+{
+	int first;
+
+	first = CustomData_get_layer_index(data, type);
+
+	if((first != -1) &&
+	   (first + offset < data->totlayer) &&
+	   (data->layers[first + offset].type == type))
+		return data->layers[first + offset].name;
+	else
+		return NULL;
+}
+
 void CustomData_set_layer_active(CustomData *data, int type, int n)
 {
 	int i;
@@ -1182,6 +1213,17 @@ void CustomData_set_layer_flag(struct CustomData *data, int type, int flag)
 			data->layers[i].flag |= flag;
 }
 
+void CustomData_set_layer_offset_flag(struct CustomData *data, int type, int offset, int flag)
+{
+	int first = CustomData_get_layer_index(data, type);
+
+	if((first != -1) &&
+	   (first+offset < data->totlayer) &&
+	   (data->layers[first+offset].type == type)) {
+		data->layers[first+offset].flag |= flag;
+	}
+}
+
 static int customData_resize(CustomData *data, int amount)
 {
 	CustomDataLayer *tmp = MEM_callocN(sizeof(*tmp)*(data->maxlayer + amount),
@@ -2374,10 +2416,204 @@ int CustomData_verify_versions(struct CustomData *data, int index)
 	return keeplayer;
 }
 
-/* Subsurf grids */
-void CustomData_set_num_grid_elements(CustomData *data, int grid_elems)
+/* Multires */
+
+int CustomData_multires_type_totfloat(int type)
+{
+	switch(type) {
+	case CD_MVERT:
+		return 3;
+	case CD_MCOL:
+		return 4;
+	case CD_PAINTMASK:
+		return 1;
+	default:
+		return 0;
+	}
+}
+
+static int customdata_multires_find_first_layer_index(CustomDataMultires *cdm, int type)
+{
+	int i;
+	for(i = 0; i < cdm->totlayer; ++i) {
+		if(cdm->layers[i].type == type)
+			return i;
+	}
+	return -1;
+}
+
+static int customdata_multires_find_named_layer_index(CustomDataMultires *cdm,
+						      int type, char *name)
+{
+	int first, i;
+
+	first = customdata_multires_find_first_layer_index(cdm, type);
+
+	if(first != -1) {
+		for(i = first; first < cdm->totlayer; ++i) {
+			if(cdm->layers[i].type != type)
+				break;
+			else if(!strcmp(name, cdm->layers[i].name))
+				return i;
+		}
+	}
+
+	return -1;
+}
+
+int CustomData_get_multires_count(CustomData *cd, int type)
+{
+	int first, tot, i;
+
+	first = CustomData_get_layer_index(cd, type);
+	tot = 0;
+
+	if(first != -1) {
+		for(i = first; i < cd->totlayer; ++i) {
+			if(cd->layers[i].type != type)
+				break;
+			if(cd->layers[i].flag & CD_FLAG_MULTIRES)
+				++tot;
+		}
+	}
+
+	return tot;
+}
+
+void *CustomData_get_multires_names(CustomData *cd, int type)
 {
-	data->grid_elems = grid_elems;
+	char (*names)[32] = NULL;
+	int count, first;
+	
+	count = CustomData_get_multires_count(cd, type);
+
+	if(count) {
+		int layer_ndx, names_ndx;
+
+		names = MEM_callocN(32*count, "CustomData_get_multires_names");
+		first = CustomData_get_layer_index(cd, type);
+		names_ndx = 0;
+
+		for(layer_ndx = first; layer_ndx < cd->totlayer; ++layer_ndx) {
+			CustomDataLayer *cdl = cd->layers + layer_ndx;
+
+			if(cdl->type != type)
+				break;
+			else if(cdl->flag & CD_FLAG_MULTIRES) {
+				BLI_strncpy(names[names_ndx], cdl->name, 32);
+				++names_ndx;
+			}
+		}
+	}
+
+	return names;
+}
+
+float *CustomData_multires_get_data(CustomDataMultires *cdm, int type,
+				    char *name)
+{
+	int layer;
+
+	layer = customdata_multires_find_named_layer_index(cdm, type, name);
+
+	if(layer == -1)
+		return NULL;
+	else
+		return cdm->layers[layer].griddata;
+}
+
+void CustomData_multires_assign_data(CustomDataMultires *cdm, int type,
+				     char *name, float *data)
+{
+	int layer;
+
+	layer = customdata_multires_find_named_layer_index(cdm, type, name);
+
+	if(layer == -1)
+		CustomData_multires_add_layer(cdm, type, name, data);
+	else {
+		if(cdm->layers[layer].griddata)
+			MEM_freeN(cdm->layers[layer].griddata);
+		cdm->layers[layer].griddata = data;
+	}
+}
+
+void CustomData_multires_add_layer(CustomDataMultires *cdm, int type,
+				   char *name, float *data)
+{
+	CustomDataMultiresLayer *old = cdm->layers;
+	int first, layer, i;
+
+	cdm->layers = MEM_callocN(sizeof(CustomDataMultiresLayer) *
+				  (cdm->totlayer + 1),
+				  "customdata multires add layer");
+
+	first = customdata_multires_find_first_layer_index(cdm, type);
+
+	/* if no layers of type yet, add new layer at end
+	   otherwise add layer at the beginning of the type's segment */
+	if(first == -1)
+		layer = cdm->totlayer;
+	else
+		layer = first;
+
+	for(i = 0; i <= cdm->totlayer; ++i) {
+		if(i < layer)
+			cdm->layers[i] = old[i];
+		else if(i == layer) {
+			cdm->layers[i].griddata = data;
+			cdm->layers[i].type = type;
+			BLI_strncpy(cdm->layers[i].name, name,
+				    sizeof(cdm->layers[i].name));
+		}
+		else if(i > layer)
+			cdm->layers[i] = old[i-1];
+	}
+
+	++cdm->totlayer;
+	if(old) MEM_freeN(old);
+}
+
+int CustomData_multires_remove_layer(CustomDataMultires *cdm, int type,
+				      char *name)
+{
+	CustomDataMultiresLayer *old = cdm->layers;
+	int layer, i;
+
+	layer = customdata_multires_find_named_layer_index(cdm, type, name);
+
+	if(layer == -1)
+		return 0;
+
+	cdm->layers = MEM_callocN(sizeof(CustomDataMultiresLayer) *
+				  (cdm->totlayer - 1),
+				  "customdata multires remove layer");
+
+	/* copy over layer data, skipping the removed layer */
+	for(i = 0; i < cdm->totlayer; ++i) {
+		if(i < layer)
+			cdm->layers[i] = old[i];
+		else if(i > layer)
+			cdm->layers[i - 1] = old[i];
+	}
+
+	--cdm->totlayer;
+	MEM_freeN(old);
+	
+	return 1;
+}
+
+void CustomData_multires_rename(CustomDataMultires *cdm, int type,
+				char *old_name, char *name)
+{
+	int layer;
+
+	layer = customdata_multires_find_named_layer_index(cdm, type, old_name);
+
+	if(layer != -1) {
+		BLI_strncpy(cdm->layers[layer].name, name,
+			    sizeof(cdm->layers[layer].name));
+	}
 }
 
 /****************************** External Files *******************************/
diff --git a/source/blender/blenkernel/intern/multires.c b/source/blender/blenkernel/intern/multires.c
index 56b9b1b2c81..1d8f9c54f73 100644
--- a/source/blender/blenkernel/intern/multires.c
+++ b/source/blender/blenkernel/intern/multires.c
@@ -298,39 +298,60 @@ static void multires_set_tot_mdisps(Mesh *me, int lvl)
 	}
 }
 
-static void multires_reallocate_mdisps(Mesh *me, MDisps *mdisps, GridKey *gridkey, int lvl)
+static void multires_sync_customdata_layer(CustomDataMultires *cdm,
+					   char *name, int type, int totelem)
 {
-	int i;
+	float *griddata;
+	int totfloat;
+
+	totfloat = CustomData_multires_type_totfloat(type);
+
+	griddata = MEM_callocN(sizeof(float) * totfloat * totelem,
+			       "sync layer griddata");
+			
+	CustomData_multires_assign_data(cdm, type, name, griddata);
+}
+
+/* ensure all the layers needed by the gridkey have
+   a matching multires layer
 
-	CustomData *cd_facegrids;
+   clear and resize the all griddata to match the level
+*/
+static void multires_sync_customdata(Mesh *me, GridKey *gridkey,
+				     int lvl)
+{
+	CustomDataMultires *cd_grids;
+	int i, j;
 
-	cd_facegrids = CustomData_get_layer(&me->fdata, CD_FACEGRID);
-	if(!cd_facegrids)
-		cd_facegrids = CustomData_add_layer(&me->fdata, CD_FACEGRID, CD_CALLOC, NULL, me->totface);
+	cd_grids = CustomData_get_layer(&me->fdata, CD_GRIDS);
+	if(!cd_grids)
+		cd_grids = CustomData_add_layer(&me->fdata, CD_GRIDS,
+						CD_CALLOC, NULL, me->totface);
 
 	for(i = 0; i < me->totface; ++i) {
 		int nvert = (me->mface[i].v4)? 4: 3;
 		int totelem = multires_grid_tot[lvl]*nvert;
-		int pmask_totlayer;
-		CustomData old, *cd = cd_facegrids + i;
-
-		/* Resize all existing layers */
-		old = *cd;
-		memset(cd, 0, sizeof(*cd));
-		CustomData_copy(&old, cd, ~0, CD_CALLOC, totelem);
-		CustomData_free(&old, 0);
-		CustomData_set_num_grid_elements(cd, totelem);
-
-		/* If multires modifier is added after mask layers were
-		   created, update the grids to have those layers as well */
-		if(gridkey) {
-			pmask_totlayer = CustomData_number_of_layers(cd, CD_PAINTMASK);
-			while(pmask_totlayer < gridkey->mask) {
-				CustomData_add_layer(cd, CD_PAINTMASK, CD_CALLOC, NULL, totelem);
-				++pmask_totlayer;
-			}
+
+		for(j = 0; j < gridkey->color; ++j) {
+			multires_sync_customdata_layer(cd_grids+i,
+						       gridkey->color_names[j],
+						       CD_MCOL, totelem);
+		}
+
+		for(j = 0; j < gridkey->mask; ++j) {
+			multires_sync_customdata_layer(cd_grids+i,
+						       gridkey->mask_names[j],
+						       CD_PAINTMASK, totelem);
 		}
+
+		cd_grids[i].totelem = totelem;
 	}
+}
+
+/* TODO: removed this in favor of sync_customdata */
+static void multires_reallocate_mdisps(Mesh *me, MDisps *mdisps, int lvl)
+{
+	int i;
 
 	/* This will be replaced when we do CD_DISPS */
 
@@ -465,15 +486,33 @@ static DerivedMesh *multires_dm_create_local(Object *ob, DerivedMesh *dm,
 	return multires_dm_create_from_derived(&mmd, 1, dm, ob, gridkey, 0, 0);
 }
 
+static void init_gridkey_from_customdata(GridKey *gridkey,
+					 CustomData *vdata,
+					 CustomData *fdata)
+{
+ 	GRIDELEM_KEY_INIT(gridkey, 1,
+			  CustomData_get_multires_count(fdata, CD_MCOL),
+			  CustomData_get_multires_count(vdata, CD_PAINTMASK),
+			  1);
+
+	gridkey->color_names = CustomData_get_multires_names(fdata,
+							     CD_MCOL);
+
+	gridkey->mask_names = CustomData_get_multires_names(vdata,
+							    CD_PAINTMASK);
+			  
+}
+
 static DerivedMesh *subsurf_dm_create_local(Object *ob, DerivedMesh *dm,
 					    GridKey *gridkey, int lvl,
 					    int simple, int optimal)
 {
+	DerivedMesh *result;
+	Mesh *me = get_mesh(ob);
 	SubsurfModifierData smd;
 	GridKey default_gridkey;
-	int color_totlayer;
-	int pmask_totlayer;
 
+	memset(&default_gridkey, 0, sizeof(GridKey));
 	memset(&smd, 0, sizeof(SubsurfModifierData));
 	smd.levels = smd.renderLevels = lvl;
 	smd.flags |= eSubsurfModifierFlag_SubsurfUv;
@@ -483,16 +522,17 @@ static DerivedMesh *subsurf_dm_create_local(Object *ob, DerivedMesh *dm,
 		smd.flags |= eSubsurfModifierFlag_ControlEdges;
 
 	if(!gridkey) {
-		/* TODO: enable/disable element types */
-		color_totlayer = CustomData_number_of_layers(&get_mesh(ob)->fdata,
-							     CD_MCOL);
-		pmask_totlayer = CustomData_number_of_layers(&get_mesh(ob)->vdata,
-							     CD_PAINTMASK);
-		GRIDELEM_KEY_INIT(&default_gridkey, 1, color_totlayer, pmask_totlayer, 1);
+		init_gridkey_from_customdata(&default_gridkey,
+					     &me->vdata, &me->fdata);
 		gridkey = &default_gridkey;
 	}
 			  
-	return subsurf_make_derived_from_derived(dm, &smd, gridkey, 0, NULL, 0, 0);
+	result = subsurf_make_derived_from_derived(dm, &smd, gridkey, 0, NULL, 0, 0);
+
+	if(default_gridkey.color_names) MEM_freeN(default_gridkey.color_names);
+	if(default_gridkey.mask_names) MEM_freeN(default_gridkey.mask_names);
+	
+	return result;
 }
 
 /* assumes no is normalized; return value's sign is negative if v is on
@@ -683,7 +723,8 @@ void multiresModifier_subdivide(MultiresModifierData *mmd, Object *ob, int updat
 		ccgSubSurf_updateLevels(ss, lvl, NULL, 0);
 
 		/* reallocate displacements */
-		multires_reallocate_mdisps(me, mdisps, NULL, totlvl); 
+		multires_reallocate_mdisps(me, mdisps, totlvl); 
+		multires_sync_customdata(me, gridkey, totlvl);
 
 		/* compute displacements */
 		multiresModifier_disp_run(highdm, me, CALC_DISPS, subGridData, totlvl);
@@ -695,8 +736,11 @@ void multiresModifier_subdivide(MultiresModifierData *mmd, Object *ob, int updat
 		MEM_freeN(subGridData);
 	}
 	else {
+		/* XXX: I think this can be safely removed, we already check in
+		   disp run for unallocated disps -- nicholas */
+
 		/* only reallocate, nothing to upsample */
-		multires_reallocate_mdisps(me, mdisps, NULL, totlvl); 
+		//multires_reallocate_mdisps(me, mdisps, totlvl); 
 	}
 
 	multires_set_tot_level(ob, mmd, totlvl);
@@ -747,13 +791,46 @@ static void debug_print_paintmask_grids(CustomData *grids, int gridsize)
 }
 #endif
 
+/* XXX - move these to blenlib? */
+void add_v4_v4v4(float v[4], float a[4], float b[4])
+{
+	v[0] = a[0] + b[0];
+	v[1] = a[1] + b[1];
+	v[2] = a[2] + b[2];
+	v[3] = a[3] + b[3];
+}
+
+void add_v4_v4(float v[4], float a[4])
+{
+	v[0] += a[0];
+	v[1] += a[1];
+	v[2] += a[2];
+	v[3] += a[3];
+}
+
+void sub_v4_v4v4(float v[4], float a[4], float b[4])
+{
+	v[0] = a[0] - b[0];
+	v[1] = a[1] - b[1];
+	v[2] = a[2] - b[2];
+	v[3] = a[3] - b[3];
+}
+
+void clamp_v4_fl(float v[4], float min, float max)
+{
+	CLAMP(v[0], min, max);
+	CLAMP(v[1], min, max);
+	CLAMP(v[2], min, max);
+	CLAMP(v[3], min, max);
+}
+
 static void multiresModifier_disp_run(DerivedMesh *dm, Mesh *me, DispOp op, DMGridData **oldGridData, int totlvl)
 {
 	CCGDerivedMesh *ccgdm = (CCGDerivedMesh*)dm;
 	DMGridData **gridData, **subGridData;
 	MFace *mface = me->mface;
 	MDisps *mdisps = CustomData_get_layer(&me->fdata, CD_MDISPS);
-	CustomData *stored_grids;
+	CustomDataMultires *stored_grids;
 	int *gridOffset;
 	GridKey *gridkey;
 	int i, numGrids, gridSize, dGridSize, dSkip;
@@ -775,8 +852,8 @@ static void multiresModifier_disp_run(DerivedMesh *dm, Mesh *me, DispOp op, DMGr
 	dGridSize = multires_side_tot[totlvl];
 	dSkip = (dGridSize-1)/(gridSize-1);
 
-	stored_grids = CustomData_get_layer(&me->fdata, CD_FACEGRID);
-	
+	stored_grids = CustomData_get_layer(&me->fdata, CD_GRIDS);
+
 	#pragma omp parallel for private(i) if(me->totface*gridSize*gridSize*4 >= CCG_OMP_LIMIT)
 	for(i = 0; i < me->totface; ++i) {
 		const int numVerts = mface[i].v4 ? 4 : 3;
@@ -786,19 +863,15 @@ static void multiresModifier_disp_run(DerivedMesh *dm, Mesh *me, DispOp op, DMGr
 		/* when adding new faces in edit mode, need to allocate disps;
 		   may need to allocate paintmask storage after adding multires as well */
 		if(!mdisp->disps ||
-		   (gridkey->mask && 
-		    (!stored_grids ||
-		     CustomData_number_of_layers(&stored_grids[i], CD_PAINTMASK) != gridkey->mask)))
+		   ((gridkey->mask || gridkey->color) && 
+		    !stored_grids /* XXX: this needs updating for non vert types */ ))
 		#pragma omp critical
 		{
-			multires_reallocate_mdisps(me, mdisps, gridkey, totlvl);
+			multires_reallocate_mdisps(me, mdisps, totlvl);
+			multires_sync_customdata(me, gridkey, totlvl);
+			stored_grids = CustomData_get_layer(&me->fdata, CD_GRIDS);
 		}
 
-		/* Check masks */
-		assert(gridkey->mask == 0 || stored_grids);
-		if(stored_grids)
-			assert(CustomData_number_of_layers(&stored_grids[i], CD_PAINTMASK) == gridkey->mask);
-
 		for(S = 0; S < numVerts; ++S, ++gIndex) {
 			DMGridData *grid = gridData[gIndex];
 			DMGridData *subgrid = subGridData[gIndex];
@@ -852,14 +925,15 @@ static void multiresModifier_disp_run(DerivedMesh *dm, Mesh *me, DispOp op, DMGr
 						add_v3_v3(data, d);
 						break;
 					}
-
 					
 					/* Paint Masks */
 					for(j = 0; j < gridkey->mask; ++j) {
 						float *mask = &GRIDELEM_MASK_AT(grid, ccgdm_offset, gridkey)[j];
 						float *smask = &GRIDELEM_MASK_AT(subgrid, ccgdm_offset, gridkey)[j];
-						float *stored_mask_layer = CustomData_get_layer_n(&stored_grids[i],
-												  CD_PAINTMASK, j);
+						float *stored_mask_layer =
+							CustomData_multires_get_data(&stored_grids[i],
+										     CD_PAINTMASK,
+										     gridkey->mask_names[j]);
 						float *stored_mask = &stored_mask_layer[stored_index];
 
 						switch(op) {
@@ -876,6 +950,31 @@ static void multiresModifier_disp_run(DerivedMesh *dm, Mesh *me, DispOp op, DMGr
 							break;
 						}
 					}
+
+					/* Colors */
+					for(j = 0; j < gridkey->color; ++j) {
+						float *color = GRIDELEM_COLOR_AT(grid, ccgdm_offset, gridkey)[j];
+						float *scolor = GRIDELEM_COLOR_AT(subgrid, ccgdm_offset, gridkey)[j];
+						float *stored_color_layer =
+							CustomData_multires_get_data(&stored_grids[i],
+										     CD_MCOL,
+										     gridkey->color_names[j]);
+						float *stored_color = &stored_color_layer[(stored_index+
+											   stored_offset)*4];
+
+						switch(op) {
+						case APPLY_DISPS:
+							add_v4_v4v4(color, scolor, stored_color);
+							clamp_v4_fl(color, 0, 1);
+							break;
+						case CALC_DISPS:
+							sub_v4_v4v4(stored_color, color, scolor);
+							break;
+						case ADD_DISPS:
+							add_v4_v4(stored_color, color);
+							break;
+						}
+					}
 				}
 			}
 		}
diff --git a/source/blender/blenkernel/intern/subsurf_ccg.c b/source/blender/blenkernel/intern/subsurf_ccg.c
index 5a41721a4ab..dfefb309d4b 100644
--- a/source/blender/blenkernel/intern/subsurf_ccg.c
+++ b/source/blender/blenkernel/intern/subsurf_ccg.c
@@ -119,6 +119,8 @@ static CCGSubSurf *_getSubSurf(CCGSubSurf *prevSS, GridKey *gridkey, int subdivL
 	ifc.vertDataSize = GRIDELEM_SIZE(gridkey);
 	ifc.finterpCount = GRIDELEM_INTERP_COUNT(gridkey);
 	ifc.gridkey = *gridkey;
+	ifc.gridkey.color_names = MEM_dupallocN(gridkey->color_names);
+	ifc.gridkey.mask_names = MEM_dupallocN(gridkey->mask_names);
 
 	if (useArena) {
 		CCGAllocatorIFC allocatorIFC;
@@ -443,7 +445,7 @@ static void ss_sync_from_derivedmesh(CCGSubSurf *ss, DerivedMesh *dm,
 	MFace *mf;
 	float *vertData;
 	GridKey *gridkey = ccgSubSurf_getGridKey(ss);
-	float (*colors)[3] = NULL;
+	float (*colors)[4] = NULL;
 	int pmask_layer_count, pmask_first_layer;
 
 	ccgSubSurf_initFullSync(ss);
@@ -460,14 +462,10 @@ static void ss_sync_from_derivedmesh(CCGSubSurf *ss, DerivedMesh *dm,
 
 	/* for editable subdivided colors, find the average mcol for each vert */
 	if(gridkey->color) {
-		int mcol_totlayer, mcol_first_layer;
 		int *users;
 		int k;
 
-		mcol_totlayer = CustomData_number_of_layers(&dm->faceData, CD_MCOL);
-		mcol_first_layer = CustomData_get_layer_index(&dm->faceData, CD_MCOL);
-
-		colors = MEM_callocN(sizeof(float)*3*mcol_totlayer*totvert,
+		colors = MEM_callocN(sizeof(float)*4*gridkey->color*totvert,
 				     "ss_sync_from_derivedmesh.colors");
 		users = MEM_callocN(sizeof(int)*totvert,
 				    "ss_sync_from_derivedmesh.users");
@@ -481,23 +479,29 @@ static void ss_sync_from_derivedmesh(CCGSubSurf *ss, DerivedMesh *dm,
 
 				++users[vndx];
 
-				for(k = 0; k < mcol_totlayer; ++k) {
-					MCol *mcol = dm->faceData.layers[mcol_first_layer+k].data;
+				for(k = 0; k < gridkey->color; ++k) {
+					MCol *mcol = CustomData_get_layer_named(&dm->faceData,
+										CD_MCOL,
+										gridkey->color_names[k]);
 
-					colors[vndx*mcol_totlayer + k][0] = mcol[i*4+j].r;
-					colors[vndx*mcol_totlayer + k][1] = mcol[i*4+j].g;
-					colors[vndx*mcol_totlayer + k][2] = mcol[i*4+j].b;
+					colors[vndx*gridkey->color + k][0] += mcol[i*4+j].b;
+					colors[vndx*gridkey->color + k][1] += mcol[i*4+j].g;
+					colors[vndx*gridkey->color + k][2] += mcol[i*4+j].r;
+					colors[vndx*gridkey->color + k][3] += mcol[i*4+j].a;
 
 				}
 			}
 		}
 
+		/* divide by number of faces sharing the corner
+		   also convert from [0,255] to [0,1] */
 		for(i = 0; i < totvert; ++i) {
-			float inv = 1.0f / users[i];
-			for(j = 0; j < mcol_totlayer; ++j) {
-				colors[i*mcol_totlayer + j][0] *= inv;
-				colors[i*mcol_totlayer + j][1] *= inv;
-				colors[i*mcol_totlayer + j][2] *= inv;
+			float inv = 1.0f / (users[i] * 255);
+			for(j = 0; j < gridkey->color; ++j) {
+				colors[i*gridkey->color + j][0] *= inv;
+				colors[i*gridkey->color + j][1] *= inv;
+				colors[i*gridkey->color + j][2] *= inv;
+				colors[i*gridkey->color + j][3] *= inv;
 			}
 		}
 
@@ -511,11 +515,11 @@ static void ss_sync_from_derivedmesh(CCGSubSurf *ss, DerivedMesh *dm,
 
 		/* copy color data */
 		for(j = 0; j < gridkey->color; ++j)
-			memcpy(&vertData[3 + 3*j], colors[i*gridkey->color + j], sizeof(float)*3);
+			memcpy(&vertData[3 + 4*j], colors[i*gridkey->color + j], sizeof(float)*4);
 
 		/* copy paint mask data */
 		for(j = 0; j < gridkey->mask; ++j)
-			vertData[3 + gridkey->color * 3 + j] = ((float*)dm->vertData.layers[pmask_first_layer+j].data)[i];
+			vertData[3 + gridkey->color*3 + j] = ((float*)dm->vertData.layers[pmask_first_layer+j].data)[i];
 
 		ccgSubSurf_syncVert(ss, SET_INT_IN_POINTER(i), vertData, 0, &v);
 
@@ -2370,6 +2374,7 @@ static struct PBVH *ccgDM_getPBVH(Object *ob, DerivedMesh *dm)
 	CCGDerivedMesh *ccgdm= (CCGDerivedMesh*)dm;
 	int gridSize, numGrids, grid_pbvh;
 	GridKey *gridkey;
+	Mesh *me;
 
 	if(!ob) {
 		ccgdm->pbvh= NULL;
@@ -2401,6 +2406,8 @@ static struct PBVH *ccgDM_getPBVH(Object *ob, DerivedMesh *dm)
 	if(ccgdm->pbvh)
 		return ccgdm->pbvh;
 
+	me = ob->data;
+
 	/* no pbvh exists yet, we need to create one. only in case of multires
 	   we build a pbvh over the modified mesh, in other cases the base mesh
 	   is being sculpted, so we build a pbvh from that. */
@@ -2414,13 +2421,11 @@ static struct PBVH *ccgDM_getPBVH(Object *ob, DerivedMesh *dm)
 		ob->paint->pbvh= ccgdm->pbvh = BLI_pbvh_new();
 		BLI_pbvh_build_grids(ccgdm->pbvh, ccgdm->gridData, ccgdm->gridAdjacency,
 				     numGrids, gridSize, gridkey, (void**)ccgdm->gridFaces,
-				     &get_mesh(ob)->vdata,
+				     &me->vdata, &me->fdata,
 				     ss ? &ss->hidden_areas : NULL);
 		ccgdm->pbvh_draw = 1;
 	}
 	else if(ob->type == OB_MESH) {
-		Mesh *me= ob->data;
-
 		ob->paint->pbvh= ccgdm->pbvh = BLI_pbvh_new();
 		BLI_pbvh_build_mesh(ccgdm->pbvh, me->mface, me->mvert,
 				    &me->vdata, &me->fdata, me->totface, me->totvert,
diff --git a/source/blender/blenlib/BLI_pbvh.h b/source/blender/blenlib/BLI_pbvh.h
index 614a7d8ca35..15871527c1c 100644
--- a/source/blender/blenlib/BLI_pbvh.h
+++ b/source/blender/blenlib/BLI_pbvh.h
@@ -76,7 +76,8 @@ void BLI_pbvh_build_mesh(PBVH *bvh, struct MFace *faces, struct MVert *verts,
 void BLI_pbvh_build_grids(PBVH *bvh, struct DMGridData **grids,
 			  struct DMGridAdjacency *gridadj, int totgrid,
 			  int gridsize, struct GridKey *gridkey, void **gridfaces,
-			  struct CustomData *vdata, ListBase *hidden_areas);
+			  struct CustomData *vdata, struct CustomData *fdata,
+			  ListBase *hidden_areas);
 void BLI_pbvh_free(PBVH *bvh);
 
 /* Hierarchical Search in the BVH, two methods:
diff --git a/source/blender/blenlib/intern/pbvh.c b/source/blender/blenlib/intern/pbvh.c
index d258412fc18..f26c6783892 100644
--- a/source/blender/blenlib/intern/pbvh.c
+++ b/source/blender/blenlib/intern/pbvh.c
@@ -678,7 +678,8 @@ void BLI_pbvh_build_mesh(PBVH *bvh, MFace *faces, MVert *verts,
 void BLI_pbvh_build_grids(PBVH *bvh, DMGridData **grids,
 			  DMGridAdjacency *gridadj,
 			  int totgrid, int gridsize, GridKey *gridkey,
-			  void **gridfaces, CustomData *vdata, ListBase *hidden_areas)
+			  void **gridfaces, CustomData *vdata,
+			  CustomData *fdata, ListBase *hidden_areas)
 {
 	bvh->grids= grids;
 	bvh->gridadj= gridadj;
@@ -687,6 +688,7 @@ void BLI_pbvh_build_grids(PBVH *bvh, DMGridData **grids,
 	bvh->gridsize= gridsize;
 	bvh->gridkey= gridkey;
 	bvh->vdata= vdata;
+	bvh->fdata= fdata;
 	bvh->leaf_limit = MAX2(LEAF_LIMIT/((gridsize-1)*(gridsize-1)), 1);
 
 	if(totgrid)
@@ -1347,7 +1349,7 @@ void BLI_pbvh_node_get_faces(PBVH *bvh, PBVHNode *node,
 {
 	if(bvh->grids) {
 		if(mface) *mface= NULL;
-		if(fdata) *fdata= NULL;
+		if(fdata) *fdata= bvh->fdata;
 		if(face_indices) *face_indices= NULL;
 		if(face_vert_indices) *face_vert_indices= NULL;
 		if(totnode) *totnode= 0;
diff --git a/source/blender/blenloader/intern/readfile.c b/source/blender/blenloader/intern/readfile.c
index 7bb19572ae2..4810d4b8c71 100644
--- a/source/blender/blenloader/intern/readfile.c
+++ b/source/blender/blenloader/intern/readfile.c
@@ -3262,9 +3262,25 @@ static void direct_link_mdisps(FileData *fd, int count, MDisps *mdisps, int exte
 	}       
 }
 
+static void direct_link_customdata_multires(FileData *fd, int count,
+					    CustomDataMultires *cdm)
+{
+	if(cdm) {
+		int i, j;
+
+		for(i = 0; i < count; ++i, ++cdm) {
+			cdm->layers = newdataadr(fd, cdm->layers);
+			for(j = 0; j < cdm->totlayer; ++j) {
+				CustomDataMultiresLayer *l = cdm->layers + j;
+				l->griddata = newdataadr(fd, l->griddata);
+			}
+		}
+	}
+}
+
 static void direct_link_customdata(FileData *fd, CustomData *data, int count)
 {
-	int i = 0, j;
+	int i = 0;
 
 	data->layers= newdataadr(fd, data->layers);
 	data->external= newdataadr(fd, data->external);
@@ -3279,11 +3295,8 @@ static void direct_link_customdata(FileData *fd, CustomData *data, int count)
 			layer->data = newdataadr(fd, layer->data);
 			if(layer->type == CD_MDISPS)
 				direct_link_mdisps(fd, count, layer->data, layer->flag & CD_FLAG_EXTERNAL);
-			if(layer->type == CD_FACEGRID) {
-				CustomData *grids = layer->data;
-				for(j = 0; j < count; ++j)
-					direct_link_customdata(fd, grids + j, grids[j].grid_elems);
-			}
+			if(layer->type == CD_GRIDS)
+				direct_link_customdata_multires(fd, count, layer->data);
 			i++;
 		}
 	}
diff --git a/source/blender/blenloader/intern/writefile.c b/source/blender/blenloader/intern/writefile.c
index cdf53119352..7b1018fc6f1 100644
--- a/source/blender/blenloader/intern/writefile.c
+++ b/source/blender/blenloader/intern/writefile.c
@@ -1484,10 +1484,28 @@ static void write_mdisps(WriteData *wd, int count, MDisps *mdlist, int external)
 	}
 }
 
-static void write_customdata(WriteData *wd, ID *id, int count, CustomData *data, int partial_type, int partial_count)
+static void write_customdata_multires(WriteData *wd, int count,
+				      CustomDataMultires *grids)
 {
 	int i, j;
 
+	writestruct(wd, DATA, "CustomDataMultires", count, grids);
+
+	for(i = 0; i < count; ++i) {
+		for(j = 0; j < grids[i].totlayer; ++j) {
+			CustomDataMultiresLayer *l = &grids[i].layers[j];
+
+			writedata(wd, DATA, sizeof(float) * grids[i].totelem *
+				  CustomData_multires_type_totfloat(l->type),
+				  l->griddata);
+		}
+	}
+}
+
+static void write_customdata(WriteData *wd, ID *id, int count, CustomData *data, int partial_type, int partial_count)
+{
+	int i;
+
 	/* write external customdata (not for undo) */
 	if(data->external && !wd->current)
 		CustomData_external_write(data, id, CD_MASK_MESH, count, 0);
@@ -1509,12 +1527,8 @@ static void write_customdata(WriteData *wd, ID *id, int count, CustomData *data,
 		else if (layer->type == CD_PAINTMASK) {
 			writedata(wd, DATA, sizeof(float)*count, layer->data);
 		}
-		else if (layer->type == CD_FACEGRID) {
-			CustomData *grids = layer->data;
-			writestruct(wd, DATA, "CustomData", count, grids);
-			for(j = 0; j < count; ++j)
-				write_customdata(wd, id, grids[j].grid_elems,
-						 grids + j, -1, 0);
+		else if (layer->type == CD_GRIDS) {
+			write_customdata_multires(wd, count, layer->data);
 		}
 		else {
 			CustomData_file_write_info(layer->type, &structname, &structnum);
diff --git a/source/blender/editors/mesh/mesh_data.c b/source/blender/editors/mesh/mesh_data.c
index e353c2bbff9..3f4ff1d5d05 100644
--- a/source/blender/editors/mesh/mesh_data.c
+++ b/source/blender/editors/mesh/mesh_data.c
@@ -469,6 +469,61 @@ void MESH_OT_vertex_color_remove(wmOperatorType *ot)
 	ot->flag= OPTYPE_REGISTER|OPTYPE_UNDO;
 }
 
+static int vertex_color_multires_toggle_exec(bContext *C, wmOperator *op)
+{
+	Object *ob= CTX_data_active_object(C);
+	Mesh *me= ob->data;
+	CustomDataMultires *cdm;
+	CustomDataLayer *cdl;
+	int active, i;
+
+	active = CustomData_get_active_layer_index(&me->fdata, CD_MCOL);
+	cdm = CustomData_get_layer(&me->fdata, CD_GRIDS);
+
+	if(active == -1)
+		return OPERATOR_FINISHED;
+
+	cdl = &me->fdata.layers[active];
+
+	if(cdm) {
+		if(cdl->flag & CD_FLAG_MULTIRES) {
+			/* delete multires data */
+			for(i = 0; i < me->totface; ++i)
+				CustomData_multires_remove_layer(cdm+i, CD_MCOL, cdl->name);
+		}
+		else {
+			/* add multires data */
+			for(i = 0; i < me->totface; ++i) {
+				float *data = MEM_callocN(sizeof(float) *
+							  CustomData_multires_type_totfloat(CD_MCOL) *
+							  cdm[i].totelem,
+							  "vertex_color_multires_toggle");
+				CustomData_multires_add_layer(cdm+i, CD_MCOL, cdl->name, data);
+			}
+		}
+	}
+
+	/* note - if there's no griddata, it can still be synced up later */
+	cdl->flag ^= CD_FLAG_MULTIRES;
+
+	return OPERATOR_FINISHED;
+}
+
+void MESH_OT_vertex_color_multiresolution_toggle(wmOperatorType *ot)
+{
+	/* identifiers */
+	ot->name= "Vertex Color Multiresolution Toggle";
+	ot->description= "Add or remove multiresolution data from this layer";
+	ot->idname= "MESH_OT_vertex_color_multiresolution_toggle";
+	
+	/* api callbacks */
+	ot->exec= vertex_color_multires_toggle_exec;
+	ot->poll= ED_mesh_layers_poll;
+
+	/* flags */
+	ot->flag= OPTYPE_REGISTER|OPTYPE_UNDO;
+}
+
 /*********************** sticky operators ************************/
 
 static int sticky_add_exec(bContext *C, wmOperator *op)
diff --git a/source/blender/editors/mesh/mesh_intern.h b/source/blender/editors/mesh/mesh_intern.h
index 6db2a3f211d..9b40b9cec0f 100644
--- a/source/blender/editors/mesh/mesh_intern.h
+++ b/source/blender/editors/mesh/mesh_intern.h
@@ -249,6 +249,7 @@ void MESH_OT_uv_texture_add(struct wmOperatorType *ot);
 void MESH_OT_uv_texture_remove(struct wmOperatorType *ot);
 void MESH_OT_vertex_color_add(struct wmOperatorType *ot);
 void MESH_OT_vertex_color_remove(struct wmOperatorType *ot);
+void MESH_OT_vertex_color_multiresolution_toggle(struct wmOperatorType *ot);
 void MESH_OT_sticky_add(struct wmOperatorType *ot);
 void MESH_OT_sticky_remove(struct wmOperatorType *ot);
 void MESH_OT_drop_named_image(struct wmOperatorType *ot);
diff --git a/source/blender/editors/mesh/mesh_ops.c b/source/blender/editors/mesh/mesh_ops.c
index e143e4859f5..6c82e422b48 100644
--- a/source/blender/editors/mesh/mesh_ops.c
+++ b/source/blender/editors/mesh/mesh_ops.c
@@ -137,6 +137,7 @@ void ED_operatortypes_mesh(void)
 	WM_operatortype_append(MESH_OT_uv_texture_remove);
 	WM_operatortype_append(MESH_OT_vertex_color_add);
 	WM_operatortype_append(MESH_OT_vertex_color_remove);
+	WM_operatortype_append(MESH_OT_vertex_color_multiresolution_toggle);
 	WM_operatortype_append(MESH_OT_sticky_add);
 	WM_operatortype_append(MESH_OT_sticky_remove);
 	WM_operatortype_append(MESH_OT_drop_named_image);
diff --git a/source/blender/editors/sculpt_paint/paint_mask.c b/source/blender/editors/sculpt_paint/paint_mask.c
index 27195bf9b87..7c6ac305cbc 100644
--- a/source/blender/editors/sculpt_paint/paint_mask.c
+++ b/source/blender/editors/sculpt_paint/paint_mask.c
@@ -29,6 +29,7 @@
 #include "BLI_listbase.h"
 #include "BLI_math.h"
 #include "BLI_pbvh.h"
+#include "BLI_string.h"
 
 #include "ED_mesh.h"
 #include "ED_sculpt.h"
@@ -405,49 +406,46 @@ static int paintmask_check_multires(bContext *C)
 	return 0;
 }
 
-/* If this is a multires mesh, update it and free the DM, then add or remove
-   a paintmask layer from the grid layer */
+/* When adding paintmasks, assume we want them subdivided for multires */
 static void paintmask_adjust_multires(bContext *C,
 				      PaintMaskLayerOp op,
-				      int layer_offset,
+				      char *layer_name,
 				      float **removed_multires_data)
 {
 	Object *ob = CTX_data_active_object(C);
 	Mesh *me = get_mesh(ob);
-	CustomData *cd = CustomData_get_layer(&me->fdata, CD_FACEGRID);
-	int pmask_first_layer, i;
+	CustomDataMultires *cdm = CustomData_get_layer(&me->fdata, CD_GRIDS);
+	float *griddata;
+	int i;
 
-	assert(cd);
+	assert(cdm && layer_name);
 
 	for(i = 0; i < me->totface; ++i) {
-		pmask_first_layer = CustomData_get_layer_index(cd + i,
-							       CD_PAINTMASK);
-
 		switch(op) {
 		case LAYER_ADDED:
 			/* Add a layer of paintmask from grids */
 
-			assert(!removed_multires_data || layer_offset != -1);
-
 			/* if restoring from undo, copy the old data
 			   back into CustomData */
 			if(removed_multires_data)
-				CustomData_add_layer_at_offset(cd + i, CD_PAINTMASK,
-							       CD_ASSIGN,
-							       removed_multires_data[i],
-							       cd[i].grid_elems,
-							       layer_offset);
-			else							     
-				CustomData_add_layer(cd + i, CD_PAINTMASK,
-						     CD_CALLOC, NULL,
-						     cd[i].grid_elems);
+				griddata = removed_multires_data[i];
+			else
+				griddata = MEM_callocN(sizeof(float) *
+						       cdm->totelem,
+						       "paintmask griddata");
+
+			CustomData_multires_add_layer(cdm + i,
+						      CD_PAINTMASK,
+						      layer_name,
+						      griddata);
+
 			break;
 
 		case LAYER_REMOVED:
 			/* Remove a layer of paintmask from grids */
-			CustomData_free_layer(cd + i, CD_PAINTMASK,
-					      cd[i].grid_elems,
-					      pmask_first_layer + layer_offset);
+			
+			CustomData_multires_remove_layer(cdm + i, CD_PAINTMASK,
+							 layer_name);
 			break;
 		}
 	}
@@ -467,9 +465,9 @@ typedef struct PaintMaskUndoNode {
 	int totface;
 } PaintMaskUndoNode;
 
-/* copy the mesh/multires data for paintmask at layer_offset into unode */
+/* copy the mesh/multires data for named paintmask layer into unode */
 static void paintmask_undo_backup_layer(bContext *C, PaintMaskUndoNode *unode,
-					int layer_offset)
+					char *layer_name)
 {
 	struct Scene *scene = CTX_data_scene(C);
 	Object *ob = CTX_data_active_object(C);
@@ -482,17 +480,17 @@ static void paintmask_undo_backup_layer(bContext *C, PaintMaskUndoNode *unode,
 	/* if removing a layer, make a copy of the mask
 	   data for restoring from undo */
 	unode->removed_data =
-		MEM_dupallocN(CustomData_get_layer_n(&me->vdata,
-						     CD_PAINTMASK,
-						     layer_offset));
+		MEM_dupallocN(CustomData_get_layer_named(&me->vdata,
+							 CD_PAINTMASK,
+							 layer_name));
+
 
-	strcpy(unode->layer_name,
-	       me->vdata.layers[pmask_first_layer+layer_offset].name);
+	/* XXX: need to deal somewhere with possibility of layer name-change */
 
 	if(paint_multires_active(scene, ob)) {
 		/* need to copy active layer of multires data too */
-		CustomData *grids = CustomData_get_layer(&me->fdata,
-							 CD_FACEGRID);
+		CustomDataMultires *grids = CustomData_get_layer(&me->fdata,
+								 CD_GRIDS);
 		int i;
 
 		unode->totface = me->totface;
@@ -502,9 +500,9 @@ static void paintmask_undo_backup_layer(bContext *C, PaintMaskUndoNode *unode,
 
 		for(i = 0; i < me->totface; ++i) {
 			unode->removed_multires_data[i] =
-				MEM_dupallocN(CustomData_get_layer_n(grids + i,
-								     CD_PAINTMASK,
-								     layer_offset));
+				MEM_dupallocN(CustomData_multires_get_data(grids + i,
+									   CD_PAINTMASK,
+									   layer_name));
 		}
 	}
 }
@@ -522,10 +520,10 @@ static void paintmask_undo_restore(bContext *C, ListBase *lb)
 
 	switch(unode->type) {
 	case LAYER_ADDED:
-		paintmask_undo_backup_layer(C, unode, unode->layer_offset);
+		paintmask_undo_backup_layer(C, unode, unode->layer_name);
 
 		if(multires)
-			paintmask_adjust_multires(C, LAYER_REMOVED, unode->layer_offset, NULL);
+			paintmask_adjust_multires(C, LAYER_REMOVED, unode->layer_name, NULL);
 
 		CustomData_free_layer(vdata, CD_PAINTMASK, me->totvert,
 				      pmask_first_layer + unode->layer_offset);
@@ -533,7 +531,7 @@ static void paintmask_undo_restore(bContext *C, ListBase *lb)
 	case LAYER_REMOVED:
 		if(multires)
 			paintmask_adjust_multires(C, LAYER_ADDED,
-						  unode->layer_offset,
+						  unode->layer_name,
 						  unode->removed_multires_data);
 		CustomData_add_layer_at_offset(vdata,
 					       CD_PAINTMASK,
@@ -542,6 +540,7 @@ static void paintmask_undo_restore(bContext *C, ListBase *lb)
 					       me->totvert,
 					       unode->layer_offset);
 
+		CustomData_set_layer_offset_flag(vdata, CD_PAINTMASK, unode->layer_offset, CD_FLAG_MULTIRES);
 		CustomData_set_layer_active(vdata, CD_PAINTMASK, unode->layer_offset);
 		strcpy(vdata->layers[pmask_first_layer+unode->layer_offset].name, unode->layer_name);
 
@@ -561,21 +560,6 @@ static void paintmask_undo_restore(bContext *C, ListBase *lb)
 	paintmask_redraw(C);
 }
 
-static void paintmask_undo_push_node(bContext *C, PaintMaskLayerOp op, int layer_offset)
-{
-	ListBase *lb= undo_paint_push_get_list(UNDO_PAINT_MESH);
-	PaintMaskUndoNode *unode = MEM_callocN(sizeof(PaintMaskUndoNode),
-					       "PaintMaskUndoNode");
-
-	unode->type = op;
-	unode->layer_offset = layer_offset;
-
-	if(op == LAYER_REMOVED)
-		paintmask_undo_backup_layer(C, unode, layer_offset);
-
-	BLI_addtail(lb, unode);
-}
-
 static void paintmask_undo_free(ListBase *lb)
 {
 	PaintMaskUndoNode *unode = lb->first;
@@ -594,14 +578,27 @@ static void paintmask_undo_free(ListBase *lb)
 }
 
 static void paintmask_undo_push(bContext *C, char *name,
-			   PaintMaskLayerOp op, int layer_offset)
+				PaintMaskLayerOp op, char *layer_name,
+				int layer_offset)
 {
+	PaintMaskUndoNode *unode;
+
 	undo_paint_push_begin(UNDO_PAINT_MESH,
 			      name,
 			      paintmask_undo_restore,
 			      paintmask_undo_free);
 
-	paintmask_undo_push_node(C, op, layer_offset);
+	unode = MEM_callocN(sizeof(PaintMaskUndoNode), "PaintMaskUndoNode");
+
+	unode->type = op;
+	BLI_strncpy(unode->layer_name, layer_name, sizeof(unode->layer_name));
+	/* store offset so a removed layer can be restored at the same place */
+	unode->layer_offset = layer_offset;
+
+	if(op == LAYER_REMOVED)
+		paintmask_undo_backup_layer(C, unode, layer_name);
+
+	BLI_addtail(undo_paint_push_get_list(UNDO_PAINT_MESH), unode);
 
 	undo_paint_push_end(UNDO_PAINT_MESH);
 }
@@ -616,18 +613,27 @@ static int mask_layer_add_exec(bContext *C, wmOperator *op)
 	Object *ob= CTX_data_active_object(C);
 	Mesh *me= ob->data;
 	int multires, top;
+	char *layer_name;
 
 	multires = paintmask_check_multires(C);
 
-	paintmask_undo_push(C, "Add paint mask", LAYER_ADDED,
-			    CustomData_number_of_layers(&me->vdata, CD_PAINTMASK));
+	top= CustomData_number_of_layers(&me->vdata, CD_PAINTMASK);
+	CustomData_add_layer(&me->vdata, CD_PAINTMASK, CD_DEFAULT,
+			     NULL, me->totvert);
+
+	CustomData_set_layer_offset_flag(&me->vdata, CD_PAINTMASK, top, CD_FLAG_MULTIRES);
+	CustomData_set_layer_active(&me->vdata, CD_PAINTMASK, top);
+
+	layer_name = CustomData_get_layer_name_at_offset(&me->vdata,
+							 CD_PAINTMASK, top);
+
+	/* now that we have correct name, update multires and do undo push */
 
 	if(multires)
-		paintmask_adjust_multires(C, LAYER_ADDED, -1, NULL);
+		paintmask_adjust_multires(C, LAYER_ADDED, layer_name, NULL);
 
-	top= CustomData_number_of_layers(&me->vdata, CD_PAINTMASK);
-	CustomData_add_layer(&me->vdata, CD_PAINTMASK, CD_DEFAULT, NULL, me->totvert);
-	CustomData_set_layer_active(&me->vdata, CD_PAINTMASK, top);
+	paintmask_undo_push(C, "Add paint mask", LAYER_ADDED,
+			    layer_name, top);
 
 	paintmask_redraw(C);
 
@@ -659,12 +665,15 @@ static int mask_layer_remove_exec(bContext *C, wmOperator *op)
 
 	if(active_offset >= 0) {
 		int multires = paintmask_check_multires(C);
+		char *layer_name = CustomData_get_layer_name_at_offset(&me->vdata,
+								       CD_PAINTMASK,
+								       active_offset);
 
 		paintmask_undo_push(C, "Remove paint mask", LAYER_REMOVED,
-				    active_offset);
+				    layer_name, active_offset);
 
 		if(multires)
-			paintmask_adjust_multires(C, LAYER_REMOVED, active_offset, NULL);
+			paintmask_adjust_multires(C, LAYER_REMOVED, layer_name, NULL);
 
 		CustomData_free_layer_active(&me->vdata, CD_PAINTMASK, me->totvert);
 
diff --git a/source/blender/editors/sculpt_paint/paint_vertex.c b/source/blender/editors/sculpt_paint/paint_vertex.c
index 314de7ab4cd..f5283795f2e 100644
--- a/source/blender/editors/sculpt_paint/paint_vertex.c
+++ b/source/blender/editors/sculpt_paint/paint_vertex.c
@@ -66,6 +66,7 @@
 #include "BKE_global.h"
 #include "BKE_mesh.h"
 #include "BKE_modifier.h"
+#include "BKE_multires.h"
 #include "BKE_object.h"
 #include "BKE_paint.h"
 #include "BKE_utildefines.h"
@@ -1935,25 +1936,35 @@ static void vpaint_stroke_update_step_old(bContext *C, PaintStroke *stroke, Poin
 /* apply paint at specified coordinate
    returns 1 if paint was applied, 0 otherwise */
 static int vpaint_paint_coord(VPaint *vp, VPaintData *vpd, float co[3],
-			      unsigned int *col, unsigned int *orig_col,
+			      float col[4], float orig_col[4],
 			      float center[3], float radius,
 			      float radius_squared)
 {
 	Brush *brush = paint_brush(&vp->paint);
-	float str, dist, dist_squared;
+	float strength, dist, dist_squared;
+	float paint_col[4];
+
+	paint_col[0] = ((MCol*)&vpd->paintcol)->b / 255.0f;
+	paint_col[1] = ((MCol*)&vpd->paintcol)->g / 255.0f;
+	paint_col[2] = ((MCol*)&vpd->paintcol)->r / 255.0f;
+	paint_col[3] = ((MCol*)&vpd->paintcol)->a / 255.0f;
 
 	dist_squared = len_squared_v3v3(center, co);
 
 	if(dist_squared < radius_squared) {
 		dist = sqrtf(dist_squared);
 
-		str = brush->alpha *
+		strength = brush->alpha *
 			brush_curve_strength(brush, dist,
 					     radius);
 
-		vpaint_blend(vp, col, orig_col,
+		//printf("orig=%f,%f,%f,%f   ", col[0], col[1], col[2], col[3]);
+		IMB_blend_color_float(col, col, paint_col, strength, IMB_BLEND_MIX /* TODO */);
+		//printf("result=%f,%f,%f,%f\n", col[0], col[1], col[2], col[3]);
+
+		/*vpaint_blend(vp, col, orig_col,
 			     vpd->paintcol,
-			     str*255);
+			     str*255);*/
 
 		return 1;
 	}
@@ -1963,7 +1974,8 @@ static int vpaint_paint_coord(VPaint *vp, VPaintData *vpd, float co[3],
 
 static void vpaint_nodes_grids(VPaint *vp, VPaintData *vpd, DMGridData **grids,
 			       GridKey *gridkey, int *grid_indices, int totgrid,
-			       int gridsize, float center[3], float radius)
+			       int gridsize, int active, float center[3],
+			       float radius)
 {
 	float radius_squared = radius*radius;
 	int i, x, y;
@@ -1977,23 +1989,13 @@ static void vpaint_nodes_grids(VPaint *vp, VPaintData *vpd, DMGridData **grids,
 							       y*gridsize+x,
 							       gridkey);
 				float *co = GRIDELEM_CO(elem, gridkey);
-				float *gridcol = GRIDELEM_COLOR(elem, gridkey);
-				MCol col, orig_col;
-
-				col.r = gridcol[0] * 255;
-				col.g = gridcol[1] * 255;
-				col.b = gridcol[2] * 255;
-
-				if(vpaint_paint_coord(vp, vpd, co,
-					(unsigned int*)(&col),
-					(unsigned int*)(&orig_col),
-					center, radius,
-					radius_squared)) {
-
-					gridcol[0] = col.r * (1.0 / 255);
-					gridcol[1] = col.g * (1.0 / 255);
-					gridcol[2] = col.b * (1.0 / 255);
-				}
+				float *gridcol = GRIDELEM_COLOR(elem, gridkey)[active];
+
+				vpaint_paint_coord(vp, vpd, co,
+						   gridcol,
+						   NULL, /* TODO */
+						   center, radius,
+						   radius_squared);
 			}
 		}
 	}
@@ -2021,15 +2023,44 @@ static void vpaint_nodes_faces(VPaint *vp, VPaintData *vpd, MFace *mface,
 			int vndx = (&f->v1)[j];
 			int cndx = face_index*4 + j;
 			float *co = mvert[vndx].co;
-			unsigned int *col = (unsigned int*)(mcol + cndx);
-			unsigned int *orig_col = (unsigned int*)(orig + cndx);
+			//unsigned int *col = (unsigned int*)(mcol + cndx);
+			//unsigned int *orig_col = (unsigned int*)(orig + cndx);
+			float fcol[4];
+
+			fcol[0] = mcol[cndx].b / 255.0f;
+			fcol[1] = mcol[cndx].g / 255.0f;
+			fcol[2] = mcol[cndx].r / 255.0f;
+			fcol[3] = mcol[cndx].a / 255.0f;
 
-			vpaint_paint_coord(vp, vpd, co, col, orig_col, center,
+			vpaint_paint_coord(vp, vpd, co, fcol, NULL /* TODO */, center,
 					   radius, radius_squared);
+
+			mcol[cndx].b = fcol[0] * 255.0f;
+			mcol[cndx].g = fcol[1] * 255.0f;
+			mcol[cndx].r = fcol[2] * 255.0f;
+			mcol[cndx].a = fcol[3] * 255.0f;
 		}
 	}
 }
 
+static int vpaint_find_gridkey_active_layer(CustomData *fdata, GridKey *gridkey)
+{
+	int active, i;
+
+	active = CustomData_get_active_layer_index(fdata, CD_MCOL);
+	
+	if(active == -1)
+		return -1;
+
+	for(i = 0; i < gridkey->color; ++i) {
+		if(!strcmp(gridkey->color_names[i],
+			   fdata->layers[active].name))
+			return i;
+	}
+
+	return -1;
+}
+
 static void vpaint_nodes(VPaint *vp, VPaintData *vpd, PBVH *pbvh,
 			 PBVHNode **nodes, int totnode,
 			 float center[3], float radius)
@@ -2039,7 +2070,7 @@ static void vpaint_nodes(VPaint *vp, VPaintData *vpd, PBVH *pbvh,
 	for(n = 0; n < totnode; ++n) {
 		MVert *mvert;
 		MFace *mface;
-		CustomData *fdata;
+		CustomData *fdata = NULL;
 		int *face_indices, totface;
 
 		DMGridData **grids;
@@ -2056,9 +2087,15 @@ static void vpaint_nodes(VPaint *vp, VPaintData *vpd, PBVH *pbvh,
 					NULL, &gridkey);
 
 		if(grids) {
-			vpaint_nodes_grids(vp, vpd, grids, gridkey,
-					   grid_indices, totgrid, gridsize,
-					   center, radius);
+			int active = vpaint_find_gridkey_active_layer(fdata,
+								      gridkey);
+
+			if(active != -1) {
+				vpaint_nodes_grids(vp, vpd, grids, gridkey,
+						   grid_indices, totgrid,
+						   gridsize, active, center,
+						   radius);
+			}
 		}
 		else {
 			vpaint_nodes_faces(vp, vpd, mface, mvert, fdata,
@@ -2168,6 +2205,8 @@ static void vpaint_stroke_update_step_new(bContext *C, PaintStroke *stroke,
 		vpaint_nodes(vp, vpd, ob->paint->pbvh, nodes, totnode,
 			     center, radius);
 
+		multires_mark_as_modified(ob);
+
 		if(nodes)
 			MEM_freeN(nodes);
 	}
diff --git a/source/blender/gpu/intern/gpu_buffers.c b/source/blender/gpu/intern/gpu_buffers.c
index 7a3f4126b67..02ef0897768 100644
--- a/source/blender/gpu/intern/gpu_buffers.c
+++ b/source/blender/gpu/intern/gpu_buffers.c
@@ -446,9 +446,9 @@ static void delete_buffer(GLuint *buf)
 
 static void float_col_to_gpu_colors(unsigned char out[3], float col[3])
 {
-	out[0] = col[2] * 255;
+	out[0] = col[0] * 255;
 	out[1] = col[1] * 255;
-	out[2] = col[0] * 255;
+	out[2] = col[2] * 255;
 }
 
 static void mcol_to_gpu_colors(unsigned char out[3], MCol *mcol)
@@ -705,11 +705,18 @@ static void gpu_update_grid_color_buffers_from_mcol(GPU_Buffers *buffers, DMGrid
 
 			for(j = 0; j < gridsize*gridsize; ++j, color_data += 3) {
 				DMGridData *elem = GRIDELEM_AT(grid, j, gridkey);
-				float v[3] = {0, 0, 0};
+				float v[3] = {1, 1, 1};
 
-				for(k = 0; k < gridkey->color; ++k)
-					add_v3_v3(v, &GRIDELEM_COLOR(elem, gridkey)[k*3]);
-				mul_v3_fl(v, 1.0f / gridkey->color);
+				for(k = 0; k < gridkey->color; ++k) {
+					float *col = GRIDELEM_COLOR(elem, gridkey)[k];
+
+					/* for now we just combine layers in order
+					   interpolating using the alpha component
+					   ("order" is ill-defined here since we
+					   don't guarantee the order of cdm data) */
+					interp_v3_v3v3(v, v, col, col[3]);
+				}
+				// clamp?
 
 				float_col_to_gpu_colors(color_data, v);
 			}
diff --git a/source/blender/makesdna/DNA_customdata_types.h b/source/blender/makesdna/DNA_customdata_types.h
index 192f5022f51..33c536d8df6 100644
--- a/source/blender/makesdna/DNA_customdata_types.h
+++ b/source/blender/makesdna/DNA_customdata_types.h
@@ -55,11 +55,27 @@ typedef struct CustomData {
 	CustomDataLayer *layers;      /* CustomDataLayers, ordered by type */
 	int totlayer, maxlayer;       /* number of layers, size of layers array */
 	int totsize;                  /* in editmode, total size of all data layers */
-	int grid_elems;               /* For grids, number of elements */
+	int pad;
 	void *pool;                   /* Bmesh: Memory pool for allocation of blocks */
 	CustomDataExternal *external; /* external file storing customdata layers */
 } CustomData;
 
+typedef struct CustomDataMultiresLayer {
+	float *griddata;  /* storage for n grids, where n is face's number of corners */
+	char name[32];    /* source layer's name, should be kept in sync */
+	int type;         /* customdata type */
+	int pad;
+} CustomDataMultiresLayer;
+
+/* some CustomData layers can be subdivided to store multires data, but since
+   multires data is always related to faces and the source layers might be
+   vertex data, the subdivided data is stored separately
+*/
+typedef struct CustomDataMultires {
+	CustomDataMultiresLayer *layers;
+	int totlayer, totelem;
+} CustomDataMultires;
+
 /* CustomData.type */
 #define CD_MVERT		0
 #define CD_MSTICKY		1
@@ -85,10 +101,9 @@ typedef struct CustomData {
 #define CD_ID_MCOL		21
 #define CD_TEXTURE_MCOL	22
 #define CD_CLOTH_ORCO	23
-#define CD_FACEGRID		24
-#define CD_DISP			25
-#define CD_PAINTMASK		26
-#define CD_NUMTYPES		27
+#define CD_GRIDS	       	24
+#define CD_PAINTMASK		25
+#define CD_NUMTYPES		26
 
 /* Bits for CustomDataMask */
 #define CD_MASK_MVERT		(1 << CD_MVERT)
@@ -113,8 +128,7 @@ typedef struct CustomData {
 #define CD_MASK_MDISPS		(1 << CD_MDISPS)
 #define CD_MASK_WEIGHT_MCOL	(1 << CD_WEIGHT_MCOL)
 #define CD_MASK_CLOTH_ORCO	(1 << CD_CLOTH_ORCO)
-#define CD_MASK_FACEGRID	(1 << CD_FACEGRID)
-#define CD_MASK_DISP		(1 << CD_DISP)
+#define CD_MASK_GRIDS		(1 << CD_GRIDS)
 #define CD_MASK_PAINTMASK	(1 << CD_PAINTMASK)
 
 /* CustomData.flag */
@@ -132,6 +146,8 @@ typedef struct CustomData {
 #define CD_FLAG_IN_MEMORY (1<<4)
 /* used to enable/disable a layer */
 #define CD_FLAG_ENABLED   (1<<5)
+/* indicates the layer is linked to subdivided data */
+#define CD_FLAG_MULTIRES  (1<<6)
 
 /* Limits */
 #define MAX_MTFACE 8
diff --git a/source/blender/makesrna/intern/rna_mesh.c b/source/blender/makesrna/intern/rna_mesh.c
index f8e1e27a905..7b7931a2b4a 100644
--- a/source/blender/makesrna/intern/rna_mesh.c
+++ b/source/blender/makesrna/intern/rna_mesh.c
@@ -1575,6 +1575,11 @@ static void rna_def_mcol(BlenderRNA *brna)
 	RNA_def_property_ui_text(prop, "Active Render", "Sets the layer as active for rendering");
 	RNA_def_property_update(prop, 0, "rna_Mesh_update_data");
 
+	prop= RNA_def_property(srna, "multiresolution", PROP_BOOLEAN, PROP_NONE);
+	RNA_def_property_boolean_sdna(prop, NULL, "flag", CD_FLAG_MULTIRES);
+	RNA_def_property_ui_text(prop, "Multiresolution", "Sets whether the layer can be edited at multiple resolutions");
+	RNA_def_property_clear_flag(prop, PROP_EDITABLE);
+
 	prop= RNA_def_property(srna, "data", PROP_COLLECTION, PROP_NONE);
 	RNA_def_property_struct_type(prop, "MeshColor");
 	RNA_def_property_ui_text(prop, "Data", "");
@@ -1845,12 +1850,12 @@ static void rna_def_mesh(BlenderRNA *brna)
 	RNA_def_property_pointer_funcs(prop, "rna_Mesh_active_vertex_color_get", "rna_Mesh_active_vertex_color_set", NULL);
 	RNA_def_property_flag(prop, PROP_EDITABLE);
 	RNA_def_property_ui_text(prop, "Active Vertex Color Layer", "Active vertex color layer");
-	RNA_def_property_update(prop, 0, "rna_Mesh_update_data");
+	//RNA_def_property_update(prop, 0, "rna_Mesh_update_data");
 
 	prop= RNA_def_property(srna, "active_vertex_color_index", PROP_INT, PROP_UNSIGNED);
 	RNA_def_property_int_funcs(prop, "rna_Mesh_active_vertex_color_index_get", "rna_Mesh_active_vertex_color_index_set", "rna_Mesh_active_vertex_color_index_range");
 	RNA_def_property_ui_text(prop, "Active Vertex Color Index", "Active vertex color index");
-	RNA_def_property_update(prop, 0, "rna_Mesh_update_data");
+	//RNA_def_property_update(prop, 0, "rna_Mesh_update_data");
 
 	prop= RNA_def_property(srna, "paint_mask_layers", PROP_COLLECTION, PROP_NONE);
 	RNA_def_property_collection_sdna(prop, NULL, "vdata.layers", "vdata.totlayer");