前言
一般需要性能优化的任务分为数据密集型任务和计算密集型任务,数据密集型任务是指需要处理和分析大量数据的任务。最常见的数据密集型任务应该是数据库操作了,这类任务的性能瓶颈一般在数据的读取和写入速度,因此对应的性能优化手段一般为:
- 使用更好的资源,比如SSD(甚至可以把数据都放在内存里)
- 使用数据结构(比如索引)来提高查询效率
- 使用缓存
而计算型密集任务是指需要需要大量计算资源的任务,这类任务的性能瓶颈通常在CPU和GPU的计算能力,这类任务最常见的优化手段就是并行计算了。并行计算区别于并发,是指多个处理器同时执行多个计算任务,而GPU的处理器数量远大于CPU,因此GPU更适合并行计算。
异构计算
GPU设计初衷是为了实现在并行执行数千个线程时达到尽可能高的性能,而CPU的设计初衷是为了执行一系列操作时达到尽可能高的性能。因此GPU将更多的晶体管用于数据计算而不是缓存或控制流
实际任务中可能会有的地方适合使用CPU,有的地方适合使用GPU,而异构计算就是使用多个不同类型的处理器,这些处理器会去承担任务的不同部分。
使用CUDA进行并行计算
那么我们使用CUDA的思路就可以被简单的抽象为三个流程
- 将数据从CPU中复制到GPU(GPU默认和CPU数据不共享)
- 通过GPU进行并行计算
- 将GPU计算后的结果复制到CPU内存中。
举个例子
将数据从CPU中复制到GPU
我们先进行数据的初始化
int main(){
int *a, *b, *c; // host point
int *d_a, *d_b. *d_c;
int size = N * sizeof(int);
// alloc space for device copies;
cudaMalloc((void **)&d_a, size);
cudaMalloc((void **)&d_a, size);
cudaMalloc((void **)&d_a, size);
// alloc space for host copies and setup input values
a = (int *)malloc(size); random_inits(a, N);
b = (int *)malloc(size); random_inits(b, N);
c = (int *)malloc(size);
}
这里有一个约定俗成的概念,通常我们将CPU中的内存称作host memory,GPU的显存称作device memory,因此当设置GPU中的指针时,我们命名用
d_作为前缀。
接下来将数据从CPU中复制到GPU中
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
以下几点需要注意:
- 这里使用了cuda提供的api
cudaMemcpy完成数据复制 - 注意到有四个参数,前两个是目标指针和源指针,第三个size 表示数据的大小。而第四个参数,标明了数据是从CPU -〉GPU。这是多余的吗?并不是,指针本身不携带任何元数据,而不同的处理器对于数据的取值未必相同,因此表明数据的流转是很有必要的。
通过GPU进行并行计算
这里以向量相加作为一个入门例子,向量相加是并行计算的一个很简单但很经典的模式,不需要考虑数据通信,不需要考虑数据同步,所有的数据使用统一的计算方式,单纯的并行计算即可。
__global__ void add(int *a, int *b, int *c){
c[blockIdx.x ] = a[blockIdx.x ] + b[blockIdx.x]; // 很多时候,未必是一维向量,因此可能需要获取全局唯一索引
}
add<<<M,N>>>(d_a, d_b, d_c);
以下几点需要注意:
-
__global__下划线global 表明函数在GPU运行且会被host code 调用 -
我们使用
nvcc编译器驱动程序,分离源代码到CPU和GPU上。a. 而三个箭头则是声明这段代码会在GPU上运行。
-
M.N 分别代表几个线程块,线程块内有几个线程(这里会在之后的文章更详细的讲解)
将GPU计算后的结果复制到CPU内存中。
cudaMemcpy(c, d_c, size, cudaMemcpDeviceToToken);
总结
本文对CUDA进行了一个简单的介绍, 并提供了一个简单的并行计算的例子,后续会进一步的系统的对CUDA进行讲解