1

我在 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内存中。

4

1 回答 1

1

OpenCL 默认情况下不允许内核中的单字节访问：内存访问需要是 4 字节的倍数，与 4 字节边界对齐。如果您的实现支持它，您可以启用逐字节内存访问。这涉及cl_khr_byte_addressable_store 扩展，您需要在内核源代码中检查并显式启用它。试一试，看看它是否能解决您的问题。

要使用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为 a__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，您可能需要通过读取属性来检查是否需要反转解包。）CL_DEVICE_ENDIAN_LITTLEclGetDeviceInfo

于 2017-08-09T10:16:55.770 回答