OpenCL - 工作組軸可交換嗎？

Question

我正在努力爲一個問題尋找最佳的工作組規模，我想出了一些我無法爲自己辯解的事情。

這是我的結果：

• GlobalWorkSize {6400 6400 1}，WorkGroupSize {64 4 1}，時間（毫秒）= 44.18
• GlobalWorkSize {6400 6400 1}，WorkGroupSize {4 64 1 }，時間（毫秒）= 24.39

交換軸引起兩倍的執行速度。爲什麼！？

順便說一下，我使用的是AMD GPU。

感謝:-)

編輯： 這是內核（一個簡單的矩陣轉置）：

__kernel void transpose(__global float *input, __global float *output, const int size){ 
    int i = get_global_id(0); 
    int j = get_global_id(1); 
    output[i*size + j] = input[j*size + i]; 
} 

Answer 1

我@Thomas同意，這很可能取決於你的內核。最有可能的是，在第二種情況下，您可以通過合併的方式訪問內存和/或充分利用內存事務。

聚結：當線程需要訪問元素在存儲器中的硬件試圖訪問在這些元件儘可能少的交易，即如果線程0和線程1具有訪問連續元素，將僅存在一個事務。

完全使用內存事務：假設您有一個在一次事務中獲取32個字節的GPU。因此，如果您有4個線程需要獲取一個int，則您只使用事務獲取的一半數據;你浪費其餘的（假設int是4個字節）。

爲了說明這一點，假設您有一個n乘n的矩陣來訪問。你的矩陣在主要行中，並且你使用在一個維度中組織的n個線程。您有兩種可能性：

1. 每個工作項都負責一個列，每次循環一個列元素。
2. 每個工作項都需要一條線，每次循環一個線條元素。

這可能是違反直覺的，但第一個解決方案將能夠使第二個解決方案不會成爲聚結點。原因是當第一個工作項需要訪問第一列中的第一個元素時，第二個工作項將訪問第二個列中的第一個元素，依此類推。這些元素在內存中是連續的。第二種解決方案並非如此。

現在，如果你採用相同的例子，並應用解決方案1，但這次你有4個工作項目，而不是n和我剛剛講過的同一個GPU，因爲你很可能將時間增加了2倍你會浪費一半的內存交易。

編輯：現在你發佈你的內核，我看到我忘了提到別的東西。

對於你的內核，似乎選擇（1,256）或（256,1）的本地大小總是一個不好的選擇。在第一種情況下，需要256個事務來讀取一列（每次讀取32個字節，其中只有4個將被使用 - 請記住前面例子中的相同GPU），而輸出中需要32個事務寫入：您可以在一個事務中寫8個浮動，因此32個事務可以寫入256個元素。

這與工作組大小爲（256,1）的問題相同，但是這次使用32個事務讀取和256個寫入。

那麼爲什麼第一個尺寸更好？這是因爲有一個緩存系統，可以減輕讀取部分的不良訪問。因此，寫入部分的大小（1,256）是有利的，並且高速緩存系統處理不太好的讀取部分，減少了必需的讀取事務的數量。

請注意事務處理數量總體上減少（考慮到NDRange中的所有工作組）。例如，第一個工作組發佈256個事務，讀取第一列的256個第一個元素。第二個工作組可能只是進入緩存以檢索第二列的元素，因爲它們是由第一個工作組發佈的事務（32字節）提取的。

現在，我幾乎可以肯定，你可以做得比（1,256）嘗試（8,32）更好。