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

implement clipFacesAndFindContactsCPU, to track down OpenCL issue on Mac OSX

Browse Source
This commit is contained in:
erwin coumans
2013-12-13 10:05:07 -08:00
parent 33889a3dc9
commit 571b14787a
4 changed files with 225 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

176
src/Bullet3Collision/NarrowPhaseCollision/shared/b3ClipFaces.h Normal file
View File

@@ -0,0 +1,176 @@
#ifndef B3_CLIP_FACES_H
#define B3_CLIP_FACES_H
#include "Bullet3Common/shared/b3Int4.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3RigidBodyData.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Collidable.h"
#include "Bullet3Collision/BroadPhaseCollision/shared/b3Aabb.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3BvhSubtreeInfoData.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3QuantizedBvhNodeData.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3ConvexPolyhedronData.h"
inline b3Float4 b3Lerp3(b3Float4ConstArg a,b3Float4ConstArg b, float t)
{
return b3MakeFloat4( a.x + (b.x - a.x) * t,
a.y + (b.y - a.y) * t,
a.z + (b.z - a.z) * t,
0.f);
}
// Clips a face to the back of a plane, return the number of vertices out, stored in ppVtxOut
int clipFaceGlobal(__global const b3Float4* pVtxIn, int numVertsIn, b3Float4ConstArg planeNormalWS,float planeEqWS, __global b3Float4* ppVtxOut)
{
int ve;
float ds, de;
int numVertsOut = 0;
//double-check next test
// if (numVertsIn < 2)
// return 0;
b3Float4 firstVertex=pVtxIn[numVertsIn-1];
b3Float4 endVertex = pVtxIn[0];
ds = b3Dot(planeNormalWS,firstVertex)+planeEqWS;
for (ve = 0; ve < numVertsIn; ve++)
{
endVertex=pVtxIn[ve];
de = b3Dot(planeNormalWS,endVertex)+planeEqWS;
if (ds<0)
{
if (de<0)
{
// Start < 0, end < 0, so output endVertex
ppVtxOut[numVertsOut++] = endVertex;
}
else
{
// Start < 0, end >= 0, so output intersection
ppVtxOut[numVertsOut++] = b3Lerp3(firstVertex, endVertex,(ds * 1.f/(ds - de)) );
}
}
else
{
if (de<0)
{
// Start >= 0, end < 0 so output intersection and end
ppVtxOut[numVertsOut++] = b3Lerp3(firstVertex, endVertex,(ds * 1.f/(ds - de)) );
ppVtxOut[numVertsOut++] = endVertex;
}
}
firstVertex = endVertex;
ds = de;
}
return numVertsOut;
}
__kernel void clipFacesAndFindContactsKernel( __global const b3Float4* separatingNormals,
__global const int* hasSeparatingAxis,
__global b3Int4* clippingFacesOut,
__global b3Float4* worldVertsA1,
__global b3Float4* worldNormalsA1,
__global b3Float4* worldVertsB1,
__global b3Float4* worldVertsB2,
int vertexFaceCapacity,
int pairIndex
)
{
// int i = get_global_id(0);
//int pairIndex = i;
int i = pairIndex;
float minDist = -1e30f;
float maxDist = 0.02f;
// if (i<numPairs)
{
if (hasSeparatingAxis[i])
{
// int bodyIndexA = pairs[i].x;
// int bodyIndexB = pairs[i].y;
int numLocalContactsOut = 0;
int capacityWorldVertsB2 = vertexFaceCapacity;
__global b3Float4* pVtxIn = &worldVertsB1[pairIndex*capacityWorldVertsB2];
__global b3Float4* pVtxOut = &worldVertsB2[pairIndex*capacityWorldVertsB2];
{
__global b3Int4* clippingFaces = clippingFacesOut;
int closestFaceA = clippingFaces[pairIndex].x;
int closestFaceB = clippingFaces[pairIndex].y;
int numVertsInA = clippingFaces[pairIndex].z;
int numVertsInB = clippingFaces[pairIndex].w;
int numVertsOut = 0;
if (closestFaceA>=0)
{
// clip polygon to back of planes of all faces of hull A that are adjacent to witness face
for(int e0=0;e0<numVertsInA;e0++)
{
const b3Float4 aw = worldVertsA1[pairIndex*capacityWorldVertsB2+e0];
const b3Float4 bw = worldVertsA1[pairIndex*capacityWorldVertsB2+((e0+1)%numVertsInA)];
const b3Float4 WorldEdge0 = aw - bw;
b3Float4 worldPlaneAnormal1 = worldNormalsA1[pairIndex];
b3Float4 planeNormalWS1 = -b3Cross(WorldEdge0,worldPlaneAnormal1);
b3Float4 worldA1 = aw;
float planeEqWS1 = -b3Dot(worldA1,planeNormalWS1);
b3Float4 planeNormalWS = planeNormalWS1;
float planeEqWS=planeEqWS1;
numVertsOut = clipFaceGlobal(pVtxIn, numVertsInB, planeNormalWS,planeEqWS, pVtxOut);
__global b3Float4* tmp = pVtxOut;
pVtxOut = pVtxIn;
pVtxIn = tmp;
numVertsInB = numVertsOut;
numVertsOut = 0;
}
b3Float4 planeNormalWS = worldNormalsA1[pairIndex];
float planeEqWS=-b3Dot(planeNormalWS,worldVertsA1[pairIndex*capacityWorldVertsB2]);
for (int i=0;i<numVertsInB;i++)
{
float depth = b3Dot(planeNormalWS,pVtxIn[i])+planeEqWS;
if (depth <=minDist)
{
depth = minDist;
}
if (depth <=maxDist)
{
b3Float4 pointInWorld = pVtxIn[i];
pVtxOut[numLocalContactsOut++] = b3MakeFloat4(pointInWorld.x,pointInWorld.y,pointInWorld.z,depth);
}
}
}
clippingFaces[pairIndex].w =numLocalContactsOut;
}
for (int i=0;i<numLocalContactsOut;i++)
pVtxIn[i] = pVtxOut[i];
}// if (hasSeparatingAxis[i])
}// if (i<numPairs)
}
#endif //B3_CLIP_FACES_H

56
src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp
View File

@@ -17,6 +17,9 @@ bool findSeparatingAxisOnGpu = true;
bool bvhTraversalKernelGPU = true;
bool findConcaveSeparatingAxisKernelGPU = true;
bool clipFacesAndFindContactsCPU = false;
///This file was written by Erwin Coumans
///Separating axis rest based on work from Pierre Terdiman, see
@@ -71,6 +74,7 @@ typedef b3AlignedObjectArray<b3Vector3> b3VertexArray;
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3BvhTraversal.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3FindConcaveSatAxis.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3ClipFaces.h"
@@ -3579,9 +3583,50 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
//clipFacesAndFindContacts
bool clipFacesAndFindContactsCPU = false;
if (clipFacesAndFindContactsCPU)
{
b3AlignedObjectArray<b3Int4> clippingFacesOutCPU;
b3AlignedObjectArray<b3Vector3> worldVertsA1CPU;
b3AlignedObjectArray<b3Vector3> worldNormalsACPU;
b3AlignedObjectArray<b3Vector3> worldVertsB1CPU;
clippingFacesOutGPU.copyToHost(clippingFacesOutCPU);
worldVertsA1GPU.copyToHost(worldVertsA1CPU);
worldNormalsAGPU.copyToHost(worldNormalsACPU);
worldVertsB1GPU.copyToHost(worldVertsB1CPU);
b3AlignedObjectArray<int>concaveHasSeparatingNormalsCPU;
m_concaveHasSeparatingNormals.copyToHost(concaveHasSeparatingNormalsCPU);
b3AlignedObjectArray<b3Vector3> concaveSepNormalsHost;
m_concaveSepNormals.copyToHost(concaveSepNormalsHost);
b3AlignedObjectArray<b3Vector3> worldVertsB2CPU;
worldVertsB2CPU.resize(worldVertsB2GPU.size());
for (int i=0;i<numConcavePairs;i++)
{
clipFacesAndFindContactsKernel( &concaveSepNormalsHost.at(0),
&concaveHasSeparatingNormalsCPU.at(0),
&clippingFacesOutCPU.at(0),
&worldVertsA1CPU.at(0),
&worldNormalsACPU.at(0),
&worldVertsB1CPU.at(0),
&worldVertsB2CPU.at(0),
vertexFaceCapacity,
i);
}
clippingFacesOutGPU.copyFromHost(clippingFacesOutCPU);
worldVertsB2GPU.copyFromHost(worldVertsB2CPU);
} else
{
@@ -3589,6 +3634,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
{
B3_PROFILE("clipFacesAndFindContacts");
//nContacts = m_totalContactsOut.at(0);
//int h = m_hasSeparatingNormals.at(0);
@@ -3596,13 +3642,11 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
b3BufferInfoCL bInfo[] = {
b3BufferInfoCL( m_concaveSepNormals.getBufferCL()),
b3BufferInfoCL( m_concaveHasSeparatingNormals.getBufferCL()),
b3BufferInfoCL( contactOut->getBufferCL()),
b3BufferInfoCL( clippingFacesOutGPU.getBufferCL()),
b3BufferInfoCL( worldVertsA1GPU.getBufferCL()),
b3BufferInfoCL( worldNormalsAGPU.getBufferCL()),
b3BufferInfoCL( worldVertsB1GPU.getBufferCL()),
b3BufferInfoCL( worldVertsB2GPU.getBufferCL()),
b3BufferInfoCL( m_totalContactsOut.getBufferCL())
b3BufferInfoCL( worldVertsB2GPU.getBufferCL())
};
b3LauncherCL launcher(m_queue, m_clipFacesAndFindContacts,"m_clipFacesAndFindContacts");
launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
@@ -3749,13 +3793,11 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
b3BufferInfoCL bInfo[] = {
b3BufferInfoCL( m_sepNormals.getBufferCL()),
b3BufferInfoCL( m_hasSeparatingNormals.getBufferCL()),
b3BufferInfoCL( contactOut->getBufferCL()),
b3BufferInfoCL( clippingFacesOutGPU.getBufferCL()),
b3BufferInfoCL( worldVertsA1GPU.getBufferCL()),
b3BufferInfoCL( worldNormalsAGPU.getBufferCL()),
b3BufferInfoCL( worldVertsB1GPU.getBufferCL()),
b3BufferInfoCL( worldVertsB2GPU.getBufferCL()),
b3BufferInfoCL( m_totalContactsOut.getBufferCL())
b3BufferInfoCL( worldVertsB2GPU.getBufferCL())
};
b3LauncherCL launcher(m_queue, m_clipFacesAndFindContacts,"m_clipFacesAndFindContacts");

2
src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl
View File

@@ -1671,13 +1671,11 @@ __kernel void findClippingFacesKernel( __global const int4* pairs,
__kernel void clipFacesAndFindContactsKernel( __global const float4* separatingNormals,
__global const int* hasSeparatingAxis,
__global struct b3Contact4Data* globalContactsOut,
__global int4* clippingFacesOut,
__global float4* worldVertsA1,
__global float4* worldNormalsA1,
__global float4* worldVertsB1,
__global float4* worldVertsB2,
counter32_t nGlobalContactsOut,
int vertexFaceCapacity,
int numPairs,
int debugMode

2
src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h
View File

@@ -1865,13 +1865,11 @@ static const char* satClipKernelsCL= \
"}\n"
"__kernel void clipFacesAndFindContactsKernel( __global const float4* separatingNormals,\n"
" __global const int* hasSeparatingAxis,\n"
" __global struct b3Contact4Data* globalContactsOut,\n"
" __global int4* clippingFacesOut,\n"
" __global float4* worldVertsA1,\n"
" __global float4* worldNormalsA1,\n"
" __global float4* worldVertsB1,\n"
" __global float4* worldVertsB2,\n"
" counter32_t nGlobalContactsOut,\n"
" int vertexFaceCapacity,\n"
" int numPairs,\n"
" int debugMode\n"