用於工作項緩存的全局內存

Question

由於本地內存限制，我需要使用全局內存作爲我的工作項的緩存。

假設我有1000個工作組，每個工作組有64個工作項目。每個項目需要4K緩存。工作項目完成後，高速緩存不需要保留。

我將分配一個單一的全局內存緩衝區，並將4K塊分配給工作項目。

（我是針對AMD的GPU）

什麼是我需要保證不會有 工作項目之間的任何併發問題的最小尺寸？

由於AMD具有< = 64 CU，那麼我的猜測是

64 * 128 * 4000個字節，並且使用（全局工作項ID％（64 * 128）） 到高速緩存塊分配給一個工作項。

Answer 1

如果每個緩存項（由global work item ID % (64*128)訪問）是一個4000個字節長的結構，並且實現不強制每個結構在4096個字節上對齊，並且緩存行大小不是4000的確切除數，並且如果全局內存銀行步長不是4000的確切除數，那麼它應該不成問題。

---

評測的這款內核codexl，（花了0.5秒爲16K的工作項）：

__kernel void test(__global float * a) 
    { 
     int i=get_global_id(0)*4096; 
     for(int j=0;j<4096;j++) 
      a[i+j]*=2.0f; 
    } 

和一些輸出：

• MEM單位停滯不前％55
• 緩存命中％45
• 存儲單元忙％99
• VALU忙％0.05

然後改變內核交織類型（在0.25秒執行）：

__kernel void test(__global float * a) 
    { 
     int i=get_global_id(0); 
     for(int j=0;j<4096;j++) 
      a[i+j*4096*4]*=2.0f; 
    } 

• MEM單元停頓％57
• 高速緩存命中％47
• MEM單位忙％84
• 值繁忙％1.5

所以交錯模式可以減少mem單元的壓力，並且更頻繁地點擊緩存，並且ALU部分更頻繁地提供並且更快地完成％50。

然後嘗試這樣做：

__kernel void test(__global float * a) 
{ 
    int i=get_global_id(0)*4100; 
    for(int j=0;j<4100;j++) 
     a[i+j]*=2.0f; 
} 

這個了0.37s，％30快於4096版本，但具有更高的MEM部檔位（終點非對準必須導致此浪費在不必要的數據的一些週期-fetches），緩存命中減少到％37。

測試GPU是R7-240

---

最後的測試與結構：

typedef struct test_struct 
{ 
    float test_field[4096]; 
}strr; 
__kernel void test(__global strr * a) 
{ 
    int i=get_global_id(0); 
    for(int j=0;j<4096;j++) 
    a[i].test_field[j]*=2.0f; 
} 

這0.53秒內完成，並有類似的分析數據，在開始跨入內核。

空內核在0.25秒內執行，所以它沒有用這個加載整個結構。只讀需要的元素。再次

typedef struct test_struct 
{ 
    float test_field[4096]; 
}strr; 
__kernel void test(__global strr * a) 
{ 
    int iLocal=get_local_id(0); 
    int iGroup=get_group_id(0); 
    for(int j=0;j<64;j++) 
    a[iGroup].test_field[iLocal+j*64]*=2.0f; 
} 

0.25秒所以它是儘可能快，因爲它可以：

---

分析信息交錯組爲中心的全球訪問。

Cache命中：44％ 紀念品單位忙道：％82 紀念品單位停滯不前：％67 VALU忙：0.9％

因此它具有所有的最好的條件，即使沒有高速緩存。