Passing matrices and vectors to the GPU

Content

RAII storage objects for the GPU
Passing matrices and vectors to the GPU
Simple Jacobi solver for the GPU
Sample solution
Working with a 2-dimensional grid of 2-dimensional blocks
Sample solution

We cannot pass anything by reference to a kernel function as it is not helpful to pass the address of an object in host memory to the device which cannot be accessed. Hence, all parameters must be passed by value. There is no exception to this.

But we cannot pass a DeviceGeMatrix object by value as proper RAII objects either inhibit copying (this is the approach of GeMatrix) or interpret it as cloning (which would not help us).

The solution is to pass a non-RAII object as value. Do we have such objects? Yes, indeed all our view objects like GeMatrixView are non-RAII objects. Similarly, we have DeviceGeMatrixView. Such views can be easily passed by value and the viewer classes DeviceGeMatrixView and DeviceGeMatrixConstView support copying:

template<
   template<typename> class Matrix,
   typename T,
   Require< DeviceGe<Matrix<T>>, DeviceView<Matrix<T>> > = true
>
__global__ void f(Matrix<T> A) {
   // ...
}

/* ... */
   GeMatrix<double> A(M, N, Order::RowMajor);
   // fill A...
   DeviceGeMatrix<double> devA(A.numRows(), A.numCols(), Order::RowMajor);
   copy(A, devA); // copy A to devA
   f<<< /* kernel config */ >>>(devA.view());
   copy(devA, A); // copy devA to A