diff --git a/opensubdiv/osd/clKernelBundle.cpp b/opensubdiv/osd/clKernelBundle.cpp
index 3e507b70..35a3408c 100644
--- a/opensubdiv/osd/clKernelBundle.cpp
+++ b/opensubdiv/osd/clKernelBundle.cpp
@@ -140,8 +140,7 @@ OsdCLKernelBundle::Compile(cl_context clContext,
 
     cl_int ciErrNum;
 
-    _numVertexElements = numVertexElements;
-    _numVaryingElements = numVaryingElements;
+    _vdesc.Set( numVertexElements, numVaryingElements );
 
     char constantDefine[256];
     snprintf(constantDefine, sizeof(constantDefine),
diff --git a/opensubdiv/osd/clKernelBundle.h b/opensubdiv/osd/clKernelBundle.h
index bd6790c7..e790aa65 100644
--- a/opensubdiv/osd/clKernelBundle.h
+++ b/opensubdiv/osd/clKernelBundle.h
@@ -60,6 +60,7 @@
 #include "../version.h"
 
 #include "../osd/nonCopyable.h"
+#include "../osd/vertexDescriptor.h"
 
 #if defined(__APPLE__)
     #include <OpenCL/opencl.h>
@@ -72,6 +73,7 @@ namespace OpenSubdiv {
 namespace OPENSUBDIV_VERSION {
 
 class OsdCLKernelBundle : OsdNonCopyable<OsdCLKernelBundle> {
+
 public:
     OsdCLKernelBundle();
     ~OsdCLKernelBundle();
@@ -100,14 +102,17 @@ public:
     cl_kernel GetVertexEditAdd() const        { return _clVertexEditAdd; }
 
     struct Match {
-    Match(int numVertexElements, int numVaryingElements) :
-        _numVertexElements(numVertexElements),
-        _numVaryingElements(numVaryingElements) {}
-        bool operator() (const OsdCLKernelBundle *kernel) {
-            return (kernel->_numVertexElements == _numVertexElements
-                    && kernel->_numVaryingElements == _numVaryingElements);
+    
+        /// Constructor
+        Match(int numVertexElements, int numVaryingElements)
+            : vdesc(numVertexElements, numVaryingElements) {
         }
-        int _numVertexElements, _numVaryingElements;
+        
+        bool operator() (OsdCLKernelBundle const *kernel) {
+            return vdesc == kernel->_vdesc;
+        }
+        
+        OsdVertexDescriptor vdesc;
     };
 
     friend struct Match;
@@ -126,8 +131,7 @@ protected:
               _clLoopVertexB,
               _clVertexEditAdd;
 
-    int _numVertexElements,
-        _numVaryingElements;
+    OsdVertexDescriptor _vdesc;
 };
 
 }  // end namespace OPENSUBDIV_VERSION
diff --git a/opensubdiv/osd/cpuComputeContext.cpp b/opensubdiv/osd/cpuComputeContext.cpp
index 33f8ae38..ad26fe17 100644
--- a/opensubdiv/osd/cpuComputeContext.cpp
+++ b/opensubdiv/osd/cpuComputeContext.cpp
@@ -168,7 +168,6 @@ OsdCpuComputeContext::OsdCpuComputeContext(FarMesh<OsdVertex> const *farMesh) {
             _editTables.push_back(new OsdCpuHEditTable(edit));
         }
     }
-    _vdesc = 0;
     _currentVertexBuffer = 0;
     _currentVaryingBuffer = 0;
 }
@@ -181,7 +180,6 @@ OsdCpuComputeContext::~OsdCpuComputeContext() {
     for (size_t i = 0; i < _editTables.size(); ++i) {
         delete _editTables[i];
     }
-    delete _vdesc;
 }
 
 const OsdCpuTable *
@@ -190,12 +188,6 @@ OsdCpuComputeContext::GetTable(int tableIndex) const {
     return _tables[tableIndex];
 }
 
-OsdVertexDescriptor *
-OsdCpuComputeContext::GetVertexDescriptor() const {
-
-    return _vdesc;
-}
-
 int
 OsdCpuComputeContext::GetNumEditTables() const {
 
diff --git a/opensubdiv/osd/cpuComputeContext.h b/opensubdiv/osd/cpuComputeContext.h
index ecc7f2ef..3a53f8ce 100644
--- a/opensubdiv/osd/cpuComputeContext.h
+++ b/opensubdiv/osd/cpuComputeContext.h
@@ -152,16 +152,14 @@ public:
 
         int numVertexElements = vertex ? vertex->GetNumElements() : 0;
         int numVaryingElements = varying ? varying->GetNumElements() : 0;
-        _vdesc = new OsdVertexDescriptor(numVertexElements, numVaryingElements);
+        _vdesc.Set(numVertexElements, numVaryingElements);
     }
 
     /// Unbinds any previously bound vertex and varying data buffers.
     void Unbind() {
         _currentVertexBuffer = 0;
         _currentVaryingBuffer = 0;
-
-        delete _vdesc;
-        _vdesc = 0;
+        _vdesc.Reset();
     }
 
     /// Returns one of the vertex refinement tables.
@@ -174,7 +172,9 @@ public:
     ///
     /// @return a descriptor for the format of the vertex data currently bound
     ///
-    OsdVertexDescriptor * GetVertexDescriptor() const;
+    OsdVertexDescriptor const & GetVertexDescriptor() const {
+        return _vdesc;
+    }
 
     /// Returns the number of hierarchical edit tables
     int GetNumEditTables() const;
@@ -201,7 +201,7 @@ private:
     float *_currentVertexBuffer, 
           *_currentVaryingBuffer;
 
-    OsdVertexDescriptor *_vdesc;
+    OsdVertexDescriptor _vdesc;
 };
 
 }  // end namespace OPENSUBDIV_VERSION
diff --git a/opensubdiv/osd/cpuKernel.cpp b/opensubdiv/osd/cpuKernel.cpp
index 1fc0edf2..93d163d3 100644
--- a/opensubdiv/osd/cpuKernel.cpp
+++ b/opensubdiv/osd/cpuKernel.cpp
@@ -64,7 +64,7 @@ namespace OpenSubdiv {
 namespace OPENSUBDIV_VERSION {
 
 void OsdCpuComputeFace(
-    const OsdVertexDescriptor *vdesc, float * vertex, float * varying,
+    OsdVertexDescriptor const &vdesc, float * vertex, float * varying,
     const int *F_IT, const int *F_ITa, int vertexOffset, int tableOffset,
     int start, int end) {
 
@@ -77,19 +77,19 @@ void OsdCpuComputeFace(
         // XXX: should use local vertex struct variable instead of
         // accumulating directly into global memory.
         int dstIndex = i + vertexOffset - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
         for (int j = 0; j < n; ++j) {
             int index = F_IT[h+j];
-            vdesc->AddWithWeight(vertex, dstIndex, index, weight);
-            vdesc->AddVaryingWithWeight(varying, dstIndex, index, weight);
+            vdesc.AddWithWeight(vertex, dstIndex, index, weight);
+            vdesc.AddVaryingWithWeight(varying, dstIndex, index, weight);
         }
     }
 
 }
 
 void OsdCpuComputeEdge(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *E_IT, const float *E_W, int vertexOffset, int tableOffset,
     int start, int end) {
 
@@ -102,25 +102,25 @@ void OsdCpuComputeEdge(
         float vertWeight = E_W[i*2+0];
 
         int dstIndex = i + vertexOffset - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
-        vdesc->AddWithWeight(vertex, dstIndex, eidx0, vertWeight);
-        vdesc->AddWithWeight(vertex, dstIndex, eidx1, vertWeight);
+        vdesc.AddWithWeight(vertex, dstIndex, eidx0, vertWeight);
+        vdesc.AddWithWeight(vertex, dstIndex, eidx1, vertWeight);
 
         if (eidx2 != -1) {
             float faceWeight = E_W[i*2+1];
 
-            vdesc->AddWithWeight(vertex, dstIndex, eidx2, faceWeight);
-            vdesc->AddWithWeight(vertex, dstIndex, eidx3, faceWeight);
+            vdesc.AddWithWeight(vertex, dstIndex, eidx2, faceWeight);
+            vdesc.AddWithWeight(vertex, dstIndex, eidx3, faceWeight);
         }
 
-        vdesc->AddVaryingWithWeight(varying, dstIndex, eidx0, 0.5f);
-        vdesc->AddVaryingWithWeight(varying, dstIndex, eidx1, 0.5f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, eidx0, 0.5f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, eidx1, 0.5f);
     }
 }
 
 void OsdCpuComputeVertexA(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *V_ITa, const float *V_W, int vertexOffset, int tableOffset,
     int start, int end, int pass) {
 
@@ -141,23 +141,23 @@ void OsdCpuComputeVertexA(
         int dstIndex = i + vertexOffset - tableOffset;
 
         if (not pass)
-            vdesc->Clear(vertex, varying, dstIndex);
+            vdesc.Clear(vertex, varying, dstIndex);
 
         if (eidx0 == -1 || (pass == 0 && (n == -1))) {
-            vdesc->AddWithWeight(vertex, dstIndex, p, weight);
+            vdesc.AddWithWeight(vertex, dstIndex, p, weight);
         } else {
-            vdesc->AddWithWeight(vertex, dstIndex, p, weight * 0.75f);
-            vdesc->AddWithWeight(vertex, dstIndex, eidx0, weight * 0.125f);
-            vdesc->AddWithWeight(vertex, dstIndex, eidx1, weight * 0.125f);
+            vdesc.AddWithWeight(vertex, dstIndex, p, weight * 0.75f);
+            vdesc.AddWithWeight(vertex, dstIndex, eidx0, weight * 0.125f);
+            vdesc.AddWithWeight(vertex, dstIndex, eidx1, weight * 0.125f);
         }
 
         if (not pass)
-            vdesc->AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
+            vdesc.AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
     }
 }
 
 void OsdCpuComputeVertexB(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *V_ITa, const int *V_IT, const float *V_W,
     int vertexOffset, int tableOffset, int start, int end) {
 
@@ -171,20 +171,20 @@ void OsdCpuComputeVertexB(
         float wv = (n-2.0f) * n * wp;
 
         int dstIndex = i + vertexOffset - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
-        vdesc->AddWithWeight(vertex, dstIndex, p, weight * wv);
+        vdesc.AddWithWeight(vertex, dstIndex, p, weight * wv);
 
         for (int j = 0; j < n; ++j) {
-            vdesc->AddWithWeight(vertex, dstIndex, V_IT[h+j*2], weight * wp);
-            vdesc->AddWithWeight(vertex, dstIndex, V_IT[h+j*2+1], weight * wp);
+            vdesc.AddWithWeight(vertex, dstIndex, V_IT[h+j*2], weight * wp);
+            vdesc.AddWithWeight(vertex, dstIndex, V_IT[h+j*2+1], weight * wp);
         }
-        vdesc->AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
     }
 }
 
 void OsdCpuComputeLoopVertexB(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *V_ITa, const int *V_IT, const float *V_W,
     int vertexOffset, int tableOffset, int start, int end) {
 
@@ -200,19 +200,19 @@ void OsdCpuComputeLoopVertexB(
         beta = (0.625f - beta) * wp;
 
         int dstIndex = i + vertexOffset - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
-        vdesc->AddWithWeight(vertex, dstIndex, p, weight * (1.0f - (beta * n)));
+        vdesc.AddWithWeight(vertex, dstIndex, p, weight * (1.0f - (beta * n)));
 
         for (int j = 0; j < n; ++j)
-            vdesc->AddWithWeight(vertex, dstIndex, V_IT[h+j], weight * beta);
+            vdesc.AddWithWeight(vertex, dstIndex, V_IT[h+j], weight * beta);
 
-        vdesc->AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
     }
 }
 
 void OsdCpuComputeBilinearEdge(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *E_IT, int vertexOffset, int tableOffset, int start, int end) {
 
     for (int i = start + tableOffset; i < end + tableOffset; i++) {
@@ -220,58 +220,58 @@ void OsdCpuComputeBilinearEdge(
         int eidx1 = E_IT[2*i+1];
 
         int dstIndex = i + vertexOffset - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
-        vdesc->AddWithWeight(vertex, dstIndex, eidx0, 0.5f);
-        vdesc->AddWithWeight(vertex, dstIndex, eidx1, 0.5f);
+        vdesc.AddWithWeight(vertex, dstIndex, eidx0, 0.5f);
+        vdesc.AddWithWeight(vertex, dstIndex, eidx1, 0.5f);
 
-        vdesc->AddVaryingWithWeight(varying, dstIndex, eidx0, 0.5f);
-        vdesc->AddVaryingWithWeight(varying, dstIndex, eidx1, 0.5f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, eidx0, 0.5f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, eidx1, 0.5f);
     }
 }
 
 void OsdCpuComputeBilinearVertex(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *V_ITa, int vertexOffset, int tableOffset, int start, int end) {
 
     for (int i = start + tableOffset; i < end + tableOffset; i++) {
         int p = V_ITa[i];
 
         int dstIndex = i + vertexOffset - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
-        vdesc->AddWithWeight(vertex, dstIndex, p, 1.0f);
-        vdesc->AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
+        vdesc.AddWithWeight(vertex, dstIndex, p, 1.0f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
     }
 }
 
 void OsdCpuEditVertexAdd(
-    const OsdVertexDescriptor *vdesc, float *vertex,
+    OsdVertexDescriptor const &vdesc, float *vertex,
     int primVarOffset, int primVarWidth, int vertexOffset, int tableOffset,
     int start, int end,
     const unsigned int *editIndices, const float *editValues) {
 
     for (int i = start+tableOffset; i < end+tableOffset; i++) {
-        vdesc->ApplyVertexEditAdd(vertex,
-                                  primVarOffset,
-                                  primVarWidth,
-                                  editIndices[i] + vertexOffset,
-                                  &editValues[i*primVarWidth]);
+        vdesc.ApplyVertexEditAdd(vertex,
+                                 primVarOffset,
+                                 primVarWidth,
+                                 editIndices[i] + vertexOffset,
+                                 &editValues[i*primVarWidth]);
     }
 }
 
 void OsdCpuEditVertexSet(
-    const OsdVertexDescriptor *vdesc, float *vertex,
+    OsdVertexDescriptor const &vdesc, float *vertex,
     int primVarOffset, int primVarWidth, int vertexOffset, int tableOffset,
     int start, int end,
     const unsigned int *editIndices, const float *editValues) {
 
     for (int i = start+tableOffset; i < end+tableOffset; i++) {
-        vdesc->ApplyVertexEditSet(vertex,
-                                  primVarOffset,
-                                  primVarWidth,
-                                  editIndices[i] + vertexOffset,
-                                  &editValues[i*primVarWidth]);
+        vdesc.ApplyVertexEditSet(vertex,
+                                 primVarOffset,
+                                 primVarWidth,
+                                 editIndices[i] + vertexOffset,
+                                 &editValues[i*primVarWidth]);
     }
 }
 
diff --git a/opensubdiv/osd/cpuKernel.h b/opensubdiv/osd/cpuKernel.h
index a2f9e39b..077d6ab9 100644
--- a/opensubdiv/osd/cpuKernel.h
+++ b/opensubdiv/osd/cpuKernel.h
@@ -66,57 +66,57 @@ namespace OPENSUBDIV_VERSION {
 
 struct OsdVertexDescriptor;
 
-void OsdCpuComputeFace(const OsdVertexDescriptor *vdesc,
+void OsdCpuComputeFace(OsdVertexDescriptor const &vdesc,
                        float * vertex, float * varying,
                        const int *F_IT, const int *F_ITa,
                        int vertexOffset, int tableOffset,
                        int start, int end);
 
-void OsdCpuComputeEdge(const OsdVertexDescriptor *vdesc,
+void OsdCpuComputeEdge(OsdVertexDescriptor const &vdesc,
                        float *vertex, float * varying,
                        const int *E_IT, const float *E_ITa,
                        int vertexOffset, int tableOffset,
                        int start, int end);
 
-void OsdCpuComputeVertexA(const OsdVertexDescriptor *vdesc,
+void OsdCpuComputeVertexA(OsdVertexDescriptor const &vdesc,
                           float *vertex, float * varying,
                           const int *V_ITa, const float *V_IT,
                           int vertexOffset, int tableOffset,
                           int start, int end, int pass);
 
-void OsdCpuComputeVertexB(const OsdVertexDescriptor *vdesc,
+void OsdCpuComputeVertexB(OsdVertexDescriptor const &vdesc,
                           float *vertex, float * varying,
                           const int *V_ITa, const int *V_IT, const float *V_W,
                           int vertexOffset, int tableOffset,
                           int start, int end);
 
-void OsdCpuComputeLoopVertexB(const OsdVertexDescriptor *vdesc,
+void OsdCpuComputeLoopVertexB(OsdVertexDescriptor const &vdesc,
                               float *vertex, float * varying,
                               const int *V_ITa, const int *V_IT,
                               const float *V_W,
                               int vertexOffset, int tableOffset,
                               int start, int end);
 
-void OsdCpuComputeBilinearEdge(const OsdVertexDescriptor *vdesc,
+void OsdCpuComputeBilinearEdge(OsdVertexDescriptor const &vdesc,
                                float *vertex, float * varying,
                                const int *E_IT,
                                int vertexOffset, int tableOffset,
                                int start, int end);
 
-void OsdCpuComputeBilinearVertex(const OsdVertexDescriptor *vdesc,
+void OsdCpuComputeBilinearVertex(OsdVertexDescriptor const &vdesc,
                                  float *vertex, float * varying,
                                  const int *V_ITa,
                                  int vertexOffset, int tableOffset,
                                  int start, int end);
 
-void OsdCpuEditVertexAdd(const OsdVertexDescriptor *vdesc, float *vertex,
+void OsdCpuEditVertexAdd(OsdVertexDescriptor const &vdesc, float *vertex,
                          int primVarOffset, int primVarWidth,
                          int vertexOffset, int tableOffset,
                          int start, int end,
                          const unsigned int *editIndices,
                          const float *editValues);
 
-void OsdCpuEditVertexSet(const OsdVertexDescriptor *vdesc, float *vertex,
+void OsdCpuEditVertexSet(OsdVertexDescriptor const &vdesc, float *vertex,
                          int primVarOffset, int primVarWidth,
                          int vertexOffset, int tableOffset,
                          int start, int end,
diff --git a/opensubdiv/osd/cudaComputeContext.cpp b/opensubdiv/osd/cudaComputeContext.cpp
index 7d55d48f..96c62122 100644
--- a/opensubdiv/osd/cudaComputeContext.cpp
+++ b/opensubdiv/osd/cudaComputeContext.cpp
@@ -201,18 +201,6 @@ OsdCudaComputeContext::GetCurrentVaryingBuffer() const {
     return _currentVaryingBuffer;
 }
 
-int
-OsdCudaComputeContext::GetCurrentVertexNumElements() const {
-
-    return _numVertexElements;
-}
-
-int
-OsdCudaComputeContext::GetCurrentVaryingNumElements() const {
-
-    return _numVaryingElements;
-}
-
 OsdCudaComputeContext *
 OsdCudaComputeContext::Create(FarMesh<OsdVertex> const *farmesh) {
 
diff --git a/opensubdiv/osd/cudaComputeContext.h b/opensubdiv/osd/cudaComputeContext.h
index 19e47426..1be66563 100644
--- a/opensubdiv/osd/cudaComputeContext.h
+++ b/opensubdiv/osd/cudaComputeContext.h
@@ -61,6 +61,7 @@
 
 #include "../far/vertexEditTables.h"
 #include "../osd/vertex.h"
+#include "../osd/vertexDescriptor.h"
 #include "../osd/nonCopyable.h"
 
 #include <stdlib.h>
@@ -146,18 +147,18 @@ public:
 
         if (vertex) {
             _currentVertexBuffer = static_cast<float*>(vertex->BindCudaBuffer());
-            _numVertexElements = vertex->GetNumElements();
+            _vdesc.numVertexElements = vertex->GetNumElements();
         } else {
             _currentVertexBuffer = 0;
-            _numVertexElements = 0;
+            _vdesc.numVertexElements = 0;
         }
 
         if (varying) {
             _currentVaryingBuffer = static_cast<float*>(varying->BindCudaBuffer());
-            _numVertexElements = varying->GetNumElements();
+            _vdesc.numVertexElements = varying->GetNumElements();
         } else {
             _currentVaryingBuffer = 0;
-            _numVaryingElements = 0;
+            _vdesc.numVaryingElements = 0;
         }
     }
 
@@ -185,11 +186,17 @@ public:
     /// Returns a pointer to the vertex-interpolated data
     float * GetCurrentVertexBuffer() const;
 
+    /// Returns a pointer to the varying-interpolated data
     float * GetCurrentVaryingBuffer() const;
 
-    int GetCurrentVertexNumElements() const;
+    /// Returns an OsdVertexDescriptor if vertex buffers have been bound.
+    ///
+    /// @return a descriptor for the format of the vertex data currently bound
+    ///
+    OsdVertexDescriptor const & GetVertexDescriptor() const {
+        return _vdesc;
+    }
 
-    int GetCurrentVaryingNumElements() const;
 
 protected:
     explicit OsdCudaComputeContext(FarMesh<OsdVertex> const *farMesh);
@@ -202,7 +209,7 @@ private:
     float *_currentVertexBuffer, // cuda buffers
           *_currentVaryingBuffer;
 
-    int _numVertexElements, _numVaryingElements;
+    OsdVertexDescriptor _vdesc;
 };
 
 }  // end namespace OPENSUBDIV_VERSION
diff --git a/opensubdiv/osd/cudaComputeController.cpp b/opensubdiv/osd/cudaComputeController.cpp
index b82c06a4..8b6ca88e 100644
--- a/opensubdiv/osd/cudaComputeController.cpp
+++ b/opensubdiv/osd/cudaComputeController.cpp
@@ -129,8 +129,8 @@ OsdCudaComputeController::ApplyBilinearFaceVerticesKernel(
     OsdCudaComputeFace(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(F_IT->GetCudaMemory()),
         static_cast<int*>(F_ITa->GetCudaMemory()),
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@@ -150,8 +150,8 @@ OsdCudaComputeController::ApplyBilinearEdgeVerticesKernel(
     OsdCudaComputeBilinearEdge(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(E_IT->GetCudaMemory()),
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
 }
@@ -170,8 +170,8 @@ OsdCudaComputeController::ApplyBilinearVertexVerticesKernel(
     OsdCudaComputeBilinearVertex(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(V_ITa->GetCudaMemory()),
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
 }
@@ -192,8 +192,8 @@ OsdCudaComputeController::ApplyCatmarkFaceVerticesKernel(
     OsdCudaComputeFace(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(F_IT->GetCudaMemory()),
         static_cast<int*>(F_ITa->GetCudaMemory()),
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@@ -215,8 +215,8 @@ OsdCudaComputeController::ApplyCatmarkEdgeVerticesKernel(
     OsdCudaComputeEdge(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(E_IT->GetCudaMemory()),
         static_cast<float*>(E_W->GetCudaMemory()),
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@@ -240,8 +240,8 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelB(
     OsdCudaComputeVertexB(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(V_ITa->GetCudaMemory()),
         static_cast<int*>(V_IT->GetCudaMemory()),
         static_cast<float*>(V_W->GetCudaMemory()),
@@ -264,8 +264,8 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelA1(
     OsdCudaComputeVertexA(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(V_ITa->GetCudaMemory()),
         static_cast<float*>(V_W->GetCudaMemory()),
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
@@ -287,8 +287,8 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelA2(
     OsdCudaComputeVertexA(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(V_ITa->GetCudaMemory()),
         static_cast<float*>(V_W->GetCudaMemory()),
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
@@ -310,8 +310,8 @@ OsdCudaComputeController::ApplyLoopEdgeVerticesKernel(
     OsdCudaComputeEdge(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(E_IT->GetCudaMemory()),
         static_cast<float*>(E_W->GetCudaMemory()),
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@@ -335,8 +335,8 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelB(
     OsdCudaComputeLoopVertexB(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(V_ITa->GetCudaMemory()),
         static_cast<int*>(V_IT->GetCudaMemory()),
         static_cast<float*>(V_W->GetCudaMemory()),
@@ -359,8 +359,8 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelA1(
     OsdCudaComputeVertexA(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(V_ITa->GetCudaMemory()),
         static_cast<float*>(V_W->GetCudaMemory()),
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
@@ -382,8 +382,8 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelA2(
     OsdCudaComputeVertexA(
         context->GetCurrentVertexBuffer(),
         context->GetCurrentVaryingBuffer(),
-        context->GetCurrentVertexNumElements()-3,
-        context->GetCurrentVaryingNumElements(),
+        context->GetVertexDescriptor().numVertexElements-3,
+        context->GetVertexDescriptor().numVaryingElements,
         static_cast<int*>(V_ITa->GetCudaMemory()),
         static_cast<float*>(V_W->GetCudaMemory()),
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
@@ -406,7 +406,7 @@ OsdCudaComputeController::ApplyVertexEdits(
     if (edit->GetOperation() == FarVertexEdit::Add) {
         OsdCudaEditVertexAdd(
             context->GetCurrentVertexBuffer(),
-            context->GetCurrentVertexNumElements()-3,
+            context->GetVertexDescriptor().numVertexElements-3,
             edit->GetPrimvarOffset(),
             edit->GetPrimvarWidth(),
             batch.GetVertexOffset(),
diff --git a/opensubdiv/osd/d3d11ComputeContext.cpp b/opensubdiv/osd/d3d11ComputeContext.cpp
index 498bbd14..9aa18267 100644
--- a/opensubdiv/osd/d3d11ComputeContext.cpp
+++ b/opensubdiv/osd/d3d11ComputeContext.cpp
@@ -258,18 +258,6 @@ OsdD3D11ComputeContext::GetCurrentVaryingBufferUAV() const {
     return _currentVaryingBufferUAV;
 }
 
-int
-OsdD3D11ComputeContext::GetNumCurrentVertexElements() const {
-
-    return _numVertexElements;
-}
-
-int
-OsdD3D11ComputeContext::GetNumCurrentVaryingElements() const {
-    
-    return _numVaryingElements;
-}
-
 OsdD3D11ComputeKernelBundle *
 OsdD3D11ComputeContext::GetKernelBundle() const {
 
diff --git a/opensubdiv/osd/d3d11ComputeContext.h b/opensubdiv/osd/d3d11ComputeContext.h
index 45370a3e..82da7689 100644
--- a/opensubdiv/osd/d3d11ComputeContext.h
+++ b/opensubdiv/osd/d3d11ComputeContext.h
@@ -157,8 +157,8 @@ public:
         _currentVertexBufferUAV = vertex ? vertex->BindD3D11UAV(_deviceContext) : 0;
         _currentVaryingBufferUAV = varying ? varying->BindD3D11UAV(_deviceContext) : 0;
 
-        _numVertexElements = vertex ? vertex->GetNumElements() : 0;
-        _numVaryingElements = varying ? varying->GetNumElements() : 0;
+        _vdesc.numVertexElements = vertex ? vertex->GetNumElements() : 0;
+        _vdesc.numVaryingElements = varying ? varying->GetNumElements() : 0;
 
         bindShaderStorageBuffers();
     }
@@ -192,9 +192,13 @@ public:
     /// Returns a handle to the varying-interpolated buffer
     ID3D11UnorderedAccessView * GetCurrentVaryingBufferUAV() const;
 
-    int GetNumCurrentVertexElements() const;
-
-    int GetNumCurrentVaryingElements() const;
+    /// Returns an OsdVertexDescriptor if vertex buffers have been bound.
+    ///
+    /// @return a descriptor for the format of the vertex data currently bound
+    ///
+    OsdVertexDescriptor const & GetVertexDescriptor() const {
+        return _vdesc;
+    }
 
     OsdD3D11ComputeKernelBundle * GetKernelBundle() const;
 
@@ -221,8 +225,7 @@ private:
 
     ID3D11DeviceContext *_deviceContext;
 
-    int _numVertexElements;
-    int _numVaryingElements;
+    OsdVertexDescriptor _vdesc;
 
     ID3D11UnorderedAccessView * _currentVertexBufferUAV,
                               * _currentVaryingBufferUAV;
diff --git a/opensubdiv/osd/d3d11KernelBundle.cpp b/opensubdiv/osd/d3d11KernelBundle.cpp
index ebdaafb5..ce395191 100644
--- a/opensubdiv/osd/d3d11KernelBundle.cpp
+++ b/opensubdiv/osd/d3d11KernelBundle.cpp
@@ -115,8 +115,7 @@ bool
 OsdD3D11ComputeKernelBundle::Compile(int numVertexElements,
                                      int numVaryingElements) {
 
-    _numVertexElements = numVertexElements;
-    _numVaryingElements = numVaryingElements;
+    _vdesc.Set( numVertexElements, numVaryingElements );
 
     DWORD dwShaderFlags = D3DCOMPILE_ENABLE_STRICTNESS;
 #ifdef _DEBUG
diff --git a/opensubdiv/osd/d3d11KernelBundle.h b/opensubdiv/osd/d3d11KernelBundle.h
index 3cbf9e2e..a58d855c 100644
--- a/opensubdiv/osd/d3d11KernelBundle.h
+++ b/opensubdiv/osd/d3d11KernelBundle.h
@@ -71,95 +71,100 @@ struct ID3D11DeviceContext;
 namespace OpenSubdiv {
 namespace OPENSUBDIV_VERSION {
 
-    class OsdD3D11ComputeKernelBundle
-        : OsdNonCopyable<OsdD3D11ComputeKernelBundle> {
-    public:
-        OsdD3D11ComputeKernelBundle(ID3D11DeviceContext *deviceContext);
-        ~OsdD3D11ComputeKernelBundle();
+class OsdD3D11ComputeKernelBundle : OsdNonCopyable<OsdD3D11ComputeKernelBundle> {
 
-        bool Compile(int numVertexElements, int numVaryingElements);
+public:
+    /// Constructor
+    OsdD3D11ComputeKernelBundle(ID3D11DeviceContext *deviceContext);
+    
+    /// Destructor
+    ~OsdD3D11ComputeKernelBundle();
 
-        void ApplyBilinearFaceVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    bool Compile(int numVertexElements, int numVaryingElements);
 
-        void ApplyBilinearEdgeVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyBilinearFaceVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyBilinearVertexVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyBilinearEdgeVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkFaceVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyBilinearVertexVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkEdgeVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyCatmarkFaceVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkVertexVerticesKernelB(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyCatmarkEdgeVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkVertexVerticesKernelA(
-            int vertexOffset, int tableOffset, int start, int end, bool pass);
+    void ApplyCatmarkVertexVerticesKernelB(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyLoopEdgeVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyCatmarkVertexVerticesKernelA(
+        int vertexOffset, int tableOffset, int start, int end, bool pass);
 
-        void ApplyLoopVertexVerticesKernelB(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyLoopEdgeVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyLoopVertexVerticesKernelA(
-            int vertexOffset, int tableOffset, int start, int end, bool pass);
+    void ApplyLoopVertexVerticesKernelB(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyEditAdd(int primvarOffset, int primvarWidth,
-                          int vertexOffset, int tableOffset, int start, int end);
+    void ApplyLoopVertexVerticesKernelA(
+        int vertexOffset, int tableOffset, int start, int end, bool pass);
 
-        struct Match {
-        Match(int numVertexElements, int numVaryingElements) :
-            _numVertexElements(numVertexElements),
-                _numVaryingElements(numVaryingElements) {}
-            bool operator() (const OsdD3D11ComputeKernelBundle *kernel) {
-                return (kernel->_numVertexElements == _numVertexElements
-                        && kernel->_numVaryingElements == _numVaryingElements);
-            }
-            int _numVertexElements, _numVaryingElements;
-        };
+    void ApplyEditAdd(int primvarOffset, int primvarWidth,
+                      int vertexOffset, int tableOffset, int start, int end);
 
-        friend struct Match;
+    struct Match {
 
-    protected:
-        struct KernelCB;
-        void dispatchCompute(ID3D11ClassInstance * kernel, KernelCB const & args);
+        /// Constructor
+        Match(int numVertexElements, int numVaryingElements)
+            : vdesc(numVertexElements, numVaryingElements) {
+        }
 
-        ID3D11DeviceContext * _deviceContext;
+        bool operator() (OsdD3D11ComputeKernelBundle const *kernel) {
+            return vdesc == kernel->_vdesc;
+        }
 
-        ID3D11ComputeShader * _computeShader;
-        ID3D11ClassLinkage * _classLinkage;
-
-        // constant buffer
-        ID3D11Buffer * _kernelCB;
-
-        // general face-vertex kernel (all schemes)
-        ID3D11ClassInstance * _kernelComputeFace;
-        // edge-vertex kernel (catmark + loop schemes)
-        ID3D11ClassInstance * _kernelComputeEdge;
-        // edge-vertex kernel (bilinear scheme)
-        ID3D11ClassInstance * _kernelComputeBilinearEdge;
-        // vertex-vertex kernel (bilinear scheme)
-        ID3D11ClassInstance * _kernelComputeVertex;
-        // vertex-vertex kernel A (catmark + loop schemes)
-        ID3D11ClassInstance * _kernelComputeVertexA;
-        // vertex-vertex kernel B (catmark scheme)
-        ID3D11ClassInstance * _kernelComputeCatmarkVertexB;
-        // vertex-vertex kernel B (loop scheme)
-        ID3D11ClassInstance * _kernelComputeLoopVertexB;
-        // hedit kernel (add)
-        ID3D11ClassInstance * _kernelEditAdd;
-
-        int _numVertexElements,
-            _numVaryingElements;
-
-        int _workGroupSize;
+        OsdVertexDescriptor vdesc;
     };
 
+    friend struct Match;
+
+protected:
+    struct KernelCB;
+    void dispatchCompute(ID3D11ClassInstance * kernel, KernelCB const & args);
+
+    ID3D11DeviceContext * _deviceContext;
+
+    ID3D11ComputeShader * _computeShader;
+    ID3D11ClassLinkage * _classLinkage;
+
+
+    ID3D11Buffer * _kernelCB; // constant buffer
+
+
+    ID3D11ClassInstance * _kernelComputeFace; // general face-vertex kernel (all schemes)
+
+    ID3D11ClassInstance * _kernelComputeEdge; // edge-vertex kernel (catmark + loop schemes)
+
+    ID3D11ClassInstance * _kernelComputeBilinearEdge; // edge-vertex kernel (bilinear scheme)
+
+    ID3D11ClassInstance * _kernelComputeVertex; // vertex-vertex kernel (bilinear scheme)
+
+    ID3D11ClassInstance * _kernelComputeVertexA; // vertex-vertex kernel A (catmark + loop schemes)
+
+    ID3D11ClassInstance * _kernelComputeCatmarkVertexB; // vertex-vertex kernel B (catmark scheme)
+
+    ID3D11ClassInstance * _kernelComputeLoopVertexB; // vertex-vertex kernel B (loop scheme)
+
+    ID3D11ClassInstance * _kernelEditAdd; // hedit kernel (add)
+
+    int _workGroupSize;
+
+    OsdVertexDescriptor _vdesc;
+};
+
 }  // end namespace OPENSUBDIV_VERSION
 using namespace OPENSUBDIV_VERSION;
 
diff --git a/opensubdiv/osd/drawContext.cpp b/opensubdiv/osd/drawContext.cpp
index 5063a3be..ea75a396 100644
--- a/opensubdiv/osd/drawContext.cpp
+++ b/opensubdiv/osd/drawContext.cpp
@@ -87,6 +87,51 @@ OsdDrawContext::ConvertPatchArrays(FarPatchTables::PatchArrayVector const &farPa
             osdPatchArrays.push_back(PatchArray(desc, parray.GetArrayRange()));
         }
     }
+/*    
+#if defined(GL_ES_VERSION_2_0)
+        // XXX: farmesh should have FarDensePatchTable for dense mesh indices.
+        //      instead of GetFaceVertices().
+        const FarSubdivisionTables<OsdVertex> *tables = farMesh->GetSubdivisionTables();
+        int level = tables->GetMaxLevel();
+        const std::vector<int> &indices = farMesh->GetFaceVertices(level-1);
+
+        int numIndices = (int)indices.size();
+
+        // Allocate and fill index buffer.
+        glGenBuffers(1, &patchIndexBuffer);
+        glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, patchIndexBuffer);
+        glBufferData(GL_ELEMENT_ARRAY_BUFFER,
+                     numIndices * sizeof(unsigned int), &(indices[0]), GL_STATIC_DRAW);
+
+
+        // OpenGLES 2 supports only triangle topologies for filled
+        // primitives i.e. not QUADS or PATCHES or LINES_ADJACENCY
+        // For the convenience of clients build build a triangles
+        // index buffer by splitting quads.
+        int numQuads = indices.size() / 4;
+        int numTrisIndices = numQuads * 6;
+
+        std::vector<short> trisIndices;
+        trisIndices.reserve(numTrisIndices);
+        for (int i=0; i<numQuads; ++i) {
+            const int * quad = &indices[i*4];
+            trisIndices.push_back(short(quad[0]));
+            trisIndices.push_back(short(quad[1]));
+            trisIndices.push_back(short(quad[2]));
+
+            trisIndices.push_back(short(quad[2]));
+            trisIndices.push_back(short(quad[3]));
+            trisIndices.push_back(short(quad[0]));
+        }
+
+        // Allocate and fill triangles index buffer.
+        glGenBuffers(1, &patchTrianglesIndexBuffer);
+        glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, patchTrianglesIndexBuffer);
+        glBufferData(GL_ELEMENT_ARRAY_BUFFER,
+                     numTrisIndices * sizeof(short), &(trisIndices[0]), GL_STATIC_DRAW);
+#endif
+*/    
+    
 }
 
 } // end namespace OPENSUBDIV_VERSION
diff --git a/opensubdiv/osd/glDrawContext.cpp b/opensubdiv/osd/glDrawContext.cpp
index 86343e12..e72977b2 100644
--- a/opensubdiv/osd/glDrawContext.cpp
+++ b/opensubdiv/osd/glDrawContext.cpp
@@ -155,50 +155,26 @@ OsdGLDrawContext::create(FarPatchTables const * patchTables, bool requireFVarDat
                  ptables.size() * sizeof(unsigned int), &ptables[0], GL_STATIC_DRAW);
     glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, 0);
     
-    int patchOffset=0;
-    
-    // Process PatchArrays
-    if (IsAdaptive()) {
-        OsdDrawContext::ConvertPatchArrays(patchTables->GetPatchArrayVector(),
-            patchArrays, patchTables->GetMaxValence(), 0);
-
-        // allocate and initialize additional buffer data
-        FarPatchTables::VertexValenceTable const &
-            valenceTable = patchTables->GetVertexValenceTable();
-
-        // create vertex valence buffer and vertex texture
-        if (not valenceTable.empty()) {
-            vertexValenceTextureBuffer = createTextureBuffer(valenceTable, GL_R32I);
-
-            // also create vertex texture buffer (will be updated in UpdateVertexTexture())
-            glGenTextures(1, &vertexTextureBuffer);
-        }
-
-        // create quad offset table buffer
-        FarPatchTables::QuadOffsetTable const &
-            quadOffsetTable = patchTables->GetQuadOffsetTable();
-
-        if (not quadOffsetTable.empty())
-            quadOffsetTextureBuffer = createTextureBuffer(quadOffsetTable, GL_R32I);
-            
-    } else {
-    
-        // Create a single OsdPatchArray from the FarPatchTables::PatchArray
-        // corresponding to the highest level of subdivision
-        
-        // XXXX manuelk : client code will want control over this
-        FarPatchTables::PatchArray const & fpa = 
-            *(patchTables->GetPatchArrayVector().rbegin());
-
-        OsdDrawContext::PatchDescriptor desc( 
-            fpa.GetDescriptor(), patchTables->GetMaxValence(), 0, 0);
-
-        patchArrays.push_back( 
-            OsdDrawContext::PatchArray( desc, fpa.GetArrayRange() ) );
-
-        patchOffset = fpa.GetPatchIndex();
+    OsdDrawContext::ConvertPatchArrays(patchTables->GetPatchArrayVector(),
+        patchArrays, patchTables->GetMaxValence(), 0);
 
+/*    
 #if defined(GL_ES_VERSION_2_0)
+        // XXX: farmesh should have FarDensePatchTable for dense mesh indices.
+        //      instead of GetFaceVertices().
+        const FarSubdivisionTables<OsdVertex> *tables = farMesh->GetSubdivisionTables();
+        int level = tables->GetMaxLevel();
+        const std::vector<int> &indices = farMesh->GetFaceVertices(level-1);
+
+        int numIndices = (int)indices.size();
+
+        // Allocate and fill index buffer.
+        glGenBuffers(1, &patchIndexBuffer);
+        glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, patchIndexBuffer);
+        glBufferData(GL_ELEMENT_ARRAY_BUFFER,
+                     numIndices * sizeof(unsigned int), &(indices[0]), GL_STATIC_DRAW);
+
+
         // OpenGLES 2 supports only triangle topologies for filled
         // primitives i.e. not QUADS or PATCHES or LINES_ADJACENCY
         // For the convenience of clients build build a triangles
@@ -225,21 +201,44 @@ OsdGLDrawContext::create(FarPatchTables const * patchTables, bool requireFVarDat
         glBufferData(GL_ELEMENT_ARRAY_BUFFER,
                      numTrisIndices * sizeof(short), &(trisIndices[0]), GL_STATIC_DRAW);
 #endif
+*/    
+    
+    // allocate and initialize additional buffer data
+
+    // create vertex valence buffer and vertex texture
+    FarPatchTables::VertexValenceTable const &
+        valenceTable = patchTables->GetVertexValenceTable();
+
+    if (not valenceTable.empty()) {
+        vertexValenceTextureBuffer = createTextureBuffer(valenceTable, GL_R32I);
+
+        // also create vertex texture buffer (will be updated in UpdateVertexTexture())
+        glGenTextures(1, &vertexTextureBuffer);
     }
 
+
+    // create quad offset table buffer
+    FarPatchTables::QuadOffsetTable const &
+        quadOffsetTable = patchTables->GetQuadOffsetTable();
+
+    if (not quadOffsetTable.empty())
+        quadOffsetTextureBuffer = createTextureBuffer(quadOffsetTable, GL_R32I);
+
+
     // create ptex coordinate buffer
     FarPatchTables::PtexCoordinateTable const &
         ptexCoordTables = patchTables->GetPtexCoordinatesTable();
 
     if (not ptexCoordTables.empty())
-        ptexCoordinateTextureBuffer = createTextureBuffer(ptexCoordTables, GL_RG32I, patchOffset);
+        ptexCoordinateTextureBuffer = createTextureBuffer(ptexCoordTables, GL_RG32I);
+
 
     // create fvar data buffer if requested
     FarPatchTables::FVarDataTable const &
         fvarTables = patchTables->GetFVarDataTable();
 
     if (requireFVarData and not fvarTables.empty())
-        fvarDataTextureBuffer = createTextureBuffer(fvarTables, GL_R32F, patchOffset);
+        fvarDataTextureBuffer = createTextureBuffer(fvarTables, GL_R32F);
 
     glBindBuffer(GL_TEXTURE_BUFFER, 0);
 
diff --git a/opensubdiv/osd/glslComputeContext.cpp b/opensubdiv/osd/glslComputeContext.cpp
index 15f99c95..80f01326 100644
--- a/opensubdiv/osd/glslComputeContext.cpp
+++ b/opensubdiv/osd/glslComputeContext.cpp
@@ -238,18 +238,6 @@ OsdGLSLComputeContext::GetCurrentVaryingBuffer() const {
     return _currentVaryingBuffer;
 }
 
-int
-OsdGLSLComputeContext::GetNumCurrentVertexElements() const {
-
-    return _numVertexElements;
-}
-
-int
-OsdGLSLComputeContext::GetNumCurrentVaryingElements() const {
-    
-    return _numVaryingElements;
-}
-
 OsdGLSLComputeKernelBundle *
 OsdGLSLComputeContext::GetKernelBundle() const {
 
diff --git a/opensubdiv/osd/glslComputeContext.h b/opensubdiv/osd/glslComputeContext.h
index 5eca5c2d..d8b32727 100644
--- a/opensubdiv/osd/glslComputeContext.h
+++ b/opensubdiv/osd/glslComputeContext.h
@@ -77,6 +77,7 @@
 
 #include "../far/vertexEditTables.h"
 #include "../osd/vertex.h"
+#include "../osd/vertexDescriptor.h"
 #include "../osd/nonCopyable.h"
 
 #include <vector>
@@ -142,6 +143,7 @@ private:
 /// compute device.
 ///
 class OsdGLSLComputeContext {
+
 public:
     /// Creates an OsdGLSLComputeContext instance
     ///
@@ -166,8 +168,8 @@ public:
         _currentVertexBuffer = vertex ? vertex->BindVBO() : 0;
         _currentVaryingBuffer = varying ? varying->BindVBO() : 0;
 
-        _numVertexElements = vertex ? vertex->GetNumElements() : 0;
-        _numVaryingElements = varying ? varying->GetNumElements() : 0;
+        _vdesc.numVertexElements = vertex ? vertex->GetNumElements() : 0;
+        _vdesc.numVaryingElements = varying ? varying->GetNumElements() : 0;
 
         bindShaderStorageBuffers();
     }
@@ -201,9 +203,13 @@ public:
     /// Returns a handle to the varying-interpolated buffer
     GLuint GetCurrentVaryingBuffer() const;
 
-    int GetNumCurrentVertexElements() const;
-
-    int GetNumCurrentVaryingElements() const;
+    /// Returns an OsdVertexDescriptor if vertex buffers have been bound.
+    ///
+    /// @return a descriptor for the format of the vertex data currently bound
+    ///
+    OsdVertexDescriptor const & GetVertexDescriptor() const {
+        return _vdesc;
+    }
 
     OsdGLSLComputeKernelBundle * GetKernelBundle() const;
 
@@ -249,8 +255,7 @@ private:
     GLuint _vertexTexture,
            _varyingTexture;
 
-    int _numVertexElements,
-        _numVaryingElements;
+    OsdVertexDescriptor _vdesc;
 
     GLuint _currentVertexBuffer, 
            _currentVaryingBuffer;
diff --git a/opensubdiv/osd/glslKernelBundle.cpp b/opensubdiv/osd/glslKernelBundle.cpp
index 97ce288b..c2876df3 100644
--- a/opensubdiv/osd/glslKernelBundle.cpp
+++ b/opensubdiv/osd/glslKernelBundle.cpp
@@ -95,11 +95,9 @@ OsdGLSLComputeKernelBundle::~OsdGLSLComputeKernelBundle() {
 }
 
 bool
-OsdGLSLComputeKernelBundle::Compile(int numVertexElements,
-                                    int numVaryingElements) {
+OsdGLSLComputeKernelBundle::Compile(int numVertexElements, int numVaryingElements) {
 
-    _numVertexElements = numVertexElements;
-    _numVaryingElements = numVaryingElements;
+    _vdesc.Set(numVertexElements, numVaryingElements );
 
     if (_program) {
         glDeleteProgram(_program);
diff --git a/opensubdiv/osd/glslKernelBundle.h b/opensubdiv/osd/glslKernelBundle.h
index f0f051e2..9cf4836f 100644
--- a/opensubdiv/osd/glslKernelBundle.h
+++ b/opensubdiv/osd/glslKernelBundle.h
@@ -75,113 +75,116 @@
 #endif
 
 #include "../version.h"
+#include "../far/subdivisionTables.h"
 #include "../osd/nonCopyable.h"
 #include "../osd/vertex.h"
-#include "../far/subdivisionTables.h"
+#include "../osd/vertexDescriptor.h"
+
 
 namespace OpenSubdiv {
 namespace OPENSUBDIV_VERSION {
 
-    class OsdGLSLComputeKernelBundle
-        : OsdNonCopyable<OsdGLSLComputeKernelBundle> {
-    public:
-        OsdGLSLComputeKernelBundle();
-        ~OsdGLSLComputeKernelBundle();
+class OsdGLSLComputeKernelBundle : OsdNonCopyable<OsdGLSLComputeKernelBundle> {
+public:
+    OsdGLSLComputeKernelBundle();
+    ~OsdGLSLComputeKernelBundle();
 
-        bool Compile(int numVertexElements, int numVaryingElements);
+    bool Compile(int numVertexElements, int numVaryingElements);
 
-        void ApplyBilinearFaceVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyBilinearFaceVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyBilinearEdgeVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyBilinearEdgeVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyBilinearVertexVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyBilinearVertexVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkFaceVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyCatmarkFaceVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkEdgeVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyCatmarkEdgeVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkVertexVerticesKernelB(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyCatmarkVertexVerticesKernelB(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkVertexVerticesKernelA(
-            int vertexOffset, int tableOffset, int start, int end, bool pass);
+    void ApplyCatmarkVertexVerticesKernelA(
+        int vertexOffset, int tableOffset, int start, int end, bool pass);
 
-        void ApplyLoopEdgeVerticesKernel(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyLoopEdgeVerticesKernel(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyLoopVertexVerticesKernelB(
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyLoopVertexVerticesKernelB(
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyLoopVertexVerticesKernelA(
-            int vertexOffset, int tableOffset, int start, int end, bool pass);
+    void ApplyLoopVertexVerticesKernelA(
+        int vertexOffset, int tableOffset, int start, int end, bool pass);
 
-        void ApplyEditAdd(int primvarOffset, int primvarWidth,
-                          int vertexOffset, int tableOffset, int start, int end);
+    void ApplyEditAdd(int primvarOffset, int primvarWidth,
+                      int vertexOffset, int tableOffset, int start, int end);
 
-        void UseProgram() const;
+    void UseProgram() const;
 
-        GLuint GetTableUniformLocation(int tableIndex) const {
-            return _tableUniforms[tableIndex];
+    GLuint GetTableUniformLocation(int tableIndex) const {
+        return _tableUniforms[tableIndex];
+    }
+
+    struct Match {
+
+        /// Constructor
+        Match(int numVertexElements, int numVaryingElements)
+            : vdesc(numVertexElements, numVaryingElements) {
         }
 
-        struct Match {
-        Match(int numVertexElements, int numVaryingElements) :
-            _numVertexElements(numVertexElements),
-                _numVaryingElements(numVaryingElements) {}
-            bool operator() (const OsdGLSLComputeKernelBundle *kernel) {
-                return (kernel->_numVertexElements == _numVertexElements
-                        && kernel->_numVaryingElements == _numVaryingElements);
-            }
-            int _numVertexElements, _numVaryingElements;
-        };
+        bool operator() (OsdGLSLComputeKernelBundle const *kernel) {
+            return vdesc == kernel->_vdesc;
+        }
 
-        friend struct Match;
-
-    protected:
-        void dispatchCompute(int vertexOffset, int tableOffset, int start, int end) const;
-
-        GLuint _program;
-
-        // uniform locations for compute
-        GLuint _tableUniforms[FarSubdivisionTables<OsdVertex>::TABLE_TYPES_COUNT];
-        GLuint _uniformVertexPass;
-        GLuint _uniformVertexOffset;
-        GLuint _uniformTableOffset;
-        GLuint _uniformIndexStart;
-        GLuint _uniformIndexEnd;
-
-        // uniform locations for vertex edit
-        GLuint _uniformEditPrimVarOffset;
-        GLuint _uniformEditPrimVarWidth;
-
-        // general face-vertex kernel (all schemes)
-        GLuint _subComputeFace;
-        // edge-vertex kernel (catmark + loop schemes)
-        GLuint _subComputeEdge;
-        // edge-vertex kernel (bilinear scheme)
-        GLuint _subComputeBilinearEdge;
-        // vertex-vertex kernel (bilinear scheme)
-        GLuint _subComputeVertex;
-        // vertex-vertex kernel A (catmark + loop schemes)
-        GLuint _subComputeVertexA;
-        // vertex-vertex kernel B (catmark scheme)
-        GLuint _subComputeCatmarkVertexB;
-        // vertex-vertex kernel B (loop scheme)
-        GLuint _subComputeLoopVertexB;
-        // hedit kernel (add)
-        GLuint _subEditAdd;
-
-        int _numVertexElements,
-            _numVaryingElements;
-
-        int _workGroupSize;
+        OsdVertexDescriptor vdesc;
     };
 
+    friend struct Match;
+
+protected:
+    void dispatchCompute(int vertexOffset, int tableOffset, int start, int end) const;
+
+    GLuint _program;
+
+    // uniform locations for compute
+    GLuint _tableUniforms[FarSubdivisionTables<OsdVertex>::TABLE_TYPES_COUNT];
+    GLuint _uniformVertexPass;
+    GLuint _uniformVertexOffset;
+    GLuint _uniformTableOffset;
+    GLuint _uniformIndexStart;
+    GLuint _uniformIndexEnd;
+
+    // uniform locations for vertex edit
+    GLuint _uniformEditPrimVarOffset;
+    GLuint _uniformEditPrimVarWidth;
+
+    
+    GLuint _subComputeFace; // general face-vertex kernel (all schemes)
+
+    GLuint _subComputeEdge; // edge-vertex kernel (catmark + loop schemes)
+
+    GLuint _subComputeBilinearEdge; // edge-vertex kernel (bilinear scheme)
+
+    GLuint _subComputeVertex; // vertex-vertex kernel (bilinear scheme)
+
+    GLuint _subComputeVertexA; // vertex-vertex kernel A (catmark + loop schemes)
+
+    GLuint _subComputeCatmarkVertexB; // vertex-vertex kernel B (catmark scheme)
+
+    GLuint _subComputeLoopVertexB; // vertex-vertex kernel B (loop scheme)
+
+    GLuint _subEditAdd; // hedit kernel (add)
+
+    int _workGroupSize;
+
+    OsdVertexDescriptor _vdesc;
+};
+
 }  // end namespace OPENSUBDIV_VERSION
 using namespace OPENSUBDIV_VERSION;
 
diff --git a/opensubdiv/osd/glslTransformFeedbackComputeContext.cpp b/opensubdiv/osd/glslTransformFeedbackComputeContext.cpp
index b54f6a34..d21b049d 100644
--- a/opensubdiv/osd/glslTransformFeedbackComputeContext.cpp
+++ b/opensubdiv/osd/glslTransformFeedbackComputeContext.cpp
@@ -243,18 +243,6 @@ OsdGLSLTransformFeedbackComputeContext::GetCurrentVaryingBuffer() const {
     return _currentVaryingBuffer;
 }
 
-int
-OsdGLSLTransformFeedbackComputeContext::GetNumCurrentVertexElements() const {
-
-    return _numVertexElements;
-}
-
-int
-OsdGLSLTransformFeedbackComputeContext::GetNumCurrentVaryingElements() const {
-    
-    return _numVaryingElements;
-}
-
 OsdGLSLTransformFeedbackKernelBundle *
 OsdGLSLTransformFeedbackComputeContext::GetKernelBundle() const {
 
diff --git a/opensubdiv/osd/glslTransformFeedbackComputeContext.h b/opensubdiv/osd/glslTransformFeedbackComputeContext.h
index b638bbbf..600cc866 100644
--- a/opensubdiv/osd/glslTransformFeedbackComputeContext.h
+++ b/opensubdiv/osd/glslTransformFeedbackComputeContext.h
@@ -77,6 +77,7 @@
 
 #include "../far/vertexEditTables.h"
 #include "../osd/vertex.h"
+#include "../osd/vertexDescriptor.h"
 #include "../osd/nonCopyable.h"
 
 #include <vector>
@@ -164,8 +165,8 @@ public:
         _currentVertexBuffer = vertex ? vertex->BindVBO() : 0;
         _currentVaryingBuffer = varying ? varying->BindVBO() : 0;
 
-        _numVertexElements = vertex ? vertex->GetNumElements() : 0;
-        _numVaryingElements = varying ? varying->GetNumElements() : 0;
+        _vdesc.numVertexElements = vertex ? vertex->GetNumElements() : 0;
+        _vdesc.numVaryingElements = varying ? varying->GetNumElements() : 0;
 
         bindTextures();
     }
@@ -198,9 +199,13 @@ public:
     /// Returns a handle to the varying-interpolated buffer
     GLuint GetCurrentVaryingBuffer() const;
 
-    int GetNumCurrentVertexElements() const;
-
-    int GetNumCurrentVaryingElements() const;
+    /// Returns an OsdVertexDescriptor if vertex buffers have been bound.
+    ///
+    /// @return a descriptor for the format of the vertex data currently bound
+    ///
+    OsdVertexDescriptor const & GetVertexDescriptor() const {
+        return _vdesc;
+    }
 
     OsdGLSLTransformFeedbackKernelBundle * GetKernelBundle() const;
 
@@ -228,8 +233,7 @@ private:
     GLuint _vertexTexture,
            _varyingTexture;
 
-    int _numVertexElements,
-        _numVaryingElements;
+    OsdVertexDescriptor _vdesc;
 
     GLuint _currentVertexBuffer, 
            _currentVaryingBuffer;
diff --git a/opensubdiv/osd/glslTransformFeedbackComputeController.cpp b/opensubdiv/osd/glslTransformFeedbackComputeController.cpp
index 0ccf34a3..982c5a4f 100644
--- a/opensubdiv/osd/glslTransformFeedbackComputeController.cpp
+++ b/opensubdiv/osd/glslTransformFeedbackComputeController.cpp
@@ -119,9 +119,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyBilinearFaceVerticesKernel(
 
     kernelBundle->ApplyBilinearFaceVerticesKernel(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
 }
 
@@ -137,9 +137,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyBilinearEdgeVerticesKernel(
 
     kernelBundle->ApplyBilinearEdgeVerticesKernel(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
 }
 
@@ -155,9 +155,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyBilinearVertexVerticesKernel(
 
     kernelBundle->ApplyBilinearVertexVerticesKernel(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
 }
 
@@ -173,9 +173,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkFaceVerticesKernel(
 
     kernelBundle->ApplyCatmarkFaceVerticesKernel(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
 }
 
@@ -193,9 +193,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkEdgeVerticesKernel(
 
     kernelBundle->ApplyCatmarkEdgeVerticesKernel(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
 }
 
@@ -211,9 +211,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkVertexVerticesKernelB(
 
     kernelBundle->ApplyCatmarkVertexVerticesKernelB(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
 }
 
@@ -229,9 +229,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkVertexVerticesKernelA1(
 
     kernelBundle->ApplyCatmarkVertexVerticesKernelA(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
 }
 
@@ -247,9 +247,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkVertexVerticesKernelA2(
 
     kernelBundle->ApplyCatmarkVertexVerticesKernelA(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
 }
 
@@ -265,9 +265,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyLoopEdgeVerticesKernel(
 
     kernelBundle->ApplyLoopEdgeVerticesKernel(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
 }
 
@@ -283,9 +283,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyLoopVertexVerticesKernelB(
 
     kernelBundle->ApplyLoopVertexVerticesKernelB(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
 }
 
@@ -301,9 +301,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyLoopVertexVerticesKernelA1(
 
     kernelBundle->ApplyLoopVertexVerticesKernelA(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
 }
 
@@ -319,9 +319,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyLoopVertexVerticesKernelA2(
 
     kernelBundle->ApplyLoopVertexVerticesKernelA(
         context->GetCurrentVertexBuffer(),
-        context->GetNumCurrentVertexElements(),
+        context->GetVertexDescriptor().numVertexElements,
         context->GetCurrentVaryingBuffer(),
-        context->GetNumCurrentVaryingElements(),
+        context->GetVertexDescriptor().numVaryingElements,
         batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
 }
 
@@ -346,9 +346,9 @@ OsdGLSLTransformFeedbackComputeController::ApplyVertexEdits(
     if (edit->GetOperation() == FarVertexEdit::Add) {
         kernelBundle->ApplyEditAdd(
             context->GetCurrentVertexBuffer(),
-            context->GetNumCurrentVertexElements(),
+            context->GetVertexDescriptor().numVertexElements,
             context->GetCurrentVaryingBuffer(),
-            context->GetNumCurrentVaryingElements(),
+            context->GetVertexDescriptor().numVaryingElements,
             primvarOffset, primvarWidth,
             batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
     } else {
diff --git a/opensubdiv/osd/glslTransformFeedbackKernelBundle.cpp b/opensubdiv/osd/glslTransformFeedbackKernelBundle.cpp
index fb9d2cee..dcd662f8 100644
--- a/opensubdiv/osd/glslTransformFeedbackKernelBundle.cpp
+++ b/opensubdiv/osd/glslTransformFeedbackKernelBundle.cpp
@@ -108,8 +108,8 @@ OsdGLSLTransformFeedbackKernelBundle::Compile(int numVertexElements, int numVary
 
     assert(numVertexElements >= 3); // at least xyz required (for performance reason)
 
-    _numVertexElements = numVertexElements;
-    _numVaryingElements = numVaryingElements;
+    _vdesc.Set(numVertexElements, numVaryingElements);
+    
     _program = glCreateProgram();
 
     GLuint shader = glCreateShader(GL_VERTEX_SHADER);
diff --git a/opensubdiv/osd/glslTransformFeedbackKernelBundle.h b/opensubdiv/osd/glslTransformFeedbackKernelBundle.h
index 6f8dbb77..dafa989f 100644
--- a/opensubdiv/osd/glslTransformFeedbackKernelBundle.h
+++ b/opensubdiv/osd/glslTransformFeedbackKernelBundle.h
@@ -76,156 +76,160 @@
 #include "../version.h"
 #include "../osd/nonCopyable.h"
 #include "../osd/vertex.h"
+#include "../osd/vertexDescriptor.h"
 #include "../far/subdivisionTables.h"
 
 namespace OpenSubdiv {
 namespace OPENSUBDIV_VERSION {
 
-    class OsdGLSLTransformFeedbackKernelBundle
-        : OsdNonCopyable<OsdGLSLTransformFeedbackKernelBundle> {
-    public:
-        OsdGLSLTransformFeedbackKernelBundle();
-        ~OsdGLSLTransformFeedbackKernelBundle();
+class OsdGLSLTransformFeedbackKernelBundle : OsdNonCopyable<OsdGLSLTransformFeedbackKernelBundle> {
+public:
+    /// Constructor
+    OsdGLSLTransformFeedbackKernelBundle();
+    
+    ~OsdGLSLTransformFeedbackKernelBundle();
 
-        bool Compile(int numVertexElements, int numVaryingElements);
+    bool Compile(int numVertexElements, int numVaryingElements);
 
-        void ApplyBilinearFaceVerticesKernel(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyBilinearFaceVerticesKernel(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyBilinearEdgeVerticesKernel(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyBilinearEdgeVerticesKernel(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyBilinearVertexVerticesKernel(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyBilinearVertexVerticesKernel(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkFaceVerticesKernel(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyCatmarkFaceVerticesKernel(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkEdgeVerticesKernel(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyCatmarkEdgeVerticesKernel(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkVertexVerticesKernelB(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyCatmarkVertexVerticesKernelB(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyCatmarkVertexVerticesKernelA(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end, bool pass);
+    void ApplyCatmarkVertexVerticesKernelA(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end, bool pass);
 
-        void ApplyLoopEdgeVerticesKernel(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyLoopEdgeVerticesKernel(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyLoopVertexVerticesKernelB(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyLoopVertexVerticesKernelB(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void ApplyLoopVertexVerticesKernelA(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end, bool pass);
+    void ApplyLoopVertexVerticesKernelA(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end, bool pass);
 
-        void ApplyEditAdd(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int primvarOffset, int primvarWidth,
-            int vertexOffset, int tableOffset, int start, int end);
+    void ApplyEditAdd(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int primvarOffset, int primvarWidth,
+        int vertexOffset, int tableOffset, int start, int end);
 
-        void UseProgram() const;
+    void UseProgram() const;
 
-        GLuint GetTableUniformLocation(int tableIndex) const {
-            return _uniformTables[tableIndex];
-        }
-        GLuint GetVertexUniformLocation() const {
-            return _uniformVertexBuffer;
-        }
-        GLuint GetVaryingUniformLocation() const {
-            return _uniformVaryingBuffer;
-        }
-        GLuint GetEditIndicesUniformLocation() const {
-            return _uniformEditIndices;
-        }
-        GLuint GetEditValuesUniformLocation() const {
-            return _uniformEditValues;
-        }
-        GLuint GetVertexBufferImageUniformLocation() const {
-            return _uniformVertexBufferImage;
+    GLuint GetTableUniformLocation(int tableIndex) const {
+        return _uniformTables[tableIndex];
+    }
+    GLuint GetVertexUniformLocation() const {
+        return _uniformVertexBuffer;
+    }
+    GLuint GetVaryingUniformLocation() const {
+        return _uniformVaryingBuffer;
+    }
+    GLuint GetEditIndicesUniformLocation() const {
+        return _uniformEditIndices;
+    }
+    GLuint GetEditValuesUniformLocation() const {
+        return _uniformEditValues;
+    }
+    GLuint GetVertexBufferImageUniformLocation() const {
+        return _uniformVertexBufferImage;
+    }
+
+    struct Match {
+
+        /// Constructor
+        Match(int numVertexElements, int numVaryingElements)
+            : vdesc(numVertexElements, numVaryingElements) {
         }
 
-        struct Match {
-        Match(int numVertexElements, int numVaryingElements) :
-            _numVertexElements(numVertexElements),
-                _numVaryingElements(numVaryingElements) {}
-            bool operator() (const OsdGLSLTransformFeedbackKernelBundle *kernel) {
-                return (kernel->_numVertexElements == _numVertexElements
-                        && kernel->_numVaryingElements == _numVaryingElements);
-            }
-            int _numVertexElements, _numVaryingElements;
-        };
+        bool operator() (OsdGLSLTransformFeedbackKernelBundle const *kernel) {
+            return vdesc == kernel->_vdesc;
+        }
 
-        friend struct Match;
-
-    protected:
-        void transformGpuBufferData(
-            GLuint vertexBuffer, int numVertexElements,
-            GLuint varyingBuffer, int numVaryingElements,
-            int vertexOffset, int tableOffset, int start, int end) const;
-
-        GLuint _program;
-
-        // uniform locations
-        GLuint _uniformTables[FarSubdivisionTables<OsdVertex>::TABLE_TYPES_COUNT];
-        GLuint _uniformVertexPass;
-        GLuint _uniformVertexOffset;
-        GLuint _uniformTableOffset;
-        GLuint _uniformIndexStart;
-
-        GLuint _uniformVertexBuffer;
-        GLuint _uniformVaryingBuffer;
-
-        GLuint _uniformEditPrimVarOffset;
-        GLuint _uniformEditPrimVarWidth;
-
-        GLuint _uniformEditIndices;
-        GLuint _uniformEditValues;
-        GLuint _uniformVertexBufferImage;
-
-        // subroutines
-        // general face-vertex kernel (all schemes)
-        GLuint _subComputeFace;
-        // edge-vertex kernel (catmark + loop schemes)
-        GLuint _subComputeEdge;
-        // edge-vertex kernel (bilinear scheme)
-        GLuint _subComputeBilinearEdge;
-        // vertex-vertex kernel (bilinear scheme)
-        GLuint _subComputeVertex;
-        // vertex-vertex kernel A (catmark + loop schemes)
-        GLuint _subComputeVertexA;
-        // vertex-vertex kernel B (catmark scheme)
-        GLuint _subComputeCatmarkVertexB;
-        // vertex-vertex kernel B (loop scheme)
-        GLuint _subComputeLoopVertexB;
-        // hedit kernel (add)
-        GLuint _subEditAdd;
-
-        int _numVertexElements,
-            _numVaryingElements;
+        OsdVertexDescriptor vdesc;
     };
 
+    friend struct Match;
+
+protected:
+    void transformGpuBufferData(
+        GLuint vertexBuffer, int numVertexElements,
+        GLuint varyingBuffer, int numVaryingElements,
+        int vertexOffset, int tableOffset, int start, int end) const;
+
+    GLuint _program;
+
+    // uniform locations
+    GLuint _uniformTables[FarSubdivisionTables<OsdVertex>::TABLE_TYPES_COUNT];
+    GLuint _uniformVertexPass;
+    GLuint _uniformVertexOffset;
+    GLuint _uniformTableOffset;
+    GLuint _uniformIndexStart;
+
+    GLuint _uniformVertexBuffer;
+    GLuint _uniformVaryingBuffer;
+
+    GLuint _uniformEditPrimVarOffset;
+    GLuint _uniformEditPrimVarWidth;
+
+    GLuint _uniformEditIndices;
+    GLuint _uniformEditValues;
+    GLuint _uniformVertexBufferImage;
+
+    // subroutines
+
+    GLuint _subComputeFace; // general face-vertex kernel (all schemes)
+
+    GLuint _subComputeEdge; // edge-vertex kernel (catmark + loop schemes)
+
+    GLuint _subComputeBilinearEdge; // edge-vertex kernel (bilinear scheme)
+
+    GLuint _subComputeVertex; // vertex-vertex kernel (bilinear scheme)
+
+    GLuint _subComputeVertexA; // vertex-vertex kernel A (catmark + loop schemes)
+
+    GLuint _subComputeCatmarkVertexB;// vertex-vertex kernel B (catmark scheme)
+
+    GLuint _subComputeLoopVertexB; // vertex-vertex kernel B (loop scheme)
+
+    GLuint _subEditAdd; // hedit kernel (add)
+
+    OsdVertexDescriptor _vdesc;
+};
+
 }  // end namespace OPENSUBDIV_VERSION
 using namespace OPENSUBDIV_VERSION;
 
diff --git a/opensubdiv/osd/ompKernel.cpp b/opensubdiv/osd/ompKernel.cpp
index b146283b..785a0653 100644
--- a/opensubdiv/osd/ompKernel.cpp
+++ b/opensubdiv/osd/ompKernel.cpp
@@ -65,7 +65,7 @@ namespace OpenSubdiv {
 namespace OPENSUBDIV_VERSION {
 
 void OsdOmpComputeFace(
-    const OsdVertexDescriptor *vdesc, float * vertex, float * varying,
+    OsdVertexDescriptor const &vdesc, float * vertex, float * varying,
     const int *F_IT, const int *F_ITa, int offset, int tableOffset, int start, int end) {
 
 #pragma omp parallel for
@@ -78,18 +78,18 @@ void OsdOmpComputeFace(
         // XXX: should use local vertex struct variable instead of
         // accumulating directly into global memory.
         int dstIndex = offset + i - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
         for (int j = 0; j < n; ++j) {
             int index = F_IT[h+j];
-            vdesc->AddWithWeight(vertex, dstIndex, index, weight);
-            vdesc->AddVaryingWithWeight(varying, dstIndex, index, weight);
+            vdesc.AddWithWeight(vertex, dstIndex, index, weight);
+            vdesc.AddVaryingWithWeight(varying, dstIndex, index, weight);
         }
     }
 }
 
 void OsdOmpComputeEdge(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *E_IT, const float *E_W, int offset, int tableOffset, int start, int end) {
 
 #pragma omp parallel for
@@ -102,25 +102,25 @@ void OsdOmpComputeEdge(
         float vertWeight = E_W[i*2+0];
 
         int dstIndex = offset + i - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
-        vdesc->AddWithWeight(vertex, dstIndex, eidx0, vertWeight);
-        vdesc->AddWithWeight(vertex, dstIndex, eidx1, vertWeight);
+        vdesc.AddWithWeight(vertex, dstIndex, eidx0, vertWeight);
+        vdesc.AddWithWeight(vertex, dstIndex, eidx1, vertWeight);
 
         if (eidx2 != -1) {
             float faceWeight = E_W[i*2+1];
 
-            vdesc->AddWithWeight(vertex, dstIndex, eidx2, faceWeight);
-            vdesc->AddWithWeight(vertex, dstIndex, eidx3, faceWeight);
+            vdesc.AddWithWeight(vertex, dstIndex, eidx2, faceWeight);
+            vdesc.AddWithWeight(vertex, dstIndex, eidx3, faceWeight);
         }
 
-        vdesc->AddVaryingWithWeight(varying, dstIndex, eidx0, 0.5f);
-        vdesc->AddVaryingWithWeight(varying, dstIndex, eidx1, 0.5f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, eidx0, 0.5f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, eidx1, 0.5f);
     }
 }
 
 void OsdOmpComputeVertexA(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *V_ITa, const float *V_W,
     int offset, int tableOffset, int start, int end, int pass) {
 
@@ -141,23 +141,23 @@ void OsdOmpComputeVertexA(
 
         int dstIndex = offset + i - tableOffset;
         if (not pass)
-            vdesc->Clear(vertex, varying, dstIndex);
+            vdesc.Clear(vertex, varying, dstIndex);
 
         if (eidx0 == -1 || (pass == 0 && (n == -1))) {
-            vdesc->AddWithWeight(vertex, dstIndex, p, weight);
+            vdesc.AddWithWeight(vertex, dstIndex, p, weight);
         } else {
-            vdesc->AddWithWeight(vertex, dstIndex, p, weight * 0.75f);
-            vdesc->AddWithWeight(vertex, dstIndex, eidx0, weight * 0.125f);
-            vdesc->AddWithWeight(vertex, dstIndex, eidx1, weight * 0.125f);
+            vdesc.AddWithWeight(vertex, dstIndex, p, weight * 0.75f);
+            vdesc.AddWithWeight(vertex, dstIndex, eidx0, weight * 0.125f);
+            vdesc.AddWithWeight(vertex, dstIndex, eidx1, weight * 0.125f);
         }
 
         if (not pass)
-            vdesc->AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
+            vdesc.AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
     }
 }
 
 void OsdOmpComputeVertexB(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *V_ITa, const int *V_IT, const float *V_W,
     int offset, int tableOffset, int start, int end) {
 
@@ -172,20 +172,20 @@ void OsdOmpComputeVertexB(
         float wv = (n-2.0f) * n * wp;
 
         int dstIndex = offset + i - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
-        vdesc->AddWithWeight(vertex, dstIndex, p, weight * wv);
+        vdesc.AddWithWeight(vertex, dstIndex, p, weight * wv);
 
         for (int j = 0; j < n; ++j) {
-            vdesc->AddWithWeight(vertex, dstIndex, V_IT[h+j*2], weight * wp);
-            vdesc->AddWithWeight(vertex, dstIndex, V_IT[h+j*2+1], weight * wp);
+            vdesc.AddWithWeight(vertex, dstIndex, V_IT[h+j*2], weight * wp);
+            vdesc.AddWithWeight(vertex, dstIndex, V_IT[h+j*2+1], weight * wp);
         }
-        vdesc->AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
     }
 }
 
 void OsdOmpComputeLoopVertexB(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *V_ITa, const int *V_IT, const float *V_W,
     int vertexOffset, int tableOffset, int start, int end) {
 
@@ -202,19 +202,19 @@ void OsdOmpComputeLoopVertexB(
         beta = (0.625f - beta) * wp;
 
         int dstIndex = i + vertexOffset - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
-        vdesc->AddWithWeight(vertex, dstIndex, p, weight * (1.0f - (beta * n)));
+        vdesc.AddWithWeight(vertex, dstIndex, p, weight * (1.0f - (beta * n)));
 
         for (int j = 0; j < n; ++j)
-            vdesc->AddWithWeight(vertex, dstIndex, V_IT[h+j], weight * beta);
+            vdesc.AddWithWeight(vertex, dstIndex, V_IT[h+j], weight * beta);
 
-        vdesc->AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
     }
 }
 
 void OsdOmpComputeBilinearEdge(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *E_IT, int vertexOffset, int tableOffset, int start, int end) {
 
 #pragma omp parallel for
@@ -223,18 +223,18 @@ void OsdOmpComputeBilinearEdge(
         int eidx1 = E_IT[2*i+1];
 
         int dstIndex = i + vertexOffset - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
-        vdesc->AddWithWeight(vertex, dstIndex, eidx0, 0.5f);
-        vdesc->AddWithWeight(vertex, dstIndex, eidx1, 0.5f);
+        vdesc.AddWithWeight(vertex, dstIndex, eidx0, 0.5f);
+        vdesc.AddWithWeight(vertex, dstIndex, eidx1, 0.5f);
 
-        vdesc->AddVaryingWithWeight(varying, dstIndex, eidx0, 0.5f);
-        vdesc->AddVaryingWithWeight(varying, dstIndex, eidx1, 0.5f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, eidx0, 0.5f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, eidx1, 0.5f);
     }
 }
 
 void OsdOmpComputeBilinearVertex(
-    const OsdVertexDescriptor *vdesc, float *vertex, float *varying,
+    OsdVertexDescriptor const &vdesc, float *vertex, float *varying,
     const int *V_ITa, int vertexOffset, int tableOffset, int start, int end) {
 
 #pragma omp parallel for
@@ -242,42 +242,42 @@ void OsdOmpComputeBilinearVertex(
         int p = V_ITa[i];
 
         int dstIndex = i + vertexOffset - tableOffset;
-        vdesc->Clear(vertex, varying, dstIndex);
+        vdesc.Clear(vertex, varying, dstIndex);
 
-        vdesc->AddWithWeight(vertex, dstIndex, p, 1.0f);
-        vdesc->AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
+        vdesc.AddWithWeight(vertex, dstIndex, p, 1.0f);
+        vdesc.AddVaryingWithWeight(varying, dstIndex, p, 1.0f);
     }
 }
 
 void OsdOmpEditVertexAdd(
-    const OsdVertexDescriptor *vdesc, float *vertex,
+    OsdVertexDescriptor const &vdesc, float *vertex,
     int primVarOffset, int primVarWidth, int vertexOffset, int tableOffset,
     int start, int end,
     const unsigned int *editIndices, const float *editValues) {
 
 #pragma omp parallel for
     for (int i = start+tableOffset; i < end+tableOffset; i++) {
-        vdesc->ApplyVertexEditAdd(vertex,
-                                  primVarOffset,
-                                  primVarWidth,
-                                  editIndices[i] + vertexOffset,
-                                  &editValues[i*primVarWidth]);
+        vdesc.ApplyVertexEditAdd(vertex,
+                                 primVarOffset,
+                                 primVarWidth,
+                                 editIndices[i] + vertexOffset,
+                                 &editValues[i*primVarWidth]);
     }
 }
 
 void OsdOmpEditVertexSet(
-    const OsdVertexDescriptor *vdesc, float *vertex,
+    OsdVertexDescriptor const &vdesc, float *vertex,
     int primVarOffset, int primVarWidth, int vertexOffset, int tableOffset,
     int start, int end,
     const unsigned int *editIndices, const float *editValues) {
 
 #pragma omp parallel for
     for (int i = start+tableOffset; i < end+tableOffset; i++) {
-        vdesc->ApplyVertexEditSet(vertex,
-                                  primVarOffset,
-                                  primVarWidth,
-                                  editIndices[i] + vertexOffset,
-                                  &editValues[i*primVarWidth]);
+        vdesc.ApplyVertexEditSet(vertex,
+                                 primVarOffset,
+                                 primVarWidth,
+                                 editIndices[i] + vertexOffset,
+                                 &editValues[i*primVarWidth]);
     }
 }
 
diff --git a/opensubdiv/osd/ompKernel.h b/opensubdiv/osd/ompKernel.h
index dde778f5..b8fe25c3 100644
--- a/opensubdiv/osd/ompKernel.h
+++ b/opensubdiv/osd/ompKernel.h
@@ -64,57 +64,57 @@ namespace OPENSUBDIV_VERSION {
 
 struct OsdVertexDescriptor;
 
-void OsdOmpComputeFace(const OsdVertexDescriptor *vdesc,
+void OsdOmpComputeFace(OsdVertexDescriptor const &vdesc,
                        float * vertex, float * varying,
                        const int *F_IT, const int *F_ITa,
                        int vertexOffset, int tableOffset,
                        int start, int end);
 
-void OsdOmpComputeEdge(const OsdVertexDescriptor *vdesc,
+void OsdOmpComputeEdge(OsdVertexDescriptor const &vdesc,
                        float *vertex, float * varying,
                        const int *E_IT, const float *E_ITa,
                        int vertexOffset, int tableOffset,
                        int start, int end);
 
-void OsdOmpComputeVertexA(const OsdVertexDescriptor *vdesc,
+void OsdOmpComputeVertexA(OsdVertexDescriptor const &vdesc,
                           float *vertex, float * varying,
                           const int *V_ITa, const float *V_IT,
                           int vertexOffset, int tableOffset,
                           int start, int end, int pass);
 
-void OsdOmpComputeVertexB(const OsdVertexDescriptor *vdesc,
+void OsdOmpComputeVertexB(OsdVertexDescriptor const &vdesc,
                           float *vertex, float * varying,
                           const int *V_ITa, const int *V_IT, const float *V_W,
                           int vertexOffset, int tableOffset,
                           int start, int end);
 
-void OsdOmpComputeLoopVertexB(const OsdVertexDescriptor *vdesc,
+void OsdOmpComputeLoopVertexB(OsdVertexDescriptor const &vdesc,
                               float *vertex, float * varying,
                               const int *V_ITa, const int *V_IT,
                               const float *V_W,
                               int vertexOffset, int tableOffset,
                               int start, int end);
 
-void OsdOmpComputeBilinearEdge(const OsdVertexDescriptor *vdesc,
+void OsdOmpComputeBilinearEdge(OsdVertexDescriptor const &vdesc,
                                float *vertex, float * varying,
                                const int *E_IT,
                                int vertexOffset, int tableOffset,
                                int start, int end);
 
-void OsdOmpComputeBilinearVertex(const OsdVertexDescriptor *vdesc,
+void OsdOmpComputeBilinearVertex(OsdVertexDescriptor const &vdesc,
                                  float *vertex, float * varying,
                                  const int *V_ITa,
                                  int vertexOffset, int tableOffset,
                                  int start, int end);
 
-void OsdOmpEditVertexAdd(const OsdVertexDescriptor *vdesc, float *vertex,
+void OsdOmpEditVertexAdd(OsdVertexDescriptor const &vdesc, float *vertex,
                          int primVarOffset, int primVarWidth,
                          int vertexOffset, int tableOffset,
                          int start, int end,
                          const unsigned int *editIndices,
                          const float *editValues);
 
-void OsdOmpEditVertexSet(const OsdVertexDescriptor *vdesc, float *vertex,
+void OsdOmpEditVertexSet(OsdVertexDescriptor const &vdesc, float *vertex,
                          int primVarOffset, int primVarWidth,
                          int vertexOffset, int tableOffset,
                          int start, int end,
diff --git a/opensubdiv/osd/vertexDescriptor.h b/opensubdiv/osd/vertexDescriptor.h
index b9684285..72e306cc 100644
--- a/opensubdiv/osd/vertexDescriptor.h
+++ b/opensubdiv/osd/vertexDescriptor.h
@@ -64,10 +64,40 @@ namespace OPENSUBDIV_VERSION {
 
 struct OsdVertexDescriptor {
 
+    /// Constructor
+    OsdVertexDescriptor() : numVertexElements(0), numVaryingElements(0) {}
+
+    /// Constructor
+    ///
+    /// @param numVertexElem   number of vertex-interpolated data elements (floats)
+    ///
+    /// @param numVaryingElem  number of varying-interpolated data elements (floats)
+    ///
     OsdVertexDescriptor(int numVertexElem, int numVaryingElem)
         : numVertexElements(numVertexElem),
         numVaryingElements(numVaryingElem) { }
 
+    /// Sets descriptor
+    ///
+    /// @param numVertexElem   number of vertex-interpolated data elements (floats)
+    ///
+    /// @param numVaryingElem  number of varying-interpolated data elements (floats)
+    ///
+    void Set(int numVertexElem, int numVaryingElem) {
+        numVertexElements = numVertexElem;
+        numVaryingElements = numVaryingElem;
+    }
+    
+    /// Resets the descriptor
+    void Reset() {
+        numVertexElements = numVaryingElements = 0;
+    }
+
+    bool operator == (OsdVertexDescriptor const & other) {
+        return (numVertexElements == other.numVertexElements and
+                numVaryingElements == other.numVaryingElements);
+    }
+
     /// Resets the contents of vertex & varying primvar data buffers for a given
     /// vertex.
     ///
@@ -90,7 +120,7 @@ struct OsdVertexDescriptor {
                 varying[index*numVaryingElements+i] = 0.0f;
         }
     }
-
+    
     /// Applies "dst += src*weight" to "vertex" primvar data in a vertex buffer.
     ///
     /// @param vertex The VertexData buffer