【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的前一半或者后一半，这是内存访问包括本期后面将要讨论的合并内存访问的一个重要概念。
从前面的讨论中我们可以得出以下几点有用信息：
多处理器资源如共享内存是有限的、宝贵的。
为一系列支持CUDA的设备配置，有效管理有限的多处理器资源(如共享内存)，已经成为CUDA开发人员日常生活的一部分。
请注意条件操作(例如，if语句)对内核的运行有深刻影响。
CUDA使用率计算工具和nvcc编译器是值得我们学习和使用的重要工具――尤其在探索执行配置时。
在设备中执行的线程可以通过下图展示的内存类型来访问全局内存和芯片内存：
每个多处理器，见上图的块(0, 0)和块(1, 0)，均包含以下四种内存类型：
每线程一组本地寄存器。
一个并行数据缓存或共享内存。被所有线程共享;能实现共享内存空间。
一个只读的常量缓存。被所有线程共享;能加快读取常量内存空间的速度;被实现为设备内存的只读区(在后面的专栏中将介绍常量内存。在那之前，如需更多信息，请参阅CUDA编程指南的第5.1.2.2节)。
一个只读的纹理缓存。被所有处理器共享;能加快读取纹理内存空间的速度;被实现为设备内存的只读区(在后面的专栏中将介绍纹理内存。在那之前，如需更多信息，请参见CUDA编程指南的第5.1.2.3节)。
不要被上图中的多处理器中标有“本地内存”的块所迷惑。“本地内存”指的是每个线程范围内的本地内存。它是一个内存抽象，而不是多处理器的实际硬件组件。实际上，本地内存由编译器在全局内存中分配，与其他全局内存区提供的性能相同。本地内存基本上被编译器用来保存一些程序员视为线程本地的而又由于某种原因不适合保存到速度更快的内存中的信息。通常在内核中声明的自动变量驻留在提供快速访问的寄存器中。在某些情况下，编译器可能选择在本地内存中存放这些变量，其中一种情况就是寄存器变量太多，一个数组包含四个以上元素，某些结构或数组会消耗太多的寄存器空间，另外一种情况就是编译器不能确定一个数组是否通过恒量进行索引。
请谨慎使用，因为本地内存可能会降低性能。检查ptx汇编代码(可以通过编译–ptx或-keep选项获得)可以知道一个变量是否在最初编译阶段被放置在本地内存，因为变量是通过.local助记符来声明的，是通过ld.local和st.local助记符来访问的。如果在最初阶段没有将它放置到本地内存，那么在随后的编译阶段如果发现它可能会消耗目标架构太多的寄存器空间也可能会决定将它放置在那里。
在阅读下一期之前，我建议利用使用率计算工具来深入了解执行模性和内核启动执行配置对寄存器和共享内存的数量有什么影响。
改变共享内存的数量会影响可用寄存器的数量吗?
如何取得活动块数量和可用多处理器资源之间的平衡?
其他资源
