2015-07-20 58 views
0

我正在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.*/ 
    ........................................................ 

} 

但上面的代碼,因爲這會降低性能的條件語句像發散的問題。

可以對代碼進行哪些修改,以便整個數組可以有效地複製到每個工作組的共享內存中?

任何建議,非常感謝。

回答

0

根據您正在運行的設備,您很可能會完全忽略本地內存。如果你使用的是桌面GPU,他們使用幾乎沒有任何緩存,這使得使用本地內存非常重要,但是現在它們有相當的數量。如果你在gpu上的內存中找到相同的部分,它們全都在緩存中(它與共享內存大小一樣),它與本地內存一樣快(它們是同一塊內存,只是分割) 。手動將其拷貝到本地內存可能會額外強加一個小的性能損失

如果你不是在桌面GPU(ARM /等),或者你的要求使這種不切實際的,async_work_group_copy可能是你正在尋找

在一個不相關的說明中,上面的代碼只需要做一個屏障(CLK_LOCAL_MEM_FENCE),因爲你大概沒有修改你的輸入