假设我有一个包含多个固定数值的数组，这些数值将被同一块内的多个线程多次访问，如果我将这些值存储在以下位置，那么在访问速度和内存使用方面有哪些优缺点：
线程本地内存：double x[3] = {1,2,3};
共享内存：__shared__ double x[3] = {1,2,3};
数字文字：直接在它们出现的表达式中硬编码这些值
谢谢！
【问题讨论】：
__constant__内存缓存。
TL;DR
使用__constant__ double x[3]; // ... initialization ...
在你的问题中：
- 线程局部内存：double x[3] = {1,2,3};
这是不精确的。取决于您的代码如何访问 x[]、x[] 可以驻留在寄存器或本地内存中。
由于没有类型限定符，编译器会尽量把东西放在寄存器中，
在设备代码中声明的没有本节中描述的任何 __device__、__shared__ 和 __constant__ 限定符的自动变量通常驻留在寄存器中。但是在某些情况下，编译器可能会选择将其放在本地内存中，
但是当它不能时，它会将它们放入本地内存：
无法确定是否使用常数索引的数组，
会占用过多寄存器空间的大型结构或数组，
内核使用的寄存器多于可用寄存器时的任何变量（这也称为寄存器溢出）。
您真的不希望x 位于本地内存，它很慢。在你的情况下，
具有多个固定数值的数组，将被同一块内的多个线程多次访问
__constant__ 和 __shared__ 都是不错的选择。
有关此主题的完整说明，请查看：CUDA Toolkit Documentation: variable-type-qualifiers
硬编码
该编号将嵌入说明中。您可能会期待一些性能改进。在执行此操作之前和之后更好地对您的程序进行基准测试。
注册
它很快，但很少。考虑一个有 16x16 个线程的块，每个块最多有 64k 个寄存器，每个线程可以使用 256 个寄存器。 （嗯，也许不是那么稀缺，对于大多数内核来说应该足够了）
本地内存
很慢。但是，一个线程最多可以使用 512KB 本地内存。
本地内存空间驻留在设备内存中，因此本地内存访问与全局内存访问具有相同的高延迟和低带宽...
共享内存
它很快，但很少。通常每个块 48KB（小于寄存器！）。
因为它是片上的，所以共享内存比本地或全局内存具有更高的带宽和更低的延迟。
恒定内存
它以不同的方式快速（见下文），它高度依赖缓存，而缓存是稀缺的。每个多处理器通常有 8KB ~ 10KB 缓存。
常量内存空间驻留在设备内存中，缓存在Compute Capability 2.x中提到的常量缓存中。
然后，一个请求会被拆分为与初始请求中不同的内存地址一样多的单独请求，从而将吞吐量降低与单独请求数相等的因子。
然后在缓存命中的情况下以常量缓存的吞吐量为结果请求提供服务，否则以设备内存的吞吐量提供服务。
【讨论】：