【CUDA 基础】5.4 合并的全局内存访问
title: 【CUDA 基础】5.4 合并的全局内存访问
categories:
- CUDA
- Freshman
tags:
- 合并
- 转置
toc: true
date: 2018-06-04 21:34:22
Abstract: 本文介绍使用共享内存进行矩阵转置以减少内存的交叉访问
Keywords: 合并，转置
开篇废话
没废话，看以前的废话感觉自己像个傻瓜。。就像以后看我正在写的文字一样。
还记得我们矩阵转置的例子么，在全局内存部分介绍的：4.4核函数可达到的带宽
在4.4中我们当时只有共享内存这一种工具可以使用，为了达到最高效率，我们要配合一级缓存，二级缓存进行编程，来提高转置的效率，因为转置只能在行读取列写入或者列读取行写入之间选择一个，这样就必然会引发非合并的访问，虽然我们利用一级缓存的性质可以提高性能，但是我们今天会介绍我们的新工具共享内存，在共享内存中完成转置后写入全局内存，这样就可以避免交叉访问了。
基准转置内核
在介绍我们的神奇共享内存之前，我们最好先研究出来一下我们的问题的极限在哪，换句话说，我们需要清楚的知道我们最慢的情况（最简单的方式能达到的速度）以及最快的理论速度，理论速度可能会达不到，但是可以接近，最慢速度肯定可以超越，你永远可以写出更慢的程序，所以我们用最简单的方法作为下界，而用正行读取，然后不经变换的写入来作为上限，这一招我们在前面使用过，就是在4.4中，那次我们突破极限了（哈哈，很有可能是计时有问题），但是正常来讲，极限是最好的参考值。
完整的代码在github:https://github.com/Tony-Tan/CUDA_Freshman（欢迎随手star? ）
上限：
__global__ void copyRow(float * in,float * out,int nx,int ny)
{
int ix=threadIdx.x+blockDim.x*blockIdx.x;
int iy=threadIdx.y+blockDim.y*blockIdx.y;
int idx=ix+iy*nx;
if (ix<nx && iy<ny)
{
out[idx]=in[idx];
}
}
下限是我们的too young too naive版本，就是最常规的方法：
__global__ void transformNaiveRow(float * in,float * out,int nx,int ny)
{
int ix=threadIdx.x+blockDim.x*blockIdx.x;
int iy=threadIdx.y+blockDim.y*blockIdx.y;
int idx_row=ix+iy*nx;
int idx_col=ix*ny+iy;
if (ix<nx && iy<ny)
{
out[idx_col]=in[idx_row];
}
}
这两段代码中第一段并没有转置的功能，只是为了测试上限，第二段是naive的转置，前面也讲过，这里就直接贴结果了
copyRow的cpu计时和nvprof结果：
transformNaiveRow的cpu计时和nvprof结果：
我们可以得到下表：
| 核函数 | CPU计时 | nvprof计时 |
|---|---|---|
| copyRow | 0.001442 s | 1.4859 ms |
| transformNaiveRow | 0.003964 s | 3.9640 ms |
可以看出计cpu计时还是比较准的，在数据量比较大情况下，我们现在的矩阵大小是 212×2122^{12}\times 2^{12}212×212 的大小。
然后是加载和存储全局内存请求的平均事务数（越少越好）
copyRow:
transformNaiveRow:
接着我们就开始用共享内存进行操作了。
使用共享内存的矩阵转置
完整内容 https://face2ai.com/CUDA-F-5-4-合并的全局内存访问/
最新文章
- 【玩转单片机系列001】 08接口双色LED显示屏驱动方式探索
- 找不到mysql.sock，mysql.sock丢失问题解决方法
- 15款免费的 HTML5/CSS3 响应式网页模板
- PHP isset() 检测变量是否设置
- 给Android程序员的六个建议
- swift 基于SDK8.0 获取当前时间
- 选项切换条--第三方开源--SHSegmentControl
- msssql 用numberic(38)替代int去解决int不够的问题
- 我的第一个BAE python应用
- qt之treeview例子
- js使用for in遍历时的细节问题
- react——一个todolist的demo
- Web程序-----批量生成二维码并形成一张图片
- day46 前端基础HTML5+CSS3
- bzoj4361 isn (dp+树状数组+容斥)
- IOS初级：NSTimer
- c# 跨域api
- 华为5G折叠屏幕适配
- DSP 程序的执行时间
- c语言for循环等语句详解