CUDA学习100天

我是在GPU世界论坛学习到的，吸收了一些知识，在这里作一下记录，以供以后温故而知新。
地址：https://bbs.gpuworld.cn/index.php?board=65.0;sort=last_post

CUDA软件架构

原文链接：https://blog.csdn.net/dcrmg/article/details/54867507

• CUDA的软件架构由网格（Grid）、线程块（Block）和线程（Thread）组成，相当于把GPU上的计算单元分为若干（2~3）个网格，每个网格内包含若干（65535）个线程块，每个线程块包含若干（512）个线程
  thread：一个CUDA的并行程序会被以许多个threads来执行。
  block：数个threads会被群组成一个block，同一个block中的threads可以同步，也可以通过shared memory通信。
  grid：多个blocks则会再构成grid。
  
• CUDA中可以创建的网格数量跟GPU的计算能力有关，可创建的Grid、Block和Thread的最大数量参看以下表格：
  

线程索引的计算公式

1、 grid划分成1维，block划分为1维

    int threadId = blockIdx.x *blockDim.x + threadIdx.x;  
    
  
2、 grid划分成1维，block划分为2维  

    int threadId = blockIdx.x * blockDim.x * blockDim.y+ threadIdx.y * blockDim.x + threadIdx.x;  
  
  
3、 grid划分成1维，block划分为3维  

    int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z  
                       + threadIdx.z * blockDim.y * blockDim.x  
                       + threadIdx.y * blockDim.x + threadIdx.x;  

  
4、 grid划分成2维，block划分为1维  

    int blockId = blockIdx.y * gridDim.x + blockIdx.x;  
    int threadId = blockId * blockDim.x + threadIdx.x;  
   
  
5、 grid划分成2维，block划分为2维 

    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int threadId = blockId * (blockDim.x * blockDim.y)  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;  
    
  
6、 grid划分成2维，block划分为3维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  
                       + (threadIdx.z * (blockDim.x * blockDim.y))  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;  
   
  
7、 grid划分成3维，block划分为1维 

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * blockDim.x + threadIdx.x;  
   
  
8、 grid划分成3维，block划分为2维  

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * (blockDim.x * blockDim.y)  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;  
   
  
9、 grid划分成3维，block划分为3维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  
                       + (threadIdx.z * (blockDim.x * blockDim.y))  
                       + (threadIdx.y * blockDim.x) + threadIdx.x; 

常用的GPU内存函数

cudaMalloc()

1. 函数原型： cudaError_t cudaMalloc (void **devPtr, size_t size)。
2. 函数用处：与C语言中的malloc函数一样，只是此函数在GPU的内存你分配内存。
3. 注意事项：
   3.1. 可以将cudaMalloc()分配的指针传递给在设备上执行的函数；
   3.2. 可以在设备代码中使用cudaMalloc()分配的指针进行设备内存读写操作；
   3.3. 可以将cudaMalloc()分配的指针传递给在主机上执行的函数；
   3.4. 不可以在主机代码中使用cudaMalloc()分配的指针进行主机内存读写操作（即不能进行解引用）。

cudaMemcpy()

1. 函数原型：cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)。
2. 函数作用：与c语言中的memcpy函数一样，只是此函数可以在主机内存和GPU内存之间互相拷贝数据。
3. 函数参数：cudaMemcpyKind kind表示数据拷贝方向，如果kind赋值为cudaMemcpyDeviceToHost表示数据从设备内存拷贝到主机内存。
4. 与C中的memcpy()一样，以同步方式执行，即当函数返回时，复制操作就已经完成了，并且在输出缓冲区中包含了复制进去的内容。
5. 相应的有个异步方式执行的函数cudaMemcpyAsync()，这个函数详解请看下面的流一节有关内容。

cudaFree()

1. 函数原型：cudaError_t cudaFree ( void* devPtr )。
2. 函数作用：与c语言中的free()函数一样，只是此函数释放的是cudaMalloc()分配的内存。
   下面实例用于解释上面三个函数

DAY5:阅读 CUDA C编程接口之CUDA C runtime

原文地址（点击）

• 所有的入口函数（也叫导出函数）都具有cuda前缀。（例如我们常说的cudaMemcpy就是这样的）。CUDA分成两部分，runtime api前缀都是cuda，driver api前缀都是cu（其他的扩展库具有更多其他前缀）。请注意driver api的前缀只有cuda的前两个字母（cu）。遇到cu开头就知道是Driver API的函数，而遇到cuda就知道是runtime api的。
• Call Stack一般总是翻译成“调用栈”，指函数调用时候保存返回地址之类的参数信息的。

DAY10:阅读CUDA异步并发执行中的Streams

• 回调函数不能调用任何CUDA API函数，无论是直接的，还是间接的调用。因为如果在回调函数中这样做了，调用CUDA函数的回调函数将自己等待自己，造成死锁。其实这很显然的，流中的下一个任务将需要等待流中的之前任务完成才能继续，因为CUDA Stream是顺序执行的， 而如果你一个流中的某回调函数，继续给某流发布了一个任务，很有可能该回调函数永远也等待不完下一个任务完成，因为等待下一个任务完成首先需要这回调函数先结束，而回调函数却在等待下一个任务完成…于是就死锁了。

DAY30：阅读CPU和GPU间的数据传输

• 这章节主要说了如何优化Host和Device间的数据传输
• 首先章节说, 应当尽量尝试能减少传输量就要减少.有的时候虽然这种改写可能很渣, 并行度不高, 但只要是合算的, 就应当使用这种GPU代码
• 在Linux下, 应当考虑使用unified memory
• 但在Windows下面不建议使用(因为Windows下面的Unified memory性能很低)，Linux下的unified memory可以看成是zero-copy这里的高级版本.进化过了
• 现在多了TK1/TX1/TX2这种设备, 他们的GPU也是集成的, 也没有独立的显存.在这种设备上应当考虑使用zero-copy/unified memory以便减少无辜的复制传输

DAY34：阅读算术指令

• 6.1的Pascal(例如GTX1080), 单精度性能是128指令每SM每周期, 但是double, 却只有4指令/SM/周期,也就是double性能对于6.1的卡只有1/32,你就知道6.1的1080这种卡, 不适合双精度运算.因为它的双精度性能只有单精度性能的3%左右(1/32)
• 而你会发现, 6.0的卡(GP100), 却具有32 / 64 = 50%的相对峰值单精度性能的双精度性能, 因此这卡适合双精度运算(目前最高的双精度就是50%性能–基准是用的单精度性能)
• 很多时候GPU和CPU的结果不同, 往往是GPU的结果更加正确(和手工计算的精确结果相比)，也就是, 目前N卡算的又快又好.而并非是很多人想象中的, CPU结果更正确.

DAY35：阅读控制流指令

• 在GPU上, 源代码中使用任何流程控制语句(if, while, for)这种, 如果使用的不当, 可能会影响性能.
• 流程控制语句导致最终运行的时候, 产生的分支和跳转发生在warp内部,在以前的阅读章节中, 你知道SIMT结构只是构造成了线程可以自由执行的假象, 而实际上它们是按照warp一组了执行的,任何在warp内部的分支都将严重的影响性能
• 一个能正确应用了这三点(自行翻看原文)的(或者其他点, 以后说)的CUDA C的Kernel,还是可以能在代码中大量的出现if, for, while的, 而不会影响性能.所以很多人认为我不能在GPU上使用if, for, while, 是错误的. 该用就用.用的好了没啥问题.

DAY36：阅读C语言

• 本章节说了"CUDA C的执行空间"扩展修饰符：__global__和__device__, 以及, 不常用的__host__
• 而CUDA C默认的Runtime API风格的编译, 允许这个函数依然保留在普通的源代码文件中, 只需要加上一点点处理, 就可以在GPU上运行了.这一点点处理就是:
  (1)原本的老CPU函数, 需要加上__global__修饰符前缀
  (2)原本直接对老CPU函数的调用, 需要改成<<<>>>语法
  (3)老版本以前能直接返回结果的, 现在必须是void了, 不能返回了. 但可以通过写入global memory, 然后通过cudaMemcpy回来.
• 今天本章节说的, 则只是第一点中的__global__这种扩展修饰符的用法, 而暂时不涉及其他
• 它们实际上不仅仅指定了有这两个前缀的函数将在GPU上执行,也同时指定了CUDA C编译器遇到这两个前缀后, 会将有这些特殊前缀的函数, 生成GPU代码,而其他源文件中的剩余部分, 没有这两个前缀的函数, CUDA C编译器自动跳过, 调用你本机上的原来的CPU编译器, 继续编译剩下的部分.
• 通过这种方式, 用户看来, 它只需要将源文件改成.cu扩展名, 将一些函数添加上__global__之类的扩展, 就自动能在GPU上执行了.不需要考虑如何将代码传递给GPU, 也不需要考虑如果通知GPU开始执行.很方便的.
• 本章节说, __global__前缀的函数, 编译完成后, 可以从Host端调用它(通过某种<<<>>>语法), 它从Device上运行.这就是我们常说的"启动一个kernel"的过程.请注意是CPU(host)代码中, 你要求启动Kernel, 而Kernel是在GPU(device)上运行的.也就是通过__global__前缀 + <<<>>>语法, 此时你将在异种设备上开始了代码执行, 或者说计算.俗称"异构计算"
• 如果你的CPU代码有子函数一样, 将所有的需要的代码放置在一个__global__开头的函数中, 可能会过于庞大复杂了,此时引入了__device__前缀, 你可以用它来写一些只能在GPU上运行的子函数,然后常规的能被调用的kernel(__global__开头的函数)可以调用这些小片段(以__device__开头)嵌入到自己内部, 这样被反复使用的一些代码可以被抽取出来, 做成小片段. 很方便的.因为被设计成__global__是被能跨越CPU/GPU边界调用的函数, 而__device__是被设计成只能在GPU上调用运行的函数,__device__虽然失去了从Host上调用它的能力,但却多了可以直接返回函数值的功能.
  (1)__device__前缀的函数只能从GPU上运行, 但可以更像正常的C函数一样的返回结果.
  (2)__global__前缀的能从Host上调用, 然后从GPU上运行; 但不能直接返回任何结果(可以通过其他变通手段)
• 从GPU上启动kernel和从CPU上启动kernel相比, 具有很多好处, 同时能进一步的降低通讯延迟(因为这个是GPU给GPU自己发布任务, 而不是遥远的跨PCI-E那端的CPU老大)
• 然后还有另外一种, 叫__host__修饰, 这种不常用.单一的__host__修饰等于没有修饰(常规的CPU函数)
• 可以指定__host__前缀和__device__前缀同时存在，可以同时只写一次, 想同时给CPU上的普通函数, 和GPU上的__global__的kernel用。这种写法等效于:编译器自动当成2个函数看, 一个只有__host__, 另外一个只有__device__，名字一样而已.所以这种能同时从CPU和GPU上调用.(因为编译了两次, 生成了两种代码)
• __CUDA_ARCH__宏, 这个很重要，可以让你的代码中, 有选择的根据计算能力不同, 来编译出来不同的代码.
• 请注意使用了inline"可能"会生成较好的代码(因为取消了函数间的流程转移, 参数传递, 结果返回之类的开销)。但也可能会造成性能下降(因为都将小函数内嵌到大函数中了, 可能会有多个副本存在, 增大了代价。是否inline和具体的实际代码有关.

DAY41：阅读同步功能

• 主要说__syncthreads*()家族的系列函数，和__syncwarp()。
• __syncthreads*()主要是在一个block的范围内进行控制（代码执行位置的同步）。 而__syncwarp()作为新增内容, 是CUDA 9引入为了适应新卡的, 范围则缩减为warp内部。
• __syncthreads*()可以完成局部同步, 也就是block内部的同步, 同时附带memory fence效果.往往用于在shared memory上的数据交换操作.
• 扩展的__syncthreads_count(), __syncthreads_and(), __synchtreads_or()。