Kernel launches

void ParticleSystem::runParticleSimulation(float dt)
{
    //Switch around buffers
    dVelocityCurrent = dVelocityOld;
    dVelocityOld = dVelocityNew;
    dVelocityNew = dVelocityCurrent;


    //Set cells to 0
    checkCudaError(cudaMemset((void*) cellCountArray, 0, sizeof(size_t) * params.gridCellCount));
    checkCudaError(cudaDeviceSynchronize());

    //Count particles in cells
    cellCount<<<(numParticles + 255) / 256, 256>>>(numParticles, dPosition, cellCountArray);
    checkCudaError(cudaDeviceSynchronize());

    //Calculate prefix sum
    thrust::device_ptr<size_t> ptr = thrust::device_pointer_cast(cellCountArray);
    thrust::exclusive_scan(ptr, ptr + params.gridCellCount, ptr);
    checkCudaError(cudaDeviceSynchronize());

    //Populate cellGrid
    fillCellGrid<<<(numParticles + 255) / 256, 256>>>(numParticles, dPosition, cellCountArray, cellGrid);
    checkCudaError(cudaDeviceSynchronize());

    //Calculate densities
    calculateDensities<<<(numParticles + 255) / 256, 256>>>(numParticles, dPosition, cellGrid, cellCountArray, densityArray);
    checkCudaError(cudaDeviceSynchronize());

    //Calculate Velocity
    calculateVellocity<<<(numParticles + 255) / 256, 256>>>(numParticles, dPosition, dVelocityCurrent, dVelocityNew, cellGrid, cellCountArray, densityArray, dt);
    checkCudaError(cudaDeviceSynchronize());

    //Calculate Position
    calculatePosition<<<(numParticles + 255) / 256, 256>>>(numParticles, dPosition, dVelocityNew, dt);
    checkCudaError(cudaDeviceSynchronize());

    dVelocityCurrent = dVelocityNew;
}