cuda
Komunikacja między blokami

Uwagi

Strażnik ostatniego bloku

Zastanów się nad siatką działającą nad jakimś zadaniem, np. Redukcją równoległą. Początkowo każdy blok może wykonywać swoją pracę niezależnie, powodując częściowy wynik. Na koniec jednak częściowe wyniki należy połączyć i połączyć. Typowym przykładem jest algorytm redukcji dużych zbiorów danych.

Typowe podejście polega na wywołaniu dwóch jąder, jednego do częściowego obliczenia, a drugiego do scalenia. Jeśli jednak scalanie można wykonać skutecznie za pomocą jednego bloku, wymagane jest tylko jedno wywołanie jądra. Osiąga się to poprzez zabezpieczenie lastBlock zdefiniowane jako:

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

Przy takiej osłonie gwarantuje się, że ostatni blok zobaczy wszystkie wyniki wygenerowane przez wszystkie inne bloki i będzie mógł wykonać scalenie.

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

Założenia:

• Licznik musi być globalnym wskaźnikiem pamięci, zainicjowanym na 0 przed wywołaniem jądra.
• Funkcja lastBlock jest wywoływana jednakowo przez wszystkie wątki we wszystkich blokach
• Jądro jest wywoływane w siatce jednowymiarowej (dla uproszczenia przykładu)
• T nazywa dowolny typ, który lubisz, ale ten przykład nie ma być szablonem w sensie C ++

Globalna kolejka pracy

Rozważ szereg elementów pracy. Czas potrzebny do ukończenia każdego elementu pracy jest bardzo różny. Aby zrównoważyć rozkład pracy między blokami, rozsądne może być, aby każdy blok pobierał następny element tylko po zakończeniu poprzedniego. Jest to sprzeczne z przypisywaniem elementów do bloków z góry.

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

Założenia:

• Obiekt WorkQueue, a także tablica gItem znajdują się w pamięci globalnej
• Żadne nowe elementy pracy nie są dodawane do obiektu WorkQueue w jądrze, które się z niego pobiera
• WorkItem to mała reprezentacja przydziału pracy, np. Wskaźnik do innego obiektu
• Statyczna funkcja WorkItem::none() tworzy obiekt WorkItem który w ogóle nie reprezentuje pracy
• WorkQueue::fetch() muszą być wywoływane jednakowo przez wszystkie wątki w bloku
• Nie ma 2 wywołań WorkQueue::fetch() bez kolejnego __syncthreads() pomiędzy nimi. W przeciwnym razie pojawią się warunki wyścigu!

Przykład nie obejmuje sposobu inicjowania lub zapełniania WorkQueue . Odbywa się to za pomocą innego jądra lub kodu procesora i powinno być dość proste.