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

move parts of collision pipeline to shared header files (work-in-progress)

Browse Source
This commit is contained in:
erwin coumans
2013-12-12 11:03:55 -08:00
parent ba2ba28a89
commit c155e126d0
17 changed files with 969 additions and 120 deletions
Download Patch File Download Diff File Expand all files Collapse all files

219
src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp
View File

@@ -13,6 +13,10 @@ subject to the following restrictions:
3. This notice may not be removed or altered from any source distribution.
*/
bool findSeparatingAxisOnGpu = true;
bool bvhTraversalKernelGPU = true;
bool findConcaveSeparatingAxisKernelGPU = false;//true;
///This file was written by Erwin Coumans
///Separating axis rest based on work from Pierre Terdiman, see
@@ -20,7 +24,7 @@ subject to the following restrictions:
//#define B3_DEBUG_SAT_FACE
//#define CHECK_ON_HOST
#define CHECK_ON_HOST
#ifdef CHECK_ON_HOST
//#define PERSISTENT_CONTACTS_HOST
@@ -65,6 +69,11 @@ typedef b3AlignedObjectArray<b3Vector3> b3VertexArray;
#endif
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3BvhTraversal.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3FindConcaveSatAxis.h"
#define dot3F4 b3Dot
GpuSatCollision::GpuSatCollision(cl_context ctx,cl_device_id device, cl_command_queue q )
@@ -1197,7 +1206,7 @@ int clipHullHullSingle(
int numPoints = 0;
{
B3_PROFILE("extractManifold");
// B3_PROFILE("extractManifold");
numPoints = extractManifold(contactsOut, numContactsOut, normalOnSurfaceB, &contactIdx);
}
@@ -2723,6 +2732,9 @@ int computeContactConvexConvex2(
}
void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>* pairs, int nPairs,
const b3OpenCLArray<b3RigidBodyCL>* bodyBuf,
b3OpenCLArray<b3Contact4>* contactOut, int& nContacts,
@@ -2898,15 +2910,13 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
hostCollidables[collidableIndexB].m_shapeType == SHAPE_CONVEX_HULL)
{
//printf("hostPairs[i].z=%d\n",hostPairs[i].z);
int contactIndex = computeContactConvexConvex2(i,bodyIndexA,bodyIndexB,collidableIndexA,collidableIndexB,hostBodyBuf,
hostCollidables,hostConvexData,hostVertices,hostUniqueEdges,hostIndices,hostFaces,hostContacts,nContacts,maxContactCapacity,oldHostContacts);
//int contactIndex = computeContactConvexConvex(hostPairs,i,bodyIndexA,bodyIndexB,collidableIndexA,collidableIndexB,hostBodyBuf,
// hostCollidables,hostConvexData,hostVertices,hostUniqueEdges,hostIndices,hostFaces,hostContacts,nContacts,maxContactCapacity,
// oldHostContacts);
int contactIndex = 0;//computeContactConvexConvex2(i,bodyIndexA,bodyIndexB,collidableIndexA,collidableIndexB,hostBodyBuf, hostCollidables,hostConvexData,hostVertices,hostUniqueEdges,hostIndices,hostFaces,hostContacts,nContacts,maxContactCapacity,oldHostContacts);
//int contactIndex = computeContactConvexConvex(hostPairs,i,bodyIndexA,bodyIndexB,collidableIndexA,collidableIndexB,hostBodyBuf,hostCollidables,hostConvexData,hostVertices,hostUniqueEdges,hostIndices,hostFaces,hostContacts,nContacts,maxContactCapacity,oldHostContacts);
if (contactIndex>=0)
{
// printf("convex convex contactIndex = %d\n",contactIndex);
hostPairs[i].z = contactIndex;
}
// printf("plane-convex\n");
@@ -2932,7 +2942,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
contactOut->resize(0);
}
return;
m_totalContactsOut.copyFromHostPointer(&nContacts,1,0,true);
#else
{
@@ -2996,7 +3007,6 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
int numCompoundPairs = 0;
bool findSeparatingAxisOnGpu = true;//false;
int numConcavePairs =0;
{
@@ -3038,65 +3048,172 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
if (treeNodesGPU->size() && treeNodesGPU->size())
{
B3_PROFILE("m_bvhTraversalKernel");
if (bvhTraversalKernelGPU)
{
B3_PROFILE("m_bvhTraversalKernel");
numConcavePairs = m_numConcavePairsOut.at(0);
numConcavePairs = m_numConcavePairsOut.at(0);
b3LauncherCL launcher(m_queue, m_bvhTraversalKernel,"m_bvhTraversalKernel");
launcher.setBuffer( pairs->getBufferCL());
launcher.setBuffer( bodyBuf->getBufferCL());
launcher.setBuffer( gpuCollidables.getBufferCL());
launcher.setBuffer( clAabbsWorldSpace.getBufferCL());
launcher.setBuffer( triangleConvexPairsOut.getBufferCL());
launcher.setBuffer( m_numConcavePairsOut.getBufferCL());
launcher.setBuffer( subTreesGPU->getBufferCL());
launcher.setBuffer( treeNodesGPU->getBufferCL());
launcher.setBuffer( bvhInfo->getBufferCL());
b3LauncherCL launcher(m_queue, m_bvhTraversalKernel,"m_bvhTraversalKernel");
launcher.setBuffer( pairs->getBufferCL());
launcher.setBuffer( bodyBuf->getBufferCL());
launcher.setBuffer( gpuCollidables.getBufferCL());
launcher.setBuffer( clAabbsWorldSpace.getBufferCL());
launcher.setBuffer( triangleConvexPairsOut.getBufferCL());
launcher.setBuffer( m_numConcavePairsOut.getBufferCL());
launcher.setBuffer( subTreesGPU->getBufferCL());
launcher.setBuffer( treeNodesGPU->getBufferCL());
launcher.setBuffer( bvhInfo->getBufferCL());
launcher.setConst( nPairs );
launcher.setConst( maxTriConvexPairCapacity);
int num = nPairs;
launcher.launch1D( num);
clFinish(m_queue);
numConcavePairs = m_numConcavePairsOut.at(0);
//printf("numConcavePairs=%d (max = %d\n",numConcavePairs,maxTriConvexPairCapacity);
launcher.setConst( nPairs );
launcher.setConst( maxTriConvexPairCapacity);
int num = nPairs;
launcher.launch1D( num);
clFinish(m_queue);
numConcavePairs = m_numConcavePairsOut.at(0);
} else
{
b3AlignedObjectArray<b3Int4> hostPairs;
pairs->copyToHost(hostPairs);
b3AlignedObjectArray<b3RigidBodyCL> hostBodyBuf;
bodyBuf->copyToHost(hostBodyBuf);
b3AlignedObjectArray<b3Collidable> hostCollidables;
gpuCollidables.copyToHost(hostCollidables);
b3AlignedObjectArray<b3Aabb> hostAabbsWorldSpace;
clAabbsWorldSpace.copyToHost(hostAabbsWorldSpace);
//int maxTriConvexPairCapacity,
b3AlignedObjectArray<b3Int4> triangleConvexPairsOutHost;
triangleConvexPairsOutHost.resize(maxTriConvexPairCapacity);
int numTriConvexPairsOutHost=0;
numConcavePairs = 0;
//m_numConcavePairsOut
b3AlignedObjectArray<b3QuantizedBvhNode> treeNodesCPU;
treeNodesGPU->copyToHost(treeNodesCPU);
b3AlignedObjectArray<b3BvhSubtreeInfo> subTreesCPU;
subTreesGPU->copyToHost(subTreesCPU);
b3AlignedObjectArray<b3BvhInfo> bvhInfoCPU;
bvhInfo->copyToHost(bvhInfoCPU);
//compute it...
volatile int hostNumConcavePairsOut=0;
//
for (int i=0;i<nPairs;i++)
{
b3BvhTraversal( &hostPairs.at(0),
&hostBodyBuf.at(0),
&hostCollidables.at(0),
&hostAabbsWorldSpace.at(0),
&triangleConvexPairsOutHost.at(0),
&hostNumConcavePairsOut,
&subTreesCPU.at(0),
&treeNodesCPU.at(0),
&bvhInfoCPU.at(0),
nPairs,
maxTriConvexPairCapacity,
i);
}
numConcavePairs = hostNumConcavePairsOut;
if (hostNumConcavePairsOut)
{
triangleConvexPairsOutHost.resize(hostNumConcavePairsOut);
triangleConvexPairsOut.copyFromHost(triangleConvexPairsOutHost);
}
//
m_numConcavePairsOut.resize(0);
m_numConcavePairsOut.push_back(numConcavePairs);
}
//printf("numConcavePairs=%d (max = %d\n",numConcavePairs,maxTriConvexPairCapacity);
if (numConcavePairs > maxTriConvexPairCapacity)
{
static int exceeded_maxTriConvexPairCapacity_count = 0;
b3Error("Rxceeded %d times the maxTriConvexPairCapacity (found %d but max is %d)\n", exceeded_maxTriConvexPairCapacity_count++,
numConcavePairs,maxTriConvexPairCapacity);
b3Error("Exceeded the maxTriConvexPairCapacity (found %d but max is %d, it happened %d times)\n",
numConcavePairs,maxTriConvexPairCapacity,exceeded_maxTriConvexPairCapacity_count++);
numConcavePairs = maxTriConvexPairCapacity;
}
triangleConvexPairsOut.resize(numConcavePairs);
if (numConcavePairs)
{
//now perform a SAT test for each triangle-convex element (stored in triangleConvexPairsOut)
B3_PROFILE("findConcaveSeparatingAxisKernel");
b3BufferInfoCL bInfo[] = {
b3BufferInfoCL( triangleConvexPairsOut.getBufferCL() ),
b3BufferInfoCL( bodyBuf->getBufferCL(),true),
b3BufferInfoCL( gpuCollidables.getBufferCL(),true),
b3BufferInfoCL( convexData.getBufferCL(),true),
b3BufferInfoCL( gpuVertices.getBufferCL(),true),
b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
b3BufferInfoCL( gpuFaces.getBufferCL(),true),
b3BufferInfoCL( gpuIndices.getBufferCL(),true),
b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
b3BufferInfoCL( m_concaveSepNormals.getBufferCL())
};
if (findConcaveSeparatingAxisKernelGPU)
{
//now perform a SAT test for each triangle-convex element (stored in triangleConvexPairsOut)
B3_PROFILE("findConcaveSeparatingAxisKernel");
b3BufferInfoCL bInfo[] = {
b3BufferInfoCL( triangleConvexPairsOut.getBufferCL() ),
b3BufferInfoCL( bodyBuf->getBufferCL(),true),
b3BufferInfoCL( gpuCollidables.getBufferCL(),true),
b3BufferInfoCL( convexData.getBufferCL(),true),
b3BufferInfoCL( gpuVertices.getBufferCL(),true),
b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
b3BufferInfoCL( gpuFaces.getBufferCL(),true),
b3BufferInfoCL( gpuIndices.getBufferCL(),true),
b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
b3BufferInfoCL( m_concaveSepNormals.getBufferCL())
};
b3LauncherCL launcher(m_queue, m_findConcaveSeparatingAxisKernel,"m_findConcaveSeparatingAxisKernel");
launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
b3LauncherCL launcher(m_queue, m_findConcaveSeparatingAxisKernel,"m_findConcaveSeparatingAxisKernel");
launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
launcher.setConst( numConcavePairs );
launcher.setConst( numConcavePairs );
int num = numConcavePairs;
launcher.launch1D( num);
clFinish(m_queue);
int num = numConcavePairs;
launcher.launch1D( num);
clFinish(m_queue);
} else
{
b3AlignedObjectArray<b3Int4> triangleConvexPairsOutHost;
triangleConvexPairsOut.copyToHost(triangleConvexPairsOutHost);
//triangleConvexPairsOutHost.resize(maxTriConvexPairCapacity);
b3AlignedObjectArray<b3RigidBodyCL> hostBodyBuf;
bodyBuf->copyToHost(hostBodyBuf);
b3AlignedObjectArray<b3Collidable> hostCollidables;
gpuCollidables.copyToHost(hostCollidables);
b3AlignedObjectArray<b3Aabb> hostAabbsWorldSpace;
clAabbsWorldSpace.copyToHost(hostAabbsWorldSpace);
b3AlignedObjectArray<b3ConvexPolyhedronCL> hostConvexData;
convexData.copyToHost(hostConvexData);
b3AlignedObjectArray<b3Vector3> hostVertices;
gpuVertices.copyToHost(hostVertices);
b3AlignedObjectArray<b3Vector3> hostUniqueEdges;
gpuUniqueEdges.copyToHost(hostUniqueEdges);
b3AlignedObjectArray<b3GpuFace> hostFaces;
gpuFaces.copyToHost(hostFaces);
b3AlignedObjectArray<int> hostIndices;
gpuIndices.copyToHost(hostIndices);
b3AlignedObjectArray<b3GpuChildShape> cpuChildShapes;
gpuChildShapes.copyToHost(cpuChildShapes);
//numConcavePairs
//b3BufferInfoCL( triangleConvexPairsOut.getBufferCL() ),
//b3BufferInfoCL( bodyBuf->getBufferCL(),true),
//b3BufferInfoCL( gpuCollidables.getBufferCL(),true),
// b3BufferInfoCL( convexData.getBufferCL(),true),
//b3BufferInfoCL( gpuVertices.getBufferCL(),true),
//b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
//b3BufferInfoCL( gpuFaces.getBufferCL(),true),
//b3BufferInfoCL( gpuIndices.getBufferCL(),true),
//b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
//b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
//b3BufferInfoCL( m_concaveSepNormals.getBufferCL())
b3AlignedObjectArray<b3Vector3> concaveSepNormalsHost;
m_concaveSepNormals.copyToHost(concaveSepNormalsHost);
}
// b3AlignedObjectArray<b3Vector3> cpuCompoundSepNormals;
// m_concaveSepNormals.copyToHost(cpuCompoundSepNormals);
// b3AlignedObjectArray<b3Int4> cpuConcavePairs;

36
src/Bullet3OpenCL/NarrowphaseCollision/b3QuantizedBvh.h
View File

@@ -41,6 +41,9 @@ class b3Serializer;
#define b3QuantizedBvhDataName "b3QuantizedBvhFloatData"
#endif
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3QuantizedBvhNodeData.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3BvhSubtreeInfoData.h"
//http://msdn.microsoft.com/library/default.asp?url=/library/en-us/vclang/html/vclrf__m128.asp
@@ -55,16 +58,10 @@ class b3Serializer;
///b3QuantizedBvhNode is a compressed aabb node, 16 bytes.
///Node can be used for leafnode or internal node. Leafnodes can point to 32-bit triangle index (non-negative range).
B3_ATTRIBUTE_ALIGNED16 (struct) b3QuantizedBvhNode
B3_ATTRIBUTE_ALIGNED16 (struct) b3QuantizedBvhNode : public b3QuantizedBvhNodeData
{
B3_DECLARE_ALIGNED_ALLOCATOR();
//12 bytes
unsigned short int m_quantizedAabbMin[3];
unsigned short int m_quantizedAabbMax[3];
//4 bytes
int m_escapeIndexOrTriangleIndex;
bool isLeafNode() const
{
//skipindex is negative (internal node), triangleindex >=0 (leafnode)
@@ -116,20 +113,11 @@ B3_ATTRIBUTE_ALIGNED16 (struct) b3OptimizedBvhNode
///b3BvhSubtreeInfo provides info to gather a subtree of limited size
B3_ATTRIBUTE_ALIGNED16(class) b3BvhSubtreeInfo
B3_ATTRIBUTE_ALIGNED16(class) b3BvhSubtreeInfo : public b3BvhSubtreeInfoData
{
public:
B3_DECLARE_ALIGNED_ALLOCATOR();
//12 bytes
unsigned short int m_quantizedAabbMin[3];
unsigned short int m_quantizedAabbMax[3];
//4 bytes, points to the root of the subtree
int m_rootNodeIndex;
//4 bytes
int m_subtreeSize;
int m_padding[3];
b3BvhSubtreeInfo()
{
//memset(&m_padding[0], 0, sizeof(m_padding));
@@ -501,14 +489,6 @@ private:
;
struct b3BvhSubtreeInfoData
{
int m_rootNodeIndex;
int m_subtreeSize;
unsigned short m_quantizedAabbMin[3];
unsigned short m_quantizedAabbMax[3];
};
struct b3OptimizedBvhNodeFloatData
{
b3Vector3FloatData m_aabbMinOrg;
@@ -530,12 +510,6 @@ struct b3OptimizedBvhNodeDoubleData
};
struct b3QuantizedBvhNodeData
{
unsigned short m_quantizedAabbMin[3];
unsigned short m_quantizedAabbMax[3];
int m_escapeIndexOrTriangleIndex;
};
struct b3QuantizedBvhFloatData
{

27
src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.cl
View File

@@ -34,33 +34,6 @@ typedef struct
} b3BvhInfo;
/*
bool isLeafNode() const
{
//skipindex is negative (internal node), triangleindex >=0 (leafnode)
return (m_escapeIndexOrTriangleIndex >= 0);
}
int getEscapeIndex() const
{
btAssert(!isLeafNode());
return -m_escapeIndexOrTriangleIndex;
}
int getTriangleIndex() const
{
btAssert(isLeafNode());
unsigned int x=0;
unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS);
// Get only the lower bits where the triangle index is stored
return (m_escapeIndexOrTriangleIndex&~(y));
}
int getPartId() const
{
btAssert(isLeafNode());
// Get only the highest bits where the part index is stored
return (m_escapeIndexOrTriangleIndex>>(31-MAX_NUM_PARTS_IN_BITS));
}
*/
int getTriangleIndex(const btQuantizedBvhNode* rootNode)
{
unsigned int x=0;

10
src/Bullet3OpenCL/NarrowphaseCollision/kernels/sat.cl
View File

@@ -401,7 +401,7 @@ bool findSeparatingAxisLocalA( const ConvexPolyhedronCL* hullA, __global const C
float4* sep,
float* dmin)
{
int i = get_global_id(0);
float4 posA = posA1;
posA.w = 0.f;
@@ -452,7 +452,7 @@ bool findSeparatingAxisLocalB( __global const ConvexPolyhedronCL* hullA, const
float4* sep,
float* dmin)
{
int i = get_global_id(0);
float4 posA = posA1;
posA.w = 0.f;
@@ -505,7 +505,7 @@ bool findSeparatingAxisEdgeEdgeLocalA( const ConvexPolyhedronCL* hullA, __global
float4* sep,
float* dmin)
{
int i = get_global_id(0);
float4 posA = posA1;
posA.w = 0.f;
@@ -607,7 +607,7 @@ bool findSeparatingAxis( __global const ConvexPolyhedronCL* hullA, __global cons
float4* sep,
float* dmin)
{
int i = get_global_id(0);
float4 posA = posA1;
posA.w = 0.f;
@@ -666,7 +666,7 @@ bool findSeparatingAxisEdgeEdge( __global const ConvexPolyhedronCL* hullA, __glo
float4* sep,
float* dmin)
{
int i = get_global_id(0);
float4 posA = posA1;
posA.w = 0.f;