Puis-je appeler une "macro fonctionnelle" dans un fichier d'en-tête d'une fonction CUDA __global__?

Question

Cela fait partie de mon fichier d'en-tête aes_locl.h:

. 
. 
# define SWAP(x) (_lrotl(x, 8) & 0x00ff00ff | _lrotr(x, 8) & 0xff00ff00) 
# define GETU32(p) SWAP(*((u32 *)(p))) 
# define PUTU32(ct, st) { *((u32 *)(ct)) = SWAP((st)); } 
. 
. 

Maintenant, à partir du fichier .cu je l'ai déclaré une fonction __ global__ et inclus le fichier d'en-tête comme ceci:

#include "aes_locl.h" 
..... 
__global__ void cudaEncryptKern(u32* _Te0, u32* _Te1, u32* _Te2, u32* _Te3, unsigned char* in, u32* rdk, unsigned long* length) 
{ 
    u32 *rk = rdk; 
    u32 s0, s1, s2, s3, t0, t1, t2, t3; 

    s0 = GETU32(in + threadIdx.x*(i))^rk[0]; 
} 

Cela me conduit à ce qui suit message d'erreur:

error: calling a host function from a __ device__/__ global__ function is only allowed in device emulation mode

J'ai l'exemple de code où le pro grammer appelle la macro exactement de cette manière. Est-ce que je peux l'appeler de cette façon, ou est-ce que ce n'est pas possible du tout? Si ce n'est pas, j'apprécierai quelques conseils de ce qui serait la meilleure approche pour réécrire les macros et assigner la valeur désirée à S0.

merci beaucoup d'avance !!!

Answer 1

Le matériel n'a pas d'instruction de rotation intégrée, et donc il n'y a pas d'intrinsèque pour l'exposer (vous ne pouvez pas exposer quelque chose qui n'existe pas!).

Il est assez simple à mettre en œuvre des changements et des masques si, par exemple, si x est de 32 bits puis à rotation à gauche huit bits que vous pouvez faire:

((x << 8) | (x >> 24)) 

Où x << 8 va pousser tout reste huit bits (c.-à- en supprimant les huit bits les plus à gauche), x >> 24 poussera tout droit à 34 bits (c'est-à-dire en supprimant tous les huit bits sauf les plus à gauche), et en les combinant OR donne le résultat dont vous avez besoin.

// # define SWAP(x) (_lrotl(x, 8) & 0x00ff00ff | _lrotr(x, 8) & 0xff00ff00) 
# define SWAP(x) (((x << 8) | (x >> 24)) & 0x00ff00ff | ((x >> 8) | (x << 24)) & 0xff00ff00) 

Vous pouvez bien sûr rendre plus efficace en reconnaissant que ce qui précède est surpuissant:

# define SWAP(x) (((x & 0xff00ff00) >> 8) | ((x & 0x00ff00ff) << 8)) 

Answer 2

L'erreur indique quel est réellement le problème. Vous appelez une fonction/macro définie dans un autre fichier (qui appartient au code CPU) depuis la fonction CUDA. C'est impossible!

Vous ne pouvez pas appeler CPU functions/macros/code à partir d'une fonction GPU.

Vous devez mettre vos définitions (ne _lrotl() existent dans CUDA?) Dans le même fichier qui sera compilé par nvcc.

Answer 3

Je pense que le problème n'est pas les macros elles-mêmes - le processus de compilation utilisé par nvcc pour le code CUDA exécute le préprocesseur C de la manière habituelle et donc utiliser les fichiers d'en-tête de cette manière devrait être bon. Je crois que le problème est dans vos appels à _lrotl et _lrotr.

Vous devriez être en mesure de vérifier que c'est effectivement le problème en supprimant temporairement ces appels.

Vous devriez consulter le guide de programmation CUDA pour voir les fonctionnalités dont vous avez besoin pour remplacer ces appels sur le GPU.