5
votes

Comment transformer chaque bit en octet

J'ai le code suivant pour transformer un peu en octet.

__device__ UINT64 bitToByte(const UINT8 input) {
    UINT64 b = ((0x8040201008040201ULL * input) >> 7) & 0x0101010101010101ULL; 
    //reverse the byte order <<-- this step is missing
    return b;
}

Cependant les octets sont dans le mauvais ordre, l'endianité est inversée. Sur le processeur, je peux simplement utiliser un bswap reg, reg pour résoudre ce problème, mais que dois-je faire sur le GPU?

Sinon, quelle astuce similaire puis-je utiliser pour que le les octets sont placés dans le bon sens, c'est-à-dire que le bit le plus significatif va à l'octet le plus significatif, de sorte que je n'ai pas besoin d'une astuce bswap.

c++ cuda bit-manipulation

4 commentaires

Une pensée naïve peut-être, mais ne pouvez-vous pas faire d'abord l'étape bswap reg, reg , pour que l'endianness soit correcte, puis faire la conversion bit à octet.

@JoeyMallone, je n'ai pas d'instruction bswap, tout mon code fonctionne sur le GPU. D'où la balise CUDA et l'annotation __device__ sur la fonction.

L'instruction prmt PTX permet de faire cela. Je n'ai pas le temps de vous trouver les liens maintenant, mais je les ai utilisés dans cette réponse , ce qui pourrait donner vous avez juste assez d'informations pour le résoudre vous-même.

@tera, ah yess, prmt correspond à l'instruction intrinsèque __byte_perm (a32, b32, swap) . Ça marchera.

3 Réponses :

2
votes

Au lieu d'inverser le résultat, vous pouvez inverser input , avec l'une des astuces expliquées ici . Par exemple, en utilisant l'approche de cette réponse :

static UINT8 lookup[16] = {
    0x0, 0x8, 0x4, 0xc, 0x2, 0xa, 0x6, 0xe,
    0x1, 0x9, 0x5, 0xd, 0x3, 0xb, 0x7, 0xf, };

UINT8 reverse(UINT8 n) {
    return (lookup[n & 0xF] << 4) | lookup[n >> 4];
}

__device__ UINT64 bitToByte(const UINT8 input) {
    UINT64 b = ((0x8040201008040201ULL * reverse(input)) >> 7) & 0x0101010101010101ULL; 
    return b;
}

1 commentaires

Je préfère vraiment ne pas utiliser de table de consultation.

3
votes

Merci à @tera, voici la réponse:

//Expand every bit into a byte
__device__ static UINT64 Add012(const UINT8 input) {
    const UINT32 reversed = (__brev(input) >> 24);
    return ((0x8040201008040201ULL * reversed) >> 7) & 0x0101010101010101ULL; //extract every bit into a byte
}

Le __byte_perm remplace l'instruction bswap .

Alternativement, l'entrée peut être inversée à l'aide de __brev (bit-reverse) intrinsèque :

//Expand every bit into a byte
__device__ static UINT64 Add012(const UINT8 input) {
    const UINT64 b = ((0x8040201008040201ULL * input) >> 7) & 0x0101010101010101ULL; //extract every bit into a byte
    //unfortunatly this returns the wrong byte order
    UINT32* const b2 = (UINT32*)&b;
    UINT64 Result;
    UINT32* const Result2 = (UINT32*)&Result;
    Result2[0] = __byte_perm(b2[0]/*LSB*/, b2[1], 0x4567);  //swap the bytes around, the MSB's go into the LSB in reverse order
    Result2[1] = __byte_perm(b2[0]/*LSB*/, b2[1], 0x0123);  //and the LSB -> MSB reversed.
    return Result;
}

La deuxième version semble plus simple.

0 commentaires

2
votes

Pour inverser l'ordre des octets, l'extraction des bits peut être effectuée avec la même astuce, mais en échangeant les coefficients qui effectuent le décalage de la multiplication. Cependant, pour éviter les conflits dans la multiplication, il faut le faire en deux étapes, pour les bits pairs et impairs. De cette façon, 2 octets sont libres de contenir le résultat de chaque multiplication, ce qui est suffisant pour assurer l'intégrité du résultat.

__device__ UINT64 bitToByte(const UINT8 input) {
  UINT64 b =  ( ((0x0002000800200080ULL * input) & 0x0080008000800080ULL) 
              | ((0x0100040010004000ULL * input) & 0x8000800080008000ULL) )
                >> 7 ;
    return b;
}

Comme indiqué dans les commentaires, pour optimiser, les décalages peuvent être factorisés .

__device__ UINT64 bitToByte(const UINT8 input) {
  UINT64 b = ( ((0x0002000800200080ULL * input) >> 7) & 0x0001000100010001ULL) 
          |  ( ((0x0100040010004000ULL * input) >> 7) & 0x0100010001000100ULL);

    return b;
}

4 commentaires

Agréable! J'ai essayé cela, mais j'ai rencontré des conflits avec les multiplicandes débordant dans l'octet adjacent. Cela résout cela.

Logique. Dans 0x8040201008040201 , les 1 bits sont séparés par 9, mais dans 0x0102040810204080 ils ne sont séparés que de 7. De cette façon, vous les séparez de 14 bits.

Vous pouvez éliminer deux opérations en factorisant le «>> 7» dans les constantes ici. Mais l'optimiseur pourrait également le repérer.

@MSalters, c'est vrai! Mais les masques doivent être modifiés et je doute que l'optimiseur puisse le faire. Mise à jour de la réponse avec ces suggestions.