在OpenCL中使用本地內存的奇怪行爲

我目前正在使用NVIDIA Tesla C1060（驅動程序版本195.17）上的OpenCL項目。不過，我收到了一些我無法解釋的奇怪行爲。這裏是一個困擾我的代碼（減少清晰度和測試目的）：在OpenCL中使用本地內存的奇怪行爲

kernel void TestKernel(global const int* groupOffsets, global  float* result,  
         local  int* tmpData,    const int itemcount) 
{ 
    unsigned int groupid = get_group_id(0); 
    unsigned int globalsize = get_global_size(0); 
    unsigned int groupcount = get_num_groups(0); 

    for(unsigned int id = get_global_id(0); id < itemcount; id += globalsize, groupid += groupcount) 
    { 
     barrier(CLK_LOCAL_MEM_FENCE); 
     if(get_local_id(0) == 0) 
     tmpData[0] = groupOffsets[groupid]; 
     barrier(CLK_LOCAL_MEM_FENCE); 
     int offset = tmpData[0]; 
     result[id] = (float) offset; 
    } 
}

此代碼應加載每個工作組到本地內存的偏移量，然後讀回，並將其寫入到相應的outputvector條目。對於大多數工作項目而言，這是可行的，但是對於每個工作組而言，本地ID爲1到31的工作項會讀取不正確的值。我的輸出向量（對於workgroupsize = 128）是如下：

index  0: 0 
index 1- 31: 470400 
index 32-127: 0 
index  128: 640 
index 129-159: 471040 
index 160-255: 640 
index  256: 1280 
index 257-287: 471680 
index 288-511: 1280 
...

我期望的輸出將是

index 0-127: 0 
index 128-255: 640 
index 256-511: 1280 
...

奇怪的是：只有當我使用較少然後ITEMCOUNT工作項出現該問題（因此，當globalsize> = itemcount時，它會按預期工作，這意味着每個工作項只處理一個條目）。所以我猜測它與循環有關。有誰知道我在做什麼錯，如何解決它？

更新：我發現，好像如果我改變工作

if(get_local_id(0) == 0) 
    tmpData[0] = groupOffsets[groupid];

到

if(get_local_id(0) < 32) 
    tmpData[0] = groupOffsets[groupid];

哪個更使我感到驚訝，所以雖然它可能會解決問題，我不要以這種方式修復它（因爲它可能會打破其他時間）。除此之外，我寧願避免在Geforce 8xxx類硬件上運行時由於額外的（根據我的理解，未合併）內存訪問而損失性能。所以問題依然存在。

來源

2010-01-30 Grizzly

首先，重要的是，您需要注意itemcount是本地工作規模的倍數，以避免在執行屏障時發生偏差。

在處理器上執行內核的工作組中的所有工作項必須先執行此函數，然後才允許任何超出屏障的執行繼續執行。執行內核的工作組中的所有工作項都必須遇到該函數。

您可以實現此如下：

unsigned int itemcountrounded = get_local_size(0) * ((itemcount + get_local_size(0) - 1)/get_local_size(0)); 
for(unsigned int id = get_global_id(0); id < itemcountrounded; id += globalsize, groupid += groupcount) 
{ 
    // ... 
    if (id < itemcount) 
     result[id] = (float) offset; 
}

你說的代碼減少爲簡單起見，如果你運行你發佈什麼，會發生什麼？只是想知道你是否需要把障礙放在全球記憶中。

來源

2010-02-01 20:49:21 Tom

我爲了簡單而減少的含義是，我真正想要的不是將groupid公佈到每個vektor條目中。我發佈的結果是運行psted內核的結果（至少一次運行，不同運行之間的錯誤條目似乎有所不同）。我已經確保itemcount是localworksize的倍數，但是從我的測試中，無論哪種方式都沒有關係（因爲無論itemcount是否可以被本地工作大小整除，行爲基本相同） – Grizzly 2010-02-02 16:52:00

您是否嘗試過也把全球記憶的障礙？即「屏障（CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE）」 – Tom 2010-02-02 21:15:06

在OpenCL中使用本地內存的奇怪行爲

回答

相關問題