Bart/bullet3
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

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

Browse Source
This commit is contained in:
erwin.coumans
2009-05-09 21:25:42 +00:00
parent 9da45d5751
commit 0a360904b6
14 changed files with 312 additions and 251 deletions
Download Patch File Download Diff File Expand all files Collapse all files

32
Demos/Box2dDemo/CMakeLists.txt Normal file
View File

@@ -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
)

3
Demos/Box2dDemo/Jamfile Normal file
View File

@@ -0,0 +1,3 @@
SubDir TOP Demos Box2dDemo ;
BulletDemo Box2dDemo : [ Wildcard *.h *.cpp ] ;

4
Demos/CMakeLists.txt
View File

@@ -1,7 +1,7 @@
if (CMAKE_SIZEOF_VOID_P MATCHES "8") if (CMAKE_SIZEOF_VOID_P MATCHES "8")
SUBDIRS( OpenGL AllBulletDemos ConvexDecompositionDemo Benchmarks HelloWorld SUBDIRS( OpenGL AllBulletDemos ConvexDecompositionDemo Benchmarks HelloWorld
CcdPhysicsDemo ConstraintDemo SliderConstraintDemo GenericJointDemo Raytracer CcdPhysicsDemo ConstraintDemo SliderConstraintDemo GenericJointDemo Raytracer
RagdollDemo ForkLiftDemo BasicDemo BspDemo MovingConcaveDemo VehicleDemo RagdollDemo ForkLiftDemo BasicDemo Box2dDemo Gpu2dDemo BspDemo MovingConcaveDemo VehicleDemo
ColladaDemo UserCollisionAlgorithm CharacterDemo SoftDemo HeightFieldFluidDemo ColladaDemo UserCollisionAlgorithm CharacterDemo SoftDemo HeightFieldFluidDemo
CollisionInterfaceDemo ConcaveConvexcastDemo SimplexDemo DynamicControlDemo CollisionInterfaceDemo ConcaveConvexcastDemo SimplexDemo DynamicControlDemo
DoublePrecisionDemo ConcaveDemo CollisionDemo DoublePrecisionDemo ConcaveDemo CollisionDemo
@@ -10,7 +10,7 @@ MultiMaterialDemo)
else (CMAKE_SIZEOF_VOID_P MATCHES "8") else (CMAKE_SIZEOF_VOID_P MATCHES "8")
SUBDIRS( OpenGL AllBulletDemos ConvexDecompositionDemo Benchmarks HelloWorld SUBDIRS( OpenGL AllBulletDemos ConvexDecompositionDemo Benchmarks HelloWorld
MultiThreadedDemo CcdPhysicsDemo ConstraintDemo SliderConstraintDemo Raytracer MultiThreadedDemo CcdPhysicsDemo ConstraintDemo SliderConstraintDemo Raytracer
GenericJointDemo RagdollDemo ForkLiftDemo BasicDemo BspDemo MovingConcaveDemo GenericJointDemo RagdollDemo ForkLiftDemo BasicDemo Box2dDemo Gpu2dDemo BspDemo MovingConcaveDemo
VehicleDemo ColladaDemo UserCollisionAlgorithm CharacterDemo SoftDemo VehicleDemo ColladaDemo UserCollisionAlgorithm CharacterDemo SoftDemo
HeightFieldFluidDemo HeightFieldFluidDemo
CollisionInterfaceDemo ConcaveConvexcastDemo SimplexDemo DynamicControlDemo CollisionInterfaceDemo ConcaveConvexcastDemo SimplexDemo DynamicControlDemo

16
Demos/Gpu2dDemo/CMakeLists.txt
View File

@@ -19,8 +19,20 @@ LINK_LIBRARIES(
OpenGLSupport BulletDynamics BulletCollision LinearMath ${GLUT_glut_LIBRARY} ${OPENGL_gl_LIBRARY} ${OPENGL_glu_LIBRARY} OpenGLSupport BulletDynamics BulletCollision LinearMath ${GLUT_glut_LIBRARY} ${OPENGL_gl_LIBRARY} ${OPENGL_glu_LIBRARY}
) )
ADD_EXECUTABLE(AppBasicDemo ADD_EXECUTABLE(AppGpu2dDemo
main.cpp 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
) )

4
Demos/Gpu2dDemo/Jamfile
View File

@@ -1,3 +1,3 @@
SubDir TOP Demos BasicDemo ; SubDir TOP Demos Gpu2dDemo ;
BulletDemo BasicDemo : [ Wildcard *.h *.cpp ] ; BulletDemo Gpu2dDemo : [ Wildcard *.h *.cpp ] ;

45
Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h
View File

@@ -19,6 +19,7 @@ subject to the following restrictions:
#define FRICTION_BOX_GROUND_FACT 0.05f #define FRICTION_BOX_GROUND_FACT 0.05f
#define FRICTION_BOX_BOX_FACT 0.05f #define FRICTION_BOX_BOX_FACT 0.05f
#define USE_CENTERS 1 #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); float dist = BT_GPU_dot(del, del);
dist = sqrtf(dist); dist = sqrtf(dist);
float maxD = radA + radB; float maxD = radA + radB;
if(dist > maxD) if(dist > maxD)
{ {
return; 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); normal = BT_GPU_make_float3(1.f, 0.f, 0.f);
} }
// float3 contact = (bPos + aPos + normal * (radB - radA)) * 0.5f; // 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 // now add point
int numPoints = 0; int numPoints = 0;
for(int i = 0; i < MAX_VTX_PER_OBJ; i++) 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++) for(i = 0; i < numSphA; i++)
{ {
float3 va = aPos; float3 va = aPos;
va += ai * shapeA[i].x; float3 tmp = ai * shapeA[i].x;
va += aj * shapeA[i].y; float3 tmp2 = aj * shapeA[i].y;
va += tmp;
va += tmp2;
float radA = shapeA[i].w; float radA = shapeA[i].w;
for(j = 0; j < numSphB; j++) for(j = 0; j < numSphB; j++)
{ {
float3 vb = bPos; float3 vb = bPos;
vb += bi * shapeB[j].x; float3 tmp =bi * shapeB[j].x;
vb += bj * shapeB[j].y; float3 tmp2 = bj * shapeB[j].y;
vb += tmp;
vb += tmp2;
float radB = shapeB[j].w; float radB = shapeB[j].w;
testSphSph(va, vb, radA, radB, pOut); testSphSph(va, vb, radA, radB, pOut);
} }
@@ -174,7 +184,7 @@ BT_GPU___device__ float computeImpulse1(float3 rVel,
if(positionConstraint > 0) if(positionConstraint > 0)
return lambdaDt; return lambdaDt;
positionConstraint = min(0.0f,positionConstraint+penetrationError); positionConstraint = btMin(0.0f,positionConstraint+penetrationError);
lambdaDt = -(BT_GPU_dot(cNormal,rVel)*(1+collisionConstant)); lambdaDt = -(BT_GPU_dot(cNormal,rVel)*(1+collisionConstant));
lambdaDt -= (baumgarteConstant/dt*positionConstraint); 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); float3 aVel = BT_GPU_make_float3(vel[idx].x, vel[idx].y, vel[idx].z);
float aAngVel = angVel[idx]; float aAngVel = angVel[idx];
float3 rerVertex = ai * shape[iVtx].x; float3 rerVertex = ai * shape[iVtx].x;
rerVertex += aj * shape[iVtx].y; float3 tmp = aj * shape[iVtx].y;
rerVertex += tmp;
float3 vPos = aPos + rerVertex; float3 vPos = aPos + rerVertex;
float rad = shape[iVtx].w; float rad = shape[iVtx].w;
float3 vVel =aVel+BT_GPU_cross(BT_GPU_make_float3(0.0f,0.0f,aAngVel),rerVertex); 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_len = sqrtf(lat_vel_len);
lat_vel *= 1.f/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 #endif //USE_FRICTION
vel[idx] += BT_GPU_make_float42(impulse,0.0f); float4 tmp = BT_GPU_make_float42(impulse,0.0f);
angVel[idx] += BT_GPU_cross(rerVertex,impulse).z; 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), BT_GPU_make_float3(1.0f,0.0f,0.0f),
dt); 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; 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), BT_GPU_make_float3(-1.0f,0.0f,0.0f),
dt); 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; 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 rLambdaDt=lambdaDtBox[(MAX_VTX_PER_OBJ)*(2*constrId)+iVtx];
float pLambdaDt=rLambdaDt; float pLambdaDt=rLambdaDt;
rLambdaDt=max(pLambdaDt+lambdaDt,0.0f); rLambdaDt=btMax(pLambdaDt+lambdaDt,0.0f);
lambdaDt=rLambdaDt-pLambdaDt; lambdaDt=rLambdaDt-pLambdaDt;
lambdaDtBox[(MAX_VTX_PER_OBJ)*(2*constrId)+iVtx]=rLambdaDt; 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_len = sqrtf(lat_vel_len);
lat_vel *= 1.f/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 #endif //USE_FRICTION

2
Demos/Jamfile
View File

@@ -81,6 +81,8 @@ SubInclude TOP Demos EPAPenDepthDemo ;
SubInclude TOP Demos HelloWorld ; SubInclude TOP Demos HelloWorld ;
SubInclude TOP Demos BspDemo ; SubInclude TOP Demos BspDemo ;
SubInclude TOP Demos BasicDemo ; SubInclude TOP Demos BasicDemo ;
SubInclude TOP Demos Box2dDemo ;
SubInclude TOP Demos Gpu2dDemo ;
SubInclude TOP Demos ConvexDecompositionDemo ; SubInclude TOP Demos ConvexDecompositionDemo ;
SubInclude TOP Demos ColladaDemo ; SubInclude TOP Demos ColladaDemo ;
SubInclude TOP Demos CharacterDemo ; SubInclude TOP Demos CharacterDemo ;

3
Demos/OpenGL/GLDebugFont.cpp
View File

@@ -16,8 +16,9 @@ subject to the following restrictions:
#include "GLDebugFont.h" #include "GLDebugFont.h"
#include "GlutStuff.h" #include "GlutStuff.h"
#include <stdio.h> #include <stdio.h>
#include <string>
unsigned char sFontData[]; extern unsigned char sFontData[];
static GLuint sTexture = -1; static GLuint sTexture = -1;
static int sScreenWidth = 0; static int sScreenWidth = 0;
static int sScreenHeight = 0; static int sScreenHeight = 0;

2
Demos/OpenGL/GL_ShapeDrawer.cpp
View File

@@ -368,7 +368,7 @@ void renderSquareA(float x, float y, float z)
glEnd(); 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) void GL_ShapeDrawer::drawOpenGL(btScalar* m, const btCollisionShape* shape, const btVector3& color,int debugMode,const btVector3& worldBoundsMin,const btVector3& worldBoundsMax)

2
mk/msvcgen/control.tlib
View File

@@ -219,7 +219,7 @@ projtypes = {
}, },
library => library =>
{ {
defines = ['_LIB', '_WINDOWS','_CRT_SECURE_NO_DEPRECATE','_CRT_NONSTDC_NO_DEPRECATE'] defines = ['_LIB', '_CONSOLE','_CRT_SECURE_NO_DEPRECATE','_CRT_NONSTDC_NO_DEPRECATE']
} }
}; };

12
mk/msvcgen/projectx7.tlib
View File

@@ -68,13 +68,13 @@
[% MACRO compiler_release BLOCK -%] [% MACRO compiler_release BLOCK -%]
EnableIntrinsicFunctions="true" EnableIntrinsicFunctions="true"
EnableEnhancedInstructionSet="2" EnableEnhancedInstructionSet="1"
FloatingPointModel="2" FloatingPointModel="2"
Optimization="2" Optimization="2"
StringPooling="TRUE" StringPooling="TRUE"
EnableFunctionLevelLinking="TRUE" EnableFunctionLevelLinking="TRUE"
RuntimeLibrary="0" RuntimeLibrary="0"
DebugInformationFormat="3" DebugInformationFormat="1"
BufferSecurityCheck="FALSE" BufferSecurityCheck="FALSE"
[% END %] [% END %]
[% MACRO compiler_release_dll BLOCK -%] [% MACRO compiler_release_dll BLOCK -%]
@@ -88,21 +88,21 @@
[% MACRO compiler_debug BLOCK -%] [% MACRO compiler_debug BLOCK -%]
Optimization="0" Optimization="0"
MinimalRebuild="TRUE" MinimalRebuild="TRUE"
DebugInformationFormat="4" DebugInformationFormat="3"
RuntimeTypeInfo="FALSE" RuntimeTypeInfo="FALSE"
RuntimeLibrary="1" RuntimeLibrary="1"
[% END %] [% END %]
[% MACRO compiler_debug_dll BLOCK -%] [% MACRO compiler_debug_dll BLOCK -%]
Optimization="0" Optimization="0"
MinimalRebuild="TRUE" MinimalRebuild="TRUE"
DebugInformationFormat="4" DebugInformationFormat="3"
RuntimeTypeInfo="FALSE" RuntimeTypeInfo="FALSE"
RuntimeLibrary="3" RuntimeLibrary="3"
[% END %] [% END %]
[% MACRO compiler_debug_dbl BLOCK -%] [% MACRO compiler_debug_dbl BLOCK -%]
Optimization="0" Optimization="0"
MinimalRebuild="TRUE" MinimalRebuild="TRUE"
DebugInformationFormat="4" DebugInformationFormat="3"
RuntimeTypeInfo="FALSE" RuntimeTypeInfo="FALSE"
RuntimeLibrary="1" RuntimeLibrary="1"
[% END %] [% END %]
@@ -110,7 +110,7 @@
LinkIncremental="1" LinkIncremental="1"
OptimizeReferences="2" OptimizeReferences="2"
EnableCOMDATFolding="2" EnableCOMDATFolding="2"
GenerateDebugInformation="TRUE" GenerateDebugInformation="FALSE"
IgnoreDefaultLibraryNames="LIBC,LIBCD" IgnoreDefaultLibraryNames="LIBC,LIBCD"
[% END %] [% END %]
[% MACRO linker_debug BLOCK -%] [% MACRO linker_debug BLOCK -%]

2
mk/msvcgen/projectxsn71.tlib
View File

@@ -77,7 +77,7 @@
[% MACRO compiler_debug BLOCK -%] [% MACRO compiler_debug BLOCK -%]
Optimization="0" Optimization="0"
MinimalRebuild="TRUE" MinimalRebuild="TRUE"
DebugInformationFormat="4" DebugInformationFormat="3"
RuntimeTypeInfo="FALSE" RuntimeTypeInfo="FALSE"
RuntimeLibrary="3" RuntimeLibrary="3"
[% END %] [% END %]

1
src/BulletMultiThreaded/btGpu3DGridBroadphase.cpp
View File

@@ -27,6 +27,7 @@ subject to the following restrictions:
#include "btGpu3DGridBroadphaseSharedDefs.h" #include "btGpu3DGridBroadphaseSharedDefs.h"
#include "btGpu3DGridBroadphase.h" #include "btGpu3DGridBroadphase.h"
#include <string> //for memset
//-------------------------------------------------------------------------- //--------------------------------------------------------------------------

435
src/BulletMultiThreaded/btGpuDefines.h
View File

@@ -1,221 +1,214 @@
/* /*
Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org
Copyright (C) 2006, 2009 Sony Computer Entertainment Inc. Copyright (C) 2006, 2009 Sony Computer Entertainment Inc.
This software is provided 'as-is', without any express or implied warranty. 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. 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, Permission is granted to anyone to use this software for any purpose,
including commercial applications, and to alter it and redistribute it freely, including commercial applications, and to alter it and redistribute it freely,
subject to the following restrictions: 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. 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. 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. 3. This notice may not be removed or altered from any source distribution.
*/ */
//-------------------------------------------------------------------------- //--------------------------------------------------------------------------
// definitions for "GPU on CPU" code // definitions for "GPU on CPU" code
//-------------------------------------------------------------------------- //--------------------------------------------------------------------------
#ifndef BT_GPU_DEFINES_H #ifndef BT_GPU_DEFINES_H
#define BT_GPU_DEFINES_H #define BT_GPU_DEFINES_H
typedef unsigned int uint; typedef unsigned int uint;
struct int2 struct int2
{ {
int x, y; int x, y;
}; };
struct uint2 struct uint2
{ {
unsigned int x, y; unsigned int x, y;
}; };
struct int3 struct int3
{ {
int x, y, z; int x, y, z;
}; };
struct uint3 struct uint3
{ {
unsigned int x, y, z; unsigned int x, y, z;
}; };
struct float4 struct float4
{ {
float x, y, z, w; float x, y, z, w;
}; };
struct float3 struct float3
{ {
float x, y, z; float x, y, z;
}; };
#define BT_GPU___device__ inline #define BT_GPU___device__ inline
#define BT_GPU___devdata__ #define BT_GPU___devdata__
#define BT_GPU___constant__ #define BT_GPU___constant__
#define BT_GPU_max(a, b) ((a) > (b) ? (a) : (b)) #define BT_GPU_max(a, b) ((a) > (b) ? (a) : (b))
#define BT_GPU_min(a, b) ((a) < (b) ? (a) : (b)) #define BT_GPU_min(a, b) ((a) < (b) ? (a) : (b))
#define BT_GPU_params s3DGridBroadphaseParams #define BT_GPU_params s3DGridBroadphaseParams
#define BT_GPU___mul24(a, b) ((a)*(b)) #define BT_GPU___mul24(a, b) ((a)*(b))
#define BT_GPU___global__ inline #define BT_GPU___global__ inline
#define BT_GPU___shared__ static #define BT_GPU___shared__ static
#define BT_GPU___syncthreads() #define BT_GPU___syncthreads()
#define CUDART_PI_F SIMD_PI #define CUDART_PI_F SIMD_PI
static inline uint2 bt3dGrid_make_uint2(unsigned int x, unsigned int y) static inline uint2 bt3dGrid_make_uint2(unsigned int x, unsigned int y)
{ {
uint2 t; t.x = x; t.y = y; return t; uint2 t; t.x = x; t.y = y; return t;
} }
#define BT_GPU_make_uint2(x, y) bt3dGrid_make_uint2(x, y) #define BT_GPU_make_uint2(x, y) bt3dGrid_make_uint2(x, y)
static inline int3 bt3dGrid_make_int3(int x, int y, int z) 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; 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) #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) 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; 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) #define BT_GPU_make_float3(x, y, z) bt3dGrid_make_float3(x, y, z)
static inline float3 bt3dGrid_make_float34(float4 f) static inline float3 bt3dGrid_make_float34(float4 f)
{ {
float3 t; t.x = f.x; t.y = f.y; t.z = f.z; return t; 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) #define BT_GPU_make_float34(f) bt3dGrid_make_float34(f)
static inline float3 bt3dGrid_make_float31(float f) static inline float3 bt3dGrid_make_float31(float f)
{ {
float3 t; t.x = t.y = t.z = f; return t; float3 t; t.x = t.y = t.z = f; return t;
} }
#define BT_GPU_make_float31(x) bt3dGrid_make_float31(x) #define BT_GPU_make_float31(x) bt3dGrid_make_float31(x)
static inline float4 bt3dGrid_make_float42(float3 v, float f) 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; 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) #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) 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; 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) #define BT_GPU_make_float44(a, b, c, d) bt3dGrid_make_float44(a, b, c, d)
inline int3 operator+(int3 a, int3 b) inline int3 operator+(int3 a, int3 b)
{ {
return bt3dGrid_make_int3(a.x + b.x, a.y + b.y, a.z + b.z); return bt3dGrid_make_int3(a.x + b.x, a.y + b.y, a.z + b.z);
} }
inline float4 operator+(float4& a, float4& b) 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; 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) 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; 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) inline float4 operator*(float fact, float4& a)
{ {
return (a * fact); return (a * fact);
} }
inline float4& operator*=(float4& a, float fact) inline float4& operator*=(float4& a, float fact)
{ {
a = fact * a; a = fact * a;
return a; return a;
} }
inline float4& operator+=(float4& a, float4& b) inline float4& operator+=(float4& a, float4& b)
{ {
a = a + b; a = a + b;
return a; return a;
} }
inline float3 operator+(float3& a, float3& b) 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; 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) 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; 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) static inline float bt3dGrid_dot(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; return a.x*b.x+a.y*b.y+a.z*b.z;
} }
static inline float bt3dGrid_dot(float3& a, float3& b) #define BT_GPU_dot(a,b) bt3dGrid_dot(a,b)
{
return a.x*b.x+a.y*b.y+a.z*b.z; static inline float bt3dGrid_dot4(float4& a, float4& b)
} {
#define BT_GPU_dot(a,b) bt3dGrid_dot(a,b) return a.x*b.x+a.y*b.y+a.z*b.z+a.w*b.w;
}
static inline float bt3dGrid_dot4(float4& a, float4& b) #define BT_GPU_dot4(a,b) bt3dGrid_dot4(a,b)
{
return a.x*b.x+a.y*b.y+a.z*b.z+a.w*b.w; static inline float3 bt3dGrid_cross(const float3& a, const float3& b)
} {
#define BT_GPU_dot4(a,b) bt3dGrid_dot4(a,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;
}
static inline float3 bt3dGrid_cross(float3& a, float3& b) #define BT_GPU_cross(a,b) bt3dGrid_cross(a,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;
} inline float3 operator*(const float3& a, float fact)
#define BT_GPU_cross(a,b) bt3dGrid_cross(a,b) {
float3 r; r.x = a.x*fact; r.y = a.y*fact; r.z = a.z*fact; return r;
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+=(float3& a, float3& b)
inline float3 operator*(const float3& a, float fact) {
{ a = a + b;
float3 r; r.x = a.x*fact; r.y = a.y*fact; r.z = a.z*fact; return r; return a;
} }
inline float3& operator-=(float3& a, float3& b)
{
inline float3& operator+=(float3& a, float3& b) a = a - b;
{ return a;
a = a + b; }
return a; inline float3& operator*=(float3& a, float fact)
} {
inline float3& operator-=(float3& a, float3& b) a = a * fact;
{ return a;
a = a - b; }
return a; inline float3 operator-(const float3& v)
} {
inline float3& operator*=(float3& a, float fact) float3 r; r.x = -v.x; r.y = -v.y; r.z = -v.z; return r;
{ }
a = a * fact;
return a;
} #define BT_GPU_FETCH(a, b) a[b]
inline float3 operator-(const float3& v) #define BT_GPU_FETCH4(a, b) a[b]
{ #define BT_GPU_PREF(func) btGpu_##func
float3 r; r.x = -v.x; r.y = -v.y; r.z = -v.z; return r; #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_FETCH(a, b) a[b] #define BT_GPU_UnbindTexture(a)
#define BT_GPU_FETCH4(a, b) a[b]
#define BT_GPU_PREF(func) btGpu_##func static uint2 s_blockIdx, s_blockDim, s_threadIdx;
#define BT_GPU_SAFE_CALL(func) func #define BT_GPU_blockIdx s_blockIdx
#define BT_GPU_Memset memset #define BT_GPU_blockDim s_blockDim
#define BT_GPU_MemcpyToSymbol(a, b, c) memcpy(a, b, c) #define BT_GPU_threadIdx s_threadIdx
#define BT_GPU_BindTexture(a, b, c, d) #define BT_GPU_EXECKERNEL(numb, numt, kfunc, args) {s_blockDim.x=numt;for(int nb=0;nb<numb;nb++){s_blockIdx.x=nb;for(int nt=0;nt<numt;nt++){s_threadIdx.x=nt;kfunc args;}}}
#define BT_GPU_UnbindTexture(a)
#define BT_GPU_CHECK_ERROR(s)
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<numb;nb++){s_blockIdx.x=nb;for(int nt=0;nt<numt;nt++){s_threadIdx.x=nt;kfunc args;}}}
#endif //BT_GPU_DEFINES_H
#define BT_GPU_CHECK_ERROR(s)
//--------------------------------------------------------------------------
//--------------------------------------------------------------------------
//--------------------------------------------------------------------------
#endif //BT_GPU_DEFINES_H