added gpudemo

build/premake4.lua

@@ -97,7 +97,7 @@
 		include "../btgui/GwenOpenGLTest"
 		include "../btgui/OpenGLTrueTypeFont"
 		include "../btgui/OpenGLWindow"
-		
+		include "../demo/gpudemo"
 		
 		
 		

demo/donttouch/btCpuDynamicsWorld.cpp

@@ -0,0 +1,17 @@
+#include "btCpuDynamicsWorld.h"
+
+#include "btBulletDynamicsCommon.h"
+
+btCpuDynamicsWorld::btCpuDynamicsWorld()
+	:btDiscreteDynamicsWorld(
+			new btCollisionDispatcher(new btDefaultCollisionConfiguration()),
+			new btDbvtBroadphase(),new btSequentialImpulseConstraintSolver(),
+			new btDefaultCollisionConfiguration()//todo: remove this!
+			)
+{
+}
+	
+btCpuDynamicsWorld::~btCpuDynamicsWorld()
+{
+	
+}

demo/donttouch/btCpuDynamicsWorld.h

@@ -0,0 +1,24 @@
+
+#ifndef BT_CPU_DYNAMICS_WORLD_H
+#define BT_CPU_DYNAMICS_WORLD_H
+
+class btDefaultCollisionConfiguration;
+class btCollisionDispatcher;
+struct btDbvtBroadphase;
+class btSequentialImpulseConstraintSolver;
+
+#include "BulletDynamics/Dynamics/btDiscreteDynamicsWorld.h"
+
+class btCpuDynamicsWorld : public btDiscreteDynamicsWorld
+{
+	
+public:
+	
+	btCpuDynamicsWorld();
+	
+	virtual ~btCpuDynamicsWorld();
+
+
+};
+
+#endif //BT_CPU_DYNAMICS_WORLD_H

demo/donttouch/btGpuDynamicsWorld.cpp

@@ -0,0 +1,300 @@
+#include "btGpuDynamicsWorld.h"
+#include "BulletDynamics/Dynamics/btRigidBody.h"
+
+#include "../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.h"
+#include "../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.h"
+#include "BulletCollision/CollisionShapes/btPolyhedralConvexShape.h"
+#include "BulletCollision/CollisionShapes/btBvhTriangleMeshShape.h"
+#include "BulletCollision/CollisionShapes/btCompoundShape.h"
+#include "BulletCollision/CollisionShapes/btSphereShape.h"
+#include "BulletCollision/CollisionShapes/btStaticPlaneShape.h"
+
+#include "LinearMath/btQuickprof.h"
+
+
+#ifdef _WIN32
+	#include <wiNdOws.h>
+#endif
+
+
+
+btGpuDynamicsWorld::btGpuDynamicsWorld(int preferredOpenCLPlatformIndex,int preferredOpenCLDeviceIndex)
+:btDynamicsWorld(0,0,0),
+m_gravity(0,-10,0),
+m_once(true)
+{
+	m_gpuPhysics = new CLPhysicsDemo(512*1024, MAX_CONVEX_BODIES_CL);
+	bool useInterop = false;
+	///platform and device are swapped, todo: fix this and make it consistent
+	m_gpuPhysics->init(preferredOpenCLDeviceIndex,preferredOpenCLPlatformIndex,useInterop);
+}
+
+btGpuDynamicsWorld::~btGpuDynamicsWorld()
+{
+	delete m_gpuPhysics;
+}
+
+void btGpuDynamicsWorld::exitOpenCL()
+{
+}
+
+
+
+
+
+
+int		btGpuDynamicsWorld::stepSimulation( btScalar timeStep,int maxSubSteps, btScalar fixedTimeStep)
+{
+#ifndef BT_NO_PROFILE
+//	CProfileManager::Reset();
+#endif //BT_NO_PROFILE
+
+	BT_PROFILE("stepSimulation");
+
+	//convert all shapes now, and if any change, reset all (todo)
+	
+	if (m_once)
+	{
+		m_once = false;
+		m_gpuPhysics->writeBodiesToGpu();
+	}
+
+	m_gpuPhysics->stepSimulation();
+
+	{
+		{
+			BT_PROFILE("readbackBodiesToCpu");
+			//now copy info back to rigid bodies....
+			m_gpuPhysics->readbackBodiesToCpu();
+		}
+
+		
+		{
+			BT_PROFILE("scatter transforms into rigidbody (CPU)");
+			for (int i=0;i<this->m_collisionObjects.size();i++)
+			{
+				btVector3 pos;
+				btQuaternion orn;
+				m_gpuPhysics->getObjectTransformFromCpu(&pos[0],&orn[0],i);
+				btTransform newTrans;
+				newTrans.setOrigin(pos);
+				newTrans.setRotation(orn);
+				this->m_collisionObjects[i]->setWorldTransform(newTrans);
+			}
+		}
+	}
+
+
+#ifndef BT_NO_PROFILE
+	//CProfileManager::Increment_Frame_Counter();
+#endif //BT_NO_PROFILE
+
+
+	return 1;
+}
+
+
+void	btGpuDynamicsWorld::setGravity(const btVector3& gravity)
+{
+}
+
+int btGpuDynamicsWorld::findOrRegisterCollisionShape(const btCollisionShape* colShape)
+{
+	int index = m_uniqueShapes.findLinearSearch(colShape);
+	if (index==m_uniqueShapes.size())
+	{
+		if (colShape->isPolyhedral())
+		{
+			m_uniqueShapes.push_back(colShape);
+			
+			btPolyhedralConvexShape* convex = (btPolyhedralConvexShape*)colShape;
+			int numVertices=convex->getNumVertices();
+			
+			int strideInBytes=sizeof(btVector3);
+			btAlignedObjectArray<btVector3> tmpVertices;
+			tmpVertices.resize(numVertices);
+			for (int i=0;i<numVertices;i++)
+				convex->getVertex(i,tmpVertices[i]);
+			const float scaling[4]={1,1,1,1};
+			bool noHeightField=true;
+			
+			int gpuShapeIndex = m_gpuPhysics->registerConvexPolyhedron(&tmpVertices[0].getX(), strideInBytes, numVertices, scaling, noHeightField);
+			m_uniqueShapeMapping.push_back(gpuShapeIndex);
+		} else
+		{
+			if (colShape->getShapeType()==TRIANGLE_MESH_SHAPE_PROXYTYPE)
+			{
+				m_uniqueShapes.push_back(colShape);
+				
+				btBvhTriangleMeshShape* trimesh = (btBvhTriangleMeshShape*) colShape;
+				btStridingMeshInterface* meshInterface = trimesh->getMeshInterface();
+				btAlignedObjectArray<btVector3> vertices;
+				btAlignedObjectArray<int> indices;
+				
+				btVector3 trimeshScaling(1,1,1);
+				for (int partId=0;partId<meshInterface->getNumSubParts();partId++)
+				{
+					
+					const unsigned char *vertexbase = 0;
+					int numverts = 0;
+					PHY_ScalarType type = PHY_INTEGER;
+					int stride = 0;
+					const unsigned char *indexbase = 0;
+					int indexstride = 0;
+					int numfaces = 0;
+					PHY_ScalarType indicestype = PHY_INTEGER;
+					//PHY_ScalarType indexType=0;
+					
+					btVector3 triangleVerts[3];
+					meshInterface->getLockedReadOnlyVertexIndexBase(&vertexbase,numverts,	type,stride,&indexbase,indexstride,numfaces,indicestype,partId);
+					btVector3 aabbMin,aabbMax;
+					
+					for (int triangleIndex = 0 ; triangleIndex < numfaces;triangleIndex++)
+					{
+						unsigned int* gfxbase = (unsigned int*)(indexbase+triangleIndex*indexstride);
+						
+						for (int j=2;j>=0;j--)
+						{
+							
+							int graphicsindex = indicestype==PHY_SHORT?((unsigned short*)gfxbase)[j]:gfxbase[j];
+							if (type == PHY_FLOAT)
+							{
+								float* graphicsbase = (float*)(vertexbase+graphicsindex*stride);
+								triangleVerts[j] = btVector3(
+															 graphicsbase[0]*trimeshScaling.getX(),
+															 graphicsbase[1]*trimeshScaling.getY(),
+															 graphicsbase[2]*trimeshScaling.getZ());
+							}
+							else
+							{
+								double* graphicsbase = (double*)(vertexbase+graphicsindex*stride);
+								triangleVerts[j] = btVector3( btScalar(graphicsbase[0]*trimeshScaling.getX()),
+															 btScalar(graphicsbase[1]*trimeshScaling.getY()),
+															 btScalar(graphicsbase[2]*trimeshScaling.getZ()));
+							}
+						}
+						vertices.push_back(triangleVerts[0]);
+						vertices.push_back(triangleVerts[1]);
+						vertices.push_back(triangleVerts[2]);
+						indices.push_back(indices.size());
+						indices.push_back(indices.size());
+						indices.push_back(indices.size());
+					}
+				}
+				//GraphicsShape* gfxShape = 0;//btBulletDataExtractor::createGraphicsShapeFromWavefrontObj(objData);
+				
+				//GraphicsShape* gfxShape = btBulletDataExtractor::createGraphicsShapeFromConvexHull(&sUnitSpherePoints[0],MY_UNITSPHERE_POINTS);
+				float meshScaling[4] = {1,1,1,1};
+				//int shapeIndex = renderer.registerShape(gfxShape->m_vertices,gfxShape->m_numvertices,gfxShape->m_indices,gfxShape->m_numIndices);
+				float groundPos[4] = {0,0,0,0};
+				
+				//renderer.registerGraphicsInstance(shapeIndex,groundPos,rotOrn,color,meshScaling);
+				if (vertices.size() && indices.size())
+				{
+					int gpuShapeIndex = m_gpuPhysics->registerConcaveMesh(&vertices,&indices, meshScaling);
+					m_uniqueShapeMapping.push_back(gpuShapeIndex);
+				} else
+				{
+					printf("Error: no vertices in mesh in btGpuDynamicsWorld::addRigidBody\n");
+					index = -1;
+					btAssert(0);
+				}
+				
+				
+			} else
+			{
+				if (colShape->getShapeType()==COMPOUND_SHAPE_PROXYTYPE)
+				{
+					
+					btCompoundShape* compound = (btCompoundShape*) colShape;
+					btAlignedObjectArray<btGpuChildShape> childShapes;
+					
+					for (int i=0;i<compound->getNumChildShapes();i++)
+					{
+						//for now, only support polyhedral child shapes
+						btAssert(compound->getChildShape(i)->isPolyhedral());
+						btGpuChildShape child;
+						child.m_shapeIndex = findOrRegisterCollisionShape(compound->getChildShape(i));
+						btVector3 pos = compound->getChildTransform(i).getOrigin();
+						btQuaternion orn = compound->getChildTransform(i).getRotation();
+						for (int v=0;v<4;v++)
+						{
+							child.m_childPosition[v] = pos[v];
+							child.m_childOrientation[v] = orn[v];
+						}
+						childShapes.push_back(child);
+					}
+					index = m_uniqueShapes.size();
+					m_uniqueShapes.push_back(colShape);
+					
+					int gpuShapeIndex = m_gpuPhysics->registerCompoundShape(&childShapes);
+					m_uniqueShapeMapping.push_back(gpuShapeIndex);
+
+					
+					
+					
+					/*printf("Error: unsupported compound type (%d) in btGpuDynamicsWorld::addRigidBody\n",colShape->getShapeType());
+					index = -1;
+					btAssert(0);
+					 */
+				} else
+				{
+					if (colShape->getShapeType()==SPHERE_SHAPE_PROXYTYPE)
+					{
+						m_uniqueShapes.push_back(colShape);
+						btSphereShape* sphere = (btSphereShape*)colShape;
+						
+						int gpuShapeIndex = m_gpuPhysics->registerSphereShape(sphere->getRadius());
+						m_uniqueShapeMapping.push_back(gpuShapeIndex);
+					} else
+					{
+						if (colShape->getShapeType()==STATIC_PLANE_PROXYTYPE)
+						{
+							m_uniqueShapes.push_back(colShape);
+							btStaticPlaneShape* plane = (btStaticPlaneShape*)colShape;
+						
+							int gpuShapeIndex = m_gpuPhysics->registerPlaneShape(plane->getPlaneNormal(),plane->getPlaneConstant());
+							m_uniqueShapeMapping.push_back(gpuShapeIndex);
+						} else
+						{
+							printf("Error: unsupported shape type (%d) in btGpuDynamicsWorld::addRigidBody\n",colShape->getShapeType());
+							index = -1;
+							btAssert(0);
+						}
+					}
+				}
+			}
+		}
+		
+	}
+	
+	return index;
+}
+
+void	btGpuDynamicsWorld::addRigidBody(btRigidBody* body)
+{
+
+	body->setMotionState(0);
+	
+
+	int index = findOrRegisterCollisionShape(body->getCollisionShape());
+
+	if (index>=0)
+	{
+		int gpuShapeIndex= m_uniqueShapeMapping[index];
+		float mass = body->getInvMass() ? 1.f/body->getInvMass() : 0.f;
+		btVector3 pos = body->getWorldTransform().getOrigin();
+		btQuaternion orn = body->getWorldTransform().getRotation();
+	
+		m_gpuPhysics->registerPhysicsInstance(mass,&pos.getX(),&orn.getX(),gpuShapeIndex,m_collisionObjects.size());
+
+		m_collisionObjects.push_back(body);
+	}
+}
+
+void	btGpuDynamicsWorld::removeCollisionObject(btCollisionObject* colObj)
+{
+	btDynamicsWorld::removeCollisionObject(colObj);
+}
+
+

demo/donttouch/btGpuDynamicsWorld.h

@@ -0,0 +1,109 @@
+#ifndef BT_GPU_DYNAMICS_WORLD_H
+#define BT_GPU_DYNAMICS_WORLD_H
+
+class btVector3;
+class btRigidBody;
+class btCollisionObject;
+struct btGpuInternalData;//use this struct to avoid 'leaking' all OpenCL headers into clients code base
+class CLPhysicsDemo;
+
+#include "BulletCommon/btAlignedObjectArray.h"
+//#include "BulletDynamics/Dynamics/btDynamicsWorld.h"
+
+
+class btGpuDynamicsWorld //: public btDynamicsWorld
+{
+	
+	btAlignedObjectArray<const class  btCollisionShape*> m_uniqueShapes;
+	btAlignedObjectArray<int> m_uniqueShapeMapping;
+
+
+	CLPhysicsDemo*		m_gpuPhysics;
+	btVector3			m_gravity;
+	bool	m_once;
+	
+	bool initOpenCL(int preferredDeviceIndex, int preferredPlatformIndex, bool useInterop);
+	void exitOpenCL();
+	
+	int findOrRegisterCollisionShape(const btCollisionShape* colShape);
+
+	
+public:
+	btGpuDynamicsWorld(int preferredOpenCLPlatformIndex,int preferredOpenCLDeviceIndex);
+
+	virtual ~btGpuDynamicsWorld();
+
+	virtual int		stepSimulation( btScalar timeStep,int maxSubSteps=1, btScalar fixedTimeStep=btScalar(1.)/btScalar(60.));
+
+	virtual void	synchronizeMotionStates()
+	{
+		btAssert(0);
+	}
+
+	void	debugDrawWorld() {}
+
+	void	setGravity(const btVector3& gravity);
+
+	void	addRigidBody(btRigidBody* body);
+
+	void	removeCollisionObject(btCollisionObject* colObj);
+
+	
+
+	btAlignedObjectArray<class btCollisionObject*>& getCollisionObjectArray();
+
+	const btAlignedObjectArray<class btCollisionObject*>& getCollisionObjectArray() const;
+
+	virtual void	addAction(btActionInterface* action)
+	{
+		btAssert(0);
+	}
+
+	virtual void	removeAction(btActionInterface* action)
+	{
+		btAssert(0);
+	}
+
+
+	btVector3 getGravity () const
+	{
+		return m_gravity;
+	}
+
+		virtual void	addRigidBody(btRigidBody* body, short group, short mask)
+		{
+			addRigidBody(body);
+		}
+
+		virtual void	removeRigidBody(btRigidBody* body)
+		{
+			btAssert(0);
+		}
+
+		virtual void	setConstraintSolver(btConstraintSolver* solver)
+		{
+			btAssert(0);
+		}
+
+		virtual btConstraintSolver* getConstraintSolver()
+		{
+			btAssert(0);
+			return 0;
+		}
+
+		virtual btDynamicsWorldType	getWorldType() const
+		{
+			return BT_GPU_PHYSICS_WORLD;
+		}
+
+		virtual void	clearForces()
+		{
+			btAssert(0);
+		}
+
+		
+
+};
+
+
+#endif //BT_GPU_DYNAMICS_WORLD_H

demo/gpudemo/AMD/premake4.lua

@@ -0,0 +1,132 @@
+hasCL = findOpenCL_AMD()
+	
+if (hasCL) then
+
+	project "bullet2_gpu_demo_opengl3core"
+
+	initOpenCL_AMD()
+
+	language "C++"
+			
+	kind "ConsoleApp"
+	
+	targetdir "../../../../bin"
+
+		includedirs {
+              "..",
+              "../../../../bullet2",
+              "../../../../rendering/Gwen"
+    }
+	
+
+	links { 
+		"gwen"
+	}
+	
+
+	initOpenGL()
+	initGlew()
+
+	files {
+			"../GpuDemo.cpp",
+			"../GpuDemo.h",
+			"../btGpuDynamicsWorld.cpp",
+			"../btGpuDynamicsWorld.h",
+			"../btCpuDynamicsWorld.cpp",
+			"../btCpuDynamicsWorld.h",
+			"../btGpuIntegrateTransforms.cpp",
+			"../btGpuIntegrateTransforms.h",
+			"../gwenUserInterface.cpp",
+			"../gwenUserInterface.h",
+			"../ParticleDemo.cpp",
+			"../ParticleDemo.h",
+			"../BroadphaseBenchmark.cpp",
+			"../BroadphaseBenchmark.h",
+			
+			"../main_opengl3core.cpp",
+	
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.cpp",
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.h",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.h",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.cpp",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.h",
+			"../../../../opencl/broadphase_benchmark/btFillCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btFillCL.h",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.h",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.cpp",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.h",
+			"../../../../dynamics/basic_demo/Stubs/ChNarrowphase.cpp",
+			"../../../../dynamics/basic_demo/Stubs/Solver.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.h",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.h",
+
+			"../../../../opencl/basic_initialize/btOpenCLUtils.cpp",
+			"../../../../opencl/basic_initialize/btOpenCLUtils.h",
+			"../../../../opencl/basic_initialize/btOpenCLInclude.h",
+
+
+			"../../../DemosCommon/GL_ShapeDrawer.cpp",
+			"../../../DemosCommon/GL_ShapeDrawer.h",
+			"../../../DemosCommon/OpenGL3CoreRenderer.cpp",
+			"../../../DemosCommon/OpenGL3CoreRenderer.h",
+			
+			"../../../../rendering/WavefrontObjLoader/string_extra.cpp",
+			"../../../../rendering/WavefrontObjLoader/string_extra.h",
+			"../../../../rendering/WavefrontObjLoader/objLoader.cpp",
+			"../../../../rendering/WavefrontObjLoader/objLoader.h",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.cpp",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.h",
+			"../../../../rendering/WavefrontObjLoader/list.cpp",
+			"../../../../rendering/WavefrontObjLoader/list.h",
+			
+			"../../../../rendering/rendertest/GLInstancingRenderer.cpp",
+			"../../../../rendering/rendertest/GLInstancingRenderer.h",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.cpp",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.h",
+			"../../../../rendering/rendertest/LoadShader.cpp",
+			"../../../../rendering/rendertest/LoadShader.h",
+			"../../../../rendering/rendertest/TwFonts.cpp",
+			"../../../../rendering/rendertest/TwFonts.h",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.h",									
+	}
+ 	
+	if os.is("Windows") then
+        files {
+                "../../../../rendering/rendertest/Win32OpenGLWindow.cpp",
+                "../../../../rendering/rendertest/Win32OpenGLWindow.h",
+                "../../../../rendering/rendertest/Win32Window.cpp",
+                "../../../../rendering/rendertest/Win32Window.h",
+        }
+        end
+        if os.is("Linux") then
+        files{
+                "../../../../rendering/rendertest/X11OpenGLWindow.cpp",
+                "../../../../rendering/rendertest/X11OpenGLWindow.h",
+
+        }
+        end
+
+end

demo/gpudemo/Apple/premake4.lua

@@ -0,0 +1,266 @@
+
+hasCL = findOpenCL_Apple()
+	
+if (hasCL) then
+	
+	project "bullet2_gpu_demo_opengl2_Apple"
+
+	initOpenCL_Apple()
+
+	language "C++"
+			
+	kind "ConsoleApp"
+	
+	targetdir "../../../../bin"
+
+		includedirs {
+              "..",
+	      "../../../../rendering/rendertest",
+              "../../../../bullet2",
+              "../../../../rendering/Gwen"
+    }
+	
+
+	links { 
+		"BulletSoftBody",
+		"BulletDynamics",
+		"BulletCollision",
+		"LinearMath"
+	}
+	
+
+	initOpenGL()
+	initGlew()
+
+	files {
+			"../GpuDemo.cpp",
+			"../GpuDemo.h",
+			"../btGpuDynamicsWorld.cpp",
+			"../btGpuDynamicsWorld.h",
+			"../btCpuDynamicsWorld.cpp",
+			"../btCpuDynamicsWorld.h",
+			"../btGpuIntegrateTransforms.cpp",
+			"../btGpuIntegrateTransforms.h",
+			"../main_opengl2.cpp",
+			
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.cpp",
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.h",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.h",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.cpp",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.h",
+			"../../../../opencl/broadphase_benchmark/btFillCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btFillCL.h",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.h",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.cpp",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.h",
+			"../../../../dynamics/basic_demo/Stubs/ChNarrowphase.cpp",
+			"../../../../dynamics/basic_demo/Stubs/Solver.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.h",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.h",
+
+			"../../../../opencl/basic_initialize/btOpenCLUtils.cpp",
+			"../../../../opencl/basic_initialize/btOpenCLUtils.h",
+			"../../../../opencl/basic_initialize/btOpenCLInclude.h",
+			
+			
+			
+			"../../../DemosCommon/GL_ShapeDrawer.cpp",
+			"../../../DemosCommon/GL_ShapeDrawer.h",
+			"../../../DemosCommon/OpenGL2Renderer.cpp",
+			"../../../DemosCommon/OpenGL2Renderer.h",
+			
+			"../../../../rendering/WavefrontObjLoader/string_extra.cpp",
+			"../../../../rendering/WavefrontObjLoader/string_extra.h",
+			"../../../../rendering/WavefrontObjLoader/objLoader.cpp",
+			"../../../../rendering/WavefrontObjLoader/objLoader.h",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.cpp",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.h",
+			"../../../../rendering/WavefrontObjLoader/list.cpp",
+			"../../../../rendering/WavefrontObjLoader/list.h",
+			
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.cpp",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.h",
+			"../../../../rendering/rendertest/LoadShader.cpp",
+			"../../../../rendering/rendertest/LoadShader.h",
+			"../../../../rendering/rendertest/TwFonts.cpp",
+			"../../../../rendering/rendertest/TwFonts.h",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.h",													
+	}
+
+	if os.is("windows") then
+		files
+		{
+		    "../../../../rendering/rendertest/Win32OpenGLWindow.cpp",
+                        "../../../../rendering/rendertest/Win32OpenGLWindow.h",
+                        "../../../../rendering/rendertest/Win32Window.cpp",
+                        "../../../../rendering/rendertest/Win32Window.h",
+		}	
+	end
+
+	if os.is("macosx") then
+	files
+	{
+		"../../../../rendering/rendertest/MacOpenGLWindow.mm",
+		"../../../../rendering/rendertest/MacOpenGLWindow.h"
+	}
+	links {"Cocoa.framework"}
+	end	
+
+	project "bullet2_gpu_demo_opengl3core_Apple"
+
+	initOpenCL_Apple()
+
+	language "C++"
+			
+	kind "ConsoleApp"
+	
+	targetdir "../../../../bin"
+
+		includedirs {
+              "..",
+	      "../../../../rendering/rendertest",
+              "../../../../bullet2",
+              "../../../../rendering/Gwen"
+    }
+	
+
+	links { 
+		"BulletSoftBody",
+		"BulletDynamics",
+		"BulletCollision",
+		"LinearMath",
+		"gwen",
+	}
+	
+
+	initOpenGL()
+	initGlew()
+
+	files {
+			"../GpuDemo.cpp",
+			"../GpuDemo.h",
+			"../btGpuDynamicsWorld.cpp",
+			"../btGpuDynamicsWorld.h",
+			"../btCpuDynamicsWorld.cpp",
+			"../btCpuDynamicsWorld.h",
+			"../btGpuIntegrateTransforms.cpp",
+			"../btGpuIntegrateTransforms.h",
+			"../GwenUserInterface.cpp",
+			"../GwenUserInterface.h",		
+			"../ParticleDemo.cpp",
+			"../ParticleDemo.h",
+		"../BroadphaseBenchmark.cpp",
+			"../BroadphaseBenchmark.h",
+			"../main_opengl3core.cpp",
+	
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.cpp",
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.h",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.h",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.cpp",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.h",
+			"../../../../opencl/broadphase_benchmark/btFillCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btFillCL.h",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.h",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.cpp",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.h",
+			"../../../../dynamics/basic_demo/Stubs/ChNarrowphase.cpp",
+			"../../../../dynamics/basic_demo/Stubs/Solver.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.h",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.h",
+
+			"../../../../opencl/basic_initialize/btOpenCLUtils.cpp",
+			"../../../../opencl/basic_initialize/btOpenCLUtils.h",
+			"../../../../opencl/basic_initialize/btOpenCLInclude.h",
+
+
+			"../../../DemosCommon/GL_ShapeDrawer.cpp",
+			"../../../DemosCommon/GL_ShapeDrawer.h",
+			"../../../DemosCommon/OpenGL3CoreRenderer.cpp",
+			"../../../DemosCommon/OpenGL3CoreRenderer.h",
+			
+			"../../../../rendering/WavefrontObjLoader/string_extra.cpp",
+			"../../../../rendering/WavefrontObjLoader/string_extra.h",
+			"../../../../rendering/WavefrontObjLoader/objLoader.cpp",
+			"../../../../rendering/WavefrontObjLoader/objLoader.h",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.cpp",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.h",
+			"../../../../rendering/WavefrontObjLoader/list.cpp",
+			"../../../../rendering/WavefrontObjLoader/list.h",
+			
+			"../../../../rendering/rendertest/GLInstancingRenderer.cpp",
+			"../../../../rendering/rendertest/GLInstancingRenderer.h",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.cpp",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.h",
+			"../../../../rendering/rendertest/LoadShader.cpp",
+			"../../../../rendering/rendertest/LoadShader.h",
+			"../../../../rendering/rendertest/TwFonts.cpp",
+			"../../../../rendering/rendertest/TwFonts.h",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.h",		
+	}
+	
+	if os.is("windows") then
+		files
+		{
+		    "../../../../rendering/rendertest/Win32OpenGLWindow.cpp",
+                        "../../../../rendering/rendertest/Win32OpenGLWindow.h",
+                        "../../../../rendering/rendertest/Win32Window.cpp",
+                        "../../../../rendering/rendertest/Win32Window.h",
+		}	
+	end
+
+	if os.is("macosx") then
+	files
+	{
+		"../../../../rendering/rendertest/MacOpenGLWindow.mm",
+		"../../../../rendering/rendertest/MacOpenGLWindow.h"
+	}
+	links {"Cocoa.framework"}
+	end	
+	
+end

demo/gpudemo/Apple/premake4.lua.bak

@@ -0,0 +1,265 @@
+
+hasCL = findOpenCL_Apple()
+	
+if (hasCL) then
+	
+	project "bullet2_gpu_demo_opengl2_Apple"
+
+	initOpenCL_Apple()
+
+	language "C++"
+			
+	kind "ConsoleApp"
+	
+	targetdir "../../../../bin"
+
+		includedirs {
+              "..",
+	      "../../../../rendering/rendertest",
+              "../../../../bullet2",
+              "../../../../rendering/Gwen"
+    }
+	
+
+	links { 
+		"BulletSoftBody",
+		"BulletDynamics",
+		"BulletCollision",
+		"LinearMath"
+	}
+	
+
+	initOpenGL()
+	initGlew()
+
+	files {
+			"../GpuDemo.cpp",
+			"../GpuDemo.h",
+			"../btGpuDynamicsWorld.cpp",
+			"../btGpuDynamicsWorld.h",
+			"../btCpuDynamicsWorld.cpp",
+			"../btCpuDynamicsWorld.h",
+			"../btGpuIntegrateTransforms.cpp",
+			"../btGpuIntegrateTransforms.h",
+			"../main_opengl2.cpp",
+			
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.cpp",
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.h",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.h",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.cpp",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.h",
+			"../../../../opencl/broadphase_benchmark/btFillCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btFillCL.h",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.h",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.cpp",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.h",
+			"../../../../dynamics/basic_demo/Stubs/ChNarrowphase.cpp",
+			"../../../../dynamics/basic_demo/Stubs/Solver.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.h",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.h",
+
+			"../../../../opencl/basic_initialize/btOpenCLUtils.cpp",
+			"../../../../opencl/basic_initialize/btOpenCLUtils.h",
+			"../../../../opencl/basic_initialize/btOpenCLInclude.h",
+			
+			
+			
+			"../../../DemosCommon/GL_ShapeDrawer.cpp",
+			"../../../DemosCommon/GL_ShapeDrawer.h",
+			"../../../DemosCommon/OpenGL2Renderer.cpp",
+			"../../../DemosCommon/OpenGL2Renderer.h",
+			
+			"../../../../rendering/WavefrontObjLoader/string_extra.cpp",
+			"../../../../rendering/WavefrontObjLoader/string_extra.h",
+			"../../../../rendering/WavefrontObjLoader/objLoader.cpp",
+			"../../../../rendering/WavefrontObjLoader/objLoader.h",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.cpp",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.h",
+			"../../../../rendering/WavefrontObjLoader/list.cpp",
+			"../../../../rendering/WavefrontObjLoader/list.h",
+			
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.cpp",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.h",
+			"../../../../rendering/rendertest/LoadShader.cpp",
+			"../../../../rendering/rendertest/LoadShader.h",
+			"../../../../rendering/rendertest/TwFonts.cpp",
+			"../../../../rendering/rendertest/TwFonts.h",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.h",													
+	}
+
+	if os.is("windows") then
+		files
+		{
+		    "../../../../rendering/rendertest/Win32OpenGLWindow.cpp",
+                        "../../../../rendering/rendertest/Win32OpenGLWindow.h",
+                        "../../../../rendering/rendertest/Win32Window.cpp",
+                        "../../../../rendering/rendertest/Win32Window.h",
+		}	
+	end
+
+	if os.is("macosx") then
+	files
+	{
+		"../../../../rendering/rendertest/MacOpenGLWindow.mm",
+		"../../../../rendering/rendertest/MacOpenGLWindow.h"
+	}
+	links {"Cocoa.framework"}
+	end	
+
+	project "bullet2_gpu_demo_opengl3core_Apple"
+
+	initOpenCL_Apple()
+
+	language "C++"
+			
+	kind "ConsoleApp"
+	
+	targetdir "../../../../bin"
+
+		includedirs {
+              "..",
+	      "../../../../rendering/rendertest",
+              "../../../../bullet2",
+              "../../../../rendering/Gwen"
+    }
+	
+
+	links { 
+		"BulletSoftBody",
+		"BulletDynamics",
+		"BulletCollision",
+		"LinearMath",
+		"gwen",
+	}
+	
+
+	initOpenGL()
+	initGlew()
+
+	files {
+			"../GpuDemo.cpp",
+			"../GpuDemo.h",
+			"../btGpuDynamicsWorld.cpp",
+			"../btGpuDynamicsWorld.h",
+			"../btCpuDynamicsWorld.cpp",
+			"../btCpuDynamicsWorld.h",
+			"../btGpuIntegrateTransforms.cpp",
+			"../btGpuIntegrateTransforms.h",
+			"../GwenUserInterface.cpp",
+			"../GwenUserInterface.h",		
+			"../ParticleDemo.cpp",
+			"../ParticleDemo.h",
+		
+			"../main_opengl3core.cpp",
+	
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.cpp",
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.h",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.h",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.cpp",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.h",
+			"../../../../opencl/broadphase_benchmark/btFillCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btFillCL.h",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.h",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.cpp",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.h",
+			"../../../../dynamics/basic_demo/Stubs/ChNarrowphase.cpp",
+			"../../../../dynamics/basic_demo/Stubs/Solver.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.h",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.h",
+
+			"../../../../opencl/basic_initialize/btOpenCLUtils.cpp",
+			"../../../../opencl/basic_initialize/btOpenCLUtils.h",
+			"../../../../opencl/basic_initialize/btOpenCLInclude.h",
+
+
+			"../../../DemosCommon/GL_ShapeDrawer.cpp",
+			"../../../DemosCommon/GL_ShapeDrawer.h",
+			"../../../DemosCommon/OpenGL3CoreRenderer.cpp",
+			"../../../DemosCommon/OpenGL3CoreRenderer.h",
+			
+			"../../../../rendering/WavefrontObjLoader/string_extra.cpp",
+			"../../../../rendering/WavefrontObjLoader/string_extra.h",
+			"../../../../rendering/WavefrontObjLoader/objLoader.cpp",
+			"../../../../rendering/WavefrontObjLoader/objLoader.h",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.cpp",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.h",
+			"../../../../rendering/WavefrontObjLoader/list.cpp",
+			"../../../../rendering/WavefrontObjLoader/list.h",
+			
+			"../../../../rendering/rendertest/GLInstancingRenderer.cpp",
+			"../../../../rendering/rendertest/GLInstancingRenderer.h",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.cpp",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.h",
+			"../../../../rendering/rendertest/LoadShader.cpp",
+			"../../../../rendering/rendertest/LoadShader.h",
+			"../../../../rendering/rendertest/TwFonts.cpp",
+			"../../../../rendering/rendertest/TwFonts.h",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.h",		
+	}
+	
+	if os.is("windows") then
+		files
+		{
+		    "../../../../rendering/rendertest/Win32OpenGLWindow.cpp",
+                        "../../../../rendering/rendertest/Win32OpenGLWindow.h",
+                        "../../../../rendering/rendertest/Win32Window.cpp",
+                        "../../../../rendering/rendertest/Win32Window.h",
+		}	
+	end
+
+	if os.is("macosx") then
+	files
+	{
+		"../../../../rendering/rendertest/MacOpenGLWindow.mm",
+		"../../../../rendering/rendertest/MacOpenGLWindow.h"
+	}
+	links {"Cocoa.framework"}
+	end	
+	
+end

demo/gpudemo/GpuDemo.cpp

@@ -0,0 +1,611 @@
+/*
+Bullet Continuous Collision Detection and Physics Library
+Copyright (c) 2003-2006 Erwin Coumans  http://continuousphysics.com/Bullet/
+
+This software is provided 'as-is', without any express or implied warranty.
+In no event will the authors be held liable for any damages arising from the use of this software.
+Permission is granted to anyone to use this software for any purpose, 
+including commercial applications, and to alter it and redistribute it freely, 
+subject to the following restrictions:
+
+1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.
+2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.
+3. This notice may not be removed or altered from any source distribution.
+*/
+
+
+#include "btCpuDynamicsWorld.h"
+#include "btGpuDynamicsWorld.h"
+
+
+#define SCALING 1.
+#define START_POS_X -5
+#define START_POS_Y 10
+#define START_POS_Z -3
+
+#include "LinearMath/btVector3.h"
+
+#include "GpuDemo.h"
+//#include "GlutStuff.h"
+///btBulletDynamicsCommon.h is the main Bullet include file, contains most common include files.
+//#include "btBulletDynamicsCommon.h"
+
+
+#include "BulletCollision/CollisionShapes/btTriangleMesh.h"
+#include "BulletCollision/CollisionShapes/btBvhTriangleMeshShape.h"
+#include "BulletCollision/CollisionShapes/btSphereShape.h"
+#include "BulletCollision/CollisionShapes/btConvexHullShape.h"
+#include "BulletCollision/CollisionShapes/btBoxShape.h"
+#include "BulletCollision/CollisionShapes/btCompoundShape.h"
+#include "BulletCollision/CollisionShapes/btStaticPlaneShape.h"
+
+#include "BulletDynamics/Dynamics/btRigidBody.h"
+#include "LinearMath/btDefaultMotionState.h"
+#include "LinearMath/btQuickprof.h"
+
+
+#include <stdio.h> //printf debugging
+
+
+void GpuDemo::clientMoveAndDisplay()
+{
+//	glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); 
+
+	//simple dynamics world doesn't handle fixed-time-stepping
+	float dt = getDeltaTimeInSeconds();
+	
+	///step the simulation
+	if (m_dynamicsWorld)
+	{
+		static bool once = true;
+		if (once)
+		{
+			once=false;
+			btDefaultSerializer*	serializer = new btDefaultSerializer();
+			m_dynamicsWorld->serialize(serializer);
+ 
+			FILE* file = fopen("testFile.bullet","wb");
+			fwrite(serializer->getBufferPointer(),serializer->getCurrentBufferSize(),1, file);
+			fclose(file);
+		}
+
+		m_dynamicsWorld->stepSimulation(dt);
+		static int count=0;
+		count++;
+		if (count==25)
+		{
+			//CProfileManager::dumpAll();
+		}
+	}
+		
+//	renderme(); 
+
+
+	//swapBuffers();
+
+}
+
+
+
+
+
+
+btAlignedObjectArray<btVector3> vertices;
+
+void	EmptyDemo::setupScene(const ConstructionInfo& ci)
+{
+	//empty test
+}
+
+void SpheresDemo::setupScene(const ConstructionInfo& ci)
+{
+	
+	
+	if (1)
+	{
+		btSphereShape* sphere = new btSphereShape(1);
+			m_collisionShapes.push_back(sphere);
+
+		/// Create Dynamic Objects
+		btTransform startTransform;
+		startTransform.setIdentity();
+
+	
+
+		float start_x = START_POS_X - ci.gapX*ci.arraySizeX/2;
+		float start_y = START_POS_Y;
+		float start_z = START_POS_Z - ci.gapZ*ci.arraySizeZ/2;
+
+		for (int k=0;k<ci.arraySizeY;k++)
+		{
+			int sizeX = ci.arraySizeX;
+			int startX = -sizeX/2;
+			float gapX = ci.gapX;
+
+			for (int i=0;i<sizeX;i++)
+			{
+				int sizeZ = ci.arraySizeZ;
+				int startZ = -sizeX/2;
+				float gapZ =ci.gapZ;
+				for(int j = 0;j<sizeZ;j++)
+				{
+					//btCollisionShape* shape = k==0? boxShape : colShape;
+
+					btCollisionShape* shape = sphere;
+
+					
+					btScalar	mass  = 1;
+					if (!ci.m_useConcaveMesh && k==0)
+						mass = k==0? 0.f : 1.f;
+
+					//rigidbody is dynamic if and only if mass is non zero, otherwise static
+					bool isDynamic = (mass != 0.f);
+
+					btVector3 localInertia(0,0,0);
+					if (isDynamic)
+						shape->calculateLocalInertia(mass,localInertia);
+
+					startTransform.setOrigin(SCALING*btVector3(
+										btScalar(gapX*i + start_x),
+										btScalar(ci.gapY*k + start_y),
+										btScalar(gapZ*j + start_z)));
+
+			
+					//using motionstate is recommended, it provides interpolation capabilities, and only synchronizes 'active' objects
+					btDefaultMotionState* myMotionState = new btDefaultMotionState(startTransform);
+					btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,myMotionState,shape,localInertia);
+					btRigidBody* body = new btRigidBody(rbInfo);
+					
+
+					m_dynamicsWorld->addRigidBody(body);
+				}
+			}
+		}
+	}
+
+	{
+		btVector3 planeNormal(0,1,0);
+		btScalar planeConstant=0;
+
+		btCollisionShape* shape = new btStaticPlaneShape(planeNormal,planeConstant);
+		//btBoxShape* plane = new btBoxShape(btVector3(100,1,100));
+		//plane->initializePolyhedralFeatures();
+		//btSphereShape* shape = new btSphereShape(1000);
+
+		btScalar mass(0.);
+
+		//rigidbody is dynamic if and only if mass is non zero, otherwise static
+		bool isDynamic = (mass != 0.f);
+
+		btVector3 localInertia(0,0,0);
+		btTransform groundTransform;
+		groundTransform.setIdentity();
+		groundTransform.setRotation(btQuaternion(btVector3(1,0,0),0.3));
+		groundTransform.setOrigin(btVector3(0,0,0));
+
+		//using motionstate is recommended, it provides interpolation capabilities, and only synchronizes 'active' objects
+		btDefaultMotionState* myMotionState = new btDefaultMotionState(groundTransform);
+		btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,myMotionState,shape,localInertia);
+		btRigidBody* body = new btRigidBody(rbInfo);
+
+		//add the body to the dynamics world
+		m_dynamicsWorld->addRigidBody(body);
+	}
+}
+
+
+
+void	GpuCompoundDemo::setupScene(const ConstructionInfo& ci)
+{
+		btCollisionShape* groundShape =0;
+//	btCollisionShape* groundShape = new btStaticPlaneShape(btVector3(0,1,0),50);
+
+	if (ci.m_useConcaveMesh)
+	{
+		btTriangleMesh* meshInterface = new btTriangleMesh();
+
+		btAlignedObjectArray<btVector3> concaveVertices;
+		concaveVertices.push_back(btVector3(0,-20,0));
+		concaveVertices.push_back(btVector3(80,10,80));
+		concaveVertices.push_back(btVector3(80,10,-80));
+		concaveVertices.push_back(btVector3(-80,10,-80));
+		concaveVertices.push_back(btVector3(-80,10,80));
+
+		meshInterface->addTriangle(concaveVertices[0],concaveVertices[1],concaveVertices[2],true);
+		meshInterface->addTriangle(concaveVertices[0],concaveVertices[2],concaveVertices[3],true);
+		meshInterface->addTriangle(concaveVertices[0],concaveVertices[3],concaveVertices[4],true);
+		meshInterface->addTriangle(concaveVertices[0],concaveVertices[4],concaveVertices[1],true);
+
+#if 0
+		groundShape = new btBvhTriangleMeshShape(meshInterface,true);//btStaticPlaneShape(btVector3(0,1,0),50);
+#else
+		btBoxShape* shape =new btBoxShape(btVector3(btScalar(250.),btScalar(10.),btScalar(250.)));
+		shape->initializePolyhedralFeatures();
+		groundShape = shape;
+#endif
+
+	} else
+	{
+		groundShape  = new btBoxShape(btVector3(btScalar(250.),btScalar(50.),btScalar(250.)));
+	}
+	
+	m_collisionShapes.push_back(groundShape);
+
+	btTransform groundTransform;
+	groundTransform.setIdentity();
+	groundTransform.setOrigin(btVector3(0,0,0));
+
+	//We can also use DemoApplication::localCreateRigidBody, but for clarity it is provided here:
+	if (ci.m_useConcaveMesh)
+	{
+		btScalar mass(0.);
+
+		//rigidbody is dynamic if and only if mass is non zero, otherwise static
+		bool isDynamic = (mass != 0.f);
+
+		btVector3 localInertia(0,0,0);
+		if (isDynamic)
+			groundShape->calculateLocalInertia(mass,localInertia);
+
+		//using motionstate is recommended, it provides interpolation capabilities, and only synchronizes 'active' objects
+		btDefaultMotionState* myMotionState = new btDefaultMotionState(groundTransform);
+		btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,myMotionState,groundShape,localInertia);
+		btRigidBody* body = new btRigidBody(rbInfo);
+
+		//add the body to the dynamics world
+		m_dynamicsWorld->addRigidBody(body);
+	}
+
+
+	{
+		//create a few dynamic rigidbodies
+		// Re-using the same collision is better for memory usage and performance
+
+		//vertices.push_back(btVector3(0,1,0));
+		vertices.push_back(btVector3(1,1,1));
+		vertices.push_back(btVector3(1,1,-1));
+		vertices.push_back(btVector3(-1,1,-1));
+		vertices.push_back(btVector3(-1,1,1));
+		vertices.push_back(btVector3(1,-1,1));
+		vertices.push_back(btVector3(1,-1,-1));
+		vertices.push_back(btVector3(-1,-1,-1));
+		vertices.push_back(btVector3(-1,-1,1));
+		
+#if 0
+		btPolyhedralConvexShape* colShape = new btConvexHullShape(&vertices[0].getX(),vertices.size());
+		colShape->initializePolyhedralFeatures();
+#else
+		btCompoundShape* compoundShape = 0;
+		{
+			btPolyhedralConvexShape* colShape = new btConvexHullShape(&vertices[0].getX(),vertices.size());
+			colShape->initializePolyhedralFeatures();
+			compoundShape = new btCompoundShape();
+			btTransform tr;
+			tr.setIdentity();
+			tr.setOrigin(btVector3(0,-1,0));
+			compoundShape->addChildShape(tr,colShape);
+			tr.setOrigin(btVector3(0,0,2));
+			compoundShape->addChildShape(tr,colShape);
+			tr.setOrigin(btVector3(2,0,0));
+			compoundShape->addChildShape(tr,colShape);
+		}
+		btCollisionShape* colShape = compoundShape;
+#endif
+
+
+
+		btPolyhedralConvexShape* boxShape = new btBoxShape(btVector3(SCALING*1,SCALING*1,SCALING*1));
+		boxShape->initializePolyhedralFeatures();
+		
+
+
+
+		//btCollisionShape* colShape = new btSphereShape(btScalar(1.));
+		m_collisionShapes.push_back(colShape);
+		m_collisionShapes.push_back(boxShape);
+
+		/// Create Dynamic Objects
+		btTransform startTransform;
+		startTransform.setIdentity();
+
+	
+
+		float start_x = START_POS_X - ci.arraySizeX/2;
+		float start_y = START_POS_Y;
+		float start_z = START_POS_Z - ci.arraySizeZ/2;
+
+		for (int k=0;k<ci.arraySizeY;k++)
+		{
+			int sizeX = ci.arraySizeX;
+			if (!ci.m_useConcaveMesh && k==0)
+				sizeX = 50;
+
+			int startX = !ci.m_useConcaveMesh&&k==0? -20 : 0;
+			float gapX = !ci.m_useConcaveMesh&&k==0? 3.05 : ci.gapX;
+			for (int i=0;i<sizeX;i++)
+			{
+				int sizeZ = !ci.m_useConcaveMesh&&k==0? 50 : ci.arraySizeZ;
+				int startZ = (!ci.m_useConcaveMesh)&&k==0? -20 : 0;
+				float gapZ = !ci.m_useConcaveMesh&&k==0? 3.05 : ci.gapZ;
+				for(int j = 0;j<sizeZ;j++)
+				{
+					//btCollisionShape* shape = k==0? boxShape : colShape;
+
+					btCollisionShape* shape = colShape;
+
+					
+					btScalar	mass  = 1;
+					if (!ci.m_useConcaveMesh && k==0)
+						mass = k==0? 0.f : 1.f;
+
+					//rigidbody is dynamic if and only if mass is non zero, otherwise static
+					bool isDynamic = (mass != 0.f);
+
+					btVector3 localInertia(0,0,0);
+					if (isDynamic)
+						shape->calculateLocalInertia(mass,localInertia);
+
+					startTransform.setOrigin(SCALING*btVector3(
+										btScalar(startX+gapX*i + start_x),
+										btScalar(20+ci.gapY*k + start_y),
+										btScalar(startZ+gapZ*j + start_z)));
+
+			
+					//using motionstate is recommended, it provides interpolation capabilities, and only synchronizes 'active' objects
+					btDefaultMotionState* myMotionState = new btDefaultMotionState(startTransform);
+					btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,myMotionState,shape,localInertia);
+					btRigidBody* body = new btRigidBody(rbInfo);
+					
+
+					m_dynamicsWorld->addRigidBody(body);
+				}
+			}
+		}
+	}
+}
+
+
+
+void	GpuBoxDemo::setupScene(const ConstructionInfo& ci)
+{
+		btCollisionShape* groundShape =0;
+//	btCollisionShape* groundShape = new btStaticPlaneShape(btVector3(0,1,0),50);
+
+	if (ci.m_useConcaveMesh)
+	{
+		btTriangleMesh* meshInterface = new btTriangleMesh();
+
+		btAlignedObjectArray<btVector3> concaveVertices;
+		concaveVertices.push_back(btVector3(0,-20,0));
+		concaveVertices.push_back(btVector3(80,10,80));
+		concaveVertices.push_back(btVector3(80,10,-80));
+		concaveVertices.push_back(btVector3(-80,10,-80));
+		concaveVertices.push_back(btVector3(-80,10,80));
+
+		meshInterface->addTriangle(concaveVertices[0],concaveVertices[1],concaveVertices[2],true);
+		meshInterface->addTriangle(concaveVertices[0],concaveVertices[2],concaveVertices[3],true);
+		meshInterface->addTriangle(concaveVertices[0],concaveVertices[3],concaveVertices[4],true);
+		meshInterface->addTriangle(concaveVertices[0],concaveVertices[4],concaveVertices[1],true);
+
+#if 0
+		groundShape = new btBvhTriangleMeshShape(meshInterface,true);//btStaticPlaneShape(btVector3(0,1,0),50);
+#else
+		btBoxShape* shape =new btBoxShape(btVector3(btScalar(250.),btScalar(10.),btScalar(250.)));
+		shape->initializePolyhedralFeatures();
+		groundShape = shape;
+#endif
+
+	} else
+	{
+		groundShape  = new btBoxShape(btVector3(btScalar(250.),btScalar(50.),btScalar(250.)));
+	}
+	
+	m_collisionShapes.push_back(groundShape);
+
+	btTransform groundTransform;
+	groundTransform.setIdentity();
+	groundTransform.setOrigin(btVector3(0,0,0));
+
+	//We can also use DemoApplication::localCreateRigidBody, but for clarity it is provided here:
+	if (ci.m_useConcaveMesh)
+	{
+		btScalar mass(0.);
+
+		//rigidbody is dynamic if and only if mass is non zero, otherwise static
+		bool isDynamic = (mass != 0.f);
+
+		btVector3 localInertia(0,0,0);
+		if (isDynamic)
+			groundShape->calculateLocalInertia(mass,localInertia);
+
+		//using motionstate is recommended, it provides interpolation capabilities, and only synchronizes 'active' objects
+		btDefaultMotionState* myMotionState = new btDefaultMotionState(groundTransform);
+		btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,myMotionState,groundShape,localInertia);
+		btRigidBody* body = new btRigidBody(rbInfo);
+
+		//add the body to the dynamics world
+		m_dynamicsWorld->addRigidBody(body);
+	}
+
+
+	{
+		//create a few dynamic rigidbodies
+		// Re-using the same collision is better for memory usage and performance
+
+		//vertices.push_back(btVector3(0,1,0));
+		vertices.push_back(btVector3(1,1,1));
+		vertices.push_back(btVector3(1,1,-1));
+		vertices.push_back(btVector3(-1,1,-1));
+		vertices.push_back(btVector3(-1,1,1));
+		vertices.push_back(btVector3(1,-1,1));
+		vertices.push_back(btVector3(1,-1,-1));
+		vertices.push_back(btVector3(-1,-1,-1));
+		vertices.push_back(btVector3(-1,-1,1));
+		
+#if 1
+		btPolyhedralConvexShape* colShape = new btConvexHullShape(&vertices[0].getX(),vertices.size());
+		colShape->initializePolyhedralFeatures();
+#else
+		btCompoundShape* compoundShape = 0;
+		{
+			btPolyhedralConvexShape* colShape = new btConvexHullShape(&vertices[0].getX(),vertices.size());
+			colShape->initializePolyhedralFeatures();
+			compoundShape = new btCompoundShape();
+			btTransform tr;
+			tr.setIdentity();
+			tr.setOrigin(btVector3(0,-1,0));
+			compoundShape->addChildShape(tr,colShape);
+			tr.setOrigin(btVector3(0,0,2));
+			compoundShape->addChildShape(tr,colShape);
+			tr.setOrigin(btVector3(2,0,0));
+			compoundShape->addChildShape(tr,colShape);
+		}
+		btCollisionShape* colShape = compoundShape;
+#endif
+
+
+
+		btPolyhedralConvexShape* boxShape = new btBoxShape(btVector3(SCALING*1,SCALING*1,SCALING*1));
+		boxShape->initializePolyhedralFeatures();
+		
+
+
+
+		//btCollisionShape* colShape = new btSphereShape(btScalar(1.));
+		m_collisionShapes.push_back(colShape);
+		m_collisionShapes.push_back(boxShape);
+
+		/// Create Dynamic Objects
+		btTransform startTransform;
+		startTransform.setIdentity();
+
+	
+
+		float start_x = START_POS_X - ci.arraySizeX/2;
+		float start_y = START_POS_Y;
+		float start_z = START_POS_Z - ci.arraySizeZ/2;
+
+		for (int k=0;k<ci.arraySizeY;k++)
+		{
+			int sizeX = ci.arraySizeX;
+			if (!ci.m_useConcaveMesh && k==0)
+				sizeX = 50;
+
+			int startX = !ci.m_useConcaveMesh&&k==0? -20 : 0;
+			float gapX = !ci.m_useConcaveMesh&&k==0? 3.05 : ci.gapX;
+			for (int i=0;i<sizeX;i++)
+			{
+				int sizeZ = !ci.m_useConcaveMesh&&k==0? 50 : ci.arraySizeZ;
+				int startZ = (!ci.m_useConcaveMesh)&&k==0? -20 : 0;
+				float gapZ = !ci.m_useConcaveMesh&&k==0? 3.05 : ci.gapZ;
+				for(int j = 0;j<sizeZ;j++)
+				{
+					//btCollisionShape* shape = k==0? boxShape : colShape;
+
+					btCollisionShape* shape = colShape;
+
+					
+					btScalar	mass  = 1;
+					if (!ci.m_useConcaveMesh && k==0)
+						mass = k==0? 0.f : 1.f;
+
+					//rigidbody is dynamic if and only if mass is non zero, otherwise static
+					bool isDynamic = (mass != 0.f);
+
+					btVector3 localInertia(0,0,0);
+					if (isDynamic)
+						shape->calculateLocalInertia(mass,localInertia);
+
+					startTransform.setOrigin(SCALING*btVector3(
+										btScalar(startX+gapX*i + start_x),
+										btScalar(ci.gapY*(k+0.5) + start_y),
+										btScalar(startZ+gapZ*j + start_z)));
+
+			
+					//using motionstate is recommended, it provides interpolation capabilities, and only synchronizes 'active' objects
+					btDefaultMotionState* myMotionState = new btDefaultMotionState(startTransform);
+					btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,myMotionState,shape,localInertia);
+					btRigidBody* body = new btRigidBody(rbInfo);
+					
+
+					m_dynamicsWorld->addRigidBody(body);
+				}
+			}
+		}
+	}
+}
+
+void	GpuDemo::initPhysics(const ConstructionInfo& ci)
+{
+
+//	setTexturing(true);
+	//setShadows(false);
+
+//	setCameraDistance(btScalar(SCALING*250.));
+
+	///collision configuration contains default setup for memory, collision setup
+	if (ci.useOpenCL)
+	{
+		m_dynamicsWorld = new btGpuDynamicsWorld(ci.preferredOpenCLPlatformIndex,ci.preferredOpenCLDeviceIndex);
+	} else
+	{
+		m_dynamicsWorld = new btCpuDynamicsWorld();
+	}
+
+	
+	m_dynamicsWorld->setGravity(btVector3(0,-10,0));
+
+	///create a few basic rigid bodies
+
+	setupScene(ci);
+
+
+}
+
+/*void	GpuDemo::clientResetScene()
+{
+	exitPhysics();
+	initPhysics();
+}
+	*/
+
+
+void	GpuDemo::exitPhysics()
+{
+
+	//cleanup in the reverse order of creation/initialization
+
+	//remove the rigidbodies from the dynamics world and delete them
+	int i;
+	if (m_dynamicsWorld)
+	{
+		for (i=m_dynamicsWorld->getNumCollisionObjects()-1; i>=0 ;i--)
+		{
+			btCollisionObject* obj = m_dynamicsWorld->getCollisionObjectArray()[i];
+			btRigidBody* body = btRigidBody::upcast(obj);
+			if (body && body->getMotionState())
+			{
+				delete body->getMotionState();
+			}
+			m_dynamicsWorld->removeCollisionObject( obj );
+			delete obj;
+		}
+	}
+
+	//delete collision shapes
+	for (int j=0;j<m_collisionShapes.size();j++)
+	{
+		btCollisionShape* shape = m_collisionShapes[j];
+		delete shape;
+	}
+	m_collisionShapes.clear();
+
+	delete m_dynamicsWorld;
+	m_dynamicsWorld=0;
+	
+
+	
+}
+
+
+

demo/gpudemo/GpuDemo.h

@@ -0,0 +1,184 @@
+/*
+Bullet Continuous Collision Detection and Physics Library
+Copyright (c) 2003-2006 Erwin Coumans  http://continuousphysics.com/Bullet/
+
+This software is provided 'as-is', without any express or implied warranty.
+In no event will the authors be held liable for any damages arising from the use of this software.
+Permission is granted to anyone to use this software for any purpose, 
+including commercial applications, and to alter it and redistribute it freely, 
+subject to the following restrictions:
+
+1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.
+2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.
+3. This notice may not be removed or altered from any source distribution.
+*/
+#ifndef GPU_DEMO_H
+#define GPU_DEMO_H
+
+
+#include "BulletCommon/btAlignedObjectArray.h"
+
+
+class btBroadphaseInterface;
+class btCollisionShape;
+class btOverlappingPairCache;
+class btCollisionDispatcher;
+class btConstraintSolver;
+struct btCollisionAlgorithmCreateFunc;
+class btDefaultCollisionConfiguration;
+class btDynamicsWorld;
+class GLInstancingRenderer;
+
+///GpuDemo is good starting point for learning the code base and porting.
+
+
+
+class GpuDemo
+{
+
+protected:
+
+	btDynamicsWorld*	m_dynamicsWorld;
+
+	//keep the collision shapes, for deletion/cleanup
+	btAlignedObjectArray<btCollisionShape*>	m_collisionShapes;
+
+	float getDeltaTimeInSeconds()
+	{
+		return 1./60.f;
+	}
+public:
+
+	typedef class GpuDemo* (CreateFunc)();
+
+	struct ConstructionInfo
+	{
+		bool useOpenCL;
+		int preferredOpenCLPlatformIndex;
+		int preferredOpenCLDeviceIndex;
+		int arraySizeX;
+		int arraySizeY;
+		int arraySizeZ;
+		bool m_useConcaveMesh;
+		float gapX;
+		float gapY;
+		float gapZ;
+		GLInstancingRenderer*	m_instancingRenderer;
+		ConstructionInfo()
+			:useOpenCL(false),//true),
+			preferredOpenCLPlatformIndex(-1),
+			preferredOpenCLDeviceIndex(-1),
+			arraySizeX(10),
+			arraySizeY(10 ),
+			arraySizeZ(10),
+			m_useConcaveMesh(false),
+			gapX(4.3),
+			gapY(4.0),
+			gapZ(4.3),
+			m_instancingRenderer(0)
+		{
+		}
+	};
+
+protected:
+		virtual void setupScene(const ConstructionInfo& ci)=0;
+public:
+
+	GpuDemo()
+	{
+		m_dynamicsWorld=0;
+	
+	}
+	virtual ~GpuDemo()
+	{
+		exitPhysics();
+	}
+	virtual void	initPhysics(const ConstructionInfo& ci);
+
+	virtual const char* getName()=0;
+
+	virtual void	exitPhysics();
+
+	virtual const btDynamicsWorld* getDynamicsWorld() const
+	{
+		return m_dynamicsWorld;
+	}
+
+	virtual void renderScene()
+	{
+
+	}
+
+	virtual void clientMoveAndDisplay();
+
+	
+	//virtual void	clientResetScene();
+	
+	
+
+	
+};
+
+class GpuCompoundDemo : public GpuDemo
+{
+public:
+	virtual void setupScene(const ConstructionInfo& ci);
+	virtual const char* getName()
+	{
+		return "GpuCompoundDemo";
+	}
+	static GpuDemo* CreateFunc()
+	{
+		GpuDemo* demo = new GpuCompoundDemo;
+		return demo;
+	}
+};
+
+class GpuBoxDemo : public GpuDemo
+{
+public:
+	virtual void setupScene(const ConstructionInfo& ci);
+	virtual const char* getName()
+	{
+		return "GpuBoxDemo";
+	}
+	static GpuDemo* CreateFunc()
+	{
+		GpuDemo* demo = new GpuBoxDemo;
+		return demo;
+	}
+};
+
+class EmptyDemo : public GpuDemo
+{
+public:
+	virtual void setupScene(const ConstructionInfo& ci);
+	virtual const char* getName()
+	{
+		return "EmptyDemo";
+	}
+	static GpuDemo* CreateFunc()
+	{
+		GpuDemo* demo = new EmptyDemo;
+		return demo;
+	}
+	
+};
+
+class SpheresDemo : public GpuDemo
+{
+public:
+	virtual void setupScene(const ConstructionInfo& ci);
+	virtual const char* getName()
+	{
+		return "SpheresDemo";
+	}
+	static GpuDemo* CreateFunc()
+	{
+		GpuDemo* demo = new SpheresDemo;
+		return demo;
+	}
+	
+};
+
+#endif //GPU_DEMO_H

demo/gpudemo/NVIDIA/premake4.lua

@@ -0,0 +1,234 @@
+
+hasCL = findOpenCL_NVIDIA()
+	
+if (hasCL) then
+	
+	project "bullet2_gpu_demo_opengl2_NVIDIA"
+
+	initOpenCL_NVIDIA()
+
+	language "C++"
+			
+	kind "ConsoleApp"
+	
+	targetdir "../../../../bin"
+
+		includedirs {
+              "..",
+              "../../../../bullet2",
+              "../../../../rendering/Gwen"
+    }
+	
+
+	links { 
+		"BulletSoftBody",
+		"BulletDynamics",
+		"BulletCollision",
+		"LinearMath"
+	}
+	
+
+	initOpenGL()
+	initGlew()
+
+	files {
+			"../GpuDemo.cpp",
+			"../GpuDemo.h",
+			"../btGpuDynamicsWorld.cpp",
+			"../btGpuDynamicsWorld.h",
+			"../btCpuDynamicsWorld.cpp",
+			"../btCpuDynamicsWorld.h",
+			"../btGpuIntegrateTransforms.cpp",
+			"../btGpuIntegrateTransforms.h",
+			"../main_opengl2.cpp",
+			
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.cpp",
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.h",			
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.h",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.h",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.cpp",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.h",
+			"../../../../opencl/broadphase_benchmark/btFillCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btFillCL.h",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.h",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.cpp",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.h",
+			"../../../../dynamics/basic_demo/Stubs/ChNarrowphase.cpp",
+			"../../../../dynamics/basic_demo/Stubs/Solver.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.h",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.h",
+
+			"../../../../opencl/basic_initialize/btOpenCLUtils.cpp",
+			"../../../../opencl/basic_initialize/btOpenCLUtils.h",
+			"../../../../opencl/basic_initialize/btOpenCLInclude.h",
+			
+			
+			
+			"../../../DemosCommon/GL_ShapeDrawer.cpp",
+			"../../../DemosCommon/GL_ShapeDrawer.h",
+			"../../../DemosCommon/OpenGL2Renderer.cpp",
+			"../../../DemosCommon/OpenGL2Renderer.h",
+			
+			"../../../../rendering/WavefrontObjLoader/string_extra.cpp",
+			"../../../../rendering/WavefrontObjLoader/string_extra.h",
+			"../../../../rendering/WavefrontObjLoader/objLoader.cpp",
+			"../../../../rendering/WavefrontObjLoader/objLoader.h",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.cpp",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.h",
+			"../../../../rendering/WavefrontObjLoader/list.cpp",
+			"../../../../rendering/WavefrontObjLoader/list.h",
+			
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.cpp",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.h",
+			"../../../../rendering/rendertest/Win32OpenGLWindow.cpp",
+			"../../../../rendering/rendertest/Win32OpenGLWindow.h",
+			"../../../../rendering/rendertest/Win32Window.cpp",
+			"../../../../rendering/rendertest/Win32Window.h",
+			"../../../../rendering/rendertest/LoadShader.cpp",
+			"../../../../rendering/rendertest/LoadShader.h",
+			"../../../../rendering/rendertest/TwFonts.cpp",
+			"../../../../rendering/rendertest/TwFonts.h",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.h",
+	}
+
+	project "bullet2_gpu_demo_opengl3core_NVIDIA"
+
+	initOpenCL_NVIDIA()
+
+	language "C++"
+			
+	kind "ConsoleApp"
+	
+	targetdir "../../../../bin"
+
+		includedirs {
+              "..",
+              "../../../../bullet2",
+              "../../../../rendering/Gwen"
+    }
+	
+
+	links { 
+		"BulletSoftBody",
+		"BulletDynamics",
+		"BulletCollision",
+		"LinearMath",
+		"gwen",
+	}
+	
+
+	initOpenGL()
+	initGlew()
+
+	files {
+			"../GpuDemo.cpp",
+			"../GpuDemo.h",
+			"../btGpuDynamicsWorld.cpp",
+			"../btGpuDynamicsWorld.h",
+			"../btCpuDynamicsWorld.cpp",
+			"../btCpuDynamicsWorld.h",
+			"../btGpuIntegrateTransforms.cpp",
+			"../btGpuIntegrateTransforms.h",
+			"../GwenUserInterface.cpp",
+			"../GwenUserInterface.h",
+			"../ParticleDemo.cpp",
+			"../ParticleDemo.h",
+			"../BroadphaseBenchmark.cpp",
+			"../BroadphaseBenchmark.h",
+			"../main_opengl3core.cpp",
+	
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.cpp",
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.h",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.h",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.cpp",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.h",
+			"../../../../opencl/broadphase_benchmark/btFillCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btFillCL.h",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.h",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.cpp",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.h",
+			"../../../../dynamics/basic_demo/Stubs/ChNarrowphase.cpp",
+			"../../../../dynamics/basic_demo/Stubs/Solver.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.h",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.h",
+
+			"../../../../opencl/basic_initialize/btOpenCLUtils.cpp",
+			"../../../../opencl/basic_initialize/btOpenCLUtils.h",
+			"../../../../opencl/basic_initialize/btOpenCLInclude.h",
+
+
+			"../../../DemosCommon/GL_ShapeDrawer.cpp",
+			"../../../DemosCommon/GL_ShapeDrawer.h",
+			"../../../DemosCommon/OpenGL3CoreRenderer.cpp",
+			"../../../DemosCommon/OpenGL3CoreRenderer.h",
+			
+			"../../../../rendering/WavefrontObjLoader/string_extra.cpp",
+			"../../../../rendering/WavefrontObjLoader/string_extra.h",
+			"../../../../rendering/WavefrontObjLoader/objLoader.cpp",
+			"../../../../rendering/WavefrontObjLoader/objLoader.h",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.cpp",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.h",
+			"../../../../rendering/WavefrontObjLoader/list.cpp",
+			"../../../../rendering/WavefrontObjLoader/list.h",
+			
+			"../../../../rendering/rendertest/GLInstancingRenderer.cpp",
+			"../../../../rendering/rendertest/GLInstancingRenderer.h",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.cpp",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.h",
+			"../../../../rendering/rendertest/Win32OpenGLWindow.cpp",
+			"../../../../rendering/rendertest/Win32OpenGLWindow.h",
+			"../../../../rendering/rendertest/Win32Window.cpp",
+			"../../../../rendering/rendertest/Win32Window.h",
+			"../../../../rendering/rendertest/LoadShader.cpp",
+			"../../../../rendering/rendertest/LoadShader.h",
+			"../../../../rendering/rendertest/TwFonts.cpp",
+			"../../../../rendering/rendertest/TwFonts.h",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.h",
+														
+	}
+end

demo/gpudemo/NVIDIA/premake4.lua.bak

@@ -0,0 +1,233 @@
+
+hasCL = findOpenCL_NVIDIA()
+	
+if (hasCL) then
+	
+	project "bullet2_gpu_demo_opengl2_NVIDIA"
+
+	initOpenCL_NVIDIA()
+
+	language "C++"
+			
+	kind "ConsoleApp"
+	
+	targetdir "../../../../bin"
+
+		includedirs {
+              "..",
+              "../../../../bullet2",
+              "../../../../rendering/Gwen"
+    }
+	
+
+	links { 
+		"BulletSoftBody",
+		"BulletDynamics",
+		"BulletCollision",
+		"LinearMath"
+	}
+	
+
+	initOpenGL()
+	initGlew()
+
+	files {
+			"../GpuDemo.cpp",
+			"../GpuDemo.h",
+			"../btGpuDynamicsWorld.cpp",
+			"../btGpuDynamicsWorld.h",
+			"../btCpuDynamicsWorld.cpp",
+			"../btCpuDynamicsWorld.h",
+			"../btGpuIntegrateTransforms.cpp",
+			"../btGpuIntegrateTransforms.h",
+			"../main_opengl2.cpp",
+			
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.cpp",
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.h",			
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.h",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.h",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.cpp",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.h",
+			"../../../../opencl/broadphase_benchmark/btFillCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btFillCL.h",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.h",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.cpp",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.h",
+			"../../../../dynamics/basic_demo/Stubs/ChNarrowphase.cpp",
+			"../../../../dynamics/basic_demo/Stubs/Solver.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.h",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.h",
+
+			"../../../../opencl/basic_initialize/btOpenCLUtils.cpp",
+			"../../../../opencl/basic_initialize/btOpenCLUtils.h",
+			"../../../../opencl/basic_initialize/btOpenCLInclude.h",
+			
+			
+			
+			"../../../DemosCommon/GL_ShapeDrawer.cpp",
+			"../../../DemosCommon/GL_ShapeDrawer.h",
+			"../../../DemosCommon/OpenGL2Renderer.cpp",
+			"../../../DemosCommon/OpenGL2Renderer.h",
+			
+			"../../../../rendering/WavefrontObjLoader/string_extra.cpp",
+			"../../../../rendering/WavefrontObjLoader/string_extra.h",
+			"../../../../rendering/WavefrontObjLoader/objLoader.cpp",
+			"../../../../rendering/WavefrontObjLoader/objLoader.h",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.cpp",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.h",
+			"../../../../rendering/WavefrontObjLoader/list.cpp",
+			"../../../../rendering/WavefrontObjLoader/list.h",
+			
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.cpp",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.h",
+			"../../../../rendering/rendertest/Win32OpenGLWindow.cpp",
+			"../../../../rendering/rendertest/Win32OpenGLWindow.h",
+			"../../../../rendering/rendertest/Win32Window.cpp",
+			"../../../../rendering/rendertest/Win32Window.h",
+			"../../../../rendering/rendertest/LoadShader.cpp",
+			"../../../../rendering/rendertest/LoadShader.h",
+			"../../../../rendering/rendertest/TwFonts.cpp",
+			"../../../../rendering/rendertest/TwFonts.h",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.h",
+	}
+
+	project "bullet2_gpu_demo_opengl3core_NVIDIA"
+
+	initOpenCL_NVIDIA()
+
+	language "C++"
+			
+	kind "ConsoleApp"
+	
+	targetdir "../../../../bin"
+
+		includedirs {
+              "..",
+              "../../../../bullet2",
+              "../../../../rendering/Gwen"
+    }
+	
+
+	links { 
+		"BulletSoftBody",
+		"BulletDynamics",
+		"BulletCollision",
+		"LinearMath",
+		"gwen",
+	}
+	
+
+	initOpenGL()
+	initGlew()
+
+	files {
+			"../GpuDemo.cpp",
+			"../GpuDemo.h",
+			"../btGpuDynamicsWorld.cpp",
+			"../btGpuDynamicsWorld.h",
+			"../btCpuDynamicsWorld.cpp",
+			"../btCpuDynamicsWorld.h",
+			"../btGpuIntegrateTransforms.cpp",
+			"../btGpuIntegrateTransforms.h",
+			"../GwenUserInterface.cpp",
+			"../GwenUserInterface.h",
+			"../ParticleDemo.cpp",
+			"../ParticleDemo.h",
+	
+			"../main_opengl3core.cpp",
+	
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.cpp",
+			"../../../../opencl/opengl_interop/btOpenCLGLInteropBuffer.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/CLPhysicsDemo.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btPgsJacobiSolver.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/btGpuSapBroadphase.h",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline2/ConvexHullContact.h",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btPrefixScanCL.h",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.cpp",
+			"../../../../opencl/broadphase_benchmark/btRadixSort32CL.h",
+			"../../../../opencl/broadphase_benchmark/btFillCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btFillCL.h",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btBoundSearchCL.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btConvexUtility.h",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.cpp",
+			"../../../../opencl/gpu_rigidbody_pipeline/btGpuNarrowPhaseAndSolver.h",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.cpp",
+			"../../../../dynamics/basic_demo/ConvexHeightFieldShape.h",
+			"../../../../dynamics/basic_demo/Stubs/ChNarrowphase.cpp",
+			"../../../../dynamics/basic_demo/Stubs/Solver.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.cpp",
+			"../../../../opencl/broadphase_benchmark/findPairsOpenCL.h",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.cpp",
+			"../../../../opencl/broadphase_benchmark/btGridBroadphaseCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.h",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.cpp",
+			"../../../../opencl/3dGridBroadphase/Shared/btGpu3DGridBroadphase.h",
+
+			"../../../../opencl/basic_initialize/btOpenCLUtils.cpp",
+			"../../../../opencl/basic_initialize/btOpenCLUtils.h",
+			"../../../../opencl/basic_initialize/btOpenCLInclude.h",
+
+
+			"../../../DemosCommon/GL_ShapeDrawer.cpp",
+			"../../../DemosCommon/GL_ShapeDrawer.h",
+			"../../../DemosCommon/OpenGL3CoreRenderer.cpp",
+			"../../../DemosCommon/OpenGL3CoreRenderer.h",
+			
+			"../../../../rendering/WavefrontObjLoader/string_extra.cpp",
+			"../../../../rendering/WavefrontObjLoader/string_extra.h",
+			"../../../../rendering/WavefrontObjLoader/objLoader.cpp",
+			"../../../../rendering/WavefrontObjLoader/objLoader.h",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.cpp",
+			"../../../../rendering/WavefrontObjLoader/obj_parser.h",
+			"../../../../rendering/WavefrontObjLoader/list.cpp",
+			"../../../../rendering/WavefrontObjLoader/list.h",
+			
+			"../../../../rendering/rendertest/GLInstancingRenderer.cpp",
+			"../../../../rendering/rendertest/GLInstancingRenderer.h",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.cpp",
+			"../../../../rendering/rendertest/GLPrimitiveRenderer.h",
+			"../../../../rendering/rendertest/Win32OpenGLWindow.cpp",
+			"../../../../rendering/rendertest/Win32OpenGLWindow.h",
+			"../../../../rendering/rendertest/Win32Window.cpp",
+			"../../../../rendering/rendertest/Win32Window.h",
+			"../../../../rendering/rendertest/LoadShader.cpp",
+			"../../../../rendering/rendertest/LoadShader.h",
+			"../../../../rendering/rendertest/TwFonts.cpp",
+			"../../../../rendering/rendertest/TwFonts.h",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.cpp",
+			"../../../../rendering/OpenGLTrueTypeFont/fontstash.h",
+														
+	}
+end

demo/gpudemo/OpenGL3CoreRenderer.cpp

@@ -0,0 +1,677 @@
+
+#include "OpenGL3CoreRenderer.h"
+#include "../../rendering/rendertest/GLInstancingRenderer.h"
+#include "../../rendering/rendertest/ShapeData.h"
+#include "BulletDynamics/Dynamics/btDiscreteDynamicsWorld.h"
+#include "BulletCollision/CollisionDispatch/btCollisionObject.h"
+#include "LinearMath/btQuickprof.h"
+
+#include "BulletCollision/CollisionShapes/btBvhTriangleMeshShape.h"
+#include "BulletCollision/CollisionShapes/btConvexPolyhedron.h"
+#include "BulletCollision/CollisionShapes/btConvexHullShape.h"
+#include "BulletCollision/CollisionShapes/btCollisionShape.h"
+#include "BulletCollision/CollisionShapes/btBoxShape.h"
+#include "BulletCollision/CollisionShapes/btCompoundShape.h"
+#include "BulletCollision/CollisionShapes/btSphereShape.h"
+#include "BulletCollision/CollisionShapes/btStaticPlaneShape.h"
+
+#include "../../rendering/WavefrontObjLoader/objLoader.h"
+
+OpenGL3CoreRenderer::OpenGL3CoreRenderer()
+{
+	int maxNumObjects = 2*1024*1024;
+	m_instancingRenderer = new GLInstancingRenderer(maxNumObjects);
+	m_instancingRenderer->setCameraDistance(150);
+}
+OpenGL3CoreRenderer::~OpenGL3CoreRenderer()
+{
+	delete m_instancingRenderer;
+}
+
+void OpenGL3CoreRenderer::init()
+{
+	m_instancingRenderer->InitShaders();
+}
+
+
+
+void OpenGL3CoreRenderer::reshape(int w, int h)
+{
+	m_instancingRenderer->resize(w,h);
+}
+void OpenGL3CoreRenderer::keyboardCallback(unsigned char key)
+{
+}
+
+struct GraphicsVertex
+{
+	float xyzw[4];
+	float normal[3];
+	float uv[2];
+};
+struct GraphicsShape
+{
+	const float*	m_vertices;
+	int				m_numvertices;
+	const int*		m_indices;
+	int				m_numIndices;
+	float			m_scaling[4];
+};
+
+
+
+GraphicsShape* createGraphicsShapeFromConvexHull(const btConvexPolyhedron* utilPtr)
+{
+	
+	btAlignedObjectArray<GraphicsVertex>* vertices = new btAlignedObjectArray<GraphicsVertex>;
+	{
+		int numVertices = utilPtr->m_vertices.size();
+		int numIndices = 0;
+		btAlignedObjectArray<int>* indicesPtr = new btAlignedObjectArray<int>;
+		for (int f=0;f<utilPtr->m_faces.size();f++)
+		{
+			const btFace& face = utilPtr->m_faces[f];
+			btVector3 normal(face.m_plane[0],face.m_plane[1],face.m_plane[2]);
+			if (face.m_indices.size()>2)
+			{
+				
+				GraphicsVertex vtx;
+				const btVector3& orgVertex = utilPtr->m_vertices[face.m_indices[0]];
+				vtx.xyzw[0] = orgVertex[0];vtx.xyzw[1] = orgVertex[1];vtx.xyzw[2] = orgVertex[2];vtx.xyzw[3] = 0.f;
+				vtx.normal[0] = normal[0];vtx.normal[1] = normal[1];vtx.normal[2] = normal[2];
+				vtx.uv[0] = 0.5f;vtx.uv[1] = 0.5f;
+				int newvtxindex0 = vertices->size();
+				vertices->push_back(vtx);
+			
+				for (int j=1;j<face.m_indices.size()-1;j++)
+				{
+					indicesPtr->push_back(newvtxindex0);
+					{
+						GraphicsVertex vtx;
+						const btVector3& orgVertex = utilPtr->m_vertices[face.m_indices[j]];
+						vtx.xyzw[0] = orgVertex[0];vtx.xyzw[1] = orgVertex[1];vtx.xyzw[2] = orgVertex[2];vtx.xyzw[3] = 0.f;
+						vtx.normal[0] = normal[0];vtx.normal[1] = normal[1];vtx.normal[2] = normal[2];
+						vtx.uv[0] = 0.5f;vtx.uv[1] = 0.5f;
+						int newvtxindexj = vertices->size();
+						vertices->push_back(vtx);
+						indicesPtr->push_back(newvtxindexj);
+					}
+
+					{
+						GraphicsVertex vtx;
+						const btVector3& orgVertex = utilPtr->m_vertices[face.m_indices[j+1]];
+						vtx.xyzw[0] = orgVertex[0];vtx.xyzw[1] = orgVertex[1];vtx.xyzw[2] = orgVertex[2];vtx.xyzw[3] = 0.f;
+						vtx.normal[0] = normal[0];vtx.normal[1] = normal[1];vtx.normal[2] = normal[2];
+						vtx.uv[0] = 0.5f;vtx.uv[1] = 0.5f;
+						int newvtxindexj1 = vertices->size();
+						vertices->push_back(vtx);
+						indicesPtr->push_back(newvtxindexj1);
+					}
+				}
+			}
+		}
+		
+		
+		GraphicsShape* gfxShape = new GraphicsShape;
+		gfxShape->m_vertices = &vertices->at(0).xyzw[0];
+		gfxShape->m_numvertices = vertices->size();
+		gfxShape->m_indices = &indicesPtr->at(0);
+		gfxShape->m_numIndices = indicesPtr->size();
+		for (int i=0;i<4;i++)
+			gfxShape->m_scaling[i] = 1;//bake the scaling into the vertices 
+		return gfxShape;
+	}
+}
+
+GraphicsShape* createGraphicsShapeFromCompoundShape(btCompoundShape* compound)
+{
+	GraphicsShape* gfxShape = new GraphicsShape();
+	btAlignedObjectArray<GraphicsVertex>* vertexArray = new btAlignedObjectArray<GraphicsVertex>;
+	btAlignedObjectArray<int>* indexArray = new btAlignedObjectArray<int>;
+
+
+
+	//create a graphics shape for each child, combine them into a single graphics shape using their child transforms
+	for (int i=0;i<compound->getNumChildShapes();i++)
+	{
+		btAssert(compound->getChildShape(i)->isPolyhedral());
+		if (compound->getChildShape(i)->isPolyhedral())
+		{
+			btPolyhedralConvexShape* convexHull = (btPolyhedralConvexShape*) compound->getChildShape(i);
+			btTransform tr = compound->getChildTransform(i);
+			
+			const btConvexPolyhedron* polyhedron = convexHull->getConvexPolyhedron();
+			GraphicsShape* childGfxShape = createGraphicsShapeFromConvexHull(polyhedron);
+			int baseIndex = vertexArray->size();
+
+			for (int j=0;j<childGfxShape->m_numIndices;j++)
+				indexArray->push_back(childGfxShape->m_indices[j]+baseIndex);
+			
+			GraphicsVertex* orgVerts = (GraphicsVertex*)childGfxShape->m_vertices;
+
+			for (int j=0;j<childGfxShape->m_numvertices;j++)
+			{
+				GraphicsVertex vtx;
+				btVector3 pos(orgVerts[j].xyzw[0],orgVerts[j].xyzw[1],orgVerts[j].xyzw[2]);
+				pos = tr*pos;
+				vtx.xyzw[0] = childGfxShape->m_scaling[0]*pos.x();
+				vtx.xyzw[1] = childGfxShape->m_scaling[1]*pos.y();
+				vtx.xyzw[2] = childGfxShape->m_scaling[2]*pos.z();
+				vtx.xyzw[3] = 10.f;
+				
+				vtx.uv[0] = 0.5f;
+				vtx.uv[1] = 0.5f;
+
+				btVector3 normal(orgVerts[j].normal[0],orgVerts[j].normal[1],orgVerts[j].normal[2]);
+				normal = tr.getBasis()*normal;
+				vtx.normal[0] = normal.x();
+				vtx.normal[1] = normal.y();
+				vtx.normal[2] = normal.z();
+				vertexArray->push_back(vtx);
+			}
+		}
+	}
+
+	btPolyhedralConvexShape* convexHull = (btPolyhedralConvexShape*) compound->getChildShape(0);
+	const btConvexPolyhedron* polyhedron = convexHull->getConvexPolyhedron();
+	GraphicsShape* childGfxShape = createGraphicsShapeFromConvexHull(polyhedron);
+
+	gfxShape->m_indices = &indexArray->at(0);
+	gfxShape->m_numIndices = indexArray->size();
+	gfxShape->m_vertices = &vertexArray->at(0).xyzw[0];
+	gfxShape->m_numvertices = vertexArray->size();
+	gfxShape->m_scaling[0] = 1;
+	gfxShape->m_scaling[1] = 1;
+	gfxShape->m_scaling[2] = 1;
+	gfxShape->m_scaling[3] = 1;
+	
+	return gfxShape;
+}
+
+GraphicsShape* createGraphicsShapeFromConcaveMesh(const btBvhTriangleMeshShape* trimesh)
+{
+	
+	btAlignedObjectArray<GraphicsVertex>* vertices = new btAlignedObjectArray<GraphicsVertex>;
+	btAlignedObjectArray<int>* indicesPtr = new btAlignedObjectArray<int>;
+
+	const btStridingMeshInterface* meshInterface = trimesh->getMeshInterface();
+
+	btVector3 trimeshScaling(1,1,1);
+	for (int partId=0;partId<meshInterface->getNumSubParts();partId++)
+	{
+		
+		const unsigned char *vertexbase = 0;
+		int numverts = 0;
+		PHY_ScalarType type = PHY_INTEGER;
+		int stride = 0;
+		const unsigned char *indexbase = 0;
+		int indexstride = 0;
+		int numfaces = 0;
+		PHY_ScalarType indicestype = PHY_INTEGER;
+		//PHY_ScalarType indexType=0;
+
+		btVector3 triangleVerts[3];
+		meshInterface->getLockedReadOnlyVertexIndexBase(&vertexbase,numverts,	type,stride,&indexbase,indexstride,numfaces,indicestype,partId);
+		btVector3 aabbMin,aabbMax;
+
+		for (int triangleIndex = 0 ; triangleIndex < numfaces;triangleIndex++)
+		{
+			unsigned int* gfxbase = (unsigned int*)(indexbase+triangleIndex*indexstride);
+
+			for (int j=2;j>=0;j--)
+			{
+
+				int graphicsindex = indicestype==PHY_SHORT?((unsigned short*)gfxbase)[j]:gfxbase[j];
+				if (type == PHY_FLOAT)
+				{
+					float* graphicsbase = (float*)(vertexbase+graphicsindex*stride);
+					triangleVerts[j] = btVector3(
+						graphicsbase[0]*trimeshScaling.getX(),
+						graphicsbase[1]*trimeshScaling.getY(),
+						graphicsbase[2]*trimeshScaling.getZ());
+				}
+				else
+				{
+					double* graphicsbase = (double*)(vertexbase+graphicsindex*stride);
+					triangleVerts[j] = btVector3( btScalar(graphicsbase[0]*trimeshScaling.getX()), 
+						btScalar(graphicsbase[1]*trimeshScaling.getY()), 
+						btScalar(graphicsbase[2]*trimeshScaling.getZ()));
+				}
+			}
+			btVector3 normal = (triangleVerts[2]-triangleVerts[0]).cross(triangleVerts[1]-triangleVerts[0]);
+			normal.normalize();
+
+			GraphicsVertex vtx0,vtx1,vtx2;
+			vtx0.xyzw[0] = triangleVerts[0].getX();
+			vtx0.xyzw[1] = triangleVerts[0].getY();
+			vtx0.xyzw[2] = triangleVerts[0].getZ();
+			vtx0.xyzw[3] = 0;
+			vtx0.uv[0] = 0.5f;
+			vtx0.uv[1] = 0.5f;
+			vtx0.normal[0] = normal[0];
+			vtx0.normal[1] = normal[1];
+			vtx0.normal[2] = normal[2];
+
+			vtx1.xyzw[0] = triangleVerts[1].getX();
+			vtx1.xyzw[1] = triangleVerts[1].getY();
+			vtx1.xyzw[2] = triangleVerts[1].getZ();
+			vtx1.xyzw[3] = 0;
+			vtx1.uv[0] = 0.5f;
+			vtx1.uv[1] = 0.5f;
+			vtx1.normal[0] = normal[0];
+			vtx1.normal[1] = normal[1];
+			vtx1.normal[2] = normal[2];
+
+			vtx2.xyzw[0] = triangleVerts[2].getX();
+			vtx2.xyzw[1] = triangleVerts[2].getY();
+			vtx2.xyzw[2] = triangleVerts[2].getZ();
+			vtx2.xyzw[3] = 0;
+			vtx2.uv[0] = 0.5f;
+			vtx2.uv[1] = 0.5f;
+			vtx2.normal[0] = normal[0];
+			vtx2.normal[1] = normal[1];
+			vtx2.normal[2] = normal[2];
+
+//			triangleVerts[1]
+//			triangleVerts[1]
+//			triangleVerts[2]
+			vertices->push_back(vtx0);
+			vertices->push_back(vtx1);
+			vertices->push_back(vtx2);
+			indicesPtr->push_back(indicesPtr->size());
+			indicesPtr->push_back(indicesPtr->size());
+			indicesPtr->push_back(indicesPtr->size());
+		}
+	}
+
+		
+	GraphicsShape* gfxShape = new GraphicsShape;
+	gfxShape->m_vertices = &vertices->at(0).xyzw[0];
+	gfxShape->m_numvertices = vertices->size();
+	gfxShape->m_indices = &indicesPtr->at(0);
+	gfxShape->m_numIndices = indicesPtr->size();
+	for (int i=0;i<4;i++)
+		gfxShape->m_scaling[i] = 1;//bake the scaling into the vertices 
+	return gfxShape;
+}
+
+
+GraphicsShape* createGraphicsShapeFromWavefrontObj(objLoader* obj)
+{
+	btAlignedObjectArray<GraphicsVertex>* vertices = new btAlignedObjectArray<GraphicsVertex>;
+	{
+//		int numVertices = obj->vertexCount;
+	//	int numIndices = 0;
+		btAlignedObjectArray<int>* indicesPtr = new btAlignedObjectArray<int>;
+			/*
+		for (int v=0;v<obj->vertexCount;v++)
+		{
+			vtx.xyzw[0] = obj->vertexList[v]->e[0];
+			vtx.xyzw[1] = obj->vertexList[v]->e[1];
+			vtx.xyzw[2] = obj->vertexList[v]->e[2];
+			btVector3 n(vtx.xyzw[0],vtx.xyzw[1],vtx.xyzw[2]);
+			if (n.length2()>SIMD_EPSILON)
+			{
+				n.normalize();
+				vtx.normal[0] = n[0];
+				vtx.normal[1] = n[1];
+				vtx.normal[2] = n[2];
+
+			} else
+			{
+				vtx.normal[0] = 0; //todo
+				vtx.normal[1] = 1;
+				vtx.normal[2] = 0;
+			}
+			vtx.uv[0] = 0.5f;vtx.uv[1] = 0.5f;	//todo
+			vertices->push_back(vtx);
+		}
+		*/
+
+		for (int f=0;f<obj->faceCount;f++)
+		{
+			obj_face* face = obj->faceList[f];
+			//btVector3 normal(face.m_plane[0],face.m_plane[1],face.m_plane[2]);
+			if (face->vertex_count>=3)
+			{
+				btVector3 normal(0,1,0);
+				int vtxBaseIndex = vertices->size();
+
+				if (face->vertex_count<=4)
+				{
+					indicesPtr->push_back(vtxBaseIndex);
+					indicesPtr->push_back(vtxBaseIndex+1);
+					indicesPtr->push_back(vtxBaseIndex+2);
+					
+					GraphicsVertex vtx0;
+					vtx0.xyzw[0] = obj->vertexList[face->vertex_index[0]]->e[0];
+					vtx0.xyzw[1] = obj->vertexList[face->vertex_index[0]]->e[1];
+					vtx0.xyzw[2] = obj->vertexList[face->vertex_index[0]]->e[2];
+					vtx0.uv[0] = obj->textureList[face->vertex_index[0]]->e[0];
+					vtx0.uv[1] = obj->textureList[face->vertex_index[0]]->e[1];
+
+					GraphicsVertex vtx1;
+					vtx1.xyzw[0] = obj->vertexList[face->vertex_index[1]]->e[0];
+					vtx1.xyzw[1] = obj->vertexList[face->vertex_index[1]]->e[1];
+					vtx1.xyzw[2] = obj->vertexList[face->vertex_index[1]]->e[2];
+					vtx1.uv[0] = obj->textureList[face->vertex_index[1]]->e[0];
+					vtx1.uv[1] = obj->textureList[face->vertex_index[1]]->e[1];
+
+					GraphicsVertex vtx2;
+					vtx2.xyzw[0] = obj->vertexList[face->vertex_index[2]]->e[0];
+					vtx2.xyzw[1] = obj->vertexList[face->vertex_index[2]]->e[1];
+					vtx2.xyzw[2] = obj->vertexList[face->vertex_index[2]]->e[2];
+					vtx2.uv[0] = obj->textureList[face->vertex_index[2]]->e[0];
+					vtx2.uv[1] = obj->textureList[face->vertex_index[2]]->e[1];
+
+
+					btVector3 v0(vtx0.xyzw[0],vtx0.xyzw[1],vtx0.xyzw[2]);
+					btVector3 v1(vtx1.xyzw[0],vtx1.xyzw[1],vtx1.xyzw[2]);
+					btVector3 v2(vtx2.xyzw[0],vtx2.xyzw[1],vtx2.xyzw[2]);
+
+					normal = (v1-v0).cross(v2-v0);
+					normal.normalize();
+					vtx0.normal[0] = normal[0];
+					vtx0.normal[1] = normal[1];
+					vtx0.normal[2] = normal[2];
+					vtx1.normal[0] = normal[0];
+					vtx1.normal[1] = normal[1];
+					vtx1.normal[2] = normal[2];
+					vtx2.normal[0] = normal[0];
+					vtx2.normal[1] = normal[1];
+					vtx2.normal[2] = normal[2];
+					vertices->push_back(vtx0);
+					vertices->push_back(vtx1);
+					vertices->push_back(vtx2);
+				}
+				if (face->vertex_count==4)
+				{
+
+					indicesPtr->push_back(vtxBaseIndex);
+					indicesPtr->push_back(vtxBaseIndex+1);
+					indicesPtr->push_back(vtxBaseIndex+2);
+					indicesPtr->push_back(vtxBaseIndex+3);
+//
+					GraphicsVertex vtx3;
+					vtx3.xyzw[0] = obj->vertexList[face->vertex_index[3]]->e[0];
+					vtx3.xyzw[1] = obj->vertexList[face->vertex_index[3]]->e[1];
+					vtx3.xyzw[2] = obj->vertexList[face->vertex_index[3]]->e[2];
+					vtx3.uv[0] = 0.5;
+					vtx3.uv[1] = 0.5;
+
+					vtx3.normal[0] = normal[0];
+					vtx3.normal[1] = normal[1];
+					vtx3.normal[2] = normal[2];
+
+					vertices->push_back(vtx3);
+
+				}
+			}
+		}
+		
+		
+		GraphicsShape* gfxShape = new GraphicsShape;
+		gfxShape->m_vertices = &vertices->at(0).xyzw[0];
+		gfxShape->m_numvertices = vertices->size();
+		gfxShape->m_indices = &indicesPtr->at(0);
+		gfxShape->m_numIndices = indicesPtr->size();
+		for (int i=0;i<4;i++)
+			gfxShape->m_scaling[i] = 1;//bake the scaling into the vertices 
+		return gfxShape;
+	}
+}
+
+
+
+//very incomplete conversion from physics to graphics
+void graphics_from_physics(GLInstancingRenderer& renderer, bool syncTransformsOnly, int numObjects, btCollisionObject** colObjArray)
+{
+	///@todo: we need to sort the objects based on collision shape type, so we can share instances
+	BT_PROFILE("graphics_from_physics");
+    
+	int strideInBytes = sizeof(float)*9;
+    
+	int prevGraphicsShapeIndex  = -1;
+	btCollisionShape* prevShape = 0;
+
+    
+	int numColObj = numObjects;
+    int curGraphicsIndex = 0;
+
+	float localScaling[4] = {1,1,1,1};
+	
+    
+    for (int i=0;i<numColObj;i++)
+    {
+		btCollisionObject* colObj = colObjArray[i];
+		
+        btVector3 pos = colObj->getWorldTransform().getOrigin();
+        btQuaternion orn = colObj->getWorldTransform().getRotation();
+        
+        float position[4] = {pos.getX(),pos.getY(),pos.getZ(),0.f};
+        float orientation[4] = {orn.getX(),orn.getY(),orn.getZ(),orn.getW()};
+        float color[4] = {0,0,0,1};
+        btVector3 localScaling = colObj->getCollisionShape()->getLocalScaling();
+        
+       
+        if (colObj->isStaticOrKinematicObject())
+        {
+            color[0]=1.f;
+        }else
+        {
+            color[1]=1.f;
+        }
+        
+		if (!syncTransformsOnly)
+		{
+			
+			if (prevShape != colObj->getCollisionShape())
+			{
+				if (colObj->getCollisionShape()->isPolyhedral())
+				{
+					btPolyhedralConvexShape* polyShape = (btPolyhedralConvexShape*)colObj->getCollisionShape();
+					const btConvexPolyhedron* pol = polyShape->getConvexPolyhedron();
+					GraphicsShape* gfxShape = createGraphicsShapeFromConvexHull(pol);
+
+					prevGraphicsShapeIndex = renderer.registerShape(&gfxShape->m_vertices[0],gfxShape->m_numvertices,gfxShape->m_indices,gfxShape->m_numIndices);
+					prevShape = colObj->getCollisionShape();
+					const btVector3& scaling = prevShape->getLocalScaling();
+					localScaling[0] = scaling.getX();localScaling[1] = scaling.getY();localScaling[2] = scaling.getZ();
+				} else
+				{
+					if (colObj->getCollisionShape()->getShapeType()==TRIANGLE_MESH_SHAPE_PROXYTYPE)
+					{
+						btBvhTriangleMeshShape* trimesh = (btBvhTriangleMeshShape*) colObj->getCollisionShape();
+						GraphicsShape* gfxShape = createGraphicsShapeFromConcaveMesh(trimesh);
+						prevGraphicsShapeIndex = renderer.registerShape(&gfxShape->m_vertices[0],gfxShape->m_numvertices,gfxShape->m_indices,gfxShape->m_numIndices);
+						prevShape = colObj->getCollisionShape();
+						const btVector3& scaling = prevShape->getLocalScaling();
+						localScaling[0] = scaling.getX();localScaling[1] = scaling.getY();localScaling[2] = scaling.getZ();
+					} else
+					{
+						if (colObj->getCollisionShape()->getShapeType()==COMPOUND_SHAPE_PROXYTYPE)
+						{
+							btCompoundShape* compound = (btCompoundShape*) colObj->getCollisionShape();
+							GraphicsShape* gfxShape = createGraphicsShapeFromCompoundShape(compound);
+							if (gfxShape)
+							{
+								prevGraphicsShapeIndex = renderer.registerShape(&gfxShape->m_vertices[0],gfxShape->m_numvertices,gfxShape->m_indices,gfxShape->m_numIndices);
+								prevShape = colObj->getCollisionShape();
+								const btVector3& scaling = prevShape->getLocalScaling();
+								localScaling[0] = scaling.getX();localScaling[1] = scaling.getY();localScaling[2] = scaling.getZ();
+							} else
+							{
+								prevGraphicsShapeIndex = -1;
+							}
+						} else
+						{
+							if (colObj->getCollisionShape()->getShapeType()==SPHERE_SHAPE_PROXYTYPE)
+							{
+								btSphereShape* sphere = (btSphereShape*) colObj->getCollisionShape();
+								btScalar radius = sphere->getRadius();
+								
+								//btConvexHullShape* spherePoly = new btConvexHullShape(
+								//const btConvexPolyhedron* pol = polyShape->getConvexPolyhedron();
+								
+								/*
+								objLoader loader;
+								
+								int result = loader.load("../../bin/wavefront/sphere_low.obj");
+								
+
+								GraphicsShape* gfxShape = createGraphicsShapeFromWavefrontObj(&loader);
+								
+
+								int vertexStrideInBytes = 9*sizeof(float);
+
+								
+								printf("vertices (%d):\n",gfxShape->m_numvertices);
+								for (int i=0;i<gfxShape->m_numvertices;i++)
+								{
+									gfxShape->m_vertices[i*9+4] = gfxShape->m_vertices[i*9];
+									gfxShape->m_vertices[i*9+5] = gfxShape->m_vertices[i*9+1];
+									gfxShape->m_vertices[i*9+6] = gfxShape->m_vertices[i*9+2];
+
+									printf("%f,%f,%f,%f,%f,%f,%f,%f,%f,\n",
+										gfxShape->m_vertices[i*9],
+										gfxShape->m_vertices[i*9+1],
+										gfxShape->m_vertices[i*9+2],
+										0.f,//gfxShape->m_vertices[i*9+3],
+										//gfxShape->m_vertices[i*9+4],//face normals
+										//gfxShape->m_vertices[i*9+5],
+										//gfxShape->m_vertices[i*9+6],
+
+										gfxShape->m_vertices[i*9+0],
+										gfxShape->m_vertices[i*9+1],
+										gfxShape->m_vertices[i*9+2],
+
+										gfxShape->m_vertices[i*9+7],
+										gfxShape->m_vertices[i*9+8]);
+								}
+								printf("indices (%d):\n",gfxShape->m_numIndices);
+								for (int i=0;i<gfxShape->m_numIndices/3;i++)
+								{
+									printf("%d,%d,%d,\n",gfxShape->m_indices[i*3],
+									gfxShape->m_indices[i*3+1],
+									gfxShape->m_indices[i*3+2]);
+								}
+
+								prevGraphicsShapeIndex = renderer.registerShape(&gfxShape->m_vertices[0],gfxShape->m_numvertices,gfxShape->m_indices,gfxShape->m_numIndices);
+								*/
+								
+								if (radius>=100)
+								{
+									int numVertices = sizeof(detailed_sphere_vertices)/strideInBytes;
+									int numIndices = sizeof(detailed_sphere_indices)/sizeof(int);
+									prevGraphicsShapeIndex = renderer.registerShape(&detailed_sphere_vertices[0],numVertices,detailed_sphere_indices,numIndices);
+								} else
+								{
+									bool usePointSprites = true;
+									if (usePointSprites)
+									{
+										int numVertices = sizeof(point_sphere_vertices)/strideInBytes;
+										int numIndices = sizeof(point_sphere_indices)/sizeof(int);
+										prevGraphicsShapeIndex = renderer.registerShape(&point_sphere_vertices[0],numVertices,point_sphere_indices,numIndices,BT_GL_POINTS);
+									} else
+									{
+										if (radius>=10)
+										{
+											int numVertices = sizeof(medium_sphere_vertices)/strideInBytes;
+											int numIndices = sizeof(medium_sphere_indices)/sizeof(int);
+											prevGraphicsShapeIndex = renderer.registerShape(&medium_sphere_vertices[0],numVertices,medium_sphere_indices,numIndices);
+										} else
+										{
+											int numVertices = sizeof(low_sphere_vertices)/strideInBytes;
+											int numIndices = sizeof(low_sphere_indices)/sizeof(int);
+											prevGraphicsShapeIndex = renderer.registerShape(&low_sphere_vertices[0],numVertices,low_sphere_indices,numIndices);
+										}
+									}
+								}
+								prevShape = sphere;
+								const btVector3& scaling = prevShape->getLocalScaling();
+								//assume uniform scaling, using X component
+								float sphereScale = radius*scaling.getX();
+								localScaling[0] = sphereScale;
+								localScaling[1] = sphereScale;
+								localScaling[2] = sphereScale;
+							} else
+							{
+								if (colObj->getCollisionShape()->getShapeType()==STATIC_PLANE_PROXYTYPE)
+								{
+									btStaticPlaneShape* plane= (btStaticPlaneShape*) colObj->getCollisionShape();
+									prevShape = colObj->getCollisionShape();
+
+									//plane->getPlaneNormal()
+									//plane->getPlaneConstant()
+									if (1)
+									{
+										int numVertices = sizeof(quad_vertices)/strideInBytes;
+										int numIndices = sizeof(quad_indices)/sizeof(int);
+									
+										prevGraphicsShapeIndex = renderer.registerShape(&quad_vertices[0],numVertices,quad_indices,numIndices);
+									} else
+									{
+										int numVertices = sizeof(detailed_sphere_vertices)/strideInBytes;
+										int numIndices = sizeof(detailed_sphere_indices)/sizeof(int);
+										prevGraphicsShapeIndex = renderer.registerShape(&detailed_sphere_vertices[0],numVertices,detailed_sphere_indices,numIndices);
+									}
+							
+									localScaling[0] = 100;
+									localScaling[1] = 1;
+									localScaling[2] = 100;
+								} else
+								{
+									printf("Error: unsupported collision shape type in %s %d\n", __FILE__, __LINE__);
+									prevGraphicsShapeIndex = -1;
+									btAssert(0);
+								}
+							}
+						}
+					}
+
+				}
+			}
+		}
+    
+		
+		
+		{
+                if (!syncTransformsOnly)
+                {
+					if (prevShape && prevGraphicsShapeIndex>=0)
+					{
+						
+						renderer.registerGraphicsInstance(prevGraphicsShapeIndex,position,orientation,color,localScaling);
+					}
+                }
+                else
+                {
+                    renderer.writeSingleInstanceTransformToCPU(position,orientation,curGraphicsIndex);
+                
+                }
+                curGraphicsIndex++;
+		}
+
+       
+    }
+	
+}
+
+
+
+void OpenGL3CoreRenderer::renderPhysicsWorld(int numObjects, btCollisionObject** colObjArray, bool syncOnly)
+{
+	//sync changes from physics world to render world	
+	//for now, we don't deal with adding/removing objects to the world during the simulation, to keep the rendererer simpler
+
+
+	m_instancingRenderer->writeTransforms();
+	
+	graphics_from_physics(*m_instancingRenderer,syncOnly,numObjects, colObjArray);
+	
+
+	//render
+	
+	m_instancingRenderer->RenderScene();
+}
+

demo/gpudemo/OpenGL3CoreRenderer.h

@@ -0,0 +1,26 @@
+#ifndef OPENGL3_CORE_RENDERER_H
+#define OPENGL3_CORE_RENDERER_H
+
+class btCollisionObject;
+class GLInstancingRenderer;
+
+class OpenGL3CoreRenderer
+{
+
+	GLInstancingRenderer* m_instancingRenderer;
+public:
+	OpenGL3CoreRenderer();
+	virtual ~OpenGL3CoreRenderer();
+	void init();
+	void reshape(int w, int h); 
+	void keyboardCallback(unsigned char key);
+	void renderPhysicsWorld(int numObjects, btCollisionObject** colObjArray, bool syncOnly);
+
+	GLInstancingRenderer*	getInstancingRenderer()
+	{
+		return m_instancingRenderer;
+	}
+};
+
+#endif //OPENGL3_CORE_RENDERER_H
+

demo/gpudemo/ParticleDemo.cpp

@@ -0,0 +1,524 @@
+#include "ParticleDemo.h"
+
+#include "OpenGLWindow/GLInstancingRenderer.h"
+#include "OpenGLWindow/ShapeData.h"
+#include "basic_initialize/btOpenCLUtils.h"
+
+#define MSTRINGIFY(A) #A
+static char* particleKernelsString = 
+#include "ParticleKernels.cl"
+
+#define INTEROPKERNEL_SRC_PATH "../../demos/gpudemo/ParticleKernels.cl"
+#include "BulletCommon/btVector3.h"
+#include "OpenGLWindow/OpenGLInclude.h"
+#include "OpenGLWindow/GLInstanceRendererInternalData.h"
+#include "parallel_primitives/host/btLauncherCL.h"
+//#include "../../opencl/primitives/AdlPrimitives/Math/Math.h"
+//#include "../../opencl/broadphase_benchmark/btGridBroadphaseCL.h"
+#include "gpu_broadphase/host/btGpuSapBroadphase.h"
+
+
+
+#include "BulletCommon/btQuickprof.h"
+
+//1000000 particles
+//#define NUM_PARTICLES_X 100
+//#define NUM_PARTICLES_Y 100
+//#define NUM_PARTICLES_Z 100
+
+//512k particles
+//#define NUM_PARTICLES_X 80
+//#define NUM_PARTICLES_Y 80
+//#define NUM_PARTICLES_Z 80
+
+//256k particles
+//#define NUM_PARTICLES_X 60
+//#define NUM_PARTICLES_Y 60
+//#define NUM_PARTICLES_Z 60
+
+//27k particles
+#define NUM_PARTICLES_X 30
+#define NUM_PARTICLES_Y 30
+#define NUM_PARTICLES_Z 30
+
+	
+
+ATTRIBUTE_ALIGNED16(struct) btSimParams
+{
+	BT_DECLARE_ALIGNED_ALLOCATOR();
+	btVector3	m_gravity;
+	float m_worldMin[4];
+	float m_worldMax[4];
+	
+	float m_particleRad;
+	float m_globalDamping;
+	float m_boundaryDamping;
+	float m_collisionDamping;
+
+	float m_spring;
+	float m_shear;
+	float m_attraction;
+	float m_dummy;
+
+		
+	btSimParams()
+	{
+		m_gravity.setValue(0,-0.03,0.f);
+		m_particleRad = 0.023f;
+		m_globalDamping = 1.0f;
+		m_boundaryDamping = -0.5f;
+		m_collisionDamping = 0.025f;//0.02f;
+		m_spring = 0.5f;
+		m_shear = 0.1f;
+		m_attraction = 0.001f;
+		m_worldMin[0] = -1.f;
+		m_worldMin[1] = -2*m_particleRad;
+		m_worldMin[2] =-1.f;
+
+		m_worldMax[0] = 5.f;
+		m_worldMax[1] = 5.f;
+		m_worldMax[2] = 5.f;
+
+	}
+};
+
+struct ParticleInternalData
+{
+	cl_context m_clContext;
+	cl_device_id m_clDevice;
+	cl_command_queue m_clQueue;
+	cl_kernel m_updatePositionsKernel;
+	cl_kernel m_updatePositionsKernel2;
+
+	cl_kernel m_updateAabbsKernel;
+
+	cl_kernel m_collideParticlesKernel;
+
+	btGpuSapBroadphase*	m_broadphaseGPU;
+	
+
+	cl_mem		m_clPositionBuffer;
+
+	btAlignedObjectArray<btVector3> m_velocitiesCPU;
+	btOpenCLArray<btVector3>*	m_velocitiesGPU;
+
+	btAlignedObjectArray<btSimParams>	m_simParamCPU;
+	btOpenCLArray<btSimParams>*	m_simParamGPU;
+
+	bool m_clInitialized;
+
+	ParticleInternalData()
+		:m_clInitialized(false),
+		m_clPositionBuffer(0),
+		m_velocitiesGPU(0),
+		m_simParamGPU(0),
+		m_updatePositionsKernel(0),
+		m_updatePositionsKernel2(0),
+		m_updateAabbsKernel(0),
+		m_collideParticlesKernel(0)
+	{
+		m_simParamCPU.resize(1);
+	}
+
+	char*	m_clDeviceName;
+
+};
+
+
+ParticleDemo::ParticleDemo()
+:m_instancingRenderer(0)
+{
+	m_data = new ParticleInternalData;
+}
+
+ParticleDemo::~ParticleDemo()
+{
+	exitCL();
+
+	delete m_data;
+
+}
+
+void ParticleDemo::exitCL()
+{
+	if (m_data->m_clInitialized)
+	{
+		m_data->m_clInitialized = false;
+		clReleaseCommandQueue(m_data->m_clQueue);
+		clReleaseKernel(m_data->m_updatePositionsKernel);
+		clReleaseKernel(m_data->m_updatePositionsKernel2);
+		clReleaseKernel(m_data->m_updateAabbsKernel);
+		clReleaseKernel(m_data->m_collideParticlesKernel);
+
+		clReleaseContext(m_data->m_clContext);
+	}
+}
+
+void ParticleDemo::initCL(int preferredDeviceIndex, int preferredPlatformIndex)
+{
+	void* glCtx=0;
+	void* glDC = 0;
+
+
+    
+	int ciErrNum = 0;
+//#ifdef CL_PLATFORM_INTEL
+//	cl_device_type deviceType = CL_DEVICE_TYPE_ALL;
+//#else
+	cl_device_type deviceType = CL_DEVICE_TYPE_GPU;
+//#endif
+
+	
+
+//	if (useInterop)
+//	{
+//		m_data->m_clContext = btOpenCLUtils::createContextFromType(deviceType, &ciErrNum, glCtx, glDC);
+//	} else
+	{
+		m_data->m_clContext = btOpenCLUtils::createContextFromType(deviceType, &ciErrNum, 0,0,preferredDeviceIndex, preferredPlatformIndex);
+	}
+
+
+	oclCHECKERROR(ciErrNum, CL_SUCCESS);
+
+	int numDev = btOpenCLUtils::getNumDevices(m_data->m_clContext);
+
+	if (numDev>0)
+	{
+		m_data->m_clDevice= btOpenCLUtils::getDevice(m_data->m_clContext,0);
+		m_data->m_clQueue = clCreateCommandQueue(m_data->m_clContext, m_data->m_clDevice, 0, &ciErrNum);
+		oclCHECKERROR(ciErrNum, CL_SUCCESS);
+        
+        btOpenCLUtils::printDeviceInfo(m_data->m_clDevice);
+		btOpenCLDeviceInfo info;
+		btOpenCLUtils::getDeviceInfo(m_data->m_clDevice,&info);
+		m_data->m_clDeviceName = info.m_deviceName;
+		m_data->m_clInitialized = true;
+
+	}
+
+}
+
+
+void ParticleDemo::setupScene(const ConstructionInfo& ci)
+{
+
+	initCL(ci.preferredOpenCLDeviceIndex,ci.preferredOpenCLPlatformIndex);
+	
+	int numParticles = NUM_PARTICLES_X*NUM_PARTICLES_Y*NUM_PARTICLES_Z;
+
+	btOverlappingPairCache* overlappingPairCache=0;
+	int maxObjects = NUM_PARTICLES_X*NUM_PARTICLES_Y*NUM_PARTICLES_Z+1024;
+	
+	int maxPairsSmallProxy = 32;
+	float radius = 3.f*m_data->m_simParamCPU[0].m_particleRad;
+
+	m_data->m_broadphaseGPU = new btGpuSapBroadphase(m_data->m_clContext ,m_data->m_clDevice,m_data->m_clQueue);//overlappingPairCache,btVector3(4.f, 4.f, 4.f), 128, 128, 128,maxObjects, maxObjects, maxPairsSmallProxy, 100.f, 128,
+
+	/*m_data->m_broadphaseGPU = new btGridBroadphaseCl(overlappingPairCache,btVector3(radius,radius,radius), 128, 128, 128,
+		maxObjects, maxObjects, maxPairsSmallProxy, 100.f, 128,
+			m_data->m_clContext ,m_data->m_clDevice,m_data->m_clQueue);
+			*/
+
+	m_data->m_velocitiesGPU = new btOpenCLArray<btVector3>(m_data->m_clContext,m_data->m_clQueue,numParticles);
+	m_data->m_velocitiesCPU.resize(numParticles);
+	for (int i=0;i<numParticles;i++)
+	{
+		m_data->m_velocitiesCPU[i].setValue(0,0,0);
+	}
+	m_data->m_velocitiesGPU->copyFromHost(m_data->m_velocitiesCPU);
+
+	m_data->m_simParamGPU = new btOpenCLArray<btSimParams>(m_data->m_clContext,m_data->m_clQueue,1,false);
+	m_data->m_simParamGPU->copyFromHost(m_data->m_simParamCPU);
+
+	cl_int pErrNum;
+
+	cl_program prog = btOpenCLUtils::compileCLProgramFromString(m_data->m_clContext,m_data->m_clDevice,particleKernelsString,0,"",INTEROPKERNEL_SRC_PATH);
+	m_data->m_updatePositionsKernel = btOpenCLUtils::compileCLKernelFromString(m_data->m_clContext, m_data->m_clDevice,particleKernelsString, "updatePositionsKernel" ,&pErrNum,prog);
+	oclCHECKERROR(pErrNum, CL_SUCCESS);
+	m_data->m_updatePositionsKernel2 = btOpenCLUtils::compileCLKernelFromString(m_data->m_clContext, m_data->m_clDevice,particleKernelsString, "integrateMotionKernel" ,&pErrNum,prog);
+	oclCHECKERROR(pErrNum, CL_SUCCESS);
+
+	m_data->m_updateAabbsKernel= btOpenCLUtils::compileCLKernelFromString(m_data->m_clContext, m_data->m_clDevice,particleKernelsString, "updateAabbsKernel" ,&pErrNum,prog);
+	oclCHECKERROR(pErrNum, CL_SUCCESS);
+
+	m_data->m_collideParticlesKernel = btOpenCLUtils::compileCLKernelFromString(m_data->m_clContext, m_data->m_clDevice,particleKernelsString, "collideParticlesKernel" ,&pErrNum,prog);
+	oclCHECKERROR(pErrNum, CL_SUCCESS);
+
+	m_instancingRenderer = ci.m_instancingRenderer;
+
+	int strideInBytes = 9*sizeof(float);
+	bool pointSprite = true;
+	int shapeId =-1;
+
+	if (pointSprite)
+	{
+		int numVertices = sizeof(point_sphere_vertices)/strideInBytes;
+		int numIndices = sizeof(point_sphere_indices)/sizeof(int);
+		shapeId = m_instancingRenderer->registerShape(&point_sphere_vertices[0],numVertices,point_sphere_indices,numIndices,BT_GL_POINTS);
+	} else
+	{
+		int numVertices = sizeof(low_sphere_vertices)/strideInBytes;
+		int numIndices = sizeof(low_sphere_indices)/sizeof(int);
+		shapeId = m_instancingRenderer->registerShape(&low_sphere_vertices[0],numVertices,low_sphere_indices,numIndices);
+	}
+
+	float position[4] = {0,0,0,0};
+	float quaternion[4] = {0,0,0,1};
+	float color[4]={1,0,0,1};
+	float scaling[4] = {0.023,0.023,0.023,1};
+
+	int userIndex = 0;
+	for (int x=0;x<NUM_PARTICLES_X;x++)
+	{
+		for (int y=0;y<NUM_PARTICLES_Y;y++)
+		{
+			for (int z=0;z<NUM_PARTICLES_Z;z++)
+			{
+				float rad = m_data->m_simParamCPU[0].m_particleRad;
+				position[0] = x*(rad*3);
+				position[1] = y*(rad*3);
+				position[2] = z*(rad*3);
+
+				color[0] = float(x)/float(NUM_PARTICLES_X);
+				color[1] = float(y)/float(NUM_PARTICLES_Y);
+				color[2] = float(z)/float(NUM_PARTICLES_Z);
+
+				int id = m_instancingRenderer->registerGraphicsInstance(shapeId,position,quaternion,color,scaling);
+				
+				void* userPtr = (void*)userIndex;
+				int collidableIndex = userIndex;
+				btVector3 aabbMin,aabbMax;
+				btVector3 particleRadius(rad,rad,rad);
+
+				aabbMin = btVector3(position[0],position[1],position[2])-particleRadius;
+				aabbMax = btVector3(position[0],position[1],position[2])+particleRadius;
+				m_data->m_broadphaseGPU->createProxy(aabbMin,aabbMax,collidableIndex,1,1);
+				userIndex++;
+
+			}
+		}
+	}
+	m_data->m_broadphaseGPU->writeAabbsToGpu();
+
+	float camPos[4]={1.5,0.5,2.5,0};
+	m_instancingRenderer->setCameraTargetPosition(camPos);
+	m_instancingRenderer->setCameraDistance(4);
+	m_instancingRenderer->writeTransforms();
+
+}
+
+void	ParticleDemo::initPhysics(const ConstructionInfo& ci)
+{
+	setupScene(ci);
+}
+
+void	ParticleDemo::exitPhysics()
+{
+}
+
+void	ParticleDemo::renderScene()
+{
+	
+	if (m_instancingRenderer)
+	{
+		m_instancingRenderer->RenderScene();
+	}
+
+}
+
+
+void ParticleDemo::clientMoveAndDisplay()
+{
+	int numParticles = NUM_PARTICLES_X*NUM_PARTICLES_Y*NUM_PARTICLES_Z;
+	GLuint vbo = m_instancingRenderer->getInternalData()->m_vbo;
+	glBindBuffer(GL_ARRAY_BUFFER, vbo);
+	glFlush();
+
+	int posArraySize = numParticles*sizeof(float)*4;
+
+	cl_bool blocking=  CL_TRUE;
+	char* hostPtr=  (char*)glMapBufferRange( GL_ARRAY_BUFFER,m_instancingRenderer->getMaxShapeCapacity(),posArraySize, GL_MAP_WRITE_BIT|GL_MAP_READ_BIT );//GL_READ_WRITE);//GL_WRITE_ONLY
+		GLint err = glGetError();
+    assert(err==GL_NO_ERROR);
+	glFinish();
+
+	
+
+#if 1
+
+
+
+	//do some stuff using the OpenCL buffer
+
+	bool useCpu = false;
+	if (useCpu)
+	{
+		
+
+		float* posBuffer = (float*)hostPtr;
+		
+		for (int i=0;i<numParticles;i++)
+		{
+			posBuffer[i*4+1] += 0.1;
+		}
+	}
+	else
+	{
+		cl_int ciErrNum;
+		if (!m_data->m_clPositionBuffer)
+		{
+			m_data->m_clPositionBuffer = clCreateBuffer(m_data->m_clContext, CL_MEM_READ_WRITE,
+				posArraySize, 0, &ciErrNum);
+
+			clFinish(m_data->m_clQueue);
+			oclCHECKERROR(ciErrNum, CL_SUCCESS);
+			ciErrNum = clEnqueueWriteBuffer (	m_data->m_clQueue,m_data->m_clPositionBuffer,
+ 				blocking,0,posArraySize,hostPtr,0,0,0
+			);
+			clFinish(m_data->m_clQueue);
+		}
+	
+
+		
+
+
+
+		if (0)
+		{
+			btBufferInfoCL bInfo[] = { 
+				btBufferInfoCL( m_data->m_velocitiesGPU->getBufferCL(), true ),
+				btBufferInfoCL( m_data->m_clPositionBuffer)
+			};
+			
+			btLauncherCL launcher(m_data->m_clQueue, m_data->m_updatePositionsKernel );
+
+			launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) );
+			launcher.setConst( numParticles);
+
+			launcher.launch1D( numParticles);
+			clFinish(m_data->m_clQueue);
+	
+		}
+
+
+		if (1)
+		{
+			btBufferInfoCL bInfo[] = { 
+				btBufferInfoCL( m_data->m_clPositionBuffer),
+				btBufferInfoCL( m_data->m_velocitiesGPU->getBufferCL() ),
+				btBufferInfoCL( m_data->m_simParamGPU->getBufferCL(),true)
+			};
+			
+			btLauncherCL launcher(m_data->m_clQueue, m_data->m_updatePositionsKernel2 );
+
+			launcher.setConst( numParticles);
+			launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) );
+			float timeStep = 1.f/60.f;
+			launcher.setConst( timeStep);
+
+			launcher.launch1D( numParticles);
+			clFinish(m_data->m_clQueue);
+	
+		}
+
+		{
+			btBufferInfoCL bInfo[] = { 
+				btBufferInfoCL( m_data->m_clPositionBuffer),
+				btBufferInfoCL( m_data->m_broadphaseGPU->getAabbBuffer()),
+			};
+			
+			btLauncherCL launcher(m_data->m_clQueue, m_data->m_updateAabbsKernel );
+			launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) );
+			launcher.setConst( m_data->m_simParamCPU[0].m_particleRad);
+			launcher.setConst( numParticles);
+			
+			launcher.launch1D( numParticles);
+			clFinish(m_data->m_clQueue);
+		}
+
+		//broadphase
+		int numPairsGPU=0;
+		cl_mem pairsGPU  = 0;
+
+		{
+			m_data->m_broadphaseGPU->calculateOverlappingPairs();
+			pairsGPU = m_data->m_broadphaseGPU->getOverlappingPairBuffer();
+			numPairsGPU = m_data->m_broadphaseGPU->getNumOverlap();
+		}
+
+		if (numPairsGPU)
+		{
+			btBufferInfoCL bInfo[] = { 
+				btBufferInfoCL( m_data->m_clPositionBuffer),
+				btBufferInfoCL( m_data->m_velocitiesGPU->getBufferCL() ),
+				btBufferInfoCL( m_data->m_broadphaseGPU->getOverlappingPairBuffer(),true),
+			};
+			
+			btLauncherCL launcher(m_data->m_clQueue, m_data->m_collideParticlesKernel);
+			launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) );
+			launcher.setConst( numPairsGPU);
+			launcher.launch1D( numPairsGPU);
+			clFinish(m_data->m_clQueue);
+
+			//__kernel void collideParticlesKernel(  __global float4* pPos, __global float4* pVel, __global int2* pairs, const int numPairs)
+		}
+
+
+		if (1)
+		{
+			ciErrNum = clEnqueueReadBuffer (	m_data->m_clQueue,
+				m_data->m_clPositionBuffer,
+	 			blocking,
+ 				0,
+ 				posArraySize,
+ 			hostPtr,0,0,0);
+
+			//clReleaseMemObject(clBuffer);
+			clFinish(m_data->m_clQueue);
+
+			
+		}
+	}
+	
+#endif
+
+	glUnmapBuffer( GL_ARRAY_BUFFER);
+	glFlush();
+
+	/*
+	int numParticles = NUM_PARTICLES_X*NUM_PARTICLES_Y*NUM_PARTICLES_Z;
+	for (int objectIndex=0;objectIndex<numParticles;objectIndex++)
+	{
+		float pos[4]={0,0,0,0};
+		float orn[4]={0,0,0,1};
+
+//		m_instancingRenderer->writeSingleInstanceTransformToGPU(pos,orn,i);
+		{
+			glBindBuffer(GL_ARRAY_BUFFER, m_instancingRenderer->getInternalData()->m_vbo);
+			glFlush();
+
+			char* orgBase =  (char*)glMapBuffer( GL_ARRAY_BUFFER,GL_READ_WRITE);
+			//btGraphicsInstance* gfxObj = m_graphicsInstances[k];
+			int totalNumInstances= numParticles;
+	
+
+			int POSITION_BUFFER_SIZE = (totalNumInstances*sizeof(float)*4);
+
+			char* base = orgBase;
+			int capInBytes = m_instancingRenderer->getMaxShapeCapacity();
+
+			float* positions = (float*)(base+capInBytes);
+			float* orientations = (float*)(base+capInBytes+ POSITION_BUFFER_SIZE);
+
+			positions[objectIndex*4+1] += 0.1f;
+			glUnmapBuffer( GL_ARRAY_BUFFER);
+			glFlush();
+		}
+	}
+	*/
+
+	
+}
+
+//	m_data->m_positionOffsetInBytes = demo.m_maxShapeBufferCapacity/4;

demo/gpudemo/ParticleDemo.h

@@ -0,0 +1,50 @@
+#ifndef PARTICLE_DEMO_H
+#define PARTICLE_DEMO_H
+
+#include "GpuDemo.h"
+
+class ParticleDemo : public GpuDemo
+{
+	
+protected:
+
+	struct ParticleInternalData*	m_data;
+
+	GLInstancingRenderer*	m_instancingRenderer;
+
+	void initCL(int preferredDeviceIndex, int preferredPlatformIndex);
+	void exitCL();
+
+public:
+
+	ParticleDemo();
+
+	virtual ~ParticleDemo();
+
+	virtual void setupScene(const ConstructionInfo& ci);
+
+	virtual void	initPhysics(const ConstructionInfo& ci);
+
+	virtual void	exitPhysics();
+
+	virtual const char* getName()
+	{
+		return "ParticleDemo";
+	}
+	static GpuDemo* CreateFunc()
+	{
+		GpuDemo* demo = new ParticleDemo;
+		return demo;
+	}
+
+	virtual const btDynamicsWorld* getDynamicsWorld() const
+	{
+		return 0;
+	}
+
+	virtual void renderScene();
+
+	virtual void clientMoveAndDisplay();
+};
+
+#endif //PARTICLE_DEMO_H

demo/gpudemo/ParticleKernels.cl

@@ -0,0 +1,149 @@
+MSTRINGIFY(
+
+
+typedef struct
+{
+	float4	m_gravity;
+	float4	m_worldMin;
+	float4	m_worldMax;
+	
+	float m_particleRad;
+	float m_globalDamping;
+	float m_boundaryDamping;
+	float m_collisionDamping;
+
+	float m_spring;
+	float m_shear;
+	float m_attraction;
+	float m_dummy;
+} btSimParams;
+
+
+__kernel void  updatePositionsKernel( __global float4* linearVelocities, __global float4* positions,const int numNodes)
+{
+	int nodeID = get_global_id(0);
+	float timeStep = 0.0166666;
+	
+	float BT_GPU_ANGULAR_MOTION_THRESHOLD = (0.25f * 3.14159254);
+	
+	if( nodeID < numNodes )
+	{
+		positions[nodeID] += linearVelocities[nodeID]*timeStep;
+	}
+}
+
+
+
+__kernel void integrateMotionKernel(	int numParticles,
+								__global float4* pPos, 
+								__global float4* pVel, 
+								__global const btSimParams* simParams,
+								float timeStep GUID_ARG)
+{
+    int index = get_global_id(0);
+    if(index >= numParticles)
+    {
+		return;
+	}
+    float4 pos = pPos[index];
+    float4 vel = pVel[index];
+    pos.w = 1.0f;
+    vel.w = 0.0f;
+    // apply gravity
+    float4 gravity = simParams[0].m_gravity;
+    float particleRad = simParams[0].m_particleRad;
+    float globalDamping = simParams[0].m_globalDamping;
+    float boundaryDamping = simParams[0].m_boundaryDamping;
+
+    vel += gravity * timeStep;
+    vel *= globalDamping;
+    // integrate position
+    pos += vel * timeStep;
+    // collide with world boundaries
+    float4 worldMin = simParams[0].m_worldMin;
+    float4 worldMax = simParams[0].m_worldMax;
+    
+    
+    if(pos.x < (worldMin.x + 2*particleRad))
+    {
+        pos.x = worldMin.x + 2*particleRad;
+        vel.x *= boundaryDamping;
+    }
+    if(pos.x > (worldMax.x - 2*particleRad))
+    {
+        pos.x = worldMax.x - 2*particleRad;
+        vel.x *= boundaryDamping;
+    }
+    if(pos.y < (worldMin.y + 2*particleRad))
+    {
+        pos.y = worldMin.y + 2*particleRad;
+        vel.y *= boundaryDamping;
+    }
+    if(pos.y > (worldMax.y - 2*particleRad))
+    {
+        pos.y = worldMax.y - 2*particleRad;
+        vel.y *= boundaryDamping;
+    }
+    if(pos.z < (worldMin.z + 2*particleRad))
+    {
+        pos.z = worldMin.z + 2*particleRad;
+        vel.z *= boundaryDamping;
+    }
+    if(pos.z > (worldMax.z - 2*particleRad))
+    {
+        pos.z = worldMax.z - 2*particleRad;
+        vel.z *= boundaryDamping;
+    }
+    // write back position and velocity
+    pPos[index] = pos;
+    pVel[index] = vel;
+}
+
+typedef struct 
+{
+	float			fx;
+	float			fy;
+	float			fz;
+	int	uw;
+} btAABBCL;
+
+__kernel void collideParticlesKernel(  __global float4* pPos, __global float4* pVel, __global int2* pairs, const int numPairs)
+{
+	int index = get_global_id(0);
+	if (index<numPairs)
+	{
+		pPos[pairs[index].x].x = -30*index;
+		pPos[pairs[index].x].z = -30*index;
+
+		pPos[pairs[index].y].x = 30*index;
+		pPos[pairs[index].y].z = 30*index;
+
+	}
+	
+}
+
+
+__kernel void updateAabbsKernel(  __global float4* pPos, __global btAABBCL* pAABB, float particleRadius, const int numNodes)
+{
+	int nodeID = get_global_id(0);
+		
+	if( nodeID < numNodes )
+	{
+		float4 position = pPos[nodeID];
+		float4 extent = (float4) (	particleRadius,particleRadius,particleRadius,0.f);
+			
+	
+		pAABB[nodeID*2].fx = position.x-extent.x;
+		pAABB[nodeID*2].fy = position.y-extent.y;
+		pAABB[nodeID*2].fz = position.z-extent.z;
+		pAABB[nodeID*2].uw = nodeID;
+	
+		pAABB[nodeID*2+1].fx = position.x+extent.x;
+		pAABB[nodeID*2+1].fy = position.y+extent.y;
+		pAABB[nodeID*2+1].fz = position.z+extent.z;
+		pAABB[nodeID*2+1].uw = nodeID;
+	} 
+}
+
+
+);

demo/gpudemo/gwenUserInterface.cpp

@@ -0,0 +1,346 @@
+
+#include "gwenUserInterface.h"
+#include "OpenGLWindow/GwenOpenGL3CoreRenderer.h"
+#include "OpenGLWindow/GLPrimitiveRenderer.h"
+#include "Gwen/Platform.h"
+#include "Gwen/Controls/TreeControl.h"
+#include "Gwen/Controls/RadioButtonController.h"
+#include "Gwen/Controls/VerticalSlider.h"
+#include "Gwen/Controls/HorizontalSlider.h"
+#include "Gwen/Controls/GroupBox.h"
+#include "Gwen/Controls/CheckBox.h"
+#include "Gwen/Controls/StatusBar.h"
+#include "Gwen/Controls/Button.h"
+#include "Gwen/Controls/ComboBox.h"
+#include "Gwen/Controls/MenuStrip.h"
+#include "Gwen/Controls/Property/Text.h"
+#include "Gwen/Controls/SplitterBar.h"
+#include "BulletCommon/btAlignedObjectArray.h"
+#include "Gwen/Gwen.h"
+#include "Gwen/Align.h"
+#include "Gwen/Utility.h"
+#include "Gwen/Controls/WindowControl.h"
+#include "Gwen/Controls/TabControl.h"
+#include "Gwen/Controls/ListBox.h"
+#include "Gwen/Skins/Simple.h"
+//#include "Gwen/Skins/TexturedBase.h"
+
+
+struct GwenInternalData
+{
+	struct sth_stash;
+	class GwenOpenGL3CoreRenderer*	pRenderer;
+	Gwen::Skin::Simple				skin;
+	Gwen::Controls::Canvas*			pCanvas;
+	GLPrimitiveRenderer* m_primRenderer;
+	Gwen::Controls::TabButton*		m_demoPage;
+
+	Gwen::Controls::Label* m_rightStatusBar;
+	Gwen::Controls::Label* m_leftStatusBar;
+
+	btAlignedObjectArray<struct Gwen::Event::Handler*>	m_handlers;
+	btToggleButtonCallback			m_toggleButtonCallback;
+	btComboBoxCallback				m_comboBoxCallback;
+
+};
+GwenUserInterface::GwenUserInterface()
+{
+	m_data = new GwenInternalData();
+	m_data->m_toggleButtonCallback = 0;
+	m_data->m_comboBoxCallback = 0;
+
+}
+		
+GwenUserInterface::~GwenUserInterface()
+{
+	delete m_data;
+}
+		
+
+struct MyTestMenuBar : public Gwen::Controls::MenuStrip
+{
+	
+
+
+	MyTestMenuBar(Gwen::Controls::Base* pParent)
+		:Gwen::Controls::MenuStrip(pParent)
+	{
+//		Gwen::Controls::MenuStrip* menu = new Gwen::Controls::MenuStrip( pParent );
+		{
+			Gwen::Controls::MenuItem* pRoot = AddItem( L"File" );
+		
+			pRoot = AddItem( L"View" );
+//			Gwen::Event::Handler* handler =	GWEN_MCALL(&MyTestMenuBar::MenuItemSelect );
+			pRoot->GetMenu()->AddItem( L"Profiler");//,,m_profileWindow,(Gwen::Event::Handler::Function)&MyProfileWindow::MenuItemSelect);
+
+/*			pRoot->GetMenu()->AddItem( L"New", L"test16.png", GWEN_MCALL( ThisClass::MenuItemSelect ) );
+			pRoot->GetMenu()->AddItem( L"Load", L"test16.png", GWEN_MCALL( ThisClass::MenuItemSelect ) );
+			pRoot->GetMenu()->AddItem( L"Save", GWEN_MCALL( ThisClass::MenuItemSelect ) );
+			pRoot->GetMenu()->AddItem( L"Save As..", GWEN_MCALL( ThisClass::MenuItemSelect ) );
+			pRoot->GetMenu()->AddItem( L"Quit", GWEN_MCALL( ThisClass::MenuItemSelect ) );
+			*/
+		}
+	}
+
+};
+
+void	GwenUserInterface::resize(int width, int height)
+{
+	m_data->pCanvas->SetSize(width,height);
+}
+
+
+struct MyComboBoxHander   :public Gwen::Event::Handler
+{
+	GwenInternalData*	m_data;
+	int					m_buttonId;
+
+	MyComboBoxHander  (GwenInternalData* data, int buttonId)
+		:m_data(data),
+		m_buttonId(buttonId)
+	{
+	}
+
+	void onSelect( Gwen::Controls::Base* pControl )
+	{
+		Gwen::Controls::ComboBox* but = (Gwen::Controls::ComboBox*) pControl;
+		
+		
+		
+		Gwen::String str = Gwen::Utility::UnicodeToString(	but->GetSelectedItem()->GetText());
+		
+		if (m_data->m_comboBoxCallback)
+			(*m_data->m_comboBoxCallback)(m_buttonId,str.c_str());
+	}
+
+};
+
+struct MyButtonHander   :public Gwen::Event::Handler
+{
+	GwenInternalData*	m_data;
+	int					m_buttonId;
+
+	MyButtonHander  (GwenInternalData* data, int buttonId)
+		:m_data(data),
+		m_buttonId(buttonId)
+	{
+	}
+
+	void onButtonA( Gwen::Controls::Base* pControl )
+	{
+		Gwen::Controls::Button* but = (Gwen::Controls::Button*) pControl;
+		int dep = but->IsDepressed();
+		int tog = but->GetToggleState();
+		if (m_data->m_toggleButtonCallback)
+			(*m_data->m_toggleButtonCallback)(m_buttonId,tog);
+	}
+
+};
+
+
+void	GwenUserInterface::setStatusBarMessage(const char* message, bool isLeft)
+{
+	Gwen::UnicodeString msg = Gwen::Utility::StringToUnicode(message);
+	if (isLeft)
+	{
+		m_data->m_leftStatusBar->SetText( msg);
+	} else
+	{
+		m_data->m_rightStatusBar->SetText( msg);
+		
+	}
+}
+
+void	GwenUserInterface::init(int width, int height,struct sth_stash* stash,float retinaScale)
+{
+	m_data->m_primRenderer = new GLPrimitiveRenderer(width,height);
+	m_data->pRenderer = new GwenOpenGL3CoreRenderer(m_data->m_primRenderer,stash,width,height,retinaScale);
+
+	m_data->skin.SetRender( m_data->pRenderer );
+
+	m_data->pCanvas= new Gwen::Controls::Canvas( &m_data->skin );
+	m_data->pCanvas->SetSize( width,height);
+	m_data->pCanvas->SetDrawBackground( false);
+	m_data->pCanvas->SetBackgroundColor( Gwen::Color( 150, 170, 170, 255 ) );
+
+	MyTestMenuBar* menubar = new MyTestMenuBar(m_data->pCanvas);
+	Gwen::Controls::StatusBar* bar = new Gwen::Controls::StatusBar(m_data->pCanvas);
+	m_data->m_rightStatusBar = new Gwen::Controls::Label( bar );
+	m_data->m_rightStatusBar->SetWidth(width/2);
+	//m_data->m_rightStatusBar->SetText( L"Label Added to Right" );
+	bar->AddControl( m_data->m_rightStatusBar, true );
+
+	m_data->m_leftStatusBar = new Gwen::Controls::Label( bar );
+	//m_data->m_leftStatusBar->SetText( L"Label Added to Left" );
+	m_data->m_leftStatusBar->SetWidth(width/2);
+	bar->AddControl( m_data->m_leftStatusBar,false);
+
+	/*Gwen::Controls::GroupBox* box = new Gwen::Controls::GroupBox(m_data->pCanvas);
+	box->SetText("text");
+	box->SetName("name");
+	box->SetHeight(500);
+	*/
+	Gwen::Controls::ScrollControl* windowLeft= new Gwen::Controls::ScrollControl(m_data->pCanvas);
+	windowLeft->Dock(Gwen::Pos::Right);
+	windowLeft->SetWidth(150);
+	windowLeft->SetHeight(250);
+	windowLeft->SetScroll(false,true);
+
+	/*Gwen::Controls::WindowControl* windowLeft = new Gwen::Controls::WindowControl(m_data->pCanvas);
+	windowLeft->Dock(Gwen::Pos::Left);
+	windowLeft->SetTitle("title");
+	windowLeft->SetWidth(150);
+	windowLeft->SetClosable(false);
+	windowLeft->SetShouldDrawBackground(true);
+	windowLeft->SetTabable(true);
+	*/
+
+	//windowLeft->SetSkin(
+	Gwen::Controls::TabControl* tab = new Gwen::Controls::TabControl(windowLeft);
+	
+	//tab->SetHeight(300);
+	tab->SetWidth(140);
+	tab->SetHeight(250);
+	//tab->Dock(Gwen::Pos::Left);
+	tab->Dock( Gwen::Pos::Fill );
+	//tab->SetMargin( Gwen::Margin( 2, 2, 2, 2 ) );
+
+	Gwen::UnicodeString str1(L"Main");
+	m_data->m_demoPage = tab->AddPage(str1);
+
+	
+
+	
+	Gwen::UnicodeString str2(L"OpenCL");
+	tab->AddPage(str2);
+	//Gwen::UnicodeString str3(L"page3");
+//	tab->AddPage(str3);
+	
+		
+	
+	//but->onPress.Add(handler, &MyHander::onButtonA);
+
+	
+	
+	//box->Dock(Gwen::Pos::Left);
+
+	/*Gwen::Controls::WindowControl* windowBottom = new Gwen::Controls::WindowControl(m_data->pCanvas);
+	windowBottom->SetHeight(100);
+	windowBottom->Dock(Gwen::Pos::Bottom);
+	windowBottom->SetTitle("bottom");
+	*/
+//	Gwen::Controls::Property::Text* prop = new Gwen::Controls::Property::Text(m_data->pCanvas);
+	//prop->Dock(Gwen::Pos::Bottom);
+	/*Gwen::Controls::SplitterBar* split = new Gwen::Controls::SplitterBar(m_data->pCanvas);
+	split->Dock(Gwen::Pos::Center);
+	split->SetHeight(300);
+	split->SetWidth(300);
+	*/
+	/*
+	
+	
+	*/
+}
+	
+
+void	GwenUserInterface::setToggleButtonCallback(btToggleButtonCallback callback)
+{
+	m_data->m_toggleButtonCallback = callback;
+}
+void	GwenUserInterface::registerToggleButton(int buttonId, const char* name)
+{
+	assert(m_data);
+	assert(m_data->m_demoPage);
+
+	Gwen::Controls::Button* but = new Gwen::Controls::Button(m_data->m_demoPage->GetPage());
+	
+	///some heuristic to find the button location
+	int ypos = m_data->m_handlers.size()*20;
+	but->SetPos(10, ypos );
+	but->SetWidth( 100 );
+	//but->SetBounds( 200, 30, 300, 200 );
+	
+	MyButtonHander* handler = new MyButtonHander(m_data, buttonId);
+	m_data->m_handlers.push_back(handler);
+	but->onToggle.Add(handler, &MyButtonHander::onButtonA);
+	but->SetIsToggle(true);
+	but->SetToggleState(false);
+	but->SetText(name);
+
+}
+
+void	GwenUserInterface::setComboBoxCallback(btComboBoxCallback callback)
+{
+	m_data->m_comboBoxCallback = callback;
+}
+
+void	GwenUserInterface::registerComboBox(int comboboxId, int numItems, const char** items)
+{
+	Gwen::Controls::ComboBox* combobox = new Gwen::Controls::ComboBox(m_data->m_demoPage->GetPage());
+	MyComboBoxHander* handler = new MyComboBoxHander(m_data, comboboxId);
+	m_data->m_handlers.push_back(handler);
+	
+	combobox->onSelection.Add(handler,&MyComboBoxHander::onSelect);
+	int ypos = m_data->m_handlers.size()*20;
+	combobox->SetPos(10, ypos );
+	combobox->SetWidth( 100 );
+	//box->SetPos(120,130);
+	for (int i=0;i<numItems;i++)
+		combobox->AddItem(Gwen::Utility::StringToUnicode(items[i]));
+	
+	
+}
+
+void	GwenUserInterface::draw(int width, int height)
+{
+	
+//	printf("width = %d, height=%d\n", width,height);
+	if (m_data->pCanvas)
+	{
+		m_data->pCanvas->SetSize(width,height);
+		m_data->m_primRenderer->setScreenSize(width,height);
+		m_data->pRenderer->resize(width,height);
+		m_data->pCanvas->RenderCanvas();
+		//restoreOpenGLState();
+	}
+
+}
+
+bool	GwenUserInterface::mouseMoveCallback( float x, float y)
+{
+	bool handled  = false;
+
+	static int m_lastmousepos[2] = {0,0};
+	static bool isInitialized = false;
+	if (m_data->pCanvas)
+	{
+		if (!isInitialized)
+		{
+			isInitialized = true;
+			m_lastmousepos[0] = x+1;
+			m_lastmousepos[1] = y+1;
+		}
+		handled = m_data->pCanvas->InputMouseMoved(x,y,m_lastmousepos[0],m_lastmousepos[1]);
+	}
+	return handled;
+	
+}
+bool	GwenUserInterface::mouseButtonCallback(int button, int state, float x, float y)
+{
+	bool handled = false;
+	if (m_data->pCanvas)
+	{
+		handled = m_data->pCanvas->InputMouseMoved(x,y,x, y);
+
+		if (button>=0)
+		{
+			handled = m_data->pCanvas->InputMouseButton(button,state);
+			if (handled)
+			{
+				//if (!state)
+				//	return false;
+			}
+		}
+	}
+	return handled;
+}

demo/gpudemo/gwenUserInterface.h

@@ -0,0 +1,38 @@
+#ifndef _GWEN_USER_INTERFACE_H
+#define _GWEN_USER_INTERFACE_H
+
+struct GwenInternalData;
+
+typedef void (*btComboBoxCallback) (int combobox, const char* item);
+typedef void (*btToggleButtonCallback)(int button, int state);
+
+class GwenUserInterface
+{
+	GwenInternalData*	m_data;
+
+	public:
+		
+		GwenUserInterface();
+		
+		virtual ~GwenUserInterface();
+		
+		void	init(int width, int height,struct sth_stash* stash,float retinaScale);
+		
+		void	draw(int width, int height);
+
+		void	resize(int width, int height);
+				
+		bool	mouseMoveCallback( float x, float y);
+		bool	mouseButtonCallback(int button, int state, float x, float y);
+
+		void	setToggleButtonCallback(btToggleButtonCallback callback);
+		void	registerToggleButton(int buttonId, const char* name);
+
+		void	setComboBoxCallback(btComboBoxCallback callback);
+		void	registerComboBox(int buttonId, int numItems, const char** items);
+
+		void	setStatusBarMessage(const char* message, bool isLeft=true);
+};
+
+#endif //_GWEN_USER_INTERFACE_H
+

demo/gpudemo/main_opengl3core.cpp

@@ -0,0 +1,664 @@
+
+#include "GpuDemo.h"
+
+#ifdef _WIN32
+#include <Windows.h> //for GetLocalTime/GetSystemTime
+#endif
+
+#ifdef __APPLE__
+#include "MacOpenGLWindow.h"
+#elif defined _WIN32
+#include "OpenGLWindow/Win32OpenGLWindow.h"
+#elif defined __linux
+#include "../rendering/rendertest/X11OpenGLWindow.h"
+#endif
+
+#include "OpenGLWindow/GLPrimitiveRenderer.h"
+#include "OpenGLWindow/GLInstancingRenderer.h"
+#include "OpenGL3CoreRenderer.h"
+#include "BulletCommon/btQuickprof.h"
+//#include "btGpuDynamicsWorld.h"
+#include <assert.h>
+#include <string.h>
+#include "OpenGLTrueTypeFont/fontstash.h"
+#include "OpenGLTrueTypeFont/opengl_fontstashcallbacks.h"
+#include "gwenUserInterface.h"
+#include "ParticleDemo.h"
+//#include "BroadphaseBenchmark.h"
+
+int g_OpenGLWidth=1024;
+int g_OpenGLHeight = 768;
+bool dump_timings = false;
+
+static void MyResizeCallback( float width, float height)
+{
+	g_OpenGLWidth = width;
+	g_OpenGLHeight = height;
+}
+
+btgWindowInterface* window=0;
+GwenUserInterface* gui  = 0;
+bool gPause = false;
+bool gReset = false;
+
+enum
+{
+	MYPAUSE=1,
+	MYPROFILE=2,
+	MYRESET,
+};
+
+enum
+{
+	MYCOMBOBOX1 = 1,
+};
+
+btAlignedObjectArray<const char*> demoNames;
+int selectedDemo = 0;
+GpuDemo::CreateFunc* allDemos[]=
+{
+	//BroadphaseBenchmark::CreateFunc,
+	//GpuBoxDemo::CreateFunc,
+		ParticleDemo::CreateFunc,
+	//SpheresDemo::CreateFunc,
+	//GpuCompoundDemo::CreateFunc,
+	//EmptyDemo::CreateFunc,
+};
+
+
+void	MyComboBoxCallback(int comboId, const char* item)
+{
+	int numDemos = demoNames.size();
+	for (int i=0;i<numDemos;i++)
+	{
+		if (!strcmp(demoNames[i],item))
+		{
+			if (selectedDemo != i)
+			{
+				gReset = true;
+				selectedDemo = i;
+				printf("selected demo %s!\n", item);
+			}
+		}
+	}
+
+
+}
+void	MyButtonCallback(int buttonId, int state)
+{
+	switch (buttonId)
+	{
+	case MYPAUSE:
+		{
+		gPause =!gPause;
+		break;
+		}
+	case MYPROFILE:
+		{
+		dump_timings = !dump_timings;
+		break;
+		}
+	case MYRESET:
+		{
+		gReset=!gReset;
+		break;
+		}
+	default:
+		{
+			printf("hello\n");
+		}
+	}
+}
+
+static void MyMouseMoveCallback( float x, float y)
+{
+	if (gui)
+	{
+		bool handled = gui ->mouseMoveCallback(x,y);
+		if (!handled)
+			btDefaultMouseMoveCallback(x,y);
+	}
+}
+static void MyMouseButtonCallback(int button, int state, float x, float y)
+{
+	if (gui)
+	{
+		bool handled = gui->mouseButtonCallback(button,state,x,y);
+		if (!handled)
+			btDefaultMouseButtonCallback(button,state,x,y);
+	}
+}
+
+
+void MyKeyboardCallback(int key, int state)
+{
+	if (key==BTG_ESCAPE && window)
+	{
+		window->setRequestExit();
+	}
+	btDefaultKeyboardCallback(key,state);
+}
+
+
+
+extern bool enableExperimentalCpuConcaveCollision;
+
+
+
+
+	int droidRegular, droidItalic, droidBold, droidJapanese, dejavu;
+
+sth_stash* stash=0;
+
+sth_stash* initFont(GLPrimitiveRenderer* primRender)
+{
+	GLint err;
+
+		struct sth_stash* stash = 0;
+	int datasize;
+	unsigned char* data;
+	float sx,sy,dx,dy,lh;
+	GLuint texture;
+
+	OpenGL2RenderCallbacks* renderCallbacks = new OpenGL2RenderCallbacks(primRender);
+
+	stash = sth_create(512,512,renderCallbacks);//256,256);//,1024);//512,512);
+    err = glGetError();
+    assert(err==GL_NO_ERROR);
+
+	if (!stash)
+	{
+		fprintf(stderr, "Could not create stash.\n");
+		return 0;
+	}
+
+	const char* fontPaths[]={
+	"./",
+	"../../bin/",
+	"../bin/",
+	"bin/"
+	};
+
+	int numPaths=sizeof(fontPaths)/sizeof(char*);
+
+	// Load the first truetype font from memory (just because we can).
+
+	FILE* fp = 0;
+	const char* fontPath ="./";
+	char fullFontFileName[1024];
+
+	for (int i=0;i<numPaths;i++)
+	{
+
+		fontPath = fontPaths[i];
+		//sprintf(fullFontFileName,"%s%s",fontPath,"OpenSans.ttf");//"DroidSerif-Regular.ttf");
+		sprintf(fullFontFileName,"%s%s",fontPath,"DroidSerif-Regular.ttf");//OpenSans.ttf");//"DroidSerif-Regular.ttf");
+		fp = fopen(fullFontFileName, "rb");
+		if (fp)
+			break;
+	}
+
+    err = glGetError();
+    assert(err==GL_NO_ERROR);
+
+    assert(fp);
+    if (fp)
+    {
+        fseek(fp, 0, SEEK_END);
+        datasize = (int)ftell(fp);
+        fseek(fp, 0, SEEK_SET);
+        data = (unsigned char*)malloc(datasize);
+        if (data == NULL)
+        {
+            assert(0);
+            return 0;
+        }
+        else
+            fread(data, 1, datasize, fp);
+        fclose(fp);
+        fp = 0;
+    }
+	if (!(droidRegular = sth_add_font_from_memory(stash, data)))
+    {
+        assert(0);
+        return 0;
+    }
+    err = glGetError();
+    assert(err==GL_NO_ERROR);
+
+	// Load the remaining truetype fonts directly.
+    sprintf(fullFontFileName,"%s%s",fontPath,"DroidSerif-Italic.ttf");
+
+	if (!(droidItalic = sth_add_font(stash,fullFontFileName)))
+	{
+        assert(0);
+        return 0;
+    }
+     sprintf(fullFontFileName,"%s%s",fontPath,"DroidSerif-Bold.ttf");
+
+	if (!(droidBold = sth_add_font(stash,fullFontFileName)))
+	{
+        assert(0);
+        return 0;
+    }
+    err = glGetError();
+    assert(err==GL_NO_ERROR);
+
+     sprintf(fullFontFileName,"%s%s",fontPath,"DroidSansJapanese.ttf");
+    if (!(droidJapanese = sth_add_font(stash,fullFontFileName)))
+	{
+        assert(0);
+        return 0;
+    }
+    err = glGetError();
+    assert(err==GL_NO_ERROR);
+
+	return stash;
+}
+
+
+
+
+#include "OpenGLWindow/OpenGLInclude.h"
+#include "BulletCommon/CommandLineArgs.h"
+
+void Usage()
+{
+	printf("\nprogram.exe [--cl_device=<int>] [--benchmark] [--disable_opencl] [--cl_platform=<int>]  [--x_dim=<int>] [--y_dim=<num>] [--z_dim=<int>] [--x_gap=<float>] [--y_gap=<float>] [--z_gap=<float>] [--use_concave_mesh]\n");
+};
+
+
+void	DumpSimulationTime(FILE* f)
+{
+	CProfileIterator* profileIterator = CProfileManager::Get_Iterator();
+
+	profileIterator->First();
+	if (profileIterator->Is_Done())
+		return;
+
+	float accumulated_time=0,parent_time = profileIterator->Is_Root() ? CProfileManager::Get_Time_Since_Reset() : profileIterator->Get_Current_Parent_Total_Time();
+	int i;
+	int frames_since_reset = CProfileManager::Get_Frame_Count_Since_Reset();
+
+	//fprintf(f,"%.3f,",	parent_time );
+	float totalTime = 0.f;
+
+
+
+	static bool headersOnce = true;
+
+	if (headersOnce)
+	{
+		headersOnce = false;
+		fprintf(f,"root,");
+
+			for (i = 0; !profileIterator->Is_Done(); i++,profileIterator->Next())
+		{
+			float current_total_time = profileIterator->Get_Current_Total_Time();
+			accumulated_time += current_total_time;
+			float fraction = parent_time > SIMD_EPSILON ? (current_total_time / parent_time) * 100 : 0.f;
+			const char* name = profileIterator->Get_Current_Name();
+			fprintf(f,"%s,",name);
+		}
+		fprintf(f,"\n");
+	}
+
+
+	fprintf(f,"%.3f,",parent_time);
+	profileIterator->First();
+	for (i = 0; !profileIterator->Is_Done(); i++,profileIterator->Next())
+	{
+		float current_total_time = profileIterator->Get_Current_Total_Time();
+		accumulated_time += current_total_time;
+		float fraction = parent_time > SIMD_EPSILON ? (current_total_time / parent_time) * 100 : 0.f;
+		const char* name = profileIterator->Get_Current_Name();
+		//if (!strcmp(name,"stepSimulation"))
+		{
+			fprintf(f,"%.3f,",current_total_time);
+		}
+		totalTime += current_total_time;
+		//recurse into children
+	}
+
+	fprintf(f,"\n");
+
+
+	CProfileManager::Release_Iterator(profileIterator);
+
+
+}
+extern const char* g_deviceName;
+
+int main(int argc, char* argv[])
+{
+    printf("main start");
+
+	CommandLineArgs args(argc,argv);
+	GpuDemo::ConstructionInfo ci;
+
+	if (args.CheckCmdLineFlag("help"))
+	{
+		Usage();
+		return 0;
+	}
+
+	bool benchmark=args.CheckCmdLineFlag("benchmark");
+	dump_timings=args.CheckCmdLineFlag("dump_timings");
+	ci.useOpenCL = !args.CheckCmdLineFlag("disable_opencl");
+	ci.m_useConcaveMesh = true;//args.CheckCmdLineFlag("use_concave_mesh");
+	if (ci.m_useConcaveMesh)
+	{
+		enableExperimentalCpuConcaveCollision = true;
+	}
+
+	args.GetCmdLineArgument("cl_device", ci.preferredOpenCLDeviceIndex);
+	args.GetCmdLineArgument("cl_platform", ci.preferredOpenCLPlatformIndex);
+	args.GetCmdLineArgument("x_dim", ci.arraySizeX);
+	args.GetCmdLineArgument("y_dim", ci.arraySizeY);
+	args.GetCmdLineArgument("z_dim", ci.arraySizeZ);
+	args.GetCmdLineArgument("x_gap", ci.gapX);
+	args.GetCmdLineArgument("y_gap", ci.gapY);
+	args.GetCmdLineArgument("z_gap", ci.gapZ);
+
+
+	printf("Demo settings:\n");
+	printf("x_dim=%d, y_dim=%d, z_dim=%d\n",ci.arraySizeX,ci.arraySizeY,ci.arraySizeZ);
+	printf("x_gap=%f, y_gap=%f, z_gap=%f\n",ci.gapX,ci.gapY,ci.gapZ);
+
+	printf("Preferred cl_device index %d\n", ci.preferredOpenCLDeviceIndex);
+	printf("Preferred cl_platform index%d\n", ci.preferredOpenCLPlatformIndex);
+	printf("-----------------------------------------------------\n");
+
+	#ifndef BT_NO_PROFILE
+	CProfileManager::Reset();
+#endif //BT_NO_PROFILE
+
+
+	window = new btgDefaultOpenGLWindow();
+
+	btgWindowConstructionInfo wci(g_OpenGLWidth,g_OpenGLHeight);
+
+	window->createWindow(wci);
+	window->setResizeCallback(MyResizeCallback);
+	window->setMouseMoveCallback(MyMouseMoveCallback);
+	window->setMouseButtonCallback(MyMouseButtonCallback);
+	window->setKeyboardCallback(MyKeyboardCallback);
+
+	window->setWindowTitle("MyTest");
+	printf("-----------------------------------------------------\n");
+
+
+#ifndef __APPLE__
+	glewInit();
+#endif
+
+	gui = new GwenUserInterface();
+
+    printf("started GwenUserInterface");
+
+
+	GLPrimitiveRenderer prim(g_OpenGLWidth,g_OpenGLHeight);
+
+	stash = initFont(&prim);
+
+
+	gui->init(g_OpenGLWidth,g_OpenGLHeight,stash,window->getRetinaScale());
+
+    printf("init fonts");
+
+
+	gui->setToggleButtonCallback(MyButtonCallback);
+
+	gui->registerToggleButton(MYPAUSE,"Pause");
+	gui->registerToggleButton(MYPROFILE,"Profile");
+	gui->registerToggleButton(MYRESET,"Reset");
+
+
+
+
+
+
+	int numItems = sizeof(allDemos)/sizeof(GpuDemo::CreateFunc*);
+	demoNames.clear();
+	for (int i=0;i<numItems;i++)
+	{
+		GpuDemo* demo = allDemos[i]();
+		demoNames.push_back(demo->getName());
+		delete demo;
+	}
+
+	gui->registerComboBox(MYCOMBOBOX1,numItems,&demoNames[0]);
+	gui->setComboBoxCallback(MyComboBoxCallback);
+
+
+
+	do
+	{
+		bool syncOnly = false;
+		gReset = false;
+
+
+
+
+	static bool once=true;
+
+
+
+
+	glClearColor(1,0,0,1);
+	glClear(GL_COLOR_BUFFER_BIT);
+
+	{
+		window->startRendering();
+		glFinish();
+
+
+
+
+		float color[4] = {1,1,1,1};
+		prim.drawRect(0,0,200,200,color);
+		float retinaScale = 1;
+
+		  float x = 10;
+            float y=220;
+            float  dx=0;
+            if (1)
+            {
+                BT_PROFILE("font sth_draw_text");
+
+				glEnable(GL_BLEND);
+				GLint err = glGetError();
+				assert(err==GL_NO_ERROR);
+
+				glBlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
+				err = glGetError();
+				assert(err==GL_NO_ERROR);
+
+				glDisable(GL_DEPTH_TEST);
+				err = glGetError();
+				assert(err==GL_NO_ERROR);
+
+
+				glDisable(GL_CULL_FACE);
+
+                sth_begin_draw(stash);
+                sth_flush_draw(stash);
+                sth_draw_text(stash, droidRegular,20.f, x, y, "Non-retina font rendering !@#$", &dx,g_OpenGLWidth,g_OpenGLHeight,0,1);//retinaScale);
+                if (retinaScale!=1.f)
+                    sth_draw_text(stash, droidRegular,20.f*retinaScale, x, y+20, "Retina font rendering!@#$", &dx,g_OpenGLWidth,g_OpenGLHeight,0,retinaScale);
+                sth_flush_draw(stash);
+
+                sth_end_draw(stash);
+            }
+
+		gui->draw(g_OpenGLWidth,g_OpenGLHeight);
+		window->endRendering();
+		glFinish();
+	}
+	once=false;
+
+	OpenGL3CoreRenderer render;
+
+	glClearColor(0,1,0,1);
+	glClear(GL_COLOR_BUFFER_BIT);
+
+	window->endRendering();
+
+	glFinish();
+
+
+
+	window->setWheelCallback(btDefaultWheelCallback);
+
+
+
+
+	{
+		GpuDemo* demo = allDemos[selectedDemo]();
+//		demo->myinit();
+		bool useGpu = false;
+
+
+		ci.m_instancingRenderer = render.getInstancingRenderer();
+		render.init();
+
+		demo->initPhysics(ci);
+		printf("-----------------------------------------------------\n");
+
+		FILE* f = 0;
+		if (benchmark)
+		{
+			gPause = false;
+			char fileName[1024];
+
+#ifdef _WIN32
+			SYSTEMTIME time;
+			GetLocalTime(&time);
+			char buf[1024];
+			DWORD dwCompNameLen = 1024;
+			if (0 != GetComputerName(buf, &dwCompNameLen))
+			{
+				printf("%s", buf);
+			} else
+			{
+				printf("unknown", buf);
+			}
+			sprintf(fileName,"%s_%s_%s_%d_%d_%d_date_%d-%d-%d_time_%d-%d-%d.csv",g_deviceName,buf,demoNames[selectedDemo],ci.arraySizeX,ci.arraySizeY,ci.arraySizeZ,time.wDay,time.wMonth,time.wYear,time.wHour,time.wMinute,time.wSecond);
+
+			printf("Open file %s\n", fileName);
+#else
+			sprintf(fileName,"%s_%d_%d_%d.csv",g_deviceName,ci.arraySizeX,ci.arraySizeY,ci.arraySizeZ);
+			printf("Open file %s\n", fileName);
+#endif
+
+
+			//GetSystemTime(&time2);
+
+			f=fopen(fileName,"w");
+			//if (f)
+			//	fprintf(f,"%s (%dx%dx%d=%d),\n",  g_deviceName,ci.arraySizeX,ci.arraySizeY,ci.arraySizeZ,ci.arraySizeX*ci.arraySizeY*ci.arraySizeZ);
+		}
+
+		printf("-----------------------------------------------------\n");
+		do
+		{
+			CProfileManager::Reset();
+			CProfileManager::Increment_Frame_Counter();
+
+			render.reshape(g_OpenGLWidth,g_OpenGLHeight);
+
+			window->startRendering();
+
+			char msg[1024];
+			int numInstances = 0;//ci.m_instancingRenderer->getNumInstances();
+			sprintf(msg,"Num objects = %d",numInstances);
+			gui->setStatusBarMessage(msg,true);
+
+			glClearColor(0.6,0.6,0.6,1);
+			glClear(GL_COLOR_BUFFER_BIT| GL_DEPTH_BUFFER_BIT|GL_STENCIL_BUFFER_BIT);
+			glEnable(GL_DEPTH_TEST);
+
+
+			if (!gPause)
+			{
+				BT_PROFILE("simulate");
+
+				demo->clientMoveAndDisplay();
+			}
+			else
+			{
+
+			}
+
+			{
+				BT_PROFILE("renderScene");
+				demo->renderScene();
+			}
+
+
+			/*if (demo->getDynamicsWorld() && demo->getDynamicsWorld()->getNumCollisionObjects())
+			{
+				BT_PROFILE("renderPhysicsWorld");
+				btAlignedObjectArray<btCollisionObject*> arr = demo->getDynamicsWorld()->getCollisionObjectArray();
+				btCollisionObject** colObjArray = &arr[0];
+
+				render.renderPhysicsWorld(demo->getDynamicsWorld()->getNumCollisionObjects(),colObjArray, syncOnly);
+				syncOnly = true;
+
+			}
+			*/
+			{
+				BT_PROFILE("gui->draw");
+				gui->draw(g_OpenGLWidth,g_OpenGLHeight);
+			}
+			{
+				BT_PROFILE("window->endRendering");
+				window->endRendering();
+			}
+			{
+				BT_PROFILE("glFinish");
+			}
+
+
+		if (dump_timings)
+			CProfileManager::dumpAll();
+
+		if (f)
+		{
+			static int count=0;
+
+			if (count>2 && count<102)
+			{
+				DumpSimulationTime(f);
+			}
+			if (count>=102)
+				window->setRequestExit();
+			count++;
+		}
+
+
+
+		} while (!window->requestedExit() && !gReset);
+
+
+		demo->exitPhysics();
+		delete demo;
+		if (f)
+			fclose(f);
+	}
+
+
+
+	} while (gReset);
+
+
+	gui->setComboBoxCallback(0);
+	delete gui;
+	gui=0;
+
+	window->closeWindow();
+	delete window;
+	window = 0;
+
+	return 0;
+}

demo/gpudemo/premake4.lua

@@ -0,0 +1,60 @@
+function createProject(vendor)
+	
+	hasCL = findOpenCL(vendor)
+	
+	if (hasCL) then
+
+		project ("Bullet3_OpenCL_gpu_demo_" .. vendor)
+
+		initOpenCL(vendor)
+		
+		language "C++"
+				
+		kind "ConsoleApp"
+		targetdir "../../bin"
+
+
+		initOpenGL()
+		initGlew()
+
+		includedirs {
+		 	"..",
+		 	"../../src",
+		 	"../../btgui",
+		 	"../../opencl"
+		}
+		
+		links {
+			"gwen"
+		}
+		
+		files {
+			"main_opengl3core.cpp",
+			"gwenUserInterface.cpp",
+			"gwenUserInterface.h",
+			"ParticleDemo.cpp",
+			"ParticleDemo.h",
+
+		}
+
+		if os.is("Windows") then 
+			files{  
+				"../../btgui/OpenGLWindow/Win32OpenGLWindow.cpp",
+      	"../../btgui/OpenGLWindow/Win32OpenGLWindow.h",
+      	"../../btgui/OpenGLWindow/Win32Window.cpp",
+      	"../../btgui/OpenGLWindow/Win32Window.h",
+			}
+		end
+		if os.is("Linux") then
+			files {
+				"../../btgui/OpenGLWindow/X11OpenGLWindow.cpp",
+				"../../btgui/OpenGLWindow/X11OpenGLWindows.h"
+			}
+		end
+	end
+end
+
+createProject("Apple")
+createProject("AMD")
+createProject("Intel")
+createProject("NVIDIA")