// --------------------
magma_err_t
magma_zgetvector(
magma_int_t n,
magmaDoubleComplex_const_ptr dA_src, size_t dA_offset, magma_int_t incx,
magmaDoubleComplex* hA_dst, size_t hA_offset, magma_int_t incy,
magma_queue_t queue )
{
if ( n <= 0 )
return MAGMA_SUCCESS;
cl_int err;
if(incx ==1 && incy ==1) {
err = clEnqueueReadBuffer(
queue, dA_src, CL_TRUE,
dA_offset*sizeof(magmaDoubleComplex), n*sizeof(magmaDoubleComplex),
hA_dst+hA_offset, 0, NULL, gevent);
return err;
} else {
magma_int_t ldda = incx;
magma_int_t ldha = incy;
err = magma_zgetmatrix(1, n,
dA_src, dA_offset, ldda,
hA_dst, hA_offset, ldha,
queue);
return err;
}
}
// --------------------
extern "C" void
magma_zgetvector(
magma_int_t n,
magmaDoubleComplex_const_ptr dx_src, size_t dx_offset, magma_int_t incx,
magmaDoubleComplex* hy_dst, magma_int_t incy,
magma_queue_t queue )
{
if (n <= 0)
return;
if (incx == 1 && incy == 1) {
cl_int err = clEnqueueReadBuffer(
queue, dx_src, CL_TRUE,
dx_offset*sizeof(magmaDoubleComplex), n*sizeof(magmaDoubleComplex),
hy_dst, 0, NULL, g_event);
check_error( err );
}
else {
magma_int_t ldda = incx;
magma_int_t ldhb = incy;
magma_zgetmatrix( 1, n,
dx_src, dx_offset, ldda,
hy_dst, ldhb,
queue);
}
}
magma_int_t magma_znan_inf_gpu(
magma_uplo_t uplo, magma_int_t m, magma_int_t n,
magmaDoubleComplex_const_ptr dA, magma_int_t dA_offset, magma_int_t ldda,
magma_int_t *cnt_nan,
magma_int_t *cnt_inf,
magma_queue_t queue )
{
magma_int_t info = 0;
if ( uplo != MagmaLower && uplo != MagmaUpper && uplo != MagmaFull )
info = -1;
else 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;
}
magma_int_t lda = m;
magmaDoubleComplex* A;
magma_zmalloc_cpu( &A, lda*n );
magma_zgetmatrix( m, n, dA, dA_offset, ldda, A, lda, queue );
magma_int_t cnt = magma_znan_inf( uplo, m, n, A, lda, cnt_nan, cnt_inf );
magma_free_cpu( A );
return cnt;
}
Tensor_hao<complex<double>,2> inverse_magma(const LUDecomp<complex<double>>& x)
{
magma_int_t N=x.A.rank(0); magma_int_t info;
magmaDoubleComplex_ptr d_A , dwork;
magma_int_t lda, ldwork;
lda = ((N+31)/32)*32; //round up to multiple of 32 for best GPU performance
ldwork = N*magma_get_zgetri_nb(N); // magma_get_zgetri_nb optimizes the blocksize
magma_zmalloc( &d_A, lda*N ); magma_zmalloc( &dwork, ldwork );
//copy matrix from CPU to GPU
magma_zsetmatrix( N, N, (magmaDoubleComplex* )x.A.data(), N, d_A, lda );
//calculate the inverse matrix with zgetri
magma_zgetri_gpu( N, d_A, lda, (magma_int_t*) x.ipiv.data(), dwork, ldwork, &info );
if(info<0) {cout<<"The "<<info<<"-th parameter is illegal in inverse_magma!"<<endl; exit(1);}
//copy matrix from GPU to CPU
Tensor_hao<complex<double>,2> A(N,N);
magma_zgetmatrix( N, N, d_A, lda, (magmaDoubleComplex* )A.data(), N );
magma_free(d_A); magma_free(dwork);
return A;
}
LUDecomp<complex<double>> LUconstruct_magma(const Tensor_core<complex<double>,2>& x)
{
if( x.rank(0) != x.rank(1) ) {cout<<"Input for LU is not square matrix!"<<endl; exit(1);}
//Create LU object
LUDecomp<complex<double>> y;
y.A = Tensor_hao< complex<double>, 2 > ( x.n_ptr() );
y.ipiv = Tensor_hao<int,1>( x.rank(0) );
//Prepare for zgetrf
magma_int_t M = x.rank(0), N = x.rank(1);
magma_int_t LDA = ((M+31)/32)*32;
magmaDoubleComplex_ptr d_A; magma_zmalloc(&d_A, LDA*N);
magma_int_t info;
//Transfer data and call zgetrf
magma_zsetmatrix(M, N, (magmaDoubleComplex* ) x.data(), M, d_A, LDA );
magma_zgetrf_gpu(M, N, d_A, LDA, (magma_int_t*) y.ipiv.data(), &info);
magma_zgetmatrix(M, N, d_A, LDA, (magmaDoubleComplex* ) y.A.data(), M);
y.info=info;
//Clean
magma_free(d_A);
if(y.info<0) {cout<<"The "<<y.info<<"-th parameter is illegal in LUconstruct_magma!"<<endl; exit(1);}
return y;
}
// --------------------
magma_err_t
magma_zher2k(
magma_uplo_t uplo, magma_trans_t trans,
magma_int_t n, magma_int_t k,
magmaDoubleComplex alpha, magmaDoubleComplex_const_ptr dA, size_t dA_offset, magma_int_t lda,
magmaDoubleComplex_const_ptr dB, size_t dB_offset, magma_int_t ldb,
double beta, magmaDoubleComplex_ptr dC, size_t dC_offset, magma_int_t ldc,
magma_queue_t queue)
{ // cblas wrapper
magma_int_t ka, kb;
if(trans == MagmaNoTrans){
ka = k;
kb = k;
}else{
ka = n;
kb = n;
}
magmaDoubleComplex *hA, *hB, *hC;
hA = (magmaDoubleComplex*)malloc(lda*ka*sizeof(magmaDoubleComplex));
hB = (magmaDoubleComplex*)malloc(ldb*kb*sizeof(magmaDoubleComplex));
hC = (magmaDoubleComplex*)malloc(ldc*n*sizeof(magmaDoubleComplex));
magma_zgetmatrix(lda, ka, dA, dA_offset, lda, hA, 0, lda, queue);
magma_zgetmatrix(ldb, kb, dB, dB_offset, ldb, hB, 0, ldb, queue);
magma_zgetmatrix(ldc, n, dC, dC_offset, ldc, hC, 0, ldc, queue);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zher2k(CblasColMajor, cblas_uplo_const(uplo), cblas_trans_const(trans),
n, k, (void*)&alpha, hA, lda, hB, ldb, beta, hC, ldc);
#else
cblas_zher2k(CblasColMajor, cblas_uplo_const(uplo), cblas_trans_const(trans),
n, k, alpha, hA, lda, hB, ldb, beta, hC, ldc);
#endif
magma_zsetmatrix(ldc, n, hC, 0, ldc, dC, dC_offset, ldc, queue);
free(hA);
free(hB);
free(hC);
return CL_SUCCESS;
}
void gmm_magma(const Tensor_core<complex<double>,2>& A, const Tensor_core<complex<double>,2>& B, Tensor_core<complex<double>,2>& C,
char TRANSA, char TRANSB, complex<double> alpha, complex<double> beta)
{
int AL0 = A.rank(0); int AL1 = A.rank(1);
int BL0 = B.rank(0); int BL1 = B.rank(1);
int CL0 = C.rank(0); int CL1 = C.rank(1);
magma_int_t M, N, K, LDA, LDB, LDC;
magma_trans_t transA=magma_trans_const(TRANSA), transB=magma_trans_const(TRANSB);
magmaDoubleComplex_ptr d_A, d_B, d_C;
//Set LDA, LDB, and LDC, round up to multiple of 32 for best GPU performance
LDA = ((AL0+31)/32)*32; LDB = ((BL0+31)/32)*32; LDC = ((CL0+31)/32)*32;
// Allocate memory for the matrices on GPU
magma_zmalloc(&d_A, LDA*AL1 );
magma_zmalloc(&d_B, LDB*BL1 );
magma_zmalloc(&d_C, LDC*CL1 );
// Copy data from host (CPU) to device (GPU)
magma_zsetmatrix( AL0, AL1, (magmaDoubleComplex* ) A.data(), AL0, d_A, LDA );
magma_zsetmatrix( BL0, BL1, (magmaDoubleComplex* ) B.data(), BL0, d_B, LDB );
if( abs(beta)>1e-32 ) magma_zsetmatrix( CL0, CL1, (magmaDoubleComplex* ) C.data(), CL0, d_C, LDC );
//Call magma_sgemm
M=( TRANSA=='N' || TRANSA=='n' ) ? AL0:AL1;
K=( TRANSA=='N' || TRANSA=='n' ) ? AL1:AL0;
N=( TRANSB=='N' || TRANSB=='n' ) ? BL1:BL0;
magma_zgemm(transA, transB, M, N, K, _cast_Z(alpha), d_A, LDA, d_B, LDB, _cast_Z(beta),d_C, LDC);
// Copy solution from device (GPU) to host (CPU)
magma_zgetmatrix(CL0, CL1, d_C, LDC, (magmaDoubleComplex* ) C.data(), CL0);
// Free memory on GPU
magma_free(d_A); magma_free(d_B); magma_free(d_C);
}
Tensor_hao<complex<double>,2> solve_lineq_magma(const LUDecomp<complex<double>>& x, const Tensor_core<complex<double>,2>& B, char TRANS)
{
if( x.A.rank(0) != B.rank(0) ) {cout<<"Input size for solving linear equation is not consistent!"<<endl; exit(1);}
magma_int_t N=B.rank(0); magma_int_t NRHS=B.rank(1); magma_int_t info;
magma_trans_t Trans = magma_trans_const(TRANS);
magmaDoubleComplex_ptr d_A, d_B;
magma_int_t lda, ldb;
lda = ((N+31)/32)*32;
ldb = ((N+31)/32)*32;
//allocate memory on GPU
magma_zmalloc( &d_A, lda*N );
magma_zmalloc( &d_B, ldb*NRHS );
//copy matrix from CPU to GPU
magma_zsetmatrix( N, N, (magmaDoubleComplex* )x.A.data(), N, d_A, lda );
magma_zsetmatrix( N, NRHS, (magmaDoubleComplex* )B.data(), N, d_B, ldb );
//Solve the equation
magma_zgetrs_gpu( Trans, N, NRHS, d_A, lda, (magma_int_t*)x.ipiv.data(), d_B, ldb, &info );
if(info!=0)
{
cout<<"Solve linear equation is not suceesful: "<<info<<"-th parameter is illegal!"<<endl;
exit(1);
}
//copy matrix from GPU to CPU
Tensor_hao<complex<double>,2> M(N,NRHS);
magma_zgetmatrix( N, NRHS, d_B, ldb, (magmaDoubleComplex* ) M.data(), N );
//free memory
magma_free( d_A ); magma_free( d_B );
return M;
}
/**
Purpose
-------
ZPOTRF computes the Cholesky factorization of a complex Hermitian
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
dA COMPLEX_16 array on the GPU, dimension (LDDA,N)
On entry, the Hermitian matrix dA. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_zposv_comp
********************************************************************/
extern "C" magma_int_t
magma_zpotrf_gpu(
magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magma_int_t *info )
{
#ifdef HAVE_clBLAS
#define dA(i_, j_) dA, ((i_) + (j_)*ldda + dA_offset)
#else
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#endif
/* Constants */
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const double d_one = 1.0;
const double d_neg_one = -1.0;
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
bool upper = (uplo == MagmaUpper);
magma_int_t j, jb, nb;
magmaDoubleComplex *work;
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
nb = magma_get_zpotrf_nb( n );
if (MAGMA_SUCCESS != magma_zmalloc_pinned( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
if (nb <= 1 || nb >= n) {
/* Use unblocked code. */
magma_zgetmatrix( n, n, dA(0,0), ldda, work, n, queues[0] );
lapackf77_zpotrf( uplo_, &n, work, &n, info );
magma_zsetmatrix( n, n, work, n, dA(0,0), ldda, queues[0] );
}
else {
/* Use blocked code. */
if (upper) {
//=========================================================
/* Compute the Cholesky factorization A = U'*U. */
for (j=0; j < n; j += nb) {
// apply all previous updates to diagonal block,
// then transfer it to CPU
jb = min( nb, n-j );
magma_zherk( MagmaUpper, MagmaConjTrans, jb, j,
d_neg_one, dA(0, j), ldda,
d_one, dA(j, j), ldda, queues[1] );
magma_queue_sync( queues[1] );
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, queues[0] );
// apply all previous updates to block row right of diagonal block
if (j+jb < n) {
magma_zgemm( MagmaConjTrans, MagmaNoTrans,
jb, n-j-jb, j,
c_neg_one, dA(0, j ), ldda,
dA(0, j+jb), ldda,
c_one, dA(j, j+jb), ldda, queues[1] );
}
// simultaneous with above zgemm, transfer diagonal block,
// factor it on CPU, and test for positive definiteness
magma_queue_sync( queues[0] );
lapackf77_zpotrf( MagmaUpperStr, &jb, work, &jb, info );
magma_zsetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, queues[1] );
if (*info != 0) {
*info = *info + j;
break;
}
// apply diagonal block to block row right of diagonal block
if (j+jb < n) {
magma_ztrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, n-j-jb,
c_one, dA(j, j), ldda,
dA(j, j+jb), ldda, queues[1] );
}
}
}
else {
//=========================================================
// Compute the Cholesky factorization A = L*L'.
for (j=0; j < n; j += nb) {
// apply all previous updates to diagonal block,
// then transfer it to CPU
jb = min( nb, n-j );
magma_zherk( MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dA(j, 0), ldda,
d_one, dA(j, j), ldda, queues[1] );
magma_queue_sync( queues[1] );
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, queues[0] );
// apply all previous updates to block column below diagonal block
if (j+jb < n) {
magma_zgemm( MagmaNoTrans, MagmaConjTrans,
n-j-jb, jb, j,
c_neg_one, dA(j+jb, 0), ldda,
dA(j, 0), ldda,
c_one, dA(j+jb, j), ldda, queues[1] );
}
// simultaneous with above zgemm, transfer diagonal block,
// factor it on CPU, and test for positive definiteness
magma_queue_sync( queues[0] );
lapackf77_zpotrf( MagmaLowerStr, &jb, work, &jb, info );
magma_zsetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, queues[1] );
if (*info != 0) {
*info = *info + j;
break;
}
// apply diagonal block to block column below diagonal
if (j+jb < n) {
magma_ztrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-j-jb, jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda, queues[1] );
}
}
}
}
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free_pinned( work );
return *info;
} /* magma_zpotrf_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgetrf
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R, *work;
magmaDoubleComplex_ptr d_A, dwork;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t N, n2, lda, ldda, info, lwork, ldwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex tmp;
double error, rwork[1];
magma_int_t *ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
// need looser bound (3000*eps instead of 30*eps) for tests
// TODO: should compute ||I - A*A^{-1}|| / (n*||A||*||A^{-1}||)
opts.tolerance = max( 3000., opts.tolerance );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / (N*||A||_F)\n");
printf("=================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
ldwork = N * magma_get_zgetri_nb( N );
gflops = FLOPS_ZGETRI( N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_zgetri( &N, NULL, &lda, NULL, &tmp, &lwork, &info );
if (info != 0)
printf("lapackf77_zgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = int( MAGMA_Z_REAL( tmp ));
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( work, magmaDoubleComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dwork, magmaDoubleComplex, ldwork );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
error = lapackf77_zlange( "f", &N, &N, h_A, &lda, rwork ); // norm(A)
/* Factor the matrix. Both MAGMA and LAPACK will use this factor. */
magma_zsetmatrix( N, N, h_A, lda, d_A, 0, ldda, opts.queue );
magma_zgetrf_gpu( N, N, d_A, 0, ldda, ipiv, opts.queue, &info );
magma_zgetmatrix( N, N, d_A, 0, ldda, h_A, lda, opts.queue );
if ( info != 0 )
printf("magma_zgetrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// check for exact singularity
//h_A[ 10 + 10*lda ] = MAGMA_Z_MAKE( 0.0, 0.0 );
//magma_zsetmatrix( N, N, h_A, lda, d_A, 0, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zgetri_gpu( N, d_A, 0, ldda, ipiv, dwork, 0, ldwork, opts.queues2, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgetri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_zgetmatrix( N, N, d_A, 0, ldda, h_R, lda, opts.queue );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_zgetri( &N, h_A, &lda, ipiv, work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_zaxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_zlange( "f", &N, &N, h_R, &lda, rwork ) / (N*error);
printf( "%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf( "%5d --- ( --- ) %7.2f (%7.2f) ---\n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
Solves a system of linear equations
A * X = B, A**T * X = B, or A**H * X = B
with a general N-by-N matrix A using the LU factorization computed by ZGETRF_GPU.
Arguments
---------
@param[in]
trans magma_trans_t
Specifies the form of the system of equations:
- = MagmaNoTrans: A * X = B (No transpose)
- = MagmaTrans: A**T * X = B (Transpose)
- = MagmaConjTrans: A**H * X = B (Conjugate transpose)
@param[in]
n INTEGER
The order of the matrix A. 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]
dA 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.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in]
ipiv INTEGER array, dimension (N)
The pivot indices from ZGETRF; for 1 <= i <= N, row i of the
matrix was interchanged with row IPIV(i).
@param[in,out]
dB COMPLEX_16 array on the GPU, dimension (LDB,NRHS)
On entry, the right hand side matrix B.
On exit, the solution matrix X.
@param[in]
lddb INTEGER
The leading dimension of the array B. LDB >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_zgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_zgetrs_gpu(
magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex_ptr dA, magma_int_t ldda, magma_int_t *ipiv,
magmaDoubleComplex_ptr dB, magma_int_t lddb,
magma_int_t *info)
{
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex *work = NULL;
int notran = (trans == MagmaNoTrans);
magma_int_t i1, i2, inc;
*info = 0;
if ( (! notran) &&
(trans != MagmaTrans) &&
(trans != 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 == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
i1 = 1;
i2 = n;
if (notran) {
inc = 1;
/* Solve A * X = B. */
magma_zgetmatrix( n, nrhs, dB, lddb, work, n );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, n, dB, lddb );
if ( nrhs == 1) {
magma_ztrsv(MagmaLower, MagmaNoTrans, MagmaUnit, n, dA, ldda, dB, 1 );
magma_ztrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
} else {
magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
}
} else {
inc = -1;
/* Solve A**T * X = B or A**H * X = B. */
if ( nrhs == 1) {
magma_ztrsv(MagmaUpper, trans, MagmaNonUnit, n, dA, ldda, dB, 1 );
magma_ztrsv(MagmaLower, trans, MagmaUnit, n, dA, ldda, dB, 1 );
} else {
magma_ztrsm(MagmaLeft, MagmaUpper, trans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
magma_ztrsm(MagmaLeft, MagmaLower, trans, MagmaUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
}
magma_zgetmatrix( n, nrhs, dB, lddb, work, n );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, n, dB, lddb );
}
magma_free_cpu(work);
return *info;
}
extern "C" magma_int_t
magma_zgeqrf2_2q_gpu(
magma_int_t m, magma_int_t n,
magmaDoubleComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magmaDoubleComplex *tau,
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
=======
ZGEQRF computes a QR factorization of a complex M-by-N matrix A:
A = Q * R.
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. N >= 0.
dA (input/output) COMPLEX_16 array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix dA.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
TAU (output) COMPLEX_16 array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
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.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
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-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
#define dA(a_1,a_2) dA, (dA_offset + (a_1) + (a_2)*(ldda))
#define work_ref(a_1) ( work + (a_1))
#define hwork ( work + (nb)*(m))
magmaDoubleComplex_ptr dwork;
magmaDoubleComplex *work;
magma_int_t i, k, ldwork, lddwork, old_i, old_ib, rows;
magma_int_t nbmin, nx, ib, nb;
magma_int_t lhwork, lwork;
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
k = min(m,n);
if (k == 0)
return MAGMA_SUCCESS;
nb = magma_get_zgeqrf_nb(m);
lwork = (m+n) * nb;
lhwork = lwork - (m)*nb;
if ( MAGMA_SUCCESS != magma_zmalloc( &dwork, n*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/*
if ( MAGMA_SUCCESS != magma_zmalloc_cpu( &work, lwork ) ) {
*info = MAGMA_ERR_HOST_ALLOC;
magma_free( dwork );
return *info;
}
*/
cl_mem buffer = clCreateBuffer(gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, sizeof(magmaDoubleComplex)*lwork, NULL, NULL);
work = (magmaDoubleComplex*)clEnqueueMapBuffer(queues[0], buffer, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, lwork*sizeof(magmaDoubleComplex), 0, NULL, NULL, NULL);
nbmin = 2;
nx = 2*nb;
ldwork = m;
lddwork= n;
if (nb >= nbmin && nb < k && nx < k) {
/* Use blocked code initially */
old_i = 0; old_ib = nb;
for (i = 0; i < k-nx; i += nb) {
ib = min(k-i, nb);
rows = m -i;
magma_zgetmatrix_async(rows, ib, dA(i, i), ldda, work_ref(i), ldwork, queues[0], NULL);
clFlush(queues[0]);
if (i>0){
/* Apply H' to A(i:m,i+2*ib:n) from the left */
magma_zlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, n-old_i-2*old_ib, old_ib,
dA(old_i, old_i ), ldda, dwork,0, lddwork,
dA(old_i, old_i+2*old_ib), ldda, dwork,old_ib, lddwork, queues[1]);
magma_zsetmatrix_async( old_ib, old_ib, work_ref(old_i), ldwork,
dA(old_i, old_i), ldda, queues[1], NULL);
clFlush(queues[1]);
}
magma_queue_sync(queues[0]);
lapackf77_zgeqrf(&rows, &ib, work_ref(i), &ldwork, tau+i, hwork, &lhwork, info);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_zlarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib,
work_ref(i), &ldwork, tau+i, hwork, &ib);
zpanel_to_q( MagmaUpper, ib, work_ref(i), ldwork, hwork+ib*ib );
magma_zsetmatrix( rows, ib, work_ref(i), ldwork, dA(i,i), ldda, queues[0]);
zq_to_panel( MagmaUpper, ib, work_ref(i), ldwork, hwork+ib*ib );
if (i + ib < n)
{
magma_zsetmatrix( ib, ib, hwork, ib, dwork, 0, lddwork, queues[1]);
if (i+nb < k-nx){
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left */
magma_zlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dA(i, i ), ldda, dwork,0, lddwork,
dA(i, i+ib), ldda, dwork,ib, lddwork, queues[1]);
magma_queue_sync(queues[1]);
}else {
magma_zlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n-i-ib, ib,
dA(i, i ), ldda, dwork,0, lddwork,
dA(i, i+ib), ldda, dwork,ib, lddwork, queues[1]);
magma_zsetmatrix( ib, ib, work_ref(i), ldwork, dA(i,i), ldda, queues[1]);
clFlush(queues[1]);
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
magma_free(dwork);
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
rows = m-i;
magma_zgetmatrix( rows, ib, dA(i, i), ldda, work, rows, queues[0]);
lhwork = lwork - rows*ib;
lapackf77_zgeqrf(&rows, &ib, work, &rows, tau+i, work+ib*rows, &lhwork, info);
magma_zsetmatrix( rows, ib, work, rows, dA(i, i), ldda, queues[0]);
}
clEnqueueUnmapMemObject(queues[0], buffer, work, 0, NULL, NULL);
clReleaseMemObject(buffer);
// magma_free_cpu(work);
return *info;
} /* magma_zgeqrf2_gpu */
extern "C" magma_int_t
magma_zgetrf_mgpu(magma_int_t num_gpus,
magma_int_t m, magma_int_t n,
cuDoubleComplex **d_lA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
ZGETRF 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) COMPLEX_16 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 inAT(id,i,j) (d_lAT[(id)] + (i)*nb*lddat + (j)*nb)
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t iinfo, nb, n_local[MagmaMaxGPUs];
magma_int_t maxm, mindim;
magma_int_t i, j, d, rows, cols, s, lddat, lddwork;
magma_int_t id, i_local, i_local2, nb0, nb1;
cuDoubleComplex *d_lAT[MagmaMaxGPUs];
cuDoubleComplex *d_panel[MagmaMaxGPUs], *work;
cudaStream_t streaml[4][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;
/* Function Body */
mindim = min(m, n);
nb = magma_get_zgetrf_nb(m);
if (nb <= 1 || nb >= n) {
/* Use CPU code. */
magma_zmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_zgetmatrix( m, n, d_lA[0], ldda, work, m );
lapackf77_zgetrf(&m, &n, work, &m, ipiv, info);
magma_zsetmatrix( m, n, work, m, d_lA[0], ldda );
magma_free_cpu(work);
} else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
if( num_gpus > ceil((double)n/nb) ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) num_gpus );
*info = -1;
return *info;
}
/* allocate workspace for each GPU */
lddat = ((((((n+nb-1)/nb)/num_gpus)*nb)+31)/32)*32;
lddat = (n+nb-1)/nb; /* number of block columns */
lddat = (lddat+num_gpus-1)/num_gpus; /* number of block columns per GPU */
lddat = nb*lddat; /* number of columns per GPU */
lddat = ((lddat+31)/32)*32; /* make it a multiple of 32 */
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
/* local-n and local-ld */
n_local[i] = ((n/nb)/num_gpus)*nb;
if (i < (n/nb)%num_gpus)
n_local[i] += nb;
else if (i == (n/nb)%num_gpus)
n_local[i] += n%nb;
/* workspaces */
if (MAGMA_SUCCESS != magma_zmalloc( &d_panel[i], 3*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_zmalloc( &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;
}
/* create the streams */
magma_queue_create( &streaml[i][0] );
magma_queue_create( &streaml[i][1] );
magmablasSetKernelStream(streaml[i][1]);
magmablas_ztranspose2( d_lAT[i], lddat, d_lA[i], ldda, m, n_local[i] );
}
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
cudaStreamSynchronize(streaml[i][0]);
magmablasSetKernelStream(NULL);
}
magma_setdevice(0);
/* cpu workspace */
lddwork = maxm;
if (MAGMA_SUCCESS != magma_zmalloc_pinned( &work, lddwork*nb*num_gpus )) {
for(i=0; i<num_gpus; 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 streams */
//magma_zgetrf1_mgpu( num_gpus, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
// (cudaStream_t **)streaml, info );
magma_zgetrf2_mgpu(num_gpus, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
streaml, info);
/* clean up */
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
/* save on output */
magmablas_ztranspose2( d_lA[d], ldda, d_lAT[d], lddat, n_local[d], m );
magma_device_sync();
magma_free( d_lAT[d] );
magma_free( d_panel[d] );
magma_queue_destroy( streaml[d][0] );
magma_queue_destroy( streaml[d][1] );
magmablasSetKernelStream(NULL);
} /* end of for d=1,..,num_gpus */
magma_setdevice(0);
magma_free_pinned( work );
}
return *info;
/* End of MAGMA_ZGETRF_MGPU */
}
extern "C" magma_int_t
magma_zcposv_gpu(char uplo, 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,
magmaDoubleComplex *dworkd, magmaFloatComplex *dworks,
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
=======
ZCPOSV computes the solution to a complex system of linear equations
A * X = B,
where A is an N-by-N Hermitian positive definite matrix and X and B
are N-by-NRHS matrices.
ZCPOSV 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
=========
UPLO (input) CHARACTER
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The number of linear equations, i.e., 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.
dA (input or input/output) COMPLEX_16 array on the GPU, dimension (LDDA,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 iterative refinement has been successfully used
(INFO.EQ.0 and ITER.GE.0, see description below), then A is
unchanged, if double factorization has been used
(INFO.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the factor U or L from the Cholesky
factorization A = U**T*U or A = L*L**T.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
dB (input) COMPLEX_16 array on the GPU, dimension (LDDB,NRHS)
The N-by-NRHS right hand side matrix B.
LDDB (input) INTEGER
The leading dimension of the array dB. LDDB >= max(1,N).
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).
dworkd (workspace) COMPLEX_16 array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
dworks (workspace) COMPLEX array on the GPU, dimension (N*(N+NRHS))
This array is used to store the complex single precision matrix
and the right-hand sides or solutions in single precision.
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 SPOTRF
-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
> 0: if INFO = i, the leading minor of order i of (DOUBLE
PRECISION) A is not positive definite, so the
factorization could not be completed, and the solution
has not been computed.
===================================================================== */
#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 *dR;
magmaFloatComplex *dSA, *dSX;
magmaDoubleComplex Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddsx, lddr;
/* Check arguments */
*iter = 0;
*info = 0;
if ( n < 0 )
*info = -1;
else if ( nrhs < 0 )
*info = -2;
else if ( ldda < max(1,n))
*info = -4;
else if ( lddb < max(1,n))
*info = -7;
else if ( lddx < max(1,n))
*info = -9;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
lddsa = n;
lddsx = n;
lddr = n;
dSA = dworks;
dSX = dSA + lddsa*n;
dR = dworkd;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_zlanhe('I', uplo, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
magmablas_zlag2c( n, nrhs, dB, lddb, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_zlat2c( uplo, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_cpotrf_gpu( uplo, n, dSA, lddsa, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// solve dSA*dSX = dB in single precision
magma_cpotrs_gpu( uplo, n, nrhs, dSA, lddsa, dSX, lddsx, info );
// residual dR = dB - dA*dX in double precision
magmablas_clag2z( n, nrhs, dSX, lddsx, dX, lddx, info );
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_zhemv( uplo, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zhemm( MagmaLeft, uplo, n, nrhs,
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 ( n, 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;
return *info;
REFINEMENT:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
magmablas_zlag2c( n, nrhs, dR, lddr, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// solve dSA*dSX = R in single precision
magma_cpotrs_gpu( uplo, n, nrhs, dSA, lddsa, dSX, lddsx, info );
// Add correction and setup residual
// dX += dSX [including conversion] --and--
// dR = dB
for( j=0; j < nrhs; j++ ) {
magmablas_zcaxpycp( n, dSX(0,j), dX(0,j), dB(0,j), dR(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_zhemv( uplo, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zhemm( MagmaLeft, uplo, n, nrhs,
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 ( n, 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;
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_zpotrf_gpu( uplo, n, dA, ldda, info );
if (*info == 0) {
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
magma_zpotrs_gpu( uplo, n, nrhs, dA, ldda, dX, lddx, info );
}
return *info;
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
magmaDoubleComplex c_neg_one = MAGMA_Z_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 };
magmaDoubleComplex *A, *B, *C, *C2, *LU;
magmaDoubleComplex *dA, *dB, *dC1, *dC2;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.5, 0.1 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double 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_zmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_zmalloc( &dA, size ); assert( err == 0 );
err = magma_zmalloc( &dB, size ); assert( err == 0 );
err = magma_zmalloc( &dC1, size ); assert( err == 0 );
err = magma_zmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_zlarnv( &ione, ISEED, &size, A );
lapackf77_zlarnv( &ione, ISEED, &size, B );
lapackf77_zlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test ZSWAP
// swap columns 2 and 3 of dA, then copy to C2 and compare with A
if ( n >= 3 ) {
magma_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetmatrix( m, n, A, ld, dB, ld );
magma_zswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_zswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dA, 1, dB, 1 );
magma_zgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "zswap diff %.2g\n", error );
}
else {
printf( "zswap skipped for n < 3\n" );
}
// ----- test IZAMAX
// get argmax of column of A
magma_zsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_izamax( m, dA(0,j), 1 );
int i2; // NOT magma_int_t, for cublas
cublasIzamax( handle, m, dA(0,j), 1, &i2 );
// todo need sync here?
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "izamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test ZGEMV
// 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_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetvector( maxn, B, 1, dB, 1 );
magma_zsetvector( maxn, C, 1, dC1, 1 );
magma_zsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZgemv( 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);
cublasZaxpy( handle, size, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMV( m, n ) / 1e9;
printf( "zgemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZHEMV
// 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_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetvector( m, B, 1, dB, 1 );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zhemv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZhemv( 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
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMV( m ) / 1e9;
printf( "zhemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZTRSV
// 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_zlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_zgetrf( &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_zsetmatrix( m, m, LU, ld, dA, ld );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ztrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZtrsv( 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
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "ztrsv( %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 ZGEMM
// 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_zsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_zsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zgemm( 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 );
cublasZgemm( 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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMM( m, n, k ) / 1e9;
printf( "zgemm( %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 ZHEMM
// 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_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetmatrix( m, n, B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zhemm( 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 );
cublasZhemm( 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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMM( side[is], m, n ) / 1e9;
printf( "zhemm( %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 ZHERK
// 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_zsetmatrix( n, k, A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zherk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZherk( 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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHERK( k, n ) / 1e9;
printf( "zherk( %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 ZHER2K
// 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_zsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zher2k( 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 );
cublasZher2k( 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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHER2K( k, n ) / 1e9;
printf( "zher2k( %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 ZTRMM
// 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_zsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrmm( 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 );
cublasZtrmm( 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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRMM( side[is], m, n ) / 1e9;
printf( "ztrmm( %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 ZTRSM
// 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_zsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrsm( 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 );
cublasZtrsm( 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
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( side[is], m, n ) / 1e9;
printf( "ztrsm( %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;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zlarfb_gpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t M, N, K, size, ldc, ldv, ldt, ldw, nv;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double error, work[1];
// test all combinations of input parameters
const char side[] = { MagmaLeft, MagmaRight };
const char trans[] = { MagmaConjTrans, MagmaNoTrans };
const char direct[] = { MagmaForward, MagmaBackward };
const char storev[] = { MagmaColumnwise, MagmaRowwise };
magma_opts opts;
parse_opts( argc, argv, &opts );
printf(" M N K storev side direct trans ||R||_F / ||HC||_F\n");
printf("========================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
M = opts.msize[i];
N = opts.nsize[i];
K = opts.ksize[i];
if ( M < K || N < K || K <= 0 ) {
printf( "skipping M %d, N %d, K %d; requires M >= K, N >= K, K >= 0.\n", (int) M, (int) N, (int) K );
continue;
}
for( int istor = 0; istor < 2; ++istor ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int idir = 0; idir < 2; ++idir ) {
for( int itran = 0; itran < 2; ++itran ) {
ldc = ((M+31)/32)*32;
ldt = ((K+31)/32)*32;
ldw = (side[iside] == MagmaLeft ? N : M);
// (ldv, nv) get swapped later if rowwise
ldv = (side[iside] == MagmaLeft ? M : N);
nv = K;
// Allocate memory for matrices
magmaDoubleComplex *C, *R, *V, *T, *W;
TESTING_MALLOC( C, magmaDoubleComplex, ldc*N );
TESTING_MALLOC( R, magmaDoubleComplex, ldc*N );
TESTING_MALLOC( V, magmaDoubleComplex, ldv*K );
TESTING_MALLOC( T, magmaDoubleComplex, ldt*K );
TESTING_MALLOC( W, magmaDoubleComplex, ldw*K );
magmaDoubleComplex *dC, *dV, *dT, *dW;
TESTING_DEVALLOC( dC, magmaDoubleComplex, ldc*N );
TESTING_DEVALLOC( dV, magmaDoubleComplex, ldv*K );
TESTING_DEVALLOC( dT, magmaDoubleComplex, ldt*K );
TESTING_DEVALLOC( dW, magmaDoubleComplex, ldw*K );
// C is M x N.
size = ldc*N;
lapackf77_zlarnv( &ione, ISEED, &size, C );
//printf( "C=" ); magma_zprint( M, N, C, ldc );
// V is ldv x nv. See larfb docs for description.
// if column-wise and left, M x K
// if column-wise and right, N x K
// if row-wise and left, K x M
// if row-wise and right, K x N
size = ldv*nv;
lapackf77_zlarnv( &ione, ISEED, &size, V );
if ( storev[istor] == MagmaColumnwise ) {
if ( direct[idir] == MagmaForward ) {
lapackf77_zlaset( MagmaUpperStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_zlaset( MagmaLowerStr, &K, &K, &c_zero, &c_one, &V[(ldv-K)], &ldv );
}
}
else {
// rowwise, swap V's dimensions
std::swap( ldv, nv );
if ( direct[idir] == MagmaForward ) {
lapackf77_zlaset( MagmaLowerStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_zlaset( MagmaUpperStr, &K, &K, &c_zero, &c_one, &V[(nv-K)*ldv], &ldv );
}
}
//printf( "# ldv %d, nv %d\n", ldv, nv );
//printf( "V=" ); magma_zprint( ldv, nv, V, ldv );
// T is K x K, upper triangular for forward, and lower triangular for backward
magma_int_t k1 = K-1;
size = ldt*K;
lapackf77_zlarnv( &ione, ISEED, &size, T );
if ( direct[idir] == MagmaForward ) {
lapackf77_zlaset( MagmaLowerStr, &k1, &k1, &c_zero, &c_zero, &T[1], &ldt );
}
else {
lapackf77_zlaset( MagmaUpperStr, &k1, &k1, &c_zero, &c_zero, &T[1*ldt], &ldt );
}
//printf( "T=" ); magma_zprint( K, K, T, ldt );
magma_zsetmatrix( M, N, C, ldc, dC, ldc );
magma_zsetmatrix( ldv, nv, V, ldv, dV, ldv );
magma_zsetmatrix( K, K, T, ldt, dT, ldt );
lapackf77_zlarfb( &side[iside], &trans[itran], &direct[idir], &storev[istor],
&M, &N, &K,
V, &ldv, T, &ldt, C, &ldc, W, &ldw );
//printf( "HC=" ); magma_zprint( M, N, C, ldc );
magma_zlarfb_gpu( side[iside], trans[itran], direct[idir], storev[istor],
M, N, K,
dV, ldv, dT, ldt, dC, ldc, dW, ldw );
magma_zgetmatrix( M, N, dC, ldc, R, ldc );
//printf( "dHC=" ); magma_zprint( M, N, R, ldc );
// compute relative error |HC_magma - HC_lapack| / |HC_lapack|
error = lapackf77_zlange( "Fro", &M, &N, C, &ldc, work );
size = ldc*N;
blasf77_zaxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_zlange( "Fro", &M, &N, R, &ldc, work ) / error;
printf( "%5d %5d %5d %c %c %c %c %8.2e\n",
(int) M, (int) N, (int) K,
storev[istor], side[iside], direct[idir], trans[itran], error );
TESTING_FREE( C );
TESTING_FREE( R );
TESTING_FREE( V );
TESTING_FREE( T );
TESTING_FREE( W );
TESTING_DEVFREE( dC );
TESTING_DEVFREE( dV );
TESTING_DEVFREE( dT );
TESTING_DEVFREE( dW );
}}}}
printf( "\n" );
}
TESTING_FINALIZE();
return 0;
}
extern "C" magma_int_t
magma_zunmqr_gpu(
magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaDoubleComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magmaDoubleComplex *tau,
magmaDoubleComplex_ptr dC, size_t dC_offset, magma_int_t lddc,
magmaDoubleComplex *hwork, magma_int_t lwork,
magmaDoubleComplex_ptr dT, size_t dT_offset, magma_int_t nb,
magma_queue_t queue,
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
=======
ZUNMQR_GPU overwrites the general complex M-by-N matrix C with
SIDE = 'L' SIDE = 'R'
TRANS = 'N': Q * C C * Q
TRANS = 'T': Q**H * C C * Q**H
where Q is a complex orthogonal matrix defined as the product of k
elementary reflectors
Q = H(1) H(2) . . . H(k)
as returned by ZGEQRF. Q is of order M if SIDE = 'L' and of order N
if SIDE = 'R'.
Arguments
=========
SIDE (input) CHARACTER*1
= 'L': apply Q or Q**H from the Left;
= 'R': apply Q or Q**H from the Right.
TRANS (input) CHARACTER*1
= 'N': No transpose, apply Q;
= 'T': Transpose, apply Q**H.
M (input) INTEGER
The number of rows of the matrix C. M >= 0.
N (input) INTEGER
The number of columns of the matrix C. N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = 'L', M >= K >= 0;
if SIDE = 'R', N >= K >= 0.
DA (input) COMPLEX_16 array on the GPU, dimension (LDDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
ZGEQRF in the first k columns of its array argument DA.
DA is modified by the routine but restored on exit.
LDDA (input) INTEGER
The leading dimension of the array DA.
If SIDE = 'L', LDDA >= max(1,M);
if SIDE = 'R', LDDA >= max(1,N).
TAU (input) COMPLEX_16 array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by ZGEQRF.
DC (input/output) COMPLEX_16 array on the GPU, dimension (LDDC,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.
LDDC (input) INTEGER
The leading dimension of the array DC. LDDC >= max(1,M).
HWORK (workspace/output) COMPLEX_16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, HWORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array HWORK.
LWORK >= (M-K+NB)*(N+2*NB) if SIDE = 'L',
and LWORK >= (N-K+NB)*(M+2*NB) if SIDE = 'R', where NB is the
optimal blocksize.
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, and no error
message related to LWORK is issued by XERBLA.
DT (input) COMPLEX_16 array on the GPU that is the output
(the 9th argument) of magma_zgeqrf_gpu.
NB (input) INTEGER
This is the blocking size that was used in pre-computing DT, e.g.,
the blocking size used in magma_zgeqrf_gpu.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
===================================================================== */
#define a_ref(a_1,a_2) dA, (dA_offset+(a_1)+(a_2)*(ldda))
#define c_ref(a_1,a_2) dC, (dC_offset+(a_1)+(a_2)*(lddc))
#define t_ref(a_1) dT, (dT_offset+(a_1)*nb)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex_ptr dwork;
magma_int_t i, lddwork;
magma_int_t i1, i2, i3, ib, ic, jc, mi, ni, nq, nw, ret;
int left, notran, lquery;
magma_int_t lwkopt;
*info = 0;
left = (side == MagmaLeft);
notran = (trans == MagmaNoTrans);
lquery = (lwork == -1);
if (!left || notran)
printf("zunmqr_gpu called with arguments not yet supported\n");
/* 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;
}
if ( (!left) && (side != MagmaRight) ) {
*info = -1;
} else if ( (!notran) && (trans != MagmaConjTrans) ) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (ldda < max(1,nq)) {
*info = -7;
} else if (lddc < max(1,m)) {
*info = -10;
} else if (lwork < max(1,nw) && ! lquery) {
*info = -12;
}
lwkopt = (m-k+nb)*(n+2*nb);
hwork[0] = MAGMA_Z_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) {
hwork[0] = c_one;
return *info;
}
lddwork= k;
dwork = dT;
size_t dwork_offset = 2*lddwork*nb;
if ( (left && (! notran)) || ( (!left) && notran ) ) {
i1 = 0;
i2 = k-nb;
i3 = nb;
} else {
i1 = (k - 1 - nb) / nb * nb;
i2 = 0;
i3 = -nb;
}
if (left) {
ni = n;
jc = 0;
} else {
mi = m;
ic = 0;
}
if (nb < k)
{
for (i=i1; i3<0 ? i>i2 : i<i2; i+=i3)
{
ib = min(nb, k - i);
if (left){
mi = m - i;
ic = i;
}
else {
ni = n - i;
jc = i;
}
ret = magma_zlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
mi, ni, ib,
a_ref(i, i ), ldda, t_ref(i), nb,
c_ref(ic, jc), lddc, dwork, dwork_offset, nw, queue);
if ( ret != MAGMA_SUCCESS )
return ret;
}
}
else
{
i = i1;
}
/* Use unblocked code to multiply the last or only block. */
if (i < k) {
ib = k-i;
if (left){
mi = m - i;
ic = i;
}
else {
ni = n - i;
jc = i;
}
magma_zgetmatrix(mi, ib, a_ref(i, i), ldda, hwork, mi, queue);
magma_zgetmatrix(mi, ni, c_ref(ic, jc), lddc, hwork+mi*ib, mi, queue);
magma_int_t lhwork = lwork - mi*(ib + ni);
lapackf77_zunmqr( MagmaLeftStr, MagmaConjTransStr,
&mi, &ni, &ib,
hwork, &mi, tau+i,
hwork+mi*ib, &mi,
hwork+mi*(ib+ni), &lhwork, info);
// send the updated part of c back to the GPU
magma_zsetmatrix(mi, ni, hwork+mi*ib, mi, c_ref(ic, jc), lddc, queue);
}
return *info;
/* End of MAGMA_ZUNMQR_GPU */
}
extern "C" magma_int_t
magma_zhegst(magma_int_t itype, char uplo, magma_int_t n,
magmaDoubleComplex *a, magma_int_t lda,
magmaDoubleComplex *b, magma_int_t ldb, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZHEGST reduces a complex Hermitian-definite generalized
eigenproblem to standard form.
If ITYPE = 1, the problem is A*x = lambda*B*x,
and A is overwritten by inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H)
If ITYPE = 2 or 3, the problem is A*B*x = lambda*x or
B*A*x = lambda*x, and A is overwritten by U*A*U**H or L**H*A*L.
B must have been previously factorized as U**H*U or L*L**H by ZPOTRF.
Arguments
=========
ITYPE (input) INTEGER
= 1: compute inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H);
= 2 or 3: compute U*A*U**H or L**H*A*L.
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored and B is factored as
U**H*U;
= 'L': Lower triangle of A is stored and B is factored as
L*L**H.
N (input) INTEGER
The order of the matrices A and B. N >= 0.
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 INFO = 0, the transformed matrix, stored in the
same format as A.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
B (input) COMPLEX_16 array, dimension (LDB,N)
The triangular factor from the Cholesky factorization of B,
as returned by ZPOTRF.
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
=====================================================================*/
char uplo_[2] = {uplo, 0};
magma_int_t nb;
magma_int_t k, kb, kb2;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_half = MAGMA_Z_HALF;
magmaDoubleComplex c_neg_half = MAGMA_Z_NEG_HALF;
magmaDoubleComplex *dw;
magma_int_t ldda = n;
magma_int_t lddb = n;
double d_one = 1.0;
int upper = lapackf77_lsame(uplo_, "U");
/* Test the input parameters. */
*info = 0;
if (itype<1 || itype>3){
*info = -1;
}else if ((! upper) && (! lapackf77_lsame(uplo_, "L"))) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
}else if (ldb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
if (MAGMA_SUCCESS != magma_zmalloc( &dw, 2*n*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
nb = magma_get_zhegst_nb(n);
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_zsetmatrix( n, n, A(0, 0), lda, dA(0, 0), ldda );
magma_zsetmatrix( n, n, B(0, 0), ldb, dB(0, 0), lddb );
/* Use hybrid blocked code */
if (itype==1) {
if (upper) {
/* Compute inv(U')*A*inv(U) */
for(k = 0; k<n; k+=nb){
kb = min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the upper triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[0] );
if(k+kb<n){
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k,k), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[0] );
magma_zhemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_zher2k(MagmaUpper, MagmaConjTrans,
n-k-kb, kb,
c_neg_one, dA(k,k+kb), ldda,
dB(k,k+kb), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, stream[1] );
magma_zhemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_ztrsm(MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
kb, n-k-kb,
c_one ,dB(k+kb,k+kb), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[1] );
}
}
magma_queue_sync( stream[0] );
}
else {
/* Compute inv(L)*A*inv(L') */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the lower triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[0] );
if(k+kb<n){
magma_ztrsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k,k), lddb,
dA(k+kb,k), ldda);
magma_queue_sync( stream[0] );
magma_zhemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_zher2k(MagmaLower, MagmaNoTrans,
n-k-kb, kb,
c_neg_one, dA(k+kb,k), ldda,
dB(k+kb,k), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, stream[1] );
magma_zhemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k+kb,k), ldda);
}
magma_queue_sync( stream[1] );
}
}
magma_queue_sync( stream[0] );
}
else {
if (upper) {
/* Compute U*A*U' */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, stream[0] );
/* Update the upper triangle of A(1:k+kb-1,1:k+kb-1) */
if(k>0){
magma_ztrmm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
k, kb,
c_one ,dB(0,0), lddb,
dA(0,k), ldda);
magma_zhemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_queue_sync( stream[1] );
magma_zher2k(MagmaUpper, MagmaNoTrans,
k, kb,
c_one, dA(0,k), ldda,
dB(0,k), lddb,
d_one, dA(0,0), ldda);
magma_zhemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_ztrmm(MagmaRight, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
k, kb,
c_one, dB(k,k), lddb,
dA(0,k), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_zhegst( &itype, uplo_, &kb, A(k, k), &lda, B(k, k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
else {
/* Compute L'*A*L */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, stream[0] );
/* Update the lower triangle of A(1:k+kb-1,1:k+kb-1) */
if(k>0){
magma_ztrmm(MagmaRight, MagmaLower, MagmaNoTrans, MagmaNonUnit,
kb, k,
c_one ,dB(0,0), lddb,
dA(k,0), ldda);
magma_zhemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_queue_sync( stream[1] );
magma_zher2k(MagmaLower, MagmaConjTrans,
k, kb,
c_one, dA(k,0), ldda,
dB(k,0), lddb,
d_one, dA(0,0), ldda);
magma_zhemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_ztrmm(MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit,
kb, k,
c_one, dB(k,k), lddb,
dA(k,0), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
}
magma_zgetmatrix( n, n, dA(0, 0), ldda, A(0, 0), lda );
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dw );
return *info;
} /* magma_zhegst_gpu */
/**
Purpose
-------
ZUNGQR generates an M-by-N COMPLEX_16 matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by ZGEQRF_GPU.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix Q. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix Q. M >= N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
@param[in,out]
dA COMPLEX_16 array A on the GPU, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by ZGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
@param[in]
ldda INTEGER
The first dimension of the array A. LDDA >= max(1,M).
@param[in]
tau COMPLEX_16 array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by ZGEQRF_GPU.
@param[in]
dT (workspace) COMPLEX_16 work space array on the GPU,
dimension (2*MIN(M, N) + ceil(N/32)*32 )*NB.
This must be the 6th argument of magma_zgeqrf_gpu
[ note that if N here is bigger than N in magma_zgeqrf_gpu,
the workspace requirement DT in magma_zgeqrf_gpu must be
as specified in this routine ].
@param[in]
nb INTEGER
This is the block size used in ZGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_zgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_zungqr_gpu(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex *tau,
magmaDoubleComplex_ptr dT, magma_int_t nb,
magma_int_t *info)
{
#define dA(i,j) (dA + (i) + (j)*ldda)
#define dT(j) (dT + (j)*nb)
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, lpanel;
magma_int_t i, ib, ki, kk, iinfo;
magma_int_t lddwork;
magmaDoubleComplex_ptr dV, dW;
magmaDoubleComplex *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 < 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 = 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_zmalloc_cpu( &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_zmalloc( &dV, ldda*nb )) {
magma_free_cpu( work );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
// dT workspace has:
// 2*min(m,n)*nb for T and R^{-1} matrices from geqrf
// roundup(n,32) * nb for dW larfb workspace.
lddwork = min(m,n);
dW = dT + 2*lddwork*nb;
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
// 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_zgetmatrix( m_kk, k_kk,
dA(kk, kk), ldda, panel, m_kk, queue );
lapackf77_zungqr( &m_kk, &n_kk, &k_kk,
panel, &m_kk,
&tau[kk], work, &lwork, &iinfo );
magma_zsetmatrix( m_kk, n_kk,
panel, m_kk, dA(kk, kk), ldda, queue );
// Set A(1:kk,kk+1:n) to zero.
magmablas_zlaset( MagmaFull, kk, n - kk, c_zero, c_zero, dA(0, kk), ldda, queue );
}
if (kk > 0) {
// Use blocked code
// queue: copy Aii to V --> laset --> laset --> larfb --> [next]
// CPU has no computation
for (i = ki; i >= 0; i -= nb) {
ib = min( nb, k-i );
mi = m - i;
// Copy current panel on the GPU from dA to dV
magma_zcopymatrix_async( mi, ib,
dA(i,i), ldda,
dV, ldda, queue );
// set panel to identity
magmablas_zlaset( MagmaFull, i, ib, c_zero, c_zero, dA(0, i), ldda, queue );
magmablas_zlaset( 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_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT(i), nb,
dA(i, i), ldda, dW, lddwork, queue );
}
}
}
magma_queue_sync( queue );
magma_free( dV );
magma_free_cpu( work );
magma_queue_destroy( queue );
return *info;
} /* magma_zungqr_gpu */
/**
Purpose
-------
ZHETRD reduces a complex Hermitian matrix A to real symmetric
tridiagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
Arguments
---------
@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]
A COMPLEX_16 array, dimension (LDA,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, 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 diagonal and first superdiagonal
of A are overwritten by the corresponding elements of the
tridiagonal matrix T, and the elements above the first
superdiagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the diagonal and first subdiagonal of A are over-
written by the corresponding elements of the tridiagonal
matrix T, 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]
d COMPLEX_16 array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
@param[out]
e COMPLEX_16 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]
tau COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@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 dimension of the array WORK. LWORK >= N*NB, where NB is the
optimal blocksize given by magma_get_zhetrd_nb().
\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
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(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+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-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 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_zheev_comp
********************************************************************/
extern "C" magma_int_t
magma_zhetrd(
magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda,
double *d, double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i_, j_) ( A + (i_) + (j_)*lda )
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t ldda = roundup( n, 32 );
magma_int_t nb = magma_get_zhetrd_nb( n );
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const double d_one = MAGMA_D_ONE;
magma_int_t kk, nx;
magma_int_t i, j, i_n;
magma_int_t iinfo;
magma_int_t ldw, lddw, lwkopt;
magma_int_t lquery;
*info = 0;
int upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < nb*n && ! lquery) {
*info = -9;
}
/* Determine the block size. */
ldw = n;
lddw = ldda;
lwkopt = n * nb;
if (*info == 0) {
work[0] = MAGMA_Z_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magmaDoubleComplex *dA;
#ifdef FAST_HEMV
magma_int_t ldwork2 = ldda*ceildiv(n,64);
#else
magma_int_t ldwork2 = 0;
#endif
if (MAGMA_SUCCESS != magma_zmalloc( &dA, ldda*n + 2*lddw*nb + ldwork2 )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmaDoubleComplex *dwork = dA + ldda*n;
#ifdef FAST_HEMV
magmaDoubleComplex *dwork2 = dwork + 2*lddw*nb;
#endif
//if (n < 2048)
// nx = n;
//else
// nx = 512;
nx = min( 128, n ); // nx <= n is required
// clear out dwork in case it has NANs (used as y in zhemv)
// rest of dwork (used as work in magmablas_zhemv) doesn't need to be cleared
magmablas_zlaset( MagmaFull, n, nb, c_zero, c_zero, dwork, lddw );
if (upper) {
/* Copy the matrix to the GPU */
magma_zsetmatrix( n, n, A(0, 0), lda, dA(0, 0), ldda );
/* Reduce the upper triangle of A.
Columns 1:kk are handled by the unblocked method. */
kk = n - (n - nx + nb - 1) / nb * nb;
for (i = n - nb; i >= kk; i -= nb) {
/* Reduce columns i:i+nb-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i != n-nb)
magma_zgetmatrix( i+nb, nb, dA(0, i), ldda, A(0, i), lda );
#ifdef FAST_HEMV
magma_zlatrd2( uplo, i+nb, nb, A(0, 0), lda, e, tau,
work, ldw, dA(0, 0), ldda, dwork, lddw,
dwork2, ldwork2 );
#else
magma_zlatrd( uplo, i+nb, nb, A(0, 0), lda, e, tau,
work, ldw, dA(0, 0), ldda, dwork, lddw );
#endif
/* Update the unreduced submatrix A(0:i-2,0:i-2), using an
update of the form: A := A - V*W' - W*V' */
magma_zsetmatrix( i + nb, nb, work, ldw, dwork, lddw );
magma_zher2k( uplo, MagmaNoTrans, i, nb, c_neg_one,
dA(0, i), ldda, dwork, lddw,
d_one, dA(0, 0), ldda );
/* Copy superdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+nb; ++j) {
*A(j-1,j) = MAGMA_Z_MAKE( e[j - 1], 0 );
d[j] = MAGMA_Z_REAL( *A(j, j) );
}
}
magma_zgetmatrix( kk, kk, dA(0, 0), ldda, A(0, 0), lda );
/* Use CPU code to reduce the last or only block */
lapackf77_zhetrd( uplo_, &kk, A(0, 0), &lda, d, e, tau, work, &lwork, &iinfo );
}
else {
/* Copy the matrix to the GPU */
if (1 <= n-nx)
magma_zsetmatrix( n, n, A(0,0), lda, dA(0,0), ldda );
/* Reduce the lower triangle of A */
for (i = 0; i < n-nx; i += nb) {
/* Reduce columns i:i+nb-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i != 0)
magma_zgetmatrix( n-i, nb, dA(i, i), ldda, A(i, i), lda );
#ifdef FAST_HEMV
magma_zlatrd2( uplo, n-i, nb, A(i, i), lda, &e[i], &tau[i],
work, ldw, dA(i, i), ldda, dwork, lddw,
dwork2, ldwork2 );
#else
magma_zlatrd( uplo, n-i, nb, A(i, i), lda, &e[i], &tau[i],
work, ldw, dA(i, i), ldda, dwork, lddw );
#endif
/* Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V' */
magma_zsetmatrix( n-i, nb, work, ldw, dwork, lddw );
magma_zher2k( MagmaLower, MagmaNoTrans, n-i-nb, nb, c_neg_one,
dA(i+nb, i), ldda, &dwork[nb], lddw,
d_one, dA(i+nb, i+nb), ldda );
/* Copy subdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+nb; ++j) {
*A(j+1,j) = MAGMA_Z_MAKE( e[j], 0 );
d[j] = MAGMA_Z_REAL( *A(j, j) );
}
}
/* Use CPU code to reduce the last or only block */
if (1 <= n-nx)
magma_zgetmatrix( n-i, n-i, dA(i, i), ldda, A(i, i), lda );
i_n = n-i;
lapackf77_zhetrd( uplo_, &i_n, A(i, i), &lda, &d[i], &e[i],
&tau[i], work, &lwork, &iinfo );
}
magma_free( dA );
work[0] = MAGMA_Z_MAKE( lwkopt, 0 );
return *info;
} /* magma_zhetrd */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zcgeqrsv
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time, gpu_perfd, gpu_perfs;
double error, gpu_error, cpu_error, Anorm, work[1];
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_A2, *h_B, *h_X, *h_R;
magmaDoubleComplex_ptr d_A, d_B, d_X, d_T;
magmaFloatComplex *d_SA, *d_SB;
magmaDoubleComplex *h_workd, *tau, tmp[1];
magmaFloatComplex *h_works;
magma_int_t lda, ldb, lhwork, lworkgpu;
magma_int_t ldda, lddb, lddx;
magma_int_t M, N, nrhs, qrsv_iters, info, size, min_mn, max_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
printf("Epsilon(double): %8.6e\n"
"Epsilon(single): %8.6e\n\n",
lapackf77_dlamch("Epsilon"), lapackf77_slamch("Epsilon") );
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
nrhs = opts.nrhs;
printf(" CPU Gflop/s GPU Gflop/s |b-Ax|| / (N||A||) ||dx-x||/(N||A||)\n");
printf(" M N NRHS double double single mixed Iter CPU GPU \n");
printf("=============================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
if ( M < N ) {
printf( "%5d %5d %5d skipping because M < N is not yet supported.\n", (int) M, (int) N, (int) nrhs );
continue;
}
min_mn = min(M, N);
max_mn = max(M, N);
lda = M;
ldb = max_mn;
ldda = ((M+31)/32) * 32;
lddb = ((max_mn+31)/32)*32;
lddx = ((N+31)/32) * 32;
nb = max( magma_get_zgeqrf_nb( M ), magma_get_cgeqrf_nb( M ) );
gflops = (FLOPS_ZGEQRF( M, N ) + FLOPS_ZGEQRS( M, N, nrhs )) / 1e9;
lworkgpu = (M - N + nb)*(nrhs + nb) + nrhs*nb;
// query for workspace size
lhwork = -1;
lapackf77_zgels( MagmaNoTransStr, &M, &N, &nrhs,
NULL, &lda, NULL, &ldb, tmp, &lhwork, &info );
lhwork = (magma_int_t) MAGMA_Z_REAL( tmp[0] );
lhwork = max( lhwork, lworkgpu );
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( h_A2, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( h_B, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_X, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_R, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_workd, magmaDoubleComplex, lhwork );
h_works = (magmaFloatComplex*)h_workd;
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( d_B, magmaDoubleComplex, lddb*nrhs );
TESTING_MALLOC_DEV( d_X, magmaDoubleComplex, lddx*nrhs );
TESTING_MALLOC_DEV( d_T, magmaDoubleComplex, ( 2*min_mn + (N+31)/32*32 )*nb );
/* Initialize the matrices */
size = lda*N;
lapackf77_zlarnv( &ione, ISEED, &size, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_A2, &lda );
// make random RHS
size = ldb*nrhs;
lapackf77_zlarnv( &ione, ISEED, &size, h_B );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
//=====================================================================
// Mixed Precision Iterative Refinement - GPU
//=====================================================================
gpu_time = magma_wtime();
magma_zcgeqrsv_gpu( M, N, nrhs,
d_A, ldda, d_B, lddb,
d_X, lddx, &qrsv_iters, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zcgeqrsv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute the residual
magma_zgetmatrix( N, nrhs, d_X, lddx, h_X, ldb );
blasf77_zgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A, &lda,
h_X, &ldb,
&c_one, h_R, &ldb);
Anorm = lapackf77_zlange("f", &M, &N, h_A, &lda, work);
//=====================================================================
// Double Precision Solve
//=====================================================================
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
gpu_time = magma_wtime();
magma_zgels_gpu( MagmaNoTrans, M, N, nrhs, d_A, ldda,
d_B, lddb, h_workd, lworkgpu, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perfd = gflops / gpu_time;
//=====================================================================
// Single Precision Solve
//=====================================================================
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
/* The allocation of d_SA and d_SB is done here to avoid
* to double the memory used on GPU with zcgeqrsv */
TESTING_MALLOC_DEV( d_SA, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_SB, magmaFloatComplex, lddb*nrhs );
magmablas_zlag2c( M, N, d_A, ldda, d_SA, ldda, &info );
magmablas_zlag2c( N, nrhs, d_B, lddb, d_SB, lddb, &info );
gpu_time = magma_wtime();
magma_cgels_gpu( MagmaNoTrans, M, N, nrhs, d_SA, ldda,
d_SB, lddb, h_works, lhwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perfs = gflops / gpu_time;
TESTING_FREE_DEV( d_SA );
TESTING_FREE_DEV( d_SB );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_X, &ldb );
cpu_time = magma_wtime();
lapackf77_zgels( MagmaNoTransStr, &M, &N, &nrhs,
h_A, &lda, h_X, &ldb, h_workd, &lhwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zgels returned error %d: %s.\n",
(int) info, magma_strerror( info ));
blasf77_zgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A2, &lda,
h_X, &ldb,
&c_one, h_B, &ldb );
cpu_error = lapackf77_zlange("f", &M, &nrhs, h_B, &ldb, work) / (min_mn*Anorm);
gpu_error = lapackf77_zlange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*Anorm);
// error relative to LAPACK
size = M*nrhs;
blasf77_zaxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error = lapackf77_zlange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*Anorm);
printf("%5d %5d %5d %7.2f %7.2f %7.2f %7.2f %4d %8.2e %8.2e %8.2e %s\n",
(int) M, (int) N, (int) nrhs,
cpu_perf, gpu_perfd, gpu_perfs, gpu_perf,
(int) qrsv_iters,
cpu_error, gpu_error, error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_A2 );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( h_R );
TESTING_FREE_CPU( h_workd );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_X );
TESTING_FREE_DEV( d_T );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing magma_zhemm_mgpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex calpha = MAGMA_Z_MAKE( 3.456, 5.678 );
magmaDoubleComplex cbeta = MAGMA_Z_MAKE( 1.234, 2.456 );
real_Double_t gflops, gpu_perf=0., cpu_perf=0., gpu_time=0., cpu_time=0.;
real_Double_t gpu_perf2=0., gpu_time2=0.;
double error=0., errorbis=0., work[1];
magmaDoubleComplex *hA, *hX, *hB, *hR;
magmaDoubleComplex_ptr dA[MagmaMaxGPUs], dX[MagmaMaxGPUs], dB[MagmaMaxGPUs], dwork[MagmaMaxGPUs], hwork[MagmaMaxGPUs+1];
magmaDoubleComplex_ptr dA2;
magma_int_t M, N, size, lda, ldda, msize, nb, nstream;
magma_int_t ione = 1;
magma_int_t iseed[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
// default values
nb = (opts.nb > 0 ? opts.nb : 64);
nstream = (opts.nstream > 0 ? opts.nstream : 2);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t nbcmplx = 0;
magma_buildconnection_mgpu(gnode, &nbcmplx, opts.ngpu);
printf("Initializing communication pattern... GPU-ncmplx %d\n\n", (int) nbcmplx);
for (int i=0; i < nbcmplx; ++i) {
int myngpu = gnode[i][MagmaMaxGPUs];
printf("cmplx %d has %d gpu ", i, myngpu);
for(int j=0; j < myngpu; ++j)
printf(" %d", (int) gnode[i][j]);
printf("\n");
}
magma_int_t nbevents = 2;
magma_queue_t streams[MagmaMaxGPUs][20];
magma_event_t redevents[MagmaMaxGPUs][20];
magma_event_t redevents2[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs+10];
for( int d = 0; d < opts.ngpu; ++d ) {
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_create( &streams[d][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&redevents[d][i], cudaEventDisableTiming);
cudaEventCreateWithFlags(&redevents2[d][i], cudaEventDisableTiming);
}
}
printf( "nb %d, ngpu %d, nstream %d version %d\n", (int) nb, (int) opts.ngpu, (int) nstream, (int) opts.version );
printf(" M N nb offset CPU GFlop/s (sec) GPU GFlop/s (sec) CUBLAS hemm (sec) ||R|| / ||A||*||X||\n");
printf("=========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
for( int offset = 0; offset < N; offset += min(N,nb) ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
msize = M - offset;
lda = M;
ldda = ((M + 31)/32)*32;
size = lda*M;
gflops = FLOPS_ZHEMM( MagmaLeft, (double)msize, (double)N ) / 1e9;
magma_int_t dworksiz = ldda*N*3;
magma_int_t hworksiz = lda*N;
TESTING_MALLOC_CPU( hA, magmaDoubleComplex, lda*M );
TESTING_MALLOC_CPU( hX, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( hB, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( hR, magmaDoubleComplex, lda*N );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_int_t mlocal = ((M / nb) / opts.ngpu + 1) * nb;
magma_setdevice( d );
TESTING_MALLOC_DEV( dA[d], magmaDoubleComplex, ldda*mlocal );
TESTING_MALLOC_DEV( dX[d], magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dB[d], magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dwork[d], magmaDoubleComplex, dworksiz );
TESTING_MALLOC_PIN( hwork[d], magmaDoubleComplex, hworksiz );
}
TESTING_MALLOC_PIN( hwork[opts.ngpu], magmaDoubleComplex, lda*N );
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_MALLOC_DEV( dA2, magmaDoubleComplex, ldda*M );
}
lapackf77_zlarnv( &ione, iseed, &size, hA );
magma_zmake_hermitian( M, hA, lda );
size = lda*N;
lapackf77_zlarnv( &ione, iseed, &size, hX );
lapackf77_zlarnv( &ione, iseed, &size, hB );
lapackf77_zlacpy( "Full", &M, &N, hB, &lda, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix_1D_col_bcyclic( M, M, hA, lda, dA, ldda, opts.ngpu, nb );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
//magmablasSetKernelStream( streams[ d ][ 0 ] );
magma_zsetmatrix( M, N, hX, lda, dX[d], ldda );
//if (d == 0) magma_zsetmatrix( M, N, hB, lda, dB[d], ldda ); // this is wrong coz when offset != 0 the gpu who do the beta*C may be not 0 so this should be related to stdev(starting device who own i=0 first col)
magma_zsetmatrix( M, N, hB, lda, dB[d], ldda );
}
//memset(hR, 0, lda*N*sizeof(magmaDoubleComplex));
trace_init( 1, opts.ngpu, nstream, (magma_queue_t*) streams );
//magma_int_t offset = 0; //nb;
gpu_time = magma_sync_wtime(0);
magmablas_zhemm_mgpu_com(
MagmaLeft, MagmaLower, msize, N,
calpha, dA, ldda, offset,
dX, ldda,
cbeta, dB, ldda, dwork, dworksiz, hR, lda, hwork, hworksiz,
opts.ngpu, nb, streams, nstream, redevents2, nbevents, gnode, nbcmplx);
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
#ifdef TRACING
char buf[80];
snprintf( buf, sizeof(buf), "zhemm-m%d-n%d-nb%d-stream%d-ngpu%d-run%d.svg",
(int) M, (int) N, (int) nb, (int) nstream, (int) opts.ngpu, (int) iter );
trace_finalize( buf, "trace.css" );
#endif
/* ====================================================================
Performs operation using CUBLAS
=================================================================== */
if ( opts.check && iter == 0 ) {
magma_setdevice( 0 );
magmablasSetKernelStream( 0 );
magma_zsetmatrix( M, M, hA, lda, dA2, ldda );
magma_zsetmatrix( M, N, hX, lda, dX[0], ldda );
magma_zsetmatrix( M, N, hB, lda, dwork[0], ldda );
gpu_time2 = magma_sync_wtime(0);
magma_zhemm(
MagmaLeft, MagmaLower, msize, N,
calpha, dA2+offset*ldda+offset, ldda,
dX[0], ldda,
cbeta, dwork[0], ldda );
gpu_time2 = magma_sync_wtime(0) - gpu_time2;
gpu_perf2 = gflops / gpu_time2;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.check ) {
// store ||A||*||X||
errorbis = lapackf77_zlange("fro", &msize, &msize, hA+offset*lda+offset, &lda, work );
errorbis *= lapackf77_zlange("fro", &msize, &N, hX, &lda, work );
//printf( "A =" ); magma_zprint( M, M, hA, lda );
//printf( "X =" ); magma_zprint( M, N, hX, lda );
//printf( "B =" ); magma_zprint( M, N, hB, lda );
cpu_time = magma_wtime();
blasf77_zhemm( "Left", "Lower", &msize, &N,
&calpha, hA+offset*lda+offset, &lda,
hX, &lda,
&cbeta, hB, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/*
trace_file = fopen("AJETE/C", "w");
for (int j = 0; j < N; j++)
for (int i = 0; i < siz; i++)
fprintf(trace_file, "%10d%10d%40.30e\n", i+1, j+1, hB[j*lda+i]);
fclose(trace_file);
*/
magma_int_t firstprint=0;
for(magma_int_t dev=0; dev < opts.ngpu; ++dev) {
magma_setdevice( dev );
magma_zgetmatrix( M, N, dB[dev], ldda, hR, lda );
// compute relative error ||R||/||A||*||X||, where R := B_magma - B_lapack = R - B
size = lda*N;
blasf77_zaxpy( &size, &c_neg_one, hB, &ione, hR, &ione );
error = lapackf77_zlange("fro", &msize, &N, hR, &lda, work) / errorbis;
//printf( "R =" ); magma_zprint( M, N, hR, lda );
if (firstprint == 0) {
printf( "%5d %5d %5d %5d %7.1f (%7.4f) %7.1f (%7.4f) %7.1f (%7.4f) %8.2e %s\n",
(int) M, (int) N, (int) nb, (int) offset,
cpu_perf, cpu_time,
gpu_perf, gpu_time,
gpu_perf2, gpu_time2,
error, (error < tol ? "ok" : "failed") );
}
else {
printf( "%89s %8.2e %s\n", " ",
error, (error < tol ? "ok" : "failed") );
}
status += ! (error < tol);
firstprint =1;
}
} else {
printf( "%5d %5d %5d %5d --- ( --- ) %7.1f (%7.4f) --- ( --- ) ---\n",
(int) M, (int) N, (int) nb, (int) offset,
gpu_perf, gpu_time );
}
TESTING_FREE_CPU( hA );
TESTING_FREE_CPU( hX );
TESTING_FREE_CPU( hB );
TESTING_FREE_PIN( hR );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
TESTING_FREE_DEV( dA[d] );
TESTING_FREE_DEV( dX[d] );
TESTING_FREE_DEV( dB[d] );
TESTING_FREE_DEV( dwork[d] );
TESTING_FREE_PIN( hwork[d] );
}
TESTING_FREE_PIN( hwork[opts.ngpu] );
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_FREE_DEV( dA2 );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
} // offset
printf( "\n" );
}
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[d][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
magma_event_destroy( redevents[d][i] );
magma_event_destroy( redevents2[d][i] );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgetrf
*/
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time, error;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex_ptr d_A, dwork;
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] = { 1024, 2048, 3072, 4032, 5184, 5600, 5600, 5600, 5600, 5600 };
magma_int_t ntest = 10;
magma_int_t i, info;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0, 0, 0, 1};
magmaDoubleComplex *work;
magmaDoubleComplex tmp;
double rwork[1];
magma_int_t *ipiv;
magma_int_t lwork, ldwork;
double A_norm, R_norm;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
}
if (N>0) size[0] = size[ntest-1] = N;
else exit(1);
}
else {
printf("\nUsage: \n");
printf(" testing_zgetri_gpu -N %d\n\n", 1024);
}
/* query for Lapack workspace size */
N = size[ntest-1];
lda = N;
work = &tmp;
lwork = -1;
lapackf77_zgetri( &N, h_A, &lda, ipiv, work, &lwork, &info );
if (info != 0)
printf("lapackf77_zgetri returned error %d\n", (int) info);
lwork = int( MAGMA_Z_REAL( *work ));
/* query for Magma workspace size */
ldwork = N * magma_get_zgetri_nb( N );
/* Initialize */
magma_queue_t queue;
magma_device_t device[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
/* Allocate memory */
n2 = N * N;
ldda = ((N+31)/32) * 32;
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( work, magmaDoubleComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dwork, magmaDoubleComplex, ldwork );
printf(" N CPU GFlop/s GPU GFlop/s ||R||_F / ||A||_F\n");
printf("========================================================\n");
for( i=0; i < ntest; i++ ){
N = size[i];
lda = N;
n2 = lda*N;
gflops = FLOPS_ZGETRI( (double)N ) / 1e9;
ldda = ((N+31)/32)*32;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
A_norm = lapackf77_zlange( "f", &N, &N, h_A, &lda, rwork );
/* Factor the matrix. Both MAGMA and LAPACK will use this factor. */
magma_zsetmatrix( N, N, h_A, 0, lda, d_A, 0, ldda, queue );
magma_zgetrf_gpu( N, N, d_A, 0, ldda, ipiv, &info, queue );
magma_zgetmatrix( N, N, d_A, 0, ldda, h_A, 0, lda, queue );
// check for exact singularity
//h_A[ 10 + 10*lda ] = MAGMA_Z_MAKE( 0.0, 0.0 );
//magma_zsetmatrix( N, N, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
//warm-up
magma_zgetri_gpu( N, d_A, 0, ldda, ipiv, dwork, 0, ldwork, &info, queue );
magma_zsetmatrix( N, N, h_A, 0, lda, d_A, 0, ldda, queue );
gpu_time = magma_wtime();
magma_zgetri_gpu( N, d_A, 0, ldda, ipiv, dwork, 0, ldwork, &info, queue );
gpu_time = magma_wtime()-gpu_time;
if (info != 0)
printf("magma_zgetri_gpu returned error %d\n", (int) info);
gpu_perf = gflops / gpu_time;
magma_zgetmatrix( N, N, d_A, 0, ldda, h_R, 0, lda, queue );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zgetri( &N, h_A, &lda, ipiv, work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf("lapackf77_zgetri returned error %d\n", (int) info);
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_zaxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
R_norm = lapackf77_zlange( "f", &N, &N, h_R, &lda, rwork );
printf( "%5d %6.2f %6.2f %e\n",
(int) N, cpu_perf, gpu_perf, R_norm / A_norm );
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dwork );
/* Shutdown */
magma_queue_destroy( queue );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zher2k
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cublas_perf, cublas_time, cpu_perf, cpu_time;
double cublas_error, Cnorm, work[1];
magma_int_t N, K;
magma_int_t Ak, An, Bk, Bn;
magma_int_t sizeA, sizeB, sizeC;
magma_int_t lda, ldb, ldc, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex *h_A, *h_B, *h_C, *h_Ccublas;
magmaDoubleComplex *d_A, *d_B, *d_C;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.29, -0.86 );
double beta = MAGMA_D_MAKE( -0.48, 0.38 );
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("If running lapack (option --lapack), CUBLAS error is computed\n"
"relative to CPU BLAS result.\n\n");
printf("uplo = %s, transA = %s\n",
lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA) );
printf(" N K CUBLAS Gflop/s (ms) CPU Gflop/s (ms) CUBLAS error\n");
printf("==================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.msize[itest];
K = opts.ksize[itest];
gflops = FLOPS_ZHER2K(K, N) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
lda = An = N;
Ak = K;
ldb = Bn = N;
Bk = K;
} else {
lda = An = K;
Ak = N;
ldb = Bn = K;
Bk = N;
}
ldc = N;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ((ldc+31)/32)*32;
sizeA = lda*Ak;
sizeB = ldb*Ak;
sizeC = ldc*N;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC_CPU( h_B, magmaDoubleComplex, ldb*Bk );
TESTING_MALLOC_CPU( h_C, magmaDoubleComplex, ldc*N );
TESTING_MALLOC_CPU( h_Ccublas, magmaDoubleComplex, ldc*N );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*Ak );
TESTING_MALLOC_DEV( d_B, magmaDoubleComplex, lddb*Bk );
TESTING_MALLOC_DEV( d_C, magmaDoubleComplex, lddc*N );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_zlarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_zlarnv( &ione, ISEED, &sizeC, h_C );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_zsetmatrix( An, Ak, h_A, lda, d_A, ldda );
magma_zsetmatrix( Bn, Bk, h_B, ldb, d_B, lddb );
magma_zsetmatrix( N, N, h_C, ldc, d_C, lddc );
cublas_time = magma_sync_wtime( NULL );
cublasZher2k( handle, cublas_uplo_const(opts.uplo), cublas_trans_const(opts.transA), N, K,
&alpha, d_A, ldda,
d_B, lddb,
&beta, d_C, lddc );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetmatrix( N, N, d_C, lddc, h_Ccublas, ldc );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_zher2k( lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), &N, &K,
&alpha, h_A, &lda,
h_B, &ldb,
&beta, h_C, &ldc );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & cublas, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_zlange( "M", &N, &N, h_C, &ldc, work );
blasf77_zaxpy( &sizeC, &c_neg_one, h_C, &ione, h_Ccublas, &ione );
cublas_error = lapackf77_zlange( "M", &N, &N, h_Ccublas, &ldc, work ) / Cnorm;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) K,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
cublas_error, (cublas_error < tol ? "ok" : "failed"));
status += ! (cublas_error < tol);
}
else {
printf("%5d %5d %7.2f (%7.2f) --- ( --- ) --- ---\n",
(int) N, (int) K,
cublas_perf, 1000.*cublas_time);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_Ccublas );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgetri_batched
*/
int main( int argc, char** argv)
{
TESTING_INIT();
// constants
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_Ainv, *h_R, *work;
magmaDoubleComplex_ptr d_A, d_invA;
magmaDoubleComplex_ptr *dA_array;
magmaDoubleComplex_ptr *dinvA_array;
magma_int_t **dipiv_array;
magma_int_t *dinfo_array;
magma_int_t *ipiv, *cpu_info;
magma_int_t *d_ipiv, *d_info;
magma_int_t N, n2, lda, ldda, info, info1, info2, lwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex tmp;
double error, rwork[1];
magma_int_t columns;
magma_int_t status = 0;
magma_opts opts( MagmaOptsBatched );
opts.parse_opts( argc, argv );
magma_int_t batchCount = opts.batchcount;
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% batchCount N CPU Gflop/s (ms) GPU Gflop/s (ms) ||I - A*A^{-1}||_1 / (N*cond(A))\n");
printf("%%===============================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N * batchCount;
ldda = magma_roundup( N, opts.align ); // multiple of 32 by default
// This is the correct flops but since this getri_batched is based on
// 2 trsm = getrs and to know the real flops I am using the getrs one
//gflops = (FLOPS_ZGETRF( N, N ) + FLOPS_ZGETRI( N ))/ 1e9 * batchCount;
gflops = (FLOPS_ZGETRF( N, N ) + FLOPS_ZGETRS( N, N ))/ 1e9 * batchCount;
// query for workspace size
lwork = -1;
lapackf77_zgetri( &N, NULL, &lda, NULL, &tmp, &lwork, &info );
if (info != 0) {
printf("lapackf77_zgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
lwork = magma_int_t( MAGMA_Z_REAL( tmp ));
TESTING_MALLOC_CPU( cpu_info, magma_int_t, batchCount );
TESTING_MALLOC_CPU( ipiv, magma_int_t, N * batchCount );
TESTING_MALLOC_CPU( work, magmaDoubleComplex, lwork*batchCount );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_Ainv, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N * batchCount );
TESTING_MALLOC_DEV( d_invA, magmaDoubleComplex, ldda*N * batchCount );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, N * batchCount );
TESTING_MALLOC_DEV( d_info, magma_int_t, batchCount );
TESTING_MALLOC_DEV( dA_array, magmaDoubleComplex*, batchCount );
TESTING_MALLOC_DEV( dinvA_array, magmaDoubleComplex*, batchCount );
TESTING_MALLOC_DEV( dinfo_array, magma_int_t, batchCount );
TESTING_MALLOC_DEV( dipiv_array, magma_int_t*, batchCount );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
columns = N * batchCount;
lapackf77_zlacpy( MagmaFullStr, &N, &columns, h_A, &lda, h_R, &lda );
lapackf77_zlacpy( MagmaFullStr, &N, &columns, h_A, &lda, h_Ainv, &lda );
magma_zsetmatrix( N, columns, h_R, lda, d_A, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zset_pointer( dA_array, d_A, ldda, 0, 0, ldda * N, batchCount, opts.queue );
magma_zset_pointer( dinvA_array, d_invA, ldda, 0, 0, ldda * N, batchCount, opts.queue );
magma_iset_pointer( dipiv_array, d_ipiv, 1, 0, 0, N, batchCount, opts.queue );
gpu_time = magma_sync_wtime( opts.queue );
info1 = magma_zgetrf_batched( N, N, dA_array, ldda, dipiv_array, dinfo_array, batchCount, opts.queue);
info2 = magma_zgetri_outofplace_batched( N, dA_array, ldda, dipiv_array, dinvA_array, ldda, dinfo_array, batchCount, opts.queue);
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
// check correctness of results throught "dinfo_magma" and correctness of argument throught "info"
magma_getvector( batchCount, sizeof(magma_int_t), dinfo_array, 1, cpu_info, 1, opts.queue );
for (magma_int_t i=0; i < batchCount; i++)
{
if (cpu_info[i] != 0 ) {
printf("magma_zgetrf_batched matrix %d returned error %d\n", (int) i, (int)cpu_info[i] );
}
}
if (info1 != 0) printf("magma_zgetrf_batched returned argument error %d: %s.\n", (int) info1, magma_strerror( info1 ));
if (info2 != 0) printf("magma_zgetri_batched returned argument error %d: %s.\n", (int) info2, magma_strerror( info2 ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
#if !defined (BATCHED_DISABLE_PARCPU) && defined(_OPENMP)
magma_int_t nthreads = magma_get_lapack_numthreads();
magma_set_lapack_numthreads(1);
magma_set_omp_numthreads(nthreads);
#pragma omp parallel for schedule(dynamic)
#endif
for (int i=0; i < batchCount; i++)
{
magma_int_t locinfo;
lapackf77_zgetrf(&N, &N, h_Ainv + i*lda*N, &lda, ipiv + i*N, &locinfo);
if (locinfo != 0) {
printf("lapackf77_zgetrf returned error %d: %s.\n",
(int) locinfo, magma_strerror( locinfo ));
}
lapackf77_zgetri(&N, h_Ainv + i*lda*N, &lda, ipiv + i*N, work + i*lwork, &lwork, &locinfo );
if (locinfo != 0) {
printf("lapackf77_zgetri returned error %d: %s.\n",
(int) locinfo, magma_strerror( locinfo ));
}
}
#if !defined (BATCHED_DISABLE_PARCPU) && defined(_OPENMP)
magma_set_lapack_numthreads(nthreads);
#endif
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
printf("%10d %5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) batchCount, (int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000. );
}
else {
printf("%10d %5d --- ( --- ) %7.2f (%7.2f)",
(int) batchCount, (int) N, gpu_perf, gpu_time*1000. );
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.check ) {
magma_igetvector( N*batchCount, d_ipiv, 1, ipiv, 1, opts.queue );
magma_zgetmatrix( N, N*batchCount, d_invA, ldda, h_Ainv, lda, opts.queue );
error = 0;
for (magma_int_t i=0; i < batchCount; i++)
{
for (magma_int_t k=0; k < N; k++) {
if (ipiv[i*N+k] < 1 || ipiv[i*N+k] > N )
{
printf("error for matrix %d ipiv @ %d = %d\n", (int) i, (int) k, (int) ipiv[i*N+k]);
error = -1;
}
}
if (error == -1) {
break;
}
// compute 1-norm condition number estimate, following LAPACK's zget03
double normA, normAinv, rcond, err;
normA = lapackf77_zlange( "1", &N, &N, h_A + i*lda*N, &lda, rwork );
normAinv = lapackf77_zlange( "1", &N, &N, h_Ainv + i*lda*N, &lda, rwork );
if ( normA <= 0 || normAinv <= 0 ) {
rcond = 0;
err = 1 / (tol/opts.tolerance); // == 1/eps
}
else {
rcond = (1 / normA) / normAinv;
// R = I
// R -= A*A^{-1}
// err = ||I - A*A^{-1}|| / ( N ||A||*||A^{-1}|| ) = ||R|| * rcond / N, using 1-norm
lapackf77_zlaset( "full", &N, &N, &c_zero, &c_one, h_R + i*lda*N, &lda );
blasf77_zgemm( "no", "no", &N, &N, &N, &c_neg_one,
h_A + i*lda*N, &lda,
h_Ainv + i*lda*N, &lda, &c_one,
h_R + i*lda*N, &lda );
err = lapackf77_zlange( "1", &N, &N, h_R + i*lda*N, &lda, rwork );
err = err * rcond / N;
}
if ( isnan(err) || isinf(err) ) {
error = err;
break;
}
error = max( err, error );
}
bool okay = (error < tol);
status += ! okay;
printf(" %8.2e %s\n", error, (okay ? "ok" : "failed") );
}
else {
printf("\n");
}
TESTING_FREE_CPU( cpu_info );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_Ainv );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_invA );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_info );
TESTING_FREE_DEV( dA_array );
TESTING_FREE_DEV( dinvA_array );
TESTING_FREE_DEV( dinfo_array );
TESTING_FREE_DEV( dipiv_array );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
int main( int argc, char** argv)
{
#define hA(i,j) (hA + (i) + (j)*lda)
TESTING_CUDA_INIT();
cuDoubleComplex c_zero = MAGMA_Z_ZERO;
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex *hA, *hR, *dA;
//real_Double_t gpu_time, gpu_perf;
//int ione = 1;
//int ISEED[4] = {0, 0, 0, 1};
int nsize[] = { 32, 64, 96, 256, 100, 200, 512 };
int ntest = sizeof(nsize) / sizeof(int);
int n = nsize[ntest-1];
int lda = ((n + 31)/32)*32;
int ntile, nb;
TESTING_MALLOC ( hA, cuDoubleComplex, lda*n );
TESTING_MALLOC ( hR, cuDoubleComplex, lda*n );
TESTING_DEVALLOC ( dA, cuDoubleComplex, lda*n );
for( int t = 0; t < ntest; ++t ) {
n = nsize[t];
lda = ((n + 31)/32)*32;
// initialize matrices; entries are (i.j) for A
double nf = 100.;
for( int j = 0; j < n; ++j ) {
// upper
for( int i = 0; i < j; ++i ) {
*hA(i,j) = MAGMA_Z_MAKE( (i + j/nf)/nf, 0. );
}
// lower
for( int i = j; i < n; ++i ) {
*hA(i,j) = MAGMA_Z_MAKE( i + j/nf, 0. );
}
}
printf( "A%d = ", n );
magma_zprint( n, n, hA, lda );
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize( MagmaLower, n, dA, lda );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d = ", n );
magma_zprint( n, n, hR, lda );
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize( MagmaUpper, n, dA, lda );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "U%d = ", n );
magma_zprint( n, n, hR, lda );
// -----
//lapackf77_zlaset( "u", &n, &n, &c_zero, &c_one, hA, &lda );
nb = 64;
ntile = n / nb;
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_zprint( n, n, hR, lda );
nb = 32;
ntile = n / nb;
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_zprint( n, n, hR, lda );
ntile = (n - nb < 0 ? 0 : (n - nb) / (2*nb) + 1);
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, 2*nb, nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d_2m = ", n, nb );
magma_zprint( n, n, hR, lda );
nb = 25;
ntile = n / nb;
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_zprint( n, n, hR, lda );
nb = 25;
ntile = (n - nb < 0 ? 0 : (n - nb) / (3*nb) + 1);
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, 3*nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d_3n = ", n, nb );
magma_zprint( n, n, hR, lda );
nb = 100;
ntile = n / nb;
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magmablas_zsymmetrize( MagmaLower, n%nb, &dA[ ntile*nb*(1+lda) ], lda ); // last partial block
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_zprint( n, n, hR, lda );
// -----
nb = 64;
ntile = n / nb;
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaUpper, nb, dA, lda, ntile, nb, nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "U%d_%d = ", n, nb );
magma_zprint( n, n, hR, lda );
}
TESTING_FREE( hA );
TESTING_FREE( hR );
TESTING_DEVFREE( dA );
/* Shutdown */
TESTING_CUDA_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqrs
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double gpu_error, cpu_error, matnorm, work[1];
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_A2, *h_B, *h_X, *h_R, *tau, *h_work, tmp[1];
magmaDoubleComplex *d_A, *d_B;
magma_int_t M, N, n2, nrhs, lda, ldb, ldda, lddb, min_mn, max_mn, nb, info;
magma_int_t lworkgpu, lhwork, lhwork2;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
double tol = opts.tolerance * lapackf77_dlamch("E");
nrhs = opts.nrhs;
printf(" ||b-Ax|| / (N||A||)\n");
printf(" M N NRHS CPU GFlop/s (sec) GPU GFlop/s (sec) CPU GPU \n");
printf("===============================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
if ( M < N ) {
printf( "skipping M=%d, N=%d because M < N is not yet supported.\n", (int) M, (int) N );
continue;
}
min_mn = min(M, N);
max_mn = max(M, N);
lda = M;
ldb = max_mn;
n2 = lda*N;
ldda = ((M+31)/32)*32;
lddb = ((max_mn+31)/32)*32;
nb = magma_get_zgeqrf_nb(M);
gflops = (FLOPS_ZGEQRF( M, N ) + FLOPS_ZGEQRS( M, N, nrhs )) / 1e9;
// query for workspace size
lworkgpu = (M - N + nb)*(nrhs + nb) + nrhs*nb;
lhwork = -1;
lapackf77_zgeqrf(&M, &N, h_A, &M, tau, tmp, &lhwork, &info);
lhwork2 = (magma_int_t) MAGMA_Z_REAL( tmp[0] );
lhwork = -1;
lapackf77_zunmqr( MagmaLeftStr, MagmaConjTransStr,
&M, &nrhs, &min_mn, h_A, &lda, tau,
h_X, &ldb, tmp, &lhwork, &info);
lhwork = (magma_int_t) MAGMA_Z_REAL( tmp[0] );
lhwork = max( max( lhwork, lhwork2 ), lworkgpu );
TESTING_MALLOC( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC( h_A, magmaDoubleComplex, lda*N );
TESTING_MALLOC( h_A2, magmaDoubleComplex, lda*N );
TESTING_MALLOC( h_B, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC( h_X, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC( h_R, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC( h_work, magmaDoubleComplex, lhwork );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*N );
TESTING_DEVALLOC( d_B, magmaDoubleComplex, lddb*nrhs );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_A2, &lda );
// make random RHS
n2 = M*nrhs;
lapackf77_zlarnv( &ione, ISEED, &n2, h_B );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
// make consistent RHS
//n2 = N*nrhs;
//lapackf77_zlarnv( &ione, ISEED, &n2, h_X );
//blasf77_zgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
// &c_one, h_A, &lda,
// h_X, &ldb,
// &c_zero, h_B, &ldb );
//lapackf77_zlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
gpu_time = magma_wtime();
magma_zgels3_gpu( MagmaNoTrans, M, N, nrhs, d_A, ldda,
d_B, lddb, h_work, lworkgpu, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgels returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Get the solution in h_X
magma_zgetmatrix( N, nrhs, d_B, lddb, h_X, ldb );
// compute the residual
blasf77_zgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A, &lda,
h_X, &ldb,
&c_one, h_R, &ldb);
matnorm = lapackf77_zlange("f", &M, &N, h_A, &lda, work);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_X, &ldb );
cpu_time = magma_wtime();
lapackf77_zgels( MagmaNoTransStr, &M, &N, &nrhs,
h_A, &lda, h_X, &ldb, h_work, &lhwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zgels returned error %d: %s.\n",
(int) info, magma_strerror( info ));
blasf77_zgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A2, &lda,
h_X, &ldb,
&c_one, h_B, &ldb);
cpu_error = lapackf77_zlange("f", &M, &nrhs, h_B, &ldb, work) / (min_mn*matnorm);
gpu_error = lapackf77_zlange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*matnorm);
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, (int) nrhs,
cpu_perf, cpu_time, gpu_perf, gpu_time, cpu_error, gpu_error );
printf("%s\n", (gpu_error < tol ? "" : " failed"));
status |= ! (gpu_error < tol);
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_A2 );
TESTING_FREE( h_B );
TESTING_FREE( h_X );
TESTING_FREE( h_R );
TESTING_FREE( h_work );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zpotrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex *d_A;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t N, n2, lda, ldda, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
gflops = FLOPS_ZPOTRI( N ) / 1e9;
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*N );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hpd( N, h_A, lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* factorize matrix */
magma_zsetmatrix( N, N, h_A, lda, d_A, ldda );
magma_zpotrf_gpu( opts.uplo, N, d_A, ldda, &info );
// check for exact singularity
//magma_zgetmatrix( N, N, d_A, ldda, h_R, lda );
//h_R[ 10 + 10*lda ] = MAGMA_Z_MAKE( 0.0, 0.0 );
//magma_zsetmatrix( N, N, h_R, lda, d_A, ldda );
gpu_time = magma_wtime();
magma_zpotri_gpu( opts.uplo, N, d_A, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zpotri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
lapackf77_zpotrf( &opts.uplo, &N, h_A, &lda, &info );
cpu_time = magma_wtime();
lapackf77_zpotri( &opts.uplo, &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zpotri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( N, N, d_A, ldda, h_R, lda );
error = lapackf77_zlange("f", &N, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &N, &N, h_R, &lda, work) / error;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e%s\n",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "" : " failed") );
status |= ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) ---\n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
extern "C" magma_int_t
magma_zheevx_gpu(char jobz, char range, char uplo, magma_int_t n,
magmaDoubleComplex *da, magma_int_t ldda, double vl, double vu,
magma_int_t il, magma_int_t iu, double abstol, magma_int_t *m,
double *w, magmaDoubleComplex *dz, magma_int_t lddz,
magmaDoubleComplex *wa, magma_int_t ldwa,
magmaDoubleComplex *wz, magma_int_t ldwz,
magmaDoubleComplex *work, magma_int_t lwork,
double *rwork, magma_int_t *iwork, magma_int_t *ifail, magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
ZHEEVX 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.
Arguments
=========
JOBZ (input) CHARACTER*1
= 'N': Compute eigenvalues only;
= 'V': Compute eigenvalues and eigenvectors.
RANGE (input) CHARACTER*1
= 'A': all eigenvalues will be found.
= 'V': all eigenvalues in the half-open interval (VL,VU]
will be found.
= 'I': the IL-th through IU-th eigenvalues will be found.
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
DA (device input/output) COMPLEX_16 array, dimension (LDDA, 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. If UPLO = 'L',
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, the lower triangle (if UPLO='L') or the upper
triangle (if UPLO='U') of A, including the diagonal, is
destroyed.
LDDA (input) INTEGER
The leading dimension of the array DA. LDDA >= max(1,N).
VL (input) DOUBLE PRECISION
VU (input) DOUBLE PRECISION
If RANGE='V', the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = 'A' or 'I'.
IL (input) INTEGER
IU (input) INTEGER
If RANGE='I', 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 = 'A' or 'V'.
ABSTOL (input) DOUBLE PRECISION
The absolute error tolerance for the eigenvalues.
An approximate eigenvalue is accepted as converged
when it is determined to lie in an interval [a,b]
of width less than or equal to
ABSTOL + EPS * max( |a|,|b| ) ,
where EPS is the machine precision. If ABSTOL is less than
or equal to zero, then EPS*|T| will be used in its place,
where |T| is the 1-norm of the tridiagonal matrix obtained
by reducing A to tridiagonal form.
Eigenvalues will be computed most accurately when ABSTOL is
set to twice the underflow threshold 2*DLAMCH('S'), not zero.
If this routine returns with INFO>0, indicating that some
eigenvectors did not converge, try setting ABSTOL to
2*DLAMCH('S').
See "Computing Small Singular Values of Bidiagonal Matrices
with Guaranteed High Relative Accuracy," by Demmel and
Kahan, LAPACK Working Note #3.
M (output) INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = 'A', M = N, and if RANGE = 'I', M = IU-IL+1.
W (output) DOUBLE PRECISION array, dimension (N)
On normal exit, the first M elements contain the selected
eigenvalues in ascending order.
DZ (device output) COMPLEX_16 array, dimension (LDDZ, max(1,M))
If JOBZ = 'V', then if INFO = 0, the first M columns of Z
contain the orthonormal eigenvectors of the matrix A
corresponding to the selected eigenvalues, with the i-th
column of Z holding the eigenvector associated with W(i).
If an eigenvector fails to converge, then that column of Z
contains the latest approximation to the eigenvector, and the
index of the eigenvector is returned in IFAIL.
If JOBZ = 'N', then Z is not referenced.
Note: the user must ensure that at least max(1,M) columns are
supplied in the array Z; if RANGE = 'V', the exact value of M
is not known in advance and an upper bound must be used.
********* (workspace) If FAST_HEMV is defined DZ should be (LDDZ, max(1,N)) in both cases.
LDDZ (input) INTEGER
The leading dimension of the array DZ. LDDZ >= 1, and if
JOBZ = 'V', LDDZ >= max(1,N).
WA (workspace) COMPLEX_16 array, dimension (LDWA, N)
LDWA (input) INTEGER
The leading dimension of the array WA. LDWA >= max(1,N).
WZ (workspace) COMPLEX_16 array, dimension (LDWZ, max(1,M))
LDWZ (input) INTEGER
The leading dimension of the array DZ. LDWZ >= 1, and if
JOBZ = 'V', LDWZ >= max(1,N).
WORK (workspace/output) COMPLEX_16 array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK. LWORK >= (NB+1)*N,
where NB is the max of the blocksize for ZHETRD.
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.
RWORK (workspace) DOUBLE PRECISION array, dimension (7*N)
IWORK (workspace) INTEGER array, dimension (5*N)
IFAIL (output) INTEGER array, dimension (N)
If JOBZ = 'V', then if INFO = 0, the first M elements of
IFAIL are zero. If INFO > 0, then IFAIL contains the
indices of the eigenvectors that failed to converge.
If JOBZ = 'N', then IFAIL is not referenced.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, then i eigenvectors failed to converge.
Their indices are stored in array IFAIL.
===================================================================== */
char uplo_[2] = {uplo, 0};
char jobz_[2] = {jobz, 0};
char range_[2] = {range, 0};
magma_int_t ione = 1;
char order[1];
magma_int_t indd, inde;
magma_int_t imax;
magma_int_t lopt, itmp1, indee;
magma_int_t lower, wantz;
magma_int_t i, j, jj, i__1;
magma_int_t alleig, valeig, indeig;
magma_int_t iscale, indibl;
magma_int_t indiwk, indisp, indtau;
magma_int_t indrwk, indwrk;
magma_int_t llwork, nsplit;
magma_int_t lquery;
magma_int_t iinfo;
double safmin;
double bignum;
double smlnum;
double eps, tmp1;
double anrm;
double sigma, d__1;
double rmin, rmax;
double *dwork;
/* Function Body */
lower = lapackf77_lsame(uplo_, MagmaLowerStr);
wantz = lapackf77_lsame(jobz_, MagmaVecStr);
alleig = lapackf77_lsame(range_, "A");
valeig = lapackf77_lsame(range_, "V");
indeig = lapackf77_lsame(range_, "I");
lquery = lwork == -1;
*info = 0;
if (! (wantz || lapackf77_lsame(jobz_, MagmaNoVecStr))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || lapackf77_lsame(uplo_, MagmaUpperStr))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (ldda < max(1,n)) {
*info = -6;
} else if (lddz < 1 || (wantz && lddz < n)) {
*info = -15;
} else if (ldwa < max(1,n)) {
*info = -17;
} else if (ldwz < 1 || (wantz && ldwz < n)) {
*info = -19;
} 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);
lopt = n * (nb + 1);
work[0] = MAGMA_Z_MAKE( lopt, 0 );
if (lwork < lopt && ! lquery) {
*info = -21;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
} else if (lquery) {
return *info;
}
*m = 0;
/* Check if matrix is very small then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
#ifdef ENABLE_DEBUG
printf("--------------------------------------------------------------\n");
printf(" warning matrix too small N=%d NB=%d, calling lapack on CPU \n", (int) n, (int) nb);
printf("--------------------------------------------------------------\n");
#endif
magmaDoubleComplex *a = (magmaDoubleComplex *) malloc( n * n * sizeof(magmaDoubleComplex) );
magma_zgetmatrix(n, n, da, ldda, a, n);
lapackf77_zheevx(jobz_, range_, uplo_,
&n, a, &n, &vl, &vu, &il, &iu, &abstol, m,
w, wz, &ldwz, work, &lwork,
rwork, iwork, ifail, info);
magma_zsetmatrix( n, n, a, n, da, ldda);
magma_zsetmatrix( n, *m, wz, ldwz, dz, lddz);
free(a);
return *info;
}
if (MAGMA_SUCCESS != magma_dmalloc( &dwork, n )) {
fprintf (stderr, "!!!! device memory allocation error (magma_zheevx_gpu)\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
--w;
--work;
--rwork;
--iwork;
--ifail;
/* 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('M', 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) {
d__1 = 1.;
magmablas_zlascl(uplo, 0, 0, 1., sigma, n, n, da,
ldda, info);
if (abstol > 0.) {
abstol *= sigma;
}
if (valeig) {
vl *= sigma;
vu *= sigma;
}
}
/* Call ZHETRD to reduce Hermitian matrix to tridiagonal form. */
indd = 1;
inde = indd + n;
indrwk = inde + n;
indtau = 1;
indwrk = indtau + n;
llwork = lwork - indwrk + 1;
#ifdef FAST_HEMV
magma_zhetrd2_gpu(uplo, n, da, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wa, ldwa, &work[indwrk], llwork, dz, lddz*n, &iinfo);
#else
magma_zhetrd_gpu (uplo, n, da, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wa, ldwa, &work[indwrk], llwork, &iinfo);
#endif
lopt = n + (magma_int_t)MAGMA_Z_REAL(work[indwrk]);
/* If all eigenvalues are desired and ABSTOL is less than or equal to
zero, then call DSTERF or ZUNGTR and ZSTEQR. If this fails for
some eigenvalue, then try DSTEBZ. */
if ((alleig || (indeig && il == 1 && iu == n)) && abstol <= 0.) {
blasf77_dcopy(&n, &rwork[indd], &ione, &w[1], &ione);
indee = indrwk + 2*n;
if (! wantz) {
i__1 = n - 1;
blasf77_dcopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_dsterf(&n, &w[1], &rwork[indee], info);
}
else {
lapackf77_zlacpy("A", &n, &n, wa, &ldwa, wz, &ldwz);
lapackf77_zungtr(uplo_, &n, wz, &ldwz, &work[indtau], &work[indwrk], &llwork, &iinfo);
i__1 = n - 1;
blasf77_dcopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_zsteqr(jobz_, &n, &w[1], &rwork[indee], wz, &ldwz, &rwork[indrwk], info);
if (*info == 0) {
for (i = 1; i <= n; ++i) {
ifail[i] = 0;
}
magma_zsetmatrix( n, n, wz, ldwz, dz, lddz );
}
}
if (*info == 0) {
*m = n;
}
}
/* Otherwise, call DSTEBZ and, if eigenvectors are desired, ZSTEIN. */
if (*m == 0) {
*info = 0;
if (wantz) {
*(unsigned char *)order = 'B';
} else {
*(unsigned char *)order = 'E';
}
indibl = 1;
indisp = indibl + n;
indiwk = indisp + n;
lapackf77_dstebz(range_, order, &n, &vl, &vu, &il, &iu, &abstol, &rwork[indd], &rwork[inde], m,
&nsplit, &w[1], &iwork[indibl], &iwork[indisp], &rwork[indrwk], &iwork[indiwk], info);
if (wantz) {
lapackf77_zstein(&n, &rwork[indd], &rwork[inde], m, &w[1], &iwork[indibl], &iwork[indisp],
wz, &ldwz, &rwork[indrwk], &iwork[indiwk], &ifail[1], info);
magma_zsetmatrix( n, *m, wz, ldwz, dz, lddz );
/* Apply unitary matrix used in reduction to tridiagonal
form to eigenvectors returned by ZSTEIN. */
magma_zunmtr_gpu(MagmaLeft, uplo, MagmaNoTrans, n, *m, da, ldda, &work[indtau],
dz, lddz, wa, ldwa, &iinfo);
}
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = *m;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_dscal(&imax, &d__1, &w[1], &ione);
}
/* If eigenvalues are not in order, then sort them, along with
eigenvectors. */
if (wantz) {
for (j = 1; j <= *m-1; ++j) {
i = 0;
tmp1 = w[j];
for (jj = j + 1; jj <= *m; ++jj) {
if (w[jj] < tmp1) {
i = jj;
tmp1 = w[jj];
}
}
if (i != 0) {
itmp1 = iwork[indibl + i - 1];
w[i] = w[j];
iwork[indibl + i - 1] = iwork[indibl + j - 1];
w[j] = tmp1;
iwork[indibl + j - 1] = itmp1;
magma_zswap(n, dz + (i-1)*lddz, ione, dz + (j-1)*lddz, ione);
if (*info != 0) {
itmp1 = ifail[i];
ifail[i] = ifail[j];
ifail[j] = itmp1;
}
}
}
}
/* Set WORK(1) to optimal complex workspace size. */
work[1] = MAGMA_Z_MAKE( lopt, 0 );
return *info;
} /* magma_zheevx_gpu */
/**
Purpose
-------
ZGETRF 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.
If the current stream is NULL, this version replaces it with a new
stream to overlap computation with communication.
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. N >= 0.
@param[in,out]
dA COMPLEX_16 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.
@param[in]
ldda INTEGER
The leading dimension of the array A. 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_zgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_zgetrf_gpu(
magma_int_t m, magma_int_t n,
magmaDoubleComplex_ptr dA, magma_int_t ldda, magma_int_t *ipiv,
magma_int_t *info)
{
#define dAT(i_, j_) (dAT + (i_)*nb*lddat + (j_)*nb)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, maxn, mindim;
magma_int_t i, j, rows, cols, s, lddat, ldwork;
magmaDoubleComplex *dAT, *dAP, *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_zgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_zmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_zgetmatrix( m, n, dA, ldda, work, m );
lapackf77_zgetrf(&m, &n, work, &m, ipiv, info);
magma_zsetmatrix( m, n, work, m, dA, ldda );
magma_free_cpu(work);
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
if (MAGMA_SUCCESS != magma_zmalloc( &dAP, nb*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
// square matrices can be done in place;
// rectangular requires copy to transpose
if ( m == n ) {
dAT = dA;
lddat = ldda;
magmablas_ztranspose_inplace( m, dAT, ldda );
}
else {
lddat = maxn; // N-by-M
if (MAGMA_SUCCESS != magma_zmalloc( &dAT, lddat*maxm )) {
magma_free( dAP );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_ztranspose( m, n, dA, ldda, dAT, lddat );
}
ldwork = maxm;
if (MAGMA_SUCCESS != magma_zmalloc_pinned( &work, ldwork*nb )) {
magma_free( dAP );
if ( ! (m == n))
magma_free( dAT );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
magma_queue_t stream[2];
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_create( &stream[0] );
if (orig_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = orig_stream;
}
for( j=0; j < s; j++ ) {
// download j-th panel
cols = maxm - j*nb;
magmablas_ztranspose( nb, m-j*nb, dAT(j,j), lddat, dAP, cols );
// make sure that the transpose has completed
magma_queue_sync( stream[1] );
magma_zgetmatrix_async( m-j*nb, nb, dAP, cols, work, ldwork,
stream[0]);
if ( j > 0 ) {
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (j+1)*nb, nb,
c_one, dAT(j-1,j-1), lddat,
dAT(j-1,j+1), lddat );
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+1)*nb, m-j*nb, nb,
c_neg_one, dAT(j-1,j+1), lddat,
dAT(j, j-1), lddat,
c_one, dAT(j, j+1), lddat );
}
// do the cpu part
rows = m - j*nb;
magma_queue_sync( stream[0] );
lapackf77_zgetrf( &rows, &nb, work, &ldwork, ipiv+j*nb, &iinfo);
if ( *info == 0 && iinfo > 0 )
*info = iinfo + j*nb;
// upload j-th panel
magma_zsetmatrix_async( m-j*nb, nb, work, ldwork, dAP, maxm,
stream[0]);
for( i=j*nb; i < j*nb + nb; ++i ) {
ipiv[i] += j*nb;
}
magmablas_zlaswp( n, dAT, lddat, j*nb + 1, j*nb + nb, ipiv, 1 );
magma_queue_sync( stream[0] );
magmablas_ztranspose( m-j*nb, nb, dAP, maxm, dAT(j,j), lddat );
// do the small non-parallel computations (next panel update)
if ( s > (j+1) ) {
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(j, j ), lddat,
dAT(j, j+1), lddat);
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(j+1)*nb, nb,
c_neg_one, dAT(j, j+1), lddat,
dAT(j+1, j ), lddat,
c_one, dAT(j+1, j+1), lddat );
}
else {
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(j, j ), lddat,
dAT(j, j+1), lddat);
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+1)*nb, m-(j+1)*nb, nb,
c_neg_one, dAT(j, j+1), lddat,
dAT(j+1, j ), lddat,
c_one, dAT(j+1, j+1), lddat );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
if ( nb0 > 0 ) {
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_ztranspose( nb0, rows, dAT(s,s), lddat, dAP, maxm );
magma_zgetmatrix( rows, nb0, dAP, maxm, work, ldwork );
// do the cpu part
lapackf77_zgetrf( &rows, &nb0, work, &ldwork, ipiv+s*nb, &iinfo);
if ( *info == 0 && iinfo > 0 )
*info = iinfo + s*nb;
for( i=s*nb; i < s*nb + nb0; ++i ) {
ipiv[i] += s*nb;
}
magmablas_zlaswp( n, dAT, lddat, s*nb + 1, s*nb + nb0, ipiv, 1 );
// upload j-th panel
magma_zsetmatrix( rows, nb0, work, ldwork, dAP, maxm );
magmablas_ztranspose( rows, nb0, dAP, maxm, dAT(s,s), lddat );
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s,s), lddat,
dAT(s,s)+nb0, lddat);
}
// undo transpose
if ( m == n ) {
magmablas_ztranspose_inplace( m, dAT, lddat );
}
else {
magmablas_ztranspose( n, m, dAT, lddat, dA, ldda );
magma_free( dAT );
}
magma_free( dAP );
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (orig_stream == NULL) {
magma_queue_destroy( stream[1] );
}
magmablasSetKernelStream( orig_stream );
}
return *info;
} /* magma_zgetrf_gpu */
