风辰的CUDA入门系列教程

风辰的CUDA入门系列教程





1. CUDA简介

GPU是图形处理单元(Graphic Processing Unit)的简称,最初主要用于图形渲染。自九十年代开始,GPU的发展产生了较大的变化,NVIDIA、AMD(ATI)等GPU生产商敏锐的观察到GPU天生的并行性,经过他们对硬件和软件的改进,GPU的可编程能力不断提高,GPU通用计算应运而生。由于GPU具有比CPU强大的计算能力,为科学计算的应用提供了新的选择。

最早的GPU的开发直接使用了图形学的API,将任务映射成纹理的渲染过程,使用汇编或者高级着色器语言Cg,HLSL等编写程序,然后通过图形学API执行(Direct3D和OpenGL),这样的开发不仅难度较大,而且难以优化,对开发人员的要求非常高,因此,传统的GPU计算并没有广泛应用。

2007年6月,NVIDIA公司推出了CUDA(Compute Unified Device Architecture),CUDA不需要借助图形学API,而是采用了类C语言进行开发。同时,CUDA采用了统一处理架构,降低了编程的难度,使得NVIDIA相比AMD/ATI后来居上。相比AMD的GPU,NVIDIA GPU引入了片内共享存储器,提高了效率。这两项改进使CUDA架构更加适合进行GPU通用计算。由于这些特性,CUDA推出后迅速发展,被应用于石油勘测、天文计算、流体力学模拟、分子动力学仿真、生物计算、图像处理、音视频编解码等领域。

由于采用的是C/C++编译器为前端,以C/C++语法为基础设计,因此对熟悉C系列语言的程序员来说,CUDA的语法比较容易掌握。CUDA只对ANSI C进行了最小的必要扩展,以实现其关键特性--线程按照两个层次进行组织、共享存储器和栅栏同步。

这些关键特性使得CUDA拥有了两个层次的并行:线程级并行实现的细粒度数据并行,和任务级并行实现的粗粒度并行。





2. linux下CUDA开发环境构建


前一节已经简单了说了一下CUDA,为了能够使用CUDA开放,这一节将说明怎样构建CUDA开发环境。

本节讲解在ubuntu9.04操作系统和gcc前端的基础上安装CUDA开发环境。
   
首先,要保证自己机器上的gcc能够使用,因为ubuntu缺少gcc的一些包和g++,所以这些得自己安装。安装命令:sudo apt-get install g++,待其完成后,弄个C代码试试看:);当然你得保证你的显卡支持CUDA。
   
其次,到nVidia官方网站(http://www.nvidia.cn/object/cuda_get_cn.html)上下载对应操作系统的驱动(driver)和工具包(toolkit)。

再次,转换到控制台,命令为Ctrl+Alt+F1/F2/F3/F4,关掉gdm,命令为:sudo /etc/init.d/gdm stop,要确定已经关闭,否则在安装时会提示你有x server程序在运行。
   
再次,进入driver和toolkit目录,执行安装命令,为了方便,请一定按照默认安装。
   
然后,打开个人目录下的.bashrc文件或者/etc/profile文件,在其中加入命令:PATH=${PATH}:/usr/local/cuda/bin/   export PATH,LD_LIBRARY_PATH=${LD_LIBRARY_PATH}:/usr/local/cuda/lib  export LD_LIBRARY_PATH,保存后执行source .bashrc或者/etc/profile,依据你添加PATH和LD_LIBRARY_PATH时修改了那个文件确定。
   
最后执行nvcc命令,看看,如果提示你没有输入文件,就安装完成了。
   
如果你要安装SDK的话,还得安装一些包,依据make时的提示,google和新力得应该可以搞定一切,现在你可以享受CUDA了!






3. CUDA架构

本来一直都想把自己对CUDA架构的理解写出来,但是一方面自己是个懒人,不太愿意动笔,另一方面又感觉不太好组织语言,今天终于借着这个第四期有奖讨论的机会,把我对CUDA 架构的理解写下来吧!
    
至于题目中提到的有人将CUDA架构和x86的一样理解,这一定是错误的,如果说把GPU的架构和x86架构类比,这还有点道理。
        
要理解CUDA架构,个人认为应当先理解CUDA,CUDA是compute unit device architecture的简称,它统一了NVIDIA用于通用计算的GPU的编程模式,同时引入的共享存储器也是一大特色,大大提高了速度。有时我们说CUDA是指软件,有时又是指硬件,更多时是指CUDA C语言,因此本文将分三个部分简略的说明。
        
从软件的角度上说,CUDA软件栈包含两个层次,一个是驱动层的API,这类函数以cu开关,一个是运行层的API,以cuda开头,运行层API是建立在驱动层API之上的,是对驱动层API的封装,这种封装方便了使用,但是同时也牺牲了灵活性。一般而言,我们的开发都是优先使用运行时API。要提到的是:NVIDIA自身也提供了一些建立在这两者之上的库,供大家使用。

从另外一个方面说,CUDA的软件架构也应当包含这一系列的语言支持。当然这只是个人观点,欢迎大家讨论。
        
从硬件的角度看,CUDA架构应当包含TPC-->SM-->SP结构及它们的关系。各种存储器及其组织,以及硬件对线程的创建、执行、调度等,本人博客将会专门细说这个问题,所以也就不详细说了。
        
从CUDA C语言的角度看,CUDA的架构包括函数修饰符,变量修饰符,以及一些内置函数等,这方面内容本人的博客将会详细说明,因此,此处就不多说了。






4.CUDA C语言

前一节,我已经说了怎样在ubuntu linux上构建CUDA开发环境了,对一个语言来说,最简单的,也是用得最多的当然是它的语法了,下面我简单的介绍一下CUDA的语法。

CUDA C不是C语言,而是对C语言进行扩展。CUDA对C的扩展主要包括以下四个方面:

函数类型限定符,用来确定函数是在CPU还是在GPU上执行,以及这个函数是从CPU调用还是从GPU调用。

  • __device__,__device__表示从GPU上调用,在GPU上执行;

  • __global__,__global__表示在CPU上调用,在GPU上执行,也就是所谓的内核(kernel)函数;内核主要用来执行多线程调用。

  • __host__,__host__表明在CPU上调用,在CPU上执行,这是默认时的情况,也就是传统的C函数。CUDA支持__host__和__device__的联用,表示同时为主机和设备编译。此时这个函数不能出现多线程语句。



变量类型限定符,用来规定变量存储什么位置上。在传统的CPU程序上,这个任务由编译器承担。在CUDA中,不仅要使用主机端的内存,还要使用设备端的显存和GPU片上的寄存器、共享存储器和缓存。在CUDA存储器模型中,一共抽象出来了8种不同的存储器。复杂的存储器模型使得必须要使用限定符要说明变量的存储位置。
  • __device__,__device__表明声明的数据存放在显存中,所有的线程都可以访问,而且主机也可以通过运行时库访问;

  • __shared__,__shared__表示数据存放在共享存储器在,只有在所在的块内的线程可以访问,其它块内的线程不能访问;

  • __constant__,__constant__表明数据存放在常量存储器中,可以被所有的线程访问,也可以被主机通过运行时库访问;



texture,texture表明其绑定的数据可以被纹理缓存加速存取,其实数据本身的存放位置并没有改变,纹理是来源于图形学的一介概念,CUDA使用它的原因一部分在于支持图形处理,另一方面也可以利用它的一些特殊功能。

如果在GPU上执行的函数内部的变量没有限定符,那表示它存放在寄存器或者本地存储器中,在寄存器中的数据只归线程所有,其它线程不可见。

如果SM的寄存器用完,那么编译器就会将本应放到寄存器中的变量放到本地存储器中。

执行配置运算符<<< >>>,用来传递内核函数的执行参数。执行配置有四个参数,第一个参数声明网格的大小,第二个参数声明块的大小,第三个参数声明动态分配的共享存储器大小,默认为0,最后一个参数声明执行的流,默认为0。


五个内建变量,用于在运行时获得网格和块的尺寸及线程索引等信息
  • gridDim, gridDim是一个包含三个元素x,y,z的结构体,分别表示网格在x,y,z三个方向上的尺寸,虽然其有三维,但是目前只能使用二维;

  • blockDim, blockDim也是一个包含三个元素x,y,z的结构体,分别表示块在x,y,z三个方向上的尺寸,对应于执行配置中的第一个参数,对应于执行配置的第二个参数;

  • blockIdx, blockIdx也是一个包含三个元素x,y,z的结构体,分别表示当前线程所在块在网格中x,y,z三个方向上的索引;

  • threadIdx, threadIdx也是一个包含三个元素x,y,z的结构体,分别表示当前线程在其所在块中x,y,z三个方向上的索引;

  • warpSize,warpSize表明warp的尺寸,在计算能力为1.0的设备中,这个值是24,在1.0以上的设备中,这个值是32。


其它的还有数学函数,原子函数,纹理读取、绑定函数,内建栅栏,内存fence函数等。一般而言,知道这些就应该能够写出CUDA程序了,当然要写好的话,必须知道很多其它的细节。





6. 一个例子

上一节,已经简单的说了一下CUDA C的基本语法;因而在本节,兄弟决定以一个例子为基础说明CUDA程序的基本组成部分,不过说实话兄弟选择的例子并不太好,这个例子就是采用积分法计算圆周率/π的值。其计算原理是:在[0,1]范围内积分1/(1+x*x)

首先,让我们看一下在CPU上的计算流程,其计算流程如下
  1. /*
  2. 串行计算PI的程序,基本思想为:将积分区间均分为num小块,将每小块的面积加起来。
  3. */ 

  4.    float cpuPI(int num){

  5.       float sum=0.0f;

  6.       float temp;

  7.       for(int i=0;i<num;i++){

  8.           temp=(i+0.5f)/num;

  9.       // printf("%f\n",temp);

  10.           sum+=4/(1+temp*temp);

  11.       // printf("%f\n",sum);

  12.       }
  13.       
  14.       return sum/num;

  15.    }
复制代码
很明显,我们可以将for循环分解,使用CUDA处理。

有一个问题就是:for内部对sum变量的更新是互斥的,而CUDA中并没有浮点原子函数,对于这个问题的解决方案是:将for循环内部的两个语句拆开,分成两个内核函数来做运算。内核函数英文名为kernel,就是一个能够在GPU上运算的模块。第一个内核计算计算每个小积分块面积,并将每个block内所有线程对应的积分块面积加起来,存入全局存储器;第二个内核将前一个内核存入的数据加起来。下面是kernel代码:

  1. /*
  2. 在GPU上计算PI的程序,要求块数和块内线程数都是2的幂
  3. 前一部分为计算block内归约,最后大小为块数
  4. 后一部分为单个block归约,最后存储到*pi中。
  5. */

  6. /*
  7. 在GPU上计算PI的程序,要求块数和块内线程数都是2的幂
  8. 前一部分为计算block内归约,最后大小为块数
  9. 后一部分为单个block归约,最后存储到*pi中。
  10. */

  11.    __global__ void reducePI1(float *d_sum,int num){

  12.    int id=blockIdx.x*blockDim.x+threadIdx.x;//线程索引

  13.    int gid=id;

  14.    float temp;

  15.    extern float __shared__ s_pi[];//动态分配,长度为block线程数

  16.    s_pi[threadIdx.x]=0.0f;

  17.    
  18.    while(gid<num){

  19.       temp=(gid+0.5f)/num;//当前x值

  20.       s_pi[threadIdx.x]+=4.0f/(1+temp*temp);

  21.       gid+=blockDim.x*gridDim.x;

  22.    }


  23.    for(int i=(blockDim.x>>1);i>0;i>>=1){

  24.       if(threadIdx.x<i){

  25.           s_pi[threadIdx.x]+=s_pi[threadIdx.x+i];

  26.       }

  27.       __syncthreads();

  28.    }

  29.    if(threadIdx.x==0)

  30.    d_sum[blockIdx.x]=s_pi[0];

  31.    }

  32.    

  33.    __global__ void reducePI2(float *d_sum,int num,float *d_pi){

  34.    int id=threadIdx.x;

  35.    extern float __shared__ s_sum[];

  36.    s_sum[id]=d_sum[id];

  37.    __syncthreads();

  38.    for(int i=(blockDim.x>>1);i>0;i>>=1){

  39.       if(id<i)

  40.       s_sum[id]+=s_sum[id+i];

  41.       __syncthreads();

  42.    }

  43. // printf("%d,%f\n",id,s_sum[id]);

  44.    if(id==0){

  45.    *d_pi=s_sum[0]/num;

  46. // printf("%d,%f\n",id,*pi);

  47.    }

  48.    }
复制代码
其中__syncthreads()是CUDA的内置命令,其作用是保证block内的所有线程都已经运行到调用__syncthreads()的位置。
   
由上面的代码可以看出,使用使用CUDA的主要阻碍在于数据相关性。
   
一般而言,CUDA程序的基本模式是:
  • 分配内存空间和显存空间
  • 初始化内存空间
  • 将要计算的数据从内存上复制到显存上
  • 执行kernel计算
  • 将计算后显存上的数据复制到内存上
  • 处理复制到内存上的数据


这个程序在我的机器(CPU 2.0GHZ,GPU GTX295)上的加速比超过100,不知道在你们的机器上能够加速多少?






7. CUDA编程模式

CUDA支持大量的线程级并行(Thread Level Parallel),并在硬件中动态地创建、调度和执行这些线程,在CPU中,这些操作是重量级的,但是在CUDA中,这些操作是轻量级的。CUDA编程模型将CPU作为主机(Host),而将GPU做为协处理器(Coprocessor),或者设备(Device),以CPU来控制程序整体的串行逻辑和任务调度,而让GPU来运行一些能够被高度线程化的数据并行部分。即让GPU与CPU协同工作,更确切的说是CPU控制GPU工作。GPU只有在计算高度数据并行任务时才发挥作用。

一般而言,CUDA并行程序包括串行计算部分和并行计算部分,并行计算部分称之为内核(Kernel),内核只是一个在GPU上执行的数据并行代码段。理想情况下,串行代码的作用应该只是清理上个内核函数,并启动下一个内核函数,但由于目前的GPU的功能仍然十分有限,串行部分的工作量仍然十分可观







8. CUDA线程层次

GPU线程以网格(grid)的方式组织,而每个网格中又包含若干个线程块,在G80/GT200系列中,每一个线程块最多可包含512/768/1024个线程,Fermi架构中每个线程块支持高达1536个线程。同一线程块中的众多线程拥有相同的指令地址,不仅能够并行执行,而且能够通过共享存储器(Shared memory)和栅栏(barrier)实现块内通信。这样,同一网格内的不同块之间存在不需要通信的粗粒度并行,而一个块内的线程之间又形成了允许通信的细粒度并行。这些就是CUDA的关键特性:线程按照粗粒度的线程块和细粒度的线程两个层次进行组织、在细粒度并行的层次通过共享存储器和栅栏同步实现通信,这就是CUDA的双层线程模型。

在执行时,GPU的任务分配单元(global block scheduler)将网格分配到GPU芯片上。启动CUDA 内核时,需要将网格信息从CPU传输到GPU。任务分配单元根据这些信息将块分配到SM上。任务分配单元使用的是轮询策略:轮询查看SM是否还有足够的资源来执行新的块,如果有则给SM分配一个新的块,如果没有则查看下一个SM。决定能否分配的因素有:每个块使用的共享存储器数量,每个块使用的寄存器数量,以及其它的一些限制条件。任务分配单元在SM的任务分配中保持平衡,但是程序员可以通过更改块内线程数,每个线程使用的寄存器数和共享存储器数来隐式的控制,从而保证SM之间的任务均衡。任务以这种方式划分能够使程序获得了可扩展性:由于每个子问题都能在任意一个SM上运行,CUDA程序在核心数量不同的处理器上都能正常运行,这样就隐藏了硬件差异。

对于程序员来说,他们需要将任务划分为互不相干的粗粒度子问题(最好是易并行计算),再将每个子问题划分为能够使用线程处理的问题。

同一线程块中的线程开始于相同的指令地址,理论上能够以不同的分支执行。但实际上,在块内的分支因为SM构架的原因被大大限制了。内核函数实质上是以块为单位执行的。同一线程块中的线程需要SM中的共享存储器共享数据,因此它们必须在同一个SM中发射。线程块中的每一个线程被发射到一个SP上。

任务分配单元可以为每个SM分配最多8个块。而SM中的线程调度单元又将分配到的块进行细分,将其中的线程组织成更小的结构,称为线程束(warp)。在CUDA中,warp对程序员来说是透明的,它的大小可能会随着硬件的发展发生变化,在当前版本的CUDA中,每个warp是由32个线程组成的。SM中一条指令的延迟最小为4个指令周期。8个SP采用了发射一次指令,执行4次的流水线结构。所以由32个线程组成的Warp是CUDA程序执行的最小单位,并且同一个warp是严格串行的,因此在warp内是无须同步的。在一个SM中可能同时有来自不同块的warp。当一个块中的warp在进行访存或者同步等高延迟操作时,另一个块可以占用SM中的计算资源。这样,在SM内就实现了简单的乱序执行。不同块之间的执行没有顺序,完全并行。无论是在一次只能处理一个线程块的GPU上,还是在一次能处理数十乃至上百个线程块的GPU上,这一模型都能很好的适用。

目前,某一时刻只能有一个内核函数正在执行,但是在Fermi架构中,这一限制已被解除。如果在一个内核访问数据时,另一个内核能够进行计算,则可以有效的提高设备的利用率。







9. CUDA存储器组织

CUDA的存储器由一系列不同的地址空间组成。其中,shared memory和register位于GPU片内,Texture memory和Constant memory可以由GPU片内缓存加速对片外显存的访问,而Local memory和Device memory位于GPU片外的显存中。

最靠近流处理器的是寄存器文件(register file),每个寄存器文件是32bit。对线程来说,寄存器都是私有的,不允许其它线程染指。由于更靠近流处理器,寄存器具有最快的速度,GT200的每个SM拥有64KB的寄存器文件(Register Files),故一个块内最多可分配16K个寄存器,而G80中每个SM只有32KB,故一个块最多可分配8K个寄存器。最新加入的64bit数据类型(双精度浮点和64位整数型)将占用两个相邻的寄存器单元。CUDA的运行环境能够动态的为线程块分配寄存器,而每个线程块中的线程占用的寄存器大小则是静态分配的,在线程块寿命期间都不会更改,因此一个线程占用的寄存器数目是它在运行时占用的最大数目。如果寄存器被消耗完,数据将被存储在本地存储器(local memory)。对每个线程来说,本地存储器也是私有的,但是本地存储器是显存中的一个分区,速度很慢,而且使用本地存储器过多的话,程序也会终止,因此,编程时要尽量保证不能将数据放到本地存储器中,这可以通过修改块大小,使用共享存储器等方法来解决。

共享存储器是可以被同一块中的所有线程访问的可读写存储器,它的生存期就是块的生命期。在没有冲突的情况下,访问共享存储器几乎与访问寄存器一样快,是实现线程间通信的最好方法。共享存储器可以实现许多不同的功能,如用于保存共用的计数器或者块内的公用结果(例如reduction)。在同一个块内,所有的线程都能够读写共享存储器中的数据,相比于AMD的显卡来说,共享存储器是NVIDIA显卡的一项特色。一般而言,在kernel运行时,要先将数据从全局存储器写入共享存储器;计算完成后要将共享存储器中的结果转存入全局存储器。

Tesla的每个SM拥有16KB共享存储器,用于同一个线程块内的线程间通信。为了使一个half-warp内的线程能够在一个内核周期中并行访问,共享存储器被组织成16个bank,每个bank拥有32bit的宽度,故每个bank可保存256个整形或单精度浮点数,或者说目前的bank组织成了256行16列的矩阵。如果一个half-warp中有一部分线程访问属于同一bank的数据,则会产生bank conflict,降低访存效率,在冲突最严重的情况下,速度会比全局显存还慢,但是如果half-warp的线程访问同一地址的时候,会产生一次广播,其速度反而没有下降。在不发生bank conflict时,访问共享存储器的速度与寄存器相同。在不同的块之间,共享存储器是毫不相关的。

在实现中,GPU要把显存中的数据写到共享存储器中,必须先把数据写到寄存器里,再转移到共享存储器中,在编程时,这是隐式实现的。所以如果没有块内的数据共享千万不能用共享存储器,否则会降低速度。但是如果由于寄存器使用过量,那么我们可以使用共享存储器来当寄存器使用,此时比纯使用寄存器慢一点,但是是值得的。

Tesla能够在共享存储器内进行高速的原子操作。这里的原子操作是指保证每个线程能够独占的访问存储器,即只有当一个线程完成对存储器的某个位置的操作以后,其他线程才能访问这一位置。G80只支持对global memory的原子操作。访问global memory需要很长的访存延迟(长达数百个时钟周期),性能很低。在GT200及以后的GPU上,可以支持对shared memory中的原子操作指令(其中包括CAS指令,并且支持64位)。但是,CUDA并不提供对浮点数的原子操作(只有一个赋值的浮点原子指令),而在科学计算中,浮点数的使用远比整数要多,而且在Fermi的特性列表中,也没有看到加入浮点原子指令的信息,这不能不说是一个遗憾。

除此以外,多处理器上,还有两种只读的存储器:常数存储器(constant memory)和纹理存储器(texture memory),它们是利用GPU用于图形计算的专用单元实现的。常数存储器空间较小(只有64KB),属于片外存储器,其速度比shared要慢,但是它具有缓存,并且无须考虑冲突问题,主要用来加速对常数的访问。

从物理上说,纹理存储器不是存储器,它只是利用了纹理缓存而已。纹理缓存与CPU的缓存有很大的不同。首先,CPU的缓存往往是一维的,因为大多数的架构中的存储器地址是线性的。当访问一个只有4-8Byte的数据字时,会取出一个缓存单元中所有的64B数据。根据局部性原理,CPU处理的数据往往有很强的时空相关性,因此多取出的相邻的数据极有可能会被用到。CPU处理的数据只有一维,因而其缓存也只是在一个维度上是连续的;GPU需要处理的纹理则是连续的二维图像,因此纹理缓存也必须是在两个维度上连续分布的。典型的存储器控制器会将二维的纹理存储器空间映射为一维。其次,纹理缓存是只读的,也不满足数据一致性。当纹理被修改以后,必须更新整个纹理缓存,而不是纹理缓存中被修改的一小部分。第三,纹理缓存的主要功能是为了节省带宽和功耗,而CPU的缓存则是为了实现较低的延迟。第四,纹理可以实现对数据的特殊处理,比如怎样处理越界数据,自动实现插值等。

最后是全局存储器(global memory),使用的是普通的显存。整个网格中的任意线程都能读写全局存储器的任意位置。目前对Global memory的访问没有缓存,因此显存的性能对GPU至关重要。为了能够高效的访问显存,读取和存储必须对齐,宽度为4Byte。如果没有正确的对齐,读写将被编译器拆分为多次操作,极大的影响效率。此外,多个half-warp的读写操作如果能够满足合并访问(coalesced access),那么多次访存操作会被合并成一次完成,从而提高访问效率。

G80的合并访存条件十分严格。首先,访存的开始地址必须对齐:16x32bit的合并必须对齐到64Byte(即访存起始地址必须是64Byte的整数倍);16x64bit的合并访存起始必须对齐到128Byte;16x128bit合并访存的起始地址必须对齐到128Byte,但是必须横跨连续的两个128Byte区域。其次,只有当第K个线程访问的就是第K个数据字时,才能实现合并访问,否则half warp中的16个访存指令就会被发射成16次单独的访存。

GT200不仅放宽了合并访问条件,而且还能支持对8bit和16bit数据字的合并访问(分别使用32Byte和64Byte传输)。在一次合并传输的数据中,并不要求线程编号和访问的数据字编号相同。其次,当访问128Byte数据时如果地址没有对齐到128Byte,在G80中会产生16次访存指令发射,而在GT200中只会产生两次合并访存。而且,这两次合并访存并不是两次128Byte的。例如,一次128Byte访存中有32Byte在一个区域中,另外一个区域中有96Byte,那么只会产生一次32Byte合并访存(对有32Byte数据的区域)和一次128Byte(对有96Byte数据的区域)。

除了device端存储器外,还有存在于host端的存储器,即内存。在CUDA中,主机端内存分为两种:Pageable host memory和Page-locked host memory,其中Page-locked host memory保证位在于物理内存中,并且能够通过DMA加速与显卡的通信,提高数据传输速度,但是如果主机的内存不够用的话,会减弱系统的性能,但是一般不会出现这种情况。







10.CUDA执行模式

Tesla架构的构建以一个可伸缩的流多处理器(SM)阵列为中心。当主机CPU上的CUDA程序调用内核网格时,网格的块将被枚举并分发到多处理器上。一个线程块的线程在一个多处理器上并发执行。在线程块终止时,调度单元将决定是否启动新块和启动那一个块。

为了管理运行各种不同程序的数百个线程,多处理器利用了一种称为SIMT(单指令、多线程)的新架构。多处理器会将各线程映射到一个标量处理器核心,各标量线程使用自己的指令地址和寄存器状态独立执行。多处理器SIMT单元以warp为单位来创建、管理、调度和执行线程,构成warp的各个线程在同一个程序地址一起启动,严格串行。

为一个多处理器指定了一个或多个要执行的线程块时,它会将其分成warp块,并由SIMT单元进行调度。将块分割为warp的方法总是相同的,每个warp都包含连续的线程,递增线程索引,第一个warp中包含全局线程过索引0-31。每发出一条指令时,SIMT单元都会选择一个已准备好执行的warp块,并将指令发送到该warp块的活动线程。Warp块每次执行一条通用指令,因此在warp块的全部32个线程执行同一条路径时,可达到最高效率。如果一个warp块的线程通过独立于数据的条件分支而分散,warp块将连续执行所使用的各分支路径,而禁用未在此路径上的线程,完成所有路径时,线程重新汇聚到同一执行路径下,其执行时间为各时间总和。分支仅在warp块内出现,不同的warp块总是独立执行的--无论它们执行的是通用的代码路径还是彼此无关的代码路径。

SIMT架构类似于SIMD(单指令流多数据流)向量组织方法,共同之处是使用单指令来控制多个处理元素。一项主要差别在于SIMD向量组织方法会向软件公开SIMD宽度,而SIMT指令指定单一线程的执行和分支行为。与SIMD向量机不同,SIMT允许程序员为独立、标量线程编写线程级的并行代码,还允许为协同线程编写数据并行代码。为了确保正确性,程序员可忽略SIMT行为,但通过维护很少需要使一个warp块内的线程分支的代码,即可实现显著的性能提升。

另外一个重要不同是SIMD中的向量中的元素相互之间可以自由通信,因为它们存在于相同的地址空间(例如,都在CPU的同一寄存器中),而SIMT中的每个线程的寄存器都是私有的,线程之间只能通过shared memory和同步机制进行通信。

在SIMT编程模型中如果需要控制单个线程的行为,必须使用分支,这会大大的降低效率。例如,如果一个warp中需要进行分支(即warp内的线程执行的指令指针指向不同的位置),性能将急剧的下降。如果一个warp内需要执行N个分支,那么SM就需要把每一个分支的指令发射到每一个SP上,再由SP根据线程的逻辑决定需不需要执行。这是一个串行过程,因此SIMT完成分支的时间是多个分支时间之和。
©️2020 CSDN 皮肤主题: 大白 设计师:CSDN官方博客 返回首页