================================================
Datentypen für Matrizen und Vektoren auf der GPU
================================================

Entsprechend dem RAII-Prinzip ist es sehr viel handlicher, das Belegen
und die Freigabe von GPU-Speicher einem Objekt zu überlassen.
Grundsätzlich kann die gleiche Klasse auf beiden Seiten verwendet
werden. Dazu muss für jede Methode einschließlich der Konstruktoren
und Destruktoren jeweils festgelegt werden, auf welcher Seite sie
aufgerufen werden darf. Dies wird von CUDA unterstützt. Per
Voreinstellung sind Methoden nur auf der Host-Seite aufrufbar.
Wenn `_``_device_``_` vor einer Methode steht, dann ist sie
nur von der GPU-Seite aus aufzurufen. Wenn eine Methode von
beiden Seiten aus benutzbar sein soll, dann können die
Schlüsselwörter `_``_device_``_` und `_``_host_``_`
kombiniert werden.

Grundsätzlich liegt dabei folgende Vorgehensweise nahe:

 * Konstruktoren und Destruktoren bleiben der
   CPU-Seite vorbehalten, da nur auf dieser dynamisch
   Speicher auf der GPU-Seite belegt und freigegeben werden kann.
   Beispiel:

   ---- CODE (type=cpp) -------------------------------------------------------
    DenseVector(Index length)
        : length(length), inc(1)
    {
	CHECK_CUDA(cudaMalloc, (void**)&data, length * sizeof(T));
    }

    ~DenseVector()
    {
	CHECK_CUDA(cudaFree, data);
    }
   ----------------------------------------------------------------------------

 * View-Objekte können prinzipiell auf beiden Seiten erzeugt werden,
   da sie keinen Speicher dynamisch belegen oder freigeben. Beispiel:

   ---- CODE (type=cpp) -------------------------------------------------------
    __device__ __host__
    DenseVectorView(Index length, ElementType *data, Index inc)
        : length(length), inc(inc), data(data)
    {
    }
   ----------------------------------------------------------------------------

 * Die Zugriffsmethoden auf GPU-Speicher sind der GPU-Seite
   vorbehalten. Beispiel:

   ---- CODE (type=cpp) -------------------------------------------------------
    __device__
    const ElementType &
    operator()(Index i) const
    {
        assert(i<length);
        return data[i];
    }
   ----------------------------------------------------------------------------

Um den Transfer zwischen GPU- und CPU-Objekten lesbarer zu gestalten,
empfiehlt sich die Definition entsprechender `copy`-Operationen. Hier ist
es aber sinnvoll, nur Datentypen zu unterstützen, die jeweils aus einem
einzigen zusammenhängenden Speicherbereich bestehen, der im Falle von
Matrizen auch gleich organisiert ist. Bei dem Transfer gibt es keine
effiziente Möglichkeit der Umorganisation. Stattdessen sollten die
Kopieraktionen mit einem Aufruf von `cudaMemcpy` erledigbar sein:

---- CODE (type=cpp) ----------------------------------------------------------
template<typename T, typename IndexX, typename IndexY>
void
copy(const hpc::matvec::DenseVector<T, IndexX>& x, DenseVector<T, IndexY>& y) {
   assert(x.length == y.length);
   CHECK_CUDA(cudaMemcpy, y.data, x.data, x.length * sizeof(T),
      cudaMemcpyHostToDevice);
}

template<typename T, typename IndexX, typename IndexY>
void
copy(const DenseVector<T, IndexX>& x, hpc::matvec::DenseVector<T, IndexY>& y) {
   assert(x.length == y.length);
   CHECK_CUDA(cudaMemcpy, y.data, x.data, x.length * sizeof(T),
      cudaMemcpyDeviceToHost);
}
-------------------------------------------------------------------------------

Entsprechend wird hier auf die Unterstützung von Views bewusst verzichtet.

Aufgabe
=======

Passen Sie das Beispiel für das Skalarprodukt so an, dass die
Vektorklasse aus `<hpc/cuda/densevector.h>` und die Kopier-Operationen
aus `<hpc/cuda/copy.h>` verwendet werden. Das Beispiel sollte dann
nirgends mehr `cudaMalloc`, `cudaMemcpy` oder `cudaFree` explizit
aufrufen.

Die Kernel-Funktion können Sie dabei gleich belassen. Überlegen Sie
sich aber gut, wie Sie die Adressen übergeben können. Bislang erfolgte
dies mit Hilfe der Zugriffsfunktion, etwa mit `&x(0)`. Warum ist das
hier nicht akzeptabel?

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