cuda
Inter-block kommunikation

Anmärkningar

Sista-block vakt

Tänk på ett rutnät som arbetar med någon uppgift, t.ex. en parallell reduktion. Till att börja med kan varje block göra sitt arbete oberoende, vilket ger viss delresultat. I slutet måste emellertid delresultaten kombineras och slås samman. Ett typiskt exempel är en reduktionsalgoritm för stor data.

En typisk metod är att åberopa två kärnor, en för partiell beräkning och en för sammanslagning. Men om sammanslagningen kan göras effektivt med ett enda block krävs endast ett kärnsamtal. Detta uppnås med en lastBlock skydd definierad som:

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

Med en sådan skydd är det sista blocket garanterat att se alla resultat som produceras av alla andra block och kan utföra sammanslagningen.

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

antaganden:

• Räknaren måste vara en global minnespekare, initialiserad till 0 innan kärnan startas.
• lastBlock funktionen åberopas enhetligt av alla trådar i alla block
• Kärnan åberopas i en-dimensionell rutnät (för enkelhetens exempel)
• T namnger vilken typ du vill, men exemplet är inte avsett att vara en mall i C ++ -sinne

Global arbetskö

Tänk på en mängd arbetsobjekt. Den tid som krävs för varje arbetsobjekt att slutföra varierar mycket. För att balansera arbetsfördelningen mellan block kan det vara klokt för varje block att hämta nästa objekt endast när föregående är slutfört. Detta i motsats till att a-priori tilldelar objekt till block.

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

antaganden:

• WorkQueue-objektet, liksom gItem-arrayen finns i det globala minnet
• Inga nya arbetsobjekt läggs till i WorkQueue-objektet i kärnan som hämtar från det
• WorkItem är en liten representation av arbetsuppgiften, t.ex. en pekare till ett annat objekt
• WorkItem::none() statisk medlemsfunktion skapar ett WorkItem objekt som inte representerar något arbete alls
• WorkQueue::fetch() måste kallas enhetligt av alla trådar i blocket
• Det finns inga två WorkQueue::fetch() av WorkQueue::fetch() utan ytterligare __syncthreads() däremellan. Annars kommer ett tävlingsvillkor att dyka upp!

Exemplet inkluderar inte hur initialiserar WorkQueue eller fyller det. Det görs av en annan kärna eller CPU-kod och bör vara ganska rak.