From 3ed7ecbb58814dcd159f328f76ca967ef89ccdca Mon Sep 17 00:00:00 2001 From: hujiajie Date: Fri, 11 Mar 2016 09:40:40 +0800 Subject: [PATCH] Fix ternary selection operator on float scalars in OpenCL kernels. According to Section 6.3(i) of the OpenCL 1.1 specification, the first expression of the operator cannot be a float. --- src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl | 4 ++-- src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h | 4 ++-- src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl | 4 ++-- src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h | 4 ++-- 4 files changed, 8 insertions(+), 8 deletions(-) diff --git a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl index 0ad5161bc..7f5dabe27 100644 --- a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl +++ b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl @@ -410,8 +410,8 @@ __kernel void initBatchConstraintsKernel(__global unsigned int* numConstraintRow int rbA = constraints[i].m_rbA; int rbB = constraints[i].m_rbB; - batchConstraints[i].m_bodyAPtrAndSignBit = bodies[rbA].m_invMass? rbA : -rbA; - batchConstraints[i].m_bodyBPtrAndSignBit = bodies[rbB].m_invMass? rbB : -rbB; + batchConstraints[i].m_bodyAPtrAndSignBit = bodies[rbA].m_invMass != 0.f ? rbA : -rbA; + batchConstraints[i].m_bodyBPtrAndSignBit = bodies[rbB].m_invMass != 0.f ? rbB : -rbB; batchConstraints[i].m_batchId = -1; batchConstraints[i].m_originalConstraintIndex = i; diff --git a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h index e011ac7dc..d48ecf6ea 100644 --- a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h +++ b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h @@ -325,8 +325,8 @@ static const char* solveConstraintRowsCL= \ " return;\n" " int rbA = constraints[i].m_rbA;\n" " int rbB = constraints[i].m_rbB;\n" -" batchConstraints[i].m_bodyAPtrAndSignBit = bodies[rbA].m_invMass? rbA : -rbA;\n" -" batchConstraints[i].m_bodyBPtrAndSignBit = bodies[rbB].m_invMass? rbB : -rbB;\n" +" batchConstraints[i].m_bodyAPtrAndSignBit = bodies[rbA].m_invMass != 0.f ? rbA : -rbA;\n" +" batchConstraints[i].m_bodyBPtrAndSignBit = bodies[rbB].m_invMass != 0.f ? rbB : -rbB;\n" " batchConstraints[i].m_batchId = -1;\n" " batchConstraints[i].m_originalConstraintIndex = i;\n" "}\n" diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl index 9c4ac4ce8..a21a08c3b 100644 --- a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl +++ b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl @@ -954,8 +954,8 @@ float positionConstraintCoeff Constraint4 cs; - float countA = invMassA ? (float)bodyCount[aIdx] : 1; - float countB = invMassB ? (float)bodyCount[bIdx] : 1; + float countA = invMassA != 0.f ? (float)bodyCount[aIdx] : 1; + float countB = invMassB != 0.f ? (float)bodyCount[bIdx] : 1; setConstraint4( posA, linVelA, angVelA, invMassA, invInertiaA, posB, linVelB, angVelB, invMassB, invInertiaB, &gContact[gIdx], dt, positionDrift, positionConstraintCoeff,countA,countB, diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h index fef9dd2a6..c0173ad9f 100644 --- a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h +++ b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h @@ -896,8 +896,8 @@ static const char* solverUtilsCL= \ " float invMassB = gBodies[bIdx].m_invMass;\n" " Matrix3x3 invInertiaB = gShapes[bIdx].m_invInertia;\n" " Constraint4 cs;\n" -" float countA = invMassA ? (float)bodyCount[aIdx] : 1;\n" -" float countB = invMassB ? (float)bodyCount[bIdx] : 1;\n" +" float countA = invMassA != 0.f ? (float)bodyCount[aIdx] : 1;\n" +" float countB = invMassB != 0.f ? (float)bodyCount[bIdx] : 1;\n" " setConstraint4( posA, linVelA, angVelA, invMassA, invInertiaA, posB, linVelB, angVelB, invMassB, invInertiaB,\n" " &gContact[gIdx], dt, positionDrift, positionConstraintCoeff,countA,countB,\n" " &cs );\n"