5.6 共享内存

共享内存被使用来在同一线程块内的CUDA线程间交换数据。物理上，它属于SM可见的内存，可以被非常快地访问。在速度上，共享内存的访问速度比寄存器访问速度慢十倍，但比全局内存访问快十倍。因此，共享内存通常作为一个减少CUDA内核所需要外部带宽的重要资源。

由于开发者显式分配与引用共享内存，共享内存可以被视作“手动管理缓存”或“暂存器”内存。开发者可以在内核和设备两个级别上请求不同的缓存配置：CUDADevice

SetCacheConfig()/cuCtxSetCacheConfig()为设备指定优先缓存配置，而CUDAFuncSetCacheConfig()/cuFuncSetCacheConfig()则为给定的内核指定优先缓存配置。如果两项同时指定，内核级别的请求会取得优先级，但是，在任何的情况下，内核的需求可能会重写开发者的优先请求。

使用共享内存的内核由以下过程执行写入。

·加载共享内存和__syncthreads()
· 处理共享内存和__syncthreads()

·写入结果

使用nvcc选项-Xptxas-v,abi=no，开发者可以让编译器报告给定内核使用的共享内存数量。在运行时中，内核使用的共享内存数量可以使用函数cuFuncGetAttribute(CU FUNC_ATTRIBUTE_SHARED_SIZE_BYTES)查询。

5.6.1 不定大小共享内存声明

内核声明的每一个共享内存，都会在内核启动时为线程块自动分配。如果内核包含了一个未确定大小的共享内存声明，在内核启动时，该声明所需的内存数量必须被指定。

如果存在多于一个external_shared__内存声明，它们互为别名，所以声明：

extern __shared char sharedChars[]; extern __shared int sharedInts[];

使相同的共享内存根据需要以8或32位整数寻址。一个使用这种别名的动机是，当可能读写全局内存时可使用更宽的类型，而使用更窄的类型进行内核计算。

注意 如果你有超过一个使用不定大小共享内存的内核，它们必须在独立的文件中编译。

5.6.2 束同步编码

在束同步编程中使用的共享内存变量必须声明为volatile，来保护程序代码不被编译器优化，从而避免代码错误。

5.6.3 共享内存的指针

使用指针引用共享内存是有效的，并且这很方便。使用这一方法的内核样例包括第12章的归约内核（见代码清单12-3）和第13章中的scanBlock内核（见代码清单13-3）。