2011-09-28 87 views
3

饿了一部分,当我使用太多的寄存器,基本上有3个选项我能做得到的内核:确定最注册内核

  1. 离开内核,因为它是导致低入住
  2. 集编译器使用数量较少的寄存器,它们溢出,导致糟糕的表现
  3. 重写内核

选项3,我想知道哪个内核的一部分需要记数的最大数量ERS。有没有任何工具或技术可以让我识别这部分?通过PTX代码(我在NVidia上开发)读取没有帮助,寄存器有很多数字,说实话,我能做的最好的是确定汇编代码的哪一部分映射到C代码的哪一部分。

只是注释掉一些代码并不是一个很好的方法 - 例如,我注意到如果我只是把代码放到循环中,寄存器的数量会大大增加,不仅仅是一个循环控制变量。我个人怀疑NVidia编译器是从不完善的变量生存分析中,但当然我不能用这么做:-)

回答

0

这是一个任何语言都很棘手的问题,可能没有一个正确的答案。但是这里有一些需要考虑的事情:

  • 寻找你可以找到的“最深”范围内的代码,记住大多数函数可能是由你的OpenCL编译器内联的。计算在该范围中使用的变量,然后遍历包含的范围。在每个包含范围内,计数在内部范围之前和之后使用的变量。这些可能会在内部作用域执行时生效。这个过程可以帮助您在程序的特定部分考虑实时寄存器。
  • 尝试采用“跨度”深层作用域的变量,并尽可能使其不跨越范围。举例来说,如果你有这样的事情:

    int i = func1(); 
    int j = func2(); // perhaps lots of live registers here 
    int k = func3(i,j); 
    

    你可以尝试重新排序的前两行,如果FUNC2有大量的现场登记。当func2正在运行时,这将从活动寄存器集中移除i。当然,这是一个微不足道的模式,但希望它是说明性的。

  • 想想摆脱那些仅仅围绕简单计算结果的变量。您可能可以在需要时重新计算这些值。例如,如果你有像int i = get_local_id(0)这样的东西,那么你可能只需要使用get_local_id(0)就可以使用i。

  • 想想摆脱保留在内存中的值的变量。

没有这种东西的好工具,它最终成为艺术而不是科学。但希望这其中的一些是有帮助的。

+0

嗯,我知道这个“理论”,但我宁愿寻找好的工具。我不确定编译器使用了多少优化,但是一个好的编译器可以改变很多代码。例如,使用循环展开(肯定会完成),所使用的寄存器数量会增加变量数量。 –

1

如果你在NVidia硬件上运行,你可以传递-cl-nv-verbose编译选项到clBuildProgram然后clGetProgramInfo CL_PROGRAM_BINARIES获得关于编译的可读文本。在那里它会说它使用的寄存器的数量。请注意,NVidia缓存编译,并且只在内核源实际更改时生成该寄存器信息,因此您可能需要在源代码中注入一些多余的更改以强制其执行完整分析。

如果您在AMD硬件上运行,只需设置环境变量GPU_DUMP_DEVICE_KERNEL = 1即可。它将在编译期间生成IL的文本文件。不确定它明确说明了使用的寄存器数量,但它与上面的NVidia技术相当。

看着那个输出(至少在nvidia上),你会发现它似乎使用了无数的寄存器(如果你按寄存器号码)。实际上,它在进行流程分析时,实际上以查看IL时并不明显的方式重新使用寄存器。