MindSpore(CUDA编程(六)存储单元)
CUDA的存储单元包含以下类型:
文章图片
如下表所示:名称位置用途使用方法限制备注Register寄存器GPU的SM上存储局部变量每个SM上有成千上万个一个线程最大数量为256个需要省着用线程私有,最快线程退出则失效Shared memoryGPU芯片上实现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
【MindSpore(CUDA编程(六)存储单元)】}并将前面 代码中调用矩阵乘的地方:gpu_matrix_mult<<
文章图片
修改blocksize,将其分别改为 16,8,4,再进行统计汇总:矩阵MXN(m)矩阵NXK(n)矩阵NXK(k)blocksizestop-start(ms)100100100321.83286100100100161.2736510010010081.2329210010010043.528651001001006(补测)2.199910010010012(补测)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)10001000100032265.10610001000100016228.091000100010008202.3821000100010004518.3151000100010006(补测)386.17110001000100012(补测)246.29
推荐阅读
- MindSpore:CUDA编程(三)线程层次
- MindSpore:CUDA编程(五)Event
- 编程语言|2021年2月程序员工资最新出炉,涨了!这下踏实了!
- 人工智能|人工智能、机器人、编程啥关系((科普))
- 计算机专业技术|大数据编程技术——期末复习
- 30秒接入一台PLC设备,用Shifu 快速实现工控软件编程
- 编程语言|2020 ,6 种不死的编程语言!
- 关于编程本质那些事
- 编程笔记|微信公众号开发系列教程(四)(监听关注/取消关注事件 消息接收与响应处理)
- 深度学习|深度学习环境搭建(win11+tensorflow2.7+CUDA11.6+cuDNN8.3.2)