Bart/bullet3
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Clean up plbvhCalculateOverlappingPairs kernel.

Browse Source 
Also fix overlapping pair generation with triangle mesh.
(Currently, large/concave AABBs must be the first entry in a pair.)
This commit is contained in:
Jackson Lee
2014-03-14 13:55:22 -07:00
parent f324e66f86
commit b709d6beeb
2 changed files with 20 additions and 31 deletions
Download Patch File Download Diff File Expand all files Collapse all files

26
src/Bullet3OpenCL/BroadphaseCollision/kernels/parallelLinearBvh.cl
View File

@@ -184,18 +184,12 @@ __kernel void plbvhCalculateOverlappingPairs(__global b3AabbCL* rigidAabbs,
__global int* out_numPairs, __global int4* out_overlappingPairs,
int maxPairs, int numQueryAabbs)
{
#define USE_SPATIALLY_COHERENT_INDICIES //mortonCodesAndAabbIndices[] contains rigid body indices sorted along the z-curve
#ifdef USE_SPATIALLY_COHERENT_INDICIES
int queryRigidIndex = get_group_id(0) * get_local_size(0) + get_local_id(0);
if(queryRigidIndex >= numQueryAabbs) return;
//Using get_group_id()/get_local_id() is Faster than get_global_id(0) since
//mortonCodesAndAabbIndices[] contains rigid body indices sorted along the z-curve (more spatially coherent)
int queryBvhNodeIndex = get_group_id(0) * get_local_size(0) + get_local_id(0);
if(queryBvhNodeIndex >= numQueryAabbs) return;
int queryBvhNodeIndex = queryRigidIndex;
queryRigidIndex = mortonCodesAndAabbIndices[queryRigidIndex].m_value; // fix queryRigidIndex naming for this branch
#else
int queryRigidIndex = get_global_id(0);
if(queryRigidIndex >= numQueryAabbs) return;
#endif
int queryRigidIndex = mortonCodesAndAabbIndices[queryBvhNodeIndex].m_value;
b3AabbCL queryAabb = rigidAabbs[queryRigidIndex];
int stack[B3_PLVBH_TRAVERSE_MAX_STACK_SIZE];
@@ -213,17 +207,17 @@ __kernel void plbvhCalculateOverlappingPairs(__global b3AabbCL* rigidAabbs,
//Optimization - if the BVH is structured as a binary radix tree, then
//each internal node corresponds to a contiguous range of leaf nodes(internalNodeLeafIndexRanges[]).
//This can be used to avoid testing each AABB-AABB pair twice.
//This can be used to avoid testing each AABB-AABB pair twice, including preventing each node from colliding with itself.
{
int highestLeafIndex = (isLeaf) ? bvhNodeIndex : internalNodeLeafIndexRanges[bvhNodeIndex].y;
if(highestLeafIndex < queryBvhNodeIndex) continue;
if(highestLeafIndex <= queryBvhNodeIndex) continue;
}
//bvhRigidIndex is not used if internal node
int bvhRigidIndex = (isLeaf) ? mortonCodesAndAabbIndices[bvhNodeIndex].m_value : -1;
b3AabbCL bvhNodeAabb = (isLeaf) ? rigidAabbs[bvhRigidIndex] : internalNodeAabbs[bvhNodeIndex];
if( queryRigidIndex != bvhRigidIndex && TestAabbAgainstAabb2(&queryAabb, &bvhNodeAabb) )
if( TestAabbAgainstAabb2(&queryAabb, &bvhNodeAabb) )
{
if(isLeaf)
{
@@ -388,8 +382,8 @@ __kernel void plbvhLargeAabbAabbTest(__global b3AabbCL* smallAabbs, __global b3A
if( TestAabbAgainstAabb2(&smallAabb, &largeAabb) )
{
int4 pair;
pair.x = smallAabb.m_minIndices[3];
pair.y = largeAabb.m_minIndices[3];
pair.x = largeAabb.m_minIndices[3];
pair.y = smallAabb.m_minIndices[3];
pair.z = NEW_PAIR_MARKER;
pair.w = NEW_PAIR_MARKER;

25
src/Bullet3OpenCL/BroadphaseCollision/kernels/parallelLinearBvhKernels.h
View File

@@ -170,17 +170,12 @@ static const char* parallelLinearBvhCL= \
" __global int* out_numPairs, __global int4* out_overlappingPairs, \n"
" int maxPairs, int numQueryAabbs)\n"
"{\n"
"#define USE_SPATIALLY_COHERENT_INDICIES //mortonCodesAndAabbIndices[] contains rigid body indices sorted along the z-curve\n"
"#ifdef USE_SPATIALLY_COHERENT_INDICIES\n"
" int queryRigidIndex = get_group_id(0) * get_local_size(0) + get_local_id(0);\n"
" if(queryRigidIndex >= numQueryAabbs) return;\n"
" //Using get_group_id()/get_local_id() is Faster than get_global_id(0) since\n"
" //mortonCodesAndAabbIndices[] contains rigid body indices sorted along the z-curve (more spatially coherent)\n"
" int queryBvhNodeIndex = get_group_id(0) * get_local_size(0) + get_local_id(0);\n"
" if(queryBvhNodeIndex >= numQueryAabbs) return;\n"
" \n"
" int queryBvhNodeIndex = queryRigidIndex;\n"
" queryRigidIndex = mortonCodesAndAabbIndices[queryRigidIndex].m_value; // fix queryRigidIndex naming for this branch\n"
"#else\n"
" int queryRigidIndex = get_global_id(0);\n"
" if(queryRigidIndex >= numQueryAabbs) return;\n"
"#endif\n"
" int queryRigidIndex = mortonCodesAndAabbIndices[queryBvhNodeIndex].m_value;\n"
" b3AabbCL queryAabb = rigidAabbs[queryRigidIndex];\n"
" \n"
" int stack[B3_PLVBH_TRAVERSE_MAX_STACK_SIZE];\n"
@@ -198,17 +193,17 @@ static const char* parallelLinearBvhCL= \
" \n"
" //Optimization - if the BVH is structured as a binary radix tree, then\n"
" //each internal node corresponds to a contiguous range of leaf nodes(internalNodeLeafIndexRanges[]).\n"
" //This can be used to avoid testing each AABB-AABB pair twice.\n"
" //This can be used to avoid testing each AABB-AABB pair twice, including preventing each node from colliding with itself.\n"
" {\n"
" int highestLeafIndex = (isLeaf) ? bvhNodeIndex : internalNodeLeafIndexRanges[bvhNodeIndex].y;\n"
" if(highestLeafIndex < queryBvhNodeIndex) continue;\n"
" if(highestLeafIndex <= queryBvhNodeIndex) continue;\n"
" }\n"
" \n"
" //bvhRigidIndex is not used if internal node\n"
" int bvhRigidIndex = (isLeaf) ? mortonCodesAndAabbIndices[bvhNodeIndex].m_value : -1;\n"
" \n"
" b3AabbCL bvhNodeAabb = (isLeaf) ? rigidAabbs[bvhRigidIndex] : internalNodeAabbs[bvhNodeIndex];\n"
" if( queryRigidIndex != bvhRigidIndex && TestAabbAgainstAabb2(&queryAabb, &bvhNodeAabb) )\n"
" if( TestAabbAgainstAabb2(&queryAabb, &bvhNodeAabb) )\n"
" {\n"
" if(isLeaf)\n"
" {\n"
@@ -365,8 +360,8 @@ static const char* parallelLinearBvhCL= \
" if( TestAabbAgainstAabb2(&smallAabb, &largeAabb) )\n"
" {\n"
" int4 pair;\n"
" pair.x = smallAabb.m_minIndices[3];\n"
" pair.y = largeAabb.m_minIndices[3];\n"
" pair.x = largeAabb.m_minIndices[3];\n"
" pair.y = smallAabb.m_minIndices[3];\n"
" pair.z = NEW_PAIR_MARKER;\n"
" pair.w = NEW_PAIR_MARKER;\n"
" \n"