我想了解如何在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上的處理還是有意義的,還是因爲它們共享高速緩存而造成的開銷?
問候, 斯特凡
謝謝!關於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