From 7992ff816b09945762d4357eceb3a99cb648768c Mon Sep 17 00:00:00 2001 From: erwincoumans Date: Wed, 31 Jul 2013 09:58:15 -0700 Subject: [PATCH] use int4 for broadphase pair, it allows to store persistent information in the future (contact cache, applied impulse/warm starting info etc) --- Demos3/GpuDemos/broadphase/PairBench.cpp | 6 +-- .../BroadPhaseCollision/b3OverlappingPair.h | 4 +- .../b3GpuSapBroadphase.cpp | 45 +++++++++---------- .../BroadphaseCollision/b3GpuSapBroadphase.h | 6 +-- .../BroadphaseCollision/kernels/sap.cl | 30 ++++++++----- .../BroadphaseCollision/kernels/sapFast.cl | 26 +++++++---- .../kernels/sapFastKernels.h | 26 +++++++---- .../BroadphaseCollision/kernels/sapKernels.h | 30 ++++++++----- .../b3ConvexHullContact.cpp | 4 +- .../b3ConvexHullContact.h | 2 +- .../kernels/bvhTraversal.cl | 2 +- .../kernels/bvhTraversal.h | 2 +- .../kernels/primitiveContacts.cl | 2 +- .../kernels/primitiveContacts.h | 2 +- .../NarrowphaseCollision/kernels/sat.cl | 4 +- .../kernels/satClipHullContacts.cl | 12 ++--- .../kernels/satClipHullContacts.h | 12 ++--- .../NarrowphaseCollision/kernels/satKernels.h | 4 +- .../RigidBody/b3GpuNarrowPhase.cpp | 2 +- 19 files changed, 127 insertions(+), 94 deletions(-) diff --git a/Demos3/GpuDemos/broadphase/PairBench.cpp b/Demos3/GpuDemos/broadphase/PairBench.cpp index 980316b3e..a250b74d6 100644 --- a/Demos3/GpuDemos/broadphase/PairBench.cpp +++ b/Demos3/GpuDemos/broadphase/PairBench.cpp @@ -31,7 +31,7 @@ __kernel void moveObjectsKernel(__global float4* posOrnColors, int numObjects) colors[iGID] = (float4)(0,0,1,1); } -__kernel void colorPairsKernel(__global float4* posOrnColors, int numObjects, __global const int2* pairs, int numPairs) +__kernel void colorPairsKernel(__global float4* posOrnColors, int numObjects, __global const int4* pairs, int numPairs) { int iPairId = get_global_id(0); if (iPairId>=numPairs) @@ -48,8 +48,8 @@ __kernel void sineWaveKernel( __global float4* posOrnColors, __global float* pBodyTimes,const int numNodes) { int nodeID = get_global_id(0); - float timeStepPos = 0.00166666; - float mAmplitude = 36.f; + float timeStepPos = 0.000166666; + float mAmplitude = 86.f; if( nodeID < numNodes ) { pBodyTimes[nodeID] += timeStepPos; diff --git a/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPair.h b/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPair.h index 596705c1a..4ec51ded1 100644 --- a/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPair.h +++ b/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPair.h @@ -16,10 +16,10 @@ subject to the following restrictions: #ifndef B3_OVERLAPPING_PAIR_H #define B3_OVERLAPPING_PAIR_H -#include "Bullet3Common/b3Int2.h" +#include "Bullet3Common/b3Int4.h" //typedef b3Int2 b3BroadphasePair; -struct b3BroadphasePair : public b3Int2 +struct b3BroadphasePair : public b3Int4 { explicit b3BroadphasePair(){} b3BroadphasePair(int xx,int yy) diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp index c0fce3247..f8a99b5ae 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp +++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp @@ -192,29 +192,29 @@ void b3GpuSapBroadphase::init3dSap() } -static bool b3PairCmp(const b3Int2& p, const b3Int2& q) +static bool b3PairCmp(const b3Int4& p, const b3Int4& q) { return ((p.x(const b3Int2& a,const b3Int2& b) +static bool operator>(const b3Int4& a,const b3Int4& b) { return a.x > b.x || (a.x == b.x && a.y > b.y); }; -b3AlignedObjectArray addedHostPairs; -b3AlignedObjectArray removedHostPairs; +b3AlignedObjectArray addedHostPairs; +b3AlignedObjectArray removedHostPairs; b3AlignedObjectArray preAabbs; @@ -247,7 +247,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap() m_allAabbsGPU.copyToHost(m_allAabbsCPU); } - b3AlignedObjectArray allPairs; + b3AlignedObjectArray allPairs; { B3_PROFILE("m_overlappingPairs.copyToHost"); m_overlappingPairs.copyToHost(allPairs); @@ -268,7 +268,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap() { - b3Int2 newPair; + b3Int4 newPair; newPair.x = 40; newPair.y = 53; int index = allPairs.findBinarySearch(newPair); @@ -587,7 +587,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap() if (overlap && !prevOverlap) { //add a pair - b3Int2 newPair; + b3Int4 newPair; if (i<=otherIndex) { newPair.x = i; @@ -606,7 +606,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap() { //remove a pair - b3Int2 removedPair; + b3Int4 removedPair; if (i<=otherIndex) { removedPair.x = i; @@ -664,7 +664,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap() if (overlap && !prevOverlap) { //add a pair - b3Int2 newPair; + b3Int4 newPair; if (i<=otherIndex) { newPair.x = i; @@ -684,7 +684,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap() { //if (otherIndex2&1==0) -> min? //remove a pair - b3Int2 removedPair; + b3Int4 removedPair; if (i<=otherIndex) { removedPair.x = i; @@ -727,7 +727,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap() } } - b3Int2 prevPair; + b3Int4 prevPair; prevPair.x = -1; prevPair.y = -1; @@ -739,7 +739,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap() B3_PROFILE("actual removing"); for (int i=0;i actualAddedPairs; + b3AlignedObjectArray actualAddedPairs; { B3_PROFILE("actual adding"); for (int i=0;i addedHostPairs; - //b3AlignedObjectArray removedHostPairs; { B3_PROFILE("m_overlappingPairs.copyFromHost"); m_overlappingPairs.copyFromHost(allPairs); @@ -905,7 +902,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHost(int maxPairs) } } - b3AlignedObjectArray hostPairs; + b3AlignedObjectArray hostPairs; { int numSmallAabbs = m_smallAabbsCPU.size(); @@ -918,7 +915,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHost(int maxPairs) if (TestAabbAgainstAabb2((b3Vector3&)m_smallAabbsCPU[i].m_min, (b3Vector3&)m_smallAabbsCPU[i].m_max, (b3Vector3&)m_smallAabbsCPU[j].m_min,(b3Vector3&)m_smallAabbsCPU[j].m_max)) { - b3Int2 pair; + b3Int4 pair; int a = m_smallAabbsCPU[i].m_minIndices[3]; int b = m_smallAabbsCPU[j].m_minIndices[3]; if (a<=b) @@ -949,7 +946,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHost(int maxPairs) if (TestAabbAgainstAabb2((b3Vector3&)m_smallAabbsCPU[i].m_min, (b3Vector3&)m_smallAabbsCPU[i].m_max, (b3Vector3&)m_largeAabbsCPU[j].m_min,(b3Vector3&)m_largeAabbsCPU[j].m_max)) { - b3Int2 pair; + b3Int4 pair; int a = m_largeAabbsCPU[j].m_minIndices[3]; int b = m_smallAabbsCPU[i].m_minIndices[3]; if (a<=b) @@ -1276,8 +1273,8 @@ void b3GpuSapBroadphase::calculateOverlappingPairs(int maxPairs) pairCount.setFromOpenCLBuffer(launcher.m_arrays[2]->getBufferCL(),numElements); numPairs = pairCount.at(0); //printf("overlapping pairs = %d\n",numPairs); - b3AlignedObjectArray hostOoverlappingPairs; - b3OpenCLArray tmpGpuPairs(m_context,m_queue); + b3AlignedObjectArray hostOoverlappingPairs; + b3OpenCLArray tmpGpuPairs(m_context,m_queue); tmpGpuPairs.setFromOpenCLBuffer(launcher.m_arrays[1]->getBufferCL(),numPairs ); tmpGpuPairs.copyToHost(hostOoverlappingPairs); diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h index 92c8d6169..be6a8740a 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h +++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h @@ -44,8 +44,8 @@ class b3GpuSapBroadphase b3OpenCLArray m_sortedAxisGPU2prev; - b3OpenCLArray m_addedHostPairsGPU; - b3OpenCLArray m_removedHostPairsGPU; + b3OpenCLArray m_addedHostPairsGPU; + b3OpenCLArray m_removedHostPairsGPU; b3OpenCLArray m_addedCountGPU; b3OpenCLArray m_removedCountGPU; @@ -68,7 +68,7 @@ class b3GpuSapBroadphase b3OpenCLArray m_largeAabbsGPU; b3AlignedObjectArray m_largeAabbsCPU; - b3OpenCLArray m_overlappingPairs; + b3OpenCLArray m_overlappingPairs; //temporary gpu work memory b3OpenCLArray m_gpuSmallSortData; diff --git a/src/Bullet3OpenCL/BroadphaseCollision/kernels/sap.cl b/src/Bullet3OpenCL/BroadphaseCollision/kernels/sap.cl index 4d94f9cdf..f2f273663 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/kernels/sap.cl +++ b/src/Bullet3OpenCL/BroadphaseCollision/kernels/sap.cl @@ -62,7 +62,7 @@ bool TestAabbAgainstAabb2Global(const btAabbCL* aabb1, __global const btAabbCL* } -__kernel void computePairsKernelTwoArrays( __global const btAabbCL* unsortedAabbs, __global const btAabbCL* sortedAabbs, volatile __global int2* pairsOut,volatile __global int* pairCount, int numUnsortedAabbs, int numSortedAabbs, int axis, int maxPairs) +__kernel void computePairsKernelTwoArrays( __global const btAabbCL* unsortedAabbs, __global const btAabbCL* sortedAabbs, volatile __global int4* pairsOut,volatile __global int* pairCount, int numUnsortedAabbs, int numSortedAabbs, int axis, int maxPairs) { int i = get_global_id(0); if (i>=numUnsortedAabbs) @@ -74,10 +74,19 @@ __kernel void computePairsKernelTwoArrays( __global const btAabbCL* unsortedAa if (TestAabbAgainstAabb2GlobalGlobal(&unsortedAabbs[i],&sortedAabbs[j])) { - int2 myPair; + int4 myPair; - myPair.x = unsortedAabbs[i].m_minIndices[3]; - myPair.y = sortedAabbs[j].m_minIndices[3]; + int xIndex = unsortedAabbs[i].m_minIndices[3]; + int yIndex = sortedAabbs[j].m_minIndices[3]; + if (xIndex>yIndex) + { + int tmp = xIndex; + xIndex=yIndex; + yIndex=tmp; + } + + myPair.x = xIndex; + myPair.y = yIndex; int curPair = atomic_inc (pairCount); if (curPair=numObjects) @@ -100,9 +109,10 @@ __kernel void computePairsKernelOriginal( __global const btAabbCL* aabbs, vola } if (TestAabbAgainstAabb2GlobalGlobal(&aabbs[i],&aabbs[j])) { - int2 myPair; + int4 myPair; myPair.x = aabbs[i].m_minIndices[3]; myPair.y = aabbs[j].m_minIndices[3]; + int curPair = atomic_inc (pairCount); if (curPair min? //remove a pair - int2 removedPair; + int4 removedPair; if (i<=otherIndex) { removedPair.x = i; @@ -318,7 +318,7 @@ __kernel void computePairsIncremental3dSapKernel( __global const uint2* object } //computePairsKernelBatchWrite -__kernel void computePairsKernel( __global const btAabbCL* aabbs, volatile __global int2* pairsOut,volatile __global int* pairCount, int numObjects, int axis, int maxPairs) +__kernel void computePairsKernel( __global const btAabbCL* aabbs, volatile __global int4* pairsOut,volatile __global int* pairCount, int numObjects, int axis, int maxPairs) { int i = get_global_id(0); int localId = get_local_id(0); @@ -393,7 +393,11 @@ __kernel void computePairsKernel( __global const btAabbCL* aabbs, volatile __g { for (int p=0;p min?\n" " //remove a pair\n" -" int2 removedPair;\n" +" int4 removedPair;\n" " if (i<=otherIndex)\n" " {\n" " removedPair.x = i;\n" @@ -320,7 +320,7 @@ static const char* sapFastCL= \ "}\n" "\n" "//computePairsKernelBatchWrite\n" -"__kernel void computePairsKernel( __global const btAabbCL* aabbs, volatile __global int2* pairsOut,volatile __global int* pairCount, int numObjects, int axis, int maxPairs)\n" +"__kernel void computePairsKernel( __global const btAabbCL* aabbs, volatile __global int4* pairsOut,volatile __global int* pairCount, int numObjects, int axis, int maxPairs)\n" "{\n" " int i = get_global_id(0);\n" " int localId = get_local_id(0);\n" @@ -395,7 +395,11 @@ static const char* sapFastCL= \ " {\n" " for (int p=0;p=numUnsortedAabbs)\n" @@ -76,10 +76,19 @@ static const char* sapCL= \ "\n" " if (TestAabbAgainstAabb2GlobalGlobal(&unsortedAabbs[i],&sortedAabbs[j]))\n" " {\n" -" int2 myPair;\n" +" int4 myPair;\n" " \n" -" myPair.x = unsortedAabbs[i].m_minIndices[3];\n" -" myPair.y = sortedAabbs[j].m_minIndices[3];\n" +" int xIndex = unsortedAabbs[i].m_minIndices[3];\n" +" int yIndex = sortedAabbs[j].m_minIndices[3];\n" +" if (xIndex>yIndex)\n" +" {\n" +" int tmp = xIndex;\n" +" xIndex=yIndex;\n" +" yIndex=tmp;\n" +" }\n" +" \n" +" myPair.x = xIndex;\n" +" myPair.y = yIndex;\n" "\n" " int curPair = atomic_inc (pairCount);\n" " if (curPair=numObjects)\n" @@ -102,9 +111,10 @@ static const char* sapCL= \ " }\n" " if (TestAabbAgainstAabb2GlobalGlobal(&aabbs[i],&aabbs[j]))\n" " {\n" -" int2 myPair;\n" +" int4 myPair;\n" " myPair.x = aabbs[i].m_minIndices[3];\n" " myPair.y = aabbs[j].m_minIndices[3];\n" +" \n" " int curPair = atomic_inc (pairCount);\n" " if (curPair* pairs, int nPairs, +void GpuSatCollision::computeConvexConvexContactsGPUSAT( const b3OpenCLArray* pairs, int nPairs, const b3OpenCLArray* bodyBuf, b3OpenCLArray* contactOut, int& nContacts, int maxContactCapacity, @@ -1719,7 +1719,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const b3OpenCLArray hostAabbs; clAabbsWS.copyToHost(hostAabbs); - b3AlignedObjectArray hostPairs; + b3AlignedObjectArray hostPairs; pairs->copyToHost(hostPairs); b3AlignedObjectArray hostBodyBuf; diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h index b138931ff..326acec66 100644 --- a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h +++ b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h @@ -75,7 +75,7 @@ struct GpuSatCollision virtual ~GpuSatCollision(); - void computeConvexConvexContactsGPUSAT( const b3OpenCLArray* pairs, int nPairs, + void computeConvexConvexContactsGPUSAT( const b3OpenCLArray* pairs, int nPairs, const b3OpenCLArray* bodyBuf, b3OpenCLArray* contactOut, int& nContacts, int maxContactCapacity, diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.cl b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.cl index db6257f15..adc2b5d8c 100644 --- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.cl +++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.cl @@ -192,7 +192,7 @@ void quantizeWithClamp(unsigned short* out, float4 point2,int isMax, float4 bvhA // work-in-progress -__kernel void bvhTraversalKernel( __global const int2* pairs, +__kernel void bvhTraversalKernel( __global const int4* pairs, __global const BodyData* rigidBodies, __global const btCollidableGpu* collidables, __global btAabbCL* aabbs, diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.h b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.h index 6640917cb..23f4865d8 100644 --- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.h +++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.h @@ -194,7 +194,7 @@ static const char* bvhTraversalKernelCL= \ "\n" "\n" "// work-in-progress\n" -"__kernel void bvhTraversalKernel( __global const int2* pairs, \n" +"__kernel void bvhTraversalKernel( __global const int4* pairs, \n" " __global const BodyData* rigidBodies, \n" " __global const btCollidableGpu* collidables,\n" " __global btAabbCL* aabbs,\n" diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl index 347a2a67a..1f37e3cbf 100644 --- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl +++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl @@ -802,7 +802,7 @@ void computeContactPlaneSphere(int pairIndex, } -__kernel void primitiveContactsKernel( __global const int2* pairs, +__kernel void primitiveContactsKernel( __global const int4* pairs, __global const BodyData* rigidBodies, __global const btCollidableGpu* collidables, __global const ConvexPolyhedronCL* convexShapes, diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h index e0f1d1be3..c5bdbc97a 100644 --- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h +++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h @@ -804,7 +804,7 @@ static const char* primitiveContactsKernelsCL= \ "}\n" "\n" "\n" -"__kernel void primitiveContactsKernel( __global const int2* pairs, \n" +"__kernel void primitiveContactsKernel( __global const int4* pairs, \n" " __global const BodyData* rigidBodies, \n" " __global const btCollidableGpu* collidables,\n" " __global const ConvexPolyhedronCL* convexShapes, \n" diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/sat.cl b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/sat.cl index e8e9e68a1..638d7399f 100644 --- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/sat.cl +++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/sat.cl @@ -756,7 +756,7 @@ __kernel void processCompoundPairsKernel( __global const int4* gpuCompoundPair } // work-in-progress -__kernel void findCompoundPairsKernel( __global const int2* pairs, +__kernel void findCompoundPairsKernel( __global const int4* pairs, __global const BodyData* rigidBodies, __global const btCollidableGpu* collidables, __global const ConvexPolyhedronCL* convexShapes, @@ -938,7 +938,7 @@ __kernel void findCompoundPairsKernel( __global const int2* pairs, } // work-in-progress -__kernel void findSeparatingAxisKernel( __global const int2* pairs, +__kernel void findSeparatingAxisKernel( __global const int4* pairs, __global const BodyData* rigidBodies, __global const btCollidableGpu* collidables, __global const ConvexPolyhedronCL* convexShapes, diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl index 42cde9b33..e5f8c85b4 100644 --- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl +++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl @@ -885,7 +885,7 @@ int extractManifoldSequential(const float4* p, int nPoints, float4 nearNormal, i -__kernel void extractManifoldAndAddContactKernel(__global const int2* pairs, +__kernel void extractManifoldAndAddContactKernel(__global const int4* pairs, __global const BodyData* rigidBodies, __global const float4* closestPointsWorld, __global const float4* separatingNormalsWorld, @@ -960,7 +960,7 @@ void trMul(float4 translationA, Quaternion orientationA, -__kernel void clipHullHullKernel( __global const int2* pairs, +__kernel void clipHullHullKernel( __global const int4* pairs, __global const BodyData* rigidBodies, __global const btCollidableGpu* collidables, __global const ConvexPolyhedronCL* convexShapes, @@ -1192,7 +1192,7 @@ __kernel void clipCompoundsHullHullKernel( __global const int4* gpuCompoundPai -__kernel void sphereSphereCollisionKernel( __global const int2* pairs, +__kernel void sphereSphereCollisionKernel( __global const int4* pairs, __global const BodyData* rigidBodies, __global const btCollidableGpu* collidables, __global const float4* separatingNormals, @@ -1677,7 +1677,7 @@ int clipFaces(__global float4* worldVertsA1, -__kernel void findClippingFacesKernel( __global const int2* pairs, +__kernel void findClippingFacesKernel( __global const int4* pairs, __global const BodyData* rigidBodies, __global const btCollidableGpu* collidables, __global const ConvexPolyhedronCL* convexShapes, @@ -1740,7 +1740,7 @@ __kernel void findClippingFacesKernel( __global const int2* pairs, -__kernel void clipFacesAndContactReductionKernel( __global const int2* pairs, +__kernel void clipFacesAndContactReductionKernel( __global const int4* pairs, __global const BodyData* rigidBodies, __global const float4* separatingNormals, __global const int* hasSeparatingAxis, @@ -1853,7 +1853,7 @@ __kernel void clipFacesAndContactReductionKernel( __global const int2* pairs, -__kernel void newContactReductionKernel( __global const int2* pairs, +__kernel void newContactReductionKernel( __global const int4* pairs, __global const BodyData* rigidBodies, __global const float4* separatingNormals, __global const int* hasSeparatingAxis, diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h index a4b654925..5496cc90e 100644 --- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h +++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h @@ -887,7 +887,7 @@ static const char* satClipKernelsCL= \ "\n" "\n" "\n" -"__kernel void extractManifoldAndAddContactKernel(__global const int2* pairs, \n" +"__kernel void extractManifoldAndAddContactKernel(__global const int4* pairs, \n" " __global const BodyData* rigidBodies, \n" " __global const float4* closestPointsWorld,\n" " __global const float4* separatingNormalsWorld,\n" @@ -962,7 +962,7 @@ static const char* satClipKernelsCL= \ "\n" "\n" "\n" -"__kernel void clipHullHullKernel( __global const int2* pairs, \n" +"__kernel void clipHullHullKernel( __global const int4* pairs, \n" " __global const BodyData* rigidBodies, \n" " __global const btCollidableGpu* collidables,\n" " __global const ConvexPolyhedronCL* convexShapes, \n" @@ -1194,7 +1194,7 @@ static const char* satClipKernelsCL= \ "\n" "\n" "\n" -"__kernel void sphereSphereCollisionKernel( __global const int2* pairs, \n" +"__kernel void sphereSphereCollisionKernel( __global const int4* pairs, \n" " __global const BodyData* rigidBodies, \n" " __global const btCollidableGpu* collidables,\n" " __global const float4* separatingNormals,\n" @@ -1679,7 +1679,7 @@ static const char* satClipKernelsCL= \ "\n" "\n" "\n" -"__kernel void findClippingFacesKernel( __global const int2* pairs,\n" +"__kernel void findClippingFacesKernel( __global const int4* pairs,\n" " __global const BodyData* rigidBodies,\n" " __global const btCollidableGpu* collidables,\n" " __global const ConvexPolyhedronCL* convexShapes,\n" @@ -1742,7 +1742,7 @@ static const char* satClipKernelsCL= \ "\n" "\n" "\n" -"__kernel void clipFacesAndContactReductionKernel( __global const int2* pairs,\n" +"__kernel void clipFacesAndContactReductionKernel( __global const int4* pairs,\n" " __global const BodyData* rigidBodies,\n" " __global const float4* separatingNormals,\n" " __global const int* hasSeparatingAxis,\n" @@ -1855,7 +1855,7 @@ static const char* satClipKernelsCL= \ "\n" "\n" "\n" -"__kernel void newContactReductionKernel( __global const int2* pairs,\n" +"__kernel void newContactReductionKernel( __global const int4* pairs,\n" " __global const BodyData* rigidBodies,\n" " __global const float4* separatingNormals,\n" " __global const int* hasSeparatingAxis,\n" diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h index 5e9f8b4f6..25f8fc49b 100644 --- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h +++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h @@ -758,7 +758,7 @@ static const char* satKernelsCL= \ "}\n" "\n" "// work-in-progress\n" -"__kernel void findCompoundPairsKernel( __global const int2* pairs, \n" +"__kernel void findCompoundPairsKernel( __global const int4* pairs, \n" " __global const BodyData* rigidBodies, \n" " __global const btCollidableGpu* collidables,\n" " __global const ConvexPolyhedronCL* convexShapes, \n" @@ -940,7 +940,7 @@ static const char* satKernelsCL= \ "}\n" "\n" "// work-in-progress\n" -"__kernel void findSeparatingAxisKernel( __global const int2* pairs, \n" +"__kernel void findSeparatingAxisKernel( __global const int4* pairs, \n" " __global const BodyData* rigidBodies, \n" " __global const btCollidableGpu* collidables,\n" " __global const ConvexPolyhedronCL* convexShapes, \n" diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.cpp b/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.cpp index b60828548..27926d221 100644 --- a/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.cpp +++ b/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.cpp @@ -727,7 +727,7 @@ void b3GpuNarrowPhase::computeContacts(cl_mem broadphasePairs, int numBroadphase int maxTriConvexPairCapacity = m_data->m_config.m_maxTriConvexPairCapacity; int numTriConvexPairsOut=0; - b3OpenCLArray broadphasePairsGPU(m_context,m_queue); + b3OpenCLArray broadphasePairsGPU(m_context,m_queue); broadphasePairsGPU.setFromOpenCLBuffer(broadphasePairs,numBroadphasePairs); b3OpenCLArray clAabbArray(this->m_context,this->m_queue); clAabbArray.setFromOpenCLBuffer(aabbsWS,numObjects);