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

编译，执行：

修改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