AuroraMiddleware/OpenSubdiv
1
0
Fork 0
mirror of https://github.com/PixarAnimationStudios/OpenSubdiv synced 2025-01-09 16:20:10 +00:00
Code Issues Releases Wiki Activity
a492af45e3
OpenSubdiv/opensubdiv/osd/clKernel.cl

250 lines
8.2 KiB
Common Lisp
Raw Normal View History

Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
//
Amending Apache license language & file headers. New text: Copyright 2013 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.
2013-09-26 19:04:57 +00:00
// Copyright 2013 Pixar
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
//
Amending Apache license language & file headers. New text: Copyright 2013 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.
2013-09-26 19:04:57 +00:00
// 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:
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
//
Amending Apache license language & file headers. New text: Copyright 2013 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.
2013-09-26 19:04:57 +00:00
// 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.
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
//
Amending Apache license language & file headers. New text: Copyright 2013 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.
2013-09-26 19:04:57 +00:00
// You may obtain a copy of the Apache License at
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
//
Amending Apache license language & file headers. New text: Copyright 2013 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.
2013-09-26 19:04:57 +00:00
// http://www.apache.org/licenses/LICENSE-2.0
Updating EULA
2013-07-18 21:19:50 +00:00
//
Amending Apache license language & file headers. New text: Copyright 2013 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.
2013-09-26 19:04:57 +00:00
// 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.
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
//
Landing 3.0.0.alpha Sync'ing the 'dev' branch with the 'feature_3.0dev' branch at commit 68c6d11fc36761ae1a5e6cdc3457be16f2e9704a The branch 'feature_3.0dev' is now locked and preserved for historical purposes.
2014-09-05 22:07:46 +00:00
struct Vertex {
float v[LENGTH];
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
};
Landing 3.0.0.alpha Sync'ing the 'dev' branch with the 'feature_3.0dev' branch at commit 68c6d11fc36761ae1a5e6cdc3457be16f2e9704a The branch 'feature_3.0dev' is now locked and preserved for historical purposes.
2014-09-05 22:07:46 +00:00
static void clear(struct Vertex *vertex) {
for (int i = 0; i < LENGTH; i++) {
vertex->v[i] = 0.0f;
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
}
}
Interleaved buffer support in OsdCompute. Removed OsdVertexDescriptor and replaced with OsdVertexBufferDescriptor. All kernels take offset/length/stride to apply subdivision partially in each vertex elements. Also the offset can be used for client-based VBO aggregation, without modifying index buffers. This is useful for topology sharing, in conjunction with glDrawElementsBaseVertex etc. However, gregory patch shader fetches vertex buffer via texture buffer, which index should also be offsetted too. Although gl_BaseVertexARB extension should be able to do that job, it's a relatively new extension. So we use OsdBaseVertex() call to mitigate the compatibility issue as clients can provide it in their way at least for the time being.
2014-05-09 00:20:54 +00:00
static void addWithWeight(struct Vertex *dst,
__global float *srcOrigin,
int index, float weight) {
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
osd refactor: split the vertex buffer argument into src and dst Changing all device kernels to take two buffer identifiers for source and destination separately. This change is an intermediate step toward upcoming context/controller refactoring. Previously we have a limitation that the source and destination vertex buffer has to be a single buffer, since the subdivision kernels are iteratively applied by level. With stencil tables, we don't have such a limitation any more, so we may want to apply stencils from seprate source buffer to another. To specifiy the output location within the destination buffer, we can use VertexBufferDescriptor.offset. This allows us not only configuring arbitrary batching scheme, but also relaxing the limitation that source and destination buffers are in same interleaved layout. For examples, we could include derivatives only in the destination buffer, which doesn't need to be allocated in the source buffer.
2015-05-07 23:11:00 +00:00
__global float *src = srcOrigin + index * SRC_STRIDE;
Landing 3.0.0.alpha Sync'ing the 'dev' branch with the 'feature_3.0dev' branch at commit 68c6d11fc36761ae1a5e6cdc3457be16f2e9704a The branch 'feature_3.0dev' is now locked and preserved for historical purposes.
2014-09-05 22:07:46 +00:00
for (int i = 0; i < LENGTH; ++i) {
Interleaved buffer support in OsdCompute. Removed OsdVertexDescriptor and replaced with OsdVertexBufferDescriptor. All kernels take offset/length/stride to apply subdivision partially in each vertex elements. Also the offset can be used for client-based VBO aggregation, without modifying index buffers. This is useful for topology sharing, in conjunction with glDrawElementsBaseVertex etc. However, gregory patch shader fetches vertex buffer via texture buffer, which index should also be offsetted too. Although gl_BaseVertexARB extension should be able to do that job, it's a relatively new extension. So we use OsdBaseVertex() call to mitigate the compatibility issue as clients can provide it in their way at least for the time being.
2014-05-09 00:20:54 +00:00
dst->v[i] += src[i] * weight;
}
}
static void writeVertex(__global float *dstOrigin,
int index,
struct Vertex *src) {
osd refactor: split the vertex buffer argument into src and dst Changing all device kernels to take two buffer identifiers for source and destination separately. This change is an intermediate step toward upcoming context/controller refactoring. Previously we have a limitation that the source and destination vertex buffer has to be a single buffer, since the subdivision kernels are iteratively applied by level. With stencil tables, we don't have such a limitation any more, so we may want to apply stencils from seprate source buffer to another. To specifiy the output location within the destination buffer, we can use VertexBufferDescriptor.offset. This allows us not only configuring arbitrary batching scheme, but also relaxing the limitation that source and destination buffers are in same interleaved layout. For examples, we could include derivatives only in the destination buffer, which doesn't need to be allocated in the source buffer.
2015-05-07 23:11:00 +00:00
__global float *dst = dstOrigin + index * DST_STRIDE;
Landing 3.0.0.alpha Sync'ing the 'dev' branch with the 'feature_3.0dev' branch at commit 68c6d11fc36761ae1a5e6cdc3457be16f2e9704a The branch 'feature_3.0dev' is now locked and preserved for historical purposes.
2014-09-05 22:07:46 +00:00
for (int i = 0; i < LENGTH; ++i) {
Interleaved buffer support in OsdCompute. Removed OsdVertexDescriptor and replaced with OsdVertexBufferDescriptor. All kernels take offset/length/stride to apply subdivision partially in each vertex elements. Also the offset can be used for client-based VBO aggregation, without modifying index buffers. This is useful for topology sharing, in conjunction with glDrawElementsBaseVertex etc. However, gregory patch shader fetches vertex buffer via texture buffer, which index should also be offsetted too. Although gl_BaseVertexARB extension should be able to do that job, it's a relatively new extension. So we use OsdBaseVertex() call to mitigate the compatibility issue as clients can provide it in their way at least for the time being.
2014-05-09 00:20:54 +00:00
dst[i] = src->v[i];
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
}
}
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
static void writeVertexStride(__global float *dstOrigin,
int index,
struct Vertex *src,
int stride) {
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
__global float *dst = dstOrigin + index * stride;
for (int i = 0; i < LENGTH; ++i) {
dst[i] = src->v[i];
}
}
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
__kernel void computeStencils(
__global float * src, int srcOffset,
__global float * dst, int dstOffset,
__global int * sizes,
__global int * offsets,
__global int * indices,
__global float * weights,
int batchStart, int batchEnd) {
Add OpenCL kernel and dispatcher. not to be compiled yet.
2012-06-16 00:52:36 +00:00
Landing 3.0.0.alpha Sync'ing the 'dev' branch with the 'feature_3.0dev' branch at commit 68c6d11fc36761ae1a5e6cdc3457be16f2e9704a The branch 'feature_3.0dev' is now locked and preserved for historical purposes.
2014-09-05 22:07:46 +00:00
int current = get_global_id(0) + batchStart;
Added the CATMARK_QUAD_FACE_VERTEX and CATMARK_TRI_QUAD_FACE_VERTEX kernels. * The CATMARK_QUAD_FACE_VERTEX kernel calculates the face-vertex for a quadrilateral face. It applies to every face after the first subdivision step, and may be applied for the first subdivision step of a quadrilateral coarse mesh. * The CATMARK_TRI_QUAD_FACE_VERTEX kernel calculates the face-vertex for a triangle or quadrilateral face. It may be applied for the first subdivision step of a coarse mesh composed of triangles and/or quadrilaterals. * Both kernels calculate each face-vertex using four vertex indices (triangles are specified by repeating the third index). Therefore neither kernel uses the F_ITa codex table, and instead the first vertex offset in the F_IT index table is stored in the FarKernelBatch's table offset.
2014-05-13 23:06:58 +00:00
osd refactor: split the vertex buffer argument into src and dst Changing all device kernels to take two buffer identifiers for source and destination separately. This change is an intermediate step toward upcoming context/controller refactoring. Previously we have a limitation that the source and destination vertex buffer has to be a single buffer, since the subdivision kernels are iteratively applied by level. With stencil tables, we don't have such a limitation any more, so we may want to apply stencils from seprate source buffer to another. To specifiy the output location within the destination buffer, we can use VertexBufferDescriptor.offset. This allows us not only configuring arbitrary batching scheme, but also relaxing the limitation that source and destination buffers are in same interleaved layout. For examples, we could include derivatives only in the destination buffer, which doesn't need to be allocated in the source buffer.
2015-05-07 23:11:00 +00:00
if (current>=batchEnd) {
return;
}
struct Vertex v;
clear(&v);
Added the CATMARK_QUAD_FACE_VERTEX and CATMARK_TRI_QUAD_FACE_VERTEX kernels. * The CATMARK_QUAD_FACE_VERTEX kernel calculates the face-vertex for a quadrilateral face. It applies to every face after the first subdivision step, and may be applied for the first subdivision step of a quadrilateral coarse mesh. * The CATMARK_TRI_QUAD_FACE_VERTEX kernel calculates the face-vertex for a triangle or quadrilateral face. It may be applied for the first subdivision step of a coarse mesh composed of triangles and/or quadrilaterals. * Both kernels calculate each face-vertex using four vertex indices (triangles are specified by repeating the third index). Therefore neither kernel uses the F_ITa codex table, and instead the first vertex offset in the F_IT index table is stored in the FarKernelBatch's table offset.
2014-05-13 23:06:58 +00:00
Change stencils sizes from unsigned char to int.
2015-05-19 17:16:56 +00:00
int size = sizes[current],
Landing 3.0.0.alpha Sync'ing the 'dev' branch with the 'feature_3.0dev' branch at commit 68c6d11fc36761ae1a5e6cdc3457be16f2e9704a The branch 'feature_3.0dev' is now locked and preserved for historical purposes.
2014-09-05 22:07:46 +00:00
offset = offsets[current];
Added the CATMARK_QUAD_FACE_VERTEX and CATMARK_TRI_QUAD_FACE_VERTEX kernels. * The CATMARK_QUAD_FACE_VERTEX kernel calculates the face-vertex for a quadrilateral face. It applies to every face after the first subdivision step, and may be applied for the first subdivision step of a quadrilateral coarse mesh. * The CATMARK_TRI_QUAD_FACE_VERTEX kernel calculates the face-vertex for a triangle or quadrilateral face. It may be applied for the first subdivision step of a coarse mesh composed of triangles and/or quadrilaterals. * Both kernels calculate each face-vertex using four vertex indices (triangles are specified by repeating the third index). Therefore neither kernel uses the F_ITa codex table, and instead the first vertex offset in the F_IT index table is stored in the FarKernelBatch's table offset.
2014-05-13 23:06:58 +00:00
osd refactor: split the vertex buffer argument into src and dst Changing all device kernels to take two buffer identifiers for source and destination separately. This change is an intermediate step toward upcoming context/controller refactoring. Previously we have a limitation that the source and destination vertex buffer has to be a single buffer, since the subdivision kernels are iteratively applied by level. With stencil tables, we don't have such a limitation any more, so we may want to apply stencils from seprate source buffer to another. To specifiy the output location within the destination buffer, we can use VertexBufferDescriptor.offset. This allows us not only configuring arbitrary batching scheme, but also relaxing the limitation that source and destination buffers are in same interleaved layout. For examples, we could include derivatives only in the destination buffer, which doesn't need to be allocated in the source buffer.
2015-05-07 23:11:00 +00:00
src += srcOffset;
dst += dstOffset;
Added the CATMARK_QUAD_FACE_VERTEX and CATMARK_TRI_QUAD_FACE_VERTEX kernels. * The CATMARK_QUAD_FACE_VERTEX kernel calculates the face-vertex for a quadrilateral face. It applies to every face after the first subdivision step, and may be applied for the first subdivision step of a quadrilateral coarse mesh. * The CATMARK_TRI_QUAD_FACE_VERTEX kernel calculates the face-vertex for a triangle or quadrilateral face. It may be applied for the first subdivision step of a coarse mesh composed of triangles and/or quadrilaterals. * Both kernels calculate each face-vertex using four vertex indices (triangles are specified by repeating the third index). Therefore neither kernel uses the F_ITa codex table, and instead the first vertex offset in the F_IT index table is stored in the FarKernelBatch's table offset.
2014-05-13 23:06:58 +00:00
Landing 3.0.0.alpha Sync'ing the 'dev' branch with the 'feature_3.0dev' branch at commit 68c6d11fc36761ae1a5e6cdc3457be16f2e9704a The branch 'feature_3.0dev' is now locked and preserved for historical purposes.
2014-09-05 22:07:46 +00:00
for (int i=0; i<size; ++i) {
osd refactor: split the vertex buffer argument into src and dst Changing all device kernels to take two buffer identifiers for source and destination separately. This change is an intermediate step toward upcoming context/controller refactoring. Previously we have a limitation that the source and destination vertex buffer has to be a single buffer, since the subdivision kernels are iteratively applied by level. With stencil tables, we don't have such a limitation any more, so we may want to apply stencils from seprate source buffer to another. To specifiy the output location within the destination buffer, we can use VertexBufferDescriptor.offset. This allows us not only configuring arbitrary batching scheme, but also relaxing the limitation that source and destination buffers are in same interleaved layout. For examples, we could include derivatives only in the destination buffer, which doesn't need to be allocated in the source buffer.
2015-05-07 23:11:00 +00:00
addWithWeight(&v, src, indices[offset+i], weights[offset+i]);
Added the CATMARK_QUAD_FACE_VERTEX and CATMARK_TRI_QUAD_FACE_VERTEX kernels. * The CATMARK_QUAD_FACE_VERTEX kernel calculates the face-vertex for a quadrilateral face. It applies to every face after the first subdivision step, and may be applied for the first subdivision step of a quadrilateral coarse mesh. * The CATMARK_TRI_QUAD_FACE_VERTEX kernel calculates the face-vertex for a triangle or quadrilateral face. It may be applied for the first subdivision step of a coarse mesh composed of triangles and/or quadrilaterals. * Both kernels calculate each face-vertex using four vertex indices (triangles are specified by repeating the third index). Therefore neither kernel uses the F_ITa codex table, and instead the first vertex offset in the F_IT index table is stored in the FarKernelBatch's table offset.
2014-05-13 23:06:58 +00:00
}
osd refactor: split the vertex buffer argument into src and dst Changing all device kernels to take two buffer identifiers for source and destination separately. This change is an intermediate step toward upcoming context/controller refactoring. Previously we have a limitation that the source and destination vertex buffer has to be a single buffer, since the subdivision kernels are iteratively applied by level. With stencil tables, we don't have such a limitation any more, so we may want to apply stencils from seprate source buffer to another. To specifiy the output location within the destination buffer, we can use VertexBufferDescriptor.offset. This allows us not only configuring arbitrary batching scheme, but also relaxing the limitation that source and destination buffers are in same interleaved layout. For examples, we could include derivatives only in the destination buffer, which doesn't need to be allocated in the source buffer.
2015-05-07 23:11:00 +00:00
writeVertex(dst, current, &v);
Release Candidate 1.0 : - [Feature Adaptive GPU Rendering of Catmull-Clark Surfaces](http://research.microsoft.com/en-us/um/people/cloop/tog2012.pdf). - New API architecture : we are planning to lock on to this new framework as the basis for backward compatibility, which we will enforce from Release 1.0 onward. Subsequent releases of OpenSubdiv should not break client code. - DirectX 11 support - and much more...
2012-12-11 01:15:13 +00:00
}
Osd API refactor: EvalStencils and EvalPatches Add EvalStencils and EvalPatches API for most of CPU and GPU evaluators. with this change, Eval API in the osd layer consists of following parts: - Evaluators (Cpu, Omp, Tbb, Cuda, CL, GLXFB, GLCompute, D3D11Compute) implements EvalStencils and EvalPatches(*). Both supports derivatives (not fully implemented though) - Interop vertex buffer classes (optional, same as before) Note that these classes are not necessary to use Evaluators. All evaluators have EvalStencils/Patches which take device-specific buffer objects. For example, GLXFBEvaluator can take GLuint directly for both stencil tables and input primvars. Although using these interop classes makes it easy to integrate osd into relatively simple applications. - device-dependent StencilTable and PatchTable (optional) These are also optional, but can be used simply a substitute of Far::StencilTable and Far::PatchTable for osd evaluators. - PatchArray, PatchCoord, PatchParam They are tiny structs used for GPU based patch evaluation. (*) TODO and known issues: - CLEvaluator and D3D11Evaluator's EvalPatches() have not been implemented. - GPU Gregory patch evaluation has not been implemented in EvalPatches(). - CudaEvaluator::EvalPatches() is very unstable. - All patch evaluation kernels have not been well optimized. - Currently GLXFB kernel doesn't support derivative evaluation. There's a technical difficulty for the multi-stream output.
2015-05-26 04:51:55 +00:00
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
__kernel void computeStencilsDerivatives(
__global float * src, int srcOffset,
__global float * dst, int dstOffset,
__global float * du, int duOffset, int duStride,
__global float * dv, int dvOffset, int dvStride,
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
__global float * duu, int duuOffset, int duuStride,
__global float * duv, int duvOffset, int duvStride,
__global float * dvv, int dvvOffset, int dvvStride,
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
__global int * sizes,
__global int * offsets,
__global int * indices,
__global float * weights,
__global float * duWeights,
__global float * dvWeights,
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
__global float * duuWeights,
__global float * duvWeights,
__global float * dvvWeights,
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
int batchStart, int batchEnd) {
int current = get_global_id(0) + batchStart;
if (current>=batchEnd) {
return;
}
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
struct Vertex v, vdu, vdv, vduu, vduv, vdvv;
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
clear(&v);
clear(&vdu);
clear(&vdv);
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
clear(&vduu);
clear(&vduv);
clear(&vdvv);
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
int size = sizes[current],
offset = offsets[current];
if (src) src += srcOffset;
if (dst) dst += dstOffset;
if (du) du += duOffset;
if (dv) dv += dvOffset;
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
if (duu) duu += duuOffset;
if (duv) duv += duvOffset;
if (dvv) dvv += dvvOffset;
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
for (int i=0; i<size; ++i) {
int ofs = offset + i;
int vid = indices[ofs];
if (weights) addWithWeight( &v, src, vid, weights[ofs]);
if (duWeights) addWithWeight(&vdu, src, vid, duWeights[ofs]);
if (dvWeights) addWithWeight(&vdv, src, vid, dvWeights[ofs]);
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
if (duuWeights) addWithWeight(&vduu, src, vid, duuWeights[ofs]);
if (duvWeights) addWithWeight(&vduv, src, vid, duvWeights[ofs]);
if (dvvWeights) addWithWeight(&vdvv, src, vid, dvvWeights[ofs]);
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
}
if (dst) writeVertex (dst, current, &v);
if (du) writeVertexStride(du, current, &vdu, duStride);
if (dv) writeVertexStride(dv, current, &vdv, dvStride);
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
if (duu) writeVertexStride(duu, current, &vduu, duuStride);
if (duv) writeVertexStride(duv, current, &vduv, duvStride);
if (dvv) writeVertexStride(dvv, current, &vdvv, dvvStride);
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
}
Osd API refactor: EvalStencils and EvalPatches Add EvalStencils and EvalPatches API for most of CPU and GPU evaluators. with this change, Eval API in the osd layer consists of following parts: - Evaluators (Cpu, Omp, Tbb, Cuda, CL, GLXFB, GLCompute, D3D11Compute) implements EvalStencils and EvalPatches(*). Both supports derivatives (not fully implemented though) - Interop vertex buffer classes (optional, same as before) Note that these classes are not necessary to use Evaluators. All evaluators have EvalStencils/Patches which take device-specific buffer objects. For example, GLXFBEvaluator can take GLuint directly for both stencil tables and input primvars. Although using these interop classes makes it easy to integrate osd into relatively simple applications. - device-dependent StencilTable and PatchTable (optional) These are also optional, but can be used simply a substitute of Far::StencilTable and Far::PatchTable for osd evaluators. - PatchArray, PatchCoord, PatchParam They are tiny structs used for GPU based patch evaluation. (*) TODO and known issues: - CLEvaluator and D3D11Evaluator's EvalPatches() have not been implemented. - GPU Gregory patch evaluation has not been implemented in EvalPatches(). - CudaEvaluator::EvalPatches() is very unstable. - All patch evaluation kernels have not been well optimized. - Currently GLXFB kernel doesn't support derivative evaluation. There's a technical difficulty for the multi-stream output.
2015-05-26 04:51:55 +00:00
// ---------------------------------------------------------------------------
__kernel void computePatches(__global float *src, int srcOffset,
__global float *dst, int dstOffset,
Fix OpenCL and CUDA EvalPatches kernels.
2015-05-26 18:13:30 +00:00
__global float *du, int duOffset, int duStride,
__global float *dv, int dvOffset, int dvStride,
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
__global float *duu, int duuOffset, int duuStride,
__global float *duv, int duvOffset, int duvStride,
__global float *dvv, int dvvOffset, int dvvStride,
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
__global struct OsdPatchCoord *patchCoords,
__global struct OsdPatchArray *patchArrayBuffer,
Osd API refactor: EvalStencils and EvalPatches Add EvalStencils and EvalPatches API for most of CPU and GPU evaluators. with this change, Eval API in the osd layer consists of following parts: - Evaluators (Cpu, Omp, Tbb, Cuda, CL, GLXFB, GLCompute, D3D11Compute) implements EvalStencils and EvalPatches(*). Both supports derivatives (not fully implemented though) - Interop vertex buffer classes (optional, same as before) Note that these classes are not necessary to use Evaluators. All evaluators have EvalStencils/Patches which take device-specific buffer objects. For example, GLXFBEvaluator can take GLuint directly for both stencil tables and input primvars. Although using these interop classes makes it easy to integrate osd into relatively simple applications. - device-dependent StencilTable and PatchTable (optional) These are also optional, but can be used simply a substitute of Far::StencilTable and Far::PatchTable for osd evaluators. - PatchArray, PatchCoord, PatchParam They are tiny structs used for GPU based patch evaluation. (*) TODO and known issues: - CLEvaluator and D3D11Evaluator's EvalPatches() have not been implemented. - GPU Gregory patch evaluation has not been implemented in EvalPatches(). - CudaEvaluator::EvalPatches() is very unstable. - All patch evaluation kernels have not been well optimized. - Currently GLXFB kernel doesn't support derivative evaluation. There's a technical difficulty for the multi-stream output.
2015-05-26 04:51:55 +00:00
__global int *patchIndexBuffer,
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
__global struct OsdPatchParam *patchParamBuffer) {
Osd API refactor: EvalStencils and EvalPatches Add EvalStencils and EvalPatches API for most of CPU and GPU evaluators. with this change, Eval API in the osd layer consists of following parts: - Evaluators (Cpu, Omp, Tbb, Cuda, CL, GLXFB, GLCompute, D3D11Compute) implements EvalStencils and EvalPatches(*). Both supports derivatives (not fully implemented though) - Interop vertex buffer classes (optional, same as before) Note that these classes are not necessary to use Evaluators. All evaluators have EvalStencils/Patches which take device-specific buffer objects. For example, GLXFBEvaluator can take GLuint directly for both stencil tables and input primvars. Although using these interop classes makes it easy to integrate osd into relatively simple applications. - device-dependent StencilTable and PatchTable (optional) These are also optional, but can be used simply a substitute of Far::StencilTable and Far::PatchTable for osd evaluators. - PatchArray, PatchCoord, PatchParam They are tiny structs used for GPU based patch evaluation. (*) TODO and known issues: - CLEvaluator and D3D11Evaluator's EvalPatches() have not been implemented. - GPU Gregory patch evaluation has not been implemented in EvalPatches(). - CudaEvaluator::EvalPatches() is very unstable. - All patch evaluation kernels have not been well optimized. - Currently GLXFB kernel doesn't support derivative evaluation. There's a technical difficulty for the multi-stream output.
2015-05-26 04:51:55 +00:00
int current = get_global_id(0);
Fix OpenCL and CUDA EvalPatches kernels.
2015-05-26 18:13:30 +00:00
if (src) src += srcOffset;
if (dst) dst += dstOffset;
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
if (du) du += duOffset;
if (dv) dv += dvOffset;
if (duu) duu += duuOffset;
if (duv) duv += duvOffset;
if (dvv) dvv += dvvOffset;
Osd API refactor: EvalStencils and EvalPatches Add EvalStencils and EvalPatches API for most of CPU and GPU evaluators. with this change, Eval API in the osd layer consists of following parts: - Evaluators (Cpu, Omp, Tbb, Cuda, CL, GLXFB, GLCompute, D3D11Compute) implements EvalStencils and EvalPatches(*). Both supports derivatives (not fully implemented though) - Interop vertex buffer classes (optional, same as before) Note that these classes are not necessary to use Evaluators. All evaluators have EvalStencils/Patches which take device-specific buffer objects. For example, GLXFBEvaluator can take GLuint directly for both stencil tables and input primvars. Although using these interop classes makes it easy to integrate osd into relatively simple applications. - device-dependent StencilTable and PatchTable (optional) These are also optional, but can be used simply a substitute of Far::StencilTable and Far::PatchTable for osd evaluators. - PatchArray, PatchCoord, PatchParam They are tiny structs used for GPU based patch evaluation. (*) TODO and known issues: - CLEvaluator and D3D11Evaluator's EvalPatches() have not been implemented. - GPU Gregory patch evaluation has not been implemented in EvalPatches(). - CudaEvaluator::EvalPatches() is very unstable. - All patch evaluation kernels have not been well optimized. - Currently GLXFB kernel doesn't support derivative evaluation. There's a technical difficulty for the multi-stream output.
2015-05-26 04:51:55 +00:00
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
struct OsdPatchCoord coord = patchCoords[current];
struct OsdPatchArray array = patchArrayBuffer[coord.arrayIndex];
struct OsdPatchParam param = patchParamBuffer[coord.patchIndex];
int patchType = OsdPatchParamIsRegular(param) ? array.regDesc : array.desc;
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);
Osd API refactor: EvalStencils and EvalPatches Add EvalStencils and EvalPatches API for most of CPU and GPU evaluators. with this change, Eval API in the osd layer consists of following parts: - Evaluators (Cpu, Omp, Tbb, Cuda, CL, GLXFB, GLCompute, D3D11Compute) implements EvalStencils and EvalPatches(*). Both supports derivatives (not fully implemented though) - Interop vertex buffer classes (optional, same as before) Note that these classes are not necessary to use Evaluators. All evaluators have EvalStencils/Patches which take device-specific buffer objects. For example, GLXFBEvaluator can take GLuint directly for both stencil tables and input primvars. Although using these interop classes makes it easy to integrate osd into relatively simple applications. - device-dependent StencilTable and PatchTable (optional) These are also optional, but can be used simply a substitute of Far::StencilTable and Far::PatchTable for osd evaluators. - PatchArray, PatchCoord, PatchParam They are tiny structs used for GPU based patch evaluation. (*) TODO and known issues: - CLEvaluator and D3D11Evaluator's EvalPatches() have not been implemented. - GPU Gregory patch evaluation has not been implemented in EvalPatches(). - CudaEvaluator::EvalPatches() is very unstable. - All patch evaluation kernels have not been well optimized. - Currently GLXFB kernel doesn't support derivative evaluation. There's a technical difficulty for the multi-stream output.
2015-05-26 04:51:55 +00:00
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
int indexBase = array.indexBase + array.stride *
Osd Varying and FaceVarying Patch Evaluation Implemented EvalPatchesVarying and EvalPatchesFaceVarying methods for Osd::*Evaluator classes, i.e. cpu, omp, tbb, GLXFB, GLSLCompute, OpenCL, and CUDA. Also, the GPU Kernel implementations have been updated to use the common patchBasis implementation instead of re-implementing methods to compute patch basis weights locally.
2016-09-29 16:56:15 +00:00
(coord.patchIndex - array.primitiveIdBase);
Fix OpenCL and CUDA EvalPatches kernels.
2015-05-26 18:13:30 +00:00
Osd API refactor: EvalStencils and EvalPatches Add EvalStencils and EvalPatches API for most of CPU and GPU evaluators. with this change, Eval API in the osd layer consists of following parts: - Evaluators (Cpu, Omp, Tbb, Cuda, CL, GLXFB, GLCompute, D3D11Compute) implements EvalStencils and EvalPatches(*). Both supports derivatives (not fully implemented though) - Interop vertex buffer classes (optional, same as before) Note that these classes are not necessary to use Evaluators. All evaluators have EvalStencils/Patches which take device-specific buffer objects. For example, GLXFBEvaluator can take GLuint directly for both stencil tables and input primvars. Although using these interop classes makes it easy to integrate osd into relatively simple applications. - device-dependent StencilTable and PatchTable (optional) These are also optional, but can be used simply a substitute of Far::StencilTable and Far::PatchTable for osd evaluators. - PatchArray, PatchCoord, PatchParam They are tiny structs used for GPU based patch evaluation. (*) TODO and known issues: - CLEvaluator and D3D11Evaluator's EvalPatches() have not been implemented. - GPU Gregory patch evaluation has not been implemented in EvalPatches(). - CudaEvaluator::EvalPatches() is very unstable. - All patch evaluation kernels have not been well optimized. - Currently GLXFB kernel doesn't support derivative evaluation. There's a technical difficulty for the multi-stream output.
2015-05-26 04:51:55 +00:00
struct Vertex v;
clear(&v);
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
for (int i = 0; i < nPoints; ++i) {
Osd API refactor: EvalStencils and EvalPatches Add EvalStencils and EvalPatches API for most of CPU and GPU evaluators. with this change, Eval API in the osd layer consists of following parts: - Evaluators (Cpu, Omp, Tbb, Cuda, CL, GLXFB, GLCompute, D3D11Compute) implements EvalStencils and EvalPatches(*). Both supports derivatives (not fully implemented though) - Interop vertex buffer classes (optional, same as before) Note that these classes are not necessary to use Evaluators. All evaluators have EvalStencils/Patches which take device-specific buffer objects. For example, GLXFBEvaluator can take GLuint directly for both stencil tables and input primvars. Although using these interop classes makes it easy to integrate osd into relatively simple applications. - device-dependent StencilTable and PatchTable (optional) These are also optional, but can be used simply a substitute of Far::StencilTable and Far::PatchTable for osd evaluators. - PatchArray, PatchCoord, PatchParam They are tiny structs used for GPU based patch evaluation. (*) TODO and known issues: - CLEvaluator and D3D11Evaluator's EvalPatches() have not been implemented. - GPU Gregory patch evaluation has not been implemented in EvalPatches(). - CudaEvaluator::EvalPatches() is very unstable. - All patch evaluation kernels have not been well optimized. - Currently GLXFB kernel doesn't support derivative evaluation. There's a technical difficulty for the multi-stream output.
2015-05-26 04:51:55 +00:00
int index = patchIndexBuffer[indexBase + i];
addWithWeight(&v, src, index, wP[i]);
}
writeVertex(dst, current, &v);
Fix OpenCL and CUDA EvalPatches kernels.
2015-05-26 18:13:30 +00:00
if (du) {
struct Vertex vdu;
clear(&vdu);
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
for (int i = 0; i < nPoints; ++i) {
Fix OpenCL and CUDA EvalPatches kernels.
2015-05-26 18:13:30 +00:00
int index = patchIndexBuffer[indexBase + i];
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
addWithWeight(&vdu, src, index, wDu[i]);
Fix OpenCL and CUDA EvalPatches kernels.
2015-05-26 18:13:30 +00:00
}
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
writeVertexStride(du, current, &vdu, duStride);
Fix OpenCL and CUDA EvalPatches kernels.
2015-05-26 18:13:30 +00:00
}
if (dv) {
struct Vertex vdv;
clear(&vdv);
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
for (int i = 0; i < nPoints; ++i) {
Fix OpenCL and CUDA EvalPatches kernels.
2015-05-26 18:13:30 +00:00
int index = patchIndexBuffer[indexBase + i];
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
addWithWeight(&vdv, src, index, wDv[i]);
Fix OpenCL and CUDA EvalPatches kernels.
2015-05-26 18:13:30 +00:00
}
Add GPU stencil/patch derivative kernels implements cuda, CL, GLXFB and GLCompute kernels of derivative evaluation for both limit stencils and patches.
2015-05-28 00:23:36 +00:00
writeVertexStride(dv, current, &vdv, dvStride);
Fix OpenCL and CUDA EvalPatches kernels.
2015-05-26 18:13:30 +00:00
}
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
if (duu) {
struct Vertex vduu;
clear(&vduu);
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
for (int i = 0; i < nPoints; ++i) {
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
int index = patchIndexBuffer[indexBase + i];
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
addWithWeight(&vduu, src, index, wDuu[i]);
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
}
writeVertexStride(duu, current, &vduu, duuStride);
}
if (duv) {
struct Vertex vduv;
clear(&vduv);
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
for (int i = 0; i < nPoints; ++i) {
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
int index = patchIndexBuffer[indexBase + i];
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
addWithWeight(&vduv, src, index, wDuv[i]);
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
}
writeVertexStride(duv, current, &vduv, duvStride);
}
if (dvv) {
struct Vertex vdvv;
clear(&vdvv);
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
for (int i = 0; i < nPoints; ++i) {
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
int index = patchIndexBuffer[indexBase + i];
Improved patch basis eval for Osd to match Far 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
2018-10-30 07:37:25 +00:00
addWithWeight(&vdvv, src, index, wDvv[i]);
Added missing Osd Evaluator methods for deriv eval Now that Far::LimitStencilTable and Far::PatchTable support evaluation of 1st and 2nd derivatives the Osd Evaluator API for evaluating stencils and patches has been updated to match.
2017-01-26 22:36:30 +00:00
}
writeVertexStride(dvv, current, &vdvv, dvvStride);
}
Osd API refactor: EvalStencils and EvalPatches Add EvalStencils and EvalPatches API for most of CPU and GPU evaluators. with this change, Eval API in the osd layer consists of following parts: - Evaluators (Cpu, Omp, Tbb, Cuda, CL, GLXFB, GLCompute, D3D11Compute) implements EvalStencils and EvalPatches(*). Both supports derivatives (not fully implemented though) - Interop vertex buffer classes (optional, same as before) Note that these classes are not necessary to use Evaluators. All evaluators have EvalStencils/Patches which take device-specific buffer objects. For example, GLXFBEvaluator can take GLuint directly for both stencil tables and input primvars. Although using these interop classes makes it easy to integrate osd into relatively simple applications. - device-dependent StencilTable and PatchTable (optional) These are also optional, but can be used simply a substitute of Far::StencilTable and Far::PatchTable for osd evaluators. - PatchArray, PatchCoord, PatchParam They are tiny structs used for GPU based patch evaluation. (*) TODO and known issues: - CLEvaluator and D3D11Evaluator's EvalPatches() have not been implemented. - GPU Gregory patch evaluation has not been implemented in EvalPatches(). - CudaEvaluator::EvalPatches() is very unstable. - All patch evaluation kernels have not been well optimized. - Currently GLXFB kernel doesn't support derivative evaluation. There's a technical difficulty for the multi-stream output.
2015-05-26 04:51:55 +00:00
}
Reference in New Issue Copy Permalink