“暑”你当学霸|2022 CUDA线上训练营Day 2学员笔记分享

2022-07-06 19:41:29 浏览数 (1)

7月5日，2022 CUDA on Arm Platform线上训练营开始第二天的课程。

课程大纲：

  1. 1.     CUDA编程模型---CUDA存储单元的使用与错误检测（2.1 2.2实验课）
  2. Ÿ  设备初始化
  3. Ÿ  GPU的存储单元
  4. Ÿ  GPU存储单元的分配与释放
  5. Ÿ  数据的传输
  6. Ÿ  数据与线程之间的对应关系
  7. Ÿ  CUDA应用程序运行时的错误检测
  8. Ÿ  CUDA中的事件
  9. Ÿ  利用事件进行计时
  10. Ÿ  实验课内容：编写MatrixMul程序，体验线程和数据的对应关系
  11. Ÿ  留课后作业
  13. 2.     多种CUDA存储单元详解（2.3）
  14. Ÿ  CUDA中的存储单元种类
  15. Ÿ  CUDA中的各种存储单元的使用方法
  16. Ÿ  CUDA中的各种存储单元的适用条件
  18. 3.     利用共享存储单元优化应用（2.4实验课）
  19. Ÿ  共享存储单元详解
  20. Ÿ  共享内存的Bank conflict
  21. Ÿ  利用共享存储单元进行矩阵转置和矩阵乘积
  22. Ÿ  实验课内容：编写Shared Memory优化过的矩阵乘法
  23. Ÿ 介绍shared memory原理，介绍利用shared memory 优化的多种案例
  24. Ÿ 矩阵转置

课堂问题汇总：

  1. 1.     怎么看一个sm里边几个cuda core

答：--这个不能通过API查询。只能检测计算能力后，写死。（例如7.5的计算能力，对应64个/SM）。而计算能力可以通过cudaGetDeviceProperties()获得，这样你再硬编码一个计算能力和SP数量/SM的对应关系的表格，就可以得到你的卡的每SM的SP个数了。这也是deviceQuery例子的做法，无法直接通过API得到。

  1. 2.     尖括号中，只要blocksize设置成32的倍数，warp就能最佳分配，那我设成32，64，128这样会有什么区别呢？

--不同的大小可能会导致不同的性能变化。在你的卡（Jetson Nano上），我不建议你使用低于64（不含）的数值。因为该硬件设备最大能上2048线程/SM，但最多只能同时上32个线程。这样小于64个线程/block，将影响最大驻留blocks能力（不一定会表现出来性能上的降低，但是有潜在影响）。其他的形状哪种能最佳性能，需要试验，这个我不能直接知道（你也不能），我们需要实验。

  1. 3.     cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);这里面的M是干啥用的？

--第三个参数是指的要传输（复制）的字节数。搜索《CUDA Runtime API》手册以获取更多信息。（其他不懂的函数，也可以直接快速翻阅手册得到答案，或者自学）。

  1. 4.     我的理解是 只要加了边界判断 只会导致效率低些 其他没什么影响

--没错。是这样的。过大的grid会导致一些额外的blocks上到SM上，但是这些blocks被if之类的条件给快速抑制掉了（结束掉了）。所以除了浪费一点点性能，没啥缺点

  1. 5.     怎么查看每一维最大的size数来着？

--CUDA C Programming Guide上按照计算能力给出的block最大形状，和grid最大性能。不过目前所有支持的卡都是block的3个分量，每个最大1024（同时总数量1024）限制；grid的3个分量(x,y,z), 最大分别是INT_MAX（即2^31 - 1)，64K，64K, 同时无总数量的乘积上的额外限制。所以你直接死记硬背就可以了。

  1. 6.     改成了33就分成2个warp？

--是的。超出哪怕1个线程，也会分配一个warp（浪费31/32的潜在执行能力）。

  1. 7.     warp是硬件调度吧？

——在计算能力5.0 的硬件上，warp是硬件 软件协同调度的。搜索maxwell control code（将maxwell替换成其他架构），从google获取更多信息。

  1. 8.     indedx（x,y,z）z是竖轴的，物理上也是三维的？还是说只是逻辑上虚拟出来的呢？

——实际的访存往往也是Z轴，在线性地址上的跨度最大的；或者安排x/y/z填充warp的顺序，Z轴也是最后不优先变化的。如果这些可以算成“物理上”，则你可以这样认为。其他都可以认为是“虚拟的”。

9、尖括号中的两个数是不是就是gridDim和blockDim？

--菱形配置符号里的前两个参数是这样的。不过gridDim和blockDim仅在设备代码(GPU代码)中才有效。在Host端他们是普通的两个dim3结构体。

  1. 10.  也就是说gpu的内存调度是以block调度的，不是以warp调度的是吗？

--如果你这里的内存是指的shared memory的笔误，那么的确是以block为单位分配shared memory资源的。

  1. 11.  Thread不够，forloop那个循环不太理解

--每个thread和其他thread并无本质不同，连代码都执行的是同一份。唯一让它变得特别的，是它使用的下标。所以根据这个理解：如果原本有0,1,2,3,4,5号线程去完成任务，和你只有0号线程，但是它完成了原本0号的下标负责的对应的任务后，继续下标 1变成1号线程，也一样可以的。如此类推，哪怕只有1个线程，连续变身从0到5，和直接有5个线程，并无本质区别的。有个这个理解后，再扩大一下，如果原本是1000个线程，直接上1000个可以。如果上不了，只能上300个。那么这300个完成了前300号任务后， = 总线程数，变身成为300-600号（不含），依然可以完成再300个不存在的线程的原本要负责的任务。这样，最终他们累加多次，就将前300号，中间300号，最后300号，以及剩下的100号（你需要用while或者if判断别干多了），这样都干完了。

12.  cuda里把连续128bit的数据从global memery先复制到shared memory再复制到register，和先从gmem到reg再到smem，速度有差别吗

--直接复制到shared memory, 不使用特殊的写法，直接用等号的形式，例如：

__shared__ int kachi[...];

kachi[xxx] = ptr[xxxx]; （其中ptr是一个指向显存/global memory某区域的指针）。

这种写法实际上编译器，“会自动通过寄存器中转的”，和你手工：

tmp = ptr[xxxx];

dog[xxx] = tmp;

并无本质区别。因为这两个是一回事，只是节省了你的中间变量的写法，实际代码生成是一样的，所以不会有任何差别。（真正的直接能一步到shared，需要特殊的写法，本次课程不讲述）

以下是学员学习的笔记分享：

深度学习 api vr视频解决方案

0 人点赞