2011-09-28 85 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時並不明顯的方式重新使用寄存器。