diff --git a/Demos3/GpuDemos/constraints/ConstraintsDemo.cpp b/Demos3/GpuDemos/constraints/ConstraintsDemo.cpp index 64217736d..38e4344a7 100644 --- a/Demos3/GpuDemos/constraints/ConstraintsDemo.cpp +++ b/Demos3/GpuDemos/constraints/ConstraintsDemo.cpp @@ -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]; diff --git a/build3/stringify_linux.sh b/build3/stringify_linux.sh index 95f56424e..d8d6df4b8 100755 --- a/build3/stringify_linux.sh +++ b/build3/stringify_linux.sh @@ -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 diff --git a/src/Bullet3Dynamics/ConstraintSolver/b3SolverConstraint.h b/src/Bullet3Dynamics/ConstraintSolver/b3SolverConstraint.h index 79dbaceb9..8dd0e362a 100644 --- a/src/Bullet3Dynamics/ConstraintSolver/b3SolverConstraint.h +++ b/src/Bullet3Dynamics/ConstraintSolver/b3SolverConstraint.h @@ -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; diff --git a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl index 9925b2ee4..7f7b0d422 100644 --- a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl +++ b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl @@ -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, } } } -} \ No newline at end of file +} diff --git a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h index 4e0072975..e87642108 100644 --- a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h +++ b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h @@ -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" ; diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h index cdd3b66d4..29cd712a1 100644 --- a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h +++ b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h @@ -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"