我对cuda的调度系统有几个疑问。CUDA如何安排线程
答：当我使用例如foo < < < 255,255 >>（）函数，卡内实际发生了什么？我知道每个SM从上级接收一个要安排的块，并且每个SM负责安排其传入的BLOCK，但是哪一部分执行？例如，如果我有8个SM，每个SM包含8个小CPU，则上级负责安排剩余的255 * 255 - （8 * 8）个线程？
B.什么是可以定义的最大线程的限制？我的意思是foo<<<X, Y>>>(); x，y =？
C.关于最后一个例子，一个块内可以有多少个线程？我们可以说，我们有更多的块/线程，执行速度会更快吗？
感谢您的帮助
A.计算工作分配器将从网格向SM分配一个块。 SM将通过warp转换块（所有NVIDIA GPU上的WARP_SIZE = 32）。 Fermi 2.0 GPU每个SM都有两个共享一组数据路径的warp调度器。每个周期中，每个warp调度程序都会挑选一个warp并向其中一条数据路径发出指令（请不要考虑CUDA内核）。在Fermi 2.1 GPU上，每个warp调度程序都有独立的数据路径以及一组共享数据路径。在每个周期2.1上，每个变形调度程序将挑选一个变形，并尝试为每个变形指定双重指令。
warp调度程序试图优化数据路径的使用。这意味着单个warp可能会在背对背循环中执行多条指令，或者warp调度程序可以选择每个循环从不同的warp发出。
每个SM可处理的经线/线程数在CUDA编程指南v.4.2表F-1中指定。这从768个线程扩展到2048个线程（24-64个经线）。
B.每次启动的最大线程数由最大GridDims *每块最大线程数定义。请参阅表F-1或参阅cudaGetDeviceProperties的文档。 C.请参阅与（B）相同的资源。线程/块的最佳分配由您的算法分区定义，并受占用率计算的影响。基于SM上的经纱的问题集大小以及在指令障碍（除其他事项）上阻塞的时间量，可观察到的性能影响。对于初学者，我建议每个SM至少2个街区，占用率约50％。
B.这取决于你的设备上。您可以使用cuda功能cudaGetDeviceProperties查看设备的规格。一个公共的最大数目是y =每块1024线程，x = 65535块每个网格维度。
C.通常的做法是拥有2（128,256,512等）线程/块的幂。减少大型阵列非常有效（见Reduction）。块和线程的最佳分配实际上取决于您的应用程序和硬件。我个人在TeslaM2050上使用512个线程/块进行大型稀疏线性代数计算，因为它对我的应用程序来说效率最高。