GPU编程自学7 —— 常量内存与事件

深度学习的兴起，使得多线程以及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只会产生一次读取请求并在随后将数据