【并行计算-CUDA开发】GPU 的硬体架构 <关于存储体的介绍比较好bank conflicts>

GPU 的硬体架构
 

这里我们会简单介绍，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。大致上如下图所示：

详细的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 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 memory

目前CUDA 装置中，每个multiprocessor 有16KB 的shared memory。Shared memory 分成16 个bank。如果同时每个thread 是存取不同的bank，就不会产生任何问题，存取shared memory 的速度和存取暂存器相同。不过，如果同时有两个（或更多个） threads 存取同一个bank 的资料，就会发生bank conflict，这些threads 就必须照顺序去存取，而无法同时存取shared memory 了。

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]; 
    ... 
    int number = data[16 * tid];

会造成严重的bank conflict，为了避免这个问题，可以把资料的排列方式稍加修改，把存取方式改成：

    int row = tid / 16; 
    int column = tid % 16; 
    data[row * 17 + column] = global_data[tid]; 
    ... 
    int number = data[17 * tid];

这样就不会造成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) { 
        int number = data[tid]; 
    }

虽然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; }; 
    ... 
    __global__ void func(struct vec3d* data, float* output) 
    { 
        output[tid] = data[tid].x * data[tid].x + 
            data[ tid].y * data[tid].y + 
            data[tid].z * data[tid].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) 
    { 
        output[tid] = x[tid] * x[tid] + y[tid] * y[tid] + 
            z[tid] * z[tid]; 
    }

如果因为其它原因使资料结构无法这样调整，也可以考虑利用shared memory 在GPU 上做结构的调整。例如：

    __global__ void func(struct vec3d* data, float* output) 
    { 
        __shared__ float temp[THREAD_NUM * 3]; 
        const float* fdata = (float*) data; 
        temp[tid] = fdata[tid]; 
        temp[tid + THREAD_NUM] = fdata[tid + THREAD_NUM]; 
        temp[tid + THREAD_NUM*2] = fdata[tid + THREAD_NUM*2]; 
        __syncthreads(); 
        output[tid] = temp[tid*3] * temp[tid*3] + 
            temp [tid*3+1] * temp[tid*3+1] + 
            temp[tid*3+2] * temp[tid*3+2]; 
    }

在上面的例子中，我们先用连续的方式，把资料从global memory 读到shared memory。由于shared memory 不需要担心存取顺序（但要注意bank conflict 问题，参照前一节），所以可以避开non-coalesced 读取的问题。

Texture

CUDA 支援texture。在CUDA 的kernel 程式中，可以利用显示晶片的texture 单元，读取texture 的资料。使用texture 和global memory 最大的差别在于texture 只能读取，不能写入，而且显示晶片上有一定大小的texture cache。因此，读取texture 的时候，不需要符合coalesced 的规则，也可以达到不错的效率。此外，读取texture 时，也可以利用显示晶片中的texture filtering 功能（例如bilinear filtering），也可以快速转换资料型态，例如可以直接将32 bits RGBA 的资料转换成四个32 bits 浮点数。

显示晶片上的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 ²¹⁶ < y < 2 ²¹⁸时会得到错误的结果。

此外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 locked 的记忆体。不过，如果要求太大量的page locked 的记忆体，将会影响到作业系统对记忆体的管理，可能会减低系统的效率。