常量内存
用于保存在核函数执行期间不会发生变化的数据。某些情况用常量内存替换全局内存能有效减少带宽。
修饰符:__constant__
需要使用cudaMemcpyToSymbol()将数据复制到常量内存
常量内存中读取数据可以节省带宽的两个原因:
1.对常量内存的单独读操作会将数据广播到其他的“邻近”线程,会节约15次读操作。
2.常量内存的数据将缓存起来,对相同地址的数据进行连续读操作不会产生额外内存通信量。
“邻近”是什么意思?我们需要先了解以下线程束。
线程束
它是一个线程集合,程序的每一行,线程束中的每个线程都在不同数据上执行相同的指令。处理常量内存时,NVIDIA硬件将单次读写内存操作广播到每个半线程束,如果半线程束中每个线程都要读取常量内存中相同地址的数据,那么GPU只会产生一次读取请求,将数据广播到半线程束的每个线程,这样产生的内存流量只是使用全局内存的1/16(约6%)。实际使用中会减少更多带宽,因为硬件会主动把常量内存数据缓存在GPU上,这样当其他半线程束需要数据时,将命中缓存。同时,也可能会带来负面影响。如果半线程束上的每个线程需要访问常量内存上不同的数据,那么读取操作将被串行化,但在全局内存中读取,这些请求会被同时发出。那么怎么可以知道,产生的影响是正面的还是负面的?我们可以使用GPU的事件API。
事件
cuda中的事件本质是一个GPU时间戳。获得一个时间戳需要两个步骤:
1.创建一个事件。cudaEventCreate()
2.记录一个事件。cudaEventRecord()
要统计代码执行时间,我们要创建开始事件和结束事件。但是存在一个问题,对于一些异步函数,GPU执行完之前,CPU会执行程序中的下一行代码,这将使计时工作更加复杂。采用的解决办法就是一行代码:cudaEventSynchronize(),当该函数返回结果时,可以安全的读取结束的时间戳。
由于cuda事件是在GPU上执行的,因此它不适用于在CPU和GPU上混合执行的代码。