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）Ubuntu 12.04 下 CUDA 编程 http://www.linuxidc.com/linux/2014-06/103056.htmUbuntu 12.04 安装 CUDA-5.5 http://www.linuxidc.com/Linux/2013-10/91101.htmUbuntu 11.10 上安装CUDA开发环境 http://www.linuxidc.com/Linux/2012-04/58913.htmFedora 15系统下配置CUDA环境 http://www.linuxidc.com/Linux/2011-12/49874.htmUbuntu 11.04 安装 NVIDIA CUDA 4.0 RC2 http://www.linuxidc.com/Linux/2011-10/46304.htmLinux Mint 13/Ubuntu 12.04 配置CUDA 4.2 & OpenCV 2.4.2 方法 http://www.linuxidc.com/Linux/2013-10/91102.htm

2 异构计算（Heterogeneous Computing）

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

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

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

将输入数据通过PCI总线从CPU内存拷贝到GPU的DRAM中。
从内存中加载需要执行的代码到GPU后。
数据和指令都就绪后，就可以执行了。注意，在执行的过程中，GPU会在片上缓存数据以提升性能。
计算完毕后，将结果从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！ "）;
return 0;
}上述代码编译后运行生成可执行文件cuda_test，运行cuda_test后将输出： Hello World！注意：

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

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: "）;
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 ",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 块（Blocks）

GPU是用来实现大规模并行的，如何实现呢？将上述例子扩展一下，如果我们要实现两个向量相加：add<<<1,1>>>（） ---> add<<<N,1>>>N表示同时调用N次add函数，这样就可以实现并行的向量相加了。每个被并行调用的add函数称之为一个块（block）。

块的集合称之为网格（grid）.
每个块可以使用索引值blockIdx.x

通过使用blockIdx.x作为索引，每个块可以处理数组元素中的一部分。有了这些基础后，就可以实现并行版本的向量相加了。例3:向量相加#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<cutil.h>
#include<time.h>
#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<count;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<<<N,1>>>（d_a,d_b,d_c）;
cudaMemcpy（c,d_c,size,cudaMemcpyDeviceToHost）;

#if 1
for（int i=0;i<N;i++） {
printf（"%-5d: a:%-5d b:%-5d c:%-5d ",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来定位当前执行的是循环的哪个部分。从硬件的角度看，相当于同时有多个块在并行执行：块0: c[0]=a[0]+b[0]块1: c[1]=a[1]+b[1]块2: c[2]=a[2]+b[2]块3: c[3]=a[3]+b[3]....更多详情见请继续阅读下一页的精彩内容： http://www.linuxidc.com/Linux/2014-07/104328p2.htm