AuroraMiddleware/OpenSubdiv
1
0
Fork 0
mirror of https://github.com/PixarAnimationStudios/OpenSubdiv synced 2024-09-19 14:20:00 +00:00
Code Issues Releases Wiki Activity

Merge pull request #924 from davidgyu/dev_osd_eval_deriv

Browse Source 
Updated Osd Evaluator methods for derivative eval
This commit is contained in:
George ElKoura 2017-01-27 18:38:41 -08:00 committed by GitHub
commit a410a52f4e
28 changed files with 9171 additions and 584 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

149
examples/glEvalLimit/glEvalLimit.cpp
View File

@ -121,7 +121,8 @@ enum DrawMode { kUV,
kVARYING,
kNORMAL,
kSHADE,
kFACEVARYING };
kFACEVARYING,
kMEAN_CURVATURE };
std::vector<float> g_orgPositions,
g_positions,
@ -192,8 +193,11 @@ struct Program {
GLuint uniformDrawMode;
GLuint attrPosition;
GLuint attrColor;
GLuint attrTangentU;
GLuint attrTangentV;
GLuint attrDu;
GLuint attrDv;
GLuint attrDuu;
GLuint attrDuv;
GLuint attrDvv;
GLuint attrPatchCoord;
GLuint attrFVarData;
} g_defaultProgram;
@ -221,7 +225,8 @@ public:
virtual ~EvalOutputBase() {}
virtual GLuint BindSourceData() const = 0;
virtual GLuint BindVertexData() const = 0;
virtual GLuint BindDerivatives() const = 0;
virtual GLuint Bind1stDerivatives() const = 0;
virtual GLuint Bind2ndDerivatives() const = 0;
virtual GLuint BindFaceVaryingData() const = 0;
virtual GLuint BindPatchCoords() const = 0;
virtual void UpdateData(const float *src, int startVertex, int numVertices) = 0;
@ -230,7 +235,8 @@ public:
virtual bool HasFaceVaryingData() const = 0;
virtual void Refine() = 0;
virtual void EvalPatches() = 0;
virtual void EvalPatchesWithDerivatives() = 0;
virtual void EvalPatchesWith1stDerivatives() = 0;
virtual void EvalPatchesWith2ndDerivatives() = 0;
virtual void EvalPatchesVarying() = 0;
virtual void EvalPatchesFaceVarying() = 0;
virtual void UpdatePatchCoords(
@ -264,6 +270,9 @@ public:
_fvarDesc( /*offset*/ 0, /*length*/ fvarWidth, /*stride*/ fvarWidth),
_duDesc( /*offset*/ 0, /*length*/ 3, /*stride*/ 6),
_dvDesc( /*offset*/ 3, /*length*/ 3, /*stride*/ 6),
_duuDesc( /*offset*/ 0, /*length*/ 3, /*stride*/ 9),
_duvDesc( /*offset*/ 3, /*length*/ 3, /*stride*/ 9),
_dvvDesc( /*offset*/ 6, /*length*/ 3, /*stride*/ 9),
_deviceContext(deviceContext) {
// total number of vertices = coarse points + refined points + local points
@ -273,7 +282,8 @@ public:
_srcData = SRC_VERTEX_BUFFER::Create(3, numTotalVerts, _deviceContext);
_srcVaryingData = SRC_VERTEX_BUFFER::Create(3, numTotalVerts, _deviceContext);
_vertexData = EVAL_VERTEX_BUFFER::Create(6, numParticles, _deviceContext);
_derivatives = EVAL_VERTEX_BUFFER::Create(6, numParticles, _deviceContext);
_deriv1 = EVAL_VERTEX_BUFFER::Create(6, numParticles, _deviceContext);
_deriv2 = EVAL_VERTEX_BUFFER::Create(9, numParticles, _deviceContext);
_patchTable = PATCH_TABLE::Create(patchTable, _deviceContext);
_patchCoords = NULL;
_numCoarseVerts = vertexStencils->GetNumControlVertices();
@ -307,7 +317,8 @@ public:
delete _srcVaryingData;
delete _srcFVarData;
delete _vertexData;
delete _derivatives;
delete _deriv1;
delete _deriv2;
delete _fvarData;
delete _patchTable;
delete _patchCoords;
@ -321,8 +332,11 @@ public:
virtual GLuint BindVertexData() const {
return _vertexData->BindVBO();
}
virtual GLuint BindDerivatives() const {
return _derivatives->BindVBO();
virtual GLuint Bind1stDerivatives() const {
return _deriv1->BindVBO();
}
virtual GLuint Bind2ndDerivatives() const {
return _deriv2->BindVBO();
}
virtual GLuint BindFaceVaryingData() const {
return _fvarData->BindVBO();
@ -392,14 +406,31 @@ public:
_patchCoords,
_patchTable, evalInstance, _deviceContext);
}
virtual void EvalPatchesWithDerivatives() {
virtual void EvalPatchesWith1stDerivatives() {
EVALUATOR const *evalInstance = OpenSubdiv::Osd::GetEvaluator<EVALUATOR>(
_evaluatorCache, _srcDesc, _vertexDesc, _duDesc, _dvDesc, _deviceContext);
EVALUATOR::EvalPatches(
_srcData, _srcDesc,
_vertexData, _vertexDesc,
_derivatives, _duDesc,
_derivatives, _dvDesc,
_deriv1, _duDesc,
_deriv1, _dvDesc,
_patchCoords->GetNumVertices(),
_patchCoords,
_patchTable, evalInstance, _deviceContext);
}
virtual void EvalPatchesWith2ndDerivatives() {
EVALUATOR const *evalInstance = OpenSubdiv::Osd::GetEvaluator<EVALUATOR>(
_evaluatorCache, _srcDesc, _vertexDesc,
_duDesc, _dvDesc, _duuDesc, _duvDesc, _dvvDesc,
_deviceContext);
EVALUATOR::EvalPatches(
_srcData, _srcDesc,
_vertexData, _vertexDesc,
_deriv1, _duDesc,
_deriv1, _dvDesc,
_deriv2, _duuDesc,
_deriv2, _duvDesc,
_deriv2, _dvvDesc,
_patchCoords->GetNumVertices(),
_patchCoords,
_patchTable, evalInstance, _deviceContext);
@ -446,7 +477,8 @@ private:
SRC_VERTEX_BUFFER *_srcVaryingData;
EVAL_VERTEX_BUFFER *_srcFVarData;
EVAL_VERTEX_BUFFER *_vertexData;
EVAL_VERTEX_BUFFER *_derivatives;
EVAL_VERTEX_BUFFER *_deriv1;
EVAL_VERTEX_BUFFER *_deriv2;
EVAL_VERTEX_BUFFER *_fvarData;
EVAL_VERTEX_BUFFER *_patchCoords;
PATCH_TABLE *_patchTable;
@ -458,6 +490,9 @@ private:
Osd::BufferDescriptor _fvarDesc;
Osd::BufferDescriptor _duDesc;
Osd::BufferDescriptor _dvDesc;
Osd::BufferDescriptor _duuDesc;
Osd::BufferDescriptor _duvDesc;
Osd::BufferDescriptor _dvvDesc;
int _numCoarseVerts;
int _numCoarseFVarVerts;
@ -532,9 +567,12 @@ updateGeom() {
g_evalOutput->UpdatePatchCoords(patchCoords);
// Evaluate the positions of the samples on the limit surface
if (g_drawMode == kNORMAL || g_drawMode == kSHADE) {
// evaluate positions and derivatives
g_evalOutput->EvalPatchesWithDerivatives();
if (g_drawMode == kMEAN_CURVATURE) {
// evaluate positions and 2nd derivatives
g_evalOutput->EvalPatchesWith2ndDerivatives();
} else if (g_drawMode == kNORMAL || g_drawMode == kSHADE) {
// evaluate positions and 1st derivatives
g_evalOutput->EvalPatchesWith1stDerivatives();
} else {
// evaluate positions
g_evalOutput->EvalPatches();
@ -823,8 +861,11 @@ linkDefaultProgram() {
GLSL_VERSION_DEFINE
"in vec3 position;\n"
"in vec3 color;\n"
"in vec3 tangentU;\n"
"in vec3 tangentV;\n"
"in vec3 du;\n"
"in vec3 dv;\n"
"in vec3 duu;\n"
"in vec3 duv;\n"
"in vec3 dvv;\n"
"in vec2 patchCoord;\n"
"in vec2 fvarData;\n"
"out vec4 fragColor;\n"
@ -833,7 +874,7 @@ linkDefaultProgram() {
"uniform int DrawMode;\n"
"void main() {\n"
" vec3 normal = (ModelViewMatrix * "
" vec4(normalize(cross(tangentU, tangentV)), 0)).xyz;\n"
" vec4(normalize(cross(du, dv)), 0)).xyz;\n"
" gl_Position = ProjectionMatrix * ModelViewMatrix * "
" vec4(position, 1);\n"
" if (DrawMode == 0) {\n" // UV
@ -846,6 +887,16 @@ linkDefaultProgram() {
" // generating a checkerboard pattern\n"
" int checker = int(floor(20*fvarData.r)+floor(20*fvarData.g))&1;\n"
" fragColor = vec4(fvarData.rg*checker, 1-checker, 1);\n"
" } else if (DrawMode == 5) {\n" // mean curvature
" vec3 N = normalize(cross(du, dv));\n"
" float E = dot(du, du);\n"
" float F = dot(du, dv);\n"
" float G = dot(dv, dv);\n"
" float e = dot(N, duu);\n"
" float f = dot(N, duv);\n"
" float g = dot(N, dvv);\n"
" float H = 0.5 * abs(0.5 * (E*g - 2*F*f - G*e) / (E*G - F*F));\n"
" fragColor = vec4(H, H, H, 1.0);\n"
" } else {\n" // varying
" fragColor = vec4(color, 1);\n"
" }\n"
@ -868,10 +919,13 @@ linkDefaultProgram() {
glBindAttribLocation(program, 0, "position");
glBindAttribLocation(program, 1, "color");
glBindAttribLocation(program, 2, "tangentU");
glBindAttribLocation(program, 3, "tangentV");
glBindAttribLocation(program, 4, "patchCoord");
glBindAttribLocation(program, 5, "fvarData");
glBindAttribLocation(program, 2, "du");
glBindAttribLocation(program, 3, "dv");
glBindAttribLocation(program, 4, "duu");
glBindAttribLocation(program, 5, "duv");
glBindAttribLocation(program, 6, "dvv");
glBindAttribLocation(program, 7, "patchCoord");
glBindAttribLocation(program, 8, "fvarData");
glBindFragDataLocation(program, 0, "color");
glLinkProgram(program);
@ -897,8 +951,11 @@ linkDefaultProgram() {
glGetUniformLocation(program, "DrawMode");
g_defaultProgram.attrPosition = glGetAttribLocation(program, "position");
g_defaultProgram.attrColor = glGetAttribLocation(program, "color");
g_defaultProgram.attrTangentU = glGetAttribLocation(program, "tangentU");
g_defaultProgram.attrTangentV = glGetAttribLocation(program, "tangentV");
g_defaultProgram.attrDu = glGetAttribLocation(program, "du");
g_defaultProgram.attrDv = glGetAttribLocation(program, "dv");
g_defaultProgram.attrDuu = glGetAttribLocation(program, "duu");
g_defaultProgram.attrDuv = glGetAttribLocation(program, "duv");
g_defaultProgram.attrDvv = glGetAttribLocation(program, "dvv");
g_defaultProgram.attrPatchCoord = glGetAttribLocation(program, "patchCoord");
g_defaultProgram.attrFVarData = glGetAttribLocation(program, "fvarData");
@ -920,36 +977,32 @@ drawSamples() {
glEnableVertexAttribArray(g_defaultProgram.attrPosition);
glEnableVertexAttribArray(g_defaultProgram.attrColor);
glEnableVertexAttribArray(g_defaultProgram.attrTangentU);
glEnableVertexAttribArray(g_defaultProgram.attrTangentV);
glEnableVertexAttribArray(g_defaultProgram.attrPatchCoord);
if (g_evalOutput->HasFaceVaryingData()) {
glEnableVertexAttribArray(g_defaultProgram.attrFVarData);
}
glBindBuffer(GL_ARRAY_BUFFER, g_evalOutput->BindVertexData());
glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 6, 0);
glVertexAttribPointer(1, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 6, (float*)12);
glBindBuffer(GL_ARRAY_BUFFER, g_evalOutput->BindDerivatives());
glEnableVertexAttribArray(g_defaultProgram.attrDu);
glEnableVertexAttribArray(g_defaultProgram.attrDv);
glBindBuffer(GL_ARRAY_BUFFER, g_evalOutput->Bind1stDerivatives());
glVertexAttribPointer(2, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 6, 0);
glVertexAttribPointer(3, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 6, (float*)12);
glBindBuffer(GL_ARRAY_BUFFER, g_evalOutput->BindPatchCoords());
glVertexAttribPointer(4, 2, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 5, (float*)12);
glEnableVertexAttribArray(g_defaultProgram.attrDuu);
glEnableVertexAttribArray(g_defaultProgram.attrDuv);
glEnableVertexAttribArray(g_defaultProgram.attrDvv);
glBindBuffer(GL_ARRAY_BUFFER, g_evalOutput->Bind2ndDerivatives());
glVertexAttribPointer(4, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 9, 0);
glVertexAttribPointer(5, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 9, (float*)12);
glVertexAttribPointer(6, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 9, (float*)24);
if (g_evalOutput->HasFaceVaryingData()) {
glBindBuffer(GL_ARRAY_BUFFER, g_evalOutput->BindFaceVaryingData());
glVertexAttribPointer(5, 2, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 2, 0);
}
glEnableVertexAttribArray(g_defaultProgram.attrPosition);
glEnableVertexAttribArray(g_defaultProgram.attrColor);
glEnableVertexAttribArray(g_defaultProgram.attrTangentU);
glEnableVertexAttribArray(g_defaultProgram.attrTangentV);
glEnableVertexAttribArray(g_defaultProgram.attrPatchCoord);
glBindBuffer(GL_ARRAY_BUFFER, g_evalOutput->BindPatchCoords());
glVertexAttribPointer(7, 2, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 5, (float*)12);
if (g_evalOutput->HasFaceVaryingData()) {
glEnableVertexAttribArray(g_defaultProgram.attrFVarData);
glBindBuffer(GL_ARRAY_BUFFER, g_evalOutput->BindFaceVaryingData());
glVertexAttribPointer(8, 2, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 2, 0);
}
glPointSize(2.0f);
@ -959,8 +1012,11 @@ drawSamples() {
glDisableVertexAttribArray(g_defaultProgram.attrPosition);
glDisableVertexAttribArray(g_defaultProgram.attrColor);
glDisableVertexAttribArray(g_defaultProgram.attrTangentU);
glDisableVertexAttribArray(g_defaultProgram.attrTangentV);
glDisableVertexAttribArray(g_defaultProgram.attrDu);
glDisableVertexAttribArray(g_defaultProgram.attrDv);
glDisableVertexAttribArray(g_defaultProgram.attrDuu);
glDisableVertexAttribArray(g_defaultProgram.attrDuv);
glDisableVertexAttribArray(g_defaultProgram.attrDvv);
glDisableVertexAttribArray(g_defaultProgram.attrPatchCoord);
glDisableVertexAttribArray(g_defaultProgram.attrFVarData);
@ -1323,6 +1379,7 @@ initHUD() {
g_hud.AddPullDownButton(shading_pulldown, "Normal", kNORMAL, g_drawMode==kNORMAL);
g_hud.AddPullDownButton(shading_pulldown, "Shade", kSHADE, g_drawMode==kSHADE);
g_hud.AddPullDownButton(shading_pulldown, "FaceVarying", kFACEVARYING, g_drawMode==kFACEVARYING);
g_hud.AddPullDownButton(shading_pulldown, "Mean Curvature", kMEAN_CURVATURE, g_drawMode==kMEAN_CURVATURE);
for (int i = 1; i < 11; ++i) {
char level[16];

204
opensubdiv/osd/clEvaluator.cpp
View File

@ -49,6 +49,10 @@ static const char *patchBasisSource =
template <class T> cl_mem
createCLBuffer(std::vector<T> const & src, cl_context clContext) {
if (src.empty()) {
return NULL;
}
cl_int errNum = 0;
cl_mem devicePtr = clCreateBuffer(clContext,
CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
@ -76,9 +80,11 @@ CLStencilTable::CLStencilTable(Far::StencilTable const *stencilTable,
clContext);
_weights = createCLBuffer(stencilTable->GetWeights(), clContext);
_duWeights = _dvWeights = NULL;
_duuWeights = _duvWeights = _dvvWeights = NULL;
} else {
_sizes = _offsets = _indices = _weights = NULL;
_duWeights = _dvWeights = NULL;
_duuWeights = _duvWeights = _dvvWeights = NULL;
}
}
@ -96,9 +102,16 @@ CLStencilTable::CLStencilTable(Far::LimitStencilTable const *limitStencilTable,
limitStencilTable->GetDuWeights(), clContext);
_dvWeights = createCLBuffer(
limitStencilTable->GetDvWeights(), clContext);
_duuWeights = createCLBuffer(
limitStencilTable->GetDuuWeights(), clContext);
_duvWeights = createCLBuffer(
limitStencilTable->GetDuvWeights(), clContext);
_dvvWeights = createCLBuffer(
limitStencilTable->GetDvvWeights(), clContext);
} else {
_sizes = _offsets = _indices = _weights = NULL;
_duWeights = _dvWeights = NULL;
_duuWeights = _duvWeights = _dvvWeights = NULL;
}
}
@ -109,6 +122,9 @@ CLStencilTable::~CLStencilTable() {
if (_weights) clReleaseMemObject(_weights);
if (_duWeights) clReleaseMemObject(_duWeights);
if (_dvWeights) clReleaseMemObject(_dvWeights);
if (_duuWeights) clReleaseMemObject(_duuWeights);
if (_duvWeights) clReleaseMemObject(_duvWeights);
if (_dvvWeights) clReleaseMemObject(_dvvWeights);
}
// ---------------------------------------------------------------------------
@ -130,7 +146,10 @@ bool
CLEvaluator::Compile(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const & /*duDesc*/,
BufferDescriptor const & /*dvDesc*/) {
BufferDescriptor const & /*dvDesc*/,
BufferDescriptor const & /*duuDesc*/,
BufferDescriptor const & /*duvDesc*/,
BufferDescriptor const & /*dvvDesc*/) {
if (srcDesc.length > dstDesc.length) {
Far::Error(Far::FAR_RUNTIME_ERROR,
"srcDesc length must be less than or equal to "
@ -263,6 +282,7 @@ CLEvaluator::EvalStencils(cl_mem src, BufferDescriptor const &srcDesc,
size_t globalWorkSize = (size_t)(end - start);
BufferDescriptor empty;
clSetKernelArg(_stencilDerivKernel, 0, sizeof(cl_mem), &src);
clSetKernelArg(_stencilDerivKernel, 1, sizeof(int), &srcDesc.offset);
clSetKernelArg(_stencilDerivKernel, 2, sizeof(cl_mem), &dst);
@ -273,14 +293,26 @@ CLEvaluator::EvalStencils(cl_mem src, BufferDescriptor const &srcDesc,
clSetKernelArg(_stencilDerivKernel, 7, sizeof(cl_mem), &dv);
clSetKernelArg(_stencilDerivKernel, 8, sizeof(int), &dvDesc.offset);
clSetKernelArg(_stencilDerivKernel, 9, sizeof(int), &dvDesc.stride);
clSetKernelArg(_stencilDerivKernel, 10, sizeof(cl_mem), &sizes);
clSetKernelArg(_stencilDerivKernel, 11, sizeof(cl_mem), &offsets);
clSetKernelArg(_stencilDerivKernel, 12, sizeof(cl_mem), &indices);
clSetKernelArg(_stencilDerivKernel, 13, sizeof(cl_mem), &weights);
clSetKernelArg(_stencilDerivKernel, 14, sizeof(cl_mem), &duWeights);
clSetKernelArg(_stencilDerivKernel, 15, sizeof(cl_mem), &dvWeights);
clSetKernelArg(_stencilDerivKernel, 16, sizeof(int), &start);
clSetKernelArg(_stencilDerivKernel, 17, sizeof(int), &end);
clSetKernelArg(_stencilDerivKernel, 10, sizeof(cl_mem), NULL);
clSetKernelArg(_stencilDerivKernel, 11, sizeof(int), &empty.offset);
clSetKernelArg(_stencilDerivKernel, 12, sizeof(int), &empty.stride);
clSetKernelArg(_stencilDerivKernel, 13, sizeof(cl_mem), NULL);
clSetKernelArg(_stencilDerivKernel, 14, sizeof(int), &empty.offset);
clSetKernelArg(_stencilDerivKernel, 15, sizeof(int), &empty.stride);
clSetKernelArg(_stencilDerivKernel, 16, sizeof(cl_mem), NULL);
clSetKernelArg(_stencilDerivKernel, 17, sizeof(int), &empty.offset);
clSetKernelArg(_stencilDerivKernel, 18, sizeof(int), &empty.stride);
clSetKernelArg(_stencilDerivKernel, 19, sizeof(cl_mem), &sizes);
clSetKernelArg(_stencilDerivKernel, 20, sizeof(cl_mem), &offsets);
clSetKernelArg(_stencilDerivKernel, 21, sizeof(cl_mem), &indices);
clSetKernelArg(_stencilDerivKernel, 22, sizeof(cl_mem), &weights);
clSetKernelArg(_stencilDerivKernel, 23, sizeof(cl_mem), &duWeights);
clSetKernelArg(_stencilDerivKernel, 24, sizeof(cl_mem), &dvWeights);
clSetKernelArg(_stencilDerivKernel, 25, sizeof(cl_mem), NULL);
clSetKernelArg(_stencilDerivKernel, 26, sizeof(cl_mem), NULL);
clSetKernelArg(_stencilDerivKernel, 27, sizeof(cl_mem), NULL);
clSetKernelArg(_stencilDerivKernel, 28, sizeof(int), &start);
clSetKernelArg(_stencilDerivKernel, 29, sizeof(int), &end);
cl_int errNum = clEnqueueNDRangeKernel(
_clCommandQueue, _stencilDerivKernel, 1, NULL,
@ -292,8 +324,79 @@ CLEvaluator::EvalStencils(cl_mem src, BufferDescriptor const &srcDesc,
return false;
}
if (endEvent == NULL)
{
if (endEvent == NULL) {
clFinish(_clCommandQueue);
}
return true;
}
bool
CLEvaluator::EvalStencils(cl_mem src, BufferDescriptor const &srcDesc,
cl_mem dst, BufferDescriptor const &dstDesc,
cl_mem du, BufferDescriptor const &duDesc,
cl_mem dv, BufferDescriptor const &dvDesc,
cl_mem duu, BufferDescriptor const &duuDesc,
cl_mem duv, BufferDescriptor const &duvDesc,
cl_mem dvv, BufferDescriptor const &dvvDesc,
cl_mem sizes,
cl_mem offsets,
cl_mem indices,
cl_mem weights,
cl_mem duWeights,
cl_mem dvWeights,
cl_mem duuWeights,
cl_mem duvWeights,
cl_mem dvvWeights,
int start, int end,
unsigned int numStartEvents,
const cl_event* startEvents,
cl_event* endEvent) const {
if (end <= start) return true;
size_t globalWorkSize = (size_t)(end - start);
clSetKernelArg(_stencilDerivKernel, 0, sizeof(cl_mem), &src);
clSetKernelArg(_stencilDerivKernel, 1, sizeof(int), &srcDesc.offset);
clSetKernelArg(_stencilDerivKernel, 2, sizeof(cl_mem), &dst);
clSetKernelArg(_stencilDerivKernel, 3, sizeof(int), &dstDesc.offset);
clSetKernelArg(_stencilDerivKernel, 4, sizeof(cl_mem), &du);
clSetKernelArg(_stencilDerivKernel, 5, sizeof(int), &duDesc.offset);
clSetKernelArg(_stencilDerivKernel, 6, sizeof(int), &duDesc.stride);
clSetKernelArg(_stencilDerivKernel, 7, sizeof(cl_mem), &dv);
clSetKernelArg(_stencilDerivKernel, 8, sizeof(int), &dvDesc.offset);
clSetKernelArg(_stencilDerivKernel, 9, sizeof(int), &dvDesc.stride);
clSetKernelArg(_stencilDerivKernel, 10, sizeof(cl_mem), &duu);
clSetKernelArg(_stencilDerivKernel, 11, sizeof(int), &duuDesc.offset);
clSetKernelArg(_stencilDerivKernel, 12, sizeof(int), &duuDesc.stride);
clSetKernelArg(_stencilDerivKernel, 13, sizeof(cl_mem), &duv);
clSetKernelArg(_stencilDerivKernel, 14, sizeof(int), &duvDesc.offset);
clSetKernelArg(_stencilDerivKernel, 15, sizeof(int), &duvDesc.stride);
clSetKernelArg(_stencilDerivKernel, 16, sizeof(cl_mem), &dvv);
clSetKernelArg(_stencilDerivKernel, 17, sizeof(int), &dvvDesc.offset);
clSetKernelArg(_stencilDerivKernel, 18, sizeof(int), &dvvDesc.stride);
clSetKernelArg(_stencilDerivKernel, 19, sizeof(cl_mem), &sizes);
clSetKernelArg(_stencilDerivKernel, 20, sizeof(cl_mem), &offsets);
clSetKernelArg(_stencilDerivKernel, 21, sizeof(cl_mem), &indices);
clSetKernelArg(_stencilDerivKernel, 22, sizeof(cl_mem), &weights);
clSetKernelArg(_stencilDerivKernel, 23, sizeof(cl_mem), &duWeights);
clSetKernelArg(_stencilDerivKernel, 24, sizeof(cl_mem), &dvWeights);
clSetKernelArg(_stencilDerivKernel, 25, sizeof(cl_mem), &duuWeights);
clSetKernelArg(_stencilDerivKernel, 26, sizeof(cl_mem), &duvWeights);
clSetKernelArg(_stencilDerivKernel, 27, sizeof(cl_mem), &dvvWeights);
clSetKernelArg(_stencilDerivKernel, 28, sizeof(int), &start);
clSetKernelArg(_stencilDerivKernel, 29, sizeof(int), &end);
cl_int errNum = clEnqueueNDRangeKernel(
_clCommandQueue, _stencilDerivKernel, 1, NULL,
&globalWorkSize, NULL, numStartEvents, startEvents, endEvent);
if (errNum != CL_SUCCESS) {
Far::Error(Far::FAR_RUNTIME_ERROR,
"ApplyStencilKernel (%d) ", errNum);
return false;
}
if (endEvent == NULL) {
clFinish(_clCommandQueue);
}
return true;
@ -315,6 +418,7 @@ CLEvaluator::EvalPatches(cl_mem src, BufferDescriptor const &srcDesc,
size_t globalWorkSize = (size_t)(numPatchCoords);
BufferDescriptor empty;
clSetKernelArg(_patchKernel, 0, sizeof(cl_mem), &src);
clSetKernelArg(_patchKernel, 1, sizeof(int), &srcDesc.offset);
clSetKernelArg(_patchKernel, 2, sizeof(cl_mem), &dst);
@ -325,10 +429,19 @@ CLEvaluator::EvalPatches(cl_mem src, BufferDescriptor const &srcDesc,
clSetKernelArg(_patchKernel, 7, sizeof(cl_mem), &dv);
clSetKernelArg(_patchKernel, 8, sizeof(int), &dvDesc.offset);
clSetKernelArg(_patchKernel, 9, sizeof(int), &dvDesc.stride);
clSetKernelArg(_patchKernel, 10, sizeof(cl_mem), &patchCoordsBuffer);
clSetKernelArg(_patchKernel, 11, sizeof(cl_mem), &patchArrayBuffer);
clSetKernelArg(_patchKernel, 12, sizeof(cl_mem), &patchIndexBuffer);
clSetKernelArg(_patchKernel, 13, sizeof(cl_mem), &patchParamBuffer);
clSetKernelArg(_patchKernel, 10, sizeof(cl_mem), NULL);
clSetKernelArg(_patchKernel, 11, sizeof(int), &empty.offset);
clSetKernelArg(_patchKernel, 12, sizeof(int), &empty.stride);
clSetKernelArg(_patchKernel, 13, sizeof(cl_mem), NULL);
clSetKernelArg(_patchKernel, 14, sizeof(int), &empty.offset);
clSetKernelArg(_patchKernel, 15, sizeof(int), &empty.stride);
clSetKernelArg(_patchKernel, 16, sizeof(cl_mem), NULL);
clSetKernelArg(_patchKernel, 17, sizeof(int), &empty.offset);
clSetKernelArg(_patchKernel, 18, sizeof(int), &empty.stride);
clSetKernelArg(_patchKernel, 19, sizeof(cl_mem), &patchCoordsBuffer);
clSetKernelArg(_patchKernel, 20, sizeof(cl_mem), &patchArrayBuffer);
clSetKernelArg(_patchKernel, 21, sizeof(cl_mem), &patchIndexBuffer);
clSetKernelArg(_patchKernel, 22, sizeof(cl_mem), &patchParamBuffer);
cl_int errNum = clEnqueueNDRangeKernel(
_clCommandQueue, _patchKernel, 1, NULL,
@ -340,13 +453,70 @@ CLEvaluator::EvalPatches(cl_mem src, BufferDescriptor const &srcDesc,
return false;
}
if (endEvent == NULL)
{
if (endEvent == NULL) {
clFinish(_clCommandQueue);
}
return true;
}
bool
CLEvaluator::EvalPatches(cl_mem src, BufferDescriptor const &srcDesc,
cl_mem dst, BufferDescriptor const &dstDesc,
cl_mem du, BufferDescriptor const &duDesc,
cl_mem dv, BufferDescriptor const &dvDesc,
cl_mem duu, BufferDescriptor const &duuDesc,
cl_mem duv, BufferDescriptor const &duvDesc,
cl_mem dvv, BufferDescriptor const &dvvDesc,
int numPatchCoords,
cl_mem patchCoordsBuffer,
cl_mem patchArrayBuffer,
cl_mem patchIndexBuffer,
cl_mem patchParamBuffer,
unsigned int numStartEvents,
const cl_event* startEvents,
cl_event* endEvent) const {
size_t globalWorkSize = (size_t)(numPatchCoords);
clSetKernelArg(_patchKernel, 0, sizeof(cl_mem), &src);
clSetKernelArg(_patchKernel, 1, sizeof(int), &srcDesc.offset);
clSetKernelArg(_patchKernel, 2, sizeof(cl_mem), &dst);
clSetKernelArg(_patchKernel, 3, sizeof(int), &dstDesc.offset);
clSetKernelArg(_patchKernel, 4, sizeof(cl_mem), &du);
clSetKernelArg(_patchKernel, 5, sizeof(int), &duDesc.offset);
clSetKernelArg(_patchKernel, 6, sizeof(int), &duDesc.stride);
clSetKernelArg(_patchKernel, 7, sizeof(cl_mem), &dv);
clSetKernelArg(_patchKernel, 8, sizeof(int), &dvDesc.offset);
clSetKernelArg(_patchKernel, 9, sizeof(int), &dvDesc.stride);
clSetKernelArg(_patchKernel, 10, sizeof(cl_mem), &duu);
clSetKernelArg(_patchKernel, 11, sizeof(int), &duuDesc.offset);
clSetKernelArg(_patchKernel, 12, sizeof(int), &duuDesc.stride);
clSetKernelArg(_patchKernel, 13, sizeof(cl_mem), &duv);
clSetKernelArg(_patchKernel, 14, sizeof(int), &duvDesc.offset);
clSetKernelArg(_patchKernel, 15, sizeof(int), &duvDesc.stride);
clSetKernelArg(_patchKernel, 16, sizeof(cl_mem), &dvv);
clSetKernelArg(_patchKernel, 17, sizeof(int), &dvvDesc.offset);
clSetKernelArg(_patchKernel, 18, sizeof(int), &dvvDesc.stride);
clSetKernelArg(_patchKernel, 19, sizeof(cl_mem), &patchCoordsBuffer);
clSetKernelArg(_patchKernel, 20, sizeof(cl_mem), &patchArrayBuffer);
clSetKernelArg(_patchKernel, 21, sizeof(cl_mem), &patchIndexBuffer);
clSetKernelArg(_patchKernel, 22, sizeof(cl_mem), &patchParamBuffer);
cl_int errNum = clEnqueueNDRangeKernel(
_clCommandQueue, _patchKernel, 1, NULL,
&globalWorkSize, NULL, numStartEvents, startEvents, endEvent);
if (errNum != CL_SUCCESS) {
Far::Error(Far::FAR_RUNTIME_ERROR,
"ApplyPatchKernel (%d) ", errNum);
return false;
}
if (endEvent == NULL) {
clFinish(_clCommandQueue);
}
return true;
}
/* static */

1652
opensubdiv/osd/clEvaluator.h
View File

File diff suppressed because it is too large Load Diff

54
opensubdiv/osd/clKernel.cl
View File

@ -99,12 +99,18 @@ __kernel void computeStencilsDerivatives(
__global float * dst, int dstOffset,
__global float * du, int duOffset, int duStride,
__global float * dv, int dvOffset, int dvStride,
__global float * duu, int duuOffset, int duuStride,
__global float * duv, int duvOffset, int duvStride,
__global float * dvv, int dvvOffset, int dvvStride,
__global int * sizes,
__global int * offsets,
__global int * indices,
__global float * weights,
__global float * duWeights,
__global float * dvWeights,
__global float * duuWeights,
__global float * duvWeights,
__global float * dvvWeights,
int batchStart, int batchEnd) {
int current = get_global_id(0) + batchStart;
@ -113,10 +119,13 @@ __kernel void computeStencilsDerivatives(
return;
}
struct Vertex v, vdu, vdv;
struct Vertex v, vdu, vdv, vduu, vduv, vdvv;
clear(&v);
clear(&vdu);
clear(&vdv);
clear(&vduu);
clear(&vduv);
clear(&vdvv);
int size = sizes[current],
offset = offsets[current];
@ -125,6 +134,9 @@ __kernel void computeStencilsDerivatives(
if (dst) dst += dstOffset;
if (du) du += duOffset;
if (dv) dv += dvOffset;
if (duu) duu += duuOffset;
if (duv) duv += duvOffset;
if (dvv) dvv += dvvOffset;
for (int i=0; i<size; ++i) {
int ofs = offset + i;
@ -132,11 +144,17 @@ __kernel void computeStencilsDerivatives(
if (weights) addWithWeight( &v, src, vid, weights[ofs]);
if (duWeights) addWithWeight(&vdu, src, vid, duWeights[ofs]);
if (dvWeights) addWithWeight(&vdv, src, vid, dvWeights[ofs]);
if (duuWeights) addWithWeight(&vduu, src, vid, duuWeights[ofs]);
if (duvWeights) addWithWeight(&vduv, src, vid, duvWeights[ofs]);
if (dvvWeights) addWithWeight(&vdvv, src, vid, dvvWeights[ofs]);
}
if (dst) writeVertex (dst, current, &v);
if (du) writeVertexStride(du, current, &vdu, duStride);
if (dv) writeVertexStride(dv, current, &vdv, dvStride);
if (duu) writeVertexStride(duu, current, &vduu, duuStride);
if (duv) writeVertexStride(duv, current, &vduv, duvStride);
if (dvv) writeVertexStride(dvv, current, &vdvv, dvvStride);
}
// ---------------------------------------------------------------------------
@ -205,6 +223,9 @@ __kernel void computePatches(__global float *src, int srcOffset,
__global float *dst, int dstOffset,
__global float *du, int duOffset, int duStride,
__global float *dv, int dvOffset, int dvStride,
__global float *duu, int duuOffset, int duuStride,
__global float *duv, int duvOffset, int duvStride,
__global float *dvv, int dvvOffset, int dvvStride,
__global struct PatchCoord *patchCoords,
__global struct PatchArray *patchArrayBuffer,
__global int *patchIndexBuffer,
@ -215,6 +236,9 @@ __kernel void computePatches(__global float *src, int srcOffset,
if (dst) dst += dstOffset;
if (du) du += duOffset;
if (dv) dv += dvOffset;
if (duu) duu += duuOffset;
if (duv) duv += duvOffset;
if (dvv) dvv += dvvOffset;
struct PatchCoord coord = patchCoords[current];
struct PatchArray array = patchArrayBuffer[coord.arrayIndex];
@ -274,5 +298,31 @@ __kernel void computePatches(__global float *src, int srcOffset,
}
writeVertexStride(dv, current, &vdv, dvStride);
}
if (duu) {
struct Vertex vduu;
clear(&vduu);
for (int i = 0; i < numControlVertices; ++i) {
int index = patchIndexBuffer[indexBase + i];
addWithWeight(&vduu, src, index, wDss[i]);
}
writeVertexStride(duu, current, &vduu, duuStride);
}
if (duv) {
struct Vertex vduv;
clear(&vduv);
for (int i = 0; i < numControlVertices; ++i) {
int index = patchIndexBuffer[indexBase + i];
addWithWeight(&vduv, src, index, wDst[i]);
}
writeVertexStride(duv, current, &vduv, duvStride);
}
if (dvv) {
struct Vertex vdvv;
clear(&vdvv);
for (int i = 0; i < numControlVertices; ++i) {
int index = patchIndexBuffer[indexBase + i];
addWithWeight(&vdvv, src, index, wDtt[i]);
}
writeVertexStride(dvv, current, &vdvv, dvvStride);
}
}

156
opensubdiv/osd/cpuEvaluator.cpp
View File

@ -82,6 +82,48 @@ CpuEvaluator::EvalStencils(const float *src, BufferDescriptor const &srcDesc,
return true;
}
/* static */
bool
CpuEvaluator::EvalStencils(const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
const int * sizes,
const int * offsets,
const int * indices,
const float * weights,
const float * duWeights,
const float * dvWeights,
const float * duuWeights,
const float * duvWeights,
const float * dvvWeights,
int start, int end) {
if (end <= start) return true;
if (srcDesc.length != dstDesc.length) return false;
if (srcDesc.length != duDesc.length) return false;
if (srcDesc.length != dvDesc.length) return false;
if (srcDesc.length != duuDesc.length) return false;
if (srcDesc.length != duvDesc.length) return false;
if (srcDesc.length != dvvDesc.length) return false;
CpuEvalStencils(src, srcDesc,
dst, dstDesc,
du, duDesc,
dv, dvDesc,
duu, duuDesc,
duv, duvDesc,
dvv, dvvDesc,
sizes, offsets, indices,
weights, duWeights, dvWeights,
duuWeights, duvWeights, dvvWeights,
start, end);
return true;
}
template <typename T>
struct BufferAdapter {
BufferAdapter(T *p, int length, int stride) :
@ -264,6 +306,120 @@ CpuEvaluator::EvalPatches(const float *src, BufferDescriptor const &srcDesc,
return true;
}
/* static */
bool
CpuEvaluator::EvalPatches(const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
int numPatchCoords,
const PatchCoord *patchCoords,
const PatchArray *patchArrays,
const int *patchIndexBuffer,
const PatchParam *patchParamBuffer) {
if (src) {
src += srcDesc.offset;
} else {
return false;
}
if (dst) {
if (srcDesc.length != dstDesc.length) return false;
dst += dstDesc.offset;
}
if (du) {
du += duDesc.offset;
if (srcDesc.length != duDesc.length) return false;
}
if (dv) {
dv += dvDesc.offset;
if (srcDesc.length != dvDesc.length) return false;
}
if (duu) {
duu += duuDesc.offset;
if (srcDesc.length != duuDesc.length) return false;
}
if (duv) {
duv += duvDesc.offset;
if (srcDesc.length != duvDesc.length) return false;
}
if (dvv) {
dvv += dvvDesc.offset;
if (srcDesc.length != dvvDesc.length) return false;
}
BufferAdapter<const float> srcT(src, srcDesc.length, srcDesc.stride);
BufferAdapter<float> dstT(dst, dstDesc.length, dstDesc.stride);
BufferAdapter<float> duT(du, duDesc.length, duDesc.stride);
BufferAdapter<float> dvT(dv, dvDesc.length, dvDesc.stride);
BufferAdapter<float> duuT(duu, duuDesc.length, duuDesc.stride);
BufferAdapter<float> duvT(duv, duvDesc.length, duvDesc.stride);
BufferAdapter<float> dvvT(dvv, dvvDesc.length, dvvDesc.stride);
float wP[20], wDu[20], wDv[20], wDuu[20], wDuv[20], wDvv[20];
for (int i = 0; i < numPatchCoords; ++i) {
PatchCoord const &coord = patchCoords[i];
PatchArray const &array = patchArrays[coord.handle.arrayIndex];
Far::PatchParam const & param =
patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
? Far::PatchDescriptor::REGULAR
: array.GetPatchType();
int numControlVertices = 0;
if (patchType == Far::PatchDescriptor::REGULAR) {
Far::internal::GetBSplineWeights(param,
coord.s, coord.t, wP, wDu, wDv,
wDuu, wDuv, wDvv);
numControlVertices = 16;
} else if (patchType == Far::PatchDescriptor::GREGORY_BASIS) {
Far::internal::GetGregoryWeights(param,
coord.s, coord.t, wP, wDu, wDv,
wDuu, wDuv, wDvv);
numControlVertices = 20;
} else if (patchType == Far::PatchDescriptor::QUADS) {
Far::internal::GetBilinearWeights(param,
coord.s, coord.t, wP, wDu, wDv,
wDuu, wDuv, wDvv);
numControlVertices = 4;
} else {
assert(0);
}
int indexStride = Far::PatchDescriptor(array.GetPatchType()).GetNumControlVertices();
int indexBase = array.GetIndexBase() + indexStride *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
const int *cvs = &patchIndexBuffer[indexBase];
dstT.Clear();
duT.Clear();
dvT.Clear();
duuT.Clear();
duvT.Clear();
dvvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
duT.AddWithWeight (srcT[cvs[j]], wDu[j]);
dvT.AddWithWeight (srcT[cvs[j]], wDv[j]);
duuT.AddWithWeight (srcT[cvs[j]], wDuu[j]);
duvT.AddWithWeight (srcT[cvs[j]], wDuv[j]);
dvvT.AddWithWeight (srcT[cvs[j]], wDvv[j]);
}
++dstT;
++duT;
++dvT;
++duuT;
++duvT;
++dvvT;
}
return true;
}
} // end namespace Osd

689
opensubdiv/osd/cpuEvaluator.h
View File

@ -26,11 +26,11 @@
#define OPENSUBDIV3_OSD_CPU_EVALUATOR_H
#include "../version.h"
#include <cstddef>
#include "../osd/bufferDescriptor.h"
#include "../osd/types.h"
#include <cstddef>
namespace OpenSubdiv {
namespace OPENSUBDIV_VERSION {
@ -107,7 +107,6 @@ public:
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
/// to apply for the range [start, end)
///
/// @param offsets pointer to the offsets buffer of the stencil table
///
@ -145,17 +144,17 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output U-derivative buffer
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the output buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output V-derivative buffer
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the output buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param stencilTable Far::StencilTable or equivalent
///
@ -206,15 +205,15 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output U-derivatives pointer. An offset of
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the output buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output V-derivatives pointer. An offset of
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the output buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
///
@ -245,6 +244,177 @@ public:
const float * dvWeights,
int start, int end);
/// \brief Generic static eval stencils function with derivatives.
/// This function has a same signature as other device kernels
/// have so that it can be called in the same way from OsdMesh
/// template interface.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param stencilTable Far::StencilTable or equivalent
///
/// @param instance not used in the cpu kernel
/// (declared as a typed pointer to prevent
/// undesirable template resolution)
///
/// @param deviceContext not used in the cpu kernel
///
template <typename SRC_BUFFER, typename DST_BUFFER, typename STENCIL_TABLE>
static bool EvalStencils(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
STENCIL_TABLE const *stencilTable,
const CpuEvaluator *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalStencils(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
&stencilTable->GetSizes()[0],
&stencilTable->GetOffsets()[0],
&stencilTable->GetControlIndices()[0],
&stencilTable->GetWeights()[0],
&stencilTable->GetDuWeights()[0],
&stencilTable->GetDvWeights()[0],
&stencilTable->GetDuuWeights()[0],
&stencilTable->GetDuvWeights()[0],
&stencilTable->GetDvvWeights()[0],
/*start = */ 0,
/*end = */ stencilTable->GetNumStencils());
}
/// \brief Static eval stencils function with derivatives, which takes
/// raw CPU pointers for input and output.
///
/// @param src Input primvar pointer. An offset of srcDesc
/// will be applied internally (i.e. the pointer
/// should not include the offset)
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dst Output primvar pointer. An offset of dstDesc
/// will be applied internally.
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duu Output pointer 2nd derivative wrt u. An offset of
/// duuDesc will be applied internally.
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duv Output pointer 2nd derivative wrt u and v. An offset of
/// duvDesc will be applied internally.
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvv Output pointer 2nd derivative wrt v. An offset of
/// dvvDesc will be applied internally.
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
///
/// @param offsets pointer to the offsets buffer of the stencil table
///
/// @param indices pointer to the indices buffer of the stencil table
///
/// @param weights pointer to the weights buffer of the stencil table
///
/// @param duWeights pointer to the du-weights buffer of the stencil table
///
/// @param dvWeights pointer to the dv-weights buffer of the stencil table
///
/// @param duuWeights pointer to the duu-weights buffer of the stencil table
///
/// @param duvWeights pointer to the duv-weights buffer of the stencil table
///
/// @param dvvWeights pointer to the dvv-weights buffer of the stencil table
///
/// @param start start index of stencil table
///
/// @param end end index of stencil table
///
static bool EvalStencils(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
const int * sizes,
const int * offsets,
const int * indices,
const float * weights,
const float * duWeights,
const float * dvWeights,
const float * duuWeights,
const float * duvWeights,
const float * dvvWeights,
int start, int end);
/// ----------------------------------------------------------------------
///
/// Limit evaluations with PatchTable
@ -318,13 +488,13 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output U-derivatives buffer
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output V-derivatives buffer
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
@ -354,6 +524,7 @@ public:
PATCH_TABLE *patchTable,
CpuEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
@ -373,6 +544,102 @@ public:
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function with derivatives. This function has
/// a same signature as other device kernels have so that it can be
/// called in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the cpu evaluator
///
/// @param deviceContext not used in the cpu evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatches(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
CpuEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
// XXX: PatchCoords is somewhat abusing vertex primvar buffer interop.
// ideally all buffer classes should have templated by datatype
// so that downcast isn't needed there.
// (e.g. Osd::CpuBuffer<PatchCoord> )
//
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetPatchArrayBuffer(),
patchTable->GetPatchIndexBuffer(),
patchTable->GetPatchParamBuffer());
}
/// \brief Static limit eval function. It takes an array of PatchCoord
/// and evaluate limit values on given PatchTable.
///
@ -423,15 +690,15 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output U-derivatives pointer. An offset of
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the du buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output V-derivatives pointer. An offset of
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dv buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
@ -457,6 +724,72 @@ public:
const int *patchIndexBuffer,
PatchParam const *patchParamBuffer);
/// \brief Static limit eval function. It takes an array of PatchCoord
/// and evaluate limit values on given PatchTable.
///
/// @param src Input primvar pointer. An offset of srcDesc
/// will be applied internally (i.e. the pointer
/// should not include the offset)
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dst Output primvar pointer. An offset of dstDesc
/// will be applied internally.
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duu Output pointer 2nd derivative wrt u. An offset of
/// duuDesc will be applied internally.
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duv Output pointer 2nd derivative wrt u and v. An offset of
/// duvDesc will be applied internally.
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvv Output pointer 2nd derivative wrt v. An offset of
/// dvvDesc will be applied internally.
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchArrays an array of Osd::PatchArray struct
/// indexed by PatchCoord::arrayIndex
///
/// @param patchIndexBuffer an array of patch indices
/// indexed by PatchCoord::vertIndex
///
/// @param patchParamBuffer an array of Osd::PatchParam struct
/// indexed by PatchCoord::patchIndex
///
static bool EvalPatches(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PatchCoord const *patchCoords,
PatchArray const *patchArrays,
const int *patchIndexBuffer,
PatchParam const *patchParamBuffer);
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
@ -508,6 +841,164 @@ public:
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the cpu evaluator
///
/// @param deviceContext not used in the cpu evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
CpuEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetVaryingPatchArrayBuffer(),
patchTable->GetVaryingPatchIndexBuffer(),
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the cpu evaluator
///
/// @param deviceContext not used in the cpu evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
CpuEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetVaryingPatchArrayBuffer(),
patchTable->GetVaryingPatchIndexBuffer(),
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
@ -562,6 +1053,170 @@ public:
patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param fvarChannel face-varying channel
///
/// @param instance not used in the cpu evaluator
///
/// @param deviceContext not used in the cpu evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesFaceVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
int fvarChannel,
CpuEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetFVarPatchArrayBuffer(fvarChannel),
patchTable->GetFVarPatchIndexBuffer(fvarChannel),
patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param fvarChannel face-varying channel
///
/// @param instance not used in the cpu evaluator
///
/// @param deviceContext not used in the cpu evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesFaceVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
int fvarChannel,
CpuEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetFVarPatchArrayBuffer(fvarChannel),
patchTable->GetFVarPatchIndexBuffer(fvarChannel),
patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// ----------------------------------------------------------------------
///
/// Other methods

70
opensubdiv/osd/cpuKernel.cpp
View File

@ -169,6 +169,76 @@ CpuEvalStencils(float const * src, BufferDescriptor const &srcDesc,
}
}
void
CpuEvalStencils(float const * src, BufferDescriptor const &srcDesc,
float * dst, BufferDescriptor const &dstDesc,
float * dstDu, BufferDescriptor const &dstDuDesc,
float * dstDv, BufferDescriptor const &dstDvDesc,
float * dstDuu, BufferDescriptor const &dstDuuDesc,
float * dstDuv, BufferDescriptor const &dstDuvDesc,
float * dstDvv, BufferDescriptor const &dstDvvDesc,
int const * sizes,
int const * offsets,
int const * indices,
float const * weights,
float const * duWeights,
float const * dvWeights,
float const * duuWeights,
float const * duvWeights,
float const * dvvWeights,
int start, int end) {
if (start > 0) {
sizes += start;
indices += offsets[start];
weights += offsets[start];
duWeights += offsets[start];
dvWeights += offsets[start];
duuWeights += offsets[start];
duvWeights += offsets[start];
dvvWeights += offsets[start];
}
src += srcDesc.offset;
dst += dstDesc.offset;
dstDu += dstDuDesc.offset;
dstDv += dstDvDesc.offset;
dstDuu += dstDuuDesc.offset;
dstDuv += dstDuvDesc.offset;
dstDvv += dstDvvDesc.offset;
int nOutLength = dstDesc.length + dstDuDesc.length + dstDvDesc.length
+ dstDuuDesc.length + dstDuvDesc.length + dstDvvDesc.length;
float * result = (float*)alloca(nOutLength * sizeof(float));
float * resultDu = result + dstDesc.length;
float * resultDv = resultDu + dstDuDesc.length;
float * resultDuu = resultDv + dstDuuDesc.length;
float * resultDuv = resultDuu + dstDuvDesc.length;
float * resultDvv = resultDuv + dstDvvDesc.length;
int nStencils = end - start;
for (int i = 0; i < nStencils; ++i, ++sizes) {
// clear
memset(result, 0, nOutLength * sizeof(float));
for (int j=0; j<*sizes; ++j) {
addWithWeight(result, src, *indices, *weights++, srcDesc);
addWithWeight(resultDu, src, *indices, *duWeights++, srcDesc);
addWithWeight(resultDv, src, *indices, *dvWeights++, srcDesc);
addWithWeight(resultDuu, src, *indices, *duuWeights++, srcDesc);
addWithWeight(resultDuv, src, *indices, *duvWeights++, srcDesc);
addWithWeight(resultDvv, src, *indices, *dvvWeights++, srcDesc);
++indices;
}
copy(dst, i, result, dstDesc);
copy(dstDu, i, resultDu, dstDuDesc);
copy(dstDv, i, resultDv, dstDvDesc);
copy(dstDuu, i, resultDuu, dstDuuDesc);
copy(dstDuv, i, resultDuu, dstDuvDesc);
copy(dstDvv, i, resultDvv, dstDvvDesc);
}
}
} // end namespace Osd
} // end namespace OPENSUBDIV_VERSION

19
opensubdiv/osd/cpuKernel.h
View File

@ -57,6 +57,25 @@ CpuEvalStencils(float const * src, BufferDescriptor const &srcDesc,
float const * dvWeights,
int start, int end);
void
CpuEvalStencils(float const * src, BufferDescriptor const &srcDesc,
float * dst, BufferDescriptor const &dstDesc,
float * dstDu, BufferDescriptor const &dstDuDesc,
float * dstDv, BufferDescriptor const &dstDvDesc,
float * dstDuu, BufferDescriptor const &dstDuuDesc,
float * dstDuv, BufferDescriptor const &dstDuvDesc,
float * dstDvv, BufferDescriptor const &dstDvvDesc,
int const * sizes,
int const * offsets,
int const * indices,
float const * weights,
float const * duWeights,
float const * dvWeights,
float const * duuWeights,
float const * duvWeights,
float const * dvvWeights,
int start, int end);
//
// SIMD ICC optimization of the stencil kernel
//

139
opensubdiv/osd/cudaEvaluator.cpp
View File

@ -53,9 +53,12 @@ extern "C" {
const void *patchParams);
void CudaEvalPatchesWithDerivatives(
const float *src, float *dst, float *du, float *dv,
int length,
int srcStride, int dstStride, int dvStride, int duStride,
const float *src, float *dst,
float *du, float *dv,
float *duu, float *duv, float *dvv,
int length, int srcStride, int dstStride,
int duStride, int dvStride,
int duuStride, int duvStride, int dvvStride,
int numPatchCoords,
const void *patchCoords,
const void *patchArrays,
@ -71,6 +74,10 @@ namespace Osd {
template <class T> void *
createCudaBuffer(std::vector<T> const & src) {
if (src.empty()) {
return NULL;
}
void * devicePtr = 0;
size_t size = src.size()*sizeof(T);
@ -98,9 +105,11 @@ CudaStencilTable::CudaStencilTable(Far::StencilTable const *stencilTable) {
_indices = createCudaBuffer(stencilTable->GetControlIndices());
_weights = createCudaBuffer(stencilTable->GetWeights());
_duWeights = _dvWeights = NULL;
_duuWeights = _duvWeights = _dvvWeights = NULL;
} else {
_sizes = _offsets = _indices = _weights = NULL;
_duWeights = _dvWeights = NULL;
_duuWeights = _duvWeights = _dvvWeights = NULL;
}
}
@ -113,9 +122,13 @@ CudaStencilTable::CudaStencilTable(Far::LimitStencilTable const *limitStencilTab
_weights = createCudaBuffer(limitStencilTable->GetWeights());
_duWeights = createCudaBuffer(limitStencilTable->GetDuWeights());
_dvWeights = createCudaBuffer(limitStencilTable->GetDvWeights());
_duuWeights = createCudaBuffer(limitStencilTable->GetDuuWeights());
_duvWeights = createCudaBuffer(limitStencilTable->GetDuvWeights());
_dvvWeights = createCudaBuffer(limitStencilTable->GetDvvWeights());
} else {
_sizes = _offsets = _indices = _weights = NULL;
_duWeights = _dvWeights = NULL;
_duuWeights = _duvWeights = _dvvWeights = NULL;
}
}
@ -126,6 +139,9 @@ CudaStencilTable::~CudaStencilTable() {
if (_weights) cudaFree(_weights);
if (_duWeights) cudaFree(_duWeights);
if (_dvWeights) cudaFree(_dvWeights);
if (_duuWeights) cudaFree(_duuWeights);
if (_duvWeights) cudaFree(_duvWeights);
if (_dvvWeights) cudaFree(_dvvWeights);
}
// ---------------------------------------------------------------------------
@ -197,6 +213,84 @@ CudaEvaluator::EvalStencils(const float *src, BufferDescriptor const &srcDesc,
return true;
}
/* static */
bool
CudaEvaluator::EvalStencils(const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
const int * sizes,
const int * offsets,
const int * indices,
const float * weights,
const float * duWeights,
const float * dvWeights,
const float * duuWeights,
const float * duvWeights,
const float * dvvWeights,
int start,
int end) {
// PERFORMANCE: need to combine 3 launches together
if (dst) {
CudaEvalStencils(src + srcDesc.offset,
dst + dstDesc.offset,
srcDesc.length,
srcDesc.stride,
dstDesc.stride,
sizes, offsets, indices, weights,
start, end);
}
if (du) {
CudaEvalStencils(src + srcDesc.offset,
du + duDesc.offset,
srcDesc.length,
srcDesc.stride,
duDesc.stride,
sizes, offsets, indices, duWeights,
start, end);
}
if (dv) {
CudaEvalStencils(src + srcDesc.offset,
dv + dvDesc.offset,
srcDesc.length,
srcDesc.stride,
dvDesc.stride,
sizes, offsets, indices, dvWeights,
start, end);
}
if (duu) {
CudaEvalStencils(src + srcDesc.offset,
duu + duuDesc.offset,
srcDesc.length,
srcDesc.stride,
duuDesc.stride,
sizes, offsets, indices, duuWeights,
start, end);
}
if (duv) {
CudaEvalStencils(src + srcDesc.offset,
duv + duvDesc.offset,
srcDesc.length,
srcDesc.stride,
duvDesc.stride,
sizes, offsets, indices, duvWeights,
start, end);
}
if (dvv) {
CudaEvalStencils(src + srcDesc.offset,
dvv + dvvDesc.offset,
srcDesc.length,
srcDesc.stride,
dvvDesc.stride,
sizes, offsets, indices, dvvWeights,
start, end);
}
return true;
}
/* static */
bool
CudaEvaluator::EvalPatches(const float *src,
@ -237,9 +331,42 @@ CudaEvaluator::EvalPatches(
if (dv) dv += dvDesc.offset;
CudaEvalPatchesWithDerivatives(
src, dst, du, dv,
srcDesc.length, srcDesc.stride,
dstDesc.stride, duDesc.stride, dvDesc.stride,
src, dst, du, dv, NULL, NULL, NULL,
srcDesc.length, srcDesc.stride, dstDesc.stride,
duDesc.stride, dvDesc.stride, 0, 0, 0,
numPatchCoords, patchCoords, patchArrays, patchIndices, patchParams);
return true;
}
/* static */
bool
CudaEvaluator::EvalPatches(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
int numPatchCoords,
const PatchCoord *patchCoords,
const PatchArray *patchArrays,
const int *patchIndices,
const PatchParam *patchParams) {
if (src) src += srcDesc.offset;
if (dst) dst += dstDesc.offset;
if (du) du += duDesc.offset;
if (dv) dv += dvDesc.offset;
if (duu) duu += duuDesc.offset;
if (duv) duv += duvDesc.offset;
if (dvv) dvv += dvvDesc.offset;
CudaEvalPatchesWithDerivatives(
src, dst, du, dv, duu, duv, dvv,
srcDesc.length, srcDesc.stride, dstDesc.stride,
duDesc.stride, dvDesc.stride,
duuDesc.stride, duvDesc.stride, dvvDesc.stride,
numPatchCoords, patchCoords, patchArrays, patchIndices, patchParams);
return true;
}

690
opensubdiv/osd/cudaEvaluator.h
View File

@ -73,6 +73,9 @@ public:
void *GetWeightsBuffer() const { return _weights; }
void *GetDuWeightsBuffer() const { return _duWeights; }
void *GetDvWeightsBuffer() const { return _dvWeights; }
void *GetDuuWeightsBuffer() const { return _duuWeights; }
void *GetDuvWeightsBuffer() const { return _duvWeights; }
void *GetDvvWeightsBuffer() const { return _dvvWeights; }
int GetNumStencils() const { return _numStencils; }
private:
@ -81,7 +84,10 @@ private:
* _indices,
* _weights,
* _duWeights,
* _dvWeights;
* _dvWeights,
* _duuWeights,
* _duvWeights,
* _dvvWeights;
int _numStencils;
};
@ -188,17 +194,17 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output U-derivative buffer
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the output buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output V-derivative buffer
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the output buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param stencilTable stencil table to be applied.
///
@ -249,15 +255,15 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output U-derivatives pointer. An offset of
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the output buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output V-derivatives pointer. An offset of
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the output buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
///
@ -288,6 +294,177 @@ public:
const float * dvWeights,
int start, int end);
/// \brief Generic static eval stencils function with derivatives.
/// This function has a same signature as other device kernels
/// have so that it can be called in the same way from OsdMesh
/// template interface.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCudaBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param stencilTable stencil table to be applied.
///
/// @param instance not used in the cuda kernel
/// (declared as a typed pointer to prevent
/// undesirable template resolution)
///
/// @param deviceContext not used in the cuda kernel
///
template <typename SRC_BUFFER, typename DST_BUFFER, typename STENCIL_TABLE>
static bool EvalStencils(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
STENCIL_TABLE const *stencilTable,
const CudaEvaluator *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalStencils(srcBuffer->BindCudaBuffer(), srcDesc,
dstBuffer->BindCudaBuffer(), dstDesc,
duBuffer->BindCudaBuffer(), duDesc,
dvBuffer->BindCudaBuffer(), dvDesc,
duuBuffer->BindCudaBuffer(), duuDesc,
duvBuffer->BindCudaBuffer(), duvDesc,
dvvBuffer->BindCudaBuffer(), dvvDesc,
(int const *)stencilTable->GetSizesBuffer(),
(int const *)stencilTable->GetOffsetsBuffer(),
(int const *)stencilTable->GetIndicesBuffer(),
(float const *)stencilTable->GetWeightsBuffer(),
(float const *)stencilTable->GetDuWeightsBuffer(),
(float const *)stencilTable->GetDvWeightsBuffer(),
(float const *)stencilTable->GetDuuWeightsBuffer(),
(float const *)stencilTable->GetDuvWeightsBuffer(),
(float const *)stencilTable->GetDvvWeightsBuffer(),
/*start = */ 0,
/*end = */ stencilTable->GetNumStencils());
}
/// \brief Static eval stencils function with derivatives, which takes
/// raw cuda pointers for input and output.
///
/// @param src Input primvar pointer. An offset of srcDesc
/// will be applied internally (i.e. the pointer
/// should not include the offset)
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dst Output primvar pointer. An offset of dstDesc
/// will be applied internally.
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duu Output pointer 2nd derivative wrt u. An offset of
/// duuDesc will be applied internally.
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duv Output pointer 2nd derivative wrt u and v. An offset of
/// duvDesc will be applied internally.
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvv Output pointer 2nd derivative wrt v. An offset of
/// dvvDesc will be applied internally.
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
///
/// @param offsets pointer to the offsets buffer of the stencil table
///
/// @param indices pointer to the indices buffer of the stencil table
///
/// @param weights pointer to the weights buffer of the stencil table
///
/// @param duWeights pointer to the du-weights buffer of the stencil table
///
/// @param dvWeights pointer to the dv-weights buffer of the stencil table
///
/// @param duuWeights pointer to the duu-weights buffer of the stencil table
///
/// @param duvWeights pointer to the duv-weights buffer of the stencil table
///
/// @param dvvWeights pointer to the dvv-weights buffer of the stencil table
///
/// @param start start index of stencil table
///
/// @param end end index of stencil table
///
static bool EvalStencils(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
const int * sizes,
const int * offsets,
const int * indices,
const float * weights,
const float * duWeights,
const float * dvWeights,
const float * duuWeights,
const float * duvWeights,
const float * dvvWeights,
int start, int end);
/// ----------------------------------------------------------------------
///
/// Limit evaluations with PatchTable
@ -361,13 +538,13 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output U-derivatives buffer
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output V-derivatives buffer
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
@ -410,6 +587,95 @@ public:
(const PatchParam *)patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function with derivatives. This function has
/// a same signature as other device kernels have so that it can be
/// called in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCudaBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CudaPatchTable or equivalent
///
/// @param instance not used in the cuda evaluator
///
/// @param deviceContext not used in the cuda evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatches(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
CudaEvaluator const *instance,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCudaBuffer(), srcDesc,
dstBuffer->BindCudaBuffer(), dstDesc,
duBuffer->BindCudaBuffer(), duDesc,
dvBuffer->BindCudaBuffer(), dvDesc,
duuBuffer->BindCudaBuffer(), duuDesc,
duvBuffer->BindCudaBuffer(), duvDesc,
dvvBuffer->BindCudaBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord *)patchCoords->BindCudaBuffer(),
(const PatchArray *)patchTable->GetPatchArrayBuffer(),
(const int *)patchTable->GetPatchIndexBuffer(),
(const PatchParam *)patchTable->GetPatchParamBuffer());
}
/// \brief Static limit eval function. It takes an array of PatchCoord
/// and evaluate limit values on given PatchTable.
///
@ -460,15 +726,15 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output U-derivatives pointer. An offset of
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the du buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output V-derivatives pointer. An offset of
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dv buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
@ -489,10 +755,76 @@ public:
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
int numPatchCoords,
const PatchCoord *patchCoords,
const PatchArray *patchArrays,
PatchCoord const *patchCoords,
PatchArray const *patchArrays,
const int *patchIndices,
const PatchParam *patchParams);
PatchParam const *patchParams);
/// \brief Static limit eval function. It takes an array of PatchCoord
/// and evaluate limit values on given PatchTable.
///
/// @param src Input primvar pointer. An offset of srcDesc
/// will be applied internally (i.e. the pointer
/// should not include the offset)
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dst Output primvar pointer. An offset of dstDesc
/// will be applied internally.
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duu Output pointer 2nd derivative wrt u. An offset of
/// duuDesc will be applied internally.
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duv Output pointer 2nd derivative wrt u and v. An offset of
/// duvDesc will be applied internally.
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvv Output pointer 2nd derivative wrt v. An offset of
/// dvvDesc will be applied internally.
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchArrays an array of Osd::PatchArray struct
/// indexed by PatchCoord::arrayIndex
///
/// @param patchIndices an array of patch indices
/// indexed by PatchCoord::vertIndex
///
/// @param patchParams an array of Osd::PatchParam struct
/// indexed by PatchCoord::patchIndex
///
static bool EvalPatches(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PatchCoord const *patchCoords,
PatchArray const *patchArrays,
const int *patchIndices,
PatchParam const *patchParams);
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
@ -545,6 +877,164 @@ public:
(const PatchParam *)patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCudaBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
/// must have BindCudaBuffer() method returning an
/// array of PatchCoord struct in cuda memory.
///
/// @param patchTable CudaPatchTable or equivalent
///
/// @param instance not used in the cuda evaluator
///
/// @param deviceContext not used in the cuda evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
CudaEvaluator const *instance,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCudaBuffer(), srcDesc,
dstBuffer->BindCudaBuffer(), dstDesc,
duBuffer->BindCudaBuffer(), duDesc,
dvBuffer->BindCudaBuffer(), dvDesc,
numPatchCoords,
(const PatchCoord *)patchCoords->BindCudaBuffer(),
(const PatchArray *)patchTable->GetVaryingPatchArrayBuffer(),
(const int *)patchTable->GetVaryingPatchIndexBuffer(),
(const PatchParam *)patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCudaBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
/// must have BindCudaBuffer() method returning an
/// array of PatchCoord struct in cuda memory.
///
/// @param patchTable CudaPatchTable or equivalent
///
/// @param instance not used in the cuda evaluator
///
/// @param deviceContext not used in the cuda evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
CudaEvaluator const *instance,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCudaBuffer(), srcDesc,
dstBuffer->BindCudaBuffer(), dstDesc,
duBuffer->BindCudaBuffer(), duDesc,
dvBuffer->BindCudaBuffer(), dvDesc,
duuBuffer->BindCudaBuffer(), duuDesc,
duvBuffer->BindCudaBuffer(), duvDesc,
dvvBuffer->BindCudaBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord *)patchCoords->BindCudaBuffer(),
(const PatchArray *)patchTable->GetVaryingPatchArrayBuffer(),
(const int *)patchTable->GetVaryingPatchIndexBuffer(),
(const PatchParam *)patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
@ -599,6 +1089,170 @@ public:
(const PatchParam *)patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCudaBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
/// must have BindCudaBuffer() method returning an
/// array of PatchCoord struct in cuda memory.
///
/// @param patchTable CudaPatchTable or equivalent
///
/// @param fvarChannel face-varying channel
///
/// @param instance not used in the cuda evaluator
///
/// @param deviceContext not used in the cuda evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesFaceVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
int fvarChannel,
CudaEvaluator const *instance,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCudaBuffer(), srcDesc,
dstBuffer->BindCudaBuffer(), dstDesc,
duBuffer->BindCudaBuffer(), duDesc,
dvBuffer->BindCudaBuffer(), dvDesc,
numPatchCoords,
(const PatchCoord *)patchCoords->BindCudaBuffer(),
(const PatchArray *)patchTable->GetFVarPatchArrayBuffer(fvarChannel),
(const int *)patchTable->GetFVarPatchIndexBuffer(fvarChannel),
(const PatchParam *)patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCudaBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCudaBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
/// must have BindCudaBuffer() method returning an
/// array of PatchCoord struct in cuda memory.
///
/// @param patchTable CudaPatchTable or equivalent
///
/// @param fvarChannel face-varying channel
///
/// @param instance not used in the cuda evaluator
///
/// @param deviceContext not used in the cuda evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesFaceVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
int fvarChannel,
CudaEvaluator const *instance,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCudaBuffer(), srcDesc,
dstBuffer->BindCudaBuffer(), dstDesc,
duBuffer->BindCudaBuffer(), duDesc,
dvBuffer->BindCudaBuffer(), dvDesc,
duuBuffer->BindCudaBuffer(), duuDesc,
duvBuffer->BindCudaBuffer(), duvDesc,
dvvBuffer->BindCudaBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord *)patchCoords->BindCudaBuffer(),
(const PatchArray *)patchTable->GetFVarPatchArrayBuffer(fvarChannel),
(const int *)patchTable->GetFVarPatchIndexBuffer(fvarChannel),
(const PatchParam *)patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// ----------------------------------------------------------------------
///
/// Other methods

47
opensubdiv/osd/cudaKernel.cu
View File

@ -305,8 +305,12 @@ int getNumControlVertices(int patchType) {
}
__global__ void
computePatches(const float *src, float *dst, float *dstDu, float *dstDv,
int length, int srcStride, int dstStride, int dstDuStride, int dstDvStride,
computePatches(const float *src, float *dst,
float *dstDu, float *dstDv,
float *dstDuu, float *dstDuv, float *dstDvv,
int length, int srcStride, int dstStride,
int dstDuStride, int dstDvStride,
int dstDuuStride, int dstDuvStride, int dstDvvStride,
int numPatchCoords, const PatchCoord *patchCoords,
const PatchArray *patchArrayBuffer,
const int *patchIndexBuffer,
@ -376,6 +380,30 @@ computePatches(const float *src, float *dst, float *dstDu, float *dstDv,
addWithWeight(d, srcVert, wDt[j], length);
}
}
if (dstDuu) {
float *d = dstDuu + i * dstDuuStride;
clear(d, length);
for (int j = 0; j < numControlVertices; ++j) {
const float * srcVert = src + cvs[j] * srcStride;
addWithWeight(d, srcVert, wDss[j], length);
}
}
if (dstDuv) {
float *d = dstDuv + i * dstDuvStride;
clear(d, length);
for (int j = 0; j < numControlVertices; ++j) {
const float * srcVert = src + cvs[j] * srcStride;
addWithWeight(d, srcVert, wDst[j], length);
}
}
if (dstDvv) {
float *d = dstDvv + i * dstDvvStride;
clear(d, length);
for (int j = 0; j < numControlVertices; ++j) {
const float * srcVert = src + cvs[j] * srcStride;
addWithWeight(d, srcVert, wDtt[j], length);
}
}
}
}
@ -447,14 +475,19 @@ void CudaEvalPatches(
// PERFORMANCE: not optimized at all
computePatches <<<512, 32>>>(
src, dst, NULL, NULL, length, srcStride, dstStride, 0, 0,
src, dst, NULL, NULL, NULL, NULL, NULL,
length, srcStride, dstStride, 0, 0, 0, 0, 0,
numPatchCoords, patchCoords,
patchArrayBuffer, patchIndexBuffer, patchParamBuffer);
}
void CudaEvalPatchesWithDerivatives(
const float *src, float *dst, float *dstDu, float *dstDv,
int length, int srcStride, int dstStride, int dstDuStride, int dstDvStride,
const float *src, float *dst,
float *dstDu, float *dstDv,
float *dstDuu, float *dstDuv, float *dstDvv,
int length, int srcStride, int dstStride,
int dstDuStride, int dstDvStride,
int dstDuuStride, int dstDuvStride, int dstDvvStride,
int numPatchCoords, const PatchCoord *patchCoords,
const PatchArray *patchArrayBuffer,
const int *patchIndexBuffer,
@ -463,7 +496,9 @@ void CudaEvalPatchesWithDerivatives(
// PERFORMANCE: not optimized at all
computePatches <<<512, 32>>>(
src, dst, dstDu, dstDv, length, srcStride, dstStride, dstDuStride, dstDvStride,
src, dst, dstDu, dstDv, dstDuu, dstDuv, dstDvv,
length, srcStride, dstStride,
dstDuStride, dstDvStride, dstDuuStride, dstDuvStride, dstDvvStride,
numPatchCoords, patchCoords,
patchArrayBuffer, patchIndexBuffer, patchParamBuffer);
}

16
opensubdiv/osd/d3d11ComputeEvaluator.cpp
View File

@ -172,6 +172,22 @@ D3D11ComputeEvaluator::Create(BufferDescriptor const &srcDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc,
ID3D11DeviceContext *deviceContext) {
return Create(srcDesc, dstDesc, duDesc, dvDesc,
BufferDescriptor(),
BufferDescriptor(),
BufferDescriptor(),
deviceContext);
}
D3D11ComputeEvaluator *
D3D11ComputeEvaluator::Create(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc,
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc,
ID3D11DeviceContext *deviceContext) {
(void)deviceContext; // not used
// TODO: implements derivatives

9
opensubdiv/osd/d3d11ComputeEvaluator.h
View File

@ -102,6 +102,15 @@ public:
BufferDescriptor const &dvDesc,
ID3D11DeviceContext *deviceContext);
static D3D11ComputeEvaluator * Create(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc,
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc,
ID3D11DeviceContext *deviceContext);
/// Constructor.
D3D11ComputeEvaluator();

195
opensubdiv/osd/glComputeEvaluator.cpp
View File

@ -44,6 +44,10 @@ static const char *shaderSource =
template <class T> GLuint
createSSBO(std::vector<T> const & src) {
if (src.empty()) {
return 0;
}
GLuint devicePtr = 0;
glGenBuffers(1, &devicePtr);
@ -75,9 +79,11 @@ GLStencilTableSSBO::GLStencilTableSSBO(
_indices = createSSBO(stencilTable->GetControlIndices());
_weights = createSSBO(stencilTable->GetWeights());
_duWeights = _dvWeights = 0;
_duuWeights = _duvWeights = _dvvWeights = 0;
} else {
_sizes = _offsets = _indices = _weights = 0;
_duWeights = _dvWeights = 0;
_duuWeights = _duvWeights = _dvvWeights = 0;
}
}
@ -91,9 +97,13 @@ GLStencilTableSSBO::GLStencilTableSSBO(
_weights = createSSBO(limitStencilTable->GetWeights());
_duWeights = createSSBO(limitStencilTable->GetDuWeights());
_dvWeights = createSSBO(limitStencilTable->GetDvWeights());
_duuWeights = createSSBO(limitStencilTable->GetDuuWeights());
_duvWeights = createSSBO(limitStencilTable->GetDuvWeights());
_dvvWeights = createSSBO(limitStencilTable->GetDvvWeights());
} else {
_sizes = _offsets = _indices = _weights = 0;
_duWeights = _dvWeights = 0;
_duuWeights = _duvWeights = _dvvWeights = 0;
}
}
@ -104,6 +114,9 @@ GLStencilTableSSBO::~GLStencilTableSSBO() {
if (_weights) glDeleteBuffers(1, &_weights);
if (_duWeights) glDeleteBuffers(1, &_duWeights);
if (_dvWeights) glDeleteBuffers(1, &_dvWeights);
if (_duuWeights) glDeleteBuffers(1, &_duuWeights);
if (_duvWeights) glDeleteBuffers(1, &_duvWeights);
if (_dvvWeights) glDeleteBuffers(1, &_dvvWeights);
}
// ---------------------------------------------------------------------------
@ -120,8 +133,11 @@ GLComputeEvaluator::~GLComputeEvaluator() {
static GLuint
compileKernel(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const & /* duDesc */,
BufferDescriptor const & /* dvDesc */,
BufferDescriptor const & duDesc,
BufferDescriptor const & dvDesc,
BufferDescriptor const & duuDesc,
BufferDescriptor const & duvDesc,
BufferDescriptor const & dvvDesc,
const char *kernelDefine,
int workGroupSize) {
GLuint program = glCreateProgram();
@ -139,6 +155,16 @@ compileKernel(BufferDescriptor const &srcDesc,
<< "#define WORK_GROUP_SIZE " << workGroupSize << "\n"
<< kernelDefine << "\n"
<< patchBasisShaderSourceDefine << "\n";
bool deriv1 = (duDesc.length > 0 || dvDesc.length > 0);
bool deriv2 = (duuDesc.length > 0 || duvDesc.length > 0 || dvvDesc.length > 0);
if (deriv1) {
defines << "#define OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES\n";
}
if (deriv2) {
defines << "#define OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES\n";
}
std::string defineStr = defines.str();
const char *shaderSources[4] = {"#version 430\n", 0, 0, 0};
@ -175,16 +201,23 @@ bool
GLComputeEvaluator::Compile(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc) {
BufferDescriptor const &dvDesc,
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc) {
// create a stencil kernel
if (!_stencilKernel.Compile(srcDesc, dstDesc, duDesc, dvDesc,
if (!_stencilKernel.Compile(srcDesc, dstDesc,
duDesc, dvDesc,
duuDesc, duvDesc, dvvDesc,
_workGroupSize)) {
return false;
}
// create a patch kernel
if (!_patchKernel.Compile(srcDesc, dstDesc, duDesc, dvDesc,
if (!_patchKernel.Compile(srcDesc, dstDesc,
duDesc, dvDesc,
duuDesc, duvDesc, dvvDesc,
_workGroupSize)) {
return false;
}
@ -214,6 +247,40 @@ GLComputeEvaluator::EvalStencils(
GLuint dvWeightsBuffer,
int start, int end) const {
return EvalStencils(srcBuffer, srcDesc,
dstBuffer, dstDesc,
duBuffer, duDesc,
dvBuffer, dvDesc,
0, BufferDescriptor(),
0, BufferDescriptor(),
0, BufferDescriptor(),
sizesBuffer, offsetsBuffer, indicesBuffer,
weightsBuffer,
duWeightsBuffer, dvWeightsBuffer,
0, 0, 0,
start, end);
}
bool
GLComputeEvaluator::EvalStencils(
GLuint srcBuffer, BufferDescriptor const &srcDesc,
GLuint dstBuffer, BufferDescriptor const &dstDesc,
GLuint duBuffer, BufferDescriptor const &duDesc,
GLuint dvBuffer, BufferDescriptor const &dvDesc,
GLuint duuBuffer, BufferDescriptor const &duuDesc,
GLuint duvBuffer, BufferDescriptor const &duvDesc,
GLuint dvvBuffer, BufferDescriptor const &dvvDesc,
GLuint sizesBuffer,
GLuint offsetsBuffer,
GLuint indicesBuffer,
GLuint weightsBuffer,
GLuint duWeightsBuffer,
GLuint dvWeightsBuffer,
GLuint duuWeightsBuffer,
GLuint duvWeightsBuffer,
GLuint dvvWeightsBuffer,
int start, int end) const {
if (!_stencilKernel.program) return false;
int count = end - start;
if (count <= 0) {
@ -224,6 +291,9 @@ GLComputeEvaluator::EvalStencils(
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, dstBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, duBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, dvBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 10, duuBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 11, duvBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 12, dvvBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 4, sizesBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 5, offsetsBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 6, indicesBuffer);
@ -232,6 +302,12 @@ GLComputeEvaluator::EvalStencils(
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 8, duWeightsBuffer);
if (dvWeightsBuffer)
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 9, dvWeightsBuffer);
if (duuWeightsBuffer)
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 13, duuWeightsBuffer);
if (duvWeightsBuffer)
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 14, duvWeightsBuffer);
if (dvvWeightsBuffer)
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 15, dvvWeightsBuffer);
glUseProgram(_stencilKernel.program);
@ -247,13 +323,25 @@ GLComputeEvaluator::EvalStencils(
glUniform3i(_stencilKernel.uniformDvDesc,
dvDesc.offset, dvDesc.length, dvDesc.stride);
}
if (_stencilKernel.uniformDuuDesc > 0) {
glUniform3i(_stencilKernel.uniformDuuDesc,
duuDesc.offset, duuDesc.length, duuDesc.stride);
}
if (_stencilKernel.uniformDuvDesc > 0) {
glUniform3i(_stencilKernel.uniformDuvDesc,
duvDesc.offset, duvDesc.length, duvDesc.stride);
}
if (_stencilKernel.uniformDvvDesc > 0) {
glUniform3i(_stencilKernel.uniformDvvDesc,
dvvDesc.offset, dvvDesc.length, dvvDesc.stride);
}
glDispatchCompute((count + _workGroupSize - 1) / _workGroupSize, 1, 1);
glUseProgram(0);
glMemoryBarrier(GL_TEXTURE_FETCH_BARRIER_BIT);
for (int i = 0; i < 10; ++i) {
for (int i = 0; i < 16; ++i) {
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, i, 0);
}
@ -272,12 +360,44 @@ GLComputeEvaluator::EvalPatches(
GLuint patchIndexBuffer,
GLuint patchParamsBuffer) const {
return EvalPatches(srcBuffer, srcDesc,
dstBuffer, dstDesc,
duBuffer, duDesc,
dvBuffer, dvDesc,
0, BufferDescriptor(),
0, BufferDescriptor(),
0, BufferDescriptor(),
numPatchCoords,
patchCoordsBuffer,
patchArrays,
patchIndexBuffer,
patchParamsBuffer);
}
bool
GLComputeEvaluator::EvalPatches(
GLuint srcBuffer, BufferDescriptor const &srcDesc,
GLuint dstBuffer, BufferDescriptor const &dstDesc,
GLuint duBuffer, BufferDescriptor const &duDesc,
GLuint dvBuffer, BufferDescriptor const &dvDesc,
GLuint duuBuffer, BufferDescriptor const &duuDesc,
GLuint duvBuffer, BufferDescriptor const &duvDesc,
GLuint dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
GLuint patchCoordsBuffer,
const PatchArrayVector &patchArrays,
GLuint patchIndexBuffer,
GLuint patchParamsBuffer) const {
if (!_patchKernel.program) return false;
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, srcBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, dstBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, duBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, dvBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 10, duuBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 11, duvBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 12, dvvBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 4, patchCoordsBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 5, patchIndexBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 6, patchParamsBuffer);
@ -288,8 +408,27 @@ GLComputeEvaluator::EvalPatches(
glUniform1i(_patchKernel.uniformDstOffset, dstDesc.offset);
glUniform4iv(_patchKernel.uniformPatchArray, (int)patchArrays.size(),
(const GLint*)&patchArrays[0]);
glUniform3i(_patchKernel.uniformDuDesc, duDesc.offset, duDesc.length, duDesc.stride);
glUniform3i(_patchKernel.uniformDvDesc, dvDesc.offset, dvDesc.length, dvDesc.stride);
if (_patchKernel.uniformDuDesc > 0) {
glUniform3i(_patchKernel.uniformDuDesc,
duDesc.offset, duDesc.length, duDesc.stride);
}
if (_patchKernel.uniformDvDesc > 0) {
glUniform3i(_patchKernel.uniformDvDesc,
dvDesc.offset, dvDesc.length, dvDesc.stride);
}
if (_patchKernel.uniformDuuDesc > 0) {
glUniform3i(_patchKernel.uniformDuuDesc,
duuDesc.offset, duuDesc.length, duuDesc.stride);
}
if (_patchKernel.uniformDuvDesc > 0) {
glUniform3i(_patchKernel.uniformDuvDesc,
duvDesc.offset, duvDesc.length, duvDesc.stride);
}
if (_patchKernel.uniformDvvDesc > 0) {
glUniform3i(_patchKernel.uniformDvvDesc,
dvvDesc.offset, dvvDesc.length, dvvDesc.stride);
}
glDispatchCompute((numPatchCoords + _workGroupSize - 1) / _workGroupSize, 1, 1);
@ -303,6 +442,10 @@ GLComputeEvaluator::EvalPatches(
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 5, 0);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 6, 0);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 10, 0);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 11, 0);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 12, 0);
return true;
}
// ---------------------------------------------------------------------------
@ -320,20 +463,21 @@ GLComputeEvaluator::_StencilKernel::Compile(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc,
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc,
int workGroupSize) {
// create stencil kernel
if (program) {
glDeleteProgram(program);
}
bool derivatives = (duDesc.length > 0 || dvDesc.length > 0);
const char *kernelDef = derivatives
? "#define OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_STENCILS\n"
"#define OPENSUBDIV_GLSL_COMPUTE_USE_DERIVATIVES\n"
: "#define OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_STENCILS\n";
const char * kernelDefine =
"#define OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_STENCILS\n";
program = compileKernel(srcDesc, dstDesc, duDesc, dvDesc, kernelDef,
workGroupSize);
program = compileKernel(srcDesc, dstDesc,
duDesc, dvDesc, duuDesc, duvDesc, dvvDesc,
kernelDefine, workGroupSize);
if (program == 0) return false;
// cache uniform locations (TODO: use uniform block)
@ -343,6 +487,9 @@ GLComputeEvaluator::_StencilKernel::Compile(BufferDescriptor const &srcDesc,
uniformDstOffset = glGetUniformLocation(program, "dstOffset");
uniformDuDesc = glGetUniformLocation(program, "duDesc");
uniformDvDesc = glGetUniformLocation(program, "dvDesc");
uniformDuuDesc = glGetUniformLocation(program, "duuDesc");
uniformDuvDesc = glGetUniformLocation(program, "duvDesc");
uniformDvvDesc = glGetUniformLocation(program, "dvvDesc");
return true;
}
@ -362,20 +509,21 @@ GLComputeEvaluator::_PatchKernel::Compile(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc,
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc,
int workGroupSize) {
// create stencil kernel
if (program) {
glDeleteProgram(program);
}
bool derivatives = (duDesc.length > 0 || dvDesc.length > 0);
const char *kernelDef = derivatives
? "#define OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_PATCHES\n"
"#define OPENSUBDIV_GLSL_COMPUTE_USE_DERIVATIVES\n"
: "#define OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_PATCHES\n";
const char * kernelDefine =
"#define OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_PATCHES\n";
program = compileKernel(srcDesc, dstDesc, duDesc, dvDesc, kernelDef,
workGroupSize);
program = compileKernel(srcDesc, dstDesc,
duDesc, dvDesc, duuDesc, duvDesc, dvvDesc,
kernelDefine, workGroupSize);
if (program == 0) return false;
// cache uniform locations
@ -384,6 +532,9 @@ GLComputeEvaluator::_PatchKernel::Compile(BufferDescriptor const &srcDesc,
uniformPatchArray = glGetUniformLocation(program, "patchArray");
uniformDuDesc = glGetUniformLocation(program, "duDesc");
uniformDvDesc = glGetUniformLocation(program, "dvDesc");
uniformDuuDesc = glGetUniformLocation(program, "duuDesc");
uniformDuvDesc = glGetUniformLocation(program, "duvDesc");
uniformDvvDesc = glGetUniformLocation(program, "dvvDesc");
return true;
}

1366
opensubdiv/osd/glComputeEvaluator.h
View File

File diff suppressed because it is too large Load Diff

360
opensubdiv/osd/glXFBEvaluator.cpp
View File

@ -48,6 +48,10 @@ static const char *shaderSource =
template <class T> GLuint
createGLTextureBuffer(std::vector<T> const & src, GLenum type) {
if (src.empty()) {
return 0;
}
GLint size = static_cast<int>(src.size()*sizeof(T));
void const * ptr = &src.at(0);
@ -95,9 +99,11 @@ GLStencilTableTBO::GLStencilTableTBO(
stencilTable->GetControlIndices(), GL_R32I);
_weights = createGLTextureBuffer(stencilTable->GetWeights(), GL_R32F);
_duWeights = _dvWeights = 0;
_duuWeights = _duvWeights = _dvvWeights = 0;
} else {
_sizes = _offsets = _indices = _weights = 0;
_duWeights = _dvWeights = 0;
_duuWeights = _duvWeights = _dvvWeights = 0;
}
}
@ -118,9 +124,16 @@ GLStencilTableTBO::GLStencilTableTBO(
limitStencilTable->GetDuWeights(), GL_R32F);
_dvWeights = createGLTextureBuffer(
limitStencilTable->GetDvWeights(), GL_R32F);
_duuWeights = createGLTextureBuffer(
limitStencilTable->GetDuuWeights(), GL_R32F);
_duvWeights = createGLTextureBuffer(
limitStencilTable->GetDuvWeights(), GL_R32F);
_dvvWeights = createGLTextureBuffer(
limitStencilTable->GetDvvWeights(), GL_R32F);
} else {
_sizes = _offsets = _indices = _weights = 0;
_duWeights = _dvWeights = 0;
_duuWeights = _duvWeights = _dvvWeights = 0;
}
}
@ -131,11 +144,16 @@ GLStencilTableTBO::~GLStencilTableTBO() {
if (_weights) glDeleteTextures(1, &_weights);
if (_duWeights) glDeleteTextures(1, &_duWeights);
if (_dvWeights) glDeleteTextures(1, &_dvWeights);
if (_duuWeights) glDeleteTextures(1, &_duuWeights);
if (_duvWeights) glDeleteTextures(1, &_duvWeights);
if (_dvvWeights) glDeleteTextures(1, &_dvvWeights);
}
// ---------------------------------------------------------------------------
GLXFBEvaluator::GLXFBEvaluator() : _srcBufferTexture(0) {
GLXFBEvaluator::GLXFBEvaluator(bool sharedDerivativeBuffers)
: _srcBufferTexture(0),
_sharedDerivativeBuffers(sharedDerivativeBuffers) {
}
GLXFBEvaluator::~GLXFBEvaluator() {
@ -149,7 +167,11 @@ compileKernel(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc,
const char *kernelDefine) {
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc,
const char *kernelDefine,
bool sharedDerivativeBuffers) {
GLuint program = glCreateProgram();
@ -165,8 +187,25 @@ compileKernel(BufferDescriptor const &srcDesc,
<< "#define VERTEX_SHADER\n"
<< kernelDefine << "\n"
<< patchBasisShaderSourceDefine << "\n";
std::string defineStr = defines.str();
bool deriv1 = (duDesc.length > 0 || dvDesc.length > 0);
bool deriv2 = (duuDesc.length > 0 || duvDesc.length > 0 || dvvDesc.length > 0);
if (deriv1) {
defines << "#define OPENSUBDIV_GLSL_XFB_USE_1ST_DERIVATIVES\n";
if (sharedDerivativeBuffers) {
defines <<
"#define OPENSUBDIV_GLSL_XFB_SHARED_1ST_DERIVATIVE_BUFFERS\n";
}
}
if (deriv2) {
defines << "#define OPENSUBDIV_GLSL_XFB_USE_2ND_DERIVATIVES\n";
if (sharedDerivativeBuffers) {
defines <<
"#define OPENSUBDIV_GLSL_XFB_SHARED_2ND_DERIVATIVE_BUFFERS\n";
}
}
std::string defineStr = defines.str();
const char *shaderSources[4] = {"#version 410\n", NULL, NULL, NULL};
@ -204,7 +243,6 @@ compileKernel(BufferDescriptor const &srcDesc,
outputs.push_back("gl_SkipComponents1");
}
}
if (duDesc.length) {
//
// For derivatives, we use another buffer bindings so gl_NextBuffer
// is inserted here to switch the destination of transform feedback.
@ -213,6 +251,35 @@ compileKernel(BufferDescriptor const &srcDesc,
// vertex and each derivatives. gl_NextBuffer seems still works well
// in either case.
//
// If we know that the buffers for derivatives are shared, then we
// can use fewer buffer bindings. This can be important, since most GL
// implementations will support only up to 4 transform feedback bindings.
//
if (deriv1 && sharedDerivativeBuffers) {
outputs.push_back("gl_NextBuffer");
int primvar1Offset = (duDesc.offset % duDesc.stride);
int primvar2Offset = (dvDesc.offset % dvDesc.stride);
for (int i = 0; i < primvar1Offset; ++i) {
outputs.push_back("gl_SkipComponents1");
}
for (int i = 0; i < duDesc.length; ++i) {
snprintf(attrName, sizeof(attrName), "outDeriv1Buffer[%d]", i);
outputs.push_back(attrName);
}
for (int i = primvar1Offset + duDesc.length; i < primvar2Offset; ++i) {
outputs.push_back("gl_SkipComponents1");
}
for (int i = 0; i < dvDesc.length; ++i) {
snprintf(attrName, sizeof(attrName), "outDeriv1Buffer[%d]", i+duDesc.length);
outputs.push_back(attrName);
}
for (int i = primvar2Offset + dvDesc.length; i < dvDesc.stride; ++i) {
outputs.push_back("gl_SkipComponents1");
}
} else {
if (duDesc.length) {
outputs.push_back("gl_NextBuffer");
int primvarOffset = (duDesc.offset % duDesc.stride);
for (int i = 0; i < primvarOffset; ++i) {
@ -240,6 +307,85 @@ compileKernel(BufferDescriptor const &srcDesc,
outputs.push_back("gl_SkipComponents1");
}
}
}
if (deriv2 && sharedDerivativeBuffers) {
outputs.push_back("gl_NextBuffer");
int primvar1Offset = (duuDesc.offset % duuDesc.stride);
int primvar2Offset = (duvDesc.offset % duvDesc.stride);
int primvar3Offset = (dvvDesc.offset % dvvDesc.stride);
for (int i = 0; i < primvar1Offset; ++i) {
outputs.push_back("gl_SkipComponents1");
}
for (int i = 0; i < duuDesc.length; ++i) {
snprintf(attrName, sizeof(attrName), "outDeriv2Buffer[%d]", i);
outputs.push_back(attrName);
}
for (int i = primvar1Offset + duuDesc.length; i < primvar2Offset; ++i) {
outputs.push_back("gl_SkipComponents1");
}
for (int i = 0; i < duvDesc.length; ++i) {
snprintf(attrName, sizeof(attrName), "outDeriv2Buffer[%d]", i+duuDesc.length);
outputs.push_back(attrName);
}
for (int i = primvar2Offset + duvDesc.length; i < primvar3Offset; ++i) {
outputs.push_back("gl_SkipComponents1");
}
for (int i = 0; i < dvvDesc.length; ++i) {
snprintf(attrName, sizeof(attrName), "outDeriv2Buffer[%d]", i+duuDesc.length+duvDesc.length);
outputs.push_back(attrName);
}
for (int i = primvar3Offset + dvvDesc.length; i < dvvDesc.stride; ++i) {
outputs.push_back("gl_SkipComponents1");
}
} else {
if (duuDesc.length) {
outputs.push_back("gl_NextBuffer");
int primvarOffset = (duuDesc.offset % duuDesc.stride);
for (int i = 0; i < primvarOffset; ++i) {
outputs.push_back("gl_SkipComponents1");
}
for (int i = 0; i < duuDesc.length; ++i) {
snprintf(attrName, sizeof(attrName), "outDuuBuffer[%d]", i);
outputs.push_back(attrName);
}
for (int i = primvarOffset + duuDesc.length; i < duuDesc.stride; ++i) {
outputs.push_back("gl_SkipComponents1");
}
}
if (duvDesc.length) {
outputs.push_back("gl_NextBuffer");
int primvarOffset = (duvDesc.offset % duvDesc.stride);
for (int i = 0; i < primvarOffset; ++i) {
outputs.push_back("gl_SkipComponents1");
}
for (int i = 0; i < duvDesc.length; ++i) {
snprintf(attrName, sizeof(attrName), "outDuvBuffer[%d]", i);
outputs.push_back(attrName);
}
for (int i = primvarOffset + duvDesc.length; i < duvDesc.stride; ++i) {
outputs.push_back("gl_SkipComponents1");
}
}
if (dvvDesc.length) {
outputs.push_back("gl_NextBuffer");
int primvarOffset = (dvvDesc.offset % dvvDesc.stride);
for (int i = 0; i < primvarOffset; ++i) {
outputs.push_back("gl_SkipComponents1");
}
for (int i = 0; i < dvvDesc.length; ++i) {
snprintf(attrName, sizeof(attrName), "outDvvBuffer[%d]", i);
outputs.push_back(attrName);
}
for (int i = primvarOffset + dvvDesc.length; i < dvvDesc.stride; ++i) {
outputs.push_back("gl_SkipComponents1");
}
}
}
// convert to char* array
std::vector<const char *> pOutputs;
for (size_t i = 0; i < outputs.size(); ++i) {
@ -274,13 +420,20 @@ bool
GLXFBEvaluator::Compile(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc) {
BufferDescriptor const &dvDesc,
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc) {
// create a stencil kernel
_stencilKernel.Compile(srcDesc, dstDesc, duDesc, dvDesc);
_stencilKernel.Compile(srcDesc, dstDesc, duDesc, dvDesc,
duuDesc, duvDesc, dvvDesc,
_sharedDerivativeBuffers);
// create a patch kernel
_patchKernel.Compile(srcDesc, dstDesc, duDesc, dvDesc);
_patchKernel.Compile(srcDesc, dstDesc, duDesc, dvDesc,
duuDesc, duvDesc, dvvDesc,
_sharedDerivativeBuffers);
// create a texture for input buffer
if (!_srcBufferTexture) {
@ -314,12 +467,46 @@ GLXFBEvaluator::EvalStencils(
GLuint dstBuffer, BufferDescriptor const &dstDesc,
GLuint duBuffer, BufferDescriptor const &duDesc,
GLuint dvBuffer, BufferDescriptor const &dvDesc,
GLuint sizesBuffer,
GLuint offsetsBuffer,
GLuint indicesBuffer,
GLuint weightsBuffer,
GLuint duWeightsBuffer,
GLuint dvWeightsBuffer,
int start, int end) const {
return EvalStencils(srcBuffer, srcDesc,
dstBuffer, dstDesc,
duBuffer, duDesc,
dvBuffer, dvDesc,
0, BufferDescriptor(),
0, BufferDescriptor(),
0, BufferDescriptor(),
sizesBuffer, offsetsBuffer, indicesBuffer,
weightsBuffer,
duWeightsBuffer, dvWeightsBuffer,
0, 0, 0,
start, end);
}
bool
GLXFBEvaluator::EvalStencils(
GLuint srcBuffer, BufferDescriptor const &srcDesc,
GLuint dstBuffer, BufferDescriptor const &dstDesc,
GLuint duBuffer, BufferDescriptor const &duDesc,
GLuint dvBuffer, BufferDescriptor const &dvDesc,
GLuint duuBuffer, BufferDescriptor const &duuDesc,
GLuint duvBuffer, BufferDescriptor const &duvDesc,
GLuint dvvBuffer, BufferDescriptor const &dvvDesc,
GLuint sizesTexture,
GLuint offsetsTexture,
GLuint indicesTexture,
GLuint weightsTexture,
GLuint duWeightsTexture,
GLuint dvWeightsTexture,
GLuint duuWeightsTexture,
GLuint duvWeightsTexture,
GLuint dvvWeightsTexture,
int start, int end) const {
if (!_stencilKernel.program) return false;
@ -353,6 +540,14 @@ GLXFBEvaluator::EvalStencils(
bindTexture(_stencilKernel.uniformDuWeightsTexture, duWeightsTexture, 5);
if (_stencilKernel.uniformDvWeightsTexture >= 0 && dvWeightsTexture)
bindTexture(_stencilKernel.uniformDvWeightsTexture, dvWeightsTexture, 6);
if (_stencilKernel.uniformDuWeightsTexture >= 0 && duWeightsTexture)
bindTexture(_stencilKernel.uniformDuWeightsTexture, duWeightsTexture, 5);
if (_stencilKernel.uniformDuuWeightsTexture >= 0 && duuWeightsTexture)
bindTexture(_stencilKernel.uniformDuuWeightsTexture, duuWeightsTexture, 5);
if (_stencilKernel.uniformDuvWeightsTexture >= 0 && duvWeightsTexture)
bindTexture(_stencilKernel.uniformDuvWeightsTexture, duvWeightsTexture, 6);
if (_stencilKernel.uniformDvvWeightsTexture >= 0 && dvvWeightsTexture)
bindTexture(_stencilKernel.uniformDvvWeightsTexture, dvvWeightsTexture, 6);
// set batch range
glUniform1i(_stencilKernel.uniformStart, start);
@ -392,6 +587,12 @@ GLXFBEvaluator::EvalStencils(
(duDesc.offset - (duDesc.offset % duDesc.stride)) : 0;
int dvBufferBindOffset = dvDesc.stride ?
(dvDesc.offset - (dvDesc.offset % dvDesc.stride)) : 0;
int duuBufferBindOffset = duuDesc.stride ?
(duuDesc.offset - (duuDesc.offset % duuDesc.stride)) : 0;
int duvBufferBindOffset = duvDesc.stride ?
(duvDesc.offset - (duvDesc.offset % duvDesc.stride)) : 0;
int dvvBufferBindOffset = dvvDesc.stride ?
(dvvDesc.offset - (dvvDesc.offset % dvvDesc.stride)) : 0;
// bind destination buffer
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
@ -399,6 +600,12 @@ GLXFBEvaluator::EvalStencils(
dstBufferBindOffset * sizeof(float),
count * dstDesc.stride * sizeof(float));
if ((duDesc.length > 0) && _sharedDerivativeBuffers) {
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
1, duBuffer,
duBufferBindOffset * sizeof(float),
count * duDesc.stride * sizeof(float));
} else {
if (duDesc.length > 0) {
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
1, duBuffer,
@ -412,6 +619,35 @@ GLXFBEvaluator::EvalStencils(
dvBufferBindOffset * sizeof(float),
count * dvDesc.stride * sizeof(float));
}
}
if ((duuDesc.length > 0) && _sharedDerivativeBuffers) {
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
2, duuBuffer,
duuBufferBindOffset * sizeof(float),
count * duuDesc.stride * sizeof(float));
} else {
if (duuDesc.length > 0) {
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
3, duuBuffer,
duuBufferBindOffset * sizeof(float),
count * duuDesc.stride * sizeof(float));
}
if (duvDesc.length > 0) {
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
4, duvBuffer,
duvBufferBindOffset * sizeof(float),
count * duvDesc.stride * sizeof(float));
}
if (dvvDesc.length > 0) {
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
5, dvvBuffer,
dvvBufferBindOffset * sizeof(float),
count * dvvDesc.stride * sizeof(float));
}
}
glBeginTransformFeedback(GL_POINTS);
glDrawArrays(GL_POINTS, 0, count);
@ -419,7 +655,7 @@ GLXFBEvaluator::EvalStencils(
glBindBuffer(GL_TRANSFORM_FEEDBACK_BUFFER, 0);
for (int i = 0; i < 5; ++i) {
for (int i = 0; i < 6; ++i) {
glActiveTexture(GL_TEXTURE0 + i);
glBindTexture(GL_TEXTURE_BUFFER, 0);
}
@ -448,7 +684,36 @@ GLXFBEvaluator::EvalPatches(
GLuint patchIndexTexture,
GLuint patchParamTexture) const {
bool derivatives = (duDesc.length > 0 || dvDesc.length > 0);
return EvalPatches(srcBuffer, srcDesc,
dstBuffer, dstDesc,
duBuffer, duDesc,
dvBuffer, dvDesc,
0, BufferDescriptor(),
0, BufferDescriptor(),
0, BufferDescriptor(),
numPatchCoords,
patchCoordsBuffer, patchArrays,
patchIndexTexture,
patchParamTexture);
}
bool
GLXFBEvaluator::EvalPatches(
GLuint srcBuffer, BufferDescriptor const &srcDesc,
GLuint dstBuffer, BufferDescriptor const &dstDesc,
GLuint duBuffer, BufferDescriptor const &duDesc,
GLuint dvBuffer, BufferDescriptor const &dvDesc,
GLuint duuBuffer, BufferDescriptor const &duuDesc,
GLuint duvBuffer, BufferDescriptor const &duvDesc,
GLuint dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
GLuint patchCoordsBuffer,
const PatchArrayVector &patchArrays,
GLuint patchIndexTexture,
GLuint patchParamTexture) const {
bool deriv1 = (duDesc.length > 0 || dvDesc.length > 0);
bool deriv2 = (duuDesc.length > 0 || duvDesc.length > 0 || dvvDesc.length > 0);
if (!_patchKernel.program) return false;
@ -493,6 +758,15 @@ GLXFBEvaluator::EvalPatches(
int dvBufferBindOffset = dvDesc.stride
? (dvDesc.offset - (dvDesc.offset % dvDesc.stride))
: 0;
int duuBufferBindOffset = duuDesc.stride
? (duuDesc.offset - (duuDesc.offset % duuDesc.stride))
: 0;
int duvBufferBindOffset = duvDesc.stride
? (duvDesc.offset - (duvDesc.offset % duvDesc.stride))
: 0;
int dvvBufferBindOffset = dvvDesc.stride
? (dvvDesc.offset - (dvvDesc.offset % dvvDesc.stride))
: 0;
// bind destination buffer
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
@ -500,7 +774,12 @@ GLXFBEvaluator::EvalPatches(
dstBufferBindOffset * sizeof(float),
numPatchCoords * dstDesc.stride * sizeof(float));
if (derivatives) {
if (deriv1 && _sharedDerivativeBuffers) {
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
1, duBuffer,
duBufferBindOffset * sizeof(float),
numPatchCoords * duDesc.stride * sizeof(float));
} else if (deriv1) {
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
1, duBuffer,
duBufferBindOffset * sizeof(float),
@ -510,7 +789,27 @@ GLXFBEvaluator::EvalPatches(
2, dvBuffer,
dvBufferBindOffset * sizeof(float),
numPatchCoords * dvDesc.stride * sizeof(float));
}
if (deriv2 && _sharedDerivativeBuffers) {
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
2, duuBuffer,
duuBufferBindOffset * sizeof(float),
numPatchCoords * duuDesc.stride * sizeof(float));
} else if (deriv2) {
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
3, duuBuffer,
duuBufferBindOffset * sizeof(float),
numPatchCoords * duuDesc.stride * sizeof(float));
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
4, duvBuffer,
duvBufferBindOffset * sizeof(float),
numPatchCoords * duvDesc.stride * sizeof(float));
glBindBufferRange(GL_TRANSFORM_FEEDBACK_BUFFER,
5, dvvBuffer,
dvvBufferBindOffset * sizeof(float),
numPatchCoords * dvvDesc.stride * sizeof(float));
}
glBeginTransformFeedback(GL_POINTS);
@ -520,7 +819,7 @@ GLXFBEvaluator::EvalPatches(
glBindBuffer(GL_TRANSFORM_FEEDBACK_BUFFER, 0);
// unbind textures
for (int i = 0; i < 3; ++i) {
for (int i = 0; i < 6; ++i) {
glActiveTexture(GL_TEXTURE0 + i);
glBindTexture(GL_TEXTURE_BUFFER, 0);
}
@ -536,7 +835,6 @@ GLXFBEvaluator::EvalPatches(
glBindVertexArray(0);
glDeleteVertexArrays(1, &vao);
return true;
}
@ -554,19 +852,22 @@ bool
GLXFBEvaluator::_StencilKernel::Compile(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc) {
BufferDescriptor const &dvDesc,
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc,
bool sharedDerivativeBuffers) {
// create stencil kernel
if (program) {
glDeleteProgram(program);
}
bool derivatives = (duDesc.length > 0 || dvDesc.length > 0);
const char *kernelDef = derivatives
? "#define OPENSUBDIV_GLSL_XFB_KERNEL_EVAL_STENCILS\n"
"#define OPENSUBDIV_GLSL_XFB_USE_DERIVATIVES\n"
: "#define OPENSUBDIV_GLSL_XFB_KERNEL_EVAL_STENCILS\n";
const char * kernelDefines =
"#define OPENSUBDIV_GLSL_XFB_KERNEL_EVAL_STENCILS\n";
program = compileKernel(srcDesc, dstDesc, duDesc, dvDesc, kernelDef);
program = compileKernel(srcDesc, dstDesc, duDesc, dvDesc,
duuDesc, duvDesc, dvvDesc,
kernelDefines, sharedDerivativeBuffers);
if (program == 0) return false;
// cache uniform locations (TODO: use uniform block)
@ -578,6 +879,9 @@ GLXFBEvaluator::_StencilKernel::Compile(BufferDescriptor const &srcDesc,
uniformWeightsTexture = glGetUniformLocation(program, "weights");
uniformDuWeightsTexture = glGetUniformLocation(program, "duWeights");
uniformDvWeightsTexture = glGetUniformLocation(program, "dvWeights");
uniformDuuWeightsTexture = glGetUniformLocation(program, "duuWeights");
uniformDuvWeightsTexture = glGetUniformLocation(program, "duvWeights");
uniformDvvWeightsTexture = glGetUniformLocation(program, "dvvWeights");
uniformStart = glGetUniformLocation(program, "batchStart");
uniformEnd = glGetUniformLocation(program, "batchEnd");
@ -598,19 +902,22 @@ bool
GLXFBEvaluator::_PatchKernel::Compile(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc) {
BufferDescriptor const &dvDesc,
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc,
bool sharedDerivativeBuffers) {
// create stencil kernel
if (program) {
glDeleteProgram(program);
}
bool derivatives = (duDesc.length > 0 || dvDesc.length > 0);
const char *kernelDef = derivatives
? "#define OPENSUBDIV_GLSL_XFB_KERNEL_EVAL_PATCHES\n"
"#define OPENSUBDIV_GLSL_XFB_USE_DERIVATIVES\n"
: "#define OPENSUBDIV_GLSL_XFB_KERNEL_EVAL_PATCHES\n";
const char * kernelDefines =
"#define OPENSUBDIV_GLSL_XFB_KERNEL_EVAL_PATCHES\n";
program = compileKernel(srcDesc, dstDesc, duDesc, dvDesc, kernelDef);
program = compileKernel(srcDesc, dstDesc, duDesc, dvDesc,
duuDesc, duvDesc, dvvDesc,
kernelDefines, sharedDerivativeBuffers);
if (program == 0) return false;
// cache uniform locations
@ -623,7 +930,6 @@ GLXFBEvaluator::_PatchKernel::Compile(BufferDescriptor const &srcDesc,
return true;
}
} // end namespace Osd
} // end namespace OPENSUBDIV_VERSION

1326
opensubdiv/osd/glXFBEvaluator.h
View File

File diff suppressed because it is too large Load Diff

96
opensubdiv/osd/glslComputeKernel.glsl
View File

@ -37,13 +37,22 @@ layout(binding=1) buffer dst_buffer { float dstVertexBuffer[]; };
// derivative buffers (if needed)
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_DERIVATIVES)
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES)
uniform ivec3 duDesc;
uniform ivec3 dvDesc;
layout(binding=2) buffer du_buffer { float duBuffer[]; };
layout(binding=3) buffer dv_buffer { float dvBuffer[]; };
#endif
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES)
uniform ivec3 duuDesc;
uniform ivec3 duvDesc;
uniform ivec3 dvvDesc;
layout(binding=10) buffer duu_buffer { float duuBuffer[]; };
layout(binding=11) buffer duv_buffer { float duvBuffer[]; };
layout(binding=12) buffer dvv_buffer { float dvvBuffer[]; };
#endif
// stencil buffers
#if defined(OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_STENCILS)
@ -55,11 +64,17 @@ layout(binding=5) buffer stencilOffsets { int _offsets[]; };
layout(binding=6) buffer stencilIndices { int _indices[]; };
layout(binding=7) buffer stencilWeights { float _weights[]; };
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_DERIVATIVES)
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES)
layout(binding=8) buffer stencilDuWeights { float _duWeights[]; };
layout(binding=9) buffer stencilDvWeights { float _dvWeights[]; };
#endif
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES)
layout(binding=13) buffer stencilDuuWeights { float _duuWeights[]; };
layout(binding=14) buffer stencilDuvWeights { float _duvWeights[]; };
layout(binding=15) buffer stencilDvvWeights { float _dvvWeights[]; };
#endif
#endif
// patch buffers
@ -119,7 +134,7 @@ void addWithWeight(inout Vertex v, const Vertex src, float weight) {
}
}
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_DERIVATIVES)
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES)
void writeDu(int index, Vertex du) {
int duIndex = duDesc.x + index * duDesc.z;
for (int i = 0; i < LENGTH; ++i) {
@ -135,6 +150,29 @@ void writeDv(int index, Vertex dv) {
}
#endif
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES)
void writeDuu(int index, Vertex duu) {
int duuIndex = duuDesc.x + index * duuDesc.z;
for (int i = 0; i < LENGTH; ++i) {
duuBuffer[duuIndex + i] = duu.vertexData[i];
}
}
void writeDuv(int index, Vertex duv) {
int duvIndex = duvDesc.x + index * duvDesc.z;
for (int i = 0; i < LENGTH; ++i) {
duvBuffer[duvIndex + i] = duv.vertexData[i];
}
}
void writeDvv(int index, Vertex dvv) {
int dvvIndex = dvvDesc.x + index * dvvDesc.z;
for (int i = 0; i < LENGTH; ++i) {
dvvBuffer[dvvIndex + i] = dvv.vertexData[i];
}
}
#endif
//------------------------------------------------------------------------------
#if defined(OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_STENCILS)
@ -160,7 +198,7 @@ void main() {
writeVertex(current, dst);
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_DERIVATIVES)
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES)
Vertex du, dv;
clear(du);
clear(dv);
@ -178,6 +216,29 @@ void main() {
writeDv(current, dv);
}
#endif
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES)
Vertex duu, duv, dvv;
clear(duu);
clear(duv);
clear(dvv);
for (int i=0; i<size; ++i) {
// expects the compiler optimizes readVertex out here.
Vertex src = readVertex(_indices[offset+i]);
addWithWeight(duu, src, _duuWeights[offset+i]);
addWithWeight(duv, src, _duvWeights[offset+i]);
addWithWeight(dvv, src, _dvvWeights[offset+i]);
}
if (duuDesc.y > 0) { // length
writeDuu(current, duu);
}
if (duvDesc.y > 0) {
writeDuv(current, duv);
}
if (dvvDesc.y > 0) {
writeDvv(current, dvv);
}
#endif
}
#endif
@ -260,6 +321,9 @@ void main() {
wP[i] = wP4[i];
wDs[i] = wDs4[i];
wDt[i] = wDt4[i];
wDss[i] = wDss4[i];
wDst[i] = wDst4[i];
wDtt[i] = wDtt4[i];
}
} else if (patchType == 6) {
float wP16[16], wDs16[16], wDt16[16], wDss16[16], wDst16[16], wDtt16[16];
@ -269,16 +333,22 @@ void main() {
wP[i] = wP16[i];
wDs[i] = wDs16[i];
wDt[i] = wDt16[i];
wDss[i] = wDss16[i];
wDst[i] = wDst16[i];
wDtt[i] = wDtt16[i];
}
} else if (patchType == 9) {
OsdGetGregoryPatchWeights(uv.s, uv.t, dScale, wP, wDs, wDt, wDss, wDst, wDtt);
numControlVertices = 20;
}
Vertex dst, du, dv;
Vertex dst, du, dv, duu, duv, dvv;
clear(dst);
clear(du);
clear(dv);
clear(duu);
clear(duv);
clear(dvv);
int indexStride = getNumControlVertices(array.x);
int indexBase = array.z + indexStride * (patchIndex - array.w);
@ -288,10 +358,13 @@ void main() {
addWithWeight(dst, readVertex(index), wP[cv]);
addWithWeight(du, readVertex(index), wDs[cv]);
addWithWeight(dv, readVertex(index), wDt[cv]);
addWithWeight(duu, readVertex(index), wDss[cv]);
addWithWeight(duv, readVertex(index), wDst[cv]);
addWithWeight(dvv, readVertex(index), wDtt[cv]);
}
writeVertex(current, dst);
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_DERIVATIVES)
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_1ST_DERIVATIVES)
if (duDesc.y > 0) { // length
writeDu(current, du);
}
@ -299,6 +372,17 @@ void main() {
writeDv(current, dv);
}
#endif
#if defined(OPENSUBDIV_GLSL_COMPUTE_USE_2ND_DERIVATIVES)
if (duuDesc.y > 0) { // length
writeDuu(current, duu);
}
if (duvDesc.y > 0) { // length
writeDuv(current, duv);
}
if (dvvDesc.y > 0) {
writeDvv(current, dvv);
}
#endif
}
#endif

123
opensubdiv/osd/glslXFBKernel.glsl
View File

@ -63,7 +63,22 @@ void writeVertex(Vertex v) {
//------------------------------------------------------------------------------
#if defined(OPENSUBDIV_GLSL_XFB_USE_DERIVATIVES)
#if defined(OPENSUBDIV_GLSL_XFB_USE_1ST_DERIVATIVES) && \
defined(OPENSUBDIV_GLSL_XFB_SHARED_1ST_DERIVATIVE_BUFFERS)
out float outDeriv1Buffer[2*LENGTH];
void writeDu(Vertex v) {
for(int i = 0; i < LENGTH; i++) {
outDeriv1Buffer[i] = v.vertexData[i];
}
}
void writeDv(Vertex v) {
for(int i = 0; i < LENGTH; i++) {
outDeriv1Buffer[i+LENGTH] = v.vertexData[i];
}
}
#elif defined(OPENSUBDIV_GLSL_XFB_USE_1ST_DERIVATIVES)
out float outDuBuffer[LENGTH];
out float outDvBuffer[LENGTH];
@ -80,6 +95,51 @@ void writeDv(Vertex v) {
}
#endif
#if defined(OPENSUBDIV_GLSL_XFB_USE_2ND_DERIVATIVES) && \
defined(OPENSUBDIV_GLSL_XFB_SHARED_2ND_DERIVATIVE_BUFFERS)
out float outDeriv2Buffer[3*LENGTH];
void writeDuu(Vertex v) {
for(int i = 0; i < LENGTH; i++) {
outDeriv2Buffer[i] = v.vertexData[i];
}
}
void writeDuv(Vertex v) {
for(int i = 0; i < LENGTH; i++) {
outDeriv2Buffer[i+LENGTH] = v.vertexData[i];
}
}
void writeDvv(Vertex v) {
for(int i = 0; i < LENGTH; i++) {
outDeriv2Buffer[i+2*LENGTH] = v.vertexData[i];
}
}
#elif defined(OPENSUBDIV_GLSL_XFB_USE_2ND_DERIVATIVES)
out float outDuuBuffer[LENGTH];
out float outDuvBuffer[LENGTH];
out float outDvvBuffer[LENGTH];
void writeDuu(Vertex v) {
for(int i = 0; i < LENGTH; i++) {
outDuuBuffer[i] = v.vertexData[i];
}
}
void writeDuv(Vertex v) {
for(int i = 0; i < LENGTH; i++) {
outDuvBuffer[i] = v.vertexData[i];
}
}
void writeDvv(Vertex v) {
for(int i = 0; i < LENGTH; i++) {
outDvvBuffer[i] = v.vertexData[i];
}
}
#endif
//------------------------------------------------------------------------------
#if defined(OPENSUBDIV_GLSL_XFB_KERNEL_EVAL_STENCILS)
@ -89,11 +149,17 @@ uniform isamplerBuffer offsets;
uniform isamplerBuffer indices;
uniform samplerBuffer weights;
#if defined(OPENSUBDIV_GLSL_XFB_USE_DERIVATIVES)
#if defined(OPENSUBDIV_GLSL_XFB_USE_1ST_DERIVATIVES)
uniform samplerBuffer duWeights;
uniform samplerBuffer dvWeights;
#endif
#if defined(OPENSUBDIV_GLSL_XFB_USE_2ND_DERIVATIVES)
uniform samplerBuffer duuWeights;
uniform samplerBuffer duvWeights;
uniform samplerBuffer dvvWeights;
#endif
uniform int batchStart = 0;
uniform int batchEnd = 0;
@ -104,10 +170,13 @@ void main() {
return;
}
Vertex dst, du, dv;
Vertex dst, du, dv, duu, duv, dvv;
clear(dst);
clear(du);
clear(dv);
clear(duu);
clear(duv);
clear(dvv);
int offset = texelFetch(offsets, current).x;
uint size = texelFetch(sizes, current).x;
@ -117,19 +186,32 @@ void main() {
float weight = texelFetch(weights, offset+stencil).x;
addWithWeight(dst, readVertex( index ), weight);
#if defined(OPENSUBDIV_GLSL_XFB_USE_DERIVATIVES)
#if defined(OPENSUBDIV_GLSL_XFB_USE_1ST_DERIVATIVES)
float duWeight = texelFetch(duWeights, offset+stencil).x;
float dvWeight = texelFetch(dvWeights, offset+stencil).x;
addWithWeight(du, readVertex(index), duWeight);
addWithWeight(dv, readVertex(index), dvWeight);
#endif
#if defined(OPENSUBDIV_GLSL_XFB_USE_2ND_DERIVATIVES)
float duuWeight = texelFetch(duuWeights, offset+stencil).x;
float duvWeight = texelFetch(duvWeights, offset+stencil).x;
float dvvWeight = texelFetch(dvvWeights, offset+stencil).x;
addWithWeight(duu, readVertex(index), duuWeight);
addWithWeight(duv, readVertex(index), duvWeight);
addWithWeight(dvv, readVertex(index), dvvWeight);
#endif
}
writeVertex(dst);
#if defined(OPENSUBDIV_GLSL_XFB_USE_DERIVATIVES)
#if defined(OPENSUBDIV_GLSL_XFB_USE_1ST_DERIVATIVES)
writeDu(du);
writeDv(dv);
#endif
#if defined(OPENSUBDIV_GLSL_XFB_USE_2ND_DERIVATIVES)
writeDuu(duu);
writeDuv(duv);
writeDvv(dvv);
#endif
}
#endif
@ -213,31 +295,43 @@ void main() {
int numControlVertices = 0;
if (patchType == 3) {
float wP4[4], wDs4[4], wDt4[4], wDss4[4], wDst4[4], wDtt4[4];
OsdGetBilinearPatchWeights(coord.s, coord.t, dScale, wP4, wDs4, wDt4, wDss4, wDst4, wDtt4);
OsdGetBilinearPatchWeights(coord.s, coord.t, dScale, wP4,
wDs4, wDt4, wDss4, wDst4, wDtt4);
numControlVertices = 4;
for (int i=0; i<numControlVertices; ++i) {
wP[i] = wP4[i];
wDs[i] = wDs4[i];
wDt[i] = wDt4[i];
wDss[i] = wDss4[i];
wDst[i] = wDst4[i];
wDtt[i] = wDtt4[i];
}
} else if (patchType == 6) {
float wP16[16], wDs16[16], wDt16[16], wDss16[16], wDst16[16], wDtt16[16];
OsdGetBSplinePatchWeights(coord.s, coord.t, dScale, boundary, wP16, wDs16, wDt16, wDss16, wDst16, wDtt16);
OsdGetBSplinePatchWeights(coord.s, coord.t, dScale, boundary, wP16,
wDs16, wDt16, wDss16, wDst16, wDtt16);
numControlVertices = 16;
for (int i=0; i<numControlVertices; ++i) {
wP[i] = wP16[i];
wDs[i] = wDs16[i];
wDt[i] = wDt16[i];
wDss[i] = wDss16[i];
wDst[i] = wDst16[i];
wDtt[i] = wDtt16[i];
}
} else if (patchType == 9) {
OsdGetGregoryPatchWeights(coord.s, coord.t, dScale, wP, wDs, wDt, wDss, wDst, wDtt);
OsdGetGregoryPatchWeights(coord.s, coord.t, dScale, wP,
wDs, wDt, wDss, wDst, wDtt);
numControlVertices = 20;
}
Vertex dst, du, dv;
Vertex dst, du, dv, duu, duv, dvv;
clear(dst);
clear(du);
clear(dv);
clear(duu);
clear(duv);
clear(dvv);
int indexStride = getNumControlVertices(array.x);
int indexBase = array.z + indexStride * (patchIndex - array.w);
@ -247,15 +341,22 @@ void main() {
addWithWeight(dst, readVertex(index), wP[cv]);
addWithWeight(du, readVertex(index), wDs[cv]);
addWithWeight(dv, readVertex(index), wDt[cv]);
addWithWeight(duu, readVertex(index), wDss[cv]);
addWithWeight(duv, readVertex(index), wDst[cv]);
addWithWeight(dvv, readVertex(index), wDtt[cv]);
}
writeVertex(dst);
#if defined(OPENSUBDIV_GLSL_XFB_USE_DERIVATIVES)
#if defined(OPENSUBDIV_GLSL_XFB_USE_1ST_DERIVATIVES)
writeDu(du);
writeDv(dv);
#endif
#if defined(OPENSUBDIV_GLSL_XFB_USE_2ND_DERIVATIVES)
writeDuu(duu);
writeDuv(duv);
writeDvv(dvv);
#endif
}
#endif

90
opensubdiv/osd/mesh.h
View File

@ -176,7 +176,7 @@ convertToCompatibleStencilTable<Far::StencilTable, Far::StencilTable, ID3D11Devi
// ---------------------------------------------------------------------------
// Osd evaluator cache: for the GPU backends require compiled instance
// (GLXFB, GLCompue, CL)
// (GLXFB, GLCompute, CL)
//
// note: this is just an example usage and client applications are supposed
// to implement their own structure for Evaluator instance.
@ -198,8 +198,27 @@ public:
BufferDescriptor const &duDescArg,
BufferDescriptor const &dvDescArg,
EVALUATOR *evalArg) : srcDesc(srcDescArg), dstDesc(dstDescArg),
duDesc(duDescArg), dvDesc(dvDescArg), evaluator(evalArg) {}
BufferDescriptor srcDesc, dstDesc, duDesc, dvDesc;
duDesc(duDescArg), dvDesc(dvDescArg),
duuDesc(BufferDescriptor()),
duvDesc(BufferDescriptor()),
dvvDesc(BufferDescriptor()),
evaluator(evalArg) {}
Entry(BufferDescriptor const &srcDescArg,
BufferDescriptor const &dstDescArg,
BufferDescriptor const &duDescArg,
BufferDescriptor const &dvDescArg,
BufferDescriptor const &duuDescArg,
BufferDescriptor const &duvDescArg,
BufferDescriptor const &dvvDescArg,
EVALUATOR *evalArg) : srcDesc(srcDescArg), dstDesc(dstDescArg),
duDesc(duDescArg), dvDesc(dvDescArg),
duuDesc(duuDescArg),
duvDesc(duvDescArg),
dvvDesc(dvvDescArg),
evaluator(evalArg) {}
BufferDescriptor srcDesc, dstDesc;
BufferDescriptor duDesc, dvDesc;
BufferDescriptor duuDesc, duvDesc, dvvDesc;
EVALUATOR *evaluator;
};
typedef std::vector<Entry> Evaluators;
@ -209,6 +228,9 @@ public:
BufferDescriptor const &dstDesc,
DEVICE_CONTEXT *deviceContext) {
return GetEvaluator(srcDesc, dstDesc,
BufferDescriptor(),
BufferDescriptor(),
BufferDescriptor(),
BufferDescriptor(),
BufferDescriptor(),
deviceContext);
@ -220,20 +242,43 @@ public:
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc,
DEVICE_CONTEXT *deviceContext) {
return GetEvaluator(srcDesc, dstDesc,
duDesc, dvDesc,
BufferDescriptor(),
BufferDescriptor(),
BufferDescriptor(),
deviceContext);
}
template <typename DEVICE_CONTEXT>
EVALUATOR *GetEvaluator(BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc,
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc,
DEVICE_CONTEXT *deviceContext) {
for(typename Evaluators::iterator it = _evaluators.begin();
it != _evaluators.end(); ++it) {
if (isEqual(srcDesc, it->srcDesc) &&
isEqual(dstDesc, it->dstDesc) &&
isEqual(duDesc, it->duDesc) &&
isEqual(dvDesc, it->dvDesc)) {
isEqual(dvDesc, it->dvDesc) &&
isEqual(duuDesc, it->duuDesc) &&
isEqual(duvDesc, it->duvDesc) &&
isEqual(dvvDesc, it->dvvDesc)) {
return it->evaluator;
}
}
EVALUATOR *e = EVALUATOR::Create(srcDesc, dstDesc,
duDesc, dvDesc,
duuDesc, duvDesc, dvvDesc,
deviceContext);
_evaluators.push_back(Entry(srcDesc, dstDesc, duDesc, dvDesc, e));
_evaluators.push_back(Entry(srcDesc, dstDesc,
duDesc, dvDesc,
duuDesc, duvDesc, dvvDesc, e));
return e;
}
@ -273,6 +318,25 @@ struct enable_if<false, T> { };
/// @endcond
// extract a kernel from cache if available
template <typename EVALUATOR, typename DEVICE_CONTEXT>
static EVALUATOR *GetEvaluator(
EvaluatorCacheT<EVALUATOR> *cache,
BufferDescriptor const &srcDesc,
BufferDescriptor const &dstDesc,
BufferDescriptor const &duDesc,
BufferDescriptor const &dvDesc,
BufferDescriptor const &duuDesc,
BufferDescriptor const &duvDesc,
BufferDescriptor const &dvvDesc,
DEVICE_CONTEXT deviceContext,
typename enable_if<instantiatable<EVALUATOR>::value, void>::type*t=0) {
(void)t;
if (cache == NULL) return NULL;
return cache->GetEvaluator(srcDesc, dstDesc,
duDesc, dvDesc, duuDesc, duvDesc, dvvDesc,
deviceContext);
}
template <typename EVALUATOR, typename DEVICE_CONTEXT>
static EVALUATOR *GetEvaluator(
EvaluatorCacheT<EVALUATOR> *cache,
@ -303,6 +367,22 @@ static EVALUATOR *GetEvaluator(
}
// fallback
template <typename EVALUATOR, typename DEVICE_CONTEXT>
static EVALUATOR *GetEvaluator(
EvaluatorCacheT<EVALUATOR> *,
BufferDescriptor const &,
BufferDescriptor const &,
BufferDescriptor const &,
BufferDescriptor const &,
BufferDescriptor const &,
BufferDescriptor const &,
BufferDescriptor const &,
DEVICE_CONTEXT,
typename enable_if<!instantiatable<EVALUATOR>::value, void>::type*t=0) {
(void)t;
return NULL;
}
template <typename EVALUATOR, typename DEVICE_CONTEXT>
static EVALUATOR *GetEvaluator(
EvaluatorCacheT<EVALUATOR> *,

107
opensubdiv/osd/ompEvaluator.cpp
View File

@ -197,7 +197,7 @@ OmpEvaluator::EvalPatches(
#pragma omp parallel for
for (int i = 0; i < numPatchCoords; ++i) {
float wP[20], wDs[20], wDt[20];
float wP[20], wDu[20], wDv[20];
BufferAdapter<float> dstT(dst + dstDesc.stride*i, dstDesc.length, dstDesc.stride);
BufferAdapter<float> duT(du + duDesc.stride*i, duDesc.length, duDesc.stride);
BufferAdapter<float> dvT(dv + dvDesc.stride*i, dvDesc.length, dvDesc.stride);
@ -214,15 +214,15 @@ OmpEvaluator::EvalPatches(
int numControlVertices = 0;
if (patchType == Far::PatchDescriptor::REGULAR) {
Far::internal::GetBSplineWeights(param,
coord.s, coord.t, wP, wDs, wDt);
coord.s, coord.t, wP, wDu, wDv);
numControlVertices = 16;
} else if (patchType == Far::PatchDescriptor::GREGORY_BASIS) {
Far::internal::GetGregoryWeights(param,
coord.s, coord.t, wP, wDs, wDt);
coord.s, coord.t, wP, wDu, wDv);
numControlVertices = 20;
} else if (patchType == Far::PatchDescriptor::QUADS) {
Far::internal::GetBilinearWeights(param,
coord.s, coord.t, wP, wDs, wDt);
coord.s, coord.t, wP, wDu, wDv);
numControlVertices = 4;
} else {
continue;
@ -239,8 +239,8 @@ OmpEvaluator::EvalPatches(
dvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
duT.AddWithWeight(srcT[cvs[j]], wDs[j]);
dvT.AddWithWeight(srcT[cvs[j]], wDt[j]);
duT.AddWithWeight(srcT[cvs[j]], wDu[j]);
dvT.AddWithWeight(srcT[cvs[j]], wDv[j]);
}
++dstT;
++duT;
@ -249,6 +249,101 @@ OmpEvaluator::EvalPatches(
return true;
}
/* static */
bool
OmpEvaluator::EvalPatches(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PatchCoord const *patchCoords,
PatchArray const *patchArrays,
const int *patchIndexBuffer,
PatchParam const *patchParamBuffer) {
src += srcDesc.offset;
if (dst) dst += dstDesc.offset;
if (du) du += duDesc.offset;
if (dv) dv += dvDesc.offset;
if (duu) duu += duuDesc.offset;
if (duv) duv += duvDesc.offset;
if (dvv) dvv += dvvDesc.offset;
BufferAdapter<const float> srcT(src, srcDesc.length, srcDesc.stride);
#pragma omp parallel for
for (int i = 0; i < numPatchCoords; ++i) {
float wP[20], wDu[20], wDv[20], wDuu[20], wDuv[20], wDvv[20];
BufferAdapter<float> dstT(dst + dstDesc.stride*i, dstDesc.length, dstDesc.stride);
BufferAdapter<float> duT(du + duDesc.stride*i, duDesc.length, duDesc.stride);
BufferAdapter<float> dvT(dv + dvDesc.stride*i, dvDesc.length, dvDesc.stride);
BufferAdapter<float> duuT(duu + duuDesc.stride*i, duuDesc.length, duuDesc.stride);
BufferAdapter<float> duvT(duv + duvDesc.stride*i, duvDesc.length, duvDesc.stride);
BufferAdapter<float> dvvT(dvv + dvvDesc.stride*i, dvvDesc.length, dvvDesc.stride);
PatchCoord const &coord = patchCoords[i];
PatchArray const &array = patchArrays[coord.handle.arrayIndex];
Far::PatchParam const & param =
patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
? Far::PatchDescriptor::REGULAR
: array.GetPatchType();
int numControlVertices = 0;
if (patchType == Far::PatchDescriptor::REGULAR) {
Far::internal::GetBSplineWeights(param,
coord.s, coord.t, wP,
wDu, wDv, wDuu, wDuv, wDvv);
numControlVertices = 16;
} else if (patchType == Far::PatchDescriptor::GREGORY_BASIS) {
Far::internal::GetGregoryWeights(param,
coord.s, coord.t, wP,
wDu, wDv, wDuu, wDuv, wDvv);
numControlVertices = 20;
} else if (patchType == Far::PatchDescriptor::QUADS) {
Far::internal::GetBilinearWeights(param,
coord.s, coord.t, wP,
wDu, wDv, wDuu, wDuv, wDvv);
numControlVertices = 4;
} else {
continue;
}
int indexStride = Far::PatchDescriptor(array.GetPatchType()).GetNumControlVertices();
int indexBase = array.GetIndexBase() + indexStride *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
const int *cvs = &patchIndexBuffer[indexBase];
dstT.Clear();
duT.Clear();
dvT.Clear();
duuT.Clear();
duvT.Clear();
dvvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
duT.AddWithWeight(srcT[cvs[j]], wDu[j]);
dvT.AddWithWeight(srcT[cvs[j]], wDv[j]);
duuT.AddWithWeight(srcT[cvs[j]], wDuu[j]);
duvT.AddWithWeight(srcT[cvs[j]], wDuv[j]);
dvvT.AddWithWeight(srcT[cvs[j]], wDvv[j]);
}
++dstT;
++duT;
++dvT;
++duuT;
++duvT;
++dvvT;
}
return true;
}
/* static */
void

689
opensubdiv/osd/ompEvaluator.h
View File

@ -26,11 +26,11 @@
#define OPENSUBDIV3_OSD_OMP_EVALUATOR_H
#include "../version.h"
#include <cstddef>
#include "../osd/bufferDescriptor.h"
#include "../osd/types.h"
#include <cstddef>
namespace OpenSubdiv {
namespace OPENSUBDIV_VERSION {
@ -107,7 +107,6 @@ public:
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
/// to apply for the range [start, end)
///
/// @param offsets pointer to the offsets buffer of the stencil table
///
@ -145,17 +144,17 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output U-derivative buffer
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the output buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output V-derivative buffer
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the output buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param stencilTable Far::StencilTable or equivalent
///
@ -206,15 +205,15 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output U-derivatives pointer. An offset of
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the output buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output V-derivatives pointer. An offset of
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the output buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
///
@ -245,6 +244,177 @@ public:
const float * dvWeights,
int start, int end);
/// \brief Generic static eval stencils function with derivatives.
/// This function has a same signature as other device kernels
/// have so that it can be called in the same way from OsdMesh
/// template interface.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param stencilTable Far::StencilTable or equivalent
///
/// @param instance not used in the omp kernel
/// (declared as a typed pointer to prevent
/// undesirable template resolution)
///
/// @param deviceContext not used in the omp kernel
///
template <typename SRC_BUFFER, typename DST_BUFFER, typename STENCIL_TABLE>
static bool EvalStencils(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
STENCIL_TABLE const *stencilTable,
const OmpEvaluator *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalStencils(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
&stencilTable->GetSizes()[0],
&stencilTable->GetOffsets()[0],
&stencilTable->GetControlIndices()[0],
&stencilTable->GetWeights()[0],
&stencilTable->GetDuWeights()[0],
&stencilTable->GetDvWeights()[0],
&stencilTable->GetDuuWeights()[0],
&stencilTable->GetDuvWeights()[0],
&stencilTable->GetDvvWeights()[0],
/*start = */ 0,
/*end = */ stencilTable->GetNumStencils());
}
/// \brief Static eval stencils function with derivatives, which takes
/// raw CPU pointers for input and output.
///
/// @param src Input primvar pointer. An offset of srcDesc
/// will be applied internally (i.e. the pointer
/// should not include the offset)
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dst Output primvar pointer. An offset of dstDesc
/// will be applied internally.
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duu Output pointer 2nd derivative wrt u. An offset of
/// duuDesc will be applied internally.
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duv Output pointer 2nd derivative wrt u and v. An offset of
/// duvDesc will be applied internally.
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvv Output pointer 2nd derivative wrt v. An offset of
/// dvvDesc will be applied internally.
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
///
/// @param offsets pointer to the offsets buffer of the stencil table
///
/// @param indices pointer to the indices buffer of the stencil table
///
/// @param weights pointer to the weights buffer of the stencil table
///
/// @param duWeights pointer to the du-weights buffer of the stencil table
///
/// @param dvWeights pointer to the dv-weights buffer of the stencil table
///
/// @param duuWeights pointer to the duu-weights buffer of the stencil table
///
/// @param duvWeights pointer to the duv-weights buffer of the stencil table
///
/// @param dvvWeights pointer to the dvv-weights buffer of the stencil table
///
/// @param start start index of stencil table
///
/// @param end end index of stencil table
///
static bool EvalStencils(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
const int * sizes,
const int * offsets,
const int * indices,
const float * weights,
const float * duWeights,
const float * dvWeights,
const float * duuWeights,
const float * duvWeights,
const float * dvvWeights,
int start, int end);
/// ----------------------------------------------------------------------
///
/// Limit evaluations with PatchTable
@ -318,13 +488,13 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output U-derivatives buffer
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output V-derivatives buffer
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
@ -354,6 +524,7 @@ public:
PATCH_TABLE *patchTable,
OmpEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
@ -373,6 +544,102 @@ public:
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function with derivatives. This function has
/// a same signature as other device kernels have so that it can be
/// called in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the omp evaluator
///
/// @param deviceContext not used in the omp evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatches(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
OmpEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
// XXX: PatchCoords is somewhat abusing vertex primvar buffer interop.
// ideally all buffer classes should have templated by datatype
// so that downcast isn't needed there.
// (e.g. Osd::CpuBuffer<PatchCoord> )
//
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetPatchArrayBuffer(),
patchTable->GetPatchIndexBuffer(),
patchTable->GetPatchParamBuffer());
}
/// \brief Static limit eval function. It takes an array of PatchCoord
/// and evaluate limit values on given PatchTable.
///
@ -423,15 +690,15 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output U-derivatives pointer. An offset of
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the du buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output V-derivatives pointer. An offset of
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dv buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
@ -457,6 +724,72 @@ public:
const int *patchIndexBuffer,
PatchParam const *patchParamBuffer);
/// \brief Static limit eval function. It takes an array of PatchCoord
/// and evaluate limit values on given PatchTable.
///
/// @param src Input primvar pointer. An offset of srcDesc
/// will be applied internally (i.e. the pointer
/// should not include the offset)
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dst Output primvar pointer. An offset of dstDesc
/// will be applied internally.
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duu Output pointer 2nd derivative wrt u. An offset of
/// duuDesc will be applied internally.
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duv Output pointer 2nd derivative wrt u and v. An offset of
/// duvDesc will be applied internally.
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvv Output pointer 2nd derivative wrt v. An offset of
/// dvvDesc will be applied internally.
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchArrays an array of Osd::PatchArray struct
/// indexed by PatchCoord::arrayIndex
///
/// @param patchIndexBuffer an array of patch indices
/// indexed by PatchCoord::vertIndex
///
/// @param patchParamBuffer an array of Osd::PatchParam struct
/// indexed by PatchCoord::patchIndex
///
static bool EvalPatches(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PatchCoord const *patchCoords,
PatchArray const *patchArrays,
const int *patchIndexBuffer,
PatchParam const *patchParamBuffer);
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
@ -508,6 +841,164 @@ public:
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the omp evaluator
///
/// @param deviceContext not used in the omp evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
OmpEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetVaryingPatchArrayBuffer(),
patchTable->GetVaryingPatchIndexBuffer(),
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the omp evaluator
///
/// @param deviceContext not used in the omp evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
OmpEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetVaryingPatchArrayBuffer(),
patchTable->GetVaryingPatchIndexBuffer(),
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
@ -562,6 +1053,170 @@ public:
patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param fvarChannel face-varying channel
///
/// @param instance not used in the omp evaluator
///
/// @param deviceContext not used in the omp evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesFaceVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
int fvarChannel,
OmpEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetFVarPatchArrayBuffer(fvarChannel),
patchTable->GetFVarPatchIndexBuffer(fvarChannel),
patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param fvarChannel face-varying channel
///
/// @param instance not used in the omp evaluator
///
/// @param deviceContext not used in the omp evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesFaceVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
int fvarChannel,
OmpEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetFVarPatchArrayBuffer(fvarChannel),
patchTable->GetFVarPatchIndexBuffer(fvarChannel),
patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// ----------------------------------------------------------------------
///
/// Other methods

93
opensubdiv/osd/ompKernel.cpp
View File

@ -177,6 +177,99 @@ OmpEvalStencils(float const * src, BufferDescriptor const &srcDesc,
}
void
OmpEvalStencils(float const * src, BufferDescriptor const &srcDesc,
float * dst, BufferDescriptor const &dstDesc,
float * dstDu, BufferDescriptor const &dstDuDesc,
float * dstDv, BufferDescriptor const &dstDvDesc,
float * dstDuu, BufferDescriptor const &dstDuuDesc,
float * dstDuv, BufferDescriptor const &dstDuvDesc,
float * dstDvv, BufferDescriptor const &dstDvvDesc,
int const * sizes,
int const * offsets,
int const * indices,
float const * weights,
float const * duWeights,
float const * dvWeights,
float const * duuWeights,
float const * duvWeights,
float const * dvvWeights,
int start, int end) {
start = (start > 0 ? start : 0);
src += srcDesc.offset;
dst += dstDesc.offset;
dstDu += dstDuDesc.offset;
dstDv += dstDvDesc.offset;
dstDuu += dstDuuDesc.offset;
dstDuv += dstDuvDesc.offset;
dstDvv += dstDvvDesc.offset;
int numThreads = omp_get_max_threads();
int n = end - start;
float * result = (float*)alloca(srcDesc.length * numThreads * sizeof(float));
float * resultDu = (float*)alloca(srcDesc.length * numThreads * sizeof(float));
float * resultDv = (float*)alloca(srcDesc.length * numThreads * sizeof(float));
float * resultDuu = (float*)alloca(srcDesc.length * numThreads * sizeof(float));
float * resultDuv = (float*)alloca(srcDesc.length * numThreads * sizeof(float));
float * resultDvv = (float*)alloca(srcDesc.length * numThreads * sizeof(float));
#pragma omp parallel for
for (int i = 0; i < n; ++i) {
int index = i + start; // Stencil index
// Get thread-local pointers
int const * threadIndices = indices + offsets[index];
float const * threadWeights = weights + offsets[index];
float const * threadWeightsDu = duWeights + offsets[index];
float const * threadWeightsDv = dvWeights + offsets[index];
float const * threadWeightsDuu = duuWeights + offsets[index];
float const * threadWeightsDuv = duvWeights + offsets[index];
float const * threadWeightsDvv = dvvWeights + offsets[index];
int threadId = omp_get_thread_num();
float * threadResult = result + threadId*srcDesc.length;
float * threadResultDu = resultDu + threadId*srcDesc.length;
float * threadResultDv = resultDv + threadId*srcDesc.length;
float * threadResultDuu = resultDuu + threadId*srcDesc.length;
float * threadResultDuv = resultDuv + threadId*srcDesc.length;
float * threadResultDvv = resultDvv + threadId*srcDesc.length;
clear(threadResult, dstDesc);
clear(threadResultDu, dstDuDesc);
clear(threadResultDv, dstDvDesc);
clear(threadResultDuu, dstDuuDesc);
clear(threadResultDuv, dstDuvDesc);
clear(threadResultDvv, dstDvvDesc);
for (int j=0; j<(int)sizes[index]; ++j) {
addWithWeight(threadResult, src,
threadIndices[j], threadWeights[j], srcDesc);
addWithWeight(threadResultDu, src,
threadIndices[j], threadWeightsDu[j], srcDesc);
addWithWeight(threadResultDv, src,
threadIndices[j], threadWeightsDv[j], srcDesc);
addWithWeight(threadResultDuu, src,
threadIndices[j], threadWeightsDuu[j], srcDesc);
addWithWeight(threadResultDuv, src,
threadIndices[j], threadWeightsDuv[j], srcDesc);
addWithWeight(threadResultDvv, src,
threadIndices[j], threadWeightsDvv[j], srcDesc);
}
copy(dst, i, threadResult, dstDesc);
copy(dstDu, i, threadResultDu, dstDuDesc);
copy(dstDv, i, threadResultDv, dstDvDesc);
copy(dstDuu, i, threadResultDuu, dstDuuDesc);
copy(dstDuv, i, threadResultDuv, dstDuvDesc);
copy(dstDvv, i, threadResultDvv, dstDvvDesc);
}
}
} // end namespace Osd
} // end namespace OPENSUBDIV_VERSION

19
opensubdiv/osd/ompKernel.h
View File

@ -56,6 +56,25 @@ OmpEvalStencils(float const * src, BufferDescriptor const &srcDesc,
float const * dvWeights,
int start, int end);
void
OmpEvalStencils(float const * src, BufferDescriptor const &srcDesc,
float * dst, BufferDescriptor const &dstDesc,
float * dstDu, BufferDescriptor const &dstDuDesc,
float * dstDv, BufferDescriptor const &dstDvDesc,
float * dstDuu, BufferDescriptor const &dstDuuDesc,
float * dstDuv, BufferDescriptor const &dstDuvDesc,
float * dstDvv, BufferDescriptor const &dstDvvDesc,
int const * sizes,
int const * offsets,
int const * indices,
float const * weights,
float const * duWeights,
float const * dvWeights,
float const * duuWeights,
float const * duvWeights,
float const * dvvWeights,
int start, int end);
} // end namespace Osd
} // end namespace OPENSUBDIV_VERSION

80
opensubdiv/osd/tbbEvaluator.cpp
View File

@ -75,8 +75,55 @@ TbbEvaluator::EvalStencils(
dst, dstDesc,
du, duDesc,
dv, dvDesc,
NULL, BufferDescriptor(),
NULL, BufferDescriptor(),
NULL, BufferDescriptor(),
sizes, offsets, indices,
weights, duWeights, dvWeights, NULL, NULL, NULL,
start, end);
return true;
}
/* static */
bool
TbbEvaluator::EvalStencils(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
const int * sizes,
const int * offsets,
const int * indices,
const float * weights,
const float * duWeights,
const float * dvWeights,
const float * duuWeights,
const float * duvWeights,
const float * dvvWeights,
int start, int end) {
if (end <= start) return true;
if (srcDesc.length != dstDesc.length) return false;
if (srcDesc.length != duDesc.length) return false;
if (srcDesc.length != dvDesc.length) return false;
if (srcDesc.length != duuDesc.length) return false;
if (srcDesc.length != duvDesc.length) return false;
if (srcDesc.length != dvvDesc.length) return false;
TbbEvalStencils(src, srcDesc,
dst, dstDesc,
du, duDesc,
dv, dvDesc,
duu, duuDesc,
duv, duvDesc,
dvv, dvvDesc,
sizes, offsets, indices,
weights, duWeights, dvWeights,
duuWeights, duvWeights, dvvWeights,
start, end);
return true;
@ -96,6 +143,9 @@ TbbEvaluator::EvalPatches(
if (srcDesc.length != dstDesc.length) return false;
TbbEvalPatches(src, srcDesc, dst, dstDesc,
NULL, BufferDescriptor(),
NULL, BufferDescriptor(),
NULL, BufferDescriptor(),
NULL, BufferDescriptor(),
NULL, BufferDescriptor(),
numPatchCoords, patchCoords,
@ -121,6 +171,36 @@ TbbEvaluator::EvalPatches(
TbbEvalPatches(src, srcDesc, dst, dstDesc,
du, duDesc, dv, dvDesc,
NULL, BufferDescriptor(),
NULL, BufferDescriptor(),
NULL, BufferDescriptor(),
numPatchCoords, patchCoords,
patchArrayBuffer, patchIndexBuffer, patchParamBuffer);
return true;
}
/* static */
bool
TbbEvaluator::EvalPatches(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
int numPatchCoords,
const PatchCoord *patchCoords,
const PatchArray *patchArrayBuffer,
const int *patchIndexBuffer,
const PatchParam *patchParamBuffer) {
if (srcDesc.length != dstDesc.length) return false;
TbbEvalPatches(src, srcDesc, dst, dstDesc,
du, duDesc, dv, dvDesc,
duu, duuDesc, duv, duvDesc, dvv, dvvDesc,
numPatchCoords, patchCoords,
patchArrayBuffer, patchIndexBuffer, patchParamBuffer);

744
opensubdiv/osd/tbbEvaluator.h
View File

@ -26,9 +26,8 @@
#define OPENSUBDIV3_OSD_TBB_EVALUATOR_H
#include "../version.h"
#include "../osd/types.h"
#include "../osd/bufferDescriptor.h"
#include "../far/patchTable.h"
#include "../osd/types.h"
#include <cstddef>
@ -61,7 +60,7 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param stencilTable stencil table to be applied.
/// @param stencilTable Far::StencilTable or equivalent
///
/// @param instance not used in the tbb kernel
/// (declared as a typed pointer to prevent
@ -108,7 +107,6 @@ public:
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
/// to apply for the range [start, end)
///
/// @param offsets pointer to the offsets buffer of the stencil table
///
@ -146,19 +144,19 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output U-derivative buffer
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the output buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output V-derivative buffer
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the output buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param stencilTable stencil table to be applied.
/// @param stencilTable Far::StencilTable or equivalent
///
/// @param instance not used in the tbb kernel
/// (declared as a typed pointer to prevent
@ -207,18 +205,17 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output s-derivatives pointer. An offset of
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the output buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output t-derivatives pointer. An offset of
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the output buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
/// to apply for the range [start, end)
///
/// @param offsets pointer to the offsets buffer of the stencil table
///
@ -226,9 +223,9 @@ public:
///
/// @param weights pointer to the weights buffer of the stencil table
///
/// @param duWeights pointer to the u-weights buffer of the stencil table
/// @param duWeights pointer to the du-weights buffer of the stencil table
///
/// @param dvWeights pointer to the v-weights buffer of the stencil table
/// @param dvWeights pointer to the dv-weights buffer of the stencil table
///
/// @param start start index of stencil table
///
@ -247,6 +244,177 @@ public:
const float * dvWeights,
int start, int end);
/// \brief Generic static eval stencils function with derivatives.
/// This function has a same signature as other device kernels
/// have so that it can be called in the same way from OsdMesh
/// template interface.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param stencilTable Far::StencilTable or equivalent
///
/// @param instance not used in the tbb kernel
/// (declared as a typed pointer to prevent
/// undesirable template resolution)
///
/// @param deviceContext not used in the tbb kernel
///
template <typename SRC_BUFFER, typename DST_BUFFER, typename STENCIL_TABLE>
static bool EvalStencils(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
STENCIL_TABLE const *stencilTable,
const TbbEvaluator *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalStencils(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
&stencilTable->GetSizes()[0],
&stencilTable->GetOffsets()[0],
&stencilTable->GetControlIndices()[0],
&stencilTable->GetWeights()[0],
&stencilTable->GetDuWeights()[0],
&stencilTable->GetDvWeights()[0],
&stencilTable->GetDuuWeights()[0],
&stencilTable->GetDuvWeights()[0],
&stencilTable->GetDvvWeights()[0],
/*start = */ 0,
/*end = */ stencilTable->GetNumStencils());
}
/// \brief Static eval stencils function with derivatives, which takes
/// raw CPU pointers for input and output.
///
/// @param src Input primvar pointer. An offset of srcDesc
/// will be applied internally (i.e. the pointer
/// should not include the offset)
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dst Output primvar pointer. An offset of dstDesc
/// will be applied internally.
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duu Output pointer 2nd derivative wrt u. An offset of
/// duuDesc will be applied internally.
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duv Output pointer 2nd derivative wrt u and v. An offset of
/// duvDesc will be applied internally.
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvv Output pointer 2nd derivative wrt v. An offset of
/// dvvDesc will be applied internally.
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param sizes pointer to the sizes buffer of the stencil table
///
/// @param offsets pointer to the offsets buffer of the stencil table
///
/// @param indices pointer to the indices buffer of the stencil table
///
/// @param weights pointer to the weights buffer of the stencil table
///
/// @param duWeights pointer to the du-weights buffer of the stencil table
///
/// @param dvWeights pointer to the dv-weights buffer of the stencil table
///
/// @param duuWeights pointer to the duu-weights buffer of the stencil table
///
/// @param duvWeights pointer to the duv-weights buffer of the stencil table
///
/// @param dvvWeights pointer to the dvv-weights buffer of the stencil table
///
/// @param start start index of stencil table
///
/// @param end end index of stencil table
///
static bool EvalStencils(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
const int * sizes,
const int * offsets,
const int * indices,
const float * weights,
const float * duWeights,
const float * dvWeights,
const float * duuWeights,
const float * duvWeights,
const float * dvvWeights,
int start, int end);
/// ----------------------------------------------------------------------
///
/// Limit evaluations with PatchTable
@ -273,7 +441,9 @@ public:
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable Far::PatchTable
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the cpu evaluator
///
@ -293,10 +463,8 @@ public:
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(),
srcDesc,
dstBuffer->BindCpuBuffer(),
dstDesc,
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetPatchArrayBuffer(),
@ -320,13 +488,13 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output s-derivatives buffer
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output t-derivatives buffer
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
@ -336,7 +504,9 @@ public:
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable Far::PatchTable
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the cpu evaluator
///
@ -358,8 +528,12 @@ public:
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(
srcBuffer->BindCpuBuffer(), srcDesc,
// XXX: PatchCoords is somewhat abusing vertex primvar buffer interop.
// ideally all buffer classes should have templated by datatype
// so that downcast isn't needed there.
// (e.g. Osd::CpuBuffer<PatchCoord> )
//
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
@ -370,6 +544,102 @@ public:
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function with derivatives. This function has
/// a same signature as other device kernels have so that it can be
/// called in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the cpu evaluator
///
/// @param deviceContext not used in the cpu evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatches(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
TbbEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
// XXX: PatchCoords is somewhat abusing vertex primvar buffer interop.
// ideally all buffer classes should have templated by datatype
// so that downcast isn't needed there.
// (e.g. Osd::CpuBuffer<PatchCoord> )
//
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetPatchArrayBuffer(),
patchTable->GetPatchIndexBuffer(),
patchTable->GetPatchParamBuffer());
}
/// \brief Static limit eval function. It takes an array of PatchCoord
/// and evaluate limit values on given PatchTable.
///
@ -420,15 +690,15 @@ public:
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output s-derivatives pointer. An offset of
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the du buffer
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output t-derivatives pointer. An offset of
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dv buffer
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
@ -449,10 +719,76 @@ public:
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
int numPatchCoords,
const PatchCoord *patchCoords,
const PatchArray *patchArrays,
PatchCoord const *patchCoords,
PatchArray const *patchArrays,
const int *patchIndexBuffer,
const PatchParam *patchParamBuffer);
PatchParam const *patchParamBuffer);
/// \brief Static limit eval function. It takes an array of PatchCoord
/// and evaluate limit values on given PatchTable.
///
/// @param src Input primvar pointer. An offset of srcDesc
/// will be applied internally (i.e. the pointer
/// should not include the offset)
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dst Output primvar pointer. An offset of dstDesc
/// will be applied internally.
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param du Output pointer derivative wrt u. An offset of
/// duDesc will be applied internally.
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dv Output pointer derivative wrt v. An offset of
/// dvDesc will be applied internally.
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duu Output pointer 2nd derivative wrt u. An offset of
/// duuDesc will be applied internally.
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duv Output pointer 2nd derivative wrt u and v. An offset of
/// duvDesc will be applied internally.
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvv Output pointer 2nd derivative wrt v. An offset of
/// dvvDesc will be applied internally.
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchArrays an array of Osd::PatchArray struct
/// indexed by PatchCoord::arrayIndex
///
/// @param patchIndexBuffer an array of patch indices
/// indexed by PatchCoord::vertIndex
///
/// @param patchParamBuffer an array of Osd::PatchParam struct
/// indexed by PatchCoord::patchIndex
///
static bool EvalPatches(
const float *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *du, BufferDescriptor const &duDesc,
float *dv, BufferDescriptor const &dvDesc,
float *duu, BufferDescriptor const &duuDesc,
float *duv, BufferDescriptor const &duvDesc,
float *dvv, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PatchCoord const *patchCoords,
PatchArray const *patchArrays,
const int *patchIndexBuffer,
PatchParam const *patchParamBuffer);
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
@ -474,7 +810,9 @@ public:
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable Far::PatchTable
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the cpu evaluator
///
@ -494,10 +832,166 @@ public:
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(),
srcDesc,
dstBuffer->BindCpuBuffer(),
dstDesc,
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetVaryingPatchArrayBuffer(),
patchTable->GetVaryingPatchIndexBuffer(),
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the cpu evaluator
///
/// @param deviceContext not used in the cpu evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
TbbEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetVaryingPatchArrayBuffer(),
patchTable->GetVaryingPatchIndexBuffer(),
patchTable->GetPatchParamBuffer());
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param instance not used in the cpu evaluator
///
/// @param deviceContext not used in the cpu evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
TbbEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetVaryingPatchArrayBuffer(),
@ -525,7 +1019,9 @@ public:
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable Far::PatchTable
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param fvarChannel face-varying channel
///
@ -548,10 +1044,172 @@ public:
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(),
srcDesc,
dstBuffer->BindCpuBuffer(),
dstDesc,
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetFVarPatchArrayBuffer(fvarChannel),
patchTable->GetFVarPatchIndexBuffer(fvarChannel),
patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param fvarChannel face-varying channel
///
/// @param instance not used in the cpu evaluator
///
/// @param deviceContext not used in the cpu evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesFaceVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
int fvarChannel,
TbbEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetFVarPatchArrayBuffer(fvarChannel),
patchTable->GetFVarPatchIndexBuffer(fvarChannel),
patchTable->GetFVarPatchParamBuffer(fvarChannel));
}
/// \brief Generic limit eval function. This function has a same
/// signature as other device kernels have so that it can be called
/// in the same way.
///
/// @param srcBuffer Input primvar buffer.
/// must have BindCpuBuffer() method returning a
/// const float pointer for read
///
/// @param srcDesc vertex buffer descriptor for the input buffer
///
/// @param dstBuffer Output primvar buffer
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dstDesc vertex buffer descriptor for the output buffer
///
/// @param duBuffer Output buffer derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duDesc vertex buffer descriptor for the duBuffer
///
/// @param dvBuffer Output buffer derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvDesc vertex buffer descriptor for the dvBuffer
///
/// @param duuBuffer Output buffer 2nd derivative wrt u
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duuDesc vertex buffer descriptor for the duuBuffer
///
/// @param duvBuffer Output buffer 2nd derivative wrt u and v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param duvDesc vertex buffer descriptor for the duvBuffer
///
/// @param dvvBuffer Output buffer 2nd derivative wrt v
/// must have BindCpuBuffer() method returning a
/// float pointer for write
///
/// @param dvvDesc vertex buffer descriptor for the dvvBuffer
///
/// @param numPatchCoords number of patchCoords.
///
/// @param patchCoords array of locations to be evaluated.
///
/// @param patchTable CpuPatchTable or equivalent
/// XXX: currently Far::PatchTable can't be used
/// due to interface mismatch
///
/// @param fvarChannel face-varying channel
///
/// @param instance not used in the cpu evaluator
///
/// @param deviceContext not used in the cpu evaluator
///
template <typename SRC_BUFFER, typename DST_BUFFER,
typename PATCHCOORD_BUFFER, typename PATCH_TABLE>
static bool EvalPatchesFaceVarying(
SRC_BUFFER *srcBuffer, BufferDescriptor const &srcDesc,
DST_BUFFER *dstBuffer, BufferDescriptor const &dstDesc,
DST_BUFFER *duBuffer, BufferDescriptor const &duDesc,
DST_BUFFER *dvBuffer, BufferDescriptor const &dvDesc,
DST_BUFFER *duuBuffer, BufferDescriptor const &duuDesc,
DST_BUFFER *duvBuffer, BufferDescriptor const &duvDesc,
DST_BUFFER *dvvBuffer, BufferDescriptor const &dvvDesc,
int numPatchCoords,
PATCHCOORD_BUFFER *patchCoords,
PATCH_TABLE *patchTable,
int fvarChannel,
TbbEvaluator const *instance = NULL,
void * deviceContext = NULL) {
(void)instance; // unused
(void)deviceContext; // unused
return EvalPatches(srcBuffer->BindCpuBuffer(), srcDesc,
dstBuffer->BindCpuBuffer(), dstDesc,
duBuffer->BindCpuBuffer(), duDesc,
dvBuffer->BindCpuBuffer(), dvDesc,
duuBuffer->BindCpuBuffer(), duuDesc,
duvBuffer->BindCpuBuffer(), duvDesc,
dvvBuffer->BindCpuBuffer(), dvvDesc,
numPatchCoords,
(const PatchCoord*)patchCoords->BindCpuBuffer(),
patchTable->GetFVarPatchArrayBuffer(fvarChannel),

246
opensubdiv/osd/tbbKernel.cpp
View File

@ -219,6 +219,78 @@ TbbEvalStencils(float const * src, BufferDescriptor const &srcDesc,
tbb::blocked_range<int> range(start, end, grain_size);
tbb::parallel_for(range, kernel);
}
}
void
TbbEvalStencils(float const * src, BufferDescriptor const &srcDesc,
float * dst, BufferDescriptor const &dstDesc,
float * du, BufferDescriptor const &duDesc,
float * dv, BufferDescriptor const &dvDesc,
float * duu, BufferDescriptor const &duuDesc,
float * duv, BufferDescriptor const &duvDesc,
float * dvv, BufferDescriptor const &dvvDesc,
int const * sizes,
int const * offsets,
int const * indices,
float const * weights,
float const * duWeights,
float const * dvWeights,
float const * duuWeights,
float const * duvWeights,
float const * dvvWeights,
int start, int end) {
if (src) src += srcDesc.offset;
if (dst) dst += dstDesc.offset;
if (du) du += duDesc.offset;
if (dv) dv += dvDesc.offset;
if (duu) duu += duuDesc.offset;
if (duv) duv += duvDesc.offset;
if (dvv) dvv += dvvDesc.offset;
// PERFORMANCE: need to combine 3 launches together
if (dst) {
TBBStencilKernel kernel(src, srcDesc, dst, dstDesc,
sizes, offsets, indices, weights);
tbb::blocked_range<int> range(start, end, grain_size);
tbb::parallel_for(range, kernel);
}
if (du) {
TBBStencilKernel kernel(src, srcDesc, du, duDesc,
sizes, offsets, indices, duWeights);
tbb::blocked_range<int> range(start, end, grain_size);
tbb::parallel_for(range, kernel);
}
if (dv) {
TBBStencilKernel kernel(src, srcDesc, dv, dvDesc,
sizes, offsets, indices, dvWeights);
tbb::blocked_range<int> range(start, end, grain_size);
tbb::parallel_for(range, kernel);
}
if (duu) {
TBBStencilKernel kernel(src, srcDesc, duu, duuDesc,
sizes, offsets, indices, duuWeights);
tbb::blocked_range<int> range(start, end, grain_size);
tbb::parallel_for(range, kernel);
}
if (duv) {
TBBStencilKernel kernel(src, srcDesc, duv, duvDesc,
sizes, offsets, indices, duvWeights);
tbb::blocked_range<int> range(start, end, grain_size);
tbb::parallel_for(range, kernel);
}
if (dvv) {
TBBStencilKernel kernel(src, srcDesc, dvv, dvvDesc,
sizes, offsets, indices, dvvWeights);
tbb::blocked_range<int> range(start, end, grain_size);
tbb::parallel_for(range, kernel);
}
}
// ---------------------------------------------------------------------------
@ -257,10 +329,16 @@ class TbbEvalPatchesKernel {
BufferDescriptor _dstDesc;
BufferDescriptor _dstDuDesc;
BufferDescriptor _dstDvDesc;
BufferDescriptor _dstDuuDesc;
BufferDescriptor _dstDuvDesc;
BufferDescriptor _dstDvvDesc;
float const * _src;
float * _dst;
float * _dstDu;
float * _dstDv;
float * _dstDuu;
float * _dstDuv;
float * _dstDvv;
int _numPatchCoords;
const PatchCoord *_patchCoords;
const PatchArray *_patchArrayBuffer;
@ -272,6 +350,9 @@ public:
float *dst, BufferDescriptor dstDesc,
float *dstDu, BufferDescriptor dstDuDesc,
float *dstDv, BufferDescriptor dstDvDesc,
float *dstDuu, BufferDescriptor dstDuuDesc,
float *dstDuv, BufferDescriptor dstDuvDesc,
float *dstDvv, BufferDescriptor dstDvvDesc,
int numPatchCoords,
const PatchCoord *patchCoords,
const PatchArray *patchArrayBuffer,
@ -279,7 +360,10 @@ public:
const PatchParam *patchParamBuffer) :
_srcDesc(srcDesc), _dstDesc(dstDesc),
_dstDuDesc(dstDuDesc), _dstDvDesc(dstDvDesc),
_src(src), _dst(dst), _dstDu(dstDu), _dstDv(dstDv),
_dstDuuDesc(dstDuuDesc), _dstDuvDesc(dstDuvDesc), _dstDvvDesc(dstDvvDesc),
_src(src), _dst(dst),
_dstDu(dstDu), _dstDv(dstDv),
_dstDuu(dstDuu), _dstDuv(dstDuv), _dstDvv(dstDvv),
_numPatchCoords(numPatchCoords),
_patchCoords(patchCoords),
_patchArrayBuffer(patchArrayBuffer),
@ -290,13 +374,15 @@ public:
void operator() (tbb::blocked_range<int> const &r) const {
if (_dstDu == NULL && _dstDv == NULL) {
compute(r);
} else if (_dstDuu == NULL && _dstDuv == NULL && _dstDvv == NULL) {
computeWith1stDerivative(r);
} else {
computeWithDerivative(r);
computeWith2ndDerivative(r);
}
}
void compute(tbb::blocked_range<int> const &r) const {
float wP[20], wDs[20], wDt[20];
float wP[20], wDu[20], wDv[20];
BufferAdapter<const float> srcT(_src + _srcDesc.offset,
_srcDesc.length,
_srcDesc.stride);
@ -305,12 +391,6 @@ public:
_dstDesc.length,
_dstDesc.stride);
BufferAdapter<float> dstDuT(_dstDu,
_dstDuDesc.length,
_dstDuDesc.stride);
BufferAdapter<float> dstDvT(_dstDv,
_dstDvDesc.length,
_dstDvDesc.stride);
for (int i = r.begin(); i < r.end(); ++i) {
PatchCoord const &coord = _patchCoords[i];
@ -325,15 +405,18 @@ public:
int numControlVertices = 0;
if (patchType == Far::PatchDescriptor::REGULAR) {
Far::internal::GetBSplineWeights(param,
coord.s, coord.t, wP, wDs, wDt);
coord.s, coord.t, wP,
wDu, wDv);
numControlVertices = 16;
} else if (patchType == Far::PatchDescriptor::GREGORY_BASIS) {
Far::internal::GetGregoryWeights(param,
coord.s, coord.t, wP, wDs, wDt);
coord.s, coord.t, wP,
wDu, wDv);
numControlVertices = 20;
} else if (patchType == Far::PatchDescriptor::QUADS) {
Far::internal::GetBilinearWeights(param,
coord.s, coord.t, wP, wDs, wDt);
coord.s, coord.t, wP,
wDu, wDv);
numControlVertices = 4;
} else {
assert(0);
@ -353,8 +436,8 @@ public:
}
}
void computeWithDerivative(tbb::blocked_range<int> const &r) const {
float wP[20], wDs[20], wDt[20];
void computeWith1stDerivative(tbb::blocked_range<int> const &r) const {
float wP[20], wDu[20], wDv[20];
BufferAdapter<const float> srcT(_src + _srcDesc.offset,
_srcDesc.length,
_srcDesc.stride);
@ -384,15 +467,18 @@ public:
int numControlVertices = 0;
if (patchType == Far::PatchDescriptor::REGULAR) {
Far::internal::GetBSplineWeights(param,
coord.s, coord.t, wP, wDs, wDt);
coord.s, coord.t, wP,
wDu, wDv);
numControlVertices = 16;
} else if (patchType == Far::PatchDescriptor::GREGORY_BASIS) {
Far::internal::GetGregoryWeights(param,
coord.s, coord.t, wP, wDs, wDt);
coord.s, coord.t, wP,
wDu, wDv);
numControlVertices = 20;
} else if (patchType == Far::PatchDescriptor::QUADS) {
Far::internal::GetBilinearWeights(param,
coord.s, coord.t, wP, wDs, wDt);
coord.s, coord.t,
wP, wDu, wDv);
numControlVertices = 4;
} else {
assert(0);
@ -409,14 +495,103 @@ public:
dstDvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
dstDuT.AddWithWeight(srcT[cvs[j]], wDs[j]);
dstDvT.AddWithWeight(srcT[cvs[j]], wDt[j]);
dstDuT.AddWithWeight(srcT[cvs[j]], wDu[j]);
dstDvT.AddWithWeight(srcT[cvs[j]], wDv[j]);
}
++dstT;
++dstDuT;
++dstDvT;
}
}
void computeWith2ndDerivative(tbb::blocked_range<int> const &r) const {
float wP[20], wDu[20], wDv[20], wDuu[20], wDuv[20], wDvv[20];
BufferAdapter<const float> srcT(_src + _srcDesc.offset,
_srcDesc.length,
_srcDesc.stride);
BufferAdapter<float> dstT(_dst + _dstDesc.offset
+ r.begin() * _dstDesc.stride,
_dstDesc.length,
_dstDesc.stride);
BufferAdapter<float> dstDuT(_dstDu + _dstDuDesc.offset
+ r.begin() * _dstDuDesc.stride,
_dstDuDesc.length,
_dstDuDesc.stride);
BufferAdapter<float> dstDvT(_dstDv + _dstDvDesc.offset
+ r.begin() * _dstDvDesc.stride,
_dstDvDesc.length,
_dstDvDesc.stride);
BufferAdapter<float> dstDuuT(_dstDuu + _dstDuuDesc.offset
+ r.begin() * _dstDuuDesc.stride,
_dstDuuDesc.length,
_dstDuuDesc.stride);
BufferAdapter<float> dstDuvT(_dstDuv + _dstDuvDesc.offset
+ r.begin() * _dstDuvDesc.stride,
_dstDuvDesc.length,
_dstDuvDesc.stride);
BufferAdapter<float> dstDvvT(_dstDvv + _dstDvvDesc.offset
+ r.begin() * _dstDvvDesc.stride,
_dstDvvDesc.length,
_dstDvvDesc.stride);
for (int i = r.begin(); i < r.end(); ++i) {
PatchCoord const &coord = _patchCoords[i];
PatchArray const &array = _patchArrayBuffer[coord.handle.arrayIndex];
Far::PatchParam const & param =
_patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
? Far::PatchDescriptor::REGULAR
: array.GetPatchType();
int numControlVertices = 0;
if (patchType == Far::PatchDescriptor::REGULAR) {
Far::internal::GetBSplineWeights(param,
coord.s, coord.t, wP,
wDu, wDv, wDuu, wDuv, wDvv);
numControlVertices = 16;
} else if (patchType == Far::PatchDescriptor::GREGORY_BASIS) {
Far::internal::GetGregoryWeights(param,
coord.s, coord.t, wP,
wDu, wDv, wDuu, wDuv, wDvv);
numControlVertices = 20;
} else if (patchType == Far::PatchDescriptor::QUADS) {
Far::internal::GetBilinearWeights(param,
coord.s, coord.t, wP,
wDu, wDv, wDuu, wDuv, wDvv);
numControlVertices = 4;
} else {
assert(0);
}
int indexStride = Far::PatchDescriptor(array.GetPatchType()).GetNumControlVertices();
int indexBase = array.GetIndexBase() + indexStride *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
const int *cvs = &_patchIndexBuffer[indexBase];
dstT.Clear();
dstDuT.Clear();
dstDvT.Clear();
dstDuuT.Clear();
dstDuvT.Clear();
dstDvvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
dstDuT.AddWithWeight(srcT[cvs[j]], wDu[j]);
dstDvT.AddWithWeight(srcT[cvs[j]], wDv[j]);
dstDuuT.AddWithWeight(srcT[cvs[j]], wDuu[j]);
dstDuvT.AddWithWeight(srcT[cvs[j]], wDuv[j]);
dstDvvT.AddWithWeight(srcT[cvs[j]], wDvv[j]);
}
++dstT;
++dstDuT;
++dstDvT;
++dstDuuT;
++dstDuvT;
++dstDvvT;
}
}
};
@ -433,6 +608,39 @@ TbbEvalPatches(float const *src, BufferDescriptor const &srcDesc,
TbbEvalPatchesKernel kernel(src, srcDesc, dst, dstDesc,
dstDu, dstDuDesc, dstDv, dstDvDesc,
NULL, BufferDescriptor(),
NULL, BufferDescriptor(),
NULL, BufferDescriptor(),
numPatchCoords, patchCoords,
patchArrayBuffer,
patchIndexBuffer,
patchParamBuffer);
tbb::blocked_range<int> range(0, numPatchCoords, grain_size);
tbb::parallel_for(range, kernel);
}
void
TbbEvalPatches(float const *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *dstDu, BufferDescriptor const &dstDuDesc,
float *dstDv, BufferDescriptor const &dstDvDesc,
float *dstDuu, BufferDescriptor const &dstDuuDesc,
float *dstDuv, BufferDescriptor const &dstDuvDesc,
float *dstDvv, BufferDescriptor const &dstDvvDesc,
int numPatchCoords,
const PatchCoord *patchCoords,
const PatchArray *patchArrayBuffer,
const int *patchIndexBuffer,
const PatchParam *patchParamBuffer) {
TbbEvalPatchesKernel kernel(src, srcDesc, dst, dstDesc,
dstDu, dstDuDesc, dstDv, dstDvDesc,
dstDuu, dstDuuDesc,
dstDuv, dstDuvDesc,
dstDvv, dstDvvDesc,
numPatchCoords, patchCoords,
patchArrayBuffer,
patchIndexBuffer,

33
opensubdiv/osd/tbbKernel.h
View File

@ -61,6 +61,25 @@ TbbEvalStencils(float const * src, BufferDescriptor const &srcDesc,
float const * dvWeights,
int start, int end);
void
TbbEvalStencils(float const * src, BufferDescriptor const &srcDesc,
float * dst, BufferDescriptor const &dstDesc,
float * dstDu, BufferDescriptor const &dstDuDesc,
float * dstDv, BufferDescriptor const &dstDvDesc,
float * dstDuu, BufferDescriptor const &dstDuuDesc,
float * dstDuv, BufferDescriptor const &dstDuvDesc,
float * dstDvv, BufferDescriptor const &dstDvvDesc,
int const * sizes,
int const * offsets,
int const * indices,
float const * weights,
float const * duWeights,
float const * dvWeights,
float const * duuWeights,
float const * duvWeights,
float const * dvvWeights,
int start, int end);
void
TbbEvalPatches(float const *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
@ -72,6 +91,20 @@ TbbEvalPatches(float const *src, BufferDescriptor const &srcDesc,
const int *patchIndexBuffer,
const PatchParam *patchParamBuffer);
void
TbbEvalPatches(float const *src, BufferDescriptor const &srcDesc,
float *dst, BufferDescriptor const &dstDesc,
float *dstDu, BufferDescriptor const &dstDuDesc,
float *dstDv, BufferDescriptor const &dstDvDesc,
float *dstDuu, BufferDescriptor const &dstDuuDesc,
float *dstDuv, BufferDescriptor const &dstDuvDesc,
float *dstDvv, BufferDescriptor const &dstDvvDesc,
int numPatchCoords,
const PatchCoord *patchCoords,
const PatchArray *patchArrayBuffer,
const int *patchIndexBuffer,
const PatchParam *patchParamBuffer);
} // end namespace Osd
} // end namespace OPENSUBDIV_VERSION