7月5日,2022 CUDA on Arm Platform线上训练营开始第二天的课程。
课程大纲:
课堂问题汇总:
答:--这个不能通过API查询。只能检测计算能力后,写死。(例如7.5的计算能力,对应64个/SM)。而计算能力可以通过cudaGetDeviceProperties()获得,这样你再硬编码一个计算能力和SP数量/SM的对应关系的表格,就可以得到你的卡的每SM的SP个数了。这也是deviceQuery例子的做法,无法直接通过API得到。
--不同的大小可能会导致不同的性能变化。在你的卡(Jetson Nano上),我不建议你使用低于64(不含)的数值。因为该硬件设备最大能上2048线程/SM,但最多只能同时上32个线程。这样小于64个线程/block,将影响最大驻留blocks能力(不一定会表现出来性能上的降低,但是有潜在影响)。其他的形状哪种能最佳性能,需要试验,这个我不能直接知道(你也不能),我们需要实验。
--第三个参数是指的要传输(复制)的字节数。搜索《CUDA Runtime API》手册以获取更多信息。(其他不懂的函数,也可以直接快速翻阅手册得到答案,或者自学)。
--没错。是这样的。过大的grid会导致一些额外的blocks上到SM上,但是这些blocks被if之类的条件给快速抑制掉了(结束掉了)。所以除了浪费一点点性能,没啥缺点
--CUDA C Programming Guide上按照计算能力给出的block最大形状,和grid最大性能。不过目前所有支持的卡都是block的3个分量,每个最大1024(同时总数量1024)限制;grid的3个分量(x,y,z), 最大分别是INT_MAX(即2^31 - 1),64K,64K, 同时无总数量的乘积上的额外限制。所以你直接死记硬背就可以了。
--是的。超出哪怕1个线程,也会分配一个warp(浪费31/32的潜在执行能力)。
——在计算能力5.0+的硬件上,warp是硬件+软件协同调度的。搜索maxwell + control code(将maxwell替换成其他架构),从google获取更多信息。
——实际的访存往往也是Z轴,在线性地址上的跨度最大的;或者安排x/y/z填充warp的顺序,Z轴也是最后不优先变化的。如果这些可以算成“物理上”,则你可以这样认为。其他都可以认为是“虚拟的”。
9、尖括号中的两个数是不是就是gridDim和blockDim?
--菱形配置符号里的前两个参数是这样的。不过gridDim和blockDim仅在设备代码(GPU代码)中才有效。在Host端他们是普通的两个dim3结构体。
--如果你这里的内存是指的shared memory的笔误,那么的确是以block为单位分配shared memory资源的。
--每个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,需要特殊的写法,本次课程不讲述)
以下是学员学习的笔记分享: