void neighborhood_cell::detection() { VEC3F_PTR pos = md->particleSystem()->position(); VEC3I cell3d; for (unsigned int i = 0; i < md->numParticle(); i++){ cell3d = getCellNumber(pos[i].x, pos[i].y, pos[i].z); cell_id[i] = getHash(cell3d); body_id[i] = i; } thrust::sort_by_key(cell_id, cell_id + md->numParticle(), body_id); memset(cell_start, 0xffffffff, sizeof(unsigned int) * ng); memset(cell_end, 0, sizeof(unsigned int)*ng); unsigned int begin = 0, end = 0, id = 0; bool ispass; while (end++ != md->numParticle()){ ispass = true; id = cell_id[begin]; if (id != cell_id[end]){ end - begin > 1 ? ispass = false : reorderDataAndFindCellStart(id, begin++, end); } if (!ispass){ reorderDataAndFindCellStart(id, begin, end); begin = end; } } }
void Particles::Update() { assert(initFlag); float *dPos; dPos = (float *)mapGLBufferObject(&m_cuda_posvbo_resource); setParameters(&mparams); integrateSystem(dPos, velGpu, mparams.timeStep, numParticles); calcHash(gridParticleHash, gridParticleIndex,dPos,numParticles); sortParticles(gridParticleHash,gridParticleIndex,numParticles); reorderDataAndFindCellStart(cellStart,cellEnd,sortedPos,sortedVel,gridParticleHash,gridParticleIndex,dPos,velGpu,numParticles,mparams.wholeNumCells); simFluid(velGpu, sortedPos, sortedVel, gridParticleIndex, cellStart, cellEnd, numParticles, solidPosGpu, solidVelGpu, buoyancyGpu, buoyancyAngGpu); unmapGLBufferObject(m_cuda_posvbo_resource); }
// step the simulation void ParticleSystem::update(float deltaTime) { assert(m_bInitialized); float *dPos; if (m_bUseOpenGL) { dPos = (float *) mapGLBufferObject(&m_cuda_posvbo_resource); } else { dPos = (float *) m_cudaPosVBO; } // update constants setParameters(&m_params); // integrate integrateSystem( dPos, m_dVel, deltaTime, m_numParticles); // calculate grid hash calcHash( m_dGridParticleHash, m_dGridParticleIndex, dPos, m_numParticles); // sort particles based on hash sortParticles(m_dGridParticleHash, m_dGridParticleIndex, m_numParticles); // reorder particle arrays into sorted order and // find start and end of each cell reorderDataAndFindCellStart( m_dCellStart, m_dCellEnd, m_dSortedPos, m_dSortedVel, m_dGridParticleHash, m_dGridParticleIndex, dPos, m_dVel, m_numParticles, m_numGridCells); // process collisions collide( m_dVel, m_dSortedPos, m_dSortedVel, m_dGridParticleIndex, m_dCellStart, m_dCellEnd, m_numParticles, m_numGridCells); // note: do unmap at end here to avoid unnecessary graphics/CUDA context switch if (m_bUseOpenGL) { unmapGLBufferObject(m_cuda_posvbo_resource); } }
void FluidSystem::update(float deltaTime) //整个particle或流体的update { assert(m_bInitialized); float *dPos; if (m_bUseOpenGL) { m_dPos = (float *)mapGLBufferObject(&m_cuda_posvbo_resource); m_dVel = (float *)mapGLBufferObject(&m_cuda_velvbo_resource); } else { m_dPos = (float *)m_cudaPosVBO; m_dVel = (float *)m_cudaVelVBO; } /////////////////////////////////////////////////////calculate time cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); // Start record cudaEventRecord(start, 0); /////////////////////////////////////////////////////////////////////// setParameters(&m_params); integrateSystem( m_dPos, m_dVel, m_dDen, m_dPre, deltaTime, m_numParticles, buoyancyForce); //calculateBuoyancy(); cudaEvent_t start1, stop1; cudaEventCreate(&start1); cudaEventCreate(&stop1); // Start record cudaEventRecord(start1, 0); calcHash( m_dGridParticleHash, m_dGridParticleIndex, m_dPos, m_numParticles); sortParticles(m_dGridParticleHash, m_dGridParticleIndex, m_numParticles); reorderDataAndFindCellStart( m_dCellStart, m_dCellEnd, m_dSortedPos, m_dSortedVel, m_dGridParticleHash, m_dGridParticleIndex, m_dPos, m_dVel, m_numParticles, m_numGridCells, m_dSortedDen, m_dSortedPre, m_dDen, //old density m_dPre,//old pressure m_dSortedColorf, m_dColorf); cudaEventRecord(stop1, 0); cudaEventSynchronize(stop1); float elapsedTime1; cudaEventElapsedTime(&elapsedTime1, start1, stop1); // that's our time! std::cout <<"neighbour " <<elapsedTime1 << endl; // Clean up: cudaEventDestroy(start1); cudaEventDestroy(stop1); /*colorfield( m_dSortedPos, m_dGridParticleIndex, m_dCellStart, m_dCellEnd, m_numParticles, m_numGridCells, m_dColorf, m_dSortedColorf);*/ // copyArrayFromDevice(m_hDen, m_dDen, 0, sizeof(float) * m_numParticles); /*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 ( m_hDen[i]<2.5f){ colorRamp(1, 1, 1, ptr); ptr += 3; *ptr++ = 0.1f; } else { colorRamp(0, 1, 1, ptr); ptr += 3; *ptr++ = 0.1f; } } glUnmapBufferARB(GL_ARRAY_BUFFER);*/ cudaEvent_t start2, stop2; cudaEventCreate(&start2); cudaEventCreate(&stop2); // Start record cudaEventRecord(start2, 0); calculateDensity(m_dSortedPos, m_dGridParticleIndex, m_dCellStart, m_dCellEnd, m_numParticles, m_dDen); collide( m_dVel, //new velocity m_dSortedPos, m_dSortedVel, m_dGridParticleIndex, m_dCellStart, m_dCellEnd, m_numParticles, m_numGridCells, m_dDen, //new density m_dPre, //new pressure m_dSortedDen, m_dSortedPre, m_dPos, buoyancyForce); cudaEventRecord(stop2, 0); cudaEventSynchronize(stop2); float elapsedTime2; cudaEventElapsedTime(&elapsedTime2, start2, stop2); // that's our time! std::cout << "renew " << elapsedTime2 << endl; // Clean up: cudaEventDestroy(start2); cudaEventDestroy(stop2); ////////////////////////////////////////////////////////////stop calculate time cudaEventRecord(stop, 0); cudaEventSynchronize(stop); float elapsedTime; cudaEventElapsedTime(&elapsedTime, start, stop); // that's our time! std::cout << "whole " << elapsedTime << endl; // Clean up: cudaEventDestroy(start); cudaEventDestroy(stop); ////////////////////////////////////////////////////////// if (m_bUseOpenGL) { unmapGLBufferObject(m_cuda_posvbo_resource); unmapGLBufferObject(m_cuda_velvbo_resource); } }