我有一个从for循环中调用的CUDA内核.就像是
for(i=0; i<10; i++) {
myKernel<<<1000,256>>>(A,i);
}
现在假设我有一个带有15个流多处理器(SM)的NVIDIA卡.同样假设,为了简单起见,只有一个块可以映射到SM,这基本上说大多数时候,我将在设备上执行15个块.由于内核执行是异步的,基本上i = 1的调用将在第一个内核启动后立即执行(i = 0).
我的问题是:在执行第一个内核(i = 0)的某个时刻,只有14个SM忙,然后只有13个,然后只有12个,那么只有11个等等.
只要有一个SM可用,就会将i = 1的内核发送到设备上执行,或者第二个内核的启动是否会等到所有SM完成处理第一个内核(i = 0的内核)？
还假设我在一个CUDA流中工作.
前言:假设我在CUDA中使用NVIDIA GTX480卡.该卡的理论峰值全局内存带宽为177.4 GB/s:384*2*1848/8*1E9 = 177.4 GB/s
384来自内存接口宽度,2来自内存的DDR特性,1848是内存时钟频率(以MHz为单位),8来自我希望在Bytes中获得答案的事实.
可以为共享存储器计算类似的东西:每个存储体4个字节*32个存储体*每个周期0.5个存储体*1400MHz*15个SM = 1,344 GB/s
SM数量的上述因素,即15,因此,为了达到这个最大共享内存带宽,我需要让所有 15个SM读取共享内存.
我的问题:为了达到最大全局内存带宽,只从全局内存中读取一个 SM,或者所有SM是否同时尝试从全局内存中读取？更具体地说,假设我使用一个具有32个线程的块启动内核.然后,如果我在SM-0上有唯一的warp,并且我在内核中所做的所有操作都以合并的方式从全局内存中不间断地读取,那么我将达到177.4 GB/s吗？或者我应该启动至少15个块,每个块有32个线程,以便SM-0到SM-14上的15个warp同时尝试读取？
当务之急可能是运行一个基准测试来解决这个问题.我想知道为什么会发生什么,发生.