【IT168 文档】CUDA开发人员面临的最重要的性能挑战之一就是如何充分利用本地多处理器内存资源，如共享内存、常量内存，以及寄存器。原因就是我们上一篇文章中讨论的，虽然全局内存可以提供超过60 GB/秒的速度，但这对于只获取使用一次的数据来说，仅相当于15gf/秒――要获得更高的性能则要求能够重用本地数据。CUDA软件和硬件设计师做了一些出色的工作，以隐藏全局内存的延迟和全局内存的带宽限制――但这都是以本地数据重用为前提的。

　　记得我们在第二期中提到内核启动要求执行配置信息中必须指定多少个线程组成一个块以及多少个块结合形成一个网格。要注意一个块里的线程可以通过本地多处理器资源进行互相通信，因为CUDA执行模型指定了一个块只能在一个多处理器上处理。换句话说，写入块共享内存的数据可以被同一个块里的其他线程读取，但不能被不同块的线程读取。具有这些特征的共享内存可以在硬件中非常有效地实现，并转化为CUDA开发人员所需的快速内存访问(下文将具体进行说明)。

　　现在我们有办法让支持CUDA的硬件设计师在价格和CUDA软件开发人员需求之间的取得平衡。作为开发人员，我们希望有大量的本地多处理器资源，例如寄存器和共享内存。这使得我们的工作更轻松，我们的软件更有效率。而另一方面，硬件设计师需要提供价格低廉的硬件，但不幸的是快速的本地多处理器内存是很昂贵的。我们都同意廉价的CUDA硬件是最好的，所以支持CUDA的硬件被设计成针对不同的市场价位具有不同的处理能力。由市场根据处理能力来决定合适的价格。其实这是一个很好的解决办法，因为技术发展得很快――每个新一代的支持CUDA的设备都比上一代强大，并且在跟上一代同样价位的情况下具有更多更高性能的组件。

　　等等!这听起来更像是一个软件问题而不是折衷方法，因为CUDA开发人员需要考虑所有这些不同的硬件配置并且我们都面临着有限的设备资源的挑战。为了帮助解决问题，一些辅助设计工具已经创建用于协助给不同的架构选择“最好”的高效能执行配置。我强烈推荐大家下载和使用CUDA使用率计算工具(请参见其他资源)，这个工具实际上是一个做得很好的电子表格(当传递“--ptxas-options=-v”选项时，nvcc编译器会报告电子表格所需的有关每个内核的信息，例如有多少个寄存器以及本地内存、共享内存和常量内存的使用率)。尽管如此，在论坛和文档中常见的一个建议就是：“尝试一些不同的配置并衡量对性能的影响”。这个很容易做到，因为执行配置是通过变量来指定的。事实上，很多应用程序在安装时可以有效地自动配置自己(例如，确定非常好的的执行配置)。此外，CUDA运行时函数CUDAgetdevicecount()和CUDAgetdeviceproperties()可以用来计算系统中的CUDA设备的数量并获取它们的属性。其中一个利用这些信息的方法就是执行一个表查找来确定性能较好的执行配置或启动一个自动调优工具。

　　CUDA执行模型

　　为了尽可能提高性能，每个硬件多处理器都可以同时积极处理多个块。能处理多少则取决于每个线程有多少个寄存器和某一内核需要每个块的多少共享内存。一个多处理器同时处理的块被称为活动块。具有最少资源需求的内核可以更好地利用(或占用)每个多处理器，因为多处理器的寄存器和共享内存在活动块的所有线程之间分割。使用CUDA使用率计算工具来探索线程以及活动块数量与寄存器以及共享内存数量之间的平衡点。找到合适的组合能大大提高内核的性能。如果每个多处理器可用的寄存器或共享内存不足于处理至少一个块，内核将无法启动(请参见上一期关于CUDAgetlasterror()的讨论以了解如何捕获这些失败)。

　　每个活动块被分割成线程SIMD(单指令多数据)群，称为Warp：每个Warp包含同样数量的线程，称为Warp size，被多处理器以SIMD方式执行。这意味着Warp中的每个线程传递的都是指令库中的相同指令，指导线程执行一些操作或操纵本地和/或全局内存。从硬件角度看，SIMD模型效率高而且经济有效，但从软件角度来看，很遗憾，它会使条件操作序列化(比如，使条件句的两个分支都必须被求值)。请注意条件操作对内核运行有深刻的影响。谨慎使用的话，一般是可以控制的，但也有可能引起某些问题。

　　活动Warp(比如所有活动块的所有Warp)是以时间片分配的：线程调度程序定期从一个Warp切换到另一个Warp以最大限度利用多处理器的计算资源。块之间或块中的Warp之间的执行次序是不确定的，也就是说可以是任何次序。但是，线程可以通过__syncthreads()进行同步。请注意只有执行__syncthreads()后，才能保证写入共享(和全局)内存是可见的。除非变量被声明为易失性(volatile)变量，否则编译器可以被用于优化(例如重新排序或消除)内存的读和写，以提高性能。__syncthreads()允许在一个条件句范围内被调用，但只有当条件的求值方式在整个线程块中都相同时才可以。否则的话，代码执行过程可能会被挂起或产生意想不到的副作用。可喜的是，__syncthreads()的开销很低，因为在没有线程等待任何其他线程的情况下，它仅花费4个时钟周期为一个Warp发送指令。Half-Warp指的是Warp的前一半或者后一半，这是内存访问包括本期后面将要讨论的合并内存访问的一个重要概念。