cuda
Communication inter-blocs

Remarques

Garde de dernier bloc

Considérons une grille travaillant sur une tâche, par exemple une réduction parallèle. Au départ, chaque bloc peut faire son travail de manière indépendante, produisant un résultat partiel. À la fin, les résultats partiels doivent être combinés et fusionnés. Un exemple typique est un algorithme de réduction sur un big data.

Une approche typique consiste à invoquer deux noyaux, l’un pour le calcul partiel et l’autre pour la fusion. Cependant, si la fusion peut être effectuée efficacement par un seul bloc, un seul appel au noyau est requis. Ceci est réalisé par un garde de dernier lastBlock défini comme:

__device__ bool lastBlock(int* counter) {
  __threadfence(); //ensure that partial result is visible by all blocks
  int last = 0;
  if (threadIdx.x == 0)
    last = atomicAdd(counter, 1);
  return __syncthreads_or(last == gridDim.x-1);
}

__device__ bool lastBlock(int* counter) {
  __shared__ int last;
  __threadfence(); //ensure that partial result is visible by all blocks
  if (threadIdx.x == 0) {
    last = atomicAdd(counter, 1);
  }
  __syncthreads();
  return last == gridDim.x-1;
}

Avec un tel garde, le dernier bloc est garanti pour voir tous les résultats produits par tous les autres blocs et peut effectuer la fusion.

__device__ void computePartial(T* out) { ... }
__device__ void merge(T* partialResults, T* out) { ... }

__global__ void kernel(int* counter, T* partialResults, T* finalResult) {
    computePartial(&partialResults[blockIdx.x]);
    if (lastBlock(counter)) {
      //this is executed by all threads of the last block only
      merge(partialResults,finalResult);
    }
}

Hypothèses:

• Le compteur doit être un pointeur de mémoire global initialisé à 0 avant que le noyau soit appelé.
• La fonction lastBlock est invoquée uniformément par tous les threads de tous les blocs
• Le noyau est appelé dans une grille de dimension unique (pour simplifier l'exemple)
• T nomme n'importe quel type que vous aimez, mais l'exemple n'est pas destiné à être un modèle au sens C ++

File d'attente de travail globale

Considérez un tableau d'éléments de travail. Le temps nécessaire à l'exécution de chaque tâche varie grandement. Afin d'équilibrer la répartition du travail entre les blocs, il peut être prudent pour chaque bloc de récupérer l'élément suivant uniquement lorsque le précédent est terminé. Ceci contraste avec l'attribution a priori d'éléments aux blocs.

class WorkQueue {
private:
  WorkItem* gItems;
  size_t totalSize;
  size_t current;
public:
  __device__ WorkItem& fetch() {
    __shared__ WorkItem item;
    if (threadIdx.x == 0) {
      size_t itemIdx = atomicAdd(current,1);
      if (itemIdx<totalSize)
        item = gItems[itemIdx];
      else
        item = WorkItem::none();
    }
    __syncthreads();
    return item; //returning reference to smem - ok
  }
}

Hypothèses:

• L'objet WorkQueue, ainsi que le tableau gItem résident dans la mémoire globale
• Aucun nouvel élément de travail n'est ajouté à l'objet WorkQueue dans le noyau qui en extrait
• Le WorkItem est une petite représentation de l'affectation de travail, par exemple un pointeur vers un autre objet.
• WorkItem::none() membre statique WorkItem::none() crée un objet WorkItem qui ne représente aucun travail
• WorkQueue::fetch() doit être appelé uniformément par tous les threads du bloc
• Il n'y a pas 2 invocations de WorkQueue::fetch() sans un autre __syncthreads() entre. Sinon une condition de course apparaîtra!

L'exemple n'inclut pas comment initialiser le WorkQueue ou le remplir. Cela est fait par un autre noyau ou code CPU et devrait être assez simple.