AuroraMiddleware/OpenSubdiv
1
0
Fork 0
mirror of https://github.com/PixarAnimationStudios/OpenSubdiv synced 2025-01-11 09:00:08 +00:00
Code Issues Releases Wiki Activity

Reorganize ComputeContext and ComputeController.

Browse Source 
Moved transient states (current vertex buffer etc) to controller.
ComputeContext becomes constant so that it's well suited for coarse-grain
parallelism on cpu. The prims sharing same topology (ComputeContext) can
be refined simultaneously by having mutiple compute controllers.
Client facing API doesn't change.
This commit is contained in:
Takahito Tejima 2014-05-06 08:53:36 -07:00
parent 02da793be4
commit a57dd034e7
31 changed files with 940 additions and 1383 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
opensubdiv/far/dispatcher.h
View File

@ -61,14 +61,14 @@ public:
/// @param batch a batch of kernel that need to be processed
///
template <class CONTROLLER, class CONTEXT>
static bool ApplyKernel(CONTROLLER const *controller, CONTEXT *context, FarKernelBatch const &batch);
static bool ApplyKernel(CONTROLLER *controller, CONTEXT *context, FarKernelBatch const &batch);
/// \brief Launches the processing of a vector of kernel batches
/// this is a convenient API for controllers which don't have any user defined kernels.
///
/// @param controller refinement controller implementation
/// @param controller refinement controller implementation (vertex array)
///
/// @param context refinement context implementation (vertex array and subdivision tables)
/// @param context refinement context implementation (subdivision tables)
/// passed to the controller.
///
/// @param batches batches of kernels that need to be processed
@ -76,12 +76,12 @@ public:
/// @param maxlevel process vertex batches up to this level
///
template <class CONTROLLER, class CONTEXT>
static void Refine(CONTROLLER const *controller, CONTEXT *context, FarKernelBatchVector const & batches, int maxlevel);
static void Refine(CONTROLLER *controller, CONTEXT *context, FarKernelBatchVector const & batches, int maxlevel);
};
template <class CONTROLLER, class CONTEXT> bool
FarDispatcher::ApplyKernel(CONTROLLER const *controller, CONTEXT *context, FarKernelBatch const &batch) {
FarDispatcher::ApplyKernel(CONTROLLER *controller, CONTEXT *context, FarKernelBatch const &batch) {
switch(batch.GetKernelType()) {
case FarKernelBatch::CATMARK_FACE_VERTEX:
@ -135,7 +135,7 @@ FarDispatcher::ApplyKernel(CONTROLLER const *controller, CONTEXT *context, FarKe
}
template <class CONTROLLER, class CONTEXT> void
FarDispatcher::Refine(CONTROLLER const *controller, CONTEXT *context, FarKernelBatchVector const & batches, int maxlevel) {
FarDispatcher::Refine(CONTROLLER *controller, CONTEXT *context, FarKernelBatchVector const & batches, int maxlevel) {
for (int i = 0; i < (int)batches.size(); ++i) {
const FarKernelBatch &batch = batches[i];

39
opensubdiv/osd/clComputeContext.cpp
View File

@ -101,8 +101,7 @@ OsdCLHEditTable::GetPrimvarWidth() const {
OsdCLComputeContext::OsdCLComputeContext(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables,
cl_context clContext)
: _clQueue(NULL), _kernelBundle(NULL) {
cl_context clContext) {
// allocate 5 or 7 tables
_tables.resize(subdivisionTables->GetNumTables(), 0);
@ -158,42 +157,6 @@ OsdCLComputeContext::GetEditTable(int tableIndex) const {
return _editTables[tableIndex];
}
cl_mem
OsdCLComputeContext::GetCurrentVertexBuffer() const {
return _currentVertexBuffer;
}
cl_mem
OsdCLComputeContext::GetCurrentVaryingBuffer() const {
return _currentVaryingBuffer;
}
OsdCLKernelBundle *
OsdCLComputeContext::GetKernelBundle() const {
return _kernelBundle;
}
void
OsdCLComputeContext::SetKernelBundle(OsdCLKernelBundle *kernelBundle) {
_kernelBundle = kernelBundle;
}
void
OsdCLComputeContext::SetCommandQueue(cl_command_queue queue) {
_clQueue = queue;
}
cl_command_queue
OsdCLComputeContext::GetCommandQueue() const {
return _clQueue;
}
OsdCLComputeContext *
OsdCLComputeContext::Create(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables,

48
opensubdiv/osd/clComputeContext.h
View File

@ -116,33 +116,6 @@ public:
/// Destructor
virtual ~OsdCLComputeContext();
/// Binds a vertex and a varying data buffers to the context. Binding ensures
/// that data buffers are properly inter-operated between Contexts and
/// Controllers operating across multiple devices.
///
/// @param vertex a buffer containing vertex-interpolated primvar data
///
/// @param varying a buffer containing varying-interpolated primvar data
///
/// @param clQueue OpenCL command queue associated with the primvar data
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying, cl_command_queue clQueue) {
_currentVertexBuffer = vertex ? vertex->BindCLBuffer(clQueue) : NULL;
_currentVaryingBuffer = varying ? varying->BindCLBuffer(clQueue) : NULL;
_clQueue = clQueue;
}
/// Unbinds any previously bound vertex and varying data buffers.
void Unbind() {
_currentVertexBuffer = NULL;
_currentVaryingBuffer = NULL;
_clQueue = NULL;
_kernelBundle = NULL;
}
/// Returns one of the vertex refinement tables.
///
/// @param tableIndex the type of table
@ -158,20 +131,6 @@ public:
///
const OsdCLHEditTable * GetEditTable(int tableIndex) const;
/// Returns a CL handle to the vertex-interpolated data
cl_mem GetCurrentVertexBuffer() const;
/// Returns a CL handle to the varying-interpolated data
cl_mem GetCurrentVaryingBuffer() const;
OsdCLKernelBundle * GetKernelBundle() const;
void SetKernelBundle(OsdCLKernelBundle *kernelBundle);
cl_command_queue GetCommandQueue() const;
void SetCommandQueue(cl_command_queue queue);
protected:
explicit OsdCLComputeContext(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables,
@ -180,13 +139,6 @@ protected:
private:
std::vector<OsdCLTable*> _tables;
std::vector<OsdCLHEditTable*> _editTables;
cl_mem _currentVertexBuffer,
_currentVaryingBuffer;
cl_command_queue _clQueue;
OsdCLKernelBundle *_kernelBundle;
};
} // end namespace OPENSUBDIV_VERSION

147
opensubdiv/osd/clComputeController.cpp
View File

@ -53,7 +53,9 @@ namespace OPENSUBDIV_VERSION {
OsdCLComputeController::OsdCLComputeController(cl_context clContext,
cl_command_queue queue) :
_clContext(clContext), _clQueue(queue) {
_clContext(clContext), _clQueue(queue),
_currentVertexBuffer(0), _currentVaryingBuffer(0),
_currentKernelBundle(NULL) {
}
OsdCLComputeController::~OsdCLComputeController() {
@ -92,33 +94,31 @@ OsdCLComputeController::getKernelBundle(int numVertexElements,
void
OsdCLComputeController::ApplyBilinearFaceVerticesKernel(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
ApplyCatmarkFaceVerticesKernel(batch, context);
}
void
OsdCLComputeController::ApplyBilinearEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
cl_kernel kernel = context->GetKernelBundle()->GetBilinearEdgeKernel();
cl_kernel kernel = _currentKernelBundle->GetBilinearEdgeKernel();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem E_IT = context->GetTable(FarSubdivisionTables::E_IT)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &E_IT);
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(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "bilinear edge kernel %d\n", ciErrNum);
@ -126,26 +126,24 @@ OsdCLComputeController::ApplyBilinearEdgeVerticesKernel(
void
OsdCLComputeController::ApplyBilinearVertexVerticesKernel(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
cl_kernel kernel = context->GetKernelBundle()->GetBilinearVertexKernel();
cl_kernel kernel = _currentKernelBundle->GetBilinearVertexKernel();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem V_ITa = context->GetTable(FarSubdivisionTables::V_ITa)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa);
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(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "bilinear vertex kernel 1 %d\n", ciErrNum);
@ -153,21 +151,19 @@ OsdCLComputeController::ApplyBilinearVertexVerticesKernel(
void
OsdCLComputeController::ApplyCatmarkFaceVerticesKernel(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
cl_kernel kernel = context->GetKernelBundle()->GetCatmarkFaceKernel();
cl_kernel kernel = _currentKernelBundle->GetCatmarkFaceKernel();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem F_IT = context->GetTable(FarSubdivisionTables::F_IT)->GetDevicePtr();
cl_mem F_ITa = context->GetTable(FarSubdivisionTables::F_ITa)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &F_IT);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &F_ITa);
clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr());
@ -175,7 +171,7 @@ OsdCLComputeController::ApplyCatmarkFaceVerticesKernel(
clSetKernelArg(kernel, 6, sizeof(int), batch.GetStartPtr());
clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr());
ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "face kernel %d\n", ciErrNum);
@ -183,21 +179,19 @@ OsdCLComputeController::ApplyCatmarkFaceVerticesKernel(
void
OsdCLComputeController::ApplyCatmarkEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
cl_kernel kernel = context->GetKernelBundle()->GetCatmarkEdgeKernel();
cl_kernel kernel = _currentKernelBundle->GetCatmarkEdgeKernel();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem E_IT = context->GetTable(FarSubdivisionTables::E_IT)->GetDevicePtr();
cl_mem E_W = context->GetTable(FarSubdivisionTables::E_W)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &E_IT);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &E_W);
clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr());
@ -205,7 +199,7 @@ OsdCLComputeController::ApplyCatmarkEdgeVerticesKernel(
clSetKernelArg(kernel, 6, sizeof(int), batch.GetStartPtr());
clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr());
ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "edge kernel %d\n", ciErrNum);
@ -213,22 +207,20 @@ OsdCLComputeController::ApplyCatmarkEdgeVerticesKernel(
void
OsdCLComputeController::ApplyCatmarkVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
cl_kernel kernel = context->GetKernelBundle()->GetCatmarkVertexKernelB();
cl_kernel kernel = _currentKernelBundle->GetCatmarkVertexKernelB();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem V_ITa = context->GetTable(FarSubdivisionTables::V_ITa)->GetDevicePtr();
cl_mem V_IT = context->GetTable(FarSubdivisionTables::V_IT)->GetDevicePtr();
cl_mem V_W = context->GetTable(FarSubdivisionTables::V_W)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_IT);
clSetKernelArg(kernel, 4, sizeof(cl_mem), &V_W);
@ -237,7 +229,7 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelB(
clSetKernelArg(kernel, 7, sizeof(int), batch.GetStartPtr());
clSetKernelArg(kernel, 8, sizeof(int), batch.GetEndPtr());
ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "vertex kernel 1 %d\n", ciErrNum);
@ -245,22 +237,20 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelB(
void
OsdCLComputeController::ApplyCatmarkVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
int ipass = false;
cl_kernel kernel = context->GetKernelBundle()->GetCatmarkVertexKernelA();
cl_kernel kernel = _currentKernelBundle->GetCatmarkVertexKernelA();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem V_ITa = context->GetTable(FarSubdivisionTables::V_ITa)->GetDevicePtr();
cl_mem V_W = context->GetTable(FarSubdivisionTables::V_W)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_W);
clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr());
@ -269,7 +259,7 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelA1(
clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr());
clSetKernelArg(kernel, 8, sizeof(int), &ipass);
ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "vertex kernel 2 %d\n", ciErrNum);
@ -277,22 +267,20 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelA1(
void
OsdCLComputeController::ApplyCatmarkVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
int ipass = true;
cl_kernel kernel = context->GetKernelBundle()->GetCatmarkVertexKernelA();
cl_kernel kernel = _currentKernelBundle->GetCatmarkVertexKernelA();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem V_ITa = context->GetTable(FarSubdivisionTables::V_ITa)->GetDevicePtr();
cl_mem V_W = context->GetTable(FarSubdivisionTables::V_W)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_W);
clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr());
@ -301,7 +289,7 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelA2(
clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr());
clSetKernelArg(kernel, 8, sizeof(int), &ipass);
ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "vertex kernel 2 %d\n", ciErrNum);
@ -309,21 +297,19 @@ OsdCLComputeController::ApplyCatmarkVertexVerticesKernelA2(
void
OsdCLComputeController::ApplyLoopEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
cl_kernel kernel = context->GetKernelBundle()->GetLoopEdgeKernel();
cl_kernel kernel = _currentKernelBundle->GetLoopEdgeKernel();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem E_IT = context->GetTable(FarSubdivisionTables::E_IT)->GetDevicePtr();
cl_mem E_W = context->GetTable(FarSubdivisionTables::E_W)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &E_IT);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &E_W);
clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr());
@ -331,7 +317,7 @@ OsdCLComputeController::ApplyLoopEdgeVerticesKernel(
clSetKernelArg(kernel, 6, sizeof(int), batch.GetStartPtr());
clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr());
ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "edge kernel %d\n", ciErrNum);
@ -339,22 +325,20 @@ OsdCLComputeController::ApplyLoopEdgeVerticesKernel(
void
OsdCLComputeController::ApplyLoopVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
cl_kernel kernel = context->GetKernelBundle()->GetLoopVertexKernelB();
cl_kernel kernel = _currentKernelBundle->GetLoopVertexKernelB();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem V_ITa = context->GetTable(FarSubdivisionTables::V_ITa)->GetDevicePtr();
cl_mem V_IT = context->GetTable(FarSubdivisionTables::V_IT)->GetDevicePtr();
cl_mem V_W = context->GetTable(FarSubdivisionTables::V_W)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_IT);
clSetKernelArg(kernel, 4, sizeof(cl_mem), &V_W);
@ -363,7 +347,7 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelB(
clSetKernelArg(kernel, 7, sizeof(int), batch.GetStartPtr());
clSetKernelArg(kernel, 8, sizeof(int), batch.GetEndPtr());
ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "vertex kernel 1 %d\n", ciErrNum);
@ -371,22 +355,20 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelB(
void
OsdCLComputeController::ApplyLoopVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
int ipass = false;
cl_kernel kernel = context->GetKernelBundle()->GetLoopVertexKernelA();
cl_kernel kernel = _currentKernelBundle->GetLoopVertexKernelA();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem V_ITa = context->GetTable(FarSubdivisionTables::V_ITa)->GetDevicePtr();
cl_mem V_W = context->GetTable(FarSubdivisionTables::V_W)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_W);
clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr());
@ -395,7 +377,7 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelA1(
clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr());
clSetKernelArg(kernel, 8, sizeof(int), &ipass);
ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "vertex kernel 2 %d\n", ciErrNum);
@ -403,22 +385,20 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelA1(
void
OsdCLComputeController::ApplyLoopVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
int ipass = true;
cl_kernel kernel = context->GetKernelBundle()->GetLoopVertexKernelA();
cl_kernel kernel = _currentKernelBundle->GetLoopVertexKernelA();
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
cl_mem varyingBuffer = context->GetCurrentVaryingBuffer();
cl_mem V_ITa = context->GetTable(FarSubdivisionTables::V_ITa)->GetDevicePtr();
cl_mem V_W = context->GetTable(FarSubdivisionTables::V_W)->GetDevicePtr();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &varyingBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &_currentVaryingBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &V_ITa);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &V_W);
clSetKernelArg(kernel, 4, sizeof(int), batch.GetVertexOffsetPtr());
@ -427,7 +407,7 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelA2(
clSetKernelArg(kernel, 7, sizeof(int), batch.GetEndPtr());
clSetKernelArg(kernel, 8, sizeof(int), &ipass);
ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);
CL_CHECK_ERROR(ciErrNum, "vertex kernel 2 %d\n", ciErrNum);
@ -435,13 +415,12 @@ OsdCLComputeController::ApplyLoopVertexVerticesKernelA2(
void
OsdCLComputeController::ApplyVertexEdits(
FarKernelBatch const &batch, OsdCLComputeContext *context) const {
FarKernelBatch const &batch, OsdCLComputeContext const *context) const {
assert(context);
cl_int ciErrNum;
size_t globalWorkSize[1] = { (size_t)(batch.GetEnd() - batch.GetStart()) };
cl_mem vertexBuffer = context->GetCurrentVertexBuffer();
const OsdCLHEditTable * edit = context->GetEditTable(batch.GetTableIndex());
assert(edit);
@ -455,9 +434,9 @@ OsdCLComputeController::ApplyVertexEdits(
int primvarWidth = edit->GetPrimvarWidth();
if (edit->GetOperation() == FarVertexEdit::Add) {
cl_kernel kernel = context->GetKernelBundle()->GetVertexEditAdd();
cl_kernel kernel = _currentKernelBundle->GetVertexEditAdd();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertexBuffer);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &_currentVertexBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &indices);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &values);
clSetKernelArg(kernel, 3, sizeof(int), &primvarOffset);
@ -467,7 +446,7 @@ OsdCLComputeController::ApplyVertexEdits(
clSetKernelArg(kernel, 7, sizeof(int), batch.GetStartPtr());
clSetKernelArg(kernel, 8, sizeof(int), batch.GetEndPtr());
ciErrNum = clEnqueueNDRangeKernel(context->GetCommandQueue(),
ciErrNum = clEnqueueNDRangeKernel(_clQueue,
kernel, 1, NULL, globalWorkSize,
NULL, 0, NULL, NULL);

59
opensubdiv/osd/clComputeController.h
View File

@ -80,23 +80,18 @@ public:
/// @param varyingBuffer varying-interpolated data buffer
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Refine(OsdCLComputeContext *context,
void Refine(ComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer,
VARYING_BUFFER *varyingBuffer) {
if (batches.empty()) return;
int numVertexElements = vertexBuffer ? vertexBuffer->GetNumElements() : 0;
int numVaryingElements = varyingBuffer ? varyingBuffer->GetNumElements() : 0;
context->SetKernelBundle(getKernelBundle(numVertexElements, numVaryingElements));
context->Bind(vertexBuffer, varyingBuffer, _clQueue);
bind(vertexBuffer, varyingBuffer);
FarDispatcher::Refine(this, context, batches, /*maxlevel*/-1);
context->Unbind();
unbind();
}
/// Launch subdivision kernels and apply to given vertex buffers.
@ -109,7 +104,7 @@ public:
/// @param vertexBuffer vertex-interpolated data buffer
///
template<class VERTEX_BUFFER>
void Refine(OsdCLComputeContext *context,
void Refine(ComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer) {
Refine(context, batches, vertexBuffer, (VERTEX_BUFFER*)NULL);
@ -127,43 +122,63 @@ public:
protected:
friend class FarDispatcher;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext const *context) const;
OsdCLKernelBundle * getKernelBundle(int numVertexElements,
int numVaryingElements);
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
int numVertexElements = vertex ? vertex->GetNumElements() : 0;
int numVaryingElements = varying ? varying->GetNumElements() : 0;
_currentVertexBuffer = vertex ? vertex->BindCLBuffer(_clQueue) : NULL;
_currentVaryingBuffer = varying ? varying->BindCLBuffer(_clQueue) : NULL;
_currentKernelBundle = getKernelBundle(numVertexElements, numVaryingElements);
}
void unbind() {
_currentVertexBuffer = NULL;
_currentVaryingBuffer = NULL;
_currentKernelBundle = NULL;
}
private:
cl_context _clContext;
cl_command_queue _clQueue;
std::vector<OsdCLKernelBundle *> _kernelRegistry;
cl_mem _currentVertexBuffer, _currentVaryingBuffer;
OsdCLKernelBundle *_currentKernelBundle;
};
} // end namespace OPENSUBDIV_VERSION

14
opensubdiv/osd/cpuComputeContext.cpp
View File

@ -130,8 +130,6 @@ OsdCpuComputeContext::OsdCpuComputeContext(FarSubdivisionTables const *subdivisi
_editTables.push_back(new OsdCpuHEditTable(edit));
}
}
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
}
OsdCpuComputeContext::~OsdCpuComputeContext() {
@ -162,18 +160,6 @@ OsdCpuComputeContext::GetEditTable(int tableIndex) const {
return _editTables[tableIndex];
}
float *
OsdCpuComputeContext::GetCurrentVertexBuffer() const {
return _currentVertexBuffer;
}
float *
OsdCpuComputeContext::GetCurrentVaryingBuffer() const {
return _currentVaryingBuffer;
}
OsdCpuComputeContext *
OsdCpuComputeContext::Create(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables) {

45
opensubdiv/osd/cpuComputeContext.h
View File

@ -106,46 +106,12 @@ public:
/// Destructor
virtual ~OsdCpuComputeContext();
/// Binds a vertex and a varying data buffers to the context. Binding ensures
/// that data buffers are properly inter-operated between Contexts and
/// Controllers operating across multiple devices.
///
/// @param vertex a buffer containing vertex-interpolated primvar data
///
/// @param varying a buffer containing varying-interpolated primvar data
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
_currentVertexBuffer = vertex ? vertex->BindCpuBuffer() : 0;
_currentVaryingBuffer = varying ? varying->BindCpuBuffer() : 0;
int numVertexElements = vertex ? vertex->GetNumElements() : 0;
int numVaryingElements = varying ? varying->GetNumElements() : 0;
_vdesc.Set(numVertexElements, numVaryingElements);
}
/// Unbinds any previously bound vertex and varying data buffers.
void Unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
_vdesc.Reset();
}
/// Returns one of the vertex refinement tables.
///
/// @param tableIndex the type of table
///
const OsdCpuTable * GetTable(int tableIndex) const;
/// Returns an OsdVertexDescriptor if vertex buffers have been bound.
///
/// @return a descriptor for the format of the vertex data currently bound
///
OsdVertexDescriptor const & GetVertexDescriptor() const {
return _vdesc;
}
/// Returns the number of hierarchical edit tables
int GetNumEditTables() const;
@ -155,12 +121,6 @@ public:
///
const OsdCpuHEditTable * GetEditTable(int tableIndex) const;
/// Returns a pointer to the vertex-interpolated data
float * GetCurrentVertexBuffer() const;
/// Returns a pointer to the varying-interpolated data
float * GetCurrentVaryingBuffer() const;
protected:
explicit OsdCpuComputeContext(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables);
@ -168,11 +128,6 @@ protected:
private:
std::vector<OsdCpuTable*> _tables;
std::vector<OsdCpuHEditTable*> _editTables;
float *_currentVertexBuffer,
*_currentVaryingBuffer;
OsdVertexDescriptor _vdesc;
};
} // end namespace OPENSUBDIV_VERSION

85
opensubdiv/osd/cpuComputeController.cpp
View File

@ -30,7 +30,8 @@ namespace OpenSubdiv {
namespace OPENSUBDIV_VERSION {
OsdCpuComputeController::OsdCpuComputeController() {
OsdCpuComputeController::OsdCpuComputeController() :
_currentVertexBuffer(NULL), _currentVaryingBuffer(NULL) {
}
OsdCpuComputeController::~OsdCpuComputeController() {
@ -38,14 +39,12 @@ OsdCpuComputeController::~OsdCpuComputeController() {
void
OsdCpuComputeController::ApplyBilinearFaceVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeFace(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::F_IT)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::F_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -53,42 +52,36 @@ OsdCpuComputeController::ApplyBilinearFaceVerticesKernel(
void
OsdCpuComputeController::ApplyBilinearEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeBilinearEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdCpuComputeController::ApplyBilinearVertexVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeBilinearVertex(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdCpuComputeController::ApplyCatmarkFaceVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeFace(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::F_IT)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::F_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -96,14 +89,12 @@ OsdCpuComputeController::ApplyCatmarkFaceVerticesKernel(
void
OsdCpuComputeController::ApplyCatmarkEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::E_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -111,14 +102,12 @@ OsdCpuComputeController::ApplyCatmarkEdgeVerticesKernel(
void
OsdCpuComputeController::ApplyCatmarkVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeVertexB(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::V_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
@ -127,14 +116,12 @@ OsdCpuComputeController::ApplyCatmarkVertexVerticesKernelB(
void
OsdCpuComputeController::ApplyCatmarkVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
@ -142,14 +129,12 @@ OsdCpuComputeController::ApplyCatmarkVertexVerticesKernelA1(
void
OsdCpuComputeController::ApplyCatmarkVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
@ -157,14 +142,12 @@ OsdCpuComputeController::ApplyCatmarkVertexVerticesKernelA2(
void
OsdCpuComputeController::ApplyLoopEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::E_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -172,14 +155,12 @@ OsdCpuComputeController::ApplyLoopEdgeVerticesKernel(
void
OsdCpuComputeController::ApplyLoopVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeLoopVertexB(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::V_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
@ -188,14 +169,12 @@ OsdCpuComputeController::ApplyLoopVertexVerticesKernelB(
void
OsdCpuComputeController::ApplyLoopVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
@ -203,14 +182,12 @@ OsdCpuComputeController::ApplyLoopVertexVerticesKernelA1(
void
OsdCpuComputeController::ApplyLoopVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdCpuComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
@ -218,7 +195,7 @@ OsdCpuComputeController::ApplyLoopVertexVerticesKernelA2(
void
OsdCpuComputeController::ApplyVertexEdits(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
@ -229,8 +206,8 @@ OsdCpuComputeController::ApplyVertexEdits(
const OsdCpuTable * editValues = edit->GetEditValues();
if (edit->GetOperation() == FarVertexEdit::Add) {
OsdCpuEditVertexAdd(context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
OsdCpuEditVertexAdd(_vdesc,
_currentVertexBuffer,
edit->GetPrimvarOffset(),
edit->GetPrimvarWidth(),
batch.GetVertexOffset(),
@ -240,8 +217,8 @@ OsdCpuComputeController::ApplyVertexEdits(
static_cast<unsigned int*>(primvarIndices->GetBuffer()),
static_cast<float*>(editValues->GetBuffer()));
} else if (edit->GetOperation() == FarVertexEdit::Set) {
OsdCpuEditVertexSet(context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
OsdCpuEditVertexSet(_vdesc,
_currentVertexBuffer,
edit->GetPrimvarOffset(),
edit->GetPrimvarWidth(),
batch.GetVertexOffset(),

54
opensubdiv/osd/cpuComputeController.h
View File

@ -65,18 +65,18 @@ public:
/// @param varyingBuffer varying-interpolated data buffer
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Refine(OsdCpuComputeContext *context,
void Refine(OsdCpuComputeContext const *context,
FarKernelBatchVector const & batches,
VERTEX_BUFFER *vertexBuffer,
VARYING_BUFFER *varyingBuffer) {
if (batches.empty()) return;
context->Bind(vertexBuffer, varyingBuffer);
bind(vertexBuffer, varyingBuffer);
FarDispatcher::Refine(this, context, batches, /*maxlevel*/-1);
context->Unbind();
unbind();
}
/// Launch subdivision kernels and apply to given vertex buffers.
@ -89,7 +89,7 @@ public:
/// @param vertexBuffer vertex-interpolated data buffer
///
template<class VERTEX_BUFFER>
void Refine(OsdCpuComputeContext *context,
void Refine(OsdCpuComputeContext const *context,
FarKernelBatchVector const & batches,
VERTEX_BUFFER *vertexBuffer) {
Refine(context, batches, vertexBuffer, (VERTEX_BUFFER*)0);
@ -100,34 +100,54 @@ public:
protected:
friend class FarDispatcher;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext const *context) const;
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
_currentVertexBuffer = vertex ? vertex->BindCpuBuffer() : 0;
_currentVaryingBuffer = varying ? varying->BindCpuBuffer() : 0;
int numVertexElements = vertex ? vertex->GetNumElements() : 0;
int numVaryingElements = varying ? varying->GetNumElements() : 0;
_vdesc.Set(numVertexElements, numVaryingElements);
}
void unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
_vdesc.Reset();
}
private:
float *_currentVertexBuffer, *_currentVaryingBuffer;
OsdVertexDescriptor _vdesc;
};

15
opensubdiv/osd/cudaComputeContext.cpp
View File

@ -120,8 +120,7 @@ OsdCudaHEditTable::GetPrimvarWidth() const {
// ----------------------------------------------------------------------------
OsdCudaComputeContext::OsdCudaComputeContext() :
_currentVertexBuffer(NULL), _currentVaryingBuffer(NULL) {
OsdCudaComputeContext::OsdCudaComputeContext() {
}
OsdCudaComputeContext::~OsdCudaComputeContext() {
@ -197,18 +196,6 @@ OsdCudaComputeContext::GetEditTable(int tableIndex) const {
return _editTables[tableIndex];
}
float *
OsdCudaComputeContext::GetCurrentVertexBuffer() const {
return _currentVertexBuffer;
}
float *
OsdCudaComputeContext::GetCurrentVaryingBuffer() const {
return _currentVaryingBuffer;
}
OsdCudaComputeContext *
OsdCudaComputeContext::Create(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables) {

55
opensubdiv/osd/cudaComputeContext.h
View File

@ -114,40 +114,6 @@ public:
/// Destructor
virtual ~OsdCudaComputeContext();
/// Binds a vertex and a varying data buffers to the context. Binding ensures
/// that data buffers are properly inter-operated between Contexts and
/// Controllers operating across multiple devices.
///
/// @param vertex a buffer containing vertex-interpolated primvar data
///
/// @param varying a buffer containing varying-interpolated primvar data
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
if (vertex) {
_currentVertexBuffer = static_cast<float*>(vertex->BindCudaBuffer());
_vdesc.numVertexElements = vertex->GetNumElements();
} else {
_currentVertexBuffer = 0;
_vdesc.numVertexElements = 0;
}
if (varying) {
_currentVaryingBuffer = static_cast<float*>(varying->BindCudaBuffer());
_vdesc.numVaryingElements = varying->GetNumElements();
} else {
_currentVaryingBuffer = 0;
_vdesc.numVaryingElements = 0;
}
}
/// Unbinds any previously bound vertex and varying data buffers.
void Unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
}
/// Returns one of the vertex refinement tables.
///
/// @param tableIndex the type of table
@ -163,21 +129,6 @@ public:
///
const OsdCudaHEditTable * GetEditTable(int tableIndex) const;
/// Returns a pointer to the vertex-interpolated data
float * GetCurrentVertexBuffer() const;
/// Returns a pointer to the varying-interpolated data
float * GetCurrentVaryingBuffer() const;
/// Returns an OsdVertexDescriptor if vertex buffers have been bound.
///
/// @return a descriptor for the format of the vertex data currently bound
///
OsdVertexDescriptor const & GetVertexDescriptor() const {
return _vdesc;
}
protected:
OsdCudaComputeContext();
@ -187,12 +138,6 @@ protected:
private:
std::vector<OsdCudaTable*> _tables;
std::vector<OsdCudaHEditTable*> _editTables;
float *_currentVertexBuffer, // cuda buffers
*_currentVaryingBuffer;
OsdVertexDescriptor _vdesc;
};
} // end namespace OPENSUBDIV_VERSION

105
opensubdiv/osd/cudaComputeController.cpp
View File

@ -74,7 +74,8 @@ void OsdCudaEditVertexAdd(float *vertex, int numUserVertexElements,
namespace OpenSubdiv {
namespace OPENSUBDIV_VERSION {
OsdCudaComputeController::OsdCudaComputeController() {
OsdCudaComputeController::OsdCudaComputeController() :
_currentVertexBuffer(NULL), _currentVaryingBuffer(NULL) {
}
OsdCudaComputeController::~OsdCudaComputeController() {
@ -82,7 +83,7 @@ OsdCudaComputeController::~OsdCudaComputeController() {
void
OsdCudaComputeController::ApplyBilinearFaceVerticesKernel(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -92,10 +93,8 @@ OsdCudaComputeController::ApplyBilinearFaceVerticesKernel(
assert(F_ITa);
OsdCudaComputeFace(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(F_IT->GetCudaMemory()),
static_cast<int*>(F_ITa->GetCudaMemory()),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -103,7 +102,7 @@ OsdCudaComputeController::ApplyBilinearFaceVerticesKernel(
void
OsdCudaComputeController::ApplyBilinearEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -111,17 +110,15 @@ OsdCudaComputeController::ApplyBilinearEdgeVerticesKernel(
assert(E_IT);
OsdCudaComputeBilinearEdge(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(E_IT->GetCudaMemory()),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdCudaComputeController::ApplyBilinearVertexVerticesKernel(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -129,17 +126,15 @@ OsdCudaComputeController::ApplyBilinearVertexVerticesKernel(
assert(V_ITa);
OsdCudaComputeBilinearVertex(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(V_ITa->GetCudaMemory()),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdCudaComputeController::ApplyCatmarkFaceVerticesKernel(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -149,10 +144,8 @@ OsdCudaComputeController::ApplyCatmarkFaceVerticesKernel(
assert(F_ITa);
OsdCudaComputeFace(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(F_IT->GetCudaMemory()),
static_cast<int*>(F_ITa->GetCudaMemory()),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -160,7 +153,7 @@ OsdCudaComputeController::ApplyCatmarkFaceVerticesKernel(
void
OsdCudaComputeController::ApplyCatmarkEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -170,10 +163,8 @@ OsdCudaComputeController::ApplyCatmarkEdgeVerticesKernel(
assert(E_W);
OsdCudaComputeEdge(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(E_IT->GetCudaMemory()),
static_cast<float*>(E_W->GetCudaMemory()),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -181,7 +172,7 @@ OsdCudaComputeController::ApplyCatmarkEdgeVerticesKernel(
void
OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -193,10 +184,8 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelB(
assert(V_W);
OsdCudaComputeVertexB(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(V_ITa->GetCudaMemory()),
static_cast<int*>(V_IT->GetCudaMemory()),
static_cast<float*>(V_W->GetCudaMemory()),
@ -205,7 +194,7 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelB(
void
OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -215,10 +204,8 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelA1(
assert(V_W);
OsdCudaComputeVertexA(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(V_ITa->GetCudaMemory()),
static_cast<float*>(V_W->GetCudaMemory()),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
@ -226,7 +213,7 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelA1(
void
OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -236,10 +223,8 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelA2(
assert(V_W);
OsdCudaComputeVertexA(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(V_ITa->GetCudaMemory()),
static_cast<float*>(V_W->GetCudaMemory()),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
@ -247,7 +232,7 @@ OsdCudaComputeController::ApplyCatmarkVertexVerticesKernelA2(
void
OsdCudaComputeController::ApplyLoopEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -257,10 +242,8 @@ OsdCudaComputeController::ApplyLoopEdgeVerticesKernel(
assert(E_W);
OsdCudaComputeEdge(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(E_IT->GetCudaMemory()),
static_cast<float*>(E_W->GetCudaMemory()),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -268,7 +251,7 @@ OsdCudaComputeController::ApplyLoopEdgeVerticesKernel(
void
OsdCudaComputeController::ApplyLoopVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -280,10 +263,8 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelB(
assert(V_W);
OsdCudaComputeLoopVertexB(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(V_ITa->GetCudaMemory()),
static_cast<int*>(V_IT->GetCudaMemory()),
static_cast<float*>(V_W->GetCudaMemory()),
@ -292,7 +273,7 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelB(
void
OsdCudaComputeController::ApplyLoopVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -302,10 +283,8 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelA1(
assert(V_W);
OsdCudaComputeVertexA(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(V_ITa->GetCudaMemory()),
static_cast<float*>(V_W->GetCudaMemory()),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
@ -313,7 +292,7 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelA1(
void
OsdCudaComputeController::ApplyLoopVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -323,10 +302,8 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelA2(
assert(V_W);
OsdCudaComputeVertexA(
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
context->GetVertexDescriptor().numVaryingElements,
_currentVertexBuffer, _currentVaryingBuffer,
_vdesc.numVertexElements-3, _vdesc.numVaryingElements,
static_cast<int*>(V_ITa->GetCudaMemory()),
static_cast<float*>(V_W->GetCudaMemory()),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
@ -334,7 +311,7 @@ OsdCudaComputeController::ApplyLoopVertexVerticesKernelA2(
void
OsdCudaComputeController::ApplyVertexEdits(
FarKernelBatch const &batch, OsdCudaComputeContext *context) const {
FarKernelBatch const &batch, OsdCudaComputeContext const *context) const {
assert(context);
@ -346,8 +323,8 @@ OsdCudaComputeController::ApplyVertexEdits(
if (edit->GetOperation() == FarVertexEdit::Add) {
OsdCudaEditVertexAdd(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements-3,
_currentVertexBuffer,
_vdesc.numVertexElements-3,
edit->GetPrimvarOffset(),
edit->GetPrimvarWidth(),
batch.GetVertexOffset(),

67
opensubdiv/osd/cudaComputeController.h
View File

@ -65,18 +65,18 @@ public:
/// @param varyingBuffer varying-interpolated data buffer
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Refine(OsdCudaComputeContext *context,
void Refine(OsdCudaComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer,
VARYING_BUFFER *varyingBuffer) {
if (batches.empty()) return;
context->Bind(vertexBuffer, varyingBuffer);
bind(vertexBuffer, varyingBuffer);
FarDispatcher::Refine(this, context, batches, /*maxlevel*/-1);
context->Unbind();
unbind();
}
/// Launch subdivision kernels and apply to given vertex buffers.
@ -89,7 +89,7 @@ public:
/// @param vertexBuffer vertex-interpolated data buffer
///
template<class VERTEX_BUFFER>
void Refine(OsdCudaComputeContext *context,
void Refine(OsdCudaComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer) {
Refine(context, batches, vertexBuffer, (VERTEX_BUFFER*)0);
@ -100,34 +100,67 @@ public:
protected:
friend class FarDispatcher;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext const *context) const;
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
if (vertex) {
_currentVertexBuffer = static_cast<float*>(vertex->BindCudaBuffer());
_vdesc.numVertexElements = vertex->GetNumElements();
} else {
_currentVertexBuffer = 0;
_vdesc.numVertexElements = 0;
}
if (varying) {
_currentVaryingBuffer = static_cast<float*>(varying->BindCudaBuffer());
_vdesc.numVaryingElements = varying->GetNumElements();
} else {
_currentVaryingBuffer = 0;
_vdesc.numVaryingElements = 0;
}
}
/// Unbinds any previously bound vertex and varying data buffers.
void unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
}
private:
float *_currentVertexBuffer, // cuda buffers
*_currentVaryingBuffer;
OsdVertexDescriptor _vdesc;
};
} // end namespace OPENSUBDIV_VERSION

78
opensubdiv/osd/d3d11ComputeContext.cpp
View File

@ -37,7 +37,8 @@ namespace OPENSUBDIV_VERSION {
#define SAFE_RELEASE(p) { if(p) { (p)->Release(); (p)=NULL; } }
void
OsdD3D11ComputeTable::createBuffer(int size, const void *ptr, DXGI_FORMAT format, int numElements, ID3D11DeviceContext *deviceContext) {
OsdD3D11ComputeTable::createBuffer(int size, const void *ptr, DXGI_FORMAT format, int numElements,
ID3D11DeviceContext *deviceContext) {
if (size == 0)
return;
@ -147,9 +148,7 @@ OsdD3D11ComputeHEditTable::GetPrimvarWidth() const {
OsdD3D11ComputeContext::OsdD3D11ComputeContext(
FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables,
ID3D11DeviceContext *deviceContext)
: _deviceContext(deviceContext),
_currentVertexBufferUAV(0), _currentVaryingBufferUAV(0) {
ID3D11DeviceContext *deviceContext) {
// allocate 5 or 7 tables
// XXXtakahito: Although _tables size depends on table type, F_IT is set
@ -211,41 +210,6 @@ OsdD3D11ComputeContext::GetEditTable(int tableIndex) const {
return _editTables[tableIndex];
}
ID3D11UnorderedAccessView *
OsdD3D11ComputeContext::GetCurrentVertexBufferUAV() const {
return _currentVertexBufferUAV;
}
ID3D11UnorderedAccessView *
OsdD3D11ComputeContext::GetCurrentVaryingBufferUAV() const {
return _currentVaryingBufferUAV;
}
OsdD3D11ComputeKernelBundle *
OsdD3D11ComputeContext::GetKernelBundle() const {
return _kernelBundle;
}
void
OsdD3D11ComputeContext::SetKernelBundle(
OsdD3D11ComputeKernelBundle *kernelBundle) {
_kernelBundle = kernelBundle;
}
ID3D11DeviceContext *
OsdD3D11ComputeContext::GetDeviceContext() const {
return _deviceContext;
}
void
OsdD3D11ComputeContext::SetDeviceContext(ID3D11DeviceContext *deviceContext) {
_deviceContext = deviceContext;
}
OsdD3D11ComputeContext *
OsdD3D11ComputeContext::Create(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables,
@ -255,7 +219,8 @@ OsdD3D11ComputeContext::Create(FarSubdivisionTables const *subdivisionTables,
}
void
OsdD3D11ComputeContext::BindEditShaderStorageBuffers(int editIndex) {
OsdD3D11ComputeContext::BindEditShaderStorageBuffers(int editIndex,
ID3D11DeviceContext *deviceContext) const {
const OsdD3D11ComputeHEditTable * edit = _editTables[editIndex];
const OsdD3D11ComputeTable * primvarIndices = edit->GetPrimvarIndices();
@ -265,33 +230,18 @@ OsdD3D11ComputeContext::BindEditShaderStorageBuffers(int editIndex) {
primvarIndices->GetSRV(),
editValues->GetSRV(),
};
_deviceContext->CSSetShaderResources(9, 2, SRViews); // t9-t10
deviceContext->CSSetShaderResources(9, 2, SRViews); // t9-t10
}
void
OsdD3D11ComputeContext::UnbindEditShaderStorageBuffers() {
OsdD3D11ComputeContext::UnbindEditShaderStorageBuffers(ID3D11DeviceContext *deviceContext) const {
ID3D11ShaderResourceView *SRViews[] = { 0, 0 };
_deviceContext->CSSetShaderResources(9, 2, SRViews); // t9-t10
deviceContext->CSSetShaderResources(9, 2, SRViews); // t9-t10
}
void
OsdD3D11ComputeContext::bindShaderStorageBuffers() {
// Unbind the vertexBuffer from the input assembler
ID3D11Buffer *NULLBuffer = 0;
UINT voffset = 0;
UINT vstride = 0;
_deviceContext->IASetVertexBuffers(0, 1, &NULLBuffer, &voffset, &vstride);
// Unbind the vertexBuffer from the vertex shader (gregory patch vertex srv)
ID3D11ShaderResourceView *NULLSRV = 0;
_deviceContext->VSSetShaderResources(0, 1, &NULLSRV);
if (_currentVertexBufferUAV)
_deviceContext->CSSetUnorderedAccessViews(0, 1, &_currentVertexBufferUAV, 0); // u0
if (_currentVaryingBufferUAV)
_deviceContext->CSSetUnorderedAccessViews(1, 1, &_currentVaryingBufferUAV, 0); // u1
OsdD3D11ComputeContext::BindShaderStorageBuffers(ID3D11DeviceContext *deviceContext) const {
// XXX: should be better handling for loop subdivision.
if (_tables[FarSubdivisionTables::F_IT]) {
@ -299,7 +249,7 @@ OsdD3D11ComputeContext::bindShaderStorageBuffers() {
_tables[FarSubdivisionTables::F_IT]->GetSRV(),
_tables[FarSubdivisionTables::F_ITa]->GetSRV(),
};
_deviceContext->CSSetShaderResources(2, 2, SRViews); // t2-t3
deviceContext->CSSetShaderResources(2, 2, SRViews); // t2-t3
}
ID3D11ShaderResourceView *SRViews[] = {
@ -309,16 +259,14 @@ OsdD3D11ComputeContext::bindShaderStorageBuffers() {
_tables[FarSubdivisionTables::E_W]->GetSRV(),
_tables[FarSubdivisionTables::V_W]->GetSRV(),
};
_deviceContext->CSSetShaderResources(4, 5, SRViews); // t4-t8
deviceContext->CSSetShaderResources(4, 5, SRViews); // t4-t8
}
void
OsdD3D11ComputeContext::unbindShaderStorageBuffers() {
OsdD3D11ComputeContext::UnbindShaderStorageBuffers(ID3D11DeviceContext *deviceContext) const {
ID3D11UnorderedAccessView *UAViews[] = { 0, 0 };
_deviceContext->CSSetUnorderedAccessViews(0, 2, UAViews, 0); // u0-u2
ID3D11ShaderResourceView *SRViews[] = { 0, 0, 0, 0, 0, 0, 0 };
_deviceContext->CSSetShaderResources(2, 7, SRViews); // t2-t8
deviceContext->CSSetShaderResources(2, 7, SRViews); // t2-t8
}
} // end namespace OPENSUBDIV_VERSION

67
opensubdiv/osd/d3d11ComputeContext.h
View File

@ -118,34 +118,6 @@ public:
/// Destructor
virtual ~OsdD3D11ComputeContext();
/// Binds a vertex and a varying data buffers to the context. Binding ensures
/// that data buffers are properly inter-operated between Contexts and
/// Controllers operating across multiple devices.
///
/// @param vertex a buffer containing vertex-interpolated primvar data
///
/// @param varying a buffer containing varying-interpolated primvar data
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
_currentVertexBufferUAV = vertex ? vertex->BindD3D11UAV(_deviceContext) : 0;
_currentVaryingBufferUAV = varying ? varying->BindD3D11UAV(_deviceContext) : 0;
_vdesc.numVertexElements = vertex ? vertex->GetNumElements() : 0;
_vdesc.numVaryingElements = varying ? varying->GetNumElements() : 0;
bindShaderStorageBuffers();
}
/// Unbinds any previously bound vertex and varying data buffers.
void Unbind() {
_currentVertexBufferUAV = 0;
_currentVaryingBufferUAV = 0;
unbindShaderStorageBuffers();
}
/// Returns one of the vertex refinement tables.
///
/// @param tableIndex the type of table
@ -161,53 +133,22 @@ public:
///
const OsdD3D11ComputeHEditTable * GetEditTable(int tableIndex) const;
/// Returns a handle to the vertex-interpolated buffer
ID3D11UnorderedAccessView * GetCurrentVertexBufferUAV() const;
void BindShaderStorageBuffers(ID3D11DeviceContext *deviceContext) const;
/// Returns a handle to the varying-interpolated buffer
ID3D11UnorderedAccessView * GetCurrentVaryingBufferUAV() const;
void UnbindShaderStorageBuffers(ID3D11DeviceContext *deviceContext) const;
/// Returns an OsdVertexDescriptor if vertex buffers have been bound.
///
/// @return a descriptor for the format of the vertex data currently bound
///
OsdVertexDescriptor const & GetVertexDescriptor() const {
return _vdesc;
}
void BindEditShaderStorageBuffers(int editIndex, ID3D11DeviceContext *deviceContext) const;
OsdD3D11ComputeKernelBundle * GetKernelBundle() const;
void SetKernelBundle(OsdD3D11ComputeKernelBundle *kernelBundle);
ID3D11DeviceContext * GetDeviceContext() const;
void SetDeviceContext(ID3D11DeviceContext *deviceContext);
void BindEditShaderStorageBuffers(int editIndex);
void UnbindEditShaderStorageBuffers();
void UnbindEditShaderStorageBuffers(ID3D11DeviceContext *deviceContext) const;
protected:
explicit OsdD3D11ComputeContext(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables,
ID3D11DeviceContext *deviceContext);
void bindShaderStorageBuffers();
void unbindShaderStorageBuffers();
private:
std::vector<OsdD3D11ComputeTable*> _tables;
std::vector<OsdD3D11ComputeHEditTable*> _editTables;
ID3D11DeviceContext *_deviceContext;
OsdVertexDescriptor _vdesc;
ID3D11UnorderedAccessView * _currentVertexBufferUAV,
* _currentVaryingBufferUAV;
OsdD3D11ComputeKernelBundle * _kernelBundle;
};
} // end namespace OPENSUBDIV_VERSION

121
opensubdiv/osd/d3d11ComputeController.cpp
View File

@ -38,7 +38,9 @@ namespace OPENSUBDIV_VERSION {
OsdD3D11ComputeController::OsdD3D11ComputeController(
ID3D11DeviceContext *deviceContext)
: _deviceContext(deviceContext), _query(0) {
: _deviceContext(deviceContext), _query(0),
_currentVertexBufferUAV(0), _currentVaryingBufferUAV(0),
_currentKernelBundle(NULL) {
}
OsdD3D11ComputeController::~OsdD3D11ComputeController() {
@ -88,52 +90,69 @@ OsdD3D11ComputeController::getKernels(int numVertexElements,
}
}
void
OsdD3D11ComputeController::bindShaderResources()
{
// Unbind the vertexBuffer from the input assembler
ID3D11Buffer *NULLBuffer = 0;
UINT voffset = 0;
UINT vstride = 0;
_deviceContext->IASetVertexBuffers(0, 1, &NULLBuffer, &voffset, &vstride);
// Unbind the vertexBuffer from the vertex shader (gregory patch vertex srv)
ID3D11ShaderResourceView *NULLSRV = 0;
_deviceContext->VSSetShaderResources(0, 1, &NULLSRV);
if (_currentVertexBufferUAV)
_deviceContext->CSSetUnorderedAccessViews(0, 1, &_currentVertexBufferUAV, 0); // u0
if (_currentVaryingBufferUAV)
_deviceContext->CSSetUnorderedAccessViews(1, 1, &_currentVaryingBufferUAV, 0); // u1
}
void
OsdD3D11ComputeController::unbindShaderResources()
{
ID3D11UnorderedAccessView *UAViews[] = { 0, 0 };
_deviceContext->CSSetUnorderedAccessViews(0, 2, UAViews, 0); // u0-u2
}
void
OsdD3D11ComputeController::ApplyBilinearFaceVerticesKernel(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyBilinearFaceVerticesKernel(
_currentKernelBundle->ApplyBilinearFaceVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdD3D11ComputeController::ApplyBilinearEdgeVerticesKernel(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyBilinearEdgeVerticesKernel(
_currentKernelBundle->ApplyBilinearEdgeVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdD3D11ComputeController::ApplyBilinearVertexVerticesKernel(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyBilinearVertexVerticesKernel(
_currentKernelBundle->ApplyBilinearVertexVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdD3D11ComputeController::ApplyCatmarkFaceVerticesKernel(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkFaceVerticesKernel(
_currentKernelBundle->ApplyCatmarkFaceVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
@ -141,125 +160,109 @@ OsdD3D11ComputeController::ApplyCatmarkFaceVerticesKernel(
void
OsdD3D11ComputeController::ApplyCatmarkEdgeVerticesKernel(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkEdgeVerticesKernel(
_currentKernelBundle->ApplyCatmarkEdgeVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdD3D11ComputeController::ApplyCatmarkVertexVerticesKernelB(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkVertexVerticesKernelB(
_currentKernelBundle->ApplyCatmarkVertexVerticesKernelB(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdD3D11ComputeController::ApplyCatmarkVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkVertexVerticesKernelA(
_currentKernelBundle->ApplyCatmarkVertexVerticesKernelA(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
}
void
OsdD3D11ComputeController::ApplyCatmarkVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkVertexVerticesKernelA(
_currentKernelBundle->ApplyCatmarkVertexVerticesKernelA(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
}
void
OsdD3D11ComputeController::ApplyLoopEdgeVerticesKernel(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopEdgeVerticesKernel(
_currentKernelBundle->ApplyLoopEdgeVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdD3D11ComputeController::ApplyLoopVertexVerticesKernelB(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopVertexVerticesKernelB(
_currentKernelBundle->ApplyLoopVertexVerticesKernelB(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdD3D11ComputeController::ApplyLoopVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopVertexVerticesKernelA(
_currentKernelBundle->ApplyLoopVertexVerticesKernelA(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
}
void
OsdD3D11ComputeController::ApplyLoopVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopVertexVerticesKernelA(
_currentKernelBundle->ApplyLoopVertexVerticesKernelA(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
}
void
OsdD3D11ComputeController::ApplyVertexEdits(
FarKernelBatch const &batch, OsdD3D11ComputeContext *context) const {
FarKernelBatch const &batch, OsdD3D11ComputeContext const *context) const {
assert(context);
OsdD3D11ComputeKernelBundle * kernelBundle = context->GetKernelBundle();
const OsdD3D11ComputeHEditTable * edit = context->GetEditTable(batch.GetTableIndex());
assert(edit);
context->BindEditShaderStorageBuffers(batch.GetTableIndex());
context->BindEditShaderStorageBuffers(batch.GetTableIndex(), _deviceContext);
int primvarOffset = edit->GetPrimvarOffset();
int primvarWidth = edit->GetPrimvarWidth();
if (edit->GetOperation() == FarVertexEdit::Add) {
kernelBundle->ApplyEditAdd(primvarOffset, primvarWidth,
batch.GetVertexOffset(), batch.GetTableOffset(),
batch.GetStart(), batch.GetEnd());
_currentKernelBundle->ApplyEditAdd(primvarOffset, primvarWidth,
batch.GetVertexOffset(),
batch.GetTableOffset(),
batch.GetStart(),
batch.GetEnd());
} else {
// XXX: edit SET is not implemented yet.
}
context->UnbindEditShaderStorageBuffers();
context->UnbindEditShaderStorageBuffers(_deviceContext);
}
} // end namespace OPENSUBDIV_VERSION
} // end namespace OpenSubdiv

75
opensubdiv/osd/d3d11ComputeController.h
View File

@ -76,23 +76,23 @@ public:
/// @param varyingBuffer varying-interpolated data buffer
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Refine(OsdD3D11ComputeContext *context,
void Refine(OsdD3D11ComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer,
VARYING_BUFFER *varyingBuffer) {
if (batches.empty()) return;
int numVertexElements = vertexBuffer ? vertexBuffer->GetNumElements() : 0;
int numVaryingElements = varyingBuffer ? varyingBuffer->GetNumElements() : 0;
bind(vertexBuffer, varyingBuffer);
context->BindShaderStorageBuffers(_deviceContext);
context->SetKernelBundle(getKernels(numVertexElements, numVaryingElements));
context->Bind(vertexBuffer, varyingBuffer);
FarDispatcher::Refine(this,
context,
batches,
-1);
context->Unbind();
context->UnbindShaderStorageBuffers(_deviceContext);
unbind();
}
/// Launch subdivision kernels and apply to given vertex buffers.
@ -105,7 +105,7 @@ public:
/// @param vertexBuffer vertex-interpolated data buffer
///
template<class VERTEX_BUFFER>
void Refine(OsdD3D11ComputeContext *context,
void Refine(OsdD3D11ComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer) {
Refine(context, batches, vertexBuffer, (VERTEX_BUFFER*)NULL);
@ -116,42 +116,77 @@ public:
protected:
friend class FarDispatcher;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext const *context) const;
OsdD3D11ComputeKernelBundle * getKernels(int numVertexElements,
int numVaryingElements);
void bindShaderResources();
void unbindShaderResources();
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
_currentVertexBufferUAV = vertex ? vertex->BindD3D11UAV(_deviceContext) : 0;
_currentVaryingBufferUAV = varying ? varying->BindD3D11UAV(_deviceContext) : 0;
_vdesc.numVertexElements = vertex ? vertex->GetNumElements() : 0;
_vdesc.numVaryingElements = varying ? varying->GetNumElements() : 0;
_currentKernelBundle = getKernels(_vdesc.numVertexElements,
_vdesc.numVaryingElements);
bindShaderResources();
}
void unbind() {
_currentVertexBufferUAV = 0;
_currentVaryingBufferUAV = 0;
_currentKernelBundle = 0;
unbindShaderResources();
}
private:
ID3D11DeviceContext *_deviceContext;
ID3D11Query *_query;
std::vector<OsdD3D11ComputeKernelBundle *> _kernelRegistry;
OsdVertexDescriptor _vdesc;
ID3D11UnorderedAccessView * _currentVertexBufferUAV,
* _currentVaryingBufferUAV;
OsdD3D11ComputeKernelBundle * _currentKernelBundle;
};
} // end namespace OPENSUBDIV_VERSION

85
opensubdiv/osd/gcdComputeController.cpp
View File

@ -30,20 +30,19 @@ namespace OpenSubdiv {
namespace OPENSUBDIV_VERSION {
OsdGcdComputeController::OsdGcdComputeController() {
OsdGcdComputeController::OsdGcdComputeController() :
_currentVertexBuffer(0), _currentVaryingBuffer(0) {
_gcd_queue = dispatch_get_global_queue(DISPATCH_QUEUE_PRIORITY_DEFAULT, 0);
}
void
OsdGcdComputeController::ApplyBilinearFaceVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeFace(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::F_IT)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::F_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(),
@ -52,14 +51,12 @@ OsdGcdComputeController::ApplyBilinearFaceVerticesKernel(
void
OsdGcdComputeController::ApplyBilinearEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeBilinearEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(),
_gcd_queue);
@ -67,14 +64,12 @@ OsdGcdComputeController::ApplyBilinearEdgeVerticesKernel(
void
OsdGcdComputeController::ApplyBilinearVertexVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeBilinearVertex(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(),
_gcd_queue);
@ -82,14 +77,12 @@ OsdGcdComputeController::ApplyBilinearVertexVerticesKernel(
void
OsdGcdComputeController::ApplyCatmarkFaceVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeFace(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::F_IT)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::F_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(),
@ -98,14 +91,12 @@ OsdGcdComputeController::ApplyCatmarkFaceVerticesKernel(
void
OsdGcdComputeController::ApplyCatmarkEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::E_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(),
@ -114,14 +105,12 @@ OsdGcdComputeController::ApplyCatmarkEdgeVerticesKernel(
void
OsdGcdComputeController::ApplyCatmarkVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeVertexB(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::V_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
@ -131,14 +120,12 @@ OsdGcdComputeController::ApplyCatmarkVertexVerticesKernelB(
void
OsdGcdComputeController::ApplyCatmarkVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false,
@ -147,14 +134,12 @@ OsdGcdComputeController::ApplyCatmarkVertexVerticesKernelA1(
void
OsdGcdComputeController::ApplyCatmarkVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true,
@ -163,14 +148,12 @@ OsdGcdComputeController::ApplyCatmarkVertexVerticesKernelA2(
void
OsdGcdComputeController::ApplyLoopEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::E_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(),
@ -179,14 +162,12 @@ OsdGcdComputeController::ApplyLoopEdgeVerticesKernel(
void
OsdGcdComputeController::ApplyLoopVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeLoopVertexB(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::V_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
@ -196,14 +177,12 @@ OsdGcdComputeController::ApplyLoopVertexVerticesKernelB(
void
OsdGcdComputeController::ApplyLoopVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false,
@ -212,14 +191,12 @@ OsdGcdComputeController::ApplyLoopVertexVerticesKernelA1(
void
OsdGcdComputeController::ApplyLoopVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdGcdComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true,
@ -228,7 +205,7 @@ OsdGcdComputeController::ApplyLoopVertexVerticesKernelA2(
void
OsdGcdComputeController::ApplyVertexEdits(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
@ -239,8 +216,8 @@ OsdGcdComputeController::ApplyVertexEdits(
const OsdCpuTable * editValues = edit->GetEditValues();
if (edit->GetOperation() == FarVertexEdit::Add) {
OsdGcdEditVertexAdd(context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
OsdGcdEditVertexAdd(_vdesc,
_currentVertexBuffer,
edit->GetPrimvarOffset(),
edit->GetPrimvarWidth(),
batch.GetVertexOffset(),
@ -251,8 +228,8 @@ OsdGcdComputeController::ApplyVertexEdits(
static_cast<float*>(editValues->GetBuffer()),
_gcd_queue);
} else if (edit->GetOperation() == FarVertexEdit::Set) {
OsdGcdEditVertexSet(context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
OsdGcdEditVertexSet(_vdesc,
_currentVertexBuffer,
edit->GetPrimvarOffset(),
edit->GetPrimvarWidth(),
batch.GetVertexOffset(),

54
opensubdiv/osd/gcdComputeController.h
View File

@ -65,18 +65,18 @@ public:
/// @param varyingBuffer varying-interpolated data buffer
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Refine(OsdCpuComputeContext *context,
void Refine(OsdCpuComputeContext const *context,
FarKernelBatchVector const & batches,
VERTEX_BUFFER *vertexBuffer,
VARYING_BUFFER *varyingBuffer) {
if (batches.empty()) return;
context->Bind(vertexBuffer, varyingBuffer);
bind(vertexBuffer, varyingBuffer);
FarDispatcher::Refine(this, context, batches, /*maxlevel*/-1);
context->Unbind();
unbind();
}
/// Launch subdivision kernels and apply to given vertex buffers.
@ -89,7 +89,7 @@ public:
/// @param vertexBuffer vertex-interpolated data buffer
///
template<class VERTEX_BUFFER>
void Refine(OsdCpuComputeContext *context,
void Refine(OsdCpuComputeContext const *context,
FarKernelBatchVector const & batches,
VERTEX_BUFFER *vertexBuffer) {
Refine(context, batches, vertexBuffer, (VERTEX_BUFFER*)0);
@ -100,37 +100,57 @@ public:
protected:
friend class FarDispatcher;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext const *context) const;
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
_currentVertexBuffer = vertex ? vertex->BindCpuBuffer() : 0;
_currentVaryingBuffer = varying ? varying->BindCpuBuffer() : 0;
int numVertexElements = vertex ? vertex->GetNumElements() : 0;
int numVaryingElements = varying ? varying->GetNumElements() : 0;
_vdesc.Set(numVertexElements, numVaryingElements);
}
void unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
_vdesc.Reset();
}
private:
dispatch_queue_t _gcd_queue;
float *_currentVertexBuffer, *_currentVaryingBuffer;
OsdVertexDescriptor _vdesc;
};
} // end namespace OPENSUBDIV_VERSION

49
opensubdiv/osd/glslComputeContext.cpp
View File

@ -119,8 +119,7 @@ OsdGLSLComputeHEditTable::GetPrimvarWidth() const {
OsdGLSLComputeContext::OsdGLSLComputeContext(
FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables)
: _vertexTexture(0), _varyingTexture(0) {
FarVertexEditTables const *vertexEditTables) {
// allocate 5 or 7 tables
// XXXtakahito: Although _tables size depends on table type, F_IT is set
@ -184,31 +183,6 @@ OsdGLSLComputeContext::GetEditTable(int tableIndex) const {
return _editTables[tableIndex];
}
GLuint
OsdGLSLComputeContext::GetCurrentVertexBuffer() const {
return _currentVertexBuffer;
}
GLuint
OsdGLSLComputeContext::GetCurrentVaryingBuffer() const {
return _currentVaryingBuffer;
}
OsdGLSLComputeKernelBundle *
OsdGLSLComputeContext::GetKernelBundle() const {
return _kernelBundle;
}
void
OsdGLSLComputeContext::SetKernelBundle(
OsdGLSLComputeKernelBundle *kernelBundle) {
_kernelBundle = kernelBundle;
}
OsdGLSLComputeContext *
OsdGLSLComputeContext::Create(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables) {
@ -217,7 +191,7 @@ OsdGLSLComputeContext::Create(FarSubdivisionTables const *subdivisionTables,
}
void
OsdGLSLComputeContext::BindEditShaderStorageBuffers(int editIndex) {
OsdGLSLComputeContext::BindEditShaderStorageBuffers(int editIndex) const {
const OsdGLSLComputeHEditTable * edit = _editTables[editIndex];
const OsdGLSLComputeTable * primvarIndices = edit->GetPrimvarIndices();
@ -230,24 +204,16 @@ OsdGLSLComputeContext::BindEditShaderStorageBuffers(int editIndex) {
}
void
OsdGLSLComputeContext::UnbindEditShaderStorageBuffers() {
OsdGLSLComputeContext::UnbindEditShaderStorageBuffers() const {
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 9, 0);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 10, 0);
}
void
OsdGLSLComputeContext::bindShaderStorageBuffers() {
OsdGLSLComputeContext::BindShaderStorageBuffers() const {
_kernelBundle->UseProgram();
if (_currentVertexBuffer)
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, _currentVertexBuffer);
if (_currentVaryingBuffer)
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, _currentVaryingBuffer);
// XXX: should be better handling for loop subdivision.
// 0 and 1 are reserved for vertex/varying buffer bindings.
if (_tables[FarSubdivisionTables::F_IT]) {
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2,
_tables[FarSubdivisionTables::F_IT]->GetBuffer());
@ -268,10 +234,9 @@ OsdGLSLComputeContext::bindShaderStorageBuffers() {
}
void
OsdGLSLComputeContext::unbindShaderStorageBuffers() {
OsdGLSLComputeContext::UnbindShaderStorageBuffers() const {
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, 0);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, 0);
// 0 and 1 are reserved for vertex/varying buffer bindings.
if (_tables[FarSubdivisionTables::F_IT]) {
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, 0);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, 0);

86
opensubdiv/osd/glslComputeContext.h
View File

@ -113,34 +113,6 @@ public:
/// Destructor
virtual ~OsdGLSLComputeContext();
/// Binds a vertex and a varying data buffers to the context. Binding ensures
/// that data buffers are properly inter-operated between Contexts and
/// Controllers operating across multiple devices.
///
/// @param vertex a buffer containing vertex-interpolated primvar data
///
/// @param varying a buffer containing varying-interpolated primvar data
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
_currentVertexBuffer = vertex ? vertex->BindVBO() : 0;
_currentVaryingBuffer = varying ? varying->BindVBO() : 0;
_vdesc.numVertexElements = vertex ? vertex->GetNumElements() : 0;
_vdesc.numVaryingElements = varying ? varying->GetNumElements() : 0;
bindShaderStorageBuffers();
}
/// Unbinds any previously bound vertex and varying data buffers.
void Unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
unbindShaderStorageBuffers();
}
/// Returns one of the vertex refinement tables.
///
/// @param tableIndex the type of table
@ -156,71 +128,21 @@ public:
///
const OsdGLSLComputeHEditTable * GetEditTable(int tableIndex) const;
/// Returns a handle to the vertex-interpolated buffer
GLuint GetCurrentVertexBuffer() const;
void BindShaderStorageBuffers() const;
/// Returns a handle to the varying-interpolated buffer
GLuint GetCurrentVaryingBuffer() const;
void UnbindShaderStorageBuffers() const;
/// Returns an OsdVertexDescriptor if vertex buffers have been bound.
///
/// @return a descriptor for the format of the vertex data currently bound
///
OsdVertexDescriptor const & GetVertexDescriptor() const {
return _vdesc;
}
void BindEditShaderStorageBuffers(int editIndex) const;
OsdGLSLComputeKernelBundle * GetKernelBundle() const;
void SetKernelBundle(OsdGLSLComputeKernelBundle *kernelBundle);
void BindUniformBlockBilinearFace(GLuint program, int level);
void BindUniformBlockBilinearEdge(GLuint program, int level);
void BindUniformBlockBilinearVertex(GLuint program, int level);
void BindUniformBlockCatmarkFace(GLuint program, int level);
void BindUniformBlockCatmarkEdge(GLuint program, int level);
void BindUniformBlockCatmarkVertexA0(GLuint program, int level);
void BindUniformBlockCatmarkVertexA1(GLuint program, int level);
void BindUniformBlockCatmarkVertexB(GLuint program, int level);
void BindUniformBlockLoopEdge(GLuint program, int level);
void BindUniformBlockLoopVertexA(GLuint program, int level);
void BindUniformBlockLoopVertexB(GLuint program, int level);
void BindEditShaderStorageBuffers(int editIndex);
void UnbindEditShaderStorageBuffers();
void UnbindEditShaderStorageBuffers() const;
protected:
explicit OsdGLSLComputeContext(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables);
void bindShaderStorageBuffers();
void unbindShaderStorageBuffers();
private:
std::vector<OsdGLSLComputeTable*> _tables;
std::vector<OsdGLSLComputeHEditTable*> _editTables;
GLuint _vertexTexture,
_varyingTexture;
OsdVertexDescriptor _vdesc;
GLuint _currentVertexBuffer,
_currentVaryingBuffer;
OsdGLSLComputeKernelBundle * _kernelBundle;
};
} // end namespace OPENSUBDIV_VERSION

115
opensubdiv/osd/glslComputeController.cpp
View File

@ -34,7 +34,8 @@
namespace OpenSubdiv {
namespace OPENSUBDIV_VERSION {
OsdGLSLComputeController::OsdGLSLComputeController() {
OsdGLSLComputeController::OsdGLSLComputeController()
: _currentVertexBuffer(0), _currentVaryingBuffer(0), _currentKernelBundle(NULL) {
}
OsdGLSLComputeController::~OsdGLSLComputeController() {
@ -71,51 +72,63 @@ OsdGLSLComputeController::getKernels(int numVertexElements,
}
}
void
OsdGLSLComputeController::bindBufferAndProgram() {
if (_currentVertexBuffer)
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, _currentVertexBuffer);
if (_currentVaryingBuffer)
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, _currentVaryingBuffer);
_currentKernelBundle->UseProgram();
}
void
OsdGLSLComputeController::unbindBufferAndProgram() {
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, 0);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, 0);
glUseProgram(0);
}
void
OsdGLSLComputeController::ApplyBilinearFaceVerticesKernel(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyBilinearFaceVerticesKernel(
_currentKernelBundle->ApplyBilinearFaceVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLComputeController::ApplyBilinearEdgeVerticesKernel(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyBilinearEdgeVerticesKernel(
_currentKernelBundle->ApplyBilinearEdgeVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLComputeController::ApplyBilinearVertexVerticesKernel(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyBilinearVertexVerticesKernel(
_currentKernelBundle->ApplyBilinearVertexVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLComputeController::ApplyCatmarkFaceVerticesKernel(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkFaceVerticesKernel(
_currentKernelBundle->ApplyCatmarkFaceVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
@ -123,108 +136,90 @@ OsdGLSLComputeController::ApplyCatmarkFaceVerticesKernel(
void
OsdGLSLComputeController::ApplyCatmarkEdgeVerticesKernel(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkEdgeVerticesKernel(
_currentKernelBundle->ApplyCatmarkEdgeVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLComputeController::ApplyCatmarkVertexVerticesKernelB(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkVertexVerticesKernelB(
_currentKernelBundle->ApplyCatmarkVertexVerticesKernelB(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLComputeController::ApplyCatmarkVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkVertexVerticesKernelA(
_currentKernelBundle->ApplyCatmarkVertexVerticesKernelA(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
}
void
OsdGLSLComputeController::ApplyCatmarkVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkVertexVerticesKernelA(
_currentKernelBundle->ApplyCatmarkVertexVerticesKernelA(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
}
void
OsdGLSLComputeController::ApplyLoopEdgeVerticesKernel(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopEdgeVerticesKernel(
_currentKernelBundle->ApplyLoopEdgeVerticesKernel(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLComputeController::ApplyLoopVertexVerticesKernelB(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopVertexVerticesKernelB(
_currentKernelBundle->ApplyLoopVertexVerticesKernelB(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLComputeController::ApplyLoopVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopVertexVerticesKernelA(
_currentKernelBundle->ApplyLoopVertexVerticesKernelA(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
}
void
OsdGLSLComputeController::ApplyLoopVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopVertexVerticesKernelA(
_currentKernelBundle->ApplyLoopVertexVerticesKernelA(
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
}
void
OsdGLSLComputeController::ApplyVertexEdits(
FarKernelBatch const &batch, OsdGLSLComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLComputeContext const *context) const {
assert(context);
OsdGLSLComputeKernelBundle * kernelBundle = context->GetKernelBundle();
const OsdGLSLComputeHEditTable * edit = context->GetEditTable(batch.GetTableIndex());
assert(edit);
@ -232,18 +227,18 @@ OsdGLSLComputeController::ApplyVertexEdits(
int primvarOffset = edit->GetPrimvarOffset();
int primvarWidth = edit->GetPrimvarWidth();
if (edit->GetOperation() == FarVertexEdit::Add) {
kernelBundle->ApplyEditAdd( primvarOffset,
primvarWidth,
batch.GetVertexOffset(),
batch.GetTableOffset(),
batch.GetStart(),
batch.GetEnd());
_currentKernelBundle->ApplyEditAdd( primvarOffset,
primvarWidth,
batch.GetVertexOffset(),
batch.GetTableOffset(),
batch.GetStart(),
batch.GetEnd());
} else {
// XXX: edit SET is not implemented yet.
}
context->UnbindEditShaderStorageBuffers();
}

70
opensubdiv/osd/glslComputeController.h
View File

@ -70,7 +70,7 @@ public:
/// @param varyingBuffer varying-interpolated data buffer
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Refine(OsdGLSLComputeContext *context,
void Refine(OsdGLSLComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer,
VARYING_BUFFER *varyingBuffer) {
@ -80,12 +80,15 @@ public:
int numVertexElements = vertexBuffer ? vertexBuffer->GetNumElements() : 0;
int numVaryingElements = varyingBuffer ? varyingBuffer->GetNumElements() : 0;
context->SetKernelBundle(getKernels(numVertexElements, numVaryingElements));
context->Bind(vertexBuffer, varyingBuffer);
bind(vertexBuffer, varyingBuffer, getKernels(numVertexElements, numVaryingElements));
// bind table buffers.
context->BindShaderStorageBuffers();
FarDispatcher::Refine(this, context, batches, /*maxlevel*/-1);
context->Unbind();
// unbind table buffers.
context->UnbindShaderStorageBuffers();
unbind();
}
/// Launch subdivision kernels and apply to given vertex buffers.
@ -98,7 +101,7 @@ public:
/// @param vertexBuffer vertex-interpolated data buffer
///
template<class VERTEX_BUFFER>
void Refine(OsdGLSLComputeContext *context,
void Refine(OsdGLSLComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer) {
Refine(context, batches, vertexBuffer, (VERTEX_BUFFER*)NULL);
@ -109,40 +112,71 @@ public:
protected:
friend class FarDispatcher;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext const *context) const;
OsdGLSLComputeKernelBundle * getKernels(int numVertexElements,
int numVaryingElements);
void bindBufferAndProgram();
void unbindBufferAndProgram();
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying, OsdGLSLComputeKernelBundle *kernelBundle) {
_currentVertexBuffer = vertex ? vertex->BindVBO() : 0;
_currentVaryingBuffer = varying ? varying->BindVBO() : 0;
_vdesc.numVertexElements = vertex ? vertex->GetNumElements() : 0;
_vdesc.numVaryingElements = varying ? varying->GetNumElements() : 0;
_currentKernelBundle = kernelBundle;
bindBufferAndProgram();
}
/// Unbinds any previously bound vertex and varying data buffers.
void unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
}
private:
std::vector<OsdGLSLComputeKernelBundle *> _kernelRegistry;
GLuint _currentVertexBuffer, _currentVaryingBuffer;
OsdVertexDescriptor _vdesc;
OsdGLSLComputeKernelBundle * _currentKernelBundle;
};
} // end namespace OPENSUBDIV_VERSION

111
opensubdiv/osd/glslTransformFeedbackComputeContext.cpp
View File

@ -128,8 +128,7 @@ OsdGLSLTransformFeedbackHEditTable::GetPrimvarWidth() const {
OsdGLSLTransformFeedbackComputeContext::OsdGLSLTransformFeedbackComputeContext(
FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables) :
_vertexTexture(0), _varyingTexture(0) {
FarVertexEditTables const *vertexEditTables) {
// allocate 5 or 7 tables
_tables.resize(7, 0);
@ -169,8 +168,6 @@ OsdGLSLTransformFeedbackComputeContext::~OsdGLSLTransformFeedbackComputeContext(
for (size_t i = 0; i < _editTables.size(); ++i) {
delete _editTables[i];
}
if (_vertexTexture) glDeleteTextures(1, &_vertexTexture);
if (_varyingTexture) glDeleteTextures(1, &_varyingTexture);
}
const OsdGLSLTransformFeedbackTable *
@ -191,30 +188,6 @@ OsdGLSLTransformFeedbackComputeContext::GetEditTable(int tableIndex) const {
return _editTables[tableIndex];
}
GLuint
OsdGLSLTransformFeedbackComputeContext::GetCurrentVertexBuffer() const {
return _currentVertexBuffer;
}
GLuint
OsdGLSLTransformFeedbackComputeContext::GetCurrentVaryingBuffer() const {
return _currentVaryingBuffer;
}
OsdGLSLTransformFeedbackKernelBundle *
OsdGLSLTransformFeedbackComputeContext::GetKernelBundle() const {
return _kernelBundle;
}
void
OsdGLSLTransformFeedbackComputeContext::SetKernelBundle(OsdGLSLTransformFeedbackKernelBundle *kernelBundle) {
_kernelBundle = kernelBundle;
}
OsdGLSLTransformFeedbackComputeContext *
OsdGLSLTransformFeedbackComputeContext::Create(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTables) {
@ -223,27 +196,28 @@ OsdGLSLTransformFeedbackComputeContext::Create(FarSubdivisionTables const *subdi
}
void
OsdGLSLTransformFeedbackComputeContext::BindEditTextures(int editIndex) {
OsdGLSLTransformFeedbackComputeContext::BindEditTextures(
int editIndex, OsdGLSLTransformFeedbackKernelBundle const *kernelBundle) const {
const OsdGLSLTransformFeedbackHEditTable * edit = _editTables[editIndex];
const OsdGLSLTransformFeedbackTable * primvarIndices = edit->GetPrimvarIndices();
const OsdGLSLTransformFeedbackTable * editValues = edit->GetEditValues();
bindTexture(_kernelBundle->GetEditIndicesUniformLocation(),
bindTexture(kernelBundle->GetEditIndicesUniformLocation(),
primvarIndices->GetTexture(), 9);
bindTexture(_kernelBundle->GetEditValuesUniformLocation(),
bindTexture(kernelBundle->GetEditValuesUniformLocation(),
editValues->GetTexture(), 10);
}
void
OsdGLSLTransformFeedbackComputeContext::UnbindEditTextures() {
OsdGLSLTransformFeedbackComputeContext::UnbindEditTextures() const {
unbindTexture(9);
unbindTexture(10);
}
void
OsdGLSLTransformFeedbackComputeContext::bindTexture(GLint samplerUniform, GLuint texture, int unit) {
OsdGLSLTransformFeedbackComputeContext::bindTexture(GLint samplerUniform, GLuint texture, int unit) const {
if (samplerUniform == -1) return;
glUniform1i(samplerUniform, unit);
@ -253,89 +227,42 @@ OsdGLSLTransformFeedbackComputeContext::bindTexture(GLint samplerUniform, GLuint
}
void
OsdGLSLTransformFeedbackComputeContext::unbindTexture(GLuint unit) {
OsdGLSLTransformFeedbackComputeContext::unbindTexture(GLuint unit) const {
glActiveTexture(GL_TEXTURE0 + unit);
glBindTexture(GL_TEXTURE_BUFFER, 0);
}
void
OsdGLSLTransformFeedbackComputeContext::bind() {
glEnable(GL_RASTERIZER_DISCARD);
_kernelBundle->UseProgram();
// bind vertex texture
if (_currentVertexBuffer) {
if (not _vertexTexture) glGenTextures(1, &_vertexTexture);
#if defined(GL_EXT_direct_state_access)
if (glTextureBufferEXT) {
glTextureBufferEXT(_vertexTexture, GL_TEXTURE_BUFFER, GL_R32F, _currentVertexBuffer);
} else {
#else
{
#endif
glBindTexture(GL_TEXTURE_BUFFER, _vertexTexture);
glTexBuffer(GL_TEXTURE_BUFFER, GL_R32F, _currentVertexBuffer);
glBindTexture(GL_TEXTURE_BUFFER, 0);
}
}
if (_currentVaryingBuffer) {
if (not _varyingTexture) glGenTextures(1, &_varyingTexture);
#if defined(GL_EXT_direct_state_access)
if (glTextureBufferEXT) {
glTextureBufferEXT(_varyingTexture, GL_TEXTURE_BUFFER, GL_R32F, _currentVaryingBuffer);
} else {
#else
{
#endif
glBindTexture(GL_TEXTURE_BUFFER, _varyingTexture);
glTexBuffer(GL_TEXTURE_BUFFER, GL_R32F, _currentVaryingBuffer);
glBindTexture(GL_TEXTURE_BUFFER, 0);
}
}
if (_vertexTexture)
bindTexture(_kernelBundle->GetVertexUniformLocation(), _vertexTexture, 0);
if (_varyingTexture)
bindTexture(_kernelBundle->GetVaryingUniformLocation(), _varyingTexture, 1);
OsdGLSLTransformFeedbackComputeContext::BindTableTextures(
OsdGLSLTransformFeedbackKernelBundle const *kernelBundle) const {
// XXX: loop...
if (_tables[FarSubdivisionTables::F_IT]) {
bindTexture(_kernelBundle->GetTableUniformLocation(FarSubdivisionTables::F_IT),
bindTexture(kernelBundle->GetTableUniformLocation(FarSubdivisionTables::F_IT),
_tables[FarSubdivisionTables::F_IT]->GetTexture(), 2);
bindTexture(_kernelBundle->GetTableUniformLocation(FarSubdivisionTables::F_ITa),
bindTexture(kernelBundle->GetTableUniformLocation(FarSubdivisionTables::F_ITa),
_tables[FarSubdivisionTables::F_ITa]->GetTexture(), 3);
}
bindTexture(_kernelBundle->GetTableUniformLocation(FarSubdivisionTables::E_IT),
bindTexture(kernelBundle->GetTableUniformLocation(FarSubdivisionTables::E_IT),
_tables[FarSubdivisionTables::E_IT]->GetTexture(), 4);
bindTexture(_kernelBundle->GetTableUniformLocation(FarSubdivisionTables::V_IT),
bindTexture(kernelBundle->GetTableUniformLocation(FarSubdivisionTables::V_IT),
_tables[FarSubdivisionTables::V_IT]->GetTexture(), 5);
bindTexture(_kernelBundle->GetTableUniformLocation(FarSubdivisionTables::V_ITa),
bindTexture(kernelBundle->GetTableUniformLocation(FarSubdivisionTables::V_ITa),
_tables[FarSubdivisionTables::V_ITa]->GetTexture(), 6);
bindTexture(_kernelBundle->GetTableUniformLocation(FarSubdivisionTables::E_W),
bindTexture(kernelBundle->GetTableUniformLocation(FarSubdivisionTables::E_W),
_tables[FarSubdivisionTables::E_W]->GetTexture(), 7);
bindTexture(_kernelBundle->GetTableUniformLocation(FarSubdivisionTables::V_W),
bindTexture(kernelBundle->GetTableUniformLocation(FarSubdivisionTables::V_W),
_tables[FarSubdivisionTables::V_W]->GetTexture(), 8);
// bind texture image (for edit kernel)
glUniform1i(_kernelBundle->GetVertexBufferImageUniformLocation(), 0);
glBindImageTexture(0, _vertexTexture, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_R32F);
}
void
OsdGLSLTransformFeedbackComputeContext::unbind() {
OsdGLSLTransformFeedbackComputeContext::UnbindTableTextures() const {
for (int i = 8; i >= 0; --i) {
for (int i = 8; i >= 2; --i) {
unbindTexture(i);
}
glBindImageTexture(0, 0, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_R32F);
glDisable(GL_RASTERIZER_DISCARD);
glUseProgram(0);
glActiveTexture(GL_TEXTURE0);
}
} // end namespace OPENSUBDIV_VERSION

70
opensubdiv/osd/glslTransformFeedbackComputeContext.h
View File

@ -108,33 +108,6 @@ public:
/// Destructor
virtual ~OsdGLSLTransformFeedbackComputeContext();
/// Binds a vertex and a varying data buffers to the context. Binding ensures
/// that data buffers are properly inter-operated between Contexts and
/// Controllers operating across multiple devices.
///
/// @param vertex a buffer containing vertex-interpolated primvar data
///
/// @param varying a buffer containing varying-interpolated primvar data
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
_currentVertexBuffer = vertex ? vertex->BindVBO() : 0;
_currentVaryingBuffer = varying ? varying->BindVBO() : 0;
_vdesc.numVertexElements = vertex ? vertex->GetNumElements() : 0;
_vdesc.numVaryingElements = varying ? varying->GetNumElements() : 0;
bind();
}
/// Unbinds any previously bound vertex and varying data buffers.
void Unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
unbind();
}
/// Returns one of the vertex refinement tables.
///
/// @param tableIndex the type of table
@ -150,53 +123,28 @@ public:
///
const OsdGLSLTransformFeedbackHEditTable * GetEditTable(int tableIndex) const;
/// Returns a handle to the vertex-interpolated buffer
GLuint GetCurrentVertexBuffer() const;
void BindTableTextures(
OsdGLSLTransformFeedbackKernelBundle const *kernelBundle) const;
/// Returns a handle to the varying-interpolated buffer
GLuint GetCurrentVaryingBuffer() const;
void UnbindTableTextures() const;
/// Returns an OsdVertexDescriptor if vertex buffers have been bound.
///
/// @return a descriptor for the format of the vertex data currently bound
///
OsdVertexDescriptor const & GetVertexDescriptor() const {
return _vdesc;
}
void BindEditTextures(
int editIndex,
OsdGLSLTransformFeedbackKernelBundle const *kernelBundle) const;
OsdGLSLTransformFeedbackKernelBundle * GetKernelBundle() const;
void SetKernelBundle(OsdGLSLTransformFeedbackKernelBundle *kernelBundle);
void BindEditTextures(int editIndex);
void UnbindEditTextures();
void UnbindEditTextures() const;
protected:
explicit OsdGLSLTransformFeedbackComputeContext(FarSubdivisionTables const *subdivisionTables,
FarVertexEditTables const *vertexEditTabes);
void bindTexture(GLint samplerUniform, GLuint texture, int unit);
void bindTexture(GLint samplerUniform, GLuint texture, int unit) const;
void unbindTexture(GLuint unit);
void bind();
void unbind();
void unbindTexture(GLuint unit) const;
private:
std::vector<OsdGLSLTransformFeedbackTable*> _tables;
std::vector<OsdGLSLTransformFeedbackHEditTable*> _editTables;
GLuint _vertexTexture,
_varyingTexture;
OsdVertexDescriptor _vdesc;
GLuint _currentVertexBuffer,
_currentVaryingBuffer;
OsdGLSLTransformFeedbackKernelBundle * _kernelBundle;
};
} // end namespace OPENSUBDIV_VERSION

238
opensubdiv/osd/glslTransformFeedbackComputeController.cpp
View File

@ -34,7 +34,10 @@
namespace OpenSubdiv {
namespace OPENSUBDIV_VERSION {
OsdGLSLTransformFeedbackComputeController::OsdGLSLTransformFeedbackComputeController() {
OsdGLSLTransformFeedbackComputeController::OsdGLSLTransformFeedbackComputeController() :
_vertexTexture(0), _varyingTexture(0),
_currentVertexBuffer(0), _currentVaryingBuffer(0),
_currentKernelBundle(NULL) {
}
OsdGLSLTransformFeedbackComputeController::~OsdGLSLTransformFeedbackComputeController() {
@ -44,6 +47,8 @@ OsdGLSLTransformFeedbackComputeController::~OsdGLSLTransformFeedbackComputeContr
it != _kernelRegistry.end(); ++it) {
delete *it;
}
if (_vertexTexture) glDeleteTextures(1, &_vertexTexture);
if (_varyingTexture) glDeleteTextures(1, &_varyingTexture);
}
void
@ -70,67 +75,124 @@ OsdGLSLTransformFeedbackComputeController::getKernels(int numVertexElements,
}
}
static void
bindTexture(GLint samplerUniform, GLuint texture, int unit) {
if (samplerUniform == -1) return;
glUniform1i(samplerUniform, unit);
glActiveTexture(GL_TEXTURE0 + unit);
glBindTexture(GL_TEXTURE_BUFFER, texture);
glActiveTexture(GL_TEXTURE0);
}
void
OsdGLSLTransformFeedbackComputeController::bindTextures() {
glEnable(GL_RASTERIZER_DISCARD);
_currentKernelBundle->UseProgram();
// bind vertex texture
if (_currentVertexBuffer) {
if (not _vertexTexture) glGenTextures(1, &_vertexTexture);
#if defined(GL_EXT_direct_state_access)
if (glTextureBufferEXT) {
glTextureBufferEXT(_vertexTexture, GL_TEXTURE_BUFFER, GL_R32F, _currentVertexBuffer);
} else {
#else
{
#endif
glBindTexture(GL_TEXTURE_BUFFER, _vertexTexture);
glTexBuffer(GL_TEXTURE_BUFFER, GL_R32F, _currentVertexBuffer);
glBindTexture(GL_TEXTURE_BUFFER, 0);
}
}
if (_currentVaryingBuffer) {
if (not _varyingTexture) glGenTextures(1, &_varyingTexture);
#if defined(GL_EXT_direct_state_access)
if (glTextureBufferEXT) {
glTextureBufferEXT(_varyingTexture, GL_TEXTURE_BUFFER, GL_R32F, _currentVaryingBuffer);
} else {
#else
{
#endif
glBindTexture(GL_TEXTURE_BUFFER, _varyingTexture);
glTexBuffer(GL_TEXTURE_BUFFER, GL_R32F, _currentVaryingBuffer);
glBindTexture(GL_TEXTURE_BUFFER, 0);
}
}
if (_vertexTexture)
bindTexture(_currentKernelBundle->GetVertexUniformLocation(), _vertexTexture, 0);
if (_varyingTexture)
bindTexture(_currentKernelBundle->GetVaryingUniformLocation(), _varyingTexture, 1);
// bind vertex texture image (for edit kernel)
glUniform1i(_currentKernelBundle->GetVertexBufferImageUniformLocation(), 0);
glBindImageTexture(0, _vertexTexture, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_R32F);
}
void
OsdGLSLTransformFeedbackComputeController::unbindTextures() {
glActiveTexture(GL_TEXTURE0);
glBindTexture(GL_TEXTURE_BUFFER, 0);
glActiveTexture(GL_TEXTURE1);
glBindTexture(GL_TEXTURE_BUFFER, 0);
// unbind vertex texture image
glBindImageTexture(0, 0, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_R32F);
glDisable(GL_RASTERIZER_DISCARD);
glUseProgram(0);
glActiveTexture(GL_TEXTURE0);
}
void
OsdGLSLTransformFeedbackComputeController::ApplyBilinearFaceVerticesKernel(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyBilinearFaceVerticesKernel(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyBilinearFaceVerticesKernel(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLTransformFeedbackComputeController::ApplyBilinearEdgeVerticesKernel(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyBilinearEdgeVerticesKernel(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyBilinearEdgeVerticesKernel(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLTransformFeedbackComputeController::ApplyBilinearVertexVerticesKernel(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyBilinearVertexVerticesKernel(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyBilinearVertexVerticesKernel(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLTransformFeedbackComputeController::ApplyCatmarkFaceVerticesKernel(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkFaceVerticesKernel(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyCatmarkFaceVerticesKernel(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
@ -138,154 +200,118 @@ OsdGLSLTransformFeedbackComputeController::ApplyCatmarkFaceVerticesKernel(
void
OsdGLSLTransformFeedbackComputeController::ApplyCatmarkEdgeVerticesKernel(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkEdgeVerticesKernel(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyCatmarkEdgeVerticesKernel(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLTransformFeedbackComputeController::ApplyCatmarkVertexVerticesKernelB(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkVertexVerticesKernelB(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyCatmarkVertexVerticesKernelB(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLTransformFeedbackComputeController::ApplyCatmarkVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkVertexVerticesKernelA(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyCatmarkVertexVerticesKernelA(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
}
void
OsdGLSLTransformFeedbackComputeController::ApplyCatmarkVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyCatmarkVertexVerticesKernelA(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyCatmarkVertexVerticesKernelA(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
}
void
OsdGLSLTransformFeedbackComputeController::ApplyLoopEdgeVerticesKernel(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopEdgeVerticesKernel(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyLoopEdgeVerticesKernel(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLTransformFeedbackComputeController::ApplyLoopVertexVerticesKernelB(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopVertexVerticesKernelB(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyLoopVertexVerticesKernelB(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdGLSLTransformFeedbackComputeController::ApplyLoopVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopVertexVerticesKernelA(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyLoopVertexVerticesKernelA(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
}
void
OsdGLSLTransformFeedbackComputeController::ApplyLoopVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
kernelBundle->ApplyLoopVertexVerticesKernelA(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyLoopVertexVerticesKernelA(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
}
void
OsdGLSLTransformFeedbackComputeController::ApplyVertexEdits(
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext *context) const {
FarKernelBatch const &batch, OsdGLSLTransformFeedbackComputeContext const *context) const {
assert(context);
OsdGLSLTransformFeedbackKernelBundle * kernelBundle = context->GetKernelBundle();
const OsdGLSLTransformFeedbackHEditTable * edit = context->GetEditTable(batch.GetTableIndex());
assert(edit);
context->BindEditTextures(batch.GetTableIndex());
context->BindEditTextures(batch.GetTableIndex(), _currentKernelBundle);
int primvarOffset = edit->GetPrimvarOffset();
int primvarWidth = edit->GetPrimvarWidth();
if (edit->GetOperation() == FarVertexEdit::Add) {
kernelBundle->ApplyEditAdd(
context->GetCurrentVertexBuffer(),
context->GetVertexDescriptor().numVertexElements,
context->GetCurrentVaryingBuffer(),
context->GetVertexDescriptor().numVaryingElements,
_currentKernelBundle->ApplyEditAdd(
_currentVertexBuffer, _vdesc.numVertexElements,
_currentVaryingBuffer, _vdesc.numVaryingElements,
primvarOffset, primvarWidth,
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
} else {

77
opensubdiv/osd/glslTransformFeedbackComputeController.h
View File

@ -70,23 +70,20 @@ public:
/// @param varyingBuffer varying-interpolated data buffer
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Refine(OsdGLSLTransformFeedbackComputeContext *context,
void Refine(OsdGLSLTransformFeedbackComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer,
VARYING_BUFFER *varyingBuffer) {
if (batches.empty()) return;
int numVertexElements = vertexBuffer ? vertexBuffer->GetNumElements() : 0;
int numVaryingElements = varyingBuffer ? varyingBuffer->GetNumElements() : 0;
context->SetKernelBundle(getKernels(numVertexElements, numVaryingElements));
context->Bind(vertexBuffer, varyingBuffer);
bind(vertexBuffer, varyingBuffer);
context->BindTableTextures(_currentKernelBundle);
FarDispatcher::Refine(this, context, batches, /*maxlevel*/-1);
context->Unbind();
context->UnbindTableTextures();
unbind();
}
/// Launch subdivision kernels and apply to given vertex buffers.
@ -99,7 +96,7 @@ public:
/// @param vertexBuffer vertex-interpolated data buffer
///
template<class VERTEX_BUFFER>
void Refine(OsdGLSLTransformFeedbackComputeContext *context,
void Refine(OsdGLSLTransformFeedbackComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer) {
Refine(context, batches, vertexBuffer, (VERTEX_BUFFER*)NULL);
@ -110,40 +107,76 @@ public:
protected:
friend class FarDispatcher;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext const *context) const;
OsdGLSLTransformFeedbackKernelBundle * getKernels(int numVertexElements,
int numVaryingElements);
void bindTextures();
void unbindTextures();
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
_currentVertexBuffer = vertex ? vertex->BindVBO() : 0;
_currentVaryingBuffer = varying ? varying->BindVBO() : 0;
_vdesc.numVertexElements = vertex ? vertex->GetNumElements() : 0;
_vdesc.numVaryingElements = varying ? varying->GetNumElements() : 0;
_currentKernelBundle =
getKernels(_vdesc.numVertexElements, _vdesc.numVaryingElements);
bindTextures();
}
/// Unbinds any previously bound vertex and varying data buffers.
void unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
_currentKernelBundle = NULL;
unbindTextures();
}
private:
std::vector<OsdGLSLTransformFeedbackKernelBundle *> _kernelRegistry;
GLuint _vertexTexture, _varyingTexture;
GLuint _currentVertexBuffer, _currentVaryingBuffer;
OsdVertexDescriptor _vdesc;
OsdGLSLTransformFeedbackKernelBundle * _currentKernelBundle;
};
} // end namespace OPENSUBDIV_VERSION

87
opensubdiv/osd/ompComputeController.cpp
View File

@ -34,22 +34,21 @@ namespace OpenSubdiv {
namespace OPENSUBDIV_VERSION {
OsdOmpComputeController::OsdOmpComputeController(int numThreads) {
OsdOmpComputeController::OsdOmpComputeController(int numThreads) :
_currentVertexBuffer(NULL), _currentVaryingBuffer(NULL) {
_numThreads = (numThreads == -1) ? omp_get_num_procs() : numThreads;
_numThreads = (numThreads == -1) ? omp_get_max_threads() : numThreads;
}
void
OsdOmpComputeController::ApplyBilinearFaceVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeFace(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::F_IT)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::F_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -57,42 +56,36 @@ OsdOmpComputeController::ApplyBilinearFaceVerticesKernel(
void
OsdOmpComputeController::ApplyBilinearEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeBilinearEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdOmpComputeController::ApplyBilinearVertexVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeBilinearVertex(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdOmpComputeController::ApplyCatmarkFaceVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeFace(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::F_IT)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::F_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -100,14 +93,12 @@ OsdOmpComputeController::ApplyCatmarkFaceVerticesKernel(
void
OsdOmpComputeController::ApplyCatmarkEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::E_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -115,14 +106,12 @@ OsdOmpComputeController::ApplyCatmarkEdgeVerticesKernel(
void
OsdOmpComputeController::ApplyCatmarkVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeVertexB(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::V_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
@ -131,14 +120,12 @@ OsdOmpComputeController::ApplyCatmarkVertexVerticesKernelB(
void
OsdOmpComputeController::ApplyCatmarkVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
@ -146,14 +133,12 @@ OsdOmpComputeController::ApplyCatmarkVertexVerticesKernelA1(
void
OsdOmpComputeController::ApplyCatmarkVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
@ -161,14 +146,12 @@ OsdOmpComputeController::ApplyCatmarkVertexVerticesKernelA2(
void
OsdOmpComputeController::ApplyLoopEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::E_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -176,14 +159,12 @@ OsdOmpComputeController::ApplyLoopEdgeVerticesKernel(
void
OsdOmpComputeController::ApplyLoopVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeLoopVertexB(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::V_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
@ -192,14 +173,12 @@ OsdOmpComputeController::ApplyLoopVertexVerticesKernelB(
void
OsdOmpComputeController::ApplyLoopVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
@ -207,14 +186,12 @@ OsdOmpComputeController::ApplyLoopVertexVerticesKernelA1(
void
OsdOmpComputeController::ApplyLoopVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdOmpComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
@ -222,7 +199,7 @@ OsdOmpComputeController::ApplyLoopVertexVerticesKernelA2(
void
OsdOmpComputeController::ApplyVertexEdits(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
@ -233,8 +210,8 @@ OsdOmpComputeController::ApplyVertexEdits(
const OsdCpuTable * editValues = edit->GetEditValues();
if (edit->GetOperation() == FarVertexEdit::Add) {
OsdOmpEditVertexAdd(context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
OsdOmpEditVertexAdd(_vdesc,
_currentVertexBuffer,
edit->GetPrimvarOffset(),
edit->GetPrimvarWidth(),
batch.GetVertexOffset(),
@ -244,8 +221,8 @@ OsdOmpComputeController::ApplyVertexEdits(
static_cast<unsigned int*>(primvarIndices->GetBuffer()),
static_cast<float*>(editValues->GetBuffer()));
} else if (edit->GetOperation() == FarVertexEdit::Set) {
OsdOmpEditVertexSet(context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
OsdOmpEditVertexSet(_vdesc,
_currentVertexBuffer,
edit->GetPrimvarOffset(),
edit->GetPrimvarWidth(),
batch.GetVertexOffset(),

53
opensubdiv/osd/ompComputeController.h
View File

@ -70,7 +70,7 @@ public:
/// @param varyingBuffer varying-interpolated data buffer
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Refine(OsdCpuComputeContext *context,
void Refine(OsdCpuComputeContext const *context,
FarKernelBatchVector const & batches,
VERTEX_BUFFER * vertexBuffer,
VARYING_BUFFER * varyingBuffer) {
@ -79,11 +79,11 @@ public:
omp_set_num_threads(_numThreads);
context->Bind(vertexBuffer, varyingBuffer);
bind(vertexBuffer, varyingBuffer);
FarDispatcher::Refine(this, context, batches, /*maxlevel*/-1);
context->Unbind();
unbind();
}
/// Launch subdivision kernels and apply to given vertex buffers.
@ -96,7 +96,7 @@ public:
/// @param vertexBuffer vertex-interpolated data buffer
///
template<class VERTEX_BUFFER>
void Refine(OsdCpuComputeContext *context,
void Refine(OsdCpuComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer) {
Refine(context, batches, vertexBuffer, (VERTEX_BUFFER*)0);
@ -108,34 +108,53 @@ public:
protected:
friend class FarDispatcher;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext const *context) const;
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
_currentVertexBuffer = vertex ? vertex->BindCpuBuffer() : 0;
_currentVaryingBuffer = varying ? varying->BindCpuBuffer() : 0;
int numVertexElements = vertex ? vertex->GetNumElements() : 0;
int numVaryingElements = varying ? varying->GetNumElements() : 0;
_vdesc.Set(numVertexElements, numVaryingElements);
}
void unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
_vdesc.Reset();
}
private:
float *_currentVertexBuffer, *_currentVaryingBuffer;
OsdVertexDescriptor _vdesc;
int _numThreads;
};

89
opensubdiv/osd/tbbComputeController.cpp
View File

@ -36,8 +36,11 @@ namespace OpenSubdiv {
namespace OPENSUBDIV_VERSION {
OsdTbbComputeController::OsdTbbComputeController(int numThreads) {
_numThreads = numThreads;
OsdTbbComputeController::OsdTbbComputeController(int numThreads)
: _currentVertexBuffer(NULL),
_currentVaryingBuffer(NULL),
_numThreads(numThreads) {
if(_numThreads == -1)
tbb::task_scheduler_init init;
else
@ -47,14 +50,12 @@ OsdTbbComputeController::OsdTbbComputeController(int numThreads) {
void
OsdTbbComputeController::ApplyBilinearFaceVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeFace(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::F_IT)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::F_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -62,42 +63,36 @@ OsdTbbComputeController::ApplyBilinearFaceVerticesKernel(
void
OsdTbbComputeController::ApplyBilinearEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeBilinearEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdTbbComputeController::ApplyBilinearVertexVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeBilinearVertex(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
}
void
OsdTbbComputeController::ApplyCatmarkFaceVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeFace(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::F_IT)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::F_ITa)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -105,14 +100,12 @@ OsdTbbComputeController::ApplyCatmarkFaceVerticesKernel(
void
OsdTbbComputeController::ApplyCatmarkEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::E_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -120,14 +113,12 @@ OsdTbbComputeController::ApplyCatmarkEdgeVerticesKernel(
void
OsdTbbComputeController::ApplyCatmarkVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeVertexB(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::V_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
@ -136,14 +127,12 @@ OsdTbbComputeController::ApplyCatmarkVertexVerticesKernelB(
void
OsdTbbComputeController::ApplyCatmarkVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
@ -151,14 +140,12 @@ OsdTbbComputeController::ApplyCatmarkVertexVerticesKernelA1(
void
OsdTbbComputeController::ApplyCatmarkVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
@ -166,14 +153,12 @@ OsdTbbComputeController::ApplyCatmarkVertexVerticesKernelA2(
void
OsdTbbComputeController::ApplyLoopEdgeVerticesKernel(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeEdge(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::E_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::E_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd());
@ -181,14 +166,12 @@ OsdTbbComputeController::ApplyLoopEdgeVerticesKernel(
void
OsdTbbComputeController::ApplyLoopVertexVerticesKernelB(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeLoopVertexB(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const int*)context->GetTable(FarSubdivisionTables::V_IT)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
@ -197,14 +180,12 @@ OsdTbbComputeController::ApplyLoopVertexVerticesKernelB(
void
OsdTbbComputeController::ApplyLoopVertexVerticesKernelA1(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), false);
@ -212,14 +193,12 @@ OsdTbbComputeController::ApplyLoopVertexVerticesKernelA1(
void
OsdTbbComputeController::ApplyLoopVertexVerticesKernelA2(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
OsdTbbComputeVertexA(
context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
context->GetCurrentVaryingBuffer(),
_vdesc, _currentVertexBuffer, _currentVaryingBuffer,
(const int*)context->GetTable(FarSubdivisionTables::V_ITa)->GetBuffer(),
(const float*)context->GetTable(FarSubdivisionTables::V_W)->GetBuffer(),
batch.GetVertexOffset(), batch.GetTableOffset(), batch.GetStart(), batch.GetEnd(), true);
@ -227,7 +206,7 @@ OsdTbbComputeController::ApplyLoopVertexVerticesKernelA2(
void
OsdTbbComputeController::ApplyVertexEdits(
FarKernelBatch const &batch, OsdCpuComputeContext *context) const {
FarKernelBatch const &batch, OsdCpuComputeContext const *context) const {
assert(context);
@ -238,8 +217,8 @@ OsdTbbComputeController::ApplyVertexEdits(
const OsdCpuTable * editValues = edit->GetEditValues();
if (edit->GetOperation() == FarVertexEdit::Add) {
OsdTbbEditVertexAdd(context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
OsdTbbEditVertexAdd(_vdesc,
_currentVertexBuffer,
edit->GetPrimvarOffset(),
edit->GetPrimvarWidth(),
batch.GetVertexOffset(),
@ -249,8 +228,8 @@ OsdTbbComputeController::ApplyVertexEdits(
static_cast<unsigned int*>(primvarIndices->GetBuffer()),
static_cast<float*>(editValues->GetBuffer()));
} else if (edit->GetOperation() == FarVertexEdit::Set) {
OsdTbbEditVertexSet(context->GetVertexDescriptor(),
context->GetCurrentVertexBuffer(),
OsdTbbEditVertexSet(_vdesc,
_currentVertexBuffer,
edit->GetPrimvarOffset(),
edit->GetPrimvarWidth(),
batch.GetVertexOffset(),

53
opensubdiv/osd/tbbComputeController.h
View File

@ -66,16 +66,16 @@ public:
/// @param varyingBuffer varying-interpolated data buffer
///
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void Refine(OsdCpuComputeContext *context,
void Refine(OsdCpuComputeContext const *context,
FarKernelBatchVector const & batches,
VERTEX_BUFFER * vertexBuffer,
VARYING_BUFFER * varyingBuffer) {
context->Bind(vertexBuffer, varyingBuffer);
bind(vertexBuffer, varyingBuffer);
FarDispatcher::Refine(this, context, batches, /*maxlevel*/-1);
context->Unbind();
unbind();
}
/// Launch subdivision kernels and apply to given vertex buffers.
@ -88,7 +88,7 @@ public:
/// @param vertexBuffer vertex-interpolated data buffer
///
template<class VERTEX_BUFFER>
void Refine(OsdCpuComputeContext *context,
void Refine(OsdCpuComputeContext const *context,
FarKernelBatchVector const &batches,
VERTEX_BUFFER *vertexBuffer) {
Refine(context, batches, vertexBuffer, (VERTEX_BUFFER*)0);
@ -100,34 +100,53 @@ public:
protected:
friend class FarDispatcher;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyBilinearVertexVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkFaceVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyCatmarkVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopEdgeVerticesKernel(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelB(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA1(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyLoopVertexVerticesKernelA2(FarKernelBatch const &batch, ComputeContext const *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext *context) const;
void ApplyVertexEdits(FarKernelBatch const &batch, ComputeContext const *context) const;
private:
template<class VERTEX_BUFFER, class VARYING_BUFFER>
void bind(VERTEX_BUFFER *vertex, VARYING_BUFFER *varying) {
_currentVertexBuffer = vertex ? vertex->BindCpuBuffer() : 0;
_currentVaryingBuffer = varying ? varying->BindCpuBuffer() : 0;
int numVertexElements = vertex ? vertex->GetNumElements() : 0;
int numVaryingElements = varying ? varying->GetNumElements() : 0;
_vdesc.Set(numVertexElements, numVaryingElements);
}
void unbind() {
_currentVertexBuffer = 0;
_currentVaryingBuffer = 0;
_vdesc.Reset();
}
float *_currentVertexBuffer, *_currentVaryingBuffer;
OsdVertexDescriptor _vdesc;
int _numThreads;
};