ggaaooppeenngg

为什么计算机科学是无限的但生命是有限的

CUDA 的计算模型

发表于 更新于 Disqus：

我的一点观点

我个人觉得 nvidia 的 CUDA 太封闭了，不是很能明白这样封闭的产品怎么能够长久生存，仅仅是因为家大业大么，如果有家公司推出了八成性能的 GPU，但是整套开发的生态非常友好，是不是会像 Android 取代诺基亚一样，还是 nvidia 就是苹果，就算封闭环境也能保证强劲体验，这我也不好说了。

安装

CUDA+cudnn 装起来挺麻烦的，反正如果有错误的话，可以检查 CUDA samples 里面的 deviceQuery 是否成功，如果不成功可以用 strace 看一下少了什么东西，再想办法安装上去。检查 cudnn samples 是否成功也是一样的，里面有一个 mnistDNN 的例子。

Hello World!

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
#include <stdio.h>

__global__ void helloFromGPU (void)
{
	printf("Hello world!\n");
}

int main(void)
{
	printf("Hello World! from CPU\n");

	helloFromGPU <<<1, 10>>>();
	cudaDeviceReset();
	return 0;
}

以上就是一段 GPU 的 Hello world，执行下面的代码可以看到我们在 GPU 上并行执行了十个 “Hello world!” 的打印。

1
2
3
4
5
6
7
8
9
10
11
12
13
$ nvcc -arch sm_20 hello.cu -o hello
$ ./hello
Hello World! from CPU
Hello world!
Hello world!
Hello world!
Hello world!
Hello world!
Hello world!
Hello world!
Hello world!
Hello world!
Hello world!

CUDA

CUDA 全称 Compute Unified Device Architecure，用于定义 GPU 的架构标准。GPU 的工作方式主要依赖于多核心的并行计算，CUDA 提供了方便的模型进行这种模式的编程，下面就会简单介绍一下 CUDA 的架构以及基于 GPU 的编程。

CUDA 一次启动的线程称为网格(grid)，网格中包含块，每个块包含新程，是一个二维的模式，这张图片就说明得很清晰，首先是由 Grid 组成，然后每个 Block 有 shared memory，同 block 的线程可以访问 shared memory，不同的 Block 的线程只能访问全局的内存，这种结构也方便设计和实现并行算法。

cuda-arch

CUDA 的编译器 nvcc 是 gcc 的一个扩展，支持编写运行在 GPU 上的函数，其中的，<<x, y>> 扩展就是用来指定 block 和 thread 的数量的。

比如典型的俩个向量相加的例子：

1
2
3
4
void sumArrayOnHost(float *A, float *B, float *C, const int N) {
    for (int i = 0; i < N; i++)
        C[i] = A[i] + B[i]
}

但是 GPU 的核函数怎么写呢，是这样的。

1
2
3
4
__global__ void sumArrayOnGPU(float *A, float *N, float C) {
    int i = threadIdx.x
    C[i] = A[i] + B[i]
}

然后，如果 N 是 32 的话，可以如下调用，__global__ 表示这个函数可以在 host 调用也可以在 GPU 上调用，用个 32 个线程的 block 算这个向量和，N 隐性包含在了定义中。

1
sumArrayOnGPU<<1, 32>>(float *A, float *B, float *C)

线程组织

基于 CPU 和 GPU 的异构计算平台可以优势互补，CPU 负责处理逻辑复杂的串行程序，而 GPU 重点处理数据密集型的并行计算程序，从而发挥最大功效。

线程组织主要依据网格（现在网络这个词其实比较容易混淆，可以是神经网络可以是计算机网络，这里指的是线程的组织形式）模型。主要分 grid, block, thread，组织的方式就靠索引来决定，可以通过 block 索引和 thread 索引进行线程的定位。CUDA 提供了块内线程同步的方法，但是没有提供块间同步的原语。

线程束分化

线程束 = warp

CPU 有很强的分支预测的能力，会预加载指令，如果预测正确，执行代价就很小，但是 GPU 在这方面是很弱的，因为线程束当中如果执行分支出现不同，那出现分支的线程就会被禁止执行，这也是用 GPU 做并行编程的时候需要参考的一个很重要的因素，保证并行的线程出现线程束分化的情况尽量少。

下面就是线程束分化的例子。

并行归约问题

并行归约是指如何对并行问题进行归约，比如说相邻配对和交错配对。

设备管理

管理 GPU 主要是通过 CUDA API 或者 nvidia-smi 命令来获取。

SM（流式多处理器）

上面讲的是抽象上的分层，但是实际的物理层面承载 GPU 的是 SM。最早的 GPU 架构叫 Fermi，然后是 Kepler，然后才是 Tesla。一般说的 cm_20, sm_20 就指这种计算能力和架构，新款的 GPU 计算能力要更强一点。

内存模型

CUDA 的内存模型和 CPU 是类似的也是多级的结构。线程有局部内存，块之间有全局内存。GPU 和 CPU 都是主要采用 DRAM 来做主存，CPU 的一级缓存会用 SRAM 做。CPU 的多级缓存对用户来说不是很需要考虑，尽量屏蔽了其中的细节，但是 GPU 相对来说会把这种分级结构暴露给用户，这对编程来说也是一种新的挑战。

后面

用 GPU 计算一些计算密集型的程序，速度是真快，比 CPU 块很多，所以这也是为什么大量深度学习的应用都要通过 GPU 加速的原因。