OpenCL 将字符从全局复制到本地内存答案

【问题标题】：OpenCL copy character from global to local memoryOpenCL 将字符从全局复制到本地内存
【发布时间】：2017-08-09 18:24:21
【问题描述】：

我正在 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)

我阅读了有关 async_work_group_copy() 的信息，但我不明白如何使用它 - 在文档中我找不到任何示例代码。

我已尝试使用 char character = (__private char) message[i];，但它也不起作用。

我不明白如何将最后一个参数传递给async_work_group_copy()，以及如何使用它将数据从__global 内存复制到__private 内存。

【问题讨论】：

标签： c++ c opencl opencl-c

【解决方案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

根据实现、硬件等，您可能会发现这比按字节访问要快。（如果您不确定您的所有部署平台是否都是 little-endian，您可能需要 check if you need to reverse the unpacking by reading the CL_DEVICE_ENDIAN_LITTLE property using clGetDeviceInfo。）

【讨论】：