從全局到本地內存的OpenCL副本字符

Question

我在OpenCL技術中實現sha512。我有內核函數

__kernel void _sha512(__global char *message, const uint length, __global char *hash); 

的簡單定義在主機上我已經實現並測試成功實施SHA512算法。

我有一個從message數組複製數據到臨時變量character的問題。

char character = message[i]; 

i哪裏是一個循環變量 - 在範圍從0到消息的大小。

當我試圖運行有我的計劃，我得到這個錯誤

0x00007FFD9FA03D54 (0x0000000010CD0F88 0x0000000010CD0F88 0x0000000010BAEE88 0x000000001A2942A0), nvvmCompilerProperty() + 0x26174 bytes(s) 
... 
0x00007FFDDFA70D51 (0x0000000000000000 0x0000000000000000 0x0000000000000000 0x0000000000000000), RtlUserThreadStart() + 0x21 bytes(s) 
0x00007FFDDFA70D51 (0x0000000000000000 0x0000000000000000 0x0000000000000000 0x0000000000000000), RtlUserThreadStart() + 0x21 bytes(s) 

我readed約async_work_group_copy()，但我不明白如何使用它 - 在文檔我沒有發現任何示例代碼。

我試過char character = (__private char) message[i];，但它不工作。

我不明白如何將最後一個參數傳遞到async_work_group_copy()以及如何使用它將__global內存中的數據複製到__private內存中。

Answer 1

OpenCL默認情況下不允許在內核中進行單字節訪問：內存訪問需要以4個字節的倍數，對齊到4個字節的邊界。如果您的實現支持它，則可以啓用逐字節內存訪問。這涉及cl_khr_byte_addressable_store extension，您需要檢查並在您的內核源代碼中明確啓用。試一試，看看它是否能解決你的問題。

要使用async_work_group_copy，嘗試這樣的事情：

#define LOCAL_MESSAGE_SIZE 64 // or some other suitable size for your workgroup 
__local char local_message[LOCAL_MESSAGE_SIZE]; 
event_t local_message_ready = async_work_group_copy(local_message, message, LOCAL_MESSAGE_SIZE, 0); 
// ... 

// Just before you need to use local_message's content: 
wait_group_events(1, &local_message_ready); 
// Use local_message from here onwards 

注意async_work_group_copy不是必需的;你可以直接訪問全局內存。哪個更快取決於你的內核，OpenCL實現和硬件。

另一種選擇（如果你的實現/硬件不支持cl_khr_byte_addressable_store，唯一的選擇）是以至少4個字節的塊來獲取你的數據。聲明你的message作爲__global uint*和移位和屏蔽解壓字節：

uint word = message[i]; 
char byte0 = (word & 0xff); 
char byte1 = ((word >> 8) & 0xff); 
char byte2 = ((word >> 16) & 0xff); 
char byte3 = ((word >> 24) & 0xff); 
// use byte0..byte3 in your algorithm 

根據實現方式，硬件等你會發現這是比字節訪問速度更快。 （如果不確定是否所有的部署平臺都是小端的話，您可能需要check if you need to reverse the unpacking by reading the CL_DEVICE_ENDIAN_LITTLE property using clGetDeviceInfo。）