prepare for btGpuGridBroadphase

Demos3/GpuDemos/GpuDemo.h

@@ -48,9 +48,9 @@ public:
 	arraySizeZ(10),
 	#else
 
-		arraySizeX(30),
+		arraySizeX(1),
-		arraySizeY(30),
+		arraySizeY(1),
-		arraySizeZ(30),
+		arraySizeZ(1),
 #endif
 		m_useConcaveMesh(false),
 		gapX(16.3),

Demos3/GpuDemos/broadphase/PairBench.cpp

@@ -4,6 +4,8 @@
 #include "Bullet3Common/b3Quaternion.h"
 #include "OpenGLWindow/b3gWindowInterface.h"
 #include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h"
+#include "Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h"
+
 #include "../GpuDemoInternalData.h"
 #include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
 #include "OpenGLWindow/OpenGLInclude.h"
@@ -103,7 +105,7 @@ __kernel void updateAabbSimple( __global float4* posOrnColors, const int numNode
 
 struct	PairBenchInternalData
 {
-	b3GpuSapBroadphase*	m_broadphaseGPU;
+	b3GpuBroadphaseInterface*	m_broadphaseGPU;
 
 	cl_kernel	m_moveObjectsKernel;
 	cl_kernel	m_sineWaveKernel;
@@ -176,6 +178,8 @@ void	PairBench::initPhysics(const ConstructionInfo& ci)
 	if (m_clData->m_clContext)
 	{
 		m_data->m_broadphaseGPU = new b3GpuSapBroadphase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue);
+		//m_data->m_broadphaseGPU = new b3GpuGridBroadphase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue);
+		
 		cl_program pairBenchProg=0;
 		int errNum=0;
 		m_data->m_moveObjectsKernel = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_pairBenchKernelString,"moveObjectsKernel",&errNum,pairBenchProg);
@@ -446,7 +450,7 @@ void PairBench::clientMoveAndDisplay()
 		}
 	}
 
-	bool updateOnGpu=false;
+	bool updateOnGpu=false;//true;
 
 	if (updateOnGpu)
 	{
@@ -463,7 +467,7 @@ void PairBench::clientMoveAndDisplay()
 		B3_PROFILE("updateOnCpu");
 		if (!gPairBenchFileName)
 		{
-		int allAabbs = m_data->m_broadphaseGPU->m_allAabbsCPU.size();
+		int allAabbs = m_data->m_broadphaseGPU->getAllAabbsCPU().size();
 			
 
 		b3AlignedObjectArray<b3Vector4> posOrnColorsCpu;
@@ -476,10 +480,10 @@ void PairBench::clientMoveAndDisplay()
 			{
 				b3Vector3 position = posOrnColorsCpu[nodeId];
 				b3Vector3 halfExtents = b3MakeFloat4(1.01f,1.01f,1.01f,0.f);
-				m_data->m_broadphaseGPU->m_allAabbsCPU[nodeId].m_minVec = position-halfExtents;
+				m_data->m_broadphaseGPU->getAllAabbsCPU()[nodeId].m_minVec = position-halfExtents;
-				m_data->m_broadphaseGPU->m_allAabbsCPU[nodeId].m_minIndices[3] = nodeId;
+				m_data->m_broadphaseGPU->getAllAabbsCPU()[nodeId].m_minIndices[3] = nodeId;
-				m_data->m_broadphaseGPU->m_allAabbsCPU[nodeId].m_maxVec = position+halfExtents;
+				m_data->m_broadphaseGPU->getAllAabbsCPU()[nodeId].m_maxVec = position+halfExtents;
-				m_data->m_broadphaseGPU->m_allAabbsCPU[nodeId].m_signedMaxIndices[3]= nodeId;		
+				m_data->m_broadphaseGPU->getAllAabbsCPU()[nodeId].m_signedMaxIndices[3]= nodeId;		
 			}
 		}
 		m_data->m_broadphaseGPU->writeAabbsToGpu();
@@ -506,7 +510,7 @@ void PairBench::clientMoveAndDisplay()
 			
 	if (m_data->m_gui)
 	{
-		int allAabbs = m_data->m_broadphaseGPU->m_allAabbsCPU.size();
+		int allAabbs = m_data->m_broadphaseGPU->getAllAabbsCPU().size();
 		int numOverlap = m_data->m_broadphaseGPU->getNumOverlap();
 
 		float time = dt/1000.f;

Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp

@@ -4,6 +4,8 @@
 #include "Bullet3Common/b3Quaternion.h"
 #include "OpenGLWindow/b3gWindowInterface.h"
 #include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h"
+#include "Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h"
+
 #include "../GpuDemoInternalData.h"
 #include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
 #include "OpenGLWindow/OpenGLInclude.h"
@@ -119,6 +121,7 @@ void	GpuRigidBodyDemo::initPhysics(const ConstructionInfo& ci)
 
 		b3GpuNarrowPhase* np = new b3GpuNarrowPhase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue,m_data->m_config);
 		b3GpuSapBroadphase* bp = new b3GpuSapBroadphase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue);
+		//b3GpuBroadphaseInterface* bp = new b3GpuGridBroadphase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue);
 		m_data->m_np = np;
 		m_data->m_bp = bp;
 		m_data->m_broadphaseDbvt = new b3DynamicBvhBroadphase(m_data->m_config.m_maxConvexBodies);
@@ -419,6 +422,8 @@ bool	GpuRigidBodyDemo::mouseButtonCallback(int button, int state, float x, float
 						b3Quaternion orn(0,0,0,1);
 						int fixedSphere = m_data->m_np->registerConvexHullShape(0,0,0,0);//>registerSphereShape(0.1);
 						m_data->m_pickFixedBody = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0,pos,orn,fixedSphere,0,false);
+						m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
+						m_data->m_bp->writeAabbsToGpu();
 					
 						if (m_data->m_pickGraphicsShapeIndex<0)
 						{

Demos3/GpuDemos/rigidbody/GpuRigidBodyDemoInternalData.h

@@ -16,7 +16,7 @@ struct	GpuRigidBodyDemoInternalData
 	class b3GpuRigidBodyPipeline* m_rigidBodyPipeline;
 
 	class b3GpuNarrowPhase* m_np;
-	class b3GpuSapBroadphase* m_bp;
+	class b3GpuBroadphaseInterface* m_bp;
 	class b3DynamicBvhBroadphase* m_broadphaseDbvt;
 
 	b3Vector3 m_pickPivotInA;

Demos3/ImplicitCloth/premake4.lua

@@ -8,6 +8,7 @@
 
   	includedirs {
                 ".",
+                "../../src",
                 "../../btgui"
                 }
 
@@ -39,8 +40,8 @@
 		"../../btgui/OpenGLTrueTypeFont/fontstash.h",
 		"../../btgui/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp",
  		"../../btgui/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h",
- 		"../../btgui/Bullet3Common/**.cpp",
+ 		"../../src/Bullet3Common/**.cpp",
- 		"../../btgui/Bullet3Common/**.h",
+ 		"../../src/Bullet3Common/**.h",
 		"../../btgui/Timing/b3Clock.cpp",
 		"../../btgui/Timing/b3Clock.h"
 

Demos3/SimpleOpenGL3/premake4.lua

@@ -8,6 +8,7 @@
 
   	includedirs {
                 ".",
+                "../../src",
                 "../../btgui"
                 }
 
@@ -19,6 +20,8 @@
 		files {
 		"**.cpp",
 		"**.h",
+		"../../src/Bullet3Common/**.cpp",
+ 		"../../src/Bullet3Common/**.h",
 		"../../btgui/OpenGLWindow/SimpleOpenGL3App.cpp",
 		"../../btgui/OpenGLWindow/SimpleOpenGL3App.h",
 		"../../btgui/OpenGLWindow/TwFonts.cpp",

src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h

@@ -0,0 +1,39 @@
+
+#ifndef B3_GPU_BROADPHASE_INTERFACE_H
+#define B3_GPU_BROADPHASE_INTERFACE_H
+
+#include "Bullet3OpenCL/Initialize/b3OpenCLInclude.h"
+#include "Bullet3Common/b3Vector3.h"
+#include "b3SapAabb.h"
+#include "Bullet3Common/shared/b3Int2.h"
+#include "Bullet3Common/shared/b3Int4.h"
+#include "Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h"
+
+class b3GpuBroadphaseInterface
+{
+public:
+
+
+	virtual ~b3GpuBroadphaseInterface()
+	{
+	}
+
+	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  calculateOverlappingPairs(int maxPairs)=0;
+	virtual void  calculateOverlappingPairsHost(int maxPairs)=0;
+
+	//call writeAabbsToGpu after done making all changes (createProxy etc)
+	virtual void writeAabbsToGpu()=0;
+
+	virtual cl_mem	getAabbBufferWS()=0;
+	virtual int	getNumOverlap()=0;
+	virtual cl_mem	getOverlappingPairBuffer()=0;
+
+	virtual b3OpenCLArray<b3SapAabb>&	getAllAabbsGPU()=0;
+	virtual b3AlignedObjectArray<b3SapAabb>&	getAllAabbsCPU()=0;
+
+};
+
+#endif //B3_GPU_BROADPHASE_INTERFACE_H

src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp

@@ -0,0 +1,98 @@
+
+#include "b3GpuGridBroadphase.h"
+#include "Bullet3Geometry/b3AabbUtil.h"
+
+b3GpuGridBroadphase::b3GpuGridBroadphase(cl_context ctx,cl_device_id device, cl_command_queue  q )
+:m_context(ctx),
+m_device(device),
+m_queue(q),
+m_allAabbsGPU(ctx,q),
+m_gpuPairs(ctx,q)
+{
+}
+b3GpuGridBroadphase::~b3GpuGridBroadphase()
+{
+}
+
+
+
+void b3GpuGridBroadphase::createProxy(const b3Vector3& aabbMin,  const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask)
+{
+	b3SapAabb aabb;
+	aabb.m_minVec = aabbMin;
+	aabb.m_maxVec = aabbMax;
+	aabb.m_minIndices[3] = userPtr;
+	aabb.m_signedMaxIndices[3] = userPtr;
+	m_allAabbsCPU.push_back(aabb);
+}
+void b3GpuGridBroadphase::createLargeProxy(const b3Vector3& aabbMin,  const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask)
+{
+	createProxy(aabbMin,aabbMax,userPtr,collisionFilterGroup,collisionFilterMask);
+	
+}
+
+void  b3GpuGridBroadphase::calculateOverlappingPairs(int maxPairs)
+{
+	calculateOverlappingPairsHost(maxPairs);
+}
+void  b3GpuGridBroadphase::calculateOverlappingPairsHost(int maxPairs)
+{
+	m_hostPairs.resize(0);
+		
+	for (int i=0;i<m_allAabbsCPU.size();i++)
+	{
+		for (int j=i+1;j<m_allAabbsCPU.size();j++)
+		{
+			if (b3TestAabbAgainstAabb2(m_allAabbsCPU[i].m_minVec, m_allAabbsCPU[i].m_maxVec,
+				m_allAabbsCPU[j].m_minVec,m_allAabbsCPU[j].m_maxVec))
+			{
+				b3Int4 pair;
+				int a = m_allAabbsCPU[j].m_minIndices[3];
+				int b = m_allAabbsCPU[i].m_minIndices[3];
+				if (a<=b)
+				{
+					pair.x = a; 
+					pair.y = b;//store the original index in the unsorted aabb array
+				} else
+				{
+					pair.x = b;
+					pair.y = a;//store the original index in the unsorted aabb array
+				}
+					
+				m_hostPairs.push_back(pair);
+			}
+		}
+	}
+
+
+	m_gpuPairs.copyFromHost(m_hostPairs);
+}
+
+	//call writeAabbsToGpu after done making all changes (createProxy etc)
+void b3GpuGridBroadphase::writeAabbsToGpu()
+{
+	m_allAabbsGPU.copyFromHost(m_allAabbsCPU);
+}
+
+cl_mem	b3GpuGridBroadphase::getAabbBufferWS()
+{
+	return this->m_allAabbsGPU.getBufferCL();
+}
+int	b3GpuGridBroadphase::getNumOverlap()
+{
+	return m_gpuPairs.size();
+}
+cl_mem	b3GpuGridBroadphase::getOverlappingPairBuffer()
+{
+	return m_gpuPairs.getBufferCL();
+}
+
+b3OpenCLArray<b3SapAabb>&	b3GpuGridBroadphase::getAllAabbsGPU()
+{
+	return m_allAabbsGPU;
+}
+
+b3AlignedObjectArray<b3SapAabb>&	b3GpuGridBroadphase::getAllAabbsCPU()
+{
+	return m_allAabbsCPU;
+}

src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h

@@ -0,0 +1,44 @@
+#ifndef B3_GPU_GRID_BROADPHASE_H
+#define B3_GPU_GRID_BROADPHASE_H
+
+#include "b3GpuBroadphaseInterface.h"
+
+class b3GpuGridBroadphase : public b3GpuBroadphaseInterface
+{
+protected:
+	cl_context				m_context;
+	cl_device_id			m_device;
+	cl_command_queue		m_queue;
+
+	b3OpenCLArray<b3SapAabb>	m_allAabbsGPU;
+	b3AlignedObjectArray<b3SapAabb>	m_allAabbsCPU;
+
+	b3AlignedObjectArray<b3Int4> m_hostPairs;
+	b3OpenCLArray<b3Int4>			m_gpuPairs;
+
+public:
+
+	b3GpuGridBroadphase(cl_context ctx,cl_device_id device, cl_command_queue  q );
+	virtual ~b3GpuGridBroadphase();
+
+
+
+	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  calculateOverlappingPairs(int maxPairs);
+	virtual void  calculateOverlappingPairsHost(int maxPairs);
+
+	//call writeAabbsToGpu after done making all changes (createProxy etc)
+	virtual void writeAabbsToGpu();
+
+	virtual cl_mem	getAabbBufferWS();
+	virtual int	getNumOverlap();
+	virtual cl_mem	getOverlappingPairBuffer();
+
+	virtual b3OpenCLArray<b3SapAabb>&	getAllAabbsGPU();
+	virtual b3AlignedObjectArray<b3SapAabb>&	getAllAabbsCPU();
+
+};
+
+#endif //B3_GPU_GRID_BROADPHASE_H

src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h

@@ -9,8 +9,9 @@ class b3Vector3;
 #include "b3SapAabb.h"
 #include "Bullet3Common/shared/b3Int2.h"
 
+#include "b3GpuBroadphaseInterface.h"
 
-class b3GpuSapBroadphase
+class b3GpuSapBroadphase : public b3GpuBroadphaseInterface
 {
 	
 	cl_context				m_context;
@@ -58,6 +59,15 @@ class b3GpuSapBroadphase
 	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;
@@ -79,23 +89,23 @@ class b3GpuSapBroadphase
 	b3GpuSapBroadphase(cl_context ctx,cl_device_id device, cl_command_queue  q );
 	virtual ~b3GpuSapBroadphase();
 	
-	void  calculateOverlappingPairs(int maxPairs);
+	virtual void  calculateOverlappingPairs(int maxPairs);
-	void  calculateOverlappingPairsHost(int maxPairs);
+	virtual void  calculateOverlappingPairsHost(int maxPairs);
 	
 	void  reset();
 
 	void init3dSap();
-	void calculateOverlappingPairsHostIncremental3Sap();
+	virtual void calculateOverlappingPairsHostIncremental3Sap();
 
-	void createProxy(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 ,short int collisionFilterGroup,short int collisionFilterMask);
-	void createLargeProxy(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);
 
 	//call writeAabbsToGpu after done making all changes (createProxy etc)
-	void writeAabbsToGpu();
+	virtual void writeAabbsToGpu();
 
-	cl_mem	getAabbBufferWS();
+	virtual cl_mem	getAabbBufferWS();
-	int	getNumOverlap();
+	virtual int	getNumOverlap();
-	cl_mem	getOverlappingPairBuffer();
+	virtual cl_mem	getOverlappingPairBuffer();
 };
 
 #endif //B3_GPU_SAP_BROADPHASE_H

src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp

@@ -22,7 +22,7 @@ subject to the following restrictions:
 #include "b3GpuNarrowPhase.h"
 #include "Bullet3Geometry/b3AabbUtil.h"
 #include "Bullet3OpenCL/BroadphaseCollision/b3SapAabb.h"
-#include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h"
+#include "Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h"
 #include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h"
 #include "Bullet3Dynamics/ConstraintSolver/b3PgsJacobiSolver.h"
 #include "Bullet3Collision/NarrowPhaseCollision/shared/b3UpdateAabbs.h"
@@ -43,7 +43,7 @@ subject to the following restrictions:
 	bool integrateOnCpu = true;
 
 #else
-	bool useDbvt = false;//true;
+	bool useDbvt = false;
 	bool useBullet2CpuSolver = true;
 	bool dumpContactStats = false;
 	bool calcWorldSpaceAabbOnCpu = false;//true;
@@ -70,7 +70,7 @@ subject to the following restrictions:
 #include "Bullet3Dynamics/shared/b3IntegrateTransforms.h"
 #include "Bullet3OpenCL/RigidBody/b3GpuNarrowPhaseInternalData.h"
 
-b3GpuRigidBodyPipeline::b3GpuRigidBodyPipeline(cl_context ctx,cl_device_id device, cl_command_queue  q,class b3GpuNarrowPhase* narrowphase, class b3GpuSapBroadphase* broadphaseSap , struct b3DynamicBvhBroadphase* broadphaseDbvt, const b3Config& config)
+b3GpuRigidBodyPipeline::b3GpuRigidBodyPipeline(cl_context ctx,cl_device_id device, cl_command_queue  q,class b3GpuNarrowPhase* narrowphase, class b3GpuBroadphaseInterface* broadphaseSap , struct b3DynamicBvhBroadphase* broadphaseDbvt, const b3Config& config)
 {
 	m_data = new b3GpuRigidBodyPipelineInternalData;
 	m_data->m_constraintUid=0;
@@ -497,13 +497,14 @@ void	b3GpuRigidBodyPipeline::setupGpuAabbsFull()
 				m_data->m_allAabbsGPU->copyFromHost(m_data->m_allAabbsCPU);
 			} else
 			{
-				m_data->m_broadphaseSap->m_allAabbsCPU.resize(numBodies);
+				m_data->m_broadphaseSap->getAllAabbsCPU().resize(numBodies);
 				m_data->m_narrowphase->readbackAllBodiesToCpu();
 				for (int i=0;i<numBodies;i++)
 				{
-					b3ComputeWorldAabb(  i, m_data->m_narrowphase->getBodiesCpu(), m_data->m_narrowphase->getCollidablesCpu(), m_data->m_narrowphase->getLocalSpaceAabbsCpu(),&m_data->m_broadphaseSap->m_allAabbsCPU[0]);
+					b3ComputeWorldAabb(  i, m_data->m_narrowphase->getBodiesCpu(), m_data->m_narrowphase->getCollidablesCpu(), m_data->m_narrowphase->getLocalSpaceAabbsCpu(),&m_data->m_broadphaseSap->getAllAabbsCPU()[0]);
 				}
-				m_data->m_broadphaseSap->m_allAabbsGPU.copyFromHost(m_data->m_broadphaseSap->m_allAabbsCPU);
+				m_data->m_broadphaseSap->getAllAabbsGPU().copyFromHost(m_data->m_broadphaseSap->getAllAabbsCPU());
+				//m_data->m_broadphaseSap->writeAabbsToGpu();
 			}
 		}
 	} else
@@ -631,10 +632,10 @@ int		b3GpuRigidBodyPipeline::registerPhysicsInstance(float mass, const float* po
 		{
 			if (mass)
 			{
-				m_data->m_broadphaseSap->createProxy(aabbMin,aabbMax,userIndex,1,1);//m_dispatcher);
+				m_data->m_broadphaseSap->createProxy(aabbMin,aabbMax,bodyIndex,1,1);//m_dispatcher);
 			} else
 			{
-				m_data->m_broadphaseSap->createLargeProxy(aabbMin,aabbMax,userIndex,1,1);//m_dispatcher);	
+				m_data->m_broadphaseSap->createLargeProxy(aabbMin,aabbMax,bodyIndex,1,1);//m_dispatcher);	
 			}
 		}
 	}

src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.h

@@ -32,7 +32,7 @@ protected:
 public:
 
 
-	b3GpuRigidBodyPipeline(cl_context ctx,cl_device_id device, cl_command_queue  q , class b3GpuNarrowPhase* narrowphase, class b3GpuSapBroadphase* broadphaseSap, struct b3DynamicBvhBroadphase* broadphaseDbvt, const b3Config& config);
+	b3GpuRigidBodyPipeline(cl_context ctx,cl_device_id device, cl_command_queue  q , class b3GpuNarrowPhase* narrowphase, class b3GpuBroadphaseInterface* broadphaseSap, struct b3DynamicBvhBroadphase* broadphaseDbvt, const b3Config& config);
 	virtual ~b3GpuRigidBodyPipeline();
 
 	void	stepSimulation(float deltaTime);

src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipelineInternalData.h

@@ -50,7 +50,7 @@ struct b3GpuRigidBodyPipelineInternalData
 	class b3GpuJacobiSolver* m_solver3;
 	class b3GpuRaycast* m_raycaster;
 	
-	class b3GpuSapBroadphase* m_broadphaseSap;
+	class b3GpuBroadphaseInterface* m_broadphaseSap;
 	
 	struct b3DynamicBvhBroadphase* m_broadphaseDbvt;
 	b3OpenCLArray<b3SapAabb>*	m_allAabbsGPU;