【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，请执行如下操作：
__global__ voidkernel(intaSize)
{floatsData[];
float*a, *b;=
寄存器/局部内存注意事项
1. 寄存器存储器可以透明地放入局部内存中。这可能导致性能低下。检查ptx汇编码或通过nvcc在输出中查找lmem，命令如下：“—ptxas-options=-v”。
2. 编译时已知使用常量索引的数组通常位于寄存器中，但是如果使用变量索引，则不能位于寄存器中。这给开发人员出了一个难题，因为可能需要循环展开才能在寄存器存储器(不是较慢的全局存储器)中保存数组元素。但是，展开循环可能使寄存器的使用量大大上升，这将导致在局部内存中保存变量——抵消了循环展开的诸多好处。可以使用nvcc选项“—maxrregcount=value”告诉编译器使用更多寄存器(注意：可以指定的最大寄存器数量为128)。这需要在“使用更多的寄存器”和“创建更少的线程”之间权衡利弊，有可能会妨碍隐藏存储器延迟。在某些架构中，使用该选项可能造成资源不足，从而导致内核无法启动。
共享内存内核
程序reverseArray_multiblock.cu和revereseArray_multiblock_fast.cu执行的任务相同。它们都创建一个一维整数数组h_a，数组包含整数值[0 .. dimA-1]。数组可以通过cudaMemcpy移动到设备，然后主机启动reverseArrayBlock内核就地逆转数组内容的顺序。再次使用cudaMemcpy将数据从设备传回主机，并执行一个检查，检查设备生成了正确的结果(例如[dimA-1 .. 0])。
不同之处在于reverseArray_multiblock_fast.cu使用共享内存改进了内核的性能，而reverseArray_multiblock.cu完全在全局内存中执行。您可以尝试为这两个程序计时并验证性能的差异。reverseArray_multiblock.cu也能以低效的方式访问全局内存。在将来的专栏中，我们将介绍如何使用CUDA profiler诊断和修复该性能问题，还将展示最新的10个系列架构中的改进如何减少了许多情况下对这些优化的需求。
在此处插入源代码列表reverseArray_multiblock_fast.cu 决定运行时共享内存的容量需要在主机代码和设备代码进行一些设置。在本例中，内核每个数据块中共享内存的容量(以字节为单位)都在主机的执行配置中指定(使用可选的第三个参数)。(仅当共享内存的容量在内核启动时指定后，才设置主机端。如果在编译时修复，则主机端不需要任何设置。)例如，在主机代码arrayReversal_multiblock_fast.cu中，以下代码片段为一个整数数组(包含的元素数等于数据块中的线程数)分配共享内存：
//Part 1 of 2: Compute the number of bytes of share memory needed//This is used in the kernel invocation below
intsharedMemSize =numThreadsPerBlock *sizeof(int);
查看reverseArrayBlock内核，共享内存使用以下代码声明：
extern __shared__ ints_data[];
注意，内核中没有指定大小，大小通过执行配置从主机获取。
在下一个关于配置的专栏前，我建议研究一下reverseArray_multiblock.cu。您认为访问全局内存时存在性能问题吗?如果您认为存在，请尝试修复它。
其他资源
· CUDA编程指南：位于CUDA Zone的文档部分 http://www.nvidia.com/cuda
