cuda
Kommunikation zwischen Blöcken

Bemerkungen

Last-Block-Wächter

Stellen Sie sich ein Netz vor, das an einer Aufgabe arbeitet, z. B. eine parallele Reduzierung. Anfangs kann jeder Block seine Arbeit unabhängig machen, was zu einem Teilergebnis führt. Am Ende müssen jedoch die Teilergebnisse kombiniert und zusammengeführt werden. Ein typisches Beispiel ist ein Reduktionsalgorithmus für Big Data.

Ein typischer Ansatz besteht darin, zwei Kernel aufzurufen, einen für die Teilberechnung und den anderen für das Zusammenführen. Wenn das Zusammenführen jedoch durch einen einzelnen Block effizient durchgeführt werden kann, ist nur ein Kernelaufruf erforderlich. Dies wird durch einen lastBlock Guard erreicht, der wie lastBlock definiert ist:

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

Mit einem solchen Schutz werden garantiert, dass der letzte Block alle Ergebnisse aller anderen Blöcke anzeigt und das Zusammenführen ausführen kann.

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

Annahmen:

• Der Zähler muss ein globaler Speicherzeiger sein, der vor dem Aufrufen des Kernels auf 0 initialisiert wird.
• Die lastBlock Funktion wird von allen Threads in allen Blöcken einheitlich aufgerufen
• Der Kernel wird in einem eindimensionalen Raster aufgerufen (zur Vereinfachung des Beispiels)
• T benennt jeden beliebigen Typ, aber das Beispiel soll keine Vorlage im Sinne von C ++ sein

Globale Arbeitswarteschlange

Betrachten Sie eine Reihe von Arbeitselementen. Die Zeit, die ein Workitem zum Abschluss benötigt, ist sehr unterschiedlich. Um die Arbeitsverteilung zwischen den Blöcken auszugleichen, kann es sinnvoll sein, für jeden Block den nächsten Artikel erst dann zu holen, wenn der vorherige abgeschlossen ist. Dies steht im Gegensatz zu dem Zuordnen von Elementen zu Blöcken von vornherein.

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

Annahmen:

• Das WorkQueue-Objekt sowie das gItem-Array befinden sich im globalen Speicher
• Dem WorkQueue-Objekt im Kernel, der es abruft, werden keine neuen Arbeitselemente hinzugefügt
• Das WorkItem ist eine kleine Darstellung der Arbeitszuordnung, z. B. ein Zeiger auf ein anderes Objekt
• WorkItem::none() statische Member-Funktion WorkItem::none() erstellt ein WorkItem Objekt, das überhaupt keine Arbeit darstellt
• WorkQueue::fetch() muss von allen Threads im Block einheitlich aufgerufen werden
• Es gibt keine zwei WorkQueue::fetch() von WorkQueue::fetch() ohne einen anderen __syncthreads() dazwischen. Andernfalls wird eine Rennbedingung angezeigt!

Das Beispiel beinhaltet nicht, wie WorkQueue initialisiert oder WorkQueue wird. Es wird von einem anderen Kernel- oder CPU-Code ausgeführt und sollte ziemlich geradlinig sein.