 Dr. Dobb's is part of the Informa Tech Division of Informa PLC

This site is operated by a business or businesses owned by Informa PLC and all copyright resides with them. Informa PLC's registered office is 5 Howick Place, London SW1P 1WG. Registered in England and Wales. Number 8860726.

# CUDA, Supercomputing for the Masses: Part 18

### 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.

### More Insights To upload an avatar photo, first complete your Disqus profile. | View the list of supported HTML tags you can use to style comments. | Please read our commenting policy.

## Featured Reports ## Featured Whitepapers 