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

fix many warnings

Browse Source 
remove btMultiSapBroadphase.*
make collisionFilterGroup/collisionFilterMark int (instead of short int)
This commit is contained in:
Erwin Coumans
2017-01-15 22:26:11 -08:00
parent e3df00d5f1
commit c0c4c8ba3f
146 changed files with 830 additions and 1431 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h
View File

@@ -19,8 +19,8 @@ public:
{
}
virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask)=0;
virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask)=0;
virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr , int collisionFilterGroup, int collisionFilterMask)=0;
virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr , int collisionFilterGroup, int collisionFilterMask)=0;
virtual void calculateOverlappingPairs(int maxPairs)=0;
virtual void calculateOverlappingPairsHost(int maxPairs)=0;

9
src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp
View File

@@ -38,8 +38,9 @@ m_largeAabbsMappingGPU(ctx,q),
m_gpuPairs(ctx,q),
m_hashGpu(ctx,q),
m_paramsGPU(ctx,q),
m_cellStartGpu(ctx,q)
m_cellStartGpu(ctx,q),
m_paramsGPU(ctx,q)
{
@@ -110,7 +111,7 @@ b3GpuGridBroadphase::~b3GpuGridBroadphase()
void b3GpuGridBroadphase::createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask)
void b3GpuGridBroadphase::createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr , int collisionFilterGroup, int collisionFilterMask)
{
b3SapAabb aabb;
aabb.m_minVec = aabbMin;
@@ -122,7 +123,7 @@ void b3GpuGridBroadphase::createProxy(const b3Vector3& aabbMin, const b3Vector3
m_allAabbsCPU1.push_back(aabb);
}
void b3GpuGridBroadphase::createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask)
void b3GpuGridBroadphase::createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr , int collisionFilterGroup, int collisionFilterMask)
{
b3SapAabb aabb;
aabb.m_minVec = aabbMin;

4
src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h
View File

@@ -63,8 +63,8 @@ public:
virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask);
virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask);
virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr , int collisionFilterGroup, int collisionFilterMask);
virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr , int collisionFilterGroup, int collisionFilterMask);
virtual void calculateOverlappingPairs(int maxPairs);
virtual void calculateOverlappingPairsHost(int maxPairs);

4
src/Bullet3OpenCL/BroadphaseCollision/b3GpuParallelLinearBvhBroadphase.cpp
View File

@@ -24,7 +24,7 @@ b3GpuParallelLinearBvhBroadphase::b3GpuParallelLinearBvhBroadphase(cl_context co
{
}
void b3GpuParallelLinearBvhBroadphase::createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr, short int collisionFilterGroup, short int collisionFilterMask)
void b3GpuParallelLinearBvhBroadphase::createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr, int collisionFilterGroup, int collisionFilterMask)
{
int newAabbIndex = m_aabbsCpu.size();
@@ -39,7 +39,7 @@ void b3GpuParallelLinearBvhBroadphase::createProxy(const b3Vector3& aabbMin, con
m_aabbsCpu.push_back(aabb);
}
void b3GpuParallelLinearBvhBroadphase::createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr, short int collisionFilterGroup, short int collisionFilterMask)
void b3GpuParallelLinearBvhBroadphase::createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr, int collisionFilterGroup, int collisionFilterMask)
{
int newAabbIndex = m_aabbsCpu.size();

4
src/Bullet3OpenCL/BroadphaseCollision/b3GpuParallelLinearBvhBroadphase.h
View File

@@ -36,8 +36,8 @@ public:
b3GpuParallelLinearBvhBroadphase(cl_context context, cl_device_id device, cl_command_queue queue);
virtual ~b3GpuParallelLinearBvhBroadphase() {}
virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr, short int collisionFilterGroup, short int collisionFilterMask);
virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr, short int collisionFilterGroup, short int collisionFilterMask);
virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr, int collisionFilterGroup, int collisionFilterMask);
virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr, int collisionFilterGroup, int collisionFilterMask);
virtual void calculateOverlappingPairs(int maxPairs);
virtual void calculateOverlappingPairsHost(int maxPairs);

92
src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp
View File

@@ -14,22 +14,53 @@ bool searchIncremental3dSapOnGpu = true;
#define B3_BROADPHASE_SAP_PATH "src/Bullet3OpenCL/BroadphaseCollision/kernels/sap.cl"
/*
b3OpenCLArray<int> m_pairCount;
b3OpenCLArray<b3SapAabb> m_allAabbsGPU;
b3AlignedObjectArray<b3SapAabb> m_allAabbsCPU;
virtual b3OpenCLArray<b3SapAabb>& getAllAabbsGPU()
{
return m_allAabbsGPU;
}
virtual b3AlignedObjectArray<b3SapAabb>& getAllAabbsCPU()
{
return m_allAabbsCPU;
}
b3OpenCLArray<b3Vector3> m_sum;
b3OpenCLArray<b3Vector3> m_sum2;
b3OpenCLArray<b3Vector3> m_dst;
b3OpenCLArray<int> m_smallAabbsMappingGPU;
b3AlignedObjectArray<int> m_smallAabbsMappingCPU;
b3OpenCLArray<int> m_largeAabbsMappingGPU;
b3AlignedObjectArray<int> m_largeAabbsMappingCPU;
b3OpenCLArray<b3Int4> m_overlappingPairs;
//temporary gpu work memory
b3OpenCLArray<b3SortData> m_gpuSmallSortData;
b3OpenCLArray<b3SapAabb> m_gpuSmallSortedAabbs;
class b3PrefixScanFloat4CL* m_prefixScanFloat4;
*/
b3GpuSapBroadphase::b3GpuSapBroadphase(cl_context ctx,cl_device_id device, cl_command_queue q , b3GpuSapKernelType kernelType)
:m_context(ctx),
m_device(device),
m_queue(q),
m_allAabbsGPU(ctx,q),
m_smallAabbsMappingGPU(ctx,q),
m_largeAabbsMappingGPU(ctx,q),
m_pairCount(ctx,q),
m_overlappingPairs(ctx,q),
m_gpuSmallSortData(ctx,q),
m_gpuSmallSortedAabbs(ctx,q),
m_sum(ctx,q),
m_sum2(ctx,q),
m_dst(ctx,q),
m_currentBuffer(-1),
m_objectMinMaxIndexGPUaxis0(ctx,q),
m_objectMinMaxIndexGPUaxis1(ctx,q),
m_objectMinMaxIndexGPUaxis2(ctx,q),
@@ -45,7 +76,18 @@ m_sortedAxisGPU2prev(ctx,q),
m_addedHostPairsGPU(ctx,q),
m_removedHostPairsGPU(ctx,q),
m_addedCountGPU(ctx,q),
m_removedCountGPU(ctx,q)
m_removedCountGPU(ctx,q),
m_currentBuffer(-1),
m_pairCount(ctx,q),
m_allAabbsGPU(ctx,q),
m_sum(ctx,q),
m_sum2(ctx,q),
m_dst(ctx,q),
m_smallAabbsMappingGPU(ctx,q),
m_largeAabbsMappingGPU(ctx,q),
m_overlappingPairs(ctx,q),
m_gpuSmallSortData(ctx,q),
m_gpuSmallSortedAabbs(ctx,q)
{
const char* sapSrc = sapCL;
@@ -191,7 +233,7 @@ void b3GpuSapBroadphase::init3dSap()
for (int axis=0;axis<3;axis++)
{
int totalNumAabbs = m_allAabbsCPU.size();
//int totalNumAabbs = m_allAabbsCPU.size();
int numEndPoints = m_sortedAxisCPU[axis][m_currentBuffer].size();
m_objectMinMaxIndexCPU[axis][m_currentBuffer].resize(numEndPoints);
for (int i=0;i<numEndPoints;i++)
@@ -240,7 +282,7 @@ b3AlignedObjectArray<b3SapAabb> preAabbs;
void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap()
{
static int framepje = 0;
//static int framepje = 0;
//printf("framepje=%d\n",framepje++);
@@ -371,6 +413,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap()
m_sorter->executeHost(m_sortedAxisCPU[axis][m_currentBuffer]);
}
#if 0
if (0)
{
for (int axis=0;axis<3;axis++)
@@ -378,14 +421,14 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap()
//printf("axis %d\n",axis);
for (int i=0;i<m_sortedAxisCPU[axis][m_currentBuffer].size();i++)
{
int key = m_sortedAxisCPU[axis][m_currentBuffer][i].m_key;
int value = m_sortedAxisCPU[axis][m_currentBuffer][i].m_value;
//int key = m_sortedAxisCPU[axis][m_currentBuffer][i].m_key;
//int value = m_sortedAxisCPU[axis][m_currentBuffer][i].m_value;
//printf("[%d]=%d\n",i,value);
}
}
}
#endif
{
B3_PROFILE("assign m_objectMinMaxIndexCPU");
@@ -409,7 +452,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap()
}
}
#if 0
if (0)
{
printf("==========================\n");
@@ -440,6 +483,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap()
}
}
#endif
int a = m_objectMinMaxIndexCPU[0][m_currentBuffer].size();
@@ -654,7 +698,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHostIncremental3Sap()
int otherIndex = otherIndex2/2;
if (otherIndex!=i)
{
bool otherIsMin = ((otherIndex2&1)==0);
//bool otherIsMin = ((otherIndex2&1)==0);
//if (otherIsMin)
{
//bool overlap = TestAabbAgainstAabb2((const b3Vector3&)m_allAabbsCPU[i].m_min, (const b3Vector3&)m_allAabbsCPU[i].m_max,(const b3Vector3&)m_allAabbsCPU[otherIndex].m_min,(const b3Vector3&)m_allAabbsCPU[otherIndex].m_max);
@@ -906,7 +950,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHost(int maxPairs)
for (int i=0;i<numSmallAabbs;i++)
{
b3SapAabb smallAabbi = m_allAabbsCPU[m_smallAabbsMappingCPU[i]];
float reference = smallAabbi.m_max[axis];
//float reference = smallAabbi.m_max[axis];
for (int j=i+1;j<numSmallAabbs;j++)
{
@@ -941,7 +985,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairsHost(int maxPairs)
{
b3SapAabb smallAabbi = m_allAabbsCPU[m_smallAabbsMappingCPU[i]];
float reference = smallAabbi.m_max[axis];
//float reference = smallAabbi.m_max[axis];
int numLargeAabbs = m_largeAabbsMappingCPU.size();
for (int j=0;j<numLargeAabbs;j++)
@@ -1022,7 +1066,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairs(int maxPairs)
{
bool syncOnHost = false;
//bool syncOnHost = false;
int numSmallAabbs = m_smallAabbsMappingCPU.size();
if (m_prefixScanFloat4 && numSmallAabbs)
@@ -1261,7 +1305,7 @@ void b3GpuSapBroadphase::writeAabbsToGpu()
}
void b3GpuSapBroadphase::createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask)
void b3GpuSapBroadphase::createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr , int collisionFilterGroup, int collisionFilterMask)
{
int index = userPtr;
b3SapAabb aabb;
@@ -1277,7 +1321,7 @@ void b3GpuSapBroadphase::createLargeProxy(const b3Vector3& aabbMin, const b3Vec
m_allAabbsCPU.push_back(aabb);
}
void b3GpuSapBroadphase::createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask)
void b3GpuSapBroadphase::createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr , int collisionFilterGroup, int collisionFilterMask)
{
int index = userPtr;
b3SapAabb aabb;

4
src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h
View File

@@ -133,8 +133,8 @@ public:
void init3dSap();
virtual void calculateOverlappingPairsHostIncremental3Sap();
virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask);
virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask);
virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr , int collisionFilterGroup, int collisionFilterMask);
virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr , int collisionFilterGroup, int collisionFilterMask);
//call writeAabbsToGpu after done making all changes (createProxy etc)
virtual void writeAabbsToGpu();

10
src/Bullet3OpenCL/Initialize/b3OpenCLUtils.cpp
View File

@@ -172,7 +172,7 @@ cl_platform_id b3OpenCLUtils_getPlatform(int platformIndex0, cl_int* pErrNum)
cl_uint numPlatforms;
cl_int ciErrNum = clGetPlatformIDs(0, NULL, &numPlatforms);
if (platformIndex>=0 && platformIndex<numPlatforms)
if (platformIndex<numPlatforms)
{
cl_platform_id* platforms = (cl_platform_id*) malloc (sizeof(cl_platform_id)*numPlatforms);
ciErrNum = clGetPlatformIDs(numPlatforms, platforms, NULL);
@@ -583,7 +583,7 @@ static const char* strip2(const char* name, const char* pattern)
const char * oriptr;
const char * patloc;
// find how many times the pattern occurs in the original string
for (oriptr = name; patloc = strstr(oriptr, pattern); oriptr = patloc + patlen)
for (oriptr = name; (patloc = strstr(oriptr, pattern)); oriptr = patloc + patlen)
{
patcnt++;
}
@@ -608,8 +608,9 @@ cl_program b3OpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev
char driverVersion[256];
const char* strippedName;
int fileUpToDate = 0;
#ifdef _WIN32
int binaryFileValid=0;
#endif
if (!disableBinaryCaching && clFileNameForCaching)
{
clGetDeviceInfo(device, CL_DEVICE_NAME, 256, &deviceName, NULL);
@@ -862,7 +863,8 @@ cl_program b3OpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev
int kernelSize = ftell( file );
rewind( file );
kernelSrc = (char*)malloc(kernelSize+1);
int readBytes = fread((void*)kernelSrc,1,kernelSize, file);
int readBytes;
readBytes = fread((void*)kernelSrc,1,kernelSize, file);
kernelSrc[kernelSize] = 0;
fclose(file);
kernelSource = kernelSrc;

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

@@ -105,21 +105,29 @@ GpuSatCollision::GpuSatCollision(cl_context ctx,cl_device_id device, cl_command_
:m_context(ctx),
m_device(device),
m_queue(q),
m_findSeparatingAxisKernel(0),
m_findSeparatingAxisVertexFaceKernel(0),
m_findSeparatingAxisEdgeEdgeKernel(0),
m_unitSphereDirections(m_context,m_queue),
m_totalContactsOut(m_context, m_queue),
m_sepNormals(m_context, m_queue),
m_dmins(m_context,m_queue),
m_hasSeparatingNormals(m_context, m_queue),
m_concaveSepNormals(m_context, m_queue),
m_concaveHasSeparatingNormals(m_context,m_queue),
m_numConcavePairsOut(m_context, m_queue),
m_gpuCompoundPairs(m_context, m_queue),
m_gpuCompoundSepNormals(m_context, m_queue),
m_gpuHasCompoundSepNormals(m_context, m_queue),
m_numCompoundPairsOut(m_context, m_queue),
m_dmins(m_context,m_queue),
m_unitSphereDirections(m_context,m_queue)
m_numCompoundPairsOut(m_context, m_queue)
{
m_totalContactsOut.push_back(0);
@@ -553,7 +561,7 @@ inline void project(const b3ConvexPolyhedronData& hull, const float4& pos, cons
{
//b3Vector3 pt = trans * vertices[m_vertexOffset+i];
//b3Scalar dp = pt.dot(dir);
b3Vector3 vertex = vertices[hull.m_vertexOffset+i];
//b3Vector3 vertex = vertices[hull.m_vertexOffset+i];
b3Scalar dp = dot3F4((float4&)vertices[hull.m_vertexOffset+i],localDir);
//b3Assert(dp==dpL);
if(dp < min) min = dp;
@@ -764,7 +772,7 @@ bool findSeparatingAxisEdgeEdge( __global const b3ConvexPolyhedronData* hullA, _
float4 posB = posB1;
posB.w = 0.f;
int curPlaneTests=0;
//int curPlaneTests=0;
int curEdgeEdge = 0;
// Test edges
@@ -927,7 +935,7 @@ int clipFaceAgainstHull(const float4& separatingNormal, const b3ConvexPolyhedron
b3GpuFace 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 numContacts = numWorldVertsB1;
// int numContacts = numWorldVertsB1;
int numVerticesA = polyA.m_numIndices;
for(int e0=0;e0<numVerticesA;e0++)
{
@@ -1008,7 +1016,7 @@ static int clipHullAgainstHull(const float4& separatingNormal,
B3_PROFILE("clipHullAgainstHull");
float curMaxDist=maxDist;
// float curMaxDist=maxDist;
int closestFaceB=-1;
float dmax = -FLT_MAX;
@@ -1318,7 +1326,7 @@ int clipHullHullSingle(
contact.m_frictionCoeffCmp = 45874;
contact.m_restituitionCoeffCmp = 0;
float distance = 0.f;
// float distance = 0.f;
for (int p=0;p<numPoints;p++)
{
contact.m_worldPosB[p] = contactsOut[contactIdx.s[p]];//check if it is actually on B
@@ -1362,8 +1370,8 @@ void computeContactPlaneConvex(int pairIndex,
b3Vector3 posA = rigidBodies[bodyIndexA].m_pos;
b3Quaternion ornA = rigidBodies[bodyIndexA].m_quat;
int numContactsOut = 0;
int numWorldVertsB1= 0;
// int numContactsOut = 0;
// int numWorldVertsB1= 0;
b3Vector3 planeEq = faces[collidables[collidableIndexA].m_shapeIndex].m_plane;
b3Vector3 planeNormal=b3MakeVector3(planeEq.x,planeEq.y,planeEq.z);
@@ -1518,7 +1526,7 @@ __kernel void findCompoundPairsKernel(
{
numAabbChecks=0;
maxNumAabbChecks=0;
int i = pairIndex;
// int i = pairIndex;
{
@@ -1590,10 +1598,10 @@ __kernel void findCompoundPairsKernel(
{
int startNodeIndexA = subtreeA.m_rootNodeIndex+bvhInfoCPU[bvhA].m_nodeOffset;
int endNodeIndexA = startNodeIndexA+subtreeA.m_subtreeSize;
// int endNodeIndexA = startNodeIndexA+subtreeA.m_subtreeSize;
int startNodeIndexB = subtreeB.m_rootNodeIndex+bvhInfoCPU[bvhB].m_nodeOffset;
int endNodeIndexB = startNodeIndexB+subtreeB.m_subtreeSize;
// int endNodeIndexB = startNodeIndexB+subtreeB.m_subtreeSize;
b3AlignedObjectArray<b3Int2> nodeStack;
b3Int2 node0;
@@ -1790,19 +1798,21 @@ __kernel void findCompoundPairsKernel(
{
if (1)
{
int numFacesA = convexShapes[shapeIndexA].m_numFaces;
float dmin = FLT_MAX;
// int numFacesA = convexShapes[shapeIndexA].m_numFaces;
// float dmin = FLT_MAX;
float4 posA = newPosA;
posA.w = 0.f;
float4 posB = rigidBodies[bodyIndexB].m_pos;
posB.w = 0.f;
float4 c0local = convexShapes[shapeIndexA].m_localCenter;
b3Quat ornA = newOrnA;
float4 c0 = transform(&c0local, &posA, &ornA);
float4 c0;
c0 = transform(&c0local, &posA, &ornA);
float4 c1local = convexShapes[shapeIndexB].m_localCenter;
b3Quat ornB = rigidBodies[bodyIndexB].m_quat;
float4 c1 = transform(&c1local,&posB,&ornB);
const float4 DeltaC2 = c0 - c1;
float4 c1;
c1 = transform(&c1local,&posB,&ornB);
// const float4 DeltaC2 = c0 - c1;
{
int compoundPairIdx = b3AtomicInc(numCompoundPairsOut);
@@ -1838,19 +1848,21 @@ __kernel void findCompoundPairsKernel(
if (1)
{
int numFacesA = convexShapes[shapeIndexA].m_numFaces;
float dmin = FLT_MAX;
// int numFacesA = convexShapes[shapeIndexA].m_numFaces;
// float dmin = FLT_MAX;
float4 posA = rigidBodies[bodyIndexA].m_pos;
posA.w = 0.f;
float4 posB = newPosB;
posB.w = 0.f;
float4 c0local = convexShapes[shapeIndexA].m_localCenter;
b3Quat ornA = rigidBodies[bodyIndexA].m_quat;
float4 c0 = transform(&c0local, &posA, &ornA);
float4 c0;
c0 = transform(&c0local, &posA, &ornA);
float4 c1local = convexShapes[shapeIndexB].m_localCenter;
b3Quat ornB =newOrnB;
float4 c1 = transform(&c1local,&posB,&ornB);
const float4 DeltaC2 = c0 - c1;
float4 c1;
c1 = transform(&c1local,&posB,&ornB);
// const float4 DeltaC2 = c0 - c1;
{//
int compoundPairIdx = b3AtomicInc(numCompoundPairsOut);
if (compoundPairIdx<maxNumCompoundPairsCapacity)
@@ -1948,7 +1960,7 @@ __kernel void processCompoundPairsKernel( __global const b3Int4* gpuCompoundPa
int hasSeparatingAxis = 5;
int numFacesA = convexShapes[shapeIndexA].m_numFaces;
// int numFacesA = convexShapes[shapeIndexA].m_numFaces;
float dmin = FLT_MAX;
posA.w = 0.f;
posB.w = 0.f;
@@ -2326,8 +2338,8 @@ void computeContactPlaneCompound(int pairIndex,
b3Vector3 posA = rigidBodies[bodyIndexA].m_pos;
b3Quaternion ornA = rigidBodies[bodyIndexA].m_quat;
int numContactsOut = 0;
int numWorldVertsB1= 0;
// int numContactsOut = 0;
// int numWorldVertsB1= 0;
b3Vector3 planeEq = faces[collidables[collidableIndexA].m_shapeIndex].m_plane;
b3Vector3 planeNormal=b3MakeVector3(planeEq.x,planeEq.y,planeEq.z);
@@ -2475,7 +2487,7 @@ void computeContactSphereConvex(int pairIndex,
int shapeIndex = collidables[collidableIndex].m_shapeIndex;
int numFaces = convexShapes[shapeIndex].m_numFaces;
float4 closestPnt = b3MakeVector3(0, 0, 0, 0);
float4 hitNormalWorld = b3MakeVector3(0, 0, 0, 0);
// float4 hitNormalWorld = b3MakeVector3(0, 0, 0, 0);
float minDist = -1000000.f; // TODO: What is the largest/smallest float?
bool bCollide = true;
int region = -1;
@@ -2641,7 +2653,7 @@ int computeContactConvexConvex( b3AlignedObjectArray<b3Int4>& pairs,
int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
int sz = sizeof(b3Contact4);
//int sz = sizeof(b3Contact4);
bool result2 = getClosestPoints(&gjkDetector, transA, transB,
convexShapes[shapeIndexA], convexShapes[shapeIndexB],
@@ -2709,8 +2721,8 @@ int computeContactConvexConvex( b3AlignedObjectArray<b3Int4>& pairs,
{
resultPointOnBWorld.w = distance2;
newContact.m_worldPosB[p] = resultPointOnBWorld;
b3Vector3 resultPointOnAWorld = resultPointOnBWorld+distance2*sepAxis2;
#ifdef PERSISTENT_CONTACTS_HOST
b3Vector3 resultPointOnAWorld = resultPointOnBWorld+distance2*sepAxis2;
newContact.m_localPosA[p] = transA.inverse()*resultPointOnAWorld;
newContact.m_localPosB[p] = transB.inverse()*resultPointOnBWorld;
#endif
@@ -2774,8 +2786,8 @@ int computeContactConvexConvex2(
hullB = convexShapes[colB.m_shapeIndex];
//printf("numvertsB = %d\n",hullB.m_numVertices);
int contactCapacity = MAX_VERTS;
int numContactsOut=0;
// int contactCapacity = MAX_VERTS;
//int numContactsOut=0;
#ifdef _WIN32
@@ -3488,7 +3500,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
//add contact point
int contactIndex = nGlobalContactsOut;
//int contactIndex = nGlobalContactsOut;
b3Contact4& newContact = hostContacts.at(nGlobalContactsOut);
nGlobalContactsOut++;
newContact.m_batchIdx = 0;//i;
@@ -3511,7 +3523,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
resultPointOnBWorld.w = -depth;
newContact.m_worldPosB[0] = resultPointOnBWorld;
b3Vector3 resultPointOnAWorld = resultPointOnBWorld+depth*sepAxis2;
//b3Vector3 resultPointOnAWorld = resultPointOnBWorld+depth*sepAxis2;
newContact.m_worldNormalOnB = sepAxis2;
newContact.m_worldNormalOnB.w = (b3Scalar)1;
} else
@@ -3527,7 +3539,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
int result = computeContactConvexConvex( hostPairs,
int result;
result = computeContactConvexConvex( hostPairs,
pairIndex,
bodyIndexA, bodyIndexB,
collidableIndexA, collidableIndexB,
@@ -3893,7 +3906,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
b3AlignedObjectArray<b3Int4> triangleConvexPairsOutHost;
triangleConvexPairsOutHost.resize(maxTriConvexPairCapacity);
int numTriConvexPairsOutHost=0;
//int numTriConvexPairsOutHost=0;
numConcavePairs = 0;
//m_numConcavePairsOut

12
src/Bullet3OpenCL/NarrowphaseCollision/b3GjkPairDetector.cpp
View File

@@ -97,7 +97,7 @@ inline void project(const b3ConvexPolyhedronData& hull, const float4& pos, cons
{
//b3Vector3 pt = trans * vertices[m_vertexOffset+i];
//b3Scalar dp = pt.dot(dir);
b3Vector3 vertex = vertices[hull.m_vertexOffset+i];
//b3Vector3 vertex = vertices[hull.m_vertexOffset+i];
b3Scalar dp = dot3F4((float4&)vertices[hull.m_vertexOffset+i],localDir);
//b3Assert(dp==dpL);
if(dp < min) min = dp;
@@ -113,7 +113,7 @@ inline void project(const b3ConvexPolyhedronData& hull, const float4& pos, cons
max += offset;
}
#if 0
static bool TestSepAxis(const b3ConvexPolyhedronData& hullA, const b3ConvexPolyhedronData& hullB,
const float4& posA,const b3Quaternion& ornA,
const float4& posB,const b3Quaternion& ornB,
@@ -143,7 +143,7 @@ static bool TestSepAxis(const b3ConvexPolyhedronData& hullA, const b3ConvexPolyh
return true;
}
#endif
bool getClosestPoints(b3GjkPairDetector* gjkDetector, const b3Transform& transA, const b3Transform& transB,
const b3ConvexPolyhedronData& hullA, const b3ConvexPolyhedronData& hullB,
@@ -198,11 +198,11 @@ bool getClosestPoints(b3GjkPairDetector* gjkDetector, const b3Transform& transA,
{
b3Scalar squaredDistance = B3_LARGE_FLOAT;
b3Scalar delta = -1e30f;//b3Scalar(0.);
b3Scalar prevDelta = -1e30f;//b3Scalar(0.);
// b3Scalar prevDelta = -1e30f;//b3Scalar(0.);
b3Scalar margin = marginA + marginB;
b3Scalar bestDeltaN = -1e30f;
b3Vector3 bestSepAxis= b3MakeVector3(0,0,0);
// b3Scalar bestDeltaN = -1e30f;
// b3Vector3 bestSepAxis= b3MakeVector3(0,0,0);

2
src/Bullet3OpenCL/NarrowphaseCollision/b3QuantizedBvh.h
View File

@@ -157,7 +157,7 @@ typedef b3AlignedObjectArray<b3BvhSubtreeInfo> BvhSubtreeInfoArray;
///The b3QuantizedBvh class stores an AABB tree that can be quickly traversed on CPU and Cell SPU.
///It is used by the b3BvhTriangleMeshShape as midphase, and by the b3MultiSapBroadphase.
///It is used by the b3BvhTriangleMeshShape as midphase
///It is recommended to use quantization for better performance and lower memory requirements.
B3_ATTRIBUTE_ALIGNED16(class) b3QuantizedBvh
{

2
src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp
View File

@@ -196,7 +196,7 @@ void b3GpuRaycast::castRaysHost(const b3AlignedObjectArray<b3RayInfo>& rays, b3A
{
const b3Vector3& pos = bodies[b].m_pos;
const b3Quaternion& orn = bodies[b].m_quat;
//const b3Quaternion& orn = bodies[b].m_quat;
switch (collidables[bodies[b].m_collidableIdx].m_shapeType)
{

2
src/Bullet3OpenCL/RigidBody/b3GpuGenericConstraint.cpp
View File

@@ -55,7 +55,7 @@ void getInfo2Point2Point(b3GpuGenericConstraint* constraint, b3GpuConstraintInfo
info->m_J1linearAxis[2*info->rowskip+2] = 1;
b3Vector3 a1 = trA.getBasis()*constraint->getPivotInA();
b3Vector3 a1a = b3QuatRotate(trA.getRotation(),constraint->getPivotInA());
//b3Vector3 a1a = b3QuatRotate(trA.getRotation(),constraint->getPivotInA());
{
b3Vector3* angular0 = (b3Vector3*)(info->m_J1angularAxis);

4
src/Bullet3OpenCL/RigidBody/b3GpuJacobiContactSolver.cpp
View File

@@ -576,7 +576,7 @@ void b3GpuJacobiContactSolver::solveGroupHost(b3RigidBodyData* bodies,b3InertiaD
for( i=0; i<numManifolds; i++)
{
float frictionCoeff = contactConstraints[i].getFrictionCoeff();
//float frictionCoeff = contactConstraints[i].getFrictionCoeff();
int aIdx = (int)contactConstraints[i].m_bodyA;
int bIdx = (int)contactConstraints[i].m_bodyB;
b3RigidBodyData& bodyA = bodies[aIdx];
@@ -650,7 +650,7 @@ void b3GpuJacobiContactSolver::solveGroupHost(b3RigidBodyData* bodies,b3InertiaD
}
for (int iter = 0;iter<maxIter;iter++)
{
int i=0;
//int i=0;
//solve friction

4
src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.cpp
View File

@@ -445,7 +445,7 @@ int b3GpuNarrowPhase::registerCompoundShape(b3AlignedObjectArray<b3GpuChildShap
for (int i=0;i<childShapes->size();i++)
{
int childColIndex = childShapes->at(i).m_shapeIndex;
b3Collidable& childCol = getCollidableCpu(childColIndex);
//b3Collidable& childCol = getCollidableCpu(childColIndex);
b3SapAabb aabbLoc =m_data->m_localShapeAABBCPU->at(childColIndex);
b3Vector3 childLocalAabbMin=b3MakeVector3(aabbLoc.m_min[0],aabbLoc.m_min[1],aabbLoc.m_min[2]);
@@ -832,7 +832,7 @@ void b3GpuNarrowPhase::computeContacts(cl_mem broadphasePairs, int numBroadphase
//swap buffer
m_data->m_currentContactBuffer=1-m_data->m_currentContactBuffer;
int curSize = m_data->m_pBufContactBuffersGPU[m_data->m_currentContactBuffer]->size();
//int curSize = m_data->m_pBufContactBuffersGPU[m_data->m_currentContactBuffer]->size();
int maxTriConvexPairCapacity = m_data->m_config.m_maxTriConvexPairCapacity;
int numTriConvexPairsOut=0;

18
src/Bullet3OpenCL/RigidBody/b3GpuPgsConstraintSolver.cpp
View File

@@ -84,7 +84,7 @@ struct b3GpuPgsJacobiSolverInternalData
};
/*
static b3Transform getWorldTransform(b3RigidBodyData* rb)
{
b3Transform newTrans;
@@ -98,7 +98,7 @@ static const b3Matrix3x3& getInvInertiaTensorWorld(b3InertiaData* inertia)
return inertia->m_invInertiaWorld;
}
*/
static const b3Vector3& getLinearVelocity(b3RigidBodyData* rb)
{
@@ -254,7 +254,7 @@ b3Scalar b3GpuPgsConstraintSolver::solveGroupCacheFriendlySetup(b3OpenCLArray<b3
}
}
int totalBodies = 0;
// int totalBodies = 0;
int totalNumRows = 0;
//b3RigidBody* rb0=0,*rb1=0;
//if (1)
@@ -604,9 +604,9 @@ b3Scalar b3GpuPgsConstraintSolver::solveGroupCacheFriendlySetup(b3OpenCLArray<b3
// b3ContactSolverInfo info = infoGlobal;
int numNonContactPool = m_tmpSolverNonContactConstraintPool.size();
int numConstraintPool = m_tmpSolverContactConstraintPool.size();
int numFrictionPool = m_tmpSolverContactFrictionConstraintPool.size();
// int numNonContactPool = m_tmpSolverNonContactConstraintPool.size();
// int numConstraintPool = m_tmpSolverContactConstraintPool.size();
// int numFrictionPool = m_tmpSolverContactFrictionConstraintPool.size();
return 0.f;
@@ -858,7 +858,7 @@ static b3AlignedObjectArray<int> curUsed;
inline int b3GpuPgsConstraintSolver::sortConstraintByBatch3( b3BatchConstraint* cs, int numConstraints, int simdWidth , int staticIdx, int numBodies)
{
int sz = sizeof(b3BatchConstraint);
//int sz = sizeof(b3BatchConstraint);
B3_PROFILE("sortConstraintByBatch3");
@@ -892,7 +892,7 @@ inline int b3GpuPgsConstraintSolver::sortConstraintByBatch3( b3BatchConstraint*
#endif
int numValidConstraints = 0;
int unprocessedConstraintIndex = 0;
// int unprocessedConstraintIndex = 0;
int batchIdx = 0;
@@ -1032,7 +1032,7 @@ void b3GpuPgsConstraintSolver::solveJoints(int numBodies, b3OpenCLArray<b3RigidB
b3Scalar b3GpuPgsConstraintSolver::solveGroupCacheFriendlyFinish(b3OpenCLArray<b3RigidBodyData>* gpuBodies,b3OpenCLArray<b3InertiaData>* gpuInertias,int numBodies,b3OpenCLArray<b3GpuGenericConstraint>* gpuConstraints,int numConstraints,const b3ContactSolverInfo& infoGlobal)
{
B3_PROFILE("solveGroupCacheFriendlyFinish");
int numPoolConstraints = m_tmpSolverContactConstraintPool.size();
// int numPoolConstraints = m_tmpSolverContactConstraintPool.size();
// int i,j;

67
src/Bullet3OpenCL/RigidBody/b3GpuPgsContactSolver.cpp
View File

@@ -138,7 +138,7 @@ b3GpuPgsContactSolver::b3GpuPgsContactSolver(cl_context ctx,cl_device_id device,
m_data->m_offsets = new b3OpenCLArray<unsigned int>( ctx,m_data->m_queue,B3_SOLVER_N_CELLS);
m_data->m_offsets->resize(B3_SOLVER_N_CELLS);
const char* additionalMacros = "";
const char* srcFileNameForCaching="";
//const char* srcFileNameForCaching="";
@@ -556,26 +556,8 @@ static bool b3ContactCmp(const b3Contact4& p, const b3Contact4& q)
static const int gridTable4x4[] =
{
0,1,17,16,
1,2,18,19,
17,18,32,3,
16,19,3,34
};
static const int gridTable8x8[] =
{
0, 2, 3, 16, 17, 18, 19, 1,
66, 64, 80, 67, 82, 81, 65, 83,
131,144,128,130,147,129,145,146,
208,195,194,192,193,211,210,209,
21, 22, 23, 5, 4, 6, 7, 20,
86, 85, 69, 87, 70, 68, 84, 71,
151,133,149,150,135,148,132,134,
197,27,214,213,212,199,198,196
};
@@ -583,6 +565,30 @@ static const int gridTable8x8[] =
#define USE_SPATIAL_BATCHING 1
#define USE_4x4_GRID 1
#ifndef USE_SPATIAL_BATCHING
static const int gridTable4x4[] =
{
0,1,17,16,
1,2,18,19,
17,18,32,3,
16,19,3,34
};
static const int gridTable8x8[] =
{
0, 2, 3, 16, 17, 18, 19, 1,
66, 64, 80, 67, 82, 81, 65, 83,
131,144,128,130,147,129,145,146,
208,195,194,192,193,211,210,209,
21, 22, 23, 5, 4, 6, 7, 20,
86, 85, 69, 87, 70, 68, 84, 71,
151,133,149,150,135,148,132,134,
197,27,214,213,212,199,198,196
};
#endif
void SetSortDataCPU(b3Contact4* gContact, b3RigidBodyData* gBodies, b3SortData* gSortDataOut, int nContacts,float scale,const b3Int4& nSplit,int staticIdx)
{
@@ -597,7 +603,6 @@ void SetSortDataCPU(b3Contact4* gContact, b3RigidBodyData* gBodies, b3SortData*
int bIdx = abs(bPtrAndSignBit);
bool aStatic = (aPtrAndSignBit<0) ||(aPtrAndSignBit==staticIdx);
bool bStatic = (bPtrAndSignBit<0) ||(bPtrAndSignBit==staticIdx);
#if USE_SPATIAL_BATCHING
int idx = (aStatic)? bIdx: aIdx;
@@ -609,6 +614,8 @@ void SetSortDataCPU(b3Contact4* gContact, b3RigidBodyData* gBodies, b3SortData*
int newIndex = (xIdx+yIdx*nSplit.x+zIdx*nSplit.x*nSplit.y);
#else//USE_SPATIAL_BATCHING
bool bStatic = (bPtrAndSignBit<0) ||(bPtrAndSignBit==staticIdx);
#if USE_4x4_GRID
int aa = aIdx&3;
int bb = bIdx&3;
@@ -797,7 +804,7 @@ void b3GpuPgsContactSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
const b3OpenCLArray<b3RigidBodyData>* bodyNative = bodyBuf;
//const b3OpenCLArray<b3RigidBodyData>* bodyNative = bodyBuf;
{
@@ -1033,7 +1040,7 @@ void b3GpuPgsContactSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
}
bool compareGPU = false;
// bool compareGPU = false;
if (nContacts)
{
if (!gCpuBatchContacts)
@@ -1070,7 +1077,7 @@ void b3GpuPgsContactSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
{
m_data->m_batchSizes.resize(B3_MAX_NUM_BATCHES);
int totalNumConstraints = cpuContacts.size();
int simdWidth =numBodies+1;//-1;//64;//-1;//32;
//int simdWidth =numBodies+1;//-1;//64;//-1;//32;
int numBatches = sortConstraintByBatch3( &cpuContacts[0], totalNumConstraints, totalNumConstraints+1,csCfg.m_staticIdx ,numBodies,&m_data->m_batchSizes[0]); // on GPU
maxNumBatches = b3Max(numBatches,maxNumBatches);
static int globalMaxBatch = 0;
@@ -1138,17 +1145,17 @@ void b3GpuPgsContactSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
B3_PROFILE("copyToHost");
m_data->m_pBufContactOutGPU->copyToHost(cpuContacts);
}
b3OpenCLArray<unsigned int>* countsNative = m_data->m_solverGPU->m_numConstraints;
b3OpenCLArray<unsigned int>* offsetsNative = m_data->m_solverGPU->m_offsets;
// b3OpenCLArray<unsigned int>* countsNative = m_data->m_solverGPU->m_numConstraints;
// b3OpenCLArray<unsigned int>* offsetsNative = m_data->m_solverGPU->m_offsets;
int numNonzeroGrid=0;
// int numNonzeroGrid=0;
{
m_data->m_batchSizes.resize(B3_MAX_NUM_BATCHES);
int totalNumConstraints = cpuContacts.size();
int simdWidth =numBodies+1;//-1;//64;//-1;//32;
// int simdWidth =numBodies+1;//-1;//64;//-1;//32;
int numBatches = sortConstraintByBatch3( &cpuContacts[0], totalNumConstraints, totalNumConstraints+1,csCfg.m_staticIdx ,numBodies,&m_data->m_batchSizes[0]); // on GPU
maxNumBatches = b3Max(numBatches,maxNumBatches);
static int globalMaxBatch = 0;
@@ -1420,7 +1427,7 @@ inline int b3GpuPgsContactSolver::sortConstraintByBatch2( b3Contact4* cs, int nu
idxSrc[i] = i;
int numValidConstraints = 0;
int unprocessedConstraintIndex = 0;
// int unprocessedConstraintIndex = 0;
int batchIdx = 0;
@@ -1588,7 +1595,7 @@ inline int b3GpuPgsContactSolver::sortConstraintByBatch3( b3Contact4* cs, int nu
#endif
int numValidConstraints = 0;
int unprocessedConstraintIndex = 0;
// int unprocessedConstraintIndex = 0;
int batchIdx = 0;

49
src/Bullet3OpenCL/RigidBody/b3Solver.cpp
View File

@@ -87,11 +87,12 @@ public:
b3Solver::b3Solver(cl_context ctx, cl_device_id device, cl_command_queue queue, int pairCapacity)
:m_nIterations(4),
:
m_context(ctx),
m_device(device),
m_queue(queue),
m_batchSizes(ctx,queue)
m_batchSizes(ctx,queue),
m_nIterations(4)
{
m_sort32 = new b3RadixSort32CL(ctx,device,queue);
m_scan = new b3PrefixScanCL(ctx,device,queue,B3_SOLVER_N_CELLS);
@@ -108,7 +109,7 @@ b3Solver::b3Solver(cl_context ctx, cl_device_id device, cl_command_queue queue,
m_offsets = new b3OpenCLArray<unsigned int>( ctx,queue,B3_SOLVER_N_CELLS);
m_offsets->resize(B3_SOLVER_N_CELLS);
const char* additionalMacros = "";
const char* srcFileNameForCaching="";
// const char* srcFileNameForCaching="";
@@ -356,17 +357,32 @@ void solveContact(b3GpuConstraint4& cs,
}
}
/*
b3AlignedObjectArray<b3RigidBodyData>& m_bodies;
b3AlignedObjectArray<b3InertiaData>& m_shapes;
b3AlignedObjectArray<b3GpuConstraint4>& m_constraints;
b3AlignedObjectArray<int>* m_batchSizes;
int m_cellIndex;
int m_curWgidx;
int m_start;
int m_nConstraints;
bool m_solveFriction;
int m_maxNumBatches;
*/
struct SolveTask// : public ThreadPool::Task
{
SolveTask(b3AlignedObjectArray<b3RigidBodyData>& bodies, b3AlignedObjectArray<b3InertiaData>& shapes, b3AlignedObjectArray<b3GpuConstraint4>& constraints,
int start, int nConstraints,int maxNumBatches,b3AlignedObjectArray<int>* wgUsedBodies, int curWgidx, b3AlignedObjectArray<int>* batchSizes, int cellIndex)
: m_bodies( bodies ), m_shapes( shapes ), m_constraints( constraints ), m_start( start ), m_nConstraints( nConstraints ),
m_solveFriction( true ),m_maxNumBatches(maxNumBatches),
m_curWgidx(curWgidx),
: m_bodies( bodies ), m_shapes( shapes ),
m_constraints( constraints ),
m_batchSizes(batchSizes),
m_cellIndex(cellIndex)
m_cellIndex(cellIndex),
m_curWgidx(curWgidx),
m_start( start ),
m_nConstraints( nConstraints ),
m_solveFriction( true ),
m_maxNumBatches(maxNumBatches)
{}
unsigned short int getType(){ return 0; }
@@ -388,7 +404,7 @@ struct SolveTask// : public ThreadPool::Task
float frictionCoeff = m_constraints[i].getFrictionCoeff();
int aIdx = (int)m_constraints[i].m_bodyA;
int bIdx = (int)m_constraints[i].m_bodyB;
int localBatch = m_constraints[i].m_batchIdx;
// int localBatch = m_constraints[i].m_batchIdx;
b3RigidBodyData& bodyA = m_bodies[aIdx];
b3RigidBodyData& bodyB = m_bodies[bIdx];
@@ -604,7 +620,7 @@ void b3Solver::solveContactConstraintHost( b3OpenCLArray<b3RigidBodyData>* body
}
const int start = offsetsHost[cellIdx];
int numConstraintsInCell = numConstraintsHost[cellIdx];
const int end = start + numConstraintsInCell;
// const int end = start + numConstraintsInCell;
SolveTask task( bodyNative, shapeNative, constraintNative, start, numConstraintsInCell ,maxNumBatches,usedBodies,wgIdx,batchSizes,cellIdx);
task.m_solveFriction = false;
@@ -640,7 +656,7 @@ void b3Solver::solveContactConstraintHost( b3OpenCLArray<b3RigidBodyData>* body
const int start = offsetsHost[cellIdx];
int numConstraintsInCell = numConstraintsHost[cellIdx];
const int end = start + numConstraintsInCell;
// const int end = start + numConstraintsInCell;
SolveTask task( bodyNative, shapeNative, constraintNative, start, numConstraintsInCell,maxNumBatches, 0,0,batchSizes,cellIdx);
task.m_solveFriction = true;
@@ -688,7 +704,7 @@ void checkConstraintBatch(const b3OpenCLArray<b3RigidBodyData>* bodyBuf,
int cellBatch = batchId;
const int nn = B3_SOLVER_N_CELLS;
int numWorkItems = 64*nn/B3_SOLVER_N_BATCHES;
// int numWorkItems = 64*nn/B3_SOLVER_N_BATCHES;
b3AlignedObjectArray<unsigned int> gN;
m_numConstraints->copyToHost(gN);
@@ -697,7 +713,7 @@ void checkConstraintBatch(const b3OpenCLArray<b3RigidBodyData>* bodyBuf,
int nSplitX = B3_SOLVER_N_SPLIT_X;
int nSplitY = B3_SOLVER_N_SPLIT_Y;
int bIdx = batchId;
// int bIdx = batchId;
b3AlignedObjectArray<b3GpuConstraint4> cpuConstraints;
constraint->copyToHost(cpuConstraints);
@@ -932,7 +948,7 @@ void b3Solver::convertToConstraints( const b3OpenCLArray<b3RigidBodyData>* bodyB
b3OpenCLArray<b3Contact4>* contactsIn, b3OpenCLArray<b3GpuConstraint4>* contactCOut, void* additionalData,
int nContacts, const ConstraintCfg& cfg )
{
b3OpenCLArray<b3GpuConstraint4>* constraintNative =0;
// b3OpenCLArray<b3GpuConstraint4>* constraintNative =0;
contactCOut->resize(nContacts);
struct CB
{
@@ -1099,7 +1115,6 @@ void b3Solver::sortContacts( const b3OpenCLArray<b3RigidBodyData>* bodyBuf,
}
*/
void b3Solver::batchContacts( b3OpenCLArray<b3Contact4>* contacts, int nContacts, b3OpenCLArray<unsigned int>* nNative, b3OpenCLArray<unsigned int>* offsetsNative, int staticIdx )
{
@@ -1122,7 +1137,7 @@ void b3Solver::batchContacts( b3OpenCLArray<b3Contact4>* contacts, int nContact
#if 0
b3BufferInfoCL bInfo[] = {
b3BufferInfoCL( contacts->getBufferCL() ),
b3BufferInfoCL( m_contactBuffer2->getBufferCL()),
@@ -1132,7 +1147,7 @@ void b3Solver::batchContacts( b3OpenCLArray<b3Contact4>* contacts, int nContact
, b3BufferInfoCL(&gpuDebugInfo)
#endif
};
#endif