My first CUDA program!

Note: Check out “CUDA Gets Easier” for a simpler way to create CUDA projects in Visual Studio.

I got CUDA setup and running with Visual C++ 2005 Express Edition in my previous post. Now I’ll write my first CUDA program. It’s a modification of an example program from a great series of articles on CUDA by Rob Farber published in Dr. Dobbs Journal. Rob does his examples in a make-based build environment; I’ll show how to build a CUDA program in the Visual C++ IDE.

Simple CUDA programs have a basic flow:

1. The host initializes an array with data.
2. The array is copied from the host to the memory on the CUDA device.
3. The CUDA device operates on the data in the array.
4. The array is copied back to the host.

My first CUDA program, shown below, follows this flow. It takes an array and squares each element. I can barely contain my excitement.

// example1.cpp : Defines the entry point for the console application.
//

#include "stdafx.h"

#include <stdio.h>
#include <cuda.h>

// Kernel that executes on the CUDA device
__global__ void square_array(float *a, int N)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx<N) a[idx] = a[idx] * a[idx];
}

// main routine that executes on the host
int main(void)
{
  float *a_h, *a_d;  // Pointer to host & device arrays
  const int N = 10;  // Number of elements in arrays
  size_t size = N * sizeof(float);
  a_h = (float *)malloc(size);        // Allocate array on host
  cudaMalloc((void **) &a_d, size);   // Allocate array on device
  // Initialize host array and copy it to CUDA device
  for (int i=0; i<N; i++) a_h[i] = (float)i;
  cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
  // Do calculation on device:
  int block_size = 4;
  int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);
  square_array <<< n_blocks, block_size >>> (a_d, N);
  // Retrieve result from device and store it in host array
  cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
  // Print results
  for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
  // Cleanup
  free(a_h); cudaFree(a_d);
}

Two pointers are declared on line 19 of the main routine: a_h points to the array that is stored on the host, while a_d points to the array on the CUDA device. The a_h array is allocated in the host memory on line 22 using the standard malloc subroutine, but a_d is allocated in the CUDA device memory using the cudaMalloc subroutine found in the CUDA API (line 23). (Note that a pointer to the a_d pointer is passed to cudaMalloc so it can store the address of the array in a_d.)

In order to create some values to operate upon, each element in the host array is initialized with its array index (line 25). Then the cudaMemcpy subroutine is used to copy a_h from the host into a_d on the CUDA device. (The cudaMemcpyHostToDevice flag, defined in the API, indicates the direction of the transfer.)

In lines 28-30, the host initiates the execution of the kernel function, square_array, on the CUDA device. A CUDA device contains individual processing elements, each of which can execute a thread. A number of the processing elements are grouped together to form a block, and a number of blocks constitutes a grid. In this example, the number of threads per block is set to four (line 28). Then the total number of blocks that are needed to get enough threads to square each array element is calculated on line 29. (For ten array elements, three blocks each with four threads are needed.) On line 30, the host initiates the kernel function on the CUDA device. The number of blocks and the number of threads in each block are indicated between the <<<…>>> following the kernel name. (This information is picked up by the Nvidia compiler, nvcc, and is used when generating the instructions that start the kernel on the CUDA device. More on nvcc, later.) Following that, the standard argument list to square_array contains a pointer to the array in the CUDA device memory and the number of elements in the array.

The kernel is shown on lines 10-14. The __global__ keyword indicates that this is a kernel function that should be processed by nvcc to create machine code that executes on the CUDA device, not the host. In this example, each thread will execute the same kernel function and will operate upon only a single array element. Each thread is distinguished from all the others by block and thread indices that can be used to determine the array element the thread will access. On line 12, the array index is found by multiplying the thread’s block index (blockIdx.x) by the number of threads in each block (blockDim.x) and then adding the index of the thread within the block (threadIdx.x). If the index is within the bounds of the array, then the corresponding array element is squared (line 13).

Immediately after starting the kernel, the host begins a transfer of the data from the array in the CUDA device memory back to the array in the host memory (line 32). This transfer is delayed until the CUDA device has finished executing the kernel, so there is no chance of getting data that has not been processed yet. Then the host displays the contents of the array (line 34) and frees the array memory on both itself and the CUDA device (line 36).

At this point, I have a CUDA-enabled program, but I don’t have it integrated into a Visual C++ project. It actually takes a bit of work to do that. To start, I brought up the Visual C++ 2005 Express Edition IDE and clicked on the New Project button (you can also use File→New→Project… from the menu). In the New Project window, I selected Win32 as the project type and Win32 Console Application as the template. I gave the project the creative name of example1 and set its location to the C:\llpanorama\CUDA\examples directory. After clicking OK in the New Project window, and then clicking Finish in the Win32 Application Wizard window, a window opened with a simple code skeleton. I replaced the code skeleton with the code shown above.

After saving the code, I right-clicked the example1.cpp file, selected Rename from the drop-down menu and renamed the file to example1.cu. Files with the .cu extension are intended to be processed by nvcc. nvcc will extract the kernel portion of example1.cu and compile it for execution on the CUDA device while using the Visual C++ compiler to compile the remainder of the file for execution on the host.

In its default configuration, Visual C++ doesn’t know how to compile .cu file. It has to be told explicitly how to do this using a Custom Build Step. This is done by right-clicking on the example1.cu file and selecting Properties from the drop-down menu. In the Property Pages window that appears, set the Custom Build Step command line as follows:

Configuration Properties → Custom Build Step → General:
Command Line =
“$(CUDA_BIN_PATH)\nvcc.exe” -ccbin “$(VCInstallDir)bin” -c -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/MTd -I”$(CUDA_INC_PATH)” -I./ -o $(ConfigurationName)\example1.obj example1.cu

What does this command line do? Let’s break it down piece-by-piece:

“$(CUDA_BIN_PATH)\nvcc.exe”: The location of the nvcc compiler.

-ccbin “$(VCInstallDir)bin”: The location of the Visual C++ compiler.

-c: The compilation will proceed all the way to the generation of an object file (.obj extension).

-D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS: Various macro definitions.

-Xcompiler /EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/MTd: Various options that are passed by nvcc directly to the Visual C++ compiler.

-I”$(CUDA_INC_PATH)”: Look in the CUDA include directories for needed header files.

-I./: Look in the current directory for needed header files.

-o $(ConfigurationName)\example1.obj: The location and name of the resulting object file.

example1.cu: The source file that the compiler will work on.

In addition to setting the command line for the example1.cu file, the location of the output file is specified as follows:

Configuration Properties → Custom Build Step → General:
Outputs = $(ConfigurationName)\example1.obj

After setting the file properties, the properties for the example1 project have to be modified. Here are the project property settings I used for the Debug configuration:

Configuration Properties → C/C++ → General:
Additional Include Directories = $(CUDA_INC_PATH);”C:\Program Files\NVIDIA Corporation\NVIDIA CUDA SDK\common\inc”

Configuration Properties → C/C++ → General:
Debug Information Format = Program Database (/Zi)

Configuration Properties → C/C++ → Code Generation:
Runtime Library = Multi-threaded Debug (/MTd)

Configuration Properties → Linker → General:
Enable incremental linking = No (/INCREMENTAL:NO)

Configuration Properties → Linker -> General:
Additional Library Directories = “C:\CUDA\lib”;”C:\Program Files\NVIDIA Corporation\NVIDIA CUDA SDK\common\lib”

Configuration Properties → Linker → Input:
Additional Dependencies = cudart.lib cutil32D.lib

Configuration Properties → Linker → Optimization:
Enable COMDAT folding = Do Not Remove Redundant COMDATs (/OPT:NOICF)

Now the project can be compiled and run. Here’s the result:

0 0.000000
1 1.000000
2 4.000000
3 9.000000
4 16.000000
5 25.000000
6 36.000000
7 49.000000
8 64.000000
9 81.000000

I told you it was exciting! Well, at least it’s right.

In order to compile the Release configuration, a few changes need to be made to the file and project properties. For the example1.cu file,the Custom Build Step command line has to be changed to remove the _DEBUG macro definition, enable compiler optimization, and link with the Release runtime library:

Configuration Properties → Custom Build Step → General:
Command Line =
“$(CUDA_BIN_PATH)\nvcc.exe” -ccbin “$(VCInstallDir)bin” -c -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT -I”$(CUDA_INC_PATH)” -I./ -o $(ConfigurationName)\example1.obj example1.cu

The project properties that have to be changed in the Release configuration are the linking for the runtime library and the use of the non-debug version of the CUDA utilities library:

Configuration Properties → C/C++ → Code Generation:
Runtime Library = Multi-threaded (/MT)

Configuration Properties → Linker → Input:
Additional Dependencies = cudart.lib cutil32.lib

Once those changes are made, the Release version of the example1 project can be compiled and run. It will output the same exciting result.

Here’s the source code for this example if you want to try it.

Don’t have a CUDA-capable GPU board on your PC but still want to try running this program? Easy! Just add the following option to the Custom Build Step command line: -deviceemu. This will link-in a CUDA device emulator that runs on the host. The emulator becomes the target for all the CUDA API calls and executes the kernel. The program will run just like a CUDA device is there, except slower.  (Here is the project file with the -deviceemu option.)

So I’ve written my first CUDA program and gotten it to compile using Visual C++ 2005 Express Edition. Setting up the compilation options was as much (more?) work as writing the program, so you might be interested in a CUDA template for Visual C++ 2005 written by kyzhao. The installer doesn’t work for me (maybe because I’m using the free Express Edition), but it might help you.