从全局到本地内存的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。）