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()

    还比如循环展开。