AuroraMiddleware/OpenSubdiv
1
0
Fork 0
mirror of https://github.com/PixarAnimationStudios/OpenSubdiv synced 2024-11-09 13:50:05 +00:00
Code Issues Releases Wiki Activity

Improved patch basis eval for Osd to match Far

Browse Source 
This updates the patch basis evaluation functions in Osd
to match recent changes to far/patchBasis.

This also exposes a common facility for dealing with PatchCoord,
PatchArray, and PatchParam. These are exposed as global functions
operating on struct data, since C++ style class methods are not
supported by all of the Osd shader and kernel execution envirionments.

Changes:
    - Merged far/patchBasis.cpp to osd/patchBasisCommon{,Types,Eval}.h
    - Exposed PatchCoord, PatchArray, and PatchParam to Osd kernels
    - exposed OsdEvaluatePatchBasis and OsdEvaluatePatchBasisNormalized
    - Updated CPU, TBB, Omp, CUDA, OpenCL, GLSL, HLSL, and Metal evaluators
    - Updated glFVarViewer
This commit is contained in:
David G Yu 2018-10-30 00:37:25 -07:00
parent 63c189b5e6
commit ff34c514f5
22 changed files with 2197 additions and 1036 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

34
examples/glFVarViewer/glFVarViewer.cpp
View File

@ -117,7 +117,9 @@ int g_tessLevelMin = 1;
GLuint g_transformUB = 0,
g_transformBinding = 0,
g_tessellationUB = 0,
g_tessellationBinding = 0;
g_tessellationBinding = 0,
g_fvarArrayDataUB = 0,
g_fvarArrayDataBinding = 0;
struct Transform {
float ModelViewMatrix[16];
@ -685,6 +687,11 @@ public:
if (uboIndex != GL_INVALID_INDEX)
glUniformBlockBinding(program, uboIndex, g_tessellationBinding);
g_fvarArrayDataBinding = 2;
uboIndex = glGetUniformBlockIndex(program, "FVarArrayData");
if (uboIndex != GL_INVALID_INDEX)
glUniformBlockBinding(program, uboIndex, g_fvarArrayDataBinding);
// assign texture locations
GLint loc;
glUseProgram(program);
@ -708,6 +715,9 @@ ShaderCache g_shaderCache;
//------------------------------------------------------------------------------
static void
updateUniformBlocks() {
using namespace OpenSubdiv;
if (!g_transformUB) {
glGenBuffers(1, &g_transformUB);
glBindBuffer(GL_UNIFORM_BUFFER, g_transformUB);
@ -740,6 +750,28 @@ updateUniformBlocks() {
glBindBuffer(GL_UNIFORM_BUFFER, 0);
glBindBufferBase(GL_UNIFORM_BUFFER, g_tessellationBinding, g_tessellationUB);
// Update and bind fvar patch array state
Osd::PatchArrayVector const &fvarPatchArrays =
g_mesh->GetPatchTable()->GetFVarPatchArrays();
if (! fvarPatchArrays.empty()) {
// bind patch arrays UBO (std140 struct size padded to vec4 alignment)
int patchArraySize =
sizeof(GLint) * ((sizeof(Osd::PatchArray)/sizeof(GLint) + 3) & ~3);
if (!g_fvarArrayDataUB) {
glGenBuffers(1, &g_fvarArrayDataUB);
}
glBindBuffer(GL_UNIFORM_BUFFER, g_fvarArrayDataUB);
glBufferData(GL_UNIFORM_BUFFER,
fvarPatchArrays.size()*patchArraySize, NULL, GL_STATIC_DRAW);
for (int i=0; i<(int)fvarPatchArrays.size(); ++i) {
glBufferSubData(GL_UNIFORM_BUFFER,
i*patchArraySize, sizeof(Osd::PatchArray), &fvarPatchArrays[i]);
}
glBindBufferBase(GL_UNIFORM_BUFFER,
g_fvarArrayDataBinding, g_fvarArrayDataUB);
}
}
static void

53
examples/glFVarViewer/shader.glsl
View File

@ -167,48 +167,32 @@ out block {
} outpt;
uniform isamplerBuffer OsdFVarParamBuffer;
layout(std140) uniform FVarArrayData {
OsdPatchArray fvarPatchArray[2];
};
vec2
interpolateFaceVarying(vec2 uv, int fvarOffset)
{
int patchIndex = OsdGetPatchIndex(gl_PrimitiveID);
#if defined(SHADING_FACEVARYING_SMOOTH_BSPLINE_BASIS)
float wP[16], wDs[16], wDt[16], wDss[16], wDst[16], wDtt[16];
int patchCVs = 16;
int patchStride = patchCVs;
OsdPatchArray array = fvarPatchArray[0];
ivec3 fvarPatchParam = texelFetch(OsdFVarParamBuffer, patchIndex).xyz;
int boundaryMask = OsdGetPatchBoundaryMask(fvarPatchParam);
OsdGetBSplinePatchWeights(uv.s, uv.t, 1.0f, boundaryMask, wP, wDs, wDt, wDss, wDst, wDtt);
OsdPatchParam param = OsdPatchParamInit(fvarPatchParam.x,
fvarPatchParam.y,
fvarPatchParam.z);
#elif defined(SHADING_FACEVARYING_SMOOTH_GREGORY_BASIS)
float wP[20], wDs[20], wDt[20], wDss[20], wDst[20], wDtt[20];
int patchCVs = 20;
int patchStride = patchCVs;
ivec3 fvarPatchParam = texelFetch(OsdFVarParamBuffer, patchIndex).xyz;
if (OsdGetPatchIsRegular(fvarPatchParam)) {
float wP16[16], wDs16[16], wDt16[16], wDss16[16], wDst16[16], wDtt16[16];
patchCVs = 16;
int boundaryMask = OsdGetPatchBoundaryMask(fvarPatchParam);
OsdGetBSplinePatchWeights(uv.s, uv.t, 1.0f, boundaryMask, wP16, wDs16, wDt16, wDss16, wDst16, wDtt16);
for (int i=0; i<patchCVs; ++i) {
wP[i] = wP16[i];
}
} else {
OsdGetGregoryPatchWeights(uv.s, uv.t, 1.0f, wP, wDs, wDt, wDss, wDst, wDtt);
}
int patchType = OsdPatchParamIsRegular(param) ? array.regDesc : array.desc;
#else
float wP[4], wDs[4], wDt[4], wDss[4], wDst[4], wDtt[4];
int patchCVs = 4;
int patchStride = patchCVs;
OsdGetBilinearPatchWeights(uv.s, uv.t, 1.0f, wP, wDs, wDt, wDss, wDst, wDtt);
#endif
float wP[20], wDu[20], wDv[20], wDuu[20], wDuv[20], wDvv[20];
int numPoints = OsdEvaluatePatchBasisNormalized(patchType, param,
uv.s, uv.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
int primOffset = patchIndex * patchStride;
int primOffset = patchIndex * array.stride;
vec2 result = vec2(0);
for (int i=0; i<patchCVs; ++i) {
for (int i=0; i<numPoints; ++i) {
int index = (primOffset+i)*OSD_FVAR_WIDTH + fvarOffset;
vec2 cv = vec2(texelFetch(OsdFVarDataBuffer, index).s,
texelFetch(OsdFVarDataBuffer, index + 1).s);
@ -228,8 +212,13 @@ void emit(int index, vec3 normal)
#endif
#ifdef LOOP // ----- scheme : LOOP
vec2 uv;
OSD_COMPUTE_FACE_VARYING_TRI_2(uv, /*fvarOffste=*/0, index);
vec2 trist[3] = vec2[](vec2(0,0), vec2(1,0), vec2(0,1));
#ifdef SHADING_FACEVARYING_UNIFORM_SUBDIVISION
vec2 st = trist[index];
#else
vec2 st = inpt[index].v.tessCoord;
#endif
vec2 uv = interpolateFaceVarying(st, /*fvarOffset*/0);
#else // ----- scheme : CATMARK / BILINEAR

2
opensubdiv/osd/CMakeLists.txt
View File

@ -53,7 +53,9 @@ set(PUBLIC_HEADER_FILES
)
list(APPEND KERNEL_FILES
patchBasisCommonTypes.h
patchBasisCommon.h
patchBasisCommonEval.h
)
set(DOXY_HEADER_FILES ${PUBLIC_HEADER_FILES})

14
opensubdiv/osd/clEvaluator.cpp
View File

@ -41,9 +41,15 @@ namespace Osd {
static const char *clSource =
#include "clKernel.gen.h"
;
static const char *patchBasisTypesSource =
#include "patchBasisCommonTypes.gen.h"
;
static const char *patchBasisSource =
#include "patchBasisCommon.gen.h"
;
static const char *patchBasisEvalSource =
#include "patchBasisCommonEval.gen.h"
;
// ----------------------------------------------------------------------------
@ -166,8 +172,12 @@ CLEvaluator::Compile(BufferDescriptor const &srcDesc,
<< "#define OSD_PATCH_BASIS_OPENCL\n";
std::string defineStr = defines.str();
const char *sources[] = { defineStr.c_str(), patchBasisSource, clSource };
_program = clCreateProgramWithSource(_clContext, 3, sources, 0, &errNum);
const char *sources[] = { defineStr.c_str(),
patchBasisTypesSource,
patchBasisSource,
patchBasisEvalSource,
clSource };
_program = clCreateProgramWithSource(_clContext, 5, sources, 0, &errNum);
if (errNum != CL_SUCCESS) {
Far::Error(Far::FAR_RUNTIME_ERROR,
"clCreateProgramWithSource (%d)", errNum);

123
opensubdiv/osd/clKernel.cl
View File

@ -159,66 +159,6 @@ __kernel void computeStencilsDerivatives(
// ---------------------------------------------------------------------------
struct PatchArray {
int patchType;
int numPatches;
int indexBase; // an offset within the index buffer
int primitiveIdBase; // an offset within the patch param buffer
};
struct PatchCoord {
int arrayIndex;
int patchIndex;
int vertIndex;
float s;
float t;
};
struct PatchParam {
uint field0;
uint field1;
float sharpness;
};
static int getDepth(uint patchBits) {
return (patchBits & 0xf);
}
static float getParamFraction(uint patchBits) {
bool nonQuadRoot = (patchBits >> 4) & 0x1;
int depth = getDepth(patchBits);
if (nonQuadRoot) {
return 1.0f / (float)( 1 << (depth-1) );
} else {
return 1.0f / (float)( 1 << depth );
}
}
static void normalizePatchCoord(uint patchBits, float *uv) {
float frac = getParamFraction(patchBits);
int iu = (patchBits >> 22) & 0x3ff;
int iv = (patchBits >> 12) & 0x3ff;
// top left corner
float pu = (float)iu*frac;
float pv = (float)iv*frac;
// normalize u,v coordinates
uv[0] = (uv[0] - pu) / frac;
uv[1] = (uv[1] - pv) / frac;
}
static bool isRegular(uint patchBits) {
return ((patchBits >> 5) & 0x1) != 0;
}
static int getNumControlVertices(int patchType) {
return (patchType == 3) ? 4 :
(patchType == 6) ? 16 :
(patchType == 9) ? 20 : 0;
}
__kernel void computePatches(__global float *src, int srcOffset,
__global float *dst, int dstOffset,
__global float *du, int duOffset, int duStride,
@ -226,10 +166,10 @@ __kernel void computePatches(__global float *src, int srcOffset,
__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 struct OsdPatchCoord *patchCoords,
__global struct OsdPatchArray *patchArrayBuffer,
__global int *patchIndexBuffer,
__global struct PatchParam *patchParamBuffer) {
__global struct OsdPatchParam *patchParamBuffer) {
int current = get_global_id(0);
if (src) src += srcOffset;
@ -240,41 +180,22 @@ __kernel void computePatches(__global float *src, int srcOffset,
if (duv) duv += duvOffset;
if (dvv) dvv += dvvOffset;
struct PatchCoord coord = patchCoords[current];
struct PatchArray array = patchArrayBuffer[coord.arrayIndex];
struct OsdPatchCoord coord = patchCoords[current];
struct OsdPatchArray array = patchArrayBuffer[coord.arrayIndex];
struct OsdPatchParam param = patchParamBuffer[coord.patchIndex];
uint patchBits = patchParamBuffer[coord.patchIndex].field1;
int patchType = isRegular(patchBits) ? 6 : array.patchType;
int patchType = OsdPatchParamIsRegular(param) ? array.regDesc : array.desc;
float uv[2] = {coord.s, coord.t};
normalizePatchCoord(patchBits, uv);
float dScale = (float)(1 << getDepth(patchBits));
int boundary = (patchBits >> 7) & 0x1f;
float wP[20], wDu[20], wDv[20], wDuu[20], wDuv[20], wDvv[20];
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
float wP[20], wDs[20], wDt[20], wDss[20], wDst[20], wDtt[20];
int numControlVertices = 0;
if (patchType == 3) {
OsdGetBilinearPatchWeights(uv[0], uv[1], dScale,
wP, wDs, wDt, wDss, wDst, wDtt);
numControlVertices = 4;
} else if (patchType == 6) {
OsdGetBSplinePatchWeights(uv[0], uv[1], dScale, boundary,
wP, wDs, wDt, wDss, wDst, wDtt);
numControlVertices = 16;
} else if (patchType == 9) {
OsdGetGregoryPatchWeights(uv[0], uv[1], dScale,
wP, wDs, wDt, wDss, wDst, wDtt);
numControlVertices = 20;
}
int indexStride = getNumControlVertices(array.patchType);
int indexBase = array.indexBase + indexStride *
int indexBase = array.indexBase + array.stride *
(coord.patchIndex - array.primitiveIdBase);
struct Vertex v;
clear(&v);
for (int i = 0; i < numControlVertices; ++i) {
for (int i = 0; i < nPoints; ++i) {
int index = patchIndexBuffer[indexBase + i];
addWithWeight(&v, src, index, wP[i]);
}
@ -283,45 +204,45 @@ __kernel void computePatches(__global float *src, int srcOffset,
if (du) {
struct Vertex vdu;
clear(&vdu);
for (int i = 0; i < numControlVertices; ++i) {
for (int i = 0; i < nPoints; ++i) {
int index = patchIndexBuffer[indexBase + i];
addWithWeight(&vdu, src, index, wDs[i]);
addWithWeight(&vdu, src, index, wDu[i]);
}
writeVertexStride(du, current, &vdu, duStride);
}
if (dv) {
struct Vertex vdv;
clear(&vdv);
for (int i = 0; i < numControlVertices; ++i) {
for (int i = 0; i < nPoints; ++i) {
int index = patchIndexBuffer[indexBase + i];
addWithWeight(&vdv, src, index, wDt[i]);
addWithWeight(&vdv, src, index, wDv[i]);
}
writeVertexStride(dv, current, &vdv, dvStride);
}
if (duu) {
struct Vertex vduu;
clear(&vduu);
for (int i = 0; i < numControlVertices; ++i) {
for (int i = 0; i < nPoints; ++i) {
int index = patchIndexBuffer[indexBase + i];
addWithWeight(&vduu, src, index, wDss[i]);
addWithWeight(&vduu, src, index, wDuu[i]);
}
writeVertexStride(duu, current, &vduu, duuStride);
}
if (duv) {
struct Vertex vduv;
clear(&vduv);
for (int i = 0; i < numControlVertices; ++i) {
for (int i = 0; i < nPoints; ++i) {
int index = patchIndexBuffer[indexBase + i];
addWithWeight(&vduv, src, index, wDst[i]);
addWithWeight(&vduv, src, index, wDuv[i]);
}
writeVertexStride(duv, current, &vduv, duvStride);
}
if (dvv) {
struct Vertex vdvv;
clear(&vdvv);
for (int i = 0; i < numControlVertices; ++i) {
for (int i = 0; i < nPoints; ++i) {
int index = patchIndexBuffer[indexBase + i];
addWithWeight(&vdvv, src, index, wDtt[i]);
addWithWeight(&vdvv, src, index, wDvv[i]);
}
writeVertexStride(dvv, current, &vdvv, dvvStride);
}

45
opensubdiv/osd/cpuEvaluator.cpp
View File

@ -24,7 +24,9 @@
#include "../osd/cpuEvaluator.h"
#include "../osd/cpuKernel.h"
#include "../far/patchBasis.h"
#include "../osd/patchBasisCommonTypes.h"
#include "../osd/patchBasisCommon.h"
#include "../osd/patchBasisCommonEval.h"
#include <cstdlib>
@ -177,20 +179,23 @@ CpuEvaluator::EvalPatches(const float *src, BufferDescriptor const &srcDesc,
BufferAdapter<const float> srcT(src, srcDesc.length, srcDesc.stride);
BufferAdapter<float> dstT(dst, dstDesc.length, dstDesc.stride);
float wP[20], wDs[20], wDt[20];
float wP[20];
for (int i = 0; i < numPatchCoords; ++i) {
PatchCoord const &coord = patchCoords[i];
PatchArray const &array = patchArrays[coord.handle.arrayIndex];
Far::PatchParam const & param =
Osd::PatchParam const & paramStruct =
patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
OsdPatchParam param = OsdPatchParamInit(
paramStruct.field0, paramStruct.field1, paramStruct.sharpness);
int patchType = OsdPatchParamIsRegular(param)
? array.GetPatchTypeRegular()
: array.GetPatchTypeIrregular();
int numControlVertices = Far::internal::EvaluatePatchBasis(patchType,
param, coord.s, coord.t, wP, wDs, wDt);
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, 0, 0, 0, 0, 0);
int indexBase = array.GetIndexBase() + array.GetStride() *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
@ -198,7 +203,7 @@ CpuEvaluator::EvalPatches(const float *src, BufferDescriptor const &srcDesc,
const int *cvs = &patchIndexBuffer[indexBase];
dstT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
}
++dstT;
@ -246,14 +251,17 @@ CpuEvaluator::EvalPatches(const float *src, BufferDescriptor const &srcDesc,
PatchCoord const &coord = patchCoords[i];
PatchArray const &array = patchArrays[coord.handle.arrayIndex];
Far::PatchParam const & param =
Osd::PatchParam const & paramStruct =
patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
OsdPatchParam param = OsdPatchParamInit(
paramStruct.field0, paramStruct.field1, paramStruct.sharpness);
int patchType = OsdPatchParamIsRegular(param)
? array.GetPatchTypeRegular()
: array.GetPatchTypeIrregular();
int numControlVertices = Far::internal::EvaluatePatchBasis(patchType,
param, coord.s, coord.t, wP, wDs, wDt);
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, wDs, wDt, 0, 0, 0);
int indexBase = array.GetIndexBase() + array.GetStride() *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
@ -263,7 +271,7 @@ CpuEvaluator::EvalPatches(const float *src, BufferDescriptor const &srcDesc,
dstT.Clear();
duT.Clear();
dvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
duT.AddWithWeight (srcT[cvs[j]], wDs[j]);
dvT.AddWithWeight (srcT[cvs[j]], wDt[j]);
@ -333,14 +341,17 @@ CpuEvaluator::EvalPatches(const float *src, BufferDescriptor const &srcDesc,
PatchCoord const &coord = patchCoords[i];
PatchArray const &array = patchArrays[coord.handle.arrayIndex];
Far::PatchParam const & param =
Osd::PatchParam const & paramStruct =
patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
OsdPatchParam param = OsdPatchParamInit(
paramStruct.field0, paramStruct.field1, paramStruct.sharpness);
int patchType = OsdPatchParamIsRegular(param)
? array.GetPatchTypeRegular()
: array.GetPatchTypeIrregular();
int numControlVertices = Far::internal::EvaluatePatchBasis(patchType,
param, coord.s, coord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
int indexBase = array.GetIndexBase() + array.GetStride() *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
@ -353,7 +364,7 @@ CpuEvaluator::EvalPatches(const float *src, BufferDescriptor const &srcDesc,
duuT.Clear();
duvT.Clear();
dvvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
duT.AddWithWeight (srcT[cvs[j]], wDu[j]);
dvT.AddWithWeight (srcT[cvs[j]], wDv[j]);

147
opensubdiv/osd/cudaKernel.cu
View File

@ -24,7 +24,9 @@
#include <assert.h>
#define OSD_PATCH_BASIS_CUDA
#include "../osd/patchBasisCommonTypes.h"
#include "../osd/patchBasisCommon.h"
#include "../osd/patchBasisCommonEval.h"
// -----------------------------------------------------------------------------
template<int N> struct DeviceVertex {
@ -240,70 +242,6 @@ __global__ void computeStencilsNv_v4(float const *__restrict cvs,
// -----------------------------------------------------------------------------
// Osd::PatchCoord osd/types.h
struct PatchCoord {
int arrayIndex;
int patchIndex;
int vertIndex;
float s;
float t;
};
struct PatchArray {
int patchType; // Far::PatchDescriptor::Type
int numPatches;
int indexBase; // offset in the index buffer
int primitiveIdBase; // offset in the patch param buffer
};
struct PatchParam {
unsigned int field0;
unsigned int field1;
float sharpness;
};
__device__
int getDepth(unsigned int patchBits) {
return (patchBits & 0xf);
}
__device__
float getParamFraction(unsigned int patchBits) {
bool nonQuadRoot = (patchBits >> 4) & 0x1;
int depth = getDepth(patchBits);
if (nonQuadRoot) {
return 1.0f / float( 1 << (depth-1) );
} else {
return 1.0f / float( 1 << depth );
}
}
__device__
void normalizePatchCoord(unsigned int patchBits, float *u, float *v) {
float frac = getParamFraction(patchBits);
int iu = (patchBits >> 22) & 0x3ff;
int iv = (patchBits >> 12) & 0x3ff;
// top left corner
float pu = (float)iu*frac;
float pv = (float)iv*frac;
// normalize u,v coordinates
*u = (*u - pu) / frac;
*v = (*v - pv) / frac;
}
__device__
bool isRegular(unsigned int patchBits) {
return ((patchBits >> 5) & 0x1) != 0;
}
__device__
int getNumControlVertices(int patchType) {
return (patchType == 3) ? 4 :
(patchType == 6) ? 16 :
(patchType == 9) ? 20 : 0;
}
__global__ void
computePatches(const float *src, float *dst,
float *dstDu, float *dstDv,
@ -311,97 +249,80 @@ computePatches(const float *src, float *dst,
int length, int srcStride, int dstStride,
int dstDuStride, int dstDvStride,
int dstDuuStride, int dstDuvStride, int dstDvvStride,
int numPatchCoords, const PatchCoord *patchCoords,
const PatchArray *patchArrayBuffer,
int numPatchCoords, const OsdPatchCoord *patchCoords,
const OsdPatchArray *patchArrayBuffer,
const int *patchIndexBuffer,
const PatchParam *patchParamBuffer) {
const OsdPatchParam *patchParamBuffer) {
int first = threadIdx.x + blockIdx.x * blockDim.x;
// PERFORMANCE: not yet optimized
float wP[20], wDs[20], wDt[20], wDss[20], wDst[20], wDtt[20];
for (int i = first; i < numPatchCoords; i += blockDim.x * gridDim.x) {
PatchCoord const &coord = patchCoords[i];
PatchArray const &array = patchArrayBuffer[coord.arrayIndex];
OsdPatchCoord const &coord = patchCoords[i];
int arrayIndex = coord.arrayIndex;
int patchIndex = coord.patchIndex;
unsigned int patchBits = patchParamBuffer[coord.patchIndex].field1;
int patchType = isRegular(patchBits) ? 6 : array.patchType;
OsdPatchArray const &array = patchArrayBuffer[arrayIndex];
OsdPatchParam const &param = patchParamBuffer[patchIndex];
// normalize
float s = coord.s;
float t = coord.t;
normalizePatchCoord(patchBits, &s, &t);
float dScale = (float)(1 << getDepth(patchBits));
int boundary = int((patchBits >> 7) & 0x1fU);
int patchType = OsdPatchParamIsRegular(param)
? array.regDesc : array.desc;
int numControlVertices = 0;
if (patchType == 3) {
OsdGetBilinearPatchWeights(s, t, dScale,
wP, wDs, wDt, wDss, wDst, wDtt);
numControlVertices = 4;
} else if (patchType == 6) {
OsdGetBSplinePatchWeights(s, t, dScale, boundary,
wP, wDs, wDt, wDss, wDst, wDtt);
numControlVertices = 16;
} else if (patchType == 9) {
OsdGetGregoryPatchWeights(s, t, dScale,
wP, wDs, wDt, wDss, wDst, wDtt);
numControlVertices = 20;
}
float wP[20], wDu[20], wDv[20], wDuu[20], wDuv[20], wDvv[20];
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
int indexStride = getNumControlVertices(array.patchType);
int indexBase = array.indexBase + indexStride *
(coord.patchIndex - array.primitiveIdBase);
int indexBase = array.indexBase + array.stride *
(patchIndex - array.primitiveIdBase);
const int *cvs = patchIndexBuffer + indexBase;
float * dstVert = dst + i * dstStride;
clear(dstVert, length);
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
const float * srcVert = src + cvs[j] * srcStride;
addWithWeight(dstVert, srcVert, wP[j], length);
}
if (dstDu) {
float *d = dstDu + i * dstDuStride;
clear(d, length);
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
const float * srcVert = src + cvs[j] * srcStride;
addWithWeight(d, srcVert, wDs[j], length);
addWithWeight(d, srcVert, wDu[j], length);
}
}
if (dstDv) {
float *d = dstDv + i * dstDvStride;
clear(d, length);
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
const float * srcVert = src + cvs[j] * srcStride;
addWithWeight(d, srcVert, wDt[j], length);
addWithWeight(d, srcVert, wDv[j], length);
}
}
if (dstDuu) {
float *d = dstDuu + i * dstDuuStride;
clear(d, length);
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
const float * srcVert = src + cvs[j] * srcStride;
addWithWeight(d, srcVert, wDss[j], length);
addWithWeight(d, srcVert, wDuu[j], length);
}
}
if (dstDuv) {
float *d = dstDuv + i * dstDuvStride;
clear(d, length);
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
const float * srcVert = src + cvs[j] * srcStride;
addWithWeight(d, srcVert, wDst[j], length);
addWithWeight(d, srcVert, wDuv[j], length);
}
}
if (dstDvv) {
float *d = dstDvv + i * dstDvvStride;
clear(d, length);
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
const float * srcVert = src + cvs[j] * srcStride;
addWithWeight(d, srcVert, wDtt[j], length);
addWithWeight(d, srcVert, wDvv[j], length);
}
}
}
@ -467,10 +388,10 @@ void CudaEvalStencils(
void CudaEvalPatches(
const float *src, float *dst,
int length, int srcStride, int dstStride,
int numPatchCoords, const PatchCoord *patchCoords,
const PatchArray *patchArrayBuffer,
int numPatchCoords, const OsdPatchCoord *patchCoords,
const OsdPatchArray *patchArrayBuffer,
const int *patchIndexBuffer,
const PatchParam *patchParamBuffer) {
const OsdPatchParam *patchParamBuffer) {
// PERFORMANCE: not optimized at all
@ -488,10 +409,10 @@ void CudaEvalPatchesWithDerivatives(
int length, int srcStride, int dstStride,
int dstDuStride, int dstDvStride,
int dstDuuStride, int dstDuvStride, int dstDvvStride,
int numPatchCoords, const PatchCoord *patchCoords,
const PatchArray *patchArrayBuffer,
int numPatchCoords, const OsdPatchCoord *patchCoords,
const OsdPatchArray *patchArrayBuffer,
const int *patchIndexBuffer,
const PatchParam *patchParamBuffer) {
const OsdPatchParam *patchParamBuffer) {
// PERFORMANCE: not optimized at all

30
opensubdiv/osd/glComputeEvaluator.cpp
View File

@ -122,12 +122,17 @@ GLStencilTableSSBO::~GLStencilTableSSBO() {
// ---------------------------------------------------------------------------
GLComputeEvaluator::GLComputeEvaluator() : _workGroupSize(64) {
GLComputeEvaluator::GLComputeEvaluator()
: _workGroupSize(64),
_patchArraysSSBO(0) {
memset (&_stencilKernel, 0, sizeof(_stencilKernel));
memset (&_patchKernel, 0, sizeof(_patchKernel));
}
GLComputeEvaluator::~GLComputeEvaluator() {
if (_patchArraysSSBO) {
glDeleteBuffers(1, &_patchArraysSSBO);
}
}
static GLuint
@ -222,6 +227,11 @@ GLComputeEvaluator::Compile(BufferDescriptor const &srcDesc,
return false;
}
// create a patch arrays buffer
if (!_patchArraysSSBO) {
glGenBuffers(1, &_patchArraysSSBO);
}
return true;
}
@ -398,16 +408,24 @@ GLComputeEvaluator::EvalPatches(
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);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 5, patchCoordsBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 6, patchIndexBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 7, patchParamsBuffer);
glUseProgram(_patchKernel.program);
glUniform1i(_patchKernel.uniformSrcOffset, srcDesc.offset);
glUniform1i(_patchKernel.uniformDstOffset, dstDesc.offset);
glUniform4iv(_patchKernel.uniformPatchArray, (int)patchArrays.size(),
(const GLint*)&patchArrays[0]);
int patchArraySize = sizeof(PatchArray);
glBindBuffer(GL_SHADER_STORAGE_BUFFER, _patchArraysSSBO);
glBufferData(GL_SHADER_STORAGE_BUFFER,
patchArrays.size()*patchArraySize, NULL, GL_STATIC_DRAW);
for (int i=0; i<(int)patchArrays.size(); ++i) {
glBufferSubData(GL_SHADER_STORAGE_BUFFER,
i*patchArraySize, sizeof(PatchArray), &patchArrays[i]);
}
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 4, _patchArraysSSBO);
if (_patchKernel.uniformDuDesc > 0) {
glUniform3i(_patchKernel.uniformDuDesc,

1
opensubdiv/osd/glComputeEvaluator.h
View File

@ -2101,6 +2101,7 @@ private:
} _patchKernel;
int _workGroupSize;
GLuint _patchArraysSSBO;
};
} // end namespace Osd

31
opensubdiv/osd/glXFBEvaluator.cpp
View File

@ -153,6 +153,7 @@ GLStencilTableTBO::~GLStencilTableTBO() {
GLXFBEvaluator::GLXFBEvaluator(bool interleavedDerivativeBuffers)
: _srcBufferTexture(0),
_patchArraysUBO(0),
_interleavedDerivativeBuffers(interleavedDerivativeBuffers) {
}
@ -160,6 +161,9 @@ GLXFBEvaluator::~GLXFBEvaluator() {
if (_srcBufferTexture) {
glDeleteTextures(1, &_srcBufferTexture);
}
if (_patchArraysUBO) {
glDeleteBuffers(1, &_patchArraysUBO);
}
}
static GLuint
@ -439,6 +443,9 @@ GLXFBEvaluator::Compile(BufferDescriptor const &srcDesc,
if (!_srcBufferTexture) {
glGenTextures(1, &_srcBufferTexture);
}
if (!_patchArraysUBO) {
glGenBuffers(1, &_patchArraysUBO);
}
return true;
}
@ -735,9 +742,20 @@ GLXFBEvaluator::EvalPatches(
bindTexture(_patchKernel.uniformPatchParamTexture, patchParamTexture, 1);
bindTexture(_patchKernel.uniformPatchIndexTexture, patchIndexTexture, 2);
// bind patch arrays UBO (std140 struct size padded to vec4 alignment)
int patchArraySize =
sizeof(GLint) * ((sizeof(PatchArray)/sizeof(GLint) + 3) & ~3);
glBindBuffer(GL_UNIFORM_BUFFER, _patchArraysUBO);
glBufferData(GL_UNIFORM_BUFFER,
patchArrays.size()*patchArraySize, NULL, GL_STATIC_DRAW);
for (int i=0; i<(int)patchArrays.size(); ++i) {
glBufferSubData(GL_UNIFORM_BUFFER,
i*patchArraySize, sizeof(PatchArray), &patchArrays[i]);
}
glBindBufferBase(GL_UNIFORM_BUFFER,
_patchKernel.uniformPatchArraysUBOBinding, _patchArraysUBO);
// set other uniforms
glUniform4iv(_patchKernel.uniformPatchArray, (int)patchArrays.size(),
(const GLint*)&patchArrays[0]);
glUniform1i(_patchKernel.uniformSrcOffset, srcDesc.offset);
// input patchcoords
@ -822,6 +840,10 @@ GLXFBEvaluator::EvalPatches(
glBindTexture(GL_TEXTURE_BUFFER, 0);
}
// unbind UBO
glBindBufferBase(GL_UNIFORM_BUFFER,
_patchKernel.uniformPatchArraysUBOBinding, 0);
glDisable(GL_RASTERIZER_DISCARD);
glUseProgram(0);
glActiveTexture(GL_TEXTURE0);
@ -921,10 +943,13 @@ GLXFBEvaluator::_PatchKernel::Compile(BufferDescriptor const &srcDesc,
// cache uniform locations
uniformSrcBufferTexture = glGetUniformLocation(program, "vertexBuffer");
uniformSrcOffset = glGetUniformLocation(program, "srcOffset");
uniformPatchArray = glGetUniformLocation(program, "patchArray");
uniformPatchParamTexture = glGetUniformLocation(program, "patchParamBuffer");
uniformPatchIndexTexture = glGetUniformLocation(program, "patchIndexBuffer");
uniformPatchArraysUBOBinding = 1;
int uboIndex = glGetUniformBlockIndex(program, "PatchArrays");
glUniformBlockBinding(program, uboIndex , uniformPatchArraysUBOBinding);
return true;
}

3
opensubdiv/osd/glXFBEvaluator.h
View File

@ -2129,6 +2129,7 @@ public:
private:
GLuint _srcBufferTexture;
GLuint _patchArraysUBO;
bool _interleavedDerivativeBuffers;
struct _StencilKernel {
@ -2174,7 +2175,7 @@ private:
GLint uniformSrcBufferTexture;
GLint uniformSrcOffset; // src buffer offset (in elements)
GLint uniformPatchArray;
GLint uniformPatchArraysUBOBinding;
GLint uniformPatchParamTexture;
GLint uniformPatchIndexTexture;
} _patchKernel;

145
opensubdiv/osd/glslComputeKernel.glsl
View File

@ -81,22 +81,25 @@ layout(binding=15) buffer stencilDvvWeights { float _dvvWeights[]; };
#if defined(OPENSUBDIV_GLSL_COMPUTE_KERNEL_EVAL_PATCHES)
struct PatchCoord {
int arrayIndex;
int patchIndex;
int vertIndex;
float s;
float t;
};
struct PatchParam {
uint field0;
uint field1;
float sharpness;
};
uniform ivec4 patchArray[2];
layout(binding=4) buffer patchCoord_buffer { PatchCoord patchCoords[]; };
layout(binding=5) buffer patchIndex_buffer { int patchIndexBuffer[]; };
layout(binding=6) buffer patchParam_buffer { PatchParam patchParamBuffer[]; };
layout(binding=4) buffer patchArray_buffer { OsdPatchArray patchArrayBuffer[]; };
layout(binding=5) buffer patchCoord_buffer { OsdPatchCoord patchCoords[]; };
layout(binding=6) buffer patchIndex_buffer { int patchIndexBuffer[]; };
layout(binding=7) buffer patchParam_buffer { OsdPatchParam patchParamBuffer[]; };
OsdPatchCoord GetPatchCoord(int coordIndex)
{
return patchCoords[coordIndex];
}
OsdPatchArray GetPatchArray(int arrayIndex)
{
return patchArrayBuffer[arrayIndex];
}
OsdPatchParam GetPatchParam(int patchIndex)
{
return patchParamBuffer[patchIndex];
}
#endif
@ -248,99 +251,19 @@ void main() {
// PERFORMANCE: stride could be constant, but not as significant as length
//struct PatchArray {
// int patchType;
// int numPatches;
// int indexBase; // an offset within the index buffer
// int primitiveIdBase; // an offset within the patch param buffer
//};
// # of patcharrays is 1 or 2.
uint getDepth(uint patchBits) {
return (patchBits & 0xf);
}
float getParamFraction(uint patchBits) {
uint nonQuadRoot = (patchBits >> 4) & 0x1;
uint depth = getDepth(patchBits);
if (nonQuadRoot == 1) {
return 1.0f / float( 1 << (depth-1) );
} else {
return 1.0f / float( 1 << depth );
}
}
vec2 normalizePatchCoord(uint patchBits, vec2 uv) {
float frac = getParamFraction(patchBits);
uint iu = (patchBits >> 22) & 0x3ff;
uint iv = (patchBits >> 12) & 0x3ff;
// top left corner
float pu = float(iu*frac);
float pv = float(iv*frac);
// normalize u,v coordinates
return vec2((uv.x - pu) / frac, (uv.y - pv) / frac);
}
bool isRegular(uint patchBits) {
return (((patchBits >> 5) & 0x1u) != 0);
}
int getNumControlVertices(int patchType) {
return (patchType == 3) ? 4 :
(patchType == 6) ? 16 :
(patchType == 9) ? 20 : 0;
}
void main() {
int current = int(gl_GlobalInvocationID.x);
PatchCoord coord = patchCoords[current];
int patchIndex = coord.patchIndex;
OsdPatchCoord coord = GetPatchCoord(current);
OsdPatchArray array = GetPatchArray(coord.arrayIndex);
OsdPatchParam param = GetPatchParam(coord.patchIndex);
ivec4 array = patchArray[coord.arrayIndex];
int patchType = OsdPatchParamIsRegular(param) ? array.regDesc : array.desc;
uint patchBits = patchParamBuffer[patchIndex].field1;
int patchType = isRegular(patchBits) ? 6 : array.x;
vec2 uv = normalizePatchCoord(patchBits, vec2(coord.s, coord.t));
float dScale = float(1 << getDepth(patchBits));
int boundary = int((patchBits >> 7) & 0x1fU);
float wP[20], wDs[20], wDt[20], wDss[20], wDst[20], wDtt[20];
int numControlVertices = 0;
if (patchType == 3) {
float wP4[4], wDs4[4], wDt4[4], wDss4[4], wDst4[4], wDtt4[4];
OsdGetBilinearPatchWeights(uv.s, uv.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(uv.s, uv.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(uv.s, uv.t, dScale, wP, wDs, wDt, wDss, wDst, wDtt);
numControlVertices = 20;
}
float wP[20], wDu[20], wDv[20], wDuu[20], wDuv[20], wDvv[20];
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
Vertex dst, du, dv, duu, duv, dvv;
clear(dst);
@ -350,17 +273,17 @@ void main() {
clear(duv);
clear(dvv);
int indexStride = getNumControlVertices(array.x);
int indexBase = array.z + indexStride * (patchIndex - array.w);
int indexBase = array.indexBase + array.stride *
(coord.patchIndex - array.primitiveIdBase);
for (int cv = 0; cv < numControlVertices; ++cv) {
for (int cv = 0; cv < nPoints; ++cv) {
int index = patchIndexBuffer[indexBase + cv];
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]);
addWithWeight(du, readVertex(index), wDu[cv]);
addWithWeight(dv, readVertex(index), wDv[cv]);
addWithWeight(duu, readVertex(index), wDuu[cv]);
addWithWeight(duv, readVertex(index), wDuv[cv]);
addWithWeight(dvv, readVertex(index), wDvv[cv]);
}
writeVertex(current, dst);

8
opensubdiv/osd/glslPatchShaderSource.cpp
View File

@ -34,9 +34,15 @@ namespace Osd {
static const char *commonShaderSource =
#include "glslPatchCommon.gen.h"
;
static const char *patchBasisTypesShaderSource =
#include "patchBasisCommonTypes.gen.h"
;
static const char *patchBasisShaderSource =
#include "patchBasisCommon.gen.h"
;
static const char *patchBasisEvalShaderSource =
#include "patchBasisCommonEval.gen.h"
;
static const char *bsplineShaderSource =
#include "glslPatchBSpline.gen.h"
;
@ -60,7 +66,9 @@ GLSLPatchShaderSource::GetPatchBasisShaderSource() {
#if defined(OPENSUBDIV_GREGORY_EVAL_TRUE_DERIVATIVES)
ss << "#define OPENSUBDIV_GREGORY_EVAL_TRUE_DERIVATIVES\n";
#endif
ss << std::string(patchBasisTypesShaderSource);
ss << std::string(patchBasisShaderSource);
ss << std::string(patchBasisEvalShaderSource);
return ss.str();
}

115
opensubdiv/osd/glslXFBKernel.glsl
View File

@ -223,107 +223,38 @@ void main() {
layout (location = 0) in ivec3 patchHandles;
layout (location = 1) in vec2 patchCoords;
//struct PatchArray {
// int patchType;
// int numPatches;
// int indexBase; // an offset within the index buffer
// int primitiveIdBase; // an offset within the patch param buffer
//};
// # of patcharrays is 1 or 2.
uniform ivec4 patchArray[2];
layout (std140) uniform PatchArrays {
OsdPatchArray patchArrays[2];
};
uniform isamplerBuffer patchParamBuffer;
uniform isamplerBuffer patchIndexBuffer;
uint getDepth(uint patchBits) {
return (patchBits & 0xfU);
OsdPatchArray GetPatchArray(int arrayIndex) {
return patchArrays[arrayIndex];
}
float getParamFraction(uint patchBits) {
uint nonQuadRoot = (patchBits >> 4) & 0x1U;
uint depth = getDepth(patchBits);
if (nonQuadRoot == 1) {
return 1.0f / float( 1 << (depth-1) );
} else {
return 1.0f / float( 1 << depth );
}
}
vec2 normalizePatchCoord(uint patchBits, vec2 uv) {
float frac = getParamFraction(patchBits);
uint iu = (patchBits >> 22) & 0x3ffU;
uint iv = (patchBits >> 12) & 0x3ffU;
// top left corner
float pu = float(iu*frac);
float pv = float(iv*frac);
// normalize u,v coordinates
return vec2((uv.x - pu) / frac, (uv.y - pv) / frac);
}
bool isRegular(uint patchBits) {
return (((patchBits >> 5) & 0x1u) != 0);
}
int getNumControlVertices(int patchType) {
return (patchType == 3) ? 4 :
(patchType == 6) ? 16 :
(patchType == 9) ? 20 : 0;
OsdPatchParam GetPatchParam(int patchIndex) {
ivec3 patchParamBits = texelFetch(patchParamBuffer, patchIndex).xyz;
return OsdPatchParamInit(patchParamBits.x, patchParamBits.y, patchParamBits.z);
}
void main() {
int current = gl_VertexID;
ivec3 handle = patchHandles;
int arrayIndex = handle.x;
int patchIndex = handle.y;
vec2 coord = patchCoords;
ivec4 array = patchArray[handle.x];
uint patchBits = texelFetch(patchParamBuffer, patchIndex).y;
int patchType = isRegular(patchBits) ? 6 : array.x;
OsdPatchArray array = GetPatchArray(arrayIndex);
OsdPatchParam param = GetPatchParam(patchIndex);
// normalize
coord = normalizePatchCoord(patchBits, coord);
float dScale = float(1 << getDepth(patchBits));
int boundary = int((patchBits >> 7) & 0x1fU);
int patchType = OsdPatchParamIsRegular(param) ? array.regDesc : array.desc;
float wP[20], wDs[20], wDt[20], wDss[20], wDst[20], wDtt[20];
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);
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);
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);
numControlVertices = 20;
}
float wP[20], wDu[20], wDv[20], wDuu[20], wDuv[20], wDvv[20];
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.x, coord.y, wP, wDu, wDv, wDuu, wDuv, wDvv);
Vertex dst, du, dv, duu, duv, dvv;
clear(dst);
@ -333,17 +264,17 @@ void main() {
clear(duv);
clear(dvv);
int indexStride = getNumControlVertices(array.x);
int indexBase = array.z + indexStride * (patchIndex - array.w);
int indexBase = array.indexBase + array.stride *
(patchIndex - array.primitiveIdBase);
for (int cv = 0; cv < numControlVertices; ++cv) {
for (int cv = 0; cv < nPoints; ++cv) {
int index = texelFetch(patchIndexBuffer, indexBase + cv).x;
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]);
addWithWeight(du, readVertex(index), wDu[cv]);
addWithWeight(dv, readVertex(index), wDv[cv]);
addWithWeight(duu, readVertex(index), wDuu[cv]);
addWithWeight(duv, readVertex(index), wDuv[cv]);
addWithWeight(dvv, readVertex(index), wDvv[cv]);
}
writeVertex(dst);

8
opensubdiv/osd/hlslPatchShaderSource.cpp
View File

@ -36,9 +36,15 @@ namespace Osd {
static const char *commonShaderSource =
#include "hlslPatchCommon.gen.h"
;
static const char *patchBasisTypesShaderSource =
#include "patchBasisCommonTypes.gen.h"
;
static const char *patchBasisShaderSource =
#include "patchBasisCommon.gen.h"
;
static const char *patchBasisEvalShaderSource =
#include "patchBasisCommonEval.gen.h"
;
static const char *bsplineShaderSource =
#include "hlslPatchBSpline.gen.h"
;
@ -62,7 +68,9 @@ HLSLPatchShaderSource::GetPatchBasisShaderSource() {
#if defined(OPENSUBDIV_GREGORY_EVAL_TRUE_DERIVATIVES)
ss << "#define OPENSUBDIV_GREGORY_EVAL_TRUE_DERIVATIVES\n";
#endif
ss << std::string(patchBasisTypesShaderSource);
ss << std::string(patchBasisShaderSource);
ss << std::string(patchBasisEvalShaderSource);
return ss.str();
}

135
opensubdiv/osd/mtlComputeKernel.metal
View File

@ -36,29 +36,13 @@
using namespace metal;
struct PatchCoord
{
int arrayIndex;
int patchIndex;
int vertIndex;
float s;
float t;
};
struct PatchParam
{
uint field0;
uint field1;
float sharpness;
};
struct KernelUniformArgs
{
int batchStart;
int batchEnd;
int batchStart;
int batchEnd;
int srcOffset;
int dstOffset;
int dstOffset;
int3 duDesc;
int3 dvDesc;
@ -179,7 +163,7 @@ kernel void eval_stencils(
)
{
auto current = thread_position_in_grid + args.batchStart;
if(current >= args.batchEnd)
if(current >= (unsigned int)args.batchEnd)
return;
Vertex dst;
@ -240,60 +224,11 @@ kernel void eval_stencils(
// PERFORMANCE: stride could be constant, but not as significant as length
//struct PatchArray {
// int patchType;
// int numPatches;
// int indexBase; // an offset within the index buffer
// int primitiveIdBase; // an offset within the patch param buffer
//};
// # of patcharrays is 1 or 2.
uint getDepth(uint patchBits) {
return (patchBits & 0xf);
}
float getParamFraction(uint patchBits) {
uint nonQuadRoot = (patchBits >> 4) & 0x1;
uint depth = getDepth(patchBits);
if (nonQuadRoot == 1) {
return 1.0f / float( 1 << (depth-1) );
} else {
return 1.0f / float( 1 << depth );
}
}
float2 normalizePatchCoord(uint patchBits, float2 uv) {
float frac = getParamFraction(patchBits);
uint iu = (patchBits >> 22) & 0x3ff;
uint iv = (patchBits >> 12) & 0x3ff;
// top left corner
float pu = float(iu*frac);
float pv = float(iv*frac);
// normalize u,v coordinates
return float2((uv.x - pu) / frac, (uv.y - pv) / frac);
}
bool isRegular(uint patchBits) {
return (((patchBits >> 5) & 0x1u) != 0);
}
int getNumControlVertices(int patchType) {
switch(patchType) {
case 3: return 4;
case 6: return 16;
case 9: return 20;
default: return 0;
}
}
// ---------------------------------------------------------------------------
kernel void eval_patches(
uint thread_position_in_grid [[thread_position_in_grid]],
const constant uint4* patchArrays [[buffer(PATCH_ARRAYS_BUFFER_INDEX)]],
const constant int* patchArrays [[buffer(PATCH_ARRAYS_BUFFER_INDEX)]],
const device int* patchCoords [[buffer(PATCH_COORDS_BUFFER_INDEX)]],
const device int* patchIndices [[buffer(PATCH_INDICES_BUFFER_INDEX)]],
const device uint* patchParams [[buffer(PATCH_PARAMS_BUFFER_INDEX)]],
@ -314,34 +249,29 @@ kernel void eval_patches(
auto current = thread_position_in_grid;
// unpack struct (5 ints unaligned)
PatchCoord patchCoord;
patchCoord.arrayIndex = patchCoords[current*5+0];
patchCoord.patchIndex = patchCoords[current*5+1];
patchCoord.vertIndex = patchCoords[current*5+2];
patchCoord.s = as_type<float>(patchCoords[current*5+3]);
patchCoord.t = as_type<float>(patchCoords[current*5+4]);
OsdPatchCoord patchCoord = OsdPatchCoordInit(patchCoords[current*5+0],
patchCoords[current*5+1],
patchCoords[current*5+2],
as_type<float>(patchCoords[current*5+3]),
as_type<float>(patchCoords[current*5+4]));
auto patchArray = patchArrays[patchCoord.arrayIndex];
OsdPatchArray patchArray = OsdPatchArrayInit(patchArrays[current*6+0],
patchArrays[current*6+1],
patchArrays[current*6+2],
patchArrays[current*6+3],
patchArrays[current*6+4],
patchArrays[current*6+5]);
// unpack struct (3 uints unaligned)
auto patchBits = patchParams[patchCoord.patchIndex*3+1]; // field1
auto patchType = select(patchArray.x, uint(6), isRegular(patchBits));
OsdPatchParam patchParam = OsdPatchParamInit(patchParams[current*3+0],
patchParams[current*3+1],
as_type<float>(patchParams[current*3+2]));
auto numControlVertices = getNumControlVertices(patchType);
auto uv = normalizePatchCoord(patchBits, float2(patchCoord.s, patchCoord.t));
auto dScale = float(1 << getDepth(patchBits));
auto boundary = int((patchBits >> 7) & 0x1FU);
int patchType = OsdPatchParamIsRegular(patchParam)
? patchArray.regDesc : patchArray.desc;
float wP[20], wDs[20], wDt[20], wDss[20], wDst[20], wDtt[20];
if(patchType == 3) {
OsdGetBilinearPatchWeights(uv.x, uv.y, dScale, wP, wDs, wDt, wDss, wDst, wDtt);
} else if(patchType == 6) {
OsdGetBSplinePatchWeights(uv.x, uv.y, dScale, boundary, wP, wDs, wDt, wDss, wDst, wDtt);
} else if(patchType == 9) {
OsdGetGregoryPatchWeights(uv.x, uv.y, dScale, wP, wDs, wDt, wDss, wDst, wDtt);
}
float wP[20], wDu[20], wDv[20], wDuu[20], wDuv[20], wDvv[20];
int nPoints = OsdEvaluatePatchBasis(patchType, patchParam,
patchCoord.s, patchCoord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
Vertex dst, du, dv, duu, duv, dvv;
clear(dst);
@ -351,19 +281,18 @@ kernel void eval_patches(
clear(duv);
clear(dvv);
auto indexStride = getNumControlVertices(patchArray.x);
auto indexBase = patchArray.z + indexStride * (patchCoord.patchIndex - patchArray.w);
for(auto cv = 0; cv < numControlVertices; cv++)
auto indexBase = patchArray.indexBase + patchArray.stride *
(patchCoord.patchIndex - patchArray.primitiveIdBase);
for(auto cv = 0; cv < nPoints; cv++)
{
auto index = patchIndices[indexBase + cv];
auto src = readVertex(index, srcVertexBuffer, args);
addWithWeight(dst, src, wP[cv]);
addWithWeight(du, src, wDs[cv]);
addWithWeight(dv, src, wDt[cv]);
addWithWeight(duu, src, wDss[cv]);
addWithWeight(duv, src, wDst[cv]);
addWithWeight(dvv, src, wDtt[cv]);
addWithWeight(du, src, wDu[cv]);
addWithWeight(dv, src, wDv[cv]);
addWithWeight(duu, src, wDuu[cv]);
addWithWeight(duv, src, wDuv[cv]);
addWithWeight(dvv, src, wDvv[cv]);
}
writeVertex(current, dst, dstVertexBuffer, args);

12
opensubdiv/osd/mtlPatchShaderSource.mm
View File

@ -37,9 +37,15 @@ namespace OpenSubdiv {
static std::string commonShaderSource(
#include "mtlPatchCommon.gen.h"
);
static std::string patchBasisShaderSource(
static std::string patchBasisTypesShaderSource(
#include "patchBasisCommonTypes.gen.h"
);
static std::string patchBasisShaderSource(
#include "patchBasisCommon.gen.h"
);
);
static std::string patchBasisEvalShaderSource(
#include "patchBasisCommonEval.gen.h"
);
static std::string bsplineShaderSource(
#include "mtlPatchBSpline.gen.h"
);
@ -105,7 +111,9 @@ namespace OpenSubdiv {
#if defined(OPENSUBDIV_GREGORY_EVAL_TRUE_DERIVATIVES)
ss << "#define OPENSUBDIV_GREGORY_EVAL_TRUE_DERIVATIVES 1\n";
#endif
ss << patchBasisTypesShaderSource;
ss << patchBasisShaderSource;
ss << patchBasisEvalShaderSource;
return ss.str();
}

66
opensubdiv/osd/ompEvaluator.cpp
View File

@ -24,7 +24,9 @@
#include "../osd/ompEvaluator.h"
#include "../osd/ompKernel.h"
#include "../far/patchBasis.h"
#include "../osd/patchBasisCommonTypes.h"
#include "../osd/patchBasisCommon.h"
#include "../osd/patchBasisCommonEval.h"
#include <omp.h>
namespace OpenSubdiv {
@ -178,27 +180,29 @@ OmpEvaluator::EvalPatches(
for (int i = 0; i < numPatchCoords; ++i) {
BufferAdapter<float> dstT(dst + dstDesc.stride*i, dstDesc.length, dstDesc.stride);
float wP[20], wDs[20], wDt[20];
float wP[20];
PatchCoord const &coord = patchCoords[i];
PatchArray const &array = patchArrays[coord.handle.arrayIndex];
Far::PatchParam const & param =
Osd::PatchParam const & paramStruct =
patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
? Far::PatchDescriptor::REGULAR
: array.GetPatchType();
OsdPatchParam param = OsdPatchParamInit(
paramStruct.field0, paramStruct.field1, paramStruct.sharpness);
int numControlVertices = Far::internal::EvaluatePatchBasis(patchType,
param, coord.s, coord.t, wP, wDs, wDt);
int patchType = OsdPatchParamIsRegular(param)
? array.GetPatchTypeRegular()
: array.GetPatchTypeIrregular();
int indexStride = Far::PatchDescriptor(array.GetPatchType()).GetNumControlVertices();
int indexBase = array.GetIndexBase() + indexStride *
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, 0, 0, 0, 0, 0);
int indexBase = array.GetIndexBase() + array.GetStride() *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
const int *cvs = &patchIndexBuffer[indexBase];
dstT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
}
}
@ -235,17 +239,19 @@ OmpEvaluator::EvalPatches(
PatchCoord const &coord = patchCoords[i];
PatchArray const &array = patchArrays[coord.handle.arrayIndex];
Far::PatchParam const & param =
Osd::PatchParam const & paramStruct =
patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
? Far::PatchDescriptor::REGULAR
: array.GetPatchType();
OsdPatchParam param = OsdPatchParamInit(
paramStruct.field0, paramStruct.field1, paramStruct.sharpness);
int numControlVertices = Far::internal::EvaluatePatchBasis(patchType,
param, coord.s, coord.t, wP, wDu, wDv);
int patchType = OsdPatchParamIsRegular(param)
? array.GetPatchTypeRegular()
: array.GetPatchTypeIrregular();
int indexStride = Far::PatchDescriptor(array.GetPatchType()).GetNumControlVertices();
int indexBase = array.GetIndexBase() + indexStride *
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, wDu, wDv, 0, 0, 0);
int indexBase = array.GetIndexBase() + array.GetStride() *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
const int *cvs = &patchIndexBuffer[indexBase];
@ -253,7 +259,7 @@ OmpEvaluator::EvalPatches(
dstT.Clear();
duT.Clear();
dvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
duT.AddWithWeight(srcT[cvs[j]], wDu[j]);
dvT.AddWithWeight(srcT[cvs[j]], wDv[j]);
@ -304,17 +310,19 @@ OmpEvaluator::EvalPatches(
PatchCoord const &coord = patchCoords[i];
PatchArray const &array = patchArrays[coord.handle.arrayIndex];
Far::PatchParam const & param =
Osd::PatchParam const & paramStruct =
patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
? Far::PatchDescriptor::REGULAR
: array.GetPatchType();
OsdPatchParam param = OsdPatchParamInit(
paramStruct.field0, paramStruct.field1, paramStruct.sharpness);
int numControlVertices = Far::internal::EvaluatePatchBasis(patchType,
param, coord.s, coord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
int patchType = OsdPatchParamIsRegular(param)
? array.GetPatchTypeRegular()
: array.GetPatchTypeIrregular();
int indexStride = Far::PatchDescriptor(array.GetPatchType()).GetNumControlVertices();
int indexBase = array.GetIndexBase() + indexStride *
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
int indexBase = array.GetIndexBase() + array.GetStride() *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
const int *cvs = &patchIndexBuffer[indexBase];
@ -325,7 +333,7 @@ OmpEvaluator::EvalPatches(
duuT.Clear();
duvT.Clear();
dvvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
duT.AddWithWeight(srcT[cvs[j]], wDu[j]);
dvT.AddWithWeight(srcT[cvs[j]], wDv[j]);

1577
opensubdiv/osd/patchBasisCommon.h
View File

File diff suppressed because it is too large Load Diff

196
opensubdiv/osd/patchBasisCommonEval.h Normal file
View File

@ -0,0 +1,196 @@
//
// Copyright 2018 Pixar
//
// Licensed under the Apache License, Version 2.0 (the "Apache License")
// with the following modification; you may not use this file except in
// compliance with the Apache License and the following modification to it:
// Section 6. Trademarks. is deleted and replaced with:
//
// 6. Trademarks. This License does not grant permission to use the trade
// names, trademarks, service marks, or product names of the Licensor
// and its affiliates, except as required to comply with Section 4(c) of
// the License and to reproduce the content of the NOTICE file.
//
// You may obtain a copy of the Apache License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the Apache License with the above modification is
// distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
// KIND, either express or implied. See the Apache License for the specific
// language governing permissions and limitations under the Apache License.
//
#ifndef OPENSUBDIV3_OSD_PATCH_BASIS_COMMON_EVAL_H
#define OPENSUBDIV3_OSD_PATCH_BASIS_COMMON_EVAL_H
OSD_FUNCTION_STORAGE_CLASS
// template <typename REAL>
int
OsdEvaluatePatchBasisNormalized(
int patchType, OsdPatchParam param,
OSD_REAL s, OSD_REAL t,
OSD_OUT_ARRAY(OSD_REAL, wP, 20),
OSD_OUT_ARRAY(OSD_REAL, wDs, 20),
OSD_OUT_ARRAY(OSD_REAL, wDt, 20),
OSD_OUT_ARRAY(OSD_REAL, wDss, 20),
OSD_OUT_ARRAY(OSD_REAL, wDst, 20),
OSD_OUT_ARRAY(OSD_REAL, wDtt, 20)) {
int boundaryMask = OsdPatchParamGetBoundary(param);
int nPoints = 0;
if (patchType == OSD_PATCH_DESCRIPTOR_REGULAR) {
#if OSD_ARRAY_ARG_BOUND_OPTIONAL
nPoints = Osd_EvalBasisBSpline(s, t, wP, wDs, wDt, wDss, wDst, wDtt);
if (boundaryMask != 0) {
Osd_boundBasisBSpline(
boundaryMask, wP, wDs, wDt, wDss, wDst, wDtt);
}
#else
OSD_REAL wP16[16], wDs16[16], wDt16[16],
wDss16[16], wDst16[16], wDtt16[16];
nPoints = Osd_EvalBasisBSpline(
s, t, wP16, wDs16, wDt16, wDss16, wDst16, wDtt16);
if (boundaryMask != 0) {
Osd_boundBasisBSpline(
boundaryMask, wP16, wDs16, wDt16, wDss16, wDst16, wDtt16);
}
for (int i=0; i<nPoints; ++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];
}
#endif
} else if (patchType == OSD_PATCH_DESCRIPTOR_LOOP) {
#if OSD_ARRAY_ARG_BOUND_OPTIONAL
nPoints = Osd_EvalBasisBoxSplineTri(
s, t, wP, wDs, wDt, wDss, wDst, wDtt);
if (boundaryMask != 0) {
Osd_boundBasisBoxSplineTri(
boundaryMask, wP, wDs, wDt, wDss, wDst, wDtt);
}
#else
OSD_REAL wP12[12], wDs12[12], wDt12[12],
wDss12[12], wDst12[12], wDtt12[12];
nPoints = Osd_EvalBasisBoxSplineTri(
s, t, wP12, wDs12, wDt12, wDss12, wDst12, wDtt12);
if (boundaryMask != 0) {
Osd_boundBasisBoxSplineTri(
boundaryMask, wP12, wDs12, wDt12, wDss12, wDst12, wDtt12);
}
for (int i=0; i<nPoints; ++i) {
wP[i] = wP12[i];
wDs[i] = wDs12[i]; wDt[i] = wDt12[i];
wDss[i] = wDss12[i]; wDst[i] = wDst12[i]; wDtt[i] = wDtt12[i];
}
#endif
} else if (patchType == OSD_PATCH_DESCRIPTOR_GREGORY_BASIS) {
nPoints = Osd_EvalBasisGregory(s, t, wP, wDs, wDt, wDss, wDst, wDtt);
} else if (patchType == OSD_PATCH_DESCRIPTOR_GREGORY_TRIANGLE) {
#if OSD_ARRAY_ARG_BOUND_OPTIONAL
nPoints = Osd_EvalBasisGregoryTri(s, t, wP, wDs, wDt, wDss, wDst, wDtt);
#else
OSD_REAL wP15[15], wDs15[15], wDt15[15],
wDss15[15], wDst15[15], wDtt15[15];
nPoints = Osd_EvalBasisGregoryTri(
s, t, wP15, wDs15, wDt15, wDss15, wDst15, wDtt15);
for (int i=0; i<nPoints; ++i) {
wP[i] = wP15[i];
wDs[i] = wDs15[i]; wDt[i] = wDt15[i];
wDss[i] = wDss15[i]; wDst[i] = wDst15[i]; wDtt[i] = wDtt15[i];
}
#endif
} else if (patchType == OSD_PATCH_DESCRIPTOR_QUADS) {
#if OSD_ARRAY_ARG_BOUND_OPTIONAL
nPoints = Osd_EvalBasisLinear(s, t, wP, wDs, wDt, wDss, wDst, wDtt);
#else
OSD_REAL wP4[4], wDs4[4], wDt4[4],
wDss4[4], wDst4[4], wDtt4[4];
nPoints = Osd_EvalBasisLinear(
s, t, wP4, wDs4, wDt4, wDss4, wDst4, wDtt4);
for (int i=0; i<nPoints; ++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];
}
#endif
} else if (patchType == OSD_PATCH_DESCRIPTOR_TRIANGLES) {
#if OSD_ARRAY_ARG_BOUND_OPTIONAL
nPoints = Osd_EvalBasisLinearTri(s, t, wP, wDs, wDt, wDss, wDst, wDtt);
#else
OSD_REAL wP3[3], wDs3[3], wDt3[3],
wDss3[3], wDst3[3], wDtt3[3];
nPoints = Osd_EvalBasisLinearTri(
s, t, wP3, wDs3, wDt3, wDss3, wDst3, wDtt3);
for (int i=0; i<nPoints; ++i) {
wP[i] = wP3[i];
wDs[i] = wDs3[i]; wDt[i] = wDt3[i];
wDss[i] = wDss3[i]; wDst[i] = wDst3[i]; wDtt[i] = wDtt3[i];
}
#endif
} else {
// assert(0);
}
return nPoints;
}
OSD_FUNCTION_STORAGE_CLASS
// template <typename REAL>
int
OsdEvaluatePatchBasis(
int patchType, OsdPatchParam param,
OSD_REAL s, OSD_REAL t,
OSD_OUT_ARRAY(OSD_REAL, wP, 20),
OSD_OUT_ARRAY(OSD_REAL, wDs, 20),
OSD_OUT_ARRAY(OSD_REAL, wDt, 20),
OSD_OUT_ARRAY(OSD_REAL, wDss, 20),
OSD_OUT_ARRAY(OSD_REAL, wDst, 20),
OSD_OUT_ARRAY(OSD_REAL, wDtt, 20)) {
OSD_REAL derivSign = 1.0f;
if ((patchType == OSD_PATCH_DESCRIPTOR_LOOP) ||
(patchType == OSD_PATCH_DESCRIPTOR_GREGORY_TRIANGLE) ||
(patchType == OSD_PATCH_DESCRIPTOR_TRIANGLES)) {
OSD_REAL uv[2] = OSD_ARRAY_2(OSD_REAL, s, t);
OsdPatchParamNormalizeTriangle(param, uv);
s = uv[0];
t = uv[1];
if (OsdPatchParamIsTriangleRotated(param)) {
derivSign = -1.0f;
}
} else {
OSD_REAL uv[2] = OSD_ARRAY_2(OSD_REAL, s, t);
OsdPatchParamNormalize(param, uv);
s = uv[0];
t = uv[1];
}
int nPoints = OsdEvaluatePatchBasisNormalized(
patchType, param, s, t, wP, wDs, wDt, wDss, wDst, wDtt);
if (OSD_OPTIONAL(wDs && wDt)) {
OSD_REAL d1Scale =
derivSign * OSD_REAL_CAST(1 << OsdPatchParamGetDepth(param));
for (int i = 0; i < nPoints; ++i) {
wDs[i] *= d1Scale;
wDt[i] *= d1Scale;
}
if (OSD_OPTIONAL(wDss && wDst && wDtt)) {
OSD_REAL d2Scale = derivSign * d1Scale * d1Scale;
for (int i = 0; i < nPoints; ++i) {
wDss[i] *= d2Scale;
wDst[i] *= d2Scale;
wDtt[i] *= d2Scale;
}
}
}
return nPoints;
}
#endif /* OPENSUBDIV3_OSD_PATCH_BASIS_COMMON_EVAL_H */

422
opensubdiv/osd/patchBasisCommonTypes.h Normal file
View File

@ -0,0 +1,422 @@
//
// Copyright 2018 Pixar
//
// Licensed under the Apache License, Version 2.0 (the "Apache License")
// with the following modification; you may not use this file except in
// compliance with the Apache License and the following modification to it:
// Section 6. Trademarks. is deleted and replaced with:
//
// 6. Trademarks. This License does not grant permission to use the trade
// names, trademarks, service marks, or product names of the Licensor
// and its affiliates, except as required to comply with Section 4(c) of
// the License and to reproduce the content of the NOTICE file.
//
// You may obtain a copy of the Apache License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the Apache License with the above modification is
// distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
// KIND, either express or implied. See the Apache License for the specific
// language governing permissions and limitations under the Apache License.
//
#ifndef OPENSUBDIV3_OSD_PATCH_BASIS_COMMON_TYPES_H
#define OPENSUBDIV3_OSD_PATCH_BASIS_COMMON_TYPES_H
#if defined(OSD_PATCH_BASIS_GLSL)
#define OSD_FUNCTION_STORAGE_CLASS
#define OSD_DATA_STORAGE_CLASS
#define OSD_REAL float
#define OSD_REAL_CAST float
#define OSD_ARG_ARRAY_BOUND_OPTIONAL false
#define OSD_OPTIONAL(a) true
#define OSD_OPTIONAL_INIT(a,b) b
#define OSD_IN_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_OUT_ARRAY(elementType, identifier, arraySize) \
out elementType identifier[arraySize]
#define OSD_INOUT_ARRAY(elementType, identifier, arraySize) \
inout elementType identifier[arraySize]
#define OSD_ARRAY_2(elementType,a0,a1) \
elementType[](a0,a1)
#define OSD_ARRAY_3(elementType,a0,a1,a2) \
elementType[](a0,a1,a2)
#define OSD_ARRAY_4(elementType,a0,a1,a2,a3) \
elementType[](a0,a1,a2,a3)
#define OSD_ARRAY_6(elementType,a0,a1,a2,a3,a4,a5) \
elementType[](a0,a1,a2,a3,a4,a5)
#define OSD_ARRAY_8(elementType,a0,a1,a2,a3,a4,a5,a6,a7) \
elementType[](a0,a1,a2,a3,a4,a5,a6,a7)
#define OSD_ARRAY_9(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8) \
elementType[](a0,a1,a2,a3,a4,a5,a6,a7,a8)
#define OSD_ARRAY_12(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11) \
elementType[](a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11)
#elif defined(OSD_PATCH_BASIS_HLSL)
#define OSD_FUNCTION_STORAGE_CLASS
#define OSD_DATA_STORAGE_CLASS
#define OSD_REAL float
#define OSD_REAL_CAST float
#define OSD_ARG_ARRAY_BOUND_OPTIONAL false
#define OSD_OPTIONAL(a) true
#define OSD_OPTIONAL_INIT(a,b) b
#define OSD_IN_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_OUT_ARRAY(elementType, identifier, arraySize) \
out elementType identifier[arraySize]
#define OSD_INOUT_ARRAY(elementType, identifier, arraySize) \
inout elementType identifier[arraySize]
#define OSD_ARRAY_2(elementType,a0,a1) \
{a0,a1}
#define OSD_ARRAY_3(elementType,a0,a1,a2) \
{a0,a1,a2}
#define OSD_ARRAY_4(elementType,a0,a1,a2,a3) \
{a0,a1,a2,a3}
#define OSD_ARRAY_6(elementType,a0,a1,a2,a3,a4,a5) \
{a0,a1,a2,a3,a4,a5}
#define OSD_ARRAY_8(elementType,a0,a1,a2,a3,a4,a5,a6,a7) \
{a0,a1,a2,a3,a4,a5,a6,a7}
#define OSD_ARRAY_9(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8) \
{a0,a1,a2,a3,a4,a5,a6,a7,a8}
#define OSD_ARRAY_12(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11) \
{a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11}
#elif defined(OSD_PATCH_BASIS_CUDA)
#define OSD_FUNCTION_STORAGE_CLASS __device__
#define OSD_DATA_STORAGE_CLASS
#define OSD_REAL float
#define OSD_REAL_CAST float
#define OSD_OPTIONAL(a) true
#define OSD_OPTIONAL_INIT(a,b) b
#define OSD_ARRAY_ARG_BOUND_OPTIONAL false
#define OSD_IN_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_OUT_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_INOUT_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_ARRAY_2(elementType,a0,a1) \
{a0,a1}
#define OSD_ARRAY_3(elementType,a0,a1,a2) \
{a0,a1,a2}
#define OSD_ARRAY_4(elementType,a0,a1,a2,a3) \
{a0,a1,a2,a3}
#define OSD_ARRAY_6(elementType,a0,a1,a2,a3,a4,a5) \
{a0,a1,a2,a3,a4,a5}
#define OSD_ARRAY_8(elementType,a0,a1,a2,a3,a4,a5,a6,a7) \
{a0,a1,a2,a3,a4,a5,a6,a7}
#define OSD_ARRAY_9(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8) \
{a0,a1,a2,a3,a4,a5,a6,a7,a8}
#define OSD_ARRAY_12(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11) \
{a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11}
#elif defined(OSD_PATCH_BASIS_OPENCL)
#define OSD_FUNCTION_STORAGE_CLASS static
#define OSD_DATA_STORAGE_CLASS
#define OSD_REAL float
#define OSD_REAL_CAST convert_float
#define OSD_OPTIONAL(a) true
#define OSD_OPTIONAL_INIT(a,b) b
#define OSD_ARRAY_ARG_BOUND_OPTIONAL false
#define OSD_IN_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_OUT_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_INOUT_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_ARRAY_2(elementType,a0,a1) \
{a0,a1}
#define OSD_ARRAY_3(elementType,a0,a1,a2) \
{a0,a1,a2}
#define OSD_ARRAY_4(elementType,a0,a1,a2,a3) \
{a0,a1,a2,a3}
#define OSD_ARRAY_6(elementType,a0,a1,a2,a3,a4,a5) \
{a0,a1,a2,a3,a4,a5}
#define OSD_ARRAY_8(elementType,a0,a1,a2,a3,a4,a5,a6,a7) \
{a0,a1,a2,a3,a4,a5,a6,a7}
#define OSD_ARRAY_9(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8) \
{a0,a1,a2,a3,a4,a5,a6,a7,a8}
#define OSD_ARRAY_12(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11) \
{a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11}
#elif defined(OSD_PATCH_BASIS_METAL)
#define OSD_FUNCTION_STORAGE_CLASS
#define OSD_DATA_STORAGE_CLASS
#define OSD_REAL float
#define OSD_REAL_CAST float
#define OSD_OPTIONAL(a) true
#define OSD_OPTIONAL_INIT(a,b) b
#define OSD_ARRAY_ARG_BOUND_OPTIONAL false
#define OSD_IN_ARRAY(elementType, identifier, arraySize) \
thread elementType* identifier
#define OSD_OUT_ARRAY(elementType, identifier, arraySize) \
thread elementType* identifier
#define OSD_INOUT_ARRAY(elementType, identifier, arraySize) \
thread elementType* identifier
#define OSD_ARRAY_2(elementType,a0,a1) \
{a0,a1}
#define OSD_ARRAY_3(elementType,a0,a1,a2) \
{a0,a1,a2}
#define OSD_ARRAY_4(elementType,a0,a1,a2,a3) \
{a0,a1,a2,a3}
#define OSD_ARRAY_6(elementType,a0,a1,a2,a3,a4,a5) \
{a0,a1,a2,a3,a4,a5}
#define OSD_ARRAY_8(elementType,a0,a1,a2,a3,a4,a5,a6,a7) \
{a0,a1,a2,a3,a4,a5,a6,a7}
#define OSD_ARRAY_9(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8) \
{a0,a1,a2,a3,a4,a5,a6,a7,a8}
#define OSD_ARRAY_12(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11) \
{a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11}
#else
#define OSD_FUNCTION_STORAGE_CLASS static inline
#define OSD_DATA_STORAGE_CLASS static
#define OSD_REAL float
#define OSD_REAL_CAST float
#define OSD_OPTIONAL(a) (a)
#define OSD_OPTIONAL_INIT(a,b) (a ? b : 0)
#define OSD_ARRAY_ARG_BOUND_OPTIONAL true
#define OSD_IN_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_OUT_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_INOUT_ARRAY(elementType, identifier, arraySize) \
elementType identifier[arraySize]
#define OSD_ARRAY_2(elementType,a0,a1) \
{a0,a1}
#define OSD_ARRAY_3(elementType,a0,a1,a2) \
{a0,a1,a2}
#define OSD_ARRAY_4(elementType,a0,a1,a2,a3) \
{a0,a1,a2,a3}
#define OSD_ARRAY_6(elementType,a0,a1,a2,a3,a4,a5) \
{a0,a1,a2,a3,a4,a5}
#define OSD_ARRAY_8(elementType,a0,a1,a2,a3,a4,a5,a6,a7) \
{a0,a1,a2,a3,a4,a5,a6,a7}
#define OSD_ARRAY_9(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8) \
{a0,a1,a2,a3,a4,a5,a6,a7,a8}
#define OSD_ARRAY_12(elementType,a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11) \
{a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11}
#endif
#if defined(OSD_PATCH_BASIS_OPENCL)
// OpenCL binding uses typedef to provide the required "struct" type specifier.
typedef struct OsdPatchParam OsdPatchParam;
typedef struct OsdPatchArray OsdPatchArray;
typedef struct OsdPatchCoord OsdPatchCoord;
#endif
// Osd reflection of Far::PatchDescriptor
#define OSD_PATCH_DESCRIPTOR_QUADS 3
#define OSD_PATCH_DESCRIPTOR_TRIANGLES 4
#define OSD_PATCH_DESCRIPTOR_LOOP 5
#define OSD_PATCH_DESCRIPTOR_REGULAR 6
#define OSD_PATCH_DESCRIPTOR_GREGORY_BASIS 9
#define OSD_PATCH_DESCRIPTOR_GREGORY_TRIANGLE 10
// Osd reflection of Osd::PatchCoord
struct OsdPatchCoord {
int arrayIndex;
int patchIndex;
int vertIndex;
float s;
float t;
};
OSD_FUNCTION_STORAGE_CLASS
OsdPatchCoord
OsdPatchCoordInit(
int arrayIndex, int patchIndex, int vertIndex, float s, float t)
{
OsdPatchCoord coord;
coord.arrayIndex = arrayIndex;
coord.patchIndex = patchIndex;
coord.vertIndex = vertIndex;
coord.s = s;
coord.t = t;
return coord;
}
// Osd reflection of Osd::PatchArray
struct OsdPatchArray {
int regDesc;
int desc;
int numPatches;
int indexBase;
int stride;
int primitiveIdBase;
};
OSD_FUNCTION_STORAGE_CLASS
OsdPatchArray
OsdPatchArrayInit(
int regDesc, int desc,
int numPatches, int indexBase, int stride, int primitiveIdBase)
{
OsdPatchArray array;
array.regDesc = regDesc;
array.desc = desc;
array.numPatches = numPatches;
array.indexBase = indexBase;
array.stride = stride;
array.primitiveIdBase = primitiveIdBase;
return array;
}
// Osd reflection of Osd::PatchParam
struct OsdPatchParam {
int field0;
int field1;
float sharpness;
};
OSD_FUNCTION_STORAGE_CLASS
OsdPatchParam
OsdPatchParamInit(int field0, int field1, float sharpness)
{
OsdPatchParam param;
param.field0 = field0;
param.field1 = field1;
param.sharpness = sharpness;
return param;
}
OSD_FUNCTION_STORAGE_CLASS
int
OsdPatchParamGetFaceId(OsdPatchParam param)
{
return (param.field0 & 0xfffffff);
}
OSD_FUNCTION_STORAGE_CLASS
int
OsdPatchParamGetU(OsdPatchParam param)
{
return ((param.field1 >> 22) & 0x3ff);
}
OSD_FUNCTION_STORAGE_CLASS
int
OsdPatchParamGetV(OsdPatchParam param)
{
return ((param.field1 >> 12) & 0x3ff);
}
OSD_FUNCTION_STORAGE_CLASS
int
OsdPatchParamGetTransition(OsdPatchParam param)
{
return ((param.field0 >> 28) & 0xf);
}
OSD_FUNCTION_STORAGE_CLASS
int
OsdPatchParamGetBoundary(OsdPatchParam param)
{
return ((param.field1 >> 7) & 0x1f);
}
OSD_FUNCTION_STORAGE_CLASS
int
OsdPatchParamGetNonQuadRoot(OsdPatchParam param)
{
return ((param.field1 >> 4) & 0x1);
}
OSD_FUNCTION_STORAGE_CLASS
int
OsdPatchParamGetDepth(OsdPatchParam param)
{
return (param.field1 & 0xf);
}
OSD_FUNCTION_STORAGE_CLASS
OSD_REAL
OsdPatchParamGetParamFraction(OsdPatchParam param)
{
return 1.0f / OSD_REAL_CAST(1 <<
(OsdPatchParamGetDepth(param) - OsdPatchParamGetNonQuadRoot(param)));
}
OSD_FUNCTION_STORAGE_CLASS
bool
OsdPatchParamIsRegular(OsdPatchParam param)
{
return (((param.field1 >> 5) & 0x1) != 0);
}
OSD_FUNCTION_STORAGE_CLASS
bool
OsdPatchParamIsTriangleRotated(OsdPatchParam param)
{
return ((OsdPatchParamGetU(param) + OsdPatchParamGetV(param)) >=
(1 << OsdPatchParamGetDepth(param)));
}
OSD_FUNCTION_STORAGE_CLASS
void
OsdPatchParamNormalize(
OsdPatchParam param,
OSD_INOUT_ARRAY(OSD_REAL, uv, 2))
{
OSD_REAL fracInv = 1.0f / OsdPatchParamGetParamFraction(param);
uv[0] = uv[0] * fracInv - OSD_REAL_CAST(OsdPatchParamGetU(param));
uv[1] = uv[1] * fracInv - OSD_REAL_CAST(OsdPatchParamGetV(param));
}
OSD_FUNCTION_STORAGE_CLASS
void
OsdPatchParamUnnormalize(
OsdPatchParam param,
OSD_INOUT_ARRAY(OSD_REAL, uv, 2))
{
OSD_REAL frac = OsdPatchParamGetParamFraction(param);
uv[0] = (uv[0] + OSD_REAL_CAST(OsdPatchParamGetU(param))) * frac;
uv[1] = (uv[1] + OSD_REAL_CAST(OsdPatchParamGetV(param))) * frac;
}
OSD_FUNCTION_STORAGE_CLASS
void
OsdPatchParamNormalizeTriangle(
OsdPatchParam param,
OSD_INOUT_ARRAY(OSD_REAL, uv, 2))
{
if (OsdPatchParamIsTriangleRotated(param)) {
OSD_REAL fracInv = 1.0f / OsdPatchParamGetParamFraction(param);
int depthFactor = 1 << OsdPatchParamGetDepth(param);
uv[0] = OSD_REAL_CAST(depthFactor - OsdPatchParamGetU(param)) - (uv[0] * fracInv);
uv[1] = OSD_REAL_CAST(depthFactor - OsdPatchParamGetV(param)) - (uv[1] * fracInv);
} else {
OsdPatchParamNormalize(param, uv);
}
}
OSD_FUNCTION_STORAGE_CLASS
void
OsdPatchParamUnnormalizeTriangle(
OsdPatchParam param,
OSD_INOUT_ARRAY(OSD_REAL, uv, 2))
{
if (OsdPatchParamIsTriangleRotated(param)) {
OSD_REAL frac = OsdPatchParamGetParamFraction(param);
int depthFactor = 1 << OsdPatchParamGetDepth(param);
uv[0] = (OSD_REAL_CAST(depthFactor - OsdPatchParamGetU(param)) - uv[0]) * frac;
uv[1] = (OSD_REAL_CAST(depthFactor - OsdPatchParamGetV(param)) - uv[1]) * frac;
} else {
OsdPatchParamUnnormalize(param, uv);
}
}
#endif /* OPENSUBDIV3_OSD_PATCH_BASIS_COMMON_TYPES_H */

66
opensubdiv/osd/tbbKernel.cpp
View File

@ -26,7 +26,9 @@
#include "../osd/tbbKernel.h"
#include "../osd/types.h"
#include "../osd/bufferDescriptor.h"
#include "../far/patchBasis.h"
#include "../osd/patchBasisCommonTypes.h"
#include "../osd/patchBasisCommon.h"
#include "../osd/patchBasisCommonEval.h"
#include <cassert>
#include <cstdlib>
@ -382,7 +384,7 @@ public:
}
void compute(tbb::blocked_range<int> const &r) const {
float wP[20], wDu[20], wDv[20];
float wP[20];
BufferAdapter<const float> srcT(_src + _srcDesc.offset,
_srcDesc.length,
_srcDesc.stride);
@ -396,23 +398,25 @@ public:
PatchCoord const &coord = _patchCoords[i];
PatchArray const &array = _patchArrayBuffer[coord.handle.arrayIndex];
Far::PatchParam const & param =
Osd::PatchParam const & paramStruct =
_patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
? Far::PatchDescriptor::REGULAR
: array.GetPatchType();
OsdPatchParam param = OsdPatchParamInit(
paramStruct.field0, paramStruct.field1, paramStruct.sharpness);
int numControlVertices = Far::internal::EvaluatePatchBasis(patchType,
param, coord.s, coord.t, wP, wDu, wDv);
int patchType = OsdPatchParamIsRegular(param)
? array.GetPatchTypeRegular()
: array.GetPatchTypeIrregular();
int indexStride = Far::PatchDescriptor(array.GetPatchType()).GetNumControlVertices();
int indexBase = array.GetIndexBase() + indexStride *
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, 0, 0, 0, 0, 0);
int indexBase = array.GetIndexBase() + array.GetStride() *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
const int *cvs = &_patchIndexBuffer[indexBase];
dstT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
}
++dstT;
@ -441,17 +445,19 @@ public:
PatchCoord const &coord = _patchCoords[i];
PatchArray const &array = _patchArrayBuffer[coord.handle.arrayIndex];
Far::PatchParam const & param =
Osd::PatchParam const & paramStruct =
_patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
? Far::PatchDescriptor::REGULAR
: array.GetPatchType();
OsdPatchParam param = OsdPatchParamInit(
paramStruct.field0, paramStruct.field1, paramStruct.sharpness);
int numControlVertices = Far::internal::EvaluatePatchBasis(patchType,
param, coord.s, coord.t, wP, wDu, wDv);
int patchType = OsdPatchParamIsRegular(param)
? array.GetPatchTypeRegular()
: array.GetPatchTypeIrregular();
int indexStride = Far::PatchDescriptor(array.GetPatchType()).GetNumControlVertices();
int indexBase = array.GetIndexBase() + indexStride *
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, wDu, wDv, 0, 0, 0);
int indexBase = array.GetIndexBase() + array.GetStride() *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
const int *cvs = &_patchIndexBuffer[indexBase];
@ -459,7 +465,7 @@ public:
dstT.Clear();
dstDuT.Clear();
dstDvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
dstDuT.AddWithWeight(srcT[cvs[j]], wDu[j]);
dstDvT.AddWithWeight(srcT[cvs[j]], wDv[j]);
@ -504,17 +510,19 @@ public:
PatchCoord const &coord = _patchCoords[i];
PatchArray const &array = _patchArrayBuffer[coord.handle.arrayIndex];
Far::PatchParam const & param =
Osd::PatchParam const & paramStruct =
_patchParamBuffer[coord.handle.patchIndex];
int patchType = param.IsRegular()
? Far::PatchDescriptor::REGULAR
: array.GetPatchType();
OsdPatchParam param = OsdPatchParamInit(
paramStruct.field0, paramStruct.field1, paramStruct.sharpness);
int numControlVertices = Far::internal::EvaluatePatchBasis(patchType,
param, coord.s, coord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
int patchType = OsdPatchParamIsRegular(param)
? array.GetPatchTypeRegular()
: array.GetPatchTypeIrregular();
int indexStride = Far::PatchDescriptor(array.GetPatchType()).GetNumControlVertices();
int indexBase = array.GetIndexBase() + indexStride *
int nPoints = OsdEvaluatePatchBasis(patchType, param,
coord.s, coord.t, wP, wDu, wDv, wDuu, wDuv, wDvv);
int indexBase = array.GetIndexBase() + array.GetStride() *
(coord.handle.patchIndex - array.GetPrimitiveIdBase());
const int *cvs = &_patchIndexBuffer[indexBase];
@ -525,7 +533,7 @@ public:
dstDuuT.Clear();
dstDuvT.Clear();
dstDvvT.Clear();
for (int j = 0; j < numControlVertices; ++j) {
for (int j = 0; j < nPoints; ++j) {
dstT.AddWithWeight(srcT[cvs[j]], wP[j]);
dstDuT.AddWithWeight(srcT[cvs[j]], wDu[j]);
dstDvT.AddWithWeight(srcT[cvs[j]], wDv[j]);