fix Linux version of ConstraintDemo

This commit is contained in:
erwincoumans 2013-07-09 14:41:56 -07:00
parent 4a0d2847c4
commit 282e19171f
6 changed files with 43 additions and 37 deletions

View File

@ -53,13 +53,13 @@ void GpuConstraintsDemo::setupScene(const ConstructionInfo& ci)
m_data->m_rigidBodyPipeline->writeAllInstancesToGpu();
// m_data->m_rigidBodyPipeline->setGravity(b3Vector3(4,-10,0));
float camPos[4]={ci.arraySizeX,ci.arraySizeY/2,ci.arraySizeZ,0};
float camPos[4]={ci.arraySizeX,ci.gapY*ci.arraySizeY/2,ci.arraySizeZ,0};
//float camPos[4]={1,12.5,1.5,0};
m_instancingRenderer->setCameraTargetPosition(camPos);
m_instancingRenderer->setCameraDistance(30);
m_instancingRenderer->setCameraDistance(180);
m_instancingRenderer->setCameraPitch(200);//90);
m_instancingRenderer->updateCamera();
char msg[1024];

View File

@ -14,15 +14,17 @@
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.cl" --headerfile="../src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.h" --stringname="bvhTraversalKernelCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/integrateKernel.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/integrateKernel.h" --stringname="integrateKernelCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/updateAabbsKernel.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/updateAabbsKernel.h" --stringname="updateAabbsKernelCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/solverSetup.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solverSetup.h" --stringname="solverSetupCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/solverSetup2.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solverSetup2.h" --stringname="solverSetup2CL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/batchingKernels.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/batchingKernels.h" --stringname="batchingKernelsCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/batchingKernelsNew.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/batchingKernelsNew.h" --stringname="batchingKernelsNewCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/solverUtils.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solverUtils.h" --stringname="solverUtilsCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/solveContact.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solveContact.h" --stringname="solveContactCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/solveFriction.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solveFriction.h" --stringname="solveFrictionCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/integrateKernel.cl" --headerfile="../src/Bullet3OpenCL/RigidBody/kernels/integrateKernel.h" --stringname="integrateKernelCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/updateAabbsKernel.cl" --headerfile="../src/Bullet3OpenCL/RigidBody/kernels/updateAabbsKernel.h" --stringname="updateAabbsKernelCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/solverSetup.cl" --headerfile="../src/Bullet3OpenCL/RigidBody/kernels/solverSetup.h" --stringname="solverSetupCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.cl" --headerfile="../src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.h" --stringname="solverSetup2CL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.cl" --headerfile="../src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.h" --stringname="batchingKernelsCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.cl" --headerfile="../src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.h" --stringname="batchingKernelsNewCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl" --headerfile="../src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h" --stringname="solverUtilsCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/solveContact.cl" --headerfile="../src/Bullet3OpenCL/RigidBody/kernels/solveContact.h" --stringname="solveContactCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/solveFriction.cl" --headerfile="../src/Bullet3OpenCL/RigidBody/kernels/solveFriction.h" --stringname="solveFrictionCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl" --headerfile="../src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h" --stringname="solveConstraintRowsCL" stringify
./premake4_linux --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl" --headerfile="../src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h" --stringname="rayCastKernelCL" stringify

View File

@ -40,9 +40,10 @@ B3_ATTRIBUTE_ALIGNED16 (struct) b3SolverConstraint
b3Vector3 m_angularComponentA;
b3Vector3 m_angularComponentB;
mutable b3SimdScalar m_appliedPushImpulse;
mutable b3SimdScalar m_appliedImpulse;
mutable b3Scalar m_appliedPushImpulse;
mutable b3Scalar m_appliedImpulse;
int m_padding1;
int m_padding2;
b3Scalar m_friction;
b3Scalar m_jacDiagABInv;
b3Scalar m_rhs;

View File

@ -133,9 +133,10 @@ typedef struct
float4 m_angularComponentA;
float4 m_angularComponentB;
float4 m_appliedPushImpulse;
float4 m_appliedImpulse;
float m_appliedPushImpulse;
float m_appliedImpulse;
int m_padding1;
int m_padding2;
float m_friction;
float m_jacDiagABInv;
float m_rhs;
@ -272,27 +273,27 @@ __inline void internalApplyImpulse(__global b3GpuSolverBody* body, float4 linea
void resolveSingleConstraintRowGeneric(__global b3GpuSolverBody* body1, __global b3GpuSolverBody* body2, __global b3SolverConstraint* c)
{
float deltaImpulse = c->m_rhs-c->m_appliedImpulse.x*c->m_cfm;
float deltaImpulse = c->m_rhs-c->m_appliedImpulse*c->m_cfm;
float deltaVel1Dotn = dot3F4(c->m_contactNormal,body1->m_deltaLinearVelocity) + dot3F4(c->m_relpos1CrossNormal,body1->m_deltaAngularVelocity);
float deltaVel2Dotn = -dot3F4(c->m_contactNormal,body2->m_deltaLinearVelocity) + dot3F4(c->m_relpos2CrossNormal,body2->m_deltaAngularVelocity);
deltaImpulse -= deltaVel1Dotn*c->m_jacDiagABInv;
deltaImpulse -= deltaVel2Dotn*c->m_jacDiagABInv;
float sum = c->m_appliedImpulse.x + deltaImpulse;
float sum = c->m_appliedImpulse + deltaImpulse;
if (sum < c->m_lowerLimit)
{
deltaImpulse = c->m_lowerLimit-c->m_appliedImpulse.x;
c->m_appliedImpulse.x = c->m_lowerLimit;
deltaImpulse = c->m_lowerLimit-c->m_appliedImpulse;
c->m_appliedImpulse = c->m_lowerLimit;
}
else if (sum > c->m_upperLimit)
{
deltaImpulse = c->m_upperLimit-c->m_appliedImpulse.x;
c->m_appliedImpulse.x = c->m_upperLimit;
deltaImpulse = c->m_upperLimit-c->m_appliedImpulse;
c->m_appliedImpulse = c->m_upperLimit;
}
else
{
c->m_appliedImpulse.x = sum;
c->m_appliedImpulse = sum;
}
internalApplyImpulse(body1,c->m_contactNormal*body1->m_invMass,c->m_angularComponentA,deltaImpulse);
@ -730,4 +731,4 @@ __kernel void getInfo2Kernel(__global b3SolverConstraint* solverConstraintRows,
}
}
}
}
}

View File

@ -135,9 +135,10 @@ static const char* solveConstraintRowsCL= \
" float4 m_angularComponentA;\n"
" float4 m_angularComponentB;\n"
" \n"
" float4 m_appliedPushImpulse;\n"
" float4 m_appliedImpulse;\n"
"\n"
" float m_appliedPushImpulse;\n"
" float m_appliedImpulse;\n"
" int m_padding1;\n"
" int m_padding2;\n"
" float m_friction;\n"
" float m_jacDiagABInv;\n"
" float m_rhs;\n"
@ -274,27 +275,27 @@ static const char* solveConstraintRowsCL= \
"\n"
"void resolveSingleConstraintRowGeneric(__global b3GpuSolverBody* body1, __global b3GpuSolverBody* body2, __global b3SolverConstraint* c)\n"
"{\n"
" float deltaImpulse = c->m_rhs-c->m_appliedImpulse.x*c->m_cfm;\n"
" float deltaImpulse = c->m_rhs-c->m_appliedImpulse*c->m_cfm;\n"
" float deltaVel1Dotn = dot3F4(c->m_contactNormal,body1->m_deltaLinearVelocity) + dot3F4(c->m_relpos1CrossNormal,body1->m_deltaAngularVelocity);\n"
" float deltaVel2Dotn = -dot3F4(c->m_contactNormal,body2->m_deltaLinearVelocity) + dot3F4(c->m_relpos2CrossNormal,body2->m_deltaAngularVelocity);\n"
"\n"
" deltaImpulse -= deltaVel1Dotn*c->m_jacDiagABInv;\n"
" deltaImpulse -= deltaVel2Dotn*c->m_jacDiagABInv;\n"
"\n"
" float sum = c->m_appliedImpulse.x + deltaImpulse;\n"
" float sum = c->m_appliedImpulse + deltaImpulse;\n"
" if (sum < c->m_lowerLimit)\n"
" {\n"
" deltaImpulse = c->m_lowerLimit-c->m_appliedImpulse.x;\n"
" c->m_appliedImpulse.x = c->m_lowerLimit;\n"
" deltaImpulse = c->m_lowerLimit-c->m_appliedImpulse;\n"
" c->m_appliedImpulse = c->m_lowerLimit;\n"
" }\n"
" else if (sum > c->m_upperLimit) \n"
" {\n"
" deltaImpulse = c->m_upperLimit-c->m_appliedImpulse.x;\n"
" c->m_appliedImpulse.x = c->m_upperLimit;\n"
" deltaImpulse = c->m_upperLimit-c->m_appliedImpulse;\n"
" c->m_appliedImpulse = c->m_upperLimit;\n"
" }\n"
" else\n"
" {\n"
" c->m_appliedImpulse.x = sum;\n"
" c->m_appliedImpulse = sum;\n"
" }\n"
"\n"
" internalApplyImpulse(body1,c->m_contactNormal*body1->m_invMass,c->m_angularComponentA,deltaImpulse);\n"
@ -733,4 +734,5 @@ static const char* solveConstraintRowsCL= \
" }\n"
" }\n"
"}\n"
"\n"
;

View File

@ -1,7 +1,7 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
static const char* solverUtilsCL= \
"/*\n"
"Copyright (c) 2012 Advanced Micro Devices, Inc. \n"
"Copyright (c) 2013 Advanced Micro Devices, Inc. \n"
"\n"
"This software is provided 'as-is', without any express or implied warranty.\n"
"In no event will the authors be held liable for any damages arising from the use of this software.\n"