Копирование символа OpenCL из глобальной в локальную память
Я реализую sha512 в технологии OpenCL. У меня есть простое определение функции ядра
__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 объем памяти.
Решение
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
В зависимости от реализации, оборудования и т. Д. Вы можете обнаружить, что это будет быстрее, чем байтовый доступ. (Вы можете захотеть проверьте, если вам нужно отменить распаковку, прочитав CL_DEVICE_ENDIAN_LITTLE использование имущества clGetDeviceInfo если вы не уверены, что все ваши платформы развертывания будут иметь порядок байтов.)
Другие решения
Других решений пока нет …