cuda
Comunicazione tra blocchi

Osservazioni

Last block guard

Considera una griglia che lavori su qualche compito, ad esempio una riduzione parallela. Inizialmente, ogni blocco può fare il suo lavoro in modo indipendente, producendo alcuni risultati parziali. Alla fine, tuttavia, i risultati parziali devono essere combinati e riuniti. Un tipico esempio è un algoritmo di riduzione su un big data.

Un approccio tipico consiste nel richiamare due kernel, uno per il calcolo parziale e l'altro per l'unione. Tuttavia, se la fusione può essere eseguita in modo efficiente da un singolo blocco, è necessaria solo una chiamata del kernel. Ciò è ottenuto da una guardia lastBlock definita come:

__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;
}

Con una tale guardia l'ultimo blocco è garantito per vedere tutti i risultati prodotti da tutti gli altri blocchi e può eseguire la fusione.

__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);
    }
}

ipotesi:

• Il contatore deve essere un puntatore di memoria globale, inizializzato a 0 prima che il kernel venga richiamato.
• La funzione lastBlock è invocata in modo uniforme da tutti i thread in tutti i blocchi
• Il kernel è invocato in una griglia monodimensionale (per semplicità dell'esempio)
• T identifica qualsiasi tipo che ti piace, ma l'esempio non intende essere un modello in senso C ++

Coda di lavoro globale

Considera una serie di oggetti di lavoro. Il tempo necessario per il completamento di ciascun elemento di lavoro varia notevolmente. Per bilanciare la distribuzione del lavoro tra i blocchi, può essere prudente per ogni blocco recuperare l'elemento successivo solo quando il precedente è completo. Ciò è in contrasto con l'attribuzione a priori di elementi a blocchi.

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
  }
}

ipotesi:

• L'oggetto WorkQueue e l'array gItem risiedono nella memoria globale
• Nessun nuovo oggetto di lavoro viene aggiunto all'oggetto WorkQueue nel kernel che sta recuperando da esso
• WorkItem è una piccola rappresentazione dell'assegnazione del lavoro, ad esempio un puntatore a un altro oggetto
• WorkItem::none() funzione membro statico WorkItem::none() crea un oggetto WorkItem che non rappresenta affatto lavoro
• WorkQueue::fetch() deve essere chiamato uniformemente da tutti i thread nel blocco
• Non ci sono 2 invocazioni di WorkQueue::fetch() senza un altro __syncthreads() in mezzo. Altrimenti apparirà una condizione di gara!

L'esempio non include come inizializzare WorkQueue o popolarlo. È fatto da un altro kernel o codice CPU e dovrebbe essere abbastanza semplice.