CUDA编程技术汇总

本文主要罗列了CUDA编程的基本概念、基本原理、基础框架等相关内容，同时结合笔者工作实践，总结了若干CUDA程序的优化技巧。比较适合CUDA初学者。对于想要系统地、深入地研究CUDA的朋友，建议参阅<CUDA C++ Programming Guide>、<CUDA C++ Best Practices Guide>。

限于笔者认知水平与研究深度，难免有不当指出，欢迎批评指正！

注1：文章内容会不定期更新。

注2：文中引用了<CUDA C++ Programming Guide>的图片，但未逐一标注其出处。

一、高性能计算

高性能计算(High performance computing，缩写HPC) 通常指使用很多处理器或者某一集群中组织的几台计算机（作为单个计算资源操作）的计算系统和环境来完成计算任务。

目前，各大超算中心以及国内那几个大厂所搭建的HPC集群基本上都是基于Beowulf架构及其派生架构。虽然国内也有不少公司靠此“发家致富”，但是具有软硬件核心技术的公司其实比较少。造成这种现象的一个原因，就是很多公司实施“暴力式基础建设”，但在关键的、核心的软硬件研发投入、产业化推广等方面，还存在很多问题。

二、软件层面

2.1 CUDA

CUDA是一个并行计算开发环境，主要针对基于NVIDIA GPU硬件。

除了CUDA，还有OpenCL、OpenACC、DirectCompute等类似软件开发平台。

2.2 编程模型

将系统分成主机端与设备端；设备端划分为粗细两层的并行结构；kernel函数由主机端调用，但在设备端运行。

核函数：由__global__修饰;返回值类型为void；主机端调用需要"<<<gridDim,blockDim>>>"指定执行配置。

2.3 线程模型

核函数的全部线程块构成一个网格(Grid)，线程网格由多个线程块（Block）组成；每个线程块包含多个线程。在软件层面，kernel以线程块为单位并发执行。

CUDA 两级线程模型主要是为了使同一CUDA程序可以在不同代次的GPU上自动伸缩运行。

同时，这种线程模型也提供了粗细两种粒度的并行模式（至少软件层面上，可以这样理解），支持线程块内线程同步与数据通信。

kernel函数中内建变量：a) uint3类型 blockIdx、threadIdx; b）dim3类型gridDim、blockDim

线程配置限制因素：a) gridDim.x、gridDim.y、gridDim.z大小限制; b）blockDim.x、blockDim.y、blockDim.z大小限制; c) blockDim.x*blockDim.y*blockDim.z乘积大小限制；d)线程块内线程总数通常要设置为32的倍数，可取64、128、256; e) 线程网格的维度通常由问题规模决定。

2.4 内存模型

CUDA为了降低访存延迟，进一步提高计算效率，采用了一种由多种类型内存构成的分级内存模型。这些内存主要包括寄存器、共享内存、L1缓存、L2缓存、全局内存、常量内存、纹理内存、表面内存等。

每一个线程拥有自己的私有存储器寄存器和局部存储器；每一个线程块拥有一块共享存储器（shared_memory）；最后，grid中所有的线程都可以访问同一块全局存储器（global memory）。

除此之外，还有两种可以被所有线程访问的只读存储器：常量存储器（constant memory）和纹理存储器（Texture memory），他们分别为不同的应用进行了优化。

全局存储器、常数存储器和纹理存储器中的值在一个内核函数执行完成后将被继续保持，可以被同一程序中的其他内核函数调用。

存储器	位置	拥有缓存	访问权限	变量生命周期	带宽	延迟（cycles）
register	GPU片内	N/A	device可读/写	与thread相同	~8TB/s	1
local_memory	板载显存	无	device可读/写	与thread相同
shared_memory	GPU片内	N/A	device可读/写	与block相同	~1.5TB/s	1 - 32
constant_memory	板载显存	有	device可读,host可读/写	可在程序中保持	~200GB/s	~400-600
texture_memory	板载显存	有	device可读,host可读/写	可在程序中保持	~200GB/s	~400-600
global_memory	板载显存	无	device可读/写,host可读/写	可在程序中保持	~200GB/s	~400-600
host_memory	host内存	无	host可读/写	可在程序中保持
pinned_memory	host内存	无	host可读/写	可在程序中保持

在计算能力2.x之前的设备，全局内存的访问会在L1\L2 cache上缓存；在计算能力3.x以上的设备，全局内存的访问只在L2 cache上缓存。对于L1 cache，每次按照128字节进行缓存；对于L2 cache，每次按照32字节进行缓存。

3.1 GPU硬件架构
三、硬件层面

GPU是显卡上的核心组件，显卡一般插在主板PCI-E插槽上，经北桥与CPU相连。

GPU可以简单看成由多个流多处理器(Streaming Multiprocessor, SM)构成的阵列，每个SM具有多个流处理器(Streaming Processor,SP)。

3.2 Warp

在硬件层面，线程块内的线程按照线程号划分成不同的warp，kernel以warp为单位执行。线程束(Warp)是同一线程块中相邻的warpSize个线程。对于目前所有的GPU架构warpSize是32。

四、优化策略

4.1 线程块的数量应该是SM数量的整数倍，总线程数要大于GPU计算核心数，这样可以有效地重叠计算与访存，使GPU 计算核心保持忙碌，从而最大程度的发挥GPU的计算性能。

4.2 减少CPU<->GPU的数据传输；可以的话，考虑合并kernel函数。

4.3 访问一次全局内存，将耗费400~600个cycle，成本是非常高的，所以需要考虑提升全局内存访问效率的方法。

4.3.1 当硬件检测到同一个warp中的线程访问全局存储器中连续的存储单元时，便会将这些访存合并成一个访存。合并访问可以提高DRAM的带宽利用率，使DRAM在传输数据时的速度接近全局存储器的峰值。

参考书籍

樊哲勇. CUDA编程：基础与实践. 清华大学出版社, 2020.

张舒. GPU高性能计算之CUDA. 中国水利水电出版社, 2009.

苏统华. CUDA并行程序设计：GPU编程指南. 机械出版社, 2014.

网络资源

Linux高性能计算集群 -- Beowulf集群http://www.jointforce.com.cn/page/hardware_linux.html

CUDA C++ Programming Guidehttps://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

CUDA C++ Best Practices Guidehttps://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html

周斌：NVIDIA CUDA初级教程视频https://www.bilibili.com/video/BV1kx411m7Fk?from=search&seid=11175574756932230442&spm_id_from=333.337.0.0

苏统华：CUDA并行程序设计https://www.bilibili.com/video/BV15E411x7yT?from=search&seid=11175574756932230442&spm_id_from=333.337.0.0