From 0a360904b60cab5717c2306f470798ccf23e328a Mon Sep 17 00:00:00 2001 From: "erwin.coumans" Date: Sat, 9 May 2009 21:25:42 +0000 Subject: [PATCH] fixes in Jam msvcgen -> use SSE instead of SSE2 to stay compatible with older AMD Athlon XP processors, updated build files, minor compile fixes for GPU experimental code --- Demos/Box2dDemo/CMakeLists.txt | 32 ++ Demos/Box2dDemo/Jamfile | 3 + Demos/CMakeLists.txt | 4 +- Demos/Gpu2dDemo/CMakeLists.txt | 16 +- Demos/Gpu2dDemo/Jamfile | 4 +- Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h | 45 +- Demos/Jamfile | 2 + Demos/OpenGL/GLDebugFont.cpp | 3 +- Demos/OpenGL/GL_ShapeDrawer.cpp | 2 +- mk/msvcgen/control.tlib | 2 +- mk/msvcgen/projectx7.tlib | 12 +- mk/msvcgen/projectxsn71.tlib | 2 +- .../btGpu3DGridBroadphase.cpp | 1 + src/BulletMultiThreaded/btGpuDefines.h | 435 +++++++++--------- 14 files changed, 312 insertions(+), 251 deletions(-) create mode 100644 Demos/Box2dDemo/CMakeLists.txt create mode 100644 Demos/Box2dDemo/Jamfile diff --git a/Demos/Box2dDemo/CMakeLists.txt b/Demos/Box2dDemo/CMakeLists.txt new file mode 100644 index 000000000..e817417eb --- /dev/null +++ b/Demos/Box2dDemo/CMakeLists.txt @@ -0,0 +1,32 @@ +# This is basically the overall name of the project in Visual Studio this is the name of the Solution File + + +# For every executable you have with a main method you should have an add_executable line below. +# For every add executable line you should list every .cpp and .h file you have associated with that executable. + + +# This is the variable for Windows. I use this to define the root of my directory structure. +SET(GLUT_ROOT ${BULLET_PHYSICS_SOURCE_DIR}/Glut) + +# You shouldn't have to modify anything below this line +######################################################## + +INCLUDE_DIRECTORIES( +${BULLET_PHYSICS_SOURCE_DIR}/src ${BULLET_PHYSICS_SOURCE_DIR}/Demos/OpenGL } +) + +LINK_LIBRARIES( +OpenGLSupport BulletDynamics BulletCollision LinearMath ${GLUT_glut_LIBRARY} ${OPENGL_gl_LIBRARY} ${OPENGL_glu_LIBRARY} +) + +ADD_EXECUTABLE(AppBox2dDemo + main.cpp + Box2dDemo.cpp + btBox2dBox2dCollisionAlgorithm.h + Box2dDemo.h + btBox2dShape.cpp + btBox2dBox2dCollisionAlgorithm.cpp + btBox2dShape.h +) + + diff --git a/Demos/Box2dDemo/Jamfile b/Demos/Box2dDemo/Jamfile new file mode 100644 index 000000000..fde6f54b8 --- /dev/null +++ b/Demos/Box2dDemo/Jamfile @@ -0,0 +1,3 @@ +SubDir TOP Demos Box2dDemo ; + +BulletDemo Box2dDemo : [ Wildcard *.h *.cpp ] ; diff --git a/Demos/CMakeLists.txt b/Demos/CMakeLists.txt index 56eace1e5..f12ac4a2e 100644 --- a/Demos/CMakeLists.txt +++ b/Demos/CMakeLists.txt @@ -1,7 +1,7 @@ if (CMAKE_SIZEOF_VOID_P MATCHES "8") SUBDIRS( OpenGL AllBulletDemos ConvexDecompositionDemo Benchmarks HelloWorld CcdPhysicsDemo ConstraintDemo SliderConstraintDemo GenericJointDemo Raytracer -RagdollDemo ForkLiftDemo BasicDemo BspDemo MovingConcaveDemo VehicleDemo +RagdollDemo ForkLiftDemo BasicDemo Box2dDemo Gpu2dDemo BspDemo MovingConcaveDemo VehicleDemo ColladaDemo UserCollisionAlgorithm CharacterDemo SoftDemo HeightFieldFluidDemo CollisionInterfaceDemo ConcaveConvexcastDemo SimplexDemo DynamicControlDemo DoublePrecisionDemo ConcaveDemo CollisionDemo @@ -10,7 +10,7 @@ MultiMaterialDemo) else (CMAKE_SIZEOF_VOID_P MATCHES "8") SUBDIRS( OpenGL AllBulletDemos ConvexDecompositionDemo Benchmarks HelloWorld MultiThreadedDemo CcdPhysicsDemo ConstraintDemo SliderConstraintDemo Raytracer -GenericJointDemo RagdollDemo ForkLiftDemo BasicDemo BspDemo MovingConcaveDemo +GenericJointDemo RagdollDemo ForkLiftDemo BasicDemo Box2dDemo Gpu2dDemo BspDemo MovingConcaveDemo VehicleDemo ColladaDemo UserCollisionAlgorithm CharacterDemo SoftDemo HeightFieldFluidDemo CollisionInterfaceDemo ConcaveConvexcastDemo SimplexDemo DynamicControlDemo diff --git a/Demos/Gpu2dDemo/CMakeLists.txt b/Demos/Gpu2dDemo/CMakeLists.txt index 32baf8b88..707357c58 100644 --- a/Demos/Gpu2dDemo/CMakeLists.txt +++ b/Demos/Gpu2dDemo/CMakeLists.txt @@ -19,8 +19,20 @@ LINK_LIBRARIES( OpenGLSupport BulletDynamics BulletCollision LinearMath ${GLUT_glut_LIBRARY} ${OPENGL_gl_LIBRARY} ${OPENGL_glu_LIBRARY} ) -ADD_EXECUTABLE(AppBasicDemo +ADD_EXECUTABLE(AppGpu2dDemo main.cpp - BasicDemo.cpp + BasicDemo.cpp + BasicDemo.h + btGpuDemoPairCache.cpp + btGpuDemoPairCache.h + btGpuDemo2dSharedTypes.h + btGpuDemo2dCpuFunc.cpp + btGpuDemoDynamicsWorld.cpp + btGpuDemoDynamicsWorld.h + oecakeLoader.cpp + oecakeLoader.h + btGpuDemo2dSharedCode.h + btGpuDemo2dSharedDefs.h ) + diff --git a/Demos/Gpu2dDemo/Jamfile b/Demos/Gpu2dDemo/Jamfile index 3584562d7..e2f506add 100644 --- a/Demos/Gpu2dDemo/Jamfile +++ b/Demos/Gpu2dDemo/Jamfile @@ -1,3 +1,3 @@ -SubDir TOP Demos BasicDemo ; +SubDir TOP Demos Gpu2dDemo ; -BulletDemo BasicDemo : [ Wildcard *.h *.cpp ] ; +BulletDemo Gpu2dDemo : [ Wildcard *.h *.cpp ] ; diff --git a/Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h b/Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h index d1afc9450..270db1376 100644 --- a/Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h +++ b/Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h @@ -19,6 +19,7 @@ subject to the following restrictions: #define FRICTION_BOX_GROUND_FACT 0.05f #define FRICTION_BOX_BOX_FACT 0.05f #define USE_CENTERS 1 +#include "LinearMath/btMinMax.h" //------------------------------------------------------------------------------------------------ //------------------------------------------------------------------------------------------------ @@ -58,6 +59,7 @@ BT_GPU___device__ void testSphSph(float3 aPos, float3 bPos, float radA, float ra float dist = BT_GPU_dot(del, del); dist = sqrtf(dist); float maxD = radA + radB; + if(dist > maxD) { return; @@ -76,7 +78,9 @@ BT_GPU___device__ void testSphSph(float3 aPos, float3 bPos, float radA, float ra normal = BT_GPU_make_float3(1.f, 0.f, 0.f); } // float3 contact = (bPos + aPos + normal * (radB - radA)) * 0.5f; - float3 contact = aPos - normal * radA; + float3 tmp = (normal * radA); + float3 contact = aPos - tmp; + // now add point int numPoints = 0; for(int i = 0; i < MAX_VTX_PER_OBJ; i++) @@ -143,14 +147,20 @@ BT_GPU___global__ void setConstraintDataD(int2 *constraints, for(i = 0; i < numSphA; i++) { float3 va = aPos; - va += ai * shapeA[i].x; - va += aj * shapeA[i].y; + float3 tmp = ai * shapeA[i].x; + float3 tmp2 = aj * shapeA[i].y; + + va += tmp; + va += tmp2; + float radA = shapeA[i].w; for(j = 0; j < numSphB; j++) { float3 vb = bPos; - vb += bi * shapeB[j].x; - vb += bj * shapeB[j].y; + float3 tmp =bi * shapeB[j].x; + float3 tmp2 = bj * shapeB[j].y; + vb += tmp; + vb += tmp2; float radB = shapeB[j].w; testSphSph(va, vb, radA, radB, pOut); } @@ -174,7 +184,7 @@ BT_GPU___device__ float computeImpulse1(float3 rVel, if(positionConstraint > 0) return lambdaDt; - positionConstraint = min(0.0f,positionConstraint+penetrationError); + positionConstraint = btMin(0.0f,positionConstraint+penetrationError); lambdaDt = -(BT_GPU_dot(cNormal,rVel)*(1+collisionConstant)); lambdaDt -= (baumgarteConstant/dt*positionConstraint); @@ -216,7 +226,8 @@ BT_GPU___global__ void collisionWithWallBoxD(float4 *pos, float3 aVel = BT_GPU_make_float3(vel[idx].x, vel[idx].y, vel[idx].z); float aAngVel = angVel[idx]; float3 rerVertex = ai * shape[iVtx].x; - rerVertex += aj * shape[iVtx].y; + float3 tmp = aj * shape[iVtx].y; + rerVertex += tmp; float3 vPos = aPos + rerVertex; float rad = shape[iVtx].w; float3 vVel =aVel+BT_GPU_cross(BT_GPU_make_float3(0.0f,0.0f,aAngVel),rerVertex); @@ -239,11 +250,14 @@ BT_GPU___global__ void collisionWithWallBoxD(float4 *pos, { lat_vel_len = sqrtf(lat_vel_len); lat_vel *= 1.f/lat_vel_len; - impulse -= lat_vel * BT_GPU_dot(lat_vel, vVel) * FRICTION_BOX_GROUND_FACT; + float3 tmp = lat_vel * BT_GPU_dot(lat_vel, vVel) * FRICTION_BOX_GROUND_FACT; + impulse -= tmp; } #endif //USE_FRICTION - vel[idx] += BT_GPU_make_float42(impulse,0.0f); - angVel[idx] += BT_GPU_cross(rerVertex,impulse).z; + float4 tmp = BT_GPU_make_float42(impulse,0.0f); + vel[idx] += tmp; + float tmp2 = BT_GPU_cross(rerVertex,impulse).z; + angVel[idx] += tmp2; } } @@ -257,7 +271,8 @@ BT_GPU___global__ void collisionWithWallBoxD(float4 *pos, BT_GPU_make_float3(1.0f,0.0f,0.0f), dt); - vel[idx] += BT_GPU_make_float42(impulse,0.0f); + float4 tmp = BT_GPU_make_float42(impulse,0.0f); + vel[idx] += tmp; angVel[idx] += BT_GPU_cross(rerVertex,impulse).z; } } @@ -272,7 +287,8 @@ BT_GPU___global__ void collisionWithWallBoxD(float4 *pos, BT_GPU_make_float3(-1.0f,0.0f,0.0f), dt); - vel[idx] += BT_GPU_make_float42(impulse,0.0f); + float4 tmp = BT_GPU_make_float42(impulse,0.0f); + vel[idx] += tmp; angVel[idx] += BT_GPU_cross(rerVertex,impulse).z; } } @@ -326,7 +342,7 @@ BT_GPU___device__ void collisionResolutionBox( int constrId, { float rLambdaDt=lambdaDtBox[(MAX_VTX_PER_OBJ)*(2*constrId)+iVtx]; float pLambdaDt=rLambdaDt; - rLambdaDt=max(pLambdaDt+lambdaDt,0.0f); + rLambdaDt=btMax(pLambdaDt+lambdaDt,0.0f); lambdaDt=rLambdaDt-pLambdaDt; lambdaDtBox[(MAX_VTX_PER_OBJ)*(2*constrId)+iVtx]=rLambdaDt; } @@ -340,7 +356,8 @@ BT_GPU___device__ void collisionResolutionBox( int constrId, { lat_vel_len = sqrtf(lat_vel_len); lat_vel *= 1.f/lat_vel_len; - impulse -= lat_vel * BT_GPU_dot(lat_vel , relVel) * FRICTION_BOX_BOX_FACT; + float3 tmp = lat_vel * BT_GPU_dot(lat_vel , relVel) * FRICTION_BOX_BOX_FACT; + impulse -= tmp; } } #endif //USE_FRICTION diff --git a/Demos/Jamfile b/Demos/Jamfile index b29045e04..7173c38cb 100644 --- a/Demos/Jamfile +++ b/Demos/Jamfile @@ -81,6 +81,8 @@ SubInclude TOP Demos EPAPenDepthDemo ; SubInclude TOP Demos HelloWorld ; SubInclude TOP Demos BspDemo ; SubInclude TOP Demos BasicDemo ; +SubInclude TOP Demos Box2dDemo ; +SubInclude TOP Demos Gpu2dDemo ; SubInclude TOP Demos ConvexDecompositionDemo ; SubInclude TOP Demos ColladaDemo ; SubInclude TOP Demos CharacterDemo ; diff --git a/Demos/OpenGL/GLDebugFont.cpp b/Demos/OpenGL/GLDebugFont.cpp index 721101275..e9131e270 100644 --- a/Demos/OpenGL/GLDebugFont.cpp +++ b/Demos/OpenGL/GLDebugFont.cpp @@ -16,8 +16,9 @@ subject to the following restrictions: #include "GLDebugFont.h" #include "GlutStuff.h" #include +#include -unsigned char sFontData[]; +extern unsigned char sFontData[]; static GLuint sTexture = -1; static int sScreenWidth = 0; static int sScreenHeight = 0; diff --git a/Demos/OpenGL/GL_ShapeDrawer.cpp b/Demos/OpenGL/GL_ShapeDrawer.cpp index 6a4fa8526..e90cca0dc 100644 --- a/Demos/OpenGL/GL_ShapeDrawer.cpp +++ b/Demos/OpenGL/GL_ShapeDrawer.cpp @@ -368,7 +368,7 @@ void renderSquareA(float x, float y, float z) glEnd(); } -inline void glDrawVector(btVector3& v) { glVertex3d(v[0], v[1], v[2]); } +inline void glDrawVector(const btVector3& v) { glVertex3d(v[0], v[1], v[2]); } void GL_ShapeDrawer::drawOpenGL(btScalar* m, const btCollisionShape* shape, const btVector3& color,int debugMode,const btVector3& worldBoundsMin,const btVector3& worldBoundsMax) diff --git a/mk/msvcgen/control.tlib b/mk/msvcgen/control.tlib index ff85d5305..d6b349466 100644 --- a/mk/msvcgen/control.tlib +++ b/mk/msvcgen/control.tlib @@ -219,7 +219,7 @@ projtypes = { }, library => { - defines = ['_LIB', '_WINDOWS','_CRT_SECURE_NO_DEPRECATE','_CRT_NONSTDC_NO_DEPRECATE'] + defines = ['_LIB', '_CONSOLE','_CRT_SECURE_NO_DEPRECATE','_CRT_NONSTDC_NO_DEPRECATE'] } }; diff --git a/mk/msvcgen/projectx7.tlib b/mk/msvcgen/projectx7.tlib index a2949c67e..04f112571 100644 --- a/mk/msvcgen/projectx7.tlib +++ b/mk/msvcgen/projectx7.tlib @@ -68,13 +68,13 @@ [% MACRO compiler_release BLOCK -%] EnableIntrinsicFunctions="true" - EnableEnhancedInstructionSet="2" + EnableEnhancedInstructionSet="1" FloatingPointModel="2" Optimization="2" StringPooling="TRUE" EnableFunctionLevelLinking="TRUE" RuntimeLibrary="0" - DebugInformationFormat="3" + DebugInformationFormat="1" BufferSecurityCheck="FALSE" [% END %] [% MACRO compiler_release_dll BLOCK -%] @@ -88,21 +88,21 @@ [% MACRO compiler_debug BLOCK -%] Optimization="0" MinimalRebuild="TRUE" - DebugInformationFormat="4" + DebugInformationFormat="3" RuntimeTypeInfo="FALSE" RuntimeLibrary="1" [% END %] [% MACRO compiler_debug_dll BLOCK -%] Optimization="0" MinimalRebuild="TRUE" - DebugInformationFormat="4" + DebugInformationFormat="3" RuntimeTypeInfo="FALSE" RuntimeLibrary="3" [% END %] [% MACRO compiler_debug_dbl BLOCK -%] Optimization="0" MinimalRebuild="TRUE" - DebugInformationFormat="4" + DebugInformationFormat="3" RuntimeTypeInfo="FALSE" RuntimeLibrary="1" [% END %] @@ -110,7 +110,7 @@ LinkIncremental="1" OptimizeReferences="2" EnableCOMDATFolding="2" - GenerateDebugInformation="TRUE" + GenerateDebugInformation="FALSE" IgnoreDefaultLibraryNames="LIBC,LIBCD" [% END %] [% MACRO linker_debug BLOCK -%] diff --git a/mk/msvcgen/projectxsn71.tlib b/mk/msvcgen/projectxsn71.tlib index 4896077b0..c8df68323 100644 --- a/mk/msvcgen/projectxsn71.tlib +++ b/mk/msvcgen/projectxsn71.tlib @@ -77,7 +77,7 @@ [% MACRO compiler_debug BLOCK -%] Optimization="0" MinimalRebuild="TRUE" - DebugInformationFormat="4" + DebugInformationFormat="3" RuntimeTypeInfo="FALSE" RuntimeLibrary="3" [% END %] diff --git a/src/BulletMultiThreaded/btGpu3DGridBroadphase.cpp b/src/BulletMultiThreaded/btGpu3DGridBroadphase.cpp index a18e8806b..3b8838f00 100644 --- a/src/BulletMultiThreaded/btGpu3DGridBroadphase.cpp +++ b/src/BulletMultiThreaded/btGpu3DGridBroadphase.cpp @@ -27,6 +27,7 @@ subject to the following restrictions: #include "btGpu3DGridBroadphaseSharedDefs.h" #include "btGpu3DGridBroadphase.h" +#include //for memset //-------------------------------------------------------------------------- diff --git a/src/BulletMultiThreaded/btGpuDefines.h b/src/BulletMultiThreaded/btGpuDefines.h index 65c7db7ac..c2036a54f 100644 --- a/src/BulletMultiThreaded/btGpuDefines.h +++ b/src/BulletMultiThreaded/btGpuDefines.h @@ -1,221 +1,214 @@ -/* -Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org -Copyright (C) 2006, 2009 Sony Computer Entertainment Inc. - -This software is provided 'as-is', without any express or implied warranty. -In no event will the authors be held liable for any damages arising from the use of this software. -Permission is granted to anyone to use this software for any purpose, -including commercial applications, and to alter it and redistribute it freely, -subject to the following restrictions: - -1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. -2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. -3. This notice may not be removed or altered from any source distribution. -*/ - -//-------------------------------------------------------------------------- - -// definitions for "GPU on CPU" code - -//-------------------------------------------------------------------------- -#ifndef BT_GPU_DEFINES_H -#define BT_GPU_DEFINES_H - -typedef unsigned int uint; - -struct int2 -{ - int x, y; -}; - -struct uint2 -{ - unsigned int x, y; -}; - -struct int3 -{ - int x, y, z; -}; - -struct uint3 -{ - unsigned int x, y, z; -}; - -struct float4 -{ - float x, y, z, w; -}; - -struct float3 -{ - float x, y, z; -}; - - -#define BT_GPU___device__ inline -#define BT_GPU___devdata__ -#define BT_GPU___constant__ -#define BT_GPU_max(a, b) ((a) > (b) ? (a) : (b)) -#define BT_GPU_min(a, b) ((a) < (b) ? (a) : (b)) -#define BT_GPU_params s3DGridBroadphaseParams -#define BT_GPU___mul24(a, b) ((a)*(b)) -#define BT_GPU___global__ inline -#define BT_GPU___shared__ static -#define BT_GPU___syncthreads() -#define CUDART_PI_F SIMD_PI - -static inline uint2 bt3dGrid_make_uint2(unsigned int x, unsigned int y) -{ - uint2 t; t.x = x; t.y = y; return t; -} -#define BT_GPU_make_uint2(x, y) bt3dGrid_make_uint2(x, y) - -static inline int3 bt3dGrid_make_int3(int x, int y, int z) -{ - int3 t; t.x = x; t.y = y; t.z = z; return t; -} -#define BT_GPU_make_int3(x, y, z) bt3dGrid_make_int3(x, y, z) - -static inline float3 bt3dGrid_make_float3(float x, float y, float z) -{ - float3 t; t.x = x; t.y = y; t.z = z; return t; -} -#define BT_GPU_make_float3(x, y, z) bt3dGrid_make_float3(x, y, z) - -static inline float3 bt3dGrid_make_float34(float4 f) -{ - float3 t; t.x = f.x; t.y = f.y; t.z = f.z; return t; -} -#define BT_GPU_make_float34(f) bt3dGrid_make_float34(f) - -static inline float3 bt3dGrid_make_float31(float f) -{ - float3 t; t.x = t.y = t.z = f; return t; -} -#define BT_GPU_make_float31(x) bt3dGrid_make_float31(x) - -static inline float4 bt3dGrid_make_float42(float3 v, float f) -{ - float4 t; t.x = v.x; t.y = v.y; t.z = v.z; t.w = f; return t; -} -#define BT_GPU_make_float42(a, b) bt3dGrid_make_float42(a, b) - -static inline float4 bt3dGrid_make_float44(float a, float b, float c, float d) -{ - float4 t; t.x = a; t.y = b; t.z = c; t.w = d; return t; -} -#define BT_GPU_make_float44(a, b, c, d) bt3dGrid_make_float44(a, b, c, d) - -inline int3 operator+(int3 a, int3 b) -{ - return bt3dGrid_make_int3(a.x + b.x, a.y + b.y, a.z + b.z); -} - -inline float4 operator+(float4& a, float4& b) -{ - float4 r; r.x = a.x+b.x; r.y = a.y+b.y; r.z = a.z+b.z; r.w = a.w+b.w; return r; -} -inline float4 operator*(float4& a, float fact) -{ - float4 r; r.x = a.x*fact; r.y = a.y*fact; r.z = a.z*fact; r.w = a.w*fact; return r; -} -inline float4 operator*(float fact, float4& a) -{ - return (a * fact); -} -inline float4& operator*=(float4& a, float fact) -{ - a = fact * a; - return a; -} -inline float4& operator+=(float4& a, float4& b) -{ - a = a + b; - return a; -} - -inline float3 operator+(float3& a, float3& b) -{ - float3 r; r.x = a.x+b.x; r.y = a.y+b.y; r.z = a.z+b.z; return r; -} -inline float3 operator+(const float3& a, const float3& b) -{ - float3 r; r.x = a.x+b.x; r.y = a.y+b.y; r.z = a.z+b.z; return r; -} -inline float3 operator-(float3& a, float3& b) -{ - float3 r; r.x = a.x-b.x; r.y = a.y-b.y; r.z = a.z-b.z; return r; -} -static inline float bt3dGrid_dot(float3& a, float3& b) -{ - return a.x*b.x+a.y*b.y+a.z*b.z; -} -#define BT_GPU_dot(a,b) bt3dGrid_dot(a,b) - -static inline float bt3dGrid_dot4(float4& a, float4& b) -{ - return a.x*b.x+a.y*b.y+a.z*b.z+a.w*b.w; -} -#define BT_GPU_dot4(a,b) bt3dGrid_dot4(a,b) - -static inline float3 bt3dGrid_cross(float3& a, float3& b) -{ - float3 r; r.x = a.y*b.z-a.z*b.y; r.y = -a.x*b.z+a.z*b.x; r.z = a.x*b.y-a.y*b.x; return r; -} -#define BT_GPU_cross(a,b) bt3dGrid_cross(a,b) - -inline float3 operator*(float3& a, float fact) -{ - float3 r; r.x = a.x*fact; r.y = a.y*fact; r.z = a.z*fact; return r; -} -inline float3 operator*(const float3& a, float fact) -{ - float3 r; r.x = a.x*fact; r.y = a.y*fact; r.z = a.z*fact; return r; -} - - -inline float3& operator+=(float3& a, float3& b) -{ - a = a + b; - return a; -} -inline float3& operator-=(float3& a, float3& b) -{ - a = a - b; - return a; -} -inline float3& operator*=(float3& a, float fact) -{ - a = a * fact; - return a; -} -inline float3 operator-(const float3& v) -{ - float3 r; r.x = -v.x; r.y = -v.y; r.z = -v.z; return r; -} - - -#define BT_GPU_FETCH(a, b) a[b] -#define BT_GPU_FETCH4(a, b) a[b] -#define BT_GPU_PREF(func) btGpu_##func -#define BT_GPU_SAFE_CALL(func) func -#define BT_GPU_Memset memset -#define BT_GPU_MemcpyToSymbol(a, b, c) memcpy(a, b, c) -#define BT_GPU_BindTexture(a, b, c, d) -#define BT_GPU_UnbindTexture(a) - -static uint2 s_blockIdx, s_blockDim, s_threadIdx; -#define BT_GPU_blockIdx s_blockIdx -#define BT_GPU_blockDim s_blockDim -#define BT_GPU_threadIdx s_threadIdx -#define BT_GPU_EXECKERNEL(numb, numt, kfunc, args) {s_blockDim.x=numt;for(int nb=0;nb (b) ? (a) : (b)) +#define BT_GPU_min(a, b) ((a) < (b) ? (a) : (b)) +#define BT_GPU_params s3DGridBroadphaseParams +#define BT_GPU___mul24(a, b) ((a)*(b)) +#define BT_GPU___global__ inline +#define BT_GPU___shared__ static +#define BT_GPU___syncthreads() +#define CUDART_PI_F SIMD_PI + +static inline uint2 bt3dGrid_make_uint2(unsigned int x, unsigned int y) +{ + uint2 t; t.x = x; t.y = y; return t; +} +#define BT_GPU_make_uint2(x, y) bt3dGrid_make_uint2(x, y) + +static inline int3 bt3dGrid_make_int3(int x, int y, int z) +{ + int3 t; t.x = x; t.y = y; t.z = z; return t; +} +#define BT_GPU_make_int3(x, y, z) bt3dGrid_make_int3(x, y, z) + +static inline float3 bt3dGrid_make_float3(float x, float y, float z) +{ + float3 t; t.x = x; t.y = y; t.z = z; return t; +} +#define BT_GPU_make_float3(x, y, z) bt3dGrid_make_float3(x, y, z) + +static inline float3 bt3dGrid_make_float34(float4 f) +{ + float3 t; t.x = f.x; t.y = f.y; t.z = f.z; return t; +} +#define BT_GPU_make_float34(f) bt3dGrid_make_float34(f) + +static inline float3 bt3dGrid_make_float31(float f) +{ + float3 t; t.x = t.y = t.z = f; return t; +} +#define BT_GPU_make_float31(x) bt3dGrid_make_float31(x) + +static inline float4 bt3dGrid_make_float42(float3 v, float f) +{ + float4 t; t.x = v.x; t.y = v.y; t.z = v.z; t.w = f; return t; +} +#define BT_GPU_make_float42(a, b) bt3dGrid_make_float42(a, b) + +static inline float4 bt3dGrid_make_float44(float a, float b, float c, float d) +{ + float4 t; t.x = a; t.y = b; t.z = c; t.w = d; return t; +} +#define BT_GPU_make_float44(a, b, c, d) bt3dGrid_make_float44(a, b, c, d) + +inline int3 operator+(int3 a, int3 b) +{ + return bt3dGrid_make_int3(a.x + b.x, a.y + b.y, a.z + b.z); +} + +inline float4 operator+(const float4& a, const float4& b) +{ + float4 r; r.x = a.x+b.x; r.y = a.y+b.y; r.z = a.z+b.z; r.w = a.w+b.w; return r; +} +inline float4 operator*(const float4& a, float fact) +{ + float4 r; r.x = a.x*fact; r.y = a.y*fact; r.z = a.z*fact; r.w = a.w*fact; return r; +} +inline float4 operator*(float fact, float4& a) +{ + return (a * fact); +} +inline float4& operator*=(float4& a, float fact) +{ + a = fact * a; + return a; +} +inline float4& operator+=(float4& a, float4& b) +{ + a = a + b; + return a; +} + +inline float3 operator+(const float3& a, const float3& b) +{ + float3 r; r.x = a.x+b.x; r.y = a.y+b.y; r.z = a.z+b.z; return r; +} +inline float3 operator-(const float3& a, const float3& b) +{ + float3 r; r.x = a.x-b.x; r.y = a.y-b.y; r.z = a.z-b.z; return r; +} +static inline float bt3dGrid_dot(float3& a, float3& b) +{ + return a.x*b.x+a.y*b.y+a.z*b.z; +} +#define BT_GPU_dot(a,b) bt3dGrid_dot(a,b) + +static inline float bt3dGrid_dot4(float4& a, float4& b) +{ + return a.x*b.x+a.y*b.y+a.z*b.z+a.w*b.w; +} +#define BT_GPU_dot4(a,b) bt3dGrid_dot4(a,b) + +static inline float3 bt3dGrid_cross(const float3& a, const float3& b) +{ + float3 r; r.x = a.y*b.z-a.z*b.y; r.y = -a.x*b.z+a.z*b.x; r.z = a.x*b.y-a.y*b.x; return r; +} +#define BT_GPU_cross(a,b) bt3dGrid_cross(a,b) + + +inline float3 operator*(const float3& a, float fact) +{ + float3 r; r.x = a.x*fact; r.y = a.y*fact; r.z = a.z*fact; return r; +} + + +inline float3& operator+=(float3& a, float3& b) +{ + a = a + b; + return a; +} +inline float3& operator-=(float3& a, float3& b) +{ + a = a - b; + return a; +} +inline float3& operator*=(float3& a, float fact) +{ + a = a * fact; + return a; +} +inline float3 operator-(const float3& v) +{ + float3 r; r.x = -v.x; r.y = -v.y; r.z = -v.z; return r; +} + + +#define BT_GPU_FETCH(a, b) a[b] +#define BT_GPU_FETCH4(a, b) a[b] +#define BT_GPU_PREF(func) btGpu_##func +#define BT_GPU_SAFE_CALL(func) func +#define BT_GPU_Memset memset +#define BT_GPU_MemcpyToSymbol(a, b, c) memcpy(a, b, c) +#define BT_GPU_BindTexture(a, b, c, d) +#define BT_GPU_UnbindTexture(a) + +static uint2 s_blockIdx, s_blockDim, s_threadIdx; +#define BT_GPU_blockIdx s_blockIdx +#define BT_GPU_blockDim s_blockDim +#define BT_GPU_threadIdx s_threadIdx +#define BT_GPU_EXECKERNEL(numb, numt, kfunc, args) {s_blockDim.x=numt;for(int nb=0;nb