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

further work towards sharing structures/code between C++ and OpenCL

Browse Source 
(could break the build in C++/OpenCL)
This commit is contained in:
erwincoumans
2013-08-03 10:13:13 -07:00
parent 906415429c
commit 18eed5b067
8 changed files with 1031 additions and 831 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
src/Bullet3Collision/NarrowPhaseCollision/b3Contact4.h
View File

@@ -17,9 +17,9 @@ subject to the following restrictions:
#define B3_CONTACT4_H
#include "Bullet3Common/b3Vector3.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Contact4Data.h"
B3_ATTRIBUTE_ALIGNED16(struct) b3Contact4
B3_ATTRIBUTE_ALIGNED16(struct) b3Contact4 : public b3Contact4Data
{
B3_DECLARE_ALIGNED_ALLOCATOR();

12
src/Bullet3Common/shared/b3Float4.h Normal file
View File

@@ -0,0 +1,12 @@
#ifndef B3_FLOAT4_H
#define B3_FLOAT4_H
#include "Bullet3Common/shared/b3PlatformDefinitions.h"
#ifdef __cplusplus
#define b3Float4 b3Vector3
#else//bla
typedef float4 b3Float4;
#endif
#endif //B3_FLOAT4_H

9
src/Bullet3Common/shared/b3PlatformDefinitions.h Normal file
View File

@@ -0,0 +1,9 @@
#ifndef B3_PLATFORM_DEFINITIONS_H
#define B3_PLATFORM_DEFINITIONS_H
struct MyTest
{
int bla;
};
#endif

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

@@ -1,3 +1,4 @@
/*
Bullet Continuous Collision Detection and Physics Library
Copyright (c) 2003-2006 Erwin Coumans http://continuousphysics.com/Bullet/

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

@@ -1,4 +1,4 @@
#define TRIANGLE_NUM_CONVEX_FACES 5
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Contact4Data.h"
#define SHAPE_CONVEX_HULL 3
#define SHAPE_PLANE 4
@@ -40,22 +40,6 @@ 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;
int m_unused1;
int m_unused2;
} Contact4;
typedef struct
{
@@ -361,7 +345,7 @@ void computeContactSphereConvex(int pairIndex,
__global const float4* convexVertices,
__global const int* convexIndices,
__global const btGpuFace* faces,
__global Contact4* restrict globalContactsOut,
__global b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int maxContactCapacity,
float4 spherePos2,
@@ -482,7 +466,7 @@ void computeContactSphereConvex(int pairIndex,
if (1)//dstIdx < maxContactCapacity)
{
__global Contact4* c = &globalContactsOut[dstIdx];
__global b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = normalOnSurfaceB1;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_batchIdx = pairIndex;
@@ -606,7 +590,7 @@ int computeContactPlaneConvex(int pairIndex,
__global const float4* convexVertices,
__global const int* convexIndices,
__global const btGpuFace* faces,
__global Contact4* restrict globalContactsOut,
__global b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int maxContactCapacity,
float4 posB,
@@ -708,7 +692,7 @@ int computeContactPlaneConvex(int pairIndex,
if (dstIdx < maxContactCapacity)
{
resultIndex = dstIdx;
__global Contact4* c = &globalContactsOut[dstIdx];
__global b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = planeNormalWorld;
//c->setFrictionCoeff(0.7);
//c->setRestituitionCoeff(0.f);
@@ -748,7 +732,7 @@ void computeContactPlaneSphere(int pairIndex,
__global const BodyData* rigidBodies,
__global const btCollidableGpu* collidables,
__global const btGpuFace* faces,
__global Contact4* restrict globalContactsOut,
__global b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int maxContactCapacity)
{
@@ -791,7 +775,7 @@ void computeContactPlaneSphere(int pairIndex,
if (dstIdx < maxContactCapacity)
{
__global Contact4* c = &globalContactsOut[dstIdx];
__global b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = normalOnSurfaceB1;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_batchIdx = pairIndex;
@@ -814,7 +798,7 @@ __kernel void primitiveContactsKernel( __global int4* pairs,
__global const float4* uniqueEdges,
__global const btGpuFace* faces,
__global const int* indices,
__global Contact4* restrict globalContactsOut,
__global b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int numPairs, int maxContactCapacity)
{
@@ -969,7 +953,7 @@ __kernel void primitiveContactsKernel( __global int4* pairs,
if (dstIdx < maxContactCapacity)
{
__global Contact4* c = &globalContactsOut[dstIdx];
__global b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = -normalOnSurfaceB;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_batchIdx = pairIndex;
@@ -1003,7 +987,7 @@ __kernel void processCompoundPairsPrimitivesKernel( __global const int4* gpuCo
__global const int* indices,
__global btAabbCL* aabbs,
__global const btGpuChildShape* gpuChildShapes,
__global Contact4* restrict globalContactsOut,
__global b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int numCompoundPairs, int maxContactCapacity
)
@@ -1182,7 +1166,7 @@ void computeContactSphereTriangle(int pairIndex,
__global const BodyData* rigidBodies,
__global const btCollidableGpu* collidables,
const float4* triangleVertices,
__global Contact4* restrict globalContactsOut,
__global b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int maxContactCapacity,
float4 spherePos2,
@@ -1309,7 +1293,7 @@ void computeContactSphereTriangle(int pairIndex,
if (dstIdx < maxContactCapacity)
{
__global Contact4* c = &globalContactsOut[dstIdx];
__global b3Contact4Data* c = &globalContactsOut[dstIdx];
c->m_worldNormal = normalOnSurfaceB1;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
c->m_batchIdx = pairIndex;
@@ -1341,7 +1325,7 @@ __kernel void findConcaveSphereContactsKernel( __global int4* concavePairs,
__global const btGpuFace* faces,
__global const int* indices,
__global btAabbCL* aabbs,
__global Contact4* restrict globalContactsOut,
__global b3Contact4Data* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
int numConcavePairs, int maxContactCapacity
)

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

@@ -1,6 +1,47 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
static const char* primitiveContactsKernelsCL= \
"#define TRIANGLE_NUM_CONVEX_FACES 5\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"struct MyTest\n"
"{\n"
" int bla;\n"
"};\n"
"\n"
"\n"
"\n"
"\n"
"\n"
"\n"
" typedef float4 b3Float4;\n"
"\n"
"\n"
"\n"
"\n"
"typedef struct\n"
"{\n"
" b3Float4 m_worldPos[4];\n"
" b3Float4 m_worldNormal; // w: m_nPoints\n"
" unsigned int m_coeffs;\n"
" unsigned int m_batchIdx;\n"
"\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"
"#define SHAPE_CONVEX_HULL 3\n"
"#define SHAPE_PLANE 4\n"
@@ -42,22 +83,6 @@ static const char* primitiveContactsKernelsCL= \
"\n"
"\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_worldPos[4];\n"
" float4 m_worldNormal; // w: m_nPoints\n"
" u32 m_coeffs;\n"
" u32 m_batchIdx;\n"
"\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"
"} Contact4;\n"
"\n"
"typedef struct \n"
"{\n"
@@ -363,7 +388,7 @@ static const char* primitiveContactsKernelsCL= \
" __global const float4* convexVertices,\n"
" __global const int* convexIndices,\n"
" __global const btGpuFace* faces,\n"
" __global Contact4* restrict globalContactsOut,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int maxContactCapacity,\n"
" float4 spherePos2,\n"
@@ -484,7 +509,7 @@ static const char* primitiveContactsKernelsCL= \
" \n"
" if (1)//dstIdx < maxContactCapacity)\n"
" {\n"
" __global Contact4* c = &globalContactsOut[dstIdx];\n"
" __global 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_batchIdx = pairIndex;\n"
@@ -608,7 +633,7 @@ static const char* primitiveContactsKernelsCL= \
" __global const float4* convexVertices,\n"
" __global const int* convexIndices,\n"
" __global const btGpuFace* faces,\n"
" __global Contact4* restrict globalContactsOut,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int maxContactCapacity,\n"
" float4 posB,\n"
@@ -710,7 +735,7 @@ static const char* primitiveContactsKernelsCL= \
" if (dstIdx < maxContactCapacity)\n"
" {\n"
" resultIndex = dstIdx;\n"
" __global Contact4* c = &globalContactsOut[dstIdx];\n"
" __global b3Contact4Data* c = &globalContactsOut[dstIdx];\n"
" c->m_worldNormal = planeNormalWorld;\n"
" //c->setFrictionCoeff(0.7);\n"
" //c->setRestituitionCoeff(0.f);\n"
@@ -750,7 +775,7 @@ static const char* primitiveContactsKernelsCL= \
" __global const BodyData* rigidBodies, \n"
" __global const btCollidableGpu* collidables,\n"
" __global const btGpuFace* faces,\n"
" __global Contact4* restrict globalContactsOut,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int maxContactCapacity)\n"
"{\n"
@@ -793,7 +818,7 @@ static const char* primitiveContactsKernelsCL= \
" \n"
" if (dstIdx < maxContactCapacity)\n"
" {\n"
" __global Contact4* c = &globalContactsOut[dstIdx];\n"
" __global 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_batchIdx = pairIndex;\n"
@@ -816,7 +841,7 @@ static const char* primitiveContactsKernelsCL= \
" __global const float4* uniqueEdges,\n"
" __global const btGpuFace* faces,\n"
" __global const int* indices,\n"
" __global Contact4* restrict globalContactsOut,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int numPairs, int maxContactCapacity)\n"
"{\n"
@@ -971,7 +996,7 @@ static const char* primitiveContactsKernelsCL= \
" \n"
" if (dstIdx < maxContactCapacity)\n"
" {\n"
" __global Contact4* c = &globalContactsOut[dstIdx];\n"
" __global 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_batchIdx = pairIndex;\n"
@@ -1005,7 +1030,7 @@ static const char* primitiveContactsKernelsCL= \
" __global const int* indices,\n"
" __global btAabbCL* aabbs,\n"
" __global const btGpuChildShape* gpuChildShapes,\n"
" __global Contact4* restrict globalContactsOut,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int numCompoundPairs, int maxContactCapacity\n"
" )\n"
@@ -1184,7 +1209,7 @@ static const char* primitiveContactsKernelsCL= \
" __global const BodyData* rigidBodies, \n"
" __global const btCollidableGpu* collidables,\n"
" const float4* triangleVertices,\n"
" __global Contact4* restrict globalContactsOut,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int maxContactCapacity,\n"
" float4 spherePos2,\n"
@@ -1311,7 +1336,7 @@ static const char* primitiveContactsKernelsCL= \
" \n"
" if (dstIdx < maxContactCapacity)\n"
" {\n"
" __global Contact4* c = &globalContactsOut[dstIdx];\n"
" __global 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_batchIdx = pairIndex;\n"
@@ -1343,7 +1368,7 @@ static const char* primitiveContactsKernelsCL= \
" __global const btGpuFace* faces,\n"
" __global const int* indices,\n"
" __global btAabbCL* aabbs,\n"
" __global Contact4* restrict globalContactsOut,\n"
" __global b3Contact4Data* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
" int numConcavePairs, int maxContactCapacity\n"
" )\n"