mirror of
https://github.com/bulletphysics/bullet3
synced 2024-12-13 21:30:09 +00:00
Use more shared data structures in src\Bullet3OpenCL\NarrowphaseCollision\kernels\satClipHullContacts.cl
Revert default error func (was broken)
This commit is contained in:
erwincoumans
parent
26dfaa441e
commit
1d5c651753
@ -563,11 +563,7 @@ void writeTextureToPng(int textureWidth, int textureHeight, const char* fileName
|
||||
#include "Bullet3Dynamics/ConstraintSolver/b3Generic6DofConstraint.h"
|
||||
#include "Bullet3Dynamics/ConstraintSolver/b3Point2PointConstraint.h"
|
||||
|
||||
void MyErrorFunc(const char* msg)
|
||||
{
|
||||
printf("Error: %s\n",msg);
|
||||
exit(0);
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char* argv[])
|
||||
{
|
||||
@ -580,7 +576,7 @@ int main(int argc, char* argv[])
|
||||
int sz6 = sizeof(b3Transform);
|
||||
|
||||
//b3OpenCLUtils::setCachePath("/Users/erwincoumans/develop/mycache");
|
||||
b3SetCustomErrorMessageFunc(MyErrorFunc);
|
||||
|
||||
|
||||
b3SetCustomEnterProfileZoneFunc(b3ProfileManager::Start_Profile);
|
||||
b3SetCustomLeaveProfileZoneFunc(b3ProfileManager::Stop_Profile);
|
||||
|
||||
@ -38,7 +38,7 @@ struct b3Collidable
|
||||
};
|
||||
|
||||
|
||||
|
||||
typedef struct b3GpuChildShape b3GpuChildShape_t;
|
||||
struct b3GpuChildShape
|
||||
{
|
||||
b3Float4 m_childPosition;
|
||||
|
||||
@ -7,6 +7,7 @@
|
||||
#include "Bullet3Common/shared/b3Float4.h"
|
||||
#include "Bullet3Common/shared/b3Quat.h"
|
||||
|
||||
typedef struct b3GpuFace b3GpuFace_t;
|
||||
struct b3GpuFace
|
||||
{
|
||||
b3Float4 m_plane;
|
||||
@ -16,7 +17,9 @@ struct b3GpuFace
|
||||
int m_unusedPadding2;
|
||||
};
|
||||
|
||||
B3_ATTRIBUTE_ALIGNED16(struct) b3ConvexPolyhedronData
|
||||
typedef struct b3ConvexPolyhedronData b3ConvexPolyhedronData_t;
|
||||
|
||||
struct b3ConvexPolyhedronData
|
||||
{
|
||||
b3Float4 m_localCenter;
|
||||
b3Float4 m_extents;
|
||||
|
||||
@ -136,7 +136,7 @@ m_numCompoundPairsOut(m_context, m_queue)
|
||||
m_findClippingFacesKernel = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,srcClip, "findClippingFacesKernel",&errNum,satClipContactsProg);
|
||||
b3Assert(errNum==CL_SUCCESS);
|
||||
|
||||
m_clipFacesAndContactReductionKernel = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,srcClip, "clipFacesAndContactReductionKernel",&errNum,satClipContactsProg);
|
||||
m_clipFacesAndFindContacts = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,srcClip, "clipFacesAndFindContactsKernel",&errNum,satClipContactsProg);
|
||||
b3Assert(errNum==CL_SUCCESS);
|
||||
|
||||
m_clipHullHullConcaveConvexKernel = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,srcClip, "clipHullHullConcaveConvexKernel",&errNum,satClipContactsProg);
|
||||
@ -155,7 +155,7 @@ m_numCompoundPairsOut(m_context, m_queue)
|
||||
m_clipCompoundsHullHullKernel = 0;
|
||||
m_findClippingFacesKernel = 0;
|
||||
m_newContactReductionKernel=0;
|
||||
m_clipFacesAndContactReductionKernel = 0;
|
||||
m_clipFacesAndFindContacts = 0;
|
||||
m_clipHullHullConcaveConvexKernel = 0;
|
||||
m_extractManifoldAndAddContactKernel = 0;
|
||||
}
|
||||
@ -210,8 +210,8 @@ GpuSatCollision::~GpuSatCollision()
|
||||
if (m_findClippingFacesKernel)
|
||||
clReleaseKernel(m_findClippingFacesKernel);
|
||||
|
||||
if (m_clipFacesAndContactReductionKernel)
|
||||
clReleaseKernel(m_clipFacesAndContactReductionKernel);
|
||||
if (m_clipFacesAndFindContacts)
|
||||
clReleaseKernel(m_clipFacesAndFindContacts);
|
||||
if (m_newContactReductionKernel)
|
||||
clReleaseKernel(m_newContactReductionKernel);
|
||||
if (m_primitiveContactsKernel)
|
||||
@ -3509,7 +3509,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
|
||||
b3BufferInfoCL( m_totalContactsOut.getBufferCL())
|
||||
};
|
||||
|
||||
b3LauncherCL launcher(m_queue, m_clipFacesAndContactReductionKernel,"m_clipFacesAndContactReductionKernel");
|
||||
b3LauncherCL launcher(m_queue, m_clipFacesAndFindContacts,"m_clipFacesAndFindContacts");
|
||||
launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
|
||||
launcher.setConst(vertexFaceCapacity);
|
||||
|
||||
|
||||
@ -33,7 +33,7 @@ struct GpuSatCollision
|
||||
cl_kernel m_clipHullHullKernel;
|
||||
cl_kernel m_clipCompoundsHullHullKernel;
|
||||
|
||||
cl_kernel m_clipFacesAndContactReductionKernel;
|
||||
cl_kernel m_clipFacesAndFindContacts;
|
||||
cl_kernel m_findClippingFacesKernel;
|
||||
|
||||
cl_kernel m_clipHullHullConcaveConvexKernel;
|
||||
|
||||
@ -1,11 +1,6 @@
|
||||
|
||||
#define TRIANGLE_NUM_CONVEX_FACES 5
|
||||
|
||||
#define SHAPE_CONVEX_HULL 3
|
||||
#define SHAPE_PLANE 4
|
||||
#define SHAPE_CONCAVE_TRIMESH 5
|
||||
#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6
|
||||
#define SHAPE_SPHERE 7
|
||||
|
||||
|
||||
#pragma OPENCL EXTENSION cl_amd_printf : enable
|
||||
@ -42,69 +37,15 @@ typedef unsigned int u32;
|
||||
|
||||
|
||||
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Contact4Data.h"
|
||||
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3ConvexPolyhedronData.h"
|
||||
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Collidable.h"
|
||||
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3RigidBodyData.h"
|
||||
|
||||
|
||||
///keep this in sync with btCollidable.h
|
||||
typedef struct
|
||||
{
|
||||
int m_numChildShapes;
|
||||
float m_radius;
|
||||
int m_shapeType;
|
||||
int m_shapeIndex;
|
||||
|
||||
} btCollidableGpu;
|
||||
|
||||
typedef struct
|
||||
{
|
||||
float4 m_childPosition;
|
||||
float4 m_childOrientation;
|
||||
int m_shapeIndex;
|
||||
int m_unused0;
|
||||
int m_unused1;
|
||||
int m_unused2;
|
||||
} btGpuChildShape;
|
||||
|
||||
#define GET_NPOINTS(x) (x).m_worldNormalOnB.w
|
||||
|
||||
typedef struct
|
||||
{
|
||||
float4 m_pos;
|
||||
float4 m_quat;
|
||||
float4 m_linVel;
|
||||
float4 m_angVel;
|
||||
|
||||
u32 m_collidableIdx;
|
||||
float m_invMass;
|
||||
float m_restituitionCoeff;
|
||||
float m_frictionCoeff;
|
||||
} BodyData;
|
||||
|
||||
|
||||
typedef struct
|
||||
{
|
||||
float4 m_localCenter;
|
||||
float4 m_extents;
|
||||
float4 mC;
|
||||
float4 mE;
|
||||
|
||||
float m_radius;
|
||||
int m_faceOffset;
|
||||
int m_numFaces;
|
||||
int m_numVertices;
|
||||
|
||||
int m_vertexOffset;
|
||||
int m_uniqueEdgesOffset;
|
||||
int m_numUniqueEdges;
|
||||
int m_unused;
|
||||
|
||||
} ConvexPolyhedronCL;
|
||||
|
||||
typedef struct
|
||||
{
|
||||
float4 m_plane;
|
||||
int m_indexOffset;
|
||||
int m_numIndices;
|
||||
} btGpuFace;
|
||||
|
||||
#define SELECT_UINT4( b, a, condition ) select( b,a,condition )
|
||||
|
||||
@ -341,12 +282,12 @@ int clipFace(const float4* pVtxIn, int numVertsIn, float4 planeNormalWS,float pl
|
||||
}
|
||||
|
||||
|
||||
int clipFaceAgainstHull(const float4 separatingNormal, __global const ConvexPolyhedronCL* hullA,
|
||||
int clipFaceAgainstHull(const float4 separatingNormal, __global const b3ConvexPolyhedronData_t* hullA,
|
||||
const float4 posA, const Quaternion ornA, float4* worldVertsB1, int numWorldVertsB1,
|
||||
float4* worldVertsB2, int capacityWorldVertsB2,
|
||||
const float minDist, float maxDist,
|
||||
__global const float4* vertices,
|
||||
__global const btGpuFace* faces,
|
||||
__global const b3GpuFace_t* faces,
|
||||
__global const int* indices,
|
||||
float4* contactsOut,
|
||||
int contactCapacity)
|
||||
@ -381,7 +322,7 @@ int clipFaceAgainstHull(const float4 separatingNormal, __global const ConvexPoly
|
||||
if (closestFaceA<0)
|
||||
return numContactsOut;
|
||||
|
||||
btGpuFace polyA = faces[hullA->m_faceOffset+closestFaceA];
|
||||
b3GpuFace_t polyA = faces[hullA->m_faceOffset+closestFaceA];
|
||||
|
||||
// clip polygon to back of planes of all faces of hull A that are adjacent to witness face
|
||||
int numVerticesA = polyA.m_numIndices;
|
||||
@ -442,15 +383,15 @@ int clipFaceAgainstHull(const float4 separatingNormal, __global const ConvexPoly
|
||||
|
||||
|
||||
|
||||
int clipFaceAgainstHullLocalA(const float4 separatingNormal, const ConvexPolyhedronCL* hullA,
|
||||
int clipFaceAgainstHullLocalA(const float4 separatingNormal, const b3ConvexPolyhedronData_t* hullA,
|
||||
const float4 posA, const Quaternion ornA, float4* worldVertsB1, int numWorldVertsB1,
|
||||
float4* worldVertsB2, int capacityWorldVertsB2,
|
||||
const float minDist, float maxDist,
|
||||
const float4* verticesA,
|
||||
const btGpuFace* facesA,
|
||||
const b3GpuFace_t* facesA,
|
||||
const int* indicesA,
|
||||
__global const float4* verticesB,
|
||||
__global const btGpuFace* facesB,
|
||||
__global const b3GpuFace_t* facesB,
|
||||
__global const int* indicesB,
|
||||
float4* contactsOut,
|
||||
int contactCapacity)
|
||||
@ -485,7 +426,7 @@ int clipFaceAgainstHullLocalA(const float4 separatingNormal, const ConvexPolyhed
|
||||
if (closestFaceA<0)
|
||||
return numContactsOut;
|
||||
|
||||
btGpuFace polyA = facesA[hullA->m_faceOffset+closestFaceA];
|
||||
b3GpuFace_t polyA = facesA[hullA->m_faceOffset+closestFaceA];
|
||||
|
||||
// clip polygon to back of planes of all faces of hull A that are adjacent to witness face
|
||||
int numVerticesA = polyA.m_numIndices;
|
||||
@ -545,12 +486,12 @@ int clipFaceAgainstHullLocalA(const float4 separatingNormal, const ConvexPolyhed
|
||||
}
|
||||
|
||||
int clipHullAgainstHull(const float4 separatingNormal,
|
||||
__global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB,
|
||||
__global const b3ConvexPolyhedronData_t* hullA, __global const b3ConvexPolyhedronData_t* hullB,
|
||||
const float4 posA, const Quaternion ornA,const float4 posB, const Quaternion ornB,
|
||||
float4* worldVertsB1, float4* worldVertsB2, int capacityWorldVerts,
|
||||
const float minDist, float maxDist,
|
||||
__global const float4* vertices,
|
||||
__global const btGpuFace* faces,
|
||||
__global const b3GpuFace_t* faces,
|
||||
__global const int* indices,
|
||||
float4* localContactsOut,
|
||||
int localContactCapacity)
|
||||
@ -578,7 +519,7 @@ int clipHullAgainstHull(const float4 separatingNormal,
|
||||
}
|
||||
|
||||
{
|
||||
const btGpuFace polyB = faces[hullB->m_faceOffset+closestFaceB];
|
||||
const b3GpuFace_t polyB = faces[hullB->m_faceOffset+closestFaceB];
|
||||
const int numVertices = polyB.m_numIndices;
|
||||
for(int e0=0;e0<numVertices;e0++)
|
||||
{
|
||||
@ -601,15 +542,15 @@ int clipHullAgainstHull(const float4 separatingNormal,
|
||||
|
||||
|
||||
int clipHullAgainstHullLocalA(const float4 separatingNormal,
|
||||
const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB,
|
||||
const b3ConvexPolyhedronData_t* hullA, __global const b3ConvexPolyhedronData_t* hullB,
|
||||
const float4 posA, const Quaternion ornA,const float4 posB, const Quaternion ornB,
|
||||
float4* worldVertsB1, float4* worldVertsB2, int capacityWorldVerts,
|
||||
const float minDist, float maxDist,
|
||||
const float4* verticesA,
|
||||
const btGpuFace* facesA,
|
||||
const b3GpuFace_t* facesA,
|
||||
const int* indicesA,
|
||||
__global const float4* verticesB,
|
||||
__global const btGpuFace* facesB,
|
||||
__global const b3GpuFace_t* facesB,
|
||||
__global const int* indicesB,
|
||||
float4* localContactsOut,
|
||||
int localContactCapacity)
|
||||
@ -637,7 +578,7 @@ int clipHullAgainstHullLocalA(const float4 separatingNormal,
|
||||
}
|
||||
|
||||
{
|
||||
const btGpuFace polyB = facesB[hullB->m_faceOffset+closestFaceB];
|
||||
const b3GpuFace_t polyB = facesB[hullB->m_faceOffset+closestFaceB];
|
||||
const int numVertices = polyB.m_numIndices;
|
||||
for(int e0=0;e0<numVertices;e0++)
|
||||
{
|
||||
@ -871,7 +812,7 @@ int extractManifoldSequential(const float4* p, int nPoints, float4 nearNormal, i
|
||||
|
||||
|
||||
__kernel void extractManifoldAndAddContactKernel(__global const int4* pairs,
|
||||
__global const BodyData* rigidBodies,
|
||||
__global const b3RigidBodyData_t* rigidBodies,
|
||||
__global const float4* closestPointsWorld,
|
||||
__global const float4* separatingNormalsWorld,
|
||||
__global const int* contactCounts,
|
||||
@ -946,12 +887,12 @@ void trMul(float4 translationA, Quaternion orientationA,
|
||||
|
||||
|
||||
__kernel void clipHullHullKernel( __global int4* pairs,
|
||||
__global const BodyData* rigidBodies,
|
||||
__global const btCollidableGpu* collidables,
|
||||
__global const ConvexPolyhedronCL* convexShapes,
|
||||
__global const b3RigidBodyData_t* rigidBodies,
|
||||
__global const b3Collidable_t* collidables,
|
||||
__global const b3ConvexPolyhedronData_t* convexShapes,
|
||||
__global const float4* vertices,
|
||||
__global const float4* uniqueEdges,
|
||||
__global const btGpuFace* faces,
|
||||
__global const b3GpuFace_t* faces,
|
||||
__global const int* indices,
|
||||
__global const float4* separatingNormals,
|
||||
__global const int* hasSeparatingAxis,
|
||||
@ -1048,14 +989,14 @@ __kernel void clipHullHullKernel( __global int4* pairs,
|
||||
|
||||
|
||||
__kernel void clipCompoundsHullHullKernel( __global const int4* gpuCompoundPairs,
|
||||
__global const BodyData* rigidBodies,
|
||||
__global const btCollidableGpu* collidables,
|
||||
__global const ConvexPolyhedronCL* convexShapes,
|
||||
__global const b3RigidBodyData_t* rigidBodies,
|
||||
__global const b3Collidable_t* collidables,
|
||||
__global const b3ConvexPolyhedronData_t* convexShapes,
|
||||
__global const float4* vertices,
|
||||
__global const float4* uniqueEdges,
|
||||
__global const btGpuFace* faces,
|
||||
__global const b3GpuFace_t* faces,
|
||||
__global const int* indices,
|
||||
__global const btGpuChildShape* gpuChildShapes,
|
||||
__global const b3GpuChildShape_t* gpuChildShapes,
|
||||
__global const float4* gpuCompoundSepNormalsOut,
|
||||
__global const int* gpuHasCompoundSepNormalsOut,
|
||||
__global struct b3Contact4Data* restrict globalContactsOut,
|
||||
@ -1181,8 +1122,8 @@ __kernel void clipCompoundsHullHullKernel( __global const int4* gpuCompoundPai
|
||||
|
||||
|
||||
__kernel void sphereSphereCollisionKernel( __global const int4* pairs,
|
||||
__global const BodyData* rigidBodies,
|
||||
__global const btCollidableGpu* collidables,
|
||||
__global const b3RigidBodyData_t* rigidBodies,
|
||||
__global const b3Collidable_t* collidables,
|
||||
__global const float4* separatingNormals,
|
||||
__global const int* hasSeparatingAxis,
|
||||
__global struct b3Contact4Data* restrict globalContactsOut,
|
||||
@ -1251,14 +1192,14 @@ __kernel void sphereSphereCollisionKernel( __global const int4* pairs,
|
||||
}
|
||||
|
||||
__kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
|
||||
__global const BodyData* rigidBodies,
|
||||
__global const btCollidableGpu* collidables,
|
||||
__global const ConvexPolyhedronCL* convexShapes,
|
||||
__global const b3RigidBodyData_t* rigidBodies,
|
||||
__global const b3Collidable_t* collidables,
|
||||
__global const b3ConvexPolyhedronData_t* convexShapes,
|
||||
__global const float4* vertices,
|
||||
__global const float4* uniqueEdges,
|
||||
__global const btGpuFace* faces,
|
||||
__global const b3GpuFace_t* faces,
|
||||
__global const int* indices,
|
||||
__global const btGpuChildShape* gpuChildShapes,
|
||||
__global const b3GpuChildShape_t* gpuChildShapes,
|
||||
__global const float4* separatingNormals,
|
||||
__global struct b3Contact4Data* restrict globalContactsOut,
|
||||
counter32_t nGlobalContactsOut,
|
||||
@ -1300,14 +1241,14 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
|
||||
|
||||
bool overlap = false;
|
||||
|
||||
ConvexPolyhedronCL convexPolyhedronA;
|
||||
b3ConvexPolyhedronData_t convexPolyhedronA;
|
||||
|
||||
//add 3 vertices of the triangle
|
||||
convexPolyhedronA.m_numVertices = 3;
|
||||
convexPolyhedronA.m_vertexOffset = 0;
|
||||
float4 localCenter = make_float4(0.f,0.f,0.f,0.f);
|
||||
|
||||
btGpuFace face = faces[convexShapes[shapeIndexA].m_faceOffset+f];
|
||||
b3GpuFace_t face = faces[convexShapes[shapeIndexA].m_faceOffset+f];
|
||||
|
||||
float4 verticesA[3];
|
||||
for (int i=0;i<3;i++)
|
||||
@ -1336,7 +1277,7 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
|
||||
|
||||
float4 normal = make_float4(face.m_plane.x,face.m_plane.y,face.m_plane.z,0.f);
|
||||
|
||||
btGpuFace facesA[TRIANGLE_NUM_CONVEX_FACES];
|
||||
b3GpuFace_t facesA[TRIANGLE_NUM_CONVEX_FACES];
|
||||
int indicesA[3+3+2+2+2];
|
||||
int curUsedIndices=0;
|
||||
int fidx=0;
|
||||
@ -1491,7 +1432,7 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
|
||||
|
||||
|
||||
int findClippingFaces(const float4 separatingNormal,
|
||||
__global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB,
|
||||
__global const b3ConvexPolyhedronData_t* hullA, __global const b3ConvexPolyhedronData_t* hullB,
|
||||
const float4 posA, const Quaternion ornA,const float4 posB, const Quaternion ornB,
|
||||
__global float4* worldVertsA1,
|
||||
__global float4* worldNormalsA1,
|
||||
@ -1499,7 +1440,7 @@ int findClippingFaces(const float4 separatingNormal,
|
||||
int capacityWorldVerts,
|
||||
const float minDist, float maxDist,
|
||||
__global const float4* vertices,
|
||||
__global const btGpuFace* faces,
|
||||
__global const b3GpuFace_t* faces,
|
||||
__global const int* indices,
|
||||
__global int4* clippingFaces, int pairIndex)
|
||||
{
|
||||
@ -1526,7 +1467,7 @@ int findClippingFaces(const float4 separatingNormal,
|
||||
}
|
||||
|
||||
{
|
||||
const btGpuFace polyB = faces[hullB->m_faceOffset+closestFaceB];
|
||||
const b3GpuFace_t polyB = faces[hullB->m_faceOffset+closestFaceB];
|
||||
const int numVertices = polyB.m_numIndices;
|
||||
for(int e0=0;e0<numVertices;e0++)
|
||||
{
|
||||
@ -1666,12 +1607,12 @@ int clipFaces(__global float4* worldVertsA1,
|
||||
|
||||
|
||||
__kernel void findClippingFacesKernel( __global const int4* pairs,
|
||||
__global const BodyData* rigidBodies,
|
||||
__global const btCollidableGpu* collidables,
|
||||
__global const ConvexPolyhedronCL* convexShapes,
|
||||
__global const b3RigidBodyData_t* rigidBodies,
|
||||
__global const b3Collidable_t* collidables,
|
||||
__global const b3ConvexPolyhedronData_t* convexShapes,
|
||||
__global const float4* vertices,
|
||||
__global const float4* uniqueEdges,
|
||||
__global const btGpuFace* faces,
|
||||
__global const b3GpuFace_t* faces,
|
||||
__global const int* indices,
|
||||
__global const float4* separatingNormals,
|
||||
__global const int* hasSeparatingAxis,
|
||||
@ -1728,8 +1669,8 @@ __kernel void findClippingFacesKernel( __global const int4* pairs,
|
||||
|
||||
|
||||
|
||||
__kernel void clipFacesAndContactReductionKernel( __global int4* pairs,
|
||||
__global const BodyData* rigidBodies,
|
||||
__kernel void clipFacesAndFindContactsKernel( __global int4* pairs,
|
||||
__global const b3RigidBodyData_t* rigidBodies,
|
||||
__global const float4* separatingNormals,
|
||||
__global const int* hasSeparatingAxis,
|
||||
__global struct b3Contact4Data* globalContactsOut,
|
||||
@ -1842,7 +1783,7 @@ __kernel void clipFacesAndContactReductionKernel( __global int4* pairs,
|
||||
|
||||
|
||||
__kernel void newContactReductionKernel( __global int4* pairs,
|
||||
__global const BodyData* rigidBodies,
|
||||
__global const b3RigidBodyData_t* rigidBodies,
|
||||
__global const float4* separatingNormals,
|
||||
__global const int* hasSeparatingAxis,
|
||||
__global struct b3Contact4Data* globalContactsOut,
|
||||
|
||||
@ -1,11 +1,6 @@
|
||||
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
|
||||
static const char* satClipKernelsCL= \
|
||||
"#define TRIANGLE_NUM_CONVEX_FACES 5\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"
|
||||
"#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"
|
||||
@ -103,59 +98,329 @@ static const char* satClipKernelsCL= \
|
||||
" contact->m_worldNormalOnB.w = (float)numPoints;\n"
|
||||
"};\n"
|
||||
"#endif //B3_CONTACT4DATA_H\n"
|
||||
"///keep this in sync with btCollidable.h\n"
|
||||
"typedef struct\n"
|
||||
"{\n"
|
||||
" int m_numChildShapes;\n"
|
||||
" float m_radius;\n"
|
||||
" int m_shapeType;\n"
|
||||
" int m_shapeIndex;\n"
|
||||
"#ifndef B3_CONVEX_POLYHEDRON_DATA_H\n"
|
||||
"#define B3_CONVEX_POLYHEDRON_DATA_H\n"
|
||||
"#ifndef B3_FLOAT4_H\n"
|
||||
"#ifdef __cplusplus\n"
|
||||
"#else\n"
|
||||
"#endif \n"
|
||||
"#endif //B3_FLOAT4_H\n"
|
||||
"#ifndef B3_QUAT_H\n"
|
||||
"#define B3_QUAT_H\n"
|
||||
"#ifndef B3_PLATFORM_DEFINITIONS_H\n"
|
||||
"#ifdef __cplusplus\n"
|
||||
"#else\n"
|
||||
"#endif\n"
|
||||
"#endif\n"
|
||||
"#ifndef B3_FLOAT4_H\n"
|
||||
"#ifdef __cplusplus\n"
|
||||
"#else\n"
|
||||
"#endif \n"
|
||||
"#endif //B3_FLOAT4_H\n"
|
||||
"#ifdef __cplusplus\n"
|
||||
"#else\n"
|
||||
" typedef float4 b3Quat;\n"
|
||||
" #define b3QuatConstArg const b3Quat\n"
|
||||
" \n"
|
||||
"} btCollidableGpu;\n"
|
||||
"typedef struct\n"
|
||||
"{\n"
|
||||
" float4 m_childPosition;\n"
|
||||
" float4 m_childOrientation;\n"
|
||||
" int m_shapeIndex;\n"
|
||||
" int m_unused0;\n"
|
||||
" int m_unused1;\n"
|
||||
" int m_unused2;\n"
|
||||
"} btGpuChildShape;\n"
|
||||
"#define GET_NPOINTS(x) (x).m_worldNormalOnB.w\n"
|
||||
"typedef struct\n"
|
||||
"{\n"
|
||||
" float4 m_pos;\n"
|
||||
" float4 m_quat;\n"
|
||||
" float4 m_linVel;\n"
|
||||
" float4 m_angVel;\n"
|
||||
" u32 m_collidableIdx; \n"
|
||||
" float m_invMass;\n"
|
||||
" float m_restituitionCoeff;\n"
|
||||
" float m_frictionCoeff;\n"
|
||||
"} BodyData;\n"
|
||||
"typedef struct \n"
|
||||
"{\n"
|
||||
" float4 m_localCenter;\n"
|
||||
" float4 m_extents;\n"
|
||||
" float4 mC;\n"
|
||||
" float4 mE;\n"
|
||||
" \n"
|
||||
"inline float4 b3FastNormalize4(float4 v)\n"
|
||||
"{\n"
|
||||
" v = (float4)(v.xyz,0.f);\n"
|
||||
" return fast_normalize(v);\n"
|
||||
"}\n"
|
||||
" \n"
|
||||
"inline b3Quat b3QuatMul(b3Quat a, b3Quat b);\n"
|
||||
"inline b3Quat b3QuatNormalized(b3QuatConstArg in);\n"
|
||||
"inline b3Quat b3QuatRotate(b3QuatConstArg q, b3QuatConstArg vec);\n"
|
||||
"inline b3Quat b3QuatInvert(b3QuatConstArg q);\n"
|
||||
"inline b3Quat b3QuatMul(b3QuatConstArg a, b3QuatConstArg b)\n"
|
||||
"{\n"
|
||||
" b3Quat ans;\n"
|
||||
" ans = b3Cross3( a, b );\n"
|
||||
" ans += a.w*b+b.w*a;\n"
|
||||
"// ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z);\n"
|
||||
" ans.w = a.w*b.w - b3Dot3F4(a, b);\n"
|
||||
" return ans;\n"
|
||||
"}\n"
|
||||
"inline b3Quat b3QuatNormalized(b3QuatConstArg in)\n"
|
||||
"{\n"
|
||||
" b3Quat q;\n"
|
||||
" q=in;\n"
|
||||
" //return b3FastNormalize4(in);\n"
|
||||
" float len = native_sqrt(dot(q, q));\n"
|
||||
" if(len > 0.f)\n"
|
||||
" {\n"
|
||||
" q *= 1.f / len;\n"
|
||||
" }\n"
|
||||
" else\n"
|
||||
" {\n"
|
||||
" q.x = q.y = q.z = 0.f;\n"
|
||||
" q.w = 1.f;\n"
|
||||
" }\n"
|
||||
" return q;\n"
|
||||
"}\n"
|
||||
"inline float4 b3QuatRotate(b3QuatConstArg q, b3QuatConstArg vec)\n"
|
||||
"{\n"
|
||||
" b3Quat qInv = b3QuatInvert( q );\n"
|
||||
" float4 vcpy = vec;\n"
|
||||
" vcpy.w = 0.f;\n"
|
||||
" float4 out = b3QuatMul(b3QuatMul(q,vcpy),qInv);\n"
|
||||
" return out;\n"
|
||||
"}\n"
|
||||
"inline b3Quat b3QuatInvert(b3QuatConstArg q)\n"
|
||||
"{\n"
|
||||
" return (b3Quat)(-q.xyz, q.w);\n"
|
||||
"}\n"
|
||||
"inline float4 b3QuatInvRotate(b3QuatConstArg q, b3QuatConstArg vec)\n"
|
||||
"{\n"
|
||||
" return b3QuatRotate( b3QuatInvert( q ), vec );\n"
|
||||
"}\n"
|
||||
"inline b3Float4 b3TransformPoint(b3Float4ConstArg point, b3Float4ConstArg translation, b3QuatConstArg orientation)\n"
|
||||
"{\n"
|
||||
" return b3QuatRotate( orientation, point ) + (translation);\n"
|
||||
"}\n"
|
||||
" \n"
|
||||
"#endif \n"
|
||||
"#endif //B3_QUAT_H\n"
|
||||
"typedef struct b3GpuFace b3GpuFace_t;\n"
|
||||
"struct b3GpuFace\n"
|
||||
"{\n"
|
||||
" b3Float4 m_plane;\n"
|
||||
" int m_indexOffset;\n"
|
||||
" int m_numIndices;\n"
|
||||
" int m_unusedPadding1;\n"
|
||||
" int m_unusedPadding2;\n"
|
||||
"};\n"
|
||||
"typedef struct b3ConvexPolyhedronData b3ConvexPolyhedronData_t;\n"
|
||||
"struct b3ConvexPolyhedronData\n"
|
||||
"{\n"
|
||||
" b3Float4 m_localCenter;\n"
|
||||
" b3Float4 m_extents;\n"
|
||||
" b3Float4 mC;\n"
|
||||
" b3Float4 mE;\n"
|
||||
" float m_radius;\n"
|
||||
" 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"
|
||||
"#endif //B3_CONVEX_POLYHEDRON_DATA_H\n"
|
||||
"#ifndef B3_COLLIDABLE_H\n"
|
||||
"#define B3_COLLIDABLE_H\n"
|
||||
"#ifndef B3_FLOAT4_H\n"
|
||||
"#ifdef __cplusplus\n"
|
||||
"#else\n"
|
||||
"#endif \n"
|
||||
"#endif //B3_FLOAT4_H\n"
|
||||
"#ifndef B3_QUAT_H\n"
|
||||
"#ifdef __cplusplus\n"
|
||||
"#else\n"
|
||||
"#endif \n"
|
||||
"#endif //B3_QUAT_H\n"
|
||||
"enum b3ShapeTypes\n"
|
||||
"{\n"
|
||||
" SHAPE_HEIGHT_FIELD=1,\n"
|
||||
" SHAPE_CONVEX_HULL=3,\n"
|
||||
" SHAPE_PLANE=4,\n"
|
||||
" SHAPE_CONCAVE_TRIMESH=5,\n"
|
||||
" SHAPE_COMPOUND_OF_CONVEX_HULLS=6,\n"
|
||||
" SHAPE_SPHERE=7,\n"
|
||||
" MAX_NUM_SHAPE_TYPES,\n"
|
||||
"};\n"
|
||||
"typedef struct b3Collidable b3Collidable_t;\n"
|
||||
"struct b3Collidable\n"
|
||||
"{\n"
|
||||
" union {\n"
|
||||
" int m_numChildShapes;\n"
|
||||
" int m_bvhIndex;\n"
|
||||
" };\n"
|
||||
" union\n"
|
||||
" {\n"
|
||||
" float m_radius;\n"
|
||||
" int m_compoundBvhIndex;\n"
|
||||
" };\n"
|
||||
" int m_shapeType;\n"
|
||||
" int m_shapeIndex;\n"
|
||||
"};\n"
|
||||
"typedef struct b3GpuChildShape b3GpuChildShape_t;\n"
|
||||
"struct b3GpuChildShape\n"
|
||||
"{\n"
|
||||
" b3Float4 m_childPosition;\n"
|
||||
" b3Quat m_childOrientation;\n"
|
||||
" int m_shapeIndex;\n"
|
||||
" int m_unused0;\n"
|
||||
" int m_unused1;\n"
|
||||
" int m_unused2;\n"
|
||||
"};\n"
|
||||
"struct b3CompoundOverlappingPair\n"
|
||||
"{\n"
|
||||
" int m_bodyIndexA;\n"
|
||||
" int m_bodyIndexB;\n"
|
||||
"// int m_pairType;\n"
|
||||
" int m_childShapeIndexA;\n"
|
||||
" int m_childShapeIndexB;\n"
|
||||
"};\n"
|
||||
"#endif //B3_COLLIDABLE_H\n"
|
||||
"#ifndef B3_RIGIDBODY_DATA_H\n"
|
||||
"#define B3_RIGIDBODY_DATA_H\n"
|
||||
"#ifndef B3_FLOAT4_H\n"
|
||||
"#ifdef __cplusplus\n"
|
||||
"#else\n"
|
||||
"#endif \n"
|
||||
"#endif //B3_FLOAT4_H\n"
|
||||
"#ifndef B3_QUAT_H\n"
|
||||
"#ifdef __cplusplus\n"
|
||||
"#else\n"
|
||||
"#endif \n"
|
||||
"#endif //B3_QUAT_H\n"
|
||||
"#ifndef B3_MAT3x3_H\n"
|
||||
"#define B3_MAT3x3_H\n"
|
||||
"#ifndef B3_QUAT_H\n"
|
||||
"#ifdef __cplusplus\n"
|
||||
"#else\n"
|
||||
"#endif \n"
|
||||
"#endif //B3_QUAT_H\n"
|
||||
"#ifdef __cplusplus\n"
|
||||
"#else\n"
|
||||
"typedef struct\n"
|
||||
"{\n"
|
||||
" float4 m_plane;\n"
|
||||
" int m_indexOffset;\n"
|
||||
" int m_numIndices;\n"
|
||||
"} btGpuFace;\n"
|
||||
" b3Float4 m_row[3];\n"
|
||||
"}b3Mat3x3;\n"
|
||||
"#define b3Mat3x3ConstArg const b3Mat3x3\n"
|
||||
"#define b3GetRow(m,row) (m.m_row[row])\n"
|
||||
"inline b3Mat3x3 b3QuatGetRotationMatrix(b3Quat quat)\n"
|
||||
"{\n"
|
||||
" b3Float4 quat2 = (b3Float4)(quat.x*quat.x, quat.y*quat.y, quat.z*quat.z, 0.f);\n"
|
||||
" b3Mat3x3 out;\n"
|
||||
" out.m_row[0].x=1-2*quat2.y-2*quat2.z;\n"
|
||||
" out.m_row[0].y=2*quat.x*quat.y-2*quat.w*quat.z;\n"
|
||||
" out.m_row[0].z=2*quat.x*quat.z+2*quat.w*quat.y;\n"
|
||||
" out.m_row[0].w = 0.f;\n"
|
||||
" out.m_row[1].x=2*quat.x*quat.y+2*quat.w*quat.z;\n"
|
||||
" out.m_row[1].y=1-2*quat2.x-2*quat2.z;\n"
|
||||
" out.m_row[1].z=2*quat.y*quat.z-2*quat.w*quat.x;\n"
|
||||
" out.m_row[1].w = 0.f;\n"
|
||||
" out.m_row[2].x=2*quat.x*quat.z-2*quat.w*quat.y;\n"
|
||||
" out.m_row[2].y=2*quat.y*quat.z+2*quat.w*quat.x;\n"
|
||||
" out.m_row[2].z=1-2*quat2.x-2*quat2.y;\n"
|
||||
" out.m_row[2].w = 0.f;\n"
|
||||
" return out;\n"
|
||||
"}\n"
|
||||
"inline b3Mat3x3 b3AbsoluteMat3x3(b3Mat3x3ConstArg matIn)\n"
|
||||
"{\n"
|
||||
" b3Mat3x3 out;\n"
|
||||
" out.m_row[0] = fabs(matIn.m_row[0]);\n"
|
||||
" out.m_row[1] = fabs(matIn.m_row[1]);\n"
|
||||
" out.m_row[2] = fabs(matIn.m_row[2]);\n"
|
||||
" return out;\n"
|
||||
"}\n"
|
||||
"__inline\n"
|
||||
"b3Mat3x3 mtZero();\n"
|
||||
"__inline\n"
|
||||
"b3Mat3x3 mtIdentity();\n"
|
||||
"__inline\n"
|
||||
"b3Mat3x3 mtTranspose(b3Mat3x3 m);\n"
|
||||
"__inline\n"
|
||||
"b3Mat3x3 mtMul(b3Mat3x3 a, b3Mat3x3 b);\n"
|
||||
"__inline\n"
|
||||
"b3Float4 mtMul1(b3Mat3x3 a, b3Float4 b);\n"
|
||||
"__inline\n"
|
||||
"b3Float4 mtMul3(b3Float4 a, b3Mat3x3 b);\n"
|
||||
"__inline\n"
|
||||
"b3Mat3x3 mtZero()\n"
|
||||
"{\n"
|
||||
" b3Mat3x3 m;\n"
|
||||
" m.m_row[0] = (b3Float4)(0.f);\n"
|
||||
" m.m_row[1] = (b3Float4)(0.f);\n"
|
||||
" m.m_row[2] = (b3Float4)(0.f);\n"
|
||||
" return m;\n"
|
||||
"}\n"
|
||||
"__inline\n"
|
||||
"b3Mat3x3 mtIdentity()\n"
|
||||
"{\n"
|
||||
" b3Mat3x3 m;\n"
|
||||
" m.m_row[0] = (b3Float4)(1,0,0,0);\n"
|
||||
" m.m_row[1] = (b3Float4)(0,1,0,0);\n"
|
||||
" m.m_row[2] = (b3Float4)(0,0,1,0);\n"
|
||||
" return m;\n"
|
||||
"}\n"
|
||||
"__inline\n"
|
||||
"b3Mat3x3 mtTranspose(b3Mat3x3 m)\n"
|
||||
"{\n"
|
||||
" b3Mat3x3 out;\n"
|
||||
" out.m_row[0] = (b3Float4)(m.m_row[0].x, m.m_row[1].x, m.m_row[2].x, 0.f);\n"
|
||||
" out.m_row[1] = (b3Float4)(m.m_row[0].y, m.m_row[1].y, m.m_row[2].y, 0.f);\n"
|
||||
" out.m_row[2] = (b3Float4)(m.m_row[0].z, m.m_row[1].z, m.m_row[2].z, 0.f);\n"
|
||||
" return out;\n"
|
||||
"}\n"
|
||||
"__inline\n"
|
||||
"b3Mat3x3 mtMul(b3Mat3x3 a, b3Mat3x3 b)\n"
|
||||
"{\n"
|
||||
" b3Mat3x3 transB;\n"
|
||||
" transB = mtTranspose( b );\n"
|
||||
" b3Mat3x3 ans;\n"
|
||||
" // why this doesn't run when 0ing in the for{}\n"
|
||||
" a.m_row[0].w = 0.f;\n"
|
||||
" a.m_row[1].w = 0.f;\n"
|
||||
" a.m_row[2].w = 0.f;\n"
|
||||
" for(int i=0; i<3; i++)\n"
|
||||
" {\n"
|
||||
"// a.m_row[i].w = 0.f;\n"
|
||||
" ans.m_row[i].x = b3Dot3F4(a.m_row[i],transB.m_row[0]);\n"
|
||||
" ans.m_row[i].y = b3Dot3F4(a.m_row[i],transB.m_row[1]);\n"
|
||||
" ans.m_row[i].z = b3Dot3F4(a.m_row[i],transB.m_row[2]);\n"
|
||||
" ans.m_row[i].w = 0.f;\n"
|
||||
" }\n"
|
||||
" return ans;\n"
|
||||
"}\n"
|
||||
"__inline\n"
|
||||
"b3Float4 mtMul1(b3Mat3x3 a, b3Float4 b)\n"
|
||||
"{\n"
|
||||
" b3Float4 ans;\n"
|
||||
" ans.x = b3Dot3F4( a.m_row[0], b );\n"
|
||||
" ans.y = b3Dot3F4( a.m_row[1], b );\n"
|
||||
" ans.z = b3Dot3F4( a.m_row[2], b );\n"
|
||||
" ans.w = 0.f;\n"
|
||||
" return ans;\n"
|
||||
"}\n"
|
||||
"__inline\n"
|
||||
"b3Float4 mtMul3(b3Float4 a, b3Mat3x3 b)\n"
|
||||
"{\n"
|
||||
" b3Float4 colx = b3MakeFloat4(b.m_row[0].x, b.m_row[1].x, b.m_row[2].x, 0);\n"
|
||||
" b3Float4 coly = b3MakeFloat4(b.m_row[0].y, b.m_row[1].y, b.m_row[2].y, 0);\n"
|
||||
" b3Float4 colz = b3MakeFloat4(b.m_row[0].z, b.m_row[1].z, b.m_row[2].z, 0);\n"
|
||||
" b3Float4 ans;\n"
|
||||
" ans.x = b3Dot3F4( a, colx );\n"
|
||||
" ans.y = b3Dot3F4( a, coly );\n"
|
||||
" ans.z = b3Dot3F4( a, colz );\n"
|
||||
" return ans;\n"
|
||||
"}\n"
|
||||
"#endif\n"
|
||||
"#endif //B3_MAT3x3_H\n"
|
||||
"typedef struct b3RigidBodyData b3RigidBodyData_t;\n"
|
||||
"struct b3RigidBodyData\n"
|
||||
"{\n"
|
||||
" b3Float4 m_pos;\n"
|
||||
" b3Quat m_quat;\n"
|
||||
" b3Float4 m_linVel;\n"
|
||||
" b3Float4 m_angVel;\n"
|
||||
" int m_collidableIdx;\n"
|
||||
" float m_invMass;\n"
|
||||
" float m_restituitionCoeff;\n"
|
||||
" float m_frictionCoeff;\n"
|
||||
"};\n"
|
||||
"typedef struct b3InertiaData b3InertiaData_t;\n"
|
||||
"struct b3InertiaData\n"
|
||||
"{\n"
|
||||
" b3Mat3x3 m_invInertiaWorld;\n"
|
||||
" b3Mat3x3 m_initInvInertia;\n"
|
||||
"};\n"
|
||||
"#endif //B3_RIGIDBODY_DATA_H\n"
|
||||
" \n"
|
||||
"#define GET_NPOINTS(x) (x).m_worldNormalOnB.w\n"
|
||||
"#define SELECT_UINT4( b, a, condition ) select( b,a,condition )\n"
|
||||
"#define make_float4 (float4)\n"
|
||||
"#define make_float2 (float2)\n"
|
||||
@ -350,12 +615,12 @@ static const char* satClipKernelsCL= \
|
||||
" }\n"
|
||||
" return numVertsOut;\n"
|
||||
"}\n"
|
||||
"int clipFaceAgainstHull(const float4 separatingNormal, __global const ConvexPolyhedronCL* hullA, \n"
|
||||
"int clipFaceAgainstHull(const float4 separatingNormal, __global const b3ConvexPolyhedronData_t* hullA, \n"
|
||||
" const float4 posA, const Quaternion ornA, float4* worldVertsB1, int numWorldVertsB1,\n"
|
||||
" float4* worldVertsB2, int capacityWorldVertsB2,\n"
|
||||
" const float minDist, float maxDist,\n"
|
||||
" __global const float4* vertices,\n"
|
||||
" __global const btGpuFace* faces,\n"
|
||||
" __global const b3GpuFace_t* faces,\n"
|
||||
" __global const int* indices,\n"
|
||||
" float4* contactsOut,\n"
|
||||
" int contactCapacity)\n"
|
||||
@ -387,7 +652,7 @@ static const char* satClipKernelsCL= \
|
||||
" }\n"
|
||||
" if (closestFaceA<0)\n"
|
||||
" return numContactsOut;\n"
|
||||
" btGpuFace polyA = faces[hullA->m_faceOffset+closestFaceA];\n"
|
||||
" b3GpuFace_t polyA = faces[hullA->m_faceOffset+closestFaceA];\n"
|
||||
" // clip polygon to back of planes of all faces of hull A that are adjacent to witness face\n"
|
||||
" int numVerticesA = polyA.m_numIndices;\n"
|
||||
" for(int e0=0;e0<numVerticesA;e0++)\n"
|
||||
@ -439,15 +704,15 @@ static const char* satClipKernelsCL= \
|
||||
" }\n"
|
||||
" return numContactsOut;\n"
|
||||
"}\n"
|
||||
"int clipFaceAgainstHullLocalA(const float4 separatingNormal, const ConvexPolyhedronCL* hullA, \n"
|
||||
"int clipFaceAgainstHullLocalA(const float4 separatingNormal, const b3ConvexPolyhedronData_t* hullA, \n"
|
||||
" const float4 posA, const Quaternion ornA, float4* worldVertsB1, int numWorldVertsB1,\n"
|
||||
" float4* worldVertsB2, int capacityWorldVertsB2,\n"
|
||||
" const float minDist, float maxDist,\n"
|
||||
" const float4* verticesA,\n"
|
||||
" const btGpuFace* facesA,\n"
|
||||
" const b3GpuFace_t* facesA,\n"
|
||||
" const int* indicesA,\n"
|
||||
" __global const float4* verticesB,\n"
|
||||
" __global const btGpuFace* facesB,\n"
|
||||
" __global const b3GpuFace_t* facesB,\n"
|
||||
" __global const int* indicesB,\n"
|
||||
" float4* contactsOut,\n"
|
||||
" int contactCapacity)\n"
|
||||
@ -479,7 +744,7 @@ static const char* satClipKernelsCL= \
|
||||
" }\n"
|
||||
" if (closestFaceA<0)\n"
|
||||
" return numContactsOut;\n"
|
||||
" btGpuFace polyA = facesA[hullA->m_faceOffset+closestFaceA];\n"
|
||||
" b3GpuFace_t polyA = facesA[hullA->m_faceOffset+closestFaceA];\n"
|
||||
" // clip polygon to back of planes of all faces of hull A that are adjacent to witness face\n"
|
||||
" int numVerticesA = polyA.m_numIndices;\n"
|
||||
" for(int e0=0;e0<numVerticesA;e0++)\n"
|
||||
@ -532,12 +797,12 @@ static const char* satClipKernelsCL= \
|
||||
" return numContactsOut;\n"
|
||||
"}\n"
|
||||
"int clipHullAgainstHull(const float4 separatingNormal,\n"
|
||||
" __global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
|
||||
" __global const b3ConvexPolyhedronData_t* hullA, __global const b3ConvexPolyhedronData_t* hullB, \n"
|
||||
" const float4 posA, const Quaternion ornA,const float4 posB, const Quaternion ornB, \n"
|
||||
" float4* worldVertsB1, float4* worldVertsB2, int capacityWorldVerts,\n"
|
||||
" const float minDist, float maxDist,\n"
|
||||
" __global const float4* vertices,\n"
|
||||
" __global const btGpuFace* faces,\n"
|
||||
" __global const b3GpuFace_t* faces,\n"
|
||||
" __global const int* indices,\n"
|
||||
" float4* localContactsOut,\n"
|
||||
" int localContactCapacity)\n"
|
||||
@ -561,7 +826,7 @@ static const char* satClipKernelsCL= \
|
||||
" }\n"
|
||||
" }\n"
|
||||
" {\n"
|
||||
" const btGpuFace polyB = faces[hullB->m_faceOffset+closestFaceB];\n"
|
||||
" const b3GpuFace_t polyB = faces[hullB->m_faceOffset+closestFaceB];\n"
|
||||
" const int numVertices = polyB.m_numIndices;\n"
|
||||
" for(int e0=0;e0<numVertices;e0++)\n"
|
||||
" {\n"
|
||||
@ -580,15 +845,15 @@ static const char* satClipKernelsCL= \
|
||||
" return numContactsOut;\n"
|
||||
"}\n"
|
||||
"int clipHullAgainstHullLocalA(const float4 separatingNormal,\n"
|
||||
" const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
|
||||
" const b3ConvexPolyhedronData_t* hullA, __global const b3ConvexPolyhedronData_t* hullB, \n"
|
||||
" const float4 posA, const Quaternion ornA,const float4 posB, const Quaternion ornB, \n"
|
||||
" float4* worldVertsB1, float4* worldVertsB2, int capacityWorldVerts,\n"
|
||||
" const float minDist, float maxDist,\n"
|
||||
" const float4* verticesA,\n"
|
||||
" const btGpuFace* facesA,\n"
|
||||
" const b3GpuFace_t* facesA,\n"
|
||||
" const int* indicesA,\n"
|
||||
" __global const float4* verticesB,\n"
|
||||
" __global const btGpuFace* facesB,\n"
|
||||
" __global const b3GpuFace_t* facesB,\n"
|
||||
" __global const int* indicesB,\n"
|
||||
" float4* localContactsOut,\n"
|
||||
" int localContactCapacity)\n"
|
||||
@ -612,7 +877,7 @@ static const char* satClipKernelsCL= \
|
||||
" }\n"
|
||||
" }\n"
|
||||
" {\n"
|
||||
" const btGpuFace polyB = facesB[hullB->m_faceOffset+closestFaceB];\n"
|
||||
" const b3GpuFace_t polyB = facesB[hullB->m_faceOffset+closestFaceB];\n"
|
||||
" const int numVertices = polyB.m_numIndices;\n"
|
||||
" for(int e0=0;e0<numVertices;e0++)\n"
|
||||
" {\n"
|
||||
@ -816,7 +1081,7 @@ static const char* satClipKernelsCL= \
|
||||
" }\n"
|
||||
"}\n"
|
||||
"__kernel void extractManifoldAndAddContactKernel(__global const int4* pairs, \n"
|
||||
" __global const BodyData* rigidBodies, \n"
|
||||
" __global const b3RigidBodyData_t* rigidBodies, \n"
|
||||
" __global const float4* closestPointsWorld,\n"
|
||||
" __global const float4* separatingNormalsWorld,\n"
|
||||
" __global const int* contactCounts,\n"
|
||||
@ -881,12 +1146,12 @@ static const char* satClipKernelsCL= \
|
||||
" *translationOut = transform(&translationB,&translationA,&orientationA);\n"
|
||||
"}\n"
|
||||
"__kernel void clipHullHullKernel( __global int4* pairs, \n"
|
||||
" __global const BodyData* rigidBodies, \n"
|
||||
" __global const btCollidableGpu* collidables,\n"
|
||||
" __global const ConvexPolyhedronCL* convexShapes, \n"
|
||||
" __global const b3RigidBodyData_t* rigidBodies, \n"
|
||||
" __global const b3Collidable_t* collidables,\n"
|
||||
" __global const b3ConvexPolyhedronData_t* convexShapes, \n"
|
||||
" __global const float4* vertices,\n"
|
||||
" __global const float4* uniqueEdges,\n"
|
||||
" __global const btGpuFace* faces,\n"
|
||||
" __global const b3GpuFace_t* faces,\n"
|
||||
" __global const int* indices,\n"
|
||||
" __global const float4* separatingNormals,\n"
|
||||
" __global const int* hasSeparatingAxis,\n"
|
||||
@ -969,14 +1234,14 @@ static const char* satClipKernelsCL= \
|
||||
" }// if (i<numPairs)\n"
|
||||
"}\n"
|
||||
"__kernel void clipCompoundsHullHullKernel( __global const int4* gpuCompoundPairs, \n"
|
||||
" __global const BodyData* rigidBodies, \n"
|
||||
" __global const btCollidableGpu* collidables,\n"
|
||||
" __global const ConvexPolyhedronCL* convexShapes, \n"
|
||||
" __global const b3RigidBodyData_t* rigidBodies, \n"
|
||||
" __global const b3Collidable_t* collidables,\n"
|
||||
" __global const b3ConvexPolyhedronData_t* convexShapes, \n"
|
||||
" __global const float4* vertices,\n"
|
||||
" __global const float4* uniqueEdges,\n"
|
||||
" __global const btGpuFace* faces,\n"
|
||||
" __global const b3GpuFace_t* faces,\n"
|
||||
" __global const int* indices,\n"
|
||||
" __global const btGpuChildShape* gpuChildShapes,\n"
|
||||
" __global const b3GpuChildShape_t* gpuChildShapes,\n"
|
||||
" __global const float4* gpuCompoundSepNormalsOut,\n"
|
||||
" __global const int* gpuHasCompoundSepNormalsOut,\n"
|
||||
" __global struct b3Contact4Data* restrict globalContactsOut,\n"
|
||||
@ -1092,8 +1357,8 @@ static const char* satClipKernelsCL= \
|
||||
" }// if (i<numCompoundPairs)\n"
|
||||
"}\n"
|
||||
"__kernel void sphereSphereCollisionKernel( __global const int4* pairs, \n"
|
||||
" __global const BodyData* rigidBodies, \n"
|
||||
" __global const btCollidableGpu* collidables,\n"
|
||||
" __global const b3RigidBodyData_t* rigidBodies, \n"
|
||||
" __global const b3Collidable_t* collidables,\n"
|
||||
" __global const float4* separatingNormals,\n"
|
||||
" __global const int* hasSeparatingAxis,\n"
|
||||
" __global struct b3Contact4Data* restrict globalContactsOut,\n"
|
||||
@ -1157,14 +1422,14 @@ static const char* satClipKernelsCL= \
|
||||
" }//if (i<numPairs)\n"
|
||||
"} \n"
|
||||
"__kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,\n"
|
||||
" __global const BodyData* rigidBodies, \n"
|
||||
" __global const btCollidableGpu* collidables,\n"
|
||||
" __global const ConvexPolyhedronCL* convexShapes, \n"
|
||||
" __global const b3RigidBodyData_t* rigidBodies, \n"
|
||||
" __global const b3Collidable_t* collidables,\n"
|
||||
" __global const b3ConvexPolyhedronData_t* convexShapes, \n"
|
||||
" __global const float4* vertices,\n"
|
||||
" __global const float4* uniqueEdges,\n"
|
||||
" __global const btGpuFace* faces,\n"
|
||||
" __global const b3GpuFace_t* faces,\n"
|
||||
" __global const int* indices,\n"
|
||||
" __global const btGpuChildShape* gpuChildShapes,\n"
|
||||
" __global const b3GpuChildShape_t* gpuChildShapes,\n"
|
||||
" __global const float4* separatingNormals,\n"
|
||||
" __global struct b3Contact4Data* restrict globalContactsOut,\n"
|
||||
" counter32_t nGlobalContactsOut,\n"
|
||||
@ -1202,12 +1467,12 @@ static const char* satClipKernelsCL= \
|
||||
" \n"
|
||||
" bool overlap = false;\n"
|
||||
" \n"
|
||||
" ConvexPolyhedronCL convexPolyhedronA;\n"
|
||||
" b3ConvexPolyhedronData_t convexPolyhedronA;\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"
|
||||
" btGpuFace face = faces[convexShapes[shapeIndexA].m_faceOffset+f];\n"
|
||||
" b3GpuFace_t face = faces[convexShapes[shapeIndexA].m_faceOffset+f];\n"
|
||||
" \n"
|
||||
" float4 verticesA[3];\n"
|
||||
" for (int i=0;i<3;i++)\n"
|
||||
@ -1231,7 +1496,7 @@ static const char* satClipKernelsCL= \
|
||||
" \n"
|
||||
" float4 normal = make_float4(face.m_plane.x,face.m_plane.y,face.m_plane.z,0.f);\n"
|
||||
" \n"
|
||||
" btGpuFace facesA[TRIANGLE_NUM_CONVEX_FACES];\n"
|
||||
" b3GpuFace_t facesA[TRIANGLE_NUM_CONVEX_FACES];\n"
|
||||
" int indicesA[3+3+2+2+2];\n"
|
||||
" int curUsedIndices=0;\n"
|
||||
" int fidx=0;\n"
|
||||
@ -1372,7 +1637,7 @@ static const char* satClipKernelsCL= \
|
||||
" }// if (i<numPairs)\n"
|
||||
"}\n"
|
||||
"int findClippingFaces(const float4 separatingNormal,\n"
|
||||
" __global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB,\n"
|
||||
" __global const b3ConvexPolyhedronData_t* hullA, __global const b3ConvexPolyhedronData_t* hullB,\n"
|
||||
" const float4 posA, const Quaternion ornA,const float4 posB, const Quaternion ornB,\n"
|
||||
" __global float4* worldVertsA1,\n"
|
||||
" __global float4* worldNormalsA1,\n"
|
||||
@ -1380,7 +1645,7 @@ static const char* satClipKernelsCL= \
|
||||
" int capacityWorldVerts,\n"
|
||||
" const float minDist, float maxDist,\n"
|
||||
" __global const float4* vertices,\n"
|
||||
" __global const btGpuFace* faces,\n"
|
||||
" __global const b3GpuFace_t* faces,\n"
|
||||
" __global const int* indices,\n"
|
||||
" __global int4* clippingFaces, int pairIndex)\n"
|
||||
"{\n"
|
||||
@ -1407,7 +1672,7 @@ static const char* satClipKernelsCL= \
|
||||
" }\n"
|
||||
" \n"
|
||||
" {\n"
|
||||
" const btGpuFace polyB = faces[hullB->m_faceOffset+closestFaceB];\n"
|
||||
" const b3GpuFace_t polyB = faces[hullB->m_faceOffset+closestFaceB];\n"
|
||||
" const int numVertices = polyB.m_numIndices;\n"
|
||||
" for(int e0=0;e0<numVertices;e0++)\n"
|
||||
" {\n"
|
||||
@ -1536,12 +1801,12 @@ static const char* satClipKernelsCL= \
|
||||
" return numContactsOut;\n"
|
||||
"}\n"
|
||||
"__kernel void findClippingFacesKernel( __global const int4* pairs,\n"
|
||||
" __global const BodyData* rigidBodies,\n"
|
||||
" __global const btCollidableGpu* collidables,\n"
|
||||
" __global const ConvexPolyhedronCL* convexShapes,\n"
|
||||
" __global const b3RigidBodyData_t* rigidBodies,\n"
|
||||
" __global const b3Collidable_t* collidables,\n"
|
||||
" __global const b3ConvexPolyhedronData_t* convexShapes,\n"
|
||||
" __global const float4* vertices,\n"
|
||||
" __global const float4* uniqueEdges,\n"
|
||||
" __global const btGpuFace* faces,\n"
|
||||
" __global const b3GpuFace_t* faces,\n"
|
||||
" __global const int* indices,\n"
|
||||
" __global const float4* separatingNormals,\n"
|
||||
" __global const int* hasSeparatingAxis,\n"
|
||||
@ -1594,8 +1859,8 @@ static const char* satClipKernelsCL= \
|
||||
" }// if (i<numPairs)\n"
|
||||
" \n"
|
||||
"}\n"
|
||||
"__kernel void clipFacesAndContactReductionKernel( __global int4* pairs,\n"
|
||||
" __global const BodyData* rigidBodies,\n"
|
||||
"__kernel void clipFacesAndFindContactsKernel( __global int4* pairs,\n"
|
||||
" __global const b3RigidBodyData_t* rigidBodies,\n"
|
||||
" __global const float4* separatingNormals,\n"
|
||||
" __global const int* hasSeparatingAxis,\n"
|
||||
" __global struct b3Contact4Data* globalContactsOut,\n"
|
||||
@ -1700,7 +1965,7 @@ static const char* satClipKernelsCL= \
|
||||
" \n"
|
||||
"}\n"
|
||||
"__kernel void newContactReductionKernel( __global int4* pairs,\n"
|
||||
" __global const BodyData* rigidBodies,\n"
|
||||
" __global const b3RigidBodyData_t* rigidBodies,\n"
|
||||
" __global const float4* separatingNormals,\n"
|
||||
" __global const int* hasSeparatingAxis,\n"
|
||||
" __global struct b3Contact4Data* globalContactsOut,\n"
|
||||
|
||||
@ -327,6 +327,7 @@ static const char* updateAabbsKernelCL= \
|
||||
" int m_shapeType;\n"
|
||||
" int m_shapeIndex;\n"
|
||||
"};\n"
|
||||
"typedef struct b3GpuChildShape b3GpuChildShape_t;\n"
|
||||
"struct b3GpuChildShape\n"
|
||||
"{\n"
|
||||
" b3Float4 m_childPosition;\n"
|
||||
|
||||
Loading…
Reference in New Issue
Block a user