GPU 一致性和功耗:每个 warp 仅在一个线程上运行代码是否可以节省电量?

问题描述 投票:0回答:1

在 GPU 编程中,人们非常关心“一致性”:线程在所谓的“warps”中执行(每个线程大约有 32 个线程)。扭曲的所有线程都以“锁步”方式运行,以便始终为所有线程执行相同的指令。如果某个分支仅被这些线程的一部分占用,则其他线程将被禁用,但同时不执行任何有用的操作。这称为背离。相反(所有线程总是一起执行有用的指令)称为一致性。

我的问题是关于禁用线程到底做了什么。特别是:禁用的线程与启用的线程消耗相同的功率吗?

许多 GPU 架构使用线程掩码,其中每个线程的标志存储线程是否“实际”执行该指令。我听说较新的 GPU 架构每个线程都有一个程序计数器,这在某些方面要好一些。如果禁用的线程仍然完全执行所有指令,但由于线程掩码,结果根本不会刷新,那么我假设功耗也是相同的。


为了使其更具体一些,假设我有一个工作组大小为 16x16 的计算着色器。我需要计算工作组中所有线程都相同的一些值。计算很难并行化(并且不会花费太长时间),所以我只是以正常的串行方式编写它。

在这种情况下,用

if (gl_LocalInvocationIndex == 0)
包裹该代码并将结果写入工作组共享内存有什么好处吗?或者我可以让所有线程执行代码吗?

cuda gpu vulkan gpgpu compute-shader
1个回答
0
投票

每个 warp 仅在一个线程上运行代码是否省电?

没关系,因为如果你这样做,你就是在浪费你的 GPU。如果您计划每个 warp 使用一个线程,那么最好只在 CPU 上运行代码,这样也会浪费更少的电量。

...但少一点开玩笑:

禁用的线程到底做什么...禁用的线程与启用的线程消耗相同的功率吗?

没有这样的线程。这只是一个比喻。将 GPU 扭曲视为 CPU 上的 32 通道宽 SIMD 寄存器。 (GPU 硬件不一定完全一样。)

现在,实现“线程禁用”的最简单方法是为扭曲中的每个线程设置掩码位,并使用它来禁用写入结果,也许可以避免被零除异常等。我非常怀疑是否会为了节省电量而避免无用的计算 - 尽管理论上这是可能的。

我想你可以检查一下:编写两个内核,全部执行相同类型的计算(如果可以的话,尝试使其密集,所以也许交错不同类型的操作,以便不同的扭曲保持每个核心忙碌),并安排它其中之一使 32 个扭曲“线程”中的 31 个禁用(例如

if (threadIdx.x % 32 == 0) { /* do work */ }
)。让它们运行足够长的时间以使功耗稳定下来,并测量该功耗,例如使用
nvidia-smi

© www.soinside.com 2019 - 2024. All rights reserved.