CUDA进阶：内存模型

存储单元读写速度：Register file > Shared Memory > Constant Memory > Texture Memory > Local and Global Memory

寄存器

属于每个处理单元的原生存储单元，最稀有也最快速。

共享内存 Shared memory

Shared Memory是目前最快的可以让多个线程沟通的地方，每个SM拥有一块Shared memory，一个Block共享一个shared memory。共享内存本质上是SM中的一级缓存，每个SM中的一级缓存与共享内存共享一个内存段。在计算能力2.x以上的设备上用户可以配置，在每次内核调用时可以设置它，共享内存与一级缓存的大小配置方法如下：

cudaFuncSetCacheConfig(cache_prefer,kernel_name);

• cache_prefer参数可取
  • cudaFuncCachePreferShared：表示使用48KB共享内存和16KB一级缓存
  • cudaFuncCachePreferL1：表示使用48KB一级缓存和16KB共享内存

为了克服这个多线程同时访问Shared Memory上的数据的瓶颈，Shared Memory被分成32个逻辑块（banks），在更早的架构中是16个banks。

每32-bit 访存(费米架构) 为一组(开普勒架构引入了64-bit的存储体)，被分配到不同的连续的 banks。 每个 bank 每个周期可以响应一个地址。

1. 同常量内存一样，当一个 warp 中的所有线程访问同一 地址的共享内存时，会触发一个广播（broadcast）机制到 warp 中所有线程，这是最高效的。
2. 如果同一个 half-warp/warp 中的线程访问同一个 bank 中的不同地址时将发生 bank conflict。
3. 即使同一个 warp 中的线程 随机的访问不同的 bank，只要没有访问同一个 bank 的不同地址就不会发生 bank conflict。

不定大小共享内存声明

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

如果存在多于一个extern __shared__内存声明,它们互为别名,所以声明:

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

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

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

束同步编码

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

只要保证连续的线程读取连续的shared memory就不会出现bank conflict

shared memory读写 块内同步：__syncthreads；

块间同步：__threadfence();

常量内存

常量内存其实只是全局内存的一种虚拟地址形式,并没有特殊保留的常量内存块。常量内存有两个特性,一个是配有高速缓存,另一个是它支持将单个值广播到线程束中的每个线程。
常量内存,通过名字我们就能猜到它是只读内存。这种类型的内存要么是在编译时声明为只读内存,要么是在运行时通过主机端定义为只读内存。當量只是从 GPU 内存的角度而言。常量内存的大小被限制为64K。
在编译时声明一块常量内存,需要用到__ constant __关键字,例如:

__constant__ float myarray[1024]={0.0F,1.0F,...};

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

常量内存会将部分数据缓存到一块一级缓存上，使随后的访问变得更快。

常量内存的广播机制

常量内存将数据分配或广播到线程束的每个线程中，广播能在单个周期完成。

在基于二级缓存访问机制的费米结构设备上，同样可以使用广播机制。

线程束的多个线程从常量内存读取同一个元素会触发广播机制，在读取到数据后广播到线程束中的每一个线程，这在利用多线程执行一些常见的变换时特别有用。

运行时常量内存更新

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

__host__ cudaError_t cudaMemcpyToSymbol (const void* s