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.

        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.