=============================================
Aufruf mehrerer Kernel-Funktionen auf der GPU
=============================================

Kernel-Funktionen können auch hintereinander aufgerufen werden
und entsprechend mit den bereits in der GPU liegenden Resultaten
weiterarbeiten. Dann ist es nicht mehr notwendig, die Daten bei
jedem Aufruf hin- und herzukopieren.

Aufrufe von Kernel-Funktionen werden innerhalb eines sogenannten
Streams serialisiert. Wenn bei der Konfiguration eines Kernels
kein Stream angegeben ist, wird ein voreingestellter Stream
ausgewählt. Entsprechend werden alle Kernel-Funktionsaufrufe
ohne Stream-Angabe sequentiell abgearbeitet.

Anmerkung: Es ist dabei festzuhalten, dass im Erfolgsfalle der Aufruf
einer Kernel-Funktion sofort zurückkehrt und nicht auf die
Beendigung der Kernel-Funktion wartet. Kopier-Funktionen wie
`cudaMemcpy` werden ebenfalls einem Stream zugeordnet und
entsprechend synchronisiert. Das führte dazu, dass wir mit
dem Abholen der Ergebnisse implizit auf die Abarbeitung der
zuvor aufgerufenen Kernel-Funktion gewartet haben.

Aufgabe
=======

* Schreiben Sie eine Kernel-Funktion `asum`, der blockweise
  die Teilsummen des Vektors $\vec{x}$ bestimmt und diese
  in `sums` ablegt.

   ---- CODE (type=cpp) -------------------------------------------------------
   template<typename Index, typename T>
   __global__ void asum(Index n, const T* x, T* sums) {
      /* ... */
   }
   ----------------------------------------------------------------------------

* Passen Sie die vorherige Lösung so an, dass die Gesamtsumme
  nicht mit der CPU berechnet wird, sondern mit der GPU,
  indem wiederholt `asum` aufgerufen wird, bis am Ende nur
  noch ein Block zum Einsatz kommt, der die Gesamtsumme
  berechnet. Dann ist nur noch der eine `double`-Wert an
  die CPU zurückzutransferieren.

:navigate: up    -> doc:index
           back  -> doc:session24/page04
           next  -> doc:session24/page06