深度学习的兴起,使得多线程以及GPU编程逐渐成为算法工程师无法规避的问题。这里主要记录自己的GPU自学历程。
目录
- 《GPU编程自学1 —— 引言》
- 《GPU编程自学2 —— CUDA环境配置》
- 《GPU编程自学3 —— CUDA程序初探》
- 《GPU编程自学4 —— CUDA核函数运行参数》
- 《GPU编程自学5 —— 线程协作》
- 《GPU编程自学6 —— 函数与变量类型限定符》
- 《GPU编程自学7 —— 常量内存与事件》
- 《GPU编程自学8 —— 纹理内存》
- 《GPU编程自学9 —— 原子操作》
- 《GPU编程自学10 —— 流并行》
七、 常量内存与事件
GPU通常包含大量的数学计算单元,因此性能瓶颈往往不在于芯片的数学计算吞吐量,而在于芯片的内存带宽,即有时候输入数据的速率甚至不能维持满负荷的运算。 于是我们需要一些手段来减少内存通信量。 目前的GPU均提供了64KB的常量内存,并且对常量内存采取了不同于全局内存的处理方式。 在某些场景下,使用常量内存来替换全局内存可以有效地提高通信效率。
7.1 常量内存
常量内存具有以下特点:
- 需要由 __constant__ 限定符来声明
- 只读
- 硬件上并没有特殊的常量内存块,常量内存只是只是全局内存的一种虚拟地址形式
- 目前的GPU常量内存大小都只有64K,主要是因为常量内存采用了更快速的16位地址寻址(2^16 = 65536 = 64K)
- 对于数据不太集中或者重用率不高的内存访问,尽量不要使用常量内存,否则甚至会慢于使用全局内存
- 常量内存无需cudaMalloc()来开辟,而是在声明时直接提交一个固定大小,比如 __constant__ float mdata[1000]
- 常量内存的赋值不能再用cudaMemcpy(),而是使用cudaMemcpyToSymbol()
常量内存带来性能提升的原因主要有两个:
- 对常量内存的单次读操作可以广播到其他的“邻近(nearby)”线程,这将节约15次读取操作
- 常量内存的数据将缓存起来,因此对于相同地址的连续操作将不会产生额外的内存通信量。
对于原因1,涉及到 线程束(Warp)的概念。
在CUDA架构中,线程束是指一个包含32个线程的集合,这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。 即线程束中的每个线程都将在不同数据上执行相同的指令。
当处理常量内存时,NVIDIA硬件将把单次内存读取操作广播到每个半线程束(Half-Warp)。在半线程束中包含16个线程,即线程束中线程数量的一半。如果在半线程束中的每个线程从常量内存的相同地址上读取数据,那么GPU只会产生一次读取请求并在随后将数据