基于CUDA平台的GPU并行计算技术研究终稿.docx

1、基于CUDA平台的GPU并行计算技术研究终稿CUDA架构下GPU硬件结构GPU与CPU的最大不同点在于硬件结构自身，而硬件结构不同是由于晶体管用途分布造成的，如图2.1所示。GPU比CPU在数据计算能力和存储器带框上有相对大的优势，在能耗和价格上付出的代价也相对较小，从而在异构操作协同处理运算上占领并行运算的一席之地。GPU作为图形渲染的主要硬件，因图像渲染运算的高度并行性，它可以采用添加ALU和Control Cache（存储器控制单元）的方式提高运算能力和存储器带宽。CPU的制造厂家是把更多的晶体管用于复杂的控制单元以及缓冲区，而GPU的晶体管作为ALU的居多，从何提高GPU的总体执行单元

2、效率。图2.1 CPU与GPU硬件结构比较在传统的CPU+GPU异构并行的系统中，复杂逻辑的事务计算由CPU完成，这样的要求，也促使CPU采用分配大量缓存、分支预测以及复杂控制逻辑的方式较快得获取数据和指令，以此提高CPU的运算速率；而GPU负责的是计算高度密集的图像渲染工作，它通过分配更多的ALU来运行量大但相对简单的线程（Thread），且要求较高的存储器带宽以保证整体的数据传输速率。CPU和GPU的不同设计理念造就了这两个处理器的架构差异及性能差异，具体如下：1）线程差异，CPU的多线程是一种软件粗粒度多线程，在执行的线程中断后，需要保存中断线程前后数据同时装载下一个线程的数据，这个过程

3、需要很大的时间成本；GPU的多线则不存在这样的数据交换过程，因此在并行执行过程中省去更多时间，提高效率。2）计算核心差异，主流CPU采用多核技术，3到6条执行流水线在每个计算核心中，乱序执行、预测执行和大容量缓存技术都被采用到CPU核心中意达到提高指令级并行的目的，这些技术的加入也限制了更多核心数量的集成。而GPU集成了多个流多处理器，因此每个GPU可被看成1到30个SIMD处理器，其中每个SIMD处理器包含8个ID流处理器，GPU的并行则是利用了多个流处理器间的粗粒度并行及流多处理器内的细粒度数据并行。ALU较同时期CPU多的主流GPU在单精度浮点处理能力比CPU高大约15倍。3）带宽差异，

4、因GDDR存储器存在相对位置不一致导致CPU内存的信号完整性问题考虑比GPU显存严谨，故而同等条件下显存的工作效率比内存高；其次，多存储器控制单元的GPU比存储器控制单元为双通道或者单通道的CPU在同时访问的存储器粒数量要高。图2.2为2007年以前GPU与CPU带宽比较。4）缓存差异，为降低访存延迟并节约带宽，CPU存在使用大量晶体颗粒的缓存，为了保持缓存与内存数据一致，还添加了复杂的缓存逻辑控制单元；GPU则没有如此复杂的缓存机制和替换体系，因GPU缓存只读，减少显存带宽，从而控制单元相对简单些。 1.1. CUDA软件体系1.1.1. CUDA的存储器模型CUDA架构下的GPU是一种多层

5、存储空间复杂结构体，这是GPU自身的发展历史和GPU原本的设计理念决定的。如图2.3所示，该图表明了CUDA的存储器结构包括六种存储器：全局存储器（Global Memory）、寄存器（Registers）、本地存储器（Local Memory）、共享存储器（Shared Memory）、常量存储器（Constant Memory）、纹理存储器（Texture Memory），其中前三者支持读/写操作，而后两者只具备读操作。这些存储器在性能上存在较大差别，它们的结构各异、运算速度不同、大小及存储规模不一样，也决定了CUDA在并行运算过程中存在相同并行运算过程有不同的运算加速收益。共享存储器、寄

6、存器以其独特的结构组成方式在这些存储器中在数据流访问操作速度方面是相对较快的。对于每个并行的运算线程来说，寄存器、局部存储器都是私有的，也就是说每个线程单独存在使用各自的寄存器及局部寄存器。之中，作为GPU上的高速缓存区，寄存器能够在极低延迟的情况下别每个线程的ALU所读取。数量客观的寄存器文件（Register File）构成了寄存器的基础单元，并且被分配给每个正在执行的并行线程，因此，在并行执行的过程中每个线程所拥有的寄存器大小及其有限，在编程过程中应尽量为每个线程分配足够且较少的的私有变量。局部存储器是寄存器在线程执行过程中的补充，当寄存器使用过多，或者并行使用到结构体、数组这样的大数据

7、变量，局部寄存器就充当暂时的收容所，为这些私有数据分配必要的存储空间，其区别就在于，分配到局部存储器的数据是存在于显卡显存中，而不是像寄存器在边上或者缓冲区中，所以局部存储器对数据的访问操作会比寄存器慢很多。与寄存器类似，共享存储器同样也是GPU上的高速存储器，但如同名字一样，存在于共享存储器的数据是被同一个线程块（Block）中的所有线程共享的。分配这样一种存储器的目的在于解决线程块内各个线程的通信问题，这个也是CUDA的创新之处，以往的GPU并行运算技术在执行过程中线程之间的通行总是难以解决，各个线程之间的通行过程总是会存在很长的延迟，而CUDA因为共享寄存器的存在，其在保存共用计数器或者

8、线程块内的公共结果时，省去大量时间进行线程间通行，所以相应的速率会提升很多。占有显存很大一块的全局存储器，GPU、CPU可以直接访问操作其中的全局数据，也就是说执行并行运算的任意一个线程、串行运算的部分都可以读写当中的任意区域的数据，但访问操作过程有很高的延迟。全局存储器带宽很高，在遵守访问要求前提下，避免分区冲突，如此高的带宽才能被充分的利用好。同样位于显存的常数存储器拥有缓存加速，且为只读存储器。GPU分配给常数存储器的空间只有64K，可用于存储并行运算执行过程中需要多次读取的只读变量。由于它的缓存机制，常量存储器节约了带宽，提高了读取的速率，同时也因为数据只读，因此不存在缓存不一的问题。

9、特殊的只读存储器还有纹理存储器，开辟这样的一块存储器也是由于GPU纹理渲染功能的特殊性。与常量内存类似，纹理内存存在缓存从而减少了对内存的请求并提供更高的带宽。图2.4为纹理内存原理的示意图，图中的地址并不连续，普通的CPU缓存机制下这类地址得不到缓存，而GPU就可以在该种模式下进行缓存加速，获取并行程序性能的提升。1.1.2. CUDA编程模型CUDA软件体系主要包括了CUDA C编译器NVCC以及与相关显卡关联的各类硬件开发工具。整个的CUDA软件开发体系由运行期环境（Runtime）、开发库模板（Library）、配套驱动（Drvier）三个部分组成，如图2.5所示。正如CUDA名字一样

10、，为了更好地利用GPU及CPU上的运算单元，CUDA在数据执行并行过程中采用了统一的处理架构进行处理操作。统一架构下，任意写入功能及共享存储器的存在确保了线程之间能更快速地通信。CUDA的编程模式下，程序将CPU作为主机（host），将能用于CUDA并行运算的GPU作为设备（device），随着CPU+GPU异构的发展，在CUDA并行运算中可存在多个设备，但主机只能有一个。用于cuda开发应用的程序段在程序执行运算过程中也相应地分为两类，运行在CPU上的程序代码为主机（host）代码，运行在GPU上的程序代码为设备（device）代码，kernel即为后者运行的并行程序段。并行运算研究人员以C

11、UDA C对GPU编程，CUDA C是工业标准C的一个简单的扩展。一般情况下，C函数为串行执行一次，而在调用kernel函数时，函数执行的过程在GPU上是重复的，也就是说GPU上分配了N个线程执行并行加速，则kernel就被平行地切不相依靠执行N次，这个也是并行运算的基本原理，可以理解为kernel函数只是CUDA程序中被并行执行的一个步骤，如图2.6所示，线程中执行过程对存储空间的操作及关系在2.2.1节已阐述，GUP在执行的处理器层次与CUDA线程数目相对应。Kernel函数存在双层次的并行，一个是线程格（Grid）中线程块间的并行，另一个则是线程块中线程间的并行，这个双层次的并行也是CU

12、DA重要创新之一。在对GPU进行编程前，首先要考虑的是要对程序中哪部分的运算进行并行加速。主机端对CPU的数据操作及串行处理部分代码和设备端对GPU执行并行运算的一系列操作代码构成了完整的CUDA程序，这些代码执行顺序与语句的顺序有关，执行与语句满足顺序一致性。Kernel函数启动前的数据准备及GPU设备初始化工作、kernel与kernel之间的串行操作都由CPU端代码完成。就目前看来，最理想的CPU串行代码一定程度上说是为kernel函数服务的，即体现处理并行上下文的作用。核函数Kernel的定义必须通过_global_函数类型限定符定义，通知CPU这段函数执行时是在GPU上执行，且仅能在

13、CPU端代码代用该函数，调用过程中通过运算符对kernel函数的执行参数进行声明,其中grid表示执行并行运算的线程格中线程块的数量，block表示执行并行运算的线程块中线程的数量。这两个参数的配置对执行并行运算的效率收益也有重要的影响，并行执行所操作数据量大，则着重考虑block大小，但当数据吞吐量很小的时候就要同时考虑grid和block之间的大小关系了。这个问题也是并行程序优化的一个重要内容。执行并行运算的线程总数可以由上述两个参数grid和block计算得出，等于两个参数的乘积。在线程块中，每个线程在线程块内都有自己相应身份标识索引，这些索引由线程在线程块中的相对位置以及线程与线程块的

14、维度关系确定。在一维线程块内线程的访问操作，程序员可以直接通过线程索引对线程进行寻址访问，即利用CUDA专用变量threadIdx进行寻址。当系统为二维线程块的时候，类似的可以通过线程的索引、线程块的维度大小以及线程块的索引进行寻址，对应线程块索引的专用变量为blockIdx，而维度大小则对应专有的blockDim进行访问。例如，并行程序调用的了一维线程块进行并行运算，这线程ID与线程索引相等；当调用的是二维线程块，大小为（Dim.x，Dim.y），线程的索引为（x，y），则线程ID=x+y*Dim.x；当调用的是三维的线程格时亦可类似地得出相应的线程ID。为了更好地说明CUDA编程模型，下面

15、以两个一维数组相加为例，给出CUDA计算核心的伪代码：_global_ void Add(int *X,int *Y,int *Z)Int i=threadIdx.x;Zi=Xi+Yi;Int main()/Kernel Invocation With N Threads;Add(X,Y,Z);伪代码中的Add函数在GPU平台上被调用的所有线程并行执行，但每个线程因系统给其设置了专用的私有ID进行身份标识，所以每个线程的并行运算还是能区别于其他的线程独立地进行。1.2. CUDA并行运算的优化CUDA在并行运算方便性能是优良，但优良的性能受制于很多因素。在解决同样的并行运算问题时，因存储器模型

16、多样化、算法多样化、任务划分方式多样化而导致程序编写的多样化，GPU的速率收益往往会存在很大的区别。在这些众多的区别中寻求精度上满足计算要求并尽可能多得提升运算速度，避免造成不必要的资源浪费，则需要我们在诸多CUDA问题解决方式中寻求最优解，这个过程就是CUDA并行运算的优化过程。而程序优化的最终标准应同时具备以下特点：1） 在等额的数据量下，最优算法的计算复杂度不能明显低于选定的算法。2） active block的数目不低于2，active warp的数量足以令SM满负载运行，可以在数据访问时有效使访问延迟不可见。3） 当运算出现关口时，程序执行的效率已经得到足够的优化。4） 当在内存访问

17、或者I/O出现关口时，合适的存储器应当被合适的访问模式选用于程序存储数据，以便此程序硬件拥有最大的带宽。针对上述的优化程序几个特点，我们可以从算法、数据和任务划分、显存访问方式、指令流、存储资源的整合、GPU与CPU通行方面对编写的CUDA程序进行优化，达到资源利用率最好，运算速率最高的目的。在评估并行程序运算是否有很高的效率收益，该收益是多少等指标，我们要使用到运行时间这项比较因素。CUDA程序中内核函数的运行时间有两种方式获取，一个是通过主机端测量，一个则是通过设备端进行测量，但是CUDA API如cudaMalloc()、cudaMemcpy()等函数则必须通过主机端进行测量。为了保证测

18、量时间的准确性，在选定好测量方式后应多次对核函数进行时间测量求其平均值，并确保GPU是处于工作状态而非节能模式，这样的测试结果会相对可靠。在设备端对内核函数的运行进行测时则需要GPU的计时器时戳进行计时，这也有两种方式，一种是调用clock()函数，但一般测得的是一个block在GPU中保持上下文的时间，这个测量的时间会比block实际执行kernel时间长；一种则是通过调用CUDA API的事件管理功能进行测时，就方便而言，使用CUDA API的事件管理功能计时在程序书写复杂度上比调用clock()函数简单很多。CUDA主机端的测时则是用到CPU的计时器，CUDA API的函数时异步的，若想

19、通过主机端对GPU执行程序时间进行准确的测量，则需要调cudaThreadSynchronize()对GPU和CPU线程进行同步。CUDA程序优化的第一个目标是对并行部分进行算法的优化，这部分的优化内容主要是对并行程序的算法复杂度进行分析简化，使算法的执行符合硬件最简化执行标准。针对第二个优化后程序特点，CUDA程序的优化势必会以提高SP（Stream Processor）的使用率为目标，从而避免了GPU中计算单元的浪费。在并行运算过程中最大可能地去利用计算资源、利用硬件条件优势提高执行效率，需要确保并行运算过程中能保有足量的线程参与到算法的运行中，且线程数目设定能根据问题的实际规模进行限定，

20、并非是为了增大SP的使用率而纯粹地增加线程的量，这样会出现SP空转，致使空置的大量线程等待执行并行运算的少量线程，反而造成资源的另一种浪费。另一种情况则是设定的线程相对较小，则会出现线程的切换，并行运算问题的规模以及线程设定的多少决定着线程切换的频率，当切换频率过高或者多个线程存在切换，亦会降低并行运算的效率。因此在设定并行程序的线程数目以及线程块、线程格的大小时，应合理地根据问题的规模进行设定，避免这两种极端的情况发生，从而提升执行的效率。一般情况下，处理器的内存访问带宽是比处理器的计算能力小很多的，存储器的带宽也在影响并行运算的性能。CUDA内核程序与主机程序不同，IO和事物处理不再其考虑

21、范围内，GPU内每个线程束（warp）只完成两件事情：访问存储器数据和运算。即使NVIDIA公司在GPU的显存带宽上下了很重的功夫并实现了很高的面积带宽比，但是这个有效的带宽比也只有在满足合并访问前提下才能获得很好的利用。合并访问是进行访问存储器数据的操作中，把多个数据进行合并到规定的字长之后再进行访问。以图2.7为例展示half-warp（半线程束）访存优化的操作。根据GPU的工作特点和CUDA的编程需要，半线程束的操作被利用到warp块访存中。遵循特定的字长对对齐的字段进行操作，半线程束的线程能够满足访存优化操作条件，大大提高访问的效率。如图例中16个线程对两个64Byte对齐的段和128

22、Byte对齐的段进行访问存储器操作，以固定的字长16Byte进行合并访问，实现了该16个线程只经过一次合并就可以满足自身对数据的访问需求，得到较高的带宽。在考虑完并行算法自身的可行性及复杂程度以及上诉问题的优化后，CUDA程序在GPU和CPU间通行的延迟也可以通过优化可以有效隐藏。CUDA的函数调用有两个不同的方式，一个是CPU线程在不需要等待函数返回就可以运行后续指令的异步调用，一个是与之相反的同步调用。显而易见，CUDA函数的异步调用方式能确保CPU和GPU两者的并发操作，是有效的隐藏CPU与GPU通信延迟的最好函数调用方式，图2.8是两种kernel执行方式比较图例。当CUDA使用到异步

23、调用时，CUDA控制GPU的并发执行是由流操作的，cudaMemcpyAsyn()函数可用于数据的异步调用操作，但需要指明函数所属的流，这样的要求同样也适用于调用的kernel函数。2. 基于内容的图像搜索2.1. 图像检索概况在这个获取信息可通过多个渠道的年代，图像信息是继文字信息之后又一个大的获取信息渠道。文字搜索引擎已然成熟，新兴的以“图”搜“图”既往开来，得到了XX、Google等几大搜索引擎公司的重视。这种以“图”搜“图”的检索方式即为基于内容的图像搜索。在此前，图像的搜索技术主要是基于文本的图像检索技术（Text-based Image Retrieval，简称TBIR），利用文本

24、描述的方式描述图像的特征，如绘画作品的年代、作者、尺寸、流派等。到90年代以后，出现了对图像的内容语义，如图像的纹理、颜色、布局和图像对象的形状等作为内容特征进行采集和检索的图像检索技术，即基于内容的图像检索（Content-based Image Retrieval，简称CBIR）技术。CBIR技术根据图像本身数据阵列、图像上下特征进行搜索，根据图像语义特征，在图像数据库中寻找出具有类似特征的相应图像。图像的规模一般要大于纯粹的文本信息，因此，CBIR技术在检索的速率上要求较TBIR技术要高。在早期，众多应用于实践环境的CBIR系统纷纷涌现，如由麻省理工学院研发的Photobook系统，以及

25、由IBM公司开发的最早商业化QBIC系统、哥伦比亚大学研发的WebSeek系统等，而后因自身就做网络搜索引擎的更具普遍性的Google搜图、Baidu搜图等CBIR系统也随之出现。在检索过程上，无论是CBIR还是TBIR，其基本的思路主要包括以下几个方面：首先是对用户搜索需求的采集和转变，构成可以检索索引数据库的提问；其次是收集和加工源图像，提取相应的图像特征，研究并对图片进行标引，建立图像的索引数据库；最后是依据自身设定的相似度算法，计算用户提问与索引数据库中记录的相似度大小，做满足设定阈值记录提起并将其作为结果，以设定相应相似度降序或者升序的方式输出。由上可以看出CBIR一般包括图像处理模

26、块、查询模块、对象库、特征库和知识库。把目标图片构成索引或者在构建图像索引数据库的时，需要确定索引的表示方式、索引的提取方式、索引的存储方式等等。在索引的提取方式中，就目前而言，大多数的搜索引擎采用的是自动提取，即通过一定的算法利用计算机获取。而索引的存储方式因索引的组织方式有关，当索引的维数很高时，采用传统的存储结构很难把高维索引组织起来，这便需要在数据结构研究领域中寻找合适的组织方式对索引进行存储。索引的表示因选取的图片特征不同可以分为以下几种：关键词索引、色彩特征索引、纹理特征索引、形状特征索引等，其中关键字索引属于TBIR内容，后四个为CBIR的范畴，为了带来更佳的效果，往往这几个索引

27、表示方式可以互相结合，提高用户在图像搜索这方面的满意程度。索引的表示、提取、组织之后就是对比时特征相似度量方式。1977 年，Amos Tversky 提出了著名的特征对比模型（contrast model），Tversky 认为对每个实体用一个特征集进行表示。设两个实体a和b，它们实体的特征集分别为A和B，则两个特征之间的相似性应当满足匹配性、单调性和独立性假设。在这样的假设基础上，Tversky 提出了contrast modeI定理，对于度量函数S满足上述的3个假设，那么存在一个相似度量函数S和一个非负函数f，以及两个常量，0,对于实体a、b、c、d，它们的特征集为A、B、C、D，有：

28、（3.1） （3.2）其中f是一个反映特征显著性的函数，衡量指定特征对相似性的贡献，当时，相似函数是不对称的。2.2. CBIR的结果评价在基于内容的检索中，因为搜索的算法相对较多，我们需要评判每种算法的检索结果，以比较其搜索的结果优劣性。在相同的条件下，搜索算法的结果评价方式能够寻找最优算法，令不同的搜索方式能更好地得到改进和提高，以使CBIR朝更好的方向发展。从前期的研究状况来看，对于搜索系统的响应速率和吞吐率的评价论述得较少，搜索结果的是否正确被更多地放在对检索效果评价中，主要使用的是查准率（precision）和查全率（recaII）两个指标。查全率的主要含义是在一次查询过程中，用户所

29、查到的相关图像的数目和数据库中与目标图像相关的所有图像数目之比；而查准率主要指在一次查询过程中所查到的相关图像数目同该次查到的所有图像数目之比。用户在评价算法的时候，可以预先选定含有特有目标的图像作为一组相关的图像，然后根据返回的结果计算查全率和查准率.查全率和查准率越高，说明该检索算法的效果越好。而近几年来因CBIR的快速发展，人们在满足查全率和查准率的高要求前提下，同时更加注重整个系统的响应时间以及吞吐率，即对搜索过程的等待时间以及数据库的海量要求提出更高的要求。对于图像数据来说，相比所含信息量相等的文本文件，其所占的存储空间要大得多，也就是说在对图像进行处理时，所处理的数据量相对较多。而

30、在现今对搜索速率、数据库量要求的提高，促使着人们在CBIR技术中寻求加速的突破口，利用并行算法、插值等方式提高搜索速度，以达到用户的对时间以及数据量的需求。3. CUDA与CBIR的结合实验过程3.1. 感知哈希算法本章在原有搜索引擎的最简单搜索原理基础上，利用CUDA平台对搜索过程中的灰度平均值计算部分进行加速优化，论述CUDA在图像处理领域中的应用。将目标图像在待选图像队列中进行图像搜索的原理多种多样，但基本上每种方式可以分为三个大步骤：1） 对图像进行色彩特征提取。如图3.1为对24位真彩图进行颜色特征提取的演示图。图像作为一个重要的信息传播载体，因信息的多样性，其包含的各类信息也会有所

31、不一样，在每幅图像中不一样的信息提取到存储在图像文件中的信息代码则是获取信息的另一种方式。利用计算机编码描述图像的算法也是多种多样，主要的有：SIFT（Scale-Invariant Transform）描述子、捆绑特征（Bundling Feature）算法、哈希函数（Hash Function）等。根据图像的不同，描述图像的算法可以设计得不一样，例如通过图像局部N阶矩的方式获取图像特征。2） 将获取的图像特征信息根据自身设计算法进行编码，并对海量的图像编码获得待查找表即图像的数据库。这里的编码可以根据自身的需求对分辨率较高的图像进行降采样，以减少编码的运算量。3） 对目标图像进行重复上述两

32、个步骤，即对目标图像进行特征提取和编码处理。目标图像编码与数据库内编码进行相似度匹配运算。在图像搜索引擎中的图像数据库进行全局或者局部相似度计算，在满足一定鲁棒性的前提下设定阈值，选取相似度最高的图像作为输出结果。 正如图像搜索原理中所论述，描述图像的算法多种多样，在众多的算法中，因各图像文件类型数据存储具有一般性，本文所研究的仅局限于针对24位真彩位图的图像搜索，对其他类型的图像可根据其要求进行算法调整，且因图像数量的约束，本文所获取图像数据库也具有一定的局限性。根据上述原理，本文采用的是一种相对简单的算法“感知哈希算法”（Perceptual Hash Algorithm，下面简称PHA），其属于哈希算法中的一种，作用是对目标图像生成一个指纹字符串，然后比较数据库中不同图片的指纹，结果越接近说明图片越相似。下面是实现PHA的基本过程：1） 图像缩小。将图片缩小至8*8像素的尺寸，总共得到64个像素，去除图片的细节，摒弃不同大小图片带来的差异，只保留图像结构、明暗等基本信息。2） 色彩简化。将缩小后的图像转变成灰度图，因针对的是