的OpenCL - 動態的循環

如何for循環中的OpenCL內核執行設備上時，他們的邊界是動態的執行，也就是當每個工作項的for循環執行不同數額的時間？的OpenCL - 動態的循環

AFAIK，內核是指令集（或更好地說一個流）的指令。 GPU設備是一組獨立的計算單元（流式多處理器-SM），每個計算單元包含多個計算單元（流處理器-SP）。

每個SM可以加載一個指令（對於不同的SMS，這可能是不同的指令）從內核（即指令流），並執行所加載的指令爲儘可能多的工作項，因爲在當前SM的SP（每個SP操作相同的指令但具有不同的數據 - SIMD）。

一個SM中的所有SP都必須運行相同的指令，因此在執行for-loop條件後，必須根據每個工作條件的結果做出動態決策，下一條指令是什麼將在SM 和上運行，以便爲其運行的工作項目運行。

基於這個假設，我認爲foobaz內核（見下文）會更快地執行，因爲當一個工作項完成執行時，另一個工作項可以取代它。

這是假設錯了嗎？

以下哪兩個內核的，foobar和foobaz，最終將執行得更快？性能取決於什麼？（一個元素的屬性數量可以比其他元素大一個數量級）。

foobar;

__kernel void foobar(__global int* elements,   /* size N   */ 
        __global int* element_properties, /* size N*constant */ 
        __global int* output)    /* size N   */ 
{ 
    size_t gid = get_global_id(0); 
    int reduced = 0; 

    for (size_t i=N*gid; i<N+N*gid; i++) 
     reduce += predict_future_events(reduce, element_properties[i]); 


    output[gid] = reduced; 
}

...和foobaz;

__kernel void foobaz(__global int* elements,     /* size N  */ 
        __global int* element_properties,   /* size upper-bounded */ 
        __global int2* element_properties_ranges, /* size N  */ 
        __global int* output)      /* size N  */ 
{ 
    size_t gid = get_global_id(0); 
    int reduced = 0; 

    // `range.x` = starting index in `element_properties` 
    // `range.y` = ending index in `element_properties` 
    int2 range = element_properties_ranges[gid]; 

    for (size_t i=range.x; i<range.y; i++) 
     reduce += predict_future_events(reduce, element_properties[i]); 


    output[gid] = reduced; 
}

來源

2016-07-31 sarasvati

假設它的OpenCL 1.2設備，

如果每個「predict_future_events」是在性能方面亂，你可以檢查一些「硬件優化」的轉變。你可以在同一時間丟2級不同的內核（兩個不同的全穀粒（N），如果他們可以分開/獨立的），或者你可以推着半內核（N/2）的「恆定版」和下半年的不同的內核（因爲這不會是從計算的第一個例子不同），也許驅動程序可以處理某些情況下一個內核延遲永遠但至少另一半獲取計算資源（如驅動程序可以做到這一點）。因此，更多的管道將忙於做一些事情，並最終爲內核提供更好的時機。

除此之外，具有每功能的隨機延遲使得難以預見到的功能組中的環給出什麼的總延遲時間，以便使所有線程的步驟相等數目的（如在第一實施例/常數）是更容易「假設「線程之間的平衡負載會有更大的機會。例如，對於光線跟蹤內核，1000深度折射+ 1000深度反射將足夠混亂，因此您可以給每個線程1個光線進行計算，因爲您無法知道光線是否會折射或反映在下一個表面（如果有的話）。也許分組更近的可以更頻繁地使用L1-L2緩存。

針對OpenCL 2.0裝置，您可以生成內核線程內的多個線程/組應該使這個更具活力。

來源

2016-07-31 21:38:01

您在兩種解決方案中都做了幾乎相同的事情。我敢打賭，他們幾乎必須在同一時間完成。

如果你想更快，使用使用INT4爲您參數SIMD功能，減少變量，然後調整predict_future_events函數來處理INT4值。通過這種方式，您可以獲得高達4倍的性能，因爲每條指令都並行處理4個元素。

根據您的硬件，您可以使用INT8或高達INT16。

順便說一句：我不看到分配到N個變量也不看到任何使用元件陣列組成。

來源

2016-07-31 20:20:17 user6656519

我也許應該指出，這兩個內核都只是僞代碼（所以絕對加速是不是我關心的話）。核心部分卻是，如何將*數據依賴*爲每個線程環路通行證的數量*對*迭代每個線程不變（但高數）影響性能，以及是否我的怎麼做硬件的假設（ GPU）處理for-loop是正確的。 – sarasvati

什麼部分是動態的？ range.y？ –

@huseyintugrulbuyukisik是的。嗯，其實我不知道任何事前既不顆粒'range.x'也不'range.y'，所以這兩個值是動態的（我希望我們通過「動態」意味着同樣的事情）。假設'element_properties_ranges'中的值是在執行'foobar'和'foobaz'之前由另一個內核產生的。 – sarasvati

的OpenCL - 動態的循環

回答

相關問題