CUDA：哪個內存空間是固定大小的數組存儲？

Question

當在內核設置一個固定尺寸數組，如：

int my_array[100]; 

在哪個存儲器空間並數組結束？

特別是，我想知道這樣一個數組是否可以存儲在寄存器文件或共享內存中的大於等於2.0的設備上，如果是這樣，需求是什麼。

Answer 1

對於費米（也可能是早期的架構），對於要存儲在寄存器文件中的一個數組，以下條件必須得到滿足：

1. 陣列僅與常數
2. 有可用的寄存器
3. 希望編譯器索引也做一些分析以確定對整體性能的影響

（1）的原因是寄存器索引直接在SASS指令內編碼。沒有辦法間接處理寄存器。

限制爲寄存器的數量的主要因素（2）爲：

• 的SASS指令只包含6寄存器索引，這限制了可以在一個內核被用於寄存器的數目的位64.實際人數是63人，因此其中一人留作東西。
• 一個SM有一個由所有正在運行的線程共享的寄存器塊。
• 寄存器對於保存變量也是必需的，因此編譯器必須平衡寄存器使用以獲得最佳整體性能。

（1）的潛在解決方法是循環展開。如果循環使用循環計數器作爲數組中的索引，展開循環（使用#pragma unroll或手動）將導致數組索引變爲常量，因爲現在每個數組訪問都有單獨的SASS指令。

部分基於此NVIDIA演示文稿：Local Memory and Register Spilling。該文件還詳細介紹了變量和陣列的位置如何影響性能。

Answer 2

如果沒有足夠的寄存器，內核中的本地數組（如已定義的內存數組）將分配到寄存器和本地內存中。

如果你想分配共享內存中的數組，你必須將其指定爲：

__shared__ int my_array[100];