step closer to GPU bvh traversal
This commit is contained in:
erwin coumans
109
opencl/gpu_sat/kernels/bvhTraversal.cl
Normal file
109
opencl/gpu_sat/kernels/bvhTraversal.cl
Normal file
@@ -0,0 +1,109 @@
|
||||
//keep this enum in sync with the CPU version (in btCollidable.h)
|
||||
//written by Erwin Coumans
|
||||
|
||||
#define SHAPE_CONVEX_HULL 3
|
||||
#define SHAPE_CONCAVE_TRIMESH 5
|
||||
#define TRIANGLE_NUM_CONVEX_FACES 5
|
||||
#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6
|
||||
|
||||
typedef unsigned int u32;
|
||||
|
||||
///keep this in sync with btCollidable.h
|
||||
typedef struct
|
||||
{
|
||||
int m_numChildShapes;
|
||||
int blaat2;
|
||||
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;
|
||||
|
||||
|
||||
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
|
||||
{
|
||||
union
|
||||
{
|
||||
float4 m_min;
|
||||
float m_minElems[4];
|
||||
int m_minIndices[4];
|
||||
};
|
||||
union
|
||||
{
|
||||
float4 m_max;
|
||||
float m_maxElems[4];
|
||||
int m_maxIndices[4];
|
||||
};
|
||||
} btAabbCL;
|
||||
|
||||
// work-in-progress
|
||||
__kernel void bvhTraversalKernel( __global const int2* pairs,
|
||||
__global const BodyData* rigidBodies,
|
||||
__global const btCollidableGpu* collidables,
|
||||
__global btAabbCL* aabbs,
|
||||
__global int4* concavePairsOut,
|
||||
__global volatile int* numConcavePairsOut,
|
||||
int numPairs,
|
||||
int maxNumConcavePairsCapacity
|
||||
)
|
||||
{
|
||||
|
||||
int i = get_global_id(0);
|
||||
|
||||
if (i<numPairs)
|
||||
{
|
||||
|
||||
|
||||
int bodyIndexA = pairs[i].x;
|
||||
int bodyIndexB = pairs[i].y;
|
||||
|
||||
int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
|
||||
int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;
|
||||
|
||||
int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
|
||||
int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
|
||||
|
||||
|
||||
//once the broadphase avoids static-static pairs, we can remove this test
|
||||
if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))
|
||||
{
|
||||
return;
|
||||
}
|
||||
|
||||
if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))// && (collidables[collidableIndexB].m_shapeType==SHAPE_CONVEX_HULL))
|
||||
{
|
||||
int pairIdx = atomic_inc(numConcavePairsOut);
|
||||
if (pairIdx<maxNumConcavePairsCapacity)
|
||||
{
|
||||
//int4 newPair;
|
||||
concavePairsOut[pairIdx].x = bodyIndexA;
|
||||
concavePairsOut[pairIdx].y = bodyIndexB;
|
||||
concavePairsOut[pairIdx].z = 5;
|
||||
concavePairsOut[pairIdx].w = 3;
|
||||
}
|
||||
}//SHAPE_CONCAVE_TRIMESH
|
||||
|
||||
}//i<numpairs
|
||||
}
|
||||
112
opencl/gpu_sat/kernels/bvhTraversal.h
Normal file
112
opencl/gpu_sat/kernels/bvhTraversal.h
Normal file
@@ -0,0 +1,112 @@
|
||||
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
|
||||
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"
|
||||
"\n"
|
||||
"typedef unsigned int u32;\n"
|
||||
"\n"
|
||||
"///keep this in sync with btCollidable.h\n"
|
||||
"typedef struct\n"
|
||||
"{\n"
|
||||
" int m_numChildShapes;\n"
|
||||
" int blaat2;\n"
|
||||
" int m_shapeType;\n"
|
||||
" int m_shapeIndex;\n"
|
||||
" \n"
|
||||
"} btCollidableGpu;\n"
|
||||
"\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"
|
||||
"\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"
|
||||
" {\n"
|
||||
" float4 m_min;\n"
|
||||
" float m_minElems[4];\n"
|
||||
" int m_minIndices[4];\n"
|
||||
" };\n"
|
||||
" union\n"
|
||||
" {\n"
|
||||
" float4 m_max;\n"
|
||||
" float m_maxElems[4];\n"
|
||||
" int m_maxIndices[4];\n"
|
||||
" };\n"
|
||||
"} btAabbCL;\n"
|
||||
"\n"
|
||||
"// work-in-progress\n"
|
||||
"__kernel void bvhTraversalKernel( __global const int2* pairs, \n"
|
||||
" __global const BodyData* rigidBodies, \n"
|
||||
" __global const btCollidableGpu* collidables,\n"
|
||||
" __global btAabbCL* aabbs,\n"
|
||||
" __global int4* concavePairsOut,\n"
|
||||
" __global volatile int* numConcavePairsOut,\n"
|
||||
" int numPairs,\n"
|
||||
" int maxNumConcavePairsCapacity\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"
|
||||
" 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_CONCAVE_TRIMESH))// && (collidables[collidableIndexB].m_shapeType==SHAPE_CONVEX_HULL))\n"
|
||||
" {\n"
|
||||
" int pairIdx = atomic_inc(numConcavePairsOut);\n"
|
||||
" if (pairIdx<maxNumConcavePairsCapacity)\n"
|
||||
" {\n"
|
||||
" //int4 newPair;\n"
|
||||
" concavePairsOut[pairIdx].x = bodyIndexA;\n"
|
||||
" concavePairsOut[pairIdx].y = bodyIndexB;\n"
|
||||
" concavePairsOut[pairIdx].z = 5;\n"
|
||||
" concavePairsOut[pairIdx].w = 3;\n"
|
||||
" }\n"
|
||||
" }//SHAPE_CONCAVE_TRIMESH\n"
|
||||
" \n"
|
||||
" }//i<numpairs\n"
|
||||
"}\n"
|
||||
;
|
||||
@@ -1,4 +1,7 @@
|
||||
//keep this enum in sync with the CPU version (in AdlCollisionShape.h)
|
||||
//keep this enum in sync with the CPU version (in btCollidable.h)
|
||||
//written by Erwin Coumans
|
||||
|
||||
|
||||
#define SHAPE_CONVEX_HULL 3
|
||||
#define SHAPE_CONCAVE_TRIMESH 5
|
||||
#define TRIANGLE_NUM_CONVEX_FACES 5
|
||||
|
||||
@@ -1,6 +1,9 @@
|
||||
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
|
||||
static const char* satKernelsCL= \
|
||||
"//keep this enum in sync with the CPU version (in AdlCollisionShape.h)\n"
|
||||
"//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"
|
||||
|
||||
Reference in New Issue
Block a user