AuroraMiddleware/OpenSubdiv
1
0
Fork 0
mirror of https://github.com/PixarAnimationStudios/OpenSubdiv synced 2024-12-01 15:50:07 +00:00
Code Issues Releases Wiki Activity
c4e760e722
OpenSubdiv/opensubdiv/osd/cudaKernel.cu

941 lines
39 KiB
Plaintext
Raw Normal View History

adding CUDA kernels
2012-06-09 21:22:57 +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
adding CUDA kernels
2012-06-09 21:22:57 +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:
adding CUDA kernels
2012-06-09 21:22:57 +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.
adding CUDA kernels
2012-06-09 21:22:57 +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
adding CUDA kernels
2012-06-09 21:22:57 +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
adding CUDA kernels
2012-06-09 21:22:57 +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.
Updating EULA
2013-07-18 21:19:50 +00:00
//
adding CUDA kernels
2012-06-09 21:22:57 +00:00
#include <assert.h>
template<int N> struct DeviceVertex
{
float v[N];
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
__device__ void addWithWeight(const DeviceVertex<N> *src, float weight) {
#pragma unroll
adding CUDA kernels
2012-06-09 21:22:57 +00:00
for(int i = 0; i < N; ++i){
v[i] += src->v[i] * weight;
}
}
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
adding CUDA kernels
2012-06-09 21:22:57 +00:00
__device__ void clear() {
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
#pragma unroll
adding CUDA kernels
2012-06-09 21:22:57 +00:00
for(int i = 0; i < N; ++i){
v[i] = 0.0f;
}
}
};
// Specialize DeviceVarying for N=0 to avoid compile error:
// "flexible array member in otherwise empty struct"
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
template<> struct DeviceVertex<0>
adding CUDA kernels
2012-06-09 21:22:57 +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
__device__ void addWithWeight(const DeviceVertex<0> *src, float weight) {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
__device__ void clear() {
}
};
struct DeviceTable
{
void **tables;
int *F0_IT;
int *F0_ITa;
int *E0_IT;
int *V0_IT;
int *V0_ITa;
float *E0_S;
float *V0_S;
};
__device__ void clear(float *dst, int count)
{
for(int i = 0; i < count; ++i) dst[i] = 0;
}
__device__ void addWithWeight(float *dst, float *src, float weight, int count)
{
for(int i = 0; i < count; ++i) dst[i] += src[i] * weight;
}
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
template <int NUM_VERTEX_ELEMENTS, int NUM_VARYING_ELEMENTS> __global__ void
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
computeFace(float *fVertex, float *fVaryings, int *F0_IT, int *F0_ITa, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> *vertex = (DeviceVertex<NUM_VERTEX_ELEMENTS>*)fVertex;
DeviceVertex<NUM_VARYING_ELEMENTS> *varyings = (DeviceVertex<NUM_VARYING_ELEMENTS>*)fVaryings;
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i += blockDim.x * gridDim.x) {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int h = F0_ITa[2*i];
int n = F0_ITa[2*i+1];
float weight = 1.0f/n;
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> dst;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dst.clear();
if(NUM_VARYING_ELEMENTS > 0){
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
DeviceVertex<NUM_VARYING_ELEMENTS> dstVarying;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dstVarying.clear();
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
for(int j=0; j<n; ++j){
int index = F0_IT[h+j];
dst.addWithWeight(&vertex[index], weight);
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
dstVarying.addWithWeight(&varyings[index], weight);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
vertex[offset + i - tableOffset] = dst;
varyings[offset + i - tableOffset] = dstVarying;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}else{
for(int j=0; j<n; ++j){
int index = F0_IT[h+j];
dst.addWithWeight(&vertex[index], weight);
}
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
vertex[offset + i - tableOffset] = dst;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
}
}
__global__ void
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
computeFace(float *fVertex, float *fVarying,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *F0_IT, int *F0_ITa, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +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
for (int i = start + tableOffset +threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i += blockDim.x * gridDim.x){
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int h = F0_ITa[2*i];
int n = F0_ITa[2*i+1];
float weight = 1.0f/n;
// XXX: can we use local stack like alloca?
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
float *dstVertex = fVertex + (i+offset-tableOffset)*vertexStride;
clear(dstVertex, vertexLength);
float *dstVarying = fVarying + (i+offset-tableOffset)*varyingStride;
clear(dstVarying, varyingLength);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
for(int j=0; j<n; ++j){
int index = F0_IT[h+j];
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
addWithWeight(dstVertex, fVertex + index*vertexStride, weight, vertexLength);
addWithWeight(dstVarying, fVarying + index*varyingStride, weight, varyingLength);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
}
}
Implemented CATMARK_QUAD_FACE_VERTEX and CATMARK_TRI_QUAD_FACE_VERTEX kernels for CUDA, GLSL Transform Feedback, and GLSL Compute platforms.
2014-05-28 23:46:38 +00:00
template <int NUM_VERTEX_ELEMENTS, int NUM_VARYING_ELEMENTS> __global__ void
computeQuadFace(float *fVertex, float *fVaryings, int *F0_IT, int offset, int tableOffset, int start, int end)
{
DeviceVertex<NUM_VERTEX_ELEMENTS> *vertex = (DeviceVertex<NUM_VERTEX_ELEMENTS>*)fVertex;
DeviceVertex<NUM_VARYING_ELEMENTS> *varyings = (DeviceVertex<NUM_VARYING_ELEMENTS>*)fVaryings;
for (int i = start + threadIdx.x + blockIdx.x*blockDim.x;
i < end;
i += blockDim.x * gridDim.x) {
int fidx0 = F0_IT[tableOffset + 4 * i + 0];
int fidx1 = F0_IT[tableOffset + 4 * i + 1];
int fidx2 = F0_IT[tableOffset + 4 * i + 2];
int fidx3 = F0_IT[tableOffset + 4 * i + 3];
DeviceVertex<NUM_VERTEX_ELEMENTS> dst;
dst.clear();
if(NUM_VARYING_ELEMENTS > 0){
DeviceVertex<NUM_VARYING_ELEMENTS> dstVarying;
dstVarying.clear();
dst.addWithWeight(&vertex[fidx0], 0.25f);
dst.addWithWeight(&vertex[fidx1], 0.25f);
dst.addWithWeight(&vertex[fidx2], 0.25f);
dst.addWithWeight(&vertex[fidx3], 0.25f);
dstVarying.addWithWeight(&varyings[fidx0], 0.25f);
dstVarying.addWithWeight(&varyings[fidx1], 0.25f);
dstVarying.addWithWeight(&varyings[fidx2], 0.25f);
dstVarying.addWithWeight(&varyings[fidx3], 0.25f);
vertex[offset + i] = dst;
varyings[offset + i] = dstVarying;
}else{
dst.addWithWeight(&vertex[fidx0], 0.25f);
dst.addWithWeight(&vertex[fidx1], 0.25f);
dst.addWithWeight(&vertex[fidx2], 0.25f);
dst.addWithWeight(&vertex[fidx3], 0.25f);
vertex[offset + i] = dst;
}
}
}
__global__ void
computeQuadFace(float *fVertex, float *fVarying,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
int *F0_IT, int offset, int tableOffset, int start, int end)
{
for (int i = start +threadIdx.x + blockIdx.x*blockDim.x;
i < end;
i += blockDim.x * gridDim.x){
int fidx0 = F0_IT[tableOffset + 4 * i + 0];
int fidx1 = F0_IT[tableOffset + 4 * i + 1];
int fidx2 = F0_IT[tableOffset + 4 * i + 2];
int fidx3 = F0_IT[tableOffset + 4 * i + 3];
// XXX: can we use local stack like alloca?
float *dstVertex = fVertex + (i+offset)*vertexStride;
clear(dstVertex, vertexLength);
float *dstVarying = fVarying + (i+offset)*varyingStride;
clear(dstVarying, varyingLength);
addWithWeight(dstVertex, fVertex + fidx0*vertexStride, 0.25f, vertexLength);
addWithWeight(dstVertex, fVertex + fidx1*vertexStride, 0.25f, vertexLength);
addWithWeight(dstVertex, fVertex + fidx2*vertexStride, 0.25f, vertexLength);
addWithWeight(dstVertex, fVertex + fidx3*vertexStride, 0.25f, vertexLength);
addWithWeight(dstVarying, fVarying + fidx0*varyingStride, 0.25f, varyingLength);
addWithWeight(dstVarying, fVarying + fidx1*varyingStride, 0.25f, varyingLength);
addWithWeight(dstVarying, fVarying + fidx2*varyingStride, 0.25f, varyingLength);
addWithWeight(dstVarying, fVarying + fidx3*varyingStride, 0.25f, varyingLength);
}
}
template <int NUM_VERTEX_ELEMENTS, int NUM_VARYING_ELEMENTS> __global__ void
computeTriQuadFace(float *fVertex, float *fVaryings, int *F0_IT, int offset, int tableOffset, int start, int end)
{
DeviceVertex<NUM_VERTEX_ELEMENTS> *vertex = (DeviceVertex<NUM_VERTEX_ELEMENTS>*)fVertex;
DeviceVertex<NUM_VARYING_ELEMENTS> *varyings = (DeviceVertex<NUM_VARYING_ELEMENTS>*)fVaryings;
for (int i = start + threadIdx.x + blockIdx.x*blockDim.x;
i < end;
i += blockDim.x * gridDim.x) {
int fidx0 = F0_IT[tableOffset + 4 * i + 0];
int fidx1 = F0_IT[tableOffset + 4 * i + 1];
int fidx2 = F0_IT[tableOffset + 4 * i + 2];
int fidx3 = F0_IT[tableOffset + 4 * i + 3];
bool triangle = (fidx2 == fidx3);
float weight = triangle ? 1.0f / 3.0f : 1.0f / 4.0f;
DeviceVertex<NUM_VERTEX_ELEMENTS> dst;
dst.clear();
if(NUM_VARYING_ELEMENTS > 0){
DeviceVertex<NUM_VARYING_ELEMENTS> dstVarying;
dstVarying.clear();
dst.addWithWeight(&vertex[fidx0], weight);
dst.addWithWeight(&vertex[fidx1], weight);
dst.addWithWeight(&vertex[fidx2], weight);
dstVarying.addWithWeight(&varyings[fidx0], weight);
dstVarying.addWithWeight(&varyings[fidx1], weight);
dstVarying.addWithWeight(&varyings[fidx2], weight);
if (!triangle) {
dst.addWithWeight(&vertex[fidx3], weight);
dstVarying.addWithWeight(&varyings[fidx3], 0.25f);
}
vertex[offset + i] = dst;
varyings[offset + i] = dstVarying;
}else{
dst.addWithWeight(&vertex[fidx0], weight);
dst.addWithWeight(&vertex[fidx1], weight);
dst.addWithWeight(&vertex[fidx2], weight);
if (!triangle)
dst.addWithWeight(&vertex[fidx3], weight);
vertex[offset + i] = dst;
}
}
}
__global__ void
computeTriQuadFace(float *fVertex, float *fVarying,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
int *F0_IT, int offset, int tableOffset, int start, int end)
{
for (int i = start +threadIdx.x + blockIdx.x*blockDim.x;
i < end;
i += blockDim.x * gridDim.x){
int fidx0 = F0_IT[tableOffset + 4 * i + 0];
int fidx1 = F0_IT[tableOffset + 4 * i + 1];
int fidx2 = F0_IT[tableOffset + 4 * i + 2];
int fidx3 = F0_IT[tableOffset + 4 * i + 3];
bool triangle = (fidx2 == fidx3);
float weight = triangle ? 1.0f / 3.0f : 1.0f / 4.0f;
// XXX: can we use local stack like alloca?
float *dstVertex = fVertex + (i+offset)*vertexStride;
clear(dstVertex, vertexLength);
float *dstVarying = fVarying + (i+offset)*varyingStride;
clear(dstVarying, varyingLength);
addWithWeight(dstVertex, fVertex + fidx0*vertexStride, weight, vertexLength);
addWithWeight(dstVertex, fVertex + fidx1*vertexStride, weight, vertexLength);
addWithWeight(dstVertex, fVertex + fidx2*vertexStride, weight, vertexLength);
addWithWeight(dstVarying, fVarying + fidx0*varyingStride, weight, varyingLength);
addWithWeight(dstVarying, fVarying + fidx1*varyingStride, weight, varyingLength);
addWithWeight(dstVarying, fVarying + fidx2*varyingStride, weight, varyingLength);
if (!triangle) {
addWithWeight(dstVertex, fVertex + fidx3*vertexStride, weight, vertexLength);
addWithWeight(dstVarying, fVarying + fidx3*varyingStride, weight, varyingLength);
}
}
}
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
template <int NUM_VERTEX_ELEMENTS, int NUM_VARYING_ELEMENTS> __global__ void
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
computeEdge(float *fVertex, float *fVaryings, int *E0_IT, float *E0_S, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> *vertex = (DeviceVertex<NUM_VERTEX_ELEMENTS>*)fVertex;
DeviceVertex<NUM_VARYING_ELEMENTS> *varyings = (DeviceVertex<NUM_VARYING_ELEMENTS>*)fVaryings;
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i+= blockDim.x * gridDim.x){
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int eidx0 = E0_IT[4*i+0];
int eidx1 = E0_IT[4*i+1];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
int eidx2 = E0_IT[4*i+2];
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int eidx3 = E0_IT[4*i+3];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
float vertWeight = E0_S[i*2+0];
// Fully sharp edge : vertWeight = 0.5f;
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
DeviceVertex<NUM_VERTEX_ELEMENTS> dst;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dst.clear();
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dst.addWithWeight(&vertex[eidx0], vertWeight);
dst.addWithWeight(&vertex[eidx1], vertWeight);
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
if(eidx2 > -1){
float faceWeight = E0_S[i*2+1];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dst.addWithWeight(&vertex[eidx2], faceWeight);
dst.addWithWeight(&vertex[eidx3], faceWeight);
}
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
vertex[offset+i-tableOffset] = dst;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
if(NUM_VARYING_ELEMENTS > 0){
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
DeviceVertex<NUM_VARYING_ELEMENTS> dstVarying;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dstVarying.clear();
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
dstVarying.addWithWeight(&varyings[eidx0], 0.5f);
dstVarying.addWithWeight(&varyings[eidx1], 0.5f);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
varyings[offset+i-tableOffset] = dstVarying;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
}
}
__global__ void
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
computeEdge(float *fVertex, float *fVarying,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *E0_IT, float *E0_S, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +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
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;i+= blockDim.x * gridDim.x) {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int eidx0 = E0_IT[4*i+0];
int eidx1 = E0_IT[4*i+1];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
int eidx2 = E0_IT[4*i+2];
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int eidx3 = E0_IT[4*i+3];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
float vertWeight = E0_S[i*2+0];
// Fully sharp edge : vertWeight = 0.5f;
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
float *dstVertex = fVertex + (i+offset-tableOffset)*vertexStride;
clear(dstVertex, vertexLength);
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
addWithWeight(dstVertex, fVertex + eidx0*vertexStride, vertWeight, vertexLength);
addWithWeight(dstVertex, fVertex + eidx1*vertexStride, vertWeight, vertexLength);
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
if(eidx2 > -1){
float faceWeight = E0_S[i*2+1];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
addWithWeight(dstVertex, fVertex + eidx2*vertexStride, faceWeight, vertexLength);
addWithWeight(dstVertex, fVertex + eidx3*vertexStride, faceWeight, vertexLength);
adding CUDA kernels
2012-06-09 21:22:57 +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
if (varyingLength > 0){
float *dstVarying = fVarying + (i+offset-tableOffset)*varyingStride;
clear(dstVarying, varyingLength);
adding CUDA kernels
2012-06-09 21:22:57 +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
addWithWeight(dstVarying, fVarying + eidx0*varyingStride, 0.5f, varyingLength);
addWithWeight(dstVarying, fVarying + eidx1*varyingStride, 0.5f, varyingLength);
adding CUDA kernels
2012-06-09 21:22:57 +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
template <int NUM_VERTEX_ELEMENTS, int NUM_VARYING_ELEMENTS> __global__ void
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
computeVertexA(float *fVertex, float *fVaryings, int *V0_ITa, float *V0_S, int offset, int tableOffset, int start, int end, int pass)
adding CUDA kernels
2012-06-09 21:22:57 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> *vertex = (DeviceVertex<NUM_VERTEX_ELEMENTS>*)fVertex;
DeviceVertex<NUM_VARYING_ELEMENTS> *varyings = (DeviceVertex<NUM_VARYING_ELEMENTS>*)fVaryings;
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end+tableOffset;
i += blockDim.x * gridDim.x) {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int n = V0_ITa[5*i+1];
int p = V0_ITa[5*i+2];
int eidx0 = V0_ITa[5*i+3];
int eidx1 = V0_ITa[5*i+4];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
float weight = (pass==1) ? V0_S[i] : 1.0f - V0_S[i];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
// In the case of fractional weight, the weight must be inverted since
// the value is shared with the k_Smooth kernel (statistically the
adding CUDA kernels
2012-06-09 21:22:57 +00:00
// k_Smooth kernel runs much more often than this one)
if (weight>0.0f && weight<1.0f && n > 0)
weight=1.0f-weight;
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> dst;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
if (not pass) {
dst.clear();
} else {
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
dst = vertex[i+offset-tableOffset];
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
if (eidx0==-1 || (pass==0 && (n==-1)) ) {
dst.addWithWeight(&vertex[p], weight);
} else {
dst.addWithWeight(&vertex[p], weight * 0.75f);
dst.addWithWeight(&vertex[eidx0], weight * 0.125f);
dst.addWithWeight(&vertex[eidx1], weight * 0.125f);
}
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
vertex[i+offset-tableOffset] = dst;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
if(NUM_VARYING_ELEMENTS > 0){
if(not pass){
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
DeviceVertex<NUM_VARYING_ELEMENTS> dstVarying;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dstVarying.clear();
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
dstVarying.addWithWeight(&varyings[p], 1.0f);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
varyings[i+offset-tableOffset] = dstVarying;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
}
}
}
__global__ void
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
computeVertexA(float *fVertex, float *fVaryings,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *V0_ITa, float *V0_S, int offset, int tableOffset, int start, int end, int pass)
adding CUDA kernels
2012-06-09 21:22:57 +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
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i += blockDim.x * gridDim.x){
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int n = V0_ITa[5*i+1];
int p = V0_ITa[5*i+2];
int eidx0 = V0_ITa[5*i+3];
int eidx1 = V0_ITa[5*i+4];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
float weight = (pass==1) ? V0_S[i] : 1.0f - V0_S[i];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
// In the case of fractional weight, the weight must be inverted since
// the value is shared with the k_Smooth kernel (statistically the
adding CUDA kernels
2012-06-09 21:22:57 +00:00
// k_Smooth kernel runs much more often than this one)
if (weight>0.0f && weight<1.0f && n > 0)
weight=1.0f-weight;
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
float *dstVertex = fVertex + (i+offset-tableOffset)*vertexStride;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
if (not pass) {
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
clear(dstVertex, vertexLength);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
if (eidx0==-1 || (pass==0 && (n==-1)) ) {
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
addWithWeight(dstVertex, fVertex + p*vertexStride, weight, vertexLength);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
} else {
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
addWithWeight(dstVertex, fVertex + p*vertexStride, weight*0.75f, vertexLength);
addWithWeight(dstVertex, fVertex + eidx0*vertexStride, weight*0.125f, vertexLength);
addWithWeight(dstVertex, fVertex + eidx1*vertexStride, weight*0.125f, vertexLength);
adding CUDA kernels
2012-06-09 21:22:57 +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
if(varyingLength > 0){
adding CUDA kernels
2012-06-09 21:22:57 +00:00
if(not pass){
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
float *dstVarying = fVaryings + (i+offset-tableOffset)*varyingStride;
clear(dstVarying, varyingLength);
addWithWeight(dstVarying, fVaryings + p*varyingStride, 1.0f, varyingLength);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
}
}
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
//texture <int, 1> texV0_IT;
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
template <int NUM_VERTEX_ELEMENTS, int NUM_VARYING_ELEMENTS> __global__ void
adding CUDA kernels
2012-06-09 21:22:57 +00:00
computeVertexB(float *fVertex, float *fVaryings,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
const int *V0_ITa, const int *V0_IT, const float *V0_S, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> *vertex = (DeviceVertex<NUM_VERTEX_ELEMENTS>*)fVertex;
DeviceVertex<NUM_VARYING_ELEMENTS> *varyings = (DeviceVertex<NUM_VARYING_ELEMENTS>*)fVaryings;
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i += blockDim.x * gridDim.x) {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int h = V0_ITa[5*i];
int n = V0_ITa[5*i+1];
int p = V0_ITa[5*i+2];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
float weight = V0_S[i];
float wp = 1.0f/float(n*n);
float wv = (n-2.0f) * n * wp;
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> dst;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dst.clear();
dst.addWithWeight(&vertex[p], weight * wv);
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
for (int j = 0; j < n; ++j) {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dst.addWithWeight(&vertex[V0_IT[h+j*2]], weight * wp);
dst.addWithWeight(&vertex[V0_IT[h+j*2+1]], weight * wp);
// int idx0 = tex1Dfetch(texV0_IT, h+j*2);
// int idx1 = tex1Dfetch(texV0_IT, h+j*2+1);
// dst.addWithWeight(&vertex[idx0], weight * wp);
// dst.addWithWeight(&vertex[idx1], weight * wp);
}
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
vertex[i+offset-tableOffset] = dst;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
if(NUM_VARYING_ELEMENTS > 0){
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
DeviceVertex<NUM_VARYING_ELEMENTS> dstVarying;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dstVarying.clear();
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
dstVarying.addWithWeight(&varyings[p], 1.0f);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
varyings[i+offset-tableOffset] = dstVarying;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
}
}
__global__ void
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
computeVertexB(float *fVertex, float *fVarying,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
const int *V0_ITa, const int *V0_IT, const float *V0_S, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +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
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i += blockDim.x * gridDim.x) {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int h = V0_ITa[5*i];
int n = V0_ITa[5*i+1];
int p = V0_ITa[5*i+2];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
float weight = V0_S[i];
float wp = 1.0f/float(n*n);
float wv = (n-2.0f) * n * wp;
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
float *dstVertex = fVertex + (i+offset-tableOffset)*vertexStride;
clear(dstVertex, vertexLength);
addWithWeight(dstVertex, fVertex + p*vertexStride, weight*wv, vertexLength);
adding CUDA kernels
2012-06-09 21:22:57 +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
for (int j = 0; j < n; ++j) {
addWithWeight(dstVertex, fVertex + V0_IT[h+j*2]*vertexStride, weight*wp, vertexLength);
addWithWeight(dstVertex, fVertex + V0_IT[h+j*2+1]*vertexStride, weight*wp, vertexLength);
adding CUDA kernels
2012-06-09 21:22:57 +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
if (varyingLength > 0) {
float *dstVarying = fVarying + (i+offset-tableOffset)*varyingStride;
clear(dstVarying, varyingLength);
addWithWeight(dstVarying, fVarying + p*varyingStride, 1.0f, varyingLength);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
}
}
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +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
template <int NUM_VERTEX_ELEMENTS, int NUM_VARYING_ELEMENTS> __global__ void
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
computeLoopVertexB(float *fVertex, float *fVaryings, int *V0_ITa, int *V0_IT, float *V0_S, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> *vertex = (DeviceVertex<NUM_VERTEX_ELEMENTS>*)fVertex;
DeviceVertex<NUM_VARYING_ELEMENTS> *varyings = (DeviceVertex<NUM_VARYING_ELEMENTS>*)fVaryings;
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i += blockDim.x * gridDim.x) {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int h = V0_ITa[5*i];
int n = V0_ITa[5*i+1];
int p = V0_ITa[5*i+2];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
float weight = V0_S[i];
float wp = 1.0f/float(n);
float beta = 0.25f * __cosf(float(M_PI) * 2.0f * wp) + 0.375f;
beta = beta * beta;
beta = (0.625f - beta) * wp;
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> dst;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dst.clear();
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dst.addWithWeight(&vertex[p], weight * (1.0f - (beta * n)));
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
for (int j = 0; j < n; ++j) {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dst.addWithWeight(&vertex[V0_IT[h+j]], weight * beta);
}
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
vertex[i+offset-tableOffset] = dst;
adding CUDA kernels
2012-06-09 21:22:57 +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
if (NUM_VARYING_ELEMENTS > 0) {
DeviceVertex<NUM_VARYING_ELEMENTS> dstVarying;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
dstVarying.clear();
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
dstVarying.addWithWeight(&varyings[p], 1.0f);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
varyings[i+offset-tableOffset] = dstVarying;
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
}
}
__global__ void
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
computeLoopVertexB(float *fVertex, float *fVarying,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
const int *V0_ITa, const int *V0_IT, const float *V0_S, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +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
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i += blockDim.x * gridDim.x) {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
int h = V0_ITa[5*i];
int n = V0_ITa[5*i+1];
int p = V0_ITa[5*i+2];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
adding CUDA kernels
2012-06-09 21:22:57 +00:00
float weight = V0_S[i];
float wp = 1.0f/float(n);
float beta = 0.25f * __cosf(float(M_PI) * 2.0f * wp) + 0.375f;
beta = beta * beta;
beta = (0.625f - beta) * wp;
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
float *dstVertex = fVertex + (i+offset-tableOffset)*vertexStride;
clear(dstVertex, vertexLength);
addWithWeight(dstVertex, fVertex + p*vertexStride, weight*(1.0f-(beta*n)), vertexLength);
adding CUDA kernels
2012-06-09 21:22:57 +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
for (int j = 0; j < n; ++j) {
addWithWeight(dstVertex, fVertex + V0_IT[h+j]*vertexStride, weight*beta, vertexLength);
adding CUDA kernels
2012-06-09 21:22:57 +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
if (varyingLength > 0) {
float *dstVarying = fVarying + (i+offset-tableOffset)*varyingStride;
clear(dstVarying, varyingLength);
addWithWeight(dstVarying, fVarying + p*varyingStride, 1.0f, varyingLength);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
}
}
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
template <int NUM_VERTEX_ELEMENTS, int NUM_VARYING_ELEMENTS> __global__ void
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
computeBilinearEdge(float *fVertex, float *fVaryings, int *E0_IT, int offset, int tableOffset, int start, int end)
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> *vertex = (DeviceVertex<NUM_VERTEX_ELEMENTS>*)fVertex;
DeviceVertex<NUM_VARYING_ELEMENTS> *varyings = (DeviceVertex<NUM_VARYING_ELEMENTS>*)fVaryings;
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i+= blockDim.x * gridDim.x) {
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
int eidx0 = E0_IT[2*i+0];
int eidx1 = E0_IT[2*i+1];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> dst;
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
dst.clear();
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
dst.addWithWeight(&vertex[eidx0], 0.5f);
dst.addWithWeight(&vertex[eidx1], 0.5f);
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
vertex[offset+i-tableOffset] = dst;
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
if (NUM_VARYING_ELEMENTS > 0) {
DeviceVertex<NUM_VARYING_ELEMENTS> dstVarying;
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
dstVarying.clear();
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
dstVarying.addWithWeight(&varyings[eidx0], 0.5f);
dstVarying.addWithWeight(&varyings[eidx1], 0.5f);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
varyings[offset+i-tableOffset] = dstVarying;
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
}
}
}
__global__ void
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
computeBilinearEdge(float *fVertex, float *fVarying,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *E0_IT, int offset, int tableOffset, int start, int end)
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i+= blockDim.x * gridDim.x) {
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
int eidx0 = E0_IT[2*i+0];
int eidx1 = E0_IT[2*i+1];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
float *dstVertex = fVertex + (i+offset-tableOffset)*vertexStride;
clear(dstVertex, vertexLength);
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
addWithWeight(dstVertex, fVertex + eidx0*vertexStride, 0.5f, vertexLength);
addWithWeight(dstVertex, fVertex + eidx1*vertexStride, 0.5f, vertexLength);
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
if (varyingLength > 0) {
float *dstVarying = fVarying + (i+offset-tableOffset)*varyingStride;
clear(dstVarying, varyingLength);
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
addWithWeight(dstVarying, fVarying + eidx0*varyingStride, 0.5f, varyingLength);
addWithWeight(dstVarying, fVarying + eidx1*varyingStride, 0.5f, varyingLength);
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
template <int NUM_VERTEX_ELEMENTS, int NUM_VARYING_ELEMENTS> __global__ void
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
computeBilinearVertex(float *fVertex, float *fVaryings, int *V0_ITa, int offset, int tableOffset, int start, int end)
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> *vertex = (DeviceVertex<NUM_VERTEX_ELEMENTS>*)fVertex;
DeviceVertex<NUM_VARYING_ELEMENTS> *varyings = (DeviceVertex<NUM_VARYING_ELEMENTS>*)fVaryings;
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i += blockDim.x * gridDim.x) {
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
int p = V0_ITa[i];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
DeviceVertex<NUM_VERTEX_ELEMENTS> dst;
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
dst.clear();
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
dst.addWithWeight(&vertex[p], 1.0f);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
vertex[i+offset-tableOffset] = dst;
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
if (NUM_VARYING_ELEMENTS > 0) {
DeviceVertex<NUM_VARYING_ELEMENTS> dstVarying;
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
dstVarying.clear();
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
dstVarying.addWithWeight(&varyings[p], 1.0f);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
varyings[i+offset-tableOffset] = dstVarying;
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
}
}
}
__global__ void
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
computeBilinearVertex(float *fVertex, float *fVarying,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
const int *V0_ITa, int offset, int tableOffset, int start, int end)
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i += blockDim.x * gridDim.x) {
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
int p = V0_ITa[i];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
float *dstVertex = fVertex + (i+offset-tableOffset)*vertexStride;
clear(dstVertex, vertexLength);
addWithWeight(dstVertex, fVertex + p*vertexStride, 1.0f, vertexLength);
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
if (varyingLength > 0) {
float *dstVarying = fVarying + (i+offset-tableOffset)*varyingStride;
clear(dstVarying, varyingLength);
addWithWeight(dstVarying, fVarying + p*varyingStride, 1.0f, varyingLength);
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
}
}
}
adding CUDA kernels
2012-06-09 21:22:57 +00:00
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
// --------------------------------------------------------------------------------------------
__global__ void
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
editVertexAdd(float *fVertex, int vertexLength, int vertexStride,
int primVarOffset, int primVarWidth,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int vertexOffset, int tableOffset, int start, int end,
const int *editIndices, const float *editValues)
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
for (int i = start + tableOffset + threadIdx.x + blockIdx.x*blockDim.x;
i < end + tableOffset;
i += blockDim.x * gridDim.x) {
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +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
float *dstVertex = fVertex + (editIndices[i] + vertexOffset) * vertexStride + primVarOffset;
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
for(int j = 0; j < primVarWidth; j++) {
Fix a indexing bug in hierarchical edit kernel of Cuda, GLSLcompute and GLSLTransformFeedback Closes #183
2013-06-26 23:03:48 +00:00
*dstVertex++ += editValues[i*primVarWidth + j];
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
}
}
}
adding CUDA kernels
2012-06-09 21:22:57 +00:00
// --------------------------------------------------------------------------------------------
including the correct version file so namespaces work
2012-06-12 00:02:27 +00:00
#include "../version.h"
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
// XXX: this macro usage is tentative. Since cuda kernel can't be dynamically configured,
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
// still trying to find better way to have optimized kernel..
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
#define OPT_KERNEL(NUM_VERTEX_ELEMENTS, NUM_VARYING_ELEMENTS, KERNEL, X, Y, ARG) \
if(vertexLength == NUM_VERTEX_ELEMENTS && \
varyingLength == NUM_VARYING_ELEMENTS && \
vertexStride == vertexLength && \
varyingStride == varyingLength) \
{ KERNEL<NUM_VERTEX_ELEMENTS, NUM_VARYING_ELEMENTS><<<X,Y>>>ARG; \
return; }
adding CUDA kernels
2012-06-09 21:22:57 +00:00
including the correct version file so namespaces work
2012-06-12 00:02:27 +00:00
extern "C" {
adding CUDA kernels
2012-06-09 21:22:57 +00:00
void OsdCudaComputeFace(float *vertex, float *varying,
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
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *F_IT, int *F_ITa, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +00:00
{
//computeFace<3, 0><<<512,32>>>(vertex, varying, F_IT, F_ITa, offset, start, end);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
OPT_KERNEL(0, 0, computeFace, 512, 32, (vertex, varying, F_IT, F_ITa, offset, tableOffset, start, end));
OPT_KERNEL(0, 3, computeFace, 512, 32, (vertex, varying, F_IT, F_ITa, offset, tableOffset, start, end));
OPT_KERNEL(3, 0, computeFace, 512, 32, (vertex, varying, F_IT, F_ITa, offset, tableOffset, start, end));
OPT_KERNEL(3, 3, computeFace, 512, 32, (vertex, varying, F_IT, F_ITa, offset, tableOffset, start, end));
adding CUDA kernels
2012-06-09 21:22:57 +00:00
// fallback kernel (slow)
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
computeFace<<<512, 32>>>(vertex, varying,
vertexLength, vertexStride, varyingLength, varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
F_IT, F_ITa, offset, tableOffset, start, end);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
Implemented CATMARK_QUAD_FACE_VERTEX and CATMARK_TRI_QUAD_FACE_VERTEX kernels for CUDA, GLSL Transform Feedback, and GLSL Compute platforms.
2014-05-28 23:46:38 +00:00
void OsdCudaComputeQuadFace(float *vertex, float *varying,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
int *F_IT, int offset, int tableOffset, int start, int end)
{
//computeQuadFace<3, 0><<<512,32>>>(vertex, varying, F_IT, offset, start, end);
OPT_KERNEL(0, 0, computeQuadFace, 512, 32, (vertex, varying, F_IT, offset, tableOffset, start, end));
OPT_KERNEL(0, 3, computeQuadFace, 512, 32, (vertex, varying, F_IT, offset, tableOffset, start, end));
OPT_KERNEL(3, 0, computeQuadFace, 512, 32, (vertex, varying, F_IT, offset, tableOffset, start, end));
OPT_KERNEL(3, 3, computeQuadFace, 512, 32, (vertex, varying, F_IT, offset, tableOffset, start, end));
// fallback kernel (slow)
computeQuadFace<<<512, 32>>>(vertex, varying,
vertexLength, vertexStride, varyingLength, varyingStride,
F_IT, offset, tableOffset, start, end);
}
void OsdCudaComputeTriQuadFace(float *vertex, float *varying,
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
int *F_IT, int offset, int tableOffset, int start, int end)
{
//computeTriQuadFace<3, 0><<<512,32>>>(vertex, varying, F_IT, offset, start, end);
OPT_KERNEL(0, 0, computeTriQuadFace, 512, 32, (vertex, varying, F_IT, offset, tableOffset, start, end));
OPT_KERNEL(0, 3, computeTriQuadFace, 512, 32, (vertex, varying, F_IT, offset, tableOffset, start, end));
OPT_KERNEL(3, 0, computeTriQuadFace, 512, 32, (vertex, varying, F_IT, offset, tableOffset, start, end));
OPT_KERNEL(3, 3, computeTriQuadFace, 512, 32, (vertex, varying, F_IT, offset, tableOffset, start, end));
// fallback kernel (slow)
computeTriQuadFace<<<512, 32>>>(vertex, varying,
vertexLength, vertexStride, varyingLength, varyingStride,
F_IT, offset, tableOffset, start, end);
}
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
adding CUDA kernels
2012-06-09 21:22:57 +00:00
void OsdCudaComputeEdge(float *vertex, float *varying,
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
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *E_IT, float *E_W, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +00:00
{
//computeEdge<0, 3><<<512,32>>>(vertex, varying, E_IT, E_W, offset, start, end);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
OPT_KERNEL(0, 0, computeEdge, 512, 32, (vertex, varying, E_IT, E_W, offset, tableOffset, start, end));
OPT_KERNEL(0, 3, computeEdge, 512, 32, (vertex, varying, E_IT, E_W, offset, tableOffset, start, end));
OPT_KERNEL(3, 0, computeEdge, 512, 32, (vertex, varying, E_IT, E_W, offset, tableOffset, start, end));
OPT_KERNEL(3, 3, computeEdge, 512, 32, (vertex, varying, E_IT, E_W, offset, tableOffset, start, end));
adding CUDA kernels
2012-06-09 21:22:57 +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
computeEdge<<<512, 32>>>(vertex, varying,
vertexLength, vertexStride, varyingLength, varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
E_IT, E_W, offset, tableOffset, start, end);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
void OsdCudaComputeVertexA(float *vertex, float *varying,
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
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *V_ITa, float *V_W, int offset, int tableOffset, int start, int end, int pass)
adding CUDA kernels
2012-06-09 21:22:57 +00:00
{
// computeVertexA<0, 3><<<512,32>>>(vertex, varying, V_ITa, V_W, offset, start, end, pass);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
OPT_KERNEL(0, 0, computeVertexA, 512, 32, (vertex, varying, V_ITa, V_W, offset, tableOffset, start, end, pass));
OPT_KERNEL(0, 3, computeVertexA, 512, 32, (vertex, varying, V_ITa, V_W, offset, tableOffset, start, end, pass));
OPT_KERNEL(3, 0, computeVertexA, 512, 32, (vertex, varying, V_ITa, V_W, offset, tableOffset, start, end, pass));
OPT_KERNEL(3, 3, computeVertexA, 512, 32, (vertex, varying, V_ITa, V_W, offset, tableOffset, start, end, pass));
adding CUDA kernels
2012-06-09 21:22:57 +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
computeVertexA<<<512, 32>>>(vertex, varying,
vertexLength, vertexStride, varyingLength, varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
V_ITa, V_W, offset, tableOffset, start, end, pass);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
void OsdCudaComputeVertexB(float *vertex, float *varying,
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
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *V_ITa, int *V_IT, float *V_W, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +00:00
{
// computeVertexB<0, 3><<<512,32>>>(vertex, varying, V_ITa, V_IT, V_W, offset, start, end);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
OPT_KERNEL(0, 0, computeVertexB, 512, 32, (vertex, varying, V_ITa, V_IT, V_W, offset, tableOffset, start, end));
OPT_KERNEL(0, 3, computeVertexB, 512, 32, (vertex, varying, V_ITa, V_IT, V_W, offset, tableOffset, start, end));
OPT_KERNEL(3, 0, computeVertexB, 512, 32, (vertex, varying, V_ITa, V_IT, V_W, offset, tableOffset, start, end));
OPT_KERNEL(3, 3, computeVertexB, 512, 32, (vertex, varying, V_ITa, V_IT, V_W, offset, tableOffset, start, end));
adding CUDA kernels
2012-06-09 21:22:57 +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
computeVertexB<<<512, 32>>>(vertex, varying,
vertexLength, vertexStride, varyingLength, varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
V_ITa, V_IT, V_W, offset, tableOffset, start, end);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
void OsdCudaComputeLoopVertexB(float *vertex, float *varying,
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
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *V_ITa, int *V_IT, float *V_W, int offset, int tableOffset, int start, int end)
adding CUDA kernels
2012-06-09 21:22:57 +00:00
{
// computeLoopVertexB<0, 3><<<512,32>>>(vertex, varying, V_ITa, V_IT, V_W, offset, start, end);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
OPT_KERNEL(0, 0, computeLoopVertexB, 512, 32, (vertex, varying, V_ITa, V_IT, V_W, offset, tableOffset, start, end));
OPT_KERNEL(0, 3, computeLoopVertexB, 512, 32, (vertex, varying, V_ITa, V_IT, V_W, offset, tableOffset, start, end));
OPT_KERNEL(3, 0, computeLoopVertexB, 512, 32, (vertex, varying, V_ITa, V_IT, V_W, offset, tableOffset, start, end));
OPT_KERNEL(3, 3, computeLoopVertexB, 512, 32, (vertex, varying, V_ITa, V_IT, V_W, offset, tableOffset, start, end));
adding CUDA kernels
2012-06-09 21:22:57 +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
computeLoopVertexB<<<512, 32>>>(vertex, varying,
vertexLength, vertexStride, varyingLength, varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
V_ITa, V_IT, V_W, offset, tableOffset, start, end);
adding CUDA kernels
2012-06-09 21:22:57 +00:00
}
including the correct version file so namespaces work
2012-06-12 00:02:27 +00:00
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
void OsdCudaComputeBilinearEdge(float *vertex, float *varying,
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
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *E_IT, int offset, int tableOffset, int start, int end)
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
{
//computeBilinearEdge<0, 3><<<512,32>>>(vertex, varying, E_IT, offset, start, end);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
OPT_KERNEL(0, 0, computeBilinearEdge, 512, 32, (vertex, varying, E_IT, offset, tableOffset, start, end));
OPT_KERNEL(0, 3, computeBilinearEdge, 512, 32, (vertex, varying, E_IT, offset, tableOffset, start, end));
OPT_KERNEL(3, 0, computeBilinearEdge, 512, 32, (vertex, varying, E_IT, offset, tableOffset, start, end));
OPT_KERNEL(3, 3, computeBilinearEdge, 512, 32, (vertex, varying, E_IT, offset, tableOffset, start, end));
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
computeBilinearEdge<<<512, 32>>>(vertex, varying,
vertexLength, vertexStride, varyingLength, varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
E_IT, offset, tableOffset, start, end);
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
}
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
void OsdCudaComputeBilinearVertex(float *vertex, float *varying,
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
int vertexLength, int vertexStride,
int varyingLength, int varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int *V_ITa, int offset, int tableOffset, int start, int end)
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +00:00
{
// computeBilinearVertex<0, 3><<<512,32>>>(vertex, varying, V_ITa, offset, start, end);
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
OPT_KERNEL(0, 0, computeBilinearVertex, 512, 32, (vertex, varying, V_ITa, offset, tableOffset, start, end));
OPT_KERNEL(0, 3, computeBilinearVertex, 512, 32, (vertex, varying, V_ITa, offset, tableOffset, start, end));
OPT_KERNEL(3, 0, computeBilinearVertex, 512, 32, (vertex, varying, V_ITa, offset, tableOffset, start, end));
OPT_KERNEL(3, 3, computeBilinearVertex, 512, 32, (vertex, varying, V_ITa, offset, tableOffset, start, end));
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
computeBilinearVertex<<<512, 32>>>(vertex, varying,
vertexLength, vertexStride, varyingLength, varyingStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
V_ITa, offset, tableOffset, start, end);
add bilinear computation for cuda kernel (not yet for CL, GLSL kernels)
2012-06-12 01:09:23 +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
void OsdCudaEditVertexAdd(float *vertex, int vertexLength, int vertexStride,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
int primVarOffset, int primVarWidth,
int vertexOffset, int tableOffset,
int start, int end, int *editIndices, float *editValues)
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +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
editVertexAdd<<<512, 32>>>(vertex, vertexLength, vertexStride, primVarOffset, primVarWidth,
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
vertexOffset, tableOffset, start, end,
editIndices, editValues);
Siggrpah 2012 - rolling over all of prepro work into beta 1.1
2012-08-04 02:51:27 +00:00
}
Batching stuffs: generalized kernel batches, table/dispatcher refactoring, multiMeshFactory, drawContext, etc. 2 client APIs are changed. - VertexBuffer::UpdateData() takes start vertex offset - ComputeController::Refine() takes FarKernelBatchVector Also, ComputeContext no longer holds farmesh. Client can free farmesh after OsdComputeContext is created. (but still need FarKernelBatchVector to apply subdivision kernels)
2013-03-08 01:50:15 +00:00
} /* extern "C" */
Reference in New Issue Copy Permalink