我发现了一篇文章,其中一位 CS 研究人员使用微基准测试对 Fermi 设备上的块调度程序进行逆向工程:
http://cs.rochester.edu/~sree/fermi-tbs/fermi-tbs.html http://cs.rochester.edu/%7Esree/fermi-tbs/fermi-tbs.html
我修改了他的代码以在我的 GPU 设备(GTX 1080,带有 Pascal GP104 GPU)上运行,并随机化运行时间。
Methods
每个块仅包含 1 个线程,并使用足够的共享内存启动,每个 SM 只能驻留 2 个块。内核记录其开始时间(通过获取clock64()
),然后运行一段随机时间(该任务是使用进位乘法算法生成随机数)。
GTX 1080 由 4 个图形处理集群 (GPC) 组成,每个集群有 5 个流式多处理器 (SM)。每个 GPC 都有自己的时钟,因此我使用链接中描述的相同方法来确定哪些 SM 属于哪些 GPC,然后减去固定偏移量以将所有时钟值转换为同一时区。
Results
对于一维块网格,我发现块确实是按连续顺序启动的:
我们有 40 个块立即开始(每个 SM 2 个块 * 20 个 SM),后续块在前一个块结束时开始。
对于二维网格,我发现了相同的线性顺序,blockIdx.x
是快速维度并且blockIdx.y
慢维度:
注意:我在标记这些图时犯了一个严重的拼写错误。 “threadIdx”的所有实例都应替换为“blockIdx”。
And for a 3-d block grid:
结论
对于一维网格,这些结果与 Pai 博士在链接文章中报告的结果相匹配。然而,对于二维网格,我没有找到任何块执行顺序中空间填充曲线的证据,因此这可能在费米和帕斯卡之间的某个地方发生了变化。
当然,基准测试的常见警告也适用,并且不能保证这不是特定于特定处理器型号的。
Appendix
作为参考,下面的图显示了随机运行时间与固定运行时间的结果:
事实上,我们在随机运行时看到了这种趋势,这让我更有信心这是一个真实的结果,而不仅仅是基准测试任务的一个怪癖。