============================ 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: ---- CODE (type=cpp) ---------------------------------------------------------- template __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. :navigate: up -> doc:index back -> doc:session24/page02 next -> doc:session24/page04