为了避免每次都要键入nvcc的命令，要准备一个makefile。makefile如下：

CUFLAG = -g -Xcompiler -v \
-gencode=arch=compute_20,code=sm_20\
-gencode=arch=compute_20,code=compute_20\
-O2
IFLAGS = -I$(CUDA_DIR)/include -I$(CUDA_SDK_DIR)/C/common/inc -I../include
LFLAGS = -L$(CUDA_DIR)/lib64 -L$(CUDA_SDK_DIR)/C/lib
PRG = cuda_test
$(PRG) : main.cu
nvcc main.cu -o $(PRG) $(CUFLAG) $(IFLAGS) $(LFLAGS)

宏观上看，GPU执行代码的流程如下：

例1: Hello World

#include
#include
#include
#include

__global__ void mykernel(void) {
}

int main(void) {
mykernel<<<1,1>>>();
printf("Hello World!\n");
return 0;
}

上述代码编译后运行生成可执行文件cuda_test，运行cuda_test后将输出：

Hello World!

注意：

CUDA C/C++中引入的新关键字__global__所修饰的函数有以下两方面含义：

nvcc将源代码分为设备函数和主机函数两大类：

例2: 整数相加

#include
#include
#include
#include

__global__ void integer_add(int * a, int * b, int * c) {
*c = *a + *b;
}

int main(void) {
int a,b,c;
int * d_a, * d_b, * d_c;
int size = sizeof(int);
cudaMalloc((void**)&d_a,size);
cudaMalloc((void**)&d_b,size);
cudaMalloc((void**)&d_c,size);
printf("Enter two integers with a space to separate them:\n");
scanf("%d %d",&a,&b);
cudaMemcpy(d_a,&a,size,cudaMemcpyHostToDevice);
cudaMemcpy(d_b,&b,size,cudaMemcpyHostToDevice);
integer_add<<<1,1>>>(d_a,d_b,d_c);
cudaMemcpy(&c,d_c,size,cudaMemcpyDeviceToHost);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
printf("Sum is %d\n",c);
return 0;
}

__global__修饰的integer_add函数说明：

#include
#include
#include
#include
#include

#define N 512

__global__ void vec_block_add(int * a, int * b, int * c) {
c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

void rand_ints(int * arr, int count) {
srand(time(NULL));
for(int i=0;i arr[i] = rand() % 100;
}
}

int main(void) {
int * a,* b,* c;
int * d_a, * d_b, * d_c;
int size = N * sizeof(int);
cudaMalloc((void**)&d_a,size);
cudaMalloc((void**)&d_b,size);
cudaMalloc((void**)&d_c,size);

a = (int *) malloc(size);
rand_ints(a,N);
b = (int *) malloc(size);
rand_ints(b,N);
c = (int *) malloc(size);

cudaMemcpy(d_a,a,size,cudaMemcpyHostToDevice);
cudaMemcpy(d_b,b,size,cudaMemcpyHostToDevice);
vec_block_add<<>>(d_a,d_b,d_c);
cudaMemcpy(c,d_c,size,cudaMemcpyDeviceToHost);

#if 1
for(int i=0;i printf("%-5d: a:%-5d b:%-5d c:%-5d\n",i,a[i],b[i],c[i]);
}
#endif

cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);

free(a);
free(b);
free(c);
return 0;
}

例3中最关键的代码为如下几行：

__global__ void vec_block_add(int * a, int * b, int * c) {
c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

由于函数是并行执行的，和传统的串行程序在integer_add函数中使用循环来完成加法相比，相当于由GPU这个加速器使用硬件的方式进行了循环展开，展开后便可以并行执行了。所以在编写这段代码时，需要使用blockIdx.x来定位当前执行的是循环的哪个部分。