CUDA

CUDA-杂记待归类

  1. 二维数组在存储时会被扁平化处理，所以在程序中，最好是将多为数组扁平化为等价的一维数组。CUDA对于多为数组也是按行优先存储。

  2. 算法与硬件越是契合，算法的执行效率，硬件的使用率越高。所以有两个选择

    • 可以为当前算法设计特殊的处理架构，DSA（Domain Specific Architecture）如Google TPU。
    • 根据当前硬件精心设计算法。

  3. 向量数据类型

    同时适用于Host和Device，通过make_<type name>来构造，如

    int2 i2 = make_int2(3,4)：i2向量含有3和4两个元素。
    float4 f4 = make_float4(1.0f, 2.0f, 3.0f, 4.0f)：f4是含有4个元素的数组。

    访问方式如下：

    int x = i2.x; int y = i2.y;

  4. CUDA程序调试和开发

    • 可以使用ssh登路远程含有CUDA enabled GPU的服务器。
    • 通常使用双GPU的系统开发CUDA程序，一个GPU负责显示，另一个负责计算，可以使用Nsight等工具。
    • 只有一个GPU时，在Linux系统中可以关闭桌面环境（释放桌面环境对GPU的占用），只在命令行中使用CUDA-gdb调试。（实际上，实验阶段不关闭桌面环境，也是可以正确执行的）

  5. CUDA开发的任务

    有效的数据并行算法 + 针对GPU架构特性的优化 = 获得并行最优性能

  6. OpenCL

    OpenCL使用起来繁琐，而且运行速度远远低于CUDA运行速度。OpenCL与CUDA的主要功能有着十分相似之处，一个CUDA程序员很容易掌握OpenCL编程。

  7. half-warp

    截图

  8. Streams-流

    任务并行（Task Parallelism），不同于在大量的数据上执行相同的任务（SIMD），而是同时执行多个不同的任务。

  9. CUDA数据并行原语

    啥是原语，就是这个领域的基本操作。CUDA并行原语库（CUDA Data Parallel Primitive Library, CUDPP）. 含有并行前缀和，并行排序，并行规约等。这些原语为许多数据并行算法提供了重要基础。如果你正在编写某个复杂算法，那么CUDPP很可能已经提供了这个算法。

    CUDPP 下载地址：http://code.google.com/p/cudpp

    for more information：CUDA By Example 178页

  1. 想清楚一个问题

    如果使用shared memory，每个block对应自己的shared memory，当这个block中threads的ID更新后，这个block并没有再被分配新的shared memory. 也就是说如果使用两个blocks, 两个blocks有两段shared memory, block中threads 的ID更新后, 所计算结果会写入这个thread所在block的shared memory中. thread ID变了, 但所属的block不变.

  1. 理解SIMD，SIMT

    Single Instruction Multiple Data（Thread）.

    • Single：相同的操作，kernel函数只有一个
    • Instruction：kernel函数所做的事情
    • Multiple：所处理的数据量大，要拆分为一批一批
    • Data：大的数据量
    • Threads：一个SP上的大量thread超快速切换，获得延时隐藏

  2. warp中的divergence

    已经知道，在一个warp中，所有的threads执行相同的指令，但是如果指令中含有条件分支语句，很大程度上会发生divergence。比如优化并行归约 中的描述：一个宿舍的6个学生可以是一个warp，今天有的想先上厕所，后吃饭，而有的不需要上厕所，此时所有的同学都会一起先上厕所，后一起吃饭。

    总之，分支发散使得性能明显下降。但是，注意divergence只发生在一个warp中。

    可以综合算法的上下文，将divergence的粒度变为32（warp的大小）的倍数，从而避免warp内的分支发散。比如：

    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    __global__ void kernel(float* c){
        int tid = threadIdx.x + blockDim.x * blockIdx.x;
        float a=0;
        float b=0;
        if (tid%2 == 0){  
            c[tid] = 100;
        }else{
            c[tid] = 200;
        }
    }

    偶数id的thread把100写入偶数编号的地址，奇数id的thread将200写入奇数编号的地址。会发生分支发散。如果将分支的粒度定为32，则没有了divergence [代码中tid%32 == 0有问题]：

    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    __global__ void kernel(float* c){
        int tid = threadIdx.x + blockDim.x * blockIdx.x;
        float a=0;
        float b=0;
        if ( tid%32 == 0 ){
            c[tid] = 100;
        }else{
            c[tid] = 200;
        }
    }

    但是两者的结果的是不同的，实际中，需要根据算法上下文考虑结合此方法。

    使用brand_efficiency指标来查看divergence的情况：

    $ nvprof --mereics brand_efficiency ./out, 但是CUDA编译器会进行优化，将短的，有条件的代码段的断定指令取代了分支指令。所以，会看到虽然代码中有分支，却显示分支效率100%。

  1. 数据预读取?

    在一次读取global memory的操作和使用这个数据之间，插入独立于以上数据的操作，可以隐藏访问延迟。如：

    1
    2
    3
    float m = dev_a[i];  // 1. 从global memory中读取
    float f = a * b;     // 2. 与m无关的操作 
    float f2 = m*f;      // 3. 使用m

    分析：在warp切换中，

  2. 指令优化

    GPU中执行指令时很快速的，所以通常不用太在意指令的优化。优化顺序一般是存储优化，后执行配置优化，最后可以考虑指令优化。

    比如：除以2^n，使用移位操作>>n，以2^n求模，使用&(2^n - 1)；避免从double到float的自动转换，float a = 0.0，使用float a = 0.0f。

    还比如，两种运行时数学库函数的取舍，精度高的速度低，精度低的速度高。使用-use-fast-math编译选项后，强制将速度慢的func()转化为速度快的__func()。

    还比如循环展开。