cuda
Communicatie tussen blokken

Opmerkingen

Laatste blok bewaker

Overweeg een raster dat aan een bepaalde taak werkt, bijvoorbeeld een parallelle reductie. Aanvankelijk kan elk blok zijn werk onafhankelijk doen, wat een gedeeltelijk resultaat oplevert. Op het einde moeten de gedeeltelijke resultaten echter worden gecombineerd en samengevoegd. Een typisch voorbeeld is een reductie-algoritme voor big data.

Een typische benadering is om twee kernels op te roepen, één voor de gedeeltelijke berekening en de andere voor samenvoegen. Als het samenvoegen echter efficiënt kan worden uitgevoerd door een enkel blok, is slechts één kernelaanroep vereist. Dit wordt bereikt door een lastBlock guard gedefinieerd als:

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

Met zo'n bewaker ziet het laatste blok gegarandeerd alle resultaten die door alle andere blokken zijn geproduceerd en kan het samenvoegen worden uitgevoerd.

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

Veronderstellingen:

• De teller moet een globale geheugenpointer zijn, geïnitialiseerd op 0 voordat de kernel wordt aangeroepen.
• De lastBlock functie wordt uniform opgeroepen door alle threads in alle blokken
• De kernel wordt aangeroepen in een-dimensionaal raster (voor de eenvoud van het voorbeeld)
• T benoemt elk type dat u maar wilt, maar het voorbeeld is niet bedoeld als een sjabloon in C ++ - zin

Wereldwijde werkwachtrij

Overweeg een reeks werkitems. De benodigde tijd voor het voltooien van elk werkitem varieert enorm. Om de werkverdeling tussen blokken in evenwicht te brengen, kan het verstandig zijn voor elk blok het volgende item alleen op te halen als het vorige is voltooid. Dit in tegenstelling tot het a priori toewijzen van items aan blokken.

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

Veronderstellingen:

• Het WorkQueue-object en de gItem-array bevinden zich in het wereldwijde geheugen
• Er worden geen nieuwe werkitems toegevoegd aan het WorkQueue-object in de kernel die het ophaalt
• Het WorkItem is een kleine weergave van de werkopdracht, bijvoorbeeld een aanwijzer naar een ander object
• WorkItem::none() statische WorkItem maakt een WorkItem object dat helemaal geen werk vertegenwoordigt
• WorkQueue::fetch() moet uniform worden aangeroepen door alle threads in het blok
• Er zijn geen 2 aanroepen van WorkQueue::fetch() zonder een andere __syncthreads() ertussen. Anders verschijnt er een raceconditie!

Het voorbeeld bevat niet hoe de WorkQueue geïnitialiseerd of ingevuld. Het wordt gedaan door een andere kernel of CPU-code en zou vrij eenvoudig moeten zijn.