526互联

0001笔记【并行计算】CUDA在现代C++中如何运用？看这一个就够了

发布时间 2023-04-18 17:18:13作者: 马角的逆袭

目录

SM（流多处理器）和板块（block)

一个板块会被调度到一个SM上，直到执行结束

  1. 一个block一旦被调度到某个SM上，就会一直执行，直到执行结束（gpu不存在时间片轮转），好处是不需要像cpu一样切换上下文，就也不需要保存寄存器和共享内存的开销
  2. 一个block里的一个线程，则是被调度到一个SP上执行

常用函数

cudaMalloc在显存上分配内存

cudaMallocHost在主存上分配锁页内存

  1. cudaMallocHost分配的是CPU内存，并且GPU可以直接读写，通过PCI-E总线传输（带宽约为6G/s）
  2. 注意要cudaDeviceSynchroinize()后再使用，不然数据可能错乱

cudaMemcpy在主存和显存之间拷贝数据

  1. 会隐含一次cpu和gpu同步（cudaDeviceSynchroinize())

cudaMallocManagerd统一内存

  1. 用cudaMallocManagerd分配的内存在cpu和gpu上都可以访问，
    因为显卡驱动会自动的把数据在主存和显存之间来回拷贝
  2. 需要手动同步一下，才能保证数据正确（主动写cudaDeviceSynchroinize())

优化

时间依赖和空间依赖

  1. 时间依赖：假设有代码int sum; for(i=0;i<100;i++) { sum += arr[i]; }其中sum就是有时间依赖的，第10次循环里的sum的结果依赖于第9次循环的sum，实际上就变成了串行
    如果编译器看到结果依赖于前一个值，就不会优化，导致效率不高

线程太多不行：防止寄存器打翻(register spill)

  1. 当kernel函数声明普通变量时，是存放再寄存器里的
    展开查看：示例代码 __global__ void sum(int* arr) { int i; // 存放在寄存器里 }
  2. blockDim太大容易发生寄存器打翻:
    1. 一个block有固定数量的寄存器，寄存器会均匀的分给这个block中的所有线程，如果线程太多，每个线程分到的寄存器数量就很小
    2. 如果你每个线程声明的普通变量太多，寄存器存不下，就会使用L1缓存

线程太少不行：读数据时延很高

  1. "大吞吐量掩盖高时延"这里的高时延指的就是从显存DRAW读数据很慢，所以一个线程读取数据的时候就会切换到其他线程执行
  2. 结论：1.使用寄存器少但访存使用多的函数，尽可能多的开线程（让blockDim大）如：矢量加法。2. 光线追踪尽可能少开线程（让blockDim小，但不要太小）

减少跨步才高效

  1. 例如矩阵转置，同时存在按行访问[1]和按列访问[2]，无论如何都存在跨步，只能把访存跨步转移为共享内存跨步，效率才会有所提升
    展开查看：共享内存快，所以把跨步转移到共享内存上能提高一点效率 drawing

共享内存反直觉优化

bank和__shared__Memory概念

sinf和__sinf(两个下划线)的区别

  1. sinf精度比__sinf高，但__sinf计算更快
  2. 当n接近1e25时，__sinf完全错误

thrust(gpu的仿STL库)

Gpu原子操作

问题:

  1. 有线程就可以实现并行了，那为什么要引入板块(grid和block)的概念?
    答：

    1. gpu编程常常涉及三维和二维矩阵运算，引入grid和block会比较方便(有x,y,z)
    2. 如果没有block和grid线程就都是一维的，就不方便
    3. 矩阵转置的例子：如果只有一维，会发生按行访问和按列访问，导致低效。cuda有三维天生带循环分块
      展开查看：循环分块截图02:36:03 drawing

  2. 声明和定义放在同一个文件内，真的可以提高性能吗？
    答：教程00:35:34里提到，放在同一个文件内编译器可能知道要如何优化，如果分离了声明和定义，编译器会生成call指令，从而导致效率降低

  3. "Gpu是使用大吞吐量掩盖延迟"这句话如何理解?
    答：

    1. 高延迟：指的是gpu读取显存很慢，如下图，从GlobleMemory(DRAM)中读取数据是很慢的
    2. 大吞吐量：线程足够多，当一个线程在读取数据的时候，就可以切换到其他线程进行运算
      展开查看：Gpu内存模型图截图 drawing

  4. 主存与显存分配函数区分

    1. 主机内存：malloc, free
    2. 设备内存: cudaMalloc, cudaFree
    3. 统一内存（managerd）：cudaMallocManaged, cudaFree
      统一内存虽然方便，但毕竟有一点开销，尽可能用主机内存和设备内存吧

  5. 在向量加法中，如果向量大小不能被32整除怎么办？

    1. 课程00:58::36中说可以用网格跨步循环
      展开查看： drawing
    2. 向上取整，(n+blockDim-1) / blockDim这条式子可以多申请一个block，然后用if判断threadId有没有超过向量长度即可
      展开查看：截图中红线部分错了，应该是减一 drawing

  6. cpu函数向gpu函数传参数要传值还是传引用？

    1. 传值，因为要考虑是gpu内存还是cpu内存
    2. 用lambda捕获外部vector时,最好是捕获指针vec.data()，防止发生拷贝影响性能
      展开查看：lambda捕获外部vec最好用vec.data() drawing

  7. 如何测量gpu耗时？

    1. 要多测几次，先让gpu预热(显卡驱动初始化也耗时间)再测试更准
    2. 用nsight工具测时间最准

  8. 课程中提到的"再线程局部分配一个数组，并通过动态下标访问（例如遍历BVH树时用到的模拟栈），是无论如何都会打翻到一级缓存的，因为寄存器不能动态寻址"是啥？

    展开查看：截图 drawing
    1. BVH树 2. 寄存器不能动态寻址