/* Task execution code */
static void SCHED_panel_update(Quark* quark)
{
int N;
cuDoubleComplex *A1;
int LDA;
int K2;
int *IPIV;
cuDoubleComplex *A2;
int M;
int K;
cuDoubleComplex *A3;
cuDoubleComplex *A4;
int ione=1;
cuDoubleComplex mone = MAGMA_Z_NEG_ONE;
cuDoubleComplex one = MAGMA_Z_ONE;
quark_unpack_args_10(quark, N, A1, LDA, K2, IPIV, A2, M, K, A3, A4);
lapackf77_zlaswp(&N, A1, &LDA, &ione, &K2, IPIV, &ione);
blasf77_ztrsm("l", "l", "n", "u",
&K2, &N, &one, A2, &LDA, A1, &LDA);
if (M > 0) {
blasf77_zgemm("n","n",
&M, &N, &K, &mone, A3, &LDA, A1, &LDA, &one, A4, &LDA);
}
}
extern "C" magma_int_t
magma_zgetrf2_piv(magma_int_t m, magma_int_t n, magma_int_t start, magma_int_t end,
magmaDoubleComplex *A, magma_int_t lda, magma_int_t *ipiv, magma_int_t *info)
{
magma_int_t I, k1, k2, nb, incx, minmn;
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0)
return MAGMA_ERR_ILLEGAL_VALUE;
/* Quick return if possible */
if (m == 0 || n == 0)
return MAGMA_SUCCESS;
/* initialize nb */
nb = magma_get_zgetrf_nb(m);
minmn = min( end, min(m,n) );
for( I=start; I < end-nb; I += nb ) {
incx = 1;
k1 = 1+I+nb;
k2 = minmn;
lapackf77_zlaswp(&nb, A(0,I), &lda, &k1, &k2, ipiv, &incx);
}
return MAGMA_SUCCESS;
} /* magma_zgetrf_piv */
extern "C" magma_int_t
magma_zgetrf_piv(magma_int_t m, magma_int_t n, magma_int_t NB,
magmaDoubleComplex *A, magma_int_t lda, magma_int_t *ipiv, magma_int_t *info)
{
magma_int_t I, k1, k2, incx, minmn;
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0)
return *info;
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* initialize nb */
minmn = min(m,n);
for( I=0; I < minmn-NB; I += NB ) {
k1 = 1+I+NB;
k2 = minmn;
incx = 1;
lapackf77_zlaswp(&NB, A(0,I), &lda, &k1, &k2, ipiv, &incx);
}
return *info;
} /* magma_zgetrf_piv */
/* Task execution code */
void SCHED_zlaswp(Quark* quark)
{
int N;
cuDoubleComplex *A;
int LDA;
int K2;
int *IPIV;
int ione=1;
quark_unpack_args_5(quark, N, A, LDA, K2, IPIV);
lapackf77_zlaswp(&N, A, &LDA, &ione, &K2, IPIV, &ione);
}
// On input, LU and ipiv is LU factorization of A. On output, LU is overwritten.
// Works for any m, n.
// Uses init_matrix() to re-generate original A as needed.
// Returns error in factorization, |PA - LU| / (n |A|)
// This allocates 3 more matrices to store A, L, and U.
double get_LU_error(magma_int_t M, magma_int_t N,
magmaDoubleComplex *LU, magma_int_t lda,
magma_int_t *ipiv)
{
magma_int_t min_mn = min(M,N);
magma_int_t ione = 1;
magma_int_t i, j;
magmaDoubleComplex alpha = MAGMA_Z_ONE;
magmaDoubleComplex beta = MAGMA_Z_ZERO;
magmaDoubleComplex *A, *L, *U;
double work[1], matnorm, residual;
TESTING_MALLOC_CPU( A, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( L, magmaDoubleComplex, M*min_mn );
TESTING_MALLOC_CPU( U, magmaDoubleComplex, min_mn*N );
memset( L, 0, M*min_mn*sizeof(magmaDoubleComplex) );
memset( U, 0, min_mn*N*sizeof(magmaDoubleComplex) );
// set to original A
init_matrix( M, N, A, lda );
lapackf77_zlaswp( &N, A, &lda, &ione, &min_mn, ipiv, &ione);
// copy LU to L and U, and set diagonal to 1
lapackf77_zlacpy( MagmaLowerStr, &M, &min_mn, LU, &lda, L, &M );
lapackf77_zlacpy( MagmaUpperStr, &min_mn, &N, LU, &lda, U, &min_mn );
for(j=0; j<min_mn; j++)
L[j+j*M] = MAGMA_Z_MAKE( 1., 0. );
matnorm = lapackf77_zlange("f", &M, &N, A, &lda, work);
blasf77_zgemm("N", "N", &M, &N, &min_mn,
&alpha, L, &M, U, &min_mn, &beta, LU, &lda);
for( j = 0; j < N; j++ ) {
for( i = 0; i < M; i++ ) {
LU[i+j*lda] = MAGMA_Z_SUB( LU[i+j*lda], A[i+j*lda] );
}
}
residual = lapackf77_zlange("f", &M, &N, LU, &lda, work);
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( L );
TESTING_FREE_CPU( U );
return residual / (matnorm * N);
}
/**
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_err_t
magma_zgetrs_gpu(magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magma_int_t *ipiv,
magmaDoubleComplex_ptr dB, size_t dB_offset, magma_int_t lddb,
magma_int_t *info, magma_queue_t queue)
{
/* -- clMagma (version 0.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
April 2012
Purpose
=======
Solves a system of linear equations
A * X = B or A' * X = B
with a general N-by-N matrix A using the LU factorization computed by ZGETRF_GPU.
Arguments
=========
TRANS (input) CHARACTER*1
Specifies the form of the system of equations:
= 'N': A * X = B (No transpose)
= 'T': A'* X = B (Transpose)
= 'C': A'* X = B (Conjugate transpose = Transpose)
N (input) INTEGER
The order of the matrix A. N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
A (input) COMPLEX_16 array on the GPU, dimension (LDA,N)
The factors L and U from the factorization A = P*L*U as computed
by ZGETRF_GPU.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
IPIV (input) INTEGER array, dimension (N)
The pivot indices from ZGETRF; for 1<=i<=N, row i of the
matrix was interchanged with row IPIV(i).
B (input/output) COMPLEX_16 array on the GPU, dimension (LDB,NRHS)
On entry, the right hand side matrix B.
On exit, the solution matrix X.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
HWORK (workspace) COMPLEX_16 array, dimension N*NRHS
===================================================================== */
magmaDoubleComplex z_one = MAGMA_Z_MAKE( 1.0, 0.0 );
magmaDoubleComplex *work = NULL;
magma_trans_t trans_ = trans;
long int notran = lapackf77_lsame(lapack_const(trans_), lapack_const(MagmaNoTrans));
magma_int_t i1, i2, inc;
*info = 0;
if ( (! notran) &&
(! lapackf77_lsame(lapack_const(trans_), lapack_const(MagmaTrans))) &&
(! lapackf77_lsame(lapack_const(trans_), lapack_const(MagmaConjTrans))) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (nrhs < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
} else if (lddb < max(1,n)) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (n == 0 || nrhs == 0) {
return *info;
}
work=(magmaDoubleComplex*)malloc( n*nrhs*sizeof(magmaDoubleComplex) );
if ( !work ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
i1 = 1;
i2 = n;
if (notran) {
inc = 1;
/* Solve A * X = B. */
magma_zgetmatrix( n, nrhs, dB, dB_offset, lddb, work, 0, n, queue );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, 0, n, dB, dB_offset, lddb, queue );
if ( nrhs == 1) {
chk(magma_ztrsv(MagmaLower, MagmaNoTrans, MagmaUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue));
chk(magma_ztrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue));
} else {
chk(magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
chk(magma_ztrsm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
}
} else {
inc = -1;
/* Solve A' * X = B. */
if ( nrhs == 1) {
chk(magma_ztrsv(MagmaUpper, trans, MagmaNonUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue ));
chk(magma_ztrsv(MagmaLower, trans, MagmaUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue ));
} else {
chk(magma_ztrsm(MagmaLeft, MagmaUpper, trans, MagmaNonUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
chk(magma_ztrsm(MagmaLeft, MagmaLower, trans, MagmaUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
}
magma_zgetmatrix( n, nrhs, dB, dB_offset, lddb, work, 0, n, queue );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, 0, n, dB, dB_offset, lddb, queue );
}
free(work);
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zswap, zswapblk, zlaswp, zlaswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
// OpenCL use: cl_mem , offset (two arguments);
// else use: pointer + offset (one argument).
#ifdef HAVE_clBLAS
#define d_A1(i_, j_) d_A1 , (i_) + (j_)*ldda
#define d_A2(i_, j_) d_A2 , (i_) + (j_)*ldda
#define d_ipiv(i_) d_ipiv , (i_)
#else
#define d_A1(i_, j_) (d_A1 + (i_) + (j_)*ldda)
#define d_A2(i_, j_) (d_A2 + (i_) + (j_)*ldda)
#define d_ipiv(i_) (d_ipiv + (i_))
#endif
#define h_A1(i_, j_) (h_A1 + (i_) + (j_)*lda)
#define h_A2(i_, j_) (h_A2 + (i_) + (j_)*lda)
magmaDoubleComplex *h_A1, *h_A2;
magmaDoubleComplex *h_R1, *h_R2;
magmaDoubleComplex_ptr d_A1, d_A2;
// row-major and column-major performance
real_Double_t row_perf0 = MAGMA_D_NAN, col_perf0 = MAGMA_D_NAN;
real_Double_t row_perf1 = MAGMA_D_NAN, col_perf1 = MAGMA_D_NAN;
real_Double_t row_perf2 = MAGMA_D_NAN, col_perf2 = MAGMA_D_NAN;
real_Double_t row_perf4 = MAGMA_D_NAN;
real_Double_t row_perf5 = MAGMA_D_NAN, col_perf5 = MAGMA_D_NAN;
real_Double_t row_perf6 = MAGMA_D_NAN, col_perf6 = MAGMA_D_NAN;
real_Double_t row_perf7 = MAGMA_D_NAN;
real_Double_t cpu_perf = MAGMA_D_NAN;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magmaInt_ptr d_ipiv;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
printf("%% %8s zswap zswap zswapblk zlaswp zlaswp2 zlaswpx zcopymatrix CPU (all in )\n", g_platform_str );
printf("%% N nb row-maj/col-maj row-maj/col-maj row-maj/col-maj row-maj row-maj row-maj/col-maj row-blk/col-blk zlaswp (GByte/s)\n");
printf("%%========================================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// For an N x N matrix, swap nb rows or nb columns using various methods.
// Each test is assigned one bit in the 'check' bitmask; bit=1 indicates failure.
// The variable 'shift' keeps track of which bit is for current test
magma_int_t shift = 1;
magma_int_t check = 0;
N = opts.nsize[itest];
lda = N;
ldda = magma_roundup( N, opts.align ); // multiple of 32 by default
nb = (opts.nb > 0 ? opts.nb : magma_get_zgetrf_nb( N, N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(magmaDoubleComplex) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_A2, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_R1, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_R2, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, nb );
TESTING_MALLOC_CPU( ipiv2, magma_int_t, nb );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, nb );
TESTING_MALLOC_DEV( d_A1, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( d_A2, magmaDoubleComplex, ldda*N );
// getrf always makes ipiv[j] >= j+1, where ipiv is one based and j is zero based
// some implementations (e.g., MacOS dlaswp) assume this
for( j=0; j < nb; j++ ) {
ipiv[j] = (rand() % (N-j)) + j + 1;
assert( ipiv[j] >= j+1 );
assert( ipiv[j] <= N );
}
/* =====================================================================
* cublas / clBLAS / Xeon Phi zswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1(0,0), ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2(0,0), ldda );
magmablasSetKernelStream( opts.queue ); // opts.handle also uses opts.queue
time = magma_sync_wtime( opts.queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
#ifdef HAVE_CUBLAS
cublasZswap( opts.handle, N, d_A1(0,j), 1, d_A2(0,ipiv[j]-1), 1 );
#else
magma_zswap( N, d_A1(0,j), 1, d_A2(0,ipiv[j]-1), 1, opts.queue );
#endif
}
}
time = magma_sync_wtime( opts.queue ) - time;
row_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1(0,j), &ione, h_A2(0,ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1(0,0), ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2(0,0), ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1(0,0), ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2(0,0), ldda );
time = magma_sync_wtime( opts.queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
#ifdef HAVE_CUBLAS
cublasZswap( opts.handle, N, d_A1(j,0), ldda, d_A2(ipiv[j]-1,0), ldda );
#else
magma_zswap( N, d_A1(j,0), ldda, d_A2(ipiv[j]-1,0), ldda, opts.queue );
#endif
}
}
time = magma_sync_wtime( opts.queue ) - time;
col_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_zgetmatrix( N, N, d_A1(0,0), ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2(0,0), ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* zswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1(0,0), ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2(0,0), ldda );
time = magma_sync_wtime( opts.queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_zswap( N, d_A1(0,j), 1, d_A2(0,ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( opts.queue ) - time;
row_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1(0,j), &ione, h_A2(0,ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1(0,0), ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2(0,0), ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1(0,0), ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2(0,0), ldda );
time = magma_sync_wtime( opts.queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_zswap( N, d_A1(j,0), ldda, d_A2(ipiv[j]-1,0), ldda );
}
}
time = magma_sync_wtime( opts.queue ) - time;
col_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_zgetmatrix( N, N, d_A1(0,0), ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2(0,0), ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* zswapblk, blocked version (2 matrices)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1(0,0), ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2(0,0), ldda );
time = magma_sync_wtime( opts.queue );
magmablas_zswapblk( MagmaRowMajor, N, d_A1(0,0), ldda, d_A2(0,0), ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( opts.queue ) - time;
row_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1(0,j), &ione, h_A2(0,ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1(0,0), ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2(0,0), ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1(0,0), ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2(0,0), ldda );
time = magma_sync_wtime( opts.queue );
magmablas_zswapblk( MagmaColMajor, N, d_A1(0,0), ldda, d_A2(0,0), ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( opts.queue ) - time;
col_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1(j,0), &lda, h_A2(ipiv[j]-1,0), &lda);
}
}
magma_zgetmatrix( N, N, d_A1(0,0), ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2(0,0), ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style zlaswp (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1(0,0), ldda );
time = magma_sync_wtime( opts.queue );
magmablas_zlaswp( N, d_A1(0,0), ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( opts.queue ) - time;
row_perf4 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1(0,j), &ione, h_A1(0,ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1(0,0), ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style zlaswp (1 matrix) - d_ipiv on GPU
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1(0,0), ldda );
time = magma_sync_wtime( opts.queue );
magma_setvector( nb, sizeof(magma_int_t), ipiv, 1, d_ipiv(0), 1 );
magmablas_zlaswp2( N, d_A1(0,0), ldda, 1, nb, d_ipiv(0), 1 );
time = magma_sync_wtime( opts.queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1(0,j), &ione, h_A1(0,ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1(0,0), ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style zlaswpx (extended for row- and col-major) (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1(0,0), ldda );
time = magma_sync_wtime( opts.queue );
magmablas_zlaswpx( N, d_A1(0,0), ldda, 1, 1, nb, ipiv, 1);
time = magma_sync_wtime( opts.queue ) - time;
row_perf5 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1(0,j), &ione, h_A1(0,ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1(0,0), ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* Col Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1(0,0), ldda );
time = magma_sync_wtime( opts.queue );
magmablas_zlaswpx( N, d_A1(0,0), 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( opts.queue ) - time;
col_perf5 = gbytes / time;
/* LAPACK swap on CPU for comparison */
time = magma_wtime();
lapackf77_zlaswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
magma_zgetmatrix( N, N, d_A1(0,0), ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* Copy matrix.
*/
time = magma_sync_wtime( opts.queue );
magma_zcopymatrix( N, nb, d_A1(0,0), ldda, d_A2(0,0), ldda );
time = magma_sync_wtime( opts.queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
col_perf6 = 0.5 * gbytes / time;
time = magma_sync_wtime( opts.queue );
magma_zcopymatrix( nb, N, d_A1(0,0), ldda, d_A2(0,0), ldda );
time = magma_sync_wtime( opts.queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf4, ((check & 0x040) != 0 ? '*' : ' '),
row_perf7, ((check & 0x080) != 0 ? '*' : ' '),
row_perf5, ((check & 0x100) != 0 ? '*' : ' '),
col_perf5, ((check & 0x200) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failed") );
status += ! (check == 0);
TESTING_FREE_PIN( h_A1 );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_R1 );
TESTING_FREE_PIN( h_R2 );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( ipiv2 );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_A1 );
TESTING_FREE_DEV( d_A2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
extern "C" magma_int_t
magma_zgetrs_gpu(char trans, magma_int_t n, magma_int_t nrhs,
cuDoubleComplex *dA, magma_int_t ldda,
magma_int_t *ipiv,
cuDoubleComplex *dB, magma_int_t lddb,
magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
Solves a system of linear equations
A * X = B or A' * X = B
with a general N-by-N matrix A using the LU factorization computed by ZGETRF_GPU.
Arguments
=========
TRANS (input) CHARACTER*1
Specifies the form of the system of equations:
= 'N': A * X = B (No transpose)
= 'T': A'* X = B (Transpose)
= 'C': A'* X = B (Conjugate transpose = Transpose)
N (input) INTEGER
The order of the matrix A. N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
A (input) COMPLEX_16 array on the GPU, dimension (LDA,N)
The factors L and U from the factorization A = P*L*U as computed
by ZGETRF_GPU.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
IPIV (input) INTEGER array, dimension (N)
The pivot indices from ZGETRF; for 1<=i<=N, row i of the
matrix was interchanged with row IPIV(i).
B (input/output) COMPLEX_16 array on the GPU, dimension (LDB,NRHS)
On entry, the right hand side matrix B.
On exit, the solution matrix X.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
HWORK (workspace) COMPLEX_16 array, dimension N*NRHS
===================================================================== */
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex *work = NULL;
char trans_[2] = {trans, 0};
int notran = lapackf77_lsame(trans_, "N");
magma_int_t i1, i2, inc;
*info = 0;
if ( (! notran) &&
(! lapackf77_lsame(trans_, "T")) &&
(! lapackf77_lsame(trans_, "C")) ) {
*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' * 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;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zswap, zswapblk, zpermute, zlaswp, zlaswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaDoubleComplex *h_A1, *h_A2;
magmaDoubleComplex *d_A1, *d_A2;
magmaDoubleComplex *h_R1, *h_R2;
// row-major and column-major performance
real_Double_t row_perf0, col_perf0;
real_Double_t row_perf1, col_perf1;
real_Double_t row_perf2, col_perf2;
real_Double_t row_perf3;
real_Double_t row_perf4;
real_Double_t row_perf5, col_perf5;
real_Double_t row_perf6, col_perf6;
real_Double_t row_perf7;
real_Double_t cpu_perf;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magma_int_t *d_ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" cublasZswap zswap zswapblk zlaswp zpermute zlaswp2 zlaswpx zcopymatrix CPU (all in )\n");
printf(" N nb row-maj/col-maj row-maj/col-maj row-maj/col-maj row-maj row-maj row-maj row-maj/col-maj row-blk/col-blk zlaswp (GByte/s)\n");
printf("==================================================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// For an N x N matrix, swap nb rows or nb columns using various methods.
// Each test is assigned one bit in the 'check' bitmask; bit=1 indicates failure.
// The variable 'shift' keeps track of which bit is for current test
int shift = 1;
int check = 0;
N = opts.nsize[itest];
lda = N;
ldda = ((N+31)/32)*32;
nb = (opts.nb > 0 ? opts.nb : magma_get_zgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(magmaDoubleComplex) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_A2, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_R1, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_R2, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, nb );
TESTING_MALLOC_CPU( ipiv2, magma_int_t, nb );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, nb );
TESTING_MALLOC_DEV( d_A1, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( d_A2, magmaDoubleComplex, ldda*N );
for( j=0; j < nb; j++ ) {
ipiv[j] = (magma_int_t) ((rand()*1.*N) / (RAND_MAX * 1.)) + 1;
}
/* =====================================================================
* cublasZswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasZswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasZswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda);
}
}
time = magma_sync_wtime( queue ) - time;
col_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* zswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_zswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_zswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
}
}
time = magma_sync_wtime( queue ) - time;
col_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* zswapblk, blocked version (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_zswapblk( MagmaRowMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
row_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_zswapblk( MagmaColMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
col_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* zpermute_long (1 matrix)
*/
/* Row Major */
memcpy( ipiv2, ipiv, nb*sizeof(magma_int_t) ); // zpermute updates ipiv2
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_zpermute_long2( N, d_A1, ldda, ipiv2, nb, 0 );
time = magma_sync_wtime( queue ) - time;
row_perf3 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style zlaswp (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_zlaswp( N, d_A1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf4 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style zlaswp (1 matrix) - d_ipiv on GPU
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magma_setvector( nb, sizeof(magma_int_t), ipiv, 1, d_ipiv, 1 );
magmablas_zlaswp2( N, d_A1, ldda, 1, nb, d_ipiv );
time = magma_sync_wtime( queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style zlaswpx (extended for row- and col-major) (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_zlaswpx( N, d_A1, ldda, 1, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf5 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* Col Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_zlaswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
time = magma_wtime();
lapackf77_zlaswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* Copy matrix.
*/
time = magma_sync_wtime( queue );
magma_zcopymatrix( N, nb, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
col_perf6 = 0.5 * gbytes / time;
time = magma_sync_wtime( queue );
magma_zcopymatrix( nb, N, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf3, ((check & 0x040) != 0 ? '*' : ' '),
row_perf4, ((check & 0x080) != 0 ? '*' : ' '),
row_perf7, ((check & 0x100) != 0 ? '*' : ' '),
row_perf5, ((check & 0x200) != 0 ? '*' : ' '),
col_perf5, ((check & 0x400) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failed") );
status += ! (check == 0);
TESTING_FREE_PIN( h_A1 );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_R1 );
TESTING_FREE_PIN( h_R2 );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( ipiv2 );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_A1 );
TESTING_FREE_DEV( d_A2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
