随着近年来深度学习的爆发,原来被用于图形渲染的GPU被大量用于并行加速深度学习的模型训练中,在这个过程中 CUDA 作为 NVIDIA 推出的基于GPU的一个通用并行计算平台和编程模型也得到了广泛的使用。或许你已经十分了解 现代CPU的体系架构,但是对于GPU还不甚清晰,GPU的体系架构到底和CPU有何区别,CUDA模型是什么,我们该如何使用 CUDA 实现并行计算,本文将为你扫盲祛魅,本文中使用到的所有代码可以在我的 Github 中找到。
在 CPU 算力快速提升的这五十年,人们需要的计算量也同时在迅猛发展着,从最开始的桌面互联网,到后来的移动互联网,以及5年前爆发的深度学习,无一不需要庞大的计算力。在这个过程中,仅仅依靠CPU的算力开始力有不逮,这个过程中像GPU、FPGA、DSP等异构计算单元开始得到广泛的应用。下面,我回归计算的本质,以GPU为例来分析为什么我们需要这些异构计算单元。
无论是 CPU 还是 GPU,我们可以把计算模型抽象为下面这张图,这也是典型的冯诺伊曼体系架构。
影响计算能力的4个主要因素如下:
Parallel Processing:Amount of data processed at one time
Clock Frequency:Processing speed on each data element
Memory Bandwidth:Amount of data transferred at one time
Memory Lantency:Time for each data element to be transferred
对于CPU,依次分析这几个因素:
为了提供并行处理能力,我们从单核单处理器发展到多核多处理器,每个时钟周期CPU也能够处理多条指令
因为CPU时钟频率和功率的关系 $ Power = k ClockFrequency Voltage^2 $ ,在CPU过去的发展历史中,通过提高CPU时钟频率可以变得更快,与此同时为了保持CPU功耗的正常,也需要不断降低电压。但是当主频逐渐逼近到 4GHz 时,电压已经不能再降低了,因为这已经到达了晶体管高低电平反转的极限,关于这部分的更多内容可以参考 摩尔定律 。
基于 CPU + GPU 的异构计算平台可以优势互补,CPU负责处理逻辑复杂的串行程序,GPU重点处理数据密集型的并行计算程序,从而发挥最大功效。CUDA 程序中既包含 Host 程序,又包含 Device 程序,它们分别在CPU和GPU上运行。
同时, Host 与 Device 之间通过PCIe总线交互进行数据拷贝,典型的 CUDA 程序的执行流程如下:
初始化后,将数据从 Main Memory 拷贝到 GPU Memory
CPU 调用 CUDA 的核函数
GPU 的 CUDA Core 并行执行核函数
将 Device 上的运算结果拷贝到 Host 上
GPU核心在做计算时,只能直接从显存中读写数据,程序员需要在代码中指明哪些数据需要从内存和显存之间相互拷贝。这些数据传输都是在总线上,因此总线的传输速度和带宽成了部分计算任务的瓶颈。当前最新的总线技术是NVLink,IBM的 Power CPU 和 NVIDIA 的高端显卡可以通过NVLink直接通信,Intel 的 CPU目前不支持NVLink,只能使用PCIe技术。同时,单台机器上的多张英伟达显卡也可以使用NVLink相互通信,适合多GPU卡并行计算的场景。
Streaming Multiprocessor
在 NVIDIA 的设计里,一张GPU卡有多个Streaming Multiprocessor(SM),每个 SM 中有多个计算核心,SM 是运算和调度的基本单元。下图为当前计算力最强的显卡Tesla V100,密密麻麻的绿色小格子就是GPU小核心,多个小核心一起组成了一个SM。
在进一步学习 CUDA 编程模型之前,我们首先配置好 CUDA 的运行环境,跑通 Hello World 从而对 CUDA 编程有一个直观的认识,这里使用的是腾讯云的 GPU 服务器,机器安装的是 CentOS 7 系统,CUDA 环境配置可以参考 CUDA Installation Guide Linux 。
根据上图的 NVIDIA GPU 软件栈,有了一个插上了 GPU 的服务器之后,我们首先查看机器上的 GPU,可以看到当前机器上装GPU是 Tesla P40:
$ nvidia-smi Sat Nov 21 17:09:13 2020 +-----------------------------------------------------------------------------+ | NVIDIA-SMI 455.32.00 Driver Version: 455.32.00 CUDA Version: 11.1 | |-------------------------------+----------------------+----------------------+ | GPU Name Persistence-M| Bus-Id Disp.A | Volatile Uncorr. ECC | | Fan Temp Perf Pwr:Usage/Cap| Memory-Usage | GPU-Util Compute M. | | | | MIG M. | |===============================+======================+======================| | 0 Tesla P40 Off | 00000000:00:08.0 Off | 0 | | N/A 27C P0 49W / 250W | 0MiB / 22919MiB | 3% Default | | | | N/A | +-------------------------------+----------------------+----------------------+
+-----------------------------------------------------------------------------+ | Processes: | | GPU GI CI PID Type Process name GPU Memory | | ID ID Usage | |=============================================================================| | No running processes found | +-----------------------------------------------------------------------------+
CUDA 自己提供了一系列的代码示例,可以通过下面的方法安装:
1
$ cuda-install-samples-11.1.sh <dir>
在对应目录下,我们可以看到 CUDA 提供的源代码:
1 2 3
$ ls NVIDIA_CUDA-11.1_Samples 0_Simple 2_Graphics 4_Finance 6_Advanced bin EULA.txt Makefile 1_Utilities 3_Imaging 5_Simulations 7_CUDALibraries common LICENSE
CUDA Device Query (Runtime API) version (CUDART static linking)
Detected 1 CUDA Capable device(s)
Device 0: "Tesla P40" CUDA Driver Version / Runtime Version 11.1 / 11.1 CUDA Capability Major/Minor version number: 6.1 Total amount of global memory: 22919 MBytes (24032378880 bytes) (30) Multiprocessors, (128) CUDA Cores/MP: 3840 CUDA Cores GPU Max Clock rate: 1531 MHz (1.53 GHz) Memory Clock rate: 3615 Mhz Memory Bus Width: 384-bit L2 Cache Size: 3145728 bytes Maximum Texture Dimension Size (x,y,z) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384) Maximum Layered 1D Texture Size, (num) layers 1D=(32768), 2048 layers Maximum Layered 2D Texture Size, (num) layers 2D=(32768, 32768), 2048 layers Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total shared memory per multiprocessor: 98304 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) Maximum memory pitch: 2147483647 bytes Texture alignment: 512 bytes Concurrent copy and kernel execution: Yes with 2 copy engine(s) Run time limit on kernels: No Integrated GPU sharing Host Memory: No Support host page-locked memory mapping: Yes Alignment requirement for Surfaces: Yes Device has ECC support: Enabled Device supports Unified Addressing (UVA): Yes Device supports Managed Memory: Yes Device supports Compute Preemption: Yes Supports Cooperative Kernel Launch: Yes Supports MultiDevice Co-op Kernel Launch: Yes Device PCI Domain ID / Bus ID / location ID: 0 / 0 / 8 Compute Mode: < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >
deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 11.1, CUDA Runtime Version = 11.1, NumDevs = 1 Result = PASS
到现在,CUDA Toolkit 安装完毕,接下来通过编写一个简单的 hello world 来直观感受 CUDA 编程:
// Transfer data from host to device memory cudaMemcpy(d_a, a, sizeof(float) * N, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, sizeof(float) * N, cudaMemcpyHostToDevice);
// Executing kernel vector_add<<<1,1>>>(d_out, d_a, d_b, N); // Transfer data back to host memory cudaMemcpy(out, d_out, sizeof(float) * N, cudaMemcpyDeviceToHost);
// Verification for(int i = 0; i < N; i++){ assert(fabs(out[i] - a[i] - b[i]) < MAX_ERR); }
// Transfer data from host to device memory cudaMemcpy(d_a, a, sizeof(float) * N, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, sizeof(float) * N, cudaMemcpyHostToDevice);
// Executing kernel vector_add<<<1,256>>>(d_out, d_a, d_b, N); // Transfer data back to host memory cudaMemcpy(out, d_out, sizeof(float) * N, cudaMemcpyDeviceToHost);
// Verification for(int i = 0; i < N; i++){ assert(fabs(out[i] - a[i] - b[i]) < MAX_ERR); }
// Transfer data from host to device memory cudaMemcpy(d_a, a, sizeof(float) * N, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, sizeof(float) * N, cudaMemcpyHostToDevice);
// Executing kernel int block_size = 256; int grid_size = ((N + block_size - 1) / block_size); vector_add<<<grid_size,block_size>>>(d_out, d_a, d_b, N); // Transfer data back to host memory cudaMemcpy(out, d_out, sizeof(float) * N, cudaMemcpyDeviceToHost);
// Verification for(int i = 0; i < N; i++){ assert(fabs(out[i] - a[i] - b[i]) < MAX_ERR); }