mirror of
https://github.com/PixarAnimationStudios/OpenSubdiv
synced 2024-11-23 04:00:07 +00:00
ff34c514f5
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
321 lines
11 KiB
Metal
321 lines
11 KiB
Metal
#line 0 "osd/mtlComputeKernel.metal"
|
|
|
|
//
|
|
// Copyright 2015 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.
|
|
//
|
|
|
|
#include <metal_stdlib>
|
|
|
|
#ifndef OPENSUBDIV_MTL_COMPUTE_USE_1ST_DERIVATIVES
|
|
#define OPENSUBDIV_MTL_COMPUTE_USE_1ST_DERIVATIVES 0
|
|
#endif
|
|
|
|
#ifndef OPENSUBDIV_MTL_COMPUTE_USE_2ND_DERIVATIVES
|
|
#define OPENSUBDIV_MTL_COMPUTE_USE_2ND_DERIVATIVES 0
|
|
#endif
|
|
|
|
using namespace metal;
|
|
|
|
struct KernelUniformArgs
|
|
{
|
|
int batchStart;
|
|
int batchEnd;
|
|
|
|
int srcOffset;
|
|
int dstOffset;
|
|
|
|
int3 duDesc;
|
|
int3 dvDesc;
|
|
|
|
int3 duuDesc;
|
|
int3 duvDesc;
|
|
int3 dvvDesc;
|
|
};
|
|
|
|
struct Vertex {
|
|
float vertexData[LENGTH];
|
|
};
|
|
|
|
void clear(thread Vertex& v) {
|
|
for (int i = 0; i < LENGTH; ++i) {
|
|
v.vertexData[i] = 0;
|
|
}
|
|
}
|
|
|
|
Vertex readVertex(int index, device const float* vertexBuffer, KernelUniformArgs args) {
|
|
Vertex v;
|
|
int vertexIndex = args.srcOffset + index * SRC_STRIDE;
|
|
for (int i = 0; i < LENGTH; ++i) {
|
|
v.vertexData[i] = vertexBuffer[vertexIndex + i];
|
|
}
|
|
return v;
|
|
}
|
|
|
|
void writeVertex(int index, Vertex v, device float* vertexBuffer, KernelUniformArgs args) {
|
|
int vertexIndex = args.dstOffset + index * DST_STRIDE;
|
|
for (int i = 0; i < LENGTH; ++i) {
|
|
vertexBuffer[vertexIndex + i] = v.vertexData[i];
|
|
}
|
|
}
|
|
|
|
void writeVertexSeparate(int index, Vertex v, device float* dstVertexBuffer, KernelUniformArgs args) {
|
|
int vertexIndex = args.dstOffset + index * DST_STRIDE;
|
|
for (int i = 0; i < LENGTH; ++i) {
|
|
dstVertexBuffer[vertexIndex + i] = v.vertexData[i];
|
|
}
|
|
}
|
|
|
|
void addWithWeight(thread Vertex& v, const Vertex src, float weight) {
|
|
for (int i = 0; i < LENGTH; ++i) {
|
|
v.vertexData[i] += weight * src.vertexData[i];
|
|
}
|
|
}
|
|
|
|
#if OPENSUBDIV_MTL_COMPUTE_USE_1ST_DERIVATIVES
|
|
void writeDu(int index, Vertex du, device float* duDerivativeBuffer, KernelUniformArgs args)
|
|
{
|
|
int duIndex = args.duDesc.x + index * args.duDesc.z;
|
|
for(int i = 0; i < LENGTH; i++) {
|
|
duDerivativeBuffer[duIndex + i] = du.vertexData[i];
|
|
}
|
|
}
|
|
|
|
void writeDv(int index, Vertex dv, device float* dvDerivativeBuffer, KernelUniformArgs args)
|
|
{
|
|
int dvIndex = args.dvDesc.x + index * args.dvDesc.z;
|
|
for(int i = 0; i < LENGTH; i++) {
|
|
dvDerivativeBuffer[dvIndex + i] = dv.vertexData[i];
|
|
}
|
|
}
|
|
#endif
|
|
|
|
#if OPENSUBDIV_MTL_COMPUTE_USE_2ND_DERIVATIVES
|
|
void writeDuu(int index, Vertex duu, device float* duuDerivativeBuffer, KernelUniformArgs args)
|
|
{
|
|
int duuIndex = args.duuDesc.x + index * args.duuDesc.z;
|
|
for(int i = 0; i < LENGTH; i++) {
|
|
duuDerivativeBuffer[duuIndex + i] = duu.vertexData[i];
|
|
}
|
|
}
|
|
|
|
void writeDuv(int index, Vertex duv, device float* duvDerivativeBuffer, KernelUniformArgs args)
|
|
{
|
|
int duvIndex = args.duvDesc.x + index * args.duvDesc.z;
|
|
for(int i = 0; i < LENGTH; i++) {
|
|
duvDerivativeBuffer[duvIndex + i] = duv.vertexData[i];
|
|
}
|
|
}
|
|
|
|
void writeDvv(int index, Vertex dvv, device float* dvvDerivativeBuffer, KernelUniformArgs args)
|
|
{
|
|
int dvvIndex = args.dvvDesc.x + index * args.dvvDesc.z;
|
|
for(int i = 0; i < LENGTH; i++) {
|
|
dvvDerivativeBuffer[dvvIndex + i] = dvv.vertexData[i];
|
|
}
|
|
}
|
|
#endif
|
|
|
|
// ---------------------------------------------------------------------------
|
|
|
|
kernel void eval_stencils(
|
|
uint thread_position_in_grid [[thread_position_in_grid]],
|
|
const device int* sizes [[buffer(SIZES_BUFFER_INDEX)]],
|
|
const device int* offsets [[buffer(OFFSETS_BUFFER_INDEX)]],
|
|
const device int* indices [[buffer(INDICES_BUFFER_INDEX)]],
|
|
const device float* weights [[buffer(WEIGHTS_BUFFER_INDEX)]],
|
|
const device float* srcVertices [[buffer(SRC_VERTEX_BUFFER_INDEX)]],
|
|
device float* dstVertexBuffer [[buffer(DST_VERTEX_BUFFER_INDEX)]],
|
|
#if OPENSUBDIV_MTL_COMPUTE_USE_1ST_DERIVATIVES
|
|
const device float* duWeights [[buffer(DU_WEIGHTS_BUFFER_INDEX)]],
|
|
const device float* dvWeights [[buffer(DV_WEIGHTS_BUFFER_INDEX)]],
|
|
device float* duDerivativeBuffer [[buffer(DU_DERIVATIVE_BUFFER_INDEX)]],
|
|
device float* dvDerivativeBuffer [[buffer(DV_DERIVATIVE_BUFFER_INDEX)]],
|
|
#endif
|
|
#if OPENSUBDIV_MTL_COMPUTE_USE_2ND_DERIVATIVES
|
|
const device float* duuWeights [[buffer(DUU_WEIGHTS_BUFFER_INDEX)]],
|
|
const device float* duvWeights [[buffer(DUV_WEIGHTS_BUFFER_INDEX)]],
|
|
const device float* dvvWeights [[buffer(DVV_WEIGHTS_BUFFER_INDEX)]],
|
|
device float* duuDerivativeBuffer [[buffer(DUU_DERIVATIVE_BUFFER_INDEX)]],
|
|
device float* duvDerivativeBuffer [[buffer(DUV_DERIVATIVE_BUFFER_INDEX)]],
|
|
device float* dvvDerivativeBuffer [[buffer(DVV_DERIVATIVE_BUFFER_INDEX)]],
|
|
#endif
|
|
const constant KernelUniformArgs& args [[buffer(PARAMETER_BUFFER_INDEX)]]
|
|
)
|
|
{
|
|
auto current = thread_position_in_grid + args.batchStart;
|
|
if(current >= (unsigned int)args.batchEnd)
|
|
return;
|
|
|
|
Vertex dst;
|
|
clear(dst);
|
|
|
|
|
|
auto offset = offsets[current];
|
|
auto size = sizes[current];
|
|
|
|
for(auto stencil = 0; stencil < size; stencil++)
|
|
{
|
|
auto vindex = offset + stencil;
|
|
addWithWeight(dst, readVertex(indices[vindex], srcVertices, args), weights[vindex]);
|
|
}
|
|
|
|
writeVertex(current, dst, dstVertexBuffer, args);
|
|
|
|
#if OPENSUBDIV_MTL_COMPUTE_USE_1ST_DERIVATIVES
|
|
Vertex du, dv;
|
|
clear(du);
|
|
clear(dv);
|
|
|
|
|
|
for(auto i = 0; i < size; i++)
|
|
{
|
|
auto src = readVertex(indices[offset + i], srcVertices, args);
|
|
addWithWeight(du, src, duWeights[offset + i]);
|
|
addWithWeight(dv, src, dvWeights[offset + i]);
|
|
}
|
|
|
|
writeDu(current, du, duDerivativeBuffer, args);
|
|
writeDv(current, dv, dvDerivativeBuffer, args);
|
|
#endif
|
|
|
|
#if OPENSUBDIV_MTL_COMPUTE_USE_2ND_DERIVATIVES
|
|
Vertex duu, duv, dvv;
|
|
clear(duu);
|
|
clear(duv);
|
|
clear(dvv);
|
|
|
|
|
|
for(auto i = 0; i < size; i++)
|
|
{
|
|
auto src = readVertex(indices[offset + i], srcVertices, args);
|
|
addWithWeight(duu, src, duuWeights[offset + i]);
|
|
addWithWeight(duv, src, duvWeights[offset + i]);
|
|
addWithWeight(dvv, src, dvvWeights[offset + i]);
|
|
}
|
|
|
|
writeDuu(current, duu, duuDerivativeBuffer, args);
|
|
writeDuv(current, duv, duvDerivativeBuffer, args);
|
|
writeDvv(current, dvv, dvvDerivativeBuffer, args);
|
|
#endif
|
|
}
|
|
|
|
|
|
// ---------------------------------------------------------------------------
|
|
|
|
// PERFORMANCE: stride could be constant, but not as significant as length
|
|
|
|
// ---------------------------------------------------------------------------
|
|
|
|
kernel void eval_patches(
|
|
uint thread_position_in_grid [[thread_position_in_grid]],
|
|
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)]],
|
|
const device float* srcVertexBuffer [[buffer(SRC_VERTEX_BUFFER_INDEX)]],
|
|
device float* dstVertexBuffer [[buffer(DST_VERTEX_BUFFER_INDEX)]],
|
|
#if OPENSUBDIV_MTL_COMPUTE_USE_1ST_DERIVATIVES
|
|
device float* duDerivativeBuffer [[buffer(DU_DERIVATIVE_BUFFER_INDEX)]],
|
|
device float* dvDerivativeBuffer [[buffer(DV_DERIVATIVE_BUFFER_INDEX)]],
|
|
#endif
|
|
#if OPENSUBDIV_MTL_COMPUTE_USE_2ND_DERIVATIVES
|
|
device float* duuDerivativeBuffer [[buffer(DUU_DERIVATIVE_BUFFER_INDEX)]],
|
|
device float* duvDerivativeBuffer [[buffer(DUV_DERIVATIVE_BUFFER_INDEX)]],
|
|
device float* dvvDerivativeBuffer [[buffer(DVV_DERIVATIVE_BUFFER_INDEX)]],
|
|
#endif
|
|
const constant KernelUniformArgs& args [[buffer(PARAMETER_BUFFER_INDEX)]]
|
|
)
|
|
{
|
|
auto current = thread_position_in_grid;
|
|
|
|
// unpack struct (5 ints unaligned)
|
|
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]));
|
|
|
|
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]);
|
|
|
|
OsdPatchParam patchParam = OsdPatchParamInit(patchParams[current*3+0],
|
|
patchParams[current*3+1],
|
|
as_type<float>(patchParams[current*3+2]));
|
|
|
|
int patchType = OsdPatchParamIsRegular(patchParam)
|
|
? patchArray.regDesc : patchArray.desc;
|
|
|
|
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);
|
|
clear(du);
|
|
clear(dv);
|
|
clear(duu);
|
|
clear(duv);
|
|
clear(dvv);
|
|
|
|
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, 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);
|
|
|
|
#if OPENSUBDIV_MTL_COMPUTE_USE_1ST_DERIVATIVES
|
|
if(args.duDesc.y > 0)
|
|
writeDu(current, du, duDerivativeBuffer, args);
|
|
|
|
if(args.dvDesc.y > 0)
|
|
writeDv(current, dv, dvDerivativeBuffer, args);
|
|
#endif
|
|
|
|
#if OPENSUBDIV_MTL_COMPUTE_USE_2ND_DERIVATIVES
|
|
if(args.duuDesc.y > 0)
|
|
writeDuu(current, duu, duuDerivativeBuffer, args);
|
|
|
|
if(args.duvDesc.y > 0)
|
|
writeDuv(current, duv, duvDerivativeBuffer, args);
|
|
|
|
if(args.dvvDesc.y > 0)
|
|
writeDvv(current, dvv, dvvDerivativeBuffer, args);
|
|
#endif
|
|
|
|
|
|
}
|