我正在GPU上使用Open CL實現算法。OpenCL:將全局內存複製到每個工作組的本地內存
目前,我發動內核包含128個工作項目在全局內存 .The數據僅供一個工作組的各項工作,項目被多次使用。要採取共享內存的速度優勢,我把它抄了到共享內存使用下面的代碼。
__kernel void kernel1(__global float2* input,
__global int* bin,
__global float2* DFT,
__local float2* localInput,
__const int N){
size_t itemId = get_local_id(0);
localInput[itemId] = input[itemId];
barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
........................................................
/*Remaining algo here.*/
........................................................
}
上面的代碼工作得很好,如果僅存在一個工作group.But如果有多於一個的工作組,假設有兩個工作組以相等的數量的項目在他們每個人以上內核只複製第一個工作組共享內存中的前半部分和後一箇中的後半部分。
我也嘗試了下面的內核:
__kernel void kernel1(__global float2* input,
__global int* bin,
__global float2* DFT,
__local float2* localInput,
__const int N){
size_t itemId = get_local_id(0);
if(itemId == 0){
for(int index = 0;index<N;index++){
localInput[index] = input[index];
}
}
barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
........................................................
/*Remaining algo here.*/
........................................................
}
但上面的代碼,因爲這會降低性能的條件語句像發散的問題。
可以對代碼進行哪些修改,以便整個數組可以有效地複製到每個工作組的共享內存中?
任何建議,非常感謝。