CUDA

1. Tutorial
   1. 使用 global 可以将一个普通函数转换成 kernel 函数；
      1. 猜测此处的 global 可能是一个宏；
   2. 使用 cudaMallocManaged 函数，可以实现在 GPU 上分配内存
      1. 示例
         1. float *x;
         2. cudaMallocManaged(&x, N*sizeof(float));
   3. 使用 cudaFree 可释放分配的 GPU 内存
      1. cudaFree(x);
   4. 使用三个尖括号表示对 kernel 函数的调用
      1. add<<<1, 1>>>(N, x, y);
      2. 它会创建一个 GPU 线程，来执行所调用的核函数；
      3. 由于 GPU 线程的执行，不会中断 CPU 线程，有点像是异步的；
      4. 因此如果 CPU 要使用 GPU 的计算结果，则需要显式的让 CPU 线程进行等待；
   5. 使用 cudaDeviceSynchronize() 函数，让 CPU 线程进入等待；
   6. 使用 GPU 加速的文件，后缀名需要更改为 .cu；
   7. 使用 nvcc 命令调用 CUDA C++ 编译器，对代码进行编译；
   8. 调用核函数所用的尖括号内部，原来是用来做并行计算的选项配置的；
      1. 尖括号中的第二个数字表示一个线程块中的线程数量；块的基本单位为32；因此块中的线程数量需要是32的倍数；（越高级的显卡，块中能够并行的线程数量越多，例如 P100 显卡可以达到2048个线程）；
      2. 对于 <<<1, 256>>> 它表示在1个线程周期中，使用 256个线程来做并行计算；但是，每个线程只会做一次计算；而不是把所有的计算分配到所有线程中；
      3. CUDA 可以让核函数获得线程在其所在块中的下标；
         1. threadIdx.x 存放线程在其所在块中的下标，假设命名为 index；
         2. blockDim.x 存放块中的线程数量，假设命名为 stride；
         3. 通过将循环条件设置为 for (int i = index; i < n; i += stride) ，将计算分配到每个线程中；
      4. 原理：GPU 有很多个处理器，组合成所谓的 SM（streaming Multiprocessors）流处理器；每个流处理器可以运行多个并行线程块；例如 P100显卡总共有56个流处理器，每个流处理器可以处理 2048 个线程（如果每个块按256个线程，算下来，貌似每个流处理器可以处理 64个块 block）；为了能够充分利用这些线程，对核函数的调用应该转换化多个线程块进行调用；
      5. 尖括号的第一个数字，表示线程块的数量；并行线程的块组成 grid；由于我们有 N 个元素需要处理，每个线程块有256个线程，我们只需要计算出计算 N 个元素所需要的块数量，即使用 N 除以块大小（如果 N 不能整除，需要小心）；
         1. int blockSize = 256;
         2. int numBlocks = (N + blockSize - 1) / blockSize;
         3. add<<<numBlocks, blockSize>>>(N, x, y);
      6. blockIdx.x 存放块的下标；
      7. gridDim.x 存放总共的块数量；
      8. 今天重新看了官方的 CUDA 编程文档，发现 block 只是一种方便程序员的抽象；通过 block，让线程的组织看起来更符合某种规律，有时候会更加直观；但其实背后的本质仍然是通过计算指针的位置，来实现对目标位置的取值、计算和赋值；
2. programming guide
   1. 介绍
      1. GPU 原理：CPU 给控制单元和寄存器预留了很多空间，用来处理有很多控制条件类型的运算；GPU 是相反，它的控制单元和寄存器很小，但 ALU 计算单元很多，这特别适合用来计算控制条件少，但有大量相同类型计算任务的场景；
      2. CUDA 有三个重要的抽象：线程组合层级、共享内存、屏障同步；通过这些抽象机制，让开发者能够将应用程序中的单个计算任务拆解成多个并行计算的小任务的方式；只要运行环境可以获取到流处理器的数量，代码就可以根据环境变量，自动调整，以最大化利用多个流处理器进行并行计算；
   2. 编程模型
      1. CUDA 通过将普通函数包装成核函数，将计算工作分发到不同的线程中并行处理，从而提高了计算速度；
      2. 普通函数包装成核函数后，使用 <<<>>> 两对尖括号进行调用，尖括号里面填写块数量和块线程数量两个参数；每个线程在其所有的块中，有一个下标；在包装普通函数时，需要将这个下标引入代码中，让代码中的数值计算，对应每个下标值（也即对应的线程）；
         1. <<<>>> 接受两种类型的参数，分别是 int 和 dim3 类型；
      3. 事实上，包装成核函数后，除了原来会传入的那些实参；默认还会传入另外几个参数，包括线程下标 ThreadIdx，块下标 blockIdx，块尺寸 blockDimx；这些信息可以用来改造原有的函数，使其计算工作匹配到单个线程中；
      4. 块可以有3种维度，分别可以用来对应三种场景的计算，分别是 vector, matrix, volume；
         1. 但实际上，单个块内部的线程数量由硬件决定的，因此是固定的；区别在于单个块内的这些线程，貌似会共享内存？答：是的，当形参是一个多维的时候，一次传入的实参值，刚好可以实现在一次计算中共同，不需要重复传入相同的实参；
      5. 块内部的线程，正常是并行进行的，每个线程之间没有顺序依赖关系；如果某个计算步骤有顺序的要求，需要所有线程将计算结果算出来后，才能继续向下进行，则此时可以通过 __syncthreads 函数来实现；它会让所有线程进入等待的状态，直到所有线程执行完毕，才会继续；
         1. 在 Cooperative Groups API 中，据说有更多的接口，可以用来实现线程同步的需求；
         2. 另外通过 shared memory 共享内存的接口，可以实现更快的性能；
      6. 内存层级
         1. per-thread local memory：线程内部私有；
         2. per-block shared memory：块内共享，仅块内的线程可以访问；
         3. global memory：全局共享，所有块都能够访问；
      7. 异构编程
         1. 编程模型假设 GPU 部分的代码和 CPU 部分的代码是独立进行的，二者没有依赖关系；并假设它们拥有各自独立的 DRAM 动态内存；
         2. 通过 unified memory 模块，可以实现两种内存之间的桥梁；
      8. GPU 的版本号由 Major 和 Minor 两部分组成，其中的大版本号表示 GPU 的架构类型，而架构类型决定了 GPU 的计算能力；小版本号则用来表示 GPU 的局部更新；
   3. 编程接口
      1. CUDA 对 C 语言的语法进行了小拓展，并提供了一个运行时库；
         1. 有一节专门讲所有的 C 语言语法拓展；当使用这些拓展时，需要使用专门的 nvcc 编辑器来编译源代码；
         2. 运行时库提供了一些函数，可以用来在 host 分配和释放内存，在 host 和 device 的内存之间转移数据，管理多个 device 等；完整功能在 CUDA 手册中有介绍；
         3. 运行时库是比 CUDA 驱动 API 的更高层次的封装，它可以让编程更容易，不用关心底层的驱动细节；
      2. CUDA 指令集称为 PTX；但一般通过 C 来编程会更直观；
      3. nvcc 工作流程
         1. 先将 device 源代码编译成 PTX 指令集的汇编格式，或者二进制 cubin 对象格式（它在执行期间还会进一步做 JIT 编译）；
         2. 将 host 代码中，标记有 <<<>>> 两个尖括号的地方，替换为运行时库函数，这些函数在执行时，会加载和调用第一步编译完成的 device 目标代码；
         3. 最后替换完成的源代码，交给 host 编译器完成剩下的编译工作；
         4. 引入 PTX 的机制，相当于增加了一层抽象，通过这层抽象，使得代码能够在多种不同规格和架构的 GPU 之间兼容，不会受到硬件实现细节的影响；
      4. 运行时库
         1. 通过 cudart 库实现，应用一般通过链接静态库 cudart.lib 或者 libcudart.a，或者 cudart.dll 或 libcudart.so 来调用；
         2. 模块介绍
            1. device memory 用来管理 GPU 内存；
            2. shared memory 用来设置共享内存；
            3. Graphics Interoperability 提供了多个函数来与 OpenGL 和 Direct3D 进行交互；
      5. 初始化：没有显式的初始化函数，来做初始化的动作，而是当第一次执行 device code 的时候，才会触发初始化；初始化的时候，会为每个 GPU 设备生成一个上下文环境，这个环境可以被所有 host 的应用共享；
         1. 如果显式调用 cudaDeviceReset()函数，则会销毁这个环境，然后等到下一次调用 device code 的时候，才会再次初始化一个新环境；
      6. 显存
         1. cudaMalloc 函数用来分配显存；
         2. cudaFree 用来释放显存；
         3. cudaMemcpy 用来在主存和显存之间传输数据；
         4. cudaMallocPitch 和 cudaMalloc3D 可用来分配二维和三维数组的显存，它们可以确保满足对齐要求；
            1. cudaMallocPitch 和 cudaMalloc3D 的参数格式有点奇怪，待查询下手册；
         5. cudaMemcpy2D 和 cudaMemcpy3D 则专用来复制内存中的二维和三维数据；
         6. cudaMemcpyToSymbol
         7. cudaMemcpyFromSymbol
         8. constant, device 用来访问全局变量；
      7. 共享内存
         1. 复制类对象到显存的方法，先在显存中声明一个新的类对象，值数据成员，直接手工进行赋值初始化；指针指向的数据成员，使用 cudaMalloc 加 cudaMemcpy 来实现拷贝
            1. 不知为何，写到这里，我重新看了一下代码，我感觉对象仍然是声明和存放在主存中的，只是对象的指针成员，指向了显存好像
         2. 对于改造后的 kernel 函数，它被各个 GPU 线程调用的时候，代码都是一模一样的，唯一的区别在于计算 index 的环境变量不同，因此它使得每个线程会根据下标去获取实参的不同部分进行计算；
         3. 对于矩阵乘法运算，由于其特殊的性质，可以将大矩阵分成小矩阵进行计算，然后再累加结果；当分解成多个小矩阵时，每个小矩阵由一个单块进行运算，可以实现共享内存
         4. 原理：GPU 有多个流处理器，每个流处理器都有自己的寄存器，这个寄存器的存储容易不大，但是它的读取速度很快，因此，如果能够将单个流处理器中，多个线程所需要用到的数据，都提前加载到各自的寄存器中，则可以减少对 GPU 显存的访问，大大提高了计算性能；
         5. 如何实现将数据加载到流处理器的寄存器，使得单个流处理器中的多个线程可以共享这些数据，从而不必要每个线程都对主显存发起访问？
            1. 答：使用 shared 关键字来声明变量，由于每个线程只拷贝初始化一个变量值，因此需要再做一下线程同步，才能使得共享内存全部完成初始化，之后再进入下一步使用共享数据进行计算
   4. 硬件实施
      1. 使用 deviceQuery 可以查询当前机器上面的显卡信息，包括流处理器数量，每个SM 的最大线程数，每个 block 支持的最大线程数等
   5. 性能优化
   6. 附录
3. 编译 nvcc
   1. 静态库编译步骤
      1. nvcc -rdc=true -c -o temp.o foo.cu
         1. 将 cu 文件编译成中间临时 object 文件 temp.o
         2. rdc 表示 relocatable device code，可重新定位的GPU代码的文件
         3. -rdc=true 表示生成的 GPU 代码文件不可执行，需要链接后才能执行；默认是 false ，表示直接生成可执行的GPU 代码文件；
            1. 怎么感觉这里跟 gcc 里面的 -fPIC 好像是一样的道理？
      2. nvcc -dlink -o foo.o temp.o -lcudart
         1. -dlink 表示将不可执行的 rdc 文件链接到可执行的GPU代码文件中
            1. 感觉这个选项很像 gcc 里面的 -ldl，即 dynamic link；
         2. -lcudart 表示链接 cudart 库；这个库的完整名称为 cuda_runtime
      3. ar rc libfoo.a foo.o temp.o
         1. 将两个 object 文件合并成一个库文件 libfoo
      4. ranlib libfoo.a
         1. ranlib 表示更新静态库的符号表索引
      5. g++ main.cpp -L. lfoo -o main -L/usr/local/cuda/lib64 -lcudart
         1. 编译主程序文件，需要分别链接静态库和动态库
   2. 编译选项
      1. –cuda (-cuda)
         1. 用来将 cu 文件编译成 C/C++ 格式的源文件，输出的文件可以被 host 的编译器进行处理；
      2. –device-c（-dc)
         1. 等同于原来 -rdc=true 和 –compile 的效果叠加；即将 cu 文件单独编译成 .o 对象文件；
      3. –device-link(-dlink)
         1. 将有 rdc 代码的 obj 文件和 ptx, cubin、fatbin 文件链接成一个新的 obj 文件，这个文件含有可执行代码，可传递给 host 链接器；
      4. –library(-l)
         1. 指定要在链接阶段要使用的库名（库名无须写后缀）
         2. 库的搜索路径由另外一个选项 –library-path 进行指定
      5. –library-path(-L)
         1. 指定要链接的库的搜索路径；
      6. –output-file (-o)
         1. 指定输出文件的名称
      7. –lib(-lib)
         1. 将输入的文件，编译成 .o 文件；同时还可以将结果输出到指定的库文件中；
   3. 编译流程
      1. 步骤
         1. 源文件先进行预处理，然后编译成 CUDA 二进制机器码或者 PTX 中间码；
         2. 源文件再次预处理，将上一步的 fatbinary 嵌入进去；将 CUDA 专用的 C++ 扩展转换成标准 C++ 结构，变成合成码；
         3. 主机编译器将合成码转成主机对象文件；
      2. 编译样例
         1. 生成可执行文件
            1. nvcc -arch=sm_50 –device-c a.cu b.cu
            2. nvcc -arch=sm_50 a.o b.o
               1. 注：-arch=sm_50，在官方文档中有提示这个选项的作用会决定最终可执行代码内容）
         2. 生成静态库
            1. nvcc -arch=sm_50 -dc a.cu b.cu
            2. nvcc –lib a.o b.o -o test.a
               1. 官方文档提示 device linker 仅支持处理静态库；
   4. 编译问题
      1. 当 host 文件调用了 device 文件中代码时，host 文件在链接时，都需要加入 device 文件，不然会找不到所调用的代码，出现如下报错：
         1. undefined symbol __fatbinwrap_name (原来如此，终于找到原因了）
         2. 晕，当时忘了写下如何添加 device 文件了；
            1. 详见 test/lib_test/demo8
            2. 需要在 g++ 编译时写下需要链接的库名称，例如：
               1. LIBPATH = /usr/local/cuda/lib64
               2. g++ main.o -lgpu -L. -lcudadevrt -lcudart -L$(LIBPATH)