MindSpore：CUDA编程（六）存储单元

2022/8/12 1:25:25

编程Tag： int size block CUDA 100 __ 1000 存储单元 MindSpore

本文主要是介绍MindSpore：CUDA编程（六）存储单元，对大家解决编程问题具有一定的参考价值，需要的程序猿们随着小编来一起学习吧！

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); 其余不变。

编译，执行：