OpenCL內核的優化

Question

我在內核中使用大量無符號字符工作，我用clCreateBuffer創建了內存對象。比我通過clEnqueueWriteBuffer複製一塊無符號字符到這個內存對象。而且，我在循環中調用從這個內存對象讀取的內核，執行一些邏輯並將新數據寫入同一位置（在本週期中，我沒有調用clEnqueueWriteBuffer或clEnqueueReadBuffer）。下面是內核代碼：

__kernel void test(__global unsigned char *in, unsigned int offset) { 
    int grId = get_group_id(0); 
    unsigned char msg[1024]; 
    offset *= grId; 

    // Copy from global to private memory 
    size_t i; 
    for (i = 0; i < 1024; i++) 
     msg[i] = in[ offset + i ]; 

    // Make some computation here, not complicated logic  

    // Copy from private to global memory 
    for (i = 0; i < 1024; i++) 
     in[ offset + i ] = msg[i]; 
} 

當週期完成後（循環運行CCA 1000倍），那麼我讀通過clEnqueueReadBuffer內存對象的結果。

可以優化這段代碼嗎？

Answer 1

你可以嘗試矢量版本（uchar8而不是uchar），但是編譯器可以用這種方式優化它。 最重要的是您的代碼始終在進行實驗。

編輯

好像連uchar16現在支持： http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/vectorDataTypes.html

Answer 2

東西，首先穿越心靈，是您展開循環可以幫助你跳過狀態評估。您可以使用this pragma使其更容易。

使用NVIDIA的芯片共享存儲器還可以極大地幫助（如果您的當前本地MEM不使用共享內存默認情況下）

Answer 3

對於你需要解釋你做什麼樣的計算優化。性能的最大好處可以通過將計算分組到工作組中並讓它們在本地內存上工作來獲得。您需要特別注意私有內存（最小）和本地內存（小）的大小。

你的內核多久調用一次？所有內核都使用相同的數據嗎？人們可以考慮一個本地內存緩衝區，工作組中的所有線程都會將部分數據讀入本地內存，然後再共享數據。您需要注意同步。

我建議看看SDK供應商的樣品。我只知道nVidia SDK。那裏的樣本非常複雜，但非常有趣。

像float4shall這樣的向量類型的改變適用於ATI板。據說nVidia在標量和內部編譯器優化方面效果最佳。這是稍後用分析器進行微調的情況。你可以通過內存優化獲得大量的性能。

Answer 4

幾點建議：

• 在內核的開始做單in += get_group_id(0) * offset。
• 一次讀取4個字符（在uchar4或uint上工作）。
• 如果可能，一次也處理4個字符。
• 在每個線程中都有一個1K專用陣列，工作組大小和佔用率將受到嚴重限制，運行更多處理少量字符的線程效率可能更高。
• 似乎每個組中的所有線程都會處理完全相同的數據;它可能不是你想到的。