如何在OpenCL中使用async_work_group_copy？

我想了解如何在OpenCL中正確使用async_work_group_copy（）調用。讓我們在一個簡單的例子來看看：如何在OpenCL中使用async_work_group_copy？

__kernel void test(__global float *x) { 
    __local xcopy[GROUP_SIZE]; 

    int globalid = get_global_id(0); 
    int localid = get_local_id(0); 
    event_t e = async_work_group_copy(xcopy, x+globalid-localid, GROUP_SIZE, 0); 
    wait_group_events(1, &e); 
}

參考http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/async_work_group_copy.html說：「執行從SRC num_elements gentype元素的異步複製到DST異步副本由所有工作項目在一個工作組進行，這個內置函數因此必須被執行內核的工作組中的所有工作項所遇到，並且具有相同的參數值;否則結果是不確定的。「

但是，這並不清楚我的問題......

我想知道，如果下面的假設是正確的：

的調用async_work_group_copy（）必須由所有工作執行 - 組中的項目。
調用的方式應該是，所有工作項的源地址都相同，並指向要複製的存儲區的第一個元素。
由於我的源地址是相對基於工作組中第一個工作項的全局工作項ID。所以我必須減去本地ID以使所有工作項的地址相同...
第三個參數是否是真正的元素數量（而不是字節大小）？

獎金的問題：

一個。我可以使用barrier（CLK_LOCAL_MEM_FENCE）而不是wait_group_events（）並忽略返回值嗎？如果是這樣，那可能會更快嗎？

b。本地副本對於CPU上的處理還是有意義的，還是因爲它們共享高速緩存而造成的開銷？

問候，斯特凡

來源

2013-03-21 SDwarfs

其中一個主要原因此功能是存在的，允許驅動器/內核編譯器有效，而無需進行有關硬件的假設開發商的內存複製。

您可以描述需要複製的內存，就像它是單線程副本一樣，async_work_group_copy可以使用並行硬件爲您完成。

爲了您的具體問題：

我從來沒有見過async_work_group_copy只用了一些組中的工作項目。我一直認爲這是因爲它需要。我認爲wait_group_events的阻塞性質會強制所有工作項目成爲副本的一部分。
是的。源（和目標）地址對於所有工作項目都必須相同。
你可以減去你的本地ID來得到正確的地址，但是我發現基於groupId的地址解決了這個問題。（get_group_id）
是的。最後一個參數是元素的數量，而不是字節的大小。

a。不可以。基於事件，您會發現您的屏障幾乎立即被工作項目擊中，數據不一定會被複制。這是有道理的，因爲一些opencl硬件甚至可能根本不使用計算單元來執行實際的複製操作。

b。我認爲，使用本地內存時，cpu opencl實現可能會保證L1緩存使用率。唯一可以肯定的是，如果這種方法性能更好，則可以使用各種設置對應用程序進行基準測試

來源

2013-03-21 13:10:49 mfa

謝謝！關於3.get_group_id（） - 很好的提示，但在這個特殊的需要乘以工作組的大小，即使對於已知的工作組大小爲2^n這可能會加快左移。。但是，在其他情況下，這可能是非常有用的知道它也存在！感謝您分享您的體驗！ --- Stefan – SDwarfs 2013-03-21 15:05:54

我對你對「a」的回答有一個疑問：據我所知，複製應至少由一個工作項目處理。由於工作組中的所有其他節點都需要準備好本地數據，因此我認爲它們都等待數據可用。等待事件應該做到這一點。但是障礙應該保證是一樣的，因爲小組的所有其他工作項目都會等待複製過程中的工作項目。障礙可能更簡單（每個工作組最多1個障礙，而每個工作項目最多等待1個事件），因此速度更快。 – SDwarfs 2013-03-21 15:21:14

re：'a'...未來可能存在一個處理異步拷貝而不使用任何工作項目的設備。如果至少有一個工作項目被複制操作阻止，則障礙將起作用。 – mfa 2013-03-21 18:31:32

如何在OpenCL中使用async_work_group_copy？

回答

相關問題