Optimisation d'un noyau de manipulation Bit-Wise

Question

J'ai le code suivant qui parcourt progressivement une chaîne de bits et les réorganise en blocs de 20 octets. J'utilise 32 * 8 blocs avec 40 fils par bloc. Cependant, le processus prend quelque chose comme 36ms sur mon GT630M. Y a-t-il une optimisation supplémentaire que je peux faire? Surtout en ce qui concerne la suppression de l'if-else dans la boucle la plus interne.

__global__ void test(unsigned char *data) 
{ 
    __shared__ unsigned char dataBlock[20]; 
    __shared__ int count; 
    count = 0; 

    unsigned char temp = 0x00; 

    for(count=0; count<(streamSize/8); count++) 
    { 
     for(int i=0; i<8; i++) 
     { 
      if(blockIdx.y >= i) 
       temp |= (*(data + threadIdx.x*(blockIdx.x + gridDim.x*(i+count)))&(0x01<<blockIdx.y))>>(blockIdx.y - i); 
      else 
       temp |= (*(data + threadIdx.x*(blockIdx.x + gridDim.x*(i+count)))&(0x01<<blockIdx.y))<<(i - blockIdx.y); 
     } 
     dataBlock[threadIdx.x] = temp; 
      //do something 

    } 

} 

Answer 1

Il est pas clair ce que votre code tente d'accomplir, mais un couple des opportunités évidentes sont:

1) si possible, utilisez des mots de 32 bits au lieu de char non signé.

2) utiliser des tailles de blocs qui sont des multiples de 32.

3) Le code conditionnel ne peut vous coûter autant que vous attendez. Vous pouvez vérifier en compilant avec --cubin --gpu-architecture sm_xx (où xx est la version SM de votre matériel cible), et en utilisant cuobjdump --dump-sass sur le fichier cubin résultant pour regarder l'assemblage généré. Vous devrez peut-être modifier le code source pour loft la sous-expression commune dans une variable distincte, et/ou utiliser l'opérateur ternaire? : pour suggérer au compilateur d'utiliser la prédication.