void PoiseuilleFlowSystem::setArray(ParticleArray array, const float* data, int start, int count){
assert(IsInitialized);
switch (array)
{
default:
case POSITION:
{
if (IsOpenGL) {
unregisterGLBufferObject(cuda_posvbo_resource);
glBindBuffer(GL_ARRAY_BUFFER, posVbo);
glBufferSubData(GL_ARRAY_BUFFER, start*4*sizeof(float), count*4*sizeof(float), data);
glBindBuffer(GL_ARRAY_BUFFER, 0);
registerGLBufferObject(posVbo, &cuda_posvbo_resource);
}else
{
copyArrayToDevice(cudaPosVBO, data, start*4*sizeof(float), count*4*sizeof(float));
}
}
break;
case VELOCITY:
copyArrayToDevice(dVel, data, start*4*sizeof(float), count*4*sizeof(float));
break;
case MEASURES:
copyArrayToDevice(dMeasures, data, start*4*sizeof(float), count*4*sizeof(float));
break;
case ACCELERATION:
copyArrayToDevice(dAcceleration, data, start*4*sizeof(float), count*4*sizeof(float));
break;
case VELOCITYLEAPFROG:
copyArrayToDevice(dVelLeapFrog, data, start*4*sizeof(float), count*4*sizeof(float));
break;
}
}
void
ParticleSystem::setArray(ParticleArray array, const float *data, int start, int count)
{
assert(m_bInitialized);
switch (array)
{
default:
case POSITION:
{
if (m_bUseOpenGL)
{
unregisterGLBufferObject(m_cuda_posvbo_resource);
glBindBuffer(GL_ARRAY_BUFFER, m_posVbo);
glBufferSubData(GL_ARRAY_BUFFER, start*4*sizeof(float), count*4*sizeof(float), data);
glBindBuffer(GL_ARRAY_BUFFER, 0);
registerGLBufferObject(m_posVbo, &m_cuda_posvbo_resource);
}
}
break;
case VELOCITY:
copyArrayToDevice(m_dVel, data, start*4*sizeof(float), count*4*sizeof(float));
break;
}
}
void Particles::InitSolid()
{
//init solidPos
solidPos = new float[4];
solidPos[0] = 20.0f;
solidPos[1] = 10.0f;
solidPos[2] = 150.0f;
solidPos[3] = 0.0f;
cudaMalloc((void **)&solidPosGpu, 4 * sizeof(float));
copyArrayToDevice(solidPosGpu, solidPos, 0, 4 * sizeof(float));
//init solid vel
solidVel = new float[4];
memset(solidVel, 0, sizeof(float)* 4);
allocateArray((void **)&solidVelGpu, sizeof(float)* 4);
cudaMemset(solidVelGpu, 0, sizeof(float)* 4);
//buoyancy
buoyancy = new float[4 * numParticles];
memset(buoyancy, 0, sizeof(float)* 4 * numParticles);
allocateArray((void **)&buoyancyGpu, sizeof(float)* 4 * numParticles);
cudaMemset(buoyancyGpu, 0, sizeof(float)* 4 * numParticles);
//buoyancy angular velocity
buoyancyAng = new float[4 * numParticles];
memset(buoyancyAng, 0, sizeof(float)* 4 * numParticles);
allocateArray((void **)&buoyancyAngGpu, sizeof(float)* 4 * numParticles);
cudaMemset(buoyancyAngGpu, 0, sizeof(float)* 4 * numParticles);
}
void MulC_I(std::vector<T* >& h_Imgs, T c,
int n, int nImgs, T* d_scratchI[])
{
#if 0
assert(nImgs < h_Imgs.size());
// load the first image
copyArrayToDevice(d_scratchI[0], h_Imgs[0], n);
cplVectorOpers::MulC_I(d_scratchI[0], c, n);
// load the second image
if (nImgs > 1)
copyArrayToDevice(d_scratchI[1], h_Imgs[1], n);
int i=2;
for (; i < nImgs; ++i){
copyArrayToDeviceAsync(d_scratchI[i % 3], h_Imgs[i], n, STM_H2D);
cplVectorOpers::MulC_I(d_scratchI[(i-1) % 3], c, n, STM_D2D);
copyArrayFromDeviceAsync(h_Imgs[i-2], d_scratchI[(i-2)%3], n, STM_D2H);
cudaThreadSynchronize();
}
if (nImgs> 1)
cplVectorOpers::MulC_I(d_scratchI[(i-1) % 3], c, n, STM_D2D);
copyArrayFromDeviceAsync(h_Imgs[i-2], d_scratchI[(i-2) % 3], n, STM_D2H);
cudaThreadSynchronize();
++i;
if (nImgs> 1)
copyArrayFromDeviceAsync(h_Imgs[i-2], d_scratchI[(i-2) % 3], n, STM_D2H);
#else
for (int i=0; i < nImgs + 2; ++i) {
if (i < nImgs)
copyArrayToDeviceAsync(d_scratchI[i % 3], h_Imgs[i], n, STM_H2D);
if ((i >=1) && ((i-1) < nImgs))
cplVectorOpers::MulC_I(d_scratchI[(i-1) % 3], c, n, STM_D2D);
if ((i >=2) && ((i-2) < nImgs))
copyArrayFromDeviceAsync(h_Imgs[i-2], d_scratchI[(i-2)%3], n, STM_D2H);
cudaThreadSynchronize();
}
#endif
}
void FluidSystem::setArray(FluidArray array, const float *data, int start, int count) //重新设置particle的速度和位置数组
{
assert(m_bInitialized);
switch (array)
{
default:
case POSITION:
{
if (m_bUseOpenGL)
{
unregisterGLBufferObject(m_cuda_posvbo_resource);
glBindBuffer(GL_ARRAY_BUFFER, m_posVbo);
glBufferSubData(GL_ARRAY_BUFFER, start*4*sizeof(float), count*4*sizeof(float), data);
glBindBuffer(GL_ARRAY_BUFFER, 0);
registerGLBufferObject(m_posVbo, &m_cuda_posvbo_resource);
}
}
break;
case VELOCITY:
unregisterGLBufferObject(m_cuda_velvbo_resource);
glBindBuffer(GL_ARRAY_BUFFER, m_velVBO);
glBufferSubData(GL_ARRAY_BUFFER, start * 4 * sizeof(float), count * 4 * sizeof(float), data);
glBindBuffer(GL_ARRAY_BUFFER, 0);
registerGLBufferObject(m_velVBO, &m_cuda_velvbo_resource);
break;
case DENSITY:
copyArrayToDevice(m_dDen, data, start * sizeof(float), count * sizeof(float));
break;
case PRESSURE:
copyArrayToDevice(m_dPre, data, start * sizeof(float), count * sizeof(float));
break;
case COLORFIELD:
copyArrayToDevice(m_dColorf, data, start * sizeof(float), count * sizeof(float));
break;
}
}
void Min(T* d_o, std::vector<T* >& h_Imgs,
int n, int nImgs,
cplReduceS* rd, T* d_scratchI[])
{
assert(nImgs <= h_Imgs.size());
copyArrayToDevice(d_scratchI[0], h_Imgs[0], n);
int i=1;
for (; i< h_Imgs.size(); ++i){
copyArrayToDeviceAsync(d_scratchI[i&1], h_Imgs[i], n, STM_H2D);
(i==1) ? rd->Min(d_o, d_scratchI[(i-1)&1], n, STM_D2D)
: rd->MinA(d_o, d_scratchI[(i-1)&1], n, STM_D2D);
cudaThreadSynchronize();
}
rd->MinA(d_o, d_scratchI[(i-1)&1], n);
}
Matrix4 Particles::BuildTransform()
{
copyArrayFromDevice(buoyancy, buoyancyGpu, 0, sizeof(float)* 4 * numParticles);
copyArrayFromDevice(buoyancyAng, buoyancyAngGpu, 0, sizeof(float)* 4 * numParticles);
float* force = new float[4];
for (int i = 0; i < numParticles; i++)
{
force[0] += buoyancy[i * 4];
force[1] += buoyancy[i * 4 + 1];
force[2] += buoyancy[i * 4 + 2];
force[3] = 0;
Vector3 ang = Vector3(buoyancyAng[i * 4], buoyancyAng[i * 4 + 1], buoyancyAng[i * 4 + 2]);
if (ang.x != 0 || ang.y != 0 || ang.z != 0)
{
orientation = orientation + orientation * (ang / 20000000.0f);
orientation.Normalise();
}
}
force[1] -= 9.81f * 2000.0f;
for (int i = 0; i < 4; i++)
{
solidVel[i] = solidVel[i] + force[i] * mparams.timeStep;
}
CheckEdges(solidPos, solidVel);
copyArrayToDevice(solidPosGpu, solidPos, 0, 4 * sizeof(float));
copyArrayToDevice(solidVelGpu, solidVel, 0, 4 * sizeof(float));
//orientation.Normalise();
Matrix4 m = orientation.ToMatrix();
Vector3 p = Vector3(solidPos[0], solidPos[1], solidPos[2]);
m.SetPositionVector(p);
return m;
}
void CudaSystem::evaluateForcesExt()
{
for(unsigned int i=0;i<particles.size();i++){
m_hForce[(i*3)] = 0; m_hForce[(i*3)+1] = 0; m_hForce[(i*3)+2] = 0;
}
copyArrayToDevice(FExt->m_F, m_hForce, 0, sizeof(double)*3*particles.size());
for(unsigned int i=0;i<FExt->getNbForces();i++){
if(typeid(*(FExt->getForce(i))) == typeid(ForceExt_Constante)){
ForceExt_Constante* FC = (ForceExt_Constante*) FExt->getForce(i);
evaluate_ForceExt_Constante(particles.size(), FExt->m_F, m_dMass, FC);
}
if(typeid(*(FExt->getForce(i))) == typeid(ForceExt_Trochoide)){
ForceExt_Trochoide *T = (ForceExt_Trochoide*) FExt->getForce(i);
evaluate_ForceExt_Trochoide(particles.size(), m_hPos[1], FExt->m_F, m_dMass, T);
T->setTime(T->getTime()+1);
}
}
interactionSystem(m_dPos[0], m_dVel[0], FExt->m_F, m_dInteractionRadius, m_dSpring, m_dDamping,
m_dShear, m_dAttraction, voisines, particles.size());
}
void Particles::SetArray(ParticleArray array, const float *data, int start, int count)
{
assert(initFlag);
switch (array)
{
default:
case POSITION:
{
unregisterGLBufferObject(m_cuda_posvbo_resource);
glBindBufferARB(GL_ARRAY_BUFFER_ARB, posVbo);
glBufferSubDataARB(GL_ARRAY_BUFFER_ARB, start * 4 * sizeof(float), count * 4 * sizeof(float), data);
glBindBufferARB(GL_ARRAY_BUFFER_ARB, 0);
registerGLBufferObject(posVbo, &m_cuda_posvbo_resource);
}
break;
case VELOCITY:
copyArrayToDevice(velGpu, data, start * 4 * sizeof(float), count * 4 * sizeof(float));
break;
}
}
btGridBroadphaseCl::btGridBroadphaseCl( btOverlappingPairCache* overlappingPairCache,
const btVector3& cellSize,
int gridSizeX, int gridSizeY, int gridSizeZ,
int maxSmallProxies, int maxLargeProxies, int maxPairsPerSmallProxy,
btScalar maxSmallProxySize,
int maxSmallProxiesPerCell,
cl_context context,
cl_device_id device,
cl_command_queue queue,
adl::DeviceCL* deviceCL)
:bt3dGridBroadphaseOCL(overlappingPairCache,cellSize,
gridSizeX, gridSizeY, gridSizeZ,
maxSmallProxies, maxLargeProxies, maxPairsPerSmallProxy,
maxSmallProxySize,maxSmallProxiesPerCell,
context,device,queue,deviceCL)
{
m_computeAabbKernel = m_deviceCL->getKernel(COMPUTE_AABB_KERNEL_PATH,"computeAabb","",spComputeAabbSource);
m_countOverlappingPairs = m_deviceCL->getKernel(COMPUTE_AABB_KERNEL_PATH,"countOverlappingpairs","",spComputeAabbSource);
m_squeezePairCaches = m_deviceCL->getKernel(COMPUTE_AABB_KERNEL_PATH,"squeezePairCaches","",spComputeAabbSource);
m_aabbConstBuffer = new adl::Buffer<MyAabbConstData >(m_deviceCL,1,adl::BufferBase::BUFFER_CONST);
size_t memSize = m_maxHandles * m_maxPairsPerBody * sizeof(unsigned int)*2;
cl_int ciErrNum=0;
m_dAllOverlappingPairs = clCreateBuffer(m_cxMainContext, CL_MEM_READ_WRITE, memSize, NULL, &ciErrNum);
memset(m_hAllOverlappingPairs, 0x00, sizeof(MyUint2)*m_maxHandles * m_maxPairsPerBody);
copyArrayToDevice(m_dAllOverlappingPairs, m_hAllOverlappingPairs, m_maxHandles * m_maxPairsPerBody * sizeof(MyUint2));
oclCHECKERROR(ciErrNum, CL_SUCCESS);
}
void btGridBroadphaseCl::calculateOverlappingPairs(float* positions, int numObjects)
{
btDispatcher* dispatcher=0;
// update constants
{
BT_PROFILE("setParameters");
setParameters(&m_params);
}
// prepare AABB array
{
BT_PROFILE("prepareAABB");
prepareAABB(positions, numObjects);
}
// calculate hash
{
BT_PROFILE("calcHashAABB");
calcHashAABB();
}
{
BT_PROFILE("sortHash");
// sort bodies based on hash
sortHash();
}
// find start of each cell
{
BT_PROFILE("findCellStart");
findCellStart();
}
{
BT_PROFILE("findOverlappingPairs");
// findOverlappingPairs (small/small)
findOverlappingPairs();
}
// add pairs to CPU cache
{
BT_PROFILE("computePairCacheChanges");
#if 0
computePairCacheChanges();
#else
int ciErrNum=0;
ciErrNum=clSetKernelArg((cl_kernel)m_countOverlappingPairs->m_kernel, 0, sizeof(int), (void*)&numObjects);
ciErrNum=clSetKernelArg((cl_kernel)m_countOverlappingPairs->m_kernel, 1, sizeof(cl_mem),(void*)&m_dPairBuff);
ciErrNum=clSetKernelArg((cl_kernel)m_countOverlappingPairs->m_kernel, 2, sizeof(cl_mem),(void*)&m_dPairBuffStartCurr);
ciErrNum=clSetKernelArg((cl_kernel)m_countOverlappingPairs->m_kernel, 3, sizeof(cl_mem),(void*)&m_dPairScanChanged);
ciErrNum=clSetKernelArg((cl_kernel)m_countOverlappingPairs->m_kernel, 4, sizeof(cl_mem),(void*)&m_dAABB);
size_t localWorkSize=64;
size_t numWorkItems = localWorkSize*((numObjects+ (localWorkSize)) / localWorkSize);
ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue, (cl_kernel)m_countOverlappingPairs->m_kernel, 1, NULL, &numWorkItems, &localWorkSize, 0,0,0 );
oclCHECKERROR(ciErrNum, CL_SUCCESS);
ciErrNum = clFlush(m_cqCommandQue);
#endif
}
{
BT_PROFILE("scanOverlappingPairBuff");
scanOverlappingPairBuff(false);
}
{
BT_PROFILE("squeezeOverlappingPairBuff");
//#define FORCE_CPU
#ifdef FORCE_CPU
bt3dGridBroadphaseOCL::squeezeOverlappingPairBuff();
copyArrayToDevice(m_dPairsChangedXY, m_hPairsChangedXY, sizeof( MyUint2) * m_numPrefixSum); //gSum
#else
//squeezeOverlappingPairBuff();
int ciErrNum = 0;
ciErrNum=clSetKernelArg((cl_kernel)m_squeezePairCaches->m_kernel, 0, sizeof(int), (void*)&numObjects);
ciErrNum=clSetKernelArg((cl_kernel)m_squeezePairCaches->m_kernel, 1, sizeof(cl_mem),(void*)&m_dPairBuff);
ciErrNum=clSetKernelArg((cl_kernel)m_squeezePairCaches->m_kernel, 2, sizeof(cl_mem),(void*)&m_dPairBuffStartCurr);
ciErrNum=clSetKernelArg((cl_kernel)m_squeezePairCaches->m_kernel, 3, sizeof(cl_mem),(void*)&m_dPairScanChanged);
ciErrNum=clSetKernelArg((cl_kernel)m_squeezePairCaches->m_kernel, 4, sizeof(cl_mem),(void*)&m_dAllOverlappingPairs);
ciErrNum=clSetKernelArg((cl_kernel)m_squeezePairCaches->m_kernel, 5, sizeof(cl_mem),(void*)&m_dAABB);
size_t workGroupSize = 64;
size_t numWorkItems = workGroupSize*((numObjects+ (workGroupSize)) / workGroupSize);
ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue, (cl_kernel)m_squeezePairCaches->m_kernel, 1, NULL, &numWorkItems, &workGroupSize, 0,0,0 );
oclCHECKERROR(ciErrNum, CL_SUCCESS);
// copyArrayFromDevice(m_hAllOverlappingPairs, m_dAllOverlappingPairs, sizeof(unsigned int) * m_numPrefixSum*2); //gSum
// clFinish(m_cqCommandQue);
#endif
}
return;
}
//Step the simulation
void ParticleSystem::update(float deltaTime){
assert(m_bInitialized);
setParameters(&m_params);
setParametersHost(&m_params);
//Download positions from VBO
memHandle_t pos;
if (!m_bQATest)
{
glBindBufferARB(GL_ARRAY_BUFFER, m_posVbo);
pos = (memHandle_t)glMapBufferARB(GL_ARRAY_BUFFER, GL_READ_WRITE);
copyArrayToDevice(m_dPos, pos, 0, m_numParticles * 4 * sizeof(float));
}
integrateSystem(
m_dPos,
m_dVel,
deltaTime,
m_numParticles
);
calcHash(
m_dHash,
m_dIndex,
m_dPos,
m_numParticles
);
bitonicSort(NULL, m_dHash, m_dIndex, m_dHash, m_dIndex, 1, m_numParticles, 0);
//Find start and end of each cell and
//Reorder particle data for better cache coherency
findCellBoundsAndReorder(
m_dCellStart,
m_dCellEnd,
m_dReorderedPos,
m_dReorderedVel,
m_dHash,
m_dIndex,
m_dPos,
m_dVel,
m_numParticles,
m_numGridCells
);
collide(
m_dVel,
m_dReorderedPos,
m_dReorderedVel,
m_dIndex,
m_dCellStart,
m_dCellEnd,
m_numParticles,
m_numGridCells
);
//Update buffers
if (!m_bQATest)
{
copyArrayFromDevice(pos,m_dPos, 0, m_numParticles * 4 * sizeof(float));
glUnmapBufferARB(GL_ARRAY_BUFFER);
}
}
void ParticleMesh::calculate(LevelsetCollider &LC, int size, float r, float mass)
{
float scale = 2 * r / (LC.gridStep*(LC.gridSize.x - 4) / size);
vec3<double> center = vec3<double>(0, 0, 0);
vec3<double> MaxB = LC.MaxBoundary;
vec3<double> MinB = LC.MinBoundary;
scale = CalScale(MaxB, MinB, r, size);
int index_par = 0;
par_list = new triple[20000];
sdf_list = new SDF[20000];
float step = r * 2 / scale;
for (int i = 0; i < size; i++)
{
for (int j = 0; j < size; j++)
{
for (int k = 0; k <size; k++)
{
vec3<int> gridPos = vec3<int>(i, j, k);
vec3<double> pos = vec3<double>(i + 0.5, j + 0.5, k + 0.5) * (2 * r / scale) + MinB;
int gridInd = dot(gridPos, LC.gridSizeOffset);
double sdf = 0;
LC.getGridDist(pos, sdf);
if (sdf < 0 && sdf > -r * 2 / scale * 2)
{
vec3<double> pos_par = pos - MinB;// -vec3<double>(0.5, 0.5, 0.5) * (2 * r / scale);
center += mass*pos_par * scale;
par_list[index_par].x[0] = pos_par.x * scale;
par_list[index_par].x[1] = pos_par.y * scale;
par_list[index_par].x[2] = pos_par.z * scale;
sdf_list[index_par].sdf = sdf * scale;
vec3<double> gradient;
LC.checkCollision(pos, sdf, gradient, -1);
sdf_list[index_par].gradient[0] = gradient.x;
sdf_list[index_par].gradient[1] = gradient.y;
sdf_list[index_par].gradient[2] = gradient.z;
index_par++;
}
}
}
}
center /= index_par*mass;
float3 center_f3 = make_float3(center[0], center[1], center[2]);
m_totalParticles = index_par;
m_totalConnectedPoints = LC.meshModel->m_totalConnectedPoints;
m_totalFaces = LC.meshModel->m_totalFaces;
float3 vertex,relative;
mp_vertexXYZ = new float[m_totalConnectedPoints * 3];
Relative_Vertex = new float3[m_totalConnectedPoints];
mp_vertexNorm = new float[m_totalConnectedPoints * 3];
mp_vertexRGB = new float[m_totalConnectedPoints * 3];
for (int i = 0; i < m_totalConnectedPoints; i++)
{
vertex.x = LC.meshModel->mp_vertexXYZ[i * 3];
vertex.y = LC.meshModel->mp_vertexXYZ[i * 3 + 1];
vertex.z = LC.meshModel->mp_vertexXYZ[i * 3 + 2];
//vertex *= scale;
vertex -= make_float3(MinB[0], MinB[1], MinB[2]);// +make_float3(0.5, 0.5, 0.5) * (2 * r / scale);
vertex *= scale;
relative = vertex - center_f3;
*(float3*)(mp_vertexXYZ + i * 3) = vertex;
mp_vertexRGB[i * 3] = LC.meshModel->mp_vertexRGB[i * 3];
mp_vertexRGB[i * 3 + 1] = LC.meshModel->mp_vertexRGB[i * 3 + 1];
mp_vertexRGB[i * 3 + 2] = LC.meshModel->mp_vertexRGB[i * 3 + 2];
mp_vertexNorm[i * 3] = 0; //LC.meshModel->mp_vertexNorm[i * 3];
mp_vertexNorm[i * 3 + 1] = 0; //LC.meshModel->mp_vertexNorm[i * 3 + 1];
mp_vertexNorm[i * 3 + 2] = 0;// LC.meshModel->mp_vertexNorm[i * 3 + 2];
Relative_Vertex[i] = relative;
}
FaceVerInd = new uint[3 * m_totalFaces];
//memcpy(FaceVerInd, LC.meshModel->mp_face, 9 * m_totalFaces * sizeof(uint));
float3 n;
for (int i = 0; i < m_totalFaces; i++)
{
uint i_1 = FaceVerInd[3 * i] = LC.meshModel->mp_face[3 * i];
uint i_2 = FaceVerInd[3 * i + 1] = LC.meshModel->mp_face[3 * i + 1];
uint i_3 = FaceVerInd[3 * i + 2] = LC.meshModel->mp_face[3 * i + 2];
float3 v1 = make_float3(mp_vertexXYZ[i_1 * 3], mp_vertexXYZ[i_1 * 3 + 1], mp_vertexXYZ[i_1 * 3 + 2]);
float3 v2 = make_float3(mp_vertexXYZ[i_2 * 3], mp_vertexXYZ[i_2 * 3 + 1], mp_vertexXYZ[i_2 * 3 + 2]);
float3 v3 = make_float3(mp_vertexXYZ[i_3 * 3], mp_vertexXYZ[i_3 * 3 + 1], mp_vertexXYZ[i_3 * 3 + 2]);
float3 e12 = v2 - v1;
float3 e13 = v3 - v1;
n = cross(e12, e13);
n /= length(n);
mp_vertexNorm[i_1 * 3] += n.x;
mp_vertexNorm[i_1 * 3 + 1] += n.y;
mp_vertexNorm[i_1 * 3 + 2] += n.z;
//atomicAdd(&num_tri_per_point[i_1], 1);
float3 e21 = -e12;
float3 e23 = v3 - v2;
n = cross(e23, e21);
n /= length(n);
mp_vertexNorm[i_2 * 3] += n.x;
mp_vertexNorm[i_2 * 3 + 1] += n.y;
mp_vertexNorm[i_2 * 3 + 2] += n.z;
//atomicAdd(&num_tri_per_point[i_2], 1);
float3 e31 = -e13;
float3 e32 = -e23;
n = cross(e31, e32);
n /= length(n);
mp_vertexNorm[i_3 * 3] += n.x;
mp_vertexNorm[i_3 * 3 + 1] += n.y;
mp_vertexNorm[i_3 * 3 + 2] += n.z;
}
for (int i = 0; i < m_totalConnectedPoints; i++)
{
float3 n = make_float3(mp_vertexNorm[i * 3], mp_vertexNorm[i * 3 + 1], mp_vertexNorm[i * 3 + 2]);
mp_vertexNorm[i * 3] /= length(n);
mp_vertexNorm[i * 3 + 1] /= length(n);
mp_vertexNorm[i * 3 + 2] /= length(n);
}
m_vertexVbo = createVBO(m_totalConnectedPoints * 3 * sizeof(float));
registerGLBufferObject(m_vertexVbo, &m_cuda_vertexvbo_resource);
m_indexVBO = createVBO(3 * m_totalFaces * sizeof(uint));
registerGLBufferObject(m_indexVBO, &m_cuda_indexvbo_resource);
m_colorVBO = createVBO(m_totalConnectedPoints * 3 * sizeof(float));
registerGLBufferObject(m_colorVBO, &m_cuda_colorvbo_resource);
m_NorVBO = createVBO(m_totalConnectedPoints * 3 * sizeof(float));
registerGLBufferObject(m_NorVBO, &m_cuda_normalvbo_resource);
unregisterGLBufferObject(m_cuda_vertexvbo_resource);
glBindBuffer(GL_ARRAY_BUFFER, m_vertexVbo);
glBufferSubData(GL_ARRAY_BUFFER, 0, m_totalConnectedPoints * 3 * sizeof(float), mp_vertexXYZ);
glBindBuffer(GL_ARRAY_BUFFER, 0);
registerGLBufferObject(m_vertexVbo, &m_cuda_vertexvbo_resource);
unregisterGLBufferObject(m_cuda_indexvbo_resource);
glBindBuffer(GL_ARRAY_BUFFER, m_indexVBO);
glBufferSubData(GL_ARRAY_BUFFER, 0, 3 * m_totalFaces * sizeof(uint), FaceVerInd);
glBindBuffer(GL_ARRAY_BUFFER, 0);
registerGLBufferObject(m_indexVBO, &m_cuda_indexvbo_resource);
unregisterGLBufferObject(m_cuda_colorvbo_resource);
glBindBuffer(GL_ARRAY_BUFFER, m_colorVBO);
glBufferSubData(GL_ARRAY_BUFFER, 0, m_totalConnectedPoints * 3 * sizeof(float), mp_vertexRGB);
glBindBuffer(GL_ARRAY_BUFFER, 0);
registerGLBufferObject(m_colorVBO, &m_cuda_colorvbo_resource);
unregisterGLBufferObject(m_cuda_normalvbo_resource);
glBindBuffer(GL_ARRAY_BUFFER, m_NorVBO);
glBufferSubData(GL_ARRAY_BUFFER, 0, m_totalConnectedPoints * 3 * sizeof(float), mp_vertexNorm);
glBindBuffer(GL_ARRAY_BUFFER, 0);
registerGLBufferObject(m_NorVBO, &m_cuda_normalvbo_resource);
allocateArray((void**)&dRelative_Vertex, m_totalConnectedPoints * 3 * sizeof(float));
copyArrayToDevice(dRelative_Vertex, Relative_Vertex, 0, m_totalConnectedPoints * 3 * sizeof(float));
hasMesh = true;
}
void CudaSystem::_initialize(int begin, int numParticles)
{
if(numParticles>0){
if(begin>0){
copyArrayFromDevice(m_hPos[0], m_dPos[0], 0, sizeof(double)*3*begin);
copyArrayFromDevice(m_hVel[0], m_dVel[0], 0, sizeof(double)*3*begin);
copyArrayFromDevice(m_hPos[1], m_dPos[1], 0, sizeof(double)*3*begin);
copyArrayFromDevice(m_hVel[1], m_dVel[1], 0, sizeof(double)*3*begin);
}
printf("begin:%d numParticles:%d\n",begin, numParticles);
for(int i=begin;i<numParticles;i++){
CudaParticle* p = (CudaParticle*) particles[i];
m_hPos[0][i*3] = p->getNewPos().x();
m_hPos[0][(i*3)+1] = p->getNewPos().y();
m_hPos[0][(i*3)+2] = p->getNewPos().z();
m_hVel[0][i*3] = p->getNewVel().x();
m_hVel[0][(i*3)+1] = p->getNewVel().y();
m_hVel[0][(i*3)+2] = p->getNewVel().z();
m_hPos[1][(i*3)] = p->getNewPos().x();
m_hPos[1][(i*3)+1] = p->getNewPos().y();
m_hPos[1][(i*3)+2] = p->getNewPos().z();
m_hVel[1][i*3] = p->getNewVel().x();
m_hVel[1][(i*3)+1] = p->getNewVel().y();
m_hVel[1][(i*3)+2] = p->getNewVel().z();
m_hMass[i] = p->getMass();
m_hParticleRadius[i] = p->getParticleRadius();
m_hInteractionRadius[i] = p->getInteractionRadius();
m_hSpring[i] = p->getSpring();
m_hDamping[i] = p->getDamping();
m_hShear[i] = p->getShear();
m_hAttraction[i] = p->getAttraction();
m_hColors[i*4] = p->getColor().x();
m_hColors[(i*4)+1] = p->getColor().y();
m_hColors[(i*4)+2] = p->getColor().z();
m_hColors[(i*4)+3] = 1.0;
}
copyArrayToDevice(m_dPos[0], m_hPos[0], 0, sizeof(double)*3*numParticles);
copyArrayToDevice(m_dVel[0], m_hVel[0], 0, sizeof(double)*3*numParticles);
copyArrayToDevice(m_dPos[1], m_hPos[1], 0, sizeof(double)*3*numParticles);
copyArrayToDevice(m_dVel[1], m_hVel[1], 0, sizeof(double)*3*numParticles);
copyArrayToDevice(m_dMass, m_hMass, 0, sizeof(double)*numParticles);
copyArrayToDevice(m_dParticleRadius, m_hParticleRadius, 0, sizeof(double)*numParticles);
copyArrayToDevice(m_dInteractionRadius, m_hInteractionRadius, 0, sizeof(double)*numParticles);
copyArrayToDevice(m_dSpring, m_hSpring, 0, sizeof(double)*numParticles);
copyArrayToDevice(m_dDamping, m_hDamping, 0, sizeof(double)*numParticles);
copyArrayToDevice(m_dShear, m_hShear, 0, sizeof(double)*numParticles);
copyArrayToDevice(m_dAttraction, m_hAttraction, 0, sizeof(double)*numParticles);
if(gridCreated==false){
CudaParticle* p = (CudaParticle*) particles[0];
grid = (UniformGrid*)malloc(sizeof(UniformGrid));
createGrid(0,0,0,1.5,1.5,1.5,p->getInteractionRadius(),grid);
gridCreated = true;
}
}
}
