Update Smoke Particles in a CUDA Kernel

The smoke update step is now performed on the GPU using the following method:

updateSmoke<<<gpuGridSize, gpuBlockSize>>>(gpuParticleArray, MAX_PARTICLES, interval );

to the kernel function:


/*Update smoke particle positions*/
__global__ void
updateSmoke(particle *particle, int timeInterval) {

int i = blockIdx.x * blockDim.x + threadIdx.x;

//Pos += speed.x * lifetime/interval
particle[i].x += ((particle[i].xv * (1.0f-particle[i].life))/LIFE_TIME) * timeInterval;// Greater dispersion for older particles
particle[i].y += (particle[i].yv/LIFE_TIME) * timeInterval;
particle[i].z += ((particle[i].zv * (1.0f-particle[i].life))/LIFE_TIME) * timeInterval;  // Greater dispersion for older particles

// Speed change based on gravity
particle[i].xv += particle[i].xg * timeInterval;
particle[i].yv += particle[i].yg * timeInterval;
particle[i].zv += particle[i].zg * timeInterval;

//Update the lifetime based on the fade/mass value.
particle[i].life -= particle[i].fade * timeInterval;
}

which was adapted from Encelo's smoke tutorial (http://encelo.netsons.org/programming/opengl) and modified into a CUDA kernel for parallel access of particles.

The demo can now render up to about 8,000 smoke particles (combined with the water particles and terrain) and stay around 30 frames per second. The limitation that I've encountered now is in handling the re-initialization of dispersed particles. Since the method I use relies on the standard random number generator in C, it cannot be used in a CUDA kernel. I've looked at implementing a basic random number generator kernel in CUDA, but have yet to obtain desirable results. 

Here's a shot of the particles in action:

Next week I plan on finishing the random number generator in CUDA to improve re-initialization of particles and then proceeding to fire simulation.