2013-04-27 4 views
1

J'ai un algorithme, effectuant une réduction parallèle à deux étapes sur GPU pour trouver le plus petit élément d'une chaîne. Je sais qu'il y a un indice sur la façon de le faire fonctionner plus rapidement, mais je ne sais pas ce que c'est. Des idées sur comment je peux régler ce noyau pour accélérer mon programme? Il n'est pas nécessaire de changer réellement l'algorithme, il peut y avoir d'autres astuces. Toutes les idées sont les bienvenues.boosting parallel reduction OpenCL

Merci!

__kernel 
void reduce(__global float* buffer, 
      __local float* scratch, 
      __const int length, 
      __global float* result) {  
    int global_index = get_global_id(0); 
    float accumulator = INFINITY 
     while (global_index < length) { 
      float element = buffer[global_index]; 
      accumulator = (accumulator < element) ? accumulator : element; 
      global_index += get_global_size(0); 
    } 
    int local_index = get_local_id(0); 
    scratch[local_index] = accumulator; 
    barrier(CLK_LOCAL_MEM_FENCE); 
    for(int offset = get_local_size(0)/2; 
     offset > 0; 
     offset = offset/2) { 
      if (local_index < offset) { 
       float other = scratch[local_index + offset]; 
       float mine = scratch[local_index]; 
       scratch[local_index] = (mine < other) ? mine : other; 
      } 
     barrier(CLK_LOCAL_MEM_FENCE); 
    } 
    if (local_index == 0) { 
     result[get_group_id(0)] = scratch[0]; 
    } 
} 

Répondre

0
accumulator = (accumulator < element) ? accumulator : element; 

Utilisez la fonction fmin - il est exactement ce dont vous avez besoin, et il peut en résulter un code plus rapide (appel intégré dans l'enseignement, le cas échéant, au lieu de ramification coûteux)

global_index += get_global_size(0); 

Quelle est votre typique get_global_size(0)?

Bien que votre modèle d'accès ne soit pas très mauvais (il est coalescé, morceaux de 128 octets pour 32-chaîne) - il est préférable d'accéder à la mémoire de manière séquentielle chaque fois que possible. Par exemple, l'accès séquentiel peut aider memory prefetching (notez, le code OpenCL peut être exécuté sur n'importe quel périphérique, y compris le processeur).

Tenir compte schéma suivant: chaque thread traiterait plage

[ get_global_id(0)*delta , (get_global_id(0)+1)*delta) 

Il se traduira par un accès totalement séquentiel.