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

New version of btCudaBroadphase compatible with Bullet and better performance

Browse Source
This commit is contained in:
rponom
2008-11-15 00:21:09 +00:00
parent 8abddb2400
commit afec653726
17 changed files with 2739 additions and 2412 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1820
Extras/CUDA/btCudaBroadphase.cpp
View File

File diff suppressed because it is too large Load Diff

589
Extras/CUDA/btCudaBroadphase.cu Normal file
View File

@@ -0,0 +1,589 @@
/*
* Copyright 1993-2006 NVIDIA Corporation. All rights reserved.
*
* NOTICE TO USER:
*
* This source code is subject to NVIDIA ownership rights under U.S. and
* international Copyright laws.
*
* NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
* CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
* IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH
* REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
* MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
* IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL,
* OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS
* OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE
* OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE
* OR PERFORMANCE OF THIS SOURCE CODE.
*
* U.S. Government End Users. This source code is a "commercial item" as
* that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of
* "commercial computer software" and "commercial computer software
* documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995)
* and is provided to the U.S. Government only as a commercial end item.
* Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
* 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
* source code with only those rights set forth herein.
*/
#include <cstdlib>
#include <cstdio>
#include <string.h>
#include "cutil_math.h"
#include "math_constants.h"
#if defined(__APPLE__) || defined(MACOSX)
#include <GLUT/glut.h>
#else
#include <GL/glut.h>
#endif
#include <cuda_gl_interop.h>
#include "btCudaBroadphaseKernel.h"
//#include "radixsort.cu"
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
// K E R N E L F U N C T I O N S
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
#ifdef __DEVICE_EMULATION__
#define B_CUDA_USE_TEX 0
#else
#define B_CUDA_USE_TEX 1
#endif
__device__ inline btCuda3F1U tex_fetch3F1U(float4 a) { return *((btCuda3F1U*)(&a)); }
#if B_CUDA_USE_TEX
#define FETCH(t, i) tex_fetch3F1U(tex1Dfetch(t##Tex, i))
#else
#define FETCH(t, i) t[i]
#endif
texture<uint2, 1, cudaReadModeElementType> particleHashTex;
texture<uint, 1, cudaReadModeElementType> cellStartTex;
texture<float4, 1, cudaReadModeElementType> pAABBTex;
//----------------------------------------------------------------------------------------
__constant__ btCudaBroadphaseParams params;
//----------------------------------------------------------------------------------------
// calculate position in uniform grid
__device__ int3 btCuda_calcGridPos(float4 p)
{
int3 gridPos;
gridPos.x = floor((p.x - params.m_worldOriginX) / params.m_cellSizeX);
gridPos.y = floor((p.y - params.m_worldOriginY) / params.m_cellSizeY);
gridPos.z = floor((p.z - params.m_worldOriginZ) / params.m_cellSizeZ);
return gridPos;
}
//----------------------------------------------------------------------------------------
// calculate address in grid from position (clamping to edges)
__device__ uint btCuda_calcGridHash(int3 gridPos)
{
gridPos.x = max(0, min(gridPos.x, params.m_gridSizeX - 1));
gridPos.y = max(0, min(gridPos.y, params.m_gridSizeY - 1));
gridPos.z = max(0, min(gridPos.z, params.m_gridSizeZ - 1));
return __mul24(__mul24(gridPos.z, params.m_gridSizeY), params.m_gridSizeX) + __mul24(gridPos.y, params.m_gridSizeX) + gridPos.x;
}
//----------------------------------------------------------------------------------------
// calculate grid hash value for each body using its AABB
__global__ void calcHashAABBD(btCuda3F1U* pAABB, uint2* pHash, uint numBodies)
{
int index = __mul24(blockIdx.x, blockDim.x) + threadIdx.x;
if(index >= numBodies)
{
return;
}
btCuda3F1U bbMin = pAABB[index*2];
btCuda3F1U bbMax = pAABB[index*2 + 1];
float4 pos;
pos.x = (bbMin.fx + bbMax.fx) * 0.5f;
pos.y = (bbMin.fy + bbMax.fy) * 0.5f;
pos.z = (bbMin.fz + bbMax.fz) * 0.5f;
// get address in grid
int3 gridPos = btCuda_calcGridPos(pos);
uint gridHash = btCuda_calcGridHash(gridPos);
// store grid hash and body index
pHash[index] = make_uint2(gridHash, index);
}
//----------------------------------------------------------------------------------------
__global__ void findCellStartD(uint2* pHash, uint* cellStart, uint numBodies)
{
int index = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
if(index >= numBodies)
{
return;
}
uint2 sortedData = pHash[index];
// Load hash data into shared memory so that we can look
// at neighboring body's hash value without loading
// two hash values per thread
__shared__ uint sharedHash[257];
sharedHash[threadIdx.x+1] = sortedData.x;
if((index > 0) && (threadIdx.x == 0))
{
// first thread in block must load neighbor body hash
volatile uint2 prevData = pHash[index-1];
sharedHash[0] = prevData.x;
}
__syncthreads();
if((index == 0) || (sortedData.x != sharedHash[threadIdx.x]))
{
cellStart[sortedData.x] = index;
}
}
//----------------------------------------------------------------------------------------
__device__ uint cudaTestAABBOverlap(btCuda3F1U min0, btCuda3F1U max0, btCuda3F1U min1, btCuda3F1U max1)
{
return (min0.fx <= max1.fx)&& (min1.fx <= max0.fx) &&
(min0.fy <= max1.fy)&& (min1.fy <= max0.fy) &&
(min0.fz <= max1.fz)&& (min1.fz <= max0.fz);
}
//----------------------------------------------------------------------------------------
__device__ void findPairsInCell(int3 gridPos,
uint index,
uint2* pHash,
uint* pCellStart,
btCuda3F1U* pAABB,
uint* pPairBuff,
uint2* pPairBuffStartCurr,
uint numBodies)
{
if ( (gridPos.x < 0) || (gridPos.x > params.m_gridSizeX - 1)
|| (gridPos.y < 0) || (gridPos.y > params.m_gridSizeY - 1)
|| (gridPos.z < 0) || (gridPos.z > params.m_gridSizeZ - 1))
{
return;
}
uint gridHash = btCuda_calcGridHash(gridPos);
// get start of bucket for this cell
uint bucketStart = pCellStart[gridHash];
if (bucketStart == 0xffffffff)
{
return; // cell empty
}
// iterate over bodies in this cell
uint2 sortedData = pHash[index];
uint unsorted_indx = sortedData.y;
btCuda3F1U min0 = FETCH(pAABB, unsorted_indx*2);
btCuda3F1U max0 = FETCH(pAABB, unsorted_indx*2 + 1);
uint handleIndex = min0.uw;
uint2 start_curr = pPairBuffStartCurr[handleIndex];
uint start = start_curr.x;
uint curr = start_curr.y;
uint2 start_curr_next = pPairBuffStartCurr[handleIndex+1];
uint curr_max = start_curr_next.x - start - 1;
uint bucketEnd = bucketStart + params.m_maxBodiesPerCell;
bucketEnd = (bucketEnd > numBodies) ? numBodies : bucketEnd;
for(uint index2 = bucketStart; index2 < bucketEnd; index2++)
{
uint2 cellData = pHash[index2];
if (cellData.x != gridHash)
{
break; // no longer in same bucket
}
uint unsorted_indx2 = cellData.y;
if (unsorted_indx2 < unsorted_indx) // check not colliding with self
{
btCuda3F1U min1 = FETCH(pAABB, unsorted_indx2*2);
btCuda3F1U max1 = FETCH(pAABB, unsorted_indx2*2 + 1);
if(cudaTestAABBOverlap(min0, max0, min1, max1))
{
uint handleIndex2 = min1.uw;
uint k;
for(k = 0; k < curr; k++)
{
uint old_pair = pPairBuff[start+k] & (~BT_CUDA_PAIR_ANY_FLG);
if(old_pair == handleIndex2)
{
pPairBuff[start+k] |= BT_CUDA_PAIR_FOUND_FLG;
break;
}
}
if(k == curr)
{
pPairBuff[start+curr] = handleIndex2 | BT_CUDA_PAIR_NEW_FLG;
if(curr >= curr_max)
{ // not a good solution, but let's avoid crash
break;
}
curr++;
}
}
}
}
pPairBuffStartCurr[handleIndex] = make_uint2(start, curr);
return;
}
//----------------------------------------------------------------------------------------
__global__ void
findOverlappingPairsD( btCuda3F1U* pAABB, uint2* pHash, uint* pCellStart, uint* pPairBuff,
uint2* pPairBuffStartCurr, uint numBodies)
{
int index = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
if(index >= numBodies)
{
return;
}
uint2 sortedData = pHash[index];
uint unsorted_indx = sortedData.y;
btCuda3F1U bbMin = FETCH(pAABB, unsorted_indx*2);
btCuda3F1U bbMax = FETCH(pAABB, unsorted_indx*2 + 1);
float4 pos;
pos.x = (bbMin.fx + bbMax.fx) * 0.5f;
pos.y = (bbMin.fy + bbMax.fy) * 0.5f;
pos.z = (bbMin.fz + bbMax.fz) * 0.5f;
// get address in grid
int3 gridPos = btCuda_calcGridPos(pos);
// examine only neighbouring cells
for(int z=-1; z<=1; z++) {
for(int y=-1; y<=1; y++) {
for(int x=-1; x<=1; x++) {
findPairsInCell(gridPos + make_int3(x, y, z), index, pHash, pCellStart, pAABB, pPairBuff, pPairBuffStartCurr, numBodies);
}
}
}
}
//----------------------------------------------------------------------------------------
__global__ void
findPairsLargeD( btCuda3F1U* pAABB, uint2* pHash, uint* pCellStart, uint* pPairBuff,
uint2* pPairBuffStartCurr, uint numBodies, uint numLarge)
{
int index = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
if(index >= numBodies)
{
return;
}
uint2 sortedData = pHash[index];
uint unsorted_indx = sortedData.y;
btCuda3F1U min0 = FETCH(pAABB, unsorted_indx*2);
btCuda3F1U max0 = FETCH(pAABB, unsorted_indx*2 + 1);
uint handleIndex = min0.uw;
uint2 start_curr = pPairBuffStartCurr[handleIndex];
uint start = start_curr.x;
uint curr = start_curr.y;
uint2 start_curr_next = pPairBuffStartCurr[handleIndex+1];
uint curr_max = start_curr_next.x - start - 1;
for(uint i = 0; i < numLarge; i++)
{
uint indx2 = numBodies + i;
btCuda3F1U min1 = FETCH(pAABB, indx2*2);
btCuda3F1U max1 = FETCH(pAABB, indx2*2 + 1);
if(cudaTestAABBOverlap(min0, max0, min1, max1))
{
uint k;
uint handleIndex2 = min1.uw;
for(k = 0; k < curr; k++)
{
uint old_pair = pPairBuff[start+k] & (~BT_CUDA_PAIR_ANY_FLG);
if(old_pair == handleIndex2)
{
pPairBuff[start+k] |= BT_CUDA_PAIR_FOUND_FLG;
break;
}
}
if(k == curr)
{
pPairBuff[start+curr] = handleIndex2 | BT_CUDA_PAIR_NEW_FLG;
if(curr >= curr_max)
{ // not a good solution, but let's avoid crash
break;
}
curr++;
}
}
}
pPairBuffStartCurr[handleIndex] = make_uint2(start, curr);
return;
}
//----------------------------------------------------------------------------------------
__global__ void computePairCacheChangesD(uint* pPairBuff, uint2* pPairBuffStartCurr, uint* pPairScan, btCuda3F1U* pAABB, uint numBodies)
{
int index = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
if(index >= numBodies)
{
return;
}
btCuda3F1U bbMin = pAABB[index * 2];
uint handleIndex = bbMin.uw;
uint2 start_curr = pPairBuffStartCurr[handleIndex];
uint start = start_curr.x;
uint curr = start_curr.y;
uint *pInp = pPairBuff + start;
uint num_changes = 0;
for(uint k = 0; k < curr; k++, pInp++)
{
if(!((*pInp) & BT_CUDA_PAIR_FOUND_FLG))
{
num_changes++;
}
}
pPairScan[index+1] = num_changes;
}
//----------------------------------------------------------------------------------------
__global__ void squeezeOverlappingPairBuffD(uint* pPairBuff, uint2* pPairBuffStartCurr, uint* pPairScan, uint* pPairOut, btCuda3F1U* pAABB, uint numBodies)
{
int index = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
if(index >= numBodies)
{
return;
}
btCuda3F1U bbMin = pAABB[index * 2];
uint handleIndex = bbMin.uw;
uint2 start_curr = pPairBuffStartCurr[handleIndex];
uint start = start_curr.x;
uint curr = start_curr.y;
uint* pInp = pPairBuff + start;
uint* pOut = pPairOut + pPairScan[index];
uint* pOut2 = pInp;
uint num = 0;
for(uint k = 0; k < curr; k++, pInp++)
{
if(!((*pInp) & BT_CUDA_PAIR_FOUND_FLG))
{
*pOut = *pInp;
pOut++;
}
if((*pInp) & BT_CUDA_PAIR_ANY_FLG)
{
*pOut2 = (*pInp) & (~BT_CUDA_PAIR_ANY_FLG);
pOut2++;
num++;
}
}
pPairBuffStartCurr[handleIndex] = make_uint2(start, num);
} // squeezeOverlappingPairBuffD()
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
// E N D O F K E R N E L F U N C T I O N S
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
//----------------------------------------------------------------------------------------
//! Check for CUDA error
# define CUT_CHECK_ERROR(errorMessage) do { \
cudaError_t err = cudaGetLastError(); \
if( cudaSuccess != err) { \
fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n", \
errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
btCuda_exit(EXIT_FAILURE); \
} \
err = cudaThreadSynchronize(); \
if( cudaSuccess != err) { \
fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n", \
errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
btCuda_exit(EXIT_FAILURE); \
} } while (0)
# define MY_CUDA_SAFE_CALL_NO_SYNC( call) do { \
cudaError err = call; \
if( cudaSuccess != err) { \
fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n", \
__FILE__, __LINE__, cudaGetErrorString( err) ); \
btCuda_exit(EXIT_FAILURE); \
} } while (0)
# define MY_CUDA_SAFE_CALL( call) do { \
MY_CUDA_SAFE_CALL_NO_SYNC(call); \
cudaError err = cudaThreadSynchronize(); \
if( cudaSuccess != err) { \
fprintf(stderr, "Cuda errorSync in file '%s' in line %i : %s.\n", \
__FILE__, __LINE__, cudaGetErrorString( err) ); \
btCuda_exit(EXIT_FAILURE); \
} } while (0)
extern "C"
{
void btCuda_exit(int val)
{
exit(val);
}
void btCuda_allocateArray(void** devPtr, unsigned int size)
{
MY_CUDA_SAFE_CALL(cudaMalloc(devPtr, size));
}
void btCuda_freeArray(void* devPtr)
{
MY_CUDA_SAFE_CALL(cudaFree(devPtr));
}
void btCuda_copyArrayFromDevice(void* host, const void* device, unsigned int size)
{
MY_CUDA_SAFE_CALL(cudaMemcpy(host, device, size, cudaMemcpyDeviceToHost));
}
void btCuda_copyArrayToDevice(void* device, const void* host, unsigned int size)
{
MY_CUDA_SAFE_CALL(cudaMemcpy((char*)device, host, size, cudaMemcpyHostToDevice));
}
void btCuda_setParameters(btCudaBroadphaseParams* hostParams)
{
// copy parameters to constant memory
MY_CUDA_SAFE_CALL(cudaMemcpyToSymbol(params, hostParams, sizeof(btCudaBroadphaseParams)));
}
//Round a / b to nearest higher integer value
int btCuda_iDivUp(int a, int b)
{
return (a % b != 0) ? (a / b + 1) : (a / b);
}
// compute grid and thread block size for a given number of elements
void btCuda_computeGridSize(int n, int blockSize, int &numBlocks, int &numThreads)
{
numThreads = min(blockSize, n);
numBlocks = btCuda_iDivUp(n, numThreads);
}
void btCuda_calcHashAABB(btCuda3F1U* pAABB, unsigned int* hash, unsigned int numBodies)
{
int numThreads, numBlocks;
btCuda_computeGridSize(numBodies, 256, numBlocks, numThreads);
// execute the kernel
calcHashAABBD<<< numBlocks, numThreads >>>(pAABB, (uint2*)hash, numBodies);
// check if kernel invocation generated an error
CUT_CHECK_ERROR("calcHashAABBD kernel execution failed");
}
void btCuda_findCellStart(unsigned int* hash, unsigned int* cellStart, unsigned int numBodies, unsigned int numCells)
{
int numThreads, numBlocks;
btCuda_computeGridSize(numBodies, 256, numBlocks, numThreads);
MY_CUDA_SAFE_CALL(cudaMemset(cellStart, 0xffffffff, numCells*sizeof(uint)));
findCellStartD<<< numBlocks, numThreads >>>((uint2*)hash, (uint*)cellStart, numBodies);
CUT_CHECK_ERROR("Kernel execution failed: findCellStartD");
}
void btCuda_findOverlappingPairs( btCuda3F1U* pAABB, unsigned int* pHash,
unsigned int* pCellStart,
unsigned int* pPairBuff,
unsigned int* pPairBuffStartCurr,
unsigned int numBodies)
{
#if B_CUDA_USE_TEX
MY_CUDA_SAFE_CALL(cudaBindTexture(0, pAABBTex, pAABB, numBodies * 2 * sizeof(btCuda3F1U)));
#endif
int numThreads, numBlocks;
btCuda_computeGridSize(numBodies, 64, numBlocks, numThreads);
findOverlappingPairsD<<< numBlocks, numThreads >>>(
pAABB,
(uint2*)pHash,
(uint*)pCellStart,
(uint*)pPairBuff,
(uint2*)pPairBuffStartCurr,
numBodies
);
CUT_CHECK_ERROR("Kernel execution failed: bt_CudaFindOverlappingPairsD");
#if B_CUDA_USE_TEX
MY_CUDA_SAFE_CALL(cudaUnbindTexture(pAABBTex));
#endif
} // btCuda_findOverlappingPairs()
void btCuda_findPairsLarge( btCuda3F1U* pAABB, unsigned int* pHash,
unsigned int* pCellStart,
unsigned int* pPairBuff,
unsigned int* pPairBuffStartCurr,
unsigned int numBodies,
unsigned int numLarge)
{
#if B_CUDA_USE_TEX
MY_CUDA_SAFE_CALL(cudaBindTexture(0, pAABBTex, pAABB, (numBodies+numLarge) * 2 * sizeof(btCuda3F1U)));
#endif
int numThreads, numBlocks;
btCuda_computeGridSize(numBodies, 64, numBlocks, numThreads);
findPairsLargeD<<< numBlocks, numThreads >>>(
pAABB,
(uint2*)pHash,
(uint*)pCellStart,
(uint*)pPairBuff,
(uint2*)pPairBuffStartCurr,
numBodies,
numLarge
);
CUT_CHECK_ERROR("Kernel execution failed: btCuda_findPairsLargeD");
#if B_CUDA_USE_TEX
MY_CUDA_SAFE_CALL(cudaUnbindTexture(pAABBTex));
#endif
} // btCuda_findPairsLarge()
void btCuda_computePairCacheChanges(unsigned int* pPairBuff, unsigned int* pPairBuffStartCurr,
unsigned int* pPairScan, btCuda3F1U* pAABB, unsigned int numBodies)
{
int numThreads, numBlocks;
btCuda_computeGridSize(numBodies, 256, numBlocks, numThreads);
computePairCacheChangesD<<< numBlocks, numThreads >>>(
(uint*)pPairBuff,
(uint2*)pPairBuffStartCurr,
(uint*)pPairScan,
pAABB,
numBodies
);
CUT_CHECK_ERROR("Kernel execution failed: btCudaComputePairCacheChangesD");
} // btCuda_computePairCacheChanges()
void btCuda_squeezeOverlappingPairBuff( unsigned int* pPairBuff, unsigned int* pPairBuffStartCurr, unsigned int* pPairScan,
unsigned int* pPairOut, btCuda3F1U* pAABB, unsigned int numBodies)
{
int numThreads, numBlocks;
btCuda_computeGridSize(numBodies, 256, numBlocks, numThreads);
squeezeOverlappingPairBuffD<<< numBlocks, numThreads >>>(
(uint*)pPairBuff,
(uint2*)pPairBuffStartCurr,
(uint*)pPairScan,
(uint*)pPairOut,
pAABB,
numBodies
);
CUT_CHECK_ERROR("Kernel execution failed: btCudaSqueezeOverlappingPairBuffD");
} // btCuda_squeezeOverlappingPairBuff()
} // extern "C"

206
Extras/CUDA/btCudaBroadphase.h
View File

@@ -18,169 +18,95 @@ subject to the following restrictions:
#include "BulletCollision/BroadphaseCollision/btSimpleBroadphase.h" #include "BulletCollision/BroadphaseCollision/btSimpleBroadphase.h"
#include "btCudaBroadphaseKernel.h"
///The btCudaBroadphase uses CUDA to compute overlapping pairs using a GPU. ///The btCudaBroadphase uses CUDA to compute overlapping pairs using a GPU.
class btCudaBroadphase : public btSimpleBroadphase class btCudaBroadphase : public btSimpleBroadphase
{ {
bool m_bInitialized;
unsigned int m_numBodies;
unsigned int m_numCells;
bool m_bInitialized; unsigned int m_maxPairsPerBody;
int m_numParticles; btScalar m_cellFactorAABB;
// CPU data
// CPU data unsigned int* m_hBodiesHash;
float* m_hPos; unsigned int* m_hCellStart;
float* m_hVel;
float* m_hSortedPos;
unsigned int* m_hGridCounters;
unsigned int* m_hGridCells;
unsigned int* m_hParticleHash;
unsigned int* m_hCellStart;
unsigned int* m_hPairBuffStartCurr; unsigned int* m_hPairBuffStartCurr;
float* m_hAABB; btCuda3F1U* m_hAABB;
unsigned int* m_hPairBuff; unsigned int* m_hPairBuff;
unsigned int* m_hPairScan; unsigned int* m_hPairScan;
unsigned int* m_hPairOut; unsigned int* m_hPairOut;
// GPU data // GPU data
float* m_dPos[2]; unsigned int* m_dBodiesHash[2];
float* m_dVel[2]; unsigned int* m_dCellStart;
float* m_dSortedPos;
float* m_dSortedVel;
// uniform grid data
unsigned int* m_dGridCounters; // counts number of entries per grid cell
unsigned int* m_dGridCells; // contains indices of up to "m_maxParticlesPerCell" particles per cell
unsigned int* m_dParticleHash[2];
unsigned int* m_dCellStart;
unsigned int m_posVbo[2];
unsigned int m_colorVBO;
unsigned int m_currentPosRead, m_currentVelRead;
unsigned int m_currentPosWrite, m_currentVelWrite;
// buffers on GPU
unsigned int* m_dPairBuff; unsigned int* m_dPairBuff;
unsigned int* m_dPairBuffStartCurr; unsigned int* m_dPairBuffStartCurr;
float* m_dAABB; btCuda3F1U* m_dAABB;
unsigned int* m_dPairScan; unsigned int* m_dPairScan;
unsigned int* m_dPairOut; unsigned int* m_dPairOut;
unsigned int m_maxBodiesPerCell;
// params btCudaBroadphaseParams m_params;
struct SimParams& m_simParams; btScalar m_maxRadius;
// large proxies
int m_numLargeHandles;
int m_maxLargeHandles;
unsigned int m_maxParticlesPerCell; int m_LastLargeHandleIndex;
btSimpleBroadphaseProxy* m_pLargeHandles;
void* m_pLargeHandlesRawPtr;
int m_firstFreeLargeHandle;
int allocLargeHandle()
{
btAssert(m_numLargeHandles < m_maxLargeHandles);
int freeLargeHandle = m_firstFreeLargeHandle;
m_firstFreeLargeHandle = m_pLargeHandles[freeLargeHandle].GetNextFree();
m_numLargeHandles++;
if(freeLargeHandle > m_LastLargeHandleIndex)
{
m_LastLargeHandleIndex = freeLargeHandle;
}
return freeLargeHandle;
}
void freeLargeHandle(btSimpleBroadphaseProxy* proxy)
{
int handle = int(proxy - m_pLargeHandles);
btAssert((handle >= 0) && (handle < m_maxHandles));
if(handle == m_LastLargeHandleIndex)
{
m_LastLargeHandleIndex--;
}
proxy->SetNextFree(m_firstFreeLargeHandle);
m_firstFreeLargeHandle = handle;
proxy->m_clientObject = 0;
m_numLargeHandles--;
}
bool isLargeProxy(const btVector3& aabbMin, const btVector3& aabbMax);
bool isLargeProxy(btBroadphaseProxy* proxy);
// debug // debug
unsigned int m_numPairsAdded; unsigned int m_numPairsAdded;
unsigned int m_maxPairsPerParticle; unsigned int m_numPairsRemoved;
unsigned int m_numOverflows; unsigned int m_numOverflows;
//
protected:
unsigned int createVBO(unsigned int size);
void _initialize(int numParticles);
void _finalize();
public: public:
btCudaBroadphase(const btVector3& worldAabbMin,const btVector3& worldAabbMax,
enum ParticleArray int gridSizeX, int gridSizeY, int gridSizeZ,
{ int maxSmallProxies, int maxLargeProxies, int maxPairsPerBody,
POSITION, int maxBodiesPerCell = 8,
VELOCITY, btScalar cellFactorAABB = btScalar(1.0f));
};
enum ParticleConfig
{
CONFIG_RANDOM,
CONFIG_GRID,
_NUM_CONFIGS
};
btCudaBroadphase(SimParams& simParams,int maxProxies);
virtual ~btCudaBroadphase(); virtual ~btCudaBroadphase();
void initGrid(unsigned int* size, float spacing, float jitter, unsigned int numParticles);
void reset(ParticleConfig config);
void setArray(ParticleArray array, const float* data, int start, int count);
float* getArray(ParticleArray array);
void addSphere(int start, float *pos, float *vel, int r, float spacing);
virtual void calculateOverlappingPairs(btDispatcher* dispatcher); virtual void calculateOverlappingPairs(btDispatcher* dispatcher);
unsigned int getCurrentReadBuffer() const { return m_posVbo[m_currentPosRead]; } virtual btBroadphaseProxy* createProxy(const btVector3& aabbMin, const btVector3& aabbMax,int shapeType,void* userPtr ,short int collisionFilterGroup,short int collisionFilterMask, btDispatcher* dispatcher,void* multiSapProxy);
unsigned int getColorBuffer() const { return m_colorVBO; } virtual void destroyProxy(btBroadphaseProxy* proxy,btDispatcher* dispatcher);
void dumpParticles(unsigned int start, unsigned int count); virtual void rayTest(const btVector3& rayFrom,const btVector3& rayTo, btBroadphaseRayCallback& rayCallback);
void dumpGrid();
float* copyBuffersFromDeviceToHost();
void copyBuffersFromHostToDevice();
float* getHvelPtr();
float* getHposPtr();
void quickHack(float deltaTime);
void quickHack2();
void integrate();
void findOverlappingPairs(btDispatcher* dispatcher);
int3 calcGridPosCPU(float4 p);
uint calcGridHashCPU(int3 gridPos);
void computePairCacheChangesCPU(uint* pPairBuff, uint* pPairBuffStartCurr, uint* pPairScan, uint numParticles);
void computePairCacheChangesCPU_D(uint index, uint* pPairBuff, uint2* pPairBuffStartCurr, uint* pPairScan);
void findOverlappingPairsCPU( float* pAABB,
uint* pParticleHash,
uint* pCellStart,
uint* pPairBuff,
uint* pPairBuffStartCurr,
uint numParticles);
void findOverlappingPairsCPU_D( uint index,
float4* pAABB,
uint2* pParticleHash,
uint* pCellStart,
uint* pPairBuff,
uint2* pPairBuffStartCurr,
uint numParticles);
void findPairsInCellCPU(int3 gridPos,
uint index,
uint2* pParticleHash,
uint* pCellStart,
float4* pAABB,
uint* pPairBuff,
uint2* pPairBuffStartCurr,
uint numParticles);
uint cudaTestAABBOverlapCPU(float4 min0, float4 max0, float4 min1, float4 max1);
protected:
void _initialize();
void _finalize();
void scanOverlappingPairBuffCPU(); void scanOverlappingPairBuffCPU();
void squeezeOverlappingPairBuffCPU(uint* pPairBuff, uint* pPairBuffStartCurr, uint* pPairScan, uint* pPairOut, uint numParticles);
void squeezeOverlappingPairBuffCPU_D(uint index, uint* pPairBuff, uint2* pPairBuffStartCurr, uint* pPairScan, uint* pPairOut);
void addPairsToCacheCPU(btDispatcher* dispatcher); void addPairsToCacheCPU(btDispatcher* dispatcher);
void resetOverlappingPairBuffCPU(); void addLarge2LargePairsToCache(btDispatcher* dispatcher);
}; };
#endif //CUDA_BROADPHASE_H #endif //CUDA_BROADPHASE_H

113
Extras/CUDA/btCudaBroadphase.sln
View File

@@ -1,113 +0,0 @@

Microsoft Visual Studio Solution File, Format Version 9.00
# Visual Studio 2005
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "btCudaBroadphase", "btCudaBroadphase.vcproj", "{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}"
ProjectSection(ProjectDependencies) = postProject
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE} = {6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}
{61BD1097-CF2E-B296-DAA9-73A6FE135319} = {61BD1097-CF2E-B296-DAA9-73A6FE135319}
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A} = {7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}
EndProjectSection
EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "libbulletmath", "..\..\msvc\8\libbulletmath.vcproj", "{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}"
EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "libbulletdynamics", "..\..\msvc\8\libbulletdynamics.vcproj", "{61BD1097-CF2E-B296-DAA9-73A6FE135319}"
EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "libbulletcollision", "..\..\msvc\8\libbulletcollision.vcproj", "{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}"
EndProject
Global
GlobalSection(SolutionConfigurationPlatforms) = preSolution
Debug|Win32 = Debug|Win32
Debug|x64 = Debug|x64
DebugDoublePrecision|Win32 = DebugDoublePrecision|Win32
DebugDoublePrecision|x64 = DebugDoublePrecision|x64
EmuDebug|Win32 = EmuDebug|Win32
EmuDebug|x64 = EmuDebug|x64
EmuRelease|Win32 = EmuRelease|Win32
EmuRelease|x64 = EmuRelease|x64
Release|Win32 = Release|Win32
Release|x64 = Release|x64
ReleaseDoublePrecision|Win32 = ReleaseDoublePrecision|Win32
ReleaseDoublePrecision|x64 = ReleaseDoublePrecision|x64
EndGlobalSection
GlobalSection(ProjectConfigurationPlatforms) = postSolution
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Debug|Win32.ActiveCfg = Debug|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Debug|Win32.Build.0 = Debug|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Debug|x64.ActiveCfg = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Debug|x64.Build.0 = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.DebugDoublePrecision|Win32.ActiveCfg = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.DebugDoublePrecision|x64.ActiveCfg = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.DebugDoublePrecision|x64.Build.0 = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuDebug|Win32.ActiveCfg = EmuDebug|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuDebug|Win32.Build.0 = EmuDebug|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuDebug|x64.ActiveCfg = EmuDebug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuDebug|x64.Build.0 = EmuDebug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuRelease|Win32.ActiveCfg = EmuRelease|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuRelease|Win32.Build.0 = EmuRelease|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuRelease|x64.ActiveCfg = EmuRelease|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuRelease|x64.Build.0 = EmuRelease|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Release|Win32.ActiveCfg = Release|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Release|Win32.Build.0 = Release|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Release|x64.ActiveCfg = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Release|x64.Build.0 = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.ReleaseDoublePrecision|Win32.ActiveCfg = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.ReleaseDoublePrecision|x64.ActiveCfg = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.ReleaseDoublePrecision|x64.Build.0 = Release|x64
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Debug|Win32.ActiveCfg = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Debug|Win32.Build.0 = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Debug|x64.ActiveCfg = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.DebugDoublePrecision|Win32.ActiveCfg = DebugDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.DebugDoublePrecision|Win32.Build.0 = DebugDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.DebugDoublePrecision|x64.ActiveCfg = DebugDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuDebug|Win32.ActiveCfg = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuDebug|Win32.Build.0 = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuDebug|x64.ActiveCfg = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuRelease|Win32.ActiveCfg = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuRelease|Win32.Build.0 = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuRelease|x64.ActiveCfg = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Release|Win32.ActiveCfg = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Release|Win32.Build.0 = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Release|x64.ActiveCfg = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.ReleaseDoublePrecision|Win32.ActiveCfg = ReleaseDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.ReleaseDoublePrecision|Win32.Build.0 = ReleaseDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.ReleaseDoublePrecision|x64.ActiveCfg = ReleaseDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Debug|Win32.ActiveCfg = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Debug|Win32.Build.0 = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Debug|x64.ActiveCfg = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.DebugDoublePrecision|Win32.ActiveCfg = DebugDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.DebugDoublePrecision|Win32.Build.0 = DebugDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.DebugDoublePrecision|x64.ActiveCfg = DebugDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuDebug|Win32.ActiveCfg = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuDebug|Win32.Build.0 = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuDebug|x64.ActiveCfg = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuRelease|Win32.ActiveCfg = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuRelease|Win32.Build.0 = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuRelease|x64.ActiveCfg = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Release|Win32.ActiveCfg = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Release|Win32.Build.0 = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Release|x64.ActiveCfg = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.ReleaseDoublePrecision|Win32.ActiveCfg = ReleaseDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.ReleaseDoublePrecision|Win32.Build.0 = ReleaseDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.ReleaseDoublePrecision|x64.ActiveCfg = ReleaseDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Debug|Win32.ActiveCfg = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Debug|Win32.Build.0 = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Debug|x64.ActiveCfg = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.DebugDoublePrecision|Win32.ActiveCfg = DebugDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.DebugDoublePrecision|Win32.Build.0 = DebugDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.DebugDoublePrecision|x64.ActiveCfg = DebugDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuDebug|Win32.ActiveCfg = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuDebug|Win32.Build.0 = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuDebug|x64.ActiveCfg = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuRelease|Win32.ActiveCfg = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuRelease|Win32.Build.0 = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuRelease|x64.ActiveCfg = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Release|Win32.ActiveCfg = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Release|Win32.Build.0 = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Release|x64.ActiveCfg = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.ReleaseDoublePrecision|Win32.ActiveCfg = ReleaseDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.ReleaseDoublePrecision|Win32.Build.0 = ReleaseDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.ReleaseDoublePrecision|x64.ActiveCfg = ReleaseDoublePrecision|Win32
EndGlobalSection
GlobalSection(SolutionProperties) = preSolution
HideSolutionNode = FALSE
EndGlobalSection
EndGlobal

84
Extras/CUDA/btCudaBroadphaseKernel.h Normal file
View File

@@ -0,0 +1,84 @@
/*
Bullet Continuous Collision Detection and Physics Library
Copyright (c) 2003-2008 Erwin Coumans http://continuousphysics.com/Bullet/
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.
*/
//!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
// Keep this file free from Bullet headers
// it is included into CUDA program
//!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
#ifndef CUDA_BROADPHASE_KERNEL_H
#define CUDA_BROADPHASE_KERNEL_H
#define CUDA_BROADPHASE_USE_CUDA 1
#define BT_CUDA_PAIR_FOUND_FLG (0x40000000)
#define BT_CUDA_PAIR_NEW_FLG (0x20000000)
#define BT_CUDA_PAIR_ANY_FLG (BT_CUDA_PAIR_FOUND_FLG | BT_CUDA_PAIR_NEW_FLG)
struct btCudaBroadphaseParams
{
unsigned int m_gridSizeX;
unsigned int m_gridSizeY;
unsigned int m_gridSizeZ;
unsigned int m_numCells;
float m_worldOriginX;
float m_worldOriginY;
float m_worldOriginZ;
float m_cellSizeX;
float m_cellSizeY;
float m_cellSizeZ;
unsigned int m_numBodies;
unsigned int m_maxBodiesPerCell;
};
struct btCuda3F1U
{
float fx;
float fy;
float fz;
unsigned int uw;
};
extern "C"
{
void btCuda_allocateArray(void** devPtr, unsigned int size);
void btCuda_freeArray(void* devPtr);
void btCuda_copyArrayFromDevice(void* host, const void* device, unsigned int size);
void btCuda_copyArrayToDevice(void* device, const void* host, unsigned int size);
void btCuda_setParameters(btCudaBroadphaseParams* hostParams);
void btCuda_calcHashAABB(btCuda3F1U* pAABB, unsigned int* hash, unsigned int numBodies);
void btCuda_findCellStart(unsigned int* hash, unsigned int* cellStart, unsigned int numBodies, unsigned int numCells);
void btCuda_findOverlappingPairs( btCuda3F1U* pAABB, unsigned int* pHash,
unsigned int* pCellStart,
unsigned int* pPairBuff,
unsigned int* pPairBuffStartCurr,
unsigned int numBodies);
void btCuda_findPairsLarge( btCuda3F1U* pAABB, unsigned int* pHash,
unsigned int* pCellStart,
unsigned int* pPairBuff,
unsigned int* pPairBuffStartCurr,
unsigned int numBodies,
unsigned int numLarge);
void btCuda_computePairCacheChanges(unsigned int* pPairBuff, unsigned int* pPairBuffStartCurr,
unsigned int* pPairScan, btCuda3F1U* pAABB, unsigned int numBodies);
void btCuda_squeezeOverlappingPairBuff( unsigned int* pPairBuff, unsigned int* pPairBuffStartCurr, unsigned int* pPairScan,
unsigned int* pPairOut, btCuda3F1U* pAABB, unsigned int numBodies);
}
#endif // CUDA_BROADPHASE_KERNEL_H

593
Extras/CUDA/libbulletcuda.vcproj Normal file
View File

@@ -0,0 +1,593 @@
<?xml version="1.0" encoding="Windows-1252"?>
<VisualStudioProject
ProjectType="Visual C++"
Version="8.00"
Name="libbulletcuda"
ProjectGUID="{F74E8E02-0B47-4816-BD0B-FAEAE3343165}"
RootNamespace="libbulletcuda"
>
<Platforms>
<Platform
Name="Win32"
/>
</Platforms>
<ToolFiles>
</ToolFiles>
<Configurations>
<Configuration
Name="Release|Win32"
OutputDirectory="..\..\out\release8\build\libbulletcuda\"
IntermediateDirectory="..\..\out\release8\build\libbulletcuda\"
ConfigurationType="4"
UseOfMFC="0"
ATLMinimizesCRunTimeLibraryUsage="false"
>
<Tool
Name="VCPreBuildEventTool"
/>
<Tool
Name="VCCustomBuildTool"
/>
<Tool
Name="VCXMLDataGeneratorTool"
/>
<Tool
Name="VCWebServiceProxyGeneratorTool"
/>
<Tool
Name="VCMIDLTool"
PreprocessorDefinitions="NDEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE"
MkTypLibCompatible="true"
SuppressStartupBanner="true"
TargetEnvironment="1"
TypeLibraryName="..\..\out\release8\build\libbulletcuda\libbulletcuda.tlb"
/>
<Tool
Name="VCCLCompilerTool"
AdditionalOptions=" "
Optimization="2"
AdditionalIncludeDirectories=".;..\..;..\..\src;&quot;$(CUDA_INC_PATH)&quot;"
PreprocessorDefinitions="NDEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;WIN32"
StringPooling="true"
ExceptionHandling="0"
RuntimeLibrary="0"
BufferSecurityCheck="false"
EnableFunctionLevelLinking="true"
TreatWChar_tAsBuiltInType="false"
PrecompiledHeaderFile="..\..\out\release8\build\libbulletcuda\libbulletcuda.pch"
AssemblerListingLocation="..\..\out\release8\build\libbulletcuda\"
ObjectFile="..\..\out\release8\build\libbulletcuda\"
ProgramDataBaseFileName="..\..\out\release8\build\libbulletcuda\bulletcuda.pdb"
WarningLevel="3"
SuppressStartupBanner="true"
Detect64BitPortabilityProblems="true"
DebugInformationFormat="3"
CompileAs="2"
/>
<Tool
Name="VCManagedResourceCompilerTool"
/>
<Tool
Name="VCResourceCompilerTool"
PreprocessorDefinitions="NDEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;PROJECTGEN_VERSION=8"
Culture="1033"
AdditionalIncludeDirectories=".;..\..;..\..\src"
/>
<Tool
Name="VCPreLinkEventTool"
/>
<Tool
Name="VCLibrarianTool"
OutputFile="..\..\out\release8\libs\libbulletcuda.lib"
SuppressStartupBanner="true"
/>
<Tool
Name="VCALinkTool"
/>
<Tool
Name="VCXDCMakeTool"
/>
<Tool
Name="VCBscMakeTool"
/>
<Tool
Name="VCFxCopTool"
/>
<Tool
Name="VCPostBuildEventTool"
/>
</Configuration>
<Configuration
Name="ReleaseDll|Win32"
OutputDirectory="..\..\out\release_dll8\build\libbulletcuda\"
IntermediateDirectory="..\..\out\release_dll8\build\libbulletcuda\"
ConfigurationType="4"
UseOfMFC="0"
ATLMinimizesCRunTimeLibraryUsage="false"
>
<Tool
Name="VCPreBuildEventTool"
/>
<Tool
Name="VCCustomBuildTool"
/>
<Tool
Name="VCXMLDataGeneratorTool"
/>
<Tool
Name="VCWebServiceProxyGeneratorTool"
/>
<Tool
Name="VCMIDLTool"
PreprocessorDefinitions="NDEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE"
MkTypLibCompatible="true"
SuppressStartupBanner="true"
TargetEnvironment="1"
TypeLibraryName="..\..\out\release_dll8\build\libbulletcuda\libbulletcuda.tlb"
/>
<Tool
Name="VCCLCompilerTool"
AdditionalOptions=" "
Optimization="2"
AdditionalIncludeDirectories=".;..\..;..\..\src"
PreprocessorDefinitions="NDEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;WIN32"
StringPooling="true"
ExceptionHandling="0"
RuntimeLibrary="2"
BufferSecurityCheck="false"
EnableFunctionLevelLinking="true"
TreatWChar_tAsBuiltInType="false"
PrecompiledHeaderFile="..\..\out\release_dll8\build\libbulletcuda\libbulletcuda.pch"
AssemblerListingLocation="..\..\out\release_dll8\build\libbulletcuda\"
ObjectFile="..\..\out\release_dll8\build\libbulletcuda\"
ProgramDataBaseFileName="..\..\out\release_dll8\build\libbulletcuda\bulletcuda.pdb"
WarningLevel="3"
SuppressStartupBanner="true"
Detect64BitPortabilityProblems="true"
DebugInformationFormat="3"
CompileAs="0"
/>
<Tool
Name="VCManagedResourceCompilerTool"
/>
<Tool
Name="VCResourceCompilerTool"
PreprocessorDefinitions="NDEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;PROJECTGEN_VERSION=8"
Culture="1033"
AdditionalIncludeDirectories=".;..\..;..\..\src"
/>
<Tool
Name="VCPreLinkEventTool"
/>
<Tool
Name="VCLibrarianTool"
SuppressStartupBanner="true"
/>
<Tool
Name="VCALinkTool"
/>
<Tool
Name="VCXDCMakeTool"
/>
<Tool
Name="VCBscMakeTool"
/>
<Tool
Name="VCFxCopTool"
/>
<Tool
Name="VCPostBuildEventTool"
/>
</Configuration>
<Configuration
Name="ReleaseDoublePrecision|Win32"
OutputDirectory="..\..\out\release_dbl8\build\libbulletcuda\"
IntermediateDirectory="..\..\out\release_dbl8\build\libbulletcuda\"
ConfigurationType="4"
UseOfMFC="0"
ATLMinimizesCRunTimeLibraryUsage="false"
>
<Tool
Name="VCPreBuildEventTool"
/>
<Tool
Name="VCCustomBuildTool"
/>
<Tool
Name="VCXMLDataGeneratorTool"
/>
<Tool
Name="VCWebServiceProxyGeneratorTool"
/>
<Tool
Name="VCMIDLTool"
PreprocessorDefinitions="NDEBUG;BT_USE_DOUBLE_PRECISION;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE"
MkTypLibCompatible="true"
SuppressStartupBanner="true"
TargetEnvironment="1"
TypeLibraryName="..\..\out\release_dbl8\build\libbulletcuda\libbulletcuda.tlb"
/>
<Tool
Name="VCCLCompilerTool"
AdditionalOptions=" "
AdditionalIncludeDirectories=".;..\..;..\..\src"
PreprocessorDefinitions="NDEBUG;BT_USE_DOUBLE_PRECISION;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;WIN32"
ExceptionHandling="0"
TreatWChar_tAsBuiltInType="false"
PrecompiledHeaderFile="..\..\out\release_dbl8\build\libbulletcuda\libbulletcuda.pch"
AssemblerListingLocation="..\..\out\release_dbl8\build\libbulletcuda\"
ObjectFile="..\..\out\release_dbl8\build\libbulletcuda\"
ProgramDataBaseFileName="..\..\out\release_dbl8\build\libbulletcuda\bulletcuda.pdb"
WarningLevel="3"
SuppressStartupBanner="true"
Detect64BitPortabilityProblems="true"
CompileAs="0"
/>
<Tool
Name="VCManagedResourceCompilerTool"
/>
<Tool
Name="VCResourceCompilerTool"
PreprocessorDefinitions="NDEBUG;BT_USE_DOUBLE_PRECISION;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;PROJECTGEN_VERSION=8"
Culture="1033"
AdditionalIncludeDirectories=".;..\..;..\..\src"
/>
<Tool
Name="VCPreLinkEventTool"
/>
<Tool
Name="VCLibrarianTool"
SuppressStartupBanner="true"
/>
<Tool
Name="VCALinkTool"
/>
<Tool
Name="VCXDCMakeTool"
/>
<Tool
Name="VCBscMakeTool"
/>
<Tool
Name="VCFxCopTool"
/>
<Tool
Name="VCPostBuildEventTool"
/>
</Configuration>
<Configuration
Name="Debug|Win32"
OutputDirectory="..\..\out\debug8\build\libbulletcuda\"
IntermediateDirectory="..\..\out\debug8\build\libbulletcuda\"
ConfigurationType="4"
UseOfMFC="0"
ATLMinimizesCRunTimeLibraryUsage="false"
>
<Tool
Name="VCPreBuildEventTool"
/>
<Tool
Name="VCCustomBuildTool"
/>
<Tool
Name="VCXMLDataGeneratorTool"
/>
<Tool
Name="VCWebServiceProxyGeneratorTool"
/>
<Tool
Name="VCMIDLTool"
PreprocessorDefinitions="_DEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE"
MkTypLibCompatible="true"
SuppressStartupBanner="true"
TargetEnvironment="1"
TypeLibraryName="..\..\out\debug8\build\libbulletcuda\libbulletcuda.tlb"
/>
<Tool
Name="VCCLCompilerTool"
AdditionalOptions=" "
Optimization="0"
AdditionalIncludeDirectories=".;..\..;..\..\src;&quot;$(CUDA_INC_PATH)&quot;"
PreprocessorDefinitions="_DEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;WIN32"
MinimalRebuild="true"
ExceptionHandling="0"
RuntimeLibrary="1"
TreatWChar_tAsBuiltInType="false"
RuntimeTypeInfo="false"
PrecompiledHeaderFile="..\..\out\debug8\build\libbulletmath\libbulletcuda.pch"
AssemblerListingLocation="..\..\out\debug8\build\libbulletcuda\"
ObjectFile="..\..\out\debug8\build\libbulletcuda\"
ProgramDataBaseFileName="..\..\out\debug8\build\libbulletcuda\bulletcuda.pdb"
WarningLevel="3"
SuppressStartupBanner="true"
Detect64BitPortabilityProblems="true"
DebugInformationFormat="3"
CompileAs="0"
/>
<Tool
Name="VCManagedResourceCompilerTool"
/>
<Tool
Name="VCResourceCompilerTool"
PreprocessorDefinitions="_DEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;PROJECTGEN_VERSION=8"
Culture="1033"
AdditionalIncludeDirectories=".;..\..;..\..\src"
/>
<Tool
Name="VCPreLinkEventTool"
/>
<Tool
Name="VCLibrarianTool"
OutputFile="..\..\out\debug8\libs\libbulletcuda_d.lib"
SuppressStartupBanner="true"
/>
<Tool
Name="VCALinkTool"
/>
<Tool
Name="VCXDCMakeTool"
/>
<Tool
Name="VCBscMakeTool"
/>
<Tool
Name="VCFxCopTool"
/>
<Tool
Name="VCPostBuildEventTool"
/>
</Configuration>
<Configuration
Name="DebugDll|Win32"
OutputDirectory="..\..\out\debug_dll8\build\libbulletcuda\"
IntermediateDirectory="..\..\out\debug_dll8\build\libbulletcuda\"
ConfigurationType="4"
UseOfMFC="0"
ATLMinimizesCRunTimeLibraryUsage="false"
>
<Tool
Name="VCPreBuildEventTool"
/>
<Tool
Name="VCCustomBuildTool"
/>
<Tool
Name="VCXMLDataGeneratorTool"
/>
<Tool
Name="VCWebServiceProxyGeneratorTool"
/>
<Tool
Name="VCMIDLTool"
PreprocessorDefinitions="_DEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE"
MkTypLibCompatible="true"
SuppressStartupBanner="true"
TargetEnvironment="1"
TypeLibraryName="..\..\out\debug_dll8\build\libbulletcuda\libbulletcuda.tlb"
/>
<Tool
Name="VCCLCompilerTool"
AdditionalOptions=" "
Optimization="0"
AdditionalIncludeDirectories=".;..\..;..\..\src"
PreprocessorDefinitions="_DEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;WIN32"
MinimalRebuild="true"
ExceptionHandling="0"
RuntimeLibrary="3"
TreatWChar_tAsBuiltInType="false"
RuntimeTypeInfo="false"
PrecompiledHeaderFile="..\..\out\debug_dll8\build\libbulletcuda\libbulletcuda.pch"
AssemblerListingLocation="..\..\out\debug_dll8\build\libbulletcuda\"
ObjectFile="..\..\out\debug_dll8\build\libbulletcuda\"
ProgramDataBaseFileName="..\..\out\debug_dll8\build\libbulletmath\bulletcuda.pdb"
WarningLevel="3"
SuppressStartupBanner="true"
Detect64BitPortabilityProblems="true"
DebugInformationFormat="4"
CompileAs="0"
/>
<Tool
Name="VCManagedResourceCompilerTool"
/>
<Tool
Name="VCResourceCompilerTool"
PreprocessorDefinitions="_DEBUG;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;PROJECTGEN_VERSION=8"
Culture="1033"
AdditionalIncludeDirectories=".;..\..;..\..\src"
/>
<Tool
Name="VCPreLinkEventTool"
/>
<Tool
Name="VCLibrarianTool"
SuppressStartupBanner="true"
/>
<Tool
Name="VCALinkTool"
/>
<Tool
Name="VCXDCMakeTool"
/>
<Tool
Name="VCBscMakeTool"
/>
<Tool
Name="VCFxCopTool"
/>
<Tool
Name="VCPostBuildEventTool"
/>
</Configuration>
<Configuration
Name="DebugDoublePrecision|Win32"
OutputDirectory="..\..\out\debug_dbl8\build\libbulletcuda\"
IntermediateDirectory="..\..\out\debug_dbl8\build\libbulletcuda\"
ConfigurationType="4"
UseOfMFC="0"
ATLMinimizesCRunTimeLibraryUsage="false"
>
<Tool
Name="VCPreBuildEventTool"
/>
<Tool
Name="VCCustomBuildTool"
/>
<Tool
Name="VCXMLDataGeneratorTool"
/>
<Tool
Name="VCWebServiceProxyGeneratorTool"
/>
<Tool
Name="VCMIDLTool"
PreprocessorDefinitions="_DEBUG;BT_USE_DOUBLE_PRECISION;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE"
MkTypLibCompatible="true"
SuppressStartupBanner="true"
TargetEnvironment="1"
TypeLibraryName="..\..\out\debug_dbl8\build\libbulletcuda\libbulletcuda.tlb"
/>
<Tool
Name="VCCLCompilerTool"
AdditionalOptions=" "
Optimization="0"
AdditionalIncludeDirectories=".;..\..;..\..\src"
PreprocessorDefinitions="_DEBUG;BT_USE_DOUBLE_PRECISION;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;WIN32"
MinimalRebuild="true"
ExceptionHandling="0"
RuntimeLibrary="1"
TreatWChar_tAsBuiltInType="false"
RuntimeTypeInfo="false"
PrecompiledHeaderFile="..\..\out\debug_dbl8\build\libbulletcuda\libbulletcuda.pch"
AssemblerListingLocation="..\..\out\debug_dbl8\build\libbulletcuda\"
ObjectFile="..\..\out\debug_dbl8\build\libbulletcuda\"
ProgramDataBaseFileName="..\..\out\debug_dbl8\build\libbulletmath\bulletcuda.pdb"
WarningLevel="3"
SuppressStartupBanner="true"
Detect64BitPortabilityProblems="true"
DebugInformationFormat="4"
CompileAs="0"
/>
<Tool
Name="VCManagedResourceCompilerTool"
/>
<Tool
Name="VCResourceCompilerTool"
PreprocessorDefinitions="_DEBUG;BT_USE_DOUBLE_PRECISION;_LIB;_WINDOWS;_CRT_SECURE_NO_DEPRECATE;_CRT_NONSTDC_NO_DEPRECATE;PROJECTGEN_VERSION=8"
Culture="1033"
AdditionalIncludeDirectories=".;..\..;..\..\src"
/>
<Tool
Name="VCPreLinkEventTool"
/>
<Tool
Name="VCLibrarianTool"
SuppressStartupBanner="true"
/>
<Tool
Name="VCALinkTool"
/>
<Tool
Name="VCXDCMakeTool"
/>
<Tool
Name="VCBscMakeTool"
/>
<Tool
Name="VCFxCopTool"
/>
<Tool
Name="VCPostBuildEventTool"
/>
</Configuration>
</Configurations>
<References>
</References>
<Files>
<Filter
Name="Source Files"
>
<File
RelativePath=".\btCudaBroadphase.cpp"
>
</File>
<File
RelativePath=".\btCudaBroadphase.cu"
>
<FileConfiguration
Name="Release|Win32"
>
<Tool
Name="VCCustomBuildTool"
Description="CUDA compiling"
CommandLine="&quot;$(CUDA_BIN_PATH)\nvcc.exe&quot; -arch=sm_10 -use_fast_math -ccbin &quot;$(VCInstallDir)\bin&quot; -c -I &quot;$(CUDA_INC_PATH)&quot;; -I./ -I../../Glut -o $(ConfigurationName)\btCudaBroadphase_cu.obj btCudaBroadphase.cu&#x0D;&#x0A;"
AdditionalDependencies="btCudaBroadphaseKernel.h;radixsort.cu"
Outputs="$(ConfigurationName)\btCudaBroadphase_cu.obj"
/>
</FileConfiguration>
<FileConfiguration
Name="Debug|Win32"
>
<Tool
Name="VCCustomBuildTool"
Description="CUDA compiling"
CommandLine="&quot;$(CUDA_BIN_PATH)\nvcc.exe&quot; -arch=sm_10 -ccbin &quot;$(VCInstallDir)\bin&quot; -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd -c -I &quot;$(CUDA_INC_PATH)&quot;; -I./ -I../../Glut -o $(ConfigurationName)\btCudaBroadphase_cu.obj btCudaBroadphase.cu&#x0D;&#x0A;"
AdditionalDependencies="btCudaBroadphaseKernel.h;radixsort.cu"
Outputs="$(ConfigurationName)\btCudaBroadphase_cu.obj"
/>
</FileConfiguration>
</File>
<File
RelativePath=".\btCudaBroadphaseKernel.h"
>
</File>
<File
RelativePath=".\cutil_math.h"
>
</File>
<File
RelativePath=".\radixsort.cu"
>
<FileConfiguration
Name="Release|Win32"
>
<Tool
Name="VCCustomBuildTool"
Description="CUDA compiling"
CommandLine="&quot;$(CUDA_BIN_PATH)\nvcc.exe&quot; -arch=sm_10 -use_fast_math -ccbin &quot;$(VCInstallDir)\bin&quot; -c -I &quot;$(CUDA_INC_PATH)&quot;; -I./ -I../../../Glut -o $(ConfigurationName)\radixsort_cu.obj radixsort.cu&#x0D;&#x0A;"
AdditionalDependencies="radixsort.cuh; radixsort_kernel.cu"
Outputs="$(ConfigurationName)\radixsort_cu.obj"
/>
</FileConfiguration>
<FileConfiguration
Name="Debug|Win32"
>
<Tool
Name="VCCustomBuildTool"
Description="CUDA compiling"
CommandLine="&quot;$(CUDA_BIN_PATH)\nvcc.exe&quot; -arch=sm_10 -ccbin &quot;$(VCInstallDir)\bin&quot; -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd -c -I &quot;$(CUDA_INC_PATH)&quot;; -I./ -I../../Glut -o $(ConfigurationName)\radixsort_cu.obj radixsort.cu&#x0D;&#x0A;"
AdditionalDependencies="radixsort.cuh; radixsort_kernel.cu"
Outputs="$(ConfigurationName)\radixsort_cu.obj"
/>
</FileConfiguration>
</File>
<File
RelativePath=".\radixsort.cuh"
>
</File>
<File
RelativePath=".\radixsort_kernel.cu"
>
</File>
</Filter>
<Filter
Name="Header Files"
>
<File
RelativePath=".\btCudaBroadphase.h"
>
</File>
</Filter>
</Files>
<Globals>
</Globals>
</VisualStudioProject>

813
Extras/CUDA/particleSystem.cpp
View File

@@ -32,6 +32,7 @@
#include "radixsort.cuh" #include "radixsort.cuh"
#include "particles_kernel.cuh" #include "particles_kernel.cuh"
//#include <cutil.h>
#include <assert.h> #include <assert.h>
#include <math.h> #include <math.h>
@@ -42,6 +43,7 @@
#include <GL/glew.h> #include <GL/glew.h>
#include <btBulletDynamicsCommon.h> #include <btBulletDynamicsCommon.h>
#include "../../Demos/OpenGL/GLDebugDrawer.h"
#include "btCudaBroadphase.h" #include "btCudaBroadphase.h"
@@ -50,22 +52,609 @@
#define CUDART_PI_F 3.141592654f #define CUDART_PI_F 3.141592654f
#endif #endif
#define USE_BULLET 1
#define VEL_DIR_FACT (30.0F)
#define ACC_DIR_FACT (VEL_DIR_FACT*VEL_DIR_FACT)
#define VEL_INV_FACT (1.0F/VEL_DIR_FACT)
#define ACC_INV_FACT (1.0F/ACC_DIR_FACT)
GLDebugDrawer debugDrawer;
ParticleSystem::ParticleSystem(uint numParticles, uint3 gridSize) : ParticleSystem::ParticleSystem(uint numParticles, uint3 gridSize) :
m_simulationMode(SIMULATION_BULLET_CPU)//SIMULATION_CUDA) m_bInitialized(false),
m_numParticles(numParticles),
m_hPos(0),
m_hVel(0),
m_currentPosRead(0),
m_currentVelRead(0),
m_currentPosWrite(1),
m_currentVelWrite(1),
m_gridSize(gridSize),
m_maxParticlesPerCell(4),
m_timer(0),
m_solverIterations(1),
// m_simulationMode(SIMULATION_CUDA)
m_simulationMode(SIMULATION_BULLET_CPU)
{ {
this->m_params.numBodies = numParticles; m_dPos[0] = m_dPos[1] = 0;
this->m_params.m_gridSize = gridSize; m_dVel[0] = m_dVel[1] = 0;
m_numGridCells = m_gridSize.x*m_gridSize.y*m_gridSize.z;
float3 worldSize = make_float3(2.0f, 2.0f, 2.0f);
// set simulation parameters
m_params.gridSize = m_gridSize;
m_params.numCells = m_numGridCells;
m_params.numBodies = m_numParticles;
m_params.maxParticlesPerCell = m_maxParticlesPerCell;
m_params.worldOrigin = make_float3(-1.0f, -1.0f, -1.0f);
m_params.cellSize = make_float3(worldSize.x / m_gridSize.x, worldSize.y / m_gridSize.y, worldSize.z / m_gridSize.z);
m_params.particleRadius = m_params.cellSize.x * 0.5f;
m_params.colliderPos = make_float4(0.0f, -0.7f, 0.0f, 1.0f);
m_params.colliderRadius = 0.2f;
m_params.spring = 0.5f;
m_params.damping = 0.02f;
m_params.shear = 0.1f;
m_params.attraction = 0.0f;
m_params.boundaryDamping = -0.5f;
m_params.gravity = make_float3(0.0f, -0.0003f, 0.0f);
m_params.globalDamping = 1.0f;
_initialize(numParticles);
#if USE_BULLET
initializeBullet(); initializeBullet();
#endif
} }
ParticleSystem::~ParticleSystem() ParticleSystem::~ParticleSystem()
{ {
#if USE_BULLET
finalizeBullet(); finalizeBullet();
#endif
_finalize();
m_numParticles = 0;
}
uint
ParticleSystem::createVBO(uint size)
{
GLuint vbo;
glGenBuffers(1, &vbo);
glBindBuffer(GL_ARRAY_BUFFER, vbo);
glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);
glBindBuffer(GL_ARRAY_BUFFER, 0);
registerGLBufferObject(vbo);
return vbo;
}
inline float lerp(float a, float b, float t)
{
return a + t*(b-a);
}
void colorRamp(float t, float *r)
{
const int ncolors = 7;
float c[ncolors][3] = {
{ 1.0, 0.0, 0.0, },
{ 1.0, 0.5, 0.0, },
{ 1.0, 1.0, 0.0, },
{ 0.0, 1.0, 0.0, },
{ 0.0, 1.0, 1.0, },
{ 0.0, 0.0, 1.0, },
{ 1.0, 0.0, 1.0, },
};
t = t * (ncolors-1);
int i = (int) t;
float u = t - floor(t);
r[0] = lerp(c[i][0], c[i+1][0], u);
r[1] = lerp(c[i][1], c[i+1][1], u);
r[2] = lerp(c[i][2], c[i+1][2], u);
}
void
ParticleSystem::_initialize(int numParticles)
{
assert(!m_bInitialized);
m_numParticles = numParticles;
// allocate host storage
m_hPos = new float[m_numParticles*4];
m_hVel = new float[m_numParticles*4];
memset(m_hPos, 0, m_numParticles*4*sizeof(float));
memset(m_hVel, 0, m_numParticles*4*sizeof(float));
m_hGridCounters = new uint[m_numGridCells];
m_hGridCells = new uint[m_numGridCells*m_maxParticlesPerCell];
memset(m_hGridCounters, 0, m_numGridCells*sizeof(uint));
memset(m_hGridCells, 0, m_numGridCells*m_maxParticlesPerCell*sizeof(uint));
m_hParticleHash = new uint[m_numParticles*2];
memset(m_hParticleHash, 0, m_numParticles*2*sizeof(uint));
m_hCellStart = new uint[m_numGridCells];
memset(m_hCellStart, 0, m_numGridCells*sizeof(uint));
// allocate GPU data
unsigned int memSize = sizeof(float) * 4 * m_numParticles;
m_posVbo[0] = createVBO(memSize);
m_posVbo[1] = createVBO(memSize);
allocateArray((void**)&m_dVel[0], memSize);
allocateArray((void**)&m_dVel[1], memSize);
allocateArray((void**)&m_dSortedPos, memSize);
allocateArray((void**)&m_dSortedVel, memSize);
#if USE_SORT
allocateArray((void**)&m_dParticleHash[0], m_numParticles*2*sizeof(uint));
allocateArray((void**)&m_dParticleHash[1], m_numParticles*2*sizeof(uint));
allocateArray((void**)&m_dCellStart, m_numGridCells*sizeof(uint));
#else
allocateArray((void**)&m_dGridCounters, m_numGridCells*sizeof(uint));
allocateArray((void**)&m_dGridCells, m_numGridCells*m_maxParticlesPerCell*sizeof(uint));
#endif
m_colorVBO = createVBO(m_numParticles*4*sizeof(float));
#if 1
// fill color buffer
glBindBufferARB(GL_ARRAY_BUFFER, m_colorVBO);
float *data = (float *) glMapBufferARB(GL_ARRAY_BUFFER, GL_WRITE_ONLY);
float *ptr = data;
for(uint i=0; i<m_numParticles; i++) {
float t = i / (float) m_numParticles;
#if 0
*ptr++ = rand() / (float) RAND_MAX;
*ptr++ = rand() / (float) RAND_MAX;
*ptr++ = rand() / (float) RAND_MAX;
#else
colorRamp(t, ptr);
ptr+=3;
#endif
*ptr++ = 1.0f;
}
glUnmapBufferARB(GL_ARRAY_BUFFER);
#endif
// CUT_SAFE_CALL(cutCreateTimer(&m_timer));
setParameters(&m_params);
m_bInitialized = true;
}
void
ParticleSystem::_finalize()
{
assert(m_bInitialized);
delete [] m_hPos;
delete [] m_hVel;
delete [] m_hGridCounters;
delete [] m_hGridCells;
freeArray(m_dVel[0]);
freeArray(m_dVel[1]);
freeArray(m_dSortedPos);
freeArray(m_dSortedVel);
#if USE_SORT
freeArray(m_dParticleHash[0]);
freeArray(m_dParticleHash[1]);
freeArray(m_dCellStart);
#else
freeArray(m_dGridCounters);
freeArray(m_dGridCells);
#endif
unregisterGLBufferObject(m_posVbo[0]);
unregisterGLBufferObject(m_posVbo[1]);
glDeleteBuffers(2, (const GLuint*)m_posVbo);
glDeleteBuffers(1, (const GLuint*)&m_colorVBO);
}
void
ParticleSystem::update(float deltaTime)
{
assert(m_bInitialized);
#if USE_BULLET
switch (m_simulationMode)
{
case SIMULATION_CUDA:
{
updateCuda(deltaTime);
break;
}
case SIMULATION_BULLET_CPU:
{
updateBullet(deltaTime);
break;
}
default:
{
printf("unknown simulation method\n");
}
}
#else
updateCuda(deltaTime);
#endif
}
void
ParticleSystem::updateBullet(float deltaTime)
{
float* hPos = copyBuffersFromDeviceToHost();
float* hVel = m_hVel;
for (uint i=0;i<m_params.numBodies;i++)
{
float3 pos;
pos.x = hPos[i*4];
pos.y = hPos[i*4+1];
pos.z = hPos[i*4+2];
float3 vel;
vel.x = hVel[i*4];
vel.y = hVel[i*4+1];
vel.z = hVel[i*4+2];
// if (pos.x > 1.0f - m_params.particleRadius) { pos.x = 1.0f - m_params.particleRadius; vel.x *= m_params.boundaryDamping; }
// if (pos.x < -1.0f + m_params.particleRadius) { pos.x = -1.0f + m_params.particleRadius; vel.x *= m_params.boundaryDamping;}
// if (pos.y > 1.0f - m_params.particleRadius) { pos.y = 1.0f - m_params.particleRadius; vel.y *= m_params.boundaryDamping; }
// if (pos.y < -1.0f + m_params.particleRadius) { pos.y = -1.0f + m_params.particleRadius; vel.y *= m_params.boundaryDamping;}
// if (pos.z > 1.0f - m_params.particleRadius) { pos.z = 1.0f - m_params.particleRadius; vel.z *= m_params.boundaryDamping; }
// if (pos.z < -1.0f + m_params.particleRadius) { pos.z = -1.0f + m_params.particleRadius; vel.z *= m_params.boundaryDamping;}
btTransform& trans = m_bulletParticles[i]->getWorldTransform();
trans.setOrigin(btVector3(pos.x, pos.y, pos.z));
m_bulletParticles[i]->setLinearVelocity(btVector3(vel.x, vel.y, vel.z)*btScalar(VEL_DIR_FACT));
m_bulletParticles[i]->setAngularVelocity(btVector3(0,0,0));
}
glUnmapBufferARB(GL_ARRAY_BUFFER);
std::swap(m_currentPosRead, m_currentPosWrite);
std::swap(m_currentVelRead, m_currentVelWrite);
btTransform& collTrans = m_bulletCollider->getWorldTransform();
collTrans.setOrigin(btVector3(m_params.colliderPos.x, m_params.colliderPos.y, m_params.colliderPos.z));
m_dynamicsWorld->stepSimulation(deltaTime);
glBindBufferARB(GL_ARRAY_BUFFER, m_posVbo[m_currentPosRead]);
hPos = (float *) glMapBufferARB(GL_ARRAY_BUFFER, GL_READ_WRITE);//GL_WRITE_ONLY);
//sync transform and velocity from Bullet to particle system
for (uint i=0;i<m_params.numBodies;i++)
{
btTransform& trans = m_bulletParticles[i]->getWorldTransform();
hPos[i*4] = trans.getOrigin().getX();
hPos[i*4+1] = trans.getOrigin().getY();
hPos[i*4+2] = trans.getOrigin().getZ();
hVel[i*4] = m_bulletParticles[i]->getLinearVelocity().getX() * VEL_INV_FACT;
hVel[i*4+1] = m_bulletParticles[i]->getLinearVelocity().getY() * VEL_INV_FACT;
hVel[i*4+2] = m_bulletParticles[i]->getLinearVelocity().getZ() * VEL_INV_FACT;
}
copyBuffersFromHostToDevice();
collTrans = m_bulletCollider->getWorldTransform();
m_params.colliderPos.x = collTrans.getOrigin().getX();
m_params.colliderPos.y = collTrans.getOrigin().getY();
m_params.colliderPos.z = collTrans.getOrigin().getZ();
}
void
ParticleSystem::updateCuda(float deltaTime)
{
#ifndef BT_NO_PROFILE
CProfileManager::Reset();
#endif //BT_NO_PROFILE
BT_PROFILE("update CUDA");
// update constants
setParameters(&m_params);
// integrate
{
BT_PROFILE("integrate");
integrateSystem(m_posVbo[m_currentPosRead], m_posVbo[m_currentPosWrite],
m_dVel[m_currentVelRead], m_dVel[m_currentVelWrite],
deltaTime,
m_numParticles);
}
std::swap(m_currentPosRead, m_currentPosWrite);
std::swap(m_currentVelRead, m_currentVelWrite);
#if USE_SORT
// sort and search method
// calculate hash
{
BT_PROFILE("calcHash");
calcHash(m_posVbo[m_currentPosRead],
m_dParticleHash[0],
m_numParticles);
}
#if DEBUG_GRID
copyArrayFromDevice((void *) m_hParticleHash, (void *) m_dParticleHash[0], 0, sizeof(uint)*2*m_numParticles);
printf("particle hash:\n");
for(uint i=0; i<m_numParticles; i++) {
printf("%d: %d, %d\n", i, m_hParticleHash[i*2], m_hParticleHash[i*2+1]);
}
#endif
// sort particles based on hash
{
BT_PROFILE("RadixSort");
RadixSort((KeyValuePair *) m_dParticleHash[0], (KeyValuePair *) m_dParticleHash[1], m_numParticles, 32);
}
#if DEBUG_GRID
copyArrayFromDevice((void *) m_hParticleHash, (void *) m_dParticleHash[0], 0, sizeof(uint)*2*m_numParticles);
printf("particle hash sorted:\n");
for(uint i=0; i<m_numParticles; i++) {
printf("%d: %d, %d\n", i, m_hParticleHash[i*2], m_hParticleHash[i*2+1]);
}
#endif
// reorder particle arrays into sorted order and
// find start of each cell
{
BT_PROFILE("reorder");
reorderDataAndFindCellStart(m_dParticleHash[0],
m_posVbo[m_currentPosRead],
m_dVel[m_currentVelRead],
m_dSortedPos,
m_dSortedVel,
m_dCellStart,
m_numParticles,
m_numGridCells);
}
#if DEBUG_GRID
copyArrayFromDevice((void *) m_hCellStart, (void *) m_dCellStart, 0, sizeof(uint)*m_numGridCells);
printf("cell start:\n");
for(uint i=0; i<m_numGridCells; i++) {
printf("%d: %d\n", i, m_hCellStart[i]);
}
#endif
#else
// update grid using atomics
updateGrid(m_posVbo[m_currentPosRead],
m_dGridCounters,
m_dGridCells,
m_numParticles,
m_numGridCells);
#endif
// process collisions
{
BT_PROFILE("collide");
for(uint i=0; i<m_solverIterations; i++) {
collide(m_posVbo[m_currentPosRead], m_posVbo[m_currentPosWrite],
m_dSortedPos, m_dSortedVel,
m_dVel[m_currentVelRead], m_dVel[m_currentVelWrite],
m_dGridCounters,
m_dGridCells,
m_dParticleHash[0],
m_dCellStart,
m_numParticles,
m_numGridCells,
m_maxParticlesPerCell
);
std::swap(m_currentVelRead, m_currentVelWrite);
}
}
#ifndef BT_NO_PROFILE
CProfileManager::Increment_Frame_Counter();
#endif //BT_NO_PROFILE
}
void
ParticleSystem::dumpGrid()
{
// debug
copyArrayFromDevice(m_hGridCounters, m_dGridCounters, 0, sizeof(uint)*m_numGridCells);
copyArrayFromDevice(m_hGridCells, m_dGridCells, 0, sizeof(uint)*m_numGridCells*m_maxParticlesPerCell);
uint total = 0;
uint maxPerCell = 0;
for(uint i=0; i<m_numGridCells; i++) {
if (m_hGridCounters[i] > maxPerCell)
maxPerCell = m_hGridCounters[i];
if (m_hGridCounters[i] > 0) {
printf("%d (%d): ", i, m_hGridCounters[i]);
for(uint j=0; j<m_hGridCounters[i]; j++) {
printf("%d ", m_hGridCells[i*m_maxParticlesPerCell + j]);
}
total += m_hGridCounters[i];
printf("\n");
}
}
printf("max per cell = %d\n", maxPerCell);
printf("total = %d\n", total);
}
void
ParticleSystem::dumpParticles(uint start, uint count)
{
// debug
copyArrayFromDevice(m_hPos, 0, m_posVbo[m_currentPosRead], sizeof(float)*4*count);
copyArrayFromDevice(m_hVel, m_dVel[m_currentVelRead], 0, sizeof(float)*4*count);
for(uint i=start; i<start+count; i++) {
// printf("%d: ", i);
printf("pos: (%.4f, %.4f, %.4f, %.4f)\n", m_hPos[i*4+0], m_hPos[i*4+1], m_hPos[i*4+2], m_hPos[i*4+3]);
printf("vel: (%.4f, %.4f, %.4f, %.4f)\n", m_hVel[i*4+0], m_hVel[i*4+1], m_hVel[i*4+2], m_hVel[i*4+3]);
}
}
float*
ParticleSystem::getArray(ParticleArray array)
{
assert(m_bInitialized);
float* hdata = 0;
float* ddata = 0;
unsigned int vbo = 0;
switch (array)
{
default:
case POSITION:
hdata = m_hPos;
ddata = m_dPos[m_currentPosRead];
vbo = m_posVbo[m_currentPosRead];
break;
case VELOCITY:
hdata = m_hVel;
ddata = m_dVel[m_currentVelRead];
break;
}
copyArrayFromDevice(hdata, ddata, vbo, m_numParticles*4*sizeof(float));
return hdata;
}
void
ParticleSystem::setArray(ParticleArray array, const float* data, int start, int count)
{
assert(m_bInitialized);
switch (array)
{
default:
case POSITION:
{
unregisterGLBufferObject(m_posVbo[m_currentPosRead]);
glBindBuffer(GL_ARRAY_BUFFER, m_posVbo[m_currentPosRead]);
glBufferSubData(GL_ARRAY_BUFFER, start*4*sizeof(float), count*4*sizeof(float), data);
glBindBuffer(GL_ARRAY_BUFFER, 0);
registerGLBufferObject(m_posVbo[m_currentPosRead]);
}
break;
case VELOCITY:
copyArrayToDevice(m_dVel[m_currentVelRead], data, start*4*sizeof(float), count*4*sizeof(float));
break;
}
}
inline float frand()
{
return rand() / (float) RAND_MAX;
}
void
ParticleSystem::initGrid(uint *size, float spacing, float jitter, uint numParticles)
{
srand(1973);
for(uint z=0; z<size[2]; z++) {
for(uint y=0; y<size[1]; y++) {
for(uint x=0; x<size[0]; x++) {
uint i = (z*size[1]*size[0]) + (y*size[0]) + x;
if (i < numParticles) {
m_hPos[i*4] = (spacing * x) + m_params.particleRadius - 1.0f + (frand()*2.0f-1.0f)*jitter;
m_hPos[i*4+1] = (spacing * y) + m_params.particleRadius - 1.0f + (frand()*2.0f-1.0f)*jitter;
m_hPos[i*4+2] = (spacing * z) + m_params.particleRadius - 1.0f + (frand()*2.0f-1.0f)*jitter;
m_hPos[i*4+3] = 1.0f;
m_hVel[i*4] = 0.0f;
m_hVel[i*4+1] = 0.0f;
m_hVel[i*4+2] = 0.0f;
m_hVel[i*4+3] = 0.0f;
}
}
}
}
}
void
ParticleSystem::reset(ParticleConfig config)
{
switch(config)
{
default:
case CONFIG_RANDOM:
{
int p = 0, v = 0;
for(uint i=0; i < m_numParticles; i++)
{
float point[3];
point[0] = frand();
point[1] = frand();
point[2] = frand();
m_hPos[p++] = 2 * (point[0] - 0.5f);
m_hPos[p++] = 2 * (point[1] - 0.5f);
m_hPos[p++] = 2 * (point[2] - 0.5f);
m_hPos[p++] = 1.0f; // radius
m_hVel[v++] = 0.0f;
m_hVel[v++] = 0.0f;
m_hVel[v++] = 0.0f;
m_hVel[v++] = 0.0f;
}
}
break;
case CONFIG_GRID:
{
float jitter = m_params.particleRadius*0.01f;
uint s = (int) ceilf(powf((float) m_numParticles, 1.0f / 3.0f));
uint gridSize[3];
gridSize[0] = gridSize[1] = gridSize[2] = s;
initGrid(gridSize, m_params.particleRadius*2.0f, jitter, m_numParticles);
}
break;
}
setArray(POSITION, m_hPos, 0, m_numParticles);
setArray(VELOCITY, m_hVel, 0, m_numParticles);
}
void
ParticleSystem::addSphere(int start, float *pos, float *vel, int r, float spacing)
{
uint index = start;
for(int z=-r; z<=r; z++) {
for(int y=-r; y<=r; y++) {
for(int x=-r; x<=r; x++) {
float dx = x*spacing;
float dy = y*spacing;
float dz = z*spacing;
float l = sqrtf(dx*dx + dy*dy + dz*dz);
if ((l <= m_params.particleRadius*2.0f*r) && (index < m_numParticles)) {
m_hPos[index*4] = pos[0] + dx;
m_hPos[index*4+1] = pos[1] + dy;
m_hPos[index*4+2] = pos[2] + dz;
m_hPos[index*4+3] = pos[3];
m_hVel[index*4] = vel[0];
m_hVel[index*4+1] = vel[1];
m_hVel[index*4+2] = vel[2];
m_hVel[index*4+3] = vel[3];
index++;
}
}
}
}
setArray(POSITION, m_hPos, start, index);
setArray(VELOCITY, m_hVel, start, index);
} }
#include "../../Demos/OpenGL/GLDebugDrawer.h"
GLDebugDrawer debugDrawer;
void ParticleSystem::initializeBullet() void ParticleSystem::initializeBullet()
{ {
@@ -73,9 +662,8 @@ void ParticleSystem::initializeBullet()
m_collisionConfiguration = new btDefaultCollisionConfiguration(); m_collisionConfiguration = new btDefaultCollisionConfiguration();
m_dispatcher = new btCollisionDispatcher(m_collisionConfiguration); m_dispatcher = new btCollisionDispatcher(m_collisionConfiguration);
// m_broadphase = new btDbvtBroadphase(); // m_broadphase = new btDbvtBroadphase();
//m_broadphase = new btAxisSweep3(btVector3(-3,-3,-3),btVector3(3,3,3)); // m_broadphase = new btAxisSweep3(btVector3(-3,-3,-3),btVector3(3,3,3));
m_broadphase = new btCudaBroadphase(m_params,m_params.numBodies+6); m_broadphase = new btCudaBroadphase(btVector3(-1, -1, -1), btVector3(1, 1, 1), 64, 64, 64, m_params.numBodies, 16, 64, 8, btScalar(1.0f/1.733f));
m_constraintSolver=new btSequentialImpulseConstraintSolver(); m_constraintSolver=new btSequentialImpulseConstraintSolver();
@@ -84,27 +672,67 @@ void ParticleSystem::initializeBullet()
//debugDrawer.setDebugMode(btIDebugDraw::DBG_DrawPairs); //debugDrawer.setDebugMode(btIDebugDraw::DBG_DrawPairs);
m_dynamicsWorld->setGravity(100*btVector3(m_params.gravity.x,m_params.gravity.y,m_params.gravity.z)); // m_dynamicsWorld->setGravity(100*btVector3(m_params.gravity.x,m_params.gravity.y,m_params.gravity.z));
m_dynamicsWorld->setGravity(btScalar(ACC_DIR_FACT) * btVector3(m_params.gravity.x,m_params.gravity.y,m_params.gravity.z));
m_dynamicsWorld->getSolverInfo().m_numIterations=1; m_dynamicsWorld->getSolverInfo().m_numIterations=1;
btBoxShape* worldBox = new btBoxShape(btVector3(m_params.worldSize.x/2,m_params.worldSize.y/2,m_params.worldSize.z/2));
worldBox->setMargin(0.f);
//create 6 static planes for the world cube
btStaticPlaneShape* planeShape;
btRigidBody* body; btRigidBody* body;
btVector3 worldSize();
int i; btCollisionShape* boxShape = new btBoxShape(btVector3(btScalar(1.2),btScalar(0.05),btScalar(1.2)));
// boxShape->setMargin(0.03f);
btScalar mass(0.);
btVector3 localInertia(0,0,0);
btRigidBody::btRigidBodyConstructionInfo boxRbcInfo(mass, 0, boxShape, localInertia);
boxRbcInfo.m_startWorldTransform.setIdentity();
boxRbcInfo.m_startWorldTransform.setOrigin(btVector3(0, -1.05f,0));
boxRbcInfo.m_friction = 0.0f;
body = new btRigidBody(boxRbcInfo);
m_dynamicsWorld->addRigidBody(body);
boxRbcInfo.m_startWorldTransform.setIdentity();
boxRbcInfo.m_startWorldTransform.setOrigin(btVector3(0, 1.05f,0));
boxRbcInfo.m_friction = 0.0f;
body = new btRigidBody(boxRbcInfo);
m_dynamicsWorld->addRigidBody(body);
boxRbcInfo.m_startWorldTransform.setIdentity();
boxRbcInfo.m_startWorldTransform.getBasis().setEulerZYX(0, 0, SIMD_HALF_PI);
boxRbcInfo.m_startWorldTransform.setOrigin(btVector3(-1.05f, 0, 0));
boxRbcInfo.m_friction = 0.0f;
body = new btRigidBody(boxRbcInfo);
m_dynamicsWorld->addRigidBody(body);
boxRbcInfo.m_startWorldTransform.setIdentity();
boxRbcInfo.m_startWorldTransform.getBasis().setEulerZYX(0, 0, SIMD_HALF_PI);
boxRbcInfo.m_startWorldTransform.setOrigin(btVector3(1.05f, 0, 0));
boxRbcInfo.m_friction = 0.0f;
body = new btRigidBody(boxRbcInfo);
m_dynamicsWorld->addRigidBody(body);
boxRbcInfo.m_startWorldTransform.setIdentity();
boxRbcInfo.m_startWorldTransform.getBasis().setEulerZYX(SIMD_HALF_PI, 0, 0);
boxRbcInfo.m_startWorldTransform.setOrigin(btVector3(0, 0, -1.05f));
boxRbcInfo.m_friction = 0.0f;
body = new btRigidBody(boxRbcInfo);
m_dynamicsWorld->addRigidBody(body);
boxRbcInfo.m_startWorldTransform.setIdentity();
boxRbcInfo.m_startWorldTransform.getBasis().setEulerZYX(SIMD_HALF_PI, 0, 0);
boxRbcInfo.m_startWorldTransform.setOrigin(btVector3(0, 0, 1.05f));
boxRbcInfo.m_friction = 0.0f;
body = new btRigidBody(boxRbcInfo);
m_dynamicsWorld->addRigidBody(body);
unsigned int i;
btSphereShape* particleSphere = new btSphereShape(m_params.particleRadius); btSphereShape* particleSphere = new btSphereShape(m_params.particleRadius);
particleSphere->setMargin(0.0); particleSphere->setMargin(0.0);
btVector3 localInertia;
particleSphere->calculateLocalInertia(1,localInertia); particleSphere->calculateLocalInertia(1,localInertia);
float* m_hPos = m_broadphase->getHposPtr(); reset(CONFIG_GRID);
for (i=0;i<m_params.numBodies;i++) for (i=0;i<m_params.numBodies;i++)
{ {
@@ -116,7 +744,15 @@ void ParticleSystem::initializeBullet()
m_dynamicsWorld->addRigidBody(body); m_dynamicsWorld->addRigidBody(body);
} }
reset(CONFIG_GRID); btSphereShape* colliderSphere = new btSphereShape(m_params.colliderRadius);
colliderSphere->setMargin(0.0);
colliderSphere->calculateLocalInertia(10., localInertia);
btRigidBody::btRigidBodyConstructionInfo rbci(5., 0, colliderSphere,localInertia);
rbci.m_startWorldTransform.setOrigin(btVector3(m_params.colliderPos.x, m_params.colliderPos.y, m_params.colliderPos.z));
body = new btRigidBody(rbci);
body->setActivationState(DISABLE_DEACTIVATION);
m_bulletCollider = body;
m_dynamicsWorld->addRigidBody(body);
/* for (i=0;i<6;i++) /* for (i=0;i<6;i++)
{ {
@@ -130,7 +766,6 @@ void ParticleSystem::initializeBullet()
m_dynamicsWorld->addRigidBody(body); m_dynamicsWorld->addRigidBody(body);
} }
*/ */
} }
void ParticleSystem::finalizeBullet() void ParticleSystem::finalizeBullet()
@@ -142,139 +777,29 @@ void ParticleSystem::finalizeBullet()
delete m_collisionConfiguration; delete m_collisionConfiguration;
} }
float* ParticleSystem::copyBuffersFromDeviceToHost()
void
ParticleSystem::update(float deltaTime)
{ {
assert(m_bInitialized); copyArrayFromDevice(m_hVel, m_dVel[m_currentVelRead], 0, sizeof(float)*4*m_numParticles);
// fill color buffer
switch (m_simulationMode) glBindBufferARB(GL_ARRAY_BUFFER, m_posVbo[m_currentPosRead]);
{ float* hPosData = (float *) glMapBufferARB(GL_ARRAY_BUFFER, GL_READ_WRITE);//GL_WRITE_ONLY);
case SIMULATION_CUDA: return hPosData;
{ }
m_broadphase->quickHack(deltaTime);
//todo
break;
}
case SIMULATION_BULLET_CPU:
{
m_broadphase->integrate();
///copy particles from device to main memory
{
float* hPosData = m_broadphase->copyBuffersFromDeviceToHost();
float* m_hVel = m_broadphase->getHvelPtr();
m_broadphase->copyBuffersFromHostToDevice();
//sync transform and velocity from particle system to Bullet
for (int i=0;i<m_params.numBodies;i++)
{
btTransform& trans = m_bulletParticles[i]->getWorldTransform();
trans.setOrigin(btVector3(hPosData[i*4],hPosData[i*4+1],hPosData[i*4+2]));
m_bulletParticles[i]->setLinearVelocity(btVector3(m_hVel[i*4],m_hVel[i*4+1],m_hVel[i*4+2])*10.);
}
}
m_dynamicsWorld->stepSimulation(deltaTime);
/* for (int i=0;i<m_numParticles;i++)
{
data[i*4+1] -= 0.001f;
m_hVel[i*4]=0;
m_hVel[i*4+1]=0;
m_hVel[i*4+2]=0;
}
*/
{
float* hPosData = m_broadphase->copyBuffersFromDeviceToHost();
float* m_hVel = m_broadphase->getHvelPtr();
//sync transform and velocity from Bullet to particle system
for (int i=0;i<m_params.numBodies;i++)
{
btTransform& trans = m_bulletParticles[i]->getWorldTransform();
hPosData[i*4] = trans.getOrigin().getX();
hPosData[i*4+1] = trans.getOrigin().getY();
hPosData[i*4+2] = trans.getOrigin().getZ();
m_hVel[i*4] = m_bulletParticles[i]->getLinearVelocity().getX()/10.f;
m_hVel[i*4+1] = m_bulletParticles[i]->getLinearVelocity().getY()/10.f;
m_hVel[i*4+2] = m_bulletParticles[i]->getLinearVelocity().getZ()/10.f;
}
m_broadphase->copyBuffersFromHostToDevice();
}
break;
}
default:
{
printf("unknown simulation method\n");
}
};
void ParticleSystem::copyBuffersFromHostToDevice()
{
glUnmapBufferARB(GL_ARRAY_BUFFER);
copyArrayToDevice(m_dVel[m_currentVelRead],m_hVel, 0, sizeof(float)*4*m_numParticles);
} }
float* ParticleSystem::getArray(ParticleArray array)
{
return m_broadphase->getArray((btCudaBroadphase::ParticleArray)array);
}
void ParticleSystem::debugDraw() void ParticleSystem::debugDraw()
{ {
#if USE_BULLET
glDisable(GL_DEPTH_TEST); glDisable(GL_DEPTH_TEST);
m_dynamicsWorld->debugDrawWorld(); m_dynamicsWorld->debugDrawWorld();
glEnable(GL_DEPTH_TEST); glEnable(GL_DEPTH_TEST);
#endif
} }
void ParticleSystem::reset(ParticleConfig config)
{
m_broadphase->reset((btCudaBroadphase::ParticleConfig)config);
for (int i=0;i<m_bulletParticles.size();i++)
{
m_bulletParticles[i]->setAngularVelocity(btVector3(0,0,0));
}
}
void ParticleSystem::addSphere(int start, float *pos, float *vel, int r, float spacing)
{
m_broadphase->addSphere(start,pos,vel,r,spacing);
}
unsigned int ParticleSystem::getCurrentReadBuffer() const
{
return m_broadphase->getCurrentReadBuffer();
}
unsigned int ParticleSystem::getColorBuffer() const
{
return m_broadphase->getColorBuffer();
}
void ParticleSystem::dumpGrid()
{
return m_broadphase->dumpGrid();
}
void ParticleSystem::dumpParticles(uint start, uint count)
{
m_broadphase->dumpParticles(start,count);
}
int ParticleSystem::getNumParticles() const
{
return m_params.numBodies;
}

114
Extras/CUDA/particleSystem.cu
View File

@@ -41,7 +41,7 @@
#include <cuda_gl_interop.h> #include <cuda_gl_interop.h>
#include "particles_kernel.cu" #include "particles_kernel.cu"
#include "radixsort.cu" //#include "radixsort.cu"
//! Check for CUDA error //! Check for CUDA error
# define CUT_CHECK_ERROR(errorMessage) do { \ # define CUT_CHECK_ERROR(errorMessage) do { \
@@ -77,9 +77,9 @@
} } while (0) } } while (0)
extern "C" extern "C"
{ {
void mm_exit(int val) void mm_exit(int val)
{ {
exit(val); exit(val);
@@ -87,7 +87,7 @@ void mm_exit(int val)
void cudaInit(int argc, char **argv) void cudaInit(int argc, char **argv)
{ {
//CUT_DEVICE_INIT(argc, argv); // CUT_DEVICE_INIT(argc, argv);
} }
void allocateArray(void **devPtr, size_t size) void allocateArray(void **devPtr, size_t size)
@@ -117,26 +117,6 @@ void copyArrayFromDevice(void* host, const void* device, unsigned int vbo, int s
void copyArrayToDevice(void* device, const void* host, int offset, int size) void copyArrayToDevice(void* device, const void* host, int offset, int size)
{ {
MY_CUDA_SAFE_CALL(cudaMemcpy((char *) device + offset, host, size, cudaMemcpyHostToDevice)); MY_CUDA_SAFE_CALL(cudaMemcpy((char *) device + offset, host, size, cudaMemcpyHostToDevice));
/*
cudaError_t err = cudaMemcpy((char *) device + offset, host, size, cudaMemcpyHostToDevice);
switch(err)
{
case cudaSuccess :
return;
case cudaErrorInvalidValue :
printf("\ncudaErrorInvalidValue : %d\n", err);
return;
case cudaErrorInvalidDevicePointer :
printf("\ncudaErrorInvalidDevicePointer : %d\n", err);
return;
case cudaErrorInvalidMemcpyDirection :
printf("\ncudaErrorInvalidMemcpyDirection : %d\n", err);
return;
default :
printf("\nX3 : %d\n", err);
return;
}
*/
} }
void registerGLBufferObject(uint vbo) void registerGLBufferObject(uint vbo)
@@ -280,27 +260,6 @@ reorderDataAndFindCellStart(uint* particleHash,
MY_CUDA_SAFE_CALL(cudaGLUnmapBufferObject(vboOldPos)); MY_CUDA_SAFE_CALL(cudaGLUnmapBufferObject(vboOldPos));
} }
#if 1
void
findCellStart( uint* particleHash,
uint* cellStart,
uint numBodies,
uint numCells)
{
int numThreads, numBlocks;
computeGridSize(numBodies, 256, numBlocks, numThreads);
MY_CUDA_SAFE_CALL(cudaMemset(cellStart, 0xffffffff, numCells*sizeof(uint)));
findCellStartD<<< numBlocks, numThreads >>>(
(uint2 *) particleHash,
(uint *) cellStart);
CUT_CHECK_ERROR("Kernel execution failed: findCellStartD");
}
#endif
void void
collide(uint vboOldPos, uint vboNewPos, collide(uint vboOldPos, uint vboNewPos,
float* sortedPos, float* sortedVel, float* sortedPos, float* sortedVel,
@@ -374,71 +333,4 @@ collide(uint vboOldPos, uint vboNewPos,
#endif #endif
} }
void
btCudaFindOverlappingPairs( float* pAABB,
uint* pParticleHash,
uint* pCellStart,
uint* pPairBuff,
uint* pPairBuffStartCurr,
uint numParticles)
{
// cudaError err = cudaMemset(pPairBuff, 0x00, numParticles*32*4);
// if(err != cudaSuccess)
// {
// printf("\nAAAAA\n");
// }
MY_CUDA_SAFE_CALL(cudaBindTexture(0, pAABBTex, pAABB, numParticles*2*sizeof(float4)));
int numThreads, numBlocks;
// computeGridSize(numParticles, 256, numBlocks, numThreads);
computeGridSize(numParticles, 64, numBlocks, numThreads);
// numThreads = 1;
// numBlocks = 1;
btCudaFindOverlappingPairsD<<< numBlocks, numThreads >>>(
(float4 *)pAABB,
(uint2*)pParticleHash,
(uint*)pCellStart,
(uint*)pPairBuff,
(uint2*)pPairBuffStartCurr,
numParticles
);
CUT_CHECK_ERROR("Kernel execution failed: btCudaFindOverlappingPairsD");
MY_CUDA_SAFE_CALL(cudaUnbindTexture(pAABBTex));
} // btCudaFindOverlappingPairs()
void
btCudaComputePairCacheChanges(uint* pPairBuff, uint* pPairBuffStartCurr, uint* pPairScan, uint numParticles)
{
int numThreads, numBlocks;
computeGridSize(numParticles, 256, numBlocks, numThreads);
btCudaComputePairCacheChangesD<<< numBlocks, numThreads >>>(
(uint*)pPairBuff,
(uint2*)pPairBuffStartCurr,
(uint*)pPairScan
);
CUT_CHECK_ERROR("Kernel execution failed: btCudaComputePairCacheChangesD");
} // btCudaFindOverlappingPairs()
void
btCudaSqueezeOverlappingPairBuff(uint* pPairBuff, uint* pPairBuffStartCurr, uint* pPairScan, uint* pPairOut, uint numParticles)
{
int numThreads, numBlocks;
computeGridSize(numParticles, 256, numBlocks, numThreads);
btCudaSqueezeOverlappingPairBuffD<<< numBlocks, numThreads >>>(
(uint*)pPairBuff,
(uint2*)pPairBuffStartCurr,
(uint*)pPairScan,
pPairOut
);
CUT_CHECK_ERROR("Kernel execution failed: btCudaSqueezeOverlappingPairBuffD");
}
} // extern "C" } // extern "C"

21
Extras/CUDA/particleSystem.cuh
View File

@@ -42,12 +42,6 @@ reorderDataAndFindCellStart(uint* particleHash,
uint numBodies, uint numBodies,
uint numCells); uint numCells);
void
findCellStart( uint* particleHash,
uint* cellStart,
uint numBodies,
uint numCells);
void void
collide(uint vboOldPos, uint vboNewPos, collide(uint vboOldPos, uint vboNewPos,
float* sortedPos, float* sortedVel, float* sortedPos, float* sortedVel,
@@ -59,20 +53,5 @@ collide(uint vboOldPos, uint vboNewPos,
uint numBodies, uint numBodies,
uint numCells, uint numCells,
uint maxParticlesPerCell); uint maxParticlesPerCell);
void
btCudaFindOverlappingPairs( float* pAABB,
uint* pParticleHash,
uint* pCellStart,
uint* pPairBuff,
uint* pPairBuffStartCurr,
uint numParticles);
void
btCudaComputePairCacheChanges(uint* pPairBuff, uint* pPairBuffStartCurr, uint* pPairScan, uint numParticles);
void btCudaSqueezeOverlappingPairBuff(uint* pPairBuff, uint* pPairBuffStartCurr, uint* pPairScan, uint* pPairOut, uint numParticles);
} }

107
Extras/CUDA/particleSystem.h
View File

@@ -35,6 +35,7 @@
#include "particles_kernel.cuh" #include "particles_kernel.cuh"
#include "vector_functions.h" #include "vector_functions.h"
#include "LinearMath/btAlignedObjectArray.h" #include "LinearMath/btAlignedObjectArray.h"
// CUDA BodySystem: runs on the GPU // CUDA BodySystem: runs on the GPU
@@ -44,13 +45,6 @@ public:
ParticleSystem(uint numParticles, uint3 gridSize); ParticleSystem(uint numParticles, uint3 gridSize);
~ParticleSystem(); ~ParticleSystem();
enum ParticleArray
{
POSITION,
VELOCITY,
};
enum ParticleConfig enum ParticleConfig
{ {
CONFIG_RANDOM, CONFIG_RANDOM,
@@ -58,6 +52,12 @@ public:
_NUM_CONFIGS _NUM_CONFIGS
}; };
enum ParticleArray
{
POSITION,
VELOCITY,
};
enum SimulationMode enum SimulationMode
{ {
SIMULATION_CUDA, SIMULATION_CUDA,
@@ -65,33 +65,19 @@ public:
SIMULATION_NUM_MODES SIMULATION_NUM_MODES
}; };
void reset(ParticleConfig config);
void debugDraw();
///
///Bullet data
///
void initializeBullet();
void finalizeBullet();
class btDiscreteDynamicsWorld* m_dynamicsWorld;
class btDefaultCollisionConfiguration* m_collisionConfiguration;
class btCollisionDispatcher* m_dispatcher;
class btCudaBroadphase* m_broadphase;
// class btBroadphaseInterface* m_broadphase;
class btSequentialImpulseConstraintSolver* m_constraintSolver;
btAlignedObjectArray<class btRigidBody*> m_bulletParticles;
void update(float deltaTime); void update(float deltaTime);
void updateCuda(float deltaTime);
void updateBullet(float deltaTime);
void reset(ParticleConfig config);
float* getArray(ParticleArray array); float* getArray(ParticleArray array);
void setArray(ParticleArray array, const float* data, int start, int count);
int getNumParticles() const; int getNumParticles() const { return m_numParticles; }
unsigned int getCurrentReadBuffer() const; unsigned int getCurrentReadBuffer() const { return m_posVbo[m_currentPosRead]; }
unsigned int getColorBuffer() const; unsigned int getColorBuffer() const { return m_colorVBO; }
void dumpGrid(); void dumpGrid();
void dumpParticles(uint start, uint count); void dumpParticles(uint start, uint count);
@@ -127,21 +113,73 @@ public:
m_simulationMode=mode; m_simulationMode=mode;
} }
void debugDraw();
protected: // methods protected: // methods
ParticleSystem() ParticleSystem() {}
: m_simulationMode(SIMULATION_CUDA) uint createVBO(uint size);
{}
void _initialize(int numParticles);
void _finalize();
void initGrid(uint *size, float spacing, float jitter, uint numParticles); void initGrid(uint *size, float spacing, float jitter, uint numParticles);
protected:
// Bullet data
void initializeBullet();
void finalizeBullet();
class btDiscreteDynamicsWorld* m_dynamicsWorld;
class btDefaultCollisionConfiguration* m_collisionConfiguration;
class btCollisionDispatcher* m_dispatcher;
// class btCudaBroadphase* m_broadphase;
class btBroadphaseInterface* m_broadphase;
class btSequentialImpulseConstraintSolver* m_constraintSolver;
btAlignedObjectArray<class btRigidBody*> m_bulletParticles;
btRigidBody* m_bulletCollider;
float* copyBuffersFromDeviceToHost();
void copyBuffersFromHostToDevice();
protected: // data protected: // data
bool m_bInitialized; bool m_bInitialized;
uint m_numParticles;
// CPU data
float* m_hPos;
float* m_hVel;
uint* m_hGridCounters;
uint* m_hGridCells;
uint* m_hParticleHash;
uint* m_hCellStart;
// GPU data
float* m_dPos[2];
float* m_dVel[2];
float* m_dSortedPos;
float* m_dSortedVel;
// uniform grid data
uint* m_dGridCounters; // counts number of entries per grid cell
uint* m_dGridCells; // contains indices of up to "m_maxParticlesPerCell" particles per cell
uint* m_dParticleHash[2];
uint* m_dCellStart;
uint m_posVbo[2];
uint m_colorVBO;
uint m_currentPosRead, m_currentVelRead;
uint m_currentPosWrite, m_currentVelWrite;
// params // params
SimParams m_params; SimParams m_params;
uint3 m_gridSize;
uint m_numGridCells;
uint m_maxParticlesPerCell; uint m_maxParticlesPerCell;
uint m_timer; uint m_timer;
@@ -149,6 +187,7 @@ protected: // data
uint m_solverIterations; uint m_solverIterations;
SimulationMode m_simulationMode; SimulationMode m_simulationMode;
}; };
#endif // __BODYSYSTEMCUDA_H__ #endif // __BODYSYSTEMCUDA_H__

50
Extras/CUDA/particles.cpp
View File

@@ -45,7 +45,6 @@
#include <GL/glut.h> #include <GL/glut.h>
#endif #endif
#include "LinearMath/btQuickprof.h" #include "LinearMath/btQuickprof.h"
#include "particleSystem.h" #include "particleSystem.h"
@@ -62,8 +61,6 @@ float camera_rot_lag[] = {0, 0, 0};
const float inertia = 0.1; const float inertia = 0.1;
ParticleRenderer::DisplayMode displayMode = ParticleRenderer::PARTICLE_SPHERES; ParticleRenderer::DisplayMode displayMode = ParticleRenderer::PARTICLE_SPHERES;
int mode = 0; int mode = 0;
bool displayEnabled = true; bool displayEnabled = true;
bool bPause = false; bool bPause = false;
@@ -91,6 +88,9 @@ float collideAttraction = 0.0f;
ParticleSystem *psystem = 0; ParticleSystem *psystem = 0;
// fps // fps
static int fpsCount = 0;
static int fpsLimit = 1;
unsigned int timer;
ParticleRenderer *renderer = 0; ParticleRenderer *renderer = 0;
@@ -109,7 +109,7 @@ void init(int numParticles, uint3 gridSize)
renderer->setParticleRadius(psystem->getParticleRadius()); renderer->setParticleRadius(psystem->getParticleRadius());
renderer->setColorBuffer(psystem->getColorBuffer()); renderer->setColorBuffer(psystem->getColorBuffer());
// CUT_SAFE_CALL(cutCreateTimer(&timer));
} }
void initGL() void initGL()
@@ -129,7 +129,6 @@ void initGL()
void display() void display()
{ {
// update the simulation // update the simulation
if (!bPause) if (!bPause)
{ {
@@ -143,8 +142,6 @@ void display()
psystem->update(timestep); psystem->update(timestep);
renderer->setVertexBuffer(psystem->getCurrentReadBuffer(), psystem->getNumParticles()); renderer->setVertexBuffer(psystem->getCurrentReadBuffer(), psystem->getNumParticles());
float* posArray = psystem->getArray(ParticleSystem::POSITION);
renderer->setPositions(posArray,psystem->getNumParticles());
} }
// render // render
@@ -168,9 +165,7 @@ void display()
glColor3f(1.0, 1.0, 1.0); glColor3f(1.0, 1.0, 1.0);
glutWireCube(2.0); glutWireCube(2.0);
// collider
// collider
glPushMatrix(); glPushMatrix();
float4 p = psystem->getColliderPos(); float4 p = psystem->getColliderPos();
glTranslatef(p.x, p.y, p.z); glTranslatef(p.x, p.y, p.z);
@@ -195,18 +190,12 @@ void display()
psystem->debugDraw(); psystem->debugDraw();
glDisable(GL_DEPTH_TEST); glDisable(GL_DEPTH_TEST);
// glDisable(GL_LIGHTING);
// glColor3f(0, 0, 0);
float offsX = 10.f; float offsX = 10.f;
float offsY = 10.f; float offsY = 10.f;
renderer->showProfileInfo(offsX, offsY, 20.f); renderer->showProfileInfo(offsX, offsY, 20.f);
// glEnable(GL_LIGHTING);
glEnable(GL_DEPTH_TEST); glEnable(GL_DEPTH_TEST);
glutSwapBuffers(); glutSwapBuffers();
{ {
@@ -233,6 +222,7 @@ void display()
} }
glutReportErrors(); glutReportErrors();
} }
void reshape(int w, int h) void reshape(int w, int h)
@@ -375,7 +365,6 @@ inline float frand()
// commented out to remove unused parameter warnings in Linux // commented out to remove unused parameter warnings in Linux
void key(unsigned char key, int /*x*/, int /*y*/) void key(unsigned char key, int /*x*/, int /*y*/)
{ {
#ifndef BT_NO_PROFILE #ifndef BT_NO_PROFILE
if (key >= 0x31 && key < 0x37) if (key >= 0x31 && key < 0x37)
{ {
@@ -407,12 +396,9 @@ void key(unsigned char key, int /*x*/, int /*y*/)
case 'm': case 'm':
mode = M_MOVE; mode = M_MOVE;
break; break;
case 's':
psystem->setSimulationMode((ParticleSystem::SimulationMode) ((psystem->getSimulationMode() + 1) % ParticleSystem::SIMULATION_NUM_MODES));
break;
case 'p': case 'p':
displayMode = (ParticleRenderer::DisplayMode) ((displayMode + 1) % ParticleRenderer::PARTICLE_NUM_MODES); displayMode = (ParticleRenderer::DisplayMode)
((displayMode + 1) % ParticleRenderer::PARTICLE_NUM_MODES);
break; break;
case 'd': case 'd':
psystem->dumpGrid(); psystem->dumpGrid();
@@ -474,6 +460,10 @@ void key(unsigned char key, int /*x*/, int /*y*/)
case 'h': case 'h':
displaySliders = !displaySliders; displaySliders = !displaySliders;
break; break;
case 's':
psystem->setSimulationMode((ParticleSystem::SimulationMode) ((psystem->getSimulationMode() + 1) % ParticleSystem::SIMULATION_NUM_MODES));
CProfileManager::CleanupMemory();
break;
} }
glutPostRedisplay(); glutPostRedisplay();
@@ -536,16 +526,22 @@ void initMenus()
int int
main(int argc, char** argv) main(int argc, char** argv)
{ {
// numParticles =1024;//1024;//64;//16380;//32768; // numParticles = 65536*2;
numParticles =8192; // numParticles = 65536;
// numParticles = 32768;
// numParticles = 8192;
// numParticles = 4096;
numParticles = 2048;
// numParticles = 1024;
// numParticles = 256;
// numParticles = 32;
// numParticles = 2;
uint gridDim = 64; uint gridDim = 64;
numIterations = 0; numIterations = 0;
gridSize.x = gridSize.y = gridSize.z = gridDim; gridSize.x = gridSize.y = gridSize.z = gridDim;
printf("grid: %d x %d x %d = %d cells\n", gridSize.x, gridSize.y, gridSize.z, gridSize.x*gridSize.y*gridSize.z); printf("grid: %d x %d x %d = %d cells\n", gridSize.x, gridSize.y, gridSize.z, gridSize.x*gridSize.y*gridSize.z);
bool benchmark = false;
cudaInit(argc, argv); cudaInit(argc, argv);
glutInit(&argc, argv); glutInit(&argc, argv);
@@ -558,7 +554,6 @@ main(int argc, char** argv)
initParams(); initParams();
initMenus(); initMenus();
glutDisplayFunc(display); glutDisplayFunc(display);
glutReshapeFunc(reshape); glutReshapeFunc(reshape);
glutMouseFunc(mouse); glutMouseFunc(mouse);
@@ -569,7 +564,6 @@ main(int argc, char** argv)
glutMainLoop(); glutMainLoop();
if (psystem) if (psystem)
delete psystem; delete psystem;

228
Extras/CUDA/particles.sln Normal file
View File

@@ -0,0 +1,228 @@

Microsoft Visual Studio Solution File, Format Version 9.00
# Visual Studio 2005
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "particles", "particles.vcproj", "{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}"
ProjectSection(ProjectDependencies) = postProject
{F74E8E02-0B47-4816-BD0B-FAEAE3343165} = {F74E8E02-0B47-4816-BD0B-FAEAE3343165}
{7C428E76-9271-6284-20F0-9B38ED6931E3} = {7C428E76-9271-6284-20F0-9B38ED6931E3}
{61BD1097-CF2E-B296-DAA9-73A6FE135319} = {61BD1097-CF2E-B296-DAA9-73A6FE135319}
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A} = {7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE} = {6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}
EndProjectSection
EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "libbulletcollision", "..\..\msvc\8\libbulletcollision.vcproj", "{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}"
EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "libbulletmath", "..\..\msvc\8\libbulletmath.vcproj", "{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}"
EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "libbulletdynamics", "..\..\msvc\8\libbulletdynamics.vcproj", "{61BD1097-CF2E-B296-DAA9-73A6FE135319}"
EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "libbulletcuda", "libbulletcuda.vcproj", "{F74E8E02-0B47-4816-BD0B-FAEAE3343165}"
EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "libbulletopenglsupport", "..\..\msvc\8\libbulletopenglsupport.vcproj", "{7C428E76-9271-6284-20F0-9B38ED6931E3}"
EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "appBasicDemo", "..\..\msvc\8\appBasicDemo.vcproj", "{3578834A-4B06-DE6F-78AC-FE11F7226D35}"
ProjectSection(ProjectDependencies) = postProject
{F74E8E02-0B47-4816-BD0B-FAEAE3343165} = {F74E8E02-0B47-4816-BD0B-FAEAE3343165}
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE} = {6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A} = {7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}
{61BD1097-CF2E-B296-DAA9-73A6FE135319} = {61BD1097-CF2E-B296-DAA9-73A6FE135319}
{7C428E76-9271-6284-20F0-9B38ED6931E3} = {7C428E76-9271-6284-20F0-9B38ED6931E3}
EndProjectSection
EndProject
Global
GlobalSection(SolutionConfigurationPlatforms) = preSolution
Debug|Win32 = Debug|Win32
Debug|x64 = Debug|x64
DebugDll|Win32 = DebugDll|Win32
DebugDll|x64 = DebugDll|x64
DebugDoublePrecision|Win32 = DebugDoublePrecision|Win32
DebugDoublePrecision|x64 = DebugDoublePrecision|x64
EmuDebug|Win32 = EmuDebug|Win32
EmuDebug|x64 = EmuDebug|x64
EmuRelease|Win32 = EmuRelease|Win32
EmuRelease|x64 = EmuRelease|x64
Release|Win32 = Release|Win32
Release|x64 = Release|x64
ReleaseDll|Win32 = ReleaseDll|Win32
ReleaseDll|x64 = ReleaseDll|x64
ReleaseDoublePrecision|Win32 = ReleaseDoublePrecision|Win32
ReleaseDoublePrecision|x64 = ReleaseDoublePrecision|x64
EndGlobalSection
GlobalSection(ProjectConfigurationPlatforms) = postSolution
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Debug|Win32.ActiveCfg = Debug|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Debug|Win32.Build.0 = Debug|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Debug|x64.ActiveCfg = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Debug|x64.Build.0 = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.DebugDll|Win32.ActiveCfg = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.DebugDll|x64.ActiveCfg = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.DebugDll|x64.Build.0 = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.DebugDoublePrecision|Win32.ActiveCfg = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.DebugDoublePrecision|x64.ActiveCfg = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.DebugDoublePrecision|x64.Build.0 = Debug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuDebug|Win32.ActiveCfg = EmuDebug|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuDebug|Win32.Build.0 = EmuDebug|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuDebug|x64.ActiveCfg = EmuDebug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuDebug|x64.Build.0 = EmuDebug|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuRelease|Win32.ActiveCfg = EmuRelease|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuRelease|Win32.Build.0 = EmuRelease|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuRelease|x64.ActiveCfg = EmuRelease|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.EmuRelease|x64.Build.0 = EmuRelease|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Release|Win32.ActiveCfg = Release|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Release|Win32.Build.0 = Release|Win32
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Release|x64.ActiveCfg = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.Release|x64.Build.0 = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.ReleaseDll|Win32.ActiveCfg = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.ReleaseDll|x64.ActiveCfg = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.ReleaseDll|x64.Build.0 = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.ReleaseDoublePrecision|Win32.ActiveCfg = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.ReleaseDoublePrecision|x64.ActiveCfg = Release|x64
{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}.ReleaseDoublePrecision|x64.Build.0 = Release|x64
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Debug|Win32.ActiveCfg = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Debug|Win32.Build.0 = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Debug|x64.ActiveCfg = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.DebugDll|Win32.ActiveCfg = DebugDll|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.DebugDll|Win32.Build.0 = DebugDll|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.DebugDll|x64.ActiveCfg = DebugDll|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.DebugDoublePrecision|Win32.ActiveCfg = DebugDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.DebugDoublePrecision|Win32.Build.0 = DebugDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.DebugDoublePrecision|x64.ActiveCfg = DebugDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuDebug|Win32.ActiveCfg = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuDebug|Win32.Build.0 = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuDebug|x64.ActiveCfg = Debug|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuRelease|Win32.ActiveCfg = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuRelease|Win32.Build.0 = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.EmuRelease|x64.ActiveCfg = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Release|Win32.ActiveCfg = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Release|Win32.Build.0 = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.Release|x64.ActiveCfg = Release|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.ReleaseDll|Win32.ActiveCfg = ReleaseDll|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.ReleaseDll|Win32.Build.0 = ReleaseDll|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.ReleaseDll|x64.ActiveCfg = ReleaseDll|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.ReleaseDoublePrecision|Win32.ActiveCfg = ReleaseDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.ReleaseDoublePrecision|Win32.Build.0 = ReleaseDoublePrecision|Win32
{6ADA430D-009C-2ED4-A787-2AC2D6FEB8CE}.ReleaseDoublePrecision|x64.ActiveCfg = ReleaseDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Debug|Win32.ActiveCfg = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Debug|Win32.Build.0 = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Debug|x64.ActiveCfg = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.DebugDll|Win32.ActiveCfg = DebugDll|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.DebugDll|Win32.Build.0 = DebugDll|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.DebugDll|x64.ActiveCfg = DebugDll|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.DebugDoublePrecision|Win32.ActiveCfg = DebugDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.DebugDoublePrecision|Win32.Build.0 = DebugDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.DebugDoublePrecision|x64.ActiveCfg = DebugDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuDebug|Win32.ActiveCfg = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuDebug|Win32.Build.0 = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuDebug|x64.ActiveCfg = Debug|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuRelease|Win32.ActiveCfg = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuRelease|Win32.Build.0 = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.EmuRelease|x64.ActiveCfg = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Release|Win32.ActiveCfg = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Release|Win32.Build.0 = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.Release|x64.ActiveCfg = Release|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.ReleaseDll|Win32.ActiveCfg = ReleaseDll|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.ReleaseDll|Win32.Build.0 = ReleaseDll|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.ReleaseDll|x64.ActiveCfg = ReleaseDll|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.ReleaseDoublePrecision|Win32.ActiveCfg = ReleaseDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.ReleaseDoublePrecision|Win32.Build.0 = ReleaseDoublePrecision|Win32
{7D6E339F-9C2C-31DA-FDB0-5EE50973CF2A}.ReleaseDoublePrecision|x64.ActiveCfg = ReleaseDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Debug|Win32.ActiveCfg = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Debug|Win32.Build.0 = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Debug|x64.ActiveCfg = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.DebugDll|Win32.ActiveCfg = DebugDll|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.DebugDll|Win32.Build.0 = DebugDll|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.DebugDll|x64.ActiveCfg = DebugDll|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.DebugDoublePrecision|Win32.ActiveCfg = DebugDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.DebugDoublePrecision|Win32.Build.0 = DebugDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.DebugDoublePrecision|x64.ActiveCfg = DebugDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuDebug|Win32.ActiveCfg = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuDebug|Win32.Build.0 = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuDebug|x64.ActiveCfg = Debug|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuRelease|Win32.ActiveCfg = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuRelease|Win32.Build.0 = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.EmuRelease|x64.ActiveCfg = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Release|Win32.ActiveCfg = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Release|Win32.Build.0 = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.Release|x64.ActiveCfg = Release|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.ReleaseDll|Win32.ActiveCfg = ReleaseDll|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.ReleaseDll|Win32.Build.0 = ReleaseDll|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.ReleaseDll|x64.ActiveCfg = ReleaseDll|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.ReleaseDoublePrecision|Win32.ActiveCfg = ReleaseDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.ReleaseDoublePrecision|Win32.Build.0 = ReleaseDoublePrecision|Win32
{61BD1097-CF2E-B296-DAA9-73A6FE135319}.ReleaseDoublePrecision|x64.ActiveCfg = ReleaseDoublePrecision|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.Debug|Win32.ActiveCfg = Debug|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.Debug|Win32.Build.0 = Debug|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.Debug|x64.ActiveCfg = Debug|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.DebugDll|Win32.ActiveCfg = DebugDll|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.DebugDll|Win32.Build.0 = DebugDll|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.DebugDll|x64.ActiveCfg = DebugDll|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.DebugDoublePrecision|Win32.ActiveCfg = DebugDoublePrecision|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.DebugDoublePrecision|Win32.Build.0 = DebugDoublePrecision|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.DebugDoublePrecision|x64.ActiveCfg = DebugDoublePrecision|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.EmuDebug|Win32.ActiveCfg = Debug|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.EmuDebug|Win32.Build.0 = Debug|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.EmuDebug|x64.ActiveCfg = Debug|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.EmuRelease|Win32.ActiveCfg = Release|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.EmuRelease|Win32.Build.0 = Release|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.EmuRelease|x64.ActiveCfg = Release|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.Release|Win32.ActiveCfg = Release|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.Release|Win32.Build.0 = Release|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.Release|x64.ActiveCfg = Release|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.ReleaseDll|Win32.ActiveCfg = ReleaseDll|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.ReleaseDll|Win32.Build.0 = ReleaseDll|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.ReleaseDll|x64.ActiveCfg = ReleaseDll|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.ReleaseDoublePrecision|Win32.ActiveCfg = ReleaseDoublePrecision|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.ReleaseDoublePrecision|Win32.Build.0 = ReleaseDoublePrecision|Win32
{F74E8E02-0B47-4816-BD0B-FAEAE3343165}.ReleaseDoublePrecision|x64.ActiveCfg = ReleaseDoublePrecision|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.Debug|Win32.ActiveCfg = Debug|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.Debug|Win32.Build.0 = Debug|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.Debug|x64.ActiveCfg = Debug|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.DebugDll|Win32.ActiveCfg = DebugDll|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.DebugDll|Win32.Build.0 = DebugDll|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.DebugDll|x64.ActiveCfg = DebugDll|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.DebugDoublePrecision|Win32.ActiveCfg = DebugDoublePrecision|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.DebugDoublePrecision|Win32.Build.0 = DebugDoublePrecision|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.DebugDoublePrecision|x64.ActiveCfg = DebugDoublePrecision|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.EmuDebug|Win32.ActiveCfg = Debug|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.EmuDebug|Win32.Build.0 = Debug|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.EmuDebug|x64.ActiveCfg = Debug|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.EmuRelease|Win32.ActiveCfg = Release|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.EmuRelease|Win32.Build.0 = Release|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.EmuRelease|x64.ActiveCfg = Release|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.Release|Win32.ActiveCfg = Release|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.Release|Win32.Build.0 = Release|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.Release|x64.ActiveCfg = Release|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.ReleaseDll|Win32.ActiveCfg = ReleaseDll|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.ReleaseDll|Win32.Build.0 = ReleaseDll|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.ReleaseDll|x64.ActiveCfg = ReleaseDll|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.ReleaseDoublePrecision|Win32.ActiveCfg = ReleaseDoublePrecision|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.ReleaseDoublePrecision|Win32.Build.0 = ReleaseDoublePrecision|Win32
{7C428E76-9271-6284-20F0-9B38ED6931E3}.ReleaseDoublePrecision|x64.ActiveCfg = ReleaseDoublePrecision|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.Debug|Win32.ActiveCfg = Debug|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.Debug|Win32.Build.0 = Debug|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.Debug|x64.ActiveCfg = Debug|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.DebugDll|Win32.ActiveCfg = DebugDll|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.DebugDll|Win32.Build.0 = DebugDll|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.DebugDll|x64.ActiveCfg = DebugDll|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.DebugDoublePrecision|Win32.ActiveCfg = DebugDoublePrecision|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.DebugDoublePrecision|Win32.Build.0 = DebugDoublePrecision|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.DebugDoublePrecision|x64.ActiveCfg = DebugDoublePrecision|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.EmuDebug|Win32.ActiveCfg = Debug|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.EmuDebug|Win32.Build.0 = Debug|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.EmuDebug|x64.ActiveCfg = Debug|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.EmuRelease|Win32.ActiveCfg = Release|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.EmuRelease|Win32.Build.0 = Release|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.EmuRelease|x64.ActiveCfg = Release|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.Release|Win32.ActiveCfg = Release|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.Release|Win32.Build.0 = Release|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.Release|x64.ActiveCfg = Release|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.ReleaseDll|Win32.ActiveCfg = ReleaseDll|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.ReleaseDll|Win32.Build.0 = ReleaseDll|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.ReleaseDll|x64.ActiveCfg = ReleaseDll|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.ReleaseDoublePrecision|Win32.ActiveCfg = ReleaseDoublePrecision|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.ReleaseDoublePrecision|Win32.Build.0 = ReleaseDoublePrecision|Win32
{3578834A-4B06-DE6F-78AC-FE11F7226D35}.ReleaseDoublePrecision|x64.ActiveCfg = ReleaseDoublePrecision|Win32
EndGlobalSection
GlobalSection(SolutionProperties) = preSolution
HideSolutionNode = FALSE
EndGlobalSection
EndGlobal

197
Extras/CUDA/btCudaBroadphase.vcproj → Extras/CUDA/particles.vcproj
View File

@@ -2,9 +2,9 @@
<VisualStudioProject <VisualStudioProject
ProjectType="Visual C++" ProjectType="Visual C++"
Version="8.00" Version="8.00"
Name="btCudaBroadphase" Name="particles"
ProjectGUID="{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}" ProjectGUID="{AF7F45C7-8545-4CA2-B835-FDE8823B7B09}"
RootNamespace="btCudaBroadphase" RootNamespace="particles"
Keyword="Win32Proj" Keyword="Win32Proj"
> >
<Platforms> <Platforms>
@@ -45,7 +45,7 @@
Name="VCCLCompilerTool" Name="VCCLCompilerTool"
Optimization="0" Optimization="0"
AdditionalIncludeDirectories="../../Glut;&quot;$(CUDA_INC_PATH)&quot;;./;../../src;../../Demos/OpenGL" AdditionalIncludeDirectories="../../Glut;&quot;$(CUDA_INC_PATH)&quot;;./;../../src;../../Demos/OpenGL"
PreprocessorDefinitions="WIN32;_DEBUG;_CONSOLE" PreprocessorDefinitions="WIN32;_DEBUG;_CONSOLE;_CRT_SECURE_NO_WARNINGS"
MinimalRebuild="true" MinimalRebuild="true"
BasicRuntimeChecks="3" BasicRuntimeChecks="3"
RuntimeLibrary="1" RuntimeLibrary="1"
@@ -70,7 +70,7 @@
LinkIncremental="1" LinkIncremental="1"
AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../Glut" AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../Glut"
GenerateDebugInformation="true" GenerateDebugInformation="true"
ProgramDatabaseFile="$(OutDir)/btCudaBroadphase.pdb" ProgramDatabaseFile="$(OutDir)/particles.pdb"
SubSystem="1" SubSystem="1"
TargetMachine="1" TargetMachine="1"
/> />
@@ -149,11 +149,11 @@
<Tool <Tool
Name="VCLinkerTool" Name="VCLinkerTool"
AdditionalDependencies="cudart.lib cutil64D.lib glew64.lib glut64.lib" AdditionalDependencies="cudart.lib cutil64D.lib glew64.lib glut64.lib"
OutputFile="../../bin/win64/$(ConfigurationName)/btCudaBroadphase.exe" OutputFile="../../bin/win64/$(ConfigurationName)/particles.exe"
LinkIncremental="1" LinkIncremental="1"
AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../common/lib" AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../common/lib"
GenerateDebugInformation="true" GenerateDebugInformation="true"
ProgramDatabaseFile="$(OutDir)/btCudaBroadphase.pdb" ProgramDatabaseFile="$(OutDir)/particles.pdb"
SubSystem="1" SubSystem="1"
TargetMachine="17" TargetMachine="17"
/> />
@@ -209,7 +209,7 @@
<Tool <Tool
Name="VCCLCompilerTool" Name="VCCLCompilerTool"
AdditionalIncludeDirectories="../../Glut;&quot;$(CUDA_INC_PATH)&quot;;./;../../src;../../Demos/OpenGL" AdditionalIncludeDirectories="../../Glut;&quot;$(CUDA_INC_PATH)&quot;;./;../../src;../../Demos/OpenGL"
PreprocessorDefinitions="WIN32;_CONSOLE" PreprocessorDefinitions="WIN32;_CONSOLE;_CRT_SECURE_NO_WARNINGS"
RuntimeLibrary="0" RuntimeLibrary="0"
UsePrecompiledHeader="0" UsePrecompiledHeader="0"
WarningLevel="3" WarningLevel="3"
@@ -227,8 +227,8 @@
/> />
<Tool <Tool
Name="VCLinkerTool" Name="VCLinkerTool"
AdditionalDependencies="cudart.lib glew32.lib" AdditionalDependencies="cudart.lib glew32.lib"
OutputFile="../../ReleaseCudaBroadphase.exe" OutputFile="../../ReleaseCudaParticles.exe"
LinkIncremental="1" LinkIncremental="1"
AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../Glut" AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../Glut"
GenerateDebugInformation="true" GenerateDebugInformation="true"
@@ -309,7 +309,7 @@
<Tool <Tool
Name="VCLinkerTool" Name="VCLinkerTool"
AdditionalDependencies="cudart.lib cutil64.lib glew64.lib glut64.lib" AdditionalDependencies="cudart.lib cutil64.lib glew64.lib glut64.lib"
OutputFile="../../bin/win64/$(ConfigurationName)/btCudaBroadphase.exe" OutputFile="../../bin/win64/$(ConfigurationName)/particles.exe"
LinkIncremental="1" LinkIncremental="1"
AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../common/lib" AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../common/lib"
GenerateDebugInformation="true" GenerateDebugInformation="true"
@@ -370,7 +370,7 @@
<Tool <Tool
Name="VCCLCompilerTool" Name="VCCLCompilerTool"
Optimization="0" Optimization="0"
AdditionalIncludeDirectories="../../Glut;&quot;$(CUDA_INC_PATH)&quot;;./;../../src;../../Demos/OpenGL" AdditionalIncludeDirectories="../../../Glut;&quot;$(CUDA_INC_PATH)&quot;;./;../../../src;../../../Demos/OpenGL"
PreprocessorDefinitions="WIN32;_DEBUG;_CONSOLE" PreprocessorDefinitions="WIN32;_DEBUG;_CONSOLE"
MinimalRebuild="true" MinimalRebuild="true"
BasicRuntimeChecks="3" BasicRuntimeChecks="3"
@@ -391,12 +391,12 @@
/> />
<Tool <Tool
Name="VCLinkerTool" Name="VCLinkerTool"
AdditionalDependencies="cudart.lib glew32.lib" AdditionalDependencies="cudart.lib cutil32D.lib glew32.lib"
OutputFile="../../EmuDebugCudaBroadphase.exe" OutputFile="../../bin/win32/$(ConfigurationName)/particles.exe"
LinkIncremental="1" LinkIncremental="1"
AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../Glut" AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../common/lib"
GenerateDebugInformation="true" GenerateDebugInformation="true"
ProgramDatabaseFile="$(OutDir)/btCudaBroadphase.pdb" ProgramDatabaseFile="$(OutDir)/particles.pdb"
SubSystem="1" SubSystem="1"
TargetMachine="1" TargetMachine="1"
/> />
@@ -475,11 +475,11 @@
<Tool <Tool
Name="VCLinkerTool" Name="VCLinkerTool"
AdditionalDependencies="cudart.lib cutil64D.lib glew64.lib glut64.lib" AdditionalDependencies="cudart.lib cutil64D.lib glew64.lib glut64.lib"
OutputFile="../../bin/win64/$(ConfigurationName)/btCudaBroadphase.exe" OutputFile="../../bin/win64/$(ConfigurationName)/particles.exe"
LinkIncremental="1" LinkIncremental="1"
AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../common/lib" AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../common/lib"
GenerateDebugInformation="true" GenerateDebugInformation="true"
ProgramDatabaseFile="$(OutDir)/btCudaBroadphase.pdb" ProgramDatabaseFile="$(OutDir)/particles.pdb"
SubSystem="1" SubSystem="1"
TargetMachine="17" TargetMachine="17"
/> />
@@ -535,7 +535,7 @@
<Tool <Tool
Name="VCCLCompilerTool" Name="VCCLCompilerTool"
Optimization="0" Optimization="0"
AdditionalIncludeDirectories="../../Glut;&quot;$(CUDA_INC_PATH)&quot;;./;../../src;../../Demos/OpenGL" AdditionalIncludeDirectories="../../../Glut;&quot;$(CUDA_INC_PATH)&quot;;./;../../../src;../../../Demos/OpenGL"
PreprocessorDefinitions="WIN32;_CONSOLE" PreprocessorDefinitions="WIN32;_CONSOLE"
RuntimeLibrary="0" RuntimeLibrary="0"
UsePrecompiledHeader="0" UsePrecompiledHeader="0"
@@ -554,10 +554,10 @@
/> />
<Tool <Tool
Name="VCLinkerTool" Name="VCLinkerTool"
AdditionalDependencies="cudart.lib glew32.lib" AdditionalDependencies="cudart.lib cutil32.lib glew32.lib"
OutputFile="../../EmuReleaseCudaBroadphase.exe" OutputFile="../../bin/win32/$(ConfigurationName)/particles.exe"
LinkIncremental="1" LinkIncremental="1"
AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../Glut" AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../common/lib"
GenerateDebugInformation="true" GenerateDebugInformation="true"
SubSystem="1" SubSystem="1"
OptimizeReferences="2" OptimizeReferences="2"
@@ -637,7 +637,7 @@
<Tool <Tool
Name="VCLinkerTool" Name="VCLinkerTool"
AdditionalDependencies="cudart.lib cutil64.lib glew64.lib glut64.lib" AdditionalDependencies="cudart.lib cutil64.lib glew64.lib glut64.lib"
OutputFile="../../bin/win64/$(ConfigurationName)/btCudaBroadphase.exe" OutputFile="../../bin/win64/$(ConfigurationName)/particles.exe"
LinkIncremental="1" LinkIncremental="1"
AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../common/lib" AdditionalLibraryDirectories="$(CUDA_LIB_PATH);../../common/lib"
GenerateDebugInformation="true" GenerateDebugInformation="true"
@@ -681,26 +681,6 @@
Filter="cu;cpp;c;cxx;def;odl;idl;hpj;bat;asm;asmx" Filter="cu;cpp;c;cxx;def;odl;idl;hpj;bat;asm;asmx"
UniqueIdentifier="{4FC737F1-C7A5-4376-A066-2A32D752A2FF}" UniqueIdentifier="{4FC737F1-C7A5-4376-A066-2A32D752A2FF}"
> >
<File
RelativePath="..\..\Demos\OpenGL\BMF_Api.cpp"
>
</File>
<File
RelativePath="..\..\Demos\OpenGL\BMF_BitmapFont.cpp"
>
</File>
<File
RelativePath="..\..\Demos\OpenGL\BMF_font_helv10.cpp"
>
</File>
<File
RelativePath=".\btCudaBroadphase.cpp"
>
</File>
<File
RelativePath="..\..\Demos\OpenGL\GLDebugDrawer.cpp"
>
</File>
<File <File
RelativePath=".\paramgl.cpp" RelativePath=".\paramgl.cpp"
> >
@@ -793,7 +773,7 @@
> >
<Tool <Tool
Name="VCCustomBuildTool" Name="VCCustomBuildTool"
CommandLine="&quot;$(CUDA_BIN_PATH)\nvcc.exe&quot; -arch=sm_10 -ccbin &quot;$(VCInstallDir)\bin&quot; -c -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd -I &quot;$(CUDA_INC_PATH)&quot;; -I./ -Icommon/inc -I../../Glut -o $(ConfigurationName)\particleSystem_cu.obj particleSystem.cu&#x0D;&#x0A;" CommandLine="&quot;$(CUDA_BIN_PATH)\nvcc.exe&quot; -arch=sm_10 -ccbin &quot;$(VCInstallDir)\bin&quot; -c -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd -I &quot;$(CUDA_INC_PATH)&quot;; -I./ -I../../Glut -o $(ConfigurationName)\particleSystem_cu.obj particleSystem.cu&#x0D;&#x0A;"
AdditionalDependencies="particleSystem.cuh; particles_kernel.cu; radixsort.cu" AdditionalDependencies="particleSystem.cuh; particles_kernel.cu; radixsort.cu"
Outputs="$(ConfigurationName)\particleSystem_cu.obj" Outputs="$(ConfigurationName)\particleSystem_cu.obj"
/> />
@@ -813,7 +793,7 @@
> >
<Tool <Tool
Name="VCCustomBuildTool" Name="VCCustomBuildTool"
CommandLine="&quot;$(CUDA_BIN_PATH)\nvcc.exe&quot; -arch=sm_10 -use_fast_math -ccbin &quot;$(VCInstallDir)\bin&quot; -c -I &quot;$(CUDA_INC_PATH)&quot;; -I./ -Icommon/inc -I../../Glut -o $(ConfigurationName)\particleSystem_cu.obj particleSystem.cu&#x0D;&#x0A;" CommandLine="&quot;$(CUDA_BIN_PATH)\nvcc.exe&quot; -arch=sm_10 -use_fast_math -ccbin &quot;$(VCInstallDir)\bin&quot; -c -I &quot;$(CUDA_INC_PATH)&quot;; -I./ -I../../Glut -o $(ConfigurationName)\particleSystem_cu.obj particleSystem.cu&#x0D;&#x0A;"
AdditionalDependencies="particleSystem.cuh; particles_kernel.cu; radixsort.cu" AdditionalDependencies="particleSystem.cuh; particles_kernel.cu; radixsort.cu"
Outputs="$(ConfigurationName)\particleSystem_cu.obj" Outputs="$(ConfigurationName)\particleSystem_cu.obj"
/> />
@@ -877,95 +857,6 @@
RelativePath=".\particleSystem.h" RelativePath=".\particleSystem.h"
> >
</File> </File>
<File
RelativePath=".\radixsort.cu"
>
<FileConfiguration
Name="Debug|Win32"
ExcludedFromBuild="true"
>
<Tool
Name="VCCLCompilerTool"
/>
</FileConfiguration>
<FileConfiguration
Name="Release|Win32"
ExcludedFromBuild="true"
>
<Tool
Name="VCCustomBuildTool"
CommandLine=""
AdditionalDependencies=""
Outputs=""
/>
</FileConfiguration>
<FileConfiguration
Name="Release|x64"
>
<Tool
Name="VCCustomBuildTool"
CommandLine=""
AdditionalDependencies=""
Outputs=""
/>
</FileConfiguration>
<FileConfiguration
Name="EmuDebug|Win32"
ExcludedFromBuild="true"
>
<Tool
Name="VCCLCompilerTool"
/>
</FileConfiguration>
<FileConfiguration
Name="EmuRelease|Win32"
ExcludedFromBuild="true"
>
<Tool
Name="VCCLCompilerTool"
/>
</FileConfiguration>
</File>
<File
RelativePath=".\radixsort.cuh"
>
</File>
<File
RelativePath=".\radixsort_kernel.cu"
>
<FileConfiguration
Name="Debug|Win32"
ExcludedFromBuild="true"
>
<Tool
Name="VCCLCompilerTool"
/>
</FileConfiguration>
<FileConfiguration
Name="Release|Win32"
ExcludedFromBuild="true"
>
<Tool
Name="VCCLCompilerTool"
/>
</FileConfiguration>
<FileConfiguration
Name="EmuDebug|Win32"
ExcludedFromBuild="true"
>
<Tool
Name="VCCLCompilerTool"
/>
</FileConfiguration>
<FileConfiguration
Name="EmuRelease|Win32"
ExcludedFromBuild="true"
>
<Tool
Name="VCCLCompilerTool"
/>
</FileConfiguration>
</File>
<File <File
RelativePath=".\render_particles.cpp" RelativePath=".\render_particles.cpp"
> >
@@ -983,46 +874,6 @@
> >
</File> </File>
</Filter> </Filter>
<File
RelativePath="..\..\Demos\OpenGL\BMF_Api.h"
>
</File>
<File
RelativePath="..\..\Demos\OpenGL\BMF_BitmapFont.h"
>
</File>
<File
RelativePath="..\..\Demos\OpenGL\BMF_FontData.h"
>
</File>
<File
RelativePath="..\..\Demos\OpenGL\BMF_Fonts.h"
>
</File>
<File
RelativePath=".\btCudaBroadphase.h"
>
</File>
<File
RelativePath=".\VTune\btCudaBroadphase.vpj"
>
</File>
<File
RelativePath=".\cutil_math.h"
>
</File>
<File
RelativePath="..\..\Demos\OpenGL\GLDebugDrawer.h"
>
</File>
<File
RelativePath=".\param.h"
>
</File>
<File
RelativePath=".\paramgl.h"
>
</File>
</Files> </Files>
<Globals> <Globals>
</Globals> </Globals>

196
Extras/CUDA/particles_kernel.cu
View File

@@ -50,10 +50,6 @@ texture<uint, 1, cudaReadModeElementType> cellStartTex;
texture<uint, 1, cudaReadModeElementType> gridCountersTex; texture<uint, 1, cudaReadModeElementType> gridCountersTex;
texture<uint, 1, cudaReadModeElementType> gridCellsTex; texture<uint, 1, cudaReadModeElementType> gridCellsTex;
texture<float4, 1, cudaReadModeElementType> pAABBTex;
#endif #endif
__constant__ SimParams params; __constant__ SimParams params;
@@ -203,35 +199,6 @@ reorderDataAndFindCellStartD(uint2* particleHash, // particle id sorted by has
} }
__global__ void
findCellStartD(uint2* particleHash, // particle id sorted by hash
uint* cellStart)
{
int index = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
uint2 sortedData = particleHash[index];
// Load hash data into shared memory so that we can look
// at neighboring particle's hash value without loading
// two hash values per thread
__shared__ uint sharedHash[257];
sharedHash[threadIdx.x+1] = sortedData.x;
if (index > 0 && threadIdx.x == 0)
{
// first thread in block must load neighbor particle hash
volatile uint2 prevData = particleHash[index-1];
sharedHash[0] = prevData.x;
}
__syncthreads();
if (index == 0 || sortedData.x != sharedHash[threadIdx.x])
{
cellStart[sortedData.x] = index;
}
}
// collide two spheres using DEM method // collide two spheres using DEM method
__device__ float3 collideSpheres(float4 posA, float4 posB, __device__ float3 collideSpheres(float4 posA, float4 posB,
float4 velA, float4 velB, float4 velA, float4 velB,
@@ -411,167 +378,4 @@ collideD(float4* newPos, float4* newVel,
#endif #endif
} }
__device__
uint cudaTestAABBOverlap(float4 min0, float4 max0, float4 min1, float4 max1)
{
return (min0.x <= max1.x)&& (min1.x <= max0.x) &&
(min0.y <= max1.y)&& (min1.y <= max0.y) &&
(min0.z <= max1.z)&& (min1.z <= max0.z);
}
__device__
void findPairsInCell( int3 gridPos,
uint index,
uint2* pParticleHash,
uint* pCellStart,
float4* pAABB,
uint* pPairBuff,
uint2* pPairBuffStartCurr,
uint numParticles)
{
if ((gridPos.x < 0) || (gridPos.x > params.gridSize.x-1) ||
(gridPos.y < 0) || (gridPos.y > params.gridSize.y-1) ||
(gridPos.z < 0) || (gridPos.z > params.gridSize.z-1)) {
return;
}
uint gridHash = calcGridHash(gridPos);
// get start of bucket for this cell
uint bucketStart = pCellStart[gridHash];
if (bucketStart == 0xffffffff)
return; // cell empty
// iterate over particles in this cell
// float4 min0 = pAABB[index*2];
// float4 max0 = pAABB[index*2+1];
float4 min0 = FETCH(pAABB, index*2); // pAABB[index*2];
float4 max0 = FETCH(pAABB, index*2 + 1); // pAABB[index*2+1];
uint2 sortedData = pParticleHash[index];
uint unsorted_indx = sortedData.y;
uint2 start_curr = pPairBuffStartCurr[unsorted_indx];
// uint2 start_curr = pPairBuffStartCurr[index];
uint start = start_curr.x;
uint curr = start_curr.y;
uint bucketEnd = bucketStart + params.maxParticlesPerCell;
bucketEnd = (bucketEnd > numParticles) ? numParticles : bucketEnd;
for(uint index2=bucketStart; index2 < bucketEnd; index2++)
{
uint2 cellData = pParticleHash[index2];
if (cellData.x != gridHash) break; // no longer in same bucket
if (index2 != index) // check not colliding with self
{
uint unsorted_indx2 = cellData.y;
if(unsorted_indx2 < unsorted_indx)
{
continue;
}
// float4 min1 = pAABB[index2*2];
// float4 max1 = pAABB[index2*2 + 1];
float4 min1 = FETCH(pAABB, index2*2);// pAABB[index2*2];
float4 max1 = FETCH(pAABB, index2*2 + 1); // pAABB[index2*2 + 1];
if(cudaTestAABBOverlap(min0, max0, min1, max1))
{
uint k;
for(k = 0; k < curr; k++)
{
uint old_pair = pPairBuff[start+k] & (~BT_CUDA_PAIR_ANY_FLG);
if(old_pair == unsorted_indx2)
{
pPairBuff[start+k] |= BT_CUDA_PAIR_FOUND_FLG;
break;
}
}
if(k == curr)
{
pPairBuff[start+curr] = unsorted_indx2 | BT_CUDA_PAIR_NEW_FLG;
curr++;
}
}
}
}
pPairBuffStartCurr[unsorted_indx] = make_uint2(start, curr);
return;
}
__global__ void
btCudaFindOverlappingPairsD(
float4* pAABB,
uint2* pParticleHash,
uint* pCellStart,
uint* pPairBuff,
uint2* pPairBuffStartCurr,
uint numParticles
)
{
int index = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
float4 bbMin = pAABB[index*2];
float4 bbMax = pAABB[index*2+1];
float4 pos = (bbMin + bbMax) * 0.5f;
// get address in grid
int3 gridPos = calcGridPos(pos);
// examine only neighbouring cells
for(int z=-1; z<=1; z++) {
for(int y=-1; y<=1; y++) {
for(int x=-1; x<=1; x++) {
findPairsInCell(gridPos + make_int3(x, y, z), index, pParticleHash, pCellStart, pAABB, pPairBuff, pPairBuffStartCurr, numParticles);
}
}
}
}
__global__ void
btCudaComputePairCacheChangesD(uint* pPairBuff, uint2* pPairBuffStartCurr, uint* pPairScan)
{
int index = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
uint2 start_curr = pPairBuffStartCurr[index];
uint start = start_curr.x;
uint curr = start_curr.y;
uint *pInp = pPairBuff + start;
uint num_changes = 0;
for(uint k = 0; k < curr; k++, pInp++)
{
if(!((*pInp) & BT_CUDA_PAIR_FOUND_FLG))
{
num_changes++;
}
}
pPairScan[index+1] = num_changes;
}
__global__ void
btCudaSqueezeOverlappingPairBuffD(uint* pPairBuff, uint2* pPairBuffStartCurr, uint* pPairScan, uint* pPairOut)
{
int index = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
uint2 start_curr = pPairBuffStartCurr[index];
uint start = start_curr.x;
uint curr = start_curr.y;
uint* pInp = pPairBuff + start;
uint* pOut = pPairOut + pPairScan[index];
uint* pOut2 = pInp;
uint num = 0;
for(uint k = 0; k < curr; k++, pInp++)
{
if(!((*pInp) & BT_CUDA_PAIR_FOUND_FLG))
{
*pOut = *pInp;
pOut++;
}
if((*pInp) & BT_CUDA_PAIR_ANY_FLG)
{
*pOut2 = (*pInp) & (~BT_CUDA_PAIR_ANY_FLG);
pOut2++;
num++;
}
}
pPairBuffStartCurr[index] = make_uint2(start, num);
} // btCudaBroadphase::squeezeOverlappingPairBuffCPU_D()
#endif #endif

7
Extras/CUDA/particles_kernel.cuh
View File

@@ -14,12 +14,6 @@
#define FETCH(t, i) t[i] #define FETCH(t, i) t[i]
#endif #endif
#define BT_CUDA_PAIR_FOUND_FLG (0x40000000)
#define BT_CUDA_PAIR_NEW_FLG (0x20000000)
#define BT_CUDA_PAIR_ANY_FLG (BT_CUDA_PAIR_FOUND_FLG | BT_CUDA_PAIR_NEW_FLG)
#include "vector_types.h" #include "vector_types.h"
typedef unsigned int uint; typedef unsigned int uint;
@@ -47,7 +41,6 @@ struct SimParams {
float shear; float shear;
float attraction; float attraction;
float boundaryDamping; float boundaryDamping;
}; };
#endif #endif

9
Extras/CUDA/render_particles.cpp
View File

@@ -33,12 +33,13 @@
#include <assert.h> #include <assert.h>
#include <stdio.h> #include <stdio.h>
#include <paramgl.h>
#include "BMF_Api.h" #include "BMF_Api.h"
#include "LinearMath/btQuickprof.h"
#include "render_particles.h" #include "render_particles.h"
#include "shaders.h" #include "shaders.h"
#include "LinearMath/btQuickprof.h"
#include "paramgl.h"
#ifndef M_PI #ifndef M_PI
#define M_PI 3.1415926535897932384626433832795 #define M_PI 3.1415926535897932384626433832795
@@ -186,7 +187,7 @@ void ParticleRenderer::_initGL()
#endif #endif
} }
#if 1
void ParticleRenderer::showProfileInfo(float& xOffset,float& yStart, float yIncr) void ParticleRenderer::showProfileInfo(float& xOffset,float& yStart, float yIncr)
{ {
#ifndef BT_NO_PROFILE #ifndef BT_NO_PROFILE
@@ -257,5 +258,3 @@ void ParticleRenderer::displayProfileString(int xOffset,int yStart,char* message
glRasterPos3f(xOffset,yStart,0); glRasterPos3f(xOffset,yStart,0);
BMF_DrawString(BMF_GetFont(BMF_kHelvetica10),message); BMF_DrawString(BMF_GetFont(BMF_kHelvetica10),message);
} }
#endif

4
Extras/CUDA/render_particles.h
View File

@@ -30,9 +30,6 @@
#ifndef __RENDER_PARTICLES__ #ifndef __RENDER_PARTICLES__
#define __RENDER_PARTICLES__ #define __RENDER_PARTICLES__
class CProfileIterator;
class ParticleRenderer class ParticleRenderer
{ {
public: public:
@@ -62,7 +59,6 @@ public:
void displayProfileString(int xOffset,int yStart,char* message); void displayProfileString(int xOffset,int yStart,char* message);
class CProfileIterator* m_profileIterator; class CProfileIterator* m_profileIterator;
protected: // methods protected: // methods
void _initGL(); void _initGL();
void _drawPoints(); void _drawPoints();