【并行计算-CUDA开发】GPU 的硬体架构 关于存储体的介绍比较好bank conflicts
这里我们会简单介绍,NVIDIA 目前支援CUDA 的GPU,其在执行CUDA 程式的部份(基本上就是其shader 单元)的架构。这里的资料是综合NVIDIA 所公布的资讯,以及NVIDIA 在各个研讨会、学校课程等所提供的资料,因此有可能会有不正确的地方。主要的资料来源包括NVIDIA 的CUDA Programming Guide 1.1、NVIDIA 在Supercomputing '07 介绍CUDA 的session,以及UIUC 的CUDA 课程。 GPU 的基本介绍目前NVIDIA 推出的显示晶片,支援CUDA 的是G80 系列的显示晶片。其中G80 显示晶片支援CUDA 1.0 版,而G84、G86、G92、G94、G96 则支援CUDA 1.1 版。基本上,除了最早的GeForce 8800 Ultra/GTX 及320MB/640MB 版本的GeForce 8800GTS、Tesla 等显示卡是CUDA 1.0 版之外,其它GeForce 8 系列及9 系列显示卡都支援CUDA 1.1。详细情形可以参考CUDA Programming Guide 1.1 的Appendix A。 所有目前支援CUDA的NVIDIA显示晶片,其shader部份都是由多个multiprocessors组成。每个multiprocessor里包含了八个stream processors,其组成是四个四个一组,也就是说实际上可以看成是有两组4D的SIMD处理器。此外,每个multiprocessor还具有8192个暂存器,16KB的share memory,以及texture cache和constant cache。大致上如下图所示: 在CUDA 中,大部份基本的运算动作,都可以由stream processor 进行。每个stream processor 都包含一个FMA(fused-multiply-add)单元,可以进行一个乘法和一个加法。比较复杂的运算则会需要比较长的时间。 执行过程在执行CUDA 程式的时候,每个stream processor 就是对应一个thread。每个multiprocessor 则对应一个block。从之前的文章中,可以注意到一个block 经常有很多个thread(例如256 个),远超过一个multiprocessor 所有的stream processor 数目。这又是怎么回事呢? 实际上,虽然一个multiprocessor只有八个stream processor,但是由于stream processor进行各种运算都有latency,更不用提记忆体存取的latency,因此CUDA在执行程式的时候,是以warp为单位。目前的CUDA装置,一个warp里面有32个threads,分成两组16 threads的half-warp。由于stream processor的运算至少有4 cycles的latency,因此对一个4D的stream processors来说,一次至少执行16个threads(即half-warp)才能有效隐藏各种运算的latency。 由于multiprocessor 中并没有太多别的记忆体,因此每个thread 的状态都是直接保存在multiprocessor 的暂存器中。所以,如果一个multiprocessor 同时有愈多的thread 要执行,就会需要愈多的暂存器空间。例如,假设一个block 里面有256 个threads,每个thread 用到20 个暂存器,那么总共就需要256x20 = 5,120 个暂存器才能保存每个thread 的状态。目前CUDA 装置中每个multiprocessor 有8,192 个暂存器,因此,如果每个thread 使用到16 个暂存器,那就表示一个multiprocessor 同时最多只能维持512 个thread 的执行。如果同时进行的thread 数目超过这个数字,那么就会需要把一部份的资料储存在显示记忆体中,就会降低执行的效率了。 Shared memoryShared memory 是以4 bytes 为单位分成banks。因此,假设以下的资料: 因此,如果程式在存取shared memory 的时候,使用以下的方式: int number = data[base + tid]; 那就不会有任何bank conflict,可以达到最高的效率。但是,如果是以下的方式: int number = data[base + 4 * tid]; 这样不会造成bank conflict,因为所有的thread 都读取同一个位址的资料。 很多时候shared memory 的bank conflict 可以透过修改资料存放的方式来解决。例如,以下的程式: data[tid] = global_data[tid]; 会造成严重的bank conflict,为了避免这个问题,可以把资料的排列方式稍加修改,把存取方式改成: int row = tid / 16; 这样就不会造成bank conflict 了。 Global memory由于multiprocessor 并没有对global memory 做cache(如果每个multiprocessor 都有自己的global memory cache,将会需要cache coherence protocol,会大幅增加cache 的复杂度),所以global memory 存取的latency 非常的长。除此之外,前面的文章中也提到过global memory 的存取,要尽可能的连续。这是因为DRAM 存取的特性所造成的结果。 更精确的说,global memory 的存取,需要是"coalesced"。所谓的coalesced,是表示除了连续之外,而且它开始的位址,必须是每个thread 所存取的大小的16 倍。例如,如果每个thread 都读取32 bits 的资料,那么第一个thread 读取的位址,必须是16*4 = 64 bytes 的倍数。 如果有一部份的thread 没有读取记忆体,并不会影响到其它的thread 速行coalesced 的存取。例如: if(tid != 3) { 虽然thread 3 并没有读取资料,但是由于其它的thread 仍符合coalesced 的条件(假设data 的位址是64 bytes 的倍数),这样的记忆体读取仍会符合coalesced 的条件。 在目前的CUDA 1.1 装置中,每个thread 一次读取的记忆体资料量,可以是32 bits、64 bits、或128 bits。不过,32 bits 的效率是最好的。64 bits 的效率会稍差,而一次读取128 bits 的效率则比一次读取32 bits 要显著来得低(但仍比non-coalesced 的存取要好)。 如果每个thread 一次存取的资料并不是32 bits、64 bits、或128 bits,那就无法符合coalesced 的条件。例如,以下的程式: struct vec3d { float x, y, z; }; 并不是coalesced 的读取,因为vec3d 的大小是12 bytes,而非4 bytes、8 bytes、或16 bytes。要解决这个问题,可以使用__align(n)__ 的指示,例如: struct __align__(16) vec3d { float x, y, z; }; 这会让compiler 在vec3d 后面加上一个空的4 bytes,以补齐16 bytes。另一个方法,是把资料结构转换成三个连续的阵列,例如: __global__ void func(float* x, float* y, float* z, float* output) 如果因为其它原因使资料结构无法这样调整,也可以考虑利用shared memory 在GPU 上做结构的调整。例如: __global__ void func(struct vec3d* data, float* output) 在上面的例子中,我们先用连续的方式,把资料从global memory 读到shared memory。由于shared memory 不需要担心存取顺序(但要注意bank conflict 问题,参照前一节),所以可以避开non-coalesced 读取的问题。 Texture显示晶片上的texture cache 是针对一般绘图应用所设计,因此它仍最适合有区块性质的存取动作,而非随机的存取。因此,同一个warp 中的各个thread 最好是读取位址相近的资料,才能达到最高的效率。 对于已经能符合coalesced 规则的资料,使用global memory 通常会比使用texture 要来得快。 运算单元Stream processor 里的运算单元,基本上是一个浮点数的fused multiply-add 单元,也就是说它可以进行一次乘法和一次加法,如下所示: a = b * c + d; compiler 会自动把适当的加法和乘法运算,结合成一个fmad 指令。 除了浮点数的加法及乘法之外,整数的加法、位元运算、比较、取最小值、取最大值、及以型态的转换(浮点数转整数或整数转浮点数)都是可以全速进行的。整数的乘法则无法全速进行,但24 bits 的乘法则可以。在CUDA 中可以利用内建的__mul24 和__umul24 函式来进行24 bits 的整数乘法。 浮点数的除法是利用先取倒数,再相乘的方式计算,因此精确度并不能达到IEEE 754的规范(最大误差为2 ulp)。内建的__fdividef(x,y)提供更快速的除法,和一般的除法有相同的精确度,但是在2 216 < y < 2 218时会得到错误的结果。 此外CUDA 还提供了一些精确度较低的内建函式,包括__expf、__logf、__sinf、__cosf、__powf 等等。这些函式的速度较快,但精确度不如标准的函式。详细的资料可以参考CUDA Programming Guide 1.1 的Appendix B。 和主记忆体间的资料传输 |
【并行计算-CUDA开发】GPU 的硬体架构 关于存储体的介绍比较好bank conflicts相关推荐
- 【并行计算-CUDA开发】GPU 的硬体架构
GPU 的硬体架构 这里我们会简单介绍,NVIDIA 目前支援CUDA 的GPU,其在执行CUDA 程式的部份(基本上就是其shader 单元)的架构.这里的资料是综合NVIDIA 所公布的资讯, ...
- 【并行计算-CUDA开发】从零开始学习OpenCL开发(一)架构
多谢大家关注 转载本文请注明:http://blog.csdn.net/leonwei/article/details/8880012 本文将作为我<从零开始做OpenCL开发>系列文章的 ...
- 【并行计算-CUDA开发】CUDA线程、线程块、线程束、流多处理器、流处理器、网格概念的深入理解
GPU的硬件结构,也不是具体的硬件结构,就是与CUDA相关的几个概念:thread,block,grid,warp,sp,sm. sp: 最基本的处理单元,streaming processor 最 ...
- 【并行计算-CUDA开发】显卡两大生产商
ATI显卡 ATI显卡即AMD显卡.俗称A卡.搭载AMD公司出品的显示芯片.与NVIDIA齐名,同为世界两大显示芯片厂商. 不同的是AMD不是只有显卡,而且还出品CPU(处理器),其AMD处理器与In ...
- 【并行计算-CUDA开发】关于共享内存(shared memory)和存储体(bank)的事实和疑惑...
关于共享内存(shared memory)和存储体(bank)的事实和疑惑 主要是在研究访问共享内存会产生bank conflict时,自己产生的疑惑.对于这点疑惑,网上都没有相关描述, 不管是国内还 ...
- 【并行计算-CUDA开发】从熟悉到精通 英伟达显卡选购指南
举报 说到显卡,就不免令人想到英伟达和AMD两家面向个人消费级和企业级最大的显示芯片生产企业,英伟达和AMD,今天小编为大家简单的介绍一下英伟达的显卡选购方面的攻略,为一些想要购买显卡的用户提供一些参 ...
- CUDA编程——GPU架构,由sp,sm,thread,block,grid,warp说起
目录 1.从硬件看 2.从软件看 3.对应关系 4.SIMT和SIMD 掌握部分硬件知识,有助于程序员编写更好的CUDA程序,提升CUDA程序性能,本文目的是理清sp,sm,thread,block, ...
- 一文了解GPU并行计算CUDA
了解GPU并行计算CUDA 一.CUDA和GPU简介 二.GPU工作原理与结构 2.1.基础GPU架构 2.2.GPU编程模型 2.3.软件和硬件的对应关系 三.GPU应用领域 四.GPU+CPU异构 ...
- 3维线程格 gpu_基于CUDA的GPU并行优化重力三维反演
重力勘探由于其成本较低.施工方法方便等, 被广泛应用于大尺度的地质异常体勘查.大范围找矿普查.以及小比例尺密度三维地质建模等工作中.目前常用的反演方法有两种, 2.5维联合3维界面反演[和三维物性反演 ...
最新文章
- 关于git bush 中不能复制黏贴的问题
- ADS1.2安装教程
- linux nc命令用法举例
- netflow报文格式与数据处理流程分析_RTK、三维激光扫描、无人机倾斜摄影在大比例尺地形图测绘的对比分析...
- python时间计算_python datetime库使用和时间加减计算
- python字符串转浮点数_如何在Python中检查字符串是否为数字(浮点数)?
- springboot接收json参数_Springboot + Vue + shiro 实现前后端分离、权限控制
- 【毕业设计】一种多商家网络商店的设计与实现(源代码+论文)
- 比尔·盖茨不玩IT了
- 高性能计算多集群管理平台
- python列表的内置方法_python 基础之列表切片内置方法
- 阻滞增长模型求解_马尔萨斯与阻滞增长模型对于人口预测的分析
- atlas 1.1.0 的本地开发文档(编译)
- 江苏农村商业银行计算机类笔试考什么时候,2020江苏农商行笔试考试内容是什么?...
- leetcode 803.打砖块(C/C++/Java/python)
- SAP ABAP 标题左边 GOS文件菜单 自定义 CL_GUI_GOS_CONTAINER MENU Gos based menu GOS BUTTON
- apk自行修改后的操作(软件安装不了,安了打不开,闪退)
- i7 11800h参数
- arch linux格式化,用Arch linux打造自己的操作系统(一)
- ESFP型人格的爱情分析,ESFP型人格的职业分析