【并行计算-CUDA开发】GPU 的硬体架构
这里我们会简单介绍,NVIDIA 目前支援CUDA 的GPU,其在执行CUDA 程式的部份(基本上就是其shader 单元)的架构。这里的资料是综合NVIDIA 所公布的资讯,以及NVIDIA 在各个研讨会、学校课程等所提供的资料,因此有可能会有不正确的地方。主要的资料来源包括NVIDIA 的CUDA Programming Guide 1.1、NVIDIA 在Supercomputing 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 所有目前支援CUDA的NVIDIA显示晶片,其shader部份都是由多个multiprocessors组成。每个multiprocessor里包含了八个stream processors,其组成是四个四个一组,也就是说实际上可以看成是有两组4D的SIMD处理器。此外,每个multiprocessor还具有8192个暂存器,16KB的share 详细的multiprocessor资讯,都可以透过CUDA的cudaGetDeviceProperties()函式或cuDeviceGetProperties()函式取得。不过,目前还没有办法直接取得一个显示晶片中有多少multiprocessor的资讯。 在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 由于multiprocessor 中并没有太多别的记忆体,因此每个thread 的状态都是直接保存在multiprocessor 的暂存器中。所以,如果一个multiprocessor 同时有愈多的thread 要执行,就会需要愈多的暂存器空间。例如,假设一个block 里面有256 个threads,每个thread 用到20 个暂存器,那么总共就需要256x20 = 5,120 个暂存器才能保存每个thread Shared memory目前CUDA 装置中,每个multiprocessor 有16KB 的shared memory。Shared memory 分成16 个bank。如果同时每个thread 是存取不同的bank,就不会产生任何问题,存取shared memory 的速度和存取暂存器相同。不过,如果同时有两个(或更多个) threads 存取同一个bank 的资料,就会发生bank conflict,这些threads 就必须照顺序去存取,而无法同时存取shared Shared memory 是以4 bytes 为单位分成banks。因此,假设以下的资料: __shared__ int data[128]; 那么,data[0] 是bank 0、data[1] 是bank 1、data[2] 是bank 2、…、data[15] 是bank 15,而data[16] 又回到bank 0。由于warp 在执行时是以half-warp 的方式执行,因此分属于不同的half warp 的threads,不会造成bank conflict。 因此,如果程式在存取shared memory 的时候,使用以下的方式: int number = data[base + tid]; 那就不会有任何bank conflict,可以达到最高的效率。但是,如果是以下的方式: int number = data[base + 4 * tid]; 那么,thread 0 和thread 4 就会存取到同一个bank,thread 1 和thread 5 也是同样,这样就会造成bank conflict。在这个例子中,一个half warp 的16 个threads 会有四个threads 存取同一个bank,因此存取share memory 的速度会变成原来的1/4。 一个重要的例外是,当多个thread 存取到同一个shared memory 的位址时,shared memory 可以将这个位址的32 bits 资料「广播」到所有读取的threads,因此不会造成bank conflict。例如: int number = data[3]; 这样不会造成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 读取的问题。 TextureCUDA 支援texture。在CUDA 的kernel 程式中,可以利用显示晶片的texture 单元,读取texture 的资料。使用texture 和global memory 最大的差别在于texture 只能读取,不能写入,而且显示晶片上有一定大小的texture cache。因此,读取texture 的时候,不需要符合coalesced 的规则,也可以达到不错的效率。此外,读取texture 时,也可以利用显示晶片中的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 不能直接存取主记忆体,只能存取显示卡上的显示记忆体。因此,会需要将资料从主记忆体先复制到显示记忆体中,进行运算后,再将结果从显示记忆体中复制到主记忆体中。这些复制的动作会限于PCI Express 的速度。使用PCI Express x16 时,PCI Express 1.0 可以提供双向各4GB/s 的频宽,而PCI Express 2.0 则可提供8GB/s 的频宽。当然这都是理论值。 从一般的记忆体复制资料到显示记忆体的时候,由于一般的记忆体可能随时会被作业系统搬动,因此CUDA 会先将资料复制到一块内部的记忆体中,才能利用DMA 将资料复制到显示记忆体中。如果想要避免这个重复的复制动作,可以使用cudaMallocHost 函式,在主记忆体中取得一块page |
最新文章
- Node.js 教程 01 - 简介、安装及配置
- Ubuntu下安装QQ22013
- replace实现正则过滤替换非法字符
- VB6.0 为批量字体改名
- 控件包含代码块(即 <;% ... %>;),因此无法修改控件集合
- project.VERSION_NAME定义
- Hanio汉诺塔代码递归实现
- android 四大组件Broadcast Receiver
- 阿里云至 Windows Azure 的 Linux 虚拟机迁移
- python命令行解析工具argparse模块【1】
- SpringMVC handleMapping 处理器映射器 属性清单
- shiro框架的使用实例
- 框架学习之Spring(一IOC)----HelloWrod
- Docker技术应用场景(转载)
- 爬虫之xpath用法
- Mybatis之拦截器原理(jdk动态代理优化版本)
- 【TJOI2015】线性代数
- SQL语句中 chinese_prc_CS_AI_WS 以及replace用法
- BAT脚本如何自动执行 adb shell 以后的命令
- plantuml使用教程【转】
热门文章
- Codeforces Round #588 (Div. 2) C. Anadi and Domino(思维)
- 部分易错JS知识点整理(缓慢填坑)
- 图论小专题B
- Zabbix 数据库迁移
- Codeforces 940 E.Cashback (单调队列,dp)
- zookeeper系列(九)zookeeper的会话详解
- java面试题,转载自http://www.cnblogs.com/nnngu/p/8471043.html#3914167
- koa 项目实战(九)passport验证token
- centos7 开启80端口
- redis 概述及部署 安装php和python客户端