=======================================
Passing matrices and vectors to the GPU                    [TOC]
=======================================

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:

---- CODE (type=cpp) ----------------------------------------------------------
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
-------------------------------------------------------------------------------

:navigate: up    -> doc:index
           back  -> doc:session09/page01
           next  -> doc:session09/page03