在网上看到一篇比较不错的CUDA工程模板程序的详细分析文章，大家可以用来了解一个CUDA程序的基本结构。转载于http://blog.csdn.net/darkstorm2111203/archive/2008/08/22/2813480.aspx。

/* Template_Host.c 用于演示如何生成cuda工程的样本程序 */

/* 主机端，也就是cpu code*/

// includes, system
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>

// includes, project
#include <cutil.h> //调用CUDA扩展必须的头文件
// includes, kernels
#include “template_kernel.cu”
////////////////////////////////////////////////////////////////////////////////
// declaration, forward
void runTest( int argc, char** argv);

extern “C” //表示使用C编译器编译下面函数
void computeGold( float* reference, float* idata, const unsigned int len);
////////////////////////////////////////////////////////////////////////////////
// Program main
// host端程序是在CPU上运行的，可以在其中加入一些C代码，不同的是多了一些CUDA扩展
////////////////////////////////////////////////////////////////////////////////
int main( int argc, char** argv)
{
runTest( argc, argv); //经过封装的显卡计算程序
CUT_EXIT(argc, argv); //退出CUDA，停止使用显卡进行计算
}
////////////////////////////////////////////////////////////////////////////////
//! Run a simple test for CUDA
////////////////////////////////////////////////////////////////////////////////

void runTest( int argc, char** argv)
{
CUT_DEVICE_INIT(argc, argv); //初始化设备，使用多卡时应该加上设备号，或者使用cudasetdevice()函数

unsigned int timer = 0;
CUT_SAFE_CALL( cutCreateTimer( &timer)); //定义一个计时器
CUT_SAFE_CALL( cutStartTimer( timer)); //开始计时

unsigned int num_threads = 32; //定义每个block中的线程数
unsigned int mem_size = sizeof( float) * num_threads; //需要分配的存储器大小，这里每个线程只处理一个浮点数

// allocate host memory
float* h_idata = (float*) malloc( mem_size); //在主机内存上分配空间，前缀h_表示host

// initalize the memory
for( unsigned int i = 0; i < num_threads; ++i){
h_idata[i] = (float) i;  //初始化内存中的数值
}

// allocate device memory
float* d_idata; //定义指针，这个指针将指向显卡上的显存，前缀d_表示device，i表示input
CUDA_SAFE_CALL( cudaMalloc( (void**) &d_idata, mem_size)); //在显卡上分配空间

// copy host memory to device
CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, mem_size, cudaMemcpyHostToDevice) );
//将内存中h_idata中的值拷贝到显存中的d_idata里，这样就完成了主机对设备的数据写入

// allocate device memory for result
float* d_odata; //定义指针，这个指针将指向显卡上的显存，前缀d_表示device，o表示output
CUDA_SAFE_CALL( cudaMalloc( (void**) &d_odata, mem_size)); //在显卡上分配空间

// execute the kernel
testKernel<<< grid, threads, mem_size >>>( d_idata, d_odata);
/*调用核函数，其中那个夸张的<<< >>>中是调用核函数的一些设置，，而( )括号里是程序的的参数列表。<<< >>>里第一个参数是dim3类型的grid网格形状，第二个参数是block的形状，第三个参数表示给核函数分配的shared memory大小。因为我们的grid和threads实际上是一维的，因此以<<< 1, num_threads, mem_size>>>的形式调用testKernel()效果是一样的。（）中的指针必须是指向显存的，而其他参数如float, int等则不必传到显卡。*/

// check if kernel execution generated and error
CUT_CHECK_ERROR(“Kernel execution failed”); //用于检查错误

// allocate mem for the result on host side
float* h_odata = (float*) malloc( mem_size); //在内存上开辟空间准备接收显卡计算得到的结果
// copy result from device to host
CUDA_SAFE_CALL( cudaMemcpy( h_odata, d_odata, sizeof( float) * num_threads, cudaMemcpyDeviceToHost) );
//将显存中的数据拷贝回内存

CUT_SAFE_CALL( cutStopTimer( timer)); //结束计时
printf( “Processing time: %f (ms)\n”, cutGetTimerValue( timer));//显示时间
CUT_SAFE_CALL( cutDeleteTimer( timer));//删除timer

// compute reference solution
float* reference = (float*) malloc( mem_size);//在内存上分配空间，用于计算cpu结果
computeGold( reference, h_idata, num_threads);//computeGold函数在computeGold.c中，是一个与GPU计算功能相同的对照程序

// check result
if( cutCheckCmdLineFlag( argc, (const char**) argv, “regression”))
//判断gpu程序接收到的消息中是否设定需要纪录显卡计算结果，否则比较cpu与gpu计算结果
{
// write file for regression test
CUT_SAFE_CALL( cutWriteFilef( “./data/regression.dat”, h_odata, num_threads, 0.0));
}
else
{
// custom output handling when no regression test running
// in this case check if the result is equivalent to the expected soluion
CUTBoolean res = cutComparef( reference, h_odata, num_threads);
printf( “Test %s\n”, (1 == res) ? “PASSED” : “FAILED”);
}

// cleanup memory
free( h_idata);//释放程序使用的内存空间，否则会造成内存溢出
free( h_odata);
free( reference);
CUDA_SAFE_CALL(cudaFree(d_idata));//释放使用的显存空间，否则会造成显存溢出，多次运行后显卡无法工作
CUDA_SAFE_CALL(cudaFree(d_odata));
}

/*  Template_Kernel.c 用于演示如何生成cuda工程的样本程序
* 设备端，也就是gpu code */

#ifndef _TEMPLATE_KERNEL_H_
#define _TEMPLATE_KERNEL_H_

#include <stdio.h> //在emu模式下可以有stdio以便输出一些中间结果来观察，由gpu运行时是不能使用其中的函数的
#define SDATA( index) CUT_BANK_CHECKER(sdata, index)

//主要是在emu模式下检查bank访问,如果程序最终从gpu运行，这一句不是必要的，此时程序中的SDATA()应该替换成s_data[]的形式

////////////////////////////////////////////////////////////////////////////////
//! Simple test kernel for device functionality
//! @param g_idata input data in global memory
//! @param g_odata output data in global memory
////////////////////////////////////////////////////////////////////////////////

__global__ void testKernel( float* g_idata, float* g_odata)
/*对应host代码中的testKernel<<< grid, threads, mem_size >>>( d_idata, d_odata)，注意到同样的指针在host端中前缀是d，而在kernel函数中，相同的指针前缀却变成了g，这是由于观察角度不同造成的。在host端中，程序编写者是从cpu的角度思考，因此显卡只是一个协处理设备；而编写kernel程序时，我们要养成从gpu的角度思考的习惯，因此显卡上的显存也就被看成了global memory*/
{
// shared memory
// the size is determined by the host application
extern __shared__ float sdata[];
/*定义每个block中使用的shared_mem的大小，此处extern表示shared memory的大小由外部定义，也就是在host端代码调用kernel函数<<<>>>中的第三个参数mem_size*/

// access thread id
const unsigned int tid = threadIdx.x;
//定义常数 tid，储存在GPU端寄存器中，threadIdx.x是每个线程在block中的x轴编号，即（0，1，2.. blockDim.x-1）
// access number of threads in this block

const unsigned int num_threads = blockDim.x;
//定义常数， blockDim.x就是dim3 threads( num_threads, 1, 1)的第一个参数，即32

// read in input data from global memory
// use the bank checker macro to check for bank conflicts during host
// emulation
SDATA(tid) = g_idata[tid]; //从global memory（显存）拷贝到shared memory
__syncthreads(); //输入数据后同步一次，保证计算时所有数据均已到位

// perform some computations
SDATA(tid) = (float) num_threads * SDATA( tid); //实际上就是一个32*i
__syncthreads(); //输入数据后同步一次，保证输出前所有数据已经计算完成

// write data to global memory
g_odata[tid] = SDATA(tid);//将shared中计算完的数据写入g_data中
}

#endif // #ifndef _TEMPLATE_KERNEL_H_

本文来自: CUDA程序初窥