2 #ifndef B3_NEW_CONTACT_REDUCTION_H
3 #define B3_NEW_CONTACT_REDUCTION_H
5 #include "Bullet3Common/shared/b3Float4.h"
6 #include "Bullet3Collision/NarrowPhaseCollision/shared/b3RigidBodyData.h"
7 #include "Bullet3Collision/NarrowPhaseCollision/shared/b3Contact4Data.h"
9 #define GET_NPOINTS(x) (x).m_worldNormalOnB.w
11 int b3ExtractManifoldSequentialGlobal(__global const b3Float4* p, int nPoints, b3Float4ConstArg nearNormal, b3Int4* contactIdx)
22 b3Float4 center = b3MakeFloat4(0, 0, 0, 0);
24 for (int i = 0; i < nPoints; i++)
26 center /= (float)nPoints;
29 // sample 4 directions
31 b3Float4 aVector = p[0] - center;
32 b3Float4 u = b3Cross(nearNormal, aVector);
33 b3Float4 v = b3Cross(nearNormal, u);
37 //keep point with deepest penetration
49 for (int ie = 0; ie < nPoints; ie++)
57 b3Float4 r = p[ie] - center;
87 if (contactIdx[0].x != minIndex && contactIdx[0].y != minIndex && contactIdx[0].z != minIndex && contactIdx[0].w != minIndex)
89 //replace the first contact with minimum (todo: replace contact with least penetration)
90 contactIdx[0].x = minIndex;
96 __kernel void b3NewContactReductionKernel(__global b3Int4* pairs,
97 __global const b3RigidBodyData_t* rigidBodies,
98 __global const b3Float4* separatingNormals,
99 __global const int* hasSeparatingAxis,
100 __global struct b3Contact4Data* globalContactsOut,
101 __global b3Int4* clippingFaces,
102 __global b3Float4* worldVertsB2,
103 volatile __global int* nGlobalContactsOut,
104 int vertexFaceCapacity,
109 // int i = get_global_id(0);
114 contactIdx = b3MakeInt4(0, 1, 2, 3);
118 if (hasSeparatingAxis[i])
120 int nPoints = clippingFaces[pairIndex].w;
124 __global b3Float4* pointsIn = &worldVertsB2[pairIndex * vertexFaceCapacity];
125 b3Float4 normal = -separatingNormals[i];
127 int nReducedContacts = b3ExtractManifoldSequentialGlobal(pointsIn, nPoints, normal, &contactIdx);
130 dstIdx = b3AtomicInc(nGlobalContactsOut);
133 b3Assert(dstIdx < contactCapacity);
134 if (dstIdx < contactCapacity)
136 __global struct b3Contact4Data* c = &globalContactsOut[dstIdx];
137 c->m_worldNormalOnB = -normal;
138 c->m_restituitionCoeffCmp = (0.f * 0xffff);
139 c->m_frictionCoeffCmp = (0.7f * 0xffff);
140 c->m_batchIdx = pairIndex;
141 int bodyA = pairs[pairIndex].x;
142 int bodyB = pairs[pairIndex].y;
144 pairs[pairIndex].w = dstIdx;
146 c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass == 0 ? -bodyA : bodyA;
147 c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass == 0 ? -bodyB : bodyB;
148 c->m_childIndexA = -1;
149 c->m_childIndexB = -1;
151 switch (nReducedContacts)
154 c->m_worldPosB[3] = pointsIn[contactIdx.w];
156 c->m_worldPosB[2] = pointsIn[contactIdx.z];
158 c->m_worldPosB[1] = pointsIn[contactIdx.y];
160 c->m_worldPosB[0] = pointsIn[contactIdx.x];
166 GET_NPOINTS(*c) = nReducedContacts;
171 } // if (numContactsOut>0)
172 } // if (hasSeparatingAxis[i])