CUDA 学习（十二）、常量内存

一、概述

       常量内存其实只是全局内存的一种虚拟地址形式，并没有特殊保留的常量内存块。常量内存有两个特征，一个是高速缓存，另一个是它支持将单个值广播到线程束中的每个线程。

       常量内存，通过名字我们就能猜到它是只读内存。这种类型的内存要么是在编译时声明为只读内存，要么是在运行时通过主机端定义为只读内存。常量只是从GPU内存的角度而言。常量内存大小被限制为64K。

      在编译时声明一块常量内存，需要用到__constant__ 关键字。

      如果要在运行时改变常量内存区中的内容，只需在调用GPU内核之前简单地调用cudaCopyToSymbol 函数。如果在编译阶段或主机端运行阶段没有定义常量内存，那么常量内存区将未定义。

二、常量内存高速缓存

1、在计算能力为1.x 的设备

     在计算能力为1.x 的设备上，常量内存有一个特性，就是将数据缓存到一块8KB 的一级缓存上，使随后的访问变得更加快速。但前提是内存中的数据可能在程序中重复利用。此外，广播访问也被高度优化，使得访问相同内存地址的线程能够在单个期间内完成。

2、在计算能力为2.x 的设备

     在费米架构的设备以及后续发布的设备上出现了二级缓存。费米架构设备上的二级缓存可以在SM之间进行共享。所有的内存访问将自动地缓存到二级缓存中。对任何声明为常量的内存区进行基于非线程的访问都将通过常量缓存。所谓的基于非线程的访问，即对所访问的数组的索引进行计算时不需要用到 threadIdx.x。

三、常量内存广播机制

      常量内存有一个非常有用的特性，该特性主要将数据分配或广播到线程束的每个线程中。广播能在单个周期内发生，因此，该特性非常有用。

      在基于二级缓存访问机制的费米设备上，同样可以使用广播机制，通过广播机制，我们可以快速的将数据分配到线程束的多个线程中。每个线程从常量内存读取元素N，这样会触发广播机制，在读到数据之后广播到线程束中的每个线程。由于常量内存区几乎可以到达与一级缓存相同的速度，因此这类算法使用常量内存性能非常好。

      然而，需要特别注意的是，如果一个常量只是字面值，那么最好用 #define 对字面值进行定义，因为这样可以减少常量内存的使用。

四、运行时进行常量内存更新

       GPU 上的常量内存并不是真正意义上的常量内存，因为GPU上并没有专门常量内存预留特殊内存区。由于常量内存是通过16位地址进行访问的，而16位地址能够快速进行访问，因此常量内存最大限制为64KB。这样做会带来一定好处但也会带来一些问题。首先，通过调研cudaMemcpyToSymbol 函数，常量内存可以按块或片的形式进行更新，一次最多能更新64K。

      注意cudaMemcpyToSymbol 函数的工作原理。该函数可以将数据复制到GPU上任何以全局符合命名的内存区域，无论该符号是全局内存还是常量内存。因此，我们可以将一块64K大小的数据复制到一个64K 大小内存区上，从而通过常量内存缓存进行访问。当所有线程访问同一数据元素时，这种访问方式非常有用，因为我们可以借助缓存技术从常量内存区获取然后在广播到每个线程中。