CUDA： 简单入门

1. 准备makefile

为了避免每次都要键入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)

2 异构计算(Heterogeneous Computing)

以下为几个技术名词的简单介绍：

• 主机(host)：CPU及其内存(host memory)。
• 设备(device)：GPU及其内存(device memory)。
• 主机代码(host code)：运行在CPU上的（一般来说「串行执行」的）代码。
• 设备代码(device code)：运行在GPU上的并行执行的代码。
• 异构计算：由主机代码(host code)和设备代码(device code)协同执行完成的计算。

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

1. 将输入数据通过PCI总线从CPU内存拷贝到GPU的DRAM中。
2. 从内存中加载需要执行的代码到GPU后。
3. 数据和指令都就绪后，就可以执行了。注意，在执行的过程中，GPU会在片上缓存数据以提升性能。
4. 计算完毕后，将结果从GPU的DRAM中拷回CPU的Memory中。

例1: Hello World

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

__global__ void mykernel(void) {
}

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

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

Hello World!

注意：

1. 调用kernel时需要三个尖括号
2. 包含必要的头文件

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

• 此函数代码由设备执行
• 此函数由主机代码调用

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

• 设备函数由NVIDA编译器编译
• 主机函数由主机上配置的编译器编译

三个尖括号标志着一个从主机代码调用设备代码的函数，称为“启动内核”(kernel launch)

例2: 整数相加

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

__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函数说明：

• integer_add函数将在设备上执行
• integer_add函数将被主机调用

由于integer_add函数在设备上执行，所以指针a，b，c应该指向设备内存。这说明需要在设备内存中为变量开辟内存。

设备内存和主机内存在物理上是完全分开的不同电子部件：

• 设备指针指向GPU内存的某个位置。设备指针可以从主机端传给设备端或者从设备端传给主机端，但是设备指针不能在主机端解引用。
• 主机指针指向CPU内存的某个位置。主机指针可以从设备端传给主机端或者从主机端传给设备端，但是主机指针不能在设备端解引用。

CUDA API提供的用于处理设备内存的函数有cudaMalloc, cudaFree, cudaMemcpy。语义上分别对应于C语言的malloc, free, memcpy函数。这几个函数的具体使用方法如例2所示。

3 块(Bloc