Memory Hierarchy
GPU Hardware
Memory Abstraction
全局内存 global memory
- 独立于GPU核心的硬件RAM
- GPU绝大多数内存空间都是全局内存
- 全局内存的IO是GPU上最慢的IO形式（除了访问host端内存）
通过cache L2(CC>=3.0)访问，cache line 大小128 bytes ，每个线程操作尽量少的cache line，速度更快
Examples of Global Memory Accesses. Examples of Global Memory Accesses by a Warp, 4-Byte Word per Thread, and Associated Memory Transactions for Compute Capabilities 2.x and Beyond
共享内存 shared memory
- SM（SM = streaming multiprocessor）中的内存空间
- 最大48KB
- 作用域是线程块
静态分配语法
__shared__ float data[1024];
Declared in the kernel function, nothing in host code
动态分配语法
Host:
kernel<<<grid_dim, block_dim, numBytesShMem>>>(args);
Device (in kernel):
extern __shared__ float s[];
多个动态分配的变量 需要额外注意对齐
extern __shared__ int s[];
int *integerData = s; // nI ints
float *floatData = (float*)&integerData[nI]; // nF floats
char *charData = (char*)&floatData[nF]; // nC chars
myKernel<<<gridSize, blockSize, nI*sizeof(int)+nF*sizeof(float)+nC*sizeof(char)>>>(...);
共享内存块冲突
共享内存分成相同大小的内存块，实现高速并行访问，但是当多个线程的请求地址映射到同一个内存块block时，访问是串行的
步幅stride为n时 最大公约数为1，即gcd(n,32)==1 ,访问共享内存可以避免块冲突
Strided Shared Memory Accesses. Examples for devices of compute capability 3.x (in 32-bit mode) or compute capability 5.x and 6.x
Left
Linear addressing with a stride of one 32-bit word (no bank conflict).
Middle
Linear addressing with a stride of two 32-bit words (two-way bank conflict).
Right
Linear addressing with a stride of three 32-bit words (no bank conflict).
Irregular Shared Memory Accesses. Examples for devices of compute capability 3.x, 5.x, or 6.x
Left
Conflict-free access via random permutation.
Middle
Conflict-free access since threads 3, 4, 6, 7, and 9 access the same word within bank 5.
Right
Conflict-free broadcast access (threads access the same word within a bank).
本地内存 local memory
位于堆栈中，不在寄存器中的所有内容
作用域为特定线程
存储在global内存空间中，速度比寄存器慢很多
寄存器溢出 register spilling
内核使用的寄存器比可用的寄存器多,存储到local memory中
L1 cache
- 每个SM都有自己的L1 cache
- 可配置大小16KB/48KB cudaFuncSetCacheConfig
- 2.x Fermi - caches local & global memory
- 3.x 及以上 Kepler, Maxwell - only caches local memory
L2 cache
- 缓存 local and global memory
- 被所有的SM共享
- 大约为1MB
常量内存 constant memory
- 属于全局内存，大小64KB
- 线程请求同一个数据时很快，请求不同的数据时性能下降
- 在运行中不变，所有constant变量的值必须在kernel启动之前从host设置
-
__global__函数参数通过 constant memory穿的到device端， 限定4 KB，即kernel参数通过常量内存传递
__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));
常量缓存 constant cache
- 每个SM上大小8KB,CC>=5.0大小为10KB
- 把一个内存地址广播到所有的warp线程
- 可以加载静态索引数据，通过 “load uniform” (LDU)指令
纹理内存空间 texture memory
类似constant memory，是只读内存，以某种形式访问的时候可以提升性能。原本是用在OpenGL和DirectX渲染管线中的。
有用的特点：
- 不需考虑要聚合coalescing访问的问题
- 通过“CUDA Array”进行缓存的2D或3D空间的数据位置
- 在1D，2D或3D数组上进行快速插值
- 将整数转换为“unitized”浮点数
用例：
- 通过纹理缓存和CUDA数组读取输入数据，以利用空间缓存
- 利用数字纹理功能。
- 与OpenGL和通用计算机图形的交互
纹理缓存 read-only texture cache
CC ≥ 3.5 大多数的 __restrict__ 变量自动加载到纹理缓存中了
通过 __ldg函数强行加载到缓存
// 2D float texture
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
cudaArray* cuArray;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>;
cudaMallocArray(&cuArray, &channelDesc, width, height);
cudaMemcpyToArray(cuArray, 0, 0, h_data, size,
cudaMemcpyHostToDevice);
// Set texture reference parameters
texRef.addressMode[0] = cudaAddressModeWrap;
texRef.addressMode[1] = cudaAddressModeWrap;
texRef.filterMode = cudaFilterModeLinear;
texRef.normalized = true;
// Bind the array to the texture reference
cudaBindTextureToArray(texRef, cuArray, channelDesc);
cudaUnbindTexture (const textureReference *texref);
cudaFreeArray(cuArray);
