Caractère de copie OpenCL de la mémoire globale vers la mémoire locale

Question

Je suis en train d'implémenter sha512 dans la technologie OpenCL. J'ai la définition simple de la fonction du noyau

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

Sur l'hôte j'ai implémenté et testé avec succès l'implémentation de l'algorithme sha512.

J'ai un problème avec les données de copie de message tableau à variable temporaire appelé character.

char character = message[i]; 

Où i est une variable de boucle - dans la plage de 0 à la taille du message.

Quand j'ai essayé d'exécuter mon programme là-bas j'ai eu cette erreur

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) 

Je readed à propos async_work_group_copy(), mais je ne peux pas comprendre comment l'utiliser - dans docs Je ne peux pas trouver un code d'exemple.

J'ai essayé avec char character = (__private char) message[i]; mais cela ne fonctionne pas aussi.

Je ne comprends pas comment passer le dernier paramètre en async_work_group_copy() et comment l'utiliser pour copier des données de la mémoire __global dans la mémoire __private.

Answer 1

OpenCL par défaut ne permet pas l'accès à un octet dans les noyaux: l'accès à la mémoire doit être multiple de 4 octets, aligné sur les limites de 4 octets. Si votre implémentation le prend en charge, vous pouvez activer des accès mémoire par octet. Cela implique le cl_khr_byte_addressable_store extension, que vous devez vérifier et activer explicitement dans votre source de noyau. Essayez-le et voyez s'il résout votre problème.

Pour utiliser async_work_group_copy, essayer quelque chose comme ceci:

#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 

Notez que async_work_group_copy n'est pas nécessaire; vous pouvez accéder directement à la mémoire globale. Ce qui sera le plus rapide dépend de votre noyau, de l'implémentation OpenCL et du matériel.

Une autre option (la seule option si votre implémentation/matériel ne supporte pas cl_khr_byte_addressable_store) est d'extraire vos données en morceaux d'au moins 4 octets. Déclarez votre message comme __global uint* et déballer les octets en déplaçant et le masquage:

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 

Selon l'implémentation, le matériel, etc., vous trouverez peut-être ce plus rapide que l'accès par octet. (Vous voudrez peut-être check if you need to reverse the unpacking by reading the CL_DEVICE_ENDIAN_LITTLE property using clGetDeviceInfo si vous n'êtes pas sûr si toutes vos plates-formes de déploiement seront little-endian.)