【IT168 文档】上一个专栏讨论了执行模型和内核启动执行配置如何影响寄存器数量以及本地多处理器资源(比如共享内存, share memory)。现在我们继续讨论内存的性能以及共享内存在reverseArray_multiblock_fast.cu中的使用。

　　CUDA存储器性能

　　局部内存空间和全局内存空间不会缓存，这意味着每次对全局内存(或局部内存)进行访问都将导致一次实际的内存访问。那么访问(例如读取或写入)各种类型的内存的开销是多少?

　　多处理器每次交换(定义见下文)时需要4个时钟周期才能发出一条存储器指令。访问局部内存空间或全局内存将导致400到600个时钟周期的延迟。例如，以下代码片段中的赋值运算符需要4个时钟周期才能从全局内存中进行一次读取，需要4个时钟周期从共享内存进行一次写入，需要400到600个时钟周期从全局内存读取一个浮点值。注意：使用__device__变量类型限定符表示全局内存中的变量(有关其他变量特征，请参见CUDA编程指南第4.2.2.1节)。主代码不能访问变量类型__device__。

　　__shared__ float shared[32];

　　__device__ float device[32];

　　shared[threadIdx.x] = device[threadIdx.x];

　　由于访问时间相差100-150倍，因此开发人员必须最小化对全局内存的访问并在局部多处理器存储器中重用数据。CUDA设计者对线程调度程序的设计十分巧妙，大量的全局内存延迟都可以透明地隐藏起来：只需在执行配置中指定大量数据块，并尽可能在内核中使用寄存器、__shared__和__constant__存储器类型处理变量即可。

　　共享内存位于芯片上，因此访问速度明显快于全局内存，最主要的优化在于避免了存储器组冲突。共享内存速度很快(有些文档指出它与寄存器访问一样快)，但是，最近在CUBLAS和CUFFT性能方面取得的巨大改进就是因为使用寄存器替代了共享内存，所以应该尽可能使用寄存器。CUDA共享内存分为大小相等的存储器模块，这些模块称为存储器组(memory bank)。每个存储器组都保存一个连续的32位值(比如int和float)，因此连续线程进行的连续数组访问非常快。向同一个存储器组(可能是同一个地址，或者映射到同一个存储器组的多个地址)进行多个数据请求时将发生存储器组冲突。发生冲突时，硬件将有效地序列化存储器运算，强迫所有线程等待，直到完成了所有存储器请求。如果所有线程从同一个存储器地址读取，则将自动调用广播机制，不会进行序列化。共享内存广播是一个能够同时向多个线程提供数据的高效方式。使用共享内存时，这项功能很值得注意。

　　在以后的专栏中我们将详细讨论存储器组冲突。目前我们只需知道reverseArray_multiblock_fast.cu不存在存储器组冲突，因为连续线程访问连续值。

　　具有读取/写入功能的多处理器(multi-processor)本地存储器类型概述如下：

　　• 寄存器：

　　o 多处理器上最快的内存形式。

　　o 只有线程能够访问。

　　o 拥有线程的生命周期。

　　• 共享内存：

　　o 在没有存储器组冲突(从同一个地址读取)时与寄存器一样快。

　　o 数据块创建的任何线程都可以访问。

　　o 拥有线程块的生命周期。

　　• 全局内存：

　　o 可能比寄存器或共享内存慢150倍，注意非联合读取和写入(将在下一专栏中讨论)。

　　o 可从主机或设备访问。

　　o 拥有应用程序生命周期。

　　• 局部内存：

　　o 潜在的性能缺陷，位于全局内存中，可能比寄存器或共享内存慢150倍。

　　o 只有线程能够访问。

　　o 拥有线程生命周期。

　　共享内存注意事项

　　1. 当心共享内存组冲突，这可能会降低性能。

　　2. 所有在内核中动态分配的共享变量都从相同的地址开始。使用多个动态分配的共享内存数组需要手动生成偏移量。例如，如果想动态分配共享内存以包含两个数组a和b，请执行如下操作：