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

OpenCL copy character from global to local memory

本文关键字：内存 复制 全局 字符 OpenCL      更新时间：2023-10-16

我在OpenCL技术中实现了sha512。我对内核函数有简单的定义

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

在主机上，我已经实现并成功测试了sha512算法的实现。

我在将数据从数组复制到名为character的临时变量时遇到问题message。

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

默认情况下，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声明为__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

根据实现、硬件等，您可能会发现这比字节访问更快。(如果您不确定是否所有部署平台都是小端序，则可能需要通过使用clGetDeviceInfo读取CL_DEVICE_ENDIAN_LITTLE属性来检查是否需要反向解压缩。

相关文章：