CUDA的存储单元包含以下类型:
如下表所示:
名称 | 位置 | 用途 | 使用方法 | 限制 | 备注 |
---|---|---|---|---|---|
Register寄存器 | GPU的SM上 | 存储局部变量 | 每个SM上有成千上万个 一个线程最大数量为256个 需要省着用 |
线程私有,最快 线程退出则失效 |
|
Shared memory | GPU芯片上 | 实现Block内的线程通信,目前最快的多Thread沟通的地方 | __shared__修饰符 需要__syncThreads()同步 |
分为32个banks 需要省着用,会影响活动warp数量 |
可被1个block所有thread访问,次快 高带宽,低延迟 |
Local memory | 存放单线程的大型数组和变量(Register不够时用它) | 没有特定的存储单元 | 线程私有,速度较慢,速度与Global memory接近 | ||
Constant memory 常量内存 |
驻留在device memory中 | 用于同一warp的所有thread同时访问同样的常量数据,比如光线追踪 | __constant__修饰符 必须在host端使用 cudaMemcpyToSymbol初始化 |
没有特定的存储单元,但是有单独的缓存 | 只读,全局 |
Global memory | 等同于GPU显存 驻留在device memory中 |
输入数据,写入结果 | 全局,速度较慢 | ||
Texture memory 纹理内存 |
用于加速局部性访问,比如热传导模型 | 只读,全局,速度次于Shared Memory(延迟比Shared Memory高,带宽比hared Memory小) | |||
Host memory: 可分页内存 |
主机端内存 | 使用malloc访问使用free释放 | 不可以使用DMA访问 | 内存页可以置换到磁盘中 | |
另一种Host memory: 又称: Page-locked Memory,Zero-Copy Memory |
主机端内存 | 使用cudaMallocHost访问 使用cudaFreeHost释放 |
属于另一种Global memory | ||
如何使用Shared Memory优化CUDA应用呢?
Shared Memory的特点是快的时候特别快,慢的时候特别慢。
什么时候快?
同一warp中所有线程访问不同的banks
或者 同一warp中所有线程读取同一地址(通过广播)
什么时候慢?
同一warp中多个线程访问同一个bank的不同地址(此时将产生 bank conflict)
串行访问
请注意:bank conflict发生的原因就是 warp的分配和bank的分配重叠了:
如何避免bank conflict,简单的方法是Padding法(好像叫做补边):
通过增加一个空列,让bank强行错位,使得每段连续的数据被分配到不同的bank中。
具体做法很简单:
就是在设置Shared Memory的时候,不设置成 方阵BLOCK_SIZE X BLOCK_SIZE,而设置成 BLOCK_SIZE X (BLOCK_SIZE+1).
最后,我们可以使用Shared Memory优化mXn, nXk的矩阵乘 的代码,提高访存的效率。
具体方法如下:
申请两块 Shared Memory,都是BLOCK_SIZE X BLOCK_SIZE 大小。一个沿着矩阵mXn滑动,一个沿着矩阵 nXk滑动。将 子集的结果累加到 目的矩阵中:
具体的代码如下:
__global__ void gpu_matrix_mult_shared(int *d_a, int *d_b, int *d_result, int m, int n, int k) { __shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE]; __shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE]; int row = blockIdx.y * BLOCK_SIZE + threadIdx.y; int col = blockIdx.x * BLOCK_SIZE + threadIdx.x; int tmp = 0; int idx; for (int sub = 0; sub < gridDim.x; ++sub) { idx = row * n + sub * BLOCK_SIZE + threadIdx.x; tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? d_a[idx]:0; idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col; tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? d_b[idx]:0; __syncthreads(); for (int k = 0; k < BLOCK_SIZE; ++k) { tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x]; } __syncthreads(); } if(row < n && col < n) { d_result[row * n + col] = tmp; } }
并将前面 代码中调用矩阵乘的地方:gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); 改为 gpu_matrix_mult_shared<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); 其余不变。
编译,执行:
修改blocksize,将其分别改为 16,8,4,再进行统计汇总:
矩阵MXN(m) | 矩阵NXK(n) | 矩阵NXK(k) | blocksize | stop-start(ms) |
---|---|---|---|---|
100 | 100 | 100 | 32 | 1.83286 |
100 | 100 | 100 | 16 | 1.27365 |
100 | 100 | 100 | 8 | 1.23292 |
100 | 100 | 100 | 4 | 3.52865 |
100 | 100 | 100 | 6(补测) | 2.1999 |
100 | 100 | 100 | 12(补测) | 1.34755 |
从上面的结果来看,blocksize为8,16,32时好像差异不大,但是blocksize为4的时候速度降得比较厉害。
blocksize为4时,其实并没有发生bank conflict!而只是因为4X4,只有16个线程,而一个warp需要32个线程,所以相当于计算时,有一半算力被浪费掉了,进而速度慢了一倍。
专家建议,至少应该NXN>32比较好。
将 矩阵从100改为1000试试,但是发现一旦改为1000后,CPU计算可能算不过来了,需要将CPU那部分代码和后面比较的代码屏蔽掉。
再重新统计:
矩阵MXN(m) | 矩阵NXK(n) | 矩阵NXK(k) | blocksize | stop-start(ms) |
---|---|---|---|---|
1000 | 1000 | 1000 | 32 | 265.106 |
1000 | 1000 | 1000 | 16 | 228.09 |
1000 | 1000 | 1000 | 8 | 202.382 |
1000 | 1000 | 1000 | 4 | 518.315 |
1000 | 1000 | 1000 | 6(补测) | 386.171 |
1000 | 1000 | 1000 | 12(补测) | 246.29 |