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

change lcpp Lua preprocessor, to keep #defines and comments, remove empty lines

Browse Source 
remove duplicate data in b3Contact4 (now in btContact4Data shared between CPU/C++ and OpenCL)
OpenCL kernels use #include "Bullet3Collision/NarrowPhaseCollision/shared/b3Contact4Data.h"
Increase number of batches back to 250 (from 50), need to fix this hard coded number (see https://github.com/erwincoumans/bullet3/issues/12)
Work towards GJK/EPA, in addition to SAT/clipping (early on)
This commit is contained in:
erwincoumans
2013-08-08 12:24:09 -07:00
parent 46a08e3282
commit 3bf003ace1
50 changed files with 920 additions and 2731 deletions
Download Patch File Download Diff File Expand all files Collapse all files

200
src/Bullet3OpenCL/NarrowphaseCollision/b3ContactCache.cpp
View File

@@ -1,3 +1,4 @@
#if 0
/*
Bullet Continuous Collision Detection and Physics Library
@@ -18,77 +19,22 @@ subject to the following restrictions:
#include "b3ContactCache.h"
#include "Bullet3Common/b3Transform.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Contact4Data.h"
b3Scalar gContactBreakingThreshold = b3Scalar(0.02);
b3Scalar m_contactBreakingThreshold;
b3Scalar m_contactProcessingThreshold;
///gContactCalcArea3Points will approximate the convex hull area using 3 points
///when setting it to false, it will use 4 points to compute the area: it is more accurate but slower
bool gContactCalcArea3Points = true;
b3ContactCache::b3ContactCache()
:m_index1a(0)
{
}
#ifdef DEBUG_PERSISTENCY
#include <stdio.h>
void b3ContactCache::DebugPersistency()
{
int i;
printf("DebugPersistency : numPoints %d\n",m_cachedPoints);
for (i=0;i<m_cachedPoints;i++)
{
printf("m_pointCache[%d].m_userPersistentData = %x\n",i,m_pointCache[i].m_userPersistentData);
}
}
#endif //DEBUG_PERSISTENCY
void b3ContactCache::clearUserCache(btManifoldPoint& pt)
{
void* oldPtr = pt.m_userPersistentData;
if (oldPtr)
{
#ifdef DEBUG_PERSISTENCY
int i;
int occurance = 0;
for (i=0;i<m_cachedPoints;i++)
{
if (m_pointCache[i].m_userPersistentData == oldPtr)
{
occurance++;
if (occurance>1)
printf("error in clearUserCache\n");
}
}
btAssert(occurance<=0);
#endif //DEBUG_PERSISTENCY
if (pt.m_userPersistentData && gContactDestroyedCallback)
{
(*gContactDestroyedCallback)(pt.m_userPersistentData);
pt.m_userPersistentData = 0;
}
#ifdef DEBUG_PERSISTENCY
DebugPersistency();
#endif
}
}
static inline b3Scalar calcArea4Points(const btVector3 &p0,const btVector3 &p1,const btVector3 &p2,const btVector3 &p3)
static inline b3Scalar calcArea4Points(const b3Vector3 &p0,const b3Vector3 &p1,const b3Vector3 &p2,const b3Vector3 &p3)
{
// It calculates possible 3 area constructed from random 4 points and returns the biggest one.
btVector3 a[3],b[3];
b3Vector3 a[3],b[3];
a[0] = p0 - p1;
a[1] = p0 - p2;
a[2] = p0 - p3;
@@ -97,14 +43,16 @@ static inline b3Scalar calcArea4Points(const btVector3 &p0,const btVector3 &p1,c
b[2] = p1 - p2;
//todo: Following 3 cross production can be easily optimized by SIMD.
btVector3 tmp0 = a[0].cross(b[0]);
btVector3 tmp1 = a[1].cross(b[1]);
btVector3 tmp2 = a[2].cross(b[2]);
b3Vector3 tmp0 = a[0].cross(b[0]);
b3Vector3 tmp1 = a[1].cross(b[1]);
b3Vector3 tmp2 = a[2].cross(b[2]);
return btMax(btMax(tmp0.length2(),tmp1.length2()),tmp2.length2());
return b3Max(b3Max(tmp0.length2(),tmp1.length2()),tmp2.length2());
}
#if 0
int b3ContactCache::sortCachedPoints(const btManifoldPoint& pt)
//using localPointA for all points
int b3ContactCache::sortCachedPoints(const b3Vector3& pt)
{
//calculate 4 possible cases areas, and take biggest area
//also need to keep 'deepest'
@@ -129,32 +77,32 @@ int b3ContactCache::sortCachedPoints(const btManifoldPoint& pt)
{
if (maxPenetrationIndex != 0)
{
btVector3 a0 = pt.m_localPointA-m_pointCache[1].m_localPointA;
btVector3 b0 = m_pointCache[3].m_localPointA-m_pointCache[2].m_localPointA;
btVector3 cross = a0.cross(b0);
b3Vector3 a0 = pt.m_localPointA-m_pointCache[1].m_localPointA;
b3Vector3 b0 = m_pointCache[3].m_localPointA-m_pointCache[2].m_localPointA;
b3Vector3 cross = a0.cross(b0);
res0 = cross.length2();
}
if (maxPenetrationIndex != 1)
{
btVector3 a1 = pt.m_localPointA-m_pointCache[0].m_localPointA;
btVector3 b1 = m_pointCache[3].m_localPointA-m_pointCache[2].m_localPointA;
btVector3 cross = a1.cross(b1);
b3Vector3 a1 = pt.m_localPointA-m_pointCache[0].m_localPointA;
b3Vector3 b1 = m_pointCache[3].m_localPointA-m_pointCache[2].m_localPointA;
b3Vector3 cross = a1.cross(b1);
res1 = cross.length2();
}
if (maxPenetrationIndex != 2)
{
btVector3 a2 = pt.m_localPointA-m_pointCache[0].m_localPointA;
btVector3 b2 = m_pointCache[3].m_localPointA-m_pointCache[1].m_localPointA;
btVector3 cross = a2.cross(b2);
b3Vector3 a2 = pt.m_localPointA-m_pointCache[0].m_localPointA;
b3Vector3 b2 = m_pointCache[3].m_localPointA-m_pointCache[1].m_localPointA;
b3Vector3 cross = a2.cross(b2);
res2 = cross.length2();
}
if (maxPenetrationIndex != 3)
{
btVector3 a3 = pt.m_localPointA-m_pointCache[0].m_localPointA;
btVector3 b3 = m_pointCache[2].m_localPointA-m_pointCache[1].m_localPointA;
btVector3 cross = a3.cross(b3);
b3Vector3 a3 = pt.m_localPointA-m_pointCache[0].m_localPointA;
b3Vector3 b3 = m_pointCache[2].m_localPointA-m_pointCache[1].m_localPointA;
b3Vector3 cross = a3.cross(b3);
res3 = cross.length2();
}
}
@@ -176,23 +124,23 @@ int b3ContactCache::sortCachedPoints(const btManifoldPoint& pt)
res3 = calcArea4Points(pt.m_localPointA,m_pointCache[0].m_localPointA,m_pointCache[1].m_localPointA,m_pointCache[2].m_localPointA);
}
}
btVector4 maxvec(res0,res1,res2,res3);
b3Vector4 maxvec(res0,res1,res2,res3);
int biggestarea = maxvec.closestAxis4();
return biggestarea;
}
int b3ContactCache::getCacheEntry(const btManifoldPoint& newPoint) const
int b3ContactCache::getCacheEntry(const b3Vector3& newPoint) const
{
b3Scalar shortestDist = getContactBreakingThreshold() * getContactBreakingThreshold();
int size = getNumContacts();
int nearestPoint = -1;
for( int i = 0; i < size; i++ )
{
const btManifoldPoint &mp = m_pointCache[i];
const b3Vector3 &mp = m_pointCache[i];
btVector3 diffA = mp.m_localPointA- newPoint.m_localPointA;
b3Vector3 diffA = mp.m_localPointA- newPoint.m_localPointA;
const b3Scalar distToManiPoint = diffA.dot(diffA);
if( distToManiPoint < shortestDist )
{
@@ -203,12 +151,9 @@ int b3ContactCache::getCacheEntry(const btManifoldPoint& newPoint) const
return nearestPoint;
}
int b3ContactCache::addManifoldPoint(const btManifoldPoint& newPoint, bool isPredictive)
int b3ContactCache::addManifoldPoint(const b3Vector3& newPoint)
{
if (!isPredictive)
{
btAssert(validContactDistance(newPoint));
}
b3Assert(validContactDistance(newPoint));
int insertIndex = getNumContacts();
if (insertIndex == MANIFOLD_CACHE_SIZE)
@@ -230,73 +175,80 @@ int b3ContactCache::addManifoldPoint(const btManifoldPoint& newPoint, bool isPre
if (insertIndex<0)
insertIndex=0;
btAssert(m_pointCache[insertIndex].m_userPersistentData==0);
//b3Assert(m_pointCache[insertIndex].m_userPersistentData==0);
m_pointCache[insertIndex] = newPoint;
return insertIndex;
}
b3Scalar b3ContactCache::getContactBreakingThreshold() const
#endif
bool b3ContactCache::validContactDistance(const b3Vector3& pt)
{
return m_contactBreakingThreshold;
return pt.w <= gContactBreakingThreshold;
}
void b3ContactCache::refreshContactPoints(const btTransform& trA,const btTransform& trB)
void b3ContactCache::removeContactPoint(struct b3Contact4Data& newContactCache,int i)
{
int i;
#ifdef DEBUG_PERSISTENCY
printf("refreshContactPoints posA = (%f,%f,%f) posB = (%f,%f,%f)\n",
trA.getOrigin().getX(),
trA.getOrigin().getY(),
trA.getOrigin().getZ(),
trB.getOrigin().getX(),
trB.getOrigin().getY(),
trB.getOrigin().getZ());
#endif //DEBUG_PERSISTENCY
/// first refresh worldspace positions and distance
for (i=getNumContacts()-1;i>=0;i--)
int numContacts = b3Contact4Data_getNumPoints(&newContactCache);
if (i!=(numContacts-1))
{
btManifoldPoint &manifoldPoint = m_pointCache[i];
manifoldPoint.m_positionWorldOnA = trA( manifoldPoint.m_localPointA );
manifoldPoint.m_positionWorldOnB = trB( manifoldPoint.m_localPointB );
manifoldPoint.m_distance1 = (manifoldPoint.m_positionWorldOnA - manifoldPoint.m_positionWorldOnB).dot(manifoldPoint.m_normalWorldOnB);
manifoldPoint.m_lifeTime++;
b3Swap(newContactCache.m_localPosA[i],newContactCache.m_localPosA[numContacts-1]);
b3Swap(newContactCache.m_localPosB[i],newContactCache.m_localPosB[numContacts-1]);
b3Swap(newContactCache.m_worldPos[i],newContactCache.m_worldPos[numContacts-1]);
}
b3Contact4Data_setNumPoints(&newContactCache,numContacts-1);
}
void b3ContactCache::refreshContactPoints(const b3Transform& trA,const b3Transform& trB, struct b3Contact4Data& contacts)
{
int numContacts = b3Contact4Data_getNumPoints(&contacts);
int i;
/// first refresh worldspace positions and distance
for (i=numContacts-1;i>=0;i--)
{
b3Vector3 worldPosA = trA( contacts.m_localPosA[i]);
b3Vector3 worldPosB = trB( contacts.m_localPosB[i]);
contacts.m_worldPos[i] = worldPosB;
float distance = (worldPosA - worldPosB).dot(contacts.m_worldNormal);
contacts.m_worldPos[i].w = distance;
}
/// then
b3Scalar distance2d;
btVector3 projectedDifference,projectedPoint;
for (i=getNumContacts()-1;i>=0;i--)
b3Vector3 projectedDifference,projectedPoint;
for (i=numContacts-1;i>=0;i--)
{
btManifoldPoint &manifoldPoint = m_pointCache[i];
b3Vector3 worldPosA = trA( contacts.m_localPosA[i]);
b3Vector3 worldPosB = trB( contacts.m_localPosB[i]);
b3Vector3&pt = contacts.m_worldPos[i];
//contact becomes invalid when signed distance exceeds margin (projected on contactnormal direction)
if (!validContactDistance(manifoldPoint))
if (!validContactDistance(pt))
{
removeContactPoint(i);
removeContactPoint(contacts,i);
} else
{
//contact also becomes invalid when relative movement orthogonal to normal exceeds margin
projectedPoint = manifoldPoint.m_positionWorldOnA - manifoldPoint.m_normalWorldOnB * manifoldPoint.m_distance1;
projectedDifference = manifoldPoint.m_positionWorldOnB - projectedPoint;
projectedPoint = contacts.m_worldPos[i] - contacts.m_worldNormal * contacts.m_worldPos[i].w;
projectedDifference = contacts.m_worldPos[i] - projectedPoint;
distance2d = projectedDifference.dot(projectedDifference);
if (distance2d > getContactBreakingThreshold()*getContactBreakingThreshold() )
if (distance2d > gContactBreakingThreshold*gContactBreakingThreshold )
{
removeContactPoint(i);
removeContactPoint(contacts,i);
} else
{
//contact point processed callback
if (gContactProcessedCallback)
(*gContactProcessedCallback)(manifoldPoint,(void*)m_body0,(void*)m_body1);
////contact point processed callback
//if (gContactProcessedCallback)
// (*gContactProcessedCallback)(manifoldPoint,(void*)m_body0,(void*)m_body1);
}
}
}
#ifdef DEBUG_PERSISTENCY
DebugPersistency();
#endif //
}
}

173
src/Bullet3OpenCL/NarrowphaseCollision/b3ContactCache.h
View File

@@ -1,7 +1,7 @@
/*
Bullet Continuous Collision Detection and Physics Library
Copyright (c) 2003-2006 Erwin Coumans http://continuousphysics.com/Bullet/
Copyright (c) 2003-2013 Erwin Coumans http://bulletphysics.org
This software is provided 'as-is', without any express or implied warranty.
In no event will the authors be held liable for any damages arising from the use of this software.
@@ -18,23 +18,15 @@ subject to the following restrictions:
#define B3_CONTACT_CACHE_H
#include "LinearMath/btVector3.h"
#include "LinearMath/btTransform.h"
#include "btManifoldPoint.h"
class btCollisionObject;
#include "LinearMath/btAlignedAllocator.h"
#include "Bullet3Common/b3Vector3.h"
#include "Bullet3Common/b3Transform.h"
#include "Bullet3Common/b3AlignedAllocator.h"
struct btCollisionResult;
///maximum contact breaking and merging threshold
extern b3Scalar gContactBreakingThreshold;
//the enum starts at 1024 to avoid type conflicts with btTypedConstraint
enum btContactManifoldTypes
{
MIN_CONTACT_MANIFOLD_TYPE = 1024,
BT_PERSISTENT_MANIFOLD_TYPE
};
#define MANIFOLD_CACHE_SIZE 4
@@ -45,8 +37,6 @@ enum btContactManifoldTypes
///reduces the cache to 4 points, when more then 4 points are added, using following rules:
///the contact point with deepest penetration is always kept, and it tries to maximuze the area covered by the points
///note that some pairs of objects might have more then one contact manifold.
B3_ATTRIBUTE_ALIGNED16( class) b3ContactCache
{
@@ -54,163 +44,36 @@ B3_ATTRIBUTE_ALIGNED16( class) b3ContactCache
/// sort cached points so most isolated points come first
int sortCachedPoints(const btManifoldPoint& pt);
int sortCachedPoints(const b3Vector3& pt);
int findContactPoint(const btManifoldPoint* unUsed, int numUnused,const btManifoldPoint& pt);
public:
BT_DECLARE_ALIGNED_ALLOCATOR();
int m_index1a;
b3ContactCache();
b3ContactCache(const btCollisionObject* body0,const btCollisionObject* body1,int , b3Scalar contactBreakingThreshold,b3Scalar contactProcessingThreshold)
: btTypedObject(BT_PERSISTENT_MANIFOLD_TYPE),
m_body0(body0),m_body1(body1),m_cachedPoints(0),
m_contactBreakingThreshold(contactBreakingThreshold),
m_contactProcessingThreshold(contactProcessingThreshold)
{
}
B3_FORCE_INLINE const btCollisionObject* getBody0() const { return m_body0;}
B3_FORCE_INLINE const btCollisionObject* getBody1() const { return m_body1;}
void setBodies(const btCollisionObject* body0,const btCollisionObject* body1)
{
m_body0 = body0;
m_body1 = body1;
}
void clearUserCache(btManifoldPoint& pt);
#ifdef DEBUG_PERSISTENCY
void DebugPersistency();
#endif //
B3_FORCE_INLINE int getNumContacts() const { return m_cachedPoints;}
/// the setNumContacts API is usually not used, except when you gather/fill all contacts manually
void setNumContacts(int cachedPoints)
{
m_cachedPoints = cachedPoints;
}
B3_FORCE_INLINE const btManifoldPoint& getContactPoint(int index) const
{
btAssert(index < m_cachedPoints);
return m_pointCache[index];
}
B3_FORCE_INLINE btManifoldPoint& getContactPoint(int index)
{
btAssert(index < m_cachedPoints);
return m_pointCache[index];
}
B3_DECLARE_ALIGNED_ALLOCATOR();
void setContactBreakingThreshold(b3Scalar contactBreakingThreshold)
{
m_contactBreakingThreshold = contactBreakingThreshold;
}
void setContactProcessingThreshold(b3Scalar contactProcessingThreshold)
{
m_contactProcessingThreshold = contactProcessingThreshold;
}
int addManifoldPoint( const b3Vector3& newPoint);
int getCacheEntry(const btManifoldPoint& newPoint) const;
int addManifoldPoint( const btManifoldPoint& newPoint, bool isPredictive=false);
void removeContactPoint (int index)
/*void replaceContactPoint(const b3Vector3& newPoint,int insertIndex)
{
clearUserCache(m_pointCache[index]);
int lastUsedIndex = getNumContacts() - 1;
// m_pointCache[index] = m_pointCache[lastUsedIndex];
if(index != lastUsedIndex)
{
m_pointCache[index] = m_pointCache[lastUsedIndex];
//get rid of duplicated userPersistentData pointer
m_pointCache[lastUsedIndex].m_userPersistentData = 0;
m_pointCache[lastUsedIndex].m_appliedImpulse = 0.f;
m_pointCache[lastUsedIndex].m_lateralFrictionInitialized = false;
m_pointCache[lastUsedIndex].m_appliedImpulseLateral1 = 0.f;
m_pointCache[lastUsedIndex].m_appliedImpulseLateral2 = 0.f;
m_pointCache[lastUsedIndex].m_lifeTime = 0;
}
btAssert(m_pointCache[lastUsedIndex].m_userPersistentData==0);
m_cachedPoints--;
}
void replaceContactPoint(const btManifoldPoint& newPoint,int insertIndex)
{
btAssert(validContactDistance(newPoint));
#define MAINTAIN_PERSISTENCY 1
#ifdef MAINTAIN_PERSISTENCY
int lifeTime = m_pointCache[insertIndex].getLifeTime();
b3Scalar appliedImpulse = m_pointCache[insertIndex].m_appliedImpulse;
b3Scalar appliedLateralImpulse1 = m_pointCache[insertIndex].m_appliedImpulseLateral1;
b3Scalar appliedLateralImpulse2 = m_pointCache[insertIndex].m_appliedImpulseLateral2;
// bool isLateralFrictionInitialized = m_pointCache[insertIndex].m_lateralFrictionInitialized;
btAssert(lifeTime>=0);
void* cache = m_pointCache[insertIndex].m_userPersistentData;
b3Assert(validContactDistance(newPoint));
m_pointCache[insertIndex] = newPoint;
m_pointCache[insertIndex].m_userPersistentData = cache;
m_pointCache[insertIndex].m_appliedImpulse = appliedImpulse;
m_pointCache[insertIndex].m_appliedImpulseLateral1 = appliedLateralImpulse1;
m_pointCache[insertIndex].m_appliedImpulseLateral2 = appliedLateralImpulse2;
m_pointCache[insertIndex].m_appliedImpulse = appliedImpulse;
m_pointCache[insertIndex].m_appliedImpulseLateral1 = appliedLateralImpulse1;
m_pointCache[insertIndex].m_appliedImpulseLateral2 = appliedLateralImpulse2;
m_pointCache[insertIndex].m_lifeTime = lifeTime;
#else
clearUserCache(m_pointCache[insertIndex]);
m_pointCache[insertIndex] = newPoint;
#endif
}
*/
bool validContactDistance(const btManifoldPoint& pt) const
{
return pt.m_distance1 <= getContactBreakingThreshold();
}
static bool validContactDistance(const b3Vector3& pt);
/// calculated new worldspace coordinates and depth, and reject points that exceed the collision margin
void refreshContactPoints( const btTransform& trA,const btTransform& trB);
static void refreshContactPoints( const b3Transform& trA,const b3Transform& trB, struct b3Contact4Data& newContactCache);
static void removeContactPoint(struct b3Contact4Data& newContactCache,int i);
B3_FORCE_INLINE void clearManifold()
{
int i;
for (i=0;i<m_cachedPoints;i++)
{
clearUserCache(m_pointCache[i]);
}
m_cachedPoints = 0;
}
}
;
};

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

@@ -19,13 +19,14 @@ subject to the following restrictions:
///And contact clipping based on work from Simon Hobbs
//#define B3_DEBUG_SAT_FACE
//#define CHECK_ON_HOST
int b3g_actualSATPairTests=0;
#include "b3ConvexHullContact.h"
#include <string.h>//memcpy
#include "b3ConvexPolyhedronCL.h"
#include "Bullet3OpenCL/NarrowphaseCollision/b3ContactCache.h"
typedef b3AlignedObjectArray<b3Vector3> b3VertexArray;
@@ -1603,7 +1604,7 @@ int computeContactConvexConvex( b3AlignedObjectArray<b3Int4>& pairs,
transB.setRotation(rigidBodies[bodyIndexB].m_quat);
float maximumDistanceSquared = 1e30f;
b3Vector3 resultPointOnB;
b3Vector3 resultPointOnBWorld;
b3Vector3 sepAxis2(0,1,0);
b3Scalar distance2 = 1e30f;
@@ -1618,7 +1619,7 @@ int computeContactConvexConvex( b3AlignedObjectArray<b3Int4>& pairs,
maximumDistanceSquared,
sepAxis2,
distance2,
resultPointOnB);
resultPointOnBWorld);
if (result2)
@@ -1627,31 +1628,58 @@ int computeContactConvexConvex( b3AlignedObjectArray<b3Int4>& pairs,
{
contactIndex = nGlobalContactsOut;
globalContactsOut.expand();
b3Contact4& contact = globalContactsOut.at(nGlobalContactsOut);
contact.m_batchIdx = 0;//i;
contact.m_bodyAPtrAndSignBit = (rigidBodies.at(bodyIndexA).m_invMass==0)? -bodyIndexA:bodyIndexA;
contact.m_bodyBPtrAndSignBit = (rigidBodies.at(bodyIndexB).m_invMass==0)? -bodyIndexB:bodyIndexB;
b3Contact4& newContact = globalContactsOut.at(nGlobalContactsOut);
newContact.m_batchIdx = 0;//i;
newContact.m_bodyAPtrAndSignBit = (rigidBodies.at(bodyIndexA).m_invMass==0)? -bodyIndexA:bodyIndexA;
newContact.m_bodyBPtrAndSignBit = (rigidBodies.at(bodyIndexB).m_invMass==0)? -bodyIndexB:bodyIndexB;
contact.m_frictionCoeffCmp = 45874;
contact.m_restituitionCoeffCmp = 0;
newContact.m_frictionCoeffCmp = 45874;
newContact.m_restituitionCoeffCmp = 0;
int numPoints = 1;
if (pairs[pairIndex].z>=0)
int numPoints = 0;
if (0)//pairs[pairIndex].z>=0)
{
printf("add existing points?\n");
//printf("add existing points?\n");
//refresh
}
for (int p=0;p<numPoints;p++)
{
resultPointOnB.w = distance2;
int numOldPoints = oldContacts[pairs[pairIndex].z].getNPoints();
if (numOldPoints)
{
newContact = oldContacts[pairs[pairIndex].z];
//b3ContactCache::refreshContactPoints(transA,transB,newContact);
}
numPoints = b3Contact4Data_getNumPoints(&newContact);
contact.m_worldPos[p] = resultPointOnB;
contact.m_worldNormal = -sepAxis2;
}
/*
int insertIndex = m_manifoldPtr->getCacheEntry(newPt);
if (insertIndex >= 0)
{
//const btManifoldPoint& oldPoint = m_manifoldPtr->getContactPoint(insertIndex);
m_manifoldPtr->replaceContactPoint(newPt,insertIndex);
} else
{
insertIndex = m_manifoldPtr->addManifoldPoint(newPt);
}
*/
int p=numPoints;
if (numPoints<3)
{
numPoints++;
}
{
resultPointOnBWorld.w = distance2;
newContact.m_worldPos[p] = resultPointOnBWorld;
b3Vector3 resultPointOnAWorld = resultPointOnBWorld+distance2*sepAxis2;
//newContact.m_localPosA[p] = transA.inverse()*resultPointOnAWorld;
// newContact.m_localPosB[p] = transB.inverse()*resultPointOnBWorld;
newContact.m_worldNormal = sepAxis2;
}
//printf("bodyIndexA %d,bodyIndexB %d,normal=%f,%f,%f numPoints %d\n",bodyIndexA,bodyIndexB,normalOnSurfaceB.x,normalOnSurfaceB.y,normalOnSurfaceB.z,numPoints);
contact.m_worldNormal.w = (b3Scalar)numPoints;
newContact.m_worldNormal.w = (b3Scalar)numPoints;
nGlobalContactsOut++;
} else
{
@@ -1797,7 +1825,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
return;
//#define CHECK_ON_HOST
#ifdef CHECK_ON_HOST
b3AlignedObjectArray<b3YetAnotherAabb> hostAabbs;
clAabbsWS.copyToHost(hostAabbs);
@@ -1909,9 +1937,12 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
hostCollidables[collidableIndexB].m_shapeType == SHAPE_CONVEX_HULL)
{
//printf("hostPairs[i].z=%d\n",hostPairs[i].z);
int contactIndex = computeContactConvexConvex(hostPairs,i,bodyIndexA,bodyIndexB,collidableIndexA,collidableIndexB,hostBodyBuf,
hostCollidables,hostConvexData,hostVertices,hostUniqueEdges,hostIndices,hostFaces,hostContacts,nContacts,maxContactCapacity,
oldHostContacts);
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);
if (contactIndex>=0)
{

29
src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.h
View File

@@ -2,17 +2,13 @@
static const char* bvhTraversalKernelCL= \
"//keep this enum in sync with the CPU version (in btCollidable.h)\n"
"//written by Erwin Coumans\n"
"\n"
"#define SHAPE_CONVEX_HULL 3\n"
"#define SHAPE_CONCAVE_TRIMESH 5\n"
"#define TRIANGLE_NUM_CONVEX_FACES 5\n"
"#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6\n"
"#define SHAPE_SPHERE 7\n"
"\n"
"typedef unsigned int u32;\n"
"\n"
"#define MAX_NUM_PARTS_IN_BITS 10\n"
"\n"
"///btQuantizedBvhNode is a compressed aabb node, 16 bytes.\n"
"///Node can be used for leafnode or internal node. Leafnodes can point to 32-bit triangle index (non-negative range).\n"
"typedef struct\n"
@@ -23,7 +19,6 @@ static const char* bvhTraversalKernelCL= \
" //4 bytes\n"
" int m_escapeIndexOrTriangleIndex;\n"
"} btQuantizedBvhNode;\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_aabbMin;\n"
@@ -33,9 +28,7 @@ static const char* bvhTraversalKernelCL= \
" int m_numSubTrees;\n"
" int m_nodeOffset;\n"
" int m_subTreeOffset;\n"
"\n"
"} b3BvhInfo;\n"
"\n"
"/*\n"
" bool isLeafNode() const\n"
" {\n"
@@ -62,7 +55,6 @@ static const char* bvhTraversalKernelCL= \
" return (m_escapeIndexOrTriangleIndex>>(31-MAX_NUM_PARTS_IN_BITS));\n"
" }\n"
"*/\n"
"\n"
"int getTriangleIndex(const btQuantizedBvhNode* rootNode)\n"
"{\n"
" unsigned int x=0;\n"
@@ -70,7 +62,6 @@ static const char* bvhTraversalKernelCL= \
" // Get only the lower bits where the triangle index is stored\n"
" return (rootNode->m_escapeIndexOrTriangleIndex&~(y));\n"
"}\n"
"\n"
"int isLeaf(const btQuantizedBvhNode* rootNode)\n"
"{\n"
" //skipindex is negative (internal node), triangleindex >=0 (leafnode)\n"
@@ -81,7 +72,6 @@ static const char* bvhTraversalKernelCL= \
"{\n"
" return -rootNode->m_escapeIndexOrTriangleIndex;\n"
"}\n"
"\n"
"typedef struct\n"
"{\n"
" //12 bytes\n"
@@ -93,7 +83,6 @@ static const char* bvhTraversalKernelCL= \
" int m_subtreeSize;\n"
" int m_padding[3];\n"
"} btBvhSubtreeInfo;\n"
"\n"
"///keep this in sync with btCollidable.h\n"
"typedef struct\n"
"{\n"
@@ -103,7 +92,6 @@ static const char* bvhTraversalKernelCL= \
" int m_shapeIndex;\n"
" \n"
"} btCollidableGpu;\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_childPosition;\n"
@@ -113,21 +101,17 @@ static const char* bvhTraversalKernelCL= \
" int m_unused1;\n"
" int m_unused2;\n"
"} btGpuChildShape;\n"
"\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_pos;\n"
" float4 m_quat;\n"
" float4 m_linVel;\n"
" float4 m_angVel;\n"
"\n"
" u32 m_collidableIdx;\n"
" float m_invMass;\n"
" float m_restituitionCoeff;\n"
" float m_frictionCoeff;\n"
"} BodyData;\n"
"\n"
"typedef struct \n"
"{\n"
" union\n"
@@ -143,8 +127,6 @@ static const char* bvhTraversalKernelCL= \
" int m_maxIndices[4];\n"
" };\n"
"} btAabbCL;\n"
"\n"
"\n"
"int testQuantizedAabbAgainstQuantizedAabb(\n"
" const unsigned short int* aabbMin1,\n"
" const unsigned short int* aabbMax1,\n"
@@ -170,13 +152,10 @@ static const char* bvhTraversalKernelCL= \
" //overlap = ((aabbMin1[1] > aabbMax2[1]) || (aabbMax1[1] < aabbMin2[1])) ? 0 : overlap;\n"
" //return overlap;\n"
"}\n"
"\n"
"\n"
"void quantizeWithClamp(unsigned short* out, float4 point2,int isMax, float4 bvhAabbMin, float4 bvhAabbMax, float4 bvhQuantization)\n"
"{\n"
" float4 clampedPoint = max(point2,bvhAabbMin);\n"
" clampedPoint = min (clampedPoint, bvhAabbMax);\n"
"\n"
" float4 v = (clampedPoint - bvhAabbMin) * bvhQuantization;\n"
" if (isMax)\n"
" {\n"
@@ -189,10 +168,7 @@ static const char* bvhTraversalKernelCL= \
" out[1] = (unsigned short) (((unsigned short)(v.y) & 0xfffe));\n"
" out[2] = (unsigned short) (((unsigned short)(v.z) & 0xfffe));\n"
" }\n"
"\n"
"}\n"
"\n"
"\n"
"// work-in-progress\n"
"__kernel void bvhTraversalKernel( __global const int4* pairs, \n"
" __global const BodyData* rigidBodies, \n"
@@ -223,7 +199,6 @@ static const char* bvhTraversalKernelCL= \
" \n"
" if (collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH)\n"
" return;\n"
"\n"
" int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
" \n"
" if (shapeTypeB!=SHAPE_CONVEX_HULL &&\n"
@@ -231,9 +206,7 @@ static const char* bvhTraversalKernelCL= \
" shapeTypeB!=SHAPE_COMPOUND_OF_CONVEX_HULLS\n"
" )\n"
" return;\n"
"\n"
" b3BvhInfo bvhInfo = bvhInfos[collidables[collidableIndexA].m_numChildShapes];\n"
"\n"
" float4 bvhAabbMin = bvhInfo.m_aabbMin;\n"
" float4 bvhAabbMax = bvhInfo.m_aabbMax;\n"
" float4 bvhQuantization = bvhInfo.m_quantization;\n"
@@ -241,7 +214,6 @@ static const char* bvhTraversalKernelCL= \
" __global const btBvhSubtreeInfo* subtreeHeaders = &subtreeHeadersRoot[bvhInfo.m_subTreeOffset];\n"
" __global const btQuantizedBvhNode* quantizedNodes = &quantizedNodesRoot[bvhInfo.m_nodeOffset];\n"
" \n"
"\n"
" unsigned short int quantizedQueryAabbMin[3];\n"
" unsigned short int quantizedQueryAabbMax[3];\n"
" quantizeWithClamp(quantizedQueryAabbMin,aabbs[bodyIndexB].m_min,false,bvhAabbMin, bvhAabbMax,bvhQuantization);\n"
@@ -308,6 +280,5 @@ static const char* bvhTraversalKernelCL= \
" }\n"
" }\n"
" }\n"
"\n"
"}\n"
;

34
src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl
View File

@@ -345,7 +345,7 @@ void computeContactSphereConvex(int pairIndex,
__global const float4* convexVertices,
__global const int* convexIndices,
__global const btGpuFace* faces,
__global b3Contact4Data* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int maxContactCapacity,
float4 spherePos2,
@@ -466,9 +466,9 @@ void computeContactSphereConvex(int pairIndex,
if (1)//dstIdx < maxContactCapacity)
{
__global b3Contact4Data* c = &globalContactsOut[dstIdx];
__global struct b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = normalOnSurfaceB1;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = pairIndex;
c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
@@ -590,7 +590,7 @@ int computeContactPlaneConvex(int pairIndex,
__global const float4* convexVertices,
__global const int* convexIndices,
__global const btGpuFace* faces,
__global b3Contact4Data* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int maxContactCapacity,
float4 posB,
@@ -692,11 +692,11 @@ int computeContactPlaneConvex(int pairIndex,
if (dstIdx < maxContactCapacity)
{
resultIndex = dstIdx;
__global b3Contact4Data* c = &globalContactsOut[dstIdx];
__global struct b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = planeNormalWorld;
//c->setFrictionCoeff(0.7);
//c->setRestituitionCoeff(0.f);
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = pairIndex;
c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
@@ -732,7 +732,7 @@ void computeContactPlaneSphere(int pairIndex,
__global const BodyData* rigidBodies,
__global const btCollidableGpu* collidables,
__global const btGpuFace* faces,
__global b3Contact4Data* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int maxContactCapacity)
{
@@ -775,9 +775,9 @@ void computeContactPlaneSphere(int pairIndex,
if (dstIdx < maxContactCapacity)
{
__global b3Contact4Data* c = &globalContactsOut[dstIdx];
__global struct b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = normalOnSurfaceB1;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = pairIndex;
c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
@@ -798,7 +798,7 @@ __kernel void primitiveContactsKernel( __global int4* pairs,
__global const float4* uniqueEdges,
__global const btGpuFace* faces,
__global const int* indices,
__global b3Contact4Data* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int numPairs, int maxContactCapacity)
{
@@ -953,9 +953,9 @@ __kernel void primitiveContactsKernel( __global int4* pairs,
if (dstIdx < maxContactCapacity)
{
__global b3Contact4Data* c = &globalContactsOut[dstIdx];
__global struct b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = -normalOnSurfaceB;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = pairIndex;
int bodyA = pairs[pairIndex].x;
int bodyB = pairs[pairIndex].y;
@@ -987,7 +987,7 @@ __kernel void processCompoundPairsPrimitivesKernel( __global const int4* gpuCo
__global const int* indices,
__global btAabbCL* aabbs,
__global const btGpuChildShape* gpuChildShapes,
__global b3Contact4Data* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int numCompoundPairs, int maxContactCapacity
)
@@ -1166,7 +1166,7 @@ void computeContactSphereTriangle(int pairIndex,
__global const BodyData* rigidBodies,
__global const btCollidableGpu* collidables,
const float4* triangleVertices,
__global b3Contact4Data* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int maxContactCapacity,
float4 spherePos2,
@@ -1293,9 +1293,9 @@ void computeContactSphereTriangle(int pairIndex,
if (dstIdx < maxContactCapacity)
{
__global b3Contact4Data* c = &globalContactsOut[dstIdx];
__global struct b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = normalOnSurfaceB1;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = pairIndex;
c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
@@ -1325,7 +1325,7 @@ __kernel void findConcaveSphereContactsKernel( __global int4* concavePairs,
__global const btGpuFace* faces,
__global const int* indices,
__global btAabbCL* aabbs,
__global b3Contact4Data* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int numConcavePairs, int maxContactCapacity
)

281
src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h
View File

@@ -1,67 +1,62 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
static const char* primitiveContactsKernelsCL= \
"\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"#ifndef B3_CONTACT4DATA_H\n"
"#define B3_CONTACT4DATA_H\n"
"#ifndef B3_FLOAT4_H\n"
"#define B3_FLOAT4_H\n"
"#ifndef B3_PLATFORM_DEFINITIONS_H\n"
"#define B3_PLATFORM_DEFINITIONS_H\n"
"struct MyTest\n"
"{\n"
" int bla;\n"
"};\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"#endif\n"
"#ifdef __cplusplus\n"
"#else//bla\n"
" typedef float4 b3Float4;\n"
"\n"
"\n"
"\n"
"\n"
"typedef struct\n"
"#endif \n"
"#endif //B3_FLOAT4_H\n"
"typedef struct b3Contact4Data b3Contact4Data_t;\n"
"struct b3Contact4Data\n"
"{\n"
" b3Float4 m_worldPos[4];\n"
"// b3Float4 m_localPosB[4];\n"
" b3Float4 m_worldNormal; // w: m_nPoints\n"
" unsigned int m_coeffs;\n"
" unsigned int m_batchIdx;\n"
"\n"
" unsigned short m_restituitionCoeffCmp;\n"
" unsigned short m_frictionCoeffCmp;\n"
" int m_batchIdx;\n"
" int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr\n"
" int m_bodyBPtrAndSignBit;\n"
"\n"
" int m_childIndexA;\n"
" int m_childIndexB;\n"
" int m_unused1;\n"
" int m_unused2;\n"
"\n"
"} b3Contact4Data;\n"
"\n"
"\n"
"\n"
" b3Float4 m_localPosA;\n"
"};\n"
"inline int b3Contact4Data_getNumPoints(const struct b3Contact4Data* contact)\n"
"{\n"
" return (int)contact->m_worldNormal.w;\n"
"};\n"
"inline void b3Contact4Data_setNumPoints(struct b3Contact4Data* contact, int numPoints)\n"
"{\n"
" contact->m_worldNormal.w = (float)numPoints;\n"
"};\n"
"#endif //B3_CONTACT4DATA_H\n"
"#define SHAPE_CONVEX_HULL 3\n"
"#define SHAPE_PLANE 4\n"
"#define SHAPE_CONCAVE_TRIMESH 5\n"
"#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6\n"
"#define SHAPE_SPHERE 7\n"
"\n"
"\n"
"#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
"#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
"#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n"
"#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable\n"
"#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable\n"
"\n"
"#ifdef cl_ext_atomic_counters_32\n"
"#pragma OPENCL EXTENSION cl_ext_atomic_counters_32 : enable\n"
"#else\n"
"#define counter32_t volatile __global int*\n"
"#endif\n"
"\n"
"#define GET_GROUP_IDX get_group_id(0)\n"
"#define GET_LOCAL_IDX get_local_id(0)\n"
"#define GET_GLOBAL_IDX get_global_id(0)\n"
@@ -75,15 +70,9 @@ static const char* primitiveContactsKernelsCL= \
"#define AtomAdd(x, value) atom_add(&(x), value)\n"
"#define AtomCmpxhg(x, cmp, value) atom_cmpxchg( &(x), cmp, value )\n"
"#define AtomXhg(x, value) atom_xchg ( &(x), value )\n"
"\n"
"#define max2 max\n"
"#define min2 min\n"
"\n"
"typedef unsigned int u32;\n"
"\n"
"\n"
"\n"
"\n"
"typedef struct \n"
"{\n"
" union\n"
@@ -99,7 +88,6 @@ static const char* primitiveContactsKernelsCL= \
" int m_maxIndices[4];\n"
" };\n"
"} btAabbCL;\n"
"\n"
"///keep this in sync with btCollidable.h\n"
"typedef struct\n"
"{\n"
@@ -109,7 +97,6 @@ static const char* primitiveContactsKernelsCL= \
" int m_shapeIndex;\n"
" \n"
"} btCollidableGpu;\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_childPosition;\n"
@@ -119,23 +106,18 @@ static const char* primitiveContactsKernelsCL= \
" int m_unused1;\n"
" int m_unused2;\n"
"} btGpuChildShape;\n"
"\n"
"#define GET_NPOINTS(x) (x).m_worldNormal.w\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_pos;\n"
" float4 m_quat;\n"
" float4 m_linVel;\n"
" float4 m_angVel;\n"
"\n"
" u32 m_collidableIdx; \n"
" float m_invMass;\n"
" float m_restituitionCoeff;\n"
" float m_frictionCoeff;\n"
"} BodyData;\n"
"\n"
"\n"
"typedef struct \n"
"{\n"
" float4 m_localCenter;\n"
@@ -152,48 +134,37 @@ static const char* primitiveContactsKernelsCL= \
" int m_uniqueEdgesOffset;\n"
" int m_numUniqueEdges;\n"
" int m_unused;\n"
"\n"
"} ConvexPolyhedronCL;\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_plane;\n"
" int m_indexOffset;\n"
" int m_numIndices;\n"
"} btGpuFace;\n"
"\n"
"#define SELECT_UINT4( b, a, condition ) select( b,a,condition )\n"
"\n"
"#define make_float4 (float4)\n"
"#define make_float2 (float2)\n"
"#define make_uint4 (uint4)\n"
"#define make_int4 (int4)\n"
"#define make_uint2 (uint2)\n"
"#define make_int2 (int2)\n"
"\n"
"\n"
"__inline\n"
"float fastDiv(float numerator, float denominator)\n"
"{\n"
" return native_divide(numerator, denominator); \n"
"// return numerator/denominator; \n"
"}\n"
"\n"
"__inline\n"
"float4 fastDiv4(float4 numerator, float4 denominator)\n"
"{\n"
" return native_divide(numerator, denominator); \n"
"}\n"
"\n"
"\n"
"__inline\n"
"float4 cross3(float4 a, float4 b)\n"
"{\n"
" return cross(a,b);\n"
"}\n"
"\n"
"//#define dot3F4 dot\n"
"\n"
"__inline\n"
"float dot3F4(float4 a, float4 b)\n"
"{\n"
@@ -201,35 +172,23 @@ static const char* primitiveContactsKernelsCL= \
" float4 b1 = make_float4(b.xyz,0.f);\n"
" return dot(a1, b1);\n"
"}\n"
"\n"
"__inline\n"
"float4 fastNormalize4(float4 v)\n"
"{\n"
" return fast_normalize(v);\n"
"}\n"
"\n"
"\n"
"///////////////////////////////////////\n"
"// Quaternion\n"
"///////////////////////////////////////\n"
"\n"
"typedef float4 Quaternion;\n"
"\n"
"__inline\n"
"Quaternion qtMul(Quaternion a, Quaternion b);\n"
"\n"
"__inline\n"
"Quaternion qtNormalize(Quaternion in);\n"
"\n"
"__inline\n"
"float4 qtRotate(Quaternion q, float4 vec);\n"
"\n"
"__inline\n"
"Quaternion qtInvert(Quaternion q);\n"
"\n"
"\n"
"\n"
"\n"
"__inline\n"
"Quaternion qtMul(Quaternion a, Quaternion b)\n"
"{\n"
@@ -240,7 +199,6 @@ static const char* primitiveContactsKernelsCL= \
" ans.w = a.w*b.w - dot3F4(a, b);\n"
" return ans;\n"
"}\n"
"\n"
"__inline\n"
"Quaternion qtNormalize(Quaternion in)\n"
"{\n"
@@ -257,32 +215,27 @@ static const char* primitiveContactsKernelsCL= \
" float4 out = qtMul(qtMul(q,vcpy),qInv);\n"
" return out;\n"
"}\n"
"\n"
"__inline\n"
"Quaternion qtInvert(Quaternion q)\n"
"{\n"
" return (Quaternion)(-q.xyz, q.w);\n"
"}\n"
"\n"
"__inline\n"
"float4 qtInvRotate(const Quaternion q, float4 vec)\n"
"{\n"
" return qtRotate( qtInvert( q ), vec );\n"
"}\n"
"\n"
"__inline\n"
"float4 transform(const float4* p, const float4* translation, const Quaternion* orientation)\n"
"{\n"
" return qtRotate( *orientation, *p ) + (*translation);\n"
"}\n"
"\n"
"void trInverse(float4 translationIn, Quaternion orientationIn,\n"
" float4* translationOut, Quaternion* orientationOut)\n"
"{\n"
" *orientationOut = qtInvert(orientationIn);\n"
" *translationOut = qtRotate(*orientationOut, -translationIn);\n"
"}\n"
"\n"
"void trMul(float4 translationA, Quaternion orientationA,\n"
" float4 translationB, Quaternion orientationB,\n"
" float4* translationOut, Quaternion* orientationOut)\n"
@@ -290,17 +243,12 @@ static const char* primitiveContactsKernelsCL= \
" *orientationOut = qtMul(orientationA,orientationB);\n"
" *translationOut = transform(&translationB,&translationA,&orientationA);\n"
"}\n"
"\n"
"\n"
"\n"
"__inline\n"
"float4 normalize3(const float4 a)\n"
"{\n"
" float4 n = make_float4(a.x, a.y, a.z, 0.f);\n"
" return fastNormalize4( n );\n"
"}\n"
"\n"
"\n"
"__inline float4 lerp3(const float4 a,const float4 b, float t)\n"
"{\n"
" return make_float4( a.x + (b.x - a.x) * t,\n"
@@ -308,8 +256,6 @@ static const char* primitiveContactsKernelsCL= \
" a.z + (b.z - a.z) * t,\n"
" 0.f);\n"
"}\n"
"\n"
"\n"
"float signedDistanceFromPointToPlane(float4 point, float4 planeEqn, float4* closestPointOnFace)\n"
"{\n"
" float4 n = (float4)(planeEqn.x, planeEqn.y, planeEqn.z, 0);\n"
@@ -317,9 +263,6 @@ static const char* primitiveContactsKernelsCL= \
" *closestPointOnFace = point - dist * n;\n"
" return dist;\n"
"}\n"
"\n"
"\n"
"\n"
"inline bool IsPointInPolygon(float4 p, \n"
" const btGpuFace* face,\n"
" __global const float4* baseVertex,\n"
@@ -331,17 +274,14 @@ static const char* primitiveContactsKernelsCL= \
" float4 ab;\n"
" float4 ap;\n"
" float4 v;\n"
"\n"
" float4 plane = make_float4(face->m_plane.x,face->m_plane.y,face->m_plane.z,0.f);\n"
" \n"
" if (face->m_numIndices<2)\n"
" return false;\n"
"\n"
" \n"
" float4 v0 = baseVertex[convexIndices[face->m_indexOffset + face->m_numIndices-1]];\n"
" \n"
" b = v0;\n"
"\n"
" for(unsigned i=0; i != face->m_numIndices; ++i)\n"
" {\n"
" a = b;\n"
@@ -350,7 +290,6 @@ static const char* primitiveContactsKernelsCL= \
" ab = b-a;\n"
" ap = p-a;\n"
" v = cross3(ab,plane);\n"
"\n"
" if (dot(ap, v) > 0.f)\n"
" {\n"
" float ab_m2 = dot(ab, ab);\n"
@@ -375,10 +314,6 @@ static const char* primitiveContactsKernelsCL= \
" }\n"
" return true;\n"
"}\n"
"\n"
"\n"
"\n"
"\n"
"void computeContactSphereConvex(int pairIndex,\n"
" int bodyIndexA, int bodyIndexB, \n"
" int collidableIndexA, int collidableIndexB, \n"
@@ -388,7 +323,7 @@ static const char* primitiveContactsKernelsCL= \
" __global const float4* convexVertices,\n"
" __global const int* convexIndices,\n"
" __global const btGpuFace* faces,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" __global struct b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int maxContactCapacity,\n"
" float4 spherePos2,\n"
@@ -397,25 +332,19 @@ static const char* primitiveContactsKernelsCL= \
" float4 quat\n"
" )\n"
"{\n"
"\n"
" float4 invPos;\n"
" float4 invOrn;\n"
"\n"
" trInverse(pos,quat, &invPos,&invOrn);\n"
"\n"
" float4 spherePos = transform(&spherePos2,&invPos,&invOrn);\n"
"\n"
" int shapeIndex = collidables[collidableIndexB].m_shapeIndex;\n"
" int numFaces = convexShapes[shapeIndex].m_numFaces;\n"
" float4 closestPnt = (float4)(0, 0, 0, 0);\n"
" float4 hitNormalWorld = (float4)(0, 0, 0, 0);\n"
" float minDist = -1000000.f;\n"
" bool bCollide = true;\n"
"\n"
" for ( int f = 0; f < numFaces; f++ )\n"
" {\n"
" btGpuFace face = faces[convexShapes[shapeIndex].m_faceOffset+f];\n"
"\n"
" // set up a plane equation \n"
" float4 planeEqn;\n"
" float4 n1 = face.m_plane;\n"
@@ -427,21 +356,17 @@ static const char* primitiveContactsKernelsCL= \
" // compute a signed distance from the vertex in cloth to the face of rigidbody.\n"
" float4 pntReturn;\n"
" float dist = signedDistanceFromPointToPlane(spherePos, planeEqn, &pntReturn);\n"
"\n"
" // If the distance is positive, the plane is a separating plane. \n"
" if ( dist > radius )\n"
" {\n"
" bCollide = false;\n"
" break;\n"
" }\n"
"\n"
"\n"
" if (dist>0)\n"
" {\n"
" //might hit an edge or vertex\n"
" float4 out;\n"
" float4 zeroPos = make_float4(0,0,0,0);\n"
"\n"
" bool isInPoly = IsPointInPolygon(spherePos,\n"
" &face,\n"
" &convexVertices[convexShapes[shapeIndex].m_vertexOffset],\n"
@@ -488,9 +413,7 @@ static const char* primitiveContactsKernelsCL= \
" }\n"
" \n"
" }\n"
"\n"
" \n"
"\n"
" if (bCollide && minDist > -10000)\n"
" {\n"
" float4 normalOnSurfaceB1 = qtRotate(quat,-hitNormalWorld);\n"
@@ -500,35 +423,28 @@ static const char* primitiveContactsKernelsCL= \
" if (actualDepth<=0.f)\n"
" {\n"
" \n"
"\n"
" pOnB1.w = actualDepth;\n"
"\n"
" int dstIdx;\n"
" AppendInc( nGlobalContactsOut, dstIdx );\n"
" \n"
" \n"
" if (1)//dstIdx < maxContactCapacity)\n"
" {\n"
" __global b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" __global struct b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" c->m_worldNormal = normalOnSurfaceB1;\n"
" c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
" c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);\n"
" c->m_batchIdx = pairIndex;\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
" c->m_worldPos[0] = pOnB1;\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = -1;\n"
"\n"
" GET_NPOINTS(*c) = 1;\n"
" } \n"
"\n"
" }\n"
" }//if (hasCollision)\n"
"\n"
"}\n"
" \n"
"\n"
"\n"
"int extractManifoldSequential(const float4* p, int nPoints, float4 nearNormal, int4* contactIdx)\n"
"{\n"
" if( nPoints == 0 )\n"
@@ -621,9 +537,7 @@ static const char* primitiveContactsKernelsCL= \
" return 4;\n"
" \n"
"}\n"
"\n"
"#define MAX_PLANE_CONVEX_POINTS 64\n"
"\n"
"int computeContactPlaneConvex(int pairIndex,\n"
" int bodyIndexA, int bodyIndexB, \n"
" int collidableIndexA, int collidableIndexB, \n"
@@ -633,7 +547,7 @@ static const char* primitiveContactsKernelsCL= \
" __global const float4* convexVertices,\n"
" __global const int* convexIndices,\n"
" __global const btGpuFace* faces,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" __global struct b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int maxContactCapacity,\n"
" float4 posB,\n"
@@ -641,7 +555,6 @@ static const char* primitiveContactsKernelsCL= \
" )\n"
"{\n"
" int resultIndex=-1;\n"
"\n"
" int shapeIndex = collidables[collidableIndexB].m_shapeIndex;\n"
" __global const ConvexPolyhedronCL* hullB = &convexShapes[shapeIndex];\n"
" \n"
@@ -649,10 +562,8 @@ static const char* primitiveContactsKernelsCL= \
" posA = rigidBodies[bodyIndexA].m_pos;\n"
" Quaternion ornA;\n"
" ornA = rigidBodies[bodyIndexA].m_quat;\n"
"\n"
" int numContactsOut = 0;\n"
" int numWorldVertsB1= 0;\n"
"\n"
" float4 planeEq;\n"
" planeEq = faces[collidables[collidableIndexA].m_shapeIndex].m_plane;\n"
" float4 planeNormal = make_float4(planeEq.x,planeEq.y,planeEq.z,0.f);\n"
@@ -674,18 +585,13 @@ static const char* primitiveContactsKernelsCL= \
" trInverse(posB,ornB,&invPosB,&invOrnB);\n"
" trMul(invPosB,invOrnB,posA,ornA,&planeInConvexPos1,&planeInConvexOrn1); \n"
" }\n"
"\n"
" \n"
" float4 planeNormalInConvex = qtRotate(planeInConvexOrn1,-planeNormal);\n"
" float maxDot = -1e30;\n"
" int hitVertex=-1;\n"
" float4 hitVtx;\n"
"\n"
"\n"
"\n"
" float4 contactPoints[MAX_PLANE_CONVEX_POINTS];\n"
" int numPoints = 0;\n"
"\n"
" int4 contactIdx;\n"
" contactIdx=make_int4(0,1,2,3);\n"
" \n"
@@ -694,8 +600,6 @@ static const char* primitiveContactsKernelsCL= \
" {\n"
" float4 vtx = convexVertices[hullB->m_vertexOffset+i];\n"
" float curDot = dot(vtx,planeNormalInConvex);\n"
"\n"
"\n"
" if (curDot>maxDot)\n"
" {\n"
" hitVertex=i;\n"
@@ -705,7 +609,6 @@ static const char* primitiveContactsKernelsCL= \
" if (numPoints==MAX_PLANE_CONVEX_POINTS)\n"
" numPoints--;\n"
" }\n"
"\n"
" if (numPoints<MAX_PLANE_CONVEX_POINTS)\n"
" {\n"
" float4 vtxWorld = transform(&vtx, &posB, &ornB);\n"
@@ -718,34 +621,29 @@ static const char* primitiveContactsKernelsCL= \
" numPoints++;\n"
" }\n"
" }\n"
"\n"
" }\n"
"\n"
" int numReducedPoints = numPoints;\n"
" if (numPoints>4)\n"
" {\n"
" numReducedPoints = extractManifoldSequential( contactPoints, numPoints, planeNormalInConvex, &contactIdx);\n"
" }\n"
"\n"
" if (numReducedPoints>0)\n"
" {\n"
" int dstIdx;\n"
" AppendInc( nGlobalContactsOut, dstIdx );\n"
"\n"
" if (dstIdx < maxContactCapacity)\n"
" {\n"
" resultIndex = dstIdx;\n"
" __global b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" __global struct b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" c->m_worldNormal = planeNormalWorld;\n"
" //c->setFrictionCoeff(0.7);\n"
" //c->setRestituitionCoeff(0.f);\n"
" c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
" c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);\n"
" c->m_batchIdx = pairIndex;\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = -1;\n"
"\n"
" switch (numReducedPoints)\n"
" {\n"
" case 4:\n"
@@ -764,18 +662,15 @@ static const char* primitiveContactsKernelsCL= \
" GET_NPOINTS(*c) = numReducedPoints;\n"
" }//if (dstIdx < numPairs)\n"
" } \n"
"\n"
" return resultIndex;\n"
"}\n"
"\n"
"\n"
"void computeContactPlaneSphere(int pairIndex,\n"
" int bodyIndexA, int bodyIndexB, \n"
" int collidableIndexA, int collidableIndexB, \n"
" __global const BodyData* rigidBodies, \n"
" __global const btCollidableGpu* collidables,\n"
" __global const btGpuFace* faces,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" __global struct b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int maxContactCapacity)\n"
"{\n"
@@ -812,15 +707,14 @@ static const char* primitiveContactsKernelsCL= \
" float4 normalOnSurfaceB1 = qtRotate(ornA1,planeNormal1);\n"
" float4 pOnB1 = vtxInPlaneWorld1+normalOnSurfaceB1*distance;\n"
" pOnB1.w = distance;\n"
"\n"
" int dstIdx;\n"
" AppendInc( nGlobalContactsOut, dstIdx );\n"
" \n"
" if (dstIdx < maxContactCapacity)\n"
" {\n"
" __global b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" __global struct b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" c->m_worldNormal = normalOnSurfaceB1;\n"
" c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
" c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);\n"
" c->m_batchIdx = pairIndex;\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
@@ -831,8 +725,6 @@ static const char* primitiveContactsKernelsCL= \
" }//if (dstIdx < numPairs)\n"
" }//if (hasCollision)\n"
"}\n"
"\n"
"\n"
"__kernel void primitiveContactsKernel( __global int4* pairs, \n"
" __global const BodyData* rigidBodies, \n"
" __global const btCollidableGpu* collidables,\n"
@@ -841,27 +733,23 @@ static const char* primitiveContactsKernelsCL= \
" __global const float4* uniqueEdges,\n"
" __global const btGpuFace* faces,\n"
" __global const int* indices,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" __global struct b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int numPairs, int maxContactCapacity)\n"
"{\n"
"\n"
" int i = get_global_id(0);\n"
" int pairIndex = i;\n"
" \n"
" float4 worldVertsB1[64];\n"
" float4 worldVertsB2[64];\n"
" int capacityWorldVerts = 64; \n"
"\n"
" float4 localContactsOut[64];\n"
" int localContactCapacity=64;\n"
" \n"
" float minDist = -1e30f;\n"
" float maxDist = 0.02f;\n"
"\n"
" if (i<numPairs)\n"
" {\n"
"\n"
" int bodyIndexA = pairs[i].x;\n"
" int bodyIndexB = pairs[i].y;\n"
" \n"
@@ -871,7 +759,6 @@ static const char* primitiveContactsKernelsCL= \
" if (collidables[collidableIndexA].m_shapeType == SHAPE_PLANE &&\n"
" collidables[collidableIndexB].m_shapeType == SHAPE_CONVEX_HULL)\n"
" {\n"
"\n"
" float4 posB;\n"
" posB = rigidBodies[bodyIndexB].m_pos;\n"
" Quaternion ornB;\n"
@@ -881,31 +768,22 @@ static const char* primitiveContactsKernelsCL= \
" faces, globalContactsOut, nGlobalContactsOut,maxContactCapacity, posB,ornB);\n"
" if (contactIndex>=0)\n"
" pairs[pairIndex].z = contactIndex;\n"
"\n"
" return;\n"
" }\n"
"\n"
"\n"
" if (collidables[collidableIndexA].m_shapeType == SHAPE_CONVEX_HULL &&\n"
" collidables[collidableIndexB].m_shapeType == SHAPE_PLANE)\n"
" {\n"
"\n"
" float4 posA;\n"
" posA = rigidBodies[bodyIndexA].m_pos;\n"
" Quaternion ornA;\n"
" ornA = rigidBodies[bodyIndexA].m_quat;\n"
"\n"
"\n"
" int contactIndex = computeContactPlaneConvex( pairIndex, bodyIndexB,bodyIndexA, collidableIndexB,collidableIndexA, \n"
" rigidBodies,collidables,convexShapes,vertices,indices,\n"
" faces, globalContactsOut, nGlobalContactsOut,maxContactCapacity,posA,ornA);\n"
"\n"
" if (contactIndex>=0)\n"
" pairs[pairIndex].z = contactIndex;\n"
"\n"
" return;\n"
" }\n"
"\n"
" if (collidables[collidableIndexA].m_shapeType == SHAPE_PLANE &&\n"
" collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)\n"
" {\n"
@@ -913,22 +791,15 @@ static const char* primitiveContactsKernelsCL= \
" rigidBodies,collidables,faces, globalContactsOut, nGlobalContactsOut,maxContactCapacity);\n"
" return;\n"
" }\n"
"\n"
"\n"
" if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&\n"
" collidables[collidableIndexB].m_shapeType == SHAPE_PLANE)\n"
" {\n"
"\n"
"\n"
" computeContactPlaneSphere( pairIndex, bodyIndexB,bodyIndexA, collidableIndexB,collidableIndexA, \n"
" rigidBodies,collidables,\n"
" faces, globalContactsOut, nGlobalContactsOut,maxContactCapacity);\n"
"\n"
" return;\n"
" }\n"
"\n"
" \n"
"\n"
" \n"
" if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&\n"
" collidables[collidableIndexB].m_shapeType == SHAPE_CONVEX_HULL)\n"
@@ -938,14 +809,11 @@ static const char* primitiveContactsKernelsCL= \
" float sphereRadius = collidables[collidableIndexA].m_radius;\n"
" float4 convexPos = rigidBodies[bodyIndexB].m_pos;\n"
" float4 convexOrn = rigidBodies[bodyIndexB].m_quat;\n"
"\n"
" computeContactSphereConvex(pairIndex, bodyIndexA, bodyIndexB, collidableIndexA, collidableIndexB, \n"
" rigidBodies,collidables,convexShapes,vertices,indices,faces, globalContactsOut, nGlobalContactsOut,maxContactCapacity,\n"
" spherePos,sphereRadius,convexPos,convexOrn);\n"
"\n"
" return;\n"
" }\n"
"\n"
" if (collidables[collidableIndexA].m_shapeType == SHAPE_CONVEX_HULL &&\n"
" collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)\n"
" {\n"
@@ -954,7 +822,6 @@ static const char* primitiveContactsKernelsCL= \
" float sphereRadius = collidables[collidableIndexB].m_radius;\n"
" float4 convexPos = rigidBodies[bodyIndexA].m_pos;\n"
" float4 convexOrn = rigidBodies[bodyIndexA].m_quat;\n"
"\n"
" computeContactSphereConvex(pairIndex, bodyIndexB, bodyIndexA, collidableIndexB, collidableIndexA, \n"
" rigidBodies,collidables,convexShapes,vertices,indices,faces, globalContactsOut, nGlobalContactsOut,maxContactCapacity,\n"
" spherePos,sphereRadius,convexPos,convexOrn);\n"
@@ -974,7 +841,6 @@ static const char* primitiveContactsKernelsCL= \
" float radiusB = collidables[collidableIndexB].m_radius;\n"
" float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
" float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
"\n"
" float4 diff = posA-posB;\n"
" float len = length(diff);\n"
" \n"
@@ -996,9 +862,9 @@ static const char* primitiveContactsKernelsCL= \
" \n"
" if (dstIdx < maxContactCapacity)\n"
" {\n"
" __global b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" __global struct b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" c->m_worldNormal = -normalOnSurfaceB;\n"
" c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
" c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);\n"
" c->m_batchIdx = pairIndex;\n"
" int bodyA = pairs[pairIndex].x;\n"
" int bodyB = pairs[pairIndex].y;\n"
@@ -1010,15 +876,10 @@ static const char* primitiveContactsKernelsCL= \
" GET_NPOINTS(*c) = 1;\n"
" }//if (dstIdx < numPairs)\n"
" }//if ( len <= (radiusA+radiusB))\n"
"\n"
" return;\n"
" }//SHAPE_SPHERE SHAPE_SPHERE\n"
"\n"
" }// if (i<numPairs)\n"
"\n"
"}\n"
"\n"
"\n"
"// work-in-progress\n"
"__kernel void processCompoundPairsPrimitivesKernel( __global const int4* gpuCompoundPairs,\n"
" __global const BodyData* rigidBodies, \n"
@@ -1030,18 +891,16 @@ static const char* primitiveContactsKernelsCL= \
" __global const int* indices,\n"
" __global btAabbCL* aabbs,\n"
" __global const btGpuChildShape* gpuChildShapes,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" __global struct b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int numCompoundPairs, int maxContactCapacity\n"
" )\n"
"{\n"
"\n"
" int i = get_global_id(0);\n"
" if (i<numCompoundPairs)\n"
" {\n"
" int bodyIndexA = gpuCompoundPairs[i].x;\n"
" int bodyIndexB = gpuCompoundPairs[i].y;\n"
"\n"
" int childShapeIndexA = gpuCompoundPairs[i].z;\n"
" int childShapeIndexB = gpuCompoundPairs[i].w;\n"
" \n"
@@ -1087,26 +946,21 @@ static const char* primitiveContactsKernelsCL= \
" \n"
" int shapeTypeA = collidables[collidableIndexA].m_shapeType;\n"
" int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
"\n"
" int pairIndex = i;\n"
" if ((shapeTypeA == SHAPE_PLANE) && (shapeTypeB==SHAPE_CONVEX_HULL))\n"
" {\n"
"\n"
" computeContactPlaneConvex( pairIndex, bodyIndexA,bodyIndexB, collidableIndexA,collidableIndexB, \n"
" rigidBodies,collidables,convexShapes,vertices,indices,\n"
" faces, globalContactsOut, nGlobalContactsOut,maxContactCapacity,posB,ornB);\n"
" return;\n"
" }\n"
"\n"
" if ((shapeTypeA == SHAPE_CONVEX_HULL) && (shapeTypeB==SHAPE_PLANE))\n"
" {\n"
"\n"
" computeContactPlaneConvex( pairIndex, bodyIndexB,bodyIndexA, collidableIndexB,collidableIndexA, \n"
" rigidBodies,collidables,convexShapes,vertices,indices,\n"
" faces, globalContactsOut, nGlobalContactsOut,maxContactCapacity,posA,ornA);\n"
" return;\n"
" }\n"
"\n"
" if ((shapeTypeA == SHAPE_CONVEX_HULL) && (shapeTypeB == SHAPE_SPHERE))\n"
" {\n"
" float4 spherePos = rigidBodies[bodyIndexB].m_pos;\n"
@@ -1120,15 +974,12 @@ static const char* primitiveContactsKernelsCL= \
" \n"
" return;\n"
" }\n"
"\n"
" if ((shapeTypeA == SHAPE_SPHERE) && (shapeTypeB == SHAPE_CONVEX_HULL))\n"
" {\n"
"\n"
" float4 spherePos = rigidBodies[bodyIndexA].m_pos;\n"
" float sphereRadius = collidables[collidableIndexA].m_radius;\n"
" float4 convexPos = posB;\n"
" float4 convexOrn = ornB;\n"
"\n"
" \n"
" computeContactSphereConvex(pairIndex, bodyIndexA, bodyIndexB, collidableIndexA, collidableIndexB, \n"
" rigidBodies,collidables,convexShapes,vertices,indices,faces, globalContactsOut, nGlobalContactsOut,maxContactCapacity,\n"
@@ -1138,28 +989,21 @@ static const char* primitiveContactsKernelsCL= \
" }\n"
" }// if (i<numCompoundPairs)\n"
"}\n"
"\n"
"\n"
"bool pointInTriangle(const float4* vertices, const float4* normal, float4 *p )\n"
"{\n"
"\n"
" const float4* p1 = &vertices[0];\n"
" const float4* p2 = &vertices[1];\n"
" const float4* p3 = &vertices[2];\n"
"\n"
" float4 edge1; edge1 = (*p2 - *p1);\n"
" float4 edge2; edge2 = ( *p3 - *p2 );\n"
" float4 edge3; edge3 = ( *p1 - *p3 );\n"
"\n"
" \n"
" float4 p1_to_p; p1_to_p = ( *p - *p1 );\n"
" float4 p2_to_p; p2_to_p = ( *p - *p2 );\n"
" float4 p3_to_p; p3_to_p = ( *p - *p3 );\n"
"\n"
" float4 edge1_normal; edge1_normal = ( cross(edge1,*normal));\n"
" float4 edge2_normal; edge2_normal = ( cross(edge2,*normal));\n"
" float4 edge3_normal; edge3_normal = ( cross(edge3,*normal));\n"
"\n"
" \n"
" \n"
" float r1, r2, r3;\n"
@@ -1172,10 +1016,7 @@ static const char* primitiveContactsKernelsCL= \
" if ( r1 <= 0 && r2 <= 0 && r3 <= 0 ) \n"
" return true;\n"
" return false;\n"
"\n"
"}\n"
"\n"
"\n"
"float segmentSqrDistance(float4 from, float4 to,float4 p, float4* nearest) \n"
"{\n"
" float4 diff = p - from;\n"
@@ -1201,15 +1042,13 @@ static const char* primitiveContactsKernelsCL= \
" *nearest = from + t*v;\n"
" return dot(diff,diff); \n"
"}\n"
"\n"
"\n"
"void computeContactSphereTriangle(int pairIndex,\n"
" int bodyIndexA, int bodyIndexB,\n"
" int collidableIndexA, int collidableIndexB, \n"
" __global const BodyData* rigidBodies, \n"
" __global const btCollidableGpu* collidables,\n"
" const float4* triangleVertices,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" __global struct b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int maxContactCapacity,\n"
" float4 spherePos2,\n"
@@ -1219,10 +1058,8 @@ static const char* primitiveContactsKernelsCL= \
" int faceIndex\n"
" )\n"
"{\n"
"\n"
" float4 invPos;\n"
" float4 invOrn;\n"
"\n"
" trInverse(pos,quat, &invPos,&invOrn);\n"
" float4 spherePos = transform(&spherePos2,&invPos,&invOrn);\n"
" int numFaces = 3;\n"
@@ -1230,13 +1067,10 @@ static const char* primitiveContactsKernelsCL= \
" float4 hitNormalWorld = (float4)(0, 0, 0, 0);\n"
" float minDist = -1000000.f;\n"
" bool bCollide = false;\n"
"\n"
" \n"
" //////////////////////////////////////\n"
"\n"
" float4 sphereCenter;\n"
" sphereCenter = spherePos;\n"
"\n"
" const float4* vertices = triangleVertices;\n"
" float contactBreakingThreshold = 0.f;//todo?\n"
" float radiusWithThreshold = radius + contactBreakingThreshold;\n"
@@ -1252,7 +1086,6 @@ static const char* primitiveContactsKernelsCL= \
" p1ToCenter = sphereCenter - vertices[0];\n"
" \n"
" float distanceFromPlane = dot(p1ToCenter,normal);\n"
"\n"
" if (distanceFromPlane < 0.f)\n"
" {\n"
" //triangle facing the other way\n"
@@ -1260,7 +1093,6 @@ static const char* primitiveContactsKernelsCL= \
" normal *= -1.f;\n"
" }\n"
" hitNormalWorld = normal;\n"
"\n"
" bool isInsideContactPlane = distanceFromPlane < radiusWithThreshold;\n"
" \n"
" // Check for contact / intersection\n"
@@ -1284,7 +1116,6 @@ static const char* primitiveContactsKernelsCL= \
" {\n"
" float4 pa =vertices[i];\n"
" float4 pb = vertices[(i+1)%3];\n"
"\n"
" float distanceSqr = segmentSqrDistance(pa,pb,sphereCenter, &nearestOnEdge);\n"
" if (distanceSqr < contactCapsuleRadiusSqr) \n"
" {\n"
@@ -1297,10 +1128,8 @@ static const char* primitiveContactsKernelsCL= \
" }\n"
" }\n"
" }\n"
"\n"
" if (hasContact) \n"
" {\n"
"\n"
" closestPnt = contactPoint;\n"
" float4 contactToCenter = sphereCenter - contactPoint;\n"
" minDist = length(contactToCenter);\n"
@@ -1311,23 +1140,18 @@ static const char* primitiveContactsKernelsCL= \
" }\n"
" \n"
" }\n"
"\n"
"\n"
" /////////////////////////////////////\n"
"\n"
" if (bCollide && minDist > -10000)\n"
" {\n"
" \n"
" float4 normalOnSurfaceB1 = qtRotate(quat,-hitNormalWorld);\n"
" float4 pOnB1 = transform(&closestPnt,&pos,&quat);\n"
" float actualDepth = minDist-radius;\n"
"\n"
" \n"
" if (actualDepth<=0.f)\n"
" {\n"
" pOnB1.w = actualDepth;\n"
" int dstIdx;\n"
"\n"
" \n"
" float lenSqr = dot3F4(normalOnSurfaceB1,normalOnSurfaceB1);\n"
" if (lenSqr>FLT_EPSILON)\n"
@@ -1336,28 +1160,21 @@ static const char* primitiveContactsKernelsCL= \
" \n"
" if (dstIdx < maxContactCapacity)\n"
" {\n"
" __global b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" __global struct b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" c->m_worldNormal = normalOnSurfaceB1;\n"
" c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
" c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);\n"
" c->m_batchIdx = pairIndex;\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
" c->m_worldPos[0] = pOnB1;\n"
"\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = faceIndex;\n"
"\n"
" GET_NPOINTS(*c) = 1;\n"
" } \n"
" }\n"
"\n"
" }\n"
" }//if (hasCollision)\n"
"\n"
"}\n"
"\n"
"\n"
"\n"
"// work-in-progress\n"
"__kernel void findConcaveSphereContactsKernel( __global int4* concavePairs,\n"
" __global const BodyData* rigidBodies,\n"
@@ -1368,26 +1185,21 @@ static const char* primitiveContactsKernelsCL= \
" __global const btGpuFace* faces,\n"
" __global const int* indices,\n"
" __global btAabbCL* aabbs,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" __global struct b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int numConcavePairs, int maxContactCapacity\n"
" )\n"
"{\n"
"\n"
" int i = get_global_id(0);\n"
" if (i>=numConcavePairs)\n"
" return;\n"
" int pairIdx = i;\n"
"\n"
" int bodyIndexA = concavePairs[i].x;\n"
" int bodyIndexB = concavePairs[i].y;\n"
"\n"
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
"\n"
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
"\n"
" if (collidables[collidableIndexB].m_shapeType==SHAPE_SPHERE)\n"
" {\n"
" int f = concavePairs[i].z;\n"
@@ -1400,18 +1212,15 @@ static const char* primitiveContactsKernelsCL= \
" float4 vert = vertices[convexShapes[shapeIndexA].m_vertexOffset+index];\n"
" verticesA[i] = vert;\n"
" }\n"
"\n"
" float4 spherePos = rigidBodies[bodyIndexB].m_pos;\n"
" float sphereRadius = collidables[collidableIndexB].m_radius;\n"
" float4 convexPos = rigidBodies[bodyIndexA].m_pos;\n"
" float4 convexOrn = rigidBodies[bodyIndexA].m_quat;\n"
"\n"
" computeContactSphereTriangle(i, bodyIndexB, bodyIndexA, collidableIndexB, collidableIndexA, \n"
" rigidBodies,collidables,\n"
" verticesA,\n"
" globalContactsOut, nGlobalContactsOut,maxContactCapacity,\n"
" spherePos,sphereRadius,convexPos,convexOrn, f);\n"
"\n"
" return;\n"
" }\n"
"}\n"

55
src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl
View File

@@ -41,22 +41,7 @@ typedef unsigned int u32;
typedef struct
{
float4 m_worldPos[4];
float4 m_worldNormal; // w: m_nPoints
u32 m_coeffs;
u32 m_batchIdx;
int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr
int m_bodyBPtrAndSignBit;
int m_childIndexA;
int m_childIndexB;
float m_unused1;
int m_unused2;
} Contact4;
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Contact4Data.h"
///keep this in sync with btCollidable.h
@@ -891,7 +876,7 @@ __kernel void extractManifoldAndAddContactKernel(__global const int4* pairs,
__global const float4* separatingNormalsWorld,
__global const int* contactCounts,
__global const int* contactOffsets,
__global Contact4* restrict contactsOut,
__global struct b3Contact4Data* restrict contactsOut,
counter32_t nContactsOut,
int numPairs,
int pairIndex
@@ -922,9 +907,9 @@ __kernel void extractManifoldAndAddContactKernel(__global const int4* pairs,
AppendInc( nContactsOut, dstIdx );
//if ((dstIdx+nContacts) < capacity)
{
__global Contact4* c = contactsOut + dstIdx;
__global struct b3Contact4Data* c = contactsOut + dstIdx;
c->m_worldNormal = normal;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = idx;
int bodyA = pairs[pairIndex].x;
int bodyB = pairs[pairIndex].y;
@@ -970,7 +955,7 @@ __kernel void clipHullHullKernel( __global int4* pairs,
__global const int* indices,
__global const float4* separatingNormals,
__global const int* hasSeparatingAxis,
__global Contact4* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int numPairs,
int contactCapacity)
@@ -1037,9 +1022,9 @@ __kernel void clipHullHullKernel( __global int4* pairs,
{
pairs[pairIndex].z = dstIdx;
__global Contact4* c = globalContactsOut+ dstIdx;
__global struct b3Contact4Data* c = globalContactsOut+ dstIdx;
c->m_worldNormal = normal;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = pairIndex;
int bodyA = pairs[pairIndex].x;
int bodyB = pairs[pairIndex].y;
@@ -1073,7 +1058,7 @@ __kernel void clipCompoundsHullHullKernel( __global const int4* gpuCompoundPai
__global const btGpuChildShape* gpuChildShapes,
__global const float4* gpuCompoundSepNormalsOut,
__global const int* gpuHasCompoundSepNormalsOut,
__global Contact4* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int numCompoundPairs, int maxContactCapacity)
{
@@ -1170,9 +1155,9 @@ __kernel void clipCompoundsHullHullKernel( __global const int4* gpuCompoundPai
AppendInc( nGlobalContactsOut, dstIdx );
if ((dstIdx+nReducedContacts) < maxContactCapacity)
{
__global Contact4* c = globalContactsOut+ dstIdx;
__global struct b3Contact4Data* c = globalContactsOut+ dstIdx;
c->m_worldNormal = normal;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = pairIndex;
int bodyA = gpuCompoundPairs[pairIndex].x;
int bodyB = gpuCompoundPairs[pairIndex].y;
@@ -1200,7 +1185,7 @@ __kernel void sphereSphereCollisionKernel( __global const int4* pairs,
__global const btCollidableGpu* collidables,
__global const float4* separatingNormals,
__global const int* hasSeparatingAxis,
__global Contact4* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int numPairs)
{
@@ -1246,9 +1231,9 @@ __kernel void sphereSphereCollisionKernel( __global const int4* pairs,
if (dstIdx < numPairs)
{
__global Contact4* c = &globalContactsOut[dstIdx];
__global struct b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = normalOnSurfaceB;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = pairIndex;
int bodyA = pairs[pairIndex].x;
int bodyB = pairs[pairIndex].y;
@@ -1275,7 +1260,7 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
__global const int* indices,
__global const btGpuChildShape* gpuChildShapes,
__global const float4* separatingNormals,
__global Contact4* restrict globalContactsOut,
__global struct b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int numConcavePairs)
{
@@ -1479,9 +1464,9 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
AppendInc( nGlobalContactsOut, dstIdx );
//if ((dstIdx+nReducedContacts) < capacity)
{
__global Contact4* c = globalContactsOut+ dstIdx;
__global struct b3Contact4Data* c = globalContactsOut+ dstIdx;
c->m_worldNormal = normal;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = pairIndex;
int bodyA = concavePairsIn[pairIndex].x;
int bodyB = concavePairsIn[pairIndex].y;
@@ -1747,7 +1732,7 @@ __kernel void clipFacesAndContactReductionKernel( __global int4* pairs,
__global const BodyData* rigidBodies,
__global const float4* separatingNormals,
__global const int* hasSeparatingAxis,
__global Contact4* globalContactsOut,
__global struct b3Contact4Data* globalContactsOut,
__global int4* clippingFacesOut,
__global float4* worldVertsA1,
__global float4* worldNormalsA1,
@@ -1860,7 +1845,7 @@ __kernel void newContactReductionKernel( __global int4* pairs,
__global const BodyData* rigidBodies,
__global const float4* separatingNormals,
__global const int* hasSeparatingAxis,
__global Contact4* globalContactsOut,
__global struct b3Contact4Data* globalContactsOut,
__global int4* clippingFaces,
__global float4* worldVertsB2,
volatile __global int* nGlobalContactsOut,
@@ -1901,9 +1886,9 @@ __kernel void newContactReductionKernel( __global int4* pairs,
if (dstIdx < numPairs)
{
__global Contact4* c = &globalContactsOut[dstIdx];
__global struct b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = normal;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
c->m_batchIdx = pairIndex;
int bodyA = pairs[pairIndex].x;
int bodyB = pairs[pairIndex].y;

303
src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h
View File

File diff suppressed because it is too large Load Diff

171
src/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h
View File

@@ -2,17 +2,11 @@
static const char* satKernelsCL= \
"//keep this enum in sync with the CPU version (in btCollidable.h)\n"
"//written by Erwin Coumans\n"
"\n"
"\n"
"#define SHAPE_CONVEX_HULL 3\n"
"#define SHAPE_CONCAVE_TRIMESH 5\n"
"#define TRIANGLE_NUM_CONVEX_FACES 5\n"
"#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6\n"
"\n"
"\n"
"\n"
"typedef unsigned int u32;\n"
"\n"
"///keep this in sync with btCollidable.h\n"
"typedef struct\n"
"{\n"
@@ -22,7 +16,6 @@ static const char* satKernelsCL= \
" int m_shapeIndex;\n"
" \n"
"} btCollidableGpu;\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_childPosition;\n"
@@ -32,22 +25,17 @@ static const char* satKernelsCL= \
" int m_unused1;\n"
" int m_unused2;\n"
"} btGpuChildShape;\n"
"\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_pos;\n"
" float4 m_quat;\n"
" float4 m_linVel;\n"
" float4 m_angVel;\n"
"\n"
" u32 m_collidableIdx;\n"
" float m_invMass;\n"
" float m_restituitionCoeff;\n"
" float m_frictionCoeff;\n"
"} BodyData;\n"
"\n"
"\n"
"typedef struct \n"
"{\n"
" float4 m_localCenter;\n"
@@ -59,13 +47,11 @@ static const char* satKernelsCL= \
" int m_faceOffset;\n"
" int m_numFaces;\n"
" int m_numVertices;\n"
"\n"
" int m_vertexOffset;\n"
" int m_uniqueEdgesOffset;\n"
" int m_numUniqueEdges;\n"
" int m_unused;\n"
"} ConvexPolyhedronCL;\n"
"\n"
"typedef struct \n"
"{\n"
" union\n"
@@ -81,35 +67,27 @@ static const char* satKernelsCL= \
" int m_maxIndices[4];\n"
" };\n"
"} btAabbCL;\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_plane;\n"
" int m_indexOffset;\n"
" int m_numIndices;\n"
"} btGpuFace;\n"
"\n"
"#define make_float4 (float4)\n"
"\n"
"\n"
"__inline\n"
"float4 cross3(float4 a, float4 b)\n"
"{\n"
" return cross(a,b);\n"
"\n"
" \n"
"// float4 a1 = make_float4(a.xyz,0.f);\n"
"// float4 b1 = make_float4(b.xyz,0.f);\n"
"\n"
"// return cross(a1,b1);\n"
"\n"
"//float4 c = make_float4(a.y*b.z - a.z*b.y,a.z*b.x - a.x*b.z,a.x*b.y - a.y*b.x,0.f);\n"
" \n"
" // float4 c = make_float4(a.y*b.z - a.z*b.y,1.f,a.x*b.y - a.y*b.x,0.f);\n"
" \n"
" //return c;\n"
"}\n"
"\n"
"__inline\n"
"float dot3F4(float4 a, float4 b)\n"
"{\n"
@@ -117,36 +95,24 @@ static const char* satKernelsCL= \
" float4 b1 = make_float4(b.xyz,0.f);\n"
" return dot(a1, b1);\n"
"}\n"
"\n"
"__inline\n"
"float4 fastNormalize4(float4 v)\n"
"{\n"
" v = make_float4(v.xyz,0.f);\n"
" return fast_normalize(v);\n"
"}\n"
"\n"
"\n"
"///////////////////////////////////////\n"
"// Quaternion\n"
"///////////////////////////////////////\n"
"\n"
"typedef float4 Quaternion;\n"
"\n"
"__inline\n"
"Quaternion qtMul(Quaternion a, Quaternion b);\n"
"\n"
"__inline\n"
"Quaternion qtNormalize(Quaternion in);\n"
"\n"
"__inline\n"
"float4 qtRotate(Quaternion q, float4 vec);\n"
"\n"
"__inline\n"
"Quaternion qtInvert(Quaternion q);\n"
"\n"
"\n"
"\n"
"\n"
"__inline\n"
"Quaternion qtMul(Quaternion a, Quaternion b)\n"
"{\n"
@@ -157,7 +123,6 @@ static const char* satKernelsCL= \
" ans.w = a.w*b.w - dot3F4(a, b);\n"
" return ans;\n"
"}\n"
"\n"
"__inline\n"
"Quaternion qtNormalize(Quaternion in)\n"
"{\n"
@@ -174,41 +139,33 @@ static const char* satKernelsCL= \
" float4 out = qtMul(qtMul(q,vcpy),qInv);\n"
" return out;\n"
"}\n"
"\n"
"__inline\n"
"Quaternion qtInvert(Quaternion q)\n"
"{\n"
" return (Quaternion)(-q.xyz, q.w);\n"
"}\n"
"\n"
"__inline\n"
"float4 qtInvRotate(const Quaternion q, float4 vec)\n"
"{\n"
" return qtRotate( qtInvert( q ), vec );\n"
"}\n"
"\n"
"__inline\n"
"float4 transform(const float4* p, const float4* translation, const Quaternion* orientation)\n"
"{\n"
" return qtRotate( *orientation, *p ) + (*translation);\n"
"}\n"
"\n"
"\n"
"\n"
"__inline\n"
"float4 normalize3(const float4 a)\n"
"{\n"
" float4 n = make_float4(a.x, a.y, a.z, 0.f);\n"
" return fastNormalize4( n );\n"
"}\n"
"\n"
"inline void projectLocal(const ConvexPolyhedronCL* hull, const float4 pos, const float4 orn, \n"
"const float4* dir, const float4* vertices, float* min, float* max)\n"
"{\n"
" min[0] = FLT_MAX;\n"
" max[0] = -FLT_MAX;\n"
" int numVerts = hull->m_numVertices;\n"
"\n"
" const float4 localDir = qtInvRotate(orn,*dir);\n"
" float offset = dot(pos,*dir);\n"
" for(int i=0;i<numVerts;i++)\n"
@@ -228,14 +185,12 @@ static const char* satKernelsCL= \
" min[0] += offset;\n"
" max[0] += offset;\n"
"}\n"
"\n"
"inline void project(__global const ConvexPolyhedronCL* hull, const float4 pos, const float4 orn, \n"
"const float4* dir, __global const float4* vertices, float* min, float* max)\n"
"{\n"
" min[0] = FLT_MAX;\n"
" max[0] = -FLT_MAX;\n"
" int numVerts = hull->m_numVertices;\n"
"\n"
" const float4 localDir = qtInvRotate(orn,*dir);\n"
" float offset = dot(pos,*dir);\n"
" for(int i=0;i<numVerts;i++)\n"
@@ -255,7 +210,6 @@ static const char* satKernelsCL= \
" min[0] += offset;\n"
" max[0] += offset;\n"
"}\n"
"\n"
"inline bool TestSepAxisLocalA(const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
" const float4 posA,const float4 ornA,\n"
" const float4 posB,const float4 ornB,\n"
@@ -265,28 +219,19 @@ static const char* satKernelsCL= \
" float Min1,Max1;\n"
" projectLocal(hullA,posA,ornA,sep_axis,verticesA, &Min0, &Max0);\n"
" project(hullB,posB,ornB, sep_axis,verticesB, &Min1, &Max1);\n"
"\n"
" if(Max0<Min1 || Max1<Min0)\n"
" return false;\n"
"\n"
" float d0 = Max0 - Min1;\n"
" float d1 = Max1 - Min0;\n"
" *depth = d0<d1 ? d0:d1;\n"
" return true;\n"
"}\n"
"\n"
"\n"
"\n"
"\n"
"inline bool IsAlmostZero(const float4 v)\n"
"{\n"
" if(fabs(v.x)>1e-6f || fabs(v.y)>1e-6f || fabs(v.z)>1e-6f)\n"
" return false;\n"
" return true;\n"
"}\n"
"\n"
"\n"
"\n"
"bool findSeparatingAxisLocalA( const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
" const float4 posA1,\n"
" const float4 ornA,\n"
@@ -298,7 +243,6 @@ static const char* satKernelsCL= \
" const float4* uniqueEdgesA, \n"
" const btGpuFace* facesA,\n"
" const int* indicesA,\n"
"\n"
" __global const float4* verticesB, \n"
" __global const float4* uniqueEdgesB, \n"
" __global const btGpuFace* facesB,\n"
@@ -307,7 +251,6 @@ static const char* satKernelsCL= \
" float* dmin)\n"
"{\n"
" int i = get_global_id(0);\n"
"\n"
" float4 posA = posA1;\n"
" posA.w = 0.f;\n"
" float4 posB = posB1;\n"
@@ -339,7 +282,6 @@ static const char* satKernelsCL= \
" }\n"
" return true;\n"
"}\n"
"\n"
"bool findSeparatingAxisLocalB( __global const ConvexPolyhedronCL* hullA, const ConvexPolyhedronCL* hullB, \n"
" const float4 posA1,\n"
" const float4 ornA,\n"
@@ -358,7 +300,6 @@ static const char* satKernelsCL= \
" float* dmin)\n"
"{\n"
" int i = get_global_id(0);\n"
"\n"
" float4 posA = posA1;\n"
" posA.w = 0.f;\n"
" float4 posB = posB1;\n"
@@ -390,9 +331,6 @@ static const char* satKernelsCL= \
" }\n"
" return true;\n"
"}\n"
"\n"
"\n"
"\n"
"bool findSeparatingAxisEdgeEdgeLocalA( const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
" const float4 posA1,\n"
" const float4 ornA,\n"
@@ -411,36 +349,28 @@ static const char* satKernelsCL= \
" float* dmin)\n"
"{\n"
" int i = get_global_id(0);\n"
"\n"
" float4 posA = posA1;\n"
" posA.w = 0.f;\n"
" float4 posB = posB1;\n"
" posB.w = 0.f;\n"
"\n"
" int curPlaneTests=0;\n"
"\n"
" int curEdgeEdge = 0;\n"
" // Test edges\n"
" for(int e0=0;e0<hullA->m_numUniqueEdges;e0++)\n"
" {\n"
" const float4 edge0 = uniqueEdgesA[hullA->m_uniqueEdgesOffset+e0];\n"
" float4 edge0World = qtRotate(ornA,edge0);\n"
"\n"
" for(int e1=0;e1<hullB->m_numUniqueEdges;e1++)\n"
" {\n"
" const float4 edge1 = uniqueEdgesB[hullB->m_uniqueEdgesOffset+e1];\n"
" float4 edge1World = qtRotate(ornB,edge1);\n"
"\n"
"\n"
" float4 crossje = cross3(edge0World,edge1World);\n"
"\n"
" curEdgeEdge++;\n"
" if(!IsAlmostZero(crossje))\n"
" {\n"
" crossje = normalize3(crossje);\n"
" if (dot3F4(DeltaC2,crossje)<0)\n"
" crossje *= -1.f;\n"
"\n"
" float dist;\n"
" bool result = true;\n"
" {\n"
@@ -456,10 +386,8 @@ static const char* satKernelsCL= \
" float d1 = Max1 - Min0;\n"
" dist = d0<d1 ? d0:d1;\n"
" result = true;\n"
"\n"
" }\n"
" \n"
"\n"
" if(dist<*dmin)\n"
" {\n"
" *dmin = dist;\n"
@@ -467,9 +395,7 @@ static const char* satKernelsCL= \
" }\n"
" }\n"
" }\n"
"\n"
" }\n"
"\n"
" \n"
" if((dot3F4(-DeltaC2,*sep))>0.0f)\n"
" {\n"
@@ -477,8 +403,6 @@ static const char* satKernelsCL= \
" }\n"
" return true;\n"
"}\n"
"\n"
"\n"
"inline bool TestSepAxis(__global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
" const float4 posA,const float4 ornA,\n"
" const float4 posB,const float4 ornB,\n"
@@ -488,17 +412,13 @@ static const char* satKernelsCL= \
" float Min1,Max1;\n"
" project(hullA,posA,ornA,sep_axis,vertices, &Min0, &Max0);\n"
" project(hullB,posB,ornB, sep_axis,vertices, &Min1, &Max1);\n"
"\n"
" if(Max0<Min1 || Max1<Min0)\n"
" return false;\n"
"\n"
" float d0 = Max0 - Min1;\n"
" float d1 = Max1 - Min0;\n"
" *depth = d0<d1 ? d0:d1;\n"
" return true;\n"
"}\n"
"\n"
"\n"
"bool findSeparatingAxis( __global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
" const float4 posA1,\n"
" const float4 ornA,\n"
@@ -513,14 +433,12 @@ static const char* satKernelsCL= \
" float* dmin)\n"
"{\n"
" int i = get_global_id(0);\n"
"\n"
" float4 posA = posA1;\n"
" posA.w = 0.f;\n"
" float4 posB = posB1;\n"
" posB.w = 0.f;\n"
" \n"
" int curPlaneTests=0;\n"
"\n"
" {\n"
" int numFacesA = hullA->m_numFaces;\n"
" // Test normals from hullA\n"
@@ -545,8 +463,6 @@ static const char* satKernelsCL= \
" }\n"
" }\n"
" }\n"
"\n"
"\n"
" if((dot3F4(-DeltaC2,*sep))>0.0f)\n"
" {\n"
" *sep = -(*sep);\n"
@@ -554,10 +470,6 @@ static const char* satKernelsCL= \
" \n"
" return true;\n"
"}\n"
"\n"
"\n"
"\n"
"\n"
"bool findSeparatingAxisEdgeEdge( __global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
" const float4 posA1,\n"
" const float4 ornA,\n"
@@ -572,29 +484,22 @@ static const char* satKernelsCL= \
" float* dmin)\n"
"{\n"
" int i = get_global_id(0);\n"
"\n"
" float4 posA = posA1;\n"
" posA.w = 0.f;\n"
" float4 posB = posB1;\n"
" posB.w = 0.f;\n"
"\n"
" int curPlaneTests=0;\n"
"\n"
" int curEdgeEdge = 0;\n"
" // Test edges\n"
" for(int e0=0;e0<hullA->m_numUniqueEdges;e0++)\n"
" {\n"
" const float4 edge0 = uniqueEdges[hullA->m_uniqueEdgesOffset+e0];\n"
" float4 edge0World = qtRotate(ornA,edge0);\n"
"\n"
" for(int e1=0;e1<hullB->m_numUniqueEdges;e1++)\n"
" {\n"
" const float4 edge1 = uniqueEdges[hullB->m_uniqueEdgesOffset+e1];\n"
" float4 edge1World = qtRotate(ornB,edge1);\n"
"\n"
"\n"
" float4 crossje = cross3(edge0World,edge1World);\n"
"\n"
" curEdgeEdge++;\n"
" if(!IsAlmostZero(crossje))\n"
" {\n"
@@ -617,10 +522,8 @@ static const char* satKernelsCL= \
" float d1 = Max1 - Min0;\n"
" dist = d0<d1 ? d0:d1;\n"
" result = true;\n"
"\n"
" }\n"
" \n"
"\n"
" if(dist<*dmin)\n"
" {\n"
" *dmin = dist;\n"
@@ -628,9 +531,7 @@ static const char* satKernelsCL= \
" }\n"
" }\n"
" }\n"
"\n"
" }\n"
"\n"
" \n"
" if((dot3F4(-DeltaC2,*sep))>0.0f)\n"
" {\n"
@@ -638,8 +539,6 @@ static const char* satKernelsCL= \
" }\n"
" return true;\n"
"}\n"
"\n"
"\n"
"// work-in-progress\n"
"__kernel void processCompoundPairsKernel( __global const int4* gpuCompoundPairs,\n"
" __global const BodyData* rigidBodies, \n"
@@ -656,13 +555,11 @@ static const char* satKernelsCL= \
" int numCompoundPairs\n"
" )\n"
"{\n"
"\n"
" int i = get_global_id(0);\n"
" if (i<numCompoundPairs)\n"
" {\n"
" int bodyIndexA = gpuCompoundPairs[i].x;\n"
" int bodyIndexB = gpuCompoundPairs[i].y;\n"
"\n"
" int childShapeIndexA = gpuCompoundPairs[i].z;\n"
" int childShapeIndexB = gpuCompoundPairs[i].w;\n"
" \n"
@@ -711,12 +608,10 @@ static const char* satKernelsCL= \
" int shapeTypeA = collidables[collidableIndexA].m_shapeType;\n"
" int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
" \n"
"\n"
" if ((shapeTypeA != SHAPE_CONVEX_HULL) || (shapeTypeB != SHAPE_CONVEX_HULL))\n"
" {\n"
" return;\n"
" }\n"
"\n"
" int hasSeparatingAxis = 5;\n"
" \n"
" int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
@@ -737,7 +632,6 @@ static const char* satKernelsCL= \
" } else\n"
" {\n"
" bool sepB = findSeparatingAxis( &convexShapes[shapeIndexB],&convexShapes[shapeIndexA],posB,ornB,posA,ornA,DeltaC2,vertices,uniqueEdges,faces,indices,&sepNormal,&dmin);\n"
"\n"
" if (!sepB)\n"
" {\n"
" hasSeparatingAxis = 0;\n"
@@ -756,7 +650,6 @@ static const char* satKernelsCL= \
" }\n"
" \n"
"}\n"
"\n"
"// work-in-progress\n"
"__kernel void findCompoundPairsKernel( __global const int4* pairs, \n"
" __global const BodyData* rigidBodies, \n"
@@ -774,48 +667,36 @@ static const char* satKernelsCL= \
" int maxNumCompoundPairsCapacity\n"
" )\n"
"{\n"
"\n"
" int i = get_global_id(0);\n"
"\n"
" if (i<numPairs)\n"
" {\n"
" int bodyIndexA = pairs[i].x;\n"
" int bodyIndexB = pairs[i].y;\n"
"\n"
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
"\n"
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
"\n"
"\n"
" //once the broadphase avoids static-static pairs, we can remove this test\n"
" if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))\n"
" {\n"
" return;\n"
" }\n"
"\n"
" if ((collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) ||(collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))\n"
" {\n"
"\n"
" if (collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) \n"
" {\n"
"\n"
" int numChildrenA = collidables[collidableIndexA].m_numChildShapes;\n"
" for (int c=0;c<numChildrenA;c++)\n"
" {\n"
" int childShapeIndexA = collidables[collidableIndexA].m_shapeIndex+c;\n"
" int childColIndexA = gpuChildShapes[childShapeIndexA].m_shapeIndex;\n"
"\n"
" float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
" float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
" float4 childPosA = gpuChildShapes[childShapeIndexA].m_childPosition;\n"
" float4 childOrnA = gpuChildShapes[childShapeIndexA].m_childOrientation;\n"
" float4 newPosA = qtRotate(ornA,childPosA)+posA;\n"
" float4 newOrnA = qtMul(ornA,childOrnA);\n"
"\n"
" int shapeIndexA = collidables[childColIndexA].m_shapeIndex;\n"
"\n"
" if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
" {\n"
" int numChildrenB = collidables[collidableIndexB].m_numChildShapes;\n"
@@ -829,9 +710,7 @@ static const char* satKernelsCL= \
" float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
" float4 newPosB = transform(&childPosB,&posB,&ornB);\n"
" float4 newOrnB = qtMul(ornB,childOrnB);\n"
"\n"
" int shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
"\n"
" if (1)\n"
" {\n"
" int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
@@ -847,7 +726,6 @@ static const char* satKernelsCL= \
" float4 ornB =newOrnB;\n"
" float4 c1 = transform(&c1local,&posB,&ornB);\n"
" const float4 DeltaC2 = c0 - c1;\n"
"\n"
" {//\n"
" int compoundPairIdx = atomic_inc(numCompoundPairsOut);\n"
" if (compoundPairIdx<maxNumCompoundPairsCapacity)\n"
@@ -875,7 +753,6 @@ static const char* satKernelsCL= \
" float4 ornB = rigidBodies[bodyIndexB].m_quat;\n"
" float4 c1 = transform(&c1local,&posB,&ornB);\n"
" const float4 DeltaC2 = c0 - c1;\n"
"\n"
" {\n"
" int compoundPairIdx = atomic_inc(numCompoundPairsOut);\n"
" if (compoundPairIdx<maxNumCompoundPairsCapacity)\n"
@@ -902,12 +779,8 @@ static const char* satKernelsCL= \
" float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
" float4 newPosB = qtRotate(ornB,childPosB)+posB;\n"
" float4 newOrnB = qtMul(ornB,childOrnB);\n"
"\n"
" int shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
"\n"
"\n"
" //////////////////////////////////////\n"
"\n"
" if (1)\n"
" {\n"
" int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
@@ -938,7 +811,6 @@ static const char* satKernelsCL= \
" }//fi ((collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) ||(collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))\n"
" }//i<numPairs\n"
"}\n"
"\n"
"// work-in-progress\n"
"__kernel void findSeparatingAxisKernel( __global const int4* pairs, \n"
" __global const BodyData* rigidBodies, \n"
@@ -954,16 +826,13 @@ static const char* satKernelsCL= \
" int numPairs\n"
" )\n"
"{\n"
"\n"
" int i = get_global_id(0);\n"
" \n"
" if (i<numPairs)\n"
" {\n"
"\n"
" \n"
" int bodyIndexA = pairs[i].x;\n"
" int bodyIndexB = pairs[i].y;\n"
"\n"
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
" \n"
@@ -978,7 +847,6 @@ static const char* satKernelsCL= \
" return;\n"
" }\n"
" \n"
"\n"
" if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL))\n"
" {\n"
" hasSeparatingAxis[i] = 0;\n"
@@ -990,11 +858,8 @@ static const char* satKernelsCL= \
" hasSeparatingAxis[i] = 0;\n"
" return;\n"
" }\n"
"\n"
" int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
"\n"
" float dmin = FLT_MAX;\n"
"\n"
" float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
" posA.w = 0.f;\n"
" float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
@@ -1024,7 +889,6 @@ static const char* satKernelsCL= \
" DeltaC2,\n"
" vertices,uniqueEdges,faces,\n"
" indices,&sepNormal,&dmin);\n"
"\n"
" if (!sepB)\n"
" {\n"
" hasSeparatingAxis[i] = 0;\n"
@@ -1047,12 +911,7 @@ static const char* satKernelsCL= \
" }\n"
" \n"
" }\n"
"\n"
"}\n"
"\n"
"\n"
"\n"
"\n"
"// work-in-progress\n"
"__kernel void findConcaveSeparatingAxisKernel( __global int4* concavePairs,\n"
" __global const BodyData* rigidBodies,\n"
@@ -1068,30 +927,22 @@ static const char* satKernelsCL= \
" int numConcavePairs\n"
" )\n"
"{\n"
"\n"
" int i = get_global_id(0);\n"
" if (i>=numConcavePairs)\n"
" return;\n"
" int pairIdx = i;\n"
"\n"
" int bodyIndexA = concavePairs[i].x;\n"
" int bodyIndexB = concavePairs[i].y;\n"
"\n"
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
"\n"
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
"\n"
" if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL&&\n"
" collidables[collidableIndexB].m_shapeType!=SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
" {\n"
" concavePairs[pairIdx].w = -1;\n"
" return;\n"
" }\n"
"\n"
"\n"
"\n"
" int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
" int numActualConcaveConvexTests = 0;\n"
" \n"
@@ -1100,12 +951,10 @@ static const char* satKernelsCL= \
" bool overlap = false;\n"
" \n"
" ConvexPolyhedronCL convexPolyhedronA;\n"
"\n"
" //add 3 vertices of the triangle\n"
" convexPolyhedronA.m_numVertices = 3;\n"
" convexPolyhedronA.m_vertexOffset = 0;\n"
" float4 localCenter = make_float4(0.f,0.f,0.f,0.f);\n"
"\n"
" btGpuFace face = faces[convexShapes[shapeIndexA].m_faceOffset+f];\n"
" float4 triMinAabb, triMaxAabb;\n"
" btAabbCL triAabb;\n"
@@ -1122,9 +971,7 @@ static const char* satKernelsCL= \
" \n"
" triAabb.m_min = min(triAabb.m_min,vert); \n"
" triAabb.m_max = max(triAabb.m_max,vert); \n"
"\n"
" }\n"
"\n"
" overlap = true;\n"
" overlap = (triAabb.m_min.x > aabbs[bodyIndexB].m_max.x || triAabb.m_max.x < aabbs[bodyIndexB].m_min.x) ? false : overlap;\n"
" overlap = (triAabb.m_min.z > aabbs[bodyIndexB].m_max.z || triAabb.m_max.z < aabbs[bodyIndexB].m_min.z) ? false : overlap;\n"
@@ -1135,10 +982,8 @@ static const char* satKernelsCL= \
" float dmin = FLT_MAX;\n"
" int hasSeparatingAxis=5;\n"
" float4 sepAxis=make_float4(1,2,3,4);\n"
"\n"
" int localCC=0;\n"
" numActualConcaveConvexTests++;\n"
"\n"
" //a triangle has 3 unique edges\n"
" convexPolyhedronA.m_numUniqueEdges = 3;\n"
" convexPolyhedronA.m_uniqueEdgesOffset = 0;\n"
@@ -1147,8 +992,6 @@ static const char* satKernelsCL= \
" uniqueEdgesA[0] = (verticesA[1]-verticesA[0]);\n"
" uniqueEdgesA[1] = (verticesA[2]-verticesA[1]);\n"
" uniqueEdgesA[2] = (verticesA[0]-verticesA[2]);\n"
"\n"
"\n"
" convexPolyhedronA.m_faceOffset = 0;\n"
" \n"
" float4 normal = make_float4(face.m_plane.x,face.m_plane.y,face.m_plane.z,0.f);\n"
@@ -1157,7 +1000,6 @@ static const char* satKernelsCL= \
" int indicesA[3+3+2+2+2];\n"
" int curUsedIndices=0;\n"
" int fidx=0;\n"
"\n"
" //front size of triangle\n"
" {\n"
" facesA[fidx].m_indexOffset=curUsedIndices;\n"
@@ -1189,7 +1031,6 @@ static const char* satKernelsCL= \
" facesA[fidx].m_numIndices=3;\n"
" }\n"
" fidx++;\n"
"\n"
" bool addEdgePlanes = true;\n"
" if (addEdgePlanes)\n"
" {\n"
@@ -1202,7 +1043,6 @@ static const char* satKernelsCL= \
" \n"
" float4 edgeNormal = normalize(cross(normal,v1-v0));\n"
" float c = -dot(edgeNormal,v0);\n"
"\n"
" facesA[fidx].m_numIndices = 2;\n"
" facesA[fidx].m_indexOffset=curUsedIndices;\n"
" indicesA[curUsedIndices++]=i;\n"
@@ -1218,22 +1058,15 @@ static const char* satKernelsCL= \
" }\n"
" convexPolyhedronA.m_numFaces = TRIANGLE_NUM_CONVEX_FACES;\n"
" convexPolyhedronA.m_localCenter = localCenter*(1.f/3.f);\n"
"\n"
"\n"
" float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
" posA.w = 0.f;\n"
" float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
" posB.w = 0.f;\n"
"\n"
" float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
" float4 ornB =rigidBodies[bodyIndexB].m_quat;\n"
"\n"
" \n"
"\n"
"\n"
" ///////////////////\n"
" ///compound shape support\n"
"\n"
" if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
" {\n"
" int compoundChild = concavePairs[pairIdx].w;\n"
@@ -1248,14 +1081,11 @@ static const char* satKernelsCL= \
" shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
" }\n"
" //////////////////\n"
"\n"
" float4 c0local = convexPolyhedronA.m_localCenter;\n"
" float4 c0 = transform(&c0local, &posA, &ornA);\n"
" float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
" float4 c1 = transform(&c1local,&posB,&ornB);\n"
" const float4 DeltaC2 = c0 - c1;\n"
"\n"
"\n"
" bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],\n"
" posA,ornA,\n"
" posB,ornB,\n"
@@ -1276,7 +1106,6 @@ static const char* satKernelsCL= \
" vertices,uniqueEdges,faces,indices,\n"
" verticesA,uniqueEdgesA,facesA,indicesA,\n"
" &sepAxis,&dmin);\n"
"\n"
" if (!sepB)\n"
" {\n"
" hasSeparatingAxis = 0;\n"