void sparse_fully_connected_1x1_layer_tester_cuda::enqueue_forward_propagation(
cudaStream_t stream_id,
cuda_linear_buffer_device::ptr output_buffer,
const std::vector<cuda_linear_buffer_device::const_ptr>& schema_data,
const std::vector<cuda_linear_buffer_device::const_ptr>& data,
const std::vector<cuda_linear_buffer_device::const_ptr>& data_custom,
const std::vector<cuda_linear_buffer_device::const_ptr>& input_buffers,
const std::vector<cuda_linear_buffer_device::const_ptr>& persistent_working_data,
cuda_linear_buffer_device::ptr temporary_working_fixed_buffer,
cuda_linear_buffer_device::ptr temporary_working_per_entry_buffer,
unsigned int entry_count)
{
{
cusparse_safe_call(cusparseSetStream(cuda_config->get_cusparse_handle(), stream_id));
float alpha = 1.0F;
float beta = 0.0F;
cusparseMatDescr_t mat_descr;
cusparse_safe_call(cusparseCreateMatDescr(&mat_descr));
cusparse_safe_call(cusparseScsrmm(
cuda_config->get_cusparse_handle(),
CUSPARSE_OPERATION_NON_TRANSPOSE,
output_elem_count_per_entry,
entry_count,
input_elem_count_per_entry_list[0],
feature_map_connection_count,
&alpha,
mat_descr,
*data[0],
*data_custom[1],
*data_custom[0],
*input_buffers[0],
input_elem_count_per_entry_list[0],
&beta,
*output_buffer,
output_elem_count_per_entry));
}
// Add bias
{
cudnn_safe_call(cudnnSetStream(cuda_config->get_cudnn_handle(), stream_id));
cudnn_util::set_tensor_descriptor(
output_data_desc,
output_configuration_specific,
entry_count);
float alpha = 1.0F;
float beta = 1.0F;
cudnn_safe_call(cudnnAddTensor(
cuda_config->get_cudnn_handle(),
&alpha,
bias_desc,
*data[1],
&beta,
output_data_desc,
*output_buffer));
}
}
static void apply( const matrix_type & A ,
const multi_vector_type & x ,
const multi_vector_type & y ,
const std::vector<Ordinal> & col_indices )
{
CudaSparseSingleton & s = CudaSparseSingleton::singleton();
const float alpha = 1 , beta = 0 ;
const int n = A.graph.row_map.dimension_0() - 1 ;
const int nz = A.graph.entries.dimension_0();
const size_t ncol = col_indices.size();
// Copy columns of x into a contiguous vector
vector_type xx( Kokkos::allocate_without_initializing, "xx" , n * ncol );
vector_type yy( Kokkos::allocate_without_initializing, "yy" , n * ncol );
for (size_t col=0; col<ncol; col++) {
const std::pair< size_t , size_t > span( n * col , n * ( col + 1 ) );
vector_type xx_view = Kokkos::subview<vector_type>( xx , span );
vector_type x_col =
Kokkos::subview<vector_type>( x, Kokkos::ALL(), col_indices[col] );
Kokkos::deep_copy(xx_view, x_col);
}
// Sparse matrix-times-multivector
cusparseStatus_t status =
cusparseScsrmm( s.handle ,
CUSPARSE_OPERATION_NON_TRANSPOSE ,
n , ncol , n , nz ,
&alpha ,
s.descra ,
A.values.ptr_on_device() ,
A.graph.row_map.ptr_on_device() ,
A.graph.entries.ptr_on_device() ,
xx.ptr_on_device() ,
n ,
&beta ,
yy.ptr_on_device() ,
n );
if ( CUSPARSE_STATUS_SUCCESS != status ) {
throw std::runtime_error( std::string("ERROR - cusparseDcsrmv " ) );
}
// Copy columns out of continguous multivector
for (size_t col=0; col<ncol; col++) {
const std::pair< size_t , size_t > span( n * col , n * ( col + 1 ) );
vector_type yy_view = Kokkos::subview<vector_type>( yy , span );
vector_type y_col =
Kokkos::subview<vector_type>( y, Kokkos::ALL(), col_indices[col] );
Kokkos::deep_copy(y_col, yy_view );
}
}
/* ////////////////////////////////////////////////////////////////////////////
-- testing sparse matrix vector product
*/
int main( int argc, char** argv )
{
magma_int_t info = 0;
TESTING_CHECK( magma_init() );
magma_print_environment();
magma_queue_t queue=NULL;
magma_queue_create( 0, &queue );
magma_s_matrix hA={Magma_CSR}, hA_SELLP={Magma_CSR},
dA={Magma_CSR}, dA_SELLP={Magma_CSR};
magma_s_matrix hx={Magma_CSR}, hy={Magma_CSR}, dx={Magma_CSR},
dy={Magma_CSR}, hrefvec={Magma_CSR}, hcheck={Magma_CSR};
hA_SELLP.blocksize = 8;
hA_SELLP.alignment = 8;
real_Double_t start, end, res;
#ifdef MAGMA_WITH_MKL
magma_int_t *pntre=NULL;
#endif
cusparseHandle_t cusparseHandle = NULL;
cusparseMatDescr_t descr = NULL;
float c_one = MAGMA_S_MAKE(1.0, 0.0);
float c_zero = MAGMA_S_MAKE(0.0, 0.0);
float accuracy = 1e-10;
#define PRECISION_s
#if defined(PRECISION_c)
accuracy = 1e-4;
#endif
#if defined(PRECISION_s)
accuracy = 1e-4;
#endif
magma_int_t i, j;
for( i = 1; i < argc; ++i ) {
if ( strcmp("--blocksize", argv[i]) == 0 ) {
hA_SELLP.blocksize = atoi( argv[++i] );
} else if ( strcmp("--alignment", argv[i]) == 0 ) {
hA_SELLP.alignment = atoi( argv[++i] );
} else
break;
}
printf("\n# usage: ./run_sspmm"
" [ --blocksize %lld --alignment %lld (for SELLP) ] matrices\n\n",
(long long) hA_SELLP.blocksize, (long long) hA_SELLP.alignment );
while( i < argc ) {
if ( strcmp("LAPLACE2D", argv[i]) == 0 && i+1 < argc ) { // Laplace test
i++;
magma_int_t laplace_size = atoi( argv[i] );
TESTING_CHECK( magma_sm_5stencil( laplace_size, &hA, queue ));
} else { // file-matrix test
TESTING_CHECK( magma_s_csr_mtx( &hA, argv[i], queue ));
}
printf("%% matrix info: %lld-by-%lld with %lld nonzeros\n",
(long long) hA.num_rows, (long long) hA.num_cols, (long long) hA.nnz );
real_Double_t FLOPS = 2.0*hA.nnz/1e9;
// m - number of rows for the sparse matrix
// n - number of vectors to be multiplied in the SpMM product
magma_int_t m, n;
m = hA.num_rows;
n = 48;
// init CPU vectors
TESTING_CHECK( magma_svinit( &hx, Magma_CPU, m, n, c_one, queue ));
TESTING_CHECK( magma_svinit( &hy, Magma_CPU, m, n, c_zero, queue ));
// init DEV vectors
TESTING_CHECK( magma_svinit( &dx, Magma_DEV, m, n, c_one, queue ));
TESTING_CHECK( magma_svinit( &dy, Magma_DEV, m, n, c_zero, queue ));
// calling MKL with CSR
#ifdef MAGMA_WITH_MKL
TESTING_CHECK( magma_imalloc_cpu( &pntre, m + 1 ) );
pntre[0] = 0;
for (j=0; j < m; j++ ) {
pntre[j] = hA.row[j+1];
}
MKL_INT num_rows = hA.num_rows;
MKL_INT num_cols = hA.num_cols;
MKL_INT nnz = hA.nnz;
MKL_INT num_vecs = n;
MKL_INT *col;
TESTING_CHECK( magma_malloc_cpu( (void**) &col, nnz * sizeof(MKL_INT) ));
for( magma_int_t t=0; t < hA.nnz; ++t ) {
col[ t ] = hA.col[ t ];
}
MKL_INT *row;
TESTING_CHECK( magma_malloc_cpu( (void**) &row, num_rows * sizeof(MKL_INT) ));
for( magma_int_t t=0; t < hA.num_rows; ++t ) {
row[ t ] = hA.col[ t ];
}
// === Call MKL with consecutive SpMVs, using mkl_scsrmv ===
// warmp up
mkl_scsrmv( "N", &num_rows, &num_cols,
MKL_ADDR(&c_one), "GFNC", MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val),
MKL_ADDR(&c_zero), MKL_ADDR(hy.val) );
start = magma_wtime();
for (j=0; j < 10; j++ ) {
mkl_scsrmv( "N", &num_rows, &num_cols,
MKL_ADDR(&c_one), "GFNC", MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val),
MKL_ADDR(&c_zero), MKL_ADDR(hy.val) );
}
end = magma_wtime();
printf( "\n > MKL SpMVs : %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10/(end-start) );
// === Call MKL with blocked SpMVs, using mkl_scsrmm ===
char transa = 'n';
MKL_INT ldb = n, ldc=n;
char matdescra[6] = {'g', 'l', 'n', 'c', 'x', 'x'};
// warm up
mkl_scsrmm( &transa, &num_rows, &num_vecs, &num_cols, MKL_ADDR(&c_one), matdescra,
MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val), &ldb,
MKL_ADDR(&c_zero),
MKL_ADDR(hy.val), &ldc );
start = magma_wtime();
for (j=0; j < 10; j++ ) {
mkl_scsrmm( &transa, &num_rows, &num_vecs, &num_cols, MKL_ADDR(&c_one), matdescra,
MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val), &ldb,
MKL_ADDR(&c_zero),
MKL_ADDR(hy.val), &ldc );
}
end = magma_wtime();
printf( "\n > MKL SpMM : %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10.*n/(end-start) );
magma_free_cpu( row );
magma_free_cpu( col );
row = NULL;
col = NULL;
#endif // MAGMA_WITH_MKL
// copy matrix to GPU
TESTING_CHECK( magma_smtransfer( hA, &dA, Magma_CPU, Magma_DEV, queue ));
// SpMV on GPU (CSR)
start = magma_sync_wtime( queue );
for (j=0; j < 10; j++) {
TESTING_CHECK( magma_s_spmv( c_one, dA, dx, c_zero, dy, queue ));
}
end = magma_sync_wtime( queue );
printf( " > MAGMA: %.2e seconds %.2e GFLOP/s (standard CSR).\n",
(end-start)/10, FLOPS*10.*n/(end-start) );
TESTING_CHECK( magma_smtransfer( dy, &hrefvec , Magma_DEV, Magma_CPU, queue ));
magma_smfree(&dA, queue );
// convert to SELLP and copy to GPU
TESTING_CHECK( magma_smconvert( hA, &hA_SELLP, Magma_CSR, Magma_SELLP, queue ));
TESTING_CHECK( magma_smtransfer( hA_SELLP, &dA_SELLP, Magma_CPU, Magma_DEV, queue ));
magma_smfree(&hA_SELLP, queue );
magma_smfree( &dy, queue );
TESTING_CHECK( magma_svinit( &dy, Magma_DEV, dx.num_rows, dx.num_cols, c_zero, queue ));
// SpMV on GPU (SELLP)
start = magma_sync_wtime( queue );
for (j=0; j < 10; j++) {
TESTING_CHECK( magma_s_spmv( c_one, dA_SELLP, dx, c_zero, dy, queue ));
}
end = magma_sync_wtime( queue );
printf( " > MAGMA: %.2e seconds %.2e GFLOP/s (SELLP).\n",
(end-start)/10, FLOPS*10.*n/(end-start) );
TESTING_CHECK( magma_smtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue ));
res = 0.0;
for(magma_int_t k=0; k < hA.num_rows; k++ ) {
res=res + MAGMA_S_REAL(hcheck.val[k]) - MAGMA_S_REAL(hrefvec.val[k]);
}
printf("%% |x-y|_F = %8.2e\n", res);
if ( res < accuracy )
printf("%% tester spmm SELL-P: ok\n");
else
printf("%% tester spmm SELL-P: failed\n");
magma_smfree( &hcheck, queue );
magma_smfree(&dA_SELLP, queue );
// SpMV on GPU (CUSPARSE - CSR)
// CUSPARSE context //
magma_smfree( &dy, queue );
TESTING_CHECK( magma_svinit( &dy, Magma_DEV, dx.num_rows, dx.num_cols, c_zero, queue ));
//#ifdef PRECISION_d
start = magma_sync_wtime( queue );
TESTING_CHECK( cusparseCreate( &cusparseHandle ));
TESTING_CHECK( cusparseSetStream( cusparseHandle, magma_queue_get_cuda_stream(queue) ));
TESTING_CHECK( cusparseCreateMatDescr( &descr ));
TESTING_CHECK( cusparseSetMatType( descr, CUSPARSE_MATRIX_TYPE_GENERAL ));
TESTING_CHECK( cusparseSetMatIndexBase( descr, CUSPARSE_INDEX_BASE_ZERO ));
float alpha = c_one;
float beta = c_zero;
// copy matrix to GPU
TESTING_CHECK( magma_smtransfer( hA, &dA, Magma_CPU, Magma_DEV, queue) );
for (j=0; j < 10; j++) {
cusparseScsrmm(cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE,
dA.num_rows, n, dA.num_cols, dA.nnz,
&alpha, descr, dA.dval, dA.drow, dA.dcol,
dx.dval, dA.num_cols, &beta, dy.dval, dA.num_cols);
}
end = magma_sync_wtime( queue );
printf( " > CUSPARSE: %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10*n/(end-start) );
TESTING_CHECK( magma_smtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue ));
res = 0.0;
for(magma_int_t k=0; k < hA.num_rows; k++ ) {
res = res + MAGMA_S_REAL(hcheck.val[k]) - MAGMA_S_REAL(hrefvec.val[k]);
}
printf("%% |x-y|_F = %8.2e\n", res);
if ( res < accuracy )
printf("%% tester spmm cuSPARSE: ok\n");
else
printf("%% tester spmm cuSPARSE: failed\n");
magma_smfree( &hcheck, queue );
cusparseDestroyMatDescr( descr );
cusparseDestroy( cusparseHandle );
descr = NULL;
cusparseHandle = NULL;
//#endif
printf("\n\n");
// free CPU memory
magma_smfree( &hA, queue );
magma_smfree( &hx, queue );
magma_smfree( &hy, queue );
magma_smfree( &hrefvec, queue );
// free GPU memory
magma_smfree( &dx, queue );
magma_smfree( &dy, queue );
magma_smfree( &dA, queue);
#ifdef MAGMA_WITH_MKL
magma_free_cpu( pntre );
#endif
i++;
}
magma_queue_destroy( queue );
TESTING_CHECK( magma_finalize() );
return info;
}
extern "C" magma_int_t
magma_s_spmv(
float alpha,
magma_s_sparse_matrix A,
magma_s_vector x,
float beta,
magma_s_vector y,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
if ( A.memory_location != x.memory_location ||
x.memory_location != y.memory_location ) {
printf("error: linear algebra objects are not located in same memory!\n");
printf("memory locations are: %d %d %d\n",
A.memory_location, x.memory_location, y.memory_location );
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_INVALID_PTR;
}
// DEV case
if ( A.memory_location == Magma_DEV ) {
if ( A.num_cols == x.num_rows && x.num_cols == 1 ) {
if ( A.storage_type == Magma_CSR
|| A.storage_type == Magma_CSRL
|| A.storage_type == Magma_CSRU ) {
//printf("using CSR kernel for SpMV: ");
//magma_sgecsrmv( MagmaNoTrans, A.num_rows, A.num_cols, alpha,
// A.dval, A.drow, A.dcol, x.dval, beta, y.dval );
//printf("done.\n");
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
cusparseSetStream( cusparseHandle, queue );
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO);
cusparseScsrmv( cusparseHandle,CUSPARSE_OPERATION_NON_TRANSPOSE,
A.num_rows, A.num_cols, A.nnz, &alpha, descr,
A.dval, A.drow, A.dcol, x.dval, &beta, y.dval );
cusparseDestroyMatDescr( descr );
cusparseDestroy( cusparseHandle );
}
else if ( A.storage_type == Magma_ELL ) {
//printf("using ELLPACKT kernel for SpMV: ");
magma_sgeelltmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.dval, A.dcol, x.dval, beta,
y.dval, queue );
//printf("done.\n");
}
else if ( A.storage_type == Magma_ELLPACKT ) {
//printf("using ELL kernel for SpMV: ");
magma_sgeellmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.dval, A.dcol, x.dval, beta,
y.dval, queue );
//printf("done.\n");
}
else if ( A.storage_type == Magma_ELLRT ) {
//printf("using ELLRT kernel for SpMV: ");
magma_sgeellrtmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.dval, A.dcol, A.drow, x.dval,
beta, y.dval, A.alignment, A.blocksize, queue );
//printf("done.\n");
}
else if ( A.storage_type == Magma_SELLP ) {
//printf("using SELLP kernel for SpMV: ");
magma_sgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.blocksize, A.numblocks, A.alignment,
alpha, A.dval, A.dcol, A.drow, x.dval, beta, y.dval, queue );
//printf("done.\n");
}
else if ( A.storage_type == Magma_DENSE ) {
//printf("using DENSE kernel for SpMV: ");
magmablas_sgemv( MagmaNoTrans, A.num_rows, A.num_cols, alpha,
A.dval, A.num_rows, x.dval, 1, beta, y.dval,
1 );
//printf("done.\n");
}
/* else if ( A.storage_type == Magma_BCSR ) {
//printf("using CUSPARSE BCSR kernel for SpMV: ");
// CUSPARSE context //
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
cusparseSetStream( cusparseHandle, queue );
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
// end CUSPARSE context //
cusparseDirection_t dirA = CUSPARSE_DIRECTION_ROW;
int mb = (A.num_rows + A.blocksize-1)/A.blocksize;
int nb = (A.num_cols + A.blocksize-1)/A.blocksize;
cusparseSbsrmv( cusparseHandle, dirA,
CUSPARSE_OPERATION_NON_TRANSPOSE, mb, nb, A.numblocks,
&alpha, descr, A.dval, A.drow, A.dcol, A.blocksize, x.dval,
&beta, y.dval );
//printf("done.\n");
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}*/
else {
printf("error: format not supported.\n");
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_NOT_SUPPORTED;
}
}
else if ( A.num_cols < x.num_rows || x.num_cols > 1 ) {
magma_int_t num_vecs = x.num_rows / A.num_cols * x.num_cols;
if ( A.storage_type == Magma_CSR ) {
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
cusparseSetStream( cusparseHandle, queue );
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO);
if ( x.major == MagmaColMajor) {
cusparseScsrmm(cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE,
A.num_rows, num_vecs, A.num_cols, A.nnz,
&alpha, descr, A.dval, A.drow, A.dcol,
x.dval, A.num_cols, &beta, y.dval, A.num_cols);
} else if ( x.major == MagmaRowMajor) {
cusparseScsrmm2(cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE,
CUSPARSE_OPERATION_TRANSPOSE,
A.num_rows, num_vecs, A.num_cols, A.nnz,
&alpha, descr, A.dval, A.drow, A.dcol,
x.dval, A.num_cols, &beta, y.dval, A.num_cols);
}
cusparseDestroyMatDescr( descr );
cusparseDestroy( cusparseHandle );
}
else if ( A.storage_type == Magma_ELL ) {
if ( x.major == MagmaColMajor) {
magma_smgeelltmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.max_nnz_row, alpha, A.dval, A.dcol, x.dval,
beta, y.dval, queue );
}
else if ( x.major == MagmaRowMajor) {
// transpose first to col major
magma_s_vector x2;
magma_svtranspose( x, &x2, queue );
magma_smgeellmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.max_nnz_row, alpha, A.dval, A.dcol, x.dval,
beta, y.dval, queue );
magma_s_vfree(&x2, queue );
}
}
else if ( A.storage_type == Magma_ELLPACKT ) {
if ( x.major == MagmaColMajor) {
magma_smgeellmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.max_nnz_row, alpha, A.dval, A.dcol, x.dval,
beta, y.dval, queue );
}
else if ( x.major == MagmaRowMajor) {
// transpose first to col major
magma_s_vector x2;
magma_svtranspose( x, &x2, queue );
magma_smgeelltmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.max_nnz_row, alpha, A.dval, A.dcol, x.dval,
beta, y.dval, queue );
magma_s_vfree(&x2, queue );
}
} else if ( A.storage_type == Magma_SELLP ) {
if ( x.major == MagmaRowMajor) {
magma_smgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.blocksize, A.numblocks, A.alignment,
alpha, A.dval, A.dcol, A.drow, x.dval, beta, y.dval, queue );
}
else if ( x.major == MagmaColMajor) {
// transpose first to row major
magma_s_vector x2;
magma_svtranspose( x, &x2, queue );
magma_smgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.blocksize, A.numblocks, A.alignment,
alpha, A.dval, A.dcol, A.drow, x2.dval, beta, y.dval, queue );
magma_s_vfree(&x2, queue );
}
}/*
if ( A.storage_type == Magma_DENSE ) {
//printf("using DENSE kernel for SpMV: ");
magmablas_smgemv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, alpha, A.dval, A.num_rows, x.dval, 1,
beta, y.dval, 1 );
//printf("done.\n");
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}*/
else {
printf("error: format not supported.\n");
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_NOT_SUPPORTED;
}
}
}
// CPU case missing!
else {
printf("error: CPU not yet supported.\n");
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_NOT_SUPPORTED;
}
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
