From d13de109a5867ee2acf22cd0b5f3f150adcf6310 Mon Sep 17 00:00:00 2001
From: erwin coumans <erwin.coumans@gmail.com>
Date: Sun, 10 Nov 2013 23:19:51 -0800
Subject: [PATCH] work towards re-enabling GPU grid broadphase (in addition to
 GPU sap bp)

---
 .../GpuDemos/rigidbody/GpuRigidBodyDemo.cpp   |  2 +-
 .../b3GpuGridBroadphase.cpp                   | 29 ++++++++++++++++---
 .../kernels/gridBroadphase.cl                 | 12 ++++----
 .../RigidBody/b3GpuPgsJacobiSolver.cpp        |  1 +
 4 files changed, 33 insertions(+), 11 deletions(-)

diff --git a/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp b/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp
index 838ebcb5a..f2c9f4915 100644
--- a/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp
+++ b/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp
@@ -115,7 +115,7 @@ void	GpuRigidBodyDemo::initPhysics(const ConstructionInfo& ci)
 		
 		m_data->m_config.m_maxConvexBodies = b3Max(m_data->m_config.m_maxConvexBodies,ci.arraySizeX*ci.arraySizeY*ci.arraySizeZ+10);
 		m_data->m_config.m_maxConvexShapes = m_data->m_config.m_maxConvexBodies;
-		int maxPairsPerBody = 128;
+		int maxPairsPerBody = 16;
 		m_data->m_config.m_maxBroadphasePairs = maxPairsPerBody*m_data->m_config.m_maxConvexBodies;
 		m_data->m_config.m_maxContactCapacity = m_data->m_config.m_maxBroadphasePairs;
 		
diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp
index 354434f43..97c730b2d 100644
--- a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp
+++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp
@@ -115,8 +115,21 @@ void b3GpuGridBroadphase::createLargeProxy(const b3Vector3& aabbMin,  const b3Ve
 
 void  b3GpuGridBroadphase::calculateOverlappingPairs(int maxPairs)
 {
+	B3_PROFILE("b3GpuGridBroadphase::calculateOverlappingPairs");
+	/*
 	calculateOverlappingPairsHost(maxPairs);
-	return;
+	{
+
+		b3AlignedObjectArray<b3Int4> cpuPairs;
+		m_gpuPairs.copyToHost(cpuPairs);
+		printf("host m_gpuPairs.size()=%d\n",m_gpuPairs.size());
+		for (int i=0;i<m_gpuPairs.size();i++)
+		{
+			printf("host pair %d = %d,%d\n",i,cpuPairs[i].x,cpuPairs[i].y);
+		}
+	}
+	*/
+	//return;
 	int numAabbs = m_allAabbsGPU.size();
 	if (numAabbs)
 	{
@@ -204,11 +217,19 @@ void  b3GpuGridBroadphase::calculateOverlappingPairs(int maxPairs)
 
 			
 			int actualCount = pairCount.at(0);
-			//b3AlignedObjectArray<b3Int4> pairsCpu;
 			m_gpuPairs.resize(actualCount);
-			//m_gpuPairs.copyToHost(pairsCpu);
-			//printf("?!?\n");
+			/*
+			b3AlignedObjectArray<b3Int4> pairsCpu;
+			m_gpuPairs.copyToHost(pairsCpu);
+			
+			printf("m_gpuPairs.size()=%d\n",m_gpuPairs.size());
+			for (int i=0;i<m_gpuPairs.size();i++)
+			{
+				printf("pair %d = %d,%d\n",i,pairsCpu[i].x,pairsCpu[i].y);
+			}
 
+			printf("?!?\n");
+			*/
 		}
 	
 
diff --git a/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphase.cl b/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphase.cl
index 9d4daf07c..aa453ca41 100644
--- a/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphase.cl
+++ b/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphase.cl
@@ -150,12 +150,12 @@ void findPairsInCell(	int numObjects,
 				if (pairCount)
 				{
 					int handleIndex2 = as_int(min1.w);
-					//if (handleIndex<handleIndex2)
+					if (handleIndex<handleIndex2)
 					{
 						int curPair = atomic_add(pairCount,1);
 						int4 newpair;
-						newpair.x = handleIndex2;
-						newpair.y = handleIndex;
+						newpair.x = handleIndex;
+						newpair.y = handleIndex2;
 						newpair.z = -1;
 						newpair.w = -1;
 						pPairBuff2[curPair] = newpair;
@@ -226,13 +226,13 @@ __kernel void kFindOverlappingPairs(	int numObjects,
     int4 gridPosA = getGridPos(pos, pParams);
     int4 gridPosB; 
     // examine only neighbouring cells
-    for(int z=0; z<=1; z++) 
+    for(int z=-1; z<=1; z++) 
     {
 		gridPosB.z = gridPosA.z + z;
-        for(int y=0; y<=1; y++) 
+        for(int y=-1; y<=1; y++) 
         {
 			gridPosB.y = gridPosA.y + y;
-            for(int x=0; x<=1; x++) 
+            for(int x=-1; x<=1; x++) 
             {
 				gridPosB.x = gridPosA.x + x;
                 findPairsInCell(numObjects, gridPosB, index, pHash, pCellStart, pAABB, pPairBuff, pPairBuffStartCurr, pParams, pairCount,pPairBuff2);
diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuPgsJacobiSolver.cpp b/src/Bullet3OpenCL/RigidBody/b3GpuPgsJacobiSolver.cpp
index e33bc455c..50c7250e2 100644
--- a/src/Bullet3OpenCL/RigidBody/b3GpuPgsJacobiSolver.cpp
+++ b/src/Bullet3OpenCL/RigidBody/b3GpuPgsJacobiSolver.cpp
@@ -14,6 +14,7 @@ subject to the following restrictions:
 */
 //Originally written by Erwin Coumans
 
+
 bool useGpuInitSolverBodies = true;
 bool useGpuInfo1 = true;
 bool useGpuInfo2= true;