本地內存Opencl大小對速度有影響嗎？

我是OpenCL的新手，我正在嘗試計算灰度圖像的直方圖。我正在GPU nvidia GT 330M上執行此計算。本地內存Opencl大小對速度有影響嗎？

代碼

__kernel void histogram(__global struct gray * input, __global int * global_hist, __local volatile int * histogram){ 
    int local_offset = get_local_id(0) * 256; 
    int histogram_global_offset = get_global_id(0) * 256; 
    int offset = get_global_id(0) * 1920; 
    int value; 
    for(unsigned int i = 0; i < 256; i++){ 
     histogram[local_offset + i] = 0; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 

    for(unsigned int i = 0; i < 1920; i++){ 
     value = input[offset + i].i; 
     histogram[local_offset + value]++; 
    } 

    barrier(CLK_LOCAL_MEM_FENCE); 
    for(unsigned int i = 0; i < 256; i++){ 
     global_hist[histogram_global_offset + i] = histogram[local_offset + i]; 
    } 
}

此計算圖像1920 * 1080進行。

當直方圖的局部尺寸設置爲256 *的sizeof（cl_int）這種計算速度（通過NVIDIA nsight性能分析）11 675微秒，我燒的內核與

queue.enqueueNDRangeKernel(kernel_histogram, cl::NullRange, cl::NDRange(1080), cl::NDRange(1));

。

由於本地工作組大小設置爲1。我嘗試將本地工作組大小增加到8.但是，當我將直方圖的局部大小增加到256 * 8 * sizeof（cl_int）並使用本地wg大小1進行計算時，我得到85 177微秒。

所以，當我用每個工作組的8個內核啓動它時，我不會從11ms，而是從85ms加速。因此，每個Worgroup擁有8個內核的最終速度爲13 714微秒。

但是，當我創建計算錯誤，將local_offset設置爲零，並且本地直方圖的大小爲256 * sizeof（cl_int）並且每個工作組使用8個內核時，我獲得了更好的時間 - 3 854微秒。

有沒有人有一些想法來加速這種計算？

謝謝！

來源

2014-05-08 Tomas Lysek

您發佈的示例代碼似乎正在處理每個工作項圖像的一行。當你爲每個工作組啓動多個工作項目時，你還在這麼做，還是讓每個工作組合作處理一行？當你這樣做時，你能向我們展示你使用的內核和主機代碼嗎？ – jprice

感謝您的回覆。是的，我正在使用這種方法，一個線程計算一行。 –

這會導致較差的內存訪問性能。理想情況下，您希望您的內存訪問得以合併 - 也就是說，相鄰的工作項目應儘可能訪問相鄰的內存地址。 – jprice

這個答案假定你想最終將你的直方圖降低到256個整型值。您可以使用盡可能多的工作組來調用內核，因爲您的設備上具有計算單元，組大小應該（總是）是設備上CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE的倍數。

__kernel void histogram(__global struct gray * input, __global int * global_hist){ 
    int group_id = get_group_id(0); 
    int num_groups = get_num_groups(0); 
    int local_id = get_local_id(0); 
    int local_size = get_local_size(0); 

    volatile __local int histogram[256]; 

    int i; 
    for(i=local_id; i<256; i+=local_size){ 
    histogram[i] = 0; 
    } 

    int rowNum, colNum, value, global_hist_offset 

    for(rowNum = group_id; rowNum < 1080; rowNum+=num_groups){ 
    for(colNum = local_id; colNum < 1920; colNum += local_size){ 
     value = input[rowNum*1920 + colNum].i; 
     atomic_inc(histogram[input]); 
    } 
    } 

    barrier(CLK_LOCAL_MEM_FENCE); 
    global_hist_offset = group_id * 256; 
    for(i=local_id; i<256; i+=local_size){ 
    global_hist[global_hist_offset + i] = histogram[i]; 
    } 

}

每個工作組一次在圖像的一行上協同工作。然後組移動到另一行，使用num_groups值計算。無論您擁有多少個團體，這都可以很好地工作。例如，如果您有7個計算單位，則組3（第四組）將在圖像的第3行開始，然後在第7行開始。第3組將總共計算153行，其最後一行將是第1074行。在此示例中，某些工作組可能會計算多一個行組0和1。

當查看圖像的列時，工作組內將完成相同的隔行掃描。在colNum循環中，第N個工作項從第N列開始，並由local_size列向前跳過。此循環的其餘部分不應該經常播放，因爲CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE可能是1920的一個因素。嘗試所有工作組大小（1..X）* CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE，直到設備的最大工作組大小。

關於此內核的最後一點：結果與您的原始內核不完全相同。你的global_hist數組是1080 * 256的整數。我有一個需要是num_groups * 256整數。如果你想要完全減少，這會有所幫助，因爲在內核執行後添加的內容要少得多。

來源

2014-05-09 12:37:24 mfa

本地內存Opencl大小對速度有影響嗎？

回答

相關問題