我对运行在GPU上的内核吞吐量有疑问。假设它的占用率为0.5，块大小为256: 编程指南指出最好有许多块，这样它们就可以隐藏存储器延迟等。但我不明白为什么这是正确的。因为一旦内核每个流式多处理器的线程束数= 24，即3个块，它将达到峰值吞吐量。因此，拥有超过24个线程束（或3个块）对吞吐量不会产生任何影响。
我有什么遗漏吗？有人能纠正我吗？
我有什么遗漏吗？有人能纠正我吗？
2个回答
6
虽然低占用的SM不能充分隐藏延迟，但重要的是要理解这一点：高占用并不等于高吞吐量！占用只是衡量SM在任何给定时刻可选择多少工作的一种方式。拥有更多的常驻warp使SM能够在其他warp等待结果（内存访问或计算结果 - 两者都具有非零延迟）的情况下执行更多的有效工作。吞吐量是每秒完成的工作量的度量标准，尽管它可能受到延迟（因此也可能受到占用）的限制，但它也可能受到内存带宽、指令吞吐量（执行单元数量）和其他因素的限制。编程指南中提到使用多个线程块比仅使用一个大线程块更好的原因是因为有时候最好能够从不仅是其他warp而且还从其他块发出工作。以下是一个例子：假设您的大线程块必须从全局内存（高延迟）加载数据并将其存储到共享内存（低延迟），然后必须立即执行__syncthreads()。在这种情况下，当一个warp完成其数据加载并将其写入共享内存时，它必须等待所有其他块中的线程完成相同的操作。对于大块来说，这可能需要相当长的时间。但是如果有多个较小的线程块占用SM，则SM可以在等待第一个块满足__syncthreads时切换并执行其他块的工作。这可以帮助减少GPU空闲时间并提高效率。您不一定想要非常小的块（因为Fermi上的SM最多支持8个常驻块），但使用128-512个线程的块通常比使用1024个线程的块更有效。
- harrism
6
好的，我明白你的意思了。所有的都是正确的，我相信。但是有些东西还缺失了。请使用我的例子：occupancy=0.5，blocksize=256。我感觉这对于许多内核来说很常见。这意味着在GTX580中我将有24个活动warp（或3个活动块）。那么启动3个块和6个块之间的吞吐量差异是什么？造成这种差异的原因是什么？ - Zk1001
高占用并不等于高吞吐量！因此，如果不知道内核的瓶颈具体是什么，通常无法预测启动3个块和6个块之间吞吐量的差异。 现在，如果占用率被限制在0.5（由寄存器或共享内存使用引起），并且您的GPU至少有6个SM，则启动6个块而不是3个块应该可以将吞吐量提高一倍，因为这将使用更多的SM。通常，您希望至少启动与您拥有的SM数量相同的块数，如果可能的话，每个SM可以启动多个块。但我认为这不是你所问的内容... - harrism
哦，抱歉我搞错了。我想说的是每个SM有3个块和6个块（而不是总块数）。是的，我非常确定它将取决于内核的瓶颈，但是如何呢？还应该有一些一般情况，对吧？（并且假设占用率保持在0.5-实际上它不会受到块数的影响） - Zk1001
简单来说，核心占用率为0.5。这意味着每个SM可以有3个活动块。在GPU中，有16个SM。当我总共启动48个块时，每个周期的指令数（IPC）为100。但是当我总共启动64个块时，IPC为120，这种差异的原因可能是什么？ - Zk1001
我想你的意思是1.0和1.2而不是100和120，但我猜测你的块启动时间存在一些偏差，或者你的块运行时间变量，因此拥有超过48个块可以在初始的3个块完成后填补一些SM中的空洞，从而提高整体效率。 - harrism
谢谢。这正是我一直猜测的。感谢确认。我猜它看起来就像乐高游戏吧。 - Zk1001
1
如果您的CUDA启用卡上只有一个SM，那么拥有超过3个块也不会改变吞吐量。通常单个GPU中有8个或更多SM。
此外，运行在一个SM上的块数并不仅仅基于warp数量。这只是一个限制因素，还有许多其他因素。CUDA Occupancy Calculator是一个很好的工具，可以查看内核的占用情况。
- jwdmsd
1
如果我的卡上有多个SM，比如16个，那么吞吐量会发生什么变化呢？我知道有Cuda occupancy calculator，但它在这里并没有帮助。我的问题是，“拥有足够的warp（所有warp都是活动状态）和拥有超过足够的warp（你有一些活动warp，其他warp处于空闲状态）之间的区别是什么？”我敢打赌答案不会很简短。 - Zk1001
网页内容由stack overflow 提供, 点击上面的可以查看英文原文，
原文链接
原文链接