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

Merge pull request #1021 from davidgyu/osd_patch_eval

Browse Source 
This is great, David. Thanks
This commit is contained in:
Barry Fowler 2018-11-16 17:14:14 -08:00 committed by GitHub
commit d7e0449fb2
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
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]);