diff --git a/opensubdiv/far/CMakeLists.txt b/opensubdiv/far/CMakeLists.txt index d5c06333..54533774 100644 --- a/opensubdiv/far/CMakeLists.txt +++ b/opensubdiv/far/CMakeLists.txt @@ -62,6 +62,7 @@ set(PUBLIC_HEADER_FILES catmarkSubdivisionTablesFactory.h dispatcher.h kernelBatch.h + kernelBatchFactory.h loopSubdivisionTables.h loopSubdivisionTablesFactory.h meshFactory.h diff --git a/opensubdiv/far/bilinearSubdivisionTablesFactory.h b/opensubdiv/far/bilinearSubdivisionTablesFactory.h index ddc3706a..70d0f57b 100644 --- a/opensubdiv/far/bilinearSubdivisionTablesFactory.h +++ b/opensubdiv/far/bilinearSubdivisionTablesFactory.h @@ -81,6 +81,14 @@ protected: template friend class FarMeshFactory; /// Creates a FarBilinearSubdivisiontables instance. + /// + /// @param meshFactory a valid FarMeshFactory instance + /// + /// @param farMesh + /// + /// @param batches a vector of Kernel refinement batches : the factory + /// will reserve and append refinement tasks + /// static FarBilinearSubdivisionTables * Create( FarMeshFactory * meshFactory, FarMesh * farMesh, FarKernelBatchVector *batches ); }; @@ -134,13 +142,13 @@ FarBilinearSubdivisionTablesFactory::Create( FarMeshFactory * meshFact // "For each vertex, gather all the vertices from the parent face." int nFaceVertices = (int)tablesFactory._faceVertsList[level].size(); if (nFaceVertices > 0) - batches->push_back(FarKernelBatch(level, - BILINEAR_FACE_VERTEX, - /*tableIndex=*/0, - /*start=*/0, - /*end=*/nFaceVertices, - faceTableOffset, - vertexOffset)); + batches->push_back(FarKernelBatch( FarKernelBatch::BILINEAR_FACE_VERTEX, + level, + /*tableIndex=*/0, + /*start=*/0, + /*end=*/nFaceVertices, + faceTableOffset, + vertexOffset) ); vertexOffset += nFaceVertices; faceTableOffset += nFaceVertices; @@ -167,13 +175,13 @@ FarBilinearSubdivisionTablesFactory::Create( FarMeshFactory * meshFact // "Average the end-points of the parent edge" int nEdgeVertices = (int)tablesFactory._edgeVertsList[level].size(); if (nEdgeVertices > 0) - batches->push_back(FarKernelBatch(level, - BILINEAR_EDGE_VERTEX, - /*tableIndex=*/0, - /*start=*/0, - /*end=*/nEdgeVertices, - edgeTableOffset, - vertexOffset)); + batches->push_back(FarKernelBatch( FarKernelBatch::BILINEAR_EDGE_VERTEX, + level, + /*tableIndex=*/0, + /*start=*/0, + /*end=*/nEdgeVertices, + edgeTableOffset, + vertexOffset) ); vertexOffset += nEdgeVertices; edgeTableOffset += nEdgeVertices; for (int i=0; i < nEdgeVertices; ++i) { @@ -194,13 +202,13 @@ FarBilinearSubdivisionTablesFactory::Create( FarMeshFactory * meshFact // "Pass down the parent vertex" int nVertVertices = (int)tablesFactory._vertVertsList[level].size(); - batches->push_back(FarKernelBatch(level, - BILINEAR_VERT_VERTEX, - /*tableIndex=*/0, - /*start=*/0, - /*end=*/nVertVertices, - vertTableOffset, - vertexOffset)); + batches->push_back(FarKernelBatch( FarKernelBatch::BILINEAR_VERT_VERTEX, + level, + /*tableIndex=*/0, + /*start=*/0, + /*end=*/nVertVertices, + vertTableOffset, + vertexOffset)); vertexOffset += nVertVertices; vertTableOffset += nVertVertices; diff --git a/opensubdiv/far/catmarkSubdivisionTablesFactory.h b/opensubdiv/far/catmarkSubdivisionTablesFactory.h index 67c5a630..cffced74 100644 --- a/opensubdiv/far/catmarkSubdivisionTablesFactory.h +++ b/opensubdiv/far/catmarkSubdivisionTablesFactory.h @@ -65,6 +65,7 @@ #include "../far/catmarkSubdivisionTables.h" #include "../far/meshFactory.h" +#include "../far/kernelBatchFactory.h" #include "../far/subdivisionTablesFactory.h" namespace OpenSubdiv { @@ -81,6 +82,14 @@ protected: template friend class FarMeshFactory; /// Creates a FarCatmarkSubdivisiontables instance. + /// + /// @param meshFactory a valid FarMeshFactory instance + /// + /// @param farMesh + /// + /// @param batches a vector of Kernel refinement batches : the factory + /// will reserve and append refinement tasks + /// static FarCatmarkSubdivisionTables * Create( FarMeshFactory * meshFactory, FarMesh * farMesh, FarKernelBatchVector *batches ); }; @@ -144,13 +153,13 @@ FarCatmarkSubdivisionTablesFactory::Create( FarMeshFactory * meshFacto // add a batch for face vertices if (nFaceVertices > 0) // in torus case, nfacevertices could be zero - batches->push_back(FarKernelBatch(level, - CATMARK_FACE_VERTEX, - 0, - 0, - nFaceVertices, - faceTableOffset, - vertexOffset)); + batches->push_back(FarKernelBatch( FarKernelBatch::CATMARK_FACE_VERTEX, + level, + 0, + 0, + nFaceVertices, + faceTableOffset, + vertexOffset) ); vertexOffset += nFaceVertices; faceTableOffset += nFaceVertices; @@ -187,13 +196,13 @@ FarCatmarkSubdivisionTablesFactory::Create( FarMeshFactory * meshFacto // add a batch for edge vertices if (nEdgeVertices > 0) - batches->push_back(FarKernelBatch(level, - CATMARK_EDGE_VERTEX, - 0, - 0, - nEdgeVertices, - edgeTableOffset, - vertexOffset)); + batches->push_back(FarKernelBatch( FarKernelBatch::CATMARK_EDGE_VERTEX, + level, + 0, + 0, + nEdgeVertices, + edgeTableOffset, + vertexOffset) ); vertexOffset += nEdgeVertices; edgeTableOffset += nEdgeVertices; @@ -356,7 +365,7 @@ FarCatmarkSubdivisionTablesFactory::Create( FarMeshFactory * meshFacto // add batches for vert vertices if (nVertVertices > 0) - batchFactory.AppendCatmarkBatches(batches, level, vertTableOffset, vertexOffset); + batchFactory.AppendCatmarkBatches(level, vertTableOffset, vertexOffset, batches); vertexOffset += nVertVertices; vertTableOffset += nVertVertices; } diff --git a/opensubdiv/far/dispatcher.h b/opensubdiv/far/dispatcher.h index 7f735be9..a463b6f0 100644 --- a/opensubdiv/far/dispatcher.h +++ b/opensubdiv/far/dispatcher.h @@ -97,51 +97,51 @@ FarDispatcher::Refine(CONTROLLER const *controller, FarKernelBatchVector const & for (int i = 0; i < (int)batches.size(); ++i) { const FarKernelBatch &batch = batches[i]; - if (maxlevel >= 0 && batch.level >= maxlevel) continue; + if (maxlevel >= 0 && batch.GetLevel() >= maxlevel) continue; - switch(batch.kernelType) { - case CATMARK_FACE_VERTEX: - controller->ApplyCatmarkFaceVerticesKernel(batch, clientdata); - break; - case CATMARK_EDGE_VERTEX: - controller->ApplyCatmarkEdgeVerticesKernel(batch, clientdata); - break; - case CATMARK_VERT_VERTEX_B: - controller->ApplyCatmarkVertexVerticesKernelB(batch, clientdata); - break; - case CATMARK_VERT_VERTEX_A1: - controller->ApplyCatmarkVertexVerticesKernelA1(batch, clientdata); - break; - case CATMARK_VERT_VERTEX_A2: - controller->ApplyCatmarkVertexVerticesKernelA2(batch, clientdata); - break; + switch(batch.GetKernelType()) { + case FarKernelBatch::CATMARK_FACE_VERTEX: + controller->ApplyCatmarkFaceVerticesKernel(batch, clientdata); + break; + case FarKernelBatch::CATMARK_EDGE_VERTEX: + controller->ApplyCatmarkEdgeVerticesKernel(batch, clientdata); + break; + case FarKernelBatch::CATMARK_VERT_VERTEX_B: + controller->ApplyCatmarkVertexVerticesKernelB(batch, clientdata); + break; + case FarKernelBatch::CATMARK_VERT_VERTEX_A1: + controller->ApplyCatmarkVertexVerticesKernelA1(batch, clientdata); + break; + case FarKernelBatch::CATMARK_VERT_VERTEX_A2: + controller->ApplyCatmarkVertexVerticesKernelA2(batch, clientdata); + break; - case LOOP_EDGE_VERTEX: - controller->ApplyLoopEdgeVerticesKernel(batch, clientdata); - break; - case LOOP_VERT_VERTEX_B: - controller->ApplyLoopVertexVerticesKernelB(batch, clientdata); - break; - case LOOP_VERT_VERTEX_A1: - controller->ApplyLoopVertexVerticesKernelA1(batch, clientdata); - break; - case LOOP_VERT_VERTEX_A2: - controller->ApplyLoopVertexVerticesKernelA2(batch, clientdata); - break; + case FarKernelBatch::LOOP_EDGE_VERTEX: + controller->ApplyLoopEdgeVerticesKernel(batch, clientdata); + break; + case FarKernelBatch::LOOP_VERT_VERTEX_B: + controller->ApplyLoopVertexVerticesKernelB(batch, clientdata); + break; + case FarKernelBatch::LOOP_VERT_VERTEX_A1: + controller->ApplyLoopVertexVerticesKernelA1(batch, clientdata); + break; + case FarKernelBatch::LOOP_VERT_VERTEX_A2: + controller->ApplyLoopVertexVerticesKernelA2(batch, clientdata); + break; - case BILINEAR_FACE_VERTEX: - controller->ApplyBilinearFaceVerticesKernel(batch, clientdata); - break; - case BILINEAR_EDGE_VERTEX: - controller->ApplyBilinearEdgeVerticesKernel(batch, clientdata); - break; - case BILINEAR_VERT_VERTEX: - controller->ApplyBilinearVertexVerticesKernel(batch, clientdata); - break; + case FarKernelBatch::BILINEAR_FACE_VERTEX: + controller->ApplyBilinearFaceVerticesKernel(batch, clientdata); + break; + case FarKernelBatch::BILINEAR_EDGE_VERTEX: + controller->ApplyBilinearEdgeVerticesKernel(batch, clientdata); + break; + case FarKernelBatch::BILINEAR_VERT_VERTEX: + controller->ApplyBilinearVertexVerticesKernel(batch, clientdata); + break; - case HIERARCHICAL_EDIT: - controller->ApplyVertexEdits(batch, clientdata); - break; + case FarKernelBatch::HIERARCHICAL_EDIT: + controller->ApplyVertexEdits(batch, clientdata); + break; } } } @@ -152,6 +152,7 @@ FarDispatcher::Refine(CONTROLLER const *controller, FarKernelBatchVector const & template class FarComputeController { + public: void Refine(FarMesh * mesh, int maxlevel=-1) const; @@ -183,6 +184,9 @@ public: void ApplyVertexEdits(FarKernelBatch const &batch, void * clientdata) const; static FarComputeController _DefaultController; + +private: + }; template FarComputeController FarComputeController::_DefaultController; @@ -195,118 +199,224 @@ FarComputeController::Refine(FarMesh *mesh, int maxlevel) const { template void FarComputeController::ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarBilinearSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); assert(subdivision); - subdivision->computeFacePoints(batch.vertexOffset, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeFacePoints( batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarBilinearSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); assert(subdivision); - subdivision->computeEdgePoints(batch.vertexOffset, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeEdgePoints( batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarBilinearSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); + assert(subdivision); - subdivision->computeVertexPoints(batch.vertexOffset, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeVertexPoints( batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarCatmarkSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); + assert(subdivision); - subdivision->computeFacePoints(batch.vertexOffset, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeFacePoints( batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarCatmarkSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); + assert(subdivision); - subdivision->computeEdgePoints(batch.vertexOffset, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeEdgePoints( batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarCatmarkSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); + assert(subdivision); - subdivision->computeVertexPointsB(batch.vertexOffset, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeVertexPointsB( batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarCatmarkSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); + assert(subdivision); - subdivision->computeVertexPointsA(batch.vertexOffset, false, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeVertexPointsA( batch.GetVertexOffset(), + false, + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarCatmarkSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); + assert(subdivision); - subdivision->computeVertexPointsA(batch.vertexOffset, true, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeVertexPointsA( batch.GetVertexOffset(), + true, + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarLoopSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); + assert(subdivision); - subdivision->computeEdgePoints(batch.vertexOffset, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeEdgePoints( batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarLoopSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); + assert(subdivision); - subdivision->computeVertexPointsB(batch.vertexOffset, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeVertexPointsB( batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarLoopSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); + assert(subdivision); - subdivision->computeVertexPointsA(batch.vertexOffset, false, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeVertexPointsA( batch.GetVertexOffset(), + false, + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } template void FarComputeController::ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarLoopSubdivisionTables const * subdivision = dynamic_cast const *>(mesh->GetSubdivisionTables()); + assert(subdivision); - subdivision->computeVertexPointsA(batch.vertexOffset, true, batch.tableOffset, batch.start, batch.end, clientdata); + + subdivision->computeVertexPointsA( batch.GetVertexOffset(), + true, + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata); } template void FarComputeController::ApplyVertexEdits(FarKernelBatch const &batch, void * clientdata) const { + FarMesh * mesh = static_cast *>(clientdata); + FarVertexEditTables const * vertEdit = mesh->GetVertexEdit(); + if (vertEdit) - vertEdit->computeVertexEdits(batch.tableIndex, batch.vertexOffset, batch.tableOffset, batch.start, batch.end, clientdata); + vertEdit->computeVertexEdits( batch.GetTableIndex(), + batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), + clientdata ); } } // end namespace OPENSUBDIV_VERSION diff --git a/opensubdiv/far/kernelBatch.h b/opensubdiv/far/kernelBatch.h index 3d2b850e..c4a69b31 100644 --- a/opensubdiv/far/kernelBatch.h +++ b/opensubdiv/far/kernelBatch.h @@ -64,136 +64,167 @@ namespace OpenSubdiv { namespace OPENSUBDIV_VERSION { -enum KernelType { - CATMARK_FACE_VERTEX, - CATMARK_EDGE_VERTEX, - CATMARK_VERT_VERTEX_A1, - CATMARK_VERT_VERTEX_A2, - CATMARK_VERT_VERTEX_B, - LOOP_EDGE_VERTEX, - LOOP_VERT_VERTEX_A1, - LOOP_VERT_VERTEX_A2, - LOOP_VERT_VERTEX_B, - BILINEAR_FACE_VERTEX, - BILINEAR_EDGE_VERTEX, - BILINEAR_VERT_VERTEX, - HIERARCHICAL_EDIT, -}; -struct FarKernelBatch { - FarKernelBatch(int level_, - int kernelType_, - int tableIndex_, - int start_, - int end_, - int tableOffset_, - int vertexOffset_) : - level(level_), - kernelType(kernelType_), - tableIndex(tableIndex_), - start(start_), - end(end_), - tableOffset(tableOffset_), - vertexOffset(vertexOffset_) { +/// \brief A GP Compute Kernel descriptor. +/// +/// Vertex refinement through subdivision schemes requires the successive +/// application of dedicated compute kernels. OpenSubdiv groups these vertices +/// in batches based on their topology in order to minimize the number of kernel +/// switches to process a given primitive. +/// +/// [Subdivision table for kernel k] +/// ------+---------------------------------+----- +/// | Prim p, Level n | +/// ------+---------------------------------+----- +/// | || | +/// ------+---------------------------------+----- +/// ^ ^ ^ +/// tableOffset start end +/// . . +/// . . +/// . . +/// [VBO] . . +/// ------+---------------------------------+----- +/// | Prim p, Kernel k, Level n | +/// ------+---------------------------------+----- +/// | || | +/// ------+---------------------------------+----- +/// ^ ^ ^ +/// vertexOffset start end +/// +class FarKernelBatch { + +public: + + enum KernelType { + CATMARK_FACE_VERTEX, + CATMARK_EDGE_VERTEX, + CATMARK_VERT_VERTEX_A1, + CATMARK_VERT_VERTEX_A2, + CATMARK_VERT_VERTEX_B, + LOOP_EDGE_VERTEX, + LOOP_VERT_VERTEX_A1, + LOOP_VERT_VERTEX_A2, + LOOP_VERT_VERTEX_B, + BILINEAR_FACE_VERTEX, + BILINEAR_EDGE_VERTEX, + BILINEAR_VERT_VERTEX, + HIERARCHICAL_EDIT, + }; + + /// Constructor. + /// + /// @param kernelType the type of compute kernel kernel + /// + /// @param level the level of subdivision of the vertices in the batch + /// + /// @param tableIndex edit index (for the hierarchical edit kernels only) + /// + /// @param start index of the first vertex in the batch + /// + /// @param end index of the last vertex in the batch + /// + /// @param tableOffset XXXX + /// + /// @param vertexOffset XXXX + /// + FarKernelBatch( KernelType kernelType, + int level, + int tableIndex, + int start, + int end, + int tableOffset, + int vertexOffset) : + _kernelType(kernelType), + _level(level), + _tableIndex(tableIndex), + _start(start), + _end(end), + _tableOffset(tableOffset), + _vertexOffset(vertexOffset) { } - /* - [Subdivision table for kernel k] - ------+---------------------------------+----- - | Prim p, Level n | - ------+---------------------------------+----- - | || | - ------+---------------------------------+----- - ^ ^ ^ - tableOffset start end - . . - . . - . . - [VBO] . . - ------+---------------------------------+----- - | Prim p, Kernel k, Level n | - ------+---------------------------------+----- - | || | - ------+---------------------------------+----- - ^ ^ ^ - vertexOffset start end - */ + /// Returns the type of kernel to apply to the vertices in the batch. + KernelType GetKernelType() const { + return _kernelType; + } - int level; - int kernelType; - int tableIndex; // edit index (for the h-edit kernel only) - int start; - int end; - int tableOffset; - int vertexOffset; + /// Returns the subdivision level of the vertices in the batch + int GetLevel() const { + return _level; + } + + + + /// Returns the index of the first vertex in the batch + int GetStart() const { + return _start; + } + + /// Returns the index of the first vertex in the batch + const int * GetStartPtr() const { + return & _start; + } + + + + /// Returns the index of the last vertex in the batch + int GetEnd() const { + return _end; + } + + /// Returns the index of the last vertex in the batch + const int * GetEndPtr() const { + return & _end; + } + + + + /// Returns the edit index (for the hierarchical edit kernels only) + int GetTableIndex() const { + return _tableIndex; + } + + + + /// Returns + int GetTableOffset() const { + return _tableOffset; + } + + /// Returns + const int * GetTableOffsetPtr() const { + return & _tableOffset; + } + + + + /// Returns + int GetVertexOffset() const { + return _vertexOffset; + } + + /// Returns + const int * GetVertexOffsetPtr() const { + return & _vertexOffset; + } + +private: + friend class FarKernelBatchFactory; + template friend class FarMultiMeshFactory; + + KernelType _kernelType; + int _level; + int _tableIndex; + int _start; + int _end; + int _tableOffset; + int _vertexOffset; }; typedef std::vector FarKernelBatchVector; -struct FarVertexKernelBatchFactory { - FarVertexKernelBatchFactory(int a, int b) { - kernelB.first = kernelA1.first = kernelA2.first = a; - kernelB.second = kernelA1.second = kernelA2.second = b; - } - - void AddVertex( int index, int rank ) { - // expand the range of kernel batches based on vertex index and rank - if (rank<7) { - if (index < kernelB.first) - kernelB.first=index; - if (index > kernelB.second) - kernelB.second=index; - } - if ((rank>2) and (rank<8)) { - if (index < kernelA2.first) - kernelA2.first=index; - if (index > kernelA2.second) - kernelA2.second=index; - } - if (rank>6) { - if (index < kernelA1.first) - kernelA1.first=index; - if (index > kernelA1.second) - kernelA1.second=index; - } - } - - void AppendCatmarkBatches(FarKernelBatchVector *result, int level, int tableOffset, int vertexOffset) { - if (kernelB.second >= kernelB.first) - result->push_back(FarKernelBatch(level, CATMARK_VERT_VERTEX_B, 0, - kernelB.first, kernelB.second+1, - tableOffset, vertexOffset)); - if (kernelA1.second >= kernelA1.first) - result->push_back(FarKernelBatch(level, CATMARK_VERT_VERTEX_A1, 0, - kernelA1.first, kernelA1.second+1, - tableOffset, vertexOffset)); - if (kernelA2.second >= kernelA2.first) - result->push_back(FarKernelBatch(level, CATMARK_VERT_VERTEX_A2, 0, - kernelA2.first, kernelA2.second+1, - tableOffset, vertexOffset)); - } - - void AppendLoopBatches(FarKernelBatchVector *result, int level, int tableOffset, int vertexOffset) { - if (kernelB.second >= kernelB.first) - result->push_back(FarKernelBatch(level, LOOP_VERT_VERTEX_B, 0, - kernelB.first, kernelB.second+1, - tableOffset, vertexOffset)); - if (kernelA1.second >= kernelA1.first) - result->push_back(FarKernelBatch(level, LOOP_VERT_VERTEX_A1, 0, - kernelA1.first, kernelA1.second+1, - tableOffset, vertexOffset)); - if (kernelA2.second >= kernelA2.first) - result->push_back(FarKernelBatch(level, LOOP_VERT_VERTEX_A2, 0, - kernelA2.first, kernelA2.second+1, - tableOffset, vertexOffset)); - } - - std::pair kernelB; // first / last vertex vertex batch (kernel B) - std::pair kernelA1; // first / last vertex vertex batch (kernel A pass 1) - std::pair kernelA2; // first / last vertex vertex batch (kernel A pass 2) -}; - } // end namespace OPENSUBDIV_VERSION using namespace OPENSUBDIV_VERSION; diff --git a/opensubdiv/far/kernelBatchFactory.h b/opensubdiv/far/kernelBatchFactory.h new file mode 100644 index 00000000..413e3f9f --- /dev/null +++ b/opensubdiv/far/kernelBatchFactory.h @@ -0,0 +1,219 @@ +// +// Copyright (C) Pixar. All rights reserved. +// +// This license governs use of the accompanying software. If you +// use the software, you accept this license. If you do not accept +// the license, do not use the software. +// +// 1. Definitions +// The terms "reproduce," "reproduction," "derivative works," and +// "distribution" have the same meaning here as under U.S. +// copyright law. A "contribution" is the original software, or +// any additions or changes to the software. +// A "contributor" is any person or entity that distributes its +// contribution under this license. +// "Licensed patents" are a contributor's patent claims that read +// directly on its contribution. +// +// 2. Grant of Rights +// (A) Copyright Grant- Subject to the terms of this license, +// including the license conditions and limitations in section 3, +// each contributor grants you a non-exclusive, worldwide, +// royalty-free copyright license to reproduce its contribution, +// prepare derivative works of its contribution, and distribute +// its contribution or any derivative works that you create. +// (B) Patent Grant- Subject to the terms of this license, +// including the license conditions and limitations in section 3, +// each contributor grants you a non-exclusive, worldwide, +// royalty-free license under its licensed patents to make, have +// made, use, sell, offer for sale, import, and/or otherwise +// dispose of its contribution in the software or derivative works +// of the contribution in the software. +// +// 3. Conditions and Limitations +// (A) No Trademark License- This license does not grant you +// rights to use any contributor's name, logo, or trademarks. +// (B) If you bring a patent claim against any contributor over +// patents that you claim are infringed by the software, your +// patent license from such contributor to the software ends +// automatically. +// (C) If you distribute any portion of the software, you must +// retain all copyright, patent, trademark, and attribution +// notices that are present in the software. +// (D) If you distribute any portion of the software in source +// code form, you may do so only under this license by including a +// complete copy of this license with your distribution. If you +// distribute any portion of the software in compiled or object +// code form, you may only do so under a license that complies +// with this license. +// (E) The software is licensed "as-is." You bear the risk of +// using it. The contributors give no express warranties, +// guarantees or conditions. You may have additional consumer +// rights under your local laws which this license cannot change. +// To the extent permitted under your local laws, the contributors +// exclude the implied warranties of merchantability, fitness for +// a particular purpose and non-infringement. +// +#ifndef FAR_KERNEL_BATCH_FACTORY_H +#define FAR_KERNEL_BATCH_FACTORY_H + +#include "../version.h" + +#include "../far/kernelBatch.h" + +namespace OpenSubdiv { +namespace OPENSUBDIV_VERSION { + + +class FarVertexKernelBatchFactory { + +public: + + /// Constructor. + /// + /// @param start index of the first vertex in the batch + /// + /// @param end index of the last vertex in the batch + /// + FarVertexKernelBatchFactory(int start, int end) { + kernelB.start = kernelA1.start = kernelA2.start = start; + kernelB.end = kernelA1.end = kernelA2.end = end; + } + + + /// Adds a vertex-vertex to the appropriate compute batch based on "Rank". + /// Ranking is based on the interpolation required (Smooth, Dart, Crease, + /// or Corner). With semi-sharp creases, two passes of interpolation are + /// required, from which we derive a matrix of compute kernel combinations. + /// + /// The kernel combinatorial matrix : + /// + /// Rules +----+----+----+----+----+----+----+----+----+----+ + /// Pass 0 | Dt | Sm | Sm | Dt | Sm | Dt | Sm | Cr | Co | Cr | + /// Pass 1 | | | | Co | Co | Cr | Cr | Co | | | + /// Kernel +----+----+----+----+----+----+----+----+----+----+ + /// Pass 0 | B | B | B | B | B | B | B | A | A | A | + /// Pass 1 | | | | A | A | A | A | A | | | + /// +----+----+----+----+----+----+----+----+----+----+ + /// Rank | 0 | 1 | 2 | 3 | 4 | 5 | 6 | 7 | 8 | 9 | + /// +----+----+----+----+----+----+----+----+----+----+ + /// with : + /// - A : compute kernel applying k_Crease / k_Corner rules + /// - B : compute kernel applying k_Smooth / k_Dart rules + /// + /// @param index the index of the vertex + /// + /// @param rank the rank of the vertex (see + /// FarSubdivisionTables::GetMaskRanking()) + /// + void AddVertex( int index, int rank ); + + + + /// Appends a FarKernelBatch to a vector of batches for Catmark subdivision + /// + /// @param level the subdivision level of the vertices in the batch + /// + /// @param tableOffset XXXX + /// + /// @param vertexOffset XXXX + /// + /// @result the expanded batch vector + /// + void AppendCatmarkBatches(int level, int tableOffset, int vertexOffset, FarKernelBatchVector *result); + + + + /// Appends a FarKernelBatch to a vector of batches for Loop subdivision + /// + /// @param level the subdivision level of the vertices in the batch + /// + /// @param tableOffset XXXX + /// + /// @param vertexOffset XXXX + /// + /// @result the expanded batch vector + /// + void AppendLoopBatches(int level, int tableOffset, int vertexOffset, FarKernelBatchVector *result); + +private: + + struct Range { + int start, + end; + }; + + Range kernelB; // vertex batch reange (kernel B) + Range kernelA1; // vertex batch reange (kernel A pass 1) + Range kernelA2; // vertex batch reange (kernel A pass 2) +}; + +inline void +FarVertexKernelBatchFactory::AddVertex( int index, int rank ) { + + // expand the range of kernel batches based on vertex index and rank + if (rank<7) { + if (index < kernelB.start) + kernelB.start=index; + if (index > kernelB.end) + kernelB.end=index; + } + if ((rank>2) and (rank<8)) { + if (index < kernelA2.start) + kernelA2.start=index; + if (index > kernelA2.end) + kernelA2.end=index; + } + if (rank>6) { + if (index < kernelA1.start) + kernelA1.start=index; + if (index > kernelA1.end) + kernelA1.end=index; + } +} + +inline void +FarVertexKernelBatchFactory::AppendCatmarkBatches(int level, + int tableOffset, + int vertexOffset, + FarKernelBatchVector *result) { + + if (kernelB.end >= kernelB.start) + result->push_back(FarKernelBatch( FarKernelBatch::CATMARK_VERT_VERTEX_B, level, 0, + kernelB.start, kernelB.end+1, + tableOffset, vertexOffset) ); + if (kernelA1.end >= kernelA1.start) + result->push_back(FarKernelBatch( FarKernelBatch::CATMARK_VERT_VERTEX_A1, level, 0, + kernelA1.start, kernelA1.end+1, + tableOffset, vertexOffset)); + if (kernelA2.end >= kernelA2.start) + result->push_back(FarKernelBatch( FarKernelBatch::CATMARK_VERT_VERTEX_A2, level, 0, + kernelA2.start, kernelA2.end+1, + tableOffset, vertexOffset) ); +} + +inline void +FarVertexKernelBatchFactory::AppendLoopBatches(int level, + int tableOffset, + int vertexOffset, + FarKernelBatchVector *result) { + if (kernelB.end >= kernelB.start) + result->push_back(FarKernelBatch( FarKernelBatch::LOOP_VERT_VERTEX_B, level, 0, + kernelB.start, kernelB.end+1, + tableOffset, vertexOffset) ); + if (kernelA1.end >= kernelA1.start) + result->push_back(FarKernelBatch( FarKernelBatch::LOOP_VERT_VERTEX_A1, level, 0, + kernelA1.start, kernelA1.end+1, + tableOffset, vertexOffset) ); + if (kernelA2.end >= kernelA2.start) + result->push_back(FarKernelBatch( FarKernelBatch::LOOP_VERT_VERTEX_A2, level, 0, + kernelA2.start, kernelA2.end+1, + tableOffset, vertexOffset) ); +} + +} // end namespace OPENSUBDIV_VERSION +using namespace OPENSUBDIV_VERSION; + +} // end namespace OpenSubdiv + +#endif /* FAR_KERNEL_BATCH_FACTORY_H */ diff --git a/opensubdiv/far/loopSubdivisionTablesFactory.h b/opensubdiv/far/loopSubdivisionTablesFactory.h index 8e21d545..1b4065e3 100644 --- a/opensubdiv/far/loopSubdivisionTablesFactory.h +++ b/opensubdiv/far/loopSubdivisionTablesFactory.h @@ -62,6 +62,7 @@ #include "../far/loopSubdivisionTables.h" #include "../far/meshFactory.h" +#include "../far/kernelBatchFactory.h" #include "../far/subdivisionTablesFactory.h" #include @@ -81,6 +82,14 @@ protected: template friend class FarMeshFactory; /// Creates a FarLoopSubdivisiontables instance. + /// + /// @param meshFactory a valid FarMeshFactory instance + /// + /// @param farMesh + /// + /// @param batches a vector of Kernel refinement batches : the factory + /// will reserve and append refinement tasks + /// static FarLoopSubdivisionTables * Create( FarMeshFactory * meshFactory, FarMesh * farMesh, FarKernelBatchVector * batches ); }; @@ -134,13 +143,13 @@ FarLoopSubdivisionTablesFactory::Create( FarMeshFactory * meshFactory, // Edge vertices int nEdgeVertices = (int)tablesFactory._edgeVertsList[level].size(); if (nEdgeVertices > 0) - batches->push_back(FarKernelBatch(level, - LOOP_EDGE_VERTEX, - 0, - 0, - nEdgeVertices, - edgeTableOffset, - vertexOffset)); + batches->push_back(FarKernelBatch( FarKernelBatch::LOOP_EDGE_VERTEX, + level, + 0, + 0, + nEdgeVertices, + edgeTableOffset, + vertexOffset) ); vertexOffset += nEdgeVertices; edgeTableOffset += nEdgeVertices; @@ -286,7 +295,7 @@ FarLoopSubdivisionTablesFactory::Create( FarMeshFactory * meshFactory, V_ITa += nVertVertices*5; V_W += nVertVertices; - batchFactory.AppendLoopBatches(batches, level, vertTableOffset, vertexOffset); + batchFactory.AppendLoopBatches(level, vertTableOffset, vertexOffset, batches); vertexOffset += nVertVertices; vertTableOffset += nVertVertices; } diff --git a/opensubdiv/far/multiMeshFactory.h b/opensubdiv/far/multiMeshFactory.h index 08ae6397..ba4d5b74 100644 --- a/opensubdiv/far/multiMeshFactory.h +++ b/opensubdiv/far/multiMeshFactory.h @@ -332,25 +332,32 @@ FarMultiMeshFactory::spliceSubdivisionTables(FarMesh *farMesh, FarMeshV for (size_t i = 0; i < meshes.size(); ++i) { for (int j = 0; j < (int)meshes[i]->_batches.size(); ++j) { FarKernelBatch batch = meshes[i]->_batches[j]; - batch.vertexOffset += vertexOffsets[i]; + batch._vertexOffset += vertexOffsets[i]; - if (batch.kernelType == CATMARK_FACE_VERTEX or - batch.kernelType == BILINEAR_FACE_VERTEX) { - batch.tableOffset += fvOffsets[i]; - } else if (batch.kernelType == CATMARK_EDGE_VERTEX or - batch.kernelType == LOOP_EDGE_VERTEX or - batch.kernelType == BILINEAR_EDGE_VERTEX) { - batch.tableOffset += evOffsets[i]; - } else if (batch.kernelType == CATMARK_VERT_VERTEX_A1 or - batch.kernelType == CATMARK_VERT_VERTEX_A2 or - batch.kernelType == CATMARK_VERT_VERTEX_B or - batch.kernelType == LOOP_VERT_VERTEX_A1 or - batch.kernelType == LOOP_VERT_VERTEX_A2 or - batch.kernelType == LOOP_VERT_VERTEX_B or - batch.kernelType == BILINEAR_VERT_VERTEX) { - batch.tableOffset += vvOffsets[i]; - } else if (batch.kernelType == HIERARCHICAL_EDIT) { - batch.tableIndex += editTableIndexOffset; + if (batch._kernelType == FarKernelBatch::CATMARK_FACE_VERTEX or + batch._kernelType == FarKernelBatch::BILINEAR_FACE_VERTEX) { + + batch._tableOffset += fvOffsets[i]; + + } else if (batch._kernelType == FarKernelBatch::CATMARK_EDGE_VERTEX or + batch._kernelType == FarKernelBatch::LOOP_EDGE_VERTEX or + batch._kernelType == FarKernelBatch::BILINEAR_EDGE_VERTEX) { + + batch._tableOffset += evOffsets[i]; + + } else if (batch._kernelType == FarKernelBatch::CATMARK_VERT_VERTEX_A1 or + batch._kernelType == FarKernelBatch::CATMARK_VERT_VERTEX_A2 or + batch._kernelType == FarKernelBatch::CATMARK_VERT_VERTEX_B or + batch._kernelType == FarKernelBatch::LOOP_VERT_VERTEX_A1 or + batch._kernelType == FarKernelBatch::LOOP_VERT_VERTEX_A2 or + batch._kernelType == FarKernelBatch::LOOP_VERT_VERTEX_B or + batch._kernelType == FarKernelBatch::BILINEAR_VERT_VERTEX) { + + batch._tableOffset += vvOffsets[i]; + + } else if (batch._kernelType == FarKernelBatch::HIERARCHICAL_EDIT) { + + batch._tableIndex += editTableIndexOffset; } batches.push_back(batch); } diff --git a/opensubdiv/far/vertexEditTablesFactory.h b/opensubdiv/far/vertexEditTablesFactory.h index 9482b72c..d5474b1c 100644 --- a/opensubdiv/far/vertexEditTablesFactory.h +++ b/opensubdiv/far/vertexEditTablesFactory.h @@ -82,6 +82,7 @@ protected: /// Compares the number of subfaces in an edit (for sorting purposes) static bool compareEdits(HbrVertexEdit const *a, HbrVertexEdit const *b); + static void insertHEditBatch(FarKernelBatchVector *batches, int batchIndex, int batchLevel, int batchCount, int tableOffset); /// Creates a FarVertexEditTables instance. @@ -98,11 +99,14 @@ template void FarVertexEditTablesFactory::insertHEditBatch(FarKernelBatchVector *batches, int batchIndex, int batchLevel, int batchCount, int tableOffset) { FarKernelBatchVector::iterator it = batches->begin(); - while(it != batches->end()) { - if (it->level > batchLevel+1) break; + + while (it != batches->end()) { + if (it->GetLevel() > batchLevel+1) + break; ++it; } - batches->insert(it, FarKernelBatch(batchLevel+1, HIERARCHICAL_EDIT, batchIndex, 0, batchCount, tableOffset, 0)); + + batches->insert(it, FarKernelBatch( FarKernelBatch::HIERARCHICAL_EDIT, batchLevel+1, batchIndex, 0, batchCount, tableOffset, 0) ); } template FarVertexEditTables * diff --git a/opensubdiv/osd/clComputeController.cpp b/opensubdiv/osd/clComputeController.cpp index 3b52249f..0a2bf840 100644 --- a/opensubdiv/osd/clComputeController.cpp +++ b/opensubdiv/osd/clComputeController.cpp @@ -140,7 +140,7 @@ OsdCLComputeController::ApplyBilinearEdgeVerticesKernel( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; cl_kernel kernel = context->GetKernelBundle()->GetBilinearEdgeKernel(); cl_mem vertexBuffer = context->GetCurrentVertexBuffer(); @@ -150,10 +150,10 @@ OsdCLComputeController::ApplyBilinearEdgeVerticesKernel( clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer); clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer); clSetKernelArg(kernel, 2, sizeof(cl_mem), &E_IT); - clSetKernelArg(kernel, 3, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 4, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 5, sizeof(int), &batch.start); - clSetKernelArg(kernel, 6, sizeof(int), &batch.end); + clSetKernelArg(kernel, 3, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 4, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetEndPtr()); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), kernel, 1, NULL, globalWorkSize, NULL, 0, NULL, NULL); @@ -169,7 +169,7 @@ OsdCLComputeController::ApplyBilinearVertexVerticesKernel( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; cl_kernel kernel = context->GetKernelBundle()->GetBilinearVertexKernel(); cl_mem vertexBuffer = context->GetCurrentVertexBuffer(); @@ -179,10 +179,10 @@ OsdCLComputeController::ApplyBilinearVertexVerticesKernel( clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer); clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer); clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa); - clSetKernelArg(kernel, 3, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 4, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 5, sizeof(int), &batch.start); - clSetKernelArg(kernel, 6, sizeof(int), &batch.end); + clSetKernelArg(kernel, 3, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 4, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetEndPtr()); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), kernel, 1, NULL, globalWorkSize, NULL, 0, NULL, NULL); @@ -198,7 +198,7 @@ OsdCLComputeController::ApplyCatmarkFaceVerticesKernel( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; cl_kernel kernel = context->GetKernelBundle()->GetCatmarkFaceKernel(); cl_mem vertexBuffer = context->GetCurrentVertexBuffer(); @@ -210,10 +210,10 @@ OsdCLComputeController::ApplyCatmarkFaceVerticesKernel( clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer); clSetKernelArg(kernel, 2, sizeof(cl_mem), &F_IT); clSetKernelArg(kernel, 3, sizeof(cl_mem), &F_ITa); - clSetKernelArg(kernel, 4, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 5, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 6, sizeof(int), &batch.start); - clSetKernelArg(kernel, 7, sizeof(int), &batch.end); + clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr()); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), kernel, 1, NULL, globalWorkSize, @@ -230,7 +230,7 @@ OsdCLComputeController::ApplyCatmarkEdgeVerticesKernel( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; cl_kernel kernel = context->GetKernelBundle()->GetCatmarkEdgeKernel(); cl_mem vertexBuffer = context->GetCurrentVertexBuffer(); @@ -242,10 +242,10 @@ OsdCLComputeController::ApplyCatmarkEdgeVerticesKernel( clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer); clSetKernelArg(kernel, 2, sizeof(cl_mem), &E_IT); clSetKernelArg(kernel, 3, sizeof(cl_mem), &E_W); - clSetKernelArg(kernel, 4, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 5, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 6, sizeof(int), &batch.start); - clSetKernelArg(kernel, 7, sizeof(int), &batch.end); + clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr()); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), kernel, 1, NULL, globalWorkSize, @@ -262,7 +262,7 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelB( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; cl_kernel kernel = context->GetKernelBundle()->GetCatmarkVertexKernelB(); cl_mem vertexBuffer = context->GetCurrentVertexBuffer(); @@ -276,10 +276,10 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelB( clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa); clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_IT); clSetKernelArg(kernel, 4, sizeof(cl_mem), &V_W); - clSetKernelArg(kernel, 5, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 6, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 7, sizeof(int), &batch.start); - clSetKernelArg(kernel, 8, sizeof(int), &batch.end); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 7, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 8, sizeof(int), batch.GetEndPtr()); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), kernel, 1, NULL, globalWorkSize, @@ -296,7 +296,7 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelA1( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; int ipass = false; cl_kernel kernel = context->GetKernelBundle()->GetCatmarkVertexKernelA(); @@ -309,10 +309,10 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelA1( clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer); clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa); clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_W); - clSetKernelArg(kernel, 4, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 5, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 6, sizeof(int), &batch.start); - clSetKernelArg(kernel, 7, sizeof(int), &batch.end); + clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr()); clSetKernelArg(kernel, 8, sizeof(int), &ipass); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), @@ -330,7 +330,7 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelA2( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; int ipass = true; cl_kernel kernel = context->GetKernelBundle()->GetCatmarkVertexKernelA(); @@ -343,10 +343,10 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelA2( clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer); clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa); clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_W); - clSetKernelArg(kernel, 4, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 5, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 6, sizeof(int), &batch.start); - clSetKernelArg(kernel, 7, sizeof(int), &batch.end); + clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr()); clSetKernelArg(kernel, 8, sizeof(int), &ipass); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), @@ -364,7 +364,7 @@ OsdCLComputeController::ApplyLoopEdgeVerticesKernel( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; cl_kernel kernel = context->GetKernelBundle()->GetLoopEdgeKernel(); cl_mem vertexBuffer = context->GetCurrentVertexBuffer(); @@ -376,10 +376,10 @@ OsdCLComputeController::ApplyLoopEdgeVerticesKernel( clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer); clSetKernelArg(kernel, 2, sizeof(cl_mem), &E_IT); clSetKernelArg(kernel, 3, sizeof(cl_mem), &E_W); - clSetKernelArg(kernel, 4, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 5, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 6, sizeof(int), &batch.start); - clSetKernelArg(kernel, 7, sizeof(int), &batch.end); + clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr()); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), kernel, 1, NULL, globalWorkSize, @@ -396,7 +396,7 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelB( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; cl_kernel kernel = context->GetKernelBundle()->GetLoopVertexKernelB(); cl_mem vertexBuffer = context->GetCurrentVertexBuffer(); @@ -410,10 +410,10 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelB( clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa); clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_IT); clSetKernelArg(kernel, 4, sizeof(cl_mem), &V_W); - clSetKernelArg(kernel, 5, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 6, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 7, sizeof(int), &batch.start); - clSetKernelArg(kernel, 8, sizeof(int), &batch.end); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 7, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 8, sizeof(int), batch.GetEndPtr()); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), kernel, 1, NULL, globalWorkSize, @@ -430,7 +430,7 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelA1( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; int ipass = false; cl_kernel kernel = context->GetKernelBundle()->GetLoopVertexKernelA(); @@ -443,10 +443,10 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelA1( clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer); clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa); clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_W); - clSetKernelArg(kernel, 4, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 5, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 6, sizeof(int), &batch.start); - clSetKernelArg(kernel, 7, sizeof(int), &batch.end); + clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr()); clSetKernelArg(kernel, 8, sizeof(int), &ipass); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), @@ -464,7 +464,7 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelA2( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; int ipass = true; cl_kernel kernel = context->GetKernelBundle()->GetLoopVertexKernelA(); @@ -477,10 +477,10 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelA2( clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer); clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa); clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_W); - clSetKernelArg(kernel, 4, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 5, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 6, sizeof(int), &batch.start); - clSetKernelArg(kernel, 7, sizeof(int), &batch.end); + clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr()); clSetKernelArg(kernel, 8, sizeof(int), &ipass); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), @@ -498,10 +498,10 @@ OsdCLComputeController::ApplyVertexEdits( assert(context); cl_int ciErrNum; - size_t globalWorkSize[1] = { batch.end - batch.start }; + size_t globalWorkSize[1] = { batch.GetEnd() - batch.GetStart() }; cl_mem vertexBuffer = context->GetCurrentVertexBuffer(); - const OsdCLHEditTable * edit = context->GetEditTable(batch.tableIndex); + const OsdCLHEditTable * edit = context->GetEditTable(batch.GetTableIndex()); assert(edit); const OsdCLTable * primvarIndices = edit->GetPrimvarIndices(); @@ -520,16 +520,16 @@ OsdCLComputeController::ApplyVertexEdits( clSetKernelArg(kernel, 2, sizeof(cl_mem), &values); clSetKernelArg(kernel, 3, sizeof(int), &primvarOffset); clSetKernelArg(kernel, 4, sizeof(int), &primvarWidth); - clSetKernelArg(kernel, 5, sizeof(int), &batch.vertexOffset); - clSetKernelArg(kernel, 6, sizeof(int), &batch.tableOffset); - clSetKernelArg(kernel, 7, sizeof(int), &batch.start); - clSetKernelArg(kernel, 8, sizeof(int), &batch.end); + clSetKernelArg(kernel, 5, sizeof(int), batch.GetVertexOffsetPtr()); + clSetKernelArg(kernel, 6, sizeof(int), batch.GetTableOffsetPtr()); + clSetKernelArg(kernel, 7, sizeof(int), batch.GetStartPtr()); + clSetKernelArg(kernel, 8, sizeof(int), batch.GetEndPtr()); ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(), kernel, 1, NULL, globalWorkSize, NULL, 0, NULL, NULL); - CL_CHECK_ERROR(ciErrNum, "vertex edit %d %d\n", batch.tableIndex, ciErrNum); + CL_CHECK_ERROR(ciErrNum, "vertex edit %d %d\n", batch.GetTableIndex(), ciErrNum); } else if (edit->GetOperation() == FarVertexEdit::Set) { // XXXX TODO diff --git a/opensubdiv/osd/cpuComputeController.cpp b/opensubdiv/osd/cpuComputeController.cpp index ad421243..0a4f00ba 100644 --- a/opensubdiv/osd/cpuComputeController.cpp +++ b/opensubdiv/osd/cpuComputeController.cpp @@ -84,7 +84,7 @@ OsdCpuComputeController::ApplyBilinearFaceVerticesKernel( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::F_IT)->GetBuffer(), (const int*)context->GetTable(Table::F_ITa)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -100,7 +100,7 @@ OsdCpuComputeController::ApplyBilinearEdgeVerticesKernel( context->GetCurrentVertexBuffer(), context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::E_IT)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -116,7 +116,7 @@ OsdCpuComputeController::ApplyBilinearVertexVerticesKernel( context->GetCurrentVertexBuffer(), context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -133,7 +133,7 @@ OsdCpuComputeController::ApplyCatmarkFaceVerticesKernel( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::F_IT)->GetBuffer(), (const int*)context->GetTable(Table::F_ITa)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -150,7 +150,7 @@ OsdCpuComputeController::ApplyCatmarkEdgeVerticesKernel( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::E_IT)->GetBuffer(), (const float*)context->GetTable(Table::E_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -168,7 +168,7 @@ OsdCpuComputeController::ApplyCatmarkVertexVerticesKernelB( (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const int*)context->GetTable(Table::V_IT)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -185,7 +185,7 @@ OsdCpuComputeController::ApplyCatmarkVertexVerticesKernelA1( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, false); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false); } void @@ -202,7 +202,7 @@ OsdCpuComputeController::ApplyCatmarkVertexVerticesKernelA2( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, true); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true); } void @@ -219,7 +219,7 @@ OsdCpuComputeController::ApplyLoopEdgeVerticesKernel( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::E_IT)->GetBuffer(), (const float*)context->GetTable(Table::E_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -237,7 +237,7 @@ OsdCpuComputeController::ApplyLoopVertexVerticesKernelB( (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const int*)context->GetTable(Table::V_IT)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -254,7 +254,7 @@ OsdCpuComputeController::ApplyLoopVertexVerticesKernelA1( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, false); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false); } void @@ -271,7 +271,7 @@ OsdCpuComputeController::ApplyLoopVertexVerticesKernelA2( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, true); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true); } void @@ -282,7 +282,7 @@ OsdCpuComputeController::ApplyVertexEdits( static_cast(clientdata); assert(context); - const OsdCpuHEditTable *edit = context->GetEditTable(batch.tableIndex); + const OsdCpuHEditTable *edit = context->GetEditTable(batch.GetTableIndex()); assert(edit); const OsdCpuTable * primvarIndices = edit->GetPrimvarIndices(); @@ -293,10 +293,10 @@ OsdCpuComputeController::ApplyVertexEdits( context->GetCurrentVertexBuffer(), edit->GetPrimvarOffset(), edit->GetPrimvarWidth(), - batch.vertexOffset, - batch.tableOffset, - batch.start, - batch.end, + batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), static_cast(primvarIndices->GetBuffer()), static_cast(editValues->GetBuffer())); } else if (edit->GetOperation() == FarVertexEdit::Set) { @@ -304,10 +304,10 @@ OsdCpuComputeController::ApplyVertexEdits( context->GetCurrentVertexBuffer(), edit->GetPrimvarOffset(), edit->GetPrimvarWidth(), - batch.vertexOffset, - batch.tableOffset, - batch.start, - batch.end, + batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), static_cast(primvarIndices->GetBuffer()), static_cast(editValues->GetBuffer())); } diff --git a/opensubdiv/osd/cudaComputeController.cpp b/opensubdiv/osd/cudaComputeController.cpp index 26f92fa0..f695692c 100644 --- a/opensubdiv/osd/cudaComputeController.cpp +++ b/opensubdiv/osd/cudaComputeController.cpp @@ -134,7 +134,7 @@ OsdCudaComputeController::ApplyBilinearFaceVerticesKernel( context->GetCurrentVaryingNumElements(), static_cast(F_IT->GetCudaMemory()), static_cast(F_ITa->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -154,7 +154,7 @@ OsdCudaComputeController::ApplyBilinearEdgeVerticesKernel( context->GetCurrentVertexNumElements()-3, context->GetCurrentVaryingNumElements(), static_cast(E_IT->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -174,7 +174,7 @@ OsdCudaComputeController::ApplyBilinearVertexVerticesKernel( context->GetCurrentVertexNumElements()-3, context->GetCurrentVaryingNumElements(), static_cast(V_ITa->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -197,7 +197,7 @@ OsdCudaComputeController::ApplyCatmarkFaceVerticesKernel( context->GetCurrentVaryingNumElements(), static_cast(F_IT->GetCudaMemory()), static_cast(F_ITa->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -220,7 +220,7 @@ OsdCudaComputeController::ApplyCatmarkEdgeVerticesKernel( context->GetCurrentVaryingNumElements(), static_cast(E_IT->GetCudaMemory()), static_cast(E_W->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -246,7 +246,7 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelB( static_cast(V_ITa->GetCudaMemory()), static_cast(V_IT->GetCudaMemory()), static_cast(V_W->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -269,7 +269,7 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelA1( context->GetCurrentVaryingNumElements(), static_cast(V_ITa->GetCudaMemory()), static_cast(V_W->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, false); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false); } void @@ -292,7 +292,7 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelA2( context->GetCurrentVaryingNumElements(), static_cast(V_ITa->GetCudaMemory()), static_cast(V_W->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, true); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true); } void @@ -315,7 +315,7 @@ OsdCudaComputeController::ApplyLoopEdgeVerticesKernel( context->GetCurrentVaryingNumElements(), static_cast(E_IT->GetCudaMemory()), static_cast(E_W->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -341,7 +341,7 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelB( static_cast(V_ITa->GetCudaMemory()), static_cast(V_IT->GetCudaMemory()), static_cast(V_W->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -364,7 +364,7 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelA1( context->GetCurrentVaryingNumElements(), static_cast(V_ITa->GetCudaMemory()), static_cast(V_W->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, false); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false); } void @@ -387,7 +387,7 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelA2( context->GetCurrentVaryingNumElements(), static_cast(V_ITa->GetCudaMemory()), static_cast(V_W->GetCudaMemory()), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, true); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true); } void @@ -398,7 +398,7 @@ OsdCudaComputeController::ApplyVertexEdits( static_cast(clientdata); assert(context); - const OsdCudaHEditTable *edit = context->GetEditTable(batch.tableIndex); + const OsdCudaHEditTable *edit = context->GetEditTable(batch.GetTableIndex()); assert(edit); const OsdCudaTable * primvarIndices = edit->GetPrimvarIndices(); @@ -410,10 +410,10 @@ OsdCudaComputeController::ApplyVertexEdits( context->GetCurrentVertexNumElements()-3, edit->GetPrimvarOffset(), edit->GetPrimvarWidth(), - batch.vertexOffset, - batch.tableOffset, - batch.start, - batch.end, + batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), static_cast(primvarIndices->GetCudaMemory()), static_cast(editValues->GetCudaMemory())); } else if (edit->GetOperation() == FarVertexEdit::Set) { diff --git a/opensubdiv/osd/glslComputeController.cpp b/opensubdiv/osd/glslComputeController.cpp index 6dc79d62..3c1a4059 100644 --- a/opensubdiv/osd/glslComputeController.cpp +++ b/opensubdiv/osd/glslComputeController.cpp @@ -130,7 +130,7 @@ OsdGLSLComputeController::ApplyBilinearFaceVerticesKernel( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyBilinearFaceVerticesKernel( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -144,7 +144,7 @@ OsdGLSLComputeController::ApplyBilinearEdgeVerticesKernel( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyBilinearEdgeVerticesKernel( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -158,7 +158,7 @@ OsdGLSLComputeController::ApplyBilinearVertexVerticesKernel( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyBilinearVertexVerticesKernel( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -172,7 +172,7 @@ OsdGLSLComputeController::ApplyCatmarkFaceVerticesKernel( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyCatmarkFaceVerticesKernel( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } @@ -188,7 +188,7 @@ OsdGLSLComputeController::ApplyCatmarkEdgeVerticesKernel( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyCatmarkEdgeVerticesKernel( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -202,7 +202,7 @@ OsdGLSLComputeController::ApplyCatmarkVertexVerticesKernelB( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyCatmarkVertexVerticesKernelB( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -216,7 +216,7 @@ OsdGLSLComputeController::ApplyCatmarkVertexVerticesKernelA1( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyCatmarkVertexVerticesKernelA( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, false); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false); } void @@ -230,7 +230,7 @@ OsdGLSLComputeController::ApplyCatmarkVertexVerticesKernelA2( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyCatmarkVertexVerticesKernelA( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, true); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true); } void @@ -244,7 +244,7 @@ OsdGLSLComputeController::ApplyLoopEdgeVerticesKernel( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyLoopEdgeVerticesKernel( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -258,7 +258,7 @@ OsdGLSLComputeController::ApplyLoopVertexVerticesKernelB( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyLoopVertexVerticesKernelB( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -272,7 +272,7 @@ OsdGLSLComputeController::ApplyLoopVertexVerticesKernelA1( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyLoopVertexVerticesKernelA( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, false); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false); } void @@ -286,7 +286,7 @@ OsdGLSLComputeController::ApplyLoopVertexVerticesKernelA2( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); kernelBundle->ApplyLoopVertexVerticesKernelA( - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, true); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true); } void @@ -299,18 +299,21 @@ OsdGLSLComputeController::ApplyVertexEdits( OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle(); - const OsdGLSLComputeHEditTable * edit = context->GetEditTable(batch.tableIndex); + const OsdGLSLComputeHEditTable * edit = context->GetEditTable(batch.GetTableIndex()); assert(edit); - context->BindEditShaderStorageBuffers(batch.tableIndex); + context->BindEditShaderStorageBuffers(batch.GetTableIndex()); int primvarOffset = edit->GetPrimvarOffset(); int primvarWidth = edit->GetPrimvarWidth(); if (edit->GetOperation() == FarVertexEdit::Add) { - kernelBundle->ApplyEditAdd(primvarOffset, primvarWidth, - batch.vertexOffset, batch.tableOffset, - batch.start, batch.end); + kernelBundle->ApplyEditAdd( primvarOffset, + primvarWidth, + batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd()); } else { // XXX: edit SET is not implemented yet. } diff --git a/opensubdiv/osd/glslTransformFeedbackComputeController.cpp b/opensubdiv/osd/glslTransformFeedbackComputeController.cpp index de5c794a..c75ef51e 100644 --- a/opensubdiv/osd/glslTransformFeedbackComputeController.cpp +++ b/opensubdiv/osd/glslTransformFeedbackComputeController.cpp @@ -123,7 +123,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyBilinearFaceVerticesKernel( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -141,7 +141,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyBilinearEdgeVerticesKernel( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -159,7 +159,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyBilinearVertexVerticesKernel( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -177,7 +177,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkFaceVerticesKernel( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } @@ -197,7 +197,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkEdgeVerticesKernel( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -215,7 +215,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkVertexVerticesKernelB( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -233,7 +233,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkVertexVerticesKernelA1( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, false); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false); } void @@ -251,7 +251,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkVertexVerticesKernelA2( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, true); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true); } void @@ -269,7 +269,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyLoopEdgeVerticesKernel( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -287,7 +287,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyLoopVertexVerticesKernelB( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -305,7 +305,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyLoopVertexVerticesKernelA1( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, false); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false); } void @@ -323,7 +323,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyLoopVertexVerticesKernelA2( context->GetNumCurrentVertexElements(), context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, true); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true); } void @@ -336,10 +336,10 @@ OsdGLSLTransformFeedbackComputeController::ApplyVertexEdits( OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle(); - const OsdGLSLTransformFeedbackHEditTable * edit = context->GetEditTable(batch.tableIndex); + const OsdGLSLTransformFeedbackHEditTable * edit = context->GetEditTable(batch.GetTableIndex()); assert(edit); - context->BindEditTextures(batch.tableIndex); + context->BindEditTextures(batch.GetTableIndex()); int primvarOffset = edit->GetPrimvarOffset(); int primvarWidth = edit->GetPrimvarWidth(); @@ -351,7 +351,7 @@ OsdGLSLTransformFeedbackComputeController::ApplyVertexEdits( context->GetCurrentVaryingBuffer(), context->GetNumCurrentVaryingElements(), primvarOffset, primvarWidth, - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } else { // XXX: edit SET is not implemented yet. } diff --git a/opensubdiv/osd/ompComputeController.cpp b/opensubdiv/osd/ompComputeController.cpp index a994b8cb..50fbbc51 100644 --- a/opensubdiv/osd/ompComputeController.cpp +++ b/opensubdiv/osd/ompComputeController.cpp @@ -88,7 +88,7 @@ OsdOmpComputeController::ApplyBilinearFaceVerticesKernel( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::F_IT)->GetBuffer(), (const int*)context->GetTable(Table::F_ITa)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -104,7 +104,7 @@ OsdOmpComputeController::ApplyBilinearEdgeVerticesKernel( context->GetCurrentVertexBuffer(), context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::E_IT)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -120,7 +120,7 @@ OsdOmpComputeController::ApplyBilinearVertexVerticesKernel( context->GetCurrentVertexBuffer(), context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -137,7 +137,7 @@ OsdOmpComputeController::ApplyCatmarkFaceVerticesKernel( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::F_IT)->GetBuffer(), (const int*)context->GetTable(Table::F_ITa)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -154,7 +154,7 @@ OsdOmpComputeController::ApplyCatmarkEdgeVerticesKernel( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::E_IT)->GetBuffer(), (const float*)context->GetTable(Table::E_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -172,7 +172,7 @@ OsdOmpComputeController::ApplyCatmarkVertexVerticesKernelB( (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const int*)context->GetTable(Table::V_IT)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -189,7 +189,7 @@ OsdOmpComputeController::ApplyCatmarkVertexVerticesKernelA1( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, false); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false); } void @@ -206,7 +206,7 @@ OsdOmpComputeController::ApplyCatmarkVertexVerticesKernelA2( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, true); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true); } void @@ -223,7 +223,7 @@ OsdOmpComputeController::ApplyLoopEdgeVerticesKernel( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::E_IT)->GetBuffer(), (const float*)context->GetTable(Table::E_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -241,7 +241,7 @@ OsdOmpComputeController::ApplyLoopVertexVerticesKernelB( (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const int*)context->GetTable(Table::V_IT)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd()); } void @@ -258,7 +258,7 @@ OsdOmpComputeController::ApplyLoopVertexVerticesKernelA1( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, false); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false); } void @@ -275,7 +275,7 @@ OsdOmpComputeController::ApplyLoopVertexVerticesKernelA2( context->GetCurrentVaryingBuffer(), (const int*)context->GetTable(Table::V_ITa)->GetBuffer(), (const float*)context->GetTable(Table::V_W)->GetBuffer(), - batch.vertexOffset, batch.tableOffset, batch.start, batch.end, true); + batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true); } void @@ -286,7 +286,7 @@ OsdOmpComputeController::ApplyVertexEdits( static_cast(clientdata); assert(context); - const OsdCpuHEditTable *edit = context->GetEditTable(batch.tableIndex); + const OsdCpuHEditTable *edit = context->GetEditTable(batch.GetTableIndex()); assert(edit); const OsdCpuTable * primvarIndices = edit->GetPrimvarIndices(); @@ -297,10 +297,10 @@ OsdOmpComputeController::ApplyVertexEdits( context->GetCurrentVertexBuffer(), edit->GetPrimvarOffset(), edit->GetPrimvarWidth(), - batch.vertexOffset, - batch.tableOffset, - batch.start, - batch.end, + batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), static_cast(primvarIndices->GetBuffer()), static_cast(editValues->GetBuffer())); } else if (edit->GetOperation() == FarVertexEdit::Set) { @@ -308,10 +308,10 @@ OsdOmpComputeController::ApplyVertexEdits( context->GetCurrentVertexBuffer(), edit->GetPrimvarOffset(), edit->GetPrimvarWidth(), - batch.vertexOffset, - batch.tableOffset, - batch.start, - batch.end, + batch.GetVertexOffset(), + batch.GetTableOffset(), + batch.GetStart(), + batch.GetEnd(), static_cast(primvarIndices->GetBuffer()), static_cast(editValues->GetBuffer())); }