Merge branch 'release/v2_2_0'

CMakeLists.txt

@@ -131,6 +131,12 @@ if(CMAKE_COMPILER_IS_GNUCC OR CMAKE_COMPILER_IS_CLANGCC OR CMAKE_COMPILER_IS_ICC
 
         foreach (ICC_LIB iomp5 irng intlc)
 
+            if(CMAKE_SIZEOF_VOID_P MATCHES "8")
+                list(APPEND ICC_LIB_ARCH "intel64")
+            elseif(CMAKE_SIZEOF_VOID_P MATCHES "4")
+                list(APPEND ICC_LIB_ARCH "ia32")
+            endif()
+
             find_library( ICC_${ICC_LIB}
                 NAMES
                     ${ICC_LIB}
@@ -138,9 +144,7 @@ if(CMAKE_COMPILER_IS_GNUCC OR CMAKE_COMPILER_IS_CLANGCC OR CMAKE_COMPILER_IS_ICC
                     ${ICC_LOCATION}/lib
                     /opt/intel/lib/
                 PATH_SUFFIXES
-                    intel64/
-                    ia32/
-                    mic/
+                    ${ICC_LIB_ARCH}
             )
 
             if (ICC_${ICC_LIB})

documentation/CMakeLists.txt

@@ -71,6 +71,7 @@ if (DOCUTILS_FOUND AND PYTHONINTERP_FOUND)
         far_overview.rst
         getting_started.rst
         glviewer.rst
+        glstencilviewer.rst
         glbatchviewer.rst
         hbr_overview.rst
         intro.rst 

documentation/code_examples.rst

@@ -40,9 +40,10 @@ of the software.
 
    * - | `glViewer <glviewer.html>`_
        | `glBatchViewer <glbatchviewer.html>`_
+       | `glStencilViewer <glstencilviewer.html>`_
+     - | `limitEval <limiteval.html>`_
+       | `paintTest <painttest.html>`_
        | `ptexViewer <ptexviewer.html>`_
-     - | `paintTest <painttest.html>`_
-       | `limitEval <limiteval.html>`_
 
 .. list-table:: **DirectX examples**
    :class: quickref

documentation/dxviewer.rst

@@ -65,6 +65,7 @@ SEE ALSO
 `Code Examples <code_examples.html>`__, \
 `glViewer <glviewer.html>`__, \
 `glBatchViewer <glbatchviewer.html>`__, \
+`glStencilViewer <glstencilviewer.html>`__, \
 `ptexViewer <ptexviewer.html>`__, \
 `paintTest <painttest.html>`__, \
 `limitEval <limiteval.html>`__, \

documentation/far_overview.rst

@@ -50,3 +50,146 @@ Subdivision Tables
 
 Patch Tables
 ============
+
+
+Stencil Tables
+==============
+
+
+Stencils are the most direct method of evaluation of specific locations on the
+limit of a subdivision surface starting from the coarse vertices of the control
+cage.
+
+.. image:: images/far_stencil0.png
+   :align: center
+
+Sample Location
+***************
+
+Each stencil is associated with a singular parametric location on the coarse
+mesh. The paramatric location is defined as face location and local [0.0 - 1.0]
+(u,v) triplet:
+
+In the case of a non-coarse quad face, the parametric sub-face quadrant needs to
+be identified. This can be done either explicitly or implicitly by using the
+unique ptex face indices for instance.
+
+.. image:: images/far_stencil6.png
+   :align: center
+
+
+Principles
+**********
+
+Iterative subdivision algorithms such as the one used in `FarSubdivisionTables <#subdivision-tables>`__
+converge towards the limit surface by sucessively refining the vertices of the
+coarse control cage.
+
+.. image:: images/far_stencil4.png
+   :align: center
+
+Each step is dependent upon the previous subidivion step being completed, and a
+substantial number of steps may be required in order approximate the limit. Since
+each subdivision step incurs an O(4 :superscript:`n`) growing amount of
+computations, the accrued number of interpolations can be quite large.
+
+However, every intermediate subdivided vertex can be expressed as a linear
+interpolation of vertice from the previous step. So, eventually, every point at
+on the limit surface can be expressed as a weighted average of the set of coarse 
+control vertices from the one-ring surrounding the face that the point is in:
+
+.. image:: images/far_stencil3.png
+   :align: center
+   
+Where:
+
+.. image:: images/far_stencil2.png
+   :align: center
+
+Stencils are created by combining the list of control vertices of the 1-ring
+to a set of interpolation weights obtained by successive accumulation of
+subdivision interpolation weights.
+
+The weight accumulation process is made efficient by adaptively subdividing the
+control cage only around extraordinary locations, and otherwise reverting to fast
+bi-cubic bspline patch evaluation. The use of bi-cubic patches also allows the
+accumulation of analytical derivatives.
+
+API Architecture
+****************
+
+The base container for stencil data is the FarStencilTables class. As with most
+other Far entities, it has an associated FarStencilTablesFactory that requires
+an HbrMesh:
+
+.. image:: images/far_stencil5.png
+   :align: center
+
+Assuming a properly qualified HbrMesh:
+
+.. code:: c++
+
+    HMesh<OpenSubdiv::FarStencilFactoryVertex> * mesh;
+
+    FarStencilTables controlStencils;
+
+    OpenSubdiv::FarStencilTablesFactory<> factory(mesh);
+    
+    for (int i=0; i<nfaces; ++i) {
+
+        HFace * f = mesh->GetFace(i);
+
+        int nv = f->GetNumVertices();
+
+        if (nv!=4) {
+
+            // if the face is not a quad, we have to iterate over sub-quad(rants)
+            for (int j=0; j<f->GetNumVertices(); ++j) {
+
+                factory.SetCurrentFace(i,j);
+
+                factory.AppendStencils( &controlStencils, nsamples/nv, u, v, reflevel );
+            }
+        } else {
+
+            factory.SetCurrentFace(i);
+
+            factory.AppendStencils( &controlStencils, g_nsamples, u, v, reflevel );
+        }
+    }
+
+When the control vertices (controlPoints) move in space, the limit locations can 
+be very efficiently recomputed simply by applying the blending weights to the 
+series of coarse control vertices:
+
+.. code:: c++
+
+    class StencilType {
+    public:
+
+        void Clear() {
+            memset( &x, 0, sizeof(StencilType));
+        }
+
+        void AddWithWeight( StencilType const & cv, float weight  ) {
+            x += cv.x * weight;
+            y += cv.y * weight;
+            z += cv.z * weight;
+        }
+
+        float x,y,z;
+    };
+
+    std::vector<StencilType> controlPoints,
+                             points,
+                             utan,
+                             vtan;
+    
+    // Uppdate points by applying stencils
+    controlStencils.UpdateValues<StencilType>( reinterpret_cast<StencilType const *>(
+        &controlPoints[0]), &points[0] );
+
+    // Uppdate tangents by applying derivative stencils
+    controlStencils.UpdateDerivs<StencilType>( reinterpret_cast<StencilType const *>(
+        &controlPoints[0]), &utan[0], &vtan[0] );
+

documentation/glbatchviewer.rst

@@ -81,7 +81,7 @@ SEE ALSO
 
 `Code Examples <code_examples.html>`__, \
 `glViewer <glviewer.html>`__, \
-`glBatchViewer <glbatchviewer.html>`__, \
+`glStencilViewer <glstencilviewer.html>`__, \
 `ptexViewer <ptexviewer.html>`__, \
 `paintTest <painttest.html>`__, \
 `limitEval <limiteval.html>`__, \

documentation/glstencilviewer.rst

@@ -0,0 +1,74 @@
+..  
+       Copyright 2013 Pixar
+
+       Licensed under the Apache License, Version 2.0 (the "License");
+       you may not use this file except in compliance with the License
+       and the following modification to it: Section 6 Trademarks.
+       deleted and replaced with:
+
+       6. Trademarks. This License does not grant permission to use the
+       trade names, trademarks, service marks, or product names of the
+       Licensor and its affiliates, except as required for reproducing
+       the content of the NOTICE file.
+
+       You may obtain a copy of the License at
+
+       http://www.apache.org/licenses/LICENSE-2.0
+
+       Unless required by applicable law or agreed to in writing,
+       software distributed under the License is distributed on an
+       "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+       either express or implied.  See the License for the specific
+       language governing permissions and limitations under the
+       License.
+  
+
+glStencilViewer
+---------------
+
+.. contents::
+   :local:
+   :backlinks: none
+
+SYNOPSIS
+========
+
+.. parsed-literal:: 
+   :class: codefhead
+
+   **glStencilViewer** [**-d** *isolation level*] [**-f**] *objfile(s)*
+
+DESCRIPTION
+===========
+
+``glStencilViewer`` is a stand-alone application that showcases the application of 
+pre-computed stencil tables to a collection of geometric test shapes. Multiple 
+controls are available to experiment with the algorithms.
+
+.. image:: images/glstencilviewer.png 
+   :width: 400px
+   :align: center
+   :target: images/glstencilviewer.png 
+
+OPTIONS
+=======
+
+**-d** *isolation level*
+  Select the desired isolation level of adaptive feature isolation. This can be 
+  useful when trying to load large pieces of geometry.
+
+**-f**
+  Launches the application in full-screen mode (if is supported by GLFW on the
+  OS)
+
+SEE ALSO
+========
+
+`Code Examples <code_examples.html>`__, \
+`glViewer <glviewer.html>`__, \
+`glBatchViewer <glbatchviewer.html>`__, \
+`ptexViewer <ptexviewer.html>`__, \
+`paintTest <painttest.html>`__, \
+`limitEval <limiteval.html>`__, \
+`dxViewer <dxviewer.html>`__, \
+

documentation/glviewer.rst

@@ -69,6 +69,7 @@ SEE ALSO
 
 `Code Examples <code_examples.html>`__, \
 `glBatchViewer <glbatchviewer.html>`__, \
+`glStencilViewer <glstencilviewer.html>`__, \
 `ptexViewer <ptexviewer.html>`__, \
 `paintTest <painttest.html>`__, \
 `limitEval <limiteval.html>`__, \

documentation/nav_template.txt

@@ -51,6 +51,11 @@
                 <ul>
                     <li><a href="hbr_overview.html">Hbr</a></li>
                     <li><a href="far_overview.html">Far</a></li>
+                    <ul>
+                        <li><a href="far_overview.html#subdivision-tables">Subdivision Tables</a></li>
+                        <li><a href="far_overview.html#patch-tables">Patch Tables</a></li>
+                        <li><a href="far_overview.html#stencil-tables">Stencil Tables</a></li>
+                    </ul>
                     <li><a href="osd_overview.html">Osd</a></li>
                 </ul>
             </li>

documentation/painttest.rst

@@ -79,6 +79,7 @@ SEE ALSO
 `Code Examples <code_examples.html>`__, \
 `glViewer <glviewer.html>`__, \
 `glBatchViewer <glbatchviewer.html>`__, \
+`glStencilViewer <glstencilviewer.html>`__, \
 `ptexViewer <ptexviewer.html>`__, \
 `limitEval <limiteval.html>`__, \
 `dxViewer <dxviewer.html>`__, \

documentation/ptexviewer.rst

@@ -134,6 +134,7 @@ SEE ALSO
 `Code Examples <code_examples.html>`__, \
 `glViewer <glviewer.html>`__, \
 `glBatchViewer <glbatchviewer.html>`__, \
+`glStencilViewer <glstencilviewer.html>`__, \
 `paintTest <painttest.html>`__, \
 `limitEval <limiteval.html>`__, \
 `dxViewer <dxviewer.html>`__, \

documentation/release_notes.rst

@@ -32,6 +32,18 @@ Release Notes
 
 ----
 
+Release 2.2.0
+=============
+
+**New Features**
+    - Added subdivision stencil functionality (Far & OsdEval)
+
+**Bug Fixes**
+    - Fix D3D11DrawContext to check for NULL pointers
+    - Fix cpuEvalLimitController crash bug
+    - Fixed search path suffixes for ICC libs
+    - Fixed invalid initialization of glslTransformFeedback kernel.
+
 Release 2.1.0
 =============
 

examples/CMakeLists.txt

@@ -26,6 +26,7 @@
 if( OPENGL_FOUND AND (GLEW_FOUND AND GLFW_FOUND) OR (APPLE AND GLFW_FOUND))
     add_subdirectory(glViewer)
     add_subdirectory(glBatchViewer)
+    add_subdirectory(glStencilViewer)
     add_subdirectory(simpleCpu)
     add_subdirectory(limitEval)
     if(OPENGL_4_3_FOUND AND (NOT APPLE))

examples/common/simple_math.h

@@ -27,6 +27,13 @@
 
 #include <cmath>
 
+inline void
+cross(float *n, const float *v1, const float *v2) {
+    n[0] = v1[1]*v2[2]-v1[2]*v2[1];
+    n[1] = v1[2]*v2[0]-v1[0]*v2[2];
+    n[2] = v1[0]*v2[1]-v1[1]*v2[0];
+}
+
 inline void
 cross(float *n, const float *p0, const float *p1, const float *p2) {
 

examples/glStencilViewer/CMakeLists.txt

@@ -0,0 +1,56 @@
+#
+#     Copyright 2013 Pixar
+#
+#     Licensed under the Apache License, Version 2.0 (the "License");
+#     you may not use this file except in compliance with the License
+#     and the following modification to it: Section 6 Trademarks.
+#     deleted and replaced with:
+#
+#     6. Trademarks. This License does not grant permission to use the
+#     trade names, trademarks, service marks, or product names of the
+#     Licensor and its affiliates, except as required for reproducing
+#     the content of the NOTICE file.
+#
+#     You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+#     Unless required by applicable law or agreed to in writing,
+#     software distributed under the License is distributed on an
+#     "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+#     either express or implied.  See the License for the specific
+#     language governing permissions and limitations under the
+#     License.
+#
+
+# *** glStencilViewer ***
+
+set(PLATFORM_LIBRARIES
+    ${OSD_LINK_TARGET}
+    ${GLFW_LIBRARIES}
+)
+
+include_directories(
+    ${PROJECT_SOURCE_DIR}/opensubdiv
+    ${PROJECT_SOURCE_DIR}/regression
+    ${GLFW_INCLUDE_DIR}
+)
+
+if ( GLEW_FOUND )
+    include_directories(${GLEW_INCLUDE_DIR})
+    list(APPEND PLATFORM_LIBRARIES ${GLEW_LIBRARY})
+endif()
+
+_add_possibly_cuda_executable(glStencilViewer
+    main.cpp
+    ../common/font_image.cpp
+    ../common/hud.cpp
+    ../common/gl_hud.cpp
+    ${INC_FILES}
+)
+
+target_link_libraries(glStencilViewer
+    ${PLATFORM_LIBRARIES}
+)
+
+install(TARGETS glStencilViewer DESTINATION ${CMAKE_BINDIR_BASE})

opensubdiv/far/CMakeLists.txt

@@ -41,6 +41,8 @@ set(PUBLIC_HEADER_FILES
     patchMap.h
     patchTables.h
     patchTablesFactory.h
+    stencilTablesFactory.h
+    stencilTables.h
     subdivisionTables.h
     subdivisionTablesFactory.h
     vertexEditTables.h

opensubdiv/far/stencilTables.h

@@ -0,0 +1,249 @@
+//
+//     Copyright 2013 Pixar
+//
+//     Licensed under the Apache License, Version 2.0 (the "License");
+//     you may not use this file except in compliance with the License
+//     and the following modification to it: Section 6 Trademarks.
+//     deleted and replaced with:
+//
+//     6. Trademarks. This License does not grant permission to use the
+//     trade names, trademarks, service marks, or product names of the
+//     Licensor and its affiliates, except as required for reproducing
+//     the content of the NOTICE file.
+//
+//     You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+//     Unless required by applicable law or agreed to in writing,
+//     software distributed under the License is distributed on an
+//     "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+//     either express or implied.  See the License for the specific
+//     language governing permissions and limitations under the
+//     License.
+//
+
+#ifndef FAR_STENCILTABLES_H
+#define FAR_STENCILTABLES_H
+
+#include "../version.h"
+
+#include <vector>
+
+namespace OpenSubdiv {
+namespace OPENSUBDIV_VERSION {
+
+
+/// \brief Stencil descriptor
+///
+/// Allows access and manipulation of a single stencil in a FarStencilTables.
+///
+class FarStencil {
+public:
+
+    /// Returns the size of the stencil (number of control vertices)
+    int GetSize() const {
+        return *_size;
+    }
+
+    /// Returns the control vertices indices
+    int const * GetVertexIndices() const {
+        return _indices;
+    }
+
+    /// Returns the interpolation weights
+    float const * GetValueWeights() const {
+        return _point;
+    }
+
+    /// Returns U derivative interpolation weights
+    float const * GetUDerivWeights() const {
+        return _uderiv;
+    }
+
+    /// Returns V derivative interpolation weights
+    float const * GetVDerivWeights() const {
+        return _vderiv;
+    }
+
+    /// Increment to the next stencil in the array
+    /// Note : there is no array boundary check !
+    void Increment() {
+        int stride = *_size;
+        ++_size;
+        _indices += stride;
+        _point   += stride;
+        _uderiv  += stride;
+        _vderiv  += stride;
+    }
+
+private:
+
+     FarStencil( int * size,
+                 int * indices,
+                 float * point,
+                 float * uderiv,
+                 float * vderiv )
+         : _size(size),
+           _indices(indices),
+           _point(point),
+           _uderiv(uderiv),
+           _vderiv(vderiv) {
+     }
+
+    friend class FarStencilTables;
+    template <class T> friend class FarStencilTablesFactory;
+
+    int * _size,
+        * _indices;
+
+    float * _point,
+          * _uderiv,
+          * _vderiv;
+
+};
+
+
+/// \brief Table of subdivision stencils.
+///
+/// Stencils are the most direct methods of evaluation of locations on the limit
+/// of a surface. Every point of a limit surface can be computed by linearly
+/// blending a collection of coarse control vertices.
+///
+/// A stencil assigns a series of control vertex indices with a blending weight
+/// that corresponds to a unique parametric location of the limit surface. When
+/// the control vertices move in space, the limit location can be very efficiently
+/// recomputed simply by applying the blending weights to the series of coarse
+/// control vertices.
+///
+class FarStencilTables {
+
+public:
+
+    /// \brief Clears the stencils from the table
+    void Clear() {
+        _sizes.clear();
+        _offsets.clear();
+        _indices.clear();
+        _point.clear();
+        _uderiv.clear();
+        _vderiv.clear();
+    }
+
+    /// \brief Returns the number of stencils in the tables
+    int GetNumStencils() const {
+        return (int)_sizes.size();
+    }
+
+    /// \brief Updates point values based on the control values
+    ///
+    /// \note The values array is assumed to be at least as big as the
+    /// result of \c GetNumStencils().
+    template <class T>
+    void UpdateValues( T const *controlValues, T *values, int stride=0 ) const {
+        _Update( controlValues, &_point.at(0), values, stride );
+    }
+
+    /// \brief Updates derivative values based on the control values
+    ///
+    /// \note The values array is assumed to be at least as big as the
+    /// result of \c GetNumStencils().
+    template <class T>
+    void UpdateDerivs( T const *controlValues, T *uderivs,
+                                               T *vderivs, int stride=0 ) const {
+        _Update( controlValues, &_uderiv.at(0), uderivs, stride );
+        _Update( controlValues, &_vderiv.at(0), vderivs, stride );
+    }
+
+    /// \brief Returns a FarStencil at index i in the tables
+    FarStencil GetStencil(int i) const;
+
+    /// \brief Returns the number of control vertices of each stencil in the table
+    std::vector<int> const & GetSizes() const {
+        return _sizes;
+    }
+
+    /// \brief Returns the offset to a given stencil
+    std::vector<int> const & GetOffsets() const {
+        return _offsets;
+    }
+
+    /// \brief Returns the indices of the control vertices
+    std::vector<int> const & GetControlIndices() const {
+        return _indices;
+    }
+
+    /// \brief Returns the stencils interpolation weights
+    std::vector<float> const & GetWeights() const {
+        return _point;
+    }
+
+    /// \brief Returns the stencils U deriv weights
+    std::vector<float> const & GetDuWeights() const {
+        return _uderiv;
+    }
+
+    /// \brief Returns the stencils U deriv weights
+    std::vector<float> const & GetDvWeights() const {
+        return _vderiv;
+    }
+    
+private:
+
+    template <class T> friend class FarStencilTablesFactory;
+
+    // Update values by appling cached stencil weights to new control values
+    template <class T> void _Update( T const *controlValues,
+                                     float const * weights,
+                                     T *values,
+                                     int stride ) const;
+
+    std::vector<int>    _sizes;   // number of coeffiecient for each stencil
+    std::vector<int>    _offsets; // offset to the start of each stencil
+    std::vector<int>    _indices;
+
+    std::vector<float> _point,       // weight coefficients (value & derivatives)
+                       _uderiv,
+                       _vderiv;
+
+};
+
+template <class T> void
+FarStencilTables::_Update( T const *controlValues,
+                           float const * weights,
+                           T *values,
+                           int stride ) const {
+
+    int const * index = &_indices.at(0);
+
+    for (int i=0; i<GetNumStencils(); ++i) {
+
+        // Zero out the result accumulators
+        values[i].Clear();
+
+        // For each element in the array, add the coefs contribution
+        for (int j=0; j<_sizes[i]; ++j, ++index, ++weights) {
+            values[i].AddWithWeight( controlValues[*index], *weights );
+        }
+    }
+}
+
+inline FarStencil
+FarStencilTables::GetStencil(int i) const {
+
+    int ofs = _offsets[i];
+
+    return FarStencil( const_cast<int *>(&_sizes[i]),
+                       const_cast<int *>(&_indices[ofs]),
+                       const_cast<float *>(&_point[ofs]),
+                       const_cast<float *>(&_uderiv[ofs]),
+                       const_cast<float *>(&_vderiv[ofs]) );
+}
+
+
+} // end namespace OPENSUBDIV_VERSION
+using namespace OPENSUBDIV_VERSION;
+
+} // end namespace OpenSubdiv
+
+#endif // FAR_STENCILTABLES_H

opensubdiv/hbr/mesh.h

@@ -258,19 +258,16 @@ public:
     bool HasCreaseEdits() const { return hasCreaseEdits; }
 
     void Unrefine(int numCoarseVerts, int numCoarseFaces) {
-        static int oldMaxFaceID = 0;
-        if(oldMaxFaceID == 0) {
-            oldMaxFaceID = numCoarseFaces;
-        }
+
         for (int i = numCoarseFaces; i < maxFaceID; ++i) {
             HbrFace<T>* f = GetFace(i);
                 if(f and not f->IsCoarse())
                     DeleteFace(f);
         }
-        //oldMaxFaceID = maxFaceID;
+        
         maxFaceID = numCoarseFaces;
 
-        for(int i=numCoarseVerts; (int)vertices.size(); ++i ) {
+        for(int i=numCoarseVerts; i<(int)vertices.size(); ++i ) {
             HbrVertex<T>* v = GetVertex(i);
             if(v and not v->IsReferenced())
                 DeleteVertex(v);

opensubdiv/osd/CMakeLists.txt

@@ -53,6 +53,8 @@ set(CPU_SOURCE_FILES
     cpuEvalLimitContext.cpp
     cpuEvalLimitController.cpp
     cpuEvalLimitKernel.cpp
+    cpuEvalStencilsContext.cpp
+    cpuEvalStencilsController.cpp
     cpuVertexBuffer.cpp
     error.cpp
     evalLimitContext.cpp
@@ -81,6 +83,8 @@ set(PUBLIC_HEADER_FILES
     cpuComputeController.h
     cpuEvalLimitContext.h
     cpuEvalLimitController.h
+    cpuEvalStencilsContext.h
+    cpuEvalStencilsController.h
     cpuVertexBuffer.h
     error.h
     evalLimitContext.h
@@ -148,12 +152,14 @@ list(APPEND DOXY_HEADER_FILES
 set(OPENMP_PUBLIC_HEADERS
     ompKernel.h
     ompComputeController.h
+    ompEvalStencilsController.h
 )
 
 if( OPENMP_FOUND )
     list(APPEND CPU_SOURCE_FILES
         ompKernel.cpp
         ompComputeController.cpp
+        ompEvalStencilsController.cpp
     )
 
     list(APPEND PUBLIC_HEADER_FILES ${OPENMP_PUBLIC_HEADERS})
@@ -175,10 +181,12 @@ if( TBB_FOUND )
     list(APPEND CPU_SOURCE_FILES
         tbbKernel.cpp
         tbbComputeController.cpp
+        tbbEvalStencilsController.cpp
     )
     list(APPEND PUBLIC_HEADER_FILES
         tbbKernel.h
         tbbComputeController.h
+        tbbEvalStencilsController.h
     )
     if (CMAKE_COMPILER_IS_GNUCXX)
         list(APPEND PLATFORM_CPU_LIBRARIES

opensubdiv/osd/cpuEvalLimitController.cpp

@@ -66,71 +66,64 @@ OsdCpuEvalLimitController::_EvalLimitSample( OpenSubdiv::OsdEvalCoords const & c
     
         int offset = vertexData.outDesc.stride * index;
         
+        
+        if (vertexData.IsBound()) {
+            
+            float * out   = vertexData.out.GetData()+offset, 
+                  * outDu = vertexData.outDu.IsBound() ? vertexData.outDu.GetData()+offset : 0,
+                  * outDv = vertexData.outDv.IsBound() ? vertexData.outDv.GetData()+offset : 0;
+        
             // Based on patch type - go execute interpolation
             switch( parray.GetDescriptor().GetType() ) {
 
-            case FarPatchTables::REGULAR  : if (vertexData.IsBound()) {
-                                                evalBSpline( v, u, cvs,
+                case FarPatchTables::REGULAR  : evalBSpline( v, u, cvs,
                                                              vertexData.inDesc,
                                                              vertexData.in.GetData(),
                                                              vertexData.outDesc,
-                                                             vertexData.out.GetData()+offset, 
-                                                             vertexData.outDu.GetData()+offset, 
-                                                             vertexData.outDv.GetData()+offset );
-                                            } break;
+                                                             out, outDu, outDv );
+                                                break;
 
-            case FarPatchTables::BOUNDARY : if (vertexData.IsBound()) {
-                                                evalBoundary( v, u, cvs,
+                case FarPatchTables::BOUNDARY : evalBoundary( v, u, cvs,
                                                               vertexData.inDesc,
                                                               vertexData.in.GetData(),
                                                               vertexData.outDesc,
-                                                              vertexData.out.GetData()+offset, 
-                                                              vertexData.outDu.GetData()+offset, 
-                                                              vertexData.outDv.GetData()+offset );
-                                            } break;
+                                                              out, outDu, outDv );
+                                                break;
 
-            case FarPatchTables::CORNER   : if (vertexData.IsBound()) {
-                                                evalCorner( v, u, cvs,
+                case FarPatchTables::CORNER   : evalCorner( v, u, cvs,
                                                             vertexData.inDesc,
                                                             vertexData.in.GetData(),
                                                             vertexData.outDesc,
-                                                            vertexData.out.GetData()+offset,
-                                                            vertexData.outDu.GetData()+offset,
-                                                            vertexData.outDv.GetData()+offset );
-                                            } break;
+                                                            out, outDu, outDv );
+                                                break;
 
 
-            case FarPatchTables::GREGORY  : if (vertexData.IsBound()) {
-                                                evalGregory( v, u, cvs,
+                case FarPatchTables::GREGORY  : evalGregory( v, u, cvs,
                                                              &context->GetVertexValenceTable()[0],
                                                              &context->GetQuadOffsetTable()[ parray.GetQuadOffsetIndex() + handle->vertexOffset ],
                                                              context->GetMaxValence(),
                                                              vertexData.inDesc,
                                                              vertexData.in.GetData(),
                                                              vertexData.outDesc,
-                                                             vertexData.out.GetData()+offset, 
-                                                             vertexData.outDu.GetData()+offset, 
-                                                             vertexData.outDv.GetData()+offset );
-                                            } break;
+                                                             out, outDu, outDv );
+                                                break;
 
                 case FarPatchTables::GREGORY_BOUNDARY :
-                                            if (vertexData.IsBound()) {
-                                                evalGregoryBoundary(v, u, cvs,
+                                                evalGregoryBoundary( v, u, cvs,
                                                                      &context->GetVertexValenceTable()[0],
                                                                      &context->GetQuadOffsetTable()[ parray.GetQuadOffsetIndex() + handle->vertexOffset ],
                                                                      context->GetMaxValence(),
                                                                      vertexData.inDesc,
                                                                      vertexData.in.GetData(),
                                                                      vertexData.outDesc,
-                                                                    vertexData.out.GetData()+offset,
-                                                                    vertexData.outDu.GetData()+offset,
-                                                                    vertexData.outDv.GetData()+offset );
-                                            } break;
+                                                                     out, outDu, outDv );
+                                                break;
 
                 default:
                     assert(0);
             }
         }
+    }
     
     OsdCpuEvalLimitContext::VaryingData & varyingData = context->GetVaryingData();
 

opensubdiv/osd/cpuEvalStencilsContext.cpp

@@ -0,0 +1,45 @@
+//
+//     Copyright 2013 Pixar
+//
+//     Licensed under the Apache License, Version 2.0 (the "License");
+//     you may not use this file except in compliance with the License
+//     and the following modification to it: Section 6 Trademarks.
+//     deleted and replaced with:
+//
+//     6. Trademarks. This License does not grant permission to use the
+//     trade names, trademarks, service marks, or product names of the
+//     Licensor and its affiliates, except as required for reproducing
+//     the content of the NOTICE file.
+//
+//     You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+//     Unless required by applicable law or agreed to in writing,
+//     software distributed under the License is distributed on an
+//     "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+//     either express or implied.  See the License for the specific
+//     language governing permissions and limitations under the
+//     License.
+//
+
+#include "../osd/cpuEvalStencilsContext.h"
+
+namespace OpenSubdiv {
+namespace OPENSUBDIV_VERSION {
+
+OsdCpuEvalStencilsContext::OsdCpuEvalStencilsContext(FarStencilTables const *stencils) :
+    _stencils(stencils),
+    _controlData(0),
+    _outputData(0),
+    _outputUDeriv(0),
+    _outputVDeriv(0) {
+}
+
+OsdCpuEvalStencilsContext *
+OsdCpuEvalStencilsContext::Create(FarStencilTables const *stencils) {
+    return new OsdCpuEvalStencilsContext(stencils);
+}
+
+}  // end namespace OPENSUBDIV_VERSION
+}  // end namespace OpenSubdiv

opensubdiv/osd/cpuEvalStencilsContext.h

@@ -0,0 +1,193 @@
+//
+//     Copyright 2013 Pixar
+//
+//     Licensed under the Apache License, Version 2.0 (the "License");
+//     you may not use this file except in compliance with the License
+//     and the following modification to it: Section 6 Trademarks.
+//     deleted and replaced with:
+//
+//     6. Trademarks. This License does not grant permission to use the
+//     trade names, trademarks, service marks, or product names of the
+//     Licensor and its affiliates, except as required for reproducing
+//     the content of the NOTICE file.
+//
+//     You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+//     Unless required by applicable law or agreed to in writing,
+//     software distributed under the License is distributed on an
+//     "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+//     either express or implied.  See the License for the specific
+//     language governing permissions and limitations under the
+//     License.
+//
+
+#ifndef FAR_CPU_EVALSTENCILS_CONTEXT_H
+#define FAR_CPU_EVALSTENCILS_CONTEXT_H
+
+#include "../version.h"
+
+#include "../far/stencilTables.h"
+
+#include "../osd/vertexDescriptor.h"
+#include "../osd/nonCopyable.h"
+
+namespace OpenSubdiv {
+namespace OPENSUBDIV_VERSION {
+
+///
+/// \brief CPU stencils evaluation context
+///
+///
+class OsdCpuEvalStencilsContext : OsdNonCopyable<OsdCpuEvalStencilsContext> {
+
+public:
+    /// \brief Creates an OsdCpuEvalStencilsContext instance
+    ///
+    /// @param stencils  a pointer to the FarStencilTables
+    ///
+    static OsdCpuEvalStencilsContext * Create(FarStencilTables const *stencils);
+
+    /// \brief Binds control vertex data buffer
+    ///
+    /// Binding ensures that data buffers are properly inter-operated between
+    /// Contexts and Controllers operating across multiple devices.
+    ///
+    /// @param controlDataDesc
+    ///
+    /// @param controlData
+    ///
+    template<class VERTEX_BUFFER>
+    void BindControlData(OsdVertexBufferDescriptor const & controlDataDesc, VERTEX_BUFFER *controlData ) {
+
+        _controlData = controlData ? controlData->BindCpuBuffer() : 0;
+        _controlDataDesc = controlDataDesc;
+
+    }
+
+    /// \brief Binds output vertex data buffer
+    ///
+    /// Binding ensures that data buffers are properly inter-operated between
+    /// Contexts and Controllers operating across multiple devices.
+    ///
+    /// @param outputDataDesc
+    ///
+    /// @param outputData
+    ///
+    template<class VERTEX_BUFFER>
+    void BindOutputData( OsdVertexBufferDescriptor const & outputDataDesc, VERTEX_BUFFER *outputData ) {
+
+        _outputData = outputData ? outputData->BindCpuBuffer() : 0;
+        _outputDataDesc = outputDataDesc;
+    }
+    
+    /// \brief Binds output derivative vertex data buffer
+    ///
+    /// Binding ensures that data buffers are properly inter-operated between
+    /// Contexts and Controllers operating across multiple devices.
+    ///
+    /// @param controlDataDesc
+    ///
+    /// @param controlData
+    ///
+    /// @param outputDataDesc
+    ///
+    /// @param outputData
+    ///
+    template<class VERTEX_BUFFER>
+    void BindOutputDerivData( OsdVertexBufferDescriptor const & outputDuDesc, VERTEX_BUFFER *outputDu, 
+                              OsdVertexBufferDescriptor const & outputDvDesc, VERTEX_BUFFER *outputDv ) {
+                              
+        _outputUDeriv = outputDu ? outputDu ->BindCpuBuffer() : 0;
+        _outputVDeriv = outputDv ? outputDv->BindCpuBuffer() : 0;
+        _outputDuDesc = outputDuDesc;
+        _outputDvDesc = outputDvDesc;
+    }
+
+    /// \brief Unbinds any previously bound vertex and varying data buffers.
+    void Unbind() {
+       _controlData = 0;
+       _controlDataDesc.Reset();
+
+       _outputData = 0;
+       _outputDataDesc.Reset();
+      
+       _outputUDeriv = 0;
+       _outputDuDesc.Reset();
+       
+       _outputVDeriv = 0;
+       _outputDvDesc.Reset();
+    }
+
+    /// \brief returns a pointer to the control vertex data
+    float const * GetControlData() const {
+        return _controlData;
+    }
+
+    /// \brief returns a pointer to the output vertex data
+    float * GetOutputData() {
+        return _outputData;
+    }
+
+    /// \brief returns a pointer to the output u-derivative vertex data
+    float * GetOutputUDerivData() {
+        return _outputUDeriv;
+    }
+
+    /// \brief returns a pointer to the output v-derivative vertex data
+    float * GetOutputVDerivData() {
+        return _outputVDeriv;
+    }
+
+    /// \brief Returns the vertex data descriptor for the control vertices buffer
+    OsdVertexBufferDescriptor GetControlDataDescriptor() const {
+        return _controlDataDesc;
+    }
+
+    /// \brief Returns the vertex data descriptor for the output vertices buffer
+    OsdVertexBufferDescriptor GetOutputDataDescriptor() const {
+        return _outputDataDesc;
+    }
+
+    /// \brief Returns the vertex data descriptor for the U derivatives data buffer
+    OsdVertexBufferDescriptor GetDuDataDescriptor() const {
+        return _outputDuDesc;
+    }
+
+    /// \brief Returns the vertex data descriptor for the V derivatives data buffer
+    OsdVertexBufferDescriptor GetDvDataDescriptor() const {
+        return _outputDvDesc;
+    }
+    
+
+    /// \brief Returns the FarStencilTables applied
+    FarStencilTables const * GetStencilTables() const {
+        return _stencils;
+    }
+
+protected:
+
+    OsdCpuEvalStencilsContext(FarStencilTables const *stencils);
+
+private:
+    
+    FarStencilTables const * _stencils;
+
+    OsdVertexBufferDescriptor _controlDataDesc,
+                              _outputDataDesc,
+                              _outputDuDesc,
+                              _outputDvDesc;
+
+    float * _controlData,
+          * _outputData,
+          * _outputUDeriv,
+          * _outputVDeriv;
+};
+
+} // end namespace OPENSUBDIV_VERSION
+using namespace OPENSUBDIV_VERSION;
+
+} // end namespace OpenSubdiv
+
+#endif // FAR_CPU_EVALSTENCILS_CONTEXT_H

opensubdiv/osd/cpuEvalStencilsController.cpp

@@ -0,0 +1,147 @@
+//
+//     Copyright 2013 Pixar
+//
+//     Licensed under the Apache License, Version 2.0 (the "License");
+//     you may not use this file except in compliance with the License
+//     and the following modification to it: Section 6 Trademarks.
+//     deleted and replaced with:
+//
+//     6. Trademarks. This License does not grant permission to use the
+//     trade names, trademarks, service marks, or product names of the
+//     Licensor and its affiliates, except as required for reproducing
+//     the content of the NOTICE file.
+//
+//     You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+//     Unless required by applicable law or agreed to in writing,
+//     software distributed under the License is distributed on an
+//     "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+//     either express or implied.  See the License for the specific
+//     language governing permissions and limitations under the
+//     License.
+//
+
+#include "../osd/cpuEvalStencilsController.h"
+
+#include <cassert>
+
+namespace OpenSubdiv {
+namespace OPENSUBDIV_VERSION {
+
+OsdCpuEvalStencilsController::OsdCpuEvalStencilsController() {
+}
+
+OsdCpuEvalStencilsController::~OsdCpuEvalStencilsController() {
+}
+
+int 
+OsdCpuEvalStencilsController::_UpdateValues( OsdCpuEvalStencilsContext * context ) {
+
+    int result=0;
+
+    FarStencilTables const * stencils = context->GetStencilTables();
+
+    int nstencils = stencils->GetNumStencils();
+    if (not nstencils)
+        return result;
+    
+    OsdVertexBufferDescriptor ctrlDesc = context->GetControlDataDescriptor(),
+                              outDesc = context->GetOutputDataDescriptor();
+    
+    // make sure that we have control data to work with
+    if (not ctrlDesc.CanEval(outDesc))
+        return 0;
+
+    float const * ctrl = context->GetControlData() + ctrlDesc.offset;
+
+    float * out = context->GetOutputData() + outDesc.offset;
+
+    if ((not ctrl) or (not out))
+        return result;
+
+    int const * sizes = &stencils->GetSizes().at(0),
+              * index = &stencils->GetControlIndices().at(0);
+    
+    float const * weight = &stencils->GetWeights().at(0);
+
+    for (int i=0; i<nstencils; ++i) {
+
+        memset(out, 0, outDesc.length*sizeof(float));
+        
+        for (int j=0; j<sizes[i]; ++j, ++index, ++weight) {
+        
+            float const * cv = ctrl + (*index)*ctrlDesc.stride;
+
+            for (int k=0; k<outDesc.length; ++k) {
+                out[k] += cv[k] * (*weight);
+            }
+        }
+        out += outDesc.stride;
+    }
+    
+    return nstencils;
+}
+
+int 
+OsdCpuEvalStencilsController::_UpdateDerivs( OsdCpuEvalStencilsContext * context ) {
+
+    int result=0;
+
+    FarStencilTables const * stencils = context->GetStencilTables();
+
+    int nstencils = stencils->GetNumStencils();
+    if (not nstencils)
+        return result;
+    
+    OsdVertexBufferDescriptor ctrlDesc = context->GetControlDataDescriptor(),
+                              duDesc = context->GetDuDataDescriptor(),
+                              dvDesc = context->GetDvDataDescriptor();
+    
+    // make sure that we have control data to work with
+    if (not (ctrlDesc.CanEval(duDesc) and ctrlDesc.CanEval(dvDesc)))
+        return 0;
+
+    float const * ctrl = context->GetControlData() + ctrlDesc.offset;
+
+    float * du = context->GetOutputUDerivData() + duDesc.offset,
+          * dv = context->GetOutputVDerivData() + dvDesc.offset;
+
+    if ((not ctrl) or (not du) or (not dv))
+        return result;
+
+    int const * sizes = &stencils->GetSizes().at(0),
+              * index = &stencils->GetControlIndices().at(0);
+    
+    float const * duweight = &stencils->GetDuWeights().at(0),
+                * dvweight = &stencils->GetDvWeights().at(0);
+                           
+    for (int i=0; i<nstencils; ++i) {
+
+        memset(du, 0, duDesc.length*sizeof(float));
+        memset(dv, 0, dvDesc.length*sizeof(float));
+        
+        for (int j=0; j<sizes[i]; ++j, ++index, ++duweight, ++dvweight) {
+        
+            float const * cv = ctrl + (*index)*ctrlDesc.stride;
+        
+            for (int k=0; k<duDesc.length; ++k) {
+                du[k] += cv[k] * (*duweight);
+                dv[k] += cv[k] * (*dvweight);
+            }
+        }
+        du += duDesc.stride;
+        dv += dvDesc.stride;
+    }
+    
+    return nstencils;
+}
+
+void
+OsdCpuEvalStencilsController::Synchronize() {
+}
+
+
+}  // end namespace OPENSUBDIV_VERSION
+}  // end namespace OpenSubdiv

opensubdiv/osd/cpuEvalStencilsController.h

@@ -0,0 +1,144 @@
+//
+//     Copyright 2013 Pixar
+//
+//     Licensed under the Apache License, Version 2.0 (the "License");
+//     you may not use this file except in compliance with the License
+//     and the following modification to it: Section 6 Trademarks.
+//     deleted and replaced with:
+//
+//     6. Trademarks. This License does not grant permission to use the
+//     trade names, trademarks, service marks, or product names of the
+//     Licensor and its affiliates, except as required for reproducing
+//     the content of the NOTICE file.
+//
+//     You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+//     Unless required by applicable law or agreed to in writing,
+//     software distributed under the License is distributed on an
+//     "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+//     either express or implied.  See the License for the specific
+//     language governing permissions and limitations under the
+//     License.
+//
+
+#ifndef FAR_CPU_EVALSTENCILS_CONTROLLER_H
+#define FAR_CPU_EVALSTENCILS_CONTROLLER_H
+
+#include "../version.h"
+
+#include "../osd/cpuEvalStencilsContext.h"
+
+namespace OpenSubdiv {
+namespace OPENSUBDIV_VERSION {
+
+///
+/// \brief CPU stencils evaluation controller
+///
+/// OsdCpuStencilsController is a compute controller class to launch
+/// single threaded CPU stencil evalution kernels.
+///
+/// Controller entities execute requests from Context instances that they share
+/// common interfaces with. Controllers are attached to discrete compute devices
+/// and share the devices resources with Context entities.
+///
+class OsdCpuEvalStencilsController {
+public:
+
+    /// Constructor.
+    OsdCpuEvalStencilsController();
+
+    /// Destructor.
+    ~OsdCpuEvalStencilsController();
+    
+    
+    /// \brief Applies stencil weights to the control vertex data
+    ///
+    /// Applies the stencil weights to the control vertex data to evaluate the
+    /// interpolated limit positions at the parametric locations of the stencils
+    ///
+    /// @param context          the OsdCpuEvalStencilsContext with the stencil weights
+    ///
+    /// @param controlDataDesc  vertex buffer descriptor for the control vertex data 
+    ///
+    /// @param controlVertices  vertex buffer with the control vertices data
+    ///
+    /// @param outputDataDesc   vertex buffer descriptor for the output vertex data
+    ///
+    /// @param outputData       vertex buffer where the vertex data will be output
+    ///
+    template<class CONTROL_BUFFER, class OUTPUT_BUFFER>
+    int UpdateValues( OsdCpuEvalStencilsContext * context,
+                      OsdVertexBufferDescriptor const & controlDataDesc, CONTROL_BUFFER *controlVertices,
+                      OsdVertexBufferDescriptor const & outputDataDesc, OUTPUT_BUFFER *outputData ) {
+                       
+        if (not context->GetStencilTables()->GetNumStencils())
+            return 0;
+
+        context->BindControlData( controlDataDesc, controlVertices );
+
+        context->BindOutputData( outputDataDesc, outputData );
+        
+        int n = _UpdateValues( context );
+        
+        context->Unbind();
+        
+        return n;
+    }
+
+    /// \brief Applies derivative stencil weights to the control vertex data
+    ///
+    /// Computes the U and V derivative stencils to the control vertex data at
+    /// the parametric locations contained in each stencil
+    ///
+    /// @param context          the OsdCpuEvalStencilsContext with the stencil weights
+    ///
+    /// @param controlDataDesc  vertex buffer descriptor for the control vertex data 
+    ///
+    /// @param controlVertices  vertex buffer with the control vertices data
+    ///
+    /// @param outputDuDesc     vertex buffer descriptor for the U derivative output data
+    ///
+    /// @param outputDuData     output vertex buffer for the U derivative data
+    ///
+    /// @param outputDvDesc     vertex buffer descriptor for the V deriv output data
+    ///
+    /// @param outputDvData     output vertex buffer for the V derivative data
+    ///
+    template<class CONTROL_BUFFER, class OUTPUT_BUFFER>
+    int UpdateDerivs( OsdCpuEvalStencilsContext * context,
+                      OsdVertexBufferDescriptor const & controlDataDesc, CONTROL_BUFFER *controlVertices,
+                      OsdVertexBufferDescriptor const & outputDuDesc, OUTPUT_BUFFER *outputDuData, 
+                      OsdVertexBufferDescriptor const & outputDvDesc, OUTPUT_BUFFER *outputDvData ) {
+                       
+        if (not context->GetStencilTables()->GetNumStencils())
+            return 0;
+
+        context->BindControlData( controlDataDesc, controlVertices );
+
+        context->BindOutputDerivData( outputDuDesc, outputDuData, outputDvDesc, outputDvData );
+        
+        int n = _UpdateDerivs( context );
+        
+        context->Unbind();
+        
+        return n;
+    }
+    
+    /// Waits until all running subdivision kernels finish.
+    void Synchronize();
+
+private:
+
+    int _UpdateValues( OsdCpuEvalStencilsContext * context );
+    int _UpdateDerivs( OsdCpuEvalStencilsContext * context );
+
+};
+
+} // end namespace OPENSUBDIV_VERSION
+using namespace OPENSUBDIV_VERSION;
+
+} // end namespace OpenSubdiv
+
+#endif // FAR_CPU_EVALSTENCILS_CONTROLLER_H

opensubdiv/osd/d3d11DrawContext.h

@@ -86,6 +86,7 @@ public:
     ///
     template<class VERTEX_BUFFER>
     void UpdateVertexTexture(VERTEX_BUFFER *vbo, ID3D11DeviceContext *pd3d11DeviceContext) {
+        if (vbo)
             updateVertexTexture(vbo->BindD3D11Buffer(pd3d11DeviceContext),
                                 pd3d11DeviceContext,
                                 vbo->GetNumVertices(),

opensubdiv/osd/glDrawContext.h

@@ -70,6 +70,7 @@ public:
     ///
     template<class VERTEX_BUFFER>
     void UpdateVertexTexture(VERTEX_BUFFER *vbo) {
+        if (vbo)
             updateVertexTexture(vbo->BindVBO(), vbo->GetNumElements());
     }
 

opensubdiv/osd/glslTransformFeedbackKernelBundle.cpp

@@ -40,6 +40,7 @@
 #include "../osd/opengl.h"
 
 #include <cassert>
+#include <string>
 
 namespace OpenSubdiv {
 namespace OPENSUBDIV_VERSION {
@@ -94,20 +95,30 @@ OsdGLSLTransformFeedbackKernelBundle::Compile(int numVertexElements, int numVary
     glCompileShader(shader);
     glAttachShader(_program, shader);
 
-    const char *outputs[4];
-    int nOutputs = 0;
+    std::vector<std::string> outputs;
 
     // position and custom vertex data are stored same buffer whereas varying data
     // exists on another buffer. "gl_NextBuffer" identifier helps to split them.
-    if (numVertexElements > 0)
-        outputs[nOutputs++] = "outVertexData";
-    if (numVaryingElements > 0) {
-        if (nOutputs > 0)
-            outputs[nOutputs++] = "gl_NextBuffer";
-        outputs[nOutputs++] = "outVaryingData";
+    for (int i = 0; i < numVertexElements; ++i) {
+        char attrName[32];
+        snprintf(attrName, 32, "outVertexData[%d]", i);
+        outputs.push_back(attrName);
+    }
+    for (int i = 0; i < numVaryingElements; ++i) {
+        if (i == 0 and (not outputs.empty())) {
+            outputs.push_back("gl_NextBuffer");
+        }
+        char attrName[32];
+        snprintf(attrName, 32, "outVaryingData[%d]", i);
+        outputs.push_back(attrName);
+    }
+    std::vector<const char *> pOutputs;
+    for (size_t i = 0; i < outputs.size(); ++i) {
+        pOutputs.push_back(&outputs[i][0]);
     }
 
-    glTransformFeedbackVaryings(_program, nOutputs, outputs, GL_INTERLEAVED_ATTRIBS);
+    glTransformFeedbackVaryings(_program, (GLsizei)outputs.size(),
+                                &pOutputs[0], GL_INTERLEAVED_ATTRIBS);
 
     OSD_DEBUG_CHECK_GL_ERROR("Transform feedback initialize\n");
 

opensubdiv/osd/ompEvalStencilsController.cpp

@@ -0,0 +1,152 @@
+//
+//     Copyright 2013 Pixar
+//
+//     Licensed under the Apache License, Version 2.0 (the "License");
+//     you may not use this file except in compliance with the License
+//     and the following modification to it: Section 6 Trademarks.
+//     deleted and replaced with:
+//
+//     6. Trademarks. This License does not grant permission to use the
+//     trade names, trademarks, service marks, or product names of the
+//     Licensor and its affiliates, except as required for reproducing
+//     the content of the NOTICE file.
+//
+//     You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+//     Unless required by applicable law or agreed to in writing,
+//     software distributed under the License is distributed on an
+//     "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+//     either express or implied.  See the License for the specific
+//     language governing permissions and limitations under the
+//     License.
+//
+
+#include "../osd/ompEvalStencilsController.h"
+
+#include <cassert>
+
+namespace OpenSubdiv {
+namespace OPENSUBDIV_VERSION {
+
+OsdOmpEvalStencilsController::OsdOmpEvalStencilsController(int numThreads) {
+
+    _numThreads = (numThreads == -1) ? omp_get_num_procs() : numThreads;
+}
+
+OsdOmpEvalStencilsController::~OsdOmpEvalStencilsController() {
+}
+
+int
+OsdOmpEvalStencilsController::_UpdateValues( OsdCpuEvalStencilsContext * context ) {
+
+    int result=0;
+
+    FarStencilTables const * stencils = context->GetStencilTables();
+
+    int nstencils = stencils->GetNumStencils();
+    if (not nstencils)
+        return result;
+
+    OsdVertexBufferDescriptor ctrlDesc = context->GetControlDataDescriptor(),
+                              outDesc = context->GetOutputDataDescriptor();
+
+    // make sure that we have control data to work with
+    if (not ctrlDesc.CanEval(outDesc))
+        return 0;
+
+    float const * ctrl = context->GetControlData() + ctrlDesc.offset;
+
+    if (not ctrl)
+        return result;
+
+#pragma omp parallel for
+    for (int i=0; i<nstencils; ++i) {
+
+        int size = stencils->GetSizes()[i],
+            offset = stencils->GetOffsets()[i];
+
+        int const * index = &stencils->GetControlIndices().at(offset);
+
+        float const * weight = &stencils->GetWeights().at(offset);
+
+        float * out = context->GetOutputData() + i * outDesc.stride + outDesc.offset;
+
+        memset(out, 0, outDesc.length*sizeof(float));
+
+        for (int j=0; j<size; ++j, ++index, ++weight) {
+
+            float const * cv = ctrl + (*index)*ctrlDesc.stride;
+
+            for (int k=0; k<outDesc.length; ++k) {
+                out[k] += cv[k] * (*weight);
+            }
+        }
+    }
+
+    return nstencils;
+}
+
+int
+OsdOmpEvalStencilsController::_UpdateDerivs( OsdCpuEvalStencilsContext * context ) {
+
+    int result=0;
+
+    FarStencilTables const * stencils = context->GetStencilTables();
+
+    int nstencils = stencils->GetNumStencils();
+    if (not nstencils)
+        return result;
+
+    OsdVertexBufferDescriptor ctrlDesc = context->GetControlDataDescriptor(),
+                              duDesc = context->GetDuDataDescriptor(),
+                              dvDesc = context->GetDvDataDescriptor();
+
+    // make sure that we have control data to work with
+    if (not (ctrlDesc.CanEval(duDesc) and ctrlDesc.CanEval(dvDesc)))
+        return 0;
+
+    float const * ctrl = context->GetControlData() + ctrlDesc.offset;
+
+    if (not ctrl)
+        return result;
+
+#pragma omp parallel for
+    for (int i=0; i<nstencils; ++i) {
+
+        int size = stencils->GetSizes()[i],
+            offset = stencils->GetOffsets()[i];
+
+        int const * index = &stencils->GetControlIndices().at(offset);
+
+        float const * duweight = &stencils->GetDuWeights().at(offset),
+                    * dvweight = &stencils->GetDvWeights().at(offset);
+
+        float * du = context->GetOutputUDerivData() + i * duDesc.stride + duDesc.offset,
+              * dv = context->GetOutputVDerivData() + i * dvDesc.stride + dvDesc.offset;
+
+        memset(du, 0, duDesc.length*sizeof(float));
+        memset(dv, 0, dvDesc.length*sizeof(float));
+
+        for (int j=0; j<size; ++j, ++index, ++duweight, ++dvweight) {
+
+            float const * cv = ctrl + (*index)*ctrlDesc.stride;
+
+            for (int k=0; k<duDesc.length; ++k) {
+                du[k] += cv[k] * (*duweight);
+                dv[k] += cv[k] * (*dvweight);
+            }
+        }
+    }
+
+    return nstencils;
+}
+
+void
+OsdOmpEvalStencilsController::Synchronize() {
+}
+
+
+}  // end namespace OPENSUBDIV_VERSION
+}  // end namespace OpenSubdiv

opensubdiv/osd/ompEvalStencilsController.h

@@ -0,0 +1,154 @@
+//
+//     Copyright 2013 Pixar
+//
+//     Licensed under the Apache License, Version 2.0 (the "License");
+//     you may not use this file except in compliance with the License
+//     and the following modification to it: Section 6 Trademarks.
+//     deleted and replaced with:
+//
+//     6. Trademarks. This License does not grant permission to use the
+//     trade names, trademarks, service marks, or product names of the
+//     Licensor and its affiliates, except as required for reproducing
+//     the content of the NOTICE file.
+//
+//     You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+//     Unless required by applicable law or agreed to in writing,
+//     software distributed under the License is distributed on an
+//     "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+//     either express or implied.  See the License for the specific
+//     language governing permissions and limitations under the
+//     License.
+//
+
+#ifndef FAR_OMP_EVALSTENCILS_CONTROLLER_H
+#define FAR_OMP_EVALSTENCILS_CONTROLLER_H
+
+#include "../version.h"
+
+#include "../osd/cpuEvalStencilsContext.h"
+
+#ifdef OPENSUBDIV_HAS_OPENMP
+    #include <omp.h>
+#endif
+
+namespace OpenSubdiv {
+namespace OPENSUBDIV_VERSION {
+
+///
+/// \brief CPU stencils evaluation controller
+///
+/// OsdCpuStencilsController is a compute controller class to launch
+/// single threaded CPU stencil evalution kernels.
+///
+/// Controller entities execute requests from Context instances that they share
+/// common interfaces with. Controllers are attached to discrete compute devices
+/// and share the devices resources with Context entities.
+///
+class OsdOmpEvalStencilsController {
+public:
+
+    /// \brief Constructor.
+    ///
+    /// @param numThreads specifies how many openmp parallel threads to use.
+    ///                   -1 attempts to use all available processors.
+    ///
+    OsdOmpEvalStencilsController(int numThreads=-1);
+
+    /// \brief Destructor.
+    ~OsdOmpEvalStencilsController();
+
+
+    /// \brief Applies stencil weights to the control vertex data
+    ///
+    /// Applies the stencil weights to the control vertex data to evaluate the
+    /// interpolated limit positions at the parametric locations of the stencils
+    ///
+    /// @param context          the OsdCpuEvalStencilsContext with the stencil weights
+    ///
+    /// @param controlDataDesc  vertex buffer descriptor for the control vertex data 
+    ///
+    /// @param controlVertices  vertex buffer with the control vertices data
+    ///
+    /// @param outputDataDesc   vertex buffer descriptor for the output vertex data
+    ///
+    /// @param outputData       output vertex buffer for the interpolated data
+    ///
+    template<class CONTROL_BUFFER, class OUTPUT_BUFFER>
+    int UpdateValues( OsdCpuEvalStencilsContext * context,
+                      OsdVertexBufferDescriptor const & controlDataDesc, CONTROL_BUFFER *controlVertices,
+                      OsdVertexBufferDescriptor const & outputDataDesc, OUTPUT_BUFFER *outputData ) {
+
+        if (not context->GetStencilTables()->GetNumStencils())
+            return 0;
+
+        omp_set_num_threads(_numThreads);
+
+        context->BindControlData( controlDataDesc, controlVertices );
+        context->BindOutputData( outputDataDesc, outputData );
+
+        int n = _UpdateValues( context );
+
+        context->Unbind();
+
+        return n;
+    }
+
+    /// \brief Applies derivative stencil weights to the control vertex data
+    ///
+    /// Computes the U and V derivative stencils to the control vertex data at
+    /// the parametric locations contained in each stencil
+    ///
+    /// @param context          the OsdCpuEvalStencilsContext with the stencil weights
+    ///
+    /// @param controlDataDesc  vertex buffer descriptor for the control vertex data 
+    ///
+    /// @param controlVertices  vertex buffer with the control vertices data
+    ///
+    /// @param outputDuDesc     vertex buffer descriptor for the U derivative output data
+    ///
+    /// @param outputDuData     output vertex buffer for the U derivative data
+    ///
+    /// @param outputDvDesc     vertex buffer descriptor for the V deriv output data
+    ///
+    /// @param outputDvData     output vertex buffer for the V derivative data
+    ///
+    template<class CONTROL_BUFFER, class OUTPUT_BUFFER>
+    int UpdateDerivs( OsdCpuEvalStencilsContext * context,
+                      OsdVertexBufferDescriptor const & controlDataDesc, CONTROL_BUFFER *controlVertices,
+                      OsdVertexBufferDescriptor const & outputDuDesc, OUTPUT_BUFFER *outputDuData, 
+                      OsdVertexBufferDescriptor const & outputDvDesc, OUTPUT_BUFFER *outputDvData ) {
+                       
+        if (not context->GetStencilTables()->GetNumStencils())
+            return 0;
+
+        context->BindControlData( controlDataDesc, controlVertices );
+
+        context->BindOutputDerivData( outputDuDesc, outputDuData, outputDvDesc, outputDvData );
+        
+        int n = _UpdateDerivs( context );
+        
+        context->Unbind();
+        
+        return n;
+    }
+
+    /// Waits until all running subdivision kernels finish.
+    void Synchronize();
+
+private:
+
+    int _UpdateValues( OsdCpuEvalStencilsContext * context );
+    int _UpdateDerivs( OsdCpuEvalStencilsContext * context );
+
+    int _numThreads;
+};
+
+} // end namespace OPENSUBDIV_VERSION
+using namespace OPENSUBDIV_VERSION;
+
+} // end namespace OpenSubdiv
+
+#endif // FAR_OMP_EVALSTENCILS_CONTROLLER_H

opensubdiv/osd/tbbEvalStencilsController.cpp

@@ -0,0 +1,197 @@
+//
+//     Copyright 2013 Pixar
+//
+//     Licensed under the Apache License, Version 2.0 (the "License");
+//     you may not use this file except in compliance with the License
+//     and the following modification to it: Section 6 Trademarks.
+//     deleted and replaced with:
+//
+//     6. Trademarks. This License does not grant permission to use the
+//     trade names, trademarks, service marks, or product names of the
+//     Licensor and its affiliates, except as required for reproducing
+//     the content of the NOTICE file.
+//
+//     You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+//     Unless required by applicable law or agreed to in writing,
+//     software distributed under the License is distributed on an
+//     "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+//     either express or implied.  See the License for the specific
+//     language governing permissions and limitations under the
+//     License.
+//
+
+#include "../osd/tbbEvalStencilsController.h"
+
+#include <tbb/parallel_for.h>
+#include <tbb/task_scheduler_init.h>
+
+#include <cassert>
+
+namespace OpenSubdiv {
+namespace OPENSUBDIV_VERSION {
+
+#define grain_size  200
+
+OsdTbbEvalStencilsController::OsdTbbEvalStencilsController(int numThreads) {
+
+    _numThreads = numThreads > 0 ? numThreads : tbb::task_scheduler_init::automatic;
+
+   tbb::task_scheduler_init init(numThreads);
+}
+
+OsdTbbEvalStencilsController::~OsdTbbEvalStencilsController() {
+}
+
+
+class StencilKernel {
+
+public:
+    enum Mode { UNDEFINED, POINT, U_DERIV, V_DERIV };
+
+    StencilKernel( FarStencilTables const * stencils,
+                   OsdVertexBufferDescriptor ctrlDesc,
+                   float const * ctrlData ) :
+        _stencils(stencils),
+        _mode(UNDEFINED),
+        _ctrlDesc(ctrlDesc),
+        _length(0),
+        _outStride(0),
+        _outData(0) {
+
+        _ctrlData = ctrlData + ctrlDesc.offset;
+    }
+
+    bool SetOutput(Mode mode, OsdVertexBufferDescriptor outDesc, float * outData) {
+
+        if (_ctrlDesc.CanEval(outDesc)) {
+            _mode = mode;
+            _length = outDesc.length;
+            _outStride = outDesc.stride;
+            _outData = outData + outDesc.offset;
+            return true;
+        }
+        return false;
+    }
+
+    void operator() (tbb::blocked_range<int> const &r) const {
+
+        assert(_stencils and _ctrlData and _length and _outStride and _outData);
+
+        int offset = _stencils->GetOffsets()[r.begin()];
+
+        int const * sizes = &_stencils->GetSizes()[r.begin()],
+                  * index = &_stencils->GetControlIndices()[offset];
+
+        float const * weight;
+
+        switch (_mode) {
+            case POINT   : weight = &_stencils->GetWeights()[offset]; break;
+            case U_DERIV : weight = &_stencils->GetDuWeights()[offset]; break;
+            case V_DERIV : weight = &_stencils->GetDvWeights()[offset]; break;
+            default:
+                return;
+        }
+        assert( weight);
+
+        float * out = _outData + r.begin() * _outStride;
+
+        for (int i=r.begin(); i<r.end(); ++i, ++sizes) {
+
+            memset( out, 0, _length * sizeof(float) );
+
+            for (int j=0; j<(*sizes); ++j, ++index, ++weight) {
+
+                float const * cv = _ctrlData + (*index)*_ctrlDesc.stride;
+
+                for (int k=0; k<_length; ++k) {
+                    out[k] += cv[k] * (*weight);
+                }
+            }
+            out+=_outStride;
+        }
+    }
+
+private:
+    FarStencilTables const * _stencils;
+
+    Mode _mode;
+
+    OsdVertexBufferDescriptor _ctrlDesc;
+    float const * _ctrlData;
+
+    int _length,
+        _outStride;
+
+    float * _outData;
+};
+
+int
+OsdTbbEvalStencilsController::_UpdateValues( OsdCpuEvalStencilsContext * context ) {
+
+    FarStencilTables const * stencils = context->GetStencilTables();
+    if (not stencils)
+        return 0;
+
+    int nstencils = stencils->GetNumStencils();
+    if (not nstencils)
+        return 0;
+
+    StencilKernel kernel( stencils, context->GetControlDataDescriptor(),
+                                    context->GetControlData() );
+
+
+    if (not kernel.SetOutput( StencilKernel::POINT,
+                              context->GetOutputDataDescriptor(),
+                              context->GetOutputData() ))
+        return 0;
+
+    tbb::blocked_range<int> range(0, nstencils, grain_size);
+
+    tbb::parallel_for(range, kernel);
+
+    return nstencils;
+}
+
+int
+OsdTbbEvalStencilsController::_UpdateDerivs( OsdCpuEvalStencilsContext * context ) {
+
+    FarStencilTables const * stencils = context->GetStencilTables();
+    if (not stencils)
+        return 0;
+
+    int nstencils = stencils->GetNumStencils();
+    if (not nstencils)
+        return 0;
+
+    tbb::blocked_range<int> range(0, nstencils, grain_size);
+
+    StencilKernel kernel( stencils, context->GetControlDataDescriptor(),
+                                    context->GetControlData() );
+
+    if (not kernel.SetOutput( StencilKernel::U_DERIV,
+                              context->GetDuDataDescriptor(),
+                              context->GetOutputUDerivData() ) )
+        return 0;
+
+    tbb::parallel_for(range, kernel);
+
+    if (not kernel.SetOutput( StencilKernel::V_DERIV,
+                              context->GetDvDataDescriptor(),
+                              context->GetOutputVDerivData() ) )
+        return 0;
+
+    tbb::parallel_for(range, kernel);
+
+    return nstencils;
+}
+
+void
+OsdTbbEvalStencilsController::Synchronize() {
+}
+
+
+}  // end namespace OPENSUBDIV_VERSION
+}  // end namespace OpenSubdiv

opensubdiv/osd/tbbEvalStencilsController.h

@@ -0,0 +1,149 @@
+//
+//     Copyright 2013 Pixar
+//
+//     Licensed under the Apache License, Version 2.0 (the "License");
+//     you may not use this file except in compliance with the License
+//     and the following modification to it: Section 6 Trademarks.
+//     deleted and replaced with:
+//
+//     6. Trademarks. This License does not grant permission to use the
+//     trade names, trademarks, service marks, or product names of the
+//     Licensor and its affiliates, except as required for reproducing
+//     the content of the NOTICE file.
+//
+//     You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+//     Unless required by applicable law or agreed to in writing,
+//     software distributed under the License is distributed on an
+//     "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND,
+//     either express or implied.  See the License for the specific
+//     language governing permissions and limitations under the
+//     License.
+//
+
+#ifndef FAR_TBB_EVALSTENCILS_CONTROLLER_H
+#define FAR_TBB_EVALSTENCILS_CONTROLLER_H
+
+#include "../version.h"
+
+#include "../osd/cpuEvalStencilsContext.h"
+
+
+namespace OpenSubdiv {
+namespace OPENSUBDIV_VERSION {
+
+///
+/// \brief CPU stencils evaluation controller
+///
+/// OsdCpuStencilsController is a compute controller class to launch
+/// single threaded CPU stencil evalution kernels.
+///
+/// Controller entities execute requests from Context instances that they share
+/// common interfaces with. Controllers are attached to discrete compute devices
+/// and share the devices resources with Context entities.
+///
+class OsdTbbEvalStencilsController {
+public:
+
+    /// \brief Constructor.
+    ///
+    /// @param numThreads specifies how many openmp parallel threads to use.
+    ///                   -1 attempts to use all available processors.
+    ///
+    OsdTbbEvalStencilsController(int numThreads=-1);
+
+    /// \brief Destructor.
+    ~OsdTbbEvalStencilsController();
+
+
+    /// \brief Applies stencil weights to the control vertex data
+    ///
+    /// Applies the stencil weights to the control vertex data to evaluate the
+    /// interpolated limit positions at the parametric locations of the stencils
+    ///
+    /// @param context          the OsdCpuEvalStencilsContext with the stencil weights
+    ///
+    /// @param controlDataDesc  vertex buffer descriptor for the control vertex data
+    ///
+    /// @param controlVertices  vertex buffer with the control vertices data
+    ///
+    /// @param outputDataDesc   vertex buffer descriptor for the output vertex data
+    ///
+    /// @param outputData       output vertex buffer for the interpolated data
+    ///
+    template<class CONTROL_BUFFER, class OUTPUT_BUFFER>
+    int UpdateValues( OsdCpuEvalStencilsContext * context,
+                      OsdVertexBufferDescriptor const & controlDataDesc, CONTROL_BUFFER *controlVertices,
+                      OsdVertexBufferDescriptor const & outputDataDesc, OUTPUT_BUFFER *outputData ) {
+
+        if (not context->GetStencilTables()->GetNumStencils())
+            return 0;
+
+        context->BindControlData( controlDataDesc, controlVertices );
+        context->BindOutputData( outputDataDesc, outputData );
+
+        int n = _UpdateValues( context );
+
+        context->Unbind();
+
+        return n;
+    }
+
+    /// \brief Applies derivative stencil weights to the control vertex data
+    ///
+    /// Computes the U and V derivative stencils to the control vertex data at
+    /// the parametric locations contained in each stencil
+    ///
+    /// @param context          the OsdCpuEvalStencilsContext with the stencil weights
+    ///
+    /// @param controlDataDesc  vertex buffer descriptor for the control vertex data
+    ///
+    /// @param controlVertices  vertex buffer with the control vertices data
+    ///
+    /// @param outputDuDesc     vertex buffer descriptor for the U derivative output data
+    ///
+    /// @param outputDuData     output vertex buffer for the U derivative data
+    ///
+    /// @param outputDvDesc     vertex buffer descriptor for the V deriv output data
+    ///
+    /// @param outputDvData     output vertex buffer for the V derivative data
+    ///
+    template<class CONTROL_BUFFER, class OUTPUT_BUFFER>
+    int UpdateDerivs( OsdCpuEvalStencilsContext * context,
+                      OsdVertexBufferDescriptor const & controlDataDesc, CONTROL_BUFFER *controlVertices,
+                      OsdVertexBufferDescriptor const & outputDuDesc, OUTPUT_BUFFER *outputDuData,
+                      OsdVertexBufferDescriptor const & outputDvDesc, OUTPUT_BUFFER *outputDvData ) {
+
+        if (not context->GetStencilTables()->GetNumStencils())
+            return 0;
+
+        context->BindControlData( controlDataDesc, controlVertices );
+
+        context->BindOutputDerivData( outputDuDesc, outputDuData, outputDvDesc, outputDvData );
+
+        int n = _UpdateDerivs( context );
+
+        context->Unbind();
+
+        return n;
+    }
+
+    /// Waits until all running subdivision kernels finish.
+    void Synchronize();
+
+private:
+
+    int _UpdateValues( OsdCpuEvalStencilsContext * context );
+    int _UpdateDerivs( OsdCpuEvalStencilsContext * context );
+
+    int _numThreads;
+};
+
+} // end namespace OPENSUBDIV_VERSION
+using namespace OPENSUBDIV_VERSION;
+
+} // end namespace OpenSubdiv
+
+#endif // FAR_TBB_EVALSTENCILS_CONTROLLER_H

opensubdiv/osd/vertexDescriptor.h

@@ -191,16 +191,15 @@ struct OsdVertexBufferDescriptor {
 
     /// True if the descriptor values are internally consistent
     bool IsValid() const {
-        return (length>0) and (offset<length) and (stride>=length);
+        return ((length>0) and (offset<stride) and (length<=stride-offset));
     }
 
     /// True if the 'other' descriptor can be used as a destination for
     /// data evaluations.
     bool CanEval( OsdVertexBufferDescriptor const & other ) const {
-        return IsValid() and 
+        return (IsValid() and
                 other.IsValid() and
-               (length==other.length) and 
-               (other.length <= (stride-offset));
+                (length==other.length));
     }
 
     /// Resets the descriptor to default

opensubdiv/version.h

@@ -26,6 +26,6 @@
 #ifndef OPENSUBDIV_VERSION_H
 #define OPENSUBDIV_VERSION_H
 
-#define OPENSUBDIV_VERSION v2_1_0
+#define OPENSUBDIV_VERSION v2_2_0
 
 #endif /* OPENSUBDIV_VERSION_H */

regression/osd_regression/CMakeLists.txt

@@ -51,7 +51,7 @@ if ( OPENCL_FOUND )
 endif()
 
 
-add_executable(osd_regression
+_add_possibly_cuda_executable(osd_regression
     ${SOURCE_FILES}
 )