OpenCL：選擇最優的global_work_size值

Question

假設我想添加兩個大小各爲1億的向量，應該如何最好選擇global_work_size？

• 是否必須global_work_size=100e6，或
• 是它更好地使用尺寸10e6 x 10e6的2D ND-範圍，或
• 任何其他策略？

顯然，我的GPU卡遠遠沒有同時提供1億個工作項目，甚至處理元素。

更具體地說，應該是global_work_size一般等於要並行處理的數據數量嗎？

感謝

Answer 1

我你的情況，你應該專注於選擇爲您的設備的最佳工作組大小，並讓每個工作項目過程中的多個元素。讓全局工作項數量等於要處理的元素數量通常會更簡單，但是對於GPU而言，您的問題非常大，特別是某些低端GPU。

我懷疑2D範圍會幫助你，因爲你的工作本質上是一維的。

例如，使用組大小或256，其中每個工作項處理256個元素。一個小組然後將負責65536個元素。總數將是1526個由390656個人工作項目組成的小組。

從my answer to your other question的簡單內核將正常工作。 n仍然是你的元素總數。工作項ID和工作組ID都未使用。

__kernel void vecAdd( __global double *a, __global double *b, __global double *c, const unsigned int n) 
{           
    //Get our global thread ID and global size 
    int gid = get_global_id(0);    
    int gsize = get_global_size(0);    

    //check vs n using for-loop condition 
    for(int i=gid; i<n; i+= gsize){ 
    c[i] = a[i] + b[i];    
    } 
} 

下面的內核將處理每個工作組的65536個元素塊。這具有共享全局讀取的優點，因爲它們比以前的內核更經常地與相鄰的存儲器地址相鄰。您應該使用ELEMENTS_PER_GROUP，以及您的工作組大小來查找硬件的最佳值。

#define ELEMENTS_PER_GROUP 65536 //max # of elements to be processes by a group 

__kernel void vecAddLocalized( __global double *a, __global double *b, __global double *c, const unsigned int n) 
{           
    int wgId = get_group_id(0); 
    int wgSize = get_local_size(0); 
    int itemId = get_local_id(0); 

    int startIndex = wgId * ELEMENTS_PER_GROUP; 
    int endIndex = startIndex + ELEMENTS_PER_GROUP; 
    if(endIndex > n){ 
    endIndex = n; 
    } 

    for(int i=startIndex + itemId; i<endIndex; i+= wgSize){ 
    c[i] = a[i] + b[i];    
    } 
} 

Answer 2

除非我有不同的設備，我可以分解的工作，我會一次排入整個線性範圍。 OpenCL驅動程序將連續地處理通過硬件並行處理的數據塊，而不是我們所能做到的。此外，他們也有可能以這種方式來實現最佳內存訪問行爲。

關於選擇2D vs 1D範圍，我不會將我的問題更改爲2D，只是爲了排列更小的範圍。然而，如果你的問題本質上是二維的（例如二維圖像或二維矩陣），那麼以這種方式來表示它是有意義的。

編輯：關於內核循環的mfa的觀點是一個非常有效的方法來將更大的範圍壓縮成更小的範圍。如果你的內核非常短（只有幾條語句），這也會非常有用，因爲它可以降低所有這些工作硬件線程的啓動成本。