2017-08-01 14:30:58 +02:00
|
|
|
|
|
|
|
#ifndef B3_NEW_CONTACT_REDUCTION_H
|
|
|
|
#define B3_NEW_CONTACT_REDUCTION_H
|
|
|
|
|
|
|
|
#include "Bullet3Common/shared/b3Float4.h"
|
|
|
|
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3RigidBodyData.h"
|
|
|
|
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Contact4Data.h"
|
|
|
|
|
|
|
|
#define GET_NPOINTS(x) (x).m_worldNormalOnB.w
|
|
|
|
|
|
|
|
int b3ExtractManifoldSequentialGlobal(__global const b3Float4* p, int nPoints, b3Float4ConstArg nearNormal, b3Int4* contactIdx)
|
|
|
|
{
|
2019-01-03 14:26:51 +01:00
|
|
|
if (nPoints == 0)
|
|
|
|
return 0;
|
|
|
|
|
|
|
|
if (nPoints <= 4)
|
|
|
|
return nPoints;
|
|
|
|
|
|
|
|
if (nPoints > 64)
|
|
|
|
nPoints = 64;
|
|
|
|
|
|
|
|
b3Float4 center = b3MakeFloat4(0, 0, 0, 0);
|
2017-08-01 14:30:58 +02:00
|
|
|
{
|
2019-01-03 14:26:51 +01:00
|
|
|
for (int i = 0; i < nPoints; i++)
|
2017-08-01 14:30:58 +02:00
|
|
|
center += p[i];
|
|
|
|
center /= (float)nPoints;
|
|
|
|
}
|
2019-01-03 14:26:51 +01:00
|
|
|
|
2017-08-01 14:30:58 +02:00
|
|
|
// sample 4 directions
|
2019-01-03 14:26:51 +01:00
|
|
|
|
|
|
|
b3Float4 aVector = p[0] - center;
|
|
|
|
b3Float4 u = b3Cross(nearNormal, aVector);
|
|
|
|
b3Float4 v = b3Cross(nearNormal, u);
|
|
|
|
u = b3Normalized(u);
|
|
|
|
v = b3Normalized(v);
|
|
|
|
|
|
|
|
//keep point with deepest penetration
|
|
|
|
float minW = FLT_MAX;
|
|
|
|
|
|
|
|
int minIndex = -1;
|
|
|
|
|
|
|
|
b3Float4 maxDots;
|
|
|
|
maxDots.x = FLT_MIN;
|
|
|
|
maxDots.y = FLT_MIN;
|
|
|
|
maxDots.z = FLT_MIN;
|
|
|
|
maxDots.w = FLT_MIN;
|
|
|
|
|
|
|
|
// idx, distance
|
|
|
|
for (int ie = 0; ie < nPoints; ie++)
|
|
|
|
{
|
|
|
|
if (p[ie].w < minW)
|
|
|
|
{
|
|
|
|
minW = p[ie].w;
|
|
|
|
minIndex = ie;
|
|
|
|
}
|
|
|
|
float f;
|
|
|
|
b3Float4 r = p[ie] - center;
|
|
|
|
f = b3Dot(u, r);
|
|
|
|
if (f < maxDots.x)
|
|
|
|
{
|
|
|
|
maxDots.x = f;
|
|
|
|
contactIdx[0].x = ie;
|
|
|
|
}
|
|
|
|
|
|
|
|
f = b3Dot(-u, r);
|
|
|
|
if (f < maxDots.y)
|
|
|
|
{
|
|
|
|
maxDots.y = f;
|
|
|
|
contactIdx[0].y = ie;
|
|
|
|
}
|
|
|
|
|
|
|
|
f = b3Dot(v, r);
|
|
|
|
if (f < maxDots.z)
|
|
|
|
{
|
|
|
|
maxDots.z = f;
|
|
|
|
contactIdx[0].z = ie;
|
|
|
|
}
|
|
|
|
|
|
|
|
f = b3Dot(-v, r);
|
|
|
|
if (f < maxDots.w)
|
|
|
|
{
|
|
|
|
maxDots.w = f;
|
|
|
|
contactIdx[0].w = ie;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
if (contactIdx[0].x != minIndex && contactIdx[0].y != minIndex && contactIdx[0].z != minIndex && contactIdx[0].w != minIndex)
|
|
|
|
{
|
|
|
|
//replace the first contact with minimum (todo: replace contact with least penetration)
|
|
|
|
contactIdx[0].x = minIndex;
|
|
|
|
}
|
|
|
|
|
|
|
|
return 4;
|
2017-08-01 14:30:58 +02:00
|
|
|
}
|
|
|
|
|
2019-01-03 14:26:51 +01:00
|
|
|
__kernel void b3NewContactReductionKernel(__global b3Int4* pairs,
|
|
|
|
__global const b3RigidBodyData_t* rigidBodies,
|
|
|
|
__global const b3Float4* separatingNormals,
|
|
|
|
__global const int* hasSeparatingAxis,
|
|
|
|
__global struct b3Contact4Data* globalContactsOut,
|
|
|
|
__global b3Int4* clippingFaces,
|
|
|
|
__global b3Float4* worldVertsB2,
|
|
|
|
volatile __global int* nGlobalContactsOut,
|
|
|
|
int vertexFaceCapacity,
|
|
|
|
int contactCapacity,
|
|
|
|
int numPairs,
|
|
|
|
int pairIndex)
|
2017-08-01 14:30:58 +02:00
|
|
|
{
|
2019-01-03 14:26:51 +01:00
|
|
|
// int i = get_global_id(0);
|
2017-08-01 14:30:58 +02:00
|
|
|
//int pairIndex = i;
|
|
|
|
int i = pairIndex;
|
|
|
|
|
2019-01-03 14:26:51 +01:00
|
|
|
b3Int4 contactIdx;
|
|
|
|
contactIdx = b3MakeInt4(0, 1, 2, 3);
|
|
|
|
|
|
|
|
if (i < numPairs)
|
2017-08-01 14:30:58 +02:00
|
|
|
{
|
|
|
|
if (hasSeparatingAxis[i])
|
|
|
|
{
|
|
|
|
int nPoints = clippingFaces[pairIndex].w;
|
2019-01-03 14:26:51 +01:00
|
|
|
|
|
|
|
if (nPoints > 0)
|
|
|
|
{
|
|
|
|
__global b3Float4* pointsIn = &worldVertsB2[pairIndex * vertexFaceCapacity];
|
|
|
|
b3Float4 normal = -separatingNormals[i];
|
|
|
|
|
|
|
|
int nReducedContacts = b3ExtractManifoldSequentialGlobal(pointsIn, nPoints, normal, &contactIdx);
|
|
|
|
|
|
|
|
int dstIdx;
|
|
|
|
dstIdx = b3AtomicInc(nGlobalContactsOut);
|
|
|
|
|
|
|
|
//#if 0
|
|
|
|
b3Assert(dstIdx < contactCapacity);
|
2017-08-01 14:30:58 +02:00
|
|
|
if (dstIdx < contactCapacity)
|
|
|
|
{
|
|
|
|
__global struct b3Contact4Data* c = &globalContactsOut[dstIdx];
|
|
|
|
c->m_worldNormalOnB = -normal;
|
2019-01-03 14:26:51 +01:00
|
|
|
c->m_restituitionCoeffCmp = (0.f * 0xffff);
|
|
|
|
c->m_frictionCoeffCmp = (0.7f * 0xffff);
|
2017-08-01 14:30:58 +02:00
|
|
|
c->m_batchIdx = pairIndex;
|
|
|
|
int bodyA = pairs[pairIndex].x;
|
|
|
|
int bodyB = pairs[pairIndex].y;
|
|
|
|
|
|
|
|
pairs[pairIndex].w = dstIdx;
|
|
|
|
|
2019-01-03 14:26:51 +01:00
|
|
|
c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass == 0 ? -bodyA : bodyA;
|
|
|
|
c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass == 0 ? -bodyB : bodyB;
|
|
|
|
c->m_childIndexA = -1;
|
|
|
|
c->m_childIndexB = -1;
|
|
|
|
|
|
|
|
switch (nReducedContacts)
|
|
|
|
{
|
|
|
|
case 4:
|
|
|
|
c->m_worldPosB[3] = pointsIn[contactIdx.w];
|
|
|
|
case 3:
|
|
|
|
c->m_worldPosB[2] = pointsIn[contactIdx.z];
|
|
|
|
case 2:
|
|
|
|
c->m_worldPosB[1] = pointsIn[contactIdx.y];
|
|
|
|
case 1:
|
|
|
|
c->m_worldPosB[0] = pointsIn[contactIdx.x];
|
|
|
|
default:
|
|
|
|
{
|
|
|
|
}
|
|
|
|
};
|
|
|
|
|
2017-08-01 14:30:58 +02:00
|
|
|
GET_NPOINTS(*c) = nReducedContacts;
|
2019-01-03 14:26:51 +01:00
|
|
|
}
|
|
|
|
|
|
|
|
//#endif
|
|
|
|
|
|
|
|
} // if (numContactsOut>0)
|
|
|
|
} // if (hasSeparatingAxis[i])
|
|
|
|
} // if (i<numPairs)
|
2017-08-01 14:30:58 +02:00
|
|
|
}
|
|
|
|
#endif
|