extern "C" magma_int_t
magma_smfree(
magma_s_matrix *A,
magma_queue_t queue )
{
if ( A->memory_location == Magma_CPU ) {
if ( A->storage_type == Magma_ELL || A->storage_type == Magma_ELLPACKT ){
magma_free_cpu( A->val );
magma_free_cpu( A->col );
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if (A->storage_type == Magma_ELLD ) {
magma_free_cpu( A->val );
magma_free_cpu( A->col );
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_ELLRT ) {
magma_free_cpu( A->val );
magma_free_cpu( A->row );
magma_free_cpu( A->col );
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_SELLP ) {
magma_free_cpu( A->val );
magma_free_cpu( A->row );
magma_free_cpu( A->col );
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_CSRLIST ) {
magma_free_cpu( A->val );
magma_free_cpu( A->row );
magma_free_cpu( A->col );
magma_free_cpu( A->list );
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_CSR || A->storage_type == Magma_CSC
|| A->storage_type == Magma_CSRD
|| A->storage_type == Magma_CSRL
|| A->storage_type == Magma_CSRU ) {
magma_free_cpu( A->val );
magma_free_cpu( A->col );
magma_free_cpu( A->row );
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_CSRCOO ) {
magma_free_cpu( A->val );
magma_free_cpu( A->col );
magma_free_cpu( A->row );
magma_free_cpu( A->rowidx );
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_BCSR ) {
magma_free_cpu( A->val );
magma_free_cpu( A->col );
magma_free_cpu( A->row );
magma_free_cpu( A->blockinfo );
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
A->blockinfo = 0;
}
if ( A->storage_type == Magma_DENSE ) {
magma_free_cpu( A->val );
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
A->val = NULL;
A->col = NULL;
A->row = NULL;
A->rowidx = NULL;
A->blockinfo = NULL;
A->diag = NULL;
A->dval = NULL;
A->dcol = NULL;
A->drow = NULL;
A->drowidx = NULL;
A->ddiag = NULL;
A->dlist = NULL;
A->list = NULL;
}
if ( A->memory_location == Magma_DEV ) {
if ( A->storage_type == Magma_ELL || A->storage_type == Magma_ELLPACKT ){
if ( magma_free( A->dval ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->dcol ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_ELLD ) {
if ( magma_free( A->dval ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->dcol ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_ELLRT ) {
if ( magma_free( A->dval ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->drow ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->dcol ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_SELLP ) {
if ( magma_free( A->dval ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->drow ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->dcol ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_CSRLIST ) {
if ( magma_free( A->dval ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->drow ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->dcol ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->dlist ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_CSR || A->storage_type == Magma_CSC
|| A->storage_type == Magma_CSRD
|| A->storage_type == Magma_CSRL
|| A->storage_type == Magma_CSRU ) {
if ( magma_free( A->dval ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->drow ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->dcol ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_CSRCOO ) {
if ( magma_free( A->dval ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->drow ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->dcol ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->drowidx ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_BCSR ) {
if ( magma_free( A->dval ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->drow ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
if ( magma_free( A->dcol ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
magma_free_cpu( A->blockinfo );
A->blockinfo = NULL;
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
if ( A->storage_type == Magma_DENSE ) {
if ( magma_free( A->dval ) != MAGMA_SUCCESS ) {
printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
A->num_rows = 0;
A->num_cols = 0;
A->nnz = 0; A->true_nnz = 0;
}
A->val = NULL;
A->col = NULL;
A->row = NULL;
A->rowidx = NULL;
A->blockinfo = NULL;
A->diag = NULL;
A->dval = NULL;
A->dcol = NULL;
A->drow = NULL;
A->drowidx = NULL;
A->ddiag = NULL;
A->dlist = NULL;
A->list = NULL;
}
else {
// printf("Memory Free Error.\n");
return MAGMA_ERR_INVALID_PTR;
}
return MAGMA_SUCCESS;
}
/* ////////////////////////////////////////////////////////////////////////////
-- testing sparse matrix vector product
*/
int main( int argc, char** argv )
{
magma_int_t info = 0;
TESTING_INIT();
magma_queue_t queue=NULL;
magma_queue_create( 0, &queue );
magma_c_matrix hA={Magma_CSR}, hA_SELLP={Magma_CSR}, hA_ELL={Magma_CSR},
dA={Magma_CSR}, dA_SELLP={Magma_CSR}, dA_ELL={Magma_CSR};
magma_c_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;
magmaFloatComplex c_one = MAGMA_C_MAKE(1.0, 0.0);
magmaFloatComplex c_zero = MAGMA_C_MAKE(0.0, 0.0);
float accuracy = 1e-10;
#define PRECISION_c
#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_cspmm"
" [ --blocksize %d --alignment %d (for SELLP) ]"
" matrices \n\n", int(hA_SELLP.blocksize), int(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] );
CHECK( magma_cm_5stencil( laplace_size, &hA, queue ));
} else { // file-matrix test
CHECK( magma_c_csr_mtx( &hA, argv[i], queue ));
}
printf("%% matrix info: %d-by-%d with %d nonzeros\n",
int(hA.num_rows), int(hA.num_cols), int(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
CHECK( magma_cvinit( &hx, Magma_CPU, m, n, c_one, queue ));
CHECK( magma_cvinit( &hy, Magma_CPU, m, n, c_zero, queue ));
// init DEV vectors
CHECK( magma_cvinit( &dx, Magma_DEV, m, n, c_one, queue ));
CHECK( magma_cvinit( &dy, Magma_DEV, m, n, c_zero, queue ));
// calling MKL with CSR
#ifdef MAGMA_WITH_MKL
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_MALLOC_CPU( col, MKL_INT, nnz );
for( magma_int_t t=0; t < hA.nnz; ++t ) {
col[ t ] = hA.col[ t ];
}
MKL_INT *row;
TESTING_MALLOC_CPU( row, MKL_INT, num_rows );
for( magma_int_t t=0; t < hA.num_rows; ++t ) {
row[ t ] = hA.col[ t ];
}
// === Call MKL with consecutive SpMVs, using mkl_ccsrmv ===
// warmp up
mkl_ccsrmv( "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_ccsrmv( "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_ccsrmm ===
char transa = 'n';
MKL_INT ldb = n, ldc=n;
char matdescra[6] = {'g', 'l', 'n', 'c', 'x', 'x'};
// warm up
mkl_ccsrmm( &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_ccsrmm( &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) );
TESTING_FREE_CPU( row );
TESTING_FREE_CPU( col );
row = NULL;
col = NULL;
#endif // MAGMA_WITH_MKL
// copy matrix to GPU
CHECK( magma_cmtransfer( hA, &dA, Magma_CPU, Magma_DEV, queue ));
// SpMV on GPU (CSR)
start = magma_sync_wtime( queue );
for (j=0; j<10; j++)
CHECK( magma_c_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) );
CHECK( magma_cmtransfer( dy, &hrefvec , Magma_DEV, Magma_CPU, queue ));
magma_cmfree(&dA, queue );
// convert to SELLP and copy to GPU
CHECK( magma_cmconvert( hA, &hA_SELLP, Magma_CSR, Magma_SELLP, queue ));
CHECK( magma_cmtransfer( hA_SELLP, &dA_SELLP, Magma_CPU, Magma_DEV, queue ));
magma_cmfree(&hA_SELLP, queue );
magma_cmfree( &dy, queue );
CHECK( magma_cvinit( &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++)
CHECK( magma_c_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) );
CHECK( magma_cmtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue ));
res = 0.0;
for(magma_int_t k=0; k<hA.num_rows; k++ )
res=res + MAGMA_C_REAL(hcheck.val[k]) - MAGMA_C_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_cmfree( &hcheck, queue );
magma_cmfree(&dA_SELLP, queue );
// SpMV on GPU (CUSPARSE - CSR)
// CUSPARSE context //
magma_cmfree( &dy, queue );
CHECK( magma_cvinit( &dy, Magma_DEV, dx.num_rows, dx.num_cols, c_zero, queue ));
//#ifdef PRECISION_d
start = magma_sync_wtime( queue );
CHECK_CUSPARSE( cusparseCreate( &cusparseHandle ));
CHECK_CUSPARSE( cusparseSetStream( cusparseHandle, queue->cuda_stream() ));
CHECK_CUSPARSE( cusparseCreateMatDescr( &descr ));
CHECK_CUSPARSE( cusparseSetMatType( descr, CUSPARSE_MATRIX_TYPE_GENERAL ));
CHECK_CUSPARSE( cusparseSetMatIndexBase( descr, CUSPARSE_INDEX_BASE_ZERO ));
magmaFloatComplex alpha = c_one;
magmaFloatComplex beta = c_zero;
// copy matrix to GPU
CHECK( magma_cmtransfer( hA, &dA, Magma_CPU, Magma_DEV, queue) );
for (j=0; j<10; j++)
cusparseCcsrmm(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) );
CHECK( magma_cmtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue ));
res = 0.0;
for(magma_int_t k=0; k<hA.num_rows; k++ )
res=res + MAGMA_C_REAL(hcheck.val[k]) - MAGMA_C_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_cmfree( &hcheck, queue );
cusparseDestroyMatDescr( descr );
cusparseDestroy( cusparseHandle );
descr = NULL;
cusparseHandle = NULL;
//#endif
printf("\n\n");
// free CPU memory
magma_cmfree(&hA, queue );
magma_cmfree(&hx, queue );
magma_cmfree(&hy, queue );
magma_cmfree(&hrefvec, queue );
// free GPU memory
magma_cmfree(&dx, queue );
magma_cmfree(&dy, queue );
magma_cmfree(&dA, queue);
i++;
}
cleanup:
#ifdef MAGMA_WITH_MKL
magma_free_cpu(pntre);
#endif
cusparseDestroyMatDescr( descr );
cusparseDestroy( cusparseHandle );
magma_cmfree(&hA, queue );
magma_cmfree(&dA, queue );
magma_cmfree(&hA_ELL, queue );
magma_cmfree(&dA_ELL, queue );
magma_cmfree(&hA_SELLP, queue );
magma_cmfree(&dA_SELLP, queue );
magma_queue_destroy( queue );
TESTING_FINALIZE();
return info;
}
/**
Purpose
-------
DSGEQRSV solves the least squares problem
min || A*X - B ||,
where A is an M-by-N matrix and X and B are M-by-NRHS matrices.
DSGEQRSV first attempts to factorize the matrix in real SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with real DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
real DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio real SINGLE PRECISION performance over real DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. M >= N >= 0.
@param[in]
nrhs INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
@param[in,out]
dA DOUBLE PRECISION array on the GPU, dimension (LDDA,N)
On entry, the M-by-N coefficient matrix A.
On exit, if iterative refinement has been successfully used
(info.EQ.0 and ITER.GE.0, see description below), A is
unchanged. If double precision factorization has been used
(info.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the QR factorization of A as returned by
function DGEQRF_GPU.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
@param[in,out]
dB DOUBLE PRECISION array on the GPU, dimension (LDDB,NRHS)
The M-by-NRHS right hand side matrix B.
May be overwritten (e.g., if refinement fails).
@param[in]
lddb INTEGER
The leading dimension of the array dB. LDDB >= max(1,M).
@param[out]
dX DOUBLE PRECISION array on the GPU, dimension (LDDX,NRHS)
If info = 0, the N-by-NRHS solution matrix X.
@param[in]
lddx INTEGER
The leading dimension of the array dX. LDDX >= max(1,N).
@param[out]
iter INTEGER
- < 0: iterative refinement has failed, double precision
factorization has been performed
+ -1 : the routine fell back to full precision for
implementation- or machine-specific reasons
+ -2 : narrowing the precision induced an overflow,
the routine fell back to full precision
+ -3 : failure of SGEQRF
+ -31: stop the iterative refinement after the 30th iteration
- > 0: iterative refinement has been successfully used.
Returns the number of iterations
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if info = -i, the i-th argument had an illegal value
@ingroup magma_dgels_driver
********************************************************************/
extern "C" magma_int_t
magma_dsgeqrsv_gpu(
magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaDouble_ptr dA, magma_int_t ldda,
magmaDouble_ptr dB, magma_int_t lddb,
magmaDouble_ptr dX, magma_int_t lddx,
magma_int_t *iter,
magma_int_t *info)
{
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
#define dSX(i,j) (dSX + (i) + (j)*lddsx)
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
magma_int_t ione = 1;
double *hworkd;
float *hworks;
double *tau;
float *stau;
magmaDouble_ptr dworkd;
magmaFloat_ptr dworks;
magmaDouble_ptr dR, dT;
magmaFloat_ptr dSA, dSX, dST;
double Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddsx, lddr, nb, lhwork, minmn, size, ldworkd;
/* Check arguments */
*iter = 0;
*info = 0;
if ( m < 0 )
*info = -1;
else if ( n < 0 || n > m )
*info = -2;
else if ( nrhs < 0 )
*info = -3;
else if ( ldda < max(1,m))
*info = -5;
else if ( lddb < max(1,m))
*info = -7;
else if ( lddx < max(1,n))
*info = -9;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( m == 0 || n == 0 || nrhs == 0 )
return *info;
nb = magma_get_sgeqrf_nb(m);
minmn= min(m, n);
/* dSX contains both B and X, so must be max(m or lddb,n). */
lddsa = ldda;
lddsx = max(lddb,n);
lddr = lddb;
/*
* Allocate temporary buffers
*/
/* dworks(dSA + dSX + dST) */
size = lddsa*n + lddsx*nrhs + ( 2*minmn + ((n+31)/32)*32 )*nb;
if (MAGMA_SUCCESS != magma_smalloc( &dworks, size )) {
fprintf(stderr, "Allocation of dworks failed (%d)\n", (int) size);
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dSA = dworks;
dSX = dSA + lddsa*n;
dST = dSX + lddsx*nrhs;
/* dworkd(dR) = lddr*nrhs */
ldworkd = lddr*nrhs;
if (MAGMA_SUCCESS != magma_dmalloc( &dworkd, ldworkd )) {
magma_free( dworks );
fprintf(stderr, "Allocation of dworkd failed\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dR = dworkd;
/* hworks(workspace for cgeqrs + stau) = min(m,n) + lhworks */
lhwork = (m - n + nb)*(nrhs + nb) + nrhs*nb;
size = lhwork + minmn;
magma_smalloc_cpu( &hworks, size );
if ( hworks == NULL ) {
magma_free( dworks );
magma_free( dworkd );
fprintf(stderr, "Allocation of hworks failed\n");
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
stau = hworks + lhwork;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_dlange(MagmaInfNorm, m, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
magmablas_dlag2s( m, nrhs, dB, lddb, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_dlag2s( m, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_sgeqrf_gpu( m, n, dSA, lddsa, stau, dST, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// solve dSA*dSX = dB in single precision
magma_sgeqrs_gpu( m, n, nrhs, dSA, lddsa, stau, dST, dSX, lddsx, hworks, lhwork, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// residual dR = dB - dA*dX in double precision
magmablas_slag2d( n, nrhs, dSX, lddsx, dX, lddx, info );
magmablas_dlacpy( MagmaUpperLower, m, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_dgemv( MagmaNoTrans, m, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( MagmaNoTrans, MagmaNoTrans, m, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_D_ABS( dX(i,j) ) instead of dlange?
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( m, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
/* Free workspaces */
magma_free( dworks );
magma_free( dworkd );
magma_free_cpu( hworks );
return *info;
REFINEMENT:
/* TODO: this iterative refinement algorithm works only for compatibile
* systems (B in colspan of A).
* See Matrix Computations (3rd ed) p. 267 for correct algorithm. */
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
magmablas_dlag2s( m, nrhs, dR, lddr, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// solve dSA*dSX = R in single precision
magma_sgeqrs_gpu( m, n, nrhs, dSA, lddsa, stau, dST, dSX, lddsx, hworks, lhwork, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Add correction and setup residual
// dX += dSX [including conversion] --and--
// dR[1:n] = dB[1:n] (only n rows, not whole m rows! -- useless if m > n)
for( j=0; j < nrhs; j++ ) {
magmablas_dsaxpycp( n, dSX(0,j), dX(0,j), dB(0,j), dR(0,j) );
}
// dR = dB (whole m rows)
magmablas_dlacpy( MagmaUpperLower, m, nrhs, dB, lddb, dR, lddr );
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_dgemv( MagmaNoTrans, m, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( MagmaNoTrans, MagmaNoTrans, m, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER > 0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( m, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
/* Free workspaces */
magma_free( dworks );
magma_free( dworkd );
magma_free_cpu( hworks );
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_free( dworks );
magma_free_cpu( hworks );
/*
* Allocate temporary buffers
*/
/* dworkd = dT for dgeqrf */
nb = magma_get_dgeqrf_nb( m );
size = (2*min(m, n) + (n+31)/32*32 )*nb;
if ( size > ldworkd ) {
magma_free( dworkd );
if (MAGMA_SUCCESS != magma_dmalloc( &dworkd, size )) {
fprintf(stderr, "Allocation of dworkd2 failed\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
dT = dworkd;
/* hworkd(dtau + workspace for dgeqrs) = min(m,n) + lhwork */
size = lhwork + minmn;
magma_dmalloc_cpu( &hworkd, size );
if ( hworkd == NULL ) {
magma_free( dworkd );
fprintf(stderr, "Allocation of hworkd2 failed\n");
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
tau = hworkd + lhwork;
magma_dgeqrf_gpu( m, n, dA, ldda, tau, dT, info );
if (*info == 0) {
// if m > n, then dB won't fit in dX, so solve with dB and copy n rows to dX
magma_dgeqrs_gpu( m, n, nrhs, dA, ldda, tau, dT, dB, lddb, hworkd, lhwork, info );
magmablas_dlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
}
magma_free( dworkd );
magma_free_cpu( hworkd );
return *info;
}
extern "C" magma_int_t
magma_sgetrf_msub(
magma_trans_t trans, magma_int_t num_subs, magma_int_t num_gpus,
magma_int_t m, magma_int_t n,
magmaFloat_ptr *d_lA, size_t dlA_offset, magma_int_t ldda,
magma_int_t *ipiv,
magma_queue_t *queues,
magma_int_t *info)
{
/* -- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
SGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Arguments
=========
NUM_GPUS
(input) INTEGER
The number of GPUS to be used for the factorization.
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) REAL array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
> 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define d_lAT(id,i,j) d_lAT[(id)], (((i)*nb)*lddat + (j)*nb)
#define d_lA( id,i,j) d_lA[(id)], (((i)*nb)+ldda * (j)*nb)
magma_int_t maxm, tot_subs = num_subs*num_gpus;
magma_int_t i, j, d, lddat;
/* submatrix info */
magma_int_t nb, n_local[ MagmaMaxSubs * MagmaMaxGPUs ];
magmaFloat_ptr d_lAT[ MagmaMaxSubs * MagmaMaxGPUs ];
/* local workspace per GPU */
magmaFloat_ptr d_panel[ MagmaMaxGPUs ];
magmaFloat_ptr d_lAP[ MagmaMaxGPUs ];
float *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -2;
else if (n < 0)
*info = -3;
else if (trans == MagmaTrans && ldda < max(1,m))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
nb = magma_get_sgetrf_nb(m);
if (nb <= 1 || nb >= n) {
/* Use CPU code. */
magma_smalloc_cpu( &work, m * n );
if (work == NULL) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
printf( "trans %c, m %d, n %d\n", lapacke_trans_const(trans), m, n );
magma_sgetmatrix( m, n, d_lA[0], 0, ldda, work, m, queues[0] );
lapackf77_sgetrf( &m, &n, work, &m, ipiv, info );
magma_ssetmatrix( m, n, work, m, d_lA[0], 0, ldda, queues[0] );
magma_free_cpu( work );
} else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
if (tot_subs > ceil((float)n/nb)) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) tot_subs );
*info = -1;
return *info;
}
/* allocate workspace for each GPU */
lddat = n/nb; /* number of block columns */
lddat = lddat/tot_subs; /* number of block columns per GPU */
lddat = nb*lddat; /* number of columns per GPU */
if (lddat * tot_subs < n) {
/* left over */
if (n-lddat*tot_subs >= nb) {
lddat += nb;
} else {
lddat += (n-lddat*tot_subs)%nb;
}
}
lddat = ((lddat+31)/32)*32; /* make it a multiple of 32 */
/* allocating workspace */
for (d=0; d < num_gpus; d++) {
//#define SINGLE_GPU_PER_CONTEXT
#ifdef SINGLE_GPU_PER_CONTEXT
if ((MAGMA_SUCCESS != magma_smalloc_mgpu( d, &d_panel[d], (2+num_gpus)*nb*maxm )) ||
(MAGMA_SUCCESS != magma_smalloc_mgpu( d, &d_lAP[d], (2+num_gpus)*nb*maxm )) ) {
#else
if ((MAGMA_SUCCESS != magma_smalloc( &d_panel[d], (2+num_gpus)*nb*maxm )) ||
(MAGMA_SUCCESS != magma_smalloc( &d_lAP[d], (2+num_gpus)*nb*maxm )) ) {
#endif
for( i=0; i < d; i++ ) {
magma_free( d_panel[i] );
magma_free( d_lAP[i] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
/* transposing the local matrix */
for (i=0; i < tot_subs; i++) {
/* local-n and local-ld */
n_local[i] = ((n/nb)/tot_subs)*nb;
if (i < (n/nb)%tot_subs)
n_local[i] += nb;
else if (i == (n/nb)%tot_subs)
n_local[i] += n%nb;
/* local-matrix storage */
if (trans == MagmaNoTrans) {
d_lAT[i] = d_lA[i];
} else {
if ( m == n_local[i] ) {
d_lAT[i] = d_lA[i];
magmablas_stranspose_inplace( m, d_lA[i], 0, ldda, queues[2*(i%num_gpus)+1] );
} else {
#ifdef SINGLE_GPU_PER_CONTEXT
if (MAGMA_SUCCESS != magma_smalloc_mgpu( i%num_gpus, &d_lAT[i], lddat*maxm )) {
#else
if (MAGMA_SUCCESS != magma_smalloc( &d_lAT[i], lddat*maxm )) {
#endif
for (j=0; j <= i; j++) {
magma_free( d_panel[j] );
magma_free( d_lAP[j] );
}
for (j=0; j < i; j++) {
if (d_lAT[j] != d_lA[j]) magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_stranspose( m, n_local[i], d_lA[i], 0, ldda, d_lAT[i], 0, lddat, queues[2*(i%num_gpus)+1]);
}
}
}
if (trans == MagmaNoTrans) {
for (d=0; d < num_gpus; d++){
magma_queue_sync(queues[2*d+1]);
}
}
/* cpu workspace */
#ifdef USE_PINNED_CLMEMORY
cl_mem buffer = clCreateBuffer(gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, sizeof(float)*maxm*nb*(1+num_gpus), NULL, NULL);
for (d=0; d < num_gpus; d++) {
work = (float*)clEnqueueMapBuffer(queues[2*d], buffer, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0,
sizeof(float)*maxm*nb*(1+num_gpus), 0, NULL, NULL, NULL);
}
#else
if (MAGMA_SUCCESS != magma_smalloc_cpu( &work, maxm*nb*(1+num_gpus) )) {
for(d=0; d < num_gpus; d++ ) magma_free( d_panel[d] );
for(d=0; d < tot_subs; d++ ) {
if( d_lAT[d] != d_lA[d] ) magma_free( d_lAT[d] );
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
/* calling multi-gpu interface with allocated workspaces and streams */
magma_sgetrf2_msub(num_subs, num_gpus, m, n, nb, 0, d_lAT, 0, lddat, ipiv, d_lAP, d_panel, 0, work, maxm,
queues, info);
/* save on output */
for (d=0; d < tot_subs; d++) {
if (trans == MagmaNoTrans) {
//magma_scopymatrix( n_local[d], m, d_lAT[d], 0, lddat, d_lA[d], 0, ldda, queues[2*d+1] );
} else {
if (d_lAT[d] == d_lA[d]) {
magmablas_stranspose_inplace( m, d_lA[d], 0, ldda, queues[2*(d%num_gpus)+1] );
} else {
magmablas_stranspose( n_local[d], m, d_lAT[d], 0, lddat, d_lA[d], 0, ldda, queues[2*(d%num_gpus)+1] );
}
}
}
/* clean up */
for (d=0; d < num_gpus; d++) {
magma_queue_sync(queues[2*d+1]);
magma_free( d_panel[d] );
magma_free( d_lAP[d] );
d_panel[d] = d_lAP[d] = NULL;
}
for (d=0; d < tot_subs; d++) {
if (d_lAT[d] != d_lA[d]) {
magma_free( d_lAT[d] );
d_lAT[d] = NULL;
}
}
#ifdef USE_PINNED_CLMEMORY
for (d=0; d < num_gpus; d++) {
clEnqueueUnmapMemObject(queues[2*d], buffer, work, 0, NULL, NULL);
}
clReleaseMemObject( buffer );
#else
magma_free_cpu( work );
#endif
work = NULL;
}
return *info;
/* End of MAGMA_SGETRF_MSUB */
}
static void magma_ctile_bulge_applyQ(
magma_int_t core_id, magma_side_t side, magma_int_t n_loc, magma_int_t n, magma_int_t nb, magma_int_t Vblksiz,
magmaFloatComplex *E, magma_int_t lde,
magmaFloatComplex *V, magma_int_t ldv,
magmaFloatComplex *TAU, magmaFloatComplex *T, magma_int_t ldt)
//, magma_int_t* info)
{
//%===========================
//% local variables
//%===========================
magma_int_t firstcolj;
magma_int_t bg, rownbm;
magma_int_t st,ed,fst,vlen,vnb,colj;
magma_int_t vpos,tpos;
magma_int_t cur_blksiz,avai_blksiz, ncolinvolvd;
magma_int_t nbgr, colst, coled;
if (n <= 0)
return;
if (n_loc <= 0)
return;
//info = 0;
magma_int_t INFO=0;
magma_int_t nbGblk = magma_ceildiv(n-1, Vblksiz);
/*
* version v1: for each chunck it apply all the V's then move to
* the other chunck. the locality here inside each
* chunck meaning that thread t apply V_k then move
* to V_k+1 which overlap with V_k meaning that the
* E_k+1 overlap with E_k. so here is the
* locality however thread t had to read V_k+1 and
* T_k+1 at each apply. note that all thread if they
* run at same speed they might reading the same V_k
* and T_k at the same time.
* */
magma_int_t nb_loc = 128; //$$$$$$$$
magma_int_t lwork = 2*nb_loc*max(Vblksiz,64);
magmaFloatComplex *work, *work2;
magma_cmalloc_cpu(&work, lwork);
magma_cmalloc_cpu(&work2, lwork);
magma_int_t nbchunk = magma_ceildiv(n_loc, nb_loc);
/* SIDE LEFT meaning apply E = Q*E = (q_1*q_2*.....*q_n) * E ==> so traverse Vs in reverse order (forward) from q_n to q_1
* each q_i consist of applying V to a block of row E(row_i,:) and applies are overlapped meaning
* that q_i+1 overlap a portion of the E(row_i, :).
* IN parallel E is splitten in vertical block over the threads */
/* SIDE RIGHT meaning apply E = E*Q = E * (q_1*q_2*.....*q_n) ==> so tarverse Vs in normal order (forward) from q_1 to q_n
* each q_i consist of applying V to a block of col E(:, col_i,:) and the applies are overlapped meaning
* that q_i+1 overlap a portion of the E(:, col_i).
* IN parallel E is splitten in horizontal block over the threads */
#ifdef ENABLE_DEBUG
if ((core_id == 0) || (core_id == 1))
printf(" APPLY Q2_cpu cbulge_back N %d N_loc %d nbchunk %d NB %d Vblksiz %d SIDE %c \n", n, n_loc, nbchunk, nb, Vblksiz, side);
#endif
for (magma_int_t i = 0; i < nbchunk; i++) {
magma_int_t ib_loc = min(nb_loc, (n_loc - i*nb_loc));
if (side == MagmaLeft) {
for (bg = nbGblk; bg > 0; bg--) {
firstcolj = (bg-1)*Vblksiz + 1;
rownbm = magma_ceildiv((n-(firstcolj+1)),nb);
if (bg == nbGblk) rownbm = magma_ceildiv((n-(firstcolj)),nb); // last blk has size=1 used for complex to handle A(N,N-1)
for (magma_int_t j = rownbm; j > 0; j--) {
vlen = 0;
vnb = 0;
colj = (bg-1)*Vblksiz; // for k=0; I compute the fst and then can remove it from the loop
fst = (rownbm -j)*nb+colj +1;
for (magma_int_t k=0; k < Vblksiz; k++) {
colj = (bg-1)*Vblksiz + k;
st = (rownbm -j)*nb+colj +1;
ed = min(st+nb-1,n-1);
if (st > ed)
break;
if ((st == ed) && (colj != n-2))
break;
vlen = ed-fst+1;
vnb = k+1;
}
colst = (bg-1)*Vblksiz;
magma_bulge_findVTpos(n, nb, Vblksiz, colst, fst, ldv, ldt, &vpos, &tpos);
if ((vlen > 0) && (vnb > 0)) {
lapackf77_clarfb( "L", "N", "F", "C", &vlen, &ib_loc, &vnb, V(vpos), &ldv, T(tpos), &ldt, E(fst,i*nb_loc), &lde, work, &ib_loc);
}
if (INFO != 0)
printf("ERROR CUNMQR INFO %d \n", (int) INFO);
}
}
} else if (side == MagmaRight) {
rownbm = magma_ceildiv((n-1),nb);
for (magma_int_t k = 1; k <= rownbm; k++) {
ncolinvolvd = min(n-1, k*nb);
avai_blksiz = min(Vblksiz,ncolinvolvd);
nbgr = magma_ceildiv(ncolinvolvd,avai_blksiz);
for (magma_int_t j = 1; j <= nbgr; j++) {
vlen = 0;
vnb = 0;
cur_blksiz = min(ncolinvolvd-(j-1)*avai_blksiz, avai_blksiz);
colst = (j-1)*avai_blksiz;
coled = colst + cur_blksiz -1;
fst = (rownbm -k)*nb+colst +1;
for (colj=colst; colj <= coled; colj++) {
st = (rownbm -k)*nb+colj +1;
ed = min(st+nb-1,n-1);
if (st > ed)
break;
if ((st == ed) && (colj != n-2))
break;
vlen = ed-fst+1;
vnb = vnb+1;
}
magma_bulge_findVTpos(n, nb, Vblksiz, colst, fst, ldv, ldt, &vpos, &tpos);
if ((vlen > 0) && (vnb > 0)) {
lapackf77_clarfb( "R", "N", "F", "C", &ib_loc, &vlen, &vnb, V(vpos), &ldv, T(tpos), &ldt, E(i*nb_loc,fst), &lde, work, &ib_loc);
}
}
}
} else {
printf("ERROR SIDE %d \n",side);
}
} // END loop over the chunks
magma_free_cpu(work);
magma_free_cpu(work2);
}
extern "C" magma_int_t magma_chetrd_hb2st(magma_int_t threads, char uplo, magma_int_t n, magma_int_t nb, magma_int_t Vblksiz,
magmaFloatComplex *A, magma_int_t lda, float *D, float *E,
magmaFloatComplex *V, magma_int_t ldv, magmaFloatComplex *TAU, magma_int_t compT, magmaFloatComplex *T, magma_int_t ldt)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
Arguments
=========
THREADS (input) INTEGER
Specifies the number of pthreads used.
THREADS > 0
UPLO (input) CHARACTER*1
= 'U': Upper triangles of A is stored;
= 'L': Lower triangles of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
NB (input) INTEGER
The order of the band matrix A. N >= NB >= 0.
VBLKSIZ (input) INTEGER
The size of the block of householder vectors applied at once.
A (input/workspace) COMPLEX array, dimension (LDA, N)
On entry the band matrix stored in the following way:
LDA (input) INTEGER
The leading dimension of the array A. LDA >= 2*NB.
D (output) DOUBLE array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
E (output) DOUBLE array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = 'U', E(i) = A(i+1,i) if UPLO = 'L'.
V (output) COMPLEX array, dimension (BLKCNT, LDV, VBLKSIZ)
On exit it contains the blocks of householder reflectors
BLKCNT is the number of block and it is returned by the funtion MAGMA_BULGE_GET_BLKCNT.
LDV (input) INTEGER
The leading dimension of V.
LDV > NB + VBLKSIZ + 1
TAU (output) COMPLEX dimension(BLKCNT, VBLKSIZ)
???
COMPT (input) INTEGER
if COMPT = 0 T is not computed
if COMPT = 1 T is computed
T (output) COMPLEX dimension(LDT *)
if COMPT = 1 on exit contains the matrices T needed for Q2
if COMPT = 0 T is not referenced
LDT (input) INTEGER
The leading dimension of T.
LDT > Vblksiz
INFO (output) INTEGER ????????????????????????????????????????????????????????????????????????????????????
= 0: successful exit
===================================================================== */
#ifdef ENABLE_TIMER
real_Double_t timeblg=0.0;
#endif
//char uplo_[2] = {uplo, 0};
magma_int_t mklth = threads;
magma_int_t INgrsiz=1;
magma_int_t blkcnt = magma_bulge_get_blkcnt(n, nb, Vblksiz);
magma_int_t nbtiles = magma_ceildiv(n, nb);
memset(T, 0, blkcnt*ldt*Vblksiz*sizeof(magmaFloatComplex));
memset(TAU, 0, blkcnt*Vblksiz*sizeof(magmaFloatComplex));
memset(V, 0, blkcnt*ldv*Vblksiz*sizeof(magmaFloatComplex));
magma_int_t* prog;
magma_malloc_cpu((void**) &prog, (2*nbtiles+threads+10)*sizeof(magma_int_t));
memset(prog, 0, (2*nbtiles+threads+10)*sizeof(magma_int_t));
magma_cbulge_id_data* arg;
magma_malloc_cpu((void**) &arg, threads*sizeof(magma_cbulge_id_data));
pthread_t* thread_id;
magma_malloc_cpu((void**) &thread_id, threads*sizeof(pthread_t));
pthread_attr_t thread_attr;
magma_setlapack_numthreads(1);
magma_cbulge_data data_bulge(threads, n, nb, nbtiles, INgrsiz, Vblksiz, compT,
A, lda, V, ldv, TAU, T, ldt, prog);
// Set one thread per core
pthread_attr_init(&thread_attr);
pthread_attr_setscope(&thread_attr, PTHREAD_SCOPE_SYSTEM);
pthread_setconcurrency(threads);
//timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime();
#endif
// Launch threads
for (magma_int_t thread = 1; thread < threads; thread++)
{
arg[thread] = magma_cbulge_id_data(thread, &data_bulge);
pthread_create(&thread_id[thread], &thread_attr, magma_chetrd_hb2st_parallel_section, &arg[thread]);
}
arg[0] = magma_cbulge_id_data(0, &data_bulge);
magma_chetrd_hb2st_parallel_section(&arg[0]);
// Wait for completion
for (magma_int_t thread = 1; thread < threads; thread++)
{
void *exitcodep;
pthread_join(thread_id[thread], &exitcodep);
}
// timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime()-timeblg;
printf(" time BULGE+T = %f \n" ,timeblg);
#endif
magma_free_cpu(thread_id);
magma_free_cpu(arg);
magma_free_cpu(prog);
magma_setlapack_numthreads(mklth);
/*================================================
* store resulting diag and lower diag D and E
* note that D and E are always real
*================================================*/
/* Make diagonal and superdiagonal elements real,
* storing them in D and E
*/
/* In complex case, the off diagonal element are
* not necessary real. we have to make off-diagonal
* elements real and copy them to E.
* When using HouseHolder elimination,
* the CLARFG give us a real as output so, all the
* diagonal/off-diagonal element except the last one are already
* real and thus we need only to take the abs of the last
* one.
* */
#if defined(PRECISION_z) || defined(PRECISION_c)
if(uplo==MagmaLower){
for (magma_int_t i=0; i < n-1 ; i++)
{
D[i] = MAGMA_C_REAL(A[i*lda ]);
E[i] = MAGMA_C_REAL(A[i*lda+1]);
}
D[n-1] = MAGMA_C_REAL(A[(n-1)*lda]);
} else { /* MagmaUpper not tested yet */
for (magma_int_t i=0; i<n-1; i++)
{
D[i] = MAGMA_C_REAL(A[i*lda+nb]);
E[i] = MAGMA_C_REAL(A[i*lda+nb-1]);
}
D[n-1] = MAGMA_C_REAL(A[(n-1)*lda+nb]);
} /* end MagmaUpper */
#else
if( uplo == MagmaLower ){
for (magma_int_t i=0; i < n-1; i++) {
D[i] = A[i*lda]; // diag
E[i] = A[i*lda+1]; //lower diag
}
D[n-1] = A[(n-1)*lda];
} else {
for (magma_int_t i=0; i < n-1; i++) {
D[i] = A[i*lda+nb]; // diag
E[i] = A[i*lda+nb-1]; //lower diag
}
D[n-1] = A[(n-1)*lda+nb];
}
#endif
return MAGMA_SUCCESS;
}
template<typename Ty> magma_int_t
magma_ungqr_gpu(
magma_int_t m, magma_int_t n, magma_int_t k,
cl_mem dA, size_t dA_offset, magma_int_t ldda,
Ty *tau,
cl_mem dT, size_t dT_offset, magma_int_t nb,
magma_queue_t queue,
magma_int_t *info)
{
#define dA(i,j) (dA), ((i) + (j)*ldda)
#define dT(j) (dT), ((j)*nb)
static const Ty c_zero = magma_zero<Ty>();
static const Ty c_one = magma_one<Ty>();
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, lpanel;
magma_int_t i, ib, ki, kk;
magma_int_t lddwork;
cl_mem dV;
Ty *work, *panel;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (ldda < std::max(1,m)) {
*info = -5;
}
if (*info != 0) {
//magma_xerbla( __func__, -(*info));
return *info;
}
if (n <= 0) {
return *info;
}
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = std::min(k, ki+nb);
} else {
ki = 0;
kk = 0;
}
// Allocate CPU work space
// n*nb for zungqr workspace
// (m - kk)*(n - kk) for last block's panel
lwork = n*nb;
lpanel = (m - kk)*(n - kk);
magma_malloc_cpu<Ty>(&work, lwork + lpanel);
if (work == NULL) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
panel = work + lwork;
// Allocate work space on GPU
if (MAGMA_SUCCESS != magma_malloc<Ty>(&dV, ldda*nb)) {
magma_free_cpu(work);
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
// dT workspace has:
// 2*std::min(m,n)*nb for T and R^{-1} matrices from geqrf
// ((n+31)/32*32)*nb for dW larfb workspace.
lddwork = std::min(m,n);
cl_mem dW;
magma_malloc<Ty>(&dW, (((n+31)/32)*32)*nb);
cpu_lapack_ungqr_work_func<Ty> cpu_lapack_ungqr;
// Use unblocked code for the last or only block.
if (kk < n) {
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
magma_getmatrix<Ty>(m_kk, k_kk,
dA(kk, kk), ldda, panel, m_kk, queue);
LAPACKE_CHECK(cpu_lapack_ungqr(
m_kk, n_kk, k_kk,
panel, m_kk,
&tau[kk], work, lwork));
magma_setmatrix<Ty>(m_kk, n_kk,
panel, m_kk, dA(kk, kk), ldda, queue);
// Set A(1:kk,kk+1:n) to zero.
magmablas_laset<Ty>(MagmaFull, kk, n - kk, c_zero, c_zero, dA(0, kk), ldda, queue);
}
if (kk > 0) {
// Use blocked code
// stream: copy Aii to V --> laset --> laset --> larfb --> [next]
// CPU has no computation
for (i = ki; i >= 0; i -= nb) {
ib = std::min(nb, k-i);
mi = m - i;
// Copy current panel on the GPU from dA to dV
magma_copymatrix<Ty>(mi, ib,
dA(i,i), ldda,
dV, 0, ldda, queue);
// set panel to identity
magmablas_laset<Ty>(MagmaFull, i, ib, c_zero, c_zero, dA(0, i), ldda, queue);
magmablas_laset<Ty>(MagmaFull, mi, ib, c_zero, c_one, dA(i, i), ldda, queue);
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_larfb_gpu<Ty>(MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, 0, ldda, dT(i), nb,
dA(i, i), ldda, dW, 0, lddwork, queue);
}
}
}
magma_free(dV);
magma_free(dW);
magma_free_cpu(work);
return *info;
}
extern "C" int calc_numerical_range(magmaFloatComplex *M, magma_int_t M_lead_dim, float _from, float _step, magma_int_t _steps, magmaFloatComplex *pts)
{
magma_int_t idx = 0, rslt = 0;
magmaFloatComplex p, scalar;
std::complex<float> vtmp;
float j;
magmaFloatComplex *dA = nullptr;
magmaFloatComplex *dAth = NULL, *dAthT = NULL,
*dX = NULL, *dY = NULL;
float *dE = NULL;
//float *hE = NULL;
//magma_int_t *ipiv = NULL;
magma_int_t lda = M_lead_dim;
//magma_int_t ldx = lda;
magma_int_t info = 0;
magma_int_t nb = 0;
//magma_vec_t jobvl;
//magma_vec_t jobvr;
magmaFloatComplex *work = nullptr;
magma_int_t lwork = 0;
float *rwork = nullptr;
magma_int_t lrwork = 0;
magma_int_t *iwork = nullptr;
magma_int_t liwork = 0;
nb = magma_get_cgehrd_nb( M_lead_dim );
lwork = 2 * max(M_lead_dim + M_lead_dim*nb, 2 * M_lead_dim + M_lead_dim*M_lead_dim); // MagmaVec
lrwork = 1 + 5 * M_lead_dim + 2 * M_lead_dim*M_lead_dim; // MagmaVec
liwork = (3 + 5 * M_lead_dim); // MagmaVec
magma_imalloc_cpu(&iwork, liwork);
magma_smalloc_cpu(&rwork, lrwork);
magma_cmalloc_pinned(&work, lwork);
magma_cmalloc_pinned(&dA, lda*M_lead_dim);
magma_cmalloc_pinned(&dAth, lda*M_lead_dim);
magma_cmalloc_pinned(&dAthT, lda*M_lead_dim);
magma_smalloc_pinned(&dE, M_lead_dim);
//magma_smalloc_cpu(&hE, M_lead_dim);
magma_cmalloc_pinned(&dX, M_lead_dim);
magma_cmalloc_pinned(&dY, M_lead_dim);
magma_csetmatrix(M_lead_dim, M_lead_dim, M, lda, dA, M_lead_dim, queue);
// th=[0:resolution:2*pi]
j = _from;
for (idx = 0; idx < _steps; idx++)
{
//scalar = exp( 1im * -j);
vtmp.real( 0.0f );
vtmp.imag( -j );
//vtmp = _FCbuild(0.0f, -j);
//printf("vtmp = %f + i%f\n", vtmp._Val[0], vtmp._Val[1]);
vtmp = exp(vtmp);
scalar.x = vtmp.real();
scalar.y = vtmp.imag();
//printf("scalar = %f + i%f\n", scalar.x, scalar.y);
magma_ccopy(lda * M_lead_dim, dA, 1, dAth, 1, queue);
// Ath = exp(1im * -j) * As
magma_cscal(lda * M_lead_dim, scalar, dAth, 1, queue);
//magma_cprint_gpu(N, N, dA, lda);
//magma_cprint_gpu(N, N, dAth, lda);
// AthT = (Ath + Ath')
magmablas_ctranspose_conj(M_lead_dim, M_lead_dim, dAth, M_lead_dim, dAthT, M_lead_dim, queue);
magmablas_cgeadd(M_lead_dim, M_lead_dim, MAGMA_C_MAKE(1.0f, 0.0f), dAth, M_lead_dim, dAthT, M_lead_dim, queue);
// AthT = AthT / 2
magma_cscal(lda*M_lead_dim, MAGMA_C_MAKE(0.5f, 0.0f), dAthT, 1, queue);
magma_sync_wtime(queue);
//magma_cprint_gpu(M_lead_dim, M_lead_dim, dAthT, lda);
// e, r = eig(AthT)
rslt = magma_cheevd(MagmaVec, MagmaLower,
M_lead_dim,
dAthT, lda,
dE,
work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
magma_sync_wtime(queue);
//printf("magma_cheevd info=%d\n", info);
//magma_cprint_gpu(M_lead_dim, M_lead_dim, dAthT, lda);
//magma_sprint_gpu(M_lead_dim, 1, dE, M_lead_dim);
//magma_sgetvector(M_lead_dim, dE, 1, hE, 1, queue);
//printf("%f %f\n", hE[0], hE[1]);
// p = r[:,s]' * A * r[:,s]
// r = r[:,s]
magma_ccopy(
M_lead_dim,
dAthT + (M_lead_dim*(M_lead_dim-1)), 1, // dAthT + (N), where (N) is a column offset
dX, 1,
queue);
magma_sync_wtime(queue);
//magma_cprint_gpu(M_lead_dim, 1, dX, M_lead_dim);
// pp = A * r[:,s]
magma_cgemv(MagmaNoTrans,
M_lead_dim, M_lead_dim,
MAGMA_C_MAKE(1.0f, 0.0f),
dA, lda,
dX, 1,
MAGMA_C_MAKE(0.0f, 0.0f),
dY, 1, queue);
magma_sync_wtime(queue);
//magma_cprint_gpu(M_lead_dim, 1, dY, M_lead_dim);
// p = r' * pp
p = magma_cdotc(M_lead_dim, dX, 1, dY, 1, queue);
magma_sync_wtime(queue);
pts[idx] = p;
//printf("p = %f %fi\n", p.x, p.y);
j += _step;
} // end of for (idx = 0; idx < _steps; idx++)
magma_free_pinned(dY);
magma_free_pinned(dX);
//magma_free_cpu(hE);
magma_free_pinned(dE);
magma_free_pinned(dAthT);
magma_free_pinned(dAth);
magma_free_pinned(dA);
magma_free_pinned(work);
magma_free_cpu(rwork);
magma_free_cpu(iwork);
//magma_free_cpu(w);
//magma_free_cpu(A);
return rslt;
}
extern "C" int calc_bounding_box(magmaFloatComplex *M, magma_int_t M_lead_dim, float *wReEig, float *wImEig)
{
magma_int_t rslt = 0;
//magmaFloatComplex *AT = nullptr;
magmaFloatComplex *dA = nullptr, *dAT = nullptr,
*dreA = nullptr, *dimA = nullptr;
float *dreEig = nullptr;
float *dimEig = nullptr;
//magma_int_t *ipiv = NULL;
magma_int_t lda = M_lead_dim;
//magma_int_t ldx = lda;
magma_int_t info = 0;
magma_int_t nb = 0;
//magma_vec_t jobvl;
//magma_vec_t jobvr;
magmaFloatComplex *work = nullptr;
magma_int_t lwork = 0;
float *rwork = nullptr;
magma_int_t lrwork = 0;
magma_int_t *iwork = nullptr;
magma_int_t liwork = 0;
nb = magma_get_cgehrd_nb( M_lead_dim );
lwork = 2 * (M_lead_dim + M_lead_dim*nb); // MagmaNoVec
//lwork = 2 * max(M_lead_dim + M_lead_dim*nb, 2*M_lead_dim + M_lead_dim*M_lead_dim); // MagmaVec
lrwork = M_lead_dim; // MagmaNoVec
//lrwork = 1 + 5 * M_lead_dim + 2*M_lead_dim*M_lead_dim; // MagmaVec
liwork = 1; // MagmaNoVec
//liwork = 3 + 5*M_lead_dim; // MagmaVec
magma_imalloc_cpu(&iwork, liwork);
magma_smalloc_cpu(&rwork, lrwork);
//magma_cmalloc_cpu(&A, lda*M_lead_dim);
//magma_cmalloc_cpu(&AT, lda*M_lead_dim);
//magma_smalloc_cpu(&reEig, M_lead_dim);
//magma_smalloc_cpu(&imEig, M_lead_dim);
magma_cmalloc_pinned(&dA, lda*M_lead_dim);
magma_cmalloc_pinned(&dAT, lda*M_lead_dim);
magma_cmalloc_pinned(&dreA, lda*M_lead_dim);
magma_cmalloc_pinned(&dimA, lda*M_lead_dim);
//magma_cmalloc_pinned(&VL, lda*M_lead_dim);
//magma_cmalloc_pinned(&VR, lda*M_lead_dim);
magma_cmalloc_pinned(&work, lwork);
magma_smalloc_pinned(&dreEig, M_lead_dim);
magma_smalloc_pinned(&dimEig, M_lead_dim);
//matrix_fillzero(AT, M_lead_dim);
//vector_fillzero(reEig, M_lead_dim);
//vector_fillzero(imEig, M_lead_dim);
//prepare_matrix_2(M);
magma_csetmatrix(M_lead_dim, M_lead_dim, M, lda, dA, M_lead_dim, queue);
//magma_csetmatrix(M_lead_dim, M_lead_dim, AT, lda, dAT, M_lead_dim, queue);
//magma_ssetvector(M_lead_dim, wReEig, 1, dreEig, 1, queue);
//magma_ssetvector(M_lead_dim, wImEig, 1, dimEig, 1, queue);
//magma_cprint_gpu(M_lead_dim, M_lead_dim, dA, lda);
// reA = ( (A + A')/2.0 )
// A'
magmablas_ctranspose(M_lead_dim, M_lead_dim, dA, M_lead_dim, dAT, M_lead_dim, queue);
//magma_cprint_gpu(M_lead_dim, M_lead_dim, dAT, lda);
// AT = A + A'
magmablas_cgeadd(M_lead_dim, M_lead_dim, MAGMA_C_MAKE(1.0f, 0.0f), dA, M_lead_dim, dAT, M_lead_dim, queue);
//magma_cprint_gpu(M_lead_dim, M_lead_dim, dAT, lda);
// AT=AT*0.5
magma_cscal(lda*M_lead_dim, MAGMA_C_MAKE(0.5f, 0.0f), dAT, 1, queue);
//magma_cprint_gpu(M_lead_dim, M_lead_dim, dAT, lda);
// reA = AT
magma_ccopy(lda*M_lead_dim, dAT, 1, dreA, 1, queue);
//magma_cprint_gpu(M_lead_dim, M_lead_dim, dreA, lda);
magma_sync_wtime(queue);
// imA = ( -1im*(A - A')/2.0 )
// A'
magmablas_ctranspose(M_lead_dim, M_lead_dim, dA, M_lead_dim, dAT, M_lead_dim, queue);
//magma_cprint_gpu(M_lead_dim, M_lead_dim, dAT, lda);
// AT = A + A'
magmablas_cgeadd(M_lead_dim, M_lead_dim, MAGMA_C_MAKE(-1.0f, 0.0f), dAT, M_lead_dim, dA, M_lead_dim, queue);
// A=A*-1j*0.5
magma_cscal(lda*M_lead_dim, MAGMA_C_MAKE(0.0f, -0.5f), dA, 1, queue);
// imA = A
magma_ccopy(lda*M_lead_dim, dA, 1, dimA, 1, queue);
magma_sync_wtime(queue);
//magma_cprint_gpu(M_lead_dim, M_lead_dim, dreA, lda);
//magma_cprint_gpu(M_lead_dim, M_lead_dim, dimA, lda);
// reEig::Vector=eigvals(reA)
rslt = magma_cheevd(MagmaNoVec, MagmaLower,
M_lead_dim,
dreA, lda,
dreEig,
work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
// imEig::Vector=eigvals(imA)
rslt = magma_cheevd(MagmaNoVec, MagmaLower,
M_lead_dim,
dimA, lda,
dimEig,
work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
//magma_sprint_gpu(M_lead_dim, 1, dreEig, M_lead_dim);
//magma_sprint_gpu(M_lead_dim, 1, dimEig, M_lead_dim);
magma_sgetvector(M_lead_dim, dreEig, 1, wReEig, 1, queue);
//magma_sync_wtime(queue);
magma_sgetvector(M_lead_dim, dimEig, 1, wImEig, 1, queue);
//magma_sync_wtime(queue);
/*
maxReIdx = magma_isamax(M_lead_dim, dreEig, 1, queue) - 1;
minReIdx = magma_isamin(M_lead_dim, dreEig, 1, queue) - 1;
maxImIdx = magma_isamax(M_lead_dim, dimEig, 1, queue) - 1;
minImIdx = magma_isamin(M_lead_dim, dimEig, 1, queue) - 1;
printf("max re idx = %d\nmin re idx = %d\n", maxReIdx, minReIdx);
printf("%f %f\n", wReEig[maxReIdx], wReEig[minReIdx]);
printf("max im idx = %d\nmin im idx = %d\n", maxImIdx, minImIdx);
printf("%f %f\n", wImEig[maxImIdx], wImEig[minImIdx]);
*/
//printf("test wReEig: %f %f\n", wReEig[0], wReEig[1]);
//printf("test wImEig: %f %f\n", wImEig[0], wImEig[1]);
magma_free_cpu(iwork);
magma_free_cpu(rwork);
//magma_free_cpu(AT);
magma_free_pinned(dA);
magma_free_pinned(dAT);
magma_free_pinned(dreA);
magma_free_pinned(dimA);
magma_free_pinned(work);
magma_free_pinned(dreEig);
magma_free_pinned(dimEig);
return rslt;
}
extern "C" magma_int_t
magma_dgehrd(magma_int_t n, magma_int_t ilo, magma_int_t ihi,
double *A, magma_int_t lda,
double *tau,
double *work, magma_int_t lwork,
double *dT,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
DGEHRD reduces a DOUBLE_PRECISION general matrix A to upper Hessenberg form H by
an orthogonal similarity transformation: Q' * A * Q = H . This version
stores the triangular matrices used in the factorization so that they can
be applied directly (i.e., without being recomputed) later. As a result,
the application of Q is much faster.
Arguments
=========
N (input) INTEGER
The order of the matrix A. N >= 0.
ILO (input) INTEGER
IHI (input) INTEGER
It is assumed that A is already upper triangular in rows
and columns 1:ILO-1 and IHI+1:N. ILO and IHI are normally
set by a previous call to DGEBAL; otherwise they should be
set to 1 and N respectively. See Further Details.
1 <= ILO <= IHI <= N, if N > 0; ILO=1 and IHI=0, if N=0.
A (input/output) DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the N-by-N general matrix to be reduced.
On exit, the upper triangle and the first subdiagonal of A
are overwritten with the upper Hessenberg matrix H, and the
elements below the first subdiagonal, with the array TAU,
represent the orthogonal matrix Q as a product of elementary
reflectors. See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) DOUBLE_PRECISION array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details). Elements 1:ILO-1 and IHI:N-1 of TAU are set to
zero.
WORK (workspace/output) DOUBLE_PRECISION array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK. LWORK >= max(1,N).
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
dT (output) DOUBLE_PRECISION array on the GPU, dimension NB*N,
where NB is the optimal blocksize. It stores the NB*NB blocks
of the triangular T matrices used in the reduction.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value.
Further Details
===============
The matrix Q is represented as a product of (ihi-ilo) elementary
reflectors
Q = H(ilo) H(ilo+1) . . . H(ihi-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0, v(i+1) = 1 and v(ihi+1:n) = 0; v(i+2:ihi) is stored on
exit in A(i+2:ihi,i), and tau in TAU(i).
The contents of A are illustrated by the following example, with
n = 7, ilo = 2 and ihi = 6:
on entry, on exit,
( a a a a a a a ) ( a a h h h h a )
( a a a a a a ) ( a h h h h a )
( a a a a a a ) ( h h h h h h )
( a a a a a a ) ( v2 h h h h h )
( a a a a a a ) ( v2 v3 h h h h )
( a a a a a a ) ( v2 v3 v4 h h h )
( a ) ( a )
where a denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
This version stores the T matrices in dT, for later use in magma_dorghr.
===================================================================== */
#define A( i, j ) ( A + (i) + (j)*lda)
#define dA( i, j ) (dA + (i) + (j-ilo)*ldda)
double c_one = MAGMA_D_ONE;
double c_zero = MAGMA_D_ZERO;
magma_int_t nb = magma_get_dgehrd_nb(n);
magma_int_t ldda = n; // assumed in dlahru
magma_int_t nh, iws;
magma_int_t iinfo;
magma_int_t ldwork;
magma_int_t lquery;
*info = 0;
iws = n*nb;
MAGMA_D_SET2REAL( work[0], (double) iws );
lquery = lwork == -1;
if (n < 0) {
*info = -1;
} else if (ilo < 1 || ilo > max(1,n)) {
*info = -2;
} else if (ihi < min(ilo,n) || ihi > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
// Adjust from 1-based indexing
ilo -= 1;
// Quick return if possible
nh = ihi - ilo;
if (nh <= 1) {
work[0] = c_one;
return *info;
}
// GPU workspace is:
// nb*ldda for dwork for dlahru
// nb*ldda for dV
// n*ldda for dA
double *dwork;
if (MAGMA_SUCCESS != magma_dmalloc( &dwork, 2*nb*ldda + n*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
double *dV = dwork + nb*ldda;
double *dA = dwork + nb*ldda*2;
ldwork = n;
magma_int_t i;
double *T, *dTi;
magma_dmalloc_cpu( &T, nb*nb );
if ( T == NULL ) {
magma_free( dwork );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
// zero first block of V, which is lower triangular
dzero_nbxnb_block(nb, dV, ldda);
// Set elements 0:ILO-1 and IHI-1:N-2 of TAU to zero
for(i = 0; i < ilo; ++i)
tau[i] = c_zero;
for(i = max(0,ihi-1); i < n-1; ++i)
tau[i] = c_zero;
for(i=0; i < nb*nb; i += 4)
T[i] = T[i+1] = T[i+2] = T[i+3] = c_zero;
magmablas_dlaset( 'F', nb, n, dT, nb );
// If not enough workspace, use unblocked code
if ( lwork < iws ) {
nb = 1;
}
if (nb == 1 || nb > nh) {
// Use unblocked code below
i = ilo;
}
else {
// Use blocked code
// Copy the matrix to the GPU
magma_dsetmatrix( n, n-ilo, A(0,ilo), lda, dA, ldda );
for (i = ilo; i < ihi-1 - nb; i += nb) {
// Reduce columns i:i+nb-1 to Hessenberg form, returning the
// matrices V and T of the block reflector H = I - V*T*V'
// which performs the reduction, and also the matrix Y = A*V*T
// Get the current panel (no need for the 1st iteration)
magma_dgetmatrix( ihi-i, nb,
dA(i,i), ldda,
A (i,i), lda );
// add 1 to i for 1-based index
magma_dlahr2( ihi, i+1, nb,
dA(0,i),
dV,
A (0,i), lda,
&tau[i], T, nb, work, ldwork);
// Copy T from the CPU to dT on the GPU
dTi = dT + (i - ilo)*nb;
magma_dsetmatrix( nb, nb, T, nb, dTi, nb );
magma_dlahru( n, ihi, i, nb,
A (0,i), lda,
dA(0,i), // dA
dA(i,i), // dY, stored over current panel
dV, dTi, dwork );
}
// Copy remainder to host
magma_dgetmatrix( n, n-i,
dA(0,i), ldda,
A (0,i), lda );
}
// Use unblocked code to reduce the rest of the matrix
// add 1 to i for 1-based index
i += 1;
lapackf77_dgehd2(&n, &i, &ihi, A, &lda, tau, work, &iinfo);
MAGMA_D_SET2REAL( work[0], (double) iws );
magma_free( dwork );
magma_free_cpu( T );
return *info;
} /* magma_dgehrd */
/***************************************************************************//**
Purpose
-------
SSYEVDX computes selected eigenvalues and, optionally, eigenvectors
of a real symmetric matrix A. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm makes very mild assumptions about
floating point arithmetic. It will work on machines with a guard
digit in add/subtract, or on those binary machines without guard
digits which subtract like the Cray X-MP, Cray Y-MP, Cray C-90, or
Cray-2. It could conceivably fail on hexadecimal or decimal machines
without guard digits, but we know of none.
Arguments
---------
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA REAL array on the GPU,
dimension (LDDA, N).
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, if JOBZ = MagmaVec, then if INFO = 0, the first mout columns
of A contains the required
orthonormal eigenvectors of the matrix A.
If JOBZ = MagmaNoVec, then on exit the lower triangle (if UPLO=MagmaLower)
or the upper triangle (if UPLO=MagmaUpper) of A, including the
diagonal, is destroyed.
@param[in]
ldda INTEGER
The leading dimension of the array DA. LDDA >= max(1,N).
@param[in]
vl REAL
@param[in]
vu REAL
If RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
If RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[out]
mout INTEGER
The total number of eigenvalues found. 0 <= MOUT <= N.
If RANGE = MagmaRangeAll, MOUT = N, and if RANGE = MagmaRangeI, MOUT = IU-IL+1.
@param[out]
w REAL array, dimension (N)
If INFO = 0, the required mout eigenvalues in ascending order.
@param
wA (workspace) REAL array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
- If N <= 1, LWORK >= 1.
- If JOBZ = MagmaNoVec and N > 1, LWORK >= 2*N + N*NB.
- If JOBZ = MagmaVec and N > 1, LWORK >= max( 2*N + N*NB, 1 + 6*N + 2*N**2 ).
NB can be obtained through magma_get_ssytrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK and IWORK
arrays, returns these values as the first entries of the WORK
and IWORK arrays, and no error message related to LWORK or
LIWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK[0] returns the optimal LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK.
- If N <= 1, LIWORK >= 1.
- If JOBZ = MagmaNoVec and N > 1, LIWORK >= 1.
- If JOBZ = MagmaVec and N > 1, LIWORK >= 3 + 5*N.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK and
IWORK arrays, returns these values as the first entries of
the WORK and IWORK arrays, and no error message related to
LWORK or LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i and JOBZ = MagmaNoVec, then the algorithm failed
to converge; i off-diagonal elements of an intermediate
tridiagonal form did not converge to zero;
if INFO = i and JOBZ = MagmaVec, then the algorithm failed
to compute an eigenvalue while working on the submatrix
lying in rows and columns INFO/(N+1) through
mod(INFO,N+1).
Further Details
---------------
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
Modified description of INFO. Sven, 16 Feb 05.
@ingroup magma_heevdx
*******************************************************************************/
extern "C" magma_int_t
magma_ssyevdx_gpu(
magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo,
magma_int_t n,
magmaFloat_ptr dA, magma_int_t ldda,
float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *mout, float *w,
float *wA, magma_int_t ldwa,
float *work, magma_int_t lwork,
#ifdef COMPLEX
float *rwork, magma_int_t lrwork,
#endif
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
magma_int_t ione = 1;
float d__1;
float eps;
magma_int_t inde;
float anrm;
float rmin, rmax;
float sigma;
magma_int_t iinfo, lwmin;
magma_int_t lower;
magma_int_t wantz;
magma_int_t indwk2, llwrk2;
magma_int_t iscale;
float safmin;
float bignum;
magma_int_t indtau;
magma_int_t indwrk, liwmin;
magma_int_t llwork;
float smlnum;
magma_int_t lquery;
magma_int_t alleig, valeig, indeig;
magmaFloat_ptr dwork;
magma_int_t lddc = ldda;
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1 || liwork == -1);
*info = 0;
if (! (wantz || (jobz == MagmaNoVec))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (ldda < max(1,n)) {
*info = -6;
} else if (ldwa < max(1,n)) {
*info = -14;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_ssytrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( 2*n + n*nb, 1 + 6*n + 2*n*n );
liwmin = 3 + 5*n;
}
else {
lwmin = 2*n + n*nb;
liwmin = 1;
}
work[0] = magma_smake_lwork( lwmin );
iwork[0] = liwmin;
if ((lwork < lwmin) && !lquery) {
*info = -16;
} else if ((liwork < liwmin) && ! lquery) {
*info = -18;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
/* If matrix is very small, then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
magma_int_t lda = n;
float *A;
magma_smalloc_cpu( &A, lda*n );
magma_sgetmatrix( n, n, dA, ldda, A, lda, queue );
// TODO: deal with range and mout. Use lapack ssyevx? (There is no lapack ssyevdx.)
lapackf77_ssyevd( lapack_vec_const(jobz), lapack_uplo_const(uplo),
&n, A, &lda,
w, work, &lwork,
iwork, &liwork, info );
*mout = n;
magma_ssetmatrix( n, n, A, lda, dA, ldda, queue );
magma_free_cpu( A );
magma_queue_destroy( queue );
return *info;
}
// ssytrd2_gpu requires ldda*ceildiv(n,64) + 2*ldda*nb
// sormtr_gpu requires lddc*n
// slansy requires n
magma_int_t ldwork = max( ldda*magma_ceildiv(n,64) + 2*ldda*nb, lddc*n );
ldwork = max( ldwork, n );
if ( wantz ) {
// sstedx requires 3n^2/2
ldwork = max( ldwork, 3*n*(n/2 + 1) );
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork, ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Get machine constants. */
safmin = lapackf77_slamch("Safe minimum");
eps = lapackf77_slamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_ssqrt( smlnum );
rmax = magma_ssqrt( bignum );
/* Scale matrix to allowable range, if necessary. */
anrm = magmablas_slansy( MagmaMaxNorm, uplo, n, dA, ldda, dwork, ldwork, queue );
iscale = 0;
sigma = 1;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
magmablas_slascl( uplo, 0, 0, 1., sigma, n, n, dA, ldda, queue, info );
}
/* Call SSYTRD to reduce symmetric matrix to tridiagonal form. */
// ssytrd work: e (n) + tau (n) + llwork (n*nb) ==> 2n + n*nb
// sstedx work: e (n) + tau (n) + z (n*n) + llwrk2 (1 + 4*n + n^2) ==> 1 + 6n + 2n^2
inde = 0;
indtau = inde + n;
indwrk = indtau + n;
indwk2 = indwrk + n*n;
llwork = lwork - indwrk;
llwrk2 = lwork - indwk2;
magma_timer_t time=0;
timer_start( time );
#ifdef FAST_SYMV
magma_ssytrd2_gpu( uplo, n, dA, ldda, w, &work[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
dwork, ldwork, &iinfo );
#else
magma_ssytrd_gpu( uplo, n, dA, ldda, w, &work[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
&iinfo );
#endif
timer_stop( time );
timer_printf( "time ssytrd = %6.2f\n", time );
/* For eigenvalues only, call SSTERF. For eigenvectors, first call
SSTEDC to generate the eigenvector matrix, WORK(INDWRK), of the
tridiagonal matrix, then call SORMTR to multiply it to the Householder
transformations represented as Householder vectors in A. */
if (! wantz) {
lapackf77_ssterf( &n, w, &work[inde], info );
magma_smove_eig( range, n, w, &il, &iu, vl, vu, mout );
}
else {
timer_start( time );
magma_sstedx( range, n, vl, vu, il, iu, w, &work[inde],
&work[indwrk], n, &work[indwk2],
llwrk2, iwork, liwork, dwork, info );
timer_stop( time );
timer_printf( "time sstedx = %6.2f\n", time );
timer_start( time );
magma_smove_eig( range, n, w, &il, &iu, vl, vu, mout );
magma_ssetmatrix( n, *mout, &work[indwrk + n* (il-1) ], n, dwork, lddc, queue );
magma_sormtr_gpu( MagmaLeft, uplo, MagmaNoTrans, n, *mout, dA, ldda, &work[indtau],
dwork, lddc, wA, ldwa, &iinfo );
magma_scopymatrix( n, *mout, dwork, lddc, dA, ldda, queue );
timer_stop( time );
timer_printf( "time sormtr + copy = %6.2f\n", time );
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
d__1 = 1. / sigma;
blasf77_sscal( &n, &d__1, w, &ione );
}
work[0] = magma_smake_lwork( lwmin );
iwork[0] = liwmin;
magma_queue_destroy( queue );
magma_free( dwork );
return *info;
} /* magma_ssyevd_gpu */
extern "C" magma_int_t
magma_cgehrd2(magma_int_t n, magma_int_t ilo, magma_int_t ihi,
magmaFloatComplex *a, magma_int_t lda,
magmaFloatComplex *tau, magmaFloatComplex *work,
magma_int_t lwork, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGEHRD2 reduces a COMPLEX general matrix A to upper Hessenberg form H by
an orthogonal similarity transformation: Q' * A * Q = H .
Arguments
=========
N (input) INTEGER
The order of the matrix A. N >= 0.
ILO (input) INTEGER
IHI (input) INTEGER
It is assumed that A is already upper triangular in rows
and columns 1:ILO-1 and IHI+1:N. ILO and IHI are normally
set by a previous call to CGEBAL; otherwise they should be
set to 1 and N respectively. See Further Details.
1 <= ILO <= IHI <= N, if N > 0; ILO=1 and IHI=0, if N=0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the N-by-N general matrix to be reduced.
On exit, the upper triangle and the first subdiagonal of A
are overwritten with the upper Hessenberg matrix H, and the
elements below the first subdiagonal, with the array TAU,
represent the orthogonal matrix Q as a product of elementary
reflectors. See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) COMPLEX array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details). Elements 1:ILO-1 and IHI:N-1 of TAU are set to
zero.
WORK (workspace/output) COMPLEX array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK. LWORK >= max(1,N).
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value.
Further Details
===============
The matrix Q is represented as a product of (ihi-ilo) elementary
reflectors
Q = H(ilo) H(ilo+1) . . . H(ihi-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0, v(i+1) = 1 and v(ihi+1:n) = 0; v(i+2:ihi) is stored on
exit in A(i+2:ihi,i), and tau in TAU(i).
The contents of A are illustrated by the following example, with
n = 7, ilo = 2 and ihi = 6:
on entry, on exit,
( a a a a a a a ) ( a a h h h h a )
( a a a a a a ) ( a h h h h a )
( a a a a a a ) ( h h h h h h )
( a a a a a a ) ( v2 h h h h h )
( a a a a a a ) ( v2 v3 h h h h )
( a a a a a a ) ( v2 v3 v4 h h h )
( a ) ( a )
where a denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
===================================================================== */
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magma_int_t nb = magma_get_cgehrd_nb(n);
magma_int_t N = n, ldda = n;
magma_int_t ib;
magma_int_t nh, iws;
magma_int_t nbmin, iinfo;
magma_int_t ldwork;
magma_int_t lquery;
--tau;
*info = 0;
MAGMA_C_SET2REAL( work[0], (float) n * nb );
lquery = lwork == -1;
if (n < 0) {
*info = -1;
} else if (ilo < 1 || ilo > max(1,n)) {
*info = -2;
} else if (ihi < min(ilo,n) || ihi > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
nh = ihi - ilo + 1;
if (nh <= 1) {
work[0] = c_one;
return *info;
}
magmaFloatComplex *da;
if (MAGMA_SUCCESS != magma_cmalloc( &da, N*ldda + 2*N*nb + nb*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmaFloatComplex *d_A = da;
magmaFloatComplex *d_work = da + (N+nb)*ldda;
magma_int_t i__;
magmaFloatComplex *t, *d_t;
magma_cmalloc_cpu( &t, nb*nb );
if ( t == NULL ) {
magma_free( da );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
d_t = d_work + nb * ldda;
czero_nbxnb_block(nb, d_A+N*ldda, ldda);
/* Set elements 1:ILO-1 and IHI:N-1 of TAU to zero */
for (i__ = 1; i__ < ilo; ++i__)
tau[i__] = c_zero;
for (i__ = max(1,ihi); i__ < n; ++i__)
tau[i__] = c_zero;
for(i__=0; i__< nb*nb; i__+=4)
t[i__] = t[i__+1] = t[i__+2] = t[i__+3] = c_zero;
nbmin = 2;
iws = 1;
if (nb > 1 && nb < nh) {
/* Determine when to cross over from blocked to unblocked code
(last block is always handled by unblocked code) */
if (nb < nh) {
/* Determine if workspace is large enough for blocked code */
iws = n * nb;
if (lwork < iws) {
/* Not enough workspace to use optimal NB: determine the
minimum value of NB, and reduce NB or force use of
unblocked code */
nbmin = nb;
if (lwork >= n * nbmin)
nb = lwork / n;
else
nb = 1;
}
}
}
ldwork = n;
if (nb < nbmin || nb >= nh) {
/* Use unblocked code below */
i__ = ilo;
}
else {
/* Use blocked code */
/* Copy the matrix to the GPU */
magma_csetmatrix( N, N-ilo+1, a+(ilo-1)*(lda), lda, d_A, ldda );
for (i__ = ilo; i__ < ihi - nb; i__ += nb) {
/* Computing MIN */
ib = min(nb, ihi - i__);
/* Reduce columns i:i+ib-1 to Hessenberg form, returning the
matrices V and T of the block reflector H = I - V*T*V'
which performs the reduction, and also the matrix Y = A*V*T */
/* Get the current panel (no need for the 1st iteration) */
magma_cgetmatrix( ihi-i__+1, ib,
d_A + (i__ - ilo)*ldda + i__ - 1, ldda,
a + (i__ - 1 )*lda + i__ - 1, lda );
magma_clahr2(ihi, i__, ib,
d_A + (i__ - ilo)*ldda,
d_A + N*ldda + 1,
a + (i__ - 1 )*(lda) , lda,
&tau[i__], t, nb, work, ldwork);
/* Copy T from the CPU to D_T on the GPU */
magma_csetmatrix( nb, nb, t, nb, d_t, nb );
magma_clahru(n, ihi, i__ - 1, ib,
a + (i__ - 1 )*(lda), lda,
d_A + (i__ - ilo)*ldda,
d_A + (i__ - ilo)*ldda + i__ - 1,
d_A + N*ldda, d_t, d_work);
}
}
/* Use unblocked code to reduce the rest of the matrix */
if (!(nb < nbmin || nb >= nh)) {
magma_cgetmatrix( n, n-i__+1,
d_A+ (i__-ilo)*ldda, ldda,
a + (i__-1)*(lda), lda );
}
lapackf77_cgehd2(&n, &i__, &ihi, a, &lda, &tau[1], work, &iinfo);
MAGMA_C_SET2REAL( work[0], (float) iws );
magma_free( da );
magma_free_cpu(t);
return *info;
} /* magma_cgehrd2 */
/**
Purpose
-------
CUNMLQ overwrites the general complex M-by-N matrix C with
@verbatim
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q * C C * Q
TRANS = Magma_ConjTrans: Q**H * C C * Q**H
@endverbatim
where Q is a complexunitary matrix defined as the product of k
elementary reflectors
Q = H(k)**H . . . H(2)**H H(1)**H
as returned by CGELQF. Q is of order M if SIDE = MagmaLeft and of order N
if SIDE = MagmaRight.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply Q or Q**H from the Left;
- = MagmaRight: apply Q or Q**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q;
- = Magma_ConjTrans: Conjugate transpose, apply Q**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = MagmaLeft, M >= K >= 0;
if SIDE = MagmaRight, N >= K >= 0.
@param[in]
A COMPLEX array, dimension
(LDA,M) if SIDE = MagmaLeft,
(LDA,N) if SIDE = MagmaRight.
The i-th row must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
CGELQF in the first k rows of its array argument A.
A is modified by the routine but restored on exit.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,K).
@param[in]
tau COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGELQF.
@param[in,out]
C COMPLEX array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDC >= max(1,M).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK.
If SIDE = MagmaLeft, LWORK >= max(1,N);
if SIDE = MagmaRight, LWORK >= max(1,M).
For optimum performance
if SIDE = MagmaLeft, LWORK >= N*NB;
if SIDE = MagmaRight, LWORK >= M*NB,
where NB is the optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cgelqf_comp
********************************************************************/
extern "C" magma_int_t
magma_cunmlq(
magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *C, magma_int_t ldc,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define dC(i_,j_) (dC + (i_) + (j_)*lddc)
#define dV(i_,j_) (dV + (i_) + (j_)*ib)
#define dT(i_,j_) (dT + (i_) + (j_)*ib)
#define dwork(i_) (dwork + (i_))
magmaFloatComplex *T, *T2;
magma_int_t i, i1, i2, ib, ic, jc, nb, mi, ni, nq, nq_i, nw, step;
magma_int_t iinfo, ldwork, lwkopt;
magma_int_t left, notran, lquery;
magma_trans_t transt;
*info = 0;
left = (side == MagmaLeft);
notran = (trans == MagmaNoTrans);
lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
/* Test the input arguments */
if (! left && side != MagmaRight) {
*info = -1;
} else if (! notran && trans != Magma_ConjTrans) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (lda < max(1,k)) {
*info = -7;
} else if (ldc < max(1,m)) {
*info = -10;
} else if (lwork < max(1,nw) && ! lquery) {
*info = -12;
}
if (*info == 0) {
nb = magma_get_cgelqf_nb( min( m, n ));
lwkopt = max(1,nw)*nb;
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
work[0] = MAGMA_C_ONE;
return *info;
}
ldwork = nw;
if (nb >= k) {
/* Use CPU code */
lapackf77_cunmlq( lapack_side_const(side), lapack_trans_const(trans),
&m, &n, &k, A, &lda, tau, C, &ldc, work, &lwork, &iinfo);
}
else {
/* Use hybrid CPU-GPU code */
/* Allocate work space on the GPU.
* nw*nb for dwork (m or n) by nb
* nq*nb for dV (n or m) by nb
* nb*nb for dT
* lddc*n for dC.
*/
magma_int_t lddc = ((m+31)/32)*32;
magmaFloatComplex_ptr dwork, dV, dT, dC;
magma_cmalloc( &dwork, (nw + nq + nb)*nb + lddc*n );
if ( dwork == NULL ) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dV = dwork + nw*nb;
dT = dV + nq*nb;
dC = dT + nb*nb;
/* work space on CPU.
* nb*nb for T
* nb*nb for T2, used to save and restore diagonal block of panel */
magma_cmalloc_cpu( &T, 2*nb*nb );
if ( T == NULL ) {
magma_free( dwork );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
T2 = T + nb*nb;
/* Copy matrix C from the CPU to the GPU */
magma_csetmatrix( m, n, C, ldc, dC(0,0), lddc );
if ( (left && notran) || (! left && ! notran) ) {
i1 = 0;
i2 = k;
step = nb;
} else {
i1 = ((k - 1) / nb)*nb;
i2 = 0;
step = -nb;
}
// silence "uninitialized" warnings
mi = 0;
ni = 0;
if (left) {
ni = n;
jc = 0;
} else {
mi = m;
ic = 0;
}
if (notran) {
transt = Magma_ConjTrans;
} else {
transt = MagmaNoTrans;
}
for (i = i1; (step < 0 ? i >= i2 : i < i2); i += step) {
ib = min(nb, k - i);
/* Form the triangular factor of the block reflector
H = H(i) H(i + 1) . . . H(i + ib-1) */
nq_i = nq - i;
lapackf77_clarft("Forward", "Rowwise", &nq_i, &ib,
A(i,i), &lda, &tau[i], T, &ib);
/* 1) set upper triangle of panel in A to identity,
2) copy the panel from A to the GPU, and
3) restore A */
cpanel_to_q( MagmaLower, ib, A(i,i), lda, T2 );
magma_csetmatrix( ib, nq_i, A(i,i), lda, dV(0,0), ib );
cq_to_panel( MagmaLower, ib, A(i,i), lda, T2 );
if (left) {
/* H or H**H is applied to C(i:m,1:n) */
mi = m - i;
ic = i;
}
else {
/* H or H**H is applied to C(1:m,i:n) */
ni = n - i;
jc = i;
}
/* Apply H or H**H; First copy T to the GPU */
magma_csetmatrix( ib, ib, T, ib, dT(0,0), ib );
magma_clarfb_gpu( side, transt, MagmaForward, MagmaRowwise,
mi, ni, ib,
dV(0,0), ib,
dT(0,0), ib,
dC(ic,jc), lddc,
dwork(0), ldwork );
}
magma_cgetmatrix( m, n, dC(0,0), lddc, C, ldc );
magma_free( dwork );
magma_free_cpu( T );
}
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
return *info;
} /* magma_cunmlq */
extern "C" magma_int_t
magma_sbicgstab_merge3(
magma_s_matrix A, magma_s_matrix b,
magma_s_matrix *x, magma_s_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_BICGSTABMERGE;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// solver variables
float alpha, beta, omega, rho_old, rho_new, *skp_h={0};
float nom, nom0, betanom, nomb;
// some useful variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE;
magma_int_t dofs = A.num_rows;
// workspace
magma_s_matrix q={Magma_CSR}, r={Magma_CSR}, rr={Magma_CSR}, p={Magma_CSR}, v={Magma_CSR}, s={Magma_CSR}, t={Magma_CSR};
float *d1=NULL, *d2=NULL, *skp=NULL;
d1 = NULL;
d2 = NULL;
skp = NULL;
CHECK( magma_smalloc( &d1, dofs*(2) ));
CHECK( magma_smalloc( &d2, dofs*(2) ));
// array for the parameters
CHECK( magma_smalloc( &skp, 8 ));
// skp = [alpha|beta|omega|rho_old|rho|nom|tmp1|tmp2]
CHECK( magma_svinit( &q, Magma_DEV, dofs*6, 1, c_zero, queue ));
// q = rr|r|p|v|s|t
rr.memory_location = Magma_DEV; rr.dval = NULL; rr.num_rows = rr.nnz = dofs; rr.num_cols = 1; rr.storage_type = Magma_DENSE;
r.memory_location = Magma_DEV; r.dval = NULL; r.num_rows = r.nnz = dofs; r.num_cols = 1; r.storage_type = Magma_DENSE;
p.memory_location = Magma_DEV; p.dval = NULL; p.num_rows = p.nnz = dofs; p.num_cols = 1; p.storage_type = Magma_DENSE;
v.memory_location = Magma_DEV; v.dval = NULL; v.num_rows = v.nnz = dofs; v.num_cols = 1; v.storage_type = Magma_DENSE;
s.memory_location = Magma_DEV; s.dval = NULL; s.num_rows = s.nnz = dofs; s.num_cols = 1; s.storage_type = Magma_DENSE;
t.memory_location = Magma_DEV; t.dval = NULL; t.num_rows = t.nnz = dofs; t.num_cols = 1; t.storage_type = Magma_DENSE;
rr.dval = q(0);
r.dval = q(1);
p.dval = q(2);
v.dval = q(3);
s.dval = q(4);
t.dval = q(5);
// solver setup
CHECK( magma_sresidualvec( A, b, *x, &r, &nom0, queue));
magma_scopy( dofs, r.dval, 1, q(0), 1, queue ); // rr = r
magma_scopy( dofs, r.dval, 1, q(1), 1, queue ); // q = r
betanom = nom0;
nom = nom0*nom0;
rho_new = magma_sdot( dofs, r.dval, 1, r.dval, 1, queue ); // rho=<rr,r>
rho_old = omega = alpha = MAGMA_S_MAKE( 1.0, 0. );
beta = rho_new;
solver_par->init_res = nom0;
// array on host for the parameters
CHECK( magma_smalloc_cpu( &skp_h, 8 ));
nomb = magma_snrm2( dofs, b.dval, 1, queue );
if ( nomb == 0.0 ){
nomb=1.0;
}
solver_par->final_res = solver_par->init_res;
solver_par->iter_res = solver_par->init_res;
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = nom0;
solver_par->timing[0] = 0.0;
}
skp_h[0]=alpha;
skp_h[1]=beta;
skp_h[2]=omega;
skp_h[3]=rho_old;
skp_h[4]=rho_new;
skp_h[5]=MAGMA_S_MAKE(nom, 0.0);
magma_ssetvector( 8, skp_h, 1, skp, 1, queue );
CHECK( magma_s_spmv( c_one, A, r, c_zero, v, queue )); // z = A r
nomb = magma_snrm2( dofs, b.dval, 1, queue );
if( nom0 < solver_par->atol ||
nom0/nomb < solver_par->rtol ){
info = MAGMA_SUCCESS;
goto cleanup;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
solver_par->numiter = 0;
solver_par->spmv_count = 0;
// start iteration
do
{
solver_par->numiter++;
// computes p=r+beta*(p-omega*v)
CHECK( magma_sbicgmerge1( dofs, skp, v.dval, r.dval, p.dval, queue ));
CHECK( magma_s_spmv( c_one, A, p, c_zero, v, queue )); // v = Ap
solver_par->spmv_count++;
CHECK( magma_smdotc( dofs, 1, q.dval, v.dval, d1, d2, skp, queue ));
CHECK( magma_sbicgmerge4( 1, skp, queue ));
CHECK( magma_sbicgmerge2( dofs, skp, r.dval, v.dval, s.dval, queue )); // s=r-alpha*v
CHECK( magma_s_spmv( c_one, A, s, c_zero, t, queue )); // t=As
solver_par->spmv_count++;
CHECK( magma_smdotc( dofs, 2, q.dval+4*dofs, t.dval, d1, d2, skp+6, queue ));
CHECK( magma_sbicgmerge4( 2, skp, queue ));
CHECK( magma_sbicgmerge_xrbeta( dofs, d1, d2, q.dval, r.dval, p.dval,
s.dval, t.dval, x->dval, skp, queue ));
// check stopping criterion
magma_sgetvector_async( 1 , skp+5, 1, skp_h+5, 1, queue );
betanom = sqrt(MAGMA_S_REAL(skp_h[5]));
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( betanom < solver_par->atol ||
betanom/nomb < solver_par->rtol ) {
break;
}
}
while ( solver_par->numiter+1 <= solver_par->maxiter );
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
CHECK( magma_sresidualvec( A, b, *x, &r, &residual, queue));
solver_par->iter_res = betanom;
solver_par->final_res = residual;
if ( solver_par->numiter < solver_par->maxiter ) {
info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
info = MAGMA_SLOW_CONVERGENCE;
if( solver_par->iter_res < solver_par->atol ||
solver_par->iter_res/solver_par->init_res < solver_par->rtol ){
info = MAGMA_SUCCESS;
}
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
info = MAGMA_DIVERGENCE;
}
cleanup:
magma_smfree(&q, queue ); // frees all vectors
magma_free(d1);
magma_free(d2);
magma_free( skp );
magma_free_cpu( skp_h );
solver_par->info = info;
return info;
} /* sbicgstab_merge */
/**
Purpose
-------
ZGEHRD reduces a COMPLEX_16 general matrix A to upper Hessenberg form H by
an orthogonal similarity transformation: Q' * A * Q = H . This version
stores the triangular matrices used in the factorization so that they can
be applied directly (i.e., without being recomputed) later. As a result,
the application of Q is much faster.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
ilo INTEGER
@param[in]
ihi INTEGER
It is assumed that A is already upper triangular in rows
and columns 1:ILO-1 and IHI+1:N. ILO and IHI are normally
set by a previous call to ZGEBAL; otherwise they should be
set to 1 and N respectively. See Further Details.
1 <= ILO <= IHI <= N, if N > 0; ILO=1 and IHI=0, if N=0.
@param[in,out]
A COMPLEX_16 array, dimension (LDA,N)
On entry, the N-by-N general matrix to be reduced.
On exit, the upper triangle and the first subdiagonal of A
are overwritten with the upper Hessenberg matrix H, and the
elements below the first subdiagonal, with the array TAU,
represent the orthogonal matrix Q as a product of elementary
reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details). Elements 1:ILO-1 and IHI:N-1 of TAU are set to
zero.
@param[out]
work (workspace) COMPLEX_16 array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK. LWORK >= max(1,N).
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
dT COMPLEX_16 array on the GPU, dimension NB*N,
where NB is the optimal blocksize. It stores the NB*NB blocks
of the triangular T matrices used in the reduction.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value.
Further Details
---------------
The matrix Q is represented as a product of (ihi-ilo) elementary
reflectors
Q = H(ilo) H(ilo+1) . . . H(ihi-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0, v(i+1) = 1 and v(ihi+1:n) = 0; v(i+2:ihi) is stored on
exit in A(i+2:ihi,i), and tau in TAU(i).
The contents of A are illustrated by the following example, with
n = 7, ilo = 2 and ihi = 6:
@verbatim
on entry, on exit,
( a a a a a a a ) ( a a h h h h a )
( a a a a a a ) ( a h h h h a )
( a a a a a a ) ( h h h h h h )
( a a a a a a ) ( v2 h h h h h )
( a a a a a a ) ( v2 v3 h h h h )
( a a a a a a ) ( v2 v3 v4 h h h )
( a ) ( a )
@endverbatim
where a denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
This version stores the T matrices in dT, for later use in magma_zunghr.
@ingroup magma_zgeev_comp
********************************************************************/
extern "C" magma_int_t
magma_zgehrd(
magma_int_t n, magma_int_t ilo, magma_int_t ihi,
magmaDoubleComplex *A, magma_int_t lda,
magmaDoubleComplex *tau,
magmaDoubleComplex *work, magma_int_t lwork,
magmaDoubleComplex_ptr dT,
magma_int_t *info)
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magma_int_t nb = magma_get_zgehrd_nb(n);
magma_int_t ldda = ((n+31)/32)*32;
magma_int_t i, nh, iws;
magma_int_t iinfo;
magma_int_t lquery;
*info = 0;
iws = n*nb;
work[0] = MAGMA_Z_MAKE( iws, 0 );
lquery = (lwork == -1);
if (n < 0) {
*info = -1;
} else if (ilo < 1 || ilo > max(1,n)) {
*info = -2;
} else if (ihi < min(ilo,n) || ihi > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
// Adjust from 1-based indexing
ilo -= 1;
// Quick return if possible
nh = ihi - ilo;
if (nh <= 1) {
work[0] = c_one;
return *info;
}
// If not enough workspace, use unblocked code
if ( lwork < iws ) {
nb = 1;
}
if (nb == 1 || nb > nh) {
// Use unblocked code below
i = ilo;
}
else {
// Use blocked code
// GPU workspace is:
// nb*ldda for dwork for zlahru
// nb*ldda for dV
// n*ldda for dA
magmaDoubleComplex *dwork;
if (MAGMA_SUCCESS != magma_zmalloc( &dwork, 2*nb*ldda + n*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmaDoubleComplex *dV = dwork + nb*ldda;
magmaDoubleComplex *dA = dwork + nb*ldda*2;
magmaDoubleComplex *dTi;
magmaDoubleComplex *T;
magma_zmalloc_cpu( &T, nb*nb );
if ( T == NULL ) {
magma_free( dwork );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
// zero first block of V, which is lower triangular
magmablas_zlaset( MagmaFull, nb, nb, c_zero, c_zero, dV, ldda );
// Set elements 0:ILO-1 and IHI-1:N-2 of TAU to zero
for (i = 0; i < ilo; ++i)
tau[i] = c_zero;
for (i = max(0,ihi-1); i < n-1; ++i)
tau[i] = c_zero;
assert( nb % 4 == 0 );
for (i=0; i < nb*nb; i += 4)
T[i] = T[i+1] = T[i+2] = T[i+3] = c_zero;
magmablas_zlaset( MagmaFull, nb, n, c_zero, c_zero, dT, nb );
// Copy the matrix to the GPU
magma_zsetmatrix( n, n-ilo, A(0,ilo), lda, dA, ldda );
for (i = ilo; i < ihi-1 - nb; i += nb) {
// Reduce columns i:i+nb-1 to Hessenberg form, returning the
// matrices V and T of the block reflector H = I - V*T*V'
// which performs the reduction, and also the matrix Y = A*V*T
// Get the current panel (no need for the 1st iteration)
magma_zgetmatrix( ihi-i, nb,
dA(i,i-ilo), ldda,
A(i,i), lda );
// add 1 to i for 1-based index
magma_zlahr2( ihi, i+1, nb,
dA(0,i-ilo), ldda,
dV, ldda,
A(0,i), lda,
&tau[i], T, nb, work, n);
// Copy T from the CPU to dT on the GPU
dTi = dT + (i - ilo)*nb;
magma_zsetmatrix( nb, nb, T, nb, dTi, nb );
magma_zlahru( n, ihi, i, nb,
A(0,i), lda,
dA(0,i-ilo), ldda, // dA
dA(i,i-ilo), ldda, // dY, stored over current panel
dV, ldda,
dTi, dwork );
}
// Copy remainder to host
magma_zgetmatrix( n, n-i,
dA(0,i-ilo), ldda,
A(0,i), lda );
magma_free( dwork );
magma_free_cpu( T );
}
// Use unblocked code to reduce the rest of the matrix
// add 1 to i for 1-based index
i += 1;
lapackf77_zgehd2(&n, &i, &ihi, A, &lda, tau, work, &iinfo);
work[0] = MAGMA_Z_MAKE( iws, 0 );
return *info;
} /* magma_zgehrd */
/**
Purpose
-------
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangles of A is stored;
- = MagmaLower: Lower triangles of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
nb INTEGER
The order of the band matrix A. N >= NB >= 0.
@param[in]
Vblksiz INTEGER
The size of the block of householder vectors applied at once.
@param[in]
A (workspace) DOUBLE_PRECISION array, dimension (LDA, N)
On entry the band matrix stored in the following way:
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= 2*NB.
@param[out]
d DOUBLE array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
@param[out]
e DOUBLE array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = MagmaUpper, E(i) = A(i+1,i) if UPLO = MagmaLower.
@param[out]
V DOUBLE_PRECISION array, dimension (BLKCNT, LDV, VBLKSIZ)
On exit it contains the blocks of householder reflectors
BLKCNT is the number of block and it is returned by the funtion MAGMA_BULGE_GET_BLKCNT.
@param[in]
ldv INTEGER
The leading dimension of V.
LDV > NB + VBLKSIZ + 1
@param[out]
TAU DOUBLE_PRECISION dimension(BLKCNT, VBLKSIZ)
???
@param[in]
compT INTEGER
if COMPT = 0 T is not computed
if COMPT = 1 T is computed
@param[out]
T DOUBLE_PRECISION dimension(LDT *)
if COMPT = 1 on exit contains the matrices T needed for Q2
if COMPT = 0 T is not referenced
@param[in]
ldt INTEGER
The leading dimension of T.
LDT > Vblksiz
@ingroup magma_dsyev_2stage
********************************************************************/
extern "C" magma_int_t
magma_dsytrd_sb2st(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb, magma_int_t Vblksiz,
double *A, magma_int_t lda, double *d, double *e,
double *V, magma_int_t ldv, double *TAU,
magma_int_t compT, double *T, magma_int_t ldt)
{
#ifdef ENABLE_TIMER
real_Double_t timeblg=0.0;
#endif
magma_int_t threads = magma_get_parallel_numthreads();
magma_int_t mklth = magma_get_lapack_numthreads();
magma_set_lapack_numthreads(1);
//const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t INgrsiz=1;
magma_int_t blkcnt = magma_bulge_get_blkcnt(n, nb, Vblksiz);
magma_int_t nbtiles = magma_ceildiv(n, nb);
memset(T, 0, blkcnt*ldt*Vblksiz*sizeof(double));
memset(TAU, 0, blkcnt*Vblksiz*sizeof(double));
memset(V, 0, blkcnt*ldv*Vblksiz*sizeof(double));
volatile magma_int_t* prog;
magma_malloc_cpu((void**) &prog, (2*nbtiles+threads+10)*sizeof(magma_int_t));
memset((void *) prog, 0, (2*nbtiles+threads+10)*sizeof(magma_int_t));
magma_dbulge_id_data* arg;
magma_malloc_cpu((void**) &arg, threads*sizeof(magma_dbulge_id_data));
pthread_t* thread_id;
magma_malloc_cpu((void**) &thread_id, threads*sizeof(pthread_t));
pthread_attr_t thread_attr;
magma_dbulge_data data_bulge;
magma_dbulge_data_init(&data_bulge, threads, n, nb, nbtiles, INgrsiz, Vblksiz, compT,
A, lda, V, ldv, TAU, T, ldt, prog);
// Set one thread per core
pthread_attr_init(&thread_attr);
pthread_attr_setscope(&thread_attr, PTHREAD_SCOPE_SYSTEM);
pthread_setconcurrency(threads);
//timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime();
#endif
// Launch threads
for (magma_int_t thread = 1; thread < threads; thread++) {
magma_dbulge_id_data_init(&(arg[thread]), thread, &data_bulge);
pthread_create(&thread_id[thread], &thread_attr, magma_dsytrd_sb2st_parallel_section, &arg[thread]);
}
magma_dbulge_id_data_init(&(arg[0]), 0, &data_bulge);
magma_dsytrd_sb2st_parallel_section(&arg[0]);
// Wait for completion
for (magma_int_t thread = 1; thread < threads; thread++) {
void *exitcodep;
pthread_join(thread_id[thread], &exitcodep);
}
// timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime()-timeblg;
printf(" time BULGE+T = %f\n", timeblg);
#endif
magma_free_cpu(thread_id);
magma_free_cpu(arg);
magma_free_cpu((void *) prog);
magma_dbulge_data_destroy(&data_bulge);
magma_set_lapack_numthreads(mklth);
/*================================================
* store resulting diag and lower diag d and e
* note that d and e are always real
*================================================*/
/* Make diagonal and superdiagonal elements real,
* storing them in d and e
*/
/* In real case, the off diagonal element are
* not necessary real. we have to make off-diagonal
* elements real and copy them to e.
* When using HouseHolder elimination,
* the DLARFG give us a real as output so, all the
* diagonal/off-diagonal element except the last one are already
* real and thus we need only to take the abs of the last
* one.
* */
#if defined(PRECISION_z) || defined(PRECISION_c)
if (uplo == MagmaLower) {
for (magma_int_t i=0; i < n-1; i++) {
d[i] = MAGMA_D_REAL( A[i*lda ] );
e[i] = MAGMA_D_REAL( A[i*lda+1] );
}
d[n-1] = MAGMA_D_REAL(A[(n-1)*lda]);
} else { /* MagmaUpper not tested yet */
for (magma_int_t i=0; i < n-1; i++) {
d[i] = MAGMA_D_REAL( A[i*lda+nb] );
e[i] = MAGMA_D_REAL( A[i*lda+nb-1] );
}
d[n-1] = MAGMA_D_REAL(A[(n-1)*lda+nb]);
} /* end MagmaUpper */
#else
if ( uplo == MagmaLower ) {
for (magma_int_t i=0; i < n-1; i++) {
d[i] = A[i*lda]; // diag
e[i] = A[i*lda+1]; // lower diag
}
d[n-1] = A[(n-1)*lda];
} else {
for (magma_int_t i=0; i < n-1; i++) {
d[i] = A[i*lda+nb]; // diag
e[i] = A[i*lda+nb-1]; // lower diag
}
d[n-1] = A[(n-1)*lda+nb];
}
#endif
return MAGMA_SUCCESS;
}
/**
Purpose
-------
DLATRD2 reduces NB rows and columns of a real symmetric matrix A to
symmetric tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = MagmaUpper, DLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, DLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by DSYTRD2_GPU. It uses an
accelerated HEMV that needs extra memory.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the upper or lower triangular part of the
symmetric matrix A is stored:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
nb INTEGER
The number of rows and columns to be reduced.
@param[in,out]
A DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
- if UPLO = MagmaUpper, the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
- if UPLO = MagmaLower, the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= (1,N).
@param[out]
e DOUBLE_PRECISION array, dimension (N-1)
If UPLO = MagmaUpper, E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = MagmaLower, E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
@param[out]
tau DOUBLE_PRECISION array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = MagmaUpper, and in TAU(1:nb) if UPLO = MagmaLower.
See Further Details.
@param[out]
W DOUBLE_PRECISION array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
@param[in]
ldw INTEGER
The leading dimension of the array W. LDW >= max(1,N).
@param
dA TODO: dimension (ldda, n) ??
@param
ldda TODO: ldda >= n ??
@param
dW TODO: dimension (lddw, 2*nb) ??
@param
lddw TODO: lddw >= n ??
@param
dwork TODO: dimension (ldwork) ??
@param
ldwork TODO: ldwork >= ceil(n/64)*ldda ??
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a symmetric rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
@ingroup magma_dsyev_aux
********************************************************************/
extern "C" magma_int_t
magma_dlatrd2(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
double *A, magma_int_t lda,
double *e, double *tau,
double *W, magma_int_t ldw,
magmaDouble_ptr dA, magma_int_t ldda,
magmaDouble_ptr dW, magma_int_t lddw,
magmaDouble_ptr dwork, magma_int_t ldwork)
{
#define A(i_, j_) (A + (i_) + (j_)*lda)
#define W(i_, j_) (W + (i_) + (j_)*ldw)
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#define dW(i_, j_) (dW + (i_) + (j_)*lddw)
const double c_neg_one = MAGMA_D_NEG_ONE;
const double c_one = MAGMA_D_ONE;
const double c_zero = MAGMA_D_ZERO;
const magma_int_t ione = 1;
double alpha, value;
magma_int_t i, i_n, i_1, iw;
/* Check arguments */
magma_int_t info = 0;
if ( uplo != MagmaLower && uplo != MagmaUpper ) {
info = -1;
} else if ( n < 0 ) {
info = -2;
} else if ( nb < 1 ) {
info = -3;
} else if ( lda < max(1,n) ) {
info = -5;
} else if ( ldw < max(1,n) ) {
info = -9;
} else if ( ldda < max(1,n) ) {
info = -11;
} else if ( lddw < max(1,n) ) {
info = -13;
} else if ( ldwork < ldda*ceildiv(n,64) ) {
info = -15;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
/* Quick return if possible */
if (n == 0) {
return info;
}
magma_queue_t stream;
magma_queue_create( &stream );
double *f;
magma_dmalloc_cpu( &f, n );
if ( f == NULL ) {
info = MAGMA_ERR_HOST_ALLOC;
return info;
}
if (uplo == MagmaUpper) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb; --i) {
i_1 = i + 1;
i_n = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i_n, W(i, iw+1), &ldw );
#endif
blasf77_dgemv( "No transpose", &i_1, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i, iw+1), &ldw, &c_one, A(0, i), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i_n, W(i, iw+1), &ldw );
lapackf77_dlacgv( &i_n, A(i, i+1), &lda );
#endif
blasf77_dgemv( "No transpose", &i_1, &i_n, &c_neg_one, W(0, iw+1), &ldw,
A(i, i+1), &lda, &c_one, A(0, i), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i_n, A(i, i+1), &lda );
#endif
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_dlarfg( &i, &alpha, A(0, i), &ione, &tau[i - 1] );
e[i-1] = MAGMA_D_REAL( alpha );
*A(i-1,i) = MAGMA_D_ONE;
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_dsetvector_async( i, A(0, i), 1, dA(0, i), 1, stream );
magmablas_dsymv_work( MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione,
dwork, ldwork, stream );
// 2. Start getting the result back (asynchronously)
magma_dgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw), ldw, stream );
if (i < n-1) {
blasf77_dgemv( MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione );
}
// 3. Here we need dsymv result W(0, iw)
magma_queue_sync( stream );
if (i < n-1) {
blasf77_dgemv( "No transpose", &i, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione );
blasf77_dgemv( MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione );
blasf77_dgemv( "No transpose", &i, &i_n, &c_neg_one, W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione );
}
blasf77_dscal( &i, &tau[i - 1], W(0, iw), &ione );
value = magma_cblas_ddot( i, W(0,iw), ione, A(0,i), ione );
alpha = tau[i - 1] * -0.5f * value;
blasf77_daxpy( &i, &alpha, A(0, i), &ione,
W(0, iw), &ione );
}
}
}
else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
i_n = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i, W(i, 0), &ldw );
#endif
blasf77_dgemv( "No transpose", &i_n, &i, &c_neg_one, A(i, 0), &lda,
W(i, 0), &ldw, &c_one, A(i, i), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i, W(i, 0), &ldw );
lapackf77_dlacgv( &i, A(i, 0), &lda );
#endif
blasf77_dgemv( "No transpose", &i_n, &i, &c_neg_one, W(i, 0), &ldw,
A(i, 0), &lda, &c_one, A(i, i), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i, A(i, 0), &lda );
#endif
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
i_n = n - i - 1;
alpha = *A(i+1, i);
lapackf77_dlarfg( &i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i] );
e[i] = MAGMA_D_REAL( alpha );
*A(i+1,i) = MAGMA_D_ONE;
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_dsetvector_async( i_n, A(i+1, i), 1, dA(i+1, i), 1, stream );
magmablas_dsymv_work( MagmaLower, i_n, c_one, dA(i+1, i+1), ldda,
dA(i+1, i), ione, c_zero, dW(i+1, i), ione,
dwork, ldwork, stream );
// 2. Start getting the result back (asynchronously)
magma_dgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_dgemv( MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione );
blasf77_dgemv( "No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione );
blasf77_dgemv( MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione );
// 3. Here we need dsymv result W(i+1, i)
magma_queue_sync( stream );
if (i != 0)
blasf77_daxpy( &i_n, &c_one, f, &ione, W(i+1, i), &ione );
blasf77_dgemv( "No transpose", &i_n, &i, &c_neg_one, W(i+1, 0), &ldw,
W(0, i), &ione, &c_one, W(i+1, i), &ione );
blasf77_dscal( &i_n, &tau[i], W(i+1,i), &ione );
value = magma_cblas_ddot( i_n, W(i+1,i), ione, A(i+1,i), ione );
alpha = tau[i] * -0.5f * value;
blasf77_daxpy( &i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione );
}
}
}
magma_free_cpu( f );
magma_queue_destroy( stream );
return info;
} /* magma_dlatrd */
extern "C" magma_int_t
magma_sgeqrf_expert_batched(
magma_int_t m, magma_int_t n,
float **dA_array, magma_int_t ldda,
float **dR_array, magma_int_t lddr,
float **dT_array, magma_int_t lddt,
float **dtau_array, magma_int_t provide_RT,
magma_int_t *info_array, magma_int_t batchCount, magma_queue_t queue)
{
#define dA(i, j) (dA + (i) + (j)*ldda) // A(i, j) means at i row, j column
/* Local Parameter */
magma_int_t nb = magma_get_sgeqrf_batched_nb(m);
magma_int_t nnb = 8;
magma_int_t min_mn = min(m, n);
/* Check arguments */
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
magma_int_t arginfo = 0;
if (m < 0)
arginfo = -1;
else if (n < 0)
arginfo = -2;
else if (ldda < max(1,m))
arginfo = -4;
else if (lddr < min_mn && provide_RT == 1)
arginfo = -6;
else if (lddr < min(min_mn, nb))
arginfo = -6;
else if (lddt < min(min_mn, nb))
arginfo = -8;
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
/* Quick return if possible */
if (m == 0 || n == 0)
if (min_mn == 0 ) return arginfo;
if ( m > 2048 || n > 2048 ) {
printf("=========================================================================================\n");
printf(" WARNING batched routines are designed for small sizes it might be better to use the\n Native/Hybrid classical routines if you want performance\n");
printf("=========================================================================================\n");
}
magma_int_t i, k, ib=nb, jb=nnb, offset_RT=0, use_stream;
magma_int_t ldw, offset;
float **dW0_displ = NULL;
float **dW1_displ = NULL;
float **dW2_displ = NULL;
float **dW3_displ = NULL;
float **dW4_displ = NULL;
float **dW5_displ = NULL;
float **dR_displ = NULL;
float **dT_displ = NULL;
float *dwork = NULL;
float **cpuAarray = NULL;
float **cpuTarray = NULL;
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dW5_displ, batchCount * sizeof(*dW5_displ));
magma_malloc((void**)&dR_displ, batchCount * sizeof(*dR_displ));
magma_malloc((void**)&dT_displ, batchCount * sizeof(*dT_displ));
magma_smalloc(&dwork, (2 * nb * n) * batchCount);
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(float*));
magma_malloc_cpu((void**) &cpuTarray, batchCount*sizeof(float*));
/* check allocation */
if ( dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL ||
dW3_displ == NULL || dW4_displ == NULL || dW5_displ == NULL ||
dR_displ == NULL || dT_displ == NULL || dwork == NULL ||
cpuAarray == NULL || cpuTarray == NULL ) {
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dW5_displ);
magma_free(dR_displ);
magma_free(dT_displ);
magma_free(dwork);
magma_free_cpu(cpuAarray);
magma_free_cpu(cpuTarray);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magma_sdisplace_pointers(dR_displ, dR_array, lddr, 0, 0, batchCount, queue);
magma_sdisplace_pointers(dT_displ, dT_array, lddt, 0, 0, batchCount, queue);
// set dR and dT to zero. if provide_RT == 0 only a tile of size nbxnb is used and overwritten at each step
magmablas_slaset_batched( MagmaFull, lddr, (provide_RT > 0 ? n:min(min_mn,nb)), MAGMA_S_ZERO, MAGMA_S_ZERO, dR_displ, lddr, batchCount, queue );
magmablas_slaset_batched( MagmaFull, lddt, (provide_RT > 0 ? n:min(min_mn,nb)), MAGMA_S_ZERO, MAGMA_S_ZERO, dT_displ, lddt, batchCount, queue );
/*
if ( provide_RT > 0 )
{
magmablas_slaset_q( MagmaFull, lddr, n*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dR, lddr, queue );
magmablas_slaset_q( MagmaFull, lddt, n*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dT, lddt, queue );
}
else
{
magmablas_slaset_q( MagmaFull, lddr, nb*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dR, lddr, queue );
magmablas_slaset_q( MagmaFull, lddt, nb*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dT, lddt, queue );
}
*/
magma_int_t streamid;
const magma_int_t nbstreams=10;
magma_queue_t queues[nbstreams];
for (i=0; i < nbstreams; i++) {
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[i] );
}
magma_getvector( batchCount, sizeof(float*), dA_array, 1, cpuAarray, 1, queue);
magma_getvector( batchCount, sizeof(float*), dT_array, 1, cpuTarray, 1, queue);
for (i=0; i < min_mn; i += nb)
{
ib = min(nb, min_mn-i);
//===============================================
// panel factorization
//===============================================
magma_sdisplace_pointers(dW0_displ, dA_array, ldda, i, i, batchCount, queue);
magma_sdisplace_pointers(dW2_displ, dtau_array, 1, i, 0, batchCount, queue);
if ( provide_RT > 0 )
{
offset_RT = i;
magma_sdisplace_pointers(dR_displ, dR_array, lddr, (provide_RT == 1 ? offset_RT:0), offset_RT, batchCount, queue);
magma_sdisplace_pointers(dT_displ, dT_array, lddt, 0, offset_RT, batchCount, queue);
}
//dwork is used in panel factorization and trailing matrix update
//dW4_displ, dW5_displ are used as workspace and configured inside
magma_sgeqrf_panel_batched(m-i, ib, jb,
dW0_displ, ldda,
dW2_displ,
dT_displ, lddt,
dR_displ, lddr,
dW1_displ,
dW3_displ,
dwork,
dW4_displ, dW5_displ,
info_array,
batchCount, queue);
//===============================================
// end of panel
//===============================================
//===============================================
// update trailing matrix
//===============================================
if ( (n-ib-i) > 0)
{
//dwork is used in panel factorization and trailing matrix update
//reset dW4_displ
ldw = nb;
magma_sset_pointer( dW4_displ, dwork, 1, 0, 0, ldw*n, batchCount, queue );
offset = ldw*n*batchCount;
magma_sset_pointer( dW5_displ, dwork + offset, 1, 0, 0, ldw*n, batchCount, queue );
// set the diagonal of v as one and the upper triangular part as zero already set inside geqrf_panel
//magmablas_slaset_batched( MagmaUpper, ib, ib, MAGMA_S_ZERO, MAGMA_S_ONE, dW0_displ, ldda, batchCount, queue );
//magma_sdisplace_pointers(dW2_displ, dtau_array, 1, i, 0, batchCount, queue);
// it is faster since it is using BLAS-3 GEMM routines, different from lapack implementation
magma_slarft_batched(m-i, ib, 0,
dW0_displ, ldda,
dW2_displ,
dT_displ, lddt,
dW4_displ, nb*lddt,
batchCount, queue);
// perform C = (I-V T^H V^H) * C, C is the trailing matrix
//-------------------------------------------
// USE STREAM GEMM
//-------------------------------------------
use_stream = magma_srecommend_cublas_gemm_stream(MagmaNoTrans, MagmaNoTrans, m-i-ib, n-i-ib, ib);
if ( use_stream )
{
magma_queue_sync(queue);
for (k=0; k < batchCount; k++)
{
streamid = k%nbstreams;
// the queue gemm must take cpu pointer
magma_slarfb_gpu_gemm( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
cpuAarray[k] + i + i * ldda, ldda,
cpuTarray[k] + offset_RT*lddt, lddt,
cpuAarray[k] + i + (i+ib) * ldda, ldda,
dwork + nb * n * k, -1,
dwork + nb * n * batchCount + nb * n * k, -1, queues[streamid] );
}
// need to synchronise to be sure that panel does not start before
// finishing the update at least of the next panel
// if queue is NULL, no need to sync
if ( queue != NULL ) {
for (magma_int_t s=0; s < nbstreams; s++)
magma_queue_sync(queues[s]);
}
}
//-------------------------------------------
// USE BATCHED GEMM
//-------------------------------------------
else
{
//direct trailing matrix in dW1_displ
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magma_slarfb_gemm_batched(
MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
(const float**)dW0_displ, ldda,
(const float**)dT_displ, lddt,
dW1_displ, ldda,
dW4_displ, ldw,
dW5_displ, ldw,
batchCount, queue );
}
}// update the trailing matrix
//===============================================
// copy dR back to V after the trailing matrix update,
// only when provide_RT=0 otherwise the nbxnb block of V is set to diag=1/0
// The upper portion of V could be set totaly to 0 here
if ( provide_RT == 0 )
{
magmablas_slacpy_batched( MagmaUpper, ib, ib, dR_displ, lddr, dW0_displ, ldda, batchCount, queue );
}
}
magma_queue_sync(queue);
for (k=0; k < nbstreams; k++) {
magma_queue_destroy( queues[k] );
}
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dW5_displ);
magma_free(dR_displ);
magma_free(dT_displ);
magma_free(dwork);
magma_free_cpu(cpuAarray);
magma_free_cpu(cpuTarray);
return arginfo;
}
static void magma_ctile_bulge_parallel(magma_int_t my_core_id, magma_int_t cores_num, magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *V, magma_int_t ldv, magmaFloatComplex *TAU, magma_int_t n, magma_int_t nb, magma_int_t nbtiles,
magma_int_t grsiz, magma_int_t Vblksiz, volatile magma_int_t *prog)
{
magma_int_t sweepid, myid, shift, stt, st, ed, stind, edind;
magma_int_t blklastind, colpt;
magma_int_t stepercol;
magma_int_t i,j,m,k;
magma_int_t thgrsiz, thgrnb, thgrid, thed;
magma_int_t coreid;
magma_int_t colblktile,maxrequiredcores,colpercore,mycoresnb;
magma_int_t fin;
magmaFloatComplex *work;
if(n<=0)
return ;
if(grsiz<=0)
return ;
//printf("=================> my core id %d of %d \n",my_core_id, cores_num);
/* As I store V in the V vector there are overlap between
* tasks so shift is now 4 where group need to be always
* multiple of 2, because as example if grs=1 task 2 from
* sweep 2 can run with task 6 sweep 1., but task 2 sweep 2
* will overwrite the V of tasks 5 sweep 1 which are used by
* task 6, so keep in mind that group need to be multiple of 2,
* and thus tasks 2 sweep 2 will never run with task 6 sweep 1.
* However, when storing V in A, shift could be back to 3.
* */
magma_cmalloc_cpu(&work, n);
mycoresnb = cores_num;
shift = 5;
if(grsiz==1)
colblktile=1;
else
colblktile=grsiz/2;
maxrequiredcores = nbtiles/colblktile;
if(maxrequiredcores<1)maxrequiredcores=1;
colpercore = colblktile*nb;
if(mycoresnb > maxrequiredcores)
mycoresnb = maxrequiredcores;
thgrsiz = n;
stepercol = magma_ceildiv(shift, grsiz);
thgrnb = magma_ceildiv(n-1, thgrsiz);
#ifdef ENABLE_DEBUG
if(my_core_id==0){
if(cores_num > maxrequiredcores) {
printf("==================================================================================\n");
printf(" WARNING only %3d threads are required to run this test optimizing cache reuse\n",maxrequiredcores);
printf("==================================================================================\n");
}
printf(" Static bulgechasing version v9_9col threads %4d N %5d NB %5d grs %4d thgrsiz %4d \n",cores_num, n, nb, grsiz,thgrsiz);
}
#endif
for (thgrid = 1; thgrid<=thgrnb; thgrid++){
stt = (thgrid-1)*thgrsiz+1;
thed = min( (stt + thgrsiz -1), (n-1));
for (i = stt; i <= n-1; i++){
ed=min(i,thed);
if(stt>ed)break;
for (m = 1; m <=stepercol; m++){
st=stt;
for (sweepid = st; sweepid <=ed; sweepid++){
for (k = 1; k <=grsiz; k++){
myid = (i-sweepid)*(stepercol*grsiz) +(m-1)*grsiz + k;
if(myid%2 ==0){
colpt = (myid/2)*nb+1+sweepid-1;
stind = colpt-nb+1;
edind = min(colpt,n);
blklastind = colpt;
if(stind>=edind){
printf("ERROR---------> st>=ed %d %d \n\n", (int) stind, (int) edind);
exit(-10);
}
}else{
colpt = ((myid+1)/2)*nb + 1 +sweepid -1 ;
stind = colpt-nb+1;
edind = min(colpt,n);
if( (stind>=edind-1) && (edind==n) )
blklastind=n;
else
blklastind=0;
if(stind>edind){
printf("ERROR---------> st>=ed %d %d \n\n", (int) stind, (int) edind);
exit(-10);
}
}
coreid = (stind/colpercore)%mycoresnb;
if(my_core_id==coreid)
{
fin=0;
while(fin==0)
{
if(myid==1)
{
if( (prog[myid+shift-1]== (sweepid-1)) )
{
magma_ctrdtype1cbHLsym_withQ_v2(n, nb, A, lda, V, ldv, TAU, stind, edind, sweepid, Vblksiz, work);
fin=1;
prog[myid]= sweepid;
if(blklastind >= (n-1))
{
for (j = 1; j <= shift; j++)
prog[myid+j]=sweepid;
}
} // END progress condition
}else{
if( (prog[myid-1]==sweepid) && (prog[myid+shift-1]== (sweepid-1)) )
{
if(myid%2 == 0)
magma_ctrdtype2cbHLsym_withQ_v2(n, nb, A, lda, V, ldv, TAU, stind, edind, sweepid, Vblksiz, work);
else
magma_ctrdtype3cbHLsym_withQ_v2(n, nb, A, lda, V, ldv, TAU, stind, edind, sweepid, Vblksiz, work);
fin=1;
prog[myid]= sweepid;
if(blklastind >= (n-1))
{
for (j = 1; j <= shift+mycoresnb; j++)
prog[myid+j]=sweepid;
}
} // END progress condition
} // END if myid==1
} // END while loop
} // END if my_core_id==coreid
if(blklastind >= (n-1))
{
stt=stt+1;
break;
}
} // END for k=1:grsiz
} // END for sweepid=st:ed
} // END for m=1:stepercol
} // END for i=1:n-1
} // END for thgrid=1:thgrnb
magma_free_cpu(work);
} // END FUNCTION
/**
Purpose
-------
SGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
d_lA REAL array of pointers on the GPU, dimension (ngpu).
On entry, the M-by-N matrix A distributed over GPUs
(d_lA[d] points to the local matrix on d-th GPU).
It uses 1D block column cyclic format with the block size of nb,
and each local matrix is stored by column.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array d_lA. LDDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_sgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_sgetrf_mgpu(
magma_int_t ngpu,
magma_int_t m, magma_int_t n,
magmaFloat_ptr d_lA[], magma_int_t ldda, magma_int_t *ipiv,
magma_int_t *info)
{
magma_int_t nb, n_local[MagmaMaxGPUs];
magma_int_t maxm;
magma_int_t i, j, d, lddat, lddwork;
float *d_lAT[MagmaMaxGPUs];
float *d_panel[MagmaMaxGPUs], *work;
magma_queue_t queues[MagmaMaxGPUs][2];
/* Check arguments */
*info = 0;
if (m < 0)
*info = -2;
else if (n < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* create the queues */
for( d=0; d < ngpu; d++ ) {
magma_queue_create( d, &queues[d][0] );
magma_queue_create( d, &queues[d][1] );
}
/* Function Body */
nb = magma_get_sgetrf_nb( m, n );
if (nb <= 1 || nb >= n) {
/* Use CPU code. */
magma_smalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_sgetmatrix( m, n, d_lA[0], ldda, work, m, queues[0][0] );
lapackf77_sgetrf(&m, &n, work, &m, ipiv, info);
magma_ssetmatrix( m, n, work, m, d_lA[0], ldda, queues[0][0] );
magma_free_cpu(work);
} else {
/* Use hybrid blocked code. */
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
maxm = magma_roundup( m, 32 );
if ( ngpu > ceil((float)n/nb) ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) ngpu );
*info = -1;
return *info;
}
/* allocate workspace for each GPU */
lddat = magma_roundup( ((magma_ceildiv( n, nb )/ngpu)*nb), 32 );
lddat = magma_ceildiv( n, nb ); /* number of block columns */
lddat = magma_ceildiv( lddat, ngpu ); /* number of block columns per GPU */
lddat = nb*lddat; /* number of columns per GPU */
lddat = magma_roundup( lddat, 32 ); /* make it a multiple of 32 */
for (i=0; i < ngpu; i++) {
magma_setdevice(i);
/* local-n and local-ld */
n_local[i] = ((n/nb)/ngpu)*nb;
if (i < (n/nb)%ngpu)
n_local[i] += nb;
else if (i == (n/nb)%ngpu)
n_local[i] += n%nb;
/* workspaces */
if (MAGMA_SUCCESS != magma_smalloc( &d_panel[i], (3+ngpu)*nb*maxm )) {
for( j=0; j <= i; j++ ) {
magma_setdevice(j);
}
for( j=0; j < i; j++ ) {
magma_setdevice(j);
magma_free( d_panel[j] );
magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* local-matrix storage */
if (MAGMA_SUCCESS != magma_smalloc( &d_lAT[i], lddat*maxm )) {
for( j=0; j <= i; j++ ) {
magma_setdevice(j);
magma_free( d_panel[j] );
}
for( j=0; j < i; j++ ) {
magma_setdevice(j);
magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_stranspose( m, n_local[i], d_lA[i], ldda, d_lAT[i], lddat, queues[i][1] );
}
for (i=0; i < ngpu; i++) {
magma_setdevice(i);
magma_queue_sync(queues[i][0]);
}
magma_setdevice(0);
/* cpu workspace */
lddwork = maxm;
if (MAGMA_SUCCESS != magma_smalloc_pinned( &work, lddwork*nb*ngpu )) {
for (i=0; i < ngpu; i++ ) {
magma_setdevice(i);
magma_free( d_panel[i] );
magma_free( d_lAT[i] );
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* calling multi-gpu interface with allocated workspaces and queues */
magma_sgetrf2_mgpu(ngpu, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
queues, info);
/* clean up */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
/* save on output */
magmablas_stranspose( n_local[d], m, d_lAT[d], lddat, d_lA[d], ldda, queues[d][0] );
magma_queue_sync(queues[d][0]);
magma_queue_sync(queues[d][1]);
magma_free( d_lAT[d] );
magma_free( d_panel[d] );
} /* end of for d=1,..,ngpu */
magma_setdevice( orig_dev );
magma_free_pinned( work );
}
/* clean up */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_destroy( queues[d][0] );
magma_queue_destroy( queues[d][1] );
}
return *info;
}
extern "C" magma_int_t
magma_zlatrd2(char uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *a, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *w, magma_int_t ldw,
magmaDoubleComplex *da, magma_int_t ldda,
magmaDoubleComplex *dw, magma_int_t lddw,
magmaDoubleComplex *dwork, magma_int_t ldwork)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZLATRD2 reduces NB rows and columns of a complex Hermitian matrix A to
Hermitian tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = 'U', ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = 'L', ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD2_GPU. It uses an
accelerated HEMV that needs extra memory.
Arguments
=========
UPLO (input) CHARACTER*1
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
= 'U': Upper triangular
= 'L': Lower triangular
N (input) INTEGER
The order of the matrix A.
NB (input) INTEGER
The number of rows and columns to be reduced.
A (input/output) COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
if UPLO = 'U', the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
if UPLO = 'L', the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= (1,N).
E (output) COMPLEX_16 array, dimension (N-1)
If UPLO = 'U', E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = 'L', E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
TAU (output) COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = 'U', and in TAU(1:nb) if UPLO = 'L'.
See Further Details.
W (output) COMPLEX_16 array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
LDW (input) INTEGER
The leading dimension of the array W. LDW >= max(1,N).
Further Details
===============
If UPLO = 'U', the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = 'L', the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a Hermitian rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = 'U': if UPLO = 'L':
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
===================================================================== */
char uplo_[2] = {uplo, 0};
magma_int_t i;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex value = MAGMA_Z_ZERO;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
magmaDoubleComplex alpha;
magmaDoubleComplex *f;
if (n <= 0) {
return 0;
}
magma_queue_t stream;
magma_queue_create( &stream );
magma_zmalloc_cpu( &f, n );
assert( f != NULL ); // TODO return error, or allocate outside zlatrd
if (lapackf77_lsame(uplo_, "U")) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb ; --i) {
i_1 = i + 1;
i_n = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_n, W(i, iw+1), &ldw);
#endif
blasf77_zgemv("No transpose", &i_1, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i, iw+1), &ldw, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_n, W(i, iw+1), &ldw);
lapackf77_zlacgv(&i_n, A(i, i+1), &ldw);
#endif
blasf77_zgemv("No transpose", &i_1, &i_n, &c_neg_one, W(0, iw+1), &ldw,
A(i, i+1), &lda, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_n, A(i, i+1), &ldw);
#endif
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_zlarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_Z_REAL( alpha );
MAGMA_Z_SET2REAL(*A(i-1, i), 1.);
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_zsetvector( i, A(0, i), 1, dA(0, i), 1 );
#if (GPUSHMEM < 200)
magma_zhemv(MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione);
#else
magmablas_zhemv2(MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione,
dwork, ldwork);
#endif
// 2. Start putting the result back (asynchronously)
magma_zgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw) /*test*/, ldw, stream );
if (i < n-1) {
blasf77_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
}
// 3. Here is where we need it // TODO find the right place
magma_queue_sync( stream );
if (i < n-1) {
blasf77_zgemv("No transpose", &i, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
blasf77_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_zgemv("No transpose", &i, &i_n, &c_neg_one, W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
}
blasf77_zscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -0.5f * value;
blasf77_zaxpy(&i, &alpha, A(0, i), &ione,
W(0, iw), &ione);
}
}
}
else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
i_n = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, W(i, 0), &ldw);
#endif
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(i, 0), &lda,
W(i, 0), &ldw, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, W(i, 0), &ldw);
lapackf77_zlacgv(&i, A(i ,0), &lda);
#endif
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, W(i, 0), &ldw,
A(i, 0), &lda, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, A(i, 0), &lda);
#endif
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
i_n = n - i - 1;
alpha = *A(i+1, i);
lapackf77_zlarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
e[i] = MAGMA_Z_REAL( alpha );
MAGMA_Z_SET2REAL(*A(i+1, i), 1.);
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_zsetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1 );
#if (GPUSHMEM < 200)
magma_zhemv(MagmaLower, i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
dW(i+1, i), ione);
#else
magmablas_zhemv2('L', i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
dW(i+1, i), ione,
dwork, ldwork);
#endif
// 2. Start putting the result back (asynchronously)
magma_zgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
// 3. Here is where we need it
magma_queue_sync( stream );
if (i!=0)
blasf77_zaxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, W(i+1, 0), &ldw,
W(0, i), &ione, &c_one, W(i+1, i), &ione);
blasf77_zscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_zdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i] * -0.5f * value;
blasf77_zaxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
}
}
}
magma_free_cpu(f);
magma_queue_destroy( stream );
return 0;
} /* zlatrd */
/***************************************************************************//**
Purpose
-------
DGELS solves the overdetermined, least squares problem
min || A*X - C ||
using the QR factorization A.
The underdetermined problem (m < n) is not currently handled.
Arguments
---------
@param[in]
trans magma_trans_t
- = MagmaNoTrans: the linear system involves A.
Only TRANS=MagmaNoTrans is currently handled.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. M >= N >= 0.
@param[in]
nrhs INTEGER
The number of columns of the matrix C. NRHS >= 0.
@param[in,out]
A DOUBLE PRECISION array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, A is overwritten by details of its QR
factorization as returned by DGEQRF.
@param[in]
lda INTEGER
The leading dimension of the array A, LDA >= M.
@param[in,out]
B DOUBLE PRECISION array, dimension (LDDB,NRHS)
On entry, the M-by-NRHS matrix C.
On exit, the N-by-NRHS solution matrix X.
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= M.
@param[out]
hwork (workspace) DOUBLE PRECISION array, dimension MAX(1,LWORK).
On exit, if INFO = 0, HWORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array HWORK,
LWORK >= max( N*NB, 2*NB*NB ),
where NB is the blocksize given by magma_get_dgeqrf_nb( M, N ).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the HWORK array, returns
this value as the first entry of the HWORK array.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_gels
*******************************************************************************/
extern "C" magma_int_t
magma_dgels(
magma_trans_t trans, magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaDouble_ptr A, magma_int_t lda,
magmaDouble_ptr B, magma_int_t ldb,
double *hwork, magma_int_t lwork,
magma_int_t *info)
{
/* Constants */
const double c_one = MAGMA_D_ONE;
/* Local variables */
double *tau;
magma_int_t min_mn;
magma_int_t nb = magma_get_dgeqrf_nb( m, n );
magma_int_t lwkopt = max( n*nb, 2*nb*nb ); // (m - n + nb)*(nrhs + nb) + nrhs*nb;
bool lquery = (lwork == -1);
hwork[0] = magma_dmake_lwork( lwkopt );
*info = 0;
/* For now, N is the only case working */
if ( trans != MagmaNoTrans )
*info = -1;
else if (m < 0)
*info = -2;
else if (n < 0 || m < n) /* LQ is not handle for now */
*info = -3;
else if (nrhs < 0)
*info = -4;
else if (lda < max(1,m))
*info = -6;
else if (ldb < max(1,m))
*info = -8;
else if (lwork < lwkopt && ! lquery)
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
min_mn = min(m,n);
if (min_mn == 0) {
hwork[0] = c_one;
return *info;
}
magma_dmalloc_cpu( &tau, min_mn );
if ( tau == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_dgeqrf( m, n, A, lda, tau, hwork, lwork, info );
if ( *info == 0 ) {
// B := Q' * B
lapackf77_dormqr( MagmaLeftStr, MagmaTransStr, &m, &nrhs, &n,
A, &lda, tau, B, &ldb, hwork, &lwork, info );
// Solve R*X = B(1:n,:)
blasf77_dtrsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&n, &nrhs, &c_one, A, &lda, B, &ldb );
}
magma_free_cpu( tau );
return *info;
}
/**
Purpose
-------
DLABRD reduces the first NB rows and columns of a real general
m by n matrix A to upper or lower bidiagonal form by an orthogonal
transformation Q' * A * P, and returns the matrices X and Y which
are needed to apply the transformation to the unreduced part of A.
If m >= n, A is reduced to upper bidiagonal form; if m < n, to lower
bidiagonal form.
This is an auxiliary routine called by DGEBRD.
Arguments
---------
@param[in]
m INTEGER
The number of rows in the matrix A.
@param[in]
n INTEGER
The number of columns in the matrix A.
@param[in]
nb INTEGER
The number of leading rows and columns of A to be reduced.
@param[in,out]
A DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the m by n general matrix to be reduced.
On exit, the first NB rows and columns of the matrix are
overwritten; the rest of the array is unchanged.
If m >= n, elements on and below the diagonal in the first NB
columns, with the array TAUQ, represent the orthogonal
matrix Q as a product of elementary reflectors; and
elements above the diagonal in the first NB rows, with the
array TAUP, represent the orthogonal matrix P as a product
of elementary reflectors.
\n
If m < n, elements below the diagonal in the first NB
columns, with the array TAUQ, represent the orthogonal
matrix Q as a product of elementary reflectors, and
elements on and above the diagonal in the first NB rows,
with the array TAUP, represent the orthogonal matrix P as
a product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[in,out]
dA DOUBLE_PRECISION array, dimension (LDDA,N)
Copy of A on GPU.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
@param[out]
d DOUBLE_PRECISION array, dimension (NB)
The diagonal elements of the first NB rows and columns of
the reduced matrix. D(i) = A(i,i).
@param[out]
e DOUBLE_PRECISION array, dimension (NB)
The off-diagonal elements of the first NB rows and columns of
the reduced matrix.
@param[out]
tauq DOUBLE_PRECISION array dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix Q. See Further Details.
@param[out]
taup DOUBLE_PRECISION array, dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix P. See Further Details.
@param[out]
X DOUBLE_PRECISION array, dimension (LDX,NB)
The m-by-nb matrix X required to update the unreduced part
of A.
@param[in]
ldx INTEGER
The leading dimension of the array X. LDX >= M.
@param[out]
dX DOUBLE_PRECISION array, dimension (LDDX,NB)
Copy of X on GPU.
@param[in]
lddx INTEGER
The leading dimension of the array dX. LDDX >= M.
@param[out]
Y DOUBLE_PRECISION array, dimension (LDY,NB)
The n-by-nb matrix Y required to update the unreduced part
of A.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[out]
dY DOUBLE_PRECISION array, dimension (LDDY,NB)
Copy of Y on GPU.
@param[in]
lddy INTEGER
The leading dimension of the array dY. LDDY >= N.
Further Details
---------------
The matrices Q and P are represented as products of elementary
reflectors:
Q = H(1) H(2) . . . H(nb) and P = G(1) G(2) . . . G(nb)
Each H(i) and G(i) has the form:
H(i) = I - tauq * v * v' and G(i) = I - taup * u * u'
where tauq and taup are real scalars, and v and u are real vectors.
If m >= n, v(1:i-1) = 0, v(i) = 1, and v(i:m) is stored on exit in
A(i:m,i); u(1:i) = 0, u(i+1) = 1, and u(i+1:n) is stored on exit in
A(i,i+1:n); tauq is stored in TAUQ(i) and taup in TAUP(i).
If m < n, v(1:i) = 0, v(i+1) = 1, and v(i+1:m) is stored on exit in
A(i+2:m,i); u(1:i-1) = 0, u(i) = 1, and u(i:n) is stored on exit in
A(i,i+1:n); tauq is stored in TAUQ(i) and taup in TAUP(i).
The elements of the vectors v and u together form the m-by-nb matrix
V and the nb-by-n matrix U' which are needed, with X and Y, to apply
the transformation to the unreduced part of the matrix, using a block
update of the form: A := A - V*Y' - X*U'.
The contents of A on exit are illustrated by the following examples
with nb = 2:
@verbatim
m = 6 and n = 5 (m > n): m = 5 and n = 6 (m < n):
( 1 1 u1 u1 u1 ) ( 1 u1 u1 u1 u1 u1 )
( v1 1 1 u2 u2 ) ( 1 1 u2 u2 u2 u2 )
( v1 v2 a a a ) ( v1 1 a a a a )
( v1 v2 a a a ) ( v1 v2 a a a a )
( v1 v2 a a a ) ( v1 v2 a a a a )
( v1 v2 a a a )
@endverbatim
where a denotes an element of the original matrix which is unchanged,
vi denotes an element of the vector defining H(i), and ui an element
of the vector defining G(i).
@ingroup magma_dgesvd_aux
********************************************************************/
extern "C" magma_int_t
magma_dlabrd_gpu( magma_int_t m, magma_int_t n, magma_int_t nb,
double *A, magma_int_t lda,
double *dA, magma_int_t ldda,
double *d, double *e, double *tauq, double *taup,
double *X, magma_int_t ldx,
double *dX, magma_int_t lddx,
double *Y, magma_int_t ldy,
double *dY, magma_int_t lddy)
{
#define A(i_,j_) (A + (i_) + (j_)*lda)
#define X(i_,j_) (X + (i_) + (j_)*ldx)
#define Y(i_,j_) (Y + (i_) + (j_)*ldy)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dY(i_,j_) (dY + (i_) + (j_)*lddy)
#define dX(i_,j_) (dX + (i_) + (j_)*lddx)
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
double c_zero = MAGMA_D_ZERO;
magma_int_t ione = 1;
magma_int_t i__2, i__3;
magma_int_t i;
double alpha;
A -= 1 + lda;
X -= 1 + ldx;
dX -= 1 + lddx;
Y -= 1 + ldy;
dY -= 1 + lddy;
--d;
--e;
--tauq;
--taup;
/* Quick return if possible */
if (m <= 0 || n <= 0) {
return 0;
}
double *f;
magma_queue_t stream;
magma_queue_create( &stream );
magma_dmalloc_cpu( &f, max(n,m) );
assert( f != NULL ); // TODO return error, or allocate outside dlatrd
if (m >= n) {
/* Reduce to upper bidiagonal form */
for (i = 1; i <= nb; ++i) {
/* Update A(i:m,i) */
i__2 = m - i + 1;
i__3 = i - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
A(i,1), &lda,
Y(i,1), &ldy, &c_one,
A(i,i), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
X(i,1), &ldx,
A(1,i), &ione, &c_one,
A(i,i), &ione );
/* Generate reflection Q(i) to annihilate A(i+1:m,i) */
alpha = *A(i,i);
i__2 = m - i + 1;
i__3 = i + 1;
lapackf77_dlarfg( &i__2, &alpha, A(min(i__3,m),i), &ione, &tauq[i] );
d[i] = MAGMA_D_REAL( alpha );
if (i < n) {
*A(i,i) = c_one;
/* Compute Y(i+1:n,i) */
i__2 = m - i + 1;
i__3 = n - i;
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_dsetvector( i__2,
A(i,i), 1,
dA(i-1,i-1), 1 );
// 2. Multiply ---------------------------------------------
magma_dgemv( MagmaTrans, i__2, i__3, c_one,
dA(i-1,i), ldda,
dA(i-1,i-1), ione, c_zero,
dY(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__3, 1,
dY(i+1,i), lddy,
Y(i+1,i), ldy, stream );
i__2 = m - i + 1;
i__3 = i - 1;
blasf77_dgemv( MagmaTransStr, &i__2, &i__3, &c_one,
A(i,1), &lda,
A(i,i), &ione, &c_zero,
Y(1,i), &ione );
i__2 = n - i;
i__3 = i - 1;
blasf77_dgemv( "N", &i__2, &i__3, &c_neg_one,
Y(i+1,1), &ldy,
Y(1,i), &ione, &c_zero,
f, &ione );
i__2 = m - i + 1;
i__3 = i - 1;
blasf77_dgemv( MagmaTransStr, &i__2, &i__3, &c_one,
X(i,1), &ldx,
A(i,i), &ione, &c_zero,
Y(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i__3 != 0) {
i__2 = n - i;
blasf77_daxpy( &i__2, &c_one, f, &ione, Y(i+1,i), &ione );
}
i__2 = i - 1;
i__3 = n - i;
blasf77_dgemv( MagmaTransStr, &i__2, &i__3, &c_neg_one,
A(1,i+1), &lda,
Y(1,i), &ione, &c_one,
Y(i+1,i), &ione );
i__2 = n - i;
blasf77_dscal( &i__2, &tauq[i], Y(i+1,i), &ione );
/* Update A(i,i+1:n) */
i__2 = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__2, A(i,i+1), &lda );
lapackf77_dlacgv( &i, A(i,1), &lda );
#endif
blasf77_dgemv( "No transpose", &i__2, &i, &c_neg_one,
Y(i+1,1), &ldy,
A(i,1), &lda, &c_one,
A(i,i+1), &lda );
i__2 = i - 1;
i__3 = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i, A(i,1), &lda );
lapackf77_dlacgv( &i__2, X(i,1), &ldx );
#endif
blasf77_dgemv( MagmaTransStr, &i__2, &i__3, &c_neg_one,
A(1,i+1), &lda,
X(i,1), &ldx, &c_one,
A(i,i+1), &lda );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__2, X(i,1), &ldx );
#endif
/* Generate reflection P(i) to annihilate A(i,i+2:n) */
i__2 = n - i;
i__3 = i + 2;
alpha = *A(i,i+1);
lapackf77_dlarfg( &i__2, &alpha, A(i,min(i__3,n)), &lda, &taup[i] );
e[i] = MAGMA_D_REAL( alpha );
*A(i,i+1) = c_one;
/* Compute X(i+1:m,i) */
i__2 = m - i;
i__3 = n - i;
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_dsetvector( i__3,
A(i,i+1), lda,
dA(i-1,i), ldda );
// 2. Multiply ---------------------------------------------
//magma_dcopy( i__3, dA(i-1,i), ldda, dY(1,1), 1 );
magma_dgemv( MagmaNoTrans, i__2, i__3, c_one,
dA(i,i), ldda,
dA(i-1,i), ldda,
//dY(1,1), 1,
c_zero,
dX(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__2, 1,
dX(i+1,i), lddx,
X(i+1,i), ldx, stream );
i__2 = n - i;
blasf77_dgemv( MagmaTransStr, &i__2, &i, &c_one,
Y(i+1,1), &ldy,
A(i,i+1), &lda, &c_zero,
X(1,i), &ione );
i__2 = m - i;
blasf77_dgemv( "N", &i__2, &i, &c_neg_one,
A(i+1,1), &lda,
X(1,i), &ione, &c_zero,
f, &ione );
i__2 = i - 1;
i__3 = n - i;
blasf77_dgemv( "N", &i__2, &i__3, &c_one,
A(1,i+1), &lda,
A(i,i+1), &lda, &c_zero,
X(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i != 0) {
i__2 = m - i;
blasf77_daxpy( &i__2, &c_one, f, &ione, X(i+1,i), &ione );
}
i__2 = m - i;
i__3 = i - 1;
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
X(i+1,1), &ldx,
X(1,i), &ione, &c_one,
X(i+1,i), &ione );
i__2 = m - i;
blasf77_dscal( &i__2, &taup[i], X(i+1,i), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
i__2 = n - i;
lapackf77_dlacgv( &i__2, A(i,i+1), &lda );
// 4. Send the block reflector A(i+1:m,i) to the GPU after DLACGV()
magma_dsetvector( i__2,
A(i,i+1), lda,
dA(i-1,i), ldda );
#endif
}
}
}
else {
/* Reduce to lower bidiagonal form */
for (i = 1; i <= nb; ++i) {
/* Update A(i,i:n) */
i__2 = n - i + 1;
i__3 = i - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__2, A(i,i), &lda );
lapackf77_dlacgv( &i__3, A(i,1), &lda );
#endif
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
Y(i,1), &ldy,
A(i,1), &lda, &c_one,
A(i,i), &lda );
i__2 = i - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__3, A(i,1), &lda );
lapackf77_dlacgv( &i__3, X(i,1), &ldx );
#endif
i__3 = n - i + 1;
blasf77_dgemv( MagmaTransStr, &i__2, &i__3, &c_neg_one,
A(1,i), &lda,
X(i,1), &ldx, &c_one,
A(i,i), &lda );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__2, X(i,1), &ldx );
#endif
/* Generate reflection P(i) to annihilate A(i,i+1:n) */
i__2 = n - i + 1;
i__3 = i + 1;
alpha = *A(i,i);
lapackf77_dlarfg( &i__2, &alpha, A(i,min(i__3,n)), &lda, &taup[i] );
d[i] = MAGMA_D_REAL( alpha );
if (i < m) {
*A(i,i) = c_one;
/* Compute X(i+1:m,i) */
i__2 = m - i;
i__3 = n - i + 1;
// 1. Send the block reflector A(i,i+1:n) to the GPU ------
magma_dsetvector( i__3,
A(i,i), lda,
dA(i-1,i-1), ldda );
// 2. Multiply ---------------------------------------------
//magma_dcopy( i__3, dA(i-1,i-1), ldda, dY(1,1), 1 );
magma_dgemv( MagmaNoTrans, i__2, i__3, c_one,
dA(i,i-1), ldda,
dA(i-1,i-1), ldda,
//dY(1,1), 1,
c_zero,
dX(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__2, 1,
dX(i+1,i), lddx,
X(i+1,i), ldx, stream );
i__2 = n - i + 1;
i__3 = i - 1;
blasf77_dgemv( MagmaTransStr, &i__2, &i__3, &c_one,
Y(i,1), &ldy,
A(i,i), &lda, &c_zero,
X(1,i), &ione );
i__2 = m - i;
i__3 = i - 1;
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
A(i+1,1), &lda,
X(1,i), &ione, &c_zero,
f, &ione );
i__2 = i - 1;
i__3 = n - i + 1;
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_one,
A(1,i), &lda,
A(i,i), &lda, &c_zero,
X(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i__2 != 0) {
i__3 = m - i;
blasf77_daxpy( &i__3, &c_one, f, &ione, X(i+1,i), &ione );
}
i__2 = m - i;
i__3 = i - 1;
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
X(i+1,1), &ldx,
X(1,i), &ione, &c_one,
X(i+1,i), &ione );
i__2 = m - i;
blasf77_dscal( &i__2, &taup[i], X(i+1,i), &ione );
i__2 = n - i + 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__2, A(i,i), &lda );
magma_dsetvector( i__2,
A(i,i), lda,
dA(i-1,i-1), ldda );
#endif
/* Update A(i+1:m,i) */
i__2 = m - i;
i__3 = i - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
A(i+1,1), &lda,
Y(i,1), &ldy, &c_one,
A(i+1,i), &ione );
i__2 = m - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_dgemv( "No transpose", &i__2, &i, &c_neg_one,
X(i+1,1), &ldx,
A(1,i), &ione, &c_one,
A(i+1,i), &ione );
/* Generate reflection Q(i) to annihilate A(i+2:m,i) */
i__2 = m - i;
i__3 = i + 2;
alpha = *A(i+1,i);
lapackf77_dlarfg( &i__2, &alpha, A(min(i__3,m),i), &ione, &tauq[i] );
e[i] = MAGMA_D_REAL( alpha );
*A(i+1,i) = c_one;
/* Compute Y(i+1:n,i) */
i__2 = m - i;
i__3 = n - i;
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_dsetvector( i__2,
A(i+1,i), 1,
dA(i,i-1), 1 );
// 2. Multiply ---------------------------------------------
magma_dgemv( MagmaTrans, i__2, i__3, c_one,
dA(i,i), ldda,
dA(i,i-1), ione, c_zero,
dY(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__3, 1,
dY(i+1,i), lddy,
Y(i+1,i), ldy, stream );
i__2 = m - i;
i__3 = i - 1;
blasf77_dgemv( MagmaTransStr, &i__2, &i__3, &c_one,
A(i+1,1), &lda,
A(i+1,i), &ione, &c_zero,
Y(1,i), &ione );
i__2 = n - i;
i__3 = i - 1;
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
Y(i+1,1), &ldy,
Y(1,i), &ione, &c_zero,
f, &ione );
i__2 = m - i;
blasf77_dgemv( MagmaTransStr, &i__2, &i, &c_one,
X(i+1,1), &ldx,
A(i+1,i), &ione, &c_zero,
Y(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i__3 != 0) {
i__2 = n - i;
blasf77_daxpy( &i__2, &c_one, f, &ione, Y(i+1,i), &ione );
}
i__2 = n - i;
blasf77_dgemv( MagmaTransStr, &i, &i__2, &c_neg_one,
A(1,i+1), &lda,
Y(1,i), &ione, &c_one,
Y(i+1,i), &ione );
i__2 = n - i;
blasf77_dscal( &i__2, &tauq[i], Y(i+1,i), &ione );
}
#if defined(PRECISION_z) || defined(PRECISION_c)
else {
i__2 = n - i + 1;
lapackf77_dlacgv( &i__2, A(i,i), &lda );
magma_dsetvector( i__2,
A(i,i), lda,
dA(i-1,i-1), ldda );
}
#endif
}
}
magma_queue_destroy( stream );
magma_free_cpu( f );
return MAGMA_SUCCESS;
} /* magma_dlabrd_gpu */
extern "C" magma_err_t
magma_zgetrs_gpu(magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magma_int_t *ipiv,
magmaDoubleComplex_ptr dB, size_t dB_offset, magma_int_t lddb,
magma_int_t *info, magma_queue_t queue)
{
/* -- clMagma (version 0.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date January 2014
Purpose
=======
Solves a system of linear equations
A * X = B or A' * X = B
with a general N-by-N matrix A using the LU factorization computed by ZGETRF_GPU.
Arguments
=========
TRANS (input) CHARACTER*1
Specifies the form of the system of equations:
= 'N': A * X = B (No transpose)
= 'T': A'* X = B (Transpose)
= 'C': A'* X = B (Conjugate transpose = Transpose)
N (input) INTEGER
The order of the matrix A. N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
A (input) COMPLEX_16 array on the GPU, dimension (LDA,N)
The factors L and U from the factorization A = P*L*U as computed
by ZGETRF_GPU.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
IPIV (input) INTEGER array, dimension (N)
The pivot indices from ZGETRF; for 1<=i<=N, row i of the
matrix was interchanged with row IPIV(i).
B (input/output) COMPLEX_16 array on the GPU, dimension (LDB,NRHS)
On entry, the right hand side matrix B.
On exit, the solution matrix X.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
HWORK (workspace) COMPLEX_16 array, dimension N*NRHS
===================================================================== */
magmaDoubleComplex z_one = MAGMA_Z_MAKE( 1.0, 0.0 );
magmaDoubleComplex *work = NULL;
magma_trans_t trans_ = trans;
long int notran = lapackf77_lsame(lapack_const(trans_), lapack_const(MagmaNoTrans));
magma_int_t i1, i2, inc;
*info = 0;
if ( (! notran) &&
(! lapackf77_lsame(lapack_const(trans_), lapack_const(MagmaTrans))) &&
(! lapackf77_lsame(lapack_const(trans_), lapack_const(MagmaConjTrans))) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (nrhs < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
} else if (lddb < max(1,n)) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (n == 0 || nrhs == 0) {
return *info;
}
magma_zmalloc_cpu( &work, n*nrhs );
if ( !work ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
i1 = 1;
i2 = n;
if (notran) {
inc = 1;
/* Solve A * X = B. */
magma_zgetmatrix( n, nrhs, dB, dB_offset, lddb, work, 0, n, queue );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, 0, n, dB, dB_offset, lddb, queue );
if ( nrhs == 1) {
chk(magma_ztrsv(MagmaLower, MagmaNoTrans, MagmaUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue));
chk(magma_ztrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue));
} else {
chk(magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
chk(magma_ztrsm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
}
} else {
inc = -1;
/* Solve A' * X = B. */
if ( nrhs == 1) {
chk(magma_ztrsv(MagmaUpper, trans, MagmaNonUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue ));
chk(magma_ztrsv(MagmaLower, trans, MagmaUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue ));
} else {
chk(magma_ztrsm(MagmaLeft, MagmaUpper, trans, MagmaNonUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
chk(magma_ztrsm(MagmaLeft, MagmaLower, trans, MagmaUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
}
magma_zgetmatrix( n, nrhs, dB, dB_offset, lddb, work, 0, n, queue );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, 0, n, dB, dB_offset, lddb, queue );
}
magma_free_cpu(work);
return *info;
}
extern "C" magma_int_t
magma_cbulge_back(
magma_uplo_t uplo,
magma_int_t n, magma_int_t nb,
magma_int_t ne, magma_int_t Vblksiz,
magmaFloatComplex *Z, magma_int_t ldz,
magmaFloatComplex_ptr dZ, magma_int_t lddz,
magmaFloatComplex *V, magma_int_t ldv,
magmaFloatComplex *TAU,
magmaFloatComplex *T, magma_int_t ldt,
magma_int_t* info)
{
magma_int_t threads = magma_get_parallel_numthreads();
magma_int_t mklth = magma_get_lapack_numthreads();
magma_set_lapack_numthreads(1);
real_Double_t timeaplQ2=0.0;
float f= 1.;
magma_int_t n_gpu = ne;
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
//#ifdef REAL
//float gpu_cpu_perf = 50; // gpu over cpu performance //100% ev // SandyB. - Kepler (K20c)
//float gpu_cpu_perf = 16; // gpu over cpu performance //100% ev // SandyB. - Fermi (M2090)
//#else
// float gpu_cpu_perf = 27.5; // gpu over cpu performance //100% ev // Westmere - Fermi (M2090)
//float gpu_cpu_perf = 37; // gpu over cpu performance //100% ev // SandyB. - Kepler (K20c)
// float gpu_cpu_perf = 130; // gpu over cpu performance //100% ev // Bulldozer - Kepler (K20X)
//#endif
magma_int_t gpu_cpu_perf = magma_get_cbulge_gcperf();
if (threads > 1) {
f = 1. / (1. + (float)(threads-1)/ ((float)gpu_cpu_perf) );
n_gpu = (magma_int_t)(f*ne);
}
/****************************************************
* apply V2 from left to the eigenvectors Z. dZ = (I-V2*T2*V2')*Z
* **************************************************/
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
n_gpu=ne;
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
timeaplQ2 = magma_wtime();
/*============================
* use GPU+CPU's
*==========================*/
if (n_gpu < ne) {
// define the size of Q to be done on CPU's and the size on GPU's
// note that GPU use Q(1:N_GPU) and CPU use Q(N_GPU+1:N)
#ifdef ENABLE_DEBUG
printf("---> calling GPU + CPU(if N_CPU > 0) to apply V2 to Z with NE %d N_GPU %d N_CPU %d\n",ne, n_gpu, ne-n_gpu);
#endif
magma_capplyQ_data data_applyQ;
magma_capplyQ_data_init(&data_applyQ, threads, n, ne, n_gpu, nb, Vblksiz, Z, ldz, V, ldv, TAU, T, ldt, dZ, lddz);
magma_capplyQ_id_data* arg;
magma_malloc_cpu((void**) &arg, threads*sizeof(magma_capplyQ_id_data));
pthread_t* thread_id;
magma_malloc_cpu((void**) &thread_id, threads*sizeof(pthread_t));
pthread_attr_t thread_attr;
// ===============================
// relaunch thread to apply Q
// ===============================
// Set one thread per core
pthread_attr_init(&thread_attr);
pthread_attr_setscope(&thread_attr, PTHREAD_SCOPE_SYSTEM);
pthread_setconcurrency(threads);
// Launch threads
for (magma_int_t thread = 1; thread < threads; thread++) {
magma_capplyQ_id_data_init(&(arg[thread]), thread, &data_applyQ);
pthread_create(&thread_id[thread], &thread_attr, magma_capplyQ_parallel_section, &arg[thread]);
}
magma_capplyQ_id_data_init(&(arg[0]), 0, &data_applyQ);
magma_capplyQ_parallel_section(&arg[0]);
// Wait for completion
for (magma_int_t thread = 1; thread < threads; thread++) {
void *exitcodep;
pthread_join(thread_id[thread], &exitcodep);
}
magma_free_cpu(thread_id);
magma_free_cpu(arg);
magma_capplyQ_data_destroy(&data_applyQ);
magma_csetmatrix( n, ne-n_gpu, Z + n_gpu*ldz, ldz, dZ + n_gpu*ldz, lddz, queue );
/*============================
* use only GPU
*==========================*/
} else {
magma_csetmatrix( n, ne, Z, ldz, dZ, lddz, queue );
magma_cbulge_applyQ_v2(MagmaLeft, ne, n, nb, Vblksiz, dZ, lddz, V, ldv, T, ldt, info);
}
timeaplQ2 = magma_wtime()-timeaplQ2;
magma_queue_destroy( queue );
magma_set_lapack_numthreads(mklth);
return MAGMA_SUCCESS;
}
/**
Purpose
-------
ZHEEVDX_GPU computes selected eigenvalues and, optionally, eigenvectors
of a complex Hermitian matrix A. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm makes very mild assumptions about
floating point arithmetic. It will work on machines with a guard
digit in add/subtract, or on those binary machines without guard
digits which subtract like the Cray X-MP, Cray Y-MP, Cray C-90, or
Cray-2. It could conceivably fail on hexadecimal or decimal machines
without guard digits, but we know of none.
Arguments
---------
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX_16 array on the GPU,
dimension (LDDA, N).
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, if JOBZ = MagmaVec, then if INFO = 0, the first mout columns
of A contains the required
orthonormal eigenvectors of the matrix A.
If JOBZ = MagmaNoVec, then on exit the lower triangle (if UPLO=MagmaLower)
or the upper triangle (if UPLO=MagmaUpper) of A, including the
diagonal, is destroyed.
@param[in]
ldda INTEGER
The leading dimension of the array DA. LDDA >= max(1,N).
@param[in]
vl DOUBLE PRECISION
@param[in]
vu DOUBLE PRECISION
If RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
If RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[out]
mout INTEGER
The total number of eigenvalues found. 0 <= MOUT <= N.
If RANGE = MagmaRangeAll, MOUT = N, and if RANGE = MagmaRangeI, MOUT = IU-IL+1.
@param[out]
w DOUBLE PRECISION array, dimension (N)
If INFO = 0, the required mout eigenvalues in ascending order.
@param
wA (workspace) COMPLEX_16 array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param[out]
work (workspace) COMPLEX_16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( N + N*NB, 2*N + N**2 ).
NB can be obtained through magma_get_zhetrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK, RWORK and
IWORK arrays, returns these values as the first entries of
the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
rwork (workspace) DOUBLE PRECISION array, dimension (LRWORK)
On exit, if INFO = 0, RWORK[0] returns the optimal LRWORK.
@param[in]
lrwork INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LRWORK >= N.
If JOBZ = MagmaVec and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
\n
If LRWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK[0] returns the optimal LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LIWORK >= 1.
If JOBZ = MagmaVec and N > 1, LIWORK >= 3 + 5*N.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i and JOBZ = MagmaNoVec, then the algorithm failed
to converge; i off-diagonal elements of an intermediate
tridiagonal form did not converge to zero;
if INFO = i and JOBZ = MagmaVec, then the algorithm failed
to compute an eigenvalue while working on the submatrix
lying in rows and columns INFO/(N+1) through
mod(INFO,N+1).
Further Details
---------------
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
Modified description of INFO. Sven, 16 Feb 05.
@ingroup magma_zheev_driver
********************************************************************/
extern "C" magma_int_t
magma_zheevdx_gpu(
magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo,
magma_int_t n,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
double vl, double vu, magma_int_t il, magma_int_t iu,
magma_int_t *mout, double *w,
magmaDoubleComplex *wA, magma_int_t ldwa,
magmaDoubleComplex *work, magma_int_t lwork,
#ifdef COMPLEX
double *rwork, magma_int_t lrwork,
#endif
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
magma_int_t ione = 1;
double d__1;
double eps;
magma_int_t inde;
double anrm;
magma_int_t imax;
double rmin, rmax;
double sigma;
magma_int_t iinfo, lwmin;
magma_int_t lower;
magma_int_t llrwk;
magma_int_t wantz;
//magma_int_t indwk2;
magma_int_t iscale;
double safmin;
double bignum;
magma_int_t indtau;
magma_int_t indrwk, indwrk, liwmin;
magma_int_t lrwmin, llwork;
double smlnum;
magma_int_t lquery;
magma_int_t alleig, valeig, indeig;
magmaDouble_ptr dwork;
magmaDoubleComplex_ptr dC;
magma_int_t lddc = ldda;
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1 || lrwork == -1 || liwork == -1);
*info = 0;
if (! (wantz || (jobz == MagmaNoVec))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (ldda < max(1,n)) {
*info = -6;
} else if (ldwa < max(1,n)) {
*info = -14;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_zhetrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
lrwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( n + n*nb, 2*n + n*n );
lrwmin = 1 + 5*n + 2*n*n;
liwmin = 3 + 5*n;
}
else {
lwmin = n + n*nb;
lrwmin = n;
liwmin = 1;
}
// multiply by 1+eps (in Double!) to ensure length gets rounded up,
// if it cannot be exactly represented in floating point.
real_Double_t one_eps = 1. + lapackf77_dlamch("Epsilon");
work[0] = MAGMA_Z_MAKE( lwmin * one_eps, 0 );
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
if ((lwork < lwmin) && !lquery) {
*info = -16;
} else if ((lrwork < lrwmin) && ! lquery) {
*info = -18;
} else if ((liwork < liwmin) && ! lquery) {
*info = -20;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* If matrix is very small, then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
magma_int_t lda = n;
magmaDoubleComplex *A;
magma_zmalloc_cpu( &A, lda*n );
magma_zgetmatrix( n, n, dA, ldda, A, lda );
lapackf77_zheevd( jobz_, uplo_,
&n, A, &lda,
w, work, &lwork,
rwork, &lrwork,
iwork, &liwork, info );
magma_zsetmatrix( n, n, A, lda, dA, ldda );
magma_free_cpu( A );
*mout = n;
return *info;
}
magma_queue_t stream;
magma_queue_create( &stream );
// dC and dwork are never used together, so use one buffer for both;
// unfortunately they're different types (complex and double).
// (this is easier in dsyevd_gpu where everything is double.)
// zhetrd2_gpu requires ldda*ceildiv(n,64) + 2*ldda*nb, in double-complex.
// zunmtr_gpu requires lddc*n, in double-complex.
// zlanhe requires n, in double.
magma_int_t ldwork = max( ldda*ceildiv(n,64) + 2*ldda*nb, lddc*n );
magma_int_t ldwork_real = max( ldwork*2, n );
if ( wantz ) {
// zstedx requrise 3n^2/2, in double
ldwork_real = max( ldwork_real, 3*n*(n/2 + 1) );
}
if (MAGMA_SUCCESS != magma_dmalloc( &dwork, ldwork_real )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dC = (magmaDoubleComplex*) dwork;
/* Get machine constants. */
safmin = lapackf77_dlamch("Safe minimum");
eps = lapackf77_dlamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_dsqrt( smlnum );
rmax = magma_dsqrt( bignum );
/* Scale matrix to allowable range, if necessary. */
anrm = magmablas_zlanhe( MagmaMaxNorm, uplo, n, dA, ldda, dwork );
iscale = 0;
sigma = 1;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
magmablas_zlascl( uplo, 0, 0, 1., sigma, n, n, dA, ldda, info );
}
/* Call ZHETRD to reduce Hermitian matrix to tridiagonal form. */
// zhetrd rwork: e (n)
// zstedx rwork: e (n) + llrwk (1 + 4*N + 2*N**2) ==> 1 + 5n + 2n^2
inde = 0;
indrwk = inde + n;
llrwk = lrwork - indrwk;
// zhetrd work: tau (n) + llwork (n*nb) ==> n + n*nb
// zstedx work: tau (n) + z (n^2)
// zunmtr work: tau (n) + z (n^2) + llwrk2 (n or n*nb) ==> 2n + n^2, or n + n*nb + n^2
indtau = 0;
indwrk = indtau + n;
//indwk2 = indwrk + n*n;
llwork = lwork - indwrk;
//llwrk2 = lwork - indwk2;
magma_timer_t time=0;
timer_start( time );
#ifdef FAST_HEMV
magma_zhetrd2_gpu( uplo, n, dA, ldda, w, &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
dC, ldwork, &iinfo );
#else
magma_zhetrd_gpu ( uplo, n, dA, ldda, w, &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
&iinfo );
#endif
timer_stop( time );
timer_printf( "time zhetrd_gpu = %6.2f\n", time );
/* For eigenvalues only, call DSTERF. For eigenvectors, first call
ZSTEDC to generate the eigenvector matrix, WORK(INDWRK), of the
tridiagonal matrix, then call ZUNMTR to multiply it to the Householder
transformations represented as Householder vectors in A. */
if (! wantz) {
lapackf77_dsterf( &n, w, &rwork[inde], info );
magma_dmove_eig( range, n, w, &il, &iu, vl, vu, mout );
}
else {
timer_start( time );
magma_zstedx( range, n, vl, vu, il, iu, w, &rwork[inde],
&work[indwrk], n, &rwork[indrwk],
llrwk, iwork, liwork, dwork, info );
timer_stop( time );
timer_printf( "time zstedx = %6.2f\n", time );
timer_start( time );
magma_dmove_eig( range, n, w, &il, &iu, vl, vu, mout );
magma_zsetmatrix( n, *mout, &work[indwrk + n * (il-1) ], n, dC, lddc );
magma_zunmtr_gpu( MagmaLeft, uplo, MagmaNoTrans, n, *mout, dA, ldda, &work[indtau],
dC, lddc, wA, ldwa, &iinfo );
magma_zcopymatrix( n, *mout, dC, lddc, dA, ldda );
timer_stop( time );
timer_printf( "time zunmtr_gpu + copy = %6.2f\n", time );
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = n;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_dscal( &imax, &d__1, w, &ione );
}
work[0] = MAGMA_Z_MAKE( lwmin * one_eps, 0 ); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
magma_queue_destroy( stream );
magma_free( dwork );
return *info;
} /* magma_zheevdx_gpu */
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_trans_t trans[] = { MagmaNoTrans, MagmaConjTrans, MagmaTrans };
magma_uplo_t uplo [] = { MagmaLower, MagmaUpper };
magma_diag_t diag [] = { MagmaUnit, MagmaNonUnit };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magmaFloatComplex *A, *B, *C, *C2, *LU;
magmaFloatComplex_ptr dA, dB, dC1, dC2;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.5, 0.1 );
magmaFloatComplex beta = MAGMA_C_MAKE( 0.7, 0.2 );
float dalpha = 0.6;
float dbeta = 0.8;
float work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_int_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int itest = 0; itest < opts.ntest; ++itest ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
printf("=========================================================================\n");
printf( "m=%d, n=%d, k=%d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = max( 1, maxn );
size = ld*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_cmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_cmalloc( &dA, size ); assert( err == 0 );
err = magma_cmalloc( &dB, size ); assert( err == 0 );
err = magma_cmalloc( &dC1, size ); assert( err == 0 );
err = magma_cmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_clarnv( &ione, ISEED, &size, A );
lapackf77_clarnv( &ione, ISEED, &size, B );
lapackf77_clarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test CSWAP
// swap columns 2 and 3 of dA, then copy to C2 and compare with A
if ( n >= 3 ) {
magma_csetmatrix( m, n, A, ld, dA, ld );
magma_csetmatrix( m, n, A, ld, dB, ld );
magma_cswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_cswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dA, 1, dB, 1 );
magma_cgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_clange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "cswap diff %.2g\n", error );
}
else {
printf( "cswap skipped for n < 3\n" );
}
// ----- test ICAMAX
// get argmax of column of A
magma_csetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_icamax( m, dA(0,j), 1 );
int i2; // NOT magma_int_t, for cublas
cublasIcamax( opts.handle, m, dA(0,j), 1, &i2 );
// todo need sync here?
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (float)m * k / 1e9;
printf( "icamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test CGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_csetmatrix( m, n, A, ld, dA, ld );
magma_csetvector( maxn, B, 1, dB, 1 );
magma_csetvector( maxn, C, 1, dC1, 1 );
magma_csetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_cgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCgemv( opts.handle, cublas_trans_const(trans[ia]),
m, n, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == MagmaNoTrans ? m : n);
cublasCaxpy( opts.handle, size, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CGEMV( m, n ) / 1e9;
printf( "cgemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test CHEMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_csetmatrix( m, m, A, ld, dA, ld );
magma_csetvector( m, B, 1, dB, 1 );
magma_csetvector( m, C, 1, dC1, 1 );
magma_csetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_chemv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasChemv( opts.handle, cublas_uplo_const(uplo[iu]),
m, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHEMV( m ) / 1e9;
printf( "chemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test CTRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_clacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_cgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_csetmatrix( m, m, LU, ld, dA, ld );
magma_csetvector( m, C, 1, dC1, 1 );
magma_csetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ctrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCtrsv( opts.handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
cublas_diag_const(diag[id]), m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "ctrsv( %c, %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]), lapacke_diag_const(diag[id]),
error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test CGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == MagmaNoTrans);
bool ntb = (trans[ib] == MagmaNoTrans);
magma_csetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_csetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCgemm( opts.handle, cublas_trans_const(trans[ia]), cublas_trans_const(trans[ib]),
m, n, k, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CGEMM( m, n, k ) / 1e9;
printf( "cgemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), lapacke_trans_const(trans[ib]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CHEMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_csetmatrix( m, m, A, ld, dA, ld );
magma_csetmatrix( m, n, B, ld, dB, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_chemm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasChemm( opts.handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
m, n, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHEMM( side[is], m, n ) / 1e9;
printf( "chemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_side_const(side[is]), lapacke_uplo_const(uplo[iu]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CHERK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_csetmatrix( n, k, A, ld, dA, ld );
magma_csetmatrix( n, n, C, ld, dC1, ld );
magma_csetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cherk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCherk( opts.handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &dalpha, dA, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHERK( k, n ) / 1e9;
printf( "cherk( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CHER2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == MagmaNoTrans);
magma_csetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_csetmatrix( n, n, C, ld, dC1, ld );
magma_csetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cher2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCher2k( opts.handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &alpha, dA, ld, dB, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHER2K( k, n ) / 1e9;
printf( "cher2k( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CTRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_csetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ctrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
// note cublas does trmm out-of-place (i.e., adds output matrix C),
// but allows C=B to do in-place.
t2 = magma_sync_wtime( 0 );
cublasCtrmm( opts.handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRMM( side[is], m, n ) / 1e9;
printf( "ctrmm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test CTRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_csetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ctrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCtrsm( opts.handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRSM( side[is], m, n ) / 1e9;
printf( "ctrsm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
fflush( stdout );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
int status = (total_error != 0.);
return status;
}
int main( int argc, char** argv )
{
magma_init();
cublasHandle_t handle;
cudaSetDevice( 0 );
cublasCreate( &handle );
magmaFloatComplex *A, *B, *C;
magmaFloatComplex *dA, *dB, *dC;
float error, work[1];
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = { 1, 2, 3, 4 };
magma_int_t n = 10;
magma_int_t lda = n;
magma_int_t ldda = ((n+31)/32)*32;
magma_int_t size = lda*n;
magma_int_t info;
magma_cmalloc_cpu( &A, lda*n );
magma_cmalloc_cpu( &B, lda*n );
magma_cmalloc_cpu( &C, lda*n );
magma_cmalloc( &dA, ldda*n );
magma_cmalloc( &dB, ldda*n );
magma_cmalloc( &dC, ldda*n );
// initialize matrices
lapackf77_clarnv( &ione, ISEED, &size, A );
lapackf77_clarnv( &ione, ISEED, &size, B );
lapackf77_clarnv( &ione, ISEED, &size, C );
// increase diagonal to be SPD
for( int i=0; i < n; ++i ) {
C[i+i*lda] = MAGMA_C_ADD( C[i+i*lda], MAGMA_C_MAKE( n*n, 0 ));
}
magma_csetmatrix( n, n, A, lda, dA, ldda );
magma_csetmatrix( n, n, B, lda, dB, ldda );
magma_csetmatrix( n, n, C, lda, dC, ldda );
// compute with cublas
cublasCgemm( handle, CUBLAS_OP_N, CUBLAS_OP_N, n, n, n,
&c_neg_one, dA, ldda, dB, ldda, &c_one, dC, ldda );
magma_cpotrf_gpu( MagmaLower, n, dC, ldda, &info );
if (info != 0)
printf("magma_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute with LAPACK
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &n, &n, &n,
&c_neg_one, A, &lda, B, &lda, &c_one, C, &lda );
lapackf77_cpotrf( MagmaLowerStr, &n, C, &lda, &info );
if (info != 0)
printf("lapackf77_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute difference
magma_cgetmatrix( n, n, dC, ldda, A, lda );
blasf77_caxpy( &size, &c_neg_one, C, &ione, A, &ione );
error = lapackf77_clange( "F", &n, &n, A, &lda, work );
printf( "n %d, error %8.2e\n", (int) n, error );
magma_free( dA );
magma_free( dB );
magma_free( dC );
magma_free_cpu( A );
magma_free_cpu( B );
magma_free_cpu( C );
cublasDestroy( handle );
magma_finalize();
return 0;
}
extern "C" magma_int_t
magma_zcgeqrsv_gpu(magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *dB, magma_int_t lddb,
magmaDoubleComplex *dX, magma_int_t lddx,
magma_int_t *iter, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZCGEQRSV solves the least squares problem
min || A*X - B ||,
where A is an M-by-N matrix and X and B are M-by-NRHS matrices.
ZCGEQRSV first attempts to factorize the matrix in complex SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with complex DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
complex DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio complex SINGLE PRECISION performance over complex DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. M >= N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
dA (input or input/output) COMPLEX_16 array on the GPU, dimension (LDDA,N)
On entry, the M-by-N coefficient matrix A.
On exit, if iterative refinement has been successfully used
(info.EQ.0 and ITER.GE.0, see description below), A is
unchanged. If double precision factorization has been used
(info.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the QR factorization of A as returned by
function DGEQRF_GPU.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
dB (input or input/output) COMPLEX_16 array on the GPU, dimension (LDDB,NRHS)
The M-by-NRHS right hand side matrix B.
May be overwritten (e.g., if refinement fails).
LDDB (input) INTEGER
The leading dimension of the array dB. LDDB >= max(1,M).
dX (output) COMPLEX_16 array on the GPU, dimension (LDDX,NRHS)
If info = 0, the N-by-NRHS solution matrix X.
LDDX (input) INTEGER
The leading dimension of the array dX. LDDX >= max(1,N).
ITER (output) INTEGER
< 0: iterative refinement has failed, double precision
factorization has been performed
-1 : the routine fell back to full precision for
implementation- or machine-specific reasons
-2 : narrowing the precision induced an overflow,
the routine fell back to full precision
-3 : failure of SGEQRF
-31: stop the iterative refinement after the 30th iteration
> 0: iterative refinement has been successfully used.
Returns the number of iterations
INFO (output) INTEGER
= 0: successful exit
< 0: if info = -i, the i-th argument had an illegal value
===================================================================== */
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
#define dSX(i,j) (dSX + (i) + (j)*lddsx)
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
magmaDoubleComplex *dworkd, *hworkd;
magmaFloatComplex *dworks, *hworks;
magmaDoubleComplex *dR, *tau, *dT;
magmaFloatComplex *dSA, *dSX, *dST, *stau;
magmaDoubleComplex Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddsx, lddr, nb, lhwork, minmn, size, ldworkd;
/* Check arguments */
*iter = 0;
*info = 0;
if ( m < 0 )
*info = -1;
else if ( n < 0 || n > m )
*info = -2;
else if ( nrhs < 0 )
*info = -3;
else if ( ldda < max(1,m))
*info = -5;
else if ( lddb < max(1,m))
*info = -7;
else if ( lddx < max(1,n))
*info = -9;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( m == 0 || n == 0 || nrhs == 0 )
return *info;
nb = magma_get_cgeqrf_nb(m);
minmn= min(m, n);
/* dSX contains both B and X, so must be max(m or lddb,n). */
lddsa = ldda;
lddsx = max(lddb,n);
lddr = lddb;
/*
* Allocate temporary buffers
*/
/* dworks(dSA + dSX + dST) */
size = lddsa*n + lddsx*nrhs + ( 2*minmn + ((n+31)/32)*32 )*nb;
if (MAGMA_SUCCESS != magma_cmalloc( &dworks, size )) {
fprintf(stderr, "Allocation of dworks failed (%d)\n", (int) size);
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dSA = dworks;
dSX = dSA + lddsa*n;
dST = dSX + lddsx*nrhs;
/* dworkd(dR) = lddr*nrhs */
ldworkd = lddr*nrhs;
if (MAGMA_SUCCESS != magma_zmalloc( &dworkd, ldworkd )) {
magma_free( dworks );
fprintf(stderr, "Allocation of dworkd failed\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dR = dworkd;
/* hworks(workspace for cgeqrs + stau) = min(m,n) + lhworks */
lhwork = (m - n + nb)*(nrhs + nb) + nrhs*nb;
size = lhwork + minmn;
magma_cmalloc_cpu( &hworks, size );
if ( hworks == NULL ) {
magma_free( dworks );
magma_free( dworkd );
fprintf(stderr, "Allocation of hworks failed\n");
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
stau = hworks + lhwork;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_zlange('I', m, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
magmablas_zlag2c( m, nrhs, dB, lddb, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_zlag2c( m, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_cgeqrf_gpu( m, n, dSA, lddsa, stau, dST, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// solve dSA*dSX = dB in single precision
magma_cgeqrs_gpu( m, n, nrhs, dSA, lddsa, stau, dST, dSX, lddsx, hworks, lhwork, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// residual dR = dB - dA*dX in double precision
magmablas_clag2z( n, nrhs, dSX, lddsx, dX, lddx, info );
magmablas_zlacpy( MagmaUpperLower, m, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_zgemv( MagmaNoTrans, m, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_Z_ABS( dX(i,j) ) instead of zlange?
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( m, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
/* Free workspaces */
magma_free( dworks );
magma_free( dworkd );
magma_free_cpu( hworks );
return *info;
REFINEMENT:
/* TODO: this iterative refinement algorithm works only for compatibile
* systems (B in colspan of A).
* See Matrix Computations (3rd ed) p. 267 for correct algorithm. */
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
magmablas_zlag2c( m, nrhs, dR, lddr, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// solve dSA*dSX = R in single precision
magma_cgeqrs_gpu( m, n, nrhs, dSA, lddsa, stau, dST, dSX, lddsx, hworks, lhwork, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Add correction and setup residual
// dX += dSX [including conversion] --and--
// dR[1:n] = dB[1:n] (only n rows, not whole m rows! -- useless if m > n)
for( j=0; j < nrhs; j++ ) {
magmablas_zcaxpycp( n, dSX(0,j), dX(0,j), dB(0,j), dR(0,j) );
}
// dR = dB (whole m rows)
magmablas_zlacpy( MagmaUpperLower, m, nrhs, dB, lddb, dR, lddr );
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_zgemv( MagmaNoTrans, m, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER>0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( m, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
/* Free workspaces */
magma_free( dworks );
magma_free( dworkd );
magma_free_cpu( hworks );
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_free( dworks );
magma_free_cpu( hworks );
/*
* Allocate temporary buffers
*/
/* dworkd = dT for zgeqrf */
nb = magma_get_zgeqrf_nb( m );
size = (2*min(m, n) + (n+31)/32*32 )*nb;
if ( size > ldworkd ) {
magma_free( dworkd );
if (MAGMA_SUCCESS != magma_zmalloc( &dworkd, size )) {
fprintf(stderr, "Allocation of dworkd2 failed\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
dT = dworkd;
/* hworkd(dtau + workspace for zgeqrs) = min(m,n) + lhwork */
size = lhwork + minmn;
magma_zmalloc_cpu( &hworkd, size );
if ( hworkd == NULL ) {
magma_free( dworkd );
fprintf(stderr, "Allocation of hworkd2 failed\n");
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
tau = hworkd + lhwork;
magma_zgeqrf_gpu( m, n, dA, ldda, tau, dT, info );
if (*info == 0) {
// if m > n, then dB won't fit in dX, so solve with dB and copy n rows to dX
magma_zgeqrs_gpu( m, n, nrhs, dA, ldda, tau, dT, dB, lddb, hworkd, lhwork, info );
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
}
magma_free( dworkd );
magma_free_cpu( hworkd );
return *info;
}
/**
Purpose
-------
Solves the overdetermined, least squares problem
min || A*X - C ||
using the QR factorization A.
The underdetermined problem (m < n) is not currently handled.
Arguments
---------
@param[in]
trans magma_trans_t
- = MagmaNoTrans: the linear system involves A.
Only TRANS=MagmaNoTrans is currently handled.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. M >= N >= 0.
@param[in]
nrhs INTEGER
The number of columns of the matrix C. NRHS >= 0.
@param[in,out]
dA REAL array on the GPU, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, A is overwritten by details of its QR
factorization as returned by SGEQRF.
@param[in]
ldda INTEGER
The leading dimension of the array A, LDDA >= M.
@param[in,out]
dB REAL array on the GPU, dimension (LDDB,NRHS)
On entry, the M-by-NRHS matrix C.
On exit, the N-by-NRHS solution matrix X.
@param[in]
lddb INTEGER
The leading dimension of the array dB. LDDB >= M.
@param[out]
hwork (workspace) REAL array, dimension MAX(1,LWORK).
On exit, if INFO = 0, HWORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array HWORK,
LWORK >= (M - N + NB)*(NRHS + NB) + NRHS*NB,
where NB is the blocksize given by magma_get_sgeqrf_nb( M ).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the HWORK array, returns
this value as the first entry of the HWORK array.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_sgels_driver
********************************************************************/
extern "C" magma_int_t
magma_sgels_gpu( magma_trans_t trans, magma_int_t m, magma_int_t n, magma_int_t nrhs,
float *dA, magma_int_t ldda,
float *dB, magma_int_t lddb,
float *hwork, magma_int_t lwork,
magma_int_t *info)
{
float *dT;
float *tau;
magma_int_t k;
magma_int_t nb = magma_get_sgeqrf_nb(m);
magma_int_t lwkopt = (m - n + nb)*(nrhs + nb) + nrhs*nb;
int lquery = (lwork == -1);
hwork[0] = MAGMA_S_MAKE( (float)lwkopt, 0. );
*info = 0;
/* For now, N is the only case working */
if ( trans != MagmaNoTrans )
*info = -1;
else if (m < 0)
*info = -2;
else if (n < 0 || m < n) /* LQ is not handle for now*/
*info = -3;
else if (nrhs < 0)
*info = -4;
else if (ldda < max(1,m))
*info = -6;
else if (lddb < max(1,m))
*info = -8;
else if (lwork < lwkopt && ! lquery)
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
hwork[0] = MAGMA_S_ONE;
return *info;
}
/*
* Allocate temporary buffers
*/
int ldtwork = ( 2*k + ((n+31)/32)*32 )*nb;
if (nb < nrhs)
ldtwork = ( 2*k + ((n+31)/32)*32 )*nrhs;
if (MAGMA_SUCCESS != magma_smalloc( &dT, ldtwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_smalloc_cpu( &tau, k );
if ( tau == NULL ) {
magma_free( dT );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_sgeqrf_gpu( m, n, dA, ldda, tau, dT, info );
if ( *info == 0 ) {
magma_sgeqrs_gpu( m, n, nrhs,
dA, ldda, tau, dT,
dB, lddb, hwork, lwork, info );
}
magma_free( dT );
magma_free_cpu(tau);
return *info;
}
