OpenCL get_local_id 速度/寄存器成本

Question

在 OpenCL 内核中，多次调用 get_local_id(0)，而不是在顶部调用一次，然后从局部变量访问该值是否更慢？

我是否通过不将 get_local_id(0) 存储到局部变量而是多次调用它来减少使用一个寄存器？

Answer 1

在大多数 GPU 架构上，本地工作项 ID 将保存在一组特殊的寄存器中，这些寄存器由硬件在分派(dispatch)线程时设置。这意味着当您执行 get_local_id(0) 时，您实际上只是从寄存器中读取数据，即没有速度损失“调用”get_local_id(0)。

编译器应该注意到您是否将 get_local_id(0) 分配给变量以供后续使用，并生成与每次都使用 get_local_id(0) 相同的代码。然而，根据我的经验，情况并非总是如此，您会注意到通过不使用变量来存储本地 ID 而只保存了一个寄存器。

如果您对寄存器很紧张，我建议您每次都调用 get_local_id(0)。您甚至可以设置一个宏来隐藏是否正在调用 get_local_id(0) 或正在使用寄存器。如果您有备用寄存器(您的探查器会告诉您这一点)或更喜欢您的代码使用变量，那么请使用变量。