ICode9

精准搜索请尝试: 精确搜索
首页 > 编程语言> 文章详细

MindSpore:CUDA编程(六)存储单元

2022-08-11 18:32:20  阅读:196  来源: 互联网

标签:__ MindSpore int CUDA 存储单元 100 SIZE BLOCK 1000


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)blocksizestop-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)blocksizestop-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


标签:__,MindSpore,int,CUDA,存储单元,100,SIZE,BLOCK,1000
来源: https://www.cnblogs.com/skytier/p/16577097.html

本站声明: 1. iCode9 技术分享网(下文简称本站)提供的所有内容,仅供技术学习、探讨和分享;
2. 关于本站的所有留言、评论、转载及引用,纯属内容发起人的个人观点,与本站观点和立场无关;
3. 关于本站的所有言论和文字,纯属内容发起人的个人观点,与本站观点和立场无关;
4. 本站文章均是网友提供,不完全保证技术分享内容的完整性、准确性、时效性、风险性和版权归属;如您发现该文章侵犯了您的权益,可联系我们第一时间进行删除;
5. 本站为非盈利性的个人网站,所有内容不会用来进行牟利,也不会利用任何形式的广告来间接获益,纯粹是为了广大技术爱好者提供技术内容和技术思想的分享性交流网站。

专注分享技术,共同学习,共同进步。侵权联系[81616952@qq.com]

Copyright (C)ICode9.com, All Rights Reserved.

ICode9版权所有