A Simple CUDA Kernel

The first example kernel draws a time-varying sinusoidal surface using a simple adaptation of the NVIDIA kernel provided in the simpleGL SDK example. Essentially, each height value stored in pos is updated based on position in the mesh, a frequency and the animation time passed via the parameter time.

    // calculate simple sine wave pattern
    float freq = 4.0f;
    float w = sinf(u*freq + time) * cosf(v*freq + time) * 0.5f;

The position and height information is stored in the float4 position array:

    // write output vertex
    pos[y*width+x] = make_float4(u, w, v, 1.0f);

Similarly, the colors are calculated based on position in the mesh and the animation time:

    // write the color
    colorPos[y*width+x].w = 0;
    colorPos[y*width+x].x = 255.f *0.5*(1.f+sinf(w+x));
    colorPos[y*width+x].y = 255.f *0.5*(1.f+sinf(x)*cosf(y));
    colorPos[y*width+x].z = 255.f *0.5*(1.f+sinf(w+time/10.f));

The complete listing for kernelVBO is as follows:

// kernelVBO.cpp (Rob Farber)
// Simple kernel to modify vertex positions in sine wave pattern
__global__ void kernel(float4* pos, uchar4 *colorPos,
		       unsigned int width, unsigned int height, float time)
{
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    // calculate uv coordinates
    float u = x / (float) width;
    float v = y / (float) height;
    u = u*2.0f - 1.0f;
    v = v*2.0f - 1.0f;

    // calculate simple sine wave pattern
    float freq = 4.0f;
    float w = sinf(u*freq + time) * cosf(v*freq + time) * 0.5f;

    // write output vertex
    pos[y*width+x] = make_float4(u, w, v, 1.0f);
    colorPos[y*width+x].w = 0;
    colorPos[y*width+x].x = 255.f *0.5*(1.f+sinf(w+x));
    colorPos[y*width+x].y = 255.f *0.5*(1.f+sinf(x)*cosf(y));
    colorPos[y*width+x].z = 255.f *0.5*(1.f+sinf(w+time/10.f));
}

// Wrapper for the __global__ call that sets up the kernel call
extern "C" void launch_kernel(float4* pos, uchar4* colorPos,
			      unsigned int mesh_width, unsigned int mesh_height, float time)
{
  // execute the kernel
  dim3 block(8, 8, 1);
  dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);
  kernel<<< grid, block>>>(pos, colorPos, mesh_width, mesh_height, time);
}

The program can be compiled under Linux with the following after the complete source of each of the files has been saved in correctly named files:

#/bin/bash
SDK_PATH= PATH_TO_NVIDIA_GPU_Computing_SDK

nvcc -O3 -L $SDK_PATH/C/lib -I $SDK_PATH/C/common/inc simpleGLmain.cpp simpleVBO.cpp callbacksVBO.cpp kernelVBO.cu -lglut -lGLEW -lcutil_x86_64 -o testGL

Microsoft users should consult the GPUcomputing.net website for information about building the examples with Microsoft Visual Studio.

Figures 10, 11, and 12 are images created with this program. You will likely see something similar, but the colors and shape of the surface do dynamically change with time. Also, the orientation of the surface in three-space also affects what is shown on the screen.

Figure 10: Example surface sinusoidal surface.

Figure 11: Wire-frame image.

Figure 12: Points drawn in 3D.