CUDA入门教程
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 编程
Ubuntu 12.04 安装 CUDA-5.5
Ubuntu 11.10 上安装CUDA开发环境
Fedora 15系统下配置CUDA环境
Ubuntu 11.04 安装 NVIDIA CUDA 4.0 RC2
Linux Mint 13/Ubuntu 12.04 配置CUDA 4.2 & OpenCV 2.4.2 方法
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!\n");
return 0;
}
上述代码编译后运行生成可执行文件cuda_test，运行cuda_test后将输出：
Hello World!
注意：
- 调用kernel时需要三个尖括号
- 包含必要的头文件
CUDA C/C++中引入的新关键字__global__所修饰的函数有以下两方面含义：
- 此函数代码由设备执行
- 此函数由主机代码调用
nvcc将源代码分为设备函数和主机函数两大类：
- 设备函数由NVIDA编译器编译
- 主机函数由主机上配置的编译器编译
例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函数将被主机调用
- 设备指针指向GPU内存的某个位置。设备指针可以从主机端传给设备端或者从设备端传给主机端，但是设备指针不能在主机端解引用。
- 主机指针指向CPU内存的某个位置。主机指针可以从设备端传给主机端或者从主机端传给设备端，但是主机指针不能在设备端解引用。
3 块(Blocks)
GPU是用来实现大规模并行的，如何实现呢？将上述例子扩展一下，如果我们要实现两个向量相加： add<<<1,1>>>() ---> add<<<N,1>>> N表示同时调用N次add函数，这样就可以实现并行的向量相加了。 每个被并行调用的add函数称之为一个块(block)。- 块的集合称之为网格(grid).
- 每个块可以使用索引值blockIdx.x
#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\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来定位当前执行的是循环的哪个部分。
从硬件的角度看，相当于同时有多个块在并行执行： 块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] ....更多详情见请继续阅读下一页的精彩内容：
|
评论暂时关闭