要实现CUDA高性能编程，就必须对GPU内存结构有深刻的了解。
全局内存(off-chip)
就是我们常说的显存，其容量最大、带宽最小、延迟最高。
常量内存(off-chip)
存储在片下存储的设备内存上，但是通过特殊的常量内存缓存进行缓存读取，常量内存为只读内存，只有64KB。由于有缓存，常量内存的访问速度比全局内存高。
使用常量内存的方法是在核函数外面用 __constant__ 定义变量，并用函数 cudaMemcpyToSymbol 将数据从主机端复制到设备的常量内存后供核函数使用。
纹理内存和表面内存(off-chip)
纹理内存和表面内存类似于常量内存，也是一种具有缓存的全局内存，有相同的可见范围和生命周期，而且一般仅可读(表面内存也可写)。不同的是，纹理内存和表面内存容量更大，而且使用方式和常量内存也不一样。
共享内存(on-chip)
共享内存存在于芯片上，具有仅次于寄存器的读写速度，数量也有限。一个使用共享内存的变量可以 __shared__ 修饰符来定义。该变量对block内的所有线程可见。
寄存器(on-chip)
寄存器是一个线程能独立访问的资源，它所在的位置与局部内存不一样，是在片上（on chip）的存储，用来存储当前线程的一些暂存数据。寄存器的速度是访问中最快的，但是它的容量较小。
在核函数中定义的不加任何限定符的变量一般来说就存放于寄存器(register)中。各种内建变量，如 gridDim、blockDim、blockIdx、 threadIdx 及 warpSize 都保存在特殊的寄存器中，以便高效访问。举例如下：
1 | const int n = blockDim.x * blockIdx.x + threadIdx.x; |
n 也是一个寄存器变量，当只能被当前线程访问。
局部内存(off-chip)
局部内存和寄存器几乎一样，核函数中定义的不加任何限定符的变量有可能在寄存器中，也有可能在局部内存中。寄存器中放不下的变量，以及索引值不能在编译时就确定的数组，都有可能放在局部内存中。
虽然局部内存在用法上类似于寄存器，但从硬件来看，局部内存只是全局内存的一部分。所以，局部内存的延迟也很高。每个线程最多能使用高达512KB的局部内存，但使用过多会降低程序的性能。
总结
