diff --git a/Demos3/GpuDemos/ParticleDemo.cpp b/Demos3/GpuDemos/ParticleDemo.cpp index a12071128..7221d88e2 100644 --- a/Demos3/GpuDemos/ParticleDemo.cpp +++ b/Demos3/GpuDemos/ParticleDemo.cpp @@ -17,7 +17,7 @@ static char* particleKernelsString = //#include "../../opencl/broadphase_benchmark/b3GridBroadphaseCL.h" #include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h" #include "GpuDemoInternalData.h" - +#include "Bullet3Common/b3Random.h" //1000000 particles @@ -31,14 +31,14 @@ static char* particleKernelsString = //#define NUM_PARTICLES_Z 80 //256k particles -//#define NUM_PARTICLES_X 60 -//#define NUM_PARTICLES_Y 60 -//#define NUM_PARTICLES_Z 60 +#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 +//#define NUM_PARTICLES_X 30 +//#define NUM_PARTICLES_Y 30 +//#define NUM_PARTICLES_Z 30 @@ -62,10 +62,10 @@ B3_ATTRIBUTE_ALIGNED16(struct) b3SimParams b3SimParams() { - m_gravity.setValue(0,-0.03,0.f); - m_particleRad = 0.023f; + m_gravity.setValue(0,-.3,0.f); + m_particleRad = 0.01f; m_globalDamping = 1.0f; - m_boundaryDamping = -0.5f; + m_boundaryDamping = -1.f; m_collisionDamping = 0.025f;//0.02f; m_spring = 0.5f; m_shear = 0.1f; @@ -165,7 +165,7 @@ void ParticleDemo::setupScene(const ConstructionInfo& ci) 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; + float radius = m_data->m_simParamCPU[0].m_particleRad; m_data->m_broadphaseGPU = new b3GpuSapBroadphase(m_clData->m_clContext ,m_clData->m_clDevice,m_clData->m_clQueue);//overlappingPairCache,b3Vector3(4.f, 4.f, 4.f), 128, 128, 128,maxObjects, maxObjects, maxPairsSmallProxy, 100.f, 128, @@ -219,43 +219,61 @@ void ParticleDemo::setupScene(const ConstructionInfo& ci) 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}; + + float scaling[4] = {radius,radius,radius,1}; int userIndex = 0; - for (int x=0;xm_simParamCPU[0].m_particleRad; - position[0] = x*(rad*3); - position[1] = y*(rad*3); - position[2] = z*(rad*3); + + float arg = b3RandRange(-B3_PI,B3_PI); + + float rad = m_data->m_simParamCPU[0].m_particleRad; + position[0] = arg*b3Cos(arg + angle); + position[1] = 3.0f + arg; + position[2] = arg*b3Sin(arg + angle); - color[0] = float(x)/float(NUM_PARTICLES_X); - color[1] = float(y)/float(NUM_PARTICLES_Y); - color[2] = float(z)/float(NUM_PARTICLES_Z); + b3Vector4 color = colors[curColor]; + curColor++; + curColor&=3; - int id = m_instancingRenderer->registerGraphicsInstance(shapeId,position,quaternion,color,scaling); + + int id = m_instancingRenderer->registerGraphicsInstance(shapeId,position,quaternion,color,scaling); - void* userPtr = (void*)userIndex; - int collidableIndex = userIndex; - b3Vector3 aabbMin,aabbMax; - b3Vector3 particleRadius=b3MakeVector3(rad,rad,rad); + void* userPtr = (void*)userIndex; + int collidableIndex = userIndex; + b3Vector3 aabbMin,aabbMax; + b3Vector3 particleRadius=b3MakeVector3(rad,rad,rad); - aabbMin = b3MakeVector3(position[0],position[1],position[2])-particleRadius; - aabbMax = b3MakeVector3(position[0],position[1],position[2])+particleRadius; - m_data->m_broadphaseGPU->createProxy(aabbMin,aabbMax,collidableIndex,1,1); - userIndex++; - - } + aabbMin = b3MakeVector3(position[0],position[1],position[2])-particleRadius; + aabbMax = b3MakeVector3(position[0],position[1],position[2])+particleRadius; + //m_data->m_broadphaseGPU->createProxy(aabbMin,aabbMax,collidableIndex,1,1); + userIndex++; + angle += b3RandRange(-(float)B3_PI, (float)B3_PI); } } + + m_data->m_broadphaseGPU->writeAabbsToGpu(); - float camPos[4]={1.5,0.5,2.5,0}; + //float camPos[4]={1.5,0.5,2.5,0}; + float camPos[4]={0,0.5,0,0}; m_instancingRenderer->setCameraTargetPosition(camPos); m_instancingRenderer->setCameraDistance(4); m_instancingRenderer->writeTransforms(); @@ -376,6 +394,7 @@ void ParticleDemo::clientMoveAndDisplay() } + if (0) { b3BufferInfoCL bInfo[] = { b3BufferInfoCL( m_data->m_clPositionBuffer), @@ -440,39 +459,6 @@ void ParticleDemo::clientMoveAndDisplay() glUnmapBuffer( GL_ARRAY_BUFFER); glFlush(); - /* - int numParticles = NUM_PARTICLES_X*NUM_PARTICLES_Y*NUM_PARTICLES_Z; - for (int objectIndex=0;objectIndexwriteSingleInstanceTransformToGPU(pos,orn,i); - { - glBindBuffer(GL_ARRAY_BUFFER, m_instancingRenderer->getInternalData()->m_vbo); - glFlush(); - - char* orgBase = (char*)glMapBuffer( GL_ARRAY_BUFFER,GL_READ_WRITE); - //b3GraphicsInstance* 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; diff --git a/Demos3/GpuDemos/ParticleKernels.cl b/Demos3/GpuDemos/ParticleKernels.cl index 8a044dc93..34fed45c0 100644 --- a/Demos3/GpuDemos/ParticleKernels.cl +++ b/Demos3/GpuDemos/ParticleKernels.cl @@ -63,7 +63,7 @@ __kernel void integrateMotionKernel( int numParticles, 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; @@ -74,11 +74,13 @@ __kernel void integrateMotionKernel( int numParticles, 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; + vel.y *= -1.f;//1000*boundaryDamping; } + /* if(pos.y > (worldMax.y - 2*particleRad)) { pos.y = worldMax.y - 2*particleRad; @@ -94,6 +96,7 @@ __kernel void integrateMotionKernel( int numParticles, pos.z = worldMax.z - 2*particleRad; vel.z *= boundaryDamping; } + */ // write back position and velocity pPos[index] = pos; pVel[index] = vel; diff --git a/Demos3/GpuDemos/main_opengl3core.cpp b/Demos3/GpuDemos/main_opengl3core.cpp index ba84360fa..025f388d1 100644 --- a/Demos3/GpuDemos/main_opengl3core.cpp +++ b/Demos3/GpuDemos/main_opengl3core.cpp @@ -97,7 +97,7 @@ GpuDemo::CreateFunc* allDemos[]= // ConcaveSphereScene::MyCreateFunc, - + ConcaveScene::MyCreateFunc, GpuBoxPlaneScene::MyCreateFunc, @@ -140,7 +140,7 @@ GpuDemo::CreateFunc* allDemos[]= - //ParticleDemo::MyCreateFunc, + ParticleDemo::MyCreateFunc, @@ -763,7 +763,8 @@ int main(int argc, char* argv[]) bool useGpu = false; - int maxObjectCapacity=128*1024; + //int maxObjectCapacity=128*1024; + int maxObjectCapacity=1024*1024; maxObjectCapacity = b3Max(maxObjectCapacity,ci.arraySizeX*ci.arraySizeX*ci.arraySizeX+10); { diff --git a/btgui/OpenGLWindow/Shaders/pointSpritePS.glsl b/btgui/OpenGLWindow/Shaders/pointSpritePS.glsl index e69407485..9e8190be4 100644 --- a/btgui/OpenGLWindow/Shaders/pointSpritePS.glsl +++ b/btgui/OpenGLWindow/Shaders/pointSpritePS.glsl @@ -23,7 +23,7 @@ void main(void) N.xy = gl_PointCoord.st*vec2(2.0, -2.0) + vec2(-1.0, 1.0); float mag = dot(N.xy, N.xy); if (mag > 1.0) discard; - vec4 texel = vec4(1,0,0,1);//fragment.color*texture(Diffuse,vert.texcoord);//fragment.color; + vec4 texel = fragment.color;//vec4(1,0,0,1);//fragment.color*texture(Diffuse,vert.texcoord);//fragment.color; vec3 ct; float at,af; af = 1.0; diff --git a/btgui/OpenGLWindow/Shaders/pointSpritePS.h b/btgui/OpenGLWindow/Shaders/pointSpritePS.h index c6d73928d..e0d57e238 100644 --- a/btgui/OpenGLWindow/Shaders/pointSpritePS.h +++ b/btgui/OpenGLWindow/Shaders/pointSpritePS.h @@ -18,7 +18,7 @@ static const char* pointSpriteFragmentShader= \ " N.xy = gl_PointCoord.st*vec2(2.0, -2.0) + vec2(-1.0, 1.0);\n" " float mag = dot(N.xy, N.xy);\n" " if (mag > 1.0) discard; \n" -" vec4 texel = vec4(1,0,0,1);//fragment.color*texture(Diffuse,vert.texcoord);//fragment.color;\n" +" vec4 texel = fragment.color;//vec4(1,0,0,1);//fragment.color*texture(Diffuse,vert.texcoord);//fragment.color;\n" " vec3 ct;\n" " float at,af;\n" " af = 1.0;\n" diff --git a/src/Bullet3Common/b3Random.h b/src/Bullet3Common/b3Random.h index 884a11420..dc040f156 100644 --- a/src/Bullet3Common/b3Random.h +++ b/src/Bullet3Common/b3Random.h @@ -17,6 +17,8 @@ subject to the following restrictions: #ifndef B3_GEN_RANDOM_H #define B3_GEN_RANDOM_H +#include "b3Scalar.h" + #ifdef MT19937 #include @@ -38,5 +40,11 @@ B3_FORCE_INLINE unsigned int b3rand() { return rand(); } #endif +inline b3Scalar b3RandRange(b3Scalar minRange, b3Scalar maxRange) +{ + return (b3rand() / (b3Scalar(B3_RAND_MAX) + b3Scalar(1.0))) * (maxRange - minRange) + minRange; +} + + #endif //B3_GEN_RANDOM_H diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp index f779358ed..233da4dd1 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp +++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp @@ -1000,6 +1000,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairs(int maxPairs) //if (m_currentBuffer>=0) // return calculateOverlappingPairsHostIncremental3Sap(); + B3_PROFILE("GPU 1-axis SAP calculateOverlappingPairs"); int axis = 0; @@ -1149,6 +1150,7 @@ void b3GpuSapBroadphase::calculateOverlappingPairs(int maxPairs) clFinish(m_queue); } + if (m_gpuSmallSortData.size()) { B3_PROFILE("gpu radix sort\n"); m_sorter->execute(m_gpuSmallSortData); diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp index f293d23f9..f8fd380e4 100644 --- a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp +++ b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp @@ -3056,6 +3056,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray* launcher.launch1D( num); clFinish(m_queue); numConcavePairs = m_numConcavePairsOut.at(0); + //printf("numConcavePairs=%d (max = %d\n",numConcavePairs,maxTriConvexPairCapacity); if (numConcavePairs > maxTriConvexPairCapacity) {