/**
Purpose
-------
ZSSSSM applies the LU factorization update from a complex
matrix formed by a lower triangular IB-by-K tile L1 on top of a
M2-by-K tile L2 to a second complex matrix formed by a M1-by-N1
tile A1 on top of a M2-by-N2 tile A2 (N1 == N2).
This is the right-looking Level 2.5 BLAS version of the algorithm.
Arguments
---------
@param[in]
m1 INTEGER
The number of rows of the matrix A1. M1 >= 0.
@param[in]
n1 INTEGER
The number of columns of the matrix A1. N1 >= 0.
@param[in]
m2 INTEGER
The number of rows of the matrix A2. M2 >= 0.
@param[in]
n2 INTEGER
The number of columns of the matrix A2. N2 >= 0.
@param[in]
k INTEGER
The number of columns of the matrix L1 and L2. K >= 0.
@param[in]
ib INTEGER
The inner-blocking size. IB >= 0.
@param[in,out]
dA1 COMPLEX_16 array, dimension(LDDA1, N), on gpu.
On entry, the M1-by-N1 tile dA1.
On exit, dA1 is updated by the application of dL (dL1 dL2).
@param[in]
ldda1 INTEGER
The leading dimension of the array dA1. LDDA1 >= max(1,M1).
@param[in,out]
dA2 COMPLEX_16 array, dimension(LDDA2, N), on gpu.
On entry, the M2-by-N2 tile dA2.
On exit, dA2 is updated by the application of dL (dL1 dL2).
@param[in]
ldda2 INTEGER
The leading dimension of the array dA2. LDDA2 >= max(1,M2).
@param[in]
dL1 COMPLEX_16 array, dimension(LDDL1, K), on gpu.
The inverse of the IB-by-K lower triangular tile as returned by
ZTSTRF.
@param[in]
lddl1 INTEGER
The leading dimension of the array L1. LDDL1 >= max(1,2*IB).
@param[in]
dL2 COMPLEX_16 array, dimension(LDDL2, K)
The M2-by-K tile as returned by ZTSTRF.
@param[in]
lddl2 INTEGER
The leading dimension of the array L2. LDDL2 >= max(1,M2).
@param[in]
ipiv INTEGER array on the cpu.
The pivot indices array of size K as returned by ZTSTRF
@ingroup magma_zgesv_tile
********************************************************************/
extern "C" magma_int_t
magma_zssssm_gpu(
magma_order_t order, magma_int_t m1, magma_int_t n1,
magma_int_t m2, magma_int_t n2, magma_int_t k, magma_int_t ib,
magmaDoubleComplex_ptr dA1, magma_int_t ldda1,
magmaDoubleComplex_ptr dA2, magma_int_t ldda2,
magmaDoubleComplex_ptr dL1, magma_int_t lddl1,
magmaDoubleComplex_ptr dL2, magma_int_t lddl2,
magma_int_t *ipiv,
magma_int_t *info)
{
#define A1T(i,j) (dA1T + (i)*ldda1 + (j))
#define A2T(i,j) (dA2T + (i)*ldda2 + (j))
#define L1(i) (dL1 + (i)*lddl1 )
#define L2(i,j) (dL2 + (i)*lddl2i + (j)*lddl2j)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
int ip, ii, sb;
magmaDoubleComplex_ptr dA1T, dA2T;
magma_trans_t transL;
int lddl2i, lddl2j;
MAGMA_UNUSED( ip ); // used only if NOSWAPBLK
/* Check input arguments */
*info = 0;
if (m1 < 0) {
*info = -1;
}
else if (n1 < 0) {
*info = -2;
}
else if (m2 < 0) {
*info = -3;
}
else if (n2 < 0) {
*info = -4;
}
else if (k < 0) {
*info = -5;
}
else if (ib < 0) {
*info = -6;
}
else if (ldda1 < max(1,m1)) {
*info = -8;
}
else if (ldda2 < max(1,m2)) {
*info = -10;
}
else if (lddl1 < max(1,ib)) {
*info = -12;
}
else if (lddl2 < max(1,m2)) {
*info = -14;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ((m1 == 0) || (n1 == 0) || (m2 == 0) || (n2 == 0) || (k == 0) || (ib == 0))
return *info;
if ( order == MagmaColMajor ) {
magmablas_zgetmo_in( dA1, dA1T, ldda1, m1, n1 );
magmablas_zgetmo_in( dA2, dA2T, ldda2, m2, n2 );
transL = MagmaTrans;
lddl2i = 1; lddl2j = lddl2;
} else {
dA1T = dA1;
dA2T = dA2;
transL = MagmaNoTrans;
lddl2i = lddl2; lddl2j = 1;
}
ip = 0;
for( ii=0; ii < k; ii += ib ) {
sb = min( k-ii, ib);
#ifndef NOSWAPBLK
magmablas_zswapblk( MagmaRowMajor, n1,
A1T(0, 0), ldda1,
A2T(0, 0), ldda2,
ii+1, ii+ib, ipiv, 1, m1 );
#else
{
int im;
for (i=0; i < ib; i++) {
im = ipiv[ip]-1;
if (im != (ii+i)) {
im = im - m1;
assert( (im >= 0) && (im < m1) && (im < m2) );
magmablas_zswap( n1, A1T(ii+i, 0), 1, A2T(im, 0), 1 );
}
ip++;
}
}
#endif
#ifndef WITHOUTTRTRI
/* Lower, Trans, because L1 is not transposed */
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n1, sb,
c_one, L1( ii), lddl1,
A1T(ii, 0), ldda1);
#else
/* Lower, Trans, because L1 is not transposed */
magma_ztrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n1, sb,
c_one, L1( ii), lddl1,
A1T(ii, 0), ldda1);
#endif
/* Second parameter is trans because L2 is not transposed */
magma_zgemm( MagmaNoTrans, transL,
n2, m2, sb,
c_neg_one, A1T(ii, 0), ldda1,
L2( 0, ii), lddl2,
c_one, A2T(0, 0 ), ldda2 );
}
if ( order == MagmaColMajor ) {
magmablas_zgetmo_out( dA1, dA1T, ldda1, m1, n1 );
magmablas_zgetmo_out( dA2, dA2T, ldda2, m2, n2 );
}
return *info;
}
/**
Purpose
-------
ZSSSSM applies the LU factorization update from a complex
matrix formed by a lower triangular IB-by-K tile L1 on top of a
M2-by-K tile L2 to a second complex matrix formed by a M1-by-N1
tile A1 on top of a M2-by-N2 tile A2 (N1 == N2).
This is the right-looking Level 2.5 BLAS version of the algorithm.
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]
ib INTEGER
The inner-blocking size. IB >= 0.
@param[in]
NB INTEGER
The blocking size. NB >= 0.
@param[in,out]
hU COMPLEX_16 array, dimension(LDHU, N), on cpu.
On entry, the NB-by-N upper triangular tile hU.
On exit, the content is incomplete. Shouldn't be used.
@param[in]
ldhu INTEGER
The leading dimension of the array hU. LDHU >= max(1,NB).
@param[in,out]
dU COMPLEX_16 array, dimension(LDDU, N), on gpu.
On entry, the NB-by-N upper triangular tile dU identical to hU.
On exit, the new factor U from the factorization.
@param[in]
lddu INTEGER
The leading dimension of the array dU. LDDU >= max(1,NB).
@param[in,out]
hA COMPLEX_16 array, dimension(LDHA, N), on cpu.
On entry, only the M-by-IB first panel needs to be identical to dA(1..M, 1..IB).
On exit, the content is incomplete. Shouldn't be used.
@param[in]
ldha INTEGER
The leading dimension of the array hA. LDHA >= max(1,M).
@param[in,out]
dA COMPLEX_16 array, dimension(LDDA, N), on gpu.
On entry, the M-by-N tile to be factored.
On exit, the factor L from the factorization
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
@param[out]
hL COMPLEX_16 array, dimension(LDHL, K), on vpu.
On exit, contains in the upper part the IB-by-K lower triangular tile,
and in the lower part IB-by-K the inverse of the top part.
@param[in]
ldhl INTEGER
The leading dimension of the array hL. LDHL >= max(1,2*IB).
@param[out]
dL COMPLEX_16 array, dimension(LDDL, K), on gpu.
On exit, contains in the upper part the IB-by-K lower triangular tile,
and in the lower part IB-by-K the inverse of the top part.
@param[in]
lddl INTEGER
The leading dimension of the array dL. LDDL >= max(1,2*IB).
@param[out]
hWORK COMPLEX_16 array, dimension(LDHWORK, 2*IB), on cpu.
Workspace.
@param[in]
ldhwork INTEGER
The leading dimension of the array hWORK. LDHWORK >= max(NB, 1).
@param[out]
dWORK COMPLEX_16 array, dimension(LDDWORK, 2*IB), on gpu.
Workspace.
@param[in]
lddwork INTEGER
The leading dimension of the array dWORK. LDDWORK >= max(NB, 1).
@param[out]
ipiv INTEGER array on the cpu.
The pivot indices array of size K as returned by ZTSTRF
@param[out]
info INTEGER
- PLASMA_SUCCESS successful exit
- < 0 if INFO = -k, the k-th argument had an illegal value
- > 0 if INFO = k, U(k,k) 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_tile
********************************************************************/
extern "C" magma_int_t
magma_ztstrf_gpu(
magma_order_t order, magma_int_t m, magma_int_t n,
magma_int_t ib, magma_int_t nb,
magmaDoubleComplex *hU, magma_int_t ldhu,
magmaDoubleComplex_ptr dU, magma_int_t lddu,
magmaDoubleComplex *hA, magma_int_t ldha,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex *hL, magma_int_t ldhl,
magmaDoubleComplex_ptr dL, magma_int_t lddl,
magma_int_t *ipiv,
magmaDoubleComplex *hwork, magma_int_t ldhwork,
magmaDoubleComplex_ptr dwork, magma_int_t lddwork,
magma_int_t *info)
{
#define UT(i,j) (dUT + (i)*ib*lddu + (j)*ib )
#define AT(i,j) (dAT + (i)*ib*ldda + (j)*ib )
#define L(i) (dL + (i)*ib*lddl )
#define L2(i) (dL2 + (i)*ib*lddl )
#define hU(i,j) (hU + (j)*ib*ldhu + (i)*ib )
#define hA(i,j) (hA + (j)*ib*ldha + (i)*ib )
#define hL(i) (hL + (i)*ib*ldhl )
#define hL2(i) (hL2 + (i)*ib*ldhl )
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
int iinfo = 0;
int maxm, mindim;
int i, j, im, s, ip, ii, sb, p = 1;
magmaDoubleComplex_ptr dAT, dUT;
magmaDoubleComplex_ptr dAp, dUp;
#ifndef WITHOUTTRTRI
magmaDoubleComplex_ptr dL2 = dL + ib;
magmaDoubleComplex *hL2 = hL + ib;
p = 2;
#endif
/* Check input arguments */
*info = 0;
if (m < 0) {
*info = -1;
}
else if (n < 0) {
*info = -2;
}
else if (ib < 0) {
*info = -3;
}
else if ((lddu < max(1,m)) && (m > 0)) {
*info = -6;
}
else if ((ldda < max(1,m)) && (m > 0)) {
*info = -8;
}
else if ((lddl < max(1,ib)) && (ib > 0)) {
*info = -10;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* quick return */
if ((m == 0) || (n == 0) || (ib == 0))
return *info;
ip = 0;
/* Function Body */
mindim = min(m, n);
s = mindim / ib;
if ( ib >= mindim ) {
/* Use CPU code. */
CORE_ztstrf(m, n, ib, nb,
(PLASMA_Complex64_t*)hU, ldhu,
(PLASMA_Complex64_t*)hA, ldha,
(PLASMA_Complex64_t*)hL, ldhl,
ipiv,
(PLASMA_Complex64_t*)hwork, ldhwork,
info);
#ifndef WITHOUTTRTRI
CORE_zlacpy( PlasmaUpperLower, mindim, mindim,
(PLASMA_Complex64_t*)hL, ldhl,
(PLASMA_Complex64_t*)hL2, ldhl );
CORE_ztrtri( PlasmaLower, PlasmaUnit, mindim,
(PLASMA_Complex64_t*)hL2, ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
#endif
if ( order == MagmaRowMajor ) {
magma_zsetmatrix( m, n, hU, ldhu, dwork, lddwork );
magmablas_ztranspose( m, n, dwork, lddwork, dU, lddu );
magma_zsetmatrix( m, n, hA, ldha, dwork, lddwork );
magmablas_ztranspose( m, n, dwork, lddwork, dA, ldda );
} else {
magma_zsetmatrix( m, n, hU, ldhu, dU, lddu );
magma_zsetmatrix( m, n, hA, ldha, dA, ldda );
}
magma_zsetmatrix( p*ib, n, hL, ldhl, dL, lddl );
}
else {
/* Use hybrid blocked code. */
maxm = magma_roundup( m, 32 );
if ( order == MagmaColMajor ) {
magmablas_zgetmo_in( dU, dUT, lddu, m, n );
magmablas_zgetmo_in( dA, dAT, ldda, m, n );
} else {
dUT = dU; dAT = dA;
}
dAp = dwork;
dUp = dAp + ib*lddwork;
ip = 0;
for( i=0; i < s; i++ ) {
ii = i * ib;
sb = min(mindim-ii, ib);
if ( i > 0 ) {
// download i-th panel
magmablas_ztranspose( sb, ii, UT(0,i), lddu, dUp, lddu );
magmablas_ztranspose( sb, m, AT(0,i), ldda, dAp, ldda );
magma_zgetmatrix( ii, sb, dUp, lddu, hU(0, i), ldhu );
magma_zgetmatrix( m, sb, dAp, ldda, hA(0, i), ldha );
// make sure that gpu queue is empty
//magma_device_sync();
#ifndef WITHOUTTRTRI
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-(ii+sb), ib,
c_one, L2(i-1), lddl,
UT(i-1, i+1), lddu);
#else
magma_ztrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-(ii+sb), ib,
c_one, L(i-1), lddl,
UT(i-1, i+1), lddu);
#endif
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n-(ii+sb), m, ib,
c_neg_one, UT(i-1, i+1), lddu,
AT(0, i-1), ldda,
c_one, AT(0, i+1), ldda );
}
// do the cpu part
CORE_ztstrf(m, sb, ib, nb,
(PLASMA_Complex64_t*)hU(i, i), ldhu,
(PLASMA_Complex64_t*)hA(0, i), ldha,
(PLASMA_Complex64_t*)hL(i), ldhl,
ipiv+ii,
(PLASMA_Complex64_t*)hwork, ldhwork,
info);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + ii;
// Need to swap betw U and A
#ifndef NOSWAPBLK
magmablas_zswapblk( MagmaRowMajor, n-(ii+sb),
UT(i, i+1), lddu,
AT(0, i+1), ldda,
1, sb, ipiv+ii, 1, nb );
for (j=0; j < ib; j++) {
im = ipiv[ip]-1;
if ( im == j ) {
ipiv[ip] += ii;
}
ip++;
}
#else
for (j=0; j < ib; j++) {
im = ipiv[ip]-1;
if ( im != (j) ) {
im = im - nb;
assert( (im >= 0) && (im < m) );
magmablas_zswap( n-(ii+sb), UT(i, i+1)+j*lddu, 1, AT(0, i+1)+im*ldda, 1 );
} else {
ipiv[ip] += ii;
}
ip++;
}
#endif
#ifndef WITHOUTTRTRI
CORE_zlacpy( PlasmaUpperLower, sb, sb,
(PLASMA_Complex64_t*)hL(i), ldhl,
(PLASMA_Complex64_t*)hL2(i), ldhl );
CORE_ztrtri( PlasmaLower, PlasmaUnit, sb,
(PLASMA_Complex64_t*)hL2(i), ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
#endif
// upload i-th panel
magma_zsetmatrix( sb, sb, hU(i, i), ldhu, dUp, lddu );
magma_zsetmatrix( m, sb, hA(0, i), ldha, dAp, ldda );
magma_zsetmatrix( p*ib, sb, hL(i), ldhl, L(i), lddl );
magmablas_ztranspose( sb, sb, dUp, lddu, UT(i,i), lddu );
magmablas_ztranspose( m, sb, dAp, ldda, AT(0,i), ldda );
// make sure that gpu queue is empty
//magma_device_sync();
// do the small non-parallel computations
if ( s > (i+1) ) {
#ifndef WITHOUTTRTRI
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
sb, sb,
c_one, L2(i), lddl,
UT(i, i+1), lddu);
#else
magma_ztrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
sb, sb,
c_one, L(i), lddl,
UT(i, i+1), lddu);
#endif
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
sb, m, sb,
c_neg_one, UT(i, i+1), lddu,
AT(0, i ), ldda,
c_one, AT(0, i+1), ldda );
}
else {
#ifndef WITHOUTTRTRI
magma_ztrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-mindim, sb,
c_one, L2(i), lddl,
UT(i, i+1), lddu);
#else
magma_ztrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-mindim, sb,
c_one, L(i), lddl,
UT(i, i+1), lddu);
#endif
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n-mindim, m, sb,
c_neg_one, UT(i, i+1), lddu,
AT(0, i ), ldda,
c_one, AT(0, i+1), ldda );
}
}
if ( order == MagmaColMajor ) {
magmablas_zgetmo_out( dU, dUT, lddu, m, n );
magmablas_zgetmo_out( dA, dAT, ldda, m, n );
}
}
return *info;
}
extern "C" magma_int_t
magma_zlaqps_gpu(magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaDoubleComplex *A, magma_int_t lda,
magma_int_t *jpvt, magmaDoubleComplex *tau,
double *vn1, double *vn2,
magmaDoubleComplex *auxv,
magmaDoubleComplex *F, magma_int_t ldf)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZLAQPS computes a step of QR factorization with column pivoting
of a complex M-by-N matrix A by using Blas-3. It tries to factorize
NB columns from A starting from the row OFFSET+1, and updates all
of the matrix with Blas-3 xGEMM.
In some cases, due to catastrophic cancellations, it cannot
factorize NB columns. Hence, the actual number of factorized
columns is returned in KB.
Block A(1:OFFSET,1:N) is accordingly pivoted, but not factorized.
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
OFFSET (input) INTEGER
The number of rows of A that have been factorized in
previous steps.
NB (input) INTEGER
The number of columns to factorize.
KB (output) INTEGER
The number of columns actually factorized.
A (input/output) COMPLEX*16 array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, block A(OFFSET+1:M,1:KB) is the triangular
factor obtained and block A(1:OFFSET,1:N) has been
accordingly pivoted, but no factorized.
The rest of the matrix, block A(OFFSET+1:M,KB+1:N) has
been updated.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
JPVT (input/output) INTEGER array, dimension (N)
JPVT(I) = K <==> Column K of the full matrix A has been
permuted into position I in AP.
TAU (output) COMPLEX*16 array, dimension (KB)
The scalar factors of the elementary reflectors.
VN1 (input/output) DOUBLE PRECISION array, dimension (N)
The vector with the partial column norms.
VN2 (input/output) DOUBLE PRECISION array, dimension (N)
The vector with the exact column norms.
AUXV (input/output) COMPLEX*16 array, dimension (NB)
Auxiliar vector.
F (input/output) COMPLEX*16 array, dimension (LDF,NB)
Matrix F' = L*Y'*A.
LDF (input) INTEGER
The leading dimension of the array F. LDF >= max(1,N).
===================================================================== */
#define A(i, j) (A + (i) + (j)*(lda ))
#define F(i, j) (F + (i) + (j)*(ldf ))
magmaDoubleComplex c_zero = MAGMA_Z_MAKE( 0.,0.);
magmaDoubleComplex c_one = MAGMA_Z_MAKE( 1.,0.);
magmaDoubleComplex c_neg_one = MAGMA_Z_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
//double d__1;
magmaDoubleComplex z__1;
//magma_int_t j;
magma_int_t k, rk;
//magmaDoubleComplex Akk;
magmaDoubleComplex *Aks;
magmaDoubleComplex tauk;
magma_int_t pvt;
//double temp, temp2;
double tol3z;
magma_int_t itemp;
double lsticc, *lsticcs;
magma_int_t lastrk;
magma_dmalloc( &lsticcs, 1+256*(n+255)/256 );
lastrk = min( m, n + offset );
tol3z = magma_dsqrt( lapackf77_dlamch("Epsilon"));
lsticc = 0;
k = 0;
magma_zmalloc( &Aks, nb );
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// Fortran: pvt, k, idamax are all 1-based; subtract 1 from k.
// C: pvt, k, idamax are all 0-based; don't subtract 1.
pvt = k - 1 + magma_idamax( n-k, &vn1[k], ione );
if (pvt != k) {
/*if (pvt >= nb) {
// 1. Start copy from GPU
magma_zgetmatrix_async( m - offset - nb, 1,
dA(offset + nb, pvt), ldda,
A (offset + nb, pvt), lda, stream );
}*/
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
/*if (pvt < nb){
// no need of transfer if pivot is within the panel
blasf77_zswap( &m, A(0, pvt), &ione, A(0, k), &ione );
}
else {
// 1. Finish copy from GPU
magma_queue_sync( stream );
// 2. Swap as usual on CPU
blasf77_zswap(&m, A(0, pvt), &ione, A(0, k), &ione);
// 3. Restore the GPU
magma_zsetmatrix_async( m - offset - nb, 1,
A (offset + nb, pvt), lda,
dA(offset + nb, pvt), ldda, stream);
}*/
magmablas_zswap( m, A(0, pvt), ione, A(0, k), ione );
//blasf77_zswap( &i__1, F(pvt,0), &ldf, F(k,0), &ldf );
magmablas_zswap( i__1, F(pvt, 0), ldf, F(k, 0), ldf);
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
//vn1[pvt] = vn1[k];
//vn2[pvt] = vn2[k];
#if defined(PRECISION_d) || defined(PRECISION_z)
//magma_dswap( 1, &vn1[pvt], 1, &vn1[k], 1 );
//magma_dswap( 1, &vn2[pvt], 1, &vn2[k], 1 );
magma_dswap( 2, &vn1[pvt], n+offset, &vn1[k], n+offset );
#else
//magma_sswap( 1, &vn1[pvt], 1, &vn1[k], 1 );
//magma_sswap( 1, &vn2[pvt], 1, &vn2[k], 1 );
magma_sswap(2, &vn1[pvt], n+offset, &vn1[k], n+offset);
#endif
}
/* Apply previous Householder reflectors to column K:
A(RK:M,K) := A(RK:M,K) - A(RK:M,1:K-1)*F(K,1:K-1)'.
Optimization: multiply with beta=0; wait for vector and subtract */
if (k > 0) {
/*#if (defined(PRECISION_c) || defined(PRECISION_z))
for (j = 0; j < k; ++j){
*F(k,j) = MAGMA_Z_CNJG( *F(k,j) );
}
#endif*/
//#define RIGHT_UPDATE
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_zgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, A(offset+nb, 0), lda,
F(k, 0), ldf,
c_one, A(offset+nb, k), ione );
#else
i__1 = m - rk;
i__2 = k;
/*blasf77_zgemv( MagmaNoTransStr, &i__1, &i__2,
&c_neg_one, A(rk, 0), &lda,
F(k, 0), &ldf,
&c_one, A(rk, k), &ione );*/
magma_zgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, A(rk, 0), lda,
F(k, 0), ldf,
c_one, A(rk, k), ione );
#endif
/*#if (defined(PRECISION_c) || defined(PRECISION_z))
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_Z_CNJG( *F(k,j) );
}
#endif*/
}
/* Generate elementary reflector H(k). */
magma_zlarfg_gpu(m-rk, A(rk, k), A(rk + 1, k), &tau[k], &vn1[k], &Aks[k]);
//Akk = *A(rk, k);
//*A(rk, k) = c_one;
//magma_zgetvector( 1, &Aks[k], 1, &Akk, 1 );
/* needed to avoid the race condition */
if (k == 0) magma_zsetvector( 1, &c_one, 1, A(rk, k), 1 );
else magma_zcopymatrix( 1, 1, A(offset, 0), 1, A(rk, k), 1 );
/* Compute Kth column of F:
Compute F(K+1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) on the GPU */
if (k < n-1 || k > 0) magma_zgetvector( 1, &tau[k], 1, &tauk, 1 );
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Send the vector to the GPU */
//magma_zsetmatrix( i__1, 1, A(rk, k), lda, dA(rk,k), ldda );
/* Multiply on GPU */
// was CALL ZGEMV( 'Conjugate transpose', M-RK+1, N-K,
// TAU( K ), A( RK, K+1 ), LDA,
// A( RK, K ), 1,
// CZERO, F( K+1, K ), 1 )
//magma_zgetvector( 1, &tau[k], 1, &tauk, 1 );
magma_zgemv( MagmaConjTrans, m-rk, n-k-1,
tauk, A( rk, k+1 ), lda,
A( rk, k ), 1,
c_zero, F( k+1, k ), 1 );
//magma_zscal( m-rk, tau[k], F( k+1, k), 1 );
//magma_int_t i__3 = nb-k-1;
//magma_int_t i__4 = i__2 - i__3;
//magma_int_t i__5 = nb-k;
//magma_zgemv( MagmaConjTrans, i__1 - i__5, i__2 - i__3,
// tau[k], dA(rk +i__5, k+1+i__3), ldda,
// dA(rk +i__5, k ), ione,
// c_zero, dF(k+1+i__3, k ), ione );
//magma_zgetmatrix_async( i__2-i__3, 1,
// dF(k + 1 +i__3, k), i__2,
// F (k + 1 +i__3, k), i__2, stream );
//blasf77_zgemv( MagmaConjTransStr, &i__1, &i__3,
// &tau[k], A(rk, k+1), &lda,
// A(rk, k ), &ione,
// &c_zero, F(k+1, k ), &ione );
//magma_queue_sync( stream );
//blasf77_zgemv( MagmaConjTransStr, &i__5, &i__4,
// &tau[k], A(rk, k+1+i__3), &lda,
// A(rk, k ), &ione,
// &c_one, F(k+1+i__3, k ), &ione );
}
/* Padding F(1:K,K) with zeros.
for (j = 0; j <= k; ++j) {
magma_zsetvector( 1, &c_zero, 1, F(j, k), 1 );
}*/
/* Incremental updating of F:
F(1:N,K) := F(1:N,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K).
F(1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K)
:= tau(K)(A(RK:M,K+1:N)' - F(1:N,1:K-1)*A(RK:M,1:K-1)') A(RK:M,K)
so, F is (updated A)*V */
//if (k > 0 && k<n-1) {
if (k > 0) {
//magma_zgetvector( 1, &tau[k], 1, &tauk, 1 );
z__1 = MAGMA_Z_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_zgemv( MagmaConjTrans, i__1, i__2,
z__1, A(offset+nb, 0), lda,
A(offset+nb, k), ione,
c_zero, auxv, ione );
i__1 = k;
magma_zgemv( MagmaNoTrans, n-k-1, i__1,
c_one, F(k+1,0), ldf,
auxv, ione,
c_one, F(k+1,k), ione );
#else
i__1 = m - rk;
i__2 = k;
//blasf77_zgemv( MagmaConjTransStr, &i__1, &i__2,
// &z__1, A(rk, 0), &lda,
// A(rk, k), &ione,
// &c_zero, auxv, &ione );
magma_zgemv( MagmaConjTrans, i__1, i__2,
z__1, A(rk, 0), lda,
A(rk, k), ione,
c_zero, auxv, ione );
//i__1 = k;
//blasf77_zgemv( MagmaNoTransStr, &n, &i__1,
// &c_one, F(0,0), &ldf,
// auxv, &ione,
// &c_one, F(0,k), &ione );
/*magma_zgemv( MagmaNoTrans, n, i__1,
c_one, F(0,0), ldf,
auxv, ione,
c_one, F(0,k), ione );*/
/* I think we only need stricly lower-triangular part :) */
magma_zgemv( MagmaNoTrans, n-k-1, i__2,
c_one, F(k+1,0), ldf,
auxv, ione,
c_one, F(k+1,k), ione );
#endif
}
/* Optimization: On the last iteration start sending F back to the GPU */
/* Update the current row of A:
A(RK,K+1:N) := A(RK,K+1:N) - A(RK,1:K)*F(K+1:N,1:K)'. */
if (k < n-1) {
i__1 = n - k - 1;
i__2 = k + 1;
//blasf77_zgemm( MagmaNoTransStr, MagmaConjTransStr, &ione, &i__1, &i__2,
// &c_neg_one, A(rk, 0 ), &lda,
// F(k+1,0 ), &ldf,
// &c_one, A(rk, k+1), &lda );
#ifdef RIGHT_UPDATE
/* right-looking update of rows, */
magma_zgemm( MagmaNoTrans, MagmaConjTrans, nb-k, i__1, ione,
c_neg_one, A(rk, k ), lda,
F(k+1, k ), ldf,
c_one, A(rk, k+1), lda );
#else
/* left-looking update of rows, *
* since F=A'v with original A, so no right-looking */
magma_zgemm( MagmaNoTrans, MagmaConjTrans, ione, i__1, i__2,
c_neg_one, A(rk, 0 ), lda,
F(k+1,0 ), ldf,
c_one, A(rk, k+1), lda );
#endif
}
/* Update partial column norms. */
if (rk < min(m, n+offset)-1 ) {
magmablas_dznrm2_row_check_adjust(n-k-1, tol3z, &vn1[k+1], &vn2[k+1], A(rk,k+1), lda, lsticcs);
magma_device_sync();
#if defined(PRECISION_d) || defined(PRECISION_z)
magma_dgetvector( 1, &lsticcs[0], 1, &lsticc, 1 );
#else
magma_sgetvector( 1, &lsticcs[0], 1, &lsticc, 1 );
#endif
}
/*if (rk < lastrk) {
for (j = k + 1; j < n; ++j) {
if (vn1[j] != 0.) {
// NOTE: The following 4 lines follow from the analysis in
// Lapack Working Note 176.
temp = MAGMA_Z_ABS( *A(rk,j) ) / vn1[j];
temp = max( 0., ((1. + temp) * (1. - temp)) );
d__1 = vn1[j] / vn2[j];
temp2 = temp * (d__1 * d__1);
if (temp2 <= tol3z) {
vn2[j] = (double) lsticc;
lsticc = j;
} else {
vn1[j] *= magma_dsqrt(temp);
}
}
}
}*/
//*A(rk, k) = Akk;
//magma_zsetvector( 1, &Akk, 1, A(rk, k), 1 );
//magma_zswap( 1, &Aks[k], 1, A(rk, k), 1 );
++k;
}
magma_zcopymatrix( 1, k, Aks, 1, A(offset, 0), lda+1 );
// leave k as the last column done
--k;
*kb = k + 1;
rk = offset + *kb - 1;
/* Apply the block reflector to the rest of the matrix:
A(OFFSET+KB+1:M,KB+1:N) := A(OFFSET+KB+1:M,KB+1:N) - A(OFFSET+KB+1:M,1:KB)*F(KB+1:N,1:KB)' */
if (*kb < min(n, m - offset)) {
i__1 = m - rk - 1;
i__2 = n - *kb;
/* Send F to the GPU
magma_zsetmatrix( i__2, *kb,
F (*kb, 0), ldf,
dF(*kb, 0), i__2 );*/
magma_zgemm( MagmaNoTrans, MagmaConjTrans, i__1, i__2, *kb,
c_neg_one, A(rk+1, 0 ), lda,
F(*kb, 0 ), ldf,
c_one, A(rk+1, *kb), lda );
}
/* Recomputation of difficult columns. */
if( lsticc > 0 ) {
printf( " -- recompute dnorms --\n" );
magmablas_dznrm2_check(m-rk-1, n-*kb, A(rk+1,*kb), lda,
&vn1[*kb], lsticcs);
#if defined(PRECISION_d) || defined(PRECISION_z)
magma_dcopymatrix( n-*kb, 1, &vn1[*kb], *kb, &vn2[*kb], *kb);
#else
magma_scopymatrix( n-*kb, 1, &vn1[*kb], *kb, &vn2[*kb], *kb);
#endif
/*while( lsticc > 0 ) {
itemp = (magma_int_t)(vn2[lsticc] >= 0. ? floor(vn2[lsticc] + .5) : -floor(.5 - vn2[lsticc]));
i__1 = m - rk - 1;
if (lsticc <= nb)
vn1[lsticc] = cblas_dznrm2(i__1, A(rk + 1, lsticc), ione);
else {
// Where is the data, CPU or GPU ?
double r1, r2;
r1 = cblas_dznrm2(nb-k, A(rk + 1, lsticc), ione);
r2 = magma_dznrm2(m-offset-nb, dA(offset + nb + 1, lsticc), ione);
vn1[lsticc] = magma_dsqrt(r1*r1+r2*r2);
}
// NOTE: The computation of VN1( LSTICC ) relies on the fact that
// SNRM2 does not fail on vectors with norm below the value of SQRT(DLAMCH('S'))
vn2[lsticc] = vn1[lsticc];
lsticc = itemp;*/
}
magma_free(Aks);
magma_free(lsticcs);
return MAGMA_SUCCESS;
} /* magma_zlaqps */
/* ////////////////////////////////////////////////////////////////////////////
-- 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;
}
/**
Purpose
-------
ZGETRI computes the inverse of a matrix using the LU factorization
computed by ZGETRF. This method inverts U and then computes inv(A) by
solving the system inv(A)*L = inv(U) for inv(A).
Note that it is generally both faster and more accurate to use ZGESV,
or ZGETRF and ZGETRS, to solve the system AX = B, rather than inverting
the matrix and multiplying to form X = inv(A)*B. Only in special
instances should an explicit inverse be computed with this routine.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX_16 array on the GPU, dimension (LDDA,N)
On entry, the factors L and U from the factorization
A = P*L*U as computed by ZGETRF_GPU.
On exit, if INFO = 0, the inverse of the original matrix A.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= 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[out]
dwork (workspace) COMPLEX_16 array on the GPU, dimension (MAX(1,LWORK))
@param[in]
lwork INTEGER
The dimension of the array DWORK. LWORK >= N*NB, where NB is
the optimal blocksize returned by magma_get_zgetri_nb(n).
\n
Unlike LAPACK, this version does not currently support a
workspace query, because the workspace is on the GPU.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, U(i,i) is exactly zero; the matrix is
singular and its cannot be computed.
@ingroup magma_zgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_zgetri_gpu( magma_int_t n, magmaDoubleComplex *dA, magma_int_t ldda,
magma_int_t *ipiv, magmaDoubleComplex *dwork, magma_int_t lwork,
magma_int_t *info )
{
#define dA(i, j) (dA + (i) + (j)*ldda)
#define dL(i, j) (dL + (i) + (j)*lddl)
/* Local variables */
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *dL = dwork;
magma_int_t lddl = n;
magma_int_t nb = magma_get_zgetri_nb(n);
magma_int_t j, jmax, jb, jp;
*info = 0;
if (n < 0)
*info = -1;
else if (ldda < max(1,n))
*info = -3;
else if ( lwork < n*nb )
*info = -6;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if ( n == 0 )
return *info;
/* Invert the triangular factor U */
magma_ztrtri_gpu( MagmaUpper, MagmaNonUnit, n, dA, ldda, info );
if ( *info != 0 )
return *info;
jmax = ((n-1) / nb)*nb;
for( j = jmax; j >= 0; j -= nb ) {
jb = min( nb, n-j );
// copy current block column of A to work space dL
// (only needs lower trapezoid, but we also copy upper triangle),
// then zero the strictly lower trapezoid block column of A.
magmablas_zlacpy( MagmaFull, n-j, jb,
dA(j,j), ldda,
dL(j,0), lddl );
magmablas_zlaset( MagmaLower, n-j-1, jb, c_zero, c_zero, dA(j+1,j), ldda );
// compute current block column of Ainv
// Ainv(:, j:j+jb-1)
// = ( U(:, j:j+jb-1) - Ainv(:, j+jb:n) L(j+jb:n, j:j+jb-1) )
// * L(j:j+jb-1, j:j+jb-1)^{-1}
// where L(:, j:j+jb-1) is stored in dL.
if ( j+jb < n ) {
magma_zgemm( MagmaNoTrans, MagmaNoTrans, n, jb, n-j-jb,
c_neg_one, dA(0,j+jb), ldda,
dL(j+jb,0), lddl,
c_one, dA(0,j), ldda );
}
// TODO use magmablas work interface
magma_ztrsm( MagmaRight, MagmaLower, MagmaNoTrans, MagmaUnit,
n, jb, c_one,
dL(j,0), lddl,
dA(0,j), ldda );
}
// Apply column interchanges
for( j = n-2; j >= 0; --j ) {
jp = ipiv[j] - 1;
if ( jp != j ) {
magmablas_zswap( n, dA(0,j), 1, dA(0,jp), 1 );
}
}
return *info;
}
/**
@deprecated
Purpose
-------
ZLAQPS computes a step of QR factorization with column pivoting
of a complex M-by-N matrix A by using Blas-3. It tries to factorize
NB columns from A starting from the row OFFSET+1, and updates all
of the matrix with Blas-3 xGEMM.
In some cases, due to catastrophic cancellations, it cannot
factorize NB columns. Hence, the actual number of factorized
columns is returned in KB.
Block A(1:OFFSET,1:N) is accordingly pivoted, but not factorized.
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]
offset INTEGER
The number of rows of A that have been factorized in
previous steps.
@param[in]
nb INTEGER
The number of columns to factorize.
@param[out]
kb INTEGER
The number of columns actually factorized.
@param[in,out]
dA COMPLEX_16 array, dimension (LDDA,N), on the GPU.
On entry, the M-by-N matrix A.
On exit, block A(OFFSET+1:M,1:KB) is the triangular
factor obtained and block A(1:OFFSET,1:N) has been
accordingly pivoted, but no factorized.
The rest of the matrix, block A(OFFSET+1:M,KB+1:N) has
been updated.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[in,out]
jpvt INTEGER array, dimension (N)
JPVT(I) = K <==> Column K of the full matrix A has been
permuted into position I in AP.
@param[out]
tau COMPLEX_16 array, dimension (KB)
The scalar factors of the elementary reflectors.
@param[in,out]
vn1 DOUBLE PRECISION array, dimension (N)
The vector with the partial column norms.
@param[in,out]
vn2 DOUBLE PRECISION array, dimension (N)
The vector with the exact column norms.
@param[in,out]
dauxv COMPLEX_16 array, dimension (NB), on the GPU
Auxiliary vector.
@param[in,out]
dF COMPLEX_16 array, dimension (LDDF,NB), on the GPU
Matrix F' = L*Y'*A.
@param[in]
lddf INTEGER
The leading dimension of the array F. LDDF >= max(1,N).
@ingroup magma_zgeqp3_aux
********************************************************************/
extern "C" magma_int_t
magma_zlaqps_gpu(
magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magma_int_t *jpvt, magmaDoubleComplex *tau,
double *vn1, double *vn2,
magmaDoubleComplex_ptr dauxv,
magmaDoubleComplex_ptr dF, magma_int_t lddf)
{
#define dA(i, j) (dA + (i) + (j)*(ldda))
#define dF(i, j) (dF + (i) + (j)*(lddf))
magmaDoubleComplex c_zero = MAGMA_Z_MAKE( 0.,0.);
magmaDoubleComplex c_one = MAGMA_Z_MAKE( 1.,0.);
magmaDoubleComplex c_neg_one = MAGMA_Z_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
//double d__1;
magmaDoubleComplex z__1;
//magma_int_t j;
magma_int_t k, rk;
//magmaDoubleComplex Akk;
magmaDoubleComplex_ptr dAks;
magmaDoubleComplex tauk = MAGMA_Z_ZERO;
magma_int_t pvt;
//double temp, temp2;
double tol3z;
magma_int_t itemp;
double lsticc;
magmaDouble_ptr dlsticcs;
magma_dmalloc( &dlsticcs, 1+256*(n+255)/256 );
//lastrk = min( m, n + offset );
tol3z = magma_dsqrt( lapackf77_dlamch("Epsilon"));
lsticc = 0;
k = 0;
magma_zmalloc( &dAks, nb );
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// subtract 1 from Fortran/CUBLAS idamax; pvt, k are 0-based.
pvt = k + magma_idamax( n-k, &vn1[k], ione ) - 1;
if (pvt != k) {
/*if (pvt >= nb) {
// 1. Start copy from GPU
magma_zgetmatrix_async( m - offset - nb, 1,
dA(offset + nb, pvt), ldda,
A (offset + nb, pvt), lda, stream );
}*/
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
/*if (pvt < nb) {
// no need of transfer if pivot is within the panel
blasf77_zswap( &m, A(0, pvt), &ione, A(0, k), &ione );
}
else {
// 1. Finish copy from GPU
magma_queue_sync( stream );
// 2. Swap as usual on CPU
blasf77_zswap(&m, A(0, pvt), &ione, A(0, k), &ione);
// 3. Restore the GPU
magma_zsetmatrix_async( m - offset - nb, 1,
A (offset + nb, pvt), lda,
dA(offset + nb, pvt), ldda, stream);
}*/
magmablas_zswap( m, dA(0, pvt), ione, dA(0, k), ione );
//blasf77_zswap( &i__1, F(pvt,0), &ldf, F(k,0), &ldf );
magmablas_zswap( i__1, dF(pvt, 0), lddf, dF(k, 0), lddf);
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
//vn1[pvt] = vn1[k];
//vn2[pvt] = vn2[k];
#if defined(PRECISION_d) || defined(PRECISION_z)
//magma_dswap( 1, &vn1[pvt], 1, &vn1[k], 1 );
//magma_dswap( 1, &vn2[pvt], 1, &vn2[k], 1 );
magma_dswap( 2, &vn1[pvt], n+offset, &vn1[k], n+offset );
#else
//magma_sswap( 1, &vn1[pvt], 1, &vn1[k], 1 );
//magma_sswap( 1, &vn2[pvt], 1, &vn2[k], 1 );
magma_sswap(2, &vn1[pvt], n+offset, &vn1[k], n+offset);
#endif
}
/* Apply previous Householder reflectors to column K:
A(RK:M,K) := A(RK:M,K) - A(RK:M,1:K-1)*F(K,1:K-1)'.
Optimization: multiply with beta=0; wait for vector and subtract */
if (k > 0) {
/*#if (defined(PRECISION_c) || defined(PRECISION_z))
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_Z_CNJG( *F(k,j) );
}
#endif*/
//#define RIGHT_UPDATE
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_zgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, A(offset+nb, 0), lda,
F(k, 0), ldf,
c_one, A(offset+nb, k), ione );
#else
i__1 = m - rk;
i__2 = k;
/*blasf77_zgemv( MagmaNoTransStr, &i__1, &i__2,
&c_neg_one, A(rk, 0), &lda,
F(k, 0), &ldf,
&c_one, A(rk, k), &ione ); */
magma_zgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, dA(rk, 0), ldda,
dF(k, 0), lddf,
c_one, dA(rk, k), ione );
#endif
/*#if (defined(PRECISION_c) || defined(PRECISION_z))
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_Z_CNJG( *F(k,j) );
}
#endif*/
}
/* Generate elementary reflector H(k). */
magma_zlarfg_gpu( m-rk, dA(rk, k), dA(rk + 1, k), &tau[k], &vn1[k], &dAks[k]);
//Akk = *A(rk, k);
//*A(rk, k) = c_one;
//magma_zgetvector( 1, &dAks[k], 1, &Akk, 1 );
/* needed to avoid the race condition */
if (k == 0) magma_zsetvector( 1, &c_one, 1, dA(rk, k), 1 );
else magma_zcopymatrix( 1, 1, dA(offset, 0), 1, dA(rk, k), 1 );
/* Compute Kth column of F:
Compute F(K+1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) on the GPU */
if (k < n-1 || k > 0) magma_zgetvector( 1, &tau[k], 1, &tauk, 1 );
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Send the vector to the GPU */
//magma_zsetmatrix( i__1, 1, A(rk, k), lda, dA(rk,k), ldda );
/* Multiply on GPU */
// was CALL ZGEMV( 'Conjugate transpose', M-RK+1, N-K,
// TAU( K ), A( RK, K+1 ), LDA,
// A( RK, K ), 1,
// CZERO, F( K+1, K ), 1 )
//magma_zgetvector( 1, &tau[k], 1, &tauk, 1 );
magma_zgemv( MagmaConjTrans, m-rk, n-k-1,
tauk, dA( rk, k+1 ), ldda,
dA( rk, k ), 1,
c_zero, dF( k+1, k ), 1 );
//magma_zscal( m-rk, tau[k], F( k+1, k), 1 );
//magma_int_t i__3 = nb-k-1;
//magma_int_t i__4 = i__2 - i__3;
//magma_int_t i__5 = nb-k;
//magma_zgemv( MagmaConjTrans, i__1 - i__5, i__2 - i__3,
// tau[k], dA(rk +i__5, k+1+i__3), ldda,
// dA(rk +i__5, k ), ione,
// c_zero, dF(k+1+i__3, k ), ione );
//magma_zgetmatrix_async( i__2-i__3, 1,
// dF(k + 1 +i__3, k), i__2,
// F (k + 1 +i__3, k), i__2, stream );
//blasf77_zgemv( MagmaConjTransStr, &i__1, &i__3,
// &tau[k], A(rk, k+1), &lda,
// A(rk, k ), &ione,
// &c_zero, F(k+1, k ), &ione );
//magma_queue_sync( stream );
//blasf77_zgemv( MagmaConjTransStr, &i__5, &i__4,
// &tau[k], A(rk, k+1+i__3), &lda,
// A(rk, k ), &ione,
// &c_one, F(k+1+i__3, k ), &ione );
}
/* Padding F(1:K,K) with zeros.
for (j = 0; j <= k; ++j) {
magma_zsetvector( 1, &c_zero, 1, F(j, k), 1 );
}*/
/* Incremental updating of F:
F(1:N,K) := F(1:N,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K).
F(1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K)
:= tau(K)(A(RK:M,K+1:N)' - F(1:N,1:K-1)*A(RK:M,1:K-1)') A(RK:M,K)
so, F is (updated A)*V */
//if (k > 0 && k < n-1) {
if (k > 0) {
//magma_zgetvector( 1, &tau[k], 1, &tauk, 1 );
z__1 = MAGMA_Z_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_zgemv( MagmaConjTrans, i__1, i__2,
z__1, dA(offset+nb, 0), lda,
dA(offset+nb, k), ione,
c_zero, dauxv, ione );
i__1 = k;
magma_zgemv( MagmaNoTrans, n-k-1, i__1,
c_one, F(k+1,0), ldf,
dauxv, ione,
c_one, F(k+1,k), ione );
#else
i__1 = m - rk;
i__2 = k;
//blasf77_zgemv( MagmaConjTransStr, &i__1, &i__2,
// &z__1, A(rk, 0), &lda,
// A(rk, k), &ione,
// &c_zero, auxv, &ione );
magma_zgemv( MagmaConjTrans, i__1, i__2,
z__1, dA(rk, 0), ldda,
dA(rk, k), ione,
c_zero, dauxv, ione );
//i__1 = k;
//blasf77_zgemv( MagmaNoTransStr, &n, &i__1,
// &c_one, F(0,0), &ldf,
// auxv, &ione,
// &c_one, F(0,k), &ione );
/*magma_zgemv( MagmaNoTrans, n, i__1,
c_one, F(0,0), ldf,
auxv, ione,
c_one, F(0,k), ione ); */
/* I think we only need stricly lower-triangular part :) */
magma_zgemv( MagmaNoTrans, n-k-1, i__2,
c_one, dF(k+1,0), lddf,
dauxv, ione,
c_one, dF(k+1,k), ione );
#endif
}
/* Optimization: On the last iteration start sending F back to the GPU */
/* Update the current row of A:
A(RK,K+1:N) := A(RK,K+1:N) - A(RK,1:K)*F(K+1:N,1:K)'. */
if (k < n-1) {
i__1 = n - k - 1;
i__2 = k + 1;
//blasf77_zgemm( MagmaNoTransStr, MagmaConjTransStr, &ione, &i__1, &i__2,
// &c_neg_one, A(rk, 0 ), &lda,
// F(k+1,0 ), &ldf,
// &c_one, A(rk, k+1), &lda );
#ifdef RIGHT_UPDATE
/* right-looking update of rows, */
magma_zgemm( MagmaNoTrans, MagmaConjTrans, nb-k, i__1, ione,
c_neg_one, dA(rk, k ), ldda,
dF(k+1, k ), lddf,
c_one, dA(rk, k+1), ldda );
#else
/* left-looking update of rows, *
* since F=A'v with original A, so no right-looking */
magma_zgemm( MagmaNoTrans, MagmaConjTrans, ione, i__1, i__2,
c_neg_one, dA(rk, 0 ), ldda,
dF(k+1,0 ), lddf,
c_one, dA(rk, k+1), ldda );
#endif
}
/* Update partial column norms. */
if (rk < min(m, n+offset)-1 ) {
magmablas_dznrm2_row_check_adjust(n-k-1, tol3z, &vn1[k+1], &vn2[k+1], dA(rk,k+1), ldda, dlsticcs);
magma_device_sync();
#if defined(PRECISION_d) || defined(PRECISION_z)
magma_dgetvector( 1, &dlsticcs[0], 1, &lsticc, 1 );
#else
magma_sgetvector( 1, &dlsticcs[0], 1, &lsticc, 1 );
#endif
}
/*if (rk < lastrk) {
for (j = k + 1; j < n; ++j) {
if (vn1[j] != 0.) {
// NOTE: The following 4 lines follow from the analysis in
// Lapack Working Note 176.
temp = MAGMA_Z_ABS( *A(rk,j) ) / vn1[j];
temp = max( 0., ((1. + temp) * (1. - temp)) );
d__1 = vn1[j] / vn2[j];
temp2 = temp * (d__1 * d__1);
if (temp2 <= tol3z) {
vn2[j] = (double) lsticc;
lsticc = j;
} else {
vn1[j] *= magma_dsqrt(temp);
}
}
}
}*/
//*A(rk, k) = Akk;
//magma_zsetvector( 1, &Akk, 1, A(rk, k), 1 );
//magma_zswap( 1, &dAks[k], 1, A(rk, k), 1 );
++k;
}
magma_zcopymatrix( 1, k, dAks, 1, dA(offset, 0), ldda+1 );
// leave k as the last column done
--k;
*kb = k + 1;
rk = offset + *kb - 1;
/* Apply the block reflector to the rest of the matrix:
A(OFFSET+KB+1:M,KB+1:N) := A(OFFSET+KB+1:M,KB+1:N) - A(OFFSET+KB+1:M,1:KB)*F(KB+1:N,1:KB)' */
if (*kb < min(n, m - offset)) {
i__1 = m - rk - 1;
i__2 = n - *kb;
/* Send F to the GPU
magma_zsetmatrix( i__2, *kb,
F (*kb, 0), ldf,
dF(*kb, 0), i__2 ); */
magma_zgemm( MagmaNoTrans, MagmaConjTrans, i__1, i__2, *kb,
c_neg_one, dA(rk+1, 0 ), ldda,
dF(*kb, 0 ), lddf,
c_one, dA(rk+1, *kb), ldda );
}
/* Recomputation of difficult columns. */
if ( lsticc > 0 ) {
// printf( " -- recompute dnorms --\n" );
magmablas_dznrm2_check( m-rk-1, n-*kb, dA(rk+1,*kb), ldda,
&vn1[*kb], dlsticcs );
magma_dcopymatrix( n-*kb, 1, &vn1[*kb], *kb, &vn2[*kb], *kb );
/*while( lsticc > 0 ) {
itemp = (magma_int_t)(vn2[lsticc] >= 0. ? floor(vn2[lsticc] + .5) : -floor(.5 - vn2[lsticc]));
i__1 = m - rk - 1;
if (lsticc <= nb)
vn1[lsticc] = magma_cblas_dznrm2( i__1, A(rk+1,lsticc), ione );
else {
// Where is the data, CPU or GPU ?
double r1, r2;
r1 = magma_cblas_dznrm2( nb-k, A(rk+1,lsticc), ione );
r2 = magma_dznrm2(m-offset-nb, dA(offset + nb + 1, lsticc), ione);
vn1[lsticc] = magma_dsqrt(r1*r1+r2*r2);
}
// NOTE: The computation of VN1( LSTICC ) relies on the fact that
// SNRM2 does not fail on vectors with norm below the value of SQRT(DLAMCH('S'))
vn2[lsticc] = vn1[lsticc];
lsticc = itemp; */
}
magma_free(dAks);
magma_free(dlsticcs);
return MAGMA_SUCCESS;
} /* magma_zlaqps */
/**
@deprecated
Purpose
-------
ZLAQPS computes a step of QR factorization with column pivoting
of a complex M-by-N matrix A by using Blas-3. It tries to factorize
NB columns from A starting from the row OFFSET+1, and updates all
of the matrix with Blas-3 xGEMM.
In some cases, due to catastrophic cancellations, it cannot
factorize NB columns. Hence, the actual number of factorized
columns is returned in KB.
Block A(1:OFFSET,1:N) is accordingly pivoted, but not factorized.
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]
offset INTEGER
The number of rows of A that have been factorized in
previous steps.
@param[in]
nb INTEGER
The number of columns to factorize.
@param[out]
kb INTEGER
The number of columns actually factorized.
@param[in,out]
dA COMPLEX_16 array, dimension (LDDA,N), on the GPU.
On entry, the M-by-N matrix A.
On exit, block A(OFFSET+1:M,1:KB) is the triangular
factor obtained and block A(1:OFFSET,1:N) has been
accordingly pivoted, but no factorized.
The rest of the matrix, block A(OFFSET+1:M,KB+1:N) has
been updated.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[in,out]
jpvt INTEGER array, dimension (N)
JPVT(I) = K <==> Column K of the full matrix A has been
permuted into position I in AP.
@param[out]
tau COMPLEX_16 array, dimension (KB)
The scalar factors of the elementary reflectors.
@param[in,out]
vn1 DOUBLE PRECISION array, dimension (N)
The vector with the partial column norms.
@param[in,out]
vn2 DOUBLE PRECISION array, dimension (N)
The vector with the exact column norms.
@param[in,out]
dauxv COMPLEX_16 array, dimension (NB), on the GPU
Auxiliary vector.
@param[in,out]
dF COMPLEX_16 array, dimension (LDDF,NB), on the GPU
Matrix F' = L*Y'*A.
@param[in]
lddf INTEGER
The leading dimension of the array F. LDDF >= max(1,N).
@ingroup magma_zgeqp3_aux
********************************************************************/
extern "C" magma_int_t
magma_zlaqps_gpu(
magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magma_int_t *jpvt, magmaDoubleComplex *tau,
double *vn1, double *vn2,
magmaDoubleComplex_ptr dauxv,
magmaDoubleComplex_ptr dF, magma_int_t lddf)
{
#define dA(i, j) (dA + (i) + (j)*(ldda))
#define dF(i, j) (dF + (i) + (j)*(lddf))
magmaDoubleComplex c_zero = MAGMA_Z_MAKE( 0.,0.);
magmaDoubleComplex c_one = MAGMA_Z_MAKE( 1.,0.);
magmaDoubleComplex c_neg_one = MAGMA_Z_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
magmaDoubleComplex z__1;
magma_int_t k, rk;
magmaDoubleComplex_ptr dAks;
magmaDoubleComplex tauk = MAGMA_Z_ZERO;
magma_int_t pvt;
double tol3z;
magma_int_t itemp;
double lsticc;
magmaDouble_ptr dlsticcs;
magma_dmalloc( &dlsticcs, 1+256*(n+255)/256 );
tol3z = magma_dsqrt( lapackf77_dlamch("Epsilon"));
lsticc = 0;
k = 0;
magma_zmalloc( &dAks, nb );
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// subtract 1 from Fortran/CUBLAS idamax; pvt, k are 0-based.
pvt = k + magma_idamax( n-k, &vn1[k], ione, queue ) - 1;
if (pvt != k) {
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
magmablas_zswap( m, dA(0, pvt), ione, dA(0, k), ione, queue );
magmablas_zswap( i__1, dF(pvt, 0), lddf, dF(k, 0), lddf, queue );
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
magma_dswap( 2, &vn1[pvt], n+offset, &vn1[k], n+offset, queue );
}
/* Apply previous Householder reflectors to column K:
A(RK:M,K) := A(RK:M,K) - A(RK:M,1:K-1)*F(K,1:K-1)'.
Optimization: multiply with beta=0; wait for vector and subtract */
if (k > 0) {
//#define RIGHT_UPDATE
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_zgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, A(offset+nb, 0), lda,
F(k, 0), ldf,
c_one, A(offset+nb, k), ione, queue );
#else
i__1 = m - rk;
i__2 = k;
magma_zgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, dA(rk, 0), ldda,
dF(k, 0), lddf,
c_one, dA(rk, k), ione, queue );
#endif
}
/* Generate elementary reflector H(k). */
magma_zlarfg_gpu( m-rk, dA(rk, k), dA(rk + 1, k), &tau[k], &vn1[k], &dAks[k], queue );
/* needed to avoid the race condition */
if (k == 0) magma_zsetvector( 1, &c_one, 1, dA(rk, k), 1, queue );
else magma_zcopymatrix( 1, 1, dA(offset, 0), 1, dA(rk, k), 1, queue );
/* Compute Kth column of F:
Compute F(K+1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) on the GPU */
if (k < n-1 || k > 0) magma_zgetvector( 1, &tau[k], 1, &tauk, 1, queue );
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Multiply on GPU */
magma_zgemv( MagmaConjTrans, m-rk, n-k-1,
tauk, dA( rk, k+1 ), ldda,
dA( rk, k ), 1,
c_zero, dF( k+1, k ), 1, queue );
}
/* Incremental updating of F:
F(1:N,K) := F(1:N,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K).
F(1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K)
:= tau(K)(A(RK:M,K+1:N)' - F(1:N,1:K-1)*A(RK:M,1:K-1)') A(RK:M,K)
so, F is (updated A)*V */
if (k > 0) {
z__1 = MAGMA_Z_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_zgemv( MagmaConjTrans, i__1, i__2,
z__1, dA(offset+nb, 0), lda,
dA(offset+nb, k), ione,
c_zero, dauxv, ione, queue );
i__1 = k;
magma_zgemv( MagmaNoTrans, n-k-1, i__1,
c_one, F(k+1,0), ldf,
dauxv, ione,
c_one, F(k+1,k), ione, queue );
#else
i__1 = m - rk;
i__2 = k;
magma_zgemv( MagmaConjTrans, i__1, i__2,
z__1, dA(rk, 0), ldda,
dA(rk, k), ione,
c_zero, dauxv, ione, queue );
/* I think we only need stricly lower-triangular part :) */
magma_zgemv( MagmaNoTrans, n-k-1, i__2,
c_one, dF(k+1,0), lddf,
dauxv, ione,
c_one, dF(k+1,k), ione, queue );
#endif
}
/* Optimization: On the last iteration start sending F back to the GPU */
/* Update the current row of A:
A(RK,K+1:N) := A(RK,K+1:N) - A(RK,1:K)*F(K+1:N,1:K)'. */
if (k < n-1) {
i__1 = n - k - 1;
i__2 = k + 1;
#ifdef RIGHT_UPDATE
/* right-looking update of rows, */
magma_zgemm( MagmaNoTrans, MagmaConjTrans, nb-k, i__1, ione,
c_neg_one, dA(rk, k ), ldda,
dF(k+1, k ), lddf,
c_one, dA(rk, k+1), ldda, queue );
#else
/* left-looking update of rows, *
* since F=A'v with original A, so no right-looking */
magma_zgemm( MagmaNoTrans, MagmaConjTrans, ione, i__1, i__2,
c_neg_one, dA(rk, 0 ), ldda,
dF(k+1,0 ), lddf,
c_one, dA(rk, k+1), ldda, queue );
#endif
}
/* Update partial column norms. */
if (rk < min(m, n+offset)-1 ) {
magmablas_dznrm2_row_check_adjust( n-k-1, tol3z, &vn1[k+1], &vn2[k+1],
dA(rk,k+1), ldda, dlsticcs, queue );
//magma_device_sync();
magma_dgetvector( 1, &dlsticcs[0], 1, &lsticc, 1, queue );
}
++k;
}
magma_zcopymatrix( 1, k, dAks, 1, dA(offset, 0), ldda+1, queue );
// leave k as the last column done
--k;
*kb = k + 1;
rk = offset + *kb - 1;
/* Apply the block reflector to the rest of the matrix:
A(OFFSET+KB+1:M,KB+1:N) := A(OFFSET+KB+1:M,KB+1:N) - A(OFFSET+KB+1:M,1:KB)*F(KB+1:N,1:KB)' */
if (*kb < min(n, m - offset)) {
i__1 = m - rk - 1;
i__2 = n - *kb;
magma_zgemm( MagmaNoTrans, MagmaConjTrans, i__1, i__2, *kb,
c_neg_one, dA(rk+1, 0 ), ldda,
dF(*kb, 0 ), lddf,
c_one, dA(rk+1, *kb), ldda, queue );
}
/* Recomputation of difficult columns. */
if ( lsticc > 0 ) {
// printf( " -- recompute dnorms --\n" );
magmablas_dznrm2_check( m-rk-1, n-*kb, dA(rk+1,*kb), ldda,
&vn1[*kb], dlsticcs, queue );
magma_dcopymatrix( n-*kb, 1, &vn1[*kb], *kb, &vn2[*kb], *kb, queue );
}
magma_free( dAks );
magma_free( dlsticcs );
magma_queue_destroy( queue );
return MAGMA_SUCCESS;
} /* magma_zlaqps */
/* ////////////////////////////////////////////////////////////////////////////
-- 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;
}
