Optimisation des échanges d'éléments vectoriels en utilisant CUDA

Question

Depuis que je suis nouveau à cuda .. J'ai besoin de votre aide aimable J'ai ce vecteur long, pour chaque groupe de 24 éléments, je dois faire ce qui suit: pour les 12 premiers éléments, les éléments pairs sont multipliés par -1, pour les 12 éléments suivants, les éléments impairs sont multipliés par -1, puis l'échange suivant a lieu:

Graphique: parce que je n'ai pas encore assez de points, ne pouvait pas afficher l'image si elle est ici:

http://www.freeimagehosting.net/image.php?e4b88fb666.png

J'ai écrit ce morceau de code, et je me demande si vous pourriez me aider à optimiser davantage pour résoudre les conflits de divergence bancaires ..

// est un multiple sous-vecteur de 24, sont Nds SDM et la mémoire partagée

____shared____ double Mds [sous-vecteur];

____shared____ double Nds [sous-vecteur];

int tx = threadIdx.x; 
int tx_mod = tx^0x0001; 
int basex = __umul24(blockDim.x, blockIdx.x); 

Mds[tx] = M.elements[basex + tx]; 
__syncthreads(); 

// flip the signs 
if (tx < (tx/24)*24 + 12) 
{ 
    //if < 12 and even 
    if ((tx & 0x0001)==0) 
    Mds[tx] = -Mds[tx]; 
} 
else 
if (tx < (tx/24)*24 + 24) 
{ 
    //if >12 and < 24 and odd 
    if ((tx & 0x0001)==1) 
    Mds[tx] = -Mds[tx]; 
} 

__syncthreads();

if (tx < (tx/24)*24 + 6) 
{ 
//for the first 6 elements .. swap with last six in the 24elements group (see graph) 
    Nds[tx] = Mds[tx_mod + 18]; 
    Mds [tx_mod + 18] = Mds [tx]; 
    Mds[tx] = Nds[tx]; 
} 
else 
if (tx < (tx/24)*24 + 12) 
{ 
    // for the second 6 elements .. swp with next adjacent group (see graph) 
    Nds[tx] = Mds[tx_mod + 6]; 
    Mds [tx_mod + 6] = Mds [tx]; 
    Mds[tx] = Nds[tx]; 
} 

__syncthreads();

Merci à l'avance ..

Answer 1

paul vous a donné de très bons points de départ vous des questions précédentes.

quelques choses à surveiller: vous faites une division non-base 2 qui est chère. Essayez plutôt d'utiliser la nature multidimensionnelle du bloc de threads. Par exemple, faites la dimension x de la taille 24, ce qui éliminera le besoin de division.

En général, essayez d'ajuster les dimensions du bloc de filetage afin de refléter les dimensions de vos données. Simplifier le basculement de signe: par exemple, si vous ne voulez pas inverser le signe, vous pouvez toujours le multiplier par l'identité 1. Déterminer comment mapper les nombres pairs/impairs à 1 et -1 en utilisant juste arithmétique: par exemple sign = (even*2+1) - 2 où même est 1 ou 0.