OpenSubdiv: Add OpenSubdiv files which are related on the CCGSubSurf and GPU

Those files are still not in use (SCons will tyr to compile new CCGSubSurf files
but no code will be in use at all because those new files are fully wrapped by
ifdef WITH_OPENSUBDIV check).

3 files changed, 1485 insertions, 0 deletions

diff --git a/source/blender/blenkernel/intern/CCGSubSurf_opensubdiv.c b/source/blender/blenkernel/intern/CCGSubSurf_opensubdiv.c
new file mode 100644
index 00000000000..d3898d9a543
--- /dev/null
+++ b/source/blender/blenkernel/intern/CCGSubSurf_opensubdiv.c
@@ -0,0 +1,883 @@
+/*
+ * ***** BEGIN GPL LICENSE BLOCK *****
+ *
+ * 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.
+ *
+ * ***** END GPL LICENSE BLOCK *****
+ */
+
+/** \file blender/blenkernel/intern/CCGSubSurf_opensubdiv.c
+ *  \ingroup bke
+ */
+
+#ifdef WITH_OPENSUBDIV
+
+#include "MEM_guardedalloc.h"
+#include "BLI_sys_types.h" // for intptr_t support
+
+#include "BLI_utildefines.h" /* for BLI_assert */
+#include "BLI_math.h"
+
+#include "CCGSubSurf.h"
+#include "CCGSubSurf_intern.h"
+
+#include "BKE_DerivedMesh.h"
+
+#include "DNA_userdef_types.h"
+
+#include "opensubdiv_capi.h"
+#include "opensubdiv_converter_capi.h"
+
+#include "GL/glew.h"
+
+#define OSD_LOG if (false) printf
+
+static bool compare_ccg_derivedmesh_topology(CCGSubSurf *ss, DerivedMesh *dm)
+{
+	const int num_verts = dm->getNumVerts(dm);
+	const int num_edges = dm->getNumEdges(dm);
+	const int num_polys = dm->getNumPolys(dm);
+	const MEdge *medge = dm->getEdgeArray(dm);
+	const MLoop *mloop = dm->getLoopArray(dm);
+	const MPoly *mpoly = dm->getPolyArray(dm);
+
+	/* Quick preliminary tests based on the number of verts and facces. */
+	{
+		if (num_verts != ss->vMap->numEntries ||
+		    num_edges != ss->eMap->numEntries ||
+		    num_polys != ss->fMap->numEntries)
+		{
+			return false;
+		}
+	}
+
+	/* Rather slow check for faces topology change. */
+	{
+		CCGFaceIterator ccg_face_iter;
+		for (ccgSubSurf_initFaceIterator(ss, &ccg_face_iter);
+		     !ccgFaceIterator_isStopped(&ccg_face_iter);
+		     ccgFaceIterator_next(&ccg_face_iter))
+		{
+			/*const*/ CCGFace *ccg_face = ccgFaceIterator_getCurrent(&ccg_face_iter);
+			const int poly_index = GET_INT_FROM_POINTER(ccgSubSurf_getFaceFaceHandle(ccg_face));
+			const MPoly *mp = &mpoly[poly_index];
+			int corner;
+			if (ccg_face->numVerts != mp->totloop) {
+				return false;
+			}
+			for (corner = 0; corner < ccg_face->numVerts; corner++) {
+				/*const*/ CCGVert *ccg_vert = FACE_getVerts(ccg_face)[corner];
+				const int vert_index = GET_INT_FROM_POINTER(ccgSubSurf_getVertVertHandle(ccg_vert));
+				if (vert_index != mloop[mp->loopstart + corner].v) {
+					return false;
+				}
+			}
+		}
+	}
+
+	/* Check for edge topology change. */
+	{
+		CCGEdgeIterator ccg_edge_iter;
+		for (ccgSubSurf_initEdgeIterator(ss, &ccg_edge_iter);
+		     !ccgEdgeIterator_isStopped(&ccg_edge_iter);
+		     ccgEdgeIterator_next(&ccg_edge_iter))
+		{
+			/* const */ CCGEdge *ccg_edge = ccgEdgeIterator_getCurrent(&ccg_edge_iter);
+			/* const */ CCGVert *ccg_vert1 = ccg_edge->v0;
+			/* const */ CCGVert *ccg_vert2 = ccg_edge->v1;
+			const int ccg_vert1_index = GET_INT_FROM_POINTER(ccgSubSurf_getVertVertHandle(ccg_vert1));
+			const int ccg_vert2_index = GET_INT_FROM_POINTER(ccgSubSurf_getVertVertHandle(ccg_vert2));
+			const int edge_index = GET_INT_FROM_POINTER(ccgSubSurf_getEdgeEdgeHandle(ccg_edge));
+			const MEdge *me = &medge[edge_index];
+			if (me->v1 != ccg_vert1_index || me->v2 != ccg_vert2_index) {
+				return false;
+			}
+		}
+	}
+
+	/* TODO(sergey): Crease topology changes detection. */
+	{
+		CCGEdgeIterator ccg_edge_iter;
+		for (ccgSubSurf_initEdgeIterator(ss, &ccg_edge_iter);
+		     !ccgEdgeIterator_isStopped(&ccg_edge_iter);
+		     ccgEdgeIterator_next(&ccg_edge_iter))
+		{
+			/* const */ CCGEdge *ccg_edge = ccgEdgeIterator_getCurrent(&ccg_edge_iter);
+			const int edge_index = GET_INT_FROM_POINTER(ccgSubSurf_getEdgeEdgeHandle(ccg_edge));
+			if (ccg_edge->crease != medge[edge_index].crease) {
+				return false;
+			}
+		}
+	}
+
+	return true;
+}
+
+static bool compare_osd_derivedmesh_topology(CCGSubSurf *ss, DerivedMesh *dm)
+{
+	const OpenSubdiv_TopologyRefinerDescr *topology_refiner;
+	OpenSubdiv_Converter converter;
+	bool result;
+	if (ss->osd_mesh == NULL && ss->osd_topology_refiner == NULL) {
+		return true;
+	}
+	/* TODO(sergey): De-duplicate with topology counter at the bottom of
+	 * the file.
+	 */
+	if (ss->osd_topology_refiner != NULL) {
+		topology_refiner = ss->osd_topology_refiner;
+	}
+	else {
+		topology_refiner = openSubdiv_getGLMeshTopologyRefiner(ss->osd_mesh);
+	}
+	ccgSubSurf_converter_setup_from_derivedmesh(ss, dm, &converter);
+	result = openSubdiv_topologyRefnerCompareConverter(topology_refiner,
+	                                                   &converter);
+	ccgSubSurf_converter_free(&converter);
+	return result;
+}
+
+static bool opensubdiv_is_topology_changed(CCGSubSurf *ss, DerivedMesh *dm)
+{
+	if (ss->osd_compute != U.opensubdiv_compute_type) {
+		return true;
+	}
+	if (ss->osd_topology_refiner != NULL) {
+		int levels = openSubdiv_topologyRefinerGetSubdivLevel(
+		        ss->osd_topology_refiner);
+		BLI_assert(ss->osd_mesh_invalid == true);
+		if (levels != ss->subdivLevels) {
+			return true;
+		}
+	}
+	if (ss->osd_mesh != NULL && ss->osd_mesh_invalid == false) {
+		const OpenSubdiv_TopologyRefinerDescr *topology_refiner =
+		        openSubdiv_getGLMeshTopologyRefiner(ss->osd_mesh);
+		int levels = openSubdiv_topologyRefinerGetSubdivLevel(topology_refiner);
+		BLI_assert(ss->osd_topology_refiner == NULL);
+		if (levels != ss->subdivLevels) {
+			return true;
+		}
+	}
+	if (ss->skip_grids == false) {
+		return compare_ccg_derivedmesh_topology(ss, dm) == false;
+	}
+	else {
+		return compare_osd_derivedmesh_topology(ss, dm) == false;
+	}
+	return false;
+}
+
+void ccgSubSurf_checkTopologyChanged(CCGSubSurf *ss, DerivedMesh *dm)
+{
+	if (opensubdiv_is_topology_changed(ss, dm)) {
+		/* ** Make sure both GPU and CPU backends are properly reset. ** */
+
+		ss->osd_coarse_coords_invalid = true;
+		ss->osd_uvs_invalid = true;
+
+		/* Reset GPU part. */
+		ss->osd_mesh_invalid = true;
+		if (ss->osd_topology_refiner != NULL) {
+			openSubdiv_deleteTopologyRefinerDescr(ss->osd_topology_refiner);
+			ss->osd_topology_refiner = NULL;
+		}
+
+		/* Reste CPU side. */
+		if (ss->osd_evaluator != NULL) {
+			openSubdiv_deleteEvaluatorDescr(ss->osd_evaluator);
+			ss->osd_evaluator = NULL;
+		}
+	}
+}
+
+static void ccgSubSurf__updateGLMeshCoords(CCGSubSurf *ss)
+{
+	BLI_assert(ss->meshIFC.numLayers == 3);
+	openSubdiv_osdGLMeshUpdateVertexBuffer(ss->osd_mesh,
+	                                       (float *) ss->osd_coarse_coords,
+	                                       0,
+	                                       ss->osd_num_coarse_coords);
+}
+
+bool ccgSubSurf_prepareGLMesh(CCGSubSurf *ss, bool use_osd_glsl)
+{
+	int compute_type;
+
+	switch (U.opensubdiv_compute_type) {
+#define CHECK_COMPUTE_TYPE(type) \
+		case USER_OPENSUBDIV_COMPUTE_ ## type: \
+			compute_type = OPENSUBDIV_EVALUATOR_ ## type; \
+			break;
+		CHECK_COMPUTE_TYPE(CPU)
+		CHECK_COMPUTE_TYPE(OPENMP)
+		CHECK_COMPUTE_TYPE(OPENCL)
+		CHECK_COMPUTE_TYPE(CUDA)
+		CHECK_COMPUTE_TYPE(GLSL_TRANSFORM_FEEDBACK)
+		CHECK_COMPUTE_TYPE(GLSL_COMPUTE)
+#undef CHECK_COMPUTE_TYPE
+	}
+
+	if (ss->osd_vao == 0) {
+		glGenVertexArrays(1, &ss->osd_vao);
+	}
+
+	if (ss->osd_mesh_invalid) {
+		if (ss->osd_mesh != NULL) {
+			openSubdiv_deleteOsdGLMesh(ss->osd_mesh);
+			ss->osd_mesh = NULL;
+		}
+		ss->osd_mesh_invalid = false;
+	}
+
+	if (ss->osd_mesh == NULL) {
+		ss->osd_mesh = openSubdiv_createOsdGLMeshFromTopologyRefiner(
+		        ss->osd_topology_refiner,
+		        compute_type,
+		        ss->subdivLevels,
+		        ss->osd_subsurf_uv);
+		ss->osd_topology_refiner = NULL;
+
+		if (UNLIKELY(ss->osd_mesh == NULL)) {
+			/* Most likely compute device is not available. */
+			return false;
+		}
+
+		ccgSubSurf__updateGLMeshCoords(ss);
+
+		openSubdiv_osdGLMeshRefine(ss->osd_mesh);
+		openSubdiv_osdGLMeshSynchronize(ss->osd_mesh);
+
+		glBindVertexArray(ss->osd_vao);
+		glBindBuffer(GL_ARRAY_BUFFER,
+		             openSubdiv_getOsdGLMeshVertexBuffer(ss->osd_mesh));
+
+		glEnableVertexAttribArray(0);
+		glEnableVertexAttribArray(1);
+		glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE,
+		                      sizeof(GLfloat) * 6, 0);
+		glVertexAttribPointer(1, 3, GL_FLOAT, GL_FALSE,
+		                      sizeof(GLfloat) * 6, (float *)12);
+
+		glBindBuffer(GL_ARRAY_BUFFER, 0);
+	}
+	else if (ss->osd_coarse_coords_invalid) {
+		ccgSubSurf__updateGLMeshCoords(ss);
+		openSubdiv_osdGLMeshRefine(ss->osd_mesh);
+		openSubdiv_osdGLMeshSynchronize(ss->osd_mesh);
+		ss->osd_coarse_coords_invalid = false;
+	}
+
+	openSubdiv_osdGLMeshDisplayPrepare(use_osd_glsl, ss->osd_uv_index);
+
+	return true;
+}
+
+void ccgSubSurf_drawGLMesh(CCGSubSurf *ss, bool fill_quads,
+                           int start_partition, int num_partitions)
+{
+	if (LIKELY(ss->osd_mesh != NULL)) {
+		glBindBuffer(GL_ELEMENT_ARRAY_BUFFER,
+		             openSubdiv_getOsdGLMeshPatchIndexBuffer(ss->osd_mesh));
+
+		openSubdiv_osdGLMeshBindVertexBuffer(ss->osd_mesh);
+		glBindVertexArray(ss->osd_vao);
+		openSubdiv_osdGLMeshDisplay(ss->osd_mesh, fill_quads,
+		                            start_partition, num_partitions);
+		glBindVertexArray(0);
+		glBindBuffer(GL_ARRAY_BUFFER, 0);
+		glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, 0);
+	}
+}
+
+void ccgSubSurf_setSkipGrids(CCGSubSurf *ss, bool skip_grids)
+{
+	ss->skip_grids = skip_grids;
+}
+
+bool ccgSubSurf_needGrids(CCGSubSurf *ss)
+{
+	return ss->skip_grids == false;
+}
+
+BLI_INLINE void ccgSubSurf__mapGridToFace(int S, float grid_u, float grid_v,
+                                          float *face_u, float *face_v)
+{
+	float u, v;
+
+	/* - Each grid covers half of the face along the edges.
+	 * - Grid's (0, 0) starts from the middle of the face.
+	 */
+	u = 0.5f - 0.5f * grid_u;
+	v = 0.5f - 0.5f * grid_v;
+
+	if (S == 0) {
+		*face_u = v;
+		*face_v = u;
+	}
+	else if (S == 1) {
+		*face_u = 1.0f - u;
+		*face_v = v;
+	}
+	else if (S == 2) {
+		*face_u = 1.0f - v;
+		*face_v = 1.0f - u;
+	}
+	else {
+		*face_u = u;
+		*face_v = 1.0f - v;
+	}
+}
+
+BLI_INLINE void ccgSubSurf__mapEdgeToFace(int S,
+                                          int edge_segment,
+                                          bool inverse_edge,
+                                          int edgeSize,
+                                          float *face_u, float *face_v)
+{
+	int t = inverse_edge ? edgeSize - edge_segment - 1 : edge_segment;
+	if (S == 0) {
+		*face_u = (float) t / (edgeSize - 1);
+		*face_v = 0.0f;
+	}
+	else if (S == 1) {
+		*face_u = 1.0f;
+		*face_v = (float) t / (edgeSize - 1);
+	}
+	else if (S == 2) {
+		*face_u = 1.0f - (float) t / (edgeSize - 1);
+		*face_v = 1.0f;
+	}
+	else {
+		*face_u = 0.0f;
+		*face_v = 1.0f - (float) t / (edgeSize - 1);
+	}
+}
+
+void ccgSubSurf_evaluatorSetFVarUV(CCGSubSurf *ss,
+                                   DerivedMesh *dm,
+                                   int layer_index)
+{
+	MPoly *mpoly = dm->getPolyArray(dm);
+	MLoopUV *mloopuv = CustomData_get_layer_n(&dm->loopData, CD_MLOOPUV, layer_index);
+	int num_polys = dm->getNumPolys(dm);
+	int index, poly;
+	BLI_assert(ss->osd_evaluator != NULL);
+	for (poly = 0, index = 0; poly < num_polys; poly++) {
+		int loop;
+		MPoly *mp = &mpoly[poly];
+		for (loop = 0; loop < mp->totloop; loop++, index++) {
+			MLoopUV *mluv = &mloopuv[loop + mp->loopstart];
+			(void)mluv;
+			/* TODO(sergey): Send mluv->uv to the evaluator's face varying
+			 * buffer.
+			 */
+		}
+	}
+	(void)ss;
+}
+
+void ccgSubSurf_evaluatorFVarUV(CCGSubSurf *ss,
+                                int face_index, int S,
+                                float grid_u, float grid_v,
+                                float uv[2])
+{
+	float face_u, face_v;
+	ccgSubSurf__mapGridToFace(S,
+	                          grid_u, grid_v,
+	                          &face_u, &face_v);
+	(void)ss;
+	(void)face_index;
+	/* TODO(sergey): Evaluate face varying coordinate. */
+	zero_v2(uv);
+}
+
+static bool opensubdiv_createEvaluator(CCGSubSurf *ss)
+{
+	OpenSubdiv_Converter converter;
+	OpenSubdiv_TopologyRefinerDescr *topology_refiner;
+	ccgSubSurf_converter_setup_from_ccg(ss, &converter);
+	topology_refiner = openSubdiv_createTopologyRefinerDescr(&converter);
+	ccgSubSurf_converter_free(&converter);
+	ss->osd_evaluator =
+	        openSubdiv_createEvaluatorDescr(topology_refiner,
+	                                        ss->subdivLevels);
+	return ss->osd_evaluator != NULL;
+}
+
+static bool opensubdiv_ensureEvaluator(CCGSubSurf *ss)
+{
+	if (ss->osd_evaluator == NULL) {
+		OSD_LOG("Allocating new evaluator, %d verts\n", ss->vMap->numEntries);
+		opensubdiv_createEvaluator(ss);
+	}
+	return ss->osd_evaluator != NULL;
+}
+
+static void opensubdiv_updateEvaluatorCoarsePositions(CCGSubSurf *ss)
+{
+	float (*positions)[3];
+	int vertDataSize = ss->meshIFC.vertDataSize;
+	int num_basis_verts = ss->vMap->numEntries;
+	int i;
+
+	/* TODO(sergey): Avoid allocation on every update. We could either update
+	 * coordinates in chunks of 1K vertices (which will only use stack memory)
+	 * or do some callback magic for OSD evaluator can invoke it and fill in
+	 * buffer directly.
+	 */
+	if (ss->meshIFC.numLayers == 3) {
+		/* If all the components are to be initialized, no need to memset the
+		 * new memory block.
+		 */
+		positions = MEM_mallocN(3 * sizeof(float) * num_basis_verts,
+		                        "OpenSubdiv coarse points");
+	}
+	else {
+		/* Calloc in order to have z component initialized to 0 for Uvs */
+		positions = MEM_callocN(3 * sizeof(float) * num_basis_verts,
+		                        "OpenSubdiv coarse points");
+	}
+#pragma omp parallel for
+	for (i = 0; i < ss->vMap->curSize; i++) {
+		CCGVert *v = (CCGVert *) ss->vMap->buckets[i];
+		for (; v; v = v->next) {
+			float *co = VERT_getCo(v, 0);
+			BLI_assert(v->osd_index < ss->vMap->numEntries);
+			VertDataCopy(positions[v->osd_index], co, ss);
+			OSD_LOG("Point %d has value %f %f %f\n",
+			        v->osd_index,
+			        positions[v->osd_index][0],
+			        positions[v->osd_index][1],
+			        positions[v->osd_index][2]);
+		}
+	}
+
+	openSubdiv_setEvaluatorCoarsePositions(ss->osd_evaluator,
+	                                       (float *)positions,
+	                                       0,
+	                                       num_basis_verts);
+
+	MEM_freeN(positions);
+}
+
+static void opensubdiv_evaluateQuadFaceGrids(CCGSubSurf *ss,
+                                             CCGFace *face,
+                                             const int osd_face_index)
+{
+	int normalDataOffset = ss->normalDataOffset;
+	int subdivLevels = ss->subdivLevels;
+	int gridSize = ccg_gridsize(subdivLevels);
+	int edgeSize = ccg_edgesize(subdivLevels);
+	int vertDataSize = ss->meshIFC.vertDataSize;
+	int S;
+	bool do_normals = ss->meshIFC.numLayers == 3;
+
+#pragma omp parallel for
+	for (S = 0; S < face->numVerts; S++) {
+		int x, y, k;
+		CCGEdge *edge = NULL;
+		bool inverse_edge;
+
+		for (x = 0; x < gridSize; x++) {
+			for (y = 0; y < gridSize; y++) {
+				float *co = FACE_getIFCo(face, subdivLevels, S, x, y);
+				float *no = FACE_getIFNo(face, subdivLevels, S, x, y);
+				float grid_u = (float) x / (gridSize - 1),
+				      grid_v = (float) y / (gridSize - 1);
+				float face_u, face_v;
+				float P[3], dPdu[3], dPdv[3];
+
+				ccgSubSurf__mapGridToFace(S, grid_u, grid_v, &face_u, &face_v);
+
+				/* TODO(sergey): Need proper port. */
+				openSubdiv_evaluateLimit(ss->osd_evaluator, osd_face_index,
+				                         face_u, face_v,
+				                         P,
+				                         do_normals ? dPdu : NULL,
+				                         do_normals ? dPdv : NULL);
+
+				OSD_LOG("face=%d, corner=%d, grid_u=%f, grid_v=%f, face_u=%f, face_v=%f, P=(%f, %f, %f)\n",
+				        osd_face_index, S, grid_u, grid_v, face_u, face_v, P[0], P[1], P[2]);
+
+				VertDataCopy(co, P, ss);
+				if (do_normals) {
+					cross_v3_v3v3(no, dPdu, dPdv);
+					normalize_v3(no);
+				}
+
+				if (x == gridSize - 1 && y == gridSize - 1) {
+					float *vert_co = VERT_getCo(FACE_getVerts(face)[S], subdivLevels);
+					VertDataCopy(vert_co, co, ss);
+					if (do_normals) {
+						float *vert_no = VERT_getNo(FACE_getVerts(face)[S], subdivLevels);
+						VertDataCopy(vert_no, no, ss);
+					}
+				}
+				if (S == 0 && x == 0 && y == 0) {
+					float *center_co = (float *)FACE_getCenterData(face);
+					VertDataCopy(center_co, co, ss);
+					if (do_normals) {
+						float *center_no = (float *)((byte *)FACE_getCenterData(face) + normalDataOffset);
+						VertDataCopy(center_no, no, ss);
+					}
+				}
+			}
+		}
+
+		for (x = 0; x < gridSize; x++) {
+			VertDataCopy(FACE_getIECo(face, subdivLevels, S, x),
+			             FACE_getIFCo(face, subdivLevels, S, x, 0), ss);
+			if (do_normals) {
+				VertDataCopy(FACE_getIENo(face, subdivLevels, S, x),
+				             FACE_getIFNo(face, subdivLevels, S, x, 0), ss);
+			}
+		}
+
+		for (k = 0; k < face->numVerts; k++) {
+			CCGEdge *current_edge = FACE_getEdges(face)[k];
+			CCGVert **face_verts = FACE_getVerts(face);
+			if (current_edge->v0 == face_verts[S] &&
+			    current_edge->v1 == face_verts[(S + 1) % face->numVerts])
+			{
+				edge = current_edge;
+				inverse_edge = false;
+				break;
+			}
+			if (current_edge->v1 == face_verts[S] &&
+			    current_edge->v0 == face_verts[(S + 1) % face->numVerts])
+			{
+				edge = current_edge;
+				inverse_edge = true;
+				break;
+			}
+		}
+
+		BLI_assert(edge != NULL);
+
+		for (x = 0; x < edgeSize; x++) {
+			float u = 0, v = 0;
+			float *co = EDGE_getCo(edge, subdivLevels, x);
+			float *no = EDGE_getNo(edge, subdivLevels, x);
+			float P[3], dPdu[3], dPdv[3];
+			ccgSubSurf__mapEdgeToFace(S, x,
+			                          inverse_edge,
+			                          edgeSize,
+			                          &u, &v);
+
+			/* TODO(sergey): Ideally we will re-use grid here, but for now
+			 * let's just re-evaluate for simplicity.
+			 */
+			/* TODO(sergey): Need proper port. */
+			openSubdiv_evaluateLimit(ss->osd_evaluator, osd_face_index, u, v, P, dPdu, dPdv);
+			VertDataCopy(co, P, ss);
+			if (do_normals) {
+				cross_v3_v3v3(no, dPdu, dPdv);
+				normalize_v3(no);
+			}
+		}
+	}
+}
+
+static void opensubdiv_evaluateNGonFaceGrids(CCGSubSurf *ss,
+                                             CCGFace *face,
+                                             const int osd_face_index)
+{
+	CCGVert **all_verts = FACE_getVerts(face);
+	int normalDataOffset = ss->normalDataOffset;
+	int subdivLevels = ss->subdivLevels;
+	int gridSize = ccg_gridsize(subdivLevels);
+	int edgeSize = ccg_edgesize(subdivLevels);
+	int vertDataSize = ss->meshIFC.vertDataSize;
+	int S;
+	bool do_normals = ss->meshIFC.numLayers == 3;
+
+	/* Note about handling non-quad faces.
+	 *
+	 * In order to deal with non-quad faces we need to split them
+	 * into a quads in the following way:
+	 *
+	 *                                                     |
+	 *                                                (vert_next)
+	 *                                                     |
+	 *                                                     |
+	 *                                                     |
+	 *                  (face_center) ------------------- (v2)
+	 *                         | (o)-------------------->  |
+	 *                         |  |                     v  |
+	 *                         |  |                        |
+	 *                         |  |                        |
+	 *                         |  |                        |
+	 *                         |  |                   y ^  |
+	 *                         |  |                     |  |
+	 *                         |  v  u             x    |  |
+	 *                         |                   <---(o) |
+	 * ---- (vert_prev) ---- (v1)  --------------------  (vert)
+	 *
+	 * This is how grids are expected to be stored and it's how
+	 * OpenSubdiv deals with non-quad faces using ptex face indices.
+	 * We only need to convert ptex (x, y) to grid (u, v) by some
+	 * simple flips and evaluate the ptex face.
+	 */
+
+	/* Evaluate face grids. */
+#pragma omp parallel for
+	for (S = 0; S < face->numVerts; S++) {
+		int x, y;
+		for (x = 0; x < gridSize; x++) {
+			for (y = 0; y < gridSize; y++) {
+				float *co = FACE_getIFCo(face, subdivLevels, S, x, y);
+				float *no = FACE_getIFNo(face, subdivLevels, S, x, y);
+				float u = 1.0f - (float) y / (gridSize - 1),
+				      v = 1.0f - (float) x / (gridSize - 1);
+				float P[3], dPdu[3], dPdv[3];
+
+				/* TODO(sergey): Need proper port. */
+				openSubdiv_evaluateLimit(ss->osd_evaluator, osd_face_index + S, u, v, P, dPdu, dPdv);
+
+				OSD_LOG("face=%d, corner=%d, u=%f, v=%f, P=(%f, %f, %f)\n",
+				        osd_face_index + S, S, u, v, P[0], P[1], P[2]);
+
+				VertDataCopy(co, P, ss);
+				if (do_normals) {
+					cross_v3_v3v3(no, dPdu, dPdv);
+					normalize_v3(no);
+				}
+
+				/* TODO(sergey): De-dpuplicate with the quad case. */
+				if (x == gridSize - 1 && y == gridSize - 1) {
+					float *vert_co = VERT_getCo(FACE_getVerts(face)[S], subdivLevels);
+					VertDataCopy(vert_co, co, ss);
+					if (do_normals) {
+						float *vert_no = VERT_getNo(FACE_getVerts(face)[S], subdivLevels);
+						VertDataCopy(vert_no, no, ss);
+					}
+				}
+				if (S == 0 && x == 0 && y == 0) {
+					float *center_co = (float *)FACE_getCenterData(face);
+					VertDataCopy(center_co, co, ss);
+					if (do_normals) {
+						float *center_no = (float *)((byte *)FACE_getCenterData(face) + normalDataOffset);
+						VertDataCopy(center_no, no, ss);
+					}
+				}
+			}
+		}
+		for (x = 0; x < gridSize; x++) {
+			VertDataCopy(FACE_getIECo(face, subdivLevels, S, x),
+			             FACE_getIFCo(face, subdivLevels, S, x, 0), ss);
+			if (do_normals) {
+				VertDataCopy(FACE_getIENo(face, subdivLevels, S, x),
+				             FACE_getIFNo(face, subdivLevels, S, x, 0), ss);
+			}
+		}
+	}
+
+	/* Evaluate edges. */
+	for (S = 0; S < face->numVerts; S++) {
+		CCGEdge *edge = FACE_getEdges(face)[S];
+		int x, S0, S1;
+		bool flip;
+
+		for (x = 0; x < face->numVerts; ++x) {
+			if (all_verts[x] == edge->v0) {
+				S0 = x;
+			}
+			else if (all_verts[x] == edge->v1) {
+				S1 = x;
+			}
+		}
+		if (S == face->numVerts - 1) {
+			flip = S0 > S1;
+		}
+		else {
+			flip = S0 < S1;
+		}
+
+		for (x = 0; x <= edgeSize / 2; x++) {
+			float *edge_co = EDGE_getCo(edge, subdivLevels, x);
+			float *edge_no = EDGE_getNo(edge, subdivLevels, x);
+			float *face_edge_co;
+			float *face_edge_no;
+			if (flip) {
+				face_edge_co = FACE_getIFCo(face, subdivLevels, S0, gridSize - 1, gridSize - 1 - x);
+				face_edge_no = FACE_getIFNo(face, subdivLevels, S0, gridSize - 1, gridSize - 1 - x);
+			}
+			else {
+				face_edge_co = FACE_getIFCo(face, subdivLevels, S0, gridSize - 1 - x, gridSize - 1);
+				face_edge_no = FACE_getIFNo(face, subdivLevels, S0, gridSize - 1 - x, gridSize - 1);
+			}
+			VertDataCopy(edge_co, face_edge_co, ss);
+			if (do_normals) {
+				VertDataCopy(edge_no, face_edge_no, ss);
+			}
+		}
+		for (x = edgeSize / 2 + 1; x < edgeSize; x++) {
+			float *edge_co = EDGE_getCo(edge, subdivLevels, x);
+			float *edge_no = EDGE_getNo(edge, subdivLevels, x);
+			float *face_edge_co;
+			float *face_edge_no;
+			if (flip) {
+				face_edge_co = FACE_getIFCo(face, subdivLevels, S1, x - edgeSize / 2, gridSize - 1);
+				face_edge_no = FACE_getIFNo(face, subdivLevels, S1, x - edgeSize / 2, gridSize - 1);
+			}
+			else {
+				face_edge_co = FACE_getIFCo(face, subdivLevels, S1, gridSize - 1, x - edgeSize / 2);
+				face_edge_no = FACE_getIFNo(face, subdivLevels, S1, gridSize - 1, x - edgeSize / 2);
+			}
+			VertDataCopy(edge_co, face_edge_co, ss);
+			if (do_normals) {
+				VertDataCopy(edge_no, face_edge_no, ss);
+			}
+		}
+	}
+}
+
+static void opensubdiv_evaluateGrids(CCGSubSurf *ss)
+{
+	int i;
+	for (i = 0; i < ss->fMap->curSize; i++) {
+		CCGFace *face = (CCGFace *) ss->fMap->buckets[i];
+		for (; face; face = face->next) {
+			if (face->numVerts == 4) {
+				/* For quads we do special magic with converting face coords
+				 * into corner coords and interpolating grids from it.
+				 */
+				opensubdiv_evaluateQuadFaceGrids(ss, face, face->osd_index);
+			}
+			else {
+				/* NGons and tris are split into separate osd faces which
+				 * evaluates onto grids directly.
+				 */
+				opensubdiv_evaluateNGonFaceGrids(ss, face, face->osd_index);
+			}
+		}
+	}
+}
+
+CCGError ccgSubSurf_initOpenSubdivSync(CCGSubSurf *ss)
+{
+	if (ss->syncState != eSyncState_None) {
+		return eCCGError_InvalidSyncState;
+	}
+	ss->syncState = eSyncState_OpenSubdiv;
+	return eCCGError_None;
+}
+
+void ccgSubSurf_prepareTopologyRefiner(CCGSubSurf *ss, DerivedMesh *dm)
+{
+	if (ss->osd_mesh == NULL || ss->osd_mesh_invalid) {
+		OpenSubdiv_Converter converter;
+		ccgSubSurf_converter_setup_from_derivedmesh(ss, dm, &converter);
+		/* TODO(sergey): Remove possibly previously allocated refiner. */
+		ss->osd_topology_refiner = openSubdiv_createTopologyRefinerDescr(&converter);
+		ccgSubSurf_converter_free(&converter);
+	}
+
+	/* Update number of grids, needed for things like final faces
+	 * counter, used by display drawing.
+	 */
+	{
+		const int num_polys = dm->getNumPolys(dm);
+		const MPoly *mpoly = dm->getPolyArray(dm);
+		int poly;
+		ss->numGrids = 0;
+		for (poly = 0; poly < num_polys; ++poly) {
+			ss->numGrids += mpoly[poly].totloop;
+		}
+	}
+
+	{
+		const int num_verts = dm->getNumVerts(dm);
+		const MVert *mvert = dm->getVertArray(dm);
+		int vert;
+		if (ss->osd_coarse_coords != NULL &&
+		    num_verts != ss->osd_num_coarse_coords)
+		{
+			MEM_freeN(ss->osd_coarse_coords);
+			ss->osd_coarse_coords = NULL;
+		}
+		if (ss->osd_coarse_coords == NULL) {
+			ss->osd_coarse_coords = MEM_mallocN(sizeof(float) * 6 * num_verts, "osd coarse positions");
+		}
+		for (vert = 0; vert < num_verts; vert++) {
+			copy_v3_v3(ss->osd_coarse_coords[vert * 2 + 0], mvert[vert].co);
+			normal_short_to_float_v3(ss->osd_coarse_coords[vert * 2 + 1], mvert[vert].no);
+		}
+		ss->osd_num_coarse_coords = num_verts;
+		ss->osd_coarse_coords_invalid = true;
+	}
+}
+
+void ccgSubSurf__sync_opensubdiv(CCGSubSurf *ss)
+{
+	BLI_assert(ss->meshIFC.numLayers == 2 || ss->meshIFC.numLayers == 3);
+
+	/* Common synchronization steps */
+	ss->osd_compute = U.opensubdiv_compute_type;
+
+	if (ss->skip_grids == false) {
+		/* Make sure OSD evaluator is up-to-date. */
+		if (opensubdiv_ensureEvaluator(ss)) {
+			/* Update coarse points in the OpenSubdiv evaluator. */
+			opensubdiv_updateEvaluatorCoarsePositions(ss);
+
+			/* Evaluate opensubdiv mesh into the CCG grids. */
+			opensubdiv_evaluateGrids(ss);
+		}
+		else {
+			BLI_assert(!"OpenSubdiv initializetion failed, should not happen.");
+		}
+	}
+	else {
+		BLI_assert(ss->meshIFC.numLayers == 3);
+	}
+
+#ifdef DUMP_RESULT_GRIDS
+	ccgSubSurf__dumpCoords(ss);
+#endif
+}
+
+static const OpenSubdiv_TopologyRefinerDescr *get_effective_refiner(
+        const CCGSubSurf *ss)
+{
+	if (ss->osd_topology_refiner) {
+		return ss->osd_topology_refiner;
+	}
+	return openSubdiv_getGLMeshTopologyRefiner(ss->osd_mesh);
+}
+
+int ccgSubSurf__getNumOsdBaseVerts(const CCGSubSurf *ss)
+{
+	const OpenSubdiv_TopologyRefinerDescr *topology_refiner =
+	        get_effective_refiner(ss);
+	return openSubdiv_topologyRefinerGetNumVerts(topology_refiner);
+}
+
+int ccgSubSurf__getNumOsdBaseEdges(const CCGSubSurf *ss)
+{
+	const OpenSubdiv_TopologyRefinerDescr *topology_refiner =
+	        get_effective_refiner(ss);
+	return openSubdiv_topologyRefinerGetNumEdges(topology_refiner);
+}
+
+int ccgSubSurf__getNumOsdBaseFaces(const CCGSubSurf *ss)
+{
+	const OpenSubdiv_TopologyRefinerDescr *topology_refiner =
+	        get_effective_refiner(ss);
+	return openSubdiv_topologyRefinerGetNumFaces(topology_refiner);
+}
+
+#endif  /* WITH_OPENSUBDIV */
diff --git a/source/blender/blenkernel/intern/CCGSubSurf_opensubdiv_converter.c b/source/blender/blenkernel/intern/CCGSubSurf_opensubdiv_converter.c
new file mode 100644
index 00000000000..c66425b3ba8
--- /dev/null
+++ b/source/blender/blenkernel/intern/CCGSubSurf_opensubdiv_converter.c
@@ -0,0 +1,502 @@
+/*
+ * ***** BEGIN GPL LICENSE BLOCK *****
+ *
+ * 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.
+ *
+ * ***** END GPL LICENSE BLOCK *****
+ */
+
+/** \file blender/blenkernel/intern/CCGSubSurf_opensubdiv_converter.c
+ *  \ingroup bke
+ */
+
+#ifdef WITH_OPENSUBDIV
+
+#include <stdlib.h>
+
+#include "MEM_guardedalloc.h"
+#include "BLI_sys_types.h" // for intptr_t support
+
+#include "BLI_utildefines.h" /* for BLI_assert */
+#include "BLI_math.h"
+
+#include "CCGSubSurf.h"
+#include "CCGSubSurf_intern.h"
+
+#include "BKE_DerivedMesh.h"
+
+#include "opensubdiv_capi.h"
+#include "opensubdiv_converter_capi.h"
+
+/**
+ * Converter from DerivedMesh.
+ */
+
+typedef struct ConvDMStorage {
+	CCGSubSurf *ss;
+	DerivedMesh *dm;
+} ConvDMStorage;
+
+/* TODO(sergey): Optimize this by using mesh_map, so we don't
+ * do full mesh lookup for every geometry primitive.
+ */
+
+static OpenSubdiv_SchemeType conv_dm_get_type(
+        const OpenSubdiv_Converter *converter)
+{
+	ConvDMStorage *storage = converter->user_data;
+	if (storage->ss->meshIFC.simpleSubdiv)
+		return OSD_SCHEME_BILINEAR;
+	else
+		return OSD_SCHEME_CATMARK;
+}
+
+static int conv_dm_get_num_faces(const OpenSubdiv_Converter *converter)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	return dm->getNumPolys(dm);
+}
+
+static int conv_dm_get_num_edges(const OpenSubdiv_Converter *converter)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	return dm->getNumEdges(dm);
+}
+
+static int conv_dm_get_num_verts(const OpenSubdiv_Converter *converter)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	return dm->getNumVerts(dm);
+}
+
+static int conv_dm_get_num_face_verts(const OpenSubdiv_Converter *converter,
+                                      int face)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	const MPoly *mp = dm->getPolyArray(dm);
+	const MPoly *mpoly = &mp[face];
+	return mpoly->totloop;
+}
+
+static void conv_dm_get_face_verts(const OpenSubdiv_Converter *converter,
+                                   int face,
+                                   int *face_verts)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	const MLoop *ml = dm->getLoopArray(dm);
+	const MPoly *mp = dm->getPolyArray(dm);
+	const MPoly *mpoly = &mp[face];
+	int loop;
+	for (loop = 0; loop < mpoly->totloop; loop++) {
+		face_verts[loop] = ml[mpoly->loopstart + loop].v;
+	}
+}
+
+static void conv_dm_get_face_edges(const OpenSubdiv_Converter *converter,
+                                   int face,
+                                   int *face_edges)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	const MLoop *ml = dm->getLoopArray(dm);
+	const MPoly *mp = dm->getPolyArray(dm);
+	const MPoly *mpoly = &mp[face];
+	int loop;
+	for (loop = 0; loop < mpoly->totloop; loop++) {
+		face_edges[loop] = ml[mpoly->loopstart + loop].e;
+	}
+}
+
+static void conv_dm_get_edge_verts(const OpenSubdiv_Converter *converter,
+                                   int edge,
+                                   int *edge_verts)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	const MEdge *me = dm->getEdgeArray(dm);
+	const MEdge *medge = &me[edge];
+	edge_verts[0] = medge->v1;
+	edge_verts[1] = medge->v2;
+}
+
+static int conv_dm_get_num_edge_faces(const OpenSubdiv_Converter *converter,
+                                      int edge)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	const MLoop *ml = dm->getLoopArray(dm);
+	const MPoly *mp = dm->getPolyArray(dm);
+	int num = 0, poly;
+	for (poly = 0; poly < dm->getNumPolys(dm); poly++) {
+		const MPoly *mpoly = &mp[poly];
+		int loop;
+		for (loop = 0; loop < mpoly->totloop; loop++) {
+			const MLoop *mloop = &ml[mpoly->loopstart + loop];
+			if (mloop->e == edge) {
+				++num;
+				break;
+			}
+		}
+	}
+	return num;
+}
+
+static void conv_dm_get_edge_faces(const OpenSubdiv_Converter *converter,
+                                   int edge,
+                                   int *edge_faces)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	const MLoop *ml = dm->getLoopArray(dm);
+	const MPoly *mp = dm->getPolyArray(dm);
+	int num = 0, poly;
+	for (poly = 0; poly < dm->getNumPolys(dm); poly++) {
+		const MPoly *mpoly = &mp[poly];
+		int loop;
+		for (loop = 0; loop < mpoly->totloop; loop++) {
+			const MLoop *mloop = &ml[mpoly->loopstart + loop];
+			if (mloop->e == edge) {
+				edge_faces[num++] = poly;
+				break;
+			}
+		}
+	}
+}
+
+static float conv_dm_get_edge_sharpness(const OpenSubdiv_Converter *converter,
+                                        int edge)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	CCGSubSurf *ss = storage->ss;
+	const MEdge *medge = dm->getEdgeArray(dm);
+	return (float)medge[edge].crease / 255.0f * ss->subdivLevels;
+}
+
+static int conv_dm_get_num_vert_edges(const OpenSubdiv_Converter *converter,
+                                      int vert)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	const MEdge *me = dm->getEdgeArray(dm);
+	int num = 0, edge;
+	for (edge = 0; edge < dm->getNumEdges(dm); edge++) {
+		const MEdge *medge = &me[edge];
+		if (medge->v1 == vert || medge->v2 == vert) {
+			++num;
+		}
+	}
+	return num;
+}
+
+static void conv_dm_get_vert_edges(const OpenSubdiv_Converter *converter,
+                                   int vert,
+                                   int *vert_edges)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	const MEdge *me = dm->getEdgeArray(dm);
+	int num = 0, edge;
+	for (edge = 0; edge < dm->getNumEdges(dm); edge++) {
+		const MEdge *medge = &me[edge];
+		if (medge->v1 == vert || medge->v2 == vert) {
+			vert_edges[num++] = edge;
+		}
+	}
+}
+
+static int conv_dm_get_num_vert_faces(const OpenSubdiv_Converter *converter,
+                                      int vert)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	const MLoop *ml = dm->getLoopArray(dm);
+	const MPoly *mp = dm->getPolyArray(dm);
+	int num = 0, poly;
+	for (poly = 0; poly < dm->getNumPolys(dm); poly++) {
+		const MPoly *mpoly = &mp[poly];
+		int loop;
+		for (loop = 0; loop < mpoly->totloop; loop++) {
+			const MLoop *mloop = &ml[mpoly->loopstart + loop];
+			if (mloop->v == vert) {
+				++num;
+				break;
+			}
+		}
+	}
+	return num;
+}
+
+static void conv_dm_get_vert_faces(const OpenSubdiv_Converter *converter,
+                                   int vert,
+                                   int *vert_faces)
+{
+	ConvDMStorage *storage = converter->user_data;
+	DerivedMesh *dm = storage->dm;
+	const MLoop *ml = dm->getLoopArray(dm);
+	const MPoly *mp = dm->getPolyArray(dm);
+	int num = 0, poly;
+	for (poly = 0; poly < dm->getNumPolys(dm); poly++) {
+		const MPoly *mpoly = &mp[poly];
+		int loop;
+		for (loop = 0; loop < mpoly->totloop; loop++) {
+			const MLoop *mloop = &ml[mpoly->loopstart + loop];
+			if (mloop->v == vert) {
+				vert_faces[num++] = poly;
+				break;
+			}
+		}
+	}
+}
+
+static void conv_dm_free_user_data(const OpenSubdiv_Converter *converter)
+{
+	MEM_freeN(converter->user_data);
+}
+
+void ccgSubSurf_converter_setup_from_derivedmesh(
+        CCGSubSurf *ss,
+        DerivedMesh *dm,
+        OpenSubdiv_Converter *converter)
+{
+	ConvDMStorage *user_data;
+
+	converter->get_type = conv_dm_get_type;
+
+	converter->get_num_faces = conv_dm_get_num_faces;
+	converter->get_num_edges = conv_dm_get_num_edges;
+	converter->get_num_verts = conv_dm_get_num_verts;
+
+	converter->get_num_face_verts = conv_dm_get_num_face_verts;
+	converter->get_face_verts = conv_dm_get_face_verts;
+	converter->get_face_edges = conv_dm_get_face_edges;
+
+	converter->get_edge_verts = conv_dm_get_edge_verts;
+	converter->get_num_edge_faces = conv_dm_get_num_edge_faces;
+	converter->get_edge_faces = conv_dm_get_edge_faces;
+	converter->get_edge_sharpness = conv_dm_get_edge_sharpness;
+
+	converter->get_num_vert_edges = conv_dm_get_num_vert_edges;
+	converter->get_vert_edges = conv_dm_get_vert_edges;
+	converter->get_num_vert_faces = conv_dm_get_num_vert_faces;
+	converter->get_vert_faces = conv_dm_get_vert_faces;
+
+	user_data = MEM_mallocN(sizeof(ConvDMStorage), __func__);
+	user_data->ss = ss;
+	user_data->dm = dm;
+	converter->free_user_data = conv_dm_free_user_data;
+	converter->user_data = user_data;
+}
+
+/**
+ * Converter from CCGSubSurf
+ */
+
+static OpenSubdiv_SchemeType conv_ccg_get_bilinear_type(
+        const OpenSubdiv_Converter *converter)
+{
+	CCGSubSurf *ss = converter->user_data;
+	if (ss->meshIFC.simpleSubdiv) {
+		return OSD_SCHEME_BILINEAR;
+	}
+	else {
+		return OSD_SCHEME_CATMARK;
+	}
+}
+
+static int conv_ccg_get_num_faces(const OpenSubdiv_Converter *converter)
+{
+	CCGSubSurf *ss = converter->user_data;
+	return ss->fMap->numEntries;
+}
+
+static int conv_ccg_get_num_edges(const OpenSubdiv_Converter *converter)
+{
+	CCGSubSurf *ss = converter->user_data;
+	return ss->eMap->numEntries;
+}
+
+static int conv_ccg_get_num_verts(const OpenSubdiv_Converter *converter)
+{
+	CCGSubSurf *ss = converter->user_data;
+	return ss->vMap->numEntries;
+}
+
+static int conv_ccg_get_num_face_verts(const OpenSubdiv_Converter *converter,
+                                       int face)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGFace *ccg_face = ccgSubSurf_getFace(ss, SET_INT_IN_POINTER(face));
+	return ccgSubSurf_getFaceNumVerts(ccg_face);
+}
+
+static void conv_ccg_get_face_verts(const OpenSubdiv_Converter *converter,
+                                    int face,
+                                    int *face_verts)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGFace *ccg_face = ccgSubSurf_getFace(ss, SET_INT_IN_POINTER(face));
+	int num_face_verts = ccgSubSurf_getFaceNumVerts(ccg_face);
+	int loop;
+	for (loop = 0; loop < num_face_verts; loop++) {
+		CCGVert *ccg_vert = ccgSubSurf_getFaceVert(ccg_face, loop);
+		face_verts[loop] = GET_INT_FROM_POINTER(ccgSubSurf_getVertVertHandle(ccg_vert));
+	}
+}
+
+static void conv_ccg_get_face_edges(const OpenSubdiv_Converter *converter,
+                                    int face,
+                                    int *face_edges)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGFace *ccg_face = ccgSubSurf_getFace(ss, SET_INT_IN_POINTER(face));
+	int num_face_verts = ccgSubSurf_getFaceNumVerts(ccg_face);
+	int loop;
+	for (loop = 0; loop < num_face_verts; loop++) {
+		CCGEdge *ccg_edge = ccgSubSurf_getFaceEdge(ccg_face, loop);
+		face_edges[loop] = GET_INT_FROM_POINTER(ccgSubSurf_getEdgeEdgeHandle(ccg_edge));
+	}
+}
+
+static void conv_ccg_get_edge_verts(const OpenSubdiv_Converter *converter,
+                                    int edge,
+                                    int *edge_verts)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGEdge *ccg_edge = ccgSubSurf_getEdge(ss, SET_INT_IN_POINTER(edge));
+	CCGVert *ccg_vert0 = ccgSubSurf_getEdgeVert0(ccg_edge);
+	CCGVert *ccg_vert1 = ccgSubSurf_getEdgeVert1(ccg_edge);
+	edge_verts[0] = GET_INT_FROM_POINTER(ccgSubSurf_getVertVertHandle(ccg_vert0));
+	edge_verts[1] = GET_INT_FROM_POINTER(ccgSubSurf_getVertVertHandle(ccg_vert1));
+}
+
+static int conv_ccg_get_num_edge_faces(const OpenSubdiv_Converter *converter,
+                                       int edge)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGEdge *ccg_edge = ccgSubSurf_getEdge(ss, SET_INT_IN_POINTER(edge));
+	return ccgSubSurf_getEdgeNumFaces(ccg_edge);
+}
+
+static void conv_ccg_get_edge_faces(const OpenSubdiv_Converter *converter,
+                                    int edge,
+                                    int *edge_faces)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGEdge *ccg_edge = ccgSubSurf_getEdge(ss, SET_INT_IN_POINTER(edge));
+	int num_edge_faces = ccgSubSurf_getEdgeNumFaces(ccg_edge);
+	int face;
+	for (face = 0; face < num_edge_faces; face++) {
+		CCGFace *ccg_face = ccgSubSurf_getEdgeFace(ccg_edge, face);
+		edge_faces[face] = GET_INT_FROM_POINTER(ccgSubSurf_getFaceFaceHandle(ccg_face));
+	}
+}
+
+static float conv_ccg_get_edge_sharpness(const OpenSubdiv_Converter *converter,
+                                         int edge)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGEdge *ccg_edge = ccgSubSurf_getEdge(ss, SET_INT_IN_POINTER(edge));
+	/* TODO(sergey): Multiply by subdivision level once CPU evaluator
+	 * is switched to uniform subdivision type.
+	 */
+	return ccg_edge->crease;
+}
+
+static int conv_ccg_get_num_vert_edges(const OpenSubdiv_Converter *converter,
+                                       int vert)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGVert *ccg_vert = ccgSubSurf_getVert(ss, SET_INT_IN_POINTER(vert));
+	return ccgSubSurf_getVertNumEdges(ccg_vert);
+}
+
+static void conv_ccg_get_vert_edges(const OpenSubdiv_Converter *converter,
+                                    int vert,
+                                    int *vert_edges)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGVert *ccg_vert = ccgSubSurf_getVert(ss, SET_INT_IN_POINTER(vert));
+	int num_vert_edges = ccgSubSurf_getVertNumEdges(ccg_vert);
+	int edge;
+	for (edge = 0; edge < num_vert_edges; edge++) {
+		CCGEdge *ccg_edge = ccgSubSurf_getVertEdge(ccg_vert, edge);
+		vert_edges[edge] = GET_INT_FROM_POINTER(ccgSubSurf_getEdgeEdgeHandle(ccg_edge));
+	}
+}
+
+static int conv_ccg_get_num_vert_faces(const OpenSubdiv_Converter *converter,
+                                       int vert)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGVert *ccg_vert = ccgSubSurf_getVert(ss, SET_INT_IN_POINTER(vert));
+	return ccgSubSurf_getVertNumFaces(ccg_vert);
+}
+
+static void conv_ccg_get_vert_faces(const OpenSubdiv_Converter *converter,
+                                    int vert,
+                                    int *vert_faces)
+{
+	CCGSubSurf *ss = converter->user_data;
+	CCGVert *ccg_vert = ccgSubSurf_getVert(ss, SET_INT_IN_POINTER(vert));
+	int num_vert_faces = ccgSubSurf_getVertNumFaces(ccg_vert);
+	int face;
+	for (face = 0; face < num_vert_faces; face++) {
+		CCGFace *ccg_face = ccgSubSurf_getVertFace(ccg_vert, face);
+		vert_faces[face] = GET_INT_FROM_POINTER(ccgSubSurf_getFaceFaceHandle(ccg_face));
+	}
+}
+
+void ccgSubSurf_converter_setup_from_ccg(CCGSubSurf *ss,
+                                         OpenSubdiv_Converter *converter)
+{
+	converter->get_type = conv_ccg_get_bilinear_type;
+
+	converter->get_num_faces = conv_ccg_get_num_faces;
+	converter->get_num_edges = conv_ccg_get_num_edges;
+	converter->get_num_verts = conv_ccg_get_num_verts;
+
+	converter->get_num_face_verts = conv_ccg_get_num_face_verts;
+	converter->get_face_verts = conv_ccg_get_face_verts;
+	converter->get_face_edges = conv_ccg_get_face_edges;
+
+	converter->get_edge_verts = conv_ccg_get_edge_verts;
+	converter->get_num_edge_faces = conv_ccg_get_num_edge_faces;
+	converter->get_edge_faces = conv_ccg_get_edge_faces;
+	converter->get_edge_sharpness = conv_ccg_get_edge_sharpness;
+
+	converter->get_num_vert_edges = conv_ccg_get_num_vert_edges;
+	converter->get_vert_edges = conv_ccg_get_vert_edges;
+	converter->get_num_vert_faces = conv_ccg_get_num_vert_faces;
+	converter->get_vert_faces = conv_ccg_get_vert_faces;
+
+	converter->free_user_data = NULL;
+	converter->user_data = ss;
+}
+
+void ccgSubSurf_converter_free(
+        struct OpenSubdiv_Converter *converter)
+{
+	if (converter->free_user_data) {
+		converter->free_user_data(converter);
+	}
+}
+
+#endif  /* WITH_OPENSUBDIV */
diff --git a/source/blender/gpu/shaders/gpu_shader_geometry.glsl b/source/blender/gpu/shaders/gpu_shader_geometry.glsl
new file mode 100644
index 00000000000..a0ae96a1f72
--- /dev/null
+++ b/source/blender/gpu/shaders/gpu_shader_geometry.glsl
@@ -0,0 +1,100 @@
+uniform int PrimitiveIdBase;
+uniform int osd_active_uv_offset;
+
+varying vec3 varnormal;
+varying vec3 varposition;
+
+in block {
+	VertexData v;
+} inpt[4];
+
+uniform bool osd_flat_shading;
+uniform int osd_fvar_count;
+
+#define INTERP_FACE_VARYING_2(result, fvarOffset, tessCoord)  \
+	{ \
+		vec2 v[4]; \
+		int primOffset = (gl_PrimitiveID + PrimitiveIdBase) * 4; \
+		for (int i = 0; i < 4; ++i) { \
+			int index = (primOffset + i) * osd_fvar_count + fvarOffset; \
+			v[i] = vec2(texelFetch(FVarDataBuffer, index).s, \
+			            texelFetch(FVarDataBuffer, index + 1).s); \
+		} \
+		result = mix(mix(v[0], v[1], tessCoord.s), \
+		             mix(v[3], v[2], tessCoord.s), \
+		             tessCoord.t); \
+	}
+
+uniform samplerBuffer FVarDataBuffer;
+
+out block {
+	VertexData v;
+} outpt;
+
+void set_mtface_vertex_attrs(vec2 st);
+
+void emit_flat(int index, vec3 normal)
+{
+	outpt.v.position = inpt[index].v.position;
+	outpt.v.normal = normal;
+
+	/* Compatibility */
+	varnormal = outpt.v.normal;
+	varposition = outpt.v.position.xyz;
+
+	/* TODO(sergey): Only uniform subdivisions atm. */
+	vec2 quadst[4] = vec2[](vec2(0,0), vec2(1,0), vec2(1,1), vec2(0,1));
+	vec2 st = quadst[index];
+
+	INTERP_FACE_VARYING_2(outpt.v.uv, osd_active_uv_offset, st);
+
+	set_mtface_vertex_attrs(st);
+
+	gl_Position = gl_ProjectionMatrix * inpt[index].v.position;
+	EmitVertex();
+}
+
+void emit_smooth(int index)
+{
+	outpt.v.position = inpt[index].v.position;
+	outpt.v.normal = inpt[index].v.normal;
+
+	/* Compatibility */
+	varnormal = outpt.v.normal;
+	varposition = outpt.v.position.xyz;
+
+	/* TODO(sergey): Only uniform subdivisions atm. */
+	vec2 quadst[4] = vec2[](vec2(0,0), vec2(1,0), vec2(1,1), vec2(0,1));
+	vec2 st = quadst[index];
+
+	INTERP_FACE_VARYING_2(outpt.v.uv, osd_active_uv_offset, st);
+
+	set_mtface_vertex_attrs(st);
+
+	gl_Position = gl_ProjectionMatrix * inpt[index].v.position;
+	EmitVertex();
+}
+
+void main()
+{
+	gl_PrimitiveID = gl_PrimitiveIDIn;
+
+	if (osd_flat_shading) {
+		vec3 A = (inpt[0].v.position - inpt[1].v.position).xyz;
+		vec3 B = (inpt[3].v.position - inpt[1].v.position).xyz;
+		vec3 flat_normal = normalize(cross(B, A));
+		emit_flat(0, flat_normal);
+		emit_flat(1, flat_normal);
+		emit_flat(3, flat_normal);
+		emit_flat(2, flat_normal);
+	}
+	else {
+		emit_smooth(0);
+		emit_smooth(1);
+		emit_smooth(3);
+		emit_smooth(2);
+	}
+	EndPrimitive();
+}
+
+void set_mtface_vertex_attrs(vec2 st) {