我已经使用CUDA几周了,但我对块/ warps/thread的分配有些怀疑. 我从教学的角度(大学项目)研究建筑,所以达到最佳表现并不是我关注的问题.
首先,我想了解我是否直截了当地得到了这些事实:
程序员编写内核,并在线程块网格中组织其执行.
每个块都分配给一个流式多处理器(SM).一旦分配,它就无法迁移到另一个SM.
每个SM将其自己的块拆分为Warps(当前最大大小为32个线程).warp中的所有线程在SM的资源上并发执行.
线程的实际执行由SM中包含的CUDA核执行.线程和核心之间没有特定的映射.
如果warp包含20个线程,但目前只有16个可用核心,则warp将不会运行.
另一方面,如果一个块包含48个线程,它将被分成2个warp并且它们将并行执行,前提是有足够的内存可用.
如果线程在核心上启动,则它会因内存访问或长时间浮点操作而停止,其执行可以在不同的核心上恢复.
他们是对的吗？
现在,我有一个GeForce 560 Ti,因此根据规格,它配备了8个SM,每个包含48个CUDA核心(总共384个核心).
我的目标是确保架构的每个核心都执行相同的SAME指令.假设我的代码不需要比每个SM中可用的代码更多的寄存器,我想象了不同的方法:
我创建了8个块,每个48个线程,因此每个SM有1个块来执行.在这种情况下,48个线程将在SM中并行执行(利用它们可用的所有48个内核)？
如果我推出64个6个线程的块,有什么区别吗？(假设它们将在SM之间平均映射)
如果我在预定的工作中"淹没"GPU(例如,创建每个1024个线程的1024个块),可以合理地假设所有核心将在某个点使用,并且将执行相同的计算(假设线程永远不会失速)？
有没有办法使用Profiler检查这些情况？
这个东西有没有参考？我阅读了"编程大规模并行处理器"和"CUDA应用程序设计与开发"中的CUDA编程指南和专用于硬件架构的章节; 但我无法得到准确的答案.
有关如何选择#blocks和blockSize的讨论很多,但我仍然遗漏了一些东西.我的许多问题都解决了这个问题:CUDA阻止/变形/线程如何映射到CUDA核心？ (为了简化讨论,有足够的perThread和perBlock内存.内存限制在这里不是问题.)
kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal);
1)为了让SM尽可能忙,我应该设置nThreads为倍数warpSize.真正？
2)SM一次只能执行一个内核.这就是SM的所有HWcores只执行kernelA.(不是某些HWcores运行kernelA,而其他运行kernelB.)因此,如果我只有一个线程可以运行,那么我就"浪费"了其他的HWcores.真正？
3)如果warp-scheduler以warpSize(32个线程)为单位发出工作,并且每个SM有32个HWcore,则SM将被充分利用.SM有48个HWcores时会发生什么？当调度程序以32块为单位发布工作时,如何充分利用所有48个核心？(如果上一段是真的,如果调度程序以HWcore大小为单位发布工作会不会更好？)
4)看起来warp-scheduler一次排队2个任务.因此当当前正在执行的内核停止或阻塞时,第二个内核被交换.(目前尚不清楚,但我猜这里的队列深度超过2个内核.)这是正确的吗？
5)如果我的HW的上限为每块512个线程(nThreadsMax),这并不意味着具有512个线程的内核将在一个块上运行得最快.(同样,mem也不是问题.)如果我将512线程内核分布在多个块中,而不仅仅是一个块,那么我很有可能获得更好的性能.该块在一个或多个SM上执行.真正？
5a)我认为越小越好,但是我做得多小也一样重要nBlocks吗？问题是,如何选择nBlocks那个体面的价值？(不一定是最优的.)是否有选择的数学方法nBlocks,或者它只是试验性的.
我阅读了NVIDIA Fermi白皮书，并在计算SP内核，调度程序的数量时感到困惑。
根据白皮书，每个SM中都有两个Warp调度程序和两个指令分派单元，从而允许同时发布和执行两个Warp。SM中有32个SP内核，每个内核具有完全流水线化的ALU和FPU，用于执行线程的指令
众所周知，一个warp由32个线程组成，如果我们仅在每个周期内发出一个warp，则意味着该warp中的所有线程将占用所有SP内核，并在一个周期内完成执行（假设没有任何停顿） ）。
但是，NVIDIA设计了双调度程序，该调度程序选择两个扭曲，然后从每个扭曲向一个16个核，16个加载/存储单元或4个SFU的组发出一条指令。
NVIDIA表示，这种设计可带来最高的硬件性能。最高的硬件性能可能来自对不同指令的交错执行，从而充分利用了硬件资源。
我的问题如下（假设没有内存停滞并且所有操作数都可用）：
每个warp是否需要两个周期来完成执行，并且每个warp调度程序将所有32个SP内核分为两组？
ld / st和SFU单元由所有经线共享（看起来像来自双重调度程序的经线的统一服）？
如果将经纱分为两个部分，请先安排哪一部分？有调度程序吗？或只是随机选择一个要执行的部分。
这种设计的优点是什么？只是最大限度地利用硬件？
我试图了解CUDA扭曲和线程分歧.假设我有一个朴素矩阵乘法内核来乘以nxn矩阵.
__global__ void matrix_multiply(float* a, float* b, float* c, int n)
{
int row = blockIdx.y + blockDim.y + threadIdx.y;
int col = blockIdx.x + blockDim.x + threadIdx.x;
if(row < n && col < n) {
float tmp = 0.0f;
for(int i = 0; i < n; ++i)
tmp += a[row * n + i] * b[i * n + col];
c[row * n + col] = tmp;
}
}
如果我启动一个网格大小为32乘32且块大小为16乘16并且矩阵为500乘500的内核,那么有多少warp会遇到会遇到线程分歧的线程？
由于矩阵右边缘的每个线程块都有线程发散,因此线程发散的经线数不应该是256吗？