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. 

Smoke Particle Effect

The current particle smoke calculations are performed on a single processor and limited to 2000 particles per smoke instance.  After 10,000 particles are on screen the performance starts to decrease much below 30 fps. I'm currently working on implementing a 3D CUDA function that will perform the smoke step operation, but I'm not getting proper data out of it yet:

pos.x += speed.x * lifetime/interval
pos.y += speed.y * interval
pos.z += speed.z * lifetime/interval

and then update the lifetime based on the fade/mass value.

I currently have a few GLUI controls set up for changing parameters of the smoke. They are fade/mass, size and colour. Below are some screenshots with different settings:

I'm hoping for the ability to render up to 30,000 particles (or more) with CUDA step operation compared to the current frame-rate using single instructions. There are also many more GLUI controls that could be added to change the affect of gravity on the smoke and its speed.

CUDA Smoke Simulation

The main differences in a fluid solver for smoke compared to water are as follows:

• Decrease in viscosity
• Addition of density, mass and temperature factors
• Buoyancy based on the additional factors

Obviously buoyancy is related to water as well, but since the water I rendered was on a fixed plane in 2D this factor was not necessary. Buoyancy will be needed for smoke movement affecting its position along the Y-axis. GPU Gems (Chpt. 38)  refers to the buoyancy force as follows: 

where k is the mass factor and d is the smoke density - s  a constant time scale factor and T the temperature compared to To the ambient temperature.

I plan to implement this function in my current CUDA kernel by adding the value at each fragment to the velocity grid in the Y-direction. This should simulate the effect of gravity based on the weight of the smoke.

GLSL Bump Mapping

Following the 'Swiftless' GLSL tutorial at http://www.swiftless.com/tutorials/glsl/8_bump_mapping.html bump mapping is now working on the water while stationary. Generating a normal map while animating still gives a frame rate less than 1 however, making the demo unusable.

While stationary, the bump map is convincing:

I'm thinking it would be possible to calculate and output the normal map only when saving the actual textures to disk. That way the animation would be less choppy while interacting with it.

At this point I've added  GLUI controls to adjust the water level and colour, as well as assigning record and print keys for animation. I'm going to move onto smoke simulation next, and likely work out the normal map output at some point.

GLSL Multitexturing

After a bit of a refresh on OpenGL 2's GLSL, I managed to get a multitexturing fragment shader working for water. The fragment code is as follows:


varying vec3 lightDir,normal;
uniform sampler2D color_texture;
uniform sampler2D color_texture2;

void main() {
float diffuse, alpha;

diffuse = max(dot(lightDir,normalize(normal)),0.0);

diffuse = diffuse * gl_FrontMaterial.ambient.rgb;
alpha = gl_FrontMaterial.diffuse.a;

vec3 color = diffuse * (texture2D(color_texture, gl_TexCoord[0].st).rgb +
           texture2D(color_texture2, gl_TexCoord[0].st).rgb));

gl_FragColor = vec4(color, alpha);
}

The results for various coloured textures are as follows: