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:

CHECK_CUDA(cudaMalloc, (void**)&data, length * sizeof(T));

}

~DenseVector() {

CHECK_CUDA(cudaFree, data);

}

----------------------------------------------------------------------------

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:

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?