CUDA的global内存访问的问题

http://blog.csdn.net/OpenHero/article/details/3520578

关于CUDA的global内存访问的问题,怎么是访问的冲突,怎样才能更好的访问内存,达到更高的速度。下面先看几张图,这些图都是CUDA编程手册上的图,然后分别对这些图做解释,来理解硬件1.0,1.1 以及现在最新的硬件的访问内存的区别。

我们在这里再深入的讲解一下global内存对齐的问题,每次执行一条明命令的时候,都是会按照32个thread为一个warp,一起来执行,但是在执行的时候,又会按照硬件的条件(这里有两个限制条件,一个是内存访问的时钟和执行core的时钟不一样,第二个是为了细粒度的分支的问题)然后就会把16个thread组成的half-warp来一次访问global内存才能让访问内存的性能高一些,这个可以理解;

就像手册上说的那样,如果16个thread(half-warp)访问内存的时候,如果每一个thread访问32bits就是4个字节,那么就可以合并为一个64bytes的访问,手册上这点写得有点让人咋一看不太明白~4bytes(32bits)*16 = 64bytes,就是这么来的,如果每一个thread访问64bits(8个bytes),那么就可以合并为128bytes的访问;这里啦,由于合并访问的最大限制是128bytes,所以最大也按照128bytes一次访问来合并,如果超过,就得多次访问,或者如果没有按照这样的方式对齐访问,也会多次访问;下面是1.2device之前的访问的几个图,这里要把1.2device以前和以后的分开,是因为这里在对齐访问的方式的时候,有不同的策略;先看1.2device以前的能合并为一次访问的情况:

下图是编程手册上的图:

clip_image002

这里的每一个thread都是访问的对应的地址,是对齐的,所以可以合并为一个存储event;

下面这个图是没有对齐访问,就造成了non-coalesced访问的问题,下面可以看图说话:

clip_image003

左边的那个好理解,中间对应的thread访问的地址交叉了,thread3和thread4交叉访问了,在硬件1.2版本之前的都会造成Non-Coalesced访问;

详细的需要说明的是右边的为什么也造成了Non-Coalesced(非对齐)访问,这个是基础问题,大家理解的内存对齐是怎么样的?按照固定思路,或者教材上强调的都是中间过程的内存访问的对齐,但是内存是从offset 0x00000000位置开始的,就是偏移量0开始的,如果要真的满足内存对齐,严格的说起来就需要从内存的0地址开始算起,再加上我们知道的global内存的对齐方式有几种,4位bytes,8bytes,16bytes,这里说的对齐方式,注意区别关系;再来看看右边的那个图:thread0开始,从address128的位置开始向下便宜的位置是?132-128=4 偏移了4,16个threads整体访问的是16*4=64,是从132开始的,从0算起来,132-0 =132; 132/16 = 8…4,从整体上讲,从0偏移位置开始,偏移了4个位置,这里的就造成了访问的未对齐,这个是从整体角度上讲的,和左边的图比较一下,那个是按照局部对齐来说的,注意理解;

继续看图说话:

clip_image004

左边的图看看,算一下,局部的时候偏移了,从局部和整体来说,都会引起未对齐访问;

右边的图自己算一下,是不是超出了刚才我说的范围;所以造成了内存访问的未对齐情况;

前面我们看的图都是1.2版本前的硬件的情况下的内存访问情况,现在看看1.2版本以后的硬件;

这里解释一下,什么叫1.2版本的硬件,或许有些朋友也不太了解,g80架构的都是1.0或者1.1的硬件架构,现在的gtx200系列的都是1.3的架构,其实1.2的硬件架构,或许是Nvidia的一个内部的,没有推出产品,可能准备提供给低端的产品,但是我想没有推出低端的产品,直接就上1.3device了,市场需求吧~~如果下一步GTX的架构还是按照老路子,不改进的话,或许Intel的Lrb上来以后,对Nvidia的产品,就是一个很大的竞争了;

不说废话了,先看图:

clip_image005

1.2以后的硬件版本,弱化了threads之间交叉访问的时候,没对齐的情况,只要大家都在一次访问的64bytes的一个段里面,或者128bytes的一个段里面面,这样的段访问,那就可以不用多次访问,当然如果你16个threads分别跨过了16个段,那就得产生16个存储event~记住几个段的定义,这里说的段,就是我们常常理解的对齐的方式,全局的内存访问对齐方式,8个bits的是按照32bytes对齐,16bits的是按照64bytes对齐,32bits和64bits都是按照128bytes对齐;

在优化代码的时候,这个地方是一个值得注意的部分;

API函数里面有对齐访问的接口,会按照对齐的方式分配global内存给你,不过注意其中的一个offset值的使用,这个是为了解决全局情况下的对齐偏移的问题:)cudaMallocPitch,这个函数,注意使用~

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/news/313167.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

C# 8 新特性 - 异步流 Asynchronous Streams

异步流 Asynchronous Streams例子 这是一个很简单的控制台程序。它有一个NumberFactory&#xff0c;它可以根据传递的参数来产生一串数字&#xff08;IEnumerable<int>&#xff09;。然后在这个程序中把每个数字都打印出来&#xff0c;同时在前边显示出当前的线程ID。 这…

__syncthreads()

http://www.cnblogs.com/dwdxdy/p/3215136.html __syncthreads()是cuda的内建函数&#xff0c;用于块内线程通信. __syncthreads() is you garden variety thread barrier. Any thread reaching the barrier waits until all of the other threads in that block also reach i…

互联网50周年!这有它的一张“出生证明”

2019 年 10 月 29 日是互联网的 50 周年&#xff0c;50 年前(1969 年 10 月 29 日)&#xff0c;加州大学洛杉矶分校的计算机将一个只有两个字母(LO)的数据包发送到斯坦福研究所的计算机上&#xff0c;这是互联网史上的第一个数据包&#xff0c;从此开启互联网时代的第一步。 当…

Eltwise_layer简介

http://www.voidcn.com/blog/thy_2014/article/p-6117416.html common_layer&#xff1a; ArgMaxLayer类&#xff1b; ConcatLayer类&#xff1a; EltwiseLayer类&#xff1b; FlattenLayer类&#xff1b; InnerProductLayer类&#xff1b; MVNLayer类&#xff1b; SilenceLaye…

PowerBI 秒级实时大屏展示方案 全面助力双十一

双十一来了&#xff0c;你准备好了吗&#xff1f;不管你是否准备完毕&#xff0c;我们带来了全网首发的 PowerBI 秒级实时大屏展示方案&#xff0c;你可以直接用来展示双十一的实时状况。我们一步步来说明这个套件模板教程。真实效果功能如下&#xff1a;全实时展示 双十一 当天…

caffe基本函数

http://blog.csdn.net/seven_first/article/details/47378697目录 主要函数 caffe_cpu_gemm 函数caffe_cpu_gemv 函数caffe_axpy 函数caffe_set 函数caffe_add_scalar 函数caffe_copy 函数caffe_scal 函数caffeine_cup_axpby 函数caffe_add caffe_sub caffe_mul caffe_div 函数…

优化 .net core 应用的 dockerfile

优化 .net core 应用的 dockerfileIntro在给 .net core 应用的写 dockerfile 的时候一直有个苦恼&#xff0c;就是如果有很多个项目&#xff0c;在 dockerfile 里写起来就会很繁琐&#xff0c;有很多项目文件要 copy&#xff0c;dockerfile 还不支持直接批量复制项目文件&#…

Ubuntu 上不了网

http://askubuntu.com/questions/441619/how-to-successfully-restart-a-network-without-reboot-over-ssh 指令&#xff1a; sudo service network-manager restart

C# 8 新特性 - 静态本地方法

从C# 8 开始&#xff0c;本地方法就可以是静态的了。 与其他的本地方法不同&#xff0c;静态的本地方法无法捕获任何本地状态量。 直接看例子&#xff1a; 这段代码里有两个本地方法&#xff0c;他们分别对实例的一个字段和方法里的一个本地变量进行了修改操作&#xff0c;也就…

Caffe使用其他人代码CUDA不兼容问题

只要将别人的上层代码使用即可&#xff0c;底层CUDA还是用自己的版本。

​.NET手撸2048小游戏

前言2048是一款益智小游戏&#xff0c;得益于其规则简单&#xff0c;又和 2的倍数有关&#xff0c;因此广为人知&#xff0c;特别是广受程序员的喜爱。本文将再次使用我自制的“准游戏引擎” FlysEngine&#xff0c;从空白窗口开始&#xff0c;演示如何“手撸” 2048小游戏&…

cannot import caffe

遇到问题“cannot import caffe”首先make pycaffe&#xff0c; 如果不行&#xff0c;分别尝试&#xff1a; 1. export PYHTONPATH/path/to/caffe/python:$PYTHONPATH 注&#xff1a;以上方法只能在当前terminal中生效&#xff0c;如需长久生效&#xff0c;需修改~/.bashrc …

自行实现高性能MVC

wcf虽然功能多、扩展性强但是也面临配置忒多&#xff0c;而且restful的功能相当怪异&#xff0c;并且目前没法移植。asp.net core虽然支持webapi&#xff0c;但是功能也相对繁多、配置复杂。就没有一个能让码农们安安心心的写webapi&#xff0c;无需考虑性能、配置、甚至根据问…

caffe matio问题

http://blog.csdn.net/houqiqi/article/details/46469981 注&#xff1a;如果指令行模式实在解决不了/lib/libcaffe.so: undefined reference to Mat_VarReadDataLinear问题&#xff0c;可以尝试在QT下进行训练和测试。 1&#xff0c; 下载matio(http://sourceforge.NET/pro…

技术管理者怎样跳出“泥潭”

近几年面试了不少新人&#xff0c;当问到职业规划时&#xff0c;大多都会说先积累技术&#xff0c;然后往架构师的方向发展。这可能是技术人的一个特质&#xff0c;喜欢跟机器相处&#xff0c;沉浸在代码之中&#xff0c;而不喜欢跟人打交道。现实的情况是&#xff0c;一些中小…

你或许以为你不需要领域驱动设计

作者&#xff1a;邹溪源&#xff0c;长沙资深互联网从业者&#xff0c;架构师社区合伙人&#xff01;一犹记得刚刚参加工作时&#xff0c;是地图厂商四维图新集团旗下的一家子公司&#xff0c;主要从事规划测绘相关软件研发的公司。当时我的项目是为勘测设计院提供相对应的应用…

undefined reference to Mat_VarCreate'

我们之前所写的方法&#xff0c;发现只能在QT编译模式下解决这个问题&#xff0c;而在指令行模式下不能编译通过&#xff0c;因为QT模式基于cmake&#xff0c;而make all基于Makefile 因而我们参照DeepLab的Makefile文件&#xff1a; 发现只要将其中关于matio的内容加到我们的…

redis为什么这么火该怎么用

最近一些人在介绍方案时&#xff0c;经常会出现redis这个词&#xff0c;于是很多小伙伴百度完redis也就觉得它是一个缓存&#xff0c;然后项目里面把数据丢进去完事&#xff0c;甚至有例如将实体属性拆分塞进redis hash里面的奇怪用法等等&#xff01;原因是什么呢&#xff1f;…

Deeplab训练时候unexpected label

先说一些有意义的答案&#xff1a; https://groups.google.com/forum/#!topic/caffe-users/CLF4IZ2Tkqk You probably need to save your labels as 1 Channel png images. Where the labels are encoded with a 0-based index --> in case of voc12 you would needs 21 l…

.Net Core实现健康检查

ASP.NET Core 提供运行状况检查中间件和库&#xff0c;以用于报告应用基础结构组件的运行状况。运行状况探测可以由容器业务流程协调程和负载均衡器用于检查应用的状态。例如&#xff0c;容器业务流程协调程序可以通过停止滚动部署或重新启动容器来响应失败的运行状况检查。负载…