Synchronisierte Aggregierung

Innerhalb eines Blocks können Ergebnisse synchronisiert aggregiert werden. Dies ist möglich durch

die Verwendung gemeinsamer Datenstrukturen, die mit __shared__ deklariert werden und
die Verwendung der Synchronisierungsoperation __syncthreads(), die den Aufrufer blockiert, bis alle Threads des gleichen Blocks __syncthreads() aufgerufen haben.

Wenn jeder Thread einen Wert berechnet, wovon die Gesamtsumme gewünscht ist,

kann im ersten Schritt jeder zweite Prozess die Summe des eigenen Resultats mit der seines Nachbarn verknüpfen,
im zweiten Schritt jeder vierter Prozess die im ersten Schritt aggregierte Summen bei sich und dem übernächsten Nachbarn verknüpfen
usw.

bis am Ende nur noch eine Summe pro Block bleibt. Dies funktioniert, wenn die Zahl der Threads eine Zweier-Potenz ist (z.B. 128 oder 256).

Die blockübergreifende Aggregierung kann dann nur über atomare Operationen oder außerhalb des Kernels erfolgen. Letzteres ist möglich entweder auf der CPU oder durch den Aufruf einer weiteren Kernel-Funktion.

Aufgabe

Schreiben Sie eine Kernel-Funktion für das Skalarprodukt, bei dem blockweise aggregiert wird:

template<typename Index, typename TX, typename TY, typename T>
__global__ void dot(Index n,
      const TX* x, Index incX, TY* y, Index incY, T* sums) {
   /* ... */
}

Das Array sums ist entsprechend der Zahl der Blöcke zu dimensionieren, d.h. der \(i\)-te Block legt seine Zwischensumme in sums[i] ab. Berechnen Sie die Gesamtsumme anschließend mit der CPU.