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

removed Particles OpenCL demo, all OpenCL work is moved to Bullet 3.x see

Browse Source 
http://github.com/erwincoumans/bullet3
Fixes Issue 741
This commit is contained in:
erwin.coumans
2013-10-01 17:28:56 +00:00
parent 858c890c96
commit dece911679
19 changed files with 30 additions and 3294 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
Demos/CMakeLists.txt
View File

@@ -26,7 +26,7 @@ IF (USE_GLUT)
ENDIF(BUILD_CPU_DEMOS)
IF(BUILD_MULTITHREADING)
SUBDIRS( MultiThreadedDemo ParticlesOpenCL OpenCLClothDemo )
SUBDIRS( MultiThreadedDemo OpenCLClothDemo )
ENDIF(BUILD_MULTITHREADING)
SUBDIRS(

30
Demos/MovingConcaveDemo/ConcavePhysicsDemo.cpp
View File

@@ -24,7 +24,7 @@ subject to the following restrictions:
#include "BulletCollision/Gimpact/btGImpactShape.h"
#include "BulletCollision/Gimpact/btGImpactCollisionAlgorithm.h"
#include "GLDebugFont.h"
#include "BulletCollision/Gimpact/btCompoundFromGimpact.h"
#include "GLDebugDrawer.h"
@@ -1632,10 +1632,19 @@ void ConcaveDemo::initGImpactCollision()
3*sizeof(int),
NUM_VERTICES,(REAL*) &gVertices[0],sizeof(REAL)*3);
{
btGImpactMeshShape * trimesh = new btGImpactMeshShape(indexVertexArrays);
trimesh->setLocalScaling(btVector3(4.f,4.f,4.f));
trimesh->updateBound();
#define USE_COMPOUND
#ifdef USE_COMPOUND
m_trimeshShape = btCreateCompoundFromGimpactShape(trimesh,1);
delete trimesh;
trimesh=0;
#else
m_trimeshShape = trimesh;
#endif
}
//register algorithm
@@ -1656,7 +1665,7 @@ void ConcaveDemo::initPhysics()
btConstraintSolver* constraintSolver = new btSequentialImpulseConstraintSolver();
m_dynamicsWorld = new btDiscreteDynamicsWorld(dispatcher,broadphase,constraintSolver,collisionConfiguration);
m_dynamicsWorld ->setGravity(btVector3(0,0,0));
//create trimesh model and shape
initGImpactCollision();
@@ -1695,7 +1704,7 @@ void ConcaveDemo::initPhysics()
//enable custom material callback
staticBody->setCollisionFlags(staticBody->getCollisionFlags()|btCollisionObject::CF_CUSTOM_MATERIAL_CALLBACK);
#if 0
//static plane
btVector3 normal(0.4,1.5,-0.4);
normal.normalize();
@@ -1715,10 +1724,10 @@ void ConcaveDemo::initPhysics()
localCreateRigidBody(1, startTransform,boxShape);
}
}
#endif
shootTrimesh(btVector3(0,10,0),btVector3(0,10,0));
shootTrimesh(btVector3(0,10,0),btVector3(0,0,0));
shootTrimesh(btVector3(0,20,0),btVector3(0,10,0));
shootTrimesh(btVector3(0,10,0),btVector3(0,10,0));
//m_debugMode |= btIDebugDraw::DBG_DrawWireframe;
@@ -1878,8 +1887,11 @@ void ConcaveDemo::shootTrimesh(const btVector3& startPosition,const btVector3& d
btRigidBody* body = this->localCreateRigidBody(mass, startTransform,m_trimeshShape);
btVector3 linVel(destination[0]-startPosition[0],destination[1]-startPosition[1],destination[2]-startPosition[2]);
if (linVel.length2()>SIMD_EPSILON)
{
linVel.normalize();
linVel*=m_ShootBoxInitialSpeed*0.25;
}
body->getWorldTransform().setOrigin(startPosition);
body->getWorldTransform().setRotation(btQuaternion(0,0,0,1));
@@ -1894,7 +1906,11 @@ void ConcaveDemo::clientMoveAndDisplay()
float dt = float(getDeltaTimeMicroseconds()) * 0.000001f;
m_dynamicsWorld->stepSimulation(dt);
extern int MyTTcound;
m_dynamicsWorld->stepSimulation(1./60.,0);//dt,0,1./60.);
CProfileManager::dumpAll();
printf("MyTTcound=%d\n",MyTTcound);
MyTTcound=0;
//optional but useful: debug drawing
m_dynamicsWorld->debugDrawWorld();

86
Demos/ParticlesOpenCL/AMD/CMakeLists.txt
View File

@@ -1,86 +0,0 @@
INCLUDE_DIRECTORIES(
${BULLET_PHYSICS_SOURCE_DIR}/src
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL
${BULLET_PHYSICS_SOURCE_DIR}/Demos/OpenGL
${AMD_OPENCL_INCLUDES}
)
ADD_DEFINITIONS(-DUSE_AMD_OPENCL)
ADD_DEFINITIONS(-DCL_PLATFORM_AMD)
IF(WIN32)
ADD_DEFINITIONS(-DGLEW_STATIC)
IF (CMAKE_CL_64)
SET(CMAK_GLEW_LIBRARY
${BULLET_PHYSICS_SOURCE_DIR}/Glut/glew64s.lib )
ELSE(CMAKE_CL_64)
SET(CMAK_GLEW_LIBRARY ${BULLET_PHYSICS_SOURCE_DIR}/Glut/glew32s.lib )
ENDIF(CMAKE_CL_64)
else()
SET(CMAK_GLEW_LIBRARY GLEW)
endIF(WIN32)
IF (USE_GLUT)
LINK_LIBRARIES(
OpenGLSupport
BulletDynamics
BulletCollision
BulletMultiThreaded
LinearMath
${GLUT_glut_LIBRARY}
${OPENGL_gl_LIBRARY}
${OPENGL_glu_LIBRARY}
${CMAK_GLEW_LIBRARY}
${CMAKE_ATISTREAMSDK_LIBRARY}
)
ADD_EXECUTABLE(AppParticlesOCL_AMD
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesDynamicsWorld.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesSharedDefs.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesSharedTypes.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesDemo.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/shaders.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLUtils.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLUtils.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLInclude.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesDemoDynamicsWorld.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/main.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesDemo.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/shaders.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl
)
ELSE (USE_GLUT)
ENDIF (USE_GLUT)
IF(WIN32)
IF (CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_AMD POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/glut64.dll ${CMAKE_CURRENT_BINARY_DIR} )
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_AMD POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl ${CMAKE_CURRENT_BINARY_DIR})
ENDIF()
ELSE(CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_AMD POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/GLUT32.DLL ${CMAKE_CURRENT_BINARY_DIR})
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_AMD POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl ${CMAKE_CURRENT_BINARY_DIR})
ENDIF()
ENDIF(CMAKE_CL_64)
ENDIF(WIN32)
IF (UNIX)
TARGET_LINK_LIBRARIES(AppParticlesOCL_AMD pthread)
ENDIF(UNIX)
IF (INTERNAL_ADD_POSTFIX_EXECUTABLE_NAMES)
SET_TARGET_PROPERTIES(AppParticlesOCL_AMD PROPERTIES DEBUG_POSTFIX "_Debug")
SET_TARGET_PROPERTIES(AppParticlesOCL_AMD PROPERTIES MINSIZEREL_POSTFIX "_MinsizeRel")
SET_TARGET_PROPERTIES(AppParticlesOCL_AMD PROPERTIES RELWITHDEBINFO_POSTFIX "_RelWithDebugInfo")
ENDIF(INTERNAL_ADD_POSTFIX_EXECUTABLE_NAMES)

84
Demos/ParticlesOpenCL/Apple/CMakeLists.txt
View File

@@ -1,84 +0,0 @@
INCLUDE_DIRECTORIES(
${BULLET_PHYSICS_SOURCE_DIR}/src
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL
${BULLET_PHYSICS_SOURCE_DIR}/Demos/OpenGL
)
IF (APPLE)
FIND_LIBRARY(OPENCL_LIBRARY OpenCL DOC "OpenCL lib for OSX")
FIND_PATH(OPENCL_INCLUDE_DIR OpenCL/cl.h DOC "Include for OpenCL on OSX")
ENDIF (APPLE)
#ADD_DEFINITIONS(-DUSE_MINICL)
IF(WIN32)
IF (CMAKE_CL_64)
SET(CMAK_GLEW_LIBRARY ${BULLET_PHYSICS_SOURCE_DIR}/Glut/glew64.lib )
ELSE(CMAKE_CL_64)
SET(CMAK_GLEW_LIBRARY ${BULLET_PHYSICS_SOURCE_DIR}/Glut/glew32.lib )
ENDIF(CMAKE_CL_64)
ENDIF(WIN32)
IF (USE_GLUT)
LINK_LIBRARIES(
OpenGLSupport
BulletDynamics
BulletCollision
LinearMath
${OPENCL_LIBRARY}
${GLUT_glut_LIBRARY}
${OPENGL_gl_LIBRARY}
${OPENGL_glu_LIBRARY}
${CMAK_GLEW_LIBRARY}
)
ADD_EXECUTABLE(AppParticlesOCL_Apple
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesDynamicsWorld.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesSharedDefs.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesSharedTypes.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesDemo.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/shaders.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLUtils.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLUtils.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLInclude.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesDemoDynamicsWorld.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/main.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesDemo.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/shaders.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl
)
ELSE (USE_GLUT)
ENDIF (USE_GLUT)
IF(WIN32)
IF (CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Apple POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/glut64.dll ${CMAKE_CURRENT_BINARY_DIR} )
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Apple POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/GLEW64.DLL ${CMAKE_CURRENT_BINARY_DIR})
ENDIF()
ELSE(CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Apple POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/GLUT32.DLL ${CMAKE_CURRENT_BINARY_DIR})
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Apple POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/GLEW32.DLL ${CMAKE_CURRENT_BINARY_DIR})
ENDIF()
ENDIF(CMAKE_CL_64)
ENDIF(WIN32)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Apple POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl ${CMAKE_CURRENT_BINARY_DIR})
IF (UNIX)
TARGET_LINK_LIBRARIES(AppParticlesOCL_Apple pthread)
ENDIF(UNIX)

19
Demos/ParticlesOpenCL/CMakeLists.txt
View File

@@ -1,19 +0,0 @@
IF(BUILD_MINICL_OPENCL_DEMOS)
SUBDIRS( MiniCL )
ENDIF()
IF(BUILD_INTEL_OPENCL_DEMOS)
SUBDIRS(Intel)
ENDIF()
IF(BUILD_AMD_OPENCL_DEMOS)
SUBDIRS(AMD)
ENDIF()
IF(BUILD_NVIDIA_OPENCL_DEMOS)
SUBDIRS(NVidia)
ENDIF()
IF(APPLE)
SUBDIRS(Apple)
ENDIF()

86
Demos/ParticlesOpenCL/Intel/CMakeLists.txt
View File

@@ -1,86 +0,0 @@
INCLUDE_DIRECTORIES(
${BULLET_PHYSICS_SOURCE_DIR}/src
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL
${BULLET_PHYSICS_SOURCE_DIR}/Demos/OpenGL
${INTEL_OPENCL_INCLUDES}
)
ADD_DEFINITIONS(-DUSE_INTEL_OPENCL)
ADD_DEFINITIONS(-DCL_PLATFORM_INTEL)
IF(WIN32)
ADD_DEFINITIONS(-DGLEW_STATIC)
ENDIF(WIN32)
IF (CMAKE_CL_64)
SET(CMAK_GLEW_LIBRARY
${BULLET_PHYSICS_SOURCE_DIR}/Glut/glew64s.lib )
ELSE(CMAKE_CL_64)
SET(CMAK_GLEW_LIBRARY ${BULLET_PHYSICS_SOURCE_DIR}/Glut/glew32s.lib )
ENDIF(CMAKE_CL_64)
IF (USE_GLUT)
LINK_LIBRARIES(
OpenGLSupport
BulletDynamics
BulletCollision
BulletMultiThreaded
LinearMath
${GLUT_glut_LIBRARY}
${OPENGL_gl_LIBRARY}
${OPENGL_glu_LIBRARY}
${CMAK_GLEW_LIBRARY}
${INTEL_OPENCL_LIBRARIES}
)
ADD_EXECUTABLE(AppParticlesOCL_Intel
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesDynamicsWorld.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesSharedDefs.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesSharedTypes.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesDemo.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/shaders.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLUtils.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLUtils.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLInclude.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesDemoDynamicsWorld.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/main.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesDemo.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/shaders.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl
)
ELSE (USE_GLUT)
ENDIF (USE_GLUT)
IF(WIN32)
IF (CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Intel POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/glut64.dll ${CMAKE_CURRENT_BINARY_DIR} )
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Intel POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl ${CMAKE_CURRENT_BINARY_DIR})
ENDIF()
ELSE(CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Intel POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/GLUT32.DLL ${CMAKE_CURRENT_BINARY_DIR})
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Intel POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl ${CMAKE_CURRENT_BINARY_DIR})
ENDIF()
ENDIF(CMAKE_CL_64)
ENDIF(WIN32)
IF (UNIX)
TARGET_LINK_LIBRARIES(AppParticlesOCL_Intel pthread)
ENDIF(UNIX)
IF (INTERNAL_ADD_POSTFIX_EXECUTABLE_NAMES)
SET_TARGET_PROPERTIES(AppParticlesOCL_Intel PROPERTIES DEBUG_POSTFIX "_Debug")
SET_TARGET_PROPERTIES(AppParticlesOCL_Intel PROPERTIES MINSIZEREL_POSTFIX "_MinsizeRel")
SET_TARGET_PROPERTIES(AppParticlesOCL_Intel PROPERTIES RELWITHDEBINFO_POSTFIX "_RelWithDebugInfo")
ENDIF(INTERNAL_ADD_POSTFIX_EXECUTABLE_NAMES)

94
Demos/ParticlesOpenCL/MiniCL/CMakeLists.txt
View File

@@ -1,94 +0,0 @@
INCLUDE_DIRECTORIES(
${BULLET_PHYSICS_SOURCE_DIR}/src
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL
${BULLET_PHYSICS_SOURCE_DIR}/Demos/OpenGL
)
ADD_DEFINITIONS(-DUSE_MINICL)
#to get this to compile under linux, you could try setting CMAK_GLEW_LIBRARY
#to /usr/lib/libGLEW.so
IF(WIN32)
ADD_DEFINITIONS(-DGLEW_STATIC)
ENDIF(WIN32)
IF(WIN32)
IF (CMAKE_CL_64)
SET(CMAK_GLEW_LIBRARY ${BULLET_PHYSICS_SOURCE_DIR}/Glut/glew64s.lib )
ELSE(CMAKE_CL_64)
SET(CMAK_GLEW_LIBRARY ${BULLET_PHYSICS_SOURCE_DIR}/Glut/glew32s.lib )
ENDIF(CMAKE_CL_64)
ENDIF(WIN32)
IF(UNIX AND NOT APPLE)
SET(CMAK_GLEW_LIBRARY /usr/lib/libGLEW.so )
ENDIF(UNIX AND NOT APPLE)
IF (USE_GLUT)
LINK_LIBRARIES(
OpenGLSupport
BulletDynamics
BulletCollision
MiniCL
BulletMultiThreaded
LinearMath
${GLUT_glut_LIBRARY}
${OPENGL_gl_LIBRARY}
${OPENGL_glu_LIBRARY}
${CMAK_GLEW_LIBRARY}
)
ADD_EXECUTABLE(AppParticlesOCL_Mini
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesDynamicsWorld.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesSharedDefs.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesSharedTypes.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesDemo.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesDemoDynamicsWorld.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/main.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesDemo.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/shaders.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/shaders.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/MiniCL/MiniCLTaskWrap.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLUtils.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLUtils.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLInclude.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl
)
ELSE (USE_GLUT)
ENDIF (USE_GLUT)
IF(WIN32)
IF (CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Mini POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/glut64.dll ${CMAKE_CURRENT_BINARY_DIR} )
ENDIF()
ELSE(CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Mini POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/GLUT32.DLL ${CMAKE_CURRENT_BINARY_DIR})
ENDIF()
ENDIF(CMAKE_CL_64)
ENDIF(WIN32)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Mini POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl ${CMAKE_CURRENT_BINARY_DIR})
IF (UNIX)
TARGET_LINK_LIBRARIES(AppParticlesOCL_Mini pthread)
ENDIF(UNIX)
IF (INTERNAL_ADD_POSTFIX_EXECUTABLE_NAMES)
SET_TARGET_PROPERTIES(AppParticlesOCL_Mini PROPERTIES DEBUG_POSTFIX "_Debug")
SET_TARGET_PROPERTIES(AppParticlesOCL_Mini PROPERTIES MINSIZEREL_POSTFIX "_MinsizeRel")
SET_TARGET_PROPERTIES(AppParticlesOCL_Mini PROPERTIES RELWITHDEBINFO_POSTFIX "_RelWithDebugInfo")
ENDIF(INTERNAL_ADD_POSTFIX_EXECUTABLE_NAMES)

33
Demos/ParticlesOpenCL/MiniCL/MiniCLTaskWrap.cpp
View File

@@ -1,33 +0,0 @@
/*
Bullet Continuous Collision Detection and Physics Library
Copyright (c) 2003-2007 Erwin Coumans http://bulletphysics.com
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 <MiniCL/cl_MiniCL_Defs.h>
#define MSTRINGIFY(A) A
#define LOCAL_SIZE_MAX 1024U
#include "../ParticlesOCL.cl"
MINICL_REGISTER(kComputeCellId)
MINICL_REGISTER(kClearCellStart)
MINICL_REGISTER(kFindCellStart)
MINICL_REGISTER(kIntegrateMotion)
MINICL_REGISTER(kCollideParticles)
MINICL_REGISTER(kBitonicSortCellIdLocal)
MINICL_REGISTER(kBitonicSortCellIdLocal1)
MINICL_REGISTER(kBitonicSortCellIdMergeGlobal)
MINICL_REGISTER(kBitonicSortCellIdMergeLocal)

95
Demos/ParticlesOpenCL/NVidia/CMakeLists.txt
View File

@@ -1,95 +0,0 @@
ADD_DEFINITIONS(-DUSE_NVIDIA_OPENCL)
ADD_DEFINITIONS(-DCL_PLATFORM_NVIDIA)
INCLUDE_DIRECTORIES(
${BULLET_PHYSICS_SOURCE_DIR}/src
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL
${BULLET_PHYSICS_SOURCE_DIR}/Demos/OpenGL
${NVIDIA_OPENCL_INCLUDES}
)
IF(WIN32)
ADD_DEFINITIONS(-DGLEW_STATIC)
ENDIF(WIN32)
IF(WIN32)
IF (CMAKE_CL_64)
SET(CMAK_GLEW_LIBRARY ${BULLET_PHYSICS_SOURCE_DIR}/Glut/glew64s.lib )
ELSE(CMAKE_CL_64)
SET(CMAK_GLEW_LIBRARY ${BULLET_PHYSICS_SOURCE_DIR}/Glut/glew32s.lib )
ENDIF(CMAKE_CL_64)
ELSE()
IF(UNIX AND NOT APPLE)
FIND_LIBRARY(CMAK_GLEW_LIBRARY GLEW PATH /usr/lib /usr/local/lib /usr/lib64 /usr/local/lib64 )
ENDIF()
ENDIF(WIN32)
IF(CMAK_GLEW_LIBRARY)
IF (USE_GLUT)
LINK_LIBRARIES(
OpenGLSupport
BulletDynamics
BulletCollision
LinearMath
${GLUT_glut_LIBRARY}
${OPENGL_gl_LIBRARY}
${OPENGL_glu_LIBRARY}
${CMAK_GLEW_LIBRARY}
${NVIDIA_OPENCL_LIBRARIES}
)
ADD_EXECUTABLE(AppParticlesOCL_Nv
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesDynamicsWorld.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesSharedDefs.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesSharedTypes.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesDemo.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/shaders.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLUtils.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLUtils.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOpenCLInclude.h
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/btParticlesDemoDynamicsWorld.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/main.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesDemo.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/shaders.cpp
${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl
)
ELSE (USE_GLUT)
ENDIF (USE_GLUT)
IF(WIN32)
IF (CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Nv POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/glut64.dll ${CMAKE_CURRENT_BINARY_DIR} )
ENDIF()
ELSE(CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Nv POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/GLUT32.DLL ${CMAKE_CURRENT_BINARY_DIR})
ENDIF()
ENDIF(CMAKE_CL_64)
IF (NOT INTERNAL_CREATE_DISTRIBUTABLE_MSVC_PROJECTFILES)
ADD_CUSTOM_COMMAND( TARGET AppParticlesOCL_Nv POST_BUILD
COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${BULLET_PHYSICS_SOURCE_DIR}/Demos/ParticlesOpenCL/ParticlesOCL.cl ${CMAKE_CURRENT_BINARY_DIR})
ENDIF()
ENDIF(WIN32)
IF (UNIX)
TARGET_LINK_LIBRARIES(AppParticlesOCL_Nv pthread)
ENDIF(UNIX)
ENDIF(CMAK_GLEW_LIBRARY)
IF (INTERNAL_ADD_POSTFIX_EXECUTABLE_NAMES)
SET_TARGET_PROPERTIES(AppParticlesOCL_Nv PROPERTIES DEBUG_POSTFIX "_Debug")
SET_TARGET_PROPERTIES(AppParticlesOCL_Nv PROPERTIES MINSIZEREL_POSTFIX "_MinsizeRel")
SET_TARGET_PROPERTIES(AppParticlesOCL_Nv PROPERTIES RELWITHDEBINFO_POSTFIX "_RelWithDebugInfo")
ENDIF(INTERNAL_ADD_POSTFIX_EXECUTABLE_NAMES)

651
Demos/ParticlesOpenCL/ParticlesDemo.cpp
View File

@@ -1,651 +0,0 @@
/*
Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org
Copyright (C) 2006 - 2009 Sony Computer Entertainment Inc.
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.
*/
#define START_POS_X btScalar(0.f)
#define START_POS_Y btScalar(0.f)
#define START_POS_Z btScalar(0.f)
//#define START_POS_Y btScalar(40.f)
//#define START_POS_Z btScalar(40.f)
//#define START_POS_Y btScalar(0.4f)
//#define START_POS_Z btScalar(0.4f)
#define ARRAY_SIZE_X 32
#define ARRAY_SIZE_Y 32
//#define ARRAY_SIZE_Y 16
#define ARRAY_SIZE_Z 32
//16
//#define ARRAY_SIZE_Z 1
//#define DIST btScalar(2.f)
#define DIST (DEF_PARTICLE_RADIUS * 2.f)
#define STRESS_X 20
//#define STRESS_Y 200
#define STRESS_Y 640
///The 3 following lines include the CPU implementation of the kernels, keep them in this order.
#include "BulletMultiThreaded/btGpuDefines.h"
#include "BulletMultiThreaded/btGpuUtilsSharedDefs.h"
#include "BulletMultiThreaded/btGpuUtilsSharedCode.h"
#ifndef __APPLE__
#include <GL/glew.h>
#endif
#include "GL_DialogDynamicsWorld.h"
#include "GL_DialogWindow.h"
#include "BulletCollision/CollisionDispatch/btEmptyCollisionAlgorithm.h"
#include "BulletCollision/CollisionDispatch/btSimulationIslandManager.h"
#include "GLDebugFont.h"
#include "GlutStuff.h"
///btBulletDynamicsCommon.h is the main Bullet include file, contains most common include files.
#include "btBulletDynamicsCommon.h"
#include <stdio.h> //printf debugging
#include "shaders.h"
#include "ParticlesDemo.h"
btScalar gTimeStep = 0.5f;//btScalar(1./60.);
#define SCALING btScalar(1.f)
void ParticlesDemo::clientMoveAndDisplay()
{
updateCamera();
glDisable(GL_LIGHTING);
glColor3f(1.f, 1.f, 1.f);
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
glDisable(GL_TEXTURE_2D); // we always draw wireframe in this demo
//simple dynamics world doesn't handle fixed-time-stepping
float ms = getDeltaTimeMicroseconds();
renderme();
if (m_dialogDynamicsWorld)
m_dialogDynamicsWorld->draw(gTimeStep);
///step the simulation
if (m_dynamicsWorld)
{
m_dynamicsWorld->stepSimulation(gTimeStep,0);//ms / 1000000.f);
//optional but useful: debug drawing
m_dynamicsWorld->debugDrawWorld();
}
ms = getDeltaTimeMicroseconds();
glFlush();
glutSwapBuffers();
}
void ParticlesDemo::displayCallback(void) {
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
renderme();
//optional but useful: debug drawing to detect problems
if (m_dynamicsWorld)
m_dynamicsWorld->debugDrawWorld();
//if (m_dialogDynamicsWorld)
// m_dialogDynamicsWorld->draw(gTimeStep);
glFlush();
glutSwapBuffers();
}
class btNullBroadphase : public btBroadphaseInterface
{
public:
btNullBroadphase()
{
}
virtual ~btNullBroadphase()
{
}
virtual btBroadphaseProxy* createProxy( const btVector3& aabbMin, const btVector3& aabbMax,int shapeType,void* userPtr, short int collisionFilterGroup,short int collisionFilterMask, btDispatcher* dispatcher,void* multiSapProxy)
{
return NULL;
}
virtual void destroyProxy(btBroadphaseProxy* proxy,btDispatcher* dispatcher)
{
}
virtual void setAabb(btBroadphaseProxy* proxy,const btVector3& aabbMin,const btVector3& aabbMax, btDispatcher* dispatcher)
{
}
virtual void getAabb(btBroadphaseProxy* proxy,btVector3& aabbMin, btVector3& aabbMax ) const
{
}
virtual void rayTest(const btVector3& rayFrom,const btVector3& rayTo, btBroadphaseRayCallback& rayCallback, const btVector3& aabbMin=btVector3(0,0,0), const btVector3& aabbMax = btVector3(0,0,0))
{
}
virtual void calculateOverlappingPairs(btDispatcher* dispatcher)
{
}
virtual btOverlappingPairCache* getOverlappingPairCache()
{
return NULL;
}
virtual const btOverlappingPairCache* getOverlappingPairCache() const
{
return NULL;
}
virtual void getBroadphaseAabb(btVector3& aabbMin,btVector3& aabbMax) const
{
}
virtual void resetPool(btDispatcher* dispatcher)
{
}
virtual void printStats()
{
}
virtual void aabbTest(const btVector3& aabbMin, const btVector3& aabbMax, btBroadphaseAabbCallback& callback)
{
}
};
void ParticlesDemo::initPhysics()
{
setTexturing(false);
setShadows(false);
// setCameraDistance(80.f);
setCameraDistance(3.0f);
// m_cameraTargetPosition.setValue(50, 10, 0);
m_cameraTargetPosition.setValue(0, 0, 0);
// m_azi = btScalar(0.f);
// m_ele = btScalar(0.f);
m_azi = btScalar(45.f);
m_ele = btScalar(30.f);
setFrustumZPlanes(0.1f, 10.f);
///collision configuration contains default setup for memory, collision setup
btDefaultCollisionConstructionInfo dci;
dci.m_defaultMaxPersistentManifoldPoolSize=50000;
dci.m_defaultMaxCollisionAlgorithmPoolSize=50000;
m_collisionConfiguration = new btDefaultCollisionConfiguration(dci);
///use the default collision dispatcher. For parallel processing you can use a diffent dispatcher (see Extras/BulletMultiThreaded)
m_dispatcher = new btCollisionDispatcher(m_collisionConfiguration);
m_pairCache = new (btAlignedAlloc(sizeof(btHashedOverlappingPairCache),16))btHashedOverlappingPairCache();
// m_broadphase = new btDbvtBroadphase(m_pairCache);
m_broadphase = new btNullBroadphase();
///the default constraint solver
m_solver = new btSequentialImpulseConstraintSolver();
m_pWorld = new btParticlesDynamicsWorld(m_dispatcher,m_broadphase,m_solver,m_collisionConfiguration, 65536);
m_dialogDynamicsWorld = new GL_DialogDynamicsWorld();
GL_DialogWindow* settings = m_dialogDynamicsWorld->createDialog(50,0,280,280,"CPU fallback");
m_pWorld->m_useCpuControls[0] = 0;
GL_ToggleControl* ctrl = 0;
m_pWorld->m_useCpuControls[SIMSTAGE_INTEGRATE_MOTION] = m_dialogDynamicsWorld->createToggle(settings,"Integrate Motion");
m_pWorld->m_useCpuControls[SIMSTAGE_COMPUTE_CELL_ID] = m_dialogDynamicsWorld->createToggle(settings,"Compute Cell ID");
m_pWorld->m_useCpuControls[SIMSTAGE_SORT_CELL_ID] = m_dialogDynamicsWorld->createToggle(settings,"Sort Cell ID");
m_pWorld->m_useCpuControls[SIMSTAGE_FIND_CELL_START] = m_dialogDynamicsWorld->createToggle(settings,"Find Cell Start");
m_pWorld->m_useCpuControls[SIMSTAGE_COLLIDE_PARTICLES] = m_dialogDynamicsWorld->createToggle(settings,"Collide Particles");
for(int i = 1; i < SIMSTAGE_TOTAL; i++)
{
m_pWorld->m_useCpuControls[i]->m_active = false;
}
#if defined(CL_PLATFORM_MINI_CL)
// these kernels use barrier()
m_pWorld->m_useCpuControls[SIMSTAGE_SORT_CELL_ID]->m_active = true;
m_pWorld->m_useCpuControls[SIMSTAGE_FIND_CELL_START]->m_active = true;
#endif
#if defined(CL_PLATFORM_AMD)
// these kernels use barrier()
m_pWorld->m_useCpuControls[SIMSTAGE_SORT_CELL_ID]->m_active = true;
m_pWorld->m_useCpuControls[SIMSTAGE_FIND_CELL_START]->m_active = true;
#endif
m_dynamicsWorld = m_pWorld;
m_pWorld->getSimulationIslandManager()->setSplitIslands(true);
m_pWorld->setGravity(btVector3(0,-10.,0));
m_pWorld->getSolverInfo().m_numIterations = 4;
{
// btCollisionShape* colShape = new btSphereShape(btScalar(1.0f));
/*
btCollisionShape* colShape = new btSphereShape(DEF_PARTICLE_RADIUS);
m_collisionShapes.push_back(colShape);
btTransform startTransform;
startTransform.setIdentity();
btScalar mass(1.f);
btVector3 localInertia(0,0,0);
colShape->calculateLocalInertia(mass,localInertia);
float start_x = START_POS_X - ARRAY_SIZE_X * DIST * btScalar(0.5f);
float start_y = START_POS_Y - ARRAY_SIZE_Y * DIST * btScalar(0.5f);
float start_z = START_POS_Z - ARRAY_SIZE_Z * DIST * btScalar(0.5f);
startTransform.setOrigin(btVector3(start_x, start_y, start_z));
btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,0,colShape,localInertia);
rbInfo.m_startWorldTransform = startTransform;
btRigidBody* body = new btRigidBody(rbInfo);
m_pWorld->addRigidBody(body);
*/
init_scene_directly();
}
clientResetScene();
m_pWorld->initDeviceData();
}
inline float frand(void){
return (float)rand() / (float)RAND_MAX;
}
void ParticlesDemo::init_scene_directly()
{
srand(1969);
float start_x = -1+DEF_PARTICLE_RADIUS;//START_POS_X - ARRAY_SIZE_X * DIST * btScalar(0.5f);
float start_y = -1+DEF_PARTICLE_RADIUS;//START_POS_Y - ARRAY_SIZE_Y * DIST * btScalar(0.5f);
float start_z = -1+DEF_PARTICLE_RADIUS;//START_POS_Z - ARRAY_SIZE_Z * DIST * btScalar(0.5f);
int numParticles = ARRAY_SIZE_X * ARRAY_SIZE_Y * ARRAY_SIZE_Z;
m_pWorld->m_hPos.resize(numParticles);
m_pWorld->m_hVel.resize(numParticles);
btScalar spacing = 2 * DEF_PARTICLE_RADIUS;
for(int z=0; z<ARRAY_SIZE_Z; z++)
{
for(int y=0; y<ARRAY_SIZE_Y; y++)
{
for(int x=0; x<ARRAY_SIZE_X; x++)
{
int i = (z * ARRAY_SIZE_Y * ARRAY_SIZE_X) + (y * ARRAY_SIZE_X) + x;
if (i < numParticles)
{
btVector3 jitter = 0.01f * 0.03f * btVector3(frand(), frand(), frand());
m_pWorld->m_hVel[i]= btVector3(0,0,0);
m_pWorld->m_hPos[i].setX((spacing * x) + 2*DEF_PARTICLE_RADIUS -WORLD_SIZE+jitter.getX());
m_pWorld->m_hPos[i].setY((spacing * y) + 2*DEF_PARTICLE_RADIUS -WORLD_SIZE+jitter.getY());
m_pWorld->m_hPos[i].setZ((spacing * z) + 2*DEF_PARTICLE_RADIUS -WORLD_SIZE+jitter.getZ());
}
}
}
}
m_pWorld->m_numParticles = numParticles;
}
void ParticlesDemo::clientResetScene()
{
static bool bFirstCall = true;
DemoApplication::clientResetScene();
init_scene_directly();
if(bFirstCall)
{
bFirstCall = false;
}
else
{
m_pWorld->grabSimulationData();
}
}
void ParticlesDemo::exitPhysics()
{
delete m_dialogDynamicsWorld;
m_dialogDynamicsWorld = 0;
//cleanup in the reverse order of creation/initialization
int i;
//remove the rigidbodies from the dynamics world and delete them
for (i=m_pWorld->getNumCollisionObjects()-1; i>=0 ;i--)
{
btCollisionObject* obj = m_pWorld->getCollisionObjectArray()[i];
btRigidBody* body = btRigidBody::upcast(obj);
if (body && body->getMotionState())
{
delete body->getMotionState();
}
m_pWorld->removeCollisionObject( obj );
delete obj;
}
//delete collision shapes
for (int j=0;j<m_collisionShapes.size();j++)
{
btCollisionShape* shape = m_collisionShapes[j];
delete shape;
}
delete m_pWorld;
delete m_solver;
delete m_broadphase;
delete m_dispatcher;
delete m_collisionConfiguration;
}
void ParticlesDemo::keyboardCallback(unsigned char key, int x, int y)
{
(void)x;
(void)y;
switch (key)
{
case 'G' :
{
m_drawGridMode++;
m_drawGridMode %= 3;
}
break;
case 'q' :
exitPhysics();
exit(0);
break;
default :
{
DemoApplication::keyboardCallback(key, x, y);
}
break;
}
if(key == ' ')
{
}
}
void ParticlesDemo::renderme()
{
glColor3f(1.0, 1.0, 1.0);
glutWireCube(2*WORLD_SIZE);
glPointSize(5.0f);
glEnable(GL_POINT_SPRITE_ARB);
glTexEnvi(GL_POINT_SPRITE_ARB, GL_COORD_REPLACE_ARB, GL_TRUE);
#ifndef __APPLE__
// glEnable(GL_VERTEX_PROGRAM_POINT_SIZE_NV);
glEnable(GL_VERTEX_PROGRAM_POINT_SIZE);
#endif //__APPLE__
glDepthMask(GL_TRUE);
glEnable(GL_DEPTH_TEST);
glUseProgram(m_shaderProgram);
btScalar dist = (m_glutScreenWidth > m_glutScreenHeight) ? m_glutScreenHeight : m_glutScreenWidth;
glUniform1f( glGetUniformLocation(m_shaderProgram, "pointScale"), dist );
// glUniform1f( glGetUniformLocation(m_shaderProgram, "pointRadius"), 0.5f );
int numParticles = m_pWorld->getNumParticles();
int col_vbo = m_pWorld->m_colVbo;
int curr_vbo = m_pWorld->m_vbo;
float sphere_rad = m_pWorld->m_particleRad;
glUniform1f( glGetUniformLocation(m_shaderProgram, "pointRadius"), sphere_rad );
glColor3f(1, 1, 1);
// render from the vbo
glBindBuffer(GL_ARRAY_BUFFER, curr_vbo);
glVertexPointer(4, GL_FLOAT, 0, 0);
glEnableClientState(GL_VERTEX_ARRAY);
if(col_vbo)
{
glBindBufferARB(GL_ARRAY_BUFFER_ARB, col_vbo);
glColorPointer(4, GL_FLOAT, 0, 0);
glEnableClientState(GL_COLOR_ARRAY);
}
glDrawArrays(GL_POINTS, 0, numParticles);
glDisableClientState(GL_VERTEX_ARRAY);
glDisableClientState(GL_COLOR_ARRAY);
glUseProgram(0);
glDisable(GL_POINT_SPRITE_ARB);
glBindBufferARB(GL_ARRAY_BUFFER,0);
if(m_drawGridMode)
{
btVector3& wmin = m_pWorld->m_worldMin;
btVector3& wmax = m_pWorld->m_worldMax;
glBegin(GL_LINE_LOOP);
glVertex3f(wmin[0], wmin[1], wmin[2]);
glVertex3f(wmin[0], wmax[1], wmin[2]);
glVertex3f(wmax[0], wmax[1], wmin[2]);
glVertex3f(wmax[0], wmin[1], wmin[2]);
glVertex3f(wmax[0], wmin[1], wmax[2]);
glVertex3f(wmax[0], wmax[1], wmax[2]);
glVertex3f(wmin[0], wmax[1], wmax[2]);
glVertex3f(wmin[0], wmin[1], wmax[2]);
glEnd();
glBegin(GL_LINES);
glVertex3f(wmin[0], wmin[1], wmin[2]);
glVertex3f(wmax[0], wmin[1], wmin[2]);
glVertex3f(wmin[0], wmin[1], wmax[2]);
glVertex3f(wmax[0], wmin[1], wmax[2]);
glVertex3f(wmin[0], wmax[1], wmin[2]);
glVertex3f(wmin[0], wmax[1], wmax[2]);
glVertex3f(wmax[0], wmax[1], wmin[2]);
glVertex3f(wmax[0], wmax[1], wmax[2]);
glEnd();
if(m_drawGridMode == 2)
{
int szx = m_pWorld->m_simParams.m_gridSize[0];
int szy = m_pWorld->m_simParams.m_gridSize[1];
glBegin(GL_LINES);
for(int i = 1; i < (szx-1); i++)
{
float wgt = (float)i / (float)(szx-1);
btVector3 vtx = wmax * wgt + wmin * (1.0f - wgt);
glVertex3f(vtx[0], wmin[1], wmin[2]);
glVertex3f(vtx[0], wmax[1], wmin[2]);
}
for(int i = 1; i < (szy-1); i++)
{
float wgt = (float)i / (float)(szy-1);
btVector3 vtx = wmax * wgt + wmin * (1.0f - wgt);
glVertex3f(wmin[0], vtx[1], wmin[2]);
glVertex3f(wmax[0], vtx[1], wmin[2]);
}
glEnd();
}
}
if ((m_debugMode & btIDebugDraw::DBG_NoHelpText)==0)
{
setOrthographicProjection();
int xOffset = 10.f;
int yStart = 20.f;
int yIncr = 20.f;
showProfileInfo(xOffset, yStart, yIncr);
outputDebugInfo(xOffset, yStart, yIncr);
resetPerspectiveProjection();
}
}
void ParticlesDemo::outputDebugInfo(int & xOffset,int & yStart, int yIncr)
{
char buf[124];
glDisable(GL_LIGHTING);
glColor3f(0, 0, 0);
sprintf(buf,"mouse move+buttons to interact");
GLDebugDrawString(xOffset,yStart,buf);
yStart += yIncr;
sprintf(buf,"space to reset");
GLDebugDrawString(xOffset,yStart,buf);
yStart += yIncr;
sprintf(buf,"cursor keys and z,x to navigate");
GLDebugDrawString(xOffset,yStart,buf);
yStart += yIncr;
sprintf(buf,"i to toggle simulation, s single step");
GLDebugDrawString(xOffset,yStart,buf);
yStart += yIncr;
sprintf(buf,"q to quit");
GLDebugDrawString(xOffset,yStart,buf);
yStart += yIncr;
sprintf(buf,"h to toggle help text");
GLDebugDrawString(xOffset,yStart,buf);
yStart += yIncr;
sprintf(buf,"p to toggle profiling (+results to file)");
GLDebugDrawString(xOffset,yStart,buf);
yStart += yIncr;
sprintf(buf,"j to toggle between demos (integration/OECake2D/OECake3D)");
GLDebugDrawString(xOffset,yStart,buf);
yStart += yIncr;
{
sprintf(buf,"G to draw broadphase grid");
GLDebugDrawString(xOffset,yStart,buf);
yStart += yIncr;
sprintf(buf,"D and U to toggle between GPU and CPU");
GLDebugDrawString(xOffset,yStart,buf);
yStart += yIncr;
}
}
GLuint _compileProgram(const char *vsource, const char *fsource)
{
GLuint vertexShader = glCreateShader(GL_VERTEX_SHADER);
GLuint fragmentShader = glCreateShader(GL_FRAGMENT_SHADER);
glShaderSource(vertexShader, 1, &vsource, 0);
glShaderSource(fragmentShader, 1, &fsource, 0);
glCompileShader(vertexShader);
glCompileShader(fragmentShader);
GLuint program = glCreateProgram();
glAttachShader(program, vertexShader);
glAttachShader(program, fragmentShader);
glLinkProgram(program);
// check if program linked
GLint success = 0;
glGetProgramiv(program, GL_LINK_STATUS, &success);
if (!success) {
char temp[256];
glGetProgramInfoLog(program, 256, 0, temp);
printf("Failed to link program:\n%s\n", temp);
glDeleteProgram(program);
program = 0;
}
return program;
}
void ParticlesDemo::myinit()
{
DemoApplication::myinit();
#ifndef __APPLE__
glewInit();
if (!glewIsSupported("GL_VERSION_2_0 GL_VERSION_1_5 GL_ARB_multitexture GL_ARB_vertex_buffer_object")) {
fprintf(stderr, "Required OpenGL extensions missing.");
exit(-1);
}
#endif //__APPLE__
m_shaderProgram = _compileProgram(vertexShader, spherePixelShader);
m_pWorld->initCLKernels(m_argc, m_argv);
}
void ParticlesDemo::mouseFunc(int button, int state, int x, int y)
{
if (!m_dialogDynamicsWorld->mouseFunc(button,state,x,y))
{
DemoApplication::mouseFunc(button,state,x,y);
}
}
void ParticlesDemo::mouseMotionFunc(int x,int y)
{
m_dialogDynamicsWorld->mouseMotionFunc(x,y);
DemoApplication::mouseMotionFunc(x,y);
}
void ParticlesDemo::reshape(int w, int h)
{
if (m_dialogDynamicsWorld)
m_dialogDynamicsWorld->setScreenSize(w,h);
GlutDemoApplication::reshape(w,h);
}

128
Demos/ParticlesOpenCL/ParticlesDemo.h
View File

@@ -1,128 +0,0 @@
/*
Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org
Copyright (C) 2006 - 2009 Sony Computer Entertainment Inc.
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 SPHERES_GRID_DEMO_H
#define SPHERES_GRID_DEMO_H
#define USE_BULLET_BODIES 0
///enable one or both options (NVidia profiler doesn't support multi-OpenCL context)
#define INTEGRATION_DEMO 1
#define SPHERES_DEMO 1
#include "DemoApplication.h"
#include "LinearMath/btAlignedObjectArray.h"
#include "BulletDynamics/Dynamics/btDiscreteDynamicsWorld.h"
#include "btParticlesDynamicsWorld.h"
class btBroadphaseInterface;
class btCollisionShape;
class btOverlappingPairCache;
class btCollisionDispatcher;
class btConstraintSolver;
struct btCollisionAlgorithmCreateFunc;
class btDefaultCollisionConfiguration;
#include "../OpenGL/GlutDemoApplication.h"
///BasicDemo is good starting point for learning the code base and porting.
class ParticlesDemo : public GlutDemoApplication
{
class GL_DialogDynamicsWorld* m_dialogDynamicsWorld;
//keep the collision shapes, for deletion/cleanup
btAlignedObjectArray<btCollisionShape*> m_collisionShapes;
btBroadphaseInterface* m_broadphase;
btCollisionDispatcher* m_dispatcher;
btConstraintSolver* m_solver;
btDefaultCollisionConfiguration* m_collisionConfiguration;
btOverlappingPairCache* m_pairCache;
int m_mouseButtons;
int m_mouseOldX;
int m_mouseOldY;
int m_argc;
char** m_argv;
public:
int m_drawGridMode;
int m_GpuCpuTogglePtr;
btParticlesDynamicsWorld* m_pWorld;
// shader
GLuint m_shaderProgram;
ParticlesDemo(int argc, char** argv)
{
m_argc = argc;
m_argv = argv;
m_drawGridMode = 0;
m_GpuCpuTogglePtr = SIMSTAGE_NONE;
m_dialogDynamicsWorld = 0;
}
virtual ~ParticlesDemo()
{
exitPhysics();
}
void initPhysics();
void exitPhysics();
virtual void clientMoveAndDisplay();
virtual void displayCallback();
virtual void keyboardCallback(unsigned char key, int x, int y);
virtual void clientResetScene();
virtual void mouseFunc(int button, int state, int x, int y);
virtual void mouseMotionFunc(int x,int y);
virtual void reshape(int w, int h);
/*
static DemoApplication* Create()
{
SpheresGridDemo* demo = new SpheresGridDemo;
demo->myinit();
demo->initPhysics();
demo->m_mouseButtons = 0;
demo->m_mouseOldX = 0;
demo->m_mouseOldY = 0;
return demo;
}
*/
void outputDebugInfo(int & xOffset,int & yStart, int yIncr);
virtual void renderme();
virtual void myinit();
void init_scene_directly();
};
#endif // PARTICLES_DEMO_H

467
Demos/ParticlesOpenCL/ParticlesOCL.cl
View File

@@ -1,467 +0,0 @@
MSTRINGIFY(
/*
Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org
Copyright (C) 2006 - 2009 Sony Computer Entertainment Inc.
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.
*/
int4 getGridPos(float4 worldPos, __global float4* pParams)
{
int4 gridPos;
gridPos.x = (int)floor((worldPos.x - pParams[1].x) / pParams[3].x);
gridPos.y = (int)floor((worldPos.y - pParams[1].y) / pParams[3].y);
gridPos.z = (int)floor((worldPos.z - pParams[1].z) / pParams[3].z);
return gridPos;
}
unsigned int getPosHash(int4 gridPos, __global float4* pParams)
{
int4 gridDim = *((__global int4*)(pParams + 4));
if(gridPos.x < 0) gridPos.x = 0;
if(gridPos.x >= gridDim.x) gridPos.x = gridDim.x - 1;
if(gridPos.y < 0) gridPos.y = 0;
if(gridPos.y >= gridDim.y) gridPos.y = gridDim.y - 1;
if(gridPos.z < 0) gridPos.z = 0;
if(gridPos.z >= gridDim.z) gridPos.z = gridDim.z - 1;
unsigned int hash = gridPos.z * gridDim.y * gridDim.x + gridPos.y * gridDim.x + gridPos.x;
return hash;
}
__kernel void kComputeCellId( int numParticles,
__global float4* pPos,
__global int2* pPosHash,
__global float4* pParams GUID_ARG)
{
int index = get_global_id(0);
if(index >= numParticles)
{
return;
}
float4 pos = pPos[index];
int4 gridPos = getGridPos(pos, pParams);
unsigned int hash = getPosHash(gridPos, pParams);
pPosHash[index].x = hash;
pPosHash[index].y = index;
}
__kernel void kClearCellStart( int numCells,
__global int* pCellStart GUID_ARG)
{
int index = get_global_id(0);
if(index >= numCells)
{
return;
}
pCellStart[index] = -1;
}
__kernel void kFindCellStart( int numParticles,
__global int2* pHash,
__global int* cellStart,
__global float4* pPos,
__global float4* pVel,
__global float4* pSortedPos,
__global float4* pSortedVel GUID_ARG)
{
int index = get_global_id(0);
__local int sharedHash[1025];//maximum workgroup size 1024
int2 sortedData;
if(index < numParticles)
{
sortedData = pHash[index];
// Load hash data into shared memory so that we can look
// at neighboring body's hash value without loading
// two hash values per thread
sharedHash[get_local_id(0) + 1] = sortedData.x;
if((index > 0) && (get_local_id(0) == 0))
{
// first thread in block must load neighbor body hash
sharedHash[0] = pHash[index-1].x;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if(index < numParticles)
{
if((index == 0) || (sortedData.x != sharedHash[get_local_id(0)]))
{
cellStart[sortedData.x] = index;
}
int unsortedIndex = sortedData.y;
float4 pos = pPos[unsortedIndex];
float4 vel = pVel[unsortedIndex];
pSortedPos[index] = pos;
pSortedVel[index] = vel;
}
}
__kernel void kIntegrateMotion( int numParticles,
__global float4* pPos,
__global float4* pVel,
__global float4* pParams,
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 = *((__global float4*)(pParams + 0));
float particleRad = pParams[5].x;
float globalDamping = pParams[5].y;
float boundaryDamping = pParams[5].z;
vel += gravity * timeStep;
vel *= globalDamping;
// integrate position
pos += vel * timeStep;
// collide with world boundaries
float4 worldMin = *((__global float4*)(pParams + 1));
float4 worldMax = *((__global float4*)(pParams + 2));
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;
}
float4 collideTwoParticles(
float4 posA,
float4 posB,
float4 velA,
float4 velB,
float radiusA,
float radiusB,
float spring,
float damping,
float shear,
float attraction
)
{
//Calculate relative position
float4 relPos = posB - posA; relPos.w = 0.f;
float dist = sqrt(relPos.x * relPos.x + relPos.y * relPos.y + relPos.z * relPos.z);
float collideDist = radiusA + radiusB;
float4 force = (float4)0.f;
if(dist < collideDist){
float4 norm = relPos * (1.f / dist); norm.w = 0.f;
//Relative velocity
float4 relVel = velB - velA; relVel.w = 0.f;
//Relative tangential velocity
float relVelDotNorm = relVel.x * norm.x + relVel.y * norm.y + relVel.z * norm.z;
float4 tanVel = relVel - norm * relVelDotNorm; tanVel.w = 0.f;
//Spring force (potential)
float springFactor = -spring * (collideDist - dist);
force = springFactor * norm + damping * relVel + shear * tanVel + attraction * relPos;
force.w = 0.f;
}
return force;
}
__kernel void kCollideParticles(int numParticles,
__global float4* pVel, //output: new velocity
__global const float4* pSortedPos, //input: reordered positions
__global const float4* pSortedVel, //input: reordered velocities
__global const int2 *pPosHash, //input: reordered particle indices
__global const int *pCellStart, //input: cell boundaries
__global float4* pParams GUID_ARG)
{
int index = get_global_id(0);
if(index >= numParticles)
{
return;
}
float4 posA = pSortedPos[index];
float4 velA = pSortedVel[index];
float4 force = (float4)0.f;
float particleRad = pParams[5].x;
float collisionDamping = pParams[5].w;
float spring = pParams[6].x;
float shear = pParams[6].y;
float attraction = pParams[6].z;
int unsortedIndex = pPosHash[index].y;
//Get address in grid
int4 gridPosA = getGridPos(posA, pParams);
//Accumulate surrounding cells
int4 gridPosB;
for(int z = -1; z <= 1; z++)
{
gridPosB.z = gridPosA.z + z;
for(int y = -1; y <= 1; y++)
{
gridPosB.y = gridPosA.y + y;
for(int x = -1; x <= 1; x++)
{
gridPosB.x = gridPosA.x + x;
//Get start particle index for this cell
uint hashB = getPosHash(gridPosB, pParams);
int startI = pCellStart[hashB];
//Skip empty cell
if(startI < 0)
{
continue;
}
//Iterate over particles in this cell
int endI = startI + 32;
if(endI >= numParticles)
endI = numParticles ;
for(int j = startI; j < endI; j++)
{
uint hashC = pPosHash[j].x;
if(hashC != hashB)
{
break;
}
if(j == index)
{
continue;
}
float4 posB = pSortedPos[j];
float4 velB = pSortedVel[j];
//Collide two spheres
force += collideTwoParticles( posA, posB, velA, velB, particleRad, particleRad,
spring, collisionDamping, shear, attraction);
}
}
}
}
//Write new velocity back to original unsorted location
pVel[unsortedIndex] = velA + force;
}
/*
* Copyright 1993-2009 NVIDIA Corporation. All rights reserved.
*
* NVIDIA Corporation and its licensors retain all intellectual property and
* proprietary rights in and to this software and related documentation.
* Any use, reproduction, disclosure, or distribution of this software
* and related documentation without an express license agreement from
* NVIDIA Corporation is strictly prohibited.
*
* Please refer to the applicable NVIDIA end user license agreement (EULA)
* associated with this source code for terms and conditions that govern
* your use of this NVIDIA software.
*
*/
inline void ComparatorPrivate(int2* keyA, int2* keyB, uint dir)
{
if((keyA[0].x > keyB[0].x) == dir)
{
int2 tmp = *keyA;
*keyA = *keyB;
*keyB = tmp;
}
}
inline void ComparatorLocal(__local int2* keyA, __local int2* keyB, uint dir)
{
if((keyA[0].x > keyB[0].x) == dir)
{
int2 tmp = *keyA;
*keyA = *keyB;
*keyB = tmp;
}
}
////////////////////////////////////////////////////////////////////////////////
// Monolithic bitonic sort kernel for short arrays fitting into local memory
////////////////////////////////////////////////////////////////////////////////
__kernel void kBitonicSortCellIdLocal(__global int2* pKey, uint arrayLength, uint dir GUID_ARG)
{
__local int2 l_key[LOCAL_SIZE_MAX];
int localSizeLimit = get_local_size(0) * 2;
//Offset to the beginning of subbatch and load data
pKey += get_group_id(0) * localSizeLimit + get_local_id(0);
l_key[get_local_id(0) + 0] = pKey[ 0];
l_key[get_local_id(0) + (localSizeLimit / 2)] = pKey[(localSizeLimit / 2)];
for(uint size = 2; size < arrayLength; size <<= 1)
{
//Bitonic merge
uint ddd = dir ^ ( (get_local_id(0) & (size / 2)) != 0 );
for(uint stride = size / 2; stride > 0; stride >>= 1)
{
barrier(CLK_LOCAL_MEM_FENCE);
uint pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
ComparatorLocal(&l_key[pos + 0], &l_key[pos + stride], ddd);
}
}
//ddd == dir for the last bitonic merge step
{
for(uint stride = arrayLength / 2; stride > 0; stride >>= 1)
{
barrier(CLK_LOCAL_MEM_FENCE);
uint pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
ComparatorLocal(&l_key[pos + 0], &l_key[pos + stride], dir);
}
}
barrier(CLK_LOCAL_MEM_FENCE);
pKey[ 0] = l_key[get_local_id(0) + 0];
pKey[(localSizeLimit / 2)] = l_key[get_local_id(0) + (localSizeLimit / 2)];
}
////////////////////////////////////////////////////////////////////////////////
// Bitonic sort kernel for large arrays (not fitting into local memory)
////////////////////////////////////////////////////////////////////////////////
//Bottom-level bitonic sort
//Almost the same as bitonicSortLocal with the only exception
//of even / odd subarrays (of LOCAL_SIZE_LIMIT points) being
//sorted in opposite directions
__kernel void kBitonicSortCellIdLocal1(__global int2* pKey GUID_ARG)
{
__local int2 l_key[LOCAL_SIZE_MAX];
uint localSizeLimit = get_local_size(0) * 2;
//Offset to the beginning of subarray and load data
pKey += get_group_id(0) * localSizeLimit + get_local_id(0);
l_key[get_local_id(0) + 0] = pKey[ 0];
l_key[get_local_id(0) + (localSizeLimit / 2)] = pKey[(localSizeLimit / 2)];
uint comparatorI = get_global_id(0) & ((localSizeLimit / 2) - 1);
for(uint size = 2; size < localSizeLimit; size <<= 1)
{
//Bitonic merge
uint ddd = (comparatorI & (size / 2)) != 0;
for(uint stride = size / 2; stride > 0; stride >>= 1)
{
barrier(CLK_LOCAL_MEM_FENCE);
uint pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
ComparatorLocal(&l_key[pos + 0], &l_key[pos + stride], ddd);
}
}
//Odd / even arrays of localSizeLimit elements
//sorted in opposite directions
{
uint ddd = (get_group_id(0) & 1);
for(uint stride = localSizeLimit / 2; stride > 0; stride >>= 1)
{
barrier(CLK_LOCAL_MEM_FENCE);
uint pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
ComparatorLocal(&l_key[pos + 0], &l_key[pos + stride], ddd);
}
}
barrier(CLK_LOCAL_MEM_FENCE);
pKey[ 0] = l_key[get_local_id(0) + 0];
pKey[(localSizeLimit / 2)] = l_key[get_local_id(0) + (localSizeLimit / 2)];
}
//Bitonic merge iteration for 'stride' >= LOCAL_SIZE_LIMIT
__kernel void kBitonicSortCellIdMergeGlobal(__global int2* pKey, uint arrayLength, uint size, uint stride, uint dir GUID_ARG)
{
uint global_comparatorI = get_global_id(0);
uint comparatorI = global_comparatorI & (arrayLength / 2 - 1);
//Bitonic merge
uint ddd = dir ^ ( (comparatorI & (size / 2)) != 0 );
uint pos = 2 * global_comparatorI - (global_comparatorI & (stride - 1));
int2 keyA = pKey[pos + 0];
int2 keyB = pKey[pos + stride];
ComparatorPrivate(&keyA, &keyB, ddd);
pKey[pos + 0] = keyA;
pKey[pos + stride] = keyB;
}
//Combined bitonic merge steps for
//'size' > LOCAL_SIZE_LIMIT and 'stride' = [1 .. LOCAL_SIZE_LIMIT / 2]
__kernel void kBitonicSortCellIdMergeLocal(__global int2* pKey, uint arrayLength, uint stride, uint size, uint dir GUID_ARG)
{
__local int2 l_key[LOCAL_SIZE_MAX];
int localSizeLimit = get_local_size(0) * 2;
pKey += get_group_id(0) * localSizeLimit + get_local_id(0);
l_key[get_local_id(0) + 0] = pKey[ 0];
l_key[get_local_id(0) + (localSizeLimit / 2)] = pKey[(localSizeLimit / 2)];
//Bitonic merge
uint comparatorI = get_global_id(0) & ((arrayLength / 2) - 1);
uint ddd = dir ^ ( (comparatorI & (size / 2)) != 0 );
for(; stride > 0; stride >>= 1)
{
barrier(CLK_LOCAL_MEM_FENCE);
uint pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
ComparatorLocal(&l_key[pos + 0], &l_key[pos + stride], ddd);
}
barrier(CLK_LOCAL_MEM_FENCE);
pKey[ 0] = l_key[get_local_id(0) + 0];
pKey[(localSizeLimit / 2)] = l_key[get_local_id(0) + (localSizeLimit / 2)];
}
);

1177
Demos/ParticlesOpenCL/btParticlesDemoDynamicsWorld.cpp
View File

File diff suppressed because it is too large Load Diff

183
Demos/ParticlesOpenCL/btParticlesDynamicsWorld.h
View File

@@ -1,183 +0,0 @@
/*
Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org
Copyright (C) 2006 - 2009 Sony Computer Entertainment Inc.
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 BT_PARTICLES_DYNAMICS_WORLD_H
#define BT_PARTICLES_DYNAMICS_WORLD_H
#ifdef USE_MINICL
#include <MiniCL/cl.h>
#include <MiniCL/cl_gl.h>
#else
#ifdef __APPLE__
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#include <CL/cl_gl.h>
#endif __APPLE__
#endif
#include "BulletDynamics/Dynamics/btDiscreteDynamicsWorld.h"
#include "BulletDynamics/ConstraintSolver/btTypedConstraint.h"
#include "BulletDynamics/ConstraintSolver/btPoint2PointConstraint.h"
#include "btParticlesSharedDefs.h"
#include "btParticlesSharedTypes.h"
#define PARTICLES_MAX_NEIGHBORS (32)
#define DEF_PARTICLE_RADIUS (0.023f)
//#define WORLD_SIZE 1.9f
#define WORLD_SIZE 1.f
enum
{
PARTICLES_KERNEL_INTEGRATE_MOTION = 0,
PARTICLES_KERNEL_COMPUTE_CELL_ID,
PARTICLES_KERNEL_CLEAR_CELL_START,
PARTICLES_KERNEL_FIND_CELL_START,
PARTICLES_KERNEL_COLLIDE_PARTICLES,
PARTICLES_KERNEL_BITONIC_SORT_CELL_ID_LOCAL,
PARTICLES_KERNEL_BITONIC_SORT_CELL_ID_LOCAL_1,
PARTICLES_KERNEL_BITONIC_SORT_CELL_ID_MERGE_GLOBAL,
PARTICLES_KERNEL_BITONIC_SORT_CELL_ID_MERGE_LOCAL,
PARTICLES_KERNEL_TOTAL
};
enum
{
SIMSTAGE_NONE = 0,
SIMSTAGE_INTEGRATE_MOTION,
SIMSTAGE_COMPUTE_CELL_ID,
SIMSTAGE_SORT_CELL_ID,
SIMSTAGE_FIND_CELL_START,
SIMSTAGE_COLLIDE_PARTICLES,
SIMSTAGE_TOTAL
};
struct btKernelInfo
{
int m_Id;
cl_kernel m_kernel;
const char* m_name;
int m_workgroupSize;
};
class btParticlesDynamicsWorld : public btDiscreteDynamicsWorld
{
public:
int m_numParticles;
int m_usedDevice;
btScalar m_particleRad;
struct GL_ToggleControl* m_useCpuControls[SIMSTAGE_TOTAL];
protected:
int m_hashSize; // power of 2 >= m_numSpheres;
int m_numGridCells;
int m_maxNeighbors;
int m_numSolverIterations;
// CPU side data
public:
btAlignedObjectArray<btVector3> m_hPos;
btAlignedObjectArray<btVector3> m_hVel;
btAlignedObjectArray<btVector3> m_hSortedPos;
btAlignedObjectArray<btVector3> m_hSortedVel;
protected:
btAlignedObjectArray<btInt2> m_hPosHash;
btAlignedObjectArray<int> m_hCellStart;
// GPU side data
cl_mem m_dPos;
cl_mem m_dVel;
cl_mem m_dPosHash;
cl_mem m_dCellStart;
cl_mem m_dSimParams; // copy of m_simParams : global simulation paramerers such as gravity, etc.
cl_mem m_dSortedPos;
cl_mem m_dSortedVel;
// OpenCL
public:
cl_context m_cxMainContext;
cl_device_id m_cdDevice;
cl_command_queue m_cqCommandQue;
cl_program m_cpProgram;
protected:
btKernelInfo m_kernels[PARTICLES_KERNEL_TOTAL];
btVector3 m_cellSize;
public:
btVector3 m_worldMin;
btVector3 m_worldMax;
// vbo variables
GLuint m_vbo;
unsigned int m_posVbo;
unsigned int m_colVbo;
btSimParams m_simParams;
float m_timeStep;
int getNumParticles() { return m_numParticles; }
float* getPosBuffer() { return (float*)&(m_hPos[0]); }
btParticlesDynamicsWorld(btDispatcher* dispatcher,btBroadphaseInterface* pairCache,btConstraintSolver* constraintSolver,btCollisionConfiguration* collisionConfiguration,
int maxObjs , int maxNeighbors = PARTICLES_MAX_NEIGHBORS)
: btDiscreteDynamicsWorld(dispatcher, pairCache, constraintSolver, collisionConfiguration)
{
m_cxMainContext = 0;
m_usedDevice = 1;
// m_particleRad = btScalar(0.5f);
m_particleRad = DEF_PARTICLE_RADIUS;
m_simParams.m_gravity[0] = 0.f;
m_simParams.m_gravity[1] = -10.f;
m_simParams.m_gravity[2] = 0.f;
m_simParams.m_gravity[3] = 0.f;
m_numSolverIterations = 4;
}
virtual ~btParticlesDynamicsWorld();
virtual int stepSimulation( btScalar timeStep,int maxSubSteps=1, btScalar fixedTimeStep=btScalar(1.)/btScalar(60.));
void initDeviceData();
void initCLKernels(int argc, char** argv);
void createVBO();
void postInitDeviceData();
void getShapeData();
void allocateBuffers();
void grabSimulationData();
void adjustGrid();
void runIntegrateMotionKernel();
void runComputeCellIdKernel();
void runSortHashKernel();
void runFindCellStartKernel();
void runCollideParticlesKernel();
void initKernel(int kernelId, const char* pName);
void runKernelWithWorkgroupSize(int kernelId, int globalSize);
void bitonicSortNv(cl_mem pKey, unsigned int batch, unsigned int arrayLength, unsigned int dir);
void scanExclusiveLocal1(cl_mem d_Dst, cl_mem d_Src, unsigned int n, unsigned int size);
void scanExclusiveLocal2(cl_mem d_Buffer, cl_mem d_Dst, cl_mem d_Src, unsigned int n, unsigned int size);
void uniformUpdate(cl_mem d_Dst, cl_mem d_Buffer, unsigned int n);
void scanExclusive(cl_mem d_Dst, cl_mem d_Src, unsigned int arrayLength);
};
#endif //BT_PARTICLES_DYNAMICS_WORLD_H

14
Demos/ParticlesOpenCL/btParticlesSharedDefs.h
View File

@@ -1,14 +0,0 @@
/*
Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org
Copyright (C) 2006 - 2009 Sony Computer Entertainment Inc.
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.
*/

54
Demos/ParticlesOpenCL/btParticlesSharedTypes.h
View File

@@ -1,54 +0,0 @@
/*
Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org
Copyright (C) 2006 - 2009 Sony Computer Entertainment Inc.
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 BT_SPHERES_GRID_DEMO_SHARED_TYPES
#define BT_SPHERES_GRID_DEMO_SHARED_TYPES
struct btSimParams
{
float m_gravity[4];
float m_worldMin[4];
float m_worldMax[4];
float m_cellSize[4];
int m_gridSize[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;
};
struct btInt2
{
int x;
int y;
};
struct btInt4
{
int x;
int y;
int z;
int w;
};
#endif

51
Demos/ParticlesOpenCL/main.cpp
View File

@@ -1,51 +0,0 @@
/*
Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org
Copyright (C) 2006 - 2009 Sony Computer Entertainment Inc.
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 "ParticlesDemo.h"
#include "GlutStuff.h"
#include "GLDebugDrawer.h"
#include "btBulletDynamicsCommon.h"
#include "LinearMath/btHashMap.h"
// standard utility and system includes
//#include <oclUtils.h>
// Extra CL/GL include
//#include <CL/cl_gl.h>
GLDebugDrawer gDebugDrawer;
int main(int argc,char** argv)
{
// start logs
// oclSetLogFileName ("appSpheresGrid.txt");
// oclLog(LOGBOTH, 0.0, "appSpheresGrid.exe Starting...\n\n");
ParticlesDemo ccdDemo(argc, argv);
ccdDemo.initPhysics();
ccdDemo.getDynamicsWorld()->setDebugDrawer(&gDebugDrawer);
#ifdef CHECK_MEMORY_LEAKS
ccdDemo.exitPhysics();
#else
return glutmain(argc, argv,640,480,"Bullet Physics Demo. http://bulletphysics.com",&ccdDemo);
#endif
//default glut doesn't return from mainloop
return 0;
}

56
Demos/ParticlesOpenCL/shaders.cpp
View File

@@ -1,56 +0,0 @@
#define STRINGIFY(A) #A
//use most up-to-date AMD Radeon drivers to make point sprites work
//see also http://forums.amd.com/devforum/messageview.cfm?catid=392&threadid=129431
// vertex shader
const char *vertexShader = STRINGIFY(
uniform float pointRadius; // point size in world space
uniform float pointScale; // scale to calculate size in pixels
uniform float densityScale;
uniform float densityOffset;
varying vec3 posEye;
void main()
{
// calculate window-space point size
posEye = vec3(gl_ModelViewMatrix * vec4(gl_Vertex.xyz, 1.0));
float dist = length(posEye);
gl_PointSize = pointRadius * (pointScale / dist);
gl_TexCoord[0] = gl_MultiTexCoord0;
gl_Position = gl_ModelViewProjectionMatrix * vec4(gl_Vertex.xyz, 1.0);
gl_FrontColor = gl_Color;
}
);
// pixel shader for rendering points as shaded spheres
const char *spherePixelShader = STRINGIFY(
uniform float pointRadius; // point size in world space
varying vec3 posEye; // position of center in eye space
void main()
{
const vec3 lightDir = vec3(0.577, 0.577, 0.577);
const float shininess = 40.0;
// calculate normal from texture coordinates
vec3 N;
N.xy = gl_TexCoord[0].xy*vec2(2.0, -2.0) + vec2(-1.0, 1.0);
float mag = dot(N.xy, N.xy);
if (mag > 1.0) discard; // kill pixels outside circle
N.z = sqrt(1.0-mag);
// point on surface of sphere in eye space
vec3 spherePosEye = posEye + N*pointRadius;
// calculate lighting
float diffuse = max(0.0, dot(lightDir, N));
// gl_FragColor = gl_Color * diffuse;
vec3 v = normalize(-spherePosEye);
vec3 h = normalize(lightDir + v);
float specular = pow(max(0.0, dot(N, h)), shininess);
gl_FragColor = gl_Color * diffuse + specular;
}
);

2
Demos/ParticlesOpenCL/shaders.h
View File

@@ -1,2 +0,0 @@
extern const char *vertexShader;
extern const char *spherePixelShader;