fix some out-of-bounds error in the OpenCL rigid body pipeline

re-enable an OpenCL/gpu rigid body example (box-box stack)
This commit is contained in:
erwincoumans committed 2015-05-01 21:34:27 -07:00
1 parent 7ca20f529a
commit 01d14f538a
20 files changed
+2710 -12

No files matched your search

@@ -28,7 +28,14 @@
#include "../Experiments/ImplicitCloth/ImplicitClothExample.h" #include "../Experiments/ImplicitCloth/ImplicitClothExample.h"
#include "../Importers/ImportBullet/SerializeSetup.h" #include "../Importers/ImportBullet/SerializeSetup.h"
#include "../Raycast/RaytestDemo.h" #include "../Raycast/RaytestDemo.h"
#ifdef B3_USE_CLEW
#include "../OpenCL/broadphase/PairBench.h" #include "../OpenCL/broadphase/PairBench.h"
#include "../OpenCL/rigidbody/GpuConvexScene.h"
#endif //B3_USE_CLEW
struct ExampleEntry struct ExampleEntry
@@ -175,7 +182,10 @@ static ExampleEntry gDefaultExamples[]=
static ExampleEntry gOpenCLExamples[]= static ExampleEntry gOpenCLExamples[]=
{ {
ExampleEntry(0,"OpenCL (experimental)"), ExampleEntry(0,"OpenCL (experimental)"),
ExampleEntry(1,"Box-Box", "Full OpenCL implementation of the entire physics and collision detection pipeline, showing box-box rigid body",
OpenCLBoxBoxCreateFunc),
ExampleEntry(1,"Pair Bench", "Benchmark of overlapping pair search using OpenCL.", PairBenchOpenCLCreateFunc), ExampleEntry(1,"Pair Bench", "Benchmark of overlapping pair search using OpenCL.", PairBenchOpenCLCreateFunc),
}; };
#endif // #endif //
static btAlignedObjectArray<ExampleEntry> gAdditionalRegisteredExamples; static btAlignedObjectArray<ExampleEntry> gAdditionalRegisteredExamples;
@@ -69,6 +69,14 @@ static bool pauseSimulation=false;
int midiBaseIndex = 176; int midiBaseIndex = 176;
extern bool gDisableDeactivation; extern bool gDisableDeactivation;
///some quick test variable for the OpenCL examples
int gPreferredOpenCLDeviceIndex=-1;
int gPreferredOpenCLPlatformIndex=-1;
int gGpuArraySizeX=25;
int gGpuArraySizeY=25;
int gGpuArraySizeZ=25;
//#include <float.h> //#include <float.h>
//unsigned int fp_control_state = _controlfp(_EM_INEXACT, _MCW_EM); //unsigned int fp_control_state = _controlfp(_EM_INEXACT, _MCW_EM);
+4 -1
View File
@@ -59,6 +59,7 @@
"../Raycast/*", "../Raycast/*",
"../MultiBody/MultiDofDemo.cpp", "../MultiBody/MultiDofDemo.cpp",
"../MultiBody/TestJointTorqueSetup.cpp", "../MultiBody/TestJointTorqueSetup.cpp",
"../ThirdPartyLibs/stb_image/*",
"../ThirdPartyLibs/Wavefront/tiny_obj_loader.*", "../ThirdPartyLibs/Wavefront/tiny_obj_loader.*",
"../ThirdPartyLibs/tinyxml/*", "../ThirdPartyLibs/tinyxml/*",
"../Utils/b3Clock.*", "../Utils/b3Clock.*",
@@ -92,7 +93,9 @@
if (hasCL) then if (hasCL) then
files { files {
"../OpenCL/broadphase/*", "../OpenCL/broadphase/*",
"../OpenCL/CommonOpenCL/*" "../OpenCL/CommonOpenCL/*",
"../OpenCL/rigidbody/GpuConvexScene.cpp",
"../OpenCL/rigidbody/GpuRigidBodyDemo.cpp",
} }
end end
+5 -4
View File
@@ -20,6 +20,9 @@
#include "pairsKernel.h" #include "pairsKernel.h"
extern int gPreferredOpenCLDeviceIndex;
extern int gPreferredOpenCLPlatformIndex;
#include "../CommonInterfaces/CommonExampleInterface.h" #include "../CommonInterfaces/CommonExampleInterface.h"
#include "../CommonInterfaces/CommonGUIHelperInterface.h" #include "../CommonInterfaces/CommonGUIHelperInterface.h"
@@ -207,10 +210,8 @@ void PairBench::initPhysics()
int startItem = 0; int startItem = 0;
int preferredOpenCLDeviceIndex=-1;
int preferredOpenCLPlatformIndex=-1; initCL(gPreferredOpenCLDeviceIndex,gPreferredOpenCLPlatformIndex);
initCL(preferredOpenCLDeviceIndex,preferredOpenCLPlatformIndex);
if (m_clData->m_clContext) if (m_clData->m_clContext)
{ {
+758
View File
@@ -0,0 +1,758 @@
#include "ConcaveScene.h"
#include "GpuRigidBodyDemo.h"
#include "OpenGLWindow/ShapeData.h"
#include "OpenGLWindow/GLInstancingRenderer.h"
#include "Bullet3Common/b3Quaternion.h"
#include "OpenGLWindow/b3gWindowInterface.h"
#include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h"
#include "../GpuDemoInternalData.h"
#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
#include "OpenGLWindow/OpenGLInclude.h"
#include "OpenGLWindow/GLInstanceRendererInternalData.h"
#include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h"
#include "Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.h"
#include "Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.h"
#include "Bullet3Collision/NarrowPhaseCollision/b3Config.h"
#include "GpuRigidBodyDemoInternalData.h"
#include"../../Wavefront/tiny_obj_loader.h"
#include "Bullet3Common/b3Transform.h"
#include "Bullet3Collision/NarrowPhaseCollision/b3ConvexUtility.h"
#include "Bullet3AppSupport/gwenUserInterface.h"
#include "OpenGLWindow/GLInstanceGraphicsShape.h"
#define CONCAVE_GAPX 14
#define CONCAVE_GAPY 5
#define CONCAVE_GAPZ 14
GLInstanceGraphicsShape* createGraphicsShapeFromWavefrontObj(std::vector<tinyobj::shape_t>& shapes)
{
b3AlignedObjectArray<GLInstanceVertex>* vertices = new b3AlignedObjectArray<GLInstanceVertex>;
{
// int numVertices = obj->vertexCount;
// int numIndices = 0;
b3AlignedObjectArray<int>* indicesPtr = new b3AlignedObjectArray<int>;
for (int s=0;s<shapes.size();s++)
{
tinyobj::shape_t& shape = shapes[s];
int faceCount = shape.mesh.indices.size();
for (int f=0;f<faceCount;f+=3)
{
//b3Vector3 normal(face.m_plane[0],face.m_plane[1],face.m_plane[2]);
if (1)
{
b3Vector3 normal=b3MakeVector3(0,1,0);
int vtxBaseIndex = vertices->size();
indicesPtr->push_back(vtxBaseIndex);
indicesPtr->push_back(vtxBaseIndex+1);
indicesPtr->push_back(vtxBaseIndex+2);
GLInstanceVertex vtx0;
vtx0.xyzw[0] = shape.mesh.positions[shape.mesh.indices[f]*3+0];
vtx0.xyzw[1] = shape.mesh.positions[shape.mesh.indices[f]*3+1];
vtx0.xyzw[2] = shape.mesh.positions[shape.mesh.indices[f]*3+2];
vtx0.xyzw[3] = 0.f;
vtx0.uv[0] = 0.5f;//shape.mesh.positions[shape.mesh.indices[f]*3+2];?
vtx0.uv[1] = 0.5f;
GLInstanceVertex vtx1;
vtx1.xyzw[0] = shape.mesh.positions[shape.mesh.indices[f+1]*3+0];
vtx1.xyzw[1] = shape.mesh.positions[shape.mesh.indices[f+1]*3+1];
vtx1.xyzw[2] = shape.mesh.positions[shape.mesh.indices[f+1]*3+2];
vtx1.xyzw[3]= 0.f;
vtx1.uv[0] = 0.5f;//obj->textureList[face->vertex_index[1]]->e[0];
vtx1.uv[1] = 0.5f;//obj->textureList[face->vertex_index[1]]->e[1];
GLInstanceVertex vtx2;
vtx2.xyzw[0] = shape.mesh.positions[shape.mesh.indices[f+2]*3+0];
vtx2.xyzw[1] = shape.mesh.positions[shape.mesh.indices[f+2]*3+1];
vtx2.xyzw[2] = shape.mesh.positions[shape.mesh.indices[f+2]*3+2];
vtx2.xyzw[3] = 0.f;
vtx2.uv[0] = 0.5f;
vtx2.uv[1] = 0.5f;
b3Vector3 v0=b3MakeVector3(vtx0.xyzw[0],vtx0.xyzw[1],vtx0.xyzw[2]);
b3Vector3 v1=b3MakeVector3(vtx1.xyzw[0],vtx1.xyzw[1],vtx1.xyzw[2]);
b3Vector3 v2=b3MakeVector3(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);
}
}
}
GLInstanceGraphicsShape* gfxShape = new GLInstanceGraphicsShape;
gfxShape->m_vertices = vertices;
gfxShape->m_numvertices = vertices->size();
gfxShape->m_indices = indicesPtr;
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;
}
}
void ConcaveScene::createConcaveMesh(const ConstructionInfo& ci, const char* fileName, const b3Vector3& shift, const b3Vector3& scaling)
{
char relativeFileName[1024];
const char* prefix[]={"./data/","../data/","../../data/","../../../data/","../../../../data/"};
int prefixIndex=-1;
{
int numPrefixes = sizeof(prefix)/sizeof(char*);
for (int i=0;i<numPrefixes;i++)
{
FILE* f = 0;
sprintf(relativeFileName,"%s%s",prefix[i],fileName);
f = fopen(relativeFileName,"r");
if (f)
{
fclose(f);
prefixIndex = i;
break;
}
}
}
if (prefixIndex<0)
return;
int index=10;
{
std::vector<tinyobj::shape_t> shapes;
std::string err = tinyobj::LoadObj(shapes, relativeFileName, prefix[prefixIndex]);
GLInstanceGraphicsShape* shape = createGraphicsShapeFromWavefrontObj(shapes);
b3AlignedObjectArray<b3Vector3> verts;
for (int i=0;i<shape->m_numvertices;i++)
{
for (int j=0;j<3;j++)
shape->m_vertices->at(i).xyzw[j] += shift[j];
b3Vector3 vtx=b3MakeVector3(shape->m_vertices->at(i).xyzw[0],
shape->m_vertices->at(i).xyzw[1],
shape->m_vertices->at(i).xyzw[2]);
verts.push_back(vtx*scaling);
}
int colIndex = m_data->m_np->registerConcaveMesh(&verts,shape->m_indices,b3MakeVector3(1,1,1));
{
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int shapeId = ci.m_instancingRenderer->registerShape(&shape->m_vertices->at(0).xyzw[0], shape->m_numvertices, &shape->m_indices->at(0), shape->m_numIndices);
b3Quaternion orn(0,0,0,1);
b3Vector4 color=b3MakeVector4(0.3,0.3,1,1.f);//0.5);//1.f
{
float mass = 0.f;
b3Vector3 position=b3MakeVector3(0,0,0);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index,false);
index++;
}
delete shape->m_indices;
delete shape->m_vertices;
delete shape;
}
}
}
void ConcaveScene::setupScene(const ConstructionInfo& ci)
{
if (1)
{
//char* fileName = "slopedPlane100.obj";
//char* fileName = "plane100.obj";
// char* fileName = "plane100.obj";
//char* fileName = "teddy.obj";//"plane.obj";
// char* fileName = "sponza_closed.obj";//"plane.obj";
//char* fileName = "leoTest1.obj";
const char* fileName = "samurai_monastry.obj";
// char* fileName = "teddy2_VHACD_CHs.obj";
b3Vector3 shift1=b3MakeVector3(0,0,0);//0,230,80);//150,-100,-120);
b3Vector4 scaling=b3MakeVector4(10,10,10,1);
// createConcaveMesh(ci,"plane100.obj",shift1,scaling);
//createConcaveMesh(ci,"plane100.obj",shift,scaling);
// b3Vector3 shift2(0,0,0);//0,230,80);//150,-100,-120);
// createConcaveMesh(ci,"teddy.obj",shift2,scaling);
// b3Vector3 shift3(130,-150,-75);//0,230,80);//150,-100,-120);
// createConcaveMesh(ci,"leoTest1.obj",shift3,scaling);
createConcaveMesh(ci,fileName,shift1,scaling);
} else
{
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int group=1;
int mask=1;
int index=0;
{
b3Vector4 scaling=b3MakeVector4(400,1.,400,1);
int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
b3Vector3 position=b3MakeVector3(0,-2,0);
b3Quaternion orn(0,0,0,1);
b3Vector4 color=b3MakeVector4(0,0,1,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0.f,position,orn,colIndex,index,false);
}
}
createDynamicObjects(ci);
m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
float camPos[4]={0,0,0,0};//65.5,4.5,65.5,0};
//float camPos[4]={1,12.5,1.5,0};
m_instancingRenderer->setCameraPitch(45);
m_instancingRenderer->setCameraTargetPosition(camPos);
m_instancingRenderer->setCameraDistance(355);
char msg[1024];
int numInstances = m_data->m_rigidBodyPipeline->getNumBodies();
sprintf(msg,"Num objects = %d",numInstances);
if (ci.m_gui)
ci.m_gui->setStatusBarMessage(msg,true);
}
void ConcaveScene::createDynamicObjects(const ConstructionInfo& ci)
{
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int group=1;
int mask=1;
int index=0;
if (1)
{
int curColor = 0;
b3Vector4 colors[4] =
{
b3MakeVector4(1,1,1,1),
b3MakeVector4(1,1,0.3,1),
b3MakeVector4(0.3,1,1,1),
b3MakeVector4(0.3,0.3,1,1),
};
b3ConvexUtility* utilPtr = new b3ConvexUtility();
b3Vector4 scaling=b3MakeVector4(1,1,1,1);
{
b3AlignedObjectArray<b3Vector3> verts;
unsigned char* vts = (unsigned char*) cube_vertices;
for (int i=0;i<numVertices;i++)
{
float* vertex = (float*) &vts[i*strideInBytes];
verts.push_back(b3MakeVector3(vertex[0]*scaling[0],vertex[1]*scaling[1],vertex[2]*scaling[2]));
}
bool merge = true;
if (numVertices)
{
utilPtr->initializePolyhedralFeatures(&verts[0],verts.size(),merge);
}
}
// int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
int colIndex=-1;
if (ci.m_useInstancedCollisionShapes)
colIndex = m_data->m_np->registerConvexHullShape(utilPtr);
for (int i=0;i<ci.arraySizeX;i++)
{
for (int j=0;j<ci.arraySizeY;j++)
{
for (int k=0;k<ci.arraySizeZ;k++)
{
if (!ci.m_useInstancedCollisionShapes)
colIndex = m_data->m_np->registerConvexHullShape(utilPtr);
float mass = 1;
//b3Vector3 position(-2*ci.gapX+i*ci.gapX,25+j*ci.gapY,-2*ci.gapZ+k*ci.gapZ);
b3Vector3 position=b3MakeVector3(-(ci.arraySizeX/2)*CONCAVE_GAPX+i*CONCAVE_GAPX,
23+j*CONCAVE_GAPY,
-(ci.arraySizeZ/2)*CONCAVE_GAPZ+k*CONCAVE_GAPZ);
b3Quaternion orn(0,0,0,1);
b3Vector4 color = colors[curColor];
curColor++;
curColor&=3;
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index,false);
index++;
}
}
}
}
}
void ConcaveCompoundScene::setupScene(const ConstructionInfo& ci)
{
ConcaveScene::setupScene(ci);
float camPos[4]={0,50,0,0};//65.5,4.5,65.5,0};
//float camPos[4]={1,12.5,1.5,0};
m_instancingRenderer->setCameraPitch(45);
m_instancingRenderer->setCameraTargetPosition(camPos);
m_instancingRenderer->setCameraDistance(40);
}
void ConcaveCompound2Scene::createDynamicObjects(const ConstructionInfo& ci)
{
const char* fileName = "teddy2_VHACD_CHs.obj";
//char* fileName = "cube_offset.obj";
b3Vector3 shift=b3MakeVector3(0,0,0);//0,230,80);//150,-100,-120);
b3Vector4 scaling=b3MakeVector4(1,1,1,1);
const char* prefix[]={"./data/","../data/","../../data/","../../../data/","../../../../data/"};
int prefixIndex=-1;
char relativeFileName[1024];
{
int numPrefixes = sizeof(prefix)/sizeof(char*);
for (int i=0;i<numPrefixes;i++)
{
sprintf(relativeFileName,"%s%s",prefix[i],fileName);
FILE* f = 0;
f = fopen(relativeFileName,"r");
if (f)
{
prefixIndex = i;
fclose(f);
break;
}
}
}
if (prefixIndex<0)
return;
std::vector<tinyobj::shape_t> shapes;
std::string err = tinyobj::LoadObj(shapes, relativeFileName, prefix[prefixIndex]);
if (shapes.size()>0)
{
int strideInBytes = 9*sizeof(float);
b3AlignedObjectArray<GLInstanceVertex> vertexArray;
b3AlignedObjectArray<int> indexArray;
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int group=1;
int mask=1;
int index=0;
int colIndex = 0;
b3AlignedObjectArray<GLInstanceVertex> vertices;
int stride2 = sizeof(GLInstanceVertex);
b3Assert(stride2 == strideInBytes);
{
b3AlignedObjectArray<b3GpuChildShape> childShapes;
int numChildShapes = shapes.size();
for (int i=0;i<numChildShapes;i++)
// int i=4;
{
tinyobj::shape_t& shape = shapes[i];
int numVertices = shape.mesh.positions.size()/3;
int numFaces = shape.mesh.indices.size()/3;
//for now, only support polyhedral child shapes
b3GpuChildShape child;
b3Vector3 pos=b3MakeVector3(0,0,0);
b3Quaternion orn(0,0,0,1);
for (int v=0;v<4;v++)
{
child.m_childPosition[v] = pos[v];
child.m_childOrientation[v] = orn[v];
}
b3Transform tr;
tr.setIdentity();
tr.setOrigin(pos);
tr.setRotation(orn);
int baseIndex = vertexArray.size();
for (int f=0;f<numFaces;f++)
{
for (int i=0;i<3;i++)
{
indexArray.push_back(baseIndex+shape.mesh.indices[f*3+i]);
}
}
b3Vector3 center=b3MakeVector3(0,0,0);
b3AlignedObjectArray<GLInstanceVertex> tmpVertices;
//add transformed graphics vertices and indices
b3Vector3 myScaling=b3MakeVector3(50,50,50);//300,300,300);
for (int v=0;v<numVertices;v++)
{
GLInstanceVertex vert;
vert.uv[0] = 0.5f;
vert.uv[1] = 0.5f;
vert.normal[0]=0.f;
vert.normal[1]=1.f;
vert.normal[2]=0.f;
b3Vector3 vertPos;
vertPos[0] = shape.mesh.positions[v*3+0]*myScaling[0];
vertPos[1] = shape.mesh.positions[v*3+1]*myScaling[1];
vertPos[2] = shape.mesh.positions[v*3+2]*myScaling[2];
vertPos[3] =0.f;
center+=vertPos;
}
center/=numVertices;
for (int v=0;v<numVertices;v++)
{
GLInstanceVertex vert;
vert.uv[0] = 0.5f;
vert.uv[1] = 0.5f;
vert.normal[0]=0.f;
vert.normal[1]=1.f;
vert.normal[2]=0.f;
b3Vector3 vertPos;
vertPos[0] = shape.mesh.positions[v*3+0]*myScaling[0];
vertPos[1] = shape.mesh.positions[v*3+1]*myScaling[1];
vertPos[2] = shape.mesh.positions[v*3+2]*myScaling[2];
vertPos[3] =0.f;
// vertPos-=center;
vert.xyzw[0] = vertPos[0];
vert.xyzw[1] = vertPos[1];
vert.xyzw[2] = vertPos[2];
tmpVertices.push_back(vert);
b3Vector3 newPos = tr*vertPos;
vert.xyzw[0] = newPos[0];
vert.xyzw[1] = newPos[1];
vert.xyzw[2] = newPos[2];
vert.xyzw[3] = 0.f;
vertexArray.push_back(vert);
}
int childColIndex = m_data->m_np->registerConvexHullShape(&tmpVertices[0].xyzw[0],strideInBytes,numVertices, scaling);
child.m_shapeIndex = childColIndex;
childShapes.push_back(child);
colIndex = childColIndex;
}
colIndex= m_data->m_np->registerCompoundShape(&childShapes);
}
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int shapeId = ci.m_instancingRenderer->registerShape(&vertexArray[0].xyzw[0],vertexArray.size(),&indexArray[0],indexArray.size());
b3Vector4 colors[4] =
{
b3MakeVector4(1,0,0,1),
b3MakeVector4(0,1,0,1),
b3MakeVector4(0,0,1,1),
b3MakeVector4(0,1,1,1),
};
int curColor = 0;
for (int i=0;i<1;i++)//ci.arraySizeX;i++)
{
for (int j=0;j<4;j++)
{
// for (int k=0;k<ci.arraySizeZ;k++)
int k=0;
{
float mass = 1;//j==0? 0.f : 1.f;
//b3Vector3 position(i*10*ci.gapX,j*ci.gapY,k*10*ci.gapZ);
b3Vector3 position=b3MakeVector3(i*10*ci.gapX,10+j*10*ci.gapY,k*10*ci.gapZ);
// b3Quaternion orn(0,0,0,1);
b3Quaternion orn(b3MakeVector3(0,0,1),1.8);
b3Vector4 color = colors[curColor];
curColor++;
curColor&=3;
b3Vector4 scaling=b3MakeVector4(1,1,1,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index,false);
index++;
}
}
}
}
}
void ConcaveCompoundScene::createDynamicObjects(const ConstructionInfo& ci)
{
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
b3AlignedObjectArray<GLInstanceVertex> vertexArray;
b3AlignedObjectArray<int> indexArray;
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int group=1;
int mask=1;
int index=0;
float scaling[4] = {1,1,1,1};
int colIndex = 0;
GLInstanceVertex* cubeVerts = (GLInstanceVertex*)&cube_vertices[0];
int stride2 = sizeof(GLInstanceVertex);
b3Assert(stride2 == strideInBytes);
{
int childColIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
b3Vector3 childPositions[3] = {
b3MakeVector3(0,-2,0),
b3MakeVector3(0,0,0),
b3MakeVector3(0,0,2)
};
b3AlignedObjectArray<b3GpuChildShape> childShapes;
int numChildShapes = 3;
for (int i=0;i<numChildShapes;i++)
{
//for now, only support polyhedral child shapes
b3GpuChildShape child;
child.m_shapeIndex = childColIndex;
b3Vector3 pos = childPositions[i];
b3Quaternion orn(0,0,0,1);
for (int v=0;v<4;v++)
{
child.m_childPosition[v] = pos[v];
child.m_childOrientation[v] = orn[v];
}
childShapes.push_back(child);
b3Transform tr;
tr.setIdentity();
tr.setOrigin(pos);
tr.setRotation(orn);
int baseIndex = vertexArray.size();
for (int j=0;j<numIndices;j++)
indexArray.push_back(cube_indices[j]+baseIndex);
//add transformed graphics vertices and indices
for (int v=0;v<numVertices;v++)
{
GLInstanceVertex vert = cubeVerts[v];
b3Vector3 vertPos=b3MakeVector3(vert.xyzw[0],vert.xyzw[1],vert.xyzw[2]);
b3Vector3 newPos = tr*vertPos;
vert.xyzw[0] = newPos[0];
vert.xyzw[1] = newPos[1];
vert.xyzw[2] = newPos[2];
vert.xyzw[3] = 0.f;
vertexArray.push_back(vert);
}
}
colIndex= m_data->m_np->registerCompoundShape(&childShapes);
}
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int shapeId = ci.m_instancingRenderer->registerShape(&vertexArray[0].xyzw[0],vertexArray.size(),&indexArray[0],indexArray.size());
b3Vector4 colors[4] =
{
b3MakeVector4(1,0,0,1),
b3MakeVector4(0,1,0,1),
b3MakeVector4(0,0,1,1),
b3MakeVector4(0,1,1,1),
};
int curColor = 0;
for (int i=0;i<ci.arraySizeX;i++)
{
for (int j=0;j<ci.arraySizeY;j++)
{
for (int k=0;k<ci.arraySizeZ;k++)
{
float mass = 1;//j==0? 0.f : 1.f;
b3Vector3 position=b3MakeVector3((-ci.arraySizeX/2+i)*ci.gapX,50+j*ci.gapY,(-ci.arraySizeZ/2+k)*ci.gapZ);
//b3Quaternion orn(0,0,0,1);
b3Quaternion orn(b3MakeVector3(1,0,0),0.7);
b3Vector4 color = colors[curColor];
curColor++;
curColor&=3;
b3Vector4 scaling=b3MakeVector4(1,1,1,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index,false);
index++;
}
}
}
}
void ConcaveSphereScene::setupScene(const ConstructionInfo& ci)
{
ConcaveScene::setupScene(ci);
float camPos[4]={0,50,0,0};//65.5,4.5,65.5,0};
//float camPos[4]={1,12.5,1.5,0};
m_instancingRenderer->setCameraPitch(45);
m_instancingRenderer->setCameraTargetPosition(camPos);
m_instancingRenderer->setCameraDistance(40);
}
void ConcaveSphereScene::createDynamicObjects(const ConstructionInfo& ci)
{
b3Vector4 colors[4] =
{
b3MakeVector4(1,0,0,1),
b3MakeVector4(0,1,0,1),
b3MakeVector4(0,1,1,1),
b3MakeVector4(1,1,0,1),
};
int index=0;
int curColor = 0;
float radius = 1;
//int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
int colIndex = m_data->m_np->registerSphereShape(radius);//>registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
int prevGraphicsShapeIndex = registerGraphicsSphereShape(ci,radius,false);
for (int i=0;i<ci.arraySizeX;i++)
{
for (int j=0;j<ci.arraySizeY;j++)
{
for (int k=0;k<ci.arraySizeZ;k++)
{
float mass = 1.f;
b3Vector3 position=b3MakeVector3(-(ci.arraySizeX/2)*8+i*8,50+j*8,-(ci.arraySizeZ/2)*8+k*8);
//b3Vector3 position(0,-41,0);//0,0,0);//i*radius*3,-41+j*radius*3,k*radius*3);
b3Quaternion orn(0,0,0,1);
b3Vector4 color = colors[curColor];
curColor++;
curColor&=3;
b3Vector4 scaling=b3MakeVector4(radius,radius,radius,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(prevGraphicsShapeIndex,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index,false);
index++;
}
}
}
}
+101
View File
@@ -0,0 +1,101 @@
#ifndef CONCAVE_SCENE_H
#define CONCAVE_SCENE_H
#include "GpuRigidBodyDemo.h"
#include "Bullet3Common/b3Vector3.h"
class ConcaveScene : public GpuRigidBodyDemo
{
public:
ConcaveScene(){}
virtual ~ConcaveScene(){}
virtual const char* getName()
{
return "BoxTrimesh";
}
static GpuDemo* MyCreateFunc()
{
GpuDemo* demo = new ConcaveScene;
return demo;
}
virtual void setupScene(const ConstructionInfo& ci);
virtual void createDynamicObjects(const ConstructionInfo& ci);
virtual void createConcaveMesh(const ConstructionInfo& ci, const char* fileName, const b3Vector3& shift, const b3Vector3& scaling);
};
class ConcaveSphereScene : public ConcaveScene
{
public:
ConcaveSphereScene(){}
virtual ~ConcaveSphereScene(){}
virtual const char* getName()
{
return "SphereTrimesh";
}
static GpuDemo* MyCreateFunc()
{
GpuDemo* demo = new ConcaveSphereScene;
return demo;
}
virtual void setupScene(const ConstructionInfo& ci);
virtual void createDynamicObjects(const ConstructionInfo& ci);
};
class ConcaveCompoundScene : public ConcaveScene
{
public:
ConcaveCompoundScene(){}
virtual ~ConcaveCompoundScene(){}
virtual const char* getName()
{
return "CompoundConcave";
}
static GpuDemo* MyCreateFunc()
{
GpuDemo* demo = new ConcaveCompoundScene;
return demo;
}
virtual void setupScene(const ConstructionInfo& ci);
virtual void createDynamicObjects(const ConstructionInfo& ci);
};
class ConcaveCompound2Scene : public ConcaveCompoundScene
{
public:
ConcaveCompound2Scene(){}
virtual ~ConcaveCompound2Scene(){}
virtual const char* getName()
{
return "GRBConcave2Compound";
}
static GpuDemo* MyCreateFunc()
{
GpuDemo* demo = new ConcaveCompound2Scene;
return demo;
}
virtual void createDynamicObjects(const ConstructionInfo& ci);
};
#endif //CONCAVE_SCENE_H
@@ -0,0 +1,273 @@
#include "GpuCompoundScene.h"
#include "GpuRigidBodyDemo.h"
#include "OpenGLWindow/ShapeData.h"
#include "OpenGLWindow/GLInstancingRenderer.h"
#include "Bullet3Common/b3Quaternion.h"
#include "OpenGLWindow/b3gWindowInterface.h"
#include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h"
#include "../GpuDemoInternalData.h"
#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
#include "OpenGLWindow/OpenGLInclude.h"
#include "OpenGLWindow/GLInstanceRendererInternalData.h"
#include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h"
#include "Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.h"
#include "Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.h"
#include "Bullet3Collision/NarrowPhaseCollision/b3Config.h"
#include "GpuRigidBodyDemoInternalData.h"
#include "Bullet3Common/b3Transform.h"
#include "OpenGLWindow/GLInstanceGraphicsShape.h"
#define NUM_COMPOUND_CHILDREN_X 4
#define NUM_COMPOUND_CHILDREN_Y 4
#define NUM_COMPOUND_CHILDREN_Z 4
void GpuCompoundScene::setupScene(const ConstructionInfo& ci)
{
createStaticEnvironment(ci);
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
float scaling[4] = {1,1,1,1};
GLInstanceVertex* cubeVerts = (GLInstanceVertex*)&cube_vertices[0];
int stride2 = sizeof(GLInstanceVertex);
b3Assert(stride2 == strideInBytes);
int index=0;
int colIndex = -1;
b3AlignedObjectArray<GLInstanceVertex> vertexArray;
b3AlignedObjectArray<int> indexArray;
{
int childColIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
/* b3Vector3 childPositions[3] = {
b3Vector3(0,-2,0),
b3Vector3(0,0,0),
b3Vector3(0,0,2)
};
*/
b3AlignedObjectArray<b3GpuChildShape> childShapes;
for (int x=0;x<NUM_COMPOUND_CHILDREN_X;x++)
for (int y=0;y<NUM_COMPOUND_CHILDREN_Y;y++)
for (int z=0;z<NUM_COMPOUND_CHILDREN_Z;z++)
{
int blax = x!=0 ?1 : 0;
int blay = y!=0 ?1 : 0;
int blaz = z!=0 ?1 : 0;
int bla=blax+blay+blaz;
if (bla!=1)
continue;
//for now, only support polyhedral child shapes
b3GpuChildShape child;
child.m_shapeIndex = childColIndex;
b3Vector3 pos=b3MakeVector3((x-NUM_COMPOUND_CHILDREN_X/2.f)*2,(y-NUM_COMPOUND_CHILDREN_X/2.f)*2,(z-NUM_COMPOUND_CHILDREN_X/2.f)*2);//childPositions[i];
b3Quaternion orn(0,0,0,1);
for (int v=0;v<4;v++)
{
child.m_childPosition[v] = pos[v];
child.m_childOrientation[v] = orn[v];
}
childShapes.push_back(child);
b3Transform tr;
tr.setIdentity();
tr.setOrigin(pos);
tr.setRotation(orn);
int baseIndex = vertexArray.size();
for (int j=0;j<numIndices;j++)
indexArray.push_back(cube_indices[j]+baseIndex);
//add transformed graphics vertices and indices
for (int v=0;v<numVertices;v++)
{
GLInstanceVertex vert = cubeVerts[v];
b3Vector3 vertPos=b3MakeVector3(vert.xyzw[0],vert.xyzw[1],vert.xyzw[2]);
b3Vector3 newPos = tr*vertPos;
vert.xyzw[0] = newPos[0];
vert.xyzw[1] = newPos[1];
vert.xyzw[2] = newPos[2];
vert.xyzw[3] = 0.f;
vertexArray.push_back(vert);
}
}
colIndex= m_data->m_np->registerCompoundShape(&childShapes);
}
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int shapeId = ci.m_instancingRenderer->registerShape(&vertexArray[0].xyzw[0],vertexArray.size(),&indexArray[0],indexArray.size());
b3Vector4 colors[4] =
{
b3MakeVector4(1,0,0,1),
b3MakeVector4(0,1,0,1),
b3MakeVector4(0,0,1,1),
b3MakeVector4(0,1,1,1),
};
int curColor = 0;
for (int i=0;i<ci.arraySizeX;i++)
{
for (int j=0;j<ci.arraySizeY;j++)
{
for (int k=0;k<ci.arraySizeZ;k++)
{
float mass = 1;//j==0? 0.f : 1.f;
b3Vector3 position=b3MakeVector3((i-ci.arraySizeX/2.)*ci.gapX,35+j*3*ci.gapY,(k-ci.arraySizeZ/2.f)*ci.gapZ);
//b3Quaternion orn(0,0,0,1);
b3Quaternion orn(b3MakeVector3(1,0,0),0.7);
b3Vector4 color = colors[curColor];
curColor++;
curColor&=3;
b3Vector4 scaling=b3MakeVector4(1,1,1,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index,false);
index++;
}
}
}
m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
float camPos[4]={0,0,0};//65.5,4.5,65.5,0};
//float camPos[4]={1,12.5,1.5,0};
m_instancingRenderer->setCameraTargetPosition(camPos);
m_instancingRenderer->setCameraDistance(320);
}
void GpuCompoundScene::createStaticEnvironment(const ConstructionInfo& ci)
{
int strideInBytes = 9*sizeof(float);
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int group=1;
int mask=1;
int index=0;
int colIndex = 0;
{
if (1)
{
float radius = 41;
int prevGraphicsShapeIndex = -1;
{
if (radius>=100)
{
int numVertices = sizeof(detailed_sphere_vertices)/strideInBytes;
int numIndices = sizeof(detailed_sphere_indices)/sizeof(int);
prevGraphicsShapeIndex = ci.m_instancingRenderer->registerShape(&detailed_sphere_vertices[0],numVertices,detailed_sphere_indices,numIndices);
} else
{
bool usePointSprites = false;
if (usePointSprites)
{
int numVertices = sizeof(point_sphere_vertices)/strideInBytes;
int numIndices = sizeof(point_sphere_indices)/sizeof(int);
prevGraphicsShapeIndex = ci.m_instancingRenderer->registerShape(&point_sphere_vertices[0],numVertices,point_sphere_indices,numIndices,B3_GL_POINTS);
} else
{
if (radius>=10)
{
int numVertices = sizeof(medium_sphere_vertices)/strideInBytes;
int numIndices = sizeof(medium_sphere_indices)/sizeof(int);
prevGraphicsShapeIndex = ci.m_instancingRenderer->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 = ci.m_instancingRenderer->registerShape(&low_sphere_vertices[0],numVertices,low_sphere_indices,numIndices);
}
}
}
}
b3Vector4 colors[4] =
{
b3MakeVector4(1,0,0,1),
b3MakeVector4(0,1,0,1),
b3MakeVector4(0,1,1,1),
b3MakeVector4(1,1,0,1),
};
int curColor = 1;
//int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
int colIndex = m_data->m_np->registerSphereShape(radius);//>registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
float mass = 0.f;
//b3Vector3 position((j&1)+i*2.2,1+j*2.,(j&1)+k*2.2);
b3Vector3 position=b3MakeVector3(0,-41,0);
b3Quaternion orn(0,0,0,1);
b3Vector4 color = colors[curColor];
curColor++;
curColor&=3;
b3Vector4 scaling=b3MakeVector4(radius,radius,radius,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(prevGraphicsShapeIndex,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index,false);
index++;
}
}
}
void GpuCompoundPlaneScene::createStaticEnvironment(const ConstructionInfo& ci)
{
int index=0;
b3Vector3 normal=b3MakeVector3(0,1,0);
float constant=0.f;
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
b3Vector4 scaling=b3MakeVector4(400,1.,400,1);
//int colIndex = m_data->m_np->registerPlaneShape(normal,constant);//>registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
b3Vector3 position=b3MakeVector3(0,0,0);
b3Quaternion orn(0,0,0,1);
// b3Quaternion orn(b3Vector3(1,0,0),0.3);
b3Vector4 color=b3MakeVector4(0,0,1,1);
int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0.f,position,orn,colIndex,index,false);
}
@@ -0,0 +1,50 @@
#ifndef GPU_COMPOUND_SCENE_H
#define GPU_COMPOUND_SCENE_H
#include "GpuRigidBodyDemo.h"
class GpuCompoundScene : public GpuRigidBodyDemo
{
public:
GpuCompoundScene(){}
virtual ~GpuCompoundScene(){}
virtual const char* getName()
{
return "CompoundOnSphere";
}
static GpuDemo* MyCreateFunc()
{
GpuDemo* demo = new GpuCompoundScene;
return demo;
}
virtual void setupScene(const ConstructionInfo& ci);
virtual void createStaticEnvironment(const ConstructionInfo& ci);
};
class GpuCompoundPlaneScene : public GpuCompoundScene
{
public:
GpuCompoundPlaneScene(){}
virtual ~GpuCompoundPlaneScene(){}
virtual const char* getName()
{
return "CompoundOnPlane";
}
static GpuDemo* MyCreateFunc()
{
GpuDemo* demo = new GpuCompoundPlaneScene;
return demo;
}
virtual void createStaticEnvironment(const ConstructionInfo& ci);
};
#endif //GPU_COMPOUND_SCENE_H
@@ -0,0 +1,651 @@
#include "GpuConvexScene.h"
#include "GpuRigidBodyDemo.h"
#include "../OpenGLWindow/ShapeData.h"
#include "../OpenGLWindow/GLInstancingRenderer.h"
#include "Bullet3Common/b3Quaternion.h"
#include "../CommonInterfaces/CommonWindowInterface.h"
#include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h"
#include "../CommonOpenCL/GpuDemoInternalData.h"
#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
#include "../OpenGLWindow/OpenGLInclude.h"
#include "../OpenGLWindow/GLInstanceRendererInternalData.h"
#include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h"
#include "Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.h"
#include "Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.h"
#include "Bullet3Collision/NarrowPhaseCollision/b3Config.h"
#include "GpuRigidBodyDemoInternalData.h"
#include "Bullet3Dynamics/ConstraintSolver/b3Point2PointConstraint.h"
#include "../OpenGLWindow/GLPrimitiveRenderer.h"
#include "Bullet3OpenCL/Raycast/b3GpuRaycast.h"
#include "Bullet3Collision/NarrowPhaseCollision/b3ConvexUtility.h"
#include "Bullet3Dynamics/ConstraintSolver/b3FixedConstraint.h"
#include "../OpenGLWindow/GLRenderToTexture.h"
static bool gUseInstancedCollisionShapes = true;
extern int gGpuArraySizeX;
extern int gGpuArraySizeY;
extern int gGpuArraySizeZ;
#include "GpuRigidBodyDemo.h"
#include "Bullet3Common/b3AlignedObjectArray.h"
#include "Bullet3Collision/NarrowPhaseCollision/b3RaycastInfo.h"
class GpuConvexScene : public GpuRigidBodyDemo
{
protected:
class b3GpuRaycast* m_raycaster;
public:
GpuConvexScene(GUIHelperInterface* helper)
:GpuRigidBodyDemo(helper), m_raycaster(0)
{
}
virtual ~GpuConvexScene(){}
virtual const char* getName()
{
return "Tetrahedra";
}
virtual void setupScene();
virtual void destroyScene();
virtual int createDynamicsObjects();
virtual int createDynamicsObjects2(const float* vertices, int numVertices, const int* indices,int numIndices);
virtual void createStaticEnvironment();
};
class GpuConvexPlaneScene : public GpuConvexScene
{
public:
GpuConvexPlaneScene(GUIHelperInterface* helper)
:GpuConvexScene(helper){}
virtual ~GpuConvexPlaneScene(){}
virtual const char* getName()
{
return "ConvexOnPlane";
}
virtual void createStaticEnvironment();
};
class GpuBoxPlaneScene : public GpuConvexPlaneScene
{
public:
GpuBoxPlaneScene(GUIHelperInterface* helper):GpuConvexPlaneScene(helper){}
virtual ~GpuBoxPlaneScene(){}
virtual const char* getName()
{
return "BoxBox";
}
virtual int createDynamicsObjects();
};
class GpuTetraScene : public GpuConvexScene
{
protected:
void createFromTetGenData(const char* ele, const char* node);
public:
virtual const char* getName()
{
return "TetraBreakable";
}
virtual int createDynamicsObjects();
};
b3Vector4 colors[4] =
{
b3MakeVector4(1,0,0,1),
b3MakeVector4(0,1,0,1),
b3MakeVector4(0,1,1,1),
b3MakeVector4(1,1,0,1),
};
void GpuConvexScene::setupScene()
{
m_raycaster = new b3GpuRaycast(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue);
int index=0;
createStaticEnvironment();
index+=createDynamicsObjects();
m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
float camPos[4]={0,0,0,0};//ci.arraySizeX,ci.arraySizeY/2,ci.arraySizeZ,0};
//float camPos[4]={1,12.5,1.5,0};
m_instancingRenderer->getActiveCamera()->setCameraTargetPosition(camPos[0],camPos[1],camPos[2]);
m_instancingRenderer->getActiveCamera()->setCameraDistance(150);
//m_instancingRenderer->setCameraYaw(85);
m_instancingRenderer->getActiveCamera()->setCameraYaw(30);
m_instancingRenderer->getActiveCamera()->setCameraPitch(225);
m_instancingRenderer->updateCamera();
char msg[1024];
int numInstances = index;
sprintf(msg,"Num objects = %d",numInstances);
b3Printf(msg);
//if (ci.m_gui)
// ci.m_gui->setStatusBarMessage(msg,true);
}
void GpuConvexScene::destroyScene()
{
delete m_raycaster;
m_raycaster = 0;
}
int GpuConvexScene::createDynamicsObjects()
{
int strideInBytes = 9*sizeof(float);
/*int numVertices = sizeof(barrel_vertices)/strideInBytes;
int numIndices = sizeof(barrel_indices)/sizeof(int);
return createDynamicsObjects2(ci,barrel_vertices,numVertices,barrel_indices,numIndices);
*/
int numVertices = sizeof(tetra_vertices)/strideInBytes;
int numIndices = sizeof(tetra_indices)/sizeof(int);
return createDynamicsObjects2(tetra_vertices,numVertices,tetra_indices,numIndices);
}
int GpuBoxPlaneScene::createDynamicsObjects()
{
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
return createDynamicsObjects2(cube_vertices_textured,numVertices,cube_indices,numIndices);
}
int GpuConvexScene::createDynamicsObjects2( const float* vertices, int numVertices, const int* indices, int numIndices)
{
int strideInBytes = 9*sizeof(float);
int textureIndex = -1;
if (0)
{
int width,height,n;
const char* filename = "data/cube.png";
const unsigned char* image=0;
const char* prefix[]={"./","../","../../","../../../","../../../../"};
int numprefix = sizeof(prefix)/sizeof(const char*);
for (int i=0;!image && i<numprefix;i++)
{
char relativeFileName[1024];
sprintf(relativeFileName,"%s%s",prefix[i],filename);
image = loadImage(relativeFileName,width,height,n);
}
b3Assert(image);
if (image)
{
textureIndex = m_instancingRenderer->registerTexture(image,width,height);
}
}
int shapeId = m_instancingRenderer->registerShape(&vertices[0],numVertices,indices,numIndices,B3_GL_TRIANGLES,textureIndex);
int group=1;
int mask=1;
int index=0;
{
int curColor = 0;
float scaling[4] = {1,1,1,1};
int prevBody = -1;
int insta = 0;
b3ConvexUtility* utilPtr = new b3ConvexUtility();
{
b3AlignedObjectArray<b3Vector3> verts;
unsigned char* vts = (unsigned char*) vertices;
for (int i=0;i<numVertices;i++)
{
float* vertex = (float*) &vts[i*strideInBytes];
verts.push_back(b3MakeVector3(vertex[0]*scaling[0],vertex[1]*scaling[1],vertex[2]*scaling[2]));
}
bool merge = true;
if (numVertices)
{
utilPtr->initializePolyhedralFeatures(&verts[0],verts.size(),merge);
}
}
int colIndex=-1;
if (gUseInstancedCollisionShapes)
colIndex = m_data->m_np->registerConvexHullShape(utilPtr);
//int colIndex = m_data->m_np->registerSphereShape(1);
for (int i=0;i<gGpuArraySizeX;i++)
{
//printf("%d of %d\n", i, ci.arraySizeX);
for (int j=0;j<gGpuArraySizeY;j++)
{
for (int k=0;k<gGpuArraySizeZ;k++)
{
//int colIndex = m_data->m_np->registerConvexHullShape(&vertices[0],strideInBytes,numVertices, scaling);
if (!gUseInstancedCollisionShapes)
colIndex = m_data->m_np->registerConvexHullShape(utilPtr);
float mass = 1.f;
if (j==0)//ci.arraySizeY-1)
{
//mass=0.f;
}
b3Vector3 position = b3MakeVector3(((j+1)&1)+i*2.2,1+j*2.,((j+1)&1)+k*2.2);
//b3Vector3 position = b3MakeVector3(i*2,1+j*2,k*2);
//b3Vector3 position=b3MakeVector3(1,0.9,1);
b3Quaternion orn(0,0,0,1);
b3Vector4 color = colors[curColor];
curColor++;
curColor&=3;
b3Vector4 scalin=b3MakeVector4(1,1,1,1);
int id = m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index,false);
if (prevBody>=0)
{
//b3Point2PointConstraint* p2p = new b3Point2PointConstraint(pid,prevBody,b3Vector3(0,-1.1,0),b3Vector3(0,1.1,0));
// m_data->m_rigidBodyPipeline->addConstraint(p2p);//,false);
}
prevBody = pid;
index++;
}
}
}
delete utilPtr;
}
return index;
}
void GpuConvexScene::createStaticEnvironment()
{
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int shapeId = m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int group=1;
int mask=1;
int index=0;
{
b3Vector4 scaling=b3MakeVector4(400,400,400,1);
int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
b3Vector3 position=b3MakeVector3(0,-400,0);
b3Quaternion orn(0,0,0,1);
b3Vector4 color=b3MakeVector4(0,0,1,1);
int id = m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0.f,position,orn,colIndex,index,false);
}
}
void GpuConvexPlaneScene::createStaticEnvironment()
{
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int shapeId = m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int group=1;
int mask=1;
int index=0;
{
b3Vector4 scaling=b3MakeVector4(400,400,400,1);
int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
b3Vector3 position=b3MakeVector3(0,-400,0);
b3Quaternion orn(0,0,0,1);
b3Vector4 color=b3MakeVector4(0,0,1,1);
int id = m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0.f,position,orn,colIndex,index,false);
}
}
/*
void GpuConvexPlaneScene::createStaticEnvironment(const ConstructionInfo& ci)
{
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int group=1;
int mask=1;
int index=0;
{
b3Vector4 scaling=b3MakeVector4(100,0.001,100,1);
//int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
b3Vector3 normal=b3MakeVector3(0,1,0);
float constant=0.f;
int colIndex = m_data->m_np->registerPlaneShape(normal,constant);//>registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
b3Vector3 position=b3MakeVector3(0,0,0);
b3Quaternion orn(0,0,0,1);
b3Vector4 color=b3MakeVector4(0,0,1,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0.f,position,orn,colIndex,index,false);
}
}
*/
struct TetraBunny
{
#include "bunny.inl"
};
struct TetraCube
{
#include "cube.inl"
};
static int nextLine(const char* buffer)
{
int numBytesRead=0;
while (*buffer != '\n')
{
buffer++;
numBytesRead++;
}
if (buffer[0]==0x0a)
{
buffer++;
numBytesRead++;
}
return numBytesRead;
}
static float mytetra_vertices[] =
{
-1.f, 0, -1.f, 0.5f, 0, 1,0, 0,0,
-1.f, 0, 1.f, 0.5f, 0, 1,0, 1,0,
1.f, 0, 1.f, 0.5f, 0, 1,0, 1,1,
1.f, 0, -1.f, 0.5f, 0, 1,0, 0,1
};
static int mytetra_indices[]=
{
0,1,2,
3,1,2,3,2,0,
3,0,1
};
/* Create from TetGen .ele, .face, .node data */
void GpuTetraScene::createFromTetGenData(const char* ele,
const char* node)
{
b3Scalar scaling(10);
b3AlignedObjectArray<b3Vector3> pos;
int nnode=0;
int ndims=0;
int nattrb=0;
int hasbounds=0;
int result = sscanf(node,"%d %d %d %d",&nnode,&ndims,&nattrb,&hasbounds);
result = sscanf(node,"%d %d %d %d",&nnode,&ndims,&nattrb,&hasbounds);
node += nextLine(node);
//b3AlignedObjectArray<b3Vector3> rigidBodyPositions;
//b3AlignedObjectArray<int> rigidBodyIds;
pos.resize(nnode);
for(int i=0;i<pos.size();++i)
{
int index=0;
//int bound=0;
float x,y,z;
sscanf(node,"%d %f %f %f",&index,&x,&y,&z);
// sn>>index;
// sn>>x;sn>>y;sn>>z;
node += nextLine(node);
//for(int j=0;j<nattrb;++j)
// sn>>a;
//if(hasbounds)
// sn>>bound;
pos[index].setX(b3Scalar(x)*scaling);
pos[index].setY(b3Scalar(y)*scaling);
pos[index].setZ(b3Scalar(z)*scaling);
}
if(ele&&ele[0])
{
int ntetra=0;
int ncorner=0;
int neattrb=0;
sscanf(ele,"%d %d %d",&ntetra,&ncorner,&neattrb);
ele += nextLine(ele);
//se>>ntetra;se>>ncorner;se>>neattrb;
for(int i=0;i<ntetra;++i)
{
int index=0;
int ni[4];
//se>>index;
//se>>ni[0];se>>ni[1];se>>ni[2];se>>ni[3];
sscanf(ele,"%d %d %d %d %d",&index,&ni[0],&ni[1],&ni[2],&ni[3]);
ele+=nextLine(ele);
b3Vector3 average=b3MakeVector3(0,0,0);
for (int v=0;v<4;v++)
{
average+=pos[ni[v]];
}
average/=4;
for (int v=0;v<4;v++)
{
b3Vector3 shiftedPos = pos[ni[v]]-average;
mytetra_vertices[0+v*9] = shiftedPos.getX();
mytetra_vertices[1+v*9] = shiftedPos.getY();
mytetra_vertices[2+v*9] = shiftedPos.getZ();
}
//todo: subtract average
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(mytetra_vertices)/strideInBytes;
int numIndices = sizeof(mytetra_indices)/sizeof(int);
int shapeId = m_instancingRenderer->registerShape(&mytetra_vertices[0],numVertices,mytetra_indices,numIndices);
int group=1;
int mask=1;
{
b3Vector4 scaling=b3MakeVector4(1,1,1,1);
int colIndex = m_data->m_np->registerConvexHullShape(&mytetra_vertices[0],strideInBytes,numVertices, scaling);
b3Vector3 position=b3MakeVector3(0,150,0);
// position+=average;//*1.2;//*2;
position+=average*1.2;//*2;
//rigidBodyPositions.push_back(position);
b3Quaternion orn(0,0,0,1);
static int curColor=0;
b3Vector4 color = colors[curColor++];
curColor&=3;
int id = m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(1.f,position,orn,colIndex,0,false);
//rigidBodyIds.push_back(pid);
}
//for(int j=0;j<neattrb;++j)
// se>>a;
//psb->appendTetra(ni[0],ni[1],ni[2],ni[3]);
}
// printf("Nodes: %u\r\n",psb->m_nodes.size());
// printf("Links: %u\r\n",psb->m_links.size());
// printf("Faces: %u\r\n",psb->m_faces.size());
// printf("Tetras: %u\r\n",psb->m_tetras.size());
}
m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
m_data->m_np->writeAllBodiesToGpu();
m_data->m_bp->writeAabbsToGpu();
m_data->m_rigidBodyPipeline->setupGpuAabbsFull();
m_data->m_bp->calculateOverlappingPairs(m_data->m_config.m_maxBroadphasePairs);
int numPairs = m_data->m_bp->getNumOverlap();
cl_mem pairs = m_data->m_bp->getOverlappingPairBuffer();
b3OpenCLArray<b3Int2> clPairs(m_clData->m_clContext,m_clData->m_clQueue);
clPairs.setFromOpenCLBuffer(pairs,numPairs);
b3AlignedObjectArray<b3Int2> allPairs;
clPairs.copyToHost(allPairs);
for (int p=0;p<allPairs.size();p++)
{
b3Vector3 posA,posB;
b3Quaternion ornA,ornB;
int bodyIndexA = allPairs[p].x;
int bodyIndexB = allPairs[p].y;
m_data->m_np->getObjectTransformFromCpu(posA,ornA,bodyIndexA);
m_data->m_np->getObjectTransformFromCpu(posB,ornB,bodyIndexB);
b3Vector3 pivotWorld = (posA+posB)*0.5f;
b3Transform transA,transB;
transA.setIdentity();
transA.setOrigin(posA);
transA.setRotation(ornA);
transB.setIdentity();
transB.setOrigin(posB);
transB.setRotation(ornB);
b3Vector3 pivotInA = transA.inverse()*pivotWorld;
b3Vector3 pivotInB = transB.inverse()*pivotWorld;
b3Transform frameInA,frameInB;
frameInA.setIdentity();
frameInB.setIdentity();
frameInA.setOrigin(pivotInA);
frameInB.setOrigin(pivotInB);
b3Quaternion relTargetAB = frameInA.getRotation()*frameInB.getRotation().inverse();
//c = new b3FixedConstraint(pid,prevBody,frameInA,frameInB);
float breakingThreshold = 45;//37.f;
//c->setBreakingImpulseThreshold(37.1);
bool useGPU = true;
if (useGPU)
{
int cid = m_data->m_rigidBodyPipeline->createFixedConstraint(bodyIndexA,bodyIndexB,pivotInA,pivotInB,relTargetAB,breakingThreshold);
} else
{
b3FixedConstraint* c = new b3FixedConstraint(bodyIndexA,bodyIndexB,frameInA,frameInB);
c->setBreakingImpulseThreshold(breakingThreshold);
m_data->m_rigidBodyPipeline->addConstraint(c);
}
}
printf("numPairs = %d\n",numPairs);
}
int GpuTetraScene::createDynamicsObjects()
{
//createFromTetGenData(TetraCube::getElements(),TetraCube::getNodes());
createFromTetGenData(TetraBunny::getElements(),TetraBunny::getNodes());
return 0;
}
class CommonExampleInterface* OpenCLBoxBoxCreateFunc(struct CommonExampleOptions& options)
{
return new GpuBoxPlaneScene(options.m_guiHelper);
}
@@ -0,0 +1,7 @@
#ifndef GPU_CONVEX_SCENE_H
#define GPU_CONVEX_SCENE_H
class CommonExampleInterface* OpenCLBoxBoxCreateFunc(struct CommonExampleOptions& options);
#endif //GPU_CONVEX_SCENE_H
@@ -0,0 +1,495 @@
#include "GpuRigidBodyDemo.h"
#include "../OpenGLWindow/ShapeData.h"
#include "../OpenGLWindow/GLInstancingRenderer.h"
#include "Bullet3Common/b3Quaternion.h"
#include "../CommonInterfaces/CommonWindowInterface.h"
#include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h"
#include "Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h"
#include "../CommonOpenCL/GpuDemoInternalData.h"
#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
#include "../OpenGLWindow/OpenGLInclude.h"
#include "../OpenGLWindow/GLInstanceRendererInternalData.h"
#include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h"
#include "Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.h"
#include "Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.h"
#include "Bullet3Collision/NarrowPhaseCollision/b3Config.h"
#include "GpuRigidBodyDemoInternalData.h"
#include "Bullet3Collision/BroadPhaseCollision/b3DynamicBvhBroadphase.h"
#include "Bullet3Collision/NarrowPhaseCollision/shared/b3RigidBodyData.h"
#include "Bullet3OpenCL/RigidBody/b3GpuNarrowPhaseInternalData.h"
#include "stb_image/stb_image.h"
#include "../OpenGLWindow/GLPrimitiveRenderer.h"
extern int gPreferredOpenCLDeviceIndex;
extern int gPreferredOpenCLPlatformIndex;
extern int gGpuArraySizeX;
extern int gGpuArraySizeY;
extern int gGpuArraySizeZ;
static b3KeyboardCallback oldCallback = 0;
extern bool gReset;
bool useUniformGrid = false;
bool convertOnCpu = false;
#define MSTRINGIFY(A) #A
static const char* s_rigidBodyKernelString = MSTRINGIFY(
typedef struct
{
float4 m_pos;
float4 m_quat;
float4 m_linVel;
float4 m_angVel;
unsigned int m_collidableIdx;
float m_invMass;
float m_restituitionCoeff;
float m_frictionCoeff;
} Body;
__kernel void
copyTransformsToVBOKernel( __global Body* gBodies, __global float4* posOrnColor, const int numNodes)
{
int nodeID = get_global_id(0);
if( nodeID < numNodes )
{
posOrnColor[nodeID] = (float4) (gBodies[nodeID].m_pos.xyz,1.0);
posOrnColor[nodeID + numNodes] = gBodies[nodeID].m_quat;
}
}
);
GpuRigidBodyDemo::GpuRigidBodyDemo(GUIHelperInterface* helper)
:CommonOpenCLBase(helper),
m_instancingRenderer(0),
m_window(0)
{
m_instancingRenderer = (GLInstancingRenderer*)helper->getRenderInterface();
m_window = helper->getAppInterface()->m_window;
m_data = new GpuRigidBodyDemoInternalData;
}
GpuRigidBodyDemo::~GpuRigidBodyDemo()
{
delete m_data;
}
static void PairKeyboardCallback(int key, int state)
{
if (key=='R' && state)
{
gReset = true;
}
//b3DefaultKeyboardCallback(key,state);
oldCallback(key,state);
}
void GpuRigidBodyDemo::setupScene()
{
}
void GpuRigidBodyDemo::initPhysics()
{
initCL(gPreferredOpenCLDeviceIndex,gPreferredOpenCLPlatformIndex);
if (m_clData->m_clContext)
{
int errNum=0;
cl_program rbProg=0;
m_data->m_copyTransformsToVBOKernel = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_rigidBodyKernelString,"copyTransformsToVBOKernel",&errNum,rbProg);
m_data->m_config.m_maxConvexBodies = b3Max(m_data->m_config.m_maxConvexBodies,gGpuArraySizeX*gGpuArraySizeY*gGpuArraySizeZ+10);
m_data->m_config.m_maxConvexShapes = m_data->m_config.m_maxConvexBodies;
int maxPairsPerBody = 16;
m_data->m_config.m_maxBroadphasePairs = maxPairsPerBody*m_data->m_config.m_maxConvexBodies;
m_data->m_config.m_maxContactCapacity = m_data->m_config.m_maxBroadphasePairs;
b3GpuNarrowPhase* np = new b3GpuNarrowPhase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue,m_data->m_config);
b3GpuBroadphaseInterface* bp =0;
if (useUniformGrid)
{
bp = new b3GpuGridBroadphase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue);
} else
{
bp = new b3GpuSapBroadphase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue);
}
m_data->m_np = np;
m_data->m_bp = bp;
m_data->m_broadphaseDbvt = new b3DynamicBvhBroadphase(m_data->m_config.m_maxConvexBodies);
m_data->m_rigidBodyPipeline = new b3GpuRigidBodyPipeline(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue, np, bp,m_data->m_broadphaseDbvt,m_data->m_config);
setupScene();
m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
np->writeAllBodiesToGpu();
bp->writeAabbsToGpu();
}
m_instancingRenderer->writeTransforms();
}
void GpuRigidBodyDemo::exitPhysics()
{
destroyScene();
delete m_data->m_instancePosOrnColor;
delete m_data->m_rigidBodyPipeline;
delete m_data->m_broadphaseDbvt;
delete m_data->m_np;
m_data->m_np = 0;
delete m_data->m_bp;
m_data->m_bp = 0;
exitCL();
}
void GpuRigidBodyDemo::renderScene()
{
m_instancingRenderer->renderScene();
}
void GpuRigidBodyDemo::stepSimulation(float deltaTime)
{
bool animate=true;
int numObjects= m_data->m_rigidBodyPipeline->getNumBodies();
//printf("numObjects=%d\n",numObjects);
if (numObjects > m_instancingRenderer->getInstanceCapacity())
{
static bool once = true;
if (once)
{
once=false;
b3Assert(0);
b3Error("m_instancingRenderer out-of-memory\n");
}
numObjects = m_instancingRenderer->getInstanceCapacity();
}
GLint err = glGetError();
assert(err==GL_NO_ERROR);
b3Vector4* positions = 0;
if (animate && numObjects)
{
B3_PROFILE("gl2cl");
if (!m_data->m_instancePosOrnColor)
{
GLuint vbo = m_instancingRenderer->getInternalData()->m_vbo;
int arraySizeInBytes = numObjects * (3)*sizeof(b3Vector4);
glBindBuffer(GL_ARRAY_BUFFER, vbo);
cl_bool blocking= CL_TRUE;
positions= (b3Vector4*)glMapBufferRange( GL_ARRAY_BUFFER,m_instancingRenderer->getMaxShapeCapacity(),arraySizeInBytes, GL_MAP_READ_BIT );//GL_READ_WRITE);//GL_WRITE_ONLY
GLint err = glGetError();
assert(err==GL_NO_ERROR);
m_data->m_instancePosOrnColor = new b3OpenCLArray<b3Vector4>(m_clData->m_clContext,m_clData->m_clQueue);
m_data->m_instancePosOrnColor->resize(3*numObjects);
m_data->m_instancePosOrnColor->copyFromHostPointer(positions,3*numObjects,0);
glUnmapBuffer( GL_ARRAY_BUFFER);
err = glGetError();
assert(err==GL_NO_ERROR);
}
}
{
B3_PROFILE("stepSimulation");
m_data->m_rigidBodyPipeline->stepSimulation(1./60.f);
}
if (numObjects)
{
if (convertOnCpu)
{
b3GpuNarrowPhaseInternalData* npData = m_data->m_np->getInternalData();
npData->m_bodyBufferGPU->copyToHost(*npData->m_bodyBufferCPU);
b3AlignedObjectArray<b3Vector4> vboCPU;
m_data->m_instancePosOrnColor->copyToHost(vboCPU);
for (int i=0;i<numObjects;i++)
{
b3Vector4 pos = (const b3Vector4&)npData->m_bodyBufferCPU->at(i).m_pos;
b3Quat orn = npData->m_bodyBufferCPU->at(i).m_quat;
pos.w = 1.f;
vboCPU[i] = pos;
vboCPU[i + numObjects] = (b3Vector4&)orn;
}
m_data->m_instancePosOrnColor->copyFromHost(vboCPU);
} else
{
B3_PROFILE("cl2gl_convert");
int ciErrNum = 0;
cl_mem bodies = m_data->m_rigidBodyPipeline->getBodyBuffer();
b3LauncherCL launch(m_clData->m_clQueue,m_data->m_copyTransformsToVBOKernel,"m_copyTransformsToVBOKernel");
launch.setBuffer(bodies);
launch.setBuffer(m_data->m_instancePosOrnColor->getBufferCL());
launch.setConst(numObjects);
launch.launch1D(numObjects);
oclCHECKERROR(ciErrNum, CL_SUCCESS);
}
}
if (animate && numObjects)
{
B3_PROFILE("cl2gl_upload");
GLint err = glGetError();
assert(err==GL_NO_ERROR);
GLuint vbo = m_instancingRenderer->getInternalData()->m_vbo;
int arraySizeInBytes = numObjects * (3)*sizeof(b3Vector4);
glBindBuffer(GL_ARRAY_BUFFER, vbo);
cl_bool blocking= CL_TRUE;
positions= (b3Vector4*)glMapBufferRange( GL_ARRAY_BUFFER,m_instancingRenderer->getMaxShapeCapacity(),arraySizeInBytes, GL_MAP_WRITE_BIT );//GL_READ_WRITE);//GL_WRITE_ONLY
err = glGetError();
assert(err==GL_NO_ERROR);
m_data->m_instancePosOrnColor->copyToHostPointer(positions,3*numObjects,0);
glUnmapBuffer( GL_ARRAY_BUFFER);
err = glGetError();
assert(err==GL_NO_ERROR);
}
}
b3Vector3 GpuRigidBodyDemo::getRayTo(int x,int y)
{
if (!m_instancingRenderer)
return b3MakeVector3(0,0,0);
float top = 1.f;
float bottom = -1.f;
float nearPlane = 1.f;
float tanFov = (top-bottom)*0.5f / nearPlane;
float fov = b3Scalar(2.0) * b3Atan(tanFov);
b3Vector3 camPos,camTarget;
m_instancingRenderer->getActiveCamera()->getCameraPosition(camPos);
m_instancingRenderer->getActiveCamera()->getCameraTargetPosition(camTarget);
b3Vector3 rayFrom = camPos;
b3Vector3 rayForward = (camTarget-camPos);
rayForward.normalize();
float farPlane = 10000.f;
rayForward*= farPlane;
b3Vector3 rightOffset;
b3Vector3 m_cameraUp=b3MakeVector3(0,1,0);
b3Vector3 vertical = m_cameraUp;
b3Vector3 hor;
hor = rayForward.cross(vertical);
hor.normalize();
vertical = hor.cross(rayForward);
vertical.normalize();
float tanfov = tanf(0.5f*fov);
hor *= 2.f * farPlane * tanfov;
vertical *= 2.f * farPlane * tanfov;
b3Scalar aspect;
float width = m_instancingRenderer->getScreenWidth();
float height = m_instancingRenderer->getScreenHeight();
aspect = width / height;
hor*=aspect;
b3Vector3 rayToCenter = rayFrom + rayForward;
b3Vector3 dHor = hor * 1.f/width;
b3Vector3 dVert = vertical * 1.f/height;
b3Vector3 rayTo = rayToCenter - 0.5f * hor + 0.5f * vertical;
rayTo += b3Scalar(x) * dHor;
rayTo -= b3Scalar(y) * dVert;
return rayTo;
}
unsigned char* GpuRigidBodyDemo::loadImage(const char* fileName, int& width, int& height, int& n)
{
unsigned char *data = stbi_load(fileName, &width, &height, &n, 0);
return data;
}
bool GpuRigidBodyDemo::keyboardCallback(int key, int state)
{
if (m_data)
{
if (key==B3G_ALT )
{
m_data->m_altPressed = state;
}
if (key==B3G_CONTROL )
{
m_data->m_controlPressed = state;
}
}
return false;
}
bool GpuRigidBodyDemo::mouseMoveCallback(float x,float y)
{
if (m_data->m_altPressed!=0 || m_data->m_controlPressed!=0)
return false;
if (m_data->m_pickBody>=0 && m_data->m_pickConstraint>=0)
{
m_data->m_rigidBodyPipeline->removeConstraintByUid(m_data->m_pickConstraint);
b3Vector3 newRayTo = getRayTo(x,y);
b3Vector3 rayFrom;
b3Vector3 oldPivotInB = m_data->m_pickPivotInB;
b3Vector3 newPivotB;
m_instancingRenderer->getActiveCamera()->getCameraPosition(rayFrom);
b3Vector3 dir = newRayTo-rayFrom;
dir.normalize();
dir *= m_data->m_pickDistance;
newPivotB = rayFrom + dir;
m_data->m_pickPivotInB = newPivotB;
m_data->m_rigidBodyPipeline->copyConstraintsToHost();
m_data->m_pickConstraint = m_data->m_rigidBodyPipeline->createPoint2PointConstraint(m_data->m_pickBody,m_data->m_pickFixedBody,m_data->m_pickPivotInA,m_data->m_pickPivotInB,1e30);
m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
return true;
}
return false;
}
bool GpuRigidBodyDemo::mouseButtonCallback(int button, int state, float x, float y)
{
if (state==1)
{
if(button==0 && (m_data->m_altPressed==0 && m_data->m_controlPressed==0))
{
b3AlignedObjectArray<b3RayInfo> rays;
b3AlignedObjectArray<b3RayHit> hitResults;
b3Vector3 camPos;
m_instancingRenderer->getActiveCamera()->getCameraPosition(camPos);
b3RayInfo ray;
ray.m_from = camPos;
ray.m_to = getRayTo(x,y);
rays.push_back(ray);
b3RayHit hit;
hit.m_hitFraction = 1.f;
hitResults.push_back(hit);
m_data->m_rigidBodyPipeline->castRays(rays,hitResults);
if (hitResults[0].m_hitFraction<1.f)
{
int hitBodyA = hitResults[0].m_hitBody;
if (m_data->m_np->getBodiesCpu()[hitBodyA].m_invMass)
{
//printf("hit!\n");
m_data->m_np->readbackAllBodiesToCpu();
m_data->m_pickBody = hitBodyA;
//pivotInA
b3Vector3 pivotInB;
pivotInB.setInterpolate3(ray.m_from,ray.m_to,hitResults[0].m_hitFraction);
b3Vector3 posA;
b3Quaternion ornA;
m_data->m_np->getObjectTransformFromCpu(posA,ornA,hitBodyA);
b3Transform tr;
tr.setOrigin(posA);
tr.setRotation(ornA);
b3Vector3 pivotInA = tr.inverse()*pivotInB;
if (m_data->m_pickFixedBody<0)
{
b3Vector3 pos=b3MakeVector3(0,0,0);
b3Quaternion orn(0,0,0,1);
int fixedSphere = m_data->m_np->registerConvexHullShape(0,0,0,0);//>registerSphereShape(0.1);
m_data->m_pickFixedBody = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0,pos,orn,fixedSphere,0,false);
m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
m_data->m_bp->writeAabbsToGpu();
if (m_data->m_pickGraphicsShapeIndex<0)
{
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(point_sphere_vertices)/strideInBytes;
int numIndices = sizeof(point_sphere_indices)/sizeof(int);
m_data->m_pickGraphicsShapeIndex = m_instancingRenderer->registerShape(&point_sphere_vertices[0],numVertices,point_sphere_indices,numIndices,B3_GL_POINTS);
float color[4] ={1,0,0,1};
float scaling[4]={1,1,1,1};
m_data->m_pickGraphicsShapeInstance = m_instancingRenderer->registerGraphicsInstance(m_data->m_pickGraphicsShapeIndex,pivotInB,orn,color,scaling);
m_instancingRenderer->writeTransforms();
delete m_data->m_instancePosOrnColor;
m_data->m_instancePosOrnColor=0;
} else
{
m_instancingRenderer->writeSingleInstanceTransformToCPU(pivotInB,orn,m_data->m_pickGraphicsShapeInstance);
m_instancingRenderer->writeSingleInstanceTransformToGPU(pivotInB,orn,m_data->m_pickGraphicsShapeInstance);
m_data->m_np->setObjectTransformCpu(pos,orn,m_data->m_pickFixedBody);
}
}
pivotInB.w = 0.f;
m_data->m_pickPivotInA = pivotInA;
m_data->m_pickPivotInB = pivotInB;
m_data->m_rigidBodyPipeline->copyConstraintsToHost();
m_data->m_pickConstraint = m_data->m_rigidBodyPipeline->createPoint2PointConstraint(hitBodyA,m_data->m_pickFixedBody,pivotInA,pivotInB,1e30);//hitResults[0].m_hitResult0
m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
m_data->m_np->writeAllBodiesToGpu();
m_data->m_pickDistance = (pivotInB-camPos).length();
return true;
}
}
}
} else
{
if (button==0)
{
if (m_data->m_pickConstraint>=0)
{
m_data->m_rigidBodyPipeline->removeConstraintByUid(m_data->m_pickConstraint);
m_data->m_pickConstraint=-1;
}
}
}
//printf("button=%d, state=%d\n",button,state);
return false;
}
@@ -0,0 +1,47 @@
#ifndef GPU_RIGID_BODY_DEMO_H
#define GPU_RIGID_BODY_DEMO_H
#include "Bullet3Common/b3Vector3.h"
#include "../CommonOpenCL/CommonOpenCLBase.h"
class GpuRigidBodyDemo : public CommonOpenCLBase
{
protected:
class GLInstancingRenderer* m_instancingRenderer;
class GLPrimitiveRenderer* m_primRenderer;
class CommonWindowInterface* m_window;
struct GpuRigidBodyDemoInternalData* m_data;
public:
GpuRigidBodyDemo(GUIHelperInterface* helper);
virtual ~GpuRigidBodyDemo();
virtual void initPhysics();
virtual void setupScene();
virtual void destroyScene(){};
virtual void exitPhysics();
virtual void renderScene();
virtual void stepSimulation(float deltaTime);
//for picking
b3Vector3 getRayTo(int x,int y);
virtual bool mouseMoveCallback(float x,float y);
virtual bool mouseButtonCallback(int button, int state, float x, float y);
virtual bool keyboardCallback(int key, int state);
unsigned char* loadImage(const char* fileName, int& width, int& height, int& n);
};
#endif //GPU_RIGID_BODY_DEMO_H
@@ -0,0 +1,55 @@
#ifndef GPU_RIGIDBODY_INTERNAL_DATA_H
#define GPU_RIGIDBODY_INTERNAL_DATA_H
#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
#include "Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h"
#include "Bullet3Common/b3Vector3.h"
#include "Bullet3Collision/NarrowPhaseCollision/b3Config.h"
struct GpuRigidBodyDemoInternalData
{
cl_kernel m_copyTransformsToVBOKernel;
b3OpenCLArray<b3Vector4>* m_instancePosOrnColor;
class b3GpuRigidBodyPipeline* m_rigidBodyPipeline;
class b3GpuNarrowPhase* m_np;
class b3GpuBroadphaseInterface* m_bp;
class b3DynamicBvhBroadphase* m_broadphaseDbvt;
b3Vector3 m_pickPivotInA;
b3Vector3 m_pickPivotInB;
float m_pickDistance;
int m_pickBody;
int m_pickConstraint;
int m_altPressed;
int m_controlPressed;
int m_pickFixedBody;
int m_pickGraphicsShapeIndex;
int m_pickGraphicsShapeInstance;
b3Config m_config;
GpuRigidBodyDemoInternalData()
:m_instancePosOrnColor(0),
m_copyTransformsToVBOKernel(0), m_rigidBodyPipeline(0),
m_np(0),
m_bp(0),
m_broadphaseDbvt(0),
m_pickConstraint(-1),
m_pickFixedBody(-1),
m_pickGraphicsShapeIndex(-1),
m_pickGraphicsShapeInstance(-1),
m_pickBody(-1),
m_altPressed(0),
m_controlPressed(0)
{
}
};
#endif//GPU_RIGIDBODY_INTERNAL_DATA_H
@@ -0,0 +1,204 @@
#include "GpuSphereScene.h"
#include "GpuRigidBodyDemo.h"
#include "OpenGLWindow/ShapeData.h"
#include "OpenGLWindow/GLInstancingRenderer.h"
#include "Bullet3Common/b3Quaternion.h"
#include "OpenGLWindow/b3gWindowInterface.h"
#include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h"
#include "../GpuDemoInternalData.h"
#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
#include "OpenGLWindow/OpenGLInclude.h"
#include "OpenGLWindow/GLInstanceRendererInternalData.h"
#include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h"
#include "Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.h"
#include "Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.h"
#include "Bullet3Collision/NarrowPhaseCollision/b3Config.h"
#include "GpuRigidBodyDemoInternalData.h"
#include "Bullet3AppSupport/gwenUserInterface.h"
void GpuSphereScene::setupScene(const ConstructionInfo& ci)
{
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_indices)/sizeof(int);
//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int group=1;
int mask=1;
int index=0;
bool writeInstanceToGpu = false;
if (0)
{
float radius = 60;
int prevGraphicsShapeIndex = -1;
{
if (1)//radius>=100)
{
int numVertices = sizeof(detailed_sphere_vertices)/strideInBytes;
int numIndices = sizeof(detailed_sphere_indices)/sizeof(int);
prevGraphicsShapeIndex = ci.m_instancingRenderer->registerShape(&detailed_sphere_vertices[0],numVertices,detailed_sphere_indices,numIndices);
} else
{
bool usePointSprites = false;
if (usePointSprites)
{
int numVertices = sizeof(point_sphere_vertices)/strideInBytes;
int numIndices = sizeof(point_sphere_indices)/sizeof(int);
prevGraphicsShapeIndex = ci.m_instancingRenderer->registerShape(&point_sphere_vertices[0],numVertices,point_sphere_indices,numIndices,B3_GL_POINTS);
} else
{
if (radius>=10)
{
int numVertices = sizeof(medium_sphere_vertices)/strideInBytes;
int numIndices = sizeof(medium_sphere_indices)/sizeof(int);
prevGraphicsShapeIndex = ci.m_instancingRenderer->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 = ci.m_instancingRenderer->registerShape(&low_sphere_vertices[0],numVertices,low_sphere_indices,numIndices);
}
}
}
}
b3Vector4 colors[4] =
{
b3MakeVector4(1,0,0,1),
b3MakeVector4(0,1,0,1),
b3MakeVector4(0,1,1,1),
b3MakeVector4(1,1,0,1),
};
int curColor = 0;
//int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
int colIndex = m_data->m_np->registerSphereShape(radius);//>registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
float mass = 0.f;
//b3Vector3 position((j&1)+i*2.2,1+j*2.,(j&1)+k*2.2);
b3Vector3 position=b3MakeVector3(0,0,0);
b3Quaternion orn(0,0,0,1);
b3Vector4 color = colors[curColor];
curColor++;
curColor&=3;
b3Vector4 scaling=b3MakeVector4(radius,radius,radius,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(prevGraphicsShapeIndex,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index, writeInstanceToGpu);
index++;
}
b3Vector4 colors[4] =
{
b3MakeVector4(1,0,0,1),
b3MakeVector4(0,1,0,1),
b3MakeVector4(0,1,1,1),
b3MakeVector4(1,1,0,1),
};
int curColor = 0;
float radius = 61;
//int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
int colIndex = m_data->m_np->registerSphereShape(radius);//>registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
int prevGraphicsShapeIndex = registerGraphicsSphereShape(ci,radius,false);
//for (int i=0;i<ci.arraySizeX;i++)
{
// for (int j=0;j<ci.arraySizeY;j++)
{
// for (int k=0;k<ci.arraySizeZ;k++)
{
int i=0,j=0,k=0;
float mass = 0.f;
b3Vector3 position=b3MakeVector3(0,0,0);
//b3Vector3 position((j&1)+i*142.2,-51+j*142.,(j&1)+k*142.2);
//b3Vector3 position(0,-41,0);//0,0,0);//i*radius*3,-41+j*radius*3,k*radius*3);
b3Quaternion orn(0,0,0,1);
b3Vector4 color = colors[curColor];
curColor++;
curColor&=3;
b3Vector4 scaling=b3MakeVector4(radius,radius,radius,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(prevGraphicsShapeIndex,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index, writeInstanceToGpu);
index++;
}
}
}
if (1)
{
int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
b3Vector4 scaling=b3MakeVector4(0.5,0.5,0.5,1);//1,1,1,1);//0.1,0.1,0.1,1);
int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
b3Vector3 normal=b3MakeVector3(0,-1,0);
float constant=2;
for (int j=-10;j<10;j++)
for (int i=-10;i<10;i++)
for (int k=0;k<30;k++)
//int i=0;int j=0;
{
//int colIndex = m_data->m_np->registerPlaneShape(normal,constant);//>registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
b3Vector4 position=b3MakeVector4(2*i,70+k*2,2*j+8,0);
//b3Quaternion orn(0,0,0,1);
b3Quaternion orn(b3MakeVector3(1,0,0),0.3);
b3Vector4 color=b3MakeVector4(0,0,1,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(1.f,position,orn,colIndex,index,false);
index++;
}
}
if (!writeInstanceToGpu)
{
m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
}
float camPos[4]={ci.arraySizeX,ci.arraySizeY/2,ci.arraySizeZ,0};
//float camPos[4]={1,12.5,1.5,0};
m_instancingRenderer->setCameraTargetPosition(camPos);
m_instancingRenderer->setCameraDistance(130);
char msg[1024];
int numInstances = index;
sprintf(msg,"Num objects = %d",numInstances);
ci.m_gui->setStatusBarMessage(msg,true);
}
@@ -0,0 +1,27 @@
#ifndef GPU_SPHERE_SCENE_H
#define GPU_SPHERE_SCENE_H
#include "GpuRigidBodyDemo.h"
class GpuSphereScene : public GpuRigidBodyDemo
{
public:
GpuSphereScene(){}
virtual ~GpuSphereScene(){}
virtual const char* getName()
{
return "BoxOnSphere";
}
static GpuDemo* MyCreateFunc()
{
GpuDemo* demo = new GpuSphereScene;
return demo;
}
virtual void setupScene(const ConstructionInfo& ci);
};
#endif //GPU_SPHERE_SCENE_H
+4
View File
@@ -0,0 +1,4 @@
static const char* getNodes() { return(
"454 3 0 0\n"" 0 -2.0063499999999999 0.79804200000000003 0.373554\n"" 1 -2.1011299999999999 0.90212400000000004 -0.88661400000000001\n"" 2 -1.4052100000000001 2.0628700000000002 -1.0458400000000001\n"" 3 -1.2015499999999999 1.7112400000000001 0.56249400000000005\n"" 4 0.02112 2.8512499999999998 -0.95618999999999998\n"" 5 0.011136 2.5152199999999998 0.59149200000000002\n"" 6 -1.51681 0.56199600000000005 1.42523\n"" 7 -0.977406 1.4278999999999999 1.24143\n"" 8 0.0051900000000000002 1.9088499999999999 1.41222\n"" 9 -2.4877400000000001 0.98449799999999998 -1.6695199999999999\n"" 10 -1.57328 2.1440000000000001 -1.7594799999999999\n"" 11 0.027768000000000001 2.8961600000000001 -2.0317599999999998\n"" 12 -2.4129700000000001 0.80116799999999999 -2.6594799999999998\n"" 13 -1.46269 1.9456500000000001 -2.6205799999999999\n"" 14 0.031758000000000002 2.6255500000000001 -2.7499899999999999\n"" 15 -2.0392999999999999 -0.24690000000000001 -4.0092699999999999\n"" 16 -1.4902899999999999 1.53495 -3.7649599999999999\n"" 17 0.037566000000000002 2.2566199999999998 -3.78904\n"" 18 -1.29721 -0.760656 -5.3216200000000002\n"" 19 -1.0264500000000001 0.069264000000000006 -5.2883199999999997\n"" 20 -1.0864400000000001 0.58933800000000003 -4.8886700000000003\n"" 21 -0.71934600000000004 1.3127200000000001 -4.5609200000000003\n"" 22 -0.47337000000000001 1.6606799999999999 -4.4356900000000001\n"" 23 0.040806000000000002 1.86575 -4.4139699999999999\n"" 24 -1.01305 0.61432200000000003 -5.5222899999999999\n"" 25 -0.62443199999999999 1.0636699999999999 -5.7151800000000001\n"" 26 -0.77822400000000003 1.2710900000000001 -5.0200699999999996\n"" 27 -0.59924999999999995 1.8655900000000001 -4.7962899999999999\n"" 28 0.043422000000000002 2.1701199999999998 -4.7666300000000001\n"" 29 -0.467478 1.5525199999999999 -5.5298400000000001\n"" 30 0.047742 1.6934499999999999 -5.5900800000000004\n"" 31 -1.5133300000000001 -3.3024100000000001 -3.3468599999999999\n"" 32 -1.6057999999999999 -3.6205099999999999 -3.9312499999999999\n"" 33 -2.6810299999999998 -0.71110200000000001 -2.7969499999999998\n"" 34 -2.7569300000000001 -0.55810199999999999 -1.86805\n"" 35 -2.2238699999999998 -0.60064799999999996 -0.95672400000000002\n"" 36 -2.2319 -0.55194600000000005 -0.066264000000000003\n"" 37 -1.97367 -0.58961399999999997 0.53195400000000004\n"" 38 -1.6947099999999999 -0.11317199999999999 1.8690100000000001\n"" 39 -2.11442 -0.78804600000000002 0.86941199999999996\n"" 40 -2.1847599999999998 -1.2017899999999999 1.2143299999999999\n"" 41 -1.7037899999999999 -1.3912100000000001 0.142008\n"" 42 -1.79366 -0.93130800000000002 2.2183000000000002\n"" 43 -1.7627200000000001 -0.73113600000000001 2.51458\n"" 44 -1.74098 -1.74478 0.64694399999999996\n"" 45 -1.5849899999999999 -1.6370899999999999 0.218082\n"" 46 -1.3714 -2.2354400000000001 1.7418499999999999\n"" 47 -1.14259 -1.7219800000000001 2.5314999999999999\n"" 48 -1.1462300000000001 -1.44584 3.04271\n"" 49 -1.72363 -1.65859 -0.39266400000000001\n"" 50 -1.7740499999999999 -1.6849099999999999 -1.0451999999999999\n"" 51 -2.3972199999999999 -1.87879 -2.2570000000000001\n"" 52 -2.3559999999999999 -2.9314900000000002 -2.5649600000000001\n"" 53 -1.9900100000000001 -3.4211100000000001 -2.7963200000000001\n"" 54 -1.69374 -3.7088399999999999 -3.53532\n"" 55 -2.2454299999999998 -3.5692900000000001 -1.9397899999999999\n"" 56 -2.2864300000000002 -3.77834 -2.10466\n"" 57 -2.2926700000000002 -3.7443599999999999 -1.3294600000000001\n"" 58 -1.6362099999999999 -3.39913 1.55494\n"" 59 -1.5402100000000001 -3.9804400000000002 1.7164699999999999\n"" 60 -1.6856899999999999 -2.5701499999999999 0.33473999999999998\n"" 61 -1.1098399999999999 -3.0533600000000001 1.9575899999999999\n"" 62 -1.6798299999999999 -3.1615099999999998 2.3719100000000002\n"" 63 -1.6955899999999999 -3.9803600000000001 2.4744700000000002\n"" 64 -1.1299699999999999 -2.8505600000000002 2.5077199999999999\n"" 65 -1.5803 -3.9803799999999998 2.79962\n"" 66 -1.2543800000000001 -3.9803600000000001 3.05606\n"" 67 -0.012264000000000001 -1.91774 3.31847\n"" 68 -0.0075960000000000003 -2.22756 2.4797799999999999\n"" 69 -1.3185199999999999 -2.03857 -0.24552599999999999\n"" 70 -1.54192 -1.6950700000000001 -0.038094000000000003\n"" 71 -1.6302399999999999 -1.5652900000000001 0.010584\n"" 72 -0.54714600000000002 -2.5150999999999999 -0.27547199999999999\n"" 73 -0.68966400000000005 -2.5785100000000001 -0.74843400000000004\n"" 74 -0.68381999999999998 -2.29792 -1.13124\n"" 75 -1.45807 -2.78965 -1.4571000000000001\n"" 76 -1.88889 -3.03241 Line truncated
static const char* getElements() { return(
"1374 4 0\n"" 0 395 372 399 371\n"" 1 96 46 93 44\n"" 2 135 140 137 163\n"" 3 318 326 354 313\n"" 4 133 127 128 126\n"" 5 76 83 55 218\n"" 6 193 188 198 176\n"" 7 450 443 451 1\n"" 8 437 442 438 441\n"" 9 160 336 320 144\n"" 10 79 50 51 75\n"" 11 408 378 387 273\n"" 12 94 38 47 68\n"" 13 323 308 322 311\n"" 14 339 135 163 341\n"" 15 152 135 153 145\n"" 16 333 166 160 147\n"" 17 348 346 347 342\n"" 18 446 426 451 444\n"" 19 46 92 94 93\n"" 20 208 21 26 207\n"" 21 38 360 48 413\n"" 22 229 80 243 245\n"" 23 241 229 227 242\n"" 24 34 50 35 51\n"" 25 9 33 35 34\n"" 26 103 122 121 206\n"" 27 257 260 259 440\n"" 28 142 318 163 144\n"" 29 446 443 444 451\n"" 30 272 251 252 274\n"" 31 193 199 176 198\n"" 32 165 149 147 166\n"" 33 447 3 8 449\n"" 34 29 208 209 30\n"" 35 300 285 286 299\n"" 36 100 104 119 105\n"" 37 39 44 40 41\n"" 38 315 318 313 316\n"" 39 422 417 421 419\n"" 40 382 273 264 383\n"" 41 145 143 150 146\n"" 42 406 268 405 404\n"" 43 427 411 426 428\n"" 44 323 290 308 311\n"" 45 232 18 230 227\n"" 46 243 79 80 240\n"" 47 391 412 390 389\n"" 48 360 423 107 362\n"" 49 389 71 393 392\n"" 50 142 162 163 145\n"" 51 245 247 261 80\n"" 52 376 373 375 92\n"" 53 324 352 325 362\n"" 54 119 104 106 105\n"" 55 165 149 166 151\n"" 56 428 389 411 442\n"" 57 366 351 331 352\n"" 58 166 318 316 163\n"" 59 368 281 280 350\n"" 60 315 168 313 166\n"" 61 417 422 423 419\n"" 62 279 203 205 453\n"" 63 80 243 245 226\n"" 64 416 41 420 45\n"" 65 319 330 317 335\n"" 66 15 442 225 224\n"" 67 116 131 117 129\n"" 68 335 318 337 319\n"" 69 188 197 193 198\n"" 70 326 314 325 313\n"" 71 313 315 314 325\n"" 72 451 3 449 450\n"" 73 451 449 452 450\n"" 74 299 286 300 301\n"" 75 450 1 3 2\n"" 76 107 103 120 149\n"" 77 76 82 81 79\n"" 78 81 83 218 214\n"" 79 14 10 2 11\n"" 80 22 436 437 23\n"" 81 318 144 316 163\n"" 82 420 426 425 412\n"" 83 450 14 11 445\n"" 84 262 254 429 274\n"" 85 443 442 441 444\n"" 86 356 353 364 363\n"" 87 102 122 206 101\n"" 88 315 168 325 313\n"" 89 281 331 280 350\n"" 90 107 155 120 104\n"" 91 443 426 444 451\n"" 92 123 118 354 113\n"" 93 84 220 86 83\n"" 94 268 406 405 269\n"" 95 287 293 289 286\n"" 96 17 16 443 13\n"" 97 315 168 166 325\n"" 98 283 296 297 284\n"" 99 333 144 336 160\n"" 100 450 13 1 2\n"" 101 414 423 418 419\n"" 102 331 352 324 362\n"" 103 144 315 333 316\n"" 104 21 442 16 15\n"" 105 45 72 91 69\n"" 106 437 208 21 433\n"" 107 288 310 290 306\n"" 108 213 217 214 31\n"" 109 335 327 328 341\n"" 110 109 67 48 111\n"" 111 316 314 315 312\n"" 112 45 60 91 44\n"" 113 128 118 115 116\n"" 114 423 45 38 448\n"" 115 155 362 166 325\n"" 116 130 136 132 133\n"" 117 92 72 89 91\n"" 118 145 143 142 150\n"" 119 153 140 163 142\n"" 120 317 313 312 319\n"" 121 347 118 126 354\n"" 122 148 165 147 143\n"" 123 170 157 171 173\n"" 124 78 50 74 240\n"" 125 381 273 385 378\n"" 126 337 330 335 317\n"" 127 267 273 379 268\n"" 128 80 247 261 250\n"" 129 80 225 261 245\n"" 130 80 442 261 225\n"" 131 211 24 212 20\n"" 132 398 395 399 371\n"" 133 390 389 377 391\n"" 134 308 292 290 291\n"" 135 176 191 177 199\n"" 136 318 166 313 163\n"" 137 23 16 436 17\n"" 138 354 353 325 326\n"" 139 423 362 366 279\n"" 140 415 400 375 403\n"" 141 217 239 216 213\n"" 142 377 Line truncated
+4
View File
@@ -0,0 +1,4 @@
static const char* getNodes() { return(
"400 3 0 0\n"" 0 1 1 -1\n"" 1 1 -1 -1\n"" 2 -1 -1 -1\n"" 3 1 1 1\n"" 4 -1 -1 1\n"" 5 -1 1 -1\n"" 6 1 -1 1\n"" 7 -1 1 1\n"" 8 -1 -1 -0.000116\n"" 9 0.00039500000000000006 -1 1\n"" 10 1 -1 -0.00065799999999999995\n"" 11 1 0.00061700000000000004 -1\n"" 12 0.00098999999999999999 1 -1\n"" 13 -1 1 0.00045400000000000008\n"" 14 0.00052800000000008396 1 1\n"" 15 1 1 -0.00038699999999991519\n"" 16 -1 -5.9999999999504894e-006 1\n"" 17 -1 0.00021000000000004349 -1\n"" 18 1 0.00067400000000006344 1\n"" 19 1.0863393114549872e-007 3.0682525964387253e-008 1\n"" 20 0.00097299999999989062 -1 -1\n"" 21 -1.4065544284846965e-008 1 5.0429866114086973e-008\n"" 22 1 -7.0792689266241382e-008 -4.4728299621228391e-008\n"" 23 8.4069120565846353e-008 3.8392827594879719e-007 -1\n"" 24 1 0.49980668112692161 -0.4996917965899465\n"" 25 1 0.50006700000000004 -1\n"" 26 1 1 -0.50034809999999996\n"" 27 1 1 0.50025040000000009\n"" 28 0.50020930000000008 1 1\n"" 29 -1 0.50010129999999997 1\n"" 30 -0.49958379999999991 1 1\n"" 31 -1 9.2038384777827427e-008 -1.100605984127867e-008\n"" 32 -0.5001201999999999 -1 1\n"" 33 1.0486788869007099e-007 -1 -3.3620911388254626e-008\n"" 34 -1 -1 0.50008699999999995\n"" 35 1 -0.50010220000000005 -1\n"" 36 1 -1 -0.49985400000000002\n"" 37 1 -0.4996526 1\n"" 38 1 -1 0.49990490000000004\n"" 39 -0.49930649999999999 1 -1\n"" 40 -1 1 -0.50002420000000003\n"" 41 1 -0.24974260000000004 -0.50015432157969908\n"" 42 0.50015422769290296 -0.24974260000000001 -1\n"" 43 0.50009680402228707 1 0.24993170000000009\n"" 44 1 0.50009675187018254 0.24993170000000009\n"" 45 0.24987467313200107 0.24987468273492527 -0.24987454967853764\n"" 46 0.50006970000000006 1 -1\n"" 47 -1 1 0.50062509999999982\n"" 48 -0.24952789999999991 1 0.50013191885453001\n"" 49 -0.24952789999999991 0.50013187836130635 1\n"" 50 -1 -0.50000239999999996 -1\n"" 51 -0.49917350000000005 -1 -1\n"" 52 -1 -0.37500110896930483 0.49999770999384097\n"" 53 -0.49999768418870116 -0.37500109958532857 1\n"" 54 0.50016842960649455 -0.24948929999999997 1\n"" 55 1 -0.24948929999999997 0.50016837821653382\n"" 56 -1 0.50011355427218929 -0.24978510000000004\n"" 57 -0.50011349993113807 1 -0.24978510000000009\n"" 58 0.49983560886453382 -1 -0.25025600000000003\n"" 59 1 -0.24959280000000006 -1\n"" 60 0.46861653427920624 0.31279234072878337 -1\n"" 61 0.46874721181012263 1 -0.31263609558485356\n"" 62 1 1 0.24986720000000007\n"" 63 1 0.50004500000000007 1\n"" 64 0.41666006216462159 0.41668749191836219 1\n"" 65 1 0.62492318088559184 0.62497446433589154\n"" 66 0.62493687874554793 1 0.62494713479954256\n"" 67 -0.32510780807963724 -0.32510780489834346 -0.32510793388507692\n"" 68 1 0.12495484505171366 -0.3748231366627352\n"" 69 -0.24915825000000003 0.50024732783000192 -1\n"" 70 1 0.74957571448653726 -0.74971600694021778\n"" 71 0.64874656072370729 1 -0.64882932411306993\n"" 72 0.6488668035834122 0.64886599996407779 -1\n"" 73 -1 0.41673079803079871 0.41664307505523179\n"" 74 -0.62468209179782908 1 0.62494186195929313\n"" 75 -0.62476903429567798 0.62489835940819183 1\n"" 76 -1 0.75020764999999989 1\n"" 77 0.12520010648567331 0.45856412436943739 0.37477910185628149\n"" 78 -0.24942439999999994 1 1\n"" 79 -0.37474797737976173 1 0.74996984889996166\n"" 80 -0.3748348177892612 0.74992638728967431 1\n"" 81 -1 -1 -0.49994339999999993\n"" 82 -0.50002894420427402 -1 0.24998549999999997\n"" 83 -1 -0.50045000000000006 1\n"" 84 -0.015617775048046156 -0.5000062299955631 1\n"" 85 -0.062520154989433377 -1 0.50002475010607028\n"" 86 0.50024029999999997 -1 1\n"" 87 1 -0.24929619999999997 1\n"" 88 0.37505270907528421 -0.3750526949043349 0.12522791934472338\n"" 89 -1 0.5001348000000001 -1\n"" 90 -1 1 -0.24999450000000001\n"" 91 1 -0.62477132796083668 -0.25025600000000003\n"" 92 0.50052129999999995 -1 -1\n"" 93 1 0.018238422269637886 -0.6785517711937521\n"" 94 0.750039598406532 -0.12448790000000004 -1\n"" 95 0.32367160517680327 0.022514689090561368 -0.65701963759824777\n"" 96 0.32369468988306127 0.65698780035674154 -0.022395524628340657\n"" 97 1 1 -0.25020424999999996\n"" 98 1 0.62501462613012171 -0.12492844204372336\n"" 99 0.16205966081319645 0.60689608104137205 -1\n"" 100 0.23307615322925276 1 -0.65509826790501591\n"" 101 0.6349049175761754 0.105723024682608 1\n"" 102 1 0.18290638148392097 0.59627500837254621\n"" 103 0.67708420758319754 0.67702605426027662 1\n"" 104 1 0.7498625000000001 1Line truncated
static const char* getElements() { return(
"1580 4 0\n"" 0 120 148 13 240\n"" 1 77 251 185 268\n"" 2 126 194 66 198\n"" 3 297 298 159 357\n"" 4 13 148 147 240\n"" 5 219 272 102 303\n"" 6 134 137 37 136\n"" 7 302 360 88 361\n"" 8 23 264 95 391\n"" 9 307 311 238 364\n"" 10 57 309 254 326\n"" 11 345 357 233 363\n"" 12 176 343 325 384\n"" 13 275 332 67 343\n"" 14 141 245 228 318\n"" 15 108 368 323 377\n"" 16 7 47 74 123\n"" 17 73 188 184 270\n"" 18 56 235 148 313\n"" 19 46 72 117 179\n"" 20 66 105 28 194\n"" 21 205 275 204 344\n"" 22 66 194 28 198\n"" 23 96 221 45 375\n"" 24 93 160 94 161\n"" 25 225 317 245 397\n"" 26 121 186 74 188\n"" 27 213 280 52 322\n"" 28 52 322 280 376\n"" 29 96 195 77 221\n"" 30 8 257 256 331\n"" 31 358 380 330 381\n"" 32 99 181 166 266\n"" 33 120 123 73 188\n"" 34 122 189 187 340\n"" 35 141 292 224 300\n"" 36 313 326 240 377\n"" 37 203 274 173 390\n"" 38 140 299 221 319\n"" 39 44 197 107 272\n"" 40 171 326 108 377\n"" 41 71 166 46 179\n"" 42 69 364 334 391\n"" 43 26 114 97 115\n"" 44 228 253 159 358\n"" 45 221 272 202 289\n"" 46 123 184 73 188\n"" 47 22 220 140 231\n"" 48 250 324 322 325\n"" 49 270 324 213 338\n"" 50 147 240 57 336\n"" 51 69 265 178 391\n"" 52 100 249 165 365\n"" 53 296 316 224 371\n"" 54 74 188 186 189\n"" 55 21 171 165 249\n"" 56 23 297 225 298\n"" 57 191 196 64 290\n"" 58 87 139 138 288\n"" 59 126 198 196 199\n"" 60 223 328 141 342\n"" 61 293 371 153 396\n"" 62 238 308 108 313\n"" 63 49 187 80 189\n"" 64 143 300 224 361\n"" 65 111 256 130 343\n"" 66 62 107 106 273\n"" 67 64 196 126 219\n"" 68 314 154 315 369\n"" 69 50 205 204 344\n"" 70 0 70 26 71\n"" 71 95 158 23 265\n"" 72 37 138 87 139\n"" 73 125 193 43 195\n"" 74 56 312 238 313\n"" 75 257 259 172 332\n"" 76 25 70 0 72\n"" 77 141 300 228 358\n"" 78 32 207 208 350\n"" 79 101 201 200 219\n"" 80 13 147 121 240\n"" 81 67 275 176 357\n"" 82 116 249 180 266\n"" 83 197 219 200 272\n"" 84 152 247 42 371\n"" 85 95 158 60 160\n"" 86 45 230 116 375\n"" 87 88 282 142 283\n"" 88 96 162 61 249\n"" 89 209 280 211 350\n"" 90 84 353 226 399\n"" 91 33 258 172 381\n"" 92 118 285 213 378\n"" 93 196 271 77 289\n"" 94 282 287 218 388\n"" 95 196 198 125 199\n"" 96 74 7 76 30\n"" 97 95 230 181 264\n"" 98 297 359 225 390\n"" 99 280 347 261 350\n"" 100 253 299 170 380\n"" 101 125 269 193 335\n"" 102 245 293 224 300\n"" 103 41 143 68 161\n"" 104 58 295 244 316\n"" 105 83 348 211 349\n"" 106 75 29 76 123\n"" 107 171 249 96 255\n"" 108 124 196 125 199\n"" 109 61 248 96 249\n"" 110 221 255 45 327\n"" 111 254 326 171 336\n"" 112 194 198 126 199\n"" 113 208 278 9 279\n"" 114 94 93 11 160\n"" 115 275 357 345 363\n"" 116 127 173 109 203\n"" 117 134 139 55 218\n"" 118 45 255 253 327\n"" 119 154 35 315 370\n"" 120 283 284 142 353\n"" 121 186 336 268 377\n"" 122 59 93 11 94\n"" 123 110 203 51 277\n"" 124 59 246 94 247\n"" 125 116 248 61 249\n"" 126 41 229 143 296\n"" 127 91 154 36 316\n"" 128 71 179 116 180\n"" 129 45 299 221 327\n"" 130 319 373 182 374\n"" 131 64 219 101 290\n"" 132 116 301 179 398\n"" 133 122 323 188 377\n"" 134 182 373 68 374\n"" 135 97 115 114 248\n"" 136 60 158 95 181\n"" 137 185 268 48 335\n"" 138 66 125 43 195\n"" 139 86 215 136 287\n"" 140 106 202 163 273\n"" 141 43 106 66 195\n"" 142 271 Line truncated
@@ -1031,9 +1031,9 @@ void b3GpuSapBroadphase::calculateOverlappingPairs(int maxPairs)
if (m_dst.size()!=(numSmallAabbs+1)) if (m_dst.size()!=(numSmallAabbs+1))
{ {
m_dst.resize(numSmallAabbs+1); m_dst.resize(numSmallAabbs+128);
m_sum.resize(numSmallAabbs+1); m_sum.resize(numSmallAabbs+128);
m_sum2.resize(numSmallAabbs+1); m_sum2.resize(numSmallAabbs+128);
m_sum.at(numSmallAabbs)=b3MakeVector3(0,0,0); //slow? m_sum.at(numSmallAabbs)=b3MakeVector3(0,0,0); //slow?
m_sum2.at(numSmallAabbs)=b3MakeVector3(0,0,0); //slow? m_sum2.at(numSmallAabbs)=b3MakeVector3(0,0,0); //slow?
} }
@@ -1044,8 +1044,8 @@ void b3GpuSapBroadphase::calculateOverlappingPairs(int maxPairs)
launcher.setBuffer(m_smallAabbsMappingGPU.getBufferCL()); launcher.setBuffer(m_smallAabbsMappingGPU.getBufferCL());
launcher.setBuffer(m_sum.getBufferCL()); launcher.setBuffer(m_sum.getBufferCL());
launcher.setBuffer(m_sum2.getBufferCL()); launcher.setBuffer(m_sum2.getBufferCL());
launcher.setConst( numSmallAabbs+1 ); launcher.setConst( numSmallAabbs );
int num = numSmallAabbs+1; int num = numSmallAabbs;
launcher.launch1D( num); launcher.launch1D( num);
@@ -377,7 +377,7 @@ __kernel void scatterKernel( __global const btAabbCL* allAabbs, __global const
__kernel void prepareSumVarianceKernel( __global const btAabbCL* allAabbs, __global const int* smallAabbMapping, __global float4* sum, __global float4* sum2,int numAabbs) __kernel void prepareSumVarianceKernel( __global const btAabbCL* allAabbs, __global const int* smallAabbMapping, __global float4* sum, __global float4* sum2,int numAabbs)
{ {
int i = get_global_id(0); int i = get_global_id(0);
if (i>numAabbs) if (i>=numAabbs)
return; return;
btAabbCL smallAabb = allAabbs[smallAabbMapping[i]]; btAabbCL smallAabb = allAabbs[smallAabbMapping[i]];
@@ -329,7 +329,7 @@ static const char* sapCL= \
"__kernel void prepareSumVarianceKernel( __global const btAabbCL* allAabbs, __global const int* smallAabbMapping, __global float4* sum, __global float4* sum2,int numAabbs)\n" "__kernel void prepareSumVarianceKernel( __global const btAabbCL* allAabbs, __global const int* smallAabbMapping, __global float4* sum, __global float4* sum2,int numAabbs)\n"
"{\n" "{\n"
" int i = get_global_id(0);\n" " int i = get_global_id(0);\n"
" if (i>numAabbs)\n" " if (i>=numAabbs)\n"
" return;\n" " return;\n"
" \n" " \n"
" btAabbCL smallAabb = allAabbs[smallAabbMapping[i]];\n" " btAabbCL smallAabb = allAabbs[smallAabbMapping[i]];\n"